[Openmp-dev] shared on combined target teams partially suppresses implicit map clause?

[Jose Monsalve Diaz via Openmp-dev]

We at the V&V SOLLVE project have also noticed this.

However, consider that in OMP 4.5 this is illegal.

Map restrictions
• A list item cannot appear in both a map clause and a data-sharing
attribute clause on the same construct.

For 5.0 I believe this is fixed, but it may have not been implemented yet.
I don't know.

On Tue, Mar 17, 2020 at 12:49 PM Joel E. Denny via Openmp-dev <
openmp-dev at lists.llvm.org> wrote:

> Hi,
>
> 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.
>
> # Example
>
> ```
> $ cat test.c
> #include <stdio.h>
> int main() {
>   int arr[] = {10, 20, 30, 40, 50};
>   #pragma omp target data map(tofrom:arr[1:3])
>   #pragma omp target teams num_teams(1) shared(arr)
>   arr[1] += 1000;
>   printf("%d\n", arr[1]);
>   return 0;
> }
> $ clang -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda test.c
> $ ./a.out
> 1020
> ```
>
> # Diagnostic suppression
>
> 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:
>
> ```
> test.c:6:3: error: original storage of expression in data environment is
> shared but data environment do not fully contain mapped expression storage
>   arr[1] += 1000;
>   ^~~
> test.c:4:38: note: used here
>   #pragma omp target data map(tofrom:arr[1:3])
>                                      ^~~~~~~~
> 1 error generated.
> ```
>
> 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.
>
> # OMPMapClause suppression in -ast-dump
>
> The `shared` clause also suppresses the implicit `OMPMapClause` for the
> `OMPTargetTeamsDirective` in the `-ast-dump` output for the example.
>
> # No suppression in codegen
>
> 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).
>
> # Build
>
> I'm testing LLVM built from 378b1e60809d (a master commit from yesterday)
> with subprojects clang and openmp.
>
> # Summary
>
> 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?
>
> Thanks.
>
> Joel
> _______________________________________________
> Openmp-dev mailing list
> Openmp-dev at lists.llvm.org
> https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20200317/2aefb48c/attachment.html>