offload/test/mapping/target_data_array_extension_at_exit.c

   1 // --------------------------------------------------
   2 // Check extends before
   3 // --------------------------------------------------
   4
   5 // RUN: %libomptarget-compile-generic \
   6 // RUN:   -DEXTENDS=BEFORE
   7 // RUN: %libomptarget-run-generic 2>&1 \
   8 // RUN: | %fcheck-generic
   9
  10 // --------------------------------------------------
  11 // Check extends after
  12 // --------------------------------------------------
  13
  14 // RUN: %libomptarget-compile-generic \
  15 // RUN:   -DEXTENDS=AFTER
  16 // RUN: %libomptarget-run-generic 2>&1 \
  17 // RUN: | %fcheck-generic
  18
  19 // END.
  20
  21 #include <stdio.h>
  22
  23 #define BEFORE 0
  24 #define AFTER 1
  25
  26 #define SIZE 100
  27
  28 #if EXTENDS == BEFORE
  29 #define SMALL_BEG (SIZE - 2)
  30 #define SMALL_END SIZE
  31 #define LARGE_BEG 0
  32 #define LARGE_END SIZE
  33 #elif EXTENDS == AFTER
  34 #define SMALL_BEG 0
  35 #define SMALL_END 2
  36 #define LARGE_BEG 0
  37 #define LARGE_END SIZE
  38 #else
  39 #error EXTENDS undefined
  40 #endif
  41
  42 #define SMALL                                                                  \
  43   SMALL_BEG:                                                                   \
  44   (SMALL_END - SMALL_BEG)
  45 #define LARGE                                                                  \
  46   LARGE_BEG:                                                                   \
  47   (LARGE_END - LARGE_BEG)
  48
  49 void check_not_present() {
  50   int arr[SIZE];
  51
  52   for (int i = 0; i < SIZE; ++i)
  53     arr[i] = 99;
  54
  55   // CHECK-LABEL: checking not present
  56   fprintf(stderr, "checking not present\n");
  57
  58   // arr[LARGE] isn't (fully) present at the end of the target data region, so
  59   // the device-to-host transfer should not be performed, or it might fail.
  60 #pragma omp target data map(tofrom : arr[LARGE])
  61   {
  62 #pragma omp target exit data map(delete : arr[LARGE])
  63 #pragma omp target enter data map(alloc : arr[SMALL])
  64 #pragma omp target map(alloc : arr[SMALL])
  65     for (int i = SMALL_BEG; i < SMALL_END; ++i)
  66       arr[i] = 88;
  67   }
  68
  69   // CHECK-NOT: omptarget
  70   // CHECK-NOT: error
  71   for (int i = 0; i < SIZE; ++i) {
  72     if (arr[i] != 99)
  73       fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
  74   }
  75 }
  76
  77 void check_is_present() {
  78   int arr[SIZE];
  79
  80   for (int i = 0; i < SIZE; ++i)
  81     arr[i] = 99;
  82
  83   // CHECK-LABEL: checking is present
  84   fprintf(stderr, "checking is present\n");
  85
  86   // arr[SMALL] is (fully) present at the end of the target data region, and the
  87   // device-to-host transfer should be performed only for it even though more
  88   // of the array is then present.
  89 #pragma omp target data map(tofrom : arr[SMALL])
  90   {
  91 #pragma omp target exit data map(delete : arr[SMALL])
  92 #pragma omp target enter data map(alloc : arr[LARGE])
  93 #pragma omp target map(alloc : arr[LARGE])
  94     for (int i = LARGE_BEG; i < LARGE_END; ++i)
  95       arr[i] = 88;
  96   }
  97
  98   // CHECK-NOT: omptarget
  99   // CHECK-NOT: error
 100   for (int i = 0; i < SIZE; ++i) {
 101     if (SMALL_BEG <= i && i < SMALL_END) {
 102       if (arr[i] != 88)
 103         fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
 104     } else if (arr[i] != 99) {
 105       fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
 106     }
 107   }
 108 }
 109
 110 int main() {
 111   check_not_present();
 112   check_is_present();
 113   return 0;
 114 }