[PATCH] D44747: Set calling convention for CUDA kernel

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

Let's revert it for now. I will create a review after fixing it and commit it again with the fix.

Thanks.

Sam 

-----Original Message-----
From: Artem Belevich via Phabricator [mailto:reviews at reviews.llvm.org] 
Sent: Tuesday, April 03, 2018 2:09 PM
To: Liu, Yaxun (Sam) <Yaxun.Liu at amd.com>; rjmccall at gmail.com; Arsenault, Matthew <Matthew.Arsenault at amd.com>
Cc: jlebar at google.com; 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
Subject: [PATCH] D44747: Set calling convention for CUDA kernel

tra added a comment.

In https://reviews.llvm.org/D44747#1055877, @yaxunl wrote:

> 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


Apparently it's not always the case.
Here's a bit more elaborate example demonstrating inconsistency in this behavior. Calling convention is ignored for regular functions, but not for function templates.

  __global__ void EmptyKernel(void) { }
  
  template <typename T>
  __global__ void EmptyKernelT(void) { }
  
  struct Dummy {
    /// Type definition of the EmptyKernel kernel entry point
    typedef void (*EmptyKernelPtr)();
    EmptyKernelPtr Empty() { return EmptyKernel; } // this one works
    EmptyKernelPtr EmptyT() { return EmptyKernelT<void>; } // this one errors out.
  };

Do you think this is something you can fix quickly or do you want to unroll the change while you figure out what's going on?


Repository:
  rL LLVM

https://reviews.llvm.org/D44747