diff options
-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; +} |