Index: libomptarget/deviceRTLs/nvptx/src/interface.h =================================================================== --- libomptarget/deviceRTLs/nvptx/src/interface.h +++ libomptarget/deviceRTLs/nvptx/src/interface.h @@ -488,7 +488,8 @@ 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, Index: libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu =================================================================== --- libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu +++ libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu @@ -163,7 +163,7 @@ } } -EXTERN void __kmpc_spmd_kernel_deinit() { +EXTERN __attribute__((deprecated)) void __kmpc_spmd_kernel_deinit() { // We're not going to pop the task descr stack of each thread since // there are no more parallel regions in SPMD mode. __syncthreads(); @@ -185,6 +185,28 @@ } } +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 (!RequiresOMPRuntime) { + if (threadId == 0) { + // Enqueue omp state object for use by another team. + int slot = usedSlotIdx; + omptarget_nvptx_device_simpleState[slot].Enqueue( + omptarget_nvptx_simpleThreadPrivateContext); + } + return; + } + if (threadId == 0) { + // Enqueue omp state object for use by another team. + int slot = usedSlotIdx; + omptarget_nvptx_device_State[slot].Enqueue( + omptarget_nvptx_threadPrivateContext); + } +} + // Return true if the current target region is executed in SPMD mode. EXTERN int8_t __kmpc_is_spmd_exec_mode() { return isSPMDMode();