121 lines · cpp
1// clang-format off2// RUN: %libomptarget-compilexx-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic3// clang-format on4 5// REQUIRES: libomptarget-debug6 7// UNSUPPORTED: nvptx64-nvidia-cuda8// UNSUPPORTED: nvptx64-nvidia-cuda-LTO9 10#include <stdio.h>11#include <stdlib.h>12 13struct 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 23int 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 explicitly by the user. The second mapping by36 /// the user must not overwrite the mapping set up for the pointer fields37 /// when mapping the struct happens after the mapping of the pointers.38 39 // clang-format off40 // CHECK: omptarget --> Entry 0: Base=[[DAT_HST_PTR_BASE:0x.*]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x0, Name=unknown41 // CHECK: omptarget --> Entry 1: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x1000000000001, Name=unknown42 // CHECK: omptarget --> Entry 2: Base=[[DAT_HST_PTR_BASE]], Begin=[[DATUM_HST_PTR_BASE:0x.*]], Size=40, Type=0x1000000000011, Name=unknown43 // CHECK: omptarget --> Entry 3: Base=[[MORE_DATUM_HST_PTR_BASE:0x.*]], Begin=[[MORE_DATUM_HST_PTR_BEGIN:0x.*]], Size=80, Type=0x1000000000011, Name=unknown44 // clang-format on45 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 off51 // 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=unknown54 // CHECK: omptarget --> Moving 288 bytes (hst:[[DAT_HST_PTR_BASE]]) -> (tgt:[[DAT_DEVICE_PTR_BASE]])55 // clang-format on56 57 /// Second argument is dat.datum:58 // clang-format off59 // 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=unknown62 // CHECK: omptarget --> Moving 40 bytes (hst:[[DATUM_HST_PTR_BASE]]) -> (tgt:[[DATUM_DEVICE_PTR_BASE]])63 // clang-format on64 65 /// Third argument is dat.more_datum:66 // clang-format off67 // 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=unknown70 // CHECK: omptarget --> Moving 80 bytes (hst:[[MORE_DATUM_HST_PTR_BEGIN]]) -> (tgt:[[MORE_DATUM_DEVICE_PTR_BEGIN]])71 // clang-format on72 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 off78 // CHECK: omptarget --> Entry 0: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x223, Name=unknown79 // CHECK: omptarget --> Mapping exists (implicit) with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=2 (incremented), HoldRefCount=0, Name=unknown80 // CHECK: omptarget --> Obtained target argument [[DAT_DEVICE_PTR_BASE]] from host pointer [[DAT_HST_PTR_BASE]]81 // clang-format on82 83#pragma omp target84 {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 off96 // CHECK: omptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1 (decremented), HoldRefCount=097 // clang-format on98 99#pragma omp target exit data map(from : dat)100 101 /// Target data end checks:102 // clang-format off103 // CHECK: omptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0104 // CHECK: omptarget --> Moving 288 bytes (tgt:[[DAT_DEVICE_PTR_BASE]]) -> (hst:[[DAT_HST_PTR_BASE]])105 // clang-format on106 107 // CHECK: dat.xi = 4108 // CHECK: dat.val_datum = 8109 // CHECK: dat.val_more_datum = 18110 // CHECK: dat.datum[dat.arr[0][0]] = 0111 // CHECK: dat.val_arr = 4112 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}121