次のコードスニペットがあります。
#include <stdio.h>
struct Nonsense {
float3 group;
float other;
};
__global__ void coalesced(float4* float4Array, Nonsense* nonsenseArray) {
float4 someCoordinate = float4Array[threadIdx.x];
someCoordinate.x = 5;
float4Array[threadIdx.x] = someCoordinate;
Nonsense nonsenseValue = nonsenseArray[threadIdx.x];
nonsenseValue.other = 3;
nonsenseArray[threadIdx.x] = nonsenseValue;
}
int main() {
float4* float4Array;
cudaMalloc(&float4Array, 32 * sizeof(float4));
cudaMemset(float4Array, 32 * sizeof(float4), 0);
Nonsense* nonsenseArray;
cudaMalloc(&nonsenseArray, 32 * sizeof(Nonsense));
cudaMemset(nonsenseArray, 32 * sizeof(Nonsense), 0);
coalesced<<<1, 32>>>(float4Array, nonsenseArray);
cudaDeviceSynchronize();
return 0;
}
これをNsightのNvidiaプロファイラーで実行し、グローバルメモリアクセスパターンを見ると、float4Arrayには完全に合体した読み取りと書き込みがあります。一方、ナンセンス配列のアクセスパターンは不十分です（構造体の配列であるため）。
NVCCは、概念的には構造体の配列であるfloat4配列を、メモリアクセスパターンを改善するために配列の構造体に自動的に変換しますか？
いいえ、配列の構造体には変換されません。これを注意深く考えれば、コンパイラがこのようにデータを再編成することはほぼ不可能であると結論付けることができると思います。結局のところ、渡されているのはポインタです。
配列は1つだけであり、その配列の要素には、同じ順序で構造体要素があります。
float address (i.e. index): 0 1 2 3 4 5 ...
array element : a[0].x a[0].y a[0].z a[0].w a[1].x a[1].y ...
ただし、float4コンパイラはスレッドごとに1つの16バイトのロードを生成するため、配列の方がパターンが優れています。これは、float4スレッドごとにベクトル（この場合）をロードするため、「ベクトルロード」と呼ばれることもあります。したがって、隣接するスレッドは引き続き隣接するデータを読み取っており、理想的な合体動作があります。上記の例では、スレッド0が読み出さなりa[0].x、a[0].y、a[0].zおよびa[0].w、スレッド1が読み取ることになるa[1].x、a[1].y等これのすべては、単一で起こるであろう要求（すなわちSASS命令）が、複数に分割することができる取引。リクエストを複数のトランザクションに分割しても、効率が低下することはありません（この場合）。
Nonsense構造体の場合、コンパイラはその構造体も同様の方法でロードできることを認識しないため、内部ではスレッドごとに3つまたは4つのロードを生成する必要があります。
float3 groupfloat3 groupfloat otherおそらく上の図を使用して、スレッドごとに上記の負荷をマップすると、各負荷にストライド（スレッドごとにロードされたアイテム間の未使用の要素）が含まれるため、効率が低下することがわかります。
構造体で注意深い型キャストまたは共用体定義を使用することにより、コンパイラーにNonsense1回のロードで構造体をロードさせることができます。
この回答には、AoS-> SoA変換および関連する効率の向上に関連するいくつかのアイデアも含まれています。
この回答は、ベクトル負荷の詳細をカバーしています。
この記事はインターネットから収集されたものであり、転載の際にはソースを示してください。
侵害の場合は、連絡してください[email protected]
コメントを追加