| Age | Commit message (Collapse) | Author | Files | Lines |
|---|
|
This is a merge of patch:
https://gcc.gnu.org/pipermail/gcc-patches/2021-May/571515.html
v2 patch already merged at 18dd4f283e15894a26c9a105c4f87d9a585f93c5,
this commit only consists of the v2-to-v3 diff part in above URL.
This v3 adds a small bug fix, where the initialization of the refcount didn't
handle all cases, fixed by using gomp_refcount_increment here (more consistent).
libgomp/ChangeLog:
* target.c (gomp_map_vars_internal): For new key entries, set
k->refcount to 0, remove initialization of k->structelem_refcount,
use gomp_increment_refcount to consistently handle all increment cases.
(manual cherry pick of e7073707bab79ceceaa0e2d25a632d03ac98d0fd)
|
|
This patch reworks indirect struct handling in gimplify.c (i.e. for
struct components mapped with "mystruct->a[0:n]", "mystruct->b", etc.),
for OpenACC. The key observation leading to these changes was that
component mappings of references-to-structures is already implemented
and working, and indirect struct component handling via a pointer can
work quite similarly. That lets us remove some earlier, special-case
handling for mapping indirect struct component accesses for OpenACC,
which required the pointed-to struct to be manually mapped before the
indirect component mapping.
With this patch, you can map struct components directly (e.g. an array
slice "mystruct->a[0:n]") just like you can map a non-indirect struct
component slice ("mystruct.a[0:n]"). Both references-to-pointers (with
the former syntax) and references to structs (with the latter syntax)
work now.
For Fortran class pointers, we no longer re-use GOMP_MAP_TO_PSET for the
class metadata (the structure that points to the class data and vptr)
-- it is instead treated as any other struct.
For C++, the struct handling also works for class members ("this->foo"),
without having to explicitly map "this[:1]" first.
For OpenACC, we permit chained indirect component references
("mystruct->a->b[0:n]"), though only the last part of such mappings will
trigger an attach/detach operation. To properly use such a construct
on the target, you must still manually map "mystruct->a[:1]" first --
but there's no need to map "mystruct[:1]" explicitly before that.
This version of the patch avoids altering code paths for OpenMP,
where possible.
2021-05-19 Julian Brown <julian@codesourcery.com>
gcc/fortran/
* trans-openmp.c (gfc_trans_omp_clauses): Don't create GOMP_MAP_TO_PSET
mappings for class metadata, nor GOMP_MAP_POINTER mappings for
POINTER_TYPE_P decls.
gcc/
* gimplify.c (extract_base_bit_offset): Add BASE_IND and OPENMP
parameters. Handle pointer-typed indirect references for OpenACC
alongside reference-typed ones.
(strip_components_and_deref, aggregate_base_p): New functions.
(build_struct_group): Add pointer type indirect ref handling,
including chained references, for OpenACC. Also handle references to
structs for OpenACC. Conditionalise bits for OpenMP only where
appropriate.
(gimplify_scan_omp_clauses): Rework pointer-type indirect structure
access handling to work more like the reference-typed handling for
OpenACC only.
* omp-low.c (scan_sharing_clauses): Handle pointer-type indirect struct
references, and references to pointers to structs also.
gcc/testsuite/
* g++.dg/goacc/member-array-acc.C: New test.
* g++.dg/gomp/member-array-omp.C: New test.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/deep-copy-15.c: New test.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c: New test.
* testsuite/libgomp.oacc-c++/deep-copy-17.C: New test.
|
|
This patch is a second attempt at refactoring struct component mapping
handling for OpenACC/OpenMP during gimplification, after the patch I
posted here:
https://gcc.gnu.org/pipermail/gcc-patches/2018-November/510503.html
And improved here, post-review:
https://gcc.gnu.org/pipermail/gcc-patches/2019-November/533394.html
This patch goes further, in that the struct-handling code is outlined
into its own function (to create the "GOMP_MAP_STRUCT" node and the
sorted list of nodes immediately following it, from a set of mappings
of components of a given struct or derived type). I've also gone through
the list-handling code and attempted to add comments documenting how it
works to the best of my understanding, and broken out a couple of helper
functions in order to (hopefully) have the code self-document better also.
2021-05-19 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.c (insert_struct_comp_map): Refactor function into...
(build_struct_comp_nodes): This new function. Remove list handling
and improve self-documentation.
(insert_node_after, move_node_after, move_nodes_after,
move_concat_nodes_after): New helper functions.
(build_struct_group): New function to build up GOMP_MAP_STRUCT node
groups to map struct components. Outlined from...
(gimplify_scan_omp_clauses): Here. Call above function.
|
|
For historical reasons, it seems that extract_base_bit_offset
unnecessarily used two different ways to strip ARRAY_REF/INDIRECT_REF
nodes from component accesses. I verified that the two ways of performing
the operation gave the same results across the whole testsuite (and
several additional benchmarks).
The code was like this since an earlier "mechanical" refactoring by me,
first posted here:
https://gcc.gnu.org/pipermail/gcc-patches/2018-November/510503.html
It was never clear to me if there was an important semantic
difference between the two ways of stripping the base before calling
get_inner_reference, but it appears that there is not, so one can go away.
2021-05-11 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.c (extract_base_bit_offset): Unify ARRAY_REF/INDIRECT_REF
stripping code in first call/subsequent call cases.
|
|
It never makes sense for a GOMP_MAP_ATTACH_DETACH mapping to survive
beyond gimplify.c, so this patch rewrites such mappings to GOMP_MAP_ATTACH
or GOMP_MAP_DETACH unconditionally (rather than checking for a list
of types of OpenACC or OpenMP constructs), in cases where it hasn't
otherwise been done already in the preceding code.
2021-05-19 Julian Brown <julian@codesourcery.com>
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Simplify condition
for changing GOMP_MAP_ATTACH_DETACH to GOMP_MAP_ATTACH or
GOMP_MAP_DETACH.
|
|
adjustments.
This patch largely implements three pieces of functionality:
(1) Per discussion and clarification on the omp-lang mailing list, standards
conforming behavior for mapping array sections should *NOT* also map the
base-pointer. This patch adjusts OpenMP map clause behavior to do this.
(2) Fixes in libgomp/target.c to not overwrite attached pointers when handling
device<->host copying.
(3) Changes to the C/C++ front-ends to extend the allowed component access syntax
in map clauses.
2021-05-11 Chung-Lin Tang <cltang@codesourcery.com>
gcc/c/ChangeLog:
* c-parser.c (struct omp_dim): New struct type for use inside
c_parser_omp_variable_list.
(c_parser_omp_variable_list): Allow multiple levels of array and
component accesses in array section base-pointer expression.
(c_parser_omp_clause_to): Set 'allow_deref' to true in call to
c_parser_omp_var_list_parens.
(c_parser_omp_clause_from): Likewise.
* c-typeck.c (handle_omp_array_sections_1): Extend allowed range
of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and
POINTER_PLUS_EXPR.
(c_finish_omp_clauses): Extend allowed ranged of expressions
involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR.
gcc/cp/ChangeLog:
* parser.c (struct omp_dim): New struct type for use inside
cp_parser_omp_var_list_no_open.
(cp_parser_omp_var_list_no_open): Allow multiple levels of array and
component accesses in array section base-pointer expression.
(cp_parser_omp_all_clauses): Set 'allow_deref' to true in call to
cp_parser_omp_var_list for to/from clauses.
* semantics.c (handle_omp_array_sections_1): Extend allowed range
of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and
POINTER_PLUS_EXPR.
(handle_omp_array_sections): Adjust pointer map generation of
references.
(finish_omp_clauses): Extend allowed ranged of expressions
involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR.
gcc/fortran/ChangeLog:
* trans-openmp.c (gfc_trans_omp_array_section): Do not generate
GOMP_MAP_ALWAYS_POINTER map for main array maps of ARRAY_TYPE type.
gcc/ChangeLog:
* gimplify.c (extract_base_bit_offset): Add 'tree *offsetp' parameter,
accomodate case where 'offset' return of get_inner_reference is
non-NULL.
(is_or_contains_p): Further robustify conditions.
(omp_target_reorder_clauses): In alloc/to/from sorting phase, also
move following GOMP_MAP_ALWAYS_POINTER maps along. Add new sorting
phase where we make sure pointers with an attach/detach map are ordered
correctly.
(gimplify_scan_omp_clauses): Add modifications to avoid creating
GOMP_MAP_STRUCT and associated alloc map for attach/detach maps.
gcc/testsuite/ChangeLog:
* c-c++-common/goacc/deep-copy-arrayofstruct.c: Adjust testcase.
* c-c++-common/gomp/target-enter-data-1.c: New testcase.
libgomp/ChangeLog:
* target.c (gomp_map_vars_existing): Make sure attached pointer is
not overwritten during cross-host/device copying.
(gomp_update): Likewise.
(gomp_exit_data): Likewise.
* testsuite/libgomp.c++/target-11.C: Adjust testcase.
* testsuite/libgomp.c++/target-12.C: Likewise.
* testsuite/libgomp.c++/target-15.C: Likewise.
* testsuite/libgomp.c++/target-16.C: Likewise.
* testsuite/libgomp.c++/target-17.C: Likewise.
* testsuite/libgomp.c++/target-21.C: Likewise.
* testsuite/libgomp.c++/target-23.C: Likewise.
* testsuite/libgomp.c/target-23.c: Likewise.
* testsuite/libgomp.c/target-29.c: Likewise.
|
|
Follow-up to recent og10 commit a70b5b1aa8b3d32f6728dbfcfc00b0cff8c5219d
"OpenMP 5.0: Implement relaxation of implicit map vs. existing device
mappings".
gcc/testsuite/
* c-c++-common/goacc/firstprivate-mappings-1.c: Fix up for C,
non-LP64.
|
|
This patch implements relaxing the requirements when a map with the implicit
attribute encounters an overlapping existing map. As the OpenMP 5.0 spec
describes on page 320, lines 18-27 (and 5.1 spec, page 352, lines 13-22):
"If a single contiguous part of the original storage of a list item with an
implicit data-mapping attribute has corresponding storage in the device data
environment prior to a task encountering the construct that is associated with
the map clause, only that part of the original storage will have corresponding
storage in the device data environment as a result of the map clause."
Also tracked in the OpenMP spec context as issue #1463:
https://github.com/OpenMP/spec/issues/1463
2021-05-05 Chung-Lin Tang <cltang@codesourcery.com>
include/ChangeLog:
* gomp-constants.h (GOMP_MAP_IMPLICIT): New special map kind bits value.
(GOMP_MAP_FLAG_SPECIAL_BITS): Define helper mask for whole set of
special map kind bits.
(GOMP_MAP_NONCONTIG_ARRAY_P): Adjust test for non-contiguous array map
kind bits to be more specific.
(GOMP_MAP_IMPLICIT_P): New predicate macro for implicit map kinds.
gcc/ChangeLog:
* tree.h (OMP_CLAUSE_MAP_IMPLICIT_P): New access macro for 'implicit'
bit, using 'base.deprecated_flag' field of tree_node.
* tree-pretty-print.c (dump_omp_clause): Add support for printing
implicit attribute in tree dumping.
* gimplify.c (gimplify_adjust_omp_clauses_1):
Set OMP_CLAUSE_MAP_IMPLICIT_P to 1 if map clause is implicitly created.
(gimplify_adjust_omp_clauses): Adjust place of adding implicitly created
clauses, from simple append, to starting of list, after non-map clauses.
* omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kind
values passed to libgomp for implicit maps.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/target-implicit-map-1.c: New test.
* c-c++-common/goacc/combined-reduction.c: Adjust scan test pattern.
* c-c++-common/goacc/firstprivate-mappings-1.c: Likewise.
* c-c++-common/goacc/mdc-1.c: Likewise.
* 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-8.c: Likewise.
* g++.dg/goacc/firstprivate-mappings-1.C: Likewise.
* g++.dg/gomp/target-lambda-1.C: Likewise.
* g++.dg/gomp/target-this-3.C: Likewise.
* g++.dg/gomp/target-this-4.C: Likewise.
* gfortran.dg/goacc/common-block-3.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.
libgomp/ChangeLog:
* target.c (gomp_map_vars_existing): Add 'bool implicit' parameter, add
implicit map handling to allow a "superset" existing map as valid case.
(get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return value.
(get_implicit): New function to extract implicit status.
(gomp_map_fields_existing): Adjust arguments in calls to
gomp_map_vars_existing, and add uses of get_implicit.
(gomp_map_vars_internal): Likewise.
* testsuite/libgomp.c-c++-common/target-implicit-map-1.c: New test.
|
|
We have seen an ICE both on trunk and devel/omp/gcc-10 branches which can
be reprodued with this simple testcase. It occurs if an OpenACC loop has
a collapse clause and any of the loop being collapsed uses GT or GE
condition. This issue is specific to OpenACC.
int main (void)
{
int ix, iy;
int dim_x = 16, dim_y = 16;
{
for (iy = dim_y - 1; iy > 0; --iy)
for (ix = dim_x - 1; ix > 0; --ix)
;
}
}
The problem is caused by a failing assertion in expand_oacc_collapse_init.
It checks that cond_code for fd->loop should be same as cond_code for all
the loops that are being collapsed. As the cond_code for fd->loop is
LT_EXPR with collapse clause (set at the end of omp_extract_for_data),
this assertion forces that all the loop in collapse clause should use
< operator.
There does not seem to be anything in the code which demands this
condition as loop with > condition works ok otherwise. I digged old
mailing list a bit but could not find any discussion on this change.
Looking at the code, expand_oacc_for checks that fd->loop->cond_code is
either LT_EXPR or GT_EXPR. I guess the original intention was to have
similar checks on the loop which are being collapsed. But the way check
was written does not acheive that.
I have fixed it by modifying the check in the assertion to be same as
check on fd->loop->cond_code.
I tested goacc and libgomp (with nvptx offloading) and did not see any
regression. I have added new tests to check collapse with GT/GE condition.
PR middle-end/98088
gcc/
* omp-expand.c (expand_oacc_collapse_init): Update condition in
a gcc_assert.
gcc/testsuite/
* c-c++-common/goacc/collapse-2.c: New.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Add check
for loop with GT/GE condition.
* testsuite/libgomp.oacc-c-c++-common/collapse-3.c: Likewise.
|
|
This is a fix-up for og10 commit c89b23b73edeeb7e3d8cbad278e505c2d6d770c4
"'libgomp.oacc-fortran/derivedtypes-arrays-1.f90' OpenACC 'serial' construct
diagnostic for nvptx offloading".
We're missing in og10 a few patches related to diagnostics location
tracking/checking, both compiler-side and testsuite-side.
libgomp/
* testsuite/libgomp.oacc-fortran/derivedtypes-arrays-1.f90: Adjust
for og10.
|
|
gcc/
* config/gcn/gcn-opts.h (enum processor_type): Add PROCESSOR_GFX908.
* config/gcn/gcn.c (gcn_omp_device_kind_arch_isa): Add gfx908.
(output_file_start): Add gfx908.
* config/gcn/gcn.opt (gpu_type): Add gfx908.
* config/gcn/t-gcn-hsa (MULTILIB_OPTIONS): Add march=gfx908.
(MULTILIB_DIRNAMES): Add gfx908.
* config/gcn/mkoffload.c (EF_AMDGPU_MACH_AMDGCN_GFX908): New define.
(main): Recognize gfx908.
* config/gcn/t-omp-device: Add gfx908.
libgomp/
* plugin/plugin-gcn.c (EF_AMDGPU_MACH): Add
EF_AMDGPU_MACH_AMDGCN_GFX908.
(gcn_gfx908_s): New constant string.
(isa_hsa_name): Add gfx908.
(isa_code): Add gfx908.
(cherry picked from commit 3535402e20118655b2ad4085a6e1d4f1b9c46e92)
|
|
'libhsa-runtime64.so'
For unknown reasons, this had gotten added for the libgomp HSA plugin in commit
b8d89b03db5f212919e4571671ebb4f5f8b1e19d (r242749) "Remove build dependence on
HSA run-time", and later propagated into the GCN plugin.
libgomp/
* plugin/plugin-hsa.c (init_enviroment_variables): Don't prepend
the 'HSA_RUNTIME_LIB' path to 'libhsa-runtime64.so'.
* plugin/plugin-gcn.c (init_environment_variables): Likewise.
* plugin/configfrag.ac (HSA_RUNTIME_LIB): Clean up.
* config.h.in: Regenerate.
* configure: Likewise.
(cherry picked from commit 7c1e856bedb4ae190c420ec2d2ca5e08730cf21d)
|
|
diagnostic for nvptx offloading
Fixup for recent commit d28f3da11d8c0aed9b746689d723022a9b5ec04c "openacc: Fix
lowering for derived-type mappings through array elements". With nvptx
offloading we see the usual:
[...]/libgomp.oacc-fortran/derivedtypes-arrays-1.f90: In function 'MAIN__._omp_fn.0':
[...]/libgomp.oacc-fortran/derivedtypes-arrays-1.f90:90:40: warning: using vector_length (32), ignoring 1
libgomp/
* testsuite/libgomp.oacc-fortran/derivedtypes-arrays-1.f90:
OpenACC 'serial' construct diagnostic for nvptx offloading.
(cherry picked from commit 8bafce1be11a301c2421483736c634b8bf330e69)
|
|
Fix-up for og10 commit c2e4a17adc0989f216c7fc3f93f150c66adba23a "OpenMP 5.0:
requires directive".
libgomp: while loading libgomp-plugin-hsa.so.1: [...]/libgomp-plugin-hsa.so.1: undefined symbol: GOMP_OFFLOAD_supported_features
libgomp/
* plugin/plugin-hsa.c (GOMP_OFFLOAD_supported_features): New
function.
|
|
plugin build
Fix-up for og10 commit c2e4a17adc0989f216c7fc3f93f150c66adba23a "OpenMP 5.0:
requires directive".
The GCC offloading target configurations don't build/use
'crtoffloadbegin.o'/'crtoffloadtable.o'/'crtoffloadend.o'
('libgcc/offloadstuff.c'), but the libgomp IntelMIC plugin still does link
against libgomp, and the latter unconditionally refers to
'__requires_mask_table', '__requires_mask_table_end':
make[3]: Entering directory '[...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/x86_64-intelmicemul-linux-gnu/liboffloadmic/plugin'
[...]/build-gcc-offload-x86_64-intelmicemul-linux-gnu/./gcc/xg++ [...] -loffloadmic_target -lcoi_device -lgomp -rdynamic ../ofldbegin.o offload_target_main.o ../ofldend.o -o offload_target_main
./../../libgomp/.libs/libgomp.so: undefined reference to `__requires_mask_table_end'
./../../libgomp/.libs/libgomp.so: undefined reference to `__requires_mask_table'
collect2: error: ld returned 1 exit status
Makefile:806: recipe for target 'offload_target_main' failed
make[3]: *** [offload_target_main] Error 1
I have not researched what a proper fix would look like.
libgomp/
* target.c (__requires_mask_table, __requires_mask_table_end): Add
'__attribute__((weak))'.
|
|
Testcase added in og10 commit 835160b024ce34bc6f258e7df366bb5c15e12a4b "OpenMP:
Handle order(concurrent) clause in gfortran" (cherry picked from commit
d8140b9ed3c0fed041aedaff3fa4a603984ca10f), but og10 doesn't have commit
4f2ab6b89e170f1343f935761481c3745fe603b1 "[OpenMP, gimplifier] 'inform' after
'error' diagnostic".
gcc/testsuite/
* gfortran.dg/gomp/order-4.f90: Adjust.
|
|
This patch implements proper lambda capturing of pointer and reference variables
as specified in OpenMP 5.0. We map the entire closure object as a to-map,
attach pointers to zero-length array sections, and perform mapping of
references.
2021-03-18 Chung-Lin Tang <cltang@codesourcery.com>
gcc/cp/ChangeLog:
* cp-tree.h (set_omp_target_this_expr): Delete.
(finish_omp_target_clauses): New prototype.
* lambda.c (lambda_expr_this_capture): Remove call to
set_omp_target_this_expr.
* parser.c (cp_parser_omp_target): Likewise.
* pt.c (tsubst_expr): Add call to finish_omp_target_clauses for target
directives.
* semantics.c (omp_target_this_expr): Delete.
(omp_target_ptr_members_accessed): Delete.
(finish_non_static_data_member): Remove call to
set_omp_target_this_expr. Remove use of omp_target_ptr_members_accessed.
(finish_this_expr): Remove call to set_omp_target_this_expr.
(struct omp_target_walk_data): New struct for walking over
target-directive tree body.
(finish_omp_target_clauses_r): New function for tree walk.
(finish_omp_target_clauses): New function, with code factored out from
finish_omp_target. Add lambda object handling case.
(finish_omp_target): Factor code out and adjust to use
finish_omp_target_clauses.
(finish_omp_clauses): Revert prior "Adjustments to allow '*ptr' and
'ptr->member' cases in map clausess.", since not needed with new
organization of target-directive clause processing.
gcc/testsuite/ChangeLog:
* g++.dg/gomp/target-lambda-1.C: New test.
libgomp/testsuite/ChangeLog:
* libgomp.c++/target-lambda-1.C: New test.
|
|
Prior patches for C++ non-static member access had problems under template
classes, due to re-calling of finish_omp_clauses after finish_omp_target
created the implicit maps required, but not of allowed form in finish_omp_clauses.
This patch solves this by slightly relaxing the allowed expressions in
finish_omp_clauses.
2021-03-11 Chung-Lin Tang <cltang@codesourcery.com>
gcc/cp/ChangeLog:
* semantics.c (finish_omp_clauses): Adjustments to allow '*ptr' and
'ptr->member' cases in map clausess.
(finish_omp_target): Use INDIRECT_REF instead of MEM_REF in created
clauses, add processing_template_decl handling.
gcc/ChangeLog:
* gimplify.c (gimplify_scan_omp_clauses): Under !DECL_P case of
GOMP_CLAUSE_MAP handling, add STRIP_NOPS for indir_p case, add to
struct_deref_set for map(*ptr_to_struct) cases.
gcc/testsuite/ChangeLog:
* g++.dg/gomp/target-this-3.C: Adjust scan test.
* g++.dg/gomp/target-this-4.C: Likewise.
* g++.dg/gomp/target-this-5.C: New test.
libgomp/ChangeLog:
* testsuite/libgomp.c++/target-this-5.C: New test.
|
|
This patch merges some of the equivalent changes already done for the C++
front-end to the C parts.
2021-03-08 Chung-Lin Tang <cltang@codesourcery.com>
gcc/c/ChangeLog:
* c-parser.c (c_parser_omp_clause_map): Set 'allow_deref' argument in
call to c_parser_omp_variable_list to 'true'.
* c-typeck.c (handle_omp_array_sections_1): Add strip of MEM_REF in
array base handling.
(c_finish_omp_clauses): Handle 'A->member' case in map clauses.
gcc/ChangeLog:
* gimplify.c (gimplify_scan_omp_clauses): Add MEM_REF case when
handling component_ref_p case. Add unshare_expr and gimplification
when created GOMP_MAP_STRUCT is not a DECL. Add code to add
firstprivate pointer for *pointer-to-struct case.
gcc/testsuite/ChangeLog:
* gcc.dg/gomp/target-3.c: New test.
|
|
The relocation symbols were inadvertantly wiped when the type was set in
mkoffload.
gcc/ChangeLog
* config/gcn/mkoffload.c (copy_early_debug_info): Don't wipe
relocation symbols.
(cherry picked from commit d24a4c8c4256cc287ebf4ad80368b4f1edb1733e)
|
|
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.
This patch replaces 808bdf1bb29 and fdcb23540a2. It should be squashed with
those before being posted upstream.
gcc/
* gcc/dwarf2out.c (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 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).
2021-03-01 Kwok Cheung Yeung <kcy@codesourcery.com>
gcc/
* omp-expand.c (expand_oacc_for): Convert .tile variable to
diff_type before multiplying.
* omp-general.c (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/
* testsuite/libgomp.c-c++-common/collapse-4.c: New.
* testsuite/libgomp.fortran/collapse5.f90: New.
|
|
This patch is a merge of:
https://gcc.gnu.org/legacy-ml/gcc-patches/2020-01/msg01246.html
Static members in general disqualify a C++ class from being target mappable,
but static constexprs are inline optimized away, so should not interfere.
OpenMP 5.0 in general lifts the static member limitation, so this
patch will probably further adjusted later.
2021-03-03 Chung-Lin Tang <cltang@codesourcery.com>
gcc/cp/ChangeLog:
* decl2.c (cp_omp_mappable_type_1): Allow fields with
DECL_DECLARED_CONSTEXPR_P to be mapped.
gcc/testsuite/ChangeLog:
* g++.dg/goacc/static-constexpr-1.C: New test.
* g++.dg/gomp/static-constexpr-1.C: New test.
|
|
The stacks are not supposed to be deleted every time memory is allocated, only
when there is insufficient memory. The unconditional call here seems to be in
error, and is causing a costly reallocation of the stacks before every launch.
libgomp/
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_alloc): Remove early call to
nvptx_stacks_free.
|
|
gcc/fortran/ChangeLog:
PR fortran/97242
* expr.c (gfc_is_not_contiguous): Fix check.
(gfc_check_pointer_assign): Use it.
gcc/testsuite/ChangeLog:
PR fortran/97242
* gfortran.dg/contiguous_11.f90: New test.
* gfortran.dg/contiguous_4.f90: Update.
* gfortran.dg/contiguous_7.f90: Update.
(cherry picked from commit 65167982efa4dbb96698d026e6d7e17acb513f0a)
|
|
This should be squashed with 808bdf1bb29 and fdcb23540a2 to go to mainline.
gcc/
* dwarf2out.c (gen_subprogram_die): Replace existing low/high PC
attributes, rather than ICE.
|
|
(INDIRECT_REF ...)) map clauses".
This patch was reverted in commit ea43db9372133e84f95bdb2c6934a5cc3db3d530
due to a regression. Now re-applying to devel/omp/gcc-10 after regression fixed.
This patch tries to allow map(A->ptr) to be properly handled the same way as
map(B.ptr) expressions. map(struct:*A) clauses are now produced during
gimplify.
This patch, as of time of commit, is only pushed to devel/omp/gcc-10, not yet
submitted as mainline patch to upstream.
2021-02-26 Chung-Lin Tang <cltang@codesourcery.com>
gcc/ChangeLog:
* gimplify.c ("tree-hash-traits.h"): Add include.
(gimplify_scan_omp_clauses): Change struct_map_to_clause to type
hash_map<tree_operand, tree> *. Adjust struct map handling to handle
cases of *A and A->B expressions.
(gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for
exit data directives code to earlier position.
gcc/testsuite/ChangeLog:
* g++.dg/gomp/target-3.C: Adjust testcase gimple scanning.
* g++.dg/gomp/target-this-2.C: Likewise.
* g++.dg/gomp/target-this-3.C: Likewise.
* g++.dg/gomp/target-this-4.C: Likewise.
libgomp/ChangeLog:
* testsuite/libgomp.c++/target-23.C: New testcase.
|
|
Fixed a regression of array members not working in OpenMP map clauses after
commit bf8605f14ec33ea31233a3567f3184fee667b695.
This patch itself probably should be considered a fix for commit aadfc9843.
2021-02-26 Chung-Lin Tang <cltang@codesourcery.com>
gcc/cp/ChangeLog:
* semantics.c (handle_omp_array_sections): Adjust position of making
COMPONENT_REF from FIELD_DECL to earlier position.
|
|
This patch disallows selecting components of array sections in update
directives for OpenACC, as specified in OpenACC 3.0, "2.14.4. Update
Directive":
In Fortran, members of variables of derived type may appear, including
a subarray of a member. Members of subarrays of derived type may
not appear.
The diagnostic for attempting to use the same construct on other
directives has also been improved.
gcc/fortran/
* openmp.c (resolve_omp_clauses): Disallow selecting components
of arrays of derived type.
gcc/testsuite/
* gfortran.dg/goacc/array-with-dt-2.f90: Remove expected errors.
* gfortran.dg/goacc/array-with-dt-6.f90: New test.
* gfortran.dg/goacc/mapping-tests-2.f90: Update expected error.
* gfortran.dg/goacc/ref_inquiry.f90: Update expected errors.
* gfortran.dg/gomp/ref_inquiry.f90: Likewise.
libgomp/
* testsuite/libgomp.oacc-fortran/array-stride-dt-1.f90: Remove
expected errors.
(cherry picked from commit 366cf1127a547ff77024a551abb01bb1a6e963cd)
|
|
This patch fixes lowering of derived-type mappings which select elements
of arrays of derived types, and similar. These would previously lead
to ICEs.
With this change, OpenACC directives can pass through constructs that
are no longer recognized by the gimplifier, hence alterations are needed
there also.
gcc/fortran/
* trans-openmp.c (gfc_trans_omp_clauses): Handle element selection
for arrays of derived types.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Handle ATTACH_DETACH
for non-decls.
gcc/testsuite/
* gfortran.dg/goacc/array-with-dt-1.f90: New test.
* gfortran.dg/goacc/array-with-dt-3.f90: Likewise.
* gfortran.dg/goacc/array-with-dt-4.f90: Likewise.
* gfortran.dg/goacc/array-with-dt-5.f90: Likewise.
* gfortran.dg/goacc/derived-chartypes-1.f90: Re-enable test.
* gfortran.dg/goacc/derived-chartypes-2.f90: Likewise.
* gfortran.dg/goacc/derived-classtypes-1.f95: Uncomment
previously-broken directives.
libgomp/
* testsuite/libgomp.oacc-fortran/derivedtypes-arrays-1.f90: New test.
* testsuite/libgomp.oacc-fortran/update-dt-array.f90: Likewise.
(cherry picked from commit d28f3da11d8c0aed9b746689d723022a9b5ec04c)
|
|
gcc/fortran/ChangeLog:
* expr.c (gfc_is_simplify_contiguous): Handle REF_INQUIRY, i.e.
%im and %re which are EXPR_VARIABLE.
* openmp.c (resolve_omp_clauses): Diagnose %re/%im explicitly.
gcc/testsuite/ChangeLog:
* gfortran.dg/goacc/ref_inquiry.f90: New test.
* gfortran.dg/gomp/ref_inquiry.f90: New test.
(cherry picked from commit 799478b8914c438f7a33eb319efbae69c81f2111)
|
|
gcc/fortran/ChangeLog:
* openmp.c (resolve_omp_clauses): Explicitly diagnose
substrings as not permitted.
gcc/testsuite/ChangeLog:
* gfortran.dg/goacc/substring.f90: New test.
* gfortran.dg/gomp/substring.f90: New test.
(cherry picked from commit f0e618faeb619ec02dabbef203a5575fca44a7f7)
|
|
This patch adds some tests for character types that are components
of derived types used in OpenACC data-movement clauses (some of which
currently fail and are thus XFAILed), and tests (also XFAILed) mixing
arrays and derived types. The XFAILs are addressed by follow-on patches.
Originally a combination of several mainline patches.
(cherry picked from commit b2d84e9f9cccbe4ee662f7002b83105629d09939)
(cherry picked from commit 9a4d32f85ccebc0ee4b24e6d9d7a4f11c04d7146)
(cherry picked from commit b0fb2720d88d680af18981a2097397196b505a1f)
(cherry picked from commit f7fb2f662fe12f327ece8b034ab76b36fdca4696)
gcc/testsuite/
* gfortran.dg/goacc/array-with-dt-2.f90: New test.
* gfortran.dg/goacc/derived-chartypes-1.f90: Likewise.
* gfortran.dg/goacc/derived-chartypes-2.f90: Likewise.
* gfortran.dg/goacc/derived-chartypes-3.f90: Likewise.
* gfortran.dg/goacc/derived-chartypes-4.f90: Likewise.
libgomp/
* testsuite/libgomp.oacc-fortran/array-stride-dt-1.f90: New test.
|
|
Elsewhere in the Fortran front-end, the class_pointer attribute is
used for BT_CLASS entities instead of the pointer attribute. This patch
follows suit for OpenACC. I couldn't actually come up with a test case
where this makes a difference (i.e., where "class_pointer" and "pointer"
have different values at this point in the code), but this may nonetheless
fix a latent bug.
gcc/fortran/
* trans-openmp.c (gfc_trans_omp_clauses): Use class_pointer attribute
for BT_CLASS.
(cherry picked from commit f743fe231663e32d52db987650d0ec3381a777af)
|
|
The stanza in gfc_trans_omp_clauses that handles derived type members
that are themselves derived type pointers or class pointers now adds
an explicit dereference only for the latter. The former is already
dereferenced transparently in gfc_conv_component_ref.
gcc/fortran/
* trans-openmp.c (gfc_trans_omp_clauses): Fix dereferencing for
BT_DERIVED members.
gcc/testsuite/
* gfortran.dg/goacc/derived-classtypes-1.f95: New test.
(cherry picked from commit cff6e8db880b6e262730b1ce9a9cb00c1f5571e2)
|
|
Followup to commit 383400a6078d75bbfa1216c9af2c37f7e88740c9
gcc/ChangeLog
* doc/invoke.texi (nvptx's -misa): Update default to sm_35.
(cherry picked from commit 91e4e16b550540723cca824b9674c7d8c43f4849)
|
|
The nvptx-as assembler verifies the ptx code using ptxas, if there's any
in the PATH.
The default in the nvptx port for -misa=sm_xx is sm_30, but the ptxas of the
latest cuda release (11.1) no longer supports sm_30.
Consequently we cannot build gcc against that release (although we should
still be able to build without any cuda release).
Fix this by setting -misa=sm_35 by default.
Tested check-gcc on nvptx.
Tested libgomp on x86_64-linux with nvpx accelerator.
Both build again cuda 9.1.
gcc/ChangeLog:
2020-10-09 Tom de Vries <tdevries@suse.de>
PR target/97348
* config/nvptx/nvptx.h (ASM_SPEC): Also pass -m to nvptx-as if
default is used.
* config/nvptx/nvptx.opt (misa): Init with PTX_ISA_SM35.
(cherry picked from commit 383400a6078d75bbfa1216c9af2c37f7e88740c9)
|
|
(INDIRECT_REF ...)) map clauses."
This reverts commit bf8605f14ec33ea31233a3567f3184fee667b695.
Regressions were reported for BZ 9574 - C++ class member mapping.
|
|
gcc/fortran/ChangeLog:
PR fortran/99043
* trans-expr.c (gfc_conv_procedure_call): Don't reset
rank of assumed-rank array.
gcc/testsuite/ChangeLog:
PR fortran/99043
* gfortran.dg/assumed_rank_20.f90: New test.
(cherry picked from commit f699e0b16578cdc1be8b90691ef8b0964af32d2f)
|
|
Backport from mainline; also fixes PR fortran/99045
2020-12-27 Paul Thomas <pault@gcc.gnu.org>
gcc/fortran
PR fortran/97694
PR fortran/97723
* check.c (allocatable_check): Select rank temporaries are
permitted even though they are treated as associate variables.
* resolve.c (gfc_resolve_code): Break on select rank as well as
select type so that the block os resolved.
* trans-stmt.c (trans_associate_var): Class associate variables
that are optional dummies must use the backend_decl.
gcc/testsuite/
PR fortran/97694
PR fortran/97723
* gfortran.dg/select_rank_5.f90: New test.
(cherry picked from commit c4a678981572c12d158709ace0d3f23dd04cf217)
|
|
The loops in the testcase are non-rectangular and have 0 iterations
(the outer loop iterates, but the inner one never). In this case we
just have the overall number of iterations computed (0), and don't have
factor and other values computed. We never need to map logical iterations
to the individual iterations in that case, and we were crashing during
expansion of that code.
2020-11-18 Jakub Jelinek <jakub@redhat.com>
PR middle-end/97862
* omp-expand.c (expand_omp_for_init_vars): Don't use the sqrt path
if number of iterations is constant 0.
* c-c++-common/gomp/pr97862.c: New test.
(cherry picked from commit ba009860aec4619f2424f5bdee812f14572dc3cc)
|
|
This propagates needed values from the point where number of iterations
is calculated on composite loops to the places where that information
is needed to use the more efficient square root discovery to compute
the starting iterator values from the logical iteration number.
2020-10-13 Jakub Jelinek <jakub@redhat.com>
* omp-low.c (add_taskreg_looptemp_clauses): For triangular loops
with non-constant number of iterations add another 4 _looptemp_
clauses before the (optional) one for lastprivate.
(lower_omp_for_lastprivate): Skip those clauses when looking for
the lastprivate clause.
(lower_omp_for): For triangular loops with non-constant number of
iterations add another 4 _looptemp_ clauses.
* omp-expand.c (expand_omp_for_init_counts): For triangular loops
with non-constant number of iterations set counts[0],
fd->first_inner_iterations, fd->factor and fd->adjn1 from the newly
added _looptemp_ clauses.
(expand_omp_for_init_vars): Initialize the newly added _looptemp_
clauses.
(find_lastprivate_looptemp): New function.
(expand_omp_for_static_nochunk, expand_omp_for_static_chunk,
expand_omp_taskloop_for_outer): Use it instead of manually skipping
_looptemp_ clauses.
(cherry picked from commit 14707c896a207606f13886d3b3251e8db1f3c9c0)
|
|
2020-08-13 Jakub Jelinek <jakub@redhat.com>
* gimplify.c (gimplify_omp_taskloop_expr): New function.
(gimplify_omp_for): Use it. For OMP_FOR_NON_RECTANGULAR
loops adjust in outer taskloop the var-outer decls.
* omp-expand.c (expand_omp_taskloop_for_inner): Handle non-rectangular
loops.
(expand_omp_for): Don't reject non-rectangular taskloop.
* omp-general.c (omp_extract_for_data): Don't assert that
non-rectangular loops have static schedule, instead treat loop->m1
or loop->m2 as if loop->n1 or loop->n2 is non-constant.
* testsuite/libgomp.c/loop-22.c (main): Add some further tests.
* testsuite/libgomp.c/loop-23.c (main): Likewise.
* testsuite/libgomp.c/loop-24.c: New test.
(cherry picked from commit 2e47c8c6eac405ceb599bf5e31ac3717c22a008c)
|
|
The number of loops computation and logical iteration -> actual iterator values
computations can now be done separately even on composite constructs (though
for triangular loops it would still be more efficient to propagate a few values
through, will handle that incrementally).
simd and taskloop are still unhandled.
2020-08-05 Jakub Jelinek <jakub@redhat.com>
* omp-expand.c (expand_omp_for): Don't disallow combined non-rectangular
loops.
* testsuite/libgomp.c/loop-22.c: New test.
* testsuite/libgomp.c/loop-23.c: New test.
(cherry picked from commit 9f3abfb84e2a7ca115b0550c32308b5ada3e6a46)
|
|
iterations is computed at runtime
For triangular loops use more efficient logical iteration number
to actual iterator values computation even for non-rectangular loops
where number of loop iterations could not be computed at compile time.
2020-08-05 Jakub Jelinek <jakub@redhat.com>
* omp-expand.c (expand_omp_for_init_counts): Remember
first_inner_iterations, factor and n1o from the number of iterations
computation in *fd.
(expand_omp_for_init_vars): Use more efficient logical iteration number
to actual iterator values computation even for non-rectangular loops
where number of loop iterations could not be computed at compile time.
(cherry picked from commit 325714b4968050c927fe72f4a4ba860da2bf4ee2)
|
|
some non-rectangular loops
2020-08-04 Jakub Jelinek <jakub@redhat.com>
* omp-expand.c (expand_omp_for_init_counts): For triangular loops
compute number of iterations at runtime more efficiently.
(expand_omp_for_init_vars): Adjust immediate dominators.
(extract_omp_for_update_vars): Likewise.
(cherry picked from commit 29e0ad452cdf001bccd12c1cb5f4797df1114a3a)
|
|
I've missed
+FAIL: libgomp.c/loop-21.c execution test
during testing of the recent patch. The problem is that while
for the number of iterations computation it doesn't matter if we compute
min_inner_iterations as (m2 * first + n2 + (adjusted step) + m1 * first + n1) / step
or (m2 * last + n2 + (adjusted step) + m1 * last + n1) / step provided that
in the second case we use as factor (m1 - m2) * ostep / step rather than
(m2 - m1) * ostep / step, for the logical to actual iterator values computation
it does matter and in my hand written C implementations of all the cases (outer
vs. inner loop with increasing vs. decreasing iterator) I'm using the same computation
and it worked well for all the pseudo-random iterators testing it was doing.
It also means min_inner_iterations is misnamed, because it is not really
minimum number of inner iterations, whether the first or last outer iteration
results in the smaller or larger value of this can be (sometimes) only
determined at runtime.
So this patch also renames it to first_inner_iterations.
2020-07-15 Jakub Jelinek <jakub@redhat.com>
PR libgomp/96198
* omp-general.h (struct omp_for_data): Rename min_inner_iterations
member to first_inner_iterations, adjust comment.
* omp-general.c (omp_extract_for_data): Adjust for the above change.
Always use n1first and n2first to compute it, rather than depending
on single_nonrect_cond_code. Similarly, always compute factor
as (m2 - m1) * outer_step / inner_step rather than sometimes m1 - m2
depending on single_nonrect_cond_code.
* omp-expand.c (expand_omp_for_init_vars): Rename min_inner_iterations
to first_inner_iterations and min_inner_iterationsd to
first_inner_iterationsd.
(cherry picked from commit 79c12969ec3e9185fdbb90d3b1699d64b1cd0901)
|
|
In loops like:
#pragma omp parallel for collapse(2)
for (i = -4; i < 8; i++)
for (j = 3 * i; j > 2 * i; j--)
for some outer loop iterations there are no inner loop iterations at all,
the condition is false. In order to use Summæ Potestate to count number
of iterations or to transform the logical iteration number to actual
iterator values using quadratic non-equation root discovery the outer
iterator range needs to be adjusted, such that the inner loop has at least
one iteration for each of the outer loop iterator value in the reduced
range. Sometimes this adjustment is done at the start of the range,
at other times at the end.
This patch implements it during the compile time number of loop computation
(if all expressions are compile time constants).
2020-07-14 Jakub Jelinek <jakub@redhat.com>
* omp-general.h (struct omp_for_data): Add adjn1 member.
* omp-general.c (omp_extract_for_data): For non-rect loop, punt on
count computing if n1, n2 or step are not INTEGER_CST earlier.
Narrow the outer iterator range if needed so that non-rect loop
has at least one iteration for each outer range iteration. Compute
adjn1.
* omp-expand.c (expand_omp_for_init_vars): Use adjn1 if non-NULL
instead of the outer loop's n1.
* testsuite/libgomp.c/loop-21.c: New test.
(cherry picked from commit f418bd4b92a03ee0ec0fe4cfcd896e86e11ac2cf)
|
|
computation using search for quadratic equation root(s)
This patch implements the optimized logical to actual iterators
computation for triangular loops.
I have a rough implementation using integers, but this one uses floating
point. There is a small problem that -fopenmp programs aren't linked with
-lm, so it does it only if the hw has sqrt optab (and uses ifn rather than
__builtin_sqrt because it obviously doesn't need errno handling etc.).
Do you think it is ok this way, or should I use the integral computation
using inlined isqrt (we have inequation of the form
start >= x * t10 + t11 * (((x - 1) * x) / 2)
where t10 and t11 are signed long long values and start unsigned long long,
and the division by 2 actually is a problem for accuracy in some cases, so
if we do it in integral, we need to do actually
long long t12 = 2 * t10 - t11;
unsigned long long t13 = t12 * t12 + start * 8 * t11;
unsigned long long isqrt_ = isqrtull (t13);
long long x = (((long long) isqrt_ - t12) / t11) >> 1;
with careful overflow checking on all the computations before isqrtull
(and on overflows use the fallback implementation).
2020-07-09 Jakub Jelinek <jakub@redhat.com>
* omp-general.h (struct omp_for_data): Add min_inner_iterations
and factor members.
* omp-general.c (omp_extract_for_data): Initialize them and remember
them in OMP_CLAUSE_COLLAPSE_COUNT if needed and restore from there.
* omp-expand.c (expand_omp_for_init_counts): Fix up computation of
counts[fd->last_nonrect] if fd->loop.n2 is INTEGER_CST.
(expand_omp_for_init_vars): For
fd->first_nonrect + 1 == fd->last_nonrect loops with for now
INTEGER_CST fd->loop.n2 find quadratic equation roots instead of
using fallback method when possible.
* testsuite/libgomp.c/loop-19.c: New test.
* testsuite/libgomp.c/loop-20.c: New test.
(cherry picked from commit 5acef69f9d3d9f3c537b5e5157519edf02f86c4d)
|
|
THe OpenMP 5 standard requires that if some loop in OpenMP loop nest refers
to some outer loop's iterator variable, then the subtraction of the multiplication
factors for the outer iterator multiplied by the outer increment modulo the
inner increment is 0. For loops with non-constants in any of these we can't
diagnose it, it would be a task for something like -fsanitize=openmp,
but if all these are constant, we can diagnose it.
2020-07-02 Jakub Jelinek <jakub@redhat.com>
* omp-expand.c (expand_omp_for): Diagnose non-rectangular loops with
invalid steps - ((m2 - m1) * incr_outer) % incr must be 0 in valid
OpenMP non-rectangular loops. Use XALLOCAVEC.
* c-c++-common/gomp/loop-7.c: New test.
(cherry picked from commit 9d50112acfc01f85fe0fb6d88b329e6122e817b3)
|
