blob: 01909f854b69711c800dce804f38583b8d8c3961 (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
|
// 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);
}
}
|
