Author: Ayokunle Amodu Date: 2026-09-11T21:16:18-04:00 New Revision: fe36473cb62cca28c14b7174c0797cb74bcdfc7e
URL: https://github.com/llvm/llvm-project/commit/fe36473cb62cca28c14b7174c0797cb74bcdfc7e DIFF: https://github.com/llvm/llvm-project/commit/fe36473cb62cca28c14b7174c0797cb74bcdfc7e.diff LOG: [CIR][AMDGPU] Let amdgcn_lerp and ds_swizzle use the generic codegen path (#222204) __builtin_amdgcn_lerp and __builtin_amdgcn_ds_swizzle both map 1:1 onto concrete (non-overloaded) intrinsics and carry the ClangBuiltin<> marker in IntrinsicsAMDGPU.td, so EmitBuiltinExpr already lowers them automatically by forwarding the arguments to the intrinsic. Added: Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp clang/lib/CodeGen/TargetBuiltins/AMDGPU.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 9f0cd515607b4..31bffd1ed96c4 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: { 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 713bc10700cdd..c5f6d1997c1a1 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -348,6 +348,15 @@ __device__ void test_class_f64(bool* out, double a, int b) { *out = __builtin_amdgcn_class(a, b); } +// 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); +} + // CIR-LABEL: @_Z9test_ubfePjjjj // CIR: cir.call_llvm_intrinsic "amdgcn.ubfe" {{.*}} : (!u32i, !u32i, !u32i) -> !u32i // LLVM: define{{.*}} void @_Z9test_ubfePjjjj _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
