From patchwork Fri May 6 11:19:55 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Marcel Vollweiler X-Patchwork-Id: 1627521 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 (2048 bits) server-digest SHA256) (No client certificate requested) by bilbo.ozlabs.org (Postfix) with ESMTPS id 4Kvp5R0jzqz9sGJ for ; Fri, 6 May 2022 21:20:25 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 7AD8A3952033 for ; Fri, 6 May 2022 11:20:21 +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 B6078394D8A1; Fri, 6 May 2022 11:20:01 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org B6078394D8A1 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.91,203,1647331200"; d="diff'?scan'208";a="78040220" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 06 May 2022 03:20:00 -0800 IronPort-SDR: TgIWORaF4CxKjb9/I+mDl7fVhDopkGlBYpP/t13lMnNGNAZcQPM5zi6Qa7DOlkqShln0JIkTMD tDTcC42++QHRUq2VSqJvMh4zuaanh+v2pIjnhIpkC1p+celam5WmIsF9jzjCIFBQFvL/xw+ngK SWJEK4AC2hThgVJQsKuSSFB1ueSGBPTM+C8rhOajdA/+u+UUOodOFTNkBPsh0cWWDtUC04C19m 5IsxDI7qRcL/Lk5Cg/iFLynVB7+Lrdj7BWO6WIel47DayiFc/tk17n1Hs1hF5/8lAmzssnpWBv 8Ow= Message-ID: <8306cf91-a7c7-aea9-4c5e-412315e38237@codesourcery.com> Date: Fri, 6 May 2022 13:19:55 +0200 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (Windows NT 10.0; Win64; x64; rv:91.0) Gecko/20100101 Thunderbird/91.7.0 From: Marcel Vollweiler Subject: [PATCH] OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible. To: X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-05.mgc.mentorg.com (139.181.222.5) To svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) X-Spam-Status: No, score=-12.7 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, 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: , Cc: Jakub Jelinek , fortran@gcc.gnu.org Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Hi, This is a follow up patch of the patch that adds the OpenMP runtime routine omp_target_is_accessible: https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591601.html It considers now also unified shared memory (usm) that was submitted recently (but not yet approved/committed): https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591349.html 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: Handle unified shared memory in omp_target_is_accessible. libgomp/ChangeLog: * target.c (omp_target_is_accessible): Handle unified shared memory. * testsuite/libgomp.c-c++-common/target-is-accessible-1.c: Updated. * testsuite/libgomp.fortran/target-is-accessible-1.f90: Updated. * testsuite/libgomp.c-c++-common/target-is-accessible-2.c: New test. * testsuite/libgomp.fortran/target-is-accessible-2.f90: New test. diff --git a/libgomp/target.c b/libgomp/target.c index 74a031f..e6d00c5 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -3909,9 +3909,13 @@ omp_target_is_accessible (const void *ptr, size_t size, int device_num) if (devicep == NULL) return false; - /* TODO: Unified shared memory must be handled when available. */ + if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return true; - return devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM; + if (devicep->is_usm_ptr_func && devicep->is_usm_ptr_func ((void *) ptr)) + return true; + + return false; } int 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 index 7c2cf62..e3f494b 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c +++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c @@ -23,23 +23,28 @@ main () 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 a[128]; for (int d = 0; d <= omp_get_num_devices (); d++) { + /* SHARED_MEM is 1 if and only if host and device share the same memory. + OMP_TARGET_IS_ACCESSIBLE should not return 0 for shared memory. */ 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) + + if (shared_mem && !omp_target_is_accessible (p, sizeof (int), d)) + __builtin_abort (); + + /* USM is disabled by default. Hence OMP_TARGET_IS_ACCESSIBLE should + return 0 if shared_mem is false. */ + if (!shared_mem && omp_target_is_accessible (p, sizeof (int), d)) __builtin_abort (); - if (omp_target_is_accessible (a, 128 * sizeof (int), d) != shared_mem) + if (shared_mem && !omp_target_is_accessible (a, 128 * sizeof (int), d)) __builtin_abort (); for (int i = 0; i < 128; i++) - if (omp_target_is_accessible (&a[i], sizeof (int), d) != shared_mem) + if (shared_mem && !omp_target_is_accessible (&a[i], sizeof (int), d)) __builtin_abort (); } diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c new file mode 100644 index 0000000..24af51f --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c @@ -0,0 +1,22 @@ +/* { dg-do run } */ +/* { dg-skip-if "USM is only implemented for nvptx." { ! offload_target_nvptx } } */ + +#include +#include + +#pragma omp requires unified_shared_memory + +int +main () +{ + int *a = (int *) omp_alloc (sizeof(int), ompx_unified_shared_mem_alloc); + if (!a) + __builtin_abort (); + + for (int d = 0; d <= omp_get_num_devices (); d++) + if (!omp_target_is_accessible (a, sizeof (int), d)) + __builtin_abort (); + + omp_free(a, ompx_unified_shared_mem_alloc); + return 0; +} diff --git a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 index 2611855..015f74a 100644 --- a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 +++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 @@ -1,3 +1,5 @@ +! { dg-do run } + program main use omp_lib use iso_c_binding @@ -25,24 +27,28 @@ program main 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. do d = 0, omp_get_num_devices () + ! SHARED_MEM is 1 if and only if host and device share the same memory. + ! OMP_TARGET_IS_ACCESSIBLE should not return 0 for shared memory. 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) & + if (shared_mem == 1 .and. omp_target_is_accessible (p, c_sizeof (d), d) == 0) & stop 5; - if (omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) /= shared_mem) & + ! USM is disabled by default. Hence OMP_TARGET_IS_ACCESSIBLE should + ! return 0 if shared_mem is false. + if (shared_mem == 0 .and. omp_target_is_accessible (p, c_sizeof (d), d) /= 0) & stop 6; + if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) == 0) & + stop 7; + do i = 1, 128 - if (omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) /= shared_mem) & - stop 7; + if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) == 0) & + stop 8; end do end do diff --git a/libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90 b/libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90 new file mode 100644 index 0000000..5c08564 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90 @@ -0,0 +1,20 @@ +! { dg-do run } +! { dg-skip-if "USM is only implemented for nvptx." { ! offload_target_nvptx } } + +program main + use omp_lib + use iso_c_binding + implicit none (external, type) + integer :: d + type(c_ptr) :: p + + p = omp_alloc (sizeof(d), ompx_unified_shared_mem_alloc) + if (.not. c_associated (p)) stop 1 + + do d = 0, omp_get_num_devices () + if (omp_target_is_accessible (p, c_sizeof (d), d) == 0) & + stop 2; + end do + + call omp_free (p, ompx_unified_shared_mem_alloc); +end program main