CUDA shared array not getting values?

Question

I am trying to implement simple parallel reduction. I am using the code from the CUDA SDK. But somehow there is a problem in my kernel as the shared array is not getting values of the global array and its all zeroes.

extern __ shared __ float4 sdata[];

// each thread loads one element from global to shared mem

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

sdata[tid] = dev_src[i];

__syncthreads();

// do reduction in shared mem

for(unsigned int s = 1; s < blockDim.x; s *= 2) {
    if(tid % (2*s) == 0){
        sdata[tid] += sdata[tid + s];
    }
    __syncthreads();
}

// write result for this block to global mem
if(tid == 0)
    out[blockIdx.x] = sdata[0];

Edit:

Ok I got it working by removing the extern keyword and making the shared array a constant size like 512. I am in good shape now. Maybe someone can explain why it was not working with the extern keyword.

so dev_src has the correct values but sdata somehow is not getting dev_src values — elasolova
– elasolova, Commented Feb 12, 2012 at 0:51

paleonix · Accepted Answer · 2023-08-09 16:45:44Z

2

I think I know why this is happening as I have faced this before. How are you launching the kernel?

Remember in the launch kernel<<<blocks,threads,sharedMemory>>> the sharedMemory should be the size of the shared memory in bytes. So, if you are declaring for 512 elements, the third parameter should be 512 * sizeof(float4). I think you are just calling as below, which is wrong

kernel<<<blocks,threads,512>>>   // this is wrong

edited Aug 9, 2023 at 16:45

paleonix

3,3255 gold badges20 silver badges42 bronze badges

answered Feb 12, 2012 at 6:00

Programmer

6,78331 gold badges82 silver badges130 bronze badges

Sign up to request clarification or add additional context in comments.

1 Comment

harrism Over a year ago

To clarify this answer, since it is mostly correct: if you specify a shared variable as `extern shared" then you must specify the amount of shared memory (in bytes) to allocate at kernel launch using the third parameter of the execution configuration (the <<<>>>), as Programmer describes.

Collectives™ on Stack Overflow

CUDA shared array not getting values?

1 Answer 1

1 Comment

Your Answer

Hot Network Questions

Collectives™ on Stack Overflow

1 Answer 1

1 Comment

Your Answer

Sign up or log in

Post as a guest

Related