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

Why can't member variables be shared?

本文关键字:变量 成员 共享 不能 为什么      更新时间:2023-10-16

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

但是,在尝试编译以下代码时,我收到错误:

attribute "shared" does not apply here

(NVCC 版本 4.2)。

class SharedSomething {
public:
    __shared__ int i; // this is not allowed
};
__global__ void run() {
    SharedSomething something;
}

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

标记为

__shared__的对象驻留在每个线程块专用的共享内存中。它的大小有限,并且具有与线程块相同的生存期。

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

有关详细信息,请参阅 CUDA 编程指南。

Rost解释了限制背后的理由。为了回答问题的第二部分,一个简单的解决方法是让内核声明共享内存,并初始化一个指向类拥有的它的指针,例如在类构造函数中。例。

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: %fn", sharedPointer[threadIdx.x]); }
private:
  int *sharedPointer;
};
__global__ void example(int *gData) 
{
  __shared__ int sData[BLOCKDIM];
  Foo f(sData, gData);
  f.useSharedData();
}

警告:用浏览器编写的代码,未经验证,未经测试(这是一个微不足道的例子,但这个概念延伸到真正的代码 - 我自己也使用过这种技术)。