1 // RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty -check-prefix=DEBUG
2 // REQUIRES: libomptarget-debug
3
4 #include <cstdio>
5 #include <cstdlib>
6 #include <cassert>
7
8 // Data structure definitions copied from OpenMP RTL.
9 struct __tgt_target_non_contig {
10 int64_t offset;
11 int64_t width;
12 int64_t stride;
13 };
14
15 enum tgt_map_type {
16 OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000
17 };
18
19 // OpenMP RTL interfaces
20 #ifdef __cplusplus
21 extern "C" {
22 #endif
23 void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
24 void **args_base, void **args, int64_t *arg_sizes,
25 int64_t *arg_types);
26 #ifdef __cplusplus
27 }
28 #endif
29
main()30 int main() {
31 // case 1
32 // int arr[3][4][5][6];
33 // #pragma omp target update to(arr[0:2][1:3][1:2][:])
34 // set up descriptor
35 __tgt_target_non_contig non_contig[5] = {
36 {0, 2, 480}, {1, 3, 120}, {1, 2, 24}, {0, 6, 4}, {0, 1, 4}};
37 int64_t size = 4, type = OMP_TGT_MAPTYPE_NON_CONTIG;
38
39 void *base;
40 void *begin = &non_contig;
41 int64_t *sizes = &size;
42 int64_t *types = &type;
43
44 // The below diagram is the visualization of the non-contiguous transfer after
45 // optimization. Note that each element represent the innermost dimension
46 // (unit size = 24) since the stride * count of last dimension is equal to the
47 // stride of second last dimension.
48 //
49 // OOOOO OOOOO OOOOO
50 // OXXOO OXXOO OOOOO
51 // OXXOO OXXOO OOOOO
52 // OXXOO OXXOO OOOOO
53 __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,
54 sizes, types);
55 // DEBUG: offset 144
56 // DEBUG: offset 264
57 // DEBUG: offset 384
58 // DEBUG: offset 624
59 // DEBUG: offset 744
60 // DEBUG: offset 864
61
62 // case 2
63 // double darr[3][4][5];
64 // #pragma omp target update to(darr[0:2:2][2:2][:2:2])
65 // set up descriptor
66 __tgt_target_non_contig non_contig_2[4] = {
67 {0, 2, 320}, {2, 2, 40}, {0, 2, 16}, {0, 1, 8}};
68 int64_t size_2 = 4, type_2 = OMP_TGT_MAPTYPE_NON_CONTIG;
69
70 void *base_2;
71 void *begin_2 = &non_contig_2;
72 int64_t *sizes_2 = &size_2;
73 int64_t *types_2 = &type_2;
74
75 // The below diagram is the visualization of the non-contiguous transfer after
76 // optimization. Note that each element represent the innermost dimension
77 // (unit size = 24) since the stride * count of last dimension is equal to the
78 // stride of second last dimension.
79 //
80 // OOOOO OOOOO OOOOO
81 // OOOOO OOOOO OOOOO
82 // XOXOO OOOOO XOXOO
83 // XOXOO OOOOO XOXOO
84 __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base_2, &begin_2,
85 sizes_2, types_2);
86 // DEBUG: offset 80
87 // DEBUG: offset 96
88 // DEBUG: offset 120
89 // DEBUG: offset 136
90 // DEBUG: offset 400
91 // DEBUG: offset 416
92 // DEBUG: offset 440
93 // DEBUG: offset 456
94 return 0;
95 }
96
97