1

合体グローバルメモリを理解しようとしています。
奇数の float のセットをグローバル メモリにロードしたいとします。各スレッドは、3 つの float のセットを処理します。これらのフロートが A、B、および C であるとします。 

A0,  B0,  C0
A1,  B1,  C1
A2,  B2,  C2
..          
A19, B19, C19

したがって、スレッドは次のようにデータを取得します。

Thread 0:  A0,  B0,  C0  
Thread 1:  A1,  B1,  C1  
Thread 2:  A2,  B2,  C2
..
Thread 19:  A19, B19, C19

最初のアプローチ: 次
の 3 つの配列をロードできます:float A[20]; float B[20]; floatC[20];データをグローバル メモリにロードするには、cudaMemcpy() を 3 回実行する必要があります。このアプローチは、おそらくうまく合体しないでしょう。

2 番目のアプローチ:
より良いアプローチは次のようなものです。

struct {float A, float B, float C} dataPt;
dataPt data[20];

1 つの cudaMemcpy() でデータをロードできましたが、メモリ アクセスがうまく合体するかどうかはわかりません。

3 番目のアプローチ:

struct {float A, float B, float C, float padding} dataPt2;
dataPt2 data2[20];

また

struct __align__(16){float A, float B, float C} dataPt3;
dataPt3 data3[20];

単一の cudaMemcpy() でデータをグローバル メモリにロードでき、データへのスレッド アクセスが結合されます。(無駄なグローバル メモリを犠牲にして。)

1) 各スレッドが入力データをロードするためにおそらく 3 つのバス サイクルを必要とするため、最初のアプローチは合体しません。
2) 2 番目のアプローチは多くのスレッドを結合しますが、入力データを取得するために 2 つのバス サイクルを必要とするスレッドがいくつかあります。
3) 3 番目のアプローチでは、すべてのスレッドが合体します。

これは正確ですか？2番目と3番目のアプローチに大きな違いはありますか? 3 つのスレッド ディメンション (threadIdx.x、threadIdx.y、threadIdx.z) を使用するアプローチはありますか?

4

2 に答える 2

2

@talonmiesの答えを増幅するだけです。カーネルが次のようになっていると仮定しましょう。

__global__ void kern(float *a, float *b, float *c){

  float local_a, local_b, local_c;
  int idx = threadIdx.x + (blockDim.x * blockIdx.x);

  local_a = a[idx];
  local_b = b[idx];
  local_c = c[idx];
}

最適化を無視し（カーネルが空になる）、32スレッドの1ブロックを起動すると仮定します。

  kern<<<1, 32>>>(d_a, d_b, d_c);

次に、32のスレッド（1つのワープ）がロックステップで実行されます。つまり、各スレッドは次のカーネルコード行を処理します。

  local_a = a[idx];

まったく同時に。（グローバルメモリからの）合体ロードの定義は、ワープがグローバルメモリ（CC 2.0デバイスの場合）の単一の128バイト整列境界内にある一連のデータ項目をロードする場合です。100％の帯域幅使用率で完全に合体した負荷は、各スレッドがその128バイトの整列領域内で1つの一意の32ビット量を使用していることを意味します。スレッド0がa[0]をロードし、スレッド1がa [1]をロードする場合、これは合体したロードの典型的な例である可能性があります。

したがって、最初のケースでは、a []配列はすべて隣接して整列されており、a [0..31]はグローバルメモリの128バイト整列領域内に収まるため、合体した負荷が発生します。スレッド0はa[0]を読み取り、スレッド1はa[1]を読み取ります。

2番目のケースでは、a[0]はa[1]と隣接しておらず、さらに要素a [0..31]（すべて同じコード行にロードされます）は、グローバルメモリ。3番目のケースで何が起こるかを解析させますが、2番目のケースと同様に、要素a [0..31]は隣接しておらず、グローバルメモリ内の単一の128バイト整列領域内に含まれていると言えば十分です。 。ある程度の合体を達成するために隣接するデータ項目を持つ必要はありませんが、32スレッドワープからの100％の帯域幅使用率（「完全に」）合体負荷は、各スレッドが一意の32ビット項目を使用していることを意味します。連続しており、グローバルメモリ内の単一の128バイト整列シーケンス内に含まれています。

便利なメンタルモデルは、Arrary of Structures（AoS）（ケース2と3に対応）とStructure of Arrays（SoA）（基本的に最初のケース）を対比することです。SoAは通常、AoSよりも合体の可能性が高くなります。nvidiaのウェビナーページから、このプレゼンテーション、特にスライド11〜22などがおもしろいと 思うかもしれません。

于 2012-10-22T21:33:19.937 に答える
0

ベスト プラクティス ガイドのその他の関連情報:

コンピューティング機能 2.x のデバイスの場合、要件は非常に簡単に要約できます。ワープのスレッドの同時アクセスは、ワープのすべてのスレッドにサービスを提供するために必要なキャッシュ ラインの数に等しい数のトランザクションに結合されます。 . デフォルトでは、すべてのアクセスは L1 (128 バイト ライン) を介してキャッシュされます。散在するアクセス パターンの場合、オーバーフェッチを減らすために、短い 32 バイトのセグメントをキャッシュする L2 のみにキャッシュすると便利な場合があります (『CUDA C プログラミング ガイド』を参照)。

コンパイラ フラグ: -Xptxas -dlcm=cg L1 キャッシュを無効にします。つまり、合体が不十分なデータには L2 のみを使用します。

于 2012-10-24T21:05:50.900 に答える