6

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

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

4

2 回答 2

7

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

所以这就是你不能将类成员声明为共享的原因——它们的生命周期不是由类实例管理，而是由线程块管理。可能static班级成员可以共享，但没有检查。

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

于 2012-10-03T11:41:15.477 回答
6

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: %f\n", sharedPointer[threadIdx.x]); }

private:
  int *sharedPointer;
};

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

  Foo f(sData, gData);

  f.useSharedData();
}

警告：在浏览器中编写的代码，未经验证，未经测试（这是一个微不足道的例子，但这个概念扩展到实际代码——我自己也使用过这种技术）。

于 2012-10-03T11:54:30.063 回答