invalid device symbol cudaMemcpyFromSymbol CUDA

Question

I want to calculate the sum of all elements of an array in CUDA. I came up with this code. It compiles without any error. But the result is always zero. I've got the invalid device symbol from cudaMemcpyFromSymbol. I cannot use any libraries like Thrust or Cublas.

#define TRIALS_PER_THREAD 4096
#define NUM_BLOCKS 256
#define NUM_THREADS 256
double *dev;
__device__ volatile double pi_gpu = 0;

__global__ void ArraySum(double *array)

{
unsigned int tid = threadIdx.x + blockDim.x * blockIdx.x;
pi_gpu = pi_gpu + array[tid];
__syncthreads();
}

int main (int argc, char *argv[]) {
cudaMalloc((void **) &dev, NUM_BLOCKS * NUM_THREADS * sizeof(double));
    double pi_gpu_h;

ArraySum<<<NUM_BLOCKS, NUM_THREADS>>>(dev);
cudaDeviceSynchronize();
cudaError err = cudaMemcpyFromSymbol(&pi_gpu_h, &pi_gpu, sizeof(double), cudaMemcpyDeviceToHost);
if( cudaSuccess != err )
{
    fprintf( stderr, "cudaMemcpyFromSymbolfailed : %s\n", cudaGetErrorString( err ) );
    exit( -1 );
}
return pi_gpu_h; // this is always zero!!!
}

Answer 1

Your code is not thread safe. Writing to global variable from multiple threads is not safe at all. This example of how reduction kernel may be:

//Untested code
global_void plus_reduce(int *input, int N, int *total){
    int tid = threadIdx.x;
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    // Each block loads its elements into shared memory
    _shared_ int x[blocksize];
    x[tid] = (i<N) ? input[i] : 0; // last block may pad with 0’s
    _syncthreads();
    // Build summation tree over elements.
    for(int s=blockDim.x/2; s>0; s=s/2){
        if(tid < s) x[tid] += x[tid + s];
    _syncthreads();
}
// Thread 0 adds the partial sum to the total sum
if( tid == 0 ) 
    atomicAdd(total, x[tid]);    
}

Source

Answer 2

The symbol argument in the copy from symbol call is incorrect. It should look like this:

cudaMemcpyFromSymbol(&pi_gpu_h, pi_gpu, sizeof(double), 0, cudaMemcpyDeviceToHost)