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つの要素をだけを処理するプログラムとなっている。

スレッドブロックとグリッドを2次元にして、スレッド数を増やす。そのため、各スレッドは1つの要素を処理するだけになる

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

2次元のグリッドとブロックとしたことにより、単精度では4.5倍、倍精度では4.06倍に性能が向上した

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

2次元版transposeの資源利用率。Compute側の占有率が低く、メモリバンド幅が制約になっているという指摘である