maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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

llvm-svn: 347714
This commit is contained in:
Gheorghe-Teodor Bercea 2018-11-27 21:23:40 +00:00
parent 229eee49fc
commit 31c1589ab0
2 changed files with 8 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
openmp/libomptarget/deviceRTLs/nvptx/src/interface.h
View File

@ -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,

8
openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu
View File

@ -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;