https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/214246
>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 1/2] [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) {} >From 5a17e67fd7786430df8330a9b1c36535984a3ec2 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Wed, 5 Aug 2026 17:42:59 +0200 Subject: [PATCH 2/2] fix integerattr The program address space DLTI entry is stored as an unsigned integer attribute, but was read with getInt(), which asserts on signless types. --- clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp b/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp index d7d563d79a0fa..5ce0c8f8e7a2c 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDataLayout.cpp @@ -26,7 +26,7 @@ void CIRDataLayout::reset(mlir::DataLayoutSpecInterface spec) { if (mlir::DataLayoutEntryInterface entry = spec.getSpecForIdentifier(addrSpKey)) if (auto val = llvm::dyn_cast<mlir::IntegerAttr>(entry.getValue())) - programAddrSpace = val.getInt(); + programAddrSpace = val.getUInt(); } } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
