public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] Revert "OpenMP 5.0: Implement relaxation of implicit map vs. existing device mappings"
@ 2021-05-31 9:02 Chung-Lin Tang
0 siblings, 0 replies; only message in thread
From: Chung-Lin Tang @ 2021-05-31 9:02 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:d6b9d4351d376b5298bfeff6df8f17ebc9033143
commit d6b9d4351d376b5298bfeff6df8f17ebc9033143
Author: Chung-Lin Tang <cltang@codesourcery.com>
Date: Wed May 26 19:11:35 2021 +0800
Revert "OpenMP 5.0: Implement relaxation of implicit map vs. existing device mappings"
This reverts commit 9a4a899134779cef4553f0c616e9577b379783ac.
Diff:
---
gcc/gimplify.c | 11 +--
gcc/omp-low.c | 13 ----
.../c-c++-common/goacc/combined-reduction.c | 4 +-
.../c-c++-common/goacc/firstprivate-mappings-1.c | 6 +-
gcc/testsuite/c-c++-common/goacc/mdc-1.c | 2 +-
gcc/testsuite/c-c++-common/goacc/reduction-1.c | 4 +-
gcc/testsuite/c-c++-common/goacc/reduction-2.c | 4 +-
gcc/testsuite/c-c++-common/goacc/reduction-3.c | 4 +-
gcc/testsuite/c-c++-common/goacc/reduction-4.c | 4 +-
gcc/testsuite/c-c++-common/goacc/reduction-8.c | 10 ++-
.../c-c++-common/gomp/target-implicit-map-1.c | 39 -----------
.../g++.dg/goacc/firstprivate-mappings-1.C | 2 +-
gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 6 +-
gcc/testsuite/g++.dg/gomp/target-this-3.C | 4 +-
gcc/testsuite/g++.dg/gomp/target-this-4.C | 4 +-
gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 | 8 +--
gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 | 2 +-
.../goacc/private-explicit-kernels-1.f95 | 4 +-
.../goacc/private-predetermined-kernels-1.f95 | 4 +-
gcc/tree-pretty-print.c | 3 -
gcc/tree.h | 4 --
include/ChangeLog.omp | 9 ---
include/gomp-constants.h | 17 +----
libgomp/target.c | 78 ++++++----------------
.../libgomp.c-c++-common/target-implicit-map-1.c | 31 ---------
25 files changed, 57 insertions(+), 220 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index d9d574bee09..8720f876359 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10817,7 +10817,6 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
gcc_unreachable ();
}
OMP_CLAUSE_SET_MAP_KIND (clause, kind);
- OMP_CLAUSE_MAP_IMPLICIT_P (clause) = 1;
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -11446,15 +11445,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
list_p = &OMP_CLAUSE_CHAIN (c);
}
- /* Add in any implicit data sharing. Implicit clauses are added at the start
- of the clause list, but after any non-map clauses. */
+ /* Add in any implicit data sharing. */
struct gimplify_adjust_omp_clauses_data data;
- tree *implicit_add_list_p = orig_list_p;
- while (*implicit_add_list_p
- && OMP_CLAUSE_CODE (*implicit_add_list_p) != OMP_CLAUSE_MAP)
- implicit_add_list_p = &OMP_CLAUSE_CHAIN (*implicit_add_list_p);
-
- data.list_p = implicit_add_list_p;
+ data.list_p = list_p;
data.pre_p = pre_p;
splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &data);
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 9302d6ab03e..a16603ab441 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -13143,19 +13143,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
else if (integer_nonzerop (s))
tkind_zero = tkind;
}
- if (tkind_zero == tkind
- && OMP_CLAUSE_MAP_IMPLICIT_P (c)
- && (((tkind & GOMP_MAP_FLAG_SPECIAL_BITS)
- & ~GOMP_MAP_IMPLICIT)
- == 0))
- {
- /* If this is an implicit map, and the GOMP_MAP_IMPLICIT
- bits are not interfered by other special bit encodings,
- then turn the GOMP_IMPLICIT_BIT flag on for the runtime
- to see. */
- tkind |= GOMP_MAP_IMPLICIT;
- tkind_zero = tkind;
- }
break;
case OMP_CLAUSE_FIRSTPRIVATE:
gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
index fa67e085c86..ecf23f59d66 100644
--- a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
+++ b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
@@ -23,7 +23,7 @@ main ()
return 0;
}
-/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction.+:v1. firstprivate.n. map.tofrom:v1" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction.+:v1. map.tofrom:v1" 1 "gimple" } } */
/* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_tofrom:n .len: 4..implicit.. map.force_tofrom:v1 .len: 4..implicit.." 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_tofrom:n .len: 4.. map.force_tofrom:v1 .len: 4.." 1 "gimple" } } */
/* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
index f43e4b46cb6..7987beaed9a 100644
--- a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
@@ -419,12 +419,12 @@ vla (int array_li)
copyout (array_so)
/* The gimplifier has created an implicit 'firstprivate' clause for the array
length.
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(from:array_so \[len: 4\]\)} omplower { target { ! c++ } } } }
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower { target { c++ } } } }
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\)} omplower { target { ! c++ } } } }
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower { target { c++ } } } }
(C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */
/* For C, non-LP64, the gimplifier has also created a mapping for the array
itself; PR90859.
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\[implicit\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) map\(from:array_so \[len: 4\]\) \[} omplower { target { c && { ! lp64 } } } } } */
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) \[} omplower { target { c && { ! lp64 } } } } } */
{
array_so = sizeof array;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
index 9f43de4f776..337c1f7cc77 100644
--- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -45,7 +45,7 @@ t1 ()
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.tofrom:s .len: 32..implicit.. map.attach:s.e .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-1.c b/gcc/testsuite/c-c++-common/goacc/reduction-1.c
index d9e3c380b8e..35bfc868708 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-1.c
@@ -68,5 +68,5 @@ main(void)
}
/* Check that default copy maps are generated for loop reductions. */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 7 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 7 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-2.c b/gcc/testsuite/c-c++-common/goacc/reduction-2.c
index 18dc03c93ac..9dba035adb6 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-2.c
@@ -50,5 +50,5 @@ main(void)
}
/* Check that default copy maps are generated for loop reductions. */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 4 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-3.c b/gcc/testsuite/c-c++-common/goacc/reduction-3.c
index 2311d4b0adb..669cd438113 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-3.c
@@ -50,5 +50,5 @@ main(void)
}
/* Check that default copy maps are generated for loop reductions. */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 4 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-4.c b/gcc/testsuite/c-c++-common/goacc/reduction-4.c
index 57823f8898f..5c3dfb19172 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-4.c
@@ -38,5 +38,5 @@ main(void)
}
/* Check that default copy maps are generated for loop reductions. */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-8.c b/gcc/testsuite/c-c++-common/goacc/reduction-8.c
index 8494e594665..8a0283f4ac3 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-8.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-8.c
@@ -87,10 +87,8 @@ main(void)
/* Check that default copy maps are generated for loop reductions. */
/* { dg-final { scan-tree-dump-times "reduction..:result. map.tofrom:result .len: 4.." 1 "gimple" } } */
-
-/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .len: 4..implicit.." 1 "gimple" } } */
-
-/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. firstprivate.result." 3 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. map.tofrom:result .len: 4.." 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. map.force_tofrom:result .len: 4..implicit.." 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .len: 4.." 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. firstprivate.result." 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map.tofrom:result .len: 4.. map.tofrom:array .len: 4000.." 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. map.force_tofrom:result .len: 4.." 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c
deleted file mode 100644
index 52944fdc65a..00000000000
--- a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c
+++ /dev/null
@@ -1,39 +0,0 @@
-/* { dg-do compile } */
-/* { dg-additional-options "-fdump-tree-gimple" } */
-#ifdef __cplusplus
-extern "C"
-#else
-extern
-#endif
-void abort (void);
-
-int
-main (void)
-{
- #define N 5
- int array[N][N];
-
- for (int i = 0; i < N; i++)
- {
- #pragma omp target enter data map(alloc: array[i:1][0:N])
-
- #pragma omp target
- for (int j = 0; j < N; j++)
- array[i][j] = i * 10 + j;
-
- #pragma omp target exit data map(from: array[i:1][0:N])
- }
-
- for (int i = 0; i < N; i++)
- for (int j = 0; j < N; j++)
- if (array[i][j] != i + j)
- abort ();
-
- return 0;
-}
-
-/* { dg-final { scan-tree-dump {#pragma omp target enter data map\(alloc:array\[[^]]+\]\[0\] \[len: [0-9]+\]\)} "gimple" } } */
-
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(i\) map\(tofrom:array \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */
-
-/* { dg-final { scan-tree-dump {#pragma omp target exit data map\(from:array\[[^]]+\]\[0\] \[len: [0-9]+\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
index 99a3bd472f7..1b1badb1a90 100644
--- a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
+++ b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
@@ -416,7 +416,7 @@ vla (int &array_li)
copyout (array_so)
/* The gimplifier has created an implicit 'firstprivate' clause for the array
length.
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower } }
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower } }
(C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */
{
array_so = sizeof array;
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
index e5a24d7abc4..7dceef80f47 100644
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -87,8 +87,8 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C
index 2755b4b58bd..08568f9284c 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C
@@ -100,6 +100,6 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
index 3703762f45a..3b2d5811350 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-4.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -102,6 +102,6 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
index e9f169f9517..e43d376fc28 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
@@ -33,10 +33,10 @@ end program main
! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } }
! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\\[implicit\\\]\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\[implicit\\\]\\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } }
! Expecting no mapping of un-referenced common-blocks variables
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
index 4cdfc5556b7..150f9304e46 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
@@ -44,4 +44,4 @@ end program test
! { dg-final { scan-tree-dump-times "private\\(m\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "reduction\\(\\+:sum\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
index fef512612bd..0c47045df9c 100644
--- a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
@@ -83,7 +83,7 @@ program test
!$acc kernels ! Explicit "private(i2_2_s)" clause cannot be specified here.
! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
- ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
do i2_2_s = 1, 100
!$acc loop private(j2_2_s) independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
@@ -234,7 +234,7 @@ program test
!$acc kernels ! Explicit "private(i3_5_s)" clause cannot be specified here.
! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
- ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
do i3_5_s = 1, 100
!$acc loop private(j3_5_s) independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
index 38459cfadf3..3357a20263e 100644
--- a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
@@ -83,7 +83,7 @@ program test
!$acc kernels
! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
- ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
do i2_2_s = 1, 100
!$acc loop independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
@@ -234,7 +234,7 @@ program test
!$acc kernels
! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
- ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
do i3_5_s = 1, 100
!$acc loop independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 667bdc20e4b..7d386fca8ee 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -996,9 +996,6 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
spc, flags, false);
pp_right_bracket (pp);
}
- if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_IMPLICIT_P (clause))
- pp_string (pp, "[implicit]");
pp_right_paren (pp);
break;
diff --git a/gcc/tree.h b/gcc/tree.h
index ec01f6593ed..520bfa671cc 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1665,10 +1665,6 @@ class auto_suppress_location_wrappers
map clause. */
#define OMP_CLAUSE_MAP_IMPLICIT(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.default_def_flag)
-/* Nonzero if this map clause was created through implicit data-mapping
- rules. */
-#define OMP_CLAUSE_MAP_IMPLICIT_P(NODE) \
- (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.deprecated_flag)
/* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
clause. */
diff --git a/include/ChangeLog.omp b/include/ChangeLog.omp
index c4b3e729533..a460e639d22 100644
--- a/include/ChangeLog.omp
+++ b/include/ChangeLog.omp
@@ -1,12 +1,3 @@
-2021-05-05 Chung-Lin Tang <cltang@codesourcery.com>
-
- * gomp-constants.h (GOMP_MAP_IMPLICIT): New special map kind bits value.
- (GOMP_MAP_FLAG_SPECIAL_BITS): Define helper mask for whole set of
- special map kind bits.
- (GOMP_MAP_NONCONTIG_ARRAY_P): Adjust test for non-contiguous array map
- kind bits to be more specific.
- (GOMP_MAP_IMPLICIT_P): New predicate macro for implicit map kinds.
-
2018-10-04 Cesar Philippidis <cesar@codesourcery.com>
Julian Brown <julian@codesourcery.com>
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index a995ad9379f..d9b80ac9732 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -46,16 +46,6 @@
| GOMP_MAP_FLAG_SPECIAL_0)
#define GOMP_MAP_DEEP_COPY (GOMP_MAP_FLAG_SPECIAL_4 \
| GOMP_MAP_FLAG_SPECIAL_2)
-/* This value indicates the map was created implicitly according to
- OpenMP rules. */
-#define GOMP_MAP_IMPLICIT (GOMP_MAP_FLAG_SPECIAL_3 \
- | GOMP_MAP_FLAG_SPECIAL_4)
-/* Mask for entire set of special map kind bits. */
-#define GOMP_MAP_FLAG_SPECIAL_BITS (GOMP_MAP_FLAG_SPECIAL_0 \
- | GOMP_MAP_FLAG_SPECIAL_1 \
- | GOMP_MAP_FLAG_SPECIAL_2 \
- | GOMP_MAP_FLAG_SPECIAL_3 \
- | GOMP_MAP_FLAG_SPECIAL_4)
/* Flag to force a specific behavior (or else, trigger a run-time error). */
#define GOMP_MAP_FLAG_FORCE (1 << 7)
@@ -235,12 +225,7 @@ enum gomp_map_kind
(GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM))
#define GOMP_MAP_NONCONTIG_ARRAY_P(X) \
- (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_NONCONTIG_ARRAY \
- || (X) == GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT)
-
-#define GOMP_MAP_IMPLICIT_P(X) \
- (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_IMPLICIT)
-
+ ((X) & GOMP_MAP_NONCONTIG_ARRAY)
/* Asynchronous behavior. Keep in sync with
libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */
diff --git a/libgomp/target.c b/libgomp/target.c
index 7a8fa87bd85..2633ab1e43f 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -510,7 +510,7 @@ static inline void
gomp_map_vars_existing (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, splay_tree_key oldn,
splay_tree_key newn, struct target_var_desc *tgt_var,
- unsigned char kind, bool always_to_flag, bool implicit,
+ unsigned char kind, bool always_to_flag,
struct gomp_coalesce_buf *cbuf,
htab_t *refcount_set)
{
@@ -522,22 +522,11 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind);
tgt_var->is_attach = false;
tgt_var->offset = newn->host_start - oldn->host_start;
-
- /* For implicit maps, old contained in new is valid. */
- bool implicit_subset = (implicit
- && newn->host_start <= oldn->host_start
- && oldn->host_end <= newn->host_end);
- if (implicit_subset)
- tgt_var->length = oldn->host_end - oldn->host_start;
- else
- tgt_var->length = newn->host_end - newn->host_start;
+ tgt_var->length = newn->host_end - newn->host_start;
if ((kind & GOMP_MAP_FLAG_FORCE)
- /* For implicit maps, old contained in new is valid. */
- || !(implicit_subset
- /* Otherwise, new contained inside old is considered valid. */
- || (oldn->host_start <= newn->host_start
- && newn->host_end <= oldn->host_end)))
+ || oldn->host_start > newn->host_start
+ || oldn->host_end < newn->host_end)
{
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("Trying to map into device [%p..%p) object when "
@@ -547,17 +536,11 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
}
if (GOMP_MAP_ALWAYS_TO_P (kind) || always_to_flag)
- {
- /* Implicit + always should not happen. If this does occur, below
- address/length adjustment is a TODO. */
- assert (!implicit_subset);
-
- gomp_copy_host2dev (devicep, aq,
- (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
- + newn->host_start - oldn->host_start),
- (void *) newn->host_start,
- newn->host_end - newn->host_start, false, cbuf);
- }
+ gomp_copy_host2dev (devicep, aq,
+ (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
+ + newn->host_start - oldn->host_start),
+ (void *) newn->host_start,
+ newn->host_end - newn->host_start, false, cbuf);
gomp_increment_refcount (oldn, refcount_set);
}
@@ -565,24 +548,8 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
static int
get_kind (bool short_mapkind, void *kinds, int idx)
{
- int val = (short_mapkind
- ? ((unsigned short *) kinds)[idx]
- : ((unsigned char *) kinds)[idx]);
-
- if (GOMP_MAP_IMPLICIT_P (val))
- val &= ~GOMP_MAP_IMPLICIT;
- return val;
-}
-
-
-static bool
-get_implicit (bool short_mapkind, void *kinds, int idx)
-{
- int val = (short_mapkind
- ? ((unsigned short *) kinds)[idx]
- : ((unsigned char *) kinds)[idx]);
-
- return GOMP_MAP_IMPLICIT_P (val);
+ return short_mapkind ? ((unsigned short *) kinds)[idx]
+ : ((unsigned char *) kinds)[idx];
}
static void
@@ -645,7 +612,6 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
struct splay_tree_s *mem_map = &devicep->mem_map;
struct splay_tree_key_s cur_node;
int kind;
- bool implicit;
const bool short_mapkind = true;
const int typemask = short_mapkind ? 0xff : 0x7;
@@ -653,14 +619,12 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
cur_node.host_end = cur_node.host_start + sizes[i];
splay_tree_key n2 = splay_tree_lookup (mem_map, &cur_node);
kind = get_kind (short_mapkind, kinds, i);
- implicit = get_implicit (short_mapkind, kinds, i);
if (n2
&& n2->tgt == n->tgt
&& n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, implicit, cbuf,
- refcount_set);
+ kind & typemask, false, cbuf, refcount_set);
return;
}
if (sizes[i] == 0)
@@ -676,8 +640,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
== n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, implicit, cbuf,
- refcount_set);
+ kind & typemask, false, cbuf, refcount_set);
return;
}
}
@@ -689,8 +652,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
&& n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, implicit, cbuf,
- refcount_set);
+ kind & typemask, false, cbuf, refcount_set);
return;
}
}
@@ -936,7 +898,6 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
for (i = 0; i < mapnum; i++)
{
int kind = get_kind (short_mapkind, kinds, i);
- bool implicit = get_implicit (short_mapkind, kinds, i);
if (hostaddrs[i] == NULL
|| (kind & typemask) == GOMP_MAP_FIRSTPRIVATE_INT)
{
@@ -1143,8 +1104,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
}
gomp_map_vars_existing (devicep, aq, n, &cur_node, &tgt->list[i],
- kind & typemask, always_to_cnt > 0, implicit,
- NULL, refcount_set);
+ kind & typemask, always_to_cnt > 0, NULL,
+ refcount_set);
i += always_to_cnt;
}
else
@@ -1221,7 +1182,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
{
assert (n->refcount != REFCOUNT_LINK);
gomp_map_vars_existing (devicep, aq, n, &cur_node, row_desc,
- kind & typemask, false, false,
+ kind & typemask, false,
/* TODO: cbuf? */ NULL, refcount_set);
}
else
@@ -1351,7 +1312,6 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
else if (tgt->list[i].key == NULL)
{
int kind = get_kind (short_mapkind, kinds, i);
- bool implicit = get_implicit (short_mapkind, kinds, i);
if (hostaddrs[i] == NULL)
continue;
switch (kind & typemask)
@@ -1523,7 +1483,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
splay_tree_key n = splay_tree_lookup (mem_map, k);
if (n && n->refcount != REFCOUNT_LINK)
gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i],
- kind & typemask, false, implicit, cbufp,
+ kind & typemask, false, cbufp,
refcount_set);
else
{
@@ -1742,7 +1702,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
{
assert (k->refcount != REFCOUNT_LINK);
gomp_map_vars_existing (devicep, aq, k, &cur_node, row_desc,
- kind & typemask, false, false,
+ kind & typemask, false,
cbufp, refcount_set);
}
else
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c
deleted file mode 100644
index f2e72936862..00000000000
--- a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c
+++ /dev/null
@@ -1,31 +0,0 @@
-#ifdef __cplusplus
-extern "C"
-#else
-extern
-#endif
-void abort (void);
-
-int
-main (void)
-{
- #define N 5
- int array[N][N];
-
- for (int i = 0; i < N; i++)
- {
- #pragma omp target enter data map(alloc: array[i:1][0:N])
-
- #pragma omp target
- for (int j = 0; j < N; j++)
- array[i][j] = i + j;
-
- #pragma omp target exit data map(from: array[i:1][0:N])
- }
-
- for (int i = 0; i < N; i++)
- for (int j = 0; j < N; j++)
- if (array[i][j] != i + j)
- abort ();
-
- return 0;
-}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-31 9:02 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-31 9:02 [gcc/devel/omp/gcc-11] Revert "OpenMP 5.0: Implement relaxation of implicit map vs. existing device mappings" Chung-Lin Tang
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).