[Openmp-commits] [PATCH] D74092: Changed omp_get_max_threads() implementation to more closely match spec description.

[Jon Chesterfield via Phabricator via Openmp-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rG190a11148b75: Changed omp_get_max_threads() implementation to more closely match spec… (authored by estewart08, committed by JonChesterfield).

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D74092/new/

https://reviews.llvm.org/D74092

Files:
  openmp/libomptarget/deviceRTLs/common/src/omptarget.cu
  openmp/libomptarget/deviceRTLs/nvptx/test/api/get_max_threads.c
  openmp/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c


Index: openmp/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c
===================================================================
--- openmp/libomptarget/deviceRTLs/nvptx/test/api/max_threads.c
+++ 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);
Index: openmp/libomptarget/deviceRTLs/nvptx/test/api/get_max_threads.c
===================================================================
--- /dev/null
+++ openmp/libomptarget/deviceRTLs/nvptx/test/api/get_max_threads.c
@@ -0,0 +1,22 @@
+// RUN: %compile-run-and-check
+#include <omp.h>
+#include <stdio.h>
+
+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;
+}
Index: openmp/libomptarget/deviceRTLs/common/src/omptarget.cu
===================================================================
--- openmp/libomptarget/deviceRTLs/common/src/omptarget.cu
+++ 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;
 }
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D74092.244292.patch
Type: text/x-patch
Size: 2349 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20200212/4b637f32/attachment-0001.bin>