arrays - CUDA: 固定サイズの配列はどのメモリ空間に格納されますか?

Question

次のように、カーネルで固定サイズの配列を設定する場合:

int my_array[100];

配列はどのメモリ空間に配置されますか?

特に、そのような配列が >= 2.0 デバイスのレジスタファイルまたは共有メモリに格納されるかどうか、およびその場合の要件は何かを調べたいと思います。

score 8 · Accepted Answer

Fermi (およびおそらくそれ以前のアーキテクチャ) の場合、配列をレジスタファイルに格納するには、次の条件を満たす必要があります。

(1) の理由は、レジスタインデックスが SASS 命令内で直接エンコードされるためです。レジスタを間接的にアドレス指定する方法はありません。

(2) のレジスタ数を制限する主な要因は次のとおりです。

SASS 命令にはレジスタインデックス用に 6 ビットしか含まれていないため、カーネルで使用できるレジスタの数は 64 に制限されます。実際の数は 63 であるため、1 つが何かのために予約されています。
SM には、同時に実行中のすべてのスレッドが共有するレジスタのブロックがあります。
レジスターは変数を保持するためにも必要であるため、コンパイラーは全体的なパフォーマンスを最適化するためにレジスターの使用量を調整する必要があります。

(1) の潜在的な回避策は、ループ展開です。ループが配列へのインデックスとしてループカウンターを使用する場合、ループを展開すると (#pragma unrollまたは手動で)、配列アクセスごとに個別の SASS 命令が存在するため、配列インデックスが定数になります。

この NVIDIA プレゼンテーションに部分的に基づいています: Local Memory and Register Spilling。このドキュメントでは、変数と配列の場所がパフォーマンスに与える影響についても詳しく説明しています。

score 3 · Accepted Answer

定義したカーネル内のローカル配列は、十分なレジスタがない場合、レジスタとローカルメモリに割り当てられます。

配列を共有メモリに割り当てる場合は、次のように指定する必要があります。

__shared__ int my_array[100];

2 に答える 2