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], ...
| ^
OpenMP specifies that if length is absent, it should be:
⌈(size − lower-bound)/stride⌉
(note the round up)
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], ...
| ^
OK – C/C++ defines this differently than Fortran - in Fortran, it is
a bit like 'for (i = lower; i <= upper; i += stride)' but in OpenMP
it is for C/C++:
{ lower-bound,
lower-bound + stride,
lower-bound + 2 * stride,
... ,
lower-bound + ((length - 1) * stride) }
Thus, replace in my example:
// Still valid and still a parse error:
#pragma omp target update from(x[: :4][: :3][: :2], (([10][8][6])p)[: :4][:
:3][: :2])
// WRONG:
// #pragma omp target update from(x[: 10 :4][: 8 :3][: 6 :2],
(([10][8][6])p)[: 10 :4][: 8 :3][: 6 :2])
// SHOULD BE
#pragma omp target update from(x[: (10+3)/4 :4][: (8+2)/3 :3][: (6+1)/2 :2],
(([10][8][6])p)[: (10+3)/4 :4][: (8+2)/3 :3][: (6+1)/2 :2])
* * *
However, I still run into the same run-time problem:
SEGFAULT in cuMemcpy3D_v2
... but the values seem to be fine at a glance:
#0 omp_target_memcpy_rect_worker (dst=0x7fffffffc4e8, src=0x7fffd97009f8, element_size=4,
num_dims=num_dims@entry=3, volume=0x4ad1d0 <main::.omp_len.31>, strides=0x4ad1f0
<main::.omp_stride.32>,
dst_offsets=0x5e0280 <main::.omp_index.30>, src_offsets=0x5e0280 <main::.omp_index.30>,
dst_dimensions=0x4ad1b0 <main::.omp_dim.29>, src_dimensions=0x4ad1b0 <main::.omp_dim.29>,
dst_devicep=0x0, src_devicep=0x627010,
tmp_size=0x7fffffffc440, tmp=0x7fffffffc448) at
/home/tob/repos/gcc/libgomp/target.c:5283
5283 {
(gdb) p volume[0]
$17 = 3
(gdb) p volume[1]
$18 = 3
(gdb) p volume[2]
$19 = 3
(gdb) p stride[0]
No symbol "stride" in current context.
(gdb) p strides[0]
$20 = 4
(gdb) p strides[1]
$21 = 3
(gdb) p strides[2]
$22 = 2
(gdb) p src_offset
No symbol "src_offset" in current context.
(gdb) p src_offsets
$23 = (const size_t *) 0x5e0280 <main::.omp_index.30>
(gdb) p src_offsets[0]
$24 = 0
(gdb) p src_offsets[1]
$25 = 0
(gdb) p src_offsets[2]
$26 = 0
(gdb) p dst_offsets[2]
$27 = 0
(gdb) p dst_offsets[1]
$28 = 0
(gdb) p dst_offsets[0]
$29 = 0
(gdb) p src_dimensions[0]
$30 = 10
(gdb) p src_dimensions[1]
$31 = 8
(gdb) p src_dimensions[2]
$32 = 6
(gdb) p dst_dimensions[2]
$33 = 6
(gdb) p dst_dimensions[1]
$34 = 8
(gdb) p dst_dimensions[0]
$35 = 10
This then calls:
GOMP_OFFLOAD_memcpy3d (dst_ord=-1, src_ord=0, dim2_size=12, dim1_len=3,
dim0_len=3, dst=0x7fffffffc4e8, dst_offset2_size=0, dst_offset1_len=0,
dst_offset0_len=0, dst_dim2_size=24, dst_dim1_len=8, src=0x7fffd97009f8,
src_offset2_size=0, src_offset1_len=0, src_offset0_len=0, src_dim2_size=24,
src_dim1_len=8) at /home/tob/repos/gcc/libgomp/plugin/plugin-nvptx.c:2512
2512 {
where I don't immediately see whether that's okay or not.
* * *
BTW: The following program fails with both the C and the C++ compiler:
void f()
{
int x[10];
int y[10][10][10];
#pragma omp target update to(x[0:0:-1]) // invalid - missing diagnostic:
stride must evaluate to a positive integer.
#pragma omp target update to(x[0:0:0]) // invalid - missing diagnostic:
stride must evaluate to a positive integer.
#pragma omp target update to(y[0:0:0]) // ICE (segfault)
}
Tobias
