offload/test/offloading/non_contiguous_update.cpp

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