グリッドを大きくする方法

5.カーネルの並列度を上げる

NVVPはグリッドが小さすぎると言っている。これは並列度が低いという指摘である。並列度が低いと、飛行中のC5が少なく、バンド幅を有効に利用できない。

どのようにすればグリッドを大きくできるのであろうか? これには、スレッドあたりの仕事量を減らし、スレッド数を増やせば良い。前にあげた例では1つのループを並列化しただけであるが、2つのループを並列化してみよう。

グリッドとブロックを2次元にしてスレッド数を増やした転置のプログラムは次のようになる。

int i, j;
i = blockIdx.x * blockDim.x + threadIdx.x;
j = blockIdx.y * blockDim.y + threadIdx.y;
out[i * rows + j] = in[j * cols + i];

ここで、rowsとcolsは入力の行列inの行数と列数である。このコードはinのi,j要素を読み、outのj,i要素に格納するという、1つの要素をだけを処理するプログラムとなっている。

次の表はオリジナルの1次元のtranspose1とこの2次元のtranspose2プログラムを実行した場合に性能を比較したものである。transpose2では、単精度(Float)の場合は76.23GB/sとなり、これはtranspose1の4.5倍の性能である。また、倍精度(Double)の場合は137.878GB/sとなり、これは4.06倍の性能改善となっている。

前と同様に資源の利用率を見てみると、メモリ系の占有率は75%程度に向上している。しかし、Compute側の占有率は27%程度と低く、メモリシステムのバンド幅が制約になっている可能性が高いと指摘されている。