[clang] fe36473 - [CIR][AMDGPU] Let amdgcn_lerp and ds_swizzle use the generic codegen path (#222204)

[via cfe-commits]

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