From 458db51c101bc3372e96b71bda7ca0f5ba2ae431 Mon Sep 17 00:00:00 2001 From: Shilei Tian Date: Wed, 29 Dec 2021 23:22:37 -0500 Subject: [PATCH] [OpenMP] Add missing `tt_hidden_helper_task_encountered` along with `tt_found_proxy_tasks` 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. Reviewed By: AndreyChurbanov Differential Revision: https://reviews.llvm.org/D107316 --- .../test/offloading/target_nowait_target.cpp | 31 +++++++++++++++++++ openmp/runtime/src/kmp_barrier.cpp | 6 ++-- openmp/runtime/src/kmp_csupport.cpp | 3 +- openmp/runtime/src/kmp_runtime.cpp | 3 +- openmp/runtime/src/kmp_taskdeps.cpp | 6 ++-- openmp/runtime/src/kmp_tasking.cpp | 18 ++++++++++- 6 files changed, 60 insertions(+), 7 deletions(-) create mode 100644 openmp/libomptarget/test/offloading/target_nowait_target.cpp diff --git a/openmp/libomptarget/test/offloading/target_nowait_target.cpp b/openmp/libomptarget/test/offloading/target_nowait_target.cpp new file mode 100644 index 000000000000..24a83c300524 --- /dev/null +++ b/openmp/libomptarget/test/offloading/target_nowait_target.cpp @@ -0,0 +1,31 @@ +// RUN: %libomptarget-compilexx-and-run-generic + +// UNSUPPORTED: amdgcn-amd-amdhsa + +#include + +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; +} diff --git a/openmp/runtime/src/kmp_barrier.cpp b/openmp/runtime/src/kmp_barrier.cpp index 97bf9811bcd0..ee05bb3587ca 100644 --- a/openmp/runtime/src/kmp_barrier.cpp +++ b/openmp/runtime/src/kmp_barrier.cpp @@ -2037,8 +2037,10 @@ static int __kmp_barrier_template(enum barrier_type bt, int gtid, int is_split, } #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 == + TRUE); __kmp_task_team_wait(this_thr, team USE_ITT_BUILD_ARG(itt_sync_obj)); __kmp_task_team_setup(this_thr, team, 0); diff --git a/openmp/runtime/src/kmp_csupport.cpp b/openmp/runtime/src/kmp_csupport.cpp index e95c2f072509..e263558517d0 100644 --- a/openmp/runtime/src/kmp_csupport.cpp +++ b/openmp/runtime/src/kmp_csupport.cpp @@ -531,7 +531,8 @@ void __kmpc_end_serialized_parallel(ident_t *loc, kmp_int32 global_tid) { 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(); diff --git a/openmp/runtime/src/kmp_runtime.cpp b/openmp/runtime/src/kmp_runtime.cpp index 6efc26df8de3..7af970803a30 100644 --- a/openmp/runtime/src/kmp_runtime.cpp +++ b/openmp/runtime/src/kmp_runtime.cpp @@ -4106,7 +4106,8 @@ void __kmp_unregister_root_current_thread(int gtid) { 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; diff --git a/openmp/runtime/src/kmp_taskdeps.cpp b/openmp/runtime/src/kmp_taskdeps.cpp index 7d2774a738fb..501830eaa758 100644 --- a/openmp/runtime/src/kmp_taskdeps.cpp +++ b/openmp/runtime/src/kmp_taskdeps.cpp @@ -829,8 +829,10 @@ void __kmpc_omp_wait_deps(ident_t *loc_ref, kmp_int32 gtid, kmp_int32 ndeps, 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) { diff --git a/openmp/runtime/src/kmp_tasking.cpp b/openmp/runtime/src/kmp_tasking.cpp index d956df1b2a37..d6665a7ccfb4 100644 --- a/openmp/runtime/src/kmp_tasking.cpp +++ b/openmp/runtime/src/kmp_tasking.cpp @@ -3074,6 +3074,18 @@ static inline int __kmp_execute_tasks_template( 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 && @@ -3478,6 +3490,7 @@ static kmp_task_team_t *__kmp_allocate_task_team(kmp_info_t *thread, 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); @@ -3640,6 +3653,7 @@ void __kmp_task_team_setup(kmp_info_t *this_thr, kmp_team_t *team, int always) { 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); @@ -3732,8 +3746,10 @@ void __kmp_task_team_wait( "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();