diff options
| author | Artem Belevich <tra@google.com> | 2019-10-22 14:07:15 -0700 |
|---|---|---|
| committer | Artem Belevich <tra@google.com> | 2019-10-28 13:55:30 -0700 |
| commit | d9972f848294b06807c8764615852ba2bc1e8a74 (patch) | |
| tree | 41a29fa0aaafbdb271496f9ba6a1f7682d313286 /llvm/lib/Target/NVPTX | |
| parent | 7b9d6a22569c69c230b1584829e3acbdf852b221 (diff) | |
| download | bcm5719-llvm-d9972f848294b06807c8764615852ba2bc1e8a74.tar.gz bcm5719-llvm-d9972f848294b06807c8764615852ba2bc1e8a74.zip | |
[NVPTX] Added llvm.nvvm.mma.m8n8k4.* intrinsics
Differential Revision: https://reviews.llvm.org/D69324
Diffstat (limited to 'llvm/lib/Target/NVPTX')
| -rw-r--r-- | llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 42 |
1 files changed, 28 insertions, 14 deletions
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index c52195fb044..76a4a1d4030 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -7400,7 +7400,9 @@ class WMMA_REGINFO<WMMA_REGS r> // u4/s4/b1 -> s32 @ m8n8k32 (u4/s4), m8n8k128(b1) !or(!eq(geom,"m8n8k128"), - !eq(geom,"m8n8k32")) : [hasSM75, hasPTX63]); + !eq(geom,"m8n8k32")) : [hasSM75, hasPTX63], + + !eq(geom, "m8n8k4") : [hasSM70, hasPTX64]); // template DAGs for instruction inputs/output. dag Outs = !dag(outs, ptx_regs, reg_names); @@ -7546,25 +7548,37 @@ class WMMA_MMA<WMMA_REGINFO FragA, WMMA_REGINFO FragB, let OutOperandList = FragD.Outs; let InOperandList = !con(Args, (ins MmaCode:$ptx)); string TypeList = !cond( + !eq(FragD.geom, "m8n8k4") : "." # FragD.ptx_elt_type + # ".f16.f16." + # FragC.ptx_elt_type, !eq(FragD.ptx_elt_type, "s32") : ".s32" # "." # FragA.ptx_elt_type # "." # FragB.ptx_elt_type # ".s32", 1: "." # FragD.ptx_elt_type # "." # FragC.ptx_elt_type, ); - let AsmString = "wmma.mma" - # !if(!eq(FragA.ptx_elt_type, "b1"), ".xor.popc", "") - # ".sync" - # "${ptx:aligned}" - # "." # ALayout - # "." # BLayout - # "." # FragA.geom - # TypeList - # !if(Satfinite, ".satfinite", "") # "\n\t\t" - # FragD.regstring # ",\n\t\t" - # FragA.regstring # ",\n\t\t" - # FragB.regstring # ",\n\t\t" - # FragC.regstring # ";"; + let AsmString = !if(!eq(FragA.geom, "m8n8k4"), + "mma.sync.aligned.m8n8k4" + # "." # ALayout + # "." # BLayout + # TypeList # "\n\t\t" + # FragD.regstring # ",\n\t\t" + # FragA.regstring # ",\n\t\t" + # FragB.regstring # ",\n\t\t" + # FragC.regstring # ";", + "wmma.mma" + # !if(!eq(FragA.ptx_elt_type, "b1"), ".xor.popc", "") + # ".sync" + # "${ptx:aligned}" + # "." # ALayout + # "." # BLayout + # "." # FragA.geom + # TypeList + # !if(Satfinite, ".satfinite", "") # "\n\t\t" + # FragD.regstring # ",\n\t\t" + # FragA.regstring # ",\n\t\t" + # FragB.regstring # ",\n\t\t" + # FragC.regstring # ";"); } defset list<WMMA_INSTR> MMAs = { |
