Commit Graph

165 Commits

Author SHA1 Message Date
Tom de Vries
ab3f4b27ab [omp, ftracer] Don't duplicate blocks in SIMT region
When running the libgomp testsuite on x86_64-linux with nvptx accelerator on
the test-case included in this patch, we run into:
...
FAIL: libgomp.fortran/pr95654.f90 -O3 -fomit-frame-pointer -funroll-loops \
  -fpeel-loops -ftracer -finline-functions  execution test
...

The test-case is a minimal version of this FAIL:
...
FAIL: libgomp.fortran/pr66199-5.f90 -O3 -fomit-frame-pointer -funroll-loops \
  -fpeel-loops -ftracer -finline-functions  execution test
...
but that one has stopped failing at commit c2ebf4f10d "openmp: Add support
for non-rect simd and improve collapsed simd support".

The problem is that ftracer duplicates a block containing GOMP_SIMT_VOTE_ANY.

That is, before ftracer we have (dropping the GOMP_SIMT_ prefix):
...
bb4(ENTER_ALLOC)
*----------+
|           \
|            \
|             v
|             *
v             bb8
*<------------*
bb5(VOTE_ANY)
*-------------+
|             |
|             |
|             |
|             |
|             v
|             *
v             bb7(XCHG_IDX)
*<------------*
bb6(EXIT)
...

The XCHG_IDX internal-fn does inter-SIMT-lane communication, which for nvptx
maps onto shfl, an operator which has the requirement that the warp executing
the operator is convergent.  The warp diverges at bb4, and
reconverges at bb5, and does not diverge by going to bb7, so the shfl is
indeed executed by a convergent warp.

After ftracer, we have:
...
bb4(ENTER_ALLOC)
*----------+
|           \
|            \
|             \
|              \
v               v
*               *
bb5(VOTE_ANY)   bb8(VOTE_ANY)
*               *
|\             /|
| \  +--------+ |
|  \/           |
|  /\           |
| /  +----------v
|/              *
v               bb7(XCHG_IDX)
*<--------------*
bb6(EXIT)
...

The warp diverges again at bb5, but does not reconverge again before bb6, so
the shfl is executed by a divergent warp, which causes the FAIL.

Fix this by making ftracer ignore blocks containing ENTER_ALLOC, VOTE_ANY and
EXIT, effectively treating the SIMT region conservatively.

An argument can be made that the test needs to be added in a more
generic place, like gimple_can_duplicate_bb_p or some such, and that ftracer
then needs to use the generic test.  But that's a discussion with a much
broader scope, so I'm leaving that for another patch.

Bootstrapped and reg-tested on x86_64-linux.

Build on x86_64-linux with nvptx accelerator, tested with libgomp.

gcc/ChangeLog:

	PR fortran/95654
	* tracer.c (ignore_bb_p): Ignore GOMP_SIMT_ENTER_ALLOC,
	GOMP_SIMT_VOTE_ANY and GOMP_SIMT_EXIT.

libgomp/ChangeLog:

2020-10-05  Tom de Vries  <tdevries@suse.de>

	PR fortran/95654
	* testsuite/libgomp.fortran/pr95654.f90: New test.
2020-10-05 08:53:11 +02:00
Tobias Burnus
8b0a63e47c OpenMP: Add implicit declare target for nested procedures
gcc/ChangeLog:

	* omp-offload.c (omp_discover_implicit_declare_target): Also
	handled nested functions.

libgomp/ChangeLog:

	* testsuite/libgomp.fortran/declare-target-3.f90: New test.
2020-09-30 14:59:27 +02:00
Tobias Burnus
f74c87f85f libgomp.fortran/pr66199-5.f90: Make stop codes unique
libgomp/ChangeLog:

	PR fortran/95654
	* testsuite/libgomp.fortran/pr66199-5.f90: Make stop codes unique.
2020-09-22 19:16:34 +02:00
Tobias Burnus
972da55746 OpenMP/Fortran: Fix (re)mapping of allocatable/pointer arrays [PR96668]
gcc/cp/ChangeLog:

	PR fortran/96668
	* cp-gimplify.c (cxx_omp_finish_clause): Add bool openacc arg.
	* cp-tree.h (cxx_omp_finish_clause): Likewise
	* semantics.c (handle_omp_for_class_iterator): Update call.

gcc/fortran/ChangeLog:

	PR fortran/96668
	* trans.h (gfc_omp_finish_clause): Add bool openacc arg.
	* trans-openmp.c (gfc_omp_finish_clause): Ditto. Use
	GOMP_MAP_ALWAYS_POINTER with PSET for pointers.
	(gfc_trans_omp_clauses): Like the latter and also if the always
	modifier is used.

gcc/ChangeLog:

	PR fortran/96668
	* gimplify.c (gimplify_omp_for): Add 'bool openacc' argument;
	update omp_finish_clause calls.
	(gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses,
	gimplify_expr, gimplify_omp_loop): Update omp_finish_clause
	and/or gimplify_for calls.
	* langhooks-def.h (lhd_omp_finish_clause): Add bool openacc arg.
	* langhooks.c (lhd_omp_finish_clause): Likewise.
	* langhooks.h (lhd_omp_finish_clause): Likewise.
	* omp-low.c (scan_sharing_clauses): Keep GOMP_MAP_TO_PSET cause for
	'declare target' vars.

include/ChangeLog:

	PR fortran/96668
	* gomp-constants.h (GOMP_MAP_ALWAYS_POINTER_P): Define.

libgomp/ChangeLog:

	PR fortran/96668
	* libgomp.h (struct target_var_desc): Add has_null_ptr_assoc member.
	* target.c (gomp_map_vars_existing): Add always_to_flag flag.
	(gomp_map_vars_existing): Update call to it.
	(gomp_map_fields_existing): Likewise
	(gomp_map_vars_internal): Update PSET handling such that if a nullptr is
	now allocated or if GOMP_MAP_POINTER is used PSET is updated and pointer
	remapped.
	(GOMP_target_enter_exit_data): Hanlde GOMP_MAP_ALWAYS_POINTER like
	GOMP_MAP_POINTER.
	* testsuite/libgomp.fortran/map-alloc-ptr-1.f90: New test.
	* testsuite/libgomp.fortran/map-alloc-ptr-2.f90: New test.
2020-09-15 09:24:47 +02:00
Tobias Burnus
656218ab98 Fortran: Fix OpenMP's 'if(simd:' etc. conditions
gcc/fortran/ChangeLog:

	* openmp.c (gfc_match_omp_clauses): Re-order 'if' clause pasing
	to avoid creating spurious symbols.

libgomp/ChangeLog:

	* testsuite/libgomp.fortran/lastprivate-conditional-10.f90: New test.
2020-08-20 13:33:40 +02:00
Jakub Jelinek
676b5525e8 openmp: Handle clauses with gimple sequences in convert_nonlocal_omp_clauses properly
If the walk_body on the various sequences of reduction, lastprivate and/or linear
clauses needs to create a temporary variable, we should declare that variable
in that sequence rather than outside, where it would need to be privatized inside of
the construct.

2020-08-08  Jakub Jelinek  <jakub@redhat.com>

	PR fortran/93553
	* tree-nested.c (convert_nonlocal_omp_clauses): For
	OMP_CLAUSE_REDUCTION, OMP_CLAUSE_LASTPRIVATE and OMP_CLAUSE_LINEAR
	save info->new_local_var_chain around walks of the clause gimple
	sequences and declare_vars if needed into the sequence.

2020-08-08  Tobias Burnus  <tobias@codesourcery.com>

	PR fortran/93553
	* testsuite/libgomp.fortran/pr93553.f90: New test.
2020-08-08 11:10:30 +02:00
Tobias Burnus
ade6e7204c critical-hint-*.{c,f90}: Move from gcc/testsuite to libgomp/testsuite
libgomp/ChangeLog:

	* testsuite/libgomp.c-c++-common/critical-hint-1.c: New; moved from
	gcc/testsuite/c-c++-common/gomp/.
	* testsuite/libgomp.c-c++-common/critical-hint-2.c: Likewise.
	* testsuite/libgomp.fortran/critical-hint-1.f90: New; moved
	from gcc/testsuite/gfortran.dg/gomp/.
	* testsuite/libgomp.fortran/critical-hint-2.f90: Likewise.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/critical-hint-1.c: Moved to libgomp/.
	* c-c++-common/gomp/critical-hint-2.c: Moved to libgomp/.
	* gfortran.dg/gomp/critical-hint-1.f90: Moved to libgomp/.
	* gfortran.dg/gomp/critical-hint-2.f90: Moved to libgomp/.
2020-07-22 12:14:22 +02:00
Tobias Burnus
51542d9254 libgomp.fortran/alloc-1.F90: Fix testcase for 32bit size_t
libgomp/ChangeLog:

	* testsuite/libgomp.fortran/alloc-1.F90: Use c_size_t to
	avoid conversion on 32bit systems from 32bit to 64bit due
	to -fdefault-integer-8.
2020-07-15 17:23:04 +02:00
Tobias Burnus
e0685fadb6 libgomp.fortran/struct-elem-map-1.f90: Add char kind=4 tests
As the Fortran PR 95837 has been fixed, the test could be be added.

libgomp/ChangeLog:

	* testsuite/libgomp.fortran/struct-elem-map-1.f90: Remove unused
	variables; add character(kind=4) tests; update TODO comment.
2020-07-15 12:34:03 +02:00
Tobias Burnus
fff15bad1a libgomp: Add Fortran routine support for allocators
libgomp/ChangeLog:

	* allocator.c: Add ialias for omp_init_allocator and
	omp_destroy_allocator.
	* configure.ac: Set INTPTR_T_KIND.
	* configure: Regenerate.
	* Makefile.in: Regenerate.
	* testsuite/Makefile.in: Regenerate.
	* fortran.c (omp_init_allocator_, omp_destroy_allocator_,
	omp_set_default_allocator_, omp_get_default_allocator_): New
	functions and ialias_redirect.
	* icv.c: Add ialias for omp_set_default_allocator and
	omp_get_default_allocator.
	* libgomp.map (OMP_5.0.1): Add omp_init_allocator_,
	omp_destroy_allocator_, omp_set_default_allocator_ and
	omp_get_default_allocator_.
	* omp_lib.f90.in: Add allocator traits parameters, declare
	allocator routines and add related kind parameters.
	* omp_lib.h.in: Likewise.
	* testsuite/libgomp.c-c++-common/alloc-2.c: Fix sizeof.
	* testsuite/libgomp.fortran/alloc-1.F90: New test.
	* testsuite/libgomp.fortran/alloc-2.F90: New test.
	* testsuite/libgomp.fortran/alloc-3.F: New test.
	* testsuite/libgomp.fortran/alloc-4.f90: New test.
	* testsuite/libgomp.fortran/alloc-5.f90: New test.
2020-07-15 08:33:20 +02:00
Tobias Burnus
102502e32e [OpenMP, Fortran] Add structure/derived-type element mapping
gcc/fortran/ChangeLog:

	* openmp.c (gfc_match_omp_clauses): Match also derived-type
	component refs in OMP_CLAUSE_MAP.
	(resolve_omp_clauses): Resolve those.
	* trans-openmp.c (gfc_trans_omp_array_section, gfc_trans_omp_clauses):
	Handle OpenMP structure-element mapping.
	(gfc_trans_oacc_construct, gfc_trans_oacc_executable_directive,
	(gfc_trans_oacc_combined_directive, gfc_trans_oacc_declare): Update
	add openacc=true in gfc_trans_omp_clauses call.

gcc/testsuite/ChangeLog:

	* gfortran.dg/goacc/finalize-1.f: Update dump scan pattern.
	* gfortran.dg/gomp/map-1.f90: Update dg-error.
	* gfortran.dg/gomp/map-2.f90: New test.

libgomp/ChangeLog:

	* testsuite/libgomp.fortran/struct-elem-map-1.f90: New test.
2020-07-14 13:39:46 +02:00
Tobias Burnus
174e79bf73 [Fortran, OpenMP] Fix allocatable-components check (PR67311)
gcc/fortran/ChangeLog:

	PR fortran/67311
	* trans-openmp.c (gfc_has_alloc_comps): Return false also for
	pointers to arrays.

libgomp/ChangeLog:

	PR fortran/67311
	* testsuite/libgomp.fortran/target-map-1.f90: New test.
2020-07-14 12:55:53 +02:00
Thomas Schwinge
5864930754 Add 'dg-do run' to 'libgomp.fortran/use_device_ptr-optional-3.f90' [PR94848]
Fix-up for r279858/commit f760c0c77f "Fortran]
OpenMP/OpenACC – fix more issues with OPTIONAL".

With offloading enabled, we then saw:

    PASS: libgomp.fortran/use_device_ptr-optional-3.f90   -O0  (test for excess errors)
    PASS: libgomp.fortran/use_device_ptr-optional-3.f90   -O0  execution test
    PASS: libgomp.fortran/use_device_ptr-optional-3.f90   -O1  (test for excess errors)
    PASS: libgomp.fortran/use_device_ptr-optional-3.f90   -O1  execution test
    FAIL: libgomp.fortran/use_device_ptr-optional-3.f90   -O2  (test for excess errors)
    UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90   -O2  compilation failed to produce executable
    FAIL: libgomp.fortran/use_device_ptr-optional-3.f90   -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions  (test for excess errors)
    UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90   -O3 -fomit-frame-pointer -funroll-loops -fpeel-loops -ftracer -finline-functions  compilation failed to produce executable
    FAIL: libgomp.fortran/use_device_ptr-optional-3.f90   -O3 -g  (test for excess errors)
    UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90   -O3 -g  compilation failed to produce executable
    FAIL: libgomp.fortran/use_device_ptr-optional-3.f90   -Os  (test for excess errors)
    UNRESOLVED: libgomp.fortran/use_device_ptr-optional-3.f90   -Os  compilation failed to produce executable

 ... due to:

    /tmp/cciVc43I.o:(.gnu.offload_vars+0x10): undefined reference to `A.12.4064'
    [...]

..., but after the recent PR94848, PR95551 changes, that problem is now gone.

	libgomp/
	PR lto/94848
	* testsuite/libgomp.fortran/use_device_ptr-optional-3.f90: Add
	'dg-do run'.
2020-06-18 00:14:46 +02:00
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
Thomas Schwinge
bd194a51d4 Add missing results check in 'libgomp.fortran/allocatable3.f90'
libgomp/
	* testsuite/libgomp.fortran/allocatable3.f90: Add missing results
	check.

From-SVN: r272449
2019-06-19 00:14:43 +02:00
Thomas Schwinge
4017da8d1c [PR90743] Fortran 'allocatable' with OpenACC data/OpenMP 'target' 'map' clauses
Test what OpenMP 5.0 has to say on this topic.  And, do the same for OpenACC.

	libgomp/
	PR fortran/90743
	* oacc-parallel.c (GOACC_parallel_keyed): Handle NULL mapping
	case.
	* testsuite/libgomp.fortran/target-allocatable-1-1.f90: New file.
	* testsuite/libgomp.fortran/target-allocatable-1-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/allocatable-1-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/allocatable-1-2.f90: Likewise.

From-SVN: r272447
2019-06-19 00:14:24 +02:00
Jakub Jelinek
211b7533bf re PR middle-end/90779 (Fortran array initialization in offload regions)
PR middle-end/90779
	* gimplify.c: Include omp-offload.h and context.h.
	(gimplify_bind_expr): Add "omp declare target" attributes
	to static block scope variables inside of target region or target
	functions.

	* c-c++-common/goacc/routine-5.c (func2): Don't expect error for
	static block scope variable in #pragma acc routine.

	* testsuite/libgomp.c/pr90779.c: New test.
	* testsuite/libgomp.fortran/pr90779.f90: New test.

From-SVN: r272322
2019-06-15 09:09:04 +02:00