このパートの主なテーマは、GPUのプログラミング時にグローバルメモリを操作する最適化です。
GPUには多くの機能がありますが、グローバルメモリを使用するとパフォーマンスの損失が複数発生する可能性があることを無視しています。 しかし、すべての微妙な点を考慮すると、非常に効果的なCUDAプログラムを取得できます。
降りる。
グローバルメモリの何が問題になっていますか？
グローバルメモリのボリュームは、すべての種類のメモリの中で最大ですが、同時に、このメモリは技術的特性（読み取り速度と書き込み速度）の点で最も低速です。
前のパートでは、行列の転置の例を調べました。 パフォーマンスを改善するために、共有メモリバッファが使用されました。これにより、パフォーマンスがほぼ4倍になりました。 しかし、追加のメディエーターでこの増加を見るのは十分に奇妙でした。 その秘密は、グローバルメモリの正しい使用にあります。
グローバルメモリの操作を最適化するには、使用する型のサイズを調整する方法と、クエリを組み合わせて使用する方法の2つがあります。
使用されるタイプのサイズの調整
データ型の配置により、クエリをグローバルメモリにコンパイルして1つのGPUコマンドにすることができます。コンパイルしない場合、コンパイラーは追加のコードを生成し、パフォーマンスを大幅に低下させる可能性があります。 最適なパフォーマンスを得るには、データ型を4、8、または16バイトにする必要があります。
型サイズが4、8、または16バイトと一致しない場合、より大きな次元の型を使用するか、キーワード__align __（位置合わせサイズ）を使用して位置合わせすることをお勧めします。
組み込みのCUDAタイプを使用した最適化の例。
タイプint3のサイズは12バイトです。メモリへのアクセスは最適ではありません。
__device__ int3 data[512];
__global__ void initData()
{
int idx = threadIdx.x
data[idx] = make_int3(idx, idx, idx);
};
* This source code was highlighted with Source Code Highlighter .
4番目のコンポーネントが不要な場合でも、int4型（16バイト）を使用することをお勧めします。
__device__ int4 data[512];
__global__ void initData()
{
int idx = threadIdx.x
data[idx] = make_int4(idx, idx, idx, 0);
};
* This source code was highlighted with Source Code Highlighter .
構造体を使用する場合は、__ align__キーワードを使用する必要があります。これにより、指定したサイズに型を揃えることができます。
構造のサイズのアライメントの例。
整列前のvector3構造体のサイズは12バイトです。
struct vector3
{
float x;
float y;
float z;
};
int main()
{
printf( "%i\n" , sizeof (vector3));
return 0;
};
* This source code was highlighted with Source Code Highlighter .
数字の12がコンソールに表示されます。
整列後、vector3のサイズは16バイトです。
struct __align__(16) vector3
{
float x;
float y;
float z;
};
int main()
{
printf( "%i\n" , sizeof (vector3));
return 0;
};
* This source code was highlighted with Source Code Highlighter .
数字の16がコンソールに表示されます。
フェデレーションクエリを使用する
グローバルメモリへの多数のリクエストを1つに結合することで、パフォーマンスを大幅に向上させることができます（リクエストはトランザクションと呼ばれることもあります）。 nVidiaのドキュメントでは、これを合体グローバルメモリアクセスと呼びます。 ただし、要求をメモリに結合するために必要なものの直接の議論に進む前に、GPUの動作に関するいくつかの追加事項を知る必要があります。
GPUは、いわゆるワープを使用してスレッドの実行を制御します。 ソフトウェアの観点から見ると、ワープはスレッドプールを表します。 カーネルが呼び出されたときに要求されたスレッドの並列操作が発生するのはこのワープ内であり、スレッドが相互に対話できるのはワープです。 すべてのGPUのワープのサイズは32です。つまり、ワープで実行されるスレッドは32個だけです。 同時に、いくつかのワープをGPUで起動できます。この量は、使用可能なレジスタと共有メモリのサイズによって決まります。 もう1つの興味深い機能は、ハーフワープを使用してメモリにアクセスすることです。つまり、最初は最初の16スレッドがメモリにアクセスし、次に16スレッドの後半がアクセスされます。 なぜこのようにアクセスするのか、確かに言うことはできませんが、これはGPUの主要なタスクであるグラフィックス処理によるものであるとしか考えられません。
次に、クエリをグローバルメモリに結合するために必要な要件を検討します。 メモリアクセスはハーフワープによって発生することを忘れないでください。
メモリにアクセスするときに結合するために必要な条件は、Compute Capabilityのバージョンによって異なります。バージョン1.0および1.1の場合は、それらを指定します。詳細については、nVidiaのドキュメントを参照してください。
- スレッドは、1つの128バイトブロック（トランザクション）を与えながら、1つの64バイトブロック（トランザクション）になる32ビットワードまたは64ビットワードにアクセスする必要があります。
- 128ビットワードの呼び出しが使用される場合、結果として2つのトランザクションが実行され、それぞれが128バイトの情報を返します。
- スレッドはメモリ要素に順番にアクセスする必要があり、次の各スレッドはメモリ内の次のワードに対応する必要があります（一部のスレッドは対応するワードにまったく対応しない場合があります）
- 16ワードはすべて、アクセスされるメモリブロック内にある必要があります。
条件に関する注意事項：
- 言葉はあらゆる種類のデータを意味します。主なことは、必要な次元への準拠です。
- ワードの次元はビット単位で、結果のデータブロックの次元はバイト単位で与えられます。
図 1.メモリにアクセスするときに結合を提供するクエリ
図 図1は、単一のトランザクションを提供するグローバルメモリクエリの例を示しています。 左側では、すべての条件が満たされています。ハーフワープの各スレッドは順番に対応する32ビットワードにアクセスし、メモリの先頭のアドレスはトランザクションブロックのサイズに合わせられます（16スレッド* 4バイト= 64バイト）。 右側は、ブロックからの一部のフローがメモリ内の対応するワードにまったくアクセスしない場合の例です。
図 2.メモリにアクセスするときにユニオンを提供しないクエリ
図 図2は、グローバルメモリにアクセスするときに結合を与えない例を示しています。 左側では、メモリ内の対応する単語へのスレッドの循環の条件は満たされていません。 右側では、メモリアドレスをブロックサイズに合わせるための条件が満たされていません。 その結果、1つの統合されたトランザクションの代わりに、各ハーフワープスレッドに1つずつ、16の個別のトランザクションを取得します。
配列構造または構造の配列？
構造での作業の問題と生産性の向上を達成する方法について、いくつかの言葉を与える必要があります。 構造体の配列を使用する必要がある場合は、構造体コンポーネントの個別の配列を作成することをお勧めします。これにより、関連付けによるグローバルメモリへの要求の数が削減されます。
例を考えてみましょう。
グローバルメモリを使用した無効な作業：
struct __align__(16) vec3
{
float x;
float y;
float z;
};
__device__ vec3 data[SIZE];
__global__ void initData()
{
int idx = blockDim.x * blockIdx.x + threadIdx.x;
data[idx].x = idx;
data[idx].y = idx * 2;
data[idx].z = idx * 3;
};
* This source code was highlighted with Source Code Highlighter .
個別の配列を使用する方が効率的です。
__device__ float x[SIZE];
__device__ float y[SIZE];
__device__ float z[SIZE];
__global__ void initArr()
{
int idx = blockDim.x * blockIdx.x + threadIdx.x;
x[idx] = idx;
y[idx] = idx * 2;
z[idx] = idx * 3;
};
* This source code was highlighted with Source Code Highlighter .
最初のケースでは、ベクトルの配列を使用して構造体の各フィールドにアクセスし、メモリへの個別の要求が必要です。2番目のケースでは、各ハーフワープに対して3つのクエリを組み合わせることで十分です。 このアプローチにより、平均して生産性を2倍に高めることができます。
おわりに
上記のすべての結論として、CUDAでメモリを操作する際に最も重要なアドバイスをしたいと思います。
複数のスレッドで1つのメモリセルの値を同時に変更しようとしないでください。
これは、マルチスレッドプログラミングで最もよくある間違いです。 実際、CUDAは特定のメモリ領域への各スレッドのアトミックアクセスを保証しないため、結果が期待どおりに正確にならない場合があります。 CUDAにはアトミック操作が存在しますが、不変データの概念を使用し、計算の結果を新しいオブジェクトに保存して、次の計算ステージに転送することをお勧めします。