<div dir="ltr">Hi,<div><br></div><div>I'm working on this implementation. Any suggestions will be greatly appreciated!</div><div><br></div><div><p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><b><span style="font-family:"Times New Roman",serif">Background</span></b></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">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.</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">Example 1:</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">struct S {</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New""> int len;</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New""> int *d;</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">};</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp declare mapper(struct
S s) map(s, s.d[0:s.len])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New""> </span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">struct S s;</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target map(s)</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">{}</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif"> </span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">The </span><span style="font-family:"Courier New"">map(s)</span><span style="font-family:"Times New Roman",serif">clause
above is essentially equivalent to </span><span style="font-family:"Courier New"">map(s,
s.d[0:s.len])</span><span style="font-family:"Times New Roman",serif"> in this
example. </span><span style="font-family:"Times New Roman",serif">Details can be found in OpenMP 5.0 specification. </span><span style="font-family:"Times New Roman",serif;font-size:11pt">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. </span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><b><span style="font-family:"Times New Roman",serif">Implementation</span></b></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">The
implementation of user defined mapper includes 3 major components.</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">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:</span></p>
<p class="gmail-MsoListParagraphCxSpFirst" style="margin:0cm 0cm 0.0001pt 36pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">1.<span style="font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:7pt;line-height:normal;font-family:"Times New Roman"">
</span></span><span style="font-family:"Times New Roman",serif">Check
the clause type of declare mapper directive, and there must be at least one map
clause.</span></p>
<p class="gmail-MsoListParagraphCxSpMiddle" style="margin:0cm 0cm 0.0001pt 36pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">2.<span style="font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:7pt;line-height:normal;font-family:"Times New Roman"">
</span></span><span style="font-family:"Times New Roman",serif">Check
the mapped type, which must be struct, union or class in C/C++.</span></p>
<p class="gmail-MsoListParagraphCxSpLast" style="margin:0cm 0cm 8pt 36pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">3.<span style="font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:7pt;line-height:normal;font-family:"Times New Roman"">
</span></span><span style="font-family:"Times New Roman",serif">Check
if there is redeclaration of the same mapper identifier for the same or
compatible types (same canonical type).</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">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.</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">3) The
code generation of map clause with a mapper associated.</span></p>
<p class="gmail-MsoListParagraphCxSpFirst" style="margin:0cm 0cm 0.0001pt 36pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">1.<span style="font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:7pt;line-height:normal;font-family:"Times New Roman"">
</span></span><span style="font-family:"Times New Roman",serif">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.</span></p>
<p class="gmail-MsoListParagraphCxSpMiddle" style="margin:0cm 0cm 0.0001pt 36pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">2.<span style="font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:7pt;line-height:normal;font-family:"Times New Roman"">
</span></span><span style="font-family:"Times New Roman",serif">When
we encounter a mapped variable or array section with a mapper:</span></p>
<p class="gmail-MsoListParagraphCxSpMiddle" style="margin:0cm 0cm 0.0001pt 72pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">a.<span style="font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:7pt;line-height:normal;font-family:"Times New Roman"">
</span></span><span style="font-family:"Times New Roman",serif">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.</span></p>
<p class="gmail-MsoListParagraphCxSpMiddle" style="margin:0cm 0cm 0.0001pt 72pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">b.<span style="font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:7pt;line-height:normal;font-family:"Times New Roman"">
</span></span><span style="font-family:"Times New Roman",serif">If
it is an array section,</span></p>
<p class="gmail-MsoListParagraphCxSpMiddle" style="margin:0cm 0cm 0.0001pt 108pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">i.<span style="font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:7pt;line-height:normal;font-family:"Times New Roman""> </span></span><span style="font-family:"Times New Roman",serif">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.</span></p>
<p class="gmail-MsoListParagraphCxSpMiddle" style="margin:0cm 0cm 0.0001pt 108pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">ii.<span style="font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:7pt;line-height:normal;font-family:"Times New Roman""> </span></span><span style="font-family:"Times New Roman",serif">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).</span></p>
<p class="gmail-MsoListParagraphCxSpLast" style="margin:0cm 0cm 8pt 36pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">3.<span style="font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:7pt;line-height:normal;font-family:"Times New Roman"">
</span></span><span style="font-family:"Times New Roman",serif">Repeat
Step 2 if in the expanded map clauses, another map clause with mapper is found.
</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">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.</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">3 examples
are shown below.</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">Example 2:</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp declare mapper(struct
S s) map(s, s.d[0:s.len])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">struct S *p;</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target map(p[0:N])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">// The code above is translated
into the code equivalent to</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target enter data
map(alloc:p[0:N])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">for (int i = 0; i < N; i++)</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New""> #pragma omp target enter data map(to:p[i],
p[i].d[0:p[i].len])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">{}</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">for (int i = N-1; i >= 0;
i--)</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New""> #pragma omp target exit data map(from:p[i],
p[i].d[0:p[i].len])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target exit data
map(release:p[0:N])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif"> </span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">Example 3:</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp declare mapper(struct
S s) map(s) map(alloc:s.d[0:s.len])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target map(p[0:N])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">// The code above is translated
into the code equivalent to</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target enter data
map(alloc:p[0:N])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">for (int i = 0; i < N; i++)</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New""> #pragma omp target enter data map(to:p[i]) \\</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New""> map(alloc:p[i].d[0:p[i].len])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">{}</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">for (int i = N-1; i >= 0;
i--)</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New""> #pragma omp target exit data map(from:p[i])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target exit data
map(release:p[0:N])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif"> </span></p>
<p class="MsoNormal" style="margin:0cm 0cm 8pt;line-height:107%;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Times New Roman",serif">Example 4:</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp declare mapper(struct
S a) map(a.len, a.d[0:a.len])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp declare mapper(struct
SS a) map(a.k)</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">// Assume that s has type
struct S, and s.d[x] has type struct SS.</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target enter data map(s)</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">// The code above is translated
into the code equivalent to</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target enter data map(s.len)
map(s.d[0:s.len])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">// The code above is translated
into the code equivalent to</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">#pragma omp target enter data map(s.len)
map(alloc: s.d[0:s.len])</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New"">for (int i = 0; i < s.len;
i++)</span></p>
<p class="MsoNormal" style="margin:0cm 0cm 0.0001pt;line-height:normal;font-size:11pt;font-family:Calibri,sans-serif"><span style="font-family:"Courier New""> #pragma omp target enter data map(s.d[i].k)</span><span style="font-family:"Times New Roman",serif"></span></p></div><div><br></div><div>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.</div><div><br></div><div>Thanks a lot,</div><div>Lingda Li</div></div>