https://github.com/skc7 created https://github.com/llvm/llvm-project/pull/196011
Upstreaming clangIR PR: https://github.com/llvm/clangir/pull/2052 This PR adds support for lowering of _builtin_amdgcn_ds_swizzle* amdgpu builtin to clangIR. >From 1b6cc529808d2ecc3a286e75c8e4b5587db67093 Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Wed, 6 May 2026 12:17:31 +0530 Subject: [PATCH] [CIR][AMDGPU] Add lowering for amdgcn ds swizzle builtin. --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 11 ++++++++++- clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 8 ++++++++ 2 files changed, 18 insertions(+), 1 deletion(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 929cdf8e88789..7e722d4e654d3 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -161,7 +161,16 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, .getResult(); return result; } - case AMDGPU::BI__builtin_amdgcn_ds_swizzle: + case AMDGPU::BI__builtin_amdgcn_ds_swizzle: { + mlir::Value src0 = emitScalarExpr(expr->getArg(0)); + mlir::Value src1 = emitScalarExpr(expr->getArg(1)); + mlir::Value result = cir::LLVMIntrinsicCallOp::create( + builder, getLoc(expr->getExprLoc()), + builder.getStringAttr("amdgcn.ds.swizzle"), + src0.getType(), {src0, src1}) + .getResult(); + return result; + } 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/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index d374479e6182e..4a61fde7aa90c 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -63,3 +63,11 @@ __device__ void test_div_fmas_f32(double* out, float a, float b, float c, int d) __device__ void test_div_fmas_f64(double* out, double a, double b, double c, int d) { *out = __builtin_amdgcn_div_fmas(a, b, c, d); } + +// CIR-LABEL: @_Z15test_ds_swizzlePii +// CIR: cir.call_llvm_intrinsic "amdgcn.ds.swizzle" {{.*}} : (!s32i, !s32i) -> !s32i +// LLVM: define{{.*}} void @_Z15test_ds_swizzlePii +// LLVM: call{{.*}} i32 @llvm.amdgcn.ds.swizzle(i32 %{{.*}}, i32 32) +__device__ void test_ds_swizzle(int* out, int a) { + *out = __builtin_amdgcn_ds_swizzle(a, 32); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
