Author: djasper
Date: Tue Aug 11 06:02:09 2015
New Revision: 244583
URL: http://llvm.org/viewvc/llvm-project?rev=244583&view=rev
Log:
Revert "[CUDA] Add implicit __attribute__((used)) to all __global__ functions."
This is breaking internal test. I'll provide a reproduction.
Modified:
cfe/trunk/lib/Sema/SemaDeclAttr.cpp
cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=244583&r1=244582&r2=244583&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Tue Aug 11 06:02:09 2015
@@ -3350,10 +3350,6 @@ static void handleGlobalAttr(Sema &S, De
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) {
Modified: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu?rev=244583&r1=244582&r2=244583&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu Tue Aug 11 06:02:09 2015
@@ -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}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
