Index: libomptarget/deviceRTLs/nvptx/src/libcall.cu =================================================================== --- libomptarget/deviceRTLs/nvptx/src/libcall.cu +++ libomptarget/deviceRTLs/nvptx/src/libcall.cu @@ -101,18 +101,13 @@ } EXTERN int omp_in_parallel(void) { - int rc = 0; - if (isRuntimeUninitialized()) { - ASSERT0(LT_FUSSY, isSPMDMode(), - "Expected SPMD mode only with uninitialized runtime."); - rc = 1; // SPMD mode is always in parallel. - } else { - omptarget_nvptx_TaskDescr *currTaskDescr = - getMyTopTaskDescriptor(isSPMDMode()); - if (currTaskDescr->InParallelRegion()) { - rc = 1; - } - } + // The effect of the omp_in_parallel routine is to return true if the current + // task is enclosed by an active parallel region, and the parallel region is + // enclosed by the outermost initial task region on the device; otherwise it + // returns false. + // Active parallel region - A parallel region that is executed by a team + // consisting of more than one thread. + int rc = omp_get_num_threads() > 1 ? 1 : 0; PRINT(LD_IO, "call omp_in_parallel() returns %d\n", rc); return rc; } Index: libomptarget/deviceRTLs/nvptx/test/parallel/nested.c =================================================================== --- libomptarget/deviceRTLs/nvptx/test/parallel/nested.c +++ libomptarget/deviceRTLs/nvptx/test/parallel/nested.c @@ -29,7 +29,7 @@ // Expecting serialized parallel region. #pragma omp parallel { - // Expected to be 1. + // Expected to be 0. int nestedInParallel = omp_in_parallel(); // Expected to be 1. int nestedNumThreads = omp_get_num_threads(); @@ -60,8 +60,8 @@ // Check serialized parallel region. if (i < NumThreads) { - if (check2[i] != 2) { - printf("invalid: check2[%d] should be 2, is %d\n", i, check2[i]); + if (check2[i] != 1) { + printf("invalid: check2[%d] should be 1, is %d\n", i, check2[i]); } } else if (check2[i] != 0) { printf("invalid: check2[%d] should be 0, is %d\n", i, check2[i]);