llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Ayokunle Amodu (ayokunle321) <details> <summary>Changes</summary> Adds codegen support for the following AMDGPU division fixup builtins: - __builtin_amdgcn_div_fixup (double) - __builtin_amdgcn_div_fixupf (float) - __builtin_amdgcn_div_fixuph (half) These are lowered to the corresponding `llvm.amdgcn.div.fixup` intrinsic. --- Full diff: https://github.com/llvm/llvm-project/pull/197468.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+9-4) - (modified) clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip (+16) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 04ab1c29b0d63..8f5bbfe881252 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -195,10 +195,15 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_div_fixup: case AMDGPU::BI__builtin_amdgcn_div_fixupf: case AMDGPU::BI__builtin_amdgcn_div_fixuph: { - 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 result = + LLVMIntrinsicCallOp::create(builder, getLoc(expr->getExprLoc()), + builder.getStringAttr("amdgcn.div.fixup"), + src0.getType(), {src0, src1, src2}) + .getResult(); + return result; } case AMDGPU::BI__builtin_amdgcn_trig_preop: case AMDGPU::BI__builtin_amdgcn_trig_preopf: { diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 4a61fde7aa90c..715c431fd113e 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -71,3 +71,19 @@ __device__ void test_div_fmas_f64(double* out, double a, double b, double c, int __device__ void test_ds_swizzle(int* out, int a) { *out = __builtin_amdgcn_ds_swizzle(a, 32); } + +// CIR-LABEL: @_Z18test_div_fixup_f32Pffff +// CIR: cir.call_llvm_intrinsic "amdgcn.div.fixup" {{.*}} : (!cir.float, !cir.float, !cir.float) -> !cir.float +// LLVM: define{{.*}} void @_Z18test_div_fixup_f32Pffff +// LLVM: call{{.*}} float @llvm.amdgcn.div.fixup.f32(float %{{.+}}, float %{{.+}}, float %{{.+}}) +__device__ void test_div_fixup_f32(float* out, float a, float b, float c) { + *out = __builtin_amdgcn_div_fixupf(a, b, c); +} + +// CIR-LABEL: @_Z18test_div_fixup_f64Pdddd +// CIR: cir.call_llvm_intrinsic "amdgcn.div.fixup" {{.*}} : (!cir.double, !cir.double, !cir.double) -> !cir.double +// LLVM: define{{.*}} void @_Z18test_div_fixup_f64Pdddd +// LLVM: call{{.*}} double @llvm.amdgcn.div.fixup.f64(double %{{.+}}, double %{{.+}}, double %{{.+}}) +__device__ void test_div_fixup_f64(double* out, double a, double b, double c) { + *out = __builtin_amdgcn_div_fixup(a, b, c); +} `````````` </details> https://github.com/llvm/llvm-project/pull/197468 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
