aboutsummaryrefslogtreecommitdiff
path: root/libgomp
AgeCommit message (Collapse)AuthorFilesLines
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-16Daily bump.GCC Administrator1-0/+4
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-15Daily bump.GCC Administrator1-0/+13
2025-05-14OpenMP: Fix mapping of zero-sized arrays with non-literal size: ↵Tobias Burnus6-16/+288
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-13Daily bump.GCC Administrator1-0/+15
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-10Daily bump.GCC Administrator1-0/+21
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-08Daily bump.GCC Administrator1-0/+6
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-07Daily bump.GCC Administrator1-0/+4
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-06Daily bump.GCC Administrator1-0/+4
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-02Daily bump.GCC Administrator1-0/+4
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-26Daily bump.GCC Administrator1-0/+39
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 Schwinge6-3/+315
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-25Daily bump.GCC Administrator1-0/+50
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-18Daily bump.GCC Administrator1-0/+12
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-17libgomp.texi: For HIP interop, mention cpp defines to setTobias Burnus1-0/+6
The HIP header files recognize the used compiler, defaulting to either AMD or Nvidia/CUDA; thus, the alternative way of explicitly defining a macro is less prominently documented. With GCC, the user has to define the preprocessor macro manually. Hence, as a service to the user, mention __HIP_PLATFORM_AMD__ and __HIP_PLATFORM_NVIDIA__ in the interop documentation, even though it has only indirectly to do with GCC and its interop support. Note to commit-log readers, only: For Fortran, the hipfort modules can be used; when compiling the hipfort package (defaults to use gfortran), it generates the module (*.mod) files in include/hipfort/{amdgcn,nvidia}/ such that the choice is made by setting the respective include path. libgomp/ChangeLog: * libgomp.texi (gcn interop, nvptx interop): For HIP with C/C++, add a note about setting a preprocessor define.
2025-04-17Daily bump.GCC Administrator1-0/+17
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-16Daily bump.GCC Administrator1-0/+35
2025-04-15libgomp.texi (gcn, nvptx): Mention self_maps alongside USMTobias Burnus1-2/+2
libgomp/ChangeLog: * libgomp.texi (gcn, nvptx): Mention self_maps clause besides unified_shared_memory in the requirements item.
2025-04-15OpenMP: omp.h omp::allocator C++ Allocator interfacewaffl3x3-0/+422
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 Burnus9-1/+2159
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-15Daily bump.GCC Administrator1-0/+90
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.
2025-04-14Add 'std::bad_cast' exception, dead code test cases for GCN, nvptx target ↵Thomas Schwinge2-0/+66
and OpenACC, OpenMP 'target' offloading gcc/testsuite/ * g++.target/gcn/exceptions-bad_cast-3.C: New. * g++.target/nvptx/exceptions-bad_cast-3.C: Likewise. libgomp/ * testsuite/libgomp.c++/target-exceptions-bad_cast-3.C: New. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-3.C: Likewise.
2025-04-14Add 'std::bad_cast' exception, caught test cases for GCN, nvptx target and ↵Thomas Schwinge6-0/+134
OpenACC, OpenMP 'target' offloading gcc/testsuite/ * g++.target/gcn/exceptions-bad_cast-2.C: New. * g++.target/nvptx/exceptions-bad_cast-2.C: Likewise. libgomp/ * testsuite/libgomp.c++/target-exceptions-bad_cast-2.C: New. * testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C: Likewise.
2025-04-14Add 'std::bad_cast' exception test cases for GCN, nvptx target and OpenACC, ↵Thomas Schwinge2-0/+79
OpenMP 'target' offloading gcc/testsuite/ * g++.target/gcn/exceptions-bad_cast-1.C: New. * g++.target/nvptx/exceptions-bad_cast-1.C: Likewise. libgomp/ * testsuite/libgomp.c++/target-exceptions-bad_cast-1.C: New. * testsuite/libgomp.oacc-c++/exceptions-bad_cast-1.C: Likewise.
2025-04-14Add test cases for exception handling constructs in dead code for GCN, nvptx ↵Thomas Schwinge3-0/+104
target and OpenMP 'target' offloading [PR118794] PR target/118794 gcc/testsuite/ * g++.target/gcn/exceptions-pr118794-1.C: New. * g++.target/nvptx/exceptions-pr118794-1.C: Likewise. libgomp/ * testsuite/libgomp.c++/target-exceptions-pr118794-1.C: New. * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C: Likewise. * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C: Likewise.
generated by cgit v1.1 at 2025-09-18 16:17:59 +0000