AMD GCN isa と NVIDA SASS でのアドレス演算のチューニング

ちょっとしたきっかけがあって、AMD用のGPUとNVIDIA用のGPU両方で高速化作業を行いました。

そのときに得られた知見を書いておきます。

AMD GCN isa と NVIDIA SASS

GPUでプログラミングを行うときは、OpenCL、CUDA などを使うと思いますが、これらはより低レベルなGPU用の機械語にコンパイルされます。

AMD のGPUでは、これに”ISA”、NVIDIA の GPU では、これを”SASS”という名前が付けられています。ISAは一般的な単語なので、別に正しい名前があるかもしれませんが、ビルド時のオプションなどで指定する場合は、–isa 等を指定しているので、ここでは、GCN isa とします。

通常のGPUプログラミングでこれらを見ることはないかもしれませんが、レジスタがわずかに足りない場合、性能がわずかに足りない場合などは、これを調べると改善できる可能性があります。

今回の作業で得られた経験をもとに、SASS、GCN isaレベルでのチューニング手法に書いていきます。

SASS、GCN isa の確認方法

SASS は、CUDA Toolkit に含まれる、nvcc, nvdisas を使うと確認できます。

(昔似たような文章を書いたことがあるので http://int.main.jp/txt/sass/も参照してください)

  1. $ nvcc –cubin -arch=sm_20 などで cubin を出す
  2. $ nvdisasm file.cubin のようにして SASSを出す

GCN isa は AMD CodeXL(http://developer.amd.com/tools-and-sdks/opencl-zone/codexl/)に含まれる、CodeXLAnalyzer を使うと出力できます。

$ CodeXLAnalyzer -s CL <input.cl> -k <kernel名> –isa <output.isa> -c <GPU名>

GPU 名は、一覧を書いた公式ドキュメントが見当たらないのでわからないですが、clGetDeviceInfo(CL_DEVICE_NAME) で確認できる名前を指定すればよさそうです。

共通のチューニング

いくらか見た範囲だと、アドレス計算はどちらのGPUでも無駄な命令が出やすいようです。これには以下の理由が考えられます

  • 64bit環境ではアドレスは64bit整数になるが、どちらのGPUも32bit整数命令しか無い
  • コンパイラがそれほど良い最適化をしてくれない
  • 高性能なCPUと違ってスーパースカラの幅が狭いのでアドレス演算がパイプラインを埋めてしまう

たとえば、次のプログラムを

nvcc + sm_30 でコンパイルすると ループ部分は ↓ のようなSASSが出ます。

ST.E(どこにも説明書いてないですが、おそらくメモリへのストアだと思います) が四回出ていることから、おそらく4回アンロールされているのだと思いますが、4回のストアをするだけのループにしては、22命令は多すぎではないかという気がします

まじめに検証したわけではないですが、NVIDIA、AMDどちらのGPUのコンパイラも、ポインタに関しては、かなりナイーブなコードを出すようです。最適化しやすい形に変形してあげる必要があるでしょう。

上の場合なら、

こんな感じですね。

9命令になりました。このぐらいなら妥当に見えます。

このように、アドレス演算は配慮するかしないかで、命令量は結構変わります。先日作業していたプログラムは、演算律速の問題だったので 1.5倍くらい変わったと思います。

また、これは、GPUで画像を扱う場合等に、(x,y)とGPUのスレッドを1:1で割り当てるのはあまり良くないということを示しています。

例えば、

こういうコードを書いてしまうと、各画素ごとにy*step + xの計算が入ってしまいます。もう少し大きい粒度でスレッドを割り当てて、

このように修正することで性能が改善する可能性があります。

GCN isa のチューニング

GCN isa はwavefront中の64個のワークアイテムが同時に複数のデータに対して命令を実行するSIMD命令と、wavefront1個で実行されるスカラ命令があります。

CodeXLAnalyzer で確認できる命令列のうち、v_ で始まるものが、SIMD命令、s_ で始まるものがスカラ命令になります。

このSIMD命令とスカラ命令は、パイプラインが別になっており、十分な数のワークアイテム(おそらく8wavefront = 512)を投入すれば、2個のwavefrontから64個のSIMD命令と1個のスカラ命令を同時に実行することができます。

全ワークアイテムで共通するアドレス計算をスカラ命令にすることで、データ演算と並行してアドレス計算を実行できるようになります。

どのような演算がスカラ命令になるかはよく調べていないですが、コンパイラは以下の条件が成立する値をスカラ値とみなし、演算の入力が全てスカラ値だった場合にスカラ命令が出るようです。

  1. カーネル引数
  2. コンパイル時定数
  3. __constant からのロードで、アドレスがスカラ値
  4. スカラ演算の結果

スカラ値ではない値で分岐した中で決まった値もスカラ値にはなりません。

以上を踏まえて、

こういうコードでスカラ命令を使うようにすると、

こんな感じですね。まあ今度は s_ のほうがボトルネックになっているので、この場合はあまり褒められたチューニングではないです。

命令数最小にするなら

こんな感じでしょうか。あまり考えずに色々試して出ただけなので一般化は難しいですが…。

SASSのチューニング

SASS は、アドレスのオペランドが定数オフセットしか使えないようなので、x86やGCN isaと比べるとアドレス演算の問題が表面化しやすいように見えます。

こういうコードは、

こういう感じに、使っているアドレスごとにIADDを使うループになります。

これは結構どうしようもない問題で、ポインタの数だけ加算が出てしまうのを回避できないのですが、
__shared__ 変数は、一個の配列にすることで、ポインタの数を減らすことができます。

上のプログラムなら、

IADD を一個削減できています

コメントを残す

メールアドレスが公開されることはありません。 * が付いている欄は必須項目です