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.
Similar to how the other 'mkoffload's got changed in
recent commit 683f118439
"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.
These should compile and link and execute in all configurations; host-fallback
execution, which we may actually verify.
Follow-up to recent commit 683f118439
"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.
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 683f118439
"OpenMP: Move omp requires checks to libgomp".
libgomp/
* testsuite/libgomp.c-c++-common/requires-3.c: Adjust.
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>
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.
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.
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.
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.
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.
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>
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.
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.
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.
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.
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.
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.
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.
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.
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.
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.
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.
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.
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.
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>
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.
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.
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.
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.
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.
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.
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.
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.
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.
