llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Ayokunle Amodu (ayokunle321) <details> <summary>Changes</summary> 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 --- Full diff: https://github.com/llvm/llvm-project/pull/223226.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+5-6) - (modified) clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx11.hip (+64) ``````````diff 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); +} `````````` </details> https://github.com/llvm/llvm-project/pull/223226 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
