[PATCH] D12122: [CUDA] Add appropriate host/device attribute to target-specific builtins.

[Artem Belevich via cfe-commits]

tra created this revision.
tra added reviewers: eliben, echristo.
tra added a subscriber: cfe-commits.

The patch adds appropriate __host__ or __device__ attributes to target-specific builtins 
so we can properly check whether they may or may not be called from particular context.


http://reviews.llvm.org/D12122

Files:
  include/clang/Basic/Builtins.h
  lib/Sema/SemaDecl.cpp
  test/SemaCUDA/builtins.cu
  test/SemaCUDA/implicit-intrinsic.cu

Index: test/SemaCUDA/implicit-intrinsic.cu
===================================================================
--- test/SemaCUDA/implicit-intrinsic.cu
+++ test/SemaCUDA/implicit-intrinsic.cu
@@ -1,10 +1,10 @@
-// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fcuda-is-device -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 is inferred to
-  // be __host__ __device__ and thus callable from device code.
+  // This shouldn't produce an error, since __nvvm_membar_sys should be
+  // __device__ and thus callable from device code.
   __nvvm_membar_sys();
 }
Index: test/SemaCUDA/builtins.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/builtins.cu
@@ -0,0 +1,35 @@
+// 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 at -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 at -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
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -11161,6 +11161,13 @@
       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>()) {
+      if (getLangOpts().CUDAIsDevice)
+        FD->addAttr(CUDADeviceAttr::CreateImplicit(Context, FD->getLocation()));
+      else
+        FD->addAttr(CUDAHostAttr::CreateImplicit(Context, FD->getLocation()));
+    }
   }
 
   IdentifierInfo *Name = FD->getIdentifier();
Index: include/clang/Basic/Builtins.h
===================================================================
--- include/clang/Basic/Builtins.h
+++ include/clang/Basic/Builtins.h
@@ -81,6 +81,11 @@
     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 {


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D12122.32463.patch
Type: text/x-patch
Size: 3680 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20150818/1b5a0be1/attachment.bin>