From: Thomas Schwinge <tschwinge@baylibre.com>
To: Chung-Lin Tang <cltang@baylibre.com>, gcc-patches@gcc.gnu.org
Cc: Tobias Burnus <tburnus@baylibre.com>, fortran@gcc.gnu.org
Subject: OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing (was: [PATCH, OpenACC 2.7, v2] readonly modifier support in front-ends)
Date: Thu, 14 Mar 2024 16:09:58 +0100 [thread overview]
Message-ID: <87sf0s6e9l.fsf@euler.schwinge.ddns.net> (raw)
In-Reply-To: <87le6mebri.fsf@euler.schwinge.ddns.net>
[-- Attachment #1: Type: text/plain, Size: 23603 bytes --]
Hi!
On 2024-03-13T10:12:17+0100, I wrote:
> On 2024-03-07T17:02:02+0900, Chung-Lin Tang <cltang@pllab.cs.nthu.edu.tw> 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" } }
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-OpenACC-2.7-front-end-support-for-readonly-modifier-.patch --]
[-- Type: text/x-diff, Size: 3887 bytes --]
From 38958ac987dc3e6162e2ddaba3c7e7f41381e079 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tschwinge@baylibre.com>
Date: Thu, 14 Mar 2024 15:01:01 +0100
Subject: [PATCH] OpenACC 2.7: front-end support for readonly modifier: Add
basic OpenACC 'declare' testing
... to complement commit ddf852dac2abaca317c10b8323f338123b0585c8
"OpenACC 2.7: front-end support for readonly modifier".
gcc/testsuite/
* c-c++-common/goacc/readonly-1.c: Add basic OpenACC 'declare'
testing.
* gfortran.dg/goacc/readonly-1.f90: Likewise.
---
gcc/testsuite/c-c++-common/goacc/readonly-1.c | 5 +++++
gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 | 6 ++++++
2 files changed, 11 insertions(+)
diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
index 34fc92c24d5..300464c92e3 100644
--- a/gcc/testsuite/c-c++-common/goacc/readonly-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
@@ -8,12 +8,15 @@ struct S
int a[32], b[32];
#pragma acc declare copyin(readonly: a) copyin(b)
+/* Not visible in 'original' dump; handled via 'offload_vars'. */
int main (void)
{
int x[32], y[32];
struct S s = {x, 0};
+ #pragma acc declare copyin(readonly: x/*[:32]*/, s/*.ptr[:16]*/) copyin(y/*[:32]*/)
+
#pragma acc parallel copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
{
#pragma acc cache (readonly: x[:32])
@@ -43,6 +46,8 @@ int main (void)
return 0;
}
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc declare map\\(to:y\\) map\\(readonly,to:s\\) map\\(readonly,to:x\\)" 1 "original" } } */
+
/* { 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 } } } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
index 696ebd08321..fc1e2719e67 100644
--- a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
@@ -3,6 +3,9 @@
subroutine foo (a, n)
integer :: n, a(:)
integer :: i, b(n), c(n)
+ !!$acc declare copyin(readonly: a(:), b(:n)) copyin(c(:))
+ !$acc declare copyin(readonly: b) copyin(c)
+
!$acc parallel copyin(readonly: a(:), b(:n)) copyin(c(:))
do i = 1,32
!$acc cache (readonly: a(:), b(:n))
@@ -74,6 +77,9 @@ program main
end program main
+! The front end turns OpenACC 'declare' into OpenACC 'data'.
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc data 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:g\\) map\\(to:h\\)" 1 "original" } }
! { 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" } }
--
2.34.1
next prev parent reply other threads:[~2024-03-14 15:10 UTC|newest]
Thread overview: 18+ messages / expand[flat|nested] mbox.gz Atom feed top
2023-07-10 18:33 [PATCH, OpenACC 2.7] readonly modifier support in front-ends Chung-Lin Tang
2023-07-11 7:00 ` Tobias Burnus
2023-07-20 13:33 ` Thomas Schwinge
2023-07-20 15:08 ` Tobias Burnus
2023-08-07 13:58 ` [PATCH, OpenACC 2.7, v2] " Chung-Lin Tang
2023-10-26 9:43 ` Thomas Schwinge
2024-03-07 8:02 ` Chung-Lin Tang
2024-03-13 9:12 ` Thomas Schwinge
2024-03-14 15:09 ` Thomas Schwinge [this message]
2024-03-14 16:55 ` OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing Tobias Burnus
2024-03-14 16:55 ` Tobias Burnus
2023-07-25 15:52 ` [PATCH, OpenACC 2.7] Connect readonly modifier to points-to analysis Chung-Lin Tang
2023-10-27 14:28 ` Thomas Schwinge
2023-10-30 12:46 ` Richard Biener
2024-04-03 11:50 ` Chung-Lin Tang
2024-04-11 14:29 ` Thomas Schwinge
2024-04-12 6:17 ` Richard Biener
2024-05-16 12:36 ` Richard Biener
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=87sf0s6e9l.fsf@euler.schwinge.ddns.net \
--to=tschwinge@baylibre.com \
--cc=cltang@baylibre.com \
--cc=fortran@gcc.gnu.org \
--cc=gcc-patches@gcc.gnu.org \
--cc=tburnus@baylibre.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).