| Age | Commit message (Collapse) | Author | Files | Lines |
|---|
|
gcc/cp/
* cvt.c (type_promotes_to): Handle char8_t promotion.
* decl.c (grokdeclarator): Handle invalid type specifier
combinations involving char8_t.
* lex.c (init_reswords): Add char8_t as a reserved word.
* mangle.c (write_builtin_type): Add name mangling for char8_t (Du).
* parser.c (cp_keyword_starts_decl_specifier_p)
(cp_parser_simple_type_specifier): Recognize char8_t as a simple
type specifier.
(cp_parser_string_literal): Use char8_array_type_node for the type
of CPP_UTF8STRING.
(cp_parser_set_decl_spec_type): Tolerate char8_t typedefs in system
headers.
* rtti.c (emit_support_tinfos): type_info support for char8_t.
* tree.c (char_type_p): Recognize char8_t as a character type.
* typeck.c (string_conv_p): Handle conversions of u8 string
literals of char8_t type.
(check_literal_operator_args): Handle UDLs with u8 string literals
of char8_t type.
* typeck2.c (ordinary_char_type_p): New.
(digest_init_r): Disallow initializing a char array with a u8 string
literal.
gcc/c-family/
* c-common.c (c_common_reswords): Add char8_t.
(fix_string_type): Use char8_t for the type of u8 string literals.
(c_common_get_alias_set): char8_t doesn't alias.
(c_common_nodes_and_builtins): Define char8_t as a builtin type in
C++.
(c_stddef_cpp_builtins): Add __CHAR8_TYPE__.
(keyword_begins_type_specifier): Add RID_CHAR8.
* c-common.h (rid): Add RID_CHAR8.
(c_tree_index): Add CTI_CHAR8_TYPE and CTI_CHAR8_ARRAY_TYPE.
Define D_CXX_CHAR8_T and D_CXX_CHAR8_T_FLAGS.
Define char8_type_node and char8_array_type_node.
* c-cppbuiltin.c (cpp_atomic_builtins): Predefine
__GCC_ATOMIC_CHAR8_T_LOCK_FREE.
(c_cpp_builtins): Predefine __cpp_char8_t.
* c-lex.c (lex_string): Use char8_array_type_node as the type of
CPP_UTF8STRING.
(lex_charconst): Use char8_type_node as the type of CPP_UTF8CHAR.
* c-opts.c: If not otherwise specified, enable -fchar8_t when
targeting C++2a.
* c.opt: Add the -fchar8_t command line option.
libiberty/
* cp-demangle.c (cplus_demangle_builtin_types)
(cplus_demangle_type): Add name demangling for char8_t (Du).
* cp-demangle.h: Increase D_BUILTIN_TYPE_COUNT to accommodate the
new char8_t type.
From-SVN: r267923
|
|
gcc/c-family/ChangeLog:
PR target/88638
* c-attribs.c (positional_argument): Call valid_format_string_type_p
and issue errors if it fails.
* c-common.h (valid_format_string_type_p): Declare.
* c-format.c (valid_stringptr_type_p): Rename...
(valid_format_string_type_p): ...to this and make extern.
(handle_format_arg_attribute): Adjust to new name.
(check_format_string): Same.
gcc/testsuite/ChangeLog:
PR target/88638
* gcc.dg/format/attr-8.c: New test.
* gcc.dg/darwin-cfstring-format-1.c: Adjust diagnostics.
* gcc.dg/format/attr-3.c: Same.
* obj-c++.dg/fsf-nsstring-format-1.mm: Same.
* objc.dg/fsf-nsstring-format-1.m: Same.
gcc/ChangeLog:
PR target/88638
* doc/extend.texi (Darwin Format Checks): Clarify.
From-SVN: r267922
|
|
2019-01-14 Martin Liska <mliska@suse.cz>
PR gcov-profile/88263
* decl2.c (get_tls_wrapper_fn): Use DECL_SOURCE_LOCATION
as location of the TLS wrapper.
2019-01-14 Martin Liska <mliska@suse.cz>
PR gcov-profile/88263
* g++.dg/gcov/pr88263-2.C: New test.
From-SVN: r267921
|
|
gcc/ChangeLog:
* invoke.texi (-Wmemset-transposed-args): Fix typos, adjust wording.
From-SVN: r267920
|
|
in (simplify ...) or (match ...) context.
2019-01-14 Richard Biener <rguenther@suse.de>
* genmatch.c (dt_simplify::gen_1): Change dumping dependent on
whether we are in (simplify ...) or (match ...) context.
From-SVN: r267917
|
|
PR rtl-optimization/88796
* emit-rtl.h (struct rtl_data): Add stack_protect_guard_decl field.
* cfgexpand.c (stack_protect_prologue): Initialize
crtl->stack_protect_guard_decl.
* function.c (stack_protect_epilogue): Use it instead of calling
targetm.stack_protect_guard again.
* dse.c (check_mem_read_rtx): Ignore MEM_VOLATILE_P reads from
MEMs with MEM_EXPR equal to crtl->stack_protect_guard or
crtl->stack_protect_guard_decl.
* config/i386/i386.c (ix86_stack_protect_guard): Set TREE_THIS_VOLATILE
on the returned MEM_EXPR.
* gcc.target/i386/pr88796.c: New test.
From-SVN: r267916
|
|
From-SVN: r267914
|
|
One fix in the asm statement parser to stop parsing if the end of the
statement has been reached, and moves all inline asm tests to gdc.dg.
These being adjusted where necessary to test the GCC style instead.
gcc/testsuite/ChangeLog:
2019-01-14 Iain Buclaw <ibuclaw@gdcproject.org>
* gdc.dg/asm1.d: New test.
* gdc.dg/asm2.d: New test.
* gdc.dg/asm3.d: New test.
* gdc.dg/asm4.d: New test.
* lib/gdc.exp (gdc_init): Set gcc_error_prefix and gcc_warning_prefix.
From-SVN: r267913
|
|
2019-01-13 Jerry DeLisle <jvdelisle@gcc.gnu.org>
PR libfortran/88776
* io/open.c (newunit): Free format buffer if the unit specified is for
stdin, stdout, or stderr.
* gfortran.dg/namelist_96.f90: New test.
From-SVN: r267910
|
|
From-SVN: r267909
|
|
@xref)
2019-01-13 Dominique d'Humieres <dominiq@gcc.gnu.org>
PR fortran/88803
* gfortran.texi: Replace @xref with @ref and adjust the sentence.
From-SVN: r267906
|
|
contiguous functions results as actual arguments)
2019-01-13 Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/57992
* trans-array.c (gfc_conv_array_parameter): Do not pack/unpack
functions with contiguous results.
2019-01-13 Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/59345
* gfortran.dg/internal_pack_18.f90: New test.
From-SVN: r267905
|
|
gcc/c-family/
* c-warn.c (warn_for_address_or_pointer_of_packed_member):
Replace "may may" with "may" in warning message.
gcc/testsuite/
* gcc.dg/pr51628-20.c: Updated.
* gcc.dg/pr51628-21.c: Likewise.
* gcc.dg/pr51628-25.c: Likewise.
From-SVN: r267904
|
|
2019-01-13 Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/59345
* trans-array.c (gfc_conv_array_parameter): Remove TODO. Do not
pack/unpack results of functions which return an explicit-shaped
or allocatable array.
2019-01-13 Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/59345
* gfortran.dg/internal_pack_17.f90: New test.
* gfortran.dg/alloc_comp_auto_array_3.f90: Adjust number of calls
to builtin_free.
From-SVN: r267903
|
|
2019-01-12 Steven G. Kargl <kargl@gcc.gnu.org>
PR fortran/61765
* resolve.c (gfc_verify_binding_labels): Break if-elseif-elseif
structure into independent if's with a return to simplify logic.
Avoid a check for ENTRY name with bind(c).
2019-01-12 Steven G. Kargl <kargl@gcc.gnu.org>
PR fortran/61765
* gfortran.dg/pr61765.f90: New test.
From-SVN: r267902
|
|
From-SVN: r267901
|
|
Enable setting vector length using -fopenacc-dim, f.i. -fopenacc-dim=::128.
2019-01-12 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.c (nvptx_goacc_validate_dims_1): Alow setting
vector length using -fopenacc-dim.
* plugin/plugin-nvptx.c (nvptx_exec): Update error message.
From-SVN: r267896
|
|
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
|
|
/cp
2019-01-12 Paolo Carlini <paolo.carlini@oracle.com>
* decl.c (cp_finish_decl): Improve error location.
* decl2.c (grokfield): Likewise, improve two locations.
/testsuite
2019-01-12 Paolo Carlini <paolo.carlini@oracle.com>
* g++.dg/cpp0x/pr62101.C: Test locations too.
* g++.dg/inherit/pure1.C: Likewise.
From-SVN: r267888
|
|
/cp
2019-01-12 Paolo Carlini <paolo.carlini@oracle.com>
* decl.c (cp_finish_decl): Improve error location.
* decl2.c (grokfield): Likewise, improve two locations.
/testsuite
2019-01-12 Paolo Carlini <paolo.carlini@oracle.com>
* g++.dg/cpp0x/pr62101.C: Test locations too.
* g++.dg/inherit/pure1.C: Likewise.
From-SVN: r267887
|
|
* config/i386/gnu.h (TARGET_THREAD_SSP_OFFSET): Define.
(TARGET_CAN_SPLIT_STACK): Define.
(TARGET_THREAD_SPLIT_STACK_OFFSET): Define.
From-SVN: r267886
|
|
gcc/d/ChangeLog:
* README.gcc: New file.
libphobos/ChangeLog:
* README.gcc: New file.
From-SVN: r267885
|
|
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
|
|
From-SVN: r267883
|
|
* tree-ssa-loop-ivopts.c (find_inv_vars): Fix a comment typo.
* c-typeck.c (convert_for_assignment): Fix a comment typo.
From-SVN: r267882
|
|
interoperability with C.
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
|
|
which shouldn't be executable.
From-SVN: r267873
|
|
From-SVN: r267872
|
|
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
|
|
* typeck.c (maybe_warn_pessimizing_move): Return if ARG isn't
ADDR_EXPR.
* g++.dg/cpp0x/Wredundant-move5.C: New test.
* g++.dg/cpp0x/Wredundant-move6.C: New test.
From-SVN: r267862
|
|
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
|
|
PR C++/8811
* decl2.c (maybe_emit_vtables): If needed, generate code for
the destructor of an abstract class.
(mark_used): Update comment for older function-name change.
PR C++/88114
* g++.dg/cpp0x/defaulted61.C: New.
* g++.dg/cpp0x/defaulted62.C: New.
From-SVN: r267855
|
|
setup_live_pseudos_and_spill_after_risky_transforms on aarch64 big-endian)
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
|
|
hppa-linux-gnu)
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
|
|
From-SVN: r267850
|
|
2019-01-11 Steve Ellcey <sellcey@marvell.com>
* config/aarch64/aarch64.c (aarch64_simd_call_p): New function.
(aarch64_hard_regno_call_part_clobbered): Add insn argument.
(aarch64_return_call_with_max_clobbers): New function.
(TARGET_RETURN_CALL_WITH_MAX_CLOBBERS): New macro.
* config/avr/avr.c (avr_hard_regno_call_part_clobbered): Add insn
argument.
* config/i386/i386.c (ix86_hard_regno_call_part_clobbered): Ditto.
* config/mips/mips.c (mips_hard_regno_call_part_clobbered): Ditto.
* config/rs6000/rs6000.c (rs6000_hard_regno_call_part_clobbered): Ditto.
* config/s390/s390.c (s390_hard_regno_call_part_clobbered): Ditto.
* cselib.c (cselib_process_insn): Add argument to
targetm.hard_regno_call_part_clobbered call.
* ira-conflicts.c (ira_build_conflicts): Ditto.
* ira-costs.c (ira_tune_allocno_costs): Ditto.
* lra-constraints.c (inherit_reload_reg): Ditto.
* lra-int.h (struct lra_reg): Add call_insn field, remove call_p field.
* lra-lives.c (check_pseudos_live_through_calls): Add call_insn
argument. Call targetm.return_call_with_max_clobbers.
Add argument to targetm.hard_regno_call_part_clobbered call.
(calls_have_same_clobbers_p): New function.
(process_bb_lives): Add call_insn and last_call_insn variables.
Pass call_insn to check_pseudos_live_through_calls.
Modify if stmt to check targetm.return_call_with_max_clobbers.
Update setting of flush variable.
(lra_create_live_ranges_1): Set call_insn to NULL instead of call_p
to false.
* lra.c (initialize_lra_reg_info_element): Set call_insn to NULL.
* regcprop.c (copyprop_hardreg_forward_1): Add argument to
targetm.hard_regno_call_part_clobbered call.
* reginfo.c (choose_hard_reg_mode): Ditto.
* regrename.c (check_new_reg_p): Ditto.
* reload.c (find_equiv_reg): Ditto.
* reload1.c (emit_reload_insns): Ditto.
* sched-deps.c (deps_analyze_insn): Ditto.
* sel-sched.c (init_regs_for_mode): Ditto.
(mark_unavailable_hard_regs): Ditto.
* targhooks.c (default_dwarf_frame_reg_mode): Ditto.
* target.def (hard_regno_call_part_clobbered): Add insn argument.
(return_call_with_max_clobbers): New target function.
* doc/tm.texi: Regenerate.
* doc/tm.texi.in (TARGET_RETURN_CALL_WITH_MAX_CLOBBERS): New hook.
* hooks.c (hook_bool_uint_mode_false): Change to
hook_bool_insn_uint_mode_false.
* hooks.h (hook_bool_uint_mode_false): Ditto.
From-SVN: r267848
|
|
2019-01-11 Steve Ellcey <sellcey@marvell.com>
* config/aarch64/aarch64.c (aarch64_simd_call_p): New function.
(aarch64_remove_extra_call_preserved_regs): New function.
(TARGET_REMOVE_EXTRA_CALL_PRESERVED_REGS): New macro.
* doc/tm.texi.in (TARGET_REMOVE_EXTRA_CALL_PRESERVED_REGS): New hook.
* doc/tm.texi: Regenerate.
* final.c (get_call_reg_set_usage): Call new hook.
* target.def (remove_extra_call_preserved_regs): New hook.
* targhooks.c (default_remove_extra_call_preserved_regs): New function.
* targhooks.h (default_remove_extra_call_preserved_regs): New function.
From-SVN: r267846
|
|
gcc/testsuite/ChangeLog:
2019-01-11 Tamar Christina <tamar.christina@arm.com>
* gcc.target/aarch64/advsimd-intrinsics/vector-complex_f16.c: Require neon
and add options.
From-SVN: r267843
|
|
2019-01-11 Martin Liska <mliska@suse.cz>
PR middle-end/88758
* g++.dg/lto/pr88758_0.C: New test.
* g++.dg/lto/pr88758_1.C: New test.
From-SVN: r267840
|
|
PR bootstrap/88714
* passes.c (finish_optimization_passes): Call print_combine_total_stats
inside of pass_combine_1 dump rather than pass_profile_1.
From-SVN: r267839
|
|
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
|
