[clang] [CIR][SPIR-V] Set spir_kernel calling convention for AMDGCN-flavored HIP kernels (PR #214246)

[via cfe-commits]

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