[clang] [CIR][AMDGPU] Add lowering for amdgcn ds swizzle builtin. (PR #196011)
via cfe-commits
cfe-commits at lists.llvm.org
Wed May 6 22:16:43 PDT 2026
https://github.com/skc7 updated https://github.com/llvm/llvm-project/pull/196011
>From 1b6cc529808d2ecc3a286e75c8e4b5587db67093 Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Wed, 6 May 2026 12:17:31 +0530
Subject: [PATCH 1/2] [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);
+}
>From 77d2dc9877988257221de5108217bf4ee3b88bde Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Thu, 7 May 2026 10:45:48 +0530
Subject: [PATCH 2/2] use emitIntrinsicCallOp from builder
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 9 +++------
1 file changed, 3 insertions(+), 6 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 7e722d4e654d3..04ab1c29b0d63 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -164,12 +164,9 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
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;
+ return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()),
+ "amdgcn.ds.swizzle", src0.getType(),
+ mlir::ValueRange{src0, src1});
}
case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
case AMDGPU::BI__builtin_amdgcn_mov_dpp:
More information about the cfe-commits
mailing list