From patchwork Tue Dec 13 16:12:22 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Marcel Vollweiler X-Patchwork-Id: 32871 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:adf:e747:0:0:0:0:0 with SMTP id c7csp213811wrn; Tue, 13 Dec 2022 08:13:02 -0800 (PST) X-Google-Smtp-Source: AA0mqf7XbEObLYqvlfTSuf4E6S/0MG68TKogAJbQB+TLlG0WfyEGomGTB4b6v8CBmmmlhMK/a/iR X-Received: by 2002:a17:906:828f:b0:7c1:652:d109 with SMTP id h15-20020a170906828f00b007c10652d109mr17475420ejx.35.1670947982545; Tue, 13 Dec 2022 08:13:02 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1670947982; cv=none; d=google.com; s=arc-20160816; b=VBd6O45Xy3nMdweTYUKmLYHjS+W4oPB9ag1oRUsLb0CfqzQtJp0LxjZGQTD5xmBicK ks1Hs7ipCwfXm5y8dEC61Hi8GqC0UJmN+WVh7NlcQ49ADTlYBM+CcyEQfoNtKvolybES 3rXzgPlpCApvxzbLAkF9w1ifPfP+50Ih3hjp67lM1Bvsf8LKOaU6fS1GC48oHaZ4jLBT 9qY9LPdGuANVmAnqbjBEsGvRWwmeOsPRD/gOkJKBG5MpDcAQrcjiyU31c1OGtQJgBEkq qF6+ygnYcj/cPsvZ+Mrka4RLtEz7+blvYzvdk5AbbRe73c1c7Sv74C1zoijX3oZs9k2E lKYQ== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:to:subject:from:user-agent :mime-version:date:message-id:ironport-sdr:dmarc-filter:delivered-to; bh=L2xnJPXl7LEUZqoclzywMGXd4dS+etsxdBkMVvFVRbw=; b=n6PD7jgV5F1wH8DovaPZsQgytm2+UlatmFlOETOiTkkhQYQuBMfuktRNs7+Kty4xmu Ygmf9iDdQS66suytpVPGOJhbNEeis4aMGhRmhusmTFo3PReVTBrKj3TImSIqIW8xN8xt vaInLsBXOKAVBMEOzd9xkfEHJDTll/VmzRcdq3/o0hu2vh/Vm2zhbLpY6/uQOALEHGQM eVETK6Z3VXkuMkf6iS4ZQVAiJlW+Jt8s+43Vi8ZbaK5ZKkeruCy8vI1004GjGL3nMyIJ fkuW/swlGv4phYtsYFlw+V9rsf4F3nowWkJ0D9WZ8whu1YIENqJvDjs6tFpJGwI2U73m estQ== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from sourceware.org (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id hv22-20020a17090760d600b007808f3f4cbcsi9762152ejc.239.2022.12.13.08.13.02 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Tue, 13 Dec 2022 08:13:02 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) client-ip=2620:52:3:1:0:246e:9693:128c; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id ABE26383B6F3 for ; Tue, 13 Dec 2022 16:12:56 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 53E2B384D6F3 for ; Tue, 13 Dec 2022 16:12:31 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 53E2B384D6F3 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.96,241,1665475200"; d="diff'?scan'208";a="89539052" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 13 Dec 2022 08:12:29 -0800 IronPort-SDR: 2U9kbGP1HWGst5fd4uyXV01n/CnXGmLv00C/TCGPdudtHu2FZgsX8B1Df8uWSep6AVB7cj48jI Mq0XO2Y9j2cijFPBHPqepvLNu6mlrULWSOrgqR8FIP9LWH4m6cIqngkgIaqfVmkca0C29kQ5Q0 9NW7A7tf1RBtUrKepOXPAtg1W1gTirvoqhDkWBnWAIk64UETYHNQ3O5wfchiISBBwoYHyprP2A UMwMHs+gx6daJr5icDu/IrEKszr+GEXH+WzrHG0kw5BbVJR1ZVt1BleK7rWNDtPdRU3jbzDrtl 0To= Message-ID: <0ff30ea1-ee5f-7f10-dcbc-bea85e2bfa81@codesourcery.com> Date: Tue, 13 Dec 2022 17:12:22 +0100 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (Windows NT 10.0; Win64; x64; rv:102.0) Gecko/20100101 Thunderbird/102.3.0 From: Marcel Vollweiler Subject: [OG12][committed] OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible. To: X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-14.mgc.mentorg.com (139.181.222.14) 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 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.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: =?utf-8?q?INBOX?= X-GMAIL-THRID: =?utf-8?q?1752115951543382683?= X-GMAIL-MSGID: =?utf-8?q?1752115951543382683?= This patch handles Unified Shared Memory (USM) in the OpenMP runtime routine omp_target_is_accessible implementation. A previous patch was submitted some months ago (https://gcc.gnu.org/pipermail/gcc-patches/2022-May/594187.html) but not yet reviewed due to dependencies on the Unified Shared Memory implementation. Although USM is not yet in mainline, the corresponding patches were already committed to OG12. I rebased, updated, and committed my patch to OG12 (devel/omp/gcc-12 branch). I tested the patch with nvptx offloading (x86_64-linux and PowerPC) without regressions. Since USM is not supported for all gcn targets, I tested gcn with offloading for x86_64-linux on various targets (gfx90a, gfx908, gfx906, gfx803) - also without 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 commit 9044b7efb3518de180a5b3168615b7e12d93eea8 Author: Marcel Vollweiler Date: Tue Dec 13 12:04:48 2022 +0000 OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible This patch handles Unified Shared Memory (USM) in the OpenMP runtime routine 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/ChangeLog.omp b/libgomp/ChangeLog.omp index 32bcc84..a0d0271 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,11 @@ +2022-12-13 Marcel Vollweiler + + * 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. + 2022-12-12 Tobias Burnus Backported from master: diff --git a/libgomp/target.c b/libgomp/target.c index 50709f0..2cd8e2a 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -5067,9 +5067,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 2e75c63..e7f9cf2 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 @@ -1,3 +1,5 @@ +/* { dg-do run } */ + #include int @@ -6,7 +8,8 @@ main () int d = omp_get_default_device (); int id = omp_get_initial_device (); int n = omp_get_num_devices (); - void *p; + int i = 42; + void *p = &i; if (d < 0 || d >= n) d = id; @@ -26,23 +29,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..0917365 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c @@ -0,0 +1,21 @@ +/* { dg-do run } */ +/* { dg-require-effective-target omp_usm } */ + +#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 150df6f..0df43aae 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 @@ -28,24 +30,28 @@ program main if (omp_target_is_accessible (p, c_sizeof (d), n + 1) /= 0) & stop 5 - ! 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 6; - 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 7; + if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) == 0) & + stop 8; + do i = 1, 128 - if (omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) /= shared_mem) & - stop 8; + if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) == 0) & + stop 9; 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..624d1ef --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90 @@ -0,0 +1,22 @@ +! { dg-do run } +! { dg-require-effective-target omp_usm } + +program main + use omp_lib + use iso_c_binding + implicit none (external, type) + integer :: d + type(c_ptr) :: p + + !$omp requires unified_shared_memory + + 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