[OpenMP] Implement omp_is_initial_device() as builtin

[Hahnfeld]

Summary

This allows to return the static value that we know at compile time.

[ABataev]

LG

[grokos]

Now that this issue has been addressed and regressions tests pass, should we re-enable Cmake to build libomptarget by default?

[Hahnfeld]

In D38968#898951, @grokos wrote:

Now that this issue has been addressed and regressions tests pass, should we re-enable Cmake to build libomptarget by default?

Yes, I already have a local patch which also takes care of restricting the tests to Clang versions newer than 6.0.0. I will post it for review once I've committed this revision.

[JonChesterfield]

@jdoerfert suggested replacing this with a context selector. One less special case in clang.

Related, if this is a builtin (or in header), I think we can and should delete it from the devicertl. It's probably used as a branch condition so replacing with a 0/1 asap is winning for compilation time.

include/clang/Basic/Builtins.h

 enum LanguageID {
   GNU_LANG = 0x1,     // builtin requires GNU mode.
   C_LANG = 0x2,       // builtin for c only.
   CXX_LANG = 0x4,     // builtin for cplusplus only.
   OBJC_LANG = 0x8,    // builtin for objective-c and objective-c++
   MS_LANG = 0x10,     // builtin requires MS mode.
   OCLC20_LANG = 0x20, // builtin for OpenCL C 2.0 only.
   OCLC1X_LANG = 0x40, // builtin for OpenCL C 1.x only.
+  OMP_LANG = 0x80,    // builtin requires OpenMP.
   ALL_LANGUAGES = C_LANG | CXX_LANG | OBJC_LANG, // builtin for all languages.
   ALL_GNU_LANGUAGES = ALL_LANGUAGES | GNU_LANG,  // builtin requires GNU mode.
   ALL_MS_LANGUAGES = ALL_LANGUAGES | MS_LANG,    // builtin requires MS mode.
   ALL_OCLC_LANGUAGES = OCLC1X_LANG | OCLC20_LANG // builtin for OCLC languages.
 };
 
 namespace Builtin {
 enum ID {

include/clang/Basic/Builtins.def

 LANGBUILTIN(__builtin_store_halff, "vfh*", "n", ALL_OCLC_LANGUAGES)
 LANGBUILTIN(__builtin_load_half, "dhC*", "nc", ALL_OCLC_LANGUAGES)
 LANGBUILTIN(__builtin_load_halff, "fhC*", "nc", ALL_OCLC_LANGUAGES)
 
 // Builtins for os_log/os_trace
 BUILTIN(__builtin_os_log_format_buffer_size, "zcC*.", "p:0:nut")
 BUILTIN(__builtin_os_log_format, "v*v*cC*.", "p:0:nt")
 
+// OpenMP 4.0
+LANGBUILTIN(omp_is_initial_device, "i", "nc", OMP_LANG)
+
 // Builtins for XRay
 BUILTIN(__xray_customevent, "vcC*z", "")
 
 // Win64-compatible va_list functions
 BUILTIN(__builtin_ms_va_start, "vc*&.", "nt")
 BUILTIN(__builtin_ms_va_end, "vc*&", "n")
 BUILTIN(__builtin_ms_va_copy, "vc*&c*&", "n")
 
 #undef BUILTIN
 #undef LIBBUILTIN
 #undef LANGBUILTIN

lib/AST/ExprConstant.cpp

if (Size.isPowerOfTwo()) {
           return Success(1, E);
         }
       }
     }
 
     return BuiltinOp == Builtin::BI__atomic_always_lock_free ?
         Success(0, E) : Error(E);
   }
+  case Builtin::BIomp_is_initial_device:
+    // We can decide statically which value the runtime would return if called.
+    return Success(Info.getLangOpts().OpenMPIsDevice ? 0 : 1, E);
   }
 }
 
 static bool HasSameBase(const LValue &A, const LValue &B) {
   if (!A.getLValueBase())
     return !B.getLValueBase();
   if (!B.getLValueBase())
     return false;

lib/Basic/Builtins.cpp

bool MSModeUnsupported =
       !LangOpts.MicrosoftExt && (BuiltinInfo.Langs & MS_LANG);
   bool ObjCUnsupported = !LangOpts.ObjC1 && BuiltinInfo.Langs == OBJC_LANG;
   bool OclC1Unsupported = (LangOpts.OpenCLVersion / 100) != 1 &&
                           (BuiltinInfo.Langs & ALL_OCLC_LANGUAGES ) ==  OCLC1X_LANG;
   bool OclC2Unsupported = LangOpts.OpenCLVersion != 200 &&
                           (BuiltinInfo.Langs & ALL_OCLC_LANGUAGES) == OCLC20_LANG;
   bool OclCUnsupported = !LangOpts.OpenCL &&
                          (BuiltinInfo.Langs & ALL_OCLC_LANGUAGES);
+  bool OpenMPUnsupported = !LangOpts.OpenMP && BuiltinInfo.Langs == OMP_LANG;
   return !BuiltinsUnsupported && !MathBuiltinsUnsupported && !OclCUnsupported &&
-         !OclC1Unsupported && !OclC2Unsupported &&
+         !OclC1Unsupported && !OclC2Unsupported && !OpenMPUnsupported &&
          !GnuModeUnsupported && !MSModeUnsupported && !ObjCUnsupported;
 }
 
 /// initializeBuiltins - Mark the identifiers for all the builtins with their
 /// appropriate builtin ID # and mark any non-portable builtin identifiers as
 /// such.
 void Builtin::Context::initializeBuiltins(IdentifierTable &Table,
                                           const LangOptions& LangOpts) {

test/OpenMP/is_initial_device.c

+// REQUIRES: powerpc-registered-target
+
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown \
+// RUN:            -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x ir -triple powerpc64le-unknown-unknown -emit-llvm \
+// RUN:             %t-ppc-host.bc -o - | FileCheck %s -check-prefixes HOST,OUTLINED
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple powerpc64le-unknown-unknown -emit-llvm -fopenmp-is-device \
+// RUN:             %s -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s -check-prefixes DEVICE,OUTLINED
+
+// expected-no-diagnostics
+int check() {
+  int host = omp_is_initial_device();
+  int device;
+#pragma omp target map(tofrom: device)
+  {
+    device = omp_is_initial_device();
+  }
+
+  return host + device;
+}
+
+// The host should get a value of 1:
+// HOST: define{{.*}} @check()
+// HOST: [[HOST:%.*]] = alloca i32
+// HOST: store i32 1, i32* [[HOST]]
+
+// OUTLINED: define{{.*}} @{{.*}}omp_offloading{{.*}}(i32*{{.*}} [[DEVICE_ARGUMENT:%.*]])
+// OUTLINED: [[DEVICE_ADDR_STORAGE:%.*]] = alloca i32*
+// OUTLINED: store i32* [[DEVICE_ARGUMENT]], i32** [[DEVICE_ADDR_STORAGE]]
+// OUTLINED: [[DEVICE_ADDR:%.*]] = load i32*, i32** [[DEVICE_ADDR_STORAGE]]
+
+// The outlined function that is called as fallback also runs on the host:
+// HOST: store i32 1, i32* [[DEVICE_ADDR]]
+
+// The device should get a value of 0:
+// DEVICE: store i32 0, i32* [[DEVICE_ADDR]]