Hi!
On 2021-08-20T09:51:36+0200, Richard Biener <richard.guent...@gmail.com> wrote:
> On Thu, Aug 19, 2021 at 10:14 PM Thomas Schwinge
> <tho...@codesourcery.com> wrote:
>> Richard, maybe you have an opinion here, in particular about my
>> "SLP vectorizer" comment below? Please see
>> <87r1f2puss.fsf@euler.schwinge.homeip.net">http://mid.mail-archive.com/87r1f2puss.fsf@euler.schwinge.homeip.net>
>> for the full context.
>>
>> On 2021-08-16T10:21:04+0200, Jakub Jelinek <ja...@redhat.com> wrote:
>> > On Mon, Aug 16, 2021 at 10:08:42AM +0200, Thomas Schwinge wrote:
>> >> /* 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
>From eedea7a1041720ac4da9938716145c02918dd45e Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@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
>From 2c99df97a60970f9d853aacc80a1485ad6f07052 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@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
