[PATCH] D144911: adding bf16 support to NVPTX

[Kushan Ahmadian via Phabricator via llvm-commits]

kushanam added a comment.

In D144911#4389187 <https://reviews.llvm.org/D144911#4389187>, @manishucsd wrote:

> 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]                                        ^
>
>   

Could you try again, I rebased and refactored the code


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D144911