Why can't I use a single thread to initialize shared memory?

Question

This seems like it should be simple, but I can't find any references, so I'm asking here.

I have the following CUDA kernel, which I am launching in a grid of 2D thread blocks:

__global__ void kernel(){

    if (threadIdx.x == 0 && threadIdx.y == 0) {
        __shared__ int test = 100;
    }
    __syncthreads();

    // Do more stuff
}

When I try to compile, I get the error

initializer not allowed for shared variable

What am I doing wrong? It seems to me like I have just one thread doing the initialization...

Thanks!

Answer 1

Do this instead:

__global__ void kernel(){
    __shared__ int test;
    if (threadIdx.x == 0 && threadIdx.y == 0) {
        test = 100;
    }
    __syncthreads();

    // Do more stuff
}

The declaration of the __shared___ variable must be separate from code that manipulates it.