diff mbox series

Host and offload targets have no common meaning of address spaces

Message ID 877db4djc1.fsf@euler.schwinge.homeip.net
State New
Headers show
Series Host and offload targets have no common meaning of address spaces | expand

Commit Message

Thomas Schwinge Jan. 13, 2022, 10:24 a.m. UTC 
Hi!

Jakub, I'd still like your comment on the two "should we" questions cited
below.

On 2021-08-24T13:43:38+0200, Richard Biener via Gcc-patches <gcc-patches@gcc.gnu.org> wrote:
> On Tue, Aug 24, 2021 at 12:23 PM Thomas Schwinge <thomas@codesourcery.com> wrote:
>> 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'
>> > [...]:
>> >
>> >     (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?

> I'd
> say I agree that any host address-space should go away when the corresponding
> data is offloaded

Pushed to master branch commit 9fcc3a1dd2372deea8856c55d25337b06e201203
"Host and offload targets have no common meaning of address spaces", see
attached.


>> 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
diff mbox series

Patch

From 9fcc3a1dd2372deea8856c55d25337b06e201203 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@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                       | 7 ++++++-
 libgomp/testsuite/libgomp.c/address-space-1.c | 4 ----
 3 files changed, 8 insertions(+), 5 deletions(-)

diff --git a/gcc/tree-streamer-in.c b/gcc/tree-streamer-in.c
index adaf624bda7..0d5108e36a0 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 8742bf09c6a..23d15a50670 100644
--- a/gcc/tree-streamer-out.c
+++ b/gcc/tree-streamer-out.c
@@ -119,7 +119,12 @@  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);
+      if (lto_stream_offload_p)
+	/* Host and offload targets have no common meaning of address
+	   spaces.  */
+	;
+      else
+	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.34.1