4

CUDA プログラムの NVPTX(LLVM IR) には、0 から 5 までのメモリ アドレス空間の識別子があります (下の表を参照)。

ここに画像の説明を入力

同じ LLVM IR プログラムで見たのですが、メモリ アドレスは、図に示すように「ジェネリック」またはその他のタイプとして識別されます。

「汎用」の場合 (デフォルトでは、識別子はありません): ここに画像の説明を入力

共有対象'： ここに画像の説明を入力

私の質問は、汎用メモリ アドレス空間の場合、データは実際にはハードウェア、オフチップ、オンチップ メモリ、またはローカル レジスタのどこにあるのでしょうか? 一般的なタイプのアドレス空間が最終的にどのように管理されるかを誰か説明できますか?

4

1 に答える 1

11

答えは簡単です。汎用アドレス空間にはハードウェア表現がありません。

汎用アドレス空間 (AS) は、他の各 AS が組み合わされた論理 AS として見ることができます。例: 次のカーネル呼び出しと、ポインターを受け入れるデバイス関数。

__device__ void bar(int* x){
   *x = *x + 1;
}

__global__ void foo(int* x){
   __shared__ int y[1];
   bar(x); 
   bar(y);
}

関数には任意のポインターを渡すことができます。言語の観点からは、ポインターが AS 1 (グローバル) または AS 3 (共有) のどちらにあるかは関係ありません。C++ (および CUDA C/C++) では、AS を明示的に指定する必要はありません。たとえば、OpenCL < 2.0 では、各ポインターに修飾子を明示的に追加しbar、特定の AS ポインターを受け取る関数を提供する必要があります。

LLVM IR で何が起こるかというと、ポインター witch が関数に渡されaddresspacecast、汎用 AS への命令を介してキャストされます。PTXでは、次の命令addresspacecastで表されます。cvta

// convert const, global, local, or shared address to generic address
cvta.space.size  p, a;        // source address in register a
cvta.space.size  p, var;      // get generic address of var
cvta.space.size  p, var+imm;  // generic address of var+offset

// convert generic address to const, global, local, or shared address
cvta.to.space.size  p, a;

.space = { .const, .global, .local, .shared };
.size  = { .u32, .u64 };

ジェネリック ポインターは、他の AS 用に予約されたアドレス領域内にない限り、グローバル メモリにマップされます。ハードウェアは、汎用ポインタから AS の開始アドレスを差し引いて、正しいメモリ領域を決定します。

アトミックは良い例です:

atom{.space}.op.type  d, [a], b;
atom{.space}.op.type  d, [a], b, c;

アドレス空間を指定するか、ハードウェアに選択させることができます。ポインター減算オーバーヘッドなしで正しいアトミック命令を生成したい場合、バックエンドはポインターを正しいアドレス空間にキャストする責任があります。

于 2015-09-09T05:45:38.720 に答える