Hi Sandra, hi all,
I spend a bit more time on this to try to draft a better C testcase.
For some reasons, it seems to be difficult to create a testcase that does:
- copy two chunks of data to the device, the pointee x[i].ptr[:2]
- copy the struct x[0] + x[2] to the device - or just all: x[:4]
- copy the pointer 'x' to the device
- pointer attach x[:4] to x
- pointer attach x[i].ptr[:2] to x[i].ptr
Especially, when mapping the pointee 'x[i].ptr[:2]', GCC always also
struct maps 'x[i].ptr' - which fails if mapping two 'i' as OpenMP
requires 'i' to be the same in this case as 'x[:4]' is not present
on the device (cf. quote at the very end). All other fails are because
of something like that or because pointer attachment is not done when
already present.
While 'g' works (in GCC) I am not quite sure whether compilers are
required to do the pointer attachment in this case or not.
In any case, I also tried 'g' version (plus several others) with with
HPE/Cray's CCE, but usually it failed with "Variable not found in
present table" - i.e. with a runtime-library fail.
With AMD's LLVM compiler, there is no mapping fail but it still segfault
at runtime ("Write access to a read-only page") which likely implies that mapping
worked - but no pointer attachment was done. Neither seems to support
'iterator', i.e. only the manual version was tested ('g'). I wonder
whether it is possible to write this such that it works with all
compilers, whether the current version if valid OpenMP and, if not, if
one can write one - and whether OpenMP's ref_ptee, ref_ptr, ref_ptr_ptee
would help or not. It feels as if it should.
* * *
Back to GCC:
1. The 'this_ices' function gives an ICE
- this is unrelated to 'iterator'
- same error with gcc 13 to gcc 16 (while gcc 7 gives an error
that 'x' appears multiple times)
That's for 'map(to: x[0].ptr[:2], x[2].ptr[:2])'
Interestingly, it stops failing with 'x' added explicitly. (Otherwise,
'x' should appear implicitly.)
and might be related to an implicit 'map(tofrom: x)' as it works
in 'target data' but not in 'target'.
2. The 'g' function [gives an ICE with GCC 13 but]
works with GCC 14+, including mainline + your 6/11 patch.
:-)
3. For 'f' and 'h', with '-Wall' gcc prints several 'not initialized but used'
warnings for the iterator variable 'i'.
4. 'h' fail at runtime with: 'libgomp: pointer target not mapped for attach'
5. 'f' fails at runtime with:
'libgomp: Trying to map into device [0x7ffc96d5ece0..0x7ffc96d5ece0)
structure
element when other mapped elements from the same structure weren't mapped
together with it'
Item (3) seems to be a bogus iterator warning that should be silenced.
It feels as if (3) does the same statically as (4) + (5) and as if: if 'g' works
than also should 'f' and 'g' should. Given that I cannot get 'g' to work with
other compilers nor variants of 'g', I am not sure I can claim with confidence
that 'g' is valid.
The ICE (1) is clearly wrong by definition; it seems to be also valid (if
x[0:3]'
is previously mapped); evem is not iterator code.
* * *
I have a light confidence that the code is right, however, it is easy to get
lost with the fine print of map clauses. In any case, the following restriction
of OpenMP 5.0 (later version likely same/similar) applies and I think we
honor it? Or don't we?
"If multiple list items are explicitly mapped on the same construct and
have the same containing array or have base pointers that share original
storage, and if any of the list items do not have corresponding list items
that are present in the device data environment prior to a task encountering
the construct, then the list items must refer to the same array elements of
either the containing array or the implicit array of the base pointers."
* * *
If you can, I think it makes sense to have a look at all of the issues,
esp. the -Wall diagnostic. - As some issues seem to be unrelated, we could
also go ahead and file a problem report for them or fix them either later in
the iterator series. Besides -Wall, I think there was some oddness in the
Fortran part and 'malloc/free' could be combined.
In any case, it seems as if landing of a lightly updated version of this
patch makes more sense than to defer it until all issues are fixed.
Tobias
/* f - iterator dynamic bounds
h - iterator static bounds
g - explicit mapping */
struct array_ptr
{
int *ptr;
};
#if 1
void
f (struct array_ptr *x, int low, int high, int stride)
{
__builtin_printf ("f: host %p / %p\n", x[0].ptr, x[2].ptr);
#pragma omp target data map(to: x[:4])
#pragma omp target data map(iterator(i=low:high:stride), to: x[i].ptr[:2])
#pragma omp target map(to: x) map(iterator(i=low:high:stride), to: x[i].ptr)
{
__builtin_printf ("f: dev %p / %p\n", x[0].ptr, x[2].ptr);
if (x[0].ptr[0] != 1 || x[0].ptr[1] != 2) __builtin_abort ();
if (x[2].ptr[0] != 111 || x[2].ptr[1] != 222) __builtin_abort ();
}
}
void
h (struct array_ptr *x)
{
__builtin_printf ("h: host %p / %p\n", x[0].ptr, x[2].ptr);
#pragma omp target data map(to: x[:4])
#pragma omp target data map(iterator(i=0:3:2), to: x[i].ptr[:2])
#pragma omp target map(to: x) map(iterator(i=0:3:2), to: x[i].ptr)
{
__builtin_printf ("h: dev %p / %p\n", x[0].ptr, x[2].ptr);
if (x[0].ptr[0] != 1 || x[0].ptr[1] != 2) __builtin_abort ();
if (x[2].ptr[0] != 111 || x[2].ptr[1] != 222) __builtin_abort ();
}
}
#endif
void
g (struct array_ptr *x)
{
__builtin_printf ("g: host %p / %p\n", x[0].ptr, x[2].ptr);
/* Not quite clear whether that can be put into one directive,
but that's at least one method that works in GCC; with CCE,
all tried variants failed. */
#pragma omp target data map(to: x[:4])
#pragma omp target data map(to: x[0].ptr[:2], x[2].ptr[:2])
#pragma omp target map(to: x, x[0].ptr, x[2].ptr)
{
__builtin_printf ("g: dev %p / %p\n", x[0].ptr, x[2].ptr);
if (x[0].ptr[0] != 1 || x[0].ptr[1] != 2) __builtin_abort ();
if (x[2].ptr[0] != 111 || x[2].ptr[1] != 222) __builtin_abort ();
}
}
#if 0
void
this_ices (struct array_ptr *x)
{
// Works: 'target data'
// #pragma omp target data map(to: x[0].ptr[:2], x[2].ptr[:2])
// ICE with 'target':
// internal compiler error: base pointer cycle detected
#pragma omp target map(to: x[0].ptr[:2], x[2].ptr[:2])
;
// works after adding 'x' explicitly, cf. 'g' above.
}
#endif
int
main ()
{
int x1[] = {1, 2};
int x2[] = {11, 22};
int x3[] = {111, 222};
int x4[] = {1111, 2222};
struct array_ptr var[4];
var[0].ptr = x1;
var[1].ptr = x2;
var[2].ptr = x3;
var[3].ptr = x4;
g(var);
// h(var);
f(var, 0, 3, 2);
}
