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
Jakub Jelinek