From: Jakub Jelinek <jakub@redhat.com>
To: Julian Brown <julian@codesourcery.com>
Cc: gcc-patches@gcc.gnu.org,
Chung-Lin Tang <chunglin_tang@mentor.com>,
Thomas Schwinge <thomas@codesourcery.com>,
Catherine_Moore@mentor.com
Subject: Re: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)
Date: Fri, 07 Dec 2018 13:50:00 -0000 [thread overview]
Message-ID: <20181207135019.GI12380@tucnak> (raw)
In-Reply-To: <1543578069-386-1-git-send-email-julian@codesourcery.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
next prev parent reply other threads:[~2018-12-07 13:50 UTC|newest]
Thread overview: 81+ messages / expand[flat|nested] mbox.gz Atom feed top
2018-11-10 17:11 [PATCH 0/3] " Julian Brown
2018-11-10 17:11 ` [PATCH 1/3] Host-to-device transfer coalescing & magic offset value self-documentation Julian Brown
2018-12-21 10:56 ` libgomp/target.c magic constants self-documentation Thomas Schwinge
2019-05-29 14:48 ` Thomas Schwinge
2018-11-10 17:11 ` [PATCH 2/3] Factor out duplicate code in gimplify_scan_omp_clauses Julian Brown
2018-12-18 14:16 ` Julian Brown
2018-12-18 14:50 ` Jakub Jelinek
2018-11-10 17:12 ` [PATCH 3/3] OpenACC 2.6 manual deep copy support (attach/detach) Julian Brown
2018-11-11 17:04 ` Bernhard Reutner-Fischer
2018-11-30 11:41 ` [PATCH] " Julian Brown
2018-12-03 17:03 ` Julian Brown
2018-12-07 13:50 ` Jakub Jelinek [this message]
2018-12-10 19:42 ` Julian Brown
2018-12-13 10:57 ` Jakub Jelinek
2018-12-14 19:00 ` Julian Brown
2018-12-18 12:25 ` Jakub Jelinek
2018-12-22 13:37 ` Thomas Schwinge
2019-10-18 17:20 ` Thomas Schwinge
2019-11-06 18:44 ` Julian Brown
2019-11-22 23:54 ` Julian Brown
2019-11-25 10:53 ` Tobias Burnus
2019-11-26 2:54 ` Julian Brown
2019-12-17 12:16 ` Thomas Schwinge
2019-12-17 17:28 ` [WIP] OpenACC 'acc_attach*', 'acc_detach*' runtime library routines (was: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)) Thomas Schwinge
2019-12-18 6:03 ` [PATCH 00/13] OpenACC 2.6 manual deep copy support Julian Brown
2019-12-18 6:03 ` [PATCH 01/13] Use aux struct in libgomp for infrequently-used/API-specific data Julian Brown
2019-12-18 6:03 ` [PATCH 02/13] OpenACC reference count overhaul Julian Brown
2020-05-19 15:42 ` Thomas Schwinge
2020-06-04 18:13 ` [OpenACC] Use 'tgt' returned from 'gomp_map_vars' (was: [PATCH 02/13] OpenACC reference count overhaul) Thomas Schwinge
2020-05-19 15:49 ` [PATCH 02/13] OpenACC reference count overhaul Thomas Schwinge
2020-05-19 15:58 ` Thomas Schwinge
2020-06-25 11:03 ` Thomas Schwinge
2020-07-03 15:29 ` Thomas Schwinge
2019-12-18 6:03 ` [PATCH 03/13] OpenACC reference count consistency checking Julian Brown
2019-12-18 6:04 ` [PATCH 08/13] OpenACC 2.6 deep copy: middle-end parts Julian Brown
2019-12-21 21:51 ` Thomas Schwinge
2019-12-18 6:04 ` [PATCH 09/13] OpenACC 2.6 deep copy: C and C++ front-end parts Julian Brown
2019-12-24 5:05 ` Thomas Schwinge
2019-12-26 19:04 ` Jason Merrill
2021-06-10 11:03 ` Thomas Schwinge
2019-12-18 6:04 ` [PATCH 06/13] OpenACC 2.6 deep copy: attach/detach API routines Julian Brown
2019-12-18 6:04 ` [PATCH 05/13] Factor out duplicate code in gimplify_scan_omp_clauses Julian Brown
2019-12-18 6:04 ` [PATCH 04/13] Use gomp_map_val for OpenACC host-to-device address translation Julian Brown
2019-12-18 6:05 ` [PATCH 13/13] Fortran polymorphic class-type support for OpenACC Julian Brown
2019-12-18 6:05 ` [PATCH 11/13] OpenACC 2.6 deep copy: C and C++ execution tests Julian Brown
2020-06-04 18:43 ` Fix 'sizeof' usage in 'libgomp.oacc-c-c++-common/deep-copy-{7, 8}.c' (was: [PATCH 11/13] OpenACC 2.6 deep copy: C and C++ execution tests) Thomas Schwinge
2023-10-31 14:00 ` Add OpenACC 'acc_map_data' variant to 'libgomp.oacc-c-c++-common/deep-copy-8.c' " Thomas Schwinge
2019-12-18 6:05 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Julian Brown
2019-12-21 23:37 ` Thomas Schwinge
2020-01-03 12:26 ` Julian Brown
2020-05-20 9:37 ` Thomas Schwinge
2020-06-05 16:23 ` [OpenACC 'exit data'] Simplify 'GOMP_MAP_STRUCT' handling (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) Thomas Schwinge
2020-06-05 16:36 ` [OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings " Thomas Schwinge
2020-05-20 14:52 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Thomas Schwinge
2020-05-20 19:11 ` Julian Brown
2020-06-04 18:35 ` [OpenACC] Repair/restore 'is_tgt_unmapped' checking (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) Thomas Schwinge
2020-06-04 18:53 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Thomas Schwinge
2020-06-05 10:39 ` Thomas Schwinge
2020-06-05 20:28 ` Julian Brown
2020-06-05 11:17 ` Thomas Schwinge
2020-06-05 20:31 ` Julian Brown
2020-06-09 10:41 ` OpenACC 'attach'/'detach' has no business affecting user-visible reference counting (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts) Thomas Schwinge
2020-06-09 12:23 ` Julian Brown
2020-06-18 18:21 ` Julian Brown
2020-07-16 8:35 ` OpenACC 'attach'/'detach' has no business affecting user-visible reference counting Thomas Schwinge
2020-06-26 9:20 ` [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts Thomas Schwinge
2020-07-16 9:35 ` Thomas Schwinge
2020-07-16 21:21 ` Julian Brown
2020-07-17 9:12 ` Thomas Schwinge
2020-06-30 15:58 ` Thomas Schwinge
2019-12-18 6:05 ` [PATCH 12/13] OpenACC 2.6 deep copy: Fortran execution tests Julian Brown
2019-12-18 7:20 ` [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts Julian Brown
2019-12-18 23:30 ` Tobias Burnus
2019-12-20 12:25 ` [committed] Improve is-coindexed check for OpenACC/OpenMP (was: [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts) Tobias Burnus
2019-12-20 13:25 ` [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts Tobias Burnus
2019-12-20 10:08 ` [patch,committed] Fix testsuite-fallout of OpenACC deep-copy patch (was: [PATCH 10/13] OpenACC 2.6 deep copy: Fortran front-end parts) Tobias Burnus
2019-12-18 18:24 ` [PATCH 00/13] OpenACC 2.6 manual deep copy support Thomas Schwinge
2019-12-20 1:21 ` Julian Brown
2019-12-20 14:36 ` OpenACC regression and development pace Thomas Koenig
2020-06-04 18:07 ` [OpenACC] XFAIL behavior of over-eager 'finalize' clause (was: [PATCH 00/13] OpenACC 2.6 manual deep copy support) Thomas Schwinge
2019-12-17 16:53 ` In 'libgomp/target.c', 'struct splay_tree_key_s', use 'struct splay_tree_aux' for infrequently-used or API-specific data (was: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)) Thomas Schwinge
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=20181207135019.GI12380@tucnak \
--to=jakub@redhat.com \
--cc=Catherine_Moore@mentor.com \
--cc=chunglin_tang@mentor.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=julian@codesourcery.com \
--cc=thomas@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
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).