openmp/libomptarget/test/mapping/padding_not_mapped.c

   1 // RUN: %libomptarget-compile-generic -fopenmp-version=51
   2 // RUN: %libomptarget-run-fail-generic 2>&1 \
   3 // RUN: | %fcheck-generic
   4
   5 // The host memory layout for the following program looks like this:
   6 //
   7 //   | 4 bytes | 4 bytes | 8 bytes |
   8 //   |   s.x   |   s.y   |   s.z   |
   9 //   `-----------------------------'
  10 //
  11 // s is always at least 8-byte aligned in host memory due to s.z, so
  12 // libomptarget's device padding for map(s.y,s.z) always maps to host memory
  13 // that includes s.x.  At one time, s.x appeared to be mapped as a result, but
  14 // libomptarget has since been fixed not to consider device padding as mapped to
  15 // host memory.
  16
  17 #include <omp.h>
  18 #include <stdio.h>
  19
  20 int main() {
  21   struct S { int x; int y; double z; } s = {1, 2, 3};
  22
  23   // CHECK: &s.x = 0x[[#%x,HOST_ADDR:]], size = [[#%u,SIZE:]]
  24   fprintf(stderr, "&s = %p\n", &s);
  25   fprintf(stderr, "&s.x = %p, size = %ld\n", &s.x, sizeof s.x);
  26   fprintf(stderr, "&s.y = %p\n", &s.y);
  27   fprintf(stderr, "&s.z = %p\n", &s.z);
  28
  29   // CHECK: s.x is present: 0
  30   // CHECK: s.x = 1{{$}}
  31   #pragma omp target enter data map(alloc: s.y, s.z)
  32   int dev = omp_get_default_device();
  33   fprintf(stderr, "s.x is present: %d\n", omp_target_is_present(&s.x, dev));
  34   #pragma omp target update from(s.x) // should have no effect
  35   fprintf(stderr, "s.x = %d\n", s.x);
  36
  37   // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
  38   // CHECK: Libomptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier).
  39   // CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
  40   #pragma omp target enter data map(present, alloc: s.x)
  41
  42   return 0;
  43 }