Add some test-cases using vector_length 64.
2019-01-12 Tom de Vries <tdevries@suse.de>
* testsuite/libgomp.oacc-c-c++-common/vector-length-64-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/vector-length-64-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/vector-length-64-3.c: New test.
From-SVN: r267895
Add test-cases for "[nvptx] Force vl32 if calling vector-partitionable
routines".
2019-01-12 Tom de Vries <tdevries@suse.de>
PR target/85486
* testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/pr85486.c: New test.
From-SVN: r267894
Add test-cases for PR85381.
2019-01-12 Tom de Vries <tdevries@suse.de>
PR target/85381
* testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: New test.
* testsuite/libgomp.oacc-c-c++-common/pr85381.c: New test.
From-SVN: r267893
Add various reduction test-cases with vector length 128.
2019-01-12 Tom de Vries <tdevries@suse.de>
* testsuite/libgomp.oacc-c-c++-common/vred2d-128.c: New test.
* testsuite/libgomp.oacc-fortran/gemm.f90: New test.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-10.c: New test.
From-SVN: r267892
Add various test-cases with vector length 128.
2019-01-12 Tom de Vries <tdevries@suse.de>
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: New test.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c: New test.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c: New test.
From-SVN: r267891
Update message in nvptx libgomp plugin about insufficient resources to launch
kernel, to accommodate for the fact the vector_length can now be variable.
2019-01-12 Tom de Vries <tdevries@suse.de>
* plugin/plugin-nvptx.c (nvptx_exec): Update insufficient hardware
resources diagnostic.
From-SVN: r267890
Allow vector_length clauses to accept values larger than warp size. Note that
this does not enable setting vector_length to values larger than warp size using
-fopenacc-dim.
2019-01-12 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.c (nvptx_goacc_validate_dims): Take larger vector
lengths into account.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Expect
vector length to be 128.
* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Expect vector
length 2097152 to be reduced to 1024 instead of 32.
From-SVN: r267889
2019-01-12 Paul Thomas <pault@gcc.gnu.org>
* gfortran.dg/ISO_Fortran_binding_2.f90 : Remove because of
reports of ICEs.
* gfortran.dg/ISO_Fortran_binding_2.c : Ditto.
From-SVN: r267884
2019-01-12 Paul Thomas <pault@gcc.gnu.org>
* gfortran.texi : Add description in sections on TS 29113 and
further interoperability with C.
* trans-array.c (gfc_conv_descriptor_attribute): New function.
(gfc_get_dataptr_offset): Remove static function attribute.
* trans-array.h : Add prototypes for above functions.
* trans-decl.c : Add declarations for the library functions
cfi_desc_to_gfc_desc and gfc_desc_to_cfi_desc.
* trans-expr.c (gfc_conv_gfc_desc_to_cfi_desc): New function.
(gfc_conv_procedure_call): Call it for scalar and array actual
arguments, when the formal arguments are bind_c with assumed
shape or assumed rank.
* trans.h : External declarations for gfor_fndecl_cfi_to_gfc
and gfor_fndecl_gfc_to_cfi.
2019-01-12 Paul Thomas <pault@gcc.gnu.org>
* gfortran.dg/ISO_Fortran_binding_1.f90 : New test.
* gfortran.dg/ISO_Fortran_binding_1.c : Auxilliary file for test.
* gfortran.dg/ISO_Fortran_binding_2.f90 : New test.
* gfortran.dg/ISO_Fortran_binding_2.c : Auxilliary file for test.
* gfortran.dg/bind_c_array_params_2.f90 : Change search string
for dump tree scan.
2019-01-12 Paul Thomas <pault@gcc.gnu.org>
* ISO_Fortran_binding.h : New file.
* Makefile.am : Include ISO_Fortran_binding.c in the list of
files to compile.
* Makefile.in : Regenerated.
* gfortran.map : Add _gfortran_cfi_desc_to_gfc_desc,
_gfortran_gfc_desc_to_cfi_desc and the CFI API functions.
* runtime/ISO_Fortran_binding.c : New file containing the new
functions added to the map.
From-SVN: r267881
Make "[nvptx] Force vl32 if calling vector-partitionable routines" work as well
if vector length is set by modifying PTX_DEFAULT_VECTOR_LENGTH.
2019-01-12 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.c (nvptx_goacc_validate_dims_1): In offloading
region calling vector-partitionable routine, set default_vector_length
to WARP_SIZE.
From-SVN: r267879
In nvptx_goacc_validate_dims_1, allow oacc_default_dims[DIM_VECTOR] to be
overridden, by assigning it to a new variable default_vector_length at the
start, and using it at the end.
2019-01-12 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.c (nvptx_goacc_validate_dims_1): Add an use new
variable default_vector_length.
From-SVN: r267878
There's a problem in oacc_validate_dims that when f.i. the worker dimension
is set using -fopenacc-dim=:32, and the vector_length is set using a
"vector_length (128)" clause, the compiler combines, accepts and emits the
values, while the combination of the two is invalid.
The reason for this is that while oacc_validate_dims validates the dimensions
using targetm.goacc.validate_dims before applying default or minimum values,
it does not do so afterwards.
Work around this in the nvptx port by applying the defaults from
oacc_default_dims at the end of nvptx_goacc_validate_dims_1, as
oacc_validate_dims would do it, and then apply the dimensions limits.
2019-01-12 Tom de Vries <tdevries@suse.de>
PR middle-end/88703
* config/nvptx/nvptx.c (nvptx_goacc_validate_dims_1): Apply defaults
from oacc_default_dims, as oacc_validate_dims would do it, and apply
dimensions limits.
From-SVN: r267877
Add a used parameter to TARGET_GOACC_VALIDATE_DIMS, allowing a target to make
decisions in the hook implementation based on whether a dimension is used or
not.
2019-01-12 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.c (nvptx_goacc_validate_dims_1)
(nvptx_goacc_validate_dims): Add used parameter.
* doc/tm.texi: Regenerate.
* omp-offload.c (oacc_parse_default_dims, oacc_validate_dims): Add
argument to call to targetm.goacc.validate_dims.
(default_goacc_validate_dims): Add used
parameter.
* target.def (validate_dims): Add used parameter in DEFHOOK.
* targhooks.h (default_goacc_validate_dims): Add used parameter.
From-SVN: r267876
PR libstdc++/88811
PR libstdc++/83306
* src/filesystem/path.cc: Fix typo. If first path is empty, show []
before second path.
* testsuite/experimental/filesystem/filesystem_error/cons.cc: New
test.
From-SVN: r267868
* doc/xml/manual/intro.xml: Include new section.
* doc/xml/manual/status_cxx2017.xml: Document more
implementation-defined properties of the library.
* doc/xml/manual/status_cxx2020.xml: Document C++2a status.
* doc/html/*: Regenerate.
From-SVN: r267867
This patch implements the C++2a proposal to allow incomplete types in
std::reference_wrapper, which was previously undefined.
The change cannot be implemented for earlier standards, because prior to
C++2a std::reference_wrapper has a weak result type, so must inspect the
template argument to see if it defines a nested result_type member. That
is deprecated (but still required) in C++17, and removed from C++2a.
The removal of the base class from reference_wrapper is a potential ABI
change, as it could alter the layout of a type which derives from
reference_wrapper<T> and from an empty type with _Weak_result_type<T> as
a base class. Previously the repeated _Weak_result_type<T> base class
would have prevented the empty base-class optimization, but if
reference_wrapper<T> no longer derives from it, the empty class could be
placed at the same address as the reference_wrapper<T> base. In
practice, the only types which derive from _Weak_result_type or from
_Reference_wrapper_base_memfun or any of its base classes are non-empty
types defined in libstdc++: std::reference_wrapper, std::function, and
std::_Bind. As they are non-empty types, they are not eligible for EBO
anyway.
* include/bits/refwrap.h [__cplusplus > 201703L]
(_Refwrap_base_arg1, _Refwrap_base_arg2, _Reference_wrapper_base)
(_Reference_wrapper_base_memfun): Do not define for C++2a.
(reference_wrapper): Do not derive from _Reference_wrapper_base_memfun
for C++2a.
(reference_wrapper::operator()): Add static assertion.
* testsuite/20_util/reference_wrapper/incomplete.cc: New test.
From-SVN: r267866
This paper has been included in the C++20 draft, but the changes to add
noexcept can be made unconditionally, to apply for C++11 too.
* include/std/chrono (duration_values::zero(), duration_values::min())
(duration_values::max()): Add noexcept.
(duration::zero(), duration::min(), duration::max()): Likewise.
(time_point::zero(), time_point::min(), time_point::max()): Likewise.
* testsuite/20_util/duration/requirements/noexcept.cc: New test.
* testsuite/20_util/time_point/requirements/noexcept.cc: New test.
From-SVN: r267865
2019-01-11 Steven G. Kargl <kargl@gcc.gnu.org>
PR fortran/35031
* decl.c (gfc_match_entry): Check for F2018:C1546. Fix nearby
mis-indentation.
2019-01-11 Steven G. Kargl <kargl@gcc.gnu.org>
PR fortran/35031
* gfortran.dg/pr35031.f90: new test.
From-SVN: r267864
For a struct with zero-sized last field, the address of the
field falls out of the object boundary, which confuses the
garbage collector. Pad an extra byte in this case.
Reviewed-on: https://go-review.googlesource.com/c/157557
From-SVN: r267861
The standard doesn't really talk about an expression depending on the number
of elements of a pack, but that's definitely an important form of template
argument dependence.
* pt.c (instantiation_dependent_r): A template non-type parameter
pack is instantiation-dependent.
From-SVN: r267860
The issue here was that we were cp_folding a location wrapper around a
lambda capture proxy before it had been mark_rvalue_used. I considered
adding mark_rvalue_use calls to build_new_op_1, but it seems appropriate to
have them in cp_fold_maybe_rvalue when we know we're trying to produce an
rvalue.
The change to mark_use is for a related issue: when we change the operand of
the location wrapper from VAR_DECL to INTEGER_CST, we need the TREE_CODE of
the location wrapper to change as well, from VIEW_CONVERT_EXPR to
NON_LVALUE_EXPR.
* expr.c (mark_use): Fix location wrapper handling.
* cp-gimplify.c (cp_fold_maybe_rvalue): Call mark_rvalue_use.
From-SVN: r267859
PR middle-end/85956
PR lto/88733
* tree-inline.h (struct copy_body_data): Add adjust_array_error_bounds
field.
* tree-inline.c (remap_type_1): Formatting fix. If TYPE_MAX_VALUE of
ARRAY_TYPE's TYPE_DOMAIN is newly error_mark_node, replace it with
a dummy "omp dummy var" variable if id->adjust_array_error_bounds.
* omp-low.c (new_omp_context): Set cb.adjust_array_error_bounds.
fortran/
* trans-openmp.c: Include attribs.h.
(gfc_walk_alloc_comps, gfc_omp_clause_linear_ctor): Handle
VAR_DECL max bound with "omp dummy var" attribute like NULL or
error_mark_node - recompute number of elts independently.
testsuite/
* c-c++-common/gomp/pr85956.c: New test.
* g++.dg/gomp/pr88733.C: New test.
From-SVN: r267858
2019-01-11 Vladimir Makarov <vmakarov@redhat.com>
PR rtl-optimization/87305
* lra-assigns.c
(setup_live_pseudos_and_spill_after_risky_transforms): Add code
for little endian pseudos used as paradoxical subreg.
From-SVN: r267854
PR tree-optimization/88693
* tree-ssa-strlen.c (get_min_string_length): Don't set *full_string_p
for STRING_CSTs that don't contain any NUL characters in the first
TREE_STRING_LENGTH bytes.
* gcc.c-torture/execute/pr88693.c: New test.
From-SVN: r267852
PR 88777
PR 88614
* genattrtab.c (min_fn): Don't translate values.
(min_attr_value): Return INT_MAX when the value can't be calculated.
Return minimum among any values that can be calculated.
(max_attr_value): Adjust.
From-SVN: r267851
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
Tom de Vries