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 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