From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 7845) id B4E393873857; Tue, 13 Dec 2022 14:34:06 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org B4E393873857 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1670942046; bh=1e3A1PKc6o9ozGZZKAs5kkz7/mFjaASA2wr31hXjCkU=; h=From:To:Subject:Date:From; b=wj099FXdfiJ7sXrJrca+0+VUiW44Noe1qMB6mmCQNNrxppIZMf/86RTEKQiOf7KwD f6vvk/Ta0rxvWK3yq8ErHHoHL/bQgLK3jiSqd+YnFlaQej0BAPzMpnLgN7r0gR5Kha IkD5phma1tRsSdaojw2IIQhk3E84oP9udVc7ZT1Y= Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Marcel Vollweiler To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible X-Act-Checkin: gcc X-Git-Author: Marcel Vollweiler X-Git-Refname: refs/heads/devel/omp/gcc-12 X-Git-Oldrev: daf3af9e5b6da664115207a6dae3706d11cc4614 X-Git-Newrev: 9044b7efb3518de180a5b3168615b7e12d93eea8 Message-Id: <20221213143406.B4E393873857@sourceware.org> Date: Tue, 13 Dec 2022 14:34:06 +0000 (GMT) List-Id: https://gcc.gnu.org/g:9044b7efb3518de180a5b3168615b7e12d93eea8 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: --- libgomp/ChangeLog.omp | 8 ++++++++ libgomp/target.c | 8 ++++++-- .../libgomp.c-c++-common/target-is-accessible-1.c | 22 +++++++++++++++------- .../libgomp.c-c++-common/target-is-accessible-2.c | 21 +++++++++++++++++++++ .../libgomp.fortran/target-is-accessible-1.f90 | 20 +++++++++++++------- .../libgomp.fortran/target-is-accessible-2.f90 | 22 ++++++++++++++++++++++ 6 files changed, 85 insertions(+), 16 deletions(-) diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 32bcc842af8..a0d0271f63b 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 50709f0677d..2cd8e2a65dd 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 2e75c6300ae..e7f9cf27a42 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 00000000000..0917365379b --- /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 150df6f8a4f..0df43aae095 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 00000000000..624d1efa08e --- /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