summaryrefslogtreecommitdiffstats
path: root/clang/test/CodeGenCUDA
diff options
context:
space:
mode:
authorDaniel Jasper <djasper@google.com>2015-08-11 11:02:09 +0000
committerDaniel Jasper <djasper@google.com>2015-08-11 11:02:09 +0000
commit3b0f87d289b3cd432218cc8e5568e9eb1f77119d (patch)
tree7ac41a4e628b92c33551e54425444cb99d07f81c /clang/test/CodeGenCUDA
parentf6bc15d763539b401ad740f0120ae96efb6ca4da (diff)
downloadbcm5719-llvm-3b0f87d289b3cd432218cc8e5568e9eb1f77119d.tar.gz
bcm5719-llvm-3b0f87d289b3cd432218cc8e5568e9eb1f77119d.zip
Revert "[CUDA] Add implicit __attribute__((used)) to all __global__ functions."
This is breaking internal test. I'll provide a reproduction. llvm-svn: 244583
Diffstat (limited to 'clang/test/CodeGenCUDA')
-rw-r--r--clang/test/CodeGenCUDA/ptx-kernels.cu15
1 files changed, 0 insertions, 15 deletions
diff --git a/clang/test/CodeGenCUDA/ptx-kernels.cu b/clang/test/CodeGenCUDA/ptx-kernels.cu
index bf3b14d3a28..658b3488fc1 100644
--- a/clang/test/CodeGenCUDA/ptx-kernels.cu
+++ b/clang/test/CodeGenCUDA/ptx-kernels.cu
@@ -1,16 +1,7 @@
-// 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() {}
@@ -22,10 +13,4 @@ __global__ void global_function() {
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}
OpenPOWER on IntegriCloud