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

[via cfe-commits]

https://github.com/skc7 created https://github.com/llvm/llvm-project/pull/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.

>From 0332daec6121ae89dfab5d6be75d3e5b7dd68c74 Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Mon, 27 Apr 2026 15:47:38 +0530
Subject: [PATCH] [CIR][AMDGPU] Add lowering for amdgcn div fmas builtins

---
 clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 14 ++++++++++----
 clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 16 ++++++++++++++++
 2 files changed, 26 insertions(+), 4 deletions(-)

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);
+}