[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:16:11 PDT 2026
https://github.com/ayokunle321 created https://github.com/llvm/llvm-project/pull/223226
Adds codegen for the following AMDGCN s_sendmsg_rtn builtins:
- __builtin_amdgcn_s_sendmsg_rtn
- __builtin_amdgcn_s_sendmsg_rtnl
These are lowered to the `llvm.amdgcn.s.sendmsg.rtn` intrinsic, which is mangled using the return type only, so the result type is passed explicitly rather than taken from the first argument.
The test also picks up the gfx11 targets OGCG already covers (`amdgpu11.54`, `amdgpu11.70`, `amdgpu11.71` and `amdgpu11.72`). OGCG's `spirv64-amd-amdhsa` run is not ported, since it relies on `__builtin_amdgcn_is_invocable`, which CIR does not support yet.
Assisted by: Claude Opus 5
>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] [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);
+}
More information about the cfe-commits
mailing list