[cfe-dev] [RFC] The implementation of OpenMP 5.0 declare mapper directive

[Lingda Li via cfe-dev]

Hi,

I'm working on this implementation. Any suggestions will be greatly
appreciated!

*Background*

This proposal is to implement the user defined mapper, i.e., declare mapper
directive, a new feature in OpenMP 5.0. User defined mapper is introduced
to extend existing map clauses for the purpose of simplifying the copy of
complex data structures between host and device (i.e., deep copy). An
example is shown below.

Example 1:

struct S {

  int len;

  int *d;

};

#pragma omp declare mapper(struct S s) map(s, s.d[0:s.len])



struct S s;

#pragma omp target map(s)

{}



The map(s)clause above is essentially equivalent to map(s, s.d[0:s.len]) in
this example. Details can be found in OpenMP 5.0 specification. We can see
that the programming burden is alleviated using user defined mappers,
especially when there are many instances of such complex data structure
mapping.

*Implementation*

The implementation of user defined mapper includes 3 major components.

1) The parsing and semantic analysis of declare mapper directives.
Parsing/sema generates a user defined mapper declaration whenever a declare
target directive is encountered. The mapper declaration inherits from
ValueDecl and DeclContext, and is associated with the corresponding type
(e.g., struct S in Example 1). The mapper variable (e.g., s in Example 1)
is declared in this environment so it is also a DeclContext. Data sharing
attribute (DSA) stack is required to parsing associated map clauses.
Necessary checking is needed in this phase:

1.      Check the clause type of declare mapper directive, and there must
be at least one map clause.

2.      Check the mapped type, which must be struct, union or class in
C/C++.

3.      Check if there is redeclaration of the same mapper identifier for
the same or compatible types (same canonical type).

2) The parsing and sematic analysis of map clauses (as well as to and from
clauses, we treat them as map clauses for the rest part of this proposal)
to deal with potential associated mappers. Add a field Mapper in
OMPMapClause to represent the associated mapper with it. If
mapper(mapper-identifier) modifier is present, look for the specified
mapper with a certain type. Report error if it cannot be found. Otherwise,
look for the default mapper. If no eligible mapper is found, the Mapper
field is set to nullptr.

3) The code generation of map clause with a mapper associated.

1.      First, if this is a target directive (not data, enter data, exit
data, update) which having any map clause associated with mappers, the code
generation will pretend this is a target data directive with a target
directive inside. The imaginary target data directive will inherit all map
clauses from the original target directive. We will explain why this is
necessary to map array sections with user define mapper later.

2.      When we encounter a mapped variable or array section with a mapper:

a.      If it is not an array section, expand the original map clause using
the associated mapper. I.e., the corresponding map clauses in the mapper
will substitute the original map clause.

b.      If it is an array section,

i.     First, a normal mapping of the array section is performed. E.g., we
allocate space for the array section in target data enter directive, and
release space in target data exit directive. Nothing needs to be done for
target update directive here.

ii.     Then, a for loop is generated to map each individual element in the
array section. In the loop body, a corresponding target data function
(e.g., enter, exit, update) is emitted. The mapping of individual element
is the same as Step 2.a. This loop can be either after the array section
normal mapping mentioned in Step 2.b.i (target enter data), or before it
(target exit data).

3.      Repeat Step 2 if in the expanded map clauses, another map clause
with mapper is found.

We also need to identify the map type of the expanded map clauses, e.g.,
to, from, alloc, etc. This is done by combining the map type that the user
specifies, and the map type specified in the corresponding mapper
declaration. Step 1 has to be performed for target directive, because there
is no way to insert a loop after data mapping and before the execution for
target directive.

3 examples are shown below.

Example 2:

#pragma omp declare mapper(struct S s) map(s, s.d[0:s.len])

struct S *p;

#pragma omp target map(p[0:N])

// The code above is translated into the code equivalent to

#pragma omp target enter data map(alloc:p[0:N])

for (int i = 0; i < N; i++)

  #pragma omp target enter data map(to:p[i], p[i].d[0:p[i].len])

#pragma omp target

{}

for (int i = N-1; i >= 0; i--)

  #pragma omp target exit data map(from:p[i], p[i].d[0:p[i].len])

#pragma omp target exit data map(release:p[0:N])



Example 3:

#pragma omp declare mapper(struct S s) map(s) map(alloc:s.d[0:s.len])

#pragma omp target map(p[0:N])

// The code above is translated into the code equivalent to

#pragma omp target enter data map(alloc:p[0:N])

for (int i = 0; i < N; i++)

  #pragma omp target enter data map(to:p[i]) \\

      map(alloc:p[i].d[0:p[i].len])

#pragma omp target

{}

for (int i = N-1; i >= 0; i--)

  #pragma omp target exit data map(from:p[i])

#pragma omp target exit data map(release:p[0:N])



Example 4:

#pragma omp declare mapper(struct S a) map(a.len, a.d[0:a.len])

#pragma omp declare mapper(struct SS a) map(a.k)

// Assume that s has type struct S, and s.d[x] has type struct SS.

#pragma omp target enter data map(s)

// The code above is translated into the code equivalent to

#pragma omp target enter data map(s.len) map(s.d[0:s.len])

// The code above is translated into the code equivalent to

#pragma omp target enter data map(s.len) map(alloc: s.d[0:s.len])

for (int i = 0; i < s.len; i++)

  #pragma omp target enter data map(s.d[i].k)

Each of the 3 parts mentioned above will have a corresponding patch. By
far, I almost finish the first part and the patch should be out soon.

Thanks a lot,
Lingda Li
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20181206/5e9f308e/attachment.html>