Operazioni atomiche in CUDA? Quale file di intestazione includere?

Question

Per utilizzare le operazioni atomiche in CUDA, è necessario includere alcuni file di intestazione CUDA? La guida alla programmazione CUDA sembra essere stretta su questo.

Il codice glmax.cu fornito di seguito mi dà il seguente errore di compilazione.

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

Ecco il codice. In pratica sta calcolando il valore massimo di un array sulla GPU utilizzando l'operazione atomica atomicMax. Dato che sono nuovo di CUDA, questo è sicuramente un codice piuttosto ingenuo, ma l'ho scritto per aiutarmi a capire le operazioni atomiche.

#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

non credo un #include è necessario. Le operazioni atomiche non sono disponibili sui dispositivi "Compute Capability" 1.0 (sm_10), che è quello che stai chiedendo di compilare per nvcc (per impostazione predefinita).

Per utilizzare atomicMax nel codice, specificare almeno -arch=sm_11 sulla riga di comando:

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

Per riferimento futuro, è possibile consultare l'appendice F della Guida CUDA C programmazione per informazioni su quali operazioni atomiche sono disponibili su piattaforme di una particolare capacità di calcolo.

Ovviamente, per eseguire il codice è necessaria una GPU compatibile con lo sm_11. La mia impressione è che questi sono ormai comuni.