OpenMP, libgomp: Add new runtime routine omp_get_mapped_ptr.
This patch adds the OpenMP runtime routine "omp_get_mapped_ptr" which was introduced in OpenMP 5.1. gcc/ChangeLog: * omp-low.cc (omp_runtime_api_call): Added get_mapped_ptr to omp_runtime_apis array. libgomp/ChangeLog: * libgomp.map: Added omp_get_mapped_ptr. * libgomp.texi: Tagged omp_get_mapped_ptr as supported. * omp.h.in: Added omp_get_mapped_ptr. * omp_lib.f90.in: Added interface for omp_get_mapped_ptr. * omp_lib.h.in: Likewise. * target.c (omp_get_mapped_ptr): Added implementation of omp_get_mapped_ptr. * testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c: New test. * testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c: New test. * testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c: New test. * testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c: New test. * testsuite/libgomp.fortran/get-mapped-ptr-1.f90: New test. * testsuite/libgomp.fortran/get-mapped-ptr-2.f90: New test. * testsuite/libgomp.fortran/get-mapped-ptr-3.f90: New test. * testsuite/libgomp.fortran/get-mapped-ptr-4.f90: New test.
This commit is contained in:
Marcel Vollweiler
parent
404edfce68
commit
941cdc8b6d
|
|
@ -3982,6 +3982,7 @@ omp_runtime_api_call (const_tree fndecl)
|
|||
"alloc",
|
||||
"calloc",
|
||||
"free",
|
||||
"get_mapped_ptr",
|
||||
"realloc",
|
||||
"target_alloc",
|
||||
"target_associate_ptr",
|
||||
|
|
|
|||
|
|
@ -226,6 +226,11 @@ OMP_5.1 {
|
|||
omp_get_teams_thread_limit_;
|
||||
} OMP_5.0.2;
|
||||
|
||||
OMP_5.1.1 {
|
||||
global:
|
||||
omp_get_mapped_ptr;
|
||||
} OMP_5.1;
|
||||
|
||||
GOMP_1.0 {
|
||||
global:
|
||||
GOMP_atomic_end;
|
||||
|
|
|
|||
|
|
@ -314,7 +314,7 @@ The OpenMP 4.5 specification is fully supported.
|
|||
@item @code{omp_target_is_accessible} runtime routine @tab N @tab
|
||||
@item @code{omp_target_memcpy_async} and @code{omp_target_memcpy_rect_async}
|
||||
runtime routines @tab N @tab
|
||||
@item @code{omp_get_mapped_ptr} runtime routine @tab N @tab
|
||||
@item @code{omp_get_mapped_ptr} runtime routine @tab Y @tab
|
||||
@item @code{omp_calloc}, @code{omp_realloc}, @code{omp_aligned_alloc} and
|
||||
@code{omp_aligned_calloc} runtime routines @tab Y @tab
|
||||
@item @code{omp_alloctrait_key_t} enum: @code{omp_atv_serialized} added,
|
||||
|
|
|
|||
|
|
@ -282,6 +282,7 @@ extern int omp_target_memcpy_rect (void *, const void *, __SIZE_TYPE__, int,
|
|||
extern int omp_target_associate_ptr (const void *, const void *, __SIZE_TYPE__,
|
||||
__SIZE_TYPE__, int) __GOMP_NOTHROW;
|
||||
extern int omp_target_disassociate_ptr (const void *, int) __GOMP_NOTHROW;
|
||||
extern void *omp_get_mapped_ptr (const void *, int) __GOMP_NOTHROW;
|
||||
|
||||
extern void omp_set_affinity_format (const char *) __GOMP_NOTHROW;
|
||||
extern __SIZE_TYPE__ omp_get_affinity_format (char *, __SIZE_TYPE__)
|
||||
|
|
|
|||
|
|
@ -835,6 +835,15 @@
|
|||
end function omp_target_disassociate_ptr
|
||||
end interface
|
||||
|
||||
interface
|
||||
function omp_get_mapped_ptr (ptr, device_num) bind(c)
|
||||
use, intrinsic :: iso_c_binding, only : c_ptr, c_int
|
||||
type(c_ptr) :: omp_get_mapped_ptr
|
||||
type(c_ptr), value :: ptr
|
||||
integer(c_int), value :: device_num
|
||||
end function omp_get_mapped_ptr
|
||||
end interface
|
||||
|
||||
#if _OPENMP >= 201811
|
||||
!GCC$ ATTRIBUTES DEPRECATED :: omp_get_nested, omp_set_nested
|
||||
#endif
|
||||
|
|
|
|||
|
|
@ -416,3 +416,12 @@
|
|||
integer(c_int), value :: device_num
|
||||
end function omp_target_disassociate_ptr
|
||||
end interface
|
||||
|
||||
interface
|
||||
function omp_get_mapped_ptr (ptr, device_num) bind(c)
|
||||
use, intrinsic :: iso_c_binding, only : c_ptr, c_int
|
||||
type(c_ptr) :: omp_get_mapped_ptr
|
||||
type(c_ptr), value :: ptr
|
||||
integer(c_int), value :: device_num
|
||||
end function omp_get_mapped_ptr
|
||||
end interface
|
||||
|
|
|
|||
|
|
@ -3665,6 +3665,44 @@ omp_target_disassociate_ptr (const void *ptr, int device_num)
|
|||
return ret;
|
||||
}
|
||||
|
||||
void *
|
||||
omp_get_mapped_ptr (const void *ptr, int device_num)
|
||||
{
|
||||
if (device_num < 0 || device_num > gomp_get_num_devices ())
|
||||
return NULL;
|
||||
|
||||
if (device_num == omp_get_initial_device ())
|
||||
return (void *) ptr;
|
||||
|
||||
struct gomp_device_descr *devicep = resolve_device (device_num);
|
||||
if (devicep == NULL)
|
||||
return NULL;
|
||||
|
||||
if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
|
||||
|| devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|
||||
return (void *) ptr;
|
||||
|
||||
gomp_mutex_lock (&devicep->lock);
|
||||
|
||||
struct splay_tree_s *mem_map = &devicep->mem_map;
|
||||
struct splay_tree_key_s cur_node;
|
||||
void *ret = NULL;
|
||||
|
||||
cur_node.host_start = (uintptr_t) ptr;
|
||||
cur_node.host_end = cur_node.host_start;
|
||||
splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);
|
||||
|
||||
if (n)
|
||||
{
|
||||
uintptr_t offset = cur_node.host_start - n->host_start;
|
||||
ret = (void *) (n->tgt->tgt_start + n->tgt_offset + offset);
|
||||
}
|
||||
|
||||
gomp_mutex_unlock (&devicep->lock);
|
||||
|
||||
return ret;
|
||||
}
|
||||
|
||||
int
|
||||
omp_pause_resource (omp_pause_resource_t kind, int device_num)
|
||||
{
|
||||
|
|
|
|||
|
|
@ -0,0 +1,41 @@
|
|||
#include <omp.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
int
|
||||
main ()
|
||||
{
|
||||
int d = omp_get_default_device ();
|
||||
int id = omp_get_initial_device ();
|
||||
void *p , *q;
|
||||
|
||||
if (d < 0 || d >= omp_get_num_devices ())
|
||||
d = id;
|
||||
|
||||
p = omp_target_alloc (sizeof (int), d);
|
||||
if (p == NULL)
|
||||
return 0;
|
||||
|
||||
if (omp_target_associate_ptr (q, p, sizeof (int), 0, d) != 0)
|
||||
return 0;
|
||||
|
||||
if (omp_get_mapped_ptr (q, -1) != NULL)
|
||||
abort ();
|
||||
|
||||
if (omp_get_mapped_ptr (q, omp_get_num_devices () + 1) != NULL)
|
||||
abort ();
|
||||
|
||||
if (omp_get_mapped_ptr (q, id) != q)
|
||||
abort ();
|
||||
|
||||
if (omp_get_mapped_ptr (q, d) != p)
|
||||
abort ();
|
||||
|
||||
if (omp_target_disassociate_ptr (q, d) != 0)
|
||||
abort ();
|
||||
|
||||
if (omp_get_mapped_ptr (q, d) != NULL)
|
||||
abort ();
|
||||
|
||||
omp_target_free (p, d);
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,106 @@
|
|||
#include <omp.h>
|
||||
#include <stdlib.h>
|
||||
#include <stdint.h>
|
||||
|
||||
int
|
||||
main ()
|
||||
{
|
||||
int d = omp_get_default_device ();
|
||||
int id = omp_get_initial_device ();
|
||||
int a = 42;
|
||||
int b[] = { 24, 42 };
|
||||
int c[] = { 47, 11 };
|
||||
int e[128];
|
||||
int *q = &a;
|
||||
void *p1 = NULL, *p2 = NULL, *p3 = NULL;
|
||||
void *devptrs[128];
|
||||
|
||||
if (d < 0 || d >= omp_get_num_devices ())
|
||||
d = id;
|
||||
|
||||
for (int i = 0; i < 128; i++)
|
||||
e[i] = i;
|
||||
|
||||
#pragma omp target data map(alloc: a, b, c[1], e[32:64]) device(d)
|
||||
{
|
||||
#pragma omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c[1], e[32:64]) device(d)
|
||||
{
|
||||
p1 = &a;
|
||||
p2 = &b;
|
||||
p3 = &c[1];
|
||||
for (int i = 32; i < 96; i++)
|
||||
devptrs[i] = &e[i];
|
||||
}
|
||||
|
||||
if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : p1)
|
||||
|| omp_get_mapped_ptr (q, d) != (d == id ? q : p1)
|
||||
|| omp_get_mapped_ptr (b, d) != (d == id ? b : p2)
|
||||
|| omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : p2)
|
||||
|| omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : p3)
|
||||
|| omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
|
||||
abort ();
|
||||
|
||||
for (int i = 0; i < 32; i++)
|
||||
if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
|
||||
abort ();
|
||||
for (int i = 32; i < 96; i++)
|
||||
if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : devptrs[i]))
|
||||
abort ();
|
||||
for (int i = 96; i < 128; i++)
|
||||
if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
|
||||
abort ();
|
||||
}
|
||||
|
||||
if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : NULL)
|
||||
|| omp_get_mapped_ptr (q, d) != (d == id ? q : NULL)
|
||||
|| omp_get_mapped_ptr (b, d) != (d == id ? b : NULL)
|
||||
|| omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : NULL)
|
||||
|| omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : NULL)
|
||||
|| omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
|
||||
abort ();
|
||||
for (int i = 0; i < 128; i++)
|
||||
if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
|
||||
abort ();
|
||||
|
||||
#pragma omp target enter data map (alloc: a, b, c[1], e[32:64]) device (d)
|
||||
#pragma omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c[1], e[32:64]) device(d)
|
||||
{
|
||||
p1 = &a;
|
||||
p2 = &b;
|
||||
p3 = &c[1];
|
||||
for (int i = 32; i < 96; i++)
|
||||
devptrs[i] = &e[i];
|
||||
}
|
||||
|
||||
if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : p1)
|
||||
|| omp_get_mapped_ptr (q, d) != (d == id ? q : p1)
|
||||
|| omp_get_mapped_ptr (b, d) != (d == id ? b : p2)
|
||||
|| omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : p2)
|
||||
|| omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : p3)
|
||||
|| omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
|
||||
abort ();
|
||||
for (int i = 0; i < 32; i++)
|
||||
if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
|
||||
abort ();
|
||||
for (int i = 32; i < 96; i++)
|
||||
if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : devptrs[i]))
|
||||
abort ();
|
||||
for (int i = 96; i < 128; i++)
|
||||
if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
|
||||
abort ();
|
||||
|
||||
#pragma omp target exit data map (delete: a, b, c[1], e[32:64]) device (d)
|
||||
|
||||
if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : NULL)
|
||||
|| omp_get_mapped_ptr (q, d) != (d == id ? q : NULL)
|
||||
|| omp_get_mapped_ptr (b, d) != (d == id ? b : NULL)
|
||||
|| omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : NULL)
|
||||
|| omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : NULL)
|
||||
|| omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
|
||||
abort ();
|
||||
for (int i = 0; i < 128; i++)
|
||||
if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
|
||||
abort ();
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,51 @@
|
|||
#include <omp.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
int
|
||||
main ()
|
||||
{
|
||||
int d = omp_get_default_device ();
|
||||
int id = omp_get_initial_device ();
|
||||
int a[0];
|
||||
int b[] = { 24, 42 };
|
||||
void *p1 = NULL, *p2 = NULL;
|
||||
|
||||
if (d < 0 || d >= omp_get_num_devices ())
|
||||
d = id;
|
||||
|
||||
void *p = omp_target_alloc (sizeof (int), d);
|
||||
if (p == NULL)
|
||||
return 0;
|
||||
|
||||
if (omp_target_associate_ptr (a, p, sizeof (int), 0, d) != 0)
|
||||
return 0;
|
||||
|
||||
if (omp_get_mapped_ptr (a, d) != (d == id ? a : p))
|
||||
abort ();
|
||||
|
||||
if (omp_target_disassociate_ptr (a, d) != 0)
|
||||
abort ();
|
||||
|
||||
if (omp_get_mapped_ptr (a, d) != (d == id ? a : NULL))
|
||||
abort ();
|
||||
|
||||
#pragma omp target data map(alloc: a, b[1:0]) device(d)
|
||||
{
|
||||
#pragma omp target map(from: p1, p2) map(alloc: a, b[1:0]) device(d)
|
||||
{
|
||||
p1 = &a;
|
||||
p2 = &b[1];
|
||||
}
|
||||
|
||||
/* This is probably expected to be p1/p2 instead of NULL. Zero-length arrays
|
||||
as list items of the map clause are currently not inserted into the mem
|
||||
map ?! However by returning NULL, omp_get_mapped_ptr is consistent with
|
||||
omp_target_is_present. */
|
||||
if (omp_get_mapped_ptr (a, d) != NULL
|
||||
|| omp_get_mapped_ptr (&b[1], d) != NULL)
|
||||
abort ();
|
||||
}
|
||||
|
||||
omp_target_free (p, d);
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,49 @@
|
|||
#include <omp.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
int
|
||||
main ()
|
||||
{
|
||||
int d = omp_get_default_device ();
|
||||
int id = omp_get_initial_device ();
|
||||
struct s_t { int m1; char m2; } s;
|
||||
void *p1 = NULL, *p2 = NULL;
|
||||
|
||||
if (d < 0 || d >= omp_get_num_devices ())
|
||||
d = id;
|
||||
|
||||
#pragma omp target data map(alloc: s, s.m2) device(d)
|
||||
{
|
||||
#pragma omp target map(from: p1, p2) map(alloc: s, s.m2) device(d)
|
||||
{
|
||||
p1 = &s;
|
||||
p2 = &s.m2;
|
||||
}
|
||||
if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : p1)
|
||||
|| omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : p2))
|
||||
abort ();
|
||||
}
|
||||
|
||||
if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : NULL)
|
||||
|| omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : NULL))
|
||||
abort ();
|
||||
|
||||
#pragma omp target enter data map(alloc: s, s.m2) device (d)
|
||||
#pragma omp target map(from: p1, p2) map(alloc: s, s.m2) device(d)
|
||||
{
|
||||
p1 = &s;
|
||||
p2 = &s.m2;
|
||||
}
|
||||
|
||||
if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : p1)
|
||||
|| omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : p2))
|
||||
abort ();
|
||||
|
||||
#pragma omp target exit data map (delete: s, s.m2) device (d)
|
||||
|
||||
if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : NULL)
|
||||
|| omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : NULL))
|
||||
abort ();
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,43 @@
|
|||
program main
|
||||
use omp_lib
|
||||
use iso_c_binding
|
||||
implicit none (external, type)
|
||||
integer :: d, id
|
||||
type(c_ptr) :: p
|
||||
integer, target :: q
|
||||
|
||||
d = omp_get_default_device ()
|
||||
id = omp_get_initial_device ()
|
||||
|
||||
if (d < 0 .or. d >= omp_get_num_devices ()) &
|
||||
d = id
|
||||
|
||||
p = omp_target_alloc (c_sizeof (q), d)
|
||||
if (.not. c_associated (p)) &
|
||||
stop 0 ! okay
|
||||
|
||||
if (omp_target_associate_ptr (c_loc (q), p, c_sizeof (q), &
|
||||
0_c_size_t, d) == 0) then
|
||||
|
||||
if(c_associated (omp_get_mapped_ptr (c_loc (q), -1))) &
|
||||
stop 1
|
||||
|
||||
if(c_associated (omp_get_mapped_ptr (c_loc (q), &
|
||||
omp_get_num_devices () + 1))) &
|
||||
stop 2
|
||||
|
||||
if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), id), c_loc (q))) &
|
||||
stop 3
|
||||
|
||||
if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), d), p)) &
|
||||
stop 4
|
||||
|
||||
if (omp_target_disassociate_ptr (c_loc (q), d) /= 0) &
|
||||
stop 5
|
||||
|
||||
if(c_associated (omp_get_mapped_ptr (c_loc (q), d))) &
|
||||
stop 6
|
||||
end if
|
||||
|
||||
call omp_target_free (p, d)
|
||||
end program main
|
||||
|
|
@ -0,0 +1,175 @@
|
|||
program main
|
||||
use omp_lib
|
||||
use iso_c_binding
|
||||
implicit none (external, type)
|
||||
integer :: d, id, i, j
|
||||
integer, target :: a, b(1:2), c(1:2), e(0:127)
|
||||
type(c_ptr) :: p1, p2, p3, q, devptrs(0:63)
|
||||
|
||||
a = 42;
|
||||
q = c_loc (a);
|
||||
e = [(i, i = 0, 127)]
|
||||
|
||||
d = omp_get_default_device ()
|
||||
id = omp_get_initial_device ()
|
||||
|
||||
if (d < 0 .or. d >= omp_get_num_devices ()) &
|
||||
d = id
|
||||
|
||||
if (d /= id) then
|
||||
!$omp target data map(alloc: a, b, c(2), e(32:95)) device(d)
|
||||
!$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
|
||||
p1 = c_loc (a);
|
||||
p2 = c_loc (b);
|
||||
p3 = c_loc (c(2))
|
||||
devptrs = [(c_loc (e(i)), i = 32, 95)]
|
||||
!$omp end target
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), p2) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
|
||||
stop 0
|
||||
|
||||
do j = 0, 31
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
|
||||
stop 1
|
||||
end do
|
||||
do j = 32, 95
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) &
|
||||
stop 2
|
||||
end do
|
||||
do j = 96, 128
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
|
||||
stop 3
|
||||
end do
|
||||
!$omp end target data
|
||||
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (q, d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
|
||||
stop 4
|
||||
do j = 0, 127
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
|
||||
stop 5
|
||||
end do
|
||||
|
||||
!$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d)
|
||||
!$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
|
||||
p1 = c_loc (a);
|
||||
p2 = c_loc (b);
|
||||
p3 = c_loc (c(2))
|
||||
devptrs = [(c_loc (e(i)), i = 32, 95)]
|
||||
!$omp end target
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
|
||||
stop 6
|
||||
|
||||
do j = 0, 31
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
|
||||
stop 7
|
||||
end do
|
||||
do j = 32, 95
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) &
|
||||
stop 8
|
||||
end do
|
||||
do j = 96, 128
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
|
||||
stop 9
|
||||
end do
|
||||
!$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d)
|
||||
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (q, d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d))) &
|
||||
stop 10
|
||||
do j = 0, 127
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
|
||||
stop 11
|
||||
end do
|
||||
|
||||
else ! d == id
|
||||
|
||||
!$omp target data map(alloc: a, b, c(2), e(32:95)) device(d)
|
||||
!$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
|
||||
p1 = c_loc (a);
|
||||
p2 = c_loc (b);
|
||||
p3 = c_loc (c(2))
|
||||
devptrs = [(c_loc (e(i)), i = 32, 95)]
|
||||
!$omp end target
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
|
||||
stop 12
|
||||
|
||||
do j = 0, 127
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
|
||||
stop 13
|
||||
end do
|
||||
!$omp end target data
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
|
||||
stop 14
|
||||
do j = 0, 127
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
|
||||
stop 15
|
||||
end do
|
||||
|
||||
!$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d)
|
||||
!$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
|
||||
p1 = c_loc (a);
|
||||
p2 = c_loc (b);
|
||||
p3 = c_loc (c(2))
|
||||
devptrs = [(c_loc (e(i)), i = 32, 95)]
|
||||
!$omp end target
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
|
||||
stop 16
|
||||
|
||||
do j = 0, 127
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
|
||||
stop 17
|
||||
end do
|
||||
!$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d)
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
|
||||
stop 18
|
||||
|
||||
do j = 0, 127
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
|
||||
stop 19
|
||||
end do
|
||||
end if
|
||||
|
||||
end program main
|
||||
|
|
@ -0,0 +1,48 @@
|
|||
program main
|
||||
use omp_lib
|
||||
use iso_c_binding
|
||||
implicit none (external, type)
|
||||
integer :: d, id
|
||||
type(c_ptr) :: p, p1, p2
|
||||
integer, target :: a(1:0), b(1:2)
|
||||
|
||||
d = omp_get_default_device ()
|
||||
id = omp_get_initial_device ()
|
||||
|
||||
if (d < 0 .or. d >= omp_get_num_devices ()) &
|
||||
d = id
|
||||
|
||||
p = omp_target_alloc (c_sizeof (c_int), d)
|
||||
if (.not. c_associated (p)) &
|
||||
stop 0 ! okay
|
||||
|
||||
if (omp_target_associate_ptr (c_loc (a), p, c_sizeof (c_int), &
|
||||
0_c_size_t, d) == 0) then
|
||||
|
||||
if(.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p)) &
|
||||
stop 1
|
||||
|
||||
if (omp_target_disassociate_ptr (c_loc (a), d) /= 0) &
|
||||
stop 2
|
||||
|
||||
if(c_associated (omp_get_mapped_ptr (c_loc (a), d))) &
|
||||
stop 3
|
||||
|
||||
!$omp target data map(alloc: a) device(d)
|
||||
!$omp target map(from: p1) map(alloc: a) device(d)
|
||||
p1 = c_loc (a);
|
||||
!$omp end target
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (a), d))) &
|
||||
stop 4
|
||||
!$omp end target data
|
||||
|
||||
!$omp target data map(alloc: b(1:0)) device(d)
|
||||
!$omp target map(from: p2) map(alloc: b(1:0)) device(d)
|
||||
p2 = c_loc (b(1));
|
||||
!$omp end target
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (b(1)), d))) &
|
||||
stop 5
|
||||
!$omp end target data
|
||||
end if
|
||||
call omp_target_free (p, d)
|
||||
end program main
|
||||
|
|
@ -0,0 +1,84 @@
|
|||
program main
|
||||
use omp_lib
|
||||
use iso_c_binding
|
||||
implicit none (external, type)
|
||||
integer :: d, id
|
||||
type(c_ptr) :: p1, p2
|
||||
|
||||
type t
|
||||
integer :: m1, m2
|
||||
end type t
|
||||
type(t), target :: s
|
||||
|
||||
d = omp_get_default_device ()
|
||||
id = omp_get_initial_device ()
|
||||
|
||||
if (d < 0 .or. d >= omp_get_num_devices ()) &
|
||||
d = id
|
||||
|
||||
if (d /= id) then
|
||||
!$omp target data map(alloc: s, s%m2) device(d)
|
||||
!$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
|
||||
p1 = c_loc (s);
|
||||
p2 = c_loc (s%m2);
|
||||
!$omp end target
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) &
|
||||
stop 0
|
||||
!$omp end target data
|
||||
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) &
|
||||
stop 1
|
||||
|
||||
!$omp target enter data map (alloc: s, s%m2) device (d)
|
||||
!$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
|
||||
p1 = c_loc (s);
|
||||
p2 = c_loc (s%m2);
|
||||
!$omp end target
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) &
|
||||
stop 2
|
||||
!$omp target exit data map (delete: s, s%m2) device (d)
|
||||
|
||||
if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) &
|
||||
.or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) &
|
||||
stop 3
|
||||
|
||||
else ! d == id
|
||||
|
||||
!$omp target data map(alloc: s, s%m2) device(d)
|
||||
!$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
|
||||
p1 = c_loc (s);
|
||||
p2 = c_loc (s%m2);
|
||||
!$omp end target
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
|
||||
stop 4
|
||||
!$omp end target data
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
|
||||
stop 5
|
||||
|
||||
!$omp target enter data map (alloc: s, s%m2) device (d)
|
||||
!$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
|
||||
p1 = c_loc (s);
|
||||
p2 = c_loc (s%m2);
|
||||
!$omp end target
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
|
||||
stop 6
|
||||
|
||||
!$omp target exit data map (delete: s, s%m2) device (d)
|
||||
|
||||
if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
|
||||
.or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
|
||||
stop 7
|
||||
end if
|
||||
|
||||
end program main
|
||||
Loading…
Reference in New Issue