r329099 - Revert "Set calling convention for CUDA kernel"
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Tue Apr 3 11:29:31 PDT 2018
Author: tra
Date: Tue Apr 3 11:29:31 2018
New Revision: 329099
URL: http://llvm.org/viewvc/llvm-project?rev=329099&view=rev
Log:
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.
Removed:
cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu
Modified:
cfe/trunk/include/clang/Basic/Specifiers.h
cfe/trunk/lib/AST/ItaniumMangle.cpp
cfe/trunk/lib/AST/Type.cpp
cfe/trunk/lib/AST/TypePrinter.cpp
cfe/trunk/lib/CodeGen/CGCall.cpp
cfe/trunk/lib/CodeGen/CGDebugInfo.cpp
cfe/trunk/lib/CodeGen/TargetInfo.cpp
cfe/trunk/lib/CodeGen/TargetInfo.h
cfe/trunk/lib/Sema/SemaExpr.cpp
cfe/trunk/lib/Sema/SemaOverload.cpp
cfe/trunk/lib/Sema/SemaType.cpp
cfe/trunk/tools/libclang/CXType.cpp
Modified: cfe/trunk/include/clang/Basic/Specifiers.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Specifiers.h?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Specifiers.h (original)
+++ cfe/trunk/include/clang/Basic/Specifiers.h Tue Apr 3 11:29:31 2018
@@ -231,24 +231,23 @@ namespace clang {
/// \brief CallingConv - Specifies the calling convention that a function uses.
enum CallingConv {
- CC_C, // __attribute__((cdecl))
- CC_X86StdCall, // __attribute__((stdcall))
- CC_X86FastCall, // __attribute__((fastcall))
- CC_X86ThisCall, // __attribute__((thiscall))
+ CC_C, // __attribute__((cdecl))
+ CC_X86StdCall, // __attribute__((stdcall))
+ CC_X86FastCall, // __attribute__((fastcall))
+ CC_X86ThisCall, // __attribute__((thiscall))
CC_X86VectorCall, // __attribute__((vectorcall))
- CC_X86Pascal, // __attribute__((pascal))
- CC_Win64, // __attribute__((ms_abi))
- CC_X86_64SysV, // __attribute__((sysv_abi))
- CC_X86RegCall, // __attribute__((regcall))
- CC_AAPCS, // __attribute__((pcs("aapcs")))
- CC_AAPCS_VFP, // __attribute__((pcs("aapcs-vfp")))
- CC_IntelOclBicc, // __attribute__((intel_ocl_bicc))
- CC_SpirFunction, // default for OpenCL functions on SPIR target
- CC_OpenCLKernel, // inferred for OpenCL kernels
- CC_Swift, // __attribute__((swiftcall))
- CC_PreserveMost, // __attribute__((preserve_most))
- CC_PreserveAll, // __attribute__((preserve_all))
- CC_CUDAKernel, // inferred for CUDA kernels
+ CC_X86Pascal, // __attribute__((pascal))
+ CC_Win64, // __attribute__((ms_abi))
+ CC_X86_64SysV, // __attribute__((sysv_abi))
+ CC_X86RegCall, // __attribute__((regcall))
+ CC_AAPCS, // __attribute__((pcs("aapcs")))
+ CC_AAPCS_VFP, // __attribute__((pcs("aapcs-vfp")))
+ CC_IntelOclBicc, // __attribute__((intel_ocl_bicc))
+ CC_SpirFunction, // default for OpenCL functions on SPIR target
+ CC_OpenCLKernel, // inferred for OpenCL kernels
+ CC_Swift, // __attribute__((swiftcall))
+ CC_PreserveMost, // __attribute__((preserve_most))
+ CC_PreserveAll, // __attribute__((preserve_all))
};
/// \brief Checks whether the given calling convention supports variadic
Modified: cfe/trunk/lib/AST/ItaniumMangle.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ItaniumMangle.cpp?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ItaniumMangle.cpp (original)
+++ cfe/trunk/lib/AST/ItaniumMangle.cpp Tue Apr 3 11:29:31 2018
@@ -2628,7 +2628,6 @@ StringRef CXXNameMangler::getCallingConv
case CC_OpenCLKernel:
case CC_PreserveMost:
case CC_PreserveAll:
- case CC_CUDAKernel:
// FIXME: we should be mangling all of the above.
return "";
Modified: cfe/trunk/lib/AST/Type.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/Type.cpp?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/lib/AST/Type.cpp (original)
+++ cfe/trunk/lib/AST/Type.cpp Tue Apr 3 11:29:31 2018
@@ -2748,7 +2748,6 @@ StringRef FunctionType::getNameForCallCo
case CC_Swift: return "swiftcall";
case CC_PreserveMost: return "preserve_most";
case CC_PreserveAll: return "preserve_all";
- case CC_CUDAKernel: return "cuda_kernel";
}
llvm_unreachable("Invalid calling convention.");
Modified: cfe/trunk/lib/AST/TypePrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/TypePrinter.cpp?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/lib/AST/TypePrinter.cpp (original)
+++ cfe/trunk/lib/AST/TypePrinter.cpp Tue Apr 3 11:29:31 2018
@@ -780,10 +780,6 @@ void TypePrinter::printFunctionAfter(con
case CC_OpenCLKernel:
// Do nothing. These CCs are not available as attributes.
break;
- case CC_CUDAKernel:
- // ToDo: print this before the function.
- OS << " __global__";
- break;
case CC_Swift:
OS << " __attribute__((swiftcall))";
break;
Modified: cfe/trunk/lib/CodeGen/CGCall.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGCall.cpp?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGCall.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGCall.cpp Tue Apr 3 11:29:31 2018
@@ -64,7 +64,6 @@ unsigned CodeGenTypes::ClangCallConvToLL
case CC_PreserveMost: return llvm::CallingConv::PreserveMost;
case CC_PreserveAll: return llvm::CallingConv::PreserveAll;
case CC_Swift: return llvm::CallingConv::Swift;
- case CC_CUDAKernel: return CGM.getTargetCodeGenInfo().getCUDAKernelCallingConv();
}
}
Modified: cfe/trunk/lib/CodeGen/CGDebugInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDebugInfo.cpp?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGDebugInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGDebugInfo.cpp Tue Apr 3 11:29:31 2018
@@ -1022,9 +1022,6 @@ static unsigned getDwarfCC(CallingConv C
return llvm::dwarf::DW_CC_LLVM_PreserveAll;
case CC_X86RegCall:
return llvm::dwarf::DW_CC_LLVM_X86RegCall;
- case CC_CUDAKernel:
- // ToDo: Add llvm::dwarf::DW_CC_LLVM_CUDAKernel;
- return 0;
}
return 0;
}
Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.cpp?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.cpp Tue Apr 3 11:29:31 2018
@@ -431,10 +431,6 @@ unsigned TargetCodeGenInfo::getOpenCLKer
return llvm::CallingConv::SPIR_KERNEL;
}
-unsigned TargetCodeGenInfo::getCUDAKernelCallingConv() const {
- return llvm::CallingConv::C;
-}
-
llvm::Constant *TargetCodeGenInfo::getNullPointer(const CodeGen::CodeGenModule &CGM,
llvm::PointerType *T, QualType QT) const {
return llvm::ConstantPointerNull::get(T);
@@ -7639,7 +7635,6 @@ public:
void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const override;
unsigned getOpenCLKernelCallingConv() const override;
- unsigned getCUDAKernelCallingConv() const override;
llvm::Constant *getNullPointer(const CodeGen::CodeGenModule &CGM,
llvm::PointerType *T, QualType QT) const override;
@@ -7727,10 +7722,6 @@ unsigned AMDGPUTargetCodeGenInfo::getOpe
return llvm::CallingConv::AMDGPU_KERNEL;
}
-unsigned AMDGPUTargetCodeGenInfo::getCUDAKernelCallingConv() const {
- return llvm::CallingConv::AMDGPU_KERNEL;
-}
-
// Currently LLVM assumes null pointers always have value 0,
// which results in incorrectly transformed IR. Therefore, instead of
// emitting null pointers in private and local address spaces, a null
Modified: cfe/trunk/lib/CodeGen/TargetInfo.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.h?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.h (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.h Tue Apr 3 11:29:31 2018
@@ -223,9 +223,6 @@ public:
/// Get LLVM calling convention for OpenCL kernel.
virtual unsigned getOpenCLKernelCallingConv() const;
- /// Get LLVM calling convention for CUDA kernel.
- virtual unsigned getCUDAKernelCallingConv() const;
-
/// Get target specific null pointer.
/// \param T is the LLVM type of the null pointer.
/// \param QT is the clang QualType of the null pointer.
Modified: cfe/trunk/lib/Sema/SemaExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExpr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExpr.cpp Tue Apr 3 11:29:31 2018
@@ -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 @@ Sema::BuildDeclRefExpr(ValueDecl *D, Qua
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 =
Modified: cfe/trunk/lib/Sema/SemaOverload.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOverload.cpp?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOverload.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOverload.cpp Tue Apr 3 11:29:31 2018
@@ -1481,6 +1481,7 @@ bool Sema::IsFunctionConversion(QualType
.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.
Modified: cfe/trunk/lib/Sema/SemaType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaType.cpp (original)
+++ cfe/trunk/lib/Sema/SemaType.cpp Tue Apr 3 11:29:31 2018
@@ -3316,18 +3316,6 @@ getCCForDeclaratorChunk(Sema &S, Declara
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
Removed: cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu?rev=329098&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu (removed)
@@ -1,29 +0,0 @@
-// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck %s
-#include "Inputs/cuda.h"
-
-// CHECK: define amdgpu_kernel void @_ZN1A6kernelEv
-class A {
-public:
- static __global__ void kernel(){}
-};
-
-// 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 @_Z15template_kernelI1AEvT_
-template<class T>
-__global__ void template_kernel(T x) {}
-
-void launch(void *f);
-
-int main() {
- launch((void*)A::kernel);
- launch((void*)kernel);
- launch((void*)template_kernel<A>);
- return 0;
-}
Modified: cfe/trunk/tools/libclang/CXType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/tools/libclang/CXType.cpp?rev=329099&r1=329098&r2=329099&view=diff
==============================================================================
--- cfe/trunk/tools/libclang/CXType.cpp (original)
+++ cfe/trunk/tools/libclang/CXType.cpp Tue Apr 3 11:29:31 2018
@@ -626,7 +626,6 @@ CXCallingConv clang_getFunctionTypeCalli
TCALLINGCONV(PreserveAll);
case CC_SpirFunction: return CXCallingConv_Unexposed;
case CC_OpenCLKernel: return CXCallingConv_Unexposed;
- case CC_CUDAKernel: return CXCallingConv_Unexposed;
break;
}
#undef TCALLINGCONV
More information about the cfe-commits
mailing list