When using a compiler build with:
...
+#define PTX_DEFAULT_VECTOR_LENGTH PTX_CTA_SIZE
...
consider a test-case:
...
int
main (void)
{
#pragma acc parallel vector_length (64)
#pragma acc loop worker
for (unsigned int i = 0; i < 32; i++)
#pragma acc loop vector
for (unsigned int j = 0; j < 64; j++)
;
return 0;
}
...
If num_workers is 16, either because:
- we add a "num_workers (16)" clause on the parallel directive, or
- we set "GOMP_OPENACC_DIM=:16:", or
- the libgomp plugin chooses 16 num_workers
we run into an illegal instruction at runtime, because a bar.sync instruction
tries to use a barrier 16. The instruction is illegal, because ptx supports
only 16 barriers per CTA, and the valid range is 0..15.
The problem is that with a warp-multiple vector length, we use a code generation
scheme with a per-worker barrier. And because barrier zero is reserved for
per-cta barrier, only the remaining 15 barriers can be used as per-worker
barrier, and consequently we can't use num_workers larger than 15.
This problem occurs only for vector_length 64. For vector_length 32, we use a
different code generation scheme, and for vector_length >= 96, the maximum
num_workers is not big enough not to trigger this problem.
Also, this problem only occurs for num_workers 16. As explained above,
num_workers 15 is safe to use, and 16 is already the maximum num_workers for
vector_length 64.
This patch fixes the problem in both the compiler (handling "num_workers (16)")
and in the libgomp nvptx plugin (with and without "GOMP_OPENACC_DIM=:16:").
2019-01-11 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.c (PTX_CTA_NUM_BARRIERS, PTX_PER_CTA_BARRIER)
(PTX_NUM_PER_CTA_BARRIER, PTX_FIRST_PER_WORKER_BARRIER)
(PTX_NUM_PER_WORKER_BARRIERS): Define.
(nvptx_apply_dim_limits): Prevent vector_length 64 and
num_workers 16.
* plugin/plugin-nvptx.c (nvptx_exec): Prevent vector_length 64 and
num_workers 16.
From-SVN: r267838
Move the defition of PTX_CTA_SIZE up in nvptx.c.
2019-01-11 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.c (PTX_CTA_SIZE): Move up.
From-SVN: r267837
Before the commit "[libgomp, testsuite, openacc] Don't use const int for
dimensions", the "const int" construct was used to set launch dimensions in
reductions-[1-5].c. In the case of -xc -O0, the const int is implemented as a
variable by the C front-end. Consequently, the nvptx back-end generated
warnings that vector_length was overridden to be hard-coded, rather than left to
be set at runtime. The test-cases silenced these warnings by switching off all
warnings in the accelerator compiler using "-foffload=-w".
Given that no warnings occur anymore, remove the "-foffload=-w" setting.
2019-01-11 Tom de Vries <tdevries@suse.de>
* testsuite/libgomp.oacc-c-c++-common/reduction-1.c: Remove
-foffload=-w.
* testsuite/libgomp.oacc-c-c++-common/reduction-2.c: Same.
* testsuite/libgomp.oacc-c-c++-common/reduction-3.c: Same.
* testsuite/libgomp.oacc-c-c++-common/reduction-4.c: Same.
* testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Same.
From-SVN: r267836
Add a test-case that tests the "insufficient resources" fatal in the nvptx
libgomp plugin.
2019-01-11 Tom de Vries <tdevries@suse.de>
* testsuite/libgomp.oacc-c-c++-common/insufficient-resources.c: New
test.
From-SVN: r267835
For 64-bit these should not be emitted without suffix in AT&T mode (as
being ambiguous that way); the suffixes are benign for 32-bit. For
consistency also omit the suffix in Intel mode for {,V}CVTSI2SxQ.
The omission has originally (prior to rev 260691) lead to wrong code
being generated for the 64-bit unsigned-to-float/double conversions (as
gas guesses an L suffix instead of the required Q one when the operand
is in memory). In all remaining cases (being changed here) the omission
would "just" lead to warnings with future gas versions.
As a result, arrange to check for the L suffixes in 32-bit test cases.
In order for related test cases to actually test what they're supposed
to test, add (seemingly unrelated) a few empty "asm volatile()".
Presumably there are more where constant propagation voids the intended
effect of the tests, but these are ones helping make sure the assembler
actually still assembles correctly the output after the changes here.
From-SVN: r267833
The recent changes to support operator<<(nullptr_t) changed the glob
patterns for existing operator<<(T) overloads, but did so incorrectly so
they still matched the new symbols. That broke Solaris bootstrap. This
patch replaces each of the existing globs by two more precise ones,
which match the old symbols but not the new ones.
* config/abi/pre/gnu.ver (GLIBCXX_3.4): Correct recent changes to
basic_ostream::operator<< patterns.
From-SVN: r267832
2019-01-11 Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/59345
* trans-array.c (gfc_conv_parameter_array): Temporary
arrays generated for expressions do not need to be repacked.
2019-01-11 Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/59345
* gfortran.dg/internal_pack_16.f90: New test.
From-SVN: r267829
2019-01-10 Vladimir Makarov <vmakarov@redhat.com>
PR rtl-optimization/87305
* lra-assigns.c
(setup_live_pseudos_and_spill_after_risky_transforms): Check
allocation for big endian pseudos used as paradoxical subregs and
spill them if it is wrong.
* lra-constraints.c (lra_constraints): Add a comment.
2019-01-10 Vladimir Makarov <vmakarov@redhat.com>
PR rtl-optimization/87305
* gcc.target/aarch64/pr87305.c: New.
From-SVN: r267823
2019-01-10 Richard Biener <rguenther@suse.de>
PR tree-optimization/88792
* tree-ssa-pre.c (get_representative_for): Do not return a
value-number here.
* gcc.dg/torture/pr88792.c: New testcase.
From-SVN: r267821
2019-01-10 Steven G. Kargl <kargl@gcc.gnu.org>
PR fortran/86322
* decl.c (top_var_list): Set locus of expr.
(gfc_match_data): Detect pointer on non-rightmost part-refs.
2019-01-10 Steven G. Kargl <kargl@gcc.gnu.org>
PR fortran/86322
* gfortran.dg/pr86322_1.f90: New test.
* gfortran.dg/pr86322_2.f90: Ditto.
* gfortran.dg/pr86322_3.f90: Ditto.
From-SVN: r267820
Currently Return Address Signing is only supported in lp64. Thus the
tests that I added recently (that enables return address signing by the
mbranch-protection=standard option), should also be exempted from testing in
ilp32. This patch adds the needed dg-require-effective-target directive in the
tests.
*** gcc/testsuite/ChangeLog ***
2019-01-10 Sudakshina Das <sudi.das@arm.com>
* gcc.target/aarch64/bti-1.c: Exempt for ilp32.
* gcc.target/aarch64/bti-2.c: Likewise.
* gcc.target/aarch64/bti-3.c: Likewise.
Committed as obvious.
From-SVN: r267818
The C++2a draft specifies the value 201811L for this, but as an
extension we return the number of elements erased. This is expected to
be standardised, so the macro has the value 201900L until a proper value
is specified in the draft.
* include/bits/erase_if.h: Define __cpp_lib_erase_if.
* include/std/deque: Likewise.
* include/std/forward_list: Likewise.
* include/std/list: Likewise.
* include/std/string: Likewise.
* include/std/vector: Likewise.
* include/std/version: Likewise.
* testsuite/21_strings/basic_string/erasure.cc: Test macro.
* testsuite/23_containers/deque/erasure.cc: Likewise.
* testsuite/23_containers/forward_list/erasure.cc: Likewise.
* testsuite/23_containers/list/erasure.cc: Likewise.
* testsuite/23_containers/map/erasure.cc: Likewise.
* testsuite/23_containers/set/erasure.cc: Likewise.
* testsuite/23_containers/unordered_map/erasure.cc: Likewise.
* testsuite/23_containers/unordered_set/erasure.cc: Likewise.
* testsuite/23_containers/vector/erasure.cc: Likewise.
From-SVN: r267810
The AI_NUMERICSERV constant is missing from old Darwin systems, so only
use it if it's supported.
* include/experimental/internet [AI_NUMERICSERV]
(resolver_base::numeric_service): Define conditionally.
* testsuite/experimental/net/internet/resolver/base.cc: Test it
conditionally.
* testsuite/experimental/net/internet/resolver/ops/lookup.cc:
Likewise.
From-SVN: r267809
2019-01-10 Ville Voutilainen <ville.voutilainen@gmail.com>
Jonathan Wakely <jwakely@redhat.com>
Implement LWG 2221
* config/abi/pre/gnu.ver (GLIBCXX_3.4): Tighten patterns.
(GLIBCXX_3.4.26): Add new exports.
* include/Makefile.am: Add ostream-inst.cc. Move string-inst.cc to
correct list of sources.
* include/Makefile.in: Regenerate.
* include/std/ostream (operator<<(nullptr_t)): New member function.
* src/c++17/ostream-inst.cc: New file.
* testsuite/27_io/basic_ostream/inserters_other/char/lwg2221.cc: New
test.
Co-Authored-By: Jonathan Wakely <jwakely@redhat.com>
From-SVN: r267808
This allows to use unified asm syntax when compiling for the
ARM instruction. This matches documentation and seems what the
initial patch was intended doing when the flag got added.
2019-01-10 Stefan Agner <stefan@agner.ch>
PR target/88648
* config/arm/arm.c (arm_option_override_internal): Force
opts->x_inline_asm_unified to true only if TARGET_THUMB2_P.
* gcc.target/arm/pr88648-asm-syntax-unified.c: Add test to
check if -masm-syntax-unified gets applied properly.
From-SVN: r267804
Also fix some tests that were not cleaning up after themselves, as
identified by the change to nonexistent_path.
* testsuite/util/testsuite_fs.h (nonexistent_path): Include name
of the source file containing the caller.
* testsuite/27_io/filesystem/iterators/directory_iterator.cc: Remove
directories created by test.
* testsuite/27_io/filesystem/iterators/recursive_directory_iterator.cc:
Likewise.
* testsuite/experimental/filesystem/iterators/directory_iterator.cc:
Likewise.
* testsuite/experimental/filesystem/iterators/
recursive_directory_iterator.cc: Likewise.
From-SVN: r267801
PR tree-optimization/88775
* include/bits/stl_function.h (greater<_Tp*>::operator(),
less<_Tp*>::operator(), greater_equal<_Tp*>::operator(),
less_equal<_Tp*>::operator()): Use __builtin_is_constant_evaluated
instead of __builtin_constant_p if available. Don't bother with
the pointer comparison in C++11 and earlier.
From-SVN: r267800
2019-01-09 Steven G. Kargl <kargl@gcc.gnu.org>
PR fortran/88376
* resolve.c (is_illegal_recursion): Remove an assert().
2019-01-09 Steven G. Kargl <kargl@gcc.gnu.org>
PR fortran/88376
* gfortran.dg/pr88376.f90: New test.
From-SVN: r267793
PR go/86343
* go-gcc.cc (Gcc_backend::set_placeholder_struct_type): Go back to
build_distinct_type_copy, but copy the fields so that they have
the right DECL_CONTEXT.
From-SVN: r267789
This removes updates the removal date of all deprecations in phobos.
Many of the marked functions have passed their end dates, and are now
absent in upstream.
Reviewed-on: https://github.com/dlang/phobos/pull/6828
From-SVN: r267788
2019-01-09 Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/68426
* simplify.c (gfc_simplify_spread): Also simplify if the
type of source is an EXPR_STRUCTURE.
2019-01-09 Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/68426
* gfortran.dg/spread_simplify_1.f90: New test.
From-SVN: r267781
PR target/84010
* config/sparc/sparc.c (sparc_legitimize_tls_address): Only use Pmode
consistently in TLS address generation and adjust code to the renaming
of patterns. Mark calls to __tls_get_addr as const.
* config/sparc/sparc.md (tgd_hi22): Turn into...
(tgd_hi22<P:mode>): ...this and use Pmode throughout.
(tgd_lo10): Turn into...
(tgd_lo10<P:mode>): ...this and use Pmode throughout.
(tgd_add32): Merge into...
(tgd_add64): Likewise.
(tgd_add<P:mode>): ...this and use Pmode throughout.
(tldm_hi22): Turn into...
(tldm_hi22<P:mode>): ...this and use Pmode throughout.
(tldm_lo10): Turn into...
(tldm_lo10<P:mode>): ...this and use Pmode throughout.
(tldm_add32): Merge into...
(tldm_add64): Likewise.
(tldm_add<P:mode>): ...this and use Pmode throughout.
(tldm_call32): Merge into...
(tldm_call64): Likewise.
(tldm_call<P:mode>): ...this and use Pmode throughout.
(tldo_hix22): Turn into...
(tldo_hix22<P:mode>): ...this and use Pmode throughout.
(tldo_lox10): Turn into...
(tldo_lox10<P:mode>): ...this and use Pmode throughout.
(tldo_add32): Merge into...
(tldo_add64): Likewise.
(tldo_add<P:mode>): ...this and use Pmode throughout.
(tie_hi22): Turn into...
(tie_hi22<P:mode>): ...this and use Pmode throughout.
(tie_lo10): Turn into...
(tie_lo10<P:mode>): ...this and use Pmode throughout.
(tie_ld64): Use DImode throughout.
(tie_add32): Merge into...
(tie_add64): Likewise.
(tie_add<P:mode>): ...this and use Pmode throughout.
(tle_hix22_sp32): Merge into...
(tle_hix22_sp64): Likewise.
(tle_hix22<P:mode>): ...this and use Pmode throughout.
(tle_lox22_sp32): Merge into...
(tle_lox22_sp64): Likewise.
(tle_lox22<P:mode>): ...this and use Pmode throughout.
(*tldo_ldub_sp32): Merge into...
(*tldo_ldub_sp64): Likewise.
(*tldo_ldub<P:mode>): ...this and use Pmode throughout.
(*tldo_ldub1_sp32): Merge into...
(*tldo_ldub1_sp64): Likewise.
(*tldo_ldub1<P:mode>): ...this and use Pmode throughout.
(*tldo_ldub2_sp32): Merge into...
(*tldo_ldub2_sp64): Likewise.
(*tldo_ldub2<P:mode>): ...this and use Pmode throughout.
(*tldo_ldsb1_sp32): Merge into...
(*tldo_ldsb1_sp64): Likewise.
(*tldo_ldsb1<P:mode>): ...this and use Pmode throughout.
(*tldo_ldsb2_sp32): Merge into...
(*tldo_ldsb2_sp64): Likewise.
(*tldo_ldsb2<P:mode>): ...this and use Pmode throughout.
(*tldo_ldub3_sp64): Use DImode throughout.
(*tldo_ldsb3_sp64): Likewise.
(*tldo_lduh_sp32): Merge into...
(*tldo_lduh_sp64): Likewise.
(*tldo_lduh<P:mode>): ...this and use Pmode throughout.
(*tldo_lduh1_sp32): Merge into...
(*tldo_lduh1_sp64): Likewise.
(*tldo_lduh1<P:mode>): ...this and use Pmode throughout.
(*tldo_ldsh1_sp32): Merge into...
(*tldo_ldsh1_sp64): Likewise.
(*tldo_ldsh1<P:mode>): ...this and use Pmode throughout.
(*tldo_lduh2_sp64): Use DImode throughout.
(*tldo_ldsh2_sp64): Likewise.
(*tldo_lduw_sp32): Merge into...
(*tldo_lduw_sp64): Likewise.
(*tldo_lduw<P:mode>): ...this and use Pmode throughout.
(*tldo_lduw1_sp64): Use DImode throughout.
(*tldo_ldsw1_sp64): Likewise.
(*tldo_ldx_sp64): Likewise.
(*tldo_stb_sp32): Merge into...
(*tldo_stb_sp64): Likewise.
(*tldo_stb<P:mode>): ...this and use Pmode throughout.
(*tldo_sth_sp32): Merge into...
(*tldo_sth_sp64): Likewise.
(*tldo_sth<P:mode>): ...this and use Pmode throughout.
(*tldo_stw_sp32): Merge into...
(*tldo_stw_sp64): Likewise.
(*tldo_stw<P:mode>): ...this and use Pmode throughout.
(*tldo_stx_sp64): Use DImode throughout.
From-SVN: r267771
This patch is part of a series that enables ARMv8.5-A in GCC and
adds Branch Target Identification Mechanism.
This patch is adding a new configure option for enabling BTI and
Return Address Signing by default.
*** gcc/ChangeLog ***
2018-01-09 Sudakshina Das <sudi.das@arm.com>
* config/aarch64/aarch64.c (aarch64_override_options): Add case to
check configure option to set BTI and Return Address Signing.
* configure.ac: Add --enable-standard-branch-protection and
--disable-standard-branch-protection.
* configure: Regenerated.
* doc/install.texi: Document the same.
*** gcc/testsuite/ChangeLog ***
2018-01-09 Sudakshina Das <sudi.das@arm.com>
* gcc.target/aarch64/bti-1.c: Update test to not add command line
option when configure with bti.
* gcc.target/aarch64/bti-2.c: Likewise.
* lib/target-supports.exp
(check_effective_target_default_branch_protection):
Add configure check for --enable-standard-branch-protection.
From-SVN: r267770
This patch is part of a series that enables ARMv8.5-A in GCC and
adds Branch Target Identification Mechanism.
This patch adds a new pass called "bti" which is triggered by the command
line argument -mbranch-protection whenever "bti" is turned on.
The pass iterates through the instructions and adds appropriated BTI
instructions based on the following:
* Add a new "BTI C" at the beginning of a function, unless its already
protected by a "PACIASP". We exempt the functions that are only called
directly.
* Add a new "BTI J" for every target of an indirect jump, jump table
targets, non-local goto targets or labels that might be referenced by
variables, constant pools, etc (NOTE_INSN_DELETED_LABEL).
Since we have already changed the use of indirect tail calls to only x16 and
x17, we do not have to use "BTI JC".
(check patch 3/6).
*** gcc/ChangeLog ***
2018-01-09 Sudakshina Das <sudi.das@arm.com>
Ramana Radhakrishnan <ramana.radhakrishnan@arm.com>
* config.gcc (aarch64*-*-*): Add aarch64-bti-insert.o.
* gcc/config/aarch64/aarch64.h: Update comment for TRAMPOLINE_SIZE.
* config/aarch64/aarch64.c (aarch64_asm_trampoline_template): Update
if bti is enabled.
* config/aarch64/aarch64-bti-insert.c: New file.
* config/aarch64/aarch64-passes.def (INSERT_PASS_BEFORE): Insert bti
pass.
* config/aarch64/aarch64-protos.h (make_pass_insert_bti): Declare the
new bti pass.
* config/aarch64/aarch64.md (unspecv): Add UNSPECV_BTI_NOARG,
UNSPECV_BTI_C, UNSPECV_BTI_J and UNSPECV_BTI_JC.
(bti_noarg, bti_j, bti_c, bti_jc): New define_insns.
* config/aarch64/t-aarch64: Add rule for aarch64-bti-insert.o.
*** gcc/testsuite/ChangeLog ***
2018-01-09 Sudakshina Das <sudi.das@arm.com>
* gcc.target/aarch64/bti-1.c: New test.
* gcc.target/aarch64/bti-2.c: New test.
* gcc.target/aarch64/bti-3.c: New test.
* lib/target-supports.exp
(check_effective_target_aarch64_bti_hw): Add new check for BTI hw.
Co-Authored-By: Ramana Radhakrishnan <ramana.radhakrishnan@arm.com>
From-SVN: r267769
This patch is part of a series that enables ARMv8.5-A in GCC and
adds Branch Target Identification Mechanism.
This pass updates the CLI of -mbranch-protection to add "bti" as a new
type of branch protection and also add it its definition of "none" and
"standard". The option does not really do anything functional.
The functional changes are in the next patch. I am initializing the target
variable aarch64_enable_bti to 2 since I am also adding a configure option
in a later patch and a value different from 0 and 1 would help identify if its
already been updated.
*** gcc/ChangeLog ***
2018-01-09 Sudakshina Das <sudi.das@arm.com>
* config/aarch64/aarch64-protos.h (aarch64_bti_enabled): Declare.
* config/aarch64/aarch64.c (aarch64_handle_no_branch_protection):
Disable bti for -mbranch-protection=none.
(aarch64_handle_standard_branch_protection): Enable bti for
-mbranch-protection=standard.
(aarch64_handle_bti_protection): Enable bti for "bti" in the string to
-mbranch-protection.
(aarch64_bti_enabled): Check if bti is enabled.
* config/aarch64/aarch64.opt: Declare target variable.
* doc/invoke.texi: Add bti to the -mbranch-protection documentation.
From-SVN: r267768
This patch is part of a series that enables ARMv8.5-A in GCC and
adds Branch Target Identification Mechanism.
This patch changes the registers that are allowed for indirect tail calls.
We are choosing to restrict these to only x16 or x17.
Indirect tail calls are special in a way that they convert a call statement
(BLR instruction) to a jump statement (BR instruction). For the best possible
use of Branch Target Identification Mechanism, we would like to place a
"BTI C" (call) at the beginning of the function which is only
compatible with BLRs and BR X16/X17. In order to make indirect tail calls
compatible with this scenario, we are restricting the TAILCALL_ADDR_REGS.
In order to use x16/x17 for this purpose, we also had to change the use
of these registers in the epilogue/prologue handling. For this purpose
we are now using x12 and x13 named as EP0_REGNUM and EP1_REGNUM as
scratch registers for epilogue and prologue.
*** gcc/ChangeLog***
2018-01-09 Sudakshina Das <sudi.das@arm.com>
* config/aarch64/aarch64.c (aarch64_expand_prologue): Use new
epilogue/prologue scratch registers EP0_REGNUM and EP1_REGNUM.
(aarch64_expand_epilogue): Likewise.
(aarch64_output_mi_thunk): Likewise
* config/aarch64/aarch64.h (REG_CLASS_CONTENTS): Change
TAILCALL_ADDR_REGS to x16 and x17.
* config/aarch64/aarch64.md: Define EP0_REGNUM and EP1_REGNUM.
*** gcc/testsuite/ChangeLog ***
2018-01-09 Sudakshina Das <sudi.das@arm.com>
* gcc.target/aarch64/test_frame_17.c: Update to check for EP0_REGNUM
instead of IP0_REGNUM and add test case.
From-SVN: r267767
