HMMA.884ってなんだ

CUDA Advent Calendar 2019 3日目の記事です.
妄想多めです.

HMMA.884って何?

WMMA APIで計算できる行列積の行列の大きさ(m, n, k)は(16, 16, 16)や(32, 8, 16),(8, 32, 16)などと決まっています.
一方Tensorコア自体は4x4行列積和を計算するための回路です.
そのためこのAPIでは複数または複数回のTensorコアを用いてこれらの大きさの行列積和を計算していることになります.
ではアセンブリ(SASS)レベルで4x4行列積を複数回呼んでいるかというと,そうではありません(多分).
このSASSレベルで呼んでいる命令こそがHMMA.884です.

妄想

ここまではランタイムで走っている命令を眺めるとわかるのですが,この命令が実際に何をやているかは妄想するしかありません.
じつは同じようなことを考えている人はいて,arXivに論文を出しています.
Zhe Jia, Marco Maggioni, Benjamin Staiger, Daniele P. Scarpazza - Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking arXiv:1804.06826
41ページからがTensorコアのSASSについての記述です.
が,本当にこれが正しいのか不思議でたまりません.
彼らの説明では884という数字の意味がよくわかりません.
加えてPTX ISA 6.4でPTXのmma命令のm8n8k4というShapeが登場したこと,他のm16n16k16等のShapeがm8n8k4を複数回呼べば計算できることを考えると,HMMA.884はm8n8k4を計算するための命令で,その他のShapeではこれを複数回呼んでいるのではないかと妄想してしまうのです.
上の論文ではm16n16k16はHMMA.884を16回呼んでいるとのことでいるとのことでしたが,m8n8k4を16回呼ぶとm16n16k16が計算できる気もしますよね.
全部妄想です,忘れてください.

おわり

中の人教えて

追記 .0

mma.m8n8k4命令のニモニックはCUDA TOOLKIT DOCUMENTATIONによると

mma.sync.aligned.m8n8k4.alayout.blayout.dtype.f16.f16.ctype d, a, b, c;
mma.sync.aligned.m16n8k8.row.col.dtype.f16.f16.ctype d, a, b, c;

.alayout = {.row, .col};
.blayout = {.row, .col};
.ctype   = {.f16, .f32};
.dtype   = {.f16, .f32};
なんですけど,dtypeとctypeに囲まれているf16.f16がatype,btypeと仮定するなら将来的にf16以外もサポートできるようにあえてニモニック中にA,Bの型を書いているんですかね.
期待してしまいます.

追記 .1

上の妄想ってmma.m8n8k4のSASSを見れば確認できることですよね.
こんなコードをnvccでオブジェクトファイルにしてcuobjdumpでSASSを覗きます.
asmでmovしているのは最適化で命令が消されないようにするためです.

で,こいつのSASSはこんなかんじ.
mma.m8n8k4はsm_70ようにチューニングされているとのことだったのでsm_70です.

HMMA.884.F16.F16.STEP0とHMMA.884.F16.F16.STEP1に分かれていますが,概ね妄想どおりということでしょうか?
実は上の論文ではm16n16k16でHMMA.884はSTEP0-3の4つあるとのことですが,私は0,1の2つしか確認できたことがないんですよね...

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