Commit Graph

915 Commits

Author SHA1 Message Date
Tobias Burnus
c28712beb4 libgomp plugin - init string
libgomp/
2019-09-13  Tobias Burnus  <tobias@codesourcery.com>

        * plugin/plugin-hsa.c (hsa_warn, hsa_fatal, hsa_error): Ensure
        string is initialized.

From-SVN: r275703
2019-09-13 20:14:02 +02:00
Florian Weimer
3c086f8dec Fix GCC_LINUX_FUTEX to work with C99 compilers
Without this change, libstdc++ is built without futex symbols if GCC
rejects implicit function declarations by default.

From-SVN: r275454
2019-09-06 12:27:36 +02:00
Chung-Lin Tang
c6c2d1bc9b re PR other/79543 (Inappropriate "ld --version" checking)
2019-09-03  Chung-Lin Tang <cltang@codesourcery.com>

	libatomic/
	PR other/79543
	* acinclude.m4 (LIBAT_CHECK_LINKER_FEATURES): Fix GNU ld --version
	scanning to conform to the GNU Coding Standards.
	* configure: Regenerate.

	libffi/
	PR other/79543
	* acinclude.m4 (LIBAT_CHECK_LINKER_FEATURES): Fix GNU ld --version
	scanning to conform to the GNU Coding Standards.
	* configure: Regenerate.

	libgomp/
	PR other/79543
	* acinclude.m4 (LIBGOMP_CHECK_LINKER_FEATURES): Fix GNU ld --version
	scanning to conform to the GNU Coding Standards.
	* configure: Regenerate.

	libitm/
	PR other/79543
	* acinclude.m4 (LIBITM_CHECK_LINKER_FEATURES): Fix GNU ld --version
	scanning to conform to the GNU Coding Standards.
	* configure: Regenerate.

	libstdc++-v3/
	PR other/79543
	* acinclude.m4 (GLIBCXX_CHECK_LINKER_FEATURES): Fix GNU ld --version
	scanning to conform to the GNU Coding Standards.
	* configure: Regenerate.

From-SVN: r275341
2019-09-03 14:10:26 +00:00
Jakub Jelinek
5cb72d83bb re PR libgomp/91530 (Several libgomp.*/scan-* tests FAIL without avx_runtime)
PR libgomp/91530
	* config/i386/sse.md (vec_shl_<mode>, vec_shr_<mode>): Use
	V_128 iterator instead of VI_128.

	* testsuite/libgomp.c/scan-21.c: New test.
	* testsuite/libgomp.c/scan-22.c: New test.

From-SVN: r274984
2019-08-28 12:12:11 +02:00
Jakub Jelinek
0ad7981cb4 re PR libgomp/91530 (Several libgomp.*/scan-* tests FAIL without avx_runtime)
PR libgomp/91530
	* testsuite/libgomp.c/scan-11.c: Add -msse2 option for sse2_runtime
	targets.
	* testsuite/libgomp.c/scan-12.c: Likewise.
	* testsuite/libgomp.c/scan-13.c: Likewise.
	* testsuite/libgomp.c/scan-14.c: Likewise.
	* testsuite/libgomp.c/scan-15.c: Likewise.
	* testsuite/libgomp.c/scan-16.c: Likewise.
	* testsuite/libgomp.c/scan-17.c: Likewise.
	* testsuite/libgomp.c/scan-18.c: Likewise.
	* testsuite/libgomp.c/scan-19.c: Likewise.
	* testsuite/libgomp.c/scan-20.c: Likewise.
	* testsuite/libgomp.c++/scan-9.C: Likewise.
	* testsuite/libgomp.c++/scan-10.C: Likewise.
	* testsuite/libgomp.c++/scan-11.C: Likewise.
	* testsuite/libgomp.c++/scan-12.C: Likewise.
	* testsuite/libgomp.c++/scan-14.C: Likewise.
	* testsuite/libgomp.c++/scan-15.C: Likewise.
	* testsuite/libgomp.c++/scan-13.C: Likewise.  Use sse2_runtime
	instead of i?86-*-* x86_64-*-* as target for scan-tree-dump-times.
	* testsuite/libgomp.c++/scan-16.C: Likewise.

From-SVN: r274947
2019-08-27 12:45:55 +02:00
Thomas Koenig
1e67491a0d re PR libgomp/91473 (Test case libgomp.fortran/appendix-a/a.28.5.f90 is invalid)
2019-08-17  Thomas Koenig  <tkoenig@gcc.gnu.org>

    PR fortran/91473
    * testsuite/libgomp.fortran/appendix-a/a.28.5.f90: Add
    -std=legacy so invalid code in the test case is accepted.

From-SVN: r274602
2019-08-17 11:57:25 +00:00
Thomas Koenig
393fdeb1e4 re PR fortran/91424 (Extend warnings about DO loops)
2019-08-12  Thomas Koenig  <tkoenig@gcc.gnu.org>

	PR fortran/91424
	* frontend-passes.c (do_subscript): Do not warn for an
	expression a second time.  Do not warn about a zero-trip loop.
	(doloop_warn): Also look at contained namespaces.

2019-08-12  Thomas Koenig  <tkoenig@gcc.gnu.org>

	PR fortran/91424
	* gfortran.dg/do_subscript_3.f90: New test.
	* gfortran.dg/do_subscript_4.f90: New test.
	* gfortran.dg/pr70754.f90: Use indices that to not overflow.

2019-08-12  Thomas Koenig  <tkoenig@gcc.gnu.org>

	PR fortran/91422
	* testsuite/libgomp.oacc-fortran/routine-7.f90: Correct array
	dimension.

From-SVN: r274320
2019-08-12 20:21:37 +00:00
Jakub Jelinek
8860d2706d gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT for VLA helper variables on target data even if...
* gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT
	for VLA helper variables on target data even if not GOVD_FIRSTPRIVATE.
	(gimplify_scan_omp_clauses): For OMP_CLAUSE_USE_DEVICE_* use just
	GOVD_EXPLICIT flags.
	(gimplify_omp_workshare): For OMP_TARGET_DATA move all
	OMP_CLAUSE_USE_DEVICE_* clauses to the end of clauses chain.
	* omp-low.c (scan_sharing_clauses): For OMP_CLAUSE_USE_DEVICE_*
	call install_var_field with mask 11 instead of 3.
	(lower_omp_target): For OMP_CLAUSE_USE_DEVICE_* use pass
	(splay_tree_key) &DECL_UID (var) to build_sender_ref instead of var.
gcc/c/
	* c-typeck.c (c_finish_omp_clauses): For C_ORT_OMP
	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
	instead of generic_head to track duplicates.
gcc/cp/
	* semantics.c (finish_omp_clauses): For C_ORT_OMP
	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
	instead of generic_head to track duplicates.
libgomp/
	* target.c (gomp_map_vars_internal): For GOMP_MAP_USE_DEVICE_PTR
	perform the lookup in the first loop only if !not_found_cnt, otherwise
	perform lookups for it in the second loop guarded with
	if (not_found_cnt || has_firstprivate).
	* testsuite/libgomp.c/target-37.c: New test.
	* testsuite/libgomp.c++/target-22.C: New test.

From-SVN: r274206
2019-08-08 08:39:02 +02:00
Jakub Jelinek
398e3feb8a tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR OpenMP description.
* tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR
	OpenMP description.  Add OMP_CLAUSE_USE_DEVICE_ADDR clause.
	* tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries
	for OMP_CLAUSE_USE_DEVICE_ADDR clause.
	(walk_tree_1): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Likewise.
	* gimplify.c (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses):
	Likewise.
	* omp-low.c (scan_sharing_clauses, lower_omp_target): Likewise.
	Treat OMP_CLAUSE_USE_DEVICE_ADDR like OMP_CLAUSE_USE_DEVICE_PTR
	clause with array or reference to array types, no matter what type
	except for reference it has.
gcc/c-family/
	* c-pragma.h (enum pragma_omp_clause): Add
	PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.  Set PRAGMA_OACC_CLAUSE_USE_DEVICE
	equal to PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR instead of being a separate
	enumeration value.
gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause.
	(c_parser_omp_clause_use_device_addr): New function.
	(c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(c_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
	like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
	map or use_device_* clauses.
	* c-typeck.c (c_finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
	in OpenMP, require pointer type rather than pointer or array type.
	Handle OMP_CLAUSE_USE_DEVICE_ADDR.
gcc/cp/
	* parser.c (cp_parser_omp_clause_name): Parse use_device_addr clause.
	(cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(cp_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
	like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
	map or use_device_* clauses.
	* semantics.c (finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
	in OpenMP, require pointer or reference to pointer type rather than
	pointer or array or reference to pointer or array type. Handle
	OMP_CLAUSE_USE_DEVICE_ADDR.
	* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
gcc/testsuite/
	* c-c++-common/gomp/target-data-1.c (foo): Use use_device_addr clause
	instead of use_device_ptr clause where required by OpenMP 5.0, add
	further tests for both use_device_ptr and use_device_addr clauses.
libgomp/
	* testsuite/libgomp.c/target-18.c (struct S): New type.
	(foo): Use use_device_addr clause instead of use_device_ptr clause
	where required by OpenMP 5.0, add further tests for both use_device_ptr
	and use_device_addr clauses.
	* testsuite/libgomp.c++/target-9.C (struct S): New type.
	(foo): Use use_device_addr clause instead of use_device_ptr clause
	where required by OpenMP 5.0, add further tests for both use_device_ptr
	and use_device_addr clauses.  Add t and u arguments.
	(main): Adjust caller.

From-SVN: r274159
2019-08-07 09:27:10 +02:00
Jakub Jelinek
d81ab49d05 tree.h (OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV): Rename to ...
* tree.h (OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV): Rename to ...
	(OMP_CLAUSE_LASTPRIVATE_LOOP_IV): ... this.  Adjust comment.
	* gimplify.c (gimple_add_tmp_var): In SIMD contexts, turn addressable
	new vars into GOVD_PRIVATE rather than GOVD_LOCAL.
	(gimplify_omp_for): Don't do C++ random access iterator clause
	adjustments on combined constructs from OMP_LOOP.  For OMP_LOOP,
	don't predetermine the artificial iterator in case of C++ random
	access iterators as lastprivate, but private.  For OMP_LOOP, force
	bind expr around simd body and force for_pre_body before the
	construct.  Use OMP_CLAUSE_LASTPRIVATE_LOOP_IV instead of
	OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV.
	(gimplify_omp_loop): Add firstprivate clauses on OMP_PARALLEL for
	diff var of C++ random access iterators.  Handle
	OMP_CLAUSE_FIRSTPRIVATE.  For OMP_CLAUSE_LASTPRIVATE_LOOP_IV, if
	not outermost also add OMP_CLAUSE_FIRSTPRIVATE, and in both cases
	clear OMP_CLAUSE_LASTPRIVATE_LOOP_IV on the lastprivate clause
	on the OMP_FOR and OMP_DISTRIBUTE constructs if any.
	* omp-low.c (lower_rec_input_clauses): For
	OMP_CLAUSE_LASTPRIVATE_LOOP_IV on simd copy construct the private
	variables instead of default constructing them.
	(lower_lastprivate_clauses): Use OMP_CLAUSE_LASTPRIVATE_LOOP_IV
	instead of OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV and move the
	is_taskloop_ctx check from the assert to the guarding condition.
gcc/cp/
	* parser.c (cp_parser_omp_for_loop): For OMP_LOOP, ignore parallel
	clauses and predetermine iterator as lastprivate.
	* semantics.c (handle_omp_for_class_iterator): Use
	OMP_CLAUSE_LASTPRIVATE_LOOP_IV instead of
	OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV, set it for lastprivate also
	on OMP_LOOP construct.  If a clause is missing for class iterator
	on OMP_LOOP, add firstprivate clause, and if there is private
	clause, turn it into firstprivate too.
	(finish_omp_for): Formatting fix.  For OMP_LOOP, adjust
	OMP_CLAUSE_LASTPRIVATE_LOOP_IV clause CP_CLAUSE_INFO, so that it
	uses copy ctor instead of default ctor.
	* cp-gimplify.c (cp_gimplify_expr): Handle OMP_LOOP like
	OMP_DISTRIBUTE etc.
	(cp_fold_r): Likewise.
	(cp_genericize_r): Likewise.
	(cxx_omp_finish_clause): Also finish lastprivate clause with
	OMP_CLAUSE_LASTPRIVATE_LOOP_IV flag.
	* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_BIND.
	(tsubst_omp_for_iterator): For OMP_LOOP, ignore parallel
	clauses and predetermine iterator as lastprivate.
	* constexpr.c (potential_constant_expression_1): Handle OMP_LOOP
	like OMP_DISTRIBUTE etc.
libgomp/
	* testsuite/libgomp.c++/loop-13.C: New test.
	* testsuite/libgomp.c++/loop-14.C: New test.
	* testsuite/libgomp.c++/loop-15.C: New test.

From-SVN: r274138
2019-08-06 09:26:32 +02:00
Jakub Jelinek
c3ac76aa40 re PR middle-end/91301 (ICE in omp_add_variable on random access iterator distribute parallel for private (iterator))
PR middle-end/91301
	* gimplify.c (gimplify_omp_for): If for class iterator on
	distribute parallel for there is no data sharing clause
	on inner_for_stmt, look for private clause on combined
	parallel too and if found, move it to inner_for_stmt.

	* testsuite/libgomp.c++/for-27.C: New test.

From-SVN: r273922
2019-07-31 09:49:56 +02:00
Steven G. Kargl
8dc63166e0 arith.c (gfc_convert_integer, [...]): Move to ...
2019-07-23  Steven G. Kargl  <kargl@gcc.gnu.org>

	* arith.c (gfc_convert_integer, gfc_convert_real, gfc_convert_complex):
	Move to ...
	* primary.c (convert_integer, convert_real, convert_complex): ... here.
	Rename and make static functions.
	(match_integer_constant): Use convert_integer
	(match_real_constant): Use convert_real.
	(match_complex_constant: Use convert_complex.
	* arith.h (gfc_convert_integer, gfc_convert_real, gfc_convert_complex):
	Remove prototypes.
	* array.c (match_array_cons_element): A BOZ cannot be a data 
	statement value.  Jump to a common exit point.
	* check.c (gfc_invalid_boz): New function.  Emit error or warning
	for a BOZ in an invalid context.
	(boz_args_check): Move to top of file to prevent need of forward
	declaration.
	(is_boz_constant): New function.  Check that BOZ expr is constant.
	(gfc_b	z2real): New function. In-place conversion of BOZ literal
	constant to REAL in accordance to F2018.
	(gfc_boz2int): New function. In-place conversion of BOZ literal
 	onstant to INTEGER in accordance to F2018.
	(gfc_check_achar, gfc_check_char, gfc_check_float): Use gfc_invalid_boz.
	Convert BOZ as needed.
	(gfc_check_bge_bgt_ble_blt): Enforce F2018 requirements on BGE, 
	BGT, BLE, and BLT intrinsic functions.
	(gfc_check_cmplx): Re-organize to check kind, if present, first.
	Convert BOZ real and/or imaginary parts as needed in accordance to
	F2018.
	(gfc_check_complex):  Use gfc_invalid_boz.  Convert BOZ as needed.
	(gfc_check_dcmplx, gfc_check_dble ): Convert BOZ as needed.
	(gfc_check_dshift):  Make dshift[lr] conform to F2018 standard.
	 gfc_check_float (gfc_expr *a)
	(gfc_check_iand_ieor_ior):  Make IAND, IEOR, and IOR conform to 
	F2018 standard.
	(gfc_check_int): Conform to F2018 standard.
	(gfc_check_intconv): Deprecate SHORT and LONG aliases for INT2 and
	INT.  Simply return for a BOZ argument. See gfc_simplify_intconv.
	(gfc_check_merge_bits): Make MERGE_BITS conform to Fortran 2018
	standard.
	(gfc_check_real): Remove incorrect comment. Check kind, if present,
	first.  Simply return for a BOZ argument. See gfc_simplify_real.
	(gfc_check_and): Re-do error handling for BOZ arguments.  Remove
	special casing ts.type != BT_INTEGER or BT_LOGICAL.
	* decl.c (match_old_style_init): Check for BOZ in old-style
	initialization.  Issue error or warning depending on
	-fallow-invalid-boz option.  Issue error if variable is not an
	INTEGER or REAL and the value is BOZ.
	* expr.c (gfc_copy_expr): Copy a BT_BOZ gfc_expr.
	(gfc_check_assign): Re-do error handling for a BOZ in an assignment
	statement.  Do in-place conversion of RHS based on LHS type of
	INTEGER or REAL.
	* gfortran.h (gfc_expr): Add a boz component.  Remove is_boz component.
	(gfc_boz2int, gfc_boz2real, gfc_invalid_boz): New prototypes.
	* interface.c (gfc_extend_assign): Guard against replacing an 
	intrinsic involving a BOZ literal constant on RHS.
	* invoke.texi: Doument -fallow-invalid-boz.
	* lang.opt: New option. -fallow-invalid-boz.
	* libgfortran.h (bt): Elevate BOZ to a basic type.
	* misc.c (gfc_basic_typename, gfc_typename): Translate BT_BOZ to BOZ.
	* primary.c (convert_integer, convert_real, convert_complex): to here.
	Rename and make static functions.
	* primary.c(match_boz_constant): Rewrite parsing of a BOZ. Re-do
	error handling.  Deprecate 'X' for hexidecimal and postfix notation.
	Use -fallow-invalid-boz and gfc_invalid_boz to accept deprecated code.
	* resolve.c (resolve_ordinary_assign): Rework a RHS that is a
	BOZ literal constant.  Use gfc_invalid_boz to allow previous
	nonstandard behavior.  Remove range checking of BOZ conversion.
	* simplify.c (convert_boz): Remove function.
	(simplify_cmplx): Remove conversion of BOZ constants, because
	conversion is done in gfc_check_cmplx.
	(gfc_simplify_float): Remove conversion of BOZ constant, because
	conversion is done in gfc_check_float.
	(simplify_intconv): Use gfc_boz2int to convert BOZ to INTEGER.
	Remove range checking for BOZ conversion.
	(gfc_simplify_real): Use k, if present, to determine kind.  Convert
	BOZ to REAL.  Remove range checking for BOZ conversion.
	target-memory.c (gfc_convert_boz): Rewrite to deal with convert of
	a BOZ to a REAL value.

2019-07-23  Steven G. Kargl  <kargl@gcc.gnu.org>

	* gfortran.dg/achar_5.f90: Fix for new BOZ handling.
	* arithmetic_overflow_1.f90: Ditto.
	* gfortran.dg/boz_11.f90: Ditto.
	* gfortran.dg/boz_12.f90: Ditto.
	* gfortran.dg/boz_4.f90: Ditto.
	* gfortran.dg/boz_5.f90: Ditto.
	* gfortran.dg/boz_6.f90: Ditto.
	* gfortran.dg/boz_7.f90: Ditto.
	* gfortran.dg/boz_8.f90: Ditto.
	* gfortran.dg/dec_structure_6.f90: Ditto.
	* gfortran.dg/dec_union_1.f90: Ditto.
	* gfortran.dg/dec_union_2.f90: Ditto.
	* gfortran.dg/dec_union_5.f90: Ditto.
	* gfortran.dg/dshift_3.f90: Ditto.
	* gfortran.dg/gnu_logical_2.f90: Ditto.
	* gfortran.dg/int_conv_1.f90: Ditto.
	* gfortran.dg/ishft_1.f90: Ditto.
	* gfortran.dg/nan_4.f90: Ditto.
	* gfortran.dg/no_range_check_3.f90: Ditto.
	* gfortran.dg/pr16433.f: Ditto.
	* gfortran.dg/pr44491.f90: Ditto.
	* gfortran.dg/pr58027.f90: Ditto.
	* gfortran.dg/pr81509_2.f90: Ditto.
	* gfortran.dg/unf_io_convert_1.f90: Ditto.
	* gfortran.dg/unf_io_convert_2.f90: Ditto.
	* gfortran.fortran-torture/execute/intrinsic_fraction_exponent.f90:
	Ditto.
	* gfortran.fortran-torture/execute/intrinsic_mvbits.f90: Ditto.
	* gfortran.fortran-torture/execute/intrinsic_nearest.f90: Ditto.
	* gfortran.fortran-torture/execute/seq_io.f90: Ditto.
	* gfortran.dg/gnu_logical_1.F: Delete test.
	* gfortran.dg/merge_bits_3.f90: New test.
	* gfortran.dg/merge_bits_3.f90: Ditto.
	* gfortran.dg/boz_int.f90: Ditto.
	* gfortran.dg/boz_bge.f90: Ditto.
	* gfortran.dg/boz_complex_1.f90: Ditto.
	* gfortran.dg/boz_complex_2.f90: Ditto.
	* gfortran.dg/boz_complex_3.f90: Ditto.
	* gfortran.dg/boz_dble.f90: Ditto.
	* gfortran.dg/boz_dshift_1.f90: Ditto.
	* gfortran.dg/boz_dshift_2.f90: Ditto.
	* gfortran.dg/boz_float_1.f90: Ditto.
	* gfortran.dg/boz_float_2.f90: Ditto.
	* gfortran.dg/boz_float_3.f90: Ditto.
	* gfortran.dg/boz_iand_1.f90: Ditto.
	* gfortran.dg/boz_iand_2.f90: Ditto.

2019-07-23  Steven G. Kargl  <kargl@gcc.gnu.org>

	* testsuite/libgomp.fortran/reduction4.f90: Update BOZ usage
	* testsuite/libgomp.fortran/reduction5.f90: Ditto.

From-SVN: r273747
2019-07-23 21:43:21 +00:00
Jakub Jelinek
554a530ff8 tree.def (OMP_LOOP): New tree code.
* tree.def (OMP_LOOP): New tree code.
	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_BIND.
	(enum omp_clause_bind_kind): New enum.
	(struct tree_omp_clause): Add subcode.bind_kind.
	* tree.h (OMP_LOOP_CHECK): Rename to ...
	(OMP_LOOPING_CHECK): ... this.
	(OMP_FOR_BODY, OMP_FOR_CLAUSES, OMP_FOR_INIT, OMP_FOR_COND,
	OMP_FOR_INCR, OMP_FOR_PRE_BODY, OMP_FOR_ORIG_DECLS): Use
	OMP_LOOPING_CHECK instead of OMP_LOOP_CHECK.
	(OMP_CLAUSE_BIND_KIND): Define.
	* tree.c (omp_clause_num_ops, omp_clause_code_name): Add
	bind clause entries.
	(walk_tree_1): Handle OMP_CLAUSE_BIND.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	(dump_generic_node): Handle OMP_LOOP.
	* gimplify.c (enum omp_region_type): Add ORT_IMPLICIT_TARGET.
	(in_omp_construct): New variable.
	(is_gimple_stmt): Handle OMP_LOOP.
	(gimplify_scan_omp_clauses): For lastprivate don't set
	check_non_private if code == OMP_LOOP.  For reduction clause
	on OMP_LOOP combined with parallel or teams propagate as shared
	on the combined construct.  Handle OMP_CLAUSE_BIND.
	(gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_BIND.
	(gimplify_omp_for): Pass OMP_LOOP instead of OMP_{FOR,DISTRIBUTE}
	for constructs from a loop construct to gimplify_scan_omp_clauses.
	Don't predetermine iterator linear on OMP_SIMD from loop construct.
	(replace_reduction_placeholders, gimplify_omp_loop): New functions.
	(gimplify_omp_workshare): Use ORT_IMPLICIT_TARGET instead of trying
	to match the implicit ORT_TARGET construct around whole body.
	Temporarily clear in_omp_construct when processing body.
	(gimplify_expr): Handle OMP_LOOP.  For OMP_MASTER, OMP_TASKGROUP
	etc. temporarily set in_omp_construct when processing body.
	(gimplify_body): Create ORT_IMPLICIT_TARGET instead of ORT_TARGET.
	* omp-low.c (struct omp_context): Add loop_p.
	(build_outer_var_ref): Treat ctx->loop_p similarly to simd construct
	in that the original var might be private.
	(scan_sharing_clauses): Handle OMP_CLAUSE_BIND.
	(check_omp_nesting_restrictions): Adjust nesting restrictions for
	addition of loop construct.
	(scan_omp_1_stmt): Allow setjmp inside of loop construct.
gcc/c-family/
	* c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_LOOP.
	(enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_BIND.
	* c-pragma.c (omp_pragmas_simd): Add PRAGMA_OMP_LOOP entry.
	* c-common.h (enum c_omp_clause_split): Add C_OMP_CLAUSE_SPLIT_LOOP.
	* c-omp.c (c_omp_split_clauses): Add support for 4 new combined
	constructs with the loop construct.
gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Handle bind clause.
	(c_parser_omp_clause_bind): New function.
	(c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_BIND.
	(OMP_LOOP_CLAUSE_MASK): Define.
	(c_parser_omp_loop): New function.
	(c_parser_omp_parallel, c_parser_omp_teams): Handle parsing of
	loop combined with parallel or teams.
	(c_parser_omp_construct): Handle PRAGMA_OMP_LOOP.
	* c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_BIND.
gcc/cp/
	* cp-tree.h (OMP_FOR_GIMPLIFYING_P): Use OMP_LOOPING_CHECK
	instead of OMP_LOOP_CHECK.
	* parser.c (cp_parser_omp_clause_name): Handle bind clause.
	(cp_parser_omp_clause_bind): New function.
	(cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_BIND.
	(OMP_LOOP_CLAUSE_MASK): Define.
	(cp_parser_omp_loop): New function.
	(cp_parser_omp_parallel, cp_parser_omp_teams): Handle parsing of
	loop combined with parallel or teams.
	(cp_parser_omp_construct): Handle PRAGMA_OMP_LOOP.
	(cp_parser_pragma): Likewise.
	* pt.c (tsubst_expr): Handle OMP_LOOP.
	* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_BIND.
gcc/testsuite/
	* c-c++-common/gomp/cancel-1.c: Adjust expected diagnostic wording.
	* c-c++-common/gomp/clauses-1.c (foo, baz, bar): Add order(concurrent)
	clause where allowed.  Add combined constructs with loop with all
	possible clauses.
	(qux): New function.
	* c-c++-common/gomp/loop-1.c: New test.
	* c-c++-common/gomp/loop-2.c: New test.
	* c-c++-common/gomp/loop-3.c: New test.
	* c-c++-common/gomp/loop-4.c: New test.
	* c-c++-common/gomp/loop-5.c: New test.
	* c-c++-common/gomp/order-3.c: Adjust expected diagnostic wording.
	* c-c++-common/gomp/simd-setjmp-1.c: New test.
	* c-c++-common/gomp/teams-2.c: Adjust expected diagnostic wording.
libgomp/
	* testsuite/libgomp.c-c++-common/loop-1.c: New test.

From-SVN: r273621
2019-07-20 13:21:42 +02:00
Jakub Jelinek
6cda84b509 tree-vect-stmts.c (scan_operand_equal_p): Look through MEM_REF with SSA_NAME address of POINTER_PLUS_EXPR.
* tree-vect-stmts.c (scan_operand_equal_p): Look through MEM_REF
	with SSA_NAME address of POINTER_PLUS_EXPR.  Handle MULT_EXPR
	and casts in offset when different, both through gimple stmts
	and through trees.  Rewritten using loops to minimize code duplication
	for each operand.

	* g++.dg/vect/simd-6.cc: Replace xfail with target x86.
	* g++.dg/vect/simd-9.cc: Likewise.

	* testsuite/libgomp.c++/scan-13.C: Replace xfail with target x86.
	* testsuite/libgomp.c++/scan-16.C: Likewise.

From-SVN: r273249
2019-07-09 00:11:59 +02:00
Jakub Jelinek
6f67abcdb0 omp-low.c (lower_rec_input_clauses): For lastprivate clauses in ctx->for_simd_scan_phase simd copy the outer var to...
* omp-low.c (lower_rec_input_clauses): For lastprivate clauses in
	ctx->for_simd_scan_phase simd copy the outer var to the privatized
	variable(s).  For conditional lastprivate look through outer
	GIMPLE_OMP_SCAN context.
	(lower_omp_1): For conditional lastprivate look through outer
	GIMPLE_OMP_SCAN context.

	* testsuite/libgomp.c/scan-19.c: New test.
	* testsuite/libgomp.c/scan-20.c: New test.

From-SVN: r273169
2019-07-06 23:58:01 +02:00
Jakub Jelinek
1f52d1a8b5 omp-low.c (struct omp_context): Add for_simd_scan_phase member.
* omp-low.c (struct omp_context): Add for_simd_scan_phase member.
	(maybe_lookup_ctx): Add forward declaration.
	(omp_find_scan): Likewise.  Walk into body of simd if composited
	with worksharing loop.
	(scan_omp_simd_scan): New function.
	(scan_omp_1_stmt): Call it.
	(lower_rec_simd_input_clauses): Don't create rvar nor rvar2 if
	ctx->for_simd_scan_phase.
	(lower_rec_input_clauses): Do much less work for inscan reductions
	in ctx->for_simd_scan_phase is_simd regions.
	(lower_omp_scan): Set is_simd also on simd constructs composited
	with worksharing loop, unless ctx->for_simd_scan_phase.  Never emit
	a sorry message.  Don't change GIMPLE_OMP_SCAN stmts into nops and
	emit their body after in simd constructs composited with worksharing
	loop.
	(lower_omp_for_scan): Handle worksharing loop composited with simd.

	* c-c++-common/gomp/scan-4.c: Don't expect sorry message.

	* testsuite/libgomp.c/scan-11.c: New test.
	* testsuite/libgomp.c/scan-12.c: New test.
	* testsuite/libgomp.c/scan-13.c: New test.
	* testsuite/libgomp.c/scan-14.c: New test.
	* testsuite/libgomp.c/scan-15.c: New test.
	* testsuite/libgomp.c/scan-16.c: New test.
	* testsuite/libgomp.c/scan-17.c: New test.
	* testsuite/libgomp.c/scan-18.c: New test.
	* testsuite/libgomp.c++/scan-9.C: New test.
	* testsuite/libgomp.c++/scan-10.C: New test.
	* testsuite/libgomp.c++/scan-11.C: New test.
	* testsuite/libgomp.c++/scan-12.C: New test.
	* testsuite/libgomp.c++/scan-13.C: New test.
	* testsuite/libgomp.c++/scan-14.C: New test.
	* testsuite/libgomp.c++/scan-15.C: New test.
	* testsuite/libgomp.c++/scan-16.C: New test.

From-SVN: r273157
2019-07-06 09:53:48 +02:00
Jakub Jelinek
2f03073ff2 omp-expand.c (expand_omp_for_static_nochunk): Don't emit GOMP_loop_start at the start of second worksharing loop in a scan.
* omp-expand.c (expand_omp_for_static_nochunk): Don't emit
	GOMP_loop_start at the start of second worksharing loop in a scan.
	For nowait, don't emit GOMP_loop_end_nowait at the end of first
	worksharing loop in a scan even if there are conditional lastprivates,
	and do emit GOMP_loop_end_nowait at the end of second worksharing loop.

	* testsuite/libgomp.c/scan-9.c: New test.
	* testsuite/libgomp.c/scan-10.c: New test.

From-SVN: r273095
2019-07-04 23:40:56 +02:00
Jakub Jelinek
2f6bb511d1 tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__SCANTEMP_ clause.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__SCANTEMP_
	clause.
	* tree.h (OMP_CLAUSE_DECL): Use OMP_CLAUSE__SCANTEMP_ instead of
	OMP_CLAUSE__CONDTEMP_ as range's upper bound.
	(OMP_CLAUSE__SCANTEMP__ALLOC, OMP_CLAUSE__SCANTEMP__CONTROL): Define.
	* tree.c (omp_clause_num_ops, omp_clause_code_name): Add
	OMP_CLAUSE__SCANTEMP_ entry.
	(walk_tree_1): Handle OMP_CLAUSE__SCANTEMP_.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Likewise.
	* omp-general.h (struct omp_for_data): Add have_scantemp and
	have_nonctrl_scantemp members.
	* omp-general.c (omp_extract_for_data): Initialize them.
	* omp-low.c (struct omp_context): Add scan_exclusive member.
	(scan_omp_1_stmt): Don't unnecessarily mask gimple_omp_for_kind
	result again with GF_OMP_FOR_KIND_MASK.  Initialize also
	ctx->scan_exclusive.
	(lower_rec_simd_input_clauses): Use ctx->scan_exclusive instead
	of !ctx->scan_inclusive.
	(lower_rec_input_clauses): Simplify gimplification of dtors using
	gimplify_and_add.  For non-is_simd test OMP_CLAUSE_REDUCTION_INSCAN
	rather than rvarp.  Handle OMP_CLAUSE_REDUCTION_INSCAN in worksharing
	loops.  Don't add barrier for reduction_omp_orig_ref if
	ctx->scan_??xclusive.
	(lower_reduction_clauses): Don't do anything for ctx->scan_??xclusive.
	(lower_omp_scan): Use ctx->scan_exclusive instead
	of !ctx->scan_inclusive.  Handle worksharing loops with inscan
	reductions.  Use new_vard != new_var instead of repeated
	omp_is_reference calls.
	(omp_find_scan, lower_omp_for_scan): New functions.
	(lower_omp_for): Call lower_omp_for_scan for worksharing loops with
	inscan reductions.
	* omp-expand.c (expand_omp_scantemp_alloc): New function.
	(expand_omp_for_static_nochunk): Handle fd->have_nonctrl_scantemp
	and fd->have_scantemp.

	* c-c++-common/gomp/scan-3.c (f1): Don't expect a sorry message.
	* c-c++-common/gomp/scan-5.c (foo): Likewise.

	* testsuite/libgomp.c++/scan-1.C: New test.
	* testsuite/libgomp.c++/scan-2.C: New test.
	* testsuite/libgomp.c++/scan-3.C: New test.
	* testsuite/libgomp.c++/scan-4.C: New test.
	* testsuite/libgomp.c++/scan-5.C: New test.
	* testsuite/libgomp.c++/scan-6.C: New test.
	* testsuite/libgomp.c++/scan-7.C: New test.
	* testsuite/libgomp.c++/scan-8.C: New test.
	* testsuite/libgomp.c/scan-1.c: New test.
	* testsuite/libgomp.c/scan-2.c: New test.
	* testsuite/libgomp.c/scan-3.c: New test.
	* testsuite/libgomp.c/scan-4.c: New test.
	* testsuite/libgomp.c/scan-5.c: New test.
	* testsuite/libgomp.c/scan-6.c: New test.
	* testsuite/libgomp.c/scan-7.c: New test.
	* testsuite/libgomp.c/scan-8.c: New test.

From-SVN: r272958
2019-07-03 07:03:58 +02:00
Thomas Schwinge
85fca03a09 Test cases to verify OpenACC 'firstprivate' mappings
gcc/testsuite/
	* c-c++-common/goacc/firstprivate-mappings-1.c: New file.
	* g++.dg/goacc/firstprivate-mappings-1.C: Likewise.
	libgomp/
	* testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C: New file.
	* testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c:
	Likewise.

From-SVN: r272451
2019-06-19 00:15:03 +02:00
Thomas Schwinge
bd194a51d4 Add missing results check in 'libgomp.fortran/allocatable3.f90'
libgomp/
	* testsuite/libgomp.fortran/allocatable3.f90: Add missing results
	check.

From-SVN: r272449
2019-06-19 00:14:43 +02:00
Cesar Philippidis
6652161ef3 Add 'libgomp.oacc-fortran/allocatable-array-1.f90'
libgomp/
	* testsuite/libgomp.oacc-fortran/allocatable-array-1.f90: New
	file.

From-SVN: r272448
2019-06-19 00:14:34 +02:00
Thomas Schwinge
4017da8d1c [PR90743] Fortran 'allocatable' with OpenACC data/OpenMP 'target' 'map' clauses
Test what OpenMP 5.0 has to say on this topic.  And, do the same for OpenACC.

	libgomp/
	PR fortran/90743
	* oacc-parallel.c (GOACC_parallel_keyed): Handle NULL mapping
	case.
	* testsuite/libgomp.fortran/target-allocatable-1-1.f90: New file.
	* testsuite/libgomp.fortran/target-allocatable-1-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/allocatable-1-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/allocatable-1-2.f90: Likewise.

From-SVN: r272447
2019-06-19 00:14:24 +02:00
Thomas Schwinge
6f7c1f6502 [PR90861] Document status quo for OpenACC 'declare' not cleaning up for VLAs
gcc/testsuite/
	PR testsuite/90861
	* c-c++-common/goacc/declare-pr90861.c: New file.
	libgomp/
	PR testsuite/90861
	* testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Update.

From-SVN: r272446
2019-06-19 00:14:14 +02:00
Thomas Schwinge
3a37d6f68c [PR90862] OpenACC 'declare' ICE when nested inside another construct
gcc/
	PR middle-end/90862
	* omp-low.c (check_omp_nesting_restrictions): Handle
	GF_OMP_TARGET_KIND_OACC_DECLARE.
	gcc/testsuite/
	PR middle-end/90862
	* c-c++-common/goacc/declare-1.c: Update.
	* c-c++-common/goacc/declare-2.c: Likewise.
	libgomp/
	PR middle-end/90862
	* testsuite/libgomp.oacc-c-c++-common/declare-1.c: Update.

From-SVN: r272444
2019-06-19 00:13:54 +02:00
Tom de Vries
f45ce17d98 [openacc, parloops] Fix SIGSEGV in oacc_entry_exit_ok_1
When compiling the test-case with r268755, we run into a SIGSEGV in
oacc_entry_exit_ok_1 when trying to dereference a NULL red:
...
                      struct reduction_info *red;
                      red = reduction_phi (reduction_list, use_stmt);
                      tree val = PHI_RESULT (red->keep_res);
...

Fix this by handling ref == NULL.

Bootstrapped and reg-tested on x86_64.
Build and reg-tested on x86_64 with nvptx accelerator.

2019-06-16  Tom de Vries  <tdevries@suse.de>

	PR tree-optimization/89376
	* tree-parloops.c (oacc_entry_exit_ok_1): Handle red == NULL.

	* testsuite/libgomp.oacc-c-c++-common/pr89376.c: New test.

From-SVN: r272338
2019-06-16 07:47:15 +00:00
Tom de Vries
00908992f2 [nvptx, libgomp] Update pr85381-{2,4}.c test-cases
After the fix for "PR tree-optimization/89713 - Assume loop with an exit is
finite" ( r272234 ) empty oacc loops are removed before expand.

Update pr85381-{2,4}.c accordingly.

2019-06-15  Tom de Vries  <tdevries@suse.de>

	PR tree-optimization/89713
	* testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Expect no bar.sync.
	* testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Same.

From-SVN: r272324
2019-06-15 12:01:30 +00:00
Jakub Jelinek
211b7533bf re PR middle-end/90779 (Fortran array initialization in offload regions)
PR middle-end/90779
	* gimplify.c: Include omp-offload.h and context.h.
	(gimplify_bind_expr): Add "omp declare target" attributes
	to static block scope variables inside of target region or target
	functions.

	* c-c++-common/goacc/routine-5.c (func2): Don't expect error for
	static block scope variable in #pragma acc routine.

	* testsuite/libgomp.c/pr90779.c: New test.
	* testsuite/libgomp.fortran/pr90779.f90: New test.

From-SVN: r272322
2019-06-15 09:09:04 +02:00
Tom de Vries
120a01d160 [openacc] Disable pass_thread_jumps for IFN_UNIQUE
If we compile the openacc testcase with -fopenacc -O2, we run into a SIGSEGV
or assert.  The root cause for this is that pass_thread_jumps breaks the
invariant that OACC_FORK and OACC_JOIN mark the start and end of a
single-entry-single-exit region.

Fix this by bailing out when encountering an IFN_UNIQUE in
thread_jumps::profitable_jump_thread_path.

Bootstrapped and reg-tested on x86_64.
Build and reg-tested libgomp on x86_64 with nvptx accelerator.

2019-06-15  Tom de Vries  <tdevries@suse.de>

	PR tree-optimization/90009
	* tree-ssa-threadbackward.c (thread_jumps::profitable_jump_thread_path):
	Return NULL if bb contains IFN_UNIQUE.

	* testsuite/libgomp.oacc-c-c++-common/pr90009.c: New test.

From-SVN: r272321
2019-06-15 07:06:19 +00:00
Feng Xue
c29c92c789 PR tree-optimization/89713 - Assume loop with an exit is finite
gcc/ChangeLog:

        * doc/invoke.texi (-ffinite-loops): Document new option.
        * common.opt (-ffinite-loops): New option.
        * tree-ssa-dce.c (mark_stmt_if_obviously_necessary): Mark
        IFN_GOACC_LOOP calls as necessary.
        * tree-ssa-loop-niter.c (finite_loop_p): Assume loop with an exit
        is finite.
        * omp-offload.c (oacc_xform_loop): Skip lowering if return value of
        IFN_GOACC_LOOP call is not used.
        * opts.c (default_options_table): Enable -ffinite-loops at -O2+.

gcc/testsuite/ChangeLog:

        * g++.dg/tree-ssa/empty-loop.C: New test.
        * gcc.dg/tree-ssa/dce-2.c: New test.
        * gcc.dg/const-1.c: Add -fno-finite-loops option.
        * gcc.dg/graphite/graphite.exp: Likewise.
        * gcc.dg/loop-unswitch-1.c: Likewise.
        * gcc.dg/predict-9.c: Likewise.
        * gcc.dg/pure-2.c: Likewise.
        * gcc.dg/tree-ssa/20040211-1.c: Likewise.
        * gcc.dg/tree-ssa/loop-10.c: Likewise.
        * gcc.dg/tree-ssa/split-path-6.c: Likewise.
        * gcc.dg/tree-ssa/ssa-thread-12.c: Likewise.

libgomp/ChangeLog:

        * testsuite/libgomp.oacc-c-c++-common/pr84955-1.c: New test.

From-SVN: r272234
2019-06-13 04:17:42 +00:00
Jakub Jelinek
ce9c4ec3c5 re PR target/90811 ([nvptx] ptxas error on OpenMP offloaded code)
PR target/90811
	* config/nvptx/nvptx.c (nvptx_output_softstack_switch): Use and.b%d
	instead of and.u%d.

	* testsuite/libgomp.c/pr90811.c: New test.

From-SVN: r272161
2019-06-11 18:40:10 +02:00
Jakub Jelinek
28b3a77ca0 omp-low.c (lower_rec_input_clauses): For lastprivate conditional references...
* omp-low.c (lower_rec_input_clauses): For lastprivate conditional
	references, lookup in in hash map MEM_REF operand instead of the
	MEM_REF itself.
	(lower_omp_1): When looking for lastprivate conditional assignments,
	handle MEM_REFs with REFERENCE_TYPE operands.

	* testsuite/libgomp.c++/lastprivate-conditional-1.C: New test.
	* testsuite/libgomp.c++/lastprivate-conditional-2.C: New test.

From-SVN: r271948
2019-06-05 09:37:40 +02:00
Jakub Jelinek
7855700e63 gimplify.c (gimplify_scan_omp_clauses): Don't sorry_at on lastprivate conditional on combined for simd.
* gimplify.c (gimplify_scan_omp_clauses): Don't sorry_at on lastprivate
	conditional on combined for simd.
	* omp-low.c (struct omp_context): Add combined_into_simd_safelen0
	member.
	(lower_rec_input_clauses): For gimple_omp_for_combined_into_p max_vf 1
	constructs, don't remove lastprivate_conditional_map, but instead set
	ctx->combined_into_simd_safelen0 and adjust hash_map, so that it points
	to parent construct temporaries.
	(lower_lastprivate_clauses): Handle ctx->combined_into_simd_safelen0
	like !ctx->lastprivate_conditional_map.
	(lower_omp_1) <case GIMPLE_ASSIGN>: If up->combined_into_simd_safelen0,
	use up->outer context instead of up.
	* omp-expand.c (expand_omp_for_generic): Perform cond_var bump even if
	gimple_omp_for_combined_p.
	(expand_omp_for_static_nochunk): Likewise.
	(expand_omp_for_static_chunk): Add forgotten cond_var bump that was
	probably moved over into expand_omp_for_generic rather than being copied
	there.
gcc/cp/
	* cp-tree.h (CP_OMP_CLAUSE_INFO): Allow for any clauses up to _condvar_
	instead of only up to linear.
gcc/testsuite/
	* c-c++-common/gomp/lastprivate-conditional-2.c (foo): Don't expect
	a sorry_at on any of the clauses.
libgomp/
	* testsuite/libgomp.c-c++-common/lastprivate-conditional-7.c: New test.
	* testsuite/libgomp.c-c++-common/lastprivate-conditional-8.c: New test.
	* testsuite/libgomp.c-c++-common/lastprivate-conditional-9.c: New test.
	* testsuite/libgomp.c-c++-common/lastprivate-conditional-10.c: New test.

From-SVN: r271907
2019-06-04 14:49:03 +02:00
Rainer Orth
a7155c2e0b Generalize getconf _NPROCESSORS_ONLN
libgomp:
	* configure.ac: Call AX_COUNT_CPUS.
	Substitute CPU_COUNT.
	* testsuite/Makefile.am (check-am): Use CPU_COUNT as processor
	count fallback.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.
	* Makefile.in, testsuite/Makefile.in: Regenerate.

	config:
	* ax_count_cpus.m4: New file.

From-SVN: r271769
2019-05-30 09:06:48 +00:00
Jakub Jelinek
7e47198b80 gimplify.c (struct gimplify_omp_ctx): Add clauses member.
* gimplify.c (struct gimplify_omp_ctx): Add clauses member.
	(gimplify_scan_omp_clauses): Initialize ctx->clauses.
	(gimplify_adjust_omp_clauses_1): Transform lastprivate conditional
	explicit clause on combined parallel into implicit shared clause.
	(gimplify_adjust_omp_clauses): Move lastprivate conditional clause
	and firstprivate if the decl has one too from combined parallel to
	the worksharing construct.
gcc/testsuite/
	* c-c++-common/gomp/lastprivate-conditional-2.c (foo): Don't expect
	sorry on lastprivate conditional on parallel for.
	* c-c++-common/gomp/lastprivate-conditional-3.c (foo): Add tests for
	lastprivate conditional warnings on parallel for constructs.
	* c-c++-common/gomp/lastprivate-conditional-4.c: New test.
libgomp/
	* testsuite/libgomp.c-c++-common/lastprivate_conditional_4.c: Rename
	to ...
	* testsuite/libgomp.c-c++-common/lastprivate-conditional-4.c: ... this.
	* testsuite/libgomp.c-c++-common/lastprivate-conditional-5.c: New test.
	* testsuite/libgomp.c-c++-common/lastprivate-conditional-6.c: New test.

From-SVN: r271733
2019-05-29 09:51:43 +02:00
Jakub Jelinek
8e7757ba17 gimplify.c (gimplify_scan_omp_clauses): Allow lastprivate conditional on sections construct.
* gimplify.c (gimplify_scan_omp_clauses): Allow lastprivate conditional
	on sections construct.
	* omp-low.c (lower_lastprivate_conditional_clauses): Handle sections
	construct.
	(lower_omp_sections): Handle lastprivate conditional.
	(lower_omp_1) <case GIMPLE_ASSIGN>: Handle sections construct with
	lastprivate_conditional_map.
	* omp-expand.c (expand_omp_sections): Handle lastprivate conditional.
libgomp/
	* testsuite/libgomp.c-c++-common/lastprivate_conditional_4.c: New test.

From-SVN: r271673
2019-05-27 23:33:37 +02:00
Jakub Jelinek
36c7a3fff9 omp-low.c (lower_omp_1): Look through ordered...
* omp-low.c (lower_omp_1) <case GIMPLE_ASSIGN>: Look through ordered,
	critical, taskgroup and section regions when looking for a region
	with non-NULL lastprivate_conditional_map.

	* testsuite/libgomp.c-c++-common/lastprivate-conditional-3.c: New test.

From-SVN: r271672
2019-05-27 23:31:40 +02:00
Jakub Jelinek
fcfb80325f re PR libgomp/90641 (libgomp.c-c++-common/lastprivate-conditional-1.c etc FAIL)
PR libgomp/90641
	* work.c (gomp_init_work_share): Instead of aligning final ordered
	value to multiples of long long alignment, align to that the
	first part (ordered team ids) and if inline_ordered_team_ids
	is not on a long long alignment boundary within the structure,
	use __alignof__ (long long) - 1 pad size always.
	* loop.c (GOMP_loop_start): Fix *mem computation if
	inline_ordered_team_ids is not aligned on long long alignment boundary
	within the structure.
	* loop-ull.c (GOMP_loop_ull_start): Likewise.
	* sections.c (GOMP_sections2_start): Likewise.

From-SVN: r271671
2019-05-27 23:27:00 +02:00
Jakub Jelinek
6c7ae8c56f tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__CONDTEMP_.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__CONDTEMP_.
	* tree.h (OMP_CLAUSE_DECL): Use OMP_CLAUSE__CONDTEMP_ instead of
	OMP_CLAUSE__REDUCTEMP_.
	* tree.c (omp_clause_num_ops, omp_clause_code_name): Add
	OMP_CLAUSE__CONDTEMP_.
	(walk_tree_1): Handle OMP_CLAUSE__CONDTEMP_.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Likewise.
	* gimplify.c (enum gimplify_omp_var_data): Use hexadecimal constants
	instead of decimal.  Add GOVD_LASTPRIVATE_CONDITIONAL.
	(gimplify_scan_omp_clauses): Don't reject lastprivate conditional
	on OMP_FOR.
	(gimplify_omp_for): Warn and disable conditional modifier from
	lastprivate on loop iterators.
	* omp-general.h (struct omp_for_data): Add lastprivate_conditional
	member.
	* omp-general.c (omp_extract_for_data): Initialize it.
	* omp-low.c (struct omp_context): Add lastprivate_conditional_map
	member.
	(delete_omp_context): Delete it.
	(lower_lastprivate_conditional_clauses): New function.
	(lower_lastprivate_clauses): Add BODY_P and CSTMT_LIST arguments,
	handle lastprivate conditional clauses.
	(lower_reduction_clauses): Add CLIST argument, emit it into
	the critical section if any.
	(lower_omp_sections): Adjust lower_lastprivate_clauses and
	lower_reduction_clauses callers.
	(lower_omp_for_lastprivate): Add CLIST argument, pass it through
	to lower_lastprivate_clauses.
	(lower_omp_for): Call lower_lastprivate_conditional_clauses, adjust
	lower_omp_for_lastprivate and lower_reduction_clauses callers, emit
	clist into a critical section if not emitted there already by
	lower_reduction_clauses.
	(lower_omp_taskreg, lower_omp_teams): Adjust lower_reduction_clauses
	callers.
	(lower_omp_1): Handle GIMPLE_ASSIGNs storing into lastprivate
	conditional variables.
	* omp-expand.c (determine_parallel_type): Punt if OMP_CLAUSE__CONDTEMP_
	clause is present.
	(expand_omp_for_generic, expand_omp_for_static_nochunk,
	expand_omp_for_static_chunk): Handle lastprivate conditional.
	(expand_omp_for): Handle fd.lastprivate_conditional like
	fd.have_reductemp.
gcc/testsuite/
	* c-c++-common/gomp/lastprivate-conditional-2.c (foo): Don't expect
	sorry for omp for.
	* c-c++-common/gomp/lastprivate-conditional-3.c: New test.
libgomp/
	* testsuite/libgomp.c-c++-common/lastprivate-conditional-1.c: New test.
	* testsuite/libgomp.c-c++-common/lastprivate-conditional-2.c: New test.

From-SVN: r271610
2019-05-24 23:31:59 +02:00
Jakub Jelinek
b5c26449f3 re PR libgomp/90585 (libgomp hsa plugin ftbfs in the x32 multilib variant)
PR libgomp/90585
	* plugin/plugin-hsa.c: Include gstdint.h.  Include inttypes.h only if
	HAVE_INTTYPES_H is defined.
	(print_uint64_t): New typedef.
	(PRIu64): Define if HAVE_INTTYPES_H is not defined.
	(print_kernel_dispatch, run_kernel): Use PRIu64 macro instead of
	"lu", cast uint64_t HSA_DEBUG and fprintf arguments to print_uint64_t.
	(release_kernel_dispatch): Likewise.  Cast shadow->debug to uintptr_t
	before casting to void *.
	* plugin/plugin-nvptx.c: Include gstdint.h instead of stdint.h.
	* oacc-mem.c: Don't include config.h nor stdint.h.
	* target.c: Don't include config.h.
	* oacc-cuda.c: Likewise.
	* oacc-host.c: Don't include stdint.h.

From-SVN: r271597
2019-05-24 10:59:37 +02:00
Jakub Jelinek
3e03ed6626 re PR libgomp/90527 (alloc.c:72:7: error: implicit declaration of function ‘posix_memalign’)
PR libgomp/90527
	* alloc.c (_GNU_SOURCE): Define.

From-SVN: r271438
2019-05-20 23:29:17 +02:00
Thomas Schwinge
5fae049dc2 OpenACC Profiling Interface (incomplete)
libgomp/
	* acc_prof.h: New file.
	* oacc-profiling.c: Likewise.
	* Makefile.am (nodist_libsubinclude_HEADERS, libgomp_la_SOURCES):
	Add these, respectively.
	* Makefile.in: Regenerate.
	* env.c (initialize_env): Call goacc_profiling_initialize.
	* oacc-plugin.c (GOMP_PLUGIN_goacc_thread)
	(GOMP_PLUGIN_goacc_profiling_dispatch): New functions.
	* oacc-plugin.h (GOMP_PLUGIN_goacc_thread)
	(GOMP_PLUGIN_goacc_profiling_dispatch): Declare.
	* libgomp.map (OACC_2.5.1): Add acc_prof_lookup,
	acc_prof_register, acc_prof_unregister, and acc_register_library.
	(GOMP_PLUGIN_1.3): Add GOMP_PLUGIN_goacc_profiling_dispatch, and
	GOMP_PLUGIN_goacc_thread.
	* oacc-int.h (struct goacc_thread): Add prof_info, api_info,
	prof_callbacks_enabled members.
	(goacc_prof_enabled, goacc_profiling_initialize)
	(_goacc_profiling_dispatch_p, _goacc_profiling_setup_p)
	(goacc_profiling_dispatch): Declare.
	(GOACC_PROF_ENABLED, GOACC_PROFILING_DISPATCH_P)
	(GOACC_PROFILING_SETUP_P): Define.
	* oacc-async.c (acc_async_test, acc_async_test_all, acc_wait)
	(acc_wait_async, acc_wait_all, acc_wait_all_async): Update for
	OpenACC Profiling Interface.
	* oacc-cuda.c (acc_get_current_cuda_device)
	(acc_get_current_cuda_context, acc_get_cuda_stream)
	(acc_set_cuda_stream): Likewise.
	* oacc-init.c (acc_init_1, goacc_attach_host_thread_to_device)
	(acc_init, acc_set_device_type, acc_get_device_type)
	(acc_get_device_num, goacc_lazy_initialize): Likewise.
	* oacc-mem.c (acc_malloc, acc_free, memcpy_tofrom_device)
	(acc_deviceptr, acc_hostptr, acc_is_present, acc_map_data)
	(acc_unmap_data, present_create_copy, delete_copyout)
	(update_dev_host): Likewise.
	* oacc-parallel.c (GOACC_parallel_keyed, GOACC_data_start)
	(GOACC_data_end, GOACC_enter_exit_data, GOACC_update, GOACC_wait):
	Likewise.
	* plugin/plugin-nvptx.c (nvptx_exec, nvptx_alloc, nvptx_free)
	(GOMP_OFFLOAD_openacc_exec, GOMP_OFFLOAD_openacc_async_exec):
	Likewise.
	* libgomp.texi: Update.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-dispatch-1.c: New
	file.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-valid_bytes-1.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-version-1.c:
	Likewise.

From-SVN: r271346
2019-05-17 21:13:36 +02:00
Chung-Lin Tang
1f4c5b9bb2 2019-05-13 Chung-Lin Tang <cltang@codesourcery.com>
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

	libgomp/
	* libgomp-plugin.h (struct goacc_asyncqueue): Declare.
	(struct goacc_asyncqueue_list): Likewise.
	(goacc_aq): Likewise.
	(goacc_aq_list): Likewise.
	(GOMP_OFFLOAD_openacc_register_async_cleanup): Remove.
	(GOMP_OFFLOAD_openacc_async_test): Remove.
	(GOMP_OFFLOAD_openacc_async_test_all): Remove.
	(GOMP_OFFLOAD_openacc_async_wait): Remove.
	(GOMP_OFFLOAD_openacc_async_wait_async): Remove.
	(GOMP_OFFLOAD_openacc_async_wait_all): Remove.
	(GOMP_OFFLOAD_openacc_async_wait_all_async): Remove.
	(GOMP_OFFLOAD_openacc_async_set_async): Remove.
	(GOMP_OFFLOAD_openacc_exec): Adjust declaration.
	(GOMP_OFFLOAD_openacc_cuda_get_stream): Likewise.
	(GOMP_OFFLOAD_openacc_cuda_set_stream): Likewise.
	(GOMP_OFFLOAD_openacc_async_exec): Declare.
	(GOMP_OFFLOAD_openacc_async_construct): Declare.
	(GOMP_OFFLOAD_openacc_async_destruct): Declare.
	(GOMP_OFFLOAD_openacc_async_test): Declare.
	(GOMP_OFFLOAD_openacc_async_synchronize): Declare.
	(GOMP_OFFLOAD_openacc_async_serialize): Declare.
	(GOMP_OFFLOAD_openacc_async_queue_callback): Declare.
	(GOMP_OFFLOAD_openacc_async_host2dev): Declare.
	(GOMP_OFFLOAD_openacc_async_dev2host): Declare.

	* libgomp.h (struct acc_dispatch_t): Define 'async' sub-struct.
	(gomp_acc_insert_pointer): Adjust declaration.
	(gomp_copy_host2dev): New declaration.
	(gomp_copy_dev2host): Likewise.
	(gomp_map_vars_async): Likewise.
	(gomp_unmap_tgt): Likewise.
	(gomp_unmap_vars_async): Likewise.
	(gomp_fini_device): Likewise.

	* oacc-async.c (get_goacc_thread): New function.
	(get_goacc_thread_device): New function.
	(lookup_goacc_asyncqueue): New function.
	(get_goacc_asyncqueue): New function.
	(acc_async_test): Adjust code to use new async design.
	(acc_async_test_all): Likewise.
	(acc_wait): Likewise.
	(acc_wait_async): Likewise.
	(acc_wait_all): Likewise.
	(acc_wait_all_async): Likewise.
	(goacc_async_free): New function.
	(goacc_init_asyncqueues): Likewise.
	(goacc_fini_asyncqueues): Likewise.
	* oacc-cuda.c (acc_get_cuda_stream): Adjust code to use new async
	design.
	(acc_set_cuda_stream): Likewise.
	* oacc-host.c (host_openacc_exec): Adjust parameters, remove 'async'.
	(host_openacc_register_async_cleanup): Remove.
	(host_openacc_async_exec): New function.
	(host_openacc_async_test): Adjust parameters.
	(host_openacc_async_test_all): Remove.
	(host_openacc_async_wait): Remove.
	(host_openacc_async_wait_async): Remove.
	(host_openacc_async_wait_all): Remove.
	(host_openacc_async_wait_all_async): Remove.
	(host_openacc_async_set_async): Remove.
	(host_openacc_async_synchronize): New function.
	(host_openacc_async_serialize): New function.
	(host_openacc_async_host2dev): New function.
	(host_openacc_async_dev2host): New function.
	(host_openacc_async_queue_callback): New function.
	(host_openacc_async_construct): New function.
	(host_openacc_async_destruct): New function.
	(struct gomp_device_descr host_dispatch): Remove initialization of old
	interface, add intialization of new async sub-struct.
	* oacc-init.c (acc_shutdown_1): Adjust to use gomp_fini_device.
	(goacc_attach_host_thread_to_device): Remove old async code usage.
	* oacc-int.h (goacc_init_asyncqueues): New declaration.
	(goacc_fini_asyncqueues): Likewise.
	(goacc_async_copyout_unmap_vars): Likewise.
	(goacc_async_free): Likewise.
	(get_goacc_asyncqueue): Likewise.
	(lookup_goacc_asyncqueue): Likewise.

	* oacc-mem.c (memcpy_tofrom_device): Adjust code to use new async
	design.
	(present_create_copy): Adjust code to use new async design.
	(delete_copyout): Likewise.
	(update_dev_host): Likewise.
	(gomp_acc_insert_pointer): Add async parameter, adjust code to use new
	async design.
	(gomp_acc_remove_pointer): Adjust code to use new async design.
	* oacc-parallel.c (GOACC_parallel_keyed): Adjust code to use new async
	design.
	(GOACC_enter_exit_data): Likewise.
	(goacc_wait): Likewise.
	(GOACC_update): Likewise.
	* oacc-plugin.c (GOMP_PLUGIN_async_unmap_vars): Change to assert fail
	when called, warn as obsolete in comment.

	* target.c (goacc_device_copy_async): New function.
	(gomp_copy_host2dev): Remove 'static', add goacc_asyncqueue parameter,
	add goacc_device_copy_async case.
	(gomp_copy_dev2host): Likewise.
	(gomp_map_vars_existing): Add goacc_asyncqueue parameter, adjust code.
	(gomp_map_pointer): Likewise.
	(gomp_map_fields_existing): Likewise.
	(gomp_map_vars_internal): New always_inline function, renamed from
	gomp_map_vars.
	(gomp_map_vars): Implement by calling gomp_map_vars_internal.
	(gomp_map_vars_async): Implement by calling gomp_map_vars_internal,
	passing goacc_asyncqueue argument.
	(gomp_unmap_tgt): Remove static, add attribute_hidden.
	(gomp_unref_tgt): New function.
	(gomp_unmap_vars_internal): New always_inline function, renamed from
	gomp_unmap_vars.
	(gomp_unmap_vars): Implement by calling gomp_unmap_vars_internal.
	(gomp_unmap_vars_async): Implement by calling
	gomp_unmap_vars_internal, passing goacc_asyncqueue argument.
	(gomp_fini_device): New function.
	(gomp_exit_data): Adjust gomp_copy_dev2host call.
	(gomp_load_plugin_for_device): Remove old interface, adjust to load
	new async interface.
	(gomp_target_fini): Adjust code to call gomp_fini_device.

	* plugin/plugin-nvptx.c (struct cuda_map): Remove.
	(struct ptx_stream): Remove.
	(struct nvptx_thread): Remove current_stream field.
	(cuda_map_create): Remove.
	(cuda_map_destroy): Remove.
	(map_init): Remove.
	(map_fini): Remove.
	(map_pop): Remove.
	(map_push): Remove.
	(struct goacc_asyncqueue): Define.
	(struct nvptx_callback): Define.
	(struct ptx_free_block): Define.
	(struct ptx_device): Remove null_stream, active_streams, async_streams,
	stream_lock, and next fields.
	(enum ptx_event_type): Remove.
	(struct ptx_event): Remove.
	(ptx_event_lock): Remove.
	(ptx_events): Remove.
	(init_streams_for_device): Remove.
	(fini_streams_for_device): Remove.
	(select_stream_for_async): Remove.
	(nvptx_init): Remove ptx_events and ptx_event_lock references.
	(nvptx_attach_host_thread_to_device): Remove CUDA_ERROR_NOT_PERMITTED
	case.
	(nvptx_open_device): Add free_blocks initialization, remove
	init_streams_for_device call.
	(nvptx_close_device): Remove fini_streams_for_device call, add
	free_blocks destruct code.
	(event_gc): Remove.
	(event_add): Remove.
	(nvptx_exec): Adjust parameters and code.
	(nvptx_free): Likewise.
	(nvptx_host2dev): Remove.
	(nvptx_dev2host): Remove.
	(nvptx_set_async): Remove.
	(nvptx_async_test): Remove.
	(nvptx_async_test_all): Remove.
	(nvptx_wait): Remove.
	(nvptx_wait_async): Remove.
	(nvptx_wait_all): Remove.
	(nvptx_wait_all_async): Remove.
	(nvptx_get_cuda_stream): Remove.
	(nvptx_set_cuda_stream): Remove.
	(GOMP_OFFLOAD_alloc): Adjust code.
	(GOMP_OFFLOAD_free): Likewise.
	(GOMP_OFFLOAD_openacc_register_async_cleanup): Remove.
	(GOMP_OFFLOAD_openacc_exec): Adjust parameters and code.
	(GOMP_OFFLOAD_openacc_async_test_all): Remove.
	(GOMP_OFFLOAD_openacc_async_wait): Remove.
	(GOMP_OFFLOAD_openacc_async_wait_async): Remove.
	(GOMP_OFFLOAD_openacc_async_wait_all): Remove.
	(GOMP_OFFLOAD_openacc_async_wait_all_async): Remove.
	(GOMP_OFFLOAD_openacc_async_set_async): Remove.
	(cuda_free_argmem): New function.
	(GOMP_OFFLOAD_openacc_async_exec): New plugin hook function.
	(GOMP_OFFLOAD_openacc_create_thread_data): Adjust code.
	(GOMP_OFFLOAD_openacc_cuda_get_stream): Adjust code.
	(GOMP_OFFLOAD_openacc_cuda_set_stream): Adjust code.
	(GOMP_OFFLOAD_openacc_async_construct): New plugin hook function.
	(GOMP_OFFLOAD_openacc_async_destruct): New plugin hook function.
	(GOMP_OFFLOAD_openacc_async_test): Remove and re-implement.
	(GOMP_OFFLOAD_openacc_async_synchronize): New plugin hook function.
	(GOMP_OFFLOAD_openacc_async_serialize): New plugin hook function.
	(GOMP_OFFLOAD_openacc_async_queue_callback): New plugin hook function.
	(cuda_callback_wrapper): New function.
	(cuda_memcpy_sanity_check): New function.
	(GOMP_OFFLOAD_host2dev): Remove and re-implement.
	(GOMP_OFFLOAD_dev2host): Remove and re-implement.
	(GOMP_OFFLOAD_openacc_async_host2dev): New plugin hook function.
	(GOMP_OFFLOAD_openacc_async_dev2host): New plugin hook function.

From-SVN: r271128
2019-05-13 13:32:00 +00:00
Thomas Schwinge
da2d30c199 Address compiler diagnostics in libgomp.oacc-c-c++-common/pr87835.c
source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c: In function 'main':
    source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c:45: warning: ignoring #pragma loop gang [-Wunknown-pragmas]
       45 |     #pragma loop gang
          |
    source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c:19:7: warning: unused variable 'b' [-Wunused-variable]
       19 |   int b[n];
          |       ^

	libgomp/
	PR target/87835
	* testsuite/libgomp.oacc-c-c++-common/pr87835.c: Update.

From-SVN: r271004
2019-05-08 12:01:30 +02:00
Thomas Schwinge
2bbbfa4e28 Clean up libgomp GCC 5 legacy support
libgomp/
	* oacc-parallel.c: Add comments to legacy entry points (GCC 5).

From-SVN: r270901
2019-05-06 10:49:55 +02:00
Kevin Buettner
bbf1efe1b4 team.c (gomp_team_start): Initialize pool->threads[0].
libgomp/ChangeLog:

	* team.c (gomp_team_start): Initialize pool->threads[0].

From-SVN: r269971
2019-03-27 18:30:44 +00:00
Thomas Schwinge
b03d721a62 [libgomp] In OpenACC testing, by default only build for the offload target that we're actually going to test
... to avoid compilation overhead, and to keep simple '-foffload=[...]'
handling in test cases.

	libgomp/
	* testsuite/libgomp.oacc-c++/c++.exp: Specify
	"-foffload=$offload_target".
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
	* testsuite/lib/libgomp.exp
	(check_effective_target_openacc_nvidia_accel_configured): Remove,
	as (conceptually) merged into
	check_effective_target_openacc_nvidia_accel_selected.  Adjust all
	users.

From-SVN: r269109
2019-02-22 11:51:35 +01:00
Thomas Schwinge
0a0384b43a [libgomp] In OpenACC testing, cycle though all offload targets
... instead of through offload plugins.

	libgomp/
	* plugin/configfrag.ac: Populate and AC_SUBST offload_targets.
	* testsuite/libgomp-test-support.exp.in: Adjust.
	* testsuite/lib/libgomp.exp: Likewise.  Don't populate
	openacc_device_types_s.
	(offload_target_to_openacc_device_type): New proc.
	* testsuite/libgomp.oacc-c++/c++.exp: Adjust.
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
	* Makefile.in: Regenerate.
	* configure: Likewise.
	* testsuite/Makefile.in: Likewise.

From-SVN: r269108
2019-02-22 11:51:20 +01:00
Thomas Schwinge
ee332b4a9a [libgomp] Clarify difference between offload target, offload plugin, and OpenACC device type
libgomp/
	* plugin/configfrag.ac: Populate and AC_SUBST offload_plugins
	instead of offload_targets, and AC_DEFINE_UNQUOTED OFFLOAD_PLUGINS
	instead of OFFLOAD_TARGETS.
	* target.c (gomp_target_init): Adjust.
	* testsuite/libgomp-test-support.exp.in: Likewise.
	* testsuite/lib/libgomp.exp: Likewise.  Populate
	openacc_device_types_s instead of offload_targets_s_openacc.
	(check_effective_target_openacc_nvidia_accel_selected)
	(check_effective_target_openacc_host_selected): Adjust.
	* testsuite/libgomp.oacc-c++/c++.exp: Likewise.
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
	* Makefile.in: Regenerate.
	* config.h.in: Likewise.
	* configure: Likewise.
	* testsuite/Makefile.in: Likewise.

From-SVN: r269107
2019-02-22 11:51:05 +01:00
Thomas Schwinge
1241136c71 [libgomp] In OpenACC offloading testing, be more explicit in what is supported, and what is not, or why not
libgomp/
	* testsuite/lib/libgomp.exp: Error out for unknown offload target.
	* testsuite/libgomp.oacc-c++/c++.exp: Likewise.  Report if
	"offloading: supported, but hardware not accessible".
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.

From-SVN: r269106
2019-02-22 11:50:49 +01:00
Chung-Lin Tang
19695f4d99 re PR c/87924 (OpenACC wait clauses without async-arguments)
2019-02-19  Chung-Lin Tang <cltang@codesourcery.com>

	PR c/87924
	gcc/c/
	* c-parser.c (c_parser_oacc_clause_wait): Add representation of wait
	clause without argument as 'wait (GOMP_ASYNC_NOVAL)', adjust comments.

	gcc/cp/
	* parser.c (cp_parser_oacc_clause_wait): Add representation of wait
	clause without argument as 'wait (GOMP_ASYNC_NOVAL)', adjust comments.

	gcc/fortran/
	* openmp.c (gfc_match_omp_clauses): Add representation of wait clause
	without argument as 'wait (GOMP_ASYNC_NOVAL)'.

	libgomp/
	* oacc-parallel.c (GOACC_parallel_keyed): Remove condition on call to
	goacc_wait().
	(goacc_wait): Handle ACC_ASYNC_NOVAL case, remove goacc_thread() call
	and related adjustment.

	Reviewed-by: Thomas Schwinge  <thomas@codesourcery.com>

From-SVN: r269016
2019-02-19 14:10:15 +00:00