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); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
