The libgomp.c/alloc-pinned*.c test have
/* { dg-skip-if "Pinning not implemented on this host" { ! *-*-linux-gnu* } } */
so they are only run on Linux targets right now. Duplicating the tests or
reworking them into headers looked like too much work for me right now this
late in stage4, so I've just #ifdefed the uses at least for now.
2025-04-17 Jakub Jelinek <jakub@redhat.com>
PR libgomp/119849
* testsuite/libgomp.c++/allocator-1.C (test_inequality, main): Guard
ompx::allocator::gnu_pinned_mem uses with #ifdef __gnu_linux__.
* testsuite/libgomp.c++/allocator-2.C (main): Likewise.
The HIP header files recognize the used compiler, defaulting to either AMD
or Nvidia/CUDA; thus, the alternative way of explicitly defining a macro is
less prominently documented. With GCC, the user has to define the
preprocessor macro manually. Hence, as a service to the user, mention
__HIP_PLATFORM_AMD__ and __HIP_PLATFORM_NVIDIA__ in the interop documentation,
even though it has only indirectly to do with GCC and its interop support.
Note to commit-log readers, only: For Fortran, the hipfort modules can be
used; when compiling the hipfort package (defaults to use gfortran), it
generates the module (*.mod) files in include/hipfort/{amdgcn,nvidia}/ such
that the choice is made by setting the respective include path.
libgomp/ChangeLog:
* libgomp.texi (gcn interop, nvptx interop): For HIP with C/C++, add
a note about setting a preprocessor define.
With commit ca9cffe737
"For nvptx offloading, make sure to emit C++ constructor, destructor aliases [PR97106]",
we're able to remove the 'ALWAYS_INLINE' workaround added in
commit fe283dba77
"GCN, nvptx: Support '-mfake-exceptions', and use it for offloading compilation [PR118794]".
libgomp/
* testsuite/libgomp.c++/target-exceptions-pr118794-1.C: Remove
'ALWAYS_INLINE' workaround.
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.
