diff options
Diffstat (limited to 'libgomp')
-rw-r--r-- | libgomp/ChangeLog | 9 | ||||
-rw-r--r-- | libgomp/target.c | 35 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.c++/target-22.C | 99 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.c/target-37.c | 71 |
4 files changed, 201 insertions, 13 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 01440f7..5356d00 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,12 @@ +2019-08-08 Jakub Jelinek <jakub@redhat.com> + + * target.c (gomp_map_vars_internal): For GOMP_MAP_USE_DEVICE_PTR + perform the lookup in the first loop only if !not_found_cnt, otherwise + perform lookups for it in the second loop guarded with + if (not_found_cnt || has_firstprivate). + * testsuite/libgomp.c/target-37.c: New test. + * testsuite/libgomp.c++/target-22.C: New test. + 2019-08-07 Jakub Jelinek <jakub@redhat.com> * testsuite/libgomp.c/target-18.c (struct S): New type. diff --git a/libgomp/target.c b/libgomp/target.c index 9416401..1c9ca68 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -580,20 +580,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, } else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR) { - cur_node.host_start = (uintptr_t) hostaddrs[i]; - cur_node.host_end = cur_node.host_start; - splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); - if (n == NULL) + tgt->list[i].key = NULL; + if (!not_found_cnt) { - gomp_mutex_unlock (&devicep->lock); - gomp_fatal ("use_device_ptr pointer wasn't mapped"); } - cur_node.host_start -= n->host_start; - hostaddrs[i] - = (void *) (n->tgt->tgt_start + n->tgt_offset - + cur_node.host_start); - tgt->list[i].key = NULL; - tgt->list[i].offset = ~(uintptr_t) 0; + else + tgt->list[i].offset = 0; continue; } else if ((kind & typemask) == GOMP_MAP_STRUCT) @@ -791,9 +783,26 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, tgt_size += len; continue; case GOMP_MAP_FIRSTPRIVATE_INT: - case GOMP_MAP_USE_DEVICE_PTR: case GOMP_MAP_ZERO_LEN_ARRAY_SECTION: continue; + case GOMP_MAP_USE_DEVICE_PTR: + if (tgt->list[i].offset == 0) + { + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = cur_node.host_start; + n = gomp_map_lookup (mem_map, &cur_node); + if (n == NULL) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("use_device_ptr pointer wasn't mapped"); + } + cur_node.host_start -= n->host_start; + hostaddrs[i] + = (void *) (n->tgt->tgt_start + n->tgt_offset + + cur_node.host_start); + tgt->list[i].offset = ~(uintptr_t) 0; + } + continue; case GOMP_MAP_STRUCT: first = i + 1; last = i + sizes[i]; diff --git a/libgomp/testsuite/libgomp.c++/target-22.C b/libgomp/testsuite/libgomp.c++/target-22.C new file mode 100644 index 0000000..9d9dea0 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-22.C @@ -0,0 +1,99 @@ +extern "C" void abort (void); +struct S { int e, f; }; + +void +foo (int *&p, int (&s)[5], int &t, S &u, int n) +{ + int a[4] = { 7, 8, 9, 10 }, b[n], c[3] = { 20, 21, 22 }; + int *r = a + 1, *q = p - 1, i, err; + int v = 27; + S w = { 28, 29 }; + for (i = 0; i < n; i++) + b[i] = 9 + i; + #pragma omp target data map(to:a) use_device_ptr(r) map(from:err) + #pragma omp target is_device_ptr(r) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (r[i - 1] != 7 + i) + err = 1; + } + if (err) + abort (); + #pragma omp target data use_device_ptr(p) map(from:err) map(to:q[:4]) + #pragma omp target is_device_ptr(p) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (p[i - 1] != i) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:b) use_device_addr(b) map(from:err) + #pragma omp target is_device_ptr(b) private(i) map(from:err) + { + err = 0; + for (i = 0; i < n; i++) + if (b[i] != 9 + i) + err = 1; + } + if (err) + abort (); + #pragma omp target data use_device_addr(c) map(to:c) map(from:err) + #pragma omp target is_device_ptr(c) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 3; i++) + if (c[i] != 20 + i) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:s[:5]) use_device_addr(s) map(from:err) + #pragma omp target is_device_ptr(s) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 5; i++) + if (s[i] != 17 + i) + err = 1; + } + if (err) + abort (); + #pragma omp target data use_device_addr (v) map(to: v) map(to:u) use_device_addr (u) map(from:err) + { + int *z = &v; + S *x = &u; + #pragma omp target is_device_ptr (z, x) map(from:err) + { + err = 0; + if (*z != 27 || x->e != 25 || x->f != 26) + err = 1; + } + } + if (err) + abort (); + #pragma omp target data map(to: t) use_device_addr (t, w) map (to: w) map(from:err) + { + int *z = &t; + S *x = &w; + #pragma omp target is_device_ptr (z) is_device_ptr (x) map(from:err) + { + err = 0; + if (*z != 24 || x->e != 28 || x->f != 29) + err = 1; + } + } + if (err) + abort (); +} + +int +main () +{ + int a[4] = { 0, 1, 2, 3 }, b[5] = { 17, 18, 19, 20, 21 }; + int *p = a + 1; + int t = 24; + S u = { 25, 26 }; + foo (p, b, t, u, 9); +} diff --git a/libgomp/testsuite/libgomp.c/target-37.c b/libgomp/testsuite/libgomp.c/target-37.c new file mode 100644 index 0000000..b3cc6a2 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-37.c @@ -0,0 +1,71 @@ +extern void abort (void); +struct S { int e, f; }; + +void +foo (int n) +{ + int a[4] = { 0, 1, 2, 3 }, b[n], c = 4; + struct S d = { 5, 6 }; + int *p = a + 1, i, err; + for (i = 0; i < n; i++) + b[i] = 9 + i; + #pragma omp target data use_device_ptr(p) map(from:err) map(to:a) + #pragma omp target is_device_ptr(p) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (p[i - 1] != i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < 4; i++) + a[i] = 23 + i; + #pragma omp target data map(to:a) use_device_addr(a) map(from:err) + #pragma omp target is_device_ptr(a) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (a[i] != 23 + i) + err = 1; + } + if (err) + abort (); + #pragma omp target data use_device_addr(b) map(from:err) map(to:b) + #pragma omp target is_device_ptr(b) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (b[i] != 9 + i) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:c) use_device_addr(c) map(from:err) + { + int *q = &c; + #pragma omp target is_device_ptr(q) map(from:err) + { + err = *q != 4; + } + } + if (err) + abort (); + #pragma omp target data use_device_addr(d) map(to:d) map(from:err) + { + struct S *r = &d; + #pragma omp target is_device_ptr(r) map(from:err) + { + err = r->e != 5 || r->f != 6; + } + } + if (err) + abort (); +} + +int +main () +{ + foo (9); + return 0; +} |