89 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 11int x[10];12 13void f1() {14 int *p;15 p = &x[0];16 p[0] = 111;17 p[1] = 222;18 p[2] = 333;19 p[3] = 444;20 21#pragma omp target enter data map(to : p)22#pragma omp target enter data map(to : p[0 : 5])23 24 int **p_mappedptr = (int **)omp_get_mapped_ptr(&p, omp_get_default_device());25 int *x0_mappedptr =26 (int *)omp_get_mapped_ptr(&x[0], omp_get_default_device());27 int *x0_hostaddr = &x[0];28 29 printf("p_mappedptr %s null\n", p_mappedptr == (int **)NULL ? "==" : "!=");30 printf("x0_mappedptr %s null\n", x0_mappedptr == (int *)NULL ? "==" : "!=");31 32// CHECK: p_mappedptr != null33// CHECK: x0_mappedptr != null34 35// p is predetermined firstprivate, so its address will be different from36// the mapped address for this construct. So, any changes to p within the37// region will not be visible after the construct.38#pragma omp target map(*p) map(to : p_mappedptr, x0_mappedptr, x0_hostaddr)39 {40 printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[0],41 x0_hostaddr == &p[0]);42 // CHECK: 111 0 1 043 p++;44 }45 46// For the remaining constructs, p is not firstprivate, so its address will47// be the same as the mapped address, and changes to p will be visible to any48// subsequent regions.49#pragma omp target map(to : *p, p) \50 map(to : p_mappedptr, x0_mappedptr, x0_hostaddr)51 {52 printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[0],53 x0_hostaddr == &p[0]);54 // EXPECTED: 111 1 1 055 // CHECK: 111 0 1 056 p++;57 }58 59#pragma omp target map(to : p, *p) \60 map(to : p_mappedptr, x0_mappedptr, x0_hostaddr)61 {62 printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-1],63 x0_hostaddr == &p[-1]);64 // EXPECTED: 222 1 1 065 // CHECK: {{[0-9]+}} 0 0 066 p++;67 }68 69#pragma omp target map(present, alloc : p) \70 map(to : p_mappedptr, x0_mappedptr, x0_hostaddr)71 {72 printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-2],73 x0_hostaddr == &p[-2]);74 // EXPECTED: 333 1 1 075 // CHECK: 111 1 0 076 }77 78 // The following map(from:p) should not bring back p, because p is an79 // attached pointer. So, it should still point to the same original80 // location, &x[0], on host.81#pragma omp target exit data map(always, from : p)82 printf("%d %d\n", p[0], p == &x[0]);83 // CHECK: 111 184 85#pragma omp target exit data map(delete : p[0 : 5], p)86}87 88int main() { f1(); }89