Age | Commit message (Collapse) | Author | Files | Lines |
|
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.
|
|
'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.
|
|
'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.
|
|
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.
|
|
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.
|
|
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.
|
|
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.
|
|
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.
|
|
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.
|
|
The new SVE tests didn't explicitly force SVE to be enabled,
which meant that they wouldn't work on targets that aren't
configured for SVE by default. The least invasive way of
fixing that is to add a pragma, which works for most tests.
However, for udr-sve.c, the global:
#pragma omp declare reduction (+:svint32_t: omp_out = svadd_s32_z (svptrue_b32(), omp_in, omp_out)) \
initializer (omp_priv = svindex_s32 (0, 0))
does not work with an earlier:
#pragma GCC target "+sve"
which is interesting, and maybe worthy of a PR if there isn't one
already. It seems we have to force SVE (and thus an architecture)
on the command line instead.
However, with that fixed, udr-sve.c fails execution. One problem
seems to be a missing accumulation in for_reduction. Fixing that
is enough to reach the final inscan_reduction_incl, but that fails
for reasons I haven't investigated yet. I would need to read up
more to understand what the loop is doing.
It also looks like there might be a missing "+" in simd_reduction:
#pragma omp simd reduction (+:va, i)
for (j = 0; j < 16; j++)
va = svld1_s32 (svptrue_b32 (), a);
res = svaddv_s32 (svptrue_b32 (), va);
if (res != 8)
__builtin_abort ();
since AFAICT the loop is not doing a reduction as things stand.
But perhaps that's deliberate, since it does match the != 8 test.
libgomp/
* testsuite/libgomp.c-target/aarch64/firstprivate.c: Add +sve pragma.
* testsuite/libgomp.c-target/aarch64/lastprivate.c: Likewise.
* testsuite/libgomp.c-target/aarch64/private.c: Likewise.
* testsuite/libgomp.c-target/aarch64/shared.c: Likewise.
* testsuite/libgomp.c-target/aarch64/simd-aligned.c: Likewise.
* testsuite/libgomp.c-target/aarch64/simd-nontemporal.c: Likewise.
* testsuite/libgomp.c-target/aarch64/threadprivate.c: Likewise.
* testsuite/libgomp.c-target/aarch64/udr-sve.c: Add an -march option.
(for_reduction): Use "+=" in the reduction loop.
|
|
At tree level, the addr ref is also required for array dummy arguments,
contrary to C; the GOMP_interop calls in modify_call_for_omp_dispatch
were updated accordingly (using build_fold_addr_expr).
As the GOMP_interop calls had no location data associated with them,
the init call happened as soon as executing the previous line of code,
which was confusing; solution: use the location data of the function
call itself.
PR middle-end/119662
gcc/ChangeLog:
* gimplify.cc (modify_call_for_omp_dispatch): Fix GOMP_interop
arg passing; add location info to function calls.
libgomp/ChangeLog:
* testsuite/libgomp.c/append-args-fr-1.c: New test.
* testsuite/libgomp.c/append-args-fr.h: New test.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/append-args-interop.c: Update for fixed
GOMP_interop call.
* g++.dg/gomp/append-args-8.C: Likewise.
* gfortran.dg/gomp/append-args-interop.f90: Likewise.
|
|
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.
|
|
..., 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.
|
|
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.
|
|
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.
|
|
This is a fix for the GOMP_interop commit r15-8654-g99e2906ae255fc that
added GOMP_DEVICE_DEFAULT_OMP_61 alias omp_default_device, which is a
conforming device number. But that test used -5 as check for a
non-conforming device number.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/get-mapped-ptr-1.f90: Use -6
not -5 as non-conforming device number.
|
|
The interop directive operates on an opaque object that represents a
foreign runtime. This commit adds support for
this to the two offloading plugins.
For nvptx, it supports cuda, cuda_driver and hip; the latter is AMD's
version of CUDA which for Nvidia devices boils down to normal CUDA.
Thus, at the end for this limited use, cuda/cuda_driver/hip are all
the same - and for plugin-nvptx.c, the they differ only in terms of
what gets fr_id, fr_name and get_interop_type_desc return.
For gcn, it supports hip and hsa.
Regarding get-mapped-ptr-1.c: That's actually a fix for the
GOMP_interop commit r15-8654-g99e2906ae255fc that added
GOMP_DEVICE_DEFAULT_OMP_61 alias omp_default_device, which is
a conforming device number. But that test used -5 as check for a
non-conforming device number.
libgomp/ChangeLog:
* plugin/plugin-gcn.c (_LIBGOMP_PLUGIN_INCLUDE): Define.
(struct hsa_runtime_fn_info): Add two queue functions.
(hipError_t, hipCtx_t, hipStream_s, hipStream_t): New types.
(struct hip_runtime_fn_info): New.
(hip_runtime_lib, hip_fns): New global vars.
(init_environment_variables): Handle hip_runtime_lib.
(init_hsa_runtime_functions): Load the two queue functions.
(init_hip_runtime_functions, GOMP_OFFLOAD_interop,
GOMP_OFFLOAD_get_interop_int, GOMP_OFFLOAD_get_interop_ptr,
GOMP_OFFLOAD_get_interop_str,
GOMP_OFFLOAD_get_interop_type_desc): New.
* plugin/plugin-nvptx.c (_LIBGOMP_PLUGIN_INCLUDE): Define.
(GOMP_OFFLOAD_interop, GOMP_OFFLOAD_get_interop_int,
GOMP_OFFLOAD_get_interop_ptr, GOMP_OFFLOAD_get_interop_str,
GOMP_OFFLOAD_get_interop_type_desc): New.
* testsuite/libgomp.c/interop-fr-1.c: New test.
* testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c: Use -6
not -5 as non-conforming device number.
|
|
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>
|
|
libgomp.exp added -fno-diagnostics-show-caret and -fdiagnostics-color=never
as 'additional_flags' for compilation. However, it turned out that this now
is insufficient as the [...] part of diagnostics have a hyperlink URL.
Solution: Use the -fdiagnostics-plain-output flag instead, added in commit
r11-2701-g129a1319c0ab73. This flag currently implies the following flags:
-fno-diagnostics-show-caret
-fno-diagnostics-show-line-numbers
-fdiagnostics-color=never
-fdiagnostics-urls=never
-fdiagnostics-path-format=separate-events
-fdiagnostics-text-art-charset=none
-fno-diagnostics-show-event-links
libgomp/ChangeLog:
* testsuite/lib/libgomp.exp (libgomp_init): Add
-fdiagnostics-plain-output to additional_flags; remove
-fno-diagnostics-show-caret and -fdiagnostics-color=never.
|
|
The test is a supposed to be a compile-only test but as dg-additional-sources
does not work with dg-compile (due to compiling with '-o'), dg-link had to
be used; as the code actually compiles (no diagnostic error), the linker
is actually invoked, which fails unless the system compiler by chance
provides the required files. Solution: move the test to libgomp.
PR fortran/115271
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/declare-variant-mod-1-use.f90: Move to
libgomp/testsuite/libgomp.fortran/.
* gfortran.dg/gomp/declare-variant-mod-1.f90: Likewise.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/declare-variant-mod-1-use.f90: Moved
from gcc/testsuite/gfortran.dg/gomp/.
* testsuite/libgomp.fortran/declare-variant-mod-1.f90: Likewise.
|
|
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.
|
|
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>
|
|
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>
|
|
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>
|
|
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.
|
|
..., 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.
|
|
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.
|
|
|
|
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.
|
|
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.
|
|
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.
|
|
__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.
|
|
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.
|
|
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;'.
|
|
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.
|
|
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.
|
|
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>
|
|
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.
|
|
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.
|
|
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.
|
|
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.
|
|
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.
|
|
'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'.
|
|
__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.
|
|
__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.
|
|
__builtin_is_initial_device
It turned out that 'if (omp_is_initial_device() .eqv. true)' gave an ICE
due to comparing 'int' with 'logical(4)'. When digging deeper, it also
turned out that when the procedure pointer is needed, the builtin cannot
be used, either. (Follow up to r15-2799-gf1bfba3a9b3f31 )
Extend the code to also use the builtin acc_on_device with OpenACC,
which was previously only used in C/C++. Additionally, fix folding
when offloading is not enabled.
Fixes additionally the BT_BOOL data type, which was 'char'/integer(1)
instead of bool, backing the booleaness; use bool_type_node as the rest
of GCC.
gcc/fortran/ChangeLog:
* gfortran.h (gfc_option_t): Add disable_acc_on_device.
* options.cc (gfc_handle_option): Handle -fno-builtin-acc_on_device.
* trans-decl.cc (gfc_get_extern_function_decl): Move
__builtin_omp_is_initial_device handling to ...
* trans-expr.cc (get_builtin_fn): ... this new function.
(conv_function_val): Call it.
(update_builtin_function): New.
(gfc_conv_procedure_call): Call it.
* types.def (BT_BOOL): Fix type by using bool_type_node.
gcc/ChangeLog:
* gimple-fold.cc (gimple_fold_builtin_acc_on_device): Also fold
when offloading is not configured.
libgomp/ChangeLog:
* libgomp.texi (TR13): Fix minor typos.
(omp_is_initial_device): Improve wording.
(acc_on_device): Note how to disable the builtin.
* testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Remove TODO.
* testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise.
Add -fno-builtin-acc_on_device.
* testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: Update
dg- as !offloading_enabled now compile-time expands acc_on_device.
* testsuite/libgomp.fortran/target-is-initial-device-3.f90: New test.
* testsuite/libgomp.oacc-fortran/acc_on_device-2.f90: New test.
|
|
The testcase was turned into a 'dg-do run' check to check for the alignment,
but this only works in testsuite/gfortran.dg, causing link errors for
out-of-tree testing. The test was added in r15-4104-ga8caeaacf499d5.
gcc/testsuite/:
* gfortran.dg/gomp/allocate-static.f90: Move to libgomp/testsuite/.
libgomp/:
* testsuite/libgomp.fortran/allocate-static.f90: Moved from
gcc/testsuite/ as it is a dg-do run test; use real omp_lib_kinds
instead of local definition
|
|
Loop [PR107637]
The following patch implements the C++23 P2718R0 paper
- Wording for P2644R1 Fix for Range-based for Loop.
The patch introduces a new option, -f{,no-}range-for-ext-temps so that
user can control the behavior even in older C++ versions.
The option is on by default in C++23 and later (-fno-range-for-ext-temps
is an error in that case) and in the -std=gnu++11 ... -std=gnu++20 modes
(one can use -fno-range-for-ext-temps to request previous behavior in that
case), and is not enabled by default in -std=c++11 ... -std=c++20 modes
but one can explicitly enable it with -frange-for-ext-temps.
As all the temporaries from __for_range initialization should have life
extended until the end of __for_range scope, this patch disables (for
-frange-for-ext-temps and if !processing_template_decl) CLEANUP_POINT_EXPR wrapping
of the __for_range declaration, also disables -Wdangling-reference warning
as well as the rest of extend_ref_init_temps (we know the __for_range temporary
is not TREE_STATIC and as all the temporaries from the initializer will be life
extended, we shouldn't try to handle temporaries referenced by references any
differently) and adds an extra push_stmt_list/pop_stmt_list before
cp_finish_decl of __for_range and after end of the for body and wraps all
that into CLEANUP_POINT_EXPR.
I had to repeat that also for OpenMP range loops because those are handled
differently.
2024-09-24 Jakub Jelinek <jakub@redhat.com>
PR c++/107637
gcc/
* omp-general.cc (find_combined_omp_for, find_nested_loop_xform):
Handle CLEANUP_POINT_EXPR like TRY_FINALLY_EXPR.
* doc/invoke.texi (frange-for-ext-temps): Document. Add
-fconcepts to the C++ option list.
gcc/c-family/
* c.opt (frange-for-ext-temps): New option.
* c-opts.cc (c_common_post_options): Set flag_range_for_ext_temps
for C++23 or later or for C++11 or later in !flag_iso mode if
the option wasn't set by user.
* c-cppbuiltin.cc (c_cpp_builtins): Change __cpp_range_based_for
value for flag_range_for_ext_temps from 201603L to 202212L in C++17
or later.
* c-omp.cc (c_find_nested_loop_xform_r): Handle CLEANUP_POINT_EXPR
like TRY_FINALLY_EXPR.
gcc/cp/
* cp-tree.h: Implement C++23 P2718R0 - Wording for P2644R1 Fix for
Range-based for Loop.
(cp_convert_omp_range_for): Add bool tmpl_p argument.
(find_range_for_decls): Declare.
* parser.cc (cp_convert_range_for): For flag_range_for_ext_temps call
push_stmt_list () before cp_finish_decl for range_temp and save it
temporarily to FOR_INIT_STMT.
(cp_convert_omp_range_for): Add tmpl_p argument. If set, remember
DECL_NAME of range_temp and for cp_finish_decl call restore it before
clearing it again, if unset, don't adjust DECL_NAME of range_temp at
all.
(cp_parser_omp_loop_nest): For flag_range_for_ext_temps range for add
CLEANUP_POINT_EXPR around sl. Call find_range_for_decls and adjust
DECL_NAMEs for range fors if not processing_template_decl. Adjust
cp_convert_omp_range_for caller. Remove superfluous backslash at the
end of line.
* decl.cc (initialize_local_var): For flag_range_for_ext_temps
temporarily clear stmts_are_full_exprs_p rather than set for
for_range__identifier decls.
* call.cc (extend_ref_init_temps): For flag_range_for_ext_temps return
init early for for_range__identifier decls.
* semantics.cc (find_range_for_decls): New function.
(finish_for_stmt): Use it. For flag_range_for_ext_temps if
cp_convert_range_for set FOR_INIT_STMT, pop_stmt_list it and wrap
into CLEANUP_POINT_EXPR.
* pt.cc (tsubst_omp_for_iterator): Adjust tsubst_omp_for_iterator
caller.
(tsubst_stmt) <case OMP_FOR>: For flag_range_for_ext_temps if there
are any range fors in the loop nest, add push_stmt_list starting
before the initializations, pop_stmt_list it after the body and wrap
into CLEANUP_POINT_EXPR. Change DECL_NAME of range for temps from
NULL to for_range_identifier.
gcc/testsuite/
* g++.dg/cpp23/range-for1.C: New test.
* g++.dg/cpp23/range-for2.C: New test.
* g++.dg/cpp23/range-for3.C: New test.
* g++.dg/cpp23/range-for4.C: New test.
* g++.dg/cpp23/range-for5.C: New test.
* g++.dg/cpp23/range-for6.C: New test.
* g++.dg/cpp23/range-for7.C: New test.
* g++.dg/cpp23/range-for8.C: New test.
* g++.dg/cpp23/feat-cxx2b.C (__cpp_range_based_for): Check for
202212L rather than 201603L.
* g++.dg/cpp26/feat-cxx26.C (__cpp_range_based_for): Likewise.
* g++.dg/warn/Wdangling-reference4.C: Don't expect warning for C++23
or newer. Use dg-additional-options rather than dg-options.
libgomp/
* testsuite/libgomp.c++/range-for-1.C: New test.
* testsuite/libgomp.c++/range-for-2.C: New test.
* testsuite/libgomp.c++/range-for-3.C: New test.
* testsuite/libgomp.c++/range-for-4.C: New test.
* testsuite/libgomp.c++/range-for-5.C: New test.
|
|
If requires unified_shared_memory or self_maps is set, make
'declare target link' variables to point initially to the host pointer.
libgomp/ChangeLog:
* target.c (gomp_load_image_to_device): For requires
unified_shared_memory, update 'link' vars to point to the host var.
* testsuite/libgomp.c-c++-common/target-link-3.c: New test.
* testsuite/libgomp.c-c++-common/target-link-4.c: New test.
|
|
'self_maps' implies 'unified_shared_memory', except that the latter
also permits that explicit maps copy data to device memory while
self_maps does not. In GCC, currently, both are handled identical.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_requires): Handle self_maps clause.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_requires): Handle self_maps clause.
gcc/fortran/ChangeLog:
* gfortran.h (enum gfc_omp_requires_kind): Add OMP_REQ_SELF_MAPS.
(gfc_namespace): Enlarge omp_requires bitfield.
* module.cc (enum ab_attribute, attr_bits): Add AB_OMP_REQ_SELF_MAPS.
(mio_symbol_attribute): Handle it.
* openmp.cc (gfc_check_omp_requires, gfc_match_omp_requires): Handle
self_maps clause.
* parse.cc (gfc_parse_file): Handle self_maps clause.
gcc/ChangeLog:
* lto-cgraph.cc (output_offload_tables, omp_requires_to_name): Handle
self_maps clause.
* omp-general.cc (struct omp_ts_info, omp_context_selector_matches):
Likewise for the associated trait.
* omp-general.h (enum omp_requires): Add OMP_REQUIRES_SELF_MAPS.
* omp-selectors.h (enum omp_ts_code): Add
OMP_TRAIT_IMPLEMENTATION_SELF_MAPS.
include/ChangeLog:
* gomp-constants.h (GOMP_REQUIRES_SELF_MAPS): #define.
libgomp/ChangeLog:
* plugin/plugin-gcn.c (GOMP_OFFLOAD_get_num_devices):
Accept self_maps clause.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_num_devices):
Likewise.
* libgomp.texi (TR13 Impl. Status): Set to 'Y'.
* target.c (gomp_requires_to_name, GOMP_offload_register_ver,
gomp_target_init): Handle self_maps clause.
* testsuite/libgomp.fortran/self_maps.f90: New test.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/declare-variant-1.c: Add self_maps test.
* c-c++-common/gomp/requires-4.c: Likewise.
* gfortran.dg/gomp/declare-variant-3.f90: Likewise.
* c-c++-common/gomp/requires-2.c: Update dg-error msg.
* gfortran.dg/gomp/requires-2.f90: Likewise.
* gfortran.dg/gomp/requires-self-maps-aux.f90: New.
* gfortran.dg/gomp/requires-self-maps.f90: New.
|