Revert r245496 "[CUDA] Add appropriate host/device attribute to builtins."

It's breaking internal test.

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

diff --git a/include/clang/Basic/Builtins.h b/include/clang/Basic/Builtins.h
index 87e2ac70790bd6e6efd8d564dd91656f7394ed8c..554143d6be7d7d0f728e759d34ec935a1920967f 100644 (file)
--- a/include/clang/Basic/Builtins.h
+++ b/include/clang/Basic/Builtins.h
@@ -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 {

diff --git a/lib/Sema/SemaChecking.cpp b/lib/Sema/SemaChecking.cpp
index e2940c7f57ef67f1d4a66ea2d4ae09209372318b..a8a7009ccf5bfcef34cddd5f84121b07015fd99b 100644 (file)
--- a/lib/Sema/SemaChecking.cpp
+++ b/lib/Sema/SemaChecking.cpp
@@ -525,7 +525,7 @@ Sema::CheckBuiltinFunctionCall(FunctionDecl *FDecl, unsigned BuiltinID,
 
   // 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:

diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp
index d79d60c9773bce52b4d7bca2488f6e91c6d068dd..a8d1e1203e45a1fce7afb2a1d55cf2a1238dce6c 100644 (file)
--- a/lib/Sema/SemaDecl.cpp
+++ b/lib/Sema/SemaDecl.cpp
@@ -11187,17 +11187,6 @@ void Sema::AddKnownFunctionAttributes(FunctionDecl *FD) {
       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();

diff --git a/test/SemaCUDA/builtins.cu b/test/SemaCUDA/builtins.cu
deleted file mode 100644 (file)
index 80b9d69..0000000
--- a/test/SemaCUDA/builtins.cu
+++ /dev/null
@@ -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

diff --git a/test/SemaCUDA/implicit-intrinsic.cu b/test/SemaCUDA/implicit-intrinsic.cu
index 7414a66b03cc9a2ed76f70423499c5c713dad06c..3d24aa719e57d902624925291e807e7ed686cd6e 100644 (file)
--- a/test/SemaCUDA/implicit-intrinsic.cu
+++ b/test/SemaCUDA/implicit-intrinsic.cu
@@ -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();
 }