Author: Ayokunle Amodu Date: 2026-10-08T07:42:46-04:00 New Revision: 6a3f98ef319db8cba1fb5ff45d6bc09a7bb16f27
URL: https://github.com/llvm/llvm-project/commit/6a3f98ef319db8cba1fb5ff45d6bc09a7bb16f27 DIFF: https://github.com/llvm/llvm-project/commit/6a3f98ef319db8cba1fb5ff45d6bc09a7bb16f27.diff LOG: [CIR][AMDGPU] Add support for AMDGCN s_sendmsg_rtn builtins (#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. Assisted by: Claude Opus 5 Co-authored-by: Claude Opus 5 <[email protected]> Added: Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 89bfb4862c5c6..0ee8f18551e41 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -1051,12 +1051,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(), 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
