offload/test/mapping/present/target_data_array_extension.c

   1 // --------------------------------------------------
   2 // Check extends before
   3 // --------------------------------------------------
   4
   5 // RUN: %libomptarget-compile-generic \
   6 // RUN:   -DEXTENDS=BEFORE
   7 // RUN: %libomptarget-run-fail-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-fail-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_SIZE (SMALL_END - SMALL_BEG)
  43 #define LARGE_SIZE (LARGE_END - LARGE_BEG)
  44
  45 #define SMALL                                                                  \
  46   SMALL_BEG:                                                                   \
  47   SMALL_SIZE
  48 #define LARGE                                                                  \
  49   LARGE_BEG:                                                                   \
  50   LARGE_SIZE
  51
  52 int main() {
  53   int arr[SIZE];
  54
  55   // CHECK: addr=0x[[#%x,SMALL_ADDR:]], size=[[#%u,SMALL_BYTES:]]
  56   fprintf(stderr, "addr=%p, size=%ld\n", &arr[SMALL_BEG],
  57           SMALL_SIZE * sizeof arr[0]);
  58
  59   // CHECK: addr=0x[[#%x,LARGE_ADDR:]], size=[[#%u,LARGE_BYTES:]]
  60   fprintf(stderr, "addr=%p, size=%ld\n", &arr[LARGE_BEG],
  61           LARGE_SIZE * sizeof arr[0]);
  62
  63   // CHECK-NOT: omptarget
  64 #pragma omp target data map(alloc : arr[LARGE])
  65   {
  66 #pragma omp target data map(present, tofrom : arr[SMALL])
  67     ;
  68   }
  69
  70   // CHECK: arr is present
  71   fprintf(stderr, "arr is present\n");
  72
  73   // CHECK: omptarget message: explicit extension not allowed: host address specified is 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0*}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes)
  74   // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes)
  75   // CHECK: omptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier).
  76   // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
  77 #pragma omp target data map(alloc : arr[SMALL])
  78   {
  79 #pragma omp target data map(present, tofrom : arr[LARGE])
  80     ;
  81   }
  82
  83   // CHECK-NOT: arr is present
  84   fprintf(stderr, "arr is present\n");
  85
  86   return 0;
  87 }