mirror of
https://forge.sourceware.org/marek/gcc.git
synced 2026-02-22 03:47:02 -05:00
reflection
353 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
Thomas Schwinge
|
3dc9eedd95 |
libgomp: Add a few more OpenMP/USM test cases
... where there are clear differences in behavior for OpenMP/USM run-time configurations. We shall further clarify all the intended semantics, once the implementation begins to differentiate OpenMP 'requires unified_shared_memory' vs. 'requires self_maps'. libgomp/ * testsuite/libgomp.c-c++-common/map-arrayofstruct-2-usm.c: New. * testsuite/libgomp.c-c++-common/map-arrayofstruct-3-usm.c: Likewise. * testsuite/libgomp.c-c++-common/struct-elem-5-usm.c: Likewise. * testsuite/libgomp.c-c++-common/target-present-1-usm.c: Likewise. * testsuite/libgomp.c-c++-common/target-present-2-usm.c: Likewise. * testsuite/libgomp.c-c++-common/target-present-3-usm.c: Likewise. * testsuite/libgomp.fortran/map-subarray-5-usm.f90: Likewise. * testsuite/libgomp.fortran/map-subarray-6-usm.f90: Likewise. * testsuite/libgomp.fortran/map-subarray-7-usm.f90: Likewise. * testsuite/libgomp.fortran/target-allocatable-1-1-usm.f90: Likewise. * testsuite/libgomp.fortran/target-allocatable-1-2-usm.f90: Likewise. * testsuite/libgomp.fortran/target-enter-data-2-usm.F90: Likewise. * testsuite/libgomp.fortran/target-present-1-usm.f90: Likewise. * testsuite/libgomp.fortran/target-present-2-usm.f90: Likewise. * testsuite/libgomp.fortran/target-present-3-usm.f90: Likewise. * testsuite/libgomp.fortran/target-allocatable-1-1.f90: Adjust. * testsuite/libgomp.fortran/target-allocatable-1-2.f90: Likewise. * testsuite/libgomp.fortran/target-present-1.f90: Likewise. * testsuite/libgomp.fortran/target-present-2.f90: Likewise. * testsuite/libgomp.fortran/target-present-3.f90: Likewise. |
||
|
Thomas Schwinge
|
105fddf356 |
amdgcn: Adjust failure mode for gfx908 USM: 'libgomp.fortran/map-alloc-comp-9-usm.f90'
The change/rationale that commit
|
||
|
Thomas Schwinge
|
954f804b73 |
openmp: Bump Version from 4.5 to 5.2 (2/4): Some more '-Wno-deprecated-openmp'
These changes should've been included in
commit
|
||
|
supers1ngular
|
cdb15b3900 |
openmp: Improve Fortran Diagnostics for Linear Clause
This patch improves diagnostics for the linear clause, providing a more accurate and intuitive recommendation for remediation if the deprecated syntax is used. Additionally updates the relevant test to reflect the changed verbiage of the warning. gcc/fortran/ChangeLog: * openmp.cc (gfc_match_omp_clauses): New diagnostic logic. libgomp/ChangeLog: * testsuite/libgomp.fortran/pr84418-1.f90: Fix verbiage of dg-warning to reflect updated warning. |
||
|
Tobias Burnus
|
50e2c80545 |
libgomp.fortran/uses_allocators_1.f90: Fix dg-error for r16-6273
Missed to commit dg-error changes for the new diagnostic due to commit r16-6273-g7044071f07d763 OpenMP: uses_allocators with ';'-separated list libgomp/ChangeLog: * testsuite/libgomp.fortran/uses_allocators_1.f90: Update dg-error. |
||
|
Tobias Burnus
|
7044071f07 |
OpenMP: uses_allocators with ';'-separated list
OpenMP 6.0 has the following wording for the uses_allocators clause:
"More than one clause-argument-specification may be specified";
this permits ';' lists. While that's pointless for predefined
allocators, for user-defined allocators it saves redundant
') uses_allocators(' by permitting:
uses_allocators( traits(t1): alloc1 ; traits(t2): alloc2 )
Additionally, the order in the tree dump has been changed to
place the modifiers before the allocator variable, matching
the input syntax.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_clause_uses_allocators): Accept
multiple clause-argument-specifications separated by ';'.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_clause_uses_allocators): Accept
multiple clause-argument-specifications separated by ';'.
gcc/fortran/ChangeLog:
* openmp.cc (gfc_match_omp_clause_uses_allocators): Accept
multiple clause-argument-specifications separated by ';'.
gcc/ChangeLog:
* tree-pretty-print.cc (dump_omp_clause): For uses_allocators,
print modifier before allocator variable.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/uses_allocators-7.f90: Add ';' test.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/uses_allocators-8.c: New test.
|
||
|
Tobias Burnus
|
350794cb07 |
Fix libgomp.fortran/dep-uses-allocators.f90
libgomp/ChangeLog: * testsuite/libgomp.fortran/dep-uses-allocators.f90: Properly escape '(..)' in dg-warning. |
||
|
Tobias Burnus
|
d02eebea7e |
OpenMP/Fortran: uses_allocators - suggest 5.2 format in the warning
Actually mention how the new 5.2+ syntax looks like when outputting the deprecation warning for 'uses_allocators'. gcc/fortran/ChangeLog: * openmp.cc (gfc_match_omp_clause_uses_allocators): Mention new syntax in deprecation warning. libgomp/ChangeLog: * testsuite/libgomp.fortran/dep-uses-allocators.f90: Update dg-warning. |
||
|
Chung-Lin Tang
|
0a01b42b22 |
OpenMP: Add uses_allocators parser support to C/C++
This is the parser part for C/C++, including early middle end bits, but then stops with a 'sorry, unimplemented'. It also adds support for omp_null_alloctor (6.0 clarificiation, is to be ignored). As predefined allocators do not require any special handling in GCC, those are ignored. Therefore, this patch fully supports uses_allocators that only use predefined allocators - only printing a sorry for those that use the (implicit) traits/memspace modifer. (The parsing support for Fortran was added before; this patch just adds omp_null_allocator support to Fortran. The sorry message for Fortran is also still in the FE and not in gimplify.cc, but that only make a difference for the original dump.) Except for some minor fixes, this is the same patch as https://gcc.gnu.org/pipermail/gcc-patches/2025-November/700345.html with the middle-end + libgomp handling excluded. That patch in turn is based on previous patches, the latest previous one was https://gcc.gnu.org/pipermail/gcc-patches/2023-November/637415.html and, in particular, the C/C++ parser style was updated following the review comments. Also, more C++ template-handling fixes have been applied. gcc/c-family/ChangeLog: * c-omp.cc (c_omp_split_clauses): Hande uses_allocators. * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_clause_uses_allocators): New function. (c_parser_omp_clause_name, c_parser_omp_all_clauses, OMP_TARGET_CLAUSE_MASK): Handle uses_allocators. * c-typeck.cc (c_finish_omp_clauses): Likewise. gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_clause_uses_allocators): New function. (cp_parser_omp_clause_name, cp_parser_omp_all_clauses, OMP_TARGET_CLAUSE_MASK): Handle uses_allocators. * semantics.cc (finish_omp_clauses): Likewise. * pt.cc (tsubst_omp_clauses): Likewise. gcc/fortran/ChangeLog: * openmp.cc (resolve_omp_clauses): Handle omp_null_allocator. * trans-openmp.cc (gfc_trans_omp_clauses): Mention it in a comment. gcc/ChangeLog: * gimplify.cc (gimplify_scan_omp_clauses): Handle uses_allocators by printing a 'sorry, unimplemented' and removing it. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_USES_ALLOCATORS. * tree.cc (omp_clause_num_ops, omp_clause_code_name): Likewise. * tree-pretty-print.cc (dump_omp_clause): Handle it. * tree.h (OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR, OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE, OMP_CLAUSE_USES_ALLOCATORS_TRAITS): New. libgomp/ChangeLog: * testsuite/libgomp.fortran/uses_allocators_1.f90: Add check for omp_null_allocator. * testsuite/libgomp.fortran/uses_allocators-7.f90: New test. gcc/testsuite/ChangeLog: * c-c++-common/gomp/uses_allocators-1.c: New test. * c-c++-common/gomp/uses_allocators-2.c: New test. * c-c++-common/gomp/uses_allocators-4.c: New test. * c-c++-common/gomp/uses_allocators-7.c: New test. * g++.dg/gomp/deprecate-2.C: New test. * g++.dg/gomp/uses_allocators-1.C: New test. * gcc.dg/gomp/deprecate-2.c: New test. Co-authored-by: Tobias Burnus <tburnus@baylibre.com> Co-authored-by: Andrew Stubbs <ams@baylibre.com> |
||
|
supers1ngular
|
52d7f5b103 |
openmp: Bump Version from 4.5 to 5.2 (3/4)
Implements the OpenMP 5.2 Fortran deprecations. Uses the warning established in patch 1/4, -Wdeprecated-openmp, for said deprecations. Similarly, we do not implement the relaxing of constraints for the interop construct since it is not a deprecation. However, the deprecation for 'uses_allocators' is implemented, since support exists in Fortran mainline. Additionally implements the Fortran-specific deprecation for executable allocate directives, and adds new tests. gcc/fortran/ChangeLog: * openmp.cc (gfc_match_omp_clause_reduction): Deprecate '-' operator for reductions. (gfc_match_omp_clause_uses_allocators): Deprecate allocator(traits) pattern for 'uses_allocators'. (gfc_match_omp_clauses): Deprecate 'sink' and 'source' for 'depend' clause. Deprecate list items as arguments with 'linear' clause. Deprecate non-comma-separated modifiers for the map clause. Deprecate 'to' clause with declare target. (gfc_match_omp_declare_target): Whitespace. (match_omp_metadirective): Deprecate 'default' clause on metadirectives. (resolve_omp_clauses): Deprecate executable allocate directives. libgomp/ChangeLog: * testsuite/libgomp.fortran/allocate-8a.f90: Suppress warnings. * testsuite/libgomp.fortran/allocators-1.f90: Ditto. * testsuite/libgomp.fortran/allocators-2.f90: Ditto. * testsuite/libgomp.fortran/allocators-4.f90: Ditto. * testsuite/libgomp.fortran/declare-target-1.f90: Ditto. * testsuite/libgomp.fortran/declare-target-2.f90: Ditto. * testsuite/libgomp.fortran/declare-target-indirect-1.f90: Ditto. * testsuite/libgomp.fortran/declare-target-indirect-2.f90: Ditto. * testsuite/libgomp.fortran/doacross1.f90: Ditto. * testsuite/libgomp.fortran/doacross2.f90: Ditto. * testsuite/libgomp.fortran/doacross3.f90: Ditto. * testsuite/libgomp.fortran/map-alloc-ptr-2.f90: Ditto. * testsuite/libgomp.fortran/pr84418-1.f90: Ditto. * testsuite/libgomp.fortran/pr84418-2.f90: Ditto. * testsuite/libgomp.fortran/reduction1.f90: Ditto. * testsuite/libgomp.fortran/udr11.f90: Ditto. * testsuite/libgomp.fortran/uses_allocators_1.f90: Ditto. * testsuite/libgomp.fortran/uses_allocators_2.f90: Ditto. * testsuite/libgomp.fortran/dep-uses-allocators.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/allocate-14.f90: Suppress warnings. * gfortran.dg/gomp/allocate-16.f90: Ditto. * gfortran.dg/gomp/allocate-5.f90: Ditto. * gfortran.dg/gomp/allocate-6.f90: Ditto. * gfortran.dg/gomp/allocate-7.f90: Ditto. * gfortran.dg/gomp/allocators-3.f90: Ditto. * gfortran.dg/gomp/declare-simd-2.f90: Ditto. * gfortran.dg/gomp/declare-simd-6.f90: Ditto. * gfortran.dg/gomp/declare-target-1.f90: Ditto. * gfortran.dg/gomp/declare-target-2.f90: Ditto. * gfortran.dg/gomp/declare-target-4.f90: Ditto. * gfortran.dg/gomp/declare-target-5.f90: Ditto. * gfortran.dg/gomp/declare-target-indirect-1.f90: Ditto. * gfortran.dg/gomp/declare-target-indirect-2.f90: Ditto. * gfortran.dg/gomp/declare-variant-10.f90: Ditto. * gfortran.dg/gomp/declare-variant-8.f90: Ditto. * gfortran.dg/gomp/implicit-save.f90: Ditto. * gfortran.dg/gomp/linear-1.f90: Ditto. * gfortran.dg/gomp/linear-2.f90: Ditto. * gfortran.dg/gomp/linear-3.f90: Ditto. * gfortran.dg/gomp/linear-4.f90: Ditto. * gfortran.dg/gomp/linear-6.f90: Ditto. * gfortran.dg/gomp/map-12.f90: Ditto. * gfortran.dg/gomp/map-6.f90: Ditto. * gfortran.dg/gomp/map-7.f90: Ditto. * gfortran.dg/gomp/map-8.f90: Ditto. * gfortran.dg/gomp/order-8.f90: Ditto. * gfortran.dg/gomp/pr83977.f90: Ditto. * gfortran.dg/gomp/reduction1.f90: Ditto. * gfortran.dg/gomp/schedule-modifiers-2.f90: Ditto. * gfortran.dg/gomp/workshare-reduction-55.f90: Ditto. * gfortran.dg/gomp/workshare-reduction-56.f90: Ditto. * gfortran.dg/gomp/workshare-reduction-57.f90: Ditto. * gfortran.dg/gomp/workshare-reduction-58.f90: Ditto. * gfortran.dg/gomp/52-deps.f90: New test. |
||
|
supers1ngular
|
a57b0e3ecb |
[PATCH v2 1/4] openmp: Bump Version from 4.5 to 5.2 (1/4)
Bumps OpenMP from 4.5 (201511) to 5.2 (202111), with deprecation and
test support to 5.1 (202011). Adds new tests and a new warning.
Suppresses deprecation warnings in all relevant tests and removes
suppression pragmas visible outside of the testsuite. Additionally
implements new warning in the relevant frontends. Otherwise, cleans
up some whitespace and fixed a misspelled pragma in a testcase. Also
fixes an indentation error.
gcc/c-family/ChangeLog:
* c-cppbuiltin.cc (c_cpp_builtins): Bump _OPENMP version.
* c.opt (Wdeprecated-openmp): Add warning.
* c.opt.urls: Regenerated.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_clause_proc_bind): Deprecate master
affinity.
(c_parser_omp_master): Deprecate master construct.
(c_parser_transaction): Whitespace.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_postfix_expression): Whitespace.
(cp_parser_builtin_c23_va_start): Ditto.
(cp_parser_omp_clause_proc_bind): Deprecate master affinity.
(cp_parser_omp_master): Deprecate master construct.
gcc/ChangeLog:
* doc/invoke.texi: Update docs for '-Wdeprecated-openmp'.
gcc/fortran/ChangeLog:
* cpp.cc (cpp_define_builtins): Bump _OPENMP version.
* invoke.texi: Update docs for '-Wdeprecated-openmp'.
* lang.opt (Wdeprecated-openmp): Add warning.
* lang.opt.urls: Regenerated.
* openmp.cc (gfc_match_omp_clauses): Deprecate master affinity
policy.
(gfc_match_omp_parallel_master): Deprecate master construct.
(gfc_match_omp_parallel_master_taskloop): Ditto.
(gfc_match_omp_parallel_master_taskloop_simd): Ditto.
(gfc_match_omp_master): Ditto.
(gfc_match_omp_master_taskloop): Ditto.
(gfc_match_omp_master_taskloop_simd): Ditto.
(resolve_omp_clauses): Warn for deprecated use of
{use,is}_device_ptr.
libgomp/ChangeLog:
* env.c (omp_display_env): Bump _OPENMP version.
* fortran.c (ialias_redirect): Remove suppression pragmas.
(omp_set_dynamic_8_): Ditto.
(omp_set_nested_8_): Ditto.
(omp_get_nested_): Ditto.
* icv.c (omp_get_dynamic): Ditto.
(omp_get_nested): Ditto.
(ialias): Ditto.
* omp_lib.f90.in: Bump openmp_version.
* omp_lib.h.in: Ditto.
* testsuite/libgomp.c++/affinity-1.C: Suppress deprecation
warnings.
* testsuite/libgomp.c++/ctor-1.C: Ditto.
* testsuite/libgomp.c++/ctor-11.C: Ditto.
* testsuite/libgomp.c++/ctor-13.C: Ditto.
* testsuite/libgomp.c++/ctor-2.C: Ditto.
* testsuite/libgomp.c++/ctor-5.C: Ditto.
* testsuite/libgomp.c++/ctor-7.C: Ditto.
* testsuite/libgomp.c++/depend-iterator-1.C: Ditto.
* testsuite/libgomp.c++/loop-13.C: Ditto.
* testsuite/libgomp.c++/master-1.C: Ditto.
* testsuite/libgomp.c++/pr26943.C: Ditto.
* testsuite/libgomp.c++/pr81130.C: Ditto.
* testsuite/libgomp.c++/pr81314.C: Ditto.
* testsuite/libgomp.c++/target-in-reduction-1.C: Ditto.
* testsuite/libgomp.c++/target-in-reduction-2.C: Ditto.
* testsuite/libgomp.c++/task-1.C: Ditto.
* testsuite/libgomp.c++/task-2.C: Ditto.
* testsuite/libgomp.c++/task-6.C: Ditto.
* testsuite/libgomp.c++/task-reduction-7.C: Ditto.
* testsuite/libgomp.c++/task-reduction-9.C: Ditto.
* testsuite/libgomp.c++/taskloop-reduction-1.C: Ditto.
* testsuite/libgomp.c-c++-common/cancel-taskgroup-4.c: Ditto.
* testsuite/libgomp.c-c++-common/depend-inoutset-1.c: Ditto.
* testsuite/libgomp.c-c++-common/depend-iterator-1.c: Ditto.
* testsuite/libgomp.c-c++-common/master-combined-1.c: Ditto.
* testsuite/libgomp.c-c++-common/target-in-reduction-1.c: Ditto.
* testsuite/libgomp.c-c++-common/target-in-reduction-2.c: Ditto.
* testsuite/libgomp.c-c++-common/task-detach-12.c: Ditto.
* testsuite/libgomp.c-c++-common/task-reduction-15.c: Ditto.
* testsuite/libgomp.c-c++-common/task-reduction-5.c: Ditto.
* testsuite/libgomp.c-c++-common/task-reduction-6.c: Ditto.
* testsuite/libgomp.c-c++-common/task-reduction-8.c: Ditto.
* testsuite/libgomp.c-c++-common/taskloop-reduction-1.c: Ditto.
* testsuite/libgomp.c-c++-common/taskloop-reduction-3.c: Ditto.
* testsuite/libgomp.c-c++-common/taskloop-reduction-4.c: Ditto.
* testsuite/libgomp.c/affinity-1.c: Remove extraneous dg
instruction and add suppression.
* testsuite/libgomp.c/critical-2.c: Suppress deprecation
warnings.
* testsuite/libgomp.c/debug-1.c: Ditto.
* testsuite/libgomp.c/lib-1.c: Ditto.
* testsuite/libgomp.c/loop-24.c: Ditto.
* testsuite/libgomp.c/nestedfn-2.c: Ditto.
* testsuite/libgomp.c/nestedfn-3.c: Ditto.
* testsuite/libgomp.c/pr104385.c: Ditto.
* testsuite/libgomp.c/target-31.c: Ditto.
* testsuite/libgomp.c/target-34.c: Ditto.
* testsuite/libgomp.c/target-critical-1.c: Ditto.
* testsuite/libgomp.c/task-1.c: Ditto.
* testsuite/libgomp.c/task-3.c: Ditto.
* testsuite/libgomp.c/task-6.c: Ditto.
* testsuite/libgomp.c/task-reduction-1.c: Ditto.
* testsuite/libgomp.c/task-reduction-2.c: Ditto.
* testsuite/libgomp.c/teams-1.c: Ditto.
* testsuite/libgomp.c/vla-1.c: Ditto.
* testsuite/libgomp.fortran/crayptr1.f90: Ditto.
* testsuite/libgomp.fortran/depend-inoutset-1.f90: Ditto.
* testsuite/libgomp.fortran/is_device_ptr-1.f90: Ditto.
* testsuite/libgomp.fortran/is_device_ptr-2.f90: Ditto.
* testsuite/libgomp.fortran/lib1.f90: Ditto.
* testsuite/libgomp.fortran/lib2.f: Ditto.
* testsuite/libgomp.fortran/lib3.f: Ditto.
* testsuite/libgomp.fortran/omp_parse2.f90: Ditto.
* testsuite/libgomp.fortran/openmp_version-1.f: Bump OMP version.
* testsuite/libgomp.fortran/openmp_version-2.f90: Ditto.
* testsuite/libgomp.fortran/parallel-master.f90: Suppress
warnings.
* testsuite/libgomp.fortran/pointer2.f90: Ditto.
* testsuite/libgomp.fortran/reduction6.f90: Ditto.
* testsuite/libgomp.fortran/target-firstprivate-1.f90: Ditto.
* testsuite/libgomp.fortran/use_device_addr-1.f90: Ditto.
* testsuite/libgomp.fortran/use_device_addr-2.f90: Ditto.
* testsuite/libgomp.fortran/use_device_addr-3.f90: Ditto.
* testsuite/libgomp.fortran/use_device_addr-4.f90: Ditto.
* testsuite/libgomp.fortran/use_device_addr-5.f90: Ditto.
* testsuite/libgomp.fortran/use_device_ptr-1.f90: Ditto.
* testsuite/libgomp.fortran/use_device_ptr-3.f90: Ditto.
* testsuite/libgomp.fortran/use_device_ptr-4.f90: Ditto.
* testsuite/libgomp.fortran/use_device_ptr-optional-1.f90: Ditto.
* testsuite/libgomp.fortran/use_device_ptr-optional-2.f90: Ditto.
* testsuite/libgomp.fortran/use_device_ptr-optional-3.f90: Ditto.
* testsuite/libgomp.c-c++-common/omp-atv-seq-dep.c: New test.
* testsuite/libgomp.c-c++-common/omp-lock-hint-contended-dep.c:
New test.
* testsuite/libgomp.c-c++-common/omp-lock-hint-none-dep.c: New test.
* testsuite/libgomp.c-c++-common/omp-lock-hint-speculative-dep.c:
New test.
* testsuite/libgomp.c-c++-common/omp-lock-hint-uncontended-dep.c:
New test.
* testsuite/libgomp.c/omp-proc-bind-master-dep.c: New test.
* testsuite/libgomp.fortran/omp-atv-seq-dep.f90: New test.
* testsuite/libgomp.fortran/omp-lock-hint-contended-dep.f90: New
test.
* testsuite/libgomp.fortran/omp-lock-hint-none-dep.f90: New test.
* testsuite/libgomp.fortran/omp-lock-hint-speculative-dep.f90: New
test.
* testsuite/libgomp.fortran/omp-lock-hint-uncontended-dep.f90: New
test.
gcc/testsuite/ChangeLog:
* c-c++-common/cpp/openmp-define-3.c: Bump OMP version.
* c-c++-common/gomp/Wparentheses-1.c: Suppress deprecation
warnings.
* c-c++-common/gomp/Wparentheses-3.c: Ditto.
* c-c++-common/gomp/affinity-3.c: Ditto.
* c-c++-common/gomp/allocate-18.c: Ditto.
* c-c++-common/gomp/cancel-1.c: Ditto.
* c-c++-common/gomp/clause-dups-1.c: Ditto.
* c-c++-common/gomp/clauses-1.c: Suppress deprecation
warnings and fix misspelled directive. Add
'-Wunknown-pragmas'.
* c-c++-common/gomp/clauses-6.c: Suppress deprecation warnings.
* c-c++-common/gomp/declare-variant-1.c: Ditto.
* c-c++-common/gomp/declare-variant-2.c: Ditto.
* c-c++-common/gomp/depend-iterator-1.c: Ditto.
* c-c++-common/gomp/lastprivate-conditional-1.c: Ditto.
* c-c++-common/gomp/loop-1.c: Ditto.
* c-c++-common/gomp/loop-2.c: Ditto.
* c-c++-common/gomp/loop-3.c: Ditto.
* c-c++-common/gomp/loop-4.c: Ditto.
* c-c++-common/gomp/master-combined-1.c: Ditto.
* c-c++-common/gomp/master-combined-2.c: Ditto.
* c-c++-common/gomp/nesting-2.c: Ditto.
* c-c++-common/gomp/pr100902-1.c: Ditto.
* c-c++-common/gomp/pr61486-2.c: Ditto.
* c-c++-common/gomp/pr85696.c: Ditto.
* c-c++-common/gomp/pr85956.c: Ditto.
* c-c++-common/gomp/pr98187.c: Ditto.
* c-c++-common/gomp/pr99928-1.c: Ditto.
* c-c++-common/gomp/pr99928-10.c: Ditto.
* c-c++-common/gomp/pr99928-11.c: Ditto.
* c-c++-common/gomp/pr99928-12.c: Ditto.
* c-c++-common/gomp/pr99928-13.c: Ditto.
* c-c++-common/gomp/pr99928-14.c: Ditto.
* c-c++-common/gomp/pr99928-2.c: Ditto.
* c-c++-common/gomp/pr99928-3.c: Ditto.
* c-c++-common/gomp/pr99928-4.c: Ditto.
* c-c++-common/gomp/pr99928-5.c: Ditto.
* c-c++-common/gomp/pr99928-6.c: Ditto.
* c-c++-common/gomp/pr99928-7.c: Ditto.
* c-c++-common/gomp/pr99928-8.c: Ditto.
* c-c++-common/gomp/pr99928-9.c: Ditto.
* c-c++-common/gomp/task-detach-1.c: Ditto.
* c-c++-common/gomp/teams-2.c: Ditto.
* g++.dg/gomp/attrs-1.C: Ditto.
* g++.dg/gomp/attrs-2.C: Ditto.
* g++.dg/gomp/attrs-4.C: Ditto.
* g++.dg/gomp/block-0.C: Ditto.
* g++.dg/gomp/block-10.C: Ditto.
* g++.dg/gomp/block-5.C: Ditto.
* g++.dg/gomp/block-9.C: Ditto.
* g++.dg/gomp/depend-iterator-1.C: Ditto.
* g++.dg/gomp/master-1.C: Ditto.
* g++.dg/gomp/master-2.C: Ditto.
* g++.dg/gomp/master-3.C: Ditto.
* g++.dg/gomp/method-1.C: Ditto.
* g++.dg/gomp/pr29965-3.C: Ditto.
* g++.dg/gomp/pr29965-9.C: Ditto.
* g++.dg/gomp/pr78363-4.C: Ditto.
* g++.dg/gomp/pr78363-6.C: Ditto.
* g++.dg/gomp/pr79664.C: Ditto.
* g++.dg/gomp/pr94477.C: Ditto.
* g++.dg/gomp/pr94512.C: Ditto.
* g++.dg/gomp/tpl-master-1.C: Ditto.
* gcc.dg/gomp/appendix-a/a.12.1.c: Ditto.
* gcc.dg/gomp/appendix-a/a.33.2.c: Ditto.
* gcc.dg/gomp/attrs-1.c: Ditto.
* gcc.dg/gomp/attrs-2.c: Ditto.
* gcc.dg/gomp/attrs-4.c: Ditto.
* gcc.dg/gomp/block-10.c: Ditto.
* gcc.dg/gomp/block-5.c: Ditto.
* gcc.dg/gomp/block-9.c: Ditto.
* gcc.dg/gomp/master-1.c: Ditto.
* gcc.dg/gomp/master-2.c: Ditto.
* gcc.dg/gomp/master-3.c: Ditto.
* gcc.dg/gomp/nesting-1.c: Ditto.
* gcc.dg/gomp/pr104517.c: Ditto.
* gcc.dg/gomp/pr29965-3.c: Ditto.
* gcc.dg/gomp/pr35818.c: Ditto.
* gcc.dg/gomp/pr91216.c: Ditto.
* gcc.dg/gomp/sharing-2.c: Ditto.
* gfortran.dg/gomp/adjust-args-10.f90: Ditto.
* gfortran.dg/gomp/affinity-1.f90: Ditto.
* gfortran.dg/gomp/allocate-clause.f90: Ditto.
* gfortran.dg/gomp/appendix-a/a.12.1.f90: Ditto.
* gfortran.dg/gomp/appendix-a/a.33.2.f90: Ditto.
* gfortran.dg/gomp/c_ptr_tests_20.f90: Ditto.
* gfortran.dg/gomp/c_ptr_tests_21.f90: Ditto.
* gfortran.dg/gomp/cancel-1.f90: Ditto.
* gfortran.dg/gomp/clauses-1.f90: Ditto.
* gfortran.dg/gomp/declare-variant-1.f90: Ditto.
* gfortran.dg/gomp/depend-iterator-1.f90: Ditto.
* gfortran.dg/gomp/depend-iterator-2.f90: Ditto.
* gfortran.dg/gomp/is_device_ptr-1.f90: Ditto.
* gfortran.dg/gomp/is_device_ptr-2.f90: Ditto.
* gfortran.dg/gomp/is_device_ptr-3.f90: Ditto.
* gfortran.dg/gomp/lastprivate-conditional-1.f90: Ditto.
* gfortran.dg/gomp/loop-4.f90: Ditto.
* gfortran.dg/gomp/loop-exit.f90: Ditto.
* gfortran.dg/gomp/map-3.f90: Ditto.
* gfortran.dg/gomp/nesting-2.f90: Ditto.
* gfortran.dg/gomp/nesting-3.f90: Ditto.
* gfortran.dg/gomp/nowait-2.f90: Ditto.
* gfortran.dg/gomp/nowait-4.f90: Ditto.
* gfortran.dg/gomp/nowait-5.f90: Ditto.
* gfortran.dg/gomp/openmp-simd-2.f90: Ditto.
* gfortran.dg/gomp/openmp-simd-3.f90: Ditto.
* gfortran.dg/gomp/parallel-master-1.f90: Ditto.
* gfortran.dg/gomp/parallel-master-2.f90: Ditto.
* gfortran.dg/gomp/pr107214-8.f90: Ditto.
* gfortran.dg/gomp/pr48117.f90: Ditto.
* gfortran.dg/gomp/pr94672.f90: Ditto.
* gfortran.dg/gomp/pr99928-1.f90: Suppression + fix whitespace.
* gfortran.dg/gomp/pr99928-11.f90: Suppression.
* gfortran.dg/gomp/pr99928-2.f90: Suppression + fix whitespace.
* gfortran.dg/gomp/pr99928-3.f90: Ditto.
* gfortran.dg/gomp/pr99928-4.f90: Ditto.
* gfortran.dg/gomp/pr99928-5.f90: Ditto.
* gfortran.dg/gomp/pr99928-6.f90: Ditto.
* gfortran.dg/gomp/pr99928-8.f90: Ditto.
* gfortran.dg/gomp/sharing-3.f90: Suppress deprecation warnings.
* gfortran.dg/gomp/strictly-structured-block-1.f90: Ditto.
* gfortran.dg/gomp/strictly-structured-block-2.f90: Ditto.
* gfortran.dg/gomp/target1.f90: Ditto.
* gfortran.dg/gomp/taskloop-1.f90: Ditto.
* gfortran.dg/gomp/taskloop-2.f90: Ditto.
* gfortran.dg/openmp-define-3.f90: Bump expected version.
* c-c++-common/gomp/master-construct-dep.c: New test.
* gfortran.dg/gomp/master-construct-dep.f90: New test.
|
||
|
Andrew Stubbs
|
1cf9fda493 |
amdgcn: Adjust failure mode for gfx908 USM
Unified Shared Memory does not appear to work well on gfx908, which is why we disabled xnack by default. For this reason it makes sense to inform the user as compile time, but this is causing trouble in the testsuite which assumes that USM only fails at runtime. This patch changes the gfx908 compile time message to a warning only (in case some other target does this differently), and prevents the tests from attempting to run in host-fallback mode (given that that is not what they are trying to test). It also changes the existing warning to only fire once. The patch assumes that effective target "omp_usm" also implies self-maps. gcc/ChangeLog: * config/gcn/gcn.cc (gcn_init_cumulative_args): Only warn once. Use "required" instead of "enabled" in the warning. * config/gcn/mkoffload.cc (process_asm): Warn, don't error. Use "required" instead of "on" in the warning. libgomp/ChangeLog: * testsuite/lib/libgomp.exp (check_effective_target_omp_usm): New. * testsuite/libgomp.c++/target-std__array-concurrent-usm.C: Require working Unified Shared Memory to run the test. * testsuite/libgomp.c++/target-std__bitset-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__deque-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__forward_list-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__list-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__map-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__multimap-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__multiset-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__set-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__span-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__valarray-concurrent-usm.C: Likewise. * testsuite/libgomp.c++/target-std__vector-concurrent-usm.C: Likewise. * testsuite/libgomp.c-c++-common/target-implicit-map-4.c: Likewise. * testsuite/libgomp.c-c++-common/target-link-3.c: Likewise. * testsuite/libgomp.c-c++-common/target-link-4.c: Likewise. * testsuite/libgomp.fortran/self_maps.f90: Likewise. |
||
|
Andrew Stubbs
|
723b18ce3d |
libgomp, amdgcn: Implement Managed Memory
This patch implements "managed" memory for AMD GCN GPUs in OpenMP. It builds on the support added to the NVPTX libgomp for CUDA Managed Memory, a week or two ago. These features were first posted here a few years ago, as part of a larger Unified Shared Memory patch series, and then in a slightly changed version just over a year ago. Hopefully this time the controversial bits have been removed. Since we do not use HIP we cannot use hipMallocManaged, so this patch attempts to replicate the same effect by setting the appropriate attributes. This works on more devices than support proper USM, but still I cannot be sure that the settings are correct for every device out there (I have tested on gfx900, gfx906, gfx908, gfx90a, and gfx1100). The HSA header file update uses the most recent files relicensed for us by AMD, at the time of the first patch posting. Those files have certainly moved on in the upstream sources, but I did not ask to get those relicensed. include/ChangeLog: * hsa.h: Import newer version. * hsa_ext_amd.h: Likewise. * hsa_ext_image.h: Likewise. libgomp/ChangeLog: * Makefile.in: Regenerate. * libgomp-plugin.h (gomp_simple_alloc_init_context): New prototype. (gomp_simple_alloc_register_memory): New prototype. (gomp_simple_alloc): New prototype. (gomp_simple_free): New prototype. (gomp_simple_realloc): New prototype. * libgomp.h (gomp_simple_alloc_init_context): Move to libgomp-plugin.h. (gomp_simple_alloc_register_memory): Likewise. (gomp_simple_alloc): Likewise. (gomp_simple_free): Likewise. (gomp_simple_realloc): Likewise. * libgomp.texi: Update AMD managed memory description. * plugin/Makefrag.am (libgomp_plugin_gcn_la_SOURCES): Add simple-allocator.c and plugin/mutex.c. * plugin/plugin-gcn.c: Include sys/mman.h and unistd.h. (struct hsa_runtime_fn_info): Add hsa_amd_svm_attributes_set_fn. (dump_hsa_system_info): Add HSA_AMD_SYSTEM_INFO_SVM_SUPPORTED and HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT to the GCN_DEBUG output. (init_hsa_runtime_functions): Add hsa_amd_svm_attributes_set. (isa_matches_agent): Add a new error message for the case where the ISA doesn't match but the name does. (managed_ctx): New variable. (managed_heap_create): New function. (GOMP_OFFLOAD_managed_alloc): Likewise. (GOMP_OFFLOAD_managed_free): Likewise. * simple-allocator.c (gomp_fatal): New macro. * testsuite/lib/libgomp.exp (check_effective_target_omp_managedmem): Add amdgcn support checker. (check_effective_target_offload_target_amdgcn_with_xnack): New. * testsuite/libgomp.c-c++-common/requires-4.c: Ignore xnack warning. * testsuite/libgomp.c-c++-common/requires-4a.c: Ignore xnack warning. * testsuite/libgomp.c-c++-common/requires-5.c: Ignore xnack warning. * testsuite/libgomp.c++/alloc-managed-1.C: Add -mxnack=on, if needed. * testsuite/libgomp.c/alloc-managed-1.c: Likewise. * testsuite/libgomp.c/alloc-managed-2.c: Likewise. * testsuite/libgomp.c/alloc-managed-3.c: Likewise. * testsuite/libgomp.c/alloc-managed-4.c: Likewise. * testsuite/libgomp.fortran/alloc-managed-1.f90: Likewise. * plugin/mutex.c: New file. |
||
|
Paul-Antoine Arras
|
05c2ad4a2e |
OpenMP/Fortran: Allow explicit map followed by implicit deep mapping [PR120505]
Consider the following source code, assuming tiles is allocatable:
```
!$omp target enter data map(var%tiles(1)%den1, var%tiles(1)%den2) ! (1)
[...]
!$omp target ! implicitly maps var, which triggers deep mapping of tiles (2)
```
Each omp directive causes a run-time error in libgomp:
(1) libgomp: Mapped array elements must be the same (0x14d729c0 vs 0x14d72a18)
(2) libgomp: Trying to map into device [0x3704ca50..0x3704cb00) object when
[0x3704ca50..0x3704caa8) is already mapped
Regarding (1), the OpenMP spec has the following restriction: "If multiple list
items are explicitly mapped on the same construct and have the same containing
array or have base pointers that share original storage, and if any of the list
items do not have corresponding list items that are present in the device data
environment prior to a task encountering the construct, then the list items must
refer to *the same array elements* of either the containing array or the
implicit array of the base pointers."
Because tiles is allocatable, we cannot prove at compile time that array
elements are the same, so the check is deferred to libgomp. But there the
condition enforcing that all addresses are the same is too strict, so this patch
relaxes it to only check that addresses are sorted in increasing order.
The OpenMP spec allows (2) as long as it is implicit, without extending the
original mapping. So this patch sets the GOMP_MAP_IMPLICIT flag appropriately
on deep maps at compile time to let libgomp know that it is fine.
This patch ensures that such user code is accepted by:
(1) Setting the GOMP_MAP_IMPLICIT flag appropriately on deep maps;
(2) Relaxing the restriction on struct mapping from different containing arrays,
so that the element index need not be the same, instead addresses must be sorted
in increasing order.
This fixes the two errors currently seen when running SPEC HPC clvleaf
benchmark. However, further mapping issues prevent the benchmark from running to
completion.
PR fortran/120505
gcc/ChangeLog:
* omp-low.cc (lower_omp_target): Set GOMP_MAP_IMPLICIT flag.
libgomp/ChangeLog:
* target.c (gomp_map_vars_internal): Allow struct mapping from different
containing array elements as long as adresses are in increasing order.
* testsuite/libgomp.c-c++-common/map-arrayofstruct-2.c: Adjust
dg-output.
* testsuite/libgomp.c-c++-common/map-arrayofstruct-3.c: Likewise.
* testsuite/libgomp.fortran/map-subarray-5.f90: Likewise.
* testsuite/libgomp.fortran/map-subarray-10.f90: New test.
* testsuite/libgomp.fortran/map-subarray-9.f90: New test.
|
||
|
Andrew Stubbs
|
62174ec27b |
openmp, nvptx: ompx_gnu_managed_mem_alloc
This adds support for using Cuda Managed Memory with omp_alloc. AMD support will be added in a future patch. There is one new predefined allocator, "ompx_gnu_managed_mem_alloc", plus a corresponding memory space, which can be used to allocate memory in the "managed" space. The nvptx plugin is modified to make the necessary Cuda calls, via two new (optional) plugin interfaces. gcc/fortran/ChangeLog: * openmp.cc (is_predefined_allocator): Use GOMP_OMP_PREDEF_ALLOC_MAX and GOMP_OMPX_PREDEF_ALLOC_MIN/MAX instead of hardcoded values in the comment. include/ChangeLog: * cuda/cuda.h (cuMemAllocManaged): Add declaration and related CU_MEM_ATTACH_GLOBAL flag. * gomp-constants.h (GOMP_OMPX_PREDEF_ALLOC_MAX): Update to 201. (GOMP_OMP_PREDEF_MEMSPACE_MAX): New constant. (GOMP_OMPX_PREDEF_MEMSPACE_MIN): New constant. (GOMP_OMPX_PREDEF_MEMSPACE_MAX): New constant. libgomp/ChangeLog: * allocator.c (ompx_gnu_max_predefined_alloc): Update to ompx_gnu_managed_mem_alloc. (_Static_assert): Fix assertion messages for allocators and add new assertions for memspace constants. (omp_max_predefined_mem_space): New define. (ompx_gnu_min_predefined_mem_space): New define. (ompx_gnu_max_predefined_mem_space): New define. (MEMSPACE_ALLOC): Add check for non-standard memspaces. (MEMSPACE_CALLOC): Likewise. (MEMSPACE_REALLOC): Likewise. (MEMSPACE_VALIDATE): Likewise. (predefined_ompx_gnu_alloc_mapping): Add ompx_gnu_managed_mem_space. (omp_init_allocator): Add ompx_gnu_managed_mem_space validation. * config/gcn/allocator.c (gcn_memspace_alloc): Add check for non-standard memspaces. (gcn_memspace_calloc): Likewise. (gcn_memspace_realloc): Likewise. (gcn_memspace_validate): Update to validate standard vs non-standard memspaces. * config/linux/allocator.c (linux_memspace_alloc): Add managed memory space handling. (linux_memspace_calloc): Likewise. (linux_memspace_free): Likewise. (linux_memspace_realloc): Likewise (returns NULL for fallback). * config/nvptx/allocator.c (nvptx_memspace_alloc): Add check for non-standard memspaces. (nvptx_memspace_calloc): Likewise. (nvptx_memspace_realloc): Likewise. (nvptx_memspace_validate): Update to validate standard vs non-standard memspaces. * env.c (parse_allocator): Add ompx_gnu_managed_mem_alloc, ompx_gnu_managed_mem_space, and some static asserts so I don't forget them again. * libgomp-plugin.h (GOMP_OFFLOAD_managed_alloc): New declaration. (GOMP_OFFLOAD_managed_free): New declaration. * libgomp.h (gomp_managed_alloc): New declaration. (gomp_managed_free): New declaration. (struct gomp_device_descr): Add managed_alloc_func and managed_free_func fields. * libgomp.texi: Document ompx_gnu_managed_mem_alloc and ompx_gnu_managed_mem_space, add C++ template documentation, and describe NVPTX and AMD support. * omp.h.in: Add ompx_gnu_managed_mem_space and ompx_gnu_managed_mem_alloc enumerators, and gnu_managed_mem C++ allocator template. * omp_lib.f90.in: Add Fortran bindings for new allocator and memory space. * omp_lib.h.in: Likewise. * plugin/cuda-lib.def: Add cuMemAllocManaged. * plugin/plugin-nvptx.c (nvptx_alloc): Add managed parameter to support cuMemAllocManaged. (GOMP_OFFLOAD_alloc): Move contents to ... (cleanup_and_alloc): ... this new function, and add managed support. (GOMP_OFFLOAD_managed_alloc): New function. (GOMP_OFFLOAD_managed_free): New function. * target.c (gomp_managed_alloc): New function. (gomp_managed_free): New function. (gomp_load_plugin_for_device): Load optional managed_alloc and managed_free plugin APIs. * testsuite/lib/libgomp.exp: Add check_effective_target_omp_managedmem. * testsuite/libgomp.c++/alloc-managed-1.C: New test. * testsuite/libgomp.c/alloc-managed-1.c: New test. * testsuite/libgomp.c/alloc-managed-2.c: New test. * testsuite/libgomp.c/alloc-managed-3.c: New test. * testsuite/libgomp.c/alloc-managed-4.c: New test. * testsuite/libgomp.fortran/alloc-managed-1.f90: New test. Co-authored-by: Kwok Cheung Yeung <kcyeung@baylibre.com> Co-authored-by: Thomas Schwinge <tschwinge@baylibre.com> |
||
|
Tobias Burnus
|
2de6462c38 |
libgomp.{c-c++-common,fortran}/target-is-accessible-1.c: Fix testcases for omp_default_device [P119677]
Commit r16-5188-g5da963d988e8ea added omp_default_device such that -5 became a conforming device number, but the tests used them to test for as non-conforming number; now -6 is used. libgomp/ChangeLog: PR libgomp/119677 * testsuite/libgomp.c-c++-common/target-is-accessible-1.c: Modify test as -5 is now a conforming device number. * testsuite/libgomp.fortran/target-is-accessible-1.f90: Likewise. |
||
|
Tobias Burnus
|
5da963d988 |
OpenMP: Add omp_default_device named constant [PR119677]
OpenMP TR 14 (OpenMP 6.1) adds omp_default_device < -1 as named constant alongside omp_initial_device and omp_default_device. GCC supports it already internally via GOMP_DEVICE_DEFAULT_OMP_61, but this patch now adds the omp_default_device enum/PARAMETER to omp.h / omp_lib. Note that PR119677 requests some cleanups, which still have to be done. PR libgomp/119677 gcc/fortran/ChangeLog: * intrinsic.texi (OpenMP Modules): Add omp_default_device. * openmp.cc (gfc_resolve_omp_context_selector): Accept omp_default_device as conforming device number. libgomp/ChangeLog: * omp.h.in (omp_default_device): New enum value. * omp_lib.f90.in: New parameter. * omp_lib.h.in: Likewise * target.c (gomp_get_default_device): New. Split off from ... (resolve_device): ... here; call it. (omp_target_alloc, omp_target_free, omp_target_is_present, omp_target_memcpy_check, omp_target_memset, omp_target_memset_async, omp_target_associate_ptr, omp_get_mapped_ptr, omp_target_is_accessible, omp_pause_resource, omp_get_uid_from_device): Handle omp_default_device. * testsuite/libgomp.c/device_uid.c: Likewise. * testsuite/libgomp.fortran/device_uid.f90: Likewise. * testsuite/libgomp.c-c++-common/omp-default-device.c: New test. * testsuite/libgomp.fortran/omp-default-device.f90: New test. |
||
|
Tobias Burnus
|
28d20a591d |
libgomp.fortran/omp_target_memset.f90 - Avoid implicit mapping by an uninit size [PR122543]
In OpenMP, pointers are implicitly mapped - which means for Fortran that their pointer target is also mapped. However, for uninitialized memory, this means that some random pointee with some random amount of memory is copied - in the good case, size == 0, but if not, odd things can happen. Solution: Use 'fptr => null()' before the target mapping or - as done here - declare the pointer inside the region. libgomp/ChangeLog: PR libgomp/122543 * testsuite/libgomp.fortran/omp_target_memset.f90: Move fptr inside the target to avoid implicit mapping of its uninit pointee. * testsuite/libgomp.fortran/omp_target_memset-2.f90: Likewise. |
||
|
Tobias Burnus
|
4e47e2f833 |
libgomp: Add OpenMP's omp_target_memset/omp_target_memset_async
PR libgomp/120444
include/ChangeLog:
* cuda/cuda.h (cuMemsetD8, cuMemsetD8Async): Declare.
libgomp/ChangeLog:
* libgomp-plugin.h (GOMP_OFFLOAD_memset): Declare.
* libgomp.h (struct gomp_device_descr): Add memset_func.
* libgomp.map (GOMP_6.0.1): Add omp_target_memset{,_async}.
* libgomp.texi (Device Memory Routines): Document them.
* omp.h.in (omp_target_memset, omp_target_memset_async): Declare.
* omp_lib.f90.in (omp_target_memset, omp_target_memset_async):
Add interfaces.
* omp_lib.h.in (omp_target_memset, omp_target_memset_async): Likewise.
* plugin/cuda-lib.def: Add cuMemsetD8.
* plugin/plugin-gcn.c (struct hsa_runtime_fn_info): Add
hsa_amd_memory_fill_fn.
(init_hsa_runtime_functions): DLSYM_OPT_FN load it.
(GOMP_OFFLOAD_memset): New.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_memset): New.
* target.c (omp_target_memset_int, omp_target_memset,
omp_target_memset_async_helper, omp_target_memset_async): New.
(gomp_load_plugin_for_device): Add DLSYM (memset).
* testsuite/libgomp.c-c++-common/omp_target_memset.c: New test.
* testsuite/libgomp.c-c++-common/omp_target_memset-2.c: New test.
* testsuite/libgomp.c-c++-common/omp_target_memset-3.c: New test.
* testsuite/libgomp.fortran/omp_target_memset.f90: New test.
* testsuite/libgomp.fortran/omp_target_memset-2.f90: New test.
|
||
|
Tobias Burnus
|
5b2e7afb13 |
libgomp.fortran/metadirective-1.f90: Expect 'error:' for nvptx compile [PR118694]
This should have been part of commit r16-838-gb3d07ec7ac2ccd or r16-883-g5d6ed6d604ff94 - all showing the same issue: '!$omp target' followed by a metadirective with 'teams'; if the metadirective cannot be early resolved, a diagnostic error is shown about using directives between 'target' and 'teams'. While the message is misleading, the problem is that the host invokes 'target' differently when 'teams' is present; in this case, host fallback + amdgcn offload require the no-teams case, nvptx offload the teams case such that it only can be resolved at runtime. Mark the error as 'dg-bogus + xfail' to silence the FAIL, when nvptx offloading is compiled for. (If not, the metadirective can be resolved early during compilation.) libgomp/ChangeLog: PR middle-end/118694 * testsuite/libgomp.fortran/metadirective-1.f90: xfail when compiling (also) for nvptx offloading as an error is then expected. |
||
|
Tobias Burnus
|
f99017c312 |
OpenMP/Fortran: Fix allocatable-component mapping of derived-type array comps
The check whether the location expression in map clause has allocatable components was failing for some derived-type array expressions such as map(var%tiles(1)) as the compiler produced _4 = var.tiles; MEMREF(_4, _5); This commit now also handles this case. gcc/fortran/ChangeLog: * trans-openmp.cc (gfc_omp_deep_mapping_do): Handle SSA_NAME if a def_stmt is available. libgomp/ChangeLog: * testsuite/libgomp.fortran/alloc-comp-4.f90: New test. |
||
|
Tobias Burnus
|
94e6341047 |
libgomp.{c,fortran}/interop-{hip,cuda}: Fix dg-run target selection
While the tests checked whether the CUDA/HIP runtime is available
before processing them, the execution was then done unconditionally,
leading to FAIL when the default device was the host (or the wrong
offload device).
Now the test is only executed ('run') when the default device is an
Nvidia or AMD GPU (depending on the test case, cf. the test file name).
Otherwise, only a 'link' test is done. (Except when the effective-target
check cannot find the runtime lib - then the test is skipped [as before].)
Note: The cublas/hipblas tests use variant functions and iterate over
all devices, such that the cublas or hipblas, respectively, is only
called when the active device is an AMD or Nvidia device, respectively,
while for the host and other device types the fallback is called.
libgomp/ChangeLog:
* testsuite/libgomp.c/interop-cuda-full.c: Use 'link' instead
of 'run' when the default device is "! offload_device_nvptx".
* testsuite/libgomp.c/interop-cuda-libonly.c: Likewise.
* testsuite/libgomp.c/interop-hip-nvidia-full.c: Likewise.
* testsuite/libgomp.c/interop-hip-nvidia-no-headers.c: Likewise.
* testsuite/libgomp.c/interop-hip-nvidia-no-hip-header.c: Likewise.
* testsuite/libgomp.fortran/interop-hip-nvidia-full.F90: Likewise.
* testsuite/libgomp.fortran/interop-hip-nvidia-no-module.F90: Likewise.
* testsuite/libgomp.c/interop-hip-amd-full.c: Use 'link' instead
of 'run' when the default device is "! offload_device_gcn".
* testsuite/libgomp.c/interop-hip-amd-no-hip-header.c: Likewise.
* testsuite/libgomp.fortran/interop-hip-amd-full.F90: Likewise.
* testsuite/libgomp.fortran/interop-hip-amd-no-module.F90: Likewise.
|
||
|
Tobias Burnus
|
9565076f9b |
libgomp.fortran/map-alloc-comp-9{,-usm}.f90: Add unified_shared_memory variant
When host memory is device accessible - independent whether mapping is done or not (i.e. self map), the 'vtab' pointer becomes accessible, which stores the dynamic type's type and size information. In principle, we want to test: USM available but mapping is still done, but as there is no simple + reliable not-crashing way to test for this, those checks are skipped in the (pre)existing test file map-alloc-comp-9.f90. Or rather: those are only active with self-maps, which is currently only true for the host. This commit adds map-alloc-comp-9-usm.f90 which runs the same test with 'omp requires unified_shared_memory'. While OpenMP permits both actual mapping and self maps with this flag, it in theory covers the missing cases. However, currently, GCC always uses self maps with USM. Still, having a device-run self-maps check is better than nothing, even if it misses the most interesting case. libgomp/ChangeLog: * testsuite/libgomp.fortran/map-alloc-comp-9.f90: Process differently when USE_USM_REQUIREMENT is set. * testsuite/libgomp.fortran/map-alloc-comp-9-usm.f90: New test. |
||
|
Tobias Burnus
|
08ce1b9f67 |
OpenMP: Restore lost Fortran testcase for 'omp allocate'
This testcase, which is present on the OG13 and OG14 branches, was
overlooked when the Fortran support for 'omp allocate' was added to
mainline (commit
|
||
|
Tobias Burnus
|
515d9be794 |
libgomp: Add additional OpenMP interop runtime tests
Add checks for nowait/depend and for checks that the returned CUDA, CUDA_DRIVER and HIP interop objects actually work. While the CUDA/CUDA_DRIVER ones are only for Nvidia GPUs, HIP works on both AMD and Nvidia GPUs; on Nvidia GPUs, it is a very thin wrapper around CUDA. For Fortran, only a HIP test has been added - using hipfort. While libgomp.c-c++-common/interop-2.c always works - even without GPU - and checks for depend / nowait, all others require that runtime libraries are found at link (and execution) time: For Nvidia GPUs, libcuda + libcudart or libcublas, For AMD GPUs, libamdhip64 or libhipblas. The header files and hipfort modules do not need to be present as a fallback has been implemented, but if they are, they get used. Due to the combinations, the basic 1x C/C++, 4x C and 1x Fortran tests yield 1x C/C++, 14x C and 4 Fortran run-test files. libgomp/ChangeLog: * testsuite/lib/libgomp.exp (check_effective_target_openacc_cublas, check_effective_target_openacc_cudart): Update description as the check requires more. (check_effective_target_openacc_libcuda, check_effective_target_openacc_libcublas, check_effective_target_openacc_libcudart, check_effective_target_gomp_hip_header_amd, check_effective_target_gomp_hip_header_nvidia, check_effective_target_gomp_hipfort_module, check_effective_target_gomp_libamdhip64, check_effective_target_gomp_libhipblas): New. * testsuite/libgomp.c-c++-common/interop-2.c: New test. * testsuite/libgomp.c/interop-cublas-full.c: New test. * testsuite/libgomp.c/interop-cublas-libonly.c: New test. * testsuite/libgomp.c/interop-cuda-full.c: New test. * testsuite/libgomp.c/interop-cuda-libonly.c: New test. * testsuite/libgomp.c/interop-hip-amd-full.c: New test. * testsuite/libgomp.c/interop-hip-amd-no-hip-header.c: New test. * testsuite/libgomp.c/interop-hip-nvidia-full.c: New test. * testsuite/libgomp.c/interop-hip-nvidia-no-headers.c: New test. * testsuite/libgomp.c/interop-hip-nvidia-no-hip-header.c: New test. * testsuite/libgomp.c/interop-hip.h: New test. * testsuite/libgomp.c/interop-hipblas-amd-full.c: New test. * testsuite/libgomp.c/interop-hipblas-amd-no-hip-header.c: New test. * testsuite/libgomp.c/interop-hipblas-nvidia-full.c: New test. * testsuite/libgomp.c/interop-hipblas-nvidia-no-headers.c: New test. * testsuite/libgomp.c/interop-hipblas-nvidia-no-hip-header.c: New test. * testsuite/libgomp.c/interop-hipblas.h: New test. * testsuite/libgomp.fortran/interop-hip-amd-full.F90: New test. * testsuite/libgomp.fortran/interop-hip-amd-no-module.F90: New test. * testsuite/libgomp.fortran/interop-hip-nvidia-full.F90: New test. * testsuite/libgomp.fortran/interop-hip-nvidia-no-module.F90: New test. * testsuite/libgomp.fortran/interop-hip.h: New test. |
||
|
Tobias Burnus
|
c9a8f2f9d3 |
OpenMP: Add libgomp.fortran/target-enter-data-8.f90
Add another testcase for Fortran deep mapping of allocatable components. libgomp/ChangeLog: * testsuite/libgomp.fortran/target-enter-data-8.f90: New test. |
||
|
Tobias Burnus
|
99cd28c473 |
Fortran/OpenMP: Support automatic mapping allocatable components (deep mapping)
When mapping an allocatable variable (or derived-type component), explicitly or implicitly, all its allocated allocatable components will automatically be mapped. The patch implements the target hooks, added for this feature to omp-low.cc with commit r15-3895-ge4a58b6f28383c. Namely, there is a check whether there are allocatable components at all: gfc_omp_deep_mapping_p. Then gfc_omp_deep_mapping_cnt, counting the number of required mappings; this is a dynamic value as it depends on array bounds and whether an allocatable is allocated or not. And, finally, the actual mapping: gfc_omp_deep_mapping. Polymorphic variables are partially supported: the mapping of the _data component is fully supported, but only components of the declared type are processed for additional allocatables. Additionally, _vptr is not touched. This means that everything needing _vtab information requires unified shared memory; in particular, _size data is required when accessing elements of polymorphic arrays. However, for scalar arrays, accessing components of the declare type should work just fine. As polymorphic variables are not (really) supported and OpenMP 6 explicitly disallows them, there is now a warning (-Wopenmp) when they are encountered. Unlimited polymorphics are rejected (error). Additionally, PRIVATE and FIRSTPRIVATE are not quite supported for allocatable components, polymorphic components and as polymorphic variable. Thus, those are now rejected as well. gcc/fortran/ChangeLog: * f95-lang.cc (LANG_HOOKS_OMP_DEEP_MAPPING, LANG_HOOKS_OMP_DEEP_MAPPING_P, LANG_HOOKS_OMP_DEEP_MAPPING_CNT): Define. * openmp.cc (gfc_match_omp_clause_reduction): Fix location setting. (resolve_omp_clauses): Permit allocatable components, reject them and polymorphic variables in PRIVATE/FIRSTPRIVATE. * trans-decl.cc (add_clause): Set clause location. * trans-openmp.cc (gfc_has_alloc_comps): Add ptr_ok and shallow_alloc_only Boolean arguments. (gfc_omp_replace_alloc_by_to_mapping): New. (gfc_omp_private_outer_ref, gfc_walk_alloc_comps, gfc_omp_clause_default_ctor, gfc_omp_clause_copy_ctor, gfc_omp_clause_assign_op, gfc_omp_clause_dtor): Update call to it. (gfc_omp_finish_clause): Minor cleanups, improve location data, handle allocatable components. (gfc_omp_deep_mapping_map, gfc_omp_deep_mapping_item, gfc_omp_deep_mapping_comps, gfc_omp_gen_simple_loop, gfc_omp_get_array_size, gfc_omp_elmental_loop, gfc_omp_deep_map_kind_p, gfc_omp_deep_mapping_int_p, gfc_omp_deep_mapping_p, gfc_omp_deep_mapping_do, gfc_omp_deep_mapping_cnt, gfc_omp_deep_mapping): New. (gfc_trans_omp_array_section): Save array descriptor in case deep-mapping lang hook will need it. (gfc_trans_omp_clauses): Likewise; use better clause location data. * trans.h (gfc_omp_deep_mapping_p, gfc_omp_deep_mapping_cnt, gfc_omp_deep_mapping): Add function prototypes. libgomp/ChangeLog: * libgomp.texi (5.0 Impl. Status): Mark mapping alloc comps as 'Y'. * testsuite/libgomp.fortran/allocatable-comp.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-3.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-4.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-5.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-6.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-7.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-8.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-9.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/map-alloc-comp-1.f90: Remove dg-error. * gfortran.dg/gomp/polymorphic-mapping-2.f90: Update warn wording. * gfortran.dg/gomp/polymorphic-mapping.f90: Change expected diagnostic; some tests moved to ... * gfortran.dg/gomp/polymorphic-mapping-1.f90: ... here as new test. * gfortran.dg/gomp/polymorphic-mapping-3.f90: New test. * gfortran.dg/gomp/polymorphic-mapping-4.f90: New test. * gfortran.dg/gomp/polymorphic-mapping-5.f90: New test. |
||
|
Tobias Burnus
|
c264df142a |
libgomp.fortran/get-mapped-ptr-1.f90: Use -6 for non-conf dev number
This is a fix for the GOMP_interop commit r15-8654-g99e2906ae255fc that added GOMP_DEVICE_DEFAULT_OMP_61 alias omp_default_device, which is a conforming device number. But that test used -5 as check for a non-conforming device number. libgomp/ChangeLog: * testsuite/libgomp.fortran/get-mapped-ptr-1.f90: Use -6 not -5 as non-conforming device number. |
||
|
Tobias Burnus
|
2d5c1e5149 |
Move gfortran.dg/gomp/declare-variant-mod-1*.f90 to libgomp.fortran/ [PR115271]
The test is a supposed to be a compile-only test but as dg-additional-sources does not work with dg-compile (due to compiling with '-o'), dg-link had to be used; as the code actually compiles (no diagnostic error), the linker is actually invoked, which fails unless the system compiler by chance provides the required files. Solution: move the test to libgomp. PR fortran/115271 gcc/testsuite/ChangeLog: * gfortran.dg/gomp/declare-variant-mod-1-use.f90: Move to libgomp/testsuite/libgomp.fortran/. * gfortran.dg/gomp/declare-variant-mod-1.f90: Likewise. libgomp/ChangeLog: * testsuite/libgomp.fortran/declare-variant-mod-1-use.f90: Moved from gcc/testsuite/gfortran.dg/gomp/. * testsuite/libgomp.fortran/declare-variant-mod-1.f90: Likewise. |
||
|
Sandra Loosemore
|
8fbccdb342 |
OpenMP: Fortran support for metadirectives and dynamic selectors
gcc/fortran/ChangeLog PR middle-end/112779 PR middle-end/113904 * decl.cc (gfc_match_end): Handle COMP_OMP_BEGIN_METADIRECTIVE and COMP_OMP_METADIRECTIVE. * dump-parse-tree.cc (show_omp_node): Handle EXEC_OMP_METADIRECTIVE. (show_code_node): Likewise. * gfortran.h (enum gfc_statement): Add ST_OMP_METADIRECTIVE, ST_OMP_BEGIN_METADIRECTIVE, and ST_OMP_END_METADIRECTIVE. (struct gfc_omp_clauses): Rename target_first_st_is_teams to target_first_st_is_teams_or_meta. (struct gfc_omp_variant): New. (gfc_get_omp_variant): New. (struct gfc_st_label): Add omp_region field. (enum gfc_exec_op): Add EXEC_OMP_METADIRECTIVE. (struct gfc_code): Add omp_variants fields. (gfc_free_omp_variants): Declare. (match_omp_directive): Declare. (is_omp_declarative_stmt): Declare. * io.cc (format_asterisk): Adjust initializer. * match.h (gfc_match_omp_begin_metadirective): Declare. (gfc_match_omp_metadirective): Declare. * openmp.cc (gfc_omp_directives): Uncomment metadirective. (gfc_match_omp_eos): Adjust to match context selectors. (gfc_free_omp_variants): New. (gfc_match_omp_clauses): Remove context_selector parameter and adjust to use gfc_match_omp_eos instead. (match_omp): Adjust call to gfc_match_omp_clauses. (gfc_match_omp_context_selector): Add metadirective_p parameter and adjust error-checking. Adjust matching of simd clauses. (gfc_match_omp_context_selector_specification): Adjust parameters so it can be used for metadirective as well as declare variant. (match_omp_metadirective): New. (gfc_match_omp_begin_metadirective): New. (gfc_match_omp_metadirective): New. (resolve_omp_metadirective): New. (resolve_omp_target): Handle metadirectives. (gfc_resolve_omp_directive): Handle EXEC_OMP_METADIRECTIVE. * parse.cc (gfc_matching_omp_context_selector): New. (gfc_in_omp_metadirective_body): New. (gfc_omp_region_count): New. (decode_omp_directive): Handle ST_OMP_BEGIN_METADIRECTIVE and ST_OMP_METADIRECTIVE. (match_omp_directive): New. (case_omp_structured_block): Define. (case_omp_do): Define. (gfc_ascii_statement): Handle ST_OMP_BEGIN_METADIRECTIVE, ST_OMP_END_METADIRECTIVE, and ST_OMP_METADIRECTIVE. (accept_statement): Handle ST_OMP_METADIRECTIVE and ST_OMP_BEGIN_METADIRECTIVE. (gfc_omp_end_stmt): New, split from... (parse_omp_do): ...here, and... (parse_omp_structured_block): ...here. Handle metadirectives, plus "allocate", "atomic", and "dispatch" which were missing. (parse_omp_oacc_atomic): Handle "end metadirective". (parse_openmp_allocate_block): Likewise. (parse_omp_dispatch): Likewise. (parse_omp_metadirective_body): New. (parse_executable): Handle metadirective. Use new case macros defined above. (gfc_parse_file): Initialize metadirective state. (is_omp_declarative_stmt): New. * parse.h (enum gfc_compile_state): Add COMP_OMP_METADIRECTIVE and COMP_OMP_BEGIN_METADIRECTIVE. (gfc_omp_end_stmt): Declare. (gfc_matching_omp_context_selector): Declare. (gfc_in_omp_metadirective_body): Declare. (gfc_omp_metadirective_region_count): Declare. * resolve.cc (gfc_resolve_code): Handle EXEC_OMP_METADIRECTIVE. * st.cc (gfc_free_statement): Likewise. * symbol.cc (compare_st_labels): Handle labels within a metadirective body. (gfc_get_st_label): Likewise. * trans-decl.cc (gfc_get_label_decl): Encode the metadirective region in the label_name. * trans-openmp.cc (gfc_trans_omp_directive): Handle EXEC_OMP_METADIRECTIVE. (gfc_trans_omp_set_selector): New, split/adapted from code.... (gfc_trans_omp_declare_variant): ...here. (gfc_trans_omp_metadirective): New. * trans-stmt.h (gfc_trans_omp_metadirective): Declare. * trans.cc (trans_code): Handle EXEC_OMP_METADIRECTIVE. gcc/testsuite/ChangeLog PR middle-end/112779 PR middle-end/113904 * gfortran.dg/gomp/metadirective-1.f90: New. * gfortran.dg/gomp/metadirective-10.f90: New. * gfortran.dg/gomp/metadirective-11.f90: New. * gfortran.dg/gomp/metadirective-12.f90: New. * gfortran.dg/gomp/metadirective-13.f90: New. * gfortran.dg/gomp/metadirective-2.f90: New. * gfortran.dg/gomp/metadirective-3.f90: New. * gfortran.dg/gomp/metadirective-4.f90: New. * gfortran.dg/gomp/metadirective-5.f90: New. * gfortran.dg/gomp/metadirective-6.f90: New. * gfortran.dg/gomp/metadirective-7.f90: New. * gfortran.dg/gomp/metadirective-8.f90: New. * gfortran.dg/gomp/metadirective-9.f90: New. * gfortran.dg/gomp/metadirective-construct.f90: New. * gfortran.dg/gomp/metadirective-no-score.f90: New. * gfortran.dg/gomp/pure-1.f90 (func_metadirective): New. (func_metadirective_2): New. (func_metadirective_3): New. * gfortran.dg/gomp/pure-2.f90 (func_metadirective): Delete. libgomp/ChangeLog PR middle-end/112779 PR middle-end/113904 * testsuite/libgomp.fortran/metadirective-1.f90: New. * testsuite/libgomp.fortran/metadirective-2.f90: New. * testsuite/libgomp.fortran/metadirective-3.f90: New. * testsuite/libgomp.fortran/metadirective-4.f90: New. * testsuite/libgomp.fortran/metadirective-5.f90: New. * testsuite/libgomp.fortran/metadirective-6.f90: New. Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com> Co-Authored-By: Sandra Loosemore <sandra@codesourcery.com> Co-Authored-By: Tobias Burnus <tobias@codesourcery.com> Co-Authored-By: Paul-Antoine Arras <pa@codesourcery.com> |
||
|
Paul-Antoine Arras
|
655a8a024d |
Add missing target directive in OpenMP dispatch Fortran runtime test
Without the target directive, the test would run on the host but still try to use device pointers, which causes a segfault. libgomp/ChangeLog: * testsuite/libgomp.fortran/dispatch-1.f90: Add missing target directive. |
||
|
Paul-Antoine Arras
|
bca8b13bd7 |
OpenMP: Fortran front-end support for dispatch + adjust_args
This patch adds support for the `dispatch` construct and the `adjust_args` clause to the Fortran front-end. Handling of `adjust_args` across translation units is missing due to PR115271. Minor modifications to the C++ FE and the ME are also folded into this patch as a side effect of the Fortran work. gcc/c-family/ChangeLog: * c-attribs.cc: (c_common_gnu_attributes): Rename "omp declare variant variant adjust_args" into "omp declare variant variant args" to also accommodate append_args. gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_dispatch): Handle INDIRECT_REF. gcc/fortran/ChangeLog: * dump-parse-tree.cc (show_omp_clauses): Handle novariants and nocontext clauses. (show_omp_node): Handle EXEC_OMP_DISPATCH. (show_code_node): Likewise. * frontend-passes.cc (gfc_code_walker): Handle novariants and nocontext. * gfortran.h (enum gfc_statement): Add ST_OMP_DISPATCH. (symbol_attribute): Add omp_declare_variant_need_device_ptr. (gfc_omp_clauses): Add novariants and nocontext. (gfc_omp_declare_variant): Add need_device_ptr_arg_list. (enum gfc_exec_op): Add EXEC_OMP_DISPATCH. * match.h (gfc_match_omp_dispatch): Declare. * openmp.cc (gfc_free_omp_clauses): Free novariants and nocontext clauses. (gfc_free_omp_declare_variant_list): Free need_device_ptr_arg_list namelist. (enum omp_mask2): Add OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (gfc_match_omp_clauses): Handle OMP_CLAUSE_NOVARIANTS and OMP_CLAUSE_NOCONTEXT. (OMP_DISPATCH_CLAUSES): Define. (gfc_match_omp_dispatch): New function. (gfc_match_omp_declare_variant): Parse adjust_args. (resolve_omp_clauses): Handle adjust_args, novariants and nocontext. Adjust handling of OMP_LIST_IS_DEVICE_PTR. (icode_code_error_callback): Handle EXEC_OMP_DISPATCH. (omp_code_to_statement): Likewise. (resolve_omp_dispatch): New function. (gfc_resolve_omp_directive): Handle EXEC_OMP_DISPATCH. * parse.cc (decode_omp_directive): Match dispatch. (next_statement): Handle ST_OMP_DISPATCH. (gfc_ascii_statement): Likewise. (parse_omp_dispatch): New function. (parse_executable): Handle ST_OMP_DISPATCH. * resolve.cc (gfc_resolve_blocks): Handle EXEC_OMP_DISPATCH. * st.cc (gfc_free_statement): Likewise. * trans-decl.cc (create_function_arglist): Declare. (gfc_get_extern_function_decl): Call it. * trans-openmp.cc (gfc_trans_omp_clauses): Handle novariants and nocontext. (replace_omp_dispatch_call): New function. (gfc_trans_omp_dispatch): New function. (gfc_trans_omp_directive): Handle EXEC_OMP_DISPATCH. (gfc_trans_omp_declare_variant): Handle adjust_args. * trans.cc (trans_code): Handle EXEC_OMP_DISPATCH:. gcc/ChangeLog: * gimplify.cc (gimplify_call_expr): Fix handling of need_device_ptr for type(c_ptr). Fix handling of nested function calls in a dispatch region. (find_ifn_gomp_dispatch): Return the IFN without stripping it. (gimplify_omp_dispatch): Keep IFN_GOMP_DISPATCH until gimplify_call_expr. libgomp/ChangeLog: * testsuite/libgomp.fortran/declare-variant-2-aux.f90: New test. * testsuite/libgomp.fortran/declare-variant-2.f90: New test (xfail). * testsuite/libgomp.fortran/dispatch-1.f90: New test. * testsuite/libgomp.fortran/dispatch-2.f90: New test. * testsuite/libgomp.fortran/dispatch-3.f90: New test. gcc/testsuite/ChangeLog: * g++.dg/gomp/dispatch-3.C: Update scan dumps. * gfortran.dg/gomp/declare-variant-2.f90: Update dg-error. * gfortran.dg/gomp/adjust-args-1.f90: New test. * gfortran.dg/gomp/adjust-args-2.f90: New test. * gfortran.dg/gomp/adjust-args-2a.f90: New test. * gfortran.dg/gomp/adjust-args-3.f90: New test. * gfortran.dg/gomp/adjust-args-4.f90: New test. * gfortran.dg/gomp/adjust-args-5.f90: New test. * gfortran.dg/gomp/adjust-args-6.f90: New test. * gfortran.dg/gomp/adjust-args-7.f90: New test. * gfortran.dg/gomp/adjust-args-8.f90: New test. * gfortran.dg/gomp/adjust-args-9.f90: New test. * gfortran.dg/gomp/dispatch-1.f90: New test. * gfortran.dg/gomp/dispatch-2.f90: New test. * gfortran.dg/gomp/dispatch-3.f90: New test. * gfortran.dg/gomp/dispatch-4.f90: New test. * gfortran.dg/gomp/dispatch-5.f90: New test. * gfortran.dg/gomp/dispatch-6.f90: New test. * gfortran.dg/gomp/dispatch-7.f90: New test. * gfortran.dg/gomp/dispatch-8.f90: New test. * gfortran.dg/gomp/dispatch-9.f90: New test. * gfortran.dg/gomp/dispatch-9a.f90: New test. * gfortran.dg/gomp/dispatch-10.f90: New test. |
||
|
Tobias Burnus
|
3269a722b7 |
Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' __builtin_is_initial_device
It turned out that 'if (omp_is_initial_device() .eqv. true)' gave an ICE due to comparing 'int' with 'logical(4)'. When digging deeper, it also turned out that when the procedure pointer is needed, the builtin cannot be used, either. (Follow up to r15-2799-gf1bfba3a9b3f31 ) Extend the code to also use the builtin acc_on_device with OpenACC, which was previously only used in C/C++. Additionally, fix folding when offloading is not enabled. Fixes additionally the BT_BOOL data type, which was 'char'/integer(1) instead of bool, backing the booleaness; use bool_type_node as the rest of GCC. gcc/fortran/ChangeLog: * gfortran.h (gfc_option_t): Add disable_acc_on_device. * options.cc (gfc_handle_option): Handle -fno-builtin-acc_on_device. * trans-decl.cc (gfc_get_extern_function_decl): Move __builtin_omp_is_initial_device handling to ... * trans-expr.cc (get_builtin_fn): ... this new function. (conv_function_val): Call it. (update_builtin_function): New. (gfc_conv_procedure_call): Call it. * types.def (BT_BOOL): Fix type by using bool_type_node. gcc/ChangeLog: * gimple-fold.cc (gimple_fold_builtin_acc_on_device): Also fold when offloading is not configured. libgomp/ChangeLog: * libgomp.texi (TR13): Fix minor typos. (omp_is_initial_device): Improve wording. (acc_on_device): Note how to disable the builtin. * testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Remove TODO. * testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise. Add -fno-builtin-acc_on_device. * testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: Update dg- as !offloading_enabled now compile-time expands acc_on_device. * testsuite/libgomp.fortran/target-is-initial-device-3.f90: New test. * testsuite/libgomp.oacc-fortran/acc_on_device-2.f90: New test. |
||
|
Tobias Burnus
|
b95ad25f9c |
Move gfortran.dg/gomp/allocate-static.f90 to libgomp.fortran/
The testcase was turned into a 'dg-do run' check to check for the alignment, but this only works in testsuite/gfortran.dg, causing link errors for out-of-tree testing. The test was added in r15-4104-ga8caeaacf499d5. gcc/testsuite/: * gfortran.dg/gomp/allocate-static.f90: Move to libgomp/testsuite/. libgomp/: * testsuite/libgomp.fortran/allocate-static.f90: Moved from gcc/testsuite/ as it is a dg-do run test; use real omp_lib_kinds instead of local definition |
||
|
Tobias Burnus
|
b752eed3e3 |
OpenMP: Add support for 'self_maps' to the 'require' directive
'self_maps' implies 'unified_shared_memory', except that the latter also permits that explicit maps copy data to device memory while self_maps does not. In GCC, currently, both are handled identical. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_requires): Handle self_maps clause. gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_requires): Handle self_maps clause. gcc/fortran/ChangeLog: * gfortran.h (enum gfc_omp_requires_kind): Add OMP_REQ_SELF_MAPS. (gfc_namespace): Enlarge omp_requires bitfield. * module.cc (enum ab_attribute, attr_bits): Add AB_OMP_REQ_SELF_MAPS. (mio_symbol_attribute): Handle it. * openmp.cc (gfc_check_omp_requires, gfc_match_omp_requires): Handle self_maps clause. * parse.cc (gfc_parse_file): Handle self_maps clause. gcc/ChangeLog: * lto-cgraph.cc (output_offload_tables, omp_requires_to_name): Handle self_maps clause. * omp-general.cc (struct omp_ts_info, omp_context_selector_matches): Likewise for the associated trait. * omp-general.h (enum omp_requires): Add OMP_REQUIRES_SELF_MAPS. * omp-selectors.h (enum omp_ts_code): Add OMP_TRAIT_IMPLEMENTATION_SELF_MAPS. include/ChangeLog: * gomp-constants.h (GOMP_REQUIRES_SELF_MAPS): #define. libgomp/ChangeLog: * plugin/plugin-gcn.c (GOMP_OFFLOAD_get_num_devices): Accept self_maps clause. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_num_devices): Likewise. * libgomp.texi (TR13 Impl. Status): Set to 'Y'. * target.c (gomp_requires_to_name, GOMP_offload_register_ver, gomp_target_init): Handle self_maps clause. * testsuite/libgomp.fortran/self_maps.f90: New test. gcc/testsuite/ChangeLog: * c-c++-common/gomp/declare-variant-1.c: Add self_maps test. * c-c++-common/gomp/requires-4.c: Likewise. * gfortran.dg/gomp/declare-variant-3.f90: Likewise. * c-c++-common/gomp/requires-2.c: Update dg-error msg. * gfortran.dg/gomp/requires-2.f90: Likewise. * gfortran.dg/gomp/requires-self-maps-aux.f90: New. * gfortran.dg/gomp/requires-self-maps.f90: New. |
||
|
Tobias Burnus
|
cdb9aa0f62 |
OpenMP: Fix omp_get_device_from_uid, minor cleanup
In Fortran, omp_get_device_from_uid can also accept substrings, which are then not NUL terminated. Fixed by introducing a fortran.c wrapper function. Additionally, in case of a fail the plugin functions now return NULL instead of failing fatally such that a fall-back UID is generated. gcc/ChangeLog: * omp-general.cc (omp_runtime_api_procname): Strip "omp_" from string; move get_device_from_uid as now a '_' suffix exists. libgomp/ChangeLog: * fortran.c (omp_get_device_from_uid_): New function. * libgomp.map (GOMP_6.0): Add it. * oacc-host.c (host_dispatch): Init '.uid' and '.get_uid_func'. * omp_lib.f90.in: Make it used by removing bind(C). * omp_lib.h.in: Likewise. * target.c (omp_get_device_from_uid): Ensure the device is initialized. * plugin/plugin-gcn.c (GOMP_OFFLOAD_get_uid): Add function comment; return NULL in case of an error. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_uid): Likewise. * testsuite/libgomp.fortran/device_uid.f90: Update to test substrings. |
||
|
Tobias Burnus
|
bf4a5efa80 |
OpenMP: Add get_device_from_uid/omp_get_uid_from_device routines
Those TR13/OpenMP 6.0 routines permit a reproducible offloading to a specific device by mapping an OpenMP device number to a unique ID (UID). The GPU device UIDs should be universally unique, the one for the host is not. gcc/ChangeLog: * omp-general.cc (omp_runtime_api_procname): Add get_device_from_uid and omp_get_uid_from_device routines. include/ChangeLog: * cuda/cuda.h (cuDeviceGetUuid): Declare. (cuDeviceGetUuid_v2): Add prototype. libgomp/ChangeLog: * config/gcn/target.c (omp_get_uid_from_device, omp_get_device_from_uid): Add stub implementation. * config/nvptx/target.c (omp_get_uid_from_device, omp_get_device_from_uid): Likewise. * fortran.c (omp_get_uid_from_device_, omp_get_uid_from_device_8_): New functions. * libgomp-plugin.h (GOMP_OFFLOAD_get_uid): Add prototype. * libgomp.h (struct gomp_device_descr): Add 'uid' and 'get_uid_func'. * libgomp.map (GOMP_6.0): New, includind the new UID routines. * libgomp.texi (OpenMP Technical Report 13): Mark UID routines as 'Y'. (Device Information Routines): Document new UID routines. (Offload-Target Specifics): Document UID format. * omp.h.in (omp_get_device_from_uid, omp_get_uid_from_device): New prototype. * omp_lib.f90.in (omp_get_device_from_uid, omp_get_uid_from_device): New interface. * omp_lib.h.in: Likewise. * plugin/cuda-lib.def: Add cuDeviceGetUuid and cuDeviceGetUuid_v2 via CUDA_ONE_CALL_MAYBE_NULL. * plugin/plugin-gcn.c (GOMP_OFFLOAD_get_uid): New. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_uid): New. * target.c (str_omp_initial_device): New static var. (STR_OMP_DEV_PREFIX): Define. (gomp_get_uid_for_device, omp_get_uid_from_device, omp_get_device_from_uid): New. (gomp_load_plugin_for_device): DLSYM_OPT the function 'get_uid'. (gomp_target_init): Set the device's 'uid' field to NULL. * testsuite/libgomp.c/device_uid.c: New test. * testsuite/libgomp.fortran/device_uid.f90: New test. |
||
|
Tobias Burnus
|
0beac1db38 |
libgomp: Add interop types and routines to OpenMP's headers and module
This commit adds OpenMP 5.1+'s interop enumeration, type and routine declarations to the C/C++ header file and, new in OpenMP TR13, also to the Fortran module and omp_lib.h header file. While a stub implementation is provided, only with foreign runtime support by the libgomp GPU plugins and with the 'interop' directive, this becomes really useful. libgomp/ChangeLog: * fortran.c (omp_get_interop_str_, omp_get_interop_name_, omp_get_interop_type_desc_, omp_get_interop_rc_desc_): Add. * libgomp.map (GOMP_5.1.3): New; add interop routines. * omp.h.in: Add interop typedefs, enum and prototypes. (__GOMP_DEFAULT_NULL): Define. (omp_target_memcpy_async, omp_target_memcpy_rect_async): Use it for the optional depend argument. * omp_lib.f90.in: Add paramters and interfaces for interop. * omp_lib.h.in: Likewise; move F90 '&' to column 81 for -ffree-length-80. * target.c (omp_get_num_interop_properties, omp_get_interop_int, omp_get_interop_ptr, omp_get_interop_str, omp_get_interop_name, omp_get_interop_type_desc, omp_get_interop_rc_desc): Add. * config/gcn/target.c (omp_get_num_interop_properties, omp_get_interop_int, omp_get_interop_ptr, omp_get_interop_str, omp_get_interop_name, omp_get_interop_type_desc, omp_get_interop_rc_desc): Add. * config/nvptx/target.c (omp_get_num_interop_properties, omp_get_interop_int, omp_get_interop_ptr, omp_get_interop_str, omp_get_interop_name, omp_get_interop_type_desc, omp_get_interop_rc_desc): Add. * testsuite/libgomp.c-c++-common/interop-routines-1.c: New test. * testsuite/libgomp.c-c++-common/interop-routines-2.c: New test. * testsuite/libgomp.fortran/interop-routines-1.F90: New test. * testsuite/libgomp.fortran/interop-routines-2.F90: New test. * testsuite/libgomp.fortran/interop-routines-3.F: New test. * testsuite/libgomp.fortran/interop-routines-4.F: New test. * testsuite/libgomp.fortran/interop-routines-5.F: New test. * testsuite/libgomp.fortran/interop-routines-6.F: New test. * testsuite/libgomp.fortran/interop-routines-7.F90: New test. |
||
|
Thomas Schwinge
|
9f5d22e3e2 |
OpenMP: Constructors and destructors for "declare target" static aggregates: Fix effective-target keyword in test cases
(Most of) the tests added in commit
|
||
|
Tobias Burnus
|
f1bfba3a9b |
OpenMP: Constructors and destructors for "declare target" static aggregates
This commit also compile-time expands (__builtin_)omp_is_initial_device for both Fortran and C/C++ (unless, -fno-builtin-omp_is_initial_device is used). But the main change is: This commit adds support for running constructors and destructors for static (file-scope) aggregates for C++ objects which are marked with "declare target" directives on OpenMP offload targets. Before this commit, space is allocated on the target for such aggregates, but nothing ever constructs them properly, so they end up zero-initialised. (See the new test static-aggr-constructor-destructor-3.C for a reason why running constructors on the target is preferable to e.g. constructing on the host and then copying the resulting object to the target.) 2024-08-07 Julian Brown <julian@codesourcery.com> Tobias Burnus <tobias@baylibre.com> gcc/ChangeLog: * builtins.def (DEF_GOMP_BUILTIN_COMPILER): Define DEF_GOMP_BUILTIN_COMPILER to handle the non-prefix version. * gimple-fold.cc (gimple_fold_builtin_omp_is_initial_device): New. (gimple_fold_builtin): Call it. * omp-builtins.def (BUILT_IN_OMP_IS_INITIAL_DEVICE): Define. * tree.cc (get_file_function_name): Support names for on-target constructor/destructor functions. gcc/cp/ * decl2.cc (tree-inline.h): Include. (static_init_fini_fns): Bump to four entries. Update comment. (start_objects, start_partial_init_fini_fn): Add 'omp_target' parameter. Support "declare target" decls. Update forward declaration. (emit_partial_init_fini_fn): Add 'host_fn' parameter. Return tree for the created function. Support "declare target". (OMP_SSDF_IDENTIFIER): New macro. (partition_vars_for_init_fini): Support partitioning "declare target" variables also. (generate_ctor_or_dtor_function): Add 'omp_target' parameter. Support "declare target" decls. (c_parse_final_cleanups): Support constructors/destructors on OpenMP offload targets. gcc/fortran/ChangeLog: * gfortran.h (gfc_option_t): Add disable_omp_is_initial_device. * lang.opt (fbuiltin-): Add. * options.cc (gfc_handle_option): Handle -fno-builtin-omp_is_initial_device. * f95-lang.cc (gfc_init_builtin_functions): Handle DEF_GOMP_BUILTIN_COMPILER. * trans-decl.cc (gfc_get_extern_function_decl): Add code to use DEF_GOMP_BUILTIN_COMPILER for 'omp_is_initial_device'. libgomp/ChangeLog: * testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C: New test. * testsuite/libgomp.c-c++-common/target-is-initial-host.c: New test. * testsuite/libgomp.c-c++-common/target-is-initial-host-2.c: New test. * testsuite/libgomp.fortran/target-is-initial-host.f: New test. * testsuite/libgomp.fortran/target-is-initial-host.f90: New test. * testsuite/libgomp.fortran/target-is-initial-host-2.f90: New test. Co-authored-by: Tobias Burnus <tobias@baylibre.com> |
||
|
Tobias Burnus
|
c99cdcab4f |
omp-offload.cc: Fix value-expr handling of 'declare target link' vars [PR115637]
As the PR and included testcase shows, replacing 'arr2' by its value expression
'*arr2$13$linkptr' failed for
MEM <uint128_t> [(c_char * {ref-all})&arr2]
which left 'arr2' in the code as unknown symbol. Now expand the value expression
already in pass_omp_target_link::execute's process_link_var_op walk_gimple_stmt
walk - and don't rely on gimple_regimplify_operands.
PR middle-end/115637
gcc/ChangeLog:
* gimplify.cc (gimplify_body): Fix macro name in the comment.
* omp-offload.cc (find_link_var_op): Rename to ...
(process_link_var_op): ... this. Replace value expr.
(pass_omp_target_link::execute): Update walk_gimple_stmt call.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/declare-target-link.f90: Uncomment
now working code.
Co-authored-by: Richard Biener <rguenther@suse.de
|
||
|
Tobias Burnus
|
29b1587e7d |
OpenMP/Fortran: Fix handling of 'declare target' with 'link' clause [PR115559]
Contrary to a normal 'declare target', the 'declare target link' attribute also needs to set node->offloadable and push the offload_vars in the front end. Linked variables require that the data is mapped. For module variables, this can happen anywhere. For variables in an external subprograms or the main programm, this can only happen in the either that program itself or in an internal subprogram. - Whether a variable is just normally mapped or linked then becomes relevant if a device routine exists that can access that variable, i.e. an internal procedure has then to be marked as declare target. PR fortran/115559 gcc/fortran/ChangeLog: * trans-common.cc (build_common_decl): Add 'omp declare target' and 'omp declare target link' variables to offload_vars. * trans-decl.cc (add_attributes_to_decl): Likewise; update args and call decl_attributes. (get_proc_pointer_decl, gfc_get_extern_function_decl, build_function_decl): Update calls. (gfc_get_symbol_decl): Likewise; move after 'DECL_STATIC (t)=1' to avoid errors with symtab_node::get_create. libgomp/ChangeLog: * testsuite/libgomp.fortran/declare-target-link.f90: New test. |
||
|
Andrew Stubbs
|
64001441ec |
libgomp, openmp: Add ompx_gnu_pinned_mem_alloc
This creates a new predefined allocator as a shortcut for using pinned memory with OpenMP. This is not in the OpenMP standard so it uses the "ompx" namespace and an independent enum baseline of 200 (selected to not clash with other known implementations). The allocator is equivalent to using a custom allocator with the pinned trait and the null fallback trait. One motivation for having this feature is for use by the (planned) -foffload-memory=pinned feature. gcc/fortran/ChangeLog: * openmp.cc (is_predefined_allocator): Update valid ranges to incorporate ompx_gnu_pinned_mem_alloc. libgomp/ChangeLog: * allocator.c (ompx_gnu_min_predefined_alloc): New. (ompx_gnu_max_predefined_alloc): New. (predefined_alloc_mapping): Rename to ... (predefined_omp_alloc_mapping): ... this. (predefined_ompx_gnu_alloc_mapping): New. (_Static_assert): Adjust for the new name, and add a new assert for the new table. (predefined_allocator_p): New. (predefined_alloc_mapping): New. (omp_aligned_alloc): Support ompx_gnu_pinned_mem_alloc. Use predefined_allocator_p and predefined_alloc_mapping. (omp_free): Likewise. (omp_alligned_calloc): Likewise. (omp_realloc): Likewise. * env.c (parse_allocator): Add ompx_gnu_pinned_mem_alloc. * libgomp.texi: Document ompx_gnu_pinned_mem_alloc. * omp.h.in (omp_allocator_handle_t): Add ompx_gnu_pinned_mem_alloc. * omp_lib.f90.in: Add ompx_gnu_pinned_mem_alloc. * omp_lib.h.in: Add ompx_gnu_pinned_mem_alloc. * testsuite/libgomp.c/alloc-pinned-5.c: New test. * testsuite/libgomp.c/alloc-pinned-6.c: New test. * testsuite/libgomp.fortran/alloc-pinned-1.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/allocate-pinned-1.f90: New test. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com> |
||
|
Thomas Schwinge
|
3a4775d440 |
nvptx, libgfortran: Switch out of "minimal" mode
..., in order to enable (portions of) Fortran I/O, for example. libgfortran/ * configure.ac: No longer set 'LIBGFOR_MINIMAL' for nvptx. * configure: Regenerate. libgomp/ * libgomp.texi (nvptx): Update. * testsuite/libgomp.fortran/target-print-1-nvptx.f90: Remove. * testsuite/libgomp.fortran/target-print-1.f90: Adjust. * testsuite/libgomp.oacc-fortran/error_stop-2-nvptx.f: New. * testsuite/libgomp.oacc-fortran/error_stop-2.f: Adjust. * testsuite/libgomp.oacc-fortran/print-1-nvptx.f90: Adjust. * testsuite/libgomp.oacc-fortran/print-1.f90: Adjust. * testsuite/libgomp.oacc-fortran/stop-2-nvptx.f: New. * testsuite/libgomp.oacc-fortran/stop-2.f: Adjust. Co-authored-by: Andrew Stubbs <ams@gcc.gnu.org> |
||
|
Jakub Jelinek
|
804c0f35a6 |
openmp: OpenMP loop transformation support
This patch is largely rewritten version of the https://gcc.gnu.org/pipermail/gcc-patches/2023-October/631764.html patch set which I've promissed to adjust the way I'd like it but didn't get to it until now. The previous series together in diffstat was 176 files changed, 12107 insertions(+), 298 deletions(-) This patch is 197 files changed, 10843 insertions(+), 212 deletions(-) and diff between the old series and new patch is 268 files changed, 8053 insertions(+), 9231 deletions(-) Only the 5.1/5.2 tile/unroll constructs are supported, in various places some preparations for the other 6.0 loop transformations constructs (interchange/reverse/fuse) are done, but certainly not complete and not everywhere. The important difference is that because tile/unroll partial map 1:1 the original loops to generated canonical loops and add another set of generated loops without canonical form inside of it, the tile/unroll partial constructs are terminal for the generated loop, one can't have some loops from the tile or unroll partial and some further loops from inside the body of that construct. The GENERIC representation attempts to match what the standard specifies, so there are separate OMP_TILE and OMP_UNROLL trees. If for a particular loop in a loop nest of some OpenMP loop it awaits a generated loop from a nested loop, or if in OMP_LOOPXFORM_LOWERED OMP_TILE/UNROLL construct a generated loop has been moved to some surrounding construct, that particular loop is represented by all NULL_TREEs in the OMP_FOR_{INIT,COND,INCR,ORIG_DECLS} vector. The lowering of the loop transforming constructs is done at gimplification time, at the start of gimplify_omp_for. I think this way it is more maintainable over magic clauses with various loop depths on the other looping constructs or the magic OMP_LOOP_TRANS construct. Though, I admit I'm still undecided how to represent the OpenMP 6.0 loop transformation case of say: #pragma omp for collapse (4) for (int i = 0; i < 32; ++i) #pragma omp interchange permutation (2, 1) #pragma omp reverse for (int j = 0; j < 32; ++j) #pragma omp reverse for (int k = 0; k < 32; ++k) for (int l = 0; l < 32; ++l) ; Surely the i loop would go to first vector elements of OMP_FOR_* of the work-sharing loop, then 2 loops are expecting generated loops from interchange which would be inside of the body. But the innermost l loop isn't part of the interchange, so the question is where to put it. One possibility is to have it in the 4th loop of the OMP_FOR, another possibility would be to add some artificial construct inside of the OMP_INTERCHANGE and 2 OMP_REVERSE bodies which would contain the inner loop(s), e.g. it could be OMP_INTERCHANGE without permutation clause or some artificial ones or whatever. I've recently raised various unclear things in the 5.1/5.2/TRs versions regarding loop transformations, in particular https://github.com/OpenMP/spec/issues/3908 https://github.com/OpenMP/spec/issues/3909 (sorry, private links unless you have OpenMP membership). Until those are resolved, I have a sorry on trying to mix generated loops with non-rectangular loops (way too many questions need to be answered before that can be done) and similarly for mixing non-perfectly nested loops with generated loops (again, it can be implemented somehow, but is way too unclear). The second issue is mostly about data sharing, which is ambiguous, the patch makes the artificial iterators of the loops effectively private in the associated constructs (more like local), but for user iterators doesn't do anything in particular, so for now one needs to use explicit data sharing clauses on the non-loop transformation OpenMP looping constructs or surrounding parallel/task/target etc. 2024-06-05 Jakub Jelinek <jakub@redhat.com> Frederik Harwath <frederik@codesourcery.com> Sandra Loosemore <sandra@codesourcery.com> gcc/ * tree.def (OMP_TILE, OMP_UNROLL): New tree codes. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_PARTIAL, OMP_CLAUSE_FULL and OMP_CLAUSE_SIZES. * tree.h (OMP_LOOPXFORM_CHECK): Define. (OMP_LOOPXFORM_LOWERED): Define. (OMP_CLAUSE_PARTIAL_EXPR): Define. (OMP_CLAUSE_SIZES_LIST): Define. * tree.cc (omp_clause_num_ops, omp_clause_code_name): Add entries for OMP_CLAUSE_{PARTIAL,FULL,SIZES}. * tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_{PARTIAL,FULL,SIZES}. (dump_generic_node): Handle OMP_TILE and OMP_UNROLL. Skip printing loops with NULL OMP_FOR_INIT (node) vector element. * gimplify.cc (is_gimple_stmt): Handle OMP_TILE and OMP_UNROLL. (gimplify_omp_taskloop_expr): For SAVE_EXPR use gimplify_save_expr. (gimplify_omp_loop_xform): New function. (gimplify_omp_for): Call omp_maybe_apply_loop_xforms and if that reshuffles what the passed pointer points to, retry or return GS_OK. Handle OMP_TILE and OMP_UNROLL. (gimplify_omp_loop): Call omp_maybe_apply_loop_xforms and if that reshuffles what the passed pointer points to, return GS_OK. (gimplify_expr): Handle OMP_TILE and OMP_UNROLL. * omp-general.h (omp_loop_number_of_iterations, omp_maybe_apply_loop_xforms): Declare. * omp-general.cc (omp_adjust_for_condition): For LE_EXPR and GE_EXPR with pointers, don't add/subtract one, but the size of what the pointer points to. (omp_loop_number_of_iterations, omp_apply_tile, find_nested_loop_xform, omp_maybe_apply_loop_xforms): New functions. gcc/c-family/ * c-common.h (c_omp_find_generated_loop): Declare. * c-gimplify.cc (c_genericize_control_stmt): Handle OMP_TILE and OMP_UNROLL. * c-omp.cc (c_finish_omp_for): Handle generated loops. (c_omp_is_loop_iterator): Likewise. (c_find_nested_loop_xform_r, c_omp_find_generated_loop): New functions. (c_omp_check_loop_iv): Handle generated loops. For now sorry on mixing non-rectangular loop with generated loops. (c_omp_check_loop_binding_exprs): For now sorry on mixing imperfect loops with generated loops. (c_omp_directives): Uncomment tile and unroll entries. * c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_TILE and PRAGMA_OMP_UNROLL, change PRAGMA_OMP__LAST_ to the latter. (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_FULL and PRAGMA_OMP_CLAUSE_PARTIAL. * c-pragma.cc (omp_pragmas_simd): Add tile and unroll omp pragmas. gcc/c/ * c-parser.cc (c_parser_skip_std_attribute_spec_seq): New function. (check_omp_intervening_code): Reject imperfectly nested tile. (c_parser_compound_statement_nostart): If want_nested_loop, use c_parser_omp_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. (c_parser_omp_clause_name): Handle full and partial clause names. (c_parser_omp_clause_allocate): Remove spurious semicolon. (c_parser_omp_clause_full, c_parser_omp_clause_partial): New functions. (c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_FULL and PRAGMA_OMP_CLAUSE_PARTIAL. (c_parser_omp_next_tokens_can_be_canon_loop): New function. (c_parser_omp_loop_nest): Parse C23 attributes. Handle tile/unroll constructs. Use c_parser_omp_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. Only add_stmt (body) if it is non-NULL. (c_parser_omp_for_loop): Rename tiling variable to oacc_tiling. For OMP_CLAUSE_SIZES set collapse to list length of OMP_CLAUSE_SIZES_LIST. Use c_parser_omp_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. Remove spurious semicolon. Don't call c_omp_check_loop_binding_exprs if stmt is NULL. Skip generated loops. (c_parser_omp_tile_sizes, c_parser_omp_tile): New functions. (OMP_UNROLL_CLAUSE_MASK): Define. (c_parser_omp_unroll): New function. (c_parser_omp_construct): Handle PRAGMA_OMP_TILE and PRAGMA_OMP_UNROLL. * c-typeck.cc (c_finish_omp_clauses): Adjust wording of some of the conflicting clause diagnostic messages to include word clause. Handle OMP_CLAUSE_{FULL,PARTIAL,SIZES} and diagnose full vs. partial conflict. gcc/cp/ * cp-tree.h (dependent_omp_for_p): Add another tree argument. * parser.cc (check_omp_intervening_code): Reject imperfectly nested tile. (cp_parser_statement_seq_opt): If want_nested_loop, use cp_parser_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. (cp_parser_omp_clause_name): Handle full and partial clause names. (cp_parser_omp_clause_full, cp_parser_omp_clause_partial): New functions. (cp_parser_omp_all_clauses): Formatting fix. Handle PRAGMA_OMP_CLAUSE_PARTIAL and PRAGMA_OMP_CLAUSE_FULL. (cp_parser_next_tokens_can_be_canon_loop): New function. (cp_parser_omp_loop_nest): Parse C++11 attributes. Handle tile/unroll constructs. Use cp_parser_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. Only add_stmt cp_parser_omp_loop_nest result if it is non-NULL. (cp_parser_omp_for_loop): Rename tiling variable to oacc_tiling. For OMP_CLAUSE_SIZES set collapse to list length of OMP_CLAUSE_SIZES_LIST. Use cp_parser_next_tokens_can_be_canon_loop instead of just checking for RID_FOR keyword. Remove spurious semicolon. Don't call c_omp_check_loop_binding_exprs if stmt is NULL. Skip and/or handle generated loops. Remove spurious ()s around & operands. (cp_parser_omp_tile_sizes, cp_parser_omp_tile): New functions. (OMP_UNROLL_CLAUSE_MASK): Define. (cp_parser_omp_unroll): New function. (cp_parser_omp_construct): Handle PRAGMA_OMP_TILE and PRAGMA_OMP_UNROLL. (cp_parser_pragma): Likewise. * semantics.cc (finish_omp_clauses): Don't call fold_build_cleanup_point_expr for cases which obviously won't need it, like checked INTEGER_CSTs. Handle OMP_CLAUSE_{FULL,PARTIAL,SIZES} and diagnose full vs. partial conflict. Adjust wording of some of the conflicting clause diagnostic messages to include word clause. (finish_omp_for): Use decl equal to global_namespace as a marker for generated loop. Pass also body to dependent_omp_for_p. Skip generated loops. (finish_omp_for_block): Skip generated loops. * pt.cc (tsubst_omp_clauses): Handle OMP_CLAUSE_{FULL,PARTIAL,SIZES}. (tsubst_stmt): Handle OMP_TILE and OMP_UNROLL. Handle or skip generated loops. (dependent_omp_for_p): Add body argument. If declv vector element is NULL, find generated loop. * cp-gimplify.cc (cp_gimplify_expr): Handle OMP_TILE and OMP_UNROLL. (cp_fold_r): Likewise. (cp_genericize_r): Likewise. Skip generated loops. gcc/fortran/ * gfortran.h (enum gfc_statement): Add ST_OMP_UNROLL, ST_OMP_END_UNROLL, ST_OMP_TILE and ST_OMP_END_TILE. (struct gfc_omp_clauses): Add sizes_list, partial, full and erroneous members. (enum gfc_exec_op): Add EXEC_OMP_UNROLL and EXEC_OMP_TILE. (gfc_expr_list_len): Declare. * match.h (gfc_match_omp_tile, gfc_match_omp_unroll): Declare. * openmp.cc (gfc_get_location): Declare. (gfc_free_omp_clauses): Free sizes_list. (match_oacc_expr_list): Rename to ... (match_omp_oacc_expr_list): ... this. Add is_omp argument and change diagnostic wording if it is true. (enum omp_mask2): Add OMP_CLAUSE_{FULL,PARTIAL,SIZES}. (gfc_match_omp_clauses): Parse full, partial and sizes clauses. (gfc_match_oacc_wait): Use match_omp_oacc_expr_list instead of match_oacc_expr_list. (OMP_UNROLL_CLAUSES, OMP_TILE_CLAUSES): Define. (gfc_match_omp_tile, gfc_match_omp_unroll): New functions. (resolve_omp_clauses): Diagnose full vs. partial clause conflict. Resolve sizes clause arguments. (find_nested_loop_in_chain): Use switch instead of series of ifs. Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (gfc_resolve_omp_do_blocks): Set omp_current_do_collapse to list length of sizes_list if present. (gfc_resolve_do_iterator): Return for EXEC_OMP_TILE or EXEC_OMP_UNROLL. (restructure_intervening_code): Remove spurious ()s around & operands. (is_outer_iteration_variable): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (check_nested_loop_in_chain): Likewise. (expr_is_invariant): Likewise. (resolve_omp_do): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. Diagnose tile without sizes clause. Use sizes_list length for count if non-NULL. Set code->ext.omp_clauses->erroneous on loops where we've reported diagnostics. Sorry for mixing non-rectangular loops with generated loops. (omp_code_to_statement): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (gfc_resolve_omp_directive): Likewise. * parse.cc (decode_omp_directive): Parse end tile, end unroll, tile and unroll. Move nothing entry alphabetically. (case_exec_markers): Add ST_OMP_TILE and ST_OMP_UNROLL. (gfc_ascii_statement): Handle ST_OMP_END_TILE, ST_OMP_END_UNROLL, ST_OMP_TILE and ST_OMP_UNROLL. (parse_omp_do): Add nested argument. Handle ST_OMP_TILE and ST_OMP_UNROLL. (parse_omp_structured_block): Adjust parse_omp_do caller. (parse_executable): Likewise. Handle ST_OMP_TILE and ST_OMP_UNROLL. * resolve.cc (gfc_resolve_blocks): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (gfc_resolve_code): Likewise. * st.cc (gfc_free_statement): Likewise. * trans.cc (trans_code): Likewise. * trans-openmp.cc (gfc_trans_omp_clauses): Handle full, partial and sizes clauses. Use tree_cons + nreverse instead of temporary vector and build_tree_list_vec for tile_list handling. (gfc_expr_list_len): New function. (gfc_trans_omp_do): Rename tile to oacc_tile. Handle sizes clause. Don't assert code->op is EXEC_DO. Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (gfc_trans_omp_directive): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. * dump-parse-tree.cc (show_omp_clauses): Dump full, partial and sizes clauses. (show_omp_node): Handle EXEC_OMP_TILE and EXEC_OMP_UNROLL. (show_code_node): Likewise. gcc/testsuite/ * c-c++-common/gomp/attrs-tile-1.c: New test. * c-c++-common/gomp/attrs-tile-2.c: New test. * c-c++-common/gomp/attrs-tile-3.c: New test. * c-c++-common/gomp/attrs-tile-4.c: New test. * c-c++-common/gomp/attrs-tile-5.c: New test. * c-c++-common/gomp/attrs-tile-6.c: New test. * c-c++-common/gomp/attrs-unroll-1.c: New test. * c-c++-common/gomp/attrs-unroll-2.c: New test. * c-c++-common/gomp/attrs-unroll-3.c: New test. * c-c++-common/gomp/attrs-unroll-inner-1.c: New test. * c-c++-common/gomp/attrs-unroll-inner-2.c: New test. * c-c++-common/gomp/attrs-unroll-inner-3.c: New test. * c-c++-common/gomp/attrs-unroll-inner-4.c: New test. * c-c++-common/gomp/attrs-unroll-inner-5.c: New test. * c-c++-common/gomp/imperfect-attributes.c: Adjust expected diagnostics. * c-c++-common/gomp/imperfect-loop-nest.c: New test. * c-c++-common/gomp/ordered-5.c: New test. * c-c++-common/gomp/scan-7.c: New test. * c-c++-common/gomp/tile-1.c: New test. * c-c++-common/gomp/tile-2.c: New test. * c-c++-common/gomp/tile-3.c: New test. * c-c++-common/gomp/tile-4.c: New test. * c-c++-common/gomp/tile-5.c: New test. * c-c++-common/gomp/tile-6.c: New test. * c-c++-common/gomp/tile-7.c: New test. * c-c++-common/gomp/tile-8.c: New test. * c-c++-common/gomp/tile-9.c: New test. * c-c++-common/gomp/tile-10.c: New test. * c-c++-common/gomp/tile-11.c: New test. * c-c++-common/gomp/tile-12.c: New test. * c-c++-common/gomp/tile-13.c: New test. * c-c++-common/gomp/tile-14.c: New test. * c-c++-common/gomp/tile-15.c: New test. * c-c++-common/gomp/unroll-1.c: New test. * c-c++-common/gomp/unroll-2.c: New test. * c-c++-common/gomp/unroll-3.c: New test. * c-c++-common/gomp/unroll-4.c: New test. * c-c++-common/gomp/unroll-5.c: New test. * c-c++-common/gomp/unroll-6.c: New test. * c-c++-common/gomp/unroll-7.c: New test. * c-c++-common/gomp/unroll-8.c: New test. * c-c++-common/gomp/unroll-9.c: New test. * c-c++-common/gomp/unroll-inner-1.c: New test. * c-c++-common/gomp/unroll-inner-2.c: New test. * c-c++-common/gomp/unroll-inner-3.c: New test. * c-c++-common/gomp/unroll-non-rect-1.c: New test. * c-c++-common/gomp/unroll-non-rect-2.c: New test. * c-c++-common/gomp/unroll-non-rect-3.c: New test. * c-c++-common/gomp/unroll-simd-1.c: New test. * gcc.dg/gomp/attrs-4.c: Adjust expected diagnostics. * gcc.dg/gomp/for-1.c: Likewise. * gcc.dg/gomp/for-11.c: Likewise. * g++.dg/gomp/attrs-4.C: Likewise. * g++.dg/gomp/for-1.C: Likewise. * g++.dg/gomp/pr94512.C: Likewise. * g++.dg/gomp/tile-1.C: New test. * g++.dg/gomp/tile-2.C: New test. * g++.dg/gomp/unroll-1.C: New test. * g++.dg/gomp/unroll-2.C: New test. * g++.dg/gomp/unroll-3.C: New test. * gfortran.dg/gomp/inner-loops-1.f90: New test. * gfortran.dg/gomp/inner-loops-2.f90: New test. * gfortran.dg/gomp/pure-1.f90: Add tests for !$omp unroll and !$omp tile. * gfortran.dg/gomp/pure-2.f90: Remove those tests from here. * gfortran.dg/gomp/scan-9.f90: New test. * gfortran.dg/gomp/tile-1.f90: New test. * gfortran.dg/gomp/tile-2.f90: New test. * gfortran.dg/gomp/tile-3.f90: New test. * gfortran.dg/gomp/tile-4.f90: New test. * gfortran.dg/gomp/tile-5.f90: New test. * gfortran.dg/gomp/tile-6.f90: New test. * gfortran.dg/gomp/tile-7.f90: New test. * gfortran.dg/gomp/tile-8.f90: New test. * gfortran.dg/gomp/tile-9.f90: New test. * gfortran.dg/gomp/tile-10.f90: New test. * gfortran.dg/gomp/tile-imperfect-nest-1.f90: New test. * gfortran.dg/gomp/tile-imperfect-nest-2.f90: New test. * gfortran.dg/gomp/tile-inner-loops-1.f90: New test. * gfortran.dg/gomp/tile-inner-loops-2.f90: New test. * gfortran.dg/gomp/tile-inner-loops-3.f90: New test. * gfortran.dg/gomp/tile-inner-loops-4.f90: New test. * gfortran.dg/gomp/tile-inner-loops-5.f90: New test. * gfortran.dg/gomp/tile-inner-loops-6.f90: New test. * gfortran.dg/gomp/tile-inner-loops-7.f90: New test. * gfortran.dg/gomp/tile-inner-loops-8.f90: New test. * gfortran.dg/gomp/tile-non-rectangular-1.f90: New test. * gfortran.dg/gomp/tile-non-rectangular-2.f90: New test. * gfortran.dg/gomp/tile-non-rectangular-3.f90: New test. * gfortran.dg/gomp/tile-unroll-1.f90: New test. * gfortran.dg/gomp/tile-unroll-2.f90: New test. * gfortran.dg/gomp/unroll-1.f90: New test. * gfortran.dg/gomp/unroll-2.f90: New test. * gfortran.dg/gomp/unroll-3.f90: New test. * gfortran.dg/gomp/unroll-4.f90: New test. * gfortran.dg/gomp/unroll-5.f90: New test. * gfortran.dg/gomp/unroll-6.f90: New test. * gfortran.dg/gomp/unroll-7.f90: New test. * gfortran.dg/gomp/unroll-8.f90: New test. * gfortran.dg/gomp/unroll-9.f90: New test. * gfortran.dg/gomp/unroll-10.f90: New test. * gfortran.dg/gomp/unroll-11.f90: New test. * gfortran.dg/gomp/unroll-12.f90: New test. * gfortran.dg/gomp/unroll-13.f90: New test. * gfortran.dg/gomp/unroll-inner-loop-1.f90: New test. * gfortran.dg/gomp/unroll-inner-loop-2.f90: New test. * gfortran.dg/gomp/unroll-no-clause-1.f90: New test. * gfortran.dg/gomp/unroll-non-rect-1.f90: New test. * gfortran.dg/gomp/unroll-non-rect-2.f90: New test. * gfortran.dg/gomp/unroll-simd-1.f90: New test. * gfortran.dg/gomp/unroll-simd-2.f90: New test. * gfortran.dg/gomp/unroll-simd-3.f90: New test. * gfortran.dg/gomp/unroll-tile-1.f90: New test. * gfortran.dg/gomp/unroll-tile-2.f90: New test. * gfortran.dg/gomp/unroll-tile-inner-1.f90: New test. libgomp/ * testsuite/libgomp.c-c++-common/imperfect-transform-1.c: New test. * testsuite/libgomp.c-c++-common/imperfect-transform-2.c: New test. * testsuite/libgomp.c-c++-common/matrix-1.h: New test. * testsuite/libgomp.c-c++-common/matrix-constant-iter.h: New test. * testsuite/libgomp.c-c++-common/matrix-helper.h: New test. * testsuite/libgomp.c-c++-common/matrix-no-directive-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-no-directive-unroll-full-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-distribute-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-parallel-masked-taskloop-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-parallel-masked-taskloop-simd-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-target-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-target-teams-distribute-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-taskloop-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-teams-distribute-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-simd-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-transform-variants-1.h: New test. * testsuite/libgomp.c-c++-common/target-imperfect-transform-1.c: New test. * testsuite/libgomp.c-c++-common/target-imperfect-transform-2.c: New test. * testsuite/libgomp.c-c++-common/unroll-1.c: New test. * testsuite/libgomp.c-c++-common/unroll-non-rect-1.c: New test. * testsuite/libgomp.c++/matrix-no-directive-unroll-full-1.C: New test. * testsuite/libgomp.c++/tile-2.C: New test. * testsuite/libgomp.c++/tile-3.C: New test. * testsuite/libgomp.c++/unroll-1.C: New test. * testsuite/libgomp.c++/unroll-2.C: New test. * testsuite/libgomp.c++/unroll-full-tile.C: New test. * testsuite/libgomp.fortran/imperfect-transform-1.f90: New test. * testsuite/libgomp.fortran/imperfect-transform-2.f90: New test. * testsuite/libgomp.fortran/inner-1.f90: New test. * testsuite/libgomp.fortran/nested-fn.f90: New test. * testsuite/libgomp.fortran/target-imperfect-transform-1.f90: New test. * testsuite/libgomp.fortran/target-imperfect-transform-2.f90: New test. * testsuite/libgomp.fortran/tile-1.f90: New test. * testsuite/libgomp.fortran/tile-2.f90: New test. * testsuite/libgomp.fortran/tile-unroll-1.f90: New test. * testsuite/libgomp.fortran/tile-unroll-2.f90: New test. * testsuite/libgomp.fortran/tile-unroll-3.f90: New test. * testsuite/libgomp.fortran/tile-unroll-4.f90: New test. * testsuite/libgomp.fortran/unroll-1.f90: New test. * testsuite/libgomp.fortran/unroll-2.f90: New test. * testsuite/libgomp.fortran/unroll-3.f90: New test. * testsuite/libgomp.fortran/unroll-4.f90: New test. * testsuite/libgomp.fortran/unroll-5.f90: New test. * testsuite/libgomp.fortran/unroll-6.f90: New test. * testsuite/libgomp.fortran/unroll-7a.f90: New test. * testsuite/libgomp.fortran/unroll-7b.f90: New test. * testsuite/libgomp.fortran/unroll-7c.f90: New test. * testsuite/libgomp.fortran/unroll-7.f90: New test. * testsuite/libgomp.fortran/unroll-8.f90: New test. * testsuite/libgomp.fortran/unroll-simd-1.f90: New test. * testsuite/libgomp.fortran/unroll-tile-1.f90: New test. * testsuite/libgomp.fortran/unroll-tile-2.f90: New test. |
||
|
Kwok Cheung Yeung
|
637e76b90e |
openmp: Change to using a hashtab to lookup offload target addresses for indirect function calls
A splay-tree was previously used to lookup equivalent target addresses for a given host address on offload targets. However, as splay-trees can modify their structure on lookup, they are not suitable for concurrent access from separate teams/threads without some form of locking. This patch changes the lookup data structure to a hashtab instead, which does not have these issues. The call to build_indirect_map to initialize the data structure is now called from just the first thread of the first team to avoid redundant calls to this function. 2024-03-22 Kwok Cheung Yeung <kcyeung@baylibre.com> libgomp/ * config/accel/target-indirect.c: Include string.h and hashtab.h. Remove include of splay-tree.h. Update comments. (splay_tree_prefix, splay_tree_c): Delete. (struct indirect_map_t): New. (hash_entry_type, htab_alloc, htab_free, htab_hash, htab_eq): New. (GOMP_INDIRECT_ADD_MAP): Remove volatile qualifier. (USE_SPLAY_TREE_LOOKUP): Rename to... (USE_HASHTAB_LOOKUP): ..this. (indirect_map, indirect_array): Delete. (indirect_htab): New. (build_indirect_map): Remove locking. Build indirect map using hashtab. (GOMP_target_map_indirect_ptr): Use indirect_htab to lookup target address. (GOMP_target_map_indirect_ptr): Remove volatile qualifier. * config/gcn/team.c (gomp_gcn_enter_kernel): Call build_indirect_map from first thread of first team only. * config/nvptx/team.c (gomp_nvptx_main): Likewise. * testsuite/libgomp.c-c++-common/declare-target-indirect-2.c (main): Add missing break statements. * testsuite/libgomp.fortran/declare-target-indirect-2.f90: Remove xfail. |
||
|
Tobias Burnus
|
c5037fcee2 |
OpenMP/Fortran: Fix defaultmap(none) issue with dummy procedures [PR114283]
Dummy procedures look similar to variables but aren't - neither in Fortran nor in OpenMP. As the middle end sees PARM_DECLs, mark them as predetermined firstprivate for mapping (as already done in gfc_omp_predetermined_sharing). This does not address the isses related to procedure pointers, which are still discussed on spec level [see PR]. PR fortran/114283 gcc/fortran/ChangeLog: * trans-openmp.cc (gfc_omp_predetermined_mapping): Map dummy procedures as firstprivate. libgomp/ChangeLog: * testsuite/libgomp.fortran/declare-target-indirect-4.f90: New test. |
||
|
Kwok Cheung Yeung
|
451bb58660 |
openmp, fortran: Add Fortran support for indirect clause on the declare target directive
2024-02-15 Kwok Cheung Yeung <kcyeung@baylibre.com> gcc/fortran/ * dump-parse-tree.cc (show_attr): Handle omp_declare_target_indirect attribute. * f95-lang.cc (gfc_gnu_attributes): Add entry for 'omp declare target indirect'. * gfortran.h (symbol_attribute): Add omp_declare_target_indirect field. (struct gfc_omp_clauses): Add indirect field. * openmp.cc (omp_mask2): Add OMP_CLAUSE_INDIRECT. (gfc_match_omp_clauses): Match indirect clause. (OMP_DECLARE_TARGET_CLAUSES): Add OMP_CLAUSE_INDIRECT. (gfc_match_omp_declare_target): Check omp_device_type and apply omp_declare_target_indirect attribute to symbol if indirect clause active. Show warning if there are only device_type and/or indirect clauses on the directive. * trans-decl.cc (add_attributes_to_decl): Add 'omp declare target indirect' attribute if symbol has indirect attribute set. gcc/testsuite/ * gfortran.dg/gomp/declare-target-4.f90 (f1): Update expected warning. * gfortran.dg/gomp/declare-target-indirect-1.f90: New. * gfortran.dg/gomp/declare-target-indirect-2.f90: New. libgomp/ * testsuite/libgomp.fortran/declare-target-indirect-1.f90: New. * testsuite/libgomp.fortran/declare-target-indirect-2.f90: New. * testsuite/libgomp.fortran/declare-target-indirect-3.f90: New. |
||
|
John David Anglin
|
c9c507ff8b |
Increase timeout by 2 in libgomp.fortran/alloc-comp-3.f90 on hppa*-*-*
2024-01-20 John David Anglin <danglin@gcc.gnu.org> libgomp/ChangeLog: * testsuite/libgomp.fortran/alloc-comp-3.f90: Increase timeout by 2 on hppa*-*-*. |
||
|
Julian Brown
|
144c531fe2 |
OpenMP/OpenACC: Reorganise OMP map clause handling in gimplify.cc
This patch has been separated out from the C++ "declare mapper" support patch. It contains just the gimplify.cc rearrangement work, mostly moving gimplification from gimplify_scan_omp_clauses to gimplify_adjust_omp_clauses for map clauses. The motivation for doing this was that we don't know if we need to instantiate mappers implicitly until the body of an offload region has been scanned, i.e. in gimplify_adjust_omp_clauses, but we also need the un-gimplified form of clauses to sort by base-pointer dependencies after mapper instantiation has taken place. The patch also reimplements the "present" clause sorting code to avoid another sorting pass on mapping nodes. This version of the patch is based on the version posted for og13, and additionally incorporates a follow-on fix for DECL_VALUE_EXPR handling in gimplify_adjust_omp_clauses: "OpenMP/OpenACC: Reorganise OMP map clause handling in gimplify.cc" https://gcc.gnu.org/pipermail/gcc-patches/2023-June/622223.html Parts of: "OpenMP: OpenMP 5.2 semantics for pointers with unmapped target" https://gcc.gnu.org/pipermail/gcc-patches/2023-June/623351.html 2023-12-16 Julian Brown <julian@codesourcery.com> gcc/ * gimplify.cc (omp_segregate_mapping_groups): Handle "present" groups. (gimplify_scan_omp_clauses): Use mapping group functionality to iterate through mapping nodes. Remove most gimplification of OMP_CLAUSE_MAP nodes from here, but still populate ctx->variables splay tree. (gimplify_adjust_omp_clauses): Move most gimplification of OMP_CLAUSE_MAP nodes here. libgomp/ * testsuite/libgomp.fortran/target-enter-data-6.f90: Remove XFAIL. |