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