This revision was automatically updated to reflect the committed changes.
Closed by commit rL268299: [CUDA] Make sure device-side __global__ functions
are always visible. (authored by tra).
Changed prior to commit:
http://reviews.llvm.org/D19748?vs=55674&id=55885#toc
Repository:
rL LLVM
http://reviews.llvm.org/D19748
Files:
cfe/trunk/lib/AST/ASTContext.cpp
cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
Index: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
+++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
@@ -19,8 +19,17 @@
// Make sure host-instantiated kernels are preserved on device side.
template <typename T> __global__ void templated_kernel(T param) {}
-// CHECK-LABEL: define weak_odr void @_Z16templated_kernelIiEvT_
-void host_function() { templated_kernel<<<0,0>>>(0); }
+// CHECK-DAG: define weak_odr void @_Z16templated_kernelIiEvT_(
+
+namespace {
+__global__ void anonymous_ns_kernel() {}
+// CHECK-DAG: define weak_odr void @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv(
+}
+
+void host_function() {
+ templated_kernel<<<0, 0>>>(0);
+ anonymous_ns_kernel<<<0,0>>>();
+}
// CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}
// CHECK: !{{[0-9]+}} = !{void (i32)* @_Z16templated_kernelIiEvT_, !"kernel",
i32 1}
Index: cfe/trunk/lib/AST/ASTContext.cpp
===================================================================
--- cfe/trunk/lib/AST/ASTContext.cpp
+++ cfe/trunk/lib/AST/ASTContext.cpp
@@ -8418,22 +8418,29 @@
return GVA_DiscardableODR;
}
-static GVALinkage adjustGVALinkageForAttributes(GVALinkage L, const Decl *D) {
+static GVALinkage adjustGVALinkageForAttributes(const ASTContext &Context,
+ 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>() || D->hasAttr<CUDAGlobalAttr>()) {
+ } else if (D->hasAttr<DLLExportAttr>()) {
if (L == GVA_DiscardableODR)
return GVA_StrongODR;
+ } else if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice
&&
+ D->hasAttr<CUDAGlobalAttr>()) {
+ // Device-side functions with __global__ attribute must always be
+ // visible externally so they can be launched from host.
+ if (L == GVA_DiscardableODR || L == GVA_Internal)
+ return GVA_StrongODR;
}
return L;
}
GVALinkage ASTContext::GetGVALinkageForFunction(const FunctionDecl *FD) const {
- return adjustGVALinkageForAttributes(basicGVALinkageForFunction(*this, FD),
- FD);
+ return adjustGVALinkageForAttributes(
+ *this, basicGVALinkageForFunction(*this, FD), FD);
}
static GVALinkage basicGVALinkageForVariable(const ASTContext &Context,
@@ -8490,8 +8497,8 @@
}
GVALinkage ASTContext::GetGVALinkageForVariable(const VarDecl *VD) {
- return adjustGVALinkageForAttributes(basicGVALinkageForVariable(*this, VD),
- VD);
+ return adjustGVALinkageForAttributes(
+ *this, basicGVALinkageForVariable(*this, VD), VD);
}
bool ASTContext::DeclMustBeEmitted(const Decl *D) {
Index: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
+++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
@@ -19,8 +19,17 @@
// Make sure host-instantiated kernels are preserved on device side.
template <typename T> __global__ void templated_kernel(T param) {}
-// CHECK-LABEL: define weak_odr void @_Z16templated_kernelIiEvT_
-void host_function() { templated_kernel<<<0,0>>>(0); }
+// CHECK-DAG: define weak_odr void @_Z16templated_kernelIiEvT_(
+
+namespace {
+__global__ void anonymous_ns_kernel() {}
+// CHECK-DAG: define weak_odr void @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv(
+}
+
+void host_function() {
+ templated_kernel<<<0, 0>>>(0);
+ anonymous_ns_kernel<<<0,0>>>();
+}
// CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}
// CHECK: !{{[0-9]+}} = !{void (i32)* @_Z16templated_kernelIiEvT_, !"kernel", i32 1}
Index: cfe/trunk/lib/AST/ASTContext.cpp
===================================================================
--- cfe/trunk/lib/AST/ASTContext.cpp
+++ cfe/trunk/lib/AST/ASTContext.cpp
@@ -8418,22 +8418,29 @@
return GVA_DiscardableODR;
}
-static GVALinkage adjustGVALinkageForAttributes(GVALinkage L, const Decl *D) {
+static GVALinkage adjustGVALinkageForAttributes(const ASTContext &Context,
+ 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>() || D->hasAttr<CUDAGlobalAttr>()) {
+ } else if (D->hasAttr<DLLExportAttr>()) {
if (L == GVA_DiscardableODR)
return GVA_StrongODR;
+ } else if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice &&
+ D->hasAttr<CUDAGlobalAttr>()) {
+ // Device-side functions with __global__ attribute must always be
+ // visible externally so they can be launched from host.
+ if (L == GVA_DiscardableODR || L == GVA_Internal)
+ return GVA_StrongODR;
}
return L;
}
GVALinkage ASTContext::GetGVALinkageForFunction(const FunctionDecl *FD) const {
- return adjustGVALinkageForAttributes(basicGVALinkageForFunction(*this, FD),
- FD);
+ return adjustGVALinkageForAttributes(
+ *this, basicGVALinkageForFunction(*this, FD), FD);
}
static GVALinkage basicGVALinkageForVariable(const ASTContext &Context,
@@ -8490,8 +8497,8 @@
}
GVALinkage ASTContext::GetGVALinkageForVariable(const VarDecl *VD) {
- return adjustGVALinkageForAttributes(basicGVALinkageForVariable(*this, VD),
- VD);
+ return adjustGVALinkageForAttributes(
+ *this, basicGVALinkageForVariable(*this, VD), VD);
}
bool ASTContext::DeclMustBeEmitted(const Decl *D) {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
