[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
GitOrigin-RevId: 3da61ddae7fe77f71b89ce20cf6b5febd68d216a
diff --git a/libomptarget/test/api/is_initial_device.c b/libomptarget/test/api/is_initial_device.c
new file mode 100644
index 0000000..78980d6
--- /dev/null
+++ b/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;
+}
diff --git a/runtime/src/include/omp.h.var b/runtime/src/include/omp.h.var
index 28e9259..c269fa6 100644
--- a/runtime/src/include/omp.h.var
+++ b/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
