* Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' [not found] ` <fe7bc7ae-6807-ff00-1b3a-e3c7ac41b723@mentor.com> @ 2021-08-09 14:16 ` Thomas Schwinge 2021-08-16 8:08 ` [ping] " Thomas Schwinge 2022-02-22 17:00 ` Get rid of 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref' (was: Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref') Thomas Schwinge 0 siblings, 2 replies; 15+ messages in thread From: Thomas Schwinge @ 2021-08-09 14:16 UTC (permalink / raw) To: Jakub Jelinek, gcc-patches; +Cc: Kwok Cheung Yeung, Julian Brown [-- Attachment #1: Type: text/plain, Size: 9016 bytes --] [from internal] Hi! This concerns a class of ICEs seen as of og10 branch with the "openacc: Middle-end worker-partitioning support" and "amdgcn: Enable OpenACC worker partitioning for AMD GCN" changes applied: On 2020-06-06T16:07:36+0100, Kwok Cheung Yeung <kwok_yeung@mentor.com> wrote: > On 01/06/2020 8:48 pm, Kwok Cheung Yeung wrote: >> On 21/05/2020 10:23 pm, Kwok Cheung Yeung wrote: >>> These all have the same failure mode: >>> >>> during RTL pass: expand >>> [...]/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90: In function 'MAIN__._omp_fn.1': >>> [...]/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90:86: internal compiler error: in convert_memory_address_addr_space_1, at explow.c:302 >>> 0xc29f20 convert_memory_address_addr_space_1(scalar_int_mode, rtx_def*, unsigned char, bool, bool) >>> [...]/gcc/explow.c:302 >>> 0xc29f57 convert_memory_address_addr_space(scalar_int_mode, rtx_def*, unsigned char) >>> [...]/gcc/explow.c:404 >>> [...] >>> This occurs if the -ftree-slp-vectorize flag is specified (default at -O3). >> The problematic bit of Gimple code is this: >> >> .oacc_worker_o.44._120 = gangs_min_472; >> .oacc_worker_o.44._122 = workers_min_473; >> .oacc_worker_o.44._124 = vectors_min_474; >> .oacc_worker_o.44._126 = gangs_max_475; >> .oacc_worker_o.44._128 = workers_max_476; >> .oacc_worker_o.44._130 = vectors_max_477; >> .oacc_worker_o.44._132 = 0; >> >> With SLP vectorization enabled, it becomes this: >> >> _40 = {gangs_min_472, workers_min_473, vectors_min_474, gangs_max_475}; >> ... >> MEM <vector(4) int> [(int *)&.oacc_worker_o.44] = _40; >> .oacc_worker_o.44._128 = workers_max_476; >> .oacc_worker_o.44._130 = vectors_max_477; >> .oacc_worker_o.44._132 = 0; >> >> The optimization is trying to transform 4 separate assignments into a single >> memory operation. The trouble is that &o.acc_worker_o is an SImode pointer in >> AS4 (LDS), while the memory expression appears to be in the default memory >> space. The 'to' expression of the assignment is: >> >> <mem_ref 0x7ffff74c61e0 >> type <vector_type 0x7ffff7470498 >> type <integer_type 0x7ffff73195e8 int public SI >> size <integer_cst 0x7ffff7318bb8 constant 32> >> unit-size <integer_cst 0x7ffff7318bd0 constant 4> >> align:32 warn_if_not_align:0 symtab:0 alias-set 1 canonical-type 0x7ffff73195e8 precision:32 min <integer_cst 0x7ffff7318b70 -2147483648> max <integer_cst 0x7ffff7318b88 2147483647> >> pointer_to_this <pointer_type 0x7ffff73209d8> reference_to_this <reference_type 0x7ffff73d9d20>> >> TI >> size <integer_cst 0x7ffff7318ca8 constant 128> >> unit-size <integer_cst 0x7ffff7318cc0 constant 16> >> align:128 warn_if_not_align:0 symtab:0 alias-set 1 structural-equality nunits:4 >> pointer_to_this <pointer_type 0x7ffff7470540>> >> >> arg:0 <addr_expr 0x7ffff74cdb80 >> type <pointer_type 0x7ffff73209d8 type <integer_type 0x7ffff73195e8 int> >> public unsigned DI >> size <integer_cst 0x7ffff7318978 constant 64> >> unit-size <integer_cst 0x7ffff7318990 constant 8> >> align:64 warn_if_not_align:0 symtab:0 alias-set 2 structural-equality> >> constant >> arg:0 <var_decl 0x7ffff7477f30 .oacc_worker_o.44 type <record_type 0x7ffff73eb888 .oacc_ws_data_s.21 address-space-4> >> addressable used static ignored BLK [...]/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90:86:0 >> >> size <integer_cst 0x7ffff746ce70 constant 224> >> unit-size <integer_cst 0x7ffff746ce40 constant 28> >> align:128 warn_if_not_align:0 >> (mem/c:BLK (symbol_ref:SI (".oacc_worker_o.44.14") [flags 0x2] <var_decl 0x7ffff7477f30 .oacc_worker_o.44>) [9 .oacc_worker_o.44+0 S28 A128 AS4])>> >> arg:1 <integer_cst 0x7ffff73ff078 type <pointer_type 0x7ffff73209d8> constant 0>> >> >> In convert_memory_address_addr_space_1: >> >> #ifndef POINTERS_EXTEND_UNSIGNED >> gcc_assert (GET_MODE (x) == to_mode || GET_MODE (x) == VOIDmode); >> return x; >> #else /* defined(POINTERS_EXTEND_UNSIGNED) */ >> >> POINTERS_EXTEND_UNSIGNED is not defined, so it hits the assert. The expected >> to_mode is DI_mode, but x is SI_mode, so the assert fires. > I now have a fix for this. > > > MEM <vector(4) int> [(int *)&.oacc_worker_o.44] = _40; > > The ICE occurs because the SLP vectorization pass creates the new statement > using the type of the expression '&.oacc_worker_o.44', which is a pointer to a > component ref in the default address space. The expand pass gets confused > because it is handed an SImode pointer (for LDS) when it is expecting a DImode > pointer (for flat/global space). > > The underlying problem is that although .oacc_worker_o is in the correct address > space, the component ref .oacc_worker_o is not. I fixed this by propagating the > address space of .oacc_worker_o when the component ref is created. > static tree > oacc_build_component_ref (tree obj, tree field) > { > - tree ret = build3 (COMPONENT_REF, TREE_TYPE (field), obj, field, NULL); > + tree field_type = TREE_TYPE (field); > + tree obj_type = TREE_TYPE (obj); > + if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) > + field_type = build_qualified_type > + (field_type, > + KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); > + > + tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); > if (TREE_THIS_VOLATILE (field)) > TREE_THIS_VOLATILE (ret) |= 1; > if (TREE_READONLY (field)) This code change has been included in the recent master branch commit e2a58ed6dc5293602d0d168475109caa81ad0f0d "openacc: Middle-end worker-partitioning support", which thus includes a 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref' that is slightly different from 'gcc/omp-low.c:omp_build_component_ref'. I'm confirming that with this reverted, we're seeing ICEs as follows: +FAIL: libgomp.oacc-fortran/gemm-2.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) +FAIL: libgomp.oacc-fortran/gemm-2.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) +FAIL: libgomp.oacc-fortran/gemm.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) +FAIL: libgomp.oacc-fortran/gemm.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) +FAIL: libgomp.oacc-fortran/optional-reduction.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) +FAIL: libgomp.oacc-fortran/optional-reduction.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) +FAIL: libgomp.oacc-fortran/private-variables.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) +FAIL: libgomp.oacc-fortran/private-variables.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) +FAIL: libgomp.oacc-fortran/reduction-1.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) +FAIL: libgomp.oacc-fortran/reduction-1.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) +FAIL: libgomp.oacc-fortran/reduction-5.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) +FAIL: libgomp.oacc-fortran/reduction-5.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) +FAIL: libgomp.oacc-fortran/reduction-6.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) +FAIL: libgomp.oacc-fortran/reduction-6.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) 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). Would it make sense to "Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref'", see attached? 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 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Re-unify-omp_build_component_ref-and-oacc_build_comp.patch --] [-- Type: text/x-diff, Size: 4595 bytes --] From caee66cf2abd0bea3ee99b460a108ae0d69d599f Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Fri, 30 Jul 2021 16:15:25 +0200 Subject: [PATCH] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' gcc/ * omp-general.c (omp_build_component_ref): New function, renamed/moved from... * omp-oacc-neuter-broadcast.cc (oacc_build_component_ref): ... here. (build_receiver_ref, build_sender_ref): Update. * omp-low.c (omp_build_component_ref): Remove function. * omp-general.h (omp_build_component_ref): Declare function. --- gcc/omp-general.c | 21 +++++++++++++++++++++ gcc/omp-general.h | 2 ++ gcc/omp-low.c | 15 --------------- gcc/omp-oacc-neuter-broadcast.cc | 26 ++------------------------ 4 files changed, 25 insertions(+), 39 deletions(-) diff --git a/gcc/omp-general.c b/gcc/omp-general.c index b46a537e281..67a0b752f62 100644 --- a/gcc/omp-general.c +++ b/gcc/omp-general.c @@ -2815,4 +2815,25 @@ oacc_get_ifn_dim_arg (const gimple *stmt) return (int) axis; } +/* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it + as appropriate. */ + +tree +omp_build_component_ref (tree obj, tree field) +{ + tree field_type = TREE_TYPE (field); + tree obj_type = TREE_TYPE (obj); + if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) + field_type + = build_qualified_type (field_type, + KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); + + tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); + if (TREE_THIS_VOLATILE (field)) + TREE_THIS_VOLATILE (ret) |= 1; + if (TREE_READONLY (field)) + TREE_READONLY (ret) |= 1; + return ret; +} + #include "gt-omp-general.h" diff --git a/gcc/omp-general.h b/gcc/omp-general.h index 5c3e0f0e205..6525175832c 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -145,4 +145,6 @@ get_openacc_privatization_dump_flags () return l_dump_flags; } +extern tree omp_build_component_ref (tree obj, tree field); + #endif /* GCC_OMP_GENERAL_H */ diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 926087da701..1640321c445 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -613,21 +613,6 @@ omp_copy_decl_1 (tree var, omp_context *ctx) return omp_copy_decl_2 (var, DECL_NAME (var), TREE_TYPE (var), ctx); } -/* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it - as appropriate. */ -/* See also 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref'. */ - -static tree -omp_build_component_ref (tree obj, tree field) -{ - tree ret = build3 (COMPONENT_REF, TREE_TYPE (field), obj, field, NULL); - if (TREE_THIS_VOLATILE (field)) - TREE_THIS_VOLATILE (ret) |= 1; - if (TREE_READONLY (field)) - TREE_READONLY (ret) |= 1; - return ret; -} - /* Build tree nodes to access the field for VAR on the receiver side. */ static tree diff --git a/gcc/omp-oacc-neuter-broadcast.cc b/gcc/omp-oacc-neuter-broadcast.cc index f8555380451..720cf74f12f 100644 --- a/gcc/omp-oacc-neuter-broadcast.cc +++ b/gcc/omp-oacc-neuter-broadcast.cc @@ -936,28 +936,6 @@ worker_single_simple (basic_block from, basic_block to, update_stmt (acc_bar); } -/* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it - as appropriate. */ -/* Adapted from 'gcc/omp-low.c:omp_build_component_ref'. */ - -static tree -oacc_build_component_ref (tree obj, tree field) -{ - tree field_type = TREE_TYPE (field); - tree obj_type = TREE_TYPE (obj); - if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) - field_type = build_qualified_type - (field_type, - KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); - - tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); - if (TREE_THIS_VOLATILE (field)) - TREE_THIS_VOLATILE (ret) |= 1; - if (TREE_READONLY (field)) - TREE_READONLY (ret) |= 1; - return ret; -} - static tree build_receiver_ref (tree record_type, tree var, tree receiver_decl) { @@ -965,7 +943,7 @@ build_receiver_ref (tree record_type, tree var, tree receiver_decl) tree x = build_simple_mem_ref (receiver_decl); tree field = *fields->get (var); TREE_THIS_NOTRAP (x) = 1; - x = oacc_build_component_ref (x, field); + x = omp_build_component_ref (x, field); return x; } @@ -974,7 +952,7 @@ build_sender_ref (tree record_type, tree var, tree sender_decl) { field_map_t *fields = *field_map->get (record_type); tree field = *fields->get (var); - return oacc_build_component_ref (sender_decl, field); + return omp_build_component_ref (sender_decl, field); } static int -- 2.30.2 ^ permalink raw reply [flat|nested] 15+ messages in thread
* [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' 2021-08-09 14:16 ` Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' Thomas Schwinge @ 2021-08-16 8:08 ` Thomas Schwinge 2021-08-16 8:21 ` Jakub Jelinek 2022-02-22 17:00 ` Get rid of 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref' (was: Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref') Thomas Schwinge 1 sibling, 1 reply; 15+ messages in thread From: Thomas Schwinge @ 2021-08-16 8:08 UTC (permalink / raw) To: Jakub Jelinek, gcc-patches; +Cc: Kwok Cheung Yeung, Julian Brown [-- Attachment #1: Type: text/plain, Size: 9272 bytes --] Hi! Ping. On 2021-08-09T16:16:51+0200, I wrote: > [from internal] > > > Hi! > > This concerns a class of ICEs seen as of og10 branch with the > "openacc: Middle-end worker-partitioning support" and "amdgcn: > Enable OpenACC worker partitioning for AMD GCN" changes applied: > > On 2020-06-06T16:07:36+0100, Kwok Cheung Yeung <kwok_yeung@mentor.com> wrote: >> On 01/06/2020 8:48 pm, Kwok Cheung Yeung wrote: >>> On 21/05/2020 10:23 pm, Kwok Cheung Yeung wrote: >>>> These all have the same failure mode: >>>> >>>> during RTL pass: expand >>>> [...]/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90: In function 'MAIN__._omp_fn.1': >>>> [...]/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90:86: internal compiler error: in convert_memory_address_addr_space_1, at explow.c:302 >>>> 0xc29f20 convert_memory_address_addr_space_1(scalar_int_mode, rtx_def*, unsigned char, bool, bool) >>>> [...]/gcc/explow.c:302 >>>> 0xc29f57 convert_memory_address_addr_space(scalar_int_mode, rtx_def*, unsigned char) >>>> [...]/gcc/explow.c:404 >>>> [...] > >>>> This occurs if the -ftree-slp-vectorize flag is specified (default at -O3). > >>> The problematic bit of Gimple code is this: >>> >>> .oacc_worker_o.44._120 = gangs_min_472; >>> .oacc_worker_o.44._122 = workers_min_473; >>> .oacc_worker_o.44._124 = vectors_min_474; >>> .oacc_worker_o.44._126 = gangs_max_475; >>> .oacc_worker_o.44._128 = workers_max_476; >>> .oacc_worker_o.44._130 = vectors_max_477; >>> .oacc_worker_o.44._132 = 0; >>> >>> With SLP vectorization enabled, it becomes this: >>> >>> _40 = {gangs_min_472, workers_min_473, vectors_min_474, gangs_max_475}; >>> ... >>> MEM <vector(4) int> [(int *)&.oacc_worker_o.44] = _40; >>> .oacc_worker_o.44._128 = workers_max_476; >>> .oacc_worker_o.44._130 = vectors_max_477; >>> .oacc_worker_o.44._132 = 0; >>> >>> The optimization is trying to transform 4 separate assignments into a single >>> memory operation. The trouble is that &o.acc_worker_o is an SImode pointer in >>> AS4 (LDS), while the memory expression appears to be in the default memory >>> space. The 'to' expression of the assignment is: >>> >>> <mem_ref 0x7ffff74c61e0 >>> type <vector_type 0x7ffff7470498 >>> type <integer_type 0x7ffff73195e8 int public SI >>> size <integer_cst 0x7ffff7318bb8 constant 32> >>> unit-size <integer_cst 0x7ffff7318bd0 constant 4> >>> align:32 warn_if_not_align:0 symtab:0 alias-set 1 canonical-type 0x7ffff73195e8 precision:32 min <integer_cst 0x7ffff7318b70 -2147483648> max <integer_cst 0x7ffff7318b88 2147483647> >>> pointer_to_this <pointer_type 0x7ffff73209d8> reference_to_this <reference_type 0x7ffff73d9d20>> >>> TI >>> size <integer_cst 0x7ffff7318ca8 constant 128> >>> unit-size <integer_cst 0x7ffff7318cc0 constant 16> >>> align:128 warn_if_not_align:0 symtab:0 alias-set 1 structural-equality nunits:4 >>> pointer_to_this <pointer_type 0x7ffff7470540>> >>> >>> arg:0 <addr_expr 0x7ffff74cdb80 >>> type <pointer_type 0x7ffff73209d8 type <integer_type 0x7ffff73195e8 int> >>> public unsigned DI >>> size <integer_cst 0x7ffff7318978 constant 64> >>> unit-size <integer_cst 0x7ffff7318990 constant 8> >>> align:64 warn_if_not_align:0 symtab:0 alias-set 2 structural-equality> >>> constant >>> arg:0 <var_decl 0x7ffff7477f30 .oacc_worker_o.44 type <record_type 0x7ffff73eb888 .oacc_ws_data_s.21 address-space-4> >>> addressable used static ignored BLK [...]/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90:86:0 >>> >>> size <integer_cst 0x7ffff746ce70 constant 224> >>> unit-size <integer_cst 0x7ffff746ce40 constant 28> >>> align:128 warn_if_not_align:0 >>> (mem/c:BLK (symbol_ref:SI (".oacc_worker_o.44.14") [flags 0x2] <var_decl 0x7ffff7477f30 .oacc_worker_o.44>) [9 .oacc_worker_o.44+0 S28 A128 AS4])>> >>> arg:1 <integer_cst 0x7ffff73ff078 type <pointer_type 0x7ffff73209d8> constant 0>> >>> >>> In convert_memory_address_addr_space_1: >>> >>> #ifndef POINTERS_EXTEND_UNSIGNED >>> gcc_assert (GET_MODE (x) == to_mode || GET_MODE (x) == VOIDmode); >>> return x; >>> #else /* defined(POINTERS_EXTEND_UNSIGNED) */ >>> >>> POINTERS_EXTEND_UNSIGNED is not defined, so it hits the assert. The expected >>> to_mode is DI_mode, but x is SI_mode, so the assert fires. > >> I now have a fix for this. >> >> > MEM <vector(4) int> [(int *)&.oacc_worker_o.44] = _40; >> >> The ICE occurs because the SLP vectorization pass creates the new statement >> using the type of the expression '&.oacc_worker_o.44', which is a pointer to a >> component ref in the default address space. The expand pass gets confused >> because it is handed an SImode pointer (for LDS) when it is expecting a DImode >> pointer (for flat/global space). >> >> The underlying problem is that although .oacc_worker_o is in the correct address >> space, the component ref .oacc_worker_o is not. I fixed this by propagating the >> address space of .oacc_worker_o when the component ref is created. > >> static tree >> oacc_build_component_ref (tree obj, tree field) >> { >> - tree ret = build3 (COMPONENT_REF, TREE_TYPE (field), obj, field, NULL); >> + tree field_type = TREE_TYPE (field); >> + tree obj_type = TREE_TYPE (obj); >> + if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) >> + field_type = build_qualified_type >> + (field_type, >> + KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); >> + >> + tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); >> if (TREE_THIS_VOLATILE (field)) >> TREE_THIS_VOLATILE (ret) |= 1; >> if (TREE_READONLY (field)) > > This code change has been included in the recent master branch commit > e2a58ed6dc5293602d0d168475109caa81ad0f0d "openacc: Middle-end > worker-partitioning support", which thus includes a > 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref' that is > slightly different from 'gcc/omp-low.c:omp_build_component_ref'. > > I'm confirming that with this reverted, we're seeing ICEs as follows: > > +FAIL: libgomp.oacc-fortran/gemm-2.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/gemm-2.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/gemm.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/gemm.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/optional-reduction.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/optional-reduction.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/private-variables.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/private-variables.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-1.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-1.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-5.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-5.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-6.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-6.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > 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). > > Would it make sense to "Re-unify 'omp_build_component_ref' and > 'oacc_build_component_ref'", see attached? > > > 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 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Re-unify-omp_build_component_ref-and-oacc_build_comp.patch --] [-- Type: text/x-diff, Size: 4595 bytes --] From caee66cf2abd0bea3ee99b460a108ae0d69d599f Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Fri, 30 Jul 2021 16:15:25 +0200 Subject: [PATCH] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' gcc/ * omp-general.c (omp_build_component_ref): New function, renamed/moved from... * omp-oacc-neuter-broadcast.cc (oacc_build_component_ref): ... here. (build_receiver_ref, build_sender_ref): Update. * omp-low.c (omp_build_component_ref): Remove function. * omp-general.h (omp_build_component_ref): Declare function. --- gcc/omp-general.c | 21 +++++++++++++++++++++ gcc/omp-general.h | 2 ++ gcc/omp-low.c | 15 --------------- gcc/omp-oacc-neuter-broadcast.cc | 26 ++------------------------ 4 files changed, 25 insertions(+), 39 deletions(-) diff --git a/gcc/omp-general.c b/gcc/omp-general.c index b46a537e281..67a0b752f62 100644 --- a/gcc/omp-general.c +++ b/gcc/omp-general.c @@ -2815,4 +2815,25 @@ oacc_get_ifn_dim_arg (const gimple *stmt) return (int) axis; } +/* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it + as appropriate. */ + +tree +omp_build_component_ref (tree obj, tree field) +{ + tree field_type = TREE_TYPE (field); + tree obj_type = TREE_TYPE (obj); + if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) + field_type + = build_qualified_type (field_type, + KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); + + tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); + if (TREE_THIS_VOLATILE (field)) + TREE_THIS_VOLATILE (ret) |= 1; + if (TREE_READONLY (field)) + TREE_READONLY (ret) |= 1; + return ret; +} + #include "gt-omp-general.h" diff --git a/gcc/omp-general.h b/gcc/omp-general.h index 5c3e0f0e205..6525175832c 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -145,4 +145,6 @@ get_openacc_privatization_dump_flags () return l_dump_flags; } +extern tree omp_build_component_ref (tree obj, tree field); + #endif /* GCC_OMP_GENERAL_H */ diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 926087da701..1640321c445 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -613,21 +613,6 @@ omp_copy_decl_1 (tree var, omp_context *ctx) return omp_copy_decl_2 (var, DECL_NAME (var), TREE_TYPE (var), ctx); } -/* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it - as appropriate. */ -/* See also 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref'. */ - -static tree -omp_build_component_ref (tree obj, tree field) -{ - tree ret = build3 (COMPONENT_REF, TREE_TYPE (field), obj, field, NULL); - if (TREE_THIS_VOLATILE (field)) - TREE_THIS_VOLATILE (ret) |= 1; - if (TREE_READONLY (field)) - TREE_READONLY (ret) |= 1; - return ret; -} - /* Build tree nodes to access the field for VAR on the receiver side. */ static tree diff --git a/gcc/omp-oacc-neuter-broadcast.cc b/gcc/omp-oacc-neuter-broadcast.cc index f8555380451..720cf74f12f 100644 --- a/gcc/omp-oacc-neuter-broadcast.cc +++ b/gcc/omp-oacc-neuter-broadcast.cc @@ -936,28 +936,6 @@ worker_single_simple (basic_block from, basic_block to, update_stmt (acc_bar); } -/* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it - as appropriate. */ -/* Adapted from 'gcc/omp-low.c:omp_build_component_ref'. */ - -static tree -oacc_build_component_ref (tree obj, tree field) -{ - tree field_type = TREE_TYPE (field); - tree obj_type = TREE_TYPE (obj); - if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) - field_type = build_qualified_type - (field_type, - KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); - - tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); - if (TREE_THIS_VOLATILE (field)) - TREE_THIS_VOLATILE (ret) |= 1; - if (TREE_READONLY (field)) - TREE_READONLY (ret) |= 1; - return ret; -} - static tree build_receiver_ref (tree record_type, tree var, tree receiver_decl) { @@ -965,7 +943,7 @@ build_receiver_ref (tree record_type, tree var, tree receiver_decl) tree x = build_simple_mem_ref (receiver_decl); tree field = *fields->get (var); TREE_THIS_NOTRAP (x) = 1; - x = oacc_build_component_ref (x, field); + x = omp_build_component_ref (x, field); return x; } @@ -974,7 +952,7 @@ build_sender_ref (tree record_type, tree var, tree sender_decl) { field_map_t *fields = *field_map->get (record_type); tree field = *fields->get (var); - return oacc_build_component_ref (sender_decl, field); + return omp_build_component_ref (sender_decl, field); } static int -- 2.30.2 ^ permalink raw reply [flat|nested] 15+ messages in thread
* Re: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' 2021-08-16 8:08 ` [ping] " Thomas Schwinge @ 2021-08-16 8:21 ` Jakub Jelinek 2021-08-19 20:13 ` Thomas Schwinge 0 siblings, 1 reply; 15+ messages in thread From: Jakub Jelinek @ 2021-08-16 8:21 UTC (permalink / raw) To: Thomas Schwinge; +Cc: gcc-patches, Kwok Cheung Yeung, Julian Brown On Mon, Aug 16, 2021 at 10:08:42AM +0200, Thomas Schwinge wrote: > --- a/gcc/omp-general.c > +++ b/gcc/omp-general.c > @@ -2815,4 +2815,25 @@ oacc_get_ifn_dim_arg (const gimple *stmt) > return (int) axis; > } > > +/* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it > + as appropriate. */ > + > +tree > +omp_build_component_ref (tree obj, tree field) > +{ > + tree field_type = TREE_TYPE (field); > + tree obj_type = TREE_TYPE (obj); > + if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) > + field_type > + = build_qualified_type (field_type, > + KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); Are you sure this can't trigger? Say extern int __seg_fs a; void foo (void) { #pragma omp parallel private (a) a = 2; } 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? > + > + tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); > + if (TREE_THIS_VOLATILE (field)) > + TREE_THIS_VOLATILE (ret) |= 1; > + if (TREE_READONLY (field)) > + TREE_READONLY (ret) |= 1; When touching these two, shouldn't it be better written as = 1; instead of |= 1; ? For a bitfield... Jakub ^ permalink raw reply [flat|nested] 15+ messages in thread
* Re: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' 2021-08-16 8:21 ` Jakub Jelinek @ 2021-08-19 20:13 ` Thomas Schwinge 2021-08-20 7:51 ` Richard Biener ` (2 more replies) 0 siblings, 3 replies; 15+ messages in thread From: Thomas Schwinge @ 2021-08-19 20:13 UTC (permalink / raw) To: Jakub Jelinek, Richard Biener Cc: gcc-patches, Kwok Cheung Yeung, Julian Brown [-- Attachment #1: Type: text/plain, Size: 10556 bytes --] Hi! Richard, maybe you have an opinion here, in particular about my "SLP vectorizer" comment below? Please see <http://mid.mail-archive.com/87r1f2puss.fsf@euler.schwinge.homeip.net> for the full context. 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: >> /* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it >> as appropriate. */ >> >> tree >> omp_build_component_ref (tree obj, tree field) >> { >> + tree field_type = TREE_TYPE (field); >> + tree obj_type = TREE_TYPE (obj); >> + if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) >> + field_type >> + = build_qualified_type (field_type, >> + KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); (For later reference: "Kwok's new code" here is to propagate to 'field_type' any non-generic address space of 'obj_type'.) |> 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'm attaching an altered and extended variant that does, "Add 'libgomp.c/address-space-1.c'". OK to push to master branch? 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 size <integer_cst 0x7ffff7540f30 constant 32> unit-size <integer_cst 0x7ffff7540f48 constant 4> align:32 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686498 precision:32 min <integer_cst 0x7ffff7540ee8 -2147483648> max <integer_cst 0x7ffff7540f00 2147483647> pointer_to_this <pointer_type 0x7ffff7686b28>> unsigned DI size <integer_cst 0x7ffff7540cf0 type <integer_type 0x7ffff75590a8 bitsizetype> constant 64> unit-size <integer_cst 0x7ffff7540d08 type <integer_type 0x7ffff7559000 sizetype> constant 8> align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686b28> (gdb) call debug_tree(obj_type) <record_type 0x7ffff7686bd0 .omp_data_t.0 readonly DI size <integer_cst 0x7ffff7540cf0 type <integer_type 0x7ffff75590a8 bitsizetype> constant 64> unit-size <integer_cst 0x7ffff7540d08 type <integer_type 0x7ffff7559000 sizetype> constant 8> align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686bd0 fields <field_decl 0x7ffff7568428 a type <pointer_type 0x7ffff7686b28 type <integer_type 0x7ffff7686498 int address-space-1> unsigned DI size <integer_cst 0x7ffff7540cf0 64> unit-size <integer_cst 0x7ffff7540d08 8> align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686b28> unsigned DI /home/thomas/shared/gcc/omp/as.c:4:14 size <integer_cst 0x7ffff7540cf0 64> unit-size <integer_cst 0x7ffff7540d08 8> align:64 warn_if_not_align:0 offset_align 128 offset <integer_cst 0x7ffff7540d20 constant 0> bit-offset <integer_cst 0x7ffff7540d68 constant 0> context <record_type 0x7ffff7686540 .omp_data_t.0>> reference_to_this <reference_type 0x7ffff7686c78>> The case that Kwok's new code handles, however, is when 'obj_type' has a non-generic address space, and then propagates that one to 'field_type'. For a similar OpenACC example, 'omp_build_component_ref' called via GCN offloading compilation 'pass_omp_oacc_neuter_broadcast', we've got without Kwok's new code: (gdb) call debug_tree(field_type) <boolean_type 0x7ffff7550b28 bool public unsigned QI size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7550b28 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> (gdb) call debug_tree(obj_type) <record_type 0x7ffff7631000 .oacc_ws_data_s.0 address-space-4 QI size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7631000 fields <field_decl 0x7ffff762e260 _52 type <boolean_type 0x7ffff7550b28 bool public unsigned QI size <integer_cst 0x7ffff754fa80 8> unit-size <integer_cst 0x7ffff754fa98 1> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7550b28 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> unsigned QI <built-in>:0:0 size <integer_cst 0x7ffff754fa80 8> unit-size <integer_cst 0x7ffff754fa98 1> align:8 warn_if_not_align:0 offset_align 64 offset <integer_cst 0x7ffff754f9c0 constant 0> bit-offset <integer_cst 0x7ffff754fa08 constant 0> context <record_type 0x7ffff7631000 .oacc_ws_data_s.0 address-space-4>> pointer_to_this <pointer_type 0x7ffff7631498>> ..., and with Kwok's new code the 'address-space-4' of 'obj_type' is propagated to 'field_type': (gdb) call debug_tree(field_type) <boolean_type 0x7ffff7631540 bool address-space-4 unsigned QI size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7631540 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> I'm not familiar enough with these bits to tell whether Kwok's new code is the right solution to this problem -- or if, for example, the problem is rather in the SLP vectorizer, where the ICE seems to ultimately emerge? Without (ICEs later) vs. with (works) Kwok's new code, we see the 'a.xamdgcn-amdhsa.mkoffload.175t.slp1' dump change as follows (word-diff, only additional '<address-space-4>', occasionally): [...] {+<address-space-4>+} vector(2) long int * vectp.58; {+<address-space-4>+} vector(2) long int * vectp_.oacc_worker_o.57; {+<address-space-4>+} vector(2) int * vectp.56; {+<address-space-4>+} vector(2) int * vectp_.oacc_worker_o.55; [...] {+<address-space-4>+} long int * _104; [...] {+<address-space-4>+} long int * _108; [...] <address-space-4> void * _350; [...] _350 = __builtin_gcn_single_copy_start (&.oacc_worker_o.6); [...] MEM <{+<address-space-4>+} vector(2) long int> [(long int *)&.oacc_worker_o.6] = _101; _108 = &.oacc_worker_o.6._22 + 16; MEM <{+<address-space-4>+} vector(2) long int> [(long int *)_108] = _100; _104 = &.oacc_worker_o.6._22 + 32; [...] For example, with Kwok's new code, '_108' ('<address-space-4> long int *') is cast into '(long int *)' -- presumably synthesized in the SLP vectorizer? Is that correct or shouldn't that cast also include '<address-space-4>'? I see a similar issue has been fixed a while ago: r245772 (Git commit c7d97b2846c5647a81548caa3264d77c0a595010) for PR79723 "Another case of dropped gs: prefix", changing 'gcc/tree-vect-stmts.c:get_vectype_for_scalar_type_and_size' as follows: + /* Re-attach the address-space qualifier if we canonicalized the scalar + type. */ + if (TYPE_ADDR_SPACE (orig_scalar_type) != TYPE_ADDR_SPACE (vectype)) + return build_qualified_type + (vectype, KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (orig_scalar_type))); + return vectype; (It looks a bit like the address space handling is quite fragile in GCC's 'tree' types/interfaces? Do we have ideas about how to make that more robust, less "bolt-on"?) I did add a few 'assert's for non-generic address space to 'gcc/tree-vect*', but have not yet located where things may be going wrong. > 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', but in the case of 'omp_build_component_ref' called via GCN offloading compilation 'pass_omp_oacc_neuter_broadcast', 'obj_type' has a non-generic address space. 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 the attached "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, but maybe it's problematic for other cases? (This is, however, a separate issue from what I'm discussing here.) >> + tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); >> + if (TREE_THIS_VOLATILE (field)) >> + TREE_THIS_VOLATILE (ret) |= 1; >> + if (TREE_READONLY (field)) >> + TREE_READONLY (ret) |= 1; > > When touching these two, shouldn't it be better written as > = 1; instead of |= 1; ? For a bitfield... Yes, that was just copied from the original 'gcc/omp-general.c:omp_build_component_ref' -- but happy to simplify that, of course. 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 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Add-libgomp.c-address-space-1.c.patch --] [-- Type: text/x-diff, Size: 1250 bytes --] From 6799543a4ff56360f3002931de1766a4448062a0 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Thu, 19 Aug 2021 15:14:51 +0200 Subject: [PATCH] Add 'libgomp.c/address-space-1.c' libgomp/ * testsuite/libgomp.c/address-space-1.c: New file. Co-authored-by: Jakub Jelinek <jakub@redhat.com> --- libgomp/testsuite/libgomp.c/address-space-1.c | 24 +++++++++++++++++++ 1 file changed, 24 insertions(+) create mode 100644 libgomp/testsuite/libgomp.c/address-space-1.c diff --git a/libgomp/testsuite/libgomp.c/address-space-1.c b/libgomp/testsuite/libgomp.c/address-space-1.c new file mode 100644 index 00000000000..90244db03b1 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/address-space-1.c @@ -0,0 +1,24 @@ +/* Verify OMP instances of variables with address space. */ + +/* { dg-do run { target i?86-*-* x86_64-*-* } } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +#include <assert.h> + +int __seg_fs a; + +int +main (void) +{ + // a = 123; // SIGSEGV + int b; +#pragma omp target map(alloc: a) map(from: b) + { + a = 321; // no SIGSEGV (given 'offload_device_nonshared_as') + asm volatile ("" : : : "memory"); + b = a; + } + assert (b == 321); + + return 0; +} -- 2.30.2 ^ permalink raw reply [flat|nested] 15+ messages in thread
* Re: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' 2021-08-19 20:13 ` Thomas Schwinge @ 2021-08-20 7:51 ` Richard Biener 2021-08-23 14:30 ` Thomas Schwinge 2021-08-20 14:49 ` Jakub Jelinek 2021-08-24 10:23 ` Host and offload targets have no common meaning of address spaces " Thomas Schwinge 2 siblings, 1 reply; 15+ messages in thread From: Richard Biener @ 2021-08-20 7:51 UTC (permalink / raw) To: Thomas Schwinge Cc: Jakub Jelinek, GCC Patches, Kwok Cheung Yeung, Julian Brown On Thu, Aug 19, 2021 at 10:14 PM Thomas Schwinge <thomas@codesourcery.com> wrote: > > Hi! > > Richard, maybe you have an opinion here, in particular about my > "SLP vectorizer" comment below? Please see > <http://mid.mail-archive.com/87r1f2puss.fsf@euler.schwinge.homeip.net> > for the full context. > > 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: > >> /* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it > >> as appropriate. */ > >> > >> tree > >> omp_build_component_ref (tree obj, tree field) > >> { > >> + tree field_type = TREE_TYPE (field); > >> + tree obj_type = TREE_TYPE (obj); > >> + if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) > >> + field_type > >> + = build_qualified_type (field_type, > >> + KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); > > (For later reference: "Kwok's new code" here is to propagate to > 'field_type' any non-generic address space of 'obj_type'.) > > |> 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'm attaching an altered and extended variant that does, > "Add 'libgomp.c/address-space-1.c'". OK to push to master branch? > > 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 > size <integer_cst 0x7ffff7540f30 constant 32> > unit-size <integer_cst 0x7ffff7540f48 constant 4> > align:32 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686498 precision:32 min <integer_cst 0x7ffff7540ee8 -2147483648> max <integer_cst 0x7ffff7540f00 2147483647> > pointer_to_this <pointer_type 0x7ffff7686b28>> > unsigned DI > size <integer_cst 0x7ffff7540cf0 type <integer_type 0x7ffff75590a8 bitsizetype> constant 64> > unit-size <integer_cst 0x7ffff7540d08 type <integer_type 0x7ffff7559000 sizetype> constant 8> > align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686b28> > > (gdb) call debug_tree(obj_type) > <record_type 0x7ffff7686bd0 .omp_data_t.0 readonly DI > size <integer_cst 0x7ffff7540cf0 type <integer_type 0x7ffff75590a8 bitsizetype> constant 64> > unit-size <integer_cst 0x7ffff7540d08 type <integer_type 0x7ffff7559000 sizetype> constant 8> > align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686bd0 > fields <field_decl 0x7ffff7568428 a > type <pointer_type 0x7ffff7686b28 type <integer_type 0x7ffff7686498 int address-space-1> > unsigned DI size <integer_cst 0x7ffff7540cf0 64> unit-size <integer_cst 0x7ffff7540d08 8> > align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686b28> > unsigned DI /home/thomas/shared/gcc/omp/as.c:4:14 size <integer_cst 0x7ffff7540cf0 64> unit-size <integer_cst 0x7ffff7540d08 8> > align:64 warn_if_not_align:0 offset_align 128 > offset <integer_cst 0x7ffff7540d20 constant 0> > bit-offset <integer_cst 0x7ffff7540d68 constant 0> context <record_type 0x7ffff7686540 .omp_data_t.0>> reference_to_this <reference_type 0x7ffff7686c78>> > > The case that Kwok's new code handles, however, is when 'obj_type' has a > non-generic address space, and then propagates that one to 'field_type'. > > For a similar OpenACC example, 'omp_build_component_ref' called via GCN > offloading compilation 'pass_omp_oacc_neuter_broadcast', we've got > without Kwok's new code: > > (gdb) call debug_tree(field_type) > <boolean_type 0x7ffff7550b28 bool public unsigned QI > size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> > unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> > align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7550b28 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> > > (gdb) call debug_tree(obj_type) > <record_type 0x7ffff7631000 .oacc_ws_data_s.0 address-space-4 QI > size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> > unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> > align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7631000 > fields <field_decl 0x7ffff762e260 _52 > type <boolean_type 0x7ffff7550b28 bool public unsigned QI size <integer_cst 0x7ffff754fa80 8> unit-size <integer_cst 0x7ffff754fa98 1> > align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7550b28 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> > unsigned QI <built-in>:0:0 size <integer_cst 0x7ffff754fa80 8> unit-size <integer_cst 0x7ffff754fa98 1> > align:8 warn_if_not_align:0 offset_align 64 > offset <integer_cst 0x7ffff754f9c0 constant 0> > bit-offset <integer_cst 0x7ffff754fa08 constant 0> context <record_type 0x7ffff7631000 .oacc_ws_data_s.0 address-space-4>> > pointer_to_this <pointer_type 0x7ffff7631498>> > > ..., and with Kwok's new code the 'address-space-4' of 'obj_type' is > propagated to 'field_type': > > (gdb) call debug_tree(field_type) > <boolean_type 0x7ffff7631540 bool address-space-4 unsigned QI > size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> > unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> > align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7631540 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> > > I'm not familiar enough with these bits to tell whether Kwok's new code > is the right solution to this problem -- or if, for example, the problem > is rather in the SLP vectorizer, where the ICE seems to ultimately > emerge? > > Without (ICEs later) vs. with (works) Kwok's new code, we see the > 'a.xamdgcn-amdhsa.mkoffload.175t.slp1' dump change as follows (word-diff, > only additional '<address-space-4>', occasionally): > > [...] > {+<address-space-4>+} vector(2) long int * vectp.58; > {+<address-space-4>+} vector(2) long int * vectp_.oacc_worker_o.57; > {+<address-space-4>+} vector(2) int * vectp.56; > {+<address-space-4>+} vector(2) int * vectp_.oacc_worker_o.55; > [...] > {+<address-space-4>+} long int * _104; > [...] > {+<address-space-4>+} long int * _108; > [...] > <address-space-4> void * _350; > [...] > _350 = __builtin_gcn_single_copy_start (&.oacc_worker_o.6); > [...] > MEM <{+<address-space-4>+} vector(2) long int> [(long int *)&.oacc_worker_o.6] = _101; > _108 = &.oacc_worker_o.6._22 + 16; > MEM <{+<address-space-4>+} vector(2) long int> [(long int *)_108] = _100; > _104 = &.oacc_worker_o.6._22 + 32; > [...] > > For example, with Kwok's new code, '_108' ('<address-space-4> long int *') > is cast into '(long int *)' -- presumably synthesized in the SLP > vectorizer? Is that correct or shouldn't that cast also include > '<address-space-4>'? > > I see a similar issue has been fixed a while ago: r245772 (Git commit > c7d97b2846c5647a81548caa3264d77c0a595010) for PR79723 > "Another case of dropped gs: prefix", changing > 'gcc/tree-vect-stmts.c:get_vectype_for_scalar_type_and_size' as follows: > > + /* Re-attach the address-space qualifier if we canonicalized the scalar > + type. */ > + if (TYPE_ADDR_SPACE (orig_scalar_type) != TYPE_ADDR_SPACE (vectype)) > + return build_qualified_type > + (vectype, KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (orig_scalar_type))); > + > return vectype; > > (It looks a bit like the address space handling is quite fragile in GCC's > 'tree' types/interfaces? Do we have ideas about how to make that more > robust, less "bolt-on"?) If in doubt always look at what RTL expansion does - it looks like set_mem_attributes expects the address-space qualifier to be present on the type or in case it is passed an object, on the type of the base, or in case of a dereference, on the pointed-to type of the pointer (and yes, that does look somewhat fragile). So it looks like the patch you refer to shouldn't fix anything and > + /* Re-attach the address-space qualifier if we canonicalized the scalar > + type. */ > + if (TYPE_ADDR_SPACE (orig_scalar_type) != TYPE_ADDR_SPACE (vectype)) > + return build_qualified_type > + (vectype, KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (orig_scalar_type))); looks incomplete. What you'd need to look for is MEM_REFs built by the vectorizer and the address-space information on the pointers, like generated from vect_create_data_ref_ptr. It might also be that data-ref analysis / SCEV looks through address-space qualifier changing casts and thus we pick up the wrong address-space in the end. What's the testcase that ICEs on trunk? > I did add a few 'assert's for non-generic address space to > 'gcc/tree-vect*', but have not yet located where things may be going > wrong. > > > > 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', but in the case of 'omp_build_component_ref' called via > GCN offloading compilation 'pass_omp_oacc_neuter_broadcast', 'obj_type' > has a non-generic address space. > > 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 the attached "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, but > maybe it's problematic for other cases? (This is, however, a separate > issue from what I'm discussing here.) > > > >> + tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); > >> + if (TREE_THIS_VOLATILE (field)) > >> + TREE_THIS_VOLATILE (ret) |= 1; > >> + if (TREE_READONLY (field)) > >> + TREE_READONLY (ret) |= 1; > > > > When touching these two, shouldn't it be better written as > > = 1; instead of |= 1; ? For a bitfield... > > Yes, that was just copied from the original > 'gcc/omp-general.c:omp_build_component_ref' -- but happy to simplify > that, of course. > > > 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 ^ permalink raw reply [flat|nested] 15+ messages in thread
* Re: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' 2021-08-20 7:51 ` Richard Biener @ 2021-08-23 14:30 ` Thomas Schwinge 2021-08-24 7:43 ` Richard Biener 0 siblings, 1 reply; 15+ messages in thread From: Thomas Schwinge @ 2021-08-23 14:30 UTC (permalink / raw) To: Richard Biener Cc: Jakub Jelinek, gcc-patches, Kwok Cheung Yeung, Julian Brown, Andrew Stubbs [-- Attachment #1: Type: text/plain, Size: 20385 bytes --] Hi! On 2021-08-20T09:51:36+0200, Richard Biener <richard.guenther@gmail.com> wrote: > On Thu, Aug 19, 2021 at 10:14 PM Thomas Schwinge > <thomas@codesourcery.com> wrote: >> Richard, maybe you have an opinion here, in particular about my >> "SLP vectorizer" comment below? Please see >> <http://mid.mail-archive.com/87r1f2puss.fsf@euler.schwinge.homeip.net> >> for the full context. >> >> 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: >> >> /* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it >> >> as appropriate. */ >> >> >> >> tree >> >> omp_build_component_ref (tree obj, tree field) >> >> { >> >> + tree field_type = TREE_TYPE (field); >> >> + tree obj_type = TREE_TYPE (obj); >> >> + if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) >> >> + field_type >> >> + = build_qualified_type (field_type, >> >> + KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); >> >> (For later reference: "Kwok's new code" here is to propagate to >> 'field_type' any non-generic address space of 'obj_type'.) >> >> |> 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'm attaching an altered and extended variant that does, >> "Add 'libgomp.c/address-space-1.c'". OK to push to master branch? >> >> 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 >> size <integer_cst 0x7ffff7540f30 constant 32> >> unit-size <integer_cst 0x7ffff7540f48 constant 4> >> align:32 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686498 precision:32 min <integer_cst 0x7ffff7540ee8 -2147483648> max <integer_cst 0x7ffff7540f00 2147483647> >> pointer_to_this <pointer_type 0x7ffff7686b28>> >> unsigned DI >> size <integer_cst 0x7ffff7540cf0 type <integer_type 0x7ffff75590a8 bitsizetype> constant 64> >> unit-size <integer_cst 0x7ffff7540d08 type <integer_type 0x7ffff7559000 sizetype> constant 8> >> align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686b28> >> >> (gdb) call debug_tree(obj_type) >> <record_type 0x7ffff7686bd0 .omp_data_t.0 readonly DI >> size <integer_cst 0x7ffff7540cf0 type <integer_type 0x7ffff75590a8 bitsizetype> constant 64> >> unit-size <integer_cst 0x7ffff7540d08 type <integer_type 0x7ffff7559000 sizetype> constant 8> >> align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686bd0 >> fields <field_decl 0x7ffff7568428 a >> type <pointer_type 0x7ffff7686b28 type <integer_type 0x7ffff7686498 int address-space-1> >> unsigned DI size <integer_cst 0x7ffff7540cf0 64> unit-size <integer_cst 0x7ffff7540d08 8> >> align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686b28> >> unsigned DI /home/thomas/shared/gcc/omp/as.c:4:14 size <integer_cst 0x7ffff7540cf0 64> unit-size <integer_cst 0x7ffff7540d08 8> >> align:64 warn_if_not_align:0 offset_align 128 >> offset <integer_cst 0x7ffff7540d20 constant 0> >> bit-offset <integer_cst 0x7ffff7540d68 constant 0> context <record_type 0x7ffff7686540 .omp_data_t.0>> reference_to_this <reference_type 0x7ffff7686c78>> >> >> The case that Kwok's new code handles, however, is when 'obj_type' has a >> non-generic address space, and then propagates that one to 'field_type'. >> >> For a similar OpenACC example, 'omp_build_component_ref' called via GCN >> offloading compilation 'pass_omp_oacc_neuter_broadcast', we've got >> without Kwok's new code: >> >> (gdb) call debug_tree(field_type) >> <boolean_type 0x7ffff7550b28 bool public unsigned QI >> size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> >> unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> >> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7550b28 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> >> >> (gdb) call debug_tree(obj_type) >> <record_type 0x7ffff7631000 .oacc_ws_data_s.0 address-space-4 QI >> size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> >> unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> >> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7631000 >> fields <field_decl 0x7ffff762e260 _52 >> type <boolean_type 0x7ffff7550b28 bool public unsigned QI size <integer_cst 0x7ffff754fa80 8> unit-size <integer_cst 0x7ffff754fa98 1> >> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7550b28 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> >> unsigned QI <built-in>:0:0 size <integer_cst 0x7ffff754fa80 8> unit-size <integer_cst 0x7ffff754fa98 1> >> align:8 warn_if_not_align:0 offset_align 64 >> offset <integer_cst 0x7ffff754f9c0 constant 0> >> bit-offset <integer_cst 0x7ffff754fa08 constant 0> context <record_type 0x7ffff7631000 .oacc_ws_data_s.0 address-space-4>> >> pointer_to_this <pointer_type 0x7ffff7631498>> >> >> ..., and with Kwok's new code the 'address-space-4' of 'obj_type' is >> propagated to 'field_type': >> >> (gdb) call debug_tree(field_type) >> <boolean_type 0x7ffff7631540 bool address-space-4 unsigned QI >> size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> >> unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> >> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7631540 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> >> >> I'm not familiar enough with these bits to tell whether Kwok's new code >> is the right solution to this problem -- or if, for example, the problem >> is rather in the SLP vectorizer, where the ICE seems to ultimately >> emerge? >> >> Without (ICEs later) vs. with (works) Kwok's new code, we see the >> 'a.xamdgcn-amdhsa.mkoffload.175t.slp1' dump change as follows (word-diff, >> only additional '<address-space-4>', occasionally): >> >> [...] >> {+<address-space-4>+} vector(2) long int * vectp.58; >> {+<address-space-4>+} vector(2) long int * vectp_.oacc_worker_o.57; >> {+<address-space-4>+} vector(2) int * vectp.56; >> {+<address-space-4>+} vector(2) int * vectp_.oacc_worker_o.55; >> [...] >> {+<address-space-4>+} long int * _104; >> [...] >> {+<address-space-4>+} long int * _108; >> [...] >> <address-space-4> void * _350; >> [...] >> _350 = __builtin_gcn_single_copy_start (&.oacc_worker_o.6); >> [...] >> MEM <{+<address-space-4>+} vector(2) long int> [(long int *)&.oacc_worker_o.6] = _101; >> _108 = &.oacc_worker_o.6._22 + 16; >> MEM <{+<address-space-4>+} vector(2) long int> [(long int *)_108] = _100; >> _104 = &.oacc_worker_o.6._22 + 32; >> [...] >> >> For example, with Kwok's new code, '_108' ('<address-space-4> long int *') >> is cast into '(long int *)' -- presumably synthesized in the SLP >> vectorizer? Is that correct or shouldn't that cast also include >> '<address-space-4>'? >> >> I see a similar issue has been fixed a while ago: r245772 (Git commit >> c7d97b2846c5647a81548caa3264d77c0a595010) for PR79723 >> "Another case of dropped gs: prefix", changing >> 'gcc/tree-vect-stmts.c:get_vectype_for_scalar_type_and_size' as follows: >> >> + /* Re-attach the address-space qualifier if we canonicalized the scalar >> + type. */ >> + if (TYPE_ADDR_SPACE (orig_scalar_type) != TYPE_ADDR_SPACE (vectype)) >> + return build_qualified_type >> + (vectype, KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (orig_scalar_type))); >> + >> return vectype; >> >> (It looks a bit like the address space handling is quite fragile in GCC's >> 'tree' types/interfaces? Do we have ideas about how to make that more >> robust, less "bolt-on"?) > > If in doubt always look at what RTL expansion does - it looks like > set_mem_attributes expects the address-space qualifier to be > present on the type or in case it is passed an object, on the > type of the base, or in case of a dereference, on the pointed-to > type of the pointer (and yes, that does look somewhat fragile). > > So it looks like the patch you refer to shouldn't fix anything and > >> + /* Re-attach the address-space qualifier if we canonicalized the scalar >> + type. */ >> + if (TYPE_ADDR_SPACE (orig_scalar_type) != TYPE_ADDR_SPACE (vectype)) >> + return build_qualified_type >> + (vectype, KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (orig_scalar_type))); > > looks incomplete. What you'd need to look for is MEM_REFs built > by the vectorizer and the address-space information on the pointers, > like generated from vect_create_data_ref_ptr. It might also be that > data-ref analysis / SCEV looks through address-space qualifier changing > casts and thus we pick up the wrong address-space in the end. Aah, more GCC pieces to learn about ;-) -- thanks for the pointers! > What's the testcase that ICEs on trunk? You'll need a GCN offloading build with the attached "[WIP] Reproduce GCN address space vs. SLP vectorization ICEs", run 'make check-target-libgomp', and observe a number of ICEs like: during RTL pass: expand [...]/libgomp.oacc-c-c++-common/loop-red-gwv-1.c: In function 'main._omp_fn.0': [...]/libgomp.oacc-c-c++-common/loop-red-gwv-1.c:19:9: internal compiler error: in convert_memory_address_addr_space_1, at explow.c:301 [...] mkoffload: fatal error: build-gcc/gcc/x86_64-pc-linux-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status '-O1 -ftree-slp-vectorize' would be sufficient to trigger that one. Run with '-save-temps -v', see the '[...]/build-gcc-offload-amdgcn-amdhsa/gcc/lto1' command ICE: #0 fancy_abort (file=file@entry=0x182e418 "[...]/source-gcc/gcc/explow.c", line=line@entry=301, function=function@entry=0x182e960 <convert_memory_address_addr_space_1(scalar_int_mode, rtx_def*, unsigned char, bool, bool)::__FUNCTION__> "convert_memory_address_addr_space_1") at [...]/source-gcc/gcc/diagnostic.c:1961 #1 0x00000000007ef690 in convert_memory_address_addr_space_1 (to_mode=..., x=x@entry=0x7ffff764fa08, as=as@entry=0 '\000', in_const=in_const@entry=false, no_emit=no_emit@entry=false) at [...]/source-gcc/gcc/explow.c:301 #2 0x00000000007ef6cb in convert_memory_address_addr_space (to_mode=..., x=0x7ffff764fa08, as=as@entry=0 '\000') at [...]/source-gcc/gcc/explow.c:423 #3 0x0000000000812f48 in expand_expr_addr_expr (modifier=EXPAND_SUM, tmode=E_DImode, target=0x0, exp=0x7ffff764a520) at [...]/source-gcc/gcc/expr.c:8535 #4 expand_expr_real_1 (exp=0x7ffff764a520, target=<optimized out>, tmode=<optimized out>, modifier=EXPAND_SUM, alt_rtl=0x0, inner_reference_p=<optimized out>) at [...]/source-gcc/gcc/expr.c:11741 #5 0x0000000000813139 in expand_expr (modifier=EXPAND_SUM, mode=E_VOIDmode, target=0x0, exp=0x7ffff764a520) at [...]/source-gcc/gcc/expr.h:301 #6 expand_expr_real_1 (exp=0x7ffff7649d48, target=<optimized out>, tmode=E_VOIDmode, modifier=EXPAND_WRITE, alt_rtl=0x0, inner_reference_p=<optimized out>) at [...]/source-gcc/gcc/expr.c:10887 #7 0x000000000082475a in expand_expr (modifier=EXPAND_WRITE, mode=E_VOIDmode, target=0x0, exp=0x7ffff7649d48) at [...]/source-gcc/gcc/expr.h:301 #8 expand_assignment (to=to@entry=0x7ffff7649d48, from=from@entry=0x7ffff763a7e0, nontemporal=<optimized out>) at [...]/source-gcc/gcc/expr.c:5732 #9 0x00000000006c807d in expand_gimple_stmt_1 (stmt=stmt@entry=0x7ffff7646aa0) at [...]/source-gcc/gcc/cfgexpand.c:3944 #10 0x00000000006c95c7 in expand_gimple_stmt (stmt=stmt@entry=0x7ffff7646aa0) at [...]/source-gcc/gcc/cfgexpand.c:4040 #11 0x00000000006ce884 in expand_gimple_basic_block (bb=0x7ffff7635dd0, disable_tail_calls=disable_tail_calls@entry=false) at [...]/source-gcc/gcc/cfgexpand.c:6082 #12 0x00000000006d13de in (anonymous namespace)::pass_expand::execute (this=<optimized out>, fun=<optimized out>) at [...]/source-gcc/gcc/cfgexpand.c:6808 [...] (gdb) up #1 0x00000000007ef690 in convert_memory_address_addr_space_1 (to_mode=..., x=x@entry=0x7ffff764fa08, as=as@entry=0 '\000', in_const=in_const@entry=false, no_emit=no_emit@entry=false) at [...]/source-gcc/gcc/explow.c:301 301 gcc_assert (GET_MODE (x) == to_mode || GET_MODE (x) == VOIDmode); (gdb) list 296 rtx x, addr_space_t as ATTRIBUTE_UNUSED, 297 bool in_const ATTRIBUTE_UNUSED, 298 bool no_emit ATTRIBUTE_UNUSED) 299 { 300 #ifndef POINTERS_EXTEND_UNSIGNED 301 gcc_assert (GET_MODE (x) == to_mode || GET_MODE (x) == VOIDmode); 302 return x; 303 #else /* defined(POINTERS_EXTEND_UNSIGNED) */ 304 scalar_int_mode pointer_mode, address_mode, from_mode; 305 rtx temp; (gdb) call debug_rtx(x) (symbol_ref:SI (".oacc_worker_o.13.6") [flags 0x2] <var_decl 0x7ffff7637d80 .oacc_worker_o.13>) (gdb) print x->mode $1 = E_SImode (gdb) print to_mode $2 = {m_mode = E_DImode} (gdb) up #2 0x00000000007ef6cb in convert_memory_address_addr_space (to_mode=..., x=0x7ffff764fa08, as=as@entry=0 '\000') at [...]/source-gcc/gcc/explow.c:423 423 return convert_memory_address_addr_space_1 (to_mode, x, as, false, false); (gdb) up #3 0x0000000000812f48 in expand_expr_addr_expr (modifier=EXPAND_SUM, tmode=E_DImode, target=0x0, exp=0x7ffff764a520) at [...]/source-gcc/gcc/expr.c:8535 8535 result = convert_memory_address_addr_space (new_tmode, result, as); (gdb) call debug_tree(exp) <addr_expr 0x7ffff764a520 type <pointer_type 0x7ffff7557888 type <integer_type 0x7ffff75505e8 int public SI size <integer_cst 0x7ffff754fbd0 constant 32> unit-size <integer_cst 0x7ffff754fbe8 constant 4> align:32 warn_if_not_align:0 symtab:0 alias-set 4 canonical-type 0x7ffff75505e8 precision:32 min <integer_cst 0x7ffff754fb88 -2147483648> max <integer_cst 0x7ffff754fba0 2147483647> pointer_to_this <pointer_type 0x7ffff7557888>> public unsigned DI size <integer_cst 0x7ffff754f990 constant 64> unit-size <integer_cst 0x7ffff754f9a8 constant 8> align:64 warn_if_not_align:0 symtab:0 alias-set 1 structural-equality> constant arg:0 <var_decl 0x7ffff7637d80 .oacc_worker_o.13 type <record_type 0x7ffff76215e8 .oacc_ws_data_s.0 address-space-4 no-force-blk BLK size <integer_cst 0x7ffff754f990 64> unit-size <integer_cst 0x7ffff754f9a8 8> align:32 warn_if_not_align:0 symtab:0 alias-set 5 canonical-type 0x7ffff76215e8 fields <field_decl 0x7ffff76317b8 t> pointer_to_this <pointer_type 0x7ffff76219d8>> addressable used static ignored BLK source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c:19:9 size <integer_cst 0x7ffff754f990 64> unit-size <integer_cst 0x7ffff754f9a8 8> align:128 warn_if_not_align:0 (mem/c:BLK (symbol_ref:SI (".oacc_worker_o.13.6") [flags 0x2] <var_decl 0x7ffff7637d80 .oacc_worker_o.13>) [5 .oacc_worker_o.13+0 S8 A128 AS4])>> In 'arg:0' of 'exp' note 'address-space-4' (expected): 'ADDR_SPACE_LDS' (per 'gcc/config/gcn/gcn.h:gcn_address_spaces'). With the attached "[WIP] [GCN] '+#define POINTERS_EXTEND_UNSIGNED 1'", we instead fail as follows: ./a.xamdgcn-amdhsa.mkoffload.2.s:92:23: error: invalid modifier 'rel32@lo' (no symbols present) s_add_u32 s2, s2, 32@rel32@lo+4 ^ ./a.xamdgcn-amdhsa.mkoffload.2.s:92:23: error: failed parsing operand. s_add_u32 s2, s2, 32@rel32@lo+4 ^ ./a.xamdgcn-amdhsa.mkoffload.2.s:93:24: error: invalid modifier 'rel32@hi' (no symbols present) s_addc_u32 s3, s3, 32@rel32@hi+4 ^ ./a.xamdgcn-amdhsa.mkoffload.2.s:93:24: error: failed parsing operand. s_addc_u32 s3, s3, 32@rel32@hi+4 ^ mkoffload: fatal error: build-gcc/gcc/x86_64-pc-linux-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status ..., so it's not that simple. (I have no clue whether 'POINTERS_EXTEND_UNSIGNED' would make sense for GCN -- but thought it was worth a quick try.) Grüße Thomas >> I did add a few 'assert's for non-generic address space to >> 'gcc/tree-vect*', but have not yet located where things may be going >> wrong. >> >> >> > 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', but in the case of 'omp_build_component_ref' called via >> GCN offloading compilation 'pass_omp_oacc_neuter_broadcast', 'obj_type' >> has a non-generic address space. >> >> 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 the attached "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, but >> maybe it's problematic for other cases? (This is, however, a separate >> issue from what I'm discussing here.) >> >> >> >> + tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); >> >> + if (TREE_THIS_VOLATILE (field)) >> >> + TREE_THIS_VOLATILE (ret) |= 1; >> >> + if (TREE_READONLY (field)) >> >> + TREE_READONLY (ret) |= 1; >> > >> > When touching these two, shouldn't it be better written as >> > = 1; instead of |= 1; ? For a bitfield... >> >> Yes, that was just copied from the original >> 'gcc/omp-general.c:omp_build_component_ref' -- but happy to simplify >> that, of course. >> >> >> 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 ----------------- 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 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-WIP-Reproduce-GCN-address-space-vs.-SLP-vectorizatio.patch --] [-- Type: text/x-diff, Size: 6960 bytes --] From eedea7a1041720ac4da9938716145c02918dd45e Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Fri, 20 Aug 2021 12:36:25 +0200 Subject: [PATCH] [WIP] Reproduce GCN address space vs. SLP vectorization ICEs --- gcc/omp-oacc-neuter-broadcast.cc | 2 ++ libgomp/testsuite/libgomp.c++/c++.exp | 3 +++ libgomp/testsuite/libgomp.c/c.exp | 3 +++ libgomp/testsuite/libgomp.fortran/fortran.exp | 3 +++ libgomp/testsuite/libgomp.graphite/graphite.exp | 3 +++ libgomp/testsuite/libgomp.oacc-c++/c++.exp | 11 +++++++++++ libgomp/testsuite/libgomp.oacc-c/c.exp | 11 +++++++++++ libgomp/testsuite/libgomp.oacc-fortran/fortran.exp | 11 +++++++++++ 8 files changed, 47 insertions(+) diff --git a/gcc/omp-oacc-neuter-broadcast.cc b/gcc/omp-oacc-neuter-broadcast.cc index d48627a6940..12b4d004c71 100644 --- a/gcc/omp-oacc-neuter-broadcast.cc +++ b/gcc/omp-oacc-neuter-broadcast.cc @@ -942,10 +942,12 @@ oacc_build_component_ref (tree obj, tree field) { tree field_type = TREE_TYPE (field); tree obj_type = TREE_TYPE (obj); +#if 0 // thus, 'oacc_build_component_ref' == 'gcc/omp-low.c:omp_build_component_ref' if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) field_type = build_qualified_type (field_type, KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); +#endif tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); if (TREE_THIS_VOLATILE (field)) diff --git a/libgomp/testsuite/libgomp.c++/c++.exp b/libgomp/testsuite/libgomp.c++/c++.exp index f4884e2ffa7..50448544084 100644 --- a/libgomp/testsuite/libgomp.c++/c++.exp +++ b/libgomp/testsuite/libgomp.c++/c++.exp @@ -1,3 +1,6 @@ +#TODO +return + load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp diff --git a/libgomp/testsuite/libgomp.c/c.exp b/libgomp/testsuite/libgomp.c/c.exp index 31bdd5795dc..a440a4c35b6 100644 --- a/libgomp/testsuite/libgomp.c/c.exp +++ b/libgomp/testsuite/libgomp.c/c.exp @@ -1,3 +1,6 @@ +#TODO +return + if [info exists lang_library_path] then { unset lang_library_path unset lang_link_flags diff --git a/libgomp/testsuite/libgomp.fortran/fortran.exp b/libgomp/testsuite/libgomp.fortran/fortran.exp index eb701311b6a..d7dea846afa 100644 --- a/libgomp/testsuite/libgomp.fortran/fortran.exp +++ b/libgomp/testsuite/libgomp.fortran/fortran.exp @@ -1,3 +1,6 @@ +#TODO +return + load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp load_gcc_lib gfortran-dg.exp diff --git a/libgomp/testsuite/libgomp.graphite/graphite.exp b/libgomp/testsuite/libgomp.graphite/graphite.exp index 4b01222bbc4..0d12a8323bf 100644 --- a/libgomp/testsuite/libgomp.graphite/graphite.exp +++ b/libgomp/testsuite/libgomp.graphite/graphite.exp @@ -1,3 +1,6 @@ +#TODO +return + # Copyright (C) 2009-2021 Free Software Foundation, Inc. # This program is free software; you can redistribute it and/or modify diff --git a/libgomp/testsuite/libgomp.oacc-c++/c++.exp b/libgomp/testsuite/libgomp.oacc-c++/c++.exp index 42e0395f9a5..f5cd9ff7513 100644 --- a/libgomp/testsuite/libgomp.oacc-c++/c++.exp +++ b/libgomp/testsuite/libgomp.oacc-c++/c++.exp @@ -89,9 +89,15 @@ if { $lang_test_file_found } { continue } host { + #TODO + continue + set acc_mem_shared 1 } nvidia { + #TODO + continue + if { ![check_effective_target_openacc_nvidia_accel_present] } { # Don't bother; execution testing is going to FAIL. untested "$subdir $offload_target offloading: supported, but hardware not accessible" @@ -107,11 +113,14 @@ if { $lang_test_file_found } { set acc_mem_shared 0 } radeon { + #TODO + if { 0 } { if { ![check_effective_target_openacc_radeon_accel_present] } { # Don't bother; execution testing is going to FAIL. untested "$subdir $offload_target offloading: supported, but hardware not accessible" continue } + } set acc_mem_shared 0 } @@ -144,6 +153,8 @@ if { $lang_test_file_found } { set-torture-options [list \ { -O0 } \ { -O2 } ] + #TODO + set-torture-options [list { -O2 -ftree-slp-vectorize } ] } } diff --git a/libgomp/testsuite/libgomp.oacc-c/c.exp b/libgomp/testsuite/libgomp.oacc-c/c.exp index 4bb2b2ac494..dd621568d2b 100644 --- a/libgomp/testsuite/libgomp.oacc-c/c.exp +++ b/libgomp/testsuite/libgomp.oacc-c/c.exp @@ -52,9 +52,15 @@ foreach offload_target [concat [split $offload_targets ","] "disable"] { continue } host { + #TODO + continue + set acc_mem_shared 1 } nvidia { + #TODO + continue + if { ![check_effective_target_openacc_nvidia_accel_present] } { # Don't bother; execution testing is going to FAIL. untested "$subdir $offload_target offloading: supported, but hardware not accessible" @@ -70,11 +76,14 @@ foreach offload_target [concat [split $offload_targets ","] "disable"] { set acc_mem_shared 0 } radeon { + #TODO + if { 0 } { if { ![check_effective_target_openacc_radeon_accel_present] } { # Don't bother; execution testing is going to FAIL. untested "$subdir $offload_target offloading: supported, but hardware not accessible" continue } + } set acc_mem_shared 0 } @@ -107,6 +116,8 @@ foreach offload_target [concat [split $offload_targets ","] "disable"] { set-torture-options [list \ { -O0 } \ { -O2 } ] + #TODO + set-torture-options [list { -O2 -ftree-slp-vectorize } ] } } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp index 7365b320668..85e5eb6f9d0 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp +++ b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp @@ -83,9 +83,15 @@ if { $lang_test_file_found } { continue } host { + #TODO + continue + set acc_mem_shared 1 } nvidia { + #TODO + continue + if { ![check_effective_target_openacc_nvidia_accel_present] } { # Don't bother; execution testing is going to FAIL. untested "$subdir $offload_target offloading: supported, but hardware not accessible" @@ -95,11 +101,14 @@ if { $lang_test_file_found } { set acc_mem_shared 0 } radeon { + #TODO + if { 0 } { if { ![check_effective_target_openacc_radeon_accel_present] } { # Don't bother; execution testing is going to FAIL. untested "$subdir $offload_target offloading: supported, but hardware not accessible" continue } + } set acc_mem_shared 0 } @@ -119,6 +128,8 @@ if { $lang_test_file_found } { # For Fortran we're doing torture testing, as Fortran has far more tests # with arrays etc. that testing just -O0 or -O2 is insufficient, that is # typically not the case for C/C++. + #TODO + set-torture-options [list { -O2 -ftree-slp-vectorize } ] gfortran-dg-runtest $tests "$tagopt" "" } unset offload_target -- 2.25.1 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #3: 0001-WIP-GCN-define-POINTERS_EXTEND_UNSIGNED-1.patch --] [-- Type: text/x-diff, Size: 4691 bytes --] From 2c99df97a60970f9d853aacc80a1485ad6f07052 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Fri, 20 Aug 2021 10:50:46 +0200 Subject: [PATCH] [WIP] [GCN] '+#define POINTERS_EXTEND_UNSIGNED 1' Doesn't change libgomp.oacc compile-time results. TODO But not yet execution-tested. Without 'gcn_addr_space_valid_pointer_mode', that runs into ICEs during libgomp build (only? -- because of '__lds' usage, I suppose?): /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/libgomp/config/gcn/../../task.c:2408:1: internal compiler error: in convert_debug_memory_address, at cfgexpand.c:4256 2408 | omp_in_final (void) | ^~~~~~~~~~~~ 0x7d64de ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/cfgexpand.c:4256 0x7e04b8 ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/cfgexpand.c:4701 0x7df2ba ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/cfgexpand.c:5305 0x7de3fc ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/cfgexpand.c:4489 0x7df2ba ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/cfgexpand.c:5305 0x7ecb0d ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/cfgexpand.c:5642 0xc5110f ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/passes.c:2567 0xc51a77 ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/passes.c:2656 0xc51ad4 ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/passes.c:2667 0x8359ec ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/cgraphunit.c:1828 0x8371fc ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/cgraphunit.c:1992 0x83a99d ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/cgraphunit.c:2269 0xd5bcaf ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/toplev.c:483 0x6305cd ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/toplev.c:2233 0x6331c6 ??? /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/main.c:39 Etc. That's: gcc_assert (targetm.addr_space.valid_pointer_mode (mode, as)); ... which (assumedly) is 'default_addr_space_valid_pointer_mode': return targetm.valid_pointer_mode (mode); ... which (assumedly) is 'default_valid_pointer_mode': return (mode == ptr_mode || mode == Pmode); #1 0x00000000007d64df in convert_debug_memory_address (mode=..., x=x@entry=0x7ffff75464d0, as=as@entry=4 '\004') at /home/thomas/tmp/source/gcc/build/queue-slim-omp/source-gcc/gcc/cfgexpand.c:4256 4256 gcc_assert (targetm.addr_space.valid_pointer_mode (mode, as)); (gdb) print mode $1 = {m_mode = E_SImode} (gdb) print as $2 = 4 '\004' --- gcc/config/gcn/gcn.c | 18 ++++++++++++++++++ gcc/config/gcn/gcn.h | 2 ++ 2 files changed, 20 insertions(+) diff --git a/gcc/config/gcn/gcn.c b/gcc/config/gcn/gcn.c index f2612803dac..9d7a6a2679a 100644 --- a/gcc/config/gcn/gcn.c +++ b/gcc/config/gcn/gcn.c @@ -1528,6 +1528,22 @@ gcn_addr_space_debug (addr_space_t as) gcc_unreachable (); } +/* TODO */ + +static bool +gcn_addr_space_valid_pointer_mode (scalar_int_mode mode, + addr_space_t as) +{ +#if 1 + if (as == ADDR_SPACE_LDS) + return mode == SImode; + else + return default_addr_space_valid_pointer_mode (mode, as); +#else //TODO + return mode == gcn_addr_space_pointer_mode (as); +#endif +} + /* Implement REGNO_MODE_CODE_OK_FOR_BASE_P via gcn.h @@ -6452,6 +6468,8 @@ gcn_dwarf_register_span (rtx rtl) #define TARGET_ADDR_SPACE_SUBSET_P gcn_addr_space_subset_p #undef TARGET_ADDR_SPACE_CONVERT #define TARGET_ADDR_SPACE_CONVERT gcn_addr_space_convert +#undef TARGET_ADDR_SPACE_VALID_POINTER_MODE +#define TARGET_ADDR_SPACE_VALID_POINTER_MODE gcn_addr_space_valid_pointer_mode #undef TARGET_ARG_PARTIAL_BYTES #define TARGET_ARG_PARTIAL_BYTES gcn_arg_partial_bytes #undef TARGET_ASM_ALIGNED_DI_OP diff --git a/gcc/config/gcn/gcn.h b/gcc/config/gcn/gcn.h index 5822ec34aa7..c0bd5565a49 100644 --- a/gcc/config/gcn/gcn.h +++ b/gcc/config/gcn/gcn.h @@ -66,6 +66,8 @@ #define CASE_VECTOR_MODE DImode #define FUNCTION_MODE QImode +#define POINTERS_EXTEND_UNSIGNED 1 + #define DATA_ALIGNMENT(TYPE,ALIGN) ((ALIGN) > 128 ? (ALIGN) : 128) #define LOCAL_ALIGNMENT(TYPE,ALIGN) ((ALIGN) > 64 ? (ALIGN) : 64) #define STACK_SLOT_ALIGNMENT(TYPE,MODE,ALIGN) ((ALIGN) > 64 ? (ALIGN) : 64) -- 2.25.1 ^ permalink raw reply [flat|nested] 15+ messages in thread
* Re: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' 2021-08-23 14:30 ` Thomas Schwinge @ 2021-08-24 7:43 ` Richard Biener 0 siblings, 0 replies; 15+ messages in thread From: Richard Biener @ 2021-08-24 7:43 UTC (permalink / raw) To: Thomas Schwinge Cc: Jakub Jelinek, GCC Patches, Kwok Cheung Yeung, Julian Brown, Andrew Stubbs On Mon, Aug 23, 2021 at 4:30 PM Thomas Schwinge <thomas@codesourcery.com> wrote: > > Hi! > > On 2021-08-20T09:51:36+0200, Richard Biener <richard.guenther@gmail.com> wrote: > > On Thu, Aug 19, 2021 at 10:14 PM Thomas Schwinge > > <thomas@codesourcery.com> wrote: > >> Richard, maybe you have an opinion here, in particular about my > >> "SLP vectorizer" comment below? Please see > >> <http://mid.mail-archive.com/87r1f2puss.fsf@euler.schwinge.homeip.net> > >> for the full context. > >> > >> 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: > >> >> /* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it > >> >> as appropriate. */ > >> >> > >> >> tree > >> >> omp_build_component_ref (tree obj, tree field) > >> >> { > >> >> + tree field_type = TREE_TYPE (field); > >> >> + tree obj_type = TREE_TYPE (obj); > >> >> + if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) > >> >> + field_type > >> >> + = build_qualified_type (field_type, > >> >> + KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); > >> > >> (For later reference: "Kwok's new code" here is to propagate to > >> 'field_type' any non-generic address space of 'obj_type'.) > >> > >> |> 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'm attaching an altered and extended variant that does, > >> "Add 'libgomp.c/address-space-1.c'". OK to push to master branch? > >> > >> 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 > >> size <integer_cst 0x7ffff7540f30 constant 32> > >> unit-size <integer_cst 0x7ffff7540f48 constant 4> > >> align:32 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686498 precision:32 min <integer_cst 0x7ffff7540ee8 -2147483648> max <integer_cst 0x7ffff7540f00 2147483647> > >> pointer_to_this <pointer_type 0x7ffff7686b28>> > >> unsigned DI > >> size <integer_cst 0x7ffff7540cf0 type <integer_type 0x7ffff75590a8 bitsizetype> constant 64> > >> unit-size <integer_cst 0x7ffff7540d08 type <integer_type 0x7ffff7559000 sizetype> constant 8> > >> align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686b28> > >> > >> (gdb) call debug_tree(obj_type) > >> <record_type 0x7ffff7686bd0 .omp_data_t.0 readonly DI > >> size <integer_cst 0x7ffff7540cf0 type <integer_type 0x7ffff75590a8 bitsizetype> constant 64> > >> unit-size <integer_cst 0x7ffff7540d08 type <integer_type 0x7ffff7559000 sizetype> constant 8> > >> align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686bd0 > >> fields <field_decl 0x7ffff7568428 a > >> type <pointer_type 0x7ffff7686b28 type <integer_type 0x7ffff7686498 int address-space-1> > >> unsigned DI size <integer_cst 0x7ffff7540cf0 64> unit-size <integer_cst 0x7ffff7540d08 8> > >> align:64 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7686b28> > >> unsigned DI /home/thomas/shared/gcc/omp/as.c:4:14 size <integer_cst 0x7ffff7540cf0 64> unit-size <integer_cst 0x7ffff7540d08 8> > >> align:64 warn_if_not_align:0 offset_align 128 > >> offset <integer_cst 0x7ffff7540d20 constant 0> > >> bit-offset <integer_cst 0x7ffff7540d68 constant 0> context <record_type 0x7ffff7686540 .omp_data_t.0>> reference_to_this <reference_type 0x7ffff7686c78>> > >> > >> The case that Kwok's new code handles, however, is when 'obj_type' has a > >> non-generic address space, and then propagates that one to 'field_type'. > >> > >> For a similar OpenACC example, 'omp_build_component_ref' called via GCN > >> offloading compilation 'pass_omp_oacc_neuter_broadcast', we've got > >> without Kwok's new code: > >> > >> (gdb) call debug_tree(field_type) > >> <boolean_type 0x7ffff7550b28 bool public unsigned QI > >> size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> > >> unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> > >> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7550b28 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> > >> > >> (gdb) call debug_tree(obj_type) > >> <record_type 0x7ffff7631000 .oacc_ws_data_s.0 address-space-4 QI > >> size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> > >> unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> > >> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7631000 > >> fields <field_decl 0x7ffff762e260 _52 > >> type <boolean_type 0x7ffff7550b28 bool public unsigned QI size <integer_cst 0x7ffff754fa80 8> unit-size <integer_cst 0x7ffff754fa98 1> > >> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7550b28 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> > >> unsigned QI <built-in>:0:0 size <integer_cst 0x7ffff754fa80 8> unit-size <integer_cst 0x7ffff754fa98 1> > >> align:8 warn_if_not_align:0 offset_align 64 > >> offset <integer_cst 0x7ffff754f9c0 constant 0> > >> bit-offset <integer_cst 0x7ffff754fa08 constant 0> context <record_type 0x7ffff7631000 .oacc_ws_data_s.0 address-space-4>> > >> pointer_to_this <pointer_type 0x7ffff7631498>> > >> > >> ..., and with Kwok's new code the 'address-space-4' of 'obj_type' is > >> propagated to 'field_type': > >> > >> (gdb) call debug_tree(field_type) > >> <boolean_type 0x7ffff7631540 bool address-space-4 unsigned QI > >> size <integer_cst 0x7ffff754fa80 type <integer_type 0x7ffff75500a8 bitsizetype> constant 8> > >> unit-size <integer_cst 0x7ffff754fa98 type <integer_type 0x7ffff7550000 sizetype> constant 1> > >> align:8 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type 0x7ffff7631540 precision:1 min <integer_cst 0x7ffff754fcd8 0> max <integer_cst 0x7ffff754fd08 1>> > >> > >> I'm not familiar enough with these bits to tell whether Kwok's new code > >> is the right solution to this problem -- or if, for example, the problem > >> is rather in the SLP vectorizer, where the ICE seems to ultimately > >> emerge? > >> > >> Without (ICEs later) vs. with (works) Kwok's new code, we see the > >> 'a.xamdgcn-amdhsa.mkoffload.175t.slp1' dump change as follows (word-diff, > >> only additional '<address-space-4>', occasionally): > >> > >> [...] > >> {+<address-space-4>+} vector(2) long int * vectp.58; > >> {+<address-space-4>+} vector(2) long int * vectp_.oacc_worker_o.57; > >> {+<address-space-4>+} vector(2) int * vectp.56; > >> {+<address-space-4>+} vector(2) int * vectp_.oacc_worker_o.55; > >> [...] > >> {+<address-space-4>+} long int * _104; > >> [...] > >> {+<address-space-4>+} long int * _108; > >> [...] > >> <address-space-4> void * _350; > >> [...] > >> _350 = __builtin_gcn_single_copy_start (&.oacc_worker_o.6); > >> [...] > >> MEM <{+<address-space-4>+} vector(2) long int> [(long int *)&.oacc_worker_o.6] = _101; > >> _108 = &.oacc_worker_o.6._22 + 16; > >> MEM <{+<address-space-4>+} vector(2) long int> [(long int *)_108] = _100; > >> _104 = &.oacc_worker_o.6._22 + 32; > >> [...] > >> > >> For example, with Kwok's new code, '_108' ('<address-space-4> long int *') > >> is cast into '(long int *)' -- presumably synthesized in the SLP > >> vectorizer? Is that correct or shouldn't that cast also include > >> '<address-space-4>'? > >> > >> I see a similar issue has been fixed a while ago: r245772 (Git commit > >> c7d97b2846c5647a81548caa3264d77c0a595010) for PR79723 > >> "Another case of dropped gs: prefix", changing > >> 'gcc/tree-vect-stmts.c:get_vectype_for_scalar_type_and_size' as follows: > >> > >> + /* Re-attach the address-space qualifier if we canonicalized the scalar > >> + type. */ > >> + if (TYPE_ADDR_SPACE (orig_scalar_type) != TYPE_ADDR_SPACE (vectype)) > >> + return build_qualified_type > >> + (vectype, KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (orig_scalar_type))); > >> + > >> return vectype; > >> > >> (It looks a bit like the address space handling is quite fragile in GCC's > >> 'tree' types/interfaces? Do we have ideas about how to make that more > >> robust, less "bolt-on"?) > > > > If in doubt always look at what RTL expansion does - it looks like > > set_mem_attributes expects the address-space qualifier to be > > present on the type or in case it is passed an object, on the > > type of the base, or in case of a dereference, on the pointed-to > > type of the pointer (and yes, that does look somewhat fragile). > > > > So it looks like the patch you refer to shouldn't fix anything and > > > >> + /* Re-attach the address-space qualifier if we canonicalized the scalar > >> + type. */ > >> + if (TYPE_ADDR_SPACE (orig_scalar_type) != TYPE_ADDR_SPACE (vectype)) > >> + return build_qualified_type > >> + (vectype, KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (orig_scalar_type))); > > > > looks incomplete. What you'd need to look for is MEM_REFs built > > by the vectorizer and the address-space information on the pointers, > > like generated from vect_create_data_ref_ptr. It might also be that > > data-ref analysis / SCEV looks through address-space qualifier changing > > casts and thus we pick up the wrong address-space in the end. > > Aah, more GCC pieces to learn about ;-) -- thanks for the pointers! > > > What's the testcase that ICEs on trunk? > > You'll need a GCN offloading build with the attached > "[WIP] Reproduce GCN address space vs. SLP vectorization ICEs", > run 'make check-target-libgomp', and observe a number of ICEs like: Eh, OK ;) Too much for a quick look - if you got sth that ICEs / shows missing address-spaces and that is reproducible with a cc1 cross to nvptx/gcn and a C testcase then I'm in to debug where the vectorizer is at fault ;) Richard. > during RTL pass: expand > [...]/libgomp.oacc-c-c++-common/loop-red-gwv-1.c: In function 'main._omp_fn.0': > [...]/libgomp.oacc-c-c++-common/loop-red-gwv-1.c:19:9: internal compiler error: in convert_memory_address_addr_space_1, at explow.c:301 > [...] > mkoffload: fatal error: build-gcc/gcc/x86_64-pc-linux-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status > > '-O1 -ftree-slp-vectorize' would be sufficient to trigger that one. > Run with '-save-temps -v', see the > '[...]/build-gcc-offload-amdgcn-amdhsa/gcc/lto1' command ICE: > > #0 fancy_abort (file=file@entry=0x182e418 "[...]/source-gcc/gcc/explow.c", line=line@entry=301, function=function@entry=0x182e960 <convert_memory_address_addr_space_1(scalar_int_mode, rtx_def*, unsigned char, bool, bool)::__FUNCTION__> "convert_memory_address_addr_space_1") at [...]/source-gcc/gcc/diagnostic.c:1961 > #1 0x00000000007ef690 in convert_memory_address_addr_space_1 (to_mode=..., x=x@entry=0x7ffff764fa08, as=as@entry=0 '\000', in_const=in_const@entry=false, no_emit=no_emit@entry=false) at [...]/source-gcc/gcc/explow.c:301 > #2 0x00000000007ef6cb in convert_memory_address_addr_space (to_mode=..., x=0x7ffff764fa08, as=as@entry=0 '\000') at [...]/source-gcc/gcc/explow.c:423 > #3 0x0000000000812f48 in expand_expr_addr_expr (modifier=EXPAND_SUM, tmode=E_DImode, target=0x0, exp=0x7ffff764a520) at [...]/source-gcc/gcc/expr.c:8535 > #4 expand_expr_real_1 (exp=0x7ffff764a520, target=<optimized out>, tmode=<optimized out>, modifier=EXPAND_SUM, alt_rtl=0x0, inner_reference_p=<optimized out>) at [...]/source-gcc/gcc/expr.c:11741 > #5 0x0000000000813139 in expand_expr (modifier=EXPAND_SUM, mode=E_VOIDmode, target=0x0, exp=0x7ffff764a520) at [...]/source-gcc/gcc/expr.h:301 > #6 expand_expr_real_1 (exp=0x7ffff7649d48, target=<optimized out>, tmode=E_VOIDmode, modifier=EXPAND_WRITE, alt_rtl=0x0, inner_reference_p=<optimized out>) at [...]/source-gcc/gcc/expr.c:10887 > #7 0x000000000082475a in expand_expr (modifier=EXPAND_WRITE, mode=E_VOIDmode, target=0x0, exp=0x7ffff7649d48) at [...]/source-gcc/gcc/expr.h:301 > #8 expand_assignment (to=to@entry=0x7ffff7649d48, from=from@entry=0x7ffff763a7e0, nontemporal=<optimized out>) at [...]/source-gcc/gcc/expr.c:5732 > #9 0x00000000006c807d in expand_gimple_stmt_1 (stmt=stmt@entry=0x7ffff7646aa0) at [...]/source-gcc/gcc/cfgexpand.c:3944 > #10 0x00000000006c95c7 in expand_gimple_stmt (stmt=stmt@entry=0x7ffff7646aa0) at [...]/source-gcc/gcc/cfgexpand.c:4040 > #11 0x00000000006ce884 in expand_gimple_basic_block (bb=0x7ffff7635dd0, disable_tail_calls=disable_tail_calls@entry=false) at [...]/source-gcc/gcc/cfgexpand.c:6082 > #12 0x00000000006d13de in (anonymous namespace)::pass_expand::execute (this=<optimized out>, fun=<optimized out>) at [...]/source-gcc/gcc/cfgexpand.c:6808 > [...] > (gdb) up > #1 0x00000000007ef690 in convert_memory_address_addr_space_1 (to_mode=..., x=x@entry=0x7ffff764fa08, as=as@entry=0 '\000', in_const=in_const@entry=false, no_emit=no_emit@entry=false) at [...]/source-gcc/gcc/explow.c:301 > 301 gcc_assert (GET_MODE (x) == to_mode || GET_MODE (x) == VOIDmode); > (gdb) list > 296 rtx x, addr_space_t as ATTRIBUTE_UNUSED, > 297 bool in_const ATTRIBUTE_UNUSED, > 298 bool no_emit ATTRIBUTE_UNUSED) > 299 { > 300 #ifndef POINTERS_EXTEND_UNSIGNED > 301 gcc_assert (GET_MODE (x) == to_mode || GET_MODE (x) == VOIDmode); > 302 return x; > 303 #else /* defined(POINTERS_EXTEND_UNSIGNED) */ > 304 scalar_int_mode pointer_mode, address_mode, from_mode; > 305 rtx temp; > (gdb) call debug_rtx(x) > (symbol_ref:SI (".oacc_worker_o.13.6") [flags 0x2] <var_decl 0x7ffff7637d80 .oacc_worker_o.13>) > (gdb) print x->mode > $1 = E_SImode > (gdb) print to_mode > $2 = {m_mode = E_DImode} > (gdb) up > #2 0x00000000007ef6cb in convert_memory_address_addr_space (to_mode=..., x=0x7ffff764fa08, as=as@entry=0 '\000') at [...]/source-gcc/gcc/explow.c:423 > 423 return convert_memory_address_addr_space_1 (to_mode, x, as, false, false); > (gdb) up > #3 0x0000000000812f48 in expand_expr_addr_expr (modifier=EXPAND_SUM, tmode=E_DImode, target=0x0, exp=0x7ffff764a520) at [...]/source-gcc/gcc/expr.c:8535 > 8535 result = convert_memory_address_addr_space (new_tmode, result, as); > (gdb) call debug_tree(exp) > <addr_expr 0x7ffff764a520 > type <pointer_type 0x7ffff7557888 > type <integer_type 0x7ffff75505e8 int public SI > size <integer_cst 0x7ffff754fbd0 constant 32> > unit-size <integer_cst 0x7ffff754fbe8 constant 4> > align:32 warn_if_not_align:0 symtab:0 alias-set 4 canonical-type 0x7ffff75505e8 precision:32 min <integer_cst 0x7ffff754fb88 -2147483648> max <integer_cst 0x7ffff754fba0 2147483647> > pointer_to_this <pointer_type 0x7ffff7557888>> > public unsigned DI > size <integer_cst 0x7ffff754f990 constant 64> > unit-size <integer_cst 0x7ffff754f9a8 constant 8> > align:64 warn_if_not_align:0 symtab:0 alias-set 1 structural-equality> > constant > arg:0 <var_decl 0x7ffff7637d80 .oacc_worker_o.13 > type <record_type 0x7ffff76215e8 .oacc_ws_data_s.0 address-space-4 no-force-blk BLK size <integer_cst 0x7ffff754f990 64> unit-size <integer_cst 0x7ffff754f9a8 8> > align:32 warn_if_not_align:0 symtab:0 alias-set 5 canonical-type 0x7ffff76215e8 fields <field_decl 0x7ffff76317b8 t> > pointer_to_this <pointer_type 0x7ffff76219d8>> > addressable used static ignored BLK source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c:19:9 size <integer_cst 0x7ffff754f990 64> unit-size <integer_cst 0x7ffff754f9a8 8> > align:128 warn_if_not_align:0 > (mem/c:BLK (symbol_ref:SI (".oacc_worker_o.13.6") [flags 0x2] <var_decl 0x7ffff7637d80 .oacc_worker_o.13>) [5 .oacc_worker_o.13+0 S8 A128 AS4])>> > > In 'arg:0' of 'exp' note 'address-space-4' (expected): 'ADDR_SPACE_LDS' > (per 'gcc/config/gcn/gcn.h:gcn_address_spaces'). > > > With the attached "[WIP] [GCN] '+#define POINTERS_EXTEND_UNSIGNED 1'", we > instead fail as follows: > > ./a.xamdgcn-amdhsa.mkoffload.2.s:92:23: error: invalid modifier 'rel32@lo' (no symbols present) > s_add_u32 s2, s2, 32@rel32@lo+4 > ^ > ./a.xamdgcn-amdhsa.mkoffload.2.s:92:23: error: failed parsing operand. > s_add_u32 s2, s2, 32@rel32@lo+4 > ^ > ./a.xamdgcn-amdhsa.mkoffload.2.s:93:24: error: invalid modifier 'rel32@hi' (no symbols present) > s_addc_u32 s3, s3, 32@rel32@hi+4 > ^ > ./a.xamdgcn-amdhsa.mkoffload.2.s:93:24: error: failed parsing operand. > s_addc_u32 s3, s3, 32@rel32@hi+4 > ^ > mkoffload: fatal error: build-gcc/gcc/x86_64-pc-linux-gnu-accel-amdgcn-amdhsa-gcc returned 1 exit status > > ..., so it's not that simple. (I have no clue whether > 'POINTERS_EXTEND_UNSIGNED' would make sense for GCN -- but thought it was > worth a quick try.) > > > Grüße > Thomas > > > >> I did add a few 'assert's for non-generic address space to > >> 'gcc/tree-vect*', but have not yet located where things may be going > >> wrong. > >> > >> > >> > 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', but in the case of 'omp_build_component_ref' called via > >> GCN offloading compilation 'pass_omp_oacc_neuter_broadcast', 'obj_type' > >> has a non-generic address space. > >> > >> 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 the attached "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, but > >> maybe it's problematic for other cases? (This is, however, a separate > >> issue from what I'm discussing here.) > >> > >> > >> >> + tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); > >> >> + if (TREE_THIS_VOLATILE (field)) > >> >> + TREE_THIS_VOLATILE (ret) |= 1; > >> >> + if (TREE_READONLY (field)) > >> >> + TREE_READONLY (ret) |= 1; > >> > > >> > When touching these two, shouldn't it be better written as > >> > = 1; instead of |= 1; ? For a bitfield... > >> > >> Yes, that was just copied from the original > >> 'gcc/omp-general.c:omp_build_component_ref' -- but happy to simplify > >> that, of course. > >> > >> > >> 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 > > > ----------------- > 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 ^ permalink raw reply [flat|nested] 15+ messages in thread
* Re: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' 2021-08-19 20:13 ` Thomas Schwinge 2021-08-20 7:51 ` Richard Biener @ 2021-08-20 14:49 ` Jakub Jelinek 2021-08-23 15:55 ` Add 'libgomp.c/address-space-1.c' (was: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref') Thomas Schwinge 2021-08-24 10:23 ` Host and offload targets have no common meaning of address spaces " Thomas Schwinge 2 siblings, 1 reply; 15+ messages in thread From: Jakub Jelinek @ 2021-08-20 14:49 UTC (permalink / raw) To: Thomas Schwinge Cc: Richard Biener, gcc-patches, Kwok Cheung Yeung, Julian Brown On Thu, Aug 19, 2021 at 10:13:56PM +0200, Thomas Schwinge wrote: > libgomp/ > * testsuite/libgomp.c/address-space-1.c: New file. > > Co-authored-by: Jakub Jelinek <jakub@redhat.com> > --- > libgomp/testsuite/libgomp.c/address-space-1.c | 24 +++++++++++++++++++ > 1 file changed, 24 insertions(+) > create mode 100644 libgomp/testsuite/libgomp.c/address-space-1.c > > diff --git a/libgomp/testsuite/libgomp.c/address-space-1.c b/libgomp/testsuite/libgomp.c/address-space-1.c > new file mode 100644 > index 00000000000..90244db03b1 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c/address-space-1.c > @@ -0,0 +1,24 @@ > +/* Verify OMP instances of variables with address space. */ > + > +/* { dg-do run { target i?86-*-* x86_64-*-* } } */ > +/* { dg-require-effective-target offload_device_nonshared_as } */ > + > +#include <assert.h> > + > +int __seg_fs a; > + > +int > +main (void) > +{ > + // a = 123; // SIGSEGV > + int b; > +#pragma omp target map(alloc: a) map(from: b) > + { > + a = 321; // no SIGSEGV (given 'offload_device_nonshared_as') > + asm volatile ("" : : : "memory"); Maybe better asm volatile ("" : : "g" (&a) : "memory"); so that the compiler doesn't think it could optimize it away to just b = 321; Ok with that change. > + b = a; > + } > + assert (b == 321); > + > + return 0; > +} > -- > 2.30.2 > Jakub ^ permalink raw reply [flat|nested] 15+ messages in thread
* Add 'libgomp.c/address-space-1.c' (was: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref') 2021-08-20 14:49 ` Jakub Jelinek @ 2021-08-23 15:55 ` Thomas Schwinge 0 siblings, 0 replies; 15+ messages in thread From: Thomas Schwinge @ 2021-08-23 15:55 UTC (permalink / raw) To: Jakub Jelinek, gcc-patches Cc: Richard Biener, Kwok Cheung Yeung, Julian Brown [-- Attachment #1: Type: text/plain, Size: 2192 bytes --] Hi! On 2021-08-20T16:49:25+0200, Jakub Jelinek <jakub@redhat.com> wrote: >> --- /dev/null >> +++ b/libgomp/testsuite/libgomp.c/address-space-1.c >> @@ -0,0 +1,24 @@ >> +/* Verify OMP instances of variables with address space. */ >> + >> +/* { dg-do run { target i?86-*-* x86_64-*-* } } */ >> +/* { dg-require-effective-target offload_device_nonshared_as } */ I also added: +/* 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 } } */ Might this be a symptom related to my earlier 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 the attached "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, but | maybe it's problematic for other cases? ..., or it's yet another problem, of course. ;-) >> +#include <assert.h> >> + >> +int __seg_fs a; >> + >> +int >> +main (void) >> +{ >> + // a = 123; // SIGSEGV >> + int b; >> +#pragma omp target map(alloc: a) map(from: b) >> + { >> + a = 321; // no SIGSEGV (given 'offload_device_nonshared_as') >> + asm volatile ("" : : : "memory"); > > Maybe better asm volatile ("" : : "g" (&a) : "memory"); > so that the compiler doesn't think it could optimize it away to > just b = 321; Thanks. >> + b = a; >> + } >> + assert (b == 321); >> + >> + return 0; >> +} Pushed "Add 'libgomp.c/address-space-1.c'" to master branch in commit 29c355f76ceeb4639c21acaf52c50d35c8472720, see attached. 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 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Add-libgomp.c-address-space-1.c.patch --] [-- Type: text/x-diff, Size: 1518 bytes --] From 29c355f76ceeb4639c21acaf52c50d35c8472720 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Thu, 19 Aug 2021 15:14:51 +0200 Subject: [PATCH] Add 'libgomp.c/address-space-1.c' Intel MIC (emulated) offloading execution failure remains to be analyzed. libgomp/ * testsuite/libgomp.c/address-space-1.c: New file. Co-authored-by: Jakub Jelinek <jakub@redhat.com> --- libgomp/testsuite/libgomp.c/address-space-1.c | 28 +++++++++++++++++++ 1 file changed, 28 insertions(+) create mode 100644 libgomp/testsuite/libgomp.c/address-space-1.c diff --git a/libgomp/testsuite/libgomp.c/address-space-1.c b/libgomp/testsuite/libgomp.c/address-space-1.c new file mode 100644 index 00000000000..6ad57deec42 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/address-space-1.c @@ -0,0 +1,28 @@ +/* Verify OMP instances of variables with address space. */ + +/* { 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; + +int +main (void) +{ + // a = 123; // SIGSEGV + int b; +#pragma omp target map(alloc: a) map(from: b) + { + a = 321; // no SIGSEGV (given 'offload_device_nonshared_as') + asm volatile ("" : : "g" (&a) : "memory"); + b = a; + } + assert (b == 321); + + return 0; +} -- 2.25.1 ^ permalink raw reply [flat|nested] 15+ messages in thread
* Host and offload targets have no common meaning of address spaces (was: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref') 2021-08-19 20:13 ` Thomas Schwinge 2021-08-20 7:51 ` Richard Biener 2021-08-20 14:49 ` Jakub Jelinek @ 2021-08-24 10:23 ` Thomas Schwinge 2021-08-24 11:43 ` Richard Biener 2021-09-10 8:03 ` Thomas Schwinge 2 siblings, 2 replies; 15+ messages in thread From: Thomas Schwinge @ 2021-08-24 10:23 UTC (permalink / raw) To: Jakub Jelinek, Richard Biener, gcc-patches Cc: Kwok Cheung Yeung, Julian Brown, Andrew Stubbs [-- Attachment #1: Type: text/plain, Size: 4266 bytes --] 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 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Host-and-offload-targets-have-no-common-meaning-of-a.patch --] [-- Type: text/x-diff, Size: 2653 bytes --] 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 ^ permalink raw reply [flat|nested] 15+ messages in thread
* 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') 2021-08-24 10:23 ` Host and offload targets have no common meaning of address spaces " Thomas Schwinge @ 2021-08-24 11:43 ` Richard Biener 2021-09-03 11:42 ` Andrew Stubbs 2022-01-13 10:24 ` Host and offload targets have no common meaning of address spaces Thomas Schwinge 2021-09-10 8:03 ` Thomas Schwinge 1 sibling, 2 replies; 15+ messages in thread From: Richard Biener @ 2021-08-24 11:43 UTC (permalink / raw) To: Thomas Schwinge Cc: Jakub Jelinek, GCC Patches, Kwok Cheung Yeung, Julian Brown, Andrew Stubbs 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 ^ permalink raw reply [flat|nested] 15+ messages in thread
* 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') 2021-08-24 11:43 ` Richard Biener @ 2021-09-03 11:42 ` Andrew Stubbs 2022-01-13 10:24 ` Host and offload targets have no common meaning of address spaces Thomas Schwinge 1 sibling, 0 replies; 15+ messages in thread From: Andrew Stubbs @ 2021-09-03 11:42 UTC (permalink / raw) To: Richard Biener, Thomas Schwinge; +Cc: Jakub Jelinek, Julian Brown, GCC Patches 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 ^ permalink raw reply [flat|nested] 15+ messages in thread
* Host and offload targets have no common meaning of address spaces 2021-08-24 11:43 ` Richard Biener 2021-09-03 11:42 ` Andrew Stubbs @ 2022-01-13 10:24 ` Thomas Schwinge 1 sibling, 0 replies; 15+ messages in thread From: Thomas Schwinge @ 2022-01-13 10:24 UTC (permalink / raw) To: Jakub Jelinek, gcc-patches [-- Attachment #1: Type: text/plain, Size: 4998 bytes --] 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 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Host-and-offload-targets-have-no-common-meaning-of-a.patch --] [-- Type: text/x-diff, Size: 2672 bytes --] 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 ^ permalink raw reply [flat|nested] 15+ messages in thread
* Re: Host and offload targets have no common meaning of address spaces 2021-08-24 10:23 ` Host and offload targets have no common meaning of address spaces " Thomas Schwinge 2021-08-24 11:43 ` Richard Biener @ 2021-09-10 8:03 ` Thomas Schwinge 1 sibling, 0 replies; 15+ messages in thread From: Thomas Schwinge @ 2021-09-10 8:03 UTC (permalink / raw) To: Jakub Jelinek, Richard Biener, gcc-patches; +Cc: Julian Brown, Andrew Stubbs [-- Attachment #1: Type: text/plain, Size: 4585 bytes --] 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 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Host-and-offload-targets-have-no-common-meaning-of-a.patch --] [-- Type: text/x-diff, Size: 2653 bytes --] 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 ^ permalink raw reply [flat|nested] 15+ messages in thread
* Get rid of 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref' (was: Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref') 2021-08-09 14:16 ` Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' Thomas Schwinge 2021-08-16 8:08 ` [ping] " Thomas Schwinge @ 2022-02-22 17:00 ` Thomas Schwinge 1 sibling, 0 replies; 15+ messages in thread From: Thomas Schwinge @ 2022-02-22 17:00 UTC (permalink / raw) To: Jakub Jelinek, gcc-patches Cc: Kwok Cheung Yeung, Julian Brown, Richard Biener, Andrew Stubbs [-- Attachment #1: Type: text/plain, Size: 9761 bytes --] Hi! On 2021-08-09T16:16:51+0200, I wrote: > This concerns a class of ICEs seen as of og10 branch with the > "openacc: Middle-end worker-partitioning support" and "amdgcn: > Enable OpenACC worker partitioning for AMD GCN" changes applied: I've determined that as of commit 2a3f9f6532bb21d8ab6f16fbe9ee603f6b1405f2 "openacc: Shared memory layout optimisation", we're no longer running into the vectorizer ICEs for '!ADDR_SPACE_GENERIC_P'. I have not researched if they've just gone latent (again), or whether that commit really changed something to avoid those (bug fix). Anyway: pushed to master branch commit 54f745023276e5025e34b2cc22530c78423a93cb "Get rid of 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref'", see attached. Grüße Thomas > On 2020-06-06T16:07:36+0100, Kwok Cheung Yeung <kwok_yeung@mentor.com> wrote: >> On 01/06/2020 8:48 pm, Kwok Cheung Yeung wrote: >>> On 21/05/2020 10:23 pm, Kwok Cheung Yeung wrote: >>>> These all have the same failure mode: >>>> >>>> during RTL pass: expand >>>> [...]/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90: In function 'MAIN__._omp_fn.1': >>>> [...]/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90:86: internal compiler error: in convert_memory_address_addr_space_1, at explow.c:302 >>>> 0xc29f20 convert_memory_address_addr_space_1(scalar_int_mode, rtx_def*, unsigned char, bool, bool) >>>> [...]/gcc/explow.c:302 >>>> 0xc29f57 convert_memory_address_addr_space(scalar_int_mode, rtx_def*, unsigned char) >>>> [...]/gcc/explow.c:404 >>>> [...] > >>>> This occurs if the -ftree-slp-vectorize flag is specified (default at -O3). > >>> The problematic bit of Gimple code is this: >>> >>> .oacc_worker_o.44._120 = gangs_min_472; >>> .oacc_worker_o.44._122 = workers_min_473; >>> .oacc_worker_o.44._124 = vectors_min_474; >>> .oacc_worker_o.44._126 = gangs_max_475; >>> .oacc_worker_o.44._128 = workers_max_476; >>> .oacc_worker_o.44._130 = vectors_max_477; >>> .oacc_worker_o.44._132 = 0; >>> >>> With SLP vectorization enabled, it becomes this: >>> >>> _40 = {gangs_min_472, workers_min_473, vectors_min_474, gangs_max_475}; >>> ... >>> MEM <vector(4) int> [(int *)&.oacc_worker_o.44] = _40; >>> .oacc_worker_o.44._128 = workers_max_476; >>> .oacc_worker_o.44._130 = vectors_max_477; >>> .oacc_worker_o.44._132 = 0; >>> >>> The optimization is trying to transform 4 separate assignments into a single >>> memory operation. The trouble is that &o.acc_worker_o is an SImode pointer in >>> AS4 (LDS), while the memory expression appears to be in the default memory >>> space. The 'to' expression of the assignment is: >>> >>> <mem_ref 0x7ffff74c61e0 >>> type <vector_type 0x7ffff7470498 >>> type <integer_type 0x7ffff73195e8 int public SI >>> size <integer_cst 0x7ffff7318bb8 constant 32> >>> unit-size <integer_cst 0x7ffff7318bd0 constant 4> >>> align:32 warn_if_not_align:0 symtab:0 alias-set 1 canonical-type 0x7ffff73195e8 precision:32 min <integer_cst 0x7ffff7318b70 -2147483648> max <integer_cst 0x7ffff7318b88 2147483647> >>> pointer_to_this <pointer_type 0x7ffff73209d8> reference_to_this <reference_type 0x7ffff73d9d20>> >>> TI >>> size <integer_cst 0x7ffff7318ca8 constant 128> >>> unit-size <integer_cst 0x7ffff7318cc0 constant 16> >>> align:128 warn_if_not_align:0 symtab:0 alias-set 1 structural-equality nunits:4 >>> pointer_to_this <pointer_type 0x7ffff7470540>> >>> >>> arg:0 <addr_expr 0x7ffff74cdb80 >>> type <pointer_type 0x7ffff73209d8 type <integer_type 0x7ffff73195e8 int> >>> public unsigned DI >>> size <integer_cst 0x7ffff7318978 constant 64> >>> unit-size <integer_cst 0x7ffff7318990 constant 8> >>> align:64 warn_if_not_align:0 symtab:0 alias-set 2 structural-equality> >>> constant >>> arg:0 <var_decl 0x7ffff7477f30 .oacc_worker_o.44 type <record_type 0x7ffff73eb888 .oacc_ws_data_s.21 address-space-4> >>> addressable used static ignored BLK [...]/libgomp/testsuite/libgomp.oacc-fortran/parallel-dims.f90:86:0 >>> >>> size <integer_cst 0x7ffff746ce70 constant 224> >>> unit-size <integer_cst 0x7ffff746ce40 constant 28> >>> align:128 warn_if_not_align:0 >>> (mem/c:BLK (symbol_ref:SI (".oacc_worker_o.44.14") [flags 0x2] <var_decl 0x7ffff7477f30 .oacc_worker_o.44>) [9 .oacc_worker_o.44+0 S28 A128 AS4])>> >>> arg:1 <integer_cst 0x7ffff73ff078 type <pointer_type 0x7ffff73209d8> constant 0>> >>> >>> In convert_memory_address_addr_space_1: >>> >>> #ifndef POINTERS_EXTEND_UNSIGNED >>> gcc_assert (GET_MODE (x) == to_mode || GET_MODE (x) == VOIDmode); >>> return x; >>> #else /* defined(POINTERS_EXTEND_UNSIGNED) */ >>> >>> POINTERS_EXTEND_UNSIGNED is not defined, so it hits the assert. The expected >>> to_mode is DI_mode, but x is SI_mode, so the assert fires. > >> I now have a fix for this. >> >> > MEM <vector(4) int> [(int *)&.oacc_worker_o.44] = _40; >> >> The ICE occurs because the SLP vectorization pass creates the new statement >> using the type of the expression '&.oacc_worker_o.44', which is a pointer to a >> component ref in the default address space. The expand pass gets confused >> because it is handed an SImode pointer (for LDS) when it is expecting a DImode >> pointer (for flat/global space). >> >> The underlying problem is that although .oacc_worker_o is in the correct address >> space, the component ref .oacc_worker_o is not. I fixed this by propagating the >> address space of .oacc_worker_o when the component ref is created. > >> static tree >> oacc_build_component_ref (tree obj, tree field) >> { >> - tree ret = build3 (COMPONENT_REF, TREE_TYPE (field), obj, field, NULL); >> + tree field_type = TREE_TYPE (field); >> + tree obj_type = TREE_TYPE (obj); >> + if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) >> + field_type = build_qualified_type >> + (field_type, >> + KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); >> + >> + tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); >> if (TREE_THIS_VOLATILE (field)) >> TREE_THIS_VOLATILE (ret) |= 1; >> if (TREE_READONLY (field)) > > This code change has been included in the recent master branch commit > e2a58ed6dc5293602d0d168475109caa81ad0f0d "openacc: Middle-end > worker-partitioning support", which thus includes a > 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref' that is > slightly different from 'gcc/omp-low.c:omp_build_component_ref'. > > I'm confirming that with this reverted, we're seeing ICEs as follows: > > +FAIL: libgomp.oacc-fortran/gemm-2.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/gemm-2.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/gemm.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/gemm.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/optional-reduction.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/optional-reduction.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/private-variables.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/private-variables.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-1.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-1.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-5.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-5.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-6.f90 [...] -foffload=amdgcn-amdhsa -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (internal compiler error) > > +FAIL: libgomp.oacc-fortran/reduction-6.f90 [...] -foffload=amdgcn-amdhsa -O3 -g (internal compiler error) > > 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). > > Would it make sense to "Re-unify 'omp_build_component_ref' and > 'oacc_build_component_ref'", see attached? > > > 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 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Get-rid-of-gcc-omp-oacc-neuter-broadcast.cc-oacc_bui.patch --] [-- Type: text/x-diff, Size: 4626 bytes --] From 54f745023276e5025e34b2cc22530c78423a93cb Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Fri, 30 Jul 2021 16:15:25 +0200 Subject: [PATCH] Get rid of 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref' Clean-up for commit e2a58ed6dc5293602d0d168475109caa81ad0f0d "openacc: Middle-end worker-partitioning support": as of commit 2a3f9f6532bb21d8ab6f16fbe9ee603f6b1405f2 "openacc: Shared memory layout optimisation", we're no longer running into the vectorizer ICEs for '!ADDR_SPACE_GENERIC_P'. gcc/ * omp-low.cc (omp_build_component_ref): Move function... * omp-general.cc (omp_build_component_ref): ... here. Remove 'static'. * omp-general.h (omp_build_component_ref): Declare function. * omp-oacc-neuter-broadcast.cc (oacc_build_component_ref): Remove function. (build_receiver_ref, build_sender_ref): Call 'omp_build_component_ref' instead. --- gcc/omp-general.cc | 14 ++++++++++++++ gcc/omp-general.h | 2 ++ gcc/omp-low.cc | 15 --------------- gcc/omp-oacc-neuter-broadcast.cc | 26 ++------------------------ 4 files changed, 18 insertions(+), 39 deletions(-) diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc index 19f40dc0b1d..a406c578f33 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -2980,4 +2980,18 @@ oacc_get_ifn_dim_arg (const gimple *stmt) return (int) axis; } +/* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it + as appropriate. */ + +tree +omp_build_component_ref (tree obj, tree field) +{ + tree ret = build3 (COMPONENT_REF, TREE_TYPE (field), obj, field, NULL); + if (TREE_THIS_VOLATILE (field)) + TREE_THIS_VOLATILE (ret) |= 1; + if (TREE_READONLY (field)) + TREE_READONLY (ret) |= 1; + return ret; +} + #include "gt-omp-general.h" diff --git a/gcc/omp-general.h b/gcc/omp-general.h index c0cf5f014cd..7a94831e8f5 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -149,4 +149,6 @@ get_openacc_privatization_dump_flags () return l_dump_flags; } +extern tree omp_build_component_ref (tree obj, tree field); + #endif /* GCC_OMP_GENERAL_H */ diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 77176efe715..2294456b27d 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -621,21 +621,6 @@ omp_copy_decl_1 (tree var, omp_context *ctx) return omp_copy_decl_2 (var, DECL_NAME (var), TREE_TYPE (var), ctx); } -/* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it - as appropriate. */ -/* See also 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref'. */ - -static tree -omp_build_component_ref (tree obj, tree field) -{ - tree ret = build3 (COMPONENT_REF, TREE_TYPE (field), obj, field, NULL); - if (TREE_THIS_VOLATILE (field)) - TREE_THIS_VOLATILE (ret) |= 1; - if (TREE_READONLY (field)) - TREE_READONLY (ret) |= 1; - return ret; -} - /* Build tree nodes to access the field for VAR on the receiver side. */ static tree diff --git a/gcc/omp-oacc-neuter-broadcast.cc b/gcc/omp-oacc-neuter-broadcast.cc index 314161e38f5..81e3223a94c 100644 --- a/gcc/omp-oacc-neuter-broadcast.cc +++ b/gcc/omp-oacc-neuter-broadcast.cc @@ -937,35 +937,13 @@ worker_single_simple (basic_block from, basic_block to, } } -/* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it - as appropriate. */ -/* Adapted from 'gcc/omp-low.cc:omp_build_component_ref'. */ - -static tree -oacc_build_component_ref (tree obj, tree field) -{ - tree field_type = TREE_TYPE (field); - tree obj_type = TREE_TYPE (obj); - if (!ADDR_SPACE_GENERIC_P (TYPE_ADDR_SPACE (obj_type))) - field_type = build_qualified_type - (field_type, - KEEP_QUAL_ADDR_SPACE (TYPE_QUALS (obj_type))); - - tree ret = build3 (COMPONENT_REF, field_type, obj, field, NULL); - if (TREE_THIS_VOLATILE (field)) - TREE_THIS_VOLATILE (ret) |= 1; - if (TREE_READONLY (field)) - TREE_READONLY (ret) |= 1; - return ret; -} - static tree build_receiver_ref (tree var, tree receiver_decl, field_map_t *fields) { tree x = build_simple_mem_ref (receiver_decl); tree field = *fields->get (var); TREE_THIS_NOTRAP (x) = 1; - x = oacc_build_component_ref (x, field); + x = omp_build_component_ref (x, field); return x; } @@ -975,7 +953,7 @@ build_sender_ref (tree var, tree sender_decl, field_map_t *fields) if (POINTER_TYPE_P (TREE_TYPE (sender_decl))) sender_decl = build_simple_mem_ref (sender_decl); tree field = *fields->get (var); - return oacc_build_component_ref (sender_decl, field); + return omp_build_component_ref (sender_decl, field); } static int -- 2.34.1 ^ permalink raw reply [flat|nested] 15+ messages in thread
end of thread, other threads:[~2022-02-22 17:00 UTC | newest]
Thread overview: 15+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
[not found] <992c7c29-5773-45b6-6fb7-ffb71299a98f@mentor.com>
[not found] ` <f2129af9-8314-6eb6-d342-8fd725d0a028@mentor.com>
[not found] ` <ea33e1e0-71ac-1561-01e9-67d875292904@mentor.com>
[not found] ` <fe7bc7ae-6807-ff00-1b3a-e3c7ac41b723@mentor.com>
2021-08-09 14:16 ` Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref' Thomas Schwinge
2021-08-16 8:08 ` [ping] " Thomas Schwinge
2021-08-16 8:21 ` Jakub Jelinek
2021-08-19 20:13 ` Thomas Schwinge
2021-08-20 7:51 ` Richard Biener
2021-08-23 14:30 ` Thomas Schwinge
2021-08-24 7:43 ` Richard Biener
2021-08-20 14:49 ` Jakub Jelinek
2021-08-23 15:55 ` Add 'libgomp.c/address-space-1.c' (was: [ping] Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref') Thomas Schwinge
2021-08-24 10:23 ` Host and offload targets have no common meaning of address spaces " Thomas Schwinge
2021-08-24 11:43 ` Richard Biener
2021-09-03 11:42 ` Andrew Stubbs
2022-01-13 10:24 ` Host and offload targets have no common meaning of address spaces Thomas Schwinge
2021-09-10 8:03 ` Thomas Schwinge
2022-02-22 17:00 ` Get rid of 'gcc/omp-oacc-neuter-broadcast.cc:oacc_build_component_ref' (was: Re-unify 'omp_build_component_ref' and 'oacc_build_component_ref') Thomas Schwinge
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).