[clang] [CIR][AMDGPU] Add support for AMDGCN s_sendmsg_rtn builtins (PR #223226)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Sun Sep 13 03:24:51 PDT 2026
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 <ayokunle321 at gmail.com>
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 <noreply at anthropic.com>
---
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 at +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);
-}
More information about the cfe-commits
mailing list