diff --git a/openmp/libomptarget/deviceRTLs/common/src/omptarget.cu b/openmp/libomptarget/deviceRTLs/common/src/omptarget.cu --- a/openmp/libomptarget/deviceRTLs/common/src/omptarget.cu +++ b/openmp/libomptarget/deviceRTLs/common/src/omptarget.cu @@ -68,7 +68,7 @@ // set number of threads and thread limit in team to started value omptarget_nvptx_TaskDescr *currTaskDescr = omptarget_nvptx_threadPrivateContext->GetTopLevelTaskDescr(threadId); - nThreads = GetNumberOfWorkersInTeam(); + nThreads = GetNumberOfThreadsInBlock(); threadLimit = ThreadLimit; } diff --git a/openmp/libomptarget/deviceRTLs/nvptx/test/api/get_max_threads.c b/openmp/libomptarget/deviceRTLs/nvptx/test/api/get_max_threads.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/deviceRTLs/nvptx/test/api/get_max_threads.c @@ -0,0 +1,22 @@ +// RUN: %compile-run-and-check +#include +#include + +int main(){ + int max_threads = -1; + int num_threads = -1; + + #pragma omp target map(tofrom: max_threads) + max_threads = omp_get_max_threads(); + + #pragma omp target parallel map(tofrom: num_threads) + { + #pragma omp master + num_threads = omp_get_num_threads(); + } + + // CHECK: Max Threads: 128, Num Threads: 128 + printf("Max Threads: %d, Num Threads: %d\n", max_threads, num_threads); + + return 0; +} diff --git a/openmp/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c b/openmp/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c --- a/openmp/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c +++ b/openmp/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c @@ -19,7 +19,14 @@ { MaxThreadsL2 = omp_get_max_threads(); } } - // CHECK: Non-SPMD MaxThreadsL1 = 32 + //FIXME: This Non-SPMD kernel will have 32 active threads due to + // thread_limit. However, Non-SPMD MaxThreadsL1 is the total number of + // threads in block (64 in this case), which translates to worker + // threads + WARP_SIZE for Non-SPMD kernels and worker threads for SPMD + // kernels. According to the spec, omp_get_max_threads must return the + // max active threads possible between the two kernel types. + + // CHECK: Non-SPMD MaxThreadsL1 = 64 printf("Non-SPMD MaxThreadsL1 = %d\n", MaxThreadsL1); // CHECK: Non-SPMD MaxThreadsL2 = 1 printf("Non-SPMD MaxThreadsL2 = %d\n", MaxThreadsL2);