bitwise atomic operator atomicOr data type size_t

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello.

I need a 64bit unsigned integer data type atomicOr, especially size_t data type.

but CUDA supports only unsigned long long int data type.

I can’t Be compatible between size_t and unsigned long long type.

So I want to use like this code:

__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                                          (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;
    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed, 
                        __double_as_longlong(val + 
                        __longlong_as_double(assumed)));
    } while (assumed != old);
    return __longlong_as_double(old);
}

CUDA API hasn’t double data type atomicAdd.
Can use size_t atomicOr like that?

2

Add)

and How to use atomicOr unsigned long long int example?
unsigned int(32bit) and int(32bit) are compiled, but unsigned long long int(64bit) has occurred… I don’t know the reason.

3

What is this supposed to mean? On all platforms supported by CUDA, both ‘size_t’ and ‘unsigned long long int’ are unsigned 64-bit integer types, and the former might even just be a typedef of the latter. So casting in either direction should be perfectly fine.

4

but when I use the size_t in atomicOr, it occurred this error message
error: no instance of overloaded function “atomicOr” matches the argument list argument types are: (size_t *, size_t)

and when I use the unsigned long long in atomicOr, it occurred this error message
error: no instance of overloaded function “atomicOr” matches the argument list argument types are: (unsigned long long *, unsigned long long)

but, when I use the int or unsigned int in atomicOr, there was no any compile error.

5

That’s a different issue than the one I understood you were raising in your earlier post. Not all functions are necessarily overloaded for all data types. There may be limitations due to GPU architecture. From the CUDA Programming Guide:

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicor

You would want to specify the architecture of your GPU when building the code. The compiler likely defaults to a build target of compute capability 3.0.

6

I already read that.
my GPU CC is over 6.0 (titan xp)

I can understand your reply little bit.
so, I should give the compiler setting parameter.
I confuse the parameter.
SM_60 something like that? right?

7

A simple build command setting the target architecture explicitly would be

nvcc -arch=sm_61 -o app.exe app.cu

Please consult the compiler documentation on how to specify target architectures.

8

thank you so much!!!

have an awesome day!

9

I am really sorry…

I found the page
https://stackoverflow.com/questions/35656294/cuda-how-to-use-arch-and-code-and-sm-vs-compute

but I can’t understand how to use the gencode and arch…

and I check the NVIDIA CUDA programming guide, but I can’t find the arch section.

Could you give me the information?

Thank you for your warm hand

10

Latest NVCC documentation is here:

[url]https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html[/url]