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:revi...@reviews.llvm.org] 
Sent: Tuesday, April 03, 2018 2:09 PM
To: Liu, Yaxun (Sam) <yaxun....@amd.com>; rjmcc...@gmail.com; Arsenault, 
Matthew <matthew.arsena...@amd.com>
Cc: jle...@google.com; llvm-comm...@lists.llvm.org; t...@google.com; 
Zhuravlyov, Konstantin <konstantin.zhuravl...@amd.com>; wei.di...@amd.com; 
Stuttard, David <david.stutt...@amd.com>; tpr.l...@botech.co.uk; Tye, Tony 
<tony....@amd.com>; cfe-commits@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



_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to