CUDA共有メモリ割り当ての定数を宣言する場所

Question

私は1024行列の適合関数を実行しています。各行列は独自のブロックを持ち、同じサイズです。各ブロックはn*nのスレッド（行列の次元）を持ち、n*n共有メモリを持つ必要がありますので、簡単に合計を減らすことができます。しかし、すべての行列の次元nは、実行時に変更可能です（つまり、手動で変更できますが、常に2の累乗であるため、合計は簡単です）。ここで問題となるのは、共有メモリは定数を使って割り当てなければならないが、ホストからカーネルに渡すための値も必要であるということです。サイズnはどこで（カーネルに渡すために）CPUに認識され、（カーネル内で）共有メモリのサイズを宣言するために使用できるように宣言しますか？私のコードは、このような構造になって

：

const int num_states = 1024 
const int dimension = 4 

fitness <<< num_states, dimension * dimension >>> (device_array_of_states, dimension, num_states, device_fitness_return); 

、その後 kernel.cuに私が持っている：： main.cuから

私はカーネルを呼び出す

__global__ void fitness(
    int *numbers, 
    int dimension, 
    int num_states, 
    int *fitness_return) { 
    __shared__ int fitness[16]; <<-- needs to be dimension * dimension 
    //code 
} 

numbersは1024の行列を表す配列であります、dimensionは行と列の長さです。num_statesは1024です。fitness_return ■各行列の適合度を保持する長さ1024の配列。カーネルでは、共有メモリはdimensionの2乗でハードコーディングされています（dimensionはこの例では4です）。

dimensionはどこで共有メモリを割り当てるだけでなく、カーネルを呼び出すために使うことができるのですか？私はdimensionを1か所で更新すればよいのですか？ご協力いただきありがとうございます。

Answer 1

割り当てられた共有メモリの量は、すべてのブロックにわたって一定です。それぞれのブロックで異なる量の共有メモリを実際に使用するかもしれませんが、それでも利用可能です。また、共有メモリの量は関係なく制限されているので、n * n要素はmaximum amount of space（通常は48KiB）を超えることはできません。 float型要素（それぞれ4バイト）の場合は、n < 340程度です。

共有メモリを割り当てるには、静的と動的の2つの方法があります。

静的割り当てが動作しないと思われる、あなたは一例として与えたものである。これらの例で

__shared__ int fitness[16]; 

、サイズはデバイス側コードコンパイル時に（コンパイル時に知られなければなりません時間） - これはあなたには当てはまりません。 ダイナミック共有メモリ割り当てで

、あなたはカーネルコードでサイズを指定しないでん- あなたleave it emptyとプリペンドextern：

extern __shared__ int fitness[]; 

カーネルを起動するときに代わりに、あなたは、金額を指定異なるブロックのスレッドは、それが何であるかを必ずしも知っていない。

あなたのケースでは、スレッドはは何ですかnを知る必要があります。さて、カーネルの引数として渡してください。だから、

__global__ void fitness(
    int *numbers, 
    int dimension, 
    int num_states, 
    int *fitness_return, 
    unsigned short fitness_matrix_order /* that's your n*/) 
{ 
    extern __shared__ int fitness[]; 
    /* ... etc ... */ 
} 

NVIDIAのParallel-for-allブログは、具体的には、静的および動的な共有メモリの割り当てをカバーして共有メモリを使用するより詳細な導入、とa nice postを持っています。