diff options
| author | xiaoleis-nv <99947620+xiaoleis-nv@users.noreply.github.com> | 2025-01-13 17:33:05 +0800 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2025-01-13 15:03:05 +0530 |
| commit | d03f35f9b6d031d6a9375d90ccf7cc285f8e4b79 (patch) | |
| tree | 92e4b0123c5915e28e70512693e1bcb9885553e3 /llvm/lib/Target/RISCV/Disassembler/RISCVDisassembler.cpp | |
| parent | 1b199d19902a752433c397377567ff381261e94a (diff) | |
| download | llvm-d03f35f9b6d031d6a9375d90ccf7cc285f8e4b79.zip llvm-d03f35f9b6d031d6a9375d90ccf7cc285f8e4b79.tar.gz llvm-d03f35f9b6d031d6a9375d90ccf7cc285f8e4b79.tar.bz2 | |
[MLIR][NVVM] Fix the datatype error for nvvm.mma.sync when the operand is bf16 (#122664)
The PR fixes the datatype error for `nvvm.mma.sync` when the operand is
`bf16`. This operation originally requires the A/B type to be `f16x2`
for the `bf16` MMA. However, it violates the NVVM intrinsic
[[here](https://github.com/xiaoleis-nv/llvm-project/blob/372044ee09d39942925824f8f335aef40bfe92f0/llvm/include/llvm/IR/IntrinsicsNVVM.td#L119)],
where the A/B operand type should be `i32`. This is a bug, and there are
no tests in MLIR that cover this datatype.
```
// mma bf16 -> s32 @ m16n8k16/m16n8k8
!eq(gft,"m16n8k16:a:bf16") : !listsplat(llvm_i32_ty, 4),
!eq(gft,"m16n8k16:b:bf16") : !listsplat(llvm_i32_ty, 2),
!eq(gft,"m16n8k8:a:bf16") : !listsplat(llvm_i32_ty, 2),
!eq(gft,"m16n8k8:b:bf16") : [llvm_i32_ty],
```
This PR addresses this bug and adds tests to guarantee correctness.
Co-authored-by: Xiaolei Shi <xiaoleis@nvidia.com>
Diffstat (limited to 'llvm/lib/Target/RISCV/Disassembler/RISCVDisassembler.cpp')
0 files changed, 0 insertions, 0 deletions