+2019-10-14 Jakub Jelinek <jakub@redhat.com>
+
+ PR libgomp/92081
+ * testsuite/libgomp.fortran/target-simd.f90: Iterate from 1 rather
+ than 0.
+
+2019-10-11 Tobias Burnus <tobias@codesourcery.com>
+
+ * testsuite/libgomp.fortran/use_device_addr-1.f90: New.
+ * testsuite/libgomp.fortran/use_device_addr-2.f90: New.
+
+2019-10-09 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR middle-end/92036
+ * testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: New
+ file.
+
+2019-10-09 Tobias Burnus <tobias@codesourcery.com>
+
+ PR testsuite/91884
+ * testsuite/libgomp.fortran/fortran.exp: Conditionally
+ add -lquadmath.
+ * testsuite/libgomp.oacc-fortran/fortran.exp: Ditto.
+
+2019-10-09 Jakub Jelinek <jakub@redhat.com>
+
+ PR libgomp/92028
+ * target.c (gomp_map_vars_internal): Readd the previous
+ GOMP_MAP_USE_DEVICE_PTR handling code in the first loop,
+ though do that just in the !not_found_cnt case.
+
+2019-10-08 Tobias Burnus <tobias@codesourcery.com>
+
+ * gfortran.dg/gomp/target-simd.f90: New.
+
+2019-10-02 Julian Brown <julian@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * libgomp.h (OFFSET_INLINED, OFFSET_POINTER, OFFSET_STRUCT): Define.
+ * target.c (FIELD_TGT_EMPTY): Define.
+ (gomp_map_val): Use OFFSET_* macros instead of magic constants. Write
+ as switch instead of list of ifs.
+ (gomp_map_vars_internal): Use OFFSET_* and FIELD_TGT_EMPTY macros.
+
+2019-10-02 Andreas Tobler <andreast@gcc.gnu.org>
+
+ * testsuite/libgomp.oacc-c-c++-common/loop-default.h: Remove alloca.h
+ include. Replace alloca () with __builtin_alloca ().
+ * testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c: Likewise.
+
+2019-10-01 Jakub Jelinek <jakub@redhat.com>
+
+ * configure.ac: Remove GCC_HEADER_STDINT(gstdint.h).
+ * libgomp.h: Include <stdint.h> instead of "gstdint.h".
+ * oacc-parallel.c: Don't include "libgomp_g.h".
+ * plugin/plugin-hsa.c: Include <stdint.h> instead of "gstdint.h".
+ * plugin/plugin-nvptx.c: Don't include "gstdint.h".
+ * aclocal.m4: Regenerated.
+ * config.h.in: Regenerated.
+ * configure: Regenerated.
+ * Makefile.in: Regenerated.
+
+2019-09-30 Kwok Cheung Yeung <kcy@codesourcery.com>
+
+ * libgomp_g.h: Include stdint.h instead of gstdint.h.
+
+2019-09-27 Maciej W. Rozycki <macro@wdc.com>
+
+ * configure: Regenerate.
+
+2019-09-13 Tobias Burnus <tobias@codesourcery.com>
+
+ * plugin/plugin-hsa.c (hsa_warn, hsa_fatal, hsa_error): Ensure
+ string is initialized.
+
+2019-09-06 Florian Weimer <fweimer@redhat.com>
+
+ * configure: Regenerate.
+
+2019-09-03 Chung-Lin Tang <cltang@codesourcery.com>
+
+ PR other/79543
+ * acinclude.m4 (LIBGOMP_CHECK_LINKER_FEATURES): Fix GNU ld --version
+ scanning to conform to the GNU Coding Standards.
+ * configure: Regenerate.
+
+2019-08-28 Jakub Jelinek <jakub@redhat.com>
+
+ PR libgomp/91530
+ * testsuite/libgomp.c/scan-21.c: New test.
+ * testsuite/libgomp.c/scan-22.c: New test.
+
+2019-08-27 Jakub Jelinek <jakub@redhat.com>
+
+ PR libgomp/91530
+ * testsuite/libgomp.c/scan-11.c: Add -msse2 option for sse2_runtime
+ targets.
+ * testsuite/libgomp.c/scan-12.c: Likewise.
+ * testsuite/libgomp.c/scan-13.c: Likewise.
+ * testsuite/libgomp.c/scan-14.c: Likewise.
+ * testsuite/libgomp.c/scan-15.c: Likewise.
+ * testsuite/libgomp.c/scan-16.c: Likewise.
+ * testsuite/libgomp.c/scan-17.c: Likewise.
+ * testsuite/libgomp.c/scan-18.c: Likewise.
+ * testsuite/libgomp.c/scan-19.c: Likewise.
+ * testsuite/libgomp.c/scan-20.c: Likewise.
+ * testsuite/libgomp.c++/scan-9.C: Likewise.
+ * testsuite/libgomp.c++/scan-10.C: Likewise.
+ * testsuite/libgomp.c++/scan-11.C: Likewise.
+ * testsuite/libgomp.c++/scan-12.C: Likewise.
+ * testsuite/libgomp.c++/scan-14.C: Likewise.
+ * testsuite/libgomp.c++/scan-15.C: Likewise.
+ * testsuite/libgomp.c++/scan-13.C: Likewise. Use sse2_runtime
+ instead of i?86-*-* x86_64-*-* as target for scan-tree-dump-times.
+ * testsuite/libgomp.c++/scan-16.C: Likewise.
+
+2019-08-17 Thomas Koenig <tkoenig@gcc.gnu.org>
+
+ PR fortran/91473
+ * testsuite/libgomp.fortran/appendix-a/a.28.5.f90: Add
+ -std=legacy so invalid code in the test case is accepted.
+
+2019-08-12 Thomas Koenig <tkoenig@gcc.gnu.org>
+
+ PR fortran/91422
+ * testsuite/libgomp.oacc-fortran/routine-7.f90: Correct array
+ dimension.
+
+2019-08-08 Jakub Jelinek <jakub@redhat.com>
+
+ * target.c (gomp_map_vars_internal): For GOMP_MAP_USE_DEVICE_PTR
+ perform the lookup in the first loop only if !not_found_cnt, otherwise
+ perform lookups for it in the second loop guarded with
+ if (not_found_cnt || has_firstprivate).
+ * testsuite/libgomp.c/target-37.c: New test.
+ * testsuite/libgomp.c++/target-22.C: New test.
+
+2019-08-07 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c/target-18.c (struct S): New type.
+ (foo): Use use_device_addr clause instead of use_device_ptr clause
+ where required by OpenMP 5.0, add further tests for both use_device_ptr
+ and use_device_addr clauses.
+ * testsuite/libgomp.c++/target-9.C (struct S): New type.
+ (foo): Use use_device_addr clause instead of use_device_ptr clause
+ where required by OpenMP 5.0, add further tests for both use_device_ptr
+ and use_device_addr clauses. Add t and u arguments.
+ (main): Adjust caller.
+
+2019-08-06 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c++/loop-13.C: New test.
+ * testsuite/libgomp.c++/loop-14.C: New test.
+ * testsuite/libgomp.c++/loop-15.C: New test.
+
+2019-07-31 Jakub Jelinek <jakub@redhat.com>
+
+ PR middle-end/91301
+ * testsuite/libgomp.c++/for-27.C: New test.
+
+2019-07-23 Steven G. Kargl <kargl@gcc.gnu.org>
+
+ * testsuite/libgomp.fortran/reduction4.f90: Update BOZ usage.
+ * testsuite/libgomp.fortran/reduction5.f90: Ditto.
+
+2019-07-20 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c-c++-common/loop-1.c: New test.
+
+2019-07-08 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c++/scan-13.C: Replace xfail with target x86.
+ * testsuite/libgomp.c++/scan-16.C: Likewise.
+
+2019-07-06 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c/scan-19.c: New test.
+ * testsuite/libgomp.c/scan-20.c: New test.
+
+ * testsuite/libgomp.c/scan-11.c: New test.
+ * testsuite/libgomp.c/scan-12.c: New test.
+ * testsuite/libgomp.c/scan-13.c: New test.
+ * testsuite/libgomp.c/scan-14.c: New test.
+ * testsuite/libgomp.c/scan-15.c: New test.
+ * testsuite/libgomp.c/scan-16.c: New test.
+ * testsuite/libgomp.c/scan-17.c: New test.
+ * testsuite/libgomp.c/scan-18.c: New test.
+ * testsuite/libgomp.c++/scan-9.C: New test.
+ * testsuite/libgomp.c++/scan-10.C: New test.
+ * testsuite/libgomp.c++/scan-11.C: New test.
+ * testsuite/libgomp.c++/scan-12.C: New test.
+ * testsuite/libgomp.c++/scan-13.C: New test.
+ * testsuite/libgomp.c++/scan-14.C: New test.
+ * testsuite/libgomp.c++/scan-15.C: New test.
+ * testsuite/libgomp.c++/scan-16.C: New test.
+
+2019-07-04 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c/scan-9.c: New test.
+ * testsuite/libgomp.c/scan-10.c: New test.
+
+2019-07-03 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c++/scan-1.C: New test.
+ * testsuite/libgomp.c++/scan-2.C: New test.
+ * testsuite/libgomp.c++/scan-3.C: New test.
+ * testsuite/libgomp.c++/scan-4.C: New test.
+ * testsuite/libgomp.c++/scan-5.C: New test.
+ * testsuite/libgomp.c++/scan-6.C: New test.
+ * testsuite/libgomp.c++/scan-7.C: New test.
+ * testsuite/libgomp.c++/scan-8.C: New test.
+ * testsuite/libgomp.c/scan-1.c: New test.
+ * testsuite/libgomp.c/scan-2.c: New test.
+ * testsuite/libgomp.c/scan-3.c: New test.
+ * testsuite/libgomp.c/scan-4.c: New test.
+ * testsuite/libgomp.c/scan-5.c: New test.
+ * testsuite/libgomp.c/scan-6.c: New test.
+ * testsuite/libgomp.c/scan-7.c: New test.
+ * testsuite/libgomp.c/scan-8.c: New test.
+
+2019-06-18 Thomas Schwinge <thomas@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C: New file.
+ * testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c:
+ Likewise.
+
+ * testsuite/libgomp.fortran/allocatable3.f90: Add missing results
+ check.
+
+2019-06-18 Cesar Philippidis <cesar@codesourcery.com>
+
+ * testsuite/libgomp.oacc-fortran/allocatable-array-1.f90: New
+ file.
+
+2019-06-18 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR fortran/90743
+ * oacc-parallel.c (GOACC_parallel_keyed): Handle NULL mapping
+ case.
+ * testsuite/libgomp.fortran/target-allocatable-1-1.f90: New file.
+ * testsuite/libgomp.fortran/target-allocatable-1-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/allocatable-1-1.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/allocatable-1-2.f90: Likewise.
+
+ PR testsuite/90861
+ * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Update.
+
+ PR middle-end/90862
+ * testsuite/libgomp.oacc-c-c++-common/declare-1.c: Update.
+
+2019-06-16 Tom de Vries <tdevries@suse.de>
+
+ PR tree-optimization/89376
+ * testsuite/libgomp.oacc-c-c++-common/pr89376.c: New test.
+
+2019-06-15 Tom de Vries <tdevries@suse.de>
+
+ PR tree-optimization/89713
+ * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Expect no bar.sync.
+ * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Same.
+
+2019-06-15 Jakub Jelinek <jakub@redhat.com>
+
+ PR middle-end/90779
+ * testsuite/libgomp.c/pr90779.c: New test.
+ * testsuite/libgomp.fortran/pr90779.f90: New test.
+
+2019-06-15 Tom de Vries <tdevries@suse.de>
+
+ PR tree-optimization/90009
+ * testsuite/libgomp.oacc-c-c++-common/pr90009.c: New test.
+
+2019-06-13 Feng Xue <fxue@os.amperecomputing.com>
+
+ PR tree-optimization/89713
+ * testsuite/libgomp.oacc-c-c++-common/pr84955-1.c: New test.
+
+2019-06-11 Jakub Jelinek <jakub@redhat.com>
+
+ PR target/90811
+ * testsuite/libgomp.c/pr90811.c: New test.
+
+2019-06-05 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c++/lastprivate-conditional-1.C: New test.
+ * testsuite/libgomp.c++/lastprivate-conditional-2.C: New test.
+
+2019-06-04 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c-c++-common/lastprivate-conditional-7.c: New test.
+ * testsuite/libgomp.c-c++-common/lastprivate-conditional-8.c: New test.
+ * testsuite/libgomp.c-c++-common/lastprivate-conditional-9.c: New test.
+ * testsuite/libgomp.c-c++-common/lastprivate-conditional-10.c: New test.
+
+2019-05-30 Rainer Orth <ro@CeBiTec.Uni-Bielefeld.DE>
+
+ * configure.ac: Call AX_COUNT_CPUS.
+ Substitute CPU_COUNT.
+ * testsuite/Makefile.am (check-am): Use CPU_COUNT as processor
+ count fallback.
+ * aclocal.m4: Regenerate.
+ * configure: Regenerate.
+ * Makefile.in, testsuite/Makefile.in: Regenerate.
+
+2019-05-29 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c-c++-common/lastprivate_conditional_4.c: Rename
+ to ...
+ * testsuite/libgomp.c-c++-common/lastprivate-conditional-4.c: ... this.
+ * testsuite/libgomp.c-c++-common/lastprivate-conditional-5.c: New test.
+ * testsuite/libgomp.c-c++-common/lastprivate-conditional-6.c: New test.
+
+2019-05-27 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c-c++-common/lastprivate_conditional_4.c: New test.
+
+ * testsuite/libgomp.c-c++-common/lastprivate-conditional-3.c: New test.
+
+ PR libgomp/90641
+ * work.c (gomp_init_work_share): Instead of aligning final ordered
+ value to multiples of long long alignment, align to that the
+ first part (ordered team ids) and if inline_ordered_team_ids
+ is not on a long long alignment boundary within the structure,
+ use __alignof__ (long long) - 1 pad size always.
+ * loop.c (GOMP_loop_start): Fix *mem computation if
+ inline_ordered_team_ids is not aligned on long long alignment boundary
+ within the structure.
+ * loop-ull.c (GOMP_loop_ull_start): Likewise.
+ * sections.c (GOMP_sections2_start): Likewise.
+
+2019-05-24 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c-c++-common/lastprivate-conditional-1.c: New test.
+ * testsuite/libgomp.c-c++-common/lastprivate-conditional-2.c: New test.
+
+ PR libgomp/90585
+ * plugin/plugin-hsa.c: Include gstdint.h. Include inttypes.h only if
+ HAVE_INTTYPES_H is defined.
+ (print_uint64_t): New typedef.
+ (PRIu64): Define if HAVE_INTTYPES_H is not defined.
+ (print_kernel_dispatch, run_kernel): Use PRIu64 macro instead of
+ "lu", cast uint64_t HSA_DEBUG and fprintf arguments to print_uint64_t.
+ (release_kernel_dispatch): Likewise. Cast shadow->debug to uintptr_t
+ before casting to void *.
+ * plugin/plugin-nvptx.c: Include gstdint.h instead of stdint.h.
+ * oacc-mem.c: Don't include config.h nor stdint.h.
+ * target.c: Don't include config.h.
+ * oacc-cuda.c: Likewise.
+ * oacc-host.c: Don't include stdint.h.
+
+2019-05-20 Jakub Jelinek <jakub@redhat.com>
+
+ PR libgomp/90527
+ * alloc.c (_GNU_SOURCE): Define.
+
+2019-05-17 Thomas Schwinge <thomas@codesourcery.com>
+
+ * acc_prof.h: New file.
+ * oacc-profiling.c: Likewise.
+ * Makefile.am (nodist_libsubinclude_HEADERS, libgomp_la_SOURCES):
+ Add these, respectively.
+ * Makefile.in: Regenerate.
+ * env.c (initialize_env): Call goacc_profiling_initialize.
+ * oacc-plugin.c (GOMP_PLUGIN_goacc_thread)
+ (GOMP_PLUGIN_goacc_profiling_dispatch): New functions.
+ * oacc-plugin.h (GOMP_PLUGIN_goacc_thread)
+ (GOMP_PLUGIN_goacc_profiling_dispatch): Declare.
+ * libgomp.map (OACC_2.5.1): Add acc_prof_lookup,
+ acc_prof_register, acc_prof_unregister, and acc_register_library.
+ (GOMP_PLUGIN_1.3): Add GOMP_PLUGIN_goacc_profiling_dispatch, and
+ GOMP_PLUGIN_goacc_thread.
+ * oacc-int.h (struct goacc_thread): Add prof_info, api_info,
+ prof_callbacks_enabled members.
+ (goacc_prof_enabled, goacc_profiling_initialize)
+ (_goacc_profiling_dispatch_p, _goacc_profiling_setup_p)
+ (goacc_profiling_dispatch): Declare.
+ (GOACC_PROF_ENABLED, GOACC_PROFILING_DISPATCH_P)
+ (GOACC_PROFILING_SETUP_P): Define.
+ * oacc-async.c (acc_async_test, acc_async_test_all, acc_wait)
+ (acc_wait_async, acc_wait_all, acc_wait_all_async): Update for
+ OpenACC Profiling Interface.
+ * oacc-cuda.c (acc_get_current_cuda_device)
+ (acc_get_current_cuda_context, acc_get_cuda_stream)
+ (acc_set_cuda_stream): Likewise.
+ * oacc-init.c (acc_init_1, goacc_attach_host_thread_to_device)
+ (acc_init, acc_set_device_type, acc_get_device_type)
+ (acc_get_device_num, goacc_lazy_initialize): Likewise.
+ * oacc-mem.c (acc_malloc, acc_free, memcpy_tofrom_device)
+ (acc_deviceptr, acc_hostptr, acc_is_present, acc_map_data)
+ (acc_unmap_data, present_create_copy, delete_copyout)
+ (update_dev_host): Likewise.
+ * oacc-parallel.c (GOACC_parallel_keyed, GOACC_data_start)
+ (GOACC_data_end, GOACC_enter_exit_data, GOACC_update, GOACC_wait):
+ Likewise.
+ * plugin/plugin-nvptx.c (nvptx_exec, nvptx_alloc, nvptx_free)
+ (GOMP_OFFLOAD_openacc_exec, GOMP_OFFLOAD_openacc_async_exec):
+ Likewise.
+ * libgomp.texi: Update.
+ * testsuite/libgomp.oacc-c-c++-common/acc_prof-dispatch-1.c: New
+ file.
+ * testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_prof-valid_bytes-1.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_prof-version-1.c:
+ Likewise.
+
+2019-05-13 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * libgomp-plugin.h (struct goacc_asyncqueue): Declare.
+ (struct goacc_asyncqueue_list): Likewise.
+ (goacc_aq): Likewise.
+ (goacc_aq_list): Likewise.
+ (GOMP_OFFLOAD_openacc_register_async_cleanup): Remove.
+ (GOMP_OFFLOAD_openacc_async_test): Remove.
+ (GOMP_OFFLOAD_openacc_async_test_all): Remove.
+ (GOMP_OFFLOAD_openacc_async_wait): Remove.
+ (GOMP_OFFLOAD_openacc_async_wait_async): Remove.
+ (GOMP_OFFLOAD_openacc_async_wait_all): Remove.
+ (GOMP_OFFLOAD_openacc_async_wait_all_async): Remove.
+ (GOMP_OFFLOAD_openacc_async_set_async): Remove.
+ (GOMP_OFFLOAD_openacc_exec): Adjust declaration.
+ (GOMP_OFFLOAD_openacc_cuda_get_stream): Likewise.
+ (GOMP_OFFLOAD_openacc_cuda_set_stream): Likewise.
+ (GOMP_OFFLOAD_openacc_async_exec): Declare.
+ (GOMP_OFFLOAD_openacc_async_construct): Declare.
+ (GOMP_OFFLOAD_openacc_async_destruct): Declare.
+ (GOMP_OFFLOAD_openacc_async_test): Declare.
+ (GOMP_OFFLOAD_openacc_async_synchronize): Declare.
+ (GOMP_OFFLOAD_openacc_async_serialize): Declare.
+ (GOMP_OFFLOAD_openacc_async_queue_callback): Declare.
+ (GOMP_OFFLOAD_openacc_async_host2dev): Declare.
+ (GOMP_OFFLOAD_openacc_async_dev2host): Declare.
+
+ * libgomp.h (struct acc_dispatch_t): Define 'async' sub-struct.
+ (gomp_acc_insert_pointer): Adjust declaration.
+ (gomp_copy_host2dev): New declaration.
+ (gomp_copy_dev2host): Likewise.
+ (gomp_map_vars_async): Likewise.
+ (gomp_unmap_tgt): Likewise.
+ (gomp_unmap_vars_async): Likewise.
+ (gomp_fini_device): Likewise.
+
+ * oacc-async.c (get_goacc_thread): New function.
+ (get_goacc_thread_device): New function.
+ (lookup_goacc_asyncqueue): New function.
+ (get_goacc_asyncqueue): New function.
+ (acc_async_test): Adjust code to use new async design.
+ (acc_async_test_all): Likewise.
+ (acc_wait): Likewise.
+ (acc_wait_async): Likewise.
+ (acc_wait_all): Likewise.
+ (acc_wait_all_async): Likewise.
+ (goacc_async_free): New function.
+ (goacc_init_asyncqueues): Likewise.
+ (goacc_fini_asyncqueues): Likewise.
+ * oacc-cuda.c (acc_get_cuda_stream): Adjust code to use new async
+ design.
+ (acc_set_cuda_stream): Likewise.
+ * oacc-host.c (host_openacc_exec): Adjust parameters, remove 'async'.
+ (host_openacc_register_async_cleanup): Remove.
+ (host_openacc_async_exec): New function.
+ (host_openacc_async_test): Adjust parameters.
+ (host_openacc_async_test_all): Remove.
+ (host_openacc_async_wait): Remove.
+ (host_openacc_async_wait_async): Remove.
+ (host_openacc_async_wait_all): Remove.
+ (host_openacc_async_wait_all_async): Remove.
+ (host_openacc_async_set_async): Remove.
+ (host_openacc_async_synchronize): New function.
+ (host_openacc_async_serialize): New function.
+ (host_openacc_async_host2dev): New function.
+ (host_openacc_async_dev2host): New function.
+ (host_openacc_async_queue_callback): New function.
+ (host_openacc_async_construct): New function.
+ (host_openacc_async_destruct): New function.
+ (struct gomp_device_descr host_dispatch): Remove initialization of old
+ interface, add intialization of new async sub-struct.
+ * oacc-init.c (acc_shutdown_1): Adjust to use gomp_fini_device.
+ (goacc_attach_host_thread_to_device): Remove old async code usage.
+ * oacc-int.h (goacc_init_asyncqueues): New declaration.
+ (goacc_fini_asyncqueues): Likewise.
+ (goacc_async_copyout_unmap_vars): Likewise.
+ (goacc_async_free): Likewise.
+ (get_goacc_asyncqueue): Likewise.
+ (lookup_goacc_asyncqueue): Likewise.
+ * oacc-mem.c (memcpy_tofrom_device): Adjust code to use new async
+ design.
+ (present_create_copy): Adjust code to use new async design.
+ (delete_copyout): Likewise.
+ (update_dev_host): Likewise.
+ (gomp_acc_insert_pointer): Add async parameter, adjust code to use new
+ async design.
+ (gomp_acc_remove_pointer): Adjust code to use new async design.
+ * oacc-parallel.c (GOACC_parallel_keyed): Adjust code to use new async
+ design.
+ (GOACC_enter_exit_data): Likewise.
+ (goacc_wait): Likewise.
+ (GOACC_update): Likewise.
+ * oacc-plugin.c (GOMP_PLUGIN_async_unmap_vars): Change to assert fail
+ when called, warn as obsolete in comment.
+ * target.c (goacc_device_copy_async): New function.
+ (gomp_copy_host2dev): Remove 'static', add goacc_asyncqueue parameter,
+ add goacc_device_copy_async case.
+ (gomp_copy_dev2host): Likewise.
+ (gomp_map_vars_existing): Add goacc_asyncqueue parameter, adjust code.
+ (gomp_map_pointer): Likewise.
+ (gomp_map_fields_existing): Likewise.
+ (gomp_map_vars_internal): New always_inline function, renamed from
+ gomp_map_vars.
+ (gomp_map_vars): Implement by calling gomp_map_vars_internal.
+ (gomp_map_vars_async): Implement by calling gomp_map_vars_internal,
+ passing goacc_asyncqueue argument.
+ (gomp_unmap_tgt): Remove static, add attribute_hidden.
+ (gomp_unref_tgt): New function.
+ (gomp_unmap_vars_internal): New always_inline function, renamed from
+ gomp_unmap_vars.
+ (gomp_unmap_vars): Implement by calling gomp_unmap_vars_internal.
+ (gomp_unmap_vars_async): Implement by calling
+ gomp_unmap_vars_internal, passing goacc_asyncqueue argument.
+ (gomp_fini_device): New function.
+ (gomp_exit_data): Adjust gomp_copy_dev2host call.
+ (gomp_load_plugin_for_device): Remove old interface, adjust to load
+ new async interface.
+ (gomp_target_fini): Adjust code to call gomp_fini_device.
+
+ * plugin/plugin-nvptx.c (struct cuda_map): Remove.
+ (struct ptx_stream): Remove.
+ (struct nvptx_thread): Remove current_stream field.
+ (cuda_map_create): Remove.
+ (cuda_map_destroy): Remove.
+ (map_init): Remove.
+ (map_fini): Remove.
+ (map_pop): Remove.
+ (map_push): Remove.
+ (struct goacc_asyncqueue): Define.
+ (struct nvptx_callback): Define.
+ (struct ptx_free_block): Define.
+ (struct ptx_device): Remove null_stream, active_streams, async_streams,
+ stream_lock, and next fields.
+ (enum ptx_event_type): Remove.
+ (struct ptx_event): Remove.
+ (ptx_event_lock): Remove.
+ (ptx_events): Remove.
+ (init_streams_for_device): Remove.
+ (fini_streams_for_device): Remove.
+ (select_stream_for_async): Remove.
+ (nvptx_init): Remove ptx_events and ptx_event_lock references.
+ (nvptx_attach_host_thread_to_device): Remove CUDA_ERROR_NOT_PERMITTED
+ case.
+ (nvptx_open_device): Add free_blocks initialization, remove
+ init_streams_for_device call.
+ (nvptx_close_device): Remove fini_streams_for_device call, add
+ free_blocks destruct code.
+ (event_gc): Remove.
+ (event_add): Remove.
+ (nvptx_exec): Adjust parameters and code.
+ (nvptx_free): Likewise.
+ (nvptx_host2dev): Remove.
+ (nvptx_dev2host): Remove.
+ (nvptx_set_async): Remove.
+ (nvptx_async_test): Remove.
+ (nvptx_async_test_all): Remove.
+ (nvptx_wait): Remove.
+ (nvptx_wait_async): Remove.
+ (nvptx_wait_all): Remove.
+ (nvptx_wait_all_async): Remove.
+ (nvptx_get_cuda_stream): Remove.
+ (nvptx_set_cuda_stream): Remove.
+ (GOMP_OFFLOAD_alloc): Adjust code.
+ (GOMP_OFFLOAD_free): Likewise.
+ (GOMP_OFFLOAD_openacc_register_async_cleanup): Remove.
+ (GOMP_OFFLOAD_openacc_exec): Adjust parameters and code.
+ (GOMP_OFFLOAD_openacc_async_test_all): Remove.
+ (GOMP_OFFLOAD_openacc_async_wait): Remove.
+ (GOMP_OFFLOAD_openacc_async_wait_async): Remove.
+ (GOMP_OFFLOAD_openacc_async_wait_all): Remove.
+ (GOMP_OFFLOAD_openacc_async_wait_all_async): Remove.
+ (GOMP_OFFLOAD_openacc_async_set_async): Remove.
+ (cuda_free_argmem): New function.
+ (GOMP_OFFLOAD_openacc_async_exec): New plugin hook function.
+ (GOMP_OFFLOAD_openacc_create_thread_data): Adjust code.
+ (GOMP_OFFLOAD_openacc_cuda_get_stream): Adjust code.
+ (GOMP_OFFLOAD_openacc_cuda_set_stream): Adjust code.
+ (GOMP_OFFLOAD_openacc_async_construct): New plugin hook function.
+ (GOMP_OFFLOAD_openacc_async_destruct): New plugin hook function.
+ (GOMP_OFFLOAD_openacc_async_test): Remove and re-implement.
+ (GOMP_OFFLOAD_openacc_async_synchronize): New plugin hook function.
+ (GOMP_OFFLOAD_openacc_async_serialize): New plugin hook function.
+ (GOMP_OFFLOAD_openacc_async_queue_callback): New plugin hook function.
+ (cuda_callback_wrapper): New function.
+ (cuda_memcpy_sanity_check): New function.
+ (GOMP_OFFLOAD_host2dev): Remove and re-implement.
+ (GOMP_OFFLOAD_dev2host): Remove and re-implement.
+ (GOMP_OFFLOAD_openacc_async_host2dev): New plugin hook function.
+ (GOMP_OFFLOAD_openacc_async_dev2host): New plugin hook function.
+
+2019-05-07 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR target/87835
+ * testsuite/libgomp.oacc-c-c++-common/pr87835.c: Update.
+
+2019-05-06 Thomas Schwinge <thomas@codesourcery.com>
+
+ * oacc-parallel.c: Add comments to legacy entry points (GCC 5).
+
+2019-03-27 Kevin Buettner <kevinb@redhat.com>
+
+ * team.c (gomp_team_start): Initialize pool->threads[0].
+
+2019-02-22 Thomas Schwinge <thomas@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c++/c++.exp: Specify
+ "-foffload=$offload_target".
+ * testsuite/libgomp.oacc-c/c.exp: Likewise.
+ * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
+ * testsuite/lib/libgomp.exp
+ (check_effective_target_openacc_nvidia_accel_configured): Remove,
+ as (conceptually) merged into
+ check_effective_target_openacc_nvidia_accel_selected. Adjust all
+ users.
+
+ * plugin/configfrag.ac: Populate and AC_SUBST offload_targets.
+ * testsuite/libgomp-test-support.exp.in: Adjust.
+ * testsuite/lib/libgomp.exp: Likewise. Don't populate
+ openacc_device_types_s.
+ (offload_target_to_openacc_device_type): New proc.
+ * testsuite/libgomp.oacc-c++/c++.exp: Adjust.
+ * testsuite/libgomp.oacc-c/c.exp: Likewise.
+ * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
+ * Makefile.in: Regenerate.
+ * configure: Likewise.
+ * testsuite/Makefile.in: Likewise.
+
+ * plugin/configfrag.ac: Populate and AC_SUBST offload_plugins
+ instead of offload_targets, and AC_DEFINE_UNQUOTED OFFLOAD_PLUGINS
+ instead of OFFLOAD_TARGETS.
+ * target.c (gomp_target_init): Adjust.
+ * testsuite/libgomp-test-support.exp.in: Likewise.
+ * testsuite/lib/libgomp.exp: Likewise. Populate
+ openacc_device_types_s instead of offload_targets_s_openacc.
+ (check_effective_target_openacc_nvidia_accel_selected)
+ (check_effective_target_openacc_host_selected): Adjust.
+ * testsuite/libgomp.oacc-c++/c++.exp: Likewise.
+ * testsuite/libgomp.oacc-c/c.exp: Likewise.
+ * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
+ * Makefile.in: Regenerate.
+ * config.h.in: Likewise.
+ * configure: Likewise.
+ * testsuite/Makefile.in: Likewise.
+
+ * testsuite/lib/libgomp.exp: Error out for unknown offload target.
+ * testsuite/libgomp.oacc-c++/c++.exp: Likewise. Report if
+ "offloading: supported, but hardware not accessible".
+ * testsuite/libgomp.oacc-c/c.exp: Likewise.
+ * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
+
+2019-02-19 Chung-Lin Tang <cltang@codesourcery.com>
+
+ PR c/87924
+ * oacc-parallel.c (GOACC_parallel_keyed): Remove condition on call to
+ goacc_wait().
+ (goacc_wait): Handle ACC_ASYNC_NOVAL case, remove goacc_thread() call
+ and related adjustment.
+
+2019-01-30 Jakub Jelinek <jakub@redhat.com>
+
+ PR c++/88988
+ * testsuite/libgomp.c++/pr88988.C: New test.
+
+2019-01-28 Jakub Jelinek <jakub@redhat.com>
+
+ PR middle-end/89002
+ * testsuite/libgomp.c/pr89002.c: New test.
+
+2019-01-28 Richard Biener <rguenther@suse.de>
+
+ PR testsuite/89064
+ PR tree-optimization/86865
+ * testsuite/libgomp.graphite/force-parallel-5.c: XFAIL.
+
+2019-01-24 Tom de Vries <tdevries@suse.de>
+
+ * plugin/plugin-nvptx.c (GOMP_OFFLOAD_fini_device): Free ptx_devices
+ once instantiated_devices drops to 0.
+
+2019-01-23 Tom de Vries <tdevries@suse.de>
+
+ PR target/PR88946
+ * plugin/plugin-nvptx.c (cuda_map_destroy): Use CUDA_CALL_NOCHECK for
+ cuMemFree.
+ (nvptx_exec): Don't call map_push if mapnum == 0.
+ * testsuite/libgomp.oacc-c-c++-common/pr88946.c: New test.
+
+2019-01-23 Tom de Vries <tdevries@suse.de>
+
+ PR target/88941
+ PR target/88939
+ * plugin/plugin-nvptx.c (cuda_map_destroy): Handle map->active case.
+ (map_fini): Remove "assert (!s->map->active)".
+ * testsuite/libgomp.oacc-c-c++-common/pr88941.c: New test.
+
+2019-01-23 Tom de Vries <tdevries@suse.de>
+
+ PR target/87835
+ * plugin/plugin-nvptx.c (map_push): Fix adding of allocated element.
+ * testsuite/libgomp.oacc-c-c++-common/pr87835.c: New test.
+
+2019-01-15 Tom de Vries <tdevries@suse.de>
+
+ PR target/80547
+ * testsuite/libgomp.oacc-c-c++-common/gang-reduction-var-assignment.c:
+ New test.
+
+2019-01-12 Tom de Vries <tdevries@suse.de>
+
+ * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c: New test.
+ * testsuite/libgomp.oacc-fortran/gemm-2.f90: New test.
+
+2019-01-12 Tom de Vries <tdevries@suse.de>
+
+ * plugin/plugin-nvptx.c (nvptx_exec): Update error message.
+
+2019-01-12 Tom de Vries <tdevries@suse.de>
+
+ * testsuite/libgomp.oacc-c-c++-common/vector-length-64-1.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/vector-length-64-2.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/vector-length-64-3.c: New test.
+
+2019-01-12 Tom de Vries <tdevries@suse.de>
+
+ PR target/85486
+ * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/pr85486.c: New test.
+
+2019-01-12 Tom de Vries <tdevries@suse.de>
+
+ PR target/85381
+ * testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/pr85381.c: New test.
+
+2019-01-12 Tom de Vries <tdevries@suse.de>
+
+ * testsuite/libgomp.oacc-c-c++-common/vred2d-128.c: New test.
+ * testsuite/libgomp.oacc-fortran/gemm.f90: New test.
+ * testsuite/libgomp.oacc-c-c++-common/vector-length-128-10.c: New test.
+
+2019-01-12 Tom de Vries <tdevries@suse.de>
+
+ * testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c: New test.
+
+2019-01-12 Tom de Vries <tdevries@suse.de>
+
+ * plugin/plugin-nvptx.c (nvptx_exec): Update insufficient hardware
+ resources diagnostic.
+
+2019-01-12 Tom de Vries <tdevries@suse.de>
+
+ * testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Expect
+ vector length to be 128.
+ * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Expect vector
+ length 2097152 to be reduced to 1024 instead of 32.
+
+2019-01-11 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * libgomp.texi: Better distinguish OpenACC and OpenMP "Runtime
+ Library Routines", and "Environment Variables".
+
+2019-01-11 Tom de Vries <tdevries@suse.de>
+
+ * plugin/plugin-nvptx.c (nvptx_exec): Prevent vector_length 64 and
+ num_workers 16.
+
+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.
+
+2019-01-11 Tom de Vries <tdevries@suse.de>
+
+ * testsuite/libgomp.oacc-c-c++-common/insufficient-resources.c: New
+ test.
+
+2019-01-10 Nathan Sidwell <nathan@acm.org>
+ Julian Brown <julian@codesourcery.com>
+
+ PR lto/71959
+ * testsuite/libgomp.oacc-c++/pr71959-aux.cc: New.
+ * testsuite/libgomp.oacc-c++/pr71959.C: New.
+
+2019-01-09 Sebastian Huber <sebastian.huber@embedded-brains.de>
+
+ * config/rtems/bar.c: Include "../linux/bar.c" and delete copy
+ and paste code.
+
+2019-01-09 Sebastian Huber <sebastian.huber@embedded-brains.de>
+
+ * config/rtems/affinity-fmt.c: New file. Include affinity-fmt.c,
+ undefining HAVE_GETPID and HAVE_GETHOSTNAME, and mapping fwrite to
+ write.
+
+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.
+
+2019-01-09 Tom de Vries <tdevries@suse.de>
+
+ * plugin/plugin-nvptx.c (nvptx_exec): Make sure to launch with at least
+ one worker.
+
2019-01-07 Tom de Vries <tdevries@suse.de>
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: Fix
projects / thirdparty / gcc.git / blobdiff