[CUDA] Set LLVM calling convention for CUDA kernel

[yaxunl]

Summary

Some targets need special LLVM calling convention for CUDA kernel.
This patch does that through a TargetCodeGenInfo hook.

It only affects amdgcn target.

Patch by Greg Rodgers.
Revised and lit tests added by Yaxun Liu.

[rjmccall]

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.

[yaxunl]

In 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 we 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?

[rjmccall]

Yes, I'm sorry, I think you're right. I had misunderstood the language problem when I suggested going down this road.

[yaxunl]

In D45223#1071358, @rjmccall wrote:

Yes, I'm sorry, I think you're right. I had misunderstood the language problem when I suggested going down this road.

In D45223#1071358, @rjmccall wrote:

Yes, I'm sorry, I think you're right. I had misunderstood the language problem when I suggested going down this road.

Never mind. I will update the diff for CodeGen approach.

[yaxunl]

Use CodeGen approach.

[tra]

AFAICT this is the replacement for D44747. LGTM.

[yaxunl]

In D45223#1071452, @tra wrote:

AFAICT this is the replacement for D44747. LGTM.

Yes. Thanks.

lib/Sema/SemaOverload.cpp

   bool Changed = false;
 
   // Drop 'noreturn' if not present in target type.
   if (FromEInfo.getNoReturn() && !ToEInfo.getNoReturn()) {
     FromFn = Context.adjustFunctionType(FromFn, FromEInfo.withNoReturn(false));
     Changed = true;
   }
 
-  // Drop 'noexcept' if not present in target type.
   if (const auto *FromFPT = dyn_cast<FunctionProtoType>(FromFn)) {
     const auto *ToFPT = cast<FunctionProtoType>(ToFn);
+
+    // Drop 'noexcept' if not present in target type.
     if (FromFPT->isNothrow(Context) && !ToFPT->isNothrow(Context)) {
       FromFn = cast<FunctionType>(
           Context.getFunctionTypeWithExceptionSpec(QualType(FromFPT, 0),
                                                    EST_None)
                  .getTypePtr());
       Changed = true;
     }
+
+    // Drop cuda_kernel calling convention since it is invisible in AST.
+    if (FromFPT->getCallConv() == CC_CUDAKernel &&
+        FromFPT->getCallConv() != ToFPT->getCallConv()) {
+      FromFn = Context.adjustFunctionType(
+          FromFn, FromEInfo.withCallingConv(ToFPT->getCallConv()));
+      Changed = true;
+    }
+
     // Convert FromFPT's ExtParameterInfo if necessary. The conversion is valid
     // only if the ExtParameterInfo lists of the two function prototypes can be
     // merged and the merged list is identical to ToFPT's ExtParameterInfo list.
     SmallVector<FunctionProtoType::ExtParameterInfo, 4> NewParamInfos;
     bool CanUseToFPT, CanUseFromFPT;
     if (Context.mergeExtParameterInfo(ToFPT, FromFPT, CanUseToFPT,
                                       CanUseFromFPT, NewParamInfos) &&
         CanUseToFPT && !CanUseFromFPT) {

test/CodeGenCUDA/kernel-amdgcn.cu

 // CHECK: define void @_Z10non_kernelv
 __device__ void non_kernel(){}
 
 // CHECK: define amdgpu_kernel void @_Z6kerneli
 __global__ void kernel(int x) {
   non_kernel();
 }
 
+// CHECK: define amdgpu_kernel void @_Z11EmptyKernelIvEvv
+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>; } 
+};
+
 // CHECK: define amdgpu_kernel void @_Z15template_kernelI1AEvT_
 template<class T>
 __global__ void template_kernel(T x) {}
 
 void launch(void *f);
 
 int main() {
+  Dummy D;
   launch((void*)A::kernel);
   launch((void*)kernel);
   launch((void*)template_kernel<A>);
+  launch((void*)D.Empty());
   return 0;
 }