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

[tra]

Summary

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.

[eliben]

include/clang/Basic/Builtins.h

85

You can also use it in SemaChecking now

lib/Sema/SemaDecl.cpp

11166

It is not immediately clear why you would mark target-specific builtins as host in host compilation mode. So for example __builtin_ptx_read_tid_x would be callable from a host function when compiling in host mode? Can you clarify this with a comment here, and also add a relevant test?

[tra]

used isTSBuiltin in SemaChecking.cpp

[tra]

Added a comment explaining reasoning behind attribute choice for target-specific builtins.

[tra]

lib/Sema/SemaDecl.cpp

11166

Target triple used for particular compilation pass is assumed to be valid for this particular compilation mode. Builtins provided by the target are therefore assumed to be intended to be used in that compilation mode. builtins.cu already includes test cases for using target-specific builtins from different targets and and non-target-specific builtins in host and device functions in host and device compilation.

Your example scenario (builtin_ptx_read_tid_x getting host attribute) is possible only if you compile in host mode *and* use --triple nvptx-unknown-cuda. IMO, host__ would be an appropriate attribute to assign to builtins in this scenario, even though I don't think we can currently do anything useful with NVPTX as the host at the moment. If you attempt to use __builtin_ptx_read_tid_x() during host compilation using non-NVPTX target (which is a typical scenario), then compilation will fail because that particular builtin would not be available at all.

That said, another corner case would be compiling CUDA for device with the same architecture as host. Again, it's not a practical scenario at the moment. If/when we get to the point when we may want to do that, it should be easy to check for it and treat builtins as host device which would reflect their intended use domain.

[eliben]

lgtm

[tra]

Reverted in r245592 due to breaking internal tests.

include/clang/Basic/Builtins.h

const char *getName(unsigned ID) const {
     return getRecord(ID).Name;
   }
 
   /// \brief Get the type descriptor string for the specified builtin.
   const char *getTypeString(unsigned ID) const {
     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 {
     return strchr(getRecord(ID).Attributes, 'c') != nullptr;
   }
 
   /// \brief Return true if we know this builtin never throws an exception.
   bool isNoThrow(unsigned ID) const {

lib/Sema/SemaChecking.cpp

case Builtin::BI__GetExceptionInfo:
 
     TheCall->setType(Context.VoidPtrTy);
     break;
 
   }
 
   // Since the target specific builtins for each arch overlap, only check those
   // of the arch we are compiling for.
-  if (BuiltinID >= Builtin::FirstTSBuiltin) {
+  if (Context.BuiltinInfo.isTSBuiltin(BuiltinID)) {
     switch (Context.getTargetInfo().getTriple().getArch()) {
       case llvm::Triple::arm:
       case llvm::Triple::armeb:
       case llvm::Triple::thumb:
       case llvm::Triple::thumbeb:
         if (CheckARMBuiltinFunctionCall(BuiltinID, TheCall))
           return ExprError();
         break;

lib/Sema/SemaDecl.cpp

if (unsigned BuiltinID = FD->getBuiltinID()) {
     if (Context.BuiltinInfo.isReturnsTwice(BuiltinID) &&
         !FD->hasAttr<ReturnsTwiceAttr>())
       FD->addAttr(ReturnsTwiceAttr::CreateImplicit(Context,
                                          FD->getLocation()));
     if (Context.BuiltinInfo.isNoThrow(BuiltinID) && !FD->hasAttr<NoThrowAttr>())
       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();
   if (!Name)
     return;
   if ((!getLangOpts().CPlusPlus &&
        FD->getDeclContext()->isTranslationUnit()) ||
       (isa<LinkageSpecDecl>(FD->getDeclContext()) &&

test/SemaCUDA/builtins.cu

+// 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

test/SemaCUDA/implicit-intrinsic.cu

-// 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();
 }