[RFC] The implementation of OpenMP 5.0 declare mapper directive

classic Classic list List threaded Threaded
1 message Options
Reply | Threaded
Open this post in threaded view
|

[RFC] The implementation of OpenMP 5.0 declare mapper directive

Richard Smith 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

_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev