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

Joel E. Denny via Openmp-dev openmp-dev at lists.llvm.org
Tue Mar 17 10:49:43 PDT 2020


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
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/openmp-dev/attachments/20200317/6c77948d/attachment.html>


More information about the Openmp-dev mailing list