2025-05-14 Tobias Burnus Backported from master: 2025-05-14 Tobias Burnus * target.c (gomp_attach_pointer): Return bool; accept additional bool to optionally silence the fatal pointee-not-found error. (gomp_map_vars_internal): If the pointee could not be found, check whether it was mapped as GOMP_MAP_ZERO_LEN_ARRAY_SECTION. * libgomp.h (gomp_attach_pointer): Update prototype. * oacc-mem.c (acc_attach_async, goacc_enter_data_internal): Update calls. * testsuite/libgomp.c/target-map-zero-sized.c: New test. * testsuite/libgomp.c/target-map-zero-sized-2.c: New test. * testsuite/libgomp.c/target-map-zero-sized-3.c: New test. 2025-04-25 Thomas Schwinge Backported from trunk: 2025-04-25 Thomas Schwinge PR target/119853 PR target/119854 * target-cxa-dso-dtor.c: New. * config/accel/target-cxa-dso-dtor.c: Likewise. * Makefile.am (libgomp_la_SOURCES): Add it. * Makefile.in: Regenerate. * testsuite/libgomp.c++/target-cdtor-1.C: New. * testsuite/libgomp.c++/target-cdtor-2.C: Likewise. Backported from trunk: 2025-04-25 Thomas Schwinge * testsuite/libgomp.c-c++-common/target-cdtor-1.c: New. Backported from trunk: 2025-04-25 Andrew Pinski Thomas Schwinge PR target/119737 * testsuite/libgomp.c++/target-exceptions-throw-1.C: Remove PR119737 XFAILing. * testsuite/libgomp.c++/target-exceptions-throw-2.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-throw-1.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-throw-2.C: Likewise. Backported from trunk: 2025-04-25 Thomas Schwinge PR target/118794 * testsuite/libgomp.c++/target-exceptions-pr118794-1.C: Adjust for 'targetm.arm_eabi_unwinder'. * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C: Likewise. Backported from trunk: 2024-08-09 Thomas Schwinge * testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: Fix effective-target keyword. * testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: Likewise. * testsuite/libgomp.c-c++-common/target-is-initial-host-2.c: Likewise. * testsuite/libgomp.c-c++-common/target-is-initial-host.c: Likewise. * testsuite/libgomp.fortran/target-is-initial-host-2.f90: Likewise. * testsuite/libgomp.fortran/target-is-initial-host.f: Likewise. * testsuite/libgomp.fortran/target-is-initial-host.f90: Likewise. Backported from trunk: 2024-08-08 Tobias Burnus * libgomp.texi (omp_is_initial_device): Mention -fno-builtin-omp_is_initial_device and folding by default. Backported from trunk: 2024-08-08 Tobias Burnus * testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: Split scan-tree-dump into with and without target offload_target_any. * testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: Likewise. Backported from trunk: 2024-08-07 Julian Brown Tobias Burnus * testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C: New test. * testsuite/libgomp.c-c++-common/target-is-initial-host.c: New test. * testsuite/libgomp.c-c++-common/target-is-initial-host-2.c: New test. * testsuite/libgomp.fortran/target-is-initial-host.f: New test. * testsuite/libgomp.fortran/target-is-initial-host.f90: New test. * testsuite/libgomp.fortran/target-is-initial-host-2.f90: New test. Revert: 2023-05-12 Julian Brown * testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C: New test. 2025-04-24 Tobias Burnus Backported from master: 2025-04-24 Tobias Burnus * testsuite/lib/libgomp.exp (check_effective_target_gomp_hip_header_nvidia): Compile with "-Wno-deprecated-declarations". * testsuite/libgomp.c/interop-hip-nvidia-full.c: Likewise. * testsuite/libgomp.c/interop-hipblas-nvidia-full.c: Likewise. * testsuite/libgomp.c/interop-hipblas.h: Add workarounds when using the HIP headers with __HIP_PLATFORM_NVIDIA__. 2025-04-24 Tobias Burnus Backported from master: 2025-04-24 Tobias Burnus * testsuite/lib/libgomp.exp (check_effective_target_openacc_cublas, check_effective_target_openacc_cudart): Update description as the check requires more. (check_effective_target_openacc_libcuda, check_effective_target_openacc_libcublas, check_effective_target_openacc_libcudart, check_effective_target_gomp_hip_header_amd, check_effective_target_gomp_hip_header_nvidia, check_effective_target_gomp_hipfort_module, check_effective_target_gomp_libamdhip64, check_effective_target_gomp_libhipblas): New. * testsuite/libgomp.c-c++-common/interop-2.c: New test. * testsuite/libgomp.c/interop-cublas-full.c: New test. * testsuite/libgomp.c/interop-cublas-libonly.c: New test. * testsuite/libgomp.c/interop-cuda-full.c: New test. * testsuite/libgomp.c/interop-cuda-libonly.c: New test. * testsuite/libgomp.c/interop-hip-amd-full.c: New test. * testsuite/libgomp.c/interop-hip-amd-no-hip-header.c: New test. * testsuite/libgomp.c/interop-hip-nvidia-full.c: New test. * testsuite/libgomp.c/interop-hip-nvidia-no-headers.c: New test. * testsuite/libgomp.c/interop-hip-nvidia-no-hip-header.c: New test. * testsuite/libgomp.c/interop-hip.h: New test. * testsuite/libgomp.c/interop-hipblas-amd-full.c: New test. * testsuite/libgomp.c/interop-hipblas-amd-no-hip-header.c: New test. * testsuite/libgomp.c/interop-hipblas-nvidia-full.c: New test. * testsuite/libgomp.c/interop-hipblas-nvidia-no-headers.c: New test. * testsuite/libgomp.c/interop-hipblas-nvidia-no-hip-header.c: New test. * testsuite/libgomp.c/interop-hipblas.h: New test. * testsuite/libgomp.fortran/interop-hip-amd-full.F90: New test. * testsuite/libgomp.fortran/interop-hip-amd-no-module.F90: New test. * testsuite/libgomp.fortran/interop-hip-nvidia-full.F90: New test. * testsuite/libgomp.fortran/interop-hip-nvidia-no-module.F90: New test. * testsuite/libgomp.fortran/interop-hip.h: New test. 2025-04-17 Kwok Cheung Yeung * testsuite/libgomp.fortran/allocatable-comp-iterators.f90: Add test for non-const iterator boundaries. 2025-04-17 Kwok Cheung Yeung * testsuite/libgomp.fortran/allocatable-comp-iterators.f90: New. 2025-04-17 Kwok Cheung Yeung * testsuite/libgomp.fortran/mapper-iterators-1.f90: New test. * testsuite/libgomp.fortran/mapper-iterators-2.f90: New test. * testsuite/libgomp.fortran/mapper-iterators-3.f90: New test. * testsuite/libgomp.fortran/mapper-iterators-4.f90: New test. 2025-04-17 Kwok Cheung Yeung * testsuite/libgomp.c-c++-common/mapper-iterators-1.c: New test. * testsuite/libgomp.c-c++-common/mapper-iterators-2.c: New test. * testsuite/libgomp.c-c++-common/mapper-iterators-3.c: New test. 2025-04-17 Kwok Cheung Yeung * testsuite/libgomp.c-c++-common/target-map-iterators-4.c: New. * testsuite/libgomp.c-c++-common/target-map-iterators-5.c: New. * testsuite/libgomp.c-c++-common/target-update-iterators-4.c: New. * testsuite/libgomp.fortran/target-map-iterators-4.f90: New. * testsuite/libgomp.fortran/target-map-iterators-5.f90: New. * testsuite/libgomp.fortran/target-update-iterators-4.f90: New. 2025-04-17 Kwok Cheung Yeung * testsuite/libgomp.fortran/target-update-iterators-1.f90: New. * testsuite/libgomp.fortran/target-update-iterators-2.f90: New. * testsuite/libgomp.fortran/target-update-iterators-3.f90: New. 2025-04-17 Kwok Cheung Yeung * target.c (kind_to_name): Handle GOMP_MAP_STRUCT and GOMP_MAP_STRUCT_UNORD. (gomp_add_map): New. (gomp_merge_iterator_maps): Expand fields of a struct mapping breadth-first. * testsuite/libgomp.fortran/target-map-iterators-1.f90: New. * testsuite/libgomp.fortran/target-map-iterators-2.f90: New. * testsuite/libgomp.fortran/target-map-iterators-3.f90: New. 2025-04-17 Kwok Cheung Yeung * target.c (gomp_update): Call gomp_merge_iterator_maps. Free allocated variables. * testsuite/libgomp.c-c++-common/target-update-iterators-1.c: New. * testsuite/libgomp.c-c++-common/target-update-iterators-2.c: New. * testsuite/libgomp.c-c++-common/target-update-iterators-3.c: New. 2025-04-17 Kwok Cheung Yeung * target.c (kind_to_name): New. (gomp_merge_iterator_maps): New. (gomp_map_vars_internal): Call gomp_merge_iterator_maps. Copy address of only the first iteration to target vars. Free allocated variables. * testsuite/libgomp.c-c++-common/target-map-iterators-1.c: New. * testsuite/libgomp.c-c++-common/target-map-iterators-2.c: New. * testsuite/libgomp.c-c++-common/target-map-iterators-3.c: New. 2025-04-17 Thomas Schwinge * testsuite/libgomp.oacc-c++/exceptions-bad_cast-3.C: Adjust. * testsuite/libgomp.oacc-c++/exceptions-throw-3.C: Likewise. * testsuite/libgomp.oacc-c++/pr119692-1-1.C: Likewise. * testsuite/libgomp.oacc-c++/pr119692-1-2.C: Likewise. * testsuite/libgomp.oacc-c++/pr119692-1-3.C: Likewise. Backported from trunk: 2025-04-16 Thomas Schwinge * testsuite/libgomp.c++/target-exceptions-pr118794-1.C: Remove 'ALWAYS_INLINE' workaround. Backported from trunk: 2025-04-16 Thomas Schwinge PR target/106445 * testsuite/libgomp.c++/pr106445-1.C: New. * testsuite/libgomp.c++/pr106445-1-O0.C: Likewise. Backported from trunk: 2025-04-16 Thomas Schwinge PR target/97106 * testsuite/libgomp.c++/pr96390.C: Un-XFAIL nvptx offloading. * testsuite/libgomp.c-c++-common/pr96390.c: Adjust. Backported from trunk: 2025-04-14 Thomas Schwinge PR target/118794 * testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C: Set '-foffload-options=-mno-fake-exceptions'. * testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C: Likewise. * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C: Likewise. * testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C: Likewise. * testsuite/libgomp.c++/target-exceptions-bad_cast-2.C: Adjust. * testsuite/libgomp.c++/target-exceptions-pr118794-1.C: Likewise. * testsuite/libgomp.c++/target-exceptions-throw-2.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-throw-2.C: Likewise. * testsuite/libgomp.c++/target-exceptions-throw-2-O0.C: New. Backported from trunk: 2025-04-14 Thomas Schwinge * testsuite/libgomp.c++/target-exceptions-throw-3.C: New. * testsuite/libgomp.oacc-c++/exceptions-throw-3.C: Likewise. Backported from trunk: 2025-04-14 Thomas Schwinge * testsuite/libgomp.c++/target-exceptions-throw-2.C: New. * testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-throw-2.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C: Likewise. Backported from trunk: 2025-04-14 Thomas Schwinge * testsuite/libgomp.c++/target-exceptions-throw-1.C: New. * testsuite/libgomp.c++/target-exceptions-throw-1-O0.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-throw-1.C: Likewise. Backported from trunk: 2025-04-14 Thomas Schwinge * testsuite/libgomp.c++/target-exceptions-bad_cast-3.C: New. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-3.C: Likewise. Backported from trunk: 2025-04-14 Thomas Schwinge * testsuite/libgomp.c++/target-exceptions-bad_cast-2.C: New. * testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C: Likewise. Backported from trunk: 2025-04-14 Thomas Schwinge * testsuite/libgomp.c++/target-exceptions-bad_cast-1.C: New. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-1.C: Likewise. Backported from trunk: 2025-04-14 Thomas Schwinge PR target/118794 * testsuite/libgomp.c++/target-exceptions-pr118794-1.C: New. * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C: Likewise. Backported from trunk: 2025-04-14 Thomas Schwinge PR c++/119692 * testsuite/libgomp.c++/pr119692-1-1.C: New. * testsuite/libgomp.c++/pr119692-1-2.C: Likewise. * testsuite/libgomp.c++/pr119692-1-3.C: Likewise. * testsuite/libgomp.c++/pr119692-1-4.C: Likewise. * testsuite/libgomp.c++/pr119692-1-5.C: Likewise. * testsuite/libgomp.oacc-c++/pr119692-1-1.C: Likewise. * testsuite/libgomp.oacc-c++/pr119692-1-2.C: Likewise. * testsuite/libgomp.oacc-c++/pr119692-1-3.C: Likewise. 2025-04-08 Tobias Burnus Backported from master: 2025-04-08 Tobias Burnus PR middle-end/119662 * testsuite/libgomp.c/append-args-fr-1.c: New test. * testsuite/libgomp.c/append-args-fr.h: New test. 2025-03-27 Thomas Schwinge Backported from trunk: 2025-03-26 Thomas Schwinge PR driver/101544 * testsuite/libgomp.c++/pr101544-1-O0.C: Remove '-foffload-options=-lstdc++'. * testsuite/libgomp.c++/pr101544-1.C: Likewise. * testsuite/libgomp.oacc-c++/pr101544-1.C: Likewise. 2025-03-25 Thomas Schwinge Backported from trunk: 2025-03-24 Thomas Schwinge PR libgomp/96835 * testsuite/libgomp.c++/pr96835-1.C: New. * testsuite/libgomp.c++/pr96835-1-O0.C: Likewise. * testsuite/libgomp.oacc-c++/pr96835-1.C: Likewise. Backported from trunk: 2025-03-24 Thomas Schwinge PR target/101544 * testsuite/libgomp.c++/pr101544-1.C: New. * testsuite/libgomp.c++/pr101544-1-O0.C: Likewise. * testsuite/libgomp.oacc-c++/pr101544-1.C: Likewise. 2025-03-24 Tobias Burnus Backported from master: 2025-03-24 Tobias Burnus * target.c (gomp_interop_internal): Set the 'device_num' member when initializing an interop object. 2025-03-24 Tobias Burnus Backported from master: 2025-03-24 Tobias Burnus * plugin/plugin-nvptx.c (GOMP_OFFLOAD_interop): Set context for stream creation to use the specified device. 2025-03-21 Tobias Burnus Backported from master: 2025-03-21 Tobias Burnus * testsuite/libgomp.fortran/get-mapped-ptr-1.f90: Use -6 not -5 as non-conforming device number. 2025-03-21 Tobias Burnus Backported from master: 2025-03-21 Tobias Burnus * plugin/plugin-gcn.c (_LIBGOMP_PLUGIN_INCLUDE): Define. (struct hsa_runtime_fn_info): Add two queue functions. (hipError_t, hipCtx_t, hipStream_s, hipStream_t): New types. (struct hip_runtime_fn_info): New. (hip_runtime_lib, hip_fns): New global vars. (init_environment_variables): Handle hip_runtime_lib. (init_hsa_runtime_functions): Load the two queue functions. (init_hip_runtime_functions, GOMP_OFFLOAD_interop, GOMP_OFFLOAD_get_interop_int, GOMP_OFFLOAD_get_interop_ptr, GOMP_OFFLOAD_get_interop_str, GOMP_OFFLOAD_get_interop_type_desc): New. * plugin/plugin-nvptx.c (_LIBGOMP_PLUGIN_INCLUDE): Define. (GOMP_OFFLOAD_interop, GOMP_OFFLOAD_get_interop_int, GOMP_OFFLOAD_get_interop_ptr, GOMP_OFFLOAD_get_interop_str, GOMP_OFFLOAD_get_interop_type_desc): New. * testsuite/libgomp.c/interop-fr-1.c: New test. * testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c: Use -6 not -5 as non-conforming device number. 2025-03-21 Paul-Antoine Arras Backported from master: 2025-03-21 Paul-Antoine Arras Tobias Burnus * icv-device.c (omp_set_default_device): Check GOMP_DEVICE_DEFAULT_OMP_61. * libgomp-plugin.h (struct interop_obj_t): New. (enum gomp_interop_flag): New. (GOMP_OFFLOAD_interop): Declare. (GOMP_OFFLOAD_get_interop_int): Declare. (GOMP_OFFLOAD_get_interop_ptr): Declare. (GOMP_OFFLOAD_get_interop_str): Declare. (GOMP_OFFLOAD_get_interop_type_desc): Declare. * libgomp.h (_LIBGOMP_OMP_LOCK_DEFINED): Define. (struct gomp_device_descr): Add interop_func, get_interop_int_func, get_interop_ptr_func, get_interop_str_func, get_interop_type_desc_func. * libgomp.map: Add GOMP_interop. * libgomp_g.h (GOMP_interop): Declare. * target.c (resolve_device): Handle GOMP_DEVICE_DEFAULT_OMP_61. (omp_get_interop_int): Replace stub with actual implementation. (omp_get_interop_ptr): Likewise. (omp_get_interop_str): Likewise. (omp_get_interop_type_desc): Likewise. (struct interop_data_t): Define. (gomp_interop_internal): New function. (GOMP_interop): Likewise. (gomp_load_plugin_for_device): Load symbols for get_interop_int, get_interop_ptr, get_interop_str and get_interop_type_desc. * testsuite/libgomp.c-c++-common/interop-1.c: New test. 2025-03-18 Tobias Burnus Backported from master: 2025-03-17 Tobias Burnus PR fortran/115271 * testsuite/libgomp.fortran/declare-variant-mod-1-use.f90: Moved from gcc/testsuite/gfortran.dg/gomp/. * testsuite/libgomp.fortran/declare-variant-mod-1.f90: Likewise. 2025-01-30 Tobias Burnus Backported from master: 2025-01-30 Tobias Burnus * libgomp.texi (Impl. Status): Update for accumpulated changes related to 'dispatch' and interop. 2025-01-27 Tobias Burnus Backported from master: 2024-12-19 Tobias Burnus * libgomp.texi (OpenMP Context Selectors): Document that 'kind' also accepts 'cpu'/'any' on host and 'any'/'nohost' on 'nohost' devices. 2025-01-27 Paul-Antoine Arras Backported from master: 2025-01-13 Paul-Antoine Arras * testsuite/libgomp.fortran/dispatch-1.f90: Add missing target directive. 2025-01-27 Paul-Antoine Arras Backported from master: 2025-01-03 Paul-Antoine Arras * libgomp.texi: Update dispatch and adjust_args status. 2025-01-27 Paul-Antoine Arras Backported from master: 2025-01-02 Paul-Antoine Arras * testsuite/libgomp.fortran/declare-variant-2-aux.f90: New test. * testsuite/libgomp.fortran/declare-variant-2.f90: New test (xfail). * testsuite/libgomp.fortran/dispatch-1.f90: New test. * testsuite/libgomp.fortran/dispatch-2.f90: New test. * testsuite/libgomp.fortran/dispatch-3.f90: New test. 2025-01-27 Paul-Antoine Arras Backported from master: 2024-11-20 Paul-Antoine Arras * testsuite/libgomp.c-c++-common/dispatch-1.c: New test. * testsuite/libgomp.c-c++-common/dispatch-2.c: New test. 2025-01-23 Tobias Burnus Backported from master: 2024-11-11 Tobias Burnus * testsuite/libgomp.c-c++-common/pr109062.c: Update dg-output to also accept GOMP_SPINCOUNT = 1 for x86-64. 2025-01-23 Tobias Burnus Backported from master: 2024-11-07 Tobias Burnus * libgomp.texi (OpenMP Technical Report 13): Remove 'iterator' in 'map' clause of 'declare mapper' as it is already the list above. (Interoperability Routines): Add. (omp_target_memcpy_async, omp_target_memcpy_rect_async): Document that depobj_list may be omitted in C++ and Fortran. 2025-01-23 Tobias Burnus Backported from master: 2024-09-23 Tobias Burnus * fortran.c (omp_get_device_from_uid_): New function. * libgomp.map (GOMP_6.0): Add it. * oacc-host.c (host_dispatch): Init '.uid' and '.get_uid_func'. * omp_lib.f90.in: Make it used by removing bind(C). * omp_lib.h.in: Likewise. * target.c (omp_get_device_from_uid): Ensure the device is initialized. * plugin/plugin-gcn.c (GOMP_OFFLOAD_get_uid): Add function comment; return NULL in case of an error. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_uid): Likewise. * testsuite/libgomp.fortran/device_uid.f90: Update to test substrings. 2025-01-23 Tobias Burnus Backported from master: 2024-09-20 Tobias Burnus * config/gcn/target.c (omp_get_uid_from_device, omp_get_device_from_uid): Add stub implementation. * config/nvptx/target.c (omp_get_uid_from_device, omp_get_device_from_uid): Likewise. * fortran.c (omp_get_uid_from_device_, omp_get_uid_from_device_8_): New functions. * libgomp-plugin.h (GOMP_OFFLOAD_get_uid): Add prototype. * libgomp.h (struct gomp_device_descr): Add 'uid' and 'get_uid_func'. * libgomp.map (GOMP_6.0): New, includind the new UID routines. * libgomp.texi (OpenMP Technical Report 13): Mark UID routines as 'Y'. (Device Information Routines): Document new UID routines. (Offload-Target Specifics): Document UID format. * omp.h.in (omp_get_device_from_uid, omp_get_uid_from_device): New prototype. * omp_lib.f90.in (omp_get_device_from_uid, omp_get_uid_from_device): New interface. * omp_lib.h.in: Likewise. * plugin/cuda-lib.def: Add cuDeviceGetUuid and cuDeviceGetUuid_v2 via CUDA_ONE_CALL_MAYBE_NULL. * plugin/plugin-gcn.c (GOMP_OFFLOAD_get_uid): New. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_uid): New. * target.c (str_omp_initial_device): New static var. (STR_OMP_DEV_PREFIX): Define. (gomp_get_uid_for_device, omp_get_uid_from_device, omp_get_device_from_uid): New. (gomp_load_plugin_for_device): DLSYM_OPT the function 'get_uid'. (gomp_target_init): Set the device's 'uid' field to NULL. * testsuite/libgomp.c/device_uid.c: New test. * testsuite/libgomp.fortran/device_uid.f90: New test. 2025-01-23 Tobias Burnus Backported from master: 2024-08-28 Tobias Burnus * fortran.c (omp_get_interop_str_, omp_get_interop_name_, omp_get_interop_type_desc_, omp_get_interop_rc_desc_): Add. * libgomp.map (GOMP_5.1.3): New; add interop routines. * omp.h.in: Add interop typedefs, enum and prototypes. (__GOMP_DEFAULT_NULL): Define. (omp_target_memcpy_async, omp_target_memcpy_rect_async): Use it for the optional depend argument. * omp_lib.f90.in: Add paramters and interfaces for interop. * omp_lib.h.in: Likewise; move F90 '&' to column 81 for -ffree-length-80. * target.c (omp_get_num_interop_properties, omp_get_interop_int, omp_get_interop_ptr, omp_get_interop_str, omp_get_interop_name, omp_get_interop_type_desc, omp_get_interop_rc_desc): Add. * config/gcn/target.c (omp_get_num_interop_properties, omp_get_interop_int, omp_get_interop_ptr, omp_get_interop_str, omp_get_interop_name, omp_get_interop_type_desc, omp_get_interop_rc_desc): Add. * config/nvptx/target.c (omp_get_num_interop_properties, omp_get_interop_int, omp_get_interop_ptr, omp_get_interop_str, omp_get_interop_name, omp_get_interop_type_desc, omp_get_interop_rc_desc): Add. * testsuite/libgomp.c-c++-common/interop-routines-1.c: New test. * testsuite/libgomp.c-c++-common/interop-routines-2.c: New test. * testsuite/libgomp.fortran/interop-routines-1.F90: New test. * testsuite/libgomp.fortran/interop-routines-2.F90: New test. * testsuite/libgomp.fortran/interop-routines-3.F: New test. * testsuite/libgomp.fortran/interop-routines-4.F: New test. * testsuite/libgomp.fortran/interop-routines-5.F: New test. * testsuite/libgomp.fortran/interop-routines-6.F: New test. * testsuite/libgomp.fortran/interop-routines-7.F90: New test. 2024-12-19 Thomas Schwinge PR target/65181 * testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Adjust. 2024-12-18 Thomas Schwinge Backported from trunk: 2024-12-06 Thomas Schwinge * testsuite/libgomp.c/declare-variant-3-sm89.c: New. * testsuite/libgomp.c/declare-variant-3.h: Adjust. Backported from trunk: 2024-12-06 Thomas Schwinge * testsuite/libgomp.c/declare-variant-3-sm52.c: New. * testsuite/libgomp.c/declare-variant-3.h: Adjust. Backported from trunk: 2024-12-06 Thomas Schwinge * testsuite/libgomp.c/declare-variant-3-sm37.c: New. * testsuite/libgomp.c/declare-variant-3.h: Adjust. Backported from trunk: 2024-12-10 Tobias Burnus * plugin/plugin-gcn.c (GOMP_OFFLOAD_dev2dev, GOMP_OFFLOAD_async_run): Handle omp_async_queue == NULL after call to maybe_init_omp_async. (GOMP_OFFLOAD_openacc_async_construct): Use error not fatal error, partially reverting r15-5392. Backported from trunk: 2024-11-18 Tobias Burnus * plugin/plugin-gcn.c (GOMP_OFFLOAD_openacc_async_construct): In case of an error, call GOMP_PLUGIN_fatal not ..._error; use NULL not false in return. Backported from trunk: 2024-08-07 Tobias Burnus * testsuite/libgomp.c-c++-common/target-link-2.c: Reset variable value to handle multi-device tests. Backported from trunk: 2024-11-18 Tobias Burnus PR libgomp/117626 * plugin/plugin-nvptx.c (nvptx_open_device): Use 'CUDA_CALL_ERET' with 'NULL' as error return instead of 'CUDA_CALL' that returns false. 2024-08-01 Tobias Burnus Backported from master: 2024-08-01 Tobias Burnus Richard Biener Backported from master: 2024-07-29 Tobias Burnus PR fortran/115559 * testsuite/libgomp.fortran/declare-target-link.f90: New test. 2024-07-29 Tobias Burnus Backported from master: 2024-07-29 Tobias Burnus PR middle-end/116107 * target.c (gomp_map_vars_internal): Honor array mapping offsets with declare-target 'link' variables. * testsuite/libgomp.c-c++-common/target-link-2.c: New test. 2024-07-19 Thomas Schwinge Backported from trunk: 2024-07-19 Thomas Schwinge * config/gcn/target.c (GOMP_teams4): Document. * config/nvptx/target.c (GOMP_teams4): Likewise. * target.c (GOMP_teams4): Likewise. Backported from trunk: 2024-07-19 Thomas Schwinge * config/gcn/libgomp-gcn.h (GOMP_TEAM_NUM): Inject. * config/gcn/target.c (GOMP_teams4): Handle. * config/gcn/team.c (gomp_gcn_enter_kernel): Initialize. * config/gcn/teams.c (omp_get_team_num): Adjust. 2024-04-16 Andrew Pinski * gfortran.dg/gomp/atomic-21.f90: Update testcase for the removal of `;`. 2024-05-02 Jakub Jelinek * testsuite/libgomp.c/declare-variant-4.h (gfx90c, gfx1036, gfx1103): New functions. (f): Add #pragma omp declare variant directives for those. * testsuite/libgomp.c/declare-variant-4-gfx90c.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx1036.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx1103.c: New test. 2024-05-04 Sandra Loosemore * libgomp.texi (OpenMP 5.0): Mark metadirective and declare variant as implemented. (OpenMP 5.1): Mark target_device as supported. Add changed interaction between declare target and OpenMP context and dynamic selector support. (OpenMP 5.2): Mark otherwise clause as supported, note that default is also still accepted. 2024-05-04 Sandra Loosemore * testsuite/libgomp.fortran/metadirective-1.f90: New. * testsuite/libgomp.fortran/metadirective-2.f90: New. * testsuite/libgomp.fortran/metadirective-3.f90: New. * testsuite/libgomp.fortran/metadirective-4.f90: New. * testsuite/libgomp.fortran/metadirective-5.f90: New. * testsuite/libgomp.fortran/metadirective-6.f90: New. 2024-05-04 Sandra Loosemore * testsuite/libgomp.c-c++-common/metadirective-1.c: New. * testsuite/libgomp.c-c++-common/metadirective-2.c: New. * testsuite/libgomp.c-c++-common/metadirective-3.c: New. * testsuite/libgomp.c-c++-common/metadirective-4.c: New. * testsuite/libgomp.c-c++-common/metadirective-5.c: New. 2024-05-04 Sandra Loosemore * testsuite/libgomp.c++/metadirective-template-1.C: New. * testsuite/libgomp.c++/metadirective-template-2.C: New. * testsuite/libgomp.c++/metadirective-template-3.C: New. 2024-05-04 Sandra Loosemore * Makefile.am (libgomp_la_SOURCES): Add selector.c. * Makefile.in: Regenerate. * config/gcn/selector.c: New. * config/linux/selector.c: New. * config/linux/x86/selector.c: New. * config/nvptx/selector.c: New. * libgomp-plugin.h (GOMP_OFFLOAD_evaluate_device): New. * libgomp.h (struct gomp_device_descr): Add evaluate_device_func field. * libgomp.map (GOMP_5.1.3): New, add GOMP_evaluate_target_device. * libgomp.texi (OpenMP Context Selectors): Document dynamic selector matching of kind/arch/isa. * libgomp_g.h (GOMP_evaluate_current_device): New. (GOMP_evaluate_target_device): New. * oacc-host.c (host_evaluate_device): New. (host_openacc_exec): Initialize evaluate_device_func field to host_evaluate_device. * plugin/plugin-gcn.c (gomp_match_selectors): New. (gomp_match_isa): New. (GOMP_OFFLOAD_evaluate_device): New. * plugin/plugin-nvptx.c (struct ptx_device): Add compute_major and compute_minor fields. (nvptx_open_device): Read compute capability information from device. (gomp_match_selectors): New. (gomp_match_selector): New. (CHECK_ISA): New macro. (GOMP_OFFLOAD_evaluate_device): New. * selector.c: New. * target.c (GOMP_evaluate_target_device): New. (gomp_load_plugin_for_device): Load evaluate_device plugin function. 2023-10-30 Tobias Burnus * testsuite/libgomp.fortran/allocate-8a.f90: New test. 2023-10-26 Tobias Burnus * libgomp.texi (OpenMP Impl. Status): Document that 'omp allocate' is now supported for C++ stack/automatic variables. * testsuite/libgomp.c-c++-common/allocate-4.c: Renamed from ... * testsuite/libgomp.c/allocate-4.c: ... this. * testsuite/libgomp.c-c++-common/allocate-5.c: Renamed from ... * testsuite/libgomp.c/allocate-5.c: ... this. * testsuite/libgomp.c-c++-common/allocate-6.c: Renamed from ... * testsuite/libgomp.c/allocate-6.c: ... this. * testsuite/libgomp.c++/allocate-2.C: New test. 2023-09-19 Julian Brown * target.c (omp_target_memcpy_rect_worker): Add 1D strided transfer support. 2023-08-10 Julian Brown * testsuite/libgomp.c-c++-common/declare-mapper-18.c: New test. * testsuite/libgomp.fortran/declare-mapper-25.f90: New test. * testsuite/libgomp.fortran/declare-mapper-28.f90: New test. 2023-07-14 Julian Brown * testsuite/libgomp.c-c++-common/array-shaping-14.c: New test. 2023-09-05 Julian Brown * testsuite/libgomp.c/array-shaping-1.c: New test. * testsuite/libgomp.c/array-shaping-2.c: New test. * testsuite/libgomp.c/array-shaping-3.c: New test. * testsuite/libgomp.c/array-shaping-4.c: New test. * testsuite/libgomp.c/array-shaping-5.c: New test. * testsuite/libgomp.c/array-shaping-6.c: New test. 2023-07-03 Julian Brown * libgomp.h (omp_noncontig_array_desc): Add span field. * target.c (omp_target_memcpy_rect_worker): Add span parameter. Update forward declaration. Handle span != element_size. (gomp_update): Handle bias in descriptor's size slot. Update calls to omp_target_memcpy_rect_worker. * testsuite/libgomp.fortran/noncontig-updates-1.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-2.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-3.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-4.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-5.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-6.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-7.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-8.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-9.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-10.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-11.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-12.f90: New test. * testsuite/libgomp.fortran/noncontig-updates-13.f90: New test. 2023-07-03 Julian Brown * libgomp.h (omp_noncontig_array_desc): New struct. * target.c (omp_target_memcpy_rect_worker): Add stride array parameter. Forward declare. Add STRIDES parameter and strided update support. (gomp_update): Add noncontiguous (strided/shaped) update support. * testsuite/libgomp.c++/array-shaping-1.C: New test. * testsuite/libgomp.c++/array-shaping-2.C: New test. * testsuite/libgomp.c++/array-shaping-3.C: New test. * testsuite/libgomp.c++/array-shaping-4.C: New test. * testsuite/libgomp.c++/array-shaping-5.C: New test. * testsuite/libgomp.c++/array-shaping-6.C: New test. * testsuite/libgomp.c++/array-shaping-7.C: New test. * testsuite/libgomp.c++/array-shaping-8.C: New test. * testsuite/libgomp.c++/array-shaping-9.C: New test. * testsuite/libgomp.c++/array-shaping-10.C: New test. * testsuite/libgomp.c++/array-shaping-11.C: New test. * testsuite/libgomp.c++/array-shaping-12.C: New test. * testsuite/libgomp.c++/array-shaping-13.C: New test. 2023-08-10 Julian Brown * testsuite/libgomp.fortran/declare-mapper-30.f90: New test. * testsuite/libgomp.fortran/declare-mapper-4.f90: Adjust test for new lookup behaviour. 2023-07-12 Julian Brown * testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c: Add expected warning. * testsuite/libgomp.oacc-fortran/declare-create-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/declare-create-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/declare-create-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-2.f90: Likewise. 2023-06-30 Julian Brown * testsuite/libgomp.fortran/declare-mapper-2.f90: New test. * testsuite/libgomp.fortran/declare-mapper-3.f90: New test. * testsuite/libgomp.fortran/declare-mapper-4.f90: New test. * testsuite/libgomp.fortran/declare-mapper-6.f90: New test. * testsuite/libgomp.fortran/declare-mapper-7.f90: New test. * testsuite/libgomp.fortran/declare-mapper-8.f90: New test. * testsuite/libgomp.fortran/declare-mapper-9.f90: New test. * testsuite/libgomp.fortran/declare-mapper-10.f90: New test. * testsuite/libgomp.fortran/declare-mapper-11.f90: New test. * testsuite/libgomp.fortran/declare-mapper-12.f90: New test. * testsuite/libgomp.fortran/declare-mapper-13.f90: New test. * testsuite/libgomp.fortran/declare-mapper-15.f90: New test. * testsuite/libgomp.fortran/declare-mapper-17.f90: New test. * testsuite/libgomp.fortran/declare-mapper-18.f90: New test. * testsuite/libgomp.fortran/declare-mapper-19.f90: New test. * testsuite/libgomp.fortran/declare-mapper-20.f90: New test. * testsuite/libgomp.fortran/declare-mapper-21.f90: New test. 2023-06-30 Julian Brown * testsuite/libgomp.c-c++-common/declare-mapper-9.c: Enable for C. * testsuite/libgomp.c-c++-common/declare-mapper-10.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-11.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-12.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-13.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-14.c: Likewise. 2023-06-30 Julian Brown * testsuite/libgomp.c++/declare-mapper-1.C: New test. * testsuite/libgomp.c++/declare-mapper-2.C: New test. * testsuite/libgomp.c++/declare-mapper-3.C: New test. * testsuite/libgomp.c++/declare-mapper-4.C: New test. * testsuite/libgomp.c++/declare-mapper-5.C: New test. * testsuite/libgomp.c++/declare-mapper-6.C: New test. * testsuite/libgomp.c++/declare-mapper-7.C: New test. * testsuite/libgomp.c++/declare-mapper-8.C: New test. * testsuite/libgomp.c-c++-common/declare-mapper-9.c: New test (only enabled for C++ for now). * testsuite/libgomp.c-c++-common/declare-mapper-10.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-11.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-12.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-13.c: Likewise. * testsuite/libgomp.c-c++-common/declare-mapper-14.c: Likewise. 2023-06-19 Julian Brown * testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c: New test. 2023-06-19 Julian Brown * testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-1.f90: New test. * testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-2.f90: New test. 2023-06-19 Julian Brown * testsuite/libgomp.oacc-fortran/declare-create-1.f90: New test. * testsuite/libgomp.oacc-fortran/declare-create-2.f90: New test. * testsuite/libgomp.oacc-fortran/declare-create-3.f90: New test. 2023-06-19 Julian Brown * testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr70828-2.c: Likewise. * testsuite/libgomp.oacc-fortran/pr70828.f90: Likewise. * testsuite/libgomp.oacc-fortran/pr70828-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/pr70828-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/pr70828-4.f90: Likewise. * testsuite/libgomp.oacc-fortran/pr70828-5.f90: Likewise. * testsuite/libgomp.oacc-fortran/pr70828-6.f90: Likewise. 2023-05-12 Julian Brown * testsuite/libgomp.c++/static-aggr-constructor-destructor-1.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-2.C: New test. * testsuite/libgomp.c++/static-aggr-constructor-destructor-3.C: New test. 2023-04-03 Thomas Schwinge PR other/76739 * target.c (gomp_map_vars_internal): Pass pre-allocated 'ptrblock' to 'goacc_noncontig_array_create_ptrblock'. * oacc-parallel.c (goacc_noncontig_array_create_ptrblock): Adjust. * oacc-int.h (goacc_noncontig_array_create_ptrblock): Adjust. * libgomp.texi (AMD Radeon, nvptx): Document OpenMP 'pinned' memory. 2024-06-06 Jakub Jelinek * libgomp.texi (OpenMP 5.1 status): Mark Loop transformation constructs as implemented. 2024-06-05 Jakub Jelinek Frederik Harwath Sandra Loosemore * testsuite/libgomp.c-c++-common/imperfect-transform-1.c: New test. * testsuite/libgomp.c-c++-common/imperfect-transform-2.c: New test. * testsuite/libgomp.c-c++-common/matrix-1.h: New test. * testsuite/libgomp.c-c++-common/matrix-constant-iter.h: New test. * testsuite/libgomp.c-c++-common/matrix-helper.h: New test. * testsuite/libgomp.c-c++-common/matrix-no-directive-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-no-directive-unroll-full-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-distribute-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-parallel-masked-taskloop-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-parallel-masked-taskloop-simd-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-target-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-target-teams-distribute-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-taskloop-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-omp-teams-distribute-parallel-for-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-simd-1.c: New test. * testsuite/libgomp.c-c++-common/matrix-transform-variants-1.h: New test. * testsuite/libgomp.c-c++-common/target-imperfect-transform-1.c: New test. * testsuite/libgomp.c-c++-common/target-imperfect-transform-2.c: New test. * testsuite/libgomp.c-c++-common/unroll-1.c: New test. * testsuite/libgomp.c-c++-common/unroll-non-rect-1.c: New test. * testsuite/libgomp.c++/matrix-no-directive-unroll-full-1.C: New test. * testsuite/libgomp.c++/tile-2.C: New test. * testsuite/libgomp.c++/tile-3.C: New test. * testsuite/libgomp.c++/unroll-1.C: New test. * testsuite/libgomp.c++/unroll-2.C: New test. * testsuite/libgomp.c++/unroll-full-tile.C: New test. * testsuite/libgomp.fortran/imperfect-transform-1.f90: New test. * testsuite/libgomp.fortran/imperfect-transform-2.f90: New test. * testsuite/libgomp.fortran/inner-1.f90: New test. * testsuite/libgomp.fortran/nested-fn.f90: New test. * testsuite/libgomp.fortran/target-imperfect-transform-1.f90: New test. * testsuite/libgomp.fortran/target-imperfect-transform-2.f90: New test. * testsuite/libgomp.fortran/tile-1.f90: New test. * testsuite/libgomp.fortran/tile-2.f90: New test. * testsuite/libgomp.fortran/tile-unroll-1.f90: New test. * testsuite/libgomp.fortran/tile-unroll-2.f90: New test. * testsuite/libgomp.fortran/tile-unroll-3.f90: New test. * testsuite/libgomp.fortran/tile-unroll-4.f90: New test. * testsuite/libgomp.fortran/unroll-1.f90: New test. * testsuite/libgomp.fortran/unroll-2.f90: New test. * testsuite/libgomp.fortran/unroll-3.f90: New test. * testsuite/libgomp.fortran/unroll-4.f90: New test. * testsuite/libgomp.fortran/unroll-5.f90: New test. * testsuite/libgomp.fortran/unroll-6.f90: New test. * testsuite/libgomp.fortran/unroll-7a.f90: New test. * testsuite/libgomp.fortran/unroll-7b.f90: New test. * testsuite/libgomp.fortran/unroll-7c.f90: New test. * testsuite/libgomp.fortran/unroll-7.f90: New test. * testsuite/libgomp.fortran/unroll-8.f90: New test. * testsuite/libgomp.fortran/unroll-simd-1.f90: New test. * testsuite/libgomp.fortran/unroll-tile-1.f90: New test. * testsuite/libgomp.fortran/unroll-tile-2.f90: New test. 2023-03-24 Thomas Schwinge * target.c (gomp_target_rev): Instead of 'dev_to_host_cpy', 'host_to_dev_cpy', 'token', take a single 'goacc_asyncqueue'. * libgomp.h (gomp_target_rev): Adjust. * libgomp-plugin.c (GOMP_PLUGIN_target_rev): Adjust. * libgomp-plugin.h (GOMP_PLUGIN_target_rev): Adjust. * plugin/plugin-gcn.c (process_reverse_offload): Adjust. * plugin/plugin-nvptx.c (rev_off_dev_to_host_cpy) (rev_off_host_to_dev_cpy): Remove. (GOMP_OFFLOAD_run): Adjust. * target.c (gomp_unmap_vars_internal): Queue splay-tree keys for removal after main loop. PR other/76739 * oacc-parallel.c (GOACC_parallel_keyed): Given OpenACC 'async', defer 'free' of non-contiguous array support data structures. * target.c (gomp_map_vars_internal): Likewise. 2023-03-23 Tobias Burnus * testsuite/libgomp.fortran/map-alloc-comp-8.f90: New test. 2023-03-10 Thomas Schwinge PR other/76739 * libgomp.h (goacc_map_vars): Add 'struct goacc_ncarray_info *' formal parameter. (gomp_map_vars_openacc): Remove. * target.c (goacc_map_vars): Adjust. (gomp_map_vars_openacc): Remove. * oacc-mem.c (acc_map_data, goacc_enter_datum) (goacc_enter_data_internal): Adjust. * oacc-parallel.c (GOACC_parallel_keyed, GOACC_data_start): Adjust. 2024-05-31 Thomas Schwinge * libgomp.texi (nvptx): Update. * testsuite/libgomp.fortran/target-print-1-nvptx.f90: Remove. * testsuite/libgomp.fortran/target-print-1.f90: Adjust. * testsuite/libgomp.oacc-fortran/error_stop-2-nvptx.f: New. * testsuite/libgomp.oacc-fortran/error_stop-2.f: Adjust. * testsuite/libgomp.oacc-fortran/print-1-nvptx.f90: Adjust. * testsuite/libgomp.oacc-fortran/print-1.f90: Adjust. * testsuite/libgomp.oacc-fortran/stop-2-nvptx.f: New. * testsuite/libgomp.oacc-fortran/stop-2.f: Adjust. 2024-05-31 Thomas Schwinge * plugin/cuda-lib.def (cuCtxSetLimit): Add. * plugin/plugin-nvptx.c (nvptx_open_device): Handle 'GOMP_NVPTX_NATIVE_GPU_THREAD_STACK_SIZE' environment variable. 2024-06-05 Thomas Schwinge * plugin/plugin-nvptx.c (nvptx_do_global_cdtors): New. (nvptx_close_device, GOMP_OFFLOAD_load_image) (GOMP_OFFLOAD_unload_image): Call it. 2024-06-05 Thomas Schwinge * config/nvptx/error.c (exit): Don't override. * testsuite/libgomp.oacc-fortran/error_stop-1.f: Update. * testsuite/libgomp.oacc-fortran/error_stop-2.f: Likewise. * testsuite/libgomp.oacc-fortran/error_stop-3.f: Likewise. * testsuite/libgomp.oacc-fortran/stop-1.f: Likewise. * testsuite/libgomp.oacc-fortran/stop-2.f: Likewise. * testsuite/libgomp.oacc-fortran/stop-3.f: Likewise. 2022-12-06 Paul-Antoine Arras * config/gcn/selector.c (GOMP_evaluate_current_device): Recognise 'amdgcn' as arch, and '-march' values (as well as 'gfx803') as isa traits. * testsuite/libgomp.c-c++-common/metadirective-6.c: New test. 2022-11-02 Tobias Burnus * testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90: Adjust. * testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90: Likewise. * testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1.f90: New. * testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90: Adjust. * testsuite/libgomp.oacc-fortran/declare-allocatable-1-runtime.f90: Likewise. * testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90: Likewise. 2022-11-02 Tobias Burnus * testsuite/libgomp.fortran/target-enter-data-3a.f90: New test. 2022-11-02 Tobias Burnus * testsuite/libgomp.fortran/target-13.f90: Update test. 2023-11-19 Tobias Burnus Chung-Lin Tang * testsuite/libgomp.c++/c++.exp (check_effective_target_c, check_effective_target_c++): Add. * testsuite/libgomp.c/c.exp (check_effective_target_c, check_effective_target_c++): Add. * testsuite/libgomp.fortran/uses_allocators_2.f90: Remove 'sorry'. * testsuite/libgomp.c-c++-common/uses_allocators-1.c: New test. * testsuite/libgomp.c-c++-common/uses_allocators-2.c: New test. * testsuite/libgomp.c-c++-common/uses_allocators-3.c: New test. * testsuite/libgomp.c-c++-common/uses_allocators-4.c: New test. * testsuite/libgomp.fortran/uses_allocators_3.f90: New test. * testsuite/libgomp.fortran/uses_allocators_4.f90: New test. * testsuite/libgomp.fortran/uses_allocators_5.f90: New test. * testsuite/libgomp.fortran/uses_allocators_6.f90: New test. 2023-08-23 Andrew Stubbs * Makefile.am (libgomp_la_SOURCES): Add usmpin-allocator.c. * Makefile.in: Regenerate. * config/linux/allocator.c: Include unistd.h. (pin_ctx): New variable. (ctxlock): New variable. (linux_init_pin_ctx): New function. (linux_memspace_alloc): Use usmpin-allocator for pinned memory. (linux_memspace_free): Likewise. (linux_memspace_realloc): Likewise. * libgomp.h (usmpin_init_context): New prototype. (usmpin_register_memory): New prototype. (usmpin_alloc): New prototype. (usmpin_free): New prototype. (usmpin_realloc): New prototype. * testsuite/libgomp.c/alloc-pinned-1.c: Adjust for new behaviour. * testsuite/libgomp.c/alloc-pinned-2.c: Likewise. * testsuite/libgomp.c/alloc-pinned-5.c: Likewise. * testsuite/libgomp.c/alloc-pinned-8.c: New test. * usmpin-allocator.c: New file. 2023-08-23 Andrew Stubbs * config/linux/allocator.c: Include assert.h. (using_device_for_page_locked): New variable. (linux_memspace_alloc): Add init0 parameter. Support device pinning. (linux_memspace_calloc): Set init0 to true. (linux_memspace_free): Support device pinning. (linux_memspace_realloc): Support device pinning. (MEMSPACE_ALLOC): Set init0 to false. * libgomp-plugin.h (GOMP_OFFLOAD_page_locked_host_alloc): New prototype. (GOMP_OFFLOAD_page_locked_host_free): Likewise. * libgomp.h (gomp_page_locked_host_alloc): Likewise. (gomp_page_locked_host_free): Likewise. (struct gomp_device_descr): Add page_locked_host_alloc_func and page_locked_host_free_func. * libgomp_g.h (GOMP_enable_pinned_mode): New prototype. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_page_locked_host_alloc): New function. (GOMP_OFFLOAD_page_locked_host_free): Likewise. * target.c (device_for_page_locked): New variable. (get_device_for_page_locked): New function. (gomp_page_locked_host_alloc): Likewise. (gomp_page_locked_host_free): Likewise. (gomp_load_plugin_for_device): Add page_locked_host_alloc and page_locked_host_free. * testsuite/libgomp.c/alloc-pinned-1.c: Change expectations for NVPTX devices. * testsuite/libgomp.c/alloc-pinned-2.c: Likewise. * testsuite/libgomp.c/alloc-pinned-3.c: Likewise. * testsuite/libgomp.c/alloc-pinned-4.c: Likewise. * testsuite/libgomp.c/alloc-pinned-5.c: Likewise. * testsuite/libgomp.c/alloc-pinned-6.c: Likewise. 2023-08-23 Andrew Stubbs * config/linux/allocator.c (always_pinned_mode): New variable. (GOMP_enable_pinned_mode): New function. (linux_memspace_alloc): Disable pinning when always_pinned_mode set. (linux_memspace_calloc): Likewise. (linux_memspace_free): Likewise. (linux_memspace_realloc): Likewise. * libgomp.map: Add GOMP_enable_pinned_mode. * testsuite/libgomp.c/alloc-pinned-7.c: New test. * testsuite/libgomp.c-c++-common/alloc-pinned-1.c: New test. 2023-08-23 Andrew Stubbs * allocator.c (omp_max_predefined_alloc): Update. (predefined_alloc_mapping): Add ompx_pinned_mem_alloc entry. (omp_aligned_alloc): Support ompx_pinned_mem_alloc. (omp_free): Likewise. (omp_aligned_calloc): Likewise. (omp_realloc): Likewise. * omp.h.in (omp_allocator_handle_t): Add ompx_pinned_mem_alloc. * omp_lib.f90.in: Add ompx_pinned_mem_alloc. * testsuite/libgomp.c/alloc-pinned-5.c: New test. * testsuite/libgomp.c/alloc-pinned-6.c: New test. * testsuite/libgomp.fortran/alloc-pinned-1.f90: New test. 2022-03-01 Tobias Burnus * testsuite/libgomp.fortran/allocatable-comp.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-3.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-4.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-5.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-6.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-7.f90: New test. 2022-06-21 Kwok Cheung Yeung * plugin/cuda-lib.def (cuMemAllocManaged): Add new call. (cuPointerGetAttribute): Likewise. 2021-10-21 Tobias Burnus * testsuite/libgomp.oacc-fortran/declare-allocatable-1.f90: Compile with -Wopenacc-parallelism. * testsuite/libgomp.oacc-fortran/declare-allocatable-3.f90: Likewise. 2021-03-01 Kwok Cheung Yeung * testsuite/libgomp.c-c++-common/collapse-4.c: New. * testsuite/libgomp.fortran/collapse5.f90: New. 2021-08-03 Andrew Stubbs * config/gcn/bar.h (gomp_barrier_init): Limit thread count to the actual physical number. * config/gcn/team.c (gomp_team_start): Don't attempt to set up threads that do not exist. 2021-02-23 Andrew Stubbs * plugin/plugin-nvptx.c (GOMP_OFFLOAD_alloc): Remove early call to nvptx_stacks_free. 2021-01-13 Julian Brown * testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: Adjust for loop lowering changes. * testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c: Likewise. 2020-07-16 Tobias Burnus * testsuite/libgomp.oacc-fortran/firstprivate-int.f90: Use highest available integer kind instead of assuming that kind=16 exists. 2020-03-24 Kwok Cheung Yeung * testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: XFAIL execution test. 2019-09-17 Julian Brown Kwok Cheung Yeung * testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c (main): Fix async-safety issue. Increase number of iterations. 2018-10-04 Cesar Philippidis Julian Brown * libgomp.h (gomp_acc_declare_allocate): Remove prototype. * oacc-mem.c (gomp_acc_declare_allocate): New function. * oacc-parallel.c (goacc_enter_data_internal): Handle GOMP_MAP_DECLARE_ALLOCATE. Pass new pointer argument to gomp_acc_declare_allocate. (goacc_exit_data_internal): Handle GOMP_MAP_DECLARE_DEALLOCATE. Unlock device mutex around gomp_acc_declare_allocate call. Pass new pointer argument. Handle group pointer mapping for deallocate. (find_group_last): Handle GOMP_MAP_DECLARE_ALLOCATE and GOMP_MAP_DECLARE_DEALLOCATE groupings. * testsuite/libgomp.oacc-fortran/allocatable-scalar.f90: New test. * testsuite/libgomp.oacc-fortran/declare-allocatable-2.f90: New test. * testsuite/libgomp.oacc-fortran/declare-allocatable-3.f90: New test. * testsuite/libgomp.oacc-fortran/declare-allocatable-4.f90: New test. 2019-09-05 Julian Brown * testsuite/libgomp.oacc-fortran/lib-13.f90: End data region after wait API calls. 2019-08-08 Julian Brown * plugin/plugin-gcn.c (GOMP_OFFLOAD_openacc_exec_params, GOMP_OFFLOAD_openacc_async_exec_params): New functions. 2019-07-31 Julian Brown * config/nvptx/gomp_print.c (gomp_print_string, gomp_print_integer, gomp_print_double): New. 2019-01-23 Thomas Schwinge * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Update. 2018-12-20 Maciej W. Rozycki * oacc-init.c (get_property_any): Add profiling code. 2017-02-28 Thomas Schwinge * Makefile.am (libgomp_la_SOURCES): Add oacc-profiling-acc_register_library.c. * Makefile.in: Regenerate. * libgomp.texi: Remove paragraph about acc_register_library. * oacc-parallel.c (GOACC_parallel_keyed_internal): Set device_api for profiling. * oacc-profiling-acc_register_library.c: New file. * oacc-profiling.c (goacc_profiling_initialize): Call acc_register_library. Avoid duplicate registration. (acc_register_library): Remove. * config/nvptx/oacc-profiling-acc_register_library.c: New empty file. * config/nvptx/oacc-profiling.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_prof-dispatch-1.c: Remove call to acc_register_library. * 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-01-09 Julian Brown * libgomp.texi: Update mentions of OpenACC version to 2.6. Update section numbers to match version 2.6 of the spec. 2019-03-19 Julian Brown * testsuite/libgomp.oacc-c-c++-common/lib-93.c: Adjust target selector. 2018-09-05 Cesar Philippidis Chung-Lin Tang * testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c: New test. 2018-09-20 Cesar Philippidis * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust test case to conform to the new behavior of the auto clause in OpenACC 2.5. 2018-12-22 Cesar Philippidis Julian Brown * oacc-parallel.c (GOACC_parallel_keyed): Handle GOMP_MAP_FIRSTPRIVATE_INT host addresses. * plugin/plugin-nvptx.c (nvptx_exec): Handle GOMP_MAP_FIRSTPRIVATE_INT host addresses. * testsuite/libgomp.oacc-c++/firstprivate-int.C: New test. * testsuite/libgomp.oacc-c-c++-common/firstprivate-int.c: New test. * testsuite/libgomp.oacc-fortran/firstprivate-int.f90: New test. 2018-10-05 Nathan Sidwell Tom de Vries Thomas Schwinge Julian Brown * testsuite/libgomp.oacc-c-c++-common/loop-default-compile.c: New. * testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c: New. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise. 2018-10-22 James Norris Cesar Philippidis Tom de Vries * testsuite/libgomp.oacc-fortran/data-3.f90: Update parallel regions to denote variables copyied in via acc enter data as present. * testsuite/libgomp.oacc-c-c++-common/subr.h: Reimplement. * testsuite/libgomp.oacc-c-c++-common/subr.ptx: Regenerated PTX. * testsuite/libgomp.oacc-c-c++-common/timer.h: Removed. * testsuite/libgomp.oacc-c-c++-common/lib-69.c: Change async checks. * testsuite/libgomp.oacc-c-c++-common/lib-70.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-72.c: Rework kernel i/f and change async checks. * testsuite/libgomp.oacc-c-c++-common/lib-73.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-74.c: Rework kernel i/f and timing checks. * testsuite/libgomp.oacc-c-c++-common/lib-75.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-76.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-78.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-79.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-81.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-93.c: New test. 2018-12-13 Cesar Philippidis Nathan Sidwell Julian Brown * testsuite/libgomp.oacc-c-c++-common/par-reduction-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c: New test. * testsuite/libgomp.oacc-fortran/reduction-9.f90: New test. 2018-06-29 Cesar Philippidis James Norris * oacc-parallel.c (GOACC_parallel_keyed): Handle Fortran deviceptr clause. (GOACC_data_start): Likewise. * testsuite/libgomp.oacc-fortran/deviceptr-1.f90: New test. 2019-02-12 Julian Brown * oacc-cuda.c (acc_set_cuda_stream): Return 0 on error/invalid arguments. * testsuite/libgomp.oacc-c-c++-common/lib-84.c: Handle unnumbered async stream being an alias for a numbered async stream. * testsuite/libgomp.oacc-c-c++-common/lib-85.c: Likewise. 2020-04-19 Chung-Lin Tang PR other/76739 * libgomp_g.h (GOACC_data_start): Add variadic '...' to declaration. * libgomp.h (gomp_map_vars_openacc): New function declaration. * oacc-int.h (struct goacc_ncarray_dim): New struct declaration. (struct goacc_ncarray_descr_type): Likewise. (struct goacc_ncarray): Likewise. (struct goacc_ncarray_info): Likewise. (goacc_noncontig_array_create_ptrblock): New function declaration. * oacc-parallel.c (goacc_noncontig_array_count_rows): New function. (goacc_noncontig_array_compute_sizes): Likewise. (goacc_noncontig_array_fill_rows_1): Likewise. (goacc_noncontig_array_fill_rows): Likewise. (goacc_process_noncontiguous_arrays): Likewise. (goacc_noncontig_array_create_ptrblock): Likewise. (GOACC_parallel_keyed): Use goacc_process_noncontiguous_arrays to handle non-contiguous array descriptors at end of varargs, adjust to use gomp_map_vars_openacc. (GOACC_data_start): Likewise. Adjust function type to accept varargs. * target.c (gomp_map_vars_internal): Add struct goacc_ncarray_info * nca_info parameter, add handling code for non-contiguous arrays. (gomp_map_vars_openacc): Add new function for specialization of gomp_map_vars_internal for OpenACC structured region usage. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h: Support header for new tests.