[Openmp-commits] [PATCH] D54969: [OpenMP][libomptarget] Add new version of SPMD deinit kernel function with argument

[Gheorghe-Teodor Bercea via Phabricator via Openmp-commits]

gtbercea created this revision.
gtbercea added reviewers: ABataev, kkwli0, caomhin.
Herald added subscribers: openmp-commits, guansong.

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.


Repository:
  rOMP OpenMP

https://reviews.llvm.org/D54969

Files:
  libomptarget/deviceRTLs/nvptx/src/interface.h
  libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu


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();
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,


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D54969.175565.patch
Type: text/x-patch
Size: 2291 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20181127/a812b4eb/attachment.bin>