From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id CA39738582B0 for ; Mon, 17 Oct 2022 07:35:53 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org CA39738582B0 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.95,190,1661846400"; d="scan'208";a="87666627" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 16 Oct 2022 23:35:53 -0800 IronPort-SDR: +DmHTvXQtWYicZxNFAl173nkWzcG4QrKLVcZzbs7wetgy6AiuC2fKD6hhkh3C0lgM4WcHHg3M7 aeUVTj40pBipRZW7FUbB5anWKFeviXtUNKJdgCXTBIdVaawy575t8HWhy51QiNfxc5CI+b+1gb npQV78hF5zpSKyaczMCnKgSUz1GcBjy78rARtPATzNCuoDICriIuh6g4xs4LNIzTBjBF5QRiNy aYM+21XtCS0L1HbbD0JCPviaSKOJG26EK95BujuBmuVu3/B9wa3NWiAuFdu73zYixMVG+rmlDo gu4= Message-ID: <6d4a11c0-9a2e-6899-dbdd-97c7ce6f83b9@codesourcery.com> Date: Mon, 17 Oct 2022 09:35:43 +0200 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:102.0) Gecko/20100101 Thunderbird/102.3.3 Subject: *ping* / Re: [Patch][v5] libgomp/nvptx: Prepare for reverse-offload callback handling Content-Language: en-US From: Tobias Burnus To: Alexander Monakov , Jakub Jelinek , gcc-patches References: <57b3ae5e-8f15-8bea-fa09-39bccbaa2414@codesourcery.com> <3f0fc49f-b07f-bee2-51a8-a5d03f1c33ed@codesourcery.com> <30e3ed49-0d14-8015-57ef-3d70b1dea69a@codesourcery.com> <3ebce406-46e4-8f98-8c53-83b61423644e@codesourcery.com> <798d7ee1-2ffa-a591-38cb-a9ad421265d0@codesourcery.com> <832946f-bb12-23d7-7d64-47b85c95125@ispras.ru> <1e56b27f-3c74-8bc2-028b-9091fa1ae7b1@codesourcery.com> In-Reply-To: <1e56b27f-3c74-8bc2-028b-9091fa1ae7b1@codesourcery.com> Content-Type: text/plain; charset="UTF-8"; format=flowed Content-Transfer-Encoding: quoted-printable X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-08.mgc.mentorg.com (139.181.222.8) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-5.4 required=5.0 tests=BAYES_00,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,SPF_HELO_PASS,SPF_PASS,TXREP autolearn=no autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: On 12.10.22 10:55, Tobias Burnus wrote: > On 11.10.22 13:12, Alexander Monakov wrote: >> My understanding is such trickery should not be necessary with >> the barrier-based approach, i.e. the sequence of PTX instructions >> >> st % plain store >> membar.sys >> st.volatile >> >> should be enough to guarantee that the former store is visible on the >> host >> before the latter, and work all the way back to sm_20. > > If I understand it correctly, you mean: > > GOMP_REV_OFFLOAD_VAR->dev_num =3D GOMP_ADDITIONAL_ICVS.device_num; > > __sync_synchronize (); /* membar.sys */ > asm volatile ("st.volatile.global.u64 [%0], %1;" > : : "r"(addr_struct_fn), "r" (fn) : "memory"); > > > And then directly followed by the busy wait: > > while (__atomic_load_n (&GOMP_REV_OFFLOAD_VAR->fn, __ATOMIC_ACQUIRE) > !=3D 0) > ; /* spin */ > > which GCC expands to: > > /* ld.global.u64 %r64,[__gomp_rev_offload_var]; > ld.u64 %r36,[%r64]; > membar.sys; */ > > The such updated patch is attached. > > (This is the only change + removing the mkoffload.cc part is the only > larger change. Otherwise, it only handles the minor comments by Jakub. > The now removed CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT was used > until commit r10-304-g1f4c5b9bb2eb81880e2bc725435d596fcd2bdfef i.e. > it is a really old left over!) > > Otherwise, tested* to work with sm_30 (error by mkoffload, unchanged), > sm_35 and sm_70. > > Tobias > > *With some added code; until GOMP_OFFLOAD_get_num_devices accepts > GOMP_REQUIRES_UNIFIED_SHARED_MEMORY and GOMP_OFFLOAD_load_image > gets passed a non-NULL for rev_fn_table, the current patch is a no op. > > Planned next is the related GCN patch =E2=80=93 and the actual change > in libgomp/target.c (+ accepting USM in GOMP_OFFLOAD_get_num_devices) ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstra=C3=9Fe 201= , 80634 M=C3=BCnchen; Gesellschaft mit beschr=C3=A4nkter Haftung; Gesch=C3= =A4ftsf=C3=BChrer: Thomas Heurung, Frank Th=C3=BCrauf; Sitz der Gesellschaf= t: M=C3=BCnchen; Registergericht M=C3=BCnchen, HRB 106955