From patchwork Tue Aug 13 05:36:41 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Prathamesh Kulkarni X-Patchwork-Id: 1971762 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=Nvidia.com header.i=@Nvidia.com header.a=rsa-sha256 header.s=selector2 header.b=IoGk/Y0O; 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 4WjgBJ0qcRz1ybZ for ; Tue, 13 Aug 2024 15:37:08 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 5201A3858D3C for ; Tue, 13 Aug 2024 05:37:05 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from NAM10-BN7-obe.outbound.protection.outlook.com (mail-bn7nam10on20604.outbound.protection.outlook.com [IPv6:2a01:111:f403:2009::604]) by sourceware.org (Postfix) with ESMTPS id 7E5A33858D20 for ; Tue, 13 Aug 2024 05:36:46 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 7E5A33858D20 Authentication-Results: sourceware.org; dmarc=fail (p=reject dis=none) header.from=nvidia.com Authentication-Results: sourceware.org; spf=fail smtp.mailfrom=nvidia.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 7E5A33858D20 Authentication-Results: server2.sourceware.org; arc=pass smtp.remote-ip=2a01:111:f403:2009::604 ARC-Seal: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1723527408; cv=pass; b=tK1HxUcEYCHKoN/perDkW4T07Y73esakbYrYwml7xmd0sCIQNvu4VdHiAqHCVI/rT38cqYHtrDDBm0Q3B78f3Nx5spo2sDImZqmMUwaGhaP5HM/MGVsbX2Nkro1h0L2fBj0EiJifI1QZqDYeSVFPbZ4gGLrFFa2Cd2hXbWVzW8Y= ARC-Message-Signature: i=2; a=rsa-sha256; d=sourceware.org; s=key; t=1723527408; c=relaxed/simple; bh=91QPGR7TgF7gN5ExXcCYnIqIjsydgm9E2G9cBWLp0Gg=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=KHe31Qa42qT/96FDSYMqgFUR8wwlEoCdFRdMJyaEO5Ls5JxWxHc4vIjqPESedauzinPtiXv6gq9PWmzAc7O+gJYRn85HI13VNAYOjHtHdMStRc5SRDk3F+KjEz1d/KxlBIrU1fxHuZzayetQ4yUv3vxKLJazrH+9o22fmFKsWWA= ARC-Authentication-Results: i=2; server2.sourceware.org ARC-Seal: i=1; a=rsa-sha256; s=arcselector10001; d=microsoft.com; cv=none; b=QP4iENdHMnByl9byyUHa+aNfhuie+KgOsUyU7PdA/qXGwYIQCirmnB1f3kB8KNd9lGr2erHSj6kzCIEOAIXRDUhnXp9cXXt2LXuBhSLMnOxCODH5l2Bt9iUsWYl2YjX1jpcm6+kMI6x9ggeyMC1+v0CVDn3f3AXAnxb5vfsqwm8nkmnDcjmaUEdtFKY5mVSIVIb8djZY7Et4dhnBJI1YVbXu1Arlff1CXlxHqbOMcjqV61CRzJcAIt57z+B7wSPsfc6I2ZSg5wQ2m5iACbLK8VYenElbQqw3Fyev19YRjESHy8Ob7v9JGbSq60vwoX3CiSMjbtj0zCUZjyWWHpvI5w== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=microsoft.com; s=arcselector10001; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-AntiSpam-MessageData-ChunkCount:X-MS-Exchange-AntiSpam-MessageData-0:X-MS-Exchange-AntiSpam-MessageData-1; bh=ucvJPZ42iRsZ9AEgDVkKwI4a42FOgXnL6GKIrzd6HtI=; b=t3h6Zhkbp8Qzd+pbSEyHYekSwPrmLY5PZsyUi3mbbgr3EtA7nkjb8jL+gtao4hScpBFs/80ysomLnvxvGCKSpgCqWhRieM5c7HjMPQntb6StX5u4q/DiSxzxXh4b/2JZeJXieTXykXe2rA1DOU1MCrroEehM+psoKq5xrdvgGbNsBS9/gHjkQYQeF5wAVW/xnnczW439FpbqnhZK+GJ+a3P9oG07JvIS+3fnChzGbzeL3iex7QVP70+4rZ86kWRSNs4oOyLCcLuCkCAu0SdZwcHunx/QC/NEJW1OMXbjhZshf5Xy3MjoqRfUJoCzYIvfEhAKpQXR+6sbI7Ay6EXzMA== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass smtp.mailfrom=nvidia.com; dmarc=pass action=none header.from=nvidia.com; dkim=pass header.d=nvidia.com; arc=none DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=Nvidia.com; s=selector2; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-SenderADCheck; bh=ucvJPZ42iRsZ9AEgDVkKwI4a42FOgXnL6GKIrzd6HtI=; b=IoGk/Y0O+qvAay4K0qLJqqXKiBxsx28cozM+eHuc+N6HxsUyqJBk9cI7XD3l1Gswa4FcUti0HMeAEBYNHhUv1pdsdTM+7wbd/b+6ewEJ2NYsaUZbFcTwqPLamGGFruRDjTsOh4nMXXu7E0g1ebnAgiAfWxpCyhvKyTZFazVlr1+EUY7VpFlKNyqqDVzTNx/bfjANgJo2HAKR9H6U0yGONQUGGXNQxPHdaQnjEfM0mNm4tucskd2dKpzd3ZnMNnsMeVquHxflPfbzMe/whPO2wDKZ+UuIrWCsk8yyroS6AvSd2y1B0YtHN0XKmMc1tsOnrV1w+76mZhRUg1mjA/8R1A== Received: from DS0PR12MB9037.namprd12.prod.outlook.com (2603:10b6:8:f1::13) by MW3PR12MB4409.namprd12.prod.outlook.com (2603:10b6:303:2d::23) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.7849.22; Tue, 13 Aug 2024 05:36:42 +0000 Received: from DS0PR12MB9037.namprd12.prod.outlook.com ([fe80::d839:f55f:dcaa:db06]) by DS0PR12MB9037.namprd12.prod.outlook.com ([fe80::d839:f55f:dcaa:db06%3]) with mapi id 15.20.7849.019; Tue, 13 Aug 2024 05:36:42 +0000 From: Prathamesh Kulkarni To: "gcc-patches@gcc.gnu.org" , "rguenther@suse.de" , Andrew Pinski , Thomas Schwinge Subject: [optc-save-gen.awk] Fix streaming of command line options for offloading Thread-Topic: [optc-save-gen.awk] Fix streaming of command line options for offloading Thread-Index: AdrtIGMe/bAyPaafQmyA3d0jVuyImA== Date: Tue, 13 Aug 2024 05:36:41 +0000 Message-ID: Accept-Language: en-US Content-Language: en-US X-MS-Has-Attach: yes X-MS-TNEF-Correlator: authentication-results: dkim=none (message not signed) header.d=none;dmarc=none action=none header.from=nvidia.com; x-ms-publictraffictype: Email x-ms-traffictypediagnostic: DS0PR12MB9037:EE_|MW3PR12MB4409:EE_ x-ms-office365-filtering-correlation-id: fae79711-dcf9-48de-24ed-08dcbb59e945 x-ms-exchange-senderadcheck: 1 x-ms-exchange-antispam-relay: 0 x-microsoft-antispam: BCL:0; ARA:13230040|366016|376014|1800799024|38070700018; x-microsoft-antispam-message-info: OP67QsGW2GFzMXCcrKPOIE/IPVO5b9qF73K6VG7xHXr1BJz72uElcSJDwvnWz9niqOrEBjb+ZFfohxhaCfy2tiUVlWXxl4oIgHANvC2jvN2z9NhikBQxkNI6WZtuPh2KhoESexjyymdDKY8iPnovqrATkh59t1GW06sAK1eXkL5DRSt/faEvF/n1bCyR3xFvcVYYyks/s4Q3SjZRVCepRaDXKx5U4yDjQ49riOTLxQDCZinkk5lYaaPYV+GXewtT0zi7kbMoAesJKQFlql3s16j/VjeHcMHarYMFvJnlQgwlMbMR/7v9zhXCYpWu1gbIYXCM7wAXK5rUHNgJeD5c3HbLA4N4uUwQBQL6w4cVFy8HHPPsbXP40MgrMDb/uAjId1IxvVqVq+ajM3YcegVJ5V1DF/mE52dX6rEZ9YzhemVGs3ttAT81eO3GhJh3qQmE3ochCoVxnMGIzQnV9uR71MCQUbhJb6xn8LFl5oX3yVIy4YUn4RIRc4dfrzb9r27CvU+ayU4hw/nddoksNa1dwzAeFrfsKrzy4IVCmWzGsrWOskAe8L0pUnja/JV7aF0shNtOIgYgvUOGD7eqMXu+QMLPsgDKQ4pLQn+rT9CjtHUpoPS5XFYy4m0IUtpt5xcuixcapLjdtY7qPm/jKfZf7GyuK5T7oVAZSL3325l1zmMkIIr/Vc/kX6Lnq8XCesZLfzvvAly919Vy79ZmTIQOdX2hTA6KYnADkLOH9kea2jvL7AonpfQ/bnhCQowpDLPHIaiTT+/f7Kt1U5IapOoYx6loYbKIK7pmt/nYbHm740X74UVsr9fJfncXQtZwlnYH2KvlcH5M9KgFo8q+E+w82lY/wikzd4/5CxRKyKee3KzjYsh/VabugHsQXWuIfIXWV8LyEZ7/QVbN2+j3gt00qOjNB8VoWvCOhmA96qpAAnwwt6hKnPiVCJQBUdt1c/kSdPa/x1D0WVqmASYhKuxN84Q4VVU5DVzYvcyyOgmDXxozoRDM448Z8XIaXR1KAh0EWnaHzBv4n+hdNlWjofeEVfPRwVjYeKuOfUoPoseULQKtusX7iOWGDnU+7bTtF5PFfOgoA9mSkQ1juWXbxNpWvoo58fXt7NElmxQ2V77A5FLGJMIvp6RS3UreQOpSv0uTESdjU6x4b09tP7HUTdiyDMuL94ssqKDtkwK38Xm/EhsTgxJyK6OdurowZUhGEV6MO0z1cjDdfsxiFow71Uxz3ki4T+s3i1HudDeWpawV/NoYi7l3Cgh8utGYAslUP6ui8biwcG2+toYptfZ2MF4BJfvg1OEztsiX42IVe6l6707i1n0D4KK6zD5I9DpLejkCD5lRYXuhITG3ECzF8ApQt/XXUdhDl5m1jlJINRmxo6C28xwi+2aWPjmgUVv2Q+e3DgsV+Lcwd9aLzUdwNDW1Sg== x-forefront-antispam-report: CIP:255.255.255.255; CTRY:; LANG:en; SCL:1; SRV:; IPV:NLI; SFV:NSPM; H:DS0PR12MB9037.namprd12.prod.outlook.com; PTR:; CAT:NONE; SFS:(13230040)(366016)(376014)(1800799024)(38070700018); DIR:OUT; SFP:1101; x-ms-exchange-antispam-messagedata-chunkcount: 1 x-ms-exchange-antispam-messagedata-0: Jy4fL4inpEvP9oXiP2Xas0ObjU50cknyHodX/B03QDAxkrTsaFRmSrAmijEoiAZH7nExBAHQv27eBlfZ2vgQLkhJLaec/sz1OtGuC/mSgeNfePFoCs3Uy1dKgWaZzzKeEGHagMLmoXdUmEHZg3He9vFXgNT7x8REYx7bwLNlzr6Lu+BCc/4T/xCVfcxnz2uvJ2oOv8VgQ943V7gevJabtpEz8kLgck4GiklvfBVtHJ9o8aVeJ9pXAL4JabiP/VgmLoFeOQee5aPS8ltvVzcFaBGw3vXInI47n0P2GCu4SP9Xqx2m6aDeQbl63n1d+iFaBF7Zng0VZb6p4HCuFdhb3rxr/3S9CCQVsTI91vT/UtqsGgwTfO4AQl7ExbTzEjsCZpSmw91+mFyHKCbA/LuP9LsS4mrl+0XFKoIyWz4eu8uHrp6aBtwWueOnOY+1/gq/Nflbuc/ubeElC8G4yPjCwsEtvnip4R2tQdmst5fFT2FFYpmK121GHGfuItnVuWykach55CAboAfM0RBU8CjZd5AbcH1WRpQUQs4O3+ihcST8RBSdUhpdw1+y305LCX5xnAyzJPeC70c8V6+q28ELOYdoGzh5vNZgHfBdqMEyS+OjfwAALIT2AOLVvCjeKAGPkht8r7iIuXaJkjRdxa7JLDLL/B44Z7Rj26m5y2jjspue9fWkQYU3O9HP5voD8IdpEGqUBu+Q5HwGC7hA2wDwKkC3460sfR+Dq/HHY+U5MirI0BssZ1eqyMUNdgn4l6spsrjTqplZ/h9yS5o8ynUe43pU3nDX/ds51b+CBGfWHCfCclEsU+/jzR9Y3RY+HJyjRaKDth4COzZyQjlfaCAEb1RzB3wbWixWFSslgJ+Tq6ScyYK2e1URmuVYBSlhjoxtJ8Fy/FtK+8J6KBVvrP0bCFhLKT3kgFYWSIsSHkjko6I1hJFEKX8NtpMI8zLVV65TOjDh/a0wsPlDSjHfgno/tt1cX25FyAi0eJFk/q2FmE00Y96oSpU5z6PD7nLq8P6p+f7jN9xWUi+PVd2MSTaLpklmmdWTy8HcXTeD4I55agauK/npu0MZECBuVaD0+DOfNEzf7gWseiCKwFIcvr9oriqcSD2pHBBpIuyXaQSJxccOf8Z25MLmssutKFyCZiydPRQa8Viivd3WhTx60p7UDEKgLo86nFJabhArlDd8qZNlCz+O1yxKmHM7Bnbu4R9C6jC1/wWcqoY+sbtBJsYHwbHAOJKUb2aNRRZLLC/ugnti2xFnYlxO5ETDniDsXDVgq7FIXPpnIPStdKcE38Rb1yJSTkYz8J/O+7KmlvJ4Lw2aeuJOcz/bbYsMOK5kVUAD325kMlqMckNkIPqvKl+aX+fJ2nhSjhfJTfBZb2kY9kD+d80YdyyWf0T4KfRNQ3W4zpLeHngeYEuLFAarDIz5kGBYLDWI7LdAxzhWihw5kkpHBoEj0oLr6qKeqQsu1Ok55LY4rzF5wxzFPHSxm+vR/4K4q7ueKKORm3BlRQPV6mSR5xjL5g44eaiT1y46r1R1Eay/7qEIsK6A5PrN5PsI2hst6onbqQoo3uS9ZaYPXkMVVB4h0UzZOGc7NgDOWKoK MIME-Version: 1.0 X-OriginatorOrg: Nvidia.com X-MS-Exchange-CrossTenant-AuthAs: Internal X-MS-Exchange-CrossTenant-AuthSource: DS0PR12MB9037.namprd12.prod.outlook.com X-MS-Exchange-CrossTenant-Network-Message-Id: fae79711-dcf9-48de-24ed-08dcbb59e945 X-MS-Exchange-CrossTenant-originalarrivaltime: 13 Aug 2024 05:36:41.9647 (UTC) X-MS-Exchange-CrossTenant-fromentityheader: Hosted X-MS-Exchange-CrossTenant-id: 43083d15-7273-40c1-b7db-39efd9ccc17a X-MS-Exchange-CrossTenant-mailboxtype: HOSTED X-MS-Exchange-CrossTenant-userprincipalname: QjjoV5Yql36j1kg4dIPdufVtbKkU16Wmwe0uacR52nGg9HGE+OBfaO8ufNk915Hws1YMQLgQo5ea/oAMN03oIA== X-MS-Exchange-Transport-CrossTenantHeadersStamped: MW3PR12MB4409 X-Spam-Status: No, score=-11.3 required=5.0 tests=BAYES_00, DKIMWL_WL_HIGH, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, FORGED_SPF_HELO, GIT_PATCH_0, KAM_SHORT, SPF_HELO_PASS, SPF_NONE, 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, As mentioned in: https://gcc.gnu.org/pipermail/gcc/2024-August/244581.html AArch64 cl_optimization_stream_out streams out target-specific optimization options like flag_aarch64_early_ldp_fusion, aarch64_early_ra etc, which breaks AArch64/nvptx offloading, since nvptx cl_optimization_stream_in doesn't have corresponding stream-in for these options and ends up setting invalid values for ptr->explicit_mask (and subsequent data structures). This makes even a trivial test like the following to cause ICE in lto_read_decls with -O3 -fopenmp -foffload=nvptx-none: int main() { int x; #pragma omp target map(x) x; } The attached patch modifies optc-save-gen.awk to generate if (!lto_stream_offload_p) check before streaming out target-specific opt in cl_optimization_stream_out, which fixes the issue. cl_optimization_stream_out after patch (last few entries): bp_pack_var_len_int (bp, ptr->x_flag_wrapv_pointer); bp_pack_var_len_int (bp, ptr->x_debug_nonbind_markers_p); if (!lto_stream_offload_p) bp_pack_var_len_int (bp, ptr->x_flag_aarch64_early_ldp_fusion); if (!lto_stream_offload_p) bp_pack_var_len_int (bp, ptr->x_aarch64_early_ra); if (!lto_stream_offload_p) bp_pack_var_len_int (bp, ptr->x_flag_aarch64_late_ldp_fusion); if (!lto_stream_offload_p) bp_pack_var_len_int (bp, ptr->x_flag_mlow_precision_div); if (!lto_stream_offload_p) bp_pack_var_len_int (bp, ptr->x_flag_mrecip_low_precision_sqrt); if (!lto_stream_offload_p) bp_pack_var_len_int (bp, ptr->x_flag_mlow_precision_sqrt); for (size_t i = 0; i < ARRAY_SIZE (ptr->explicit_mask); i++) bp_pack_value (bp, ptr->explicit_mask[i], 64); For target-specific options, streaming out is gated on !lto_stream_offload_p check. The patch also fixes failures due to same issue with x86_64->nvptx offloading for target-print-1.f90 (and couple more). Does the patch look OK ? Signed-off-by: Prathamesh Kulkarni Thanks, Prathamesh [optc-save-gen.awk] Fix streaming of command line options for offloading. The patch modifies optc-save-gen.awk to generate if (!lto_stream_offload_p) check before streaming out target-specific opt in cl_optimization_stream_out, when offloading is enabled. gcc/ChangeLog: * gcc/optc-save-gen.awk: New array var_target_opt. Use it to generate if (!lto_stream_offload_p) check in cl_optimization_stream_out. Signed-off-by: Prathamesh Kulkarni Signed-off-by: Prathamesh Kulkarni Signed-off-by: Prathamesh Kulkarni diff --git a/gcc/optc-save-gen.awk b/gcc/optc-save-gen.awk index a3af88e3776..228efe2accd 100644 --- a/gcc/optc-save-gen.awk +++ b/gcc/optc-save-gen.awk @@ -1307,6 +1307,11 @@ for (i = 0; i < n_opts; i++) { var_opt_optimize_init[n_opt_val] = init; } + # Mark options that are annotated with both Optimization and + # Target so we can avoid streaming out target-specifc opts when + # offloading is enabled. + if (flag_set_p("Target", flags[i])) + var_target_opt[n_opt_val] = 1; n_opt_val++; } } @@ -1384,6 +1389,10 @@ for (i = 0; i < n_opt_val; i++) { } else { sgn = "int"; } + # Do not stream out target-specifc opts if offloading is + # enabled. + if (var_target_opt[i]) + print " if (!lto_stream_offload_p)" # If applicable, encode the streamed value. if (var_opt_optimize_init[i]) { print " if (" var_opt_optimize_init[i] " > (" var_opt_val_type[i] ") 10)";