PR fortran/92756
* trans-openmp.c (gfc_trans_omp_teams): Wrap OMP_TEAMS body into a
BIND_EXPR with a forced BLOCK.
* gfortran.dg/gomp/teams1.f90: New test.
* testsuite/libgomp.fortran/teams1.f90: New test.
* testsuite/libgomp.fortran/teams2.f90: New test.
From-SVN: r278956
Check that function arguments of type acc_device_t
are valid enumeration values in all publicly visible
functions from oacc-init.c.
2019-12-03 Frederik Harwath <frederik@codesourcery.com>
libgomp/
* oacc-init.c (acc_known_device_type): Add function.
(unknown_device_type_error): Add function.
(name_of_acc_device_t): Change to call unknown_device_type_error
on unknown type.
(resolve_device): Use acc_known_device_type.
(acc_init): Fail if acc_device_t argument is not valid.
(acc_shutdown): Likewise.
(acc_get_num_devices): Likewise.
(acc_set_device_type): Likewise.
(acc_get_device_num): Likewise.
(acc_set_device_num): Likewise.
(acc_on_device): Add comment that argument validity is not checked.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
From-SVN: r278937
2019-12-03 Andrew Stubbs <ams@codesourcery.com>
libgomp/
* testsuite/lib/libgomp.exp (offload_target_to_openacc_device_type):
Recognize amdgcn.
(check_effective_target_openacc_amdgcn_accel_present): New proc.
(check_effective_target_openacc_amdgcn_accel_selected): New proc.
* testsuite/libgomp.oacc-c++/c++.exp: Add support for amdgcn.
* testsuite/libgomp.oacc-c/c.exp: Likewise.
* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
From-SVN: r278935
PR c++/60228
* parser.c (cp_parser_omp_declare_reduction_exprs): If
processing_template_decl, wrap the combiner or initializer
into EXPR_STMT.
* decl.c (start_preparsed_function): Don't start a lambda scope
for DECL_OMP_DECLARE_REDUCTION_P functions.
(finish_function): Don't finish a lambda scope for
DECL_OMP_DECLARE_REDUCTION_P functions, nor cp_fold_function
them nor cp_genericize them.
* mangle.c (decl_mangling_context): Look through
DECL_OMP_DECLARE_REDUCTION_P functions.
* semantics.c (expand_or_defer_fn_1): For DECL_OMP_DECLARE_REDUCTION_P
functions, use tentative linkage, don't keep their bodies with
-fkeep-inline-functions and return false at the end.
* g++.dg/gomp/openmp-simd-2.C: Don't expect bodies for
DECL_OMP_DECLARE_REDUCTION_P functions.
* testsuite/libgomp.c++/udr-20.C: New test.
* testsuite/libgomp.c++/udr-21.C: New test.
From-SVN: r278831
libgomp/
* testsuite/lib/libgomp.exp
(check_effective_target_offload_target_nvptx): New proc.
* testsuite/libgomp.fortran/target-print-1.f90: Use it with
'dg-skip-if'.
* testsuite/libgomp.oacc-fortran/print-1.f90: Likewise.
* testsuite/libgomp.fortran/target-print-1-nvptx.f90: New file.
* testsuite/libgomp.oacc-fortran/print-1-nvptx.f90: Likewise.
From-SVN: r278779
PR libgomp/92511
libgomp/
* oacc-mem.c (present_create_copy): Fix device pointer return value in
case of "present" subarray. Use tgt->tgt_start instead of tgt->to_free
in non-present/create case.
(delete_copyout): Change error condition to fail only on copies outside
of mapped block. Adjust error message accordingly.
* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/lib-20.c: Adjust expected error
message.
* testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-22.c: Allow test to pass now.
* testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
From-SVN: r278514
Add flags to find libatomic in build-tree testing, fixing a catastrophic
libgomp testsuite failure with `riscv*-*-linux*' targets, which imply
`-latomic' with the `-pthread' GCC option, implied in turn by the
`-fopenacc' and `-fopenmp' options, removing failures like:
.../bin/riscv64-linux-gnu-ld: cannot find -latomic
collect2: error: ld returned 1 exit status
compiler exited with status 1
FAIL: libgomp.c/../libgomp.c-c++-common/atomic-18.c (test for excess errors)
Excess errors:
.../bin/riscv64-linux-gnu-ld: cannot find -latomic
UNRESOLVED: libgomp.c/../libgomp.c-c++-common/atomic-18.c compilation failed to produce executable
and bringing overall test results for the `riscv64-linux-gnu' target
(here with the `x86_64-linux-gnu' host and RISC-V QEMU in the Linux user
emulation mode as the target board) from:
=== libgomp Summary ===
# of expected passes 90
# of unexpected failures 3267
# of expected failures 2
# of unresolved testcases 3247
# of unsupported tests 548
to:
=== libgomp Summary ===
# of expected passes 6834
# of unexpected failures 4
# of expected failures 4
# of unsupported tests 518
libgomp/
* testsuite/lib/libgomp.exp (libgomp_init): Add flags to find
libatomic in build-tree testing.
From-SVN: r278505
2019-11-15 Andrew Stubbs <ams@codesourcery.com>
libgomp/
* testsuite/libgomp.c/target-print-1.c: New file.
* testsuite/libgomp.fortran/target-print-1.f90: New file.
* testsuite/libgomp.oacc-c/print-1.c: New file.
* testsuite/libgomp.oacc-fortran/print-1.f90: New file.
From-SVN: r278284
2019-11-13 Andrew Stubbs <ams@codesourcery.com>
libgomp/
* config/gcn/team.c (gomp_gcn_enter_kernel): Set up the team arena
and use team_malloc variants.
(gomp_gcn_exit_kernel): Use team_free.
* libgomp.h (TEAM_ARENA_SIZE): Define.
(TEAM_ARENA_START): Define.
(TEAM_ARENA_FREE): Define.
(TEAM_ARENA_END): Define.
(team_malloc): New function.
(team_malloc_cleared): New function.
(team_free): New function.
* team.c (gomp_new_team): Initialize and use team_malloc.
(free_team): Use team_free.
(gomp_free_thread): Use team_free.
(gomp_pause_host): Use team_free.
* work.c (gomp_init_work_share): Use team_malloc.
(gomp_fini_work_share): Use team_free.
From-SVN: r278136
The `serial' construct (cf. section 2.5.3 of the OpenACC 2.6 standard)
is equivalent to a `parallel' construct with clauses `num_gangs(1)
num_workers(1) vector_length(1)' implied.
These clauses are therefore not supported with the `serial'
construct. All the remaining clauses accepted with `parallel' are also
accepted with `serial'.
The `serial' construct is implemented like `parallel', except for
hardcoding dimensions rather than taking them from the relevant
clauses, in `expand_omp_target'.
Separate codes are used to denote the `serial' construct throughout the
middle end, even though the mapping of `serial' to an equivalent
`parallel' construct could have been done in the individual language
frontends. In particular, this allows to distinguish between compute
constructs in warnings, error messages, dumps etc.
2019-11-12 Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Frederik Harwath <frederik@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
gcc/
* gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_SERIAL
enumeration constant.
(is_gimple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
(is_gimple_omp_offloaded): Likewise.
* gimplify.c (omp_region_type): Add ORT_ACC_SERIAL enumeration
constant. Adjust the value of ORT_NONE accordingly.
(is_gimple_stmt): Handle OACC_SERIAL.
(oacc_default_clause): Handle ORT_ACC_SERIAL.
(gomp_needs_data_present): Likewise.
(gimplify_adjust_omp_clauses): Likewise.
(gimplify_omp_workshare): Handle OACC_SERIAL.
(gimplify_expr): Likewise.
* omp-expand.c (expand_omp_target):
Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
(build_omp_regions_1, omp_make_gimple_edges): Likewise.
* omp-low.c (is_oacc_parallel): Rename function to...
(is_oacc_parallel_or_serial): ... this.
Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
(scan_sharing_clauses): Adjust accordingly.
(scan_omp_for): Likewise.
(lower_oacc_head_mark): Likewise.
(convert_from_firstprivate_int): Likewise.
(lower_omp_target): Likewise.
(check_omp_nesting_restrictions): Handle
GF_OMP_TARGET_KIND_OACC_SERIAL.
(lower_oacc_reductions): Likewise.
(lower_omp_target): Likewise.
* tree.def (OACC_SERIAL): New tree code.
* tree-pretty-print.c (dump_generic_node): Handle OACC_SERIAL.
* doc/generic.texi (OpenACC): Document OACC_SERIAL.
gcc/c-family/
* c-pragma.h (pragma_kind): Add PRAGMA_OACC_SERIAL enumeration
constant.
* c-pragma.c (oacc_pragmas): Add "serial" entry.
gcc/c/
* c-parser.c (OACC_SERIAL_CLAUSE_MASK): New macro.
(c_parser_oacc_kernels_parallel): Rename function to...
(c_parser_oacc_compute): ... this. Handle PRAGMA_OACC_SERIAL.
(c_parser_omp_construct): Update accordingly.
gcc/cp/
* constexpr.c (potential_constant_expression_1): Handle
OACC_SERIAL.
* parser.c (OACC_SERIAL_CLAUSE_MASK): New macro.
(cp_parser_oacc_kernels_parallel): Rename function to...
(cp_parser_oacc_compute): ... this. Handle PRAGMA_OACC_SERIAL.
(cp_parser_omp_construct): Update accordingly.
(cp_parser_pragma): Handle PRAGMA_OACC_SERIAL. Fix alphabetic
order.
* pt.c (tsubst_expr): Handle OACC_SERIAL.
gcc/fortran/
* gfortran.h (gfc_statement): Add ST_OACC_SERIAL_LOOP,
ST_OACC_END_SERIAL_LOOP, ST_OACC_SERIAL and ST_OACC_END_SERIAL
enumeration constants.
(gfc_exec_op): Add EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL
enumeration constants.
* match.h (gfc_match_oacc_serial): New prototype.
(gfc_match_oacc_serial_loop): Likewise.
* dump-parse-tree.c (show_omp_node, show_code_node): Handle
EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL.
* match.c (match_exit_cycle): Handle EXEC_OACC_SERIAL_LOOP.
* openmp.c (OACC_SERIAL_CLAUSES): New macro.
(gfc_match_oacc_serial_loop): New function.
(gfc_match_oacc_serial): Likewise.
(oacc_is_loop): Handle EXEC_OACC_SERIAL_LOOP.
(resolve_omp_clauses): Handle EXEC_OACC_SERIAL.
(oacc_code_to_statement): Handle EXEC_OACC_SERIAL and
EXEC_OACC_SERIAL_LOOP.
(gfc_resolve_oacc_directive): Likewise.
* parse.c (decode_oacc_directive) <'s'>: Add case for "serial"
and "serial loop".
(next_statement): Handle ST_OACC_SERIAL_LOOP and ST_OACC_SERIAL.
(gfc_ascii_statement): Likewise. Handle ST_OACC_END_SERIAL_LOOP
and ST_OACC_END_SERIAL.
(parse_oacc_structured_block): Handle ST_OACC_SERIAL.
(parse_oacc_loop): Handle ST_OACC_SERIAL_LOOP and
ST_OACC_END_SERIAL_LOOP.
(parse_executable): Handle ST_OACC_SERIAL_LOOP and
ST_OACC_SERIAL.
(is_oacc): Handle EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL.
* resolve.c (gfc_resolve_blocks, gfc_resolve_code): Likewise.
* st.c (gfc_free_statement): Likewise.
* trans-openmp.c (gfc_trans_oacc_construct): Handle
EXEC_OACC_SERIAL.
(gfc_trans_oacc_combined_directive): Handle
EXEC_OACC_SERIAL_LOOP.
(gfc_trans_oacc_directive): Handle EXEC_OACC_SERIAL_LOOP and
EXEC_OACC_SERIAL.
* trans.c (trans_code): Likewise.
gcc/testsuite/
* c-c++-common/goacc/parallel-dims.c: New test.
* gfortran.dg/goacc/parallel-dims.f90: New test.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: New test.
* testsuite/libgomp.oacc-fortran/parallel-dims-aux.c: New test.
* testsuite/libgomp.oacc-fortran/parallel-dims.f89: New test.
* testsuite/libgomp.oacc-fortran/parallel-dims-2.f90: New test.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Frederik Harwath <frederik@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>
From-SVN: r278082
OpenACC (cf. OpenACC 2.7, section 2.9.11. "reduction clause";
this was first clarified by OpenACC 2.6) requires that, if a
variable is used in reduction clauses on two nested loops, then
there must be reduction clauses for that variable on all loops
that are nested in between the two loops and all these reduction
clauses must use the same operator.
This commit introduces a check for that property which reports
warnings if it is violated.
2019-11-06 Gergö Barany <gergo@codesourcery.com>
Frederik Harwath <frederik@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
gcc/
* omp-low.c (struct omp_context): New fields
local_reduction_clauses, outer_reduction_clauses.
(new_omp_context): Initialize these.
(scan_sharing_clauses): Record reduction clauses on OpenACC constructs.
(scan_omp_for): Check reduction clauses for incorrect nesting.
gcc/testsuite/
* c-c++-common/goacc/nested-reductions-warn.c: New test.
* c-c++-common/goacc/nested-reductions.c: New test.
* gfortran.dg/goacc/nested-reductions-warn.f90: New test.
* gfortran.dg/goacc/nested-reductions.f90: New test.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c:
Add expected warnings about missing reduction clauses.
* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c:
Likewise.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
From-SVN: r277875
2019-10-25 Cesar Philippidis <cesar@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
gcc/fortran/
* openmp.c (gfc_match_omp_map_clause): Add and pass allow_commons
argument.
(gfc_match_omp_clauses): Update calls to permit common blocks for
OpenACC's copy/copyin/copyout, create/delete, host,
pcopy/pcopy_in/pcopy_out, present_or_copy, present_or_copy_in,
present_or_copy_out, present_or_create and self.
gcc/
* gimplify.c (oacc_default_clause): Privatize fortran common blocks.
(omp_notice_variable): Defer the expansion of DECL_VALUE_EXPR for
common block decls.
gcc/testsuite/
* gfortran.dg/goacc/common-block-1.f90: New test.
* gfortran.dg/goacc/common-block-2.f90: New test.
* gfortran.dg/goacc/common-block-3.f90: New test.
libgomp/
* testsuite/libgomp.oacc-fortran/common-block-1.f90: New test.
* testsuite/libgomp.oacc-fortran/common-block-2.f90: New test.
* testsuite/libgomp.oacc-fortran/common-block-3.f90: New test.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>
From-SVN: r277451
PR libgomp/92028
* target.c (gomp_map_vars_internal): Readd the previous
GOMP_MAP_USE_DEVICE_PTR handling code in the first loop,
though do that just in the !not_found_cnt case.
From-SVN: r276753
2019-10-02 Julian Brown <julian@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
libgomp/
* libgomp.h (OFFSET_INLINED, OFFSET_POINTER, OFFSET_STRUCT): Define.
* target.c (FIELD_TGT_EMPTY): Define.
(gomp_map_val): Use OFFSET_* macros instead of magic constants. Write
as switch instead of list of ifs.
(gomp_map_vars_internal): Use OFFSET_* and FIELD_TGT_EMPTY macros.
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
From-SVN: r276519
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-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-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
* 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
* 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
* 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
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-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
* 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
* 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
* 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
* 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
* 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
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
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
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
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
* 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
* 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
* 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
* 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
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
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
... 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
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
