From: Andrew Stubbs <ams@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>
Subject: [PATCH 04/17] openmp, nvptx: low-lat memory access traits
Date: Thu, 7 Jul 2022 11:34:35 +0100 [thread overview]
Message-ID: <2810723bd4e98723e5b9eca476eb7e981590c81a.1657188329.git.ams@codesourcery.com> (raw)
In-Reply-To: <cover.1657188329.git.ams@codesourcery.com>
[-- Attachment #1: Type: text/plain, Size: 1135 bytes --]
The NVPTX low latency memory is not accessible outside the team that allocates
it, and therefore should be unavailable for allocators with the access trait
"all". This change means that the omp_low_lat_mem_alloc predefined
allocator now implicitly implies the "pteam" trait.
libgomp/ChangeLog:
* allocator.c (MEMSPACE_VALIDATE): New macro.
(omp_aligned_alloc): Use MEMSPACE_VALIDATE.
(omp_aligned_calloc): Likewise.
(omp_realloc): Likewise.
* config/nvptx/allocator.c (nvptx_memspace_validate): New function.
(MEMSPACE_VALIDATE): New macro.
* testsuite/libgomp.c/allocators-4.c (main): Add access trait.
* testsuite/libgomp.c/allocators-6.c (main): Add access trait.
* testsuite/libgomp.c/allocators-7.c: New test.
---
libgomp/allocator.c | 15 +++++
libgomp/config/nvptx/allocator.c | 11 ++++
libgomp/testsuite/libgomp.c/allocators-4.c | 7 ++-
libgomp/testsuite/libgomp.c/allocators-6.c | 7 ++-
libgomp/testsuite/libgomp.c/allocators-7.c | 68 ++++++++++++++++++++++
5 files changed, 102 insertions(+), 6 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c/allocators-7.c
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0004-openmp-nvptx-low-lat-memory-access-traits.patch --]
[-- Type: text/x-patch; name="0004-openmp-nvptx-low-lat-memory-access-traits.patch", Size: 6096 bytes --]
diff --git a/libgomp/allocator.c b/libgomp/allocator.c
index 029d0d40a36..48ab0782e6b 100644
--- a/libgomp/allocator.c
+++ b/libgomp/allocator.c
@@ -54,6 +54,9 @@
#define MEMSPACE_FREE(MEMSPACE, ADDR, SIZE, PIN) \
(PIN ? NULL : free (ADDR))
#endif
+#ifndef MEMSPACE_VALIDATE
+#define MEMSPACE_VALIDATE(MEMSPACE, ACCESS) 1
+#endif
/* Map the predefined allocators to the correct memory space.
The index to this table is the omp_allocator_handle_t enum value. */
@@ -438,6 +441,10 @@ retry:
if (__builtin_add_overflow (size, new_size, &new_size))
goto fail;
+ if (allocator_data
+ && !MEMSPACE_VALIDATE (allocator_data->memspace, allocator_data->access))
+ goto fail;
+
if (__builtin_expect (allocator_data
&& allocator_data->pool_size < ~(uintptr_t) 0, 0))
{
@@ -733,6 +740,10 @@ retry:
if (__builtin_add_overflow (size_temp, new_size, &new_size))
goto fail;
+ if (allocator_data
+ && !MEMSPACE_VALIDATE (allocator_data->memspace, allocator_data->access))
+ goto fail;
+
if (__builtin_expect (allocator_data
&& allocator_data->pool_size < ~(uintptr_t) 0, 0))
{
@@ -964,6 +975,10 @@ retry:
goto fail;
old_size = data->size;
+ if (allocator_data
+ && !MEMSPACE_VALIDATE (allocator_data->memspace, allocator_data->access))
+ goto fail;
+
if (__builtin_expect (allocator_data
&& allocator_data->pool_size < ~(uintptr_t) 0, 0))
{
diff --git a/libgomp/config/nvptx/allocator.c b/libgomp/config/nvptx/allocator.c
index f740b97f6ac..0102680b717 100644
--- a/libgomp/config/nvptx/allocator.c
+++ b/libgomp/config/nvptx/allocator.c
@@ -358,6 +358,15 @@ nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr,
return realloc (addr, size);
}
+static inline int
+nvptx_memspace_validate (omp_memspace_handle_t memspace, unsigned access)
+{
+ /* Disallow use of low-latency memory when it must be accessible by
+ all threads. */
+ return (memspace != omp_low_lat_mem_space
+ || access != omp_atv_all);
+}
+
#define MEMSPACE_ALLOC(MEMSPACE, SIZE, PIN) \
nvptx_memspace_alloc (MEMSPACE, SIZE)
#define MEMSPACE_CALLOC(MEMSPACE, SIZE, PIN) \
@@ -366,5 +375,7 @@ nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr,
nvptx_memspace_realloc (MEMSPACE, ADDR, OLDSIZE, SIZE)
#define MEMSPACE_FREE(MEMSPACE, ADDR, SIZE, PIN) \
nvptx_memspace_free (MEMSPACE, ADDR, SIZE)
+#define MEMSPACE_VALIDATE(MEMSPACE, ACCESS) \
+ nvptx_memspace_validate (MEMSPACE, ACCESS)
#include "../../allocator.c"
diff --git a/libgomp/testsuite/libgomp.c/allocators-4.c b/libgomp/testsuite/libgomp.c/allocators-4.c
index 9fa6aa1624f..cae27ea33c1 100644
--- a/libgomp/testsuite/libgomp.c/allocators-4.c
+++ b/libgomp/testsuite/libgomp.c/allocators-4.c
@@ -23,10 +23,11 @@ main ()
#pragma omp target
{
/* Ensure that the memory we get *is* low-latency with a null-fallback. */
- omp_alloctrait_t traits[1]
- = { { omp_atk_fallback, omp_atv_null_fb } };
+ omp_alloctrait_t traits[2]
+ = { { omp_atk_fallback, omp_atv_null_fb },
+ { omp_atk_access, omp_atv_pteam } };
omp_allocator_handle_t lowlat = omp_init_allocator (omp_low_lat_mem_space,
- 1, traits);
+ 2, traits);
int size = 4;
diff --git a/libgomp/testsuite/libgomp.c/allocators-6.c b/libgomp/testsuite/libgomp.c/allocators-6.c
index 90bf73095ef..c03233df582 100644
--- a/libgomp/testsuite/libgomp.c/allocators-6.c
+++ b/libgomp/testsuite/libgomp.c/allocators-6.c
@@ -23,10 +23,11 @@ main ()
#pragma omp target
{
/* Ensure that the memory we get *is* low-latency with a null-fallback. */
- omp_alloctrait_t traits[1]
- = { { omp_atk_fallback, omp_atv_null_fb } };
+ omp_alloctrait_t traits[2]
+ = { { omp_atk_fallback, omp_atv_null_fb },
+ { omp_atk_access, omp_atv_pteam } };
omp_allocator_handle_t lowlat = omp_init_allocator (omp_low_lat_mem_space,
- 1, traits);
+ 2, traits);
int size = 16;
diff --git a/libgomp/testsuite/libgomp.c/allocators-7.c b/libgomp/testsuite/libgomp.c/allocators-7.c
new file mode 100644
index 00000000000..a0a738b1d1d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/allocators-7.c
@@ -0,0 +1,68 @@
+/* { dg-do run } */
+
+/* { dg-require-effective-target offload_device } */
+/* { dg-xfail-if "not implemented" { ! offload_target_nvptx } } */
+
+/* Test that GPU low-latency allocation is limited to team access. */
+
+#include <stddef.h>
+#include <omp.h>
+
+#pragma omp requires dynamic_allocators
+
+int
+main ()
+{
+ #pragma omp target
+ {
+ /* Ensure that the memory we get *is* low-latency with a null-fallback. */
+ omp_alloctrait_t traits[2]
+ = { { omp_atk_fallback, omp_atv_null_fb },
+ { omp_atk_access, omp_atv_pteam } };
+ omp_allocator_handle_t lowlat = omp_init_allocator (omp_low_lat_mem_space,
+ 2, traits);
+
+ omp_alloctrait_t traits_all[2]
+ = { { omp_atk_fallback, omp_atv_null_fb },
+ { omp_atk_access, omp_atv_all } };
+ omp_allocator_handle_t lowlat_all
+ = omp_init_allocator (omp_low_lat_mem_space, 2, traits_all);
+
+ omp_alloctrait_t traits_default[1]
+ = { { omp_atk_fallback, omp_atv_null_fb } };
+ omp_allocator_handle_t lowlat_default
+ = omp_init_allocator (omp_low_lat_mem_space, 1, traits_default);
+
+ void *a = omp_alloc(1, lowlat); // good
+ void *b = omp_alloc(1, lowlat_all); // bad
+ void *c = omp_alloc(1, lowlat_default); // bad
+
+ if (!a || b || c)
+ __builtin_abort ();
+
+ omp_free (a, lowlat);
+
+
+ a = omp_calloc(1, 1, lowlat); // good
+ b = omp_calloc(1, 1, lowlat_all); // bad
+ c = omp_calloc(1, 1, lowlat_default); // bad
+
+ if (!a || b || c)
+ __builtin_abort ();
+
+ omp_free (a, lowlat);
+
+
+ a = omp_realloc(NULL, 1, lowlat, lowlat); // good
+ b = omp_realloc(NULL, 1, lowlat_all, lowlat_all); // bad
+ c = omp_realloc(NULL, 1, lowlat_default, lowlat_default); // bad
+
+ if (!a || b || c)
+ __builtin_abort ();
+
+ omp_free (a, lowlat);
+ }
+
+return 0;
+}
+
next prev parent reply other threads:[~2022-07-07 10:35 UTC|newest]
Thread overview: 30+ messages / expand[flat|nested] mbox.gz Atom feed top
2022-07-07 10:34 [PATCH 00/17] openmp, nvptx, amdgcn: 5.0 Memory Allocators Andrew Stubbs
2022-07-07 10:34 ` [PATCH 01/17] libgomp, nvptx: low-latency memory allocator Andrew Stubbs
2022-12-08 11:40 ` Jakub Jelinek
2022-07-07 10:34 ` [PATCH 02/17] libgomp: pinned memory Andrew Stubbs
2022-12-08 12:11 ` Jakub Jelinek
2022-12-08 12:51 ` Andrew Stubbs
2022-12-08 14:02 ` Tobias Burnus
2022-12-08 14:35 ` Andrew Stubbs
2022-12-08 15:02 ` Tobias Burnus
2022-07-07 10:34 ` [PATCH 03/17] libgomp, openmp: Add ompx_pinned_mem_alloc Andrew Stubbs
2022-07-07 10:34 ` Andrew Stubbs [this message]
2022-07-07 10:34 ` [PATCH 05/17] openmp, nvptx: ompx_unified_shared_mem_alloc Andrew Stubbs
2022-07-07 10:34 ` [PATCH 06/17] openmp: Add -foffload-memory Andrew Stubbs
2022-07-07 10:34 ` [PATCH 07/17] openmp: allow requires unified_shared_memory Andrew Stubbs
2022-07-07 10:34 ` [PATCH 08/17] openmp: -foffload-memory=pinned Andrew Stubbs
2022-07-07 11:54 ` Tobias Burnus
2022-07-07 22:18 ` Andrew Stubbs
2022-07-08 9:00 ` Tobias Burnus
2022-07-08 9:55 ` Andrew Stubbs
2022-07-08 9:57 ` Tobias Burnus
2023-02-20 14:59 ` Prototype 'GOMP_enable_pinned_mode' (was: [PATCH 08/17] openmp: -foffload-memory=pinned) Thomas Schwinge
2022-07-07 10:34 ` [PATCH 09/17] openmp: Use libgomp memory allocation functions with unified shared memory Andrew Stubbs
2022-07-07 10:34 ` [PATCH 10/17] Add parsing support for allocate directive (OpenMP 5.0) Andrew Stubbs
2022-07-07 10:34 ` [PATCH 11/17] Translate " Andrew Stubbs
2022-07-07 10:34 ` [PATCH 12/17] Handle cleanup of omp allocated variables " Andrew Stubbs
2022-07-07 10:34 ` [PATCH 13/17] Gimplify allocate directive " Andrew Stubbs
2022-07-07 10:34 ` [PATCH 14/17] Lower " Andrew Stubbs
2022-07-07 10:34 ` [PATCH 15/17] amdgcn: Support XNACK mode Andrew Stubbs
2022-07-07 10:34 ` [PATCH 16/17] amdgcn, openmp: Auto-detect USM mode and set HSA_XNACK Andrew Stubbs
2022-07-07 10:34 ` [PATCH 17/17] amdgcn: libgomp plugin USM implementation Andrew Stubbs
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=2810723bd4e98723e5b9eca476eb7e981590c81a.1657188329.git.ams@codesourcery.com \
--to=ams@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
/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).