@@ -51,6 +51,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. */
@@ -279,6 +282,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))
{
@@ -505,6 +512,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))
{
@@ -672,6 +683,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))
{
@@ -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"
@@ -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;
@@ -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;
new file mode 100644
@@ -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;
+}
+