When using a compiler build with:
...
+#define PTX_DEFAULT_VECTOR_LENGTH PTX_CTA_SIZE
...
consider a test-case:
...
int
main (void)
{
#pragma acc parallel vector_length (64)
#pragma acc loop worker
for (unsigned int i = 0; i < 32; i++)
#pragma acc loop vector
for (unsigned int j = 0; j < 64; j++)
;
return 0;
}
...
If num_workers is 16, either because:
- we add a "num_workers (16)" clause on the parallel directive, or
- we set "GOMP_OPENACC_DIM=:16:", or
- the libgomp plugin chooses 16 num_workers
we run into an illegal instruction at runtime, because a bar.sync instruction
tries to use a barrier 16. The instruction is illegal, because ptx supports
only 16 barriers per CTA, and the valid range is 0..15.
The problem is that with a warp-multiple vector length, we use a code generation
scheme with a per-worker barrier. And because barrier zero is reserved for
per-cta barrier, only the remaining 15 barriers can be used as per-worker
barrier, and consequently we can't use num_workers larger than 15.
This problem occurs only for vector_length 64. For vector_length 32, we use a
different code generation scheme, and for vector_length >= 96, the maximum
num_workers is not big enough not to trigger this problem.
Also, this problem only occurs for num_workers 16. As explained above,
num_workers 15 is safe to use, and 16 is already the maximum num_workers for
vector_length 64.
This patch fixes the problem in both the compiler (handling "num_workers (16)")
and in the libgomp nvptx plugin (with and without "GOMP_OPENACC_DIM=:16:").
2019-01-11 Tom de Vries <tdevries@suse.de>
* config/nvptx/nvptx.c (PTX_CTA_NUM_BARRIERS, PTX_PER_CTA_BARRIER)
(PTX_NUM_PER_CTA_BARRIER, PTX_FIRST_PER_WORKER_BARRIER)
(PTX_NUM_PER_WORKER_BARRIERS): Define.
(nvptx_apply_dim_limits): Prevent vector_length 64 and
num_workers 16.
* plugin/plugin-nvptx.c (nvptx_exec): Prevent vector_length 64 and
num_workers 16.
From-SVN: r267838
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
libgomp/
* config/rtems/affinity-fmt.c: New file. Include affinity-fmt.c,
undefining HAVE_GETPID and HAVE_GETHOSTNAME, and mapping fwrite to
write.
From-SVN: r267751
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
When using a compiler build with:
...
+#define PTX_DEFAULT_VECTOR_LENGTH PTX_CTA_SIZE
+#define PTX_MAX_VECTOR_LENGTH PTX_CTA_SIZE
...
and running the libgomp testsuite, we run into an execution failure in
parallel-loop-1.c, due to a cuda launch failure:
...
nvptx_exec: kernel f6_none_none$_omp_fn$0: launch gangs=480, workers=0, \
vectors=1024
libgomp: cuLaunchKernel error: invalid argument
...
because workers == 0.
The workers variable is set to 0 here in nvptx_exec:
...
workers = blocks / actual_vectors;
...
because actual_vectors is 1024, and blocks is 768:
...
cuOccupancyMaxPotentialBlockSize: grid = 10, block = 768
...
Fix this by ensuring that workers is at least one.
2019-01-09 Tom de Vries <tdevries@suse.de>
* plugin/plugin-nvptx.c (nvptx_exec): Make sure to launch with at least
one worker.
From-SVN: r267746
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
... so that we're then able to use this for other flags in addition to
"GOACC_FLAG_HOST_FALLBACK".
gcc/
* omp-expand.c (expand_omp_target): Restructure OpenACC vs. OpenMP
code paths. Update for libgomp OpenACC entry points change.
include/
* gomp-constants.h (GOACC_FLAG_HOST_FALLBACK)
(GOACC_FLAGS_MARSHAL_OP, GOACC_FLAGS_UNMARSHAL): Define.
libgomp/
* oacc-parallel.c (GOACC_parallel_keyed, GOACC_parallel)
(GOACC_data_start, GOACC_enter_exit_data, GOACC_update)
(GOACC_declare): Redefine the "device" argument to "flags".
From-SVN: r267448
libgomp/
* target.c (struct gomp_coalesce_chunk): New structure.
(struct gomp_coalesce_buf): Update the chunks member to use that
type. Adjust all users.
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
From-SVN: r267446
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
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
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
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
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
Libgomp test-case libgomp.c/target-5.c is failing to compile when building for
x86_64 with nvptx accelerator due to missing:
- getpid
- gethostname
- isatty (pulled in by fwrite)
in the nvptx newlib.
This patch fixes the build failure by:
- adding a function gomp_print_string which limits the use of fwrite to a single
location (in affinity-fmt.c), and
- creating an nvptx version of affinity-fmt.c, which:
- overrides the configure test results HAVE_GETPID and HAVE_GETHOSTNAME, and
- implements fwrite using write.
Build and reg-tested on x86_64 with nvptx accelerator.
2018-12-13 Tom de Vries <tdevries@suse.de>
* affinity-fmt.c (gomp_print_string): New function, factored out of ...
(omp_display_affinity, gomp_display_affinity_thread): ... here, and ...
* fortran.c (omp_display_affinity_): ... here.
* libgomp.h (gomp_print_string): Declare.
* config/nvptx/affinity-fmt.c: New file. Include affinity-fmt.c,
undefining HAVE_GETPID and HAVE_GETHOSTNAME, and mapping fwrite to
write.
From-SVN: r267100
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
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-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
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
* 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
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
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-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
* 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 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
* affinity-fmt.c: Include inttypes.h if HAVE_INTTYPES_H.
(gomp_display_affinity): Use __builtin_choose_expr to handle
properly handle argument having integral, or pointer or some other
type. If inttypes.h is available and PRIx64 is defined, use PRIx64
with uint64_t type instead of %llx and unsigned long long.
From-SVN: r265985
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
* builtin-types.def (BT_FN_VOID_BOOL, BT_FN_VOID_SIZE_SIZE_PTR,
BT_FN_UINT_UINT_PTR_PTR, BT_FN_UINT_OMPFN_PTR_UINT_UINT,
BT_FN_BOOL_UINT_LONGPTR_LONG_LONG_LONGPTR_LONGPTR_PTR_PTR,
BT_FN_BOOL_UINT_ULLPTR_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR,
BT_FN_BOOL_LONG_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR_PTR_PTR,
BT_FN_BOOL_BOOL_ULL_ULL_ULL_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR): New.
* gengtype.c (open_base_files): Add omp-general.h.
* gimple.c (gimple_build_omp_critical):
(gimple_build_omp_taskgroup): Add CLAUSES argument. Call
gimple_omp_taskgroup_set_clauses.
(gimple_build_omp_atomic_load): Add mo argument, call
gimple_omp_atomic_set_memory_order.
(gimple_build_omp_atomic_store): Likewise.
(gimple_copy): Adjust handling of GIMPLE_OMP_TASKGROUP.
* gimple.def (GIMPLE_OMP_TASKGROUP): Use GSS_OMP_SINGLE_LAYOUT
instead of GSS_OMP.
(GIMPLE_OMP_TEAMS): Use GSS_OMP_PARALLEL_LAYOUT instead
of GSS_OMP_SINGLE_LAYOUT, adjust comments.
* gimple.h (enum gf_mask): Add GF_OMP_TEAMS_HOST, GF_OMP_TASK_TASKWAIT
and GF_OMP_ATOMIC_MEMORY_ORDER. Remove GF_OMP_ATOMIC_SEQ_CST, use
different value for GF_OMP_ATOMIC_NEED_VALUE.
(struct gimple_statement_omp_taskreg): Add GIMPLE_OMP_TEAMS to
comments.
(struct gimple_statement_omp_single_layout): And remove here.
(struct gomp_teams): Inherit from gimple_statement_omp_taskreg rather
than gimple_statement_omp_single_layout.
(is_a_helper <gimple_statement_omp_taskreg *>::test): Allow
GIMPLE_OMP_TEAMS.
(is_a_helper <const gimple_statement_omp_taskreg *>::test): Likewise.
(gimple_omp_subcode): Formatting fix.
(gimple_omp_teams_child_fn, gimple_omp_teams_child_fn_ptr,
gimple_omp_teams_set_child_fn, gimple_omp_teams_data_arg,
gimple_omp_teams_data_arg_ptr, gimple_omp_teams_set_data_arg,
gimple_omp_teams_host, gimple_omp_teams_set_host,
gimple_omp_task_taskwait_p, gimple_omp_task_set_taskwait_p,
gimple_omp_taskgroup_clauses, gimple_omp_taskgroup_clauses_ptr,
gimple_omp_taskgroup_set_clauses): New inline functions.
(gimple_build_omp_atomic_load): Add enum omp_memory_order argument.
(gimple_build_omp_atomic_store): Likewise.
(gimple_omp_atomic_seq_cst_p): Remove.
(gimple_omp_atomic_memory_order): New function.
(gimple_omp_atomic_set_seq_cst): Remove.
(gimple_omp_atomic_set_memory_order): New function.
(gimple_build_omp_taskgroup): Add clauses argument.
* gimple-pretty-print.c (dump_gimple_omp_taskgroup): New function.
(dump_gimple_omp_task): Print taskwait with depend clauses.
(dump_gimple_omp_atomic_load, dump_gimple_omp_atomic_store): Use
dump_omp_atomic_memory_order.
(pp_gimple_stmt_1): Handle GIMPLE_OMP_TASKGROUP.
* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_ALLOC_ONLY,
GOVD_MAP_FROM_ONLY and GOVD_NONTEMPORAL.
(enum omp_region_type): Reserve bits 1 and 2 for auxiliary flags,
renumber values of most of ORT_* enumerators, add ORT_HOST_TEAMS,
ORT_COMBINED_HOST_TEAMS, ORT_TASKGROUP, ORT_TASKLOOP and
ORT_UNTIED_TASKLOOP enumerators.
(enum gimplify_defaultmap_kind): New.
(struct gimplify_omp_ctx): Remove target_map_scalars_firstprivate and
target_map_pointers_as_0len_arrays members, add defaultmap.
(new_omp_context): Initialize defaultmap member.
(gimple_add_tmp_var): Handle ORT_TASKGROUP like ORT_WORKSHARE.
(maybe_fold_stmt): Don't fold even in host teams regions.
(omp_firstprivatize_variable): Handle ORT_TASKGROUP like
ORT_WORKSHARE. Test ctx->defaultmap[GDMK_SCALAR] instead of
ctx->omp_firstprivatize_variable.
(omp_add_variable): Don't add private/firstprivate for VLAs in
ORT_TASKGROUP.
(omp_default_clause): Print "taskloop" rather than "task" if
ORT_*TASKLOOP.
(omp_notice_variable): Handle ORT_TASKGROUP like ORT_WORKSHARE.
Handle new defaultmap clause kinds.
(omp_is_private): Handle ORT_TASKGROUP like ORT_WORKSHARE. Allow simd
iterator to be lastprivate or private. Fix up diagnostics if linear
is used on collapse>1 simd iterator.
(omp_check_private): Handle ORT_TASKGROUP like ORT_WORKSHARE.
(gimplify_omp_depend): New function.
(gimplify_scan_omp_clauses): Add shared clause on parallel for
combined parallel master taskloop{, simd} if taskloop has
firstprivate, lastprivate or reduction clause. Handle
OMP_CLAUSE_REDUCTION_TASK diagnostics. Adjust tests for
ORT_COMBINED_TEAMS. Gimplify depend clauses with iterators. Handle
cancel and simd OMP_CLAUSE_IF_MODIFIERs. Handle
OMP_CLAUSE_NONTEMPORAL. Handle new defaultmap clause kinds. Handle
OMP_CLAUSE_{TASK,IN}_REDUCTION. Diagnose invalid conditional
lastprivate.
(gimplify_adjust_omp_clauses_1): Ignore GOVD_NONTEMPORAL. Handle
GOVD_MAP_ALLOC_ONLY and GOVD_MAP_FROM_ONLY.
(gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_NONTEMPORAL. Handle
OMP_CLAUSE_{TASK,IN}_REDUCTION.
(gimplify_omp_task): Handle taskwait with depend clauses.
(gimplify_omp_for): Add shared clause on parallel for combined
parallel master taskloop{, simd} if taskloop has firstprivate,
lastprivate or reduction clause. Use ORT_TASKLOOP or
ORT_UNTIED_TASKLOOP instead of ORT_TASK or ORT_UNTIED_TASK. Adjust
tests for ORT_COMBINED_TEAMS. Handle C++ range for loops with
NULL TREE_PURPOSE in OMP_FOR_ORIG_DECLS. Firstprivatize
__for_end and __for_range temporaries on OMP_PARALLEL for
distribute parallel for{, simd}. Move OMP_CLAUSE_REDUCTION
and OMP_CLAUSE_IN_REDUCTION from taskloop to the task construct
sandwiched in between two taskloops.
(computable_teams_clause): Test ctx->defaultmap[GDMK_SCALAR]
instead of ctx->omp_firstprivatize_variable.
(gimplify_omp_workshare): Set ort to ORT_HOST_TEAMS or
ORT_COMBINED_HOST_TEAMS if not inside of target construct. If
host teams, use gimplify_and_return_first etc. for body like
for target or target data constructs, and at the end call
gimple_omp_teams_set_host on the GIMPLE_OMP_TEAMS object.
(gimplify_omp_atomic): Use OMP_ATOMIC_MEMORY_ORDER instead
of OMP_ATOMIC_SEQ_CST, pass it as new argument to
gimple_build_omp_atomic_load and gimple_build_omp_atomic_store, remove
gimple_omp_atomic_set_seq_cst calls.
(gimplify_expr) <case OMP_TASKGROUP>: Move handling into a separate
case, handle taskgroup clauses.
* lto-streamer-out.c (hash_tree): Handle
OMP_CLAUSE_{TASK,IN}_REDUCTION.
* Makefile.in (GTFILES): Add omp-general.h.
* omp-builtins.def (BUILT_IN_GOMP_TASKWAIT_DEPEND,
BUILT_IN_GOMP_LOOP_NONMONOTONIC_RUNTIME_START,
BUILT_IN_GOMP_LOOP_MAYBE_NONMONOTONIC_RUNTIME_START,
BUILT_IN_GOMP_LOOP_START, BUILT_IN_GOMP_LOOP_ORDERED_START,
BUILT_IN_GOMP_LOOP_DOACROSS_START,
BUILT_IN_GOMP_LOOP_NONMONOTONIC_RUNTIME_NEXT,
BUILT_IN_GOMP_LOOP_MAYBE_NONMONOTONIC_RUNTIME_NEXT,
BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_RUNTIME_START,
BUILT_IN_GOMP_LOOP_ULL_MAYBE_NONMONOTONIC_RUNTIME_START,
BUILT_IN_GOMP_LOOP_ULL_START, BUILT_IN_GOMP_LOOP_ULL_ORDERED_START,
BUILT_IN_GOMP_LOOP_ULL_DOACROSS_START,
BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_RUNTIME_NEXT,
BUILT_IN_GOMP_LOOP_ULL_MAYBE_NONMONOTONIC_RUNTIME_NEXT,
BUILT_IN_GOMP_PARALLEL_LOOP_NONMONOTONIC_RUNTIME,
BUILT_IN_GOMP_PARALLEL_LOOP_MAYBE_NONMONOTONIC_RUNTIME,
BUILT_IN_GOMP_PARALLEL_REDUCTIONS, BUILT_IN_GOMP_SECTIONS2_START,
BUILT_IN_GOMP_TEAMS_REG, BUILT_IN_GOMP_TASKGROUP_REDUCTION_REGISTER,
BUILT_IN_GOMP_TASKGROUP_REDUCTION_UNREGISTER,
BUILT_IN_GOMP_TASK_REDUCTION_REMAP,
BUILT_IN_GOMP_WORKSHARE_TASK_REDUCTION_UNREGISTER): New builtins.
* omp-expand.c (workshare_safe_to_combine_p): Return false for
non-worksharing loops.
(omp_adjust_chunk_size): Don't adjust anything if chunk_size is zero.
(determine_parallel_type): Don't combine parallel with worksharing
which has _reductemp_ clause.
(expand_parallel_call): Emit the GOMP_*nonmonotonic_runtime* or
GOMP_*maybe_nonmonotonic_runtime* builtins instead of GOMP_*runtime*
if there is nonmonotonic modifier or if there is no modifier and no
ordered clause. For dynamic and guided schedule without monotonic
and nonmonotonic modifier, default to nonmonotonic.
(expand_omp_for): Likewise. Adjust expand_omp_for_generic caller, use
GOMP_loop{,_ull}{,_ordered,_doacross}_start builtins if there are
task reductions.
(expand_task_call): Add GOMP_TASK_FLAG_REDUCTION flag to flags if
there are any reduction clauses.
(expand_taskwait_call): New function.
(expand_teams_call): New function.
(expand_omp_taskreg): Allow GIMPLE_OMP_TEAMS and call
expand_teams_call for it. Formatting fix. Handle taskwait with
depend clauses.
(expand_omp_for_generic): Add SCHED_ARG argument. Handle expansion
of worksharing loops with task reductions.
(expand_omp_for_static_nochunk, expand_omp_for_static_chunk): Handle
expansion of worksharing loops with task reductions.
(expand_omp_sections): Handle expansion of sections with task
reductions.
(expand_omp_synch): For host teams call expand_omp_taskreg.
(omp_memory_order_to_memmodel): New function.
(expand_omp_atomic_load, expand_omp_atomic_store,
expand_omp_atomic_fetch_op): Use it and gimple_omp_atomic_memory_order
instead of gimple_omp_atomic_seq_cst_p.
(build_omp_regions_1, omp_make_gimple_edges): Treat taskwait with
depend clauses as a standalone directive.
* omp-general.c (enum omp_requires): New variable.
(omp_extract_for_data): Initialize have_reductemp member. Allow
NE_EXPR even in OpenMP loops, transform them into LT_EXPR or
GT_EXPR loops depending on incr sign. Formatting fixes.
* omp-general.h (struct omp_for_data): Add have_reductemp member.
(enum omp_requires): New enum.
(omp_requires_mask): Declare.
* omp-grid.c (grid_eliminate_combined_simd_part): Formatting fix.
Fix comment typos.
* omp-low.c (struct omp_context): Add task_reductions and
task_reduction_map fields.
(is_host_teams_ctx): New function.
(is_taskreg_ctx): Return true also if is_host_teams_ctx.
(use_pointer_for_field): Use is_global_var instead of
TREE_STATIC || DECL_EXTERNAL, and apply only if not privatized
in outer contexts.
(build_outer_var_ref): Ignore taskgroup outer contexts.
(delete_omp_context): Release task_reductions and task_reduction_map.
(scan_sharing_clauses): Don't add any fields for reduction clause on
taskloop. Handle OMP_CLAUSE__REDUCTEMP_. Handle
OMP_CLAUSE_{IN,TASK}_REDUCTION and OMP_CLAUSE_REDUCTION with task
modifier. Don't ignore shared clauses in is_host_teams_ctx contexts.
Handle OMP_CLAUSE_NONTEMPORAL.
(add_taskreg_looptemp_clauses): Add OMP_CLAUSE__REDUCTEMP_ clause if
needed.
(scan_omp_parallel): Add _reductemp_ clause if there are any reduction
clauses with task modifier.
(scan_omp_task): Handle taskwait with depend clauses.
(finish_taskreg_scan): Move field corresponding to _reductemp_ clause
first. Move also OMP_CLAUSE__REDUCTEMP_ clause in front if present.
Handle GIMPLE_OMP_TEAMS like GIMPLE_OMP_PARALLEL.
(scan_omp_for): Fix comment formatting.
(scan_omp_teams): Handle host teams constructs.
(check_omp_nesting_restrictions): Allow teams with no outer
OpenMP context. Adjust diagnostics for teams strictly nested into
some explicit OpenMP construct other than target. Allow OpenMP atomics
inside of simd regions.
(scan_omp_1_stmt): Call scan_sharing_clauses for taskgroups.
(scan_omp_1_stmt) <case GIMPLE_OMP_TEAMS>: Temporarily bump
taskreg_nesting_level while scanning host teams construct.
(task_reduction_read): New function.
(lower_rec_input_clauses): Handle OMP_CLAUSE_REDUCTION on taskloop
construct. Handle OMP_CLAUSE_IN_REDUCTION and OMP_CLAUSE__REDUCTEMP_
clauses. Handle OMP_CLAUSE_REDUCTION with task modifier. Remove
second argument create_tmp_var if it is NULL. Don't ignore shared
clauses in is_host_teams_ctx contexts. Handle
OMP_CLAUSE_FIRSTPRIVATE_NO_REFERENCE on OMP_CLAUSE_FIRSTPRIVATE
clauses.
(lower_reduction_clauses): Ignore reduction clauses with task
modifier. Remove second argument create_tmp_var if it is NULL.
Initialize OMP_ATOMIC_MEMORY_ORDER to relaxed.
(lower_send_clauses): Ignore reduction clauses with task modifier.
Handle OMP_CLAUSE__REDUCTEMP_. Don't send anything for
OMP_CLAUSE_REDUCTION on taskloop. Handle OMP_CLAUSE_IN_REDUCTION.
(maybe_add_implicit_barrier_cancel): Add OMP_RETURN argument, don't
rely that it is the last stmt in body so far. Ignore outer taskgroup
contexts.
(omp_task_reductions_find_first, omp_task_reduction_iterate,
lower_omp_task_reductions): New functions.
(lower_omp_sections): Handle reduction clauses with taskgroup
modifiers. Adjust maybe_add_implicit_barrier_cancel caller.
(lower_omp_single): Adjust maybe_add_implicit_barrier_cancel caller.
(lower_omp_for): Likewise. Handle reduction clauses with taskgroup
modifiers.
(lower_omp_taskgroup): Handle taskgroup reductions.
(create_task_copyfn): Copy over OMP_CLAUSE__REDUCTEMP_ pointer.
Handle OMP_CLAUSE_IN_REDUCTION and OMP_CLAUSE_REDUCTION clauses.
(lower_depend_clauses): If there are any
OMP_CLAUSE_DEPEND_DEPOBJ or OMP_CLAUSE_DEPEND_MUTEXINOUTSET
depend clauses, use a new array format. If OMP_CLAUSE_DEPEND_LAST is
seen, assume lowering is done already and return early. Set kind
on artificial depend clause to OMP_CLAUSE_DEPEND_LAST.
(lower_omp_taskreg): Handle reduction clauses with task modifier on
parallel construct. Handle reduction clause on taskloop construct.
Handle taskwait with depend clauses.
(lower_omp_1): Use lower_omp_taskreg instead of lower_omp_teams
for host teams constructs.
* tree.c (omp_clause_num_ops): Add in_reduction, task_reduction,
nontemporal and _reductemp_ clause entries.
(omp_clause_code_name): Likewise.
(walk_tree_1): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION,
OMP_CLAUSE_NONTEMPORAL and OMP_CLAUSE__REDUCTEMP_.
* tree-core.h (enum omp_clause_code): Add
OMP_CLAUSE_{{IN,TASK}_REDUCTION,NONTEMPORAL,_REDUCTEMP_}.
(enum omp_clause_defaultmap_kind, enum omp_memory_order): New.
(struct tree_base): Add omp_atomic_memory_order field into union.
Remove OMP_ATOMIC_SEQ_CST comment.
(enum omp_clause_depend_kind): Add OMP_CLAUSE_DEPEND_MUTEXINOUTSET
and OMP_CLAUSE_DEPEND_DEPOBJ.
(struct tree_omp_clause): Add subcode.defaultmap_kind.
* tree.def (OMP_TASKGROUP): Add another operand, move next to other
OpenMP constructs with body and clauses operands.
* tree.h (OMP_BODY): Use OMP_MASTER instead of OMP_TASKGROUP.
(OMP_CLAUSES): Use OMP_TASKGROUP instead of OMP_SINGLE.
(OMP_TASKGROUP_CLAUSES): Define.
(OMP_CLAUSE_DECL): Use OMP_CLAUSE__REDUCTEMP_ instead of
OMP_CLAUSE__LOOPTEMP_.
(OMP_ATOMIC_SEQ_CST): Remove.
(OMP_ATOMIC_MEMORY_ORDER, OMP_CLAUSE_FIRSTPRIVATE_NO_REFERENCE,
OMP_CLAUSE_LASTPRIVATE_CONDITIONAL): Define.
(OMP_CLAUSE_REDUCTION_CODE, OMP_CLAUSE_REDUCTION_INIT,
OMP_CLAUSE_REDUCTION_MERGE, OMP_CLAUSE_REDUCTION_PLACEHOLDER,
OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER,
OMP_CLAUSE_REDUCTION_OMP_ORIG_REF): Handle
OMP_CLAUSE_{,IN_,TASK_}REDUCTION.
(OMP_CLAUSE_REDUCTION_TASK, OMP_CLAUSE_REDUCTION_INSCAN,
OMP_CLAUSE_DEFAULTMAP_KIND, OMP_CLAUSE_DEFAULTMAP_CATEGORY,
OMP_CLAUSE_DEFAULTMAP_BEHAVIOR, OMP_CLAUSE_DEFAULTMAP_SET_KIND):
Define.
* tree-inline.c (remap_gimple_stmt): Remap taskgroup clauses.
* tree-nested.c (convert_nonlocal_omp_clauses): Handle
OMP_CLAUSE__REDUCTEMP_, OMP_CLAUSE_NONTEMPORAL.
(convert_local_omp_clauses): Likewise. Remove useless test.
* tree-parloops.c (create_call_for_reduction_1): Pass
OMP_MEMORY_ORDER_RELAXED as new argument to
dump_gimple_omp_atomic_load and dump_gimple_omp_atomic_store.
* tree-pretty-print.c (dump_omp_iterators): New function.
(dump_omp_clause): Handle OMP_CLAUSE__REDUCTEMP_,
OMP_CLAUSE_NONTEMPORAL, OMP_CLAUSE_{TASK,IN}_REDUCTION. Print
reduction modifiers. Handle OMP_CLAUSE_DEPEND_DEPOBJ and
OMP_CLAUSE_DEPEND_MUTEXINOUTSET. Print iterators in depend clauses.
Print __internal__ for OMP_CLAUSE_DEPEND_LAST. Handle cancel and
simd OMP_CLAUSE_IF_MODIFIERs. Handle new kinds of
OMP_CLAUSE_DEFAULTMAP. Print conditional: for
OMP_CLAUSE_LASTPRIVATE_CONDITIONAL.
(dump_omp_atomic_memory_order): New function.
(dump_generic_node): Use it. Print taskgroup clauses. Print
taskwait with depend clauses.
* tree-pretty-print.h (dump_omp_atomic_memory_order): Declare.
* tree-streamer-in.c (unpack_ts_omp_clause_value_fields):
Handle OMP_CLAUSE_{TASK,IN}_REDUCTION.
* tree-streamer-out.c (pack_ts_omp_clause_value_fields,
write_ts_omp_clause_tree_pointers): Likewise.
gcc/c-family/
* c-common.h (c_finish_omp_taskgroup): Add CLAUSES argument.
(c_finish_omp_atomic): Replace bool SEQ_CST argument with
enum omp_memory_order MEMORY_ORDER.
(c_finish_omp_flush): Add MO argument.
(c_omp_depend_t_p, c_finish_omp_depobj): Declare.
(c_finish_omp_for): Add FINAL_P argument.
* c-omp.c: Include memmodel.h.
(c_finish_omp_taskgroup): Add CLAUSES argument. Set
OMP_TASKGROUP_CLAUSES to it.
(c_finish_omp_atomic): Replace bool SEQ_CST argument with
enum omp_memory_order MEMORY_ORDER. Set OMP_ATOMIC_MEMORY_ORDER
instead of OMP_ATOMIC_SEQ_CST.
(c_omp_depend_t_p, c_finish_omp_depobj): New functions.
(c_finish_omp_flush): Add MO argument, if not MEMMODEL_LAST, emit
__atomic_thread_fence call with the given value.
(check_omp_for_incr_expr): Formatting fixes.
(c_finish_omp_for): Add FINAL_P argument. Allow NE_EXPR
even in OpenMP loops, diagnose if NE_EXPR and incr expression
is not constant expression 1 or -1. Transform NE_EXPR loops
with iterators pointers to VLA into LT_EXPR or GT_EXPR loops.
(c_omp_check_loop_iv_r): Look for orig decl of C++ range for
loops too.
(c_omp_split_clauses): Add support for combined
#pragma omp parallel master and
#pragma omp {,parallel }master taskloop{, simd} constructs.
Handle OMP_CLAUSE_IN_REDUCTION. Handle OMP_CLAUSE_REDUCTION_TASK.
Handle OMP_CLAUSE_NONTEMPORAL. Handle splitting OMP_CLAUSE_IF
also to OMP_SIMD. Copy OMP_CLAUSE_LASTPRIVATE_CONDITIONAL.
(c_omp_predetermined_sharing): Don't return
OMP_CLAUSE_DEFAULT_SHARED for const qualified decls.
* c-pragma.c (omp_pragmas): Add PRAGMA_OMP_DEPOBJ and
PRAGMA_OMP_REQUIRES.
* c-pragma.h (enum pragma_kind): Likewise.
(enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_NONTEMPORAL
and PRAGMA_OMP_CLAUSE_{IN,TASK}_REDUCTION.
gcc/c/
* c-parser.c: Include memmode.h.
(c_parser_omp_depobj, c_parser_omp_requires): New functions.
(c_parser_pragma): Handle PRAGMA_OMP_DEPOBJ and PRAGMA_OMP_REQUIRES.
(c_parser_omp_clause_name): Handle nontemporal, in_reduction and
task_reduction clauses.
(c_parser_omp_variable_list): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION.
For OMP_CLAUSE_DEPEND, parse clause operands as either an array
section, or lvalue assignment expression.
(c_parser_omp_clause_if): Handle cancel and simd modifiers.
(c_parser_omp_clause_lastprivate): Parse optional
conditional: modifier.
(c_parser_omp_clause_hint): Require constant integer expression rather
than just integer expression.
(c_parser_omp_clause_defaultmap): Parse new kinds of defaultmap
clause.
(c_parser_omp_clause_reduction): Add IS_OMP and KIND arguments.
Parse reduction modifiers. Pass KIND to c_parser_omp_variable_list.
(c_parser_omp_clause_nontemporal, c_parser_omp_iterators): New
functions.
(c_parser_omp_clause_depend): Parse iterator modifier and handle
iterators. Parse mutexinoutset and depobj kinds.
(c_parser_oacc_all_clauses): Adjust c_parser_omp_clause_reduction
callers.
(c_parser_omp_all_clauses): Likewise. Handle
PRAGMA_OMP_CLAUSE_NONTEMPORAL and
PRAGMA_OMP_CLAUSE_{IN,TASK}_REDUCTION.
(c_parser_omp_atomic): Parse hint and memory order clauses. Handle
default memory order from requires directive if any. Adjust
c_finish_omp_atomic caller.
(c_parser_omp_critical): Allow comma in between (name) and hint clause.
(c_parser_omp_flush): Parse flush with memory-order-clause.
(c_parser_omp_for_loop): Allow NE_EXPR even in
OpenMP loops, adjust c_finish_omp_for caller.
(OMP_SIMD_CLAUSE_MASK): Add if and nontemporal clauses.
(c_parser_omp_master): Add p_name, mask and cclauses arguments.
Allow to be called while parsing combined parallel master.
Parse combined master taskloop{, simd}.
(c_parser_omp_parallel): Parse combined
parallel master{, taskloop{, simd}} constructs.
(OMP_TASK_CLAUSE_MASK): Add in_reduction clause.
(OMP_TASKGROUP_CLAUSE_MASK): Define.
(c_parser_omp_taskgroup): Add LOC argument. Parse taskgroup clauses.
(OMP_TASKWAIT_CLAUSE_MASK): Define.
(c_parser_omp_taskwait): Handle taskwait with depend clauses.
(c_parser_omp_teams): Force a BIND_EXPR with BLOCK
around teams body. Use SET_EXPR_LOCATION.
(c_parser_omp_target_data): Allow target data
with only use_device_ptr clauses.
(c_parser_omp_target): Use SET_EXPR_LOCATION. Set
OMP_REQUIRES_TARGET_USED bit in omp_requires_mask.
(c_parser_omp_requires): New function.
(c_finish_taskloop_clauses): New function.
(OMP_TASKLOOP_CLAUSE_MASK): Add reduction and in_reduction clauses.
(c_parser_omp_taskloop): Use c_finish_taskloop_clauses. Add forward
declaration. Disallow in_reduction clause when combined with parallel
master.
(c_parser_omp_construct): Adjust c_parser_omp_master and
c_parser_omp_taskgroup callers.
* c-typeck.c (c_finish_omp_cancel): Diagnose if clause with modifier
other than cancel.
(handle_omp_array_sections_1): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION
like OMP_CLAUSE_REDUCTION.
(handle_omp_array_sections): Likewise. Call save_expr on array
reductions before calling build_index_type. Handle depend clauses
with iterators.
(struct c_find_omp_var_s): New type.
(c_find_omp_var_r, c_omp_finish_iterators): New functions.
(c_finish_omp_clauses): Don't diagnose nonmonotonic clause
with static, runtime or auto schedule kinds. Call save_expr for whole
array reduction sizes. Diagnose reductions with zero sized elements
or variable length structures. Diagnose nogroup clause used with
reduction clause(s). Handle depend clause with
OMP_CLAUSE_DEPEND_DEPOBJ. Diagnose bit-fields. Require
omp_depend_t type for OMP_CLAUSE_DEPEND_DEPOBJ kinds and
some different type for other kinds. Use build_unary_op with
ADDR_EXPR and build_indirect_ref instead of c_mark_addressable.
Handle depend clauses with iterators. Remove no longer needed special
case that predetermined const qualified vars may be specified in
firstprivate clause. Complain if const qualified vars are mentioned
in data-sharing clauses other than firstprivate or shared. Use
error_at with OMP_CLAUSE_LOCATION (c) as first argument instead of
error. Formatting fix. Handle OMP_CLAUSE_NONTEMPORAL and
OMP_CLAUSE_{IN,TASK}_REDUCTION. Allow any lvalue as
OMP_CLAUSE_DEPEND operand (besides array section), adjust diagnostics.
gcc/cp/
* constexpr.c (potential_constant_expression_1): Handle OMP_DEPOBJ.
* cp-gimplify.c (cp_genericize_r): Handle
OMP_CLAUSE_{IN,TASK}_REDUCTION.
(cxx_omp_predetermined_sharing_1): Don't return
OMP_CLAUSE_DEFAULT_SHARED for const qualified decls with no mutable
member. Return OMP_CLAUSE_DEFAULT_FIRSTPRIVATE for this pointer.
* cp-objcp-common.c (cp_common_init_ts): Handle OMP_DEPOBJ.
* cp-tree.def (OMP_DEPOBJ): New tree code.
* cp-tree.h (OMP_ATOMIC_DEPENDENT_P): Return true also for first
argument being OMP_CLAUSE.
(OMP_DEPOBJ_DEPOBJ, OMP_DEPOBJ_CLAUSES): Define.
(cp_convert_omp_range_for, cp_finish_omp_range_for): Declare.
(finish_omp_atomic): Add LOC, CLAUSES and MO arguments. Remove
SEQ_CST argument.
(finish_omp_for_block): Declare.
(finish_omp_flush): Add MO argument.
(finish_omp_depobj): Declare.
* cxx-pretty-print.c (cxx_pretty_printer::statement): Handle
OMP_DEPOBJ.
* dump.c (cp_dump_tree): Likewise.
* lex.c (cxx_init): Likewise.
* parser.c: Include memmodel.h.
(cp_parser_for): Pass false as new is_omp argument to
cp_parser_range_for.
(cp_parser_range_for): Add IS_OMP argument, return before finalizing
if it is true.
(cp_parser_omp_clause_name): Handle nontemporal, in_reduction and
task_reduction clauses.
(cp_parser_omp_var_list_no_open): Handle
OMP_CLAUSE_{IN,TASK}_REDUCTION. For OMP_CLAUSE_DEPEND, parse clause
operands as either an array section, or lvalue assignment expression.
(cp_parser_omp_clause_if): Handle cancel and simd modifiers.
(cp_parser_omp_clause_defaultmap): Parse new kinds of defaultmap
clause.
(cp_parser_omp_clause_reduction): Add IS_OMP and KIND arguments.
Parse reduction modifiers. Pass KIND to c_parser_omp_variable_list.
(cp_parser_omp_clause_lastprivate, cp_parser_omp_iterators): New
functions.
(cp_parser_omp_clause_depend): Parse iterator modifier and handle
iterators. Parse mutexinoutset and depobj kinds.
(cp_parser_oacc_all_clauses): Adjust cp_parser_omp_clause_reduction
callers.
(cp_parser_omp_all_clauses): Likewise. Handle
PRAGMA_OMP_CLAUSE_NONTEMPORAL and
PRAGMA_OMP_CLAUSE_{IN,TASK}_REDUCTION. Call
cp_parser_omp_clause_lastprivate for OpenMP lastprivate clause.
(cp_parser_omp_atomic): Pass pragma_tok->location as
LOC to finish_omp_atomic. Parse hint and memory order clauses.
Handle default memory order from requires directive if any. Adjust
finish_omp_atomic caller.
(cp_parser_omp_critical): Allow comma in between (name) and hint
clause.
(cp_parser_omp_depobj): New function.
(cp_parser_omp_flush): Parse flush with memory-order-clause.
(cp_parser_omp_for_cond): Allow NE_EXPR even in OpenMP loops.
(cp_convert_omp_range_for, cp_finish_omp_range_for): New functions.
(cp_parser_omp_for_loop): Parse C++11 range for loops among omp
loops. Handle OMP_CLAUSE_IN_REDUCTION like OMP_CLAUSE_REDUCTION.
(OMP_SIMD_CLAUSE_MASK): Add if and nontemporal clauses.
(cp_parser_omp_simd, cp_parser_omp_for): Call keep_next_level before
begin_omp_structured_block and call finish_omp_for_block on
finish_omp_structured_block result.
(cp_parser_omp_master): Add p_name, mask and cclauses arguments.
Allow to be called while parsing combined parallel master.
Parse combined master taskloop{, simd}.
(cp_parser_omp_parallel): Parse combined
parallel master{, taskloop{, simd}} constructs.
(cp_parser_omp_single): Use SET_EXPR_LOCATION.
(OMP_TASK_CLAUSE_MASK): Add in_reduction clause.
(OMP_TASKWAIT_CLAUSE_MASK): Define.
(cp_parser_omp_taskwait): Handle taskwait with depend clauses.
(OMP_TASKGROUP_CLAUSE_MASK): Define.
(cp_parser_omp_taskgroup): Parse taskgroup clauses, adjust
c_finish_omp_taskgroup caller.
(cp_parser_omp_distribute): Call keep_next_level before
begin_omp_structured_block and call finish_omp_for_block on
finish_omp_structured_block result.
(cp_parser_omp_teams): Force a BIND_EXPR with BLOCK around teams
body.
(cp_parser_omp_target_data): Allow target data with only
use_device_ptr clauses.
(cp_parser_omp_target): Set OMP_REQUIRES_TARGET_USED bit in
omp_requires_mask.
(cp_parser_omp_requires): New function.
(OMP_TASKLOOP_CLAUSE_MASK): Add reduction and in_reduction clauses.
(cp_parser_omp_taskloop): Add forward declaration. Disallow
in_reduction clause when combined with parallel master. Call
keep_next_level before begin_omp_structured_block and call
finish_omp_for_block on finish_omp_structured_block result.
(cp_parser_omp_construct): Adjust cp_parser_omp_master caller.
(cp_parser_pragma): Handle PRAGMA_OMP_DEPOBJ and PRAGMA_OMP_REQUIRES.
* pt.c (tsubst_omp_clause_decl): Add iterators_cache argument.
Adjust recursive calls. Handle iterators.
(tsubst_omp_clauses): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION and
OMP_CLAUSE_NONTEMPORAL. Adjust tsubst_omp_clause_decl callers.
(tsubst_decomp_names):
(tsubst_omp_for_iterator): Change orig_declv into a reference.
Handle range for loops. Move orig_declv handling after declv/initv
handling.
(tsubst_expr): Force a BIND_EXPR with BLOCK around teams body.
Adjust finish_omp_atomic caller. Call keep_next_level before
begin_omp_structured_block. Call cp_finish_omp_range_for for range
for loops and use {begin,finish}_omp_structured_block instead of
{push,pop}_stmt_list if there are any range for loops. Call
finish_omp_for_block on finish_omp_structured_block result.
Handle OMP_DEPOBJ. Handle taskwait with depend clauses. For
OMP_ATOMIC call tsubst_omp_clauses on clauses if any, adjust
finish_omp_atomic caller. Use OMP_ATOMIC_MEMORY_ORDER rather
than OMP_ATOMIC_SEQ_CST. Handle clauses on OMP_TASKGROUP.
(dependent_omp_for_p): Always return true for range for loops if
processing_template_decl. Return true if class type iterator
does not have INTEGER_CST increment.
* semantics.c: Include memmodel.h.
(handle_omp_array_sections_1): Handle OMP_CLAUSE_{IN,TASK}_REDUCTION
like OMP_CLAUSE_REDUCTION.
(handle_omp_array_sections): Likewise. Call save_expr on array
reductions before calling build_index_type. Handle depend clauses
with iterators.
(finish_omp_reduction_clause): Call save_expr for whole array
reduction sizes. Don't mark OMP_CLAUSE_DECL addressable if it has
reference type. Do mark decl_placeholder addressable if needed.
Use error_at with OMP_CLAUSE_LOCATION (c) as first argument instead
of error.
(cp_omp_finish_iterators): New function.
(finish_omp_clauses): Don't diagnose nonmonotonic clause with static,
runtime or auto schedule kinds. Diagnose nogroup clause used with
reduction clause(s). Handle depend clause with
OMP_CLAUSE_DEPEND_DEPOBJ. Diagnose bit-fields. Require
omp_depend_t type for OMP_CLAUSE_DEPEND_DEPOBJ kinds and
some different type for other kinds. Use cp_build_addr_expr
and cp_build_indirect_ref instead of cxx_mark_addressable.
Handle depend clauses with iterators. Only handle static data members
in the special case that const qualified vars may be specified in
firstprivate clause. Complain if const qualified vars without mutable
members are mentioned in data-sharing clauses other than firstprivate
or shared. Use error_at with OMP_CLAUSE_LOCATION (c) as first
argument instead of error. Diagnose more than one nontemporal clause
refering to the same variable. Use error_at rather than error for
priority and hint clause diagnostics. Fix pasto for hint clause.
Diagnose hint expression that doesn't fold into INTEGER_CST.
Diagnose if clause with modifier other than cancel. Handle
OMP_CLAUSE_{IN,TASK}_REDUCTION like OMP_CLAUSE_REDUCTION. Allow any
lvalue as OMP_CLAUSE_DEPEND operand (besides array section), adjust
diagnostics.
(handle_omp_for_class_iterator): Don't create a new TREE_LIST if one
has been created already for range for, just fill TREE_PURPOSE and
TREE_VALUE. Call cp_fully_fold on incr.
(finish_omp_for): Don't check cond/incr if cond is global_namespace.
Pass to c_omp_check_loop_iv_exprs orig_declv if non-NULL. Don't
use IS_EMPTY_STMT on NULL pre_body. Adjust c_finish_omp_for caller.
(finish_omp_for_block): New function.
(finish_omp_atomic): Add LOC argument, pass it through
to c_finish_omp_atomic and set it as location of OMP_ATOMIC* trees.
Remove SEQ_CST argument. Add CLAUSES and MO arguments. Adjust
c_finish_omp_atomic caller. Stick clauses if any into first argument
of wrapping OMP_ATOMIC.
(finish_omp_depobj): New function.
(finish_omp_flush): Add MO argument, if not
MEMMODEL_LAST, emit __atomic_thread_fence call with the given value.
(finish_omp_cancel): Diagnose if clause with modifier other than
cancel.
gcc/fortran/
* trans-openmp.c (gfc_trans_omp_clauses): Use
OMP_CLAUSE_DEFAULTMAP_SET_KIND.
(gfc_trans_omp_atomic): Set OMP_ATOMIC_MEMORY_ORDER
rather than OMP_ATOMIC_SEQ_CST.
(gfc_trans_omp_taskgroup): Build OMP_TASKGROUP using
make_node instead of build1_loc.
* types.def (BT_FN_VOID_BOOL, BT_FN_VOID_SIZE_SIZE_PTR,
BT_FN_UINT_UINT_PTR_PTR, BT_FN_UINT_OMPFN_PTR_UINT_UINT,
BT_FN_BOOL_UINT_LONGPTR_LONG_LONG_LONGPTR_LONGPTR_PTR_PTR,
BT_FN_BOOL_UINT_ULLPTR_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR,
BT_FN_BOOL_LONG_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR_PTR_PTR,
BT_FN_BOOL_BOOL_ULL_ULL_ULL_LONG_ULL_ULLPTR_ULLPTR_PTR_PTR): New.
(BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR): Formatting fix.
gcc/testsuite/
* c-c++-common/gomp/atomic-17.c: New test.
* c-c++-common/gomp/atomic-18.c: New test.
* c-c++-common/gomp/atomic-19.c: New test.
* c-c++-common/gomp/atomic-20.c: New test.
* c-c++-common/gomp/atomic-21.c: New test.
* c-c++-common/gomp/atomic-22.c: New test.
* c-c++-common/gomp/clauses-1.c (r2): New variable.
(foo): Add ntm argument and test if and nontemporal clauses on
constructs with simd.
(bar): Put taskloop simd inside of taskgroup with task_reduction,
use in_reduction clause instead of reduction. Add another
taskloop simd without nogroup clause, but with reduction clause and
a new in_reduction. Add ntm and i3 arguments. Test if and
nontemporal clauses on constructs with simd. Change if clauses on
some constructs from specific to the particular constituents to one
without a modifier. Add new tests for combined host teams and for
new parallel master and {,parallel }master taskloop{, simd} combined
constructs.
(baz): New function with host teams tests.
* gcc.dg/gomp/combined-1.c: Moved to ...
* c-c++-common/gomp/combined-1.c: ... here. Adjust expected library
call.
* c-c++-common/gomp/combined-2.c: New test.
* c-c++-common/gomp/combined-3.c: New test.
* c-c++-common/gomp/critical-1.c: New test.
* c-c++-common/gomp/critical-2.c: New test.
* c-c++-common/gomp/default-1.c: New test.
* c-c++-common/gomp/defaultmap-1.c: New test.
* c-c++-common/gomp/defaultmap-2.c: New test.
* c-c++-common/gomp/defaultmap-3.c: New test.
* c-c++-common/gomp/depend-5.c: New test.
* c-c++-common/gomp/depend-6.c: New test.
* c-c++-common/gomp/depend-iterator-1.c: New test.
* c-c++-common/gomp/depend-iterator-2.c: New test.
* c-c++-common/gomp/depobj-1.c: New test.
* c-c++-common/gomp/flush-1.c: New test.
* c-c++-common/gomp/flush-2.c: New test.
* c-c++-common/gomp/for-1.c: New test.
* c-c++-common/gomp/for-2.c: New test.
* c-c++-common/gomp/for-3.c: New test.
* c-c++-common/gomp/for-4.c: New test.
* c-c++-common/gomp/for-5.c: New test.
* c-c++-common/gomp/for-6.c: New test.
* c-c++-common/gomp/for-7.c: New test.
* c-c++-common/gomp/if-1.c (foo): Add some further tests.
* c-c++-common/gomp/if-2.c (foo): Likewise. Expect slightly different
diagnostics wording in one case.
* c-c++-common/gomp/if-3.c: New test.
* c-c++-common/gomp/master-combined-1.c: New test.
* c-c++-common/gomp/master-combined-2.c: New test.
* c-c++-common/gomp/nontemporal-1.c: New test.
* c-c++-common/gomp/nontemporal-2.c: New test.
* c-c++-common/gomp/reduction-task-1.c: New test.
* c-c++-common/gomp/reduction-task-2.c: New test.
* c-c++-common/gomp/requires-1.c: New test.
* c-c++-common/gomp/requires-2.c: New test.
* c-c++-common/gomp/requires-3.c: New test.
* c-c++-common/gomp/requires-4.c: New test.
* c-c++-common/gomp/schedule-modifiers-1.c (bar): Don't expect
diagnostics for nonmonotonic modifier with static, runtime or auto
schedule kinds.
* c-c++-common/gomp/simd7.c: New test.
* c-c++-common/gomp/target-data-1.c: New test.
* c-c++-common/gomp/taskloop-reduction-1.c: New test.
* c-c++-common/gomp/taskwait-depend-1.c: New test.
* c-c++-common/gomp/teams-1.c: New test.
* c-c++-common/gomp/teams-2.c: New test.
* gcc.dg/gomp/appendix-a/a.24.1.c: Update from OpenMP examples. Add
shared(c) clause.
* gcc.dg/gomp/atomic-5.c (f1): Add another expected error.
* gcc.dg/gomp/clause-1.c: Adjust expected diagnostics for const
qualified vars without mutable member no longer being predeterined
shared.
* gcc.dg/gomp/sharing-1.c: Likewise.
* g++.dg/gomp/clause-3.C: Likewise.
* g++.dg/gomp/member-2.C: Likewise.
* g++.dg/gomp/predetermined-1.C: Likewise.
* g++.dg/gomp/private-1.C: Likewise.
* g++.dg/gomp/sharing-1.C: Likewise.
* g++.dg/gomp/sharing-2.C: Likewise. Add a few tests with aggregate
const static data member without mutable elements.
* gcc.dg/gomp/for-4.c: Expected nonmonotonic functions in the dumps.
* gcc.dg/gomp/for-5.c: Likewise.
* gcc.dg/gomp/for-6.c: Change expected library call.
* gcc.dg/gomp/pr39495-2.c (foo): Don't expect errors on !=.
* gcc.dg/gomp/reduction-2.c: New test.
* gcc.dg/gomp/simd-1.c: New test.
* gcc.dg/gomp/teams-1.c: Adjust expected diagnostic lines.
* g++.dg/gomp/atomic-18.C: New test.
* g++.dg/gomp/atomic-19.C: New test.
* g++.dg/gomp/atomic-5.C (f1): Adjust expected lines of read-only
variable messages. Add another expected error.
* g++.dg/gomp/critical-3.C: New test.
* g++.dg/gomp/depend-iterator-1.C: New test.
* g++.dg/gomp/depend-iterator-2.C: New test.
* g++.dg/gomp/depobj-1.C: New test.
* g++.dg/gomp/doacross-1.C: New test.
* g++.dg/gomp/for-21.C: New test.
* g++.dg/gomp/for-4.C: Expected nonmonotonic functions in the dumps.
* g++.dg/gomp/for-5.C: Likewise.
* g++.dg/gomp/for-6.C: Change expected library call.
* g++.dg/gomp/loop-4.C: New test.
* g++.dg/gomp/pr33372-1.C: Adjust location of the expected
diagnostics.
* g++.dg/gomp/pr33372-3.C: Likewise.
* g++.dg/gomp/pr39495-2.C (foo): Don't expect errors on !=.
* g++.dg/gomp/simd-2.C: New test.
* g++.dg/gomp/tpl-atomic-2.C: Adjust expected diagnostic lines.
include/
* gomp-constants.h (GOMP_TASK_FLAG_REDUCTION,
GOMP_DEPEND_IN, GOMP_DEPEND_OUT, GOMP_DEPEND_INOUT,
GOMP_DEPEND_MUTEXINOUTSET): Define.
libgomp/
* affinity.c (gomp_display_affinity_place): New function.
* affinity-fmt.c: New file.
* alloc.c (gomp_aligned_alloc, gomp_aligned_free): New functions.
* config/linux/affinity.c (gomp_display_affinity_place): New function.
* config/nvptx/icv-device.c (omp_get_num_teams, omp_get_team_num):
Move these functions to ...
* config/nvptx/teams.c: ... here. New file.
* config/nvptx/target.c (omp_pause_resource, omp_pause_resource_all):
New functions.
* config/nvptx/team.c (gomp_team_start, gomp_pause_host): New
functions.
* configure.ac: Check for aligned_alloc, posix_memalign, memalign
and _aligned_malloc.
(HAVE_UNAME, HAVE_GETHOSTNAME, HAVE_GETPID): Add new tests.
* configure.tgt: Add -DUSING_INITIAL_EXEC_TLS to XCFLAGS for Linux.
* env.c (gomp_display_affinity_var, gomp_affinity_format_var,
gomp_affinity_format_len): New variables.
(parse_schedule): Parse monotonic and nonmonotonic modifiers in
OMP_SCHEDULE variable. Set GFS_MONOTONIC for monotonic schedules.
(handle_omp_display_env): Display monotonic/nonmonotonic schedule
modifiers. Display (non-default) chunk sizes. Print
OMP_DISPLAY_AFFINITY and OMP_AFFINITY_FORMAT.
(initialize_env): Don't call pthread_attr_setdetachstate. Handle
OMP_DISPLAY_AFFINITY and OMP_AFFINITY_FORMAT env vars.
* fortran.c: Include stdio.h and string.h.
(omp_pause_resource, omp_pause_resource_all): Add ialias_redirect.
(omp_get_schedule_, omp_get_schedule_8_): Mask off GFS_MONOTONIC bit.
(omp_set_affinity_format_, omp_get_affinity_format_,
omp_display_affinity_, omp_capture_affinity_, omp_pause_resource_,
omp_pause_resource_all_): New functions.
* icv.c (omp_set_schedule): Mask off omp_sched_monotonic bit in
switch.
* icv-device.c (omp_get_num_teams, omp_get_team_num): Move these
functions to ...
* teams.c: ... here. New file.
* libgomp_g.h: Include gstdint.h.
(GOMP_loop_nonmonotonic_runtime_start,
GOMP_loop_maybe_nonmonotonic_runtime_start, GOMP_loop_start,
GOMP_loop_ordered_start, GOMP_loop_nonmonotonic_runtime_next,
GOMP_loop_maybe_nonmonotonic_runtime_next, GOMP_loop_doacross_start,
GOMP_parallel_loop_nonmonotonic_runtime,
GOMP_parallel_loop_maybe_nonmonotonic_runtime,
GOMP_loop_ull_nonmonotonic_runtime_start,
GOMP_loop_ull_maybe_nonmonotonic_runtime_start, GOMP_loop_ull_start,
GOMP_loop_ull_ordered_start, GOMP_loop_ull_nonmonotonic_runtime_next,
GOMP_loop_ull_maybe_nonmonotonic_runtime_next,
GOMP_loop_ull_doacross_start, GOMP_parallel_reductions,
GOMP_taskwait_depend, GOMP_taskgroup_reduction_register,
GOMP_taskgroup_reduction_unregister, GOMP_task_reduction_remap,
GOMP_workshare_task_reduction_unregister, GOMP_sections2_start,
GOMP_teams_reg): Declare.
* libgomp.h (GOMP_HAVE_EFFICIENT_ALIGNED_ALLOC): Define unless
gomp_aligned_alloc uses fallback implementation.
(gomp_aligned_alloc, gomp_aligned_free): Declare.
(enum gomp_schedule_type): Add GFS_MONOTONIC.
(struct gomp_doacross_work_share): Add extra field.
(struct gomp_work_share): Add task_reductions field.
(struct gomp_taskgroup): Add workshare and reductions fields.
(GOMP_NEEDS_THREAD_HANDLE): Define if needed.
(gomp_thread_handle): New typedef.
(gomp_display_affinity_place, gomp_set_affinity_format,
gomp_display_string, gomp_display_affinity,
gomp_display_affinity_thread): Declare.
(gomp_doacross_init, gomp_doacross_ull_init): Add size_t argument.
(gomp_parallel_reduction_register, gomp_workshare_taskgroup_start,
gomp_workshare_task_reduction_register): Declare.
(gomp_team_start): Add taskgroup argument.
(gomp_pause_host): Declare.
(gomp_init_work_share, gomp_work_share_start): Change bool argument
to size_t.
(gomp_thread_self, gomp_thread_to_pthread_t): New inline functions.
* libgomp.map (GOMP_5.0): Export GOMP_loop_start,
GOMP_loop_ordered_start, GOMP_loop_doacross_start,
GOMP_loop_ull_start, GOMP_loop_ull_ordered_start,
GOMP_loop_ull_doacross_start,
GOMP_workshare_task_reduction_unregister, GOMP_sections2_start,
GOMP_loop_maybe_nonmonotonic_runtime_next,
GOMP_loop_maybe_nonmonotonic_runtime_start,
GOMP_loop_nonmonotonic_runtime_next,
GOMP_loop_nonmonotonic_runtime_start,
GOMP_loop_ull_maybe_nonmonotonic_runtime_next,
GOMP_loop_ull_maybe_nonmonotonic_runtime_start,
GOMP_loop_ull_nonmonotonic_runtime_next,
GOMP_loop_ull_nonmonotonic_runtime_start,
GOMP_parallel_loop_maybe_nonmonotonic_runtime,
GOMP_parallel_loop_nonmonotonic_runtime, GOMP_parallel_reductions,
GOMP_taskgroup_reduction_register,
GOMP_taskgroup_reduction_unregister, GOMP_task_reduction_remap,
GOMP_teams_reg and GOMP_taskwait_depend.
(OMP_5.0): Export omp_pause_resource{,_all}{,_},
omp_{capture,display}_affinity{,_}, and
omp_[gs]et_affinity_format{,_}.
* loop.c: Include string.h.
(GOMP_loop_runtime_next): Add ialias.
(GOMP_taskgroup_reduction_register): Add ialias_redirect.
(gomp_loop_static_start, gomp_loop_dynamic_start,
gomp_loop_guided_start, gomp_loop_ordered_static_start,
gomp_loop_ordered_dynamic_start, gomp_loop_ordered_guided_start,
gomp_loop_doacross_static_start, gomp_loop_doacross_dynamic_start,
gomp_loop_doacross_guided_start): Adjust gomp_work_share_start
or gomp_doacross_init callers.
(gomp_adjust_sched, GOMP_loop_start, GOMP_loop_ordered_start,
GOMP_loop_doacross_start): New functions.
(GOMP_loop_runtime_start, GOMP_loop_ordered_runtime_start,
GOMP_loop_doacross_runtime_start, GOMP_parallel_loop_runtime_start):
Mask off GFS_MONOTONIC bit.
(GOMP_loop_maybe_nonmonotonic_runtime_next,
GOMP_loop_maybe_nonmonotonic_runtime_start,
GOMP_loop_nonmonotonic_runtime_next,
GOMP_loop_nonmonotonic_runtime_start,
GOMP_parallel_loop_maybe_nonmonotonic_runtime,
GOMP_parallel_loop_nonmonotonic_runtime): New aliases or wrapper
functions.
(gomp_parallel_loop_start): Pass NULL as taskgroup to
gomp_team_start.
* loop_ull.c: Include string.h.
(GOMP_loop_ull_runtime_next): Add ialias.
(GOMP_taskgroup_reduction_register): Add ialias_redirect.
(gomp_loop_ull_static_start, gomp_loop_ull_dynamic_start,
gomp_loop_ull_guided_start, gomp_loop_ull_ordered_static_start,
gomp_loop_ull_ordered_dynamic_start,
gomp_loop_ull_ordered_guided_start,
gomp_loop_ull_doacross_static_start,
gomp_loop_ull_doacross_dynamic_start,
gomp_loop_ull_doacross_guided_start): Adjust gomp_work_share_start
and gomp_doacross_ull_init callers.
(gomp_adjust_sched, GOMP_loop_ull_start, GOMP_loop_ull_ordered_start,
GOMP_loop_ull_doacross_start): New functions.
(GOMP_loop_ull_runtime_start,
GOMP_loop_ull_ordered_runtime_start,
GOMP_loop_ull_doacross_runtime_start): Mask off GFS_MONOTONIC bit.
(GOMP_loop_ull_maybe_nonmonotonic_runtime_next,
GOMP_loop_ull_maybe_nonmonotonic_runtime_start,
GOMP_loop_ull_nonmonotonic_runtime_next,
GOMP_loop_ull_nonmonotonic_runtime_start): Likewise.
* Makefile.am (libgomp_la_SOURCES): Add teams.c and affinity-fmt.c.
* omp.h.in (enum omp_sched_t): Add omp_sched_monotonic.
(omp_pause_resource_t, omp_depend_t): New typedefs.
(enum omp_lock_hint_t): Renamed to ...
(enum omp_sync_hint_t): ... this. Define omp_sync_hint_*
enumerators using numbers and omp_lock_hint_* as their aliases.
(omp_lock_hint_t): New typedef. Rename to ...
(omp_sync_hint_t): ... this.
(omp_init_lock_with_hint, omp_init_nest_lock_with_hint): Use
omp_sync_hint_t instead of omp_lock_hint_t.
(omp_pause_resource, omp_pause_resource_all, omp_set_affinity_format,
omp_get_affinity_format, omp_display_affinity, omp_capture_affinity):
Declare.
(omp_target_is_present, omp_target_disassociate_ptr):
Change first argument from void * to const void *.
(omp_target_memcpy, omp_target_memcpy_rect): Change second argument
from void * to const void *.
(omp_target_associate_ptr): Change first and second arguments from
void * to const void *.
* omp_lib.f90.in (omp_pause_resource_kind, omp_pause_soft,
omp_pause_hard): New parameters.
(omp_pause_resource, omp_pause_resource_all, omp_set_affinity_format,
omp_get_affinity_format, omp_display_affinity, omp_capture_affinity):
New interfaces.
* omp_lib.h.in (omp_pause_resource_kind, omp_pause_soft,
omp_pause_hard): New parameters.
(omp_pause_resource, omp_pause_resource_all, omp_set_affinity_format,
omp_get_affinity_format, omp_display_affinity, omp_capture_affinity):
New externals.
* ordered.c (gomp_doacross_init, gomp_doacross_ull_init): Add
EXTRA argument. If not needed to prepare array, if extra is 0,
clear ws->doacross, otherwise allocate just doacross structure and
extra payload. If array is needed, allocate also extra payload.
(GOMP_doacross_post, GOMP_doacross_wait, GOMP_doacross_ull_post,
GOMP_doacross_ull_wait): Handle doacross->array == NULL like
doacross == NULL.
* parallel.c (GOMP_parallel_start): Pass NULL as taskgroup to
gomp_team_start.
(GOMP_parallel): Likewise. Formatting fix.
(GOMP_parallel_reductions): New function.
(GOMP_cancellation_point): If taskgroup has workshare
flag set, check cancelled of prev taskgroup if any.
(GOMP_cancel): If taskgroup has workshare flag set, set cancelled
on prev taskgroup if any.
* sections.c: Include string.h.
(GOMP_taskgroup_reduction_register): Add ialias_redirect.
(GOMP_sections_start): Adjust gomp_work_share_start caller.
(GOMP_sections2_start): New function.
(GOMP_parallel_sections_start, GOMP_parallel_sections):
Pass NULL as taskgroup to gomp_team_start.
* single.c (GOMP_single_start, GOMP_single_copy_start): Adjust
gomp_work_share_start callers.
* target.c (GOMP_target_update_ext, GOMP_target_enter_exit_data):
If taskgroup has workshare flag set, check cancelled on prev
taskgroup if any. Guard all cancellation tests with
gomp_cancel_var test.
(omp_target_is_present, omp_target_disassociate_ptr):
Change ptr argument from void * to const void *.
(omp_target_memcpy): Change src argument from void * to const void *.
(omp_target_memcpy_rect): Likewise.
(omp_target_memcpy_rect_worker): Likewise. Use const char * casts
instead of char * where needed.
(omp_target_associate_ptr): Change host_ptr and device_ptr arguments
from void * to const void *.
(omp_pause_resource, omp_pause_resource_all): New functions.
* task.c (gomp_task_handle_depend): Handle new depend array format
in addition to the old. Handle mutexinoutset kinds the same as
inout for now, handle unspecified kinds.
(gomp_create_target_task): If taskgroup has workshare flag set, check
cancelled on prev taskgroup if any. Guard all cancellation tests with
gomp_cancel_var test. Handle new depend array format count in
addition to the old.
(GOMP_task): Likewise. Adjust function comment.
(gomp_task_run_pre): If taskgroup has workshare flag set, check
cancelled on prev taskgroup if any. Guard all cancellation tests with
gomp_cancel_var test.
(GOMP_taskwait_depend): New function.
(gomp_task_maybe_wait_for_dependencies): Handle new depend array
format in addition to the old. Handle mutexinoutset kinds the same as
inout for now, handle unspecified kinds. Fix a function comment typo.
(gomp_taskgroup_init): New function.
(GOMP_taskgroup_start): Use it.
(gomp_reduction_register, gomp_create_artificial_team,
GOMP_taskgroup_reduction_register,
GOMP_taskgroup_reduction_unregister, GOMP_task_reduction_remap,
gomp_parallel_reduction_register,
gomp_workshare_task_reduction_register,
gomp_workshare_taskgroup_start,
GOMP_workshare_task_reduction_unregister): New functions.
* taskloop.c (GOMP_taskloop): If taskgroup has workshare flag set,
check cancelled on prev taskgroup if any. Guard all cancellation
tests with gomp_cancel_var test. Handle GOMP_TASK_FLAG_REDUCTION flag
by calling GOMP_taskgroup_reduction_register.
* team.c (gomp_thread_attr): Remove comment.
(struct gomp_thread_start_data): Add handle field.
(gomp_thread_start): Call pthread_detach.
(gomp_new_team): Adjust gomp_init_work_share caller.
(gomp_free_pool_helper): Call pthread_detach.
(gomp_team_start): Add taskgroup argument, initialize implicit
tasks' taskgroup field to that. Don't call
pthread_attr_setdetachstate. Handle OMP_DISPLAY_AFFINITY env var.
(gomp_team_end): Determine nesting by thr->ts.level != 0
rather than thr->ts.team != NULL.
(gomp_pause_pool_helper, gomp_pause_host): New functions.
* work.c (alloc_work_share): Use gomp_aligned_alloc instead of
gomp_malloc if GOMP_HAVE_EFFICIENT_ALIGNED_ALLOC is defined.
(gomp_init_work_share): Change ORDERED argument from bool to size_t,
if more than 1 allocate also extra payload at the end of array. Never
keep ordered_team_ids NULL, set it to inline_ordered_team_ids instead.
(gomp_work_share_start): Change ORDERED argument from bool to size_t,
return true instead of ws.
* Makefile.in: Regenerated.
* configure: Regenerated.
* config.h.in: Regenerated.
* testsuite/libgomp.c/cancel-for-2.c (foo): Use cancel modifier
in some cases.
* testsuite/libgomp.c-c++-common/cancel-parallel-1.c: New test.
* testsuite/libgomp.c-c++-common/cancel-taskgroup-3.c: New test.
* testsuite/libgomp.c-c++-common/depend-iterator-1.c: New test.
* testsuite/libgomp.c-c++-common/depend-iterator-2.c: New test.
* testsuite/libgomp.c-c++-common/depend-mutexinout-1.c: New test.
* testsuite/libgomp.c-c++-common/depend-mutexinout-2.c: New test.
* testsuite/libgomp.c-c++-common/depobj-1.c: New test.
* testsuite/libgomp.c-c++-common/display-affinity-1.c: New test.
* testsuite/libgomp.c-c++-common/for-10.c: New test.
* testsuite/libgomp.c-c++-common/for-11.c: New test.
* testsuite/libgomp.c-c++-common/for-12.c: New test.
* testsuite/libgomp.c-c++-common/for-13.c: New test.
* testsuite/libgomp.c-c++-common/for-14.c: New test.
* testsuite/libgomp.c-c++-common/for-15.c: New test.
* testsuite/libgomp.c-c++-common/for-2.h: If CONDNE macro is defined,
define a different N(test), don't define N(f0) to N(f14), but instead
define N(f20) to N(f34) using != comparisons.
* testsuite/libgomp.c-c++-common/for-7.c: New test.
* testsuite/libgomp.c-c++-common/for-8.c: New test.
* testsuite/libgomp.c-c++-common/for-9.c: New test.
* testsuite/libgomp.c-c++-common/master-combined-1.c: New test.
* testsuite/libgomp.c-c++-common/pause-1.c: New test.
* testsuite/libgomp.c-c++-common/pause-2.c: New test.
* testsuite/libgomp.c-c++-common/pr66199-10.c: New test.
* testsuite/libgomp.c-c++-common/pr66199-11.c: New test.
* testsuite/libgomp.c-c++-common/pr66199-12.c: New test.
* testsuite/libgomp.c-c++-common/pr66199-13.c: New test.
* testsuite/libgomp.c-c++-common/pr66199-14.c: New test.
* testsuite/libgomp.c-c++-common/simd-1.c: New test.
* testsuite/libgomp.c-c++-common/taskloop-reduction-1.c: New test.
* testsuite/libgomp.c-c++-common/taskloop-reduction-2.c: New test.
* testsuite/libgomp.c-c++-common/taskloop-reduction-3.c: New test.
* testsuite/libgomp.c-c++-common/taskloop-reduction-4.c: New test.
* testsuite/libgomp.c-c++-common/task-reduction-11.c: New test.
* testsuite/libgomp.c-c++-common/task-reduction-12.c: New test.
* testsuite/libgomp.c-c++-common/task-reduction-1.c: New test.
* testsuite/libgomp.c-c++-common/task-reduction-2.c: New test.
* testsuite/libgomp.c-c++-common/task-reduction-3.c: New test.
* testsuite/libgomp.c-c++-common/task-reduction-4.c: New test.
* testsuite/libgomp.c-c++-common/task-reduction-5.c: New test.
* testsuite/libgomp.c-c++-common/task-reduction-6.c: New test.
* testsuite/libgomp.c-c++-common/task-reduction-7.c: New test.
* testsuite/libgomp.c-c++-common/task-reduction-8.c: New test.
* testsuite/libgomp.c-c++-common/task-reduction-9.c: New test.
* testsuite/libgomp.c-c++-common/taskwait-depend-1.c: New test.
* testsuite/libgomp.c++/depend-1.C: New test.
* testsuite/libgomp.c++/depend-iterator-1.C: New test.
* testsuite/libgomp.c++/depobj-1.C: New test.
* testsuite/libgomp.c++/for-16.C: New test.
* testsuite/libgomp.c++/for-21.C: New test.
* testsuite/libgomp.c++/for-22.C: New test.
* testsuite/libgomp.c++/for-23.C: New test.
* testsuite/libgomp.c++/for-24.C: New test.
* testsuite/libgomp.c++/for-25.C: New test.
* testsuite/libgomp.c++/for-26.C: New test.
* testsuite/libgomp.c++/taskloop-reduction-1.C: New test.
* testsuite/libgomp.c++/taskloop-reduction-2.C: New test.
* testsuite/libgomp.c++/taskloop-reduction-3.C: New test.
* testsuite/libgomp.c++/taskloop-reduction-4.C: New test.
* testsuite/libgomp.c++/task-reduction-10.C: New test.
* testsuite/libgomp.c++/task-reduction-11.C: New test.
* testsuite/libgomp.c++/task-reduction-12.C: New test.
* testsuite/libgomp.c++/task-reduction-13.C: New test.
* testsuite/libgomp.c++/task-reduction-14.C: New test.
* testsuite/libgomp.c++/task-reduction-15.C: New test.
* testsuite/libgomp.c++/task-reduction-16.C: New test.
* testsuite/libgomp.c++/task-reduction-17.C: New test.
* testsuite/libgomp.c++/task-reduction-18.C: New test.
* testsuite/libgomp.c++/task-reduction-19.C: New test.
* testsuite/libgomp.c/task-reduction-1.c: New test.
* testsuite/libgomp.c++/task-reduction-1.C: New test.
* testsuite/libgomp.c/task-reduction-2.c: New test.
* testsuite/libgomp.c++/task-reduction-2.C: New test.
* testsuite/libgomp.c++/task-reduction-3.C: New test.
* testsuite/libgomp.c++/task-reduction-4.C: New test.
* testsuite/libgomp.c++/task-reduction-5.C: New test.
* testsuite/libgomp.c++/task-reduction-6.C: New test.
* testsuite/libgomp.c++/task-reduction-7.C: New test.
* testsuite/libgomp.c++/task-reduction-8.C: New test.
* testsuite/libgomp.c++/task-reduction-9.C: New test.
* testsuite/libgomp.c/teams-1.c: New test.
* testsuite/libgomp.c/teams-2.c: New test.
* testsuite/libgomp.c/thread-limit-4.c: New test.
* testsuite/libgomp.c/thread-limit-5.c: New test.
* testsuite/libgomp.fortran/display-affinity-1.f90: New test.
From-SVN: r265930
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
libgomp/
* oacc-mem.c (memcpy_tofrom_device): New function, combined from
acc_memcpy_to/from_device functions, now with async parameter.
(acc_memcpy_to_device): Modify to use memcpy_tofrom_device.
(acc_memcpy_from_device): Likewise.
(acc_memcpy_to_device_async): New API function.
(acc_memcpy_from_device_async): Likewise.
(present_create_copy): Add async parameter and async setting/unsetting.
(acc_create): Adjust present_create_copy call.
(acc_copyin): Likewise.
(acc_present_or_create): Likewise.
(acc_present_or_copyin): Likewise.
(acc_create_async): New API function.
(acc_copyin_async): New API function.
(delete_copyout): Add async parameter and async setting/unsetting.
(acc_delete): Adjust delete_copyout call.
(acc_copyout): Likewise.
(acc_delete_async): New API function.
(acc_copyout_async): Likewise.
(update_dev_host): Add async parameter and async setting/unsetting.
(acc_update_device): Adjust update_dev_host call.
(acc_update_self): Likewise.
(acc_update_device_async): New API function.
(acc_update_self_async): Likewise.
* openacc.h (acc_copyin_async): Declare new API function.
(acc_create_async): Likewise.
(acc_copyout_async): Likewise.
(acc_delete_async): Likewise.
(acc_update_device_async): Likewise.
(acc_update_self_async): Likewise.
(acc_memcpy_to_device_async): Likewise.
(acc_memcpy_from_device_async): Likewise.
* openacc_lib.h (acc_copyin_async_32_h): New subroutine.
(acc_copyin_async_64_h): New subroutine.
(acc_copyin_async_array_h): New subroutine.
(acc_create_async_32_h): New subroutine.
(acc_create_async_64_h): New subroutine.
(acc_create_async_array_h): New subroutine.
(acc_copyout_async_32_h): New subroutine.
(acc_copyout_async_64_h): New subroutine.
(acc_copyout_async_array_h): New subroutine.
(acc_delete_async_32_h): New subroutine.
(acc_delete_async_64_h): New subroutine.
(acc_delete_async_array_h): New subroutine.
(acc_update_device_async_32_h): New subroutine.
(acc_update_device_async_64_h): New subroutine.
(acc_update_device_async_array_h): New subroutine.
(acc_update_self_async_32_h): New subroutine.
(acc_update_self_async_64_h): New subroutine.
(acc_update_self_async_array_h): New subroutine.
* openacc.f90 (acc_copyin_async_32_h): New subroutine.
(acc_copyin_async_64_h): New subroutine.
(acc_copyin_async_array_h): New subroutine.
(acc_create_async_32_h): New subroutine.
(acc_create_async_64_h): New subroutine.
(acc_create_async_array_h): New subroutine.
(acc_copyout_async_32_h): New subroutine.
(acc_copyout_async_64_h): New subroutine.
(acc_copyout_async_array_h): New subroutine.
(acc_delete_async_32_h): New subroutine.
(acc_delete_async_64_h): New subroutine.
(acc_delete_async_array_h): New subroutine.
(acc_update_device_async_32_h): New subroutine.
(acc_update_device_async_64_h): New subroutine.
(acc_update_device_async_array_h): New subroutine.
(acc_update_self_async_32_h): New subroutine.
(acc_update_self_async_64_h): New subroutine.
(acc_update_self_async_array_h): New subroutine.
* libgomp.map (OACC_2.5): Add acc_copyin_async*, acc_copyout_async*,
acc_copyout_finalize_async*, acc_create_async*, acc_delete_async*,
acc_delete_finalize_async*, acc_memcpy_from_device_async*,
acc_memcpy_to_device_async*, acc_update_device_async*, and
acc_update_self_async* entries.
* testsuite/libgomp.oacc-c-c++-common/lib-94.c: New test.
* testsuite/libgomp.oacc-c-c++-common/lib-95.c: New test.
* testsuite/libgomp.oacc-fortran/lib-16.f90: New test.
From-SVN: r265842
This patch updates GCC to use autoconf 2.69 and automake 1.15.1.
(That's not the latest automake version, but it's the one used by
binutils-gdb, with which consistency is desirable, and in any case
seems a useful incremental update that should make a future update to
1.16.1 easier.)
The changes are generally similar to the binutils-gdb ones, and are
copied from there where shared files and directories are involved
(there are some further changes to such shared directories, however,
which I'd expect to apply to binutils-gdb once this patch is in GCC).
Largely, obsolete AC_PREREQ calls are removed, while many
AC_LANG_SOURCE calls are added to avoid warnings from aclocal and
autoconf. Multilib support is no longer included in core automake,
meaning that multilib.am needs copying from automake's contrib
directory into the GCC source tree. Autoconf 2.69 has Go support, so
local copies of that support are removed. I hope the D support will
soon be submitted to upstream autoconf so the local copy of that can
be removed in a future update. Changes to how automake generates
runtest calls mean quotes are removed from RUNTEST definitions in five
lib*/testsuite/Makefile.am files (libatomic, libgomp, libitm,
libphobos, libvtv; some others have RUNTEST definitions without
quotes, which are still OK); libgo and libphobos also get
-Wno-override added to AM_INIT_AUTOMAKE so those overrides of RUNTEST
do not generate automake warnings.
Note that the regeneration did not include regeneration of
fixincludes/config.h.in (attempting such regeneration resulted in all
the USED_FOR_TARGET conditionals disappearing; and I don't see
anything in the fixincludes/ directory that would result in such
conditionals being generated, unlike in the gcc/ directory). Also
note that libvtv/testsuite/other-tests/Makefile.in was not
regenerated; that directory is not listed as a subdirectory for which
Makefile.in gets regenerated by calling "automake" in libvtv/, so I'm
not sure how it's meant to be regenerated.
While I mostly fixed warnings should running aclocal / automake /
autoconf, there were various such warnings from automake in the
libgfortran, libgo, libgomp, liboffloadmic, libsanitizer, libphobos
directories that I did not fix, preferring to leave those to the
relevant subsystem maintainers. Specifically, most of those warnings
were of the following form (example from libgfortran):
Makefile.am:48: warning: source file 'caf/single.c' is in a subdirectory,
Makefile.am:48: but option 'subdir-objects' is disabled
automake: warning: possible forward-incompatibility.
automake: At least a source file is in a subdirectory, but the 'subdir-objects'
automake: automake option hasn't been enabled. For now, the corresponding output
automake: object file(s) will be placed in the top-level directory. However,
automake: this behaviour will change in future Automake versions: they
will
automake: unconditionally cause object files to be placed in the same subdirectory
automake: of the corresponding sources.
automake: You are advised to start using 'subdir-objects' option throughout your
automake: project, to avoid future incompatibilities.
I think it's best for the relevant maintainers to add subdir-objects
and do any other associated Makefile.am changes needed. In some cases
the paths in the warnings involved ../; I don't know if that adds any
extra complications to the use of subdir-objects.
I've tested this with native, cross and Canadian cross builds. The
risk of any OS-specific issues should I hope be rather lower than if a
libtool upgrade were included (we *should* do such an upgrade at some
point, but it's more complicated - it involves identifying all our
local libtool changes to see if any aren't included in the upstream
version we update to, and reverting an upstream libtool patch that's
inappropriate for use in GCC); I think it would be better to get this
update into GCC so that people can test in different configurations
and we can fix any issues found, rather than to try to get more and
more testing done before it goes in.
top level:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* multilib.am: New file. From automake.
Merge from binutils-gdb:
2018-06-19 Simon Marchi <simon.marchi@ericsson.com>
* libtool.m4: Use AC_LANG_SOURCE.
* configure.ac: Remove AC_PREREQ, use AC_LANG_SOURCE.
* ar-lib: New file.
* test-driver: New file.
* configure: Re-generate.
config:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* math.m4, tls.m4: Use AC_LANG_SOURCE.
Merge from binutils-gdb:
2018-06-19 Simon Marchi <simon.marchi@ericsson.com>
* override.m4 (_GCC_AUTOCONF_VERSION): Bump from 2.64 to 2.69.
fixincludes:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* configure.ac: Remove AC_PREREQ.
* aclocal.m4, configure: Regenerate.
gcc:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* configure.ac: Remove AC_PREREQ. Use AC_LANG_SOURCE. Use single
line for second argument of AC_DEFINE_UNQUOTED.
* doc/install.texi (Tools/packages necessary for modifying GCC):
Update to autoconf 2.69 and automake 1.15.1.
* aclocal.m4, config.in, configure: Regenerate.
gnattools:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* configure.ac: Remove AC_PREREQ.
* configure: Regenerate.
gotools:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* config/go.m4: Remove file.
* Makefile.am (ACLOCAL_AMFLAGS): Do not use -I ./config.
* configure.ac: Remove AC_PREREQ. Do not include config/go.m4.
* Makefile.in, aclocal.m4, configure: Regenerate.
intl:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
Merge from binutils-gdb:
2018-06-19 Simon Marchi <simon.marchi@ericsson.com>
* configure.ac: Add AC_USE_SYSTEM_EXTENSIONS, remove AC_PREREQ.
* configure: Re-generate.
* config.h.in: Re-generate.
* aclocal.m4: Re-generate.
libada:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* configure.ac: Remove AC_PREREQ.
* configure: Regenerate.
libatomic:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
* acinclude.m4: Use AC_LANG_SOURCE.
* configure.ac: Remove AC_PREREQ.
* testsuite/Makefile.am (RUNTEST): Remove quotes.
* Makefile.in, aclocal.m4, configure, testsuite/Makefile.in:
Regenerate.
libbacktrace:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
* configure.ac: Remove AC_PREREQ. Use AC_LANG_SOURCE.
* Makefile.in, aclocal.m4, config.h.in, configure: Regenerate.
libcc1:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* configure.ac: Remove AC_PREREQ.
* Makefile.in, aclocal.m4, configure: Regenerate.
libcpp:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* configure.ac: Remove AC_PREREQ. Use AC_LANG_SOURCE.
* aclocal.m4, config.in, configure: Regenerate.
libdecnumber:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
Merge from binutils-gdb:
2018-06-19 Simon Marchi <simon.marchi@ericsson.com>
* configure.ac: Remove AC_PREREQ.
* configure: Re-generate.
* aclocal.m4.
libffi:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
(AUTOMAKE_OPTIONS): Add info-in-builddir.
(CLEANFILES): Remove doc/libffi.info.
* configure.ac: Remove AC_PREREQ.
* Makefile.in, aclocal.m4, configure, fficonfig.h.in,
include/Makefile.in, man/Makefile.in, testsuite/Makefile.in:
Regenerate.
libgcc:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* configure.ac: Remove AC_PREREQ. Use AC_LANG_SOURCE.
* configure: Regenerate.
libgfortran:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
* configure.ac: Remove AC_PREREQ.
* Makefile.in, aclocal.m4, config.h.in, configure: Regenerate.
libgo [logically part of this change but omitted from the commit]:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
* config/go.m4: Remove file.
* config/libtool.m4: Use AC_LANG_SOURCE.
* configure.ac: Remove AC_PREREQ. Use AC_LANG_SOURCE. Use
-Wno-override in AM_INIT_AUTOMAKE call.
* Makefile.in, aclocal.m4, configure, testsuite/Makefile.in:
Regenerate.
libgomp:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am
(AUTOMAKE_OPTIONS): Add info-in-builddir.
(CLEANFILES): Remove libgomp.info.
* configure.ac: Remove AC_PREREQ.
* testsuite/Makefile.am (RUNTEST): Remove quotes.
* Makefile.in, aclocal.m4, configure, testsuite/Makefile.in:
Regenerate.
libhsail-rt:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* configure.ac: Remove AC_PREREQ.
* Makefile.in, aclocal.m4, configure: Regenerate.
libiberty:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
Merge from binutils-gdb:
2018-06-19 Simon Marchi <simon.marchi@ericsson.com>
* configure.ac: Remove AC_PREREQ.
* configure: Re-generate.
* config.in: Re-generate.
libitm:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
(AUTOMAKE_OPTIONS): Add info-in-builddir.
(CLEANFILES): Remove libitm.info.
* configure.ac: Remove AC_PREREQ.
* testsuite/Makefile.am (RUNTEST): Remove quotes.
* Makefile.in, aclocal.m4, configure, testsuite/Makefile.in:
Regenerate.
libobjc:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* configure.ac: Remove AC_PREREQ.
* aclocal.m4, config.h.in, configure: Regenerate.
liboffloadmic:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
* configure.ac: Remove AC_PREREQ.
* plugin/Makefile.am: Include multilib.am.
* plugin/configure.ac: Remove AC_PREREQ.
* Makefile.in, aclocal.m4, configure, plugin/Makefile.in,
plugin/aclocal.m4, plugin/configure: Regenerate.
libphobos:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
* configure.ac: Remove AC_PREREQ. Use -Wno-override in
AM_INIT_AUTOMAKE call.
* m4/autoconf.m4: Add extra argument to AC_LANG_DEFINE call.
* m4/druntime/os.m4: Use AC_LANG_SOURCE.
* testsuite/Makefile.am (RUNTEST): Remove quotes.
* Makefile.in, aclocal.m4, configure, libdruntime/Makefile.in,
src/Makefile.in, testsuite/Makefile.in: Regenerate.
libquadmath:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
(AUTOMAKE_OPTIONS): Remove 1.8. Add info-in-builddir.
(all-local): Define outside conditional code.
(CLEANFILES): Remove libquadmath.info.
* configure.ac: Remove AC_PREREQ.
* Makefile.in, aclocal.m4, config.h.in, configure: Regenerate.
libsanitizer:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
* configure.ac: Remove AC_PREREQ. Use AC_LANG_SOURCE.
* Makefile.in, aclocal.m4, asan/Makefile.in, configure,
interception/Makefile.in, libbacktrace/Makefile.in,
lsan/Makefile.in, sanitizer_common/Makefile.in, tsan/Makefile.in,
ubsan/Makefile.in: Regenerate.
libssp:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
(AUTOMAKE_OPTIONS): Remove 1.9.5.
* configure.ac: Remove AC_PREREQ. Quote argument to
AC_RUN_IFELSE.
* Makefile.in, aclocal.m4, configure: Regenerate.
libstdc++-v3:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
* configure.ac: Remove AC_PREREQ.
* Makefile.in, aclocal.m4, configure, doc/Makefile.in,
include/Makefile.in, libsupc++/Makefile.in, po/Makefile.in,
python/Makefile.in, src/Makefile.in, src/c++11/Makefile.in,
src/c++17/Makefile.in, src/c++98/Makefile.in,
src/filesystem/Makefile.in, testsuite/Makefile.in: Regenerate.
libvtv:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
* configure.ac: Remove AC_PREREQ.
* testsuite/Makefile.am (RUNTEST): Remove quotes.
* Makefile.in, aclocal.m4, configure, testsuite/Makefile.in:
Regenerate.
lto-plugin:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* configure.ac: Remove AC_PREREQ. Use AC_LANG_SOURCE.
* Makefile.in, aclocal.m4, config.h.in, configure: Regenerate.
zlib:
2018-10-31 Joseph Myers <joseph@codesourcery.com>
PR bootstrap/82856
* Makefile.am: Include multilib.am.
Merge from binutils-gdb:
2018-06-19 Simon Marchi <simon.marchi@ericsson.com>
* configure.ac: Modernize AC_INIT call, remove AC_PREREQ.
* Makefile.am (AUTOMAKE_OPTIONS): Remove 1.8, cygnus, add foreign.
* Makefile.in: Re-generate.
* aclocal.m4: Re-generate.
* configure: Re-generate.
From-SVN: r265695
2018-10-29 Joseph Myers <joseph@codesourcery.com>
Julian Brown <julian@codesourcery.com>
* semantics.c (handle_omp_array_sections_1): Allow array sections with
"this" pointer for OpenACC.
Co-Authored-By: Julian Brown <julian@codesourcery.com>
From-SVN: r265591
2018-09-09 Cesar Philippidis <cesar@codesourcery.com>
Julian Brown <julian@codesourcery.com>
PR middle-end/86336
gcc/cp/
* semantics.c (finish_omp_clauses): Treat C++ references the same in
OpenACC as OpenMP.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Set
target_firstprivatize_array_bases in OpenACC parallel and kernels
region contexts. Remove GOMP_MAP_FIRSTPRIVATE_REFERENCE clauses from
OpenACC data regions.
libgomp/
* testsuite/libgomp.oacc-c++/non-scalar-data.C: Remove XFAIL.
Co-Authored-By: Julian Brown <julian@codesourcery.com>
From-SVN: r264244
2018-08-21 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* gfortran.texi: Add description of asynchronous I/O.
* trans-decl.c (gfc_finish_var_decl): Treat asynchronous variables
as volatile.
* trans-io.c (gfc_build_io_library_fndecls): Rename st_wait to
st_wait_async and change argument spec from ".X" to ".w".
(gfc_trans_wait): Pass ID argument via reference.
2018-08-21 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* gfortran.dg/f2003_inquire_1.f03: Add write statement.
* gfortran.dg/f2003_io_1.f03: Add wait statement.
2018-08-21 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* Makefile.am: Add async.c to gfor_io_src.
Add async.h to gfor_io_headers.
* Makefile.in: Regenerated.
* gfortran.map: Add _gfortran_st_wait_async.
* io/async.c: New file.
* io/async.h: New file.
* io/close.c: Include async.h.
(st_close): Call async_wait for an asynchronous unit.
* io/file_pos.c (st_backspace): Likewise.
(st_endfile): Likewise.
(st_rewind): Likewise.
(st_flush): Likewise.
* io/inquire.c: Add handling for asynchronous PENDING
and ID arguments.
* io/io.h (st_parameter_dt): Add async bit.
(st_parameter_wait): Correct.
(gfc_unit): Add au pointer.
(st_wait_async): Add prototype.
(transfer_array_inner): Likewise.
(st_write_done_worker): Likewise.
* io/open.c: Include async.h.
(new_unit): Initialize asynchronous unit.
* io/transfer.c (async_opt): New struct.
(wrap_scalar_transfer): New function.
(transfer_integer): Call wrap_scalar_transfer to do the work.
(transfer_real): Likewise.
(transfer_real_write): Likewise.
(transfer_character): Likewise.
(transfer_character_wide): Likewise.
(transfer_complex): Likewise.
(transfer_array_inner): New function.
(transfer_array): Call transfer_array_inner.
(transfer_derived): Call wrap_scalar_transfer.
(data_transfer_init): Check for asynchronous I/O.
Perform a wait operation on any pending asynchronous I/O
if the data transfer is synchronous. Copy PDT and enqueue
thread for data transfer.
(st_read_done_worker): New function.
(st_read_done): Enqueue transfer or call st_read_done_worker.
(st_write_done_worker): New function.
(st_write_done): Enqueue transfer or call st_read_done_worker.
(st_wait): Document as no-op for compatibility reasons.
(st_wait_async): New function.
* io/unit.c (insert_unit): Use macros LOCK, UNLOCK and TRYLOCK;
add NOTE where necessary.
(get_gfc_unit): Likewise.
(init_units): Likewise.
(close_unit_1): Likewise. Call async_close if asynchronous.
(close_unit): Use macros LOCK and UNLOCK.
(finish_last_advance_record): Likewise.
(newunit_alloc): Likewise.
* io/unix.c (find_file): Likewise.
(flush_all_units_1): Likewise.
(flush_all_units): Likewise.
* libgfortran.h (generate_error_common): Add prototype.
* runtime/error.c: Include io.h and async.h.
(generate_error_common): New function.
2018-08-21 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* testsuite/libgomp.fortran/async_io_1.f90: New test.
* testsuite/libgomp.fortran/async_io_2.f90: New test.
* testsuite/libgomp.fortran/async_io_3.f90: New test.
* testsuite/libgomp.fortran/async_io_4.f90: New test.
* testsuite/libgomp.fortran/async_io_5.f90: New test.
* testsuite/libgomp.fortran/async_io_6.f90: New test.
* testsuite/libgomp.fortran/async_io_7.f90: New test.
Co-Authored-By: Thomas Koenig <tkoenig@gcc.gnu.org>
From-SVN: r263750
The CUDA driver API starting version 6.5 offers a set of runtime functions to
calculate several occupancy-related measures, as a replacement for the occupancy
calculator spreadsheet.
This patch adds a heuristic for default runtime launch geometry, based on the
new runtime function cuOccupancyMaxPotentialBlockSize.
Build on x86_64 with nvptx accelerator and ran libgomp testsuite.
2018-08-13 Cesar Philippidis <cesar@codesourcery.com>
Tom de Vries <tdevries@suse.de>
PR target/85590
* plugin/cuda/cuda.h (CUoccupancyB2DSize): New typedef.
(cuOccupancyMaxPotentialBlockSize): Declare.
* plugin/cuda-lib.def (cuOccupancyMaxPotentialBlockSize): New
CUDA_ONE_CALL_MAYBE_NULL.
* plugin/plugin-nvptx.c (CUDA_VERSION < 6050): Define
CUoccupancyB2DSize and declare
cuOccupancyMaxPotentialBlockSize.
(nvptx_exec): Use cuOccupancyMaxPotentialBlockSize to set the
default num_gangs and num_workers when the driver supports it.
Co-Authored-By: Tom de Vries <tdevries@suse.de>
From-SVN: r263505
Cuda driver api functions cuLinkAddData and cuLinkCreate are available starting
version 5.5. In version 6.5, they are remapped onto _v2 versions.
The dlopen interface of the libgomp nvptx plugin uses the _v2 versions, so it
won't work with a cuda driver with driver api version lower than 6.5.
This patch fixes the problem by testing for the presence of the _v2 versions,
and falling back to the original versions in case of absence of the _v2
versions.
Build on x86_64 with nvptx accelerator and reg-tested libgomp, both with and
without --without-cuda-driver.
2018-08-08 Tom de Vries <tdevries@suse.de>
* plugin/cuda-lib.def (cuLinkAddData_v2, cuLinkCreate_v2): Declare using
CUDA_ONE_CALL_MAYBE_NULL.
* plugin/plugin-nvptx.c (cuLinkAddData, cuLinkCreate): Undef and declare.
(cuLinkAddData_v2, cuLinkCreate_v2): Declare.
(link_ptx): Fall back to cuLinkAddData/cuLinkCreate if the _v2 versions
are not found.
From-SVN: r263408
Cuda driver api function cuGetErrorString is available in version 6.0 and
higher.
Currently, when the driver that is used does not contain this function, the
libgomp nvptx plugin will not build (PLUGIN_NVPTX_DYNAMIC == 0) or run
(PLUGIN_NVPTX_DYNAMIC == 1).
This patch fixes this problem by testing for the presence of the function, and
handling absence.
Build on x86_64 with nvptx accelerator and reg-tested libgomp, both with and
without --without-cuda-driver.
2018-08-08 Tom de Vries <tdevries@suse.de>
* plugin/cuda-lib.def (cuGetErrorString): Use CUDA_ONE_CALL_MAYBE_NULL.
* plugin/plugin-nvptx.c (cuda_error): Handle if cuGetErrorString is not
present.
From-SVN: r263407
CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_MULTIPROCESSOR is defined in cuda driver
api version 6.0 and higher.
Currently nvptx_open_device uses a hard-coded constant instead.
This patch fixes that by:
- defining CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_MULTIPROCESSOR to the hardcoded
constant at toplevel, if not present in cuda.h, and
- using CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_MULTIPROCESSOR in nvptx_open_device
Build on x86_64 with nvptx accelerator and reg-tested libgomp.
2018-08-08 Tom de Vries <tdevries@suse.de>
* plugin/plugin-nvptx.c
(CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_MULTIPROCESSOR): Define.
(nvptx_open_device): Use
CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_MULTIPROCESSOR.
From-SVN: r263406
Cuda driver api function cuGetErrorString is available in version 6.0 and
higher.
This patch:
- removes a comment saying the declaration is not available in cuda.h 6.0
- fixes the presence test to use CUDA_VERSION < 6000
- moves the declaration to toplevel
Build on x86_64 with nvptx accelerator and reg-tested libgomp.
2018-08-08 Tom de Vries <tdevries@suse.de>
* plugin/plugin-nvptx.c (cuda_error): Move declaration of cuGetErrorString ...
(cuGetErrorString): ... here. Guard with CUDA_VERSION < 6000.
From-SVN: r263405
This patch adds handling of functions that may not be present in the cuda
driver.
Such a function can be declared using CUDA_ONE_CALL_MAYBE_NULL in cuda-lib.def,
it can be called with the usual convenience macros, but before calling its
presence needs to be tested using new macro CUDA_CALL_EXISTS.
When using the dlopen interface (PLUGIN_NVPTX_DYNAMIC == 1), we allow
non-present functions by allowing dlsym to return NULL. Otherwise
(PLUGIN_NVPTX_DYNAMIC == 0) we declare the non-present function to be weak.
Build and reg-tested libgomp on x86_64 with nvidia accelerator, with and without
--disable-cuda-driver, in combination with a trigger patch that adds a
non-existing function foo to cuda-lib.def:
...
CUDA_ONE_CALL_MAYBE_NULL (foo)
...
and declares it in plugin-nvptx.c:
...
CUresult foo (void);
...
and then uses it in nvptx_init after the init_cuda_lib call:
...
if (CUDA_CALL_EXISTS (foo))
CUDA_CALL (foo);
...
Also build and reg-tested on x86_64 with nvidia accelerator, with and without
--disable-cuda-driver, in combination with a trigger patch that replaces all
CUDA_ONE_CALLs in cuda-lib.def with CUDA_ONE_CALL_MAYBE_NULL, and guards two
CUDA_CALLs with CUDA_CALL_EXISTS, one for a regular fn, and one for a fn that is
a define in cuda/cuda.h.
2018-08-07 Tom de Vries <tdevries@suse.de>
* plugin/plugin-nvptx.c (DO_PRAGMA): Define.
(struct cuda_lib_s): Add def/undef of CUDA_ONE_CALL_MAYBE_NULL.
(init_cuda_lib): Add new param to CUDA_ONE_CALL_1. Add arg to
corresponding call in CUDA_ONE_CALL. Add def/undef of
CUDA_ONE_CALL_MAYBE_NULL.
(CUDA_CALL_EXISTS): Define.
From-SVN: r263346
This patch makes sure that the lifetimes of the CUDA_ONE_CALL macro (which is
defined twice in plugin-nvptx.c) are minimized, to make it obvious that the
definitions are used only in the lib-cuda.def include.
Build on x86_64 with nvptx accelerator and reg-tested libgomp.
2018-08-07 Tom de Vries <tdevries@suse.de>
* plugin/plugin-nvptx.c (struct cuda_lib_s, init_cuda_lib): Put
CUDA_ONE_CALL defines right before the cuda-lib.def include, and the
corresponding undefs right after.
From-SVN: r263345
Using libgomp configure option --with-cuda-driver=<dir> we can indicate what
cuda driver to use to build the libgomp nvptx plugin. Without such an option,
the system cuda driver is used, if available. If not availabe, a dlopen
interface is used instead.
However, when we use --without-cuda-driver (or the equivalent
--with-cuda-driver=no) the system cuda driver is still used if available.
This patch fixes that, making sure that --without-cuda-driver selects the dlopen
interface.
Build on x86_64 with nvptx accelerator and tested libgomp testsuite, with and
without option --without-cuda-driver.
2018-08-04 Tom de Vries <tdevries@suse.de>
* plugin/configfrag.ac: For --without-cuda-driver, set
CUDA_DRIVER_INCLUDE and CUDA_DRIVER_LIB to no. Handle
CUDA_DRIVER_INCLUDE == no and CUDA_DRIVER_LIB == no.
* configure: Regenerate.
From-SVN: r263310
The nvptx port can't support exceptions using sjlj, because ptx does not
support sjlj. However, default_except_unwind_info still returns UI_SJLJ, even
even if we configure with --disable-sjlj-exceptions, because UI_SJLJ is the
fallback option.
The reason default_except_unwind_info doesn't return UI_DWARF2 is because
DWARF2_UNWIND_INFO is not defined in defaults.h, because
INCOMING_RETURN_ADDR_RTX is not defined, because there's no ptx equivalent.
Testcase libgomp.c++/for-15.C currently doesn't compile unless fno-exceptions
is added because:
- it tries to generate sjlj exception handling code, and
- it tries to generate exception tables using label-addressed .byte sequence.
Ptx doesn't support generating random data at a label, nor being able to
load/write data relative to a label.
This patch fixes the first problem by using UI_TARGET for nvptx.
The second problem is worked around by generating all .byte sequences commented
out. It would be better to have a narrower workaround, and define
TARGET_ASM_BYTE_OP to "error: .byte unsupported " or some such.
This patch does not enable exceptions for nvptx, it merely allows c++ programs
to run correctly if they do no use exception handling.
Build and reg-tested on x86_64 with nvptx accelerator.
2018-08-02 Tom de Vries <tdevries@suse.de>
PR target/86660
* common/config/nvptx/nvptx-common.c (nvptx_except_unwind_info): New
function. Return UI_TARGET unconditionally.
(TARGET_EXCEPT_UNWIND_INFO): Redefine to nvptx_except_unwind_info.
* config/nvptx/nvptx.c (TARGET_ASM_BYTE_OP): Emit commented out '.byte'.
* testsuite/libgomp.oacc-c++/routine-1-auto.C: Remove -fno-exceptions.
* testsuite/libgomp.oacc-c++/routine-1-template-auto.C: Same.
* testsuite/libgomp.oacc-c++/routine-1-template-trailing-return-type.C:
Same.
* testsuite/libgomp.oacc-c++/routine-1-template.C: Same.
* testsuite/libgomp.oacc-c++/routine-1-trailing-return-type.C: Same.
* testsuite/libgomp.oacc-c-c++-common/routine-1.c: Same.
From-SVN: r263265
2018-08-01 Tom de Vries <tdevries@suse.de>
* plugin/cuda-lib.def: New file. Factor out of ...
* plugin/plugin-nvptx.c (CUDA_CALLS): ... here.
(struct cuda_lib_s, init_cuda_lib): Include cuda-lib.def instead of
using CUDA_CALLS.
From-SVN: r263208
2018-07-31 Andre Vieira <andre.simoesdiasvieira@arm.com>
Revert 'AsyncI/O patch committed'
2018-07-25 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* gfortran.texi: Add description of asynchronous I/O.
* trans-decl.c (gfc_finish_var_decl): Treat asynchronous variables
as volatile.
* trans-io.c (gfc_build_io_library_fndecls): Rename st_wait to
st_wait_async and change argument spec from ".X" to ".w".
(gfc_trans_wait): Pass ID argument via reference.
2018-07-31 Andre Vieira <andre.simoesdiasvieira@arm.com>
Revert 'AsyncI/O patch committed'
2018-07-25 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* gfortran.dg/f2003_inquire_1.f03: Add write statement.
* gfortran.dg/f2003_io_1.f03: Add wait statement.
2018-07-31 Andre Vieira <andre.simoesdiasvieira@arm.com>
Revert 'AsyncI/O patch committed'
2018-07-25 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* Makefile.am: Add async.c to gfor_io_src.
Add async.h to gfor_io_headers.
* Makefile.in: Regenerated.
* gfortran.map: Add _gfortran_st_wait_async.
* io/async.c: New file.
* io/async.h: New file.
* io/close.c: Include async.h.
(st_close): Call async_wait for an asynchronous unit.
* io/file_pos.c (st_backspace): Likewise.
(st_endfile): Likewise.
(st_rewind): Likewise.
(st_flush): Likewise.
* io/inquire.c: Add handling for asynchronous PENDING
and ID arguments.
* io/io.h (st_parameter_dt): Add async bit.
(st_parameter_wait): Correct.
(gfc_unit): Add au pointer.
(st_wait_async): Add prototype.
(transfer_array_inner): Likewise.
(st_write_done_worker): Likewise.
* io/open.c: Include async.h.
(new_unit): Initialize asynchronous unit.
* io/transfer.c (async_opt): New struct.
(wrap_scalar_transfer): New function.
(transfer_integer): Call wrap_scalar_transfer to do the work.
(transfer_real): Likewise.
(transfer_real_write): Likewise.
(transfer_character): Likewise.
(transfer_character_wide): Likewise.
(transfer_complex): Likewise.
(transfer_array_inner): New function.
(transfer_array): Call transfer_array_inner.
(transfer_derived): Call wrap_scalar_transfer.
(data_transfer_init): Check for asynchronous I/O.
Perform a wait operation on any pending asynchronous I/O
if the data transfer is synchronous. Copy PDT and enqueue
thread for data transfer.
(st_read_done_worker): New function.
(st_read_done): Enqueue transfer or call st_read_done_worker.
(st_write_done_worker): New function.
(st_write_done): Enqueue transfer or call st_read_done_worker.
(st_wait): Document as no-op for compatibility reasons.
(st_wait_async): New function.
* io/unit.c (insert_unit): Use macros LOCK, UNLOCK and TRYLOCK;
add NOTE where necessary.
(get_gfc_unit): Likewise.
(init_units): Likewise.
(close_unit_1): Likewise. Call async_close if asynchronous.
(close_unit): Use macros LOCK and UNLOCK.
(finish_last_advance_record): Likewise.
(newunit_alloc): Likewise.
* io/unix.c (find_file): Likewise.
(flush_all_units_1): Likewise.
(flush_all_units): Likewise.
* libgfortran.h (generate_error_common): Add prototype.
* runtime/error.c: Include io.h and async.h.
(generate_error_common): New function.
2018-07-31 Andre Vieira <andre.simoesdiasvieira@arm.com>
Revert 'AsyncI/O patch committed'.
2018-07-25 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* testsuite/libgomp.fortran/async_io_1.f90: New test.
* testsuite/libgomp.fortran/async_io_2.f90: New test.
* testsuite/libgomp.fortran/async_io_3.f90: New test.
* testsuite/libgomp.fortran/async_io_4.f90: New test.
* testsuite/libgomp.fortran/async_io_5.f90: New test.
* testsuite/libgomp.fortran/async_io_6.f90: New test.
* testsuite/libgomp.fortran/async_io_7.f90: New test.
From-SVN: r263082
Currently parallel-loop-1.c fails at -O0 on a Quadro M1200, because one of the
kernel launch configurations exceeds the resources available in the device, due
to the default dimensions chosen by the runtime.
This patch fixes that by taking the per-function max_threads_per_block into
account when using the default dimensions.
2018-07-30 Tom de Vries <tdevries@suse.de>
* plugin/plugin-nvptx.c (MIN, MAX): Redefine.
(nvptx_exec): Ensure worker and vector default dims don't exceed
targ_fn->max_threads_per_block.
From-SVN: r263062
The default dimensions are calculated using per-device properties, but
initialized once and used on all devices.
This patch fixes this problem by introducing per-device default dimensions.
2018-07-30 Tom de Vries <tdevries@suse.de>
* plugin/plugin-nvptx.c (struct ptx_device): Add default_dims field.
(nvptx_open_device): Init default_dims for device.
(nvptx_exec): Use default_dims from device.
From-SVN: r263061
PR testsuite/86660
* testsuite/libgomp.c++/for-15.C (results): Include it in
omp declare target region.
(main): Use map (always, tofrom: results) instead of
map (tofrom: results).
From-SVN: r263011
PR middle-end/86660
* omp-low.c (scan_sharing_clauses): Don't ignore map clauses for
declare target to variables if they have always,{to,from,tofrom} map
kinds.
* testsuite/libgomp.c/pr86660.c: New test.
From-SVN: r263010
Currently, when a kernel is lauched with too many workers, it results in a cuda
launch failure. This is triggered f.i. for parallel-loop-1.c at -O0 on a Quadro
M1200.
This patch detects this situation, and errors out with a hint on how to fix it.
Build and reg-tested on x86_64 with nvptx accelerator.
2018-07-26 Cesar Philippidis <cesar@codesourcery.com>
Tom de Vries <tdevries@suse.de>
* plugin/plugin-nvptx.c (nvptx_exec): Error if the hardware doesn't have
sufficient resources to launch a kernel, and give a hint on how to fix
it.
Co-Authored-By: Tom de Vries <tdevries@suse.de>
From-SVN: r262997
Move sampling of device properties from nvptx_exec to nvptx_open, and assume
the sampling always succeeds. This simplifies the default dimension
initialization code in nvptx_open.
2018-07-26 Cesar Philippidis <cesar@codesourcery.com>
Tom de Vries <tdevries@suse.de>
* plugin/plugin-nvptx.c (struct ptx_device): Add warp_size,
max_threads_per_block and max_threads_per_multiprocessor fields.
(nvptx_open_device): Initialize new fields.
(nvptx_exec): Use num_sms, and new fields.
Co-Authored-By: Tom de Vries <tdevries@suse.de>
From-SVN: r262996
In testcase lib-12.f90, all acc_async_test calls are placed in a location
where they are not guaranteed to succeed, which explains why there's an xfail
for the lower optimization levels.
This patch fixes the problem by moving the acc_async_test calls to the correct
locations.
Reg-tested on x86_64 with nvptx accelerator.
2018-07-26 Tom de Vries <tdevries@suse.de>
* testsuite/libgomp.oacc-fortran/lib-12.f90: Move acc_async_test calls
to correct locations. Remove xfail.
From-SVN: r262990
The purpose of the lib-13.f90 test-case is to test acc_wait_all_async. The
test indeed calls acc_wait_all_async, but then subsequentlys calls
acc_wait_all, so the acc_wait_all_async functionality is not tested.
Furthermore, all acc_async_test calls are placed in a location where they are
not guaranteed to succeed, which explains why there's an xfail for the lower
optimization levels.
This patch fixes the problems by replacing acc_wait_all with an acc_wait on
the async id used for the acc_wait_all_async call, and moving the
acc_async_test calls to the correct locations.
Reg-tested on x86_64 with nvptx accelerator.
2018-07-26 Tom de Vries <tdevries@suse.de>
* testsuite/libgomp.oacc-fortran/lib-13.f90: Replace acc_wait_all with
acc_wait. Move acc_async_test calls to correct locations. Remove
xfail.
From-SVN: r262989
2018-07-25 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* gfortran.texi: Add description of asynchronous I/O.
* trans-decl.c (gfc_finish_var_decl): Treat asynchronous variables
as volatile.
* trans-io.c (gfc_build_io_library_fndecls): Rename st_wait to
st_wait_async and change argument spec from ".X" to ".w".
(gfc_trans_wait): Pass ID argument via reference.
2018-07-25 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* gfortran.dg/f2003_inquire_1.f03: Add write statement.
* gfortran.dg/f2003_io_1.f03: Add wait statement.
2018-07-25 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* Makefile.am: Add async.c to gfor_io_src.
Add async.h to gfor_io_headers.
* Makefile.in: Regenerated.
* gfortran.map: Add _gfortran_st_wait_async.
* io/async.c: New file.
* io/async.h: New file.
* io/close.c: Include async.h.
(st_close): Call async_wait for an asynchronous unit.
* io/file_pos.c (st_backspace): Likewise.
(st_endfile): Likewise.
(st_rewind): Likewise.
(st_flush): Likewise.
* io/inquire.c: Add handling for asynchronous PENDING
and ID arguments.
* io/io.h (st_parameter_dt): Add async bit.
(st_parameter_wait): Correct.
(gfc_unit): Add au pointer.
(st_wait_async): Add prototype.
(transfer_array_inner): Likewise.
(st_write_done_worker): Likewise.
* io/open.c: Include async.h.
(new_unit): Initialize asynchronous unit.
* io/transfer.c (async_opt): New struct.
(wrap_scalar_transfer): New function.
(transfer_integer): Call wrap_scalar_transfer to do the work.
(transfer_real): Likewise.
(transfer_real_write): Likewise.
(transfer_character): Likewise.
(transfer_character_wide): Likewise.
(transfer_complex): Likewise.
(transfer_array_inner): New function.
(transfer_array): Call transfer_array_inner.
(transfer_derived): Call wrap_scalar_transfer.
(data_transfer_init): Check for asynchronous I/O.
Perform a wait operation on any pending asynchronous I/O
if the data transfer is synchronous. Copy PDT and enqueue
thread for data transfer.
(st_read_done_worker): New function.
(st_read_done): Enqueue transfer or call st_read_done_worker.
(st_write_done_worker): New function.
(st_write_done): Enqueue transfer or call st_read_done_worker.
(st_wait): Document as no-op for compatibility reasons.
(st_wait_async): New function.
* io/unit.c (insert_unit): Use macros LOCK, UNLOCK and TRYLOCK;
add NOTE where necessary.
(get_gfc_unit): Likewise.
(init_units): Likewise.
(close_unit_1): Likewise. Call async_close if asynchronous.
(close_unit): Use macros LOCK and UNLOCK.
(finish_last_advance_record): Likewise.
(newunit_alloc): Likewise.
* io/unix.c (find_file): Likewise.
(flush_all_units_1): Likewise.
(flush_all_units): Likewise.
* libgfortran.h (generate_error_common): Add prototype.
* runtime/error.c: Include io.h and async.h.
(generate_error_common): New function.
2018-07-25 Nicolas Koenig <koenigni@gcc.gnu.org>
Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/25829
* testsuite/libgomp.fortran/async_io_1.f90: New test.
* testsuite/libgomp.fortran/async_io_2.f90: New test.
* testsuite/libgomp.fortran/async_io_3.f90: New test.
* testsuite/libgomp.fortran/async_io_4.f90: New test.
* testsuite/libgomp.fortran/async_io_5.f90: New test.
* testsuite/libgomp.fortran/async_io_6.f90: New test.
* testsuite/libgomp.fortran/async_io_7.f90: New test.
Co-Authored-By: Thomas Koenig <tkoenig@gcc.gnu.org>
From-SVN: r262978
PR middle-end/86542
* omp-low.c (create_task_copyfn): Copy over also fields corresponding
to _looptemp_ clauses, other than the first two.
* testsuite/libgomp.c++/pr86542.C: New test.
From-SVN: r262815
PR middle-end/86539
* gimplify.c (gimplify_omp_for): Ensure taskloop firstprivatized init
and cond temporaries don't have reference type if iterator has
pointer type. For init use &for_pre_body instead of pre_p if
for_pre_body is non-empty.
* testsuite/libgomp.c++/pr86539.C: New test.
From-SVN: r262776
PR c++/86443
* gimplify.c (find_combined_omp_for): Add DATA argument, in addition
to finding the inner OMP_FOR/OMP_SIMD stmt find non-trivial wrappers,
BLOCKs with BLOCK_VARs, OMP_PARALLEL in between, OMP_FOR in between.
(gimplify_omp_for): For composite loops, move outer
OMP_{DISTRIBUTE,TASKLOOP,FOR,PARALLEL} right around innermost
OMP_FOR/OMP_SIMD if there are any non-trivial wrappers. For class
iterators add any needed clauses. Allow OMP_FOR_ORIG_DECLS to contain
TREE_LIST for both the original class iterator and the "last" helper
var. Gimplify OMP_FOR_PRE_BODY before the outermost composite
loop, remember has_decl_expr from outer composite loops for the
innermost OMP_SIMD in TREE_PRIVATE bit on OMP_FOR_INIT.
gcc/c-family/
* c-omp.c (c_omp_check_loop_iv_r, c_omp_check_loop_iv): Allow declv
to contain TREE_LIST for both the original class iterator and the
"last" helper var.
gcc/cp/
* semantics.c (handle_omp_for_class_iterator): Remove lastp argument,
instead of setting *lastp turn orig_declv elt into a TREE_LIST.
(finish_omp_for): Adjust handle_omp_for_class_iterator caller.
* pt.c (tsubst_omp_for_iterator): Allow OMP_FOR_ORIG_DECLS to contain
TREE_LIST for both the original class iterator and the "last" helper
var.
libgomp/
* testsuite/libgomp.c++/for-15.C: New test.
From-SVN: r262534
2018-05-09 Tom de Vries <tom@codesourcery.com>
PR libgomp/82901
* oacc-parallel.c (GOACC_declare): Use GOMP_ASYNC_SYNC as async argument
to GOACC_enter_exit_data.
From-SVN: r260085
2018-05-07 Tom de Vries <tom@codesourcery.com>
PR testsuite/85677
* testsuite/lib/libgomp.exp (libgomp_init): Move inclusion of top-level
include directory in ALWAYS_CFLAGS out of $blddir != "" condition.
From-SVN: r259992
2018-05-03 Tom de Vries <tom@codesourcery.com>
PR testsuite/85106
* lib/scanoffloadtree.exp: New file.
* testsuite/lib/libgomp-dg.exp (libgomp-dg-test): Add save-temps to
extra_tool_flags if it contains an -foffload=-fdump-* flag.
* testsuite/lib/libgomp.exp: Include scanoffloadtree.exp.
* testsuite/libgomp.oacc-c/vec.c: Use scan-offload-tree-dump.
* doc/sourcebuild.texi (Commands for use in dg-final, Scan optimization
dump files): Add offload-tree.
From-SVN: r259892
2018-05-02 Tom de Vries <tom@codesourcery.com>
PR testsuite/85106
* gcc.dg/ipa/ipa-icf-38.c: Use scan-ltrans-tree-dump.
* lib/scanltranstree.exp: New file.
* lib/target-supports.exp (scan-ltrans-tree-dump_required_options)
(scan-ltrans-tree-dump-times_required_options)
(scan-ltrans-tree-dump-not_required_options)
(scan-ltrans-tree-dump-dem_required_options)
(scan-ltrans-tree-dump-dem-not_required_options): New proc.
* lib/gcc-dg.exp: Include scanltranstree.exp.
* testsuite/lib/libatomic.exp: Include scanltranstree.exp.
* testsuite/lib/libgomp.exp: Include scanltranstree.exp.
* testsuite/lib/libitm.exp: Include scanltranstree.exp.
* testsuite/lib/libvtv.exp: Include scanltranstree.exp.
* doc/sourcebuild.texi (Commands for use in dg-final, Scan optimization
dump files): Add ltrans-tree.
From-SVN: r259838
2018-05-02 Tom de Vries <tom@codesourcery.com>
PR testsuite/85106
* gcc.dg/ipa/ipa-icf-38.c: New test.
* gcc.dg/ipa/ipa-icf-38a.c: New test.
* lib/scandump.exp (dump-base): New proc.
(scan-dump, scan-dump-times, scan-dump-not, scan-dump-dem)
(scan-dump-dem-not): Add and handle parameter for suffix of the dump
base.
* lib/scanipa.exp: Add "" argument to scan-dump calls.
* lib/scanlang.exp: Same.
* lib/scanrtl.exp: Same.
* lib/scantree.exp: Same.
* lib/scanwpaipa.exp: New file.
* lib/gcc-dg.exp: Include scanwpaipa.exp.
* testsuite/lib/libatomic.exp: Include scanwpaipa.exp.
* testsuite/lib/libgomp.exp: Include scanwpaipa.exp.
* testsuite/lib/libitm.exp: Include scanwpaipa.exp.
* testsuite/lib/libvtv.exp: Include scanwpaipa.exp.
* doc/sourcebuild.texi (Commands for use in dg-final, Scan optimization
dump files): Add wpa-ipa.
From-SVN: r259837
2018-04-29 Julian Brown <julian@codesourcery.com>
Tom de Vries <tom@codesourcery.com>
PR testsuite/85527
* testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c: Allow
arbitrary order for iterations of atomic subtract check.
Co-Authored-By: Tom de Vries <tom@codesourcery.com>
From-SVN: r259748
2018-04-28 Tom de Vries <tom@codesourcery.com>
PR testsuite/85527
* testsuite/libgomp.oacc-fortran/atomic_capture-1.f90 (main): Store
atomic capture results obtained in parallel loop to an array, instead of
to a scalar.
From-SVN: r259733
2018-04-26 Richard Biener <rguenther@suse.de>
Tom de Vries <tom@codesourcery.com>
PR lto/85422
* lto-streamer-out.c (output_function): Fixup loops if required to match
discovery done in the reader.
* testsuite/libgomp.oacc-c-c++-common/pr85422.c: New test.
Co-Authored-By: Tom de Vries <tom@codesourcery.com>
From-SVN: r259675
* config/cet.m4 (GCC_CET_FLAGS): Default to --disable-cet, replace
--enable-cet=default with --enable-cet=auto.
* doc/install.texi: Document --disable-cet being the default and
--enable-cet=auto.
* configure: Regenerated.
From-SVN: r259487
PR jit/85384
* acx.m4 (GCC_BASE_VER): Remove \$\$ from sed expression.
* configure.ac (gcc-driver-name.h): Honor --with-gcc-major-version
by using gcc_base_ver to generate a gcc_driver_version, and use
it when generating GCC_DRIVER_NAME.
* configure: Regenerate.
* configure: Regenerate.
From-SVN: r259462
2018-04-16 Cesar Philippidis <cesar@codesourcery.com>
Tom de Vries <tom@codesourcery.com>
PR middle-end/84955
* omp-expand.c (expand_oacc_for): Add dummy false branch for
tiled basic blocks without omp continue statements.
* testsuite/libgomp.oacc-c-c++-common/pr84955.c: New test.
* testsuite/libgomp.oacc-fortran/pr84955.f90: New test.
Co-Authored-By: Tom de Vries <tom@codesourcery.com>
From-SVN: r259406
2018-04-12 Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/83064
PR testsuite/85346
* trans-stmt.c (gfc_trans_forall_loop): Use annot_expr_ivdep_kind
for annotation and remove dependence on -ftree-parallelize-loops.
2018-04-12 Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/83064
PR testsuite/85346
* gfortran.dg/do_concurrent_5.f90: Dynamically allocate main work
array and move test to libgomp/testsuite/libgomp.fortran.
* gfortran.dg/do_concurrent_6.f90: New test.
2018-04-12 Thomas Koenig <tkoenig@gcc.gnu.org>
PR fortran/83064
PR testsuite/85346
* testsuite/libgomp.fortran/do_concurrent_5.f90: Move modified
test from gfortran.dg to here.
From-SVN: r259359
2018-04-05 Tom de Vries <tom@codesourcery.com>
PR target/85204
* config/nvptx/nvptx.c (nvptx_single): Fix neutering of bb with only
cond jump.
* testsuite/libgomp.oacc-c-c++-common/broadcast-1.c: New test.
From-SVN: r259125
2018-03-26 Tom de Vries <tom@codesourcery.com>
PR tree-optimization/85063
* omp-general.c (offloading_function_p): New function. Factor out
of ...
* omp-offload.c (pass_omp_target_link::gate): ... here.
* omp-general.h (offloading_function_p): Declare.
* tree-switch-conversion.c (build_one_array): Mark CSWTCH.x variable
with attribute omp declare target for offloading functions.
* testsuite/libgomp.c/switch-conversion-2.c: New test.
* testsuite/libgomp.c/switch-conversion.c: New test.
* testsuite/libgomp.oacc-c-c++-common/switch-conversion-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/switch-conversion.c: New test.
From-SVN: r258852
ENDBR32 and RDSSPD are multi-byte NOPs on x86-64 processors and
newer x86 processors, starting Pentium Pro. They are UD on older
32-bit processors. Detect this at configure time and adjust the
default value for enable_cet. GCC will enable CET in 32-bit run-time
libraries in any case if --enable-cet is used to configure GCC.
PR target/84148
* config/cet.m4: Check if target support multi-byte NOPS (SSE).
* libatomic/configure: Regenerate.
* libbacktrace/configure: Likewise.
* libgcc/configure: Likewise.
* libgfortran/configure: Likewise.
* libgomp/configure: Likewise.
* libitm/configure: Likewise.
* libmpx/configure: Likewise.
* libobjc/configure: Likewise.
* libquadmath/configure: Likewise.
* libsanitizer/configure: Likewise.
* libssp/configure: Likewise.
* libstdc++-v3/configure: Likewise.
* libvtv/configure: Likewise.
From-SVN: r257809
PR fortran/84418
* trans-openmp.c (gfc_trans_omp_clauses): For OMP_CLAUSE_LINEAR_REF
kind set OMP_CLAUSE_LINEAR_STEP to TYPE_SIZE_UNIT times last_step.
* libgomp.fortran/pr84418-1.f90: New test.
* libgomp.fortran/pr84418-2.f90: New test.
From-SVN: r257771
2018-02-08 Martin Jambor <mjambor@suse.cz>
* hsa-gen.c (get_symbol_for_decl): Set program allocation for
static local variables.
libgomp/
* testsuite/libgomp.hsa.c/staticvar.c: New test.
From-SVN: r257484
2018-02-07 Tom de Vries <tom@codesourcery.com>
PR libgomp/84217
* omp-expand.c (expand_oacc_collapse_init): Ensure diff_type is large
enough.
* c-c++-common/goacc/pr84217.c: New test.
* gfortran.dg/goacc/pr84217.f90: New test.
* testsuite/libgomp.oacc-c-c++-common/pr84217.c: New test.
From-SVN: r257443
2018-01-25 Tom de Vries <tom@codesourcery.com>
PR target/84028
* config/nvptx/nvptx.c (nvptx_single): Add exit insn after noreturn call
for neutered workers.
* testsuite/libgomp.oacc-fortran/pr84028.f90: New test.
From-SVN: r257046
2018-01-24 Tom de Vries <tom@codesourcery.com>
PR target/83589
* config/nvptx/nvptx.c (WORKAROUND_PTXJIT_BUG_2): Define to 1.
(nvptx_pc_set, nvptx_condjump_label): New function. Copy from jump.c.
Add strict parameter.
(prevent_branch_around_nothing): Insert dummy insn between branch to
label and label with no ptx insn inbetween.
* config/nvptx/nvptx.md (define_insn "fake_nop"): New insn.
* testsuite/libgomp.oacc-c-c++-common/pr83589.c: New test.
From-SVN: r257016
2018-01-24 Tom de Vries <tom@codesourcery.com>
PR target/81352
* config/nvptx/nvptx.c (nvptx_single): Add exit insn after noreturn call
for neutered threads in warp.
* config/nvptx/nvptx.md (define_insn "exit"): New insn.
* testsuite/libgomp.oacc-fortran/pr81352.f90: New test.
From-SVN: r257014
2018-01-19 Tom de Vries <tom@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
PR target/83920
* config/nvptx/nvptx.c (nvptx_single): Fix jit workaround.
* testsuite/libgomp.oacc-c-c++-common/pr83920.c: New test.
* testsuite/libgomp.oacc-fortran/pr83920.f90: New test.
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
From-SVN: r256894
2017-12-30 Tom de Vries <tom@codesourcery.com>
PR libgomp/83046
* omp-expand.c (expand_omp_target): If in_lto_p, mark offload_funcs with
DECL_PRESERVE_P.
* lto-streamer-out.c (prune_offload_funcs): New function. Remove
offload_funcs entries that no longer have a corresponding cgraph_node.
Mark the remaining ones as DECL_PRESERVE_P.
(output_lto): Call prune_offload_funcs.
* testsuite/libgomp.oacc-c-c++-common/pr83046.c: New test.
* testsuite/libgomp.c-c++-common/pr83046.c: New test.
From-SVN: r256045
2017-12-27 Tom de Vries <tom@codesourcery.com>
PR c++/83046
* testsuite/libgomp.oacc-c-c++-common/gang-static-2.c (test_static)
(test_nonstatic): Fix return type to workaround PR83046.
From-SVN: r256008
PR fortran/81304
* trans-openmp.c (gfc_trans_omp_array_reduction_or_udr): Set
attr.implicit_type in intrinsic_sym to avoid undesirable warning.
* testsuite/libgomp.fortran/pr81304.f90: New test.
From-SVN: r255144
2017-11-15 Tom de Vries <tom@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c: New test, copied
from asyncwait-1.f90. Rewrite into C. Rewrite from float to int.
* testsuite/libgomp.oacc-c-c++-common/f-asyncwait-2.c: New test, copied
from asyncwait-2.f90. Rewrite into C. Rewrite from float to int.
* testsuite/libgomp.oacc-c-c++-common/f-asyncwait-3.c: New test, copied
from asyncwait-3.f90. Rewrite into C. Rewrite from float to int.
From-SVN: r254769
2017-11-14 Tom de Vries <tom@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Allow to run for
non-nvidia devices.
From-SVN: r254723
PR c++/82835
* cp-gimplify.c (cxx_omp_clause_apply_fn): For methods pass i - 1 to
convert_default_arg instead of i.
* testsuite/libgomp.c++/pr82835.C: New test.
From-SVN: r254511
2017-10-31 Tom de Vries <tom@codesourcery.com>
* plugin/plugin-hsa.c (HSA_LOG): Remove semicolon after
"do {} while (false)".
(init_single_kernel, GOMP_OFFLOAD_async_run): Add missing semicolon
after HSA_DEBUG call.
From-SVN: r254264
* target.c (struct gomp_coalesce_buf): New type.
(MAX_COALESCE_BUF_SIZE, MAX_COALESCE_BUF_GAP): Define.
(gomp_coalesce_buf_add, gomp_to_device_kind_p): New functions.
(gomp_copy_host2dev): Add CBUF argument, if copying into
the cached ranges, memcpy into buffer instead of copying
into device.
(gomp_map_vars_existing, gomp_map_pointer, gomp_map_fields_existing):
Add CBUF argument, pass it through to other calls.
(gomp_map_vars): Aggregate copies from host to device if small enough
and with small enough gaps in between into memcpy into a buffer and
fewer host to device copies from the buffer.
(gomp_update): Adjust gomp_copy_host2dev caller.
From-SVN: r254194
2017-10-09 Martin Jambor <mjambor@suse.cz>
PR hsa/82416
gcc/
* hsa-common.h (hsa_op_with_type): New method extend_int_to_32bit.
* hsa-gen.c (hsa_extend_inttype_to_32bit): New function.
(hsa_type_for_scalar_tree_type): Use it. Always force min32int for
COMPLEX types.
(hsa_fixup_mov_insn_type): New function.
(hsa_op_with_type::get_in_type): Use it.
(hsa_build_append_simple_mov): Likewise. Allow sub-32bit
immediates in an assert.
(hsa_op_with_type::extend_int_to_32bit): New method.
(gen_hsa_insns_for_bitfield): Fixup instruction and intermediary
types. Convert to dest type if necessary.
(gen_hsa_insns_for_bitfield_load): Fixup load type if necessary.
(reg_for_gimple_ssa): Pass false as min32int to
hsa_type_for_scalar_tree_type.
(gen_hsa_addr): Fixup type when creating addresable temporary.
(gen_hsa_cmp_insn_from_gimple): Extend operands if necessary.
(gen_hsa_unary_operation): Extend operands and convert to dest type if
necessary. Call hsa_fixup_mov_insn_type.
(gen_hsa_binary_operation): Changed operand types to hsa_op_with_type,
extend operands and convert to dest type if necessary.
(gen_hsa_insns_for_operation_assignment): Extend operands and convert
to dest type if necessary.
(set_output_in_type): Call hsa_fixup_mov_insn_type. Just ude dest
if conversion nt necessary and size matches.
(gen_hsa_insns_for_load): Call hsa_fixup_mov_insn_type, convert
to dest type if necessary.
(gen_hsa_insns_for_store): Call hsa_fixup_mov_insn_type.
(gen_hsa_insns_for_switch_stmt): Likewise. Also extend operands if
necessary.
(gen_hsa_clrsb): Likewise.
(gen_hsa_ffs): Likewise.
(gen_hsa_divmod): Extend operands and convert to dest type if
necessary.
(gen_hsa_atomic_for_builtin): Change type of op to hsa_op_with_type.
libgomp/
* testsuite/libgomp.hsa.c/pr82416.c: New test.
From-SVN: r253538
2017-10-04 Tom de Vries <tom@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c
(main): Reduce sum of arr elements. Assert that hres is exactly
representable in 32-bit floating point.
* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c
(main): Reduce sum of arr elements. Assert that hres and hmres are
exactly representable in 32-bit floating point.
* testsuite/libgomp.oacc-c-c++-common/reduction-7.c (gwv_np_4): Same.
From-SVN: r253398
