93 lines · c
1// RUN: %libomptarget-compile-run-and-check-generic2 3// REQUIRES: libc4//5// FIXME: https://github.com/llvm/llvm-project/issues/1612656// XFAIL: gpu7 8#include <omp.h>9#include <stdio.h>10 11typedef struct {12 short x;13 int *p;14 long y;15} S;16 17void f1() {18 S s[10], *ps;19 ps = &s[0];20 s[0].x = 111;21 s[1].x = 222;22 s[2].x = 333;23 s[3].x = 444;24 25#pragma omp target enter data map(to : s)26#pragma omp target enter data map(to : ps, ps->x)27 28 S **ps_mappedptr = (S **)omp_get_mapped_ptr(&ps, omp_get_default_device());29 short *s0_mappedptr =30 (short *)omp_get_mapped_ptr(&s[0].x, omp_get_default_device());31 short *s0_hostaddr = &s[0].x;32 33 printf("ps_mappedptr %s null\n", ps_mappedptr == (S **)NULL ? "==" : "!=");34 printf("s0_mappedptr %s null\n", s0_mappedptr == (short *)NULL ? "==" : "!=");35 36// CHECK: ps_mappedptr != null37// CHECK: s0_mappedptr != null38 39// ps is predetermined firstprivate, so its address will be different from40// the mapped address for this construct. So, any changes to p within the41// region will not be visible after the construct.42#pragma omp target map(ps->x) map(to : ps_mappedptr, s0_mappedptr, s0_hostaddr)43 {44 printf("%d %d %d %d\n", ps->x, ps_mappedptr == &ps, s0_mappedptr == &ps->x,45 s0_hostaddr == &ps->x);46 // CHECK: 111 0 1 047 ps++;48 }49 50// For the remaining constructs, ps is not firstprivate, so its address will51// be the same as the mapped address, and changes to ps will be visible to any52// subsequent regions.53#pragma omp target map(to : ps->x, ps) \54 map(to : ps_mappedptr, s0_mappedptr, s0_hostaddr)55 {56 printf("%d %d %d %d\n", ps->x, ps_mappedptr == &ps, s0_mappedptr == &ps->x,57 s0_hostaddr == &ps->x);58 // EXPECTED: 111 1 1 059 // CHECK: 111 0 1 060 ps++;61 }62 63#pragma omp target map(to : ps, ps->x) \64 map(to : ps_mappedptr, s0_mappedptr, s0_hostaddr)65 {66 printf("%d %d %d %d\n", ps->x, ps_mappedptr == &ps,67 s0_mappedptr == &ps[-1].x, s0_hostaddr == &ps[-1].x);68 // EXPECTED: 222 1 1 069 // CHECK: 111 0 0 070 ps++;71 }72 73#pragma omp target map(present, alloc : ps) \74 map(to : ps_mappedptr, s0_mappedptr, s0_hostaddr)75 {76 printf("%d %d %d %d\n", ps->x, ps_mappedptr == &ps,77 s0_mappedptr == &ps[-2].x, s0_hostaddr == &ps[-2].x);78 // EXPECTED: 333 1 1 079 // CHECK: 111 1 0 080 }81 82 // The following map(from:ps) should not bring back ps, because ps is an83 // attached pointer. So, it should still point to the same original84 // location, &s[0], on host.85#pragma omp target exit data map(always, from : ps)86 printf("%d %d\n", ps->x, ps == &s[0]);87 // CHECK: 111 188 89#pragma omp target exit data map(delete : ps, s)90}91 92int main() { f1(); }93