summaryrefslogtreecommitdiff
path: root/libgomp
Commit message (Collapse)AuthorAgeFilesLines
* Structure element mapping for OpenMP 5.0 v3devel/omp/gcc-10Chung-Lin Tang2021-05-311-2/+6
| | | | | | | | | | | | | | | | | | | 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)
* [og10] Rework indirect struct handling for OpenACC in gimplify.cJulian Brown2021-05-193-0/+400
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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.
* Remove array section base-pointer mapping semantics, and other front-end ↵Chung-Lin Tang2021-05-1111-73/+190
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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.
* OpenMP 5.0: Implement relaxation of implicit map vs. existing device mappingsChung-Lin Tang2021-05-052-19/+90
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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.
* [OpenACC] Fix an ICE where a loop with GT condition is collapsed.Hafiz Abid Qadeer2021-04-113-4/+38
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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.
* Adjust 'libgomp.oacc-fortran/derivedtypes-arrays-1.f90' for og10Thomas Schwinge2021-03-262-1/+6
| | | | | | | | | | | | | 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.
* amdgcn: Add gfx908 supportAndrew Stubbs2021-03-252-0/+18
| | | | | | | | | | | | | | | | | | | | | | | | 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)
* libgomp HSA/GCN plugins: don't prepend the 'HSA_RUNTIME_LIB' path to ↵Thomas Schwinge2021-03-256-22/+9
| | | | | | | | | | | | | | | | | | '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)
* 'libgomp.oacc-fortran/derivedtypes-arrays-1.f90' OpenACC 'serial' construct ↵Thomas Schwinge2021-03-252-0/+4
| | | | | | | | | | | | | | | | | 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)
* OpenMP 5.0: requires directive: adjust libgomp HSA pluginThomas Schwinge2021-03-252-0/+11
| | | | | | | | | | | 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.
* [WIP] OpenMP 5.0: requires directive: workaround to fix libgomp IntelMIC ↵Thomas Schwinge2021-03-252-0/+7
| | | | | | | | | | | | | | | | | | | | | | | | | | | 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))'.
* Lambda capturing of pointers and references in target directivesChung-Lin Tang2021-03-191-0/+86
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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.
* Fix template case of non-static member access inside member functionsChung-Lin Tang2021-03-111-0/+30
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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.
* openmp: Scale type precision of collapsed iterator variableKwok Cheung Yeung2021-03-053-0/+51
| | | | | | | | | | | | | | | | | | | 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.
* nvptx: remove erroneous stack deletionAndrew Stubbs2021-03-021-2/+0
| | | | | | | | | | | 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.
* Recommit "Enable gimplify GOMP_MAP_STRUCT handling of (COMPONENT_REF ↵Chung-Lin Tang2021-02-261-0/+34
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | (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.
* [og10] openacc: Strided array sections and components of derived-type arraysJulian Brown2021-02-242-3/+9
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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)
* [og10] openacc: Fix lowering for derived-type mappings through array elementsJulian Brown2021-02-243-0/+169
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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)
* [og10] openacc: Character types and mixed arrays/derived type testsJulian Brown2021-02-242-0/+51
| | | | | | | | | | | | | | | | | | | | | | | 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.
* Revert "Enable gimplify GOMP_MAP_STRUCT handling of (COMPONENT_REF ↵Catherine Moore2021-02-131-34/+0
| | | | | | | | (INDIRECT_REF ...)) map clauses." This reverts commit bf8605f14ec33ea31233a3567f3184fee667b695. Regressions were reported for BZ 9574 - C++ class member mapping.
* openmp: Add support for non-rectangular loops in taskloop constructJakub Jelinek2021-02-094-0/+903
| | | | | | | | | | | | | | | | | | | | 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)
* openmp: Handle even some combined non-rectangular loopsJakub Jelinek2021-02-093-0/+386
| | | | | | | | | | | | | | | | | | 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)
* openmp: Adjust outer bounds of non-rect loopsJakub Jelinek2021-02-092-0/+237
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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)
* openmp: Optimize triangular loop logical iterator to actual iterators ↵Jakub Jelinek2021-02-093-0/+178
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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)
* openmp: Non-rectangular loop support for non-composite worksharing loops and ↵Jakub Jelinek2021-02-093-0/+442
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | distribute This implements the fallback mentioned in https://gcc.gnu.org/pipermail/gcc/2020-June/232874.html Special cases for triangular loops etc. to follow later, also composite constructs not supported yet (need to check the passing of temporaries around) and lastprivate might not give the same answers as serial loop if the last innermost body iteration isn't the last one for some of the outer loops (that will need to be solved separately together with rectangular loops that have no innermost body iterations, but some of the outer loops actually iterate). Also, simd needs work. 2020-06-27 Jakub Jelinek <jakub@redhat.com> * omp-general.h (struct omp_for_data_loop): Add non_rect_referenced member, move outer member. (struct omp_for_data): Add first_nonrect and last_nonrect members. * omp-general.c (omp_extract_for_data): Initialize first_nonrect, last_nonrect and non_rect_referenced members. * omp-expand.c (expand_omp_for_init_counts): Handle non-rectangular loops. (expand_omp_for_init_vars): Add nonrect_bounds parameter. Handle non-rectangular loops. (extract_omp_for_update_vars): Likewise. (expand_omp_for_generic, expand_omp_for_static_nochunk, expand_omp_for_static_chunk, expand_omp_simd, expand_omp_taskloop_for_outer, expand_omp_taskloop_for_inner): Adjust expand_omp_for_init_vars and extract_omp_for_update_vars callers. (expand_omp_for): Don't sorry on non-composite worksharing-loop or distribute. * testsuite/libgomp.c/loop-17.c: New test. * testsuite/libgomp.c/loop-18.c: New test. (cherry picked from commit aed3ab253dada2b7d2ed63cc6a8e15e263d5dd35)
* Enable gimplify GOMP_MAP_STRUCT handling of (COMPONENT_REF (INDIRECT_REF ↵Chung-Lin Tang2021-02-081-0/+34
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | ...)) map clauses. 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-08 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.
* OpenMP 5.0: requires directiveChung-Lin Tang2021-02-0210-0/+134
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This is a merge of: https://gcc.gnu.org/pipermail/gcc-patches/2021-January/563393.html This patch completes more of the reverse_offload, unified_address, and unified_shared_memory clauses for the OpenMP 5.0 requires directive, including runtime verification of the offload target. (currently no offload devices actually support above features, only warning messages are emitted) This may possibly reverted/updated when a final patch is approved for mainline. 2021-02-02 Chung-Lin Tang <cltang@codesourcery.com> gcc/c/ChangeLog: * c-parser.c (c_parser_declaration_or_fndef): Set OMP_REQUIRES_TARGET_USED in omp_requires_mask if function has "omp declare target" attribute. (c_parser_omp_target_data): Set OMP_REQUIRES_TARGET_USED in omp_requires_mask. (c_parser_omp_target_enter_data): Likewise. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_requires): Adjust to only mention "not implemented yet" for OMP_REQUIRES_DYNAMIC_ALLOCATORS. gcc/cp/ChangeLog: * parser.c (cp_parser_simple_declaration): Set OMP_REQUIRES_TARGET_USED in omp_requires_mask if function has "omp declare target" attribute. (cp_parser_omp_target_data): Set OMP_REQUIRES_TARGET_USED in omp_requires_mask. (cp_parser_omp_target_enter_data): Likewise. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_requires): Adjust to only mention "not implemented yet" for OMP_REQUIRES_DYNAMIC_ALLOCATORS. gcc/fortran/ChangeLog: * openmp.c (gfc_check_omp_requires): Fix REVERSE_OFFLOAD typo. (gfc_match_omp_requires): Adjust to only mention "not implemented yet" for OMP_REQUIRES_DYNAMIC_ALLOCATORS. * parse.c ("tree.h"): Add include. ("omp-general.h"): Likewise. (gfc_parse_file): Add code to merge omp_requires to omp_requires_mask. gcc/ChangeLog: * omp-offload.c (omp_finish_file): Add code to reate OpenMP requires mask variable in .gnu.gomp_requires section if needed. gcc/testsuite/ChangeLog: * c-c++-common/gomp/requires-4.c: Remove prune of "not supported yet". * gcc/testsuite/gfortran.dg/gomp/requires-4.f90: Fix REVERSE_OFFLOAD typo. * gcc/testsuite/gfortran.dg/gomp/requires-8.f90: Likewise. include/ChangeLog: * gomp-constants.h (GOMP_REQUIRES_UNIFIED_ADDRESS): New symbol. (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY): Likewise. (GOMP_REQUIRES_REVERSE_OFFLOAD): Likewise. libgcc/ChangeLog: * offloadstuff.c (__requires_mask_table): New symbol to mark start of .gnu.gomp_requires section. (__requires_mask_table_end): New symbol to mark end of .gnu.gomp_requires section. libgomp/ChangeLog: * libgomp-plugin.h (GOMP_OFFLOAD_supported_features): New declaration. * libgomp.h (struct gomp_device_descr): New 'supported_features_func' plugin hook field. * oacc-host.c (host_supported_features): New host hook function. (host_dispatch): Initialize 'supported_features_func' host hook. * plugin/plugin-gcn.c (GOMP_OFFLOAD_supported_features): New function. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_supported_features): Likewise. * target.c (<stdio.h>): Add include of standard header. (gomp_requires_mask): New static variable. (__requires_mask_table): New declaration. (__requires_mask_table_end): Likewise. (gomp_load_plugin_for_device): Add loading of 'supported_features' hook. (gomp_target_init): Add code to summarize .gnu._gomp_requires section mask values, emit error if inconsistency found. * testsuite/libgomp.c-c++-common/requires-1.c: New test. * testsuite/libgomp.c-c++-common/requires-1-aux.c: New file linked with above test. * testsuite/libgomp.c-c++-common/requires-2.c: New test. * testsuite/libgomp.c-c++-common/requires-2-aux.c: New file linked with above test. liboffloadmic/ChangeLog: * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_supported_features): New function.
* OpenMP 5.0: map this[:1] in C++ non-static member functions (PR 92120)Chung-Lin Tang2021-02-027-28/+338
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This is a merge of: https://gcc.gnu.org/pipermail/gcc-patches/2020-November/558975.html This patch creates automatic mapping of map(this[:1]) and pointer members as zero-length array sections, as specified by the OpenMP 5.0 specification. This may possibly reverted/updated when a final patch is approved for mainline. 2021-02-02 Chung-Lin Tang <cltang@codesourcery.com> PR middle-end/92120 gcc/cp/ChangeLog: * cp-tree.h (finish_omp_target): New declaration. (set_omp_target_this_expr): Likewise. * lambda.c (lambda_expr_this_capture): Add call to set_omp_target_this_expr. * parser.c (cp_parser_omp_target): Factor out code, change to call finish_omp_target, add re-initing call to set_omp_target_this_expr. * semantics.c (omp_target_this_expr): New static variable. (omp_target_ptr_members_accessed): New static hash_map for tracking accessed non-static pointer-type members. (finish_non_static_data_member): Add call to set_omp_target_this_expr. Add recording of non-static pointer-type members access. (finish_this_expr): Add call to set_omp_target_this_expr. (set_omp_target_this_expr): New function to set omp_target_this_expr. (finish_omp_target): New function with code merged from cp_parser_omp_target, plus code to implement this[:1] and __closure map clauses for OpenMP. (handle_omp_array_sections_1): Move code to peel of '*' for reference-based COMPONENT_REFs before FIELD_DECL transforming. (finish_omp_clauses): Handle 'A->member' case in map clauses. gcc/ChangeLog: * omp-low.c (lower_omp_target): Handle GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. * tree-pretty-print.c (dump_omp_clause): Likewise. gcc/testsuite/ChangeLog: * g++.dg/gomp/target-this-1.C: New testcase. * g++.dg/gomp/target-this-2.C: New testcase. * g++.dg/gomp/target-this-3.C: New testcase. * g++.dg/gomp/target-this-4.C: New testcase. include/ChangeLog: * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds. (GOMP_MAP_POINTER_P): Include GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION. libgomp/ChangeLog: * libgomp.h (gomp_attach_pointer): Add bool parameter. * oacc-mem.c (acc_attach_async): Update call to gomp_attach_pointer. (goacc_enter_data_internal): Likewise. * target.c (gomp_map_vars_existing): Update assert condition to include GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION. (gomp_map_pointer): Add 'bool allow_zero_length_array_sections' parameter, add support for mapping a pointer with NULL target. (gomp_attach_pointer): Add 'bool allow_zero_length_array_sections' parameter, add support for attaching a pointer with NULL target. (gomp_map_vars_internal): Update calls to gomp_map_pointer and gomp_attach_pointer, add handling for GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION cases. * testsuite/libgomp.c++/target-this-1.C: New testcase. * testsuite/libgomp.c++/target-this-2.C: New testcase. * testsuite/libgomp.c++/target-this-3.C: New testcase. * testsuite/libgomp.c++/target-this-4.C: New testcase.
* OpenMP 5.0 Structure element mappingChung-Lin Tang2021-01-2711-155/+688
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This is a merge of: https://gcc.gnu.org/pipermail/gcc-patches/2020-December/561139.html This patch implements the changes to the behavior of target mapping structure elements, as specified in OpenMP 5.0. This may possibly reverted/updated when a final patch is approved for mainline. libgomp/ChangeLog: * hashtab.h (htab_clear): New function with initialization code factored out from... (htab_create): ...here, adjust to use htab_clear function. * libgomp.h (REFCOUNT_SPECIAL): New symbol to denote range of special refcount values, add comments. (REFCOUNT_INFINITY): Adjust definition to use REFCOUNT_SPECIAL. (REFCOUNT_LINK): Likewise. (REFCOUNT_STRUCTELEM): New special refcount range for structure element siblings. (REFCOUNT_STRUCTELEM_P): Macro for testing for structure element sibling maps. (REFCOUNT_STRUCTELEM_FLAG_FIRST): Flag to indicate first sibling. (REFCOUNT_STRUCTELEM_FLAG_LAST): Flag to indicate last sibling. (REFCOUNT_STRUCTELEM_FIRST_P): Macro to test _FIRST flag. (REFCOUNT_STRUCTELEM_LAST_P): Macro to test _LAST flag. (struct splay_tree_key_s): Add structelem_refcount and structelem_refcount_ptr fields into a union with dynamic_refcount. Add comments. (gomp_map_vars): Delete declaration. (gomp_map_vars_async): Likewise. (gomp_unmap_vars): Likewise. (gomp_unmap_vars_async): Likewise. (goacc_map_vars): New declaration. (goacc_unmap_vars): Likewise. * oacc-mem.c (acc_map_data): Adjust to use goacc_map_vars. (goacc_enter_datum): Likewise. (goacc_enter_data_internal): Likewise. * oacc-parallel.c (GOACC_parallel_keyed): Adjust to use goacc_map_vars and goacc_unmap_vars. (GOACC_data_start): Adjust to use goacc_map_vars. (GOACC_data_end): Adjust to use goacc_unmap_vars. * target.c (hash_entry_type): New typedef. (htab_alloc): New function hook for hashtab.h. (htab_free): Likewise. (htab_hash): Likewise. (htab_eq): Likewise. (hashtab.h): Add file include. (gomp_increment_refcount): New function. (gomp_decrement_refcount): Likewise. (gomp_map_vars_existing): Add refcount_set parameter, adjust to use gomp_increment_refcount. (gomp_map_fields_existing): Add refcount_set parameter, adjust calls to gomp_map_vars_existing. (gomp_map_vars_internal): Add refcount_set parameter, add local openmp_p variable to guard OpenMP specific paths, adjust calls to gomp_map_vars_existing, add structure element sibling splay_tree_key sequence creation code, adjust Fortran map case to avoid increment under OpenMP. (gomp_map_vars): Adjust to static, add refcount_set parameter, manage local refcount_set if caller passed in NULL, adjust call to gomp_map_vars_internal. (gomp_map_vars_async): Adjust and rename into... (goacc_map_vars): ...this new function, adjust call to gomp_map_vars_internal. (gomp_remove_splay_tree_key): New function with code factored out from gomp_remove_var_internal. (gomp_remove_var_internal): Add code to handle removing multiple splay_tree_key sequence for structure elements, adjust code to use gomp_remove_splay_tree_key for splay-tree key removal. (gomp_unmap_vars_internal): Add refcount_set parameter, adjust to use gomp_decrement_refcount. (gomp_unmap_vars): Adjust to static, add refcount_set parameter, manage local refcount_set if caller passed in NULL, adjust call to gomp_unmap_vars_internal. (gomp_unmap_vars_async): Adjust and rename into... (goacc_unmap_vars): ...this new function, adjust call to gomp_unmap_vars_internal. (GOMP_target): Manage refcount_set and adjust calls to gomp_map_vars and gomp_unmap_vars. (GOMP_target_ext): Likewise. (gomp_target_data_fallback): Adjust call to gomp_map_vars. (GOMP_target_data): Likewise. (GOMP_target_data_ext): Likewise. (GOMP_target_end_data): Adjust call to gomp_unmap_vars. (gomp_exit_data): Add refcount_set parameter, adjust to use gomp_decrement_refcount, adjust to queue splay-tree keys for removal after main loop. (GOMP_target_enter_exit_data): Manage refcount_set and adjust calls to gomp_map_vars and gomp_exit_data. (gomp_target_task_fn): Likewise. * testsuite/libgomp.c-c++-common/refcount-1.c: New testcase. * testsuite/libgomp.c-c++-common/struct-elem-1.c: New testcase. * testsuite/libgomp.c-c++-common/struct-elem-2.c: New testcase. * testsuite/libgomp.c-c++-common/struct-elem-3.c: New testcase. * testsuite/libgomp.c-c++-common/struct-elem-4.c: New testcase. * testsuite/libgomp.c-c++-common/struct-elem-5.c: New testcase.
* openmp: Fix intermittent hanging of task-detach-6 libgomp tests [PR98738]Kwok Cheung Yeung2021-01-227-54/+147
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This adds support for the task detach clause to taskwait, and fixes a number of problems related to semaphores that may lead to a hang in some circumstances. 2021-01-21 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ PR libgomp/98738 * libgomp.h (enum gomp_task_kind): Add GOMP_TASK_DETACHED. * task.c (task_fulfilled_p): Check detach field as well. (GOMP_task): Add thread to debug messages. Use address of task as the event handle. (gomp_barrier_handle_tasks): Fix indentation. Use address of task as event handle. Set kind of suspended detach task to GOMP_TASK_DETACHED and decrement task_running_count. Move finish_cancelled block out of else branch. Skip decrement of task_running_count if task kind is GOMP_TASK_DETACHED. (GOMP_taskwait): Finish fulfilled detach tasks. Update comment. Queue detach tasks that have not been fulfilled. (omp_fulfill_event): Use address of task as event handle. Post to taskwait_sem and taskgroup_sem if necessary. Check task_running_count before calling gomp_team_barrier_wake. * testsuite/libgomp.c-c++-common/task-detach-5.c (main): Change data-sharing of detach events on enclosing parallel to private. * testsuite/libgomp.c-c++-common/task-detach-6.c (main): Likewise. * testsuite/libgomp.fortran/task-detach-5.f90 (task_detach_5): Likewise. * testsuite/libgomp.fortran/task-detach-6.f90 (task_detach_6): Likewise.
* libgomp: Fix up GOMP_task on s390xJakub Jelinek2021-01-222-11/+21
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | On Wed, Jan 20, 2021 at 05:04:39PM +0100, Florian Weimer wrote: > Sorry, this appears to cause OpenMP task state corruption in RPM. We > have only seen this on s390x. Haven't actually verified it, but my suspection is that this is a caller stack corruption. We play with fire with the GOMP_task API/ABI extensions, the GOMP_task function used to be: void GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *), long arg_size, long arg_align, bool if_clause, unsigned flags); and later: void GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *), long arg_size, long arg_align, bool if_clause, unsigned flags, void **depend); and later: void GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *), long arg_size, long arg_align, bool if_clause, unsigned flags, void **depend, int priority); and now: void GOMP_task (void (*fn) (void *), void *data, void (*cpyfn) (void *, void *), long arg_size, long arg_align, bool if_clause, unsigned flags, void **depend, int priority, void *detach) and which of those depend, priority and detach argument is present depends on the bits in flags. I'm afraid the compiler just decided to spill the detach = NULL store in if ((flags & GOMP_TASK_FLAG_DETACH) == 0) detach = NULL; on s390x into the argument stack slot. Not a problem if the caller passes all those 10 arguments, but if not, can clobber random stack location. This hack should fix it up. Priority doesn't need changing, but I've changed it anyway just to be safe. With the patch none of the 3 arguments are ever modified, so I'd hope gcc doesn't decide to spill something unrelated there. 2021-01-20 Jakub Jelinek <jakub@redhat.com> * task.c (GOMP_task): Rename priority argument to priority_arg, add priority automatic variable and modify that variable. Instead of clearing detach argument when GOMP_TASK_FLAG_DETACH bit is not set, check flags for that bit. (cherry picked from commit 0bb27b81a762d3c607bd25409337c749f836c0cd)
* openmp: Don't optimize shared to firstprivate on task with depend clauseJakub Jelinek2021-01-222-0/+54
| | | | | | | | | | | | | | | | | | | | | The attached testcase is miscompiled, because we optimize shared clauses to firstprivate when task body can't modify the variable even when the task has depend clause. That is wrong, because firstprivate means the variable will be copied immediately when the task is created, while with depend clause some other task might change it later before the dependencies are satisfied and the task should observe the value only after the change. 2020-12-18 Jakub Jelinek <jakub@redhat.com> * gimplify.c (struct gimplify_omp_ctx): Add has_depend member. (gimplify_scan_omp_clauses): Set it to true if OMP_CLAUSE_DEPEND appears on OMP_TASK. (gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Force GOVD_WRITTEN on shared variables if task construct has depend clause. * testsuite/libgomp.c/task-6.c: New test. (cherry picked from commit 8b60459465252c7d47b58abf83fae2aa84915b03)
* RTEMS: Fix libgomp buildSebastian Huber2021-01-222-0/+12
| | | | | | | | libgomp/ * config/rtems/sem.h (gomp_sem_getcount): New function. (cherry picked from commit 0f951b3dd34b355579b4c9a9e287d32ac771bc67)
* libgomp: Don't access gomp_sem_t as int using atomics unconditionallyJakub Jelinek2021-01-226-2/+62
| | | | | | | | | | | | | | | | | This patch introduces gomp_sem_getcount wrapper, which uses sem_getvalue for POSIX and atomic loads for linux futex and accel. rtems for now remains broken. 2021-01-18 Jakub Jelinek <jakub@redhat.com> * config/linux/sem.h (gomp_sem_getcount): New function. * config/posix/sem.h (gomp_sem_getcount): New function. * config/posix/sem.c (gomp_sem_getcount): New function. * config/accel/sem.h (gomp_sem_getcount): New function. * task.c (task_fulfilled_p): Use gomp_sem_getcount. (omp_fulfill_event): Likewise. (cherry picked from commit d3b41bde961713ff4af7e18011126434c497edba)
* openmp: Add support for the OpenMP 5.0 task detach clauseKwok Cheung Yeung2021-01-2224-22/+684
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 2021-01-16 Kwok Cheung Yeung <kcy@codesourcery.com> gcc/ * builtin-types.def (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT): Rename to... (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR): ...this. Add extra argument. * gimplify.c (omp_default_clause): Ensure that event handle is firstprivate in a task region. (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_DETACH. (gimplify_adjust_omp_clauses): Likewise. * omp-builtins.def (BUILT_IN_GOMP_TASK): Change function type to BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR. * omp-expand.c (expand_task_call): Add GOMP_TASK_FLAG_DETACH to flags if detach clause specified. Add detach argument when generating call to GOMP_task. * omp-low.c (scan_sharing_clauses): Setup data environment for detach clause. (finish_taskreg_scan): Move field for variable containing the event handle to the front of the struct. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_DETACH. Fix ordering. * tree-nested.c (convert_nonlocal_omp_clauses): Handle OMP_CLAUSE_DETACH clause. (convert_local_omp_clauses): Handle OMP_CLAUSE_DETACH clause. * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_DETACH. * tree.c (omp_clause_num_ops): Add entry for OMP_CLAUSE_DETACH. Fix ordering. (omp_clause_code_name): Add entry for OMP_CLAUSE_DETACH. Fix ordering. (walk_tree_1): Handle OMP_CLAUSE_DETACH. gcc/c-family/ * c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_DETACH. Redefine PRAGMA_OACC_CLAUSE_DETACH. gcc/c/ * c-parser.c (c_parser_omp_clause_detach): New. (c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH clause. (OMP_TASK_CLAUSE_MASK): Add mask for PRAGMA_OMP_CLAUSE_DETACH. * c-typeck.c (c_finish_omp_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH clause. Prevent use of detach with mergeable and overriding the data sharing mode of the event handle. gcc/cp/ * parser.c (cp_parser_omp_clause_detach): New. (cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH. (OMP_TASK_CLAUSE_MASK): Add mask for PRAGMA_OMP_CLAUSE_DETACH. * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_DETACH clause. * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_DETACH clause. Prevent use of detach with mergeable and overriding the data sharing mode of the event handle. gcc/fortran/ * dump-parse-tree.c (show_omp_clauses): Handle detach clause. * frontend-passes.c (gfc_code_walker): Walk detach expression. * gfortran.h (struct gfc_omp_clauses): Add detach field. (gfc_c_intptr_kind): New. * openmp.c (gfc_free_omp_clauses): Free detach clause. (gfc_match_omp_detach): New. (enum omp_mask1): Add OMP_CLAUSE_DETACH. (enum omp_mask2): Remove OMP_CLAUSE_DETACH. (gfc_match_omp_clauses): Handle OMP_CLAUSE_DETACH for OpenMP. (OMP_TASK_CLAUSES): Add OMP_CLAUSE_DETACH. (resolve_omp_clauses): Prevent use of detach with mergeable and overriding the data sharing mode of the event handle. * trans-openmp.c (gfc_trans_omp_clauses): Handle detach clause. * trans-types.c (gfc_c_intptr_kind): New. (gfc_init_kinds): Initialize gfc_c_intptr_kind. * types.def (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT): Rename to... (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR): ...this. Add extra argument. gcc/testsuite/ * c-c++-common/gomp/task-detach-1.c: New. * g++.dg/gomp/task-detach-1.C: New. * gcc.dg/gomp/task-detach-1.c: New. * gfortran.dg/gomp/task-detach-1.f90: New. include/ * gomp-constants.h (GOMP_TASK_FLAG_DETACH): New. libgomp/ * fortran.c (omp_fulfill_event_): New. * libgomp.h (struct gomp_task): Add detach and completion_sem fields. (struct gomp_team): Add task_detach_queue and task_detach_count fields. * libgomp.map (OMP_5.0.1): Add omp_fulfill_event and omp_fulfill_event_. * libgomp_g.h (GOMP_task): Add extra argument. * omp.h.in (enum omp_event_handle_t): New. (omp_fulfill_event): New. * omp_lib.f90.in (omp_event_handle_kind): New. (omp_fulfill_event): New. * omp_lib.h.in (omp_event_handle_kind): New. (omp_fulfill_event): Declare. * priority_queue.c (priority_tree_find): New. (priority_list_find): New. (priority_queue_find): New. * priority_queue.h (priority_queue_predicate): New. (priority_queue_find): New. * task.c (gomp_init_task): Initialize detach field. (task_fulfilled_p): New. (GOMP_task): Add detach argument. Ignore detach argument if GOMP_TASK_FLAG_DETACH not set in flags. Initialize completion_sem field. Copy address of completion_sem into detach argument and into the start of the data record. Wait for detach event if task not deferred. (gomp_barrier_handle_tasks): Queue tasks with unfulfilled events. Remove completed tasks and requeue dependent tasks. (omp_fulfill_event): New. * team.c (gomp_new_team): Initialize task_detach_queue and task_detach_count fields. (free_team): Free task_detach_queue field. * testsuite/libgomp.c-c++-common/task-detach-1.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-2.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-3.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-4.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-5.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-6.c: New testcase. * testsuite/libgomp.fortran/task-detach-1.f90: New testcase. * testsuite/libgomp.fortran/task-detach-2.f90: New testcase. * testsuite/libgomp.fortran/task-detach-3.f90: New testcase. * testsuite/libgomp.fortran/task-detach-4.f90: New testcase. * testsuite/libgomp.fortran/task-detach-5.f90: New testcase. * testsuite/libgomp.fortran/task-detach-6.f90: New testcase. (cherry picked from commit a6d22fb21c6f1ad7e8b6b722bfc0e7e11f50cb92)
* [og10] openacc: Adjust loop lowering for AMD GCNJulian Brown2021-01-136-8/+89
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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. 2021-01-13 Julian Brown <julian@codesourcery.com> gcc/ * omp-offload.c (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/testsuite/ * libgomp.oacc-c-c++-common/loop-gwv-1.c: Adjust for loop lowering changes. * libgomp.oacc-c-c++-common/loop-wv-1.c: Likewise. * libgomp.oacc-c-c++-common/loop-red-gwv-1.c: Likewise. * libgomp.oacc-c-c++-common/loop-red-wv-1.c: Likewise. * libgomp.oacc-c-c++-common/routine-gwv-1.c: Likewise. * libgomp.oacc-c-c++-common/routine-wv-1.c: Likewise.
* nvptx: Cache stacks block for OpenMP kernel launchJulian Brown2021-01-052-18/+115
| | | | | | | | | | | | | | | | | | | | | 2021-01-05 Julian Brown <julian@codesourcery.com> libgomp/ * plugin/plugin-nvptx.c (SOFTSTACK_CACHE_LIMIT): New define. (struct ptx_device): Add omp_stacks struct. (nvptx_open_device): Initialise cached-stacks housekeeping info. (nvptx_close_device): Free cached stacks block and mutex. (nvptx_stacks_free): New function. (nvptx_alloc): Add SUPPRESS_ERRORS parameter. (GOMP_OFFLOAD_alloc): Add strategies for freeing soft-stacks block. (nvptx_stacks_alloc): Rename to... (nvptx_stacks_acquire): This. Cache stacks block between runs if same size or smaller is required. (nvptx_stacks_free): Remove. (GOMP_OFFLOAD_run): Call nvptx_stacks_acquire and lock stacks block during kernel execution. (cherry picked from commit 6b577a17b26347e78c8b9167f24fc5c9d9724270)
* OpenMP: Add implicit declare target for nested proceduresTobias Burnus2020-12-182-0/+52
| | | | | | | | | | | | | | This is a backport from mainline of commit 8b0a63e47cd83f4e8534d0d201739bdd10f321a2. gcc/ChangeLog: * omp-offload.c (omp_discover_implicit_declare_target): Also handled nested functions. libgomp/ChangeLog: * testsuite/libgomp.fortran/declare-target-3.f90: New test.
* openmp: Implement OpenMP 5.0 base-pointer attachement and clause orderingChung-Lin Tang2020-12-074-13/+122
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This patch implements some parts of the target variable mapping changes specified in OpenMP 5.0, including base-pointer attachment/detachment behavior for array section list-items in map clauses, and ordering of map clauses according to map kind. 2020-11-10 Chung-Lin Tang <cltang@codesourcery.com> gcc/c-family/ChangeLog: * c-common.h (c_omp_adjust_map_clauses): New declaration. * c-omp.c (struct map_clause): Helper type for c_omp_adjust_map_clauses. (c_omp_adjust_map_clauses): New function. gcc/c/ChangeLog: * c-parser.c (c_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (c_parser_omp_target_enter_data): Likewise. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_target): Likewise. * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/cp/ChangeLog: * parser.c (cp_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (cp_parser_omp_target_enter_data): Likewise. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_target): Likewise. * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix interaction between reference case and attach/detach. (finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/ChangeLog: * gimplify.c (is_or_contains_p): New static helper function. (omp_target_reorder_clauses): New function. (gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to reorder clause list according to OpenMP 5.0 rules. Add handling of GOMP_MAP_ATTACH_DETACH for OpenMP cases. * omp-low.c (is_omp_target): New static helper function. (scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. (lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. gcc/testsuite/ChangeLog: * c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid. * gfortran.dg/gomp/map-2.f90: Likewise. * c-c++-common/gomp/map-5.c: New testcase. libgomp/ChangeLog: * libgomp.h (enum gomp_map_vars_kind): Adjust enum values to be bit-flag usable. * oacc-mem.c (acc_map_data): Adjust gomp_map_vars argument flags to 'GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA'. (goacc_enter_datum): Likewise for call to gomp_map_vars_async. (goacc_enter_data_internal): Likewise. * target.c (gomp_map_vars_internal): Change checks of GOMP_MAP_VARS_ENTER_DATA to use bit-and (&). Adjust use of gomp_attach_pointer for OpenMP cases. (gomp_exit_data): Add handling of GOMP_MAP_DETACH. (GOMP_target_enter_exit_data): Add handling of GOMP_MAP_ATTACH. * testsuite/libgomp.c-c++-common/ptr-attach-1.c: New testcase.
* openmp: Retire nest-var ICV for OpenMP 5.1Kwok Cheung Yeung2020-11-196-40/+132
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This removes the nest-var ICV, expressing nesting in terms of the max-active-levels-var ICV instead. The max-active-levels-var ICV is now per data environment rather than per device. This is a backport from mainline (commit 6fae7eda968db658c280ad6f94fe6906a15af0c9). 2020-11-18 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * env.c (gomp_global_icv): Remove nest_var field. Add max_active_levels_var field. (gomp_max_active_levels_var): Remove. (parse_boolean): Return true on success. (handle_omp_display_env): Express OMP_NESTED in terms of max_active_levels_var. Change format specifier for max_active_levels_var. (initialize_env): Set max_active_levels_var from OMP_MAX_ACTIVE_LEVELS, OMP_NESTED, OMP_NUM_THREADS and OMP_PROC_BIND. * icv.c (omp_set_nested): Express in terms of max_active_levels_var. (omp_get_nested): Likewise. (omp_set_max_active_levels): Use max_active_levels_var field instead of gomp_max_active_levels_var. (omp_get_max_active_levels): Likewise. * libgomp.h (struct gomp_task_icv): Remove nest_var field. Add max_active_levels_var field. (gomp_supported_active_levels): Set to UCHAR_MAX. (gomp_max_active_levels_var): Delete. * libgomp.texi (omp_get_nested): Update documentation. (omp_set_nested): Likewise. (OMP_MAX_ACTIVE_LEVELS): Likewise. (OMP_NESTED): Likewise. (OMP_NUM_THREADS): Likewise. (OMP_PROC_BIND): Likewise. * parallel.c (gomp_resolve_num_threads): Replace reference to nest_var with max_active_levels_var. Use max_active_levels_var field instead of gomp_max_active_levels_var.
* openmp: Mark deprecated symbols in OpenMP 5.0Kwok Cheung Yeung2020-11-1333-15/+121
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This is a backport from mainline (commit 10508db867934264bbc2578f1f454c19fa558fd3). 2020-11-05 Ulrich Drepper <drepper@redhat.com> Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * Makefile.am (%.mod): Add -cpp and -fopenmp to compile flags. * Makefile.in: Regenerate. * fortran.c: Wrap uses of omp_set_nested and omp_get_nested with pragmas to ignore -Wdeprecated-declarations warnings. * icv.c: Likewise. * omp.h.in (__GOMP_DEPRECATED_5_0): Define. Mark omp_lock_hint_* enum values, omp_lock_hint_t, omp_set_nested, and omp_get_nested with __GOMP_DEPRECATED_5_0. * omp_lib.f90.in: Mark omp_get_nested and omp_set_nested as deprecated. * testsuite/libgomp.c++/affinity-1.C: Add -Wno-deprecated-declarations to test options. * testsuite/libgomp.c/affinity-1.c: Likewise. * testsuite/libgomp.c/affinity-2.c: Likewise. * testsuite/libgomp.c/appendix-a/a.15.1.c: Likewise. * testsuite/libgomp.c/lib-1.c: Likewise. * testsuite/libgomp.c/nested-1.c: Likewise. * testsuite/libgomp.c/nested-2.c: Likewise. * testsuite/libgomp.c/nested-3.c: Likewise. * testsuite/libgomp.c/pr32362-1.c: Likewise. * testsuite/libgomp.c/pr32362-2.c: Likewise. * testsuite/libgomp.c/pr32362-3.c: Likewise. * testsuite/libgomp.c/pr35549.c: Likewise. * testsuite/libgomp.c/pr42942.c: Likewise. * testsuite/libgomp.c/pr61200.c: Likewise. * testsuite/libgomp.c/sort-1.c: Likewise. * testsuite/libgomp.c/target-5.c: Likewise. * testsuite/libgomp.c/target-6.c: Likewise. * testsuite/libgomp.c/teams-1.c: Likewise. * testsuite/libgomp.c/thread-limit-1.c: Likewise. * testsuite/libgomp.c/thread-limit-2.c: Likewise. * testsuite/libgomp.c/thread-limit-4.c: Likewise. * testsuite/libgomp.fortran/affinity1.f90: Likewise. * testsuite/libgomp.fortran/lib1.f90: Likewise. * testsuite/libgomp.fortran/lib2.f: Likewise. * testsuite/libgomp.fortran/nested1.f90: Likewise. * testsuite/libgomp.fortran/teams1.f90: Likewise.
* openmp: Add test for OMP_TARGET_OFFLOAD=mandatory for cases where it must ↵Jakub Jelinek2020-11-132-0/+40
| | | | | | | | | | | not fail This is a backport from mainline (commit 17c5b7e1dc47bab6e6cedbf4b2d88cef3283533e). 2020-10-22 Jakub Jelinek <jakub@redhat.com> * testsuite/libgomp.c/target-41.c: New test.
* openmp: Change omp_get_initial_device () to match OpenMP 5.1 requirementsJakub Jelinek2020-11-118-21/+97
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | > Therefore, I think until omp_get_initial_device () value is changed, we The following so far untested patch implements that change. OpenMP 4.5 said for omp_get_initial_device: The value of the device number is implementation defined. If it is between 0 and one less than omp_get_num_devices() then it is valid for use with all device constructs and routines; if it is outside that range, then it is only valid for use with the device memory routines and not in the device clause. and OpenMP 5.0 similarly, but OpenMP 5.1 says: The value of the device number is the value returned by the omp_get_num_devices routine. As the new value is compatible with what has been required earlier, I think we can change it already now. This is a backport from mainline (commit 74c9882b80bda50b37c9555498de7123c6bdb9e4). 2020-10-22 Jakub Jelinek <jakub@redhat.com> * icv.c (omp_get_initial_device): Remove including corresponding ialias. * icv-device.c (omp_get_initial_device): New function. Return gomp_get_num_devices (). Add ialias. * target.c (resolve_device): Don't fail with OMP_TARGET_OFFLOAD=mandatory if device_id is equal to gomp_get_num_devices (). (omp_target_alloc, omp_target_free, omp_target_is_present, omp_target_memcpy, omp_target_memcpy_rect, omp_target_associate_ptr, omp_target_disassociate_ptr, omp_pause_resource): Use gomp_get_num_devices () instead of GOMP_DEVICE_HOST_FALLBACK on the first use in the functions, in uses dominated by the gomp_get_num_devices call use num_devices_openmp instead. * libgomp.texi (omp_get_initial_device): Document. * config/gcn/icv-device.c (omp_get_initial_device): New function. Add ialias. * config/nvptx/icv-device.c (omp_get_initial_device): Likewise. * testsuite/libgomp.c/target-40.c: New test.
* openmp: Implement support for OMP_TARGET_OFFLOAD environment variableKwok Cheung Yeung2020-11-115-44/+192
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This implements support for the OMP_TARGET_OFFLOAD environment variable introduced in the OpenMP 5.0 standard, which controls how offloading is handled. It may be set to MANDATORY (abort if offloading cannot be performed), DISABLED (no offloading to devices) or DEFAULT (offload to device if possible, fall back to host if not). This is a backport from mainline (commits 1bfc07d150790fae93184a79a7cce897655cb37b, 35f258f4bbba7fa044f90b4f14d1bc942db58089 and 121a8812c45b3155ccbd268b000ad00a778e81e8). 2020-10-20 Kwok Cheung Yeung <kcy@codesourcery.com> Jakub Jelinek <jakub@redhat.com> libgomp/ * env.c (gomp_target_offload_var): New. (parse_target_offload): New. (handle_omp_display_env): Print value of OMP_TARGET_OFFLOAD. (initialize_env): Parse OMP_TARGET_OFFLOAD. * libgomp.h (gomp_target_offload_t): New. (gomp_target_offload_var): New. * libgomp.texi (OMP_TARGET_OFFLOAD): New section. * target.c (resolve_device): Generate error if device not found and offloading is mandatory. (gomp_target_fallback): Generate error if offloading is mandatory. (GOMP_target): Add argument in call to gomp_target_fallback. (GOMP_target_ext): Likewise. (gomp_target_data_fallback): Generate error if offloading is mandatory. (GOMP_target_data): Add argument in call to gomp_target_data_fallback. (GOMP_target_data_ext): Likewise. (gomp_target_task_fn): Add argument in call to gomp_target_fallback. (gomp_target_init): Return early if offloading is disabled. Inside of the function, use automatic variables corresponding to num_devices, num_devices_openmp and devices global variables and update the globals only at the end of the function.
* openmp: Add support for the omp_get_supported_active_levels runtime library ↵Kwok Cheung Yeung2020-11-1112-6/+100
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | routine This patch implements the omp_get_supported_active_levels runtime routine from the OpenMP 5.0 specification, which returns the maximum number of active nested parallel regions supported by this implementation. The current maximum (set using the omp_set_max_active_levels routine or the OMP_MAX_ACTIVE_LEVELS environment variable) cannot exceed this number. This is a backport from mainline (commits 8949b985dbaf07d433bd57d2883e1e5414f20e75 and 445567b22a3c535be0b1861b393e9a0b050f2b1e). 2020-10-13 Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * env.c (gomp_max_active_levels_var): Initialize to gomp_supported_active_levels. (initialize_env): Limit gomp_max_active_levels_var to be at most equal to gomp_supported_active_levels. * fortran.c (omp_get_supported_active_levels): Add ialias_redirect. (omp_get_supported_active_levels_): New. * icv.c (omp_set_max_active_levels): Limit gomp_max_active_levels_var to at most equal to gomp_supported_active_levels. (omp_get_supported_active_levels): New. * libgomp.h (gomp_supported_active_levels): New. * libgomp.map (OMP_5.0.1): Add omp_get_supported_active_levels and omp_get_supported_active_levels_. * libgomp.texi (omp_get_max_active_levels): Modify description. (omp_get_supported_active_levels): New. (omp_set_max_active_levels): Update. Add reference to omp_get_supported_active_levels. * omp.h.in (omp_get_supported_active_levels): New. * omp_lib.f90.in (omp_get_supported_active_levels): New. * omp_lib.h.in (omp_get_supported_active_levels): New. * testsuite/libgomp.c/lib-2.c (main): Check omp_get_max_active_levels against omp_get_supported_active_levels. * testsuite/libgomp.fortran/lib4.f90 (lib4): Likewise.
* OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390)Tobias Burnus2020-09-283-2/+9
| | | | | | | | | | | | | | | | | | | | Backport mainline version; updates commit ef509d1985aa53a8c0875c25ad4050ea807be10e for later review-comment changes. gcc/ChangeLog: PR middle-end/96390 * omp-offload.c (omp_discover_declare_target_tgt_fn_r): Handle alias nodes. libgomp/ChangeLog: PR middle-end/96390 * testsuite/libgomp.c++/pr96390.C: New test. * testsuite/libgomp.c-c++-common/pr96390.c: New test. (cherry picked from commit 2a10a2c0689db280ee3a94164504b7196b8370f4)
* nvptx: fix malformed patchAndrew Stubbs2020-09-191-1/+0
| | | | | | | | An edit got botched. Fixed now. libgomp/ * config/nvptx/bar.c (gomp_team_barrier_wake): Remove rogue asm.
* libgomp: disable barriers in nested teamsAndrew Stubbs2020-09-183-11/+54
| | | | | | | | | | | | | | | | | | | | | Both GCN and NVPTX allow nested parallel regions, but the barrier implementation did not allow the nested teams to run independently of each other (due to hardware limitations). This patch fixes that, under the assumption that each thread will create a new subteam of one thread, by simply not using barriers when there's no other thread to synchronise. libgomp/ChangeLog: * config/gcn/bar.c (gomp_barrier_wait_end): Skip the barrier if the total number of threads is one. (gomp_team_barrier_wake): Likewise. (gomp_team_barrier_wait_end): Likewise. (gomp_team_barrier_wait_cancel_end): Likewise. * config/nvptx/bar.c (gomp_barrier_wait_end): Likewise. (gomp_team_barrier_wake): Likewise. (gomp_team_barrier_wait_end): Likewise. (gomp_team_barrier_wait_cancel_end): Likewise. * testsuite/libgomp.c-c++-common/nested-parallel-unbalanced.c: New test.
* libgomp.c-c++-common/pr96390.c: XFAIL on nvptx.Tobias Burnus2020-09-183-1/+8
| | | | | | | libgomp/ * testsuite/libgomp.c-c++-common/pr96390.c: XFAIL on nvptx. * testsuite/libgomp.c++/pr96390.C: Likewise.
* OpenMP: Handle cpp_implicit_alias in declare-target discovery (PR96390)Tobias Burnus2020-09-173-0/+81
| | | | | | | | | | | | | | | | | Submitted to mainline at https://gcc.gnu.org/pipermail/gcc-patches/2020-September/554142.html gcc/ChangeLog: PR middle-end/96390 * omp-offload.c (omp_discover_declare_target_tgt_fn_r): Handle alias nodes. libgomp/ChangeLog: PR middle-end/96390 * testsuite/libgomp.c++/pr96390.C: New test. * testsuite/libgomp.c-c++-common/pr96390.c: New test.
View Lorry Controller Status