Error on using atomic operations in CUDA C

Question

I'm working on CUDA as a beginner and am trying to execute a pre written code the compile gives error for every atomic operation that the code contains... for example

__global__ void MarkEdgesUV(unsigned int *d_edge_flag, unsigned long long int *d_appended_uvw, unsigned int *d_size, int no_of_edges)

{

    unsigned int tid = blockIdx.x*MAX_THREADS_PER_BLOCK + threadIdx.x;

    if(tid<no_of_edges)

    {

    if(tid>0)

        {

        unsigned long long int test = INF;

        test = test << NO_OF_BITS_MOVED_FOR_VERTEX_IDS;

        test |=INF;

        unsigned long long int test1 = d_appended_uvw[tid]>>(64-(NO_OF_BITS_MOVED_FOR_VERTEX_IDS+NO_OF_BITS_MOVED_FOR_VERTEX_IDS));

        unsigned long long int test2 = d_appended_uvw[tid-1]>>(64-(NO_OF_BITS_MOVED_FOR_VERTEX_IDS+NO_OF_BITS_MOVED_FOR_VERTEX_IDS));

        if(test1>test2)

            d_edge_flag[tid]=1;

        if(test1 == test)
            * atomicMin(d_size,tid); //also to know the last element in the array, i.e. the size of new edge list

        }

    else

        d_edge_flag[tid]=1;

    }

}

gives the error: error: identifier "atomicMin" is undefined this happen to be a very reliable code...also i check out, the usage of atomics seems correct....please explain why the error has occured?

Ok I think your problem is following.. don't show any code, compile errors or anything else! How should anyone help you?! — hubs, Mar 05 '14 at 09:53
sorry for the inconvience....the 1st submission was accidental — soodankit, Mar 05 '14 at 09:58
For what compute architectur do you compute? Take a look at [this SO question](http://stackoverflow.com/questions/11532728/error-identifier-atomicadd-is-undefined-under-visual-studio-2010-cuda-4-2). Atomic operations are available since cc2.0. — hubs, Mar 05 '14 at 10:18
i'm using tesla C2075....it's cc2.0 this code originally was compiled on tesla S1070 cc 1.3 (i suppose) — soodankit, Mar 05 '14 at 11:13

Sigi · Accepted Answer · 2014-03-06T09:26:00.220

0

I guess you are compiling with nvcc only (defaulting to sm_10), without specifying the minimal needed compute capability. In fact atomicMin() on 32 bits global memory has been introduced in devices with CC1.1 (compute capability 1.1), as you can see in this table.

Try compiling with

nvcc -arch sm_11 ...

This will let the atomicMin() funcion be recognized. However best you compile for the actual architecture of your GPU.

EDITED: being it Tesla C2075, :

nvcc -arch sm_20 ...

should work as well.

EDIT: as required by the OP, adding a reference:

You can find a very detailed explanation of the meaning of -arch and -code options here.

In particular it's reported explicitly that both -arch and -code options can be omitted.

nvcc x.cu

is short hand for

nvcc x.cu -arch=compute_10 -code=sm_10,compute_10

edited Mar 06 '14 at 09:26

answered Mar 05 '14 at 10:46

Sigi

4,826
1
19
23

1

@Sigismondo your compile commands are not correct. If you specify the `-code` option, the `-arch` option [must be](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#shorthand-2) a *virtual architecture only*. You can fix this either simply by specifying `nvcc -arch sm_20 ...` or else `nvcc -arch compute_20 -code sm_20 ...` – Robert Crovella Mar 05 '14 at 12:45
Thanks...-arch sm_20 worked can you please give me some link where I can furthur read about such issue...? – soodankit Mar 06 '14 at 08:38
I've edited the answer, you can find there the link to the documentation for the nvcc options. Oh yes Robert Crovella too reported the link! – Sigi Mar 06 '14 at 09:27

Error on using atomic operations in CUDA C

1 Answers1