As done for 'GOMP_MAP_FROM', also for 'GOMP_MAP_FORCE_FROM' we should only
'gomp_copy_dev2host' if 'n->refcount == 0'.
This had gotten altered in commit 378da98fcc907d05002bcd3d6ff7951f0cf485e5
(r279621) "OpenACC reference count overhaul".
libgomp/
* oacc-mem.c (goacc_exit_data_internal): Revert always-copyfrom
behavior for 'GOMP_MAP_FORCE_FROM'.
* testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Adjust XFAIL.
These test cases use directives similar to:
/* { dg-additional-options "-save-temps" } */
/* { dg-final { scan-assembler-times "bar.sync" 2 } } */
This expects to scan the PTX offloading compilation assembler code (not host
code!), expecting that nvptx offloading code assembly is produced after the
host code, and thus overwrites the latter file. (Yes, that's certainly
ugly/fragile...)
..., and this broke with recent commit 1dedc12d186a110854537e1279b4e6c29f2df35a
"revamp dump and aux output names" plus fix-up commit commit
efc16503ca10bc0e934e0bace5777500e4dc757a "handle dumpbase in offloading, adjust
testsuite" (short summary: file names changed), so let's finally make that
robust.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Replace fragile
'scan-assembler' with 'scan-offload-rtl'.
* testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381.c: Likewise.
This implements the fallback mentioned in
https://gcc.gnu.org/pipermail/gcc/2020-June/232874.html
Special cases for triangular loops etc. to follow later, also composite
constructs not supported yet (need to check the passing of temporaries around)
and lastprivate might not give the same answers as serial loop if the last
innermost body iteration isn't the last one for some of the outer loops
(that will need to be solved separately together with rectangular loops that have no
innermost body iterations, but some of the outer loops actually iterate).
Also, simd needs work.
2020-06-27 Jakub Jelinek <jakub@redhat.com>
* omp-general.h (struct omp_for_data_loop): Add non_rect_referenced
member, move outer member.
(struct omp_for_data): Add first_nonrect and last_nonrect members.
* omp-general.c (omp_extract_for_data): Initialize first_nonrect,
last_nonrect and non_rect_referenced members.
* omp-expand.c (expand_omp_for_init_counts): Handle non-rectangular
loops.
(expand_omp_for_init_vars): Add nonrect_bounds parameter. Handle
non-rectangular loops.
(extract_omp_for_update_vars): Likewise.
(expand_omp_for_generic, expand_omp_for_static_nochunk,
expand_omp_for_static_chunk, expand_omp_simd,
expand_omp_taskloop_for_outer, expand_omp_taskloop_for_inner): Adjust
expand_omp_for_init_vars and extract_omp_for_update_vars callers.
(expand_omp_for): Don't sorry on non-composite worksharing-loop or
distribute.
* testsuite/libgomp.c/loop-17.c: New test.
* testsuite/libgomp.c/loop-18.c: New test.
Since GCC 9, C++17 support is no longer experimental. It was too late
to change the default C++ dialect to C++17 in GCC 10, but I think now
it's time to pull the trigger (C++14 was made the default in GCC 6.1).
We're still missing two C++17 library features, but that shouldn't stop
us. See
<https://gcc.gnu.org/onlinedocs/libstdc++/manual/status.html#status.iso.2017>
and
<https://gcc.gnu.org/projects/cxx-status.html#cxx17>
for the C++17 status.
I won't list all C++17 features here, but just a few heads-up:
- trigraphs were removed (hardly anyone cares, unless your keyboard is
missing the # key),
- operator++(bool) was removed (so some tests now run in C++14 and down
only),
- the keyword register was removed (some legacy code might trip on
this),
- noexcept specification is now part of the type system and C++17 does
not allow dynamic exception specifications anymore (the empty throw
specification is still available, but it is deprecated),
- the evaluation order rules are different in C++17,
- static constexpr data members are now implicitly inline (which makes
them definitions),
- C++17 requires guaranteed copy elision, meaning that a copy/move
constructor call might be elided completely. That means that if
something relied on a constructor being instantiated via e.g. copying
a function parameter, it might now fail.
I'll post an update for cxx-status.html and add a new caveat to changes.html
once this is in.
gcc/ChangeLog:
* doc/invoke.texi (C Dialect Options): Adjust -std default for C++.
* doc/standards.texi (C Language): Correct the default dialect.
(C++ Language): Update the default for C++ to gnu++17.
gcc/c-family/ChangeLog:
* c-opts.c (c_common_init_options): Default to gnu++17.
gcc/testsuite/ChangeLog:
* c-c++-common/torture/vector-subscript-3.c: In C++17, define away
the keyword register.
* g++.dg/cpp1z/attributes-enum-1a.C: Only run pre-C++17.
* g++.dg/cpp1z/fold7a.C: Likewise.
* g++.dg/cpp1z/nontype3a.C: Likewise.
* g++.dg/cpp1z/utf8-2a.C: Likewise.
* g++.dg/parse/error11.C: Update expected diagnostics for C++17.
* g++.dg/torture/pr34850.C: Add -Wno-attribute-warning.
* g++.dg/torture/pr49394.C: In C++17, use noexcept(false).
* g++.dg/torture/pr82154.C: Use -std=c++14.
* lib/target-supports.exp: Set to C++17.
* obj-c++.dg/try-catch-9.mm: Use -Wno-register.
libgomp/ChangeLog:
* testsuite/libgomp.c++/atomic-3.C: Use -std=gnu++14.
Pass dumpbase on to mkoffloads and their offload-target compiler runs,
using different suffixes for different offloading targets.
Obey -save-temps in naming temporary files while at that.
Adjust the testsuite offload dump scanning machinery to look for dump
files named under the new conventions, iterating internally over all
configured offload targets, or recognizing libgomp's testsuite's own
iteration.
for gcc/ChangeLog
* collect-utils.h (dumppfx): New.
* collect-utils.c (dumppfx): Likewise.
* lto-wrapper.c (run_gcc): Set global dumppfx.
(compile_offload_image): Pass a -dumpbase on to mkoffload.
* config/nvptx/mkoffload.c (ptx_dumpbase): New.
(main): Handle incoming -dumpbase. Set ptx_dumpbase. Obey
save_temps.
(compile_native): Pass -dumpbase et al to compiler.
* config/gcn/mkoffload.c (gcn_dumpbase): New.
(main): Handle incoming -dumpbase. Set gcn_dumpbase. Obey
save_temps. Pass -dumpbase et al to offload target compiler.
(compile_native): Pass -dumpbase et al to compiler.
for gcc/testsuite/ChangeLog
* lib/scanoffload.exp: New.
* lib/scanoffloadrtl.exp: Load it. Replace ".o" with ""
globally, and use scanoffload's scoff wrapper to fill it in.
* lib/scanoffloadtree.exp: Likewise.
for libgomp/ChangeLog
* testsuite/lib/libgomp.exp: Load gcc lib scanoffload.exp.
* testsuite/lib/libgomp-dg.exp: Drop now-obsolete -save-temps.
Fix-up for r279858/commit f760c0c77fe350616da9dbeaea16442b0acfb09c "Fortran]
OpenMP/OpenACC – fix more issues with OPTIONAL".
With offloading enabled, we then saw:
PASS: libgomp.fortran/use_device_ptr-optional-3.f90 -O0 (test for excess errors)
PASS: libgomp.fortran/use_device_ptr-optional-3.f90 -O0 execution test
PASS: libgomp.fortran/use_device_ptr-optional-3.f90 -O1 (test for excess errors)
PASS: libgomp.fortran/use_device_ptr-optional-3.f90 -O1 execution test
FAIL: libgomp.fortran/use_device_ptr-optional-3.f90 -O2 (test for excess errors)
UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90 -O2 compilation failed to produce executable
FAIL: libgomp.fortran/use_device_ptr-optional-3.f90 -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions (test for excess errors)
UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90 -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions compilation failed to produce executable
FAIL: libgomp.fortran/use_device_ptr-optional-3.f90 -O3 -g (test for excess errors)
UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90 -O3 -g compilation failed to produce executable
FAIL: libgomp.fortran/use_device_ptr-optional-3.f90 -Os (test for excess errors)
UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90 -Os compilation failed to produce executable
... due to:
/tmp/cciVc43I.o:(.gnu.offload_vars+0x10): undefined reference to `A.12.4064'
[...]
..., but after the recent PR94848, PR95551 changes, that problem is now gone.
libgomp/
PR lto/94848
* testsuite/libgomp.fortran/use_device_ptr-optional-3.f90: Add
'dg-do run'.
gcc/fortran/ChangeLog
* parse.c (decode_oacc_directive): Permit 'acc routine' also
inside pure procedures.
* openmp.c (gfc_match_oacc_routine): Inside pure procedures
do not permit gang, worker or vector clauses.
libgomp/ChangeLog:
* testsuite/libgomp.oacc-fortran/routine-10.f90: New test.
gcc/testsuite/ChangeLog:
* gfortran.dg/goacc/pure-elemental-procedures-2.f90: New test.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
libgomp/
* testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
Currently, we don't at all evaluate 'copyfrom' for 'GOMP_MAP_STRUCT' entries.
Fix this by copying/adapting the corresponding non-'GOMP_MAP_STRUCT' code.
libgomp/
* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>:
Evaluate 'copyfrom' individually for each entry.
* testsuite/libgomp.oacc-c-c++-common/struct-1.c: Update.
Currently, we don't at all evaluate 'finalize' for 'GOMP_MAP_STRUCT' entries.
Fix this by copying/adapting the corresponding non-'GOMP_MAP_STRUCT' code.
libgomp/
* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>:
Evaluate 'finalize' individually for each entry.
* testsuite/libgomp.oacc-c-c++-common/struct-1.c: New file.
* testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c: Remove
file.
If, for example, GCC is configured such that 'libgomp-plugin-nvptx.so.1'
dynamically links against 'libcuda.so.1', but testing is run on a system where
there is no 'libcuda.so.1', this produces output such as:
PASS: libgomp.oacc-fortran/error_stop-1.f -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable -O0 (test for excess errors)
PASS: libgomp.oacc-fortran/error_stop-1.f -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable -O0 execution test
FAIL: libgomp.oacc-fortran/error_stop-1.f -DACC_DEVICE_TYPE_host=1 -DACC_MEM_SHARED=1 -foffload=disable -O0 output pattern test, is CheCKpOInT
libgomp: while loading libgomp-plugin-nvptx.so.1: libcuda.so.1: cannot open shared object file: No such file or directory
ERROR STOP
Error termination. Backtrace: [...]
, should match CheCKpOInT(
|
|^M)+ERROR STOP (
|
|^M)+Error termination.*
..., where after 'CheCKpOInT' we got 'libgomp: while loading [...]' injected
before the expected 'ERROR STOP'.
libgomp/
* testsuite/libgomp.oacc-fortran/error_stop-1.f: Initialize before
the checkpoint.
* testsuite/libgomp.oacc-fortran/error_stop-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/error_stop-3.f: Likewise.
* testsuite/libgomp.oacc-fortran/stop-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/stop-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/stop-3.f: Likewise.
gcc/fortran/ChangeLog:
2020-05-23 Thomas Koenig <tkoenig@gcc.gnu.org>
PR libfortran/95191
* libgfortran.h (libgfortran_error_codes): Add
LIBERROR_BAD_WAIT_ID.
libgfortran/ChangeLog:
2020-05-23 Thomas Koenig <tkoenig@gcc.gnu.org>
PR libfortran/95191
* io/async.c (async_wait_id): Generate error if ID is higher
than the highest current ID.
* runtime/error.c (translate_error): Handle LIBERROR_BAD_WAIT_ID.
libgomp/ChangeLog:
2020-05-23 Thomas Koenig <tkoenig@gcc.gnu.org>
PR libfortran/95191
* testsuite/libgomp.fortran/async_io_9.f90: New test.
This patch adds very basic allocator support (omp_{init,destroy}_allocator,
omp_{alloc,free}, omp_[sg]et_default_allocator).
The plan is to use memkind (likely dlopened) for high bandwidth memory, but
that part isn't implemented yet, probably mlock for pinned memory and see
what other options there are for other kinds of memory.
For offloading targets, we need to decide if we want to support the
dynamic allocators (and on which targets), or if e.g. all we do is at compile
time replace omp_alloc/omp_free calls with constexpr predefined allocators
with something special.
And allocate directive and allocator/uses_allocators clauses are future work
too.
2020-05-19 Jakub Jelinek <jakub@redhat.com>
* omp.h.in (omp_uintptr_t): New typedef.
(__GOMP_UINTPTR_T_ENUM): Define.
(omp_memspace_handle_t, omp_allocator_handle_t, omp_alloctrait_key_t,
omp_alloctrait_value_t, omp_alloctrait_t): New typedefs.
(__GOMP_DEFAULT_NULL_ALLOCATOR): Define.
(omp_init_allocator, omp_destroy_allocator, omp_set_default_allocator,
omp_get_default_allocator, omp_alloc, omp_free): Declare.
* libgomp.h (struct gomp_team_state): Add def_allocator field.
(gomp_def_allocator): Declare.
* libgomp.map (OMP_5.0.1): Export omp_set_default_allocator,
omp_get_default_allocator, omp_init_allocator, omp_destroy_allocator,
omp_alloc and omp_free.
* team.c (gomp_team_start): Copy over ts.def_allocator.
* env.c (gomp_def_allocator): New variable.
(parse_wait_policy): Adjust function comment.
(parse_allocator): New function.
(handle_omp_display_env): Print OMP_ALLOCATOR.
(initialize_env): Call parse_allocator.
* Makefile.am (libgomp_la_SOURCES): Add allocator.c.
* allocator.c: New file.
* icv.c (omp_set_default_allocator, omp_get_default_allocator): New
functions.
* testsuite/libgomp.c-c++-common/alloc-1.c: New test.
* testsuite/libgomp.c-c++-common/alloc-2.c: New test.
* testsuite/libgomp.c-c++-common/alloc-3.c: New test.
* Makefile.in: Regenerated.
2020-05-14 Thomas Koenig <tkoenig@gcc.gnu.org>
PR libfortran/95119
* io/close.c (close_status): Add CLOSE_INVALID.
(st_close): Return early on invalid STATUS parameter.
2020-05-14 Thomas Koenig <tkoenig@gcc.gnu.org>
PR libfortran/95119
* testsuite/libgomp.fortran/close_errors_1.f90: New test.
OpenMP 5.0 also specifies that functions referenced from target regions
(except for target regions with device(ancestor:)) are also implicitly declare target to.
This patch implements that.
2020-05-14 Jakub Jelinek <jakub@redhat.com>
* function.h (struct function): Add has_omp_target bit.
* omp-offload.c (omp_discover_declare_target_fn_r): New function,
old renamed to ...
(omp_discover_declare_target_tgt_fn_r): ... this.
(omp_discover_declare_target_var_r): Call
omp_discover_declare_target_tgt_fn_r instead of
omp_discover_declare_target_fn_r.
(omp_discover_implicit_declare_target): Also queue functions with
has_omp_target bit set, for those walk with
omp_discover_declare_target_fn_r, for declare target to functions
walk with omp_discover_declare_target_tgt_fn_r.
gcc/c/
* c-parser.c (c_parser_omp_target): Set cfun->has_omp_target.
gcc/cp/
* cp-gimplify.c (cp_genericize_r): Set cfun->has_omp_target.
gcc/fortran/
* trans-openmp.c: Include function.h.
(gfc_trans_omp_target): Set cfun->has_omp_target.
libgomp/
* testsuite/libgomp.c-c++-common/target-40.c: New test.
This attempts to implement what the OpenMP 5.0 spec in declare target section
says as ammended by the 5.1 changes so far (related to device_type(host)), except
that it doesn't have the device(ancestor: ...) handling yet because we do not
support it yet, and I've left so far out the except lambda note, because I need
that clarified.
2020-05-12 Jakub Jelinek <jakub@redhat.com>
* omp-offload.h (omp_discover_implicit_declare_target): Declare.
* omp-offload.c: Include context.h.
(omp_declare_target_fn_p, omp_declare_target_var_p,
omp_discover_declare_target_fn_r, omp_discover_declare_target_var_r,
omp_discover_implicit_declare_target): New functions.
* cgraphunit.c (analyze_functions): Call
omp_discover_implicit_declare_target.
* testsuite/libgomp.c/target-39.c: New test.
Testing on the host does not make sense for 'declare copyout' for
a same-scope stack-allocated variable. Once the copyout is done,
the variable is gone. Hence, test the variable on the device. This
can be revisit after the OpenACC semantic has been fixed; but with
that fix, the test PASSes again with devices.
PR middle-end/94120
* testsuite/libgomp.oacc-c++/declare-pr94120.C: Fix 'declare copy(out)'
test case.
gcc/c/
PR middle-end/94120
* c-decl.c (c_check_in_current_scope): New function.
* c-tree.h (c_check_in_current_scope): Declare it.
* c-parser.c (c_parser_oacc_declare): Add check that variables
are declared in the same scope as the directive. Fix handling
of namespace vars.
gcc/cp/
PR middle-end/94120
* paser.c (cp_parser_oacc_declare): Add check that variables
are declared in the same scope as the directive.
gcc/testsuite/
PR middle-end/94120
* c-c++-common/goacc/declare-pr94120.c: New.
* g++.dg/declare-pr94120.C: New.
libgomp/testsuite/
PR middle-end/94120
* libgomp.oacc-c++/declare-pr94120.C: New.
Fix a problem with commit c8e759b4215b ("libgomp/test: Fix compilation
for build sysroot") that caused a regression in some standalone test
environments where testsuite/libgomp-test-support.exp is used, but the
compiler is expected to be determined by `[find_gcc]', and set the
GCC_UNDER_TEST TCL variable in testsuite/libgomp-site-extra.exp instead.
libgomp/
* configure.ac: Add testsuite/libgomp-site-extra.exp to output
files.
* configure: Regenerate.
* testsuite/libgomp-site-extra.exp.in: New file.
* testsuite/libgomp-test-support.exp.in (GCC_UNDER_TEST): Remove
variable.
* testsuite/Makefile.am (EXTRA_DEJAGNU_SITE_CONFIG): New
variable.
* testsuite/Makefile.in: Regenerate.
In response to PR94392 commit 75efe9cb1f8938a713ce540dc3b27bc2afcd3fae
"c/94392 - only enable -ffinite-loops for C++", this reverts PR89713
commit 00908992f2a78f213d227aea8dbab014a1361df0, as apparently now again
"empty oacc loops are" no longer "removed before expand".
libgomp/
PR tree-optimization/89713
PR c/94392
* testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Again expect
'bar.sync'.
* testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise.
Without the parser.c change we were ICEing on the testcase, because while the
uses of the captured vars inside of the constructs were replaced with capture
proxy decls, we didn't do that for decls in OpenMP clauses.
With that fixed, we don't ICE anymore, but the testcase is miscompiled and FAILs
at runtime. This is because the capture proxy decls have DECL_VALUE_EXPR and
during gimplification we were gimplifying those to their DECL_VALUE_EXPRs.
That is fine for shared vars, but for privatized ones we must not do that.
So that is what the cp-gimplify.c changes do. Had to add a DECL_CONTEXT check
before calling is_capture_proxy because some VAR_DECLs don't have DECL_CONTEXT
set (yet) and is_capture_proxy relies on that being non-NULL always.
2020-03-19 Jakub Jelinek <jakub@redhat.com>
PR c++/93931
* parser.c (cp_parser_omp_var_list_no_open): Call process_outer_var_ref
on outer_automatic_var_p decls.
* cp-gimplify.c (cxx_omp_disregard_value_expr): Return true also for
capture proxy decls.
* testsuite/libgomp.c++/pr93931.C: New test.
2020-03-18 Julian Brown <julian@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
* testsuite/libgomp.oacc-fortran/atomic_capture-1.f90: Really make
it work concurrently.
* testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C: Add
#define DO_LONG_DOUBLE; set to 1, except for nvidia + gcn.
* libgomp.oacc-c-c++-common/firstprivate-mappings-1.c: Likewise.
* g++.dg/goacc/firstprivate-mappings-1.C: Only set DO_LONG_DOUBLE if
not defined; update comments.
* c-c++-common/goacc/firstprivate-mappings-1.c: Likewise.
The commit r10-6721-g8d1a1cb1b816381bf60cb1211c93b8eba1fe1472 has changed
the name of the type that is used for the return value of the Fortran
acc_get_property function without adapting the test acc_get_property.f90.
2020-02-21 Frederik Harwath <frederik@codesourcery.com>
* testsuite/libgomp.oacc-fortran/acc_get_property.f90: Adapt to
changes from 2020-02-19, i.e. use integer(c_size_t) instead of
integer(acc_device_property) for the type of the return value of
acc_get_property.
An OpenMP "nowait" clause on a target construct currently leads to
a call to GOMP_OFFLOAD_async_run in the plugin that is used for
offloading at execution time. The nvptx plugin contains only a stub
of this function that always produces a fatal error if called.
This commit changes the "nowait" implementation to ignore the clause
if the executing device's plugin does not implement GOMP_OFFLOAD_async_run.
The stub in the nvptx plugin is removed which effectively means that
programs containing "nowait" can now be executed with nvptx offloading
as if the clause had not been used.
This behavior is consistent with the OpenMP specification which says that
"[...] execution of the target task *may* be deferred" (emphasis added),
cf. OpenMP 5.0, page 172.
libgomp/
* plugin/plugin-nvptx.c: Remove GOMP_OFFLOAD_async_run stub.
* target.c (gomp_load_plugin_for_device): Make "async_run" loading
optional.
(gomp_target_task_fn): Assert "devicep->async_run_func".
(clear_unsupported_flags): New function to remove unsupported flags
(right now only GOMP_TARGET_FLAG_NOWAIT) that can be be ignored.
(GOMP_target_ext): Apply clear_unsupported_flags to flags.
* testsuite/libgomp.c/target-33.c:
Remove xfail for offload_target_nvptx.
* testsuite/libgomp.c/target-34.c: Likewise.
Add xfails for nvptx offloading because
"no GOMP_OFFLOAD_async_run implemented in plugin-nvptx.c"
(https://gcc.gnu.org/PR81688) and because
"omp target link not implemented for nvptx"
(https://gcc.gnu.org/PR81689).
libgomp/
* testsuite/libgomp.c/target-33.c: Add xfail for execution on
offload_target_nvptx, cf. https://gcc.gnu.org/PR81688.
* testsuite/libgomp.c/target-34.c: Likewise.
* testsuite/libgomp.c/target-link-1.c: Add xfail for
offload_target_nvptx, cf. https://gcc.gnu.org/PR81689.
DECL_IN_CONSTANT_POOL are shared and thus don't really get emitted in the
BLOCK where they are used, so for OpenMP target regions that have initializers
gimplified into copying from them we actually map them at runtime from host to
offload devices. This patch instead marks them as "omp declare target", so
that they are on the target device from the beginning and don't need to be
copied there.
2020-02-09 Jakub Jelinek <jakub@redhat.com>
* gimplify.c (gimplify_adjust_omp_clauses_1): Promote
DECL_IN_CONSTANT_POOL variables into "omp declare target" to avoid
copying them around between host and target.
* testsuite/libgomp.c/target-38.c: New test.
As the following testcase shows, we need to consider even target to be a construct
that forces not to use copy in/out for shared on parallel inside of the target.
E.g. for parallel nested inside another parallel or host teams, we already avoid
copy in/out and we need to treat target the same.
2020-02-06 Jakub Jelinek <jakub@redhat.com>
PR libgomp/93515
* omp-low.c (use_pointer_for_field): For nested constructs, also
look for map clauses on target construct.
(scan_omp_1_stmt) <case GIMPLE_OMP_TARGET>: Bump temporarily
taskreg_nesting_level.
* testsuite/libgomp.c-c++-common/pr93515.c: New test.
* testsuite/lib/libgomp.exp
(check_effective_target_offload_target_nvptx): Pass flags as 'options'
and not as 'source' argument to libgomp_target_compile.
