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

[hbae]

Summary

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. Full fix also requires removal of the routine from the
built-in function list in the compiler.

[jdoerfert]

Cool, can we get rid of the clang builtin while we are at it?

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

479

Do we want them to be static? I would have expected odr linkage instead.

[hbae]

Removed omp_is_initial_device() from clang built-ins.
Added a new test under openmp/libomptarget.

[hbae]

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

479

Current clang fails to link two files that include omp.h, so I think we need static here.

[jdoerfert]

LG, thanks

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

479

I would expect this to cause "unused function" warnings if you include omp.h.
That said, given that clang might be the only impl. so far for begin/end declare variant and it doesn't warn, we can go ahead.

[JonChesterfield]

Nice, thanks! I think the function in the devicertl is dead with this change, maybe remove that too?

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

479

'inline' is right for c++. Inline plus instantiation in library code right for c. Can we hit linkonce_odr from a clang attribute?

[hbae]

Added inline after static to avoid unused function warning.
Added a RUN line to the test to catch unused function warning.

[JonChesterfield]

LGTM, thanks!

clang/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)
-
 // CUDA/HIP
 LANGBUILTIN(__builtin_get_device_side_mangled_name, "cC*.", "ncT", CUDA_LANG)
 
 // Builtins for XRay
 BUILTIN(__xray_customevent, "vcC*z", "")
 BUILTIN(__xray_typedevent, "vzcC*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

clang/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);
   case Builtin::BI__builtin_add_overflow:
   case Builtin::BI__builtin_sub_overflow:
   case Builtin::BI__builtin_mul_overflow:
   case Builtin::BI__builtin_sadd_overflow:
   case Builtin::BI__builtin_uadd_overflow:
   case Builtin::BI__builtin_uaddl_overflow:
   case Builtin::BI__builtin_uaddll_overflow:
   case Builtin::BI__builtin_usub_overflow:

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

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

#   endif
     extern void __KAI_KMPC_CONVENTION omp_set_num_teams(int num_teams);
     extern int __KAI_KMPC_CONVENTION omp_get_max_teams(void);
     extern void __KAI_KMPC_CONVENTION omp_set_teams_thread_limit(int limit);
     extern int __KAI_KMPC_CONVENTION omp_get_teams_thread_limit(void);
 
     /* 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
 
     /* Warning:
        The following typedefs are not standard, deprecated and will be removed in a future release.
     */
     typedef int     omp_int_t;
     typedef double  omp_wtime_t;
 
 #   ifdef __cplusplus
     }
 #   endif
 
 #endif /* __OMP_H */