1// clang-format off
2// RUN: %libomptarget-compilexx-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic
3// clang-format on
4
5// REQUIRES: libomptarget-debug
6
7// UNSUPPORTED: nvptx64-nvidia-cuda
8// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
9
10#include <stdio.h>
11#include <stdlib.h>
12
13struct Descriptor {
14 int *datum;
15 long int x;
16 int *more_datum;
17 int xi;
18 int val_datum, val_more_datum;
19 long int arr[1][30];
20 int val_arr;
21};
22
23int main() {
24 Descriptor dat = Descriptor();
25 dat.datum = (int *)malloc(size: sizeof(int) * 10);
26 dat.more_datum = (int *)malloc(size: sizeof(int) * 20);
27 dat.xi = 3;
28 dat.arr[0][0] = 1;
29
30 dat.datum[7] = 7;
31 dat.more_datum[17] = 17;
32 dat.datum[dat.arr[0][0]] = 0;
33
34 /// The struct is mapped with type 0x0 when the pointer fields are mapped.
35 /// The struct is also map explicitely by the user. The second mapping by
36 /// the user must not overwrite the mapping set up for the pointer fields
37 /// when mapping the struct happens after the mapping of the pointers.
38
39 // clang-format off
40 // CHECK: omptarget --> Entry 0: Base=[[DAT_HST_PTR_BASE:0x.*]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x0, Name=unknown
41 // CHECK: omptarget --> Entry 1: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x1000000000001, Name=unknown
42 // CHECK: omptarget --> Entry 2: Base=[[DAT_HST_PTR_BASE]], Begin=[[DATUM_HST_PTR_BASE:0x.*]], Size=40, Type=0x1000000000011, Name=unknown
43 // CHECK: omptarget --> Entry 3: Base=[[MORE_DATUM_HST_PTR_BASE:0x.*]], Begin=[[MORE_DATUM_HST_PTR_BEGIN:0x.*]], Size=80, Type=0x1000000000011, Name=unknown
44 // clang-format on
45
46 /// The struct will be mapped in the same order as the above entries.
47
48 /// First argument is the struct itself and it will be mapped once.
49
50 // clang-format off
51 // CHECK: omptarget --> Looking up mapping(HstPtrBegin=[[DAT_HST_PTR_BASE]], Size=288)...
52 // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 288 with host pointer [[DAT_HST_PTR_BASE]].
53 // CHECK: omptarget --> Creating new map entry with HstPtrBase=[[DAT_HST_PTR_BASE]], HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtAllocBegin=[[DAT_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1, HoldRefCount=0, Name=unknown
54 // CHECK: omptarget --> Moving 288 bytes (hst:[[DAT_HST_PTR_BASE]]) -> (tgt:[[DAT_DEVICE_PTR_BASE]])
55 // clang-format on
56
57 /// Second argument is dat.datum:
58 // clang-format off
59 // CHECK: omptarget --> Looking up mapping(HstPtrBegin=[[DATUM_HST_PTR_BASE]], Size=40)...
60 // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 40 with host pointer [[DATUM_HST_PTR_BASE]].
61 // CHECK: omptarget --> Creating new map entry with HstPtrBase=[[DATUM_HST_PTR_BASE]], HstPtrBegin=[[DATUM_HST_PTR_BASE]], TgtAllocBegin=[[DATUM_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DATUM_DEVICE_PTR_BASE]], Size=40, DynRefCount=1, HoldRefCount=0, Name=unknown
62 // CHECK: omptarget --> Moving 40 bytes (hst:[[DATUM_HST_PTR_BASE]]) -> (tgt:[[DATUM_DEVICE_PTR_BASE]])
63 // clang-format on
64
65 /// Third argument is dat.more_datum:
66 // clang-format off
67 // CHECK: omptarget --> Looking up mapping(HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], Size=80)...
68 // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 80 with host pointer [[MORE_DATUM_HST_PTR_BEGIN]].
69 // CHECK: omptarget --> Creating new map entry with HstPtrBase=[[MORE_DATUM_HST_PTR_BEGIN]], HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], TgtAllocBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN:0x.*]], TgtPtrBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN]], Size=80, DynRefCount=1, HoldRefCount=0, Name=unknown
70 // CHECK: omptarget --> Moving 80 bytes (hst:[[MORE_DATUM_HST_PTR_BEGIN]]) -> (tgt:[[MORE_DATUM_DEVICE_PTR_BEGIN]])
71 // clang-format on
72
73#pragma omp target enter data map(to : dat.datum[ : 10]) \
74 map(to : dat.more_datum[ : 20]) map(to : dat)
75
76 /// Checks induced by having a target region:
77 // clang-format off
78 // CHECK: omptarget --> Entry 0: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x223, Name=unknown
79 // CHECK: omptarget --> Mapping exists (implicit) with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=2 (incremented), HoldRefCount=0, Name=unknown
80 // CHECK: omptarget --> Obtained target argument [[DAT_DEVICE_PTR_BASE]] from host pointer [[DAT_HST_PTR_BASE]]
81 // clang-format on
82
83#pragma omp target
84 {
85 dat.xi = 4;
86 dat.datum[7]++;
87 dat.more_datum[17]++;
88 dat.val_datum = dat.datum[7];
89 dat.val_more_datum = dat.more_datum[17];
90 dat.datum[dat.arr[0][0]] = dat.xi;
91 dat.val_arr = dat.datum[dat.arr[0][0]];
92 }
93
94 /// Post-target region checks:
95 // clang-format off
96 // CHECK: omptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1 (decremented), HoldRefCount=0
97 // clang-format on
98
99#pragma omp target exit data map(from : dat)
100
101 /// Target data end checks:
102 // clang-format off
103 // CHECK: omptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
104 // CHECK: omptarget --> Moving 288 bytes (tgt:[[DAT_DEVICE_PTR_BASE]]) -> (hst:[[DAT_HST_PTR_BASE]])
105 // clang-format on
106
107 // CHECK: dat.xi = 4
108 // CHECK: dat.val_datum = 8
109 // CHECK: dat.val_more_datum = 18
110 // CHECK: dat.datum[dat.arr[0][0]] = 0
111 // CHECK: dat.val_arr = 4
112
113 printf(format: "dat.xi = %d\n", dat.xi);
114 printf(format: "dat.val_datum = %d\n", dat.val_datum);
115 printf(format: "dat.val_more_datum = %d\n", dat.val_more_datum);
116 printf(format: "dat.datum[dat.arr[0][0]] = %d\n", dat.datum[dat.arr[0][0]]);
117 printf(format: "dat.val_arr = %d\n", dat.val_arr);
118
119 return 0;
120}
121

source code of offload/test/offloading/struct_mapping_with_pointers.cpp