[clang] 922d95a - [CIR][AMDGPU] Add lowering for amdgcn div fmas builtins (#194334)

[via cfe-commits]

Author: Chaitanya
Date: 2026-04-28T09:40:02+05:30
New Revision: 922d95aefd7a09661e546a1790285b20fe70719a

URL: https://github.com/llvm/llvm-project/commit/922d95aefd7a09661e546a1790285b20fe70719a
DIFF: https://github.com/llvm/llvm-project/commit/922d95aefd7a09661e546a1790285b20fe70719a.diff

LOG: [CIR][AMDGPU] Add lowering for amdgcn div fmas builtins (#194334)

Upstreaming ClangIR PR: https://github.com/llvm/clangir/pull/2051

This PR adds support for lowering of _builtin_amdgcn_div_fmas* amdgpu
builtins to clangIR.
Followed similar lowering from reference clang->llvmir in
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp.

Added: 
    

Modified: 
    clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
    clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index e2da5b6e5d88b..186c9d3cf47e7 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -78,10 +78,16 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
   }
   case AMDGPU::BI__builtin_amdgcn_div_fmas:
   case AMDGPU::BI__builtin_amdgcn_div_fmasf: {
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented AMDGPU builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
+    mlir::Value src0 = emitScalarExpr(expr->getArg(0));
+    mlir::Value src1 = emitScalarExpr(expr->getArg(1));
+    mlir::Value src2 = emitScalarExpr(expr->getArg(2));
+    mlir::Value src3 = emitScalarExpr(expr->getArg(3));
+    mlir::Value result = cir::LLVMIntrinsicCallOp::create(
+                             builder, getLoc(expr->getExprLoc()),
+                             builder.getStringAttr("amdgcn.div.fmas"),
+                             src0.getType(), {src0, src1, src2, src3})
+                             .getResult();
+    return result;
   }
   case AMDGPU::BI__builtin_amdgcn_ds_swizzle:
   case AMDGPU::BI__builtin_amdgcn_mov_dpp8:

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index fafb94be25353..d374479e6182e 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -47,3 +47,19 @@ __device__ void test_div_scale_f32_with_ptr(float* out, int* flagout, bool* flag
 {
   *out = __builtin_amdgcn_div_scalef(a, b, true, flag);
 }
+
+// CIR-LABEL: @_Z17test_div_fmas_f32Pdfffi
+// CIR: cir.call_llvm_intrinsic "amdgcn.div.fmas" {{.*}} : (!cir.float, !cir.float, !cir.float, !cir.bool) -> !cir.float
+// LLVM: define{{.*}} void @_Z17test_div_fmas_f32Pdfffi
+// LLVM: call{{.*}} float @llvm.amdgcn.div.fmas.f32(float %{{.+}}, float %{{.+}}, float %{{.+}}, i1 %{{.*}})
+__device__ void test_div_fmas_f32(double* out, float a, float b, float c, int d) {
+  *out = __builtin_amdgcn_div_fmasf(a, b, c, d);
+}
+
+// CIR-LABEL: @_Z17test_div_fmas_f64Pddddi
+// CIR: cir.call_llvm_intrinsic "amdgcn.div.fmas" {{.*}} : (!cir.double, !cir.double, !cir.double, !cir.bool) -> !cir.double
+// LLVM: define{{.*}} void @_Z17test_div_fmas_f64Pddddi
+// LLVM: call{{.*}} double @llvm.amdgcn.div.fmas.f64(double %{{.+}}, double %{{.+}}, double %{{.+}}, i1 %{{.*}})
+__device__ void test_div_fmas_f64(double* out, double a, double b, double c, int d) {
+  *out = __builtin_amdgcn_div_fmas(a, b, c, d);
+}