d7bd009ab0
79 Commits
Author | SHA1 | Message | Date | |
---|---|---|---|---|
Kwok Cheung Yeung
|
a6d22fb21c |
openmp: Add support for the OpenMP 5.0 task detach clause
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. |
||
Jakub Jelinek
|
99dee82307 | Update copyright years. | ||
Kwok Cheung Yeung
|
6fae7eda96 |
openmp: Retire nest-var ICV for OpenMP 5.1
This removes the nest-var ICV, expressing nesting in terms of the max-active-levels-var ICV instead. The max-active-levels-var ICV is now per data environment rather than per device. 2020-11-18 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * env.c (gomp_global_icv): Remove nest_var field. Add max_active_levels_var field. (gomp_max_active_levels_var): Remove. (parse_boolean): Return true on success. (handle_omp_display_env): Express OMP_NESTED in terms of max_active_levels_var. Change format specifier for max_active_levels_var. (initialize_env): Set max_active_levels_var from OMP_MAX_ACTIVE_LEVELS, OMP_NESTED, OMP_NUM_THREADS and OMP_PROC_BIND. * icv.c (omp_set_nested): Express in terms of max_active_levels_var. (omp_get_nested): Likewise. (omp_set_max_active_levels): Use max_active_levels_var field instead of gomp_max_active_levels_var. (omp_get_max_active_levels): Likewise. * libgomp.h (struct gomp_task_icv): Remove nest_var field. Add max_active_levels_var field. (gomp_supported_active_levels): Set to UCHAR_MAX. (gomp_max_active_levels_var): Delete. * libgomp.texi (omp_get_nested): Update documentation. (omp_set_nested): Likewise. (OMP_MAX_ACTIVE_LEVELS): Likewise. (OMP_NESTED): Likewise. (OMP_NUM_THREADS): Likewise. (OMP_PROC_BIND): Likewise. * parallel.c (gomp_resolve_num_threads): Replace reference to nest_var with max_active_levels_var. Use max_active_levels_var field instead of gomp_max_active_levels_var. |
||
Chung-Lin Tang
|
9e62802422 |
openmp: Implement OpenMP 5.0 base-pointer attachement and clause ordering
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. |
||
Kwok Cheung Yeung
|
1bfc07d150 |
openmp: Implement support for OMP_TARGET_OFFLOAD environment variable
This implements support for the OMP_TARGET_OFFLOAD environment variable introduced in the OpenMP 5.0 standard, which controls how offloading is handled. It may be set to MANDATORY (abort if offloading cannot be performed), DISABLED (no offloading to devices) or DEFAULT (offload to device if possible, fall back to host if not). 2020-10-20 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * env.c (gomp_target_offload_var): New. (parse_target_offload): New. (handle_omp_display_env): Print value of OMP_TARGET_OFFLOAD. (initialize_env): Parse OMP_TARGET_OFFLOAD. * libgomp.h (gomp_target_offload_t): New. (gomp_target_offload_var): New. * libgomp.texi (OMP_TARGET_OFFLOAD): New section. * target.c (resolve_device): Generate error if device not found and offloading is mandatory. (gomp_target_fallback): Generate error if offloading is mandatory. (GOMP_target): Add argument in call to gomp_target_fallback. (GOMP_target_ext): Likewise. (gomp_target_data_fallback): Generate error if offloading is mandatory. (GOMP_target_data): Add argument in call to gomp_target_data_fallback. (GOMP_target_data_ext): Likewise. (gomp_target_task_fn): Add argument in call to gomp_target_fallback. (gomp_target_init): Return early if offloading is disabled. |
||
Kwok Cheung Yeung
|
8949b985db |
openmp: Add support for the omp_get_supported_active_levels runtime library 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. |
||
Tobias Burnus
|
972da55746 |
OpenMP/Fortran: Fix (re)mapping of allocatable/pointer arrays [PR96668]
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. |
||
Julian Brown
|
bc4ed079dc |
openacc: Deep copy attach/detach should not affect reference counts
Attach and detach operations are not supposed to affect structural or dynamic reference counts for OpenACC. Previously they did so, which led to subtle problems in some circumstances. We can avoid reference-counting attach/detach operations by extending and slightly repurposing the do_detach field in target_var_desc. It is now called is_attach to better reflect its new role. 2020-07-27 Julian Brown <julian@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> libgomp/ * libgomp.h (struct target_var_desc): Rename do_detach field to is_attach. * oacc-mem.c (goacc_exit_datum_1): Add assert. Don't set finalize for GOMP_MAP_FORCE_DETACH. Update checking to use is_attach field. (goacc_enter_data_internal): Don't affect reference counts for attach mappings. (goacc_exit_data_internal): Don't affect reference counts for detach mappings. * target.c (gomp_map_vars_existing): Don't affect reference counts for attach mappings. (gomp_map_vars_internal): Set renamed is_attach flag unconditionally to mark attach mappings. (gomp_unmap_vars_internal): Use is_attach flag to prevent affecting reference count for attach mappings. * testsuite/libgomp.oacc-c-c++-common/mdc-refcount-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/mdc-refcount-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/mdc-refcount-2.c: New test. * testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: Mark test as shouldfail. * testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Adjust to fail gracefully in no-finalize mode. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com> |
||
Julian Brown
|
6f5b4b64d2 |
openacc: Adjust dynamic reference count semantics
This patch adjusts how dynamic reference counts work so that they match the semantics of the source program more closely, instead of representing "excess" reference counts beyond those that represent pointers in the internal libgomp splay-tree data structure. This allows some corner cases to be handled more gracefully. 2020-07-10 Julian Brown <julian@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> libgomp/ * libgomp.h (struct splay_tree_key_s): Change virtual_refcount to dynamic_refcount. (struct gomp_device_descr): Remove GOMP_MAP_VARS_OPENACC_ENTER_DATA. * oacc-mem.c (acc_map_data): Substitute virtual_refcount for dynamic_refcount. (acc_unmap_data): Update comment. (goacc_map_var_existing, goacc_enter_datum): Adjust for dynamic_refcount semantics. (goacc_exit_datum_1, goacc_exit_datum): Re-add some error checking. Adjust for dynamic_refcount semantics. (goacc_enter_data_internal): Implement "present" case of dynamic memory-map handling here. Update "non-present" case for dynamic_refcount semantics. (goacc_exit_data_internal): Use goacc_exit_datum_1. * target.c (gomp_map_vars_internal): Remove GOMP_MAP_VARS_OPENACC_ENTER_DATA handling. Update for dynamic_refcount handling. (gomp_unmap_vars_internal): Remove virtual_refcount handling. (gomp_load_image_to_device): Substitute dynamic_refcount for virtual_refcount. * testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Remove XFAILs. * testsuite/libgomp.oacc-c-c++-common/refcounting-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/refcounting-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/struct-3-1-1.c: New test. * testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Remove XFAILs and trace output. * testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: Remove trace output. * testsuite/libgomp.oacc-fortran/dynamic-incr-structural-1.f90: New test. * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c: Remove stale comment. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Remove XFAILs. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Adjust XFAIL. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com> |
||
Jakub Jelinek
|
800bcc8c00 |
openmp: Add basic library allocator support.
This patch adds very basic allocator support (omp_{init,destroy}_allocator, omp_{alloc,free}, omp_[sg]et_default_allocator). The plan is to use memkind (likely dlopened) for high bandwidth memory, but that part isn't implemented yet, probably mlock for pinned memory and see what other options there are for other kinds of memory. For offloading targets, we need to decide if we want to support the dynamic allocators (and on which targets), or if e.g. all we do is at compile time replace omp_alloc/omp_free calls with constexpr predefined allocators with something special. And allocate directive and allocator/uses_allocators clauses are future work too. 2020-05-19 Jakub Jelinek <jakub@redhat.com> * omp.h.in (omp_uintptr_t): New typedef. (__GOMP_UINTPTR_T_ENUM): Define. (omp_memspace_handle_t, omp_allocator_handle_t, omp_alloctrait_key_t, omp_alloctrait_value_t, omp_alloctrait_t): New typedefs. (__GOMP_DEFAULT_NULL_ALLOCATOR): Define. (omp_init_allocator, omp_destroy_allocator, omp_set_default_allocator, omp_get_default_allocator, omp_alloc, omp_free): Declare. * libgomp.h (struct gomp_team_state): Add def_allocator field. (gomp_def_allocator): Declare. * libgomp.map (OMP_5.0.1): Export omp_set_default_allocator, omp_get_default_allocator, omp_init_allocator, omp_destroy_allocator, omp_alloc and omp_free. * team.c (gomp_team_start): Copy over ts.def_allocator. * env.c (gomp_def_allocator): New variable. (parse_wait_policy): Adjust function comment. (parse_allocator): New function. (handle_omp_display_env): Print OMP_ALLOCATOR. (initialize_env): Call parse_allocator. * Makefile.am (libgomp_la_SOURCES): Add allocator.c. * allocator.c: New file. * icv.c (omp_set_default_allocator, omp_get_default_allocator): New functions. * testsuite/libgomp.c-c++-common/alloc-1.c: New test. * testsuite/libgomp.c-c++-common/alloc-2.c: New test. * testsuite/libgomp.c-c++-common/alloc-3.c: New test. * Makefile.in: Regenerated. |
||
Thomas Schwinge
|
6fc0385c0c |
OpenACC 'acc_get_property' cleanup
include/ * gomp-constants.h (enum gomp_device_property): Remove. libgomp/ * libgomp-plugin.h (enum goacc_property): New. Adjust all users to use this instead of 'enum gomp_device_property'. (GOMP_OFFLOAD_get_property): Rename to... (GOMP_OFFLOAD_openacc_get_property): ... this. Adjust all users. * libgomp.h (struct gomp_device_descr): Move 'GOMP_OFFLOAD_openacc_get_property'... (struct acc_dispatch_t): ... here. Adjust all users. * plugin/plugin-hsa.c (GOMP_OFFLOAD_get_property): Remove. liboffloadmic/ * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_property): Remove. From-SVN: r280150 |
||
Jakub Jelinek
|
91df4397a1 |
re PR libgomp/93219 (unused return value in affinity-fmt.c)
PR libgomp/93219 * libgomp.h (gomp_print_string): Change return type from void to int. * affinity-fmt.c (gomp_print_string): Likewise. Return true if not all characters have been written. From-SVN: r280137 |
||
Jakub Jelinek
|
8d9254fc8a |
Update copyright years.
From-SVN: r279813 |
||
Maciej W. Rozycki
|
6c84c8bf9b |
Add OpenACC 2.6 `acc_get_property' support
Add generic support for the OpenACC 2.6 `acc_get_property' and `acc_get_property_string' routines, as well as full handlers for the host and the NVPTX offload targets and minimal handlers for the HSA, Intel MIC, and AMD GCN offload targets. Included are C/C++ and Fortran tests that, in particular, print the property values for acc_property_vendor, acc_property_memory, acc_property_free_memory, acc_property_name, and acc_property_driver. The output looks as follows: Vendor: GNU Name: GOMP Total memory: 0 Free memory: 0 Driver: 1.0 with the host driver (where the memory related properties are not supported for the host device and yield 0, conforming to the standard) and output like: Vendor: Nvidia Total memory: 12651462656 Free memory: 12202737664 Name: TITAN V Driver: CUDA Driver 9.1 with the NVPTX driver. 2019-12-22 Maciej W. Rozycki <macro@codesourcery.com> Frederik Harwath <frederik@codesourcery.com> Thomas Schwinge <tschwinge@codesourcery.com> include/ * gomp-constants.h (gomp_device_property): New enum. libgomp/ * libgomp.h (gomp_device_descr): Add `get_property_func' member. * libgomp-plugin.h (gomp_device_property_value): New union. (gomp_device_property_value): New prototype. * openacc.h (acc_device_t): Add `acc_device_current' enumeration constant. (acc_device_property_t): New enum. (acc_get_property, acc_get_property_string): New prototypes. * oacc-init.c (acc_get_device_type): Also assert that result is not `acc_device_current'. (get_property_any, acc_get_property, acc_get_property_string): New functions. * openacc.f90 (openacc_kinds): Add `acc_device_current' and `acc_property_memory', `acc_property_free_memory', `acc_property_name', `acc_property_vendor' and `acc_property_driver' constants. Add `acc_device_property' data type. (openacc_internal): Add `acc_get_property' and `acc_get_property_string' interfaces. Add `acc_get_property_h', `acc_get_property_string_h', `acc_get_property_l' and `acc_get_property_string_l'. * oacc-host.c (host_get_property): New function. (host_dispatch): Wire it. * target.c (gomp_load_plugin_for_device): Handle `get_property'. * libgomp.map (OACC_2.6): Add `acc_get_property', `acc_get_property_h_', `acc_get_property_string' and `acc_get_property_string_h_' symbols. * libgomp.texi (OpenACC Runtime Library Routines): Add `acc_get_property'. (acc_get_property): New node. * plugin/plugin-gcn.c (GOMP_OFFLOAD_get_property): New function (stub). * plugin/plugin-hsa.c (GOMP_OFFLOAD_get_property): New function. * plugin/plugin-nvptx.c (CUDA_CALLS): Add `cuDeviceGetName', `cuDeviceTotalMem', `cuDriverGetVersion' and `cuMemGetInfo' calls. (GOMP_OFFLOAD_get_property): New function. (struct ptx_device): Add new field "name". (cuda_driver_version_s): Add new static variable ... (nvptx_init): ... and init from here. * testsuite/libgomp.oacc-c-c++-common/acc_get_property.c: New test. * testsuite/libgomp.oacc-c-c++-common/acc_get_property-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/acc_get_property-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c: New file with test helper functions. * testsuite/libgomp.oacc-fortran/acc_get_property.f90: New test. liboffloadmic/ * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_property): New function. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com> Co-Authored-By: Frederik Harwath <frederik@codesourcery.com> Co-Authored-By: Thomas Schwinge <tschwinge@codesourcery.com> From-SVN: r279710 |
||
Julian Brown
|
8e7e71ff24 |
OpenACC 2.6 deep copy: libgomp parts
include/ * gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_4, GOMP_MAP_DEEP_COPY): Define. (gomp_map_kind): Add GOMP_MAP_ATTACH, GOMP_MAP_DETACH, GOMP_MAP_FORCE_DETACH. libgomp/ * libgomp.h (struct target_var_desc): Add do_detach flag. * oacc-init.c (acc_shutdown_1): Free aux block if present. * oacc-mem.c (find_group_last): Add SIZES parameter. Support struct components. Tidy up and add some new checks. (goacc_enter_data_internal): Update call to find_group_last. (goacc_exit_data_internal): Support detach operations and GOMP_MAP_STRUCT. (GOACC_enter_exit_data): Handle initial GOMP_MAP_STRUCT or GOMP_MAP_FORCE_PRESENT in finalization detection code. Handle attach/detach in enter/exit data detection code. * target.c (gomp_map_vars_existing): Initialise do_detach field of tgt_var_desc. (gomp_map_vars_internal): Support attach. (gomp_unmap_vars_internal): Support detach. From-SVN: r279625 |
||
Julian Brown
|
5d5be7bfb5 |
OpenACC 2.6 deep copy: attach/detach API routines
libgomp/ * libgomp.h (struct splay_tree_aux): Add attach_count field. (gomp_attach_pointer, gomp_detach_pointer): Add prototypes. * libgomp.map (OACC_2.6): New section. Add acc_attach, acc_attach_async, acc_detach, acc_detach_async, acc_detach_finalize, acc_detach_finalize_async. * oacc-mem.c (acc_attach_async, acc_attach, goacc_detach_internal, acc_detach, acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): New functions. * openacc.h (acc_attach, acc_attach_async, acc_detach, (acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): Add prototypes. * target.c (gomp_attach_pointer, gomp_detach_pointer): New functions. (gomp_remove_var_internal): Free attachment counts if present. * testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c: New test. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com> From-SVN: r279624 |
||
Julian Brown
|
5bcd470bf0 |
Use gomp_map_val for OpenACC host-to-device address translation
libgomp/ * libgomp.h (gomp_map_val): Add prototype. * oacc-parallel.c (GOACC_parallel_keyed): Use gomp_map_val instead of open-coding device-address calculation. * target.c (gomp_map_val): Make global. Use OFFSET_POINTER in non-present case. Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com> From-SVN: r279622 |
||
Julian Brown
|
378da98fcc |
OpenACC reference count overhaul
libgomp/ * libgomp.h (struct splay_tree_key_s): Substitute dynamic_refcount field for virtual_refcount. (enum gomp_map_vars_kind): Add GOMP_MAP_VARS_OPENACC_ENTER_DATA. (gomp_free_memmap): Remove prototype. * oacc-init.c (acc_shutdown_1): Iteratively call gomp_remove_var instead of calling gomp_free_memmap. * oacc-mem.c (acc_map_data): Use virtual_refcount instead of dynamic_refcount. (acc_unmap_data): Open code instead of forcing target_mem_desc's to_free field to NULL then calling gomp_unmap_vars. Handle REFCOUNT_INFINITY on target blocks. (goacc_enter_data): Rename to... (goacc_enter_datum): ...this. Remove MAPNUM parameter and special handling for mapping groups. Use virtual_refcount instead of dynamic_refcount. Use GOMP_MAP_VARS_OPENACC_ENTER_DATA for map_map_vars_async call. Re-do lookup for target pointer return value. (acc_create, acc_create_async, acc_copyin, acc_copyin_async): Call renamed goacc_enter_datum function. (goacc_exit_data): Rename to... (goacc_exit_datum): ...this. Update for virtual_refcount semantics. (acc_delete, acc_delete_async, acc_delete_finalize, acc_delete_finalize_async, acc_copyout, acc_copyout_async, acc_copyout_finalize, acc_copyout_finalize_async): Call renamed goacc_exit_datum function. (gomp_acc_remove_pointer, find_pointer): Remove functions. (find_group_last, goacc_enter_data_internal, goacc_exit_data_internal): New functions. (GOACC_enter_exit_data): Use goacc_enter_data_internal and goacc_exit_data_internal helper functions. * target.c (gomp_map_vars_internal): Handle GOMP_MAP_VARS_OPENACC_ENTER_DATA. Update for virtual_refcount semantics. (gomp_unmap_vars_internal): Update for virtual_refcount semantics. (gomp_load_image_to_device, omp_target_associate_ptr): Zero-initialise virtual_refcount field instead of dynamic_refcount. (gomp_free_memmap): Remove function. * testsuite/libgomp.oacc-c-c++-common/unmap-infinity-1.c: New test. * testsuite/libgomp.c-c++-common/unmap-infinity-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Add XFAIL. From-SVN: r279621 |
||
Julian Brown
|
2a656a9359 |
Use aux struct in libgomp for infrequently-used/API-specific data
libgomp/ * libgomp.h (struct splay_tree_aux): New. (struct splay_tree_key_s): Replace link_key field with aux pointer. * target.c (gomp_map_vars_internal): Adjust for link_key being moved to aux struct. (gomp_remove_var_internal): Free aux block if present. (gomp_load_image_to_device): Zero-initialise aux field instead of link_key field. (omp_target_associate_pointer): Zero-initialise aux field. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com> From-SVN: r279620 |
||
Thomas Schwinge
|
6278b54922 |
Make 'libgomp/target.c:gomp_unmap_tgt' 'static' again
This got changed to 'attribute_hidden' in r271128, but it's not actually used outside of 'libgomp/target.c'. libgomp/ * target.c (gomp_unmap_tgt): Make it 'static'. * libgomp.h (gomp_unmap_tgt): Remove. From-SVN: r279529 |
||
Julian Brown
|
1cbd94e834 |
Fix potential race condition in OpenACC "exit data" operations
PR libgomp/92881 libgomp/ * libgomp.h (gomp_remove_var_async): Add prototype. * oacc-mem.c (delete_copyout): Call gomp_remove_var_async instead of gomp_remove_var. * target.c (gomp_unref_tgt): Change return type to bool, indicating whether target_mem_desc was unmapped. (gomp_unref_tgt_void): New. (gomp_remove_var): Reimplement in terms of... (gomp_remove_var_internal): ...this new helper function. (gomp_remove_var_async): New, implemented using above helper function. (gomp_unmap_vars_internal): Use gomp_unref_tgt_void instead of gomp_unref_tgt. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com> From-SVN: r279388 |
||
Thomas Schwinge
|
57963e3934 |
[OpenACC] Consolidate 'GOACC_enter_exit_data' and its helper functions in 'libgomp/oacc-mem.c'
libgomp/ * oacc-parallel.c (find_pointer, GOACC_enter_exit_data): Move... * oacc-mem.c: ... here. (gomp_acc_insert_pointer, gomp_acc_remove_pointer): Rename to 'goacc_insert_pointer', 'goacc_remove_pointer', and make 'static'. * libgomp.h (gomp_acc_insert_pointer, gomp_acc_remove_pointer): Remove. * libgomp_g.h: Update. From-SVN: r279233 |
||
Thomas Schwinge
|
47afc7b4dd |
[PR92116, PR92877] [OpenACC] Replace 'openacc.data_environ' by standard libgomp mechanics
libgomp/ PR libgomp/92116 PR libgomp/92877 * oacc-mem.c (lookup_dev): Reimplement. Adjust all users. * libgomp.h (struct acc_dispatch_t): Remove 'data_environ' member. Adjust all users. * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c: Remove XFAIL. * testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr92877-1.c: New file. Co-Authored-By: Julian Brown <julian@codesourcery.com> From-SVN: r279147 |
||
Andrew Stubbs
|
cee1645106 |
Optimize GCN OpenMP malloc performance
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 |
||
Andrew Stubbs
|
fa4999953d |
GCN libgomp port
2019-11-13 Andrew Stubbs <ams@codesourcery.com> Kwok Cheung Yeung <kcy@codesourcery.com> Julian Brown <julian@codesourcery.com> Tom de Vries <tom@codesourcery.com> include/ * gomp-constants.h (GOMP_DEVICE_GCN): Define. (GOMP_VERSION_GCN): Define. libgomp/ * Makefile.am (libgomp_la_SOURCES): Add oacc-target.c. * Makefile.in: Regenerate. * config.h.in (PLUGIN_GCN): Add new undef. * config/accel/openacc.f90 (acc_device_gcn): New parameter. * config/gcn/affinity-fmt.c: New file. * config/gcn/bar.c: New file. * config/gcn/bar.h: New file. * config/gcn/doacross.h: New file. * config/gcn/icv-device.c: New file. * config/gcn/oacc-target.c: New file. * config/gcn/simple-bar.h: New file. * config/gcn/target.c: New file. * config/gcn/task.c: New file. * config/gcn/team.c: New file. * config/gcn/time.c: New file. * configure.ac: Add amdgcn*-*-*. * configure: Regenerate. * configure.tgt: Add amdgcn*-*-*. * libgomp-plugin.h (offload_target_type): Add OFFLOAD_TARGET_TYPE_GCN. * libgomp.h (gcn_thrs): Add amdgcn variant. (set_gcn_thrs): Likewise. (gomp_thread): Likewise. * oacc-int.h (goacc_thread): Likewise. * oacc-target.c: New file. * openacc.f90 (acc_device_gcn): New parameter. * openacc.h (acc_device_t): Add acc_device_gcn. * team.c (gomp_free_pool_helper): Add amdgcn support. Co-Authored-By: Julian Brown <julian@codesourcery.com> Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com> Co-Authored-By: Tom de Vries <tom@codesourcery.com> From-SVN: r278135 |
||
Julian Brown
|
6c7e076b74 |
Libgomp magic offset value self-documentation
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 |
||
Jakub Jelinek
|
810f316dd6 |
configure.ac: Remove GCC_HEADER_STDINT(gstdint.h).
* configure.ac: Remove GCC_HEADER_STDINT(gstdint.h). * libgomp.h: Include <stdint.h> instead of "gstdint.h". * oacc-parallel.c: Don't include "libgomp_g.h". * plugin/plugin-hsa.c: Include <stdint.h> instead of "gstdint.h". * plugin/plugin-nvptx.c: Don't include "gstdint.h". * aclocal.m4: Regenerated. * config.h.in: Regenerated. * configure: Regenerated. * Makefile.in: Regenerated. From-SVN: r276389 |
||
Chung-Lin Tang
|
1f4c5b9bb2 |
2019-05-13 Chung-Lin Tang <cltang@codesourcery.com>
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com> libgomp/ * libgomp-plugin.h (struct goacc_asyncqueue): Declare. (struct goacc_asyncqueue_list): Likewise. (goacc_aq): Likewise. (goacc_aq_list): Likewise. (GOMP_OFFLOAD_openacc_register_async_cleanup): Remove. (GOMP_OFFLOAD_openacc_async_test): Remove. (GOMP_OFFLOAD_openacc_async_test_all): Remove. (GOMP_OFFLOAD_openacc_async_wait): Remove. (GOMP_OFFLOAD_openacc_async_wait_async): Remove. (GOMP_OFFLOAD_openacc_async_wait_all): Remove. (GOMP_OFFLOAD_openacc_async_wait_all_async): Remove. (GOMP_OFFLOAD_openacc_async_set_async): Remove. (GOMP_OFFLOAD_openacc_exec): Adjust declaration. (GOMP_OFFLOAD_openacc_cuda_get_stream): Likewise. (GOMP_OFFLOAD_openacc_cuda_set_stream): Likewise. (GOMP_OFFLOAD_openacc_async_exec): Declare. (GOMP_OFFLOAD_openacc_async_construct): Declare. (GOMP_OFFLOAD_openacc_async_destruct): Declare. (GOMP_OFFLOAD_openacc_async_test): Declare. (GOMP_OFFLOAD_openacc_async_synchronize): Declare. (GOMP_OFFLOAD_openacc_async_serialize): Declare. (GOMP_OFFLOAD_openacc_async_queue_callback): Declare. (GOMP_OFFLOAD_openacc_async_host2dev): Declare. (GOMP_OFFLOAD_openacc_async_dev2host): Declare. * libgomp.h (struct acc_dispatch_t): Define 'async' sub-struct. (gomp_acc_insert_pointer): Adjust declaration. (gomp_copy_host2dev): New declaration. (gomp_copy_dev2host): Likewise. (gomp_map_vars_async): Likewise. (gomp_unmap_tgt): Likewise. (gomp_unmap_vars_async): Likewise. (gomp_fini_device): Likewise. * oacc-async.c (get_goacc_thread): New function. (get_goacc_thread_device): New function. (lookup_goacc_asyncqueue): New function. (get_goacc_asyncqueue): New function. (acc_async_test): Adjust code to use new async design. (acc_async_test_all): Likewise. (acc_wait): Likewise. (acc_wait_async): Likewise. (acc_wait_all): Likewise. (acc_wait_all_async): Likewise. (goacc_async_free): New function. (goacc_init_asyncqueues): Likewise. (goacc_fini_asyncqueues): Likewise. * oacc-cuda.c (acc_get_cuda_stream): Adjust code to use new async design. (acc_set_cuda_stream): Likewise. * oacc-host.c (host_openacc_exec): Adjust parameters, remove 'async'. (host_openacc_register_async_cleanup): Remove. (host_openacc_async_exec): New function. (host_openacc_async_test): Adjust parameters. (host_openacc_async_test_all): Remove. (host_openacc_async_wait): Remove. (host_openacc_async_wait_async): Remove. (host_openacc_async_wait_all): Remove. (host_openacc_async_wait_all_async): Remove. (host_openacc_async_set_async): Remove. (host_openacc_async_synchronize): New function. (host_openacc_async_serialize): New function. (host_openacc_async_host2dev): New function. (host_openacc_async_dev2host): New function. (host_openacc_async_queue_callback): New function. (host_openacc_async_construct): New function. (host_openacc_async_destruct): New function. (struct gomp_device_descr host_dispatch): Remove initialization of old interface, add intialization of new async sub-struct. * oacc-init.c (acc_shutdown_1): Adjust to use gomp_fini_device. (goacc_attach_host_thread_to_device): Remove old async code usage. * oacc-int.h (goacc_init_asyncqueues): New declaration. (goacc_fini_asyncqueues): Likewise. (goacc_async_copyout_unmap_vars): Likewise. (goacc_async_free): Likewise. (get_goacc_asyncqueue): Likewise. (lookup_goacc_asyncqueue): Likewise. * oacc-mem.c (memcpy_tofrom_device): Adjust code to use new async design. (present_create_copy): Adjust code to use new async design. (delete_copyout): Likewise. (update_dev_host): Likewise. (gomp_acc_insert_pointer): Add async parameter, adjust code to use new async design. (gomp_acc_remove_pointer): Adjust code to use new async design. * oacc-parallel.c (GOACC_parallel_keyed): Adjust code to use new async design. (GOACC_enter_exit_data): Likewise. (goacc_wait): Likewise. (GOACC_update): Likewise. * oacc-plugin.c (GOMP_PLUGIN_async_unmap_vars): Change to assert fail when called, warn as obsolete in comment. * target.c (goacc_device_copy_async): New function. (gomp_copy_host2dev): Remove 'static', add goacc_asyncqueue parameter, add goacc_device_copy_async case. (gomp_copy_dev2host): Likewise. (gomp_map_vars_existing): Add goacc_asyncqueue parameter, adjust code. (gomp_map_pointer): Likewise. (gomp_map_fields_existing): Likewise. (gomp_map_vars_internal): New always_inline function, renamed from gomp_map_vars. (gomp_map_vars): Implement by calling gomp_map_vars_internal. (gomp_map_vars_async): Implement by calling gomp_map_vars_internal, passing goacc_asyncqueue argument. (gomp_unmap_tgt): Remove static, add attribute_hidden. (gomp_unref_tgt): New function. (gomp_unmap_vars_internal): New always_inline function, renamed from gomp_unmap_vars. (gomp_unmap_vars): Implement by calling gomp_unmap_vars_internal. (gomp_unmap_vars_async): Implement by calling gomp_unmap_vars_internal, passing goacc_asyncqueue argument. (gomp_fini_device): New function. (gomp_exit_data): Adjust gomp_copy_dev2host call. (gomp_load_plugin_for_device): Remove old interface, adjust to load new async interface. (gomp_target_fini): Adjust code to call gomp_fini_device. * plugin/plugin-nvptx.c (struct cuda_map): Remove. (struct ptx_stream): Remove. (struct nvptx_thread): Remove current_stream field. (cuda_map_create): Remove. (cuda_map_destroy): Remove. (map_init): Remove. (map_fini): Remove. (map_pop): Remove. (map_push): Remove. (struct goacc_asyncqueue): Define. (struct nvptx_callback): Define. (struct ptx_free_block): Define. (struct ptx_device): Remove null_stream, active_streams, async_streams, stream_lock, and next fields. (enum ptx_event_type): Remove. (struct ptx_event): Remove. (ptx_event_lock): Remove. (ptx_events): Remove. (init_streams_for_device): Remove. (fini_streams_for_device): Remove. (select_stream_for_async): Remove. (nvptx_init): Remove ptx_events and ptx_event_lock references. (nvptx_attach_host_thread_to_device): Remove CUDA_ERROR_NOT_PERMITTED case. (nvptx_open_device): Add free_blocks initialization, remove init_streams_for_device call. (nvptx_close_device): Remove fini_streams_for_device call, add free_blocks destruct code. (event_gc): Remove. (event_add): Remove. (nvptx_exec): Adjust parameters and code. (nvptx_free): Likewise. (nvptx_host2dev): Remove. (nvptx_dev2host): Remove. (nvptx_set_async): Remove. (nvptx_async_test): Remove. (nvptx_async_test_all): Remove. (nvptx_wait): Remove. (nvptx_wait_async): Remove. (nvptx_wait_all): Remove. (nvptx_wait_all_async): Remove. (nvptx_get_cuda_stream): Remove. (nvptx_set_cuda_stream): Remove. (GOMP_OFFLOAD_alloc): Adjust code. (GOMP_OFFLOAD_free): Likewise. (GOMP_OFFLOAD_openacc_register_async_cleanup): Remove. (GOMP_OFFLOAD_openacc_exec): Adjust parameters and code. (GOMP_OFFLOAD_openacc_async_test_all): Remove. (GOMP_OFFLOAD_openacc_async_wait): Remove. (GOMP_OFFLOAD_openacc_async_wait_async): Remove. (GOMP_OFFLOAD_openacc_async_wait_all): Remove. (GOMP_OFFLOAD_openacc_async_wait_all_async): Remove. (GOMP_OFFLOAD_openacc_async_set_async): Remove. (cuda_free_argmem): New function. (GOMP_OFFLOAD_openacc_async_exec): New plugin hook function. (GOMP_OFFLOAD_openacc_create_thread_data): Adjust code. (GOMP_OFFLOAD_openacc_cuda_get_stream): Adjust code. (GOMP_OFFLOAD_openacc_cuda_set_stream): Adjust code. (GOMP_OFFLOAD_openacc_async_construct): New plugin hook function. (GOMP_OFFLOAD_openacc_async_destruct): New plugin hook function. (GOMP_OFFLOAD_openacc_async_test): Remove and re-implement. (GOMP_OFFLOAD_openacc_async_synchronize): New plugin hook function. (GOMP_OFFLOAD_openacc_async_serialize): New plugin hook function. (GOMP_OFFLOAD_openacc_async_queue_callback): New plugin hook function. (cuda_callback_wrapper): New function. (cuda_memcpy_sanity_check): New function. (GOMP_OFFLOAD_host2dev): Remove and re-implement. (GOMP_OFFLOAD_dev2host): Remove and re-implement. (GOMP_OFFLOAD_openacc_async_host2dev): New plugin hook function. (GOMP_OFFLOAD_openacc_async_dev2host): New plugin hook function. From-SVN: r271128 |
||
Jakub Jelinek
|
a554497024 |
Update copyright years.
From-SVN: r267494 |
||
Tom de Vries
|
fe0827eed0 |
[libgomp, nvptx] Fix libgomp.c/target-5.c compilation
Libgomp test-case libgomp.c/target-5.c is failing to compile when building for x86_64 with nvptx accelerator due to missing: - getpid - gethostname - isatty (pulled in by fwrite) in the nvptx newlib. This patch fixes the build failure by: - adding a function gomp_print_string which limits the use of fwrite to a single location (in affinity-fmt.c), and - creating an nvptx version of affinity-fmt.c, which: - overrides the configure test results HAVE_GETPID and HAVE_GETHOSTNAME, and - implements fwrite using write. Build and reg-tested on x86_64 with nvptx accelerator. 2018-12-13 Tom de Vries <tdevries@suse.de> * affinity-fmt.c (gomp_print_string): New function, factored out of ... (omp_display_affinity, gomp_display_affinity_thread): ... here, and ... * fortran.c (omp_display_affinity_): ... here. * libgomp.h (gomp_print_string): Declare. * config/nvptx/affinity-fmt.c: New file. Include affinity-fmt.c, undefining HAVE_GETPID and HAVE_GETHOSTNAME, and mapping fwrite to write. From-SVN: r267100 |
||
Martin Sebor
|
79a2c4281c |
PR middle-end/81824 - Warn for missing attributes with function aliases
gcc/c-family/ChangeLog: PR middle-end/81824 * c-attribs.c (handle_copy_attribute): New function. gcc/cp/ChangeLog: PR middle-end/81824 * pt.c (warn_spec_missing_attributes): Move code to attribs.c. Call decls_mismatched_attributes. gcc/ChangeLog: PR middle-end/81824 * attribs.c (has_attribute): New helper function. (decls_mismatched_attributes, maybe_diag_alias_attributes): Same. * attribs.h (decls_mismatched_attributes): Declare. * cgraphunit.c (handle_alias_pairs): Call maybe_diag_alias_attributes. (maybe_diag_incompatible_alias): Use OPT_Wattribute_alias_. * common.opt (-Wattribute-alias): Take an argument. (-Wno-attribute-alias): New option. * doc/extend.texi (Common Function Attributes): Document copy. (Common Variable Attributes): Same. * doc/invoke.texi (-Wmissing-attributes): Document enhancement. (-Wattribute-alias): Document new option argument. gcc/testsuite/ChangeLog: PR middle-end/81824 * gcc.dg/Wattribute-alias.c: New test. * gcc.dg/Wmissing-attributes.c: New test. * gcc.dg/attr-copy.c: New test. * gcc.dg/attr-copy-2.c: New test. * gcc.dg/attr-copy-3.c: New test. * gcc.dg/attr-copy-4.c: New test. From-SVN: r265980 |
||
Jakub Jelinek
|
28567c40e2 |
builtin-types.def (BT_FN_VOID_BOOL, [...]): New.
* builtin-types.def (BT_FN_VOID_BOOL, BT_FN_VOID_SIZE_SIZE_PTR, BT_FN_UINT_UINT_PTR_PTR, BT_FN_UINT_OMPFN_PTR_UINT_UINT, BT_FN_BOOL_UINT_LONGPTR_LONG_LONG_LONGPTR_LONGPTR_PTR_PTR, BT_FN_BOOL_UINT_ULLPTR_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR, BT_FN_BOOL_LONG_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR_PTR_PTR, BT_FN_BOOL_BOOL_ULL_ULL_ULL_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR): New. * gengtype.c (open_base_files): Add omp-general.h. * gimple.c (gimple_build_omp_critical): (gimple_build_omp_taskgroup): Add CLAUSES argument. Call gimple_omp_taskgroup_set_clauses. (gimple_build_omp_atomic_load): Add mo argument, call gimple_omp_atomic_set_memory_order. (gimple_build_omp_atomic_store): Likewise. (gimple_copy): Adjust handling of GIMPLE_OMP_TASKGROUP. * gimple.def (GIMPLE_OMP_TASKGROUP): Use GSS_OMP_SINGLE_LAYOUT instead of GSS_OMP. (GIMPLE_OMP_TEAMS): Use GSS_OMP_PARALLEL_LAYOUT instead of GSS_OMP_SINGLE_LAYOUT, adjust comments. * gimple.h (enum gf_mask): Add GF_OMP_TEAMS_HOST, GF_OMP_TASK_TASKWAIT and GF_OMP_ATOMIC_MEMORY_ORDER. Remove GF_OMP_ATOMIC_SEQ_CST, use different value for GF_OMP_ATOMIC_NEED_VALUE. (struct gimple_statement_omp_taskreg): Add GIMPLE_OMP_TEAMS to comments. (struct gimple_statement_omp_single_layout): And remove here. (struct gomp_teams): Inherit from gimple_statement_omp_taskreg rather than gimple_statement_omp_single_layout. (is_a_helper <gimple_statement_omp_taskreg *>::test): Allow GIMPLE_OMP_TEAMS. (is_a_helper <const gimple_statement_omp_taskreg *>::test): Likewise. (gimple_omp_subcode): Formatting fix. (gimple_omp_teams_child_fn, gimple_omp_teams_child_fn_ptr, gimple_omp_teams_set_child_fn, gimple_omp_teams_data_arg, gimple_omp_teams_data_arg_ptr, gimple_omp_teams_set_data_arg, gimple_omp_teams_host, gimple_omp_teams_set_host, gimple_omp_task_taskwait_p, gimple_omp_task_set_taskwait_p, gimple_omp_taskgroup_clauses, gimple_omp_taskgroup_clauses_ptr, gimple_omp_taskgroup_set_clauses): New inline functions. (gimple_build_omp_atomic_load): Add enum omp_memory_order argument. (gimple_build_omp_atomic_store): Likewise. (gimple_omp_atomic_seq_cst_p): Remove. (gimple_omp_atomic_memory_order): New function. (gimple_omp_atomic_set_seq_cst): Remove. (gimple_omp_atomic_set_memory_order): New function. (gimple_build_omp_taskgroup): Add clauses argument. * gimple-pretty-print.c (dump_gimple_omp_taskgroup): New function. (dump_gimple_omp_task): Print taskwait with depend clauses. (dump_gimple_omp_atomic_load, dump_gimple_omp_atomic_store): Use dump_omp_atomic_memory_order. (pp_gimple_stmt_1): Handle GIMPLE_OMP_TASKGROUP. * gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_ALLOC_ONLY, GOVD_MAP_FROM_ONLY and GOVD_NONTEMPORAL. (enum omp_region_type): Reserve bits 1 and 2 for auxiliary flags, renumber values of most of ORT_* enumerators, add ORT_HOST_TEAMS, ORT_COMBINED_HOST_TEAMS, ORT_TASKGROUP, ORT_TASKLOOP and ORT_UNTIED_TASKLOOP enumerators. (enum gimplify_defaultmap_kind): New. (struct gimplify_omp_ctx): Remove target_map_scalars_firstprivate and target_map_pointers_as_0len_arrays members, add defaultmap. (new_omp_context): Initialize defaultmap member. (gimple_add_tmp_var): Handle ORT_TASKGROUP like ORT_WORKSHARE. (maybe_fold_stmt): Don't fold even in host teams regions. (omp_firstprivatize_variable): Handle ORT_TASKGROUP like ORT_WORKSHARE. Test ctx->defaultmap[GDMK_SCALAR] instead of ctx->omp_firstprivatize_variable. (omp_add_variable): Don't add private/firstprivate for VLAs in ORT_TASKGROUP. (omp_default_clause): Print "taskloop" rather than "task" if ORT_*TASKLOOP. (omp_notice_variable): Handle ORT_TASKGROUP like ORT_WORKSHARE. Handle new defaultmap clause kinds. (omp_is_private): Handle ORT_TASKGROUP like ORT_WORKSHARE. Allow simd iterator to be lastprivate or private. Fix up diagnostics if linear is used on collapse>1 simd iterator. (omp_check_private): Handle ORT_TASKGROUP like ORT_WORKSHARE. (gimplify_omp_depend): New function. (gimplify_scan_omp_clauses): Add shared clause on parallel for combined parallel master taskloop{, simd} if taskloop has firstprivate, lastprivate or reduction clause. Handle OMP_CLAUSE_REDUCTION_TASK diagnostics. Adjust tests for ORT_COMBINED_TEAMS. Gimplify depend clauses with iterators. Handle cancel and simd OMP_CLAUSE_IF_MODIFIERs. Handle OMP_CLAUSE_NONTEMPORAL. Handle new defaultmap clause kinds. Handle OMP_CLAUSE_{TASK,IN}_REDUCTION. Diagnose invalid conditional lastprivate. (gimplify_adjust_omp_clauses_1): Ignore GOVD_NONTEMPORAL. Handle GOVD_MAP_ALLOC_ONLY and GOVD_MAP_FROM_ONLY. (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_NONTEMPORAL. Handle OMP_CLAUSE_{TASK,IN}_REDUCTION. (gimplify_omp_task): Handle taskwait with depend clauses. (gimplify_omp_for): Add shared clause on parallel for combined parallel master taskloop{, simd} if taskloop has firstprivate, lastprivate or reduction clause. Use ORT_TASKLOOP or ORT_UNTIED_TASKLOOP instead of ORT_TASK or ORT_UNTIED_TASK. Adjust tests for ORT_COMBINED_TEAMS. Handle C++ range for loops with NULL TREE_PURPOSE in OMP_FOR_ORIG_DECLS. Firstprivatize __for_end and __for_range temporaries on OMP_PARALLEL for distribute parallel for{, simd}. Move OMP_CLAUSE_REDUCTION and OMP_CLAUSE_IN_REDUCTION from taskloop to the task construct sandwiched in between two taskloops. (computable_teams_clause): Test ctx->defaultmap[GDMK_SCALAR] instead of ctx->omp_firstprivatize_variable. (gimplify_omp_workshare): Set ort to ORT_HOST_TEAMS or ORT_COMBINED_HOST_TEAMS if not inside of target construct. If host teams, use gimplify_and_return_first etc. for body like for target or target data constructs, and at the end call gimple_omp_teams_set_host on the GIMPLE_OMP_TEAMS object. (gimplify_omp_atomic): Use OMP_ATOMIC_MEMORY_ORDER instead of OMP_ATOMIC_SEQ_CST, pass it as new argument to gimple_build_omp_atomic_load and gimple_build_omp_atomic_store, remove gimple_omp_atomic_set_seq_cst calls. (gimplify_expr) <case OMP_TASKGROUP>: Move handling into a separate case, handle taskgroup clauses. * lto-streamer-out.c (hash_tree): Handle OMP_CLAUSE_{TASK,IN}_REDUCTION. * Makefile.in (GTFILES): Add omp-general.h. * omp-builtins.def (BUILT_IN_GOMP_TASKWAIT_DEPEND, BUILT_IN_GOMP_LOOP_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_LOOP_MAYBE_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_LOOP_START, BUILT_IN_GOMP_LOOP_ORDERED_START, BUILT_IN_GOMP_LOOP_DOACROSS_START, BUILT_IN_GOMP_LOOP_NONMONOTONIC_RUNTIME_NEXT, BUILT_IN_GOMP_LOOP_MAYBE_NONMONOTONIC_RUNTIME_NEXT, BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_LOOP_ULL_MAYBE_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_LOOP_ULL_START, BUILT_IN_GOMP_LOOP_ULL_ORDERED_START, BUILT_IN_GOMP_LOOP_ULL_DOACROSS_START, BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_RUNTIME_NEXT, BUILT_IN_GOMP_LOOP_ULL_MAYBE_NONMONOTONIC_RUNTIME_NEXT, BUILT_IN_GOMP_PARALLEL_LOOP_NONMONOTONIC_RUNTIME, BUILT_IN_GOMP_PARALLEL_LOOP_MAYBE_NONMONOTONIC_RUNTIME, BUILT_IN_GOMP_PARALLEL_REDUCTIONS, BUILT_IN_GOMP_SECTIONS2_START, BUILT_IN_GOMP_TEAMS_REG, BUILT_IN_GOMP_TASKGROUP_REDUCTION_REGISTER, BUILT_IN_GOMP_TASKGROUP_REDUCTION_UNREGISTER, BUILT_IN_GOMP_TASK_REDUCTION_REMAP, BUILT_IN_GOMP_WORKSHARE_TASK_REDUCTION_UNREGISTER): New builtins. * omp-expand.c (workshare_safe_to_combine_p): Return false for non-worksharing loops. (omp_adjust_chunk_size): Don't adjust anything if chunk_size is zero. (determine_parallel_type): Don't combine parallel with worksharing which has _reductemp_ clause. (expand_parallel_call): Emit the GOMP_*nonmonotonic_runtime* or GOMP_*maybe_nonmonotonic_runtime* builtins instead of GOMP_*runtime* if there is nonmonotonic modifier or if there is no modifier and no ordered clause. For dynamic and guided schedule without monotonic and nonmonotonic modifier, default to nonmonotonic. (expand_omp_for): Likewise. Adjust expand_omp_for_generic caller, use GOMP_loop{,_ull}{,_ordered,_doacross}_start builtins if there are task reductions. (expand_task_call): Add GOMP_TASK_FLAG_REDUCTION flag to flags if there are any reduction clauses. (expand_taskwait_call): New function. (expand_teams_call): New function. (expand_omp_taskreg): Allow GIMPLE_OMP_TEAMS and call expand_teams_call for it. Formatting fix. Handle taskwait with depend clauses. (expand_omp_for_generic): Add SCHED_ARG argument. Handle expansion of worksharing loops with task reductions. (expand_omp_for_static_nochunk, expand_omp_for_static_chunk): Handle expansion of worksharing loops with task reductions. (expand_omp_sections): Handle expansion of sections with task reductions. (expand_omp_synch): For host teams call expand_omp_taskreg. (omp_memory_order_to_memmodel): New function. (expand_omp_atomic_load, expand_omp_atomic_store, expand_omp_atomic_fetch_op): Use it and gimple_omp_atomic_memory_order instead of gimple_omp_atomic_seq_cst_p. (build_omp_regions_1, omp_make_gimple_edges): Treat taskwait with depend clauses as a standalone directive. * omp-general.c (enum omp_requires): New variable. (omp_extract_for_data): Initialize have_reductemp member. Allow NE_EXPR even in OpenMP loops, transform them into LT_EXPR or GT_EXPR loops depending on incr sign. Formatting fixes. * omp-general.h (struct omp_for_data): Add have_reductemp member. (enum omp_requires): New enum. (omp_requires_mask): Declare. * omp-grid.c (grid_eliminate_combined_simd_part): Formatting fix. Fix comment typos. * omp-low.c (struct omp_context): Add task_reductions and task_reduction_map fields. (is_host_teams_ctx): New function. (is_taskreg_ctx): Return true also if is_host_teams_ctx. (use_pointer_for_field): Use is_global_var instead of TREE_STATIC || DECL_EXTERNAL, and apply only if not privatized in outer contexts. (build_outer_var_ref): Ignore taskgroup outer contexts. (delete_omp_context): Release task_reductions and task_reduction_map. (scan_sharing_clauses): Don't add any fields for reduction clause on taskloop. Handle OMP_CLAUSE__REDUCTEMP_. Handle OMP_CLAUSE_{IN,TASK}_REDUCTION and OMP_CLAUSE_REDUCTION with task modifier. Don't ignore shared clauses in is_host_teams_ctx contexts. Handle OMP_CLAUSE_NONTEMPORAL. (add_taskreg_looptemp_clauses): Add OMP_CLAUSE__REDUCTEMP_ clause if needed. (scan_omp_parallel): Add _reductemp_ clause if there are any reduction clauses with task modifier. (scan_omp_task): Handle taskwait with depend clauses. (finish_taskreg_scan): Move field corresponding to _reductemp_ clause first. Move also OMP_CLAUSE__REDUCTEMP_ clause in front if present. Handle GIMPLE_OMP_TEAMS like GIMPLE_OMP_PARALLEL. (scan_omp_for): Fix comment formatting. (scan_omp_teams): Handle host teams constructs. (check_omp_nesting_restrictions): Allow teams with no outer OpenMP context. Adjust diagnostics for teams strictly nested into some explicit OpenMP construct other than target. Allow OpenMP atomics inside of simd regions. (scan_omp_1_stmt): Call scan_sharing_clauses for taskgroups. (scan_omp_1_stmt) <case GIMPLE_OMP_TEAMS>: Temporarily bump taskreg_nesting_level while scanning host teams construct. (task_reduction_read): New function. (lower_rec_input_clauses): Handle OMP_CLAUSE_REDUCTION on taskloop construct. Handle OMP_CLAUSE_IN_REDUCTION and OMP_CLAUSE__REDUCTEMP_ clauses. Handle OMP_CLAUSE_REDUCTION with task modifier. Remove second argument create_tmp_var if it is NULL. Don't ignore shared clauses in is_host_teams_ctx contexts. Handle OMP_CLAUSE_FIRSTPRIVATE_NO_REFERENCE on OMP_CLAUSE_FIRSTPRIVATE clauses. (lower_reduction_clauses): Ignore reduction clauses with task modifier. Remove second argument create_tmp_var if it is NULL. Initialize OMP_ATOMIC_MEMORY_ORDER to relaxed. (lower_send_clauses): Ignore reduction clauses with task modifier. Handle OMP_CLAUSE__REDUCTEMP_. Don't send anything for OMP_CLAUSE_REDUCTION on taskloop. Handle OMP_CLAUSE_IN_REDUCTION. (maybe_add_implicit_barrier_cancel): Add OMP_RETURN argument, don't rely that it is the last stmt in body so far. Ignore outer taskgroup contexts. (omp_task_reductions_find_first, omp_task_reduction_iterate, lower_omp_task_reductions): New functions. (lower_omp_sections): Handle reduction clauses with taskgroup modifiers. Adjust maybe_add_implicit_barrier_cancel caller. (lower_omp_single): Adjust maybe_add_implicit_barrier_cancel caller. (lower_omp_for): Likewise. Handle reduction clauses with taskgroup modifiers. (lower_omp_taskgroup): Handle taskgroup reductions. (create_task_copyfn): Copy over OMP_CLAUSE__REDUCTEMP_ pointer. Handle OMP_CLAUSE_IN_REDUCTION and OMP_CLAUSE_REDUCTION clauses. (lower_depend_clauses): If there are any OMP_CLAUSE_DEPEND_DEPOBJ or OMP_CLAUSE_DEPEND_MUTEXINOUTSET depend clauses, use a new array format. If OMP_CLAUSE_DEPEND_LAST is seen, assume lowering is done already and return early. Set kind on artificial depend clause to OMP_CLAUSE_DEPEND_LAST. (lower_omp_taskreg): Handle reduction clauses with task modifier on parallel construct. Handle reduction clause on taskloop construct. Handle taskwait with depend clauses. (lower_omp_1): Use lower_omp_taskreg instead of lower_omp_teams for host teams constructs. * tree.c (omp_clause_num_ops): Add in_reduction, task_reduction, nontemporal and _reductemp_ clause entries. (omp_clause_code_name): Likewise. (walk_tree_1): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION, OMP_CLAUSE_NONTEMPORAL and OMP_CLAUSE__REDUCTEMP_. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_{{IN,TASK}_REDUCTION,NONTEMPORAL,_REDUCTEMP_}. (enum omp_clause_defaultmap_kind, enum omp_memory_order): New. (struct tree_base): Add omp_atomic_memory_order field into union. Remove OMP_ATOMIC_SEQ_CST comment. (enum omp_clause_depend_kind): Add OMP_CLAUSE_DEPEND_MUTEXINOUTSET and OMP_CLAUSE_DEPEND_DEPOBJ. (struct tree_omp_clause): Add subcode.defaultmap_kind. * tree.def (OMP_TASKGROUP): Add another operand, move next to other OpenMP constructs with body and clauses operands. * tree.h (OMP_BODY): Use OMP_MASTER instead of OMP_TASKGROUP. (OMP_CLAUSES): Use OMP_TASKGROUP instead of OMP_SINGLE. (OMP_TASKGROUP_CLAUSES): Define. (OMP_CLAUSE_DECL): Use OMP_CLAUSE__REDUCTEMP_ instead of OMP_CLAUSE__LOOPTEMP_. (OMP_ATOMIC_SEQ_CST): Remove. (OMP_ATOMIC_MEMORY_ORDER, OMP_CLAUSE_FIRSTPRIVATE_NO_REFERENCE, OMP_CLAUSE_LASTPRIVATE_CONDITIONAL): Define. (OMP_CLAUSE_REDUCTION_CODE, OMP_CLAUSE_REDUCTION_INIT, OMP_CLAUSE_REDUCTION_MERGE, OMP_CLAUSE_REDUCTION_PLACEHOLDER, OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER, OMP_CLAUSE_REDUCTION_OMP_ORIG_REF): Handle OMP_CLAUSE_{,IN_,TASK_}REDUCTION. (OMP_CLAUSE_REDUCTION_TASK, OMP_CLAUSE_REDUCTION_INSCAN, OMP_CLAUSE_DEFAULTMAP_KIND, OMP_CLAUSE_DEFAULTMAP_CATEGORY, OMP_CLAUSE_DEFAULTMAP_BEHAVIOR, OMP_CLAUSE_DEFAULTMAP_SET_KIND): Define. * tree-inline.c (remap_gimple_stmt): Remap taskgroup clauses. * tree-nested.c (convert_nonlocal_omp_clauses): Handle OMP_CLAUSE__REDUCTEMP_, OMP_CLAUSE_NONTEMPORAL. (convert_local_omp_clauses): Likewise. Remove useless test. * tree-parloops.c (create_call_for_reduction_1): Pass OMP_MEMORY_ORDER_RELAXED as new argument to dump_gimple_omp_atomic_load and dump_gimple_omp_atomic_store. * tree-pretty-print.c (dump_omp_iterators): New function. (dump_omp_clause): Handle OMP_CLAUSE__REDUCTEMP_, OMP_CLAUSE_NONTEMPORAL, OMP_CLAUSE_{TASK,IN}_REDUCTION. Print reduction modifiers. Handle OMP_CLAUSE_DEPEND_DEPOBJ and OMP_CLAUSE_DEPEND_MUTEXINOUTSET. Print iterators in depend clauses. Print __internal__ for OMP_CLAUSE_DEPEND_LAST. Handle cancel and simd OMP_CLAUSE_IF_MODIFIERs. Handle new kinds of OMP_CLAUSE_DEFAULTMAP. Print conditional: for OMP_CLAUSE_LASTPRIVATE_CONDITIONAL. (dump_omp_atomic_memory_order): New function. (dump_generic_node): Use it. Print taskgroup clauses. Print taskwait with depend clauses. * tree-pretty-print.h (dump_omp_atomic_memory_order): Declare. * tree-streamer-in.c (unpack_ts_omp_clause_value_fields): Handle OMP_CLAUSE_{TASK,IN}_REDUCTION. * tree-streamer-out.c (pack_ts_omp_clause_value_fields, write_ts_omp_clause_tree_pointers): Likewise. gcc/c-family/ * c-common.h (c_finish_omp_taskgroup): Add CLAUSES argument. (c_finish_omp_atomic): Replace bool SEQ_CST argument with enum omp_memory_order MEMORY_ORDER. (c_finish_omp_flush): Add MO argument. (c_omp_depend_t_p, c_finish_omp_depobj): Declare. (c_finish_omp_for): Add FINAL_P argument. * c-omp.c: Include memmodel.h. (c_finish_omp_taskgroup): Add CLAUSES argument. Set OMP_TASKGROUP_CLAUSES to it. (c_finish_omp_atomic): Replace bool SEQ_CST argument with enum omp_memory_order MEMORY_ORDER. Set OMP_ATOMIC_MEMORY_ORDER instead of OMP_ATOMIC_SEQ_CST. (c_omp_depend_t_p, c_finish_omp_depobj): New functions. (c_finish_omp_flush): Add MO argument, if not MEMMODEL_LAST, emit __atomic_thread_fence call with the given value. (check_omp_for_incr_expr): Formatting fixes. (c_finish_omp_for): Add FINAL_P argument. Allow NE_EXPR even in OpenMP loops, diagnose if NE_EXPR and incr expression is not constant expression 1 or -1. Transform NE_EXPR loops with iterators pointers to VLA into LT_EXPR or GT_EXPR loops. (c_omp_check_loop_iv_r): Look for orig decl of C++ range for loops too. (c_omp_split_clauses): Add support for combined #pragma omp parallel master and #pragma omp {,parallel }master taskloop{, simd} constructs. Handle OMP_CLAUSE_IN_REDUCTION. Handle OMP_CLAUSE_REDUCTION_TASK. Handle OMP_CLAUSE_NONTEMPORAL. Handle splitting OMP_CLAUSE_IF also to OMP_SIMD. Copy OMP_CLAUSE_LASTPRIVATE_CONDITIONAL. (c_omp_predetermined_sharing): Don't return OMP_CLAUSE_DEFAULT_SHARED for const qualified decls. * c-pragma.c (omp_pragmas): Add PRAGMA_OMP_DEPOBJ and PRAGMA_OMP_REQUIRES. * c-pragma.h (enum pragma_kind): Likewise. (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_NONTEMPORAL and PRAGMA_OMP_CLAUSE_{IN,TASK}_REDUCTION. gcc/c/ * c-parser.c: Include memmode.h. (c_parser_omp_depobj, c_parser_omp_requires): New functions. (c_parser_pragma): Handle PRAGMA_OMP_DEPOBJ and PRAGMA_OMP_REQUIRES. (c_parser_omp_clause_name): Handle nontemporal, in_reduction and task_reduction clauses. (c_parser_omp_variable_list): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION. For OMP_CLAUSE_DEPEND, parse clause operands as either an array section, or lvalue assignment expression. (c_parser_omp_clause_if): Handle cancel and simd modifiers. (c_parser_omp_clause_lastprivate): Parse optional conditional: modifier. (c_parser_omp_clause_hint): Require constant integer expression rather than just integer expression. (c_parser_omp_clause_defaultmap): Parse new kinds of defaultmap clause. (c_parser_omp_clause_reduction): Add IS_OMP and KIND arguments. Parse reduction modifiers. Pass KIND to c_parser_omp_variable_list. (c_parser_omp_clause_nontemporal, c_parser_omp_iterators): New functions. (c_parser_omp_clause_depend): Parse iterator modifier and handle iterators. Parse mutexinoutset and depobj kinds. (c_parser_oacc_all_clauses): Adjust c_parser_omp_clause_reduction callers. (c_parser_omp_all_clauses): Likewise. Handle PRAGMA_OMP_CLAUSE_NONTEMPORAL and PRAGMA_OMP_CLAUSE_{IN,TASK}_REDUCTION. (c_parser_omp_atomic): Parse hint and memory order clauses. Handle default memory order from requires directive if any. Adjust c_finish_omp_atomic caller. (c_parser_omp_critical): Allow comma in between (name) and hint clause. (c_parser_omp_flush): Parse flush with memory-order-clause. (c_parser_omp_for_loop): Allow NE_EXPR even in OpenMP loops, adjust c_finish_omp_for caller. (OMP_SIMD_CLAUSE_MASK): Add if and nontemporal clauses. (c_parser_omp_master): Add p_name, mask and cclauses arguments. Allow to be called while parsing combined parallel master. Parse combined master taskloop{, simd}. (c_parser_omp_parallel): Parse combined parallel master{, taskloop{, simd}} constructs. (OMP_TASK_CLAUSE_MASK): Add in_reduction clause. (OMP_TASKGROUP_CLAUSE_MASK): Define. (c_parser_omp_taskgroup): Add LOC argument. Parse taskgroup clauses. (OMP_TASKWAIT_CLAUSE_MASK): Define. (c_parser_omp_taskwait): Handle taskwait with depend clauses. (c_parser_omp_teams): Force a BIND_EXPR with BLOCK around teams body. Use SET_EXPR_LOCATION. (c_parser_omp_target_data): Allow target data with only use_device_ptr clauses. (c_parser_omp_target): Use SET_EXPR_LOCATION. Set OMP_REQUIRES_TARGET_USED bit in omp_requires_mask. (c_parser_omp_requires): New function. (c_finish_taskloop_clauses): New function. (OMP_TASKLOOP_CLAUSE_MASK): Add reduction and in_reduction clauses. (c_parser_omp_taskloop): Use c_finish_taskloop_clauses. Add forward declaration. Disallow in_reduction clause when combined with parallel master. (c_parser_omp_construct): Adjust c_parser_omp_master and c_parser_omp_taskgroup callers. * c-typeck.c (c_finish_omp_cancel): Diagnose if clause with modifier other than cancel. (handle_omp_array_sections_1): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION like OMP_CLAUSE_REDUCTION. (handle_omp_array_sections): Likewise. Call save_expr on array reductions before calling build_index_type. Handle depend clauses with iterators. (struct c_find_omp_var_s): New type. (c_find_omp_var_r, c_omp_finish_iterators): New functions. (c_finish_omp_clauses): Don't diagnose nonmonotonic clause with static, runtime or auto schedule kinds. Call save_expr for whole array reduction sizes. Diagnose reductions with zero sized elements or variable length structures. Diagnose nogroup clause used with reduction clause(s). Handle depend clause with OMP_CLAUSE_DEPEND_DEPOBJ. Diagnose bit-fields. Require omp_depend_t type for OMP_CLAUSE_DEPEND_DEPOBJ kinds and some different type for other kinds. Use build_unary_op with ADDR_EXPR and build_indirect_ref instead of c_mark_addressable. Handle depend clauses with iterators. Remove no longer needed special case that predetermined const qualified vars may be specified in firstprivate clause. Complain if const qualified vars are mentioned in data-sharing clauses other than firstprivate or shared. Use error_at with OMP_CLAUSE_LOCATION (c) as first argument instead of error. Formatting fix. Handle OMP_CLAUSE_NONTEMPORAL and OMP_CLAUSE_{IN,TASK}_REDUCTION. Allow any lvalue as OMP_CLAUSE_DEPEND operand (besides array section), adjust diagnostics. gcc/cp/ * constexpr.c (potential_constant_expression_1): Handle OMP_DEPOBJ. * cp-gimplify.c (cp_genericize_r): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION. (cxx_omp_predetermined_sharing_1): Don't return OMP_CLAUSE_DEFAULT_SHARED for const qualified decls with no mutable member. Return OMP_CLAUSE_DEFAULT_FIRSTPRIVATE for this pointer. * cp-objcp-common.c (cp_common_init_ts): Handle OMP_DEPOBJ. * cp-tree.def (OMP_DEPOBJ): New tree code. * cp-tree.h (OMP_ATOMIC_DEPENDENT_P): Return true also for first argument being OMP_CLAUSE. (OMP_DEPOBJ_DEPOBJ, OMP_DEPOBJ_CLAUSES): Define. (cp_convert_omp_range_for, cp_finish_omp_range_for): Declare. (finish_omp_atomic): Add LOC, CLAUSES and MO arguments. Remove SEQ_CST argument. (finish_omp_for_block): Declare. (finish_omp_flush): Add MO argument. (finish_omp_depobj): Declare. * cxx-pretty-print.c (cxx_pretty_printer::statement): Handle OMP_DEPOBJ. * dump.c (cp_dump_tree): Likewise. * lex.c (cxx_init): Likewise. * parser.c: Include memmodel.h. (cp_parser_for): Pass false as new is_omp argument to cp_parser_range_for. (cp_parser_range_for): Add IS_OMP argument, return before finalizing if it is true. (cp_parser_omp_clause_name): Handle nontemporal, in_reduction and task_reduction clauses. (cp_parser_omp_var_list_no_open): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION. For OMP_CLAUSE_DEPEND, parse clause operands as either an array section, or lvalue assignment expression. (cp_parser_omp_clause_if): Handle cancel and simd modifiers. (cp_parser_omp_clause_defaultmap): Parse new kinds of defaultmap clause. (cp_parser_omp_clause_reduction): Add IS_OMP and KIND arguments. Parse reduction modifiers. Pass KIND to c_parser_omp_variable_list. (cp_parser_omp_clause_lastprivate, cp_parser_omp_iterators): New functions. (cp_parser_omp_clause_depend): Parse iterator modifier and handle iterators. Parse mutexinoutset and depobj kinds. (cp_parser_oacc_all_clauses): Adjust cp_parser_omp_clause_reduction callers. (cp_parser_omp_all_clauses): Likewise. Handle PRAGMA_OMP_CLAUSE_NONTEMPORAL and PRAGMA_OMP_CLAUSE_{IN,TASK}_REDUCTION. Call cp_parser_omp_clause_lastprivate for OpenMP lastprivate clause. (cp_parser_omp_atomic): Pass pragma_tok->location as LOC to finish_omp_atomic. Parse hint and memory order clauses. Handle default memory order from requires directive if any. Adjust finish_omp_atomic caller. (cp_parser_omp_critical): Allow comma in between (name) and hint clause. (cp_parser_omp_depobj): New function. (cp_parser_omp_flush): Parse flush with memory-order-clause. (cp_parser_omp_for_cond): Allow NE_EXPR even in OpenMP loops. (cp_convert_omp_range_for, cp_finish_omp_range_for): New functions. (cp_parser_omp_for_loop): Parse C++11 range for loops among omp loops. Handle OMP_CLAUSE_IN_REDUCTION like OMP_CLAUSE_REDUCTION. (OMP_SIMD_CLAUSE_MASK): Add if and nontemporal clauses. (cp_parser_omp_simd, cp_parser_omp_for): Call keep_next_level before begin_omp_structured_block and call finish_omp_for_block on finish_omp_structured_block result. (cp_parser_omp_master): Add p_name, mask and cclauses arguments. Allow to be called while parsing combined parallel master. Parse combined master taskloop{, simd}. (cp_parser_omp_parallel): Parse combined parallel master{, taskloop{, simd}} constructs. (cp_parser_omp_single): Use SET_EXPR_LOCATION. (OMP_TASK_CLAUSE_MASK): Add in_reduction clause. (OMP_TASKWAIT_CLAUSE_MASK): Define. (cp_parser_omp_taskwait): Handle taskwait with depend clauses. (OMP_TASKGROUP_CLAUSE_MASK): Define. (cp_parser_omp_taskgroup): Parse taskgroup clauses, adjust c_finish_omp_taskgroup caller. (cp_parser_omp_distribute): Call keep_next_level before begin_omp_structured_block and call finish_omp_for_block on finish_omp_structured_block result. (cp_parser_omp_teams): Force a BIND_EXPR with BLOCK around teams body. (cp_parser_omp_target_data): Allow target data with only use_device_ptr clauses. (cp_parser_omp_target): Set OMP_REQUIRES_TARGET_USED bit in omp_requires_mask. (cp_parser_omp_requires): New function. (OMP_TASKLOOP_CLAUSE_MASK): Add reduction and in_reduction clauses. (cp_parser_omp_taskloop): Add forward declaration. Disallow in_reduction clause when combined with parallel master. Call keep_next_level before begin_omp_structured_block and call finish_omp_for_block on finish_omp_structured_block result. (cp_parser_omp_construct): Adjust cp_parser_omp_master caller. (cp_parser_pragma): Handle PRAGMA_OMP_DEPOBJ and PRAGMA_OMP_REQUIRES. * pt.c (tsubst_omp_clause_decl): Add iterators_cache argument. Adjust recursive calls. Handle iterators. (tsubst_omp_clauses): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION and OMP_CLAUSE_NONTEMPORAL. Adjust tsubst_omp_clause_decl callers. (tsubst_decomp_names): (tsubst_omp_for_iterator): Change orig_declv into a reference. Handle range for loops. Move orig_declv handling after declv/initv handling. (tsubst_expr): Force a BIND_EXPR with BLOCK around teams body. Adjust finish_omp_atomic caller. Call keep_next_level before begin_omp_structured_block. Call cp_finish_omp_range_for for range for loops and use {begin,finish}_omp_structured_block instead of {push,pop}_stmt_list if there are any range for loops. Call finish_omp_for_block on finish_omp_structured_block result. Handle OMP_DEPOBJ. Handle taskwait with depend clauses. For OMP_ATOMIC call tsubst_omp_clauses on clauses if any, adjust finish_omp_atomic caller. Use OMP_ATOMIC_MEMORY_ORDER rather than OMP_ATOMIC_SEQ_CST. Handle clauses on OMP_TASKGROUP. (dependent_omp_for_p): Always return true for range for loops if processing_template_decl. Return true if class type iterator does not have INTEGER_CST increment. * semantics.c: Include memmodel.h. (handle_omp_array_sections_1): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION like OMP_CLAUSE_REDUCTION. (handle_omp_array_sections): Likewise. Call save_expr on array reductions before calling build_index_type. Handle depend clauses with iterators. (finish_omp_reduction_clause): Call save_expr for whole array reduction sizes. Don't mark OMP_CLAUSE_DECL addressable if it has reference type. Do mark decl_placeholder addressable if needed. Use error_at with OMP_CLAUSE_LOCATION (c) as first argument instead of error. (cp_omp_finish_iterators): New function. (finish_omp_clauses): Don't diagnose nonmonotonic clause with static, runtime or auto schedule kinds. Diagnose nogroup clause used with reduction clause(s). Handle depend clause with OMP_CLAUSE_DEPEND_DEPOBJ. Diagnose bit-fields. Require omp_depend_t type for OMP_CLAUSE_DEPEND_DEPOBJ kinds and some different type for other kinds. Use cp_build_addr_expr and cp_build_indirect_ref instead of cxx_mark_addressable. Handle depend clauses with iterators. Only handle static data members in the special case that const qualified vars may be specified in firstprivate clause. Complain if const qualified vars without mutable members are mentioned in data-sharing clauses other than firstprivate or shared. Use error_at with OMP_CLAUSE_LOCATION (c) as first argument instead of error. Diagnose more than one nontemporal clause refering to the same variable. Use error_at rather than error for priority and hint clause diagnostics. Fix pasto for hint clause. Diagnose hint expression that doesn't fold into INTEGER_CST. Diagnose if clause with modifier other than cancel. Handle OMP_CLAUSE_{IN,TASK}_REDUCTION like OMP_CLAUSE_REDUCTION. Allow any lvalue as OMP_CLAUSE_DEPEND operand (besides array section), adjust diagnostics. (handle_omp_for_class_iterator): Don't create a new TREE_LIST if one has been created already for range for, just fill TREE_PURPOSE and TREE_VALUE. Call cp_fully_fold on incr. (finish_omp_for): Don't check cond/incr if cond is global_namespace. Pass to c_omp_check_loop_iv_exprs orig_declv if non-NULL. Don't use IS_EMPTY_STMT on NULL pre_body. Adjust c_finish_omp_for caller. (finish_omp_for_block): New function. (finish_omp_atomic): Add LOC argument, pass it through to c_finish_omp_atomic and set it as location of OMP_ATOMIC* trees. Remove SEQ_CST argument. Add CLAUSES and MO arguments. Adjust c_finish_omp_atomic caller. Stick clauses if any into first argument of wrapping OMP_ATOMIC. (finish_omp_depobj): New function. (finish_omp_flush): Add MO argument, if not MEMMODEL_LAST, emit __atomic_thread_fence call with the given value. (finish_omp_cancel): Diagnose if clause with modifier other than cancel. gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Use OMP_CLAUSE_DEFAULTMAP_SET_KIND. (gfc_trans_omp_atomic): Set OMP_ATOMIC_MEMORY_ORDER rather than OMP_ATOMIC_SEQ_CST. (gfc_trans_omp_taskgroup): Build OMP_TASKGROUP using make_node instead of build1_loc. * types.def (BT_FN_VOID_BOOL, BT_FN_VOID_SIZE_SIZE_PTR, BT_FN_UINT_UINT_PTR_PTR, BT_FN_UINT_OMPFN_PTR_UINT_UINT, BT_FN_BOOL_UINT_LONGPTR_LONG_LONG_LONGPTR_LONGPTR_PTR_PTR, BT_FN_BOOL_UINT_ULLPTR_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR, BT_FN_BOOL_LONG_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR_PTR_PTR, BT_FN_BOOL_BOOL_ULL_ULL_ULL_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR): New. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR): Formatting fix. gcc/testsuite/ * c-c++-common/gomp/atomic-17.c: New test. * c-c++-common/gomp/atomic-18.c: New test. * c-c++-common/gomp/atomic-19.c: New test. * c-c++-common/gomp/atomic-20.c: New test. * c-c++-common/gomp/atomic-21.c: New test. * c-c++-common/gomp/atomic-22.c: New test. * c-c++-common/gomp/clauses-1.c (r2): New variable. (foo): Add ntm argument and test if and nontemporal clauses on constructs with simd. (bar): Put taskloop simd inside of taskgroup with task_reduction, use in_reduction clause instead of reduction. Add another taskloop simd without nogroup clause, but with reduction clause and a new in_reduction. Add ntm and i3 arguments. Test if and nontemporal clauses on constructs with simd. Change if clauses on some constructs from specific to the particular constituents to one without a modifier. Add new tests for combined host teams and for new parallel master and {,parallel }master taskloop{, simd} combined constructs. (baz): New function with host teams tests. * gcc.dg/gomp/combined-1.c: Moved to ... * c-c++-common/gomp/combined-1.c: ... here. Adjust expected library call. * c-c++-common/gomp/combined-2.c: New test. * c-c++-common/gomp/combined-3.c: New test. * c-c++-common/gomp/critical-1.c: New test. * c-c++-common/gomp/critical-2.c: New test. * c-c++-common/gomp/default-1.c: New test. * c-c++-common/gomp/defaultmap-1.c: New test. * c-c++-common/gomp/defaultmap-2.c: New test. * c-c++-common/gomp/defaultmap-3.c: New test. * c-c++-common/gomp/depend-5.c: New test. * c-c++-common/gomp/depend-6.c: New test. * c-c++-common/gomp/depend-iterator-1.c: New test. * c-c++-common/gomp/depend-iterator-2.c: New test. * c-c++-common/gomp/depobj-1.c: New test. * c-c++-common/gomp/flush-1.c: New test. * c-c++-common/gomp/flush-2.c: New test. * c-c++-common/gomp/for-1.c: New test. * c-c++-common/gomp/for-2.c: New test. * c-c++-common/gomp/for-3.c: New test. * c-c++-common/gomp/for-4.c: New test. * c-c++-common/gomp/for-5.c: New test. * c-c++-common/gomp/for-6.c: New test. * c-c++-common/gomp/for-7.c: New test. * c-c++-common/gomp/if-1.c (foo): Add some further tests. * c-c++-common/gomp/if-2.c (foo): Likewise. Expect slightly different diagnostics wording in one case. * c-c++-common/gomp/if-3.c: New test. * c-c++-common/gomp/master-combined-1.c: New test. * c-c++-common/gomp/master-combined-2.c: New test. * c-c++-common/gomp/nontemporal-1.c: New test. * c-c++-common/gomp/nontemporal-2.c: New test. * c-c++-common/gomp/reduction-task-1.c: New test. * c-c++-common/gomp/reduction-task-2.c: New test. * c-c++-common/gomp/requires-1.c: New test. * c-c++-common/gomp/requires-2.c: New test. * c-c++-common/gomp/requires-3.c: New test. * c-c++-common/gomp/requires-4.c: New test. * c-c++-common/gomp/schedule-modifiers-1.c (bar): Don't expect diagnostics for nonmonotonic modifier with static, runtime or auto schedule kinds. * c-c++-common/gomp/simd7.c: New test. * c-c++-common/gomp/target-data-1.c: New test. * c-c++-common/gomp/taskloop-reduction-1.c: New test. * c-c++-common/gomp/taskwait-depend-1.c: New test. * c-c++-common/gomp/teams-1.c: New test. * c-c++-common/gomp/teams-2.c: New test. * gcc.dg/gomp/appendix-a/a.24.1.c: Update from OpenMP examples. Add shared(c) clause. * gcc.dg/gomp/atomic-5.c (f1): Add another expected error. * gcc.dg/gomp/clause-1.c: Adjust expected diagnostics for const qualified vars without mutable member no longer being predeterined shared. * gcc.dg/gomp/sharing-1.c: Likewise. * g++.dg/gomp/clause-3.C: Likewise. * g++.dg/gomp/member-2.C: Likewise. * g++.dg/gomp/predetermined-1.C: Likewise. * g++.dg/gomp/private-1.C: Likewise. * g++.dg/gomp/sharing-1.C: Likewise. * g++.dg/gomp/sharing-2.C: Likewise. Add a few tests with aggregate const static data member without mutable elements. * gcc.dg/gomp/for-4.c: Expected nonmonotonic functions in the dumps. * gcc.dg/gomp/for-5.c: Likewise. * gcc.dg/gomp/for-6.c: Change expected library call. * gcc.dg/gomp/pr39495-2.c (foo): Don't expect errors on !=. * gcc.dg/gomp/reduction-2.c: New test. * gcc.dg/gomp/simd-1.c: New test. * gcc.dg/gomp/teams-1.c: Adjust expected diagnostic lines. * g++.dg/gomp/atomic-18.C: New test. * g++.dg/gomp/atomic-19.C: New test. * g++.dg/gomp/atomic-5.C (f1): Adjust expected lines of read-only variable messages. Add another expected error. * g++.dg/gomp/critical-3.C: New test. * g++.dg/gomp/depend-iterator-1.C: New test. * g++.dg/gomp/depend-iterator-2.C: New test. * g++.dg/gomp/depobj-1.C: New test. * g++.dg/gomp/doacross-1.C: New test. * g++.dg/gomp/for-21.C: New test. * g++.dg/gomp/for-4.C: Expected nonmonotonic functions in the dumps. * g++.dg/gomp/for-5.C: Likewise. * g++.dg/gomp/for-6.C: Change expected library call. * g++.dg/gomp/loop-4.C: New test. * g++.dg/gomp/pr33372-1.C: Adjust location of the expected diagnostics. * g++.dg/gomp/pr33372-3.C: Likewise. * g++.dg/gomp/pr39495-2.C (foo): Don't expect errors on !=. * g++.dg/gomp/simd-2.C: New test. * g++.dg/gomp/tpl-atomic-2.C: Adjust expected diagnostic lines. include/ * gomp-constants.h (GOMP_TASK_FLAG_REDUCTION, GOMP_DEPEND_IN, GOMP_DEPEND_OUT, GOMP_DEPEND_INOUT, GOMP_DEPEND_MUTEXINOUTSET): Define. libgomp/ * affinity.c (gomp_display_affinity_place): New function. * affinity-fmt.c: New file. * alloc.c (gomp_aligned_alloc, gomp_aligned_free): New functions. * config/linux/affinity.c (gomp_display_affinity_place): New function. * config/nvptx/icv-device.c (omp_get_num_teams, omp_get_team_num): Move these functions to ... * config/nvptx/teams.c: ... here. New file. * config/nvptx/target.c (omp_pause_resource, omp_pause_resource_all): New functions. * config/nvptx/team.c (gomp_team_start, gomp_pause_host): New functions. * configure.ac: Check for aligned_alloc, posix_memalign, memalign and _aligned_malloc. (HAVE_UNAME, HAVE_GETHOSTNAME, HAVE_GETPID): Add new tests. * configure.tgt: Add -DUSING_INITIAL_EXEC_TLS to XCFLAGS for Linux. * env.c (gomp_display_affinity_var, gomp_affinity_format_var, gomp_affinity_format_len): New variables. (parse_schedule): Parse monotonic and nonmonotonic modifiers in OMP_SCHEDULE variable. Set GFS_MONOTONIC for monotonic schedules. (handle_omp_display_env): Display monotonic/nonmonotonic schedule modifiers. Display (non-default) chunk sizes. Print OMP_DISPLAY_AFFINITY and OMP_AFFINITY_FORMAT. (initialize_env): Don't call pthread_attr_setdetachstate. Handle OMP_DISPLAY_AFFINITY and OMP_AFFINITY_FORMAT env vars. * fortran.c: Include stdio.h and string.h. (omp_pause_resource, omp_pause_resource_all): Add ialias_redirect. (omp_get_schedule_, omp_get_schedule_8_): Mask off GFS_MONOTONIC bit. (omp_set_affinity_format_, omp_get_affinity_format_, omp_display_affinity_, omp_capture_affinity_, omp_pause_resource_, omp_pause_resource_all_): New functions. * icv.c (omp_set_schedule): Mask off omp_sched_monotonic bit in switch. * icv-device.c (omp_get_num_teams, omp_get_team_num): Move these functions to ... * teams.c: ... here. New file. * libgomp_g.h: Include gstdint.h. (GOMP_loop_nonmonotonic_runtime_start, GOMP_loop_maybe_nonmonotonic_runtime_start, GOMP_loop_start, GOMP_loop_ordered_start, GOMP_loop_nonmonotonic_runtime_next, GOMP_loop_maybe_nonmonotonic_runtime_next, GOMP_loop_doacross_start, GOMP_parallel_loop_nonmonotonic_runtime, GOMP_parallel_loop_maybe_nonmonotonic_runtime, GOMP_loop_ull_nonmonotonic_runtime_start, GOMP_loop_ull_maybe_nonmonotonic_runtime_start, GOMP_loop_ull_start, GOMP_loop_ull_ordered_start, GOMP_loop_ull_nonmonotonic_runtime_next, GOMP_loop_ull_maybe_nonmonotonic_runtime_next, GOMP_loop_ull_doacross_start, GOMP_parallel_reductions, GOMP_taskwait_depend, GOMP_taskgroup_reduction_register, GOMP_taskgroup_reduction_unregister, GOMP_task_reduction_remap, GOMP_workshare_task_reduction_unregister, GOMP_sections2_start, GOMP_teams_reg): Declare. * libgomp.h (GOMP_HAVE_EFFICIENT_ALIGNED_ALLOC): Define unless gomp_aligned_alloc uses fallback implementation. (gomp_aligned_alloc, gomp_aligned_free): Declare. (enum gomp_schedule_type): Add GFS_MONOTONIC. (struct gomp_doacross_work_share): Add extra field. (struct gomp_work_share): Add task_reductions field. (struct gomp_taskgroup): Add workshare and reductions fields. (GOMP_NEEDS_THREAD_HANDLE): Define if needed. (gomp_thread_handle): New typedef. (gomp_display_affinity_place, gomp_set_affinity_format, gomp_display_string, gomp_display_affinity, gomp_display_affinity_thread): Declare. (gomp_doacross_init, gomp_doacross_ull_init): Add size_t argument. (gomp_parallel_reduction_register, gomp_workshare_taskgroup_start, gomp_workshare_task_reduction_register): Declare. (gomp_team_start): Add taskgroup argument. (gomp_pause_host): Declare. (gomp_init_work_share, gomp_work_share_start): Change bool argument to size_t. (gomp_thread_self, gomp_thread_to_pthread_t): New inline functions. * libgomp.map (GOMP_5.0): Export GOMP_loop_start, GOMP_loop_ordered_start, GOMP_loop_doacross_start, GOMP_loop_ull_start, GOMP_loop_ull_ordered_start, GOMP_loop_ull_doacross_start, GOMP_workshare_task_reduction_unregister, GOMP_sections2_start, GOMP_loop_maybe_nonmonotonic_runtime_next, GOMP_loop_maybe_nonmonotonic_runtime_start, GOMP_loop_nonmonotonic_runtime_next, GOMP_loop_nonmonotonic_runtime_start, GOMP_loop_ull_maybe_nonmonotonic_runtime_next, GOMP_loop_ull_maybe_nonmonotonic_runtime_start, GOMP_loop_ull_nonmonotonic_runtime_next, GOMP_loop_ull_nonmonotonic_runtime_start, GOMP_parallel_loop_maybe_nonmonotonic_runtime, GOMP_parallel_loop_nonmonotonic_runtime, GOMP_parallel_reductions, GOMP_taskgroup_reduction_register, GOMP_taskgroup_reduction_unregister, GOMP_task_reduction_remap, GOMP_teams_reg and GOMP_taskwait_depend. (OMP_5.0): Export omp_pause_resource{,_all}{,_}, omp_{capture,display}_affinity{,_}, and omp_[gs]et_affinity_format{,_}. * loop.c: Include string.h. (GOMP_loop_runtime_next): Add ialias. (GOMP_taskgroup_reduction_register): Add ialias_redirect. (gomp_loop_static_start, gomp_loop_dynamic_start, gomp_loop_guided_start, gomp_loop_ordered_static_start, gomp_loop_ordered_dynamic_start, gomp_loop_ordered_guided_start, gomp_loop_doacross_static_start, gomp_loop_doacross_dynamic_start, gomp_loop_doacross_guided_start): Adjust gomp_work_share_start or gomp_doacross_init callers. (gomp_adjust_sched, GOMP_loop_start, GOMP_loop_ordered_start, GOMP_loop_doacross_start): New functions. (GOMP_loop_runtime_start, GOMP_loop_ordered_runtime_start, GOMP_loop_doacross_runtime_start, GOMP_parallel_loop_runtime_start): Mask off GFS_MONOTONIC bit. (GOMP_loop_maybe_nonmonotonic_runtime_next, GOMP_loop_maybe_nonmonotonic_runtime_start, GOMP_loop_nonmonotonic_runtime_next, GOMP_loop_nonmonotonic_runtime_start, GOMP_parallel_loop_maybe_nonmonotonic_runtime, GOMP_parallel_loop_nonmonotonic_runtime): New aliases or wrapper functions. (gomp_parallel_loop_start): Pass NULL as taskgroup to gomp_team_start. * loop_ull.c: Include string.h. (GOMP_loop_ull_runtime_next): Add ialias. (GOMP_taskgroup_reduction_register): Add ialias_redirect. (gomp_loop_ull_static_start, gomp_loop_ull_dynamic_start, gomp_loop_ull_guided_start, gomp_loop_ull_ordered_static_start, gomp_loop_ull_ordered_dynamic_start, gomp_loop_ull_ordered_guided_start, gomp_loop_ull_doacross_static_start, gomp_loop_ull_doacross_dynamic_start, gomp_loop_ull_doacross_guided_start): Adjust gomp_work_share_start and gomp_doacross_ull_init callers. (gomp_adjust_sched, GOMP_loop_ull_start, GOMP_loop_ull_ordered_start, GOMP_loop_ull_doacross_start): New functions. (GOMP_loop_ull_runtime_start, GOMP_loop_ull_ordered_runtime_start, GOMP_loop_ull_doacross_runtime_start): Mask off GFS_MONOTONIC bit. (GOMP_loop_ull_maybe_nonmonotonic_runtime_next, GOMP_loop_ull_maybe_nonmonotonic_runtime_start, GOMP_loop_ull_nonmonotonic_runtime_next, GOMP_loop_ull_nonmonotonic_runtime_start): Likewise. * Makefile.am (libgomp_la_SOURCES): Add teams.c and affinity-fmt.c. * omp.h.in (enum omp_sched_t): Add omp_sched_monotonic. (omp_pause_resource_t, omp_depend_t): New typedefs. (enum omp_lock_hint_t): Renamed to ... (enum omp_sync_hint_t): ... this. Define omp_sync_hint_* enumerators using numbers and omp_lock_hint_* as their aliases. (omp_lock_hint_t): New typedef. Rename to ... (omp_sync_hint_t): ... this. (omp_init_lock_with_hint, omp_init_nest_lock_with_hint): Use omp_sync_hint_t instead of omp_lock_hint_t. (omp_pause_resource, omp_pause_resource_all, omp_set_affinity_format, omp_get_affinity_format, omp_display_affinity, omp_capture_affinity): Declare. (omp_target_is_present, omp_target_disassociate_ptr): Change first argument from void * to const void *. (omp_target_memcpy, omp_target_memcpy_rect): Change second argument from void * to const void *. (omp_target_associate_ptr): Change first and second arguments from void * to const void *. * omp_lib.f90.in (omp_pause_resource_kind, omp_pause_soft, omp_pause_hard): New parameters. (omp_pause_resource, omp_pause_resource_all, omp_set_affinity_format, omp_get_affinity_format, omp_display_affinity, omp_capture_affinity): New interfaces. * omp_lib.h.in (omp_pause_resource_kind, omp_pause_soft, omp_pause_hard): New parameters. (omp_pause_resource, omp_pause_resource_all, omp_set_affinity_format, omp_get_affinity_format, omp_display_affinity, omp_capture_affinity): New externals. * ordered.c (gomp_doacross_init, gomp_doacross_ull_init): Add EXTRA argument. If not needed to prepare array, if extra is 0, clear ws->doacross, otherwise allocate just doacross structure and extra payload. If array is needed, allocate also extra payload. (GOMP_doacross_post, GOMP_doacross_wait, GOMP_doacross_ull_post, GOMP_doacross_ull_wait): Handle doacross->array == NULL like doacross == NULL. * parallel.c (GOMP_parallel_start): Pass NULL as taskgroup to gomp_team_start. (GOMP_parallel): Likewise. Formatting fix. (GOMP_parallel_reductions): New function. (GOMP_cancellation_point): If taskgroup has workshare flag set, check cancelled of prev taskgroup if any. (GOMP_cancel): If taskgroup has workshare flag set, set cancelled on prev taskgroup if any. * sections.c: Include string.h. (GOMP_taskgroup_reduction_register): Add ialias_redirect. (GOMP_sections_start): Adjust gomp_work_share_start caller. (GOMP_sections2_start): New function. (GOMP_parallel_sections_start, GOMP_parallel_sections): Pass NULL as taskgroup to gomp_team_start. * single.c (GOMP_single_start, GOMP_single_copy_start): Adjust gomp_work_share_start callers. * target.c (GOMP_target_update_ext, GOMP_target_enter_exit_data): If taskgroup has workshare flag set, check cancelled on prev taskgroup if any. Guard all cancellation tests with gomp_cancel_var test. (omp_target_is_present, omp_target_disassociate_ptr): Change ptr argument from void * to const void *. (omp_target_memcpy): Change src argument from void * to const void *. (omp_target_memcpy_rect): Likewise. (omp_target_memcpy_rect_worker): Likewise. Use const char * casts instead of char * where needed. (omp_target_associate_ptr): Change host_ptr and device_ptr arguments from void * to const void *. (omp_pause_resource, omp_pause_resource_all): New functions. * task.c (gomp_task_handle_depend): Handle new depend array format in addition to the old. Handle mutexinoutset kinds the same as inout for now, handle unspecified kinds. (gomp_create_target_task): If taskgroup has workshare flag set, check cancelled on prev taskgroup if any. Guard all cancellation tests with gomp_cancel_var test. Handle new depend array format count in addition to the old. (GOMP_task): Likewise. Adjust function comment. (gomp_task_run_pre): If taskgroup has workshare flag set, check cancelled on prev taskgroup if any. Guard all cancellation tests with gomp_cancel_var test. (GOMP_taskwait_depend): New function. (gomp_task_maybe_wait_for_dependencies): Handle new depend array format in addition to the old. Handle mutexinoutset kinds the same as inout for now, handle unspecified kinds. Fix a function comment typo. (gomp_taskgroup_init): New function. (GOMP_taskgroup_start): Use it. (gomp_reduction_register, gomp_create_artificial_team, GOMP_taskgroup_reduction_register, GOMP_taskgroup_reduction_unregister, GOMP_task_reduction_remap, gomp_parallel_reduction_register, gomp_workshare_task_reduction_register, gomp_workshare_taskgroup_start, GOMP_workshare_task_reduction_unregister): New functions. * taskloop.c (GOMP_taskloop): If taskgroup has workshare flag set, check cancelled on prev taskgroup if any. Guard all cancellation tests with gomp_cancel_var test. Handle GOMP_TASK_FLAG_REDUCTION flag by calling GOMP_taskgroup_reduction_register. * team.c (gomp_thread_attr): Remove comment. (struct gomp_thread_start_data): Add handle field. (gomp_thread_start): Call pthread_detach. (gomp_new_team): Adjust gomp_init_work_share caller. (gomp_free_pool_helper): Call pthread_detach. (gomp_team_start): Add taskgroup argument, initialize implicit tasks' taskgroup field to that. Don't call pthread_attr_setdetachstate. Handle OMP_DISPLAY_AFFINITY env var. (gomp_team_end): Determine nesting by thr->ts.level != 0 rather than thr->ts.team != NULL. (gomp_pause_pool_helper, gomp_pause_host): New functions. * work.c (alloc_work_share): Use gomp_aligned_alloc instead of gomp_malloc if GOMP_HAVE_EFFICIENT_ALIGNED_ALLOC is defined. (gomp_init_work_share): Change ORDERED argument from bool to size_t, if more than 1 allocate also extra payload at the end of array. Never keep ordered_team_ids NULL, set it to inline_ordered_team_ids instead. (gomp_work_share_start): Change ORDERED argument from bool to size_t, return true instead of ws. * Makefile.in: Regenerated. * configure: Regenerated. * config.h.in: Regenerated. * testsuite/libgomp.c/cancel-for-2.c (foo): Use cancel modifier in some cases. * testsuite/libgomp.c-c++-common/cancel-parallel-1.c: New test. * testsuite/libgomp.c-c++-common/cancel-taskgroup-3.c: New test. * testsuite/libgomp.c-c++-common/depend-iterator-1.c: New test. * testsuite/libgomp.c-c++-common/depend-iterator-2.c: New test. * testsuite/libgomp.c-c++-common/depend-mutexinout-1.c: New test. * testsuite/libgomp.c-c++-common/depend-mutexinout-2.c: New test. * testsuite/libgomp.c-c++-common/depobj-1.c: New test. * testsuite/libgomp.c-c++-common/display-affinity-1.c: New test. * testsuite/libgomp.c-c++-common/for-10.c: New test. * testsuite/libgomp.c-c++-common/for-11.c: New test. * testsuite/libgomp.c-c++-common/for-12.c: New test. * testsuite/libgomp.c-c++-common/for-13.c: New test. * testsuite/libgomp.c-c++-common/for-14.c: New test. * testsuite/libgomp.c-c++-common/for-15.c: New test. * testsuite/libgomp.c-c++-common/for-2.h: If CONDNE macro is defined, define a different N(test), don't define N(f0) to N(f14), but instead define N(f20) to N(f34) using != comparisons. * testsuite/libgomp.c-c++-common/for-7.c: New test. * testsuite/libgomp.c-c++-common/for-8.c: New test. * testsuite/libgomp.c-c++-common/for-9.c: New test. * testsuite/libgomp.c-c++-common/master-combined-1.c: New test. * testsuite/libgomp.c-c++-common/pause-1.c: New test. * testsuite/libgomp.c-c++-common/pause-2.c: New test. * testsuite/libgomp.c-c++-common/pr66199-10.c: New test. * testsuite/libgomp.c-c++-common/pr66199-11.c: New test. * testsuite/libgomp.c-c++-common/pr66199-12.c: New test. * testsuite/libgomp.c-c++-common/pr66199-13.c: New test. * testsuite/libgomp.c-c++-common/pr66199-14.c: New test. * testsuite/libgomp.c-c++-common/simd-1.c: New test. * testsuite/libgomp.c-c++-common/taskloop-reduction-1.c: New test. * testsuite/libgomp.c-c++-common/taskloop-reduction-2.c: New test. * testsuite/libgomp.c-c++-common/taskloop-reduction-3.c: New test. * testsuite/libgomp.c-c++-common/taskloop-reduction-4.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-11.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-12.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-1.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-2.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-3.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-4.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-5.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-6.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-7.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-8.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-9.c: New test. * testsuite/libgomp.c-c++-common/taskwait-depend-1.c: New test. * testsuite/libgomp.c++/depend-1.C: New test. * testsuite/libgomp.c++/depend-iterator-1.C: New test. * testsuite/libgomp.c++/depobj-1.C: New test. * testsuite/libgomp.c++/for-16.C: New test. * testsuite/libgomp.c++/for-21.C: New test. * testsuite/libgomp.c++/for-22.C: New test. * testsuite/libgomp.c++/for-23.C: New test. * testsuite/libgomp.c++/for-24.C: New test. * testsuite/libgomp.c++/for-25.C: New test. * testsuite/libgomp.c++/for-26.C: New test. * testsuite/libgomp.c++/taskloop-reduction-1.C: New test. * testsuite/libgomp.c++/taskloop-reduction-2.C: New test. * testsuite/libgomp.c++/taskloop-reduction-3.C: New test. * testsuite/libgomp.c++/taskloop-reduction-4.C: New test. * testsuite/libgomp.c++/task-reduction-10.C: New test. * testsuite/libgomp.c++/task-reduction-11.C: New test. * testsuite/libgomp.c++/task-reduction-12.C: New test. * testsuite/libgomp.c++/task-reduction-13.C: New test. * testsuite/libgomp.c++/task-reduction-14.C: New test. * testsuite/libgomp.c++/task-reduction-15.C: New test. * testsuite/libgomp.c++/task-reduction-16.C: New test. * testsuite/libgomp.c++/task-reduction-17.C: New test. * testsuite/libgomp.c++/task-reduction-18.C: New test. * testsuite/libgomp.c++/task-reduction-19.C: New test. * testsuite/libgomp.c/task-reduction-1.c: New test. * testsuite/libgomp.c++/task-reduction-1.C: New test. * testsuite/libgomp.c/task-reduction-2.c: New test. * testsuite/libgomp.c++/task-reduction-2.C: New test. * testsuite/libgomp.c++/task-reduction-3.C: New test. * testsuite/libgomp.c++/task-reduction-4.C: New test. * testsuite/libgomp.c++/task-reduction-5.C: New test. * testsuite/libgomp.c++/task-reduction-6.C: New test. * testsuite/libgomp.c++/task-reduction-7.C: New test. * testsuite/libgomp.c++/task-reduction-8.C: New test. * testsuite/libgomp.c++/task-reduction-9.C: New test. * testsuite/libgomp.c/teams-1.c: New test. * testsuite/libgomp.c/teams-2.c: New test. * testsuite/libgomp.c/thread-limit-4.c: New test. * testsuite/libgomp.c/thread-limit-5.c: New test. * testsuite/libgomp.fortran/display-affinity-1.f90: New test. From-SVN: r265930 |
||
Chung-Lin Tang
|
829c6349e9 |
Update OpenACC data clause semantics to the 2.5 behavior
gcc/c-family/ * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_{FINALIZE,IF_PRESENT}. Remove PRAGMA_OACC_CLAUSE_PRESENT_OR_{COPY,COPYIN,COPYOUT,CREATE}. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Add support for finalize and if_present. Make present_or_{copy,copyin,copyout,create} aliases to their non-present_or_* counterparts. Make 'self' an alias to PRAGMA_OACC_CLAUSE_HOST. (c_parser_oacc_data_clause): Update GOMP mappings for PRAGMA_OACC_CLAUSE_{COPY,COPYIN,COPYOUT,CREATE,DELETE}. Remove PRAGMA_OACC_CLAUSE_{SELF,PRESENT_OR_*}. (c_parser_oacc_all_clauses): Handle finalize and if_present clauses. Remove support for present_or_* clauses. (OACC_KERNELS_CLAUSE_MASK): Remove PRESENT_OR_* clauses. (OACC_PARALLEL_CLAUSE_MASK): Likewise. (OACC_DECLARE_CLAUSE_MASK): Likewise. (OACC_DATA_CLAUSE_MASK): Likewise. (OACC_ENTER_DATA_CLAUSE_MASK): Remove PRESENT_OR_* clauses. (OACC_EXIT_DATA_CLAUSE_MASK): Add FINALIZE clause. (OACC_UPDATE_CLAUSE_MASK): Remove SELF, add IF_PRESENT. (c_parser_oacc_declare): Remove PRESENT_OR_* clauses. * c-typeck.c (c_finish_omp_clauses): Handle IF_PRESENT and FINALIZE. gcc/cp/ * parser.c (cp_parser_omp_clause_name): Add support for finalize and if_present. Make present_or_{copy,copyin,copyout,create} aliases to their non-present_or_* counterparts. Make 'self' an alias to PRAGMA_OACC_CLAUSE_HOST. (cp_parser_oacc_data_clause): Update GOMP mappings for PRAGMA_OACC_CLAUSE_{COPY,COPYIN,COPYOUT,CREATE,DELETE}. Remove PRAGMA_OACC_CLAUSE_{SELF,PRESENT_OR_*}. (cp_parser_oacc_all_clauses): Handle finalize and if_present clauses. Remove support for present_or_* clauses. (OACC_KERNELS_CLAUSE_MASK): Remove PRESENT_OR_* clauses. (OACC_PARALLEL_CLAUSE_MASK): Likewise. (OACC_DECLARE_CLAUSE_MASK): Likewise. (OACC_DATA_CLAUSE_MASK): Likewise. (OACC_ENTER_DATA_CLAUSE_MASK): Remove PRESENT_OR_* clauses. (OACC_EXIT_DATA_CLAUSE_MASK): Add FINALIZE clause. (OACC_UPDATE_CLAUSE_MASK): Remove SELF, add IF_PRESENT. (cp_parser_oacc_declare): Remove PRESENT_OR_* clauses. * pt.c (tsubst_omp_clauses): Handle IF_PRESENT and FINALIZE. * semantics.c (finish_omp_clauses): Handle IF_PRESENT and FINALIZE. gcc/fortran/ * gfortran.h (gfc_omp_clauses): Add unsigned if_present, finalize bitfields. * openmp.c (enum omp_mask2): Remove OMP_CLAUSE_PRESENT_OR_*. Add OMP_CLAUSE_{IF_PRESENT,FINALIZE}. (gfc_match_omp_clauses): Update handling of copy, copyin, copyout, create, deviceptr, present_of_*. Add support for finalize and if_present. (OACC_PARALLEL_CLAUSES): Remove PRESENT_OR_* clauses. (OACC_KERNELS_CLAUSES): Likewise. (OACC_DATA_CLAUSES): Likewise. (OACC_DECLARE_CLAUSES): Likewise. (OACC_UPDATE_CLAUSES): Add IF_PRESENT clause. (OACC_ENTER_DATA_CLAUSES): Remove PRESENT_OR_* clauses. (OACC_EXIT_DATA_CLAUSES): Add FINALIZE clause. (gfc_match_oacc_declare): Update to OpenACC 2.5 semantics. * trans-openmp.c (gfc_trans_omp_clauses): Add support for IF_PRESENT and FINALIZE. gcc/ * gimplify.c (gimplify_scan_omp_clauses): Add support for OMP_CLAUSE_{IF_PRESENT,FINALIZE}. (gimplify_adjust_omp_clauses): Likewise. (gimplify_oacc_declare_1): Add support for GOMP_MAP_RELEASE, remove support for GOMP_MAP_FORCE_{ALLOC,TO,FROM,TOFROM}. (gimplify_omp_target_update): Update handling of acc update and enter/exit data. * omp-low.c (install_var_field): Remove unused parameter base_pointers_restrict. (scan_sharing_clauses): Remove base_pointers_restrict parameter. Update call to install_var_field. Handle OMP_CLAUSE_{IF_PRESENT, FINALIZE} (omp_target_base_pointers_restrict_p): Delete. (scan_omp_target): Update call to scan_sharing_clauses. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_{IF_PRESENT, FINALIZE}. * tree-nested.c (convert_nonlocal_omp_clauses): Handle OMP_CLAUSE_{IF_PRESENT,FINALIZE}. (convert_local_omp_clauses): Likewise. * tree-pretty-print.c (dump_omp_clause): Likewise. * tree.c (omp_clause_num_ops): Add entries for OMP_CLAUSE_{IF_PRESENT, FINALIZE}. (omp_clause_code_name): Likewise. gcc/testsuite/ * c-c++-common/goacc/declare-1.c: Update test case to utilize OpenACC 2.5 data clause semantics. * c-c++-common/goacc/declare-2.c: Likewise. * c-c++-common/goacc/default-4.c: Likewise. * c-c++-common/goacc/finalize-1.c: New test. * c-c++-common/goacc/kernels-alias-2.c: Update test case to utilize OpenACC 2.5 data clause semantics. * c-c++-common/goacc/kernels-alias.c: Likewise. * c-c++-common/goacc/routine-5.c: Likewise. * c-c++-common/goacc/update-if_present-1.c: New test. * c-c++-common/goacc/update-if_present-2.c: New test. * g++.dg/goacc/template.C: Update test case to utilize OpenACC 2.5 data clause semantics. * gfortran.dg/goacc/combined-directives.f90: Likewise. * gfortran.dg/goacc/data-tree.f95: Likewise. * gfortran.dg/goacc/declare-2.f95: Likewise. * gfortran.dg/goacc/default-4.f: Likewise. * gfortran.dg/goacc/enter-exit-data.f95: Likewise. * gfortran.dg/goacc/finalize-1.f: New test. * gfortran.dg/goacc/kernels-alias-2.f95: Update test case to utilize OpenACC 2.5 data clause semantics. * gfortran.dg/goacc/kernels-alias.f95: Likewise. * gfortran.dg/goacc/kernels-tree.f95: Likewise. * gfortran.dg/goacc/nested-function-1.f90: Likewise. * gfortran.dg/goacc/parallel-tree.f95: Likewise. * gfortran.dg/goacc/reduction-promotions.f90: Likewise. * gfortran.dg/goacc/update-if_present-1.f90: New test. * gfortran.dg/goacc/update-if_present-2.f90: New test. libgomp/ * libgomp.h (struct splay_tree_key_s): Add dynamic_refcount member. (gomp_acc_remove_pointer): Update declaration. (gomp_acc_declare_allocate): Declare. (gomp_remove_var): Declare. * libgomp.map (OACC_2.5): Define. * oacc-mem.c (acc_map_data): Update refcount. (acc_unmap_data): Likewise. (present_create_copy): Likewise. (acc_create): Add FLAG_PRESENT when calling present_create_copy. (acc_copyin): Likewise. (FLAG_FINALIZE): Define. (delete_copyout): Update dynamic refcounts, add support for FINALIZE. (acc_delete_finalize): New function. (acc_delete_finalize_async): New function. (acc_copyout_finalize): New function. (acc_copyout_finalize_async): New function. (gomp_acc_insert_pointer): Update refcounts. (gomp_acc_remove_pointer): Return if data is not present on the accelerator. * oacc-parallel.c (find_pset): Rename to find_pointer. (find_pointer): Add support for GOMP_MAP_POINTER. (handle_ftn_pointers): New function. (GOACC_parallel_keyed): Update refcounts of variables. (GOACC_enter_exit_data): Add support for finalized data mappings. Add support for GOMP_MAP_{TO,ALLOC,RELESE,FROM}. Update handling of fortran arrays. (GOACC_update): Add support for GOMP_MAP_{ALWAYS_POINTER,TO,FROM}. (GOACC_declare): Add support for GOMP_MAP_RELEASE, remove support for GOMP_MAP_FORCE_FROM. * openacc.f90 (module openacc_internal): Add acc_copyout_finalize_{32_h,64_h,array_h,_l}, and acc_delete_finalize_{32_h,64_h,array_h,_l}. Add interfaces for acc_copyout_finalize and acc_delete_finalize. (acc_copyout_finalize_32_h): New subroutine. (acc_copyout_finalize_64_h): New subroutine. (acc_copyout_finalize_array_h): New subroutine. (acc_delete_finalize_32_h): New subroutine. (acc_delete_finalize_64_h): New subroutine. (acc_delete_finalize_array_h): New subroutine. * openacc.h (acc_copyout_finalize): Declare. (acc_copyout_finalize_async): Declare. (acc_delete_finalize): Declare. (acc_delete_finalize_async): Declare. * openacc_lib.h (acc_copyout_finalize): New interface. (acc_delete_finalize): New interface. * target.c (gomp_map_vars): Update dynamic_refcount. (gomp_remove_var): New function. (gomp_unmap_vars): Use it. (gomp_unload_image_from_device): Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-1.c: Update test case to utilize OpenACC 2.5 data clause semantics. * testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-32.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-83.c: Likewise. * testsuite/libgomp.oacc-fortran/data-5.f90: New test. * testsuite/libgomp.oacc-fortran/data-already-1.f: Update test case to utilize OpenACC 2.5 data clause semantics. * testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise. * testsuite/libgomp.oacc-fortran/lib-32-1.f: Likewise. * testsuite/libgomp.oacc-fortran/lib-32-2.f: Likewise. Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com> Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com> From-SVN: r261813 |
||
Tom de Vries
|
ec00d3faf4 |
[openacc] Move GOMP_OPENACC_DIM parsing out of nvptx plugin
2018-05-02 Tom de Vries <tom@codesourcery.com> PR libgomp/85411 * plugin/plugin-nvptx.c (nvptx_exec): Move parsing of GOMP_OPENACC_DIM ... * env.c (parse_gomp_openacc_dim): ... here. New function. (initialize_env): Call parse_gomp_openacc_dim. (goacc_default_dims): Define. * libgomp.h (goacc_default_dims): Declare. * oacc-plugin.c (GOMP_PLUGIN_acc_default_dim): New function. * oacc-plugin.h (GOMP_PLUGIN_acc_default_dim): Declare. * libgomp.map: New version "GOMP_PLUGIN_1.2". Add GOMP_PLUGIN_acc_default_dim. * testsuite/libgomp.oacc-c-c++-common/loop-default-runtime.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-default.h: New test. From-SVN: r259852 |
||
Jakub Jelinek
|
85ec4feb11 |
Update copyright years.
From-SVN: r256169 |
||
Thomas Schwinge
|
9b94fbc7e4 |
C/C++ OpenACC: acc_pcopyin, acc_pcreate
libgomp/ * openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead of preprocessor definitions. * libgomp.h (strong_alias): Guard by "#ifdef HAVE_ATTRIBUTE_ALIAS". * oacc-mem.c: Provide "acc_pcreate" as alias for "acc_present_or_create", and "acc_pcopyin" as alias for "acc_present_or_copyin". * libgomp.map: New version "OACC_2.0.1". (OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate". * testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging its content into... * testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file. Extend testing. From-SVN: r248410 |
||
Thomas Schwinge
|
345a8c1712 |
libgomp: Normalize the names of a few functions of the libgomp plugin API
libgomp/ * libgomp-plugin.h (GOMP_OFFLOAD_openacc_parallel): Rename to GOMP_OFFLOAD_openacc_exec. Adjust all users. (GOMP_OFFLOAD_openacc_get_current_cuda_device): Rename to GOMP_OFFLOAD_openacc_cuda_get_current_device. Adjust all users. (GOMP_OFFLOAD_openacc_get_current_cuda_context): Rename to GOMP_OFFLOAD_openacc_cuda_get_current_context. Adjust all users. (GOMP_OFFLOAD_openacc_get_cuda_stream): Rename to GOMP_OFFLOAD_openacc_cuda_get_stream. Adjust all users. (GOMP_OFFLOAD_openacc_set_cuda_stream): Rename to GOMP_OFFLOAD_openacc_cuda_set_stream. Adjust all users. From-SVN: r245125 |
||
Thomas Schwinge
|
dced339c8a |
libgomp: Provide prototypes for functions implemented by libgomp plugins
libgomp/ * libgomp-plugin.h: #include <stdbool.h>. (GOMP_OFFLOAD_get_name, GOMP_OFFLOAD_get_caps) (GOMP_OFFLOAD_get_type, GOMP_OFFLOAD_get_num_devices) (GOMP_OFFLOAD_init_device, GOMP_OFFLOAD_fini_device) (GOMP_OFFLOAD_version, GOMP_OFFLOAD_load_image) (GOMP_OFFLOAD_unload_image, GOMP_OFFLOAD_alloc, GOMP_OFFLOAD_free) (GOMP_OFFLOAD_dev2host, GOMP_OFFLOAD_host2dev) (GOMP_OFFLOAD_dev2dev, GOMP_OFFLOAD_can_run, GOMP_OFFLOAD_run) (GOMP_OFFLOAD_async_run, GOMP_OFFLOAD_openacc_parallel) (GOMP_OFFLOAD_openacc_register_async_cleanup) (GOMP_OFFLOAD_openacc_async_test) (GOMP_OFFLOAD_openacc_async_test_all) (GOMP_OFFLOAD_openacc_async_wait) (GOMP_OFFLOAD_openacc_async_wait_async) (GOMP_OFFLOAD_openacc_async_wait_all) (GOMP_OFFLOAD_openacc_async_wait_all_async) (GOMP_OFFLOAD_openacc_async_set_async) (GOMP_OFFLOAD_openacc_create_thread_data) (GOMP_OFFLOAD_openacc_destroy_thread_data) (GOMP_OFFLOAD_openacc_get_current_cuda_device) (GOMP_OFFLOAD_openacc_get_current_cuda_context) (GOMP_OFFLOAD_openacc_get_cuda_stream) (GOMP_OFFLOAD_openacc_set_cuda_stream): New prototypes. * libgomp.h (struct acc_dispatch_t, struct gomp_device_descr): Use these. * plugin/plugin-hsa.c (GOMP_OFFLOAD_load_image) (GOMP_OFFLOAD_unload_image): Fix argument types. liboffloadmic/ * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_type): Fix return type. (GOMP_OFFLOAD_load_image): Fix argument types. From-SVN: r245062 |
||
Jakub Jelinek
|
cbe34bb5ed |
Update copyright years.
From-SVN: r243994 |
||
Alexander Monakov
|
6103184e81 |
OpenMP offloading to NVPTX: libgomp changes
* Makefile.am (libgomp_la_SOURCES): Add atomic.c, icv.c, icv-device.c. * Makefile.in. Regenerate. * configure.ac [nvptx*-*-*] (libgomp_use_pthreads): Set and use it... (LIBGOMP_USE_PTHREADS): ...here; new define. * configure: Regenerate. * config.h.in: Likewise. * config/posix/affinity.c: Move to... * affinity.c: ...here (new file). Guard use of Pthreads-specific interface by LIBGOMP_USE_PTHREADS. * critical.c: Split out GOMP_atomic_{start,end} into... * atomic.c: ...here (new file). * env.c: Split out ICV definitions into... * icv.c: ...here (new file) and... * icv-device.c: ...here. New file. * config/linux/lock.c (gomp_init_lock_30): Move to generic lock.c. (gomp_destroy_lock_30): Ditto. (gomp_set_lock_30): Ditto. (gomp_unset_lock_30): Ditto. (gomp_test_lock_30): Ditto. (gomp_init_nest_lock_30): Ditto. (gomp_destroy_nest_lock_30): Ditto. (gomp_set_nest_lock_30): Ditto. (gomp_unset_nest_lock_30): Ditto. (gomp_test_nest_lock_30): Ditto. * lock.c: New. * config/nvptx/lock.c: New. * config/nvptx/bar.c: New. * config/nvptx/bar.h: New. * config/nvptx/doacross.h: New. * config/nvptx/error.c: New. * config/nvptx/icv-device.c: New. * config/nvptx/mutex.h: New. * config/nvptx/pool.h: New. * config/nvptx/proc.c: New. * config/nvptx/ptrlock.h: New. * config/nvptx/sem.h: New. * config/nvptx/simple-bar.h: New. * config/nvptx/target.c: New. * config/nvptx/task.c: New. * config/nvptx/team.c: New. * config/nvptx/time.c: New. * config/posix/simple-bar.h: New. * libgomp.h: Guard pthread.h inclusion. Include simple-bar.h. (gomp_num_teams_var): Declare. (struct gomp_thread_pool): Change threads_dock member to gomp_simple_barrier_t. [__nvptx__] (gomp_thread): New implementation. (gomp_thread_attr): Guard by LIBGOMP_USE_PTHREADS. (gomp_thread_destructor): Ditto. (gomp_init_thread_affinity): Ditto. * team.c: Guard uses of Pthreads-specific interfaces by LIBGOMP_USE_PTHREADS. Adjust all uses of threads_dock. (gomp_free_thread) [__nvptx__]: Do not call 'free'. * config/nvptx/alloc.c: Delete. * config/nvptx/barrier.c: Ditto. * config/nvptx/fortran.c: Ditto. * config/nvptx/iter.c: Ditto. * config/nvptx/iter_ull.c: Ditto. * config/nvptx/loop.c: Ditto. * config/nvptx/loop_ull.c: Ditto. * config/nvptx/ordered.c: Ditto. * config/nvptx/parallel.c: Ditto. * config/nvptx/priority_queue.c: Ditto. * config/nvptx/sections.c: Ditto. * config/nvptx/single.c: Ditto. * config/nvptx/splay-tree.c: Ditto. * config/nvptx/work.c: Ditto. * testsuite/libgomp.fortran/fortran.exp (lang_link_flags): Pass -foffload=-lgfortran in addition to -lgfortran. * testsuite/libgomp.oacc-fortran/fortran.exp (lang_link_flags): Ditto. * plugin/plugin-nvptx.c: Include <limits.h>. (struct targ_fn_descriptor): Add new fields. (struct ptx_device): Ditto. Set them... (nvptx_open_device): ...here. (nvptx_adjust_launch_bounds): New. (nvptx_host2dev): Allow NULL 'nvthd'. (nvptx_dev2host): Ditto. (GOMP_OFFLOAD_get_caps): Add GOMP_OFFLOAD_CAP_OPENMP_400. (link_ptx): Adjust log sizes. (nvptx_host2dev): Allow NULL 'nvthd'. (nvptx_dev2host): Ditto. (nvptx_set_clocktick): New. Use it... (GOMP_OFFLOAD_load_image): ...here. Set new targ_fn_descriptor fields. (GOMP_OFFLOAD_dev2dev): New. (nvptx_adjust_launch_bounds): New. (nvptx_stacks_size): New. (nvptx_stacks_alloc): New. (nvptx_stacks_free): New. (GOMP_OFFLOAD_run): New. (GOMP_OFFLOAD_async_run): New (stub). Co-Authored-By: Dmitry Melnik <dm@ispras.ru> Co-Authored-By: Jakub Jelinek <jakub@redhat.com> From-SVN: r242789 |
||
Chung-Lin Tang
|
b4557008c4 |
oacc-plugin.h (GOMP_PLUGIN_async_unmap_vars): Add int parameter.
2016-05-26 Chung-Lin Tang <cltang@codesourcery.com> libgomp/ * oacc-plugin.h (GOMP_PLUGIN_async_unmap_vars): Add int parameter. * oacc-plugin.c (GOMP_PLUGIN_async_unmap_vars): Add 'int async' parameter, use to set async stream around call to gomp_unmap_vars, call gomp_unmap_vars() with 'do_copyfrom' set to true. * plugin/plugin-nvptx.c (struct ptx_event): Add 'int val' field. (event_gc): Adjust event handling loop, collect PTX_EVT_ASYNC_CLEANUP events and call GOMP_PLUGIN_async_unmap_vars() for each of them. (event_add): Add int parameter, initialize 'val' field when adding new ptx_event struct. (nvptx_evec): Adjust event_add() call arguments. (nvptx_host2dev): Likewise. (nvptx_dev2host): Likewise. (nvptx_wait_async): Likewise. (nvptx_wait_all_async): Likewise. (GOMP_OFFLOAD_openacc_register_async_cleanup): Add async parameter, pass to event_add() call. * oacc-host.c (host_openacc_register_async_cleanup): Add 'int async' parameter. * oacc-mem.c (gomp_acc_remove_pointer): Adjust async case to call openacc.register_async_cleanup_func() hook. * oacc-parallel.c (GOACC_parallel_keyed): Likewise. * target.c (gomp_copy_from_async): Delete function. (gomp_map_vars): Remove async_refcount. (gomp_unmap_vars): Likewise. (gomp_load_image_to_device): Likewise. (omp_target_associate_ptr): Likewise. * libgomp.h (struct splay_tree_key_s): Remove async_refcount. (acc_dispatch_t.register_async_cleanup_func): Add int parameter. (gomp_copy_from_async): Remove. From-SVN: r236772 |
||
Chung-Lin Tang
|
6ce1307231 |
target.c (gomp_device_copy): New function.
libgomp/ 2016-05-26 Chung-Lin Tang <cltang@codesourcery.com> * target.c (gomp_device_copy): New function. (gomp_copy_host2dev): Likewise. (gomp_copy_dev2host): Likewise. (gomp_free_device_memory): Likewise. (gomp_map_vars_existing): Adjust to call gomp_copy_host2dev. (gomp_map_pointer): Likewise. (gomp_map_vars): Adjust to call gomp_copy_host2dev, handle NULL value from alloc_func plugin hook. (gomp_unmap_tgt): Adjust to call gomp_free_device_memory. (gomp_copy_from_async): Adjust to call gomp_copy_dev2host. (gomp_unmap_vars): Likewise. (gomp_update): Adjust to call gomp_copy_dev2host and gomp_copy_host2dev functions. (gomp_unload_image_from_device): Handle false value from unload_image_func plugin hook. (gomp_init_device): Handle false value from init_device_func plugin hook. (gomp_exit_data): Adjust to call gomp_copy_dev2host. (omp_target_free): Adjust to call gomp_free_device_memory. (omp_target_memcpy): Handle return values from host2dev_func, dev2host_func, and dev2dev_func plugin hooks. (omp_target_memcpy_rect_worker): Likewise. (gomp_target_fini): Handle false value from fini_device_func plugin hook. * libgomp.h (struct gomp_device_descr): Adjust return type of init_device_func, fini_device_func, unload_image_func, free_func, dev2host_func,host2dev_func, and dev2dev_func plugin hooks to 'bool'. * oacc-init.c (acc_shutdown_1): Handle false value from fini_device_func plugin hook. * oacc-host.c (host_init_device): Change return type to bool. (host_fini_device): Likewise. (host_unload_image): Likewise. (host_free): Likewise. (host_dev2host): Likewise. (host_host2dev): Likewise. * oacc-mem.c (acc_free): Handle plugin hook fatal error case. (acc_memcpy_to_device): Likewise. (acc_memcpy_from_device): Likewise. (delete_copyout): Add libfnname parameter, handle free_func hook fatal error case. (acc_delete): Adjust delete_copyout call. (acc_copyout): Likewise. (update_dev_host): Move gomp_mutex_unlock to after host2dev/dev2host hook calls. * plugin/plugin-hsa.c (hsa_warn): Adjust 'hsa_error' local variable to 'hsa_error_msg', for clarity. (hsa_fatal): Likewise. (hsa_error): New function. (init_hsa_context): Change return type to bool, adjust to return false on error. (GOMP_OFFLOAD_get_num_devices): Adjust to handle init_hsa_context return value. (GOMP_OFFLOAD_init_device): Change return type to bool, adjust to return false on error. (get_agent_info): Adjust to return NULL on error. (destroy_hsa_program): Change return type to bool, adjust to return false on error. (GOMP_OFFLOAD_load_image): Adjust to return -1 on error. (destroy_module): Change return type to bool, adjust to return false on error. (GOMP_OFFLOAD_unload_image): Likewise. (GOMP_OFFLOAD_fini_device): Likewise. (GOMP_OFFLOAD_alloc): Change to return NULL when called. (GOMP_OFFLOAD_free): Change to return false when called. (GOMP_OFFLOAD_dev2host): Likewise. (GOMP_OFFLOAD_host2dev): Likewise. (GOMP_OFFLOAD_dev2dev): Likewise. * plugin/plugin-nvptx.c (CUDA_CALL_ERET): New convenience macro. (CUDA_CALL): Likewise. (CUDA_CALL_ASSERT): Likewise. (map_init): Change return type to bool, use CUDA_CALL* macros. (map_fini): Likewise. (init_streams_for_device): Change return type to bool, adjust call to map_init. (fini_streams_for_device): Change return type to bool, adjust call to map_fini. (select_stream_for_async): Release stream_lock before calls to GOMP_PLUGIN_fatal, adjust call to map_init. (nvptx_init): Use CUDA_CALL* macros. (nvptx_attach_host_thread_to_device): Change return type to bool, use CUDA_CALL* macros. (nvptx_open_device): Use CUDA_CALL* macros. (nvptx_close_device): Change return type to bool, use CUDA_CALL* macros. (nvptx_get_num_devices): Use CUDA_CALL* macros. (link_ptx): Change return type to bool, use CUDA_CALL* macros. (nvptx_exec): Use CUDA_CALL* macros. (nvptx_alloc): Use CUDA_CALL* macros. (nvptx_free): Change return type to bool, use CUDA_CALL* macros. (nvptx_host2dev): Likewise. (nvptx_dev2host): Likewise. (nvptx_wait): Use CUDA_CALL* macros. (nvptx_wait_async): Likewise. (nvptx_wait_all): Likewise. (nvptx_wait_all_async): Likewise. (nvptx_set_cuda_stream): Adjust order of stream_lock acquire, use CUDA_CALL* macros, adjust call to map_fini. (GOMP_OFFLOAD_init_device): Change return type to bool, adjust code accordingly. (GOMP_OFFLOAD_fini_device): Likewise. (GOMP_OFFLOAD_load_image): Adjust calls to nvptx_attach_host_thread_to_device/link_ptx to handle errors, use CUDA_CALL* macros. (GOMP_OFFLOAD_unload_image): Change return type to bool, adjust return code. (GOMP_OFFLOAD_alloc): Adjust calls to code to handle error return. (GOMP_OFFLOAD_free): Change return type to bool, adjust calls to handle error return. (GOMP_OFFLOAD_dev2host): Likewise. (GOMP_OFFLOAD_host2dev): Likewise. (GOMP_OFFLOAD_openacc_register_async_cleanup): Use CUDA_CALL* macros. (GOMP_OFFLOAD_openacc_create_thread_data): Likewise. liboffloadmic/ 2016-05-26 Chung-Lin Tang <cltang@codesourcery.com> * plugin/libgomp-plugin-intelmic.cpp (offload): Change return type to bool, adjust return code. (GOMP_OFFLOAD_init_device): Likewise. (GOMP_OFFLOAD_fini_device): Likewise. (get_target_table): Likewise. (offload_image): Likwise. (GOMP_OFFLOAD_load_image): Adjust call to offload_image(), change to return -1 on error. (GOMP_OFFLOAD_unload_image): Change return type to bool, adjust return code. (GOMP_OFFLOAD_alloc): Likewise. (GOMP_OFFLOAD_free): Likewise. (GOMP_OFFLOAD_host2dev): Likewise. (GOMP_OFFLOAD_dev2host): Likewise. (GOMP_OFFLOAD_dev2dev): Likewise. From-SVN: r236768 |
||
Jakub Jelinek
|
21f3a2369b |
omp-low.c (lower_omp_target): Use GOMP_MAP_FIRSTPRIVATE_INT regardless whether there are depend clauses or not.
* omp-low.c (lower_omp_target): Use GOMP_MAP_FIRSTPRIVATE_INT regardless whether there are depend clauses or not. * libgomp.h (struct gomp_target_task): Remove firstprivate_copies field. * target.c (gomp_target_fallback_firstprivate, gomp_target_unshare_firstprivate): Removed. (GOMP_target_ext): Copy firstprivate vars into gomp_allocaed memory before waiting for dependencies. (gomp_target_task_fn): Don't copy firstprivate vars here. * task.c (GOMP_PLUGIN_target_task_completion): Don't free firstprivate_copies here. (gomp_create_target_task): Don't initialize firstprivate_copies field. * testsuite/libgomp.c/target-25.c (main): Use map (to:) instead of explicit/implicit firstprivate. From-SVN: r234894 |
||
Martin Jambor
|
b2b4005150 |
Merge of HSA
2016-01-19 Martin Jambor <mjambor@suse.cz> Martin Liska <mliska@suse.cz> Michael Matz <matz@suse.de> libgomp/ * plugin/Makefrag.am: Add HSA plugin requirements. * plugin/configfrag.ac (HSA_RUNTIME_INCLUDE): New variable. (HSA_RUNTIME_LIB): Likewise. (HSA_RUNTIME_CPPFLAGS): Likewise. (HSA_RUNTIME_INCLUDE): New substitution. (HSA_RUNTIME_LIB): Likewise. (HSA_RUNTIME_LDFLAGS): Likewise. (hsa-runtime): New configure option. (hsa-runtime-include): Likewise. (hsa-runtime-lib): Likewise. (PLUGIN_HSA): New substitution variable. Fill HSA_RUNTIME_INCLUDE and HSA_RUNTIME_LIB according to the new configure options. (PLUGIN_HSA_CPPFLAGS): Likewise. (PLUGIN_HSA_LDFLAGS): Likewise. (PLUGIN_HSA_LIBS): Likewise. Check that we have access to HSA run-time. * libgomp-plugin.h (offload_target_type): New element OFFLOAD_TARGET_TYPE_HSA. * libgomp.h (gomp_target_task): New fields firstprivate_copies and args. (bool gomp_create_target_task): Updated. (gomp_device_descr): Extra parameter of run_func and async_run_func, new field can_run_func. * libgomp_g.h (GOMP_target_ext): Update prototype. * oacc-host.c (host_run): Added a new parameter args. * target.c (calculate_firstprivate_requirements): New function. (copy_firstprivate_data): Likewise. (gomp_target_fallback_firstprivate): Use them. (gomp_target_unshare_firstprivate): New function. (gomp_get_target_fn_addr): Allow returning NULL for shared memory devices. (GOMP_target): Do host fallback for all shared memory devices. Do not pass any args to plugins. (GOMP_target_ext): Introduce device-specific argument parameter args. Allow host fallback if device shares memory. Do not remap data if device has shared memory. (gomp_target_task_fn): Likewise. Also treat shared memory devices like host fallback for mappings. (GOMP_target_data): Treat shared memory devices like host fallback. (GOMP_target_data_ext): Likewise. (GOMP_target_update): Likewise. (GOMP_target_update_ext): Likewise. Also pass NULL as args to gomp_create_target_task. (GOMP_target_enter_exit_data): Likewise. (omp_target_alloc): Treat shared memory devices like host fallback. (omp_target_free): Likewise. (omp_target_is_present): Likewise. (omp_target_memcpy): Likewise. (omp_target_memcpy_rect): Likewise. (omp_target_associate_ptr): Likewise. (gomp_load_plugin_for_device): Also load can_run. * task.c (GOMP_PLUGIN_target_task_completion): Free firstprivate_copies. (gomp_create_target_task): Accept new argument args and store it to ttask. * plugin/plugin-hsa.c: New file. gcc/ * Makefile.in (OBJS): Add new source files. (GTFILES): Add hsa.c. * common.opt (disable_hsa): New variable. (-Whsa): New warning. * config.in (ENABLE_HSA): New. * configure.ac: Treat hsa differently from other accelerators. (OFFLOAD_TARGETS): Define ENABLE_OFFLOADING according to $enable_offloading. (ENABLE_HSA): Define ENABLE_HSA according to $enable_hsa. * doc/install.texi (Configuration): Document --with-hsa-runtime, --with-hsa-runtime-include, --with-hsa-runtime-lib and --with-hsa-kmt-lib. * doc/invoke.texi (-Whsa): Document. (hsa-gen-debug-stores): Likewise. * lto-wrapper.c (compile_images_for_offload_targets): Do not attempt to invoke offload compiler for hsa acclerator. * opts.c (common_handle_option): Determine whether HSA offloading should be performed. * params.def (PARAM_HSA_GEN_DEBUG_STORES): New parameter. * builtin-types.def (BT_FN_VOID_UINT_PTR_INT_PTR): New. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT): Removed. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR): New. * gimple-low.c (lower_stmt): Also handle GIMPLE_OMP_GRID_BODY. * gimple-pretty-print.c (dump_gimple_omp_for): Also handle GF_OMP_FOR_KIND_GRID_LOOP. (dump_gimple_omp_block): Also handle GIMPLE_OMP_GRID_BODY. (pp_gimple_stmt_1): Likewise. * gimple-walk.c (walk_gimple_stmt): Likewise. * gimple.c (gimple_build_omp_grid_body): New function. (gimple_copy): Also handle GIMPLE_OMP_GRID_BODY. * gimple.def (GIMPLE_OMP_GRID_BODY): New. * gimple.h (enum gf_mask): Added GF_OMP_PARALLEL_GRID_PHONY, GF_OMP_FOR_KIND_GRID_LOOP, GF_OMP_FOR_GRID_PHONY and GF_OMP_TEAMS_GRID_PHONY. (gimple_statement_omp_single_layout): Updated comments. (gimple_build_omp_grid_body): New function. (gimple_has_substatements): Also handle GIMPLE_OMP_GRID_BODY. (gimple_omp_for_grid_phony): New function. (gimple_omp_for_set_grid_phony): Likewise. (gimple_omp_parallel_grid_phony): Likewise. (gimple_omp_parallel_set_grid_phony): Likewise. (gimple_omp_teams_grid_phony): Likewise. (gimple_omp_teams_set_grid_phony): Likewise. (gimple_return_set_retbnd): Also handle GIMPLE_OMP_GRID_BODY. * omp-builtins.def (BUILT_IN_GOMP_OFFLOAD_REGISTER): New. (BUILT_IN_GOMP_OFFLOAD_UNREGISTER): Likewise. (BUILT_IN_GOMP_TARGET): Updated type. * omp-low.c: Include symbol-summary.h, hsa.h and params.h. (adjust_for_condition): New function. (get_omp_for_step_from_incr): Likewise. (extract_omp_for_data): Moved parts to adjust_for_condition and get_omp_for_step_from_incr. (build_outer_var_ref): Handle GIMPLE_OMP_GRID_BODY. (fixup_child_record_type): Bail out if receiver_decl is NULL. (scan_sharing_clauses): Handle OMP_CLAUSE__GRIDDIM_. (scan_omp_parallel): Do not create child functions for phony constructs. (check_omp_nesting_restrictions): Handle GIMPLE_OMP_GRID_BODY. (scan_omp_1_op): Checking assert we are not remapping to ERROR_MARK. Also also handle GIMPLE_OMP_GRID_BODY. (parallel_needs_hsa_kernel_p): New function. (expand_parallel_call): Register apprpriate parallel child functions as HSA kernels. (grid_launch_attributes_trees): New type. (grid_attr_trees): New variable. (grid_create_kernel_launch_attr_types): New function. (grid_insert_store_range_dim): Likewise. (grid_get_kernel_launch_attributes): Likewise. (get_target_argument_identifier_1): Likewise. (get_target_argument_identifier): Likewise. (get_target_argument_value): Likewise. (push_target_argument_according_to_value): Likewise. (get_target_arguments): Likewise. (expand_omp_target): Call get_target_arguments instead of looking up for teams and thread limit. (grid_expand_omp_for_loop): New function. (grid_arg_decl_map): New type. (grid_remap_kernel_arg_accesses): New function. (grid_expand_target_kernel_body): New function. (expand_omp): Call it. (lower_omp_for): Do not emit phony constructs. (lower_omp_taskreg): Do not emit phony constructs but create for them a temporary variable receiver_decl. (lower_omp_taskreg): Do not emit phony constructs. (lower_omp_teams): Likewise. (lower_omp_grid_body): New function. (lower_omp_1): Call it. (grid_reg_assignment_to_local_var_p): New function. (grid_seq_only_contains_local_assignments): Likewise. (grid_find_single_omp_among_assignments_1): Likewise. (grid_find_single_omp_among_assignments): Likewise. (grid_find_ungridifiable_statement): Likewise. (grid_target_follows_gridifiable_pattern): Likewise. (grid_remap_prebody_decls): Likewise. (grid_copy_leading_local_assignments): Likewise. (grid_process_kernel_body_copy): Likewise. (grid_attempt_target_gridification): Likewise. (grid_gridify_all_targets_stmt): Likewise. (grid_gridify_all_targets): Likewise. (execute_lower_omp): Call grid_gridify_all_targets. (make_gimple_omp_edges): Handle GIMPLE_OMP_GRID_BODY. * tree-core.h (omp_clause_code): Added OMP_CLAUSE__GRIDDIM_. (tree_omp_clause): Added union field dimension. * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE__GRIDDIM_. * tree.c (omp_clause_num_ops): Added number of arguments of OMP_CLAUSE__GRIDDIM_. (omp_clause_code_name): Added name of OMP_CLAUSE__GRIDDIM_. (walk_tree_1): Handle OMP_CLAUSE__GRIDDIM_. * tree.h (OMP_CLAUSE_GRIDDIM_DIMENSION): New. (OMP_CLAUSE_SET_GRIDDIM_DIMENSION): Likewise. (OMP_CLAUSE_GRIDDIM_SIZE): Likewise. (OMP_CLAUSE_GRIDDIM_GROUP): Likewise. * passes.def: Schedule pass_ipa_hsa and pass_gen_hsail. * tree-pass.h (make_pass_gen_hsail): Declare. (make_pass_ipa_hsa): Likewise. * ipa-hsa.c: New file. * lto-section-in.c (lto_section_name): Add hsa section name. * lto-streamer.h (lto_section_type): Add hsa section. * timevar.def (TV_IPA_HSA): New. * hsa-brig-format.h: New file. * hsa-brig.c: New file. * hsa-dump.c: Likewise. * hsa-gen.c: Likewise. * hsa.c: Likewise. * hsa.h: Likewise. * toplev.c (compile_file): Call hsa_output_brig. * hsa-regalloc.c: New file. gcc/fortran/ * types.def (BT_FN_VOID_UINT_PTR_INT_PTR): New. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT): Removed. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR): New. gcc/lto/ * lto-partition.c: Include "hsa.h" (add_symbol_to_partition_1): Put hsa implementations into the same partition as host implementations. liboffloadmic/ * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_async_run): New unused parameter. (GOMP_OFFLOAD_run): Likewise. include/ * gomp-constants.h (GOMP_DEVICE_HSA): New macro. (GOMP_VERSION_HSA): Likewise. (GOMP_TARGET_ARG_DEVICE_MASK): Likewise. (GOMP_TARGET_ARG_DEVICE_ALL): Likewise. (GOMP_TARGET_ARG_SUBSEQUENT_PARAM): Likewise. (GOMP_TARGET_ARG_ID_MASK): Likewise. (GOMP_TARGET_ARG_NUM_TEAMS): Likewise. (GOMP_TARGET_ARG_THREAD_LIMIT): Likewise. (GOMP_TARGET_ARG_VALUE_SHIFT): Likewise. (GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES): Likewise. From-SVN: r232549 |
||
Jakub Jelinek
|
818ab71a41 |
Update copyright years.
From-SVN: r232055 |
||
Ilya Verbin
|
4a38b02b4e |
c-common.c (c_common_attribute_table): Handle "omp declare target link" attribute.
gcc/c-family/ * c-common.c (c_common_attribute_table): Handle "omp declare target link" attribute. gcc/ * cgraphunit.c (output_in_order): Do not assemble "omp declare target link" variables in ACCEL_COMPILER. * gimplify.c (gimplify_adjust_omp_clauses): Do not remove mapping of "omp declare target link" variables. * omp-low.c (scan_sharing_clauses): Do not remove mapping of "omp declare target link" variables. (add_decls_addresses_to_decl_constructor): For "omp declare target link" variables output address of the artificial pointer instead of address of the variable. Set most significant bit of the size to mark them. (pass_data_omp_target_link): New pass_data. (pass_omp_target_link): New class. (find_link_var_op): New static function. (make_pass_omp_target_link): New function. * passes.def: Add pass_omp_target_link. * tree-pass.h (make_pass_omp_target_link): Declare. * varpool.c (symbol_table::output_variables): Do not assemble "omp declare target link" variables in ACCEL_COMPILER. gcc/lto/ * lto.c: Include stringpool.h and fold-const.h. (offload_handle_link_vars): New static function. (lto_main): Call offload_handle_link_vars. libgomp/ * libgomp.h (REFCOUNT_LINK): Define. (struct splay_tree_key_s): Add link_key. * target.c (gomp_map_vars): Treat REFCOUNT_LINK objects as not mapped. Replace target address of the pointer with target address of newly mapped object in the splay tree. Set link pointer on target to the device address of the mapped object. (gomp_unmap_vars): Restore target address of the pointer in the splay tree for REFCOUNT_LINK objects after unmapping. (gomp_load_image_to_device): Set refcount to REFCOUNT_LINK for "omp declare target link" objects. (gomp_unload_image_from_device): Replace j with i. Force unmap of all "omp declare target link" objects, which were mapped for the image. (gomp_exit_data): Restore target address of the pointer in the splay tree for REFCOUNT_LINK objects after unmapping. * testsuite/libgomp.c/target-link-1.c: New file. From-SVN: r231655 |
||
Ilya Verbin
|
d84ffc0a56 |
libgomp.h (gomp_device_state): New enum.
libgomp/ * libgomp.h (gomp_device_state): New enum. (struct gomp_device_descr): Replace is_initialized with state. (gomp_fini_device): Remove declaration. * oacc-host.c (host_dispatch): Use state instead of is_initialized. * oacc-init.c (acc_init_1): Use state instead of is_initialized. (acc_shutdown_1): Likewise. Inline gomp_fini_device. (acc_set_device_type): Use state instead of is_initialized. (acc_set_device_num): Likewise. * target.c (resolve_device): Use state instead of is_initialized. Do not initialize finalized device. (gomp_map_vars): Do nothing if device is finalized. (gomp_unmap_vars): Likewise. (gomp_update): Likewise. (GOMP_offload_register_ver): Use state instead of is_initialized. (GOMP_offload_unregister_ver): Likewise. (gomp_init_device): Likewise. (gomp_unload_device): Likewise. (gomp_fini_device): Remove. (gomp_get_target_fn_addr): Do nothing if device is finalized. (GOMP_target): Go to host fallback if device is finalized. (GOMP_target_ext): Likewise. (gomp_exit_data): Do nothing if device is finalized. (gomp_target_task_fn): Go to host fallback if device is finalized. (gomp_target_fini): New static function. (gomp_target_init): Use state instead of is_initialized. Call gomp_target_fini at exit. liboffloadmic/ * plugin/libgomp-plugin-intelmic.cpp (unregister_main_image): Remove. (register_main_image): Do not call unregister_main_image at exit. (GOMP_OFFLOAD_fini_device): Allow for OpenMP. Unregister main image. From-SVN: r231623 |
||
Jakub Jelinek
|
e460634820 |
omp-low.c (lower_omp_ordered): Add argument to GOMP_SMD_ORDERED_* internal calls - 0 if...
gcc/ 2015-11-14 Jakub Jelinek <jakub@redhat.com> * omp-low.c (lower_omp_ordered): Add argument to GOMP_SMD_ORDERED_* internal calls - 0 if ordered simd and 1 for ordered threads simd. * tree-vectorizer.c (adjust_simduid_builtins): If GOMP_SIMD_ORDERED_* argument is 1, replace it with GOMP_ordered_* call instead of removing it. gcc/c/ 2015-11-14 Jakub Jelinek <jakub@redhat.com> * c-typeck.c (c_finish_omp_clauses): Don't mark GOMP_MAP_FIRSTPRIVATE_POINTER decls addressable. gcc/cp/ 2015-11-14 Jakub Jelinek <jakub@redhat.com> * semantics.c (finish_omp_clauses): Don't mark GOMP_MAP_FIRSTPRIVATE_POINTER decls addressable. libgomp/ 2015-11-14 Jakub Jelinek <jakub@redhat.com> Aldy Hernandez <aldyh@redhat.com> Ilya Verbin <ilya.verbin@intel.com> * ordered.c (gomp_doacross_init, GOMP_doacross_post, GOMP_doacross_wait, gomp_doacross_ull_init, GOMP_doacross_ull_post, GOMP_doacross_ull_wait): For GFS_GUIDED don't divide number of iterators or IV by chunk size. * parallel.c (gomp_resolve_num_threads): Don't assume that if thr->ts.team is non-NULL, then pool must be non-NULL. * libgomp-plugin.h (GOMP_PLUGIN_target_task_completion): Declare. * libgomp.map (GOMP_PLUGIN_1.1): New symbol version, export GOMP_PLUGIN_target_task_completion. * Makefile.am (libgomp_la_SOURCES): Add priority_queue.c. * Makefile.in: Regenerate. * libgomp.h: Shuffle prototypes and forward definitions around so priority queues can be defined. (enum gomp_task_kind): Add GOMP_TASK_ASYNC_RUNNING. (enum gomp_target_task_state): New enum. (struct gomp_target_task): Add state, tgt, task and team fields. (gomp_create_target_task): Change return type to bool, add state argument. (gomp_target_task_fn): Change return type to bool. (struct gomp_device_descr): Add async_run_func. (struct gomp_task): Remove children, next_child, prev_child, next_queue, prev_queue, next_taskgroup, prev_taskgroup. Add pnode field. (struct gomp_taskgroup): Remove children. Add taskgroup_queue. (struct gomp_team): Change task_queue type to a priority queue. (splay_compare): Define inline. (priority_queue_offset): New. (priority_node_to_task): New. (task_to_priority_node): New. * oacc-mem.c: Do not include splay-tree.h. * priority_queue.c: New file. * priority_queue.h: New file. * splay-tree.c: Do not include splay-tree.h. (splay_tree_foreach_internal): New. (splay_tree_foreach): New. * splay-tree.h: Become re-entrant if splay_tree_prefix is defined. (splay_tree_callback): Define typedef. * target.c (splay_compare): Move to libgomp.h. (GOMP_target): Don't adjust *thr in any way around running offloaded task. (GOMP_target_ext): Likewise. Handle target nowait. (GOMP_target_update_ext, GOMP_target_enter_exit_data): Check return value from gomp_create_target_task, if false, fallthrough as if no dependencies exist. (gomp_target_task_fn): Change return type to bool, return true if the task should have another part scheduled later. Handle target nowait. (gomp_load_plugin_for_device): Initialize async_run. * task.c (gomp_init_task): Initialize children_queue. (gomp_clear_parent_in_list): New. (gomp_clear_parent_in_tree): New. (gomp_clear_parent): Handle priorities. (GOMP_task): Likewise. (priority_queue_move_task_first, gomp_target_task_completion, GOMP_PLUGIN_target_task_completion): New functions. (gomp_create_target_task): Use priority queues. Change return type to bool, add state argument, return false if for async {{enter,exit} data,update} constructs no dependencies need to be waited for, handle target nowait. Set task->fn to NULL instead of gomp_target_task_fn. (verify_children_queue): Remove. (priority_list_upgrade_task): New. (priority_queue_upgrade_task): New. (verify_task_queue): Remove. (priority_list_downgrade_task): New. (priority_queue_downgrade_task): New. (gomp_task_run_pre): Use priority queues. Abstract code out to priority_queue_downgrade_task. (gomp_task_run_post_handle_dependers): Use priority queues. (gomp_task_run_post_remove_parent): Likewise. (gomp_task_run_post_remove_taskgroup): Likewise. (gomp_barrier_handle_tasks): Likewise. Handle target nowait target tasks specially. (GOMP_taskwait): Likewise. (gomp_task_maybe_wait_for_dependencies): Likewise. Abstract code to priority-queue_upgrade_task. (GOMP_taskgroup_start): Use priority queues. (GOMP_taskgroup_end): Likewise. Handle target nowait target tasks specially. If taskgroup is NULL, and thr->ts.level is 0, act as a barrier. * taskloop.c (GOMP_taskloop): Handle priorities. * team.c (gomp_new_team): Call priority_queue_init. (free_team): Call priority_queue_free. (gomp_free_thread): Call gomp_team_end if thr->ts.team is artificial team created for target nowait in implicit parallel region. (gomp_team_start): For nested check, test thr->ts.level instead of thr->ts.team != NULL. * testsuite/libgomp.c/doacross-3.c: New test. * testsuite/libgomp.c/ordered-5.c: New test. * testsuite/libgomp.c/priority.c: New test. * testsuite/libgomp.c/target-31.c: New test. * testsuite/libgomp.c/target-32.c: New test. * testsuite/libgomp.c/target-33.c: New test. * testsuite/libgomp.c/target-34.c: New test. liboffloadmic/ 2015-11-14 Ilya Verbin <ilya.verbin@intel.com> * runtime/offload_host.cpp (task_completion_callback): New variable. (offload_proxy_task_completed_ooo): Call task_completion_callback. (__offload_register_task_callback): New function. * runtime/offload_host.h (__offload_register_task_callback): New declaration. * plugin/libgomp-plugin-intelmic.cpp (offload): Add async_data argument, handle async offloading. (register_main_image): Call register_main_image. (GOMP_OFFLOAD_init_device, get_target_table, GOMP_OFFLOAD_alloc, GOMP_OFFLOAD_free, GOMP_OFFLOAD_host2dev, GOMP_OFFLOAD_dev2host, GOMP_OFFLOAD_dev2dev) Adjust offload callers. (GOMP_OFFLOAD_async_run): New function. (GOMP_OFFLOAD_run): Implement using GOMP_OFFLOAD_async_run. From-SVN: r230381 |
||
Nathan Sidwell
|
5c06742f6f |
libgomp.h (struct acc_dispatch_t): Remove args from exec_func.
* libgomp.h (struct acc_dispatch_t): Remove args from exec_func. * plugin/plugin-nvptx.c (nvptx_exec): Remove sizes & kinds arg. (GOMP_OFFLOAD_openacc_parallel): Likewise. * oacc-host.c (host_openacc_exec): Likewise. * oacc-parallel.c (GOACC_parallel_keyed): Adjust exec_func call. From-SVN: r229721 |
||
Jakub Jelinek
|
d9a6bd32ad |
builtin-types.def (BT_FN_BOOL_UINT_LONGPTR_LONGPTR_LONGPTR, [...]): New.
gcc/ 2015-10-13 Jakub Jelinek <jakub@redhat.com> Aldy Hernandez <aldyh@redhat.com> Ilya Verbin <ilya.verbin@intel.com> * builtin-types.def (BT_FN_BOOL_UINT_LONGPTR_LONGPTR_LONGPTR, BT_FN_BOOL_UINT_ULLPTR_ULLPTR_ULLPTR, BT_FN_BOOL_UINT_LONGPTR_LONG_LONGPTR_LONGPTR, BT_FN_BOOL_UINT_ULLPTR_ULL_ULLPTR_ULLPTR, BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR, BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT, BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_UINT_LONG_INT_LONG_LONG_LONG, BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_UINT_LONG_INT_ULL_ULL_ULL, BT_FN_VOID_LONG_VAR, BT_FN_VOID_ULL_VAR): New. (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR): Remove. * cgraph.h (enum cgraph_simd_clone_arg_type): Add SIMD_CLONE_ARG_TYPE_LINEAR_REF_CONSTANT_STEP, SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP and SIMD_CLONE_ARG_TYPE_LINEAR_VAL_CONSTANT_STEP. (struct cgraph_simd_clone_arg): Adjust comment. * coretypes.h (struct gomp_ordered): New forward decl. * gimple.c (gimple_build_omp_critical): Add CLAUSES argument, set critical clauses to it. (gimple_build_omp_ordered): Return gomp_ordered * instead of gimple *. Add CLAUSES argument, set ordered clauses to it. (gimple_copy): Unshare clauses on GIMPLE_OMP_CRITICAL and GIMPLE_OMP_ORDERED. * gimple.def (GIMPLE_OMP_ORDERED): Change from GSS_OMP to GSS_OMP_SINGLE_LAYOUT, move it after GIMPLE_OMP_TEAMS. * gimple.h (enum gf_mask): Add GF_OMP_TASK_TASKLOOP. Add another bit to GF_OMP_FOR_KIND_MASK mask. Add GF_OMP_FOR_KIND_TASKLOOP, renumber GF_OMP_FOR_KIND_CILKFOR and GF_OMP_FOR_KIND_OACC_LOOP. Adjust GF_OMP_FOR_SIMD, GF_OMP_FOR_COMBINED and GF_OMP_FOR_COMBINED_INTO. Add another bit to GF_OMP_TARGET_KIND_MASK mask. Add GF_OMP_TARGET_KIND_ENTER_DATA and GF_OMP_TARGET_KIND_EXIT_DATA, renumber GF_OMP_TARGET_KIND_OACC_{PARALLEL,KERNELS,DATA,UPDATE,ENTER_EXIT_DATA}. (gomp_critical): Add clauses field. (gomp_ordered): New struct. (is_a_helper <gomp_ordered *>::test): New inline. (gimple_build_omp_critical): Add CLAUSES argument. (gimple_build_omp_ordered): Likewise. Return gomp_ordered * instead of gimple *. (gimple_omp_critical_clauses, gimple_omp_critical_clauses_ptr, gimple_omp_critical_set_clauses, gimple_omp_ordered_clauses, gimple_omp_ordered_clauses_ptr, gimple_omp_ordered_set_clauses, gimple_omp_task_taskloop_p, gimple_omp_task_set_taskloop_p): New inline functions. * gimple-pretty-print.c (dump_gimple_omp_for): Handle taskloop. (dump_gimple_omp_target): Handle enter data and exit data. (dump_gimple_omp_block): Don't handle GIMPLE_OMP_ORDERED here. (dump_gimple_omp_critical): Print clauses. (dump_gimple_omp_ordered): New function. (dump_gimple_omp_task): Handle taskloop. (pp_gimple_stmt_1): Use dump_gimple_omp_ordered for GIMPLE_OMP_ORDERED. * gimple-walk.c (walk_gimple_op): Walk clauses on GIMPLE_OMP_CRITICAL and GIMPLE_OMP_ORDERED. * gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_0LEN_ARRAY. (enum omp_region_type): Add ORT_COMBINED_TARGET and ORT_NONE. (struct gimplify_omp_ctx): Add loop_iter_var, target_map_scalars_firstprivate, target_map_pointers_as_0len_arrays and target_firstprivatize_array_bases fields. (delete_omp_context): Release loop_iter_var. (gimplify_bind_expr): Handle ORT_NONE. (maybe_fold_stmt): Adjust check for ORT_TARGET for the addition of ORT_COMBINED_TARGET. (is_gimple_stmt): Return true for OMP_TASKLOOP, OMP_TEAMS and OMP_TARGET{,_DATA,_UPDATE,_ENTER_DATA,_EXIT_DATA}. (omp_firstprivatize_variable): Handle ORT_NONE. Adjust check for ORT_TARGET for the addition of ORT_COMBINED_TARGET. Handle ctx->target_map_scalars_firstprivate. (omp_add_variable): Handle ORT_NONE. Allow map clause together with data sharing clauses. For data sharing clause with VLA decl on omp target/target data don't add firstprivate for the pointer. Call omp_notice_variable on TYPE_SIZE_UNIT only if it is a DECL_P. (omp_notice_threadprivate_variable): Adjust check for ORT_TARGET for the addition of ORT_COMBINED_TARGET. (omp_notice_variable): Handle ORT_NONE. Adjust check for ORT_TARGET for the addition of ORT_COMBINED_TARGET. Handle implicit mapping of pointers as zero length array sections and ctx->target_map_scalars_firstprivate mapping of scalars as firstprivate data sharing. (omp_check_private): Handle omp_member_access_dummy_var vars. (find_decl_expr): New function. (gimplify_scan_omp_clauses): Add CODE argument. For OMP_CLAUSE_IF complain if OMP_CLAUSE_IF_MODIFIER is present and does not match code. Handle OMP_CLAUSE_GANG separately. Handle OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,NOGROUP,THREADS,SIMD,SIMDLEN} clauses. Diagnose linear clause on combined distribute {, parallel for} simd construct, unless it is the loop iterator. Handle struct element GOMP_MAP_FIRSTPRIVATE_POINTER. Handle map clauses with COMPONENT_REF. Initialize ctx->target_map_scalars_firstprivate, ctx->target_firstprivatize_array_bases and ctx->target_map_pointers_as_0len_arrays. Add firstprivate for linear clause even to target region if combined. Remove map clauses with GOMP_MAP_FIRSTPRIVATE_POINTER kind from OMP_TARGET_{,ENTER_,EXIT_}DATA. For GOMP_MAP_FIRSTPRIVATE_POINTER map kind with non-INTEGER_CST OMP_CLAUSE_SIZE firstprivatize the bias. Handle OMP_CLAUSE_DEPEND_{SINK,SOURCE}. Handle OMP_CLAUSE_{{USE,IS}_DEVICE_PTR,DEFAULTMAP,HINT}. For linear clause on worksharing loop combined with parallel add shared clause on the parallel. Handle OMP_CLAUSE_REDUCTION with MEM_REF OMP_CLAUSE_DECL. Set DECL_NAME on omp_member_access_dummy_var vars. Add lastprivate clause to outer taskloop if needed. (gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_0LEN_ARRAY. If gimplify_omp_ctxp->target_firstprivatize_array_bases, use GOMP_MAP_FIRSTPRIVATE_POINTER map kind instead of GOMP_MAP_POINTER. (gimplify_adjust_omp_clauses): Add CODE argument. Handle removal of GOMP_MAP_FIRSTPRIVATE_POINTER struct elements for struct not seen in target body. Handle removal of struct mapping if struct is not seen in target body. Remove GOMP_MAP_STRUCT map clause on OMP_TARGET_EXIT_DATA. Adjust check for ORT_TARGET for the addition of ORT_COMBINED_TARGET. Use GOMP_MAP_FIRSTPRIVATE_POINTER instead of GOMP_MAP_POINTER if ctx->target_firstprivatize_array_bases for VLAs. Set OMP_CLAUSE_MAP_PRIVATE if both data sharing and map clause appear together. Handle OMP_CLAUSE_{{USE,IS}_DEVICE_PTR,DEFAULTMAP,HINT}. Don't remove map clause if it has map-type-modifier always. Handle OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,NOGROUP,THREADS,SIMD,SIMDLEN} clauses. (gimplify_oacc_cache, gimplify_omp_parallel, gimplify_omp_task): Adjust gimplify_scan_omp_clauses and gimplify_adjust_omp_clauses callers. (gimplify_omp_for): Likewise. Handle OMP_TASKLOOP. Initialize loop_iter_var. Use OMP_FOR_ORIG_DECLS. Fix handling of lastprivate iterators in doacross loops. (gimplify_omp_workshare): Adjust gimplify_scan_omp_clauses and gimplify_adjust_omp_clauses callers. Use ORT_COMBINED_TARGET for OMP_TARGET_COMBINED. Adjust check for ORT_TARGET for the addition of ORT_COMBINED_TARGET. (gimplify_omp_target_update): Adjust gimplify_scan_omp_clauses and gimplify_adjust_omp_clauses callers. Handle OMP_TARGET_ENTER_DATA and OMP_TARGET_EXIT_DATA. (gimplify_omp_ordered): New function. (gimplify_expr): Handle OMP_TASKLOOP, OMP_TARGET_ENTER_DATA and OMP_TARGET_EXIT_DATA. Use gimplify_omp_ordered for OMP_ORDERED. Gimplify clauses on OMP_CRITICAL. * internal-fn.c (expand_GOMP_SIMD_ORDERED_START, expand_GOMP_SIMD_ORDERED_END): New functions. * internal-fn.def (GOMP_SIMD_ORDERED_START, GOMP_SIMD_ORDERED_END): New internal functions. * omp-builtins.def (BUILT_IN_GOMP_LOOP_DOACROSS_STATIC_START, BUILT_IN_GOMP_LOOP_DOACROSS_DYNAMIC_START, BUILT_IN_GOMP_LOOP_DOACROSS_GUIDED_START, BUILT_IN_GOMP_LOOP_DOACROSS_RUNTIME_START, BUILT_IN_GOMP_LOOP_ULL_DOACROSS_STATIC_START, BUILT_IN_GOMP_LOOP_ULL_DOACROSS_DYNAMIC_START, BUILT_IN_GOMP_LOOP_ULL_DOACROSS_GUIDED_START, BUILT_IN_GOMP_LOOP_ULL_DOACROSS_RUNTIME_START, BUILT_IN_GOMP_DOACROSS_POST, BUILT_IN_GOMP_DOACROSS_WAIT, BUILT_IN_GOMP_DOACROSS_ULL_POST, BUILT_IN_GOMP_DOACROSS_ULL_WAIT, BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA, BUILT_IN_GOMP_TASKLOOP, BUILT_IN_GOMP_TASKLOOP_ULL): New built-ins. (BUILT_IN_GOMP_TASK): Add INT argument to the end. (BUILT_IN_GOMP_TARGET): Rename from GOMP_target to GOMP_target_41, adjust type. (BUILT_IN_GOMP_TARGET_DATA): Rename from GOMP_target_data to GOMP_target_data_41, adjust type. (BUILT_IN_GOMP_TARGET_UPDATE): Rename from GOMP_target_update to GOMP_target_update_41, adjust type. * omp-low.c (struct omp_region): Adjust comments, add ord_stmt field. (struct omp_for_data): Add ordered and simd_schedule fields. (omp_member_access_dummy_var, unshare_and_remap_1, unshare_and_remap, is_taskloop_ctx): New functions. (is_taskreg_ctx): Use is_parallel_ctx and is_task_ctx. (extract_omp_for_data): Handle taskloops and doacross loops and simd schedule modifier. (omp_adjust_chunk_size): New function. (get_ws_args_for): Use it. (lookup_sfield): Change first argument to splay_tree_key, add overload with first argument tree. (maybe_lookup_field): Likewise. (use_pointer_for_field): Handle omp_member_access_dummy_var. (omp_copy_decl_2): If var is TREE_ADDRESSABLE listed in task_shared_vars, clear TREE_ADDRESSABLE on the copy. (build_outer_var_ref): Add LASTPRIVATE argument, handle taskloops and omp_member_access_dummy_var vars. (build_sender_ref): Change first argument to splay_tree_key, add overload with first argument tree. (install_var_field): For mask & 8 use &DECL_UID as key instead of the tree itself. (fixup_child_record_type): Const qualify *.omp_data_i. (scan_sharing_clauses): Handle OMP_CLAUSE_SHARED_FIRSTPRIVATE, C/C++ array reductions, OMP_CLAUSE_{IS,USE}_DEVICE_PTR clauses, OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,SIMDLEN,THREADS,SIMD} and OMP_CLAUSE_{NOGROUP,DEFAULTMAP} clauses, OMP_CLAUSE__LOOPTEMP_ clause on taskloop, GOMP_MAP_FIRSTPRIVATE_POINTER, OMP_CLAUSE_MAP_PRIVATE. (create_omp_child_function): Set TREE_READONLY on .omp_data_i. (find_combined_for): Allow searching for different GIMPLE_OMP_FOR kinds. (add_taskreg_looptemp_clauses): New function. (scan_omp_parallel): Use it. (scan_omp_task): Likewise. (finish_taskreg_scan): Handle OMP_CLAUSE_SHARED_FIRSTPRIVATE. For taskloop, move fields for the first two _LOOPTEMP_ clauses first. (check_omp_nesting_restrictions): Handle GF_OMP_TARGET_KIND_ENTER_DATA and GF_OMP_TARGET_KIND_EXIT_DATA. Formatting fixes. Allow the sandwiched taskloop constructs. Type check OMP_CLAUSE_DEPEND_{KIND,SOURCE}. Allow ordered simd inside of simd region. Diagnose depend(source) or depend(sink:...) on target constructs or task/taskloop. (handle_simd_reference): Use get_name. (lower_rec_input_clauses): Likewise. Ignore all OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE clauses on taskloop construct. Allow _LOOPTEMP_ clause on GOMP_TASK. Unshare new_var before passing it to omp_clause_{default,copy}_ctor. Handle OMP_CLAUSE_REDUCTION with MEM_REF OMP_CLAUSE_DECL. Set lastprivate_firstprivate flag for linear that needs copyin and copyout. Use BUILT_IN_ALLOCA_WITH_ALIGN instead of BUILT_IN_ALLOCA. (lower_lastprivate_clauses): For OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE on taskloop lookup decl in outer context. Pass true to build_outer_var_ref lastprivate argument. Handle OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV lastprivate if the decl is global outside of outer taskloop for. (lower_reduction_clauses): Handle OMP_CLAUSE_REDUCTION with MEM_REF OMP_CLAUSE_DECL. (lower_send_clauses): Ignore first two _LOOPTEMP_ clauses in taskloop GOMP_TASK. Handle OMP_CLAUSE_SHARED_FIRSTPRIVATE. Handle omp_member_access_dummy_var vars. Handle OMP_CLAUSE_REDUCTION with MEM_REF OMP_CLAUSE_DECL. Use new lookup_sfield overload. (lower_send_shared_vars): Ignore fields with NULL or FIELD_DECL abstract origin. Handle omp_member_access_dummy_var vars. (expand_parallel_call): Use expand_omp_build_assign. (expand_task_call): Handle taskloop construct expansion. Add REGION argument. Use GOMP_TASK_* defines instead of hardcoded integers. Add priority argument to GOMP_task* calls. Or in GOMP_TASK_FLAG_PRIORITY into flags if priority is present for GOMP_task call. (expand_omp_build_assign): Add prototype. Add AFTER argument, if true emit statements after *GSI_P and continue linking. (expand_omp_taskreg): Adjust expand_task_call caller. (expand_omp_for_init_counts): Rename zero_iter_bb argument to zero_iter1_bb and first_zero_iter to first_zero_iter1, add zero_iter2_bb and first_zero_iter2 arguments, handle computation of counts even for ordered loops. (expand_omp_for_init_vars): Handle GOMP_TASK inner_stmt. (expand_omp_ordered_source, expand_omp_ordered_sink, expand_omp_ordered_source_sink, expand_omp_for_ordered_loops): New functions. (expand_omp_for_generic): Use omp_adjust_chunk_size. Handle linear clauses on worksharing loop. Handle DOACROSS loop expansion. (expand_omp_for_static_nochunk): Handle linear clauses on worksharing loop. Adjust expand_omp_for_init_counts callers. (expand_omp_for_static_chunk): Likewise. Use omp_adjust_chunk_size. (expand_omp_simd): Handle addressable fd->loop.v. Adjust expand_omp_for_init_counts callers. (expand_omp_taskloop_for_outer, expand_omp_taskloop_for_inner): New functions. (expand_omp_for): Call expand_omp_taskloop_for_* for taskloop. Handle doacross loops. (expand_omp_target): Handle GF_OMP_TARGET_KIND_ENTER_DATA and GF_OMP_TARGET_KIND_EXIT_DATA. Pass flags and depend arguments to GOMP_target_{41,update_41,enter_exit_data} libcalls. (expand_omp): Don't expand ordered depend constructs here, record ord_stmt instead for later expand_omp_for_generic. (build_omp_regions_1): Handle GF_OMP_TARGET_KIND_ENTER_DATA and GF_OMP_TARGET_KIND_EXIT_DATA. Treat GIMPLE_OMP_ORDERED with depend clause as stand-alone directive. (lower_omp_ordered_clauses): New function. (lower_omp_ordered): Handle OMP_CLAUSE_SIMD, for OMP_CLAUSE_DEPEND don't lower anything. (lower_omp_for_lastprivate): Use last _looptemp_ clause on taskloop for comparison. (lower_omp_for): Handle taskloop constructs. Adjust OMP_CLAUSE_DECL and OMP_CLAUSE_LINEAR_STEP so that expand_omp_for_* can use it during expansion for linear adjustments. (create_task_copyfn): Handle OMP_CLAUSE_SHARED_FIRSTPRIVATE. (lower_depend_clauses): Assert not seeing sink/source depend kinds. Set TREE_ADDRESSABLE on array. Change first argument from gimple * to tree * pointing to the stmt's clauses. (lower_omp_taskreg): Adjust lower_depend_clauses caller. (lower_omp_target): Handle GF_OMP_TARGET_KIND_ENTER_DATA and GF_OMP_TARGET_KIND_EXIT_DATA, depend clauses, GOMP_MAP_{RELEASE,ALWAYS_{TO,FROM,TOFROM},FIRSTPRIVATE_POINTER,STRUCT} map kinds, OMP_CLAUSE_{FIRSTPRIVATE,PRIVATE,{IS,USE}_DEVICE_PTR clauses. Always use short kind and 8-bit align shift. (lower_omp_regimplify_p): Use IS_TYPE_OR_DECL_P macro. (struct lower_omp_regimplify_operands_data): New type. (lower_omp_regimplify_operands_p, lower_omp_regimplify_operands): New functions. (lower_omp_1): Use lower_omp_regimplify_operands instead of gimple_regimplify_operands. (make_gimple_omp_edges): Handle GF_OMP_TARGET_KIND_ENTER_DATA and GF_OMP_TARGET_KIND_EXIT_DATA. Treat GIMPLE_OMP_ORDERED with depend clause as stand-alone directive. (simd_clone_clauses_extract): Honor OMP_CLAUSE_LINEAR_KIND. (simd_clone_mangle): Mangle the various linear kinds per the new ABI. (simd_clone_adjust_argument_types): Handle SIMD_CLONE_ARG_TYPE_LINEAR_*_CONSTANT_STEP. (simd_clone_init_simd_arrays): Don't do anything for uval. (simd_clone_adjust): Handle SIMD_CLONE_ARG_TYPE_LINEAR_REF_CONSTANT_STEP like SIMD_CLONE_ARG_TYPE_LINEAR_CONSTANT_STEP. Handle SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP. * omp-low.h (omp_member_access_dummy_var): New prototype. * passes.def (pass_simduid_cleanup): Schedule another copy of the pass after all optimizations. * tree.c (omp_clause_code_name): Add entries for OMP_CLAUSE_{TO_DECLARE,LINK,{USE,IS}_DEVICE_PTR,DEFAULTMAP,HINT} and OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,NOGROUP,THREADS,SIMD}. (omp_clause_num_ops): Likewise. Bump number of OMP_CLAUSE_REDUCTION arguments to 5 and for OMP_CLAUSE_ORDERED to 1. (walk_tree_1): Adjust for OMP_CLAUSE_ORDERED having 1 argument and OMP_CLAUSE_REDUCTION 5 arguments. Handle OMP_CLAUSE_{TO_DECLARE,LINK,{USE,IS}_DEVICE_PTR,DEFAULTMAP,HINT} and OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,NOGROUP,THREADS,SIMD} clauses. * tree-core.h (enum omp_clause_linear_kind): New. (struct tree_omp_clause): Change type of map_kind from unsigned char to unsigned int. Add subcode.if_modifier and subcode.linear_kind fields. (enum omp_clause_code): Add OMP_CLAUSE_{TO_DECLARE,LINK,{USE,IS}_DEVICE_PTR,DEFAULTMAP,HINT} and OMP_CLAUSE_{PRIORITY,GRAINSIZE,NUM_TASKS,NOGROUP,THREADS,SIMD}. (OMP_CLAUSE_REDUCTION): Document OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER. (enum omp_clause_depend_kind): Add OMP_CLAUSE_DEPEND_{SOURCE,SINK}. * tree.def (OMP_FOR): Add OMP_FOR_ORIG_DECLS operand. (OMP_CRITICAL): Move before OMP_SINGLE. Add OMP_CRITICAL_CLAUSES operand. (OMP_ORDERED): Move before OMP_SINGLE. Add OMP_ORDERED_CLAUSES operand. (OMP_TASKLOOP, OMP_TARGET_ENTER_DATA, OMP_TARGET_EXIT_DATA): New tree codes. * tree.h (OMP_BODY): Replace OMP_CRITICAL with OMP_TASKGROUP. (OMP_CLAUSE_SET_MAP_KIND): Cast to unsigned int rather than unsigned char. (OMP_CRITICAL_NAME): Adjust to be 3rd operand instead of 2nd. (OMP_CLAUSE_NUM_TASKS_EXPR): Formatting fix. (OMP_STANDALONE_CLAUSES): Adjust to cover OMP_TARGET_{ENTER,EXIT}_DATA. (OMP_CLAUSE_DEPEND_SINK_NEGATIVE, OMP_TARGET_COMBINED, OMP_CLAUSE_MAP_PRIVATE, OMP_FOR_ORIG_DECLS, OMP_CLAUSE_IF_MODIFIER, OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION, OMP_CRITICAL_CLAUSES, OMP_CLAUSE_PRIVATE_TASKLOOP_IV, OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV, OMP_CLAUSE_HINT_EXPR, OMP_CLAUSE_SCHEDULE_SIMD, OMP_CLAUSE_LINEAR_KIND, OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER, OMP_CLAUSE_SHARED_FIRSTPRIVATE, OMP_ORDERED_CLAUSES, OMP_TARGET_ENTER_DATA_CLAUSES, OMP_TARGET_EXIT_DATA_CLAUSES, OMP_CLAUSE_NUM_TASKS_EXPR, OMP_CLAUSE_GRAINSIZE_EXPR, OMP_CLAUSE_PRIORITY_EXPR, OMP_CLAUSE_ORDERED_EXPR): Define. * tree-inline.c (remap_gimple_stmt): Handle clauses on GIMPLE_OMP_ORDERED and GIMPLE_OMP_CRITICAL. For IFN_GOMP_SIMD_ORDERED_{START,END} set has_simduid_loops. * tree-nested.c (convert_nonlocal_omp_clauses): Handle OMP_CLAUSE_{TO_DECLARE,LINK,{USE,IS}_DEVICE_PTR,SIMDLEN,PRIORITY,SIMD} and OMP_CLAUSE_{GRAINSIZE,NUM_TASKS,HINT,NOGROUP,THREADS,DEFAULTMAP} clauses. Handle OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER. (convert_local_omp_clauses): Likewise. * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_{TO_DECLARE,LINK,{USE,IS}_DEVICE_PTR,SIMDLEN,PRIORITY,SIMD} and OMP_CLAUSE_{GRAINSIZE,NUM_TASKS,HINT,NOGROUP,THREADS,DEFAULTMAP} clauses. Handle OMP_CLAUSE_IF_MODIFIER, OMP_CLAUSE_ORDERED_EXPR, OMP_CLAUSE_SCHEDULE_SIMD, OMP_CLAUSE_LINEAR_KIND, OMP_CLAUSE_DEPEND_{SOURCE,SINK}. Use "delete" for GOMP_MAP_FORCE_DEALLOC. Handle GOMP_MAP_{ALWAYS_{TO,FROM,TOFROM},RELEASE,FIRSTPRIVATE_POINTER,STRUCT}. (dump_generic_node): Handle OMP_TASKLOOP, OMP_TARGET_{ENTER,EXIT}_DATA and clauses on OMP_ORDERED and OMP_CRITICAL. * tree-vectorizer.c (adjust_simduid_builtins): Adjust comment. Remove IFN_GOMP_SIMD_ORDERED_{START,END}. (vectorize_loops): Adjust comments. (pass_simduid_cleanup::execute): Likewise. * tree-vect-stmts.c (vectorizable_simd_clone_call): Handle SIMD_CLONE_ARG_TYPE_LINEAR_{REF,VAL,UVAL}_CONSTANT_STEP. * wide-int.h (wi::gcd): New. gcc/c-family/ 2015-10-13 Jakub Jelinek <jakub@redhat.com> Aldy Hernandez <aldyh@redhat.com> * c-common.c (enum c_builtin_type): Define DEF_FUNCTION_TYPE_9, DEF_FUNCTION_TYPE_10 and DEF_FUNCTION_TYPE_11. (c_define_builtins): Likewise. * c-common.h (enum c_omp_clause_split): Add C_OMP_CLAUSE_SPLIT_TASKLOOP. (c_finish_omp_critical, c_finish_omp_ordered): Add CLAUSES argument. (c_finish_omp_for): Add ORIG_DECLV argument. * c-cppbuiltin.c (c_cpp_builtins): Predefine _OPENMP as 201511 instead of 201307. * c-omp.c (c_finish_omp_critical): Add CLAUSES argument, set OMP_CRITICAL_CLAUSES to it. (c_finish_omp_ordered): Add CLAUSES argument, set OMP_ORDERED_CLAUSES to it. (c_finish_omp_for): Add ORIG_DECLV argument, set OMP_FOR_ORIG_DECLS to it if OMP_FOR. Clear DECL_INITIAL on the IVs. (c_omp_split_clauses): Handle OpenMP 4.5 combined/composite constructs and new OpenMP 4.5 clauses. Clear OMP_CLAUSE_SCHEDULE_SIMD if not combined with OMP_SIMD. Add verification code. * c-pragma.c (omp_pragmas_simd): Add taskloop. * c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_TASKLOOP. (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_{DEFAULTMAP,GRAINSIZE,HINT,{IS,USE}_DEVICE_PTR} and PRAGMA_OMP_CLAUSE_{LINK,NOGROUP,NUM_TASKS,PRIORITY,SIMD,THREADS}. gcc/c/ 2015-10-13 Jakub Jelinek <jakub@redhat.com> Aldy Hernandez <aldyh@redhat.com> * c-parser.c (c_parser_pragma): Handle PRAGMA_OMP_ORDERED here. (c_parser_omp_clause_name): Handle OpenMP 4.5 clauses. (c_parser_omp_variable_list): Handle structure elements for map, to and from clauses. Handle array sections in reduction clause. Formatting fixes. (c_parser_omp_clause_if): Add IS_OMP argument, handle parsing of if clause modifiers. (c_parser_omp_clause_num_tasks, c_parser_omp_clause_grainsize, c_parser_omp_clause_priority, c_parser_omp_clause_hint, c_parser_omp_clause_defaultmap, c_parser_omp_clause_use_device_ptr, c_parser_omp_clause_is_device_ptr): New functions. (c_parser_omp_clause_ordered): Parse optional parameter. (c_parser_omp_clause_reduction): Handle array reductions. (c_parser_omp_clause_schedule): Parse optional simd modifier. (c_parser_omp_clause_nogroup, c_parser_omp_clause_orderedkind): New functions. (c_parser_omp_clause_linear): Parse linear clause modifiers. (c_parser_omp_clause_depend_sink): New function. (c_parser_omp_clause_depend): Parse source/sink depend kinds. (c_parser_omp_clause_map): Parse release/delete map kinds and optional always modifier. (c_parser_oacc_all_clauses): Adjust c_parser_omp_clause_if and c_finish_omp_clauses callers. (c_parser_omp_all_clauses): Likewise. Parse OpenMP 4.5 clauses. Parse "to" as OMP_CLAUSE_TO_DECLARE if on declare target directive. (c_parser_oacc_cache): Adjust c_finish_omp_clauses caller. (OMP_CRITICAL_CLAUSE_MASK): Define. (c_parser_omp_critical): Parse critical clauses. (c_parser_omp_for_loop): Handle doacross loops, adjust c_finish_omp_for and c_finish_omp_clauses callers. (OMP_SIMD_CLAUSE_MASK): Add simdlen clause. (c_parser_omp_simd): Allow ordered clause if it has no parameter. (OMP_FOR_CLAUSE_MASK): Add linear clause. (c_parser_omp_for): Disallow ordered clause when combined with distribute. Disallow linear clause when combined with distribute and not combined with simd. (OMP_ORDERED_CLAUSE_MASK, OMP_ORDERED_DEPEND_CLAUSE_MASK): Define. (c_parser_omp_ordered): Add CONTEXT argument, remove LOC argument, parse clauses and if depend clause is found, don't parse a body. (c_parser_omp_parallel): Disallow copyin clause on target parallel. Allow target parallel without for after it. (OMP_TASK_CLAUSE_MASK): Add priority clause. (OMP_TARGET_DATA_CLAUSE_MASK): Add use_device_ptr clause. (c_parser_omp_target_data): Diagnose no map clauses or clauses with invalid kinds. (OMP_TARGET_UPDATE_CLAUSE_MASK): Add depend and nowait clauses. (OMP_TARGET_ENTER_DATA_CLAUSE_MASK, OMP_TARGET_EXIT_DATA_CLAUSE_MASK): Define. (c_parser_omp_target_enter_data, c_parser_omp_target_exit_data): New functions. (OMP_TARGET_CLAUSE_MASK): Add depend, nowait, private, firstprivate, defaultmap and is_device_ptr clauses. (c_parser_omp_target): Parse target parallel and target simd. Set OMP_TARGET_COMBINED on combined constructs. Parse target enter data and target exit data. Diagnose invalid map kinds. (OMP_DECLARE_TARGET_CLAUSE_MASK): Define. (c_parser_omp_declare_target): Parse OpenMP 4.5 forms of this construct. (c_parser_omp_declare_reduction): Use STRIP_NOPS when checking for &omp_priv. (OMP_TASKLOOP_CLAUSE_MASK): Define. (c_parser_omp_taskloop): New function. (c_parser_omp_construct): Don't handle PRAGMA_OMP_ORDERED here, handle PRAGMA_OMP_TASKLOOP. (c_parser_cilk_for): Adjust c_finish_omp_clauses callers. * c-tree.h (c_finish_omp_clauses): Add two new arguments. * c-typeck.c (handle_omp_array_sections_1): Fix comment typo. Add IS_OMP argument, handle structure element bases, diagnose bitfields, pass IS_OMP recursively, diagnose known zero length array sections in depend clauses, handle array sections in reduction clause, diagnose negative length even for pointers. (handle_omp_array_sections): Add IS_OMP argument, use auto_vec for types, pass IS_OMP down to handle_omp_array_sections_1, handle array sections in reduction clause, set OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION if map could be zero length array section, use GOMP_MAP_FIRSTPRIVATE_POINTER for IS_OMP. (c_finish_omp_clauses): Add IS_OMP and DECLARE_SIMD arguments. Handle new OpenMP 4.5 clauses and new restrictions for the old ones. gcc/cp/ 2015-10-13 Jakub Jelinek <jakub@redhat.com> Aldy Hernandez <aldyh@redhat.com> * class.c (finish_struct_1): Call finish_omp_declare_simd_methods. * cp-gimplify.c (cp_gimplify_expr): Handle OMP_TASKLOOP. (cp_genericize_r): Likewise. (cxx_omp_finish_clause): Don't diagnose references. (cxx_omp_disregard_value_expr): New function. * cp-objcp-common.h (LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR): Redefine. * cp-tree.h (OMP_FOR_GIMPLIFYING_P): Document for OMP_TASKLOOP. (DECL_OMP_PRIVATIZED_MEMBER): Define. (finish_omp_declare_simd_methods, push_omp_privatization_clauses, pop_omp_privatization_clauses, save_omp_privatization_clauses, restore_omp_privatization_clauses, omp_privatize_field, cxx_omp_disregard_value_expr): New prototypes. (finish_omp_clauses): Add two new arguments. (finish_omp_for): Add ORIG_DECLV argument. * parser.c (cp_parser_lambda_body): Call save_omp_privatization_clauses and restore_omp_privatization_clauses. (cp_parser_omp_clause_name): Handle OpenMP 4.5 clauses. (cp_parser_omp_var_list_no_open): Handle structure elements for map, to and from clauses. Handle array sections in reduction clause. Parse this keyword. Formatting fixes. (cp_parser_omp_clause_if): Add IS_OMP argument, handle parsing of if clause modifiers. (cp_parser_omp_clause_num_tasks, cp_parser_omp_clause_grainsize, cp_parser_omp_clause_priority, cp_parser_omp_clause_hint, cp_parser_omp_clause_defaultmap): New functions. (cp_parser_omp_clause_ordered): Parse optional parameter. (cp_parser_omp_clause_reduction): Handle array reductions. (cp_parser_omp_clause_schedule): Parse optional simd modifier. (cp_parser_omp_clause_nogroup, cp_parser_omp_clause_orderedkind): New functions. (cp_parser_omp_clause_linear): Parse linear clause modifiers. (cp_parser_omp_clause_depend_sink): New function. (cp_parser_omp_clause_depend): Parse source/sink depend kinds. (cp_parser_omp_clause_map): Parse release/delete map kinds and optional always modifier. (cp_parser_oacc_all_clauses): Adjust cp_parser_omp_clause_if and finish_omp_clauses callers. (cp_parser_omp_all_clauses): Likewise. Parse OpenMP 4.5 clauses. Parse "to" as OMP_CLAUSE_TO_DECLARE if on declare target directive. (OMP_CRITICAL_CLAUSE_MASK): Define. (cp_parser_omp_critical): Parse critical clauses. (cp_parser_omp_for_incr): Use cp_tree_equal if processing_template_decl. (cp_parser_omp_for_loop_init): Return tree instead of bool. Handle non-static data member iterators. (cp_parser_omp_for_loop): Handle doacross loops, adjust finish_omp_for and finish_omp_clauses callers. (cp_omp_split_clauses): Adjust finish_omp_clauses caller. (OMP_SIMD_CLAUSE_MASK): Add simdlen clause. (cp_parser_omp_simd): Allow ordered clause if it has no parameter. (OMP_FOR_CLAUSE_MASK): Add linear clause. (cp_parser_omp_for): Disallow ordered clause when combined with distribute. Disallow linear clause when combined with distribute and not combined with simd. (OMP_ORDERED_CLAUSE_MASK, OMP_ORDERED_DEPEND_CLAUSE_MASK): Define. (cp_parser_omp_ordered): Add CONTEXT argument, return bool instead of tree, parse clauses and if depend clause is found, don't parse a body. (cp_parser_omp_parallel): Disallow copyin clause on target parallel. Allow target parallel without for after it. (OMP_TASK_CLAUSE_MASK): Add priority clause. (OMP_TARGET_DATA_CLAUSE_MASK): Add use_device_ptr clause. (cp_parser_omp_target_data): Diagnose no map clauses or clauses with invalid kinds. (OMP_TARGET_UPDATE_CLAUSE_MASK): Add depend and nowait clauses. (OMP_TARGET_ENTER_DATA_CLAUSE_MASK, OMP_TARGET_EXIT_DATA_CLAUSE_MASK): Define. (cp_parser_omp_target_enter_data, cp_parser_omp_target_exit_data): New functions. (OMP_TARGET_CLAUSE_MASK): Add depend, nowait, private, firstprivate, defaultmap and is_device_ptr clauses. (cp_parser_omp_target): Parse target parallel and target simd. Set OMP_TARGET_COMBINED on combined constructs. Parse target enter data and target exit data. Diagnose invalid map kinds. (cp_parser_oacc_cache): Adjust finish_omp_clauses caller. (OMP_DECLARE_TARGET_CLAUSE_MASK): Define. (cp_parser_omp_declare_target): Parse OpenMP 4.5 forms of this construct. (OMP_TASKLOOP_CLAUSE_MASK): Define. (cp_parser_omp_taskloop): New function. (cp_parser_omp_construct): Don't handle PRAGMA_OMP_ORDERED here, handle PRAGMA_OMP_TASKLOOP. (cp_parser_pragma): Handle PRAGMA_OMP_ORDERED here directly, handle PRAGMA_OMP_TASKLOOP, call push_omp_privatization_clauses and pop_omp_privatization_clauses around parsing calls. (cp_parser_cilk_for): Adjust finish_omp_clauses caller. * pt.c (apply_late_template_attributes): Adjust tsubst_omp_clauses and finish_omp_clauses callers. (tsubst_omp_clause_decl): Return NULL if decl is NULL. For TREE_LIST, copy over OMP_CLAUSE_DEPEND_SINK_NEGATIVE bit. Use tsubst_expr instead of tsubst_copy, undo convert_from_reference effects. (tsubst_omp_clauses): Add ALLOW_FIELDS argument. Handle new OpenMP 4.5 clauses. Use tsubst_omp_clause_decl for more clauses. If ALLOW_FIELDS, handle non-static data members in the clauses. Clear OMP_CLAUSE_LINEAR_STEP if it has been cleared before. (omp_parallel_combined_clauses): New variable. (tsubst_omp_for_iterator): Add ORIG_DECLV argument, recur on OMP_FOR_ORIG_DECLS, handle non-static data member iterators. Improve handling of clauses on combined constructs. (tsubst_expr): Call push_omp_privatization_clauses and pop_omp_privatization_clauses around instantiation of certain OpenMP constructs, improve handling of clauses on combined constructs, handle OMP_TASKLOOP, adjust tsubst_omp_for_iterator, tsubst_omp_clauses and finish_omp_for callers, handle clauses on critical and ordered, handle OMP_TARGET_{ENTER,EXIT}_DATA. (instantiate_decl): Call save_omp_privatization_clauses and restore_omp_privatization_clauses around instantiation. (dependent_omp_for_p): Fix up comment typo. Handle SCOPE_REF. * semantics.c (omp_private_member_map, omp_private_member_vec, omp_private_member_ignore_next): New variables. (finish_non_static_data_member): Return dummy decl for privatized non-static data members. (omp_clause_decl_field, omp_clause_printable_decl, omp_note_field_privatization, omp_privatize_field): New functions. (handle_omp_array_sections_1): Fix comment typo. Add IS_OMP argument, handle structure element bases, diagnose bitfields, pass IS_OMP recursively, diagnose known zero length array sections in depend clauses, handle array sections in reduction clause, diagnose negative length even for pointers. (handle_omp_array_sections): Add IS_OMP argument, use auto_vec for types, pass IS_OMP down to handle_omp_array_sections_1, handle array sections in reduction clause, set OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION if map could be zero length array section, use GOMP_MAP_FIRSTPRIVATE_POINTER for IS_OMP. (finish_omp_reduction_clause): Handle array sections and arrays. Use omp_clause_printable_decl. (finish_omp_declare_simd_methods, cp_finish_omp_clause_depend_sink): New functions. (finish_omp_clauses): Add ALLOW_FIELDS and DECLARE_SIMD arguments. Handle new OpenMP 4.5 clauses and new restrictions for the old ones, handle non-static data members, reject this keyword when not allowed. (push_omp_privatization_clauses, pop_omp_privatization_clauses, save_omp_privatization_clauses, restore_omp_privatization_clauses): New functions. (handle_omp_for_class_iterator): Handle OMP_TASKLOOP class iterators. Add collapse and ordered arguments. Fix handling of lastprivate iterators in doacross loops. (finish_omp_for): Add ORIG_DECLV argument, handle doacross loops, adjust c_finish_omp_for, handle_omp_for_class_iterator and finish_omp_clauses callers. Fill in OMP_CLAUSE_LINEAR_STEP on simd loops with non-static data member iterators. gcc/fortran/ 2015-10-13 Jakub Jelinek <jakub@redhat.com> Ilya Verbin <ilya.verbin@intel.com> * f95-lang.c (DEF_FUNCTION_TYPE_9, DEF_FUNCTION_TYPE_10, DEF_FUNCTION_TYPE_11, DEF_FUNCTION_TYPE_VAR_1): Define. * trans-openmp.c (gfc_trans_omp_clauses): Set OMP_CLAUSE_IF_MODIFIER to ERROR_MARK, OMP_CLAUSE_ORDERED_EXPR to NULL. (gfc_trans_omp_critical): Adjust for addition of clauses. (gfc_trans_omp_ordered): Likewise. * types.def (BT_FN_BOOL_UINT_LONGPTR_LONGPTR_LONGPTR, BT_FN_BOOL_UINT_ULLPTR_ULLPTR_ULLPTR, BT_FN_BOOL_UINT_LONGPTR_LONG_LONGPTR_LONGPTR, BT_FN_BOOL_UINT_ULLPTR_ULL_ULLPTR_ULLPTR, BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR, BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT, BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_UINT_LONG_INT_LONG_LONG_LONG, BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_UINT_LONG_INT_ULL_ULL_ULL, BT_FN_VOID_LONG_VAR, BT_FN_VOID_ULL_VAR): New. (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR): Remove. gcc/lto/ 2015-10-13 Jakub Jelinek <jakub@redhat.com> * lto-lang.c (DEF_FUNCTION_TYPE_9, DEF_FUNCTION_TYPE_10, DEF_FUNCTION_TYPE_11): Define. gcc/jit/ 2015-10-13 Jakub Jelinek <jakub@redhat.com> * jit-builtins.c (DEF_FUNCTION_TYPE_9, DEF_FUNCTION_TYPE_10, DEF_FUNCTION_TYPE_11): Define. * jit-builtins.h (DEF_FUNCTION_TYPE_9, DEF_FUNCTION_TYPE_10, DEF_FUNCTION_TYPE_11): Define. gcc/ada/ 2015-10-13 Jakub Jelinek <jakub@redhat.com> * gcc-interface/utils.c (DEF_FUNCTION_TYPE_9, DEF_FUNCTION_TYPE_10, DEF_FUNCTION_TYPE_11): Define. gcc/testsuite/ 2015-10-13 Jakub Jelinek <jakub@redhat.com> Aldy Hernandez <aldyh@redhat.com> * c-c++-common/gomp/cancel-1.c (f2): Add map clause to target data. * c-c++-common/gomp/clauses-1.c: New test. * c-c++-common/gomp/clauses-2.c: New test. * c-c++-common/gomp/clauses-3.c: New test. * c-c++-common/gomp/clauses-4.c: New test. * c-c++-common/gomp/declare-target-1.c: New test. * c-c++-common/gomp/declare-target-2.c: New test. * c-c++-common/gomp/depend-3.c: New test. * c-c++-common/gomp/depend-4.c: New test. * c-c++-common/gomp/doacross-1.c: New test. * c-c++-common/gomp/if-1.c: New test. * c-c++-common/gomp/if-2.c: New test. * c-c++-common/gomp/linear-1.c: New test. * c-c++-common/gomp/map-2.c: New test. * c-c++-common/gomp/map-3.c: New test. * c-c++-common/gomp/nesting-1.c (f_omp_parallel, f_omp_target_data): Add map clause to target data. * c-c++-common/gomp/nesting-warn-1.c (f_omp_target): Likewise. * c-c++-common/gomp/ordered-1.c: New test. * c-c++-common/gomp/ordered-2.c: New test. * c-c++-common/gomp/ordered-3.c: New test. * c-c++-common/gomp/pr61486-1.c (foo): Remove linear clause on non-iterator. * c-c++-common/gomp/pr61486-2.c (test, test2): Remove ordered clause and ordered construct where no longer allowed. * c-c++-common/gomp/priority-1.c: New test. * c-c++-common/gomp/reduction-1.c: New test. * c-c++-common/gomp/schedule-simd-1.c: New test. * c-c++-common/gomp/sink-1.c: New test. * c-c++-common/gomp/sink-2.c: New test. * c-c++-common/gomp/sink-3.c: New test. * c-c++-common/gomp/sink-4.c: New test. * c-c++-common/gomp/udr-1.c: New test. * c-c++-common/taskloop-1.c: New test. * c-c++-common/cpp/openmp-define-3.c: Adjust for the new value of _OPENMP macro. * c-c++-common/cilk-plus/PS/body.c (foo): Adjust expected diagnostics. * c-c++-common/goacc-gomp/nesting-fail-1.c (f_acc_parallel, f_acc_kernels, f_acc_data, f_acc_loop): Add map clause to target data. * gcc.dg/gomp/clause-1.c: * gcc.dg/gomp/reduction-1.c: New test. * gcc.dg/gomp/sink-fold-1.c: New test. * gcc.dg/gomp/sink-fold-2.c: New test. * gcc.dg/gomp/sink-fold-3.c: New test. * gcc.dg/vect/vect-simd-clone-15.c: New test. * g++.dg/gomp/clause-1.C (T::test): Remove dg-error on privatization of non-static data members. * g++.dg/gomp/clause-3.C (foo): Remove one dg-error directive. Add some linear clause tests. * g++.dg/gomp/declare-simd-3.C: New test. * g++.dg/gomp/linear-1.C: New test. * g++.dg/gomp/member-1.C: New test. * g++.dg/gomp/member-2.C: New test. * g++.dg/gomp/pr66571-2.C: New test. * g++.dg/gomp/pr67504.C (foo): Add test for ordered clause with dependent argument. * g++.dg/gomp/pr67522.C (foo): Add test for invalid array section in reduction clause. * g++.dg/gomp/reference-1.C: New test. * g++.dg/gomp/sink-1.C: New test. * g++.dg/gomp/sink-2.C: New test. * g++.dg/gomp/sink-3.C: New test. * g++.dg/gomp/task-1.C: Remove both dg-error directives. * g++.dg/gomp/this-1.C: New test. * g++.dg/gomp/this-2.C: New test. * g++.dg/vect/simd-clone-2.cc: New test. * g++.dg/vect/simd-clone-2.h: New test. * g++.dg/vect/simd-clone-3.cc: New test. * g++.dg/vect/simd-clone-4.cc: New test. * g++.dg/vect/simd-clone-4.h: New test. * g++.dg/vect/simd-clone-5.cc: New test. include/ 2015-10-13 Jakub Jelinek <jakub@redhat.com> Ilya Verbin <ilya.verbin@intel.com> * gomp-constants.h (GOMP_MAP_FLAG_ALWAYS): Define. (enum gomp_map_kind): Add GOMP_MAP_FIRSTPRIVATE, GOMP_MAP_FIRSTPRIVATE_INT, GOMP_MAP_USE_DEVICE_PTR, GOMP_MAP_ZERO_LEN_ARRAY_SECTION, GOMP_MAP_ALWAYS_TO, GOMP_MAP_ALWAYS_FROM, GOMP_MAP_ALWAYS_TOFROM, GOMP_MAP_STRUCT, GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION, GOMP_MAP_DELETE, GOMP_MAP_RELEASE, GOMP_MAP_FIRSTPRIVATE_POINTER. (GOMP_MAP_ALWAYS_TO_P, GOMP_MAP_ALWAYS_FROM_P): Define. (GOMP_TASK_FLAG_UNTIED, GOMP_TASK_FLAG_FINAL, GOMP_TASK_FLAG_MERGEABLE, GOMP_TASK_FLAG_DEPEND, GOMP_TASK_FLAG_PRIORITY, GOMP_TASK_FLAG_UP, GOMP_TASK_FLAG_GRAINSIZE, GOMP_TASK_FLAG_IF, GOMP_TASK_FLAG_NOGROUP, GOMP_TARGET_FLAG_NOWAIT, GOMP_TARGET_FLAG_EXIT_DATA, GOMP_TARGET_FLAG_UPDATE): Define. libgomp/ 2015-10-13 Jakub Jelinek <jakub@redhat.com> Aldy Hernandez <aldyh@redhat.com> Ilya Verbin <ilya.verbin@intel.com> * config/linux/affinity.c (omp_get_place_num_procs, omp_get_place_proc_ids, gomp_get_place_proc_ids_8): New functions. * config/linux/doacross.h: New file. * config/posix/affinity.c (omp_get_place_num_procs, omp_get_place_proc_ids, gomp_get_place_proc_ids_8): New functions. * config/posix/doacross.h: New file. * env.c: Include gomp-constants.h. (struct gomp_task_icv): Rename run_sched_modifier to run_sched_chunk_size. (gomp_max_task_priority_var): New variable. (parse_schedule): Rename run_sched_modifier to run_sched_chunk_size. (handle_omp_display_env): Change _OPENMP value from 201307 to 201511. Print OMP_MAX_TASK_PRIORITY. (initialize_env): Parse OMP_MAX_TASK_PRIORITY. (omp_set_schedule, omp_get_schedule): Rename modifier argument to chunk_size and run_sched_modifier to run_sched_chunk_size. (omp_get_max_task_priority, omp_get_initial_device, omp_get_num_places, omp_get_place_num, omp_get_partition_num_places, omp_get_partition_place_nums): New functions. * fortran.c (omp_set_schedule_, omp_set_schedule_8_, omp_get_schedule_, omp_get_schedule_8_): Rename modifier argument to chunk_size. (omp_get_num_places_, omp_get_place_num_procs_, omp_get_place_num_procs_8_, omp_get_place_proc_ids_, omp_get_place_proc_ids_8_, omp_get_place_num_, omp_get_partition_num_places_, omp_get_partition_place_nums_, omp_get_partition_place_nums_8_, omp_get_initial_device_, omp_get_max_task_priority_): New functions. * libgomp_g.h (GOMP_loop_doacross_static_start, GOMP_loop_doacross_dynamic_start, GOMP_loop_doacross_guided_start, GOMP_loop_doacross_runtime_start, GOMP_loop_ull_doacross_static_start, GOMP_loop_ull_doacross_dynamic_start, GOMP_loop_ull_doacross_guided_start, GOMP_loop_ull_doacross_runtime_start, GOMP_doacross_post, GOMP_doacross_wait, GOMP_doacross_ull_post, GOMP_doacross_wait, GOMP_taskloop, GOMP_taskloop_ull, GOMP_target_41, GOMP_target_data_41, GOMP_target_update_41, GOMP_target_enter_exit_data): New prototypes. (GOMP_task): Add prototype argument. * libgomp.h (_LIBGOMP_CHECKING_): Define to 0 if not yet defined. (struct gomp_doacross_work_share): New type. (struct gomp_work_share): Add doacross field. (struct gomp_task_icv): Rename run_sched_modifier to run_sched_chunk_size. (enum gomp_task_kind): Rename GOMP_TASK_IFFALSE to GOMP_TASK_UNDEFERRED. Add comments. (struct gomp_task_depend_entry): Add comments. (struct gomp_task): Likewise. (struct gomp_taskgroup): Likewise. (struct gomp_target_task): New type. (struct gomp_team): Add comment. (gomp_get_place_proc_ids_8, gomp_doacross_init, gomp_doacross_ull_init, gomp_task_maybe_wait_for_dependencies, gomp_create_target_task, gomp_target_task_fn): New prototypes. (struct target_var_desc): New type. (struct target_mem_desc): Adjust comment. Use struct target_var_desc instead of splay_tree_key for list. (REFCOUNT_INFINITY): Define. (struct splay_tree_key_s): Remove copy_from field. (struct gomp_device_descr): Add dev2dev_func field. (enum gomp_map_vars_kind): New enum. (gomp_map_vars): Add one argument. * libgomp.map (OMP_4.5): Export omp_get_max_task_priority, omp_get_max_task_priority_, omp_get_num_places, omp_get_num_places_, omp_get_place_num_procs, omp_get_place_num_procs_, omp_get_place_num_procs_8_, omp_get_place_proc_ids, omp_get_place_proc_ids_, omp_get_place_proc_ids_8_, omp_get_place_num, omp_get_place_num_, omp_get_partition_num_places, omp_get_partition_num_places_, omp_get_partition_place_nums, omp_get_partition_place_nums_, omp_get_partition_place_nums_8_, omp_get_initial_device, omp_get_initial_device_, omp_target_alloc, omp_target_free, omp_target_is_present, omp_target_memcpy, omp_target_memcpy_rect, omp_target_associate_ptr and omp_target_disassociate_ptr. (GOMP_4.0.2): Renamed to ... (GOMP_4.5): ... this. Export GOMP_target_41, GOMP_target_data_41, GOMP_target_update_41, GOMP_target_enter_exit_data, GOMP_taskloop, GOMP_taskloop_ull, GOMP_loop_doacross_dynamic_start, GOMP_loop_doacross_guided_start, GOMP_loop_doacross_runtime_start, GOMP_loop_doacross_static_start, GOMP_doacross_post, GOMP_doacross_wait, GOMP_loop_ull_doacross_dynamic_start, GOMP_loop_ull_doacross_guided_start, GOMP_loop_ull_doacross_runtime_start, GOMP_loop_ull_doacross_static_start, GOMP_doacross_ull_post and GOMP_doacross_ull_wait. * libgomp.texi: Document omp_get_max_task_priority. Rename modifier argument to chunk_size for omp_set_schedule and omp_get_schedule. Document OMP_MAX_TASK_PRIORITY env var. * loop.c (GOMP_loop_runtime_start): Adjust for run_sched_modifier to run_sched_chunk_size renaming. (GOMP_loop_ordered_runtime_start): Likewise. (gomp_loop_doacross_static_start, gomp_loop_doacross_dynamic_start, gomp_loop_doacross_guided_start, GOMP_loop_doacross_runtime_start, GOMP_parallel_loop_runtime_start): New functions. (GOMP_parallel_loop_runtime): Adjust for run_sched_modifier to run_sched_chunk_size renaming. (GOMP_loop_doacross_static_start, GOMP_loop_doacross_dynamic_start, GOMP_loop_doacross_guided_start): New functions or aliases. * loop_ull.c (GOMP_loop_ull_runtime_start): Adjust for run_sched_modifier to run_sched_chunk_size renaming. (GOMP_loop_ull_ordered_runtime_start): Likewise. (gomp_loop_ull_doacross_static_start, gomp_loop_ull_doacross_dynamic_start, gomp_loop_ull_doacross_guided_start, GOMP_loop_ull_doacross_runtime_start): New functions. (GOMP_loop_ull_doacross_static_start, GOMP_loop_ull_doacross_dynamic_start, GOMP_loop_ull_doacross_guided_start): New functions or aliases. * oacc-mem.c (acc_map_data, present_create_copy, gomp_acc_insert_pointer): Pass GOMP_MAP_VARS_OPENACC instead of false to gomp_map_vars. (gomp_acc_remove_pointer): Use copy_from from target_var_desc. * oacc-parallel.c (GOACC_data_start): Pass GOMP_MAP_VARS_OPENACC instead of false to gomp_map_vars. (GOACC_parallel_keyed): Likewise. Use copy_from from target_var_desc. * omp.h.in (omp_lock_hint_t): New type. (omp_init_lock_with_hint, omp_init_nest_lock_with_hint, omp_get_num_places, omp_get_place_num_procs, omp_get_place_proc_ids, omp_get_place_num, omp_get_partition_num_places, omp_get_partition_place_nums, omp_get_initial_device, omp_get_max_task_priority, 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): New prototypes. * omp_lib.f90.in (omp_lock_hint_kind): New parameter. (omp_lock_hint_none, omp_lock_hint_uncontended, omp_lock_hint_contended, omp_lock_hint_nonspeculative, omp_lock_hint_speculative): New parameters. (omp_init_lock_with_hint, omp_init_nest_lock_with_hint, omp_get_num_places, omp_get_place_num_procs, omp_get_place_proc_ids, omp_get_place_num, omp_get_partition_num_places, omp_get_partition_place_nums, omp_get_initial_device, omp_get_max_task_priority): New interfaces. (omp_set_schedule, omp_get_schedule): Rename modifier argument to chunk_size. * omp_lib.h.in (omp_lock_hint_kind): New parameter. (omp_lock_hint_none, omp_lock_hint_uncontended, omp_lock_hint_contended, omp_lock_hint_nonspeculative, omp_lock_hint_speculative): New parameters. (omp_init_lock_with_hint, omp_init_nest_lock_with_hint, omp_get_num_places, omp_get_place_num_procs, omp_get_place_proc_ids, omp_get_place_num, omp_get_partition_num_places, omp_get_partition_place_nums, omp_get_initial_device, omp_get_max_task_priority): New functions and subroutines. * ordered.c: Include stdarg.h and string.h. (MAX_COLLAPSED_BITS): Define. (gomp_doacross_init, GOMP_doacross_post, GOMP_doacross_wait, gomp_doacross_ull_init, GOMP_doacross_ull_post, GOMP_doacross_ull_wait): New functions. * target.c: Include errno.h. (resolve_device): If device is not initialized, call gomp_init_device on it. (gomp_map_lookup): New function. (gomp_map_vars_existing): Add tgt_var argument, fill it in. Don't bump refcount if REFCOUNT_INFINITY. Handle GOMP_MAP_ALWAYS_TO_P. (get_kind): Rename is_openacc argument to short_mapkind. (gomp_map_pointer): Use gomp_map_lookup. (gomp_map_fields_existing): New function. (gomp_map_vars): Rename is_openacc argument to short_mapkind and is_target to pragma_kind. Handle GOMP_MAP_VARS_ENTER_DATA, handle GOMP_MAP_FIRSTPRIVATE_INT, GOMP_MAP_STRUCT, GOMP_MAP_USE_DEVICE_PTR, GOMP_MAP_ZERO_LEN_ARRAY_SECTION. Adjust for tgt->list changed type and copy_from living in there. (gomp_copy_from_async): Adjust for tgt->list changed type and copy_from living in there. (gomp_unmap_vars): Likewise. (gomp_update): Likewise. Rename is_openacc argument to short_mapkind. Don't fail if object is not mapped. (gomp_load_image_to_device): Initialize refcount to REFCOUNT_INFINITY. (gomp_target_fallback): New function. (gomp_get_target_fn_addr): Likewise. (GOMP_target): Adjust gomp_map_vars caller, use gomp_get_target_fn_addr and gomp_target_fallback. (GOMP_target_41): New function. (gomp_target_data_fallback): New function. (GOMP_target_data): Use it, adjust gomp_map_vars caller. (GOMP_target_data_41): New function. (GOMP_target_update): Adjust gomp_update caller. (GOMP_target_update_41): New function. (gomp_exit_data, GOMP_target_enter_exit_data, gomp_target_task_fn, omp_target_alloc, omp_target_free, omp_target_is_present, omp_target_memcpy, omp_target_memcpy_rect_worker, omp_target_memcpy_rect, omp_target_associate_ptr, omp_target_disassociate_ptr, gomp_load_plugin_for_device): New functions. * task.c: Include gomp-constants.h. Include taskloop.c twice to get GOMP_taskloop and GOMP_taskloop_ull definitions. (gomp_task_handle_depend): New function. (GOMP_task): Use it. Add priority argument. Use gomp-constant.h constants instead of hardcoded numbers. Rename GOMP_TASK_IFFALSE to GOMP_TASK_UNDEFERRED. (gomp_create_target_task): New function. (verify_children_queue, verify_taskgroup_queue, verify_task_queue): New functions. (gomp_task_run_pre): Call verify_*_queue functions. If an upcoming tied task is about to leave the sibling or taskgroup queues in an invalid state, adjust appropriately. Remove taskgroup argument. Add comments. (gomp_task_run_post_handle_dependers): Add comments. (gomp_task_run_post_remove_parent): Likewise. (gomp_barrier_handle_tasks): Adjust gomp_task_run_pre caller. (GOMP_taskwait): Likewise. Add comments. (gomp_task_maybe_wait_for_dependencies): Fix scheduling problem such that the first non parent_depends_on task does not end up at the end of the children queue. (GOMP_taskgroup_start): Rename GOMP_TASK_IFFALSE to GOMP_TASK_UNDEFERRED. (GOMP_taskgroup_end): Adjust gomp_task_run_pre caller. * taskloop.c: New file. * testsuite/lib/libgomp.exp (check_effective_target_offload_device_nonshared_as): New proc. * testsuite/libgomp.c/affinity-2.c: New test. * testsuite/libgomp.c/doacross-1.c: New test. * testsuite/libgomp.c/doacross-2.c: New test. * testsuite/libgomp.c/examples-4/declare_target-1.c (fib_wrapper): Add map clause to target. * testsuite/libgomp.c/examples-4/declare_target-4.c (accum): Likewise. * testsuite/libgomp.c/examples-4/declare_target-5.c (accum): Likewise. * testsuite/libgomp.c/examples-4/device-1.c (main): Likewise. * testsuite/libgomp.c/examples-4/device-3.c (main): Likewise. * testsuite/libgomp.c/examples-4/target_data-3.c (gramSchmidt): Likewise. * testsuite/libgomp.c/examples-4/teams-2.c (dotprod): Likewise. * testsuite/libgomp.c/examples-4/teams-3.c (dotprod): Likewise. * testsuite/libgomp.c/examples-4/teams-4.c (dotprod): Likewise. * testsuite/libgomp.c/for-2.h (OMPTGT, OMPTO, OMPFROM): Define if not defined. Use those where needed. * testsuite/libgomp.c/for-4.c: New test. * testsuite/libgomp.c/for-5.c: New test. * testsuite/libgomp.c/for-6.c: New test. * testsuite/libgomp.c/linear-1.c: New test. * testsuite/libgomp.c/ordered-4.c: New test. * testsuite/libgomp.c/pr66199-2.c (f2): Adjust for linear clause only allowed on the loop iterator. * testsuite/libgomp.c/pr66199-3.c: New test. * testsuite/libgomp.c/pr66199-4.c: New test. * testsuite/libgomp.c/reduction-7.c: New test. * testsuite/libgomp.c/reduction-8.c: New test. * testsuite/libgomp.c/reduction-9.c: New test. * testsuite/libgomp.c/reduction-10.c: New test. * testsuite/libgomp.c/target-1.c (fn2, fn3, fn4): Add map(tofrom:s). * testsuite/libgomp.c/target-2.c (fn2, fn3, fn4): Likewise. * testsuite/libgomp.c/target-7.c (foo): Add map(h) where needed. * testsuite/libgomp.c/target-11.c: New test. * testsuite/libgomp.c/target-12.c: New test. * testsuite/libgomp.c/target-13.c: New test. * testsuite/libgomp.c/target-14.c: New test. * testsuite/libgomp.c/target-15.c: New test. * testsuite/libgomp.c/target-16.c: New test. * testsuite/libgomp.c/target-17.c: New test. * testsuite/libgomp.c/target-18.c: New test. * testsuite/libgomp.c/target-19.c: New test. * testsuite/libgomp.c/target-20.c: New test. * testsuite/libgomp.c/target-21.c: New test. * testsuite/libgomp.c/target-22.c: New test. * testsuite/libgomp.c/target-23.c: New test. * testsuite/libgomp.c/target-24.c: New test. * testsuite/libgomp.c/target-25.c: New test. * testsuite/libgomp.c/target-26.c: New test. * testsuite/libgomp.c/target-27.c: New test. * testsuite/libgomp.c/taskloop-1.c: New test. * testsuite/libgomp.c/taskloop-2.c: New test. * testsuite/libgomp.c/taskloop-3.c: New test. * testsuite/libgomp.c/taskloop-4.c: New test. * testsuite/libgomp.c++/ctor-13.C: New test. * testsuite/libgomp.c++/doacross-1.C: New test. * testsuite/libgomp.c++/examples-4/declare_target-2.C: Replace offload_device with offload_device_nonshared_as. * testsuite/libgomp.c++/for-12.C: New test. * testsuite/libgomp.c++/for-13.C: New test. * testsuite/libgomp.c++/for-14.C: New test. * testsuite/libgomp.c++/linear-1.C: New test. * testsuite/libgomp.c++/member-1.C: New test. * testsuite/libgomp.c++/member-2.C: New test. * testsuite/libgomp.c++/member-3.C: New test. * testsuite/libgomp.c++/member-4.C: New test. * testsuite/libgomp.c++/member-5.C: New test. * testsuite/libgomp.c++/ordered-1.C: New test. * testsuite/libgomp.c++/reduction-5.C: New test. * testsuite/libgomp.c++/reduction-6.C: New test. * testsuite/libgomp.c++/reduction-7.C: New test. * testsuite/libgomp.c++/reduction-8.C: New test. * testsuite/libgomp.c++/reduction-9.C: New test. * testsuite/libgomp.c++/reduction-10.C: New test. * testsuite/libgomp.c++/reference-1.C: New test. * testsuite/libgomp.c++/simd14.C: New test. * testsuite/libgomp.c++/target-2.C (fn2): Add map(tofrom: s) clause. * testsuite/libgomp.c++/target-5.C: New test. * testsuite/libgomp.c++/target-6.C: New test. * testsuite/libgomp.c++/target-7.C: New test. * testsuite/libgomp.c++/target-8.C: New test. * testsuite/libgomp.c++/target-9.C: New test. * testsuite/libgomp.c++/target-10.C: New test. * testsuite/libgomp.c++/target-11.C: New test. * testsuite/libgomp.c++/target-12.C: New test. * testsuite/libgomp.c++/taskloop-1.C: New test. * testsuite/libgomp.c++/taskloop-2.C: New test. * testsuite/libgomp.c++/taskloop-3.C: New test. * testsuite/libgomp.c++/taskloop-4.C: New test. * testsuite/libgomp.c++/taskloop-5.C: New test. * testsuite/libgomp.c++/taskloop-6.C: New test. * testsuite/libgomp.c++/taskloop-7.C: New test. * testsuite/libgomp.c++/taskloop-8.C: New test. * testsuite/libgomp.c++/taskloop-9.C: New test. * testsuite/libgomp.fortran/affinity1.f90: New test. * testsuite/libgomp.fortran/affinity2.f90: New test. liboffloadmic/ 2015-10-13 Ilya Verbin <ilya.verbin@intel.com> * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_dev2dev): New function. * plugin/offload_target_main.cpp (__offload_target_tgt2tgt): New static function, register it in liboffloadmic. From-SVN: r228777 |