シェアードメモリ中での転置で、連続アドレスへの書き込みを実現

8. 行列をタイルに分割して転置する

行列全体を2次元のタイル(例えば32×32のサイズ)に分割して考えて、1つタイルをSMに内蔵されている高速アクセスができるローカルなメモリに格納する。このメモリに格納されたタイルの中の要素を列方向に読み、行方向に書き出すのは高速で実行できる。GPUでは、プログラマがアクセスを制御でき、ローカルで高速にアクセスできるメモリとしてシェアードメモリが設けられている。このシェアードメモリの中で転置を行えば、連続アドレスへの書き込みとなり、書き込みの合体が使えるようになる。

このシェアードメモリを使う転置を図示すると、次の図のようになる。グローバルメモリからシェアードメモリへの転送は連続アドレスであり、メモリアクセス要求の合体が使える。

このシェアードメモリに格納されたタイルを、行番号と列番号を入れ替えて読み出すと転置が行える。この読み出しを順にグローバルメモリの連続アドレスに書き出すと、転置された結果が書き込まれる。この書き出しは連続アドレスであるので、ここでもメモリアクセスの合体が有効に働く。

このやり方のコードを次の図に示す。__shared__ float tile[TILE_DI][TILE_DIM];でタイルを格納する領域をシェアードメモリに定義する。そして入力の行列をinから読み込み、synchthreads( );ですべての読み込みが終わったことを確認する。そして、tileを行と列の番号を逆にして読み出してグローバルメモリのoutに書き出している。

このtranspose3プログラムの性能を次の図に示す。単精度(Float)の場合は、101.03GB/sの性能が得られ、これはtranspose2プログラムの1.32倍の性能となっている。しかし、倍精度(Double)の場合は、127.86GB/sとなり、transpose2と比較して0.92倍と性能が低下している。

倍精度では性能が低下してしまったので、NVVPに戻って原因を究明する。その結果、84行目のシェアードメモリの読み出しが、アクセス当たりのトランザクションが16回と多いのが原因と言う指摘である。

シェアードメモリは、SMの中にあるメモリで、レジスタファイルと同程度という高速のアクセスができる。そして、クロック当たり128バイトのバンド幅を持っている。そして、32バンク構成となっており、1クロックで32個の4バイトデータをアクセスすることができるようになっている。

シェアードメモリがどのように働くかを考えてみよう。なお、次の図は、スライドに収めるために、シェアードメモリは4バンクで、ワープも4スレッドとして書いている。

行列inを読み込んだとき、シェアードメモリのTile領域への要素の格納は、次の図のようになっている。バンク0にはtile[*][0]、バンク1にはtile[*][1]、バンク2にはtile[*][2]、バンク3にはtile[*][3]が入っている。

この状態で、転置した行を作るため、tile[0][0]、tile[1][0]、tile[2][0]、tile[3][0]を読み出そうとすると、すべてのアクセスがバンク0に集中してしまう。

シェアードメモリは1クロックで32個の4バイトデータを読み出すことができるが、それはアクセスがすべてのバンクに分散している場合で、同一バンクにアクセスが集中した場合は、それらはreplayで順に読み出される事になる。そして、replayを行っている間は、シェアードメモリのバンド幅のごく一部しか利用されないことになっている。

なお、同一バンクでも複数のスレッドが同じアドレスのデータを読むのは問題ないし、ワープが異なれば、同じアドレスをアクセスしても問題ない。バンクの競合が問題となるのは、1つのワープ内のスレッドの間だけである。

Replayを行っている間はワープの実行は進まず、メモリアクセス命令の発行も止まってしまう。このため、in-flightのメモリアクセス要求の数も減ってしまい、メモリの利用率も低くなってしまう。

このプログラムでは、行方向のアクセスは進行するが、列方向のアクセスではreplayが発生して実行が遅くなってしまう。これはシェアードメモリを使わないtranspose2の場合と同じように思うかも知れないが、transpose2では多数のキャッシュラインへのアクセスを処理するのにreplayが発生するのにたいして、こちらはシェアードメモリのバンクコンフリクトでreplayが発生しており、この2つはまったく異なる問題である。