diff --git a/openmp/libomptarget/test/offloading/taskloop_offload_nowait.cpp b/openmp/libomptarget/test/offloading/taskloop_offload_nowait.cpp new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/offloading/taskloop_offload_nowait.cpp @@ -0,0 +1,40 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +#include +#include +#include + +bool almost_equal(float x, float gold, float tol) { + if (std::signbit(x) != std::signbit(gold)) + x = std::abs(gold) - std::abs(x); + + return std::abs(gold) * (1 - tol) <= std::abs(x) && + std::abs(x) <= std::abs(gold) * (1 + tol); +} + +int main(int argc, char *argv[]) { + constexpr const int N0{2}; + constexpr const int N1{182}; + constexpr const float expected_value{N0 * N1}; + float counter_N0{}; + +#pragma omp target data map(tofrom : counter_N0) + { +#pragma omp taskloop shared(counter_N0) + for (int i0 = 0; i0 < N0; i0++) { +#pragma omp target teams distribute parallel for map(tofrom : counter_N0) nowait + for (int i1 = 0; i1 < N1; i1++) { +#pragma omp atomic update + counter_N0 = counter_N0 + 1.; + } + } + } + + if (!almost_equal(counter_N0, expected_value, 0.1)) { + std::cerr << "Expected: " << expected_value << " Got: " << counter_N0 + << '\n'; + return -1; + } + + return 0; +} diff --git a/openmp/runtime/src/kmp_tasking.cpp b/openmp/runtime/src/kmp_tasking.cpp --- a/openmp/runtime/src/kmp_tasking.cpp +++ b/openmp/runtime/src/kmp_tasking.cpp @@ -2563,7 +2563,8 @@ if (!taskdata->td_flags.team_serial || (thread->th.th_task_team != NULL && - thread->th.th_task_team->tt.tt_found_proxy_tasks)) { + (thread->th.th_task_team->tt.tt_found_proxy_tasks || + thread->th.th_task_team->tt.tt_hidden_helper_task_encountered))) { kmp_flag_32 flag( RCAST(std::atomic *, &(taskgroup->count)), 0U); while (KMP_ATOMIC_LD_ACQ(&taskgroup->count) != 0) {