https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/223226
>From 94ffb7a4adf46f87e88ce57cefbd18889784641d Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sun, 13 Sep 2026 12:23:23 +0200 Subject: [PATCH] [CIR][AMDGPU] Add support for AMDGCN s_sendmsg_rtn and s_prefetch builtins Adds codegen for the following AMDGCN builtins: __builtin_amdgcn_s_sendmsg_rtn (unsigned int) __builtin_amdgcn_s_sendmsg_rtnl (uint64_t) __builtin_amdgcn_s_prefetch_data __builtin_amdgcn_s_prefetch_inst These are lowered to the corresponding llvm.amdgcn intrinsics. Co-Authored-By: Claude Opus 5 <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 32 ++++------ .../CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip | 64 +++++++++++++++++++ .../CodeGenHIP/builtins-amdgcn-prefetch.hip | 51 +++++++++++++++ .../builtins-amdgcn-s-prefetch-inst-nyi.hip | 10 --- 4 files changed, 129 insertions(+), 28 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-prefetch.hip delete mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-s-prefetch-inst-nyi.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 1c23ea142bf8d..485771ca0fa1f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -974,12 +974,10 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return mlir::Value{}; } case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn: - case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.s.sendmsg.rtn", + convertType(expr->getType())) + .getValue(); case AMDGPU::BI__builtin_amdgcn_permlane16_swap: case AMDGPU::BI__builtin_amdgcn_permlane32_swap: { cgm.errorNYI(expr->getSourceRange(), @@ -1050,18 +1048,16 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, getContext().BuiltinInfo.getName(builtinId)); return mlir::Value{}; } - case AMDGPU::BI__builtin_amdgcn_s_prefetch_data: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } - case AMDGPU::BI__builtin_amdgcn_s_prefetch_inst: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + case AMDGPU::BI__builtin_amdgcn_s_prefetch_data: + return emitBuiltinWithOneOverloadedType<2>( + expr, "amdgcn.s.prefetch.data", + cir::VoidType::get(builder.getContext())) + .getValue(); + case AMDGPU::BI__builtin_amdgcn_s_prefetch_inst: + return emitBuiltinWithOneOverloadedType<2>( + expr, "amdgcn.s.prefetch.inst", + cir::VoidType::get(builder.getContext())) + .getValue(); case Builtin::BIlogbf: case Builtin::BI__builtin_logbf: return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEsingle()); diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip index 673fc93a014da..f7b5714d34809 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip @@ -31,6 +31,22 @@ // RUN: -fcuda-is-device -emit-cir %s -o %t.cir // RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s +// RUN: %clang_cc1 -triple amdgpu11.54-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgpu11.70-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgpu11.71-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgpu11.72-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + // RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -x hip -std=c++11 -fclangir \ // RUN: -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s @@ -63,6 +79,22 @@ // RUN: -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s +// RUN: %clang_cc1 -triple amdgpu11.54-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu11.70-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu11.71-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu11.72-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + // RUN: %clang_cc1 -triple amdgpu11.00-amd-amdhsa -x hip -std=c++11 \ // RUN: -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s @@ -95,6 +127,22 @@ // RUN: -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s +// RUN: %clang_cc1 -triple amdgpu11.54-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu11.70-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu11.71-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu11.72-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + #define __device__ __attribute__((device)) //===----------------------------------------------------------------------===// @@ -108,3 +156,19 @@ __device__ void test_permlane64(unsigned int* out, unsigned int a) { *out = __builtin_amdgcn_permlane64(a); } + +// CIR-LABEL: @_Z18test_s_sendmsg_rtnPj +// CIR: cir.call_llvm_intrinsic "amdgcn.s.sendmsg.rtn" {{.*}} : (!u32i) -> !u32i +// LLVM: define{{.*}} void @_Z18test_s_sendmsg_rtnPj +// LLVM: call i32 @llvm.amdgcn.s.sendmsg.rtn.i32(i32 0) +__device__ void test_s_sendmsg_rtn(unsigned int* out) { + *out = __builtin_amdgcn_s_sendmsg_rtn(0); +} + +// CIR-LABEL: @_Z19test_s_sendmsg_rtnlPm +// CIR: cir.call_llvm_intrinsic "amdgcn.s.sendmsg.rtn" {{.*}} : (!u32i) -> !u64i +// LLVM: define{{.*}} void @_Z19test_s_sendmsg_rtnlPm +// LLVM: call i64 @llvm.amdgcn.s.sendmsg.rtn.i64(i32 0) +__device__ void test_s_sendmsg_rtnl(unsigned long* out) { + *out = __builtin_amdgcn_s_sendmsg_rtnl(0); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-prefetch.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-prefetch.hip new file mode 100644 index 0000000000000..1b58d69838ef4 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-prefetch.hip @@ -0,0 +1,51 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu12.00-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +#define __device__ __attribute__((device)) + +typedef __attribute__((address_space(1))) float *global_ptr_t; +typedef __attribute__((address_space(4))) const char *constant_ptr_t; + +//===----------------------------------------------------------------------===// +// Test AMDGPU builtins +//===----------------------------------------------------------------------===// + +// CIR-LABEL: @_Z20test_s_prefetch_dataPiPU3AS1fPU3AS4Kcj +// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.data" {{.*}} : (!cir.ptr<!void>, !u32i) -> !void +// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.data" {{.*}} : (!cir.ptr<!void, target_address_space(1)>, !u32i) -> !void +// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.data" {{.*}} : (!cir.ptr<!void, target_address_space(4)>, !u32i) -> !void +// LLVM: define{{.*}} void @_Z20test_s_prefetch_dataPiPU3AS1fPU3AS4Kcj +// LLVM: call void @llvm.amdgcn.s.prefetch.data.p0(ptr %{{.*}}, i32 0) +// LLVM: call void @llvm.amdgcn.s.prefetch.data.p1(ptr addrspace(1) %{{.*}}, i32 %{{.*}}) +// LLVM: call void @llvm.amdgcn.s.prefetch.data.p4(ptr addrspace(4) %{{.*}}, i32 31) +__device__ void test_s_prefetch_data(int *fp, global_ptr_t gp, + constant_ptr_t cp, unsigned int len) { + __builtin_amdgcn_s_prefetch_data(fp, 0); + __builtin_amdgcn_s_prefetch_data(gp, len); + __builtin_amdgcn_s_prefetch_data(cp, 31); +} + +// CIR-LABEL: @_Z20test_s_prefetch_instPiPU3AS1fPU3AS4Kcj +// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.inst" {{.*}} : (!cir.ptr<!void>, !u32i) -> !void +// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.inst" {{.*}} : (!cir.ptr<!void, target_address_space(1)>, !u32i) -> !void +// CIR: cir.call_llvm_intrinsic "amdgcn.s.prefetch.inst" {{.*}} : (!cir.ptr<!void, target_address_space(4)>, !u32i) -> !void +// LLVM: define{{.*}} void @_Z20test_s_prefetch_instPiPU3AS1fPU3AS4Kcj +// LLVM: call void @llvm.amdgcn.s.prefetch.inst.p0(ptr %{{.*}}, i32 0) +// LLVM: call void @llvm.amdgcn.s.prefetch.inst.p1(ptr addrspace(1) %{{.*}}, i32 %{{.*}}) +// LLVM: call void @llvm.amdgcn.s.prefetch.inst.p4(ptr addrspace(4) %{{.*}}, i32 31) +__device__ void test_s_prefetch_inst(int *fp, global_ptr_t gp, + constant_ptr_t cp, unsigned int len) { + __builtin_amdgcn_s_prefetch_inst(fp, 0); + __builtin_amdgcn_s_prefetch_inst(gp, len); + __builtin_amdgcn_s_prefetch_inst(cp, 31); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-s-prefetch-inst-nyi.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-s-prefetch-inst-nyi.hip deleted file mode 100644 index 74d058a725857..0000000000000 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-s-prefetch-inst-nyi.hip +++ /dev/null @@ -1,10 +0,0 @@ -// REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgpu12.50-amd-amdhsa -x hip -std=c++11 -fclangir \ -// RUN: -fcuda-is-device -emit-cir %s -verify -o %t.cir - -#define __device__ __attribute__((device)) - -// expected-error@+2 {{ClangIR code gen Not Yet Implemented: unimplemented AMDGPU builtin call: __builtin_amdgcn_s_prefetch_inst}} -__device__ void test_s_prefetch_inst(const void *p, unsigned int len) { - __builtin_amdgcn_s_prefetch_inst(p, len); -} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
