94 lines · cpp
1// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty -check-prefix=DEBUG2// REQUIRES: libomptarget-debug3 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 interfaces18#ifdef __cplusplus19extern "C" {20#endif21void __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 __cplusplus25}26#endif27 28int main() {29 // case 130 // int arr[3][4][5][6];31 // #pragma omp target update to(arr[0:2][1:3][1:2][:])32 // set up descriptor33 __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 after43 // optimization. Note that each element represent the innermost dimension44 // (unit size = 24) since the stride * count of last dimension is equal to the45 // stride of second last dimension.46 //47 // OOOOO OOOOO OOOOO48 // OXXOO OXXOO OOOOO49 // OXXOO OXXOO OOOOO50 // OXXOO OXXOO OOOOO51 __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base, &begin,52 sizes, types);53 // DEBUG: offset 14454 // DEBUG: offset 26455 // DEBUG: offset 38456 // DEBUG: offset 62457 // DEBUG: offset 74458 // DEBUG: offset 86459 60 // case 261 // double darr[3][4][5];62 // #pragma omp target update to(darr[0:2:2][2:2][:2:2])63 // set up descriptor64 __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 after74 // optimization. Note that each element represent the innermost dimension75 // (unit size = 24) since the stride * count of last dimension is equal to the76 // stride of second last dimension.77 //78 // OOOOO OOOOO OOOOO79 // OOOOO OOOOO OOOOO80 // XOXOO OOOOO XOXOO81 // XOXOO OOOOO XOXOO82 __tgt_target_data_update(/*device_id*/ -1, /*arg_num*/ 1, &base_2, &begin_2,83 sizes_2, types_2);84 // DEBUG: offset 8085 // DEBUG: offset 9686 // DEBUG: offset 12087 // DEBUG: offset 13688 // DEBUG: offset 40089 // DEBUG: offset 41690 // DEBUG: offset 44091 // DEBUG: offset 45692 return 0;93}94