86 lines · c
1// RUN: %libomptarget-compilexx-run-and-check-generic2 3// REQUIRES: libc4 5#include <omp.h>6#include <stdio.h>7 8int x[10];9int *p;10 11void f1() {12 p = &x[0];13 p[0] = 111;14 p[1] = 222;15 p[2] = 333;16 p[3] = 444;17 18#pragma omp target enter data map(to : p)19#pragma omp target enter data map(to : p[0 : 5])20 21 int **p_mappedptr = (int **)omp_get_mapped_ptr(&p, omp_get_default_device());22 int *x0_mappedptr =23 (int *)omp_get_mapped_ptr(&x[0], omp_get_default_device());24 int *x0_hostaddr = &x[0];25 26 printf("p_mappedptr %s null\n", p_mappedptr == (int **)NULL ? "==" : "!=");27 printf("x0_mappedptr %s null\n", x0_mappedptr == (int *)NULL ? "==" : "!=");28 29// CHECK: p_mappedptr != null30// CHECK: x0_mappedptr != null31 32// p is predetermined firstprivate, so its address will be different from33// the mapped address for this construct. So, any changes to p within the34// region will not be visible after the construct.35#pragma omp target map(*p) map(to : p_mappedptr, x0_mappedptr, x0_hostaddr)36 {37 printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[0],38 x0_hostaddr == &p[0]);39 // CHECK: 111 0 1 040 p++;41 }42 43// For the remaining constructs, p is not firstprivate, so its address will44// be the same as the mapped address, and changes to p will be visible to any45// subsequent regions.46#pragma omp target map(to : *p, p) \47 map(to : p_mappedptr, x0_mappedptr, x0_hostaddr)48 {49 printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[0],50 x0_hostaddr == &p[0]);51 // EXPECTED: 111 1 1 052 // CHECK: 111 0 1 053 p++;54 }55 56#pragma omp target map(to : p, *p) \57 map(to : p_mappedptr, x0_mappedptr, x0_hostaddr)58 {59 printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-1],60 x0_hostaddr == &p[-1]);61 // EXPECTED: 222 1 1 062 // CHECK: {{[0-9]+}} 0 0 063 p++;64 }65 66#pragma omp target map(present, alloc : p) \67 map(to : p_mappedptr, x0_mappedptr, x0_hostaddr)68 {69 printf("%d %d %d %d\n", p[0], p_mappedptr == &p, x0_mappedptr == &p[-2],70 x0_hostaddr == &p[-2]);71 // EXPECTED: 333 1 1 072 // CHECK: 111 1 0 073 }74 75 // The following map(from:p) should not bring back p, because p is an76 // attached pointer. So, it should still point to the same original77 // location, &x[0], on host.78#pragma omp target exit data map(always, from : p)79 printf("%d %d\n", p[0], p == &x[0]);80 // CHECK: 111 181 82#pragma omp target exit data map(delete : p[0 : 5], p)83}84 85int main() { f1(); }86