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 <[email protected]> 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 <[email protected]> --- 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); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
