[PATCH] D45223: [CUDA] Fix overloading resolution failure due to kernel calling convention

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl added a comment.

In https://reviews.llvm.org/D45223#1056187, @rjmccall wrote:

> I think the appropriate place to do this is in IsStandardConversion, immediately after the call to ResolveAddressOfOverloadedFunction.  You might want to add a general utility for getting the type-of-reference of a function decl.


We may need to resolve overloaded functions with dropped calling conventions, e.g.

  __global__ void EmptyKernel(float) {}
  
  __global__ void EmptyKernel(double) {}
  
  struct Dummy {
    /// Type definition of the EmptyKernel kernel entry point
    typedef void (*EmptyKernelPtr)(float);
    EmptyKernelPtr Empty() { return EmptyKernel; } 
  };

In this case we have to drop the calling convention during the resolution.

Since the calling convention is invisible in the AST, why don't why just do not represent it in AST?

Going back to the original implementation in CodeGen:

  if ((getTriple().getArch() == llvm::Triple::amdgcn) &&
      D->hasAttr<CUDAGlobalAttr>())
    Fn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);

It is much simpler and straightforward.

Can we just reconsider implement this in CodeGen instead of Sema?


https://reviews.llvm.org/D45223