<div dir="ltr">We at the V&V SOLLVE project have also noticed this. <div><br></div><div>However, consider that in OMP 4.5 this is illegal. </div><div><br></div><div>Map restrictions</div><div>• A list item cannot appear in both a map clause and a data-sharing attribute clause on the same construct.<br></div><div><br></div><div>For 5.0 I believe this is fixed, but it may have not been implemented yet. I don't know. </div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Tue, Mar 17, 2020 at 12:49 PM Joel E. Denny via Openmp-dev <<a href="mailto:openmp-dev@lists.llvm.org">openmp-dev@lists.llvm.org</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"><div dir="ltr"><div>Hi,</div><div></div><div><br></div><div></div><div>For some purposes but not others, the `shared` clause on a combined OpenMP `target teams` directive appears to suppress the implicit `map` clause for an array when compiling with Clang.<br></div><div><br></div><div># Example<br></div><div><br></div><div>```<br></div><div><span style="font-family:monospace">$ cat test.c<br></span></div><span style="font-family:monospace">#include <stdio.h><br>int main() {<br> int arr[] = {10, 20, 30, 40, 50};<br> #pragma omp target data map(tofrom:arr[1:3])<br> #pragma omp target teams num_teams(1) shared(arr)<br> arr[1] += 1000; <br> printf("%d\n", arr[1]);<br> return 0;<br></span><div><span style="font-family:monospace">}</span></div><div><span style="font-family:monospace">$ clang -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda test.c<br></span></div><div><span style="font-family:monospace">$ ./a.out</span></div><div><span style="font-family:monospace">1020</span></div><div><span style="font-family:arial,sans-serif">```</span></div><div><span style="font-family:arial,sans-serif"><br></span></div><div><span style="font-family:monospace"><span style="font-family:arial,sans-serif"># Diagnostic suppression</span><br></span></div><div><span style="font-family:arial,sans-serif"><br></span></div><div><span style="font-family:arial,sans-serif">In the example, it seems `map(tofrom:arr)` should be implied on the combined target teams directive, and that should produce a diagnostic because the enclosing target data has `map(tofrom:arr[1:3])` instead. That is, I expect:<br></span></div><div><span style="font-family:arial,sans-serif"><br></span></div><div><div><span style="font-family:monospace">```<br></span></div><div><span style="font-family:monospace">test.c:6:3: error: original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage<br> arr[1] += 1000;<br> ^~~<br>test.c:4:38: note: used here<br> #pragma omp target data map(tofrom:arr[1:3])<br> ^~~~~~~~<br>1 error generated.<br></span></div><div><span style="font-family:monospace">```</span></div><span style="font-family:arial,sans-serif"></span></div><div><span style="font-family:arial,sans-serif"><br></span></div><div><span style="font-family:arial,sans-serif">Indeed, I see the expected diagnostic if either I remove `shared(arr)`, or I add an explicit `map(tofrom:arr)`, or I split the combined `target teams` directive into separate but equivalent directives, or I perform some combination of these changes.<br></span></div><div><span style="font-family:arial,sans-serif"><br></span></div><div><span style="font-family:arial,sans-serif"># OMPMapClause su</span>ppression in -ast-dump<br></div><div><br></div><div>The `shared` clause also suppresses the implicit `OMPMapClause` for the `OMPTargetTeamsDirective` in the `-ast-dump` output for the example.<br></div><div><br></div><div># No suppression in codegen<br></div><div><br></div><div>Nevertheless, the generated LLVM IR for the example shows that `map(tofrom:arr)` is in effect for the combined target teams directive. That is, the offload map type passed to `__tgt_target_teams` is 0x223, which is `OMP_TGT_MAPTYPE_IMPLICIT | OMP_TGT_MAPTYPE_TARGET_PARAM | OMP_TGT_MAPTYPE_FROM | OMP_TGT_MAPTYPE_TO`, and the offload size is 20 (that is, 5 ints, the full size of the array).</div><div></div><div><div></div></div><div><br></div><div><div></div><div># Build</div><div><br></div><div>I'm testing LLVM built from 378b1e60809d (a master commit from yesterday) with subprojects clang and openmp.</div><div><br></div><div># Summary</div><div><br></div><div>My understanding is that the `shared` clause should never suppress the implicit `map` clause for an array on a combined `target teams` directive, and so it should not suppress related diagnostics or AST nodes. Is that correct, or is Clang correct?<br></div><div></div></div><div><br></div><div>Thanks.</div><div><br></div><div>Joel<br></div></div>
_______________________________________________<br>
Openmp-dev mailing list<br>
<a href="mailto:Openmp-dev@lists.llvm.org" target="_blank">Openmp-dev@lists.llvm.org</a><br>
<a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev" rel="noreferrer" target="_blank">https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev</a><br>
</blockquote></div>