1

1024 行列のフィットネス関数を実行しています。各行列は独自のブロックを取得し、同じサイズです。各ブロックにはn*nスレッド (行列の次元)n*nがあり、合計を簡単に削減できるように共有メモリが必要です。ただし、すべての行列の次元nは実行前に可変です (つまり、手動で変更できますが、常に 2 のべき乗であるため、合計は単純です)。ここでの問題は、定数を使用して共有メモリを割り当てる必要があることですが、ホストからカーネルに渡す値も必要です。n(カーネルに渡すために) CPU に表示され、(カーネル内で) 共有メモリのサイズを宣言するために使用できるように、ディメンションをどこで宣言すればよいですか?

私のコードは次のように構成されています。

からmain.cuカーネルを呼び出します：

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私は持っています:

__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 の配列で、各行列の適合値を保持します。カーネルでは、共有メモリは の 2 乗でハードコーディングされていますdimension(dimensionこの例では 4 です)。

dimension共有メモリの割り当てとカーネルの呼び出しに使用できるように、どこでどのように宣言できますdimensionか?この方法では、1 か所で更新するだけで済みます。ご協力いただきありがとうございます。

4

1 に答える 1

3

割り当てられた共有メモリの量は、すべてのブロックで均一です。各ブロックで異なる量の共有メモリを使用している可能性がありますが、それでもすべて使用可能です。また、共有メモリの量はかなり制限されているため、n*n 要素が最大容量(通常は 48KiB) を超えることはできません。for float-type 要素 (それぞれ 4 バイト) は、n < 340 程度を意味します。

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

静的割り当ては、例として挙げたものであり、機能しません。

__shared__ int fitness[16];

このような場合、サイズはコンパイル時 (デバイス側のコードのコンパイル時) にわかっている必要がありますが、これは当てはまりません。

動的共有メモリ割り当てでは、カーネル コードでサイズを指定しません。空のままにして先頭に を追加し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ブログには、共有メモリの使用に関するより詳細な紹介が記載された素晴らしい投稿があり、具体的には静的および動的な共有メモリの割り当てについて説明しています。

于 2016-12-30T00:11:48.330 に答える