From patchwork Thu Mar 14 15:09:58 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1912127 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; dkim=pass (2048-bit key; unprotected) header.d=baylibre-com.20230601.gappssmtp.com header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256 header.s=20230601 header.b=vxCOQqFv; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4TwW6f0Nbgz1yWt for ; Fri, 15 Mar 2024 02:11:05 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 8DCD13857BA3 for ; Thu, 14 Mar 2024 15:11:03 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x334.google.com (mail-wm1-x334.google.com [IPv6:2a00:1450:4864:20::334]) by sourceware.org (Postfix) with ESMTPS id 87D1C3858C31 for ; Thu, 14 Mar 2024 15:10:28 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 87D1C3858C31 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 87D1C3858C31 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::334 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-x334.google.com with SMTP id 5b1f17b1804b1-413eb7cf7dfso7277985e9.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=UOeGc7kl5sqQEH8C9hDxHvCP5d0Rw1SGWg0Six9nStbb8gSnRSeY3WMAOSliH+nLyX aKbEgchsVjO8dwlhAbKawaXsViRPzkOOmWDU9uNm+Gy49v1xz2IWG+OrI6yW9Z3tZ09p 43AvbvZtrgwyOtxlVn1yZvBS/msKPmNfIUaRpYYeKBmf8J2leEJ7V9tPu4YM8OS5xO/I 6OdeM8q66mFy5JpzgOXK2RwMuAPpBTGTgB/t7+zqiUk0mBfaHfpz7YLPrkBWYj8e+bTa AeIeuRFew6PPkm6ice7EJl4WiMuEXnfmQGB5Q8cK1nt9yAN6Pw3mTC6ElWGGZihk4A/T UKNg== X-Forwarded-Encrypted: i=1; AJvYcCXoAj3wj749i6B1FTTsE3CVDPK3A1yhUqwf6lCkldgxgpveHyufgJVl7juxxy0oxPMpz+K6c/WTHutpCzyq4S+2uNkzEMvitg== X-Gm-Message-State: AOJu0YxvHdWrK1W9uRUT7TGUI+pMRijyTPV8DKEGVnlEimO9DCveSJGb NQaQgR7SgwuQLUdTsP7sTvVAat6G69ez8vJZj3kSUudx8/d5PK4K4z7XUD0YEgM= 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 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 X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org 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" } } 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