offload/test/offloading/struct_mapping_with_pointers.cpp

   1 // clang-format off
   2 // RUN: %libomptarget-compilexx-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic
   3 // clang-format on
   4
   5 // REQUIRES: libomptarget-debug
   6
   7 // UNSUPPORTED: nvptx64-nvidia-cuda
   8 // UNSUPPORTED: nvptx64-nvidia-cuda-LTO
   9
  10 #include <stdio.h>
  11 #include <stdlib.h>
  12
  13 struct Descriptor {
  14   int *datum;
  15   long int x;
  16   int *more_datum;
  17   int xi;
  18   int val_datum, val_more_datum;
  19   long int arr[1][30];
  20   int val_arr;
  21 };
  22
  23 int main() {
  24   Descriptor dat = Descriptor();
  25   dat.datum = (int *)malloc(sizeof(int) * 10);
  26   dat.more_datum = (int *)malloc(sizeof(int) * 20);
  27   dat.xi = 3;
  28   dat.arr[0][0] = 1;
  29
  30   dat.datum[7] = 7;
  31   dat.more_datum[17] = 17;
  32   dat.datum[dat.arr[0][0]] = 0;
  33
  34   /// The struct is mapped with type 0x0 when the pointer fields are mapped.
  35   /// The struct is also map explicitely by the user. The second mapping by
  36   /// the user must not overwrite the mapping set up for the pointer fields
  37   /// when mapping the struct happens after the mapping of the pointers.
  38
  39   // clang-format off
  40   // CHECK: omptarget --> Entry  0: Base=[[DAT_HST_PTR_BASE:0x.*]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x0, Name=unknown
  41   // CHECK: omptarget --> Entry  1: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x1000000000001, Name=unknown
  42   // CHECK: omptarget --> Entry  2: Base=[[DAT_HST_PTR_BASE]], Begin=[[DATUM_HST_PTR_BASE:0x.*]], Size=40, Type=0x1000000000011, Name=unknown
  43   // CHECK: omptarget --> Entry  3: Base=[[MORE_DATUM_HST_PTR_BASE:0x.*]], Begin=[[MORE_DATUM_HST_PTR_BEGIN:0x.*]], Size=80, Type=0x1000000000011, Name=unknown
  44   // clang-format on
  45
  46   /// The struct will be mapped in the same order as the above entries.
  47
  48   /// First argument is the struct itself and it will be mapped once.
  49
  50   // clang-format off
  51   // CHECK: omptarget --> Looking up mapping(HstPtrBegin=[[DAT_HST_PTR_BASE]], Size=288)...
  52   // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 288 with host pointer [[DAT_HST_PTR_BASE]].
  53   // CHECK: omptarget --> Creating new map entry with HstPtrBase=[[DAT_HST_PTR_BASE]], HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtAllocBegin=[[DAT_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1, HoldRefCount=0, Name=unknown
  54   // CHECK: omptarget --> Moving 288 bytes (hst:[[DAT_HST_PTR_BASE]]) -> (tgt:[[DAT_DEVICE_PTR_BASE]])
  55   // clang-format on
  56
  57   /// Second argument is dat.datum:
  58   // clang-format off
  59   // CHECK: omptarget --> Looking up mapping(HstPtrBegin=[[DATUM_HST_PTR_BASE]], Size=40)...
  60   // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 40 with host pointer [[DATUM_HST_PTR_BASE]].
  61   // CHECK: omptarget --> Creating new map entry with HstPtrBase=[[DATUM_HST_PTR_BASE]], HstPtrBegin=[[DATUM_HST_PTR_BASE]], TgtAllocBegin=[[DATUM_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DATUM_DEVICE_PTR_BASE]], Size=40, DynRefCount=1, HoldRefCount=0, Name=unknown
  62   // CHECK: omptarget --> Moving 40 bytes (hst:[[DATUM_HST_PTR_BASE]]) -> (tgt:[[DATUM_DEVICE_PTR_BASE]])
  63   // clang-format on
  64
  65   /// Third argument is dat.more_datum:
  66   // clang-format off
  67   // CHECK: omptarget --> Looking up mapping(HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], Size=80)...
  68   // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 80 with host pointer [[MORE_DATUM_HST_PTR_BEGIN]].
  69   // CHECK: omptarget --> Creating new map entry with HstPtrBase=[[MORE_DATUM_HST_PTR_BEGIN]], HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], TgtAllocBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN:0x.*]], TgtPtrBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN]], Size=80, DynRefCount=1, HoldRefCount=0, Name=unknown
  70   // CHECK: omptarget --> Moving 80 bytes (hst:[[MORE_DATUM_HST_PTR_BEGIN]]) -> (tgt:[[MORE_DATUM_DEVICE_PTR_BEGIN]])
  71   // clang-format on
  72
  73 #pragma omp target enter data map(to : dat.datum[ : 10])                       \
  74     map(to : dat.more_datum[ : 20]) map(to : dat)
  75
  76   /// Checks induced by having a target region:
  77   // clang-format off
  78   // CHECK: omptarget --> Entry  0: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x223, Name=unknown
  79   // CHECK: omptarget --> Mapping exists (implicit) with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=2 (incremented), HoldRefCount=0, Name=unknown
  80   // CHECK: omptarget --> Obtained target argument [[DAT_DEVICE_PTR_BASE]] from host pointer [[DAT_HST_PTR_BASE]]
  81   // clang-format on
  82
  83 #pragma omp target
  84   {
  85     dat.xi = 4;
  86     dat.datum[7]++;
  87     dat.more_datum[17]++;
  88     dat.val_datum = dat.datum[7];
  89     dat.val_more_datum = dat.more_datum[17];
  90     dat.datum[dat.arr[0][0]] = dat.xi;
  91     dat.val_arr = dat.datum[dat.arr[0][0]];
  92   }
  93
  94   /// Post-target region checks:
  95   // clang-format off
  96   // CHECK: omptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1 (decremented), HoldRefCount=0
  97   // clang-format on
  98
  99 #pragma omp target exit data map(from : dat)
 100
 101   /// Target data end checks:
 102   // clang-format off
 103   // CHECK: omptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
 104   // CHECK: omptarget --> Moving 288 bytes (tgt:[[DAT_DEVICE_PTR_BASE]]) -> (hst:[[DAT_HST_PTR_BASE]])
 105   // clang-format on
 106
 107   // CHECK: dat.xi = 4
 108   // CHECK: dat.val_datum = 8
 109   // CHECK: dat.val_more_datum = 18
 110   // CHECK: dat.datum[dat.arr[0][0]] = 0
 111   // CHECK: dat.val_arr = 4
 112
 113   printf("dat.xi = %d\n", dat.xi);
 114   printf("dat.val_datum = %d\n", dat.val_datum);
 115   printf("dat.val_more_datum = %d\n", dat.val_more_datum);
 116   printf("dat.datum[dat.arr[0][0]] = %d\n", dat.datum[dat.arr[0][0]]);
 117   printf("dat.val_arr = %d\n", dat.val_arr);
 118
 119   return 0;
 120 }