[libomptarget-nvptx] Ignore calls to dynamic API

[Hahnfeld]

Summary

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.

[Hahnfeld]

Ping

[grokos]

LG

libomptarget/deviceRTLs/nvptx/src/libcall.cu

EXTERN int omp_in_final(void) {
   // is not explicitely specified; will treat as if runtime can
   // actively decide to put a non-final task into a final one.
   int rc = 1;
   PRINT(LD_IO, "call omp_in_final() returns %d\n", rc);
   return rc;
 }
 
 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;
 }
 
 EXTERN void omp_set_nested(int flag) {
   PRINT(LD_IO, "call omp_set_nested(%d) is ignored (no nested support)\n",
         flag);
 }
if (level <= totLevel) {
       PRINT(LD_IO, "backtrack %d steps\n", steps);
       ASSERT0(LT_FUSSY, currTaskDescr,
               "do not expect fct to be called in a non-active thread");
       do {
         if (DON(LD_IOD)) {
           // 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()) {
           // found the level
           if (!steps) {
             rc = currTaskDescr->ThreadId();
             break;
           }

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

 ////////////////////////////////////////////////////////////////////////////////
 // task ICV and (implicit & explicit) task state
 
 class omptarget_nvptx_TaskDescr {
 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;
   }
   INLINE int IsParallelConstruct() {
     return items.flags & TaskDescr_IsParConstr;
   }
   INLINE int IsTaskConstruct() { return !IsParallelConstruct(); }
INLINE void CopyToWorkDescr(omptarget_nvptx_TaskDescr *masterTaskDescr,
                               uint16_t tnum);
   INLINE void CopyFromWorkDescr(omptarget_nvptx_TaskDescr *workTaskDescr);
   INLINE void CopyConvergentParent(omptarget_nvptx_TaskDescr *parentTaskDescr,
                                    uint16_t tid, uint16_t tnum);
   INLINE void SaveLoopData();
   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;
 
   struct SavedLoopDescr_items {
     int64_t loopUpperBound;
     int64_t nextLowerBound;
     int64_t chunk;

libomptarget/deviceRTLs/nvptx/test/api/ignored.c

+// 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;
+}