mirror of
https://forge.sourceware.org/marek/gcc.git
synced 2026-02-22 12:00:11 -05:00
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 no longer works (but omp_cgroup_mem_alloc still does). libgomp/ChangeLog: * allocator.c (MEMSPACE_VALIDATE): New macro. (omp_init_allocator): Use MEMSPACE_VALIDATE. (omp_aligned_alloc): Use OMP_LOW_LAT_MEM_ALLOC_INVALID. (omp_aligned_calloc): Likewise. (omp_realloc): Likewise. * config/nvptx/allocator.c (nvptx_memspace_validate): New function. (MEMSPACE_VALIDATE): New macro. (OMP_LOW_LAT_MEM_ALLOC_INVALID): New define. * libgomp.texi: Document low-latency implementation details. * testsuite/libgomp.c/omp_alloc-1.c (main): Add gnu_lowlat. * testsuite/libgomp.c/omp_alloc-2.c (main): Add gnu_lowlat. * testsuite/libgomp.c/omp_alloc-3.c (main): Add gnu_lowlat. * testsuite/libgomp.c/omp_alloc-4.c (main): Add access trait. * testsuite/libgomp.c/omp_alloc-5.c (main): Add gnu_lowlat. * testsuite/libgomp.c/omp_alloc-6.c (main): Add access trait. * testsuite/libgomp.c/omp_alloc-traits.c: New test.
67 lines
1.6 KiB
C
67 lines
1.6 KiB
C
/* { 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 ()
|
|
{
|
|
/* omp_low_lat_mem_alloc doesn't actually get low-latency memory on GPU. */
|
|
omp_allocator_handle_t gpu_lowlat = 0;
|
|
#pragma omp target map(from:gpu_lowlat)
|
|
{
|
|
omp_alloctrait_t traits[1] = { { omp_atk_access, omp_atv_cgroup } };
|
|
gpu_lowlat = omp_init_allocator (omp_low_lat_mem_space, 1, traits);
|
|
}
|
|
|
|
// 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, gpu_lowlat);
|
|
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, gpu_lowlat);
|
|
test (100000, omp_cgroup_mem_alloc);
|
|
test (100000, omp_pteam_mem_alloc);
|
|
test (100000, omp_thread_mem_alloc);
|
|
|
|
return 0;
|
|
}
|