[Openmp-commits] [PATCH] D61379: [OPENMP][NVPTX]Improve omp_get_max_threads() function.

[Alexey Bataev via Phabricator via Openmp-commits]

This revision was not accepted when it landed; it landed in state "Needs Review".
This revision was automatically updated to reflect the committed changes.
Closed by commit rOMP359792: [OPENMP][NVPTX]Improve omp_get_max_threads() function. (authored by ABataev, committed by ).

Changed prior to commit:
  https://reviews.llvm.org/D61379?vs=197546&id=197787#toc

Repository:
  rOMP OpenMP

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D61379/new/

https://reviews.llvm.org/D61379

Files:
  libomptarget/deviceRTLs/nvptx/src/libcall.cu
  libomptarget/deviceRTLs/nvptx/test/api/max_threads.c


Index: libomptarget/deviceRTLs/nvptx/src/libcall.cu
===================================================================
--- libomptarget/deviceRTLs/nvptx/src/libcall.cu
+++ libomptarget/deviceRTLs/nvptx/src/libcall.cu
@@ -54,14 +54,11 @@
 }
 
 EXTERN int omp_get_max_threads(void) {
-  if (isRuntimeUninitialized()) {
-    ASSERT0(LT_FUSSY, isSPMDMode(),
-            "Expected SPMD mode only with uninitialized runtime.");
+  if (isSPMDMode())
     // We're already in parallel region.
     return 1;  // default is 1 thread avail
-  }
   omptarget_nvptx_TaskDescr *currTaskDescr =
-      getMyTopTaskDescriptor(isSPMDMode());
+      getMyTopTaskDescriptor(/*isSPMDExecutionMode=*/false);
   int rc = 1; // default is 1 thread avail
   if (!currTaskDescr->InParallelRegion()) {
     // Not currently in a parallel region, return what was set.
Index: libomptarget/deviceRTLs/nvptx/test/api/max_threads.c
===================================================================
--- libomptarget/deviceRTLs/nvptx/test/api/max_threads.c
+++ libomptarget/deviceRTLs/nvptx/test/api/max_threads.c
@@ -0,0 +1,46 @@
+// RUN: %compile-run-and-check
+
+#include <omp.h>
+#include <stdio.h>
+
+int main(int argc, char *argv[]) {
+  int MaxThreadsL1 = -1, MaxThreadsL2 = -1;
+
+#pragma omp declare reduction(unique:int                                       \
+                              : omp_out = (omp_in == 1 ? omp_in : omp_out))    \
+    initializer(omp_priv = -1)
+
+  // Non-SPMD mode.
+#pragma omp target teams map(MaxThreadsL1, MaxThreadsL2) thread_limit(32)      \
+    num_teams(1)
+  {
+    MaxThreadsL1 = omp_get_max_threads();
+#pragma omp parallel reduction(unique : MaxThreadsL2)
+    { MaxThreadsL2 = omp_get_max_threads(); }
+  }
+
+  // CHECK: Non-SPMD MaxThreadsL1 = 32
+  printf("Non-SPMD MaxThreadsL1 = %d\n", MaxThreadsL1);
+  // CHECK: Non-SPMD MaxThreadsL2 = 1
+  printf("Non-SPMD MaxThreadsL2 = %d\n", MaxThreadsL2);
+
+  // SPMD mode with full runtime
+  MaxThreadsL2 = -1;
+#pragma omp target parallel reduction(unique : MaxThreadsL2)
+  { MaxThreadsL2 = omp_get_max_threads(); }
+
+  // CHECK: SPMD with full runtime MaxThreadsL2 = 1
+  printf("SPMD with full runtime MaxThreadsL2 = %d\n", MaxThreadsL2);
+
+  // SPMD mode without runtime
+  MaxThreadsL2 = -1;
+#pragma omp target parallel for reduction(unique : MaxThreadsL2)
+  for (int I = 0; I < 2; ++I) {
+    MaxThreadsL2 = omp_get_max_threads();
+  }
+
+  // CHECK: SPMD without runtime MaxThreadsL2 = 1
+  printf("SPMD without runtime MaxThreadsL2 = %d\n", MaxThreadsL2);
+
+  return 0;
+}


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D61379.197787.patch
Type: text/x-patch
Size: 2563 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20190502/a676cdd4/attachment.bin>