[clang] d8aecd0 - [CIR][AMDGPU] Add support for wave_shuffle built-in (#225463)

[via cfe-commits]

Author: Mariya Podchishchaeva
Date: 2026-09-23T09:25:16+02:00
New Revision: d8aecd0b3ff64a484197e187a292f8841e4f8e93

URL: https://github.com/llvm/llvm-project/commit/d8aecd0b3ff64a484197e187a292f8841e4f8e93
DIFF: https://github.com/llvm/llvm-project/commit/d8aecd0b3ff64a484197e187a292f8841e4f8e93.diff

LOG: [CIR][AMDGPU] Add support for wave_shuffle built-in (#225463)

Added: 
    

Modified: 
    clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.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 8d23951dd64ba..971ecb18e990d 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -240,12 +240,9 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
   case AMDGPU::BI__builtin_amdgcn_readfirstlane:
     return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.readfirstlane")
         .getValue();
-  case AMDGPU::BI__builtin_amdgcn_wave_shuffle: {
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented AMDGPU builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
-  }
+  case AMDGPU::BI__builtin_amdgcn_wave_shuffle:
+    return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.wave.shuffle")
+        .getValue();
   case AMDGPU::BI__builtin_amdgcn_div_fixup:
   case AMDGPU::BI__builtin_amdgcn_div_fixupf:
   case AMDGPU::BI__builtin_amdgcn_div_fixuph: {

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index cef2d879be0d1..01afd609a9ba9 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -382,3 +382,12 @@ __device__ void test_sbfe(unsigned int* out, unsigned int src0, unsigned int src
                           unsigned int src2) {
   *out = __builtin_amdgcn_sbfe(src0, src1, src2);
 }
+
+// CIR-LABEL: @_Z17test_wave_shufflePiii
+// CIR: cir.call_llvm_intrinsic "amdgcn.wave.shuffle" {{.*}} : (!s32i, !s32i) -> !s32i
+// LLVM: define{{.*}} void @_Z17test_wave_shufflePiii
+// LLVM: {{.*}}call{{.*}} i32 @llvm.amdgcn.wave.shuffle.i32(i32 %{{.+}}, i32 %{{.+}})
+__device__ void test_wave_shuffle(int* out, int a, int b)
+{
+  *out = __builtin_amdgcn_wave_shuffle(a, b);
+}