Index: lib/Sema/SemaOverload.cpp
===================================================================
--- lib/Sema/SemaOverload.cpp
+++ lib/Sema/SemaOverload.cpp
@@ -1471,9 +1471,10 @@
     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),
@@ -1481,6 +1482,15 @@
                  .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,6 +15,16 @@
   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) {}
@@ -22,8 +32,10 @@
 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;
 }