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.
..., so that users don't manually need to specify '-foffload-options=-lstdc++'
in addition to '-lstdc++' (specified manually, or implicitly by the driver).
Do like commit 4bcb46b3ad
"driver: Forward '-lgfortran', '-lm' to offloading compilation".
PR driver/101544
gcc/
* gcc.cc (driver_handle_option): Forward host '-lstdc++' to
offloading compilation.
* config/gcn/mkoffload.cc (main): Adjust.
* config/nvptx/mkoffload.cc (main): Likewise.
libgomp/
* testsuite/libgomp.c++/pr101544-1-O0.C: Remove
'-foffload-options=-lstdc++'.
* testsuite/libgomp.c++/pr101544-1.C: Likewise.
* testsuite/libgomp.oacc-c++/pr101544-1.C: Likewise.
This is a fix for the GOMP_interop commit r15-8654-g99e2906ae255fc that
added GOMP_DEVICE_DEFAULT_OMP_61 alias omp_default_device, which is a
conforming device number. But that test used -5 as check for a
non-conforming device number.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/get-mapped-ptr-1.f90: Use -6
not -5 as non-conforming device number.
The interop directive operates on an opaque object that represents a
foreign runtime. This commit adds support for
this to the two offloading plugins.
For nvptx, it supports cuda, cuda_driver and hip; the latter is AMD's
version of CUDA which for Nvidia devices boils down to normal CUDA.
Thus, at the end for this limited use, cuda/cuda_driver/hip are all
the same - and for plugin-nvptx.c, the they differ only in terms of
what gets fr_id, fr_name and get_interop_type_desc return.
For gcn, it supports hip and hsa.
Regarding get-mapped-ptr-1.c: That's actually a fix for the
GOMP_interop commit r15-8654-g99e2906ae255fc that added
GOMP_DEVICE_DEFAULT_OMP_61 alias omp_default_device, which is
a conforming device number. But that test used -5 as check for a
non-conforming device number.
libgomp/ChangeLog:
* plugin/plugin-gcn.c (_LIBGOMP_PLUGIN_INCLUDE): Define.
(struct hsa_runtime_fn_info): Add two queue functions.
(hipError_t, hipCtx_t, hipStream_s, hipStream_t): New types.
(struct hip_runtime_fn_info): New.
(hip_runtime_lib, hip_fns): New global vars.
(init_environment_variables): Handle hip_runtime_lib.
(init_hsa_runtime_functions): Load the two queue functions.
(init_hip_runtime_functions, GOMP_OFFLOAD_interop,
GOMP_OFFLOAD_get_interop_int, GOMP_OFFLOAD_get_interop_ptr,
GOMP_OFFLOAD_get_interop_str,
GOMP_OFFLOAD_get_interop_type_desc): New.
* plugin/plugin-nvptx.c (_LIBGOMP_PLUGIN_INCLUDE): Define.
(GOMP_OFFLOAD_interop, GOMP_OFFLOAD_get_interop_int,
GOMP_OFFLOAD_get_interop_ptr, GOMP_OFFLOAD_get_interop_str,
GOMP_OFFLOAD_get_interop_type_desc): New.
* testsuite/libgomp.c/interop-fr-1.c: New test.
* testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c: Use -6
not -5 as non-conforming device number.
libgomp.exp added -fno-diagnostics-show-caret and -fdiagnostics-color=never
as 'additional_flags' for compilation. However, it turned out that this now
is insufficient as the [...] part of diagnostics have a hyperlink URL.
Solution: Use the -fdiagnostics-plain-output flag instead, added in commit
r11-2701-g129a1319c0ab73. This flag currently implies the following flags:
-fno-diagnostics-show-caret
-fno-diagnostics-show-line-numbers
-fdiagnostics-color=never
-fdiagnostics-urls=never
-fdiagnostics-path-format=separate-events
-fdiagnostics-text-art-charset=none
-fno-diagnostics-show-event-links
libgomp/ChangeLog:
* testsuite/lib/libgomp.exp (libgomp_init): Add
-fdiagnostics-plain-output to additional_flags; remove
-fno-diagnostics-show-caret and -fdiagnostics-color=never.
The test is a supposed to be a compile-only test but as dg-additional-sources
does not work with dg-compile (due to compiling with '-o'), dg-link had to
be used; as the code actually compiles (no diagnostic error), the linker
is actually invoked, which fails unless the system compiler by chance
provides the required files. Solution: move the test to libgomp.
PR fortran/115271
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/declare-variant-mod-1-use.f90: Move to
libgomp/testsuite/libgomp.fortran/.
* gfortran.dg/gomp/declare-variant-mod-1.f90: Likewise.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/declare-variant-mod-1-use.f90: Moved
from gcc/testsuite/gfortran.dg/gomp/.
* testsuite/libgomp.fortran/declare-variant-mod-1.f90: Likewise.
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.
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.
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.
The motivation of the 'gimple_fold_builtin_acc_on_device' change in
commit 3269a722b7
"Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' __builtin_is_initial_device"
is unclear, and it unnecessarily diverges GCC's (default)
'--disable-offload-targets' vs. '--enable-offload-targets=[...]'
configurations.
PR testsuite/82250
gcc/
* gimple-fold.cc (gimple_fold_builtin_acc_on_device): Revert last
change.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: Revert
last change.
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.
With commit 55e3bd376b "c: Default to -std=gnu23"
we've got:
[-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/acc_get_property-gcn.c -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O0 (test for excess errors)
[-PASS:-]{+UNRESOLVED:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/acc_get_property-gcn.c -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O0 [-execution test-]{+compilation failed to produce executable+}
[Etc.]
..., due to:
[...]/libgomp.oacc-c/../libgomp.oacc-c-c++-common/acc_get_property-gcn.c:16:13: error: two or more data types in declaration specifiers
[...]/libgomp.oacc-c/../libgomp.oacc-c-c++-common/acc_get_property-gcn.c:16:1: warning: useless type name in empty declaration
libgomp/
* testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c
[!__cplusplus]: Don't 'typedef int bool;'.
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.
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.
The test case 'libgomp.oacc-fortran/routine-nohost-1.f90' added in 2021
commit a61f6afbee "OpenACC 'nohost' clause" was
dependend on inlining being enabled, and otherwise ('-fno-inline') failed to
optimize/link:
/tmp/ccb2hsPd.o: In function `MAIN__._omp_fn.0':
routine-nohost-1.f90:(.text+0xf4): undefined reference to `fact_nohost_'
However, as of recent commit 3269a722b7
"Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' __builtin_is_initial_device",
we're now properly handling OpenACC/Fortran 'acc_on_device', and may specify
'-fno-inline', like done in 'libgomp.oacc-c-c++-common/routine-nohost-1.c'.
libgomp/
* testsuite/libgomp.oacc-fortran/routine-nohost-1.f90: Add
'-fno-inline'.
The test case 'libgomp.oacc-fortran/acc_on_device-1-1.f90' added in
commit 3269a722b7
"Fortran: Use OpenACC's acc_on_device builtin, fix OpenMP' __builtin_is_initial_device"
was missing '-fno-builtin-acc_on_device', and all
'libgomp.oacc-fortran/acc_on_device-1-*' need comments, why that option is
specified.
PR testsuite/82250
libgomp/
* testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Add
'-fno-builtin-acc_on_device'.
* testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Comment.
* testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Comment.
The test case 'libgomp.oacc-fortran/acc_on_device-2.f90' added in
commit 3269a722b7
"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.
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
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.
In Fortran, omp_get_device_from_uid can also accept substrings, which are
then not NUL terminated. Fixed by introducing a fortran.c wrapper function.
Additionally, in case of a fail the plugin functions now return NULL instead
of failing fatally such that a fall-back UID is generated.
gcc/ChangeLog:
* omp-general.cc (omp_runtime_api_procname): Strip "omp_" from
string; move get_device_from_uid as now a '_' suffix exists.
libgomp/ChangeLog:
* fortran.c (omp_get_device_from_uid_): New function.
* libgomp.map (GOMP_6.0): Add it.
* oacc-host.c (host_dispatch): Init '.uid' and '.get_uid_func'.
* omp_lib.f90.in: Make it used by removing bind(C).
* omp_lib.h.in: Likewise.
* target.c (omp_get_device_from_uid): Ensure the device is initialized.
* plugin/plugin-gcn.c (GOMP_OFFLOAD_get_uid): Add function comment;
return NULL in case of an error.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_uid): Likewise.
* testsuite/libgomp.fortran/device_uid.f90: Update to test substrings.
Those TR13/OpenMP 6.0 routines permit a reproducible offloading to
a specific device by mapping an OpenMP device number to a
unique ID (UID). The GPU device UIDs should be universally unique,
the one for the host is not.
gcc/ChangeLog:
* omp-general.cc (omp_runtime_api_procname): Add
get_device_from_uid and omp_get_uid_from_device routines.
include/ChangeLog:
* cuda/cuda.h (cuDeviceGetUuid): Declare.
(cuDeviceGetUuid_v2): Add prototype.
libgomp/ChangeLog:
* config/gcn/target.c (omp_get_uid_from_device,
omp_get_device_from_uid): Add stub implementation.
* config/nvptx/target.c (omp_get_uid_from_device,
omp_get_device_from_uid): Likewise.
* fortran.c (omp_get_uid_from_device_,
omp_get_uid_from_device_8_): New functions.
* libgomp-plugin.h (GOMP_OFFLOAD_get_uid): Add prototype.
* libgomp.h (struct gomp_device_descr): Add 'uid' and 'get_uid_func'.
* libgomp.map (GOMP_6.0): New, includind the new UID routines.
* libgomp.texi (OpenMP Technical Report 13): Mark UID routines as 'Y'.
(Device Information Routines): Document new UID routines.
(Offload-Target Specifics): Document UID format.
* omp.h.in (omp_get_device_from_uid, omp_get_uid_from_device):
New prototype.
* omp_lib.f90.in (omp_get_device_from_uid, omp_get_uid_from_device):
New interface.
* omp_lib.h.in: Likewise.
* plugin/cuda-lib.def: Add cuDeviceGetUuid and cuDeviceGetUuid_v2 via
CUDA_ONE_CALL_MAYBE_NULL.
* plugin/plugin-gcn.c (GOMP_OFFLOAD_get_uid): New.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_uid): New.
* target.c (str_omp_initial_device): New static var.
(STR_OMP_DEV_PREFIX): Define.
(gomp_get_uid_for_device, omp_get_uid_from_device,
omp_get_device_from_uid): New.
(gomp_load_plugin_for_device): DLSYM_OPT the function 'get_uid'.
(gomp_target_init): Set the device's 'uid' field to NULL.
* testsuite/libgomp.c/device_uid.c: New test.
* testsuite/libgomp.fortran/device_uid.f90: New test.
