[Openmp-commits] [openmp] r343381 - [libomptarget-nvptx] Ignore calls to dynamic API
Jonas Hahnfeld via Openmp-commits
openmp-commits at lists.llvm.org
Sat Sep 29 09:02:25 PDT 2018
Author: hahnfeld
Date: Sat Sep 29 09:02:25 2018
New Revision: 343381
URL: http://llvm.org/viewvc/llvm-project?rev=343381&view=rev
Log:
[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
Added:
openmp/trunk/libomptarget/deviceRTLs/nvptx/test/api/
openmp/trunk/libomptarget/deviceRTLs/nvptx/test/api/ignored.c
Modified:
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu?rev=343381&r1=343380&r2=343381&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu Sat Sep 29 09:02:25 2018
@@ -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(i
// 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()) {
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h?rev=343381&r1=343380&r2=343381&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h Sat Sep 29 09:02:25 2018
@@ -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;
Added: openmp/trunk/libomptarget/deviceRTLs/nvptx/test/api/ignored.c
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/test/api/ignored.c?rev=343381&view=auto
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/test/api/ignored.c (added)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/test/api/ignored.c Sat Sep 29 09:02:25 2018
@@ -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;
+}
More information about the Openmp-commits
mailing list