From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from us-smtp-delivery-124.mimecast.com (us-smtp-delivery-124.mimecast.com [170.10.133.124]) by sourceware.org (Postfix) with ESMTPS id 58FBB3858D39 for ; Tue, 11 Jan 2022 11:53:31 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 58FBB3858D39 Received: from mimecast-mx01.redhat.com (mimecast-mx01.redhat.com [209.132.183.4]) by relay.mimecast.com with ESMTP with STARTTLS (version=TLSv1.2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id us-mta-296-gFUbTp-FORa1VaLD81IvVQ-1; Tue, 11 Jan 2022 06:53:29 -0500 X-MC-Unique: gFUbTp-FORa1VaLD81IvVQ-1 Received: from smtp.corp.redhat.com (int-mx08.intmail.prod.int.phx2.redhat.com [10.5.11.23]) (using TLSv1.2 with cipher AECDH-AES256-SHA (256/256 bits)) (No client certificate requested) by mimecast-mx01.redhat.com (Postfix) with ESMTPS id 659781853020; Tue, 11 Jan 2022 11:53:28 +0000 (UTC) Received: from tucnak.zalov.cz (unknown [10.39.195.246]) by smtp.corp.redhat.com (Postfix) with ESMTPS id EC4AF2B44F; Tue, 11 Jan 2022 11:53:27 +0000 (UTC) Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.16.1/8.16.1) with ESMTPS id 20BBrOLk1109363 (version=TLSv1.3 cipher=TLS_AES_256_GCM_SHA384 bits=256 verify=NOT); Tue, 11 Jan 2022 12:53:25 +0100 Received: (from jakub@localhost) by tucnak.zalov.cz (8.16.1/8.16.1/Submit) id 20BBrO1W1109362; Tue, 11 Jan 2022 12:53:24 +0100 Date: Tue, 11 Jan 2022 12:53:24 +0100 From: Jakub Jelinek To: Marcel Vollweiler Cc: gcc-patches@gcc.gnu.org, fortran@gcc.gnu.org Subject: Re: [PATCH] C, C++, Fortran, OpenMP: Add 'has_device_addr' clause to 'target' construct Message-ID: <20220111115324.GV2646553@tucnak> Reply-To: Jakub Jelinek References: <20211020123844.GU304296@tucnak> <25633d5a-d94d-0231-b626-97d28159237f@codesourcery.com> MIME-Version: 1.0 In-Reply-To: X-Scanned-By: MIMEDefang 2.84 on 10.5.11.23 X-Mimecast-Spam-Score: 0 X-Mimecast-Originator: redhat.com Content-Type: text/plain; charset=us-ascii Content-Disposition: inline X-Spam-Status: No, score=-5.5 required=5.0 tests=BAYES_00, DKIMWL_WL_HIGH, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, RCVD_IN_DNSWL_LOW, RCVD_IN_MSPIKE_H3, RCVD_IN_MSPIKE_WL, SPF_HELO_NONE, SPF_NONE, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: fortran@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Fortran mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Tue, 11 Jan 2022 11:53:33 -0000 On Wed, Nov 24, 2021 at 06:08:02PM +0100, Marcel Vollweiler wrote: > + case OMP_CLAUSE_HAS_DEVICE_ADDR: > + t = OMP_CLAUSE_DECL (c); > + if (TREE_CODE (t) == TREE_LIST) > + { > + if (handle_omp_array_sections (c, ort)) > + remove = true; > + else > + { > + t = OMP_CLAUSE_DECL (c); > + while (TREE_CODE (t) == ARRAY_REF) > + t = TREE_OPERAND (t, 0); > + } > + } > + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) > + bitmap_set_bit (&is_on_device_head, DECL_UID (t)); Why the OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR check? There is no goto into this block nor fallthru into it, and handle_omp_array_sections better shouldn't change OMP_CLAUSE_CODE. > goto check_dup_generic; > > + case OMP_CLAUSE_HAS_DEVICE_ADDR: > + t = OMP_CLAUSE_DECL (c); > + if (TREE_CODE (t) == TREE_LIST) > + if (handle_omp_array_sections (c, ort)) > + remove = true; > + else > + { > + t = OMP_CLAUSE_DECL (c); > + while (TREE_CODE (t) == ARRAY_REF) > + t = TREE_OPERAND (t, 0); > + } > + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) > + bitmap_set_bit (&is_on_device_head, DECL_UID (t)); Likewise. > + if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) > + cxx_mark_addressable (t); > + goto check_dup_generic_t; > + > case OMP_CLAUSE_USE_DEVICE_ADDR: > field_ok = true; > t = OMP_CLAUSE_DECL (c); > --- a/gcc/fortran/gfortran.h > +++ b/gcc/fortran/gfortran.h > @@ -1391,7 +1391,8 @@ enum > OMP_LIST_USE_DEVICE_PTR, > OMP_LIST_USE_DEVICE_ADDR, > OMP_LIST_NONTEMPORAL, > - OMP_LIST_NUM > + OMP_LIST_HAS_DEVICE_ADDR, > + OMP_LIST_NUM /* must be the last */ Capital M and . at the end. > @@ -2077,6 +2078,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, > } > break; > case 'h': > + if ((mask & OMP_CLAUSE_HAS_DEVICE_ADDR) > + && gfc_match_omp_variable_list > + ("has_device_addr (", > + &c->lists[OMP_LIST_HAS_DEVICE_ADDR], false, NULL, NULL, > + true) == MATCH_YES) Formatting, true should be IMO below &c->lists. > + continue; > if ((mask & OMP_CLAUSE_HINT) > && (m = gfc_match_dupl_check (!c->hint, "hint", true, &c->hint)) > != MATCH_NO) > @@ -2850,7 +2857,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, > if ((mask & OMP_CLAUSE_USE_DEVICE_ADDR) > && gfc_match_omp_variable_list > ("use_device_addr (", > - &c->lists[OMP_LIST_USE_DEVICE_ADDR], false) == MATCH_YES) > + &c->lists[OMP_LIST_USE_DEVICE_ADDR], false, NULL, NULL, > + true) == MATCH_YES) Likewise. > --- a/gcc/fortran/trans-openmp.c > +++ b/gcc/fortran/trans-openmp.c > @@ -1910,7 +1910,17 @@ gfc_trans_omp_variable_list (enum omp_clause_code code, > tree t = gfc_trans_omp_variable (namelist->sym, declare_simd); > if (t != error_mark_node) > { > - tree node = build_omp_clause (input_location, code); > + tree node; > + /* For HAS_DEVICE_ADDR of an array descriptor, firstprivatize the > + descriptor such that the bounds are available; its data component > + is unmodified; it is handled as device address inside target. */ > + if (code == OMP_CLAUSE_HAS_DEVICE_ADDR > + && (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (t)) > + || (POINTER_TYPE_P (TREE_TYPE (t)) > + && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (t)))))) > + node = build_omp_clause (input_location, OMP_CLAUSE_FIRSTPRIVATE); Not sure about the above, > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -10024,6 +10024,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, > flags = GOVD_EXPLICIT; > goto do_add; > > + case OMP_CLAUSE_HAS_DEVICE_ADDR: > + decl = OMP_CLAUSE_DECL (c); > + if (TREE_CODE (decl) == ARRAY_REF) > + { > + flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT; > + while (TREE_CODE (decl) == ARRAY_REF) > + decl = TREE_OPERAND (decl, 0); > + goto do_add_decl; but this looks weird. If decl after stripping the ARRAY_REFs is a var with pointer type, sure, firstprivatizing it is the way to go. But it can be also a variable with ARRAY_TYPE, can't it? Something like: int a[64]; #pragma omp target data map(a) use_device_addr(a) { #pragma omp target has_device_addr(a[3:16]) a[3] = 1; } and in this case firstprivatization of a looks wrong. use_device_addr should replace (but only at omp-low.c time I think) a used in the block with the remapped a (i.e. *device_address_of_a). Or perhaps it could be a non-static data member with array type inside of a C++ method. > + case OMP_CLAUSE_HAS_DEVICE_ADDR: > + decl = OMP_CLAUSE_DECL (c); > + if (TREE_CODE (decl) == ARRAY_REF) > + while (TREE_CODE (decl) == ARRAY_REF) > + decl = TREE_OPERAND (decl, 0); Isn't this equivalent to just the while loop without the if? Jakub