Operaciones atómicas en CUDA? ¿Qué archivo de encabezado incluir?

Question

Para usar operaciones atómicas en CUDA, ¿es necesario incluir algún archivo de encabezado CUDA? La guía de programación de CUDA parece estar cerrada en esto.

El código glmax.cu que figura a continuación me da el siguiente error de compilación.

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". 

Aquí está el código. Básicamente se trata de calcular el valor máximo de una matriz en la GPU utilizando la operación atómica atomicMax. Como soy nuevo en CUDA, estoy seguro de que soy un código bastante ingenuo, pero escribí esto para ayudarme a comprender las operaciones atómicas.

#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

No creo que sea necesario un #include. Las operaciones atómicas no están disponibles en los dispositivos "Compute Capability" 1.0 (sm_10), que es lo que está pidiendo nvcc para compilar (de forma predeterminada).

Para utilizar atomicMax en su código, especificar al menos -arch=sm_11 en la línea de comandos:

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

Para referencia futura, puede consultar el Apéndice F de la Guía CUDA C de programación para obtener información sobre qué operaciones atómica están disponibles en plataformas de una Capacidad de cálculo particular.

Por supuesto, necesitará una GPU compatible con sm_11 para ejecutar el código. Mi impresión es que estos son comunes ahora.