Sandra Loosemore [Thu, 20 Aug 2020 02:18:57 +0000 (19:18 -0700)]
Annotate inner loops in "acc kernels loop" directives (C/C++).
Normally explicit loop directives in a kernels region inhibit
automatic annotation of other loops in the same nest, on the theory
that users have indicated they want manual control over that section
of code. However there seems to be an expectation in user code that
the combined "kernels loop" directive should still allow annotation of
inner loops. This patch implements this behavior for C and C++.
Change test for CUDA callback context in nvptx_free() from using
GOMP_PLUGIN_acc_thread () into checking for CUDA_ERROR_NOT_PERMITTED,
for the former only works for OpenACC, but not OpenMP offloading.
libgomp/
* plugin/plugin-nvptx.c (nvptx_free):
Change "GOMP_PLUGIN_acc_thread () == NULL" test into check of
CUDA_ERROR_NOT_PERMITTED status for cuMemGetAddressRange. Adjust
comments.
If the walk_body on the various sequences of reduction, lastprivate and/or linear
clauses needs to create a temporary variable, we should declare that variable
in that sequence rather than outside, where it would need to be privatized inside of
the construct.
2020-08-08 Jakub Jelinek <jakub@redhat.com>
PR fortran/93553
* tree-nested.c (convert_nonlocal_omp_clauses): For
OMP_CLAUSE_REDUCTION, OMP_CLAUSE_LASTPRIVATE and OMP_CLAUSE_LINEAR
save info->new_local_var_chain around walks of the clause gimple
sequences and declare_vars if needed into the sequence.
* testsuite/libgomp.c-c++-common/critical-hint-1.c: New; moved from
gcc/testsuite/c-c++-common/gomp/.
* testsuite/libgomp.c-c++-common/critical-hint-2.c: Likewise.
* testsuite/libgomp.fortran/critical-hint-1.f90: New; moved
from gcc/testsuite/gfortran.dg/gomp/.
* testsuite/libgomp.fortran/critical-hint-2.f90: Likewise.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/critical-hint-1.c: Moved to libgomp/.
* c-c++-common/gomp/critical-hint-2.c: Moved to libgomp/.
* gfortran.dg/gomp/critical-hint-1.f90: Moved to libgomp/.
* gfortran.dg/gomp/critical-hint-2.f90: Moved to libgomp/.
* c-omp.c (c_finish_omp_critical): Check for no name but
nonzero hint provided.
gcc/c/ChangeLog:
* c-parser.c (c_parser_omp_clause_hint): Require nonnegative hint clause.
(c_parser_omp_critical): Permit hint(0) clause without named critical.
(c_parser_omp_construct): Don't assert if error_mark_node is returned.
gcc/cp/ChangeLog:
* parser.c (cp_parser_omp_clause_hint): Require nonnegative hint.
(cp_parser_omp_critical): Permit hint(0) clause without named critical.
* pt.c (tsubst_expr): Re-check the latter for templates.
gcc/fortran/ChangeLog:
* openmp.c (gfc_match_omp_critical): Fix handling hints; permit
hint clause without named critical.
(resolve_omp_clauses): Require nonnegative constant integer
for the hint clause.
(gfc_resolve_omp_directive): Check for no name but
nonzero value for hint clause.
* parse.c (parse_omp_structured_block): Fix same-name check
for critical.
* trans-openmp.c (gfc_trans_omp_critical): Handle hint clause properly.
libgomp/ChangeLog:
* omp_lib.f90.in: Add omp_sync_hint_* and omp_sync_hint_kind.
* omp_lib.h.in: Likewise.
gcc/testsuite/ChangeLog:
* g++.dg/gomp/critical-3.C: Add nameless critical with hint testcase.
* c-c++-common/gomp/critical-hint-1.c: New test.
* c-c++-common/gomp/critical-hint-2.c: New test.
* gfortran.dg/gomp/critical-hint-1.f90: New test.
* gfortran.dg/gomp/critical-hint-2.f90: New test.
PR fortran/94690
* openmp.c (OMP_DISTRIBUTE_CLAUSES): Add OMP_CLAUSE_LASTPRIVATE.
(gfc_resolve_do_iterator): Skip the private handling for SIMD as
that is handled by ME code.
* trans-openmp.c (gfc_trans_omp_do): Don't add private/lastprivate
for dovar_found == 0, unless !simple.
This adds support for __sync_val_compare_and_swap and
__sync_bool_compare_and_swap for 1-byte and 2-byte long
values, which are not natively supported on nvptx.
Julian Brown [Tue, 28 Jul 2020 13:02:50 +0000 (06:02 -0700)]
[og10] openacc: Unshare reduction temporaries for GCN
The GCN backend uses tree nodes like MEM((__lds TYPE *) <constant>)
for reduction temporaries. Unlike e.g. var decls and SSA names, these
nodes cannot be shared during gimplification, but are so in some
circumstances. This is detected when appropriate --enable-checking
options are used. This patch unshares such nodes when they are reused
more than once.
2020-07-30 Julian Brown <julian@codesourcery.com>
gcc/
* config/gcn/gcn-tree.c (gcn_goacc_get_worker_red_decl): Do not
cache/share decls for reduction temporaries between invocations.
(gcn_goacc_reduction_teardown): Unshare VAR on second use.
* config/gcn/gcn.c (gcn_init_machine_status): Do not initialise
reduc_decls.
* config/gcn/gcn.h (machine_function): Remove reduc_decls cache.
preprocessor: Teach traditional about has_include [PR95889]
Traditional cpp (used by fortran) didn;t know about the new
__has_include__ implementation. Hey, since when did traditional cpp
grow __has_include__? That wasn't in knr!
Julian Brown [Thu, 18 Jun 2020 12:11:08 +0000 (05:11 -0700)]
openacc: Deep copy attach/detach should not affect reference counts
Attach and detach operations are not supposed to affect structural or
dynamic reference counts for OpenACC. Previously they did so, which led to
subtle problems in some circumstances. We can avoid reference-counting
attach/detach operations by extending and slightly repurposing the
do_detach field in target_var_desc. It is now called is_attach to better
reflect its new role.
2020-07-27 Julian Brown <julian@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
libgomp/
* libgomp.h (struct target_var_desc): Rename do_detach field to
is_attach.
* oacc-mem.c (goacc_exit_datum_1): Add assert. Don't set finalize for
GOMP_MAP_FORCE_DETACH. Update checking to use is_attach field.
(goacc_enter_data_internal): Don't affect reference counts
for attach mappings.
(goacc_exit_data_internal): Don't affect reference counts for detach
mappings.
* target.c (gomp_map_vars_existing): Don't affect reference counts for
attach mappings.
(gomp_map_vars_internal): Set renamed is_attach flag unconditionally to
mark attach mappings.
(gomp_unmap_vars_internal): Use is_attach flag to prevent affecting
reference count for attach mappings.
* testsuite/libgomp.oacc-c-c++-common/mdc-refcount-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/mdc-refcount-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/mdc-refcount-2.c: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: Mark
test as shouldfail.
* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Adjust to fail
gracefully in no-finalize mode.
Julian Brown [Thu, 2 Jul 2020 21:18:20 +0000 (14:18 -0700)]
openacc: Remove unnecessary detach finalization
The call to gomp_detach_pointer in gomp_unmap_vars_internal does not
need to force finalization, and doing so may mask mismatched pointer
attachments/detachments. This patch removes the forcing.
2020-07-16 Julian Brown <julian@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
libgomp/
* target.c (gomp_unmap_vars_internal): Remove unnecessary forcing of
finalization for detach operation.
* testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c:
New test.
coroutines: Correct frame capture of compiler temps [PR95591+4].
When a full expression contains a co_await (or co_yield), this means
that it might be suspended; which would destroy temporary variables
(on a stack for example). However the language guarantees that such
temporaries are live until the end of the expression.
In order to preserve this, we 'promote' temporaries where necessary
so that they are saved in the coroutine frame (which allows them to
remain live potentially until the frame is destroyed). In addition,
sub-expressions that produce control flow (such as TRUTH_AND/OR_IF
or COND_EXPR) must be handled specifically to ensure that await
expressions are properly expanded.
This patch corrects two mistakes in which we were (a) failing to
promote some temporaries and (b) we we failing to sequence DTORs for
the captures properly. This manifests in a number of related (but not
exact duplicate) PRs.
The revised code collects the actions into one place and maps all the
control flow into one form - a benefit of this is that co_returns are
now expanded earlier (which provides an opportunity to address PR95517
in some future patch).
We replace a statement that contains await expression(s) with a bind
scope that has a variable list describing the temporaries that have
been 'promoted' and a statement list that contains a series of cleanup
expressions for each of those. Where we encounter nested conditional
expressions, these are wrapped in a try-finally block with a guard var
for each sub-expression variable that needs a DTOR. The guards are all
declared and initialized to false before the first conditional sub-
expression. The 'finally' block contains a series of if blocks (one
per guard variable) enclosing the relevant DTOR.
Variables listed in a bind scope in this manner are automatically moved
to a coroutine frame version by existing code (so we re-use that rather
than having a separate mechanism).
gcc/cp/ChangeLog:
PR c++/95591
PR c++/95599
PR c++/95823
PR c++/95824
PR c++/95895
* coroutines.cc (struct coro_ret_data): Delete.
(coro_maybe_expand_co_return): Delete.
(co_return_expander): Delete.
(expand_co_returns): Delete.
(co_await_find_in_subtree): Remove unused name.
(build_actor_fn): Remove unused parm, remove handling
for co_return expansion.
(register_await_info): Demote duplicate info message to a
warning.
(coro_make_frame_entry): Move closer to use site.
(struct susp_frame_data): Add fields for final suspend label
and a flag to indicate await expressions with initializers.
(captures_temporary): Delete.
(register_awaits): Remove unused code, update comments.
(find_any_await): New.
(tmp_target_expr_p): New.
(struct interesting): New.
(find_interesting_subtree): New.
(struct var_nest_node): New.
(flatten_await_stmt): New.
(handle_nested_conditionals): New.
(process_conditional): New.
(replace_statement_captures): Rename to...
(maybe_promote_temps): ... this.
(maybe_promote_captured_temps): Delete.
(analyze_expression_awaits): Check for await expressions with
initializers. Simplify handling for truth-and/or-if.
(expand_one_truth_if): Simplify (map cases that need expansion
to COND_EXPR).
(await_statement_walker): Handle CO_RETURN_EXPR. Simplify the
handling for truth-and/or-if expressions.
(register_local_var_uses): Ensure that we create names in the
implementation namespace.
(morph_fn_to_coro): Add final suspend label to suspend frame
callback data and remove it from the build_actor_fn call.
gcc/testsuite/ChangeLog:
PR c++/95591
PR c++/95599
PR c++/95823
PR c++/95824
PR c++/95895
* g++.dg/coroutines/pr95591.C: New test.
* g++.dg/coroutines/pr95599.C: New test.
* g++.dg/coroutines/pr95823.C: New test.
* g++.dg/coroutines/pr95824.C: New test.
Fix c-c++-common/goacc/routine-4.c and c-c++-common/goacc/routine-4-extern.c testcases
'Various OpenACC reduction enhancements - FE changes' (commit 6b3e1f7f05cd360bbd356b3f78511aa2ec3f40c3) introduced checks for gang
reductions on orphan loops. The checks triggered in the routine-4.c
and routine-4-extern.c testcases, requiring changes that effectively
rendered them useless as test cases.
This patch restores the original intent of the test cases, by restoring
the original tests and removing the orphan loop reductions that were
triggering the new check.
Martin Liska [Mon, 27 Jul 2020 10:30:24 +0000 (12:30 +0200)]
expr: build string_constant only for a char type
gcc/ChangeLog:
PR tree-optimization/96058
* expr.c (string_constant): Build string_constant only
for a type that has same precision as char_type_node
and is an integral type.
Jakub Jelinek [Tue, 28 Jul 2020 09:08:29 +0000 (11:08 +0200)]
expander: Fix ICE in maybe_warn_rdwr_sizes [PR96335]
The following testcase ICEs in maybe_warn_rdwr_sizes. The problem is that
the caller uses its fndecl and fntype variables to fill up rdwr_map, and
the fntype in that case is a prototype with the access attribute and all
the checks needed for that performed. But the maybe_warn_rdwr_sizes
function tries to rediscover fndecl/fntype itself and does it differently
from how the caller did (for fndecl get_callee_fndecl and fntype from that
FUNCTION_DECL, otherwise sets fntype to CALL_EXPR_FN's type).
On the testcase, get_callee_fndecl does find a FUNCTION_DECL because
it does STRIP_NOPS in between.
Instead of trying to rediscover those, this patch just passes them down,
like is done in several other functions.
2020-07-28 Jakub Jelinek <jakub@redhat.com>
PR middle-end/96335
* calls.c (maybe_warn_rdwr_sizes): Add FNDECL and FNTYPE arguments,
instead of trying to rediscover them in the body.
(initialize_argument_information): Adjust caller.
Martin Liska [Fri, 24 Jul 2020 12:33:27 +0000 (14:33 +0200)]
Use vec::reserve before vec_safe_grow_cleared is called
gcc/ChangeLog:
PR lto/45375
* symbol-summary.h: Call vec_safe_reserve before grow is called
in order to grow to a reasonable size.
* vec.h (vec_safe_reserve): Add missing function for vl_ptr
type.
Mark Eggleston [Thu, 11 Jun 2020 10:05:40 +0000 (11:05 +0100)]
Fortran : ICE in gfc_check_pointer_assign PR95612
Output an error if the right hand value is a zero sized array or
does not have a symbol tree otherwise continue checking.
2020-07-27 Steven G. Kargl <kargl@gcc.gnu.org>
gcc/fortran/
PR fortran/95612
* expr.c (gfc_check_pointer_assigb): Output an error if
rvalue is a zero sized array or output an error if rvalue
doesn't have a symbol tree.
2020-07-27 Mark Eggleston <markeggleston@gcc.gnu.org>
gcc/testsuite/
PR fortran/95612
* gfortran.dg/pr95612.f90: New test.
PR 93567, G edit descriptor uses E instead of F editing in rounding mode UP.
The switch between FMT_E and FMT_F is based on the absolute value.
Set r=0 for rounding toward zero and r = 1 otherwise.
If (exp_d - m) == 1 there is no rounding needed.
libgfortran/ChangeLog:
PR fortran/93567
* io/write_float.def (determine_en_precision): Fix switch between
FMT_E and FMT_F.
gcc/testsuite/ChangeLog:
PR fortran/93567
* gfortran.dg/round_3.f08: Add test cases.
XFAIL tests in gfortran.dg/goacc/loop-2-kernels.f95
The C-equivalent version of the test (c-c++-common/goacc/loop-2-kernels.c)
has these tests XFAILed in the commit 'Make new OpenACC kernels conversion
the default; adjust and add tests' (commit 757f56ddc43fd80bb8740222ec352111b26d66e9), so the Fortran version should
be XFAILed too.
sparc/sparc64: use crtendS.o for default-pie executables [PR96190]
In --enable-default-pie mode compiler should switch from
using crtend.o to crtendS.o. On sparc it is especially important
because crtend.o contains PIC-unfriendly code.
We use GNU_USER_TARGET_ENDFILE_SPEC as a baseline spec to get
crtendS.o instead of crtend.o in !no-pie mode.
gcc:
2020-07-14 Sergei Trofimovich <siarheit@google.com>
PR target/96190
* config/sparc/linux.h (ENDFILE_SPEC): Use GNU_USER_TARGET_ENDFILE_SPEC
to get crtendS.o for !no-pie mode.
* config/sparc/linux64.h (ENDFILE_SPEC): Ditto.
Harald Anlauf [Mon, 6 Jul 2020 16:58:23 +0000 (18:58 +0200)]
PR fortran/95980 - ICE on using sync images with -fcheck=bounds
In SELECT TYPE, the argument may be an incorrectly specified unlimited
polymorphic variable. Avoid a NULL pointer dereference for clean error
recovery.
gcc/fortran/
PR fortran/95980
* match.c (copy_ts_from_selector_to_associate, build_class_sym):
Distinguish between unlimited polymorphic and ordinary variables
to avoid NULL pointer dereference.
* resolve.c (resolve_select_type):
Distinguish between unlimited polymorphic and ordinary variables
to avoid NULL pointer dereference.
Peter Bergner [Wed, 22 Jul 2020 16:44:35 +0000 (11:44 -0500)]
rs6000: __builtin_mma_disassemble_acc() doesn't store elements correctly in LE mode
PR96236 shows a problem where we don't correctly store our 512-bit accumulators
correctly in little-endian mode. The patch below detects when we're doing a
little-endian memory access and stores to the correct memory locations.
Thomas Koenig [Sun, 19 Jul 2020 15:27:45 +0000 (17:27 +0200)]
Always use name from c_interop_kinds_table for -fc-prototypes.
When a user specified a KIND that was a parameter taking the value
of an iso_c_binding KIND, the code used the name of that parameter
to look up the type name. Corrected by always looking it up in
the table of C interop kinds (which was previously done for
non-C-interop types, anyway).
gcc/fortran/ChangeLog:
PR fortran/96220
* dump-parse-tree.c (get_c_type_name): Always use the entries from
c_interop_kinds_table to find the correct C type.
Thomas Koenig [Thu, 23 Jul 2020 15:56:09 +0000 (17:56 +0200)]
Fix handling of implicit_pure by checking if non-pure procedures are called.
Procedures are marked as implicit_pure if they fulfill the criteria of
pure procedures. In this case, a procedure was not marked as not being
implicit_pure which called another procedure, which had not yet been
marked as not being implicit_impure.
Fixed by iterating over all procedures, setting callers of procedures
which are non-pure and non-implicit_pure as non-implicit_pure and
doing this until no more procedure has been changed.
Kito Cheng [Wed, 22 Jul 2020 06:50:40 +0000 (14:50 +0800)]
PR target/96260 - KASAN should work even back-end not porting anything.
- Most KASAN function don't need any porting anything in back-end
except asan stack protection.
- However kernel will given shadow offset when enable asan stack
protection, so eveything in KASAN can work if shadow offset is given.
- Verified with x86 and risc-v.
- Verified with RISC-V linux kernel.
gcc/ChangeLog:
PR target/96260
* asan.c (asan_shadow_offset_set_p): New.
* asan.h (asan_shadow_offset_set_p): Ditto.
* toplev.c (process_options): Allow -fsanitize=kernel-address
even TARGET_ASAN_SHADOW_OFFSET not implemented, only check when
asan stack protection is enabled.
Frederik Harwath [Mon, 20 Jul 2020 09:24:21 +0000 (11:24 +0200)]
libgomp, Fortran: Fix OpenACC "gang reduction on an orphan loop" error message
According to the OpenACC standard version 2.5 and later, reductions on
orphaned gang loops are explicitly disallowed (cf. section "Changes
from Version 2.0 to 2.5"). A loop is "orphaned" if it is not
lexically contained in a compute construct (cf. section "Loop
construct" of the OpenACC standard), i.e. in either a "parallel", a
"serial", or a "kernels" construct.
This commit fixes the check for reductions on orphaned gang loops in
the Fortran frontend which (in contrast to the C, C++ frontends)
erroneously rejects reductions on gang loops that are contained in
"kernels" constructs.
2020-07-20 Frederik Harwath <frederik@codesourcery.com>
gcc/fortran/
* openmp.c (oacc_is_parallel_or_serial): Removed function.
(oacc_is_kernels): New function.
(oacc_is_compute_construct): New function.
(resolve_oacc_loop_blocks): Use "oacc_is_compute_construct"
instead of "oacc_is_parallel_or_serial" for checking that a
loop is not orphaned.
gcc/testsuite/
* gfortran.dg/goacc/orphan-reductions-2.f90: New test
verifying that the "gang reduction on an orphan loop" error message
is not emitted for non-orphaned loops.
* c-c++-common/goacc/orphan-reductions-2.c: Likewise for C and C++.
libgomp: Fix hang when profiling OpenACC programs with CUDA 9.0 nvprof
The version of nvprof in CUDA 9.0 causes a hang when used to profile an
OpenACC program. This is because it calls acc_get_device_type from
a callback called during device initialization, which then attempts
to acquire acc_device_lock while it is already taken, resulting in
deadlock. This works around the issue by returning acc_device_none
from acc_get_device_type without attempting to acquire the lock when
initialization has not completed yet.
This is a port of commit a1c022d1b9a43f85f0c451b6422fd095a704fe96 from
releases/gcc-10. Most of the patch is already in this branch, so this
commit only contains changes to existing code.
2020-07-17 Thomas Schwinge <thomas@codesourcery.com>
Kwok Cheung Yeung <kcy@codesourcery.com>
libgomp/
* oacc-init.c (acc_init_1): Move setting of acc_init_state to initializing
to the beginning of the function.
* libgomp.texi (acc_get_device_type): Update documentation.
(Implementation Status and Implementation-Defined Behavior): Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_prof-init-2.c: New.
Andrew Stubbs [Thu, 9 Jul 2020 21:48:39 +0000 (22:48 +0100)]
amdgcn: Handle early debug info in mkoffload
Forward the early debug information from the input LTO file to the output
HSACO file, in the same way lto-wrapper does. This is a little more
complicated, however, because the ELF file containing the debug needs to be
converted from x86_64 to amdgcn, and because the offloaded code will have less
content than the host program the debug info describes.
* config/gcn/mkoffload.c: Include simple-object.h and elf.h.
(EM_AMDGPU): New macro.
(ELFOSABI_AMDGPU_HSA): New macro.
(ELFABIVERSION_AMDGPU_HSA): New macro.
(EF_AMDGPU_MACH_AMDGCN_GFX803): New macro.
(EF_AMDGPU_MACH_AMDGCN_GFX900): New macro.
(EF_AMDGPU_MACH_AMDGCN_GFX906): New macro.
(R_AMDGPU_NONE): New macro.
(R_AMDGPU_ABS32_LO): New macro.
(R_AMDGPU_ABS32_HI): New macro.
(R_AMDGPU_ABS64): New macro.
(R_AMDGPU_REL32): New macro.
(R_AMDGPU_REL64): New macro.
(R_AMDGPU_ABS32): New macro.
(R_AMDGPU_GOTPCREL): New macro.
(R_AMDGPU_GOTPCREL32_LO): New macro.
(R_AMDGPU_GOTPCREL32_HI): New macro.
(R_AMDGPU_REL32_LO): New macro.
(R_AMDGPU_REL32_HI): New macro.
(reserved): New macro.
(R_AMDGPU_RELATIVE64): New macro.
(gcn_s1_name): Delete global variable.
(gcn_s2_name): Delete global variable.
(gcn_o_name): Delete global variable.
(gcn_cfile_name): Delete global variable.
(files_to_cleanup): New global variable.
(offload_abi): New global variable.
(tool_cleanup): Use files_to_cleanup, not explicit list.
(copy_early_debug_info): New function.
(main): New local variables gcn_s1_name, gcn_s2_name, gcn_o_name,
gcn_cfile_name.
Create files_to_cleanup obstack.
Recognize -march options.
Copy early debug info from input .o files.
Andrew Stubbs [Fri, 26 Jun 2020 16:10:55 +0000 (17:10 +0100)]
amdgcn: Support basic DWARF
This is enough DWARF support for "-O0 -g" to work OK, within a single frame,
using the new rocgdb debugger from AMD. Debugging with optimization enabled
also works when the values are located in SGPRs or memory.
Scalars in VGPRs are problematic, EXEC_HI and VCC_HI are unmappable, and
CFI remains unimplemented.
* config/gcn/gcn-hsa.h (DBX_REGISTER_NUMBER): New macro.
* config/gcn/gcn-protos.h (gcn_dwarf_register_number): New prototype.
* config/gcn/gcn.c (gcn_expand_prologue): Add RTX_FRAME_RELATED_P
and REG_FRAME_RELATED_EXPR to stack and frame pointer adjustments.
(gcn_dwarf_register_number): New function.
(gcn_dwarf_register_span): New function.
(TARGET_DWARF_REGISTER_SPAN): New hook macro.
libgomp/
* plugin/plugin-gcn.c (parse_target_attributes): Automatically set
the number of teams and threads if necessary.
(gcn_exec): Automatically set the number of gangs and workers if
necessary.
Co-Authored-By: Andrew Stubbs <ams@codesourcery.com>
This is an optimisation for middle-end worker-partitioning support (used
to support multiple workers on AMD GCN). At present, barriers may be
emitted in cases where they aren't needed and cannot be optimised away.
This patch stops the extraneous barriers from being emitted in the
first place.
One exception to the above (where the barrier is still needed) is for
predicated blocks of code that perform a write to gang-private shared
memory from one worker. We must execute a barrier before other workers
read that shared memory location.
2020-07-15 Julian Brown <julian@codesourcery.com>
gcc/
* config/gcn/gcn.c (gimple.h): Include.
(gcn_fork_join): Emit barrier for worker-level joins.
* omp-sese.c (find_local_vars_to_propagate): Add writes_gangprivate
bitmap parameter. Set bit for blocks containing gang-private variable
writes.
(worker_single_simple): Don't emit barrier after predicated block.
(worker_single_copy): Don't emit barrier if we're not broadcasting
anything and the block contains no gang-private writes.
(neuter_worker_single): Don't predicate blocks that only contain NOPs
or internal marker functions. Pass has_gangprivate_write argument to
worker_single_copy.
(oacc_do_neutering): Add writes_gangprivate bitmap handling.
Julian Brown [Mon, 10 Feb 2020 20:26:57 +0000 (12:26 -0800)]
amdgcn: Add waitcnt after LDS write instructions
Data-share write (ds_write) instructions do not necessarily complete
the write to LDS immediately. When a write completes, LGKM_CNT is
decremented. For now, we wait until LGKM_CNT reaches zero after each
ds_write instruction.
This fixes a race condition in the case where LDS is read immediately
after being written. This can happen with broadcast operations.
Julian Brown [Mon, 27 Jan 2020 14:14:05 +0000 (06:14 -0800)]
openacc: Shared memory layout optimisation
This patch implements an algorithm to lay out local data-share (LDS)
space. It currently works for AMD GCN. At the moment, LDS is used for
three things:
After the patch is applied, (2) and (3) are placed at preallocated
locations in LDS, and (1) continues to be handled by the backend (as it
is at present prior to this patch being applied). LDS now looks like this:
+--------------+ (gang local size + 1024, = 1536)
| free space |
| ... |
| - - - - - - -|
| worker bcast |
+--------------+
| reductions |
+--------------+ <<< -mgang-local-size=<number> (def. 512)
| gang private |
| vars |
+--------------+ (32)
| low LDS vars |
+--------------+ LDS base
So, gang-private space is fixed at a constant amount at compile time
(which can be increased with a command-line switch if necessary
for some given code). The layout algorithm takes out a slice of the
remainder of usable space for reduction vars, and uses the rest for
worker partitioning.
The partitioning algorithm works as follows.
1. An "adjacency" set is built up for each basic block that might
do a broadcast. This is calculated by starting at each such block,
and doing a recursive DFS walk over successors to find the next
block (or blocks) that *also* does a broadcast
(dfs_broadcast_reachable_1).
2. The adjacency set is inverted to get adjacent predecessor blocks also.
3. Blocks that will perform a broadcast are sorted by size of that
broadcast: the biggest blocks are handled first.
4. A splay tree structure is used to calculate the spans of LDS memory
that are already allocated by the blocks adjacent to this one
(merge_ranges{,_1}.
5. The current block's broadcast space is allocated from the first free
span not allocated in the splay tree structure calculated above
(first_fit_range). This seems to work quite nicely and efficiently
with the splay tree structure.
6. Continue with the next-biggest broadcast block until we're done.
In this way, "adjacent" broadcasts will not use the same piece of
LDS memory.
2020-07-15 Julian Brown <julian@codesourcery.com>
gcc/
* config/gcn/gcn-protos.h (gcn_goacc_adjust_private_decl): Update
prototype.
* config/gcn/gcn-tree.c (gcn_goacc_get_worker_red_decl): Use
preallocated block of LDS memory.
(gcn_goacc_create_propagation_record): Add OFFSET parameter, and return
temporary LDS space at that offset. Return pointer in "sender" case.
(gcn_goacc_adjust_private_decl): Return var.
* config/gcn/gcn.c (acc_lds_size, gangprivate_hwm, lds_allocs): New
global vars.
(ACC_LDS_SIZE): Define as acc_lds_size.
(gcn_init_machine_status): Don't initialise lds_allocated and
lds_allocs fields of machine function struct.
(gcn_option_override): Handle default size for gang-private variables
and -mgang-local-size option.
(gcn_expand_prologue): Use LDS_SIZE instead of LDS_SIZE-1 when
initialising M0_REG.
(gcn_shared_mem_layout): New function.
(gcn_print_lds_decl): Update comment. Use global lds_allocs map and
gangprivate_hwm variable.
(TARGET_GOACC_SHARED_MEM_LAYOUT): Define target hook.
* config/gcn/gcn.h (machine_function): Remove lds_allocated,
lds_allocs. Add reduction_base, reduction_limit.
* config/gcn/gcn.opt (gang_local_size_opt): New global.
(mgang-local-size=): New option.
* config/nvptx/nvptx.c (nvptx_goacc_adjust_private_decl): Return var.
* doc/tm.texi.in (TARGET_GOACC_SHARED_MEM_LAYOUT): Place documentation
hook.
* doc/tm.texi: Regenerate.
* omp-offload.c (addr_expr_rewrite_info): Change adjusted_vars to a
hash_map.
(rewrite_addr_expr): Rewrite VAR_DECLs also.
(default_goacc_create_propagation_record): Add OFFSET parameter.
(execute_oacc_gimple_workers): Calculate per-function reduction
temporary and private-variable size. Call OpenACC shared_mem_layout
hook. Move num_workers==1 handling here.
(execute_oacc_device_lower): Fix for adjusted_vars being a hash_map
rather than a hash_set.
(pass_oacc_gimple_workers::gate): Remove num_workers==1 handling from
here. Enable pass for all OpenACC routines in order to call shared
memory-layout hook.
* omp-sese.c (targhooks.h, diagnostic-core.h): Add includes.
(build_sender_ref): Handle sender_decl being pointer.
(worker_single_copy): Add PLACEMENT and ISOLATE_BROADCASTS parameters.
Pass placement argument to create_propagation_record hook invocations.
Handle sender_decl being pointer and isolate_broadcasts inserting extra
barriers.
(blk_offset_map_t): Add typedef.
(neuter_worker_single): Add BLK_OFFSET_MAP parameter. Pass
preallocated range to worker_single_copy call.
(dfs_broadcast_reachable_1): New function.
(idx_decl_pair_t, used_range_vec_t): New typedefs.
(sort_size_descending): New function.
(addr_range): New class.
(splay_tree_compare_addr_range, splay_tree_free_key, first_fit_range,
merge_ranges_1, merge_ranges): New functions.
(oacc_do_neutering): Add BOUNDS_LO, BOUNDS_HI parameters. Arrange
layout of shared memory for broadcast operations.
* omp-sese.h (oacc_do_neutering): Update prototype.
* target.def (adjust_private_decl): Change return type to tree.
(create_propagation_record): Add OFFSET parameter.
(shared_mem_layout): New hook.
* targhooks.h (default_goacc_create_propagation_record): Update
prototype.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/broadcast-many.c: New test.
Julian Brown [Mon, 20 Jan 2020 19:42:28 +0000 (11:42 -0800)]
openacc: Turn off worker partitioning if num_workers==1
This patch turns off the middle-end worker-partitioning support if the
number of workers for an outlined offload function is one. In that case,
we do not need to perform the broadcasting/neutering code transformation.
2020-07-15 Julian Brown <julian@codesourcery.com>
gcc/
* omp-offload.c (pass_oacc_gimple_workers::gate): Disable worker
partitioning if num_workers is 1.
Julian Brown [Fri, 26 Jun 2020 16:07:58 +0000 (09:07 -0700)]
openacc: Fix mkoffload SGPR/VGPR count parsing for HSACO v3
If an offload kernel uses a large number of VGPRs, AMD GCN hardware may
need to limit the number of threads/workers launched for that kernel.
The number of SGPRs/VGPRs in use is detected by mkoffload and recorded in
the processed output. The patterns emitted detailing SGPR/VGPR occupancy
changed between HSACO v2 and v3 though, so this patch updates parsing
to account for that.
2020-07-15 Julian Brown <julian@codesourcery.com>
gcc/
* config/gcn/mkoffload.c (process_asm): Initialise regcount. Update
scanning for SGPR/VGPR usage for HSACO v3.
Julian Brown [Thu, 25 Jun 2020 14:40:53 +0000 (07:40 -0700)]
openacc: Fix race condition in Fortran loop collapse tests
The gangs participating in a gang-partitioned loop are not all guaranteed
to complete before some given gang continues to execute beyond that loop.
This means that two existing test cases contain a race condition,
because a loop that may be gang-partitioned is followed immediately by
another loop. The fix is to place the loops in separate parallel regions.
The trunk versions of the patch made GNU and Advanced SIMD vectors
distinct (but inter-convertible) in all cases. However, the
traditional behaviour is that the types are distinct in template
arguments but not otherwise.
Following a suggestion from Jason, this patch puts the check
for different vector types under comparing_specializations.
In order to keep the backport as simple as possible, the patch
hard-codes the name of the attribute in the frontend rather than
adding a new branch-only target hook.
I didn't find a test that tripped the assert on the branch,
even with the --param in the PR, so instead I tested this by
forcing the hash function to only hash the tree code. That made
the static assertion in the test fail without the patch but pass
with it.
This means that the tests pass for unmodified sources even
without the patch (unless you're very unlucky).
gcc/
PR target/95726
* config/aarch64/aarch64.c (aarch64_attribute_table): Add
"Advanced SIMD type".
* config/aarch64/aarch64-builtins.c: Include stringpool.h and
attribs.h.
(aarch64_init_simd_builtin_types): Add an "Advanced SIMD type"
attribute to each Advanced SIMD type.
* config/arm/arm.c (arm_attribute_table): Add "Advanced SIMD type".
* config/arm/arm-builtins.c: Include stringpool.h and attribs.h.
(arm_init_simd_builtin_types): Add an "Advanced SIMD type"
attribute to each Advanced SIMD type.
gcc/cp/
PR target/95726
* typeck.c (structural_comptypes): When comparing template
specializations, differentiate between vectors that have and
do not have an "Advanced SIMD type" attribute.
gcc/testsuite/
PR target/95726
* g++.target/aarch64/pr95726.C: New test.
* g++.target/arm/pr95726.C: Likewise.
Jakub Jelinek [Wed, 15 Jul 2020 09:34:44 +0000 (11:34 +0200)]
fix _mm512_{,mask_}cmp*_p[ds]_mask at -O0 [PR96174]
The _mm512_{,mask_}cmp_p[ds]_mask and also _mm_{,mask_}cmp_s[ds]_mask
intrinsics have an argument which must have a constant passed to it
and so use an inline version only for ifdef __OPTIMIZE__ and have
a #define for -O0. But the _mm512_{,mask_}cmp*_p[ds]_mask intrinsics
don't need a constant argument, they are essentially the first
set with the constant added to them implicitly based on the comparison
name, and so there is no #define version for them (correctly).
But their inline versions are defined in between the first and s[ds]
set and so inside of ifdef __OPTIMIZE__, which means that with -O0
they aren't defined at all.
This patch fixes that by moving those after the #ifdef __OPTIMIZE #else
use #define #endif block.
openmp: Change omp_atv_default value and rename omp_atv_sequential to omp_atv_serialized.
While this is an OpenMP 5.1 change, it is undesirable to let people use different
values and then deal with ABI backwards compatibility in a year or two.
2020-07-09 Jakub Jelinek <jakub@redhat.com>
* omp.h.in (omp_alloctrait_value_t): Change omp_atv_default from
2 to -1. Add omp_atv_serialized and define omp_atv_sequential using
it. Remove __omp_alloctrait_value_max__.
* allocator.c (omp_init_allocator): Handle omp_atv_default for
omp_atk_alignment and omp_atk_pool_size.
projects / thirdparty / gcc.git / log
