[Openmp-commits] [openmp] 3da61dd - [OpenMP] Define omp_is_initial_device() variants in omp.h
Hansang Bae via Openmp-commits
openmp-commits at lists.llvm.org
Tue Apr 6 14:58:27 PDT 2021
Author: Hansang Bae
Date: 2021-04-06T16:58:01-05:00
New Revision: 3da61ddae7fe77f71b89ce20cf6b5febd68d216a
URL: https://github.com/llvm/llvm-project/commit/3da61ddae7fe77f71b89ce20cf6b5febd68d216a
DIFF: https://github.com/llvm/llvm-project/commit/3da61ddae7fe77f71b89ce20cf6b5febd68d216a.diff
LOG: [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
Added:
openmp/libomptarget/test/api/is_initial_device.c
Modified:
clang/include/clang/Basic/Builtins.def
clang/lib/AST/ExprConstant.cpp
openmp/runtime/src/include/omp.h.var
Removed:
clang/test/OpenMP/is_initial_device.c
################################################################################
diff --git a/clang/include/clang/Basic/Builtins.def b/clang/include/clang/Basic/Builtins.def
index 153e22f00f522..8518f3789f51a 100644
--- a/clang/include/clang/Basic/Builtins.def
+++ b/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)
diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp
index b42f3b695ec57..fe6573b8bd2c1 100644
--- a/clang/lib/AST/ExprConstant.cpp
+++ b/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:
diff --git a/clang/test/OpenMP/is_initial_device.c b/clang/test/OpenMP/is_initial_device.c
deleted file mode 100644
index 2fe93a4eb6c46..0000000000000
--- a/clang/test/OpenMP/is_initial_device.c
+++ /dev/null
@@ -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]]
diff --git a/openmp/libomptarget/test/api/is_initial_device.c b/openmp/libomptarget/test/api/is_initial_device.c
new file mode 100644
index 0000000000000..78980d6316c6d
--- /dev/null
+++ b/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;
+}
diff --git a/openmp/runtime/src/include/omp.h.var b/openmp/runtime/src/include/omp.h.var
index 28e9259482631..c269fa6359cdd 100644
--- a/openmp/runtime/src/include/omp.h.var
+++ b/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
More information about the Openmp-commits
mailing list