On Mon, 10 Sep 2018 20:31:49 -0400
Julian Brown <[email protected]> wrote:
> [...] I think the handling of references can and should match between
> the two APIs (though implementation details of the patch to make that
> so need a little work still).
Here's a new version of the patch, somewhat simplified and slightly more
obviously making the treatment of references between OpenMP and OpenACC
the same. I worried a little about the potential side-effects of making
ctx->target_firstprivatize_array_bases true for parallel and kernels
regions, but test results revealed no problems with doing that and I
think generated code may even be a little better (and more consistent)
in some cases.
For example, one case that is handled differently now is as follows:
#include <stdlib.h>
__attribute__((noinline)) int
bar (int c)
{
int arr[c];
#pragma acc parallel loop copy(arr)
for (int i = 0; i < c; i++)
arr[i] = i;
for (int i = 0; i < c; i++)
if (arr[i] != i)
abort ();
return arr[c - 1];
}
int main (int argc, char *argv[])
{
return bar (100);
}
The VLA was previously mapped as:
#pragma omp target oacc_parallel map(tofrom:*arr.1 [len: D.2607]) \
map(alloc:arr [pointer assign, bias: 0]) firstprivate(c)
and is now mapped as:
#pragma omp target oacc_parallel map(tofrom:*arr.1 [len: D.2607]) \
map(firstprivate:arr [pointer assign, bias: 0]) firstprivate(c)
Either works, but IIUC using firstprivate_pointer can be more efficient
if the pointer is dereferenced multiple times in a kernel, since a local
copy of the incoming mapped pointer is made per-thread/workitem.
Generally, array sections are already using firstprivate pointers for
their bases with OpenACC.
Re-tested with offloading to NVPTX and bootstrapped. OK, or any other
comments?
Thanks,
Julian
ChangeLog
2018-09-09 Cesar Philippidis <[email protected]>
Julian Brown <[email protected]>
PR middle-end/86336
gcc/cp/
* semantics.c (finish_omp_clauses): Treat C++ references the same in
OpenACC as OpenMP.
* gimplify.c (gimplify_scan_omp_clauses): Set
target_firstprivatize_array_bases in OpenACC parallel and kernels
region contexts. Remove GOMP_MAP_FIRSTPRIVATE_REFERENCE clauses from
OpenACC data regions.
libgomp/
* testsuite/libgomp.oacc-c++/non-scalar-data.C: Remove XFAIL.
commit 6f3d5b86b4413722c3e7ab3ca9a678d7c35b68fe
Author: Julian Brown <[email protected]>
Date: Thu Sep 6 15:32:50 2018 -0700
[OpenACC] C++ reference mapping
2018-09-09 Cesar Philippidis <[email protected]>
Julian Brown <[email protected]>
PR middle-end/86336
gcc/cp/
* semantics.c (finish_omp_clauses): Treat C++ references the same in
OpenACC as OpenMP.
* gimplify.c (gimplify_scan_omp_clauses): Set
target_firstprivatize_array_bases in OpenACC parallel and kernels
region contexts. Remove GOMP_MAP_FIRSTPRIVATE_REFERENCE clauses from
OpenACC data regions.
libgomp/
* testsuite/libgomp.oacc-c++/non-scalar-data.C: Remove XFAIL.
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index f3e5d83..bf3c63a 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6878,7 +6878,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
handle_map_references:
if (!remove
&& !processing_template_decl
- && (ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
+ && ort != C_ORT_DECLARE_SIMD
&& TYPE_REF_P (TREE_TYPE (OMP_CLAUSE_DECL (c))))
{
t = OMP_CLAUSE_DECL (c);
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index dbd0f0e..f0eb04a 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7513,6 +7513,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_EXIT_DATA:
case OACC_DECLARE:
case OACC_HOST_DATA:
+ case OACC_PARALLEL:
+ case OACC_KERNELS:
ctx->target_firstprivatize_array_bases = true;
default:
break;
@@ -8556,7 +8558,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if (code == OACC_DATA
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
remove = true;
if (remove)
*list_p = OMP_CLAUSE_CHAIN (c);
diff --git a/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
index 8e4b296..e5f8707 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/non-scalar-data.C
@@ -1,8 +1,7 @@
// Ensure that a non-scalar dummy arguments which are implicitly used inside
// offloaded regions are properly mapped using present_or_copy semantics.
-// { dg-xfail-if "TODO" { *-*-* } }
-// { dg-excess-errors "ICE" }
+// { dg-do run }
#include <cassert>
