Commit graph

1268 commits

Author SHA1 Message Date
Jakub Jelinek
5d05d496b2 libgomp: Don't test ompx::allocator::gnu_pinned_mem on non-linux targets.
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.
2025-04-17 12:14:15 +02:00
Thomas Schwinge
518efed8cb Remove 'ALWAYS_INLINE' workaround in 'libgomp.c++/target-exceptions-pr118794-1.C'
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.
2025-04-16 18:31:03 +02:00
Thomas Schwinge
0b2a2490be Add 'libgomp.c++/pr106445-1{,-O0}.C' [PR106445]
PR target/106445
	libgomp/
	* testsuite/libgomp.c++/pr106445-1.C: New.
	* testsuite/libgomp.c++/pr106445-1-O0.C: Likewise.
2025-04-16 16:07:07 +02:00
Thomas Schwinge
ca9cffe737 For nvptx offloading, make sure to emit C++ constructor, destructor aliases [PR97106]
PR target/97106
	gcc/
	* config/nvptx/nvptx.cc (nvptx_asm_output_def_from_decls)
	[ACCEL_COMPILER]: Make sure to emit C++ constructor, destructor
	aliases.
	libgomp/
	* testsuite/libgomp.c++/pr96390.C: Un-XFAIL nvptx offloading.
	* testsuite/libgomp.c-c++-common/pr96390.c: Adjust.
2025-04-16 16:03:04 +02:00
waffl3x
99835bd68e OpenMP: omp.h omp::allocator C++ Allocator interface
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>
2025-04-15 14:34:38 -06:00
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.
2025-04-15 16:42:42 +02:00
Thomas Schwinge
fe283dba77 GCN, nvptx: Support '-mfake-exceptions', and use it for offloading compilation [PR118794]
With '-mfake-exceptions' enabled, the user-visible behavior in presence of
exception handling constructs changes such that the compile-time
'sorry, unimplemented: exception handling not supported' is skipped, code
generation proceeds, and instead, exception handling constructs 'abort' at
run time.  (..., or don't, if they're in dead code.)

	PR target/118794
	gcc/
	* config/gcn/gcn.opt (-mfake-exceptions): Support.
	* config/nvptx/nvptx.opt (-mfake-exceptions): Likewise.
	* config/gcn/gcn.md (define_expand "exception_receiver"): Use it.
	* config/nvptx/nvptx.md (define_expand "exception_receiver"):
	Likewise.
	* config/gcn/mkoffload.cc (main): Set it.
	* config/nvptx/mkoffload.cc (main): Likewise.
	* config/nvptx/nvptx.cc (nvptx_assemble_integer)
	<in_section == exception_section>: Special handling for
	'SYMBOL_REF's.
	* except.cc (expand_dw2_landing_pad_for_region): Don't generate
	bogus code for (default)
	'#define EH_RETURN_DATA_REGNO(N) INVALID_REGNUM'.
	libgcc/
	* config/gcn/unwind-gcn.c (_Unwind_Resume): New.
	* config/nvptx/unwind-nvptx.c (_Unwind_Resume): Likewise.
	gcc/testsuite/
	* g++.target/gcn/exceptions-bad_cast-2.C: Set
	'-mno-fake-exceptions'.
	* g++.target/gcn/exceptions-pr118794-1.C: Likewise.
	* g++.target/gcn/exceptions-throw-2.C: Likewise.
	* g++.target/nvptx/exceptions-bad_cast-2.C: Likewise.
	* g++.target/nvptx/exceptions-pr118794-1.C: Likewise.
	* g++.target/nvptx/exceptions-throw-2.C: Likewise.
	* g++.target/gcn/exceptions-bad_cast-2_-mfake-exceptions.C: New.
	* g++.target/gcn/exceptions-pr118794-1_-mfake-exceptions.C:
	Likewise.
	* g++.target/gcn/exceptions-throw-2_-mfake-exceptions.C: Likewise.
	* g++.target/nvptx/exceptions-bad_cast-2_-mfake-exceptions.C:
	Likewise.
	* g++.target/nvptx/exceptions-pr118794-1_-mfake-exceptions.C:
	Likewise.
	* g++.target/nvptx/exceptions-throw-2_-mfake-exceptions.C:
	Likewise.
	libgomp/
	* testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C:
	Set '-foffload-options=-mno-fake-exceptions'.
	* testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C:
	Likewise.
	* testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C:
	Likewise.
	* testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C:
	Likewise.
	* testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C:
	Likewise.
	* testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C:
	Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C:
	Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C:
	Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C:
	Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C:
	Likewise.
	* testsuite/libgomp.c++/target-exceptions-bad_cast-2.C: Adjust.
	* testsuite/libgomp.c++/target-exceptions-pr118794-1.C: Likewise.
	* testsuite/libgomp.c++/target-exceptions-throw-2.C: Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C: Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-throw-2.C: Likewise.
	* testsuite/libgomp.c++/target-exceptions-throw-2-O0.C: New.
2025-04-14 23:56:05 +02:00
Thomas Schwinge
6c0ea84026 Add 'throw', dead code test cases for GCN, nvptx target and OpenACC, OpenMP 'target' offloading
gcc/testsuite/
	* g++.target/gcn/exceptions-throw-3.C: New.
	* g++.target/nvptx/exceptions-throw-3.C: Likewise.
	libgomp/
	* testsuite/libgomp.c++/target-exceptions-throw-3.C: New.
	* testsuite/libgomp.oacc-c++/exceptions-throw-3.C: Likewise.
2025-04-14 23:54:54 +02:00
Thomas Schwinge
1daf570498 Add 'throw', caught test cases for GCN, nvptx target and OpenACC, OpenMP 'target' offloading
gcc/testsuite/
	* g++.target/gcn/exceptions-throw-2.C: New.
	* g++.target/nvptx/exceptions-throw-2.C: Likewise.
	libgomp/
	* testsuite/libgomp.c++/target-exceptions-throw-2.C: New.
	* testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C: Likewise.
	* testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C: Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-throw-2.C: Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C: Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C: Likewise.
2025-04-14 23:54:54 +02:00
Thomas Schwinge
1362d9d494 Add 'throw' test cases for GCN, nvptx target and OpenACC, OpenMP 'target' offloading
gcc/testsuite/
	* g++.target/gcn/exceptions-throw-1.C: New.
	* g++.target/nvptx/exceptions-throw-1.C: Likewise.
	libgomp/
	* testsuite/libgomp.c++/target-exceptions-throw-1.C: New.
	* testsuite/libgomp.c++/target-exceptions-throw-1-O0.C: Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-throw-1.C: Likewise.
2025-04-14 23:54:53 +02:00
Thomas Schwinge
27f88cc799 Add 'std::bad_cast' exception, dead code test cases for GCN, nvptx target and OpenACC, OpenMP 'target' offloading
gcc/testsuite/
	* g++.target/gcn/exceptions-bad_cast-3.C: New.
	* g++.target/nvptx/exceptions-bad_cast-3.C: Likewise.
	libgomp/
	* testsuite/libgomp.c++/target-exceptions-bad_cast-3.C: New.
	* testsuite/libgomp.oacc-c++/exceptions-bad_cast-3.C: Likewise.
2025-04-14 23:54:53 +02:00
Thomas Schwinge
00cde164ee Add 'std::bad_cast' exception, caught test cases for GCN, nvptx target and OpenACC, OpenMP 'target' offloading
gcc/testsuite/
	* g++.target/gcn/exceptions-bad_cast-2.C: New.
	* g++.target/nvptx/exceptions-bad_cast-2.C: Likewise.
	libgomp/
	* testsuite/libgomp.c++/target-exceptions-bad_cast-2.C: New.
	* testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C: Likewise.
	* testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C: Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C: Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C: Likewise.
	* testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C: Likewise.
2025-04-14 23:54:53 +02:00
Thomas Schwinge
0e68f49db9 Add 'std::bad_cast' exception test cases for GCN, nvptx target and OpenACC, OpenMP 'target' offloading
gcc/testsuite/
	* g++.target/gcn/exceptions-bad_cast-1.C: New.
	* g++.target/nvptx/exceptions-bad_cast-1.C: Likewise.
	libgomp/
	* testsuite/libgomp.c++/target-exceptions-bad_cast-1.C: New.
	* testsuite/libgomp.oacc-c++/exceptions-bad_cast-1.C: Likewise.
2025-04-14 23:54:53 +02:00
Thomas Schwinge
aa3e72f943 Add test cases for exception handling constructs in dead code for GCN, nvptx target and OpenMP 'target' offloading [PR118794]
PR target/118794
	gcc/testsuite/
	* g++.target/gcn/exceptions-pr118794-1.C: New.
	* g++.target/nvptx/exceptions-pr118794-1.C: Likewise.
	libgomp/
	* testsuite/libgomp.c++/target-exceptions-pr118794-1.C: New.
	* testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C:
	Likewise.
	* testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C:
	Likewise.
2025-04-14 23:54:48 +02:00
Thomas Schwinge
a304c88b6f Add PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' offloading" test cases [PR119692]
... documenting the status quo.

	PR c++/119692
	gcc/testsuite/
	* g++.target/gcn/pr119692-1-1.C: New.
	* g++.target/nvptx/pr119692-1-1.C: Likewise.
	libgomp/
	* testsuite/libgomp.c++/pr119692-1-1.C: New.
	* testsuite/libgomp.c++/pr119692-1-2.C: Likewise.
	* testsuite/libgomp.c++/pr119692-1-3.C: Likewise.
	* testsuite/libgomp.c++/pr119692-1-4.C: Likewise.
	* testsuite/libgomp.c++/pr119692-1-5.C: Likewise.
	* testsuite/libgomp.oacc-c++/pr119692-1-1.C: Likewise.
	* testsuite/libgomp.oacc-c++/pr119692-1-2.C: Likewise.
	* testsuite/libgomp.oacc-c++/pr119692-1-3.C: Likewise.
2025-04-14 23:54:48 +02:00
Richard Sandiford
733a6a4c11 libgomp: Update SVE tests
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.
2025-04-10 21:09:07 +01:00
Tobias Burnus
0f77d88fdf OpenMP: Fix append_args handling in modify_call_for_omp_dispatch
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.
2025-04-08 13:47:53 +02:00
Tejas Belagod
a9bbb60b7c libgomp: Add AArch64 SVE target tests to libgomp.
Add AArch64 SVE target exectute tests to test various workshare constructs and
clauses with SVE types.

libgomp/ChangeLog:

	* testsuite/libgomp.c-target/aarch64/aarch64.exp: Test driver.
	* testsuite/libgomp.c-target/aarch64/firstprivate.c: New test.
	* 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: Likewise.
2025-04-08 11:50:50 +05:30
Thomas Schwinge
3b36e96b57 driver: Forward '-lstdc++' to offloading compilation [PR101544]
..., 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.
2025-03-26 14:20:20 +01:00
Thomas Schwinge
332a08173a Add 'libgomp.c++/pr96835-1{,-O0}.C', 'libgomp.oacc-c++/pr96835-1.C' [PR96835]
PR libgomp/96835
	libgomp/
	* testsuite/libgomp.c++/pr96835-1.C: New.
	* testsuite/libgomp.c++/pr96835-1-O0.C: Likewise.
	* testsuite/libgomp.oacc-c++/pr96835-1.C: Likewise.
2025-03-24 09:39:21 +01:00
Thomas Schwinge
62312c778a Add 'libgomp.c++/pr101544-1{,-O0}.C', 'libgomp.oacc-c++/pr101544-1.C' [PR101544]
PR target/101544
	libgomp/
	* testsuite/libgomp.c++/pr101544-1.C: New.
	* testsuite/libgomp.c++/pr101544-1-O0.C: Likewise.
	* testsuite/libgomp.oacc-c++/pr101544-1.C: Likewise.
2025-03-24 09:39:21 +01:00
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.
2025-03-22 00:36:44 +01:00
Tobias Burnus
41b9c3b848 libgomp/plugin: Add initial interop support to nvptx + gcn
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.
2025-03-21 21:39:42 +01:00
Paul-Antoine Arras
99e2906ae2 OpenMP: 'interop' construct - add ME support + target-independent libgomp
This patch partially enables use of the OpenMP interop construct by adding
middle end support, mostly in the omplower pass, and in the target-independent
part of the libgomp runtime. It follows up on previous patches for C, C++ and
Fortran front ends support. The full interop feature requires another patch to
enable foreign runtime support in libgomp plugins.

gcc/ChangeLog:

	* builtin-types.def
	(BT_FN_VOID_INT_INT_PTR_PTR_PTR_INT_PTR_INT_PTR_UINT_PTR): New.
	* gimple-low.cc (lower_stmt): Handle GIMPLE_OMP_INTEROP.
	* gimple-pretty-print.cc (dump_gimple_omp_interop): New function.
	(pp_gimple_stmt_1): Handle GIMPLE_OMP_INTEROP.
	* gimple.cc (gimple_build_omp_interop): New function.
	(gimple_copy): Handle GIMPLE_OMP_INTEROP.
	* gimple.def (GIMPLE_OMP_INTEROP): Define.
	* gimple.h (gimple_build_omp_interop): Declare.
	(gimple_omp_interop_clauses): New function.
	(gimple_omp_interop_clauses_ptr): Likewise.
	(gimple_omp_interop_set_clauses): Likewise.
	(gimple_return_set_retval): Handle GIMPLE_OMP_INTEROP.
	* gimplify.cc (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_INIT,
	OMP_CLAUSE_USE and OMP_CLAUSE_DESTROY.
	(gimplify_omp_interop): New function.
	(gimplify_expr): Replace sorry with call to gimplify_omp_interop.
	* omp-builtins.def (BUILT_IN_GOMP_INTEROP): Define.
	* omp-low.cc (scan_sharing_clauses): Handle OMP_CLAUSE_INIT,
	OMP_CLAUSE_USE and OMP_CLAUSE_DESTROY.
	(scan_omp_1_stmt): Handle GIMPLE_OMP_INTEROP.
	(lower_omp_interop_action_clauses): New function.
	(lower_omp_interop): Likewise.
	(lower_omp_1): Handle GIMPLE_OMP_INTEROP.

gcc/c/ChangeLog:

	* c-parser.cc (c_parser_omp_clause_destroy): Make addressable.
	(c_parser_omp_clause_init): Make addressable.

gcc/cp/ChangeLog:

	* parser.cc (cp_parser_omp_clause_init): Make addressable.

gcc/fortran/ChangeLog:

	* trans-openmp.cc (gfc_trans_omp_clauses): Make OMP_CLAUSE_DESTROY and
	OMP_CLAUSE_INIT addressable.
	* types.def (BT_FN_VOID_INT_INT_PTR_PTR_PTR_INT_PTR_INT_PTR_UINT_PTR):
	New.

include/ChangeLog:

	* gomp-constants.h (GOMP_DEVICE_DEFAULT_OMP_61, GOMP_INTEROP_TARGET,
	GOMP_INTEROP_TARGETSYNC, GOMP_INTEROP_FLAG_NOWAIT): Define.

libgomp/ChangeLog:

	* icv-device.c (omp_set_default_device): Check
	GOMP_DEVICE_DEFAULT_OMP_61.
	* libgomp-plugin.h (struct interop_obj_t): New.
	(enum gomp_interop_flag): New.
	(GOMP_OFFLOAD_interop): Declare.
	(GOMP_OFFLOAD_get_interop_int): Declare.
	(GOMP_OFFLOAD_get_interop_ptr): Declare.
	(GOMP_OFFLOAD_get_interop_str): Declare.
	(GOMP_OFFLOAD_get_interop_type_desc): Declare.
	* libgomp.h (_LIBGOMP_OMP_LOCK_DEFINED): Define.
	(struct gomp_device_descr): Add interop_func, get_interop_int_func,
	get_interop_ptr_func, get_interop_str_func, get_interop_type_desc_func.
	* libgomp.map: Add GOMP_interop.
	* libgomp_g.h (GOMP_interop): Declare.
	* target.c (resolve_device): Handle GOMP_DEVICE_DEFAULT_OMP_61.
	(omp_get_interop_int): Replace stub with actual implementation.
	(omp_get_interop_ptr): Likewise.
	(omp_get_interop_str): Likewise.
	(omp_get_interop_type_desc): Likewise.
	(struct interop_data_t): Define.
	(gomp_interop_internal): New function.
	(GOMP_interop): Likewise.
	(gomp_load_plugin_for_device): Load symbols for get_interop_int,
	get_interop_ptr, get_interop_str and get_interop_type_desc.
	* testsuite/libgomp.c-c++-common/interop-1.c: New test.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/interop-1.c: Remove dg-prune-output "sorry".
	* c-c++-common/gomp/interop-2.c: Likewise.
	* c-c++-common/gomp/interop-3.c: Likewise.
	* c-c++-common/gomp/interop-4.c: Remove dg-message "not supported".
	* g++.dg/gomp/interop-5.C: Likewise.
	* gfortran.dg/gomp/interop-4.f90: Likewise.
	* c-c++-common/gomp/interop-5.c: New test.
	* gfortran.dg/gomp/interop-5.f90: New test.

Co-authored-by: Tobias Burnus <tburnus@baylibre.com>
2025-03-21 19:24:16 +01:00
Tobias Burnus
12db61156f testsuite/lib/libgomp.exp: compile with -fdiagnostics-plain-output
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.
2025-03-21 13:54:49 +01:00
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.
2025-03-17 10:12:44 +01:00
Jason Merrill
556248d7d2 c++: don't default -frange-for-ext-temps in -std=gnu++20 [PR188574]
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.
2025-02-12 00:07:51 +01:00
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>
2025-01-30 19:12:34 +00:00
Sandra Loosemore
fdeceba59b OpenMP: Shared metadirective/dynamic selector tests for C and C++
gcc/testsuite/ChangeLog
	* c-c++-common/gomp/adjust-args-6.c: New.
	* c-c++-common/gomp/attrs-metadirective-1.c: New.
	* c-c++-common/gomp/attrs-metadirective-2.c: New.
	* c-c++-common/gomp/attrs-metadirective-3.c: New.
	* c-c++-common/gomp/attrs-metadirective-4.c: New.
	* c-c++-common/gomp/attrs-metadirective-5.c: New.
	* c-c++-common/gomp/attrs-metadirective-6.c: New.
	* c-c++-common/gomp/attrs-metadirective-7.c: New.
	* c-c++-common/gomp/attrs-metadirective-8.c: New.
	* c-c++-common/gomp/declare-variant-arg-exprs.c: New.
	* c-c++-common/gomp/declare-variant-dynamic-1.c: New.
	* c-c++-common/gomp/declare-variant-dynamic-2.c: New.
	* c-c++-common/gomp/metadirective-1.c: New.
	* c-c++-common/gomp/metadirective-2.c: New.
	* c-c++-common/gomp/metadirective-3.c: New.
	* c-c++-common/gomp/metadirective-4.c: New.
	* c-c++-common/gomp/metadirective-5.c: New.
	* c-c++-common/gomp/metadirective-6.c: New.
	* c-c++-common/gomp/metadirective-7.c: New.
	* c-c++-common/gomp/metadirective-8.c: New.
	* c-c++-common/gomp/metadirective-construct.c: New.
	* c-c++-common/gomp/metadirective-device.c: New.
	* c-c++-common/gomp/metadirective-no-score.c: New.
	* c-c++-common/gomp/metadirective-target-device-1.c: New.
	* c-c++-common/gomp/metadirective-target-device-2.c: New.

libgomp/ChangeLog
	* testsuite/libgomp.c-c++-common/metadirective-1.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-2.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-3.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-4.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-5.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-late-1.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-late-2.c: New.
	* testsuite/libgomp.c-c++-common/metadirective-target-device.c: New.

Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
Co-Authored-By: Sandra Loosemore <sandra@codesourcery.com>
2025-01-16 18:12:21 +00:00
Sandra Loosemore
677e452e55 OpenMP: C++ support for metadirectives and dynamic selectors.
Additional shared C/C++ testcases are included in a subsequent patch in this
series.

gcc/cp/ChangeLog
	PR middle-end/112779
	PR middle-end/113904
	* cp-tree.h (struct saved_scope): Add new field
	x_processing_omp_trait_property_expr.
	(processing_omp_trait_property_expr): New.
	* parser.cc (cp_parser_skip_to_end_of_block_or_statement): Add
	metadirective_p	parameter and handle skipping over the parentheses
	in a "for" statement.
	(struct omp_metadirective_parse_data): New.
	(mangle_metadirective_region_label): New.
	(cp_parser_label_for_labeled_statement): Mangle label names in a
	metadirective body.
	(cp_parser_jump_statement): Likewise.
	(cp_parser_omp_context_selector): Allow arbitrary expressions in
	device_num and condition properties.
	(cp_parser_omp_assumption_clauses): Handle C_OMP_DIR_META.
	(analyze_metadirective_body): New.
	(cp_parser_omp_metadirective): New.
	(cp_parser_pragma): Handle PRAGMA_OMP_METADIRECTIVE.
	* parser.h (struct cp_parser): Add omp_metadirective_state field.
	* pt.cc (tsubst_omp_context_selector): New.
	(tsubst_stmt): Handle OMP_METADIRECTIVE.
	* semantics.cc (finish_id_expression_1): Don't diagnose use of
	parameter outside function body in dynamic selector expressions here.

	gcc/testsuite/
	PR middle-end/112779
	PR middle-end/113904
	* c-c++-common/gomp/declare-variant-2.c: Adjust output for C++.
	* g++.dg/gomp/declare-variant-class-1.C: New.
	* g++.dg/gomp/declare-variant-class-2.C: New.
	* g++.dg/gomp/metadirective-template-1.C: New.

	libgomp/
	PR middle-end/112779
	PR middle-end/113904
	* testsuite/libgomp.c++/metadirective-template-1.C: New.
	* testsuite/libgomp.c++/metadirective-template-2.C: New.
	* testsuite/libgomp.c++/metadirective-template-3.C: New.

Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
Co-Authored-By: Sandra Loosemore <sandra@codesourcery.com>
2025-01-16 18:12:17 +00:00
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.
2025-01-13 16:05:14 +01:00
Thomas Schwinge
3861d362ec nvptx: PTX 'alloca' for '-mptx=7.3'+, '-march=sm_52'+ [PR65181]
..., and use it for '-mno-soft-stack': PTX "native" stacks.

	PR target/65181
	gcc/
	* config/nvptx/nvptx.cc (nvptx_get_drap_rtx): Handle
	'!TARGET_SOFT_STACK'.
	* config/nvptx/nvptx.md (define_c_enum "unspec"): Add
	'UNSPEC_STACKSAVE', 'UNSPEC_STACKRESTORE'.
	(define_expand "allocate_stack", define_expand "save_stack_block")
	(define_expand "save_stack_block"): Handle '!TARGET_SOFT_STACK',
	PTX 'alloca'.
	(define_insn "@nvptx_alloca_<mode>")
	(define_insn "@nvptx_stacksave_<mode>")
	(define_insn "@nvptx_stackrestore_<mode>"): New.
	* doc/invoke.texi (Nvidia PTX Options): Update '-msoft-stack',
	'-mno-soft-stack'.
	* doc/sourcebuild.texi (nvptx-specific attributes): Document
	'nvptx_runtime_alloca_ptx'.
	(Add Options): Document 'nvptx_alloca_ptx'.
	gcc/testsuite/
	* gcc.target/nvptx/alloca-1.c: Evolve into...
	* gcc.target/nvptx/alloca-1-O0.c: ... this, ...
	* gcc.target/nvptx/alloca-1-O1.c: ... this, and...
	* gcc.target/nvptx/alloca-1-sm_30.c: ... this.
	* gcc.target/nvptx/vla-1.c: Evolve into...
	* gcc.target/nvptx/vla-1-O0.c: ... this, ...
	* gcc.target/nvptx/vla-1-O1.c: ... this, and...
	* gcc.target/nvptx/vla-1-sm_30.c: ... this.
	* gcc.c-torture/execute/pr36321.c: Adjust.
	* gcc.target/nvptx/__builtin_alloca_0-1-O0.c: Likewise.
	* gcc.target/nvptx/__builtin_alloca_0-1-O1.c: Likewise.
	* gcc.target/nvptx/__builtin_stack_save___builtin_stack_restore-1.c:
	Likewise.
	* gcc.target/nvptx/softstack.c: Likewise.
	* gcc.target/nvptx/__builtin_stack_save___builtin_stack_restore-1-sm_30.c:
	New.
	* gcc.target/nvptx/alloca-2-O0.c: Likewise.
	* gcc.target/nvptx/alloca-3-O1.c: Likewise.
	* gcc.target/nvptx/alloca-4-O3.c: Likewise.
	* gcc.target/nvptx/alloca-5.c: Likewise.
	* lib/target-supports.exp (check_effective_target_alloca): Adjust.
	(check_nvptx_default_ptx_isa_target_architecture_at_least)
	(check_nvptx_runtime_ptx_isa_target_architecture_at_least)
	(check_effective_target_nvptx_runtime_alloca_ptx)
	(add_options_for_nvptx_alloca_ptx): New.
	libgomp/
	* fortran.c (omp_get_device_from_uid_): Adjust.
	* testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Likewise.
2025-01-09 14:17:11 +01:00
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.
2025-01-02 21:18:56 +01:00
Jakub Jelinek
6441eb6dc0 Update copyright years. 2025-01-02 11:59:57 +01:00
Thomas Schwinge
ed210c69ce nvptx: Support '-march=sm_89'
gcc/
	* config/nvptx/nvptx-sm.def: Add '89'.
	* config/nvptx/nvptx-gen.h: Regenerate.
	* config/nvptx/nvptx-gen.opt: Likewise.
	* config/nvptx/nvptx.cc (first_ptx_version_supporting_sm): Adjust.
	* config/nvptx/nvptx.opt (-march-map=sm_89, -march-map=sm_90)
	(march-map=sm_90a): Likewise.
	* config.gcc: Likewise.
	* doc/invoke.texi (Nvidia PTX Options): Document '-march=sm_89'.
	* config/nvptx/gen-multilib-matches-tests: Extend.
	gcc/testsuite/
	* gcc.target/nvptx/march-map=sm_89.c: Adjust.
	* gcc.target/nvptx/march-map=sm_90.c: Likewise.
	* gcc.target/nvptx/march-map=sm_90a.c: Likewise.
	* gcc.target/nvptx/march=sm_89.c: New.
	libgomp/
	* testsuite/libgomp.c/declare-variant-3-sm89.c: New.
	* testsuite/libgomp.c/declare-variant-3.h: Adjust.
2024-12-06 12:30:20 +01:00
Thomas Schwinge
7cff338e5a nvptx: Support '-march=sm_52'
gcc/
	* config/nvptx/nvptx-sm.def: Add '52'.
	* config/nvptx/nvptx-gen.h: Regenerate.
	* config/nvptx/nvptx-gen.opt: Likewise.
	* config/nvptx/nvptx.cc (first_ptx_version_supporting_sm): Adjust.
	* config/nvptx/nvptx.opt (-march-map=sm_52): Likewise.
	* config.gcc: Likewise.
	* doc/invoke.texi (Nvidia PTX Options): Document '-march=sm_52'.
	* config/nvptx/gen-multilib-matches-tests: Extend.
	gcc/testsuite/
	* gcc.target/nvptx/march-map=sm_52.c: Adjust.
	* gcc.target/nvptx/march=sm_52.c: New.
	libgomp/
	* testsuite/libgomp.c/declare-variant-3-sm52.c: New.
	* testsuite/libgomp.c/declare-variant-3.h: Adjust.
2024-12-06 12:26:45 +01:00
Thomas Schwinge
7151d885c4 nvptx: Support '-march=sm_37'
gcc/
	* config/nvptx/nvptx-sm.def: Add '37'.
	* config/nvptx/nvptx-gen.h: Regenerate.
	* config/nvptx/nvptx-gen.opt: Likewise.
	* config/nvptx/nvptx.cc (first_ptx_version_supporting_sm): Adjust.
	* config/nvptx/nvptx.opt (-march-map=sm_37, -march-map=sm_50):
	Likewise.
	* config.gcc: Likewise.
	* doc/invoke.texi (Nvidia PTX Options): Document '-march=sm_37'.
	* config/nvptx/gen-multilib-matches-tests: Extend.
	gcc/testsuite/
	* gcc.target/nvptx/march-map=sm_37.c: Adjust.
	* gcc.target/nvptx/march-map=sm_50.c: Likewise.
	* gcc.target/nvptx/march-map=sm_52.c: Likewise.
	* gcc.target/nvptx/march=sm_37.c: New.
	libgomp/
	* testsuite/libgomp.c/declare-variant-3-sm37.c: New.
	* testsuite/libgomp.c/declare-variant-3.h: Adjust.
2024-12-06 12:26:45 +01:00
Thomas Schwinge
ab5bd6ac68 Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' __builtin_is_initial_device: Revert 'gimple_fold_builtin_acc_on_device' change
The motivation of the 'gimple_fold_builtin_acc_on_device' change in
commit 3269a722b7
"Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' __builtin_is_initial_device"
is unclear, and it unnecessarily diverges GCC's (default)
'--disable-offload-targets' vs. '--enable-offload-targets=[...]'
configurations.

	PR testsuite/82250
	gcc/
	* gimple-fold.cc (gimple_fold_builtin_acc_on_device): Revert last
	change.
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: Revert
	last change.
2024-12-06 09:48:33 +01:00
Tobias Burnus
3d72e50cae OpenMP: 'allocate' directive - fixes for 'alignof' and [[omp::decl]]
Fixed a check to permit [[omp::decl(allocate,...)]] parsing in C.

Additionaly, we discussed that 'allocate align' should not affect
'alignof' to avoid issues like with:

  int a;
  _Alignas(_Alignof(a)) int b;
  #pragma omp allocate(a) align(128)
  _Alignas(_Alignof(a)) int c;

Thus, the alignment is no longer set in the C and Fortran front ends,
but for static variables now in varpool_node::finalize_decl.
(For stack variables, the alignment is handled in gimplify_bind_expr.)

NOTE: 'omp allocate' is not yet supported in C++.

gcc/c/ChangeLog:

	* c-parser.cc (c_parser_omp_allocate): Only check scope if
	not in_omp_decl_attribute. Remove setting the alignment.

gcc/ChangeLog:

	* cgraphunit.cc (varpool_node::finalize_decl): Set alignment
	based on OpenMP's 'omp allocate' attribute/directive.

gcc/fortran/ChangeLog:

	* trans-decl.cc (gfc_finish_var_decl): Remove setting the alignment.

libgomp/ChangeLog:

	* libgomp.texi (Memory allocation): Mention (non-)effect of 'align'
	on _Alignof.
	* testsuite/libgomp.c/allocate-7.c: New test.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/allocate-18.c: Check that alignof is unaffected
	by 'omp allocate'.
	* c-c++-common/gomp/allocate-19.c: Likewise.
2024-12-03 11:02:03 +01:00
Thomas Schwinge
bcb764ec7c Fix 'libgomp.oacc-c/../libgomp.oacc-c-c++-common/acc_get_property-gcn.c' for C23 default
With commit 55e3bd376b "c: Default to -std=gnu23"
we've got:

    [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/acc_get_property-gcn.c -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa  -O0  (test for excess errors)
    [-PASS:-]{+UNRESOLVED:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/acc_get_property-gcn.c -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa  -O0  [-execution test-]{+compilation failed to produce executable+}
    [Etc.]

..., due to:

    [...]/libgomp.oacc-c/../libgomp.oacc-c-c++-common/acc_get_property-gcn.c:16:13: error: two or more data types in declaration specifiers
    [...]/libgomp.oacc-c/../libgomp.oacc-c-c++-common/acc_get_property-gcn.c:16:1: warning: useless type name in empty declaration

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c
	[!__cplusplus]: Don't 'typedef int bool;'.
2024-11-28 15:17:36 +01:00
Rainer Orth
0f7def8549 libgomp: testsuite: Fix libgomp.c/alloc-pinned-3.c etc. for C23 on non-Linux
Since the switch to a C23 default, three libgomp tests FAIL on Solaris:

FAIL: libgomp.c/alloc-pinned-3.c (test for excess errors)
UNRESOLVED: libgomp.c/alloc-pinned-3.c compilation failed to produce executable
FAIL: libgomp.c/alloc-pinned-4.c (test for excess errors)
UNRESOLVED: libgomp.c/alloc-pinned-4.c compilation failed to produce executable
FAIL: libgomp.c/alloc-pinned-6.c (test for excess errors)
UNRESOLVED: libgomp.c/alloc-pinned-6.c compilation failed to produce executable

Excess errors:
/vol/gcc/src/hg/master/local/libgomp/testsuite/libgomp.c/alloc-pinned-3.c:104:3: error: too many arguments to function 'set_pin_limit'

Fixed by adding the missing size argument to the stub functions.

Tested on i386-pc-solaris2.11 and sparc-sun-solaris2.11.

2024-11-20  Rainer Orth  <ro@CeBiTec.Uni-Bielefeld.DE>

	libgomp:
	* testsuite/libgomp.c/alloc-pinned-3.c [!__linux__]
	(set_pin_limit): Add size arg.
	* testsuite/libgomp.c/alloc-pinned-4.c [!__linux__]
	(set_pin_limit): Likewise.
	* testsuite/libgomp.c/alloc-pinned-6.c [!__linux__]
	(set_pin_limit): Likewise.
2024-11-21 11:46:36 +01:00
Paul-Antoine Arras
377eff7c38 OpenMP: common C/C++ testcases for dispatch + adjust_args
gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/declare-variant-2.c: Adjust dg-error directives.
	* c-c++-common/gomp/adjust-args-1.c: New test.
	* c-c++-common/gomp/adjust-args-2.c: New test.
	* c-c++-common/gomp/declare-variant-dup-match-clause.c: New test.
	* c-c++-common/gomp/dispatch-1.c: New test.
	* c-c++-common/gomp/dispatch-2.c: New test.
	* c-c++-common/gomp/dispatch-3.c: New test.
	* c-c++-common/gomp/dispatch-4.c: New test.
	* c-c++-common/gomp/dispatch-5.c: New test.
	* c-c++-common/gomp/dispatch-6.c: New test.
	* c-c++-common/gomp/dispatch-7.c: New test.
	* c-c++-common/gomp/dispatch-8.c: New test.
	* c-c++-common/gomp/dispatch-9.c: New test.
	* c-c++-common/gomp/dispatch-10.c: New test.

libgomp/ChangeLog:

	* testsuite/libgomp.c-c++-common/dispatch-1.c: New test.
	* testsuite/libgomp.c-c++-common/dispatch-2.c: New test.
2024-11-20 15:31:22 +01:00
David Malcolm
c9d21e19df testsuite: move dg-test cleanup code from gcc-dg.exp to its own file
I need to use this cleanup logic for the testsuite for libdiagnostics
where it's too awkward to directly use gcc-dg.exp itself.

No functional change intended.

gcc/testsuite/ChangeLog:
	* lib/dg-test-cleanup.exp: New file, from material moved from
	lib/gcc-dg.exp.
	* lib/gcc-dg.exp: Add load_lib of dg-test-cleanup.exp.
	(cleanup-after-saved-dg-test): Move to lib/dg-test-cleanup.exp.
	(dg-test): Likewise for override.
	(initialize_prune_notes): Likewise.

libatomic/ChangeLog:
	* testsuite/lib/libatomic.exp: Add
	"load_gcc_lib dg-test-cleanup.exp".

libgomp/ChangeLog:
	* testsuite/lib/libgomp.exp: Add
	"load_gcc_lib dg-test-cleanup.exp".
libitm/ChangeLog:
	* testsuite/lib/libitm.exp: Add
	"load_gcc_lib dg-test-cleanup.exp".

libphobos/ChangeLog:
	* testsuite/lib/libphobos-dg.exp: Add
	"load_gcc_lib dg-test-cleanup.exp".

libstdc++-v3/ChangeLog:
	* testsuite/lib/libstdc++.exp: Add
	"load_gcc_lib dg-test-cleanup.exp".

libvtv/ChangeLog:
	* testsuite/lib/libvtv.exp: Add
	"load_gcc_lib dg-test-cleanup.exp".

Signed-off-by: David Malcolm <dmalcolm@redhat.com>
2024-11-18 16:49:04 -05:00
Tobias Burnus
5114789038 libgomp.c-c++-common/pr109062.c: Fix expected spin count for hybrid x86
On my system with E and P cores (hybrid) x86, the spincount is by default 1
and not 300000, cf. PR109812 and r14-4571-ge1e127de18dbee.

Hence, this commit updates the expected value of the testcase to also
accept omp_display_env showing "GOMP_SPINCOUNT = '1'" - but only for
x86-64, which might be hybrid.

libgomp/ChangeLog:

	* testsuite/libgomp.c-c++-common/pr109062.c: Update dg-output
	to also accept GOMP_SPINCOUNT = 1 for x86-64.
2024-11-11 17:22:20 +01:00
Thomas Schwinge
730f28b081 Adjust 'libgomp.c/max_vf-*.c'
For configurations where both GCN and nvptx offloading are enabled, we get:

    PASS: libgomp.c/max_vf-1.c (test for excess errors)
    PASS: libgomp.c/max_vf-1.c scan-tree-dump-times ompexp "GOMP_MAX_VF" 2
    PASS: libgomp.c/max_vf-1.c scan-tree-dump-times ompexp "__builtin_GOMP_parallel_loop_nonmonotonic_dynamic \\(.*, D\\.[0-9]*, 0\\);" 1
    PASS: libgomp.c/max_vf-1.c scan-amdgcn-amdhsa-offload-tree-dump-times optimized "__builtin_GOMP_parallel_loop_nonmonotonic_dynamic \\(.*, 64, 0\\);" 1
    FAIL: libgomp.c/max_vf-1.c scan-nvptx-none-offload-tree-dump-times optimized "__builtin_GOMP_parallel_loop_nonmonotonic_dynamic \\(.*, 64, 0\\);" 1
    FAIL: libgomp.c/max_vf-1.c scan-amdgcn-amdhsa-offload-tree-dump-times optimized "__builtin_GOMP_parallel_loop_nonmonotonic_dynamic \\(.*, 7, 0\\);" 1
    PASS: libgomp.c/max_vf-1.c scan-nvptx-none-offload-tree-dump-times optimized "__builtin_GOMP_parallel_loop_nonmonotonic_dynamic \\(.*, 7, 0\\);" 1

Avoid these FAILs via 'only_for_offload_target [...]'.  Also, for consistency
with other libgomp test cases, use effective-target specifiers of the libgomp
test suite.  Fix-up for recent commit d334f729e5
"openmp: Add testcases for omp_max_vf".

	libgomp/
	* testsuite/libgomp.c/max_vf-1.c: Adjust.
	* testsuite/libgomp.c/max_vf-2.c: Likewise.
2024-11-10 11:19:02 +01:00
Andrew Stubbs
4e91d05872 openmp: Fix max_vf testcases with -march=cascadelake
Apparently we need to explicitly disable AVX, not just enabled SSE, to
guarentee the 16-lane vectors we need for the pattern match.

libgomp/ChangeLog:

	* testsuite/libgomp.c/max_vf-1.c: Add -mno-avx.

gcc/testsuite/ChangeLog:

	* gcc.dg/gomp/max_vf-1.c: Add -mno-avx.
2024-11-07 12:00:53 +00:00
Andrew Stubbs
d334f729e5 openmp: Add testcases for omp_max_vf
Ensure that the GOMP_MAX_VF does the right thing for explicit schedules, when
offloading is enabled ("target" directives are present), and is inactive
otherwise.

libgomp/ChangeLog:

	* testsuite/libgomp.c/max_vf-1.c: New test.
	* testsuite/libgomp.c/max_vf-2.c: New test.

gcc/testsuite/ChangeLog:

	* gcc.dg/gomp/max_vf-1.c: New test.
2024-11-06 16:26:24 +00:00
Tobias Burnus
ee4fdda70f Add libgomp.oacc-fortran/acc_on_device-1-4.f
Kind of undoes r15-4315-g9f549d216c9716 by adding the original testcase back;
namely, adding acc_on_device-1-3.f as acc_on_device-1-4.f with
-fno-builtin-acc_on_device removed.

libgomp/ChangeLog:

	* testsuite/libgomp.oacc-fortran/acc_on_device-1-4.f: New test;
	same as acc_on_device-1-3.f but using the builtin function.
2024-10-16 16:15:40 +02:00
Thomas Schwinge
de0320712d OpenACC 'nohost' clause: harmonize 'libgomp.oacc-{c-c++-common,fortran}/routine-nohost-1.*'
The test case 'libgomp.oacc-fortran/routine-nohost-1.f90' added in 2021
commit a61f6afbee "OpenACC 'nohost' clause" was
dependend on inlining being enabled, and otherwise ('-fno-inline') failed to
optimize/link:

    /tmp/ccb2hsPd.o: In function `MAIN__._omp_fn.0':
    routine-nohost-1.f90:(.text+0xf4): undefined reference to `fact_nohost_'

However, as of recent commit 3269a722b7
"Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' __builtin_is_initial_device",
we're now properly handling OpenACC/Fortran 'acc_on_device', and may specify
'-fno-inline', like done in 'libgomp.oacc-c-c++-common/routine-nohost-1.c'.

	libgomp/
	* testsuite/libgomp.oacc-fortran/routine-nohost-1.f90: Add
	'-fno-inline'.
2024-10-14 18:37:48 +02:00
Thomas Schwinge
9f549d216c Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' __builtin_is_initial_device: Harmonize 'libgomp.oacc-fortran/acc_on_device-1-*'
The test case 'libgomp.oacc-fortran/acc_on_device-1-1.f90' added in
commit 3269a722b7
"Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' __builtin_is_initial_device"
was missing '-fno-builtin-acc_on_device', and all
'libgomp.oacc-fortran/acc_on_device-1-*' need comments, why that option is
specified.

	PR testsuite/82250
	libgomp/
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Add
	'-fno-builtin-acc_on_device'.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Comment.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Comment.
2024-10-14 10:34:34 +02:00