218 lines · plain
1// RUN: %libomptarget-compilexx-run-and-check-generic2 3// XFAIL: *4 5#include <cstdlib>6#include <cstdio>7#include <cassert>8#include <cstring>9 10#include <omp.h>11 12struct R {13 int d;14 int e;15 int f;16};17 18struct S {19 int a;20 int b;21 struct {22 int c;23 R r;24 R *rp;25 } sub;26 int g;27};28 29struct T {30 int a;31 int *ptr;32 int b;33};34 35int main() {36 R r;37 R *rp = new R;38 S s;39 S *sp = new S;40 T t;41 T *tp = new T;42 43 memset(&r, 0, sizeof(R));44 memset(rp, 0, sizeof(R));45 memset(&s, 0, sizeof(S));46 memset(sp, 0, sizeof(S));47 memset(&t, 0, sizeof(T));48 memset(tp, 0, sizeof(T));49 50 s.sub.rp = new R;51 sp->sub.rp = new R;52 53 memset(s.sub.rp, 0, sizeof(R));54 memset(sp->sub.rp, 0, sizeof(R));55 56 t.ptr = new int[10];57 tp->ptr = new int[10];58 59 memset(t.ptr, 0, sizeof(int)*10);60 memset(tp->ptr, 0, sizeof(int)*10);61 62#pragma omp target map(tofrom: r) map(tofrom: r.e)63{64 r.d++;65 r.e += 2;66 r.f += 3;67}68 printf ("%d\n", r.d); // CHECK: 169 printf ("%d\n", r.e); // CHECK-NEXT: 270 printf ("%d\n", r.f); // CHECK-NEXT: 371 72#pragma omp target map(tofrom: rp[:1]) map(tofrom: rp->e)73{74 rp->d++;75 rp->e += 2;76 rp->f += 3;77}78 79 printf ("%d\n", rp->d); // CHECK-NEXT: 180 printf ("%d\n", rp->e); // CHECK-NEXT: 281 printf ("%d\n", rp->f); // CHECK-NEXT: 382 83 int v;84 int *orig_addr_v = &v;85 bool separate_memory_space;86 87#pragma omp target data map(v)88 {89 void *mapped_ptr_v =90 omp_get_mapped_ptr(orig_addr_v, omp_get_default_device());91 separate_memory_space = mapped_ptr_v != (void*) orig_addr_v;92 }93 94 const char *mapping_flavour = separate_memory_space ? "separate" : "unified";95 96#pragma omp target map(to: s) map(tofrom: s.sub.r.e)97{98 s.b++;99 s.sub.r.d+=2;100 s.sub.r.e+=3;101 s.sub.r.f+=4;102}103 104 printf ("%d/%s\n", s.b, mapping_flavour);105 printf ("%d/%s\n", s.sub.r.d, mapping_flavour);106 printf ("%d/%s\n", s.sub.r.e, mapping_flavour);107 printf ("%d/%s\n", s.sub.r.f, mapping_flavour);108 109 // CHECK: {{0/separate|1/unified}}110 // CHECK-NEXT: {{0/separate|2/unified}}111 // CHECK-NEXT: 3112 // CHECK-NEXT: {{0/separate|4/unified}}113 114#pragma omp target map(to: s, s.b) map(to: s.sub.rp[:1]) map(tofrom: s.sub.rp->e)115{116 s.b++;117 s.sub.rp->d+=2;118 s.sub.rp->e+=3;119 s.sub.rp->f+=4;120}121 122 printf ("%d/%s\n", s.b, mapping_flavour);123 printf ("%d/%s\n", s.sub.rp->d, mapping_flavour);124 printf ("%d/%s\n", s.sub.rp->e, mapping_flavour);125 printf ("%d/%s\n", s.sub.rp->f, mapping_flavour);126 127 // CHECK-NEXT: {{0/separate|2/unified}}128 // CHECK-NEXT: {{0/separate|2/unified}}129 // CHECK-NEXT: 3130 // CHECK-NEXT: {{0/separate|4/unified}}131 132#pragma omp target map(to: sp[:1]) map(tofrom: sp->sub.r.e)133{134 sp->b++;135 sp->sub.r.d+=2;136 sp->sub.r.e+=3;137 sp->sub.r.f+=4;138}139 140 printf ("%d/%s\n", sp->b, mapping_flavour);141 printf ("%d/%s\n", sp->sub.r.d, mapping_flavour);142 printf ("%d/%s\n", sp->sub.r.e, mapping_flavour);143 printf ("%d/%s\n", sp->sub.r.f, mapping_flavour);144 145 // CHECK-NEXT: {{0/separate|1/unified}}146 // CHECK-NEXT: {{0/separate|2/unified}}147 // CHECK-NEXT: 3148 // CHECK-NEXT: {{0/separate|4/unified}}149 150#pragma omp target map(to: sp[:1]) map(to: sp->sub.rp[:1]) map(tofrom: sp->sub.rp->e)151{152 sp->b++;153 sp->sub.rp->d+=2;154 sp->sub.rp->e+=3;155 sp->sub.rp->f+=4;156}157 158 printf ("%d/%s\n", sp->b, mapping_flavour);159 printf ("%d/%s\n", sp->sub.rp->d, mapping_flavour);160 printf ("%d/%s\n", sp->sub.rp->e, mapping_flavour);161 printf ("%d/%s\n", sp->sub.rp->f, mapping_flavour);162 163 // CHECK-NEXT: {{0/separate|2/unified}}164 // CHECK-NEXT: {{0/separate|2/unified}}165 // CHECK-NEXT: 3166 // CHECK-NEXT: {{0/separate|4/unified}}167 168#pragma omp target map(tofrom: t) map(tofrom: t.ptr[2:1])169{170 t.a++;171 t.ptr[2]+=2;172 t.b+=3;173}174 175 printf ("%d\n", t.a); // CHECK-NEXT: 1176 printf ("%d\n", t.ptr[2]); // CHECK-NEXT: 2177 printf ("%d\n", t.b); // CHECK-NEXT: 3178 179#pragma omp target map(tofrom: t) map(tofrom: t.a)180{181 t.b++;182}183 184 printf ("%d\n", t.b); // CHECK-NEXT: 4185 186#pragma omp target map(tofrom: t) map(tofrom: t.ptr[2:1], t.a)187{188 t.a++;189 t.ptr[2]+=2;190 t.b+=3;191}192 193 printf ("%d\n", t.a); // CHECK-NEXT: 2194 printf ("%d\n", t.ptr[2]); // CHECK-NEXT: 4195 printf ("%d\n", t.b); // CHECK-NEXT: 7196 197#pragma omp target map(tofrom: t) map(tofrom: t.ptr[2:1], t.a)198{199 /* Empty */200}201 202 printf ("%d\n", t.a); // CHECK-NEXT: 2203 printf ("%d\n", t.ptr[2]); // CHECK-NEXT: 4204 printf ("%d\n", t.b); // CHECK-NEXT: 7205 206 delete s.sub.rp;207 delete sp->sub.rp;208 209 delete[] t.ptr;210 delete[] tp->ptr;211 212 delete rp;213 delete sp;214 delete tp;215 216 return 0;217}218