为什么不能共享成员变量？

Question

我想在CUDA代码中实例化一个类,它在同一个块中与其他线程共享一些成员.

但是,在尝试编译以下代码时,我收到错误:»属性"shared"在这里不适用«(nvcc版本4.2).

class SharedSomething {

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

__global__ void run() {

    SharedSomething something;
}

这背后的理由是什么？是否有解决方案来实现所需的行为(跨越一个块的类的共享成员)？

Answer 1

标记为__shared__驻留在共享内存中的对象,每个线程块专用.它的尺寸有限,与螺纹块的寿命相同.

所以这就是为什么你不能将类成员声明为共享的原因 - 它们的生命周期不是由类实例管理,而是由线程块管理.可能static会共享类成员,但没有检查它.

有关详细信息,请参阅CUDA编程指南,B.2.3节.

Answer 2

罗斯特解释了限制背后的基本原理.要回答问题的第二部分,一个简单的解决方法是让内核声明共享内存,并初始化类所拥有的指针,例如在类构造函数中.例.

class Foo 
{
public:
  __device__
  Foo(int *sPtr) : sharedPointer(sPtr, gPtr) {
    sharedPointer[threadIdx.x] = gPtr[blockIdx.x * blockDim.x + threadIdx.x];
    __syncthreads();
  }

  __device__
  void useSharedData() { printf("my data: %f\n", sharedPointer[threadIdx.x]); }

private:
  int *sharedPointer;
};

__global__ void example(int *gData) 
{
  __shared__ int sData[BLOCKDIM];

  Foo f(sData, gData);

  f.useSharedData();
}

警告:用浏览器编写的代码,未经验证,未经测试(这是一个简单的例子,但概念扩展到实际代码 - 我自己使用了这种技术).