From: Marcel Vollweiler <marcel@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>
Cc: Jakub Jelinek <jakub@redhat.com>, <fortran@gcc.gnu.org>
Subject: [PATCH] OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible.
Date: Fri, 6 May 2022 13:19:55 +0200 [thread overview]
Message-ID: <8306cf91-a7c7-aea9-4c5e-412315e38237@codesourcery.com> (raw)
[-- Attachment #1: Type: text/plain, Size: 649 bytes --]
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
[-- Attachment #2: target-is-accessible-with-usm-patch.diff --]
[-- Type: text/plain, Size: 6294 bytes --]
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 <omp.h>
+#include <stdint.h>
+
+#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
reply other threads:[~2022-05-06 11:20 UTC|newest]
Thread overview: [no followups] expand[flat|nested] mbox.gz Atom feed
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=8306cf91-a7c7-aea9-4c5e-412315e38237@codesourcery.com \
--to=marcel@codesourcery.com \
--cc=fortran@gcc.gnu.org \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).