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