[all-commits] [llvm/llvm-project] d03f35: [MLIR][NVVM] Fix the datatype error for nvvm.mma.s...

[xiaoleis-nv via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: d03f35f9b6d031d6a9375d90ccf7cc285f8e4b79
      https://github.com/llvm/llvm-project/commit/d03f35f9b6d031d6a9375d90ccf7cc285f8e4b79
  Author: xiaoleis-nv <99947620+xiaoleis-nv at users.noreply.github.com>
  Date:   2025-01-13 (Mon, 13 Jan 2025)

  Changed paths:
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    M mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
    M mlir/test/Dialect/LLVMIR/nvvm.mlir
    M mlir/test/Target/LLVMIR/nvvmir.mlir

  Log Message:
  -----------
  [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 at nvidia.com>



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications