aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAlexey Bataev <a.bataev@hotmail.com>2019-05-02 14:52:52 +0000
committerAlexey Bataev <a.bataev@hotmail.com>2019-05-02 14:52:52 +0000
commitdb75b85d572b5526020f54055f37fc6bba22e08b (patch)
tree33b8d952586c002a6684cf8fa9b553968d212617
parent573fae727555eddbc20088d0b7e5abe547bfcf39 (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.cu7
-rw-r--r--libomptarget/deviceRTLs/nvptx/test/api/max_threads.c46
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;
+}