Hi!
On 2021-08-19T22:13:56+0200, I wrote:
On 2021-08-16T10:21:04+0200, Jakub Jelinek <ja...@redhat.com> wrote:
On Mon, Aug 16, 2021 at 10:08:42AM +0200, Thomas Schwinge wrote:
|> Concerning the current 'gcc/omp-low.c:omp_build_component_ref', for the
|> current set of offloading testcases, we never see a
|> '!ADDR_SPACE_GENERIC_P' there, so the address space handling doesn't seem
|> to be necessary there (but also won't do any harm: no-op).
Are you sure this can't trigger?
Say
extern int __seg_fs a;
void
foo (void)
{
#pragma omp parallel private (a)
a = 2;
}
That test case doesn't run into 'omp_build_component_ref' at all,
but [I've pushed an altered and extended variant that does],
"Add 'libgomp.c/address-space-1.c'".
In this case, 'omp_build_component_ref' called via host compilation
'pass_lower_omp', it's the 'field_type' that has 'address-space-1', not
'obj_type', so indeed Kwok's new code is a no-op:
(gdb) call debug_tree(field_type)
<pointer_type 0x7ffff7686b28
type <integer_type 0x7ffff7686498 int address-space-1 SI
I think keeping the qual addr space here is the wrong thing to do,
it should keep the other quals and clear the address space instead,
the whole struct is going to be in generic addres space, isn't it?
Correct for 'omp_build_component_ref' called via host compilation
'pass_lower_omp'
However, regarding the former comment -- shouldn't we force generic
address space for all 'tree' types read in via LTO streaming for
offloading compilation? I assume that (in the general case) address
spaces are never compatible between host and offloading compilation?
For [...] "Add 'libgomp.c/address-space-1.c'", propagating the
'__seg_fs' address space across the offloading boundary (assuming I did
interpret the dumps correctly) doesn't seem to cause any problems
As I found later, actually the 'address-space-1' per host '__seg_fs' does
cause the "Intel MIC (emulated) offloading execution failure"
mentioned/XFAILed for 'libgomp.c/address-space-1.c': SIGSEGV, like
(expected) for host execution. For GCN offloading target, it maps to
GCN 'ADDR_SPACE_FLAT' which apparently doesn't cause any ill effects (for
that simple test case). The nvptx offloading target doesn't consider
address spaces at all.
Is the attached "Host and offload targets have no common meaning of
address spaces" OK to push?
Then, is that the way to do this, or should we add in
'gcc/tree-streamer-out.c:pack_ts_base_value_fields':
if (lto_stream_offload_p)
gcc_assert (ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (expr)));
..., and elsewhere sanitize this for offloading compilation? Jakub's
suggestion above, regarding 'gcc/omp-low.c:omp_build_component_ref':
| I think keeping the qual addr space here is the wrong thing to do,
| it should keep the other quals and clear the address space instead
But it's not obvious to me that indeed this is the one place where this
would need to be done? (It ought to work for
'libgomp.c/address-space-1.c', and any other occurrences would run into
the 'assert', so that ought to be "fine", though?)
And, should we have a new hook
'void targetm.addr_space.validate (addr_space_t as)' (better name?),
called via 'gcc/emit-rtl.c:set_mem_attrs' (only? -- assuming this is the
appropriate canonic function where address space use is observed?), to
make sure that the requested 'as' is valid for the target?
'default_addr_space_validate' would refuse everything but
'ADDR_SPACE_GENERIC_P (as)'; this hook would need implementing for all
handful of targets making use of address spaces (supposedly matching the
logic how they call 'c_register_addr_space'?). (The closest existing
hook seems to be 'targetm.addr_space.diagnose_usage', only defined for
AVR, and called from "the front ends" (C only).)
