aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite
AgeCommit message (Collapse)AuthorFilesLines
2022-07-10c: Fix location for _Pragma tokens [PR97498]Lewis Hyatt2-11/+11
The handling of #pragma GCC diagnostic uses input_location, which is not always as precise as needed; in particular the relative location of some tokens and a _Pragma directive will crucially determine whether a given diagnostic is enabled or suppressed in the desired way. PR97498 shows how the C frontend ends up with input_location pointing to the beginning of the line containing a _Pragma() directive, resulting in the wrong behavior if the diagnostic to be modified pertains to some tokens found earlier on the same line. This patch fixes that by addressing two issues: a) libcpp was not assigning a valid location to the CPP_PRAGMA token generated by the _Pragma directive. b) C frontend was not setting input_location to something reasonable. With this change, the C frontend is able to change input_location to point to the _Pragma token as needed. This is just a two-line fix (one for each of a) and b)), the testsuite changes were needed only because the location on the tested warnings has been somewhat improved, so the tests need to look for the new locations. gcc/c/ChangeLog: PR preprocessor/97498 * c-parser.cc (c_parser_pragma): Set input_location to the location of the pragma, rather than the start of the line. libcpp/ChangeLog: PR preprocessor/97498 * directives.cc (destringize_and_run): Override the location of the CPP_PRAGMA token from a _Pragma directive to the location of the expansion point, as is done for the tokens lexed from it. gcc/testsuite/ChangeLog: PR preprocessor/97498 * c-c++-common/pr97498.c: New test. * c-c++-common/gomp/pragma-3.c: Adapt for improved warning locations. * c-c++-common/gomp/pragma-5.c: Likewise. * gcc.dg/pragma-message.c: Likewise. libgomp/ChangeLog: * testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Adapt for improved warning locations. * testsuite/libgomp.oacc-c-c++-common/vred2d-128.c: Likewise.
2022-07-08Fix one issue in OpenMP 'requires' directive diagnosticsThomas Schwinge5-4/+26
Fix-up for recent commit 683f11843974f0bdf42f79cdcbb0c2b43c7b81b0 "OpenMP: Move omp requires checks to libgomp". gcc/ * lto-cgraph.cc (input_offload_tables) <LTO_symtab_edge>: Correct 'fn2' computation. libgomp/ * testsuite/libgomp.c-c++-common/requires-1.c: Add 'dg-note's. * testsuite/libgomp.c-c++-common/requires-2.c: Likewise. * testsuite/libgomp.c-c++-common/requires-3.c: Likewise. * testsuite/libgomp.c-c++-common/requires-7.c: Likewise. * testsuite/libgomp.fortran/requires-1.f90: Likewise.
2022-07-07Fix Intel MIC 'mkoffload' for OpenMP 'requires'Thomas Schwinge5-4/+9
Similar to how the other 'mkoffload's got changed in recent commit 683f11843974f0bdf42f79cdcbb0c2b43c7b81b0 "OpenMP: Move omp requires checks to libgomp". This also means finally switching Intel MIC 'mkoffload' to 'GOMP_offload_register_ver', 'GOMP_offload_unregister_ver', making 'GOMP_offload_register', 'GOMP_offload_unregister' legacy entry points. gcc/ * config/i386/intelmic-mkoffload.cc (generate_host_descr_file) (prepare_target_image, main): Handle OpenMP 'requires'. (generate_host_descr_file): Switch to 'GOMP_offload_register_ver', 'GOMP_offload_unregister_ver'. libgomp/ * target.c (GOMP_offload_register, GOMP_offload_unregister): Denote as legacy entry points. * testsuite/lib/libgomp.exp (check_effective_target_offload_target_any): New proc. * testsuite/libgomp.c-c++-common/requires-1.c: Enable for 'offload_target_any'. * testsuite/libgomp.c-c++-common/requires-3.c: Likewise. * testsuite/libgomp.c-c++-common/requires-7.c: Likewise. * testsuite/libgomp.fortran/requires-1.f90: Likewise.
2022-07-07Enhance 'libgomp.c-c++-common/requires-4.c', ↵Thomas Schwinge2-12/+31
'libgomp.c-c++-common/requires-5.c' testing These should compile and link and execute in all configurations; host-fallback execution, which we may actually verify. Follow-up to recent commit 683f11843974f0bdf42f79cdcbb0c2b43c7b81b0 "OpenMP: Move omp requires checks to libgomp". libgomp/ * testsuite/libgomp.c-c++-common/requires-4.c: Enhance testing. * testsuite/libgomp.c-c++-common/requires-5.c: Likewise.
2022-07-07Adjust 'libgomp.c-c++-common/requires-3.c'Thomas Schwinge1-1/+1
As documented, this one does "Check diagnostic by device-compiler's lto1". Indeed there are none when compiling with '-foffload=disable' with an offloading-enabled compiler, so we should use 'offload_target_[...]', as used in other similar test cases. Follow-up to recent commit 683f11843974f0bdf42f79cdcbb0c2b43c7b81b0 "OpenMP: Move omp requires checks to libgomp". libgomp/ * testsuite/libgomp.c-c++-common/requires-3.c: Adjust.
2022-07-04OpenMP: Move omp requires checks to libgompTobias Burnus15-0/+264
Handle reverse_offload, unified_address, and unified_shared_memory requirements in libgomp by saving them alongside the offload table. When the device lto1 runs, it extracts the data for mkoffload. The latter than passes the value on to GOMP_offload_register_ver. lto1 (either the host one, with -flto [+ ENABLE_OFFLOADING], or in the offload-device lto1) also does the the consistency check is done, erroring out when the 'omp requires' clause use is inconsistent. For all in-principle supported devices, if a requirement cannot be fulfilled, the device is excluded from the (supported) devices list. Currently, none of those requirements are marked as supported for any of the non-host devices. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_target_data, c_parser_omp_target_update, c_parser_omp_target_enter_data, c_parser_omp_target_exit_data): Set OMP_REQUIRES_TARGET_USED. (c_parser_omp_requires): Remove sorry. gcc/ChangeLog: * config/gcn/mkoffload.cc (process_asm): Write '#include <stdint.h>'. (process_obj): Pass omp_requires_mask to GOMP_offload_register_ver. (main): Ask lto1 to obtain omp_requires_mask and pass it on. * config/nvptx/mkoffload.cc (process, main): Likewise. * lto-cgraph.cc (omp_requires_to_name): New. (input_offload_tables): Save omp_requires_mask. (output_offload_tables): Read it, check for consistency, save value for mkoffload. * omp-low.cc (lower_omp_target): Force output_offloadtables call for OMP_REQUIRES_TARGET_USED. gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_target_data, cp_parser_omp_target_enter_data, cp_parser_omp_target_exit_data, cp_parser_omp_target_update): Set OMP_REQUIRES_TARGET_USED. (cp_parser_omp_requires): Remove sorry. gcc/fortran/ChangeLog: * openmp.cc (gfc_match_omp_requires): Remove sorry. * parse.cc (decode_omp_directive): Don't regard 'declare target' as target usage for 'omp requires'; add more flags to omp_requires_mask. include/ChangeLog: * gomp-constants.h (GOMP_VERSION): Bump to 2. (GOMP_REQUIRES_UNIFIED_ADDRESS, GOMP_REQUIRES_UNIFIED_SHARED_MEMORY, GOMP_REQUIRES_REVERSE_OFFLOAD, GOMP_REQUIRES_TARGET_USED): New defines. libgomp/ChangeLog: * libgomp-plugin.h (GOMP_OFFLOAD_get_num_devices): Add omp_requires_mask arg. * plugin/plugin-gcn.c (GOMP_OFFLOAD_get_num_devices): Likewise; return -1 when device available but omp_requires_mask != 0. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_get_num_devices): Likewise. * oacc-host.c (host_get_num_devices, host_openacc_get_property): Update call. * oacc-init.c (resolve_device, acc_init_1, acc_shutdown_1, goacc_attach_host_thread_to_device, acc_get_num_devices, acc_set_device_num, get_property_any): Likewise. * target.c (omp_requires_mask): New global var. (gomp_requires_to_name): New. (GOMP_offload_register_ver): Handle passed omp_requires_mask. (gomp_target_init): Handle omp_requires_mask. * libgomp.texi (OpenMP 5.0): Update requires impl. status. (OpenMP 5.1): Add a missed item. (OpenMP 5.2): Mark linear-clause change as supported in C/C++. * testsuite/libgomp.c-c++-common/requires-1-aux.c: New test. * testsuite/libgomp.c-c++-common/requires-1.c: New test. * testsuite/libgomp.c-c++-common/requires-2-aux.c: New test. * testsuite/libgomp.c-c++-common/requires-2.c: New test. * testsuite/libgomp.c-c++-common/requires-3-aux.c: New test. * testsuite/libgomp.c-c++-common/requires-3.c: New test. * testsuite/libgomp.c-c++-common/requires-4-aux.c: New test. * testsuite/libgomp.c-c++-common/requires-4.c: New test. * testsuite/libgomp.c-c++-common/requires-5-aux.c: New test. * testsuite/libgomp.c-c++-common/requires-5.c: New test. * testsuite/libgomp.c-c++-common/requires-6.c: New test. * testsuite/libgomp.c-c++-common/requires-7-aux.c: New test. * testsuite/libgomp.c-c++-common/requires-7.c: New test. * testsuite/libgomp.fortran/requires-1-aux.f90: New test. * testsuite/libgomp.fortran/requires-1.f90: New test. liboffloadmic/ChangeLog: * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_num_devices): Return -1 when device available but omp_requires_mask != 0. gcc/testsuite/ChangeLog: * c-c++-common/gomp/requires-4.c: Update dg-*. * c-c++-common/gomp/reverse-offload-1.c: Likewise. * c-c++-common/gomp/target-device-ancestor-2.c: Likewise. * c-c++-common/gomp/target-device-ancestor-3.c: Likewise. * c-c++-common/gomp/target-device-ancestor-4.c: Likewise. * c-c++-common/gomp/target-device-ancestor-5.c: Likewise. * gfortran.dg/gomp/target-device-ancestor-3.f90: Likewise. * gfortran.dg/gomp/target-device-ancestor-4.f90: Likewise. * gfortran.dg/gomp/target-device-ancestor-5.f90: Likewise. * gfortran.dg/gomp/target-device-ancestor-2.f90: Likewise. Move post-FE checks to ... * gfortran.dg/gomp/target-device-ancestor-2a.f90: ... this new file. * gfortran.dg/gomp/requires-8.f90: Update as we don't regard 'declare target' for the 'requires' usage requirement. Co-authored-by: Chung-Lin Tang <cltang@codesourcery.com> Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
2022-06-21libgomp: Fix up target-31.c test [PR106045]Jakub Jelinek1-1/+1
The i variable is used inside of the parallel in: #pragma omp simd safelen(32) private (v) for (i = 0; i < 64; i++) { v = 3 * i; ll[i] = u1 + v * u2[0] + u2[1] + x + y[0] + y[1] + v + h[0] + u3[i]; } where i is predetermined linear (so while inside of the body it is safe, private per SIMD lane var) the final value is written to the shared variable, and in: for (i = 0; i < 64; i++) if (ll[i] != u1 + 3 * i * u2[0] + u2[1] + x + y[0] + y[1] + 3 * i + 13 + 14 + i) #pragma omp atomic write err = 1; which is a normal loop and so it isn't in any way privatized there. So we have a data race, fixed by adding private (i) clause to the parallel. 2022-06-21 Jakub Jelinek <jakub@redhat.com> Paul Iannetta <piannetta@kalrayinc.com> PR libgomp/106045 * testsuite/libgomp.c/target-31.c: Add private (i) clause.
2022-06-15openmp: Fix up get-mapped-ptr-1.{c,f90} testsJakub Jelinek2-6/+16
On Tue, Jun 14, 2022 at 06:41:37PM +0200, Thomas Schwinge wrote: > In an offloading configuration, I'm seeing: > > PASS: libgomp.fortran/get-mapped-ptr-1.f90 -O (test for excess errors) > [-PASS:-]{+FAIL:+} libgomp.fortran/get-mapped-ptr-1.f90 -O execution test > > Does that one need similar treatment? I assume not just that but libgomp.c-c++-common/get-mapped-ptr-1.c too? It both needs the same treatment, and in the get-mapped-ptr-1.c case there is even UB, while the Fortran version was using c_loc (q) as the host pointer, in C/C++ it was using q which was value of uninitialized pointer. 2022-06-15 Jakub Jelinek <jakub@redhat.com> * testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c (main): Initialize q to ddress of an automatic variable. Use -5 instead of -1 in omp_get_mapped_ptr call. Add test with omp_initial_device. * testsuite/libgomp.fortran/get-mapped-ptr-1.f90 (main): Use -5 instead of -1 in omp_get_mapped_ptr call. Add test with omp_initial_device. Renumber stop arguments afterwards.
2022-06-13openmp: Conforming device numbers and omp_{initial,invalid}_deviceJakub Jelinek6-9/+75
OpenMP 5.2 changed once more what device numbers are allowed. In 5.1, valid device numbers were [0, omp_get_num_devices()]. 5.2 makes also -1 valid (calls it omp_initial_device), which is equivalent in behavior to omp_get_num_devices() number but has the advantage that it is a constant. And it also introduces omp_invalid_device which is also a constant with implementation defined value < -1. That value should act like sNaN, any time any device construct (GOMP_target*) or OpenMP runtime API routine is asked for such a device, the program is terminated. And if OMP_TARGET_OFFLOAD=mandatory, all non-conforming device numbers (which is all but [-1, omp_get_num_devices()] other than omp_invalid_device) must be treated like omp_invalid_device. For device constructs, we have a compatibility problem, we've historically used 2 magic negative values to mean something special. GOMP_DEVICE_ICV (-1) means device clause wasn't present, pick the omp_get_default_device () number GOMP_DEVICE_FALLBACK (-2) means the host device (this is used e.g. for #pragma omp target if (cond) where if cond is false, we pass -2 But 5.2 requires that omp_initial_device is -1 (there were discussions about it, advantage of -1 is that one can say iterate over the [-1, omp_get_num_devices()-1] range to get all devices starting with the host/initial one. And also, if user passes -2, unless it is omp_invalid_device, we need to treat it like non-conforming with OMP_TARGET_OFFLOAD=mandatory. So, the patch does on the compiler side some number remapping, user_device_num >= -2U ? user_device_num - 1 : user_device_num. This remapping is done at compile time if device clause has constant argument, otherwise at runtime, and means that for user -1 (omp_initial_device) we pass -2 to GOMP_* in the runtime library where it treats it like host fallback, while -2 is remapped to -3 (one of the non-conforming device numbers, for those it doesn't matter which one is which). omp_invalid_device is then -4. For the OpenMP device runtime APIs, no remapping is done. This patch doesn't deal with the initial default-device-var for OMP_TARGET_OFFLOAD=mandatory , the spec says that the inital ICV value for that should in that case depend on whether there are any offloading devices or not (if not, should be omp_invalid_device), but that means we can't determine the number of devices lazily (and let libraries have the possibility to register their offloading data etc.). 2022-06-13 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-expand.cc (expand_omp_target): Remap user provided device clause arguments, -1 to -2 and -2 to -3, either at compile time if constant, or at runtime. include/ * gomp-constants.h (GOMP_DEVICE_INVALID): Define. libgomp/ * omp.h.in (omp_initial_device, omp_invalid_device): New enumerators. * omp_lib.f90.in (omp_initial_device, omp_invalid_device): New parameters. * omp_lib.h.in (omp_initial_device, omp_invalid_device): Likewise. * target.c (resolve_device): Add remapped argument, handle GOMP_DEVICE_ICV only if remapped is true (and clear remapped), for negative values, treat GOMP_DEVICE_FALLBACK as fallback only if remapped, otherwise treat omp_initial_device that way. For omp_invalid_device, always emit gomp_fatal, even when OMP_TARGET_OFFLOAD isn't mandatory. (GOMP_target, GOMP_target_ext, GOMP_target_data, GOMP_target_data_ext, GOMP_target_update, GOMP_target_update_ext, GOMP_target_enter_exit_data): Pass true as remapped argument to resolve_device. (omp_target_alloc, omp_target_free, omp_target_is_present, omp_target_memcpy_check, omp_target_associate_ptr, omp_target_disassociate_ptr, omp_get_mapped_ptr, omp_target_is_accessible): Pass false as remapped argument to resolve_device. Treat omp_initial_device the same as gomp_get_num_devices (). Don't bypass resolve_device calls if device_num is negative. (omp_pause_resource): Treat omp_initial_device the same as gomp_get_num_devices (). Call resolve_device. * icv-device.c (omp_set_default_device): Always set to device_num even when it is negative. * libgomp.texi: Document that Conforming device numbers, omp_initial_device and omp_invalid_device is implemented. * testsuite/libgomp.c/target-41.c (main): Add test with omp_initial_device. * testsuite/libgomp.c/target-45.c: New test. * testsuite/libgomp.c/target-46.c: New test. * testsuite/libgomp.c/target-47.c: New test. * testsuite/libgomp.c-c++-common/target-is-accessible-1.c (main): Add test with omp_initial_device. Use -5 instead of -1 for negative value test. * testsuite/libgomp.fortran/target-is-accessible-1.f90 (main): Likewise. Reorder stop numbers.
2022-06-10libgomp nvptx plugin: Remove '--with-cuda-driver=[...]' etc. configuration ↵Thomas Schwinge3-21/+0
option That means, exposing to the user only the '--without-cuda-driver' behavior: including the GCC-shipped 'include/cuda/cuda.h' (not system <cuda.h>), and 'dlopen'ing the CUDA Driver library (not linking it). For development purposes, the libgomp nvptx plugin developer may still manually override that, to get the previous '--with-cuda-driver' behavior. libgomp/ * plugin/Makefrag.am: Evaluate 'if PLUGIN_NVPTX_DYNAMIC' to true. * plugin/configfrag.ac (--with-cuda-driver) (--with-cuda-driver-include, --with-cuda-driver-lib) (CUDA_DRIVER_INCLUDE, CUDA_DRIVER_LIB, PLUGIN_NVPTX_CPPFLAGS) (PLUGIN_NVPTX_LDFLAGS, PLUGIN_NVPTX_LIBS, PLUGIN_NVPTX_DYNAMIC): Remove. * testsuite/libgomp-test-support.exp.in (cuda_driver_include) (cuda_driver_lib): Remove. * testsuite/lib/libgomp.exp (libgomp_init): Don't consider these. * Makefile.in: Regenerate. * configure: Likewise. * testsuite/Makefile.in: Likewise.
2022-06-03OpenMP/Fortran: Add support for firstprivate and allocate clauses on scope ↵Tobias Burnus1-0/+57
construct Fortran commit to C/C++/backend commit r13-862-gf38b20d68fade5a922b9f68c4c3841e653d1b83c gcc/fortran/ChangeLog: * openmp.cc (OMP_SCOPE_CLAUSES): Add firstprivate and allocate. libgomp/ChangeLog: * libgomp.texi (OpenMP 5.2): Mark scope w/ firstprivate/allocate as Y. * testsuite/libgomp.fortran/scope-2.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/scope-5.f90: New test. * gfortran.dg/gomp/scope-6.f90: New test.
2022-06-02diagnostics: add SARIF output formatDavid Malcolm1-0/+1
This patch adds support to gcc's diagnostic subsystem for emitting diagnostics in SARIF, aka the Static Analysis Results Interchange Format: https://sarifweb.azurewebsites.net/ by extending -fdiagnostics-format= to add two new options: -fdiagnostics-format=sarif-stderr and: -fdiagnostics-format=sarif-file The patch targets SARIF v2.1.0 This is a JSON-based format suited for capturing the results of static analysis tools (like GCC's -fanalyzer), but it can also be used for plain GCC warnings and errors. SARIF supports per-event metadata in diagnostic paths such as ["acquire", "resource"] and ["release", "lock"] (specifically, the threadFlowLocation "kinds" property: SARIF v2.1.0 section 3.38.8), so the patch extends GCC"s diagnostic_event subclass with a "struct meaning" with similar purpose. The patch implements this for -fanalyzer so that the various state-machine-based warnings set these in the SARIF output. The heart of the implementation is in the new file diagnostic-format-sarif.cc. Much of the rest of the patch is interface classes, isolating the diagnostic subsystem (which has no knowledge of e.g. tree or langhook) from the "client" code in the compiler proper cc1 etc). The patch adds a langhook for specifying the SARIF v2.1.0 "artifact.sourceLanguage" property, based on the list in SARIF v2.1.0 Appendix J. The patch adds automated DejaGnu tests to our testsuite via new scan-sarif-file and scan-sarif-file-not directives (although these merely use regexps, rather than attempting to use a proper JSON parser). I've tested the patch by hand using the validator at: https://sarifweb.azurewebsites.net/Validation and the react-based viewer at: https://microsoft.github.io/sarif-web-component/ which successfully shows most of the information (although not paths, and not CWE IDs), and I've fixed all validation errors I've seen (though bugs no doubt remain). I've also tested the generated SARIF using the VS Code extension linked to from the SARIF website; I'm a novice with VS Code, but it seems to be able to handle my generated SARIF files (e.g. showing the data in the SARIF tab, and showing squiggly underlines under issues, and when I click on them, it visualizes the events in the path inline within the source window). Has anyone written an Emacs mode for SARIF files? (pretty please) gcc/ChangeLog: * Makefile.in (OBJS): Add tree-diagnostic-client-data-hooks.o and tree-logical-location.o. (OBJS-libcommon): Add diagnostic-format-sarif.o; reorder. (CFLAGS-tree-diagnostic-client-data-hooks.o): Add TARGET_NAME. * common.opt (fdiagnostics-format=): Add sarif-stderr and sarif-file. (sarif-stderr, sarif-file): New enum values. * diagnostic-client-data-hooks.h: New file. * diagnostic-format-sarif.cc: New file. * diagnostic-path.h (enum diagnostic_event::verb): New enum. (enum diagnostic_event::noun): New enum. (enum diagnostic_event::property): New enum. (struct diagnostic_event::meaning): New struct. (diagnostic_event::get_logical_location): New vfunc. (diagnostic_event::get_meaning): New vfunc. (simple_diagnostic_event::get_logical_location): New vfunc impl. (simple_diagnostic_event::get_meaning): New vfunc impl. * diagnostic.cc: Include "diagnostic-client-data-hooks.h". (diagnostic_initialize): Initialize m_client_data_hooks. (diagnostic_finish): Clean up m_client_data_hooks. (diagnostic_event::meaning::dump_to_pp): New. (diagnostic_event::meaning::maybe_get_verb_str): New. (diagnostic_event::meaning::maybe_get_noun_str): New. (diagnostic_event::meaning::maybe_get_property_str): New. (get_cwe_url): Make non-static. (diagnostic_output_format_init): Handle DIAGNOSTICS_OUTPUT_FORMAT_SARIF_STDERR and DIAGNOSTICS_OUTPUT_FORMAT_SARIF_FILE. * diagnostic.h (enum diagnostics_output_format): Add DIAGNOSTICS_OUTPUT_FORMAT_SARIF_STDERR and DIAGNOSTICS_OUTPUT_FORMAT_SARIF_FILE. (class diagnostic_client_data_hooks): New forward decl. (class logical_location): New forward decl. (diagnostic_context::m_client_data_hooks): New field. (diagnostic_output_format_init_sarif_stderr): New decl. (diagnostic_output_format_init_sarif_file): New decl. (get_cwe_url): New decl. * doc/invoke.texi (-fdiagnostics-format=): Add sarif-stderr and sarif-file. * doc/sourcebuild.texi (Scan a particular file): Add scan-sarif-file and scan-sarif-file-not. * langhooks-def.h (lhd_get_sarif_source_language): New decl. (LANG_HOOKS_GET_SARIF_SOURCE_LANGUAGE): New macro. (LANG_HOOKS_INITIALIZER): Add LANG_HOOKS_GET_SARIF_SOURCE_LANGUAGE. * langhooks.cc (lhd_get_sarif_source_language): New. * langhooks.h (lang_hooks::get_sarif_source_language): New field. * logical-location.h: New file. * plugin.cc (struct for_each_plugin_closure): New. (for_each_plugin_cb): New. (for_each_plugin): New. * plugin.h (for_each_plugin): New decl. * tree-diagnostic-client-data-hooks.cc: New file. * tree-diagnostic.cc: Include "diagnostic-client-data-hooks.h". (tree_diagnostics_defaults): Populate m_client_data_hooks. * tree-logical-location.cc: New file. * tree-logical-location.h: New file. gcc/ada/ChangeLog: * gcc-interface/misc.cc (gnat_get_sarif_source_language): New. (LANG_HOOKS_GET_SARIF_SOURCE_LANGUAGE): Redefine. gcc/analyzer/ChangeLog: * checker-path.cc (checker_event::get_meaning): New. (function_entry_event::get_meaning): New. (state_change_event::get_desc): Add dump of meaning of the event to the -fanalyzer-verbose-state-changes output. (state_change_event::get_meaning): New. (cfg_edge_event::get_meaning): New. (call_event::get_meaning): New. (return_event::get_meaning): New. (start_consolidated_cfg_edges_event::get_meaning): New. (warning_event::get_meaning): New. * checker-path.h: Include "tree-logical-location.h". (checker_event::checker_event): Construct m_logical_loc. (checker_event::get_logical_location): New. (checker_event::get_meaning): New decl. (checker_event::m_logical_loc): New. (function_entry_event::get_meaning): New decl. (state_change_event::get_meaning): New decl. (cfg_edge_event::get_meaning): New decl. (call_event::get_meaning): New decl. (return_event::get_meaning): New decl. (start_consolidated_cfg_edges_event::get_meaning): New. (warning_event::get_meaning): New decl. * pending-diagnostic.h: Include "diagnostic-path.h". (pending_diagnostic::get_meaning_for_state_change): New vfunc. * sm-file.cc (file_diagnostic::get_meaning_for_state_change): New vfunc impl. * sm-malloc.cc (malloc_diagnostic::get_meaning_for_state_change): Likewise. * sm-sensitive.cc (exposure_through_output_file::get_meaning_for_state_change): Likewise. * sm-taint.cc (taint_diagnostic::get_meaning_for_state_change): Likewise. * varargs.cc (va_list_sm_diagnostic::get_meaning_for_state_change): Likewise. gcc/c/ChangeLog: * c-lang.cc (LANG_HOOKS_GET_SARIF_SOURCE_LANGUAGE): Redefine. (c_get_sarif_source_language): New. * c-tree.h (c_get_sarif_source_language): New decl. gcc/cp/ChangeLog: * cp-lang.cc (LANG_HOOKS_GET_SARIF_SOURCE_LANGUAGE): Redefine. (cp_get_sarif_source_language): New. gcc/d/ChangeLog: * d-lang.cc (d_get_sarif_source_language): New. (LANG_HOOKS_GET_SARIF_SOURCE_LANGUAGE): Redefine. gcc/fortran/ChangeLog: * f95-lang.cc (gfc_get_sarif_source_language): New. (LANG_HOOKS_GET_SARIF_SOURCE_LANGUAGE): Redefine. gcc/go/ChangeLog: * go-lang.cc (go_get_sarif_source_language): New. (LANG_HOOKS_GET_SARIF_SOURCE_LANGUAGE): Redefine. gcc/objc/ChangeLog: * objc-act.h (objc_get_sarif_source_language): New decl. * objc-lang.cc (LANG_HOOKS_GET_SARIF_SOURCE_LANGUAGE): Redefine. (objc_get_sarif_source_language): New. gcc/testsuite/ChangeLog: * c-c++-common/diagnostic-format-sarif-file-1.c: New test. * c-c++-common/diagnostic-format-sarif-file-2.c: New test. * c-c++-common/diagnostic-format-sarif-file-3.c: New test. * c-c++-common/diagnostic-format-sarif-file-4.c: New test. * gcc.dg/analyzer/file-meaning-1.c: New test. * gcc.dg/analyzer/malloc-meaning-1.c: New test. * gcc.dg/analyzer/malloc-sarif-1.c: New test. * gcc.dg/plugin/analyzer_gil_plugin.c (gil_diagnostic::get_meaning_for_state_change): New vfunc impl. * gcc.dg/plugin/diagnostic-test-paths-5.c: New test. * gcc.dg/plugin/plugin.exp (plugin_test_list): Add diagnostic-test-paths-5.c to tests for diagnostic_plugin_test_paths.c. * lib/gcc-dg.exp: Load scansarif.exp. * lib/scansarif.exp: New test. libatomic/ChangeLog: * testsuite/lib/libatomic.exp: Add load_gcc_lib of scansarif.exp. libgomp/ChangeLog: * testsuite/lib/libgomp.exp: Add load_gcc_lib of scansarif.exp. libitm/ChangeLog: * testsuite/lib/libitm.exp: Add load_gcc_lib of scansarif.exp. libphobos/ChangeLog: * testsuite/lib/libphobos-dg.exp: Add load_gcc_lib of scansarif.exp. Signed-off-by: David Malcolm <dmalcolm@redhat.com>
2022-05-31openmp: Add support for firstprivate and allocate clauses on scope constructJakub Jelinek3-2/+117
OpenMP 5.2 adds support for firstprivate and allocate clauses on the scope construct and this patch adds that support to GCC. 5.2 unfortunately (IMNSHO mistakenly) marked scope construct as worksharing, which implies that it isn't possible to nest inside of it other scope, worksharing loop, sections, explicit barriers, single etc. which would make scope far less useful. I'm not implementing that part, keeping the 5.1 behavior here, and will file an issue to revert that for OpenMP 6.0. But, for firstprivate it keeps the restriction that is now implied from worksharing construct that listed var can't be private in outer context, where for reduction 5.1 had similar restriction explicit even for scope and 5.2 has it implicitly through worksharing construct. 2022-05-31 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-low.cc (build_outer_var_ref): For code == OMP_CLAUSE_ALLOCATE allow var to be private in the outer context. (lower_private_allocate): Pass OMP_CLAUSE_ALLOCATE as last argument to build_outer_var_ref. gcc/c/ * c-parser.cc (OMP_SCOPE_CLAUSE_MASK): Add firstprivate and allocate clauses. gcc/cp/ * parser.cc (OMP_SCOPE_CLAUSE_MASK): Add firstprivate and allocate clauses. gcc/testsuite/ * c-c++-common/gomp/scope-5.c: New test. * c-c++-common/gomp/scope-6.c: New test. * g++.dg/gomp/attrs-1.C (bar): Add firstprivate and allocate clauses to scope construct. * g++.dg/gomp/attrs-2.C (bar): Likewise. libgomp/ * testsuite/libgomp.c-c++-common/allocate-1.c (foo): Add testcase for scope construct with allocate clause. * testsuite/libgomp.c-c++-common/allocate-3.c (foo): Likewise. * testsuite/libgomp.c-c++-common/scope-2.c: New test.
2022-05-28OpenMP/Fortran: Add support for enter clause on declare targetTobias Burnus2-3/+11
Fortran version to C/C++ commit r13-797-g0ccba4ed8571c18c7015413441e971 gcc/fortran/ChangeLog: * dump-parse-tree.cc (show_omp_clauses): Handle OMP_LIST_ENTER. * gfortran.h: Add OMP_LIST_ENTER. * openmp.cc (enum omp_mask2, OMP_DECLARE_TARGET_CLAUSES): Add OMP_CLAUSE_ENTER. (gfc_match_omp_clauses, gfc_match_omp_declare_target, resolve_omp_clauses): Handle 'enter' clause. libgomp/ChangeLog: * libgomp.texi (OpenMP 5.2): Mark 'enter' clause as supported. * testsuite/libgomp.fortran/declare-target-1.f90: Extend to test explicit 'to' and 'enter' clause. * testsuite/libgomp.fortran/declare-target-2.f90: Update accordingly. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/declare-target-2.f90: Add 'enter' clause test. * gfortran.dg/gomp/declare-target-4.f90: Likewise.
2022-05-27openmp: Add support for enter clause on declare targetJakub Jelinek2-2/+2
OpenMP 5.1 and earlier had 2 different uses of to clause, one for target update construct with one semantics, and one for declare target directive with a different semantics. Under the hood we were using OMP_CLAUSE_TO_DECLARE to represent the latter. OpenMP 5.2 renamed the declare target clause to to enter, the old one is kept as a deprecated alias. As we are far from having full OpenMP 5.2 support, this patch adds support for the enter clause (and renames OMP_CLAUSE_TO_DECLARE to OMP_CLAUSE_ENTER with a flag to tell the spelling of the clause for better diagnostics), but doesn't deprecate the to clause on declare target just yet (that should be done as one of the last steps in 5.2 support). 2022-05-27 Jakub Jelinek <jakub@redhat.com> gcc/ * tree-core.h (enum omp_clause_code): Rename OMP_CLAUSE_TO_DECLARE to OMP_CLAUSE_ENTER. * tree.h (OMP_CLAUSE_ENTER_TO): Define. * tree.cc (omp_clause_num_ops, omp_clause_code_name): Rename OMP_CLAUSE_TO_DECLARE to OMP_CLAUSE_ENTER. * tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE, if OMP_CLAUSE_ENTER_TO, print "to" instead of "enter". * tree-nested.cc (convert_nonlocal_omp_clauses, convert_local_omp_clauses): Handle OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE. gcc/c-family/ * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_ENTER. gcc/c/ * c-parser.cc (c_parser_omp_clause_name): Parse enter clause. (c_parser_omp_all_clauses): For to clause on declare target, use OMP_CLAUSE_ENTER clause with OMP_CLAUSE_ENTER_TO instead of OMP_CLAUSE_TO_DECLARE clause. Handle PRAGMA_OMP_CLAUSE_ENTER. (OMP_DECLARE_TARGET_CLAUSE_MASK): Add enter clause. (c_parser_omp_declare_target): Use OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE. * c-typeck.cc (c_finish_omp_clauses): Handle OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE, to OMP_CLAUSE_ENTER_TO use "to" as clause name in diagnostics instead of omp_clause_code_name[OMP_CLAUSE_CODE (c)]. gcc/cp/ * parser.cc (cp_parser_omp_clause_name): Parse enter clause. (cp_parser_omp_all_clauses): For to clause on declare target, use OMP_CLAUSE_ENTER clause with OMP_CLAUSE_ENTER_TO instead of OMP_CLAUSE_TO_DECLARE clause. Handle PRAGMA_OMP_CLAUSE_ENTER. (OMP_DECLARE_TARGET_CLAUSE_MASK): Add enter clause. (cp_parser_omp_declare_target): Use OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE. * semantics.cc (finish_omp_clauses): Handle OMP_CLAUSE_ENTER instead of OMP_CLAUSE_TO_DECLARE, to OMP_CLAUSE_ENTER_TO use "to" as clause name in diagnostics instead of omp_clause_code_name[OMP_CLAUSE_CODE (c)]. gcc/testsuite/ * c-c++-common/gomp/clauses-3.c: Add tests with enter clause instead of to or modify some existing to clauses to enter. * c-c++-common/gomp/declare-target-1.c: Likewise. * c-c++-common/gomp/declare-target-2.c: Likewise. * c-c++-common/gomp/declare-target-3.c: Likewise. * g++.dg/gomp/attrs-9.C: Likewise. * g++.dg/gomp/declare-target-1.C: Likewise. libgomp/ * testsuite/libgomp.c-c++-common/target-40.c: Modify some existing to clauses to enter. * testsuite/libgomp.c/target-41.c: Likewise.
2022-05-25libgomp: Fix occassional hangs with taskwait nowait dependJakub Jelinek1-0/+48
Richi reported occassional hangs with taskwait-depend-nowait-1.* tests and I've finally manged to reproduce. The problem is if taskwait depend without nowait is encountered soon after taskwait depend nowait and the former depends on the latter and there is no other work to do, the taskwait depend without nowait is put to sleep, but the empty_task optimization in gomp_task_run_post_handle_dependers wouldn't wake it up in that case. gomp_task_run_post_handle_dependers normally does some wakeups because it schedules more work (another task), which is not the case of empty_task, but we need to do the wakeups that would be done upon task completion so that we awake sleeping threads when the last child is done. So, the taskwait-depend-nowait-1.* testcase is fixed with the else if (__builtin_expect (task->parent_depends_on, 0) part of the patch. The new testcase can hang on another problem, if the empty task is the last task of a taskgroup, we need to use atomic store like elsewhere to decrease the counter to 0, and wake up taskgroup end if needed. Yet another spot which can sleep is normal taskwait (without depend), but I believe nothing needs to be done for that - in that case we await solely until the children's queue has no tasks, tasks still waiting for dependencies aren't accounted in that, but the reason is that if taskwait should wait for something, there needs to be at least one active child doing something (in the children queue), which then possibly awakes some of its siblings when the dependencies are met, or in the empty task case awakes further dependencies, but in any case the child that finished is still handled as active child and will awake taskwait at the end if there is nothing further to do. Last sleeping case are barriers, but that is handled by ++ret and awaking the barrier. 2022-05-25 Jakub Jelinek <jakub@redhat.com> * task.c (gomp_task_run_post_handle_dependers): If empty_task is the last task taskwait depend depends on, wake it up. Similarly if it is the last child of a taskgroup, use atomic store instead of decrement and awak taskgroup wait if any. * testsuite/libgomp.c-c++-common/taskwait-depend-nowait-2.c: New test.
2022-05-24OpenMP: Support nowait with Fortran [PR105378]Tobias Burnus1-0/+42
Fortran part to C/C++/libgomp commit r13-724-gb43836914bdc2a37563cf31359b2c4803bfe4374 gcc/fortran/ PR c/105378 * openmp.cc (gfc_match_omp_taskwait): Accept nowait. gcc/testsuite/ PR c/105378 * gfortran.dg/gomp/taskwait-depend-nowait-1.f90: New. libgomp/ PR c/105378 * libgomp.texi (OpenMP 5.1): Set 'taskwait nowait' to 'Y'. * testsuite/libgomp.fortran/taskwait-depend-nowait-1.f90: New.
2022-05-24openmp: Add taskwait nowait depend support [PR105378]Jakub Jelinek1-0/+39
This patch adds support for (so far C/C++) #pragma omp taskwait nowait depend(...) directive, which is like #pragma omp task depend(...) ; but slightly optimized on the library side, so that it creates the task only for the purpose of dependency tracking and doesn't actually schedule it and wait for it when the dependencies are satisfied, instead makes its dependencies satisfied right away. 2022-05-24 Jakub Jelinek <jakub@redhat.com> PR c/105378 gcc/ * omp-builtins.def (BUILT_IN_GOMP_TASKWAIT_DEPEND_NOWAIT): New builtin. * gimplify.cc (gimplify_omp_task): Diagnose taskwait with nowait clause but no depend clauses. * omp-expand.cc (expand_taskwait_call): Use BUILT_IN_GOMP_TASKWAIT_DEPEND_NOWAIT rather than BUILT_IN_GOMP_TASKWAIT_DEPEND if nowait clause is present. gcc/c/ * c-parser.cc (OMP_TASKWAIT_CLAUSE_MASK): Add nowait clause. gcc/cp/ * parser.cc (OMP_TASKWAIT_CLAUSE_MASK): Add nowait clause. gcc/testsuite/ * c-c++-common/gomp/taskwait-depend-nowait-1.c: New test. libgomp/ * libgomp_g.h (GOMP_taskwait_depend_nowait): Declare. * libgomp.map (GOMP_taskwait_depend_nowait): Export at GOMP_5.1.1. * task.c (empty_task): New function. (gomp_task_run_post_handle_depend_hash): Declare earlier. (gomp_task_run_post_handle_depend): Declare. (GOMP_task): Optimize fn == empty_task if there is nothing to wait for. (gomp_task_run_post_handle_dependers): Optimize task->fn == empty_task. (GOMP_taskwait_depend_nowait): New function. * testsuite/libgomp.c-c++-common/taskwait-depend-nowait-1.c: New test.
2022-05-23OpenMP: Handle descriptors in target's firstprivate [PR104949]Tobias Burnus3-0/+170
For allocatable/pointer arrays, a firstprivate to a device not only needs to privatize the descriptor but also the actual data. This is implemented as: firstprivate(x) firstprivate(x.data) attach(x [bias: &x.data-&x) where the address of x in device memory is saved in hostaddrs[i] by libgomp and the middle end actually passes hostaddrs[i]' to attach. As side effect, has_device_addr(array_desc) had to be changed: before, it was converted to firstprivate in the front end; now it is handled in omp-low.cc as has_device_addr requires a shallow firstprivate (not touching the data pointer) while the normal firstprivate requires (now) a deep firstprivate. gcc/fortran/ChangeLog: PR fortran/104949 * f95-lang.cc (LANG_HOOKS_OMP_ARRAY_SIZE): Redefine. * trans-openmp.cc (gfc_omp_array_size): New. (gfc_trans_omp_variable_list): Never turn has_device_addr to firstprivate. * trans.h (gfc_omp_array_size): New. gcc/ChangeLog: PR fortran/104949 * langhooks-def.h (lhd_omp_array_size): New. (LANG_HOOKS_OMP_ARRAY_SIZE): Define. (LANG_HOOKS_DECLS): Add it. * langhooks.cc (lhd_omp_array_size): New. * langhooks.h (struct lang_hooks_for_decls): Add hook. * omp-low.cc (scan_sharing_clauses, lower_omp_target): Handle GOMP_MAP_FIRSTPRIVATE for array descriptors. libgomp/ChangeLog: PR fortran/104949 * target.c (gomp_map_vars_internal, copy_firstprivate_data): Support attach for GOMP_MAP_FIRSTPRIVATE. * testsuite/libgomp.fortran/target-firstprivate-1.f90: New test. * testsuite/libgomp.fortran/target-firstprivate-2.f90: New test. * testsuite/libgomp.fortran/target-firstprivate-3.f90: New test.
2022-05-20libgomp: Add new runtime routines omp_target_memcpy_async and ↵Marcel Vollweiler8-0/+615
omp_target_memcpy_rect_async This patch adds two new OpenMP runtime routines: omp_target_memcpy_async and omp_target_memcpy_rect_async. Both functions are introduced in OpenMP 5.1 as asynchronous variants of omp_target_memcpy and omp_target_memcpy_rect. In contrast to the synchronous variants, the asynchronous functions have two additional function parameters to allow the specification of task dependences: int depobj_count omp_depend_t *depobj_list integer(c_int), value :: depobj_count integer(omp_depend_kind), optional :: depobj_list(*) The implementation splits the synchronous functions into two parts: (a) check and (b) copy. Then (a) is used in the asynchronous functions for the sequential part, and the actual copy process (b) is executed in a new created task. The sequential part (a) takes into account the requirements for the return values: "The routine returns zero if successful. Otherwise, it returns a non-zero value." (omp_target_memcpy_async, OpenMP 5.1 spec, section 3.8.7) "An application can determine the number of inclusive dimensions supported by an implementation by passing NULL pointers (or C_NULL_PTR, for Fortran) for both dst and src. The routine returns the number of dimensions supported by the implementation for the specified device numbers. No copy operation is performed." (omp_target_memcpy_rect_async, OpenMP 5.1 spec, section 3.8.8) Due to asynchronicity an error is thrown if the asynchronous memcpy is not successful (in contrast to the synchronous functions which use a return value unequal to zero). gcc/ChangeLog: * omp-low.cc (omp_runtime_api_call): Added target_memcpy_async and target_memcpy_rect_async to omp_runtime_apis array. libgomp/ChangeLog: * libgomp.map: Added omp_target_memcpy_async and omp_target_memcpy_rect_async. * libgomp.texi: Both functions are now supported. * omp.h.in: Added omp_target_memcpy_async and omp_target_memcpy_rect_async. * omp_lib.f90.in: Added interfaces for both new functions. * omp_lib.h.in: Likewise. * target.c (ialias_redirect): Added for GOMP_task. (omp_target_memcpy): Restructured into check and copy part. (omp_target_memcpy_check): New helper function for omp_target_memcpy and omp_target_memcpy_async that checks requirements. (omp_target_memcpy_copy): New helper function for omp_target_memcpy and omp_target_memcpy_async that performs the memcpy. (omp_target_memcpy_async_helper): New helper function that is used in omp_target_memcpy_async for the asynchronous task. (omp_target_memcpy_async): Added. (omp_target_memcpy_rect): Restructured into check and copy part. (omp_target_memcpy_rect_check): New helper function for omp_target_memcpy_rect and omp_target_memcpy_rect_async that checks requirements. (omp_target_memcpy_rect_copy): New helper function for omp_target_memcpy_rect and omp_target_memcpy_rect_async that performs the memcpy. (omp_target_memcpy_rect_async_helper): New helper function that is used in omp_target_memcpy_rect_async for the asynchronous task. (omp_target_memcpy_rect_async): Added. * task.c (ialias): Added for GOMP_task. * testsuite/libgomp.c-c++-common/target-memcpy-async-1.c: New test. * testsuite/libgomp.c-c++-common/target-memcpy-async-2.c: New test. * testsuite/libgomp.c-c++-common/target-memcpy-rect-async-1.c: New test. * testsuite/libgomp.c-c++-common/target-memcpy-rect-async-2.c: New test. * testsuite/libgomp.fortran/target-memcpy-async-1.f90: New test. * testsuite/libgomp.fortran/target-memcpy-async-2.f90: New test. * testsuite/libgomp.fortran/target-memcpy-rect-async-1.f90: New test. * testsuite/libgomp.fortran/target-memcpy-rect-async-2.f90: New test.
2022-05-18OpenMP: Add Fortran support for inoutset depend-kindTobias Burnus4-3/+191
Fortran additions to the C/C++ + ME/libgomp commit r13-556-g2c16eb3157f86ae561468c540caf8eb326106b5f gcc/fortran/ChangeLog: * gfortran.h (enum gfc_omp_depend_op): Add OMP_DEPEND_INOUTSET. (gfc_omp_clauses): Enlarge ENUM_BITFIELD. * dump-parse-tree.cc (show_omp_namelist, show_omp_clauses): Handle 'inoutset' depend modifier. * openmp.cc (gfc_match_omp_clauses, gfc_match_omp_depobj): Likewise. * trans-openmp.cc (gfc_trans_omp_clauses, gfc_trans_omp_depobj): Likewise. libgomp/ChangeLog: * libgomp.texi (OpenMP 5.1): Set 'inoutset' to Y. (OpenMP Context Selectors): Add missing comma. * testsuite/libgomp.fortran/depend-5.f90: Add inoutset test. * testsuite/libgomp.fortran/depend-6.f90: Likewise. * testsuite/libgomp.fortran/depend-7.f90: Likewise. * testsuite/libgomp.fortran/depend-inoutset-1.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/all-memory-1.f90: Add inoutset test. * gfortran.dg/gomp/all-memory-2.f90: Likewise. * gfortran.dg/gomp/depobj-1.f90: Likewise. * gfortran.dg/gomp/depobj-2.f90: Likewise.
2022-05-17openmp: Add support for inoutset depend-kindJakub Jelinek4-3/+182
This patch adds support for inoutset depend-kind in depend clauses. It is very similar to the in depend-kind in that a task with a dependency with that depend-kind is dependent on all previously created sibling tasks with matching address unless they have the same depend-kind. In the in depend-kind case everything is dependent except for in -> in dependency, for inoutset everything is dependent except for inoutset -> inoutset dependency. mutexinoutset is also similar (everything is dependent except for mutexinoutset -> mutexinoutset dependency), but there is also the additional restriction that only one task with mutexinoutset for each address can be scheduled at once (i.e. mutual exclusitivty). For now we support mutexinoutset the same as inout/out, but the inoutset support is full. In order not to bump the ABI for dependencies each time (we've bumped it already once, the old ABI supports only inout/out and in depend-kind, the new ABI supports inout/out, mutexinoutset, in and depobj), this patch arranges for inoutset to be at least for the time being always handled as if it was specified through depobj even when it is not. So it uses the new ABI for that and inoutset are represented like depobj - pointer to a pair of pointers where the first one will be the actual address of the object mentioned in depend clause and second pointer will be (void *) GOMP_DEPEND_INOUTSET. 2022-05-17 Jakub Jelinek <jakub@redhat.com> gcc/ * tree-core.h (enum omp_clause_depend_kind): Add OMP_CLAUSE_DEPEND_INOUTSET. * tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_DEPEND_INOUTSET. * gimplify.cc (gimplify_omp_depend): Likewise. * omp-low.cc (lower_depend_clauses): Likewise. gcc/c-family/ * c-omp.cc (c_finish_omp_depobj): Handle OMP_CLAUSE_DEPEND_INOUTSET. gcc/c/ * c-parser.cc (c_parser_omp_clause_depend): Parse inoutset depend-kind. (c_parser_omp_depobj): Likewise. gcc/cp/ * parser.cc (cp_parser_omp_clause_depend): Parse inoutset depend-kind. (cp_parser_omp_depobj): Likewise. * cxx-pretty-print.cc (cxx_pretty_printer::statement): Handle OMP_CLAUSE_DEPEND_INOUTSET. gcc/testsuite/ * c-c++-common/gomp/all-memory-1.c (boo): Add test with inoutset depend-kind. * c-c++-common/gomp/all-memory-2.c (boo): Likewise. * c-c++-common/gomp/depobj-1.c (f1): Likewise. (f2): Adjusted expected diagnostics. * g++.dg/gomp/depobj-1.C (f4): Adjust expected diagnostics. include/ * gomp-constants.h (GOMP_DEPEND_INOUTSET): Define. libgomp/ * libgomp.h (struct gomp_task_depend_entry): Change is_in type from bool to unsigned char. * task.c (gomp_task_handle_depend): Handle GOMP_DEPEND_INOUTSET. Ignore dependencies where task->depend[i].is_in && task->depend[i].is_in == ent->is_in rather than just task->depend[i].is_in && ent->is_in. Remember whether GOMP_DEPEND_IN loop is needed and guard the loop with that conditional. (gomp_task_maybe_wait_for_dependencies): Handle GOMP_DEPEND_INOUTSET. Ignore dependencies where elem.is_in && elem.is_in == ent->is_in rather than just elem.is_in && ent->is_in. * testsuite/libgomp.c-c++-common/depend-1.c (test): Add task with inoutset depend-kind. * testsuite/libgomp.c-c++-common/depend-2.c (test): Likewise. * testsuite/libgomp.c-c++-common/depend-3.c (test): Likewise. * testsuite/libgomp.c-c++-common/depend-inoutset-1.c: New test.
2022-05-17OpenMP: Add omp_all_memory support to FortranTobias Burnus3-0/+360
Fortran part to the C/C++/backend implementation r13-337-g7f78783dbedca0183d193e475262ca3c489fd365 gcc/fortran/ChangeLog: * dump-parse-tree.cc (show_omp_namelist): Handle omp_all_memory. * openmp.cc (gfc_match_omp_variable_list, gfc_match_omp_depend_sink, gfc_match_omp_clauses, resolve_omp_clauses): Likewise. * trans-openmp.cc (gfc_trans_omp_clauses, gfc_trans_omp_depobj): Likewise. * resolve.cc (resolve_symbol): Reject it as symbol. libgomp/ChangeLog: * libgomp.texi (OpenMP 5.1): Set omp_all_memory to 'Y'. * testsuite/libgomp.fortran/depend-5.f90: New test. * testsuite/libgomp.fortran/depend-6.f90: New test. * testsuite/libgomp.fortran/depend-7.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/all-memory-1.f90: New test. * gfortran.dg/gomp/all-memory-2.f90: New test. * gfortran.dg/gomp/all-memory-3.f90: New test.
2022-05-16OpenMP, C++: Add template support for the has_device_addr clause.Marcel Vollweiler3-0/+113
This patch adds support for list items in the has_device_addr clause which type is given by C++ template parameters. gcc/cp/ChangeLog: * pt.cc (tsubst_omp_clauses): Added OMP_CLAUSE_HAS_DEVICE_ADDR. * semantics.cc (finish_omp_clauses): Added template decl processing. libgomp/ChangeLog: * testsuite/libgomp.c++/target-has-device-addr-7.C: New test. * testsuite/libgomp.c++/target-has-device-addr-8.C: New test. * testsuite/libgomp.c++/target-has-device-addr-9.C: New test.
2022-05-13libgomp.fortran/target-nowait-array-section.f90: Fix typoTobias Burnus1-1/+1
Fix typo as requested in the review approval. libgomp/ChangeLog: * testsuite/libgomp.fortran/target-nowait-array-section.f90: New test.
2022-05-13OpenMP/Fortran: Use firstprivat not alloc for ptr attach for arraysTobias Burnus1-0/+56
For a non-descriptor array, map(A(n:m)) was mapped as map(tofrom:A[n-1] [len: ...]) map(alloc:A [pointer assign, bias: ...]) with this patch, it is changed to map(tofrom:A[n-1] [len: ...]) map(firstprivate:A [pointer assign, bias: ...]) The latter avoids an alloc - and also avoids the race condition with nowait in the enclosed testcase. (Note: predantically, the testcase is invalid since OpenMP 5.1, violating the map clause restriction at [354:10-13]. gcc/fortran/ChangeLog: * trans-openmp.cc (gfc_trans_omp_clauses): When mapping nondescriptor array sections, use GOMP_MAP_FIRSTPRIVATE_POINTER instead of GOMP_MAP_POINTER for the pointer attachment. libgomp/ChangeLog: * testsuite/libgomp.fortran/target-nowait-array-section.f90: New test.
2022-05-12Refactor '-ldl' handling for libgomp proper and pluginsThomas Schwinge1-1/+1
Instead of implicit global 'LIBS="-ldl $LIBS"' via 'AC_CHECK_LIB', make '-ldl' explicit for libgomp proper, and clean up 'PLUGIN_GCN_LIBS', 'PLUGIN_NVPTX_LIBS' accordingly. libgomp/ * Makefile.am (libgomp_la_LIBADD): Initialize. * plugin/configfrag.ac (DL_LIBS): New. (PLUGIN_GCN_LIBS): Remove. (PLUGIN_NVPTX_LIBS): Don't set in the 'PLUGIN_NVPTX_DYNAMIC' case. * plugin/Makefrag.am (libgomp_la_LIBADD) (libgomp_plugin_gcn_la_LIBADD): Consider '$(DL_LIBS)'. (libgomp_plugin_nvptx_la_LIBADD) <PLUGIN_NVPTX_DYNAMIC>: Likewise. * Makefile.in: Regenerate. * config.h.in: Likewise. * configure: Likewise. * testsuite/Makefile.in: Likewise.
2022-05-12libgomp plugins: Don't 'AC_SUBST' and 'AC_DEFINE_UNQUOTED' for 'PLUGIN_GCN', ↵Thomas Schwinge1-2/+0
'PLUGIN_NVPTX' Nothing ever used these. libgomp/ * plugin/configfrag.ac: Don't 'AC_SUBST' and 'AC_DEFINE_UNQUOTED' for 'PLUGIN_GCN', 'PLUGIN_NVPTX'. * Makefile.in: Regenerate. * config.h.in: Likewise. * configure: Likewise. * testsuite/Makefile.in: Likewise.
2022-05-12openmp: Add omp_all_memory support (C/C++ only so far)Jakub Jelinek3-0/+329
The ugly part is that OpenMP 5.1 made omp_all_memory a reserved identifier which isn't allowed to be used anywhere but in the depend clause, this is against how everything else has been handled in OpenMP so far (where some identifiers could have special meaning in some OpenMP clauses or pragmas but not elsewhere). The patch handles it by making it a conditional keyword (for -fopenmp only) and emitting a better diagnostics when it is used in a primary expression. Having a nicer diagnostics when e.g. trying to do int omp_all_memory; or int *omp_all_memory[10]; etc. would mean changing too many spots and hooking into name lookups to reject declaring any such symbols would be too ugly and I'm afraid there are way too many spots where one can introduce a name (variables, functions, namespaces, struct, enum, enumerators, template arguments, ...). Otherwise, the handling is quite simple, normal depend clauses lower into addresses of variables being handed over to the library, for omp_all_memory I'm using NULL pointers. omp_all_memory can only be used with inout or out depend kinds and means that a task is dependent on all previously created sibling tasks that have any dependency (of any depend kind) and that any later created sibling tasks will be dependent on it if they have any dependency. 2022-05-12 Jakub Jelinek <jakub@redhat.com> gcc/ * gimplify.cc (gimplify_omp_depend): Don't build_fold_addr_expr if null_pointer_node. (gimplify_scan_omp_clauses): Likewise. * tree-pretty-print.cc (dump_omp_clause): Print null_pointer_node as omp_all_memory. gcc/c-family/ * c-common.h (enum rid): Add RID_OMP_ALL_MEMORY. * c-omp.cc (c_finish_omp_depobj): Don't build_fold_addr_expr if null_pointer_node. gcc/c/ * c-parser.cc (c_parse_init): Register omp_all_memory as keyword if flag_openmp. (c_parser_postfix_expression): Diagnose uses of omp_all_memory in postfix expressions. (c_parser_omp_variable_list): Handle omp_all_memory in depend clause. * c-typeck.cc (c_finish_omp_clauses): Handle omp_all_memory keyword in depend clause as null_pointer_node, diagnose invalid uses. gcc/cp/ * lex.cc (init_reswords): Register omp_all_memory as keyword if flag_openmp. * parser.cc (cp_parser_primary_expression): Diagnose uses of omp_all_memory in postfix expressions. (cp_parser_omp_var_list_no_open): Handle omp_all_memory in depend clause. * semantics.cc (finish_omp_clauses): Handle omp_all_memory keyword in depend clause as null_pointer_node, diagnose invalid uses. * pt.cc (tsubst_omp_clause_decl): Pass through omp_all_memory. gcc/testsuite/ * c-c++-common/gomp/all-memory-1.c: New test. * c-c++-common/gomp/all-memory-2.c: New test. * c-c++-common/gomp/all-memory-3.c: New test. * g++.dg/gomp/all-memory-1.C: New test. * g++.dg/gomp/all-memory-2.C: New test. libgomp/ * libgomp.h (struct gomp_task): Add depend_all_memory member. * task.c (gomp_init_task): Initialize depend_all_memory. (gomp_task_handle_depend): Handle omp_all_memory. (gomp_task_run_post_handle_depend_hash): Clear parent->depend_all_memory if equal to current task. (gomp_task_maybe_wait_for_dependencies): Handle omp_all_memory. * testsuite/libgomp.c-c++-common/depend-1.c: New test. * testsuite/libgomp.c-c++-common/depend-2.c: New test. * testsuite/libgomp.c-c++-common/depend-3.c: New test.
2022-05-11libgomp: Remove unused '--with-hsa-runtime', '--with-hsa-runtime-include', ↵Thomas Schwinge1-2/+0
'--with-hsa-runtime-lib' With recent commit 2e309a4eff80e55b53d32d26926a2a94eabfea21 "libgomp testsuite: Don't amend 'LD_LIBRARY_PATH' for system-provided HSA Runtime library", and commit d6adba307508c75f1ccb2121eb1a43c9ab1d4056 "libgomp GCN plugin: Clean up unused references to system-provided HSA Runtime library", the last uses of '--with-hsa-runtime' etc. are gone. gcc/ * doc/install.texi: Don't document '--with-hsa-runtime', '--with-hsa-runtime-include', '--with-hsa-runtime-lib'. libgomp/ * plugin/configfrag.ac: Remove '--with-hsa-runtime', '--with-hsa-runtime-include', '--with-hsa-runtime-lib' processing. * Makefile.in: Regenerate. * configure: Likewise. * testsuite/Makefile.in: Likewise.
2022-05-11libgomp GCN plugin: Clean up always-empty 'PLUGIN_GCN_CPPFLAGS', ↵Thomas Schwinge1-2/+0
'PLUGIN_GCN_LDFLAGS' After recent commit d6adba307508c75f1ccb2121eb1a43c9ab1d4056 "libgomp GCN plugin: Clean up unused references to system-provided HSA Runtime library", these aren't set anymore. libgomp/ * plugin/Makefrag.am (libgomp_plugin_gcn_la_CPPFLAGS): Don't consider 'PLUGIN_GCN_CPPFLAGS'. (libgomp_plugin_gcn_la_LDFLAGS): Don't consider 'PLUGIN_GCN_LDFLAGS'. * plugin/configfrag.ac (PLUGIN_GCN_CPPFLAGS, PLUGIN_GCN_LDFLAGS): Remove. * Makefile.in: Regenerate. * configure: Likewise. * testsuite/Makefile.in: Likewise.
2022-05-11libgomp testsuite: Don't amend 'LD_LIBRARY_PATH' for system-provided HSA ↵Thomas Schwinge2-5/+0
Runtime library This is only active if GCC is 'configure'd with '--with-hsa-runtime=[...]' or '--with-hsa-runtime-lib=[...]' -- which nobody really is doing, as far as I can tell. 'libgomp/testsuite/lib/libgomp.exp:libgomp_init' states: # For build-tree testing, also consider the library paths used for builing. # For installed testing, we assume all that to be provided in the sysroot. if { $blddir != "" } { [...] global hsa_runtime_lib if { $hsa_runtime_lib != "" } { append always_ld_library_path ":$hsa_runtime_lib" } } However, the libgomp GCN plugin is unconditionally built against the GCC-shipped 'include/hsa*.h' header files, and at run time does 'dlopen("libhsa-runtime64.so.1")', so there is no system-provided HSA Runtime library "used for builing". It thus doesn't make sense to amend 'LD_LIBRARY_PATH' for system-provided HSA Runtime library. libgomp/ * testsuite/lib/libgomp.exp (libgomp_init): Don't 'append always_ld_library_path ":$hsa_runtime_lib"'. * testsuite/libgomp-test-support.exp.in (hsa_runtime_lib): Don't set.
2022-05-10Fix up 'libgomp.fortran/use_device_addr-5.f90' multi-device testingThomas Schwinge1-1/+1
Fix-up for recent commit r13-116-g3f8c389fe90bf565a6221a46bb7fb745dd4c1510 "OpenMP: Fix use_device_{addr,ptr} with in-data-sharing arg", where we currently get: libgomp: use_device_ptr pointer wasn't mapped FAIL: libgomp.fortran/use_device_addr-5.f90 -O execution test libgomp/ * testsuite/libgomp.fortran/use_device_addr-5.f90: Fix up multi-device testing.
2022-05-06OpenMP, libgomp: Add new runtime routine omp_target_is_accessible.Marcel Vollweiler2-0/+97
gcc/ChangeLog: * omp-low.cc (omp_runtime_api_call): Added target_is_accessible to omp_runtime_apis array. libgomp/ChangeLog: * libgomp.map: Added omp_target_is_accessible. * libgomp.texi: Tagged omp_target_is_accessible as supported. * omp.h.in: Added omp_target_is_accessible. * omp_lib.f90.in: Added interface for omp_target_is_accessible. * omp_lib.h.in: Likewise. * target.c (omp_target_is_accessible): Added implementation of omp_target_is_accessible. * testsuite/libgomp.c-c++-common/target-is-accessible-1.c: New test. * testsuite/libgomp.fortran/target-is-accessible-1.f90: New test.
2022-05-04OpenMP: Fix use_device_{addr,ptr} with in-data-sharing argTobias Burnus1-0/+143
For array-descriptor vars, the descriptor is assigned to a temporary. However, this failed when the clause's argument was in turn in a data-sharing clause as the outer context's VALUE_EXPR wasn't used. gcc/ChangeLog: * omp-low.cc (lower_omp_target): Fix use_device_{addr,ptr} with list item that is in an outer data-sharing clause. libgomp/ChangeLog: * testsuite/libgomp.fortran/use_device_addr-5.f90: New test.
2022-05-02OpenMP, libgomp: Add new runtime routine omp_get_mapped_ptr.Marcel Vollweiler8-0/+597
This patch adds the OpenMP runtime routine "omp_get_mapped_ptr" which was introduced in OpenMP 5.1. gcc/ChangeLog: * omp-low.cc (omp_runtime_api_call): Added get_mapped_ptr to omp_runtime_apis array. libgomp/ChangeLog: * libgomp.map: Added omp_get_mapped_ptr. * libgomp.texi: Tagged omp_get_mapped_ptr as supported. * omp.h.in: Added omp_get_mapped_ptr. * omp_lib.f90.in: Added interface for omp_get_mapped_ptr. * omp_lib.h.in: Likewise. * target.c (omp_get_mapped_ptr): Added implementation of omp_get_mapped_ptr. * testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c: New test. * testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c: New test. * testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c: New test. * testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c: New test. * testsuite/libgomp.fortran/get-mapped-ptr-1.f90: New test. * testsuite/libgomp.fortran/get-mapped-ptr-2.f90: New test. * testsuite/libgomp.fortran/get-mapped-ptr-3.f90: New test. * testsuite/libgomp.fortran/get-mapped-ptr-4.f90: New test.
2022-04-28Fix up 'libgomp.oacc-fortran/print-1.f90' GCN offloading compilation [PR104717]Thomas Schwinge1-1/+29
That got broken by recent commit b2202431910e30d8505c94d1cb9341cac7080d10 "fortran: Fix up gfc_trans_oacc_construct [PR104717]". PR fortran/104717 libgomp/ * testsuite/libgomp.oacc-fortran/print-1.f90: Add OpenACC privatization scanning. For GCN offloading compilation, raise '-mgang-private-size'.
2022-04-25fortran: Fix up gfc_trans_oacc_construct [PR104717]Jakub Jelinek1-0/+7
So that move_sese_region_to_fn works properly, OpenMP/OpenACC constructs for which that function is invoked need an extra artificial BIND_EXPR around their body so that we move all variables of the bodies. The C/C++ FEs do that both for OpenMP constructs like OMP_PARALLEL, OMP_TASK or OMP_TARGET and for OpenACC constructs that behave similarly to OMP_TARGET, but the Fortran FE only does that for OpenMP constructs. The following patch does that for OpenACC constructs too. PR fortran/104717 gcc/fortran/ * trans-openmp.cc (gfc_trans_oacc_construct): Wrap construct body in an extra BIND_EXPR. gcc/testsuite/ * gfortran.dg/goacc/pr104717.f90: New test. * gfortran.dg/goacc/privatization-1-compute-loop.f90: Adjust. libgomp/ * testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Adjust. Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
2022-04-05OpenMP: Fix nested use_device_ptrChung-Lin Tang1-0/+41
This patch fixes a bug in lower_omp_target, where for Fortran arrays, the expanded sender assignment is wrongly using the variable in the current ctx, instead of the one looked-up outside, which is causing use_device_ptr/addr to fail to work when used inside an omp-parallel (where the omp child_fn is split away from the original). The fix is inside omp-low.cc, though because the omp_array_data langhook is used only by Fortran, this is essentially Fortran-specific. 2022-04-05 Chung-Lin Tang <cltang@codesourcery.com> gcc/ChangeLog: * omp-low.cc (lower_omp_target): Use outer context looked-up 'var' as argument to lang_hooks.decls.omp_array_data, instead of 'ovar' from current clause. libgomp/ChangeLog: * testsuite/libgomp.fortran/use_device_ptr-4.f90: New testcase.
2022-04-04[libgomp/testsuite] Fix libgomp.fortran/examples-4/declare_target-{1,2}.f90Tom de Vries2-24/+38
The test-cases libgomp.fortran/examples-4/declare_target-{1,2}.f90 mean to set an nvptx-specific limit using offload_target_nvptx, but also change behaviour for amd. That is, there is now a difference in behaviour between: - a compiler configured for GCN offloading, and - a compiler configured for both GCN and nvptx offloading. Fix this by using instead on_device_arch_nvptx. Tested on x86_64 with nvptx accelerator. libgomp/ChangeLog: 2022-04-04 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.fortran/examples-4/declare_target-1.f90: Use on_device_arch_nvptx instead of offload_target_nvptx. * testsuite/libgomp.fortran/examples-4/declare_target-2.f90: Same.
2022-04-01[libgomp, testsuite, nvptx] Limit recursion in declare_target-{1,2}.f90Tom de Vries2-11/+27
When running testcases libgomp.fortran/examples-4/declare_target-{1,2}.f90 on an RTX A2000 (sm_86) with driver 510.60.02 and with GOMP_NVPTX_JIT=-O0 I run into: ... FAIL: libgomp.fortran/examples-4/declare_target-1.f90 -O0 \ -DGOMP_NVPTX_JIT=-O0 execution test FAIL: libgomp.fortran/examples-4/declare_target-2.f90 -O0 \ -DGOMP_NVPTX_JIT=-O0 execution test ... Fix this by further limiting recursion depth in the test-cases for nvptx. Furthermore, make the recursion depth limiting nvptx-specific. Tested on x86_64 with nvptx accelerator. libgomp/ChangeLog: 2022-04-01 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.fortran/examples-4/declare_target-1.f90: Define and use REC_DEPTH. * testsuite/libgomp.fortran/examples-4/declare_target-2.f90: Same.
2022-04-01[libgomp, testsuite, nvptx] Fix dg-output test in vector-length-128-7.cTom de Vries1-1/+1
When running test-case libgomp.oacc-c-c++-common/vector-length-128-7.c on an RTX A2000 (sm_86) with driver 510.60.02 I run into: ... FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c \ -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 \ output pattern test ... The failing check verifies the launch dimensions: ... /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: \ launch gangs=1, workers=8, vectors=128" } */ ... which fails because (as we can see with GOMP_DEBUG=1) the actual num_workers is 6: ... nvptx_exec: kernel main$_omp_fn$0: launch gangs=1, workers=6, vectors=128 ... This is due to the result of cuOccupancyMaxPotentialBlockSize (which suggests 'a launch configuration with reasonable occupancy') printed just before: ... cuOccupancyMaxPotentialBlockSize: grid = 52, block = 768 ... [ Note: 6 * 128 == 768. ] Fix this by updating the check to allow num_workers in the range 1 to 8. Tested on x86_64 with nvptx accelerator. libgomp/ChangeLog: 2022-04-01 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: Fix num_workers check.
2022-03-25[libgomp, testsuite] Scale down some OpenACC test-casesTom de Vries3-23/+46
When a display manager is running on an nvidia card, all CUDA kernel launches get a 5 seconds watchdog timer. Consequently, when running the libgomp testsuite with nvptx accelerator and GOMP_NVPTX_JIT=-O0 we run into a few FAILs like this: ... libgomp: cuStreamSynchronize error: the launch timed out and was terminated FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims.c \ -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none -O0 \ execution test ... Fix this by scaling down the failing test-cases by default, and reverting to the original behaviour for GCC_TEST_RUN_EXPENSIVE=1. Tested on x86_64-linux with nvptx accelerator. libgomp/ChangeLog: 2022-03-25 Tom de Vries <tdevries@suse.de> PR libgomp/105042 * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Reduce execution time. * testsuite/libgomp.oacc-c-c++-common/vred2d-128.c: Same. * testsuite/libgomp.oacc-fortran/parallel-dims.f90: Same.
2022-03-23LTO: Fixes for renaming issues with offload/OpenMP [PR104285]Tobias Burnus6-0/+290
gcc/lto/ChangeLog: PR middle-end/104285 * lto-partition.cc (maybe_rewrite_identifier): Use get_identifier for the returned string to be usable as hash key. (validize_symbol_for_target): Hence, use return value directly. (privatize_symbol_name_1): Track maybe_rewrite_identifier renames. * lto.cc (offload_handle_link_vars): Move function up before ... (do_whole_program_analysis): Call it after static renamings. (lto_main): Move call after static renamings. libgomp/ChangeLog: PR middle-end/104285 * testsuite/libgomp.c++/target-same-name-2-a.C: New test. * testsuite/libgomp.c++/target-same-name-2-b.C: New test. * testsuite/libgomp.c++/target-same-name-2.C: New test. * testsuite/libgomp.c-c++-common/target-same-name-1-a.c: New test. * testsuite/libgomp.c-c++-common/target-same-name-1-b.c: New test. * testsuite/libgomp.c-c++-common/target-same-name-1.c: New test.
2022-03-22[nvptx] Add warp sync at simt exitTom de Vries1-0/+25
Consider this code (with N defined to 1024): ... float v = 0.0; #pragma omp target map(tofrom: v) #pragma omp parallel for simd for (int i = 0 ; i < N; i++) { #pragma omp atomic update v = v + 1.0; } ... It hangs when executing on target board unix/-foffload=-misa=sm_75, using drivers 470.103.01 and 510.54 on a T400 board (sm_75). I'm tentatively identifying the problem as a bug in -muniform-simt for architectures that support Independent Thread Scheduling (sm_70 and later). The problem -muniform-simt is trying to address is to make sure that a register produced outside an openmp simd region is available when used in any lane inside an simd region. The solution is to, outside an simd region, execute in all warp lanes, thus producing consistent values in result registers in each warp thread. This approach doesn't work when executing in all warp lanes multiplies the side effects from 1 to 32 separate side effects, which is the case for atomic insns. So atomic insns are rewritten to execute only in lane 0, and if there are any results, those are propagated to the other threads in the warp. [ And likewise for system calls malloc, free, vprintf. ] Now, consider a non-atomic update: ld, add, store. The store has side effects, are those multiplied or not? Pre-sm_70 we can assume that at the end of an SIMT region, any divergent control flow has reconverged, and we have a uniform warp, executing in lock step. So: - the load will load the same value into the result register across the warp, - the add will write the same value into the result register across the warp, - the store will write the same value to the same memory location, 32 times, at once, having the result of a single store. So, no side-effect multiplication (well, at least that's the observation). Starting sm_70, the threads in a warp are no longer guaranteed to reconverge after divergence. There's a "Convergence Optimizer" that can can identify that it is safe for a warp to reconverge, but that works only as long as the code does not contain "synchronizing operations". Consequently, the ld, add, store sequence can be executed by a non-uniform warp, which means the side effects can have multiplied, and the registers are no longer guarantueed to be in sync. The atomic update in the example above is translated using an atom.cas loop, which means that we have divergence (because only one thread is allowed to succeed at a time) and the "Convergence Optimizer" doesn't reconverge probably because the atom.cas counts as a "synchronizing operation". So, it seems plausible that the root cause for the mentioned hang is the problem described above. Fix this by adding an explicit warp sync at simt exit. Note that we're assuming here that the warp will stay uniform until the next SIMT region entry. Tested on x86_64 with nvptx accelerator. gcc/ChangeLog: 2022-03-09 Tom de Vries <tdevries@suse.de> PR target/104916 PR target/104783 * config/nvptx/nvptx.md (define_expand "omp_simt_exit"): Emit warp sync (or uniform warp check for mptx < 6.0). libgomp/ChangeLog: 2022-03-15 Tom de Vries <tdevries@suse.de> PR target/104916 PR target/104783 * testsuite/libgomp.c/pr104783-2.c: New test.
2022-03-18Fortran/OpenMP: Fix privatization of associated namesTobias Burnus1-0/+92
gfc_omp_predetermined_sharing cases the associate-name pointer variable to be OMP_CLAUSE_DEFAULT_FIRSTPRIVATE, which is fine. However, the associated selector is shared. Thus, the target of associate-name pointer should not get copied. (It was before but because of gfc_omp_privatize_by_reference returning false, the selector was not only wrongly copied but this was also not done properly.) gcc/fortran/ChangeLog: PR fortran/103039 * trans-openmp.cc (gfc_omp_clause_copy_ctor, gfc_omp_clause_dtor): Only privatize pointer for associate names. libgomp/ChangeLog: PR fortran/103039 * testsuite/libgomp.fortran/associate4.f90: New test.
2022-03-18[openmp] Fix SIMT reduction using TRUTH_{AND,OR}IF_EXPRTom de Vries2-0/+46
Consider test-case pr104952-1.c, included in this commit, containing: ... #pragma omp target map(tofrom:result) map(to:arr) #pragma omp simd reduction(||: result) ... When run on x86_64 with nvptx accelerator, the test-case either aborts or hangs. The reduction clause is translated by the SIMT code (active for nvptx) as a butterfly reduction loop with this butterfly shuffle / update pair: ... D.2163 = D.2163 || .GOMP_SIMT_XCHG_BFLY (D.2163, D.2164) ... in the loop body. The problem is that the butterfly shuffle is possibly not executed, while it needs to be executed unconditionally. Fix this by translating instead as: ... D.tmp_bfly = .GOMP_SIMT_XCHG_BFLY (D.2163, D.2164) D.2163 = D.2163 || D.tmp_bfly ... Tested on x86_64-linux with nvptx accelerator. gcc/ChangeLog: 2022-03-17 Tom de Vries <tdevries@suse.de> PR target/104952 * omp-low.cc (lower_rec_input_clauses): Make sure GOMP_SIMT_XCHG_BFLY is executed unconditionally. libgomp/ChangeLog: 2022-03-17 Tom de Vries <tdevries@suse.de> PR target/104952 * testsuite/libgomp.c/pr104952-1.c: New test. * testsuite/libgomp.c/pr104952-2.c: New test.
2022-03-17Enhance further testcases to verify Openacc 'kernels' decompositionThomas Schwinge3-6/+67
gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-1.c: Enhance. * c-c++-common/goacc/kernels-loop-g.c: Likewise. * c-c++-common/goacc/nesting-1.c: Likewise. * gcc.dg/goacc/nested-function-1.c: Likewise. * gfortran.dg/goacc/common-block-3.f90: Likewise. * gfortran.dg/goacc/nested-function-1.f90: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Enhance. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-g.c: Likewise. * testsuite/libgomp.oacc-fortran/if-1.f90: Likewise.
2022-03-17Enhance further testcases to verify handling of OpenACC privatization level ↵Thomas Schwinge2-65/+52
[PR90115] As originally introduced in commit 11b8286a83289f5b54e813f14ff56d730c3f3185 "[OpenACC privatization] Largely extend diagnostics and corresponding testsuite coverage [PR90115]". PR middle-end/90115 gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-1.c: Enhance. * gfortran.dg/goacc/common-block-3.f90: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Enhance. * testsuite/libgomp.oacc-fortran/if-1.f90: Likewise.
2022-03-16OpenMP, Fortran: Bugfix for omp_set_num_teams.Marcel Vollweiler1-0/+10
This patch fixes a small bug in the omp_set_num_teams implementation. libgomp/ChangeLog: * fortran.c (omp_set_num_teams_8_): Call omp_set_num_teams instead of omp_set_max_active_levels. * testsuite/libgomp.fortran/icv-8.f90: New test.
generated by cgit v1.1 at 2025-09-24 06:41:28 +0000