[OpenMP] Add missing `tt_hidden_helper_task_encountered` along with `tt_found_proxy_tasks`

[tianshilei1992]

Summary

In most cases, hidden helper task behave similar as detached tasks. That means,
for example, if we have to wait for detached tasks, we have to do the same thing
for hidden helper tasks as well. This patch adds the missing condition for hidden
helper task accordingly along with detached task.

[tianshilei1992]

rebase

[protze.joachim]

Is it necessary, to have tt_hidden_helper_task_encountered and tt_found_proxy_tasks? Would it be sufficient, to just use tt_found_proxy_tasks all the time?

[tianshilei1992]

In D107316#2924609, @protze.joachim wrote:

Is it necessary, to have tt_hidden_helper_task_encountered and tt_found_proxy_tasks? Would it be sufficient, to just use tt_found_proxy_tasks all the time?

Technically, I think it might be fine. However, I still would like to keep them separated as though proxy task and hidden helper task are similar, their execution are totally different. There is no proxy task generated during the whole execution, unlike detached task.

[tianshilei1992]

This patch can potentially fix a problem exposed in the added test case. However,
the case cannot pass now because it hangs in __kmp_execute_tasks_template.

[tianshilei1992]

rebase and fix hangs

[tianshilei1992]

openmp/runtime/src/kmp_tasking.cpp

3092

That is the reason that we can observe hangs in the past. Basically, D107496 changed the way to track incomplete children. Once we encountered a hidden helper task or detachable task, all following tasks will be tracked. In the simple case we added to this patch, we have a regular target region depending on a hidden helper task. The allocation of the regular target region increments the counter of incomplete children. Since it is a regular target region with dependences, clang lowers it to a function call to __kmpc_omp_wait_deps first followed by invoking the task directly. At this moment, current_task->td_incomplete_child_tasks is either 2 or 1, depending on whether the hidden helper task is finished, but it will never be 0. Therefore, we are trapped into the infinite loop here.
I simply change it to check if the td_incomplete_child_tasks is greater than 1 which can exclude itself. Since we don't know if the function call to __kmpc_omp_wait_deps is for taskwait or for a if0 task, I don't have any other good idea on the top of my mind to check here. We might have very tiny performance loss, caused by returning from this function and then entering it again, for the fairly rare case.

[AndreyChurbanov]

openmp/runtime/src/kmp_tasking.cpp

3092

I think the correct fix here would be addition of condition && !task_team->tt.tt_hidden_helper_task_encountered instead of changing the condition on current_task->td_incomplete_child_tasks. That should cause the thread to exit task execution loop only in the presence of a hidden helper task to check if it still needs to wait for this task or can proceed. In other cases the thread may need to execute tasks from its own queue without exiting the loop.

[tianshilei1992]

fix comments

[tianshilei1992]

openmp/runtime/src/kmp_tasking.cpp

3092

I fixed it in another way that is to add another check of the flag before set use_own_tasks. We do have two similar checks in different code path, but in the test case added to this patch, existing check cannot cover.

[AndreyChurbanov]

LGTM

openmp/libomptarget/test/offloading/target_nowait_target.cpp

+// RUN: %libomptarget-compilexx-and-run-generic
+
+// UNSUPPORTED: amdgcn-amd-amdhsa
+
+#include <cassert>
+
+int main(int argc, char *argv[]) {
+  int data[1024];
+  int sum = 0;
+
+  for (int i = 0; i < 1024; ++i)
+    data[i] = i;
+
+#pragma omp target map(tofrom: sum) map(to: data) depend(inout : data[0]) nowait
+  {
+    for (int i = 0; i < 1024; ++i) {
+      sum += data[i];
+    }
+  }
+
+#pragma omp target map(tofrom: sum) map(to: data) depend(inout : data[0])
+  {
+    for (int i = 0; i < 1024; ++i) {
+      sum += data[i];
+    }
+  }
+
+  assert(sum == 1023 * 1024);
+
+  return 0;
+}

openmp/runtime/src/kmp_barrier.cpp

 #if USE_ITT_NOTIFY
         void *itt_sync_obj = NULL;
         if (__itt_sync_create_ptr || KMP_ITT_DEBUG) {
           itt_sync_obj = __kmp_itt_barrier_object(gtid, bt, 1);
           __kmp_itt_barrier_starting(gtid, itt_sync_obj);
         }
 #endif
 
-        KMP_DEBUG_ASSERT(this_thr->th.th_task_team->tt.tt_found_proxy_tasks ==
-                         TRUE);
+        KMP_DEBUG_ASSERT(
+            this_thr->th.th_task_team->tt.tt_found_proxy_tasks == TRUE ||
+            this_thr->th.th_task_team->tt.tt_hidden_helper_task_encountered ==
         __kmp_task_team_wait(this_thr, team USE_ITT_BUILD_ARG(itt_sync_obj));
         __kmp_task_team_setup(this_thr, team, 0);
 
 #if USE_ITT_BUILD
         if (__itt_sync_create_ptr || KMP_ITT_DEBUG)
           __kmp_itt_barrier_finished(gtid, itt_sync_obj);
 #endif /* USE_ITT_BUILD */
       }

openmp/runtime/src/kmp_csupport.cpp

void __kmpc_end_serialized_parallel(ident_t *loc, kmp_int32 global_tid) {
 
   __kmp_resume_if_soft_paused();
 
   this_thr = __kmp_threads[global_tid];
   serial_team = this_thr->th.th_serial_team;
 
   kmp_task_team_t *task_team = this_thr->th.th_task_team;
   // we need to wait for the proxy tasks before finishing the thread
-  if (task_team != NULL && task_team->tt.tt_found_proxy_tasks)
+  if (task_team != NULL && (task_team->tt.tt_found_proxy_tasks ||
+                            task_team->tt.tt_hidden_helper_task_encountered))
     __kmp_task_team_wait(this_thr, serial_team USE_ITT_BUILD_ARG(NULL));
 
   KMP_MB();
   KMP_DEBUG_ASSERT(serial_team);
   KMP_ASSERT(serial_team->t.t_serialized);
   KMP_DEBUG_ASSERT(this_thr->th.th_team == serial_team);
   KMP_DEBUG_ASSERT(serial_team != this_thr->th.th_root->r.r_root_team);
   KMP_DEBUG_ASSERT(serial_team->t.t_threads);

openmp/runtime/src/kmp_runtime.cpp

void __kmp_unregister_root_current_thread(int gtid) {
 
   KMP_MB();
 
   kmp_info_t *thread = __kmp_threads[gtid];
   kmp_team_t *team = thread->th.th_team;
   kmp_task_team_t *task_team = thread->th.th_task_team;
 
   // we need to wait for the proxy tasks before finishing the thread
-  if (task_team != NULL && task_team->tt.tt_found_proxy_tasks) {
+  if (task_team != NULL && (task_team->tt.tt_found_proxy_tasks ||
+                            task_team->tt.tt_hidden_helper_task_encountered)) {
 #if OMPT_SUPPORT
     // the runtime is shutting down so we won't report any events
     thread->th.ompt_thread_info.state = ompt_state_undefined;
 #endif
     __kmp_task_team_wait(thread, team USE_ITT_BUILD_ARG(NULL));
   }
 
   __kmp_reset_root(gtid, root);

openmp/runtime/src/kmp_taskdeps.cpp

 #endif /* OMPT_SUPPORT */
 
   // We can return immediately as:
   // - dependences are not computed in serial teams (except with proxy tasks)
   // - if the dephash is not yet created it means we have nothing to wait for
   bool ignore = current_task->td_flags.team_serial ||
                 current_task->td_flags.tasking_ser ||
                 current_task->td_flags.final;
-  ignore = ignore && thread->th.th_task_team != NULL &&
-           thread->th.th_task_team->tt.tt_found_proxy_tasks == FALSE;
+  ignore =
+      ignore && thread->th.th_task_team != NULL &&
+      thread->th.th_task_team->tt.tt_found_proxy_tasks == FALSE &&
+      thread->th.th_task_team->tt.tt_hidden_helper_task_encountered == FALSE;
   ignore = ignore || current_task->td_dephash == NULL;
 
   if (ignore) {
     KA_TRACE(10, ("__kmpc_omp_wait_deps(exit): T#%d has no blocking "
                   "dependences : loc=%p\n",
                   gtid, loc_ref));
 #if OMPT_SUPPORT
     __ompt_taskwait_dep_finish(current_task, taskwait_task_data);

openmp/runtime/src/kmp_tasking.cpp

#endif
     // If this thread's task team is NULL, primary thread has recognized that
     // there are no more tasks; bail out
     if (thread->th.th_task_team == NULL) {
       KA_TRACE(15,
                ("__kmp_execute_tasks_template: T#%d no more tasks\n", gtid));
       return FALSE;
     }
 
+    // Check the flag again to see if it has already done in case to be trapped
+    // into infinite loop when a if0 task depends on a hidden helper task
+    // outside any parallel region. Detached tasks are not impacted in this case
+    // because the only thread executing this function has to execute the proxy
+    // task so it is in another code path that has the same check.
+    if (flag == NULL || (!final_spin && flag->done_check())) {
+      KA_TRACE(15,
+               ("__kmp_execute_tasks_template: T#%d spin condition satisfied\n",
+                gtid));
+      return TRUE;
+    }
+
     // We could be getting tasks from target constructs; if this is the only
     // thread, keep trying to execute tasks from own queue
     if (nthreads == 1 &&
         KMP_ATOMIC_LD_ACQ(&current_task->td_incomplete_child_tasks))
       use_own_tasks = 1;
     else {
       KA_TRACE(15,
                ("__kmp_execute_tasks_template: T#%d can't find work\n", gtid));
       return FALSE;
     }
   }
 }
#endif /* USE_ITT_BUILD && USE_ITT_NOTIFY && KMP_DEBUG */
     // Note: __kmp_allocate zeroes returned memory, othewise we would need:
     // task_team->tt.tt_threads_data = NULL;
     // task_team->tt.tt_max_threads = 0;
     // task_team->tt.tt_next = NULL;
   }
 
   TCW_4(task_team->tt.tt_found_tasks, FALSE);
   TCW_4(task_team->tt.tt_found_proxy_tasks, FALSE);
+  TCW_4(task_team->tt.tt_hidden_helper_task_encountered, FALSE);
   task_team->tt.tt_nproc = nthreads = team->t.t_nproc;
 
   KMP_ATOMIC_ST_REL(&task_team->tt.tt_unfinished_threads, nthreads);
   TCW_4(task_team->tt.tt_hidden_helper_task_encountered, FALSE);
   TCW_4(task_team->tt.tt_active, TRUE);
 
   KA_TRACE(20, ("__kmp_allocate_task_team: T#%d exiting; task_team = %p "
                 "unfinished_threads init'd to %d\n",
if (team->t.t_nproc > 1) {
     } else { // Leave the old task team struct in place for the upcoming region;
       // adjust as needed
       kmp_task_team_t *task_team = team->t.t_task_team[other_team];
       if (!task_team->tt.tt_active ||
           team->t.t_nproc != task_team->tt.tt_nproc) {
         TCW_4(task_team->tt.tt_nproc, team->t.t_nproc);
         TCW_4(task_team->tt.tt_found_tasks, FALSE);
         TCW_4(task_team->tt.tt_found_proxy_tasks, FALSE);
+        TCW_4(task_team->tt.tt_hidden_helper_task_encountered, FALSE);
         KMP_ATOMIC_ST_REL(&task_team->tt.tt_unfinished_threads,
                           team->t.t_nproc);
         TCW_4(task_team->tt.tt_active, TRUE);
       }
       // if team size has changed, the first thread to enable tasking will
       // realloc threads_data if necessary
       KA_TRACE(20, ("__kmp_task_team_setup: Primary T#%d reset next task_team "
                     "%p for team %d at parity=%d\n",
if ((task_team != NULL) && KMP_TASKING_ENABLED(task_team)) {
     // Deactivate the old task team, so that the worker threads will stop
     // referencing it while spinning.
     KA_TRACE(
         20,
         ("__kmp_task_team_wait: Primary T#%d deactivating task_team %p: "
          "setting active to false, setting local and team's pointer to NULL\n",
          __kmp_gtid_from_thread(this_thr), task_team));
     KMP_DEBUG_ASSERT(task_team->tt.tt_nproc > 1 ||
-                     task_team->tt.tt_found_proxy_tasks == TRUE);
+                     task_team->tt.tt_found_proxy_tasks == TRUE ||
+                     task_team->tt.tt_hidden_helper_task_encountered == TRUE);
     TCW_SYNC_4(task_team->tt.tt_found_proxy_tasks, FALSE);
+    TCW_SYNC_4(task_team->tt.tt_hidden_helper_task_encountered, FALSE);
     KMP_CHECK_UPDATE(task_team->tt.tt_untied_task_encountered, 0);
     TCW_SYNC_4(task_team->tt.tt_active, FALSE);
     KMP_MB();
 
     TCW_PTR(this_thr->th.th_task_team, NULL);
   }
 }