declare mapper Directive
6.9. declare mapper Directive#
The following examples show how to use the declare mapper directive to prescribe a map for later use. It is also quite useful for pre-defining partitioned and nested structure elements.
In the first example the declare mapper directive specifies that any structure of type myvec_t for which implicit data-mapping rules apply will be mapped according to its map clause. The variable v is used for referencing the structure and its elements within the map clause. Within the map clause the v variable specifies that all elements of the structure are to be mapped. Additionally, the array section v.data[0:v.len] specifies that the dynamic storage for data is to be mapped.
Within the main program the s variable is typed as myvec_t . Since the variable is found within the target region and the type has a mapping prescribed by a declare mapper directive, it will be automatically mapped according to its prescription: full structure, plus the dynamic storage of the data element.
//%compiler: clang
//%cflags: -fopenmp
/*
* name: target_mapper.1
* type: C
* version: omp_5.0
*/
#include <stdlib.h>
#include <stdio.h>
#define N 100
typedef struct myvec{
size_t len;
double *data;
} myvec_t;
#pragma omp declare mapper(myvec_t v) \
map(v, v.data[0:v.len])
void init(myvec_t *s);
int main(){
myvec_t s;
s.data = (double *)calloc(N,sizeof(double));
s.len = N;
#pragma omp target
init(&s);
printf("s.data[%d]=%lf\n",N-1,s.data[N-1]); //s.data[99]=99.000000
}
void init(myvec_t *s)
{ for(int i=0; i<s->len; i++) s->data[i]=i; }
!!%compiler: gfortran
!!%cflags: -fopenmp
! name: target_mapper.1
! type: F-free
! version: omp_5.0
module my_structures
type myvec_t
integer :: len
double precision, pointer :: data(:)
end type
end module
program main
use my_structures
integer, parameter :: N=100
!$omp declare mapper(myvec_t :: v) &
!$omp& map(v, v%data(1:v%len))
type(myvec_t) :: s
allocate(s%data(N))
s%data(1:N) = 0.0d0
s%len = N
!$omp target
call init(s)
!$omp end target
print*,"s%data(",N,")=",s%data(N) !! s%data( 100 )=100.000000000000
end program
subroutine init(s)
use my_structures
type(myvec_t) :: s
s%data = [ (i, i=1,s%len) ]
end subroutine
The next example illustrates the use of the mapper-identifier and deep copy within a structure. The structure, dzmat_t , represents a complex matrix, with separate real ( r_m ) and imaginary ( i_m ) elements. Two map identifiers are created for partitioning the dzmat_t structure.
For the C/C++ code the first identifier is named top_id and maps the top half of two matrices of type dzmat_t ; while the second identifier, bottom_id , maps the lower half of two matrices. Each identifier is applied to a different target construct, as map(mapper(top_id), tofrom: a,b) and map(mapper(bottom_id), tofrom: a,b). Each target offload is allowed to execute concurrently on two different devices ( 0 and 1 ) through the nowait clause.
The Fortran code uses the left_id and right_id map identifiers in the map(mapper(left_id),tofrom: a,b) and map(mapper(right_id),tofrom: a,b) map clauses. The array sections for these left and right contiguous portions of the matrices were defined previously in the declare mapper directive.
Note, the is and ie scalars are firstprivate by default for a target region, but are declared firstprivate anyway to remind the user of important firstprivate data-sharing properties required here.
//%compiler: clang
//%cflags: -fopenmp
/*
* name: target_mapper.2
* type: C
* version: omp_5.0
*/
#include <stdio.h>
// N MUST BE EVEN
#define N 100
typedef struct dzmat
{
double r_m[N][N];
double i_m[N][N];
} dzmat_t;
#pragma omp declare mapper( top_id: dzmat_t v) \
map(v.r_m[0:N/2][0:N], \
v.i_m[0:N/2][0:N] )
#pragma omp declare mapper(bottom_id: dzmat_t v) \
map(v.r_m[N/2:N/2][0:N], \
v.i_m[N/2:N/2][0:N] )
//initialization
void dzmat_init(dzmat_t *z, int is, int ie, int n);
//matrix add: c=a+b
void host_add( dzmat_t *a, dzmat_t *b, dzmat_t *c, int n);
int main()
{
dzmat_t a,b,c;
int is,ie;
is=0; ie=N/2-1; //top N/2 rows on device 0
#pragma omp target map(mapper(top_id), tofrom: a,b) device(0) \
firstprivate(is,ie) nowait
{
dzmat_init(&a,is,ie,N);
dzmat_init(&b,is,ie,N);
}
is=N/2; ie=N-1; //bottom N/2 rows on device 1
#pragma omp target map(mapper(bottom_id), tofrom: a,b) device(1) \
firstprivate(is,ie) nowait
{
dzmat_init(&a,is,ie,N);
dzmat_init(&b,is,ie,N);
}
#pragma omp taskwait
host_add(&a,&b,&c,N);
}
!!%compiler: gfortran
!!%cflags: -fopenmp
! name: target_mapper.2
! type: F-free
! version: omp_5.0
module complex_mats
integer, parameter :: N=100 !N must be even
type dzmat_t
double precision :: r_m(N,N), i_m(N,N)
end type
!$omp declare mapper( left_id: dzmat_t :: v) map( v%r_m(N, 1:N/2), &
!$omp& v%i_m(N, 1:N/2))
!$omp declare mapper(right_id: dzmat_t :: v) map( v%r_m(N,N/2+1:N), &
!$omp& v%i_m(N,N/2+1:N))
end module
program main
use complex_mats
type(dzmat_t) :: a,b,c
external dzmat_init, host_add !initialization and matrix add: a=b+c
integer :: is,ie
is=1; ie=N/2 !left N/2 columns on device 0
!$omp target map(mapper( left_id), tofrom: a,b) device(0) &
!$omp& firstprivate(is,ie) nowait
call dzmat_init(a,is,ie)
call dzmat_init(b,is,ie)
!$omp end target
is=N/2+1; ie=N !right N/2 columns on device 1
!$omp target map(mapper(right_id), tofrom: a,b) device(1) &
!$omp& firstprivate(is,ie) nowait
call dzmat_init(a,is,ie)
call dzmat_init(b,is,ie)
!$omp end target
!$omp taskwait
call host_add(a,b,c)
end program main
In the third example myvec structures are nested within a mypoints structure. The myvec_t type is mapped as in the first example. Following the mypoints structure declaration, the mypoints_t type is mapped by a declare mapper directive. For this structure the hostonly_data element will not be mapped; also the array section of x ( v.x[:1] ) and x will be mapped; and scratch will be allocated and used as scratch storage on the device. The default map-type mapping, tofrom, applies to the x array section, but not to scratch which is explicitly mapped with the alloc map-type. Note: the variable v is not included in the map list (otherwise the hostonly_data would be mapped) - just the elements to be mapped are listed.
The two mappers are combined when a mypoints_t structure type is mapped, because the mapper myvec_t structure type is used within a mypoints_t type structure.
//%compiler: clang
//%cflags: -fopenmp
/*
* name: target_mapper.3
* type: C
* version: omp_5.0
*/
#include <stdlib.h>
#include <stdio.h>
#define N 100
typedef struct myvec {
size_t len;
double *data;
} myvec_t;
#pragma omp declare mapper(myvec_t v) \
map(v, v.data[0:v.len])
typedef struct mypoints {
struct myvec scratch;
struct myvec *x;
double hostonly_data[500000];
} mypoints_t;
#pragma omp declare mapper(mypoints_t v) \
map(v.x, v.x[0] ) map(alloc:v.scratch)
void init_mypts_array(mypoints_t *P, int n);
void eval_mypts_array(mypoints_t *P, int n);
int main(){
mypoints_t P;
init_mypts_array(&P, N);
#pragma omp target map(P)
eval_mypts_array(&P, N);
}
!!%compiler: gfortran
!!%cflags: -fopenmp
! name: target_mapper.3
! type: F-free
! version: omp_5.0
module my_structures
type myvec_t
integer :: len
double precision, pointer :: data(:)
end type
!$omp declare mapper(myvec_t :: v) &
!$omp& map(v)
type mypoints_t
type(myvec_t) :: scratch
type(myvec_t), pointer :: x(:)
double precision :: hostonly_data(500000)
end type
!$omp declare mapper(mypoints_t :: v) &
!$omp& map(v%x, v%x(1)) map(alloc:v%scratch)
end module
program main
use my_structures
external init_mypts_array, eval_mypts_array
type(mypoints_t) :: P
call init_mypts_array(P)
!$omp target map(P)
call eval_mypts_array(P)
end program