Keplerから搭載されたレジスタ内のデータの入れ替え命令

ワープのスレッド間でデータを交換するシャッフル命令

ワープに含まれる各スレッドは自分専用のレジスタを使うので、他のスレッドとは無関係に処理を行っている。そのため、他のスレッドにデータを送る場合は、SMEMにそのデータを書き込み、受け取り側のスレッドはSMEMからデータを読み出すという手順が必要になる。ケース3の行列の転置もSMEMを経由して、書き込み側のスレッドから受け取り側のスレッドにデータを渡している。

しかし、SMEMはアクセスが速いといってもレジスタに比べると時間が掛る。そのため、Keplerでは、レジスタ内のデータの入れ替え(シャッフル)を行うSHFL命令が新設された。

そして右隣のスレッドのデータを受け取るという操作を、SMEMを使用し、確実にデータの読み書きが同期するよう"__syncthreads( )"をいれた場合、シャッフル命令を使用した場合、SMEMを使うが、"__syncthreads( )"を使わずvolatile指定だけで多少危険な実装の3種のカーネルを作って性能を比較してみたという。

性能測定は、右隣のデータを受け取るループを4096回実行するスレッッドを1024スレッドまとめてスレッドブロックとし、13SMのKeplerに26スレッドブロックを実行させている。つまり、1SMには2048スレッドが割り当てられ、8M回のデータの受け渡しが行われている。

この図に見られるように、SHFL命令を使うとSMEMを使った安全な受け渡しに比べて3倍以上の性能が得られ、SMEM領域も使用しないというメリットがある。

また、各スレッドが計算した結果の合計を求めるという処理は良く出てくるパターンで、このケースもSMEMを使うと6msかかる処理が、SHFL命令の使用で2msと3倍高速に実行できる。

詳細は省略するが、SHFL命令を使用すれば、行列の転置についてもSMEMは使わずに済み、60～70%程度性能を改善することができるとしている。