Author: tra
Date: Thu Aug 20 13:28:56 2015
New Revision: 245592
URL: http://llvm.org/viewvc/llvm-project?rev=245592&view=rev
Log:
Revert r245496 "[CUDA] Add appropriate host/device attribute to builtins."
It's breaking internal test.
Removed:
cfe/trunk/test/SemaCUDA/builtins.cu
Modified:
cfe/trunk/include/clang/Basic/Builtins.h
cfe/trunk/lib/Sema/SemaChecking.cpp
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/test/SemaCUDA/implicit-intrinsic.cu
Modified: cfe/trunk/include/clang/Basic/Builtins.h
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Builtins.h?rev=245592&r1=245591&r2=245592&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Builtins.h (original)
+++ cfe/trunk/include/clang/Basic/Builtins.h Thu Aug 20 13:28:56 2015
@@ -81,11 +81,6 @@ public:
return getRecord(ID).Type;
}
- /// \brief Return true if this function is a target-specific builtin
- bool isTSBuiltin(unsigned ID) const {
- return ID >= Builtin::FirstTSBuiltin;
- }
-
/// \brief Return true if this function has no side effects and doesn't
/// read memory.
bool isConst(unsigned ID) const {
Modified: cfe/trunk/lib/Sema/SemaChecking.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaChecking.cpp?rev=245592&r1=245591&r2=245592&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaChecking.cpp (original)
+++ cfe/trunk/lib/Sema/SemaChecking.cpp Thu Aug 20 13:28:56 2015
@@ -525,7 +525,7 @@ Sema::CheckBuiltinFunctionCall(FunctionD
// Since the target specific builtins for each arch overlap, only check those
// of the arch we are compiling for.
- if (Context.BuiltinInfo.isTSBuiltin(BuiltinID)) {
+ if (BuiltinID >= Builtin::FirstTSBuiltin) {
switch (Context.getTargetInfo().getTriple().getArch()) {
case llvm::Triple::arm:
case llvm::Triple::armeb:
Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=245592&r1=245591&r2=245592&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Thu Aug 20 13:28:56 2015
@@ -11187,17 +11187,6 @@ void Sema::AddKnownFunctionAttributes(Fu
FD->addAttr(NoThrowAttr::CreateImplicit(Context, FD->getLocation()));
if (Context.BuiltinInfo.isConst(BuiltinID) && !FD->hasAttr<ConstAttr>())
FD->addAttr(ConstAttr::CreateImplicit(Context, FD->getLocation()));
- if (getLangOpts().CUDA && Context.BuiltinInfo.isTSBuiltin(BuiltinID) &&
- !FD->hasAttr<CUDADeviceAttr>() && !FD->hasAttr<CUDAHostAttr>()) {
- // Target-specific builtins are assumed to be intended for use
- // in this particular CUDA compilation mode and should have
- // appropriate attribute set so we can enforce CUDA function
- // call restrictions.
- if (getLangOpts().CUDAIsDevice)
- FD->addAttr(CUDADeviceAttr::CreateImplicit(Context,
FD->getLocation()));
- else
- FD->addAttr(CUDAHostAttr::CreateImplicit(Context, FD->getLocation()));
- }
}
IdentifierInfo *Name = FD->getIdentifier();
Removed: cfe/trunk/test/SemaCUDA/builtins.cu
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/builtins.cu?rev=245591&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/builtins.cu (original)
+++ cfe/trunk/test/SemaCUDA/builtins.cu (removed)
@@ -1,35 +0,0 @@
-// Tests that target-specific builtins have appropriate host/device
-// attributes and that CUDA call restrictions are enforced. Also
-// verify that non-target builtins can be used from both host and
-// device functions.
-//
-// REQUIRES: x86-registered-target
-// REQUIRES: nvptx-registered-target
-// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fsyntax-only -verify %s
-// RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \
-// RUN: -fsyntax-only -verify %s
-
-
-#ifdef __CUDA_ARCH__
-// Device-side builtins are not allowed to be called from host functions.
-void hf() {
- int x = __builtin_ptx_read_tid_x(); // expected-note
{{'__builtin_ptx_read_tid_x' declared here}}
- // expected-error@-1 {{reference to __device__ function
'__builtin_ptx_read_tid_x' in __host__ function}}
- x = __builtin_abs(1);
-}
-__attribute__((device)) void df() {
- int x = __builtin_ptx_read_tid_x();
- x = __builtin_abs(1);
-}
-#else
-// Host-side builtins are not allowed to be called from device functions.
-__attribute__((device)) void df() {
- int x = __builtin_ia32_rdtsc(); // expected-note {{'__builtin_ia32_rdtsc'
declared here}}
- // expected-error@-1 {{reference to __host__ function '__builtin_ia32_rdtsc'
in __device__ function}}
- x = __builtin_abs(1);
-}
-void hf() {
- int x = __builtin_ia32_rdtsc();
- x = __builtin_abs(1);
-}
-#endif
Modified: cfe/trunk/test/SemaCUDA/implicit-intrinsic.cu
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/implicit-intrinsic.cu?rev=245592&r1=245591&r2=245592&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/implicit-intrinsic.cu (original)
+++ cfe/trunk/test/SemaCUDA/implicit-intrinsic.cu Thu Aug 20 13:28:56 2015
@@ -1,10 +1,10 @@
-// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device
-fsyntax-only -verify %s
+// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fsyntax-only
-verify %s
#include "Inputs/cuda.h"
// expected-no-diagnostics
__device__ void __threadfence_system() {
- // This shouldn't produce an error, since __nvvm_membar_sys should be
- // __device__ and thus callable from device code.
+ // This shouldn't produce an error, since __nvvm_membar_sys is inferred to
+ // be __host__ __device__ and thus callable from device code.
__nvvm_membar_sys();
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
