diff --git a/libgomp/allocator.c b/libgomp/allocator.c index 029d0d40a36..48ab0782e6b 100644 --- a/libgomp/allocator.c +++ b/libgomp/allocator.c @@ -54,6 +54,9 @@ #define MEMSPACE_FREE(MEMSPACE, ADDR, SIZE, PIN) \ (PIN ? NULL : free (ADDR)) #endif +#ifndef MEMSPACE_VALIDATE +#define MEMSPACE_VALIDATE(MEMSPACE, ACCESS) 1 +#endif /* Map the predefined allocators to the correct memory space. The index to this table is the omp_allocator_handle_t enum value. */ @@ -438,6 +441,10 @@ retry: if (__builtin_add_overflow (size, new_size, &new_size)) goto fail; + if (allocator_data + && !MEMSPACE_VALIDATE (allocator_data->memspace, allocator_data->access)) + goto fail; + if (__builtin_expect (allocator_data && allocator_data->pool_size < ~(uintptr_t) 0, 0)) { @@ -733,6 +740,10 @@ retry: if (__builtin_add_overflow (size_temp, new_size, &new_size)) goto fail; + if (allocator_data + && !MEMSPACE_VALIDATE (allocator_data->memspace, allocator_data->access)) + goto fail; + if (__builtin_expect (allocator_data && allocator_data->pool_size < ~(uintptr_t) 0, 0)) { @@ -964,6 +975,10 @@ retry: goto fail; old_size = data->size; + if (allocator_data + && !MEMSPACE_VALIDATE (allocator_data->memspace, allocator_data->access)) + goto fail; + if (__builtin_expect (allocator_data && allocator_data->pool_size < ~(uintptr_t) 0, 0)) { diff --git a/libgomp/config/nvptx/allocator.c b/libgomp/config/nvptx/allocator.c index f740b97f6ac..0102680b717 100644 --- a/libgomp/config/nvptx/allocator.c +++ b/libgomp/config/nvptx/allocator.c @@ -358,6 +358,15 @@ nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr, return realloc (addr, size); } +static inline int +nvptx_memspace_validate (omp_memspace_handle_t memspace, unsigned access) +{ + /* Disallow use of low-latency memory when it must be accessible by + all threads. */ + return (memspace != omp_low_lat_mem_space + || access != omp_atv_all); +} + #define MEMSPACE_ALLOC(MEMSPACE, SIZE, PIN) \ nvptx_memspace_alloc (MEMSPACE, SIZE) #define MEMSPACE_CALLOC(MEMSPACE, SIZE, PIN) \ @@ -366,5 +375,7 @@ nvptx_memspace_realloc (omp_memspace_handle_t memspace, void *addr, nvptx_memspace_realloc (MEMSPACE, ADDR, OLDSIZE, SIZE) #define MEMSPACE_FREE(MEMSPACE, ADDR, SIZE, PIN) \ nvptx_memspace_free (MEMSPACE, ADDR, SIZE) +#define MEMSPACE_VALIDATE(MEMSPACE, ACCESS) \ + nvptx_memspace_validate (MEMSPACE, ACCESS) #include "../../allocator.c" diff --git a/libgomp/testsuite/libgomp.c/allocators-4.c b/libgomp/testsuite/libgomp.c/allocators-4.c index 9fa6aa1624f..cae27ea33c1 100644 --- a/libgomp/testsuite/libgomp.c/allocators-4.c +++ b/libgomp/testsuite/libgomp.c/allocators-4.c @@ -23,10 +23,11 @@ main () #pragma omp target { /* Ensure that the memory we get *is* low-latency with a null-fallback. */ - omp_alloctrait_t traits[1] - = { { omp_atk_fallback, omp_atv_null_fb } }; + omp_alloctrait_t traits[2] + = { { omp_atk_fallback, omp_atv_null_fb }, + { omp_atk_access, omp_atv_pteam } }; omp_allocator_handle_t lowlat = omp_init_allocator (omp_low_lat_mem_space, - 1, traits); + 2, traits); int size = 4; diff --git a/libgomp/testsuite/libgomp.c/allocators-6.c b/libgomp/testsuite/libgomp.c/allocators-6.c index 90bf73095ef..c03233df582 100644 --- a/libgomp/testsuite/libgomp.c/allocators-6.c +++ b/libgomp/testsuite/libgomp.c/allocators-6.c @@ -23,10 +23,11 @@ main () #pragma omp target { /* Ensure that the memory we get *is* low-latency with a null-fallback. */ - omp_alloctrait_t traits[1] - = { { omp_atk_fallback, omp_atv_null_fb } }; + omp_alloctrait_t traits[2] + = { { omp_atk_fallback, omp_atv_null_fb }, + { omp_atk_access, omp_atv_pteam } }; omp_allocator_handle_t lowlat = omp_init_allocator (omp_low_lat_mem_space, - 1, traits); + 2, traits); int size = 16; diff --git a/libgomp/testsuite/libgomp.c/allocators-7.c b/libgomp/testsuite/libgomp.c/allocators-7.c new file mode 100644 index 00000000000..a0a738b1d1d --- /dev/null +++ b/libgomp/testsuite/libgomp.c/allocators-7.c @@ -0,0 +1,68 @@ +/* { dg-do run } */ + +/* { dg-require-effective-target offload_device } */ +/* { dg-xfail-if "not implemented" { ! offload_target_nvptx } } */ + +/* Test that GPU low-latency allocation is limited to team access. */ + +#include +#include + +#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; +} +