Why can't member variables be shared?

Question

I would like to instantiate a class in CUDA code, that shares some of its members with other threads in the same block.

However, when trying to compile the following code, I get the error: »attribute "shared" does not apply here« (nvcc version 4.2).

class SharedSomething {

public:
    __shared__ int i; // this is not allowed
};

__global__ void run() {

    SharedSomething something;
}

What is the rationale behind that? Is there a work-around to achieve the desired behavior (shared members of a class across one block)?

Answer 1

Objects marked as __shared__ reside in shared memory that is dedicated per thread block. It has limited size and has the same lifetime as thread block.

So this is the reason why you cannot declare class members as shared - their lifetime is not managed by class instance, but by thread block. Possibly static class members could be shared, but didn't check it.

See CUDA Programming Guide for details, section B.2.3.