| Age | Commit message (Collapse) | Author | Files | Lines |
|---|
|
This patch reprocesses expanded clauses after 'declare mapper'
instantiation -- checking things such as duplicated clauses, illegal
use of strided accesses, and so forth. Two functions are broken out
of the 'resolve_omp_clauses' function and reused in a new function
'resolve_omp_mapper_clauses', called after mapper instantiation.
This improves diagnostic output.
2023-08-10 Julian Brown <julian@codesourcery.com>
gcc/fortran/
* gfortran.h (gfc_omp_clauses): Add NS field.
* openmp.cc (verify_omp_clauses_symbol_dups,
omp_verify_map_motion_clauses): New functions, broken out of...
(resolve_omp_clauses): Here. Record namespace containing clauses.
Call above functions.
(resolve_omp_mapper_clauses): New function, using helper functions
broken out above.
(gfc_resolve_omp_directive): Add NS parameter to resolve_omp_clauses
calls.
(gfc_omp_instantiate_mappers): Call resolve_omp_mapper_clauses if we
instantiate any mappers.
gcc/testsuite/
* gfortran.dg/gomp/declare-mapper-26.f90: New test.
* gfortran.dg/gomp/declare-mapper-29.f90: New test.
|
|
This patch moves the code for explicit 'declare mapper' directive
instantiation in the Fortran front-end to openmp.cc from trans-openmp.cc.
The transformation takes place entirely in the front end's own
representation and doesn't involve middle-end trees at all. Also, having
the code in openmp.cc is more convenient for the following patch that
introduces the 'resolve_omp_mapper_clauses' function.
2023-08-10 Julian Brown <julian@codesourcery.com>
gcc/fortran/
* gfortran.h (toc_directive): Move here.
(gfc_omp_instantiate_mappers, gfc_get_location): Add prototypes.
* openmp.cc (omp_split_map_op, omp_join_map_op, omp_map_decayed_kind,
omp_basic_map_kind_name, gfc_subst_replace, gfc_subst_prepend_ref,
gfc_subst_in_expr_1, gfc_subst_in_expr, gfc_subst_mapper_var): Move
here.
(gfc_omp_instantiate_mapper, gfc_omp_instantiate_mappers): Move here
and rename.
* trans-openmp.cc (toc_directive, omp_split_map_op, omp_join_map_op,
omp_map_decayed_kind, gfc_subst_replace, gfc_subst_prepend_ref,
gfc_subst_in_expr_1, gfc_subst_in_expr, gfc_subst_mapper_var,
gfc_trans_omp_instantiate_mapper, gfc_trans_omp_instantiate_mappers):
Remove from here.
(gfc_trans_omp_target, gfc_trans_omp_target_data,
gfc_trans_omp_target_enter_data, gfc_trans_omp_target_exit_data):
Rename calls to gfc_omp_instantiate_mappers.
|
|
This patch allows 'declare mapper' mappers to be used on 'omp target
data', 'omp target enter data' and 'omp target exit data' directives.
For each of these, only explicit mappings are supported, unlike for
'omp target' directives where implicit uses of variables inside an
offload region might trigger mappers also.
Each of C, C++ and Fortran are supported.
The patch also adjusts 'map kind decay' to match OpenMP 5.2 semantics,
which is particularly important with regard to 'exit data' operations.
2023-07-06 Julian Brown <julian@codesourcery.com>
gcc/c-family/
* c-common.h (c_omp_region_type): Add C_ORT_EXIT_DATA,
C_ORT_OMP_EXIT_DATA.
(c_omp_instantiate_mappers): Add region type parameter.
* c-omp.cc (omp_split_map_kind, omp_join_map_kind,
omp_map_decayed_kind): New functions.
(omp_instantiate_mapper): Add ORT parameter. Implement map kind decay
for instantiated mapper clauses.
(c_omp_instantiate_mappers): Add ORT parameter, pass to
omp_instantiate_mapper.
gcc/c/
* c-parser.cc (c_parser_omp_target_data): Instantiate mappers for
'omp target data'.
(c_parser_omp_target_enter_data): Instantiate mappers for 'omp target
enter data'.
(c_parser_omp_target_exit_data): Instantiate mappers for 'omp target
exit data'.
(c_parser_omp_target): Add c_omp_region_type argument to
c_omp_instantiate_mappers call.
* c-tree.h (c_omp_instantiate_mappers): Remove spurious prototype.
gcc/cp/
* parser.cc (cp_parser_omp_target_data): Instantiate mappers for 'omp
target data'.
(cp_parser_omp_target_enter_data): Instantiate mappers for 'omp target
enter data'.
(cp_parser_omp_target_exit_data): Instantiate mappers for 'omp target
exit data'.
(cp_parser_omp_target): Add c_omp_region_type argument to
c_omp_instantiate_mappers call.
* pt.cc (tsubst_omp_clauses): Instantiate mappers for OMP regions other
than just C_ORT_OMP_TARGET.
(tsubst_expr): Update call to tsubst_omp_clauses for OMP_TARGET_UPDATE,
OMP_TARGET_ENTER_DATA, OMP_TARGET_EXIT_DATA stanza.
* semantics.cc (cxx_omp_map_array_section): Avoid calling
build_array_ref for non-array/non-pointer bases (error reported
already).
gcc/fortran/
* trans-openmp.cc (omp_split_map_op, omp_join_map_op,
omp_map_decayed_kind): New functions.
(gfc_trans_omp_instantiate_mapper): Add CD parameter. Implement map
kind decay.
(gfc_trans_omp_instantiate_mappers): Add CD parameter. Pass to above
function.
(gfc_trans_omp_target_data): Instantiate mappers for 'omp target data'.
(gfc_trans_omp_target_enter_data): Instantiate mappers for 'omp target
enter data'.
(gfc_trans_omp_target_exit_data): Instantiate mappers for 'omp target
exit data'.
gcc/testsuite/
* c-c++-common/gomp/declare-mapper-15.c: New test.
* c-c++-common/gomp/declare-mapper-16.c: New test.
* g++.dg/gomp/declare-mapper-1.C: Adjust expected scan output.
* gfortran.dg/gomp/declare-mapper-22.f90: New test.
* gfortran.dg/gomp/declare-mapper-23.f90: New test.
|
|
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.
|
|
This patch adds support for "declare mapper" directives (and the "mapper"
modifier on "map" clauses) for C.
gcc/c/
* c-decl.cc (c_omp_mapper_id, c_omp_mapper_decl, c_omp_mapper_lookup,
c_omp_extract_mapper_directive, c_omp_map_array_section,
c_omp_scan_mapper_bindings_r, c_omp_scan_mapper_bindings): New
functions.
* c-objc-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 for C.
* c-parser.cc (c_parser_omp_clause_map): Add KIND parameter. Handle
mapper modifier.
(c_parser_omp_all_clauses): Update call to c_parser_omp_clause_map with
new kind argument.
(c_parser_omp_target): Instantiate explicit mappers and record bindings
for implicit mappers.
(c_parser_omp_declare_mapper): Parse "declare mapper" directives.
(c_parser_omp_declare): Support "declare mapper".
* c-tree.h (c_omp_finish_mapper_clauses, c_omp_mapper_lookup,
c_omp_extract_mapper_directive, c_omp_map_array_section,
c_omp_mapper_id, c_omp_mapper_decl, c_omp_scan_mapper_bindings,
c_omp_instantiate_mappers): Add prototypes.
* c-typeck.cc (c_finish_omp_clauses): Handle GOMP_MAP_PUSH_MAPPER_NAME
and GOMP_MAP_POP_MAPPER_NAME.
(c_omp_finish_mapper_clauses): New function (langhook).
gcc/testsuite/
* c-c++-common/gomp/declare-mapper-3.c: Enable for C.
* 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: Enable for C.
* gcc.dg/gomp/declare-mapper-10.c: New test.
* gcc.dg/gomp/declare-mapper-11.c: New test.
libgomp/
* testsuite/libgomp.c-c++-common/declare-mapper-9.c: Enable for C.
* 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.
|
|
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 (c_omp_region_type): Add C_ORT_DECLARE_MAPPER and
C_ORT_OMP_DECLARE_MAPPER codes.
(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): Add prototype.
(cp_check_omp_declare_mapper): Add prototype.
(omp_instantiate_mappers): Add prototype.
(cxx_omp_finish_mapper_clauses): Add prototype.
(cxx_omp_mapper_lookup): Add prototype.
(cxx_omp_extract_mapper_directive): Add prototype.
(cxx_omp_map_array_section): Add prototype.
* 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.
|
|
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.
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/implied-copy-1.c: Likewise.
* c-c++-common/goacc/reduction-1.c: Adjust patterns.
* 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/common-block-3.f90: Likewise.
* gfortran.dg/goacc/implied-copy-1.f90: Likewise.
* gfortran.dg/goacc/loop-tree-1.f90: Likewise.
* gfortran.dg/goacc/private-explicit-kernels-1.f95: Likewise.
* gfortran.dg/goacc/private-predetermined-kernels-1.f95: 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.
Co-Authored-By: Thomas Schwinge <tschwinge@baylibre.com>
Co-Authored-By: Sandra Loosemore <sloosemore@baylibre.com>
|
|
This patch reimplements the functionality of the previously-reverted
patch "Assumed-size arrays with non-lexical data mappings". The purpose
is to support implicit uses of assumed-size arrays for Fortran when those
arrays have already been mapped on the target some other way (e.g. by
"acc enter data").
This relates to upstream OpenACC issue 489 (not yet resolved).
gcc/fortran/
* trans-openmp.cc (gfc_omp_finish_clause): Treat implicitly-mapped
assumed-size arrays as zero-sized for OpenACC, rather than an error.
gcc/testsuite/
* gfortran.dg/goacc/assumed-size.f90: Don't expect error.
libgomp/
* testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-1.f90: New
test.
* testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-2.f90: New
test.
|
|
This patch fixes a case revealed by the previous patch where a synthetic
"acc data" region created for a "declare create" variable could interact
strangely with lexical inheritance behaviour. In fact, it doesn't seem
right to create the "acc data" region for allocatable variables at all
-- doing so means that a data region is likely to be created for an
unallocated variable.
The fix is not to add such variables to the synthetic "acc data" region
at all, and defer to the code that performs "enter data"/"exit data"
for them when allocated/deallocated on the host instead. Then, "declare
create" variables are implicitly turned into "present" clauses on in-scope
offload regions.
gcc/fortran/
* trans-openmp.cc (gfc_omp_finish_clause): Handle "declare create" for
scalar allocatable variables.
(gfc_trans_omp_clauses): Don't include allocatable vars in synthetic
"acc data" region created for "declare create" variables. Mark such
variables with the "oacc declare create" attribute instead. Don't
create ALWAYS_POINTER mapping for target-to-host updates of declare
create variables.
(gfc_trans_oacc_declare): Handle empty clause list.
gcc/
* gimplify.cc (gimplify_adjust_omp_clauses_1): Handle "oacc declare
create" attribute.
gcc/testsuite/
* c-c++-common/goacc/readonly-1.c: Adjust patterns.
libgomp/
* testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90:
Remove xfails.
* testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90:
Remove xfails.
* testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90:
Remove xfails.
* testsuite/libgomp.oacc-fortran/declare-create-1.f90: New test.
* testsuite/libgomp.oacc-fortran/declare-create-2.f90: New test.
* testsuite/libgomp.oacc-fortran/declare-create-3.f90: New test.
Co-Authored-By: Paul-Antoine Arras <parras@baylibre.com>
Co-Authored-By: Sandra Loosemore <sandra@baylibre.com>
|
|
This patch reimplements "lexical inheritance" for OpenACC offload regions
inside "data" regions, allowing e.g. this to work:
int *ptr;
[...]
#pragma acc data copyin(ptr[10:2])
{
#pragma acc parallel
{ ... }
}
here, the "copyin" is mirrored on the inner "acc parallel" as
"present(ptr[10:2])" -- allowing code within the parallel to use that
section of the array even though the mapping is implicit.
In terms of implementation, this works by expanding mapping nodes for
"acc data" to include pointer mappings that might be needed by inner
offload regions. The resulting mapping group is then copied to the inner
offload region as needed, rewriting the first node to "force_present".
The pointer mapping nodes are then removed from the "acc data" later
during gimplification.
For OpenMP, pointer mapping nodes on equivalent "omp data" regions are
not needed, so remain suppressed during expansion.
gcc/c-family/
* c-omp.cc (c_omp_address_inspector::expand_array_base): Don't omit
pointer nodes for OpenACC.
gcc/
* gimplify.cc (omp_tsort_mark, omp_mapping_group): Move before
gimplify_omp_ctx. Add constructor to omp_mapping_group.
(gimplify_omp_ctx): Add DECL_DATA_CLAUSE field.
(new_omp_context, delete_omp_context): Initialise and free above field.
(omp_gather_mapping_groups_1): Use constructor for omp_mapping_group.
(gimplify_scan_omp_clauses): Record mappings that might be lexically
inherited. Don't remove
GOMP_MAP_FIRSTPRIVATE_POINTER/GOMP_MAP_FIRSTPRIVATE_REFERENCE yet.
(gomp_oacc_needs_data_present): New function.
(gimplify_adjust_omp_clauses_1): Implement lexical inheritance
behaviour for OpenACC.
(gimplify_adjust_omp_clauses): Remove
GOMP_MAP_FIRSTPRIVATE_POINTER/GOMP_MAP_FIRSTPRIVATE_REFERENCE here
instead, after lexical inheritance is done.
gcc/testsuite/
* c-c++-common/goacc/acc-data-chain.c: New test.
* gfortran.dg/goacc/pr70828.f90: Likewise.
* gfortran.dg/goacc/assumed-size.f90: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test.
* testsuite/libgomp.oacc-c-c++-common/pr70828-2.c: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-3.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-4.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-5.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-6.f90: Likewise.
|
|
(forward ported from devel/omp/gcc-12)
This is a backport of:
https://gcc.gnu.org/pipermail/gcc-patches/2023-May/619003.html
This patch implements '-fopenmp-target=acc', which enables internally handling
a subset of OpenMP target regions as OpenACC parallel regions. This basically
includes target, teams, parallel, distribute, for/do constructs, and atomics.
Essentially, we adjust the internal kinds to OpenACC type, and let OpenACC code
paths handle them, with various needed adjustments throughout middle-end and
nvptx backend. When using this "OMPACC" mode, if there are cases the patch
doesn't handle, it issues a warning, and reverts to normal processing for that
target region.
gcc/ChangeLog:
* builtins.cc (expand_builtin_omp_builtins): New function.
(expand_builtin): Add expand cases for BUILT_IN_GOMP_BARRIER,
BUILT_IN_OMP_GET_THREAD_NUM, BUILT_IN_OMP_GET_NUM_THREADS,
BUILT_IN_OMP_GET_TEAM_NUM, and BUILT_IN_OMP_GET_NUM_TEAMS using
expand_builtin_omp_builtins, enabled under -fopenmp-target=acc.
* cgraphunit.cc (analyze_functions): Add call to
omp_ompacc_attribute_tagging, enabled under -fopenmp-target=acc.
* common.opt (fopenmp-target=): Add new option and enums.
* config/nvptx/mkoffload.cc (main): Handle -fopenmp-target=.
* config/nvptx/nvptx-protos.h (nvptx_expand_omp_get_num_threads): New
prototype.
(nvptx_mem_shared_p): Likewise.
* config/nvptx/nvptx.cc (omp_num_threads_sym): New global static RTX
symbol for number of threads in team.
(omp_num_threads_align): New var for alignment of omp_num_threads_sym.
(need_omp_num_threads): New bool for if any function references
omp_num_threads_sym.
(nvptx_option_override): Initialize omp_num_threads_sym/align.
(write_as_kernel): Disable normal OpenMP kernel entry under OMPACC mode.
(nvptx_declare_function_name): Disable shim function under OMPACC mode.
Disable soft-stack under OMPACC mode. Add generation of neutering init
code under OMPACC mode.
(nvptx_output_set_softstack): Return "" under OMPACC mode.
(nvptx_expand_call): Set parallelism to vector for function calls with
"ompacc for" attached.
(nvptx_expand_oacc_fork): Set mode to GOMP_DIM_VECTOR under OMPACC mode.
(nvptx_expand_oacc_join): Likewise.
(nvptx_expand_omp_get_num_threads): New function.
(nvptx_mem_shared_p): New function.
(nvptx_mach_max_workers): Return 1 under OMPACC mode.
(nvptx_mach_vector_length): Return 32 under OMPACC mode.
(nvptx_single): Add adjustments for OMPACC mode, which have
parallel-construct fork/joins, and regions of code where neutering is
dynamically determined.
(nvptx_reorg): Enable neutering under OMPACC mode when "ompacc for"
attribute is attached to function. Disable uniform-simt when under
OMPACC mode.
(nvptx_file_end): Write __nvptx_omp_num_threads out when needed.
(nvptx_goacc_fork_join): Return true under OMPACC mode.
* config/nvptx/nvptx.h (struct GTY(()) machine_function): Add
omp_parallel_predicate and omp_fn_entry_num_threads_reg fields.
* config/nvptx/nvptx.md (unspecv): Add UNSPECV_GET_TID,
UNSPECV_GET_NTID, UNSPECV_GET_CTAID, UNSPECV_GET_NCTAID,
UNSPECV_OMP_PARALLEL_FORK, UNSPECV_OMP_PARALLEL_JOIN entries.
(nvptx_shared_mem_operand): New predicate.
(gomp_barrier): New expand pattern.
(omp_get_num_threads): New expand pattern.
(omp_get_num_teams): New insn pattern.
(omp_get_thread_num): Likewise.
(omp_get_team_num): Likewise.
(get_ntid): Likewise.
(nvptx_omp_parallel_fork): Likewise.
(nvptx_omp_parallel_join): Likewise.
* expr.cc (expand_expr_real_1): Call expand_var_decl target hook.
* flag-types.h (omp_target_mode_kind): New flag value enum.
* gimplify.cc (struct gimplify_omp_ctx): Add 'bool ompacc' field.
(gimplify_scan_omp_clauses): Handle OMP_CLAUSE__OMPACC_.
(gimplify_adjust_omp_clauses): Likewise.
(gimplify_omp_ctx_ompacc_p): New function.
(gimplify_omp_for): Handle combined loops under OMPACC.
* lto-wrapper.cc (append_compiler_options): Add OPT_fopenmp_target_.
* omp-builtins.def (BUILT_IN_OMP_GET_THREAD_NUM): Remove CONST.
(BUILT_IN_OMP_GET_NUM_THREADS): Likewise.
* omp-expand.cc (remove_exit_barrier): Disable addressable-var
processing for parallel construct child functions under OMPACC mode.
(expand_oacc_for): Add OMPACC mode handling.
(get_target_arguments): Force thread_limit clause value to 1 under
OMPACC mode.
(expand_omp): Under OMPACC mode, avoid child function expanding of
GIMPLE_OMP_PARALLEL.
* omp-general.cc (omp_extract_for_data): Adjustments for OMPACC mode.
* omp-low.cc (struct omp_context): Add 'bool ompacc_p' field.
(scan_sharing_clauses): Handle OMP_CLAUSE__OMPACC_.
(ompacc_ctx_p): New function.
(scan_omp_parallel): Handle OMPACC mode, avoid creating child function.
(scan_omp_target): Tag "ompacc"/"ompacc for" attributes for target
construct child function, remove OMP_CLAUSE__OMPACC_ clauses.
(lower_oacc_head_mark): Handle OMPACC mode cases.
(lower_omp_for): Adjust OMP_FOR kind from OpenMP to OpenACC kinds, add
vector/gang clauses as needed. Add other OMPACC handling.
(lower_omp_taskreg): Add call to lower_oacc_head_tail for OMPACC case.
(lower_omp_target): Do OpenACC gang privatization under OMPACC case.
(lower_omp_teams): Forward OpenACC privatization variables to outer
target region under OMPACC mode.
(lower_omp_1): Do OpenACC gang privatization under OMPACC case for
GIMPLE_BIND.
* omp-offload.cc (ompacc_supported_clauses_p): New function.
(struct target_region_data): New struct type for tree walk.
(scan_fndecl_for_ompacc): New function.
(scan_omp_target_region_r): New function.
(scan_omp_target_construct_r): New function.
(omp_ompacc_attribute_tagging): New function.
(oacc_dim_call): Add OMPACC case handling.
(execute_oacc_device_lower): Make parts explicitly only OpenACC enabled.
(pass_oacc_device_lower::gate): Enable pass under OMPACC mode.
* omp-offload.h (omp_ompacc_attribute_tagging): New prototype.
* opts.cc (finish_options): Only allow -fopenmp-target= when -fopenmp
and no -fopenacc.
* target-insns.def (gomp_barrier): New defined insn pattern.
(omp_get_thread_num): Likewise.
(omp_get_num_threads): Likewise.
(omp_get_team_num): Likewise.
(omp_get_num_teams): Likewise.
* tree-core.h (enum omp_clause_code): Add new OMP_CLAUSE__OMPACC_ entry
for internal clause.
* tree-nested.cc (convert_nonlocal_omp_clauses): Handle
OMP_CLAUSE__OMPACC_.
* tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE__OMPACC_.
* tree.cc (omp_clause_num_ops): Add OMP_CLAUSE__OMPACC_ entry.
(omp_clause_code_name): Likewise.
* tree.h (OMP_CLAUSE__OMPACC__FOR): New macro for OMP_CLAUSE__OMPACC_.
libgomp/ChangeLog:
* config/nvptx/team.c (__nvptx_omp_num_threads): New global variable in
shared memory.
(cherry picked from commit 5f881613fa9128edae5bbfa4e19f9752809e4bd7)
|
|
OpenMP 5.0 permits to use arrays with strides and derived
type components for the list items to the 'from'/'to' clauses
of the 'target update' directive.
Partially committed to mainline as:
6629444170f85 OpenMP/Fortran: 'target update' with DT components
This patch contains the differences to the mainline version.
gcc/fortran/ChangeLog:
* openmp.cc (resolve_omp_clauses): Apply to OpenMP target update.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/target-13.f90: Update test.
|
|
This patch disables vectorization of memory accesses to non-default address
spaces where the pointer size is different to the usual pointer size. This
condition typically occurs in OpenACC programs on amdgcn, where LDS memory is
used for broadcasting gang-private variables between threads. In particular,
see libgomp.oacc-c-c++-common/private-variables.c
The problem is that the address space information is dropped from the various
types in the middle-end and eventually it triggers an ICE trying to do an
address conversion. That ICE can be avoided by defining
POINTERS_EXTEND_UNSIGNED, but that just produces wrong RTL code later on.
A correct solution would ensure that all the vectypes have the correct address
spaces, but I don't have time for that right now.
gcc/ChangeLog:
* tree-vect-data-refs.cc (vect_analyze_data_refs): Workaround an
address-space bug.
|
|
This adds middle end support for uses_allocators, wires Fortran to use it and
add C/C++ parsing support.
gcc/ChangeLog:
* builtin-types.def (BT_FN_VOID_PTRMODE): Add.
(BT_FN_PTRMODE_PTRMODE_INT_PTR): Add.
* gimplify.cc (gimplify_bind_expr): Diagnose missing
uses_allocators clause.
(gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses,
gimplify_omp_workshare): Handle uses_allocators.
* omp-builtins.def (BUILT_IN_OMP_INIT_ALLOCATOR,
BUILT_IN_OMP_DESTROY_ALLOCATOR): Add.
* omp-low.cc (scan_sharing_clauses): Handle OMP_CLAUSE_USES_ALLOCATORS
and OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR clauses.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_USES_ALLOCATORS.
* tree.cc (omp_clause_num_ops, omp_clause_code_name): Likewise.
* tree-pretty-print.cc (dump_omp_clause): Handle it.
* tree.h (OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR,
OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE,
OMP_CLAUSE_USES_ALLOCATORS_TRAITS): New.
gcc/c-family/ChangeLog:
* c-omp.cc (c_omp_split_clauses): Hande uses_allocators.
* c-pragma.h (enum pragma_omp_clause): Add
PRAGMA_OMP_CLAUSE_USES_ALLOCATORS.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_clause_uses_allocators): New.
(c_parser_omp_clause_name, c_parser_omp_all_clauses,
OMP_TARGET_CLAUSE_MASK): Handle uses_allocators.
* c-typeck.cc (c_finish_omp_clauses): Likewise.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_clause_uses_allocators): New.
(cp_parser_omp_clause_name, cp_parser_omp_all_clauses,
OMP_TARGET_CLAUSE_MASK): Handle uses_allocators.
* semantics.cc (finish_omp_clauses): Likewise.
gcc/fortran/ChangeLog:
* trans-array.cc (gfc_conv_array_initializer): Set PURPOSE
when building constructor for get_initialized_tmp_var.
* trans-openmp.cc (gfc_trans_omp_clauses): Handle uses_allocators.
* types.def (BT_FN_VOID_PTRMODE, BT_FN_PTRMODE_PTRMODE_INT_PTR): Add.
libgomp/ChangeLog:
* testsuite/libgomp.c++/c++.exp (check_effective_target_c,
check_effective_target_c++): Add.
* testsuite/libgomp.c/c.exp (check_effective_target_c,
check_effective_target_c++): Add.
* testsuite/libgomp.fortran/uses_allocators_2.f90: Remove 'sorry'.
* testsuite/libgomp.c-c++-common/uses_allocators-1.c: New test.
* testsuite/libgomp.c-c++-common/uses_allocators-2.c: New test.
* testsuite/libgomp.c-c++-common/uses_allocators-3.c: New test.
* testsuite/libgomp.c-c++-common/uses_allocators-4.c: New test.
* testsuite/libgomp.fortran/uses_allocators_3.f90: New test.
* testsuite/libgomp.fortran/uses_allocators_4.f90: New test.
* testsuite/libgomp.fortran/uses_allocators_5.f90: New test.
* testsuite/libgomp.fortran/uses_allocators_6.f90: New test.
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/allocate-1.f90: Add uses_allocators.
* gfortran.dg/gomp/scope-6.f90: Update dg-scan-tree-dump.
* c-c++-common/gomp/uses_allocators-1.c: New test.
* c-c++-common/gomp/uses_allocators-2.c: New test.
* gfortran.dg/gomp/uses_allocators-1.f90: New test.
|
|
GIMPLE_DEBUG were put in a parallel region of its own, which is not
only pointless but also breaks -fcompare-debug. With this commit,
they are handled like simple assignments: those placed are places
into the same body as the loop such that only one parallel region
remains as without debugging. This fixes the existing testcase
libgomp.oacc-c-c++-common/kernels-loop-g.c.
Note: GIMPLE_DEBUG are only accepted with -fcompare-debug; if they
appear otherwise, decompose_kernels_region_body rejects them with
a sorry (unchanged).
Also note that there are still many xfailed tests in the
c-c++-common/goacc/kernels-decompose-pr* testcases that were added
in mainline commit c14ea6a72fb1ae66e3d32ac8329558497c6e4403.
gcc/ChangeLog
* omp-oacc-kernels-decompose.cc (top_level_omp_for_in_stmt,
decompose_kernels_region_body): Handle GIMPLE_DEBUG like
simple assignment.
gcc/testsuite/ChangeLog
* c-c++-common/goacc/kernels-decompose-pr103836-1-2.c: Adjust xfails.
* c-c++-common/goacc/kernels-decompose-pr103836-1-3.c: Likewise.
* c-c++-common/goacc/kernels-decompose-pr103836-1-4.c: Likewise.
* c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: Likewise.
* c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Likewise.
* c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise.
Co-Authored-By: Sandra Loosemore <sloosemore@baylibre.com>
|
|
https://patchwork.sourceware.org/project/gcc/list/?series=35022
Implement the -foffload-memory=pinned option such that libgomp is
instructed to enable fully-pinned memory at start-up. The option is
intended to provide a performance boost to certain offload programs without
modifying the code.
This feature only works on Linux, at present, and simply calls mlockall to
enable always-on memory pinning. It requires that the ulimit feature is
set high enough to accommodate all the program's memory usage.
In this mode the ompx_gnu_pinned_memory_alloc feature is disabled as it is not
needed and may conflict.
gcc/ChangeLog:
* omp-builtins.def (BUILT_IN_GOMP_ENABLE_PINNED_MODE): New.
* omp-low.cc (omp_enable_pinned_mode): New function.
(execute_lower_omp): Call omp_enable_pinned_mode.
libgomp/ChangeLog:
* config/linux/allocator.c (always_pinned_mode): New variable.
(GOMP_enable_pinned_mode): New function.
(linux_memspace_alloc): Disable pinning when always_pinned_mode set.
(linux_memspace_calloc): Likewise.
(linux_memspace_free): Likewise.
(linux_memspace_realloc): Likewise.
* libgomp.map: Add GOMP_enable_pinned_mode.
* testsuite/libgomp.c/alloc-pinned-7.c: New test.
* testsuite/libgomp.c-c++-common/alloc-pinned-1.c: New test.
|
|
https://patchwork.sourceware.org/project/gcc/list/?series=35022
Add a new option. It's inactive until I add some follow-up patches.
gcc/ChangeLog:
* common.opt: Add -foffload-memory and its enum values.
* coretypes.h (enum offload_memory): New.
* doc/invoke.texi: Document -foffload-memory.
|
|
This sets the type precision of the collapsed iterator variable to the
sum of the precision of the collapsed loop variables, up to a maximum of
sizeof(long long) (i.e. 64-bits).
gcc/ChangeLog
* omp-expand.cc (expand_oacc_for): Convert .tile variable to
diff_type before multiplying.
* omp-general.cc (omp_extract_for_data): Use accumulated precision
of all collapsed for-loops as precision of iteration variable, up
to the precision of a long long.
libgomp/ChangeLog
* testsuite/libgomp.c-c++-common/collapse-4.c: New.
* testsuite/libgomp.fortran/collapse5.f90: New.
|
|
This is a merge of:
https://gcc.gnu.org/pipermail/gcc-patches/2020-December/562081.html
This patch now allows multiple clauses on the same construct to map
the same variable, which was not valid in OpenMP 4.5, but now allowed
in 5.0.
This may possibly reverted/updated when a final patch is approved
for mainline.
gcc/c/ChangeLog
* c-typeck.cc (c_finish_omp_clauses): Adjust to allow duplicate
mapped variables for OpenMP.
gcc/cp/ChangeLog
* semantics.cc (finish_omp_clauses): Adjust to allow duplicate
mapped variables for OpenMP.
gcc/ChangeLog
* omp-low.cc (install_var_field): Add new 'tree key_expr = NULL_TREE'
default parameter. Set splay-tree lookup key to key_expr instead of
var if key_expr is non-NULL. Adjust call to install_parm_decl.
Update comments.
(scan_sharing_clauses): Use clause tree expression as splay-tree key
for map/to/from and OpenACC firstprivate cases when installing the
variable field into the send/receive record type.
(maybe_lookup_field_in_outer_ctx): Add code to search through
construct clauses instead of entirely based on splay-tree lookup.
(lower_oacc_reductions): Adjust to find map-clause of reduction
variable, then create receiver-ref.
(lower_omp_target): Adjust to lookup var field using clause expression.
gcc/testsuite/ChangeLog
* c-c++-common/gomp/clauses-2.c: Adjust testcase.
* c-c++-common/gomp/map-6.c: Adjust testcase.
Co-Authored-By: Paul-Antoine Arras <parras@baylibre.com>
|
|
Add DWARF address class attributes for variables that exist outside the
generic address space. In particular, this is the case for gang-private
variables in OpenACC offload kernels.
gcc/ChangeLog:
* dwarf2out.cc (add_location_or_const_value_attribute): Set
DW_AT_address_class, if appropriate.
|
|
Ensure that the parent DWARF subprograms of offload kernel functions have a
code range, and are therefore not discarded by GDB. This is only necessary
when the parent function does not actually exist in the final binary, which is
commonly the case within the offload device's binary.
gcc/
* dwarf2out.cc (notional_parents_list): New file variable.
(gen_subprogram_die): Record offload kernel functions in
notional_parents_list.
(fixup_notional_parents): New function.
(dwarf2out_finish): Call fixup_notional_parents.
(dwarf2out_c_finalize): Reset notional_parents_list.
|
|
This patch adjusts OpenACC loop lowering in the AMD GCN target compiler
in such a way that the autovectorizer can vectorize the "vector" dimension
of those loops in more cases.
Rather than generating "SIMT" code that executes a scalar instruction
stream for each lane of a vector in lockstep, for GCN we model the GPU
like a typical CPU, with separate instructions to operate on scalar and
vector data. That means that unlike other offload targets, we rely on
the autovectorizer to handle the innermost OpenACC parallelism level,
which is "vector".
Because of this, the OpenACC builtin functions to return the current
vector lane and the vector width return 0 and 1 respectively, despite
the native vector width being 64 elements wide.
This allows generated code to work with our chosen compilation model,
but the way loops are lowered in omp-offload.c:oacc_xform_loop does not
understand the discrepancy between logical (OpenACC) and physical vector
sizes correctly. That means that if a loop is partitioned over e.g. the
worker AND vector dimensions, we actually lower with unit vector size --
meaning that if we then autovectorize, we end up trying to vectorize
over the "worker" dimension rather than the vector one! Then, because
the number of workers is not fixed at compile time, that means the
autovectorizer has a hard time analysing the loop and thus vectorization
often fails entirely.
We can fix this by deducing the true vector width in oacc_xform_loop,
and using that when we are on a "non-SIMT" offload target. We can then
rearrange how loops are lowered in that function so that the loop form
fed to the autovectorizer is more amenable to vectorization -- namely,
the innermost step is set to process each loop iteration sequentially.
For some benchmarks, allowing vectorization to succeed leads to quite
impressive performance improvements -- I've observed between 2.5x and
40x on one machine/GPU combination.
The low-level builtins available to user code (__builtin_goacc_parlevel_id
and __builtin_goacc_parlevel_size) continue to return 0/1 respectively
for the vector dimension for AMD GCN, even if their containing loop is
vectorized -- that's a quirk that we might possibly want to address at
some later date.
Only non-"chunking" loops are handled at present. "Chunking" loops are
still lowered as before.
gcc/ChangeLog
* omp-offload.cc (oacc_thread_numbers): Add VF_BY_VECTORIZER parameter.
Add overloaded wrapper for previous arguments & behaviour.
(oacc_xform_loop): Lower vector loops to iterate a multiple of
omp_max_vf times over contiguous steps on non-SIMT targets.
libgomp/ChangeLog
* testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: Adjust for loop
lowering changes.
* testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c: Likewise.
|
|
This patch incorporates these commits from OG14 branch:
65be1389eeda9b3b97f6587721215c3f31bd7f98
9d43e819d88f97c7ade7f8c95c35ea3464ea7771
f2cf2b994c4d8c871fad5502ffb9aaee9ea4f4e0
2770ce41615557e595065ce0c5db71e9f3d82b0a
a29e58f4b314862a72730119f85e9125879abf0b
ffd990543f805ed448aaa355d190f37103f8f1f0
gcc/ChangeLog
* gimplify.cc (omp_group_base): Handle GOMP_MAP_DECLARE_ALLOCATE
and GOMP_MAP_DECLARE_DEALLOCATE.
(gimplify_adjust_omp_clauses): Likewise.
* 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. Use
VIEW_CONVERT also for vectors.
(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/ChangeLog
* 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_omp_check_optional_argument): Handle non-decl
case.
(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/ChangeLog
* gfortran.dg/goacc/declare-allocatable-1.f90: New test.
* gfortran.dg/goacc/declare-3.f95: Adjust expected dump output.
include/ChangeLog
* gomp-constants.h (enum gomp_map_kind): Define
GOMP_MAP_DECLARE_{ALLOCATE,DEALLOCATE} and GOMP_MAP_FLAG_SPECIAL_4.
libgomp/ChangeLog
* libgomp.h (gomp_acc_declare_allocate): Remove prototype.
* oacc-mem.c (gomp_acc_declare_allocate): New function.
(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. Unlock mutex before calling
gomp_acc_declare_allocate and relock it afterwards.
(goacc_exit_data_internal): Unlock device mutex around
gomp_acc_declare_allocate call. Pass new pointer argument. Handle
group pointer mapping for deallocate.
* testsuite/libgomp.oacc-fortran/allocatable-scalar.f90: New test.
* testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90:
Adjust.
* testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90:
Likewise.
* testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90:
Likewise.
* 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.
* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90:
Adjust.
* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90:
Likewise.
* testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1.f90:
New test.
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Paul-Antoine Arras <parras@baylibre.com>
|
|
Combination of OG14 commits
141a592bf147c91c28de7864fa12259687e827e3
8d7562192cc814c6d0d48b424d7751762871a37b
+ new testsuite fixes to add xfails for tests that already failed on OG14.
gcc/ChangeLog
* gimplify.cc (localize_reductions): Rewrite references for
OMP_CLAUSE_PRIVATE also. Do not create local variable for
privatized arrays as the size is not directly known by the type.
gcc/testsuite/ChangeLog
* gfortran.dg/goacc/privatization-1-compute-loop.f90: Add xfails.
* gfortran.dg/goacc/privatization-1-compute.f90: Likewise.
libgomp/ChangeLog
* testsuite/libgomp.oacc-c++/privatized-ref-3.C: Add xfails.
* testsuite/libgomp.oacc-fortran/optional-private.f90: Likewise.
* testsuite/libgomp.oacc-fortran/privatized-ref-1.f95: Likewise.
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>
Co-Authored-By: Sandra Loosemore <sandra@baylibre.com>
|
|
gcc/ChangeLog
* gimplify.cc (privatize_reduction): New struct.
(localize_reductions_r, localize_reductions): New functions.
(gimplify_omp_for): Call localize_reductions.
(gimplify_omp_workshare): Likewise.
* omp-low.cc (lower_oacc_reductions): Handle localized reductions.
Create fewer temp vars.
* tree-core.h (omp_clause_code): Add OMP_CLAUSE_REDUCTION_PRIVATE_DECL
documentation.
* tree.cc (omp_clause_num_ops): Bump number of ops for
OMP_CLAUSE_REDUCTION to 6.
(walk_tree_1): Adjust accordingly.
* tree.h (OMP_CLAUSE_REDUCTION_PRIVATE_DECL): Add macro.
gcc/testsuite/ChangeLog
* gfortran.dg/goacc/modules.f95: Remove xfail on bogus warnings.
libgomp/ChangeLog
* testsuite/libgomp.oacc-fortran/optional-reduction.f90: Remove
xfail on bogus warnings.
* testsuite/libgomp.oacc-fortran/parallel-reduction.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70643.f90: Likewise.
* testsuite/libgomp.oacc-fortran/reduction-5.f90: Likewise.
* testsuite/libgomp.oacc-fortran/reduction-7.f90: Likewise.
* testsuite/libgomp.oacc-fortran/reference-reductions.f90: Likewise.
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Sandra Loosemore <sloosemore@baylibre.com>
|
|
gcc/ChangeLog
* gimplify.cc (omp_add_variable): Enable firstprivate reduction
variables.
gcc/testsuite/ChangeLog
* c-c++-common/goacc/reduction-10.c: New test.
libgomp/ChangeLog
* testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c: New
test.
* testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c: New
test.
Co-Authored-By: Chung-Lin Tang <cltang@codesourcery.com>
|
|
gcc/ChangeLog
* omp-low.cc (lower_oacc_head_mark): Don't mark OpenACC auto
loops as independent inside acc parallel regions.
gcc/testsuite/ChangeLog
* c-c++-common/goacc/loop-auto-1.c: Adjust test case to conform to
the new behavior of the auto clause in OpenACC 2.5.
* c-c++-common/goacc/loop-auto-2.c: Likewise.
* gcc.dg/goacc/loop-processing-1.c: Likewise.
* c-c++-common/goacc/loop-auto-3.c: New test.
* gfortran.dg/goacc/loop-auto-1.f90: New test.
libgomp/ChangeLog
* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust test case
to conform to the new behavior of the auto clause in OpenACC 2.5.
|
|
Incorporates these commits from OG14 branch:
a743b0947593f38fd93fa1bc5f8dd3c50ba5c498
6f759e76f00d70bf1d44a25767c73ec2de855452
61d66e2e494655a34cca7289ae584d2644f1a65f
gcc/ChangeLog
* omp-low.cc (maybe_lookup_field_in_outer_ctx): New function.
(convert_to_firstprivate_int): New function.
(convert_from_firstprivate_int): New function.
(lower_omp_target): Enable GOMP_MAP_FIRSTPRIVATE_INT in OpenACC.
Remove unused variable.
libgomp/ChangeLog
* testsuite/libgomp.oacc-c++/firstprivate-int.C: New test.
* testsuite/libgomp.oacc-c-c++-common/firstprivate-int.c: New
test.
* testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: XFAIL
execution test.
* testsuite/libgomp.oacc-fortran/firstprivate-int.f90: New test.
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>
Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
|
|
This patch was previously combined with an unrelated change and testcases
that was upstreamed separately ("Add '-Wopenacc-parallelism'",
22cff118f7526bec195ed6e41452980820fdf3a8).
gcc/ChangeLog
* doc/invoke.texi (fopenacc-dim): Document syntax for using
runtime value from environment variable.
* omp-offload.cc (oacc_parse_default_dims): Implement it.
libgomp/ChangeLog
* testsuite/libgomp.oacc-c-c++-common/loop-default-compile.c: New.
Co-Authored-By: Tom de Vries <tdevries@suse.de>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Julian Brown <julian@codesourcery.com>
|
|
gcc/testsuite/ChangeLog
* g++.dg/goacc/loop-1.c: New test.
* g++.dg/goacc/loop-2.c: New test.
* g++.dg/goacc/loop-3.c: New test.
libgomp/ChangeLog
* testsuite/libgomp.oacc-fortran/data-3.f90: Update parallel
regions to denote variables copyied in via acc enter data as
present.
* testsuite/libgomp.oacc-c-c++-common/subr.h: Reimplement.
* testsuite/libgomp.oacc-c-c++-common/subr.ptx: Regenerated PTX.
* testsuite/libgomp.oacc-c-c++-common/timer.h: Removed.
* testsuite/libgomp.oacc-c-c++-common/lib-69.c: Change async checks.
* testsuite/libgomp.oacc-c-c++-common/lib-70.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-72.c: Rework kernel i/f and
change async checks.
* testsuite/libgomp.oacc-c-c++-common/lib-73.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-74.c: Rework kernel i/f and
timing checks.
* testsuite/libgomp.oacc-c-c++-common/lib-75.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-76.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-78.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-79.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-81.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-93.c: New test.
Co-Authored-By: James Norris <jnorris@codesourcery.com>
Co-Authored-By: Tom de Vries <tom@codesourcery.com>
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Tobias Burnus <tburnus@baylibre.com>
|
|
gcc/testsuite/ChangeLog
* c-c++-common/goacc/reduction-9.c: New.
* g++.dg/goacc/reductions-1.C: New.
* gcc.dg/goacc/loop-processing-1.c: Update.
libgomp/ChangeLog
* testsuite/libgomp.oacc-c-c++-common/par-reduction-3.c: New.
* testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c: New.
* testsuite/libgomp.oacc-fortran/reduction-9.f90: New.
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
Co-Authored-By: Nathan Sidwell <nathan@acm.org>
|
|
gcc/ChangeLog
* config/nvptx/nvptx.cc (nvptx_propagate_unified): New.
(nvptx_split_blocks): Call it for cond_uni insn.
(nvptx_expand_cond_uni): New.
(enum nvptx_builtins): Add NVPTX_BUILTIN_COND_UNI.
(nvptx_init_builtins): Initialize it.
(nvptx_expand_builtin): Handle NVPTX_BUILTIN_COND_UNI.
(nvptx_generate_vector_shuffle): Change integral SHIFT operand to
tree BITS operand.
(nvptx_vector_reduction): New.
(nvptx_adjust_reduction_type): New.
(nvptx_goacc_reduction_setup): Use it to adjust the type of ref_to_res.
(nvptx_goacc_reduction_init): Don't update LHS if it doesn't exist.
(nvptx_goacc_reduction_fini): Call nvptx_vector_reduction for vector.
Use it to adjust the type of ref_to_res.
(nvptx_goacc_reduction_teardown): Call nvptx_adjust_reduction_type.
* config/nvptx/nvptx.md (cond_uni): New pattern.
* gimplify.cc (gimplify_adjust_omp_clauses): Add DECL_P check
for OMP_CLAUSE_TASK_REDUCTION.
* omp-low.cc (lower_oacc_reductions): Handle
GOMP_MAP_FIRSTPRIVATE_POINTER.
* omp-offload.cc (default_goacc_reduction): Likewise.
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
|
|
gcc/c/ChangeLog
* c-parser.cc (c_parser_omp_variable_list): New c_omp_region_type
argument. Use it to specialize handling of OMP_CLAUSE_REDUCTION for
OpenACC.
(c_parser_omp_var_list_parens): Add region-type argument to call.
(c_parser_oacc_data_clause): Likewise.
(c_parser_oacc_data_clause_deviceptr): Likewise.
(c_parser_omp_clause_reduction): Change is_omp boolean parameter to
c_omp_region_type. Update call to c_parser_omp_variable_list.
(c_parser_omp_clause_map): Update call to
c_parser_omp_variable_list.
(c_parser_omp_clause_from_to): Likewise.
(c_parser_omp_clause_init): Likewise.
(c_parser_oacc_all_clauses): Update calls to
c_parser_omp_clause_reduction.
(c_parser_omp_all_clauses): Likewise.
(c_parser_oacc_cache): Update call to c_parser_omp_variable_list.
* c-typeck.cc (c_finish_omp_clauses): Emit an error on orphan OpenACC
gang reductions. Suppress user-defined reduction error for OpenACC.
gcc/cp/ChangeLog
* parser.cc (cp_parser_omp_var_list_no_open): New c_omp_region_type
argument. Use it to specialize handling of OMP_CLAUSE_REDUCTION for
OpenACC.
(cp_parser_omp_var_list): Add c_omp_region_type argument. Update call
to cp_parser_omp_var_list_no_open.
(cp_parser_oacc_data_clause): Update call to
cp_parser_omp_var_list_no_open.
(cp_parser_omp_clause_reduction): Change is_omp boolean parameter to
c_omp_region_type. Update call to cp_parser_omp_var_list_no_open.
(cp_parser_omp_clause_from_to): Update call to
cp_parser_omp_clause_var_list_no_open.
(cp_parser_omp_clause_map): Likewise.
(cp_parser_omp_clause_init): Likewise.
(cp_parser_oacc_all_clauses): Update call to
cp_parser_omp_clause_reduction.
(cp_parser_omp_all_clauses): Likewise.
* semantics.cc (finish_omp_reduction_clause): Add c_omp_region_type
argument. Suppress user-defined reduction error for OpenACC.
(finish_omp_clauses): Update call to finish_omp_reduction_clause.
gcc/fortran/ChangeLog
* trans-openmp.cc (gfc_omp_clause_copy_ctor): Permit reductions.
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
Co-Authored-By: Nathan Sidwell <nathan@acm.org>
Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
|
|
This is a merge of these OG14 commits:
987bb4c25b4076eb54f043644bdf9988378be90d
9e8395708c0027ad1de871bae870c4b0185a74fd
2adb0ec35cd47b34d47c961f6ae46089e3e02cbc
4d29174a9602e6ea783ba0e9a7b1e38fb6913db5
gcc/fortran/ChangeLog
* openmp.cc (gfc_match_omp_map_clause): Re-write handling of the
deviceptr clause. Add new common_blocks argument. Propagate it to
gfc_match_omp_variable_list.
(gfc_match_omp_clauses): Update calls to gfc_match_omp_map_clauses.
(resolve_positive_int_expr): Promote the warning to an error.
(check_array_not_assumed): Remove pointer check.
(resolve_oacc_nested_loops): Error on do concurrent loops.
* trans-openmp.cc (gfc_omp_finish_clause): Don't create pointer data
mappings for deviceptr clauses.
(gfc_trans_omp_clauses): Likewise.
gcc/ChangeLog
* gimplify.cc (enum gimplify_omp_var_data): Add GOVD_DEVICETPR.
(oacc_default_clause): Privatize fortran common blocks.
(omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate.
Defer the expansion of DECL_VALUE_EXPR for common block decls.
(gimplify_scan_omp_clauses): Add GOVD_DEVICEPTR attribute when
appropriate.
(gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for
implicit deviceptr mappings.
gcc/testsuite/ChangeLog
* c-c++-common/goacc/deviceptr-4.c: Update.
* gfortran.dg/goacc/loop-2-kernels-tile.f95: Update.
* gfortran.dg/goacc/loop-2-parallel-tile.f95: Update.
* gfortran.dg/goacc/loop-2-serial-tile.f95: Update.
* gfortran.dg/goacc/sie.f95: Update.
* gfortran.dg/goacc/tile-1.f90: Update.
* gfortran.dg/gomp/num-teams-2.f90: Update.
* gfortran.dg/gomp/pr67500.f90: Update.
* gfortran.dg/gomp/pr77516.f90: Update.
libgomp/ChangeLog
* oacc-parallel.c (GOACC_parallel_keyed): Handle Fortran deviceptr
clause.
(GOACC_data_start): Likewise.
* testsuite/libgomp.oacc-fortran/deviceptr-1.f90: New test.
Co-Authored-By: James Norris <jnorris@codesourcery.com>
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>
Co-Authored-By: Thomas Schwinge <tschwinge@baylibre.com>
|
|
This is based on OG14 commit b143c1c447945ce05903ff1360ead97774dfce4b,
which was based from v4, posted upstream here:
https://gcc.gnu.org/pipermail/gcc-patches/2020-April/543437.html
It also incorporates a number of follow-up bug and bit-rot fixes, OG14
commits
e11726d3467543de45448097dde27ba34bf04bfe
87ea4de1c4a360d5d62357491a41811213f4528c
151fc161d0ed640048444ca18f9325e3d2e03e99
628a000bdbf63252c2ede13ccab8e99a19769866
11263c048d39ab1d6a11067b18674bf8307bbbf5
8c1068bbe3e52529bede5466a43af8d98f38dac2
gcc/c/ChangeLog
PR other/76739
* 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. Reject
non-DECL base-pointer cases as unsupported.
(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
handle non-contiguous case to create dynamic array map.
gcc/cp/ChangeLog
PR other/76739
* 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. Reject
non-DECL base-pointer cases as unsupported.
(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
handle non-contiguous case to create dynamic array map.
gcc/fortran/ChangeLog
PR other/76739
* f95-lang.cc (DEF_FUNCTION_TYPE_VAR_5): New symbol.
* types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.
gcc/ChangeLog
PR other/76739
* builtin-types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.
* gimplify.cc (omp_group_base): Handle GOMP_MAP_NONCONTIG_ARRAY_*.
(gimplify_scan_omp_clauses): Handle OMP_TARGET_UPDATE.
(gimplify_adjust_omp_clauses): Skip gimplification of
OMP_CLAUSE_SIZE of non-contiguous array maps (which is a TREE_LIST).
* omp-builtins.def (BUILT_IN_GOACC_DATA_START): Adjust function type
to new BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR.
* 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.
* tree-pretty-print.cc (dump_omp_clause): Handle
GOMP_MAP_NONCONTIG_ARRAY_*.
gcc/testsuite/ChangeLog
PR other/76739
* c-c++-common/goacc/data-clause-1.c (foo): Remove expected message.
* c-c++-common/goacc/noncontig_array-1.c: New test.
* g++.dg/goacc/data-clause-1.C (foo): Remove expected message.
include/ChangeLog
PR other/76739
* 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.
libgomp/ChangeLog
PR other/76739
* libgomp.h (gomp_map_vars_openacc): New function declaration.
* libgomp_g.h (GOACC_data_start): Add variadic '...' to 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.
Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
Co-Authored-By: Paul-Antoine Arras <parras@baylibre.com>
|
|
gcc/ChangeLog
* Makefile.in (REVISION_s): Change default message.
|
|
* zh_CN.po: Update.
|
|
gcc/fortran/ChangeLog:
PR fortran/120139
* dump-parse-tree.cc (get_c_type_name): If no constant
size of an array exists, output an asterisk.
(cherry picked from commit 4f9c7b5258f2af89bba8e954c277981d2e2ee1ef)
|
|
gcc/fortran/ChangeLog:
PR fortran/120107
* dump-parse-tree.cc (write_type): Do not dump non-interoperable
types.
(cherry picked from commit fa0dff8e99e81bc7a3db1dc57d4fc340e0525b1d)
|
|
|
|
* sv.po: Update.
|
|
For a debugger to display statically-allocated[0] TLS variables the compiler
must communicate information[1] that can be used in conjunction with knowledge
of the runtime enviroment[2] to calculate a location for the variable for
each thread. That need gives rise to dw_loc_dtprel in dwarf2out, a flag tracking
whether the location description is dtprel, or relative to the
"dynamic thread pointer". Location descriptions in the .debug_info section for
TLS variables need to be relocated by the static linker accordingly, and
dw_loc_dtprel controls emission of the needed relocations.
This is further complicated by -gsplit-dwarf. -gsplit-dwarf is designed to allow
as much debugging information as possible to bypass the static linker to improve
linking performance. One of the ways that is done is by introducing a layer of
indirection for relocatable values[3]. That gives rise to addr_index_table which
ultimately results in the .debug_addr section.
While the code handling addr_index_table clearly contemplates the existence of
dtprel entries[4] resolve_addr_in_expr does not, and the result is that when
using -gsplit-dwarf the DWARF for TLS variables contains an address[5] rather
than an offset, and debuggers can't work with that.
This is visible on a trivial example. Compile
```
static __thread int tls_var;
int main(void) {
tls_var = 42;
return 0;
}
```
with -g and -g -gsplit-dwarf. Run the program under gdb. When examining the
value of tls_var before and after the assignment, -g behaves as one would
expect but -g -gsplit-dwarf does not. If the user is lucky and the miscalculated
address is not mapped, gdb will print "Cannot access memory at address ...".
If the user is unlucky and the miscalculated address is mapped, gdb will simply
give the wrong value. You can further confirm that the issue is the address
calculation by asking gdb for the address of tls_var and comparing that to what
one would expect.[6]
Thankfully this is trivial to fix by modifying resolve_addr_in_expr to propagate
the dtprel character of the location where necessary. gdb begins working as
expected and the diff in the generated assembly is clear.
```
.section .debug_addr,"",@progbits
.long 0x14
.value 0x5
.byte 0x8
.byte 0
.Ldebug_addr0:
- .quad tls_var
+ .long tls_var@dtpoff, 0
.quad .LFB0
```
[0] Referring to e.g. __thread as statically-allocated vs. e.g. a
dynamically-allocated pthread_key_create() call.
[1] Generally an offset in a TLS block.
[2] With glibc, provided by libthread_db.so.
[3] Relocatable values are moved to a table in the .debug_addr section, those
values in .debug_info are replaced with special values that look up indexes
in that table, and then the static linker elsewhere assigns a single per-CU
starting index in the .debug_addr section, allowing those special values to
remain permanently fixed and the resulting data to be ignored by the linker.
[4] ate_kind_rtx_dtprel exists, after all, and new_addr_loc_descr does produce
it where appropriate.
[5] e.g. an address in the .tbss/.tdata section.
[6] e.g. on x86-64 by examining %fsbase and the offset in the assembly
2025-05-01 Kyle Huey <me@kylehuey.com>
* dwarf2out.cc (resolve_addr_in_expr): Propagate dtprel into the address
table when appropriate.
|
|
instantiations [PR120125]
The attached testcase ICEs in maybe_thunk_body because we haven't
created a node in the cgraph for an imported explicit instantiation yet.
We in fact really shouldn't be emitting calls at all, since an imported
explicit instantiation always exists in the TU we imported it from. So
this patch adjusts DECL_NOT_REALLY_EXTERN handling to account for this.
PR c++/120125
gcc/cp/ChangeLog:
* module.cc (trees_out::write_function_def): Only set
DECL_NOT_REALLY_EXTERN if the importer might need to emit it.
* optimize.cc (maybe_thunk_body): Don't assume 'fn' has a cgraph
node created.
gcc/testsuite/ChangeLog:
* g++.dg/modules/clone-4_a.C: New test.
* g++.dg/modules/clone-4_b.C: New test.
Signed-off-by: Nathaniel Shead <nathanieloshead@gmail.com>
Reviewed-by: Jason Merrill <jason@redhat.com>
(cherry picked from commit d787bc4fd372298e9ed5b11cb3050fd3707070f6)
|
|
In r15-2799-gf1bfba3a9b3f31, a new kind of global constructor was added.
Unfortunately this broke C++20 modules, as both the host and target
constructors were given the same mangled name. This patch ensures that
only the host constructor gets the module name mangling for now, and
stops forcing the creation of the target constructor even when no such
initialization is required.
PR c++/119864
gcc/cp/ChangeLog:
* decl2.cc (start_objects): Only use module initialized for
host.
(c_parse_final_cleanups): Don't always create an OMP offload
init function in modules.
gcc/testsuite/ChangeLog:
* g++.dg/modules/openmp-1.C: New test.
Signed-off-by: Nathaniel Shead <nathanieloshead@gmail.com>
Reviewed-by: Jason Merrill <jason@redhat.com>
(cherry picked from commit 79b7e37ea3fbbc43958190f69f6da3be3d809c9c)
|
|
|
|
This bug was another case of generating a formal arglist from
an actual one where we should not have done so. The fix is
straightforward: If we have resolved the formal arglist, we should
not generare a new one.
OK for trunk and backport?
gcc/fortran/ChangeLog:
PR fortran/120163
* gfortran.h: Add formal_resolved to gfc_symbol.
* resolve.cc (gfc_resolve_formal_arglist): Set it.
(resolve_function): Do not call gfc_get_formal_from_actual_arglist
if we already resolved a formal arglist.
(resolve_call): Likewise.
gcc/testsuite/ChangeLog:
PR fortran/120163
* gfortran.dg/interface_61.f90: New test.
|
|
aliased set type
The ICE occurred because of a bug in M2GenGCC.mod:FoldBecomes which
attempted to remove the same quadruple twice. Thereafter cc1gm2
generated an erroneous error type error as PCSymBuild did not skip
an aliased set type. The type of the const was set incorrectly
(as a set type) rather than the type of the set element.
gcc/m2/ChangeLog:
PR modula2/120117
* gm2-compiler/M2GenGCC.mod (FoldBecomes): Remove the call to
RemoveQuad since this is performed by TypeCheckBecomes.
* gm2-compiler/PCSymBuild.mod (buildConstFunction): Rewrite
header comment.
Check for a set or a type aliased set and appropriately
skip type equivalences and obtain the element type.
* gm2-compiler/SymbolTable.mod (PutConst): Add call to
CheckBreak.
gcc/testsuite/ChangeLog:
PR modula2/120117
* gm2/pim/pass/highbit.mod: New test.
* gm2/pim/pass/highbit2.mod: New test.
(cherry picked from commit bb83283e5c5c55eab7493a58c5e415aa56f5940c)
Signed-off-by: Gaius Mulley <gaiusmod2@gmail.com>
|
|
starts with a directive
This bugfix is for FormatStrings to ensure that in the case of %x, %u the
procedure function PerformFormatString uses Copy rather than Slice to
avoid the case on an upper bound of zero in Slice. Oddly the %d case
had the correct code.
gcc/m2/ChangeLog:
PR modula2/119915
* gm2-libs/FormatStrings.mod (PerformFormatString): Handle
the %u and %x format specifiers in a similar way to the %d
specifier. Avoid using Slice and use Copy instead.
gcc/testsuite/ChangeLog:
PR modula2/119915
* gm2/pimlib/run/pass/format2.mod: New test.
(cherry picked from commit 7a4f7a92770db4e739e76a06175454ba38d60b22)
Signed-off-by: Gaius Mulley <gaiusmod2@gmail.com>
|
|
an unbounded array
This patch detects constants ZType, RType, CType being passed to unbounded
arrays and generates an error message highlighting the formal and
actual parameters in error.
gcc/m2/ChangeLog:
PR modula2/119914
* gm2-compiler/M2Check.mod (checkConstMeta): Add check for
Ztype, Rtype and Ctype and unbounded arrays.
(IsZRCType): New procedure function.
(isZRC): Add comment.
* gm2-compiler/M2Quads.mod:
* gm2-compiler/M2Range.mod (gdbinit): New procedure.
(BreakWhenRangeCreated): Ditto.
(CheckBreak): Ditto.
(InitRange): Call CheckBreak.
(Init): Add gdbhook and initialize interactive watch point.
* gm2-compiler/SymbolTable.def (GetNthParamAnyClosest): New
procedure function.
* gm2-compiler/SymbolTable.mod (BreakSym): Remove constant.
(BreakSym): Add Variable.
(stop): Remove.
(gdbhook): New procedure.
(BreakWhenSymCreated): Ditto.
(CheckBreak): Ditto.
(NewSym): Call CheckBreak.
(Init): Add gdbhook and initialize interactive watch point.
(MakeProcedure): Replace guarded call to stop with CheckBreak.
(GetNthParamChoice): New procedure function.
(GetNthParamOrdered): Ditto.
(GetNthParamAnyClosest): Ditto.
(GetOuterModuleScope): Ditto.
gcc/testsuite/ChangeLog:
PR modula2/119914
* gm2/pim/fail/constintarraybyte.mod: New test.
(cherry picked from commit e9a81addd5b7d018e173fa8d59aafc2f84e41d8b)
Signed-off-by: Gaius Mulley <gaiusmod2@gmail.com>
|
|
This is a cleanup patch which to use the existing plugin test
rather than check the configure build options.
gcc/testsuite/ChangeLog:
PR modula2/120188
* gm2.dg/doc/examples/plugin/fail/doc-examples-plugin-fail.exp:
Remove call to gm2-dg-frontend-configure-check and replace with
tests for whether plugin variables exist.
(cherry picked from commit f1f94e79dbcfa4b33267db27780870ea7810cd21)
Signed-off-by: Gaius Mulley <gaiusmod2@gmail.com>
|
