メンバー変数を共有できないのはなぜですか？

Question

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

しかし、次のコードをコンパイルしようとすると、エラーが表示されます。»属性「shared」はここでは適用されません。«（nvcc version 4.2）。

class SharedSomething { 

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

__global__ void run() { 

    SharedSomething something; 
} 

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

Answer 1

ロストは制限の根拠を説明しました。 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(); 
} 

警告：ブラウザで書かれたコードは、未検証、テストされていない（そして、それは簡単な例ですが、コンセプトは、実際のコード—に、私はこのテクニック自分自身を使用している拡張します）。

Answer 2

__shared__とマークされたオブジェクトは、スレッドブロックごとに専用の共有メモリに存在します。それはサイズが限られており、スレッドブロックと同じ寿命を持っています。

これは、クラスメンバを共有として宣言することができない理由です。そのメンバのライフタイムはクラスインスタンスによって管理されず、スレッドブロックによって管理されます。可能であればstaticのクラスメンバを共有できますが、チェックしませんでした。

詳細は、セクションB.2.3のCUDA Programming Guideを参照してください。