行列の転置を行う場合のメモリアクセス

3. 行列の転置の処理を考える

ここまでは前置きで、ここからは行列の転置を行う場合のメモリアクセスを説明する。行列の転置は、次の図のように、行方向の要素を、列方向に並べ直す操作で、プログラム的に言うとA[i,j]要素をB[j,i]に移す操作である。N行N列の行列の場合、メモリアクセスの回数はO(N²)であり、演算は行わない。

最も簡単な直列的実行の実装は、次の図のように、もとの行列inの要素を列方向に読み、それを出力の行列outの行方向に格納するというもので、これを2重ループで廻す。

しかし、これではN²回の読み込みと書き込みが必要である。これをCUDAで並列化したプログラムを次の図に示す。このプログラムでは、ループを使って入力行列inの1列のデータを順に読み、これを転置して出力行列outに格納して行く。1つのスレッドでは、1列のデータを転置するだけであるが、

i=blockIdx.x * blockDim.x + threadIdx.x

で処理する列番号を決めているので、列の総数のスレッドを並列に実行させれば、行列全体の転置が行われる。

原理的に、列数だけのスレッドが並列に実行されるので、その分、性能が高くなる。ただし、Nが非常に大きくて、GPUのリソースが不足する場合は、ハードウェアとして実行できるサイズに処理を分割して、分割したものを順に実行して行くということになる。

次のプログラムのグリッドは1次元で、ブロックも1次元の配列で、1ブロックのスレッド数は256である。

K40c GPUを875MHzクロックにブーストして、このプログラムを実行した場合の性能は、データが単精度(SP)の場合は16.96GB/s、倍精度(DP)の場合は33.94GB/sであった。なお、ECCはオフの状態で実行している。

実行時間の測定には、OpenMPのタイマを用いる方法と、CUDAのイベント時刻を測る方法がある。OpenMPの場合は、

double start    = omp_get_wtime( ); 開始時刻を変数startに格納する
ここで測定すべき操作を実行
double end = omp_get_wtime( );  終了時刻をendに格納する
double duration = end - start;      経過時間(秒)

CUDAのタイマを使う場合は、

cudaEventRecord(start);     開始時刻を変数startに格納する
ここで測定すべき操作を実行
cudaEventRecord(end);   終了時刻を変数endに格納する
cudaEvnetSynchronize(end);  イベントを同期する
cudaEventGetElapsedTime(&time_ms,start,end);    startからendまでの時間
    ms単位

となる。

なお、このスライドのプログラムはちょっといい加減で、OpenMPの方の終了時間の測定の関数呼び出しがomp_get_timeになっている。また、終了時刻はendに格納しているのに、cudaEventGetElapsedTimeの最後の引数がstopとなっていて、これでは正しい経過時間が求められない。

性能を上げるチューニング行う場合には、次の事項を把握しておくと役に立つ。理論的に、その計算を行なうのに、最低何回のLoadが必要か。最低何回の浮動小数点演算が必要か。これが分かれば、Byte/sec、Flops/secが計算できる。

理論的なロード回数や演算数の下限が分からない場合は、毎秒何回のアトミックオペレーションが実行されるか、毎秒何回の内積計算が行われるかなど主要なマクロオペレーションの実行回数を把握しておくと良い。

メモリのスループットを示すカウンタがあるが、その使用には注意が必要な点がある。dram_read_throughputはDRAMのリードスループット、l2_l1_read_throughputはl1キャッシュからのリクエストに対するl2キャッシュのリードスループット、l2_tex_read_throughputは、textureキャッシュからのリクエストに対するl2キャッシュのリードスループットである。ここで、注意すべきは、これらのカウンタの値は、l2キャッシュやDRAMから見たものであり、読まれたデータが有効に使われるものであるかどうかは考慮されていない。たとえば、キャッシュミスになったアクセス、読まれたが、使われなかったデータのバイト数、ECCのオーバヘッドのバイトなどもスループットに含まれている。

このため、これらのカウンタが真の性能を表しているとは限らないので注意が必要である。

なお、忍者マークのついているスライドは上級者向けである。