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

[Joel E. Denny via Openmp-dev]

Hi Jose,

Sorry, I accidentally replied off list earlier.  Trying again, but with a
little more information....

On Tue, Mar 17, 2020 at 2:11 PM Jose Monsalve Diaz <josem at udel.edu> wrote:

> 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.
>

I agree that some of the permutations I mentioned for my original example
would be illegal under 4.5.

My original example does not have an explicit map clause on the `target
teams` directive.  If the implied map clause in that example should also be
illegal under 4.5, then we're missing that diagnostic as well.  However, my
suspicion is that this was just a bug in 4.5, and we shouldn't try too hard
to make sense of it.


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

Yes, `-fopenmp-version=50` relaxes that restriction.

I just noticed another change in the spec.  The Clang source quotes the
following for the diagnostic we're discussing:

OpenMP 4.5 sec. 2.15.5.1, map Clause, Restrictions:

> If any part of the original storage of a list item has corresponding
> storage in the device data
> environment, all of the original storage must have corresponding storage
> in the device data
> environment.
>

However, that's changed in OpenMP 5.0:

OpenMP 5.0 sec. 2.19.7.1, map Clause Restrictions:

> If any part of the original storage of a list item with an explicit
> data-mapping attribute has
> corresponding storage in the device data environment prior to a task
> encountering the construct
> associated with the map clause, all of the original storage must have
> corresponding storage in the
> device data environment prior to the task encountering the construct.
>

Does the phrase "with an explicit data-mapping attribute" refer to the map
clause itself?  That is, is this restriction relaxed when the map clause is
implicit, as in my original example?

Thanks.

Joel


>
> 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/20200319/c3ef9d2f/attachment.html>