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
|
// clang-format off
// RUN: %libomptarget-compilexx-generic
// RUN: env OMPX_APU_MAPS=1 HSA_XNACK=1 LIBOMPTARGET_INFO=30 %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic -check-prefix=INFO_ZERO -check-prefix=CHECK
// RUN: %libomptarget-compilexx-generic
// RUN: env HSA_XNACK=0 LIBOMPTARGET_INFO=30 %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic -check-prefix=INFO_COPY -check-prefix=CHECK
// REQUIRES: amdgpu
// REQUIRES: unified_shared_memory
// clang-format on
#include <cstdio>
int main() {
int n = 1024;
// test various mapping types
int *a = new int[n];
int k = 3;
int b[n];
for (int i = 0; i < n; i++)
b[i] = i;
// clang-format off
// INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory
// INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory
// INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096,
// INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096,
// INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed)
// INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed)
// clang-format on
#pragma omp target teams distribute parallel for map(tofrom : a[ : n]) \
map(to : b[ : n])
for (int i = 0; i < n; i++)
a[i] = i + b[i] + k;
int err = 0;
for (int i = 0; i < n; i++)
if (a[i] != i + b[i] + k)
err++;
// CHECK: PASS
if (err == 0)
printf("PASS\n");
return err;
}
|
