[MLIR] Mark LLVM::FMAOp as legal

[chelini]

Mark LLVM::FMAOp as legal in configureGpuToNVVMConversionLegality, since we can handle intrinsic lowering in the NVPTX backend and emit fma.rn.f32.

[llvmbot]

@llvm/pr-subscribers-mlir

@llvm/pr-subscribers-mlir-gpu

Author: lorenzo chelini (chelini)

Changes

---

Full diff: https://github.com/llvm/llvm-project/pull/144671.diff

2 Files Affected:

• (modified) mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp (+1-1)
• (modified) mlir/test/Integration/GPU/CUDA/dump-ptx.mlir (+12-1)

diff --git a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
index 958d0d085fce1..ac264f4056bba 100644
--- a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
+++ b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
@@ -429,7 +429,7 @@ void mlir::configureGpuToNVVMConversionLegality(ConversionTarget &target) {
   target.addLegalDialect<::mlir::NVVM::NVVMDialect>();
   target.addIllegalDialect<gpu::GPUDialect>();
   target.addIllegalOp<LLVM::CopySignOp, LLVM::CosOp, LLVM::ExpOp, LLVM::Exp2Op,
-                      LLVM::FAbsOp, LLVM::FCeilOp, LLVM::FFloorOp, LLVM::FMAOp,
+                      LLVM::FAbsOp, LLVM::FCeilOp, LLVM::FFloorOp,
                       LLVM::FRemOp, LLVM::LogOp, LLVM::Log10Op, LLVM::Log2Op,
                       LLVM::PowOp, LLVM::RoundEvenOp, LLVM::RoundOp,
                       LLVM::SinOp, LLVM::SqrtOp>();
diff --git a/mlir/test/Integration/GPU/CUDA/dump-ptx.mlir b/mlir/test/Integration/GPU/CUDA/dump-ptx.mlir
index 0cc5d8645bb36..fdc53ff0c4160 100644
--- a/mlir/test/Integration/GPU/CUDA/dump-ptx.mlir
+++ b/mlir/test/Integration/GPU/CUDA/dump-ptx.mlir
@@ -2,7 +2,7 @@
 // RUN:  | mlir-opt -gpu-lower-to-nvvm-pipeline -debug-only=serialize-to-isa \
 // RUN:  2>&1 | FileCheck %s
 
-// CHECK: Generated by LLVM NVPTX Back-End
+// CHECK-LABEL: Generated by LLVM NVPTX Back-End
 // CHECK: .visible .func kernel_a()
 // CHECK: ret;
 gpu.module @bar {
@@ -11,3 +11,14 @@ gpu.module @bar {
     llvm.return
   }
 }
+
+// CHECK-LABEL: Generated by LLVM NVPTX Back-End
+// CHECK: fma.rn.f32
+
+gpu.module @foo {
+  llvm.func @fma(%arg0: f32, %arg1: f32) -> f32
+    attributes { gpu.kernel } {
+    %res = llvm.intr.fma (%arg0, %arg1, %arg1) : (f32, f32, f32) -> f32
+    llvm.return %res : f32
+  }
+}

[github-actions]

✅ With the latest revision this PR passed the C/C++ code formatter.

[schwarzschild-radius]

The change looks good to me, Thanks!