Commit Graph

558 Commits

Author SHA1 Message Date
Thomas Schwinge
b03d721a62 [libgomp] In OpenACC testing, by default only build for the offload target that we're actually going to test
... to avoid compilation overhead, and to keep simple '-foffload=[...]'
handling in test cases.

	libgomp/
	* testsuite/libgomp.oacc-c++/c++.exp: Specify
	"-foffload=$offload_target".
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
	* testsuite/lib/libgomp.exp
	(check_effective_target_openacc_nvidia_accel_configured): Remove,
	as (conceptually) merged into
	check_effective_target_openacc_nvidia_accel_selected.  Adjust all
	users.

From-SVN: r269109
2019-02-22 11:51:35 +01:00
Thomas Schwinge
0a0384b43a [libgomp] In OpenACC testing, cycle though all offload targets
... instead of through offload plugins.

	libgomp/
	* plugin/configfrag.ac: Populate and AC_SUBST offload_targets.
	* testsuite/libgomp-test-support.exp.in: Adjust.
	* testsuite/lib/libgomp.exp: Likewise.  Don't populate
	openacc_device_types_s.
	(offload_target_to_openacc_device_type): New proc.
	* testsuite/libgomp.oacc-c++/c++.exp: Adjust.
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
	* Makefile.in: Regenerate.
	* configure: Likewise.
	* testsuite/Makefile.in: Likewise.

From-SVN: r269108
2019-02-22 11:51:20 +01:00
Thomas Schwinge
ee332b4a9a [libgomp] Clarify difference between offload target, offload plugin, and OpenACC device type
libgomp/
	* plugin/configfrag.ac: Populate and AC_SUBST offload_plugins
	instead of offload_targets, and AC_DEFINE_UNQUOTED OFFLOAD_PLUGINS
	instead of OFFLOAD_TARGETS.
	* target.c (gomp_target_init): Adjust.
	* testsuite/libgomp-test-support.exp.in: Likewise.
	* testsuite/lib/libgomp.exp: Likewise.  Populate
	openacc_device_types_s instead of offload_targets_s_openacc.
	(check_effective_target_openacc_nvidia_accel_selected)
	(check_effective_target_openacc_host_selected): Adjust.
	* testsuite/libgomp.oacc-c++/c++.exp: Likewise.
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
	* Makefile.in: Regenerate.
	* config.h.in: Likewise.
	* configure: Likewise.
	* testsuite/Makefile.in: Likewise.

From-SVN: r269107
2019-02-22 11:51:05 +01:00
Thomas Schwinge
1241136c71 [libgomp] In OpenACC offloading testing, be more explicit in what is supported, and what is not, or why not
libgomp/
	* testsuite/lib/libgomp.exp: Error out for unknown offload target.
	* testsuite/libgomp.oacc-c++/c++.exp: Likewise.  Report if
	"offloading: supported, but hardware not accessible".
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.

From-SVN: r269106
2019-02-22 11:50:49 +01:00
Jakub Jelinek
8b44f8ec4b re PR c++/88988 (ICE: Segmentation fault (in lookup_name_real_1))
PR c++/88988
	* lambda.c (is_capture_proxy): Don't return true for
	DECL_OMP_PRIVATIZED_MEMBER artificial vars.

	* testsuite/libgomp.c++/pr88988.C: New test.

From-SVN: r268407
2019-01-31 00:28:53 +01:00
Jakub Jelinek
52bfbb69e7 re PR middle-end/89002 (ICE in scan_omp_1_op, at omp-low.c:3166)
PR middle-end/89002
	* gimplify.c (gimplify_omp_for): When adding OMP_CLAUSE_*_GIMPLE_SEQ
	for lastprivate/linear IV, push gimplify context around gimplify_assign
	and, if it needed any temporaries, pop it into a gimple bind around the
	sequence.

	* testsuite/libgomp.c/pr89002.c: New test.

From-SVN: r268346
2019-01-28 23:34:32 +01:00
Jakub Jelinek
be3a87e7b5 re PR middle-end/89002 (ICE in scan_omp_1_op, at omp-low.c:3166)
PR middle-end/89002
	* gimplify.c (gimplify_omp_for): When adding OMP_CLAUSE_*_GIMPLE_SEQ
	for lastprivate/linear IV, push gimplify context around gimplify_assign
	and, if it needed any temporaries, pop it into a gimple bind around the
	sequence.

	* testsuite/libgomp.c/pr89002.c: New test.

From-SVN: r268345
2019-01-28 23:33:33 +01:00
Richard Biener
497ef4d7f5 re PR testsuite/89064 (libgomp.graphite/force-parallel-5.c fails starting with r268257)
2019-01-28  Richard Biener  <rguenther@suse.de>

	PR testsuite/89064
	PR tree-optimization/86865
	* testsuite/libgomp.graphite/force-parallel-5.c: XFAIL.

From-SVN: r268333
2019-01-28 09:07:30 +00:00
Tom de Vries
4a75460b00 [nvptx, libgomp] Fix cuMemAlloc with size zero
Consider test-case:
...
int
main (void)
{
  #pragma acc parallel async
  ;
  #pragma acc parallel async
  ;
  #pragma acc wait

  return 0;
}
...

This fails with:
...
libgomp: cuMemAlloc error: invalid argument
Segmentation fault (core dumped)
...
The cuMemAlloc error is due to the fact that we're try to allocate 0 bytes.

Fix this by preventing calling map_push with size zero argument in nvptx_exec.

This also has the consequence that for the abort-1.c test-case, we end up
calling cuMemFree during map_fini for the struct cuda_map allocated in
map_init, which fails because an abort happened.  Fix this by calling
cuMemFree with CUDA_CALL_NOCHECK in cuda_map_destroy.

2019-01-23  Tom de Vries  <tdevries@suse.de>

	PR target/PR88946
	* plugin/plugin-nvptx.c (cuda_map_destroy): Use CUDA_CALL_NOCHECK for
	cuMemFree.
	(nvptx_exec): Don't call map_push if mapnum == 0.
	* testsuite/libgomp.oacc-c-c++-common/pr88946.c: New test.

From-SVN: r268178
2019-01-23 08:16:56 +00:00
Tom de Vries
4fef8e4d8c [nvptx, libgomp] Fix assert (!s->map->active) in map_fini
There are currently two situations where this assert triggers:
...
libgomp/plugin/plugin-nvptx.c: map_fini: Assertion `!s->map->active' failed.
...

First, in abort-1.c, a parallel region triggering an abort:
...
int
main (void)
{
  #pragma acc parallel
  abort ();

  return 0;
}
...

The abort is detected in nvptx_exec as the CUDA_ERROR_ILLEGAL_INSTRUCTION
return status of the cuStreamSynchronize call after kernel launch, which is
then handled by calling non-returning function GOMP_PLUGIN_fatal.
Consequently, the map_pop in nvptx_exec that in case of cuStreamSynchronize
success would remove or inactive the element added by the map_push earlier in
nvptx_exec, does not trigger.  With the element no longer active, but still
marked active and a member of s->map,  we run into the assert during
GOMP_OFFLOAD_fini_device, which is triggered from atexit handler
gomp_target_fini (which is triggered by the GOMP_PLUGIN_fatal mentioned above
calling exit).

Second, in pr88941.c, an async parallel region without wait:
...
int
main (void)
{
  #pragma acc parallel async
  ;

  /* no #pragma acc wait */
  return 0;
}
...

Because nvptx_exec is handling an async region, it does not call map_pop for
the element added by map_push, but schedules an kernel execution completion
event to call map_pop.  Again, we run into the assert during
GOMP_OFFLOAD_fini_device, which is triggered from atexit handler
gomp_target_fini, but the exit in this case is triggered by returning from main.
So either the kernel is still running, or the kernel has completed but the
corresponding event that is supposed to call map_pop is stuck in the event
queue, waiting for an event_gc.

Fix this by removing the assert, and skipping the freeing of device memory if
the map is still marked active (though in the async case, this is more a
workaround than an fix).

2019-01-23  Tom de Vries  <tdevries@suse.de>

	PR target/88941
	PR target/88939
	* plugin/plugin-nvptx.c (cuda_map_destroy): Handle map->active case.
	(map_fini): Remove "assert (!s->map->active)".
	* testsuite/libgomp.oacc-c-c++-common/pr88941.c: New test.

From-SVN: r268177
2019-01-23 08:16:42 +00:00
Tom de Vries
2ee6cb22c1 [nvptx, libgomp] Fix map_push
The map field of a struct ptx_stream is a FIFO.  The FIFO is implemented as a
single linked list, with pop-from-the-front semantics.

The function map_pop pops an element, either by:
- deallocating the element, if there is more than one element
- or marking the element inactive, if there's only one element

The responsibility of map_push is to push an element to the back, as well as
selecting the element to push, by:
- allocating an element, or
- reusing the element at the front if inactive and big enough, or
- dropping the element at the front if inactive and not big enough, and
  allocating one that's big enough

The current implemention gets at least the first and most basic scenario wrong:

> map = cuda_map_create (size);

We create an element, and assign it to map.

> for (t = s->map; t->next != NULL; t = t->next)
>   ;

We determine the last element in the fifo.

> t->next = map;

We append the new element.

> s->map = map;

But here, we throw away the rest of the FIFO, and declare the FIFO to be just
the new element.

This problem causes the test-case asyncwait-1.c to fail intermittently on some
systems.  The pr87835.c test-case added here is a a minimized and modified
version of asyncwait-1.c (avoiding the kernel construct) that is more likely to
fail.

Fix this by rewriting map_pop more robustly, by:
- seperating the function in two phases: select element, push element
- when reusing or dropping an element, making sure that the element is cleanly
  popped from the queue
- rewriting the push element part in such a way that it can handle all cases
  without needing if statements, such that each line is exercised for each of
  the three cases.

2019-01-23  Tom de Vries  <tdevries@suse.de>

	PR target/87835
	* plugin/plugin-nvptx.c (map_push): Fix adding of allocated element.
	* testsuite/libgomp.oacc-c-c++-common/pr87835.c: New test.

From-SVN: r268176
2019-01-23 08:16:11 +00:00
Tom de Vries
d41d952c9b [nvptx] Handle assignment to gang-level reduction variable
2019-01-15  Tom de Vries  <tdevries@suse.de>

	PR target/80547
	* config/nvptx/nvptx.c (nvptx_goacc_reduction_init): Handle
	lhs == NULL_TREE for gang-level reduction.

	* testsuite/libgomp.oacc-c-c++-common/gang-reduction-var-assignment.c:
	New test.

From-SVN: r267934
2019-01-15 10:11:16 +00:00
Tom de Vries
efb56ae82b [nvptx] Enable setting vector length using -fopenacc-dim -- testcases
Add some test-cases that set vector length using -fopenacc-dim.

2019-01-12  Tom de Vries  <tdevries@suse.de>

	* testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c: New test.
	* testsuite/libgomp.oacc-fortran/gemm-2.f90: New test.

From-SVN: r267897
2019-01-12 22:19:31 +00:00
Tom de Vries
a105775825 [nvptx] Add vector_length 64 test-cases
Add some test-cases using vector_length 64.

2019-01-12  Tom de Vries  <tdevries@suse.de>

	* testsuite/libgomp.oacc-c-c++-common/vector-length-64-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/vector-length-64-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/vector-length-64-3.c: New test.

From-SVN: r267895
2019-01-12 22:19:02 +00:00
Tom de Vries
56314b772f [nvptx] Force vl32 if calling vector-partitionable routines -- test-cases
Add test-cases for "[nvptx] Force vl32 if calling vector-partitionable
routines".

2019-01-12  Tom de Vries  <tdevries@suse.de>

	PR target/85486
	* testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/pr85486.c: New test.

From-SVN: r267894
2019-01-12 22:18:50 +00:00
Tom de Vries
b39e4366a2 [nvptx] Don't emit barriers for empty loops -- test-cases
Add test-cases for PR85381.

2019-01-12  Tom de Vries  <tdevries@suse.de>

	PR target/85381
	* testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/pr85381.c: New test.

From-SVN: r267893
2019-01-12 22:18:39 +00:00
Tom de Vries
2cb7a501ab [nvptx] Enable large vectors -- reduction testcases
Add various reduction test-cases with vector length 128.

2019-01-12  Tom de Vries  <tdevries@suse.de>

	* testsuite/libgomp.oacc-c-c++-common/vred2d-128.c: New test.
	* testsuite/libgomp.oacc-fortran/gemm.f90: New test.
	* testsuite/libgomp.oacc-c-c++-common/vector-length-128-10.c: New test.

From-SVN: r267892
2019-01-12 22:18:27 +00:00
Tom de Vries
8e77f71eda [nvptx] Enable large vectors -- test-cases
Add various test-cases with vector length 128.

2019-01-12  Tom de Vries  <tdevries@suse.de>

	* testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c: New test.

From-SVN: r267891
2019-01-12 22:18:11 +00:00
Tom de Vries
2b9d9e3937 [nvptx] Enable large vectors
Allow vector_length clauses to accept values larger than warp size.  Note that
this does not enable setting vector_length to values larger than warp size using
-fopenacc-dim.

2019-01-12  Tom de Vries  <tdevries@suse.de>

	* config/nvptx/nvptx.c (nvptx_goacc_validate_dims): Take larger vector
	lengths into account.

	* testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Expect
	vector length to be 128.
	* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Expect vector
	length 2097152 to be reduced to 1024 instead of 32.

From-SVN: r267889
2019-01-12 22:17:42 +00:00
Tom de Vries
9390f91687 [libgomp, testsuite, openacc] Remove -foffload=-w in reduction-[1-5].c
Before the commit "[libgomp, testsuite, openacc] Don't use const int for
dimensions", the "const int" construct was used to set launch dimensions in
reductions-[1-5].c.  In the case of -xc -O0, the const int is implemented as a
variable by the C front-end.  Consequently, the nvptx back-end generated
warnings that vector_length was overridden to be hard-coded, rather than left to
be set at runtime.  The test-cases silenced these warnings by switching off all
warnings in the accelerator compiler using "-foffload=-w".

Given that no warnings occur anymore, remove the "-foffload=-w" setting.

2019-01-11  Tom de Vries  <tdevries@suse.de>

	* testsuite/libgomp.oacc-c-c++-common/reduction-1.c: Remove
	-foffload=-w.
	* testsuite/libgomp.oacc-c-c++-common/reduction-2.c: Same.
	* testsuite/libgomp.oacc-c-c++-common/reduction-3.c: Same.
	* testsuite/libgomp.oacc-c-c++-common/reduction-4.c: Same.
	* testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Same.

From-SVN: r267836
2019-01-11 11:46:06 +00:00
Tom de Vries
2c3e7ad20b [nvptx, testsuite, openacc, libgomp] Add insufficient-resources.c
Add a test-case that tests the "insufficient resources" fatal in the nvptx
libgomp plugin.

2019-01-11  Tom de Vries  <tdevries@suse.de>

	* testsuite/libgomp.oacc-c-c++-common/insufficient-resources.c: New
	test.

From-SVN: r267835
2019-01-11 11:45:55 +00:00
Nathan Sidwell
e222497dcb Add testcase from PR71959
libgomp/

	PR lto/71959
	* testsuite/libgomp.oacc-c++/pr71959-aux.cc: New.
	* testsuite/libgomp.oacc-c++/pr71959.C: New.

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

From-SVN: r267806
2019-01-10 12:32:03 +00:00
Tom de Vries
5d0bc70ae4 [libgomp, testsuite, openacc] Don't use const int for dimensions
Const int is handled differently at -O0 for -xc and -xc++, which can cause noise
in testsuite/libgomp.oacc-c-c++-common test-cases (which are both run for c and
c++) if const int is used for launch dimensions.

Fix this by using #defines instead.

2019-01-09  Tom de Vries  <tdevries@suse.de>

	PR target/88756
	* testsuite/libgomp.oacc-c-c++-common/reduction-1.c (ng, nw, vl): Use
	#define instead of "const int".
	* testsuite/libgomp.oacc-c-c++-common/reduction-2.c (ng, nw, vl): Same.
	* testsuite/libgomp.oacc-c-c++-common/reduction-3.c (ng, nw, vl): Same.
	* testsuite/libgomp.oacc-c-c++-common/reduction-4.c (ng, nw, vl): Same.
	* testsuite/libgomp.oacc-c-c++-common/reduction-5.c (ng, nw, vl): Same.

From-SVN: r267747
2019-01-09 00:07:55 +00:00
Tom de Vries
43493c97a6 [nvptx] Fix libgomp.oacc-c-c++-common/vector-length-128-3.c
The vector-length-128-3.c test-case uses GOMP_OPENACC_DIM=-:-:128, but '-' is
not yet supported on trunk.  Use GOMP_OPENACC_DIM=::128 instead.

2019-01-07  Tom de Vries  <tdevries@suse.de>

	* testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: Fix
	GOMP_OPENACC_DIM argument.

From-SVN: r267624
2019-01-07 08:09:49 +00:00
Tom de Vries
5c571497e1 [nvptx] Add vector_length 128 testcases
Add a couple of test-cases using vector length 128, while checking that we
override to vector length 32.

2019-01-03  Tom de Vries  <tdevries@suse.de>

	* testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: New test.

From-SVN: r267559
2019-01-03 15:08:46 +00:00
Jakub Jelinek
a554497024 Update copyright years.
From-SVN: r267494
2019-01-01 13:31:55 +01:00
Steven G. Kargl
b3c64ed0b0 aligned1.f03: Fix invalid code that now causes an error after r267415.
2018-12-27  Steven G. Kargl  <kargl@gcc.gnu.org>

	* libgomp.fortran/aligned1.f03: Fix invalid code that now causes
	an error after r267415.

From-SVN: r267436
2018-12-27 20:57:12 +00:00
Tom de Vries
a152954ea4 [nvptx] Commit passing pr85381-*.c test-cases
Add pr85381*.c test-cases that are already passing without the fix for PR85381.

Build and reg-tested on x86_64 with nvptx accelerator.

2018-12-19  Tom de Vries  <tdevries@suse.de>

	* testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: New test.

From-SVN: r267268
2018-12-19 14:20:54 +00:00
Tom de Vries
49188cd1f2 [nvptx, libgomp] Move rtl-dump test-cases to libgomp
The goacc.exp test-cases nvptx-merged-loop.c and nvptx-sese-1.c are failing
during linking due to missing libgomp.spec.

Move them to the libgomp testsuite.

Build and reg-tested on x86_64 with nvptx accelerator.

2018-12-19  Tom de Vries  <tdevries@suse.de>

	* gcc.dg/goacc/nvptx-merged-loop.c: Move to
	libgomp/testsuite/libgomp.oacc-c-c++-common.
	* gcc.dg/goacc/nvptx-sese-1.c: Same.

	* testsuite/lib/libgomp.exp: Add load_lib of scanoffloadrtl.exp.
	* testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c: Move from
	gcc/testsuite/gcc.dg/goacc.
	* testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c: Same.

From-SVN: r267267
2018-12-19 14:20:44 +00:00
Thomas Schwinge
c759830b29 Missing changes from "Adjust copy/copyin/copyout/create for OpenACC 2.5"
Most of that patch's changes were already committed as part of r261813 "Update
OpenACC data clause semantics to the 2.5 behavior", but not all of them.

	libgomp/
	* oacc-mem.c (acc_present_or_create): Remove definition and change
	to alias of acc_create.
	(acc_present_or_copyin): Remove definition and change to alias of
	acc_copyin.
	* oacc-parallel.c (GOACC_enter_exit_data): Call acc_create instead
	of acc_present_or_create.
	* testsuite/libgomp.oacc-c-c++-common/data-already-1.c: Remove.
	* testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise.
	* testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.

Co-Authored-By: Chung-Lin Tang <cltang@codesourcery.com>

From-SVN: r267153
2018-12-14 21:43:12 +01:00
Thomas Schwinge
f847198ec3 [PR88495] An OpenACC async queue is always synchronized with itself
An OpenACC async queue is always synchronized with itself, so invocations like
"#pragma acc wait(0) async(0)", or "acc_wait_async (0, 0)" don't make a lot of
sense, but are still valid.

	libgomp/
	PR libgomp/88495
	* plugin/plugin-nvptx.c (nvptx_wait_async): Don't refuse
	"identical parameters".
	* testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: Update.
	* testsuite/libgomp.oacc-c-c++-common/lib-80.c: Remove.

From-SVN: r267152
2018-12-14 21:43:02 +01:00
Thomas Schwinge
c8ab8aab9f [PR88484] OpenACC wait directive without wait argument but with async clause
We don't correctly handle "#pragma acc wait async (a)" for "a >= 0", handling
as a no-op whereas it should enqueue the appropriate wait operations on
"async (a)".

	libgomp/
	PR libgomp/88484
	* oacc-parallel.c (GOACC_wait): Correct handling for "async >= 0".
	* testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: New file.

From-SVN: r267151
2018-12-14 21:42:50 +01:00
Thomas Schwinge
1404af62dc [PR88407] [OpenACC] Correctly handle unseen async-arguments
... which turn the operation into a no-op.

	libgomp/
	PR libgomp/88407
	* plugin/plugin-nvptx.c (nvptx_async_test, nvptx_wait)
	(nvptx_wait_async): Unseen async-argument is a no-op.
	* testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Update.
	* testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-79.c: Likewise.
	* testsuite/libgomp.oacc-fortran/lib-12.f90: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-71.c: Merge into...
	* testsuite/libgomp.oacc-c-c++-common/lib-69.c: ... this.  Update.
	* testsuite/libgomp.oacc-c-c++-common/lib-77.c: Merge into...
	* testsuite/libgomp.oacc-c-c++-common/lib-74.c: ... this.  Update

From-SVN: r267150
2018-12-14 21:42:40 +01:00
Thomas Schwinge
7de562eec2 Revise libgomp.oacc-c-c++-common/data-2-lib.c, libgomp.oacc-c-c++-common/data-2.c
These are meant to be functionally equivalent (but no longer are), just using
different means.  Also, use the OpenACC "*_async" functions recently added.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Revise.
	* testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise.

From-SVN: r267149
2018-12-14 21:42:29 +01:00
Chung-Lin Tang
17469af75b Correctly describe OpenACC async/wait dependencies
libgomp/
	* testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Adjust.
	* testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise.

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

From-SVN: r267148
2018-12-14 21:42:18 +01:00
Thomas Schwinge
18c247cc0b [PR88370] acc_get_cuda_stream/acc_set_cuda_stream: acc_async_sync, acc_async_noval
Per my reading of the OpenACC specification (and as supported by secondary
documentation, such as code examples, or presentations), it's valid to call
"acc_get_cuda_stream"/"acc_set_cuda_stream" also with "acc_async_sync",
"acc_async_noval" arguments, not just with the nonnegative values as currently
implemented.

	libgomp/
	PR libgomp/88370
	* libgomp.texi (acc_get_current_cuda_context, acc_get_cuda_stream)
	(acc_set_cuda_stream): Clarify.
	* oacc-cuda.c (acc_get_cuda_stream, acc_set_cuda_stream): Use
	"async_valid_p".
	* plugin/plugin-nvptx.c (nvptx_set_cuda_stream): Refuse "async ==
	acc_async_sync".
	* testsuite/libgomp.oacc-c-c++-common/acc_set_cuda_stream-1.c: New file.
	* testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-84.c: Update.
	* testsuite/libgomp.oacc-c-c++-common/lib-85.c: Likewise.

From-SVN: r267147
2018-12-14 21:42:08 +01:00
Tom de Vries
b0aba46ca6 [offloading] Error on missing symbols
When compiling an OpenMP or OpenACC program containing a reference in the
offloaded code to a symbol that has not been included in the offloaded code,
the offloading compiler may ICE in lto1.

Fix this by erroring out instead, mentioning the problematic symbol:
...
error: variable 'var' has been referenced in offloaded code but hasn't
  been marked to be included in the offloaded code
lto1: fatal error: errors during merging of translation units
compilation terminated.
...

Build x86_64 with nvptx accelerator and reg-tested libgomp.

Build x86_64 and reg-tested libgomp.

2018-12-14  Tom de Vries  <tdevries@suse.de>

	* lto-cgraph.c (verify_node_partition): New function.
	(input_overwrite_node, input_varpool_node): Use verify_node_partition.

	* testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c: New test.
	* testsuite/libgomp.c-c++-common/function-not-offloaded.c: New test.
	* testsuite/libgomp.c-c++-common/variable-not-offloaded.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c: New test.

From-SVN: r267134
2018-12-14 13:48:56 +00:00
Jakub Jelinek
a51f8c92e2 re PR libgomp/88460 ([nvptx] FAIL: libgomp.c++/for-24.C (internal compiler error))
PR libgomp/88460
	* testsuite/libgomp.c++/for-24.C (results): Include it in
	omp declare target region.
	(main): Use map (always, tofrom: results) instead of
	map (tofrom: results).

From-SVN: r267093
2018-12-13 13:53:19 +01:00
Jakub Jelinek
7a289b7d0a re PR fortran/88463 (Rejects conforming source, OpenMP Parallel region Default(None) reference to module parameter array, separate source)
PR fortran/88463
	* trans-openmp.c (gfc_omp_predetermined_sharing): Handle TREE_READONLY
	VAR_DECLs with DECL_EXTERNAL like those with TREE_STATIC.

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

From-SVN: r267069
2018-12-12 23:49:35 +01:00
Jakub Jelinek
a6ef2ac9d9 omp-builtins.def (BUILT_IN_GOMP_LOOP_NONMONOTONIC_RUNTIME_START, [...]): Fix up function types - remove one argument.
* omp-builtins.def (BUILT_IN_GOMP_LOOP_NONMONOTONIC_RUNTIME_START,
	BUILT_IN_GOMP_LOOP_MAYBE_NONMONOTONIC_RUNTIME_START,
	BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_RUNTIME_START,
	BUILT_IN_GOMP_LOOP_ULL_MAYBE_NONMONOTONIC_RUNTIME_START,
	BUILT_IN_GOMP_PARALLEL_LOOP_NONMONOTONIC_RUNTIME,
	BUILT_IN_GOMP_PARALLEL_LOOP_MAYBE_NONMONOTONIC_RUNTIME): Fix up
	function types - remove one argument.

	* testsuite/libgomp.c-c++-common/for-16.c: New test.

From-SVN: r267067
2018-12-12 23:47:55 +01:00
Thomas Koenig
d0cbb206e2 re PR libfortran/88411 (Random crashes for ASYNCHRONOUS writes (bad locking?))
2018-12-09  Thomas Koenig  <tkoenig@gcc.gnu.org>

	PR fortran/88411
	* io/transfer.c (dta_transfer_init): Do not treat as an
	asynchronous statement unless the statement has
	ASYNCHRONOUS="YES".
	(st_write_done): Likewise.
	(st_read_done): Do not perform async_wait for synchronous I/O
	on an async unit.
	(st_read_done): Likewise.

2018-12-09  Thomas Koenig  <tkoenig@gcc.gnu.org>

	PR fortran/88411
	* testsuite/libgomp.fortran/async_io_8.f90: New test.

From-SVN: r266929
2018-12-09 18:54:47 +00:00
Jakub Jelinek
6997628d35 re PR libgomp/87995 (libgomp.c/../libgomp.c-c++-common/cancel-taskgroup-3.c fails consistently after r265930)
PR libgomp/87995
	* testsuite/libgomp.c-c++-common/cancel-taskgroup-3.c: Require
	tls_runtime effective target.
	(t): New threadprivate variable.
	(main): Set t in threads which execute iterations of the worksharing
	loop.  Propagate that to the task after the loop and don't abort
	if the current taskgroup hasn't been cancelled.

From-SVN: r266904
2018-12-08 09:58:24 +01:00
Jakub Jelinek
4a82df9a38 tree-nested.c (convert_nonlocal_omp_clauses, [...]): Handle OMP_CLAUSE_IN_REDUCTION...
* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Handle OMP_CLAUSE_IN_REDUCTION,
	OMP_CLAUSE_TASK_REDUCTION and OMP_CLAUSE__SIMT_ clauses.
	(convert_nonlocal_reference_stmt, convert_local_reference_stmt):
	Convert clauses for GIMPLE_OMP_TASKGROUP.

	* testsuite/libgomp.c/task-reduction-3.c: New test.

From-SVN: r266723
2018-12-02 13:50:50 +01:00
Jakub Jelinek
daa8c1d763 omp-low.c (check_omp_nesting_restrictions): Allow cancel or cancellation point with taskgroup clause inside of taskloop.
* omp-low.c (check_omp_nesting_restrictions): Allow cancel or
	cancellation point with taskgroup clause inside of taskloop.  Consider
	a taskloop construct without nogroup clause as implicit taskgroup for
	diagnostics if cancel/cancellation point with taskgroup clause is
	closely nested inside of taskgroup region.

	* c-c++-common/gomp/cancel-1.c (f2): Add various taskloop related
	tests.

	* testsuite/libgomp.c-c++-common/cancel-taskgroup-4.c: New test.

From-SVN: r266722
2018-12-02 13:48:42 +01:00
Cesar Philippidis
fe570ff8d4 [PR88288, OpenACC, libgomp] Adjust offsets for present data clauses
Make libgomp respect the on device offset of subarrays which may arise in
present data clauses.

	libgomp/
	PR libgomp/88288
	* oacc-parallel.c (GOACC_parallel_keyed): Add offset to devaddrs.
	* testsuite/libgomp.oacc-c-c++-common/pr88288.c: New test.

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

From-SVN: r266688
2018-11-30 21:39:49 +01:00
Thomas Schwinge
c223608f54 Add libgomp.oacc-fortran/lib-16-2.f90
This is a copy of libgomp.oacc-fortran/lib-16.f90, but does 'include
"openacc_lib.h"' instead of 'use openacc'.

	libgomp/
	* testsuite/libgomp.oacc-fortran/lib-16-2.f90: New file.

From-SVN: r266683
2018-11-30 21:38:57 +01:00
Richard Biener
fc60283c5c re PR tree-optimization/88182 (ICE in vectorizable_reduction, at tree-vect-loop.c:6465)
2018-10-19  Richard Biener  <rguenther@suse.de>

	PR tree-optimization/88182
	* g++.dg/gomp/pr88182.C: Move from libgomp and use -fopenmp-simd.

	* testsuite/libgomp.c++/pr88182.C: Move to g++.dg/gomp.

From-SVN: r266545
2018-11-28 08:29:16 +00:00
Jakub Jelinek
ffcf3b79c2 Makefile.am (AUTOMAKE_OPTIONS): Drop dejagnu.
* testsuite/Makefile.am (AUTOMAKE_OPTIONS): Drop dejagnu.
	(RUNTEST): Don't define.
	(RUNTESTDEFAULTFLAGS): Add.
	(check-DEJAGNU, site.exp, distclean-DEJAGNU): New goals.
	(distclean-am): Depend on distclean-DEJAGNU.
	(check-am): If -j% option is present in MFLAGS and if
	`getconf _NPROCESSORS_ONLN` is more than 8, export OMP_NUM_THREADS=8.
	(.PHONY): Add check-DEJAGNU and distclean-DEJAGNU.
	* testsuite/Makefile.in: Regenerated.

From-SVN: r266482
2018-11-26 22:03:41 +01:00
Richard Biener
ef6e6914c8 re PR tree-optimization/88182 (ICE in vectorizable_reduction, at tree-vect-loop.c:6465)
2018-11-26  Richard Biener  <rguenther@suse.de>

	PR tree-optimization/88182
	* tree-vect-loop.c (vectorizable_reduction): Pick up single
	correct reduc_def_info.
	* tree-vect-slp.c (vect_analyze_slp_instance): Set
	STMT_VINFO_REDUC_DEF of the first stmt.

	libgomp/
	* testsuite/libgomp.c++/pr88182.C: New testcase.

From-SVN: r266467
2018-11-26 15:37:35 +00:00
Jakub Jelinek
207286c314 workshare-reduction-1.c: New test.
2018-11-09  Jakub Jelinek  <jakub@redhat.com>

	* gcc.dg/gomp/workshare-reduction-1.c: New test.
	* gcc.dg/gomp/workshare-reduction-2.c: New test.
	* gcc.dg/gomp/workshare-reduction-3.c: New test.
	* gcc.dg/gomp/workshare-reduction-4.c: New test.
	* gcc.dg/gomp/workshare-reduction-5.c: New test.
	* gcc.dg/gomp/workshare-reduction-6.c: New test.
	* gcc.dg/gomp/workshare-reduction-7.c: New test.
	* gcc.dg/gomp/workshare-reduction-8.c: New test.
	* gcc.dg/gomp/workshare-reduction-9.c: New test.
	* gcc.dg/gomp/workshare-reduction-10.c: New test.
	* gcc.dg/gomp/workshare-reduction-11.c: New test.
	* gcc.dg/gomp/workshare-reduction-12.c: New test.
	* gcc.dg/gomp/workshare-reduction-13.c: New test.
	* gcc.dg/gomp/workshare-reduction-14.c: New test.
	* gcc.dg/gomp/workshare-reduction-15.c: New test.
	* gcc.dg/gomp/workshare-reduction-16.c: New test.
	* gcc.dg/gomp/workshare-reduction-17.c: New test.
	* gcc.dg/gomp/workshare-reduction-18.c: New test.
	* gcc.dg/gomp/workshare-reduction-19.c: New test.
	* gcc.dg/gomp/workshare-reduction-20.c: New test.
	* gcc.dg/gomp/workshare-reduction-21.c: New test.
	* gcc.dg/gomp/workshare-reduction-22.c: New test.
	* gcc.dg/gomp/workshare-reduction-23.c: New test.
	* gcc.dg/gomp/workshare-reduction-24.c: New test.
	* gcc.dg/gomp/workshare-reduction-25.c: New test.
	* gcc.dg/gomp/workshare-reduction-26.c: New test.
	* gcc.dg/gomp/workshare-reduction-27.c: New test.
	* gcc.dg/gomp/workshare-reduction-28.c: New test.
	* gcc.dg/gomp/workshare-reduction-29.c: New test.
	* gcc.dg/gomp/workshare-reduction-30.c: New test.
	* gcc.dg/gomp/workshare-reduction-31.c: New test.
	* gcc.dg/gomp/workshare-reduction-32.c: New test.
	* gcc.dg/gomp/workshare-reduction-33.c: New test.
	* gcc.dg/gomp/workshare-reduction-34.c: New test.
	* gcc.dg/gomp/workshare-reduction-35.c: New test.
	* gcc.dg/gomp/workshare-reduction-36.c: New test.
	* gcc.dg/gomp/workshare-reduction-37.c: New test.
	* gcc.dg/gomp/workshare-reduction-38.c: New test.
	* gcc.dg/gomp/workshare-reduction-39.c: New test.
	* gcc.dg/gomp/workshare-reduction-40.c: New test.
	* gcc.dg/gomp/workshare-reduction-41.c: New test.
	* gcc.dg/gomp/workshare-reduction-42.c: New test.
	* gcc.dg/gomp/workshare-reduction-43.c: New test.
	* gcc.dg/gomp/workshare-reduction-44.c: New test.
	* gcc.dg/gomp/workshare-reduction-45.c: New test.
	* gcc.dg/gomp/workshare-reduction-46.c: New test.
	* gcc.dg/gomp/workshare-reduction-47.c: New test.
	* gcc.dg/gomp/workshare-reduction-48.c: New test.
	* gcc.dg/gomp/workshare-reduction-49.c: New test.
	* gcc.dg/gomp/workshare-reduction-50.c: New test.
	* gcc.dg/gomp/workshare-reduction-51.c: New test.
	* gcc.dg/gomp/workshare-reduction-52.c: New test.
	* gcc.dg/gomp/workshare-reduction-53.c: New test.
	* gcc.dg/gomp/workshare-reduction-54.c: New test.
	* gcc.dg/gomp/workshare-reduction-55.c: New test.
	* gcc.dg/gomp/workshare-reduction-56.c: New test.
	* gcc.dg/gomp/workshare-reduction-57.c: New test.
	* gcc.dg/gomp/workshare-reduction-58.c: New test.
libgomp/
	* testsuite/libgomp.c-c++-common/task-reduction-13.c: New test.
	* testsuite/libgomp.c-c++-common/task-reduction-14.c: New test.

From-SVN: r265967
2018-11-09 14:02:50 +01:00