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)
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:
+++ /dev/null
-// 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]]
--- /dev/null
+// 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 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