42

定数パラメーターを使用して共有メモリを割り当てようとしていますが、エラーが発生します。私のカーネルは次のようになります。

__global__ void Kernel(const int count)
{
    __shared__ int a[count];
}

そして、私はエラーが発生しています

エラー: 式には定数値が必要です

カウントは定数です！このエラーが発生するのはなぜですか? どうすればこれを回避できますか？

4

5 に答える 5

93

CUDA は動的共有メモリ割り当てをサポートしています。次のようにカーネルを定義すると:

__global__ void Kernel(const int count)
{
    extern __shared__ int a[];
}

次に、必要なバイト数をカーネル起動の 3 番目の引数として渡します

Kernel<<< gridDim, blockDim, a_size >>>(count)

その後、実行時にサイズを変更できます。ランタイムは、ブロックごとに 1 つの動的に宣言された割り当てのみをサポートすることに注意してください。さらに必要な場合は、その単一の割り当て内のオフセットへのポインターを使用する必要があります。また、ポインタを使用する場合は、共有メモリが 32 ビット ワードを使用し、共有メモリ割り当てのタイプに関係なく、すべての割り当てが 32 ビット ワードでアライメントされている必要があることに注意してください。

于 2011-04-03T18:40:42.510 に答える
39

const「定数」という意味ではなく、「読み取り専用」を意味します。

定数式は、コンパイル時にコンパイラに値が認識されるものです。

于 2011-04-03T17:36:06.693 に答える
21

オプション 1: 共有メモリを定数値で宣言します (と同じではありませんconst) 。

__global__ void Kernel(int count_a, int count_b)
{
    __shared__ int a[100];
    __shared__ int b[4];
}

オプション 2: カーネル起動構成で共有メモリを動的に宣言します。

__global__ void Kernel(int count_a, int count_b)
{
    extern __shared__ int *shared;
    int *a = &shared[0]; //a is manually set at the beginning of shared
    int *b = &shared[count_a]; //b is manually set at the end of a
}

sharedMemory = count_a*size(int) + size_b*size(int);
Kernel <<<numBlocks, threadsPerBlock, sharedMemory>>> (count_a, count_b);

注: 動的共有メモリへのポインタには、すべて同じアドレスが割り当てられます。2 つの共有メモリ配列を使用して、共有メモリに 2 つの配列を手動でセットアップする方法を説明します。

于 2011-04-04T16:57:39.480 に答える
1

このように共有変数を宣言することはできません..

__shared__ int a[count];

ただし、配列 a の最大サイズについて十分に確信がある場合は、次のように直接宣言できます

__shared__ int a[100];

ただし、この場合、プログラム内にいくつのブロックがあるかを心配する必要があります。共有メモリをブロックに固定する (そして完全に利用されない) と、グローバル メモリとのコンテキストの切り替え (高レイテンシ) が発生し、パフォーマンスが低下するためです。 ...

この問題を宣言するための優れた解決策があります 

extern __shared__ int a[];

次のようにメモリからカーネルを呼び出しながらメモリを割り当てます

Kernel<<< gridDim, blockDim, a_size >>>(count)

ただし、カーネルで割り当てているよりも多くのメモリをブロックで使用している場合、予期しない結果が生じるため、ここでも気にする必要があります。

于 2011-04-04T07:00:21.417 に答える