https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/233756
>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 1/2] [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); +} >From e83d62cc0a8a47b2b24c730ce23c0337337ab2d4 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Tue, 20 Oct 2026 12:47:23 +0200 Subject: [PATCH 2/2] comment --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 23 +++++++------------ 1 file changed, 8 insertions(+), 15 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 4e62f3f1b026cb..b30616a1fc8d70 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -822,24 +822,17 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_cluster_load_b32: + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.cluster.load.b32", + convertType(expr->getType())) + .getValue(); case AMDGPU::BI__builtin_amdgcn_cluster_load_b64: - case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: { - 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, + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.cluster.load.b64", + convertType(expr->getType())) + .getValue(); + case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: + return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.cluster.load.b128", convertType(expr->getType())) .getValue(); - } case AMDGPU::BI__builtin_amdgcn_load_to_lds: return emitBuiltinWithOneOverloadedType<5>( expr, "amdgcn.load.to.lds", _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
