[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:19:26 PDT 2026
https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/223226
>From dd2989537cf6cb964c95ac4d92b0e20daf7afa86 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Sun, 13 Sep 2026 12:10:25 +0200
Subject: [PATCH 1/2] [CIR][AMDGPU] Add support for AMDGCN s_sendmsg_rtn
builtins
Adds codegen for the following AMDGCN s_sendmsg_rtn builtins:
__builtin_amdgcn_s_sendmsg_rtn (unsigned int)
__builtin_amdgcn_s_sendmsg_rtnl (uint64_t)
These are lowered to the llvm.amdgcn.s.sendmsg.rtn intrinsic.
Co-Authored-By: Claude Opus 5 <noreply at anthropic.com>
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 11 ++--
.../CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip | 64 +++++++++++++++++++
2 files changed, 69 insertions(+), 6 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 1c23ea142bf8d..e0c2ab4b620c8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -974,12 +974,11 @@ 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:
+ // s_sendmsg_rtn is mangled using return type only.
+ 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(),
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);
+}
>From aead50eb03b18bda4b5027d5596b2ff7583d6019 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Sun, 13 Sep 2026 12:17:53 +0200
Subject: [PATCH 2/2] remove comment
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 1 -
1 file changed, 1 deletion(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index e0c2ab4b620c8..fd9fc34c7cdbd 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -975,7 +975,6 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
}
case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn:
case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl:
- // s_sendmsg_rtn is mangled using return type only.
return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.s.sendmsg.rtn",
convertType(expr->getType()))
.getValue();
More information about the cfe-commits
mailing list