llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir @llvm/pr-subscribers-clang Author: Arseniy Obolenskiy (aobolensk) <details> <summary>Changes</summary> 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 --- Full diff: https://github.com/llvm/llvm-project/pull/214246.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/Targets/SPIRV.cpp (+8-1) - (added) clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip (+23) ``````````diff 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) {} `````````` </details> https://github.com/llvm/llvm-project/pull/214246 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
