This revision was automatically updated to reflect the committed changes. Closed by commit rL248400: [CUDA] __global__ functions should always be visible externally. (authored by tra).
Changed prior to commit: http://reviews.llvm.org/D13067?vs=35412&id=35529#toc Repository: rL LLVM http://reviews.llvm.org/D13067 Files: cfe/trunk/lib/AST/ASTContext.cpp cfe/trunk/lib/Sema/SemaDeclAttr.cpp cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu Index: cfe/trunk/lib/AST/ASTContext.cpp =================================================================== --- cfe/trunk/lib/AST/ASTContext.cpp +++ cfe/trunk/lib/AST/ASTContext.cpp @@ -8284,22 +8284,22 @@ return GVA_DiscardableODR; } -static GVALinkage adjustGVALinkageForDLLAttribute(GVALinkage L, const Decl *D) { +static GVALinkage adjustGVALinkageForAttributes(GVALinkage L, const Decl *D) { // See http://msdn.microsoft.com/en-us/library/xa0d9ste.aspx // dllexport/dllimport on inline functions. if (D->hasAttr<DLLImportAttr>()) { if (L == GVA_DiscardableODR || L == GVA_StrongODR) return GVA_AvailableExternally; - } else if (D->hasAttr<DLLExportAttr>()) { + } else if (D->hasAttr<DLLExportAttr>() || D->hasAttr<CUDAGlobalAttr>()) { if (L == GVA_DiscardableODR) return GVA_StrongODR; } return L; } GVALinkage ASTContext::GetGVALinkageForFunction(const FunctionDecl *FD) const { - return adjustGVALinkageForDLLAttribute(basicGVALinkageForFunction(*this, FD), - FD); + return adjustGVALinkageForAttributes(basicGVALinkageForFunction(*this, FD), + FD); } static GVALinkage basicGVALinkageForVariable(const ASTContext &Context, @@ -8355,8 +8355,8 @@ } GVALinkage ASTContext::GetGVALinkageForVariable(const VarDecl *VD) { - return adjustGVALinkageForDLLAttribute(basicGVALinkageForVariable(*this, VD), - VD); + return adjustGVALinkageForAttributes(basicGVALinkageForVariable(*this, VD), + VD); } bool ASTContext::DeclMustBeEmitted(const Decl *D) { Index: cfe/trunk/lib/Sema/SemaDeclAttr.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaDeclAttr.cpp +++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp @@ -3351,9 +3351,6 @@ CUDAGlobalAttr(Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex())); - // Add implicit attribute((used)) so we don't eliminate kernels - // because there is nothing referencing them on device side. - D->addAttr(UsedAttr::CreateImplicit(S.Context)); } static void handleGNUInlineAttr(Sema &S, Decl *D, const AttributeList &Attr) { Index: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu =================================================================== --- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu +++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu @@ -6,11 +6,6 @@ #include "Inputs/cuda.h" -// Make sure that all __global__ functions are added to @llvm.used -// CHECK: @llvm.used = appending global -// CHECK-SAME: @global_function -// CHECK-SAME: @_Z16templated_kernelIiEvT_ - // CHECK-LABEL: define void @device_function extern "C" __device__ void device_function() {} @@ -24,7 +19,7 @@ // Make sure host-instantiated kernels are preserved on device side. template <typename T> __global__ void templated_kernel(T param) {} -// CHECK-LABEL: define linkonce_odr void @_Z16templated_kernelIiEvT_ +// CHECK-LABEL: define weak_odr void @_Z16templated_kernelIiEvT_ void host_function() { templated_kernel<<<0,0>>>(0); } // CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}
Index: cfe/trunk/lib/AST/ASTContext.cpp =================================================================== --- cfe/trunk/lib/AST/ASTContext.cpp +++ cfe/trunk/lib/AST/ASTContext.cpp @@ -8284,22 +8284,22 @@ return GVA_DiscardableODR; } -static GVALinkage adjustGVALinkageForDLLAttribute(GVALinkage L, const Decl *D) { +static GVALinkage adjustGVALinkageForAttributes(GVALinkage L, const Decl *D) { // See http://msdn.microsoft.com/en-us/library/xa0d9ste.aspx // dllexport/dllimport on inline functions. if (D->hasAttr<DLLImportAttr>()) { if (L == GVA_DiscardableODR || L == GVA_StrongODR) return GVA_AvailableExternally; - } else if (D->hasAttr<DLLExportAttr>()) { + } else if (D->hasAttr<DLLExportAttr>() || D->hasAttr<CUDAGlobalAttr>()) { if (L == GVA_DiscardableODR) return GVA_StrongODR; } return L; } GVALinkage ASTContext::GetGVALinkageForFunction(const FunctionDecl *FD) const { - return adjustGVALinkageForDLLAttribute(basicGVALinkageForFunction(*this, FD), - FD); + return adjustGVALinkageForAttributes(basicGVALinkageForFunction(*this, FD), + FD); } static GVALinkage basicGVALinkageForVariable(const ASTContext &Context, @@ -8355,8 +8355,8 @@ } GVALinkage ASTContext::GetGVALinkageForVariable(const VarDecl *VD) { - return adjustGVALinkageForDLLAttribute(basicGVALinkageForVariable(*this, VD), - VD); + return adjustGVALinkageForAttributes(basicGVALinkageForVariable(*this, VD), + VD); } bool ASTContext::DeclMustBeEmitted(const Decl *D) { Index: cfe/trunk/lib/Sema/SemaDeclAttr.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaDeclAttr.cpp +++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp @@ -3351,9 +3351,6 @@ CUDAGlobalAttr(Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex())); - // Add implicit attribute((used)) so we don't eliminate kernels - // because there is nothing referencing them on device side. - D->addAttr(UsedAttr::CreateImplicit(S.Context)); } static void handleGNUInlineAttr(Sema &S, Decl *D, const AttributeList &Attr) { Index: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu =================================================================== --- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu +++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu @@ -6,11 +6,6 @@ #include "Inputs/cuda.h" -// Make sure that all __global__ functions are added to @llvm.used -// CHECK: @llvm.used = appending global -// CHECK-SAME: @global_function -// CHECK-SAME: @_Z16templated_kernelIiEvT_ - // CHECK-LABEL: define void @device_function extern "C" __device__ void device_function() {} @@ -24,7 +19,7 @@ // Make sure host-instantiated kernels are preserved on device side. template <typename T> __global__ void templated_kernel(T param) {} -// CHECK-LABEL: define linkonce_odr void @_Z16templated_kernelIiEvT_ +// CHECK-LABEL: define weak_odr void @_Z16templated_kernelIiEvT_ void host_function() { templated_kernel<<<0,0>>>(0); } // CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits