[Openmp-commits] [openmp] r358442 - [OPENMP][NVPTX]Fix dynamic scheduling in L2+ SPMD parallel regions.
Alexey Bataev via Openmp-commits
openmp-commits at lists.llvm.org
Mon Apr 15 13:15:20 PDT 2019
Author: abataev
Date: Mon Apr 15 13:15:20 2019
New Revision: 358442
URL: http://llvm.org/viewvc/llvm-project?rev=358442&view=rev
Log:
[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
Added:
openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/spmd_parallel_regions.cpp
Modified:
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu
openmp/trunk/libomptarget/deviceRTLs/nvptx/src/parallel.cu
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=358442&r1=358441&r2=358442&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/libcall.cu Mon Apr 15 13:15:20 2019
@@ -164,7 +164,8 @@ EXTERN int omp_get_level(void) {
if (isRuntimeUninitialized()) {
ASSERT0(LT_FUSSY, isSPMDMode(),
"Expected SPMD mode only with uninitialized runtime.");
- return parallelLevel;
+ // parallelLevel starts from 0, need to add 1 for correct level.
+ return parallelLevel + 1;
}
int level = 0;
omptarget_nvptx_TaskDescr *currTaskDescr =
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu?rev=358442&r1=358441&r2=358442&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/loop.cu Mon Apr 15 13:15:20 2019
@@ -205,8 +205,12 @@ public:
INLINE static void dispatch_init(kmp_Ident *loc, int32_t threadId,
kmp_sched_t schedule, T lb, T ub, ST st,
ST chunk) {
- ASSERT0(LT_FUSSY, checkRuntimeInitialized(loc),
- "Expected non-SPMD mode + initialized runtime.");
+ if (checkRuntimeUninitialized(loc)) {
+ // In SPMD mode no need to check parallelism level - dynamic scheduling
+ // may appear only in L2 parallel regions with lightweight runtime.
+ ASSERT0(LT_FUSSY, checkSPMDMode(loc), "Expected non-SPMD mode.");
+ return;
+ }
int tid = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor(tid);
T tnum = currTaskDescr->ThreadsInTeam();
@@ -439,8 +443,15 @@ public:
INLINE static int dispatch_next(kmp_Ident *loc, int32_t gtid, int32_t *plast,
T *plower, T *pupper, ST *pstride) {
- ASSERT0(LT_FUSSY, checkRuntimeInitialized(loc),
- "Expected non-SPMD mode + initialized runtime.");
+ if (checkRuntimeUninitialized(loc)) {
+ // In SPMD mode no need to check parallelism level - dynamic scheduling
+ // may appear only in L2 parallel regions with lightweight runtime.
+ ASSERT0(LT_FUSSY, checkSPMDMode(loc), "Expected non-SPMD mode.");
+ if (*plast)
+ return DISPATCH_FINISHED;
+ *plast = 1;
+ return DISPATCH_NOTFINISHED;
+ }
// ID of a thread in its own warp
// automatically selects thread or warp ID based on selected implementation
Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/parallel.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/parallel.cu?rev=358442&r1=358441&r2=358442&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/parallel.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/parallel.cu Mon Apr 15 13:15:20 2019
@@ -407,7 +407,7 @@ EXTERN uint16_t __kmpc_parallel_level(km
if (checkRuntimeUninitialized(loc)) {
ASSERT0(LT_FUSSY, checkSPMDMode(loc),
"Expected SPMD mode with uninitialized runtime.");
- return parallelLevel;
+ return parallelLevel + 1;
}
int threadId = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
Added: openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/spmd_parallel_regions.cpp
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/spmd_parallel_regions.cpp?rev=358442&view=auto
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/spmd_parallel_regions.cpp (added)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/spmd_parallel_regions.cpp Mon Apr 15 13:15:20 2019
@@ -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;
+}
More information about the Openmp-commits
mailing list