CUDA Performance

並列リダクション（Parallel Reduction）

__shared__ float partialSum[];
// ... 共有メモリに読み込む
unsigned int t = threadIdx.x;
for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
    __syncthreads();
    if (t % (2 * stride) == 0)
        partialSum[t] += partialSum[t + stride];
}

1-2行目: 共有メモリでの要素の総和の計算

4行目: stride = 1, 2, 4, …

5行目: なぜ？

• 6-7行目:
  • 同じ共有メモリで総和を計算する
  • strideが増えると、スレッドはどうなる？

• 第1パス: スレッド1、3、5、7は何もしない
  • 本当は$n$個の要素に対して$n/2$個のスレッドのみが必要である

• 第2パス: スレッド2、6は何もしない

• 第3パス: スレッド4は何もしない

• 一般に、必要なスレッド数は各パス後に半分になる

• 実装を調整したらどうだろう？

__shared__ float partialSum[];
// ... load into shared memory
unsigned int t = threadIdx.x;
for (unsigned int stride = blockDim.x / 2; stride > 0; stride /= 2) {
    __syncthreads();
    if (t < stride)
        partialSum[t] += partialSum[t + stride];
}

• 第1パス: スレッド4、5、6、7は何もしない
  • 本当は$n$個の要素に対して$n/2$個のスレッドのみが必要である

• 第2パス: スレッド2、3は何もしない

• 第3パス: スレッド1は何もしない

• 何が違う？

• stride = 1, 2, 4, ...
  • if (t % (2 * stride) == 0)
• stride = ..., 4, 2, 1
  • if (t < stride)

Warp分割（Warp Partitioning）

• Warp分割: ひとつのブロックからのスレッドをWarpに分ける方法
• Warp分割の情報は以下で使われる
  • divergent branchesを最小化する
  • Warpを早期にリタイアさせる

• 連続的に増加するthreadIdxに基づいて分割する

• 1Dブロック
  • threadIdx.xは0から1023まで(Fermi以降)
    • threadIdx.xは0から511まで(G80/GT200)
  • $n$番目のWarp
    • $32n$番目のスレッドから始まる
    • $32(n+1)-1$番目で終わる
  • ブロックサイズが32の倍数でないならば、最後のWarpはパディングされる

• 2Dブロック
  • threadIdxの増加は以下を意味する
    • threadIdx.xは増加
    • 行はthreadIdx.y == 0から開始
• 例えば、64x8ブロックでは

• 3Dブロック
  • threadIdx.z == 0から始まる
  • 2Dブロックと同様に分割する
  • threadIdx.zを増加し、繰り返す

divergent branchesはWarp内部にある

• warpSize == 32では、以下のコードではいずれのWarpもdivergent branchを持つ

if (threadIdx.x > 15) {
    // ...
}

• warpSize > 1では、以下のコードではいずれのWarpもdivergent branchを持つ

if (threadIdx.x > warpSize - 1) {
    // ...
}

• Warp分割の情報があるとすると、どちらの並列リダクションがより良いか？
  • stride = 1, 2, 4, ...
    • if (t % (2 * stride) == 0)
  • stride = ..., 4, 2, 1
    • if (t < stride)

• warpSize == 2としてみる

• 第1パス
  • stride = 1, 2, 4, ...
    • 4つのdivergent branches
  • stride = ..., 4, 2, 1
    • divergent branchesなし

• 第2パス
  • stride = 1, 2, 4, ...
    • 2つのdivergent branches
  • stride = ..., 4, 2, 1
    • divergent branchesなし

• 第3パス
  • stride = 1, 2, 4, ...
    • 1つのdivergent branch
  • stride = ..., 4, 2, 1
    • 1つのdivergent branch

• 良い分割はWarpが早期にリタイアできるようにする
  • より良いハードウェアの利用

• 並列リダクション

• 第1パス
  • stride = 1, 2, 4, ...
    • Warpのリタイアなし
  • stride = ..., 4, 2, 1
    • 2つのWarpがリタイア

• 第2パス
  • stride = 1, 2, 4, ...
    • 2つのWarpがリタイア
  • stride = ..., 4, 2, 1
    • 1つのWarpがリタイア

Memory Coalescing

• グローバルメモリ に row-major で格納された行列があるとすると、スレッド の望ましいアクセスパターンとは？

• a) column after column
  • 個別のスレッド が増加する連続的なメモリアドレスを読む
• b) row after row
  • 隣接するスレッド が増加する連続的なメモリアドレスを読む

• グローバルメモリの帯域幅(DRAM)
  • G80 --- 86.4GB/s
  • GT200 --- 150 GB/s
  • Fermi --- 192 GB/s
  • Kepler --- 240 GB/s
• DRAMから大きく連続的な位置を要求することでピーク帯域幅を達成する
  • 無作為な位置へのアクセスはさらなる帯域幅の低下となる

• Memory coalescing --- アクセスパターンを組み替えてパフォーマンスを改善する
• こんにちでは役立つが、大きなオンチップキャッシュではあまり役に立たなくなるだろう

• GPUは full-warp での連続的読み込みを単一読み込みにcoalesceする
  • G80/GT200での half-warp (コンピュート1.x)
• 戦略: coalesce可能な方法で共有メモリにグローバルメモリを読み出す
  • 最大帯域幅で共有メモリに無作為にアクセスする
    • bank conflicts を無視すると…

Bank Conflicts

• 共有メモリ
  • ときどき並列データキャッシュと呼ばれる
    • 複数のスレッドが同時に共有メモリにアクセスできる
  • メモリは バンク《bank》 に分けられる

• バンク
  • 各バンクは2サイクルあたり1アドレスを提供できる
  • バンクごとの帯域幅: 2(シェーダクロック)サイクルあたり32ビット
  • 連続する32ビットワードは連続するバンクに割り当てられる

• バンクコンフリクト: 同じバンクだが異なるアドレスへの2つの同時アクセス
  • シリアライズされる
• G80-GT200: 16バンク、8SPで、並列実行
• Fermi以降: 32バンク、16SPで、並列実行

• ファストパス その1(Fermi以降)
  • Warp内のすべてのスレッドが異なるバンクにアクセスする

• ファストパス その2(Fermi以降)
  • Warp内のすべてのスレッドが同じアドレスにアクセスする
  • “ブロードキャスト”

• スローパス(Fermi以降)
  • half-warp内の複数スレッドが同じバンクにアクセスする
  • アクセスがシリアライズされる
  • コストはどのくらい？

__shared__ float shared[256];
// ...
float f = shared[index + s * threadIdx.x];

• sの値がどれくらいであれば、コンフリクトフリーだろうか？

• プロファイラを使わずに、どうやってバンクコンフリクトの除去で期待できるスピードアップの種類を伝えることができるか？
• Warpで1つ以上のスレッドが同じ共有メモリアドレスに書き込む場合に何が起こる？(非アトミック命令)

SM Resource Partitioning

• SMは動的に分割することを思い出そう

• 以下が持てる
  • 96スレッドを8ブロック
  • 192スレッドを4ブロック
  • ただし、192スレッドを8ブロックにはできない

• 以下が持てる(256スレッドブロックを仮定)
  • 768スレッド(3ブロック)それぞれで10レジスタを使う
  • 512スレッド(2ブロック)それぞれで11レジスタを使う

• これ以上のレジスタはスレッドレベルの並列性を損なう
  • パフォーマンスをさらに増加させることはできるか？

• パフォーマンスの崖《cliff》: リソース使用率の増加は並列性における劇的な減少を引き起こす
  • 例えば、グローバルメモリアクセスのレイテンシーを隠蔽しない限り、レジスタ数が増加する

Data Prefetching

• グローバルメモリの読み込みとその利用の間の独立した命令はメモリレイテンシを隠蔽できる

float m = Md[i];
float f = a * b + c * d;
float f2 = m * f;

• Prefetching グローバルメモリからのデータはグローバルメモリの読み込みと使用の間の独立した命令を事実上増加させる

• タイル化された行列の乗算を思い出そう:

for(...) {
    // 現在のタイルを共有メモリにロードする
    __syncthreads();
    // 内積を累積する
    __syncthreads();
}

• プリフェッチ付きタイル化された行列の乗算:

// 最初のタイルをレジスタにロードする
for(...) {
    // レジスタを共有メモリに預ける[deposit]
    __syncthreads();
    // 次のタイルをレジスタにロードする
    // 内積を累積する
    __syncthreads();
}

Instruction Mix

• Special Function Unit(SPU)
  • __sinf()、__expf()を計算するために使う
  • 4つしかないけど、クロックあたり1命令を実行できる

Loop Unrolling

for (int k = 0; k < BLOCK_SIZE; ++k) {
  Pvalue += Ms[ty][k] * Ns[k][tx];
}

• イテレーションあたりの命令
  • 浮動小数点乗算が1つ
  • 浮動小数点加算が1つ
  • ほかには？

• 他のイテレーションあたりの命令
  • ループカウンタの更新
  • 分岐
  • アドレス演算

• 命令ミックス
  • 2つの浮動小数点演算命令
  • 1つのループ分岐命令
  • 2つのアドレス演算命令
  • 1つのループカウンタ増加命令

• 1/3のみが浮動小数点計算
  • でも、完全な理論上の1TFLOPS(Fermi)が欲しい
  • loop unrolling を考慮する

• ループしない
  • ループカウントの更新なし
  • 分岐なし
  • 定数インデックス — アドレス演算命令なし

• 自動的に:
• unrollingの欠点は？

Thread Granularity

• どれだけの処理をスレッドがすべきか？
  • 並列リダクション
    • 2要素を減らす？
  • 行列の乗算
    • Pdの1要素を計算する？

• 行列の乗算
  • Pdの両要素は同じMdの行を必要とする

• 行列の乗算
  • 同じスレッドで両方のPdの要素を計算する
    • グローバルメモリアクセスを1/4減らす
    • 独立した命令数を増加させる
      • その利点は？
    • 新しいカーネルはより多くのレジスタと共有メモリを使う
      • これはどうゆう意味？