From patchwork Thu Jan 13 10:24:30 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 49958 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id B51D8389244C for ; Thu, 13 Jan 2022 10:24:56 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id B041D3857C6B for ; Thu, 13 Jan 2022 10:24:38 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org B041D3857C6B Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: 0PTVKwX1R3KKWM29ml0pWAcRB59oATVG42eL4NeAkRu4bVJrtxBLhu3X1mOg1CnRpPB+aTlv6I dneLPlgLZRg49yd38HICiPr26NpNN8m6tlx3nZEQdiehxi5bMdw5vMVsz1z6984i8InGf+8Spu BqwXvueLSOXUzyf6XytWbCchClmRrd08rcLPLZtKFVdeXQtttk+NASqdpdyrqYk7m7KN3zj2Xt Z/33gKwC+lx9JGXBymD+Jyyb5drKUa98yP6OKRkMPJSJY/G6BEDYwcSTl5I5StnLHIZPZvPqmR z1JrrgCCKnVMLgDCNiUfkADk X-IronPort-AV: E=Sophos;i="5.88,284,1635235200"; d="scan'208,223";a="70716206" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 13 Jan 2022 02:24:37 -0800 IronPort-SDR: 1scImisvdg0iI9pq7tPXPx1JbEphOfBs+JoB5HcoEpdDLThjZxdHakgbxetZh7guxSznY1UI1f Tnf0ODZKPBwGD1qfyxIKIt4+ljp5khCrMSkuhivr4C+ogCQr3sXXn7Ws3oj0Ejb2eOghcQOEu3 igMA8MWivjRI5g+FFsXtf0j8UsSqia5l8XowOQ/ISKyv76PW4C9O4GDlBnMwHCqfDRYivFwmZY Gh7jL3KL6Vfqw1pbk0raZM7XAvdUVS+u+JaFoRN+Tn1F3qjsLKDSXvEJXd4+7QkkS6kZ5AhqEA RAs= From: Thomas Schwinge To: Jakub Jelinek , Subject: Host and offload targets have no common meaning of address spaces In-Reply-To: References: <992c7c29-5773-45b6-6fb7-ffb71299a98f@mentor.com> <87r1f2puss.fsf@euler.schwinge.homeip.net> <87a6lhhkvp.fsf@euler.schwinge.homeip.net> <20210816082104.GU2380545@tucnak> <871r6pnqez.fsf@euler.schwinge.homeip.net> <87o89np2es.fsf@dem-tschwing-1.ger.mentorg.com> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Thu, 13 Jan 2022 11:24:30 +0100 Message-ID: <877db4djc1.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.1 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, 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: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hi! Jakub, I'd still like your comment on the two "should we" questions cited below. On 2021-08-24T13:43:38+0200, Richard Biener via Gcc-patches wrote: > On Tue, Aug 24, 2021 at 12:23 PM Thomas Schwinge wrote: >> On 2021-08-19T22:13:56+0200, I wrote: >> > On 2021-08-16T10:21:04+0200, Jakub Jelinek wrote: >> >> On Mon, Aug 16, 2021 at 10:08:42AM +0200, Thomas Schwinge wrote: >> > |> Concerning the current 'gcc/omp-low.c:omp_build_component_ref', for the >> > |> current set of offloading testcases, we never see a >> > |> '!ADDR_SPACE_GENERIC_P' there, so the address space handling doesn't seem >> > |> to be necessary there (but also won't do any harm: no-op). >> >> >> >> Are you sure this can't trigger? >> >> Say >> >> extern int __seg_fs a; >> >> >> >> void >> >> foo (void) >> >> { >> >> #pragma omp parallel private (a) >> >> a = 2; >> >> } >> > >> > That test case doesn't run into 'omp_build_component_ref' at all, >> > but [I've pushed an altered and extended variant that does], >> > "Add 'libgomp.c/address-space-1.c'". >> > >> > In this case, 'omp_build_component_ref' called via host compilation >> > 'pass_lower_omp', it's the 'field_type' that has 'address-space-1' >> > [...]: >> > >> > (gdb) call debug_tree(field_type) >> > > > type > >> >> I think keeping the qual addr space here is the wrong thing to do, >> >> it should keep the other quals and clear the address space instead, >> >> the whole struct is going to be in generic addres space, isn't it? >> > >> > Correct for 'omp_build_component_ref' called via host compilation >> > 'pass_lower_omp' >> >> > However, regarding the former comment -- shouldn't we force generic >> > address space for all 'tree' types read in via LTO streaming for >> > offloading compilation? I assume that (in the general case) address >> > spaces are never compatible between host and offloading compilation? >> > For [...] "Add 'libgomp.c/address-space-1.c'", propagating the >> > '__seg_fs' address space across the offloading boundary (assuming I did >> > interpret the dumps correctly) doesn't seem to cause any problems >> >> As I found later, actually the 'address-space-1' per host '__seg_fs' does >> cause the "Intel MIC (emulated) offloading execution failure" >> mentioned/XFAILed for 'libgomp.c/address-space-1.c': SIGSEGV, like >> (expected) for host execution. For GCN offloading target, it maps to >> GCN 'ADDR_SPACE_FLAT' which apparently doesn't cause any ill effects (for >> that simple test case). The nvptx offloading target doesn't consider >> address spaces at all. >> >> Is the attached "Host and offload targets have no common meaning of >> address spaces" OK to push? > I'd > say I agree that any host address-space should go away when the corresponding > data is offloaded Pushed to master branch commit 9fcc3a1dd2372deea8856c55d25337b06e201203 "Host and offload targets have no common meaning of address spaces", see attached. >> Then, is that the way to do this, or should we add in >> 'gcc/tree-streamer-out.c:pack_ts_base_value_fields': >> >> if (lto_stream_offload_p) >> gcc_assert (ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (expr))); >> >> ..., and elsewhere sanitize this for offloading compilation? Jakub's >> suggestion above, regarding 'gcc/omp-low.c:omp_build_component_ref': >> >> | I think keeping the qual addr space here is the wrong thing to do, >> | it should keep the other quals and clear the address space instead >> >> But it's not obvious to me that indeed this is the one place where this >> would need to be done? (It ought to work for >> 'libgomp.c/address-space-1.c', and any other occurrences would run into >> the 'assert', so that ought to be "fine", though?) >> >> >> And, should we have a new hook >> 'void targetm.addr_space.validate (addr_space_t as)' (better name?), >> called via 'gcc/emit-rtl.c:set_mem_attrs' (only? -- assuming this is the >> appropriate canonic function where address space use is observed?), to >> make sure that the requested 'as' is valid for the target? >> 'default_addr_space_validate' would refuse everything but >> 'ADDR_SPACE_GENERIC_P (as)'; this hook would need implementing for all >> handful of targets making use of address spaces (supposedly matching the >> logic how they call 'c_register_addr_space'?). (The closest existing >> hook seems to be 'targetm.addr_space.diagnose_usage', only defined for >> AVR, and called from "the front ends" (C only).) Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 From 9fcc3a1dd2372deea8856c55d25337b06e201203 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Tue, 24 Aug 2021 11:14:10 +0200 Subject: [PATCH] Host and offload targets have no common meaning of address spaces gcc/ * tree-streamer-out.c (pack_ts_base_value_fields): Don't pack 'TYPE_ADDR_SPACE' for offloading. * tree-streamer-in.c (unpack_ts_base_value_fields): Don't unpack 'TYPE_ADDR_SPACE' for offloading. libgomp/ * testsuite/libgomp.c/address-space-1.c: Remove 'dg-xfail-run-if' for 'offload_device_intel_mic'. --- gcc/tree-streamer-in.c | 2 ++ gcc/tree-streamer-out.c | 7 ++++++- libgomp/testsuite/libgomp.c/address-space-1.c | 4 ---- 3 files changed, 8 insertions(+), 5 deletions(-) diff --git a/gcc/tree-streamer-in.c b/gcc/tree-streamer-in.c index adaf624bda7..0d5108e36a0 100644 --- a/gcc/tree-streamer-in.c +++ b/gcc/tree-streamer-in.c @@ -146,7 +146,9 @@ unpack_ts_base_value_fields (struct bitpack_d *bp, tree expr) TYPE_REVERSE_STORAGE_ORDER (expr) = (unsigned) bp_unpack_value (bp, 1); else TYPE_SATURATING (expr) = (unsigned) bp_unpack_value (bp, 1); +#ifndef ACCEL_COMPILER TYPE_ADDR_SPACE (expr) = (unsigned) bp_unpack_value (bp, 8); +#endif } else if (TREE_CODE (expr) == BIT_FIELD_REF || TREE_CODE (expr) == MEM_REF) { diff --git a/gcc/tree-streamer-out.c b/gcc/tree-streamer-out.c index 8742bf09c6a..23d15a50670 100644 --- a/gcc/tree-streamer-out.c +++ b/gcc/tree-streamer-out.c @@ -119,7 +119,12 @@ pack_ts_base_value_fields (struct bitpack_d *bp, tree expr) bp_pack_value (bp, TYPE_REVERSE_STORAGE_ORDER (expr), 1); else bp_pack_value (bp, TYPE_SATURATING (expr), 1); - bp_pack_value (bp, TYPE_ADDR_SPACE (expr), 8); + if (lto_stream_offload_p) + /* Host and offload targets have no common meaning of address + spaces. */ + ; + else + bp_pack_value (bp, TYPE_ADDR_SPACE (expr), 8); } else if (TREE_CODE (expr) == BIT_FIELD_REF || TREE_CODE (expr) == MEM_REF) { diff --git a/libgomp/testsuite/libgomp.c/address-space-1.c b/libgomp/testsuite/libgomp.c/address-space-1.c index 6ad57deec42..39ff82c1429 100644 --- a/libgomp/testsuite/libgomp.c/address-space-1.c +++ b/libgomp/testsuite/libgomp.c/address-space-1.c @@ -3,10 +3,6 @@ /* { dg-do run { target i?86-*-* x86_64-*-* } } */ /* { dg-require-effective-target offload_device_nonshared_as } */ -/* With Intel MIC (emulated) offloading: - offload error: process on the device 0 unexpectedly exited with code 0 - { dg-xfail-run-if TODO { offload_device_intel_mic } } */ - #include int __seg_fs a; -- 2.34.1