[PATCH] D44747: Set calling convention for CUDA kernel

[Liu, Yaxun (Sam) via cfe-commits]

I will try fixing that.

The CUDA kernel calling convention should be dropped in all DRE's since it is invisible to the user.

Sam

-----Original Message-----
From: Artem Belevich via Phabricator [mailto:reviews at reviews.llvm.org] 
Sent: Tuesday, April 03, 2018 1:51 PM
To: Liu, Yaxun (Sam) <Yaxun.Liu at amd.com>; rjmccall at gmail.com; Arsenault, Matthew <Matthew.Arsenault at amd.com>
Cc: llvm-commits at lists.llvm.org; tra at google.com; Zhuravlyov, Konstantin <Konstantin.Zhuravlyov at amd.com>; wei.ding2 at amd.com; Stuttard, David <David.Stuttard at amd.com>; tpr.llvm at botech.co.uk; Tye, Tony <Tony.Tye at amd.com>; cfe-commits at lists.llvm.org; jlebar at google.com
Subject: [PATCH] D44747: Set calling convention for CUDA kernel

tra added inline comments.


================
Comment at: lib/Sema/SemaType.cpp:3319-3330
+  // Attribute AT_CUDAGlobal affects the calling convention for AMDGPU targets.
+  // This is the simplest place to infer calling convention for CUDA kernels.
+  if (S.getLangOpts().CUDA && S.getLangOpts().CUDAIsDevice) {
+    for (const AttributeList *Attr = D.getDeclSpec().getAttributes().getList();
+         Attr; Attr = Attr->getNext()) {
+      if (Attr->getKind() == AttributeList::AT_CUDAGlobal) {
+        CC = CC_CUDAKernel;
----------------
tra wrote:
> This apparently breaks compilation of some CUDA code in our internal tests. I'm working on minimizing a reproduction case. Should this code be enabled for AMD GPUs only?
Here's a small snippet of code that previously used to compile and work:

```
template <typename T>
__global__ void EmptyKernel(void) { }

struct Dummy {
  /// Type definition of the EmptyKernel kernel entry point
  typedef void (*EmptyKernelPtr)();
  EmptyKernelPtr Empty() { return EmptyKernel<void>; } }; ``` AFAICT,  it's currently impossible to apply __global__ to pointers, so there's no way to make the code above work with this patch applied.


Repository:
  rL LLVM

https://reviews.llvm.org/D44747