https://github.com/aobolensk created
https://github.com/llvm/llvm-project/pull/233756
Lower __builtin_amdgcn_load_to_lds, load_async_to_lds and
cluster_load_b{32,64,128} to their overloaded intrinsics, matching classic
codegen
>From f5898322d065bee79ecd6e3871d701fff9989314 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <[email protected]>
Date: Tue, 20 Oct 2026 11:25:17 +0200
Subject: [PATCH] [CIR][AMDGPU] Support load to LDS and cluster load builtins
Lower __builtin_amdgcn_load_to_lds, load_async_to_lds and
cluster_load_b{32,64,128} to their overloaded intrinsics, matching classic
codegen
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 35 +++++++++++++------
.../CodeGenHIP/builtins-amdgcn-gfx1250.hip | 27 ++++++++++++++
.../CIR/CodeGenHIP/builtins-amdgcn-gfx950.hip | 19 ++++++++++
3 files changed, 71 insertions(+), 10 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index cf2a424c341489..4e62f3f1b026cb 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -824,17 +824,32 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:
case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:
case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
- case AMDGPU::BI__builtin_amdgcn_load_to_lds: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ llvm::StringRef intrinsicName;
+ switch (builtinId) {
+ case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:
+ intrinsicName = "amdgcn.cluster.load.b32";
+ break;
+ case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:
+ intrinsicName = "amdgcn.cluster.load.b64";
+ break;
+ case AMDGPU::BI__builtin_amdgcn_cluster_load_b128:
+ intrinsicName = "amdgcn.cluster.load.b128";
+ break;
+ }
+ return emitBuiltinWithOneOverloadedType<3>(expr, intrinsicName,
+ convertType(expr->getType()))
+ .getValue();
}
+ case AMDGPU::BI__builtin_amdgcn_load_to_lds:
+ return emitBuiltinWithOneOverloadedType<5>(
+ expr, "amdgcn.load.to.lds",
+ cir::VoidType::get(builder.getContext()))
+ .getValue();
+ case AMDGPU::BI__builtin_amdgcn_load_async_to_lds:
+ return emitBuiltinWithOneOverloadedType<5>(
+ expr, "amdgcn.load.async.to.lds",
+ cir::VoidType::get(builder.getContext()))
+ .getValue();
case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:
case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
index 3c400aec174e23..216d0643ab2446 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
@@ -151,3 +151,30 @@ __device__ void test_permlane_xor(unsigned int* out,
unsigned int src0,
unsigned int src1, unsigned int src2) {
*out = __builtin_amdgcn_permlane_xor(src0, src1, src2);
}
+
+typedef int v2i __attribute__((ext_vector_type(2)));
+typedef int v4i __attribute__((ext_vector_type(4)));
+
+// CIR-LABEL: @_Z21test_cluster_load_b32PiPU3AS1ii
+// CIR: cir.call_llvm_intrinsic "amdgcn.cluster.load.b32" {{.*}} :
(!cir.ptr<!s32i, target_address_space(1)>, !s32i, !s32i) -> !s32i
+// LLVM: define{{.*}} void @_Z21test_cluster_load_b32PiPU3AS1ii
+// LLVM: call{{.*}} i32 @llvm.amdgcn.cluster.load.b32.i32(ptr addrspace(1)
%{{.+}}, i32 10, i32 %{{.+}})
+__device__ void test_cluster_load_b32(int* out, global int* inptr, int mask) {
+ *out = __builtin_amdgcn_cluster_load_b32(inptr, 10, mask);
+}
+
+// CIR-LABEL: @_Z21test_cluster_load_b64PDv2_iPU3AS1S_i
+// CIR: cir.call_llvm_intrinsic "amdgcn.cluster.load.b64" {{.*}} :
(!cir.ptr<!cir.vector<2 x !s32i>, target_address_space(1)>, !s32i, !s32i) ->
!cir.vector<2 x !s32i>
+// LLVM: define{{.*}} void @_Z21test_cluster_load_b64PDv2_iPU3AS1S_i
+// LLVM: call{{.*}} <2 x i32> @llvm.amdgcn.cluster.load.b64.v2i32(ptr
addrspace(1) %{{.+}}, i32 22, i32 %{{.+}})
+__device__ void test_cluster_load_b64(v2i* out, global v2i* inptr, int mask) {
+ *out = __builtin_amdgcn_cluster_load_b64(inptr, 22, mask);
+}
+
+// CIR-LABEL: @_Z22test_cluster_load_b128PDv4_iPU3AS1S_i
+// CIR: cir.call_llvm_intrinsic "amdgcn.cluster.load.b128" {{.*}} :
(!cir.ptr<!cir.vector<4 x !s32i>, target_address_space(1)>, !s32i, !s32i) ->
!cir.vector<4 x !s32i>
+// LLVM: define{{.*}} void @_Z22test_cluster_load_b128PDv4_iPU3AS1S_i
+// LLVM: call{{.*}} <4 x i32> @llvm.amdgcn.cluster.load.b128.v4i32(ptr
addrspace(1) %{{.+}}, i32 27, i32 %{{.+}})
+__device__ void test_cluster_load_b128(v4i* out, global v4i* inptr, int mask) {
+ *out = __builtin_amdgcn_cluster_load_b128(inptr, 27, mask);
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950.hip
b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950.hip
index 4fd0e12d6b0439..aec538eae5a5cd 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx950.hip
@@ -31,3 +31,22 @@ __device__ void test_bitop3_b16(global unsigned short* out,
unsigned short a,
unsigned short b, unsigned short c) {
*out = __builtin_amdgcn_bitop3_b16(a, b, c, 1);
}
+
+// CIR-LABEL: @_Z16test_load_to_ldsPjS_
+// CIR: cir.call_llvm_intrinsic "amdgcn.load.to.lds" {{.*}} :
(!cir.ptr<!void>, !cir.ptr<!void, target_address_space(3)>, !u32i, !s32i,
!u32i) -> !void
+// LLVM: define{{.*}} void @_Z16test_load_to_ldsPjS_
+// LLVM: call void @llvm.amdgcn.load.to.lds.p0(ptr %{{.+}}, ptr addrspace(3)
%{{.+}}, i32 4, i32 0, i32 0)
+__device__ void test_load_to_lds(unsigned int* src, unsigned int* dst) {
+ __builtin_amdgcn_load_to_lds(src, dst, /*size=*/4, /*offset=*/0, /*aux=*/0);
+}
+
+// CIR-LABEL: @_Z22test_load_async_to_ldsPU3AS1jPU3AS3j
+// CIR: cir.call_llvm_intrinsic "amdgcn.load.async.to.lds" {{.*}} :
(!cir.ptr<!void, target_address_space(1)>, !cir.ptr<!void,
target_address_space(3)>, !u32i, !s32i, !u32i) -> !void
+// LLVM: define{{.*}} void @_Z22test_load_async_to_ldsPU3AS1jPU3AS3j
+// LLVM: call void @llvm.amdgcn.load.async.to.lds.p1(ptr addrspace(1) %{{.+}},
ptr addrspace(3) %{{.+}}, i32 16, i32 8, i32 1)
+__device__ void
+test_load_async_to_lds(global unsigned int* src,
+ __attribute__((address_space(3))) unsigned int* dst) {
+ __builtin_amdgcn_load_async_to_lds(src, dst, /*size=*/16, /*offset=*/8,
+ /*aux=*/1);
+}
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
