[CIR][AMDGPU] Add lowering for amdgcn div fmas builtins (#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.

Author: skc7

clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp

@@ -116,7 +116,16 @@ mlir::Value CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
   }
   case AMDGPU::BI__builtin_amdgcn_div_fmas:
   case AMDGPU::BI__builtin_amdgcn_div_fmasf: {
-    llvm_unreachable("div_fmas_* NYI");
+    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 =
+        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:

clang/test/CIR/CodeGen/HIP/builtins-amdgcn.hip

@@ -255,3 +255,23 @@ __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.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 %{{.*}})
+// OGCG: define{{.*}} void @_Z17test_div_fmas_f32Pdfffi
+// OGCG: 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.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 %{{.*}})
+// OGCG: define{{.*}} void @_Z17test_div_fmas_f64Pddddi
+// OGCG: 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);
+}

clang/test/CIR/CodeGen/OpenCL/builtins_amdgcn.cl

@@ -266,3 +266,25 @@ void test_div_scale_f32_generic_ptr(global float* out, global int* flagout, floa
 {
   *out = __builtin_amdgcn_div_scalef(a, b, true, flag);
 }
+
+// CIR-LABEL: @test_div_fmas_f32
+// CIR: cir.llvm.intrinsic "amdgcn.div.fmas" {{.*}} : (!cir.float, !cir.float, !cir.float, !cir.bool) -> !cir.float
+// LLVM: define{{.*}} void @test_div_fmas_f32
+// LLVM: call float @llvm.amdgcn.div.fmas.f32(float %{{.+}}, float %{{.+}}, float %{{.+}}, i1 %{{.*}})
+// OGCG: define{{.*}} void @test_div_fmas_f32
+// OGCG: call float @llvm.amdgcn.div.fmas.f32(float %{{.+}}, float %{{.+}}, float %{{.+}}, i1 %{{.*}})
+void test_div_fmas_f32(global float* out, float a, float b, float c, int d)
+{
+  *out = __builtin_amdgcn_div_fmasf(a, b, c, d);
+}
+
+// CIR-LABEL: @test_div_fmas_f64
+// CIR: cir.llvm.intrinsic "amdgcn.div.fmas" {{.*}} : (!cir.double, !cir.double, !cir.double, !cir.bool) -> !cir.double
+// LLVM: define{{.*}} void @test_div_fmas_f64
+// LLVM: call double @llvm.amdgcn.div.fmas.f64(double %{{.+}}, double %{{.+}}, double %{{.+}}, i1 %{{.*}})
+// OGCG: define{{.*}} void @test_div_fmas_f64
+// OGCG: call double @llvm.amdgcn.div.fmas.f64(double %{{.+}}, double %{{.+}}, double %{{.+}}, i1 %{{.*}})
+void test_div_fmas_f64(global double* out, double a, double b, double c, int d)
+{
+  *out = __builtin_amdgcn_div_fmas(a, b, c, d);
+}