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
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
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
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
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
When compiling the test-case with r268755, we run into a SIGSEGV in
oacc_entry_exit_ok_1 when trying to dereference a NULL red:
...
struct reduction_info *red;
red = reduction_phi (reduction_list, use_stmt);
tree val = PHI_RESULT (red->keep_res);
...
Fix this by handling ref == NULL.
Bootstrapped and reg-tested on x86_64.
Build and reg-tested on x86_64 with nvptx accelerator.
2019-06-16 Tom de Vries <tdevries@suse.de>
PR tree-optimization/89376
* tree-parloops.c (oacc_entry_exit_ok_1): Handle red == NULL.
* testsuite/libgomp.oacc-c-c++-common/pr89376.c: New test.
From-SVN: r272338
After the fix for "PR tree-optimization/89713 - Assume loop with an exit is
finite" ( r272234 ) empty oacc loops are removed before expand.
Update pr85381-{2,4}.c accordingly.
2019-06-15 Tom de Vries <tdevries@suse.de>
PR tree-optimization/89713
* testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Expect no bar.sync.
* testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Same.
From-SVN: r272324
If we compile the openacc testcase with -fopenacc -O2, we run into a SIGSEGV
or assert. The root cause for this is that pass_thread_jumps breaks the
invariant that OACC_FORK and OACC_JOIN mark the start and end of a
single-entry-single-exit region.
Fix this by bailing out when encountering an IFN_UNIQUE in
thread_jumps::profitable_jump_thread_path.
Bootstrapped and reg-tested on x86_64.
Build and reg-tested libgomp on x86_64 with nvptx accelerator.
2019-06-15 Tom de Vries <tdevries@suse.de>
PR tree-optimization/90009
* tree-ssa-threadbackward.c (thread_jumps::profitable_jump_thread_path):
Return NULL if bb contains IFN_UNIQUE.
* testsuite/libgomp.oacc-c-c++-common/pr90009.c: New test.
From-SVN: r272321
... 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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
