6

同じブロック内の他のスレッドとそのメンバーの一部を共有するCUDAコードでクラスをインスタンス化したいと思います。

ただし、次のコードをコンパイルしようとすると、次のエラーが発生します。»属性「共有」はここには適用されません«（nvccバージョン4.2）。

class SharedSomething {

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

__global__ void run() {

    SharedSomething something;
}

その背後にある理論的根拠は何ですか？目的の動作（1つのブロックにまたがるクラスの共有メンバー）を実現するための回避策はありますか？

4

2 に答える 2

7

としてマークされたオブジェクトは、__shared__スレッド ブロックごとに専用の共有メモリに常駐します。サイズに制限があり、寿命はスレッド ブロックと同じです。

これが、クラス メンバーを共有として宣言できない理由です。それらの有効期間は、クラス インスタンスではなく、スレッド ブロックによって管理されます。おそらくstaticクラス メンバーは共有されている可能性がありますが、チェックしていません。

詳細については、 CUDA プログラミング ガイドのセクション B.2.3 を参照してください。

于 2012-10-03T11:41:15.477 に答える
6

ロスト氏は、制限の背後にある理論的根拠を説明しました。質問の 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();
}

警告: コードはブラウザで書かれ、検証もテストもされていません (これは些細な例ですが、この概念は実際のコードにも適用されます。私自身もこの手法を使用しました)。

于 2012-10-03T11:54:30.063 に答える