行列の転置では書き込みは飛び飛びアドレスになる

7. 転置の場合、列方向の書き込みは飛び飛びアドレスになってしまう

行列の転置の場合、読み込みを連続アドレスとすると、書き込みは飛び飛びのアドレスになってしまいそれぞれのアクセスが異なるキャッシュラインとなってしまうことは避けられないように思われるが、本当にそうであろうか?

2次元の行列は、次の図のように、各行のデータは順に並んでメモリに格納されている。

そして、1行目のデータに続いて2行目のデータ、その後に3行目のデータが続くというようにメモリに格納されている。

この並びは行方向にメモリをアクセスする場合は良いが、列方向にメモリをアクセスする場合は、次の図のようにN飛びのアクセスになってしまう。

次の図は、スレッド0-7のアクセスは第1のキャッシュライン、スレッド8-15は第2のキャッシュライン、スレッド16-23は第3のキャッシュラインをアクセスしており、スレッド24-31は、また、第1のキャッシュラインをアクセスするという状況を示しており、このアクセスの処理には3つのキャッシュラインが必要となり、2回のreplayが必要となる。

伝統的な最適化のアドバイスは、連続アクセス(Stride 1)を推奨している。次のプログラムは、iを0からn-1まで順に変化させ、data[i]を合計するものである。このように、連続アドレスをアクセスすればキャッシュヒットの可能性が高まり、同一キャッシュラインの利用回数も増加する。

しかし、GPUのスレッドから見ると、連続アドレスは少し違って見える。次の

for (int i=threadIdx.x; i<n; i+= blockDim.x) {
            r+=data[i];

というプログラムでは、iの値は、毎回、ブロックのx方向のサイズ(blockDim.x)だけ増加している。また、その後の

    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    int stride = blockDim.x*gridDim.x;
    for (int i=idx; i<n; i+=stride){
        r += data[i];

というプログラムではブロックのx方向のサイズとグリッドのx方向のサイズの積であるstrideだけ、iの値を増している。

このように飛び飛びにアクセスすることにより、DRAMの同じ領域を繰り返しアクセスすることを避けている。