Hi Julian,
first round of comments - I think I need a second pass as the patch is
long and complex. The depth of review also was decreasing, hence, I
assume I will spot things in later parts of the compiler.
In any case, I think the patch is a huge leap forward and very useful!
Contrary to previous iterations of the patch series, it looks as if
{1,2}/5 could be sensibly applied before {3,4,5}/5. However, those are
smaller or tiny, it should be way quicker to review them!
* * *
CROSS REFS
This patch is part of the patch set
"[PATCH v7 0/5] OpenMP/OpenACC: map clause and OMP gimplify rework"
https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627895.html
This patch set in turn is required for the follow-up patch set
https://gcc.gnu.org/pipermail/gcc-patches/2023-September/629363.html
'[PATCH 0/8] OpenMP: lvalue parsing and "declare mapper" support'
and it looks as if it would be very useful for some other WIP patches
like map+iterator.
* * *
The patch 1/5 is just an indentation patch to make the following patch
smaller - and it has been approved multiple times.
This patch can be found at
https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627897.html
in case someone wants to read the full patch or Julian's patch-intro
comments.
* * *
I do note that all 5 patches of this patch set apply cleanly with some
fuzzy line matching, but:
WARNING: at least for this patch (2/5), code gets wrongly applied
in (c-)parser.cc - see below.
On 19.08.23 00:47, Julian Brown wrote:
This patch reworks clause expansion in the C, C++ and (to a lesser
extent) Fortran front ends for OpenMP and OpenACC mapping nodes used in
GPU offloading support.
At present a single clause may be turned into several mapping nodes,
or have its mapping type changed, in several places scattered through
the front- and middle-end. The analysis relating to which particular
transformations are needed for some given expression has become quite hard
to follow. Briefly, we manipulate clause types in the following places:
[...]
[As the patch is very large, I copied the interesting bits out of the patch
and comment on them - I hope that preserves enough context to be useful.]
tree c_omp_address_inspector::get_address ()
This class member function is unused. It returns 'orig' and that member variable
is used a couple of times in the member functions - but not this function.
I think it can be removed. (I also find the name slightly misleading.)
* * *
+c_omp_address_inspector::maybe_zero_length_array_section (tree clause)
+{
+ switch (OMP_CLAUSE_MAP_KIND (clause))
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DELETE:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_FORCE_TOFROM:
+ case GOMP_MAP_FORCE_PRESENT:
Shouldn't this also include:
GOMP_MAP_ALWAYS_PRESENT_TO:
GOMP_MAP_ALWAYS_PRESENT_FROM:
GOMP_MAP_ALWAYS_PRESENT_TOFROM:
GOMP_MAP_PRESENT_ALLOC:
GOMP_MAP_PRESENT_TO:
?
* * *
+omp_expand_access_chain (tree c, tree expr, vec<omp_addr_token *> &addr_tokens,
+ unsigned *idx)
+{
+ using namespace omp_addr_tokenizer;
+ location_t loc = OMP_CLAUSE_LOCATION (c);
+ unsigned i = *idx;
+ tree c2 = NULL_TREE;
+ gomp_map_kind kind
+ = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM
+ || (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DELETE
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_RELEASE
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_FROM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_FROM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PRESENT_FROM)))
+ ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH;
GOMP_MAP_ALWAYS_PRESENT_FROM missing?
And I personally find it more readable to have
if (...
...
...)
kind = GOMP_MAP_DETACH;
else
kind = GOMP_MAP_ATTACH;
as in the assignment above, the lhs is in this case really far from the rhs.
* * *
I do see in omp_expand_access_chain, which is shared by C and C++, that the
following
accesses are handled:
ACCESS_POINTER,
ACCESS_POINTER_OFFSET,
ACCESS_INDEXED_ARRAY,
But not:
ACCESS_REF_TO_POINTER,
ACCESS_REF_TO_POINTER_OFFSET,
ACCESS_INDEXED_REF_TO_ARRAY
Do those need to be handled as well?
SIDENOTE: As Fortran allocatable components, reference-type variables in a C++
class ('class', 'struct')
are automatically be copied when doing 'map(var)' - contrary to pointer members
that need to be explicitly
be mapped (either at mapping side or via a mapper). (And contrary to Fortran
allocatables, reference-type members
are always "there" - not like with allocatables can change from unallocated to
allocated - with all compilations
of polymorphism, array sizes, string lengths, ...).
* * *
+tree
+c_omp_address_inspector::expand_array_base (tree c,
...
+ bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
+ bool implicit_p = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_IMPLICIT (c));
...
+ if (!openmp
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
How about using 'map_p' (twice), given that it is already
defined?
...
+ if (target)
+ {
+ c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ if (target
+ && !ref_to_ptr
Remove 'target &&' from the second condition.
* * *
+++ b/gcc/c/c-parser.cc
@@ -17659,7 +17659,8 @@ c_parser_omp_clause_detach (c_parser *parser, tree list)
static tree
c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
- const char *where, bool finish_p = true)
+ const char *where, bool finish_p = true,
+ bool target = false)
Here but also in gcc/c-family/c-omp.cc - I wonder whether this should be
target_p for
consistency with most other boolean variables. The code is not really
consistent,
but given that there is 'finish_p' here and map_p etc. in the other file, it
seems
to make sense to use 'target_p.
* * *
If I try the example at
https://discourse.llvm.org/t/how-do-i-map-an-array-of-pointers-using-array-notation/53640/3
it fails to compile with:
foo.c:11:34: error: ‘#pragma omp target enter data’ with map-type other than
‘to’, ‘tofrom’ or ‘alloc’ on ‘map’ clause
11 | #pragma omp target enter data map(to: a[i][0:1])
The problem is that fuzzy 'patch' mis-apply the patch for c-parser.cc's
@@ -22106,6 +22111,7 @@ c_parser_omp_target_enter_data (location_t loc,
c_parser *parser,
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
+ case GOMP_MAP_ATTACH:
break;
default:
map_seen |= 1;
@@ -22215,6 +22221,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser
*parser,
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
+ case GOMP_MAP_DETACH:
break;
default:
map_seen |= 1;
Namely after applying it, the result is:
@@ -23329,0 +23335 @@ c_parser_omp_target_data (location_t loc, c_parser
*parser, bool *if_p)
+ case GOMP_MAP_ATTACH:
@@ -23494,0 +23501 @@ c_parser_omp_target_enter_data (location_t loc, c_parser
*parser,
+ case GOMP_MAP_DETACH:
That is, the new code is added
to c_parser_omp_target_data and c_parser_omp_target_enter_data
instead of
to c_parser_omp_target_enter_data and c_parser_omp_target_exit_data.
* * *
Likewise issue for C++'s cp/parser.cc.
* * *
The first question is whether such code is also needed for 'target' and 'target
data'
besides 'target enter/exit data'. It looks as if omp_expand_access_chain should
also
apply here.
* * *
And trying the following snippet:
void f(){
int **a,i,j,n;
#pragma omp target enter data map(a[i][j:n])
#pragma omp target exit data map(a[i][j:n])
#pragma omp target data map(a[i][j:n])
{
#pragma omp target map(a[i][j:n])
;
}
}
Results - when putting the enums to enter/exit data
instead of to target data / target enter data:
g++ fails, unless I messed up, with:
fiof.c:4:36: error: ‘#pragma omp target exit data’ with map-type other than
‘from’, ‘tofrom’, ‘release’ or ‘delete’ on ‘map’ clause
4 | #pragma omp target exit data map(a[i][j:n])
| ^
fiof.c:5:31: error: ‘#pragma omp target data’ with map-type other than ‘to’,
‘from’, ‘tofrom’ or ‘alloc’ on ‘map’ clause
5 | #pragma omp target data map(a[i][j:n])
And gcc fails with:
fiof.c:3:36: error: ‘a’ appears more than once in map clauses
3 | #pragma omp target enter data map(a[i][j:n])
| ^
fiof.c:4:35: error: ‘a’ appears more than once in map clauses
4 | #pragma omp target exit data map(a[i][j:n])
| ^
fiof.c:5:30: error: ‘a’ appears more than once in map clauses
5 | #pragma omp target data map(a[i][j:n])
| ^
fiof.c:7:25: error: ‘a’ appears more than once in map clauses
7 | #pragma omp target map(a[i][j:n])
(Plus all '... data' ones: '... must contain at least one ...')
Can you fix this - and add an example in the spirit of the LLVM example above,
except that it should test all four 'target' variants.
Thanks,
Tobias
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634
München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas
Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht
München, HRB 106955
