https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/222204
>From aead53c49f3d3104b0e62a5550352269212778f7 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Wed, 9 Sep 2026 02:01:43 +0200 Subject: [PATCH 1/2] add lerp support and route swizzle through generic path --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 9 --------- clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 6 ------ clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 15 +++++++++++++++ 3 files changed, 15 insertions(+), 15 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 1c23ea142bf8d..2ae78397e6b70 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -203,9 +203,6 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_div_fmasf: return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.div.fmas") .getValue(); - case AMDGPU::BI__builtin_amdgcn_ds_swizzle: - return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.ds.swizzle") - .getValue(); case AMDGPU::BI__builtin_amdgcn_mov_dpp8: case AMDGPU::BI__builtin_amdgcn_mov_dpp: case AMDGPU::BI__builtin_amdgcn_update_dpp: { @@ -340,12 +337,6 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, getContext().BuiltinInfo.getName(builtinId)); return mlir::Value{}; } - case AMDGPU::BI__builtin_amdgcn_lerp: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } case AMDGPU::BI__builtin_amdgcn_ubfe: { cgm.errorNYI(expr->getSourceRange(), std::string("unimplemented AMDGPU builtin call: ") + diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 53e9315c85539..1b73bfe4decb8 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -630,9 +630,6 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, return Builder.CreateCall(F, {Src0, Src1, Src2, Src3ToBool}); } - case AMDGPU::BI__builtin_amdgcn_ds_swizzle: - return emitBuiltinWithOneOverloadedType<2>(*this, E, - Intrinsic::amdgcn_ds_swizzle); case AMDGPU::BI__builtin_amdgcn_mov_dpp8: case AMDGPU::BI__builtin_amdgcn_mov_dpp: case AMDGPU::BI__builtin_amdgcn_update_dpp: { @@ -783,9 +780,6 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, case AMDGPU::BI__builtin_amdgcn_fracth: return emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::amdgcn_fract); - case AMDGPU::BI__builtin_amdgcn_lerp: - return emitBuiltinWithOneOverloadedType<3>(*this, E, - Intrinsic::amdgcn_lerp); case AMDGPU::BI__builtin_amdgcn_ubfe: return emitBuiltinWithOneOverloadedType<3>(*this, E, Intrinsic::amdgcn_ubfe); diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 9bbfb6747cbf8..91a7163307140 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -331,3 +331,18 @@ __device__ void test_class_f32(bool* out, float a, int b) { __device__ void test_class_f64(bool* out, double a, int b) { *out = __builtin_amdgcn_class(a, b); } + +// __builtin_amdgcn_lerp, like __builtin_amdgcn_ds_swizzle above, is claimed by +// the generic intrinsic path in emitBuiltinExpr: its intrinsic carries +// ClangBuiltin<> and takes no overloaded type, so getIntrinsicForClangBuiltin +// resolves it. Neither needs a case in emitAMDGPUBuiltinExpr; these tests pin +// the emission down. + +// CIR-LABEL: @_Z9test_lerpPjjjj +// CIR: cir.call_llvm_intrinsic "amdgcn.lerp" {{.*}} : (!u32i, !u32i, !u32i) -> !u32i +// LLVM: define{{.*}} void @_Z9test_lerpPjjjj +// LLVM: call{{.*}} i32 @llvm.amdgcn.lerp(i32 %{{.+}}, i32 %{{.+}}, i32 %{{.+}}) +__device__ void test_lerp(unsigned int* out, unsigned int src0, unsigned int src1, + unsigned int src2) { + *out = __builtin_amdgcn_lerp(src0, src1, src2); +} >From 0b472f8322cac52bc7b7550eedad7cb7d60b6f9d Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Wed, 9 Sep 2026 02:11:38 +0200 Subject: [PATCH 2/2] remove unneccessary comment --- clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 6 ------ 1 file changed, 6 deletions(-) diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 91a7163307140..ade593ce0d9bb 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -332,12 +332,6 @@ __device__ void test_class_f64(bool* out, double a, int b) { *out = __builtin_amdgcn_class(a, b); } -// __builtin_amdgcn_lerp, like __builtin_amdgcn_ds_swizzle above, is claimed by -// the generic intrinsic path in emitBuiltinExpr: its intrinsic carries -// ClangBuiltin<> and takes no overloaded type, so getIntrinsicForClangBuiltin -// resolves it. Neither needs a case in emitAMDGPUBuiltinExpr; these tests pin -// the emission down. - // CIR-LABEL: @_Z9test_lerpPjjjj // CIR: cir.call_llvm_intrinsic "amdgcn.lerp" {{.*}} : (!u32i, !u32i, !u32i) -> !u32i // LLVM: define{{.*}} void @_Z9test_lerpPjjjj _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
