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

Question

我想在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;
}

这背后的原理是什么？有没有方法可以实现期望的行为（在一个块内共享类的成员）？

- user1716882

Answer 1

__shared__标记的对象位于每个线程块专用的共享内存中。它具有有限的大小，并且与线程块具有相同的生命周期。

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

有关详细信息，请参见CUDA Programming Guide。

- Rost

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();
}

注意：这段代码是在浏览器中编写的，未经验证、未经测试（虽然这只是一个简单的例子，但这个概念可以扩展到实际的代码——我自己也使用过这种技术）。

- harrism

---

感谢解决方法。可以通过在Foo中声明一个内部类Shared来使其更通用，该内部类保存所有共享数据。调用代码实例化共享的Foo::Shared并将其传递给Foo的构造函数。这样，如果Foo::Shared发生更改，调用代码无需更改。 - user1716882