diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 53e99aa29d9..00f8bf4376e 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -15627,7 +15627,11 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, OpenACC 2.6: no_create ( variable-list ) attach ( variable-list ) - detach ( variable-list ) */ + detach ( variable-list ) + + OpenACC 2.7: + copyin (readonly : variable-list ) + */ static tree c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, @@ -15680,11 +15684,37 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, default: gcc_unreachable (); } - tree nl, c; - nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, false); - for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + tree nl = list; + bool readonly = false; + location_t open_loc = c_parser_peek_token (parser)->location; + matching_parens parens; + if (parens.require_open (parser)) + { + /* Turn on readonly modifier parsing for copyin clause. */ + if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN) + { + c_token *token = c_parser_peek_token (parser); + if (token->type == CPP_NAME + && !strcmp (IDENTIFIER_POINTER (token->value), "readonly") + && c_parser_peek_2nd_token (parser)->type == CPP_COLON) + { + c_parser_consume_token (parser); + c_parser_consume_token (parser); + readonly = true; + } + } + nl = c_parser_omp_variable_list (parser, open_loc, OMP_CLAUSE_MAP, list, + false); + parens.skip_until_found_close (parser); + } + + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + if (readonly) + OMP_CLAUSE_MAP_READONLY (c) = 1; + } return nl; } @@ -19821,15 +19851,39 @@ c_parser_omp_structured_block (c_parser *parser, bool *if_p) /* OpenACC 2.0: # pragma acc cache (variable-list) new-line + OpenACC 2.7: + # pragma acc cache (readonly: variable-list) new-line + LOC is the location of the #pragma token. */ static tree c_parser_oacc_cache (location_t loc, c_parser *parser) { - tree stmt, clauses; + tree stmt, clauses = NULL_TREE; + bool readonly = false; + location_t open_loc = c_parser_peek_token (parser)->location; + matching_parens parens; + if (parens.require_open (parser)) + { + c_token *token = c_parser_peek_token (parser); + if (token->type == CPP_NAME + && !strcmp (IDENTIFIER_POINTER (token->value), "readonly") + && c_parser_peek_2nd_token (parser)->type == CPP_COLON) + { + c_parser_consume_token (parser); + c_parser_consume_token (parser); + readonly = true; + } + clauses = c_parser_omp_variable_list (parser, open_loc, + OMP_CLAUSE__CACHE_, NULL_TREE); + parens.skip_until_found_close (parser); + } + + if (readonly) + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE__CACHE__READONLY (c) = 1; - clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL); clauses = c_finish_omp_clauses (clauses, C_ORT_ACC); c_parser_skip_to_pragma_eol (parser); diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index e32acfc30a2..4fe27fb07b2 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -38544,7 +38544,11 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list, OpenACC 2.6: no_create ( variable-list ) attach ( variable-list ) - detach ( variable-list ) */ + detach ( variable-list ) + + OpenACC 2.7: + copyin (readonly : variable-list ) + */ static tree cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, @@ -38597,11 +38601,34 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, default: gcc_unreachable (); } - tree nl, c; - nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, false); - for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + tree nl = list; + bool readonly = false; + if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + { + /* Turn on readonly modifier parsing for copyin clause. */ + if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN) + { + cp_token *token = cp_lexer_peek_token (parser->lexer); + if (token->type == CPP_NAME + && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly") + && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON) + { + cp_lexer_consume_token (parser->lexer); + cp_lexer_consume_token (parser->lexer); + readonly = true; + } + } + nl = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, NULL, + false); + } + + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + if (readonly) + OMP_CLAUSE_MAP_READONLY (c) = 1; + } return nl; } @@ -47178,6 +47205,9 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, /* OpenACC 2.0: # pragma acc cache (variable-list) new-line + + OpenACC 2.7: + # pragma acc cache (readonly: variable-list) new-line */ static tree @@ -47187,9 +47217,28 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok) clauses. */ auto_suppress_location_wrappers sentinel; - tree stmt, clauses; + tree stmt, clauses = NULL_TREE; + bool readonly = false; + + if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + { + cp_token *token = cp_lexer_peek_token (parser->lexer); + if (token->type == CPP_NAME + && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly") + && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON) + { + cp_lexer_consume_token (parser->lexer); + cp_lexer_consume_token (parser->lexer); + readonly = true; + } + clauses = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE__CACHE_, + NULL, NULL); + } + + if (readonly) + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE__CACHE__READONLY (c) = 1; - clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE); clauses = finish_omp_clauses (clauses, C_ORT_ACC); cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer)); diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 7b154eb3ca7..db84b06289b 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1400,6 +1400,9 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) fputs (") ALLOCATE(", dumpfile); continue; } + if ((list_type == OMP_LIST_MAP || list_type == OMP_LIST_CACHE) + && n->u.map.readonly) + fputs ("readonly,", dumpfile); if (list_type == OMP_LIST_REDUCTION) switch (n->u.reduction_op) { @@ -1467,7 +1470,7 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) default: break; } else if (list_type == OMP_LIST_MAP) - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_ALLOC: fputs ("alloc:", dumpfile); break; case OMP_MAP_TO: fputs ("to:", dumpfile); break; diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index ebba2336e12..32b792f85fb 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1363,7 +1363,11 @@ typedef struct gfc_omp_namelist { gfc_omp_reduction_op reduction_op; gfc_omp_depend_doacross_op depend_doacross_op; - gfc_omp_map_op map_op; + struct + { + ENUM_BITFIELD (gfc_omp_map_op) op:8; + bool readonly; + } map; gfc_expr *align; struct { diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 38de60238c0..5c44e666eb9 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -1210,7 +1210,7 @@ gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op, { gfc_omp_namelist *n; for (n = *head; n; n = n->next) - n->u.map_op = map_op; + n->u.map.op = map_op; return true; } @@ -1524,7 +1524,7 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc, gfc_omp_namelist *p = gfc_get_omp_namelist (), **tl; p->sym = n->sym; p->where = p->where; - p->u.map_op = OMP_MAP_ALWAYS_TOFROM; + p->u.map.op = OMP_MAP_ALWAYS_TOFROM; tl = &c->lists[OMP_LIST_MAP]; while (*tl) @@ -2181,11 +2181,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, { if (openacc) { - if (gfc_match ("copyin ( ") == MATCH_YES - && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], - OMP_MAP_TO, true, - allow_derived)) - continue; + if (gfc_match ("copyin ( ") == MATCH_YES) + { + bool readonly = gfc_match ("readonly : ") == MATCH_YES; + head = NULL; + if (gfc_match_omp_variable_list ("", + &c->lists[OMP_LIST_MAP], + true, NULL, &head, true, + allow_derived) + == MATCH_YES) + { + gfc_omp_namelist *n; + for (n = *head; n; n = n->next) + { + n->u.map.op = OMP_MAP_TO; + n->u.map.readonly = readonly; + } + continue; + } + } } else if (gfc_match_omp_variable_list ("copyin (", &c->lists[OMP_LIST_COPYIN], @@ -3134,7 +3148,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, { gfc_omp_namelist *n; for (n = *head; n; n = n->next) - n->u.map_op = map_op; + n->u.map.op = map_op; continue; } gfc_current_locus = old_loc; @@ -4002,7 +4016,7 @@ gfc_match_oacc_declare (void) if (gfc_current_ns->proc_name && gfc_current_ns->proc_name->attr.flavor == FL_MODULE) { - if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO) + if (n->u.map.op != OMP_MAP_ALLOC && n->u.map.op != OMP_MAP_TO) { gfc_error ("Invalid clause in module with !$ACC DECLARE at %L", &where); @@ -4036,7 +4050,7 @@ gfc_match_oacc_declare (void) return MATCH_ERROR; } - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_FORCE_ALLOC: case OMP_MAP_ALLOC: @@ -4151,21 +4165,36 @@ gfc_match_oacc_wait (void) match gfc_match_oacc_cache (void) { + bool readonly = false; gfc_omp_clauses *c = gfc_get_omp_clauses (); /* The OpenACC cache directive explicitly only allows "array elements or subarrays", which we're currently not checking here. Either check this after the call of gfc_match_omp_variable_list, or add something like a only_sections variant next to its allow_sections parameter. */ - match m = gfc_match_omp_variable_list (" (", - &c->lists[OMP_LIST_CACHE], true, - NULL, NULL, true); + match m = gfc_match (" ( "); if (m != MATCH_YES) { gfc_free_omp_clauses(c); return m; } - if (gfc_current_state() != COMP_DO + if (gfc_match ("readonly : ") == MATCH_YES) + readonly = true; + + gfc_omp_namelist **head = NULL; + m = gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_CACHE], true, + NULL, &head, true); + if (m != MATCH_YES) + { + gfc_free_omp_clauses(c); + return m; + } + + if (readonly) + for (gfc_omp_namelist *n = *head; n; n = n->next) + n->u.map.readonly = true; + + if (gfc_current_state() != COMP_DO && gfc_current_state() != COMP_DO_CONCURRENT) { gfc_error ("ACC CACHE directive must be inside of loop %C"); @@ -8436,8 +8465,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, } if (openacc && list == OMP_LIST_MAP - && (n->u.map_op == OMP_MAP_ATTACH - || n->u.map_op == OMP_MAP_DETACH)) + && (n->u.map.op == OMP_MAP_ATTACH + || n->u.map.op == OMP_MAP_DETACH)) { symbol_attribute attr; if (n->expr) @@ -8447,7 +8476,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, if (!attr.pointer && !attr.allocatable) gfc_error ("%qs clause argument must be ALLOCATABLE or " "a POINTER at %L", - (n->u.map_op == OMP_MAP_ATTACH) ? "attach" + (n->u.map.op == OMP_MAP_ATTACH) ? "attach" : "detach", &n->where); } if (lastref @@ -8518,7 +8547,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, else if (openacc) { if (list == OMP_LIST_MAP - && n->u.map_op == OMP_MAP_FORCE_DEVICEPTR) + && n->u.map.op == OMP_MAP_FORCE_DEVICEPTR) resolve_oacc_deviceptr_clause (n->sym, n->where, name); else resolve_oacc_data_clauses (n->sym, n->where, name); @@ -8540,7 +8569,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, { case EXEC_OMP_TARGET: case EXEC_OMP_TARGET_DATA: - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_TO: case OMP_MAP_ALWAYS_TO: @@ -8567,7 +8596,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, } break; case EXEC_OMP_TARGET_ENTER_DATA: - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_TO: case OMP_MAP_ALWAYS_TO: @@ -8577,16 +8606,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, case OMP_MAP_PRESENT_ALLOC: break; case OMP_MAP_TOFROM: - n->u.map_op = OMP_MAP_TO; + n->u.map.op = OMP_MAP_TO; break; case OMP_MAP_ALWAYS_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_TO; + n->u.map.op = OMP_MAP_ALWAYS_TO; break; case OMP_MAP_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_PRESENT_TO; + n->u.map.op = OMP_MAP_PRESENT_TO; break; case OMP_MAP_ALWAYS_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_PRESENT_TO; + n->u.map.op = OMP_MAP_ALWAYS_PRESENT_TO; break; default: gfc_error ("TARGET ENTER DATA with map-type other " @@ -8596,7 +8625,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, } break; case EXEC_OMP_TARGET_EXIT_DATA: - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_FROM: case OMP_MAP_ALWAYS_FROM: @@ -8606,16 +8635,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, case OMP_MAP_DELETE: break; case OMP_MAP_TOFROM: - n->u.map_op = OMP_MAP_FROM; + n->u.map.op = OMP_MAP_FROM; break; case OMP_MAP_ALWAYS_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_FROM; + n->u.map.op = OMP_MAP_ALWAYS_FROM; break; case OMP_MAP_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_PRESENT_FROM; + n->u.map.op = OMP_MAP_PRESENT_FROM; break; case OMP_MAP_ALWAYS_PRESENT_TOFROM: - n->u.map_op = OMP_MAP_ALWAYS_PRESENT_FROM; + n->u.map.op = OMP_MAP_ALWAYS_PRESENT_FROM; break; default: gfc_error ("TARGET EXIT DATA with map-type other " diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc index 6d463036966..b7dea11461f 100644 --- a/gcc/fortran/trans-decl.cc +++ b/gcc/fortran/trans-decl.cc @@ -6744,7 +6744,7 @@ add_clause (gfc_symbol *sym, gfc_omp_map_op map_op) n = gfc_get_omp_namelist (); n->sym = sym; - n->u.map_op = map_op; + n->u.map.op = map_op; if (!module_oacc_clauses) module_oacc_clauses = gfc_get_omp_clauses (); @@ -6846,10 +6846,10 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symbol *sym, bool block) for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next) { - switch (n->u.map_op) + switch (n->u.map.op) { case OMP_MAP_DEVICE_RESIDENT: - n->u.map_op = OMP_MAP_FORCE_ALLOC; + n->u.map.op = OMP_MAP_FORCE_ALLOC; break; default: diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index a2bf15665b3..fa1bfd41380 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -3139,7 +3139,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, || (n->expr && gfc_expr_attr (n->expr).pointer))) always_modifier = true; - switch (n->u.map_op) + if (n->u.map.readonly) + OMP_CLAUSE_MAP_READONLY (node) = 1; + + switch (n->u.map.op) { case OMP_MAP_ALLOC: OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC); @@ -3266,8 +3269,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, && n->sym->attr.omp_declare_target && (always_modifier || n->sym->attr.pointer) && op != EXEC_OMP_TARGET_EXIT_DATA - && n->u.map_op != OMP_MAP_DELETE - && n->u.map_op != OMP_MAP_RELEASE) + && n->u.map.op != OMP_MAP_DELETE + && n->u.map.op != OMP_MAP_RELEASE) { gcc_assert (n->sym->ts.u.cl->backend_decl); node5 = build_omp_clause (input_location, OMP_CLAUSE_MAP); @@ -3333,7 +3336,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, { enum gomp_map_kind gmk = GOMP_MAP_POINTER; if (op == EXEC_OMP_TARGET_EXIT_DATA - && n->u.map_op == OMP_MAP_DELETE) + && n->u.map.op == OMP_MAP_DELETE) gmk = GOMP_MAP_DELETE; else if (op == EXEC_OMP_TARGET_EXIT_DATA) gmk = GOMP_MAP_RELEASE; @@ -3356,7 +3359,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, { enum gomp_map_kind gmk; if (op == EXEC_OMP_TARGET_EXIT_DATA - && n->u.map_op == OMP_MAP_DELETE) + && n->u.map.op == OMP_MAP_DELETE) gmk = GOMP_MAP_DELETE; else if (op == EXEC_OMP_TARGET_EXIT_DATA) gmk = GOMP_MAP_RELEASE; @@ -3388,18 +3391,18 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP); OMP_CLAUSE_DECL (node2) = decl; OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); - if (n->u.map_op == OMP_MAP_DELETE) + if (n->u.map.op == OMP_MAP_DELETE) map_kind = GOMP_MAP_DELETE; else if (op == EXEC_OMP_TARGET_EXIT_DATA - || n->u.map_op == OMP_MAP_RELEASE) + || n->u.map.op == OMP_MAP_RELEASE) map_kind = GOMP_MAP_RELEASE; else map_kind = GOMP_MAP_TO_PSET; OMP_CLAUSE_SET_MAP_KIND (node2, map_kind); if (op != EXEC_OMP_TARGET_EXIT_DATA - && n->u.map_op != OMP_MAP_DELETE - && n->u.map_op != OMP_MAP_RELEASE) + && n->u.map.op != OMP_MAP_DELETE + && n->u.map.op != OMP_MAP_RELEASE) { node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP); @@ -3417,7 +3420,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, = gfc_conv_descriptor_data_get (decl); OMP_CLAUSE_SIZE (node3) = size_int (0); - if (n->u.map_op == OMP_MAP_ATTACH) + if (n->u.map.op == OMP_MAP_ATTACH) { /* Standalone attach clauses used with arrays with descriptors must copy the descriptor to the @@ -3433,7 +3436,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, node3 = NULL; goto finalize_map_clause; } - else if (n->u.map_op == OMP_MAP_DETACH) + else if (n->u.map.op == OMP_MAP_DETACH) { OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH); /* Similarly to above, we don't want to unmap PTR @@ -3626,8 +3629,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, to perform a single attach/detach operation, of the pointer itself, not of the pointed-to object. */ if (openacc - && (n->u.map_op == OMP_MAP_ATTACH - || n->u.map_op == OMP_MAP_DETACH)) + && (n->u.map.op == OMP_MAP_ATTACH + || n->u.map.op == OMP_MAP_DETACH)) { OMP_CLAUSE_DECL (node) = build_fold_addr_expr (OMP_CLAUSE_DECL (node)); @@ -3656,7 +3659,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, se.string_length), TYPE_SIZE_UNIT (tmp)); gomp_map_kind kind; - if (n->u.map_op == OMP_MAP_DELETE) + if (n->u.map.op == OMP_MAP_DELETE) kind = GOMP_MAP_DELETE; else if (op == EXEC_OMP_TARGET_EXIT_DATA) kind = GOMP_MAP_RELEASE; @@ -3713,8 +3716,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, to perform a single attach/detach operation, of the pointer itself, not of the pointed-to object. */ if (openacc - && (n->u.map_op == OMP_MAP_ATTACH - || n->u.map_op == OMP_MAP_DETACH)) + && (n->u.map.op == OMP_MAP_ATTACH + || n->u.map.op == OMP_MAP_DETACH)) { OMP_CLAUSE_DECL (node) = build_fold_addr_expr (inner); @@ -3806,8 +3809,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, { /* Bare attach and detach clauses don't want any additional nodes. */ - if ((n->u.map_op == OMP_MAP_ATTACH - || n->u.map_op == OMP_MAP_DETACH) + if ((n->u.map.op == OMP_MAP_ATTACH + || n->u.map.op == OMP_MAP_DETACH) && (POINTER_TYPE_P (TREE_TYPE (inner)) || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))) { @@ -3840,8 +3843,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, map_kind = ((GOMP_MAP_ALWAYS_P (map_kind) || gfc_expr_attr (n->expr).pointer) ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO); - else if (n->u.map_op == OMP_MAP_RELEASE - || n->u.map_op == OMP_MAP_DELETE) + else if (n->u.map.op == OMP_MAP_RELEASE + || n->u.map.op == OMP_MAP_DELETE) ; else if (op == EXEC_OMP_TARGET_EXIT_DATA || op == EXEC_OACC_EXIT_DATA) @@ -4088,6 +4091,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } if (n->u.present_modifier) OMP_CLAUSE_MOTION_PRESENT (node) = 1; + if (list == OMP_LIST_CACHE && n->u.map.readonly) + OMP_CLAUSE__CACHE__READONLY (node) = 1; omp_clauses = gfc_trans_add_clause (node, omp_clauses); } break; @@ -6561,7 +6566,7 @@ gfc_add_clause_implicitly (gfc_omp_clauses *clauses_out, n2->where = n->where; n2->sym = n->sym; if (is_target) - n2->u.map_op = OMP_MAP_TOFROM; + n2->u.map.op = OMP_MAP_TOFROM; if (tail) { tail->next = n2; diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c new file mode 100644 index 00000000000..34fc92c24d5 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c @@ -0,0 +1,59 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +struct S +{ + int *ptr; + float f; +}; + +int a[32], b[32]; +#pragma acc declare copyin(readonly: a) copyin(b) + +int main (void) +{ + int x[32], y[32]; + struct S s = {x, 0}; + + #pragma acc parallel copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32]) + { + #pragma acc cache (readonly: x[:32]) + #pragma acc cache (y[:32]) + } + + #pragma acc kernels copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32]) + { + #pragma acc cache (readonly: x[:32]) + #pragma acc cache (y[:32]) + } + + #pragma acc serial copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32]) + { + #pragma acc cache (readonly: x[:32]) + #pragma acc cache (y[:32]) + } + + #pragma acc data copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32]) + { + #pragma acc cache (readonly: x[:32]) + #pragma acc cache (y[:32]) + } + + #pragma acc enter data copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32]) + + return 0; +} + +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */ + +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ + +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */ +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 new file mode 100644 index 00000000000..696ebd08321 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 @@ -0,0 +1,89 @@ +! { dg-additional-options "-fdump-tree-original" } + +subroutine foo (a, n) + integer :: n, a(:) + integer :: i, b(n), c(n) + !$acc parallel copyin(readonly: a(:), b(:n)) copyin(c(:)) + do i = 1,32 + !$acc cache (readonly: a(:), b(:n)) + !$acc cache (c(:)) + enddo + !$acc end parallel + + !$acc kernels copyin(readonly: a(:), b(:n)) copyin(c(:)) + do i = 1,32 + !$acc cache (readonly: a(:), b(:n)) + !$acc cache (c(:)) + enddo + !$acc end kernels + + !$acc serial copyin(readonly: a(:), b(:n)) copyin(c(:)) + do i = 1,32 + !$acc cache (readonly: a(:), b(:n)) + !$acc cache (c(:)) + enddo + !$acc end serial + + !$acc data copyin(readonly: a(:), b(:n)) copyin(c(:)) + do i = 1,32 + !$acc cache (readonly: a(:), b(:n)) + !$acc cache (c(:)) + enddo + !$acc end data + + !$acc enter data copyin(readonly: a(:), b(:n)) copyin(c(:)) + +end subroutine foo + +program main + integer :: g(32), h(32) + integer :: i, n = 32, a(32) + integer :: b(32), c(32) + + !$acc declare copyin(readonly: g), copyin(h) + + !$acc parallel copyin(readonly: a(:32), b(:n)) copyin(c(:)) + do i = 1,32 + !$acc cache (readonly: a(:), b(:n)) + !$acc cache (c(:)) + enddo + !$acc end parallel + + !$acc kernels copyin(readonly: a(:), b(:n)) copyin(c(:)) + do i = 1,32 + !$acc cache (readonly: a(:), b(:n)) + !$acc cache (c(:)) + enddo + !$acc end kernels + + !$acc serial copyin(readonly: a(:), b(:n)) copyin(c(:)) + do i = 1,32 + !$acc cache (readonly: a(:), b(:n)) + !$acc cache (c(:)) + enddo + !$acc end serial + + !$acc data copyin(readonly: a(:), b(:n)) copyin(c(:)) + do i = 1,32 + !$acc cache (readonly: a(:), b(:n)) + !$acc cache (c(:)) + enddo + !$acc end data + + !$acc enter data copyin(readonly: a(:), b(:n)) copyin(c(:)) + +end program main + +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } } + +! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } } diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 8a89462bd7e..d529712306d 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -1344,6 +1344,12 @@ struct GTY(()) tree_base { TYPE_READONLY in all types + OMP_CLAUSE_MAP_READONLY in + OMP_CLAUSE_MAP + + OMP_CLAUSE__CACHE__READONLY in + OMP_CLAUSE__CACHE_ + constant_flag: TREE_CONSTANT in diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 654f5247e3a..926f7e006a7 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -913,6 +913,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case OMP_CLAUSE_MAP: pp_string (pp, "map("); + if (OMP_CLAUSE_MAP_READONLY (clause)) + pp_string (pp, "readonly,"); switch (OMP_CLAUSE_MAP_KIND (clause)) { case GOMP_MAP_ALLOC: @@ -1095,6 +1097,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case OMP_CLAUSE__CACHE_: pp_string (pp, "("); + if (OMP_CLAUSE__CACHE__READONLY (clause)) + pp_string (pp, "readonly:"); dump_generic_node (pp, OMP_CLAUSE_DECL (clause), spc, flags, false); goto print_clause_size; diff --git a/gcc/tree.h b/gcc/tree.h index e1fc6c2221d..b67a37d6522 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1841,6 +1841,14 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_flag) +/* Nonzero if OpenACC 'readonly' modifier set, used for 'copyin'. */ +#define OMP_CLAUSE_MAP_READONLY(NODE) \ + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) + +/* Same as above, for use in OpenACC cache directives. */ +#define OMP_CLAUSE__CACHE__READONLY(NODE) \ + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_)) + /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present' clause. */ #define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \