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