From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from mail-wm1-x32c.google.com (mail-wm1-x32c.google.com [IPv6:2a00:1450:4864:20::32c]) by sourceware.org (Postfix) with ESMTPS id D21F53858C3A for ; Thu, 14 Mar 2024 15:10:28 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org D21F53858C3A Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=baylibre.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=baylibre.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org D21F53858C3A Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::32c ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1710429034; cv=none; b=v3SazwGqQsNEyvBmBA1S6D0NYIHBeXmt4VQjthiU4prNpzZa8nmbixunGvjhz7ccFhj2GOWQnC8G/XXPKwBU3wr0MBE9p34lG1D5m/TCXdCDguN4F1ZDS+q6GPgEsZuCv+AGd1E4m3Fq1d7lho6WTr/HKDq0nknvyWv3EvzA1Po= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1710429034; c=relaxed/simple; bh=ab3Z7XT+OaGM7pkNb9E3Pi1aA3JAHqGkC0fC7ABZPEw=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=Rf3OMcygTc1xhzPyPHnDBn3XMt9TboxH+jZN3ABUyywiMa/TCCpHxnGun7omyKwg8Rcz3m2GEMgHfMBijtlIOZ1jSFHwizgtu9TAPiMaJ4u+56+MxZRqDL49ILXwQQLmqr8SWrw+iFvt8tH0HXUoZuk0x+H+T7Ztmm8iMkPp54k= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x32c.google.com with SMTP id 5b1f17b1804b1-41312232c7aso13621125e9.0 for ; Thu, 14 Mar 2024 08:10:28 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1710429027; x=1711033827; darn=gcc.gnu.org; h=mime-version:message-id:date:user-agent:references:in-reply-to :subject:cc:to:from:from:to:cc:subject:date:message-id:reply-to; bh=e8jWsLKpuT3gRJQOM0Its6qnkZ16A64YZ39Q0pnU2bI=; b=vxCOQqFvqCS3YVMdzh4cU9bfSoWlvaNb4ZZcbwhziYpWtA5ySUr0LhXJ4tAdYTqgM5 Jcn4FA/jVCniw51jKbbVF6M3CVsq8461izvYTYB1YXVuZ8GqpWeTAhrLWR+/SlvVWVPB riQ9lJpFzW1xUFuj/5XdGFgcVaA5ClO3EyRMmlsLFmrJ29KiJwNFqGy+59tFgOLwreO+ It8OJ7uNfu/pyxxZ8QM4MN4epeynQzS9zVKJqDHE1Pp7CttitYterOwI98A9Nl6x6TXH jZKzKYO+9hL5O4m4aSivjf6gKig6NunyP6e7WzkWmxjAQpieDIUCltn0QyN4+t4St9H/ hwpQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1710429027; x=1711033827; h=mime-version:message-id:date:user-agent:references:in-reply-to :subject:cc:to:from:x-gm-message-state:from:to:cc:subject:date :message-id:reply-to; bh=e8jWsLKpuT3gRJQOM0Its6qnkZ16A64YZ39Q0pnU2bI=; b=xThk8HEyV76ut6FEcBBvpWeBrojV91XuryihU35atkNz8iFDRcA+tF8kGAxgJAj9p7 L1lEncEougqCDgC+qX3AJZ/HGj1Kmbm4tfJglwdqCL+ptCeiuT7w1yvVTE414CtZDIuR /U60RrxeQRMeLPdNw6FmYoL2o/dNLT0toC0oIvKctClKB2Lpyi2r7+N2Y0qOV9U6D3n8 7M4AKH9ZdnQbpIGWsu78nEeyPdhBJ0HS8xecyxzLu1+RDq/lm1hk3i/w/Ux9OaS1pbF4 S78HStVm2wu0+5PoLqFyakFe5MhiB0Vf/GZO81FRnjcJU+A0kakOZnn8hFPorta3Lxmd j8Ng== X-Forwarded-Encrypted: i=1; AJvYcCX9ns7PlI8eE9U+ZXVO4kCHKeVbmmRx+hjXJagRjVsu90ASJlqARXJrWt4xKw2vKz92wsZc1Od81XDCopzzdYe1YSKn X-Gm-Message-State: AOJu0YwQTHzyW1HhKdrK/2yj1eaa76Q8rRuXItFcr2E8/DPrhsyVRiyV Lp7VIUEB+SZedpF9rUn+Oh/YzlpI+8NmzeRBqtscZwc3GeRBs78VQzrN2VOOxyo= X-Google-Smtp-Source: AGHT+IFX+EVcpXxj8SkCBuKugS4cZCqCOK1IjY55B81SmgdtTDQAPXW2YzgiTzr1noG3wxzX2p5MZg== X-Received: by 2002:a05:600c:3b9a:b0:413:ee4c:a35d with SMTP id n26-20020a05600c3b9a00b00413ee4ca35dmr2149455wms.16.1710429027110; Thu, 14 Mar 2024 08:10:27 -0700 (PDT) Received: from euler.schwinge.homeip.net (p200300c8b70336000b0134869109dcb1.dip0.t-ipconnect.de. [2003:c8:b703:3600:b01:3486:9109:dcb1]) by smtp.gmail.com with ESMTPSA id o18-20020a05600c4fd200b004132f9cf053sm5829708wmq.33.2024.03.14.08.10.07 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Thu, 14 Mar 2024 08:10:17 -0700 (PDT) From: Thomas Schwinge To: Chung-Lin Tang , gcc-patches@gcc.gnu.org Cc: Tobias Burnus , 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) In-Reply-To: <87le6mebri.fsf@euler.schwinge.ddns.net> References: <87lefaaesb.fsf@euler.schwinge.homeip.net> <87ttqd7m8e.fsf@euler.schwinge.homeip.net> <2f568f7d-807b-41d1-befb-40039e2edb74@pllab.cs.nthu.edu.tw> <87le6mebri.fsf@euler.schwinge.ddns.net> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/29.1 (x86_64-pc-linux-gnu) Date: Thu, 14 Mar 2024 16:09:58 +0100 Message-ID: <87sf0s6e9l.fsf@euler.schwinge.ddns.net> MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="=-=-=" X-Spam-Status: No, score=-10.9 required=5.0 tests=BAYES_00,DKIM_SIGNED,DKIM_VALID,GIT_PATCH_0,RCVD_IN_DNSWL_NONE,SPF_HELO_NONE,SPF_PASS,TXREP,T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: --=-=-= Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: quoted-printable 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 'd= eclare' 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=C3=BC=C3=9Fe Thomas >> diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tre= e.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 =3D=3D OMP_LIST_MAP || list_type =3D=3D OMP_LIST_C= ACHE) >> + && n->u.map.readonly) >> + fputs ("readonly,", dumpfile); >> if (list_type =3D=3D 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 =3D=3D 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 =3D *head; n; n =3D n->next) >> - n->u.map_op =3D map_op; >> + n->u.map.op =3D map_op; >> return true; >> } >>=20=20 >> @@ -1524,7 +1524,7 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_c= lauses *c, bool openacc, >> gfc_omp_namelist *p =3D gfc_get_omp_namelist (), **tl; >> p->sym =3D n->sym; >> p->where =3D p->where; >> - p->u.map_op =3D OMP_MAP_ALWAYS_TOFROM; >> + p->u.map.op =3D OMP_MAP_ALWAYS_TOFROM; >>=20=20 >> tl =3D &c->lists[OMP_LIST_MAP]; >> while (*tl) >> @@ -2181,11 +2181,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, con= st omp_mask mask, >> { >> if (openacc) >> { >> - if (gfc_match ("copyin ( ") =3D=3D MATCH_YES >> - && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], >> - OMP_MAP_TO, true, >> - allow_derived)) >> - continue; >> + if (gfc_match ("copyin ( ") =3D=3D MATCH_YES) >> + { >> + bool readonly =3D gfc_match ("readonly : ") =3D=3D MATCH_YES; >> + head =3D NULL; >> + if (gfc_match_omp_variable_list ("", >> + &c->lists[OMP_LIST_MAP], >> + true, NULL, &head, true, >> + allow_derived) >> + =3D=3D MATCH_YES) >> + { >> + gfc_omp_namelist *n; >> + for (n =3D *head; n; n =3D n->next) >> + { >> + n->u.map.op =3D OMP_MAP_TO; >> + n->u.map.readonly =3D 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 =3D *head; n; n =3D n->next) >> - n->u.map_op =3D map_op; >> + n->u.map.op =3D map_op; >> continue; >> } >> gfc_current_locus =3D old_loc; >> @@ -4002,7 +4016,7 @@ gfc_match_oacc_declare (void) >> if (gfc_current_ns->proc_name >> && gfc_current_ns->proc_name->attr.flavor =3D=3D FL_MODULE) >> { >> - if (n->u.map_op !=3D OMP_MAP_ALLOC && n->u.map_op !=3D OMP_MAP_TO) >> + if (n->u.map.op !=3D OMP_MAP_ALLOC && n->u.map.op !=3D 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; >> } >>=20=20 >> - 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 =3D false; >> gfc_omp_clauses *c =3D 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 li= ke a >> only_sections variant next to its allow_sections parameter. */ >> - match m =3D gfc_match_omp_variable_list (" (", >> - &c->lists[OMP_LIST_CACHE], true, >> - NULL, NULL, true); >> + match m =3D gfc_match (" ( "); >> if (m !=3D MATCH_YES) >> { >> gfc_free_omp_clauses(c); >> return m; >> } >>=20=20 >> - if (gfc_current_state() !=3D COMP_DO=20 >> + if (gfc_match ("readonly : ") =3D=3D MATCH_YES) >> + readonly =3D true; >> + >> + gfc_omp_namelist **head =3D NULL; >> + m =3D gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_CACHE], tru= e, >> + NULL, &head, true); >> + if (m !=3D MATCH_YES) >> + { >> + gfc_free_omp_clauses(c); >> + return m; >> + } >> + >> + if (readonly) >> + for (gfc_omp_namelist *n =3D *head; n; n =3D n->next) >> + n->u.map.readonly =3D true; >> + >> + if (gfc_current_state() !=3D COMP_DO >> && gfc_current_state() !=3D 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_claus= es *omp_clauses, >> } >> if (openacc >> && list =3D=3D OMP_LIST_MAP >> - && (n->u.map_op =3D=3D OMP_MAP_ATTACH >> - || n->u.map_op =3D=3D OMP_MAP_DETACH)) >> + && (n->u.map.op =3D=3D OMP_MAP_ATTACH >> + || n->u.map.op =3D=3D OMP_MAP_DETACH)) >> { >> symbol_attribute attr; >> if (n->expr) >> @@ -8447,7 +8476,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_claus= es *omp_clauses, >> if (!attr.pointer && !attr.allocatable) >> gfc_error ("%qs clause argument must be ALLOCATABLE or " >> "a POINTER at %L", >> - (n->u.map_op =3D=3D OMP_MAP_ATTACH) ? "attach" >> + (n->u.map.op =3D=3D OMP_MAP_ATTACH) ? "attach" >> : "detach", &n->where); >> } >> if (lastref >> @@ -8518,7 +8547,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_claus= es *omp_clauses, >> else if (openacc) >> { >> if (list =3D=3D OMP_LIST_MAP >> - && n->u.map_op =3D=3D OMP_MAP_FORCE_DEVICEPTR) >> + && n->u.map.op =3D=3D 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_claus= es *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_claus= es *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_cla= uses *omp_clauses, >> case OMP_MAP_PRESENT_ALLOC: >> break; >> case OMP_MAP_TOFROM: >> - n->u.map_op =3D OMP_MAP_TO; >> + n->u.map.op =3D OMP_MAP_TO; >> break; >> case OMP_MAP_ALWAYS_TOFROM: >> - n->u.map_op =3D OMP_MAP_ALWAYS_TO; >> + n->u.map.op =3D OMP_MAP_ALWAYS_TO; >> break; >> case OMP_MAP_PRESENT_TOFROM: >> - n->u.map_op =3D OMP_MAP_PRESENT_TO; >> + n->u.map.op =3D OMP_MAP_PRESENT_TO; >> break; >> case OMP_MAP_ALWAYS_PRESENT_TOFROM: >> - n->u.map_op =3D OMP_MAP_ALWAYS_PRESENT_TO; >> + n->u.map.op =3D 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_claus= es *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_cla= uses *omp_clauses, >> case OMP_MAP_DELETE: >> break; >> case OMP_MAP_TOFROM: >> - n->u.map_op =3D OMP_MAP_FROM; >> + n->u.map.op =3D OMP_MAP_FROM; >> break; >> case OMP_MAP_ALWAYS_TOFROM: >> - n->u.map_op =3D OMP_MAP_ALWAYS_FROM; >> + n->u.map.op =3D OMP_MAP_ALWAYS_FROM; >> break; >> case OMP_MAP_PRESENT_TOFROM: >> - n->u.map_op =3D OMP_MAP_PRESENT_FROM; >> + n->u.map.op =3D OMP_MAP_PRESENT_FROM; >> break; >> case OMP_MAP_ALWAYS_PRESENT_TOFROM: >> - n->u.map_op =3D OMP_MAP_ALWAYS_PRESENT_FROM; >> + n->u.map.op =3D 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) >>=20=20 >> n =3D gfc_get_omp_namelist (); >> n->sym =3D sym; >> - n->u.map_op =3D map_op; >> + n->u.map.op =3D map_op; >>=20=20 >> if (!module_oacc_clauses) >> module_oacc_clauses =3D gfc_get_omp_clauses (); >> @@ -6846,10 +6846,10 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symb= ol *sym, bool block) >>=20=20 >> for (n =3D omp_clauses->lists[OMP_LIST_MAP]; n; n =3D n->next) >> { >> - switch (n->u.map_op) >> + switch (n->u.map.op) >> { >> case OMP_MAP_DEVICE_RESIDENT: >> - n->u.map_op =3D OMP_MAP_FORCE_ALLOC; >> + n->u.map.op =3D OMP_MAP_FORCE_ALLOC; >> break; >>=20=20 >> 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_om= p_clauses *clauses, >> || (n->expr && gfc_expr_attr (n->expr).pointer))) >> always_modifier =3D true; >>=20=20 >> - switch (n->u.map_op) >> + if (n->u.map.readonly) >> + OMP_CLAUSE_MAP_READONLY (node) =3D 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 !=3D EXEC_OMP_TARGET_EXIT_DATA >> - && n->u.map_op !=3D OMP_MAP_DELETE >> - && n->u.map_op !=3D OMP_MAP_RELEASE) >> + && n->u.map.op !=3D OMP_MAP_DELETE >> + && n->u.map.op !=3D OMP_MAP_RELEASE) >> { >> gcc_assert (n->sym->ts.u.cl->backend_decl); >> node5 =3D 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 =3D GOMP_MAP_POINTER; >> if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA >> - && n->u.map_op =3D=3D OMP_MAP_DELETE) >> + && n->u.map.op =3D=3D OMP_MAP_DELETE) >> gmk =3D GOMP_MAP_DELETE; >> else if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA) >> gmk =3D 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 =3D=3D EXEC_OMP_TARGET_EXIT_DATA >> - && n->u.map_op =3D=3D OMP_MAP_DELETE) >> + && n->u.map.op =3D=3D OMP_MAP_DELETE) >> gmk =3D GOMP_MAP_DELETE; >> else if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA) >> gmk =3D GOMP_MAP_RELEASE; >> @@ -3388,18 +3391,18 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_o= mp_clauses *clauses, >> node2 =3D build_omp_clause (input_location, OMP_CLAUSE_MAP); >> OMP_CLAUSE_DECL (node2) =3D decl; >> OMP_CLAUSE_SIZE (node2) =3D TYPE_SIZE_UNIT (type); >> - if (n->u.map_op =3D=3D OMP_MAP_DELETE) >> + if (n->u.map.op =3D=3D OMP_MAP_DELETE) >> map_kind =3D GOMP_MAP_DELETE; >> else if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA >> - || n->u.map_op =3D=3D OMP_MAP_RELEASE) >> + || n->u.map.op =3D=3D OMP_MAP_RELEASE) >> map_kind =3D GOMP_MAP_RELEASE; >> else >> map_kind =3D GOMP_MAP_TO_PSET; >> OMP_CLAUSE_SET_MAP_KIND (node2, map_kind); >>=20=20 >> if (op !=3D EXEC_OMP_TARGET_EXIT_DATA >> - && n->u.map_op !=3D OMP_MAP_DELETE >> - && n->u.map_op !=3D OMP_MAP_RELEASE) >> + && n->u.map.op !=3D OMP_MAP_DELETE >> + && n->u.map.op !=3D OMP_MAP_RELEASE) >> { >> node3 =3D build_omp_clause (input_location, >> OMP_CLAUSE_MAP); >> @@ -3417,7 +3420,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp= _clauses *clauses, >> =3D gfc_conv_descriptor_data_get (decl); >> OMP_CLAUSE_SIZE (node3) =3D size_int (0); >>=20=20 >> - if (n->u.map_op =3D=3D OMP_MAP_ATTACH) >> + if (n->u.map.op =3D=3D 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 =3D NULL; >> goto finalize_map_clause; >> } >> - else if (n->u.map_op =3D=3D OMP_MAP_DETACH) >> + else if (n->u.map.op =3D=3D 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 =3D=3D OMP_MAP_ATTACH >> - || n->u.map_op =3D=3D OMP_MAP_DETACH)) >> + && (n->u.map.op =3D=3D OMP_MAP_ATTACH >> + || n->u.map.op =3D=3D OMP_MAP_DETACH)) >> { >> OMP_CLAUSE_DECL (node) >> =3D 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 =3D=3D OMP_MAP_DELETE) >> + if (n->u.map.op =3D=3D OMP_MAP_DELETE) >> kind =3D GOMP_MAP_DELETE; >> else if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA) >> kind =3D 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 =3D=3D OMP_MAP_ATTACH >> - || n->u.map_op =3D=3D OMP_MAP_DETACH)) >> + && (n->u.map.op =3D=3D OMP_MAP_ATTACH >> + || n->u.map.op =3D=3D OMP_MAP_DETACH)) >> { >> OMP_CLAUSE_DECL (node) >> =3D 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 =3D=3D OMP_MAP_ATTACH >> - || n->u.map_op =3D=3D OMP_MAP_DETACH) >> + if ((n->u.map.op =3D=3D OMP_MAP_ATTACH >> + || n->u.map.op =3D=3D 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 =3D ((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 =3D=3D OMP_MAP_RELEASE >> - || n->u.map_op =3D=3D OMP_MAP_DELETE) >> + else if (n->u.map.op =3D=3D OMP_MAP_RELEASE >> + || n->u.map.op =3D=3D OMP_MAP_DELETE) >> ; >> else if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA >> || op =3D=3D 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) =3D 1; >> + if (list =3D=3D OMP_LIST_CACHE && n->u.map.readonly) >> + OMP_CLAUSE__CACHE__READONLY (node) =3D 1; >> omp_clauses =3D gfc_trans_add_clause (node, omp_clauses); >> } >> break; >> @@ -6561,7 +6566,7 @@ gfc_add_clause_implicitly (gfc_omp_clauses *clause= s_out, >> n2->where =3D n->where; >> n2->sym =3D n->sym; >> if (is_target) >> - n2->u.map_op =3D OMP_MAP_TOFROM; >> + n2->u.map.op =3D OMP_MAP_TOFROM; >> if (tail) >> { >> tail->next =3D n2; >> diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsu= ite/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 =3D 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 =3D 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 =3D 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 =3D 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 =3D 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 =3D 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 =3D 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 =3D 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 =3D 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\\(rea= donly,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\\(rea= donly,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\\(read= only,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\\(read= only,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\\(reado= nly,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\\(reado= nly,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\\(readonl= y,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\\(readonl= y,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\\(r= eadonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ ma= p\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } } >> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(r= eadonly,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=3D4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) = \\(readonly:\\*\\(integer\\(kind=3D4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len= : .+\\\]\\);" 8 "original" } } >> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(\\*\\(int= eger\\(kind=3D4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "orig= inal" } } --=-=-= Content-Type: text/x-diff Content-Disposition: inline; filename=0001-OpenACC-2.7-front-end-support-for-readonly-modifier-.patch >From 38958ac987dc3e6162e2ddaba3c7e7f41381e079 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge 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 --=-=-=--