diff options
author | Alexey Bataev <a.bataev@hotmail.com> | 2019-07-16 15:51:32 +0000 |
---|---|---|
committer | Alexey Bataev <a.bataev@hotmail.com> | 2019-07-16 15:51:32 +0000 |
commit | 946967528b804d3de059b9ad6fbe7e22ebbdd9bb (patch) | |
tree | 23a2141e26b3b950b8d72c87320b39fb49992db1 /libomptarget | |
parent | d4b114904a3bb0ee08a7833c52450b793d6c7781 (diff) |
[OPENMP]Fix threadid in __kmpc_omp_taskwait call for dependent target calls.
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.
Reviewers: grokos, kkwli0
Subscribers: guansong, jdoerfert, caomhin, openmp-commits
Tags: #openmp
Differential Revision: https://reviews.llvm.org/D64571
git-svn-id: https://llvm.org/svn/llvm-project/openmp/trunk@366220 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'libomptarget')
-rw-r--r-- | libomptarget/src/interface.cpp | 10 | ||||
-rw-r--r-- | libomptarget/test/offloading/target_depend_nowait.cpp | 62 |
2 files changed, 67 insertions, 5 deletions
diff --git a/libomptarget/src/interface.cpp b/libomptarget/src/interface.cpp index 52850ee..32afe3f 100644 --- a/libomptarget/src/interface.cpp +++ b/libomptarget/src/interface.cpp @@ -128,7 +128,7 @@ EXTERN void __tgt_target_data_begin_nowait(int64_t device_id, int32_t arg_num, 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); @@ -181,7 +181,7 @@ EXTERN void __tgt_target_data_end_nowait(int64_t device_id, int32_t arg_num, 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); @@ -214,7 +214,7 @@ EXTERN void __tgt_target_data_update_nowait( 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); @@ -255,7 +255,7 @@ EXTERN int __tgt_target_nowait(int64_t device_id, void *host_ptr, 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); @@ -298,7 +298,7 @@ EXTERN int __tgt_target_teams_nowait(int64_t device_id, void *host_ptr, 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); diff --git a/libomptarget/test/offloading/target_depend_nowait.cpp b/libomptarget/test/offloading/target_depend_nowait.cpp new file mode 100644 index 0000000..2c1c7e7 --- /dev/null +++ b/libomptarget/test/offloading/target_depend_nowait.cpp @@ -0,0 +1,62 @@ +// 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; +} + |