* 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.
* 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.
Tobias Burnus [Fri, 13 May 2022 18:00:34 +0000 (20:00 +0200)]
OpenMP/Fortran: Use firstprivat not alloc for ptr attach for arrays
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.
Andrew Stubbs [Mon, 20 Jun 2022 14:51:15 +0000 (15:51 +0100)]
amdgcn: libgomp plugin USM implementation
Implement the Unified Shared Memory API calls in the GCN plugin.
The allocate and free are pretty straight-forward because all "target" memory
allocations are compatible with USM, on the right hardware. However, there's
no known way to check what memory region was used, after the fact, so we use a
splay tree to record allocations so we can answer "is_usm_ptr" later.
Andrew Stubbs [Fri, 17 Jun 2022 12:07:11 +0000 (13:07 +0100)]
amdgcn, openmp: Auto-detect USM mode and set HSA_XNACK
The AMD GCN runtime must be set to the correct mode for Unified Shared Memory
to work, but this is not always clear at compile and link time due to the split
nature of the offload compilation pipeline.
This patch sets a new attribute on OpenMP offload functions to ensure that the
information is passed all the way to the backend. The backend then places a
marker in the assembler code for mkoffload to find. Finally mkoffload places a
constructor function into the final program to ensure that the HSA_XNACK
environment variable passes the correct mode to the GPU.
The HSA_XNACK variable must be set before the HSA runtime is even loaded, so
it makes more sense to have this set within the constructor than at some point
later within libgomp or the GCN plugin.
Andrew Stubbs [Fri, 10 Jun 2022 14:15:49 +0000 (15:15 +0100)]
amdgcn: Support XNACK mode
The XNACK feature allows memory load instructions to restart safely following
a page-miss interrupt. This is useful for shared-memory devices, like APUs,
and to implement OpenMP Unified Shared Memory.
To support the feature we must be able to set the appropriate meta-data and
set the load instructions to early-clobber. When the port supports scheduling
of s_waitcnt instructions there will be further requirements.
gcc/ChangeLog:
* config/gcn/gcn-hsa.h (XNACKOPT): New macro.
(ASM_SPEC): Use XNACKOPT.
* config/gcn/gcn-opts.h (enum sram_ecc_type): Rename to ...
(enum hsaco_attr_type): ... this, and generalize the names.
(TARGET_XNACK): New macro.
* config/gcn/gcn-valu.md (gather<mode>_insn_1offset<exec>):
Add xnack compatible alternatives.
(gather<mode>_insn_2offsets<exec>): Likewise.
* config/gcn/gcn.c (gcn_option_override): Permit -mxnack for devices
other than Fiji.
(gcn_expand_epilogue): Remove early-clobber problems.
(output_file_start): Emit xnack attributes.
(gcn_hsa_declare_function_name): Obey -mxnack setting.
* config/gcn/gcn.md (xnack): New attribute.
(enabled): Rework to include "xnack" attribute.
(*movbi): Add xnack compatible alternatives.
(*mov<mode>_insn): Likewise.
(*mov<mode>_insn): Likewise.
(*mov<mode>_insn): Likewise.
(*movti_insn): Likewise.
* config/gcn/gcn.opt (-mxnack): Add the "on/off/any" syntax.
(sram_ecc_type): Rename to ...
(hsaco_attr_type: ... this.)
* config/gcn/mkoffload.c (SET_XNACK_ANY): New macro.
(TEST_XNACK): Delete.
(TEST_XNACK_ANY): New macro.
(TEST_XNACK_ON): New macro.
(main): Support the new -mxnack=on/off/any syntax.
Jason Merrill [Wed, 11 May 2022 18:53:26 +0000 (14:53 -0400)]
c++: lambda template in requires [PR105541]
Since the patch for PR103408, the template parameters for the lambda in this
test have level 1 instead of 2, and we were treating null template args as 1
level of arguments, so tsubst_template_parms decided it had nothing to do.
Fixed by distinguishing between <> and no args at all, which is what we have
in our "substitution" in a requires-expression.
PR c++/105541
gcc/cp/ChangeLog:
* cp-tree.h (TMPL_ARGS_DEPTH): 0 for null args.
* parser.cc (cp_parser_enclosed_template_argument_list):
Use 0-length TREE_VEC for <>.
We were wrongly looking up the generic lambda op() in a dependent scope, and
then trying to look up its instantiation at substitution time, but lambdas
aren't instantiated, so we crashed. The fix is to not look into dependent
class scopes.
But this created trouble with wrongly trying to use a template from the
enclosing scope when we aren't actually looking at a template-argument-list,
in template/lookup18.C, so let's avoid that.
Harald Anlauf [Tue, 21 Jun 2022 21:20:18 +0000 (23:20 +0200)]
Fortran: fix simplification of INDEX(str1,str2) [PR105691]
gcc/fortran/ChangeLog:
PR fortran/105691
* simplify.cc (gfc_simplify_index): Replace old simplification
code by the equivalent of the runtime library implementation. Use
HOST_WIDE_INT instead of int for string index, length variables.
gcc/testsuite/ChangeLog:
PR fortran/105691
* gfortran.dg/index_6.f90: New test.
Harald Anlauf [Fri, 24 Jun 2022 20:21:39 +0000 (22:21 +0200)]
Fortran: fix checking of arguments to UNPACK when MASK is a variable [PR105813]
gcc/fortran/ChangeLog:
PR fortran/105813
* check.cc (gfc_check_unpack): Try to simplify MASK argument to
UNPACK so that checking of the VECTOR argument can work when MASK
is a variable.
gcc/testsuite/ChangeLog:
PR fortran/105813
* gfortran.dg/unpack_vector_1.f90: New test.
gcc/
* omp-low.cc (usm_transform): Remove unused function argument.
(This should be a fixup to 54c2d861ac62e30ebf34a4e62ee0d55478a742b9: 'Build
fix for 'openmp: Use libgomp memory allocation functions with unified shared
memory')
Tobias Burnus [Thu, 30 Jun 2022 06:30:48 +0000 (08:30 +0200)]
Build fix for 'openmp: allow requires unified_shared_memory'
OG12 commit fa65fc45972d27f2fd79a44eaba1978348177ee9 added an
error diagnostic (moved around in later commits); this diagnostic
caused bootstrap fails as %<...%> were missing. This commit adds
them.
gcc/c/
* c-parser.cc (c_parser_omp_requires): Add missing %<...%> in error.
gcc/cp/
* parser.cc (cp_parser_omp_requires): Add missing %<...%> in error.
Lulu Cheng [Mon, 27 Jun 2022 08:26:25 +0000 (16:26 +0800)]
LoongArch: Remove undefined behavior from code [PR 106097]
C++2017 and previous standard description:
The value of E1 << E2 is E1 left-shifted E2 bit positions;
vacated bits are zero-filled. If E1 has an unsigned type,
the value of the result is E1×2E2, reduced modulo one more
than the maximum value representable inthe result type.
Otherwise, if E1 has a signed type and non-negative value,
and E1×2E2 is representablein the corresponding unsigned
type of the result type, then that value, converted to the
result type, is the resulting value; otherwise, the behavior
is undefined.
The value of E1 >> E2 is E1 right-shifted E2 bit positions.
If E1 has an unsigned type or if E1 has a signed type and
a non-negative value, the value of the result is the integral
part of the quotient of E1/2E2. If E1 has a signed type and
a negative value, the resulting value is implementation-defined.
gcc/ChangeLog:
PR target/106097
* config/loongarch/loongarch.cc (loongarch_build_integer):
Remove undefined behavior from code.
Revert changes to acc_prof-init-1.c and acc_prof-parallel-1.c
There is an interaction between the commits "OpenACC profiling-interface
fixes for asynchronous operations" and "[OpenACC] Clarify sequencing of
'async' data copying vs. profiling events in
'libgomp.oacc-c-c++-common/acc_prof-{init,parallel}-1.c'", leading to an
execution test failure and hang in these two tests.
This reverts changes to the tests made by the second patch, allowing the
tests to pass.
The arguments to gfc_build_array_ref were recently updated in the commit
'fortran: Use pointer arithmetic to index arrays [PR102043]', but a call
from gfc_conv_array_ref used the old function signature. This went
unnoticed due to the use of default arguments.
This patch should be merged into 'Fortran: delinearize multi-dimensional
array accesses'.
The ICE occurs during Gimple verification after the ompexp stage because
one of the arguments to the generated builtin call is of a Gimple reg type,
but isn't a Gimple value (because it is marked addressable).
This appears to be fallout from the commit "OpenACC 'kernels' decomposition:
Mark variables used in synthesized data clauses as addressable [PR100280]".
The launch dimensions have been added to the arguments of a builtin call
by oacc_set_fn_attrib, but one of the dimensions has been marked addressable.
Fixed by forcing the added arguments to be re-gimplified.
The final tail mark has no LHS, causing code that assumes its presence to
segfault. The LHS and the assignment appear to have been removed as dead
code by the cddce1 stage.
Fixed by checking for the presence of the LHS before using it.
A change that was present in the OG11 version of
'openmp: in_reduction clause support on target construct' but
not in the mainline version resulted in non-contiguous
arrays being accepted in cache clauses, only to ICE later.
Thomas Schwinge [Thu, 5 May 2022 21:01:36 +0000 (23:01 +0200)]
Refactor '-ldl' handling for libgomp proper and plugins
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.
Jakub Jelinek [Thu, 12 May 2022 06:31:20 +0000 (08:31 +0200)]
openmp: Add omp_all_memory support (C/C++ only so far)
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.
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.
Thomas Schwinge [Wed, 6 Apr 2022 10:15:28 +0000 (12:15 +0200)]
libgomp GCN plugin: Clean up always-empty 'PLUGIN_GCN_CPPFLAGS', '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.
Thomas Schwinge [Wed, 6 Apr 2022 09:31:45 +0000 (11:31 +0200)]
libgomp GCN plugin: Clean up unused references to system-provided HSA Runtime library
This is only active if GCC is 'configure'd with '--with-hsa-runtime=[...]' or
'--with-hsa-runtime-include=[...]', '--with-hsa-runtime-lib=[...]' -- which
nobody really is doing, as far as I can tell.
Originally changed for the libgomp HSA plugin in
commit b8d89b03db5f212919e4571671ebb4f5f8b1e19d (r242749)
"Remove build dependence on HSA run-time", and later propagated into the GCN
plugin, these are no longer built against system-provided HSA Runtime library.
Instead, unconditionally built against the GCC-shipped 'include/hsa*.h' header
files, and at run time does 'dlopen("libhsa-runtime64.so.1")'. It thus doesn't
make sense to consider references to system-provided HSA Runtime library during
libgomp GCN plugin build.
Thomas Schwinge [Wed, 6 Apr 2022 08:39:56 +0000 (10:39 +0200)]
libgomp testsuite: Don't amend 'LD_LIBRARY_PATH' for system-provided HSA 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.
# 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.
Fortran: Add support for OMP non-rectangular loops.
This patch adds support for OMP 5.1 "canonical loop nest form" to the
Fortran front end, marks non-rectangular loops for processing
by the middle end, and implements missing checks in the gimplifier
for additional prohibitions on non-rectangular loops.
Note that the OMP spec also prohibits non-rectangular loops with the TILE
construct; that construct hasn't been implemented yet, so that error will
need to be filled in later.
gcc/fortran/
* gfortran.h (struct gfc_omp_clauses): Add non_rectangular bit.
* openmp.cc (is_outer_iteration_variable): New function.
(expr_is_invariant): New function.
(bound_expr_is_canonical): New function.
(resolve_omp_do): Replace existing non-rectangularity error with
check for canonical form and setting non_rectangular bit.
* trans-openmp.cc (gfc_trans_omp_do): Transfer non_rectangular
flag to generated tree structure.
gcc/
* gimplify.cc (gimplify_omp_for): Update messages for SCHEDULED
and ORDERED clause conflict errors. Add check for GRAINSIZE and
NUM_TASKS on TASKLOOP.
gcc/testsuite/
* c-c++-common/gomp/loop-6.c (f3): New function to test TASKLOOP
diagnostics.
* gfortran.dg/gomp/collapse1.f90: Update expected messages.
* gfortran.dg/gomp/pr85313.f90: Remove dg-error on non-rectangular
loops that are now accepted.
* gfortran.dg/gomp/non-rectangular-loop.f90: New file.
* gfortran.dg/gomp/canonical-loop-1.f90: New file.
* gfortran.dg/gomp/canonical-loop-2.f90: New file.
Chung-Lin Tang [Fri, 17 Jun 2022 14:22:25 +0000 (22:22 +0800)]
openmp: Implement uses_allocators clause
This is a merge of:
https://gcc.gnu.org/pipermail/gcc-patches/2022-June/596412.html
For user defined allocator handles, this allows target regions to assign
memory space and traits to allocators, and automatically calls
omp_init/destroy_allocator() in the beginning/end of the target region.
For pre-defined allocators (i.e. omp_..._mem_alloc names), this is a no-op,
such clauses are not created.
Asides from the front-end portions, the target region transforms are
done in gimplify_omp_workshare.
This patch also includes added changes to enforce the "allocate allocator
must be also in a uses_allocator clause". This is done during
gimplify_scan_omp_clauses.
* c-parser.cc (c_parser_omp_clause_name): Add case for uses_allocators
clause.
(c_parser_omp_clause_uses_allocators): New function.
(c_parser_omp_all_clauses): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS case.
(OMP_TARGET_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS to mask.
* c-typeck.cc (c_finish_omp_clauses): Add case handling for
OMP_CLAUSE_USES_ALLOCATORS.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_clause_name): Add case for uses_allocators
clause.
(cp_parser_omp_clause_uses_allocators): New function.
(cp_parser_omp_all_clauses): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS case.
(OMP_TARGET_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS to mask.
* semantics.cc (finish_omp_clauses): Add case handling for
OMP_CLAUSE_USES_ALLOCATORS.
fortran/ChangeLog:
* gfortran.h (struct gfc_omp_namelist): Add memspace_sym, traits_sym
fields.
(OMP_LIST_USES_ALLOCATORS): New list enum.
* openmp.cc (enum omp_mask2): Add OMP_CLAUSE_USES_ALLOCATORS.
(gfc_match_omp_clause_uses_allocators): New function.
(gfc_match_omp_clauses): Add case to handle OMP_CLAUSE_USES_ALLOCATORS.
(OMP_TARGET_CLAUSES): Add OMP_CLAUSE_USES_ALLOCATORS.
(resolve_omp_clauses): Add "USES_ALLOCATORS" to clause_names[].
* dump-parse-tree.cc (show_omp_namelist): Handle OMP_LIST_USES_ALLOCATORS.
(show_omp_clauses): Likewise.
* trans-array.cc (gfc_conv_array_initializer): Adjust array index
to always be a created tree expression instead of NULL_TREE when zero.
* trans-openmp.cc (gfc_trans_omp_clauses): For ALLOCATE clause, handle
using gfc_trans_omp_variable for EXPR_VARIABLE exprs.
Add handling of OMP_LIST_USES_ALLOCATORS case.
* types.def (BT_FN_VOID_PTRMODE): Define.
(BT_FN_PTRMODE_PTRMODE_INT_PTR): Define.
* gimplify.cc (gimplify_scan_omp_clauses): Add checking of OpenMP target
region allocate clauses, to require a uses_allocators clause to exist
for allocators.
(gimplify_omp_workshare): Add handling of OMP_CLAUSE_USES_ALLOCATORS
for OpenMP target regions; create calls of omp_init/destroy_allocator
around target region body.
* omp-low.cc (lower_private_allocate): Adjust receiving of allocator.
(lower_rec_input_clauses): Likewise.
(create_task_copyfn): Add dereference for allocator if needed.
* system.h (startswith): New function.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/uses_allocators-1.c: New test.
* c-c++-common/gomp/uses_allocators-2.c: New test.
* c-c++-common/gomp/uses_allocators-3.c: New test.
* gfortran.dg/gomp/allocate-1.f90: Adjust testcase.
* gfortran.dg/gomp/uses_allocators-1.f90: New test.
* gfortran.dg/gomp/uses_allocators-2.f90: New test.
* gfortran.dg/gomp/uses_allocators-3.f90: New test.
Andrew Stubbs [Thu, 24 Feb 2022 17:16:13 +0000 (17:16 +0000)]
amdgcn: Add gfx90a support
This adds architecture options and multilibs for the AMD GFX90a GPUs.
It also tidies up some of the ISA selection code, and corrects a few small
mistake in the gfx908 naming.
gcc/ChangeLog:
* config.gcc (amdgcn): Accept --with-arch=gfx908 and gfx90a.
* config/gcn/gcn-opts.h (enum gcn_isa): New.
(TARGET_GCN3): Use enum gcn_isa.
(TARGET_GCN3_PLUS): Likewise.
(TARGET_GCN5): Likewise.
(TARGET_GCN5_PLUS): Likewise.
(TARGET_CDNA1): New.
(TARGET_CDNA1_PLUS): New.
(TARGET_CDNA2): New.
(TARGET_CDNA2_PLUS): New.
(TARGET_M0_LDS_LIMIT): New.
(TARGET_PACKED_WORK_ITEMS): New.
* config/gcn/gcn.cc (gcn_isa): Change to enum gcn_isa.
(gcn_option_override): Recognise CDNA ISA variants.
(gcn_omp_device_kind_arch_isa): Support gfx90a.
(gcn_expand_prologue): Make m0 init optional.
Add support for packed work items.
(output_file_start): Support gfx90a.
(gcn_hsa_declare_function_name): Support gfx90a metadata.
* config/gcn/gcn.h (TARGET_CPU_CPP_BUILTINS):Add __CDNA1__ and
__CDNA2__.
* config/gcn/gcn.md (<su>mulsi3_highpart): Use TARGET_GCN5_PLUS.
(<su>mulsi3_highpart_imm): Likewise.
(<su>mulsidi3): Likewise.
(<su>mulsidi3_imm): Likewise.
* config/gcn/gcn.opt (gpu_type): Add gfx90a.
* config/gcn/mkoffload.cc (EF_AMDGPU_MACH_AMDGCN_GFX90a): New.
(main): Support gfx90a.
* config/gcn/t-gcn-hsa: Add gfx90a multilib.
* config/gcn/t-omp-device: Add gfx90a isa.
Tobias Burnus [Mon, 23 May 2022 08:54:32 +0000 (10:54 +0200)]
OpenMP: Handle descriptors in target's firstprivate [PR104949]
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.
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.
Tobias Burnus [Thu, 12 May 2022 08:39:58 +0000 (10:39 +0200)]
Fortran: Fix proc pointer as elemental arg handling
The vtab's _callback function calls the elemental 'cb'
cb (var(:)%comp, comp_types_vtable._callback);
which gets called in a scalarization loop as 'var' might be a
nonscalar. Without the patch, that got translated as:
D.1234 = &comp_types_vtable._callback
...
cb (&(*D.4060)[S.3 + D.4071], &D.1234);
where 'D.1234' is function_type. With the patch, it remains a pointer;
i.e. D.1234 = comp... and 'cb (..., D.1234)', avoiding ME ICE.
Note: Fortran (F2018, C15100) requires that dummy arguments are
dummy data objects, which rules out dummy procs/proc-pointer dummies,
which is enforced in resolve_fl_procedure.
Thus, this change only affects the internally generated code.
gcc/fortran/ChangeLog:
* trans-array.cc (gfc_scalar_elemental_arg_saved_as_reference):
Return true for attr.proc_pointer expressions.
gcc/testsuite/ChangeLog:
* gfortran.dg/finalize_38.f90: Compile with -Ofast.
* gfortran.dg/abstract_type_6.f03: Remove dg-error as
now hidden by other errors; copy to ...
* gfortran.dg/abstract_type_6a.f03: ... here; remove
some error to diagnose the error.
* gfortran.dg/finalize_39.f90: New test.
Fortran: Fix finalization resolution with deep copy
Follow-up patch to
"Fortran/OpenMP: Support mapping of DT with allocatable components"
https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591144.html
gcc/fortran/ChangeLog:
* resolve.cc (gfc_resolve_finalizers): Also resolve allocatable comps.
Tobias Burnus [Wed, 4 May 2022 16:18:44 +0000 (18:18 +0200)]
OpenMP: Fix use_device_{addr,ptr} with in-data-sharing arg
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.
Andrew Stubbs [Wed, 13 Apr 2022 15:55:47 +0000 (16:55 +0100)]
openmp: unified_address support
This makes "requires unified_address" work by making it eqivalent to
"requires unified_shared_memory". This is more than is strictly necessary,
but should be standard compliant.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_requires): Check requires unified_address
for conflict with -foffload-memory=shared.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_requires): Check requires unified_address
for conflict with -foffload-memory=shared.
gcc/fortran/ChangeLog:
* openmp.cc (gfc_match_omp_requires): Check requires unified_address
for conflict with -foffload-memory=shared.
gcc/ChangeLog:
* omp-low.cc: Do USM transformations for "unified_address".
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/usm-4.c: New test.
* gfortran.dg/gomp/usm-4.f90: New test.
Fix a crash due to mismatch of free and GOMP_alloc.
With allocate directive, we replace the malloc calls to GOMP_alloc if
it is associated with the allocate statement. The memory was supposed
to be free-d by the implicitely generated free calls which also get
replaced. But if user explicitely deallocated the memory using the
deallocate statement, it can cause a mismatch. This commit handles
that case and also replaces the free call generated for deallocate
clause.
Also added deallocate in the testcase and tidied it up a bit.
gcc/ChangeLog:
* omp-low.cc (lower_omp_allocate): Move allocate declaration
inside loop. Set it to false at the end of condition.
Andrew Stubbs [Fri, 11 Mar 2022 12:58:38 +0000 (12:58 +0000)]
openmp: -foffload-memory=pinned
Implement the -foffload-memory=pinned option such that libgomp is
instructed to enable fully-pinned memory at start-up. The option is
intended to provide a performance boost to certain offload programs without
modifying the code.
This feature only works on Linux, at present, and simply calls mlockall to
enable always-on memory pinning. It requires that the ulimit feature is
set high enough to accommodate all the program's memory usage.
In this mode the ompx_pinned_memory_alloc feature is disabled as it is not
needed and may conflict.
Backport of the patch posted at
https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591354.html
gcc/ChangeLog:
* omp-low.cc (omp_enable_pinned_mode): New function.
(execute_lower_omp): Call omp_enable_pinned_mode.
libgomp/ChangeLog:
* config/linux/allocator.c (always_pinned_mode): New variable.
(GOMP_enable_pinned_mode): New function.
(linux_memspace_alloc): Disable pinning when always_pinned_mode set.
(linux_memspace_calloc): Likewise.
(linux_memspace_free): Likewise.
(linux_memspace_realloc): Likewise.
* libgomp.map (GOMP_5.1.1): New version space with
GOMP_enable_pinned_mode.
* testsuite/libgomp.c/alloc-pinned-7.c: New test.
openmp: Use libgomp memory allocation functions with unified shared memory.
This patches changes calls to malloc/free/calloc/realloc/aligned_alloc and
operator new to memory allocation functions in libgomp with
allocator=ompx_unified_shared_mem_alloc. This helps existing code to benefit
from the unified shared memory. The libgomp does the correct thing with all
the mapping constructs and there is no memory copies if the pointer is pointing
to unified shared memory.
We only replace replacable new operator and not the class member or placement new.
Backport of a patch posted at
https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591353.html
gcc/ChangeLog:
* omp-low.cc (usm_transform): New function.
(make_pass_usm_transform): Likewise.
(class pass_usm_transform): New.
* passes.def: Add pass_usm_transform.
* tree-pass.h (make_pass_usm_transform): New declaration.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/usm-2.c: New test.
* c-c++-common/gomp/usm-3.c: New test.
* g++.dg/gomp/usm-1.C: New test.
* g++.dg/gomp/usm-2.C: New test.
* g++.dg/gomp/usm-3.C: New test.
* gfortran.dg/gomp/usm-2.f90: New test.
* gfortran.dg/gomp/usm-3.f90: New test.
libgomp/ChangeLog:
* testsuite/libgomp.c/usm-6.c: New test.
* testsuite/libgomp.c++/usm-1.C: Likewise.
Andrew Stubbs [Fri, 11 Mar 2022 12:37:58 +0000 (12:37 +0000)]
openmp, nvptx: ompx_unified_shared_mem_alloc
This adds support for using Cuda Managed Memory with omp_alloc. It will be
used as the underpinnings for "requires unified_shared_memory" in a later
patch.
There are two new predefined allocators, ompx_unified_shared_mem_alloc and
ompx_host_mem_alloc, plus corresponding memory spaces, which can be used to
allocate memory in the "managed" space and explicitly on the host (it is
intended that "malloc" will be intercepted by the compiler).
The nvptx plugin is modified to make the necessary Cuda calls, and libgomp
is modified to switch to shared-memory mode for USM allocated mappings.
Backport of the patch posted at
https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591352.html
libgomp/ChangeLog:
* allocator.c (omp_max_predefined_alloc): Update.
(omp_aligned_alloc): Don't fallback ompx_host_mem_alloc.
(omp_aligned_calloc): Likewise.
(omp_realloc): Likewise.
* config/linux/allocator.c (linux_memspace_alloc): Handle USM.
(linux_memspace_calloc): Handle USM.
(linux_memspace_free): Handle USM.
(linux_memspace_realloc): Handle USM.
* config/nvptx/allocator.c (nvptx_memspace_alloc): Reject
ompx_host_mem_alloc.
(nvptx_memspace_calloc): Likewise.
(nvptx_memspace_realloc): Likewise.
* libgomp-plugin.h (GOMP_OFFLOAD_usm_alloc): New prototype.
(GOMP_OFFLOAD_usm_free): New prototype.
(GOMP_OFFLOAD_is_usm_ptr): New prototype.
* libgomp.h (gomp_usm_alloc): New prototype.
(gomp_usm_free): New prototype.
(gomp_is_usm_ptr): New prototype.
(struct gomp_device_descr): Add USM functions.
* omp.h.in (omp_memspace_handle_t): Add ompx_unified_shared_mem_space
and ompx_host_mem_space.
(omp_allocator_handle_t): Add ompx_unified_shared_mem_alloc and
ompx_host_mem_alloc.
* omp_lib.f90.in: Likewise.
* plugin/plugin-nvptx.c (nvptx_alloc): Add "usm" parameter.
Call cuMemAllocManaged as appropriate.
(GOMP_OFFLOAD_alloc): Move internals to ...
(GOMP_OFFLOAD_alloc_1): ... this, and add usm parameter.
(GOMP_OFFLOAD_usm_alloc): New function.
(GOMP_OFFLOAD_usm_free): New function.
(GOMP_OFFLOAD_is_usm_ptr): New function.
* target.c (gomp_map_vars_internal): Add USM support.
(gomp_usm_alloc): New function.
(gomp_usm_free): New function.
(gomp_load_plugin_for_device): New function.
* testsuite/libgomp.c/usm-1.c: New test.
* testsuite/libgomp.c/usm-2.c: New test.
* testsuite/libgomp.c/usm-3.c: New test.
* testsuite/libgomp.c/usm-4.c: New test.
* testsuite/libgomp.c/usm-5.c: New test.
Andrew Stubbs [Fri, 11 Mar 2022 14:33:11 +0000 (14:33 +0000)]
openmp, nvptx: low-lat memory access traits
The NVPTX low latency memory is not accessible outside the team that allocates
it, and therefore should be unavailable for allocators with the access trait
"all". This change means that the omp_low_lat_mem_alloc predefined
allocator now implicitly implies the "pteam" trait.
Backport of a patch posted at
https://gcc.gnu.org/pipermail/gcc-patches/2022-January/589355.html
libgomp/ChangeLog:
* allocator.c (MEMSPACE_VALIDATE): New macro.
(omp_aligned_alloc): Use MEMSPACE_VALIDATE.
(omp_aligned_calloc): Likewise.
(omp_realloc): Likewise.
* config/nvptx/allocator.c (nvptx_memspace_validate): New function.
(MEMSPACE_VALIDATE): New macro.
* testsuite/libgomp.c/allocators-4.c (main): Add access trait.
* testsuite/libgomp.c/allocators-6.c (main): Add access trait.
* testsuite/libgomp.c/allocators-7.c: New test.
Andrew Stubbs [Fri, 11 Mar 2022 12:33:06 +0000 (12:33 +0000)]
libgomp, openmp: Add ompx_pinned_mem_alloc
This creates a new predefined allocator as a shortcut for using pinned
memory with OpenMP. The name uses the OpenMP extension space and is
intended to be consistent with other OpenMP implementations currently in
development.
The allocator is equivalent to using a custom allocator with the pinned
trait and the null fallback trait.
Backport of a patch posted at
https://gcc.gnu.org/pipermail/gcc-patches/2022-January/588951.html
* allocator.c (omp_max_predefined_alloc): Update.
(omp_aligned_alloc): Support ompx_pinned_mem_alloc.
(omp_free): Likewise.
(omp_aligned_calloc): Likewise.
(omp_realloc): Likewise.
* omp.h.in (omp_allocator_handle_t): Add ompx_pinned_mem_alloc.
* omp_lib.f90.in: Add ompx_pinned_mem_alloc.
* testsuite/libgomp.c/alloc-pinned-5.c: New test.
* testsuite/libgomp.c/alloc-pinned-6.c: New test.
* testsuite/libgomp.fortran/alloc-pinned-1.f90: New test.
Andrew Stubbs [Fri, 11 Mar 2022 12:12:39 +0000 (12:12 +0000)]
libgomp: pinned memory
Implement the OpenMP pinned memory trait on Linux hosts using the mlock
syscall. Pinned allocations are performed using mmap, not malloc, to ensure
that they can be unpinned safely when freed.
Backport of a patch posted at
https://gcc.gnu.org/pipermail/gcc-patches/2022-January/588360.html
libgomp/ChangeLog:
* allocator.c (MEMSPACE_ALLOC): Add PIN.
(MEMSPACE_CALLOC): Add PIN.
(MEMSPACE_REALLOC): Add PIN.
(MEMSPACE_FREE): Add PIN.
(xmlock): New function.
(omp_init_allocator): Don't disallow the pinned trait.
(omp_aligned_alloc): Add pinning to all MEMSPACE_* calls.
(omp_aligned_calloc): Likewise.
(omp_realloc): Likewise.
(omp_free): Likewise.
* config/linux/allocator.c: New file.
* config/nvptx/allocator.c (MEMSPACE_ALLOC): Add PIN.
(MEMSPACE_CALLOC): Add PIN.
(MEMSPACE_REALLOC): Add PIN.
(MEMSPACE_FREE): Add PIN.
* testsuite/libgomp.c/alloc-pinned-1.c: New test.
* testsuite/libgomp.c/alloc-pinned-2.c: New test.
* testsuite/libgomp.c/alloc-pinned-3.c: New test.
* testsuite/libgomp.c/alloc-pinned-4.c: New test.
Backport of a patch posted at
https://gcc.gnu.org/pipermail/gcc-patches/2022-January/588372.html
This patch looks for malloc/free calls that were generated by allocate statement
that is associated with allocate directive and replaces them with GOMP_alloc
and GOMP_free.
gcc/ChangeLog:
* omp-low.cc (scan_sharing_clauses): Handle OMP_CLAUSE_ALLOCATOR.
(scan_omp_allocate): New.
(scan_omp_1_stmt): Call it.
(lower_omp_allocate): New function.
(lower_omp_1): Call it.
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/allocate-6.f90: Add tests.
* gfortran.dg/gomp/allocate-7.f90: New test.
* gfortran.dg/gomp/allocate-8.f90: New test.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/allocate-2.f90: New test.
Handle cleanup of omp allocated variables (OpenMP 5.0).
Currently we are only handling omp allocate directive that is associated
with an allocate statement. This statement results in malloc and free calls.
The malloc calls are easy to get to as they are in the same block as allocate
directive. But the free calls come in a separate cleanup block. To help any
later passes finding them, an allocate directive is generated in the
cleanup block with kind=free. The normal allocate directive is given
kind=allocate.
Backport of a patch posted at
https://gcc.gnu.org/pipermail/gcc-patches/2022-January/588370.html
gcc/fortran/ChangeLog:
* gfortran.h (struct access_ref): Declare new members
omp_allocated and omp_allocated_end.
* openmp.cc (gfc_match_omp_allocate): Set new_st.resolved_sym to
NULL.
(prepare_omp_allocated_var_list_for_cleanup): New function.
(gfc_resolve_omp_allocate): Call it.
* trans-decl.cc (gfc_trans_deferred_vars): Process omp_allocated.
* trans-openmp.cc (gfc_trans_omp_allocate): Set kind for the stmt
generated for allocate directive.
This is backport of a patch posted in
https://gcc.gnu.org/pipermail/gcc-patches/2022-February/590597.html
An allocate clause in target region must specify an allocator
unless the compilation unit has requires construct with
dynamic_allocators clause. Current implementation of the allocate
clause did not check for this restriction. This patch fills that
gap.
gcc/ChangeLog:
* omp-low.cc (omp_maybe_offloaded_ctx): New prototype.
(scan_sharing_clauses): Check a restriction on allocate clause.
Tobias Burnus [Tue, 1 Mar 2022 15:35:08 +0000 (16:35 +0100)]
Fortran/OpenMP: Support mapping of DT with allocatable components
gcc/fortran/ChangeLog:
* class.cc (finalization_scalarizer): Mark syms as artificial.
(generate_callback_wrapper): New.
(gfc_find_derived_vtab): Call it, add _callback comp.
* f95-lang.cc (LANG_HOOKS_OMP_DEEP_MAPPING,
LANG_HOOKS_OMP_DEEP_MAPPING_P,
LANG_HOOKS_OMP_DEEP_MAPPING_CNT): Redeinfe
* gfortran.h (gfc_import_iso_c_binding_module,
GFC_CLASS_CALLBACK_DEFAULT_FLAG, GFC_CLASS_CALLBACK_VTABLE_FLAG,
GFC_CLASS_CB_ALLOCATABLE, GFC_CLASS_CB_POINTER,
GFC_CLASS_CB_PROC_POINTER, GFC_CLASS_CB_VTABLE,
GFC_CLASS_CB_VPTR): New.
* match.cc (select_type_set_tmp): Propagate allocatable property.
* module.cc (MOD_VERSION): Bump due to vtab change.
(import_iso_c_binding_module): New import_all arg.
(gfc_import_iso_c_binding_module): New.
(gfc_use_module): Update call.
* openmp.cc (resolve_omp_clauses): Accept DT with alloc comps.
* resolve.cc (gfc_resolve_formal_arglist, gfc_resolve_intrinsic,
resolve_fl_procedure, resolve_types): Permit some violations
for internal code.
* trans-array.cc (gfc_conv_descriptor_stride_get,
gfc_tree_array_size, gfc_full_array_size): Update
for GFC_TYPE_ARRAY_AKIND change.
(gfc_conv_expr_descriptor): Likewise; permit calling with tree code.
* trans-expr.cc (VTABLE_CALLBACK_FIELD): Add.
(VTAB_GET_FIELD_GEN): Use it.
(VTABLE_DEALLOCATE_FIELD): Undef at the end.
(gfc_conv_expr_reference): Fixes; avoid unneccessary temp var.
* trans-intrinsic.cc (gfc_conv_intrinsic_sizeof,
gfc_conv_associated): Fix class and comp-ref handling.
(conv_isocbinding_function): Remove buggy code.
* trans-openmp.cc (gfc_has_alloc_comps): Add ptr_ok arg.
(gfc_omp_private_outer_ref, gfc_walk_alloc_comps,
gfc_omp_clause_default_ctor, gfc_omp_clause_copy_ctor,
gfc_omp_clause_assign_op, gfc_omp_clause_dtor,
(gfc_omp_finish_clause): Update call.
(GFC_MAP_TOKEN_DATA, GFC_MAP_TOKEN_SIZES, GFC_MAP_TOKEN_KINDS,
GFC_MAP_TOKEN_DATA_OFFSET, GFC_MAP_TOKEN_OFFSET,
GFC_MAP_TOKEN_FLAGS, GFC_MAP_TOKEN_DETACH): Define.
(gfc_omp_get_token_data, gfc_omp_get_token_sizes,
gfc_omp_get_token_kinds, gfc_omp_get_token_offset_data,
gfc_omp_get_token_offset, gfc_omp_get_token_flags,
gfc_omp_get_token_detach, gfc_omp_get_map_token_type,
gfc_omp_get_cb_type, gfc_omp_gen_deep_map_fn,
gfc_omp_deep_mapping_map, gfc_omp_deep_mapping_item,
gfc_omp_deep_mapping_comps, gfc_omp_gen_simple_loop,
gfc_omp_get_array_size, gfc_omp_elmental_loop,
gfc_omp_deep_map_kind_p, gfc_omp_deep_mapping_int_p,
gfc_omp_deep_mapping_p, gfc_omp_deep_mapping_do),
gfc_omp_deep_mapping_cnt, gfc_omp_deep_mapping): New.
(gfc_trans_omp_array_section): Save clause decl to survive gimplifying.
(gfc_trans_omp_clauses): Likewise; fixes.
* trans-types.cc (gfc_build_array_type, gfc_get_derived_type,
gfc_get_array_descr_info): Update array kind to distinguish
different assumed-rank arrays.
* trans.h (gfc_class_vtab_callback_get, gfc_omp_deep_mapping_p,
gfc_omp_deep_mapping_cnt, gfc_omp_deep_mapping): New prototypes.
(enum gfc_array_kind): Additional GFC_ARRAY_ASSUMED_RANK_* entries.
gcc/ChangeLog:
* langhooks-def.h (lhd_omp_deep_mapping_p,
lhd_omp_deep_mapping_cnt, lhd_omp_deep_mapping): New.
(LANG_HOOKS_OMP_DEEP_MAPPING_P, LANG_HOOKS_OMP_DEEP_MAPPING_CNT,
LANG_HOOKS_OMP_DEEP_MAPPING): Define.
(LANG_HOOKS_DECLS): Use it.
* langhooks.cc (lhd_omp_deep_mapping_p, lhd_omp_deep_mapping_cnt,
lhd_omp_deep_mapping): New stubs.
* langhooks.h (struct lang_hooks_for_decls): Add new hooks
* omp-expand.cc (expand_omp_target): Handle dynamic-size
addr/sizes/kinds arrays.
* omp-low.cc (build_sender_ref, fixup_child_record_type,
scan_sharing_clauses, lower_omp_target): Update to handle
new hooks and dynamic-size addr/sizes/kinds arrays.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/allocatable-comp.f90: New test.
* testsuite/libgomp.fortran/map-alloc-comp-3.f90: New test.
* testsuite/libgomp.fortran/map-alloc-comp-4.f90: New test.
* testsuite/libgomp.fortran/map-alloc-comp-5.f90: New test.
* testsuite/libgomp.fortran/map-alloc-comp-6.f90: New test.
* testsuite/libgomp.fortran/map-alloc-comp-7.f90: New test.
Chung-Lin Tang [Thu, 24 Feb 2022 09:07:48 +0000 (01:07 -0800)]
openmp: Handle C/C++ array reference base-pointers in array sections
In cases where a program constructs its own deep-copying for arrays-of-pointers,
e.g:
#pragma omp target enter data map(to:level->vectors[:N])
for (i = 0; i < N; i++)
#pragma omp target enter data map(to:level->vectors[i][:N])
We need to treat the part of the array reference before the array section
as a base-pointer (here 'level->vectors[i]'), providing pointer-attachment
behavior.
This patch adds this inside handle_omp_array_sections(), tracing the whole
sequence of array dimensions, creating a whole base-pointer reference
iteratively using build_array_ref(). The conditions are that each of the
"absorbed" dimensions must be length==1, and the final reference must be
of pointer-type (so that pointer attachment makes sense).
openmp: Improve handling of nested OpenMP metadirectives in C and C++
This patch fixes a misparsing issue when encountering code like:
#pragma omp metadirective when {<selector_set>={...}: A)
#pragma omp metadirective when (<selector_set>={...}: B)
When called for the first metadirective, analyze_metadirective_body would
stop just before the colon in the second metadirective because it naively
assumes that the '}' marks the end of a code block.
The assertion for clauses to end parsing at the same point is now disabled
if a parse error has occurred during the parsing of the clause, since some
tokens may not be consumed if a parse error cuts parsing short.
gcc/c/
* c-parser.cc (c_parser_omp_construct): Move handling of
PRAGMA_OMP_METADIRECTIVE from here...
(c_parser_pragma): ...to here.
(analyze_metadirective_body): Check that the bracket nesting level
is also zero before stopping the adding of tokens on encountering a
close brace.
(c_parser_omp_metadirective): Modify function signature and update.
Do not assert on remaining tokens if there has been a parse error.
gcc/cp/
* parser.cc (cp_parser_omp_construct): Move handling of
PRAGMA_OMP_METADIRECTIVE from here...
(cp_parser_pragma): ...to here.
(analyze_metadirective_body): Check that the bracket
nesting level is also zero before stopping the adding of tokens on
encountering a close brace.
(cp_parser_omp_metadirective): Modify function signature and update.
Do not assert on remaining tokens if there has been a parse error.
gcc/testsuite/
* c-c++-common/gomp/metadirective-1.c (f): Add test for
improperly nested metadirectives.
openmp: More Fortran front-end fixes for metadirectives
This adds a check for declarative OpenMP directives in metadirective
variants (already present in the C/C++ front-ends), and fixes an
ICE when an empty metadirective (i.e. just '!$omp metadirective')
is presented.
gcc/fortran/
* gfortran.h (is_omp_declarative_stmt): New.
* openmp.cc (match_omp_metadirective): Reject declarative OpenMP
directives with 'sorry'.
* parse.cc (parse_omp_metadirective_body): Check that state stack head
is non-null before dereferencing.
(is_omp_declarative_stmt): New.
gcc/testsuite/
* gfortran.dg/gomp/metadirective-2.f90 (main): Test empty
metadirective.
openmp: Eliminate non-matching metadirective variants early in Fortran front-end
This patch checks during parsing if a metadirective selector is both
resolvable and non-matching - if so, it is removed from further
consideration. This is both more efficient, and avoids spurious
syntax errors caused by considering combinations of selectors that
lead to invalid combinations of OpenMP directives, when that
combination would never arise in the first place.
This exposes another bug - when metadirectives that are not of the
begin-end variety are nested, we might have to drill up through
multiple layers of the state stack to reach the state for the
next statement. This is now fixed.
gcc/
* omp-general.cc (DELAY_METADIRECTIVES_AFTER_LTO): Check that cfun is
non-null before derefencing.
gcc/fortran/
* decl.cc (gfc_match_end): Search for first previous state that is not
COMP_OMP_METADIRECTIVE.
* gfortran.h (gfc_skip_omp_metadirective_clause): Add prototype.
* openmp.cc (match_omp_metadirective): Skip clause if
result of gfc_skip_omp_metadirective_clause is true.
* trans-openmp.cc (gfc_trans_omp_set_selector): Add argument and
disable expression conversion if false.
(gfc_skip_omp_metadirective_clause): New.
With the OpenMP 5.0/5.1 specifications, we would expect A to be selected in
the metadirective when the target call is made, but B when f is called
directly outside of a target context. However, since GCC does not have
separate copies of f for local and target calls, and the construct selector
is static, it must be resolved one way or the other at compile-time (currently
in the favour of selecting A), which may be unexpected behaviour.
This patch attempts to detect the above situation, and will emit a warning
if found.
gcc/
* gimplify.cc (gimplify_omp_metadirective): Mark offloadable functions
containing metadirectives with 'construct={target}' in the selector.
* omp-general.cc (omp_has_target_constructor_p): New.
* omp-general.h (omp_has_target_constructor_p): New prototype.
* omp-low.cc (lower_omp_1): Emit warning if marked functions called
outside of a target context.
openmp: Add support for streaming metadirectives and resolving them after LTO
This patch adds support for streaming metadirective Gimple statements during
LTO, and adds a metadirective expansion pass that runs after LTO. This is
required for metadirectives with selectors that can only be resolved from
within the accel compiler.
openmp: Add support for resolving metadirectives during parsing and Gimplification
This adds support for resolving metadirectives according to the OpenMP 5.1
specification. The variants are sorted by score, then gathered into a list
of dynamic replacement candidates. The metadirective is then expanded into
a sequence of 'if..else' statements to test the dynamic selector and execute
the variant if the selector is satisfied.
If any of the selectors in the list are unresolvable, GCC will give up on
resolving the metadirective and try again later.
This adds a new Gimple statement type GIMPLE_OMP_METADIRECTIVE, which
represents the metadirective in Gimple. In high Gimple, the statement
contains the body of the directive variants, whereas in low Gimple, it
only contains labels to the bodies.
This patch adds support for converting metadirectives from tree to Gimple
form, and handling of the Gimple form (Gimple lowering, OpenMP lowering
and expansion, inlining, SSA handling etc).
Metadirectives should be resolved before they reach the back-end, otherwise
the compiler will crash as GCC does not know how to convert metadirective
Gimple statements to RTX.
This patch implements parsing for the OpenMP metadirective introduced in
OpenMP 5.0. Metadirectives are parsed into an OMP_METADIRECTIVE node,
with the variant clauses forming a chain accessible via
OMP_METADIRECTIVE_CLAUSES. Each clause contains the context selector
and tree for the variant.
User conditions in the selector are now permitted to be non-constant when
used in metadirectives as specified in OpenMP 5.1.
gcc/
* omp-general.cc (omp_context_selector_matches): Add extra argument.
(omp_resolve_metadirective): New stub function.
* omp-general.h (struct omp_metadirective_variant): New.
(omp_context_selector_matches): Add extra argument.
(omp_resolve_metadirective): New prototype.
* tree.def (OMP_METADIRECTIVE): New.
* tree.h (OMP_METADIRECTIVE_CLAUSES): New macro.
gcc/c/
* c-parser.cc (c_parser_skip_to_end_of_block_or_statement): Handle
parentheses in statement.
(c_parser_omp_metadirective): New prototype.
(c_parser_omp_context_selector): Add extra argument. Allow
non-constant expressions.
(c_parser_omp_context_selector_specification): Add extra argument and
propagate it to c_parser_omp_context_selector.
(analyze_metadirective_body): New.
(c_parser_omp_metadirective): New.
(c_parser_omp_construct): Handle PRAGMA_OMP_METADIRECTIVE.
gcc/c-family/
* c-common.h (enum c_omp_directive_kind): Add C_OMP_DIR_META.
(c_omp_expand_metadirective): New prototype.
* c-gimplify.cc (genericize_omp_metadirective_stmt): New.
(c_genericize_control_stmt): Handle OMP_METADIRECTIVE tree nodes.
* c-omp.cc (omp_directives): Classify metadirectives as C_OMP_DIR_META.
(c_omp_expand_metadirective): New stub function.
* c-pragma.cc (omp_pragmas): Add entry for metadirective.
* c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_METADIRECTIVE.
libgomp/
* config/nvptx/team.c (gomp_nvptx_main): Initialize shared_pool_size
to zero. Do not use dynamic_smem_size register if PTX version lower
than 4.1.
Andrew Stubbs [Fri, 3 Dec 2021 17:46:41 +0000 (17:46 +0000)]
libgomp, nvptx: low-latency memory allocator
This patch adds support for allocating low-latency ".shared" memory on
NVPTX GPU device, via the omp_low_lat_mem_space and omp_alloc. The memory
can be allocated, reallocated, and freed using a basic but fast algorithm,
is thread safe and the size of the low-latency heap can be configured using
the GOMP_NVPTX_LOWLAT_POOL environment variable.
The use of the PTX dynamic_smem_size feature means that the minimum version
requirement is now bumped to 4.1 (still old at this point).
libgomp/ChangeLog:
* allocator.c (MEMSPACE_ALLOC): New macro.
(MEMSPACE_CALLOC): New macro.
(MEMSPACE_REALLOC): New macro.
(MEMSPACE_FREE): New macro.
(dynamic_smem_size): New constants.
(omp_alloc): Use MEMSPACE_ALLOC.
Implement fall-backs for predefined allocators.
(omp_free): Use MEMSPACE_FREE.
(omp_calloc): Use MEMSPACE_CALLOC.
Implement fall-backs for predefined allocators.
(omp_realloc): Use MEMSPACE_REALLOC.
Implement fall-backs for predefined allocators.
* config/nvptx/team.c (__nvptx_lowlat_heap_root): New variable.
(__nvptx_lowlat_pool): New asm varaible.
(gomp_nvptx_main): Initialize the low-latency heap.
* plugin/plugin-nvptx.c (lowlat_pool_size): New variable.
(GOMP_OFFLOAD_init_device): Read the GOMP_NVPTX_LOWLAT_POOL envvar.
(GOMP_OFFLOAD_run): Apply lowlat_pool_size.
* config/nvptx/allocator.c: New file.
* testsuite/libgomp.c/allocators-1.c: New test.
* testsuite/libgomp.c/allocators-2.c: New test.
* testsuite/libgomp.c/allocators-3.c: New test.
* testsuite/libgomp.c/allocators-4.c: New test.
* testsuite/libgomp.c/allocators-5.c: New test.
* testsuite/libgomp.c/allocators-6.c: New test.
This updates the bundled cuda.h header file to include some new API calls and
constants that are now used in the code.
This patch should be included when the "libgomp, nvptx: low-latency memory
allocator" or "openmp: Add support for 'target_device' context selector set"
patches are upstreamed.
include/
* cuda/cuda.h (CUdevice_attribute): Add definitions for
CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR and
CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR.
(CUmemAttach_flags): New.
(CUpointer_attribute): New.
(cuMemAllocManaged): New prototype.
(cuPointerGetAttribute): New prototype.
libgomp/
* plugin/cuda-lib.def (cuMemAllocManaged): Add new call.
(cuPointerGetAttribute): Likewise.
Jakub Jelinek [Tue, 21 Jun 2022 09:40:16 +0000 (11:40 +0200)]
ifcvt: Don't introduce trapping or faulting reads in noce_try_sign_mask [PR106032]
noce_try_sign_mask as documented will optimize
if (c < 0)
x = t;
else
x = 0;
into x = (c >> bitsm1) & t;
The optimization is done if either t is unconditional
(e.g. for
x = t;
if (c >= 0)
x = 0;
) or if it is cheap. We already check that t doesn't have side-effects,
but if t is conditional, we need to punt also if it may trap or fault,
as we make it unconditional.
I've briefly skimmed other noce_try* optimizations and didn't find one that
would suffer from the same problem.
2022-06-21 Jakub Jelinek <jakub@redhat.com>
PR rtl-optimization/106032
* ifcvt.cc (noce_try_sign_mask): Punt if !t_unconditional, and
t may_trap_or_fault_p, even if it is cheap.
Jakub Jelinek [Tue, 21 Jun 2022 09:38:59 +0000 (11:38 +0200)]
expand: Fix up expand_cond_expr_using_cmove [PR106030]
If expand_cond_expr_using_cmove can't find a cmove optab for a particular
mode, it tries to promote the mode and perform the cmove in the promoted
mode.
The testcase in the patch ICEs on arm because in that case we pass temp which
has the promoted mode (SImode) as target to expand_operands where the
operands have the non-promoted mode (QImode).
Later on the function uses paradoxical subregs:
if (GET_MODE (op1) != mode)
op1 = gen_lowpart (mode, op1);
if (GET_MODE (op2) != mode)
op2 = gen_lowpart (mode, op2);
to change the operand modes.
The following patch fixes it by passing NULL_RTX as target if it has
promoted mode.
2022-06-21 Jakub Jelinek <jakub@redhat.com>
PR middle-end/106030
* expr.cc (expand_cond_expr_using_cmove): Pass NULL_RTX instead of
temp to expand_operands if mode has been promoted.
git://gcc.gnu.org / gcc.git / log