プログラミング > CUDA > Fermi Tuning Guide日本語訳

Fermi Tuning Guide日本語訳




次世代CUDA計算アーキテクチャ

Fermiは、NVIDIAの次世代CUDA計算アーキテクチャです。Fermiのホワイトペーパー[1]は、2006年にG80で導入され、後の2008年に変更されたGT200のTeslaアーキテクチャからの主要な改良点の詳細なオーバービューを与えます。

Fermi とTeslaアーキテクチャは、どちらもCUDA計算アーキテクチャのため、プログラミングモデルは、両方に対し同一であり、Teslaアーキテクチャに対するベストプラクティスに従うアプリケーションは、コードのいかなる変更もせずにFermiアーキテクチャでの速度向上が典型的には起こるはずです(特に、倍精度浮動小数点の性能に制約される計算では)。

この文書は、Fermiがこれらの速度向上をさらに加速するためのアプリケーションの最適化の方法の概略を与えます。より詳細は、CUDA C Programming Guida 3.2で文書のいたることろで見ることができます。

CUDAのベストプラクティス

CUDA C Programming Guide[2]とCUDA C Best Practices Guide[3]で述べられた性能の手引きとベストプラクティスは、すべてのCUDAアーキテクチャに適用可能です。プログラマは、最高性能を達成するために以下の推奨を第一に見据えなくてはいけません。

これらの手引きからの高い優先度の推奨は以下の通りです:

  • 逐次コードを並列化する方法を発見せよ
  • ホストとデバイス間のデータ転送を最小にせよ
  • デバイス利用を最大化するようカーネル起動設定を調整せよ
  • グローバルメモリのアクセスがコアレスされていることを確認せよ
  • 可能な時はつねにグローベルメモリのアクセスをシェアードメモリのアクセスに置き換えよ
  • 同一のwarp内で異なる実行パスを避けよ

アプリケーションの互換性

この手引きでカバーしている特定の性能最適化問題に取り組む前に、開発者は、アプリケーションがFermiと互換性が保たれる方法でコンパイルされたか確認するため、Fermi Compatibility Guide for CUDA Applicationsを参照すべきです。

Fermiのチューニング

すべての節の参考文献は、CUDA C Programming Guide 3.2での節です。

デバイス利用

compute capability 1.xのデバイスで全てのマルチプロセッサを利用する唯一の方法は、デバイスでのマルチプロセッサと少なくとも同じ数のスレッドブロックを起動することです(5.2.2節)。compute capability 2.xのデバイスでは、デバイスは複数のカーネルを並行して実行できる(3.2.7.3節)ためアプリケーションはより柔軟性があり、その結果、アプリケーションがデバイスを一つの大きなカーネルではなく、いくつかのより小さいカーネル起動で満たすことも可能になります。これはCUDAストリームを使い行います(3.2.7.5節)。

L1キャッシュ

compute capability 2.xのデバイスは、ローカルとグローバルメモリアクセスをキャッシュするために使われるL1/L2キャッシュ階層が付随しています。プログラマは、L1キャッシングに関する制御を持ちます。

  • 同じオンチップメモリは、L1とシェアードメモリの両方に使われ、そのどの程度をL1対共有メモリにささげるかは、それぞれのカーネル呼び出しに対して設定可能です(G.4.1節)。
  • L1でのグローバルメモリキャッシングは、コンパイル時に無効化できます(G.4.2節)。
  • L1でのローカルメモリキャッシングは、無効化できません(5.3.2.2節)が、プログラマは、コンパイラがローカルメモリにおそらく配置する変数の量を制限すること(5.3.2.2節)や__baunch_bounds()__属性(B.17節)もしくは-maxrregcountコンパイラオプション経由でレジスタスピルを制御することでローカルメモリ利用を制御することができます。

与えられたカーネルに対し、L1でグローバルメモリキャッシング有りもしくは無しの16KBもしくは48KBのL1キャッシュとより多いもしくはより少ないローカルメモリ利用の最適組み合わせを見つけるため実験は推奨されます。多くのローカルメモリを使用する(例:レジスタスピル(5.3.2.2節)の) カーネルは48KBのL1キャッシュから恩恵を受けるかもしれません。

compute capability 1.xのデバイス上では、あるカーネルは、通常のグローバルメモリ読み込みよりも(キャッシュされた)テクスチャフェッチを使う際に高速化が達成できます (たとえば通常の読み込みが十分にコアレスしないとき)。アドレス計算もしくはテクスチャフィルタリングのようなほかの恩恵をテクスチャフェッチが与えないなら、compute capability 2.xのデバイスではこの最適化は、しかしながらグローバルメモリ読み込みはL1でキャッシュされ、L1キャッシュはテクスチャキャッシュよりも高いバンド幅を持つため、逆効果になります。

グローバルメモリ

compute capability 1.xのデバイスでは、グローバルメモリアクセスは、半分のwarp単位で処理され、compute capability 2.xのデバイスでは、それらはwarp単位で処理さます。半分のwarpアクセス単位を仮定するカーネル起動設定を調整することで、したがって性能を向上するかもしれません。2次元のスレッドブロックは、例えば、グローバルメモリにアクセスする際単一のキャッシュラインをそれぞれのワープが指すようにするために、x次元が半分のワープサイズではなくワープサイズの倍数になるようにすべきです。

Tesla とQuadro製品群のFermiベースのGPU(例 Tesla C2050)は、ECCエラー訂正をサポートし、標準でECCは有効化されます。多くの科学アプリケーションのプロのアプリケーションに対する必要性の一方、ECCは削減されたピークメモリバンド幅になります。ECC有効化の状態では、グローバルメモリアクセスに対する実効ピークバンド幅は、約20%減ります(散在した書き込みにたいしてはより多いです)。もしECCが必要でないなら、nvidia-smiユーティリティを使い(もしくはMicrsoft Windowsシステムでコントロールパネル経由で)改善した性能のために無効化することができます。ECCオンまたはオフの切り替えは有効にするため再起動を必要とすることに注意してください。

シェアードメモリ

compute capability 2.xのデバイスのマルチプロセッサは、48KBのシェアードメモリ(G.4.1節)を持つように設定することができ、それはcompute capability 1.xのデバイスよりも3倍多いです。シェアードメモリのサイズが性能の制限となるカーネルは、マルチプロセッサあたりの駐在するブロックの数を増やすことで、もしくはスレッドブロックあたりのシェアードメモリの量を増やし、それに応じてカーネル起動設定を調整することで、これから恩恵を受けます。

compute capability 1.xのデバイスでは、シェアードメモリは16バンクを持ち、アクセスは半分のwarp単位で処理されます。compute capability 2.xのデバイスでは、シェアードメモリは、32バンクを持ち、アクセスはwarp単位で処理されます(G.4.3節)。したがってcompute capability 2.xのデバイスでバンクコンフリクトは、warpの最初の半分に属するスレッドと同じwarpの後ろ半分に属するスレッドの間で起こり、シェアードメモリでのデータレイアウト(例 パディングされた配列)は、バンクコンフリクトを避けるために調整する必要があるかもしれません。

シェアードメモリのハードウェアは、compute capability 2.xのデバイスで複数のブロードキャストワードのサポートとスレッド毎に8ビット、16ビット、64ビット、128ビットのアクセスでより少ないバンクコンフリクトを生成するように改善されます。

コンスタントキャッシュ

すべてのcompute capabilitのデバイスでサポートされるコンスタントメモリ空間(ここで、__constant__変数が与えれれている)に加え、compute capability 2.xのデバイスは、コンパイラがある状態ではコンスタントキャッシュ経由でいかなる変数をも読み込むために使われるLDU(LoaD Uniform)命令をサポートします(G.4.4節)。

32ビット整数乗算

compute capability 1.xのデバイスでは、32ビット整数乗算は、ネイティブにサポートされないために複数の命令を使い実装されます。24ビット整数乗算は__[u]mul24組み込み命令経由でネイティブにサポートされます。

compute capability 2.xのデバイスでは、しかしながら、32ビット整数乗算がネイティブにサポートされますが、24ビット整数乗算はそうではありません。したがい __[u]mul24は複数命令を使い実装され、利用すべきではありません(5.4.1節)。

IEEE 754-2008準拠

compute capability 2.xのデバイスは、特に単精度で、IEEE 754-2008浮動小数点標準からの逸脱がcompute capability 1.xのデバイスよりかなり少なくなります(G.2節)。

特に、compute capability 1.xのデバイスで加算と乗算は得てしてFMAD命令へ結合され、それはIEEE準拠ではありません。これに反し、compute capability 2.xのデバイスでそれらはFFMA命令へ結合され、それはIEEE準拠です。

デフォルトでnvccはIEEE準拠コードを生成しますが、compute capability 2.xのデバイスが以前のデバイスに対して生成されたコードに近いコードを生成するオプションを提供します。-flz=true(非正規化数はゼロにフラッシュされる)、-prec-div=false(より低精度の除算)、さらに-prec-sqrt=false(より低精度の平方根)。この方法でコンパイルされたコードは、デフォルトの設定でコンパイルされたコードより高いパフォーマンスを持つ傾向があります(5.4.1節)。

C++サポート

compute capability 2.xのデバイスは、いくつかの制限があるC++クラスをサポートし、それはprogramming guideのD.6節で詳細に記述されています。

進歩した関数

compute capability 2.xのデバイスは追加の進歩した組み込み関数をサポートします。

  • 別のメモリフェンス関数:__threadfence_system() (B.5節)
  • __syncthreads()同期関数の発展形:__syncthreads_count()、__syncthreads_and()、と__syncthreads_or() (B.6節)
  • サーフィス関数(B.9節)
  • グローバルもしくはシェアードメモリでの32ビットワードの浮動小数点のアトミック加算操作 (B.11.1.1節)
  • 別のwarp議決関数:__ballot() (B.12節)

compute capability 2.xのデバイスでアトミック関数は以前のデバイスより高いスループットを持ちます。

参考文献

[ 1 ] NVIDIA’s Next Generation CUDA Compute Architecture: Fermi http://www.nvidia.com/content/PDF/fermi_white_papers/NVIDIAFermiComputeArchitectureWhitepaper.pdf

[ 2 ] NVIDIA CUDA C Programming Guide from CUDA Toolkit 3.2

[ 3 ] NVIDIA CUDA C Best Practices Guide from CUDA Toolkit 3.2

[ 4 ] NVIDIA CUDA Reference Manual from CUDA Toolkit 3.2

(本文の訳はこれで完了)

最終更新:2011年01月04日 23:32