Re: Host and offload targets have no common meaning of address spaces (was: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref')

[Andrew Stubbs <ams@codesourcery.com>]

On 24/08/2021 12:43, Richard Biener via Gcc-patches wrote:
> On Tue, Aug 24, 2021 at 12:23 PM Thomas Schwinge
> <thomas@codesourcery.com> wrote:
>>
>> Hi!
>>
>> On 2021-08-19T22:13:56+0200, I wrote:
>>> On 2021-08-16T10:21:04+0200, Jakub Jelinek <jakub@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).)
> 
> Are address-spaces to be used in any way for OpenMP offload code?  That is,
> does the OpenMP standard talk about them and how to remap things?  I'd
> say I agree that any host address-space should go away when the corresponding
> data is offloaded and in case OpenMP allows to specify a target address-space
> that would need to be instantiated in a way so the LTO streaming knows about
> a mapping from the host to the target representation.

The new OpenMP 5 allocator features will permit allocations to different 
memories (we're planning an implementation soon). Whether that means a 
different address space may be target specific, but I would certainly 
expect that it could be. For AMD GCN there is a "flat" address space 
that covers most memories, but if you know what memory an address refers 
to then there's often a more efficient instruction you can use.

Certainly the numeric address space codes for the host system 
architecture have no meaning on the accelerator architecture.

Andrew