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

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

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

シャッフル命令には任意の入れ替え、右、左にN箇所移動、1つ置きのペアで入れ替えの4種のバリエーションがある。この図ではa-hの8個のデータであるが、実際は32スレッド分のデータがある (3)

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

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

"__syncthreads( )"付でSMEM使用，SHFL命令を使用、"__syncthreads( )"なしでSMEM使用の3種の右隣のデータを受け取るコード (3)

左がms単位の実行時間，右がSMEMの使用量。SMEMを使ったスレッド間のデータの受け渡しは約1.3msに対して、SHFL命令を使うと0.4ms弱で終わり、3倍あまり高速になっている。また、SMEMを使うと4KBを必要とするが、SHFL命令の場合はSMEMは不要である (3)

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

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

ワープ内の全スレッドの処理結果の合計を求める処理 (3)。SMEMを使う場合に比べて、SHFL命令を使って、順に2個、4個、8個、…の処理結果の合計を求めるという処理を行うと3倍性能が向上する

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

行列の転置の性能比較 (3)。SMEMを使用すると7.5ms程度の処理がSHFL命令を使用すると4.5ms程度で終わる