Commit Graph

1280 Commits

Author SHA1 Message Date
Tobias Burnus
1c0fdaf79e openmp: ensure variables in offload table are streamed out (PRs 94848 + 95551)
gcc/ChangeLog:

	PR lto/94848
	PR middle-end/95551
	* omp-offload.c (add_decls_addresses_to_decl_constructor,
	omp_finish_file): Skip removed items.
	* lto-cgraph.c (output_offload_tables): Likewise; set force_output
	to this node for variables and functions.

libgomp/ChangeLog:

	PR lto/94848
	PR middle-end/95551
	* testsuite/libgomp.fortran/target-var.f90: New test.
2020-06-08 23:24:57 +02:00
GCC Administrator
1a59f3db4a Daily bump. 2020-06-06 00:16:29 +00:00
Thomas Schwinge
1afc467256 [OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings
These are not itself necessary for OpenACC 'exit data' directives, and are
skipped over (now) in libgomp.  We might as well not emit them to start with,
in line with the equivalent OpenMP directive.  We keep the no-op handling in
libgomp for the reason of backward compatibility.

	gcc/
	* gimplify.c (gimplify_adjust_omp_clauses): Remove
	'GOMP_MAP_STRUCT' mapping from OpenACC 'exit data' directives.
	gcc/testsuite/
	* c-c++-common/goacc/struct-enter-exit-data-1.c: New file.
	libgomp/
	* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>: Explain
	special handling.

Co-Authored-By: Julian Brown <julian@codesourcery.com>
2020-06-05 18:04:13 +02:00
Thomas Schwinge
1809628fcf [OpenACC 'exit data'] Simplify 'GOMP_MAP_STRUCT' handling
libgomp/
	* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>:
	Simplify.

Co-Authored-By: Julian Brown <julian@codesourcery.com>
2020-06-05 18:04:12 +02:00
Julian Brown
9643f5bbe2 Add 'libgomp.oacc-c-c++-common/struct-copyout-{1,2}.c'
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>
2020-06-05 18:04:12 +02:00
GCC Administrator
3add342502 Daily bump. 2020-06-05 00:16:30 +00:00
Thomas Schwinge
2c838a3e4e [OpenACC 'exit data'] Evaluate 'copyfrom' individually for 'GOMP_MAP_STRUCT' entries
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.
2020-06-04 19:29:27 +02:00
Thomas Schwinge
a02f1adbfe [OpenACC 'exit data'] Evaluate 'finalize' individually for 'GOMP_MAP_STRUCT' entries
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.
2020-06-04 19:29:08 +02:00
Thomas Schwinge
db7179ec74 Fix 'sizeof' usage in 'libgomp.oacc-c-c++-common/deep-copy-{7,8}.c'
libgomp/
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c: Fix 'sizeof'
	usage.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: Likewise.
2020-06-04 18:56:37 +02:00
Thomas Schwinge
06ec61726d [OpenACC] Repair/restore 'is_tgt_unmapped' checking
libgomp/
	* oacc-mem.c (goacc_exit_datum): Repair 'is_tgt_unmapped'
	checking.
	(acc_unmap_data, goacc_exit_data_internal): Restore
	'is_tgt_unmapped' checking.
	* testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c: New
	file.
	* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Adjust.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Likewise.

Co-Authored-By: Julian Brown <julian@codesourcery.com>
2020-06-04 18:56:37 +02:00
Thomas Schwinge
2112d3242f [OpenACC] Don't open-code 'gomp_remove_var' in 'acc_unmap_data'
libgomp/
	* oacc-mem.c (acc_unmap_data): Don't open-code 'gomp_remove_var'.

Co-Authored-By: Julian Brown <julian@codesourcery.com>
2020-06-04 18:56:37 +02:00
Thomas Schwinge
4662f7fe78 [OpenACC] Remove 'tgt' reference counting from 'acc_unmap_data' [PR92854]
libgomp/
	PR libgomp/92854
	* oacc-mem.c (acc_unmap_data): Remove 'tgt' reference counting.
2020-06-04 18:56:37 +02:00
Thomas Schwinge
af8fd1a99d Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854]
libgomp/
	PR libgomp/92854
	* testsuite/libgomp.oacc-c-c++-common/pr92854-1.c: Extend some
	more.
2020-06-04 18:56:37 +02:00
Thomas Schwinge
f233418ccf [OpenACC] Use 'tgt' returned from 'gomp_map_vars'
libgomp/
	* oacc-mem.c (goacc_enter_datum): Use 'tgt' returned from
	'gomp_map_vars'.
	 (acc_map_data): Clean up accordingly.

Co-Authored-By: Julian Brown <julian@codesourcery.com>
2020-06-04 18:56:37 +02:00
Thomas Schwinge
8d7794c0a2 [OpenACC] XFAIL behavior of over-eager 'finalize' clause
libgomp/
	* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: XFAIL behavior
	of over-eager 'finalize' clause.
	* testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: New
	file.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-2.f90: Likewise.
2020-06-04 18:56:37 +02:00
Thomas Schwinge
2e24d457d8 [OpenACC] Missing unlocking on error paths in attach/detach code, part II
libgomp/
	* oacc-mem.c (goacc_exit_data_internal): Unlock on error path.
2020-06-04 18:56:37 +02:00
Julian Brown
dc9541545d [OpenACC] Missing unlocking on error paths in attach/detach code
libgomp/
	* oacc-mem.c (acc_attach_async): Add missing gomp_mutex_unlock on
	error path.
	(goacc_detach_internal): Likewise.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
2020-06-04 18:56:36 +02:00
Thomas Schwinge
1e378edd8f 'libgomp.oacc-fortran/{error_,}stop-{1,2,3}.f': initialize before the checkpoint
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.
2020-06-04 18:56:36 +02:00
GCC Administrator
a9312a7926 Daily bump. 2020-06-03 00:16:34 +00:00
Jakub Jelinek
23438370f7 openmp: Fix up build if HAVE_SYNC_BUILTINS is not defined.
2020-06-02  Jakub Jelinek  <jakub@redhat.com>

	* allocator.c (omp_free): Fix up build if HAVE_SYNC_BUILTINS is not
	defined.
2020-06-02 11:56:59 +02:00
GCC Administrator
945b53ccf6 Daily bump. 2020-05-31 00:16:22 +00:00
Jakub Jelinek
05e4db63d0 openmp: omp_alloc(0, ...) should return NULL.
2020-05-30  Jakub Jelinek  <jakub@redhat.com>

	* allocator.c (omp_alloc): For size == 0, return NULL early.

	* testsuite/libgomp.c-c++-common/alloc-4.c: New test.
2020-05-30 14:02:56 +02:00
GCC Administrator
885ef72f27 Daily bump. 2020-05-30 00:16:27 +00:00
H.J. Lu
9051b54827 Avoid nested save_CFLAGS and save_LDFLAGS
Avoid nested save_CFLAGS and save_LDFLAGS by replacing save_CFLAGS and
save_LDFLAGS with cet_save_CFLAGS and cet_save_LDFLAGS in cet.m4.

config/

	PR bootstrap/95413
	* cet.m4: Replace save_CFLAGS and save_LDFLAGS with
	cet_save_CFLAGS and cet_save_LDFLAGS.

gcc/

	PR bootstrap/95413
	* configure: Regenerated.

libatomic/

	PR bootstrap/95413
	* configure: Regenerated.

libbacktrace/

	PR bootstrap/95413
	* configure: Regenerated.

libcc1/

	PR bootstrap/95413
	* configure: Regenerated.

libcpp/

	PR bootstrap/95413
	* configure: Regenerated.

libdecnumber/

	PR bootstrap/95413
	* configure: Regenerated.

libgcc/

	PR bootstrap/95413
	* configure: Regenerated.

libgfortran/

	PR bootstrap/95413
	* configure: Regenerated.

libgomp/

	PR bootstrap/95413
	* configure: Regenerated.

libiberty/

	PR bootstrap/95413
	* configure: Regenerated.

libitm/

	PR bootstrap/95413
	* configure: Regenerated.

libobjc/

	PR bootstrap/95413
	* configure: Regenerated.

libphobos/

	PR bootstrap/95413
	* configure: Regenerated.

libquadmath/

	PR bootstrap/95413
	* configure: Regenerated.

libsanitizer/

	PR bootstrap/95413
	* configure: Regenerated.

libssp/

	PR bootstrap/95413
	* configure: Regenerated.

libstdc++-v3/

	PR bootstrap/95413
	* configure: Regenerated.

libvtv/

	PR bootstrap/95413
	* configure: Regenerated.

lto-plugin/

	PR bootstrap/95413
	* configure: Regenerated.

zlib/

	PR bootstrap/95413
	* configure: Regenerated.
2020-05-29 12:56:40 -07:00
Thomas Koenig
8df7ee67f6 Fixes a hang on an invalid ID in a WAIT statement.
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.
2020-05-23 19:01:43 +02:00
Jakub Jelinek
e107157171 openmp: Add basic library allocator support.
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>

	* allocator.c: New file.
2020-05-19 14:08:11 +02:00
Jakub Jelinek
800bcc8c00 openmp: Add basic library allocator support.
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-19 10:11:01 +02:00
H.J. Lu
4c1a5d8b71 x86: Also check if -fcf-protection works
When defaulting CET run-time support to auto, check if -fcf-protection
works.  Even if the stage1 GCC doesn't support -fcf-protection, since
the final GCC does, CET run-time support will be enabled by default if
binutils support CET.

config/

	PR bootstrap/95147
	* cet.m4 (GCC_CET_FLAGS): Also check if -fcf-protection works
	when defaulting to auto.

libatomic/

	PR bootstrap/95147
	* configure: Regenerated.

libbacktrace/

	PR bootstrap/95147
	* configure: Regenerated.

libgcc/

	PR bootstrap/95147
	* configure: Regenerated.

libgfortran/

	PR bootstrap/95147
	* configure: Regenerated.

libgomp/

	PR bootstrap/95147
	* configure: Regenerated.

libitm/

	PR bootstrap/95147
	* configure: Regenerated.

libobjc/

	PR bootstrap/95147
	* configure: Regenerated.

libphobos/

	PR bootstrap/95147
	* configure: Regenerated.

libquadmath/

	PR bootstrap/95147
	* configure: Regenerated.

libsanitizer/

	PR bootstrap/95147
	* configure: Regenerated.

libssp/

	PR bootstrap/95147
	* configure: Regenerated.

libstdc++-v3/

	PR bootstrap/95147
	* configure: Regenerated.

libvtv/

	PR bootstrap/95147
	* configure: Regenerated.

zlib/

	PR bootstrap/95147
	* configure: Regenerated.
2020-05-15 09:07:17 -07:00
Thomas Koenig
d975519ad1 Removed double ChangeLog entries from previous commit. 2020-05-14 18:37:18 +02:00
Thomas Koenig
cdc34b5057 Add early return for invalid STATUS for close.
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.
2020-05-14 18:33:24 +02:00
H.J. Lu
8d286dd118 x86: Default CET run-time support to auto
CET has been added since GCC 8.  This patch defaults CET run-time support
to auto.  It enables CET run-time support if asssembler supports CET
instructions and multi-byte NOPs are enabled via SSE2.

config/

	* cet.m4 (GCC_CET_FLAGS): Change default to auto.

gcc/

	* configure: Regenerated.

libatomic/

	* configure: Regenerated.

libbacktrace/

	* configure: Regenerated.

libcc1/

	* configure: Regenerated.

libcpp/

	* configure: Regenerated.

libdecnumber/

	* configure: Regenerated.

libgcc/

	* configure: Regenerated.

libgfortran/

	* configure: Regenerated.

libgomp/

	* configure: Regenerated.

libitm/

	* configure: Regenerated.

libobjc/

	* configure: Regenerated.

libquadmath/

	* configure: Regenerated.

libsanitizer/

	* configure: Regenerated.

libssp/

	* configure: Regenerated.

libstdc++-v3/

	* configure: Regenerated.

libvtv/

	* configure: Regenerated.

zlib/

	* configure: Regenerated.
2020-05-14 09:05:02 -07:00
Jakub Jelinek
49ddde69fc openmp: Also implicitly mark as declare target to functions mentioned in target regions
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.
2020-05-14 09:48:32 +02:00
Tobias Burnus
f884bef21c [Fortran] OpenMP - permit lastprivate in distribute + SIMD fixes (PR94690)
gcc/fortran/
2020-05-13  Tobias Burnus  <tobias@codesourcery.com>

	PR fortran/94690
        * openmp.c (OMP_DISTRIBUTE_CLAUSES): Add OMP_CLAUSE_LASTPRIVATE.
        (gfc_resolve_do_iterator): Skip the private handling for SIMD as
        that is handled by ME code.
	* trans-openmp.c (gfc_trans_omp_do): Don't add private/lastprivate
	for dovar_found == 0, unless !simple.

libgomp/
2020-05-13  Tobias Burnus  <tobias@codesourcery.com>

	PR fortran/94690
	* testsuite/libgomp.fortran/pr66199-3.f90: New.
	* testsuite/libgomp.fortran/pr66199-4.f90: New.
	* testsuite/libgomp.fortran/pr66199-5.f90: New.
	* testsuite/libgomp.fortran/pr66199-6.f90: New.
	* testsuite/libgomp.fortran/pr66199-7.f90: New.
	* testsuite/libgomp.fortran/pr66199-8.f90: New.
	* testsuite/libgomp.fortran/pr66199-9.f90: New.
2020-05-13 10:06:45 +02:00
Jakub Jelinek
dc703151d4 openmp: Implement discovery of implicit declare target to clauses
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.
2020-05-12 09:17:09 +02:00
Thomas Schwinge
a5d0bc12e1 [OpenACC] Set 'acc_device_current = -1'
There's no point in using value '-3', and even though not directly related,
value '-1' does match 'GOMP_DEVICE_ICV'.

	libgomp/
	* config/accel/openacc.f90 (acc_device_current): Set to '-1'.
	* openacc.f90 (acc_device_current): Likewise.
	* openacc.h (acc_device_current): Likewise.
	* openacc_lib.h (acc_device_current): Likewise.
2020-04-29 09:54:37 +02:00
Thomas Schwinge
7f1989249e [gcn] Set 'UI_NONE' for 'TARGET_EXCEPT_UNWIND_INFO' [PR94282]
In libgomp offloading testing, this resolves all the 'ld: error: undefined
symbol: __gxx_personality_v0' FAILs.

	gcc/
	PR target/94282
	* common/config/gcn/gcn-common.c (gcn_except_unwind_info): New
	function.
	(TARGET_EXCEPT_UNWIND_INFO): Define.
	libgomp/
	PR target/94282
	* testsuite/libgomp.c-c++-common/function-not-offloaded.c: Remove
	'dg-allow-blank-lines-in-output'.
2020-04-29 09:39:03 +02:00
Thomas Schwinge
4912a04f8b [gcn] Use 'radeon' for the environment variable 'ACC_DEVICE_TYPE'
..., per OpenACC 3.0, A.1.2. "AMD GPU Targets".

This complements commit 6687d13a87 "Rename
acc_device_gcn to acc_device_radeon".

	libgomp/
	* oacc-init.c (get_openacc_name): Handle 'gcn'.
	* testsuite/lib/libgomp.exp
	(offload_target_to_openacc_device_type) [amdgcn*]: Return
	'radeon'.  Adjust all users.
	(check_effective_target_openacc_amdgcn_accel_present): Rename
	to...
	(check_effective_target_openacc_radeon_accel_present): ... this.
	Adjust all users.
	(check_effective_target_openacc_amdgcn_accel_selected): Rename to...
	(check_effective_target_openacc_radeon_accel_selected): ... this.
	Adjust all users.
2020-04-29 09:24:07 +02:00
Thomas Schwinge
b9dc11b673 Torture testing: 'libgomp.fortran/use_device_ptr-optional-2.f90'
Fix-up for commit a2c26c5031 (r278046) "Fortran]
Support absent optional args with use_device_{ptr,addr}".

	libgomp/
	* testsuite/libgomp.fortran/use_device_ptr-optional-2.f90: Add
	'dg-do run'.
2020-04-29 09:24:07 +02:00
Andrew Stubbs
966de09be9 amdgcn: Check HSA return codes [PR94629]
Ensure that the returned status values are not ignored.  The old code was
not broken, but this is both safer and satisfies static analysis.

2020-04-23  Andrew Stubbs  <ams@codesourcery.com>

	PR other/94629

	libgomp/
	* plugin/plugin-gcn.c (init_hsa_context): Check return value from
	hsa_iterate_agents.
	(GOMP_OFFLOAD_init_device): Check return values from both calls to
	hsa_agent_iterate_regions.
2020-04-23 15:05:23 +01:00
Thomas Schwinge
3f5d94c192 Add 'dg-do run' to 'libgomp.fortran/target-enter-data-2.F90'
Fix-up for commit af557050fd "[OpenMP] Fix 'omp
exit data' for Fortran arrays (PR 94635)".

	libgomp/
	PR middle-end/94635
	* testsuite/libgomp.fortran/target-enter-data-2.F90: Add 'dg-do
	run'.
2020-04-20 23:16:40 +02:00
Tobias Burnus
85d8c05a02 Fix declare copyout in libgomp.oacc-c++/declare-pr94120.C
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.
2020-04-20 12:38:50 +02:00
Tobias Burnus
af557050fd [OpenMP] Fix 'omp exit data' for Fortran arrays (PR 94635)
PR middle-end/94635
	* gimplify.c (gimplify_scan_omp_clauses): Turn MAP_TO_PSET to
	MAP_DELETE.

	PR middle-end/94635
	* testsuite/libgomp.fortran/target-enter-data-2.F90: New.
2020-04-17 19:08:55 +02:00
Thomas Schwinge
af4c92573d Rename 'libgomp.oacc-c-c++-common/static-dynamic-lifetimes-*' to 'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-*' [PR92843]
Fix-up for commit be9862dd96 "Test cases for
mixed structured/dynamic data lifetimes with OpenACC [PR92843]": it's
"structured", not "static" data lifetimes/reference counters.

	libgomp/
	PR libgomp/92843
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c:
	... this.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c::
	Rename to...
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c:
	... this.
2020-04-13 08:56:03 +02:00
Julian Brown
be9862dd96 Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843]
libgomp/
	PR libgomp/92843
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c:
	New file.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c:
	Likewise.
2020-04-10 15:53:04 +02:00
Thomas Schwinge
6b816a5f0e Add 'dg-do run' to 'libgomp.fortran/target-enter-data-1.f90'
Fix-up for commit 689418b97e "libgomp – fix
handling of 'target enter data'".

	libgomp/
	* testsuite/libgomp.fortran/target-enter-data-1.f90: Add 'dg-do
	run'.
2020-04-10 15:44:17 +02:00
Tobias Burnus
13e41d8b9d [C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)
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.
2020-04-08 09:39:43 +02:00
Maciej W. Rozycki
749bd22ddc libgomp/test: Remove a build sysroot fix regression
Fix a problem with commit c8e759b421 ("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.
2020-04-06 23:32:45 +01:00
Thomas Schwinge
2b1e849b35 Revert "[nvptx, libgomp] Update pr85381-{2,4}.c test-cases" [PR89713, PR94392]
In response to PR94392 commit 75efe9cb1f
"c/94392 - only enable -ffinite-loops for C++", this reverts PR89713
commit 00908992f2, 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.
2020-04-03 10:10:07 +02:00
Tobias Burnus
689418b97e libgomp – fix handling of 'target enter data'
* target.c (GOMP_target_enter_exit_data): Handle PSET/MAP_POINTER.
	* testsuite/libgomp.fortran/target-enter-data-1.f90: New.
2020-03-31 20:38:38 +02:00
Tobias Burnus
c2211a60ff Fix OpenMP offload handling for target-link variables for nvptx (PR81689)
PR libgomp/81689
	* lto.c (offload_handle_link_vars): Propagate TREE_PUBLIC state.

	PR libgomp/81689
	* omp-offload.c (omp_finish_file): Fix target-link handling if
	targetm_common.have_named_sections is false.

	PR libgomp/81689
	* testsuite/libgomp.c/target-link-1.c: Remove xfail.
2020-03-24 15:13:56 +01:00
Tobias Burnus
4897bb0045 libgomp – fix declare target link handling (PR94251)
PR libgomp/94251
	* target.c (gomp_load_image_to_device): Fix link
	variable handling.
2020-03-23 15:29:05 +01:00
Jakub Jelinek
02f7334ac9 c++: Fix up handling of captured vars in lambdas in OpenMP clauses [PR93931]
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-19 12:22:47 +01:00
Tobias Burnus
bb83e069eb libgomp/testsuite: ignore blank-line output for function-not-offloaded.c
* testsuite/libgomp.c-c++-common/function-not-offloaded.c: Add
	dg-allow-blank-lines-in-output.
2020-03-19 11:42:49 +01:00
Tobias Burnus
26cbcfe5fc Fix libgomp.oacc-fortran/atomic_capture-1.f90
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.
2020-03-18 16:28:08 +01:00
Tobias Burnus
4da9288745 libgomp testsuite - disable long double for AMDGCN
* 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.
2020-03-18 12:07:54 +01:00
Jakub Jelinek
9c3cdb43c2 tree-nested: Fix handling of *reduction clauses with C array sections [PR93566]
tree-nested.c didn't handle C array sections in {,task_,in_}reduction clauses.

2020-03-14  Jakub Jelinek  <jakub@redhat.com>

	PR middle-end/93566
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Handle {,in_,task_}reduction clauses
	with C/C++ array sections.

	* testsuite/libgomp.c/pr93566.c: New test.
2020-03-15 01:27:40 +01:00
Frederik Harwath
83d45e1d71 Adapt libgomp acc_get_property.f90 test
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.
2020-02-21 15:38:48 +01:00
Tobias Burnus
8d1a1cb1b8 libgomp: Fixes + cleanup for OpenACC's Fortran module + openacc_lib.h
2020-02-19  Tobias Burnus  <tobias@codesourcery.com>

	* .gitattributes: New; whitespace handling for Fortran's openacc_lib.h.
	* config/accel/openacc.f90 (openacc_kinds): Add acc_device_current.
	(openacc_internal, acc_on_device_h): Fix argument name; minor cleanup.
	* libgomp.texi (Enabling OpenACC): No longer mark as experimental.
	(acc_set_device_num): Fix Fortran argument name, use same name for C.
	(acc_get_property): Update Fortran interface to post-OpenACC 3.0
	corrections; add note about the previous interface and named constant.
	(OpenACC library and environment variables): Fix two typos.
	* openacc.f90: Use for all procedures the argument names from the spec
	as for …_h they are user visible.
	(openacc_kinds): Rename acc_device_property to
	acc_device_property_kinds and change value to int32 ; and update users.
	Re-add acc_device_property for for backward compatibility.
	(acc_get_property_string_h): Clean up as acc_device_property_kind
	changed.
	(acc_get_property_h): Likewise and return c_size_t instead of
	acc_device_property.
	(openacc): Also export acc_device_property_kinds.
	(acc_async_test_h, acc_async_test_all_h, acc_on_device_h,
	acc_is_present_32_h, acc_is_present_64_h): Simplify logical-return-value
	handling; check against /= 0 instead of == 1 to match C.
	* openacc_lib.h: Use for all procedures the argument names from the spec
	as for …_h they are user visible. Place !GCC$ into the first column to
	be active also for fixed-form souce form.
	(acc_device_current, acc_device_property_kind, acc_device_property,
	acc_property_memory, acc_property_free_memory, acc_property_name,
	acc_property_vendor, acc_property_driver): New named constants.
	(acc_get_property, acc_get_property_string): New generic interface.
2020-02-19 09:13:44 +01:00
Frederik Harwath
2d9eb4e4ca Add ChangeLog entry for my last commit 2020-02-13 10:26:13 +01:00
Frederik Harwath
001ab12e62 openmp: ignore nowait if async execution is unsupported [PR93481]
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.
2020-02-13 10:18:31 +01:00
Frederik Harwath
fd789c816b Add xfails to libgomp tests target-{33,34}.c, target-link-1.c
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.
2020-02-10 09:16:46 +01:00
Jakub Jelinek
9bc3b95dfe openmp: Optimize DECL_IN_CONSTANT_POOL vars in target regions
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.
2020-02-09 08:17:10 +01:00
Jakub Jelinek
cb3f06480a openmp: Fix handling of non-addressable shared scalars in parallel nested inside of target [PR93515]
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.
2020-02-06 09:19:08 +01:00
Tobias Burnus
91bc3c9885 [libgomp] – Fix check_effective_target_offload_target_nvptx for remote execution
* testsuite/lib/libgomp.exp
	(check_effective_target_offload_target_nvptx): Pass flags as 'options'
	and not as 'source' argument to libgomp_target_compile.
2020-02-05 17:40:48 +01:00
Andrew Stubbs
591f869ad7 Remove gfx801 "carrizo" support
2020-02-03  Andrew Stubbs  <ams@codesourcery.com>

	gcc/
	* config.gcc: Remove "carrizo" support.
	* config/gcn/gcn-opts.h (processor_type): Likewise.
	* config/gcn/gcn.c (gcn_omp_device_kind_arch_isa): Likewise.
	* config/gcn/gcn.opt (gpu_type): Likewise.
	* config/gcn/t-omp-device: Likewise.

	libgomp/
	* plugin/plugin-gcn.c (EF_AMDGPU_MACH_AMDGCN_GFX801): Remove.
	(gcn_gfx801_s): Remove.
	(isa_hsa_name): Remove gfx801.
	(isa_gcc_name): Remove gfx801/carizzo.
	(isa_code): Remove gfx801.
2020-02-03 17:23:18 +00:00
Tobias Burnus
e464fc9035 [OpenACC] bump version for 2.6 plus libgomp.texi update
2020-02-03  Julian Brown  <julian@codesourcery.com>
            Tobias Burnus  <tobias@codesourcery.com>

	gcc/c-family/
	* c-cppbuiltin.c (c_cpp_builtins): Update _OPENACC define to 201711.

	gcc/
	* doc/invoke.texi: Update mention of OpenACC version to 2.6.

	gcc/fortran/
	* cpp.c (cpp_define_builtins): Update _OPENACC define to 201711.
	* intrinsic.texi: Update mentions of OpenACC version to 2.6.
	* gfortran.texi: Likewise. Remove experimental disclamer for OpenACC.
	* invoke.texi: Remove experimental disclamer for OpenACC.

	gcc/testsuite/
	* c-c++-common/cpp/openacc-define-3.c: Update expected value for
	_OPENACC define.
	* gfortran.dg/openacc-define-3.f90: Likewise.

	libgomp/
	* libgomp.texi (OpenACC Runtime Library Routines): Document *_async
	and *_finalize variants; document acc_attach and acc_detach; update
	references from OpenACC 2.0 to 2.6.
	* openacc.f90 (openacc_version): Update to 201711.
	* openacc_lib.h (openacc_version): Update to 201711.
	* testsuite/libgomp.oacc-fortran/openacc_version-1.f: Update expected
	openacc_version to 201711.
	* testsuite/libgomp.oacc-fortran/openacc_version-2.f90: Likewise.
2020-02-03 10:10:37 +01:00
Kwok Cheung Yeung
5a28e2727f [amdgcn] Scale number of threads/workers with VGPR usage
2020-01-31  Kwok Cheung Yeung  <kcy@codesourcery.com>

	gcc/
	* config/gcn/mkoffload.c (process_asm): Add sgpr_count and vgpr_count
	to definition of hsa_kernel_description.  Parse assembly to find SGPR
	and VGPR count of kernel and store in hsa_kernel_description.

	libgomp/
	* plugin/plugin-gcn.c (struct hsa_kernel_description): Add sgpr_count
	and vgpr_count fields.
	(struct kernel_info): Add a field for a hsa_kernel_description.
	(run_kernel): Reduce the number of threads/workers if the requested
	number would require too many VGPRs.
	(init_basic_kernel_info): Initialize description field with
	the hsa_kernel_description entry for the kernel.
2020-01-31 07:13:05 -08:00
Tobias Burnus
5ab5d81b36 Skip plugin-{gcn,hsa} for (-m)x32 (PR bootstrap/93409)
PR bootstrap/93409
        * plugin/configfrag.ac (enable_offload_targets): Skip
        HSA and GCN plugin besides -m32 also for -mx32.
        * configure: Regenerate.
2020-01-30 12:27:17 +01:00
Frederik Harwath
989a5fb3aa Add acc_device_radeon to name_of_acc_device_t function
libgomp/
	* oacc-init.c (name_of_acc_device_t): Handle acc_device_radeon.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
2020-01-29 15:52:58 +01:00
Frederik Harwath
87c3fcfa6b Adjust formatting of acc_get_property tests
libgomp/
 * testsuite/libgomp.oacc-c-c++-common/acc_get_property.c:
 Adjust to GNU coding style.
 * testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c: Likewise.
 * testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c: Likewise.
 * testsuite/libgomp.oacc-c-c++-common/acc_get_property-host.c: Likewise.
 * testsuite/libgomp.oacc-c-c++-common/acc_get_property-nvptx.c: Likewise.
2020-01-29 11:55:06 +01:00
Frederik Harwath
2e5ea57959 Add OpenACC acc_get_property support for AMD GCN
Add full support for the OpenACC 2.6 acc_get_property and
acc_get_property_string functions to the libgomp GCN plugin.

libgomp/
	* plugin-gcn.c (struct agent_info): Add fields "name" and
	"vendor_name" ...
	(GOMP_OFFLOAD_init_device): ... and init from here.
	(struct hsa_context_info): Add field "driver_version_s" ...
	(init_hsa_contest): ... and init from here.
	(GOMP_OFFLOAD_openacc_get_property): Replace stub with a proper
	implementation.
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property.c:
	Enable test execution for amdgcn and host offloading targets.
	* testsuite/libgomp.oacc-fortran/acc_get_property.f90: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c
	(expect_device_properties): Split function into ...
	(expect_device_string_properties): ... this new function ...
	(expect_device_memory): ... and this new function.
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c:
	Add test.
2020-01-29 11:54:56 +01:00
Julian Brown
278c3214b3 Don't allow mixed component and non-component accesses for OpenACC/Fortran
gcc/fortran/
	* gfortran.h (gfc_symbol): Add comp_mark bitfield.
	* openmp.c (resolve_omp_clauses): Disallow mixed component and
	full-derived-type accesses to the same variable within a single
	directive.

	libgomp/
	* testsuite/libgomp.oacc-fortran/deep-copy-2.f90: Remove test from here.
	* testsuite/libgomp.oacc-fortran/deep-copy-3.f90: Don't use mixed
	component/non-component variable refs in a single directive.
	* testsuite/libgomp.oacc-fortran/classtypes-1.f95: Likewise.

	gcc/testsuite/
	* gfortran.dg/goacc/deep-copy-2.f90: Move test here (from libgomp
	testsuite). Make a compilation test, and expect rejection of mixed
	component/non-component accesses.
	* gfortran.dg/goacc/mapping-tests-1.f90: New test.
2020-01-28 06:00:29 -08:00
Maciej W. Rozycki
e8e66971cd Add `--with-toolexeclibdir=' configuration option
Provide means, in the form of a `--with-toolexeclibdir=' configuration
option, to override the default installation directory for target
libraries, otherwise known as $toolexeclibdir.  This is so that it is
possible to get newly-built libraries, particularly the shared ones,
installed in a common place, so that they can be readily used by the
target system as their host libraries, possibly over NFS, without a need
to manually copy them over from the currently hardcoded location they
would otherwise be installed in.

In the presence of the `--enable-version-specific-runtime-libs' option
and for configurations building native GCC the option is ignored.

	config/
	* toolexeclibdir.m4: New file.

	gcc/
	* doc/install.texi (Cross-Compiler-Specific Options): Document
	`--with-toolexeclibdir' option.

	libada/
	* Makefile.in (configure_deps): Add `toolexeclibdir.m4'.
	* configure.ac: Handle `--with-toolexeclibdir='.
	* configure: Regenerate.

	libatomic/
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.
	* testsuite/Makefile.in: Regenerate.

	libffi/
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.
	* include/Makefile.in: Regenerate.
	* man/Makefile.in: Regenerate.
	* testsuite/Makefile.in: Regenerate.

	libgcc/
	* Makefile.in (configure_deps): Add `toolexeclibdir.m4'.
	* configure.ac: Handle `--with-toolexeclibdir='.
	* configure: Regenerate.

	libgfortran/
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.

	libgomp/
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.
	* testsuite/Makefile.in: Regenerate.

	libhsail-rt/
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.

	libitm/
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.
	* testsuite/Makefile.in: Regenerate.

	libobjc/
	* Makefile.in (aclocal_deps): Add `toolexeclibdir.m4'.
	* aclocal.m4: Include `toolexeclibdir.m4'.
	* configure.ac: Handle `--with-toolexeclibdir='.
	* configure: Regenerate.

	liboffloadmic/
	* plugin/configure.ac: Handle `--with-toolexeclibdir='.
	* plugin/Makefile.in: Regenerate.
	* plugin/aclocal.m4: Regenerate.
	* plugin/configure: Regenerate.
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.

	libphobos/
	* m4/druntime.m4: Handle `--with-toolexeclibdir='.
	* m4/Makefile.in: Regenerate.
	* libdruntime/Makefile.in: Regenerate.
	* src/Makefile.in: Regenerate.
	* testsuite/Makefile.in: Regenerate.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.

	libquadmath/
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.

	libsanitizer/
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.
	* asan/Makefile.in: Regenerate.
	* interception/Makefile.in: Regenerate.
	* libbacktrace/Makefile.in: Regenerate.
	* lsan/Makefile.in: Regenerate.
	* sanitizer_common/Makefile.in: Regenerate.
	* tsan/Makefile.in: Regenerate.
	* ubsan/Makefile.in: Regenerate.

	libssp/
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.

	libstdc++-v3/
	* acinclude.m4: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.
	* doc/Makefile.in: Regenerate.
	* include/Makefile.in: Regenerate.
	* libsupc++/Makefile.in: Regenerate.
	* po/Makefile.in: Regenerate.
	* python/Makefile.in: Regenerate.
	* src/Makefile.in: Regenerate.
	* src/c++11/Makefile.in: Regenerate.
	* src/c++17/Makefile.in: Regenerate.
	* src/c++98/Makefile.in: Regenerate.
	* src/filesystem/Makefile.in: Regenerate.
	* testsuite/Makefile.in: Regenerate.

	libvtv/
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.
	* testsuite/Makefile.in: Regenerate.

	zlib/
	* configure.ac: Handle `--with-toolexeclibdir='.
	* Makefile.in: Regenerate.
	* aclocal.m4: Regenerate.
	* configure: Regenerate.
2020-01-24 11:24:25 +00:00
Frederik Harwath
9e424d9749 Add missing ChangeLog entry for my last commit
libgomp/ChangeLog: Add entry for commit 4bd03ed69b.
2020-01-24 09:25:35 +01:00
Frederik Harwath
4bd03ed69b Fix expectation and types in acc_get_property tests
* Weaken expectation concerning acc_property_free_memory.
  Do not expect the value returned by CUDA since that value might have
  changed in the meantime.
* Use correct type for the results of calls to acc_get_property in tests.

libgomp/
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c
	(expect_device_properties): Remove "expected_free_mem" argument,
	change "expected_total_mem" argument type to size_t;
	change types of acc_get_property results to size_t,
	adapt format strings.
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property.c:
	Use %zu instead of %zd to print size_t values.
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property-2.c: Adapt and
	rename to ...
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property-nvptx.c: ... this.
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property-3.c: Adapt and
	rename to ...
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property-host.c: ... this.

Reviewed-by: Thomas Schwinge  <thomas@codesourcery.com>
2020-01-24 09:14:51 +01:00
Andrew Stubbs
14e5e74698 Fix libgomp plugin-gcn bug
2020-01-23  Andrew Stubbs  <ams@codesourcery.com>

	libgomp/
	* plugin/plugin-gcn.c (parse_target_attributes): Use correct mask for
	the device id.
2020-01-23 12:40:04 +00:00
Frederik Harwath
7d593fd672 Add runtime ISA check for amdgcn offloading
The HSA/ROCm runtime rejects binaries not built for the exact GPU device
present. So far, the libgomp amdgcn plugin does not verify that the GPU ISA
and the ISA specified at compile time match before handing over the binary to
the runtime.  In case of a mismatch, the user is confronted with an unhelpful
runtime error.

This commit implements a runtime ISA check. In case of an ISA mismatch, the
execution is aborted with a clear error message and a hint at the correct
compilation parameters for the GPU on which the execution has been attempted.

libgomp/
	* plugin/plugin-gcn.c (EF_AMDGPU_MACH): New enum.
	* (EF_AMDGPU_MACH_MASK): New constant.
	* (gcn_isa): New typedef.
	* (gcn_gfx801_s): New constant.
	* (gcn_gfx803_s): New constant.
	* (gcn_gfx900_s): New constant.
	* (gcn_gfx906_s): New constant.
	* (gcn_isa_name_len): New constant.
	* (elf_gcn_isa_field): New function.
	* (isa_hsa_name): New function.
	* (isa_gcc_name): New function.
	* (isa_code): New function.
	* (struct agent_info): Add field "device_isa" and remove field
	"gfx900_p".
	* (GOMP_OFFLOAD_init_device): Adapt agent init to "agent_info"
	field changes, fail if device has unknown ISA.
	* (parse_target_attributes): Replace "gfx900_p" by "device_isa".
	* (isa_matches_agent): New function ...
	* (create_and_finalize_hsa_program): ... used from here to check
	that the GPU ISA and the code-object ISA match.
2020-01-21 07:41:45 +01:00
Andrew Stubbs
09e0ad6253 Update OpenACC tests for amdgcn
2020-01-20  Andrew Stubbs  <ams@codesourcery.com>

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Skip test on gcn.
	* testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c (main):
	Adjust test dimensions for amdgcn.
	* testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c (main): Adjust
	gang/worker/vector expectations dynamically.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c
	(main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c (main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c (main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c (main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c (main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c (main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-v-1.c (main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-w-1.c (main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c (main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
	(acc_gang): Recognise acc_device_radeon.
	(acc_worker): Likewise.
	(acc_vector): Likewise.
	(main): Set expectations for amdgcn.
	* testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c
	(main): Adjust gang/worker/vector expectations dynamically.
	* testsuite/libgomp.oacc-c-c++-common/routine-v-1.c (main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-w-1.c (main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c (main): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c: Set expectations
	for amdgcn.
2020-01-20 16:51:06 +00:00
Andrew Stubbs
6687d13a87 Rename acc_device_gcn to acc_device_radeon
2020-01-17  Andrew Stubbs  <ams@codesourcery.com>

	libgomp/
	* config/accel/openacc.f90 (openacc_kinds): Rename acc_device_gcn to
	acc_device_radeon.
	(openacc): Likewise.
	* openacc.f90 (openacc_kinds): Likewise.
	(openacc): Likewise.
	* openacc.h (acc_device_t): Likewise.
	* openacc_lib.h: Likewise.
	* testsuite/lib/libgomp.exp
	(check_effective_target_openacc_amdgcn_accel_present): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c
	(cb_compute_construct_end): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
	(cb_enqueue_launch_start): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c
	(cb_enter_data_end): Likewise.
	(cb_exit_data_start): Likewise.
	(cb_exit_data_end): Likewise.
	(cb_compute_construct_end): Likewise.
	(cb_enqueue_launch_start): Likewise.
	(cb_enqueue_launch_end): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c
	(main): Likewise.
2020-01-17 18:11:52 +00:00
Thomas Schwinge
6fc0385c0c OpenACC 'acc_get_property' cleanup
include/
	* gomp-constants.h (enum gomp_device_property): Remove.
	libgomp/
	* libgomp-plugin.h (enum goacc_property): New.  Adjust all users
	to use this instead of 'enum gomp_device_property'.
	(GOMP_OFFLOAD_get_property): Rename to...
	(GOMP_OFFLOAD_openacc_get_property): ... this.  Adjust all users.
	* libgomp.h (struct gomp_device_descr): Move
	'GOMP_OFFLOAD_openacc_get_property'...
	(struct acc_dispatch_t): ... here.  Adjust all users.
	* plugin/plugin-hsa.c (GOMP_OFFLOAD_get_property): Remove.
	liboffloadmic/
	* plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_property):
	Remove.

From-SVN: r280150
2020-01-10 23:24:36 +01:00
Thomas Schwinge
b3b75e664a Further changes for the OpenACC 'if_present' clause on the 'host_data' construct
gcc/
	* tree.h (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT): New definition.
	* tree-core.h: Document it.
	* gimplify.c (gimplify_omp_workshare): Set it.
	* omp-low.c (lower_omp_target): Use it.
	* tree-pretty-print.c (dump_omp_clause): Print it.
	gcc/testsuite/
	* c-c++-common/goacc/host_data-1.c: Extend.
	* gfortran.dg/goacc/host_data-tree.f95: Likewise.

	gcc/
	* omp-low.c (lower_omp_target) <OMP_CLAUSE_USE_DEVICE_PTR etc.>:
	Assert that for OpenACC we always have 'GOMP_MAP_USE_DEVICE_PTR'.

	libgomp/
	* target.c (gomp_map_vars_internal)
	<GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT>: Clean up/elaborate code
	paths.

From-SVN: r280149
2020-01-10 23:23:44 +01:00
Jakub Jelinek
91df4397a1 re PR libgomp/93219 (unused return value in affinity-fmt.c)
PR libgomp/93219
	* libgomp.h (gomp_print_string): Change return type from void to int.
	* affinity-fmt.c (gomp_print_string): Likewise.  Return true if
	not all characters have been written.

From-SVN: r280137
2020-01-10 21:42:00 +01:00
Tobias Burnus
d5c23c6cea OpenACC – support "if" + "if_present" clauses with "host_data"
2020-01-10  Gergö Barany  <gergo@codesourcery.com>
	    Thomas Schwinge <thomas@codesourcery.com>
	    Julian Brown  <julian@codesourcery.com>
	    Tobias Burnus  <tobias@codesourcery.com>

        gcc/c/
        * c-parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF
        and PRAGMA_OACC_CLAUSE_IF_PRESENT.

        gcc/cp/
        * parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF
        and PRAGMA_OACC_CLAUSE_IF_PRESENT.

        gcc/fortran/
        * openmp.c (OACC_HOST_DATA_CLAUSES): Add PRAGMA_OACC_CLAUSE_IF
        and PRAGMA_OACC_CLAUSE_IF_PRESENT.

	gcc/
	* omp-low.c (lower_omp_target): Use GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT
	if PRAGMA_OACC_CLAUSE_IF_PRESENT exist.

	gcc/testsuite/
	* c-c++-common/goacc/host_data-1.c: Added tests of if and if_present
	clauses on host_data.
	* gfortran.dg/goacc/host_data-tree.f95: Likewise.

	include/
	* gomp-constants.h (enum gomp_map_kind): New enumeration constant
	GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT.
        
	libgomp/
	* oacc-parallel.c (GOACC_data_start): Handle
	GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT.
	* target.c (gomp_map_vars_async): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/host_data-7.c: New.
	* testsuite/libgomp.oacc-fortran/host_data-5.F90: New.

From-SVN: r280115
2020-01-10 16:08:41 +01:00
Tobias Burnus
c1030b5cd3 libgomp.texi: Fix typos, use https (actual change)
From-SVN: r280008
2020-01-08 16:00:39 +01:00
Tobias Burnus
df2c0d3620 libgomp.texi: Fix typos, use https.
From-SVN: r280007
2020-01-08 15:58:31 +01:00
Tobias Burnus
f760c0c77f Fortran] OpenMP/OpenACC – fix more issues with OPTIONAL
gcc/fortran/
        * trans-openmp.c (gfc_omp_check_optional_argument): Always return a
        Boolean expression; handle unallocated/disassociated actual arguments
        as absent if passed to nonallocatable/nonpointer dummy array arguments.
        (gfc_build_cond_assign): Change to assume a Boolean expr not a pointer.
        (gfc_omp_finish_clause, gfc_trans_omp_clauses): Assign NULL to generated
        array-data variable if the argument is absent. Simplify code as
        'present' is now a Boolean expression.

        libgomp/
        * testsuite/libgomp.fortran/optional-map.f90: Add test for
        unallocated/disassociated actual arguments to nonallocatable/nonpointer
        dummy arguments; those are/shall be regarded as absent arguments.
        * testsuite/libgomp.fortran/use_device_ptr-optional-2.f90: Ditto.
        * testsuite/libgomp.fortran/use_device_ptr-optional-3.f90: New.

From-SVN: r279858
2020-01-03 13:56:46 +01:00
Jakub Jelinek
8d9254fc8a Update copyright years.
From-SVN: r279813
2020-01-01 12:51:42 +01:00
Jakub Jelinek
7e7065b976 gcc.c (process_command): Update copyright notice dates.
* gcc.c (process_command): Update copyright notice dates.
	* gcov-dump.c (print_version): Ditto.
	* gcov.c (print_version): Ditto.
	* gcov-tool.c (print_version): Ditto.
	* gengtype.c (create_file): Ditto.
	* doc/cpp.texi: Bump @copying's copyright year.
	* doc/cppinternals.texi: Ditto.
	* doc/gcc.texi: Ditto.
	* doc/gccint.texi: Ditto.
	* doc/gcov.texi: Ditto.
	* doc/install.texi: Ditto.
	* doc/invoke.texi: Ditto.
gcc/fortran/ 
	* gfortranspec.c (lang_specific_driver): Update copyright notice
	dates.
	* gfc-internals.texi: Bump @copying's copyright year.
	* gfortran.texi: Ditto.
	* intrinsic.texi: Ditto.
	* invoke.texi: Ditto.
gcc/d/
	* gdc.texi: Bump @copyrights-d year.
gcc/go/
	* gccgo.texi: Bump @copyrights-go year.
gcc/ada/
  	* gnat_ugn.texi: Bump @copying's copyright year.
  	* gnat_rm.texi: Likewise.
libitm/
	* libitm.texi: Bump @copying's copyright year.
libgomp/
	* libgomp.texi: Bump @copying's copyright year.
libquadmath/
	* libquadmath.texi: Bump @copying's copyright year.

From-SVN: r279811
2020-01-01 12:14:37 +01:00
Ayush Mittal
2c8297996a re PR libgomp/93065 (libgomp: destructor missing to delete goacc_cleanup_key)
PR libgomp/93065
	* oacc-init.c (goacc_runtime_deinitialize): New function.

From-SVN: r279803
2019-12-31 11:34:34 +01:00
Jakub Jelinek
e9dcb75e40 re PR bootstrap/93074 (build FAIL with --enable-offload-targets=nvptx-none)
PR bootstrap/93074
	* plugin/cuda/cuda.h (cuDeviceGetName, cuDriverGetVersion): Declare.
	(cuDeviceTotalMem, cuMemGetInfo): Likewise.  Define to *_v2.

From-SVN: r279747
2019-12-28 10:26:03 +01:00
Maciej W. Rozycki
6c84c8bf9b Add OpenACC 2.6 `acc_get_property' support
Add generic support for the OpenACC 2.6 `acc_get_property' and
`acc_get_property_string' routines, as well as full handlers for the
host and the NVPTX offload targets and minimal handlers for the HSA,
Intel MIC, and AMD GCN offload targets.

Included are C/C++ and Fortran tests that, in particular, print
the property values for acc_property_vendor, acc_property_memory,
acc_property_free_memory, acc_property_name, and acc_property_driver.
The output looks as follows:

Vendor: GNU
Name: GOMP
Total memory: 0
Free memory: 0
Driver: 1.0

with the host driver (where the memory related properties are not
supported for the host device and yield 0, conforming to the standard)
and output like:

Vendor: Nvidia
Total memory: 12651462656
Free memory: 12202737664
Name: TITAN V
Driver: CUDA Driver 9.1

with the NVPTX driver.

2019-12-22  Maciej W. Rozycki  <macro@codesourcery.com>
	    Frederik Harwath  <frederik@codesourcery.com>
	    Thomas Schwinge  <tschwinge@codesourcery.com>

	include/
	* gomp-constants.h (gomp_device_property): New enum.

	libgomp/
	* libgomp.h (gomp_device_descr): Add `get_property_func' member.
	* libgomp-plugin.h (gomp_device_property_value): New union.
	(gomp_device_property_value): New prototype.
	* openacc.h (acc_device_t): Add `acc_device_current' enumeration
	constant.
	(acc_device_property_t): New enum.
	(acc_get_property, acc_get_property_string): New prototypes.
	* oacc-init.c (acc_get_device_type): Also assert that result
	is not `acc_device_current'.
	(get_property_any, acc_get_property, acc_get_property_string):
	New functions.
	* openacc.f90 (openacc_kinds): Add `acc_device_current' and
	`acc_property_memory', `acc_property_free_memory',
	`acc_property_name', `acc_property_vendor' and
	`acc_property_driver' constants.  Add `acc_device_property' data
	type.
	(openacc_internal): Add `acc_get_property' and
	`acc_get_property_string' interfaces.  Add `acc_get_property_h',
	`acc_get_property_string_h', `acc_get_property_l' and
	`acc_get_property_string_l'.
	* oacc-host.c (host_get_property): New function.
	(host_dispatch): Wire it.
	* target.c (gomp_load_plugin_for_device): Handle `get_property'.
	* libgomp.map (OACC_2.6): Add `acc_get_property', `acc_get_property_h_',
	`acc_get_property_string' and `acc_get_property_string_h_' symbols.
	* libgomp.texi (OpenACC Runtime Library Routines): Add
	`acc_get_property'.
	(acc_get_property): New node.
	* plugin/plugin-gcn.c (GOMP_OFFLOAD_get_property): New
	function (stub).
	* plugin/plugin-hsa.c (GOMP_OFFLOAD_get_property): New function.
	* plugin/plugin-nvptx.c (CUDA_CALLS): Add `cuDeviceGetName',
	`cuDeviceTotalMem', `cuDriverGetVersion' and `cuMemGetInfo'
	calls.
	(GOMP_OFFLOAD_get_property): New function.
	(struct ptx_device): Add new field "name".
	(cuda_driver_version_s): Add new static variable ...
	(nvptx_init): ... and init from here.

	* testsuite/libgomp.oacc-c-c++-common/acc_get_property.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property-3.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c: New file
	with test helper functions.

	* testsuite/libgomp.oacc-fortran/acc_get_property.f90: New test.

	liboffloadmic/
	* plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_property):
	New function.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>


Co-Authored-By: Frederik Harwath <frederik@codesourcery.com>
Co-Authored-By: Thomas Schwinge <tschwinge@codesourcery.com>

From-SVN: r279710
2019-12-22 19:54:09 +00:00
Maciej W. Rozycki
c8e759b421 libgomp/test: Fix compilation for build sysroot
Fix a problem with the libgomp testsuite using a method to determine
the compiler to use resulting in the tool being different from one the
library has been built with, and causing a catastrophic failure from the
lack of a suitable `--sysroot=' option where the `--with-build-sysroot='
configuration option has been used to build the compiler resulting in
the inability to link executables.

Address this problem by defining the compiler to use, via the
GCC_UNDER_TEST TCL variable, set in the DejaGNU configuration file from
$CC by autoconf, which will have all the required options set for the
target compiler to build executables in the environment configured,
removing failures like:

.../bin/riscv64-linux-gnu-ld: cannot find crt1.o: No such file or directory
.../bin/riscv64-linux-gnu-ld: cannot find -lm
.../bin/riscv64-linux-gnu-ld: cannot find -lpthread
.../bin/riscv64-linux-gnu-ld: cannot find -lc
.../bin/riscv64-linux-gnu-ld: cannot find -latomic
collect2: error: ld returned 1 exit status
compiler exited with status 1
FAIL: libgomp.c/../libgomp.c-c++-common/atomic-18.c (test for excess errors)
Excess errors:
.../bin/riscv64-linux-gnu-ld: cannot find crt1.o: No such file or directory
.../bin/riscv64-linux-gnu-ld: cannot find -lm
.../bin/riscv64-linux-gnu-ld: cannot find -lpthread
.../bin/riscv64-linux-gnu-ld: cannot find -lc
.../bin/riscv64-linux-gnu-ld: cannot find -latomic

UNRESOLVED: libgomp.c/../libgomp.c-c++-common/atomic-18.c compilation failed to produce executable

and bringing overall test results for the `riscv64-linux-gnu' target
(here with the `x86_64-linux-gnu' host and RISC-V QEMU in the Linux user
emulation mode as the target board) from:

		=== libgomp Summary ===

# of expected passes		44
# of unexpected failures	3274
# of unresolved testcases	3241
# of unsupported tests		548

to:

		=== libgomp Summary ===

# of expected passes		6834
# of unexpected failures	4
# of expected failures		4
# of unsupported tests		518

	libgomp/
	* testsuite/libgomp-test-support.exp.in (GCC_UNDER_TEST): New
	variable.

From-SVN: r279708
2019-12-22 00:41:36 +00:00
Thomas Schwinge
b092fb3b67 [OMP] Restore 'omp declare target link' handling
PASS: libgomp.c/target-link-1.c (test for excess errors)
    [-PASS:-]{+FAIL:+} libgomp.c/target-link-1.c execution test

We need to revert one line of code change from r279625.

	libgomp/
	* target.c (gomp_map_vars_internal): Restore 'omp declare target
	link' handling.

From-SVN: r279701
2019-12-21 23:58:43 +01:00
Julian Brown
9be3ac5d63 Fortran polymorphic class-type support for OpenACC
gcc/fortran/
	* openmp.c (resolve_oacc_data_clauses): Don't disallow allocatable
	polymorphic types for OpenACC.
	* trans-openmp.c (gfc_trans_omp_clauses): Support polymorphic class
	types.

	libgomp/
	* testsuite/libgomp.oacc-fortran/class-ptr-param.f95: New test.
	* testsuite/libgomp.oacc-fortran/classtypes-1.f95: New test.
	* testsuite/libgomp.oacc-fortran/classtypes-2.f95: New test.

From-SVN: r279631
2019-12-20 01:39:49 +00:00
Julian Brown
02817027ca OpenACC 2.6 deep copy: Fortran execution tests
libgomp/
	* testsuite/libgomp.oacc-fortran/deep-copy-1.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-2.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-3.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-4.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-5.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-7.f90: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-8.f90: New test.
	* testsuite/libgomp.oacc-fortran/derived-type-1.f90: New test.
	* testsuite/libgomp.oacc-fortran/derivedtype-1.f95: New test.
	* testsuite/libgomp.oacc-fortran/derivedtype-2.f95: New test.
	* testsuite/libgomp.oacc-fortran/multidim-slice.f95: New test.
	* testsuite/libgomp.oacc-fortran/update-2.f90: New test.

Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>

From-SVN: r279630
2019-12-20 01:39:46 +00:00
Julian Brown
c2eb021fd2 OpenACC 2.6 deep copy: C and C++ execution tests
libgomp/
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c: New test.
	* testsuite/libgomp.oacc-c++/deep-copy-12.C: New test.
	* testsuite/libgomp.oacc-c++/deep-copy-13.C: New test.

From-SVN: r279629
2019-12-20 01:39:42 +00:00
Julian Brown
8e7e71ff24 OpenACC 2.6 deep copy: libgomp parts
include/
	* gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_4, GOMP_MAP_DEEP_COPY):
	Define.
	(gomp_map_kind): Add GOMP_MAP_ATTACH, GOMP_MAP_DETACH,
	GOMP_MAP_FORCE_DETACH.

	libgomp/
	* libgomp.h (struct target_var_desc): Add do_detach flag.
	* oacc-init.c (acc_shutdown_1): Free aux block if present.
	* oacc-mem.c (find_group_last): Add SIZES parameter. Support
	struct components.  Tidy up and add some new checks.
	(goacc_enter_data_internal): Update call to find_group_last.
	(goacc_exit_data_internal): Support detach operations and
	GOMP_MAP_STRUCT.
	(GOACC_enter_exit_data): Handle initial GOMP_MAP_STRUCT or
	GOMP_MAP_FORCE_PRESENT in finalization detection code.  Handle
	attach/detach in enter/exit data detection code.
	* target.c (gomp_map_vars_existing): Initialise do_detach field of
	tgt_var_desc.
	(gomp_map_vars_internal): Support attach.
	(gomp_unmap_vars_internal): Support detach.

From-SVN: r279625
2019-12-20 01:20:30 +00:00
Julian Brown
5d5be7bfb5 OpenACC 2.6 deep copy: attach/detach API routines
libgomp/
	* libgomp.h (struct splay_tree_aux): Add attach_count field.
	(gomp_attach_pointer, gomp_detach_pointer): Add prototypes.
	* libgomp.map (OACC_2.6): New section. Add acc_attach,
	acc_attach_async, acc_detach, acc_detach_async, acc_detach_finalize,
	acc_detach_finalize_async.
	* oacc-mem.c (acc_attach_async, acc_attach, goacc_detach_internal,
	acc_detach, acc_detach_async, acc_detach_finalize,
	acc_detach_finalize_async): New functions.
	* openacc.h (acc_attach, acc_attach_async, acc_detach,
	(acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): Add
	prototypes.
	* target.c (gomp_attach_pointer, gomp_detach_pointer): New functions.
	(gomp_remove_var_internal): Free attachment counts if present.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c: New test.

Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>

From-SVN: r279624
2019-12-20 01:20:27 +00:00
Julian Brown
5bcd470bf0 Use gomp_map_val for OpenACC host-to-device address translation
libgomp/
	* libgomp.h (gomp_map_val): Add prototype.
	* oacc-parallel.c (GOACC_parallel_keyed): Use gomp_map_val instead of
	open-coding device-address calculation.
	* target.c (gomp_map_val): Make global. Use OFFSET_POINTER in
	non-present case.

Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>

From-SVN: r279622
2019-12-20 01:20:19 +00:00
Julian Brown
378da98fcc OpenACC reference count overhaul
libgomp/
	* libgomp.h (struct splay_tree_key_s): Substitute dynamic_refcount
	field for virtual_refcount.
	(enum gomp_map_vars_kind): Add GOMP_MAP_VARS_OPENACC_ENTER_DATA.
	(gomp_free_memmap): Remove prototype.
	* oacc-init.c (acc_shutdown_1): Iteratively call gomp_remove_var
	instead of calling gomp_free_memmap.
	* oacc-mem.c (acc_map_data): Use virtual_refcount instead of
	dynamic_refcount.
	(acc_unmap_data): Open code instead of forcing target_mem_desc's
	to_free field to NULL then calling gomp_unmap_vars.  Handle
	REFCOUNT_INFINITY on target blocks.
	(goacc_enter_data): Rename to...
	(goacc_enter_datum): ...this.  Remove MAPNUM parameter and special
	handling for mapping groups.  Use virtual_refcount instead of
	dynamic_refcount.  Use GOMP_MAP_VARS_OPENACC_ENTER_DATA for
	map_map_vars_async call.  Re-do lookup for target pointer return value.
	(acc_create, acc_create_async, acc_copyin, acc_copyin_async): Call
	renamed goacc_enter_datum function.
	(goacc_exit_data): Rename to...
	(goacc_exit_datum): ...this.  Update for virtual_refcount semantics.
	(acc_delete, acc_delete_async, acc_delete_finalize,
	acc_delete_finalize_async, acc_copyout, acc_copyout_async,
	acc_copyout_finalize, acc_copyout_finalize_async): Call renamed
	goacc_exit_datum function.
	(gomp_acc_remove_pointer, find_pointer): Remove functions.
	(find_group_last, goacc_enter_data_internal, goacc_exit_data_internal):
	New functions.
	(GOACC_enter_exit_data): Use goacc_enter_data_internal and
	goacc_exit_data_internal helper functions.
	* target.c (gomp_map_vars_internal): Handle
	GOMP_MAP_VARS_OPENACC_ENTER_DATA.  Update for virtual_refcount
	semantics.
	(gomp_unmap_vars_internal): Update for virtual_refcount semantics.
	(gomp_load_image_to_device, omp_target_associate_ptr): Zero-initialise
	virtual_refcount field instead of dynamic_refcount.
	(gomp_free_memmap): Remove function.
	* testsuite/libgomp.oacc-c-c++-common/unmap-infinity-1.c: New test.
	* testsuite/libgomp.c-c++-common/unmap-infinity-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Add XFAIL.

From-SVN: r279621
2019-12-20 01:20:16 +00:00
Julian Brown
2a656a9359 Use aux struct in libgomp for infrequently-used/API-specific data
libgomp/
	* libgomp.h (struct splay_tree_aux): New.
	(struct splay_tree_key_s): Replace link_key field with aux pointer.
	* target.c (gomp_map_vars_internal): Adjust for link_key being moved
	to aux struct.
	(gomp_remove_var_internal): Free aux block if present.
	(gomp_load_image_to_device): Zero-initialise aux field instead of
	link_key field.
	(omp_target_associate_pointer): Zero-initialise aux field.

Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>

From-SVN: r279620
2019-12-20 01:20:13 +00:00
Jakub Jelinek
601399c0df re PR middle-end/86416 ([OpenMP] Offloading - better lto1 error message if mode not supported on offloading target)
PR middle-end/86416
	* testsuite/libgomp.c/pr86416-1.c (main): Use L suffixes rather than
	q or none.
	* testsuite/libgomp.c/pr86416-2.c (main): Use Q suffixes rather than
	L or none.

From-SVN: r279552
2019-12-19 00:27:28 +01:00
Julian Brown
a6163563f2 Add OpenACC 2.6's no_create
The clause makes any device code use the local memory address for each
of the variables specified unless the given variable is already present
on the current device.

2019-12-19  Julian Brown  <julian@codesourcery.com>
            Maciej W. Rozycki  <macro@codesourcery.com>
            Tobias Burnus  <tobias@codesourcery.com>
            Thomas Schwinge  <thomas@codesourcery.com>

        gcc/
        * omp-low.c (lower_omp_target): Support GOMP_MAP_NO_ALLOC.
        * tree-pretty-print.c (dump_omp_clause): Likewise.

        gcc/c-family/
        * c-pragma.h (pragma_omp_clause): Add
        PRAGMA_OACC_CLAUSE_NO_CREATE.

        gcc/c/
        * c-parser.c (c_parser_omp_clause_name): Support no_create.
        (c_parser_oacc_data_clause): Likewise.
        (c_parser_oacc_all_clauses): Likewise.
        (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
        (OACC_PARALLEL_CLAUSE_MASK, OACC_SERIAL_CLAUSE_MASK): Add
        PRAGMA_OACC_CLAUSE_NO_CREATE.
        * c-typeck.c (handle_omp_array_sections): Support
        GOMP_MAP_NO_ALLOC.

        gcc/cp/
        * parser.c (cp_parser_omp_clause_name): Support no_create.
        (cp_parser_oacc_data_clause): Likewise.
        (cp_parser_oacc_all_clauses): Likewise.
        (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
        (OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_NO_CREATE.
        * semantics.c (handle_omp_array_sections): Support no_create.

        gcc/fortran/
        * gfortran.h (gfc_omp_map_op): Add OMP_MAP_NO_ALLOC.
        * openmp.c (omp_mask2): Add OMP_CLAUSE_NO_CREATE.
        (gfc_match_omp_clauses): Support no_create.
        (OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES)
        (OACC_DATA_CLAUSES): Add OMP_CLAUSE_NO_CREATE.
        * trans-openmp.c (gfc_trans_omp_clauses_1): Support
        OMP_MAP_NO_ALLOC.

        gcc/testsuite/
        * gfortran.dg/goacc/common-block-1.f90: Add no_create-clause tests.
        * gfortran.dg/goacc/common-block-1.f90: Likewise.
        * gfortran.dg/goacc/data-clauses.f95: Likewise.
        * gfortran.dg/goacc/data-tree.f95: Likewise.
        * gfortran.dg/goacc/kernels-tree.f95: Likewise.
        * gfortran.dg/goacc/parallel-tree.f95: Likewise.

        include/
        * gomp-constants.h (gomp_map_kind): Support GOMP_MAP_NO_ALLOC.

        libgomp/
        * target.c (gomp_map_vars_async): Support GOMP_MAP_NO_ALLOC.
        * testsuite/libgomp.oacc-c-c++-common/no_create-1.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/no_create-2.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/no_create-3.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/no_create-4.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/no_create-5.c: New test.
        * testsuite/libgomp.oacc-fortran/no_create-1.f90: New test.
        * testsuite/libgomp.oacc-fortran/no_create-2.f90: New test.
        * testsuite/libgomp.oacc-fortran/no_create-3.F90: New test.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

Co-Authored-By: Maciej W. Rozycki <macro@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>

From-SVN: r279551
2019-12-19 00:10:08 +01:00
Thomas Schwinge
77ce555548 [OpenACC] Refactor 'goacc_enter_data' so that it can be called from 'goacc_insert_pointer', "present" case, and simplify
No functional changes.

	libgomp/
	* oacc-mem.c (goacc_enter_data): Refactor, so that it can be
	called...
	(goacc_insert_pointer): ... from here, "present" case.
	(goacc_insert_pointer): Inline function into...
	(GOACC_enter_exit_data): ... here, and simplify.

From-SVN: r279540
2019-12-18 18:02:37 +01:00
Thomas Schwinge
9444a299fa [OpenACC] Refactor 'goacc_enter_data' so that it can be called from 'goacc_insert_pointer', "not present" case
No functional changes.

	libgomp/
	* oacc-mem.c (goacc_enter_data): Refactor, so that it can be
	called...
	(goacc_insert_pointer): ... from here, "not present" case.

From-SVN: r279539
2019-12-18 18:02:27 +01:00
Thomas Schwinge
5031b6199f [OpenACC] Refactor 'goacc_remove_pointer' interface
No functional changes.

	libgomp/
	* oacc-mem.c (goacc_remove_pointer): Refactor interface.  Adjust
	all users.

From-SVN: r279538
2019-12-18 18:02:18 +01:00
Thomas Schwinge
1a79b5788d [OpenACC] Refactor 'GOACC_enter_exit_data' to call 'goacc_enter_data', 'goacc_exit_data'
No functional changes.

	libgomp/
	* oacc-mem.c (GOACC_enter_exit_data): Refactor code to call
	'goacc_enter_data', 'goacc_exit_data'.

From-SVN: r279537
2019-12-18 18:02:10 +01:00
Thomas Schwinge
34cfe31e1d [OpenACC] Refactor 'delete_copyout' into 'goacc_exit_data'
Change 'FLAG_COPYOUT', 'FLAG_FINALIZE' into the usual map kind.

No functional changes.

	libgomp/
	* oacc-mem.c (delete_copyout): Refactor into...
	(goacc_exit_data): ... this.  Adjust all users.

From-SVN: r279536
2019-12-18 18:02:00 +01:00
Thomas Schwinge
aaf0e9d708 [OpenACC] Refactor 'present_create_copy' into 'goacc_enter_data'
Every caller passes in 'FLAG_PRESENT', 'FLAG_CREATE'.  Change the remaining
'FLAG_COPY' into the usual map kind.

No functional changes.

	libgomp/
	* oacc-mem.c (present_create_copy): Refactor into...
	(goacc_enter_data): ... this.  Adjust all users.

From-SVN: r279535
2019-12-18 18:01:51 +01:00
Thomas Schwinge
83d1d065df Assert in 'libgomp/target.c:gomp_unmap_vars_internal' that we're not unmapping 'tgt' while it's still in use
libgomp/
	* target.c (gomp_unmap_vars_internal): Add a safeguard to
	'gomp_remove_var'.

From-SVN: r279534
2019-12-18 18:01:33 +01:00
Thomas Schwinge
cc3f11f5f5 [OpenACC] In 'libgomp/target.c:gomp_to_device_kind_p', handle 'GOMP_MAP_FORCE_FROM' like 'GOMP_MAP_FROM'
Fix oversight from r254194 "Coalesce host to device transfers in libgomp".

	libgomp/
	* target.c (gomp_to_device_kind_p): Handle 'GOMP_MAP_FORCE_FROM'
	like 'GOMP_MAP_FROM'.

From-SVN: r279533
2019-12-18 18:01:22 +01:00
Thomas Schwinge
ddb25eb9ca [PR92726, PR92970, PR92984] [OpenACC] Clarify 'acc_delete' etc. for 'NULL'-in, non-present data, or size zero
PR92970 "OpenACC 2.5: 'acc_delete' etc. on non-present data is a no-op" is an
actual bug fix, and the other ones are fall-out, currently undefined behavior.

	libgomp/
	PR libgomp/92726
	PR libgomp/92970
	PR libgomp/92984
	* oacc-mem.c (delete_copyout): No-op behavior if 'lookup_host'
	fails.
	(GOACC_enter_exit_data): Simplify accordingly.
	* testsuite/libgomp.oacc-c-c++-common/pr92970-1.c: New file,
	subsuming...
	* testsuite/libgomp.oacc-c-c++-common/lib-17.c: ... this file...
	* testsuite/libgomp.oacc-c-c++-common/lib-18.c: ..., and this
	file.
	* testsuite/libgomp.oacc-c-c++-common/pr92984-1.c: New file,
	subsuming...
	* testsuite/libgomp.oacc-c-c++-common/lib-21.c: ... this file...
	* testsuite/libgomp.oacc-c-c++-common/lib-29.c: ..., and this
	file.
	* testsuite/libgomp.oacc-c-c++-common/pr92726-1.c: New file,
	subsuming...
	* testsuite/libgomp.oacc-c-c++-common/lib-28.c: ... this file.

From-SVN: r279532
2019-12-18 18:01:11 +01:00
Thomas Schwinge
32128577ae [OpenACC] Elaborate/simplify 'exit data' 'finalize' handling
No functional changes.

	gcc/
	* gimplify.c (gimplify_omp_target_update): Elaborate 'exit data'
	'finalize' handling.
	gcc/testsuite/
	* c-c++-common/goacc/finalize-1.c: Extend.
	* gfortran.dg/goacc/finalize-1.f: Likewise.
	libgomp/
	* oacc-mem.c (GOACC_enter_exit_data): Simplify 'exit data'
	'finalize' handling.

From-SVN: r279531
2019-12-18 18:00:51 +01:00
Thomas Schwinge
ba40277f6a [PR92848] [OpenACC] Use 'GOMP_MAP_VARS_ENTER_DATA' for dynamic data lifetimes
libgomp/
	PR libgomp/92848
	* oacc-mem.c (acc_map_data, present_create_copy)
	(goacc_insert_pointer): Use 'GOMP_MAP_VARS_ENTER_DATA'.
	(acc_unmap_data, delete_copyout, goacc_remove_pointer): Adjust.
	* testsuite/libgomp.oacc-c-c++-common/lib-50.c: Remove.
	* testsuite/libgomp.oacc-c-c++-common/pr92848-1-d-a.c: New file
	* testsuite/libgomp.oacc-c-c++-common/pr92848-1-d-p.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/pr92848-1-r-a.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/pr92848-1-r-p.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c:
	Remove "XFAIL"s.

From-SVN: r279530
2019-12-18 18:00:39 +01:00
Thomas Schwinge
6278b54922 Make 'libgomp/target.c:gomp_unmap_tgt' 'static' again
This got changed to 'attribute_hidden' in r271128, but it's not actually used
outside of 'libgomp/target.c'.

	libgomp/
	* target.c (gomp_unmap_tgt): Make it 'static'.
	* libgomp.h (gomp_unmap_tgt): Remove.

From-SVN: r279529
2019-12-18 18:00:28 +01:00
Tobias Burnus
c80c9e26de PR 86416 – improve lto1 diagnostic if a mode does not exist
PR middle-end/86416
        *  Makefile.in (CFLAGS-lto-streamer-in.o): Pass target_noncanonical on.
        * lto-streamer-in.c (lto_input_mode_table): Improve unsupported-mode
        diagnostic.

        PR middle-end/86416
        * testsuite/libgomp.c/pr86416-1.c: New.
        * testsuite/libgomp.c/pr86416-2.c: New.

From-SVN: r279528
2019-12-18 17:51:08 +01:00
Tobias Burnus
0b8499a480 libgomp/openacc.f90 – clean-up public/private attributes
* config/accel/openacc.f90 (module openacc_kinds): Use 'PUBLIC' to mark
        all symbols as public except for the 'use …, only' imported symbol,
        which is private. 
        (module openacc): Default to 'PRIVATE' to exclude openacc_internal; mark
        all symbols from module openacc_kinds as PUBLIC
        * openacc.f90: Add comment with crossref to that file and openmp_lib.h;
        fix comment typo.
        * openacc_lib.h (acc_device_gcn): Add this PARAMETER.

From-SVN: r279456
2019-12-17 12:19:32 +01:00
Julian Brown
1cbd94e834 Fix potential race condition in OpenACC "exit data" operations
PR libgomp/92881

	libgomp/
	* libgomp.h (gomp_remove_var_async): Add prototype.
	* oacc-mem.c (delete_copyout): Call gomp_remove_var_async instead of
	gomp_remove_var.
	* target.c (gomp_unref_tgt): Change return type to bool, indicating
	whether target_mem_desc was unmapped.
	(gomp_unref_tgt_void): New.
	(gomp_remove_var): Reimplement in terms of...
	(gomp_remove_var_internal): ...this new helper function.
	(gomp_remove_var_async): New, implemented using above helper function.
	(gomp_unmap_vars_internal): Use gomp_unref_tgt_void instead of
	gomp_unref_tgt.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

From-SVN: r279388
2019-12-13 23:14:15 +00:00
Andrew Stubbs
26b74ed022 Update OpenACC tests for amdgcn
2019-12-13  Andrew Stubbs  <ams@codesourcery.com>

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c: Handle gcn.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Disable on GCN.
	* testsuite/libgomp.oacc-c-c++-common/tile-1.c: Likewise.

From-SVN: r279378
2019-12-13 17:40:06 +00:00
Tobias Burnus
7858bb9693 libgomp/openacc.f90 – clean-up public/private attributes
libgomp/
        * openacc.f90 (module openacc_kinds): Use 'PUBLIC' to mark all symbols
        as public except for the 'use …, only' imported symbol, which is
        private.
        (module openacc): Default to 'PRIVATE' to exclude openacc_internal; mark
        all symbols from module openacc_kinds as PUBLIC; add missing PUBLIC
        attributes for acc_copyout_finalize and acc_delete_finalize.

From-SVN: r279337
2019-12-13 08:58:53 +01:00
Jakub Jelinek
f8ac7d8f35 re PR fortran/92899 ([OpenMP] ICE in gfc_trans_omp_atomic, at fortran/trans-openmp.c:3769)
PR fortran/92899
	* trans-openmp.c (gfc_trans_omp_atomic): For GFC_OMP_ATOMIC_SWAP,
	do look through conversion on expr2 if any.

	* testsuite/libgomp.fortran/atomic1.f90: New test.

From-SVN: r279266
2019-12-12 00:58:04 +01:00
Thomas Schwinge
d6e8c01cff [PR92843] [OpenACC] Fix dynamic reference counting for structured 'REFCOUNT_INFINITY'
libgomp/
	PR libgomp/92843
	* oacc-mem.c (present_create_copy, delete_copyout): Fix dynamic
	reference counting for structured 'REFCOUNT_INFINITY'.  Add some
	assertions.
	(goacc_insert_pointer, goacc_remove_pointer): Adjust accordingly.
	* testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: New file.
	* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Fix OpenACC.
	* testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.

From-SVN: r279234
2019-12-11 17:49:27 +01:00
Thomas Schwinge
57963e3934 [OpenACC] Consolidate 'GOACC_enter_exit_data' and its helper functions in 'libgomp/oacc-mem.c'
libgomp/
	* oacc-parallel.c (find_pointer, GOACC_enter_exit_data): Move...
	* oacc-mem.c: ... here.
	(gomp_acc_insert_pointer, gomp_acc_remove_pointer): Rename to
	'goacc_insert_pointer', 'goacc_remove_pointer', and make 'static'.
	* libgomp.h (gomp_acc_insert_pointer, gomp_acc_remove_pointer):
	Remove.
	* libgomp_g.h: Update.

From-SVN: r279233
2019-12-11 17:49:17 +01:00
Thomas Schwinge
c5578b56b6 [OpenACC] Consolidate 'async'/'wait' code in 'libgomp/oacc-async.c'
libgomp/
	* oacc-parallel.c (GOACC_wait, goacc_wait): Move...
	* oacc-async.c: ... here.
	* oacc-int.h (goacc_wait): Declare.
	* libgomp_g.h: Update

From-SVN: r279232
2019-12-11 17:49:08 +01:00
Thomas Schwinge
3d1b5e710e [PR92854] Add 'libgomp.oacc-c-c++-common/acc_map_data-device_already-*.c', 'libgomp.oacc-c-c++-common/acc_map_data-host_already-*.c'
... to document the status quo.

	libgomp/
	PR libgomp/92854
	* testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-1.c:
	New file.
	* testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-1.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-3.c:
	Likewise.

From-SVN: r279231
2019-12-11 17:48:59 +01:00
Thomas Schwinge
5e93943822 [OpenACC] Initialize 'dynamic_refcount' whenever we initialize 'refcount'
Cases missed in r261813 "Update OpenACC data clause semantics to the 2.5
behavior".

	libgomp/
	* target.c (gomp_load_image_to_device, omp_target_associate_ptr):
	Initialize 'dynamic_refcount' whenever we initialize 'refcount'.

Co-Authored-By: Julian Brown <julian@codesourcery.com>

From-SVN: r279230
2019-12-11 17:48:44 +01:00
Tobias Burnus
93d9021987 libgomp – spelling fixes, incl. omp_lib.h.in
* omp_lib.h.in: Fix spelling of function declaration
        omp_get_cancell(l)ation.
        * libgomp.texi (acc_is_present, acc_async_test, acc_async_test_all):
        Fix typos.
        * env.c: Fix comment typos.
        * oacc-host.c: Likewise.
        * ordered.c: Likewise.
        * task.c: Likewise.
        * team.c: Likewise.
        * config/gcn/task.c: Likewise.
        * config/gcn/team.c: Likewise.
        * config/nvptx/task.c: Likewise.
        * config/nvptx/team.c: Likewise.
        * plugin/plugin-gcn.c: Likewise.
        * testsuite/libgomp.fortran/jacobi.f: Likewise.
        * testsuite/libgomp.hsa.c/tiling-2.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c: Likewise.

From-SVN: r279218
2019-12-11 12:45:49 +01:00
Tobias Burnus
a0221aeebd [OpenMP/OpenACC/Fortran] Fix mapping of optional (present|absent) arguments
* testsuite/libgomp.oacc-fortran/optional-cache.f95: Add 'dg-do run'.
        * testsuite/libgomp.oacc-fortran/optional-reduction.f90: Remove
        unnecessary 'dg-additional-options "-w"'.

From-SVN: r279217
2019-12-11 11:40:11 +01:00
Thomas Schwinge
47afc7b4dd [PR92116, PR92877] [OpenACC] Replace 'openacc.data_environ' by standard libgomp mechanics
libgomp/
	PR libgomp/92116
	PR libgomp/92877
	* oacc-mem.c (lookup_dev): Reimplement.  Adjust all users.
	* libgomp.h (struct acc_dispatch_t): Remove 'data_environ' member.
	Adjust all users.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c:
	Remove XFAIL.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/pr92877-1.c: New file.

Co-Authored-By: Julian Brown <julian@codesourcery.com>

From-SVN: r279147
2019-12-09 23:52:56 +01:00
Thomas Schwinge
cec41816c1 [PR92503] [OpenACC] Don't silently 'acc_unmap_data' in 'acc_free'
libgomp/
	PR libgomp/92503
	* oacc-mem.c (acc_free): Error out instead of 'acc_unmap_data'.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-1.c: New
	file.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-3.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust.
	* testsuite/libgomp.oacc-c-c++-common/context-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-13.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-91.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.

From-SVN: r279146
2019-12-09 23:52:47 +01:00
Thomas Schwinge
e103542bc8 [PR92840] [OpenACC] Refuse 'acc_unmap_data' unless mapped by 'acc_map_data'
libgomp/
	PR libgomp/92840
	* oacc-mem.c (acc_map_data): Clarify reference counting behavior.
	(acc_unmap_data): Add error case for 'REFCOUNT_INFINITY'.
	* testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-1.c:
	New file.
	* testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_unmap_data-pr92840-3.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Adjust.
	* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Adjust.

From-SVN: r279145
2019-12-09 23:52:36 +01:00
Thomas Schwinge
41f36f22f3 [PR92511] More testing for OpenACC "present" subarrays
In particular, "subset subarrays".

	libgomp/
	PR libgomp/92511
	* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: Remove
	this file...
	* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c: ..., and
	this file...
	* testsuite/libgomp.oacc-c-c++-common/lib-22.c: ..., and this
	file...
	* testsuite/libgomp.oacc-c-c++-common/lib-30.c: ..., and this
	file...
	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-p.c:
	... with their content moved into, and extended in this new file.
	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-a.c:
	New file.
	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-d-p.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-1-r-a.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/subset-subarray-mappings-2.c:
	Likewise.

From-SVN: r279122
2019-12-09 12:40:36 +01:00
Thomas Schwinge
6effebe10b Add 'libgomp.oacc-c-c++-common/map-data-1.c'
libgomp/
	* testsuite/libgomp.oacc-c-c++-common/map-data-1.c: New file.

From-SVN: r279121
2019-12-09 12:40:27 +01:00
Thomas Schwinge
7d02047320 [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c'
... to document the status quo.

	libgomp/
	PR libgomp/92854
	* testsuite/libgomp.oacc-c-c++-common/pr92854-1.c: New file.

From-SVN: r279120
2019-12-09 12:40:17 +01:00
Thomas Schwinge
1e1fb715cd Add 'libgomp.oacc-c-c++-common/host_data-6.c'
libgomp/
	* testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New file.

From-SVN: r279119
2019-12-09 12:40:08 +01:00
Thomas Schwinge
b5859e4002 In 'libgomp/target.c:gomp_exit_data', remove open-coded 'gomp_remove_var'
libgomp/
	* target.c (gomp_exit_data): Use 'gomp_remove_var'.

From-SVN: r279118
2019-12-09 12:39:57 +01:00
Tobias Burnus
11586ed9c9 libgomp/testsuite/*fortran – make 'stop' values unique
* testsuite/libgomp.fortran/use_device_addr-3.f90: Make 'stop' codes
        unique.
        * testsuite/libgomp.fortran/use_device_addr-4.f90: Ditto.
        * testsuite/libgomp.fortran/use_device_ptr-optional-2.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/declare-5.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/optional-data-copyin-by-value.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/optional-firstprivate.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/optional-update-host.f90: Ditto.

From-SVN: r279117
2019-12-09 12:21:22 +01:00
Kwok Cheung Yeung
3ad4a03ac6 Fix Fortran linker errors on AMD GCN
2019-12-06  Kwok Cheung Yeung  <kcy@codesourcery.com>

	libgomp/
	* config/accel/proc.c (omp_get_num_procs): Apply ialias macro.

From-SVN: r279060
2019-12-06 18:33:19 +00:00
Tobias Burnus
6e4d01d61f [OpenMP/OpenACC/Fortran] Fix mapping of optional (present|absent) arguments
2019-12-06  Tobias Burnus  <tobias@codesourcery.com>
            Kwok Cheung Yeung <kcy@codesourcery.com>

        gcc/fortran/
        * trans-openmp.c (gfc_build_conditional_assign,
        gfc_build_conditional_assign_expr): New static functions.
        (gfc_omp_finish_clause, gfc_trans_omp_clauses): Handle mapping of
        absent optional arguments and fix mapping of present optional args.

        gcc/
        * omp-low.c (lower_omp_target): For optional arguments, deref once
        more to obtain the type.

        libgomp/
        * oacc-mem.c (update_dev_host, gomp_acc_insert_pointer): Just return
        if input it a NULL pointer.
        * testsuite/libgomp.oacc-c-c++-common/lib-43.c: Remove; dependent on
        diagnostic of NULL pointer.
        * testsuite/libgomp.oacc-c-c++-common/lib-47.c: Ditto.
        * testsuite/libgomp.fortran/optional-map.f90: New.
        * testsuite/libgomp.fortran/use_device_addr-1.f90
        (test_dummy_opt_callee_1_absent): New.
        (test_dummy_opt_call_1): Call it.
        * testsuite/libgomp.fortran/use_device_addr-2.f90: Likewise.
        * testsuite/libgomp.fortran/use_device_addr-3.f90: Likewise.
        * testsuite/libgomp.fortran/use_device_addr-4.f90: Likewise.
        * testsuite/libgomp.oacc-fortran/optional-cache.f95: New.
        * testsuite/libgomp.oacc-fortran/optional-data-copyin-by-value.f90: New.
        * testsuite/libgomp.oacc-fortran/optional-data-copyin.f90: New.
        * testsuite/libgomp.oacc-fortran/optional-data-copyout.f90: New.
        * testsuite/libgomp.oacc-fortran/optional-data-enter-exit.f90: New.
        * testsuite/libgomp.oacc-fortran/optional-declare.f90: New.
        * testsuite/libgomp.oacc-fortran/optional-firstprivate.f90: New.
        * testsuite/libgomp.oacc-fortran/optional-host_data.f90: New.
        * testsuite/libgomp.oacc-fortran/optional-nested-calls.f90: New.
        * testsuite/libgomp.oacc-fortran/optional-private.f90: New.
        * testsuite/libgomp.oacc-fortran/optional-reduction.f90: New.
        * testsuite/libgomp.oacc-fortran/optional-update-device.f90: New.
        * testsuite/libgomp.oacc-fortran/optional-update-host.f90: New.


Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>

From-SVN: r279043
2019-12-06 14:06:53 +01:00
Tobias Burnus
c3cb71ef35 OpenMP] Fix use_device_… with absent optional arg
gcc/fortran/
        * trans-openmp.c (gfc_omp_is_optional_argument,
        gfc_omp_check_optional_argument): Handle type(c_ptr),value which uses a
        hidden argument for the is-present check.

        gcc/
        * omp-low.c (lower_omp_target): For use_device_ptr/use_derice_addr
        and Fortran's optional arguments, unconditionally add the is-present
        condition before the libgomp call.

        libgomp/
        * testsuite/libgomp.fortran/use_device_ptr-optional-2.f90: Add
        'type(c_ptr), value' test case. Conditionally map the per-value
        passed arguments.

From-SVN: r279004
2019-12-05 16:18:39 +01:00
Tobias Burnus
ad80ec460b testsuite/libgomp.oacc-fortran: Update dg-output to fix GCN
* testsuite/libgomp.oacc-fortran/error_stop-1.f: Also don't
        expect dg-output of 'Error termination.' for GCN.
        * testsuite/libgomp.oacc-fortran/error_stop-2.f: Likewise.
        * testsuite/libgomp.oacc-fortran/error_stop-3.f: Likewise.

Plus: Fix date of a previous commit in:
        * gcc/fortran/ChangeLog
        * gcc/testsuite/ChangeLog

From-SVN: r278986
2019-12-05 09:38:53 +01:00
Jakub Jelinek
a8a5f4cc04 re PR fortran/92756 (ICE in lower_omp, at omp-low.c:12988)
PR fortran/92756
	* trans-openmp.c (gfc_trans_omp_teams): Wrap OMP_TEAMS body into a
	BIND_EXPR with a forced BLOCK.

	* gfortran.dg/gomp/teams1.f90: New test.

	* testsuite/libgomp.fortran/teams1.f90: New test.
	* testsuite/libgomp.fortran/teams2.f90: New test.

From-SVN: r278956
2019-12-04 09:47:13 +01:00
Frederik Harwath
d8f0024b57 Validate acc_device_t uses
Check that function arguments of type acc_device_t
are valid enumeration values in all publicly visible
functions from oacc-init.c.

2019-12-03  Frederik Harwath  <frederik@codesourcery.com>

	libgomp/
	* oacc-init.c (acc_known_device_type): Add function.
	(unknown_device_type_error): Add function.
 	(name_of_acc_device_t): Change to call unknown_device_type_error
	on unknown type.
	(resolve_device): Use acc_known_device_type.
 	(acc_init): Fail if acc_device_t argument is not valid.
 	(acc_shutdown): Likewise.
 	(acc_get_num_devices): Likewise.
 	(acc_set_device_type): Likewise.
	(acc_get_device_num): Likewise.
	(acc_set_device_num): Likewise.
 	(acc_on_device): Add comment that argument validity is not checked.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

From-SVN: r278937
2019-12-03 14:38:54 +00:00
Andrew Stubbs
83caa34e2a Enable OpenACC GCN testing.
2019-12-03  Andrew Stubbs  <ams@codesourcery.com>

	libgomp/
	* testsuite/lib/libgomp.exp (offload_target_to_openacc_device_type):
	Recognize amdgcn.
	(check_effective_target_openacc_amdgcn_accel_present): New proc.
	(check_effective_target_openacc_amdgcn_accel_selected): New proc.
	* testsuite/libgomp.oacc-c++/c++.exp: Add support for amdgcn.
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.

From-SVN: r278935
2019-12-03 12:53:53 +00:00
Szabolcs Nagy
004843655a musl: Fix invalid tls model in libgomp and libitm PR91938
Musl does not support initial-exec tls in dynamically loaded shared
libraries.

libgomp/ChangeLog:

2019-12-03  Szabolcs Nagy  <szabolcs.nagy@arm.com>

	PR libgomp/91938
	* configure.tgt: Avoid IE tls on *-*-musl*.

libitm/ChangeLog:

2019-12-03  Szabolcs Nagy  <szabolcs.nagy@arm.com>

	PR libgomp/91938
	* configure.tgt: Avoid IE tls on *-*-musl*.

From-SVN: r278932
2019-12-03 11:13:38 +00:00
Tobias Burnus
9909a05940 Fortran] OpenACC – permit common blocks in some clauses
* testsuite/libgomp.oacc-fortran/declare-5.f90: Extend by
        adding a common-block test case.

From-SVN: r278845
2019-11-29 14:59:21 +01:00
Jakub Jelinek
52702016ad re PR c++/60228 (ICE using lambda in #pragma omp declare reduction)
PR c++/60228
	* parser.c (cp_parser_omp_declare_reduction_exprs): If
	processing_template_decl, wrap the combiner or initializer
	into EXPR_STMT.
	* decl.c (start_preparsed_function): Don't start a lambda scope
	for DECL_OMP_DECLARE_REDUCTION_P functions.
	(finish_function): Don't finish a lambda scope for
	DECL_OMP_DECLARE_REDUCTION_P functions, nor cp_fold_function
	them nor cp_genericize them.
	* mangle.c (decl_mangling_context): Look through
	DECL_OMP_DECLARE_REDUCTION_P functions.
	* semantics.c (expand_or_defer_fn_1): For DECL_OMP_DECLARE_REDUCTION_P
	functions, use tentative linkage, don't keep their bodies with
	-fkeep-inline-functions and return false at the end.

	* g++.dg/gomp/openmp-simd-2.C: Don't expect bodies for
	DECL_OMP_DECLARE_REDUCTION_P functions.

	* testsuite/libgomp.c++/udr-20.C: New test.
	* testsuite/libgomp.c++/udr-21.C: New test.

From-SVN: r278832
2019-11-29 10:10:44 +01:00
Jakub Jelinek
b3f44388f3 re PR c++/60228 (ICE using lambda in #pragma omp declare reduction)
PR c++/60228
	* parser.c (cp_parser_omp_declare_reduction_exprs): If
	processing_template_decl, wrap the combiner or initializer
	into EXPR_STMT.
	* decl.c (start_preparsed_function): Don't start a lambda scope
	for DECL_OMP_DECLARE_REDUCTION_P functions.
	(finish_function): Don't finish a lambda scope for
	DECL_OMP_DECLARE_REDUCTION_P functions, nor cp_fold_function
	them nor cp_genericize them.
	* mangle.c (decl_mangling_context): Look through
	DECL_OMP_DECLARE_REDUCTION_P functions.
	* semantics.c (expand_or_defer_fn_1): For DECL_OMP_DECLARE_REDUCTION_P
	functions, use tentative linkage, don't keep their bodies with
	-fkeep-inline-functions and return false at the end.

	* g++.dg/gomp/openmp-simd-2.C: Don't expect bodies for
	DECL_OMP_DECLARE_REDUCTION_P functions.

	* testsuite/libgomp.c++/udr-20.C: New test.
	* testsuite/libgomp.c++/udr-21.C: New test.

From-SVN: r278831
2019-11-29 10:08:46 +01:00
Thomas Schwinge
b1a73b4894 Fix 'libgomp.fortran/target-print-1.f90', 'libgomp.oacc-fortran/print-1.f90' for offload target nvptx
libgomp/
	* testsuite/lib/libgomp.exp
	(check_effective_target_offload_target_nvptx): New proc.
	* testsuite/libgomp.fortran/target-print-1.f90: Use it with
	'dg-skip-if'.
	* testsuite/libgomp.oacc-fortran/print-1.f90: Likewise.
	* testsuite/libgomp.fortran/target-print-1-nvptx.f90: New file.
	* testsuite/libgomp.oacc-fortran/print-1-nvptx.f90: Likewise.

From-SVN: r278779
2019-11-27 18:50:55 +01:00
Rainer Orth
b8e724465b Fix failures on Solaris with -fno-common default
gcc/testsuite:
	* gcc.c-torture/execute/20030913-1.c: Rename glob to g.
	* gcc.c-torture/execute/960218-1.c: Rename glob to gl.
	* gcc.c-torture/execute/complex-6.c: Rename err to e.
	* gcc.dg/torture/ssa-pta-fn-1.c: Rename glob to g.

	libgomp:
	* testsuite/libgomp.c/pr39591-1.c: Rename err to e.
	* testsuite/libgomp.c/pr39591-2.c: Likewise.
	* testsuite/libgomp.c/pr39591-3.c: Likewise.
	* testsuite/libgomp.c/private-1.c: Likewise.
	* testsuite/libgomp.c/task-1.c: Likewise.
	* testsuite/libgomp.c/task-5.c: Renamed err to serr.

From-SVN: r278571
2019-11-21 16:14:21 +00:00
Julian Brown
d88b27daa1 AMD GCN libgomp plugin queue-full condition locking fix
libgomp/
	* plugin/plugin-gcn.c (wait_for_queue_nonfull): Don't lock/unlock
	aq->mutex here.
	(queue_push_launch): Lock aq->mutex before calling
	wait_for_queue_nonfull.
	(queue_push_callback): Likewise.
	(queue_push_asyncwait): Likewise.
	(queue_push_placeholder): Likewise.

Reviewed-by: Andrew Stubbs <ams@codesourcery.com>

From-SVN: r278517
2019-11-20 17:56:30 +00:00
Julian Brown
8d2f4ddfd7 Fix host-to-device copies from rodata for AMD GCN
libgomp/
	* plugin/plugin-gcn.c (hsa_memory_copy_wrapper): New.
	(copy_data, GOMP_OFFLOAD_host2dev): Use above function.
	(GOMP_OFFLOAD_dev2host, GOMP_OFFLOAD_dev2dev): Check hsa_memory_copy
	return code.

Reviewed-by: Andrew Stubbs <ams@codesourcery.com>

From-SVN: r278516
2019-11-20 17:53:31 +00:00
Julian Brown
e307b05f43 OpenACC "present" subarrays: runtime API return value and unmapping fixes
PR libgomp/92511

	libgomp/
	* oacc-mem.c (present_create_copy): Fix device pointer return value in
	case of "present" subarray.  Use tgt->tgt_start instead of tgt->to_free
	in non-present/create case.
	(delete_copyout): Change error condition to fail only on copies outside
	of mapped block.  Adjust error message accordingly.
	* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/copyin-devptr-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/lib-20.c: Adjust expected error
	message.
	* testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-22.c: Allow test to pass now.
	* testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

From-SVN: r278514
2019-11-20 17:51:09 +00:00
Maciej W. Rozycki
83115e9eb8 libgomp/test: Add flags to find libatomic in build-tree testing
Add flags to find libatomic in build-tree testing, fixing a catastrophic
libgomp testsuite failure with `riscv*-*-linux*' targets, which imply
`-latomic' with the `-pthread' GCC option, implied in turn by the
`-fopenacc' and `-fopenmp' options, removing failures like:

.../bin/riscv64-linux-gnu-ld: cannot find -latomic
collect2: error: ld returned 1 exit status
compiler exited with status 1
FAIL: libgomp.c/../libgomp.c-c++-common/atomic-18.c (test for excess errors)
Excess errors:
.../bin/riscv64-linux-gnu-ld: cannot find -latomic

UNRESOLVED: libgomp.c/../libgomp.c-c++-common/atomic-18.c compilation failed to produce executable

and bringing overall test results for the `riscv64-linux-gnu' target
(here with the `x86_64-linux-gnu' host and RISC-V QEMU in the Linux user
emulation mode as the target board) from:

		=== libgomp Summary ===

# of expected passes		90
# of unexpected failures	3267
# of expected failures		2
# of unresolved testcases	3247
# of unsupported tests		548

to:

		=== libgomp Summary ===

# of expected passes		6834
# of unexpected failures	4
# of expected failures		4
# of unsupported tests		518

	libgomp/
	* testsuite/lib/libgomp.exp (libgomp_init): Add flags to find
	libatomic in build-tree testing.

From-SVN: r278505
2019-11-20 15:49:23 +00:00
Maciej W. Rozycki
a128988785 libgomp: Regenerate `testsuite/Makefile.in' for GCC_HEADER_STDINT removal
Commit r276389 ("configure.ac: Remove GCC_HEADER_STDINT(gstdint.h)") has
not regenerated `testsuite/Makefile.in'.  Fix it.

	libgomp/
	* testsuite/Makefile.in: Regenerate.

From-SVN: r278384
2019-11-18 00:33:37 +00:00
Andrew Stubbs
8916ba874d Add tests for print from offload target.
2019-11-15  Andrew Stubbs  <ams@codesourcery.com>

	libgomp/
	* testsuite/libgomp.c/target-print-1.c: New file.
	* testsuite/libgomp.fortran/target-print-1.f90: New file.
	* testsuite/libgomp.oacc-c/print-1.c: New file.
	* testsuite/libgomp.oacc-fortran/print-1.f90: New file.

From-SVN: r278284
2019-11-15 10:49:10 +00:00
Andrew Stubbs
237957cc2c GCN Libgomp Plugin
2019-11-13  Andrew Stubbs  <ams@codesourcery.com>
	    Kwok Cheung Yeung  <kcy@codesourcery.com>
	    Julian Brown  <julian@codesourcery.com>
	    Tom de Vries  <tom@codesourcery.com>

	libgomp/
	* plugin/Makefrag.am: Add amdgcn plugin support.
	* plugin/configfrag.ac: Likewise.
	* plugin/plugin-gcn.c: New file.
	* configure: Regenerate.
	* Makefile.in: Regenerate.
	* testsuite/Makefile.in: Regenerate.

Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
Co-Authored-By: Tom de Vries <tom@codesourcery.com>

From-SVN: r278138
2019-11-13 12:38:18 +00:00
Andrew Stubbs
cee1645106 Optimize GCN OpenMP malloc performance
2019-11-13  Andrew Stubbs  <ams@codesourcery.com>

	libgomp/
	* config/gcn/team.c (gomp_gcn_enter_kernel): Set up the team arena
	and use team_malloc variants.
	(gomp_gcn_exit_kernel): Use team_free.
	* libgomp.h (TEAM_ARENA_SIZE): Define.
	(TEAM_ARENA_START): Define.
	(TEAM_ARENA_FREE): Define.
	(TEAM_ARENA_END): Define.
	(team_malloc): New function.
	(team_malloc_cleared): New function.
	(team_free): New function.
	* team.c (gomp_new_team): Initialize and use team_malloc.
	(free_team): Use team_free.
	(gomp_free_thread): Use team_free.
	(gomp_pause_host): Use team_free.
	* work.c (gomp_init_work_share): Use team_malloc.
	(gomp_fini_work_share): Use team_free.

From-SVN: r278136
2019-11-13 12:38:09 +00:00
Andrew Stubbs
fa4999953d GCN libgomp port
2019-11-13  Andrew Stubbs  <ams@codesourcery.com>
	    Kwok Cheung Yeung  <kcy@codesourcery.com>
	    Julian Brown  <julian@codesourcery.com>
	    Tom de Vries  <tom@codesourcery.com>

	include/
	* gomp-constants.h (GOMP_DEVICE_GCN): Define.
	(GOMP_VERSION_GCN): Define.

	libgomp/
	* Makefile.am (libgomp_la_SOURCES): Add oacc-target.c.
	* Makefile.in: Regenerate.
	* config.h.in (PLUGIN_GCN): Add new undef.
	* config/accel/openacc.f90 (acc_device_gcn): New parameter.
	* config/gcn/affinity-fmt.c: New file.
	* config/gcn/bar.c: New file.
	* config/gcn/bar.h: New file.
	* config/gcn/doacross.h: New file.
	* config/gcn/icv-device.c: New file.
	* config/gcn/oacc-target.c: New file.
	* config/gcn/simple-bar.h: New file.
	* config/gcn/target.c: New file.
	* config/gcn/task.c: New file.
	* config/gcn/team.c: New file.
	* config/gcn/time.c: New file.
	* configure.ac: Add amdgcn*-*-*.
	* configure: Regenerate.
	* configure.tgt: Add amdgcn*-*-*.
	* libgomp-plugin.h (offload_target_type): Add OFFLOAD_TARGET_TYPE_GCN.
	* libgomp.h (gcn_thrs): Add amdgcn variant.
	(set_gcn_thrs): Likewise.
	(gomp_thread): Likewise.
	* oacc-int.h (goacc_thread): Likewise.
	* oacc-target.c: New file.
	* openacc.f90 (acc_device_gcn): New parameter.
	* openacc.h (acc_device_t): Add acc_device_gcn.
	* team.c (gomp_free_pool_helper): Add amdgcn support.

Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
Co-Authored-By: Tom de Vries <tom@codesourcery.com>

From-SVN: r278135
2019-11-13 12:38:04 +00:00
Andrew Stubbs
d2903ce05b Add device number to GOMP_OFFLOAD_openacc_async_construct
2019-11-13  Andrew Stubbs  <ams@codesourcery.com>
	    Julian Brown  <julian@codesourcery.com>

	libgomp/
	* libgomp-plugin.h (GOMP_OFFLOAD_openacc_async_construct): Add int
	parameter.
	* oacc-async.c (lookup_goacc_asyncqueue): Pass device number to the
	queue constructor.
	* oacc-host.c (host_openacc_async_construct): Add device parameter.
	* plugin/plugin-nvptx.c (GOMP_OFFLOAD_openacc_async_construct): Add
	device parameter.

Co-Authored-By: Julian Brown <julian@codesourcery.com>

From-SVN: r278134
2019-11-13 12:37:59 +00:00
Andrew Stubbs
b3d14b3aa3 Move generic libgomp files from nvptx to accel
2019-11-13  Andrew Stubbs  <ams@codesourcery.com>

	libgomp/
	* configure.tgt (nvptx*-*-*): Add "accel" directory.
	* config/nvptx/libgomp-plugin.c: Move ...
	* config/accel/libgomp-plugin.c: ... to here.
	* config/nvptx/lock.c: Move ...
	* config/accel/lock.c: ... to here.
	* config/nvptx/mutex.c: Move ...
	* config/accel/mutex.c: ... to here.
	* config/nvptx/mutex.h: Move ...
	* config/accel/mutex.h: ... to here.
	* config/nvptx/oacc-async.c: Move ...
	* config/accel/oacc-async.c: ... to here.
	* config/nvptx/oacc-cuda.c: Move ...
	* config/accel/oacc-cuda.c: ... to here.
	* config/nvptx/oacc-host.c: Move ...
	* config/accel/oacc-host.c: ... to here.
	* config/nvptx/oacc-init.c: Move ...
	* config/accel/oacc-init.c: ... to here.
	* config/nvptx/oacc-mem.c: Move ...
	* config/accel/oacc-mem.c: ... to here.
	* config/nvptx/oacc-plugin.c: Move ...
	* config/accel/oacc-plugin.c: ... to here.
	* config/nvptx/omp-lock.h: Move ...
	* config/accel/omp-lock.h: ... to here.
	* config/nvptx/openacc.f90: Move ...
	* config/accel/openacc.f90: ... to here.
	* config/nvptx/pool.h: Move ...
	* config/accel/pool.h: ... to here.
	* config/nvptx/proc.c: Move ...
	* config/accel/proc.c: ... to here.
	* config/nvptx/ptrlock.c: Move ...
	* config/accel/ptrlock.c: ... to here.
	* config/nvptx/ptrlock.h: Move ...
	* config/accel/ptrlock.h: ... to here.
	* config/nvptx/sem.c: Move ...
	* config/accel/sem.c: ... to here.
	* config/nvptx/sem.h: Move ...
	* config/accel/sem.h: ... to here.
	* config/nvptx/thread-stacksize.h: Move ...
	* config/accel/thread-stacksize.h: ... to here.

From-SVN: r278132
2019-11-13 12:37:50 +00:00
Maciej W. Rozycki
62aee289e4 Add OpenACC 2.6 `serial' construct support
The `serial' construct (cf. section 2.5.3 of the OpenACC 2.6 standard)
    is equivalent to a `parallel' construct with clauses `num_gangs(1)
    num_workers(1) vector_length(1)' implied.
    These clauses are therefore not supported with the `serial'
    construct. All the remaining clauses accepted with `parallel' are also
    accepted with `serial'.

    The `serial' construct is implemented like `parallel', except for
    hardcoding dimensions rather than taking them from the relevant
    clauses, in `expand_omp_target'.

    Separate codes are used to denote the `serial' construct throughout the
    middle end, even though the mapping of `serial' to an equivalent
    `parallel' construct could have been done in the individual language
    frontends. In particular, this allows to distinguish between compute
    constructs in warnings, error messages, dumps etc.

    2019-11-12  Maciej W. Rozycki  <macro@codesourcery.com>
                Tobias Burnus  <tobias@codesourcery.com>
                Frederik Harwath  <frederik@codesourcery.com>
                Thomas Schwinge  <thomas@codesourcery.com>

    gcc/
    * gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_SERIAL
    enumeration constant.
    (is_gimple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
    (is_gimple_omp_offloaded): Likewise.
    * gimplify.c (omp_region_type): Add ORT_ACC_SERIAL enumeration
    constant.  Adjust the value of ORT_NONE accordingly.
    (is_gimple_stmt): Handle OACC_SERIAL.
    (oacc_default_clause): Handle ORT_ACC_SERIAL.
    (gomp_needs_data_present): Likewise.
    (gimplify_adjust_omp_clauses): Likewise.
    (gimplify_omp_workshare): Handle OACC_SERIAL.
    (gimplify_expr): Likewise.
    * omp-expand.c (expand_omp_target):
    Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
    (build_omp_regions_1, omp_make_gimple_edges): Likewise.
    * omp-low.c (is_oacc_parallel): Rename function to...
    (is_oacc_parallel_or_serial): ... this.
    Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
    (scan_sharing_clauses): Adjust accordingly.
    (scan_omp_for): Likewise.
    (lower_oacc_head_mark): Likewise.
    (convert_from_firstprivate_int): Likewise.
    (lower_omp_target): Likewise.
    (check_omp_nesting_restrictions): Handle
    GF_OMP_TARGET_KIND_OACC_SERIAL.
    (lower_oacc_reductions): Likewise.
    (lower_omp_target): Likewise.
    * tree.def (OACC_SERIAL): New tree code.
    * tree-pretty-print.c (dump_generic_node): Handle OACC_SERIAL.

    * doc/generic.texi (OpenACC): Document OACC_SERIAL.

    gcc/c-family/
    * c-pragma.h (pragma_kind): Add PRAGMA_OACC_SERIAL enumeration
    constant.
    * c-pragma.c (oacc_pragmas): Add "serial" entry.

    gcc/c/
    * c-parser.c (OACC_SERIAL_CLAUSE_MASK): New macro.
    (c_parser_oacc_kernels_parallel): Rename function to...
    (c_parser_oacc_compute): ... this.  Handle PRAGMA_OACC_SERIAL.
    (c_parser_omp_construct): Update accordingly.

    gcc/cp/
    * constexpr.c (potential_constant_expression_1): Handle
    OACC_SERIAL.
    * parser.c (OACC_SERIAL_CLAUSE_MASK): New macro.
    (cp_parser_oacc_kernels_parallel): Rename function to...
    (cp_parser_oacc_compute): ... this.  Handle PRAGMA_OACC_SERIAL.
    (cp_parser_omp_construct): Update accordingly.
    (cp_parser_pragma): Handle PRAGMA_OACC_SERIAL.  Fix alphabetic
    order.
    * pt.c (tsubst_expr): Handle OACC_SERIAL.

    gcc/fortran/
    * gfortran.h (gfc_statement): Add ST_OACC_SERIAL_LOOP,
    ST_OACC_END_SERIAL_LOOP, ST_OACC_SERIAL and ST_OACC_END_SERIAL
    enumeration constants.
    (gfc_exec_op): Add EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL
    enumeration constants.
    * match.h (gfc_match_oacc_serial): New prototype.
    (gfc_match_oacc_serial_loop): Likewise.
    * dump-parse-tree.c (show_omp_node, show_code_node): Handle
    EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL.
    * match.c (match_exit_cycle): Handle EXEC_OACC_SERIAL_LOOP.
    * openmp.c (OACC_SERIAL_CLAUSES): New macro.
    (gfc_match_oacc_serial_loop): New function.
    (gfc_match_oacc_serial): Likewise.
    (oacc_is_loop): Handle EXEC_OACC_SERIAL_LOOP.
    (resolve_omp_clauses): Handle EXEC_OACC_SERIAL.
    (oacc_code_to_statement): Handle EXEC_OACC_SERIAL and
    EXEC_OACC_SERIAL_LOOP.
    (gfc_resolve_oacc_directive): Likewise.
    * parse.c (decode_oacc_directive) <'s'>: Add case for "serial"
    and "serial loop".
    (next_statement): Handle ST_OACC_SERIAL_LOOP and ST_OACC_SERIAL.
    (gfc_ascii_statement): Likewise.  Handle ST_OACC_END_SERIAL_LOOP
    and ST_OACC_END_SERIAL.
    (parse_oacc_structured_block): Handle ST_OACC_SERIAL.
    (parse_oacc_loop): Handle ST_OACC_SERIAL_LOOP and
    ST_OACC_END_SERIAL_LOOP.
    (parse_executable): Handle ST_OACC_SERIAL_LOOP and
    ST_OACC_SERIAL.
    (is_oacc): Handle EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL.
    * resolve.c (gfc_resolve_blocks, gfc_resolve_code): Likewise.
    * st.c (gfc_free_statement): Likewise.
    * trans-openmp.c (gfc_trans_oacc_construct): Handle
    EXEC_OACC_SERIAL.
    (gfc_trans_oacc_combined_directive): Handle
    EXEC_OACC_SERIAL_LOOP.
    (gfc_trans_oacc_directive): Handle EXEC_OACC_SERIAL_LOOP and
    EXEC_OACC_SERIAL.
    * trans.c (trans_code): Likewise.

    gcc/testsuite/
    * c-c++-common/goacc/parallel-dims.c: New test.
    * gfortran.dg/goacc/parallel-dims.f90: New test.

    libgomp/
    * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: New test.
    * testsuite/libgomp.oacc-fortran/parallel-dims-aux.c: New test.
    * testsuite/libgomp.oacc-fortran/parallel-dims.f89: New test.
    * testsuite/libgomp.oacc-fortran/parallel-dims-2.f90: New test.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>


Co-Authored-By: Frederik Harwath <frederik@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>

From-SVN: r278082
2019-11-12 08:45:35 +00:00
Tobias Burnus
a2c26c5031 Fortran] Support absent optional args with use_device_{ptr,addr}
2019-11-11  Tobias Burnus  <tobias@codesourcery.com>
            Kwok Cheung Yeung  <kcy@codesourcery.com>

        gcc/
        * langhooks-def.h (LANG_HOOKS_OMP_CHECK_OPTIONAL_ARGUMENT):
        Renamed from LANG_HOOKS_OMP_IS_OPTIONAL_ARGUMENT; update define.
        (LANG_HOOKS_DECLS): Rename also here.
        * langhooks.h (lang_hooks_for_decls): Rename
        omp_is_optional_argument to omp_check_optional_argument; take
        additional bool argument.
        * omp-general.h (omp_check_optional_argument): Likewise.
        * omp-general.h (omp_check_optional_argument): Likewise.
        * omp-low.c (lower_omp_target): Update calls; handle absent
        Fortran optional arguments with USE_DEVICE_ADDR/USE_DEVICE_PTR.

        gcc/fortran/
        * trans-expr.c (gfc_conv_expr_present): Check for DECL_ARTIFICIAL
        for the VALUE hidden argument avoiding -fallow-underscore issues.
        * trans-decl.c (create_function_arglist): Also set
        GFC_DECL_OPTIONAL_ARGUMENT for per-value arguments.
        * f95-lang.c (LANG_HOOKS_OMP_CHECK_OPTIONAL_ARGUMENT):
        Renamed from LANG_HOOKS_OMP_IS_OPTIONAL_ARGUMENT; point
        to gfc_omp_check_optional_argument.
        * trans.h (gfc_omp_check_optional_argument): Subsitutes
        gfc_omp_is_optional_argument declaration.
        * trans-openmp.c (gfc_omp_is_optional_argument): Make static.
        (gfc_omp_check_optional_argument): New function.

        libgomp/
        * testsuite/libgomp.fortran/use_device_ptr-optional-1.f90: Extend.
        * testsuite/libgomp.fortran/use_device_ptr-optional-2.f90: New.


Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>

From-SVN: r278046
2019-11-11 10:19:29 +01:00
Thomas Schwinge
bfa1837b01 Torture testing: 'libgomp.fortran/target9.f90'
libgomp/
	* testsuite/libgomp.fortran/target9.f90: Specify 'dg-do run'.

From-SVN: r278045
2019-11-11 09:50:40 +01:00
Thomas Schwinge
6d4868f3a1 Torture testing: 'libgomp.fortran/use_device_addr-3.f90', 'libgomp.fortran/use_device_addr-4.f90', 'libgomp.fortran/use_device_ptr-1.f90'
libgomp/
	* testsuite/libgomp.fortran/use_device_addr-3.f90: Specify 'dg-do
	run'.
	* testsuite/libgomp.fortran/use_device_addr-4.f90: Likewise.
	* testsuite/libgomp.fortran/use_device_ptr-1.f90: Likewise.

From-SVN: r278044
2019-11-11 09:50:29 +01:00
Frederik Harwath
5d183d1740 Warn about inconsistent OpenACC nested reduction clauses
OpenACC (cf. OpenACC 2.7, section 2.9.11. "reduction clause";
	this was first clarified by OpenACC 2.6) requires that, if a
	variable is used in reduction clauses on two nested loops, then
	there must be reduction clauses for that variable on all loops
	that are nested in between the two loops and all these reduction
	clauses must use the same operator.
	This commit introduces a check for that property which reports
	warnings if it is violated.

	2019-11-06  Gergö Barany  <gergo@codesourcery.com>
	            Frederik Harwath  <frederik@codesourcery.com>
	            Thomas Schwinge  <thomas@codesourcery.com>

	gcc/
	* omp-low.c (struct omp_context): New fields
	local_reduction_clauses, outer_reduction_clauses.
	(new_omp_context): Initialize these.
	(scan_sharing_clauses): Record reduction clauses on OpenACC constructs.
	(scan_omp_for): Check reduction clauses for incorrect nesting.
	gcc/testsuite/
	* c-c++-common/goacc/nested-reductions-warn.c: New test.
	* c-c++-common/goacc/nested-reductions.c: New test.
	* gfortran.dg/goacc/nested-reductions-warn.f90: New test.
	* gfortran.dg/goacc/nested-reductions.f90: New test.
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c:
	Add expected warnings about missing reduction clauses.
	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c:
	Likewise.

	Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

From-SVN: r277875
2019-11-06 10:43:52 +00:00
Tobias Burnus
d8f03b0341 libgomp - remove dg-run if dg-options O2 is used
* testsuite/libgomp.fortran/pr66199-1.f90: Remove
	'dg-do run' (implies torture test) as 'dg-options "O2"' is used.
	* testsuite/libgomp.fortran/pr66199-2.f90: Ditto.
	* testsuite/libgomp.fortran/taskloop2.f90: Ditto.
	* testsuite/libgomp.fortran/taskloop3.f90: Ditto.
	* testsuite/libgomp.fortran/taskloop4.f90: Ditto.

From-SVN: r277776
2019-11-04 14:18:50 +01:00
Tobias Burnus
12fd2ec5ae libgomp/testsuite - use unique numbers with Fortran's 'stop'
PR fortran/92305
        * testsuite/libgomp.fortran/allocatable2.f90: Use
        unique numbers with 'stop'.
        * testsuite/libgomp.fortran/use_device_addr-1.f90: Ditto.
        * testsuite/libgomp.fortran/use_device_addr-2.f90: Ditto.
        * testsuite/libgomp.fortran/use_device_ptr-1.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/lib-15.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/pset-1.f90: Ditto.

From-SVN: r277769
2019-11-04 11:01:22 +01:00
Tobias Burnus
92e63bd2df OpenMP] use_device_addr/use_device_ptr with Fortran allocatable/pointer arrays
gcc/fortran/
        * f95-lang.c (LANG_HOOKS_OMP_ARRAY_DATA): Set to gfc_omp_array_data.
        * trans-array.c (gfc_conv_descriptor_data_get): Handle also
        REFERENCE_TYPE.
        * trans-openmp.c (gfc_omp_array_data): New.
        * trans.h (gfc_omp_array_data): New prototype.

        gcc/
        * hooks.c (hook_tree_tree_bool_null): New.
        * hooks.h (hook_tree_tree_bool_null): Declare.
        * langhooks-def.h (LANG_HOOKS_OMP_ARRAY_DATA): Define.
        (LANG_HOOKS_DECLS): Add it.
        * langhooks.h (lang_hooks_for_decls): Add omp_array_data.
        * omp-low.c (install_var_field): New mode for Fortran descriptor arrays.
        (lower_omp_target): Handle Fortran array with descriptor in
        OMP_CLAUSE_USE_DEVICE_ADDR/OMP_CLAUSE_USE_DEVICE_PTR.

        libgomp/
        * testsuite/libgomp.fortran/use_device_addr-1.f90 (test_nullptr_1,
        test_dummy_opt_nullptr_callee_1): Add present but unallocated test.
        * testsuite/libgomp.fortran/use_device_addr-2.f90: Likewise.
        * testsuite/libgomp.fortran/use_device_addr-3.f90: New.
        * testsuite/libgomp.fortran/use_device_addr-4.f90: New.
        * testsuite/testsuite/libgomp.fortran/use_device_ptr-1.f90: New.

From-SVN: r277705
2019-11-01 08:59:23 +01:00
Tobias Burnus
91ae0a46fd Fortran/OpenMP] Don't create "alloc:" for 'target exit data'
gcc/
        * gimplify.c (gimplify_scan_omp_clauses): Remove FE-generated
        GOMP_MAP_TO_PSET and GOMP_MAP_POINTER mapping for 'target update'
        and 'target exit data'.

        libgomp/
        * testsuite/libgomp.fortran/target9.f90: New.

From-SVN: r277631
2019-10-30 17:32:38 +01:00
Tobias Burnus
c425e66b27 libgomp/testsuite – use 'stop'
libgomp/
	* testsuite/libgomp.fortran/: Replace 'STOP' by 'stop'.

From-SVN: r277609
2019-10-30 12:44:54 +01:00
Tobias Burnus
01509e2f04 libgomp/testsuite – use 'stop' and 'dg-do run'
libgomp/
        * testsuite/libgomp.fortran/target-simd.f90: Use stop not abort.
        * testsuite/libgomp.fortran/use_device_ptr-optional-1.f90:
        Ditto; add 'dg-do run' for torture testing.
        * testsuite/libgomp.fortran/lastprivate1.f90:  Add 'dg-do run'.
        * testsuite/libgomp.fortran/lastprivate2.f90: Ditto.
        * testsuite/libgomp.fortran/nestedfn4.f90: Ditto.
        * testsuite/libgomp.fortran/pr25219.f90: Ditto.
        * testsuite/libgomp.fortran/pr28390.f: Ditto.
        * testsuite/libgomp.fortran/pr35130.f90: Ditto.
        * testsuite/libgomp.fortran/pr90779.f90: Ditto.
        * testsuite/libgomp.fortran/task2.f90: Ditto.
        * testsuite/libgomp.fortran/taskgroup1.f90: Ditto.
        * testsuite/libgomp.fortran/taskloop1.f90: Ditto.
        * testsuite/libgomp.fortran/use_device_addr-1.f90: Ditto.
        * testsuite/libgomp.fortran/use_device_addr-2.f90: Ditto.
        * testsuite/libgomp.fortran/workshare1.f90: Ditto.
        * testsuite/libgomp.fortran/workshare2.f90: Ditto.

From-SVN: r277606
2019-10-30 11:33:58 +01:00
Tobias Burnus
ffbdd78a4a Fortran] OpenACC – libgomp/testsuite – use 'stop' and 'dg-do run'
* testsuite/libgomp.oacc-fortran/abort-1.f90: Add 'dg-do run'.
        * testsuite/libgomp.oacc-fortran/abort-2.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/lib-1.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/common-block-1.f90:
        Use 'stop' not abort().
        * testsuite/libgomp.oacc-fortran/common-block-2.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/common-block-3.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/data-1.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/data-2.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/data-5.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/dummy-array.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/gemm-2.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/gemm.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/host_data-2.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/host_data-3.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/host_data-4.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-collapse-3.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-collapse-4.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-independent.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-loop-1.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-map-1.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-parallel-loop-data-enter-exit.f95:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-1.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-2.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-3.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-loop-gang-6.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-vector-1.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-vector-2.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-worker-1.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-worker-2.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-worker-3.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-worker-4.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-worker-5.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-worker-6.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-private-vars-worker-7.f90:
        Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-reduction-1.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/lib-12.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/lib-13.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/lib-14.f90: Ditto.
        * testsuite/libgomp.oacc-fortran/kernels-acc-loop-reduction-2.f90:
        Likewise and also add 'dg-do run'.
        * testsuite/libgomp.oacc-fortran/kernels-acc-loop-reduction.f90:
        Ditto.

From-SVN: r277503
2019-10-28 08:39:26 +01:00
Cesar Philippidis
ec0846134e [Fortran] OpenACC – permit common blocks in some clauses
2019-10-25  Cesar Philippidis <cesar@codesourcery.com>
	    Tobias Burnus  <tobias@codesourcery.com>

	gcc/fortran/
	* openmp.c (gfc_match_omp_map_clause): Add and pass allow_commons
	argument.
	(gfc_match_omp_clauses): Update calls to permit common blocks for
	OpenACC's copy/copyin/copyout, create/delete, host,
	pcopy/pcopy_in/pcopy_out, present_or_copy, present_or_copy_in,
	present_or_copy_out, present_or_create and self.

	gcc/
	* gimplify.c (oacc_default_clause): Privatize fortran common blocks.
	(omp_notice_variable): Defer the expansion of DECL_VALUE_EXPR for
	common block decls.
    
	gcc/testsuite/
	* gfortran.dg/goacc/common-block-1.f90: New test.
	* gfortran.dg/goacc/common-block-2.f90: New test.
	* gfortran.dg/goacc/common-block-3.f90: New test.
    
	libgomp/
	* testsuite/libgomp.oacc-fortran/common-block-1.f90: New test.
	* testsuite/libgomp.oacc-fortran/common-block-2.f90: New test.
	* testsuite/libgomp.oacc-fortran/common-block-3.f90: New test.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>

From-SVN: r277451
2019-10-25 16:28:40 +02:00
Jakub Jelinek
77ef339456 re PR libgomp/92081 (FAIL: libgomp.fortran/target-simd.f90 execution test)
PR libgomp/92081
	* testsuite/libgomp.fortran/target-simd.f90: Iterate from 1 rather
	than 0.

From-SVN: r276956
2019-10-14 10:48:42 +02:00
Tobias Burnus
08c14aaaab [OpenMP,Fortran] Fix several OpenMP use_device_addr/map/update errors
gcc/fortran/
	* f95-lang.c (LANG_HOOKS_OMP_IS_ALLOCATABLE_OR_PTR): Re-define to
	gfc_omp_is_allocatable_or_ptr.
	* trans-decl.c (create_function_arglist): Set GFC_DECL_OPTIONAL_ARGUMENT
	only if not passed by value.
	* trans-openmp.c (gfc_omp_is_allocatable_or_ptr): New.
	(gfc_trans_omp_clauses): For MAP, handle (present) optional arguments;
	for target update, handle allocatable/pointer scalars.
	* trans.h (gfc_omp_is_allocatable_or_ptr): Declare.

	gcc/
	* langhooks-def.h (LANG_HOOKS_OMP_IS_ALLOCATABLE_OR_PTR): Define.
	(LANG_HOOKS_DECLS): Add it.
	* langhooks.h (lang_hooks_for_decls): Add omp_is_allocatable_or_ptr;
	update comment for omp_is_optional_argument.
	* omp-general.c (omp_is_allocatable_or_ptr): New.
	* omp-general.h (omp_is_allocatable_or_ptr): Declare.
	* omp-low.c (scan_sharing_clauses, lower_omp_target): Handle
	Fortran's optional arguments and allocatable/pointer scalars
	with use_device_addr.

	libgomp/
	* testsuite/libgomp.fortran/use_device_addr-1.f90: New.
	* testsuite/libgomp.fortran/use_device_addr-2.f90: New.

From-SVN: r276875
2019-10-11 11:17:49 +02:00
Thomas Schwinge
6bbead0c5a [PR92036] Add 'libgomp.oacc-c-c++-common/data-firstprivate-1.c'
libgomp/
	PR middle-end/92036
	* testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: New
	file.

From-SVN: r276757
2019-10-09 13:31:14 +02:00
Tobias Burnus
eba3dee8eb re PR testsuite/91884 (libgomp testsuite: (not) using a specific driver for C++, Fortran)
2019-10-09  Tobias Burnus  <tobias@codesourcery.com>

        PR testsuite/91884
        * testsuite/libgomp.fortran/fortran.exp: Conditionally
        add -lquadmath.
        * testsuite/libgomp.oacc-fortran/fortran.exp: Ditto.

From-SVN: r276754
2019-10-09 10:37:44 +02:00
Jakub Jelinek
7d48e14fc6 re PR libgomp/92028 (OpenACC 'host_data' execution test regressions with nvptx offloading)
PR libgomp/92028
	* target.c (gomp_map_vars_internal): Readd the previous
	GOMP_MAP_USE_DEVICE_PTR handling code in the first loop,
	though do that just in the !not_found_cnt case.

From-SVN: r276753
2019-10-09 09:33:02 +02:00
Tobias Burnus
65b67cf390 Fortran - fix OpenMP 'target simd'
gcc/fortran/
	* parse.c (parse_executable): Add missing ST_OMP_TARGET_SIMD.

	libgomp/
	* testsuite/libgomp.fortran/target-simd.f90: New.

From-SVN: r276698
2019-10-08 14:30:44 +02:00
Julian Brown
6c7e076b74 Libgomp magic offset value self-documentation
2019-10-02  Julian Brown  <julian@codesourcery.com>
	    Cesar Philippidis  <cesar@codesourcery.com>

	libgomp/
	* libgomp.h (OFFSET_INLINED, OFFSET_POINTER, OFFSET_STRUCT): Define.
	* target.c (FIELD_TGT_EMPTY): Define.
	(gomp_map_val): Use OFFSET_* macros instead of magic constants.  Write
	as switch instead of list of ifs.
	(gomp_map_vars_internal): Use OFFSET_* and FIELD_TGT_EMPTY macros.

Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>

From-SVN: r276519
2019-10-03 17:48:36 +00:00
Andreas Tobler
d61bff850d Testsuite, remove alloca header
2019-10-02  Andreas Tobler  <andreast@gcc.gnu.org>

	* testsuite/libgomp.oacc-c-c++-common/loop-default.h: Remove alloca.h
	include. Replace alloca () with __builtin_alloca ().
	* testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c: Likewise.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

From-SVN: r276479
2019-10-02 21:05:35 +02:00
Tobias Burnus
67c259509c Fix omp target issue with Fortran optional arguments
gcc/
        * omp-low.c (lower_omp_target): Dereference optional argument
        to work with the right pointer.

        gcc/testsuite/
        * libgomp/testsuite/libgomp.fortran/use_device_ptr-optional-1.f90: New.

From-SVN: r276445
2019-10-02 12:33:42 +02:00
Jakub Jelinek
810f316dd6 configure.ac: Remove GCC_HEADER_STDINT(gstdint.h).
* configure.ac: Remove GCC_HEADER_STDINT(gstdint.h).
	* libgomp.h: Include <stdint.h> instead of "gstdint.h".
	* oacc-parallel.c: Don't include "libgomp_g.h".
	* plugin/plugin-hsa.c: Include <stdint.h> instead of "gstdint.h".
	* plugin/plugin-nvptx.c: Don't include "gstdint.h".
	* aclocal.m4: Regenerated.
	* config.h.in: Regenerated.
	* configure: Regenerated.
	* Makefile.in: Regenerated.

From-SVN: r276389
2019-10-01 09:51:46 +02:00
Kwok Cheung Yeung
d7f9ee981f libgomp_g.h: Include stdint.h instead of gstdint.h.
2019-09-30  Kwok Cheung Yeung  <kcy@codesourcery.com>

        * libgomp_g.h: Include stdint.h instead of gstdint.h.

From-SVN: r276301
2019-09-30 16:16:34 +02:00
Maciej W. Rozycki
e9085da528 Regenerate configure' scripts for uclinuxfdpiceabi' libtool.m4 update
A change made with r275564 ("[ARM/FDPIC v6 02/24] [ARM] FDPIC: Handle 
arm*-*-uclinuxfdpiceabi in configure scripts") to libtool.m4 has not 
regenerated all the `configure' scripts affected.  Fix it.

	gcc/
	* configure: Regenerate.

	libatomic/
	* configure: Regenerate.

	libbacktrace/
	* configure: Regenerate.

	libcc1/
	* configure: Regenerate.

	libffi/
	* configure: Regenerate.

	libgfortran/
	* configure: Regenerate.

	libgomp/
	* configure: Regenerate.

	libhsail-rt/
	* configure: Regenerate.

	libitm/
	* configure: Regenerate.

	libobjc/
	* configure: Regenerate.

	liboffloadmic/
	* configure: Regenerate.

	libphobos/
	* configure: Regenerate.

	libquadmath/
	* configure: Regenerate.

	libsanitizer/
	* configure: Regenerate.

	libssp/
	* configure: Regenerate.

	libstdc++-v3/
	* configure: Regenerate.

	libvtv/
	* configure: Regenerate.

	lto-plugin/
	* configure: Regenerate.

	zlib/
	* configure: Regenerate.

From-SVN: r276213
2019-09-27 21:24:42 +00:00
Tobias Burnus
c28712beb4 libgomp plugin - init string
libgomp/
2019-09-13  Tobias Burnus  <tobias@codesourcery.com>

        * plugin/plugin-hsa.c (hsa_warn, hsa_fatal, hsa_error): Ensure
        string is initialized.

From-SVN: r275703
2019-09-13 20:14:02 +02:00
Florian Weimer
3c086f8dec Fix GCC_LINUX_FUTEX to work with C99 compilers
Without this change, libstdc++ is built without futex symbols if GCC
rejects implicit function declarations by default.

From-SVN: r275454
2019-09-06 12:27:36 +02:00
Chung-Lin Tang
c6c2d1bc9b re PR other/79543 (Inappropriate "ld --version" checking)
2019-09-03  Chung-Lin Tang <cltang@codesourcery.com>

	libatomic/
	PR other/79543
	* acinclude.m4 (LIBAT_CHECK_LINKER_FEATURES): Fix GNU ld --version
	scanning to conform to the GNU Coding Standards.
	* configure: Regenerate.

	libffi/
	PR other/79543
	* acinclude.m4 (LIBAT_CHECK_LINKER_FEATURES): Fix GNU ld --version
	scanning to conform to the GNU Coding Standards.
	* configure: Regenerate.

	libgomp/
	PR other/79543
	* acinclude.m4 (LIBGOMP_CHECK_LINKER_FEATURES): Fix GNU ld --version
	scanning to conform to the GNU Coding Standards.
	* configure: Regenerate.

	libitm/
	PR other/79543
	* acinclude.m4 (LIBITM_CHECK_LINKER_FEATURES): Fix GNU ld --version
	scanning to conform to the GNU Coding Standards.
	* configure: Regenerate.

	libstdc++-v3/
	PR other/79543
	* acinclude.m4 (GLIBCXX_CHECK_LINKER_FEATURES): Fix GNU ld --version
	scanning to conform to the GNU Coding Standards.
	* configure: Regenerate.

From-SVN: r275341
2019-09-03 14:10:26 +00:00
Jakub Jelinek
e62506f362 re PR libgomp/91530 (Several libgomp.*/scan-* tests FAIL without avx_runtime)
PR libgomp/91530
	* config/i386/sse.md (vec_shl_<mode>, vec_shr_<mode>): Use
	V_128 iterator instead of VI_128.

	* testsuite/libgomp.c/scan-21.c: New test.
	* testsuite/libgomp.c/scan-22.c: New test.

From-SVN: r274985
2019-08-28 12:13:21 +02:00
Jakub Jelinek
5cb72d83bb re PR libgomp/91530 (Several libgomp.*/scan-* tests FAIL without avx_runtime)
PR libgomp/91530
	* config/i386/sse.md (vec_shl_<mode>, vec_shr_<mode>): Use
	V_128 iterator instead of VI_128.

	* testsuite/libgomp.c/scan-21.c: New test.
	* testsuite/libgomp.c/scan-22.c: New test.

From-SVN: r274984
2019-08-28 12:12:11 +02:00
Jakub Jelinek
0ad7981cb4 re PR libgomp/91530 (Several libgomp.*/scan-* tests FAIL without avx_runtime)
PR libgomp/91530
	* testsuite/libgomp.c/scan-11.c: Add -msse2 option for sse2_runtime
	targets.
	* testsuite/libgomp.c/scan-12.c: Likewise.
	* testsuite/libgomp.c/scan-13.c: Likewise.
	* testsuite/libgomp.c/scan-14.c: Likewise.
	* testsuite/libgomp.c/scan-15.c: Likewise.
	* testsuite/libgomp.c/scan-16.c: Likewise.
	* testsuite/libgomp.c/scan-17.c: Likewise.
	* testsuite/libgomp.c/scan-18.c: Likewise.
	* testsuite/libgomp.c/scan-19.c: Likewise.
	* testsuite/libgomp.c/scan-20.c: Likewise.
	* testsuite/libgomp.c++/scan-9.C: Likewise.
	* testsuite/libgomp.c++/scan-10.C: Likewise.
	* testsuite/libgomp.c++/scan-11.C: Likewise.
	* testsuite/libgomp.c++/scan-12.C: Likewise.
	* testsuite/libgomp.c++/scan-14.C: Likewise.
	* testsuite/libgomp.c++/scan-15.C: Likewise.
	* testsuite/libgomp.c++/scan-13.C: Likewise.  Use sse2_runtime
	instead of i?86-*-* x86_64-*-* as target for scan-tree-dump-times.
	* testsuite/libgomp.c++/scan-16.C: Likewise.

From-SVN: r274947
2019-08-27 12:45:55 +02:00
Thomas Koenig
1e67491a0d re PR libgomp/91473 (Test case libgomp.fortran/appendix-a/a.28.5.f90 is invalid)
2019-08-17  Thomas Koenig  <tkoenig@gcc.gnu.org>

    PR fortran/91473
    * testsuite/libgomp.fortran/appendix-a/a.28.5.f90: Add
    -std=legacy so invalid code in the test case is accepted.

From-SVN: r274602
2019-08-17 11:57:25 +00:00
Thomas Koenig
393fdeb1e4 re PR fortran/91424 (Extend warnings about DO loops)
2019-08-12  Thomas Koenig  <tkoenig@gcc.gnu.org>

	PR fortran/91424
	* frontend-passes.c (do_subscript): Do not warn for an
	expression a second time.  Do not warn about a zero-trip loop.
	(doloop_warn): Also look at contained namespaces.

2019-08-12  Thomas Koenig  <tkoenig@gcc.gnu.org>

	PR fortran/91424
	* gfortran.dg/do_subscript_3.f90: New test.
	* gfortran.dg/do_subscript_4.f90: New test.
	* gfortran.dg/pr70754.f90: Use indices that to not overflow.

2019-08-12  Thomas Koenig  <tkoenig@gcc.gnu.org>

	PR fortran/91422
	* testsuite/libgomp.oacc-fortran/routine-7.f90: Correct array
	dimension.

From-SVN: r274320
2019-08-12 20:21:37 +00:00
Jakub Jelinek
8860d2706d gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT for VLA helper variables on target data even if...
* gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT
	for VLA helper variables on target data even if not GOVD_FIRSTPRIVATE.
	(gimplify_scan_omp_clauses): For OMP_CLAUSE_USE_DEVICE_* use just
	GOVD_EXPLICIT flags.
	(gimplify_omp_workshare): For OMP_TARGET_DATA move all
	OMP_CLAUSE_USE_DEVICE_* clauses to the end of clauses chain.
	* omp-low.c (scan_sharing_clauses): For OMP_CLAUSE_USE_DEVICE_*
	call install_var_field with mask 11 instead of 3.
	(lower_omp_target): For OMP_CLAUSE_USE_DEVICE_* use pass
	(splay_tree_key) &DECL_UID (var) to build_sender_ref instead of var.
gcc/c/
	* c-typeck.c (c_finish_omp_clauses): For C_ORT_OMP
	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
	instead of generic_head to track duplicates.
gcc/cp/
	* semantics.c (finish_omp_clauses): For C_ORT_OMP
	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
	instead of generic_head to track duplicates.
libgomp/
	* target.c (gomp_map_vars_internal): For GOMP_MAP_USE_DEVICE_PTR
	perform the lookup in the first loop only if !not_found_cnt, otherwise
	perform lookups for it in the second loop guarded with
	if (not_found_cnt || has_firstprivate).
	* testsuite/libgomp.c/target-37.c: New test.
	* testsuite/libgomp.c++/target-22.C: New test.

From-SVN: r274206
2019-08-08 08:39:02 +02:00
Jakub Jelinek
398e3feb8a tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR OpenMP description.
* tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR
	OpenMP description.  Add OMP_CLAUSE_USE_DEVICE_ADDR clause.
	* tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries
	for OMP_CLAUSE_USE_DEVICE_ADDR clause.
	(walk_tree_1): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Likewise.
	* gimplify.c (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses):
	Likewise.
	* omp-low.c (scan_sharing_clauses, lower_omp_target): Likewise.
	Treat OMP_CLAUSE_USE_DEVICE_ADDR like OMP_CLAUSE_USE_DEVICE_PTR
	clause with array or reference to array types, no matter what type
	except for reference it has.
gcc/c-family/
	* c-pragma.h (enum pragma_omp_clause): Add
	PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.  Set PRAGMA_OACC_CLAUSE_USE_DEVICE
	equal to PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR instead of being a separate
	enumeration value.
gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause.
	(c_parser_omp_clause_use_device_addr): New function.
	(c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(c_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
	like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
	map or use_device_* clauses.
	* c-typeck.c (c_finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
	in OpenMP, require pointer type rather than pointer or array type.
	Handle OMP_CLAUSE_USE_DEVICE_ADDR.
gcc/cp/
	* parser.c (cp_parser_omp_clause_name): Parse use_device_addr clause.
	(cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(cp_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
	like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
	map or use_device_* clauses.
	* semantics.c (finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
	in OpenMP, require pointer or reference to pointer type rather than
	pointer or array or reference to pointer or array type. Handle
	OMP_CLAUSE_USE_DEVICE_ADDR.
	* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
gcc/testsuite/
	* c-c++-common/gomp/target-data-1.c (foo): Use use_device_addr clause
	instead of use_device_ptr clause where required by OpenMP 5.0, add
	further tests for both use_device_ptr and use_device_addr clauses.
libgomp/
	* testsuite/libgomp.c/target-18.c (struct S): New type.
	(foo): Use use_device_addr clause instead of use_device_ptr clause
	where required by OpenMP 5.0, add further tests for both use_device_ptr
	and use_device_addr clauses.
	* testsuite/libgomp.c++/target-9.C (struct S): New type.
	(foo): Use use_device_addr clause instead of use_device_ptr clause
	where required by OpenMP 5.0, add further tests for both use_device_ptr
	and use_device_addr clauses.  Add t and u arguments.
	(main): Adjust caller.

From-SVN: r274159
2019-08-07 09:27:10 +02:00
Jakub Jelinek
d81ab49d05 tree.h (OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV): Rename to ...
* tree.h (OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV): Rename to ...
	(OMP_CLAUSE_LASTPRIVATE_LOOP_IV): ... this.  Adjust comment.
	* gimplify.c (gimple_add_tmp_var): In SIMD contexts, turn addressable
	new vars into GOVD_PRIVATE rather than GOVD_LOCAL.
	(gimplify_omp_for): Don't do C++ random access iterator clause
	adjustments on combined constructs from OMP_LOOP.  For OMP_LOOP,
	don't predetermine the artificial iterator in case of C++ random
	access iterators as lastprivate, but private.  For OMP_LOOP, force
	bind expr around simd body and force for_pre_body before the
	construct.  Use OMP_CLAUSE_LASTPRIVATE_LOOP_IV instead of
	OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV.
	(gimplify_omp_loop): Add firstprivate clauses on OMP_PARALLEL for
	diff var of C++ random access iterators.  Handle
	OMP_CLAUSE_FIRSTPRIVATE.  For OMP_CLAUSE_LASTPRIVATE_LOOP_IV, if
	not outermost also add OMP_CLAUSE_FIRSTPRIVATE, and in both cases
	clear OMP_CLAUSE_LASTPRIVATE_LOOP_IV on the lastprivate clause
	on the OMP_FOR and OMP_DISTRIBUTE constructs if any.
	* omp-low.c (lower_rec_input_clauses): For
	OMP_CLAUSE_LASTPRIVATE_LOOP_IV on simd copy construct the private
	variables instead of default constructing them.
	(lower_lastprivate_clauses): Use OMP_CLAUSE_LASTPRIVATE_LOOP_IV
	instead of OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV and move the
	is_taskloop_ctx check from the assert to the guarding condition.
gcc/cp/
	* parser.c (cp_parser_omp_for_loop): For OMP_LOOP, ignore parallel
	clauses and predetermine iterator as lastprivate.
	* semantics.c (handle_omp_for_class_iterator): Use
	OMP_CLAUSE_LASTPRIVATE_LOOP_IV instead of
	OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV, set it for lastprivate also
	on OMP_LOOP construct.  If a clause is missing for class iterator
	on OMP_LOOP, add firstprivate clause, and if there is private
	clause, turn it into firstprivate too.
	(finish_omp_for): Formatting fix.  For OMP_LOOP, adjust
	OMP_CLAUSE_LASTPRIVATE_LOOP_IV clause CP_CLAUSE_INFO, so that it
	uses copy ctor instead of default ctor.
	* cp-gimplify.c (cp_gimplify_expr): Handle OMP_LOOP like
	OMP_DISTRIBUTE etc.
	(cp_fold_r): Likewise.
	(cp_genericize_r): Likewise.
	(cxx_omp_finish_clause): Also finish lastprivate clause with
	OMP_CLAUSE_LASTPRIVATE_LOOP_IV flag.
	* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_BIND.
	(tsubst_omp_for_iterator): For OMP_LOOP, ignore parallel
	clauses and predetermine iterator as lastprivate.
	* constexpr.c (potential_constant_expression_1): Handle OMP_LOOP
	like OMP_DISTRIBUTE etc.
libgomp/
	* testsuite/libgomp.c++/loop-13.C: New test.
	* testsuite/libgomp.c++/loop-14.C: New test.
	* testsuite/libgomp.c++/loop-15.C: New test.

From-SVN: r274138
2019-08-06 09:26:32 +02:00
Jakub Jelinek
c3ac76aa40 re PR middle-end/91301 (ICE in omp_add_variable on random access iterator distribute parallel for private (iterator))
PR middle-end/91301
	* gimplify.c (gimplify_omp_for): If for class iterator on
	distribute parallel for there is no data sharing clause
	on inner_for_stmt, look for private clause on combined
	parallel too and if found, move it to inner_for_stmt.

	* testsuite/libgomp.c++/for-27.C: New test.

From-SVN: r273922
2019-07-31 09:49:56 +02:00
Steven G. Kargl
8dc63166e0 arith.c (gfc_convert_integer, [...]): Move to ...
2019-07-23  Steven G. Kargl  <kargl@gcc.gnu.org>

	* arith.c (gfc_convert_integer, gfc_convert_real, gfc_convert_complex):
	Move to ...
	* primary.c (convert_integer, convert_real, convert_complex): ... here.
	Rename and make static functions.
	(match_integer_constant): Use convert_integer
	(match_real_constant): Use convert_real.
	(match_complex_constant: Use convert_complex.
	* arith.h (gfc_convert_integer, gfc_convert_real, gfc_convert_complex):
	Remove prototypes.
	* array.c (match_array_cons_element): A BOZ cannot be a data 
	statement value.  Jump to a common exit point.
	* check.c (gfc_invalid_boz): New function.  Emit error or warning
	for a BOZ in an invalid context.
	(boz_args_check): Move to top of file to prevent need of forward
	declaration.
	(is_boz_constant): New function.  Check that BOZ expr is constant.
	(gfc_b	z2real): New function. In-place conversion of BOZ literal
	constant to REAL in accordance to F2018.
	(gfc_boz2int): New function. In-place conversion of BOZ literal
 	onstant to INTEGER in accordance to F2018.
	(gfc_check_achar, gfc_check_char, gfc_check_float): Use gfc_invalid_boz.
	Convert BOZ as needed.
	(gfc_check_bge_bgt_ble_blt): Enforce F2018 requirements on BGE, 
	BGT, BLE, and BLT intrinsic functions.
	(gfc_check_cmplx): Re-organize to check kind, if present, first.
	Convert BOZ real and/or imaginary parts as needed in accordance to
	F2018.
	(gfc_check_complex):  Use gfc_invalid_boz.  Convert BOZ as needed.
	(gfc_check_dcmplx, gfc_check_dble ): Convert BOZ as needed.
	(gfc_check_dshift):  Make dshift[lr] conform to F2018 standard.
	 gfc_check_float (gfc_expr *a)
	(gfc_check_iand_ieor_ior):  Make IAND, IEOR, and IOR conform to 
	F2018 standard.
	(gfc_check_int): Conform to F2018 standard.
	(gfc_check_intconv): Deprecate SHORT and LONG aliases for INT2 and
	INT.  Simply return for a BOZ argument. See gfc_simplify_intconv.
	(gfc_check_merge_bits): Make MERGE_BITS conform to Fortran 2018
	standard.
	(gfc_check_real): Remove incorrect comment. Check kind, if present,
	first.  Simply return for a BOZ argument. See gfc_simplify_real.
	(gfc_check_and): Re-do error handling for BOZ arguments.  Remove
	special casing ts.type != BT_INTEGER or BT_LOGICAL.
	* decl.c (match_old_style_init): Check for BOZ in old-style
	initialization.  Issue error or warning depending on
	-fallow-invalid-boz option.  Issue error if variable is not an
	INTEGER or REAL and the value is BOZ.
	* expr.c (gfc_copy_expr): Copy a BT_BOZ gfc_expr.
	(gfc_check_assign): Re-do error handling for a BOZ in an assignment
	statement.  Do in-place conversion of RHS based on LHS type of
	INTEGER or REAL.
	* gfortran.h (gfc_expr): Add a boz component.  Remove is_boz component.
	(gfc_boz2int, gfc_boz2real, gfc_invalid_boz): New prototypes.
	* interface.c (gfc_extend_assign): Guard against replacing an 
	intrinsic involving a BOZ literal constant on RHS.
	* invoke.texi: Doument -fallow-invalid-boz.
	* lang.opt: New option. -fallow-invalid-boz.
	* libgfortran.h (bt): Elevate BOZ to a basic type.
	* misc.c (gfc_basic_typename, gfc_typename): Translate BT_BOZ to BOZ.
	* primary.c (convert_integer, convert_real, convert_complex): to here.
	Rename and make static functions.
	* primary.c(match_boz_constant): Rewrite parsing of a BOZ. Re-do
	error handling.  Deprecate 'X' for hexidecimal and postfix notation.
	Use -fallow-invalid-boz and gfc_invalid_boz to accept deprecated code.
	* resolve.c (resolve_ordinary_assign): Rework a RHS that is a
	BOZ literal constant.  Use gfc_invalid_boz to allow previous
	nonstandard behavior.  Remove range checking of BOZ conversion.
	* simplify.c (convert_boz): Remove function.
	(simplify_cmplx): Remove conversion of BOZ constants, because
	conversion is done in gfc_check_cmplx.
	(gfc_simplify_float): Remove conversion of BOZ constant, because
	conversion is done in gfc_check_float.
	(simplify_intconv): Use gfc_boz2int to convert BOZ to INTEGER.
	Remove range checking for BOZ conversion.
	(gfc_simplify_real): Use k, if present, to determine kind.  Convert
	BOZ to REAL.  Remove range checking for BOZ conversion.
	target-memory.c (gfc_convert_boz): Rewrite to deal with convert of
	a BOZ to a REAL value.

2019-07-23  Steven G. Kargl  <kargl@gcc.gnu.org>

	* gfortran.dg/achar_5.f90: Fix for new BOZ handling.
	* arithmetic_overflow_1.f90: Ditto.
	* gfortran.dg/boz_11.f90: Ditto.
	* gfortran.dg/boz_12.f90: Ditto.
	* gfortran.dg/boz_4.f90: Ditto.
	* gfortran.dg/boz_5.f90: Ditto.
	* gfortran.dg/boz_6.f90: Ditto.
	* gfortran.dg/boz_7.f90: Ditto.
	* gfortran.dg/boz_8.f90: Ditto.
	* gfortran.dg/dec_structure_6.f90: Ditto.
	* gfortran.dg/dec_union_1.f90: Ditto.
	* gfortran.dg/dec_union_2.f90: Ditto.
	* gfortran.dg/dec_union_5.f90: Ditto.
	* gfortran.dg/dshift_3.f90: Ditto.
	* gfortran.dg/gnu_logical_2.f90: Ditto.
	* gfortran.dg/int_conv_1.f90: Ditto.
	* gfortran.dg/ishft_1.f90: Ditto.
	* gfortran.dg/nan_4.f90: Ditto.
	* gfortran.dg/no_range_check_3.f90: Ditto.
	* gfortran.dg/pr16433.f: Ditto.
	* gfortran.dg/pr44491.f90: Ditto.
	* gfortran.dg/pr58027.f90: Ditto.
	* gfortran.dg/pr81509_2.f90: Ditto.
	* gfortran.dg/unf_io_convert_1.f90: Ditto.
	* gfortran.dg/unf_io_convert_2.f90: Ditto.
	* gfortran.fortran-torture/execute/intrinsic_fraction_exponent.f90:
	Ditto.
	* gfortran.fortran-torture/execute/intrinsic_mvbits.f90: Ditto.
	* gfortran.fortran-torture/execute/intrinsic_nearest.f90: Ditto.
	* gfortran.fortran-torture/execute/seq_io.f90: Ditto.
	* gfortran.dg/gnu_logical_1.F: Delete test.
	* gfortran.dg/merge_bits_3.f90: New test.
	* gfortran.dg/merge_bits_3.f90: Ditto.
	* gfortran.dg/boz_int.f90: Ditto.
	* gfortran.dg/boz_bge.f90: Ditto.
	* gfortran.dg/boz_complex_1.f90: Ditto.
	* gfortran.dg/boz_complex_2.f90: Ditto.
	* gfortran.dg/boz_complex_3.f90: Ditto.
	* gfortran.dg/boz_dble.f90: Ditto.
	* gfortran.dg/boz_dshift_1.f90: Ditto.
	* gfortran.dg/boz_dshift_2.f90: Ditto.
	* gfortran.dg/boz_float_1.f90: Ditto.
	* gfortran.dg/boz_float_2.f90: Ditto.
	* gfortran.dg/boz_float_3.f90: Ditto.
	* gfortran.dg/boz_iand_1.f90: Ditto.
	* gfortran.dg/boz_iand_2.f90: Ditto.

2019-07-23  Steven G. Kargl  <kargl@gcc.gnu.org>

	* testsuite/libgomp.fortran/reduction4.f90: Update BOZ usage
	* testsuite/libgomp.fortran/reduction5.f90: Ditto.

From-SVN: r273747
2019-07-23 21:43:21 +00:00
Jakub Jelinek
554a530ff8 tree.def (OMP_LOOP): New tree code.
* tree.def (OMP_LOOP): New tree code.
	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_BIND.
	(enum omp_clause_bind_kind): New enum.
	(struct tree_omp_clause): Add subcode.bind_kind.
	* tree.h (OMP_LOOP_CHECK): Rename to ...
	(OMP_LOOPING_CHECK): ... this.
	(OMP_FOR_BODY, OMP_FOR_CLAUSES, OMP_FOR_INIT, OMP_FOR_COND,
	OMP_FOR_INCR, OMP_FOR_PRE_BODY, OMP_FOR_ORIG_DECLS): Use
	OMP_LOOPING_CHECK instead of OMP_LOOP_CHECK.
	(OMP_CLAUSE_BIND_KIND): Define.
	* tree.c (omp_clause_num_ops, omp_clause_code_name): Add
	bind clause entries.
	(walk_tree_1): Handle OMP_CLAUSE_BIND.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	(dump_generic_node): Handle OMP_LOOP.
	* gimplify.c (enum omp_region_type): Add ORT_IMPLICIT_TARGET.
	(in_omp_construct): New variable.
	(is_gimple_stmt): Handle OMP_LOOP.
	(gimplify_scan_omp_clauses): For lastprivate don't set
	check_non_private if code == OMP_LOOP.  For reduction clause
	on OMP_LOOP combined with parallel or teams propagate as shared
	on the combined construct.  Handle OMP_CLAUSE_BIND.
	(gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_BIND.
	(gimplify_omp_for): Pass OMP_LOOP instead of OMP_{FOR,DISTRIBUTE}
	for constructs from a loop construct to gimplify_scan_omp_clauses.
	Don't predetermine iterator linear on OMP_SIMD from loop construct.
	(replace_reduction_placeholders, gimplify_omp_loop): New functions.
	(gimplify_omp_workshare): Use ORT_IMPLICIT_TARGET instead of trying
	to match the implicit ORT_TARGET construct around whole body.
	Temporarily clear in_omp_construct when processing body.
	(gimplify_expr): Handle OMP_LOOP.  For OMP_MASTER, OMP_TASKGROUP
	etc. temporarily set in_omp_construct when processing body.
	(gimplify_body): Create ORT_IMPLICIT_TARGET instead of ORT_TARGET.
	* omp-low.c (struct omp_context): Add loop_p.
	(build_outer_var_ref): Treat ctx->loop_p similarly to simd construct
	in that the original var might be private.
	(scan_sharing_clauses): Handle OMP_CLAUSE_BIND.
	(check_omp_nesting_restrictions): Adjust nesting restrictions for
	addition of loop construct.
	(scan_omp_1_stmt): Allow setjmp inside of loop construct.
gcc/c-family/
	* c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_LOOP.
	(enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_BIND.
	* c-pragma.c (omp_pragmas_simd): Add PRAGMA_OMP_LOOP entry.
	* c-common.h (enum c_omp_clause_split): Add C_OMP_CLAUSE_SPLIT_LOOP.
	* c-omp.c (c_omp_split_clauses): Add support for 4 new combined
	constructs with the loop construct.
gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Handle bind clause.
	(c_parser_omp_clause_bind): New function.
	(c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_BIND.
	(OMP_LOOP_CLAUSE_MASK): Define.
	(c_parser_omp_loop): New function.
	(c_parser_omp_parallel, c_parser_omp_teams): Handle parsing of
	loop combined with parallel or teams.
	(c_parser_omp_construct): Handle PRAGMA_OMP_LOOP.
	* c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_BIND.
gcc/cp/
	* cp-tree.h (OMP_FOR_GIMPLIFYING_P): Use OMP_LOOPING_CHECK
	instead of OMP_LOOP_CHECK.
	* parser.c (cp_parser_omp_clause_name): Handle bind clause.
	(cp_parser_omp_clause_bind): New function.
	(cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_BIND.
	(OMP_LOOP_CLAUSE_MASK): Define.
	(cp_parser_omp_loop): New function.
	(cp_parser_omp_parallel, cp_parser_omp_teams): Handle parsing of
	loop combined with parallel or teams.
	(cp_parser_omp_construct): Handle PRAGMA_OMP_LOOP.
	(cp_parser_pragma): Likewise.
	* pt.c (tsubst_expr): Handle OMP_LOOP.
	* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_BIND.
gcc/testsuite/
	* c-c++-common/gomp/cancel-1.c: Adjust expected diagnostic wording.
	* c-c++-common/gomp/clauses-1.c (foo, baz, bar): Add order(concurrent)
	clause where allowed.  Add combined constructs with loop with all
	possible clauses.
	(qux): New function.
	* c-c++-common/gomp/loop-1.c: New test.
	* c-c++-common/gomp/loop-2.c: New test.
	* c-c++-common/gomp/loop-3.c: New test.
	* c-c++-common/gomp/loop-4.c: New test.
	* c-c++-common/gomp/loop-5.c: New test.
	* c-c++-common/gomp/order-3.c: Adjust expected diagnostic wording.
	* c-c++-common/gomp/simd-setjmp-1.c: New test.
	* c-c++-common/gomp/teams-2.c: Adjust expected diagnostic wording.
libgomp/
	* testsuite/libgomp.c-c++-common/loop-1.c: New test.

From-SVN: r273621
2019-07-20 13:21:42 +02:00