[OPENMMP] Resolve lost LoopTripCnt for subsequent loops in same thread.

[ronlieb]

Summary

Remove loopTripCnt from threaded device stack after consuming it.
Added a libomptarget DP message to aid in future debugging and to
validate the added testcase, which only runs in Debug build.

%ld => %lu
removed extraneous braces.

[ABataev]

libomptarget/src/omptarget.cpp

738 ↗

(On Diff #210130)

Use %lu instead, it is an unsigned type.

libomptarget/test/offloading/looptripcnt.c

38 ↗

(On Diff #210130)

I think, instead of the debug messages better to check the results of the printfs.

[Hahnfeld]

Can we validate the result differently, without debug messages?

libomptarget/test/offloading/looptripcnt.c

21 ↗

(On Diff #210130)

I think the standard doesn't expect a structured block, but a for loop directly.

[ronlieb]

libomptarget/test/offloading/looptripcnt.c

38 ↗

(On Diff #210130)

I initially thought about that, but I am worried that depending on what architecture you run on it might produce different numbers of teams, if the thread limit was a lot smaller than on a gpu for example.

but what I am really trying to test is that we found the loopTripCount to pass into the runtime.

if you strongly prefer that I still do so, I will be happy to change the test. I would of course still keep the DP message for diagnostic purposes down the road.

[ABataev]

libomptarget/test/offloading/looptripcnt.c

38 ↗

(On Diff #210130)

It is always better to check the real results rather than debug messages. If you could write a test, which is able to provide a stable check, then do this.

[ronlieb]

libomptarget/test/offloading/looptripcnt.c

21 ↗

(On Diff #210130)

will fix shortly.

38 ↗

(On Diff #210130)

ok, so I spent some time trying to do this. The issue I am encountering is
that the behavior that needs to be verified is that we properly find the loop trip count in the function 'target' within omptarget.cpp. This is then passed along to function __tgt_rtl_run_target_team_region in plugins.cpp which "may" choose to modify the blocksPerGrid or not. So depending on your target, you may see different numbers of teams active if you query omp_get_num_teams(). The observed behavior on different hosts is tenuous at best in determining if loop_trip_count was properly found. The point of this test is to make sure we dont loose the loop_trip_count. In this situation i think it best to catch the problem in Debug build/test where we can explicitly observe the DP message.

[ABataev]

libomptarget/test/offloading/looptripcnt.c

38 ↗

(On Diff #210130)

Then better to check for the debug message, I think. Other opinions?

[ronlieb]

@grokos any comments here please?
hoping to get this patch in before the 9.0 branch.

[grokos]

In D64808#1589239, @ronlieb wrote:

@grokos any comments here please?
hoping to get this patch in before the 9.0 branch.

The patch should definitely be committed as it fixes a potential problem with multiple threads. Regarding the test, I'm a bit torn. I don't like the idea of relying on debug messages either, but you are right that the plugin may change the number of teams so we cannot predict consistently what omp_get_num_teams() will return on each architecture. I will accept the patch as-is so that the bugfix in merged into the 9.0 release but we will need to revise the test in the future.

[ronlieb]

Thanks folks, publishing now

openmp/trunk/libomptarget/src/omptarget.cpp

#endif
 
   assert(tgt_args.size() == tgt_offsets.size() &&
       "Size mismatch in arguments and offsets");
 
   // Pop loop trip count
   uint64_t ltc = 0;
   TblMapMtx.lock();
   auto I = Device.LoopTripCnt.find(__kmpc_global_thread_num(NULL));
-  if (I != Device.LoopTripCnt.end())
-    std::swap(ltc, I->second);
+  if (I != Device.LoopTripCnt.end()) {
+    ltc = I->second;
+    Device.LoopTripCnt.erase(I);
+    DP("loop trip count is %lu.\n", ltc);
+  }
   TblMapMtx.unlock();
 
   // Launch device execution.
   DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n",
       TargetTable->EntriesBegin[TM->Index].name,
       DPxPTR(TargetTable->EntriesBegin[TM->Index].addr), TM->Index);
   if (IsTeamConstruct) {
     rc = Device.run_team_region(TargetTable->EntriesBegin[TM->Index].addr,

openmp/trunk/libomptarget/test/offloading/looptripcnt.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=DEBUG
+// REQUIRES: libomptarget-debug
+
+/*
+  Test for looptripcount being popped from runtime stack.
+*/
+#include <stdio.h>
+#include <omp.h>
+int main()
+{
+  int N = 128;
+  int NN = 1024;
+  int num_teams[NN];
+  int num_threads[NN];
+
+  printf("#pragma omp target teams distribute parallel for thread_limit(4)\n");
+#pragma omp target teams distribute parallel for thread_limit(4)
+  for (int j = 0; j< N; j++) {
+    num_threads[j] = omp_get_num_threads();
+    num_teams[j] = omp_get_num_teams();
+  }
+  printf("num_threads %d num_teams %d\n", num_threads[0], num_teams[0]);
+// DEBUG: loop trip count is 128
+  printf("#pragma omp target teams distribute parallel for\n");
+#pragma omp target teams distribute parallel for
+  for (int j = 0; j< N; j++) {
+    num_threads[j] = omp_get_num_threads();
+    num_teams[j] = omp_get_num_teams();
+  }
+  printf("num_threads %d num_teams %d\n", num_threads[0], num_teams[0]);
+// DEBUG: loop trip count is 128
+  return 0;
+}