This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)
- From: Jakub Jelinek <jakub at redhat dot com>
- To: Julian Brown <julian at codesourcery dot com>
- Cc: gcc-patches at gcc dot gnu dot org, Chung-Lin Tang <chunglin_tang at mentor dot com>, Thomas Schwinge <thomas at codesourcery dot com>, Catherine_Moore at mentor dot com
- Date: Fri, 7 Dec 2018 14:50:19 +0100
- Subject: Re: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)
- References: <aaa47f4b99ed1cf7c54400da9e255df427da6761.1541863637.git.julian@codesourcery.com> <1543578069-386-1-git-send-email-julian@codesourcery.com>
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
On Fri, Nov 30, 2018 at 03:41:09AM -0800, Julian Brown wrote:
> gcc/c-family/
> * c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ATTACH,
> PRAGMA_OACC_CLAUSE_DETACH.
...
> @@ -11804,9 +11808,12 @@ c_parser_omp_variable_list (c_parser *parser,
> case OMP_CLAUSE_MAP:
> case OMP_CLAUSE_FROM:
> case OMP_CLAUSE_TO:
> - while (c_parser_next_token_is (parser, CPP_DOT))
> + while (c_parser_next_token_is (parser, CPP_DOT)
> + || c_parser_next_token_is (parser, CPP_DEREF))
> {
> location_t op_loc = c_parser_peek_token (parser)->location;
> + if (c_parser_next_token_is (parser, CPP_DEREF))
> + t = build_simple_mem_ref (t);
This change is not ok, if OpenACC allows it in clauses, OpenMP 4.5 does not
and OpenMP 5.0 allows arbitrary lvalues that will need to be handled
differently (still unimplemented). So, this needs to be guarded for OpenACC
only (perhaps for selected OpenACC clauses)?
> @@ -12632,6 +12631,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
> }
> t = TREE_OPERAND (t, 0);
> }
> + if (TREE_CODE (t) == MEM_REF)
> + t = TREE_OPERAND (t, 0);
Again, better guard this for OpenACC. Maybe verify that mem_ref_offset is 0?
> @@ -14163,6 +14214,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
> }
> if (remove)
> break;
> + if (TREE_CODE (t) == MEM_REF)
> + t = TREE_OPERAND (t, 0);
Guard again?
> @@ -31832,15 +31836,19 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
> case OMP_CLAUSE_MAP:
> case OMP_CLAUSE_FROM:
> case OMP_CLAUSE_TO:
> - while (cp_lexer_next_token_is (parser->lexer, CPP_DOT))
> + while (cp_lexer_next_token_is (parser->lexer, CPP_DOT)
> + || cp_lexer_next_token_is (parser->lexer, CPP_DEREF))
Ditto as for C.
> @@ -4691,6 +4690,19 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
> if (low_bound == NULL_TREE)
> low_bound = integer_zero_node;
>
> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
> + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
> + {
> + if (length != integer_one_node)
> + {
> + error_at (OMP_CLAUSE_LOCATION (c),
> + OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
> + ? "array section in %<attach%> clause"
> + : "array section in %<detach%> clause");
So, are any array sections invalid, including e.g. [0:1] or say
[5:] where size of the array is 6 elts, or what exactly is invalid?
> + if (TREE_CODE (type) != POINTER_TYPE)
> + {
> + error_at (OMP_CLAUSE_LOCATION (c),
> + OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
> + ? "expected pointer in %<attach%> clause"
> + : "expected pointer in %<detach%> clause");
Perhaps you can use %qs and omp_clause_name [OMP_CLAUSE_CODE (c)] ?
> + return true;
> + }
> + }
> +
> + return false;
> +}
> +
> /* For all elements of CLAUSES, validate them vs OpenMP constraints.
> Remove any elements from the list that are invalid. */
>
> @@ -6288,7 +6337,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
> t = OMP_CLAUSE_DECL (c);
> check_dup_generic_t:
> if (t == current_class_ptr
> - && (ort != C_ORT_OMP_DECLARE_SIMD
> + && ((ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_ACC)
> || (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_LINEAR
> && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_UNIFORM)))
> {
> @@ -6352,8 +6401,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
> handle_field_decl:
> if (!remove
> && TREE_CODE (t) == FIELD_DECL
> - && t == OMP_CLAUSE_DECL (c)
> - && ort != C_ORT_ACC)
> + && t == OMP_CLAUSE_DECL (c))
> {
> OMP_CLAUSE_DECL (c)
> = omp_privatize_field (t, (OMP_CLAUSE_CODE (c)
> @@ -6420,7 +6468,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
> omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
> else
> t = OMP_CLAUSE_DECL (c);
> - if (t == current_class_ptr)
> + if (ort != C_ORT_ACC && t == current_class_ptr)
> {
> error_at (OMP_CLAUSE_LOCATION (c),
> "%<this%> allowed in OpenMP only in %<declare simd%>"
> @@ -6907,7 +6955,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
> }
> if (t == error_mark_node)
> remove = true;
> - else if (t == current_class_ptr)
> + else if (ort != C_ORT_ACC && t == current_class_ptr)
> {
> error_at (OMP_CLAUSE_LOCATION (c),
> "%<this%> allowed in OpenMP only in %<declare simd%>"
> @@ -7037,6 +7085,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
> }
> }
> }
> + if (cp_oacc_check_attachments (c))
> + remove = true;
> break;
> }
> if (t == error_mark_node)
> @@ -7044,14 +7094,25 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
> remove = true;
> break;
> }
> + /* OpenACC attach / detach clauses must be pointers. */
> + if (cp_oacc_check_attachments (c))
> + {
> + remove = true;
> + break;
> + }
> if (REFERENCE_REF_P (t)
> && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
> {
> t = TREE_OPERAND (t, 0);
> OMP_CLAUSE_DECL (c) = t;
> }
> + if (ort == C_ORT_ACC
> + && TREE_CODE (t) == COMPONENT_REF
> + && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
> + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
> if (TREE_CODE (t) == COMPONENT_REF
> - && (ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
> + && ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
> + || ort == C_ORT_ACC)
> && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
> {
> if (type_dependent_expression_p (t))
> diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
> index d8ef35d..9f96418 100644
> --- a/gcc/fortran/gfortran.h
> +++ b/gcc/fortran/gfortran.h
> @@ -1175,10 +1175,12 @@ enum gfc_omp_depend_op
> enum gfc_omp_map_op
> {
> OMP_MAP_ALLOC,
> + OMP_MAP_ATTACH,
> OMP_MAP_TO,
> OMP_MAP_FROM,
> OMP_MAP_TOFROM,
> OMP_MAP_DELETE,
> + OMP_MAP_DETACH,
> OMP_MAP_FORCE_ALLOC,
> OMP_MAP_FORCE_TO,
> OMP_MAP_FORCE_FROM,
> diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
> index 6430e61..ebba7ca 100644
> --- a/gcc/fortran/openmp.c
> +++ b/gcc/fortran/openmp.c
> @@ -222,7 +222,8 @@ static match
> gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
> bool allow_common, bool *end_colon = NULL,
> gfc_omp_namelist ***headp = NULL,
> - bool allow_sections = false)
> + bool allow_sections = false,
> + bool allow_derived = false)
> {
> gfc_omp_namelist *head, *tail, *p;
> locus old_loc, cur_loc;
> @@ -248,7 +249,8 @@ gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
> case MATCH_YES:
> gfc_expr *expr;
> expr = NULL;
> - if (allow_sections && gfc_peek_ascii_char () == '(')
> + if ((allow_sections && gfc_peek_ascii_char () == '(')
> + || (allow_derived && gfc_peek_ascii_char () == '%'))
> {
> gfc_current_locus = cur_loc;
> m = gfc_match_variable (&expr, 0);
> @@ -785,7 +787,7 @@ enum omp_mask1
> OMP_MASK1_LAST
> };
>
> -/* OpenACC 2.0 specific clauses. */
> +/* OpenACC 2.0+ specific clauses. */
> enum omp_mask2
> {
> OMP_CLAUSE_ASYNC,
> @@ -811,6 +813,8 @@ enum omp_mask2
> OMP_CLAUSE_TILE,
> OMP_CLAUSE_IF_PRESENT,
> OMP_CLAUSE_FINALIZE,
> + OMP_CLAUSE_ATTACH,
> + OMP_CLAUSE_DETACH,
> /* This must come last. */
> OMP_MASK2_LAST
> };
> @@ -914,10 +918,12 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : omp_mask (m)
> mapping. */
>
> static bool
> -gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op)
> +gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
> + bool allow_derived = false)
> {
> gfc_omp_namelist **head = NULL;
> - if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true)
> + if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true,
> + allow_derived)
> == MATCH_YES)
> {
> gfc_omp_namelist *n;
> @@ -939,6 +945,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> {
> gfc_omp_clauses *c = gfc_get_omp_clauses ();
> locus old_loc;
> + /* Determine whether we're dealing with an OpenACC directive that permits
> + derived type member accesses. This in particular disallows
> + "!$acc declare" from using such accesses, because it's not clear if/how
> + that should work. */
> + bool allow_derived = (openacc
> + && ((mask & OMP_CLAUSE_ATTACH)
> + || (mask & OMP_CLAUSE_DETACH)
> + || (mask & OMP_CLAUSE_HOST_SELF)));
>
> gcc_checking_assert (OMP_MASK1_LAST <= 64 && OMP_MASK2_LAST <= 64);
> *cp = NULL;
> @@ -1012,6 +1026,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> needs_space = true;
> continue;
> }
> + if ((mask & OMP_CLAUSE_ATTACH)
> + && gfc_match ("attach ( ") == MATCH_YES
> + && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> + OMP_MAP_ATTACH, allow_derived))
> + continue;
> break;
> case 'c':
> if ((mask & OMP_CLAUSE_COLLAPSE)
> @@ -1039,7 +1058,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> if ((mask & OMP_CLAUSE_COPY)
> && gfc_match ("copy ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_TOFROM))
> + OMP_MAP_TOFROM, allow_derived))
> continue;
> if (mask & OMP_CLAUSE_COPYIN)
> {
> @@ -1047,7 +1066,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> {
> if (gfc_match ("copyin ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_TO))
> + OMP_MAP_TO, allow_derived))
> continue;
> }
> else if (gfc_match_omp_variable_list ("copyin (",
> @@ -1058,7 +1077,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> if ((mask & OMP_CLAUSE_COPYOUT)
> && gfc_match ("copyout ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_FROM))
> + OMP_MAP_FROM, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_COPYPRIVATE)
> && gfc_match_omp_variable_list ("copyprivate (",
> @@ -1068,7 +1087,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> if ((mask & OMP_CLAUSE_CREATE)
> && gfc_match ("create ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_ALLOC))
> + OMP_MAP_ALLOC, allow_derived))
> continue;
> break;
> case 'd':
> @@ -1104,7 +1123,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> if ((mask & OMP_CLAUSE_DELETE)
> && gfc_match ("delete ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_RELEASE))
> + OMP_MAP_RELEASE, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_DEPEND)
> && gfc_match ("depend ( ") == MATCH_YES)
> @@ -1147,6 +1166,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> else
> gfc_current_locus = old_loc;
> }
> + if ((mask & OMP_CLAUSE_DETACH)
> + && gfc_match ("detach ( ") == MATCH_YES
> + && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> + OMP_MAP_DETACH, allow_derived))
> + continue;
> if ((mask & OMP_CLAUSE_DEVICE)
> && !openacc
> && c->device == NULL
> @@ -1156,12 +1180,13 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> && openacc
> && gfc_match ("device ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_FORCE_TO))
> + OMP_MAP_FORCE_TO, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_DEVICEPTR)
> && gfc_match ("deviceptr ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_FORCE_DEVICEPTR))
> + OMP_MAP_FORCE_DEVICEPTR,
> + allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_DEVICE_RESIDENT)
> && gfc_match_omp_variable_list
> @@ -1239,7 +1264,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> if ((mask & OMP_CLAUSE_HOST_SELF)
> && gfc_match ("host ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_FORCE_FROM))
> + OMP_MAP_FORCE_FROM, allow_derived))
> continue;
> break;
> case 'i':
> @@ -1511,47 +1536,48 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> if ((mask & OMP_CLAUSE_COPY)
> && gfc_match ("pcopy ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_TOFROM))
> + OMP_MAP_TOFROM, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_COPYIN)
> && gfc_match ("pcopyin ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_TO))
> + OMP_MAP_TO, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_COPYOUT)
> && gfc_match ("pcopyout ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_FROM))
> + OMP_MAP_FROM, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_CREATE)
> && gfc_match ("pcreate ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_ALLOC))
> + OMP_MAP_ALLOC, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_PRESENT)
> && gfc_match ("present ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_FORCE_PRESENT))
> + OMP_MAP_FORCE_PRESENT,
> + allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_COPY)
> && gfc_match ("present_or_copy ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_TOFROM))
> + OMP_MAP_TOFROM, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_COPYIN)
> && gfc_match ("present_or_copyin ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_TO))
> + OMP_MAP_TO, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_COPYOUT)
> && gfc_match ("present_or_copyout ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_FROM))
> + OMP_MAP_FROM, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_CREATE)
> && gfc_match ("present_or_create ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_ALLOC))
> + OMP_MAP_ALLOC, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_PRIORITY)
> && c->priority == NULL
> @@ -1669,8 +1695,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
>
> if (gfc_match_omp_variable_list (" :",
> &c->lists[OMP_LIST_REDUCTION],
> - false, NULL, &head,
> - openacc) == MATCH_YES)
> + false, NULL, &head, openacc,
> + allow_derived) == MATCH_YES)
> {
> gfc_omp_namelist *n;
> if (rop == OMP_REDUCTION_NONE)
> @@ -1769,7 +1795,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> if ((mask & OMP_CLAUSE_HOST_SELF)
> && gfc_match ("self ( ") == MATCH_YES
> && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> - OMP_MAP_FORCE_FROM))
> + OMP_MAP_FORCE_FROM, allow_derived))
> continue;
> if ((mask & OMP_CLAUSE_SEQ)
> && !c->seq
> @@ -1927,17 +1953,17 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
> | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
> | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
> - | OMP_CLAUSE_WAIT)
> + | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
> #define OACC_KERNELS_CLAUSES \
> (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
> | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
> | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
> | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \
> - | OMP_CLAUSE_WAIT)
> + | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
> #define OACC_DATA_CLAUSES \
> (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \
> | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \
> - | OMP_CLAUSE_PRESENT)
> + | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH)
> #define OACC_LOOP_CLAUSES \
> (omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \
> | OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \
> @@ -1958,10 +1984,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
> | OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_IF_PRESENT)
> #define OACC_ENTER_DATA_CLAUSES \
> (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
> - | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE)
> + | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE | OMP_CLAUSE_ATTACH)
> #define OACC_EXIT_DATA_CLAUSES \
> (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
> - | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE | OMP_CLAUSE_FINALIZE)
> + | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE | OMP_CLAUSE_FINALIZE \
> + | OMP_CLAUSE_DETACH)
> #define OACC_WAIT_CLAUSES \
> omp_mask (OMP_CLAUSE_ASYNC)
> #define OACC_ROUTINE_CLAUSES \
> @@ -3734,9 +3761,6 @@ resolve_nonnegative_int_expr (gfc_expr *expr, const char *clause)
> static void
> check_symbol_not_pointer (gfc_symbol *sym, locus loc, const char *name)
> {
> - if (sym->ts.type == BT_DERIVED && sym->attr.pointer)
> - gfc_error ("POINTER object %qs of derived type in %s clause at %L",
> - sym->name, name, &loc);
> if (sym->ts.type == BT_DERIVED && sym->attr.cray_pointer)
> gfc_error ("Cray pointer object %qs of derived type in %s clause at %L",
> sym->name, name, &loc);
> @@ -3781,9 +3805,6 @@ check_array_not_assumed (gfc_symbol *sym, locus loc, const char *name)
> static void
> resolve_oacc_data_clauses (gfc_symbol *sym, locus loc, const char *name)
> {
> - if (sym->ts.type == BT_DERIVED && sym->attr.allocatable)
> - gfc_error ("ALLOCATABLE object %qs of derived type in %s clause at %L",
> - sym->name, name, &loc);
> if ((sym->ts.type == BT_ASSUMED && sym->attr.allocatable)
> || (sym->ts.type == BT_CLASS && CLASS_DATA (sym)
> && CLASS_DATA (sym)->attr.allocatable))
> @@ -4153,11 +4174,23 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
> && (list != OMP_LIST_REDUCTION || !openacc))
> for (n = omp_clauses->lists[list]; n; n = n->next)
> {
> - if (n->sym->mark)
> - gfc_error ("Symbol %qs present on multiple clauses at %L",
> - n->sym->name, &n->where);
> - else
> - n->sym->mark = 1;
> + bool array_only_p = true;
> + /* Disallow duplicate bare variable references and multiple
> + subarrays of the same array here, but allow multiple components of
> + the same (e.g. derived-type) variable. For the latter, duplicate
> + components are detected elsewhere. */
> + if (openacc && n->expr && n->expr->expr_type == EXPR_VARIABLE)
> + for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
> + if (ref->type != REF_ARRAY)
> + array_only_p = false;
> + if (array_only_p)
> + {
> + if (n->sym->mark)
> + gfc_error ("Symbol %qs present on multiple clauses at %L",
> + n->sym->name, &n->where);
> + else
> + n->sym->mark = 1;
> + }
> }
>
> gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1);
> @@ -4348,23 +4381,41 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
> "are allowed on ORDERED directive at %L",
> &n->where);
> }
> + gfc_ref *array_ref = NULL;
> + bool resolved = false;
> if (n->expr)
> {
> - if (!gfc_resolve_expr (n->expr)
> + array_ref = n->expr->ref;
> + resolved = gfc_resolve_expr (n->expr);
> +
> + /* Look through component refs to find last array
> + reference. */
> + while (resolved
> + && array_ref
> + && (array_ref->type == REF_COMPONENT
> + || (array_ref->type == REF_ARRAY
> + && array_ref->next
> + && array_ref->next->type == REF_COMPONENT)))
> + array_ref = array_ref->next;
I'd guard this stuff for OpenACC only, keep what it did for OpenMP.
> + }
> + if (array_ref
> + || (n->expr
> + && (!resolved || n->expr->expr_type != EXPR_VARIABLE)))
> + {
> + if (!resolved
> || n->expr->expr_type != EXPR_VARIABLE
> - || n->expr->ref == NULL
> - || n->expr->ref->next
> - || n->expr->ref->type != REF_ARRAY)
> + || array_ref->next
> + || array_ref->type != REF_ARRAY)
> gfc_error ("%qs in %s clause at %L is not a proper "
> "array section", n->sym->name, name,
> &n->where);
> - else if (n->expr->ref->u.ar.codimen)
> + else if (array_ref->u.ar.codimen)
> gfc_error ("Coarrays not supported in %s clause at %L",
> name, &n->where);
> else
> {
> int i;
> - gfc_array_ref *ar = &n->expr->ref->u.ar;
> + gfc_array_ref *ar = &array_ref->u.ar;
> for (i = 0; i < ar->dimen; i++)
> if (ar->stride[i])
> {
> + /* For OpenACC, pointers in structs should trigger an
> + attach action. */
> + if (ptr && (region_type & ORT_ACC) != 0)
> + {
> + /* Turning a GOMP_MAP_ALWAYS_POINTER clause into a
> + GOMP_MAP_ATTACH clause after we have detected a case
> + that needs a GOMP_MAP_STRUCT mapping adding. */
> + OMP_CLAUSE_SET_MAP_KIND (c,
> + (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
> + : GOMP_MAP_ATTACH);
Bad formatting, I'd suggest use a temporary with gomp_map_kind type.
> + has_attachments = true;
> + }
> if (n == NULL || (n->value & GOVD_MAP) == 0)
> {
> tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
> OMP_CLAUSE_MAP);
> - OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT);
> + OMP_CLAUSE_SET_MAP_KIND (l, attach
> + ? GOMP_MAP_FORCE_PRESENT : GOMP_MAP_STRUCT);
Likewise.
> if (!base_eq_orig_base)
> OMP_CLAUSE_DECL (l) = unshare_expr (orig_base);
> else
> OMP_CLAUSE_DECL (l) = decl;
> - OMP_CLAUSE_SIZE (l) = size_int (1);
> + OMP_CLAUSE_SIZE (l) = attach
> + ? (DECL_P (OMP_CLAUSE_DECL (l))
> + ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
> + : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))))
> + : size_int (1);
Again, bad formatting. = attach should be on a next line, the indentation
is also weird, best like:
OMP_CLAUSE_SIZE (l)
= (!attach
? size_int (1)
: DECL_P (OMP_CLAUSE_DECL (l))
? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
: TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))));
> if (struct_map_to_clause == NULL)
> struct_map_to_clause = new hash_map<tree, tree>;
> struct_map_to_clause->put (decl, l);
> @@ -8681,9 +8713,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
> flags = GOVD_MAP | GOVD_EXPLICIT;
> if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
> flags |= GOVD_SEEN;
> + if (has_attachments)
> + flags |= GOVD_MAP_HAS_ATTACHMENTS;
> goto do_add_decl;
> }
> - else
> + else if (struct_map_to_clause)
> {
> tree *osc = struct_map_to_clause->get (decl);
> tree *sc = NULL, *scp = NULL;
> @@ -8692,8 +8726,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
> sc = &OMP_CLAUSE_CHAIN (*osc);
> if (*sc != c
> && (OMP_CLAUSE_MAP_KIND (*sc)
> - == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
> + == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
> sc = &OMP_CLAUSE_CHAIN (*sc);
> + /* Here "prev_list_p" is the end of the inserted
> + alloc/release nodes after the struct node, OSC. */
> for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
> if (ptr && sc == prev_list_p)
> break;
> @@ -8752,9 +8788,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
> }
> if (remove)
> break;
> - OMP_CLAUSE_SIZE (*osc)
> - = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
> - size_one_node);
> + if (!attach)
> + OMP_CLAUSE_SIZE (*osc)
> + = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
> + size_one_node);
> if (ptr)
> {
> tree cl
> @@ -8786,11 +8823,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
> }
> if (!remove
> && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
> + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
> && OMP_CLAUSE_CHAIN (c)
> && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
> - && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
> - == GOMP_MAP_ALWAYS_POINTER))
> + && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
> + == GOMP_MAP_ALWAYS_POINTER)
> + || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
> + == GOMP_MAP_TO_PSET)))
> prev_list_p = list_p;
> +
> break;
> }
> flags = GOVD_MAP | GOVD_EXPLICIT;
> @@ -9412,6 +9453,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
> return 0;
> if ((flags & GOVD_SEEN) == 0)
> return 0;
> + if ((flags & GOVD_MAP_HAS_ATTACHMENTS) != 0)
> + return 0;
> if (flags & GOVD_DEBUG_PRIVATE)
> {
> gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_SHARED);
> @@ -11795,8 +11838,9 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
> && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
> OMP_CLAUSE_FINALIZE))
> {
> - /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
> - semantics apply to all mappings of this OpenACC directive. */
> + /* Use GOMP_MAP_DELETE, GOMP_MAP_FORCE_DETACH, and
> + GOMP_MAP_FORCE_FROM to denote that "finalize" semantics apply
> + to all mappings of this OpenACC directive. */
> bool finalize_marked = false;
> for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
> if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
> @@ -11810,10 +11854,19 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
> OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
> finalize_marked = true;
> break;
> + case GOMP_MAP_DETACH:
> + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_DETACH);
> + finalize_marked = true;
> + break;
> + case GOMP_MAP_STRUCT:
> + case GOMP_MAP_FORCE_PRESENT:
> + /* Skip over an initial struct or force_present mapping. */
> + break;
> default:
> - /* Check consistency: libgomp relies on the very first data
> - mapping clause being marked, so make sure we did that before
> - any other mapping clauses. */
> + /* Check consistency: libgomp relies on the very first
> + non-struct, non-force-present data mapping clause being
> + marked, so make sure we did that before any other mapping
> + clauses. */
> gcc_assert (finalize_marked);
> break;
> }
> diff --git a/gcc/omp-low.c b/gcc/omp-low.c
> index ca78d7a..55dbc0b 100644
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -9138,6 +9138,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
> case GOMP_MAP_FORCE_DEVICEPTR:
> case GOMP_MAP_DEVICE_RESIDENT:
> case GOMP_MAP_LINK:
> + case GOMP_MAP_ATTACH:
> + case GOMP_MAP_DETACH:
> + case GOMP_MAP_FORCE_DETACH:
> gcc_assert (is_gimple_omp_oacc (stmt));
> break;
> default:
> diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
> new file mode 100644
> index 0000000..84a44af
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
> @@ -0,0 +1,54 @@
> +/* Test OpenACC's support for manual deep copy, including the attach
> + and detach clauses. */
> +
> +/* { dg-additional-options "-fdump-tree-omplower" } */
> +
> +void
> +t1 ()
> +{
> + struct foo {
> + int *a, *b, c, d, *e;
> + } s;
> +
> + int *a, *z;
> +
> +#pragma acc enter data copyin(s)
> + {
> +#pragma acc data copy(s.a[0:10]) copy(z[0:10])
> + {
> + s.e = z;
> +#pragma acc parallel loop attach(s.e)
> + for (int i = 0; i < 10; i++)
> + s.a[i] = s.e[i];
> +
> +
> + a = s.e;
> +#pragma acc enter data attach(a)
> +#pragma acc exit data detach(a)
> + }
> +
> +#pragma acc enter data copyin(a)
> +#pragma acc acc enter data attach(s.e)
> +#pragma acc exit data detach(s.e)
> +
> +#pragma acc data attach(s.e)
> + {
> + }
> +#pragma acc exit data delete(a)
> +
> +#pragma acc exit data detach(a) finalize
> +#pragma acc exit data detach(s.a) finalize
> + }
> +}
> +
> +/* { 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 .len: 0.." 1 "omplower" } } */
> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.force_present:s .len: 32.. map.attach:s.e .len: 8.." 1 "omplower" } } */
> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .len: 8.." 1 "omplower" } } */
> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .len: 8.." 1 "omplower" } } */
> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.force_present:s .len: 32.. map.detach:s.e .len: 8.." 1 "omplower" } } */
> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.force_present:s .len: 32.. map.attach:s.e .len: 8.." 1 "omplower" } } */
> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */
> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .len: 8.." 1 "omplower" } } */
> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_present:s .len: 32.. map.force_detach:s.a .len: 8.." 1 "omplower" } } */
Aren't the lengths here heavily dependent on the target? E.g. if it depends
on sizeof (int) == 4, maybe the whole test needs to be guarded with { target int32 }
> @@ -918,8 +920,13 @@ struct splay_tree_key_s {
> uintptr_t tgt_offset;
> /* Reference count. */
> uintptr_t refcount;
> - /* Dynamic reference count. */
> - uintptr_t dynamic_refcount;
> + /* Reference counts beyond those that represent genuine references in the
> + linked splay tree key/target memory structures, e.g. for multiple OpenACC
> + "present increment" operations (via "acc enter data") refering to the same
> + host-memory block. */
> + uintptr_t virtual_refcount;
> + /* For a block with attached pointers, the attachment counters for each. */
> + unsigned short *attach_count;
> /* Pointer to the original mapping of "omp declare target link" object. */
> splay_tree_key link_key;
> };
This is something I'm worried about a lot, the nodes keep growing way too
much. Is there a way to reuse some other field if it is of certain kind?
Also, why unsigned short, can you only attach 65535 times?
Jakub