aboutsummaryrefslogtreecommitdiff
path: root/libgomp
AgeCommit message (Collapse)AuthorFilesLines
9 hoursDaily bump.GCC Administrator1-0/+52
27 hourslibgomp: Init hash table for 'indirect'-clause of 'declare target' on the ↵Tobias Burnus7-101/+295
host [PR114445, PR119857] Especially with unified-shared memory and especially with C++'s virtual functions, it is not uncommon to have on the device a function pointer that points to the host function - but has an associated device. If the pointed-to function is (explicitly or implicitly) 'declare target' with the 'indirect' clause, it is added to the lookup table. Before this commit, the conversion of the lookup table into a lookup hash table happened every time a device kernel was launched on the first team - albeit if already converted, the function immediately returned. Ignoring the overhead, there was also a race: If multiple teams were launched, it could happen that another team of the same target region already tried to use the lookup table which it was still being created. Likewise when lauching a kernel with 'nowait' and directly afterward another kernel, there could be a race of creating the table. With this commit, the creating of the kernel has been moved to the host-plugin's GOMP_OFFLOAD_load_image. The previous code stored a pointer to the host/device pointer array, which makes it hard when creating the hash table on the host (data is needed for finding the slot) - but accessing it on the device (where the lookup has to work as well). As the hash-table implementation (only) supports integral value as payload (0 and 1 having special meaning), the solution was to move to an uint128_t variable to store both the host and device address. As the host-side library is typically dynamically linked and the device-side one statically, there is the problem of backward compatibility. The current implementation permits both older binaries and newer libgomp and newer binaries with older libgomp. I could imagine us breaking the latter eventually, but for now there is up and downward compatibility. (Obviously, the race is only fixed if new + new is combined.) Code wise, on the device exist GOMP_INDIRECT_ADDR_MAP which was updated to point to the host/device-address array. Now additionally GOMP_INDIRECT_ADDR_HMAP exists, which contains the hash-table map. If the latter exists, libgomp only updates it and the former remains a NULL pointer; it is also untouched if there are no indirect functions. Being NULL therefore avoids the call to the device-side build_indirect_map. The code also currently supports to have no hash and a linear walk. I think that remained from testing; due to the backward-compat feature, it can actually be turned of on either side. libgomp/ChangeLog: PR libgomp/119857 PR libgomp/114445 * config/accel/target-indirect.c: Change to use uint128_t instead of a struct as data structure and add GOMP_INDIRECT_ADDR_HMAP as host-accessible variable. (struct indirect_map_t): Remove. (USE_HASHTAB_LOOKUP, INDIRECT_DEV_ADDR, INDIRECT_HOST_ADDR, SET_INDIRECT_HOST_ADDR, SET_INDIRECT_ADDRS): Define. (htab_free): Use __builtin_unreachable. (htab_hash, htab_eq, GOMP_target_map_indirect_ptr, build_indirect_map): Update for new representation and new pointer-to-hash variable. * config/gcn/team.c (gomp_gcn_enter_kernel): Only call build_indirect_map when GOMP_INDIRECT_ADDR_MAP. * config/nvptx/team.c (gomp_nvptx_main): Likewise. * libgomp-plugin.h (GOMP_INDIRECT_ADDR_HMAP): Define. * plugin/plugin-gcn.c: Conditionally include build-target-indirect-htab.h. (USE_HASHTAB_LOOKUP_FOR_INDIRECT): Define. (create_target_indirect_map): New prototype. (GOMP_OFFLOAD_load_image): Update to create the device's indirect-function hash table on the host. * plugin/plugin-nvptx.c: Conditionally include build-target-indirect-htab.h. (USE_HASHTAB_LOOKUP_FOR_INDIRECT): Define. (create_target_indirect_map): New prototype. (GOMP_OFFLOAD_load_image): Update to create the device's indirect-function hash table on the host. * plugin/build-target-indirect-htab.h: New file.
27 hourslibgomp: Add Fortran version of acc_copyout_finalize_async and ↵Tobias Burnus5-51/+320
acc_delete_finalize_async OpenACC 2.5 added several functions for C and Fortran; while acc_{copyout,delete}{,_finalize,_async} exist for both, for some reasons only the C version of acc_{copyout,delete}_finalize_async was actually added, even though the documentation (.texi) and the .map file listed also the auxiliar Fortran functions! OpenACC 2.5 added the Fortran version with the following odd interface: 'type, dimension(:[,:]...)'. In OpenACC 2.6, it was then updated to the Fortran 2018 syntax: 'type(*), dimension(..)', which is also used in openacc.f90 internally. This commit now also updates the documentation to the newer syntax - plus fixes a function-name typo: acc_delete_async_finalize should have the _async at the end not in the middle! libgomp/ChangeLog: * libgomp.map (OACC_2.5): Move previously unimplemented acc_{copyout,delete}_finalize_async_{32,64,array}_h_ to ... (OACC_2.6.1): ... here. * libgomp.texi (acc_copyin, acc_present_or_copyin, acc_create, acc_present_or_create, acc_copyout, acc_update_device, acc_update_self, acc_is_present): Use 'type(*), dimension(..)' instead of 'type, dimension(:[,:]...)' for Fortran. (acc_delete): Likewise; change acc_delete_async_finalize to acc_delete_finalize_async. * openacc.f90 (openacc_internal): Add interfaces for acc_{copyout,delete}_finalize_async_{{32,64,array}_h,_l}. (openacc): Add generic interfaces for acc_copyout_finalize_async and acc_delete_finalize_async. (acc_{copyout,delete}_finalize_async_{32,64,array}_h): New. * openacc_lib.h: Add generic interfaces for acc_copyout_finalize_async and acc_delete_finalize_async. * testsuite/libgomp.oacc-fortran/pr92970-1.f90: New test.
11 daysDaily bump.GCC Administrator1-0/+5
12 dayslibgomp: Use consistent formatting in <omp.h>Jakub Jelinek1-41/+42
I've noticed the new C++ part of omp.h uses libstdc++ coding conventions, while the rest of the header (and libgomp) is formatted using the normal gcc coding conventions like gcc/. This patch makes it consistent. 2025-09-06 Jakub Jelinek <jakub@redhat.com> * omp.h.in: Fix up formatting of __cplusplus >= 201103L guarded code from libstc++ style to GCC/libgomp style.
2025-08-07Daily bump.GCC Administrator1-0/+20
2025-08-06openmp: Add support for iterators in 'target update' clauses (C/C++)Kwok Cheung Yeung4-0/+204
This adds support for iterators in 'to' and 'from' clauses in the 'target update' OpenMP directive. gcc/c/ * c-parser.cc (c_parser_omp_clause_from_to): Parse 'iterator' modifier. * c-typeck.cc (c_finish_omp_clauses): Finish iterators for to/from clauses. gcc/cp/ * parser.cc (cp_parser_omp_clause_from_to): Parse 'iterator' modifier. * semantics.cc (finish_omp_clauses): Finish iterators for to/from clauses. gcc/ * gimplify.cc (remove_unused_omp_iterator_vars): Display unused variable warning for 'to' and 'from' clauses. (gimplify_scan_omp_clauses): Add argument for iterator loop sequence. Gimplify the clause decl and size into the iterator loop if iterators are used. (gimplify_omp_workshare): Add argument for iterator loops sequence in call to gimplify_scan_omp_clauses. (gimplify_omp_target_update): Call remove_unused_omp_iterator_vars and build_omp_iterators_loops. Add loop sequence as argument when calling gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses and building the Gimple statement. * tree-pretty-print.cc (dump_omp_clause): Call dump_omp_iterators for to/from clauses with iterators. * tree.cc (omp_clause_num_ops): Add extra operand for OMP_CLAUSE_FROM and OMP_CLAUSE_TO. * tree.h (OMP_CLAUSE_HAS_ITERATORS): Add check for OMP_CLAUSE_TO and OMP_CLAUSE_FROM. (OMP_CLAUSE_ITERATORS): Likewise. gcc/testsuite/ * c-c++-common/gomp/target-update-iterators-1.c: New. * c-c++-common/gomp/target-update-iterators-2.c: New. * c-c++-common/gomp/target-update-iterators-3.c: New. libgomp/ * 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-08-06openmp: Add support for iterators in map clauses (C/C++)Kwok Cheung Yeung4-7/+325
This adds preliminary support for iterators in map clauses within OpenMP 'target' constructs (which includes constructs such as 'target enter data'). Iterators with non-constant loop bounds are not currently supported. gcc/c/ * c-parser.cc (c_parser_omp_variable_list): Use location of the map expression as the clause location. (c_parser_omp_clause_map): Parse 'iterator' modifier. * c-typeck.cc (c_finish_omp_clauses): Finish iterators. Apply iterators to generated clauses. gcc/cp/ * parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier. * semantics.cc (finish_omp_clauses): Finish iterators. Apply iterators to generated clauses. gcc/ * gimple-pretty-print.cc (dump_gimple_omp_target): Print expanded iterator loops. * gimple.cc (gimple_build_omp_target): Add argument for iterator loops sequence. Initialize iterator loops field. * gimple.def (GIMPLE_OMP_TARGET): Set GSS symbol to GSS_OMP_TARGET. * gimple.h (gomp_target): Set GSS symbol to GSS_OMP_TARGET. Add extra field for iterator loops. (gimple_build_omp_target): Add argument for iterator loops sequence. (gimple_omp_target_iterator_loops): New. (gimple_omp_target_iterator_loops_ptr): New. (gimple_omp_target_set_iterator_loops): New. * gimplify.cc (find_var_decl): New. (copy_omp_iterator): New. (remap_omp_iterator_var_1): New. (remap_omp_iterator_var): New. (remove_unused_omp_iterator_vars): New. (struct iterator_loop_info_t): New type. (iterator_loop_info_map_t): New type. (build_omp_iterators_loops): New. (enter_omp_iterator_loop_context_1): New. (enter_omp_iterator_loop_context): New. (enter_omp_iterator_loop_context): New. (exit_omp_iterator_loop_context): New. (gimplify_adjust_omp_clauses): Add argument for iterator loop sequence. Gimplify the clause decl and size into the iterator loop if iterators are used. (gimplify_omp_workshare): Call remove_unused_omp_iterator_vars and build_omp_iterators_loops for OpenMP target expressions. Add loop sequence as argument when calling gimplify_adjust_omp_clauses and building the Gimple statement. * gimplify.h (enter_omp_iterator_loop_context): New prototype. (exit_omp_iterator_loop_context): New prototype. * gsstruct.def (GSS_OMP_TARGET): New. * omp-low.cc (lower_omp_map_iterator_expr): New. (lower_omp_map_iterator_size): New. (finish_omp_map_iterators): New. (lower_omp_target): Add sorry if iterators used with deep mapping. Call lower_omp_map_iterator_expr before assigning to sender ref. Call lower_omp_map_iterator_size before setting the size. Insert iterator loop sequence before the statements for the target clause. * tree-nested.cc (convert_nonlocal_reference_stmt): Walk the iterator loop sequence of OpenMP target statements. (convert_local_reference_stmt): Likewise. (convert_tramp_reference_stmt): Likewise. * tree-pretty-print.cc (dump_omp_iterators): Dump extra iterator information if present. (dump_omp_clause): Call dump_omp_iterators for iterators in map clauses. * tree.cc (omp_clause_num_ops): Add operand for OMP_CLAUSE_MAP. (walk_tree_1): Do not walk last operand of OMP_CLAUSE_MAP. * tree.h (OMP_CLAUSE_HAS_ITERATORS): New. (OMP_CLAUSE_ITERATORS): New. gcc/testsuite/ * c-c++-common/gomp/map-6.c (foo): Amend expected error message. * c-c++-common/gomp/target-map-iterators-1.c: New. * c-c++-common/gomp/target-map-iterators-2.c: New. * c-c++-common/gomp/target-map-iterators-3.c: New. * c-c++-common/gomp/target-map-iterators-4.c: New. libgomp/ * 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. Co-authored-by: Andrew Stubbs <ams@baylibre.com>
2025-07-22Daily bump.GCC Administrator1-0/+8
2025-07-21Adjust 'libgomp.c++/target-cdtor-{1,2}.C' for 'targetm.cxx.use_aeabi_atexit' ↵Thomas Schwinge2-12/+22
[PR119853, PR119854] Fix-up for commit aafe942227baf8c2bcd4cac2cb150e49a4b895a9 "GCN, nvptx offloading: Host/device compatibility: Itanium C++ ABI, DSO Object Destruction API [PR119853, PR119854]": we need to adjust for 'targetm.cxx.use_aeabi_atexit': gcc/config/arm/arm.cc:#define TARGET_CXX_USE_AEABI_ATEXIT arm_cxx_use_aeabi_atexit gcc/config/arm/arm.cc:/* The EABI says __aeabi_atexit should be used to register static gcc/config/arm/arm.cc- destructors. */ gcc/config/arm/arm.cc- gcc/config/arm/arm.cc-static bool gcc/config/arm/arm.cc:arm_cxx_use_aeabi_atexit (void) gcc/config/arm/arm.cc-{ gcc/config/arm/arm.cc- return TARGET_AAPCS_BASED; gcc/config/arm/arm.cc-} ..., which 'gcc/cp/decl.cc:get_atexit_node' then acts on: call '__aeabi_atexit' instead of '__cxa_atexit', and swap two arguments. PR target/119853 PR target/119854 libgomp/ * testsuite/libgomp.c++/target-cdtor-1.C: Adjust for 'targetm.cxx.use_aeabi_atexit'. * testsuite/libgomp.c++/target-cdtor-2.C: Likewise.
2025-07-19Daily bump.GCC Administrator1-0/+7
2025-07-18amdgcn, libgomp: Remove unused variable (PR121156)Andrew Stubbs1-6/+2
There's a new compiler warning breaking the build. This fixes it. The variable appears to be genuinely vestigial. libgomp/ChangeLog: PR target/121156 * config/gcn/bar.c (gomp_team_barrier_wait_end): Remove unused "generation" variable. (gomp_team_barrier_wait_cancel_end): Likewise.
2025-07-18Daily bump.GCC Administrator1-0/+10
2025-07-17GCN, nvptx offloading: Restrain 'WARNING: program timed out.' while in ↵Thomas Schwinge6-6/+6
'dynamic_cast' only for effective-target 'offload_device' [PR119692] In PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' offloading": > --- Comment #8 from Rainer Orth <ro at gcc dot gnu.org> --- > The last commit made things worse on sparc-sun-solaris2.11: since that one > (dg-timeout 10) I regularly get > > WARNING: libgomp.c++/target-exceptions-bad_cast-1.C (test for excess errors) > program timed out. > FAIL: libgomp.c++/target-exceptions-bad_cast-1.C (test for excess errors) > UNRESOLVED: libgomp.c++/target-exceptions-bad_cast-1.C compilation failed to produce executable > UNRESOLVED: libgomp.c++/target-exceptions-bad_cast-1.C scan-tree-dump-times optimized "gimple_call <__cxa_bad_cast, " 1 > > Before that, the test had no issue. Compiling the test on an unloaded system > usually takes less than 1 sec, but when fully loaded, times can go up. To keep things simple, let's restrict this temporary (yeah...) workaround to apply only for effective-target 'offload_device', just like the 'dg-xfail-run-if' itself. PR target/119692 libgomp/ * testsuite/libgomp.c++/pr119692-1-4.C: '{ dg-timeout 10 { target offload_device } }'. * testsuite/libgomp.c++/pr119692-1-5.C: Likewise. * testsuite/libgomp.c++/target-exceptions-bad_cast-1.C: Likewise. * testsuite/libgomp.c++/target-exceptions-bad_cast-2.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-1.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C: Likewise.
2025-06-25Daily bump.GCC Administrator1-0/+10
2025-06-24Fortran/OpenACC: Add Fortran support for acc_attach/acc_detachTobias Burnus5-20/+193
While C/++ support the routines acc_attach{,_async} and acc_detach{,_finalize}{,_async} routines since a long time, the Fortran API routines where only added in OpenACC 3.3. Unfortunately, they cannot directly be implemented in the library as GCC will introduce a temporary array descriptor in some cases, which causes the attempted attachment to the this temporary variable instead of to the original one. Therefore, those API routines are handled in a special way in the compiler. gcc/fortran/ChangeLog: * trans-stmt.cc (gfc_trans_call_acc_attach_detach): New. (gfc_trans_call): Call it. libgomp/ChangeLog: * libgomp.texi (acc_attach, acc_detach): Update for Fortran version. * openacc.f90 (acc_attach{,_async}, acc_detach{,_finalize}{,_async}): Add. * openacc_lib.h: Likewise. * testsuite/libgomp.oacc-fortran/acc-attach-detach-1.f90: New test. * testsuite/libgomp.oacc-fortran/acc-attach-detach-2.f90: New test.
2025-06-20Daily bump.GCC Administrator1-0/+17
2025-06-19libgomp/target.c: Fix buffer size for 'omp requires' diagnosticTobias Burnus1-6/+7
One of the buffers that printed the list of set 'omp requires' requirements missed the 'self' clause addition, being potentially to short when all device-affecting clauses were passed. Solved it by moving the sizeof(<string of all permitted values>" into a new '#define' just above the associated gomp_requires_to_name function. libgomp/ChangeLog: * target.c (GOMP_REQUIRES_NAME_BUF_LEN): Define. (GOMP_offload_register_ver, gomp_target_init): Use it for the char buffer size.
2025-06-19libgomp.texi: Document omp(x)::allocator::*, restructure memory allocator docTobias Burnus1-68/+113
libgomp/ChangeLog: * libgomp.texi (omp_init_allocator): Refer to 'Memory allocation' for available memory spaces. (OMP_ALLOCATOR): Move list of traits and predefined memspaces and allocators to ... (Memory allocation): ... here. Document omp(x)::allocator::*; minor wording tweaks, be more explicit about memkind, pinned and pool_size. Co-authored-by: waffl3x <waffl3x@baylibre.com>
2025-06-18Daily bump.GCC Administrator1-0/+4
2025-06-17OpenMP: Fix implicit 'declare target' for <ostream>Tobias Burnus1-0/+25
libstdc++-v3/include/std/ostream contains: namespace std _GLIBCXX_VISIBILITY(default) { ... template<typename _CharT, typename _Traits> inline basic_ostream<_CharT, _Traits>& endl(basic_ostream<_CharT, _Traits>& __os) { return flush(__os.put(__os.widen('\n'))); } ... #include <bits/ostream.tcc> and the latter, libstdc++-v3/include/bits/ostream.tcc, has: // Inhibit implicit instantiations for required instantiations, // which are defined via explicit instantiations elsewhere. #if _GLIBCXX_EXTERN_TEMPLATE extern template class basic_ostream<char>; extern template ostream& endl(ostream&); Before this commit, omp_discover_declare_target_tgt_fn_r marked 'endl' as (implicitly) declare target - but not the calls in it due to the 'extern' (DECL_EXTERNAL). Thanks to inlining and as 'endl' is (therefore) not used and, hence, discarded by the linker; hencet, it works with -O0 and -O1. However, as the (unused) function still exits, IPA CP (enabled by -O2) will try to do constant-value propagation and fails as the definition of 'widen' is not available. Solution is to still walk 'endl' despite being an 'extern(al)' decl; this has been restricted for now to DECL_DECLARED_INLINE_P. gcc/ChangeLog: * omp-offload.cc (omp_discover_declare_target_tgt_fn_r): Also walk external functions that are declare inline (and have a DECL_SAVED_TREE). libgomp/ChangeLog: * testsuite/libgomp.c++/declare_target-2.C: New test.
2025-06-11Daily bump.GCC Administrator1-0/+5
2025-06-10gcn: Add experimental MI300 (gfx942) supportTobias Burnus2-0/+16
As gfx942 and gfx950 belong to gfx9-4-generic, the latter two are also added. Note that there are no specific optimizations for MI300, yet. For none of the mentioned devices, any multilib is build by default; use '--with-multilib-list=' when configuring GCC to build them alongside. gfx942 was added in LLVM (and its mc assembler, used by GCC) in version 18, generic support in LLVM 19 and gfx950 in LLVM 20. gcc/ChangeLog: * config/gcn/gcn-devices.def: Add gfx942, gfx950 and gfx9-4-generic. * config/gcn/gcn-opts.h (TARGET_CDNA3, TARGET_CDNA3_PLUS, TARGET_GLC_NAME, TARGET_TARGET_SC_CACHE): Define. (TARGET_ARCHITECTED_FLAT_SCRATCH): Use also for CDNA3. * config/gcn/gcn.h (gcn_isa): Add ISA_CDNA3 to the enum. * config/gcn/gcn.cc (print_operand): Update 'g' to use TARGET_GLC_NAME; add 'G' to print TARGET_GLC_NAME unconditionally. * config/gcn/gcn-valu.md (scatter, gather): Use TARGET_GLC_NAME. * config/gcn/gcn.md: Use %G<num> instead of glc; use 'buffer_inv sc1' for TARGET_TARGET_SC_CACHE. * doc/invoke.texi (march): Add gfx942, gfx950 and gfx9-4-generic. * doc/install.texi (amdgcn*-*-*): Add gfx942, gfx950 and gfx9-4-generic. * config/gcn/gcn-tables.opt: Regenerate. libgomp/ChangeLog: * testsuite/libgomp.c/declare-variant-4.h (gfx942): New variant function. * testsuite/libgomp.c/declare-variant-4-gfx942.c: New test.
2025-06-07Daily bump.GCC Administrator1-0/+12
2025-06-06OpenMP: Add omp_get_initial_device/omp_get_num_devices builtinsTobias Burnus1-0/+14
By adding omp_get_initial_device and omp_get_num_devices builtins for C, C++, and Fortran, the following can be achieved: * By making them pure, multiple calls can be avoiding in some cases. * Some comparisons can be optimized at compile time. omp_get_initial_device will be converted to omp_get_num_devices for consistency; note that OpenMP 6 also permits omp_initial_device (== -1) as value. If GCC has not been configure for offloading, either intrinsic will leads to 0 - and on the offload side, -1 (= omp_initial_device) is returned for omp_initial_device. gcc/fortran/ChangeLog: * f95-lang.cc (ATTR_PURE_NOTHROW_LIST): Define. * trans-expr.cc (get_builtin_fn): Handle omp_get_num_devices and omp_get_intrinsic_device. * gfortran.h (gfc_option_t): Add disable_omp_... for them. * options.cc (gfc_handle_option): Handle them with -fno-builtin-. gcc/ChangeLog: * gimple-fold.cc (gimple_fold_builtin_omp_get_initial_device, gimple_fold_builtin_omp_get_num_devices): New. (gimple_fold_builtin): Call them. * omp-builtins.def (BUILT_IN_OMP_GET_INITIAL_DEVICE): Add (BUILT_IN_OMP_GET_NUM_DEVICES): Make uservisible + pure. libgomp/ChangeLog: * libgomp.texi (omp_get_num_devices, omp_get_intrinsic_device): Document builtin handling. gcc/testsuite/ChangeLog: * c-c++-common/gomp/omp_get_num_devices_initial_device-2.c: New test. * c-c++-common/gomp/omp_get_num_devices_initial_device.c: New test. * gfortran.dg/gomp/omp_get_num_devices_initial_device-2.f90: New test. * gfortran.dg/gomp/omp_get_num_devices_initial_device.f90: New test. Co-authored-by: Sandra Loosemore <sloosemore@baylibre.com>
2025-06-06libgomp.c/target-map-zero-sized-3.c: Fix code for non-USM offload [PR120530]Tobias Burnus1-3/+4
A mapping clause was missing, causing the code to fail with offloading when a host pointer was not device accessible. libgomp/ChangeLog: PR target/120530 * testsuite/libgomp.c/target-map-zero-sized-3.c (main): Add missing map clause; remove unused variable.
2025-06-05Daily bump.GCC Administrator1-0/+6
2025-06-04libgomp.texi (omp_interop_*): Add note about 5.2-to-6.0 incompatibilityTobias Burnus1-0/+20
GCC uses the 6.0 types - which are unfortunately not quite compatible with code expecting 5.1/5.2 data types. Therefore, this commit adds a note to hopefully reduce surprises. Namely: For C/C++: while OpenMP 5.1 and 5.2 used 'int *ret_code', OpenMP 6.0 uses 'omp_interop_rc_t *ret_code' in omp_interop_{int,ptr,str} and 'int' instead of 'omp_interop_rc_t ret_code' in omp_get_interop_rc_desc. Neither C nor C++ like passing the wrong pointer type, albeit for C, GCC < 14 and clang only warn (gcc >= r14-6037-g9715c545d33b3a has an error) and using -fpermissive turns it into a warning and -Wno-incompatible-pointer-types silences it for C. C++ also dislikes passing an int to an enum, albeit -fpermissive turns the error into a warning with g++ (but not clang++). And, here, using an enum on the caller side works with both int and enum on the callee side. libgomp/ChangeLog: * libgomp.texi (omp_interop_{int,ptr,str,rc_desc}): Add note about the 'ret_code' type change in OpenMP 6. Co-authored-by: Sandra Loosemore <sloosemore@baylibre.com>
2025-06-04Daily bump.GCC Administrator1-0/+8
2025-06-03libgomp: Fix up omp_target_memset-3.c test for C++ [PR120444]Jakub Jelinek1-2/+2
The test PASSes for C, but FAILs for C++: .../libgomp.c-c++-common/omp_target_memset-3.c: In function 'void test_it(void*, int, size_t)': .../libgomp.c-c++-common/omp_target_memset-3.c:31:7: warning: pointer of type 'void *' used in arithmetic [-Wpointer-arith] .../libgomp.c-c++-common/omp_target_memset-3.c:33:13: error: invalid conversion from 'void*' to 'int8_t*' {aka 'signed char*'} [-fpermissive] .../libgomp.c-c++-common/omp_target_memset-3.c:10:19: note: initializing argument 1 of 'void init_val(int8_t*, int, size_t)' .../libgomp.c-c++-common/omp_target_memset-3.c:37:14: error: invalid conversion from 'void*' to 'int8_t*' {aka 'signed char*'} [-fpermissive] .../libgomp.c-c++-common/omp_target_memset-3.c:17:20: note: initializing argument 1 of 'void check_val(int8_t*, int, size_t)' .../libgomp.c-c++-common/omp_target_memset-3.c:38:18: warning: pointer of type 'void *' used in arithmetic [-Wpointer-arith] .../libgomp.c-c++-common/omp_target_memset-3.c:38:18: error: invalid conversion from 'void*' to 'int8_t*' {aka 'signed char*'} [-fpermissive] .../libgomp.c-c++-common/omp_target_memset-3.c:17:20: note: initializing argument 1 of 'void check_val(int8_t*, int, size_t)' .../libgomp.c-c++-common/omp_target_memset-3.c: In function 'int main()': .../libgomp.c-c++-common/omp_target_memset-3.c:46:7: warning: pointer of type 'void *' used in arithmetic [-Wpointer-arith] The following two-liner fixes that, tested on x86_64-linux and i686-linux. 2025-06-03 Jakub Jelinek <jakub@redhat.com> PR libgomp/120444 * testsuite/libgomp.c-c++-common/omp_target_memset-3.c (test_it): Change ptr argument type from void * to int8_t *. (main): Change ptr variable type from void * to int8_t * and cast omp_target_alloc result to the latter type.
2025-06-03Daily bump.GCC Administrator1-0/+26
2025-06-02libgomp: Add OpenMP's omp_target_memset/omp_target_memset_asyncTobias Burnus16-4/+639
PR libgomp/120444 include/ChangeLog: * cuda/cuda.h (cuMemsetD8, cuMemsetD8Async): Declare. libgomp/ChangeLog: * libgomp-plugin.h (GOMP_OFFLOAD_memset): Declare. * libgomp.h (struct gomp_device_descr): Add memset_func. * libgomp.map (GOMP_6.0.1): Add omp_target_memset{,_async}. * libgomp.texi (Device Memory Routines): Document them. * omp.h.in (omp_target_memset, omp_target_memset_async): Declare. * omp_lib.f90.in (omp_target_memset, omp_target_memset_async): Add interfaces. * omp_lib.h.in (omp_target_memset, omp_target_memset_async): Likewise. * plugin/cuda-lib.def: Add cuMemsetD8. * plugin/plugin-gcn.c (struct hsa_runtime_fn_info): Add hsa_amd_memory_fill_fn. (init_hsa_runtime_functions): DLSYM_OPT_FN load it. (GOMP_OFFLOAD_memset): New. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_memset): New. * target.c (omp_target_memset_int, omp_target_memset, omp_target_memset_async_helper, omp_target_memset_async): New. (gomp_load_plugin_for_device): Add DLSYM (memset). * testsuite/libgomp.c-c++-common/omp_target_memset.c: New test. * testsuite/libgomp.c-c++-common/omp_target_memset-2.c: New test. * testsuite/libgomp.c-c++-common/omp_target_memset-3.c: New test. * testsuite/libgomp.fortran/omp_target_memset.f90: New test. * testsuite/libgomp.fortran/omp_target_memset-2.f90: New test.
2025-05-31Daily bump.GCC Administrator1-0/+126
2025-05-30Add 'libgomp.c++/target-valarray-1.C'Thomas Schwinge2-0/+201
libgomp/ * testsuite/libgomp.c++/target-std__valarray-1.C: New. * testsuite/libgomp.c++/target-std__valarray-1.output: Likewise.
2025-05-30libgomp: Add testcases for concurrent access to standard C++ containers on ↵Thomas Schwinge24-0/+92
offload targets, a number of USM variants libgomp/ * testsuite/libgomp.c++/target-std__array-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__array-concurrent.C: Adjust. * testsuite/libgomp.c++/target-std__bitset-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__bitset-concurrent.C: Adjust. * testsuite/libgomp.c++/target-std__deque-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__deque-concurrent.C: Adjust. * testsuite/libgomp.c++/target-std__forward_list-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__forward_list-concurrent.C: Adjust. * testsuite/libgomp.c++/target-std__list-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__list-concurrent.C: Adjust. * testsuite/libgomp.c++/target-std__map-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__map-concurrent.C: Adjust. * testsuite/libgomp.c++/target-std__multimap-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__multimap-concurrent.C: Adjust. * testsuite/libgomp.c++/target-std__multiset-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__multiset-concurrent.C: Adjust. * testsuite/libgomp.c++/target-std__set-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__set-concurrent.C: Adjust. * testsuite/libgomp.c++/target-std__span-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__span-concurrent.C: Adjust. * testsuite/libgomp.c++/target-std__valarray-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__valarray-concurrent.C: Adjust. * testsuite/libgomp.c++/target-std__vector-concurrent-usm.C: New. * testsuite/libgomp.c++/target-std__vector-concurrent.C: Adjust.
2025-05-30libgomp: Add testcases for concurrent access to standard C++ containers on ↵Kwok Cheung Yeung20-0/+1318
offload targets libgomp/ * testsuite/libgomp.c++/target-std__array-concurrent.C: New. * testsuite/libgomp.c++/target-std__bitset-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__deque-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__flat_map-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__flat_multimap-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__flat_multiset-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__flat_set-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__forward_list-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__list-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__map-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__multimap-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__multiset-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__set-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__span-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__unordered_map-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__unordered_multimap-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__unordered_multiset-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__unordered_set-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__valarray-concurrent.C: Likewise. * testsuite/libgomp.c++/target-std__vector-concurrent.C: Likewise. Co-authored-by: Thomas Schwinge <tschwinge@baylibre.com>
2025-05-30libgomp: Add testcases for the standard C++ math library on offload targetsKwok Cheung Yeung3-0/+608
libgomp/ * testsuite/libgomp.c++/target-std__cmath.C: New. * testsuite/libgomp.c++/target-std__complex.C: Likewise. * testsuite/libgomp.c++/target-std__numbers.C: Likewise.
2025-05-30Add 'libgomp.c++/target-flex-[...].C' test casesWaffl3x23-0/+2930
libgomp/ChangeLog: * testsuite/libgomp.c++/target-flex-10.C: New test. * testsuite/libgomp.c++/target-flex-100.C: New test. * testsuite/libgomp.c++/target-flex-101.C: New test. * testsuite/libgomp.c++/target-flex-11.C: New test. * testsuite/libgomp.c++/target-flex-12.C: New test. * testsuite/libgomp.c++/target-flex-2000.C: New test. * testsuite/libgomp.c++/target-flex-2001.C: New test. * testsuite/libgomp.c++/target-flex-2002.C: New test. * testsuite/libgomp.c++/target-flex-2003.C: New test. * testsuite/libgomp.c++/target-flex-30.C: New test. * testsuite/libgomp.c++/target-flex-300.C: New test. * testsuite/libgomp.c++/target-flex-31.C: New test. * testsuite/libgomp.c++/target-flex-32.C: New test. * testsuite/libgomp.c++/target-flex-33.C: New test. * testsuite/libgomp.c++/target-flex-41.C: New test. * testsuite/libgomp.c++/target-flex-60.C: New test. * testsuite/libgomp.c++/target-flex-61.C: New test. * testsuite/libgomp.c++/target-flex-62.C: New test. * testsuite/libgomp.c++/target-flex-70.C: New test. * testsuite/libgomp.c++/target-flex-80.C: New test. * testsuite/libgomp.c++/target-flex-81.C: New test. * testsuite/libgomp.c++/target-flex-90.C: New test. * testsuite/libgomp.c++/target-flex-common.h: New test. Co-authored-by: Thomas Schwinge <tschwinge@baylibre.com>
2025-05-30Defuse 'RESULT_DECL' check in 'pass_nrv' (for offloading compilation) [PR119835]Thomas Schwinge3-2/+7
... to avoid running into ICEs per PR119835, until that's resolved properly. PR middle-end/119835 gcc/ * tree-nrv.cc (pass_nrv::execute): Defuse 'RESULT_DECL' check. libgomp/ * testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c: '#pragma GCC optimize "-fno-inline"'. * testsuite/libgomp.c-c++-common/target-abi-struct-1.c: New. * testsuite/libgomp.c-c++-common/target-abi-struct-1-O0.c: Adjust. Co-authored-by: Richard Biener <rguenther@suse.de>
2025-05-30OpenMP: Support OpenMP 5.0 "declare mapper" directives for CJulian Brown6-6/+6
This patch adds support for "declare mapper" directives (and the "mapper" modifier on "map" clauses) for C. gcc/c/ChangeLog: * c-decl.cc (c_omp_mapper_id, c_omp_mapper_decl, c_omp_mapper_lookup, c_omp_extract_mapper_directive, c_omp_map_array_section, c_omp_scan_mapper_bindings_r, c_omp_scan_mapper_bindings): New functions. * c-objc-common.h (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define langhooks for C. * c-parser.cc (c_parser_omp_clause_map): Add declare_mapper_p parameter; handle mapper modifier. (c_parser_omp_all_clauses): Update call to c_parser_omp_clause_map. (c_parser_omp_target): Instantiate explicit mappers and record bindings for implicit mappers. (c_parser_omp_declare_mapper): Parse "declare mapper" directives. (c_parser_omp_declare): Support "declare mapper". (c_parser_omp_declare_reduction): Use inform not error_at. * c-tree.h (c_omp_finish_mapper_clauses, c_omp_mapper_lookup, c_omp_extract_mapper_directive, c_omp_map_array_section, c_omp_mapper_id, c_omp_mapper_decl, c_omp_scan_mapper_bindings, c_omp_instantiate_mappers): Add prototypes. * c-typeck.cc (c_finish_omp_clauses): Handle GOMP_MAP_PUSH_MAPPER_NAME and GOMP_MAP_POP_MAPPER_NAME. (c_omp_finish_mapper_clauses): New function (langhook). libgomp/ChangeLog: * 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. gcc/testsuite/ChangeLog: * c-c++-common/gomp/declare-mapper-3.c: Enable for C. * c-c++-common/gomp/declare-mapper-4.c: Likewise. * c-c++-common/gomp/declare-mapper-5.c: Likewise. * c-c++-common/gomp/declare-mapper-6.c: Likewise. * c-c++-common/gomp/declare-mapper-7.c: Likewise. * c-c++-common/gomp/declare-mapper-8.c: Likewise. * c-c++-common/gomp/declare-mapper-9.c: Likewise. * c-c++-common/gomp/declare-mapper-10.c: Likewise. * c-c++-common/gomp/declare-mapper-12.c: Likewise. * c-c++-common/gomp/map-6.c: Update dg-error. * gcc.dg/gomp/udr-3.c: Update for change to dg-note. * c-c++-common/gomp/declare-mapper-11.c: New. * gcc.dg/gomp/declare-mapper-10.c: New test. * gcc.dg/gomp/declare-mapper-11.c: New test. * gcc.dg/gomp/declare-mapper-13.c: New test.
2025-05-30OpenMP: C++ "declare mapper" supportJulian Brown14-0/+868
This patch adds support for OpenMP 5.0 "declare mapper" functionality for C++. I've merged it to og13 based on the last version posted upstream, with some minor changes due to the newly-added 'present' map modifier support. There's also a fix to splay-tree traversal in gimplify.cc:omp_instantiate_implicit_mappers, and this patch omits the rearrangement of gimplify.cc:gimplify_{scan,adjust}_omp_clauses that I separated out into its own patch and applied (to og13) already. gcc/c-family/ * c-common.h (c_omp_region_type): Add C_ORT_DECLARE_MAPPER and C_ORT_OMP_DECLARE_MAPPER codes. (omp_mapper_list): Add forward declaration. (c_omp_find_nested_mappers, c_omp_instantiate_mappers): Add prototypes. * c-omp.cc (c_omp_find_nested_mappers): New function. (remap_mapper_decl_info): New struct. (remap_mapper_decl_1, omp_instantiate_mapper, c_omp_instantiate_mappers): New functions. gcc/cp/ * constexpr.cc (reduced_constant_expression_p): Add OMP_DECLARE_MAPPER case. (cxx_eval_constant_expression, potential_constant_expression_1): Likewise. * cp-gimplify.cc (cxx_omp_finish_mapper_clauses): New function. * cp-objcp-common.h (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define langhooks. * cp-tree.h (lang_decl_base): Add omp_declare_mapper_p field. Recount spare bits comment. (DECL_OMP_DECLARE_MAPPER_P): New macro. (omp_mapper_id): Add prototype. (cp_check_omp_declare_mapper): Add prototype. (omp_instantiate_mappers): Add prototype. (cxx_omp_finish_mapper_clauses): Add prototype. (cxx_omp_mapper_lookup): Add prototype. (cxx_omp_extract_mapper_directive): Add prototype. (cxx_omp_map_array_section): Add prototype. * decl.cc (check_initializer): Add OpenMP declare mapper support. (cp_finish_decl): Set DECL_INITIAL for OpenMP declare mapper var decls as appropriate. * decl2.cc (mark_used): Instantiate OpenMP "declare mapper" magic var decls. * error.cc (dump_omp_declare_mapper): New function. (dump_simple_decl): Use above. * parser.cc (cp_parser_omp_clause_map): Add KIND parameter. Support "mapper" modifier. (cp_parser_omp_all_clauses): Add KIND argument to cp_parser_omp_clause_map call. (cp_parser_omp_target): Call omp_instantiate_mappers before finish_omp_clauses. (cp_parser_omp_declare_mapper): New function. (cp_parser_omp_declare): Add "declare mapper" support. * pt.cc (tsubst_decl): Adjust name of "declare mapper" magic var decls once we know their type. (tsubst_omp_clauses): Call omp_instantiate_mappers before finish_omp_clauses, for target regions. (tsubst_expr): Support OMP_DECLARE_MAPPER nodes. (instantiate_decl): Instantiate initialiser (i.e definition) for OpenMP declare mappers. * semantics.cc (gimplify.h): Include. (omp_mapper_id, omp_mapper_lookup, omp_extract_mapper_directive, cxx_omp_map_array_section, cp_check_omp_declare_mapper): New functions. (finish_omp_clauses): Delete GOMP_MAP_PUSH_MAPPER_NAME and GOMP_MAP_POP_MAPPER_NAME artificial clauses. (omp_target_walk_data): Add MAPPERS field. (finish_omp_target_clauses_r): Scan for uses of struct/union/class type variables. (finish_omp_target_clauses): Create artificial mapper binding clauses for used structs/unions/classes in offload region. gcc/fortran/ * parse.cc (tree.h, fold-const.h, tree-hash-traits.h): Add includes (for additions to omp-general.h). gcc/ * gimplify.cc (gimplify_omp_ctx): Add IMPLICIT_MAPPERS field. (new_omp_context): Initialise IMPLICIT_MAPPERS hash map. (delete_omp_context): Delete IMPLICIT_MAPPERS hash map. (instantiate_mapper_info): New structs. (remap_mapper_decl_1, omp_mapper_copy_decl, omp_instantiate_mapper, omp_instantiate_implicit_mappers): New functions. (gimplify_scan_omp_clauses): Handle MAPPER_BINDING clauses. (gimplify_adjust_omp_clauses): Instantiate implicit declared mappers. (gimplify_omp_declare_mapper): New function. (gimplify_expr): Call above function. * langhooks-def.h (lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): Add prototypes. (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define macros. (LANG_HOOK_DECLS): Add above macros. * langhooks.cc (lhd_omp_mapper_lookup, lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): New dummy functions. * langhooks.h (lang_hooks_for_decls): Add OMP_FINISH_MAPPER_CLAUSES, OMP_MAPPER_LOOKUP, OMP_EXTRACT_MAPPER_DIRECTIVE, OMP_MAP_ARRAY_SECTION hooks. * omp-general.h (omp_name_type<T>): Add templatized struct, hash type traits (for omp_name_type<tree> specialization). (omp_mapper_list<T>): Add struct. * tree-core.h (omp_clause_code): Add OMP_CLAUSE__MAPPER_BINDING_. * tree-pretty-print.cc (dump_omp_clause): Support GOMP_MAP_UNSET, GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping clauses. Support OMP_CLAUSE__MAPPER_BINDING_ and OMP_DECLARE_MAPPER. * tree.cc (omp_clause_num_ops, omp_clause_code_name): Add OMP_CLAUSE__MAPPER_BINDING_. * tree.def (OMP_DECLARE_MAPPER): New tree code. * tree.h (OMP_DECLARE_MAPPER_ID, OMP_DECLARE_MAPPER_DECL, OMP_DECLARE_MAPPER_CLAUSES): New defines. (OMP_CLAUSE__MAPPER_BINDING__ID, OMP_CLAUSE__MAPPER_BINDING__DECL, OMP_CLAUSE__MAPPER_BINDING__MAPPER): New defines. include/ * gomp-constants.h (gomp_map_kind): Add GOMP_MAP_UNSET, GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping clause types. gcc/testsuite/ * c-c++-common/gomp/map-6.c: Update error scan output. * c-c++-common/gomp/declare-mapper-3.c: New test (only enabled for C++ for now). * c-c++-common/gomp/declare-mapper-4.c: Likewise. * c-c++-common/gomp/declare-mapper-5.c: Likewise. * c-c++-common/gomp/declare-mapper-6.c: Likewise. * c-c++-common/gomp/declare-mapper-7.c: Likewise. * c-c++-common/gomp/declare-mapper-8.c: Likewise. * c-c++-common/gomp/declare-mapper-9.c: Likewise. * c-c++-common/gomp/declare-mapper-10.c: Likewise. * c-c++-common/gomp/declare-mapper-12.c: Likewise. * g++.dg/gomp/declare-mapper-1.C: New test. * g++.dg/gomp/declare-mapper-2.C: New test. * g++.dg/gomp/declare-mapper-3.C: New test. libgomp/ * 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. Co-authored-by: Tobias Burnus <tburnus@baylibre.com>
2025-05-30Daily bump.GCC Administrator1-0/+33
2025-05-29libgomp: Add OpenACC's acc_memcpy_device{,_async} routines [PR93226]Tobias Burnus13-19/+409
libgomp/ChangeLog: PR libgomp/93226 * libgomp-plugin.h (GOMP_OFFLOAD_openacc_async_dev2dev): New prototype. * libgomp.h (struct acc_dispatch_t): Add dev2dev_func. (gomp_copy_dev2dev): New prototype. * libgomp.map (OACC_2.6.1): New; add acc_memcpy_device{,_async}. * libgomp.texi (acc_memcpy_device): New. * oacc-mem.c (memcpy_tofrom_device): Change to take from/to device boolean; use memcpy not memmove; add early return if size == 0 or same device + same ptr. (acc_memcpy_to_device, acc_memcpy_to_device_async, acc_memcpy_from_device, acc_memcpy_from_device_async): Update. (acc_memcpy_device, acc_memcpy_device_async): New. * openacc.f90 (acc_memcpy_device, acc_memcpy_device_async): Add interface. * openacc_lib.h (acc_memcpy_device, acc_memcpy_device_async): Likewise. * openacc.h (acc_memcpy_device, acc_memcpy_device_async): Add prototype. * plugin/plugin-gcn.c (GOMP_OFFLOAD_openacc_async_host2dev): Update comment. (GOMP_OFFLOAD_openacc_async_dev2host): Update call. (GOMP_OFFLOAD_openacc_async_dev2dev): New. * plugin/plugin-nvptx.c (cuda_memcpy_dev_sanity_check): New. (GOMP_OFFLOAD_dev2dev): Call it. (GOMP_OFFLOAD_openacc_async_dev2dev): New. * target.c (gomp_copy_dev2dev): New. (gomp_load_plugin_for_device): Load dev2dev and async_dev2dev. * testsuite/libgomp.oacc-c-c++-common/acc_memcpy_device-1.c: New test. * testsuite/libgomp.oacc-fortran/acc_memcpy_device-1.f90: New test.
2025-05-29Daily bump.GCC Administrator1-0/+6
2025-05-28libgomp.fortran/metadirective-1.f90: Expect 'error:' for nvptx compile ↵Tobias Burnus1-1/+8
[PR118694] This should have been part of commit r16-838-gb3d07ec7ac2ccd or r16-883-g5d6ed6d604ff94 - all showing the same issue: '!$omp target' followed by a metadirective with 'teams'; if the metadirective cannot be early resolved, a diagnostic error is shown about using directives between 'target' and 'teams'. While the message is misleading, the problem is that the host invokes 'target' differently when 'teams' is present; in this case, host fallback + amdgcn offload require the no-teams case, nvptx offload the teams case such that it only can be resolved at runtime. Mark the error as 'dg-bogus + xfail' to silence the FAIL, when nvptx offloading is compiled for. (If not, the metadirective can be resolved early during compilation.) libgomp/ChangeLog: PR middle-end/118694 * testsuite/libgomp.fortran/metadirective-1.f90: xfail when compiling (also) for nvptx offloading as an error is then expected.
2025-05-24Daily bump.GCC Administrator1-0/+6
2025-05-23libgomp.c-c++-common/metadirective-1.c: Expect 'error:' for nvptx compile ↵Tobias Burnus1-1/+7
[PR118694] OpenMP's 'target teams' is strictly coupled with 'teams'; if the latter exists, the kernel is launched in directly with multiple teams. Thus, the host has to know whether the teams construct exists or not. For #pragma omp target #pragma omp metadirective when (device={arch("nvptx")}: teams loop) it is simple when 'nvptx' offloading is not supported, otherwise it depends on the default device at runtime as the user code asks for a single team for host fallback and gcn offload and multiple for nvptx offload. In any case, this commit ensures that no FAIL is printed, whatever a future solution might look like. Instead of a dg-bogus combined with an 'xfail offload_target_nvptx', one an also argue that a dg-error for 'target offload_target_nvptx' would be more appropriate. libgomp/ChangeLog: PR middle-end/118694 * testsuite/libgomp.c-c++-common/metadirective-1.c: xfail when compiling (also) for nvptx offloading as an error is then expected.
2025-05-20Daily bump.GCC Administrator1-0/+16
2025-05-19'TYPE_EMPTY_P' vs. code offloading [PR120308]Thomas Schwinge1-0/+25
We've got 'gcc/stor-layout.cc:finalize_type_size': /* Handle empty records as per the x86-64 psABI. */ TYPE_EMPTY_P (type) = targetm.calls.empty_record_p (type); (Indeed x86_64 is still the only target to define 'TARGET_EMPTY_RECORD_P', calling 'gcc/tree.cc-default_is_empty_record'.) And so it happens that for an empty struct used in code offloaded from x86_64 host (but not powerpc64le host, for example), we get to see 'TYPE_EMPTY_P' in offloading compilation (where the offload targets (currently?) don't use it themselves, and therefore aren't prepared to handle it). For nvptx offloading compilation, this causes wrong code generation: 'ptxas [...] error : Call has wrong number of parameters', as nvptx code generation for function definition doesn't pay attention to this flag (say, in 'gcc/config/nvptx/nvptx.cc:pass_in_memory', or whereever else would be appropriate to handle that), but the generic code 'gcc/calls.cc:expand_call' via 'gcc/function.cc:aggregate_value_p' does pay attention to it, and we thus get mismatching function definition vs. function call. This issue apparently isn't a problem for GCN offloading, but I don't know if that's by design or by accident. Richard Biener: > It looks like TYPE_EMPTY_P is only used during RTL expansion for ABI > purposes, so computing it during layout_type is premature as shown here. > > I would suggest to simply re-compute it at offload stream-in time. (For avoidance of doubt, the additions to 'gcc.target/nvptx/abi-struct-arg.c', 'gcc.target/nvptx/abi-struct-ret.c' are not dependent on the offload streaming code changes, but are just to mirror the changes to 'libgomp.oacc-c-c++-common/abi-struct-1.c'.) PR lto/120308 gcc/ * lto-streamer-out.cc (hash_tree): Don't handle 'TYPE_EMPTY_P' for 'lto_stream_offload_p'. * tree-streamer-in.cc (unpack_ts_type_common_value_fields): Likewise. * tree-streamer-out.cc (pack_ts_type_common_value_fields): Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c: Add empty structure testing. gcc/testsuite/ * gcc.target/nvptx/abi-struct-arg.c: Add empty structure testing. * gcc.target/nvptx/abi-struct-ret.c: Likewise.
2025-05-19Add 'libgomp.c-c++-common/target-abi-struct-1-O0.c', ↵Thomas Schwinge2-0/+99
'libgomp.oacc-c-c++-common/abi-struct-1.c' libgomp/ * testsuite/libgomp.c-c++-common/target-abi-struct-1-O0.c: New. * testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c: Likewise.
generated by cgit v1.1 at 2025-09-18 09:32:08 +0000