* [PATCH v2 0/3] libgomp: OpenMP low-latency omp_alloc
@ 2023-08-02 17:00 Andrew Stubbs
2023-08-02 17:00 ` [PATCH v2 1/3] libgomp, nvptx: low-latency memory allocator Andrew Stubbs
` (2 more replies)
0 siblings, 3 replies; 7+ messages in thread
From: Andrew Stubbs @ 2023-08-02 17:00 UTC (permalink / raw)
To: gcc-patches
This patch series is an updated and reworked version of some of the patch set
posted about a year ago (the other features will be posted soon), this
time supporting amdgcn, in addition to nvptx:
https://patchwork.sourceware.org/project/gcc/list/?series=10748&state=%2A&archive=both
The series implements device-specific allocators and adds a low-latency
allocator for both GPUs architectures.
The previous review comments have been addressed, I hope, plus a lot of
bugs have been found and fixed since the original post. With the
addition of amdgcn I have broken out the heap implementation so both
architectures can share the code.
Andrew
Andrew Stubbs (3):
libgomp, nvptx: low-latency memory allocator
openmp, nvptx: low-lat memory access traits
amdgcn, libgomp: low-latency allocator
gcc/config/gcn/gcn-builtins.def | 2 +
gcc/config/gcn/gcn.cc | 16 +-
libgomp/allocator.c | 269 +++++++++----
libgomp/basic-allocator.c | 380 ++++++++++++++++++
libgomp/config/gcn/allocator.c | 123 ++++++
libgomp/config/gcn/libgomp-gcn.h | 6 +
libgomp/config/gcn/team.c | 12 +
libgomp/config/nvptx/allocator.c | 131 ++++++
libgomp/config/nvptx/team.c | 18 +
libgomp/libgomp.h | 3 -
libgomp/plugin/plugin-gcn.c | 35 +-
libgomp/plugin/plugin-nvptx.c | 23 +-
libgomp/testsuite/libgomp.c/omp_alloc-1.c | 56 +++
libgomp/testsuite/libgomp.c/omp_alloc-2.c | 64 +++
libgomp/testsuite/libgomp.c/omp_alloc-3.c | 42 ++
libgomp/testsuite/libgomp.c/omp_alloc-4.c | 197 +++++++++
libgomp/testsuite/libgomp.c/omp_alloc-5.c | 63 +++
libgomp/testsuite/libgomp.c/omp_alloc-6.c | 118 ++++++
.../testsuite/libgomp.c/omp_alloc-traits.c | 68 ++++
19 files changed, 1528 insertions(+), 98 deletions(-)
create mode 100644 libgomp/basic-allocator.c
create mode 100644 libgomp/config/gcn/allocator.c
create mode 100644 libgomp/config/nvptx/allocator.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-1.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-2.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-3.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-4.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-5.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-6.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-traits.c
--
2.41.0
^ permalink raw reply [flat|nested] 7+ messages in thread
* [PATCH v2 1/3] libgomp, nvptx: low-latency memory allocator
2023-08-02 17:00 [PATCH v2 0/3] libgomp: OpenMP low-latency omp_alloc Andrew Stubbs
@ 2023-08-02 17:00 ` Andrew Stubbs
2023-09-08 9:04 ` Tobias Burnus
2023-08-02 17:00 ` [PATCH v2 2/3] openmp, nvptx: low-lat memory access traits Andrew Stubbs
2023-08-02 17:00 ` [PATCH v2 3/3] amdgcn, libgomp: low-latency allocator Andrew Stubbs
2 siblings, 1 reply; 7+ messages in thread
From: Andrew Stubbs @ 2023-08-02 17:00 UTC (permalink / raw)
To: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 2854 bytes --]
This patch adds support for allocating low-latency ".shared" memory on
NVPTX GPU device, via the omp_low_lat_mem_space and omp_alloc. The memory
can be allocated, reallocated, and freed using a basic but fast algorithm,
is thread safe and the size of the low-latency heap can be configured using
the GOMP_NVPTX_LOWLAT_POOL environment variable.
The use of the PTX dynamic_smem_size feature means that low-latency allocator
will not work with the PTX 3.1 multilib.
libgomp/ChangeLog:
* allocator.c (MEMSPACE_ALLOC): New macro.
(MEMSPACE_CALLOC): New macro.
(MEMSPACE_REALLOC): New macro.
(MEMSPACE_FREE): New macro.
(predefined_alloc_mapping): New array.
(omp_aligned_alloc): Use MEMSPACE_ALLOC.
Implement fall-backs for predefined allocators.
(omp_free): Use MEMSPACE_FREE.
(omp_calloc): Use MEMSPACE_CALLOC.
(omp_realloc): Use MEMSPACE_REALLOC, MEMSPACE_ALLOC, and MEMSPACE_FREE.
* config/nvptx/team.c (__nvptx_lowlat_pool): New asm variable.
(__nvptx_lowlat_init): New prototype.
(gomp_nvptx_main): Call __nvptx_lowlat_init.
* plugin/plugin-nvptx.c (lowlat_pool_size): New variable.
(GOMP_OFFLOAD_init_device): Read the GOMP_NVPTX_LOWLAT_POOL envvar.
(GOMP_OFFLOAD_run): Apply lowlat_pool_size.
* basic-allocator.c: New file.
* config/nvptx/allocator.c: New file.
* testsuite/libgomp.c/omp_alloc-1.c: New test.
* testsuite/libgomp.c/omp_alloc-2.c: New test.
* testsuite/libgomp.c/omp_alloc-3.c: New test.
* testsuite/libgomp.c/omp_alloc-4.c: New test.
* testsuite/libgomp.c/omp_alloc-5.c: New test.
* testsuite/libgomp.c/omp_alloc-6.c: New test.
Co-authored-by: Kwok Cheung Yeung <kcy@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
---
libgomp/allocator.c | 253 +++++++++-----
libgomp/basic-allocator.c | 380 ++++++++++++++++++++++
libgomp/config/nvptx/allocator.c | 120 +++++++
libgomp/config/nvptx/team.c | 18 +
libgomp/plugin/plugin-nvptx.c | 23 +-
libgomp/testsuite/libgomp.c/omp_alloc-1.c | 56 ++++
libgomp/testsuite/libgomp.c/omp_alloc-2.c | 64 ++++
libgomp/testsuite/libgomp.c/omp_alloc-3.c | 42 +++
libgomp/testsuite/libgomp.c/omp_alloc-4.c | 196 +++++++++++
libgomp/testsuite/libgomp.c/omp_alloc-5.c | 63 ++++
libgomp/testsuite/libgomp.c/omp_alloc-6.c | 117 +++++++
11 files changed, 1244 insertions(+), 88 deletions(-)
create mode 100644 libgomp/basic-allocator.c
create mode 100644 libgomp/config/nvptx/allocator.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-1.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-2.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-3.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-4.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-5.c
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-6.c
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: v2-0001-libgomp-nvptx-low-latency-memory-allocator.patch --]
[-- Type: text/x-patch; name="v2-0001-libgomp-nvptx-low-latency-memory-allocator.patch", Size: 47072 bytes --]
diff --git a/libgomp/allocator.c b/libgomp/allocator.c
index 90f2dcb60d6..fbf7b1ab061 100644
--- a/libgomp/allocator.c
+++ b/libgomp/allocator.c
@@ -37,6 +37,42 @@
#define omp_max_predefined_alloc omp_thread_mem_alloc
+/* These macros may be overridden in config/<target>/allocator.c.
+ The following definitions (ab)use comma operators to avoid unused
+ variable errors. */
+#ifndef MEMSPACE_ALLOC
+#define MEMSPACE_ALLOC(MEMSPACE, SIZE) \
+ malloc (((void)(MEMSPACE), (SIZE)))
+#endif
+#ifndef MEMSPACE_CALLOC
+#define MEMSPACE_CALLOC(MEMSPACE, SIZE) \
+ calloc (1, (((void)(MEMSPACE), (SIZE))))
+#endif
+#ifndef MEMSPACE_REALLOC
+#define MEMSPACE_REALLOC(MEMSPACE, ADDR, OLDSIZE, SIZE) \
+ realloc (ADDR, (((void)(MEMSPACE), (void)(OLDSIZE), (SIZE))))
+#endif
+#ifndef MEMSPACE_FREE
+#define MEMSPACE_FREE(MEMSPACE, ADDR, SIZE) \
+ free (((void)(MEMSPACE), (void)(SIZE), (ADDR)))
+#endif
+
+/* Map the predefined allocators to the correct memory space.
+ The index to this table is the omp_allocator_handle_t enum value.
+ When the user calls omp_alloc with a predefined allocator this
+ table determines what memory they get. */
+static const omp_memspace_handle_t predefined_alloc_mapping[] = {
+ omp_default_mem_space, /* omp_null_allocator. */
+ omp_default_mem_space, /* omp_default_mem_alloc. */
+ omp_large_cap_mem_space, /* omp_large_cap_mem_alloc. */
+ omp_const_mem_space, /* omp_const_mem_alloc. */
+ omp_high_bw_mem_space, /* omp_high_bw_mem_alloc. */
+ omp_low_lat_mem_space, /* omp_low_lat_mem_alloc. */
+ omp_low_lat_mem_space, /* omp_cgroup_mem_alloc. */
+ omp_low_lat_mem_space, /* omp_pteam_mem_alloc. */
+ omp_low_lat_mem_space, /* omp_thread_mem_alloc. */
+};
+
enum gomp_numa_memkind_kind
{
GOMP_MEMKIND_NONE = 0,
@@ -522,7 +558,7 @@ retry:
}
else
#endif
- ptr = malloc (new_size);
+ ptr = MEMSPACE_ALLOC (allocator_data->memspace, new_size);
if (ptr == NULL)
{
#ifdef HAVE_SYNC_BUILTINS
@@ -554,7 +590,13 @@ retry:
}
else
#endif
- ptr = malloc (new_size);
+ {
+ omp_memspace_handle_t memspace;
+ memspace = (allocator_data
+ ? allocator_data->memspace
+ : predefined_alloc_mapping[allocator]);
+ ptr = MEMSPACE_ALLOC (memspace, new_size);
+ }
if (ptr == NULL)
goto fail;
}
@@ -571,36 +613,38 @@ retry:
((struct omp_mem_header *) ret)[-1].allocator = allocator;
return ret;
-fail:
- if (allocator_data)
+fail:;
+ int fallback = (allocator_data
+ ? allocator_data->fallback
+ : allocator == omp_default_mem_alloc
+ ? omp_atv_null_fb
+ : omp_atv_default_mem_fb);
+ switch (fallback)
{
- switch (allocator_data->fallback)
- {
- case omp_atv_default_mem_fb:
- if ((new_alignment > sizeof (void *) && new_alignment > alignment)
+ case omp_atv_default_mem_fb:
+ if ((new_alignment > sizeof (void *) && new_alignment > alignment)
#if defined(LIBGOMP_USE_MEMKIND) || defined(LIBGOMP_USE_LIBNUMA)
- || memkind
+ || memkind
#endif
- || (allocator_data
- && allocator_data->pool_size < ~(uintptr_t) 0))
- {
- allocator = omp_default_mem_alloc;
- goto retry;
- }
- /* Otherwise, we've already performed default mem allocation
- and if that failed, it won't succeed again (unless it was
- intermittent. Return NULL then, as that is the fallback. */
- break;
- case omp_atv_null_fb:
- break;
- default:
- case omp_atv_abort_fb:
- gomp_fatal ("Out of memory allocating %lu bytes",
- (unsigned long) size);
- case omp_atv_allocator_fb:
- allocator = allocator_data->fb_data;
+ || allocator_data == NULL
+ || allocator_data->pool_size < ~(uintptr_t) 0)
+ {
+ allocator = omp_default_mem_alloc;
goto retry;
}
+ /* Otherwise, we've already performed default mem allocation
+ and if that failed, it won't succeed again (unless it was
+ intermittent). Return NULL then, as that is the fallback. */
+ break;
+ case omp_atv_null_fb:
+ break;
+ default:
+ case omp_atv_abort_fb:
+ gomp_fatal ("Out of memory allocating %lu bytes",
+ (unsigned long) size);
+ case omp_atv_allocator_fb:
+ allocator = allocator_data->fb_data;
+ goto retry;
}
return NULL;
}
@@ -633,6 +677,7 @@ void
omp_free (void *ptr, omp_allocator_handle_t allocator)
{
struct omp_mem_header *data;
+ omp_memspace_handle_t memspace = omp_default_mem_space;
if (ptr == NULL)
return;
@@ -672,10 +717,12 @@ omp_free (void *ptr, omp_allocator_handle_t allocator)
return;
}
#endif
+
+ memspace = allocator_data->memspace;
}
-#ifdef LIBGOMP_USE_MEMKIND
else
{
+#ifdef LIBGOMP_USE_MEMKIND
enum gomp_numa_memkind_kind memkind = GOMP_MEMKIND_NONE;
if (data->allocator == omp_high_bw_mem_alloc)
memkind = GOMP_MEMKIND_HBW_PREFERRED;
@@ -691,9 +738,12 @@ omp_free (void *ptr, omp_allocator_handle_t allocator)
return;
}
}
- }
#endif
- free (data->ptr);
+
+ memspace = predefined_alloc_mapping[data->allocator];
+ }
+
+ MEMSPACE_FREE (memspace, data->ptr, data->size);
}
ialias (omp_free)
@@ -820,7 +870,7 @@ retry:
}
else
#endif
- ptr = calloc (1, new_size);
+ ptr = MEMSPACE_CALLOC (allocator_data->memspace, new_size);
if (ptr == NULL)
{
#ifdef HAVE_SYNC_BUILTINS
@@ -854,7 +904,13 @@ retry:
}
else
#endif
- ptr = calloc (1, new_size);
+ {
+ omp_memspace_handle_t memspace;
+ memspace = (allocator_data
+ ? allocator_data->memspace
+ : predefined_alloc_mapping[allocator]);
+ ptr = MEMSPACE_CALLOC (memspace, new_size);
+ }
if (ptr == NULL)
goto fail;
}
@@ -871,36 +927,38 @@ retry:
((struct omp_mem_header *) ret)[-1].allocator = allocator;
return ret;
-fail:
- if (allocator_data)
+fail:;
+ int fallback = (allocator_data
+ ? allocator_data->fallback
+ : allocator == omp_default_mem_alloc
+ ? omp_atv_null_fb
+ : omp_atv_default_mem_fb);
+ switch (fallback)
{
- switch (allocator_data->fallback)
- {
- case omp_atv_default_mem_fb:
- if ((new_alignment > sizeof (void *) && new_alignment > alignment)
+ case omp_atv_default_mem_fb:
+ if ((new_alignment > sizeof (void *) && new_alignment > alignment)
#if defined(LIBGOMP_USE_MEMKIND) || defined(LIBGOMP_USE_LIBNUMA)
- || memkind
+ || memkind
#endif
- || (allocator_data
- && allocator_data->pool_size < ~(uintptr_t) 0))
- {
- allocator = omp_default_mem_alloc;
- goto retry;
- }
- /* Otherwise, we've already performed default mem allocation
- and if that failed, it won't succeed again (unless it was
- intermittent. Return NULL then, as that is the fallback. */
- break;
- case omp_atv_null_fb:
- break;
- default:
- case omp_atv_abort_fb:
- gomp_fatal ("Out of memory allocating %lu bytes",
- (unsigned long) (size * nmemb));
- case omp_atv_allocator_fb:
- allocator = allocator_data->fb_data;
+ || allocator_data == NULL
+ || allocator_data->pool_size < ~(uintptr_t) 0)
+ {
+ allocator = omp_default_mem_alloc;
goto retry;
}
+ /* Otherwise, we've already performed default mem allocation
+ and if that failed, it won't succeed again (unless it was
+ intermittent). Return NULL then, as that is the fallback. */
+ break;
+ case omp_atv_null_fb:
+ break;
+ default:
+ case omp_atv_abort_fb:
+ gomp_fatal ("Out of memory allocating %lu bytes",
+ (unsigned long) (size * nmemb));
+ case omp_atv_allocator_fb:
+ allocator = allocator_data->fb_data;
+ goto retry;
}
return NULL;
}
@@ -1090,9 +1148,10 @@ retry:
else
#endif
if (prev_size)
- new_ptr = realloc (data->ptr, new_size);
+ new_ptr = MEMSPACE_REALLOC (allocator_data->memspace, data->ptr,
+ data->size, new_size);
else
- new_ptr = malloc (new_size);
+ new_ptr = MEMSPACE_ALLOC (allocator_data->memspace, new_size);
if (new_ptr == NULL)
{
#ifdef HAVE_SYNC_BUILTINS
@@ -1140,7 +1199,13 @@ retry:
}
else
#endif
- new_ptr = realloc (data->ptr, new_size);
+ {
+ omp_memspace_handle_t memspace;
+ memspace = (allocator_data
+ ? allocator_data->memspace
+ : predefined_alloc_mapping[allocator]);
+ new_ptr = MEMSPACE_REALLOC (memspace, data->ptr, data->size, new_size);
+ }
if (new_ptr == NULL)
goto fail;
ret = (char *) new_ptr + sizeof (struct omp_mem_header);
@@ -1167,7 +1232,13 @@ retry:
}
else
#endif
- new_ptr = malloc (new_size);
+ {
+ omp_memspace_handle_t memspace;
+ memspace = (allocator_data
+ ? allocator_data->memspace
+ : predefined_alloc_mapping[allocator]);
+ new_ptr = MEMSPACE_ALLOC (memspace, new_size);
+ }
if (new_ptr == NULL)
goto fail;
}
@@ -1216,39 +1287,47 @@ retry:
return ret;
}
#endif
- free (data->ptr);
+ {
+ omp_memspace_handle_t was_memspace;
+ was_memspace = (free_allocator_data
+ ? free_allocator_data->memspace
+ : predefined_alloc_mapping[free_allocator]);
+ MEMSPACE_FREE (was_memspace, data->ptr, data->size);
+ }
return ret;
-fail:
- if (allocator_data)
+fail:;
+ int fallback = (allocator_data
+ ? allocator_data->fallback
+ : allocator == omp_default_mem_alloc
+ ? omp_atv_null_fb
+ : omp_atv_default_mem_fb);
+ switch (fallback)
{
- switch (allocator_data->fallback)
- {
- case omp_atv_default_mem_fb:
- if (new_alignment > sizeof (void *)
+ case omp_atv_default_mem_fb:
+ if (new_alignment > sizeof (void *)
#if defined(LIBGOMP_USE_MEMKIND) || defined(LIBGOMP_USE_LIBNUMA)
- || memkind
+ || memkind
#endif
- || (allocator_data
- && allocator_data->pool_size < ~(uintptr_t) 0))
- {
- allocator = omp_default_mem_alloc;
- goto retry;
- }
- /* Otherwise, we've already performed default mem allocation
- and if that failed, it won't succeed again (unless it was
- intermittent. Return NULL then, as that is the fallback. */
- break;
- case omp_atv_null_fb:
- break;
- default:
- case omp_atv_abort_fb:
- gomp_fatal ("Out of memory allocating %lu bytes",
- (unsigned long) size);
- case omp_atv_allocator_fb:
- allocator = allocator_data->fb_data;
+ || allocator_data == NULL
+ || allocator_data->pool_size < ~(uintptr_t) 0)
+ {
+ allocator = omp_default_mem_alloc;
goto retry;
}
+ /* Otherwise, we've already performed default mem allocation
+ and if that failed, it won't succeed again (unless it was
+ intermittent). Return NULL then, as that is the fallback. */
+ break;
+ case omp_atv_null_fb:
+ break;
+ default:
+ case omp_atv_abort_fb:
+ gomp_fatal ("Out of memory allocating %lu bytes",
+ (unsigned long) size);
+ case omp_atv_allocator_fb:
+ allocator = allocator_data->fb_data;
+ goto retry;
}
return NULL;
}
diff --git a/libgomp/basic-allocator.c b/libgomp/basic-allocator.c
new file mode 100644
index 00000000000..35c7439bed6
--- /dev/null
+++ b/libgomp/basic-allocator.c
@@ -0,0 +1,380 @@
+/* Copyright (C) 2023 Free Software Foundation, Inc.
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This is a basic "malloc" implementation intended for use with small,
+ low-latency memories.
+
+ To use this template, define BASIC_ALLOC_PREFIX, and then #include the
+ source file. The other configuration macros are optional.
+
+ The root heap descriptor is stored in the first bytes of the heap, and each
+ free chunk contains a similar descriptor for the next free chunk in the
+ chain.
+
+ The descriptor is two values: offset and size, which describe the
+ location of a chunk of memory available for allocation. The offset is
+ relative to the base of the heap. The special offset value 0xffffffff
+ indicates that the heap (free chain) is locked. The offset and size are
+ 32-bit values so the base alignment can be 8-bytes.
+
+ Memory is allocated to the first free chunk that fits. The free chain
+ is always stored in order of the offset to assist coalescing adjacent
+ chunks. */
+
+#include "libgomp.h"
+
+#ifndef BASIC_ALLOC_PREFIX
+#error "BASIC_ALLOC_PREFIX not defined."
+#endif
+
+#ifndef BASIC_ALLOC_YIELD
+#define BASIC_ALLOC_YIELD
+#endif
+
+#define ALIGN(VAR) (((VAR) + 7) & ~7) /* 8-byte granularity. */
+
+#define fn1(prefix, name) prefix ## _ ## name
+#define fn(prefix, name) fn1 (prefix, name)
+#define basic_alloc_init fn(BASIC_ALLOC_PREFIX,init)
+#define basic_alloc_alloc fn(BASIC_ALLOC_PREFIX,alloc)
+#define basic_alloc_calloc fn(BASIC_ALLOC_PREFIX,calloc)
+#define basic_alloc_free fn(BASIC_ALLOC_PREFIX,free)
+#define basic_alloc_realloc fn(BASIC_ALLOC_PREFIX,realloc)
+
+typedef struct {
+ uint32_t offset;
+ uint32_t size;
+} heapdesc;
+
+void
+basic_alloc_init (char *heap, size_t limit)
+{
+ if (heap == NULL)
+ return;
+
+ /* Initialize the head of the free chain. */
+ heapdesc *root = (heapdesc *) heap;
+ root->offset = ALIGN(1);
+ root->size = limit - root->offset;
+
+ /* And terminate the chain. */
+ heapdesc *next = (heapdesc *) (heap + root->offset);
+ next->offset = 0;
+ next->size = 0;
+}
+
+static void *
+basic_alloc_alloc (char *heap, size_t size)
+{
+ if (heap == NULL)
+ return NULL;
+
+ /* Memory is allocated in N-byte granularity. */
+ size = ALIGN (size);
+
+ /* Acquire a lock on the low-latency heap. */
+ heapdesc root, *root_ptr = (heapdesc *) heap;
+ do
+ {
+ root.offset = __atomic_exchange_n (&root_ptr->offset, 0xffffffff,
+ MEMMODEL_ACQUIRE);
+ if (root.offset != 0xffffffff)
+ {
+ root.size = root_ptr->size;
+ break;
+ }
+ /* Spin. */
+ BASIC_ALLOC_YIELD;
+ }
+ while (1);
+
+ /* Walk the free chain. */
+ heapdesc chunk = root;
+ heapdesc *prev_chunkptr = NULL;
+ heapdesc *chunkptr = (heapdesc *) (heap + chunk.offset);
+ heapdesc onward_chain = *chunkptr;
+ while (chunk.size != 0 && (uint32_t) size > chunk.size)
+ {
+ chunk = onward_chain;
+ prev_chunkptr = chunkptr;
+ chunkptr = (heapdesc *) (heap + chunk.offset);
+ onward_chain = *chunkptr;
+ }
+
+ void *result = NULL;
+ if (chunk.size != 0)
+ {
+ /* Allocation successful. */
+ result = chunkptr;
+
+ /* Update the free chain. */
+ heapdesc stillfree = chunk;
+ stillfree.offset += size;
+ stillfree.size -= size;
+ heapdesc *stillfreeptr = (heapdesc *) (heap + stillfree.offset);
+
+ if (stillfree.size == 0)
+ /* The whole chunk was used. */
+ stillfree = onward_chain;
+ else
+ /* The chunk was split, so restore the onward chain. */
+ *stillfreeptr = onward_chain;
+
+ /* The previous free slot or root now points to stillfree. */
+ if (prev_chunkptr)
+ *prev_chunkptr = stillfree;
+ else
+ root = stillfree;
+ }
+
+ /* Update the free chain root and release the lock. */
+ root_ptr->size = root.size;
+ __atomic_store_n (&root_ptr->offset, root.offset, MEMMODEL_RELEASE);
+
+ return result;
+}
+
+static void *
+basic_alloc_calloc (char *heap, size_t size)
+{
+ /* Memory is allocated in N-byte granularity. */
+ size = ALIGN (size);
+
+ uint64_t *result = basic_alloc_alloc (heap, size);
+ if (result)
+ /* Inline memset in which we know size is a multiple of 8. */
+ for (unsigned i = 0; i < (unsigned) size / 8; i++)
+ result[i] = 0;
+
+ return result;
+}
+
+static void
+basic_alloc_free (char *heap, void *addr, size_t size)
+{
+ /* Memory is allocated in N-byte granularity. */
+ size = ALIGN (size);
+
+ /* Acquire a lock on the low-latency heap. */
+ heapdesc root, *root_ptr = (heapdesc *) heap;
+ do
+ {
+ root.offset = __atomic_exchange_n (&root_ptr->offset, 0xffffffff,
+ MEMMODEL_ACQUIRE);
+ if (root.offset != 0xffffffff)
+ {
+ root.size = root_ptr->size;
+ break;
+ }
+ /* Spin. */
+ }
+ while (1);
+
+ /* Walk the free chain to find where to insert a new entry. */
+ heapdesc chunk = root, prev_chunk = {0};
+ heapdesc *prev_chunkptr = NULL, *prevprev_chunkptr = NULL;
+ heapdesc *chunkptr = (heapdesc *) (heap + chunk.offset);
+ heapdesc onward_chain = *chunkptr;
+ while (chunk.size != 0 && addr > (void *) chunkptr)
+ {
+ prev_chunk = chunk;
+ chunk = onward_chain;
+ prevprev_chunkptr = prev_chunkptr;
+ prev_chunkptr = chunkptr;
+ chunkptr = (heapdesc *) (heap + chunk.offset);
+ onward_chain = *chunkptr;
+ }
+
+ /* Create the new chunk descriptor. */
+ heapdesc newfreechunk;
+ newfreechunk.offset = (uint32_t) ((uintptr_t) addr - (uintptr_t) heap);
+ newfreechunk.size = (uint32_t) size;
+
+ /* Coalesce adjacent free chunks. */
+ if (newfreechunk.offset + size == chunk.offset)
+ {
+ /* Free chunk follows. */
+ newfreechunk.size += chunk.size;
+ chunk = onward_chain;
+ }
+ if (prev_chunkptr)
+ {
+ if (prev_chunk.offset + prev_chunk.size
+ == newfreechunk.offset)
+ {
+ /* Free chunk precedes. */
+ newfreechunk.offset = prev_chunk.offset;
+ newfreechunk.size += prev_chunk.size;
+ addr = heap + prev_chunk.offset;
+ prev_chunkptr = prevprev_chunkptr;
+ }
+ }
+
+ /* Update the free chain in the new and previous chunks. */
+ *(heapdesc *) addr = chunk;
+ if (prev_chunkptr)
+ *prev_chunkptr = newfreechunk;
+ else
+ root = newfreechunk;
+
+ /* Update the free chain root and release the lock. */
+ root_ptr->size = root.size;
+ __atomic_store_n (&root_ptr->offset, root.offset, MEMMODEL_RELEASE);
+
+}
+
+static void *
+basic_alloc_realloc (char *heap, void *addr, size_t oldsize,
+ size_t size)
+{
+ /* Memory is allocated in N-byte granularity. */
+ oldsize = ALIGN (oldsize);
+ size = ALIGN (size);
+
+ if (oldsize == size)
+ return addr;
+
+ /* Acquire a lock on the low-latency heap. */
+ heapdesc root, *root_ptr = (heapdesc *) heap;
+ do
+ {
+ root.offset = __atomic_exchange_n (&root_ptr->offset, 0xffffffff,
+ MEMMODEL_ACQUIRE);
+ if (root.offset != 0xffffffff)
+ {
+ root.size = root_ptr->size;
+ break;
+ }
+ /* Spin. */
+ }
+ while (1);
+
+ /* Walk the free chain. */
+ heapdesc chunk = root;
+ heapdesc *prev_chunkptr = NULL;
+ heapdesc *chunkptr = (heapdesc *) (heap + chunk.offset);
+ heapdesc onward_chain = *chunkptr;
+ while (chunk.size != 0 && (void *) chunkptr < addr)
+ {
+ chunk = onward_chain;
+ prev_chunkptr = chunkptr;
+ chunkptr = (heapdesc *) (heap + chunk.offset);
+ onward_chain = *chunkptr;
+ }
+
+ void *result = NULL;
+ if (size < oldsize)
+ {
+ /* The new allocation is smaller than the old; we can always
+ shrink an allocation in place. */
+ result = addr;
+
+ heapdesc *nowfreeptr = (heapdesc *) (addr + size);
+
+ /* Update the free chain. */
+ heapdesc nowfree;
+ nowfree.offset = (char *) nowfreeptr - heap;
+ nowfree.size = oldsize - size;
+
+ if (nowfree.offset + size == chunk.offset)
+ {
+ /* Coalesce following free chunk. */
+ nowfree.size += chunk.size;
+ *nowfreeptr = onward_chain;
+ }
+ else
+ *nowfreeptr = chunk;
+
+ /* The previous free slot or root now points to nowfree. */
+ if (prev_chunkptr)
+ *prev_chunkptr = nowfree;
+ else
+ root = nowfree;
+ }
+ else if (chunk.size != 0
+ && (char *) addr + oldsize == (char *) chunkptr
+ && chunk.size >= size-oldsize)
+ {
+ /* The new allocation is larger than the old, and we found a
+ large enough free block right after the existing block,
+ so we extend into that space. */
+ result = addr;
+
+ uint32_t delta = size-oldsize;
+
+ /* Update the free chain. */
+ heapdesc stillfree = chunk;
+ stillfree.offset += delta;
+ stillfree.size -= delta;
+ heapdesc *stillfreeptr = (heapdesc *) (heap + stillfree.offset);
+
+ if (stillfree.size == 0)
+ /* The whole chunk was used. */
+ stillfree = onward_chain;
+ else
+ /* The chunk was split, so restore the onward chain. */
+ *stillfreeptr = onward_chain;
+
+ /* The previous free slot or root now points to stillfree. */
+ if (prev_chunkptr)
+ *prev_chunkptr = stillfree;
+ else
+ root = stillfree;
+ }
+ /* Else realloc in-place has failed and result remains NULL. */
+
+ /* Update the free chain root and release the lock. */
+ root_ptr->size = root.size;
+ __atomic_store_n (&root_ptr->offset, root.offset, MEMMODEL_RELEASE);
+
+ if (result == NULL)
+ {
+ /* The allocation could not be extended in place, so we simply
+ allocate fresh memory and move the data. If we can't allocate
+ from low-latency memory then we leave the original alloaction
+ intact and return NULL.
+ We could do a fall-back to main memory, but we don't know what
+ the fall-back trait said to do. */
+ result = basic_alloc_alloc (heap, size);
+ if (result != NULL)
+ {
+ /* Inline memcpy in which we know oldsize is a multiple of 8. */
+ uint64_t *from = addr, *to = result;
+ for (unsigned i = 0; i < (unsigned) oldsize / 8; i++)
+ to[i] = from[i];
+
+ basic_alloc_free (heap, addr, oldsize);
+ }
+ }
+
+ return result;
+}
+
+#undef ALIGN
+#undef fn1
+#undef fn
+#undef basic_alloc_init
+#undef basic_alloc_alloc
+#undef basic_alloc_free
+#undef basic_alloc_realloc
diff --git a/libgomp/config/nvptx/allocator.c b/libgomp/config/nvptx/allocator.c
new file mode 100644
index 00000000000..6014fba177f
--- /dev/null
+++ b/libgomp/config/nvptx/allocator.c
@@ -0,0 +1,120 @@
+/* Copyright (C) 2023 Free Software Foundation, Inc.
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* The low-latency allocators use space reserved in .shared memory when the
+ kernel is launched. The heap is initialized in gomp_nvptx_main and all
+ allocations are forgotten when the kernel exits. Allocations to other
+ memory spaces all use the system malloc syscall.
+
+ The root heap descriptor is stored elsewhere in shared memory, and each
+ free chunk contains a similar descriptor for the next free chunk in the
+ chain.
+
+ The descriptor is two 16-bit values: offset and size, which describe the
+ location of a chunk of memory available for allocation. The offset is
+ relative to the base of the heap. The special value 0xffff, 0xffff
+ indicates that the heap is locked. The descriptor is encoded into a
+ single 32-bit integer so that it may be easily accessed atomically.
+
+ Memory is allocated to the first free chunk that fits. The free chain
+ is always stored in order of the offset to assist coalescing adjacent
+ chunks. */
+
+#include "libgomp.h"
+#include <stdlib.h>
+
+#define BASIC_ALLOC_PREFIX __nvptx_lowlat
+#include "../../basic-allocator.c"
+
+/* There should be some .shared space reserved for us. There's no way to
+ express this magic extern sizeless array in C so use asm. */
+asm (".extern .shared .u8 __nvptx_lowlat_pool[];\n");
+
+static void *
+nvptx_memspace_alloc (omp_memspace_handle_t memspace, size_t size)
+{
+ if (memspace == omp_low_lat_mem_space)
+ {
+ char *shared_pool;
+ asm ("cvta.shared.u64\t%0, __nvptx_lowlat_pool;" : "=r" (shared_pool));
+
+ return __nvptx_lowlat_alloc (shared_pool, size);
+ }
+ else
+ return malloc (size);
+}
+
+static void *
+nvptx_memspace_calloc (omp_memspace_handle_t memspace, size_t size)
+{
+ if (memspace == omp_low_lat_mem_space)
+ {
+ char *shared_pool;
+ asm ("cvta.shared.u64\t%0, __nvptx_lowlat_pool;" : "=r" (shared_pool));
+
+ return __nvptx_lowlat_calloc (shared_pool, size);
+ }
+ else
+ return calloc (1, size);
+}
+
+static void
+nvptx_memspace_free (omp_memspace_handle_t memspace, void *addr, size_t size)
+{
+ if (memspace == omp_low_lat_mem_space)
+ {
+ char *shared_pool;
+ asm ("cvta.shared.u64\t%0, __nvptx_lowlat_pool;" : "=r" (shared_pool));
+
+ __nvptx_lowlat_free (shared_pool, addr, size);
+ }
+ else
+ free (addr);
+}
+
+static void *
+nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr,
+ size_t oldsize, size_t size)
+{
+ if (memspace == omp_low_lat_mem_space)
+ {
+ char *shared_pool;
+ asm ("cvta.shared.u64\t%0, __nvptx_lowlat_pool;" : "=r" (shared_pool));
+
+ return __nvptx_lowlat_realloc (shared_pool, addr, oldsize, size);
+ }
+ else
+ return realloc (addr, size);
+}
+
+#define MEMSPACE_ALLOC(MEMSPACE, SIZE) \
+ nvptx_memspace_alloc (MEMSPACE, SIZE)
+#define MEMSPACE_CALLOC(MEMSPACE, SIZE) \
+ nvptx_memspace_calloc (MEMSPACE, SIZE)
+#define MEMSPACE_REALLOC(MEMSPACE, ADDR, OLDSIZE, SIZE) \
+ nvptx_memspace_realloc (MEMSPACE, ADDR, OLDSIZE, SIZE)
+#define MEMSPACE_FREE(MEMSPACE, ADDR, SIZE) \
+ nvptx_memspace_free (MEMSPACE, ADDR, SIZE)
+
+#include "../../allocator.c"
diff --git a/libgomp/config/nvptx/team.c b/libgomp/config/nvptx/team.c
index af5f3171a47..a78c0d86c46 100644
--- a/libgomp/config/nvptx/team.c
+++ b/libgomp/config/nvptx/team.c
@@ -36,6 +36,12 @@ int __gomp_team_num __attribute__((shared,nocommon));
static void gomp_thread_start (struct gomp_thread_pool *);
+/* There should be some .shared space reserved for us. There's no way to
+ express this magic extern sizeless array in C so use asm. */
+asm (".extern .shared .u8 __nvptx_lowlat_pool[];\n");
+
+/* Defined in basic-allocator.c via config/nvptx/allocator.c. */
+void __nvptx_lowlat_init (void *heap, size_t size);
/* This externally visible function handles target region entry. It
sets up a per-team thread pool and transfers control by calling FN (FN_DATA)
@@ -63,6 +69,18 @@ gomp_nvptx_main (void (*fn) (void *), void *fn_data)
nvptx_thrs = alloca (ntids * sizeof (*nvptx_thrs));
memset (nvptx_thrs, 0, ntids * sizeof (*nvptx_thrs));
+ /* Find the low-latency heap details .... */
+ uint32_t *shared_pool;
+ uint32_t shared_pool_size = 0;
+ asm ("cvta.shared.u64\t%0, __nvptx_lowlat_pool;" : "=r"(shared_pool));
+#if __PTX_ISA_VERSION_MAJOR__ > 4 \
+ || (__PTX_ISA_VERSION_MAJOR__ == 4 && __PTX_ISA_VERSION_MAJOR__ >= 1)
+ asm ("mov.u32\t%0, %%dynamic_smem_size;\n"
+ : "=r"(shared_pool_size));
+#endif
+ __nvptx_lowlat_init (shared_pool, shared_pool_size);
+
+ /* Initialize the thread pool. */
struct gomp_thread_pool *pool = alloca (sizeof (*pool));
pool->threads = alloca (ntids * sizeof (*pool->threads));
for (tid = 0; tid < ntids; tid++)
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 00d4241ae02..65bd430c5a6 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -339,6 +339,11 @@ struct ptx_device
static struct ptx_device **ptx_devices;
+/* OpenMP kernels reserve a small amount of ".shared" space for use by
+ omp_alloc. The size is configured using GOMP_NVPTX_LOWLAT_POOL, but the
+ default is set here. */
+static unsigned lowlat_pool_size = 8 * 1024;
+
static inline struct nvptx_thread *
nvptx_thread (void)
{
@@ -1217,6 +1222,22 @@ GOMP_OFFLOAD_init_device (int n)
instantiated_devices++;
}
+ const char *var_name = "GOMP_NVPTX_LOWLAT_POOL";
+ const char *env_var = secure_getenv (var_name);
+ notify_var (var_name, env_var);
+
+ if (env_var != NULL)
+ {
+ char *endptr;
+ unsigned long val = strtoul (env_var, &endptr, 10);
+ if (endptr == NULL || *endptr != '\0'
+ || errno == ERANGE || errno == EINVAL
+ || val > UINT_MAX)
+ GOMP_PLUGIN_error ("Error parsing %s", var_name);
+ else
+ lowlat_pool_size = val;
+ }
+
pthread_mutex_unlock (&ptx_dev_lock);
return dev != NULL;
@@ -2119,7 +2140,7 @@ GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args)
" [(teams: %u), 1, 1] [(lanes: 32), (threads: %u), 1]\n",
__FUNCTION__, fn_name, teams, threads);
r = CUDA_CALL_NOCHECK (cuLaunchKernel, function, teams, 1, 1,
- 32, threads, 1, 0, NULL, NULL, config);
+ 32, threads, 1, lowlat_pool_size, NULL, NULL, config);
if (r != CUDA_SUCCESS)
GOMP_PLUGIN_fatal ("cuLaunchKernel error: %s", cuda_error (r));
if (reverse_offload)
diff --git a/libgomp/testsuite/libgomp.c/omp_alloc-1.c b/libgomp/testsuite/libgomp.c/omp_alloc-1.c
new file mode 100644
index 00000000000..f4e594f1e98
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/omp_alloc-1.c
@@ -0,0 +1,56 @@
+/* { dg-do run } */
+
+/* Test that omp_alloc returns usable memory. */
+
+#include <omp.h>
+
+#pragma omp requires dynamic_allocators
+
+void
+test (int n, omp_allocator_handle_t allocator)
+{
+ #pragma omp target map(to:n) map(to:allocator)
+ {
+ int *a;
+ a = (int *) omp_alloc (n * sizeof (int), allocator);
+
+ #pragma omp parallel
+ for (int i = 0; i < n; i++)
+ a[i] = i;
+
+ for (int i = 0; i < n; i++)
+ if (a[i] != i)
+ {
+ __builtin_printf ("data mismatch at %i\n", i);
+ __builtin_abort ();
+ }
+
+ omp_free (a, allocator);
+ }
+}
+
+int
+main ()
+{
+ // Smaller than low-latency memory limit
+ test (10, omp_default_mem_alloc);
+ test (10, omp_large_cap_mem_alloc);
+ test (10, omp_const_mem_alloc);
+ test (10, omp_high_bw_mem_alloc);
+ test (10, omp_low_lat_mem_alloc);
+ test (10, omp_cgroup_mem_alloc);
+ test (10, omp_pteam_mem_alloc);
+ test (10, omp_thread_mem_alloc);
+
+ // Larger than low-latency memory limit
+ test (100000, omp_default_mem_alloc);
+ test (100000, omp_large_cap_mem_alloc);
+ test (100000, omp_const_mem_alloc);
+ test (100000, omp_high_bw_mem_alloc);
+ test (100000, omp_low_lat_mem_alloc);
+ test (100000, omp_cgroup_mem_alloc);
+ test (100000, omp_pteam_mem_alloc);
+ test (100000, omp_thread_mem_alloc);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/omp_alloc-2.c b/libgomp/testsuite/libgomp.c/omp_alloc-2.c
new file mode 100644
index 00000000000..e9fd1602946
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/omp_alloc-2.c
@@ -0,0 +1,64 @@
+/* { dg-do run } */
+
+/* Test concurrent and repeated allocations. */
+
+#include <omp.h>
+
+#pragma omp requires dynamic_allocators
+
+void
+test (int n, omp_allocator_handle_t allocator)
+{
+ #pragma omp target map(to:n) map(to:allocator)
+ {
+ int **a;
+ a = (int **) omp_alloc (n * sizeof (int *), allocator);
+
+ #pragma omp parallel for
+ for (int i = 0; i < n; i++)
+ {
+ /*Use 10x to ensure we do activate low-latency fall-back. */
+ a[i] = omp_alloc (sizeof (int) * 10, allocator);
+ a[i][0] = i;
+ }
+
+ for (int i = 0; i < n; i++)
+ if (a[i][0] != i)
+ {
+ __builtin_printf ("data mismatch at %i\n", i);
+ __builtin_abort ();
+ }
+
+ #pragma omp parallel for
+ for (int i = 0; i < n; i++)
+ omp_free (a[i], allocator);
+
+ omp_free (a, allocator);
+ }
+}
+
+int
+main ()
+{
+ // Smaller than low-latency memory limit
+ test (10, omp_default_mem_alloc);
+ test (10, omp_large_cap_mem_alloc);
+ test (10, omp_const_mem_alloc);
+ test (10, omp_high_bw_mem_alloc);
+ test (10, omp_low_lat_mem_alloc);
+ test (10, omp_cgroup_mem_alloc);
+ test (10, omp_pteam_mem_alloc);
+ test (10, omp_thread_mem_alloc);
+
+ // Larger than low-latency memory limit (on aggregate)
+ test (1000, omp_default_mem_alloc);
+ test (1000, omp_large_cap_mem_alloc);
+ test (1000, omp_const_mem_alloc);
+ test (1000, omp_high_bw_mem_alloc);
+ test (1000, omp_low_lat_mem_alloc);
+ test (1000, omp_cgroup_mem_alloc);
+ test (1000, omp_pteam_mem_alloc);
+ test (1000, omp_thread_mem_alloc);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/omp_alloc-3.c b/libgomp/testsuite/libgomp.c/omp_alloc-3.c
new file mode 100644
index 00000000000..792e2200f30
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/omp_alloc-3.c
@@ -0,0 +1,42 @@
+/* { dg-do run } */
+
+/* Stress-test omp_alloc/omp_malloc under concurrency. */
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#pragma omp requires dynamic_allocators
+
+#define N 1000
+
+void
+test (omp_allocator_handle_t allocator)
+{
+ #pragma omp target map(to:allocator)
+ {
+ #pragma omp parallel for
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ {
+ int *p = omp_alloc (sizeof (int), allocator);
+ omp_free (p, allocator);
+ }
+ }
+}
+
+int
+main ()
+{
+ // Smaller than low-latency memory limit
+ test (omp_default_mem_alloc);
+ test (omp_large_cap_mem_alloc);
+ test (omp_const_mem_alloc);
+ test (omp_high_bw_mem_alloc);
+ test (omp_low_lat_mem_alloc);
+ test (omp_cgroup_mem_alloc);
+ test (omp_pteam_mem_alloc);
+ test (omp_thread_mem_alloc);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/omp_alloc-4.c b/libgomp/testsuite/libgomp.c/omp_alloc-4.c
new file mode 100644
index 00000000000..66e13c09234
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/omp_alloc-4.c
@@ -0,0 +1,196 @@
+/* { dg-do run } */
+
+/* Test that low-latency free chains are sound. */
+
+#include <stddef.h>
+#include <omp.h>
+
+#pragma omp requires dynamic_allocators
+
+void
+check (int cond, const char *msg)
+{
+ if (!cond)
+ {
+ __builtin_printf ("%s\n", msg);
+ __builtin_abort ();
+ }
+}
+
+int
+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_allocator_handle_t lowlat = omp_init_allocator (omp_low_lat_mem_space,
+ 1, traits);
+
+ int size = 4;
+
+ char *a = omp_alloc (size, lowlat);
+ char *b = omp_alloc (size, lowlat);
+ char *c = omp_alloc (size, lowlat);
+ char *d = omp_alloc (size, lowlat);
+
+ /* There are headers and padding to account for. */
+ int size2 = size + (b-a);
+ int size3 = size + (c-a);
+ int size4 = size + (d-a) + 100; /* Random larger amount. */
+
+ check (a != NULL && b != NULL && c != NULL && d != NULL,
+ "omp_alloc returned NULL\n");
+
+ omp_free (a, lowlat);
+ char *p = omp_alloc (size, lowlat);
+ check (p == a, "allocate did not reuse first chunk");
+
+ omp_free (b, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == b, "allocate did not reuse second chunk");
+
+ omp_free (c, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == c, "allocate did not reuse third chunk");
+
+ omp_free (a, lowlat);
+ omp_free (b, lowlat);
+ p = omp_alloc (size2, lowlat);
+ check (p == a, "allocate did not coalesce first two chunks");
+
+ omp_free (p, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == a, "allocate did not split first chunk (1)");
+ p = omp_alloc (size, lowlat);
+ check (p == b, "allocate did not split first chunk (2)");
+
+ omp_free (b, lowlat);
+ omp_free (c, lowlat);
+ p = omp_alloc (size2, lowlat);
+ check (p == b, "allocate did not coalesce middle two chunks");
+
+ omp_free (p, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == b, "allocate did not split second chunk (1)");
+ p = omp_alloc (size, lowlat);
+ check (p == c, "allocate did not split second chunk (2)");
+
+ omp_free (b, lowlat);
+ omp_free (a, lowlat);
+ p = omp_alloc (size2, lowlat);
+ check (p == a, "allocate did not coalesce first two chunks, reverse free");
+
+ omp_free (p, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == a, "allocate did not split first chunk (1), reverse free");
+ p = omp_alloc (size, lowlat);
+ check (p == b, "allocate did not split first chunk (2), reverse free");
+
+ omp_free (c, lowlat);
+ omp_free (b, lowlat);
+ p = omp_alloc (size2, lowlat);
+ check (p == b, "allocate did not coalesce second two chunks, reverse free");
+
+ omp_free (p, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == b, "allocate did not split second chunk (1), reverse free");
+ p = omp_alloc (size, lowlat);
+ check (p == c, "allocate did not split second chunk (2), reverse free");
+
+ omp_free (a, lowlat);
+ omp_free (b, lowlat);
+ omp_free (c, lowlat);
+ p = omp_alloc (size3, lowlat);
+ check (p == a, "allocate did not coalesce first three chunks");
+
+ omp_free (p, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == a, "allocate did not split first chunk (1)");
+ p = omp_alloc (size, lowlat);
+ check (p == b, "allocate did not split first chunk (2)");
+ p = omp_alloc (size, lowlat);
+ check (p == c, "allocate did not split first chunk (3)");
+
+ omp_free (b, lowlat);
+ omp_free (c, lowlat);
+ omp_free (d, lowlat);
+ p = omp_alloc (size3, lowlat);
+ check (p == b, "allocate did not coalesce last three chunks");
+
+ omp_free (p, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == b, "allocate did not split second chunk (1)");
+ p = omp_alloc (size, lowlat);
+ check (p == c, "allocate did not split second chunk (2)");
+ p = omp_alloc (size, lowlat);
+ check (p == d, "allocate did not split second chunk (3)");
+
+ omp_free (c, lowlat);
+ omp_free (b, lowlat);
+ omp_free (a, lowlat);
+ p = omp_alloc (size3, lowlat);
+ check (p == a, "allocate did not coalesce first three chunks, reverse free");
+
+ omp_free (p, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == a, "allocate did not split first chunk (1), reverse free");
+ p = omp_alloc (size, lowlat);
+ check (p == b, "allocate did not split first chunk (2), reverse free");
+ p = omp_alloc (size, lowlat);
+ check (p == c, "allocate did not split first chunk (3), reverse free");
+
+ omp_free (d, lowlat);
+ omp_free (c, lowlat);
+ omp_free (b, lowlat);
+ p = omp_alloc (size3, lowlat);
+ check (p == b, "allocate did not coalesce second three chunks, reverse free");
+
+ omp_free (p, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == b, "allocate did not split second chunk (1), reverse free");
+ p = omp_alloc (size, lowlat);
+ check (p == c, "allocate did not split second chunk (2), reverse free");
+ p = omp_alloc (size, lowlat);
+ check (p == d, "allocate did not split second chunk (3), reverse free");
+
+ omp_free (c, lowlat);
+ omp_free (a, lowlat);
+ omp_free (b, lowlat);
+ p = omp_alloc (size3, lowlat);
+ check (p == a, "allocate did not coalesce first three chunks, mixed free");
+
+ omp_free (p, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == a, "allocate did not split first chunk (1), mixed free");
+ p = omp_alloc (size, lowlat);
+ check (p == b, "allocate did not split first chunk (2), mixed free");
+ p = omp_alloc (size, lowlat);
+ check (p == c, "allocate did not split first chunk (3), mixed free");
+
+ omp_free (d, lowlat);
+ omp_free (b, lowlat);
+ omp_free (c, lowlat);
+ p = omp_alloc (size3, lowlat);
+ check (p == b, "allocate did not coalesce second three chunks, mixed free");
+
+ omp_free (p, lowlat);
+ p = omp_alloc (size, lowlat);
+ check (p == b, "allocate did not split second chunk (1), mixed free");
+ p = omp_alloc (size, lowlat);
+ check (p == c, "allocate did not split second chunk (2), mixed free");
+ p = omp_alloc (size, lowlat);
+ check (p == d, "allocate did not split second chunk (3), mixed free");
+
+ omp_free (a, lowlat);
+ omp_free (b, lowlat);
+ omp_free (c, lowlat);
+ omp_free (d, lowlat);
+ p = omp_alloc (size4, lowlat);
+ check (p == a, "allocate did not coalesce all memory");
+ }
+
+ return 0;
+}
+
diff --git a/libgomp/testsuite/libgomp.c/omp_alloc-5.c b/libgomp/testsuite/libgomp.c/omp_alloc-5.c
new file mode 100644
index 00000000000..10805ded6d0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/omp_alloc-5.c
@@ -0,0 +1,63 @@
+/* { dg-do run } */
+
+/* Test calloc with omp_alloc. */
+
+#include <omp.h>
+
+#pragma omp requires dynamic_allocators
+
+void
+test (int n, omp_allocator_handle_t allocator)
+{
+ #pragma omp target map(to:n) map(to:allocator)
+ {
+ int *a;
+ a = (int *) omp_calloc (n, sizeof (int), allocator);
+
+ for (int i = 0; i < n; i++)
+ if (a[i] != 0)
+ {
+ __builtin_printf ("memory not zeroed at %i\n", i);
+ __builtin_abort ();
+ }
+
+ #pragma omp parallel
+ for (int i = 0; i < n; i++)
+ a[i] = i;
+
+ for (int i = 0; i < n; i++)
+ if (a[i] != i)
+ {
+ __builtin_printf ("data mismatch at %i\n", i);
+ __builtin_abort ();
+ }
+
+ omp_free (a, allocator);
+ }
+}
+
+int
+main ()
+{
+ // Smaller than low-latency memory limit
+ test (10, omp_default_mem_alloc);
+ test (10, omp_large_cap_mem_alloc);
+ test (10, omp_const_mem_alloc);
+ test (10, omp_high_bw_mem_alloc);
+ test (10, omp_low_lat_mem_alloc);
+ test (10, omp_cgroup_mem_alloc);
+ test (10, omp_pteam_mem_alloc);
+ test (10, omp_thread_mem_alloc);
+
+ // Larger than low-latency memory limit
+ test (100000, omp_default_mem_alloc);
+ test (100000, omp_large_cap_mem_alloc);
+ test (100000, omp_const_mem_alloc);
+ test (100000, omp_high_bw_mem_alloc);
+ test (100000, omp_low_lat_mem_alloc);
+ test (100000, omp_cgroup_mem_alloc);
+ test (100000, omp_pteam_mem_alloc);
+ test (100000, omp_thread_mem_alloc);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/omp_alloc-6.c b/libgomp/testsuite/libgomp.c/omp_alloc-6.c
new file mode 100644
index 00000000000..66bf69b0455
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/omp_alloc-6.c
@@ -0,0 +1,117 @@
+/* { dg-do run } */
+
+/* Test that low-latency realloc and free chains are sound. */
+
+#include <stddef.h>
+#include <omp.h>
+
+#pragma omp requires dynamic_allocators
+
+void
+check (int cond, const char *msg)
+{
+ if (!cond)
+ {
+ __builtin_printf ("%s\n", msg);
+ __builtin_abort ();
+ }
+}
+
+int
+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_allocator_handle_t lowlat = omp_init_allocator (omp_low_lat_mem_space,
+ 1, traits);
+
+ int size = 16;
+
+ char *a = (char *) omp_alloc (size, lowlat);
+ char *b = (char *) omp_alloc (size, lowlat);
+ char *c = (char *) omp_alloc (size, lowlat);
+ char *d = (char *) omp_alloc (size, lowlat);
+
+ /* There are headers and padding to account for. */
+ int size2 = size + (b-a);
+ int size3 = size + (c-a);
+ int size4 = size + (d-a) + 100; /* Random larger amount. */
+
+ check (a != NULL && b != NULL && c != NULL && d != NULL,
+ "omp_alloc returned NULL\n");
+
+ char *p = omp_realloc (b, size, lowlat, lowlat);
+ check (p == b, "realloc did not reuse same size chunk, no space after");
+
+ p = omp_realloc (b, size-8, lowlat, lowlat);
+ check (p == b, "realloc did not reuse smaller chunk, no space after");
+
+ p = omp_realloc (b, size, lowlat, lowlat);
+ check (p == b, "realloc did not reuse original size chunk, no space after");
+
+ /* Make space after b. */
+ omp_free (c, lowlat);
+
+ p = omp_realloc (b, size, lowlat, lowlat);
+ check (p == b, "realloc did not reuse same size chunk");
+
+ p = omp_realloc (b, size-8, lowlat, lowlat);
+ check (p == b, "realloc did not reuse smaller chunk");
+
+ p = omp_realloc (b, size, lowlat, lowlat);
+ check (p == b, "realloc did not reuse original size chunk");
+
+ p = omp_realloc (b, size+8, lowlat, lowlat);
+ check (p == b, "realloc did not extend in place by a little");
+
+ p = omp_realloc (b, size2, lowlat, lowlat);
+ check (p == b, "realloc did not extend into whole next chunk");
+
+ p = omp_realloc (b, size3, lowlat, lowlat);
+ check (p != b, "realloc did not move b elsewhere");
+ omp_free (p, lowlat);
+
+
+ p = omp_realloc (a, size, lowlat, lowlat);
+ check (p == a, "realloc did not reuse same size chunk, first position");
+
+ p = omp_realloc (a, size-8, lowlat, lowlat);
+ check (p == a, "realloc did not reuse smaller chunk, first position");
+
+ p = omp_realloc (a, size, lowlat, lowlat);
+ check (p == a, "realloc did not reuse original size chunk, first position");
+
+ p = omp_realloc (a, size+8, lowlat, lowlat);
+ check (p == a, "realloc did not extend in place by a little, first position");
+
+ p = omp_realloc (a, size3, lowlat, lowlat);
+ check (p == a, "realloc did not extend into whole next chunk, first position");
+
+ p = omp_realloc (a, size4, lowlat, lowlat);
+ check (p != a, "realloc did not move a elsewhere, first position");
+ omp_free (p, lowlat);
+
+
+ p = omp_realloc (d, size, lowlat, lowlat);
+ check (p == d, "realloc did not reuse same size chunk, last position");
+
+ p = omp_realloc (d, size-8, lowlat, lowlat);
+ check (p == d, "realloc did not reuse smaller chunk, last position");
+
+ p = omp_realloc (d, size, lowlat, lowlat);
+ check (p == d, "realloc did not reuse original size chunk, last position");
+
+ p = omp_realloc (d, size+8, lowlat, lowlat);
+ check (p == d, "realloc did not extend in place by d little, last position");
+
+ /* Larger than low latency memory. */
+ p = omp_realloc (d, 100000000, lowlat, lowlat);
+ check (p == NULL, "realloc did not fail on OOM");
+ }
+
+ return 0;
+}
+
^ permalink raw reply [flat|nested] 7+ messages in thread
* [PATCH v2 2/3] openmp, nvptx: low-lat memory access traits
2023-08-02 17:00 [PATCH v2 0/3] libgomp: OpenMP low-latency omp_alloc Andrew Stubbs
2023-08-02 17:00 ` [PATCH v2 1/3] libgomp, nvptx: low-latency memory allocator Andrew Stubbs
@ 2023-08-02 17:00 ` Andrew Stubbs
2023-08-02 17:00 ` [PATCH v2 3/3] amdgcn, libgomp: low-latency allocator Andrew Stubbs
2 siblings, 0 replies; 7+ messages in thread
From: Andrew Stubbs @ 2023-08-02 17:00 UTC (permalink / raw)
To: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1150 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/omp_alloc-4.c (main): Add access trait.
* testsuite/libgomp.c/omp_alloc-6.c (main): Add access trait.
* testsuite/libgomp.c/omp_alloc-traits.c: New test.
---
libgomp/allocator.c | 16 +++++
libgomp/config/nvptx/allocator.c | 11 +++
libgomp/testsuite/libgomp.c/omp_alloc-4.c | 7 +-
libgomp/testsuite/libgomp.c/omp_alloc-6.c | 7 +-
.../testsuite/libgomp.c/omp_alloc-traits.c | 68 +++++++++++++++++++
5 files changed, 103 insertions(+), 6 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c/omp_alloc-traits.c
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: v2-0002-openmp-nvptx-low-lat-memory-access-traits.patch --]
[-- Type: text/x-patch; name="v2-0002-openmp-nvptx-low-lat-memory-access-traits.patch", Size: 6154 bytes --]
diff --git a/libgomp/allocator.c b/libgomp/allocator.c
index fbf7b1ab061..35b8ec71480 100644
--- a/libgomp/allocator.c
+++ b/libgomp/allocator.c
@@ -56,6 +56,10 @@
#define MEMSPACE_FREE(MEMSPACE, ADDR, SIZE) \
free (((void)(MEMSPACE), (void)(SIZE), (ADDR)))
#endif
+#ifndef MEMSPACE_VALIDATE
+#define MEMSPACE_VALIDATE(MEMSPACE, ACCESS) \
+ (((void)(MEMSPACE), (void)(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.
@@ -507,6 +511,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))
{
@@ -817,6 +825,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))
{
@@ -1063,6 +1075,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 6014fba177f..f19ac28d32a 100644
--- a/libgomp/config/nvptx/allocator.c
+++ b/libgomp/config/nvptx/allocator.c
@@ -108,6 +108,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) \
nvptx_memspace_alloc (MEMSPACE, SIZE)
#define MEMSPACE_CALLOC(MEMSPACE, SIZE) \
@@ -116,5 +125,7 @@ nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr,
nvptx_memspace_realloc (MEMSPACE, ADDR, OLDSIZE, SIZE)
#define MEMSPACE_FREE(MEMSPACE, ADDR, SIZE) \
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/omp_alloc-4.c b/libgomp/testsuite/libgomp.c/omp_alloc-4.c
index 66e13c09234..9d169858151 100644
--- a/libgomp/testsuite/libgomp.c/omp_alloc-4.c
+++ b/libgomp/testsuite/libgomp.c/omp_alloc-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/omp_alloc-6.c b/libgomp/testsuite/libgomp.c/omp_alloc-6.c
index 66bf69b0455..b5f0a296998 100644
--- a/libgomp/testsuite/libgomp.c/omp_alloc-6.c
+++ b/libgomp/testsuite/libgomp.c/omp_alloc-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/omp_alloc-traits.c b/libgomp/testsuite/libgomp.c/omp_alloc-traits.c
new file mode 100644
index 00000000000..6294ba19c16
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/omp_alloc-traits.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;
+}
+
^ permalink raw reply [flat|nested] 7+ messages in thread
* [PATCH v2 3/3] amdgcn, libgomp: low-latency allocator
2023-08-02 17:00 [PATCH v2 0/3] libgomp: OpenMP low-latency omp_alloc Andrew Stubbs
2023-08-02 17:00 ` [PATCH v2 1/3] libgomp, nvptx: low-latency memory allocator Andrew Stubbs
2023-08-02 17:00 ` [PATCH v2 2/3] openmp, nvptx: low-lat memory access traits Andrew Stubbs
@ 2023-08-02 17:00 ` Andrew Stubbs
2 siblings, 0 replies; 7+ messages in thread
From: Andrew Stubbs @ 2023-08-02 17:00 UTC (permalink / raw)
To: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 2052 bytes --]
This implements the OpenMP low-latency memory allocator for AMD GCN using the
small per-team LDS memory (Local Data Store).
Since addresses can now refer to LDS space, the "Global" address space is
no-longer compatible. This patch therefore switches the backend to use
entirely "Flat" addressing (which supports both memories). A future patch
will re-enable "global" instructions for cases where it is known to be safe
to do so.
gcc/ChangeLog:
* config/gcn/gcn-builtins.def (DISPATCH_PTR): New built-in.
* config/gcn/gcn.cc (gcn_init_machine_status): Disable global
addressing.
(gcn_expand_builtin_1): Implement GCN_BUILTIN_DISPATCH_PTR.
libgomp/ChangeLog:
* config/gcn/libgomp-gcn.h (TEAM_ARENA_START): Move to here.
(TEAM_ARENA_FREE): Likewise.
(TEAM_ARENA_END): Likewise.
(GCN_LOWLAT_HEAP): New.
* config/gcn/team.c (LITTLEENDIAN_CPU): New, and import hsa.h.
(__gcn_lowlat_init): New prototype.
(gomp_gcn_enter_kernel): Initialize the low-latency heap.
* libgomp.h (TEAM_ARENA_START): Move to libgomp.h.
(TEAM_ARENA_FREE): Likewise.
(TEAM_ARENA_END): Likewise.
* plugin/plugin-gcn.c (lowlat_size): New variable.
(print_kernel_dispatch): Label the group_segment_size purpose.
(init_environment_variables): Read GOMP_GCN_LOWLAT_POOL.
(create_kernel_dispatch): Pass low-latency head allocation to kernel.
(run_kernel): Use shadow; don't assume values.
* testsuite/libgomp.c/omp_alloc-traits.c: Enable for amdgcn.
* config/gcn/allocator.c: New file.
---
gcc/config/gcn/gcn-builtins.def | 2 +
gcc/config/gcn/gcn.cc | 16 ++-
libgomp/config/gcn/allocator.c | 123 ++++++++++++++++++
libgomp/config/gcn/libgomp-gcn.h | 6 +
libgomp/config/gcn/team.c | 12 ++
libgomp/libgomp.h | 3 -
libgomp/plugin/plugin-gcn.c | 35 ++++-
.../testsuite/libgomp.c/omp_alloc-traits.c | 2 +-
8 files changed, 188 insertions(+), 11 deletions(-)
create mode 100644 libgomp/config/gcn/allocator.c
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: v2-0003-amdgcn-libgomp-low-latency-allocator.patch --]
[-- Type: text/x-patch; name="v2-0003-amdgcn-libgomp-low-latency-allocator.patch", Size: 11951 bytes --]
diff --git a/gcc/config/gcn/gcn-builtins.def b/gcc/config/gcn/gcn-builtins.def
index 636a8e7a1a9..471457d7c23 100644
--- a/gcc/config/gcn/gcn-builtins.def
+++ b/gcc/config/gcn/gcn-builtins.def
@@ -164,6 +164,8 @@ DEF_BUILTIN (FIRST_CALL_THIS_THREAD_P, -1, "first_call_this_thread_p", B_INSN,
_A1 (GCN_BTI_BOOL), gcn_expand_builtin_1)
DEF_BUILTIN (KERNARG_PTR, -1, "kernarg_ptr", B_INSN, _A1 (GCN_BTI_VOIDPTR),
gcn_expand_builtin_1)
+DEF_BUILTIN (DISPATCH_PTR, -1, "dispatch_ptr", B_INSN, _A1 (GCN_BTI_VOIDPTR),
+ gcn_expand_builtin_1)
DEF_BUILTIN (GET_STACK_LIMIT, -1, "get_stack_limit", B_INSN,
_A1 (GCN_BTI_VOIDPTR), gcn_expand_builtin_1)
diff --git a/gcc/config/gcn/gcn.cc b/gcc/config/gcn/gcn.cc
index 02f4dedec42..c4bf0e6ab92 100644
--- a/gcc/config/gcn/gcn.cc
+++ b/gcc/config/gcn/gcn.cc
@@ -109,7 +109,8 @@ gcn_init_machine_status (void)
f = ggc_cleared_alloc<machine_function> ();
- if (TARGET_GCN3)
+ // FIXME: re-enable global addressing with safety for LDS-flat addresses
+ //if (TARGET_GCN3)
f->use_flat_addressing = true;
return f;
@@ -4881,6 +4882,19 @@ gcn_expand_builtin_1 (tree exp, rtx target, rtx /*subtarget */ ,
}
return ptr;
}
+ case GCN_BUILTIN_DISPATCH_PTR:
+ {
+ rtx ptr;
+ if (cfun->machine->args.reg[DISPATCH_PTR_ARG] >= 0)
+ ptr = gen_rtx_REG (DImode,
+ cfun->machine->args.reg[DISPATCH_PTR_ARG]);
+ else
+ {
+ ptr = gen_reg_rtx (DImode);
+ emit_move_insn (ptr, const0_rtx);
+ }
+ return ptr;
+ }
case GCN_BUILTIN_FIRST_CALL_THIS_THREAD_P:
{
/* Stash a marker in the unused upper 16 bits of s[0:1] to indicate
diff --git a/libgomp/config/gcn/allocator.c b/libgomp/config/gcn/allocator.c
new file mode 100644
index 00000000000..151086ea225
--- /dev/null
+++ b/libgomp/config/gcn/allocator.c
@@ -0,0 +1,123 @@
+/* Copyright (C) 2023 Free Software Foundation, Inc.
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* The low-latency allocators use space reserved in LDS memory when the
+ kernel is launched. The heap is initialized in gomp_gcn_enter_kernel and
+ all allocations are forgotten when the kernel exits. Allocations to other
+ memory spaces all use the system malloc syscall.
+
+ The pointers returned are 64-bit "Flat" addresses indistinguishable from
+ regular pointers, but only compatible with the "flat_load/store"
+ instructions. The compiler has been coded to assign default address
+ spaces accordingly.
+
+ LDS memory is not visible to other teams, and therefore may only be used
+ when the memspace access trait is set accordingly. */
+
+#include "libgomp.h"
+#include <stdlib.h>
+
+#define BASIC_ALLOC_PREFIX __gcn_lowlat
+#define BASIC_ALLOC_YIELD asm ("s_sleep 1" ::: "memory")
+#include "../../basic-allocator.c"
+
+/* The low-latency heap is located in LDS memory, but we need the __flat
+ address space for compatibility reasons. */
+#define FLAT_HEAP_PTR \
+ ((void *) (uintptr_t) (void __flat *) (void __lds *) GCN_LOWLAT_HEAP)
+
+static void *
+gcn_memspace_alloc (omp_memspace_handle_t memspace, size_t size)
+{
+ if (memspace == omp_low_lat_mem_space)
+ {
+ char *shared_pool = FLAT_HEAP_PTR;
+
+ return __gcn_lowlat_alloc (shared_pool, size);
+ }
+ else
+ return malloc (size);
+}
+
+static void *
+gcn_memspace_calloc (omp_memspace_handle_t memspace, size_t size)
+{
+ if (memspace == omp_low_lat_mem_space)
+ {
+ char *shared_pool = FLAT_HEAP_PTR;
+
+ return __gcn_lowlat_calloc (shared_pool, size);
+ }
+ else
+ return calloc (1, size);
+}
+
+static void
+gcn_memspace_free (omp_memspace_handle_t memspace, void *addr, size_t size)
+{
+ if (memspace == omp_low_lat_mem_space)
+ {
+ char *shared_pool = FLAT_HEAP_PTR;
+
+ __gcn_lowlat_free (shared_pool, addr, size);
+ }
+ else
+ free (addr);
+}
+
+static void *
+gcn_memspace_realloc (omp_memspace_handle_t memspace, void *addr,
+ size_t oldsize, size_t size)
+{
+ if (memspace == omp_low_lat_mem_space)
+ {
+ char *shared_pool = FLAT_HEAP_PTR;
+
+ return __gcn_lowlat_realloc (shared_pool, addr, oldsize, size);
+ }
+ else
+ return realloc (addr, size);
+}
+
+static inline int
+gcn_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) \
+ gcn_memspace_alloc (MEMSPACE, SIZE)
+#define MEMSPACE_CALLOC(MEMSPACE, SIZE) \
+ gcn_memspace_calloc (MEMSPACE, SIZE)
+#define MEMSPACE_REALLOC(MEMSPACE, ADDR, OLDSIZE, SIZE) \
+ gcn_memspace_realloc (MEMSPACE, ADDR, OLDSIZE, SIZE)
+#define MEMSPACE_FREE(MEMSPACE, ADDR, SIZE) \
+ gcn_memspace_free (MEMSPACE, ADDR, SIZE)
+#define MEMSPACE_VALIDATE(MEMSPACE, ACCESS) \
+ gcn_memspace_validate (MEMSPACE, ACCESS)
+
+#include "../../allocator.c"
diff --git a/libgomp/config/gcn/libgomp-gcn.h b/libgomp/config/gcn/libgomp-gcn.h
index f62b7dde0e7..05b6fb60cc9 100644
--- a/libgomp/config/gcn/libgomp-gcn.h
+++ b/libgomp/config/gcn/libgomp-gcn.h
@@ -33,6 +33,12 @@
#define DEFAULT_GCN_STACK_SIZE (32*1024)
#define DEFAULT_TEAM_ARENA_SIZE (64*1024)
+/* These define the LDS location of data needed by OpenMP. */
+#define TEAM_ARENA_START 16 /* LDS offset of free pointer. */
+#define TEAM_ARENA_FREE 24 /* LDS offset of free pointer. */
+#define TEAM_ARENA_END 32 /* LDS offset of end pointer. */
+#define GCN_LOWLAT_HEAP 40 /* LDS offset of the OpenMP low-latency heap. */
+
struct heap
{
int64_t size;
diff --git a/libgomp/config/gcn/team.c b/libgomp/config/gcn/team.c
index f03207c84e3..4cbcf8fec64 100644
--- a/libgomp/config/gcn/team.c
+++ b/libgomp/config/gcn/team.c
@@ -29,6 +29,12 @@
#include <stdlib.h>
#include <string.h>
+#define LITTLEENDIAN_CPU
+#include "hsa.h"
+
+/* Defined in basic-allocator.c via config/amdgcn/allocator.c. */
+void __gcn_lowlat_init (void *heap, size_t size);
+
static void gomp_thread_start (struct gomp_thread_pool *);
/* This externally visible function handles target region entry. It
@@ -71,6 +77,12 @@ gomp_gcn_enter_kernel (void)
*arena_free = team_arena;
*arena_end = team_arena + kernargs->arena_size_per_team;
+ /* Initialize the low-latency heap. The header is the size. */
+ void __lds *lowlat = (void __lds *)GCN_LOWLAT_HEAP;
+ hsa_kernel_dispatch_packet_t *queue_ptr = __builtin_gcn_dispatch_ptr ();
+ __gcn_lowlat_init ((void*)(uintptr_t)(void __flat*)lowlat,
+ queue_ptr->group_segment_size - GCN_LOWLAT_HEAP);
+
/* Allocate and initialize the team-local-storage data. */
struct gomp_thread *thrs = team_malloc_cleared (sizeof (*thrs)
* numthreads);
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 68f20651fbf..61f3ef41be9 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -114,9 +114,6 @@ extern void gomp_aligned_free (void *);
#ifdef __AMDGCN__
#include "libgomp-gcn.h"
/* The arena is initialized in config/gcn/team.c. */
-#define TEAM_ARENA_START 16 /* LDS offset of free pointer. */
-#define TEAM_ARENA_FREE 24 /* LDS offset of free pointer. */
-#define TEAM_ARENA_END 32 /* LDS offset of end pointer. */
static inline void * __attribute__((malloc))
team_malloc (size_t size)
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index ef22d48da79..482cb966b7e 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -548,6 +548,7 @@ static size_t gcn_kernel_heap_size = DEFAULT_GCN_HEAP_SIZE;
static int team_arena_size = DEFAULT_TEAM_ARENA_SIZE;
static int stack_size = DEFAULT_GCN_STACK_SIZE;
+static int lowlat_size = -1;
/* Flag to decide whether print to stderr information about what is going on.
Set in init_debug depending on environment variables. */
@@ -1014,8 +1015,8 @@ print_kernel_dispatch (struct kernel_dispatch *dispatch, unsigned indent)
fprintf (stderr, "%*sobject: %lu\n", indent, "", dispatch->object);
fprintf (stderr, "%*sprivate_segment_size: %u\n", indent, "",
dispatch->private_segment_size);
- fprintf (stderr, "%*sgroup_segment_size: %u\n", indent, "",
- dispatch->group_segment_size);
+ fprintf (stderr, "%*sgroup_segment_size: %u (low-latency pool)\n", indent,
+ "", dispatch->group_segment_size);
fprintf (stderr, "\n");
}
@@ -1086,6 +1087,10 @@ init_environment_variables (void)
if (tmp)
stack_size = tmp;;
}
+
+ const char *lowlat = secure_getenv ("GOMP_GCN_LOWLAT_POOL");
+ if (lowlat)
+ lowlat_size = atoi (lowlat);
}
/* Return malloc'd string with name of SYMBOL. */
@@ -1903,7 +1908,25 @@ create_kernel_dispatch (struct kernel_info *kernel, int num_teams,
shadow->signal = sync_signal.handle;
shadow->private_segment_size = kernel->private_segment_size;
- shadow->group_segment_size = kernel->group_segment_size;
+
+ if (lowlat_size < 0)
+ {
+ /* Divide the LDS between the number of running teams.
+ Allocate not less than is defined in the kernel metadata. */
+ int teams_per_cu = num_teams / get_cu_count (agent);
+ int LDS_per_team = (teams_per_cu ? 65536 / teams_per_cu : 65536);
+ shadow->group_segment_size
+ = (kernel->group_segment_size > LDS_per_team
+ ? kernel->group_segment_size
+ : LDS_per_team);;
+ }
+ else if (lowlat_size < GCN_LOWLAT_HEAP+8)
+ /* Ensure that there's space for the OpenMP libgomp data. */
+ shadow->group_segment_size = GCN_LOWLAT_HEAP+8;
+ else
+ shadow->group_segment_size = (lowlat_size > 65536
+ ? 65536
+ : lowlat_size);
/* We expect kernels to request a single pointer, explicitly, and the
rest of struct kernargs, implicitly. If they request anything else
@@ -2262,9 +2285,9 @@ run_kernel (struct kernel_info *kernel, void *vars,
print_kernel_dispatch (shadow, 2);
}
- packet->private_segment_size = kernel->private_segment_size;
- packet->group_segment_size = kernel->group_segment_size;
- packet->kernel_object = kernel->object;
+ packet->private_segment_size = shadow->private_segment_size;
+ packet->group_segment_size = shadow->group_segment_size;
+ packet->kernel_object = shadow->object;
packet->kernarg_address = shadow->kernarg_address;
hsa_signal_t s;
s.handle = shadow->signal;
diff --git a/libgomp/testsuite/libgomp.c/omp_alloc-traits.c b/libgomp/testsuite/libgomp.c/omp_alloc-traits.c
index 6294ba19c16..8f6132f0280 100644
--- a/libgomp/testsuite/libgomp.c/omp_alloc-traits.c
+++ b/libgomp/testsuite/libgomp.c/omp_alloc-traits.c
@@ -1,7 +1,7 @@
/* { dg-do run } */
/* { dg-require-effective-target offload_device } */
-/* { dg-xfail-if "not implemented" { ! offload_target_nvptx } } */
+/* { dg-xfail-if "not implemented" { ! { offload_target_nvptx || offload_target_amdgcn } } } */
/* Test that GPU low-latency allocation is limited to team access. */
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [PATCH v2 1/3] libgomp, nvptx: low-latency memory allocator
2023-08-02 17:00 ` [PATCH v2 1/3] libgomp, nvptx: low-latency memory allocator Andrew Stubbs
@ 2023-09-08 9:04 ` Tobias Burnus
2023-11-29 16:25 ` Andrew Stubbs
0 siblings, 1 reply; 7+ messages in thread
From: Tobias Burnus @ 2023-09-08 9:04 UTC (permalink / raw)
To: Andrew Stubbs, gcc-patches
Hi Andrew,
some early comments. I think in general, the direction/patches are fine,
but I have some comments:
On 02.08.23 19:00, Andrew Stubbs wrote:
> This patch adds support for allocating low-latency ".shared" memory on
> NVPTX GPU device, via the omp_low_lat_mem_space and omp_alloc. The memory
> can be allocated, reallocated, and freed using a basic but fast algorithm,
> is thread safe and the size of the low-latency heap can be configured using
> the GOMP_NVPTX_LOWLAT_POOL environment variable.
>
> The use of the PTX dynamic_smem_size feature means that low-latency allocator
> will not work with the PTX 3.1 multilib.
This probably fits better to 2/3 in the series, but you really should
document the nvptx part, namely:
- that omp_low_lat_mem_space is supported on nvptx
- its limitations (access is restricted to the contention group, i.e.
all threads of a team) → implication on the supported allocators.
- the default size of this memory (8 kiB) and the GOMP_NVPTX_LOWLAT_POOL
environment variable, possibly with mentioning that there is some
internal overhead* which is worsen when using high alignment values. (*
– due to basic_allocator book keeping and for storing pointer to the
OpenMP allocator struct.)
- if I understand it correctly, our default build supports sm_30 and
uses PTX ISA version 3.1 for it. If so, I think we should mention that
nvptx GCC has to be configured with with-arch=sm_... >= sm_53 (=
supported version >=4.1) and, during compilation, no -march= < that
configure-time value may be specified. (Cf. also
https://gcc.gnu.org/install/specific.html#nvptx-x-none )
I think this best fits into
https://gcc.gnu.org/onlinedocs/libgomp/nvptx.html – but one could also
argue that it should be put elsewhere.
It probably makes sense to add a 'See also:' to
https://gcc.gnu.org/onlinedocs/libgomp/Memory-allocation.html pointing
to https://gcc.gnu.org/onlinedocs/libgomp/Offload-Target-Specifics.html
* * *
BTW: I think the following should be "...MINOR__ >= 1":
> +#if __PTX_ISA_VERSION_MAJOR__ > 4 \
> + || (__PTX_ISA_VERSION_MAJOR__ == 4 && __PTX_ISA_VERSION_MAJOR__ >= 1)
* * *
Regarding patch 2/3 and MEMSPACE_VALIDATE.
In general, I wonder how to handle memory spaces (and traits) that
aren't supported. Namely, when to return 0L and when to silently use
ignore the trait / use another memory space.
The current omp_init_allocator code only returns omp_null_allocator for
invalid value – or for pinned memory (as it is unsupported). [RFC: Shall
we keep doing so – or return omp_null_mem_alloc more often? →
https://gcc.gnu.org/PR111044 for this question, improving libmemkind
usage, and extending the allocator-related documentation.]
As we do it on the host, I think auto-fallback to omp_default_mem_space
is is also find for nvptx (and gcn), but not as done in 2/3 but slightly
different:
(a) In omp_init_allocator, there should be a check whether it is
supported, if not, we can fallback to using default memory space. (In
line with the current code host + 1/2+2/3 nvptx behaviour.)
Note: That's not the same as the current 2/3 patch. Currently, if
MEMSPACE_VALIDATE fails, a retry is attempted – but the outcome depends
on the value for 'fallback'. When changing the memory space during
omp_init_allocator, only failed 'malloc' will give abort with abort_fb.
(b) For nvptx_memspace_validate, I think an additional check should be
done based on the __PTX_ISA_VERSION* as it feels off if plugin first
claims support for it but later unconditionally uses malloc at runtime.
(c) We also need to handle omp_low_lat_mem_alloc. I think the spec
implies access:all but nvptx/gcn only support cgroup (+ pteams +
thread), potentially leading to wrong code. Example (hopefully, I got
the syntax right:
#pragma omp target uses_allocator(omp_low_lat_mem_alloc)
#pragma omp teams firstprivate(var) allocate(omp_low_lat_mem_alloc: var)
#pragma omp distribute parallel for
...
#omp atomic ...
... var ...
The current 2/3 checks in alloc/calloc/realloc only cover user-defined
allocators; if we move the check for user-defined allocators to
omp_init_allocator, we actually only need to handle predefined
allocators in alloc/calloc/realloc.
And finally: As mentioned off list, I believe that for the patch 2/3,
the pteam should be cgroup (contention group), i.e. about all threads of
a team / implicit parallel and not only the innermost parallel (pteam).
That actually matches the "access != all" check, but I think "access =
cgroup" should also be tested for in the testsuite.
* * *
3/3 patch for GCN: I think the situation is similar, except that there
is no ISA version issue and most is handled by 1/3 and 2/3 such that
only updating documentation remains.
* * *
> libgomp/ChangeLog:
>
> * allocator.c (MEMSPACE_ALLOC): New macro.
> (MEMSPACE_CALLOC): New macro.
> (MEMSPACE_REALLOC): New macro.
> (MEMSPACE_FREE): New macro.
BTW: You could (but are not required to) combine multiple macro/function
names to a single '(name1, name2, ...):' if all have the same
description, which saves (only) a few lines.
> --- a/libgomp/allocator.c
> +++ b/libgomp/allocator.c
...
> +#define MEMSPACE_CALLOC(MEMSPACE, SIZE) \
> + calloc (1, (((void)(MEMSPACE), (SIZE))))
I am not sure whether I like that there is no 'size_t nmemb' argument
(as it is always 1) or not (given that stdlib.c's calloc has size and nmemb).
(Hence, I am fine with either.)
> +/* Map the predefined allocators to the correct memory space.
> + The index to this table is the omp_allocator_handle_t enum value.
> + When the user calls omp_alloc with a predefined allocator this
> + table determines what memory they get. */
> +static const omp_memspace_handle_t predefined_alloc_mapping[] = {
> + omp_default_mem_space, /* omp_null_allocator. */
> + omp_default_mem_space, /* omp_default_mem_alloc. */
The first line is misleading: omp_null_allocator uses the allocator
associated with def-allocator-var ICV, i.e. any of the predefined
allocators might be used.
As omp_null_allocator is mapped to the def-allocator-var ICV or
(if unset) to the omp_default_mem_space, there should not be any
access to predefined_alloc_mapping[omp_null_allocator].
Still, the code is confusing. I think at a comment is required.
And: Either we still keep that superfluous line or we access
the array as
predefined_alloc_mapping[predef_alloc - 1]
which IMO requires a macro or inline function to avoid having a
puzzling "-1" in the code.
* * *
I wonder whether we should have a static assert checking for
ARRAY_SIZE (predefined_alloc_mapping) == omp_max_predefined_alloc
(or '+ 1', depending how we deal with omp_null_allocator) to ensure
better consistency. (The value is #defined in allocate.c not in omp.h,
but a static assert at least can catch one mismatch.)
[While static_assert is only in C2X alias C23, _Static_assert exists before.
(I think since C11 but GCC also accepts it with -std=c98; GCC >= 9
permits omitting the second/string argument of _Static_assert.
And with 2nd arg, it already works with GCC 7 (= oldest GCC at hand).
[I wrote ARRAY_SIZE in the sense of '#define ARRAY_SIZE(a) (sizeof (a) / sizeof ((a)[0]))',
as e.g. defined include/libiberty.h (that file is not included in libgomp/.]
* * *
> ...
> + omp_low_lat_mem_space, /* omp_cgroup_mem_alloc. */
> + omp_low_lat_mem_space, /* omp_pteam_mem_alloc. */
> + omp_low_lat_mem_space, /* omp_thread_mem_alloc. */
I think there should be a comment like: /* Implementation choice: */
Thus, when later revisiting it, it is clear that it can be changed.
I think it would make sense to document the used memory space in libgomp.texi
alias https://gcc.gnu.org/onlinedocs/libgomp/OMP_005fALLOCATOR.html
Namely replacing the dash in the table by, e.g.,
'omp_low_lat_mem_space' (implementation choice)'
or something like that.
(I personally like that the documentation makes clear
- if sensibly possible - whether a piece of information in a
compiler documentation is generic (matches the spec) or is
an implementation choice. In any case, the documentation should
match what's implemented.)
Note: That omp_low_lat_mem_space == omp_default_mem_space
is already documented at
https://gcc.gnu.org/onlinedocs/libgomp/Memory-allocation.html
Maybe the wording needs to be tweaked now as nvptx + gcn actually
handle the low-lat memory space differently. (While on the host,
a failed 'malloc' is now repeated once, which is not really observable.)
Or it is fine and "See also" to the target-specific section is enough.
(BTW: Wording/documentation suggestions and/or patches are welcome!)
* * *
> + /* Otherwise, we've already performed default mem allocation
> + and if that failed, it won't succeed again (unless it was
> + intermittent). Return NULL then, as that is the fallback. */
Thanks for adding the missing ')'. (Twice.)
And thanks for the patch set in general.
Tobias
-----------------
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
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [PATCH v2 1/3] libgomp, nvptx: low-latency memory allocator
2023-09-08 9:04 ` Tobias Burnus
@ 2023-11-29 16:25 ` Andrew Stubbs
2023-11-30 11:59 ` Tobias Burnus
0 siblings, 1 reply; 7+ messages in thread
From: Andrew Stubbs @ 2023-11-29 16:25 UTC (permalink / raw)
To: Tobias Burnus, gcc-patches
On 08/09/2023 10:04, Tobias Burnus wrote:
> Regarding patch 2/3 and MEMSPACE_VALIDATE.
>
> In general, I wonder how to handle memory spaces (and traits) that
> aren't supported. Namely, when to return 0L and when to silently use
> ignore the trait / use another memory space.
>
> The current omp_init_allocator code only returns omp_null_allocator for
> invalid value – or for pinned memory (as it is unsupported). [RFC: Shall
> we keep doing so – or return omp_null_mem_alloc more often? →
> https://gcc.gnu.org/PR111044 for this question, improving libmemkind
> usage, and extending the allocator-related documentation.]
>
> As we do it on the host, I think auto-fallback to omp_default_mem_space
> is is also find for nvptx (and gcn), but not as done in 2/3 but slightly
> different:
>
> (a) In omp_init_allocator, there should be a check whether it is
> supported, if not, we can fallback to using default memory space. (In
> line with the current code host + 1/2+2/3 nvptx behaviour.)
>
> Note: That's not the same as the current 2/3 patch. Currently, if
> MEMSPACE_VALIDATE fails, a retry is attempted – but the outcome depends
> on the value for 'fallback'. When changing the memory space during
> omp_init_allocator, only failed 'malloc' will give abort with abort_fb.
>
> (b) For nvptx_memspace_validate, I think an additional check should be
> done based on the __PTX_ISA_VERSION* as it feels off if plugin first
> claims support for it but later unconditionally uses malloc at runtime.
I have looked at moving the MEMSPACE_VALIDATE call into
omp_init_allocator so that we can't even create allocators that would be
invalid, but that changes the semantics of the fall-back traits. Here's
the example from testcase omp_alloc-traits.c:
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);
/* ... */
void *b = omp_alloc (1, lowlat_all);
With my patch as proposed, "lowlat_all" is a valid allocator, but
allocating low-latency memory fails in omp_alloc, so "b" ends up NULL
(the fall-back setting).
With the proposed change, "lowlat_all" becomes omp_null_allocator, and
"b" is non-NULL, pointing to default memory. This is probably surprising
to the user because they thought they specified "low-latency or nothing".
Another option would be to create a custom allocator that goes straight
to the fall-back somehow (we could invent an internal value
"ompx_fallback_mem_space", or some such).
What is the desired behaviour in this case? I'm not sure that what the
OpenMP spec actually says matches what the intention seems to have been
with fallbacks.
> (c) We also need to handle omp_low_lat_mem_alloc. I think the spec
> implies access:all but nvptx/gcn only support cgroup (+ pteams +
> thread), potentially leading to wrong code.
If we're not allowed to default to "cgroup" then surely
omp_low_lat_mem_alloc is useless on all GPU devices (that I am aware of)
on all toolchains? There may be some use on some specialist NUMA host
devices, but that's it.
Andrew
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [PATCH v2 1/3] libgomp, nvptx: low-latency memory allocator
2023-11-29 16:25 ` Andrew Stubbs
@ 2023-11-30 11:59 ` Tobias Burnus
0 siblings, 0 replies; 7+ messages in thread
From: Tobias Burnus @ 2023-11-30 11:59 UTC (permalink / raw)
To: Andrew Stubbs, gcc-patches
Hi Andrew,
On 29.11.23 17:25, Andrew Stubbs wrote:
> On 08/09/2023 10:04, Tobias Burnus wrote:
>
>> Regarding patch 2/3 and MEMSPACE_VALIDATE.
>>
>> In general, I wonder how to handle memory spaces (and traits) that
>> aren't supported. Namely, when to return 0L and when to silently use
>> ignore the trait / use another memory space.
>>
>> The current omp_init_allocator code only returns omp_null_allocator for
>> invalid value – or for pinned memory (as it is unsupported). [RFC: Shall
>> we keep doing so – or return omp_null_mem_alloc more often? →
>> https://gcc.gnu.org/PR111044 for this question, improving libmemkind
>> usage, and extending the allocator-related documentation.]
>>
>> As we do it on the host, I think auto-fallback to omp_default_mem_space
>> is is also find for nvptx (and gcn), but not as done in 2/3 but slightly
>> different:
>>
>> (a) In omp_init_allocator, there should be a check whether it is
>> supported, if not, we can fallback to using default memory space. (In
>> line with the current code host + 1/2+2/3 nvptx behaviour.)
>>
>> Note: That's not the same as the current 2/3 patch. Currently, if
>> MEMSPACE_VALIDATE fails, a retry is attempted – but the outcome depends
>> on the value for 'fallback'. When changing the memory space during
>> omp_init_allocator, only failed 'malloc' will give abort with abort_fb.
>>
>> (b) For nvptx_memspace_validate, I think an additional check should be
>> done based on the __PTX_ISA_VERSION* as it feels off if plugin first
>> claims support for it but later unconditionally uses malloc at runtime.
>
> I have looked at moving the MEMSPACE_VALIDATE call into
> omp_init_allocator so that we can't even create allocators that would
> be invalid, but that changes the semantics of the fall-back traits.
> Here's the example from testcase omp_alloc-traits.c:
>
> 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);
>
> /* ... */
>
> void *b = omp_alloc (1, lowlat_all);
>
> With my patch as proposed, "lowlat_all" is a valid allocator, but
> allocating low-latency memory fails in omp_alloc, so "b" ends up NULL
> (the fall-back setting).
>
> With the proposed change, "lowlat_all" becomes omp_null_allocator, and
> "b" is non-NULL, pointing to default memory. This is probably
> surprising to the user because they thought they specified
> "low-latency or nothing".
Well, a proper code would do instead:
omp_allocator_handle_t lowlat_all
= omp_init_allocator (omp_low_lat_mem_space, 2, traits_all);
if (lowlat_all == omp_null_allocator)
{
omp_allocator_handle_t lowlat_all
= omp_init_allocator (omp_low_lat_mem_space, 2, traits_all);
/* At least, preserve the traits (okay, not very useful here). */
lowlat_all
= omp_init_allocator (omp_default_mem_space, 2, traits_all);
/* Giving up: */
if (lowlat_all == omp_null_allocator)
lowlat_all = omp_default_mem_alloc;
}
OpenMP explicitly states:
"if an allocator based on the requirements cannot be created then the special omp_null_allocator handle is returned."
Thus, if the user is surprised it is their fault!
>> (c) We also need to handle omp_low_lat_mem_alloc. I think the spec
>> implies access:all but nvptx/gcn only support cgroup (+ pteams +
>> thread), potentially leading to wrong code.
> If we're not allowed to default to "cgroup" then surely
> omp_low_lat_mem_alloc is useless on all GPU devices (that I am aware
> of) on all toolchains? There may be some use on some specialist NUMA
> host devices, but that's it.
Granted, but a user-specified allocator would still work, wouldn't it?
Thus, the question is whether being a bit less safe but more useful or
being safer but slower/less useful makes more sense – in either case, an
alert user that has read the documentation would know what to do, but a
programmer might not be aware of such issues and a mere user does not
know what a programmer did.
I think I agree with Jakub that "without the user telling us that" (the
only access is within a team, "i.e. requesting cgroup access), we don't
know what it will be used for and so we need to assume worst".
BTW: Unless I missed something, LLVM does not yet support it:
KMP_WARNING(OmpNoAllocator, "omp_low_lat_mem_alloc");
Tobias
-----------------
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
^ permalink raw reply [flat|nested] 7+ messages in thread
end of thread, other threads:[~2023-11-30 11:59 UTC | newest]
Thread overview: 7+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-08-02 17:00 [PATCH v2 0/3] libgomp: OpenMP low-latency omp_alloc Andrew Stubbs
2023-08-02 17:00 ` [PATCH v2 1/3] libgomp, nvptx: low-latency memory allocator Andrew Stubbs
2023-09-08 9:04 ` Tobias Burnus
2023-11-29 16:25 ` Andrew Stubbs
2023-11-30 11:59 ` Tobias Burnus
2023-08-02 17:00 ` [PATCH v2 2/3] openmp, nvptx: low-lat memory access traits Andrew Stubbs
2023-08-02 17:00 ` [PATCH v2 3/3] amdgcn, libgomp: low-latency allocator Andrew Stubbs
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).