From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 0B91B3858D1E; Fri, 11 Mar 2022 14:13:03 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 0B91B3858D1E 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.90,174,1643702400"; d="scan'208";a="72870876" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 11 Mar 2022 06:13:03 -0800 IronPort-SDR: +8oVez/jUftMXI3qQKWF6h4wXK8mhkGCOuoSjUIaccsFfYi+N6ROLeHvG2YU52r85QgkgE4kDI n33LU/Z3EMBi4ycYguXjbHRe8s21ighpDM/ppisk7oD8rRAxv7slay4prM8zU35NYg5HI5TEqN KD4xM5R0NzlXCc4bWNHimCaYxtbnAu5Jgl+EkXOLpxsez3GfM4Xm2x+rHsgMW4dtWCqebYZli8 Mesloi14twC55YWvWMDw6XQ9cfbg0sdKFQAj+/PEPb0LVHJ1ioGsKDNCJM0zu0MGAhvaboJgkE PtM= Message-ID: <1b53e970-a0a7-66bc-4b2e-828e881cce73@codesourcery.com> Date: Fri, 11 Mar 2022 15:12:32 +0100 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:91.0) Gecko/20100101 Thunderbird/91.7.0 Subject: Re: [Patch] OpenMP, libgomp: Add new runtime routine omp_target_is_accessible. Content-Language: en-US To: Marcel Vollweiler , , References: <7fa4a70c-60e7-fa18-0fcd-98301c0b3344@codesourcery.com> From: Tobias Burnus In-Reply-To: <7fa4a70c-60e7-fa18-0fcd-98301c0b3344@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-11.mgc.mentorg.com (139.181.222.11) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-5.5 required=5.0 tests=BAYES_00, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, NICE_REPLY_A, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=no 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: Fri, 11 Mar 2022 14:13:05 -0000 Minor remark to the test: On 11.03.22 13:30, Marcel Vollweiler wrote: > + int d =3D omp_get_default_device (); ... > + int shared_mem =3D 0; > + #pragma omp target map (alloc: shared_mem) device (d) > + shared_mem =3D 1; > + if (omp_target_is_accessible (p, sizeof (int), d) !=3D shared_mem) > + __builtin_abort (); I wonder whether it makes sense to do instead for (d =3D 0; d <=3D omp_get_num_devices(); ++d) instead of just d =3D omp_get_default_device(); given that we have already found once in a while bugs when testing more than just the default device - be it because devices differed or because '0' was special. In particular, I could image having at the same time two or three devices available of type intelmic + gcn + nvptx, possibly mixing shared memory, nonshared memory and semi-shared memory* Tobias (* semi-shared: I am especially thinking of nvptx with %dynamic_smem_size, which requires some special handling. By contrast with HMM and Pascal GPUs, real USM is possible.) ----------------- 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