Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu =================================================================== --- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu +++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu @@ -122,32 +122,11 @@ } 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 @@ // 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()) { Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h =================================================================== --- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h +++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h @@ -154,13 +154,6 @@ // 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 @@ 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; Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/test/api/ignored.c =================================================================== --- openmp/trunk/libomptarget/deviceRTLs/nvptx/test/api/ignored.c +++ openmp/trunk/libomptarget/deviceRTLs/nvptx/test/api/ignored.c @@ -0,0 +1,38 @@ +// RUN: %compile-run-and-check + +#include +#include + +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; +}