aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite
AgeCommit message (Collapse)AuthorFilesLines
14 daysAdjust '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-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-24Fortran/OpenACC: Add Fortran support for acc_attach/acc_detachTobias Burnus2-0/+87
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-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-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-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-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-02libgomp: Add OpenMP's omp_target_memset/omp_target_memset_asyncTobias Burnus5-0/+310
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-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-29libgomp: Add OpenACC's acc_memcpy_device{,_async} routines [PR93226]Tobias Burnus2-0/+209
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-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-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-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.
2025-05-19Fix libgomp.oacc-fortran/lib-13.f90 async bugJulian Brown1-2/+1
libgomp/ * testsuite/libgomp.oacc-fortran/lib-13.f90: End data region after wait API calls.
2025-05-15OpenMP/Fortran: Fix allocatable-component mapping of derived-type array compsTobias Burnus1-0/+75
The check whether the location expression in map clause has allocatable components was failing for some derived-type array expressions such as map(var%tiles(1)) as the compiler produced _4 = var.tiles; MEMREF(_4, _5); This commit now also handles this case. gcc/fortran/ChangeLog: * trans-openmp.cc (gfc_omp_deep_mapping_do): Handle SSA_NAME if a def_stmt is available. libgomp/ChangeLog: * testsuite/libgomp.fortran/alloc-comp-4.f90: New test.
2025-05-14OpenMP: Fix mapping of zero-sized arrays with non-literal size: ↵Tobias Burnus3-0/+230
map(var[:n]), n = 0 For map(ptr[:0]), the used map kind is GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION and it is permitted that 'ptr' does not exist. 'ptr' is set to the device pointee if it exists or to the host value otherwise. For map(ptr[:3]), the variable is first mapped and then ptr is updated to point to the just-mapped device data; the attachment uses GOMP_MAP_ATTACH. For map(ptr[:n]), generates always a GOMP_MAP_ATTACH, but when n == 0, it was failing with: "pointer target not mapped for attach" The solution is not to fail but first to check whether it was mapped before. It turned out that for the mapping part, GCC adds a run-time check whether n == 0 - and uses GOMP_MAP_ZERO_LEN_ARRAY_SECTION for the mapping. Thus, we just have to check whether there such a mapping for the address for which the GOMP_MAP_ATTACH. was requested. And, if there was, the error diagnostic can be skipped. Unsurprisingly, this issue occurs in real-world code; it was detected in a code that distributes work via MPI and for some processes, some bounds ended up to be zero. libgomp/ChangeLog: * 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-05-12GCN, nvptx offloading: Restrain 'WARNING: program timed out.' while in ↵Thomas Schwinge6-0/+18
'dynamic_cast'" [PR119692] PR target/119692 libgomp/ * testsuite/libgomp.c++/pr119692-1-4.C: '{ dg-timeout 10 }'. * 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-05-12nvptx: Support '-march=sm_61'Thomas Schwinge2-0/+16
gcc/ * config/nvptx/nvptx-sm.def: Add '61'. * config/nvptx/nvptx-gen.h: Regenerate. * config/nvptx/nvptx-gen.opt: Likewise. * config/nvptx/nvptx.cc (first_ptx_version_supporting_sm): Adjust. * config/nvptx/nvptx.opt (-march-map=sm_61, -march-map=sm_62): Likewise. * config.gcc: Likewise. * doc/invoke.texi (Nvidia PTX Options): Document '-march=sm_61'. * config/nvptx/gen-multilib-matches-tests: Extend. gcc/testsuite/ * gcc.target/nvptx/march-map=sm_61.c: Adjust. * gcc.target/nvptx/march-map=sm_62.c: Likewise. * gcc.target/nvptx/march=sm_61.c: New. libgomp/ * testsuite/libgomp.c/declare-variant-3-sm61.c: New. * testsuite/libgomp.c/declare-variant-3.h: Adjust.
2025-05-09libgomp.{c,fortran}/interop-{hip,cuda}: Fix dg-run target selectionTobias Burnus11-0/+33
While the tests checked whether the CUDA/HIP runtime is available before processing them, the execution was then done unconditionally, leading to FAIL when the default device was the host (or the wrong offload device). Now the test is only executed ('run') when the default device is an Nvidia or AMD GPU (depending on the test case, cf. the test file name). Otherwise, only a 'link' test is done. (Except when the effective-target check cannot find the runtime lib - then the test is skipped [as before].) Note: The cublas/hipblas tests use variant functions and iterate over all devices, such that the cublas or hipblas, respectively, is only called when the active device is an AMD or Nvidia device, respectively, while for the host and other device types the fallback is called. libgomp/ChangeLog: * testsuite/libgomp.c/interop-cuda-full.c: Use 'link' instead of 'run' when the default device is "! offload_device_nvptx". * testsuite/libgomp.c/interop-cuda-libonly.c: Likewise. * testsuite/libgomp.c/interop-hip-nvidia-full.c: Likewise. * testsuite/libgomp.c/interop-hip-nvidia-no-headers.c: Likewise. * testsuite/libgomp.c/interop-hip-nvidia-no-hip-header.c: Likewise. * testsuite/libgomp.fortran/interop-hip-nvidia-full.F90: Likewise. * testsuite/libgomp.fortran/interop-hip-nvidia-no-module.F90: Likewise. * testsuite/libgomp.c/interop-hip-amd-full.c: Use 'link' instead of 'run' when the default device is "! offload_device_gcn". * testsuite/libgomp.c/interop-hip-amd-no-hip-header.c: Likewise. * testsuite/libgomp.fortran/interop-hip-amd-full.F90: Likewise. * testsuite/libgomp.fortran/interop-hip-amd-no-module.F90: Likewise.
2025-05-08diagnostics: convert HTML output test plugin to 'experimental-html' sink ↵David Malcolm1-0/+1
[PR116792] In r15-3752-g48261bd26df624 I added a test plugin that overrode the regular output, instead emitting diagnostics in crude HTML form. In r15-4760-g0b73e9382ab51c I added support for multiple kinds of diagnostic output simultaneously, adding -fdiagnostics-add-output=DIAGNOSTICS-OUTPUT-SPEC -fdiagnostics-set-output=DIAGNOSTICS-OUTPUT-SPEC for adding/changing the kind of diagnostics output, supporting "text" and "sarif" output schemes. This patch promotes the HTML output code from the test plugins so that it is available from "-fdiagnostics-add-output=", using a new "experimental-html" scheme, to allow simultaneous text, sarif and html output, and to make it easier to experiment with. The patch adds Python-based testing of the emitted HTML. The patch does not affect the generated HTML, which is still crude, and not yet ready for end-users. I hope to improve it in followups. gcc/ChangeLog: PR other/116792 * Makefile.in (OBJS-libcommon): Add diagnostic-format-html.o. * diagnostic-format-html.cc: Move here from testsuite/gcc.dg/plugin/diagnostic_plugin_xhtml_format.cc. Simplify includes. Rename "xhtml" to "html" throughout. (write_escaped_text): Drop. (class xhtml_stream_output_format): Drop. (class html_file_output_format): Reimplement using diagnostic_output_file. (diagnostic_output_format_init_xhtml): Drop. (diagnostic_output_format_init_xhtml_stderr): Drop. (diagnostic_output_format_init_xhtml_file): Drop. (diagnostic_output_format_open_html_file): New. (make_html_sink): New. (xhtml_format_selftests): Convert to... (diagnostic_format_html_cc_tests): ...this. (plugin_is_GPL_compatible): Drop. (plugin_init): Drop. * diagnostic-format-html.h: New file. * doc/invoke.texi (-fdiagnostics-add-output=): Add "experimental-html" scheme. * opts-diagnostic.cc: Include "diagnostic-format-html.h". (class html_scheme_handler): New. (output_factory::output_factory): Add html_scheme_handler. (html_scheme_handler::make_sink): New. * selftest-run-tests.cc (selftest::run_tests): Call the new selftests. * selftest.h (selftest::diagnostic_format_html_cc_tests): New decl. gcc/testsuite/ChangeLog: PR other/116792 * gcc.dg/plugin/diagnostic_plugin_xhtml_format.cc: Move to gcc/diagnostic-format-html.cc. * gcc.dg/html-output/html-output.exp: New support script. * gcc.dg/html-output/missing-semicolon.c: New test. * gcc.dg/html-output/missing-semicolon.py: New test script. * gcc.dg/plugin/diagnostic-test-xhtml-1.c: Deleted test. * gcc.dg/plugin/plugin.exp (plugin_test_list): Drop moved plugin and its deleted test. * lib/gcc-dg.exp (load_lib): Add load_lib of scanhtml.exp. * lib/htmltest.py: New support script. * lib/scanhtml.exp: New support script, based on scansarif.exp. libatomic/ChangeLog: PR other/116792 * testsuite/lib/libatomic.exp: Add load_lib of scanhtml.exp. libgomp/ChangeLog: PR other/116792 * testsuite/lib/libgomp.exp: Add load_lib of scanhtml.exp. libitm/ChangeLog: PR other/116792 * testsuite/lib/libitm.exp: Add load_lib of scanhtml.exp. libphobos/ChangeLog: PR other/116792 * testsuite/lib/libphobos-dg.exp: Add load_lib of scanhtml.exp. libvtv/ChangeLog: PR other/116792 * testsuite/lib/libvtv-dg.exp: Add load_lib of scanhtml.exp. Signed-off-by: David Malcolm <dmalcolm@redhat.com>
2025-05-07libgomp.fortran/map-alloc-comp-9{,-usm}.f90: Add unified_shared_memory variantTobias Burnus2-0/+30
When host memory is device accessible - independent whether mapping is done or not (i.e. self map), the 'vtab' pointer becomes accessible, which stores the dynamic type's type and size information. In principle, we want to test: USM available but mapping is still done, but as there is no simple + reliable not-crashing way to test for this, those checks are skipped in the (pre)existing test file map-alloc-comp-9.f90. Or rather: those are only active with self-maps, which is currently only true for the host. This commit adds map-alloc-comp-9-usm.f90 which runs the same test with 'omp requires unified_shared_memory'. While OpenMP permits both actual mapping and self maps with this flag, it in theory covers the missing cases. However, currently, GCC always uses self maps with USM. Still, having a device-run self-maps check is better than nothing, even if it misses the most interesting case. libgomp/ChangeLog: * testsuite/libgomp.fortran/map-alloc-comp-9.f90: Process differently when USE_USM_REQUIREMENT is set. * testsuite/libgomp.fortran/map-alloc-comp-9-usm.f90: New test.
2025-05-06libgomp: Update SVE testTejas Belagod1-11/+47
Fix udr-sve.c target test that to check for the correct results based on the OpenMP clauses used. The test was first written with a misunderstood functionality of the reduction clause. Tested with aarch64-linux-gnu. OK for trunk? libgomp/ChangeLog: * testsuite/libgomp.c-target/aarch64/udr-sve.c: Fix test.
2025-05-05'libgomp.c/interop-hsa.c': GCN offloading onlyThomas Schwinge1-1/+3
Fix-up for commit 8d84ea28510054fbbb8a2b7441916bd75e29163f "OpenMP, GCN: Add interop-hsa testcase", which added 'libgomp.c/interop-hsa.c'. If nvptx offloading compilation is enabled in addition to GCN, the former ICEs: during RTL pass: final [...]/libgomp.c/interop-hsa.c: In function 'get_kernel_ptr': [...]/libgomp.c/interop-hsa.c:34:1: internal compiler error: RTL check: expected code 'subreg', have 'reg' in nvptx_print_operand, at config/nvptx/nvptx.cc:3082 0x1ccdb96 internal_error(char const*, ...) [...]/gcc/diagnostic-global-context.cc:517 0x7446c3 rtl_check_failed_code1(rtx_def const*, rtx_code, char const*, int, char const*) [...]/gcc/rtl.cc:770 0x7fa533 nvptx_print_operand [...]/gcc/config/nvptx/nvptx.cc:3082 0xb25f34 output_operand(rtx_def*, int) [...]/gcc/final.cc:3641 0xb26f07 output_asm_insn(char const*, rtx_def**) [...]/gcc/final.cc:3534 0xb29d91 output_asm_insn(char const*, rtx_def**) [...]/gcc/final.cc:2639 0xb29d91 final_scan_insn_1 [...]/gcc/final.cc:2642 0xb2a59f final_scan_insn(rtx_insn*, _IO_FILE*, int, int, int*) [...]/gcc/final.cc:2892 0xb2a68c final_1 [...]/gcc/final.cc:1983 0xb2b378 rest_of_handle_final [...]/gcc/final.cc:4250 0xb2b378 execute [...]/gcc/final.cc:4328 Regardless of the issue that nvptx offloading compilation probably shouldn't ICE, the 'asm' insert clearly is valid for GCN only. libgomp/ * testsuite/libgomp.c/interop-hsa.c: GCN offloading only.
2025-05-01OpenMP: Restore lost Fortran testcase for 'omp allocate'Tobias Burnus1-0/+45
This testcase, which is present on the OG13 and OG14 branches, was overlooked when the Fortran support for 'omp allocate' was added to mainline (commit d4b6d147920b93297e621124a99ed01e7e310d92 from December 2023). libgomp/ChangeLog * testsuite/libgomp.fortran/allocate-8a.f90: New test.
2025-04-25OpenMP, GCN: Add interop-hsa testcaseAndrew Stubbs1-0/+203
This testcase ensures that the interop HSA support is sufficient to run a kernel manually on the same device. libgomp/ChangeLog: * testsuite/libgomp.c/interop-hsa.c: New test.
2025-04-25GCN, nvptx offloading: Host/device compatibility: Itanium C++ ABI, DSO ↵Thomas Schwinge2-0/+244
Object Destruction API [PR119853, PR119854] '__dso_handle' for '__cxa_atexit', '__cxa_finalize'. See <https://itanium-cxx-abi.github.io/cxx-abi/abi.html#dso-dtor>. PR target/119853 PR target/119854 libgcc/ * config/gcn/crt0.c (_fini_array): Call '__GCC_offload___cxa_finalize'. * config/nvptx/gbl-ctors.c (__static_do_global_dtors): Likewise. libgomp/ * 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.
2025-04-25Add 'libgomp.c-c++-common/target-cdtor-1.c'Thomas Schwinge1-0/+89
libgomp/ * testsuite/libgomp.c-c++-common/target-cdtor-1.c: New.
2025-04-25GCN: Properly switch sections in 'gcn_hsa_declare_function_name' [PR119737]Andrew Pinski4-12/+0
There are GCN/C++ target as well as offloading codes, where the hard-coded section names in 'gcn_hsa_declare_function_name' do not fit, and assembly thus fails: LLVM ERROR: Size expression must be absolute. This commit progresses GCN target: [-FAIL: g++.dg/init/call1.C -std=gnu++17 (internal compiler error: Aborted signal terminated program as)-] [-FAIL:-]{+PASS:+} g++.dg/init/call1.C -std=gnu++17 (test for excess errors) [-UNRESOLVED:-]{+PASS:+} g++.dg/init/call1.C -std=gnu++17 [-compilation failed to produce executable-]{+execution test+} [-FAIL: g++.dg/init/call1.C -std=gnu++26 (internal compiler error: Aborted signal terminated program as)-] [-FAIL:-]{+PASS:+} g++.dg/init/call1.C -std=gnu++26 (test for excess errors) [-UNRESOLVED:-]{+PASS:+} g++.dg/init/call1.C -std=gnu++26 [-compilation failed to produce executable-]{+execution test+} UNSUPPORTED: g++.dg/init/call1.C -std=gnu++98: exception handling not supported ..., and GCN offloading: [-XFAIL: libgomp.c++/target-exceptions-throw-1.C (internal compiler error: Aborted signal terminated program as)-] [-XFAIL: libgomp.c++/target-exceptions-throw-1.C PR119737 at line 7 (test for bogus messages, line )-] [-XFAIL:-]{+PASS:+} libgomp.c++/target-exceptions-throw-1.C (test for excess errors) [-UNRESOLVED:-]{+PASS:+} libgomp.c++/target-exceptions-throw-1.C [-compilation failed to produce executable-]{+execution test+} {+PASS: libgomp.c++/target-exceptions-throw-1.C output pattern test+} [-XFAIL: libgomp.c++/target-exceptions-throw-2.C (internal compiler error: Aborted signal terminated program as)-] [-XFAIL: libgomp.c++/target-exceptions-throw-2.C PR119737 at line 7 (test for bogus messages, line )-] [-XFAIL:-]{+PASS:+} libgomp.c++/target-exceptions-throw-2.C (test for excess errors) [-UNRESOLVED:-]{+PASS:+} libgomp.c++/target-exceptions-throw-2.C [-compilation failed to produce executable-]{+execution test+} {+PASS: libgomp.c++/target-exceptions-throw-2.C output pattern test+} [-XFAIL: libgomp.oacc-c++/exceptions-throw-1.C -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O2 (internal compiler error: Aborted signal terminated program as)-] [-XFAIL: libgomp.oacc-c++/exceptions-throw-1.C -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O2 PR119737 at line 7 (test for bogus messages, line )-] [-XFAIL:-]{+PASS:+} libgomp.oacc-c++/exceptions-throw-1.C -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O2 (test for excess errors) [-UNRESOLVED:-]{+PASS:+} libgomp.oacc-c++/exceptions-throw-1.C -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O2 [-compilation failed to produce executable-]{+execution test+} {+PASS: libgomp.oacc-c++/exceptions-throw-1.C -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O2 output pattern test+} [-XFAIL: libgomp.oacc-c++/exceptions-throw-2.C -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O2 (internal compiler error: Aborted signal terminated program as)-] [-XFAIL: libgomp.oacc-c++/exceptions-throw-2.C -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O2 PR119737 at line 9 (test for bogus messages, line )-] [-XFAIL:-]{+PASS:+} libgomp.oacc-c++/exceptions-throw-2.C -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O2 (test for excess errors) [-UNRESOLVED:-]{+PASS:+} libgomp.oacc-c++/exceptions-throw-2.C -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O2 [-compilation failed to produce executable-]{+execution test+} {+PASS: libgomp.oacc-c++/exceptions-throw-2.C -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa -O2 output pattern test+} PR target/119737 gcc/ * config/gcn/gcn.cc (gcn_hsa_declare_function_name): Properly switch sections. libgomp/ * 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. Co-authored-by: Thomas Schwinge <tschwinge@baylibre.com>
2025-04-25Adjust 'libgomp.c++/target-exceptions-pr118794-1.C' for ↵Thomas Schwinge3-6/+12
'targetm.arm_eabi_unwinder' [PR118794] Fix-up for commit aa3e72f943032e5f074b2bd2fd06d130dda8760b "Add test cases for exception handling constructs in dead code for GCN, nvptx target and OpenMP 'target' offloading [PR118794]": we need to adjust for configurations with 'targetm.arm_eabi_unwinder', as per: gcc/config/arm/arm.cc:#define TARGET_ARM_EABI_UNWINDER true gcc/config/c6x/c6x.cc:#define TARGET_ARM_EABI_UNWINDER true ..., which for ARM is conditional to '#if ARM_UNWIND_INFO' (defined in 'gcc/config/arm/bpabi.h', used for various GCC configurations), and for C6x unconditional. This gets us: --- target-exceptions-pr118794-1.C.269t.optimized +++ target-exceptions-pr118794-1.C.270t.optimized [...] __attribute__((omp declare target)) void f () [...] gimple_call <__dt_comp , NULL, &c> - gimple_call <__builtin_eh_pointer, _7, 2> - gimple_call <__builtin_unwind_resume, NULL, _7> + gimple_call <__builtin_cxa_end_cleanup, NULL> } [...] PR target/118794 libgomp/ * 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.
2025-04-24libgomp/testsuite: Fix hip_header_nvidia check, add workaround to testTobias Burnus4-4/+16
This is all about using the AMD's HIP header files with __HIP_PLATFORM_NVIDIA__ defined, i.e. HIP with Nvidia/CUDA; in that case, HIP is a thin layer on top of CUDA. First, the check_effective_target_gomp_hip_header_nvidia check failed; to fix it, -Wno-deprecated-declarations was added - and likewise to the two affected testcases that actually used the HIP headers on Nvidia. Doing so, the HIP tested was successful but the HIP-BLAS one showed two issues: * One seems to be related to include search paths as the HIP header uses #include "library_types.h" to include that CUDA header. Seemingly, it tried to included (again) the HIP header hip/library_types.h, not the CUDA one. I guess, some tweaking of -isystem vs. -I could have prevented this, but the simpler workaround was to just explicitly include the CUDA one before the HIP header files. * Once done, everything compiled but linking failed as the association between three HIP-BLAS functions and their CUDA-BLAS ones did not work. Solution: Just add three #define for mapping them. libgomp/ChangeLog: * 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-24libgomp: Add additional OpenMP interop runtime testsTobias Burnus23-2/+1397
Add checks for nowait/depend and for checks that the returned CUDA, CUDA_DRIVER and HIP interop objects actually work. While the CUDA/CUDA_DRIVER ones are only for Nvidia GPUs, HIP works on both AMD and Nvidia GPUs; on Nvidia GPUs, it is a very thin wrapper around CUDA. For Fortran, only a HIP test has been added - using hipfort. While libgomp.c-c++-common/interop-2.c always works - even without GPU - and checks for depend / nowait, all others require that runtime libraries are found at link (and execution) time: For Nvidia GPUs, libcuda + libcudart or libcublas, For AMD GPUs, libamdhip64 or libhipblas. The header files and hipfort modules do not need to be present as a fallback has been implemented, but if they are, they get used. Due to the combinations, the basic 1x C/C++, 4x C and 1x Fortran tests yield 1x C/C++, 14x C and 4 Fortran run-test files. libgomp/ChangeLog: * 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-23OpenMP: Add libgomp.fortran/target-enter-data-8.f90Tobias Burnus1-0/+532
Add another testcase for Fortran deep mapping of allocatable components. libgomp/ChangeLog: * testsuite/libgomp.fortran/target-enter-data-8.f90: New test.
2025-04-17libgomp: Don't test ompx::allocator::gnu_pinned_mem on non-linux targets.Jakub Jelinek2-0/+22
The libgomp.c/alloc-pinned*.c test have /* { dg-skip-if "Pinning not implemented on this host" { ! *-*-linux-gnu* } } */ so they are only run on Linux targets right now. Duplicating the tests or reworking them into headers looked like too much work for me right now this late in stage4, so I've just #ifdefed the uses at least for now. 2025-04-17 Jakub Jelinek <jakub@redhat.com> PR libgomp/119849 * testsuite/libgomp.c++/allocator-1.C (test_inequality, main): Guard ompx::allocator::gnu_pinned_mem uses with #ifdef __gnu_linux__. * testsuite/libgomp.c++/allocator-2.C (main): Likewise.
2025-04-16Remove 'ALWAYS_INLINE' workaround in ↵Thomas Schwinge1-6/+0
'libgomp.c++/target-exceptions-pr118794-1.C' With commit ca9cffe737d20953082333dacebb65d4261e0d0c "For nvptx offloading, make sure to emit C++ constructor, destructor aliases [PR97106]", we're able to remove the 'ALWAYS_INLINE' workaround added in commit fe283dba774be57b705a7a871b000d2894d2e553 "GCN, nvptx: Support '-mfake-exceptions', and use it for offloading compilation [PR118794]". libgomp/ * testsuite/libgomp.c++/target-exceptions-pr118794-1.C: Remove 'ALWAYS_INLINE' workaround.
2025-04-16Add 'libgomp.c++/pr106445-1{,-O0}.C' [PR106445]Thomas Schwinge2-0/+21
PR target/106445 libgomp/ * testsuite/libgomp.c++/pr106445-1.C: New. * testsuite/libgomp.c++/pr106445-1-O0.C: Likewise.
2025-04-16For nvptx offloading, make sure to emit C++ constructor, destructor aliases ↵Thomas Schwinge2-3/+1
[PR97106] PR target/97106 gcc/ * config/nvptx/nvptx.cc (nvptx_asm_output_def_from_decls) [ACCEL_COMPILER]: Make sure to emit C++ constructor, destructor aliases. libgomp/ * testsuite/libgomp.c++/pr96390.C: Un-XFAIL nvptx offloading. * testsuite/libgomp.c-c++-common/pr96390.c: Adjust.
2025-04-15OpenMP: omp.h omp::allocator C++ Allocator interfacewaffl3x2-0/+290
The implementation of each allocator is simplified by inheriting from __detail::__allocator_templ. At the moment, none of the implementations diverge in any way, simply passing in the allocator handle to be used when an allocation is made. In the future, const_mem will need special handling added to it to support constant memory space. libgomp/ChangeLog: * omp.h.in: Add omp::allocator::* and ompx::allocator::* allocators. (__detail::__allocator_templ<T, omp_allocator_handle_t>): New struct template. (null_allocator<T>): New struct template. (default_mem<T>): Likewise. (large_cap_mem<T>): Likewise. (const_mem<T>): Likewise. (high_bw_mem<T>): Likewise. (low_lat_mem<T>): Likewise. (cgroup_mem<T>): Likewise. (pteam_mem<T>): Likewise. (thread_mem<T>): Likewise. (ompx::allocator::gnu_pinned_mem<T>): Likewise. * testsuite/libgomp.c++/allocator-1.C: New test. * testsuite/libgomp.c++/allocator-2.C: New test. Signed-off-by: waffl3x <waffl3x@baylibre.com>
2025-04-15Fortran/OpenMP: Support automatic mapping allocatable components (deep mapping)Tobias Burnus8-0/+2158
When mapping an allocatable variable (or derived-type component), explicitly or implicitly, all its allocated allocatable components will automatically be mapped. The patch implements the target hooks, added for this feature to omp-low.cc with commit r15-3895-ge4a58b6f28383c. Namely, there is a check whether there are allocatable components at all: gfc_omp_deep_mapping_p. Then gfc_omp_deep_mapping_cnt, counting the number of required mappings; this is a dynamic value as it depends on array bounds and whether an allocatable is allocated or not. And, finally, the actual mapping: gfc_omp_deep_mapping. Polymorphic variables are partially supported: the mapping of the _data component is fully supported, but only components of the declared type are processed for additional allocatables. Additionally, _vptr is not touched. This means that everything needing _vtab information requires unified shared memory; in particular, _size data is required when accessing elements of polymorphic arrays. However, for scalar arrays, accessing components of the declare type should work just fine. As polymorphic variables are not (really) supported and OpenMP 6 explicitly disallows them, there is now a warning (-Wopenmp) when they are encountered. Unlimited polymorphics are rejected (error). Additionally, PRIVATE and FIRSTPRIVATE are not quite supported for allocatable components, polymorphic components and as polymorphic variable. Thus, those are now rejected as well. gcc/fortran/ChangeLog: * f95-lang.cc (LANG_HOOKS_OMP_DEEP_MAPPING, LANG_HOOKS_OMP_DEEP_MAPPING_P, LANG_HOOKS_OMP_DEEP_MAPPING_CNT): Define. * openmp.cc (gfc_match_omp_clause_reduction): Fix location setting. (resolve_omp_clauses): Permit allocatable components, reject them and polymorphic variables in PRIVATE/FIRSTPRIVATE. * trans-decl.cc (add_clause): Set clause location. * trans-openmp.cc (gfc_has_alloc_comps): Add ptr_ok and shallow_alloc_only Boolean arguments. (gfc_omp_replace_alloc_by_to_mapping): New. (gfc_omp_private_outer_ref, gfc_walk_alloc_comps, gfc_omp_clause_default_ctor, gfc_omp_clause_copy_ctor, gfc_omp_clause_assign_op, gfc_omp_clause_dtor): Update call to it. (gfc_omp_finish_clause): Minor cleanups, improve location data, handle allocatable components. (gfc_omp_deep_mapping_map, gfc_omp_deep_mapping_item, gfc_omp_deep_mapping_comps, gfc_omp_gen_simple_loop, gfc_omp_get_array_size, gfc_omp_elmental_loop, gfc_omp_deep_map_kind_p, gfc_omp_deep_mapping_int_p, gfc_omp_deep_mapping_p, gfc_omp_deep_mapping_do, gfc_omp_deep_mapping_cnt, gfc_omp_deep_mapping): New. (gfc_trans_omp_array_section): Save array descriptor in case deep-mapping lang hook will need it. (gfc_trans_omp_clauses): Likewise; use better clause location data. * trans.h (gfc_omp_deep_mapping_p, gfc_omp_deep_mapping_cnt, gfc_omp_deep_mapping): Add function prototypes. libgomp/ChangeLog: * libgomp.texi (5.0 Impl. Status): Mark mapping alloc comps as 'Y'. * 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. * testsuite/libgomp.fortran/map-alloc-comp-8.f90: New test. * testsuite/libgomp.fortran/map-alloc-comp-9.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/map-alloc-comp-1.f90: Remove dg-error. * gfortran.dg/gomp/polymorphic-mapping-2.f90: Update warn wording. * gfortran.dg/gomp/polymorphic-mapping.f90: Change expected diagnostic; some tests moved to ... * gfortran.dg/gomp/polymorphic-mapping-1.f90: ... here as new test. * gfortran.dg/gomp/polymorphic-mapping-3.f90: New test. * gfortran.dg/gomp/polymorphic-mapping-4.f90: New test. * gfortran.dg/gomp/polymorphic-mapping-5.f90: New test.
2025-04-14GCN, nvptx: Support '-mfake-exceptions', and use it for offloading ↵Thomas Schwinge16-54/+140
compilation [PR118794] With '-mfake-exceptions' enabled, the user-visible behavior in presence of exception handling constructs changes such that the compile-time 'sorry, unimplemented: exception handling not supported' is skipped, code generation proceeds, and instead, exception handling constructs 'abort' at run time. (..., or don't, if they're in dead code.) PR target/118794 gcc/ * config/gcn/gcn.opt (-mfake-exceptions): Support. * config/nvptx/nvptx.opt (-mfake-exceptions): Likewise. * config/gcn/gcn.md (define_expand "exception_receiver"): Use it. * config/nvptx/nvptx.md (define_expand "exception_receiver"): Likewise. * config/gcn/mkoffload.cc (main): Set it. * config/nvptx/mkoffload.cc (main): Likewise. * config/nvptx/nvptx.cc (nvptx_assemble_integer) <in_section == exception_section>: Special handling for 'SYMBOL_REF's. * except.cc (expand_dw2_landing_pad_for_region): Don't generate bogus code for (default) '#define EH_RETURN_DATA_REGNO(N) INVALID_REGNUM'. libgcc/ * config/gcn/unwind-gcn.c (_Unwind_Resume): New. * config/nvptx/unwind-nvptx.c (_Unwind_Resume): Likewise. gcc/testsuite/ * g++.target/gcn/exceptions-bad_cast-2.C: Set '-mno-fake-exceptions'. * g++.target/gcn/exceptions-pr118794-1.C: Likewise. * g++.target/gcn/exceptions-throw-2.C: Likewise. * g++.target/nvptx/exceptions-bad_cast-2.C: Likewise. * g++.target/nvptx/exceptions-pr118794-1.C: Likewise. * g++.target/nvptx/exceptions-throw-2.C: Likewise. * g++.target/gcn/exceptions-bad_cast-2_-mfake-exceptions.C: New. * g++.target/gcn/exceptions-pr118794-1_-mfake-exceptions.C: Likewise. * g++.target/gcn/exceptions-throw-2_-mfake-exceptions.C: Likewise. * g++.target/nvptx/exceptions-bad_cast-2_-mfake-exceptions.C: Likewise. * g++.target/nvptx/exceptions-pr118794-1_-mfake-exceptions.C: Likewise. * g++.target/nvptx/exceptions-throw-2_-mfake-exceptions.C: Likewise. libgomp/ * 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.
2025-04-14Add 'throw', dead code test cases for GCN, nvptx target and OpenACC, OpenMP ↵Thomas Schwinge2-0/+62
'target' offloading gcc/testsuite/ * g++.target/gcn/exceptions-throw-3.C: New. * g++.target/nvptx/exceptions-throw-3.C: Likewise. libgomp/ * testsuite/libgomp.c++/target-exceptions-throw-3.C: New. * testsuite/libgomp.oacc-c++/exceptions-throw-3.C: Likewise.
2025-04-14Add 'throw', caught test cases for GCN, nvptx target and OpenACC, OpenMP ↵Thomas Schwinge6-0/+134
'target' offloading gcc/testsuite/ * g++.target/gcn/exceptions-throw-2.C: New. * g++.target/nvptx/exceptions-throw-2.C: Likewise. libgomp/ * 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.
2025-04-14Add 'throw' test cases for GCN, nvptx target and OpenACC, OpenMP 'target' ↵Thomas Schwinge3-0/+94
offloading gcc/testsuite/ * g++.target/gcn/exceptions-throw-1.C: New. * g++.target/nvptx/exceptions-throw-1.C: Likewise. libgomp/ * 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.
generated by cgit v1.1 at 2025-08-03 22:34:42 +0000