The implementation of each allocator is simplified by inheriting from
__detail::__allocator_templ. At the moment, none of the implementations
diverge in any way, simply passing in the allocator handle to be used when
an allocation is made. In the future, const_mem will need special handling
added to it to support constant memory space.
libgomp/ChangeLog:
* omp.h.in: Add omp::allocator::* and ompx::allocator::* allocators.
(__detail::__allocator_templ<T, omp_allocator_handle_t>):
New struct template.
(null_allocator<T>): New struct template.
(default_mem<T>): Likewise.
(large_cap_mem<T>): Likewise.
(const_mem<T>): Likewise.
(high_bw_mem<T>): Likewise.
(low_lat_mem<T>): Likewise.
(cgroup_mem<T>): Likewise.
(pteam_mem<T>): Likewise.
(thread_mem<T>): Likewise.
(ompx::allocator::gnu_pinned_mem<T>): Likewise.
* testsuite/libgomp.c++/allocator-1.C: New test.
* testsuite/libgomp.c++/allocator-2.C: New test.
Signed-off-by: waffl3x <waffl3x@baylibre.com>
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.
The new SVE tests didn't explicitly force SVE to be enabled,
which meant that they wouldn't work on targets that aren't
configured for SVE by default. The least invasive way of
fixing that is to add a pragma, which works for most tests.
However, for udr-sve.c, the global:
#pragma omp declare reduction (+:svint32_t: omp_out = svadd_s32_z (svptrue_b32(), omp_in, omp_out)) \
initializer (omp_priv = svindex_s32 (0, 0))
does not work with an earlier:
#pragma GCC target "+sve"
which is interesting, and maybe worthy of a PR if there isn't one
already. It seems we have to force SVE (and thus an architecture)
on the command line instead.
However, with that fixed, udr-sve.c fails execution. One problem
seems to be a missing accumulation in for_reduction. Fixing that
is enough to reach the final inscan_reduction_incl, but that fails
for reasons I haven't investigated yet. I would need to read up
more to understand what the loop is doing.
It also looks like there might be a missing "+" in simd_reduction:
#pragma omp simd reduction (+:va, i)
for (j = 0; j < 16; j++)
va = svld1_s32 (svptrue_b32 (), a);
res = svaddv_s32 (svptrue_b32 (), va);
if (res != 8)
__builtin_abort ();
since AFAICT the loop is not doing a reduction as things stand.
But perhaps that's deliberate, since it does match the != 8 test.
libgomp/
* testsuite/libgomp.c-target/aarch64/firstprivate.c: Add +sve pragma.
* testsuite/libgomp.c-target/aarch64/lastprivate.c: Likewise.
* testsuite/libgomp.c-target/aarch64/private.c: Likewise.
* testsuite/libgomp.c-target/aarch64/shared.c: Likewise.
* testsuite/libgomp.c-target/aarch64/simd-aligned.c: Likewise.
* testsuite/libgomp.c-target/aarch64/simd-nontemporal.c: Likewise.
* testsuite/libgomp.c-target/aarch64/threadprivate.c: Likewise.
* testsuite/libgomp.c-target/aarch64/udr-sve.c: Add an -march option.
(for_reduction): Use "+=" in the reduction loop.
At tree level, the addr ref is also required for array dummy arguments,
contrary to C; the GOMP_interop calls in modify_call_for_omp_dispatch
were updated accordingly (using build_fold_addr_expr).
As the GOMP_interop calls had no location data associated with them,
the init call happened as soon as executing the previous line of code,
which was confusing; solution: use the location data of the function
call itself.
PR middle-end/119662
gcc/ChangeLog:
* gimplify.cc (modify_call_for_omp_dispatch): Fix GOMP_interop
arg passing; add location info to function calls.
libgomp/ChangeLog:
* testsuite/libgomp.c/append-args-fr-1.c: New test.
* testsuite/libgomp.c/append-args-fr.h: New test.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/append-args-interop.c: Update for fixed
GOMP_interop call.
* g++.dg/gomp/append-args-8.C: Likewise.
* gfortran.dg/gomp/append-args-interop.f90: Likewise.
Silence the overeager "Warning: Variable 'depobj_list' at (1) is a dummy
argument of the BIND(C) procedure ... but may not be C interoperable
[-Wc-binding-type]" when compiling omp_lib.f90(.in).
The argument is of integer kind 'omp_depend_kind = @OMP_DEPEND_KIND@'.
libgomp/ChangeLog:
* Makefile.am (%.mod): Add -Wno-c-binding-type.
* Makefile.in: Regenerate.
..., so that users don't manually need to specify '-foffload-options=-lstdc++'
in addition to '-lstdc++' (specified manually, or implicitly by the driver).
Do like commit 4bcb46b3ad
"driver: Forward '-lgfortran', '-lm' to offloading compilation".
PR driver/101544
gcc/
* gcc.cc (driver_handle_option): Forward host '-lstdc++' to
offloading compilation.
* config/gcn/mkoffload.cc (main): Adjust.
* config/nvptx/mkoffload.cc (main): Likewise.
libgomp/
* testsuite/libgomp.c++/pr101544-1-O0.C: Remove
'-foffload-options=-lstdc++'.
* testsuite/libgomp.c++/pr101544-1.C: Likewise.
* testsuite/libgomp.oacc-c++/pr101544-1.C: Likewise.
Note that this commit also updates the API interface to OpenMP 6.0;
while 5.1 and 5.2 use 'int *' for the the ret_code argument,
OpenMP 6.0 changed this to omp_interop_rc_t *; this enum also exists in
OpenMP 5.1. However, C++ does not like this change such that unless NULL
is passed (i.e. the argument is ignored), OpenMP 5.x and 6.x are not
compatible.
Note that GCC's omp.h already follows OpenMP 6.0 and is now in sync with
the documentation.
libgomp/ChangeLog:
* libgomp.texi (OpenMP 5.1): Add @ref to offload-target specifics
for 'interop'.
(OpenMP 6.0): Mark dispatch's interop clause as implemented.
(omp_get_interop_int, omp_get_interop_str,
omp_get_interop_ptr, omp_get_interop_type_desc): Add @ref to
Offload-Target Specifics; change ret_code argument type to
'omp_interop_rc_t *'.
(Offload-Target Specifics): Document the supported OpenMP
interop foreign runtimes on AMD and Nvidia GPUs.
This patch adds support for the case where #pragma omp declare variant
with append_args is used inside a #pragma omp dispatch interop that
specifies fewer interop args than required by the variant; new interop
objects are implicitly created and then destroyed around the call to the
variant, using the GOMP_interop builtin.
gcc/fortran/ChangeLog
* trans-openmp.cc (gfc_trans_omp_declare_variant): Remove accidental
redeclaration of pref.
gcc/ChangeLog
* gimplify.cc (modify_call_for_omp_dispatch): Adjust arguments.
Remove the "sorry" for the case where new interop objects must be
constructed, and add code to make it work instead.
(expand_variant_call_expr): Adjust arguments and call to
modify_call_for_omp_dispatch.
(gimplify_variant_call_expr): Simplify logic for calling
expand_variant_call_expr.
gcc/testsuite/ChangeLog
* c-c++-common/gomp/append-args-1.c: Adjust expected behavior.
* c-c++-common/gomp/append-args-interop.c: New.
* c-c++-common/gomp/dispatch-11.c: Adjust expected behavior.
* g++.dg/gomp/append-args-1.C: Likewise.
* gfortran.dg/gomp/append-args-interop.f90: New.
* gfortran.dg/gomp/declare-variant-mod-2.f90: Adjust expected behavior.
libgomp/ChangeLog
* libgomp.texi (OpenMP 5.1): Mark append_args as fully supported.
Co-Authored-By: Tobias Burnus <tburnus@baylibre.com>
The interop object (opaque object to the user, used internally in libgomp)
already had a 'device_num' member, but it was missed to actually set it.
libgomp/ChangeLog:
* target.c (gomp_interop_internal): Set the 'device_num' member
when initializing an interop object.
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.
The interop directive operates on an opaque object that represents a
foreign runtime. This commit adds support for
this to the two offloading plugins.
For nvptx, it supports cuda, cuda_driver and hip; the latter is AMD's
version of CUDA which for Nvidia devices boils down to normal CUDA.
Thus, at the end for this limited use, cuda/cuda_driver/hip are all
the same - and for plugin-nvptx.c, the they differ only in terms of
what gets fr_id, fr_name and get_interop_type_desc return.
For gcn, it supports hip and hsa.
Regarding get-mapped-ptr-1.c: That's actually 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:
* plugin/plugin-gcn.c (_LIBGOMP_PLUGIN_INCLUDE): Define.
(struct hsa_runtime_fn_info): Add two queue functions.
(hipError_t, hipCtx_t, hipStream_s, hipStream_t): New types.
(struct hip_runtime_fn_info): New.
(hip_runtime_lib, hip_fns): New global vars.
(init_environment_variables): Handle hip_runtime_lib.
(init_hsa_runtime_functions): Load the two queue functions.
(init_hip_runtime_functions, GOMP_OFFLOAD_interop,
GOMP_OFFLOAD_get_interop_int, GOMP_OFFLOAD_get_interop_ptr,
GOMP_OFFLOAD_get_interop_str,
GOMP_OFFLOAD_get_interop_type_desc): New.
* plugin/plugin-nvptx.c (_LIBGOMP_PLUGIN_INCLUDE): Define.
(GOMP_OFFLOAD_interop, GOMP_OFFLOAD_get_interop_int,
GOMP_OFFLOAD_get_interop_ptr, GOMP_OFFLOAD_get_interop_str,
GOMP_OFFLOAD_get_interop_type_desc): New.
* testsuite/libgomp.c/interop-fr-1.c: New test.
* testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c: Use -6
not -5 as non-conforming device number.
libgomp.exp added -fno-diagnostics-show-caret and -fdiagnostics-color=never
as 'additional_flags' for compilation. However, it turned out that this now
is insufficient as the [...] part of diagnostics have a hyperlink URL.
Solution: Use the -fdiagnostics-plain-output flag instead, added in commit
r11-2701-g129a1319c0ab73. This flag currently implies the following flags:
-fno-diagnostics-show-caret
-fno-diagnostics-show-line-numbers
-fdiagnostics-color=never
-fdiagnostics-urls=never
-fdiagnostics-path-format=separate-events
-fdiagnostics-text-art-charset=none
-fno-diagnostics-show-event-links
libgomp/ChangeLog:
* testsuite/lib/libgomp.exp (libgomp_init): Add
-fdiagnostics-plain-output to additional_flags; remove
-fno-diagnostics-show-caret and -fdiagnostics-color=never.
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.
Without this attribute, the building process will fail if GCC is configured
with 'CFLAGS=-DNDEBUG'.
libgomp/ChangeLog:
* oacc-mem.c (acc_unmap_data, goacc_exit_datum_1, find_group_last,
goacc_enter_data_internal): Add '__attribute__((unused))'.
* target.c (gomp_unmap_vars_internal): Likewise.
Since -frange-for-ext-temps has been causing trouble, let's not enable it
by default in pre-C++23 GNU modes for GCC 15, and also allow disabling it in
C++23 and up.
PR c++/188574
gcc/c-family/ChangeLog:
* c-opts.cc (c_common_post_options): Only enable
-frange-for-ext-temps by default in C++23.
gcc/ChangeLog:
* doc/invoke.texi: Adjust -frange-for-ext-temps documentation.
gcc/testsuite/ChangeLog:
* g++.dg/cpp23/range-for3.C: Use -frange-for-ext-temps.
* g++.dg/cpp23/range-for4.C: Adjust expected result.
libgomp/ChangeLog:
* testsuite/libgomp.c++/range-for-4.C: Adjust expected result.
libgomp/ChangeLog:
* plugin/plugin-gcn.c (ELFABIVERSION_AMDGPU_HSA_V6,
EF_AMDGPU_GENERIC_VERSION_V, EF_AMDGPU_GENERIC_VERSION_OFFSET,
GET_GENERIC_VERSION): New #define.
(elf_gcn_isa_is_generic): New.
(isa_matches_agent): Accept all generic code objects on the first
go; extend the diagnostic and handle runtime-failed case.
(create_and_finalize_hsa_program): Call it also after loading
the code failed, pass the status.
libgomp/ChangeLog
* libgomp.texi (OpenMP 5.0): Mark metadirective and declare variant
as implemented.
(OpenMP 5.1): Mark target_device as supported.
Add changed interaction between declare target and OpenMP context
and dynamic selector support.
(OpenMP 5.2): Mark otherwise clause as supported, note that
default is also still accepted.
This fixes a large number of smaller and larger issues with the append_args
clause to 'declare variant' and adds Fortran support for it; it also contains
a larger number of testcases.
In particular, for Fortran, it also handles passing allocatable, pointer,
optional arguments to an interop dummy argument with or without value
attribute. And it changes the internal representation such that dumping the
tree does not lead to an ICE.
gcc/c/ChangeLog:
* c-parser.cc (c_finish_omp_declare_variant): Modify how
append_args is saved internally.
gcc/cp/ChangeLog:
* parser.cc (cp_finish_omp_declare_variant): Modify how append_args
is saved internally.
* pt.cc (tsubst_attribute): Likewise.
(tsubst_omp_clauses): Remove C_ORT_OMP_DECLARE_SIMD from interop
handling as no longer called for it.
* decl.cc (omp_declare_variant_finalize_one): Update append_args
changes; fixes for ADL input.
gcc/fortran/ChangeLog:
* gfortran.h (gfc_omp_declare_variant): Add append_args_list.
* openmp.cc (gfc_parser_omp_clause_init_modifiers): New;
splitt of from ...
(gfc_match_omp_init): ... here; call it.
(gfc_match_omp_declare_variant): Update to handle append_args
clause; some syntax handling fixes.
* trans-openmp.cc (gfc_trans_omp_declare_variant): Handle
append_args clause; add some diagnostic.
gcc/ChangeLog:
* gimplify.cc (gimplify_call_expr): For OpenMP's append_args clause
processed by 'omp dispatch', update for internal-representation
changes; fix handling of hidden arguments, add some comments and
handle Fortran's value dummy and optional/pointer/allocatable actual
args.
libgomp/ChangeLog:
* libgomp.texi (Impl. Status): Update for accumpulated changes
related to 'dispatch' and interop.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/append-args-1.c: Update dg-*.
* c-c++-common/gomp/append-args-3.c: Likewise.
* g++.dg/gomp/append-args-1.C: Likewise.
* gfortran.dg/gomp/adjust-args-1.f90: Likewise.
* gfortran.dg/gomp/adjust-args-3.f90: Likewise.
* gfortran.dg/gomp/declare-variant-2.f90: Likewise.
* c-c++-common/gomp/append-args-6.c: New test.
* c-c++-common/gomp/append-args-7.c: New test.
* c-c++-common/gomp/append-args-8.c: New test.
* c-c++-common/gomp/append-args-9.c: New test.
* g++.dg/gomp/append-args-4.C: New test.
* g++.dg/gomp/append-args-5.C: New test.
* g++.dg/gomp/append-args-6.C: New test.
* g++.dg/gomp/append-args-7.C: New test.
* gcc.dg/gomp/append-args-1.c: New test.
* gfortran.dg/gomp/append_args-1.f90: New test.
* gfortran.dg/gomp/append_args-2.f90: New test.
* gfortran.dg/gomp/append_args-3.f90: New test.
* gfortran.dg/gomp/append_args-4.f90: New test.
