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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added reviewers: rjmccall, tra.

The following test causes overloading resolution failure in clang due to
missing kernel calling convention in the function pointer type.

  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>; } 
  };

This happens before DRE is created. The fix is to drop the kernel
calling convention when converting function types.


https://reviews.llvm.org/D45223

Files:
  lib/Sema/SemaOverload.cpp
  test/CodeGenCUDA/kernel-amdgcn.cu


Index: test/CodeGenCUDA/kernel-amdgcn.cu
===================================================================
--- test/CodeGenCUDA/kernel-amdgcn.cu
+++ test/CodeGenCUDA/kernel-amdgcn.cu
@@ -15,15 +15,27 @@
   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;
 }
Index: lib/Sema/SemaOverload.cpp
===================================================================
--- lib/Sema/SemaOverload.cpp
+++ lib/Sema/SemaOverload.cpp
@@ -1471,16 +1471,26 @@
     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.



Index: test/CodeGenCUDA/kernel-amdgcn.cu
===================================================================
--- test/CodeGenCUDA/kernel-amdgcn.cu
+++ test/CodeGenCUDA/kernel-amdgcn.cu
@@ -15,15 +15,27 @@
   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;
 }
Index: lib/Sema/SemaOverload.cpp
===================================================================
--- lib/Sema/SemaOverload.cpp
+++ lib/Sema/SemaOverload.cpp
@@ -1471,16 +1471,26 @@
     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.

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