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
20int 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, format: "&s = %p\n", &s);
25 fprintf(stderr, format: "&s.x = %p, size = %ld\n", &s.x, sizeof s.x);
26 fprintf(stderr, format: "&s.y = %p\n", &s.y);
27 fprintf(stderr, format: "&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, format: "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, format: "s.x = %d\n", s.x);
36
37 // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
38 // CHECK: omptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier).
39 // CHECK: omptarget 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}
44

source code of offload/test/mapping/padding_not_mapped.c