[OPENMP][NVPTX]Fix behavior of omp_in_parallel() function.

[ABataev]

Summary

According to the OpenMP standard, the effect of the omp_in_parallel routine is to return true if the current task is enclosed by an active parallel region, and the parallel region is enclosed by the outermost initial task region on the device; otherwise it returns false.
Active parallel region - A parallel region that is executed by a team consisting of more than one thread.

[Hahnfeld]

My understanding of enclosed is that its meaning is recursive. So omp_in_parallel is to return true iff (at least) one of the surrounding parallel regions is active. That also matches the summary in OpenMP 5.0 which says

The omp_in_parallel routine returns true if the active-levels-var ICV is greater than zero; otherwise, it returns false.

and

active-levels-var - the number of nested active parallel regions that enclose the current task such that all of the parallel regions are enclosed by the outermost initial task region on the current device. There is one copy of this ICV per data environment.

In that sense I think the test is currently right and the change is not compliant with the standard.

[ABataev]

In D61380#1486150, @Hahnfeld wrote:

My understanding of enclosed is that its meaning is recursive. So omp_in_parallel is to return true iff (at least) one of the surrounding parallel regions is active. That also matches the summary in OpenMP 5.0 which says

The omp_in_parallel routine returns true if the active-levels-var ICV is greater than zero; otherwise, it returns false.

and

active-levels-var - the number of nested active parallel regions that enclose the current task such that all of the parallel regions are enclosed by the outermost initial task region on the current device. There is one copy of this ICV per data environment.

In that sense I think the test is currently right and the change is not compliant with the standard.

Hmm, the current implementation also does not look correct in this case. It does not take into account if the enclosing parallel region is active or not. If number of threads is 1, this regions is not active. Currently, it just checks if the innermost region is a parallel region regardless the number of threads in it (plus, i does not work correctly if we have outer parallel + inner task region)

libomptarget/deviceRTLs/nvptx/src/libcall.cu

 
 EXTERN int omp_get_num_procs(void) {
   int rc = GetNumberOfProcsInDevice(isSPMDMode());
   PRINT(LD_IO, "call omp_get_num_procs() returns %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_in_parallel(void) {
-  int rc = 0;
-  if (isRuntimeUninitialized()) {
-    ASSERT0(LT_FUSSY, isSPMDMode(),
-            "Expected SPMD mode only with uninitialized runtime.");
-    rc = 1;  // SPMD mode is always in parallel.
-  } else {
-    omptarget_nvptx_TaskDescr *currTaskDescr =
-        getMyTopTaskDescriptor(isSPMDMode());
-    if (currTaskDescr->InParallelRegion()) {
-      rc = 1;
-    }
-  }
+  // The effect of the omp_in_parallel routine is to return true if the current
+  // task is enclosed by an active parallel region, and the parallel region is
+  // enclosed by the outermost initial task region on the device; otherwise it
+  // returns false.
+  // Active parallel region - A parallel region that is executed by a team
+  // consisting of more than one thread.
+  int rc = omp_get_num_threads() > 1 ? 1 : 0;
   PRINT(LD_IO, "call omp_in_parallel() returns %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_in_final(void) {
   // treat all tasks as final... Specs may expect runtime to keep
   // track more precisely if a task was actively set by users... This
   // is not explicitely specified; will treat as if runtime can

libomptarget/deviceRTLs/nvptx/test/parallel/nested.c

#pragma omp target map(inParallel, numThreads, threadNum, check1[:], check2[:])
     #pragma omp parallel num_threads(NumThreads)
     {
       int id = omp_get_thread_num();
       check1[id] += omp_get_num_threads() + omp_in_parallel();
 
       // Expecting serialized parallel region.
       #pragma omp parallel
       {
-        // Expected to be 1.
+        // Expected to be 0.
         int nestedInParallel = omp_in_parallel();
         // Expected to be 1.
         int nestedNumThreads = omp_get_num_threads();
         // Expected to be 0.
         int nestedThreadNum = omp_get_thread_num();
         #pragma omp atomic
         check2[id] += nestedInParallel + nestedNumThreads + nestedThreadNum;
       }
if (i < NumThreads) {
         printf("invalid: check1[%d] should be %d, is %d\n", i, Expected, check1[i]);
       }
     } else if (check1[i] != 0) {
       printf("invalid: check1[%d] should be 0, is %d\n", i, check1[i]);
     }
 
     // Check serialized parallel region.
     if (i < NumThreads) {
-      if (check2[i] != 2) {
-        printf("invalid: check2[%d] should be 2, is %d\n", i, check2[i]);
+      if (check2[i] != 1) {
+        printf("invalid: check2[%d] should be 1, is %d\n", i, check2[i]);
       }
     } else if (check2[i] != 0) {
       printf("invalid: check2[%d] should be 0, is %d\n", i, check2[i]);
     }
   }
 
   return 0;
 }