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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits