aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite
AgeCommit message (Collapse)AuthorFilesLines
2021-02-17openacc: Fix lowering for derived-type mappings through array elementsJulian Brown2-0/+162
This patch fixes lowering of derived-type mappings which select elements of arrays of derived types, and similar. These would previously lead to ICEs. With this change, OpenACC directives can pass through constructs that are no longer recognized by the gimplifier, hence alterations are needed there also. gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Handle element selection for arrays of derived types. gcc/ * gimplify.c (gimplify_scan_omp_clauses): Handle ATTACH_DETACH for non-decls. gcc/testsuite/ * gfortran.dg/goacc/array-with-dt-1.f90: New test. * gfortran.dg/goacc/array-with-dt-3.f90: Likewise. * gfortran.dg/goacc/array-with-dt-4.f90: Likewise. * gfortran.dg/goacc/array-with-dt-5.f90: Likewise. * gfortran.dg/goacc/derived-chartypes-1.f90: Re-enable test. * gfortran.dg/goacc/derived-chartypes-2.f90: Likewise. * gfortran.dg/goacc/derived-classtypes-1.f95: Uncomment previously-broken directives. libgomp/ * testsuite/libgomp.oacc-fortran/derivedtypes-arrays-1.f90: New test. * testsuite/libgomp.oacc-fortran/update-dt-array.f90: Likewise.
2021-02-09openacc: Add XFAILs [PR98979]Julian Brown1-2/+3
This patch adds some XFAILs for PR98979 until the patch to fix them has been approved. See: https://gcc.gnu.org/pipermail/gcc-patches/2021-February/564711.html gcc/testsuite/ PR fortran/98979 * gfortran.dg/goacc/array-with-dt-2.f90: Add expected errors. * gfortran.dg/goacc/derived-chartypes-1.f90: Skip ICEing test. * gfortran.dg/goacc/derived-chartypes-2.f90: Likewise. libgomp/ PR fortran/98979 * testsuite/libgomp.oacc-fortran/array-stride-dt-1.f90: Add expected errors.
2021-02-04openacc: Allow strided arrays in update directivesJulian Brown1-0/+44
OpenACC 3.0 ("2.14.4. Update Directive") states: Noncontiguous subarrays may appear. It is implementation-specific whether noncontiguous regions are updated by using one transfer for each contiguous subregion, or whether the non-contiguous data is packed, transferred once, and unpacked, or whether one or more larger subarrays (no larger than the smallest contiguous region that contains the specified subarray) are updated. This patch relaxes some conditions in the Fortran front-end so that strided accesses are permitted for update directives. gcc/fortran/ * openmp.c (resolve_omp_clauses): Omit OpenACC update in contiguity check and stride-specified error. gcc/testsuite/ * gfortran.dg/goacc/array-with-dt-2.f90: New test. libgomp/ * testsuite/libgomp.oacc-fortran/array-stride-dt-1.f90: New test.
2021-01-19OpenMP/Fortran: Fixes for {use,is}_device_ptrTobias Burnus1-0/+54
gcc/fortran/ChangeLog: PR fortran/98476 * openmp.c (resolve_omp_clauses): Change use_device_ptr to use_device_addr for unless type(c_ptr); check all list item for is_device_ptr. gcc/ChangeLog: PR fortran/98476 * omp-low.c (lower_omp_target): Handle nonpointer is_device_ptr. libgomp/ChangeLog: PR fortran/98476 * testsuite/libgomp.fortran/is_device_ptr-1.f90: New test. gcc/testsuite/ChangeLog: PR fortran/98476 * gfortran.dg/gomp/map-3.f90: Update expected scan-dump-tree. * gfortran.dg/gomp/is_device_ptr-2.f90: New test. * gfortran.dg/gomp/use_device_ptr-1.f90: New test.
2021-01-16openmp: Add support for the OpenMP 5.0 task detach clauseKwok Cheung Yeung12-0/+423
2021-01-16 Kwok Cheung Yeung <kcy@codesourcery.com> gcc/ * builtin-types.def (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT): Rename to... (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR): ...this. Add extra argument. * gimplify.c (omp_default_clause): Ensure that event handle is firstprivate in a task region. (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_DETACH. (gimplify_adjust_omp_clauses): Likewise. * omp-builtins.def (BUILT_IN_GOMP_TASK): Change function type to BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR. * omp-expand.c (expand_task_call): Add GOMP_TASK_FLAG_DETACH to flags if detach clause specified. Add detach argument when generating call to GOMP_task. * omp-low.c (scan_sharing_clauses): Setup data environment for detach clause. (finish_taskreg_scan): Move field for variable containing the event handle to the front of the struct. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_DETACH. Fix ordering. * tree-nested.c (convert_nonlocal_omp_clauses): Handle OMP_CLAUSE_DETACH clause. (convert_local_omp_clauses): Handle OMP_CLAUSE_DETACH clause. * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_DETACH. * tree.c (omp_clause_num_ops): Add entry for OMP_CLAUSE_DETACH. Fix ordering. (omp_clause_code_name): Add entry for OMP_CLAUSE_DETACH. Fix ordering. (walk_tree_1): Handle OMP_CLAUSE_DETACH. gcc/c-family/ * c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_DETACH. Redefine PRAGMA_OACC_CLAUSE_DETACH. gcc/c/ * c-parser.c (c_parser_omp_clause_detach): New. (c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH clause. (OMP_TASK_CLAUSE_MASK): Add mask for PRAGMA_OMP_CLAUSE_DETACH. * c-typeck.c (c_finish_omp_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH clause. Prevent use of detach with mergeable and overriding the data sharing mode of the event handle. gcc/cp/ * parser.c (cp_parser_omp_clause_detach): New. (cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH. (OMP_TASK_CLAUSE_MASK): Add mask for PRAGMA_OMP_CLAUSE_DETACH. * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_DETACH clause. * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_DETACH clause. Prevent use of detach with mergeable and overriding the data sharing mode of the event handle. gcc/fortran/ * dump-parse-tree.c (show_omp_clauses): Handle detach clause. * frontend-passes.c (gfc_code_walker): Walk detach expression. * gfortran.h (struct gfc_omp_clauses): Add detach field. (gfc_c_intptr_kind): New. * openmp.c (gfc_free_omp_clauses): Free detach clause. (gfc_match_omp_detach): New. (enum omp_mask1): Add OMP_CLAUSE_DETACH. (enum omp_mask2): Remove OMP_CLAUSE_DETACH. (gfc_match_omp_clauses): Handle OMP_CLAUSE_DETACH for OpenMP. (OMP_TASK_CLAUSES): Add OMP_CLAUSE_DETACH. (resolve_omp_clauses): Prevent use of detach with mergeable and overriding the data sharing mode of the event handle. * trans-openmp.c (gfc_trans_omp_clauses): Handle detach clause. * trans-types.c (gfc_c_intptr_kind): New. (gfc_init_kinds): Initialize gfc_c_intptr_kind. * types.def (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT): Rename to... (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR): ...this. Add extra argument. gcc/testsuite/ * c-c++-common/gomp/task-detach-1.c: New. * g++.dg/gomp/task-detach-1.C: New. * gcc.dg/gomp/task-detach-1.c: New. * gfortran.dg/gomp/task-detach-1.f90: New. include/ * gomp-constants.h (GOMP_TASK_FLAG_DETACH): New. libgomp/ * fortran.c (omp_fulfill_event_): New. * libgomp.h (struct gomp_task): Add detach and completion_sem fields. (struct gomp_team): Add task_detach_queue and task_detach_count fields. * libgomp.map (OMP_5.0.1): Add omp_fulfill_event and omp_fulfill_event_. * libgomp_g.h (GOMP_task): Add extra argument. * omp.h.in (enum omp_event_handle_t): New. (omp_fulfill_event): New. * omp_lib.f90.in (omp_event_handle_kind): New. (omp_fulfill_event): New. * omp_lib.h.in (omp_event_handle_kind): New. (omp_fulfill_event): Declare. * priority_queue.c (priority_tree_find): New. (priority_list_find): New. (priority_queue_find): New. * priority_queue.h (priority_queue_predicate): New. (priority_queue_find): New. * task.c (gomp_init_task): Initialize detach field. (task_fulfilled_p): New. (GOMP_task): Add detach argument. Ignore detach argument if GOMP_TASK_FLAG_DETACH not set in flags. Initialize completion_sem field. Copy address of completion_sem into detach argument and into the start of the data record. Wait for detach event if task not deferred. (gomp_barrier_handle_tasks): Queue tasks with unfulfilled events. Remove completed tasks and requeue dependent tasks. (omp_fulfill_event): New. * team.c (gomp_new_team): Initialize task_detach_queue and task_detach_count fields. (free_team): Free task_detach_queue field. * testsuite/libgomp.c-c++-common/task-detach-1.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-2.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-3.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-4.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-5.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-6.c: New testcase. * testsuite/libgomp.fortran/task-detach-1.f90: New testcase. * testsuite/libgomp.fortran/task-detach-2.f90: New testcase. * testsuite/libgomp.fortran/task-detach-3.f90: New testcase. * testsuite/libgomp.fortran/task-detach-4.f90: New testcase. * testsuite/libgomp.fortran/task-detach-5.f90: New testcase. * testsuite/libgomp.fortran/task-detach-6.f90: New testcase.
2021-01-04Update copyright years.Jakub Jelinek4-4/+4
2020-12-18openmp: Don't optimize shared to firstprivate on task with depend clauseJakub Jelinek1-0/+47
The attached testcase is miscompiled, because we optimize shared clauses to firstprivate when task body can't modify the variable even when the task has depend clause. That is wrong, because firstprivate means the variable will be copied immediately when the task is created, while with depend clause some other task might change it later before the dependencies are satisfied and the task should observe the value only after the change. 2020-12-18 Jakub Jelinek <jakub@redhat.com> * gimplify.c (struct gimplify_omp_ctx): Add has_depend member. (gimplify_scan_omp_clauses): Set it to true if OMP_CLAUSE_DEPEND appears on OMP_TASK. (gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Force GOVD_WRITTEN on shared variables if task construct has depend clause. * testsuite/libgomp.c/task-6.c: New test.
2020-12-08Fortran: Add 'omp scan' support of OpenMP 5.0Tobias Burnus1-0/+115
gcc/fortran/ChangeLog: * dump-parse-tree.c (show_omp_clauses, show_omp_node, show_code_node): Handle OMP SCAN. * gfortran.h (enum gfc_statement): Add ST_OMP_SCAN. (enum): Add OMP_LIST_SCAN_IN and OMP_LIST_SCAN_EX. (enum gfc_exec_op): Add EXEC_OMP_SCAN. * match.h (gfc_match_omp_scan): New prototype. * openmp.c (gfc_match_omp_scan): New. (gfc_match_omp_taskgroup): Cleanup. (resolve_omp_clauses, gfc_resolve_omp_do_blocks, omp_code_to_statement, gfc_resolve_omp_directive): Handle 'omp scan'. * parse.c (decode_omp_directive, next_statement, gfc_ascii_statement): Likewise. * resolve.c (gfc_resolve_code): Handle EXEC_OMP_SCAN. * st.c (gfc_free_statement): Likewise. * trans-openmp.c (gfc_trans_omp_clauses, gfc_trans_omp_do, gfc_split_omp_clauses): Handle 'omp scan'. libgomp/ChangeLog: * testsuite/libgomp.fortran/scan-1.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/reduction4.f90: Update; move FE some tests to ... * gfortran.dg/gomp/reduction6.f90: ... this new test and ... * gfortran.dg/gomp/reduction7.f90: ... this new test. * gfortran.dg/gomp/reduction5.f90: Add dg-error. * gfortran.dg/gomp/scan-1.f90: New test. * gfortran.dg/gomp/scan-2.f90: New test. * gfortran.dg/gomp/scan-3.f90: New test. * gfortran.dg/gomp/scan-4.f90: New test. * gfortran.dg/gomp/scan-5.f90: New test. * gfortran.dg/gomp/scan-6.f90: New test. * gfortran.dg/gomp/scan-7.f90: New test.
2020-11-25Fix templatized C++ OpenACC 'cache' directive ICEsThomas Schwinge2-1/+24
This has been broken forever, whoops... gcc/cp/ * pt.c (tsubst_omp_clauses): Handle 'OMP_CLAUSE__CACHE_'. (tsubst_expr): Handle 'OACC_CACHE'. gcc/testsuite/ * c-c++-common/goacc/cache-1.c: Update. * c-c++-common/goacc/cache-2.c: Likewise. * g++.dg/goacc/cache-1.C: New. * g++.dg/goacc/cache-2.C: Likewise. libgomp/ * testsuite/libgomp.oacc-c++/cache-1.C: New. * testsuite/libgomp.oacc-c-c++-common/cache-1.c: Update.
2020-11-25Fix atomic_capture-1.f90 testcaseAndrew Stubbs1-19/+19
The testcase had invalid assumptions about which loop iterations would run first and last. libgomp/ChangeLog * testsuite/libgomp.oacc-fortran/atomic_capture-1.f90 (main): Adjust expected results.
2020-11-24[testsuite] Avoid Tcl 8.5-specific behaviorThomas Schwinge2-0/+16
gcc/ * doc/install.texi (Prerequisites) <Tcl>: Add comment. gcc/testsuite/ * c-c++-common/goacc/kernels-decompose-1.c: Avoid Tcl 8.5-specific behavior. * c-c++-common/goacc/kernels-decompose-2.c: Likewise. * gfortran.dg/goacc/kernels-decompose-1.f95: Likewise. * gfortran.dg/goacc/kernels-decompose-2.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Avoid Tcl 8.5-specific behavior. * testsuite/libgomp.oacc-fortran/pr94358-1.f90: Likewise. Reported-by: David Edelsohn <dje.gcc@gmail.com>
2020-11-18testsuite/libgomp.c/usleep.h: Use sleep-loop also for GCNTobias Burnus1-3/+4
As typically configured, newlib's libc.a does not build 'posix' and, hence, usleep is not available. Thus, use the same fallback as for nvptx. libgomp/ * testsuite/libgomp.c/usleep.h (fallback_usleep): Renamed from nvptx_usleep; use also for device={arch(gcn)}.
2020-11-14openmp: Add support for non-VLA {,first}private allocate on omp taskJakub Jelinek2-86/+195
This patch adds support for custom allocators on private/firstprivate clauses for task (and taskloop) constructs. Private didn't need anything special, but firstprivate if it is passed by reference needs the GOMP_alloc calls in the copyfn and GOMP_free in the task body. 2020-11-14 Jakub Jelinek <jakub@redhat.com> * gimplify.c (gimplify_omp_for): Add OMP_CLAUSE_ALLOCATE_ALLOCATOR decls as firstprivate on task clauses even when allocate clause decl is not lastprivate. * omp-low.c (install_var_field): Don't dereference omp_is_reference types if mask is 33 rather than 1. (scan_sharing_clauses): Populate allocate_map even for task constructs. For now remove it back for variables mentioned in reduction and in_reduction clauses on task/taskloop constructs or on VLA task firstprivates. For firstprivate on task construct, install the var field into field_map with by_ref and 33 instead of false and 1 if mentioned in allocate clause. (lower_private_allocate): Set TREE_THIS_NOTRAP on the created MEM_REF. (lower_rec_input_clauses): Handle allocate for task firstprivatized non-VLA variables. (create_task_copyfn): Likewise. * testsuite/libgomp.c-c++-common/allocate-1.c (struct S): New type. (foo): Add tests for non-VLA private and firstprivate clauses on omp task. (bar): Likewise. Remove taking of address from private/firstprivate variables. * testsuite/libgomp.c++/allocate-1.C (struct S): New type. (foo): Add p, q, px and s arguments. Add tests for array reductions and for non-VLA private and firstprivate clauses on omp task. (bar): Removed. (main): Adjust foo caller. Don't call bar.
2020-11-13Decompose OpenACC 'kernels' constructs into parts, a sequence of compute ↵Gergö Barany5-3/+66
constructs Not yet enabled by default: for now, the current mode of OpenACC 'kernels' constructs handling still remains '-fopenacc-kernels=parloops', but that is to change later. gcc/ * omp-oacc-kernels-decompose.cc: New. * Makefile.in (OBJS): Add it. * passes.def: Instantiate it. * tree-pass.h (make_pass_omp_oacc_kernels_decompose): Declare. * flag-types.h (enum openacc_kernels): Add. * doc/invoke.texi (-fopenacc-kernels): Document. * gimple.h (enum gf_mask): Add 'GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED', 'GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE', 'GF_OMP_TARGET_KIND_OACC_DATA_KERNELS'. (is_gimple_omp_oacc, is_gimple_omp_offloaded): Handle these. * gimple-pretty-print.c (dump_gimple_omp_target): Likewise. * omp-expand.c (expand_omp_target, build_omp_regions_1) (omp_make_gimple_edges): Likewise. * omp-low.c (scan_sharing_clauses, scan_omp_for) (check_omp_nesting_restrictions, lower_oacc_reductions) (lower_oacc_head_mark, lower_omp_target): Likewise. * omp-offload.c (execute_oacc_device_lower): Likewise. gcc/c-family/ * c.opt (fopenacc-kernels): Add. gcc/fortran/ * lang.opt (fopenacc-kernels): Add. gcc/testsuite/ * c-c++-common/goacc/kernels-decompose-1.c: New. * c-c++-common/goacc/kernels-decompose-2.c: New. * c-c++-common/goacc/kernels-decompose-ice-1.c: New. * c-c++-common/goacc/kernels-decompose-ice-2.c: New. * gfortran.dg/goacc/kernels-decompose-1.f95: New. * gfortran.dg/goacc/kernels-decompose-2.f95: New. * c-c++-common/goacc/if-clause-2.c: Adjust. * gfortran.dg/goacc/kernels-tree.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c: New. * testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Adjust. * testsuite/libgomp.oacc-fortran/pr94358-1.f90: Likewise. Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
2020-11-13Add 'libgomp.oacc-fortran/pr94358-1.f90' [PR94358]Gergö Barany1-0/+34
Document status quo re PR94358 "[OMP] Privatize internal array variables introduced by the Fortran FE". libgomp/ PR fortran/94358 * testsuite/libgomp.oacc-fortran/pr94358-1.f90: New. Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
2020-11-13openmp: Support allocate for C/C++ array section reductionsJakub Jelinek1-6/+31
This adds allocate clause support for array section reductions. Furthermore, it fixes one bug that would cause inscan reductions with allocate to be rejected by C, and for now just ignores allocate for inscan/task reductions, that will need slightly more work. 2020-11-13 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-low.c (scan_sharing_clauses): For now remove for reduction clauses with inscan or task modifiers decl from allocate_map. (lower_private_allocate): Handle TYPE_P (new_var). (lower_rec_input_clauses): Handle allocate clause for C/C++ array reductions. gcc/c/ * c-typeck.c (c_finish_omp_clauses): Don't clear OMP_CLAUSE_REDUCTION_INSCAN unless reduction_seen == -2. libgomp/ * testsuite/libgomp.c-c++-common/allocate-1.c (foo): Add tests for array reductions. (main): Adjust foo callers.
2020-11-12openmp: Implement allocate clause in omp lowering.Jakub Jelinek2-0/+448
For now, task/taskloop constructs aren't handled and C/C++ array reductions and reductions with task or inscan modifiers need further work. Instead of calling omp_alloc/omp_free (where the former doesn't have alignment argument and omp_aligned_alloc is 5.1 only feature), this calls GOMP_alloc/GOMP_free, so that the library can fail if it would fall back into NULL (exception is zero length allocations). 2020-11-12 Jakub Jelinek <jakub@redhat.com> gcc/ * builtin-types.def (BT_FN_PTR_SIZE_SIZE_PTRMODE): New function type. * omp-builtins.def (BUILT_IN_GOACC_DECLARE): Move earlier. (BUILT_IN_GOMP_ALLOC, BUILT_IN_GOMP_FREE): New builtins. * gimplify.c (gimplify_scan_omp_clauses): Force allocator into a decl if it is not NULL, INTEGER_CST or decl. (gimplify_adjust_omp_clauses): Clear GOVD_EXPLICIT on explicit clauses which are being removed. Remove allocate clauses for variables not seen if they are private, firstprivate or linear too. Call omp_notice_variable on the allocator otherwise. (gimplify_omp_for): Handle iterator vars mentioned in allocate clauses similarly to non-is_gimple_reg iterators. * omp-low.c (struct omp_context): Add allocate_map field. (delete_omp_context): Delete it. (scan_sharing_clauses): Fill it from allocate clauses. Remove it if mentioned also in shared clause. (lower_private_allocate): New function. (lower_rec_input_clauses): Handle allocate clause for privatized variables, except for task/taskloop, C/C++ array reductions for now and task/inscan variables. (lower_send_shared_vars): Don't consider variables in allocate_map as shared. * omp-expand.c (expand_omp_for_generic, expand_omp_for_static_nochunk, expand_omp_for_static_chunk): Use expand_omp_build_assign instead of gimple_build_assign + gsi_insert_after. * builtins.c (builtin_fnspec): Handle BUILTIN_GOMP_ALLOC and BUILTIN_GOMP_FREE. * tree-ssa-ccp.c (evaluate_stmt): Handle BUILTIN_GOMP_ALLOC. * tree-ssa-dce.c (mark_stmt_if_obviously_necessary): Handle BUILTIN_GOMP_ALLOC. (mark_all_reaching_defs_necessary_1): Handle BUILTIN_GOMP_ALLOC and BUILTIN_GOMP_FREE. (propagate_necessity): Likewise. gcc/fortran/ * f95-lang.c (ATTR_ALLOC_WARN_UNUSED_RESULT_SIZE_2_NOTHROW_LIST): Define. (gfc_init_builtin_functions): Add alloc_size and warn_unused_result attributes to __builtin_GOMP_alloc. * types.def (BT_PTRMODE): New primitive type. (BT_FN_VOID_PTR_PTRMODE, BT_FN_PTR_SIZE_SIZE_PTRMODE): New function types. libgomp/ * libgomp.map (GOMP_alloc, GOMP_free): Export at GOMP_5.0.1. * omp.h.in (omp_alloc): Add malloc and alloc_size attributes. * libgomp_g.h (GOMP_alloc, GOMP_free): Declare. * allocator.c (omp_aligned_alloc): New for now static function, add alignment argument and handle it. (omp_alloc): Reimplement using omp_aligned_alloc. (GOMP_alloc, GOMP_free): New functions. (omp_free): Add ialias. * testsuite/libgomp.c-c++-common/allocate-1.c: New test. * testsuite/libgomp.c++/allocate-1.C: New test.
2020-11-12Adjust 'libgomp.oacc-fortran/attach-descriptor-1.f90' for improved location ↵Thomas Schwinge1-2/+1
information Fix-up for commit b71ff8c15f5a7d6b1cc1524b4d27843f0d88dbda "Fortran: improve location data for OpenACC/OpenMP directives [PR97782]". libgomp/ PR fortran/97782 * testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: Adjust.
2020-11-10openmp: Implement OpenMP 5.0 base-pointer attachement and clause orderingChung-Lin Tang1-0/+82
This patch implements some parts of the target variable mapping changes specified in OpenMP 5.0, including base-pointer attachment/detachment behavior for array section list-items in map clauses, and ordering of map clauses according to map kind. 2020-11-10 Chung-Lin Tang <cltang@codesourcery.com> gcc/c-family/ChangeLog: * c-common.h (c_omp_adjust_map_clauses): New declaration. * c-omp.c (struct map_clause): Helper type for c_omp_adjust_map_clauses. (c_omp_adjust_map_clauses): New function. gcc/c/ChangeLog: * c-parser.c (c_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (c_parser_omp_target_enter_data): Likewise. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_target): Likewise. * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/cp/ChangeLog: * parser.c (cp_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (cp_parser_omp_target_enter_data): Likewise. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_target): Likewise. * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix interaction between reference case and attach/detach. (finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/ChangeLog: * gimplify.c (is_or_contains_p): New static helper function. (omp_target_reorder_clauses): New function. (gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to reorder clause list according to OpenMP 5.0 rules. Add handling of GOMP_MAP_ATTACH_DETACH for OpenMP cases. * omp-low.c (is_omp_target): New static helper function. (scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. (lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. gcc/testsuite/ChangeLog: * c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid. * gfortran.dg/gomp/map-2.f90: Likewise. * c-c++-common/gomp/map-5.c: New testcase. libgomp/ChangeLog: * libgomp.h (enum gomp_map_vars_kind): Adjust enum values to be bit-flag usable. * oacc-mem.c (acc_map_data): Adjust gomp_map_vars argument flags to 'GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA'. (goacc_enter_datum): Likewise for call to gomp_map_vars_async. (goacc_enter_data_internal): Likewise. * target.c (gomp_map_vars_internal): Change checks of GOMP_MAP_VARS_ENTER_DATA to use bit-and (&). Adjust use of gomp_attach_pointer for OpenMP cases. (gomp_exit_data): Add handling of GOMP_MAP_DETACH. (GOMP_target_enter_exit_data): Add handling of GOMP_MAP_ATTACH. * testsuite/libgomp.c-c++-common/ptr-attach-1.c: New testcase.
2020-11-05openmp: Mark deprecated symbols in OpenMP 5.0Kwok Cheung Yeung26-1/+37
2020-11-05 Ulrich Drepper <drepper@redhat.com> Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * Makefile.am (%.mod): Add -cpp and -fopenmp to compile flags. * Makefile.in: Regenerate. * fortran.c: Wrap uses of omp_set_nested and omp_get_nested with pragmas to ignore -Wdeprecated-declarations warnings. * icv.c: Likewise. * omp.h.in (__GOMP_DEPRECATED_5_0): Define. Mark omp_lock_hint_* enum values, omp_lock_hint_t, omp_set_nested, and omp_get_nested with __GOMP_DEPRECATED_5_0. * omp_lib.f90.in: Mark omp_get_nested and omp_set_nested as deprecated. * testsuite/libgomp.c++/affinity-1.C: Add -Wno-deprecated-declarations to test options. * testsuite/libgomp.c/affinity-1.c: Likewise. * testsuite/libgomp.c/affinity-2.c: Likewise. * testsuite/libgomp.c/appendix-a/a.15.1.c: Likewise. * testsuite/libgomp.c/lib-1.c: Likewise. * testsuite/libgomp.c/nested-1.c: Likewise. * testsuite/libgomp.c/nested-2.c: Likewise. * testsuite/libgomp.c/nested-3.c: Likewise. * testsuite/libgomp.c/pr32362-1.c: Likewise. * testsuite/libgomp.c/pr32362-2.c: Likewise. * testsuite/libgomp.c/pr32362-3.c: Likewise. * testsuite/libgomp.c/pr35549.c: Likewise. * testsuite/libgomp.c/pr42942.c: Likewise. * testsuite/libgomp.c/pr61200.c: Likewise. * testsuite/libgomp.c/sort-1.c: Likewise. * testsuite/libgomp.c/target-5.c: Likewise. * testsuite/libgomp.c/target-6.c: Likewise. * testsuite/libgomp.c/teams-1.c: Likewise. * testsuite/libgomp.c/thread-limit-1.c: Likewise. * testsuite/libgomp.c/thread-limit-2.c: Likewise. * testsuite/libgomp.c/thread-limit-4.c: Likewise. * testsuite/libgomp.fortran/affinity1.f90: Likewise. * testsuite/libgomp.fortran/lib1.f90: Likewise. * testsuite/libgomp.fortran/lib2.f: Likewise. * testsuite/libgomp.fortran/nested1.f90: Likewise. * testsuite/libgomp.fortran/teams1.f90: Likewise.
2020-11-02Simplify and enhance 'libgomp.oacc-c-c++-common/pr85486*.c' [PR85486]Thomas Schwinge3-97/+20
Avoid code duplication, and better test what we expect to happen. libgomp/ PR target/85486 * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Simplify and enhance. * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise.
2020-11-02libgomp testsuite: tell warning from error diagnostics, etc. [PR80219, PR85303]Thomas Schwinge1-0/+6
This changes makes 'dg-warning', 'dg-error', 'dg-bogus', 'dg-message' behave as expected, and also enables use of relative line numbers as well as 'dg-line'. libgomp/ PR testsuite/80219 PR testsuite/85303 * testsuite/lib/libgomp.exp (libgomp_init): Set 'gcc_warning_prefix', 'gcc_error_prefix'.
2020-10-28openmp: Implicitly discover declare target for variants of declare variant callsJakub Jelinek1-0/+42
This marks all variants of declare variant also declare target if the base functions are called directly in target regions or declare target functions. 2020-10-28 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-offload.c (omp_declare_target_tgt_fn_r): Handle direct calls to declare variant base functions. libgomp/ * testsuite/libgomp.c/target-42.c: New test.
2020-10-28xfail and improve some failing libgomp tests [PR81690]Jakub Jelinek3-5/+31
With the patch I've posted today to fix up declare variant LTO handling, Tobias reported the patch still doesn't work, and there are two reasons for that. One is that when the base function is marked implicitly as declare target, we don't mark also implicitly the variants. I'll need to ask on omp-lang about details for that, but generally the compiler should do it some way. The other one is that the way base_delay is written, it will always call the usleep function, which is undesirable for nvptx. While the compiler will replace all direct calls to base_delay to nvptx_delay, the base_delay definition which calls usleep stays. 2020-10-28 Jakub Jelinek <jakub@redhat.com> Tom de Vries <tdevries@suse.de> PR testsuite/81690 * testsuite/libgomp.c/usleep.h: New file. * testsuite/libgomp.c/target-32.c: Include usleep.h. (main): Use tgt_usleep instead of usleep. * testsuite/libgomp.c/thread-limit-2.c: Include usleep.h. (main): Use tgt_usleep instead of usleep.
2020-10-28lto: LTO cgraph support for late declare variant resolution [PR96680]Jakub Jelinek1-0/+54
> I've tried to add the saving/restoring next to ipa refs saving/restoring, as > the declare variant alt stuff is kind of extension of those, unfortunately > following doesn't compile, because I need to also write or read a tree there > (ctx is a portion of DECL_ATTRIBUTES of the base function), but the ipa refs > write/read back functions don't have arguments that can be used for that. This patch adds the streaming out and in of those omp_declare_variant_alt hash table on the side data for the declare_variant_alt cgraph_nodes and treats for LTO purposes the declare_variant_alt nodes (which have no body) as if they contained a body that calls all the possible variants. After IPA all the calls to these magic declare_variant_alt calls are replaced with call to one of the variant depending on which one has the highest score in the context. 2020-10-28 Jakub Jelinek <jakub@redhat.com> PR lto/96680 gcc/ * lto-streamer.h (omp_lto_output_declare_variant_alt, omp_lto_input_declare_variant_alt): Declare variant. * symtab.c (symtab_node::get_partitioning_class): Return SYMBOL_DUPLICATE for declare_variant_alt nodes. * passes.c (ipa_write_summaries): Add declare_variant_alt to partition. * lto-cgraph.c (output_refs): Call omp_lto_output_declare_variant_alt on declare_variant_alt nodes. (input_refs): Call omp_lto_input_declare_variant_alt on declare_variant_alt nodes. * lto-streamer-out.c (output_function): Don't call collect_block_tree_leafs if DECL_INITIAL is error_mark_node. (lto_output): Call output_function even for declare_variant_alt nodes. * omp-general.c (omp_lto_output_declare_variant_alt, omp_lto_input_declare_variant_alt): New functions. gcc/lto/ * lto-common.c (lto_fixup_prevailing_decls): Don't use LTO_NO_PREVAIL on TREE_LIST's TREE_PURPOSE. * lto-partition.c (lto_balanced_map): Treat declare_variant_alt nodes like definitions. libgomp/ * testsuite/libgomp.c/declare-variant-1.c: New test.
2020-10-22openmp: Add test for OMP_TARGET_OFFLOAD=mandatory for cases where it must ↵Jakub Jelinek1-0/+33
not fail 2020-10-22 Jakub Jelinek <jakub@redhat.com> * testsuite/libgomp.c/target-41.c: New test.
2020-10-22openmp: Change omp_get_initial_device () to match OpenMP 5.1 requirementsJakub Jelinek1-0/+10
> Therefore, I think until omp_get_initial_device () value is changed, we The following so far untested patch implements that change. OpenMP 4.5 said for omp_get_initial_device: The value of the device number is implementation defined. If it is between 0 and one less than omp_get_num_devices() then it is valid for use with all device constructs and routines; if it is outside that range, then it is only valid for use with the device memory routines and not in the device clause. and OpenMP 5.0 similarly, but OpenMP 5.1 says: The value of the device number is the value returned by the omp_get_num_devices routine. As the new value is compatible with what has been required earlier, I think we can change it already now. 2020-10-22 Jakub Jelinek <jakub@redhat.com> * icv.c (omp_get_initial_device): Remove including corresponding ialias. * icv-device.c (omp_get_initial_device): New function. Return gomp_get_num_devices (). Add ialias. * target.c (resolve_device): Don't fail with OMP_TARGET_OFFLOAD=mandatory if device_id is equal to gomp_get_num_devices (). (omp_target_alloc, omp_target_free, omp_target_is_present, omp_target_memcpy, omp_target_memcpy_rect, omp_target_associate_ptr, omp_target_disassociate_ptr, omp_pause_resource): Use gomp_get_num_devices () instead of GOMP_DEVICE_HOST_FALLBACK on the first use in the functions, in uses dominated by the gomp_get_num_devices call use num_devices_openmp instead. * libgomp.texi (omp_get_initial_device): Document. * config/gcn/icv-device.c (omp_get_initial_device): New function. Add ialias. * config/nvptx/icv-device.c (omp_get_initial_device): Likewise. * testsuite/libgomp.c/target-40.c: New test.
2020-10-13openmp: Add support for the omp_get_supported_active_levels runtime library ↵Kwok Cheung Yeung2-0/+4
routine This patch implements the omp_get_supported_active_levels runtime routine from the OpenMP 5.0 specification, which returns the maximum number of active nested parallel regions supported by this implementation. The current maximum (set using the omp_set_max_active_levels routine or the OMP_MAX_ACTIVE_LEVELS environment variable) cannot exceed this number. 2020-10-13 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * env.c (gomp_max_active_levels_var): Initialize to gomp_supported_active_levels. (initialize_env): Limit gomp_max_active_levels_var to be at most equal to gomp_supported_active_levels. * fortran.c (omp_get_supported_active_levels): Add ialias_redirect. (omp_get_supported_active_levels_): New. * icv.c (omp_set_max_active_levels): Limit gomp_max_active_levels_var to at most equal to gomp_supported_active_levels. (omp_get_supported_active_levels): New. * libgomp.h (gomp_supported_active_levels): New. * libgomp.map (OMP_5.0.1): Add omp_get_supported_active_levels and omp_get_supported_active_levels_. * libgomp.texi (omp_get_supported_active_levels): New. (omp_set_max_active_levels): Update. Add reference to omp_get_supported_active_levels. * omp.h.in (omp_get_supported_active_levels): New. * omp_lib.f90.in (omp_get_supported_active_levels): New. * omp_lib.h.in (omp_get_supported_active_levels): New. * testsuite/libgomp.c/lib-2.c (main): Check omp_get_max_active_levels against omp_get_supported_active_levels. * testsuite/libgomp.fortran/lib4.f90 (lib4): Likewise.
2020-10-06[openacc, libgomp, testsuite] Xfail declare-5.f90Tom de Vries1-0/+1
We're currently running into: ... FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 execution test FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -O1 execution test FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 execution test FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -O3 -fomit-frame-pointer \ -funroll-loops -fpeel-loops -ftracer -finline-functions execution test FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -O3 -g execution test FAIL: libgomp.oacc-fortran/declare-5.f90 -DACC_DEVICE_TYPE_nvidia=1 \ -DACC_MEM_SHARED=0 -foffload=nvptx-none -Os execution test ... A PR was filed for this: PR92790 - "[OpenACC] declare device_resident - Fortran common blocks not handled / libgomp.oacc-fortran/declare-5.f90 fails" Xfail the fails. Tested on x86_64-linux with nvptx accelerator. libgomp/ChangeLog: 2020-10-06 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-fortran/declare-5.f90: Add xfail for PR92790.
2020-10-06[openacc] Fix acc declare for VLAsTom de Vries1-5/+0
Consider test-case test.c, with VLA A: ... int main (void) { int N = 1000; int A[N]; #pragma acc declare copy(A) return 0; } ... compiled using: ... $ gcc test.c -fopenacc -S -fdump-tree-all ... At original, we have: ... #pragma acc declare map(tofrom:A); ... but at gimple, we have a map (to:A.1), but not a map (from:A.1): ... int[0:D.2074] * A.1; { int A[0:D.2074] [value-expr: *A.1]; saved_stack.2 = __builtin_stack_save (); try { A.1 = __builtin_alloca_with_align (D.2078, 32); #pragma omp target oacc_declare map(to:(*A.1) [len: D.2076]) } finally { __builtin_stack_restore (saved_stack.2); } } ... This is caused by the following incompatibility. When storing the desired from clause in oacc_declare_returns, we use 'A.1' as the key: ... 10898 oacc_declare_returns->put (decl, c); (gdb) call debug_generic_expr (decl) A.1 (gdb) call debug_generic_expr (c) map(from:(*A.1)) ... but when looking it up, we use 'A' as the key: ... (gdb) 1471 tree *c = oacc_declare_returns->get (t); (gdb) call debug_generic_expr (t) A ... Fix this by extracing the 'A.1' lookup key from 'A' using the decl-expr. In addition, unshare the looked up value, to fix avoid running into an "incorrect sharing of tree nodes" error. Using these two fixes, we get our desired: ... finally { + #pragma omp target oacc_declare map(from:(*A.1)) __builtin_stack_restore (saved_stack.2); } ... Build on x86_64-linux with nvptx accelerator, tested libgomp. gcc/ChangeLog: 2020-10-06 Tom de Vries <tdevries@suse.de> PR middle-end/90861 * gimplify.c (gimplify_bind_expr): Handle lookup in oacc_declare_returns using key with decl-expr. libgomp/ChangeLog: 2020-10-06 Tom de Vries <tdevries@suse.de> PR middle-end/90861 * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Remove xfail.
2020-10-05[omp, ftracer] Don't duplicate blocks in SIMT regionTom de Vries1-0/+11
When running the libgomp testsuite on x86_64-linux with nvptx accelerator on the test-case included in this patch, we run into: ... FAIL: libgomp.fortran/pr95654.f90 -O3 -fomit-frame-pointer -funroll-loops \ -fpeel-loops -ftracer -finline-functions execution test ... The test-case is a minimal version of this FAIL: ... FAIL: libgomp.fortran/pr66199-5.f90 -O3 -fomit-frame-pointer -funroll-loops \ -fpeel-loops -ftracer -finline-functions execution test ... but that one has stopped failing at commit c2ebf4f10de "openmp: Add support for non-rect simd and improve collapsed simd support". The problem is that ftracer duplicates a block containing GOMP_SIMT_VOTE_ANY. That is, before ftracer we have (dropping the GOMP_SIMT_ prefix): ... bb4(ENTER_ALLOC) *----------+ | \ | \ | v | * v bb8 *<------------* bb5(VOTE_ANY) *-------------+ | | | | | | | | | v | * v bb7(XCHG_IDX) *<------------* bb6(EXIT) ... The XCHG_IDX internal-fn does inter-SIMT-lane communication, which for nvptx maps onto shfl, an operator which has the requirement that the warp executing the operator is convergent. The warp diverges at bb4, and reconverges at bb5, and does not diverge by going to bb7, so the shfl is indeed executed by a convergent warp. After ftracer, we have: ... bb4(ENTER_ALLOC) *----------+ | \ | \ | \ | \ v v * * bb5(VOTE_ANY) bb8(VOTE_ANY) * * |\ /| | \ +--------+ | | \/ | | /\ | | / +----------v |/ * v bb7(XCHG_IDX) *<--------------* bb6(EXIT) ... The warp diverges again at bb5, but does not reconverge again before bb6, so the shfl is executed by a divergent warp, which causes the FAIL. Fix this by making ftracer ignore blocks containing ENTER_ALLOC, VOTE_ANY and EXIT, effectively treating the SIMT region conservatively. An argument can be made that the test needs to be added in a more generic place, like gimple_can_duplicate_bb_p or some such, and that ftracer then needs to use the generic test. But that's a discussion with a much broader scope, so I'm leaving that for another patch. Bootstrapped and reg-tested on x86_64-linux. Build on x86_64-linux with nvptx accelerator, tested with libgomp. gcc/ChangeLog: PR fortran/95654 * tracer.c (ignore_bb_p): Ignore GOMP_SIMT_ENTER_ALLOC, GOMP_SIMT_VOTE_ANY and GOMP_SIMT_EXIT. libgomp/ChangeLog: 2020-10-05 Tom de Vries <tdevries@suse.de> PR fortran/95654 * testsuite/libgomp.fortran/pr95654.f90: New test.
2020-10-02libgomp: Regenerate configure files with automake 1.15.1Tobias Burnus1-4/+4
libgomp/ChangeLog: * Makefile.in: Regenerate with automake 1.15.1. * aclocal.m4: Likewise. * configure: Likewise. * testsuite/Makefile.in: Likewise.
2020-09-30OpenMP: Add implicit declare target for nested proceduresTobias Burnus1-0/+45
gcc/ChangeLog: * omp-offload.c (omp_discover_implicit_declare_target): Also handled nested functions. libgomp/ChangeLog: * testsuite/libgomp.fortran/declare-target-3.f90: New test.
2020-09-29libgomp: disable barriers in nested teamsAndrew Stubbs1-0/+31
Both GCN and NVPTX allow nested parallel regions, but the barrier implementation did not allow the nested teams to run independently of each other (due to hardware limitations). This patch fixes that, under the assumption that each thread will create a new subteam of one thread, by simply not using barriers when there's no other thread to synchronise. libgomp/ChangeLog: * config/gcn/bar.c (gomp_barrier_wait_end): Skip the barrier if the total number of threads is one. (gomp_team_barrier_wake): Likewise. (gomp_team_barrier_wait_end): Likewise. (gomp_team_barrier_wait_cancel_end): Likewise. * config/nvptx/bar.c (gomp_barrier_wait_end): Likewise. (gomp_team_barrier_wake): Likewise. (gomp_team_barrier_wait_end): Likewise. (gomp_team_barrier_wait_cancel_end): Likewise. * testsuite/libgomp.c-c++-common/nested-parallel-unbalanced.c: New test.
2020-09-28OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390)Tobias Burnus2-0/+75
gcc/ChangeLog: PR middle-end/96390 * omp-offload.c (omp_discover_declare_target_tgt_fn_r): Handle alias nodes. libgomp/ChangeLog: PR middle-end/96390 * testsuite/libgomp.c++/pr96390.C: New test. * testsuite/libgomp.c-c++-common/pr96390.c: New test.
2020-09-25openmp: Add support for non-rect simd and improve collapsed simd supportJakub Jelinek1-0/+296
The following change adds support for non-rectangular simd loops. While working on that, I've noticed we actually don't vectorize collapsed simd loops at all, because the code that I thought would be vectorizable actually is not vectorized. While in theory for the constant lower/upper bounds and constant step of all but the outermost loop we could in theory vectorize by computing the seprate iterators using vectorized division and modulo for each of them from the single iterator that increments by 1 from 0 to total iteration count in the loop nest, I think that would be fairly expensive and the chances of the loop body being vectorizable would be low e.g. because of array indices unlikely to be linear and would need scatters/gathers. This patch changes the generated code to vectorize only the innermost loop which has higher chance of being vectorized. Below is the list of tests and function names in which the patch resulted in vectorizing something that hasn't been vectorized before (ok, the first line is a new test). I've also found that the vectorizer will not vectorize loops with non-constant steps, I plan to do something about those incrementally on the omp-expand.c side (basically, compute number of iterations before the loop and use a 0 to number_of_iterations step 1 IV as the main one). I have problem with the composite simd vectorization though. The point is that each thread (or task etc.) is given only a range of consecutive iterations, so somewhere earlier it computes total number of iterations and splits the work between the workers and then the intent is to try to vectorize it. So, each thread is then given a begin ... end-1 range that it would handle. This means that from the single begin value I need to compute the individual iteration vars I should start at and then goto into the loop nest to begin iterating there (and actually compute how many iterations the innermost loop should do each time so that it stops before end). Very roughly the IL I emit is something like: int t[100][100][100]; void foo (int a, int b, int c, int d, int e, int f, int g, int h, int u, int v, int w, int x) { int i, j, k; int cnt; if (x) { i = u; j = v; k = w; goto doit; } for (i = a; i < b; i += c) for (j = d; j < e; j += f) { k = g; doit: for (; k < h; k++) t[i][j][k] += i + j + k; } } Unfortunately, some pass then turns the innermost loop to have more than 2 basic blocks and it isn't vectorized because of that. Also, I have disabled (for now) SIMTization of collapsed simd loops, because for SIMT it would be using a single thread anyway and I didn't want to bother with checking SIMT on all places I've been changing. If SIMT support is added for some or all collapsed loops, that omp-low.c change needs to be reverted. Here is that list of what hasn't been vectorized before and is now: gcc/testsuite/gcc.dg/vect/vect-simd-17.c doit gcc/testsuite/gfortran.dg/gomp/openmp-simd-6.f90 bar libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-10.c f28_taskloop_simd_normal._omp_fn.0 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-10.c _Z24f28_taskloop_simd_normalv._omp_fn.0 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-11.c f25_t_simd_normal._omp_fn.0 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-11.c f26_t_simd_normal._omp_fn.0 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-11.c f27_t_simd_normal._omp_fn.0 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-11.c f28_tpf_simd_guided32._omp_fn.1 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-11.c f28_tpf_simd_runtime._omp_fn.1 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-11.c _Z17f25_t_simd_normaliiiiiii._omp_fn.0 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-11.c _Z17f26_t_simd_normaliiiixxi._omp_fn.0 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-11.c _Z17f27_t_simd_normalv._omp_fn.0 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-11.c _Z20f28_tpf_simd_runtimev._omp_fn.1 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-11.c _Z21f28_tpf_simd_guided32v._omp_fn.1 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-2.c f7_simd_normal libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-2.c f7_simd_normal libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-2.c f8_f_simd_guided32 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-2.c f8_f_simd_guided32 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-2.c f8_f_simd_runtime libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-2.c f8_f_simd_runtime libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-2.c f8_pf_simd_guided32._omp_fn.0 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-2.c f8_pf_simd_runtime._omp_fn.0 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-2.c _Z18f8_pf_simd_runtimev._omp_fn.0 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-2.c _Z19f8_pf_simd_guided32v._omp_fn.0 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-4.c f8_taskloop_simd_normal._omp_fn.0 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-4.c _Z23f8_taskloop_simd_normalv._omp_fn.0 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-5.c f7_t_simd_normal._omp_fn.0 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-5.c f8_tpf_simd_guided32._omp_fn.1 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-5.c f8_tpf_simd_runtime._omp_fn.1 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-5.c _Z16f7_t_simd_normalv._omp_fn.0 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-5.c _Z19f8_tpf_simd_runtimev._omp_fn.1 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-5.c _Z20f8_tpf_simd_guided32v._omp_fn.1 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-8.c f25_simd_normal libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-8.c f25_simd_normal libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-8.c f26_simd_normal libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-8.c f26_simd_normal libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-8.c f27_simd_normal libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-8.c f27_simd_normal libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-8.c f28_f_simd_guided32 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-8.c f28_f_simd_guided32 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-8.c f28_f_simd_runtime libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-8.c f28_f_simd_runtime libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-8.c f28_pf_simd_guided32._omp_fn.0 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/for-8.c f28_pf_simd_runtime._omp_fn.0 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-8.c _Z19f28_pf_simd_runtimev._omp_fn.0 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/for-8.c _Z20f28_pf_simd_guided32v._omp_fn.0 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/master-combined-1.c main._omp_fn.9 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/master-combined-1.c main._omp_fn.9 libgomp/testsuite/libgomp.c++/../libgomp.c-c++-common/simd-1.c f2 libgomp/testsuite/libgomp.c/../libgomp.c-c++-common/simd-1.c f2 libgomp/testsuite/libgomp.c/pr70680-2.c f1._omp_fn.0 libgomp/testsuite/libgomp.c/pr70680-2.c f2._omp_fn.0 libgomp/testsuite/libgomp.c/pr70680-2.c f3._omp_fn.0 libgomp/testsuite/libgomp.c/pr70680-2.c f4._omp_fn.0 libgomp/testsuite/libgomp.c/simd-8.c foo libgomp/testsuite/libgomp.c/simd-9.c bar libgomp/testsuite/libgomp.c/simd-9.c foo 2020-09-25 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-low.c (scan_omp_1_stmt): Don't call scan_omp_simd for collapse > 1 loops as simt doesn't support collapsed loops yet. * omp-expand.c (expand_omp_for_init_counts, expand_omp_for_init_vars): Small tweaks to function comment. (expand_omp_simd): Rewritten collapse > 1 support to only attempt to vectorize the innermost loop and emit set of outer loops around it. For non-composite simd with collapse > 1 without broken loop don't even try to compute number of iterations first. Add support for non-rectangular simd loops. (expand_omp_for): Don't sorry_at on non-rectangular simd loops. gcc/testsuite/ * gcc.dg/vect/vect-simd-17.c: New test. libgomp/ * testsuite/libgomp.c/loop-25.c: New test.
2020-09-22libgomp.fortran/pr66199-5.f90: Make stop codes uniqueTobias Burnus1-1/+1
libgomp/ChangeLog: PR fortran/95654 * testsuite/libgomp.fortran/pr66199-5.f90: Make stop codes unique.
2020-09-16c++: local-scope OMP UDR reductions have no template headNathan Sidwell1-0/+1
This corrects the earlier problems with removing the template header from local omp reductions. And it uncovered a latent bug. When we tsubst such a decl, we immediately tsubst its body. cp_check_omp_declare_reduction gets a success return value to gate that instantiation. udr-2.C got a further error, as the omp checking machinery doesn't appear to turn the reduction into an error mark when failing. I didn't dig into that further. udr-3.C appears to have been invalid and accidentally worked. gcc/cp/ * cp-tree.h (cp_check_omp_declare_reduction): Return bool. * semantics.c (cp_check_omp_declare_reduction): Return true on for success. * pt.c (push_template_decl_real): OMP reductions do not get a template header. (tsubst_function_decl): Remove special casing for local decl omp reductions. (tsubst_expr): Call instantiate_body for a local omp reduction. (instantiate_body): Add nested_p parm, and deal with such instantiations. (instantiate_decl): Reject FUNCTION_SCOPE entities, adjust instantiate_body call. gcc/testsuite/ * g++.dg/gomp/udr-2.C: Add additional expected error. libgomp/ * testsuite/libgomp.c++/udr-3.C: Add missing ctor.
2020-09-15OpenMP/Fortran: Fix (re)mapping of allocatable/pointer arrays [PR96668]Tobias Burnus2-0/+200
gcc/cp/ChangeLog: PR fortran/96668 * cp-gimplify.c (cxx_omp_finish_clause): Add bool openacc arg. * cp-tree.h (cxx_omp_finish_clause): Likewise * semantics.c (handle_omp_for_class_iterator): Update call. gcc/fortran/ChangeLog: PR fortran/96668 * trans.h (gfc_omp_finish_clause): Add bool openacc arg. * trans-openmp.c (gfc_omp_finish_clause): Ditto. Use GOMP_MAP_ALWAYS_POINTER with PSET for pointers. (gfc_trans_omp_clauses): Like the latter and also if the always modifier is used. gcc/ChangeLog: PR fortran/96668 * gimplify.c (gimplify_omp_for): Add 'bool openacc' argument; update omp_finish_clause calls. (gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses, gimplify_expr, gimplify_omp_loop): Update omp_finish_clause and/or gimplify_for calls. * langhooks-def.h (lhd_omp_finish_clause): Add bool openacc arg. * langhooks.c (lhd_omp_finish_clause): Likewise. * langhooks.h (lhd_omp_finish_clause): Likewise. * omp-low.c (scan_sharing_clauses): Keep GOMP_MAP_TO_PSET cause for 'declare target' vars. include/ChangeLog: PR fortran/96668 * gomp-constants.h (GOMP_MAP_ALWAYS_POINTER_P): Define. libgomp/ChangeLog: PR fortran/96668 * libgomp.h (struct target_var_desc): Add has_null_ptr_assoc member. * target.c (gomp_map_vars_existing): Add always_to_flag flag. (gomp_map_vars_existing): Update call to it. (gomp_map_fields_existing): Likewise (gomp_map_vars_internal): Update PSET handling such that if a nullptr is now allocated or if GOMP_MAP_POINTER is used PSET is updated and pointer remapped. (GOMP_target_enter_exit_data): Hanlde GOMP_MAP_ALWAYS_POINTER like GOMP_MAP_POINTER. * testsuite/libgomp.fortran/map-alloc-ptr-1.f90: New test. * testsuite/libgomp.fortran/map-alloc-ptr-2.f90: New test.
2020-09-14[libgomp, nvptx] Add __sync_compare_and_swap_16Tom de Vries1-0/+1
As reported here ( https://gcc.gnu.org/pipermail/gcc-patches/2020-September/553070.html ), when running test-case libgomp.c-c++-common/reduction-16.c for powerpc host with nvptx accelerator, we run into: ... unresolved symbol __sync_val_compare_and_swap_16 ... I can reproduce the problem on x86_64 with a trigger patch that: - initializes ix86_isa_flags2 to TARGET_ISA2_CX16 - enables define_expand "atomic_load<mode>" in gcc/config/i386/sync.md for TImode The problem is that omp-expand.c generates atomic builtin calls based on checks whether those are supported on the host, which forces the target to support these, even though those checks fail for the accelerator target. Fix this by: - adding a __sync_val_compare_and_swap_16 in libgomp for nvptx, which falls back onto libatomic's __atomic_compare_and_swap_16 - adding -foffload=-latomic in the test-case Tested libgomp on x86_64-linux with nvptx accelerator. Tested libgomp with trigger patch on x86_64-linux with nvptx accelerator. libgomp/ChangeLog: * config/nvptx/atomic.c: New file. Add __sync_val_compare_and_swap_16. * testsuite/libgomp.c-c++-common/reduction-16.c: Add -latomic for target offload_target_nvptx.
2020-09-08openacc: Fix atomic_capture-2.c iteration-ordering issuesJulian Brown1-49/+43
The test case was written with assumptions about loop iteration ordering that are not guaranteed by OpenACC and do not apply on all targets, in particular AMD GCN. This patch removes those assumptions. 2020-09-08 Julian Brown <julian@codesourcery.com> libgomp/ * testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c: Remove iteration-ordering assumptions.
2020-09-08openacc: Fix race condition in Fortran loop collapse testsJulian Brown2-0/+6
The gangs participating in a gang-partitioned loop are not all guaranteed to complete before some given gang continues to execute beyond that loop. This means that two existing test cases contain a race condition, because a loop that may be gang-partitioned is followed immediately by another loop. The fix is to place the loops in separate parallel regions. 2020-09-08 Julian Brown <julian@codesourcery.com> libgomp/ * testsuite/libgomp.oacc-fortran/collapse-1.f90: Fix race condition. * testsuite/libgomp.oacc-fortran/collapse-2.f90: Likewise.
2020-08-20Fortran: Fix OpenMP's 'if(simd:' etc. conditionsTobias Burnus1-0/+63
gcc/fortran/ChangeLog: * openmp.c (gfc_match_omp_clauses): Re-order 'if' clause pasing to avoid creating spurious symbols. libgomp/ChangeLog: * testsuite/libgomp.fortran/lastprivate-conditional-10.f90: New test.
2020-08-13nvptx: Add support for subword compare-and-swapKwok Cheung Yeung1-0/+53
This adds support for __sync_val_compare_and_swap and __sync_bool_compare_and_swap for 1-byte and 2-byte long values, which are not natively supported on nvptx. Build and reg-tested on nvptx. Build and reg-tested libgomp on x86_64 with nvptx accelerator. 2020-07-16 Kwok Cheung Yeung <kcy@codesourcery.com> libgcc/ * config/nvptx/atomic.c: New. * config/nvptx/t-nvptx (LIB2ADD): Add atomic.c. gcc/testsuite/ * gcc.target/nvptx/ia64-sync-5.c: New. libgomp/ * testsuite/libgomp.c-c++-common/reduction-16.c: New.
2020-08-13openmp: Add support for non-rectangular loops in taskloop constructJakub Jelinek3-0/+894
2020-08-13 Jakub Jelinek <jakub@redhat.com> * gimplify.c (gimplify_omp_taskloop_expr): New function. (gimplify_omp_for): Use it. For OMP_FOR_NON_RECTANGULAR loops adjust in outer taskloop the var-outer decls. * omp-expand.c (expand_omp_taskloop_for_inner): Handle non-rectangular loops. (expand_omp_for): Don't reject non-rectangular taskloop. * omp-general.c (omp_extract_for_data): Don't assert that non-rectangular loops have static schedule, instead treat loop->m1 or loop->m2 as if loop->n1 or loop->n2 is non-constant. * testsuite/libgomp.c/loop-22.c (main): Add some further tests. * testsuite/libgomp.c/loop-23.c (main): Likewise. * testsuite/libgomp.c/loop-24.c: New test.
2020-08-08openmp: Handle clauses with gimple sequences in convert_nonlocal_omp_clauses ↵Jakub Jelinek1-0/+21
properly If the walk_body on the various sequences of reduction, lastprivate and/or linear clauses needs to create a temporary variable, we should declare that variable in that sequence rather than outside, where it would need to be privatized inside of the construct. 2020-08-08 Jakub Jelinek <jakub@redhat.com> PR fortran/93553 * tree-nested.c (convert_nonlocal_omp_clauses): For OMP_CLAUSE_REDUCTION, OMP_CLAUSE_LASTPRIVATE and OMP_CLAUSE_LINEAR save info->new_local_var_chain around walks of the clause gimple sequences and declare_vars if needed into the sequence. 2020-08-08 Tobias Burnus <tobias@codesourcery.com> PR fortran/93553 * testsuite/libgomp.fortran/pr93553.f90: New test.
2020-08-05openmp: Handle even some combined non-rectangular loopsJakub Jelinek2-0/+378
The number of loops computation and logical iteration -> actual iterator values computations can now be done separately even on composite constructs (though for triangular loops it would still be more efficient to propagate a few values through, will handle that incrementally). simd and taskloop are still unhandled. 2020-08-05 Jakub Jelinek <jakub@redhat.com> * omp-expand.c (expand_omp_for): Don't disallow combined non-rectangular loops. * testsuite/libgomp.c/loop-22.c: New test. * testsuite/libgomp.c/loop-23.c: New test.
2020-08-05openmp: Handle reduction clauses on host teams construct [PR96459]Jakub Jelinek3-27/+82
As the new testcase shows, we weren't actually performing reductions on host teams construct. And fixing that revealed a flaw in the for-14.c testcase. The problem is that the tests perform also initialization and checking around the calls to the functions with the OpenMP constructs. In that testcase, all the tests have been spawned from a teams construct but only the tested loops were distribute, which means the initialization and checking has been performed redundantly and racily in each team. Fixed by performing the initialization and checking outside of host teams and only do the calls to functions with the tested constructs inside of host teams. 2020-08-05 Jakub Jelinek <jakub@redhat.com> PR middle-end/96459 * omp-low.c (lower_omp_taskreg): Call lower_reduction_clauses even in for host teams. * testsuite/libgomp.c/teams-3.c: New test. * testsuite/libgomp.c-c++-common/for-2.h (OMPTEAMS): Define to nothing if not defined yet. (N(test)): Use it before all N(f*) calls. * testsuite/libgomp.c-c++-common/for-14.c (DO_PRAGMA, OMPTEAMS): Define. (main): Don't call all test_* functions from within #pragma omp teams reduction(|:err), call them directly.
2020-08-04[nvptx] Handle V2DI/V2SI mode in nvptx_gen_shuffleTom de Vries2-0/+57
With the pr96628-part1.f90 source and -ftree-slp-vectorize, we run into an ICE due to the fact that V2DI mode is not handled in nvptx_gen_shuffle. Fix this by adding handling of V2DI as well as V2SI mode in nvptx_gen_shuffle. Build and reg-tested on x86_64 with nvptx accelerator. gcc/ChangeLog: PR target/96428 * config/nvptx/nvptx.c (nvptx_gen_shuffle): Handle V2SI/V2DI. libgomp/ChangeLog: PR target/96428 * testsuite/libgomp.oacc-fortran/pr96628-part1.f90: New test. * testsuite/libgomp.oacc-fortran/pr96628-part2.f90: New test.
2020-08-03openacc: No attach/detach present/release mappings for array descriptorsJulian Brown1-10/+81
Standalone attach and detach clauses should not create present/release mappings for Fortran array descriptors (e.g. used when we have a pointer to an array), both because it is unnecessary and because those mappings will be incorrectly subject to reference counting. Simply omitting the mappings means we just use GOMP_MAP_TO_PSET and GOMP_MAP_{ATTACH,DETACH} mappings for array descriptors. That requires a tweak in gimplify.c, since we may now see GOMP_MAP_TO_PSET without a preceding data-movement mapping. 2020-08-03 Julian Brown <julian@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Don't create present/release mappings for array descriptors. gcc/ * gimplify.c (gimplify_omp_target_update): Allow GOMP_MAP_TO_PSET without a preceding data-movement mapping. gcc/testsuite/ * gfortran.dg/goacc/attach-descriptor.f90: Update pattern output. Add scanning of gimplify dump. libgomp/ * testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: Don't run for shared-memory devices. Extend with further checking. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
generated by cgit v1.1 at 2025-09-25 05:31:26 +0000