https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/214246
CIR emitted no calling convention for HIP `__global__` kernels on the `spirv64-amd-amdhsa` target, unlike generic SPIR-V, which already gets `spirv_kernel` CC >From e0289b4451833784512fa7f6f3702ea10802e536 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Wed, 5 Aug 2026 17:05:13 +0200 Subject: [PATCH] [CIR][SPIR-V] Set spir_kernel calling convention for AMDGCN-flavored SPIR-V HIP kernels --- clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 9 +++++++- .../CIR/CodeGenHIP/amdgcnspirv-kernel.hip | 23 +++++++++++++++++++ 2 files changed, 31 insertions(+), 1 deletion(-) create mode 100644 clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp index 643c635128d09..b56bbfe41e87c 100644 --- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp +++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp @@ -42,11 +42,18 @@ class SPIRVTargetCIRGenInfo : public TargetCIRGenInfo { if (!fd) return; + auto func = mlir::cast<cir::FuncOp>(global); + if (cgm.getLangOpts().OpenCL && DeviceKernelAttr::isOpenCLSpelling(fd->getAttr<DeviceKernelAttr>())) { - auto func = mlir::cast<cir::FuncOp>(global); func.setCallingConv(cir::CallingConv::SpirKernel); + return; } + + if (cgm.getLangOpts().HIP && + cgm.getTriple().getVendor() == llvm::Triple::AMD && + fd->hasAttr<CUDAGlobalAttr>()) + func.setCallingConv(cir::CallingConv::SpirKernel); } }; diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip new file mode 100644 index 0000000000000..039ab35f1c906 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip @@ -0,0 +1,23 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \ +// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir + +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ll + +// Test that HIP kernels on AMDGCN-flavored SPIR-V get the spir_kernel +// calling convention. + +#define __global__ __attribute__((global)) +#define __device__ __attribute__((device)) + +// CIR: cir.func{{.*}} @_Z13kernel_scalari{{.*}} cc(spir_kernel) +// LLVM: define spir_kernel void @_Z13kernel_scalari +__global__ void kernel_scalar(int a) {} + +// CIR: cir.func{{.*}} @_Z9device_fni +// CIR-NOT: cc(spir_kernel) +// LLVM: define{{.*}} void @_Z9device_fni +__device__ void device_fn(int a) {} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
