[PATCH] D11666: [CUDA] Make sure we emit all templated __global__ functions on device side.
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Mon Aug 10 13:57:49 PDT 2015
This revision was automatically updated to reflect the committed changes.
Closed by commit rL244501: [CUDA] Add implicit __attribute__((used)) to all __global__ functions. (authored by tra).
Changed prior to commit:
http://reviews.llvm.org/D11666?vs=31128&id=31718#toc
Repository:
rL LLVM
http://reviews.llvm.org/D11666
Files:
cfe/trunk/lib/Sema/SemaDeclAttr.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
@@ -1,7 +1,16 @@
+// Make sure that __global__ functions are emitted along with correct
+// annotations and are added to @llvm.used to prevent their elimination.
+// REQUIRES: nvptx-registered-target
+//
// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
#include "Inputs/cuda.h"
+// Make sure that all __global__ functiona 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() {}
@@ -13,4 +22,10 @@
device_function();
}
+// 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_
+void host_function() { templated_kernel<<<0,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/Sema/SemaDeclAttr.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp
@@ -3350,6 +3350,10 @@
D->addAttr(::new (S.Context)
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) {
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D11666.31718.patch
Type: text/x-patch
Size: 1942 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20150810/59d0ee83/attachment.bin>
More information about the cfe-commits
mailing list