[PATCH] D144911: adding bf16 support to NVPTX

Manish Gupta via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Thu Jun 1 16:05:11 PDT 2023


manishucsd added a comment.

I tried this patch to resolve lowering `nvvm.mma.sync` on bf16 operands for the PR 13858 <https://github.com/openxla/iree/pull/13858>

  nvvm.mma.sync A[%293, %295, %297, %299]  B[%392, %394]  C[%484, %485, %487, %488]  {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<bf16>, multiplicandBPtxType = #nvvm.mma_type<bf16>, shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (vector<2xbf16>, vector<2xbf16>, f32) -> !llvm.struct<(f32, f32, f32, f32)>

I fail to compile this patch. Please find the compilation error below:

  [build] ./llvm-project/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td:1117:40: error: Variable not defined: 'hasPTX70'
  [build]                 Requires<[useFP16Math, hasPTX70, hasSM80, Pred]>;
  [build]                                        ^

  


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D144911/new/

https://reviews.llvm.org/D144911



More information about the cfe-commits mailing list