On Tue, Aug 24, 2021 at 12:23 PM Thomas Schwinge <tho...@codesourcery.com> 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).)
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