diff options
author | Alexey Bataev <a.bataev@hotmail.com> | 2019-05-02 14:52:52 +0000 |
---|---|---|
committer | Alexey Bataev <a.bataev@hotmail.com> | 2019-05-02 14:52:52 +0000 |
commit | db75b85d572b5526020f54055f37fc6bba22e08b (patch) | |
tree | 33b8d952586c002a6684cf8fa9b553968d212617 | |
parent | 573fae727555eddbc20088d0b7e5abe547bfcf39 (diff) |
[OPENMP][NVPTX]Improve omp_get_max_threads() function.
Summary:
Function omp_get_max_threads() can always return 1 if current execution
mode is SPMD.
Reviewers: grokos, gtbercea, kkwli0
Subscribers: guansong, jdoerfert, caomhin, openmp-commits
Tags: #openmp
Differential Revision: https://reviews.llvm.org/D61379
git-svn-id: https://llvm.org/svn/llvm-project/openmp/trunk@359792 91177308-0d34-0410-b5e6-96231b3b80d8
-rw-r--r-- | libomptarget/deviceRTLs/nvptx/src/libcall.cu | 7 | ||||
-rw-r--r-- | libomptarget/deviceRTLs/nvptx/test/api/max_threads.c | 46 |
2 files changed, 48 insertions, 5 deletions
diff --git a/libomptarget/deviceRTLs/nvptx/src/libcall.cu b/libomptarget/deviceRTLs/nvptx/src/libcall.cu index 3a0c39c..26b0f4e 100644 --- a/libomptarget/deviceRTLs/nvptx/src/libcall.cu +++ b/libomptarget/deviceRTLs/nvptx/src/libcall.cu @@ -54,14 +54,11 @@ EXTERN int omp_get_num_threads(void) { } EXTERN int omp_get_max_threads(void) { - if (isRuntimeUninitialized()) { - ASSERT0(LT_FUSSY, isSPMDMode(), - "Expected SPMD mode only with uninitialized runtime."); + if (isSPMDMode()) // We're already in parallel region. return 1; // default is 1 thread avail - } omptarget_nvptx_TaskDescr *currTaskDescr = - getMyTopTaskDescriptor(isSPMDMode()); + getMyTopTaskDescriptor(/*isSPMDExecutionMode=*/false); int rc = 1; // default is 1 thread avail if (!currTaskDescr->InParallelRegion()) { // Not currently in a parallel region, return what was set. diff --git a/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c b/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c new file mode 100644 index 0000000..d0d9f31 --- /dev/null +++ b/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c @@ -0,0 +1,46 @@ +// RUN: %compile-run-and-check + +#include <omp.h> +#include <stdio.h> + +int main(int argc, char *argv[]) { + int MaxThreadsL1 = -1, MaxThreadsL2 = -1; + +#pragma omp declare reduction(unique:int \ + : omp_out = (omp_in == 1 ? omp_in : omp_out)) \ + initializer(omp_priv = -1) + + // Non-SPMD mode. +#pragma omp target teams map(MaxThreadsL1, MaxThreadsL2) thread_limit(32) \ + num_teams(1) + { + MaxThreadsL1 = omp_get_max_threads(); +#pragma omp parallel reduction(unique : MaxThreadsL2) + { MaxThreadsL2 = omp_get_max_threads(); } + } + + // CHECK: Non-SPMD MaxThreadsL1 = 32 + printf("Non-SPMD MaxThreadsL1 = %d\n", MaxThreadsL1); + // CHECK: Non-SPMD MaxThreadsL2 = 1 + printf("Non-SPMD MaxThreadsL2 = %d\n", MaxThreadsL2); + + // SPMD mode with full runtime + MaxThreadsL2 = -1; +#pragma omp target parallel reduction(unique : MaxThreadsL2) + { MaxThreadsL2 = omp_get_max_threads(); } + + // CHECK: SPMD with full runtime MaxThreadsL2 = 1 + printf("SPMD with full runtime MaxThreadsL2 = %d\n", MaxThreadsL2); + + // SPMD mode without runtime + MaxThreadsL2 = -1; +#pragma omp target parallel for reduction(unique : MaxThreadsL2) + for (int I = 0; I < 2; ++I) { + MaxThreadsL2 = omp_get_max_threads(); + } + + // CHECK: SPMD without runtime MaxThreadsL2 = 1 + printf("SPMD without runtime MaxThreadsL2 = %d\n", MaxThreadsL2); + + return 0; +} |