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

[Ayokunle Amodu via cfe-commits]

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 <ayokunle321 at gmail.com>
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 <ayokunle321 at gmail.com>
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