Hi!
Ping. Patch again attached for easy reference.
Plus, incrementally, the two "should we" questions cited below?
Grüße
Thomas
On 2021-08-24T12:23:07+0200, I wrote:
> 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).)
>
>
> Grüße
> Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634
München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas
Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht
München, HRB 106955
>From e01e06bd17bf2c7cb182d30bed02babc5edfa183 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Tue, 24 Aug 2021 11:14:10 +0200
Subject: [PATCH] Host and offload targets have no common meaning of address
spaces
gcc/
* tree-streamer-out.c (pack_ts_base_value_fields): Don't pack
'TYPE_ADDR_SPACE' for offloading.
* tree-streamer-in.c (unpack_ts_base_value_fields): Don't unpack
'TYPE_ADDR_SPACE' for offloading.
libgomp/
* testsuite/libgomp.c/address-space-1.c: Remove 'dg-xfail-run-if'
for 'offload_device_intel_mic'.
---
gcc/tree-streamer-in.c | 2 ++
gcc/tree-streamer-out.c | 4 +++-
libgomp/testsuite/libgomp.c/address-space-1.c | 4 ----
3 files changed, 5 insertions(+), 5 deletions(-)
diff --git a/gcc/tree-streamer-in.c b/gcc/tree-streamer-in.c
index e0522bf2ac1..acdc48ef09f 100644
--- a/gcc/tree-streamer-in.c
+++ b/gcc/tree-streamer-in.c
@@ -146,7 +146,9 @@ unpack_ts_base_value_fields (struct bitpack_d *bp, tree expr)
TYPE_REVERSE_STORAGE_ORDER (expr) = (unsigned) bp_unpack_value (bp, 1);
else
TYPE_SATURATING (expr) = (unsigned) bp_unpack_value (bp, 1);
+#ifndef ACCEL_COMPILER
TYPE_ADDR_SPACE (expr) = (unsigned) bp_unpack_value (bp, 8);
+#endif
}
else if (TREE_CODE (expr) == BIT_FIELD_REF || TREE_CODE (expr) == MEM_REF)
{
diff --git a/gcc/tree-streamer-out.c b/gcc/tree-streamer-out.c
index 855d1cd59b9..aac0b7ecf54 100644
--- a/gcc/tree-streamer-out.c
+++ b/gcc/tree-streamer-out.c
@@ -119,7 +119,9 @@ pack_ts_base_value_fields (struct bitpack_d *bp, tree expr)
bp_pack_value (bp, TYPE_REVERSE_STORAGE_ORDER (expr), 1);
else
bp_pack_value (bp, TYPE_SATURATING (expr), 1);
- bp_pack_value (bp, TYPE_ADDR_SPACE (expr), 8);
+ /* Host and offload targets have no common meaning of address spaces. */
+ if (!lto_stream_offload_p)
+ bp_pack_value (bp, TYPE_ADDR_SPACE (expr), 8);
}
else if (TREE_CODE (expr) == BIT_FIELD_REF || TREE_CODE (expr) == MEM_REF)
{
diff --git a/libgomp/testsuite/libgomp.c/address-space-1.c b/libgomp/testsuite/libgomp.c/address-space-1.c
index 6ad57deec42..39ff82c1429 100644
--- a/libgomp/testsuite/libgomp.c/address-space-1.c
+++ b/libgomp/testsuite/libgomp.c/address-space-1.c
@@ -3,10 +3,6 @@
/* { dg-do run { target i?86-*-* x86_64-*-* } } */
/* { dg-require-effective-target offload_device_nonshared_as } */
-/* With Intel MIC (emulated) offloading:
- offload error: process on the device 0 unexpectedly exited with code 0
- { dg-xfail-run-if TODO { offload_device_intel_mic } } */
-
#include <assert.h>
int __seg_fs a;
--
2.25.1
