diff options
| author | Gheorghe-Teodor Bercea <gheorghe-teod.bercea@ibm.com> | 2018-11-27 21:23:40 +0000 |
|---|---|---|
| committer | Gheorghe-Teodor Bercea <gheorghe-teod.bercea@ibm.com> | 2018-11-27 21:23:40 +0000 |
| commit | 61fe54cce1708855185d18f99e41a8f3e09889db (patch) | |
| tree | 3981ce55439b92bbe761eb159a7096060b6f2881 /openmp/libomptarget | |
| parent | 73a25fd5c78941f7cb752bce7140d0c4bb7e54e6 (diff) | |
[OpenMP][libomptarget] Add new version of SPMD deinit kernel function with argument
Summary: To enable the compiler to optimize parts of the function that are not needed when runtime can be omitted, a new version of the SPMD deinit kernel function is needed. This function takes the runtime required flag as an argument.
Reviewers: ABataev, kkwli0, caomhin
Reviewed By: ABataev
Subscribers: guansong, openmp-commits
Differential Revision: https://reviews.llvm.org/D54969
Diffstat (limited to 'openmp/libomptarget')
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/src/interface.h | 3 | ||||
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu | 8 |
2 files changed, 8 insertions, 3 deletions
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/interface.h b/openmp/libomptarget/deviceRTLs/nvptx/src/interface.h index 40321a624cd..2c2beae3175 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/interface.h +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/interface.h @@ -494,7 +494,8 @@ EXTERN void __kmpc_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime); EXTERN void __kmpc_kernel_deinit(int16_t IsOMPRuntimeInitialized); EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime, int16_t RequiresDataSharing); -EXTERN void __kmpc_spmd_kernel_deinit(); +EXTERN __attribute__((deprecated)) void __kmpc_spmd_kernel_deinit(); +EXTERN void __kmpc_spmd_kernel_deinit_v2(int16_t RequiresOMPRuntime); EXTERN void __kmpc_kernel_prepare_parallel(void *WorkFn, int16_t IsOMPRuntimeInitialized); EXTERN bool __kmpc_kernel_parallel(void **WorkFn, diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu index 5d95eb1c77e..b0aef62ee9b 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu @@ -162,12 +162,16 @@ EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime, } } -EXTERN void __kmpc_spmd_kernel_deinit() { +EXTERN __attribute__((deprecated)) void __kmpc_spmd_kernel_deinit() { + __kmpc_spmd_kernel_deinit_v2(isRuntimeInitialized()); +} + +EXTERN void __kmpc_spmd_kernel_deinit_v2(int16_t RequiresOMPRuntime) { // We're not going to pop the task descr stack of each thread since // there are no more parallel regions in SPMD mode. __syncthreads(); int threadId = GetThreadIdInBlock(); - if (isRuntimeUninitialized()) { + if (!RequiresOMPRuntime) { if (threadId == 0) { // Enqueue omp state object for use by another team. int slot = usedSlotIdx; |