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.
9struct __tgt_target_non_contig {
10 int64_t offset;
11 int64_t width;
12 int64_t stride;
13};
14
15enum tgt_map_type { OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000 };
16
17// OpenMP RTL interfaces
18#ifdef __cplusplus
19extern "C" {
20#endif
21void __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
28int 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 {.offset: 0, .width: 2, .stride: 480}, {.offset: 1, .width: 3, .stride: 120}, {.offset: 1, .width: 2, .stride: 24}, {.offset: 0, .width: 6, .stride: 4}, {.offset: 0, .width: 1, .stride: 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, args_base: &base, args: &begin,
52 arg_sizes: sizes, arg_types: 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 {.offset: 0, .width: 2, .stride: 320}, {.offset: 2, .width: 2, .stride: 40}, {.offset: 0, .width: 2, .stride: 16}, {.offset: 0, .width: 1, .stride: 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, args_base: &base_2, args: &begin_2,
83 arg_sizes: sizes_2, arg_types: 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

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