[OPENMP][NVPTX]Fix dynamic scheduling in L2+ SPMD parallel regions.

Summary:
If the kernel is executed in SPMD mode and the L2+ parallel for region
with the dynamic scheduling is executed, dynamic scheduling functions
are called. They expect full runtime support, but SPMD kernels may be
executed without the full runtime. It leads to the runtime crash of the
compiled program. Patch fixes this problem + fixes handling of the
parallelism level in SPMD mode, which is required as part of this patch.

Reviewers: gtbercea, kkwli0, grokos

Subscribers: guansong, jdoerfert, openmp-commits, caomhin

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D60578

llvm-svn: 358442

1 files changed, 30 insertions, 0 deletions

diff --git a/openmp/libomptarget/deviceRTLs/nvptx/test/parallel/spmd_parallel_regions.cpp b/openmp/libomptarget/deviceRTLs/nvptx/test/parallel/spmd_parallel_regions.cpp
new file mode 100644
index 00000000000..bf5843bedb2
--- /dev/null
+++ b/openmp/libomptarget/deviceRTLs/nvptx/test/parallel/spmd_parallel_regions.cpp
@@ -0,0 +1,30 @@
+// RUN: %compilexx-run-and-check
+
+#include <stdio.h>
+#include <omp.h>
+
+int main(void) {
+  int isHost = -1;
+  int ParallelLevel1, ParallelLevel2 = -1;
+
+#pragma omp target parallel map(from: isHost, ParallelLevel1, ParallelLevel2)
+  {
+    isHost = omp_is_initial_device();
+    ParallelLevel1 = omp_get_level();
+#pragma omp parallel for schedule(dynamic) lastprivate(ParallelLevel2)
+    for (int I = 0; I < 10; ++I)
+      ParallelLevel2 = omp_get_level();
+  }
+
+  if (isHost < 0) {
+    printf("Runtime error, isHost=%d\n", isHost);
+  }
+
+  // CHECK: Target region executed on the device
+  printf("Target region executed on the %s\n", isHost ? "host" : "device");
+  // CHECK: Parallel level in SPMD mode: L1 is 1, L2 is 2
+  printf("Parallel level in SPMD mode: L1 is %d, L2 is %d\n", ParallelLevel1,
+         ParallelLevel2);
+
+  return isHost;
+}