4e47e2f833
PR libgomp/120444
include/ChangeLog:
* cuda/cuda.h (cuMemsetD8, cuMemsetD8Async): Declare.
libgomp/ChangeLog:
* libgomp-plugin.h (GOMP_OFFLOAD_memset): Declare.
* libgomp.h (struct gomp_device_descr): Add memset_func.
* libgomp.map (GOMP_6.0.1): Add omp_target_memset{,_async}.
* libgomp.texi (Device Memory Routines): Document them.
* omp.h.in (omp_target_memset, omp_target_memset_async): Declare.
* omp_lib.f90.in (omp_target_memset, omp_target_memset_async):
Add interfaces.
* omp_lib.h.in (omp_target_memset, omp_target_memset_async): Likewise.
* plugin/cuda-lib.def: Add cuMemsetD8.
* plugin/plugin-gcn.c (struct hsa_runtime_fn_info): Add
hsa_amd_memory_fill_fn.
(init_hsa_runtime_functions): DLSYM_OPT_FN load it.
(GOMP_OFFLOAD_memset): New.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_memset): New.
* target.c (omp_target_memset_int, omp_target_memset,
omp_target_memset_async_helper, omp_target_memset_async): New.
(gomp_load_plugin_for_device): Add DLSYM (memset).
* testsuite/libgomp.c-c++-common/omp_target_memset.c: New test.
* testsuite/libgomp.c-c++-common/omp_target_memset-2.c: New test.
* testsuite/libgomp.c-c++-common/omp_target_memset-3.c: New test.
* testsuite/libgomp.fortran/omp_target_memset.f90: New test.
* testsuite/libgomp.fortran/omp_target_memset-2.f90: New test.
62 lines
1.6 KiB
C
62 lines
1.6 KiB
C
// PR libgomp/120444
|
|
|
|
#include <omp.h>
|
|
|
|
int main()
|
|
{
|
|
for (int dev = omp_initial_device; dev < omp_get_num_devices (); dev++)
|
|
{
|
|
char *ptr = (char *) omp_target_alloc (sizeof(int) * 1024, dev);
|
|
|
|
/* Play also around with the alignment - as hsa_amd_memory_fill operates
|
|
on multiples of 4 bytes (uint32_t). */
|
|
|
|
for (int start = 0; start < 32; start++)
|
|
for (int tail = 0; tail < 32; tail++)
|
|
{
|
|
unsigned char val = '0' + start + tail;
|
|
void *ptr2 = omp_target_memset (ptr + start, val,
|
|
1024 - start - tail, dev);
|
|
if (ptr + start != ptr2)
|
|
__builtin_abort ();
|
|
|
|
#pragma omp target device(dev) is_device_ptr(ptr)
|
|
for (int i = start; i < 1024 - start - tail; i++)
|
|
if (ptr[i] != val)
|
|
__builtin_abort ();
|
|
|
|
}
|
|
|
|
/* Check 'small' values for correctness. */
|
|
|
|
for (int start = 0; start < 32; start++)
|
|
for (int size = 0; size <= 64 + 32; size++)
|
|
{
|
|
omp_target_memset (ptr, 'a' - 2, 1024, dev);
|
|
|
|
unsigned char val = '0' + start + size % 32;
|
|
void *ptr2 = omp_target_memset (ptr + start, val, size, dev);
|
|
|
|
if (ptr + start != ptr2)
|
|
__builtin_abort ();
|
|
|
|
if (size == 0)
|
|
continue;
|
|
|
|
#pragma omp target device(dev) is_device_ptr(ptr)
|
|
{
|
|
for (int i = 0; i < start; i++)
|
|
if (ptr[i] != 'a' - 2)
|
|
__builtin_abort ();
|
|
for (int i = start; i < start + size; i++)
|
|
if (ptr[i] != val)
|
|
__builtin_abort ();
|
|
for (int i = start + size + 1; i < 1024; i++)
|
|
if (ptr[i] != 'a' - 2)
|
|
__builtin_abort ();
|
|
}
|
|
}
|
|
|
|
omp_target_free (ptr, dev);
|
|
}
|
|
}
|