diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/interface.h b/openmp/libomptarget/deviceRTLs/nvptx/src/interface.h index 40321a624cde..2c2beae31755 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 5d95eb1c77ed..b0aef62ee9bb 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;