On Mon, 10 Sep 2018 20:31:49 -0400 Julian Brown <jul...@codesourcery.com> 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 <ce...@codesourcery.com> Julian Brown <jul...@codesourcery.com> 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 <jul...@codesourcery.com> Date: Thu Sep 6 15:32:50 2018 -0700 [OpenACC] C++ reference mapping 2018-09-09 Cesar Philippidis <ce...@codesourcery.com> Julian Brown <jul...@codesourcery.com> 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>