[libomptarget-nvptx] Add tests for nested parallelism

[Hahnfeld]

Summary

Clang trunk will serialize nested parallel regions. Check that this
is correctly reflected in various API methods.

[ABataev]

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

3–4 ↗

(On Diff #164405)

Are we allowed to use these includes in tests?

[Hahnfeld]

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

3–4 ↗

(On Diff #164405)

I don't see a compelling reason why we can't: they are used in tests of libomp and libomptarget and all lit.cfg add include paths to find omp.h.

[Hahnfeld]

Ping

[ABataev]

LG

openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/level.c

+// RUN: %compile-run-and-check
+
+#include <omp.h>
+#include <stdio.h>
+
+const int MaxThreads = 1024;
+const int NumThreads = 64;
+
+int main(int argc, char *argv[]) {
+  int level = -1, activeLevel = -1;
+  int check1[MaxThreads];
+  int check2[MaxThreads];
+  for (int i = 0; i < MaxThreads; i++) {
+    check1[i] = check2[i] = 0;
+  }
+
+  #pragma omp target map(level, activeLevel, check1[:], check2[:])
+  {
+    level = omp_get_level();
+    activeLevel = omp_get_active_level();
+
+    // Expecting active parallel region.
+    #pragma omp parallel num_threads(NumThreads)
+    {
+      int id = omp_get_thread_num();
+      // Multiply return value of omp_get_level by 5 to avoid that this test
+      // passes if both API calls return wrong values.
+      check1[id] += omp_get_level() * 5 + omp_get_active_level();
+
+      // Expecting serialized parallel region.
+      #pragma omp parallel
+      {
+        #pragma omp atomic
+        check2[id] += omp_get_level() * 5 + omp_get_active_level();
+      }
+    }
+  }
+
+  // CHECK: target: level = 0, activeLevel = 0
+  printf("target: level = %d, activeLevel = %d\n", level, activeLevel);
+
+  // CHECK-NOT: invalid
+  for (int i = 0; i < MaxThreads; i++) {
+    // Check active parallel region:
+    // omp_get_level() = 1, omp_get_active_level() = 1
+    const int Expected1 = 6;
+
+    if (i < NumThreads) {
+      if (check1[i] != Expected1) {
+        printf("invalid: check1[%d] should be %d, is %d\n", i, Expected1, check1[i]);
+      }
+    } else if (check1[i] != 0) {
+      printf("invalid: check1[%d] should be 0, is %d\n", i, check1[i]);
+    }
+
+    // Check serialized parallel region:
+    // omp_get_level() = 2, omp_get_active_level() = 1
+    const int Expected2 = 11;
+    if (i < NumThreads) {
+      if (check2[i] != Expected2) {
+        printf("invalid: check2[%d] should be %d, is %d\n", i, Expected2, check2[i]);
+      }
+    } else if (check2[i] != 0) {
+      printf("invalid: check2[%d] should be 0, is %d\n", i, check2[i]);
+    }
+  }
+
+  return 0;
+}

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

+// RUN: %compile-run-and-check
+
+#include <omp.h>
+#include <stdio.h>
+
+const int MaxThreads = 1024;
+const int NumThreads = 64;
+
+int main(int argc, char *argv[]) {
+  int inParallel = -1, numThreads = -1, threadNum = -1;
+  int check1[MaxThreads];
+  int check2[MaxThreads];
+  for (int i = 0; i < MaxThreads; i++) {
+    check1[i] = check2[i] = 0;
+  }
+
+  #pragma omp target map(inParallel, numThreads, threadNum, check1[:], check2[:])
+  {
+    inParallel = omp_in_parallel();
+    numThreads = omp_get_num_threads();
+    threadNum = omp_get_thread_num();
+
+    // Expecting active parallel region.
+    #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.
+        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;
+      }
+    }
+  }
+
+  // CHECK: target: inParallel = 0, numThreads = 1, threadNum = 0
+  printf("target: inParallel = %d, numThreads = %d, threadNum = %d\n",
+         inParallel, numThreads, threadNum);
+
+  // CHECK-NOT: invalid
+  for (int i = 0; i < MaxThreads; i++) {
+    // Check that all threads reported
+    // omp_get_num_threads() = 64, omp_in_parallel() = 1.
+    int Expected = NumThreads + 1;
+    if (i < NumThreads) {
+      if (check1[i] != Expected) {
+        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]);
+      }
+    } else if (check2[i] != 0) {
+      printf("invalid: check2[%d] should be 0, is %d\n", i, check2[i]);
+    }
+  }
+
+  return 0;
+}