| Message ID | 87o89np2es.fsf@dem-tschwing-1.ger.mentorg.com |
|---|---|
| State | New |
| Headers | show |
| Series | Host and offload targets have no common meaning of address spaces (was: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref') | expand |
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. Richard. > > 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
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
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 <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).) > > > 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 <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 | 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