Atomic Operation fehlgeschlagen in CUDA

Question

Da die Rechenfähigkeit 2.1 ist, unterstützen die atomicAdd und atomicMax Operationen keine doppelte Genauigkeit, dann definiere ich beide Funktionen auf der Grundlage einiger Antworten auf Stack-Überlauf.

Es ist seltsam, dass die atomicAdd Funktion funktioniert gut, aber die atomicMax funktioniert nicht, hier ist mein Code.

Der Test meines Codes ist, Zufallszahl auf jedem Block zu generieren, und dann die Zufallszahlen auf jedem Block summieren, wir haben Blocksumme, ich möchte die atomicAdd und atomicMax auf die Blocksumme testen.

#include <iostream> 
#include <curand.h> 
#include <curand_kernel.h> 
#include <stdio.h> 
#include <stdlib.h> 


#define num_of_blocks 2 
#define threads_per_block 2 
#define tot_threads 4 


__device__ double gsum[num_of_blocks]; 

__device__ double dev_sum; 

__device__ double dev_max; 

// set seed for random number generator 
__global__ void initcuRand(curandState* globalState, unsigned long seed){ 
    int idx = threadIdx.x + blockIdx.x * blockDim.x; 
    curand_init(seed, idx, 0, &globalState[idx]); 
} 

// atomiMax for double 
__device__ double atomicMax_d(double* address, double val) 
{ 
    unsigned long long int* address_as_i = (unsigned long long int*)address; 
    unsigned long long int old = *address_as_i, assumed; 
    do { 
     assumed = old; 
     old = ::atomicCAS(address_as_i, assumed, __double_as_longlong(::fmax(val, __longlong_as_double(assumed)))); 
    } while (assumed != old); 
    return __longlong_as_double(old); 
} 

// atomicAdd for double 
__device__ double atomicAdd_d(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); 
} 

__global__ void kernel(curandState *globalState){ 
    // global id 
    int gidx = threadIdx.x + blockIdx.x * blockDim.x; 
    // local id 
    int lidx = threadIdx.x; 

    // creat shared memory to store seeds 
    __shared__ curandState localState[tot_threads]; 

    __shared__ double srandnum[threads_per_block]; 

    // copy global seed to local 
    localState[lidx] = globalState[gidx]; 

    //synchronize the local threads writing to the local memory cache 
    __syncthreads(); 

    // generate random number from normal distribution in shared memory 
    srandnum[lidx] = curand_normal(&localState[lidx]); 
    __syncthreads(); 

    if(lidx == 0){srandnum[lidx] += srandnum[lidx + 1];} // sum of each block 
    if(lidx == 0){gsum[blockIdx.x] = srandnum[lidx];}  // copy the sums back to global memory 

    __threadfence(); 

    if(gidx < num_of_blocks){ 
     atomicAdd_d(&dev_sum, gsum[gidx]); 
    } 

    if(gidx < num_of_blocks){ 
     atomicMax_d(&dev_max, gsum[gidx]); 
    } 

    if(gidx == 0){ 
     printf("Sum is: %lf\n", dev_sum); 
    } 

    if(gidx == 1){ 
     printf("Max is: %lf\n", dev_max); 
    } 
} 


int main(){ 
    // set seed on device 
    curandState *globalState; 
    cudaMalloc((void**)&globalState, tot_threads*sizeof(curandState)); 
    initcuRand<<<num_of_blocks, threads_per_block>>>(globalState, 1); 

    // launch kernel 
    kernel<<<num_of_blocks, threads_per_block>>>(globalState); 
    double randnum[num_of_blocks]; 

    cudaMemcpyFromSymbol(randnum, gsum, num_of_blocks*sizeof(double), 0, cudaMemcpyDeviceToHost); 

    std::cout << "Sum of each block:\n"; 
    for (int i = 0; i < num_of_blocks; ++i){ 
     std::cout << randnum[i] << std::endl; 
    } 

    cudaFree(globalState); 
    return 0; 
} 

Das Ergebnis, das ich bekommen ist

Sum is: -0.898329 
Max is: 0.000000 
Sum of each block: 
-0.0152994 
-0.88303 

Aus dem Ergebnis, weiß ich, dass die atomicAdd Funktion funktioniert, aber die atomicMax Funktion funktioniert nicht, ich habe keine Ahnung davon. Danke im Voraus.

Answer 1

Sie initialisieren nie dev_max oder dev_sum. Sie können diese Arten von atomaren Operationen nicht sinnvoll ausführen, wenn sie nicht mit einem bekannten Wert beginnen.

Versuchen etwas Ähnliches statt:

__device__ double dev_sum = 0.0; 

__device__ double dev_max = -1e99; 

und ich denke, Sie werden mit den Ergebnissen glücklicher sein.