• トップ
  • G-DEPについて
  • ご購入ガイド
  • サポート
  • お問い合わせ

G-DEPトップ  >  G-DEPの高速演算記  >  高速演算記 第3回 「チューニング技法その1 CUDAプログラミングガイドからピックアップ」

高速演算記 第3回 「チューニング技法その1 CUDAプログラミングガイドからピックアップ」

今回はFermi解説から一旦はずれ、実際にプログラムを開発する際のチューニングポイントを紹介したいと思います。CPUで実行していたコードをそのままGPU上で動作させても速度向上が得られない場合は以下の点を確認してみてはいかがでしょうか。

まずはじめに基本機能のおさらいをしたいと思います。
前回からご紹介している通り、Fermi内部ではマルチプロセッサと呼ばれる実行ユニットが複数存在しています。各実行ユニットは1つのスケジューラと、32個の演算コアから構成され、スレッドを2サイクルあたり2ワープ処理することが可能となっています。32スレッドをワープと呼ばれる1つの単位として扱い、GPU上でスレッドを実行する上で扱いやすい単位とされています。また、少し粗い単位として複数スレッドをブロックとして扱うことが可能です。1ブロック内の最低スレッド数はありませんので、1スレッドのブロックを構成することも可能ですが、性能を考慮すると、複数ワープ単位でブロックを構成したほうがよいことが見受けられます。

CUDAカーネルをCから実行する際にブロック数とスレッド数を指定してカーネルを起動します。1つの実行ユニットが一度に保持できる最大ブロック数と最大スレッド数は下記の通りとなっています[2]。

最大保持ブロック数 8
最大保持スレッド数 1536

TeslaC2050全体では、GPU上に存在する実行ユニット数 x 最大保持ブロック数 = 14x8 = 112となります。それ以上のブロック数を指定した場合、残りのブロック数は待機状態となり、実行ユニット内に保持されているブロックの実行が終了次第、待機ブロックと置き換わります。

各実行ユニットのスケジューラは内部で保持しているブロックから実行可能なワープを選択して実行します。コード内で演算を行ったり、メモリアクセスを行った場合、その処理が完了するまで遅延が発生するため、スケジューラはその間、実行可能な別スレッドを選択し実行するように指示をだします。


実行ユニット内に保持されるブロック数とスレッド数はどのようなパラメータに影響されるのでしょうか。実行ユニットは有限の共有メモリとレジスタを持っており、それらを実行ユニット内のブロックとスレッドに割り当てていきますので、1ブロックあたりの共有メモリの使用量、1スレッドあたりのレジスタ使用量により実行ユニットが保持できるブロック、スレッド数が決定されます。Fermiでは、実行ユニットあたりのレジスタ数と共有メモリの大きさは下記の通りとなっています。

実行ユニットあたりのレジスタ数 32768
実行ユニットあたりの共有メモリ 16kb / 48kb

ここでレジスタ数は32bitレジスタの数で、long long, double値や64ビット環境を対象としたCUDAプログラムではポインタなどが64bitとして扱われますのでレジスタを2つ使用することになります。例として、手元にあるCUDAカーネルをコンパイルした際に以下のようなリソース使用となりました。

レジスタ数 55
共有メモリ 8448 bytes

これをもとに計算しますと、実行ユニットあたり 595スレッド、1ブロック使用可能である計算になります。この構成ですと共有メモリの使用量が全体の半分ちょっととなりますので、実行ユニットが一度に保持可能なブロック数は1つのみとなります。595スレッド同時に実行したい場合、1ブロックが595スレッド全てをもつことになります。

プログラムが使用する共有メモリの量とレジスタ数を確認するには、nvccを通してCUDAアセンブラへ渡す引数 -Xptxas=-v を指定することで表示できます。また、例えばcudaランタイム関数cudaFuncGetAttributes()よりホストプログラム上で同様の情報を得ることが可能です。

1スレッドのレジスタ数を制限するにはnvccへ渡す引数 -maxrregcount で最大レジスタ数を指定できます。また、関数定義に __launch_bounds__ マクロを付与することによって関数レベルで細かく指定することが可能です。__launch_bounds__はコンパイラオプションよりも優先的に適用されます。

それでは[1]に掲載されている事項より、いくつか重要項目をピックアップして紹介していきます。


実際にレジスタ数を制限し実行可能スレッド数を増やした場合、本来レジスタにあった部分がローカルメモリに退避されます。ローカルメモリはFermiより前のアーキテクチャではキャッシュされていないグローバルメモリと同じアクセス速度のため、カーネルの複雑さによってはカーネルの実行時間に大きく影響を及ぼすことがあります。実験を通して最適な値を求めることが望ましいのではないでしょうか。

[1]の4.3章ではデータの書き込み後の読み出しは24サイクルのレイテンシがかかるので、このレイテンシを隠蔽するために必要なスレッド数はFermi では実行ユニットあたり768スレッド程度と記載があります。また、__syncthreads()のようにブロック内の同期をとるコマンドが実行されている間、他ブロックのワープを実行できるよう、各実行ユニットは2つ以上のブロックが実行可能であることが望ましく、1つの大きいブロックを構成するよりも3-4つの小さめのブロック構成でカーネルを実行したほうが効率的との記載があります。

また、カーネルがメモリバウンドであるのか、コンピュートバウンドであるのかによって、最適化するポイントが異なってきます。外部メモリを多くアクセスする場合、実行ユニットあたり少ないスレッド数を使用してもメモリ転送能力の最大値近傍でカーネルが実行されることもあります。この場合はスレッド数をそれ以上多く割り当ててもメモリからのデータ待ちとなり処理速度にはほとんど影響しません。

手元にあるカーネルをいろいろなブロック数、スレッド数の組み合わせで100回ずつ実行した結果の平均を下に示しました。
本カーネルでは1ブロックあたり256スレッド、一度に56ブロック起動することで処理時間を最小にすることが可能になります。カーネルの起動構成を変更するだけでも処理時間が変化しますので、検討してみる価値はあります。

最後に、[3]ではGPUのレジスタ数とキャッシュとの関連について触れています。Fermiでは32bitレジスタが合計2MB、L1キャッシュが1MB、L2キャッシュが768kb存在し、キャッシュはCPUから遠くなるほど小さくなるということを指摘しています(このスライドは実行ユニットが16個存在していると想定)。ただし共有メモリは他実行ユニットからは、レジスタは他スレッドからは見えないため、同じ情報はスレッドごとにコピーが存在することになります。レジスタを有効利用するには、少ないスレッドでより多くの出力を行い、スレッドあたり利用可能なレジスタ数を多くとるとのことです。

 

なお、FermiアーキテクチャはTeslaC2050, GTX480で採用されているGF100のほか、GTX460で採用されているGF104なども存在し、内部のハードウェア構成が多少変更されています。実際にどのGPUをターゲットするかによってパラメータを変更する必要があるかもしれません。アーキテクチャごとに実行されるカーネルコードを変更したい場合、コンパイラのマクロが便利です。__CUDA_ARCH__はその時コンパイルされるアーキテクチャを定義します。例えば、アーキテクチャ13とアーキテクチャ20用のコードを生成するようにnvccへ引数を渡したとします。

nvcc -gencode arch=compute_13,code=sm_13 -gencode arch=compute_20,code=sm_20

すると、カーネルのコンパイルは2度行われますが、それぞれ__CUDA_ARCH__マクロは 200と130として定義されるので、このマクロ定義を識別することによってアーキテクチャごとに異なる最適なコードをコンパイルすることが可能となります。

同じ問題に対して計算方法が複数存在するのと同様にどんなカーネルでも適用できるテクニックはありませんが、幾つかのポイントをピックアップして掲載してみました。今後のカーネル開発の参考にしてみてはいかがでしょうか。

[1] NVIDIA CUDA C Best Practices Guide Version 3.1
[2] NVIDIA CUDA C Programming Guide 3.1, Appendix G.1
[3] V. Volkov, Programming inverse memory hirerarchy: case of stencils on GPUs

(G-DEP CUDAエンジニア 田原哲雄)