From: Daniel Jasper <djasper@google.com>
Date: Tue, 11 Aug 2015 11:02:09 +0000 (+0000)
Subject: Revert "[CUDA] Add implicit __attribute__((used)) to all __global__ functions."
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=a92ab3bb56427539c2936680f075ab675be75258;p=clang

Revert "[CUDA] Add implicit __attribute__((used)) to all __global__ functions."

This is breaking internal test. I'll provide a reproduction.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@244583 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index 8e04b69f81..31ddd503af 100644
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/lib/Sema/SemaDeclAttr.cpp
@@ -3350,10 +3350,6 @@ static void handleGlobalAttr(Sema &S, Decl *D, const AttributeList &Attr) {
   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) {
diff --git a/test/CodeGenCUDA/ptx-kernels.cu b/test/CodeGenCUDA/ptx-kernels.cu
index bf3b14d3a2..658b3488fc 100644
--- a/test/CodeGenCUDA/ptx-kernels.cu
+++ b/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}