Atomic operations in CUDA? Which header file to include?

Question

To use atomic operations in CUDA, is it necessary to include a CUDA file in the header? The CUDA programming guide seems to be very concise.

The glmax.cu code below gives me the following compilation error.

gaurish108 MyPractice: nvcc glmax.cu -o glmax glmax.cu(11): error: identifier "atomicMax" is undefined 1 error detected in the compilation of "/tmp/tmpxft_000010fa_00000000-4_glmax.cpp1.ii". 

Here is the code. This is basically calculating the maximum value of an array on a GPU using the atomicMax atomic operation. Since I'm new to CUDA, I'm sure this is very naive code, but I wrote this to help myself understand atomic operations.

 #include<stdio.h> #include<stdlib.h> #include<math.h> __global__ void global_max(int* values, int* gl_max) { int i=threadIdx.x + blockDim.x * blockIdx.x; int val=values[i]; atomicMax(gl_max,val); } int main(void) { int array_size=5; int num_bytes=array_size*sizeof(int); int *device_array=0; int *host_array=0; int *device_max=0; int *host_max=0; //Allocate memory on the host host_array=(int*)malloc(num_bytes); //Allocate memory on the device cudaMalloc((void**)&device_array,num_bytes); cudaMalloc((void**)&device_max,sizeof(int)); //If either memory allocation failed, report an error message if(host_array == 0 || device_array == 0) { printf("couldn't allocate memory\n"); return 1; } //Assign a random integer in the interval [0,25] to host_array members for(int i=0;i<array_size;++i) { *(host_array+i)=rand()%26; } //Print the host array members printf("Host Array\n"); for(int i=0;i<array_size;++i) { printf("%d ",*(host_array+i)); } printf("\n"); //Copy array from host to device. cudaMemcpy(device_array,host_array,num_bytes,cudaMemcpyHostToDevice); //Configure and launch the kernel which calculates the maximum element in the device array. int grid_size=1;//Only 1 block of threads is used int block_size=5;//One block contains only 5 threads //Device array passed to the kernel as data. global_max<<<grid_size,block_size>>>(device_array,device_max); //Transfer the maximum value so calculated into the CPU and print it cudaMemcpy(host_max,device_max,sizeof(int),cudaMemcpyDeviceToHost); printf("\nMaximum value is %d\n",*host_max); // deallocate memory free(host_array); cudaFree(device_array); cudaFree(device_max); return 0; } 

Answer 1

I do not think that #include needed. Atomic operations are not available on Compute Capability 1.0 ( sm_10 ) sm_10 that you request nvcc to compile (by default).

To use atomicMax in your code, specify at least -arch=sm_11 on the command line:

 $nvcc -arch=sm_11 glmax.cu -o glmax 

Further, refer to Appendix F of the CUDA C Programming Guide for information on which atomic operations are available on specific computing capability platforms.

Of course, to run the code you will need a sm_11 GPU. My impression is that they are prevalent by now.