[libomptarget-nvptx] Ignore calls to dynamic API

There is no support and according to the OpenMP 4.5, p238:7-9:

    For implementations that do not support dynamic adjustment
    of the number of threads this routine has no effect: the
    value of dyn-var remains false.

Add a test that cancellation and nested parallelism aren't
supported either.

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

3 files changed, 44 insertions, 37 deletions

diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/libcall.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/libcall.cu
index 929f9db9962..352a4f2f0b4 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/libcall.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/libcall.cu
@@ -122,32 +122,11 @@ EXTERN int omp_in_final(void) {
 }
 
 EXTERN void omp_set_dynamic(int flag) {
-  PRINT(LD_IO, "call omp_set_dynamic(%d)\n", flag);
-  if (isRuntimeUninitialized()) {
-    ASSERT0(LT_FUSSY, isSPMDMode(),
-            "Expected SPMD mode only with uninitialized runtime.");
-    return;
-  }
-
-  omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
-  if (flag) {
-    currTaskDescr->SetDynamic();
-  } else {
-    currTaskDescr->ClearDynamic();
-  }
+  PRINT(LD_IO, "call omp_set_dynamic(%d) is ignored (no support)\n", flag);
 }
 
 EXTERN int omp_get_dynamic(void) {
   int rc = 0;
-  if (isRuntimeUninitialized()) {
-    ASSERT0(LT_FUSSY, isSPMDMode(),
-            "Expected SPMD mode only with uninitialized runtime.");
-    return rc;
-  }
-  omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor();
-  if (currTaskDescr->IsDynamic()) {
-    rc = 1;
-  }
   PRINT(LD_IO, "call omp_get_dynamic() returns %d\n", rc);
   return rc;
 }
@@ -237,14 +216,13 @@ EXTERN int omp_get_ancestor_thread_num(int level) {
           // print current state
           omp_sched_t sched = currTaskDescr->GetRuntimeSched();
           PRINT(LD_ALL,
-                "task descr %s %d: %s, in par %d, dyn %d, rt sched %d,"
+                "task descr %s %d: %s, in par %d, rt sched %d,"
                 " chunk %" PRIu64 "; tid %d, tnum %d, nthreads %d\n",
                 "ancestor", steps,
                 (currTaskDescr->IsParallelConstruct() ? "par" : "task"),
-                currTaskDescr->InParallelRegion(), currTaskDescr->IsDynamic(),
-                sched, currTaskDescr->RuntimeChunkSize(),
-                currTaskDescr->ThreadId(), currTaskDescr->ThreadsInTeam(),
-                currTaskDescr->NThreads());
+                currTaskDescr->InParallelRegion(), sched,
+                currTaskDescr->RuntimeChunkSize(), currTaskDescr->ThreadId(),
+                currTaskDescr->ThreadsInTeam(), currTaskDescr->NThreads());
         }
 
         if (currTaskDescr->IsParallelConstruct()) {
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
index 2d7f10d9cab..e0d4c1679cd 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
@@ -154,13 +154,6 @@ public:
   // methods for flags
   INLINE omp_sched_t GetRuntimeSched();
   INLINE void SetRuntimeSched(omp_sched_t sched);
-  INLINE int IsDynamic() { return items.flags & TaskDescr_IsDynamic; }
-  INLINE void SetDynamic() {
-    items.flags = items.flags | TaskDescr_IsDynamic;
-  }
-  INLINE void ClearDynamic() {
-    items.flags = items.flags & (~TaskDescr_IsDynamic);
-  }
   INLINE int InParallelRegion() { return items.flags & TaskDescr_InPar; }
   INLINE int InL2OrHigherParallelRegion() {
     return items.flags & TaskDescr_InParL2P;
@@ -196,15 +189,13 @@ public:
   INLINE void RestoreLoopData() const;
 
 private:
-  // bits for flags: (7 used, 1 free)
+  // bits for flags: (6 used, 2 free)
   //   3 bits (SchedMask) for runtime schedule
-  //   1 bit (IsDynamic) for dynamic schedule (false = static)
   //   1 bit (InPar) if this thread has encountered one or more parallel region
   //   1 bit (IsParConstr) if ICV for a parallel region (false = explicit task)
   //   1 bit (InParL2+) if this thread has encountered L2 or higher parallel
   //   region
   static const uint8_t TaskDescr_SchedMask = (0x1 | 0x2 | 0x4);
-  static const uint8_t TaskDescr_IsDynamic = 0x8;
   static const uint8_t TaskDescr_InPar = 0x10;
   static const uint8_t TaskDescr_IsParConstr = 0x20;
   static const uint8_t TaskDescr_InParL2P = 0x40;
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/test/api/ignored.c b/openmp/libomptarget/deviceRTLs/nvptx/test/api/ignored.c
new file mode 100644
index 00000000000..1fa9ae024f6
--- /dev/null
+++ b/openmp/libomptarget/deviceRTLs/nvptx/test/api/ignored.c
@@ -0,0 +1,38 @@
+// RUN: %compile-run-and-check
+
+#include <omp.h>
+#include <stdio.h>
+
+const int MaxThreads = 1024;
+
+int main(int argc, char *argv[]) {
+  int cancellation = -1, dynamic = -1, nested = -1, maxActiveLevels = -1;
+
+  #pragma omp target map(cancellation, dynamic, nested, maxActiveLevels)
+  {
+    // libomptarget-nvptx doesn't support cancellation.
+    cancellation = omp_get_cancellation();
+
+    // No support for dynamic adjustment of the number of threads.
+    omp_set_dynamic(1);
+    dynamic = omp_get_dynamic();
+
+    // libomptarget-nvptx doesn't support nested parallelism.
+    omp_set_nested(1);
+    nested = omp_get_nested();
+
+    omp_set_max_active_levels(42);
+    maxActiveLevels = omp_get_max_active_levels();
+  }
+
+  // CHECK: cancellation = 0
+  printf("cancellation = %d\n", cancellation);
+  // CHECK: dynamic = 0
+  printf("dynamic = %d\n", dynamic);
+  // CHECK: nested = 0
+  printf("nested = %d\n", nested);
+  // CHECK: maxActiveLevels = 1
+  printf("maxActiveLevels = %d\n", maxActiveLevels);
+
+  return 0;
+}