From patchwork Fri Mar 11 12:30:14 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Marcel Vollweiler X-Patchwork-Id: 1604342 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=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (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 bilbo.ozlabs.org (Postfix) with ESMTPS id 4KFQJK60Gzz9sGD for ; Fri, 11 Mar 2022 23:30:40 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 194143857820 for ; Fri, 11 Mar 2022 12:30:37 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id EC2D33858C27; Fri, 11 Mar 2022 12:30:23 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org EC2D33858C27 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,173,1643702400"; d="diff'?scan'208";a="75616038" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 11 Mar 2022 04:30:22 -0800 IronPort-SDR: AvecZ8k7jBSs80p4WTTNP0Y5Ziyh+hdX96xEgViys8mF8g8uuH9XMQGUITInvq3yWOPcgnji8S bWLUNT6KkOkUxfVIZxNaegpZFxrvbfFQGm4aH2NEaR/GvIiHoG0ovq5BCj7Z3nwgPYwTBfVsg6 sYgpwm94tR3Lmq7eghVzDdCP+stOnkh/cRSn4Mug1zf9qDI3S2UL3mRKV3Vb478/+ZU+oLMrmc iPhv9ACMftuR0iHDnSn3FWLgxOAS7DlsbkyXl8BP4aNsLCqLvueI5D1utj2HovSZ4DQs/7iSta VLs= Message-ID: <7fa4a70c-60e7-fa18-0fcd-98301c0b3344@codesourcery.com> Date: Fri, 11 Mar 2022 13:30:14 +0100 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (Windows NT 10.0; Win64; x64; rv:91.0) Gecko/20100101 Thunderbird/91.6.1 From: Marcel Vollweiler Subject: [Patch] OpenMP, libgomp: Add new runtime routine omp_target_is_accessible. To: , X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE 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+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Hi, This patch adds the OpenMP runtime routine "omp_target_is_accessible" which was introduced in OpenMP 5.1 (specification section 3.8.4): "The omp_target_is_accessible routine tests whether host memory is accessible from a given device." "This routine returns true if the storage of size bytes starting at the address given by ptr is accessible from device device_num. Otherwise, it returns false." "The value of ptr must be a valid host pointer or NULL (or C_NULL_PTR, for Fortran). The device_num argument must be greater than or equal to zero and less than or equal to the result of omp_get_num_devices()." "When called from within a target region the effect is unspecified." Currently, the only way of accessing host memory on a non-host device is via shared memory. This will change with unified shared memory (usm) that was recently submitted but not yet approved/committed. A follow-up patch for omp_target_is_accessible is planned considering usm when available. The current patch handles the basic implementation for C/C++ and Fortran and includes comments pointing to usm. Although not explicitly specified in the OpenMP 5.1 standard, the implemented function returns "true" if the given device_num is equal to "omp_get_num_devices" (i.e. the host) as it is expected that host memory can be accessed from the host device. The patch was tested on x86_64-linux and PowerPC, both with nvptx offloading. All with no regressions. Marcel ----------------- 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 OpenMP, libgomp: Add new runtime routine omp_target_is_accessible. gcc/ChangeLog: * omp-low.cc (omp_runtime_api_call): Added target_is_accessible to omp_runtime_apis array. libgomp/ChangeLog: * libgomp.map: Added omp_target_is_accessible. * libgomp.texi: Tagged omp_target_is_accessible as supported. * omp.h.in: Added omp_target_is_accessible. * omp_lib.f90.in: Added interface for omp_target_is_accessible. * omp_lib.h.in: Likewise. * target.c (omp_target_is_accessible): Added implementation of omp_target_is_accessible. * testsuite/libgomp.c-c++-common/target-is-accessible-1.c: New test. * testsuite/libgomp.fortran/target-is-accessible-1.f90: New test. diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 77176ef..bf38fad 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -3959,6 +3959,7 @@ omp_runtime_api_call (const_tree fndecl) "target_associate_ptr", "target_disassociate_ptr", "target_free", + "target_is_accessible", "target_is_present", "target_memcpy", "target_memcpy_rect", diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index 2ac5809..1764380 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -226,6 +226,11 @@ OMP_5.1 { omp_get_teams_thread_limit_; } OMP_5.0.2; +OMP_5.1.1 { + global: + omp_target_is_accessible; +} OMP_5.1; + GOMP_1.0 { global: GOMP_atomic_end; diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 161a423..58e432c 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -311,7 +311,7 @@ The OpenMP 4.5 specification is fully supported. @item @code{omp_set_num_teams}, @code{omp_set_teams_thread_limit}, @code{omp_get_max_teams}, @code{omp_get_teams_thread_limit} runtime routines @tab Y @tab -@item @code{omp_target_is_accessible} runtime routine @tab N @tab +@item @code{omp_target_is_accessible} runtime routine @tab Y @tab @item @code{omp_target_memcpy_async} and @code{omp_target_memcpy_rect_async} runtime routines @tab N @tab @item @code{omp_get_mapped_ptr} runtime routine @tab N @tab diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in index 89c5d65..1ec7415 100644 --- a/libgomp/omp.h.in +++ b/libgomp/omp.h.in @@ -282,6 +282,8 @@ extern int omp_target_memcpy_rect (void *, const void *, __SIZE_TYPE__, int, extern int omp_target_associate_ptr (const void *, const void *, __SIZE_TYPE__, __SIZE_TYPE__, int) __GOMP_NOTHROW; extern int omp_target_disassociate_ptr (const void *, int) __GOMP_NOTHROW; +extern int omp_target_is_accessible (const void *, __SIZE_TYPE__, int) + __GOMP_NOTHROW; extern void omp_set_affinity_format (const char *) __GOMP_NOTHROW; extern __SIZE_TYPE__ omp_get_affinity_format (char *, __SIZE_TYPE__) diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in index daf40dc..f369507 100644 --- a/libgomp/omp_lib.f90.in +++ b/libgomp/omp_lib.f90.in @@ -835,6 +835,16 @@ end function omp_target_disassociate_ptr end interface + interface + function omp_target_is_accessible (ptr, size, device_num) bind(c) + use, intrinsic :: iso_c_binding, only : c_ptr, c_size_t, c_int + integer(c_int) :: omp_target_is_accessible + type(c_ptr), value :: ptr + integer(c_size_t), value :: size + integer(c_int), value :: device_num + end function omp_target_is_accessible + end interface + #if _OPENMP >= 201811 !GCC$ ATTRIBUTES DEPRECATED :: omp_get_nested, omp_set_nested #endif diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in index ff857a4..5ea0366 100644 --- a/libgomp/omp_lib.h.in +++ b/libgomp/omp_lib.h.in @@ -416,3 +416,14 @@ integer(c_int), value :: device_num end function omp_target_disassociate_ptr end interface + + interface + function omp_target_is_accessible (ptr, size, device_num) & + & bind(c) + use, intrinsic :: iso_c_binding, only : c_ptr, c_size_t, c_int + integer(c_int) :: omp_target_is_accessible + type(c_ptr), value :: ptr + integer(c_size_t), value :: size + integer(c_int), value :: device_num + end function omp_target_is_accessible + end interface diff --git a/libgomp/target.c b/libgomp/target.c index 9017458..01d36d1 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -3666,6 +3666,24 @@ omp_target_disassociate_ptr (const void *ptr, int device_num) } int +omp_target_is_accessible (const void *ptr, size_t size, int device_num) +{ + if (device_num < 0 || device_num > gomp_get_num_devices ()) + return false; + + if (device_num == gomp_get_num_devices ()) + return true; + + struct gomp_device_descr *devicep = resolve_device (device_num); + if (devicep == NULL) + return false; + + /* TODO: Unified shared memory must be handled when available. */ + + return devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM; +} + +int omp_pause_resource (omp_pause_resource_t kind, int device_num) { (void) kind; diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c new file mode 100644 index 0000000..6788894 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c @@ -0,0 +1,42 @@ +#include + +int +main () +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int n = omp_get_num_devices (); + void *p; + + if (d < 0 || d >= n) + d = id; + + if (!omp_target_is_accessible (p, sizeof (int), n)) + __builtin_abort (); + + if (!omp_target_is_accessible (p, sizeof (int), id)) + __builtin_abort (); + + if (omp_target_is_accessible (p, sizeof (int), -1)) + __builtin_abort (); + + if (omp_target_is_accessible (p, sizeof (int), n + 1)) + __builtin_abort (); + + /* Currently, a host pointer is accessible if the device supports shared + memory or omp_target_is_accessible is executed on the host. This + test case must be adapted when unified shared memory is avialable. */ + int shared_mem = 0; + #pragma omp target map (alloc: shared_mem) device (d) + shared_mem = 1; + if (omp_target_is_accessible (p, sizeof (int), d) != shared_mem) + __builtin_abort (); + + int a[128]; + p = &a; + for (int i = 0; i < 128; i++) + if (omp_target_is_accessible (p, 128 * sizeof (int), d) != shared_mem) + __builtin_abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 new file mode 100644 index 0000000..eae696f --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 @@ -0,0 +1,44 @@ +program main + use omp_lib + use iso_c_binding + implicit none (external, type) + integer :: d, id, n, shared_mem, i + integer, target :: a(1:128) + type(c_ptr) :: p + + d = omp_get_default_device () + id = omp_get_initial_device () + n = omp_get_num_devices () + + if (d < 0 .or. d >= n) & + d = id + + if (omp_target_is_accessible (p, c_sizeof (d), n) /= 1) & + stop 1 + + if (omp_target_is_accessible (p, c_sizeof (d), id) /= 1) & + stop 2 + + if (omp_target_is_accessible (p, c_sizeof (d), -1) /= 0) & + stop 3 + + if (omp_target_is_accessible (p, c_sizeof (d), n + 1) /= 0) & + stop 4 + + ! Currently, a host pointer is accessible if the device supports shared + ! memory or omp_target_is_accessible is executed on the host. This + ! test case must be adapted when unified shared memory is avialable. + shared_mem = 0; + !$omp target map (alloc: shared_mem) device (d) + shared_mem = 1; + !$omp end target + if (omp_target_is_accessible (p, c_sizeof (d), d) /= shared_mem) & + stop 5; + + p = c_loc (a) + do i = 1, 128 + if (omp_target_is_accessible (p, 128 * sizeof (d), d) /= shared_mem) & + stop 6; + end do + +end program main