aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite
AgeCommit message (Collapse)AuthorFilesLines
37 hoursRemove 'ALWAYS_INLINE' workaround in ↵Thomas Schwinge1-6/+0
'libgomp.c++/target-exceptions-pr118794-1.C' With commit ca9cffe737d20953082333dacebb65d4261e0d0c "For nvptx offloading, make sure to emit C++ constructor, destructor aliases [PR97106]", we're able to remove the 'ALWAYS_INLINE' workaround added in commit fe283dba774be57b705a7a871b000d2894d2e553 "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.
40 hoursAdd 'libgomp.c++/pr106445-1{,-O0}.C' [PR106445]Thomas Schwinge2-0/+21
PR target/106445 libgomp/ * testsuite/libgomp.c++/pr106445-1.C: New. * testsuite/libgomp.c++/pr106445-1-O0.C: Likewise.
40 hoursFor nvptx offloading, make sure to emit C++ constructor, destructor aliases ↵Thomas Schwinge2-3/+1
[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.
2 daysOpenMP: omp.h omp::allocator C++ Allocator interfacewaffl3x2-0/+290
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>
3 daysFortran/OpenMP: Support automatic mapping allocatable components (deep mapping)Tobias Burnus8-0/+2158
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.
3 daysGCN, nvptx: Support '-mfake-exceptions', and use it for offloading ↵Thomas Schwinge16-54/+140
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.
3 daysAdd 'throw', dead code test cases for GCN, nvptx target and OpenACC, OpenMP ↵Thomas Schwinge2-0/+62
'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.
3 daysAdd 'throw', caught test cases for GCN, nvptx target and OpenACC, OpenMP ↵Thomas Schwinge6-0/+134
'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.
3 daysAdd 'throw' test cases for GCN, nvptx target and OpenACC, OpenMP 'target' ↵Thomas Schwinge3-0/+94
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.
3 daysAdd 'std::bad_cast' exception, dead code test cases for GCN, nvptx target ↵Thomas Schwinge2-0/+66
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.
3 daysAdd 'std::bad_cast' exception, caught test cases for GCN, nvptx target and ↵Thomas Schwinge6-0/+134
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.
3 daysAdd 'std::bad_cast' exception test cases for GCN, nvptx target and OpenACC, ↵Thomas Schwinge2-0/+79
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.
3 daysAdd test cases for exception handling constructs in dead code for GCN, nvptx ↵Thomas Schwinge3-0/+104
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.
3 daysAdd PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' ↵Thomas Schwinge8-0/+117
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.
7 dayslibgomp: Update SVE testsRichard Sandiford8-2/+16
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.
10 daysOpenMP: Fix append_args handling in modify_call_for_omp_dispatchTobias Burnus2-0/+537
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.
10 dayslibgomp: Add AArch64 SVE target tests to libgomp.Tejas Belagod9-0/+963
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-03-26driver: Forward '-lstdc++' to offloading compilation [PR101544]Thomas Schwinge3-3/+0
..., 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 4bcb46b3ade1796c5a57b294f5cca25f00671cac "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-24Add 'libgomp.c++/pr96835-1{,-O0}.C', 'libgomp.oacc-c++/pr96835-1.C' [PR96835]Thomas Schwinge3-0/+54
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-24Add 'libgomp.c++/pr101544-1{,-O0}.C', 'libgomp.oacc-c++/pr101544-1.C' [PR101544]Thomas Schwinge3-0/+94
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-22libgomp.fortran/get-mapped-ptr-1.f90: Use -6 for non-conf dev numberTobias Burnus1-1/+1
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-21libgomp/plugin: Add initial interop support to nvptx + gcnTobias Burnus2-1/+578
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-21OpenMP: 'interop' construct - add ME support + target-independent libgompPaul-Antoine Arras1-0/+43
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-21testsuite/lib/libgomp.exp: compile with -fdiagnostics-plain-outputTobias Burnus1-5/+2
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-17Move gfortran.dg/gomp/declare-variant-mod-1*.f90 to libgomp.fortran/ [PR115271]Tobias Burnus2-0/+164
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-02-12c++: don't default -frange-for-ext-temps in -std=gnu++20 [PR188574]Jason Merrill1-1/+1
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-01-30OpenMP: Fortran support for metadirectives and dynamic selectorsSandra Loosemore6-0/+276
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-16OpenMP: Shared metadirective/dynamic selector tests for C and C++Sandra Loosemore8-0/+424
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-16OpenMP: C++ support for metadirectives and dynamic selectors.Sandra Loosemore3-0/+125
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-13Add missing target directive in OpenMP dispatch Fortran runtime testPaul-Antoine Arras1-8/+13
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-09nvptx: PTX 'alloca' for '-mptx=7.3'+, '-march=sm_52'+ [PR65181]Thomas Schwinge1-10/+0
..., 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-02OpenMP: Fortran front-end support for dispatch + adjust_argsPaul-Antoine Arras5-0/+316
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-02Update copyright years.Jakub Jelinek4-4/+4
2024-12-06nvptx: Support '-march=sm_89'Thomas Schwinge2-0/+16
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-06nvptx: Support '-march=sm_52'Thomas Schwinge2-0/+16
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-06nvptx: Support '-march=sm_37'Thomas Schwinge2-0/+16
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-06Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' ↵Thomas Schwinge1-1/+2
__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 3269a722b7a03613e9c4e2862bc5088c4a17cc11 "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-03OpenMP: 'allocate' directive - fixes for 'alignof' and [[omp::decl]]Tobias Burnus1-0/+54
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-11-28Fix 'libgomp.oacc-c/../libgomp.oacc-c-c++-common/acc_get_property-gcn.c' for ↵Thomas Schwinge1-3/+0
C23 default With commit 55e3bd376b2214e200fa76d12b67ff259b06c212 "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-21libgomp: testsuite: Fix libgomp.c/alloc-pinned-3.c etc. for C23 on non-LinuxRainer Orth3-3/+3
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-20OpenMP: common C/C++ testcases for dispatch + adjust_argsPaul-Antoine Arras2-0/+160
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-18testsuite: move dg-test cleanup code from gcc-dg.exp to its own fileDavid Malcolm1-0/+1
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-11libgomp.c-c++-common/pr109062.c: Fix expected spin count for hybrid x86Tobias Burnus1-1/+5
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-10Adjust 'libgomp.c/max_vf-*.c'Thomas Schwinge2-5/+5
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 d334f729e53867b838e867375b3f475ba793d96e "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-07openmp: Fix max_vf testcases with -march=cascadelakeAndrew Stubbs1-1/+1
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-06openmp: Add testcases for omp_max_vfAndrew Stubbs2-0/+68
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-10-16Add libgomp.oacc-fortran/acc_on_device-1-4.fTobias Burnus1-0/+60
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-14OpenACC 'nohost' clause: harmonize ↵Thomas Schwinge1-0/+2
'libgomp.oacc-{c-c++-common,fortran}/routine-nohost-1.*' The test case 'libgomp.oacc-fortran/routine-nohost-1.f90' added in 2021 commit a61f6afbee370785cf091fe46e2e022748528307 "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 3269a722b7a03613e9c4e2862bc5088c4a17cc11 "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-14Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' ↵Thomas Schwinge3-2/+11
__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 3269a722b7a03613e9c4e2862bc5088c4a17cc11 "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-14Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' ↵Thomas Schwinge1-1/+1
__builtin_is_initial_device: Fix effective-target keyword in 'libgomp.oacc-fortran/acc_on_device-2.f90' The test case 'libgomp.oacc-fortran/acc_on_device-2.f90' added in commit 3269a722b7a03613e9c4e2862bc5088c4a17cc11 "Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' __builtin_is_initial_device" had a mismatch between dump file production and its scanning; the former needs to use 'offload_target_nvptx' (like 'offload_target_amdgcn'), not 'offload_device_nvptx'. PR testsuite/82250 libgomp/ * testsuite/libgomp.oacc-fortran/acc_on_device-2.f90: Fix effective-target keyword.
generated by cgit v1.1 at 2025-04-18 05:44:59 +0000