mirror of git://gcc.gnu.org/git/gcc.git
672 Commits
| Author | SHA1 | Message | Date |
|---|---|---|---|
Paul-Antoine Arras
|
3560ef45c0 |
OpenMP: 'interop' construct - add ME support + target-independent libgomp
This patch partially enables use of the OpenMP interop construct by adding
middle end support, mostly in the omplower pass, and in the target-independent
part of the libgomp runtime. It follows up on previous patches for C, C++ and
Fortran front ends support. The full interop feature requires another patch to
enable foreign runtime support in libgomp plugins.
gcc/ChangeLog:
* builtin-types.def
(BT_FN_VOID_INT_INT_PTR_PTR_PTR_INT_PTR_INT_PTR_UINT_PTR): New.
* gimple-low.cc (lower_stmt): Handle GIMPLE_OMP_INTEROP.
* gimple-pretty-print.cc (dump_gimple_omp_interop): New function.
(pp_gimple_stmt_1): Handle GIMPLE_OMP_INTEROP.
* gimple.cc (gimple_build_omp_interop): New function.
(gimple_copy): Handle GIMPLE_OMP_INTEROP.
* gimple.def (GIMPLE_OMP_INTEROP): Define.
* gimple.h (gimple_build_omp_interop): Declare.
(gimple_omp_interop_clauses): New function.
(gimple_omp_interop_clauses_ptr): Likewise.
(gimple_omp_interop_set_clauses): Likewise.
(gimple_return_set_retval): Handle GIMPLE_OMP_INTEROP.
* gimplify.cc (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_INIT,
OMP_CLAUSE_USE and OMP_CLAUSE_DESTROY.
(gimplify_omp_interop): New function.
(gimplify_expr): Replace sorry with call to gimplify_omp_interop.
* omp-builtins.def (BUILT_IN_GOMP_INTEROP): Define.
* omp-low.cc (scan_sharing_clauses): Handle OMP_CLAUSE_INIT,
OMP_CLAUSE_USE and OMP_CLAUSE_DESTROY.
(scan_omp_1_stmt): Handle GIMPLE_OMP_INTEROP.
(lower_omp_interop_action_clauses): New function.
(lower_omp_interop): Likewise.
(lower_omp_1): Handle GIMPLE_OMP_INTEROP.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_clause_destroy): Make addressable.
(c_parser_omp_clause_init): Make addressable.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_clause_init): Make addressable.
gcc/fortran/ChangeLog:
* trans-openmp.cc (gfc_trans_omp_clauses): Make OMP_CLAUSE_DESTROY and
OMP_CLAUSE_INIT addressable.
* types.def (BT_FN_VOID_INT_INT_PTR_PTR_PTR_INT_PTR_INT_PTR_UINT_PTR):
New.
include/ChangeLog:
* gomp-constants.h (GOMP_DEVICE_DEFAULT_OMP_61, GOMP_INTEROP_TARGET,
GOMP_INTEROP_TARGETSYNC, GOMP_INTEROP_FLAG_NOWAIT): Define.
libgomp/ChangeLog:
* icv-device.c (omp_set_default_device): Check
GOMP_DEVICE_DEFAULT_OMP_61.
* libgomp-plugin.h (struct interop_obj_t): New.
(enum gomp_interop_flag): New.
(GOMP_OFFLOAD_interop): Declare.
(GOMP_OFFLOAD_get_interop_int): Declare.
(GOMP_OFFLOAD_get_interop_ptr): Declare.
(GOMP_OFFLOAD_get_interop_str): Declare.
(GOMP_OFFLOAD_get_interop_type_desc): Declare.
* libgomp.h (_LIBGOMP_OMP_LOCK_DEFINED): Define.
(struct gomp_device_descr): Add interop_func, get_interop_int_func,
get_interop_ptr_func, get_interop_str_func, get_interop_type_desc_func.
* libgomp.map: Add GOMP_interop.
* libgomp_g.h (GOMP_interop): Declare.
* target.c (resolve_device): Handle GOMP_DEVICE_DEFAULT_OMP_61.
(omp_get_interop_int): Replace stub with actual implementation.
(omp_get_interop_ptr): Likewise.
(omp_get_interop_str): Likewise.
(omp_get_interop_type_desc): Likewise.
(struct interop_data_t): Define.
(gomp_interop_internal): New function.
(GOMP_interop): Likewise.
(gomp_load_plugin_for_device): Load symbols for get_interop_int,
get_interop_ptr, get_interop_str and get_interop_type_desc.
* testsuite/libgomp.c-c++-common/interop-1.c: New test.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/interop-1.c: Remove dg-prune-output "sorry".
* c-c++-common/gomp/interop-2.c: Likewise.
* c-c++-common/gomp/interop-3.c: Likewise.
* c-c++-common/gomp/interop-4.c: Remove dg-message "not supported".
* g++.dg/gomp/interop-5.C: Likewise.
* gfortran.dg/gomp/interop-4.f90: Likewise.
* c-c++-common/gomp/interop-5.c: New test.
* gfortran.dg/gomp/interop-5.f90: New test.
Co-authored-by: Tobias Burnus <tburnus@baylibre.com>
(cherry picked from commit
|
|
|
Tobias Burnus
|
e34748375a |
OpenMP: 'interop' construct - add C/C++ parser support, improve Fortran parsing
Add middle end support for the 'interop' directive and the 'init', 'use',
and 'destroy' clauses - but fail with a sorry, unimplemented in gimplify.cc.
For Fortran, generate the tree code, update the internal representation,
add some more diagnostic checks and update for newer specification changes
('fr' only takes a single value, but it integer expressions are permitted
again [like with the old syntax] not only constant identifiers).
For C and C++, this patch adds the full parser support for 'interop'.
Still missing is actually handling the directive in the middle end and
in libgomp.
The GOMP_INTEROP_IFR_* internal values have been changed to have space
for vendor specific values that are adjacent to the existing values
but negative, if needed.
gcc/c-family/ChangeLog:
* c-common.h (enum c_omp_region_type): Add C_ORT_INTEROP
and C_ORT_OMP_INTEROP.
(c_omp_interop_t_p): New prototype.
* c-omp.cc (c_omp_interop_t_p): Check whether the type is
omp_interop_t.
(c_omp_directives): Uncomment 'interop'.
* c-pragma.cc (omp_pragmas): Add 'interop'.
* c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_INTEROP.
(enum pragma_omp_clause): Add init, use, and destroy clauses.
gcc/c/ChangeLog:
* c-parser.cc (INCLUDE_STRING): Define.
(c_parser_pragma): Handle 'interop' directive.
(c_parser_omp_clause_name): Handle init, use, and destroy clauses.
(c_parser_omp_all_clauses): Likewise; use C_ORT_OMP_INTEROP, if
'use' is permitted, for c_finish_omp_clauses.
(c_parser_omp_clause_destroy, c_parser_omp_modifier_prefer_type,
c_parser_omp_clause_init, c_parser_omp_clause_use,
OMP_INTEROP_CLAUSE_MASK, c_parser_omp_interop): New.
* c-typeck.cc (c_finish_omp_clauses): Add missing OPT_Wopenmp to
a warning; handle new clauses.
gcc/cp/ChangeLog:
* parser.cc (INCLUDE_STRING): Define.
(cp_parser_omp_clause_name): Handle init, use, and destroy clauses.
(cp_parser_omp_all_clauses): Likewise; use C_ORT_OMP_INTEROP, if
'use' is permitted, for c_finish_omp_clauses.
(cp_parser_omp_modifier_prefer_type, cp_parser_omp_clause_init,
OMP_INTEROP_CLAUSE_MASK, cp_parser_omp_interop): New.
(cp_parser_pragma): Handle 'interop' directive.
* pt.cc (tsubst_omp_clauses): Handle init, use, and destroy clauses.
(tsubst_stmt): Handle OMP_INTEROP.
* semantics.cc (cp_omp_init_prefer_type_update): New.
(finish_omp_clauses): Handle init, use, and destroy clauses
and add clause check for 'depend' on 'interop'.
gcc/fortran/ChangeLog:
* gfortran.h (gfc_omp_namelist): Cleanup interop internal
representation.
* dump-parse-tree.cc (show_omp_namelist): Update for changed
internal representation.
* match.cc (gfc_free_omp_namelist): Likewise.
* openmp.cc (gfc_match_omp_prefer_type, gfc_match_omp_init):
Likewise; also handle some corner cases better and update for
newer 6.0 changes related to 'fr'.
(resolve_omp_clauses): Add type-check for interop variables.
* trans-openmp.cc (gfc_trans_omp_clauses): Handle init, use
and destroy clauses.
(gfc_trans_openmp_interop): New.
(gfc_trans_omp_directive): Call it.
gcc/ChangeLog:
* gimplify.cc (gimplify_expr): Handle OMP_INTEROP by printing
"sorry, uninplemented".
* omp-api.h (omp_get_fr_id_from_name): Change return type to
'char'.
* omp-general.cc (omp_get_fr_id_from_name): Likewise; return
GOMP_INTEROP_IFR_UNKNOWN not 0 if not found.
(omp_get_name_from_fr_id): Return "<unknown>" not NULL
if not found (used for dumps).
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_DESTROY,
OMP_CLAUSE_USE, and OMP_CLAUSE_INIT.
* tree-pretty-print.cc (dump_omp_init_prefer_type): New.
(dump_omp_clause): Handle init, use and destroy clauses.
(dump_generic_node): Handle interop directive.
* tree.cc (omp_clause_num_ops, omp_clause_code_name): Add new
init/use/destroy clauses.
* tree.def (OACC_LOOP): Fix comment.
(OMP_INTEROP): Add.
* tree.h (OMP_INTEROP_CLAUSES, OMP_CLAUSE_INIT_TARGET,
OMP_CLAUSE_INIT_TARGETSYNC, OMP_CLAUSE_INIT_PREFER_TYPE): New.
include/ChangeLog:
* gomp-constants.h (GOMP_INTEROP_IFR_NONE): Rename ...
(GOMP_INTEROP_IFR_UNKNOWN): ... to this. And change value.
(GOMP_INTEROP_IFR_SEPARATOR): Likewise.
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/interop-1.f90: Update for parser changes,
spec changes and add new tests.
* gfortran.dg/gomp/interop-2.f90: Likewise.
* gfortran.dg/gomp/interop-3.f90: Likewise.
* c-c++-common/gomp/interop-1.c: New test.
* c-c++-common/gomp/interop-2.c: New test.
* c-c++-common/gomp/interop-3.c: New test.
* c-c++-common/gomp/interop-4.c: New test.
* g++.dg/gomp/interop-5.C: New test.
* gfortran.dg/gomp/interop-4.f90: New test.
(cherry picked from commit
|
|
|
Tobias Burnus
|
66a6522caf |
OpenMP: Add get_device_from_uid/omp_get_uid_from_device routines
Those TR13/OpenMP 6.0 routines permit a reproducible offloading to
a specific device by mapping an OpenMP device number to a
unique ID (UID). The GPU device UIDs should be universally unique,
the one for the host is not.
gcc/ChangeLog:
* omp-general.cc (omp_runtime_api_procname): Add
get_device_from_uid and omp_get_uid_from_device routines.
include/ChangeLog:
* cuda/cuda.h (cuDeviceGetUuid): Declare.
(cuDeviceGetUuid_v2): Add prototype.
libgomp/ChangeLog:
* config/gcn/target.c (omp_get_uid_from_device,
omp_get_device_from_uid): Add stub implementation.
* config/nvptx/target.c (omp_get_uid_from_device,
omp_get_device_from_uid): Likewise.
* fortran.c (omp_get_uid_from_device_,
omp_get_uid_from_device_8_): New functions.
* libgomp-plugin.h (GOMP_OFFLOAD_get_uid): Add prototype.
* libgomp.h (struct gomp_device_descr): Add 'uid' and 'get_uid_func'.
* libgomp.map (GOMP_6.0): New, includind the new UID routines.
* libgomp.texi (OpenMP Technical Report 13): Mark UID routines as 'Y'.
(Device Information Routines): Document new UID routines.
(Offload-Target Specifics): Document UID format.
* omp.h.in (omp_get_device_from_uid, omp_get_uid_from_device):
New prototype.
* omp_lib.f90.in (omp_get_device_from_uid, omp_get_uid_from_device):
New interface.
* omp_lib.h.in: Likewise.
* plugin/cuda-lib.def: Add cuDeviceGetUuid and cuDeviceGetUuid_v2 via
CUDA_ONE_CALL_MAYBE_NULL.
* plugin/plugin-gcn.c (GOMP_OFFLOAD_get_uid): New.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_uid): New.
* target.c (str_omp_initial_device): New static var.
(STR_OMP_DEV_PREFIX): Define.
(gomp_get_uid_for_device, omp_get_uid_from_device,
omp_get_device_from_uid): New.
(gomp_load_plugin_for_device): DLSYM_OPT the function 'get_uid'.
(gomp_target_init): Set the device's 'uid' field to NULL.
* testsuite/libgomp.c/device_uid.c: New test.
* testsuite/libgomp.fortran/device_uid.f90: New test.
(cherry picked from commit
|
|
|
Tobias Burnus
|
2b5d2e5731 |
Fortran: Fixes to OpenMP 'interop' directive parsing support
Handle lists as argument to 'fr' and 'attr'; fix parsing corner cases.
Additionally, 'fr' values are now internally stored as integer, permitting
the diagnoses (warning) for values not defined in the OpenMP additional
definitions document.
PR fortran/116661
gcc/fortran/ChangeLog:
* gfortran.h (gfc_omp_namelist): Rename 'init' members for clarity.
* match.cc (gfc_free_omp_namelist): Handle renaming.
* dump-parse-tree.cc (show_omp_namelist): Update for new format
and features.
* openmp.cc (gfc_match_omp_prefer_type): Parse list to 'fr' and 'attr';
store 'fr' values as integer.
(gfc_match_omp_init): Rename variable names.
gcc/ChangeLog:
* omp-api.h (omp_get_fr_id_from_name, omp_get_name_from_fr_id): New
prototypes.
* omp-general.cc (omp_get_fr_id_from_name, omp_get_name_from_fr_id):
New.
include/ChangeLog:
* gomp-constants.h (GOMP_INTEROP_IFR_LAST,
GOMP_INTEROP_IFR_SEPARATOR, GOMP_INTEROP_IFR_NONE): New.
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/interop-1.f90: Extend, update dg-*.
* gfortran.dg/gomp/interop-2.f90: Update dg-error.
* gfortran.dg/gomp/interop-3.f90: Add dg-warning.
(cherry picked from commit
|
|
|
Tobias Burnus
|
84efccbc61 |
Merge branch 'releases/gcc-14' into devel/omp/gcc-14
Merge up to r14-10526-g04696df09633ba (1st August 2024). That's the 'GCC 14.2.0 released.' commit. |
|
|
Jakub Jelinek
|
04696df096 | Update ChangeLog and version files for release | |
|
Sandra Loosemore
|
85c55edc09 |
libgomp: runtime support for target_device selector
This patch implements the libgomp runtime support for the dynamic target_device selector via the GOMP_evaluate_target_device function. include/ChangeLog * cuda/cuda.h (CUdevice_attribute): Add definitions for CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR and CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR. libgomp/ChangeLog * Makefile.am (libgomp_la_SOURCES): Add selector.c. * Makefile.in: Regenerate. * config/gcn/selector.c: New. * config/linux/selector.c: New. * config/linux/x86/selector.c: New. * config/nvptx/selector.c: New. * libgomp-plugin.h (GOMP_OFFLOAD_evaluate_device): New. * libgomp.h (struct gomp_device_descr): Add evaluate_device_func field. * libgomp.map (GOMP_5.1.3): New, add GOMP_evaluate_target_device. * libgomp.texi (OpenMP Context Selectors): Document dynamic selector matching of kind/arch/isa. * libgomp_g.h (GOMP_evaluate_current_device): New. (GOMP_evaluate_target_device): New. * oacc-host.c (host_evaluate_device): New. (host_openacc_exec): Initialize evaluate_device_func field to host_evaluate_device. * plugin/plugin-gcn.c (gomp_match_selectors): New. (gomp_match_isa): New. (GOMP_OFFLOAD_evaluate_device): New. * plugin/plugin-nvptx.c (struct ptx_device): Add compute_major and compute_minor fields. (nvptx_open_device): Read compute capability information from device. (gomp_match_selectors): New. (gomp_match_selector): New. (CHECK_ISA): New macro. (GOMP_OFFLOAD_evaluate_device): New. * selector.c: New. * target.c (GOMP_evaluate_target_device): New. (gomp_load_plugin_for_device): Load evaluate_device plugin function. Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com> Co-Authored-By: Sandra Loosemore <sandra@codesourcery.com> |
|
|
Julian Brown
|
3d5bcb1a84 |
OpenMP: Support strided and shaped-array updates for C++
This patch adds support for OpenMP 5.0 strided updates and the
array-shaping operator ("([x][y][z]) foo[0:n]..."). This is mostly for
C++ only so far, though necessary changes have been made to the C FE to
adjust for changes to shared data structures.
In terms of the implementation of various bits:
- The OMP_ARRAY_SECTION tree code has been extended to take a 'stride'
argument, and changes have been made throughout semantics.cc, etc. to
take the new field into account -- including bounds checking.
- A new type of cast operator has been added to represent the OpenMP
array-shaping operator: OMP_ARRAYSHAPE_CAST_EXPR (1).
- The address tokenization mechanism from previous patches has been
extended with two new access kinds to represent noncontiguous array
updates.
- New mapping kinds have been added to represent noncontiguous updates:
those which may be subject to array shaping, or have non-unit strides.
These are processed by omp-low.cc into a kind of descriptor that is
passed to the libgomp runtime (2).
The current patch reuses an extended version of the helper code for
omp_target_memcpy_rect, which may generate very many small host-device or
device-host copies. (The "descriptor" has also been designed so reusing
that functionality is relatively straightforward.) Optimising those
multiple copies, e.g. by packing them into a single transfer when it
would be beneficial, is left as the subject of a future patch.
This patch has some adjustments to the omp-low.cc code after Chung-Lin's
patch "OpenMP 5.0: Allow multiple clauses mapping same variable"
(
|
|
|
Julian Brown
|
015cb4002d |
OpenMP: Fortran "!$omp declare mapper" support
This patch implements "omp declare mapper" functionality for Fortran,
following the equivalent support for C and C++. This version of the
patch has been merged to og13 and contains various fixes for e.g.:
* Mappers with deferred-length strings
* Array descriptors not being appropriately transferred
to the offload target (see "OMP_MAP_POINTER_ONLY" and
gimplify.cc:omp_maybe_get_descriptor_from_ptr).
2023-06-30 Julian Brown <julian@codesourcery.com>
gcc/fortran/
* dump-parse-tree.cc (show_attr): Show omp_udm_artificial_var flag.
(show_omp_namelist): Support OMP_MAP_POINTER_ONLY and OMP_MAP_UNSET.
* f95-lang.cc (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES,
LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE,
LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define language hooks.
* gfortran.h (gfc_statement): Add ST_OMP_DECLARE_MAPPER.
(symbol_attribute): Add omp_udm_artificial_var attribute.
(gfc_omp_map_op): Add OMP_MAP_POINTER_ONLY and OMP_MAP_UNSET.
(gfc_omp_namelist): Add udm pointer to u2 union.
(gfc_omp_udm): New struct.
(gfc_omp_namelist_udm): New struct.
(gfc_symtree): Add omp_udm pointer.
(gfc_namespace): Add omp_udm_root symtree. Add omp_udm_ns flag.
(gfc_free_omp_namelist): Update prototype.
(gfc_free_omp_udm, gfc_omp_udm_find, gfc_find_omp_udm,
gfc_resolve_omp_udms): Add prototypes.
* match.cc (gfc_free_omp_namelist): Change FREE_NS and FREE_ALIGN
parameters to LIST number, to handle freeing user-defined mapper
namelists safely.
* match.h (gfc_match_omp_declare_mapper): Add prototype.
* module.cc (ab_attribute): Add AB_OMP_DECLARE_MAPPER_VAR.
(attr_bits): Add OMP_DECLARE_MAPPER_VAR.
(mio_symbol_attribute): Read/write AB_OMP_DECLARE_MAPPER_VAR attribute.
Set referenced attr on read.
(omp_map_clause_ops, omp_map_cardinality): New arrays.
(load_omp_udms, check_omp_declare_mappers): New functions.
(read_module): Load and check OMP declare mappers.
(write_omp_udm, write_omp_udms): New functions.
(write_module): Write OMP declare mappers.
* openmp.cc (gfc_free_omp_clauses, gfc_match_omp_variable_list,
gfc_match_omp_to_link, gfc_match_omp_depend_sink,
gfc_match_omp_clause_reduction): Update calls to gfc_free_omp_namelist.
(gfc_free_omp_udm, gfc_find_omp_udm, gfc_omp_udm_find,
gfc_match_omp_declare_mapper): New functions.
(gfc_match_omp_clauses): Add DEFAULT_MAP_OP parameter. Update calls to
gfc_free_omp_namelist. Add declare mapper support.
(resolve_omp_clauses): Add declare mapper support. Update calls to
gfc_free_omp_namelist.
(gfc_resolve_omp_udm, gfc_resolve_omp_udms): New functions.
* parse.cc (decode_omp_directive): Add declare mapper support.
(case_omp_decl): Add ST_OMP_DECLARE_MAPPER case.
(gfc_ascii_statement): Add ST_OMP_DECLARE_MAPPER case.
* resolve.cc (resolve_types): Call gfc_resolve_omp_udms.
* st.cc (gfc_free_statement): Update call to gfc_free_omp_namelist.
* symbol.cc (free_omp_udm_tree): New function.
(gfc_free_namespace): Call above.
* trans-decl.cc (omp_declare_mapper_ns): New global.
(gfc_finish_var_decl, gfc_generate_function_code): Support declare
mappers.
(gfc_trans_deferred_vars): Ignore artificial declare-mapper vars.
* trans-openmp.cc (tree-iterator.h): Include.
(toc_directive): New enum.
(gfc_trans_omp_array_section): Change OP and OPENMP parameters to
toc_directive CD ('clause directive').
(gfc_omp_finish_mapper_clauses, gfc_omp_extract_mapper_directive,
gfc_omp_map_array_section): New functions.
(omp_clause_directive): New enum.
(gfc_trans_omp_clauses): Remove DECLARE_SIMD and OPENACC parameters.
Replace with toc_directive CD, defaulting to TOC_OPENMP. Add declare
mapper support and OMP_MAP_POINTER_ONLY support.
(gfc_trans_omp_construct, gfc_trans_oacc_executable_directive,
gfc_trans_oacc_combined_directive): Update calls to
gfc_trans_omp_clauses.
(gfc_subst_replace, gfc_subst_prepend_ref): New variables.
(gfc_subst_in_expr_1, gfc_subst_in_expr, gfc_subst_mapper_var,
gfc_trans_omp_instantiate_mapper, gfc_trans_omp_instantiate_mappers,
gfc_record_mapper_bindings_code_fn, gfc_record_mapper_bindings_expr_fn,
gfc_find_nested_mappers, gfc_record_mapper_bindings): New functions.
(gfc_typespec * hash traits): New template.
(omp_declare_mapper_ns): Extern declaration.
(gfc_trans_omp_target): Call gfc_trans_omp_instantiate_mappers and
gfc_record_mapper_bindings. Update calls to gfc_trans_omp_clauses.
(gfc_trans_omp_declare_simd, gfc_trans_omp_declare_variant): Update
calls to gfc_trans_omp_clauses.
(gfc_trans_omp_mapper_name, gfc_trans_omp_declare_mapper,
gfc_trans_omp_declare_mappers): New functions.
* trans-stmt.h (gfc_trans_omp_declare_mappers): Add prototype.
* trans.h (gfc_omp_finish_mapper_clauses,
gfc_omp_extract_mapper_directive, gfc_omp_map_array_section): Add
prototypes.
gcc/
* gimplify.cc (dwarf2out.h): Include.
(omp_maybe_get_descriptor_from_ptr): New function.
(build_omp_struct_comp_nodes): Use above function to locate array
descriptor when necessary.
(omp_mapping_group_data, omp_mapping_group_ptr,
omp_mapping_group_pset): New functions.
(omp_instantiate_mapper): Handle inlining of "declare mapper" function
bodies containing setup code (e.g. for Fortran). Handle pointers to
derived types. Handle GOMP_MAP_MAPPING_GROUPs.
* tree-pretty-print.cc (dump_omp_clause): Handle
GOMP_MAP_MAPPING_GROUP.
include/
* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_MAPPING_GROUP.
gcc/testsuite/
* gfortran.dg/gomp/declare-mapper-1.f90: New test.
* gfortran.dg/gomp/declare-mapper-5.f90: New test.
* gfortran.dg/gomp/declare-mapper-14.f90: New test.
libgomp/
* testsuite/libgomp.fortran/declare-mapper-2.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-3.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-4.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-6.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-7.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-8.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-9.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-10.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-11.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-12.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-13.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-15.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-17.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-18.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-19.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-20.f90: New test.
* testsuite/libgomp.fortran/declare-mapper-21.f90: New test.
|
|
|
Julian Brown
|
b1536b2e60 |
OpenMP: C++ "declare mapper" support
This patch adds support for OpenMP 5.0 "declare mapper" functionality
for C++. I've merged it to og13 based on the last version
posted upstream, with some minor changes due to the newly-added
'present' map modifier support. There's also a fix to splay-tree
traversal in gimplify.cc:omp_instantiate_implicit_mappers, and this patch
omits the rearrangement of gimplify.cc:gimplify_{scan,adjust}_omp_clauses
that I separated out into its own patch and applied (to og13) already.
2023-06-30 Julian Brown <julian@codesourcery.com>
gcc/c-family/
* c-common.h (omp_mapper_list): Add forward declaration.
(c_omp_find_nested_mappers, c_omp_instantiate_mappers): Add prototypes.
* c-omp.cc (c_omp_find_nested_mappers): New function.
(remap_mapper_decl_info): New struct.
(remap_mapper_decl_1, omp_instantiate_mapper,
c_omp_instantiate_mappers): New functions.
gcc/cp/
* constexpr.cc (reduced_constant_expression_p): Add OMP_DECLARE_MAPPER
case.
(cxx_eval_constant_expression, potential_constant_expression_1):
Likewise.
* cp-gimplify.cc (cxx_omp_finish_mapper_clauses): New function.
* cp-objcp-common.h (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES,
LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE,
LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define langhooks.
* cp-tree.h (lang_decl_base): Add omp_declare_mapper_p field. Recount
spare bits comment.
(DECL_OMP_DECLARE_MAPPER_P): New macro.
(omp_mapper_id, cp_check_omp_declare_mapper, omp_instantiate_mappers,
cxx_omp_finish_mapper_clauses, cxx_omp_mapper_lookup,
cxx_omp_extract_mapper_directive, cxx_omp_map_array_section: Add
prototypes.
* decl.cc (check_initializer): Add OpenMP declare mapper support.
(cp_finish_decl): Set DECL_INITIAL for OpenMP declare mapper var decls
as appropriate.
* decl2.cc (mark_used): Instantiate OpenMP "declare mapper" magic var
decls.
* error.cc (dump_omp_declare_mapper): New function.
(dump_simple_decl): Use above.
* parser.cc (cp_parser_omp_clause_map): Add KIND parameter. Support
"mapper" modifier.
(cp_parser_omp_all_clauses): Add KIND argument to
cp_parser_omp_clause_map call.
(cp_parser_omp_target): Call omp_instantiate_mappers before
finish_omp_clauses.
(cp_parser_omp_declare_mapper): New function.
(cp_parser_omp_declare): Add "declare mapper" support.
* pt.cc (tsubst_decl): Adjust name of "declare mapper" magic var decls
once we know their type.
(tsubst_omp_clauses): Call omp_instantiate_mappers before
finish_omp_clauses, for target regions.
(tsubst_expr): Support OMP_DECLARE_MAPPER nodes.
(instantiate_decl): Instantiate initialiser (i.e definition) for OpenMP
declare mappers.
* semantics.cc (gimplify.h): Include.
(omp_mapper_id, omp_mapper_lookup, omp_extract_mapper_directive,
cxx_omp_map_array_section, cp_check_omp_declare_mapper): New functions.
(finish_omp_clauses): Delete GOMP_MAP_PUSH_MAPPER_NAME and
GOMP_MAP_POP_MAPPER_NAME artificial clauses.
(omp_target_walk_data): Add MAPPERS field.
(finish_omp_target_clauses_r): Scan for uses of struct/union/class type
variables.
(finish_omp_target_clauses): Create artificial mapper binding clauses
for used structs/unions/classes in offload region.
gcc/fortran/
* parse.cc (tree.h, fold-const.h, tree-hash-traits.h): Add includes
(for additions to omp-general.h).
gcc/
* gimplify.cc (gimplify_omp_ctx): Add IMPLICIT_MAPPERS field.
(new_omp_context): Initialise IMPLICIT_MAPPERS hash map.
(delete_omp_context): Delete IMPLICIT_MAPPERS hash map.
(instantiate_mapper_info): New structs.
(remap_mapper_decl_1, omp_mapper_copy_decl, omp_instantiate_mapper,
omp_instantiate_implicit_mappers): New functions.
(gimplify_scan_omp_clauses): Handle MAPPER_BINDING clauses.
(gimplify_adjust_omp_clauses): Instantiate implicit declared mappers.
(gimplify_omp_declare_mapper): New function.
(gimplify_expr): Call above function.
* langhooks-def.h (lhd_omp_finish_mapper_clauses,
lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive,
lhd_omp_map_array_section): Add prototypes.
(LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES,
LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE,
LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define macros.
(LANG_HOOK_DECLS): Add above macros.
* langhooks.cc (lhd_omp_finish_mapper_clauses,
lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive,
lhd_omp_map_array_section): New dummy functions.
* langhooks.h (lang_hooks_for_decls): Add OMP_FINISH_MAPPER_CLAUSES,
OMP_MAPPER_LOOKUP, OMP_EXTRACT_MAPPER_DIRECTIVE, OMP_MAP_ARRAY_SECTION
hooks.
* omp-general.h (omp_name_type<T>): Add templatized struct, hash type
traits (for omp_name_type<tree> specialization).
(omp_mapper_list<T>): Add struct.
* tree-core.h (omp_clause_code): Add OMP_CLAUSE__MAPPER_BINDING_.
* tree-pretty-print.cc (dump_omp_clause): Support GOMP_MAP_UNSET,
GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping
clauses. Support OMP_CLAUSE__MAPPER_BINDING_ and OMP_DECLARE_MAPPER.
* tree.cc (omp_clause_num_ops, omp_clause_code_name): Add
OMP_CLAUSE__MAPPER_BINDING_.
* tree.def (OMP_DECLARE_MAPPER): New tree code.
* tree.h (OMP_DECLARE_MAPPER_ID, OMP_DECLARE_MAPPER_DECL,
OMP_DECLARE_MAPPER_CLAUSES): New defines.
(OMP_CLAUSE__MAPPER_BINDING__ID, OMP_CLAUSE__MAPPER_BINDING__DECL,
OMP_CLAUSE__MAPPER_BINDING__MAPPER): New defines.
include/
* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_UNSET,
GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping
clause types.
gcc/testsuite/
* c-c++-common/gomp/map-6.c: Update error scan output.
* c-c++-common/gomp/declare-mapper-3.c: New test (only enabled for C++
for now).
* c-c++-common/gomp/declare-mapper-4.c: Likewise.
* c-c++-common/gomp/declare-mapper-5.c: Likewise.
* c-c++-common/gomp/declare-mapper-6.c: Likewise.
* c-c++-common/gomp/declare-mapper-7.c: Likewise.
* c-c++-common/gomp/declare-mapper-8.c: Likewise.
* c-c++-common/gomp/declare-mapper-9.c: Likewise.
* c-c++-common/gomp/declare-mapper-12.c: Likewise.
* g++.dg/gomp/declare-mapper-1.C: New test.
* g++.dg/gomp/declare-mapper-2.C: New test.
libgomp/
* testsuite/libgomp.c++/declare-mapper-1.C: New test.
* testsuite/libgomp.c++/declare-mapper-2.C: New test.
* testsuite/libgomp.c++/declare-mapper-3.C: New test.
* testsuite/libgomp.c++/declare-mapper-4.C: New test.
* testsuite/libgomp.c++/declare-mapper-5.C: New test.
* testsuite/libgomp.c++/declare-mapper-6.C: New test.
* testsuite/libgomp.c++/declare-mapper-7.C: New test.
* testsuite/libgomp.c++/declare-mapper-8.C: New test.
* testsuite/libgomp.c-c++-common/declare-mapper-9.c: New test (only
enabled for C++ for now).
* testsuite/libgomp.c-c++-common/declare-mapper-10.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-11.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-12.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-13.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-14.c: Likewise.
|
|
|
Julian Brown
|
b918a7e4b4 |
OpenACC: Improve implicit mapping for non-lexically nested offload regions
This patch enables use of the OMP_CLAUSE_RUNTIME_IMPLICIT_P flag for
OpenACC.
This allows code like this to work correctly:
int arr[100];
[...]
#pragma acc enter data copyin(arr[20:10])
/* No explicit mapping of 'arr' here. */
#pragma acc parallel
{ /* use of arr[20:10]... */ }
#pragma acc exit data copyout(arr[20:10])
Otherwise, the implicit "copy" ("present_or_copy") on the parallel
corresponds to the whole array, and that fails at runtime when the
subarray is mapped.
The numbering of the GOMP_MAP_IMPLICIT bit clashes with the OpenACC
"non-contiguous" dynamic array support, so the GOMP_MAP_NONCONTIG_ARRAY_P
macro has been adjusted to account for that.
This behaviour relates to upstream OpenACC issue 490 (not yet resolved).
2023-06-16 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.cc (gimplify_adjust_omp_clauses_1): Set
OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P for OpenACC also.
gcc/testsuite/
* c-c++-common/goacc/combined-reduction.c: Adjust scan output.
* c-c++-common/goacc/reduction-1.c: Likewise.
* c-c++-common/goacc/reduction-2.c: Likewise.
* c-c++-common/goacc/reduction-3.c: Likewise.
* c-c++-common/goacc/reduction-4.c: Likewise.
* c-c++-common/goacc/reduction-10.c: Likewise.
* gfortran.dg/goacc/loop-tree-1.f90: Likewise.
include/
* gomp-constants.h (GOMP_MAP_NONCONTIG_ARRAY_P): Tweak condition.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c: New test.
|
|
|
Kwok Cheung Yeung
|
209d374b10 |
libgomp, nvptx: Update bundled CUDA header file
This updates the bundled cuda.h header file to include some new API calls and constants that are now used in the code. This patch should be included when the "libgomp, nvptx: low-latency memory allocator" or "openmp: Add support for 'target_device' context selector set" patches are upstreamed. 2022-06-21 Kwok Cheung Yeung <kcy@codesourcery.com> include/ * cuda/cuda.h (CUdevice_attribute): Add definitions for CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR and CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR. (CUmemAttach_flags): New. (CUpointer_attribute): New. (cuMemAllocManaged): New prototype. (cuPointerGetAttribute): New prototype. libgomp/ * plugin/cuda-lib.def (cuMemAllocManaged): Add new call. (cuPointerGetAttribute): Likewise. |
|
|
Julian Brown
|
65be1389ee |
Fortran "declare create"/allocate support for OpenACC
2018-10-04 Cesar Philippidis <cesar@codesourcery.com>
Julian Brown <julian@codesourcery.com>
gcc/
* omp-low.cc (scan_sharing_clauses): Update handling of OpenACC declare
create, declare copyin and declare deviceptr to have local lifetimes.
(convert_to_firstprivate_int): Handle pointer types.
(convert_from_firstprivate_int): Likewise. Create local storage for
the values being pointed to. Add new orig_type argument.
(lower_omp_target): Handle GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}.
Add orig_type argument to convert_from_firstprivate_int call.
Allow pointer types with GOMP_MAP_FIRSTPRIVATE_INT. Don't privatize
firstprivate VLAs.
* tree-pretty-print.cc (dump_omp_clause): Handle
GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}.
gcc/fortran/
* gfortran.h (enum gfc_omp_map_op): Add OMP_MAP_DECLARE_ALLOCATE,
OMP_MAP_DECLARE_DEALLOCATE.
(gfc_omp_clauses): Add update_allocatable.
* trans-array.cc (gfc_array_allocate): Call
gfc_trans_oacc_declare_allocate for decls that have oacc_declare_create
attribute set.
* trans-decl.cc (find_module_oacc_declare_clauses): Relax
oacc_declare_create to OMP_MAP_ALLOC, and oacc_declare_copyin to
OMP_MAP_TO, in order to match OpenACC 2.5 semantics.
* trans-openmp.cc (gfc_trans_omp_clauses): Use GOMP_MAP_ALWAYS_POINTER
(for update directive) or GOMP_MAP_FIRSTPRIVATE_POINTER (otherwise) for
allocatable scalar decls. Handle OMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}
clauses.
(gfc_trans_oacc_executable_directive): Use GOMP_MAP_ALWAYS_POINTER
for allocatable scalar data clauses inside acc update directives.
(gfc_trans_oacc_declare_allocate): New function.
* trans-stmt.cc (gfc_trans_allocate): Call
gfc_trans_oacc_declare_allocate for decls with oacc_declare_create
attribute set.
(gfc_trans_deallocate): Likewise.
* trans.h (gfc_trans_oacc_declare_allocate): Declare.
gcc/testsuite/
* gfortran.dg/goacc/declare-allocatable-1.f90: New test.
include/
* gomp-constants.h (enum gomp_map_kind): Define
GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE} and GOMP_MAP_FLAG_SPECIAL_4.
libgomp/
* oacc-mem.c (gomp_acc_declare_allocate): New function.
* oacc-parallel.c (GOACC_enter_exit_data): Handle
GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE}.
* testsuite/libgomp.oacc-fortran/allocatable-scalar.f90: New test.
* testsuite/libgomp.oacc-fortran/declare-allocatable-2.f90: New test.
* testsuite/libgomp.oacc-fortran/declare-allocatable-3.f90: New test.
* testsuite/libgomp.oacc-fortran/declare-allocatable-4.f90: New test.
2020-02-19 Julian Brown <julian@codesourcery.com>
gcc/fortran/
* trans-openmp.cc (gfc_omp_check_optional_argument): Handle non-decl
case.
gcc/
* gimplify.cc (gimplify_scan_omp_clauses): Handle
GOMP_MAP_DECLARE_ALLOCATE and GOMP_MAP_DECLARE_DEALLOCATE.
libgomp/
* libgomp.h (gomp_acc_declare_allocate): Remove prototype.
* oacc-mem.c (gomp_acc_declare_allocate): Make static. Add POINTER
argument. Use acc_delete instead of acc_free. Handle scalar
mappings.
(find_group_last): Handle GOMP_MAP_DECLARE_ALLOCATE and
GOMP_MAP_DECLARE_DEALLOCATE groupings.
(goacc_enter_data_internal): Fix kind check for
GOMP_MAP_DECLARE_ALLOCATE. Pass new pointer argument to
gomp_acc_declare_allocate.
(goacc_exit_data_internal): Unlock device mutex around
gomp_acc_declare_allocate call. Pass new pointer argument. Handle
group pointer mapping for deallocate.
2021-04-07 Kwok Cheung Yeung <kcy@codesourcery.com>
libgomp/
* oacc-mem.c (goacc_enter_data_internal): Unlock mutex before calling
gomp_acc_declare_allocate and relock it afterwards.
2023-04-17 Kwok Cheung Yeung <kcy@codesourcery.com>
* gimplify.cc (omp_group_base): Handle GOMP_MAP_DECLARE_ALLOCATE
and GOMP_MAP_DECLARE_DEALLOCATE.
|
|
|
Chung-Lin Tang
|
b143c1c447 |
Merge non-contiguous array support patches.
This version is based from v4, posted upstream here: https://gcc.gnu.org/pipermail/gcc-patches/2020-April/543437.html 2020-04-19 Chung-Lin Tang <cltang@codesourcery.com> PR other/76739 gcc/c/ * c-typeck.cc (handle_omp_array_sections_1): Add 'bool &non_contiguous' parameter, adjust recursive call site, add cases for allowing pointer based multi-dimensional arrays for OpenACC. (handle_omp_array_sections): Adjust handle_omp_array_sections_1 call, handle non-contiguous case to create dynamic array map. gcc/cp/ * semantics.cc (handle_omp_array_sections_1): Add 'bool &non_contiguous' parameter, adjust recursive call site, add cases for allowing pointer based multi-dimensional arrays for OpenACC. (handle_omp_array_sections): Adjust handle_omp_array_sections_1 call, handle non-contiguous case to create dynamic array map. gcc/fortran/ * f95-lang.cc (DEF_FUNCTION_TYPE_VAR_5): New symbol. * types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type. gcc/ * builtin-types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type. * omp-builtins.def (BUILT_IN_GOACC_DATA_START): Adjust function type to new BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR. * gimplify.cc (gimplify_scan_omp_clauses): Skip gimplification of OMP_CLAUSE_SIZE of non-contiguous array maps (which is a TREE_LIST). * omp-expand.cc (expand_omp_target): Add non-contiguous array descriptor pointers to variadic arguments. * omp-low.cc (append_field_to_record_type): New function. (create_noncontig_array_descr_type): Likewise. (create_noncontig_array_descr_init_code): Likewise. (scan_sharing_clauses): For non-contiguous array map kinds, check for supported dimension structure, and install non-contiguous array variable into current omp_context. (reorder_noncontig_array_clauses): New function. (scan_omp_target): Call reorder_noncontig_array_clauses to place non-contiguous array map clauses at beginning of clause sequence. (lower_omp_target): Add handling for non-contiguous array map kinds, add all created non-contiguous array descriptors to gimple_omp_target_data_arg. gcc/testsuite/ * c-c++-common/goacc/noncontig_array-1.c: New test. libgomp/ * libgomp_g.h (GOACC_data_start): Add variadic '...' to declaration. * libgomp.h (gomp_map_vars_openacc): New function declaration. * oacc-int.h (struct goacc_ncarray_dim): New struct declaration. (struct goacc_ncarray_descr_type): Likewise. (struct goacc_ncarray): Likewise. (struct goacc_ncarray_info): Likewise. (goacc_noncontig_array_create_ptrblock): New function declaration. * oacc-parallel.c (goacc_noncontig_array_count_rows): New function. (goacc_noncontig_array_compute_sizes): Likewise. (goacc_noncontig_array_fill_rows_1): Likewise. (goacc_noncontig_array_fill_rows): Likewise. (goacc_process_noncontiguous_arrays): Likewise. (goacc_noncontig_array_create_ptrblock): Likewise. (GOACC_parallel_keyed): Use goacc_process_noncontiguous_arrays to handle non-contiguous array descriptors at end of varargs, adjust to use gomp_map_vars_openacc. (GOACC_data_start): Likewise. Adjust function type to accept varargs. * target.c (gomp_map_vars_internal): Add struct goacc_ncarray_info * nca_info parameter, add handling code for non-contiguous arrays. (gomp_map_vars_openacc): Add new function for specialization of gomp_map_vars_internal for OpenACC structured region usage. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h: Support header for new tests. include/ * gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_3): Define. (enum gomp_map_kind): Add GOMP_MAP_NONCONTIG_ARRAY, GOMP_MAP_NONCONTIG_ARRAY_TO, GOMP_MAP_NONCONTIG_ARRAY_FROM, GOMP_MAP_NONCONTIG_ARRAY_TOFROM, GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO, GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM, GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM, GOMP_MAP_NONCONTIG_ARRAY_ALLOC, GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC, GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT. (GOMP_MAP_NONCONTIG_ARRAY_P): Define. 2023-04-18 Kwok Cheung Yeung <kcy@codesourcery.com> * gimplify.cc (omp_group_base): Handle GOMP_MAP_NONCONTIG_ARRAY_* map types. |
|
|
Jakub Jelinek
|
cd0059a197 | Update ChangeLog and version files for release | |
|
GCC Administrator
|
45532e3a92 | Daily bump. | |
|
Thomas Schwinge
|
a02d7f0edc |
GCN, nvptx: Errors during device probing are fatal
Currently, we silently disable libgomp GCN and nvptx plugins/devices in
presence of certain error conditions during device probing, thus typically
silently resorting to host-fallback execution. Make such errors fatal, similar
as for any other device access later on, so that we early and reliably notice
when things go wrong. (Keep just two cases non-fatal: (a) libgomp GCN or nvptx
plugins are available but 'libhsa-runtime64.so.1' or 'libcuda.so.1' are not,
and (b) those are available, but the corresponding devices are not.)
This resolves the issue that we've got execution test cases unexpectedly
PASSing, despite:
libgomp: GCN fatal error: Run-time could not be initialized
Runtime message: HSA_STATUS_ERROR_OUT_OF_RESOURCES: The runtime failed to allocate the necessary resources. This error may also occur when the core runtime library needs to spawn threads or create internal OS-specific events.
..., and therefore they were not offloaded to the GCN device, but ran in
host-fallback execution mode. What happend in that scenario is that in
'init_hsa_context' during the initial 'GOMP_OFFLOAD_get_num_devices' we ran
into 'HSA_STATUS_ERROR_OUT_OF_RESOURCES', but it wasn't fatal, but just
silently disabled the libgomp plugin/device.
Especially "entertaining" were cases where such unintended host-fallback
execution happened during effective-target checks like
'offload_device_available' (host-fallback execution there meaning: no offload
device available), but actual test cases then were running with an offload
device available, and therefore mis-configured.
include/
* cuda/cuda.h (CUresult): Add 'CUDA_ERROR_NO_DEVICE'.
libgomp/
* plugin/plugin-gcn.c (init_hsa_context): Add and handle
'bool probe' parameter. Adjust all users; errors during device
probing are fatal.
* plugin/plugin-nvptx.c (nvptx_get_num_devices): Aside from
'CUDA_ERROR_NO_DEVICE', errors during device probing are fatal.
|
|
|
GCC Administrator
|
b05f474c8f | Daily bump. | |
|
Tom Tromey
|
bc0e18a960 |
Fix PR libcc1/113977
PR libcc1/113977 points out a case where a simple expression is rejected with a compiler error message. The bug here is that gdb does not inform the plugin of the correct alignment -- in fact, there is no way to do that. This patch adds a new method to allow the alignment to be set, and bumps the C front end protocol version. It also includes some updates to various comments in 'include', done here to simplify the merge to binutils-gdb. include * gcc-cp-interface.h (gcc_cp_fe_context_function): Update comment. * gcc-c-interface.h (enum gcc_c_api_version) <GCC_C_FE_VERSION_2>: New constant. (gcc_c_fe_context_function): Update comment. * gcc-c-fe.def (finish_record_with_alignment): New method. Update documentation. libcc1 PR libcc1/113977 * libcc1plugin.cc (plugin_finish_record_or_union): New function. (plugin_finish_record_or_union): Rewrite. (plugin_init): Use GCC_C_FE_VERSION_2. * libcc1.cc (c_vtable): Use GCC_C_FE_VERSION_2. (gcc_c_fe_context): Check for GCC_C_FE_VERSION_2. |
|
|
GCC Administrator
|
ed5bf2080c | Daily bump. | |
|
Jakub Jelinek
|
65388b2865 |
c++, demangle: Implement https://github.com/itanium-cxx-abi/cxx-abi/issues/148 non-proposal
The following patch attempts to implement what apparently clang++ implemented for explicit object member function mangling, but nobody actually proposed in patch form in https://github.com/itanium-cxx-abi/cxx-abi/issues/148 2024-01-13 Jakub Jelinek <jakub@redhat.com> gcc/cp/ * mangle.cc (write_nested_name): Mangle explicit object member functions with H as per https://github.com/itanium-cxx-abi/cxx-abi/issues/148 non-proposal. gcc/testsuite/ * g++.dg/abi/mangle79.C: New test. include/ * demangle.h (enum demangle_component_type): Add DEMANGLE_COMPONENT_XOBJ_MEMBER_FUNCTION. libiberty/ * cp-demangle.c (FNQUAL_COMPONENT_CASE): Add case for DEMANGLE_COMPONENT_XOBJ_MEMBER_FUNCTION. (d_dump): Handle DEMANGLE_COMPONENT_XOBJ_MEMBER_FUNCTION. (d_nested_name): Parse H after N in nested name. (d_count_templates_scopes): Handle DEMANGLE_COMPONENT_XOBJ_MEMBER_FUNCTION. (d_print_mod): Likewise. (d_print_function_type): Likewise. * testsuite/demangle-expected: Add tests for explicit object member functions. |
|
|
GCC Administrator
|
73ce73fcad | Daily bump. | |
|
Jeff Law
|
9f7afa99c6 |
[committed] Adding missing prototype for __clzhi2 to xstormy port
xstormy16 has failed since the c99 transition due to a missing prototype for __clzhi2 in the implementation of stormy16_count_leading_zeros. This fixes the missing prototype. Pushed to the trunk. include/ * longlong.h (__stormy16_count_leading_zeros): Add prototype for __clzhi2. |
|
|
Jakub Jelinek
|
a945c346f5 | Update copyright years. | |
|
GCC Administrator
|
ea54b390ae | Daily bump. | |
|
Julian Brown
|
f5745dc142 |
OpenMP/OpenACC: Unordered/non-constant component offset runtime diagnostic
This patch adds support for non-constant component offsets in "map" clauses for OpenMP (and the equivalants for OpenACC), which are not able to be sorted into order at compile time. Normally struct accesses in such clauses are gathered together and sorted into increasing address order after a "GOMP_MAP_STRUCT" node: if we have variable indices, that is no longer possible. This version of the patch scales back the previously-posted version to merely add a diagnostic for incorrect usage of component accesses with variably-indexed arrays of structs: the only permitted variant is where we have multiple indices that are the same, but we could not prove so at compile time. Rather than silently producing the wrong result for cases where the indices are in fact different, we error out (e.g., "map(dtarr(i)%arrptr, dtarr(j)%arrptr(4:8))", for different i/j). For now, multiple *constant* array indices are still supported (see map-arrayofstruct-1.c). That could perhaps be addressed with a follow-up patch, if necessary. This version of the patch renumbers the GOMP_MAP_STRUCT_UNORD kind to avoid clashing with the OpenACC "non-contiguous" dynamic array support (though that is not yet applied to mainline). 2023-08-18 Julian Brown <julian@codesourcery.com> gcc/ * gimplify.cc (extract_base_bit_offset): Add VARIABLE_OFFSET parameter. (omp_get_attachment, omp_group_last, omp_group_base, omp_directive_maps_explicitly): Add GOMP_MAP_STRUCT_UNORD support. (omp_accumulate_sibling_list): Update calls to extract_base_bit_offset. Support GOMP_MAP_STRUCT_UNORD. (omp_build_struct_sibling_lists, gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses, gimplify_omp_target_update): Add GOMP_MAP_STRUCT_UNORD support. * omp-low.cc (lower_omp_target): Add GOMP_MAP_STRUCT_UNORD support. * tree-pretty-print.cc (dump_omp_clause): Likewise. include/ * gomp-constants.h (gomp_map_kind): Add GOMP_MAP_STRUCT_UNORD. libgomp/ * oacc-mem.c (find_group_last, goacc_enter_data_internal, goacc_exit_data_internal, GOACC_enter_exit_data): Add GOMP_MAP_STRUCT_UNORD support. * target.c (gomp_map_vars_internal): Add GOMP_MAP_STRUCT_UNORD support. Detect incorrect use of variable indexing of arrays of structs. (GOMP_target_enter_exit_data, gomp_target_task_fn): Add GOMP_MAP_STRUCT_UNORD support. * testsuite/libgomp.c-c++-common/map-arrayofstruct-1.c: New test. * testsuite/libgomp.c-c++-common/map-arrayofstruct-2.c: New test. * testsuite/libgomp.c-c++-common/map-arrayofstruct-3.c: New test. * testsuite/libgomp.fortran/map-subarray-5.f90: New test. |
|
|
GCC Administrator
|
4a6613e2a4 | Daily bump. | |
|
Tom Tromey
|
748766b8f6 |
Add some new DW_IDX_* constants
I've reimplemented the .debug_names code in GDB -- it was quite far from being correct, and the new implementation is much closer to what is specified by DWARF. However, the new writer in GDB needs to emit some symbol properties, so that the reader can be fully functional. This patch adds a few new DW_IDX_* constants, and tries to document the existing extensions as well. (My patch series add more documentation of these to the GDB manual as well.) include/ChangeLog 2023-12-10 Tom Tromey <tom@tromey.com> * dwarf2.def (DW_IDX_GNU_internal, DW_IDX_GNU_external): Comment. (DW_IDX_GNU_main, DW_IDX_GNU_language, DW_IDX_GNU_linkage_name): New constants. |
|
|
GCC Administrator
|
2e0f3f9759 | Daily bump. | |
|
Jason Merrill
|
c3f281a0c1 |
c++: mangle function template constraints
Per https://github.com/itanium-cxx-abi/cxx-abi/issues/24 and https://github.com/itanium-cxx-abi/cxx-abi/pull/166 We need to mangle constraints to be able to distinguish between function templates that only differ in constraints. From the latter link, we want to use the template parameter mangling previously specified for lambdas to also make explicit the form of a template parameter where the argument is not a "natural" fit for it, such as when the parameter is constrained or deduced. I'm concerned about how the latter link changes the mangling for some C++98 and C++11 patterns, so I've limited template_parm_natural_p to avoid two cases found by running the testsuite with -Wabi forced on: template <class T, T V> T f() { return V; } int main() { return f<int,42>(); } template <int i> int max() { return i; } template <int i, int j, int... rest> int max() { int sub = max<j, rest...>(); return i > sub ? i : sub; } int main() { return max<1,2,3>(); } A third C++11 pattern is changed by this patch: template <template <typename...> class TT, typename... Ts> TT<Ts...> f(); template <typename> struct A { }; int main() { f<A,int>(); } I aim to resolve these with the ABI committee before GCC 14.1. We also need to resolve https://github.com/itanium-cxx-abi/cxx-abi/issues/38 (mangling references to dependent template-ids where the name is fully resolved) as references to concepts in std:: will consistently run into this area. This is why mangle-concepts1.C only refers to concepts in the global namespace so far. The library changes are to avoid trying to mangle builtins, which fails. Demangler support and test coverage is not complete yet. gcc/cp/ChangeLog: * cp-tree.h (TEMPLATE_ARGS_TYPE_CONSTRAINT_P): New. (get_concept_check_template): Declare. * constraint.cc (combine_constraint_expressions) (finish_shorthand_constraint): Use UNKNOWN_LOCATION. * pt.cc (convert_generic_types_to_packs): Likewise. * mangle.cc (write_constraint_expression) (write_tparms_constraints, write_type_constraint) (template_parm_natural_p, write_requirement) (write_requires_expr): New. (write_encoding): Mangle trailing requires-clause. (write_name): Pass parms to write_template_args. (write_template_param_decl): Factor out from... (write_closure_template_head): ...here. (write_template_args): Mangle non-natural parms and requires-clause. (write_expression): Handle REQUIRES_EXPR. include/ChangeLog: * demangle.h (enum demangle_component_type): Add DEMANGLE_COMPONENT_CONSTRAINTS. libiberty/ChangeLog: * cp-demangle.c (d_make_comp): Handle DEMANGLE_COMPONENT_CONSTRAINTS. (d_count_templates_scopes): Likewise. (d_print_comp_inner): Likewise. (d_maybe_constraints): New. (d_encoding, d_template_args_1): Call it. (d_parmlist): Handle 'Q'. * testsuite/demangle-expected: Add some constraint tests. libstdc++-v3/ChangeLog: * include/std/bit: Avoid builtins in requires-clauses. * include/std/variant: Likewise. gcc/testsuite/ChangeLog: * g++.dg/abi/mangle10.C: Disable compat aliases. * g++.dg/abi/mangle52.C: Specify ABI 18. * g++.dg/cpp2a/class-deduction-alias3.C * g++.dg/cpp2a/class-deduction-alias8.C: Avoid builtins in requires-clauses. * g++.dg/abi/mangle-concepts1.C: New test. * g++.dg/abi/mangle-ttp1.C: New test. |
|
|
GCC Administrator
|
6c85b8a987 | Daily bump. | |
|
Jakub Jelinek
|
bf4f40cc31 |
libiberty: Use x86 HW optimized sha1
Nick has approved this patch (+ small ld change to use it for --build-id=),
so I'm commiting it to GCC as master as well.
If anyone from ARM would be willing to implement it similarly with
vsha1{cq,mq,pq,h,su0q,su1q}_u32 intrinsics, it could be a useful linker
speedup on those hosts as well, the intent in sha1.c was that
sha1_hw_process_bytes, sha1_hw_process_block functions
would be defined whenever
defined (HAVE_X86_SHA1_HW_SUPPORT) || defined (HAVE_WHATEVERELSE_SHA1_HW_SUPPORT)
but the body of sha1_hw_process_block and sha1_choose_process_bytes
would then have #elif defined (HAVE_WHATEVERELSE_SHA1_HW_SUPPORT) for
the other arch support, similarly for any target attributes on
sha1_hw_process_block if needed.
2023-11-28 Jakub Jelinek <jakub@redhat.com>
include/
* sha1.h (sha1_process_bytes_fn): New typedef.
(sha1_choose_process_bytes): Declare.
libiberty/
* configure.ac (HAVE_X86_SHA1_HW_SUPPORT): New check.
* sha1.c: If HAVE_X86_SHA1_HW_SUPPORT is defined, include x86intrin.h
and cpuid.h.
(sha1_hw_process_bytes, sha1_hw_process_block,
sha1_choose_process_bytes): New functions.
* config.in: Regenerated.
* configure: Regenerated.
|
|
|
GCC Administrator
|
c48f105685 | Daily bump. | |
|
Kwok Cheung Yeung
|
a49c7d3193 |
openmp: Add support for the 'indirect' clause in C/C++
This adds support for the 'indirect' clause in the 'declare target' directive. Functions declared as indirect may be called via function pointers passed from the host in offloaded code. Virtual calls to member functions via the object pointer in C++ are currently not supported in target regions. 2023-11-07 Kwok Cheung Yeung <kcy@codesourcery.com> gcc/c-family/ * c-attribs.cc (c_common_attribute_table): Add attribute for indirect functions. * c-pragma.h (enum parma_omp_clause): Add entry for indirect clause. gcc/c/ * c-decl.cc (c_decl_attributes): Add attribute for indirect functions. * c-lang.h (c_omp_declare_target_attr): Add indirect field. * c-parser.cc (c_parser_omp_clause_name): Handle indirect clause. (c_parser_omp_clause_indirect): New. (c_parser_omp_all_clauses): Handle indirect clause. (OMP_DECLARE_TARGET_CLAUSE_MASK): Add indirect clause to mask. (c_parser_omp_declare_target): Handle indirect clause. Emit error message if device_type or indirect clauses used alone. Emit error if indirect clause used with device_type that is not 'any'. (OMP_BEGIN_DECLARE_TARGET_CLAUSE_MASK): Add indirect clause to mask. (c_parser_omp_begin): Handle indirect clause. * c-typeck.cc (c_finish_omp_clauses): Handle indirect clause. gcc/cp/ * cp-tree.h (cp_omp_declare_target_attr): Add indirect field. * decl2.cc (cplus_decl_attributes): Add attribute for indirect functions. * parser.cc (cp_parser_omp_clause_name): Handle indirect clause. (cp_parser_omp_clause_indirect): New. (cp_parser_omp_all_clauses): Handle indirect clause. (handle_omp_declare_target_clause): Add extra parameter. Add indirect attribute for indirect functions. (OMP_DECLARE_TARGET_CLAUSE_MASK): Add indirect clause to mask. (cp_parser_omp_declare_target): Handle indirect clause. Emit error message if device_type or indirect clauses used alone. Emit error if indirect clause used with device_type that is not 'any'. (OMP_BEGIN_DECLARE_TARGET_CLAUSE_MASK): Add indirect clause to mask. (cp_parser_omp_begin): Handle indirect clause. * semantics.cc (finish_omp_clauses): Handle indirect clause. gcc/ * lto-cgraph.cc (enum LTO_symtab_tags): Add tag for indirect functions. (output_offload_tables): Write indirect functions. (input_offload_tables): read indirect functions. * lto-section-names.h (OFFLOAD_IND_FUNC_TABLE_SECTION_NAME): New. * omp-builtins.def (BUILT_IN_GOMP_TARGET_MAP_INDIRECT_PTR): New. * omp-offload.cc (offload_ind_funcs): New. (omp_discover_implicit_declare_target): Add functions marked with 'omp declare target indirect' to indirect functions list. (omp_finish_file): Add indirect functions to section for offload indirect functions. (execute_omp_device_lower): Redirect indirect calls on target by passing function pointer to BUILT_IN_GOMP_TARGET_MAP_INDIRECT_PTR. (pass_omp_device_lower::gate): Run pass_omp_device_lower if indirect functions are present on an accelerator device. * omp-offload.h (offload_ind_funcs): New. * tree-core.h (omp_clause_code): Add OMP_CLAUSE_INDIRECT. * tree.cc (omp_clause_num_ops): Add entry for OMP_CLAUSE_INDIRECT. (omp_clause_code_name): Likewise. * tree.h (OMP_CLAUSE_INDIRECT_EXPR): New. * config/gcn/mkoffload.cc (process_asm): Process offload_ind_funcs section. Count number of indirect functions. (process_obj): Emit number of indirect functions. * config/nvptx/mkoffload.cc (ind_func_ids, ind_funcs_tail): New. (process): Emit offload_ind_func_table in PTX code. Emit indirect function names and count in image. * config/nvptx/nvptx.cc (nvptx_record_offload_symbol): Mark indirect functions in PTX code with IND_FUNC_MAP. gcc/testsuite/ * c-c++-common/gomp/declare-target-7.c: Update expected error message. * c-c++-common/gomp/declare-target-indirect-1.c: New. * c-c++-common/gomp/declare-target-indirect-2.c: New. * g++.dg/gomp/attrs-21.C (v12): Update expected error message. * g++.dg/gomp/declare-target-indirect-1.C: New. * gcc.dg/gomp/attrs-21.c (v12): Update expected error message. include/ * gomp-constants.h (GOMP_VERSION): Increment to 3. (GOMP_VERSION_SUPPORTS_INDIRECT_FUNCS): New. libgcc/ * offloadstuff.c (OFFLOAD_IND_FUNC_TABLE_SECTION_NAME): New. (__offload_ind_func_table): New. (__offload_ind_funcs_end): New. (__OFFLOAD_TABLE__): Add entries for indirect functions. libgomp/ * Makefile.am (libgomp_la_SOURCES): Add target-indirect.c. * Makefile.in: Regenerate. * libgomp-plugin.h (GOMP_INDIRECT_ADDR_MAP): New define. (GOMP_OFFLOAD_load_image): Add extra argument. * libgomp.h (struct indirect_splay_tree_key_s): New. (indirect_splay_tree_node, indirect_splay_tree, indirect_splay_tree_key): New. (indirect_splay_compare): New. * libgomp.map (GOMP_5.1.1): Add GOMP_target_map_indirect_ptr. * libgomp.texi (OpenMP 5.1): Update documentation on indirect calls in target region and on indirect clause. (Other new OpenMP 5.2 features): Add entry for virtual function calls. * libgomp_g.h (GOMP_target_map_indirect_ptr): Add prototype. * oacc-host.c (host_load_image): Add extra argument. * target.c (gomp_load_image_to_device): If the GOMP_VERSION is high enough, read host indirect functions table and pass to load_image_func. * config/accel/target-indirect.c: New. * config/linux/target-indirect.c: New. * config/gcn/team.c (build_indirect_map): Add prototype. (gomp_gcn_enter_kernel): Initialize support for indirect function calls on GCN target. * config/nvptx/team.c (build_indirect_map): Add prototype. (gomp_nvptx_main): Initialize support for indirect function calls on NVPTX target. * plugin/plugin-gcn.c (struct gcn_image_desc): Add field for indirect functions count. (GOMP_OFFLOAD_load_image): Add extra argument. If the GOMP_VERSION is high enough, build address translation table and copy it to target memory. * plugin/plugin-nvptx.c (nvptx_tdata): Add field for indirect functions count. (GOMP_OFFLOAD_load_image): Add extra argument. If the GOMP_VERSION is high enough, Build address translation table and copy it to target memory. * testsuite/libgomp.c-c++-common/declare-target-indirect-1.c: New. * testsuite/libgomp.c-c++-common/declare-target-indirect-2.c: New. * testsuite/libgomp.c++/declare-target-indirect-1.C: New. |
|
|
GCC Administrator
|
f75fc1f083 | Daily bump. | |
|
Chung-Lin Tang
|
3a3596389c |
OpenACC 2.7: Implement self clause for compute constructs
This patch implements the 'self' clause for compute constructs: parallel, kernels, and serial. This clause conditionally uses the local device (the host mult-core CPU) as the executing device of the compute region. The actual implementation of the "local device" device type inside libgomp (presumably using pthreads) is still not yet completed, so the libgomp side is still implemented the exact same as host-fallback mode. (so as of now, it essentially behaves like the 'if' clause with the condition inverted) gcc/c/ChangeLog: * c-parser.cc (c_parser_oacc_compute_clause_self): New function. (c_parser_oacc_all_clauses): Add new 'bool compute_p = false' parameter, add parsing of self clause when compute_p is true. (OACC_KERNELS_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_SELF. (OACC_PARALLEL_CLAUSE_MASK): Likewise, (OACC_SERIAL_CLAUSE_MASK): Likewise. (c_parser_oacc_compute): Adjust call to c_parser_oacc_all_clauses to set compute_p argument to true. * c-typeck.cc (c_finish_omp_clauses): Add OMP_CLAUSE_SELF case. gcc/cp/ChangeLog: * parser.cc (cp_parser_oacc_compute_clause_self): New function. (cp_parser_oacc_all_clauses): Add new 'bool compute_p = false' parameter, add parsing of self clause when compute_p is true. (OACC_KERNELS_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_SELF. (OACC_PARALLEL_CLAUSE_MASK): Likewise, (OACC_SERIAL_CLAUSE_MASK): Likewise. (cp_parser_oacc_compute): Adjust call to c_parser_oacc_all_clauses to set compute_p argument to true. * pt.cc (tsubst_omp_clauses): Add OMP_CLAUSE_SELF case. * semantics.cc (c_finish_omp_clauses): Add OMP_CLAUSE_SELF case, merged with OMP_CLAUSE_IF case. gcc/fortran/ChangeLog: * gfortran.h (typedef struct gfc_omp_clauses): Add self_expr field. * openmp.cc (enum omp_mask2): Add OMP_CLAUSE_SELF. (gfc_match_omp_clauses): Add handling for OMP_CLAUSE_SELF. (OACC_PARALLEL_CLAUSES): Add OMP_CLAUSE_SELF. (OACC_KERNELS_CLAUSES): Likewise. (OACC_SERIAL_CLAUSES): Likewise. (resolve_omp_clauses): Add handling for omp_clauses->self_expr. * trans-openmp.cc (gfc_trans_omp_clauses): Add handling of clauses->self_expr and building of OMP_CLAUSE_SELF tree clause. (gfc_split_omp_clauses): Add handling of self_expr field copy. gcc/ChangeLog: * gimplify.cc (gimplify_scan_omp_clauses): Add OMP_CLAUSE_SELF case. (gimplify_adjust_omp_clauses): Likewise. * omp-expand.cc (expand_omp_target): Add OMP_CLAUSE_SELF expansion code, * omp-low.cc (scan_sharing_clauses): Add OMP_CLAUSE_SELF case. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_SELF enum. * tree-nested.cc (convert_nonlocal_omp_clauses): Add OMP_CLAUSE_SELF case. (convert_local_omp_clauses): Likewise. * tree-pretty-print.cc (dump_omp_clause): Add OMP_CLAUSE_SELF case. * tree.cc (omp_clause_num_ops): Add OMP_CLAUSE_SELF entry. (omp_clause_code_name): Likewise. * tree.h (OMP_CLAUSE_SELF_EXPR): New macro. gcc/testsuite/ChangeLog: * c-c++-common/goacc/self-clause-1.c: New test. * c-c++-common/goacc/self-clause-2.c: New test. * gfortran.dg/goacc/self.f95: New test. include/ChangeLog: * gomp-constants.h (GOACC_FLAG_LOCAL_DEVICE): New flag bit value. libgomp/ChangeLog: * oacc-parallel.c (GOACC_parallel_keyed): Add code to handle GOACC_FLAG_LOCAL_DEVICE case. * testsuite/libgomp.oacc-c-c++-common/self-1.c: New test. |
|
|
GCC Administrator
|
f9ef2e6dcd | Daily bump. | |
|
Zhang, Jun
|
e1e127de18 |
x86: set spincount 1 for x86 hybrid platform
By test, we find in hybrid platform spincount 1 is better. Use '-march=native -Ofast -funroll-loops -flto', results as follows: spec2017 speed RPL ADL 657.xz_s 0.00% 0.50% 603.bwaves_s 10.90% 26.20% 607.cactuBSSN_s 5.50% 72.50% 619.lbm_s 2.40% 2.50% 621.wrf_s -7.70% 2.40% 627.cam4_s 0.50% 0.70% 628.pop2_s 48.20% 153.00% 638.imagick_s -0.10% 0.20% 644.nab_s 2.30% 1.40% 649.fotonik3d_s 8.00% 13.80% 654.roms_s 1.20% 1.10% Geomean-int 0.00% 0.50% Geomean-fp 6.30% 21.10% Geomean-all 5.70% 19.10% omp2012 RPL ADL 350.md -1.81% -1.75% 351.bwaves 7.72% 12.50% 352.nab 14.63% 19.71% 357.bt331 -0.20% 1.77% 358.botsalgn 0.00% 0.00% 359.botsspar 0.00% 0.65% 360.ilbdc 0.00% 0.25% 362.fma3d 2.66% -0.51% 363.swim 10.44% 0.00% 367.imagick 0.00% 0.12% 370.mgrid331 2.49% 25.56% 371.applu331 1.06% 4.22% 372.smithwa 0.74% 3.34% 376.kdtree 10.67% 16.03% GEOMEAN 3.34% 5.53% include/ChangeLog: PR target/109812 * spincount.h: New file. libgomp/ChangeLog: * env.c (initialize_env): Use do_adjust_default_spincount. * config/linux/x86/spincount.h: New file. |
|
|
GCC Administrator
|
6cd8527307 | Daily bump. | |
|
Jason Merrill
|
810bcc0015 |
c++: constrained hidden friends [PR109751]
r13-4035 avoided a problem with overloading of constrained hidden friends by checking satisfaction, but checking satisfaction early is inconsistent with the usual late checking and can lead to hard errors, so let's not do that after all. We were wrongly treating the different instantiations of the same friend template as the same function because maybe_substitute_reqs_for was failing to actually substitute in the case of a non-template friend. But we don't actually need to do the substitution anyway, because [temp.friend] says that such a friend can't be the same as any other declaration. After fixing that, instead of a redefinition error we got an ambiguous overload error, fixed by allowing constrained hidden friends to coexist until overload resolution, at which point they probably won't be in the same ADL overload set anyway. And we avoid mangling collisions by following the proposed mangling for these friends as a member function with an extra 'F' before the name. I demangle this by just adding [friend] to the name of the function because it's not feasible to reconstruct the actual scope of the function since the mangling ABI doesn't distinguish between class and namespace scopes. PR c++/109751 gcc/cp/ChangeLog: * cp-tree.h (member_like_constrained_friend_p): Declare. * decl.cc (member_like_constrained_friend_p): New. (function_requirements_equivalent_p): Check it. (duplicate_decls): Check it. (grokfndecl): Check friend template constraints. * mangle.cc (decl_mangling_context): Check it. (write_unqualified_name): Check it. * pt.cc (uses_outer_template_parms_in_constraints): Fix for friends. (tsubst_friend_function): Don't check satisfaction. include/ChangeLog: * demangle.h (enum demangle_component_type): Add DEMANGLE_COMPONENT_FRIEND. libiberty/ChangeLog: * cp-demangle.c (d_make_comp): Handle DEMANGLE_COMPONENT_FRIEND. (d_count_templates_scopes): Likewise. (d_print_comp_inner): Likewise. (d_unqualified_name): Handle member-like friend mangling. * testsuite/demangle-expected: Add test. gcc/testsuite/ChangeLog: * g++.dg/cpp2a/concepts-friend11.C: Now works. Add template. * g++.dg/cpp2a/concepts-friend15.C: New test. |
|
|
GCC Administrator
|
4b92dba78d | Daily bump. | |
|
Vladimir Mezentsev
|
24552056fd
|
gprofng: a new GNU profiler
ChangeLog: * Makefile.def: Add gprofng module. * configure.ac: Add --enable-gprofng option. * Makefile.in: Regenerate. * configure: Regenerate. include/ChangeLog: * collectorAPI.h: New file. * libcollector.h: New file. * libfcollector.h: New file. |
|
|
Alan Modra
|
432c6f05b0
|
gcc-4.5 build fixes
Trying to build binutils with an older gcc currently fails. Working around these gcc bugs is not onerous so let's fix them. include/ChangeLog: * xtensa-dynconfig.h (xtensa_isa_internal): Delete unnecessary forward declaration. |
|
|
Alan Modra
|
24f5a73aa3
|
PR29961, plugin-api.h: "Could not detect architecture endianess"
Found when attempting to build binutils on sparc sunos-5.8 where sys/byteorder.h defines _BIG_ENDIAN but not any of the BYTE_ORDER variants. This patch adds the extra tests to cope with the old machine, and tidies the header a little. include/ChangeLog: * plugin-api.h: When handling non-gcc or gcc < 4.6.0 include necessary header files before testing macros. Make more use of #elif. Test _LITTLE_ENDIAN and _BIG_ENDIAN in final tests. |
|
|
GCC Administrator
|
861962eee1 | Daily bump. | |
|
Tobias Burnus
|
8b9e559fe7 |
libgomp: cuda.h and omp_target_memcpy_rect cleanup
Fixes for commit r14-2792-g25072a477a56a727b369bf9b20f4d18198ff5894
"OpenMP: Call cuMemcpy2D/cuMemcpy3D for nvptx for omp_target_memcpy_rect",
namely:
In that commit, the code was changed to handle shared-memory devices;
however, as pointed out, omp_target_memcpy_check already set the pointer
to NULL in that case. Hence, this commit reverts to the prior version.
In cuda.h, it adds cuMemcpyPeer{,Async} for symmetry for cuMemcpy3DPeer
(all currently unused) and in three structs, fixes reserved-member names
and remove a bogus 'const' in three structs.
And it changes a DLSYM to DLSYM_OPT as not all plugins support the new
functions, yet.
include/ChangeLog:
* cuda/cuda.h (CUDA_MEMCPY2D, CUDA_MEMCPY3D, CUDA_MEMCPY3D_PEER):
Remove bogus 'const' from 'const void *dst' and fix reserved-name
name in those structs.
(cuMemcpyPeer, cuMemcpyPeerAsync): Add.
libgomp/ChangeLog:
* target.c (omp_target_memcpy_rect_worker): Undo dim=1 change for
GOMP_OFFLOAD_CAP_SHARED_MEM.
(omp_target_memcpy_rect_copy): Likewise for lock condition.
(gomp_load_plugin_for_device): Use DLSYM_OPT not DLSYM for
memcpy3d/memcpy2d.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_memcpy2d,
GOMP_OFFLOAD_memcpy3d): Use memset 0 to nullify reserved and
unused src/dst fields for that mem type; remove '{src,dst}LOD = 0'.
|
|
|
GCC Administrator
|
5278cd6a45 | Daily bump. | |
|
Tobias Burnus
|
25072a477a |
OpenMP: Call cuMemcpy2D/cuMemcpy3D for nvptx for omp_target_memcpy_rect
When copying a 2D or 3D rectangular memmory block, the performance is better when using CUDA's cuMemcpy2D/cuMemcpy3D instead of copying the data one by one. That's what this commit does. Additionally, it permits device-to-device copies, if neccessary using a temporary variable on the host. include/ChangeLog: * cuda/cuda.h (CUlimit): Add CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_INVALID_HANDLE. (CUarray, CUmemorytype, CUDA_MEMCPY2D, CUDA_MEMCPY3D, CUDA_MEMCPY3D_PEER): New typdefs. (cuMemcpy2D, cuMemcpy2DAsync, cuMemcpy2DUnaligned, cuMemcpy3D, cuMemcpy3DAsync, cuMemcpy3DPeer, cuMemcpy3DPeerAsync): New prototypes. libgomp/ChangeLog: * libgomp-plugin.h (GOMP_OFFLOAD_memcpy2d, GOMP_OFFLOAD_memcpy3d): New prototypes. * libgomp.h (struct gomp_device_descr): Add memcpy2d_func and memcpy3d_func. * libgomp.texi (nvtpx): Document when cuMemcpy2D/cuMemcpy3D is used. * oacc-host.c (memcpy2d_func, .memcpy3d_func): Init with NULL. * plugin/cuda-lib.def (cuMemcpy2D, cuMemcpy2DUnaligned, cuMemcpy3D): Invoke via CUDA_ONE_CALL. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_memcpy2d, GOMP_OFFLOAD_memcpy3d): New. * target.c (omp_target_memcpy_rect_worker): (omp_target_memcpy_rect_check, omp_target_memcpy_rect_copy): Permit all device-to-device copyies; invoke new plugins for 2D and 3D copying when available. (gomp_load_plugin_for_device): DLSYM the new plugin functions. * testsuite/libgomp.c/target-12.c: Fix dimension bug. * testsuite/libgomp.fortran/target-12.f90: Likewise. * testsuite/libgomp.fortran/target-memcpy-rect-1.f90: New test. |
|
|
GCC Administrator
|
9d250bdb88 | Daily bump. | |
|
Tobias Burnus
|
38944ec2a6 |
OpenMP: Cleanups related to the 'present' modifier
Reduce number of enum values passed to libgomp as
GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} have the same semantic as
GOMP_MAP_FORCE_PRESENT (i.e. abort if not present, otherwise ignore);
that's different to GOMP_MAP_ALWAYS_PRESENT_{TO,TOFROM,FROM} which also
abort if not present but copy data when present. This is is a follow-up to
the commit r14-1579-g4ede915d5dde93 done 6 days ago.
Additionally, the commit improves a libgomp run-time and a C/C++ compile-time
error wording and extends testcases a tiny bit.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_clause_map): Reword error message for
clearness especially with 'omp target (enter/exit) data.'
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_clause_map): Reword error message for
clearness especially with 'omp target (enter/exit) data.'
* semantics.cc (handle_omp_array_sections): Handle
GOMP_MAP_{ALWAYS_,}PRESENT_{TO,TOFROM,FROM,ALLOC} enum values.
gcc/ChangeLog:
* gimplify.cc (gimplify_adjust_omp_clauses_1): Use
GOMP_MAP_FORCE_PRESENT for 'present alloc' implicit mapping.
(gimplify_adjust_omp_clauses): Change
GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} to the equivalent
GOMP_MAP_FORCE_PRESENT.
* omp-low.cc (lower_omp_target): Remove handling of no-longer valid
GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC}; update map kinds used for
to/from clauses with present modifier.
include/ChangeLog:
* gomp-constants.h (enum gomp_map_kind): Change the enum values
GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} to be compiler only.
(GOMP_MAP_PRESENT_P): Update to include also GOMP_MAP_FORCE_PRESENT.
libgomp/ChangeLog:
* target.c (gomp_to_device_kind_p, gomp_map_vars_internal): Replace
GOMP_MAP_PRESENT_{FROM,TO,TOFROM,ACLLOC} by GOMP_MAP_FORCE_PRESENT.
(gomp_map_vars_internal, gomp_update): Likewise; unify and improve
error message.
* testsuite/libgomp.c-c++-common/target-present-2.c: Update for
changed error message.
* testsuite/libgomp.fortran/target-present-1.f90: Likewise.
* testsuite/libgomp.fortran/target-present-2.f90: Likewise.
* testsuite/libgomp.oacc-c-c++-common/present-1.c: Likewise.
* testsuite/libgomp.c-c++-common/target-present-1.c: Likewise and
extend testcase to check that data is copied when needed.
* testsuite/libgomp.c-c++-common/target-present-3.c: Likewise.
* testsuite/libgomp.fortran/target-present-3.f90: Likewise.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/defaultmap-4.c: Update scan-tree-dump.
* c-c++-common/gomp/map-9.c: Likewise.
* gfortran.dg/gomp/defaultmap-8.f90: Likewise.
* gfortran.dg/gomp/map-11.f90: Likewise.
* gfortran.dg/gomp/target-update-1.f90: Likewise.
* gfortran.dg/gomp/map-12.f90: Likewise; also check original dump.
* c-c++-common/gomp/map-6.c: Update dg-error and also check
clause error with 'target (enter/exit) data'.
|