[OPENMP]Fix threadid in __kmpc_omp_taskwait call for dependent target calls.

[ABataev]

Summary

We used to call __kmpc_omp_taskwait function with global threadid set to
0. It may crash the application at the runtime if the thread executing
target region is not a master thread.

[grokos]

Looks good.

[protze.joachim]

The test fails for me with segfault in line 22.

$ gdb target_depend_nowait.cpp.tmp-x86_64-pc-linux-gnu
(gdb) run
...
Program received signal SIGSEGV, Segmentation fault.
(gdb) bt
#0  0x0000000000400fcf in .omp_outlined._debug__ (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:22
#1  0x0000000000401e8d in .omp_outlined..23 (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:18
#2  0x00002aaaab574213 in __kmp_invoke_microtask () from libomp.so
#3  0x00002aaaab51338e in __kmp_invoke_task_func () from libomp.so
#4  0x00002aaaab5126bf in __kmp_launch_thread () from libomp.so
#5  0x00002aaaab55d3d0 in __kmp_launch_worker(void*) () from libomp.so
#6  0x00002aaaabbd6e65 in start_thread () from libpthread.so.0
#7  0x00002aaaabee988d in clone () from libc.so.6

[ABataev]

In D64571#2160362, @protze.joachim wrote:

The test fails for me with segfault in line 22.

$ gdb target_depend_nowait.cpp.tmp-x86_64-pc-linux-gnu
(gdb) run
...
Program received signal SIGSEGV, Segmentation fault.
(gdb) bt
#0  0x0000000000400fcf in .omp_outlined._debug__ (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:22
#1  0x0000000000401e8d in .omp_outlined..23 (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:18
#2  0x00002aaaab574213 in __kmp_invoke_microtask () from libomp.so
#3  0x00002aaaab51338e in __kmp_invoke_task_func () from libomp.so
#4  0x00002aaaab5126bf in __kmp_launch_thread () from libomp.so
#5  0x00002aaaab55d3d0 in __kmp_launch_worker(void*) () from libomp.so
#6  0x00002aaaabbd6e65 in start_thread () from libpthread.so.0
#7  0x00002aaaabee988d in clone () from libc.so.6

It is not the problem of this patch, looks like support for mappers causes failures

[saiislam]

In D64571#2162388, @ABataev wrote:

In D64571#2160362, @protze.joachim wrote:

The test fails for me with segfault in line 22.

$ gdb target_depend_nowait.cpp.tmp-x86_64-pc-linux-gnu
(gdb) run
...
Program received signal SIGSEGV, Segmentation fault.
(gdb) bt
#0  0x0000000000400fcf in .omp_outlined._debug__ (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:22
#1  0x0000000000401e8d in .omp_outlined..23 (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:18
#2  0x00002aaaab574213 in __kmp_invoke_microtask () from libomp.so
#3  0x00002aaaab51338e in __kmp_invoke_task_func () from libomp.so
#4  0x00002aaaab5126bf in __kmp_launch_thread () from libomp.so
#5  0x00002aaaab55d3d0 in __kmp_launch_worker(void*) () from libomp.so
#6  0x00002aaaabbd6e65 in start_thread () from libpthread.so.0
#7  0x00002aaaabee988d in clone () from libc.so.6

It is not the problem of this patch, looks like support for mappers causes failures

Yes, I can confirm that this test is failing after D68100 (declare mapper runtime implementation) for me as well.

[ABataev]

In D64571#2162397, @saiislam wrote:

In D64571#2162388, @ABataev wrote:

In D64571#2160362, @protze.joachim wrote:

The test fails for me with segfault in line 22.

$ gdb target_depend_nowait.cpp.tmp-x86_64-pc-linux-gnu
(gdb) run
...
Program received signal SIGSEGV, Segmentation fault.
(gdb) bt
#0  0x0000000000400fcf in .omp_outlined._debug__ (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:22
#1  0x0000000000401e8d in .omp_outlined..23 (.global_tid.=0x2aaab96b9be0, .bound_tid.=0x2aaab96b9bd8) at target_depend_nowait.cpp:18
#2  0x00002aaaab574213 in __kmp_invoke_microtask () from libomp.so
#3  0x00002aaaab51338e in __kmp_invoke_task_func () from libomp.so
#4  0x00002aaaab5126bf in __kmp_launch_thread () from libomp.so
#5  0x00002aaaab55d3d0 in __kmp_launch_worker(void*) () from libomp.so
#6  0x00002aaaabbd6e65 in start_thread () from libpthread.so.0
#7  0x00002aaaabee988d in clone () from libc.so.6

It is not the problem of this patch, looks like support for mappers causes failures

Yes, I can confirm that this test is failing after D68100 (declare mapper runtime implementation) for me as well.

Georgios, could you check this?

[grokos]

I'm having a look. The problem actually occurs after D67833 (the clang patch), not D68100. If you see at the backtrace in gdb, control never reaches a libomptarget entry point, things fail once we enter the outlined function for the parallel region via __kmp_invoke_microtask. I doubt it's a libomptarget issue.

[protze.joachim]

Sorry, I was looking at the date and thought the commit was 2 days old instead of 1 year :)
You are right, the test started to fail after the D67833 commit.

libomptarget/src/interface.cpp

#endif
   HandleTargetOutcome(rc == OFFLOAD_SUCCESS);
 }
 
 EXTERN void __tgt_target_data_begin_nowait(int64_t device_id, int32_t arg_num,
     void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types,
     int32_t depNum, void *depList, int32_t noAliasDepNum,
     void *noAliasDepList) {
   if (depNum + noAliasDepNum > 0)
-    __kmpc_omp_taskwait(NULL, 0);
+    __kmpc_omp_taskwait(NULL, __kmpc_global_thread_num(NULL));
 
   __tgt_target_data_begin(device_id, arg_num, args_base, args, arg_sizes,
                           arg_types);
 }
 
 /// passes data from the target, releases target memory and destroys
 /// the host-target mapping (top entry from the stack of data maps)
 /// created by the last __tgt_target_data_begin.
#endif
   HandleTargetOutcome(rc == OFFLOAD_SUCCESS);
 }
 
 EXTERN void __tgt_target_data_end_nowait(int64_t device_id, int32_t arg_num,
     void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types,
     int32_t depNum, void *depList, int32_t noAliasDepNum,
     void *noAliasDepList) {
   if (depNum + noAliasDepNum > 0)
-    __kmpc_omp_taskwait(NULL, 0);
+    __kmpc_omp_taskwait(NULL, __kmpc_global_thread_num(NULL));
 
   __tgt_target_data_end(device_id, arg_num, args_base, args, arg_sizes,
                         arg_types);
 }
 
 EXTERN void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
     void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types) {
   if (IsOffloadDisabled()) return;
EXTERN void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
   HandleTargetOutcome(rc == OFFLOAD_SUCCESS);
 }
 
 EXTERN void __tgt_target_data_update_nowait(
     int64_t device_id, int32_t arg_num, void **args_base, void **args,
     int64_t *arg_sizes, int64_t *arg_types, int32_t depNum, void *depList,
     int32_t noAliasDepNum, void *noAliasDepList) {
   if (depNum + noAliasDepNum > 0)
-    __kmpc_omp_taskwait(NULL, 0);
+    __kmpc_omp_taskwait(NULL, __kmpc_global_thread_num(NULL));
 
   __tgt_target_data_update(device_id, arg_num, args_base, args, arg_sizes,
                            arg_types);
 }
 
 EXTERN int __tgt_target(int64_t device_id, void *host_ptr, int32_t arg_num,
     void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types) {
   if (IsOffloadDisabled()) return OFFLOAD_FAIL;
#endif
   return rc;
 }
 
 EXTERN int __tgt_target_nowait(int64_t device_id, void *host_ptr,
     int32_t arg_num, void **args_base, void **args, int64_t *arg_sizes,
     int64_t *arg_types, int32_t depNum, void *depList, int32_t noAliasDepNum,
     void *noAliasDepList) {
   if (depNum + noAliasDepNum > 0)
-    __kmpc_omp_taskwait(NULL, 0);
+    __kmpc_omp_taskwait(NULL, __kmpc_global_thread_num(NULL));
 
   return __tgt_target(device_id, host_ptr, arg_num, args_base, args, arg_sizes,
                       arg_types);
 }
 
 EXTERN int __tgt_target_teams(int64_t device_id, void *host_ptr,
     int32_t arg_num, void **args_base, void **args, int64_t *arg_sizes,
     int64_t *arg_types, int32_t team_num, int32_t thread_limit) {
#endif
   return rc;
 }
 
 EXTERN int __tgt_target_teams_nowait(int64_t device_id, void *host_ptr,
     int32_t arg_num, void **args_base, void **args, int64_t *arg_sizes,
     int64_t *arg_types, int32_t team_num, int32_t thread_limit, int32_t depNum,
     void *depList, int32_t noAliasDepNum, void *noAliasDepList) {
   if (depNum + noAliasDepNum > 0)
-    __kmpc_omp_taskwait(NULL, 0);
+    __kmpc_omp_taskwait(NULL, __kmpc_global_thread_num(NULL));
 
   return __tgt_target_teams(device_id, host_ptr, arg_num, args_base, args,
                             arg_sizes, arg_types, team_num, thread_limit);
 }
 
 EXTERN void __kmpc_push_target_tripcount(int64_t device_id,
     uint64_t loop_tripcount) {
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {

libomptarget/test/offloading/target_depend_nowait.cpp

+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+
+#include <omp.h>
+#include <stdio.h>
+
+#define N 1024
+
+int A[N];
+int B[N];
+int C[N];
+int main() {
+  for (int i = 0; i < N; i++)
+    A[i] = B[i] = i;
+
+#pragma omp parallel num_threads(2)
+  {
+    if (omp_get_thread_num() == 1) {
+// map data A & B and move to
+#pragma omp target enter data map(to : A, B) depend(out : A[0]) nowait
+
+// no data move since already mapped
+#pragma omp target map(A, B) depend(out : A[0]) nowait
+      {
+        for (int i = 0; i < N; i++)
+          ++A[i];
+        for (int i = 0; i < N; i++)
+          ++B[i];
+      }
+
+// no data move since already mapped
+#pragma omp target teams num_teams(1) map(A, B) depend(out : A[0]) nowait
+      {
+        for (int i = 0; i < N; i++)
+          ++A[i];
+        for (int i = 0; i < N; i++)
+          ++B[i];
+      }
+
+// A updated via update
+#pragma omp target update from(A) depend(out : A[0]) nowait
+
+// B updated via exit, A just released
+#pragma omp target exit data map(release                                       \
+                                 : A) map(from                                 \
+                                          : B) depend(out                      \
+                                                      : A[0]) nowait
+    } // if
+  }   // parallel
+
+  int Sum = 0;
+  for (int i = 0; i < N; i++)
+    Sum += A[i] + B[i];
+  // Sum is 2 * N * (2 + N - 1 + 2) / 2
+  // CHECK: Sum = 1051648.
+  printf("Sum = %d.\n", Sum);
+
+  return Sum != 2 * N * (2 + N - 1 + 2) / 2;
+}
+