Hi PA,
some first remarks; I still need to read the FE bits and come
back to some parts I was wondering about - and that I wanted to
test via the attached testcase.
In particular, I believe there are some issues:
* I think stride should be stride (offset) multiplier
(see below) - at least for Fortran, I think that's
needed - if the code I have below is valid.
* I have a parse error - and, IMGO, a bogus error message
* The IMHO valid testcase crashes - at least with nvptx
offloading - in omp_target_memcpy_rect_worker
(or actually in cuMemcpy3D_v2 invoked via GOMP_OFFLOAD_memcpy3d)
Can you check?
* * *
BTW: Possibly after the last patch in the series, we also want to
update https://gcc.gnu.org/onlinedocs/libgomp/OpenMP-5_002e0.html
From u'N'implemented to 'P'artial or even, 'Y'es - fully for the
three items: array shaping and array-sections with strides ('Y'
once C++ and C patches have landed) - and for
'Discontiguous array section with target update' either 'P'
(plus note) or 'Y'.
* * *
Paul-Antoine Arras wrote:
From: Julian Brown<[email protected]>
This patch adds support for OpenMP 5.0 strided updates and the
array-shaping operator ("([x][y][z]) foo[0:n]..."). This is mostly for
C++ only so far, though necessary changes have been made to the C FE to
adjust for changes to shared data structures.
I am wondering whether we need to take care of the following
issue that exists in Fortran:
type t2; ...; end type t2 !
type t
type(t2) :: x ! sizeof = 40
integer :: y ! sizeof = 4
end type t ! → sizeof = 44
type(t) :: arr(100)
type(t2) :: p(:)
p => arr(:)%x
Here, sizeof(p(1)) == 40
and loc(p(2)) - loc(p(1)) == 44.
And p2 => p(::2)
Here, the stride is syntax wise 2 but internally loc(p2(2)) - loc(p2(1)) = 2*44.
But that stride is not a multiple of sizeof(type(t2)) but of sizeof(type(t)).
The solution on the Fortran side is to use a stride multiplier ('sm') that is
a byte offset instead of a stride (with sm = stride * elem_size).
* * *
The question is whether we need to handle non-integer multiples
of the element size for strided updates? (And even if not, does
it make sense to be future proof?)
In particular:
* * *
Is the following valid?
!$omp target enter data map(arr) ! argument must be contiguous (*)
...
!$omp target update to(p(::)) ! or 'to(p(::2))'
The first one needs to be contiguous (quote from TR15):
"Additional restrictions to the map clause for mapping of arrays,
array sections, and structures are as follows:
* If a list item is an array or an array section, it must specify contiguous
storage."
And the wording in "12 Data-Motion Control" implies that this is
valid:
"If a list item is an array or array section then it is treated as
if it is replaced by each of its array elements in the clause."
albeit the following is slightly confusing:
* * *
If so, the stride to stride multiplier conversion of libgomp/target.c's
omp_target_memcpy_rect_worker, i.e.
+ && __builtin_mul_overflow (element_size, strides[0], &stride))
should be moved to the compiler itself.
I think comments / variable names should make clear whether they
refers to a multiple of the array-element size or a (stride)
offset/stride multiplier.
* * *
Syntax wise, Fortran permits 'var(:)%x' or 'var(::2)%x'; I think
OpenMP only permits an array section (w/ and w/o strides) when it
is the last designator, i.e. the issue only occurs when doing the
detour via either a pointer or via:
associate(p => arr(:)%x)
... to(p(:)) ! note: 'tp(p)' is invalid as not contiguous ... (**)
end associate
where associate is an alias, if the RHS counts as variable, then
it is also a variable, otherwise only the result of an expression.
(**) Namely, "12 Data-Motion Control" requires:
"* If an array appears as a list item in a data-motion clause and
it has corresponding storage in the device data environment, the
corresponding storage must correspond to a single mappable storage
block that was previously mapped.
• If a list item in a data-motion clause has corresponding storage
in the device data environment, all corresponding storage must
correspond to a single mappable storage block that was previously
mapped."
where the first one seems to apply to 'p' and the second one to
'p(:)' or 'p(4:6)' or 'p(::2)'.
* * *
TESTCASE
========
I have tried the attached testcase - but it fails to compile with
two errors, which look bogus to me:
int x[10][8][6];
...
#pragma omp target update from(x[: :4][: :3][: :2], ...)
foo.C:32:42: error: expected primary-expression before »:« token
32 | #pragma omp target update from(x[: :4][: :3][: :2], ...
| ^
And when inserting '10':
foo.C:33:66: error: length '10' with stride '4' above array section size in
'from' clause
33 | #pragma omp target update from(x[: 10 :4][: 8 :3][: 6 :2], ...
| ^
But I expect that I can access element 0, 4 and 8 using the [0:10:4] syntax?
Or did I miss something?
Replacing the literals by variables (int four = 4 etc.) makes it compile but
then it fails at runtime as mentioned above.
* * *
2026-08-13 Paul-Antoine Arras
gcc/c-family/
* c-omp.cc (omp_handle_noncontig_array): Propagate
OMP_CLAUSE_ITERATORS onto synthesised grid clauses.
gcc/cp/
* parser.cc (cp_parser_postfix_open_square_expression): Only
accept a 2nd colon (stride) for to/from clauses, not map.
(cp_parser_omp_var_list_no_open): Only count OMP_ARRAY_SECTION
layers, not every ARRAY_REF layer, in a to/from clause's base
decl.
* parser.h (cp_parser::omp_array_section_p): Replace bool with...
(cp_parser::omp_array_section_kind): ...this 3-state unsigned char.
(OMP_ARRAY_SECTION_NONE, OMP_ARRAY_SECTION_UNSTRIDED,
OMP_ARRAY_SECTION_STRIDED): Define.
* semantics.cc (handle_omp_array_sections): Handle a genuinely
discontiguous update at any access-chain depth.
(finish_omp_clauses): Splice the GOMP_MAP_TO_GRID/FROM_GRID
replacement clause in at *PC when DISCONTIGUOUS == 2.
gcc/
* gimplify.cc (gimplify_scan_omp_clauses): Handle
OMP_TARGET_UPDATE like the other target constructs.
(gimplify_adjust_omp_clauses): Fix a leaked gimplify context in
the GOMP_MAP_GRID_DIM/GRID_STRIDE case.
gimplify.cc (remove_unused_omp_iterator_vars): Exempt
GOMP_MAP_GRID_DIM/GRID_STRIDE clauses.
(gimplify_adjust_omp_clauses): Restrict the "unsupported map
expression" sorry to GOMP_MAP_GRID_DIM/GRID_STRIDE clauses.
(gimplify_omp_target_update): Reject an iterator combined with a
grid map clause.
I was wondering whether a testcase needs to be added for those?
(Including for C++26?)
* * *
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -11046,6 +11049,19 @@ omp_group_last (tree *start_p)
grp_last_p = &OMP_CLAUSE_CHAIN (c);
break;
+ case GOMP_MAP_TO_GRID:
+ case GOMP_MAP_FROM_GRID:
+ while (nc
+ && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE))
+ {
+ grp_last_p = &OMP_CLAUSE_CHAIN (c);
+ c = nc;
+ nc = OMP_CLAUSE_CHAIN (c);
+ }
Indentation issue.
* * *
+ /* The OMP_CLAUSE_DECL for GRID_DIM/GRID_STRIDE isn't necessarily
+ an lvalue -- e.g. it might be a constant. So handle it
+ specially here. */
[Off topic: OpenMP meanwhile also permits constexpr, const etc. in map
clauses – which the implementation may/has to ignore. But I guess we
want to handle them already in the front end - such that this comment
can stay.]
* * *
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1423,7 +1423,7 @@ DEFTREECODE (OMP_ATOMIC_CAPTURE_NEW,
"omp_atomic_capture_new", tcc_statement, 2)
DEFTREECODE (OMP_CLAUSE, "omp_clause", tcc_exceptional, 0)
/* An OpenMP array section. */
-DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 3)
+DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 4)
I wonder whether we should expand this, e.g. append "with strides" or
maybe even better "(decl, low bound, length, stride)" - such that one
knows what's actually stored in the four operands.
* * *
Tobias
PS: I still need to read the FE part of the patch and then come back to the
ME parts, trying to understand some fine print - but, possibly, my ME concerns
get resolved when the attached testcase works - either by fixing the compiler
or the testcase (or most likely: both).
int main() {
bool self_map = false;
#pragma omp target map(to: self_map)
self_map = true;
int x[10][8][6];
int *p = (int *) __builtin_malloc (sizeof (int) * (10*8)*8*6);
__builtin_printf("%ld\n", sizeof(x));
for (int i = 0; i < 10; i++)
for (int j = 0; j < 8; j++)
for (int k = 0; k < 6; k++)
{
x[i][j][k] = 100*i + 10*j + k;
p[i*10*8 + j*8 + k] = 100*i + 10*j + k;
}
#pragma omp target enter data map(to: x, p[:10*8*6])
#pragma omp target map(present, alloc: x, p[:10*8*6])
for (int i = 0; i < 10; i++)
for (int j = 0; j < 8; j++)
for (int k = 0; k < 6; k++)
{
x[i][j][k] *= 23;
p[i*10*8 + j*8 + k] *= 23;
}
#pragma omp target update from(x[: :4][: :3][: :2], (([10][8][6])p)[: :4][: :3][: :2])
// #pragma omp target update from(x[: 10 :4][: 8 :3][: 6 :2], (([10][8][6])p)[: 10 :4][: 8 :3][: 6 :2])
// int four = 4, three = 3, two = 2;
// #pragma omp target update from(x[: 10 :four][: 8 :three][: 6 :two], (([10][8][6])p)[: 10 :four][: 8 :three][: 6 :two])
if (self_map)
__builtin_printf("WARNING: Main tested SKIPPED because of self mapping or host fallback\n");
else
{
for (int i = 0; i < 10; i++)
for (int j = 0; j < 8; j++)
for (int k = 0; k < 6; k++)
{
bool dev_value = (i % 4 + j % 3 + i % 2) == 0;
int factor = dev_value ? 23 : 1;
if (x[i][j][k] != factor * (100*i + 10*j + k))
__builtin_abort ();
if (p[i*10*8 + j*8 + k] != factor * (100*i + 10*j + k))
__builtin_abort ();
}
}
#pragma omp target exit data map(from: x, p[:10*8*6])
for (int i = 0; i < 10; i++)
for (int j = 0; j < 8; j++)
for (int k = 0; k < 6; k++)
{
if (x[i][j][k] != 23 * (100*i + 10*j + k))
__builtin_abort ();
if (p[i*10*8 + j*8 + k] != 23 * (100*i + 10*j + k))
__builtin_abort ();
}
__builtin_free (p);
}
