Commit Graph

936 Commits

Author SHA1 Message Date
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
Jakub Jelinek 8b44f8ec4b re PR c++/88988 (ICE: Segmentation fault (in lookup_name_real_1))
PR c++/88988
	* lambda.c (is_capture_proxy): Don't return true for
	DECL_OMP_PRIVATIZED_MEMBER artificial vars.

	* testsuite/libgomp.c++/pr88988.C: New test.

From-SVN: r268407
2019-01-31 00:28:53 +01:00
Jakub Jelinek 52bfbb69e7 re PR middle-end/89002 (ICE in scan_omp_1_op, at omp-low.c:3166)
PR middle-end/89002
	* gimplify.c (gimplify_omp_for): When adding OMP_CLAUSE_*_GIMPLE_SEQ
	for lastprivate/linear IV, push gimplify context around gimplify_assign
	and, if it needed any temporaries, pop it into a gimple bind around the
	sequence.

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

From-SVN: r268346
2019-01-28 23:34:32 +01:00
Jakub Jelinek be3a87e7b5 re PR middle-end/89002 (ICE in scan_omp_1_op, at omp-low.c:3166)
PR middle-end/89002
	* gimplify.c (gimplify_omp_for): When adding OMP_CLAUSE_*_GIMPLE_SEQ
	for lastprivate/linear IV, push gimplify context around gimplify_assign
	and, if it needed any temporaries, pop it into a gimple bind around the
	sequence.

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

From-SVN: r268345
2019-01-28 23:33:33 +01:00
Richard Biener 497ef4d7f5 re PR testsuite/89064 (libgomp.graphite/force-parallel-5.c fails starting with r268257)
2019-01-28  Richard Biener  <rguenther@suse.de>

	PR testsuite/89064
	PR tree-optimization/86865
	* testsuite/libgomp.graphite/force-parallel-5.c: XFAIL.

From-SVN: r268333
2019-01-28 09:07:30 +00:00
Tom de Vries 738c56d410 [nvptx, libgomp] Fix memleak in GOMP_OFFLOAD_fini_device
I wrote a test-case:
...
int
main (void)
{
  for (unsigned i = 0; i < 128; ++i)
    {
      acc_init (acc_device_nvidia);
      acc_shutdown (acc_device_nvidia);
    }

  return 0;
}
...
and ran it under valgrind.  The only leak location reported with a frequency
of 128, was the allocation of ptx_devices in nvptx_init.

Fix this by freeing ptx_devices in GOMP_OFFLOAD_fini_device, once
instantiated_devices drops to 0.

2019-01-24  Tom de Vries  <tdevries@suse.de>

	* plugin/plugin-nvptx.c (GOMP_OFFLOAD_fini_device): Free ptx_devices
	once instantiated_devices drops to 0.

From-SVN: r268237
2019-01-24 14:12:19 +00:00
Tom de Vries 4a75460b00 [nvptx, libgomp] Fix cuMemAlloc with size zero
Consider test-case:
...
int
main (void)
{
  #pragma acc parallel async
  ;
  #pragma acc parallel async
  ;
  #pragma acc wait

  return 0;
}
...

This fails with:
...
libgomp: cuMemAlloc error: invalid argument
Segmentation fault (core dumped)
...
The cuMemAlloc error is due to the fact that we're try to allocate 0 bytes.

Fix this by preventing calling map_push with size zero argument in nvptx_exec.

This also has the consequence that for the abort-1.c test-case, we end up
calling cuMemFree during map_fini for the struct cuda_map allocated in
map_init, which fails because an abort happened.  Fix this by calling
cuMemFree with CUDA_CALL_NOCHECK in cuda_map_destroy.

2019-01-23  Tom de Vries  <tdevries@suse.de>

	PR target/PR88946
	* plugin/plugin-nvptx.c (cuda_map_destroy): Use CUDA_CALL_NOCHECK for
	cuMemFree.
	(nvptx_exec): Don't call map_push if mapnum == 0.
	* testsuite/libgomp.oacc-c-c++-common/pr88946.c: New test.

From-SVN: r268178
2019-01-23 08:16:56 +00:00
Tom de Vries 4fef8e4d8c [nvptx, libgomp] Fix assert (!s->map->active) in map_fini
There are currently two situations where this assert triggers:
...
libgomp/plugin/plugin-nvptx.c: map_fini: Assertion `!s->map->active' failed.
...

First, in abort-1.c, a parallel region triggering an abort:
...
int
main (void)
{
  #pragma acc parallel
  abort ();

  return 0;
}
...

The abort is detected in nvptx_exec as the CUDA_ERROR_ILLEGAL_INSTRUCTION
return status of the cuStreamSynchronize call after kernel launch, which is
then handled by calling non-returning function GOMP_PLUGIN_fatal.
Consequently, the map_pop in nvptx_exec that in case of cuStreamSynchronize
success would remove or inactive the element added by the map_push earlier in
nvptx_exec, does not trigger.  With the element no longer active, but still
marked active and a member of s->map,  we run into the assert during
GOMP_OFFLOAD_fini_device, which is triggered from atexit handler
gomp_target_fini (which is triggered by the GOMP_PLUGIN_fatal mentioned above
calling exit).

Second, in pr88941.c, an async parallel region without wait:
...
int
main (void)
{
  #pragma acc parallel async
  ;

  /* no #pragma acc wait */
  return 0;
}
...

Because nvptx_exec is handling an async region, it does not call map_pop for
the element added by map_push, but schedules an kernel execution completion
event to call map_pop.  Again, we run into the assert during
GOMP_OFFLOAD_fini_device, which is triggered from atexit handler
gomp_target_fini, but the exit in this case is triggered by returning from main.
So either the kernel is still running, or the kernel has completed but the
corresponding event that is supposed to call map_pop is stuck in the event
queue, waiting for an event_gc.

Fix this by removing the assert, and skipping the freeing of device memory if
the map is still marked active (though in the async case, this is more a
workaround than an fix).

2019-01-23  Tom de Vries  <tdevries@suse.de>

	PR target/88941
	PR target/88939
	* plugin/plugin-nvptx.c (cuda_map_destroy): Handle map->active case.
	(map_fini): Remove "assert (!s->map->active)".
	* testsuite/libgomp.oacc-c-c++-common/pr88941.c: New test.

From-SVN: r268177
2019-01-23 08:16:42 +00:00