[CUDA-Samples] immaTensorCoreGemm

CUDA Advent Calendar 2019 14日目の記事です.
ネタのつなぎのCUDA Samples読み回4回目です.
読んでいたもののバージョンが変わってしまうといけないのでこちらにforkしてあります.

何このサンプル

TensorコアはTuringからINT8やINT4に対応したわけですが,これの使い方に関するサンプルです.
IMMAというのはInteger Matrix Multiply and Accumulateの頭文字です.
といってもfragmentなど使い方はHMMAとほとんど同じです(サンプルも[CUDA-Samples] cudaTensorCoreGemmとほぼ同じだと思います).
IMMAでも足しこみ精度が高くなるように作られており,このサンプルだとA,Bは

wmma::fragment<wmma::matrix_a, M, N, K, uint8_t, wmma::row_major>
wmma::fragment<wmma::matrix_ab M, N, K, uint8_t, wmma::col_major>
で定義されており,Cは
wmma::fragment<wmma::accumulator, M, N, K, int>
となっています.

IMMAはSASSだとどうなるか?

HMMA.884って何だ? - 天炉48町で半精度のTensorコアの命令はHMMA.884.**となることを書きましたが,IMMAではどうなるのでしょうか?
やってみました. をsm_75でSASSにすると という感じです.
IMMA.8816なんですね.
足しこみ方向が16個ということでいいんですかね?
HMMA.884ではSTEP0とSTEP1に分かれていましたが,こちらは分かれないんですかね.

おわり

使う機会がぱっとは思いつきませんがいつか使ってみたいものです.

カテゴリー:CUDA
記事作成日:2019-12-14