[Openmp-commits] [openmp] 86bb713 - [OpenMP][FIX] Enlarge thread state array, improve test and add second
Johannes Doerfert via Openmp-commits
openmp-commits at lists.llvm.org
Sun Oct 22 17:47:12 PDT 2023
Author: Johannes Doerfert
Date: 2023-10-22T17:47:00-07:00
New Revision: 86bb713142c3cb2a895113061de6ced1f15ebb84
URL: https://github.com/llvm/llvm-project/commit/86bb713142c3cb2a895113061de6ced1f15ebb84
DIFF: https://github.com/llvm/llvm-project/commit/86bb713142c3cb2a895113061de6ced1f15ebb84.diff
LOG: [OpenMP][FIX] Enlarge thread state array, improve test and add second
Added:
openmp/libomptarget/test/offloading/thread_state_2.c
Modified:
openmp/libomptarget/DeviceRTL/src/State.cpp
openmp/libomptarget/test/offloading/thread_state_1.c
Removed:
################################################################################
diff --git a/openmp/libomptarget/DeviceRTL/src/State.cpp b/openmp/libomptarget/DeviceRTL/src/State.cpp
index 22f684b05f598d8..e929edda9502b60 100644
--- a/openmp/libomptarget/DeviceRTL/src/State.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/State.cpp
@@ -262,7 +262,8 @@ void state::enterDataEnvironment(IdentTy *Ident) {
memory::allocGlobal(sizeof(ThreadStateTy), "ThreadStates alloc"));
uintptr_t *ThreadStatesBitsPtr = reinterpret_cast<uintptr_t *>(&ThreadStates);
if (!atomic::load(ThreadStatesBitsPtr, atomic::seq_cst)) {
- uint32_t Bytes = sizeof(ThreadStates[0]) * mapping::getMaxTeamThreads();
+ uint32_t Bytes =
+ sizeof(ThreadStates[0]) * mapping::getNumberOfThreadsInBlock();
void *ThreadStatesPtr =
memory::allocGlobal(Bytes, "Thread state array allocation");
memset(ThreadStatesPtr, 0, Bytes);
diff --git a/openmp/libomptarget/test/offloading/thread_state_1.c b/openmp/libomptarget/test/offloading/thread_state_1.c
index 4e8043ea844ebf2..f3f7b32eead5645 100644
--- a/openmp/libomptarget/test/offloading/thread_state_1.c
+++ b/openmp/libomptarget/test/offloading/thread_state_1.c
@@ -1,5 +1,4 @@
-// Still broken "without optimizations"
-// XUN: %libomptarget-compile-run-and-check-generic
+// RUN: %libomptarget-compile-run-and-check-generic
// RUN: %libomptarget-compileopt-run-and-check-generic
#include <omp.h>
@@ -10,10 +9,10 @@ int main() {
int o_lvl = 111, i_lvl = 222, o_tid = 333, i_tid = 333, o_nt = 444,
i_nt = 555;
#pragma omp target teams map(tofrom : o_lvl, i_lvl, o_tid, i_tid, o_nt, i_nt) \
- num_teams(2) thread_limit(2)
+ num_teams(2) thread_limit(64)
{
if (omp_get_team_num() == 0) {
-#pragma omp parallel num_threads(128)
+#pragma omp parallel num_threads(64)
if (omp_get_thread_num() == omp_get_num_threads() - 1) {
o_lvl = omp_get_level();
o_tid = omp_get_thread_num();
@@ -27,9 +26,13 @@ int main() {
}
}
}
- // CHECK: outer: lvl: 1, tid: 1, nt: 2
- // CHECK: inner: lvl: 2, tid: 0, nt: 1
+ if (o_lvl == 1 && o_tid == o_nt - 1 && o_nt > 1 && i_lvl == 2 && i_tid == 0 &&
+ i_nt == 1) {
+ // CHECK: Success
+ printf("Success\n");
+ return 0;
+ }
printf("outer: lvl: %i, tid: %i, nt: %i\n", o_lvl, o_tid, o_nt);
printf("inner: lvl: %i, tid: %i, nt: %i\n", i_lvl, i_tid, i_nt);
- return 0;
+ return 1;
}
diff --git a/openmp/libomptarget/test/offloading/thread_state_2.c b/openmp/libomptarget/test/offloading/thread_state_2.c
new file mode 100644
index 000000000000000..6d3bf1661f46228
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/thread_state_2.c
@@ -0,0 +1,40 @@
+// This fails when optimized for now.
+// RUN: %libomptarget-compile-run-and-check-generic
+// XUN: %libomptarget-compileopt-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ // TODO: Test all ICVs on all levels
+ int o_lvl = 111, i_lvl = 222, o_tid = 333, i_tid = 333, o_nt = 444,
+ i_nt = 555;
+#pragma omp target teams map(tofrom : o_lvl, i_lvl, o_tid, i_tid, o_nt, i_nt) \
+ num_teams(2) thread_limit(64)
+ {
+ omp_set_max_active_levels(1);
+ if (omp_get_team_num() == 0) {
+#pragma omp parallel num_threads(64)
+ if (omp_get_thread_num() == omp_get_num_threads() - 1) {
+ o_lvl = omp_get_level();
+ o_tid = omp_get_thread_num();
+ o_nt = omp_get_num_threads();
+#pragma omp parallel num_threads(64)
+ if (omp_get_thread_num() == omp_get_num_threads() - 1) {
+ i_lvl = omp_get_level();
+ i_tid = omp_get_thread_num();
+ i_nt = omp_get_num_threads();
+ }
+ }
+ }
+ }
+ if (o_lvl == 1 && o_tid == o_nt - 1 && o_nt > 1 && i_lvl == 2 && i_tid == 0 &&
+ i_nt == 1) {
+ // CHECK: Success
+ printf("Success\n");
+ return 0;
+ }
+ printf("outer: lvl: %i, tid: %i, nt: %i\n", o_lvl, o_tid, o_nt);
+ printf("inner: lvl: %i, tid: %i, nt: %i\n", i_lvl, i_tid, i_nt);
+ return 1;
+}
More information about the Openmp-commits
mailing list