[OpenMP] Define omp_is_initial_device() variants in omp.h

omp_is_initial_device() is marked as a built-in function in the current
compiler, and user code guarded by this call may be optimized away,
resulting in undesired behavior in some cases. This patch provides a
possible fix for such cases by defining the routine as a variant
function and removing it from builtin list.

Differential Revision: https://reviews.llvm.org/D99447

clang/include/clang/Basic/Builtins.def

@@ -1636,9 +1636,6 @@ LANGBUILTIN(__builtin_load_halff, "fhC*", "nc", ALL_OCLC_LANGUAGES)
 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)
-
 // CUDA/HIP
 LANGBUILTIN(__builtin_get_device_side_mangled_name, "cC*.", "ncT", CUDA_LANG)
 

clang/lib/AST/ExprConstant.cpp

@@ -12010,9 +12010,6 @@ bool IntExprEvaluator::VisitBuiltinCallExpr(const CallExpr *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);
   case Builtin::BI__builtin_add_overflow:
   case Builtin::BI__builtin_sub_overflow:
   case Builtin::BI__builtin_mul_overflow:

clang/test/OpenMP/is_initial_device.c

@@ -1,41 +0,0 @@
-// 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
-
-// RUN: %clang_cc1 -verify -fopenmp-simd -x c -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-unknown-unknown -emit-llvm-bc %s -o %t-ppc-host.bc
-// RUN: %clang_cc1 -verify -fopenmp-simd -x ir -triple powerpc64le-unknown-unknown -emit-llvm %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// RUN: %clang_cc1 -verify -fopenmp-simd -x c -triple powerpc64le-unknown-unknown -emit-llvm -fopenmp-is-device %s -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s
-// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
-
-// 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]]

openmp/libomptarget/test/api/is_initial_device.c

@@ -0,0 +1,30 @@
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -DUNUSED -Wall -Werror
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int errors = 0;
+#ifdef UNUSED
+// Test if it is OK to leave the variants unused in the header
+#else // UNUSED
+  int host = omp_is_initial_device();
+  int device = 1;
+#pragma omp target map(tofrom : device)
+  { device = omp_is_initial_device(); }
+  if (!host) {
+    printf("omp_is_initial_device() returned false on host\n");
+    errors++;
+  }
+  if (device) {
+    printf("omp_is_initial_device() returned true on device\n");
+    errors++;
+  }
+#endif // UNUSED
+
+  // CHECK: PASS
+  printf("%s\n", errors ? "FAIL" : "PASS");
+
+  return errors;
+}

openmp/runtime/src/include/omp.h.var

@@ -468,6 +468,15 @@
     /* OpenMP 5.1 Display Environment */
     extern void omp_display_env(int verbose);
 
+#   if defined(_OPENMP) && _OPENMP >= 201811
+    #pragma omp begin declare variant match(device={kind(host)})
+    static inline int omp_is_initial_device(void) { return 1; }
+    #pragma omp end declare variant
+    #pragma omp begin declare variant match(device={kind(nohost)})
+    static inline int omp_is_initial_device(void) { return 0; }
+    #pragma omp end declare variant
+#   endif
+
 #   undef __KAI_KMPC_CONVENTION
 #   undef __KMP_IMP