Hi! On 2024-03-13T10:12:17+0100, I wrote: > On 2024-03-07T17:02:02+0900, Chung-Lin Tang wrote: >> Also added simple 'declare' tests, but there is not anything to scan in the 'tree-original' dump though. > > Yeah, the current OpenACC 'declare' implementation is "special". Actually -- commit 38958ac987dc3e6162e2ddaba3c7e7f41381e079 "OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing", see attached. But I realized another thing: don't we have to handle the 'readonly' modifier also in Fortran module files, that is, next to the OpenACC 'declare' 'copyin' handling in 'gcc/fortran/module.cc': 'AB_OACC_DECLARE_COPYIN' etc.? Chung-Lin, please check, via test cases. 'gfortran.dg/goacc/routine-module*', for example, should provide some guidance of how to achieve actual module file use, and then do the same 'scan-tree-dump' as in the current 'readonly' modifier test cases. I suppose the code changes would look similar to commit a61f6afbee370785cf091fe46e2e022748528307 "OpenACC 'nohost' clause", for example. By means of only emitting a tag in the module file if the 'readonly' modifier is specified, we should maintain compatibility with the current 'MOD_VERSION'. Grüße Thomas >> 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/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" } }