[OpenMP][FIX] Enlarge thread state array, improve test and add second

Author: jdoerfert

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);

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;
 }

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;
+}