From patchwork Tue Aug 24 10:23:07 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1520221 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from 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 RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4Gv4vm66hVz9sW5 for ; Tue, 24 Aug 2021 20:23:47 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 80D393858418 for ; Tue, 24 Aug 2021 10:23:44 +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 DF2A13858409 for ; Tue, 24 Aug 2021 10:23:20 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org DF2A13858409 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: BPynUPgMy42XYU8fqGRdX+GaGXTRQ5YokpNrBoRtyUWPEDZd9MBzSjL3HKe76Rh9gVOYLSkJZ7 3qlQnJxbTlqE2v3PZiZ7LpHOZRFjfjsH42KYV2XDeGhtGxzWMHasZ9TsDymMIaTVYKEEZuuUNL So4c58AP2l2pUfioHhoNg1VVtuuiGF9OQgOrTHh+H4NrPsKH1AZIEoR0KDlsE8KCE0DHOnQuGi WYXBnNZDkFoVDg2Vcee/+Fq9NbEEXjk2K+TCEf07af1/DzDHO22SjfHF5ELXpeZPMR55ojRs5u kEcVU866Vd9w7kn+uRwbNqSP X-IronPort-AV: E=Sophos;i="5.84,346,1620720000"; d="scan'208,223";a="65046495" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 24 Aug 2021 02:23:19 -0800 IronPort-SDR: h1JNs82IpeFNSNahg3FabMwk27HHx7rUpYqS2zLH6uUiYq26//5hgBDpoQQDf2m2sHEbVum+xp 7psr7J+5Or5N/tRodssa1n/b9bE2ot0pMa5BWAIgUn1q/aMcjTJdLLlRm0/ZXBQa1mwVq4VsvI zsUl/S6vV36FpD2sxfBxO3u4+oRrlXRiqrn3tByxHLU9GRLJmVded0+4JMAZG8PXO0t6SAWecp Vf2A4rOkDJ9TqCucirksuSGAfYuuNZrn5m6q4vTSLeNG39wQZ0uP9lwwxAhrxaiO0aYFh98Gfq uLU= From: Thomas Schwinge To: Jakub Jelinek , Richard Biener , Subject: Host and offload targets have no common meaning of address spaces (was: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref') In-Reply-To: <871r6pnqez.fsf@euler.schwinge.homeip.net> 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> User-Agent: Notmuch/0.29.1+93~g67ed7df (https://notmuchmail.org) Emacs/26.3 (x86_64-pc-linux-gnu) Date: Tue, 24 Aug 2021 12:23:07 +0200 Message-ID: <87o89np2es.fsf@dem-tschwing-1.ger.mentorg.com> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-09.mgc.mentorg.com (139.181.222.9) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.2 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: , Cc: Julian Brown , Andrew Stubbs Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Hi! 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', not > 'obj_type', so indeed Kwok's new code is a no-op: > > (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? 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 e01e06bd17bf2c7cb182d30bed02babc5edfa183 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 | 4 +++- libgomp/testsuite/libgomp.c/address-space-1.c | 4 ---- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/gcc/tree-streamer-in.c b/gcc/tree-streamer-in.c index e0522bf2ac1..acdc48ef09f 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 855d1cd59b9..aac0b7ecf54 100644 --- a/gcc/tree-streamer-out.c +++ b/gcc/tree-streamer-out.c @@ -119,7 +119,9 @@ 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); + /* Host and offload targets have no common meaning of address spaces. */ + if (!lto_stream_offload_p) + 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.25.1