CUDA の CRS spmv についてのメモ
Efficient Sparse Matrix-Vector Multiplication on CUDA | NVIDIA
の論文の、CRS spmv vector のメモ。
CRS spmv の場合、1行に複数のスレッドが計算したほうが効率がよい。
しかしある行で各スレッドが計算した結果を全て足し算 (parallel reduction) をする必要がある。
普通だったら、複数スレッドの reduction なので、スレッド間の同期が必要なのだが、上記の論文では half warp 単位でスレッドが動作するから、warp のスレッド数までだったら暗黙の同期が入るので明示的に同期を取る必要がない、という話。*1
なので、上の論文に記載されているカーネルコードのようにスレッド間の同期を明示的に書いてなくても問題ない。
蛇足というか思っていること
上のカーネルコード、グローバルのスレッド ID を計算して、そこから lane と row を計算しているんだけど、これってそれぞれ threadIdx.x , blockIdx.x でいいような気がするんだが... 何かあるんだろうか。
