Revert "Set calling convention for CUDA kernel"

This reverts r328795 which introduced an issue with referencing __global__
function templates. More details in the original review D44747.

llvm-svn: 329099
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index c5f581f..0e097da 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -25,7 +25,6 @@
 #include "clang/AST/ExprObjC.h"
 #include "clang/AST/ExprOpenMP.h"
 #include "clang/AST/RecursiveASTVisitor.h"
-#include "clang/AST/Type.h"
 #include "clang/AST/TypeLoc.h"
 #include "clang/Basic/PartialDiagnostic.h"
 #include "clang/Basic/SourceManager.h"
@@ -1659,16 +1658,6 @@
       isa<VarDecl>(D) &&
       NeedToCaptureVariable(cast<VarDecl>(D), NameInfo.getLoc());
 
-  // Drop CUDA kernel calling convention since it is invisible to the user
-  // in DRE.
-  if (const auto *FT = Ty->getAs<FunctionType>()) {
-    if (FT->getCallConv() == CC_CUDAKernel) {
-      FT = Context.adjustFunctionType(FT,
-                                      FT->getExtInfo().withCallingConv(CC_C));
-      Ty = QualType(FT, Ty.getQualifiers().getAsOpaqueValue());
-    }
-  }
-
   DeclRefExpr *E;
   if (isa<VarTemplateSpecializationDecl>(D)) {
     VarTemplateSpecializationDecl *VarSpec =
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index e917bcc..dea1c2e 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -1481,6 +1481,7 @@
                  .getTypePtr());
       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.
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 7bcc5b6..00bb21f 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -3316,18 +3316,6 @@
   CallingConv CC = S.Context.getDefaultCallingConvention(FTI.isVariadic,
                                                          IsCXXInstanceMethod);
 
-  // Attribute AT_CUDAGlobal affects the calling convention for AMDGPU targets.
-  // This is the simplest place to infer calling convention for CUDA kernels.
-  if (S.getLangOpts().CUDA && S.getLangOpts().CUDAIsDevice) {
-    for (const AttributeList *Attr = D.getDeclSpec().getAttributes().getList();
-         Attr; Attr = Attr->getNext()) {
-      if (Attr->getKind() == AttributeList::AT_CUDAGlobal) {
-        CC = CC_CUDAKernel;
-        break;
-      }
-    }
-  }
-
   // Attribute AT_OpenCLKernel affects the calling convention for SPIR
   // and AMDGPU targets, hence it cannot be treated as a calling
   // convention attribute. This is the simplest place to infer