| Age | Commit message (Collapse) | Author |
|---|
|
gcc/testsuite/ChangeLog
* c-c++-common/gomp/delim-declare-variant-1.c: New.
* c-c++-common/gomp/delim-declare-variant-2.c: New.
* c-c++-common/gomp/delim-declare-variant-3.c: New.
* c-c++-common/gomp/delim-declare-variant-4.c: New.
* c-c++-common/gomp/delim-declare-variant-5.c: New.
* c-c++-common/gomp/delim-declare-variant-6.c: New.
* c-c++-common/gomp/delim-declare-variant-7.c: New.
* c-c++-common/gomp/delim-declare-variant-8.c: New.
* c-c++-common/gomp/delim-declare-variant-9.c: New.
libgomp/ChangeLog
* testsuite/libgomp.c-c++-common/delim-declare-variant-1.c: New.
* testsuite/libgomp.c-c++-common/delim-declare-variant-2.c: New.
Co-Authored-By: Tobias Burnus <tburnus@baylibre.com>
|
|
This patch implements C++ support for the "begin declare variant"
construct. The OpenMP specification is hazy on interaction of this
feature with C++ language features. Variant functions in classes are
supported but must be defined as members in the class definition,
using an unqualified name for the base function which also must be
present in that class. Similarly variant functions in a namespace can
only be defined in that namespace using an unqualified name for a base
function already declared in that namespace. Variants for template
functions or inside template classes seem to (mostly) work.
gcc/c-family/ChangeLog
* c-omp.cc (c_omp_directives): Uncomment "begin declare variant"
and "end declare variant".
gcc/cp/ChangeLog
* cp-tree.h (struct cp_omp_declare_variant_attr): New.
(struct saved_scope): Add omp_declare_variant_attribute field.
* decl.cc (omp_declare_variant_finalize_one): Add logic to inject
"this" parameter for method calls.
* parser.cc (cp_parser_skip_to_pragma_omp_end_declare_variant): New.
(cp_parser_translation_unit): Handle leftover "begin declare variant"
functions.
(omp_start_variant_function): New.
(omp_finish_variant_function): New.
(omp_maybe_record_variant_base): New.
(cp_parser_init_declarator): Handle variant functions.
(cp_parser_class_specifier): Handle deferred lookup of base functions
when the entire class has been seen.
(cp_parser_member_declaration): Handle variant functions.
(cp_finish_omp_declare_variant): Merge context selectors if in
a "begin declare variant" block.
(cp_parser_omp_begin): Match "omp begin declare variant". Adjust
error messages.
(cp_parser_omp_end): Match "omp end declare variant".
* parser.h (struct omp_begin_declare_variant_map_entry): New.
(struct cp_parser): Add omp_begin_declare_variant_map field.
* semantics.cc (finish_translation_unit): Detect unmatched
"omp begin declare variant".
gcc/testsuite/ChangeLog
* g++.dg/gomp/delim-declare-variant-1.C: New.
* g++.dg/gomp/delim-declare-variant-2.C: New.
* g++.dg/gomp/delim-declare-variant-3.C: New.
* g++.dg/gomp/delim-declare-variant-4.C: New.
* g++.dg/gomp/delim-declare-variant-5.C: New.
* g++.dg/gomp/delim-declare-variant-6.C: New.
* g++.dg/gomp/delim-declare-variant-7.C: New.
* g++.dg/gomp/delim-declare-variant-40.C: New.
* g++.dg/gomp/delim-declare-variant-41.C: New.
* g++.dg/gomp/delim-declare-variant-50.C: New.
* g++.dg/gomp/delim-declare-variant-51.C: New.
* g++.dg/gomp/delim-declare-variant-52.C: New.
* g++.dg/gomp/delim-declare-variant-70.C: New.
* g++.dg/gomp/delim-declare-variant-71.C: New.
libgomp/
* testsuite/libgomp.c++/bdv_module1.C: New.
* testsuite/libgomp.c++/bdv_module1_main.C: New.
* testsuite/libgomp.c++/bdv_module2.C: New.
* testsuite/libgomp.c++/bdv_module2_impl.C: New.
* testsuite/libgomp.c++/bdv_module2_main.C: New.
* testsuite/libgomp.c++/bdv_module3.C: New.
* testsuite/libgomp.c++/bdv_module3_impl.C: New.
* testsuite/libgomp.c++/bdv_module3_main.C: New.
* testsuite/libgomp.c++/delim-declare-variant-1.C: New.
* testsuite/libgomp.c++/delim-declare-variant-2.C: New.
* testsuite/libgomp.c++/delim-declare-variant-7.C: New.
Co-Authored-By: Julian Brown <julian@codesourcery.com>
Co-Authored-By: waffl3x <waffl3x@baylibre.com>
|
|
In the testcases, the kernels scheduled on queues 11, 12, 13, 14 have
data dependencies on, respectively, 'b', 'c', 'd', and 'e', as they
write to them.
However, they also have a data dependency on 'a' and 'N', as they read
those.
Previously, the testcases exited 'a' on queue 10 and 'N' on queue 15,
meaning that it was possible for the aforementioned kernels to execute
and to have 'a' and 'N' pulled under their feet.
This patch adds waits for each of the kernels onto queue 10 before
freeing 'a', guaranteeing that 'a' outlives the kernels, and the same on
'N'.
libgomp/ChangeLog:
* testsuite/libgomp.oacc-c-c++-common/data-2-lib.c (explanatory
header): Fix typo.
(main): Insert waits on kernels reading 'a' into queue 10 before
exiting 'a', and waits on kernels reading 'N' into queue 15
before exiting 'N'.
* testsuite/libgomp.oacc-c-c++-common/data-2.c: Ditto.
|
|
Due to the if statement in ipa_compute_jump_functions_for_bb, callback
edges were never constructed for builtin functions unless LTO was
enabled. This patch corrects this behavior, allowing GCC to optimize
callbacks more broadly. It also extends our testing capabilities.
gcc/ChangeLog:
* attr-callback.cc (callback_edge_callee_has_attr): New
function.
* attr-callback.h (callback_edge_callee_has_attr): New function
decl.
* ipa-prop.cc (ipa_compute_jump_functions_for_bb): Don't skip
callback carriers when calculating jump functions.
libgomp/ChangeLog:
* testsuite/libgomp.c/ipcp-cb-spec1.c: Remove LTO requirement.
* testsuite/libgomp.c/ipcp-cb-spec2.c: Likewise.
* testsuite/libgomp.c/ipcp-cb1.c: Likewise.
Signed-off-by: Josef Melcr <josef.melcr@suse.com>
|
|
OpenMP/OpenACC array sections, generally expr[expr:expr] or
expr[expr:expr:expr] can have any of the exprs between [ and ]
omitted, low-bound (first defaults to 0, last (stride) defaults to
1 and the middle (length) for some arrays defaults to
ceil((size − lower_bound)/stride).
People have been writing this for years without spaces between [ and :
and : and ] when that expr has been omitted, but guess for C++26
one needs to add a space. I think [ :: ] isn't going to be parsed
as the same as [ : : ] either.
gcc/testsuite/
* c-c++-common/goacc/cache-3-1.c: Add dg-skip-if for c++26.
* g++.dg/goacc/data-clause-2.C: Likewise.
* g++.dg/gomp/allocate-3.C: Likewise.
* c-c++-common/gomp/affinity-2.c: Use { c || c++23_down } effective
target.
* c-c++-common/goacc/cache-3-2.c: Replace [: in OpenMP or OpenACC
pragmas or attributes with [ : and :] with : ].
* c-c++-common/goacc/data-clause-1.c: Likewise.
* c-c++-common/goacc/data-clause-2.c: Likewise.
* c-c++-common/goacc/data-clause-duplicate-1.c: Likewise.
* c-c++-common/goacc/mdc-2.c: Likewise.
* c-c++-common/goacc/readonly-1.c: Likewise.
* c-c++-common/gomp/allocate-4.c: Likewise.
* c-c++-common/gomp/clauses-3.c: Likewise.
* c-c++-common/gomp/declare-mapper-3.c: Likewise.
* c-c++-common/gomp/depend-1.c: Likewise.
* c-c++-common/gomp/depend-2.c: Likewise.
* c-c++-common/gomp/depend-3.c: Likewise.
* c-c++-common/gomp/depend-4.c: Likewise.
* c-c++-common/gomp/depend-5.c: Likewise.
* c-c++-common/gomp/depend-6.c: Likewise.
* c-c++-common/gomp/dispatch-1.c: Likewise.
* c-c++-common/gomp/loop-5.c: Likewise.
* c-c++-common/gomp/map-1.c: Likewise.
* c-c++-common/gomp/map-2.c: Likewise.
* c-c++-common/gomp/map-4.c: Likewise.
* c-c++-common/gomp/map-7.c: Likewise.
* c-c++-common/gomp/pr100902-1.c: Likewise.
* c-c++-common/gomp/pr103642.c: Likewise.
* c-c++-common/gomp/pr120180-1.c: Likewise.
* c-c++-common/gomp/pr61486-1.c: Likewise.
* c-c++-common/gomp/pr81006.c: Likewise.
* c-c++-common/gomp/pr91920.c: Likewise.
* c-c++-common/gomp/pr96867.c: Likewise.
* c-c++-common/gomp/pr99928-16.c: Likewise.
* c-c++-common/gomp/reduction-1.c: Likewise.
* c-c++-common/gomp/scan-1.c: Likewise.
* c-c++-common/gomp/target-data-1.c: Likewise.
* c-c++-common/gomp/target-enter-data-1.c: Likewise.
* c-c++-common/gomp/target-has-device-addr-1.c: Likewise.
* c-c++-common/gomp/target-implicit-map-2.c: Likewise.
* c-c++-common/gomp/target-map-iterators-1.c: Likewise.
* c-c++-common/gomp/target-map-iterators-3.c: Likewise.
* c-c++-common/gomp/target-update-iterators-1.c: Likewise.
* c-c++-common/gomp/target-update-iterators-3.c: Likewise.
* g++.dg/goacc/cache-3-1.C: Likewise.
* g++.dg/goacc/cache-3-2.C: Likewise.
* g++.dg/goacc/data-clause-1.C: Likewise.
* g++.dg/goacc/mdc.C: Likewise.
* g++.dg/gomp/array-section-2.C: Likewise.
* g++.dg/gomp/bad-array-section-10.C: Likewise.
* g++.dg/gomp/bad-array-section-11.C: Likewise.
* g++.dg/gomp/bad-array-section-9.C: Likewise.
* g++.dg/gomp/declare-mapper-1.C: Likewise.
* g++.dg/gomp/declare-mapper-2.C: Likewise.
* g++.dg/gomp/depend-1.C: Likewise.
* g++.dg/gomp/depend-2.C: Likewise.
* g++.dg/gomp/ind-base-3.C: Likewise.
* g++.dg/gomp/map-1.C: Likewise.
* g++.dg/gomp/map-2.C: Likewise.
* g++.dg/gomp/map-ptrmem-1.C: Likewise.
* g++.dg/gomp/map-ptrmem-2.C: Likewise.
* g++.dg/gomp/member-array-2.C: Likewise.
* g++.dg/gomp/target-this-3.C: Likewise.
* g++.dg/gomp/target-this-4.C: Likewise.
libgomp/
* testsuite/libgomp.c++/allocate-1.C: Replace [: in OpenMP or OpenACC
pragmas or attributes with [ : and :] with : ].
* testsuite/libgomp.c++/baseptrs-3.C: Likewise.
* testsuite/libgomp.c++/baseptrs-5.C: Likewise.
* testsuite/libgomp.c++/class-array-1.C: Likewise.
* testsuite/libgomp.c++/examples-4/target_data-5.C: Likewise.
* testsuite/libgomp.c++/lvalue-tofrom-2.C: Likewise.
* testsuite/libgomp.c++/pr101544-1.C: Likewise.
* testsuite/libgomp.c++/pr108286.C: Likewise.
* testsuite/libgomp.c++/reduction-10.C: Likewise.
* testsuite/libgomp.c++/reduction-11.C: Likewise.
* testsuite/libgomp.c++/reduction-12.C: Likewise.
* testsuite/libgomp.c++/reduction-5.C: Likewise.
* testsuite/libgomp.c++/reduction-6.C: Likewise.
* testsuite/libgomp.c++/reduction-7.C: Likewise.
* testsuite/libgomp.c++/reduction-8.C: Likewise.
* testsuite/libgomp.c++/reduction-9.C: Likewise.
* testsuite/libgomp.c++/target-18.C: Likewise.
* testsuite/libgomp.c++/target-19.C: Likewise.
* testsuite/libgomp.c++/target-2.C: Likewise.
* testsuite/libgomp.c++/target-22.C: Likewise.
* testsuite/libgomp.c++/target-23.C: Likewise.
* testsuite/libgomp.c++/target-9.C: Likewise.
* testsuite/libgomp.c++/target-flex-100.C: Likewise.
* testsuite/libgomp.c++/target-flex-101.C: Likewise.
* testsuite/libgomp.c++/target-flex-12.C: Likewise.
* testsuite/libgomp.c++/target-flex-2003.C: Likewise.
* testsuite/libgomp.c++/target-flex-30.C: Likewise.
* testsuite/libgomp.c++/target-flex-300.C: Likewise.
* testsuite/libgomp.c++/target-flex-32.C: Likewise.
* testsuite/libgomp.c++/target-flex-33.C: Likewise.
* testsuite/libgomp.c++/target-flex-41.C: Likewise.
* testsuite/libgomp.c++/target-flex-60.C: Likewise.
* testsuite/libgomp.c++/target-flex-61.C: Likewise.
* testsuite/libgomp.c++/target-flex-62.C: Likewise.
* testsuite/libgomp.c++/target-flex-80.C: Likewise.
* testsuite/libgomp.c++/target-flex-81.C: Likewise.
* testsuite/libgomp.c++/target-has-device-addr-7.C: Likewise.
* testsuite/libgomp.c++/target-in-reduction-1.C: Likewise.
* testsuite/libgomp.c++/target-in-reduction-2.C: Likewise.
* testsuite/libgomp.c++/target-lambda-1.C: Likewise.
* testsuite/libgomp.c++/target-lambda-3.C: Likewise.
* testsuite/libgomp.c++/target-map-class-1.C: Likewise.
* testsuite/libgomp.c++/target-std__array-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__bitset-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__deque-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__flat_map-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__flat_multimap-concurrent.C:
Likewise.
* testsuite/libgomp.c++/target-std__flat_multiset-concurrent.C:
Likewise.
* testsuite/libgomp.c++/target-std__flat_set-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__forward_list-concurrent.C:
Likewise.
* testsuite/libgomp.c++/target-std__list-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__map-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__multimap-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__multiset-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__set-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__span-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__unordered_map-concurrent.C:
Likewise.
* testsuite/libgomp.c++/target-std__unordered_multimap-concurrent.C:
Likewise.
* testsuite/libgomp.c++/target-std__unordered_multiset-concurrent.C:
Likewise.
* testsuite/libgomp.c++/target-std__unordered_set-concurrent.C:
Likewise.
* testsuite/libgomp.c++/target-std__valarray-1.C: Likewise.
* testsuite/libgomp.c++/target-std__valarray-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__vector-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-this-3.C: Likewise.
* testsuite/libgomp.c++/target-this-4.C: Likewise.
* testsuite/libgomp.c++/target-virtual-1.C: Likewise.
* testsuite/libgomp.c++/task-reduction-11.C: Likewise.
* testsuite/libgomp.c++/task-reduction-12.C: Likewise.
* testsuite/libgomp.c++/task-reduction-13.C: Likewise.
* testsuite/libgomp.c++/task-reduction-17.C: Likewise.
* testsuite/libgomp.c++/task-reduction-18.C: Likewise.
* testsuite/libgomp.c++/task-reduction-19.C: Likewise.
* testsuite/libgomp.c++/task-reduction-4.C: Likewise.
* testsuite/libgomp.c++/task-reduction-5.C: Likewise.
* testsuite/libgomp.c++/task-reduction-6.C: Likewise.
* testsuite/libgomp.c++/task-reduction-7.C: Likewise.
* testsuite/libgomp.c++/taskloop-reduction-2.C: Likewise.
* testsuite/libgomp.c++/taskloop-reduction-3.C: Likewise.
* testsuite/libgomp.c++/taskloop-reduction-4.C: Likewise.
* testsuite/libgomp.c-c++-common/allocate-1.c: Likewise.
* testsuite/libgomp.c-c++-common/allocate-3.c: Likewise.
* testsuite/libgomp.c-c++-common/baseptrs-2.c: Likewise.
* testsuite/libgomp.c-c++-common/dispatch-1.c: Likewise.
* testsuite/libgomp.c-c++-common/dispatch-2.c: Likewise.
* testsuite/libgomp.c-c++-common/interop-2.c: Likewise.
* testsuite/libgomp.c-c++-common/matrix-omp-target-teams-distribute-parallel-for-1.c:
Likewise.
* testsuite/libgomp.c-c++-common/ptr-attach-1.c: Likewise.
* testsuite/libgomp.c-c++-common/ptr-attach-2.c: Likewise.
* testsuite/libgomp.c-c++-common/refcount-1.c: Likewise.
* testsuite/libgomp.c-c++-common/struct-elem-4.c: Likewise.
* testsuite/libgomp.c-c++-common/target-2.c: Likewise.
* testsuite/libgomp.c-c++-common/target-has-device-addr-1.c: Likewise.
* testsuite/libgomp.c-c++-common/target-implicit-map-2.c: Likewise.
* testsuite/libgomp.c-c++-common/target-implicit-map-5.c: Likewise.
* testsuite/libgomp.c-c++-common/target-in-reduction-1.c: Likewise.
* testsuite/libgomp.c-c++-common/target-in-reduction-2.c: Likewise.
* testsuite/libgomp.c-c++-common/target-map-iterators-1.c: Likewise.
* testsuite/libgomp.c-c++-common/target-map-iterators-2.c: Likewise.
* testsuite/libgomp.c-c++-common/target-map-iterators-3.c: Likewise.
* testsuite/libgomp.c-c++-common/target-map-zlas-1.c: Likewise.
* testsuite/libgomp.c-c++-common/target-update-iterators-1.c: Likewise.
* testsuite/libgomp.c-c++-common/target-update-iterators-2.c: Likewise.
* testsuite/libgomp.c-c++-common/target-update-iterators-3.c: Likewise.
* testsuite/libgomp.c-c++-common/task-reduction-11.c: Likewise.
* testsuite/libgomp.c-c++-common/task-reduction-12.c: Likewise.
* testsuite/libgomp.c-c++-common/task-reduction-16.c: Likewise.
* testsuite/libgomp.c-c++-common/task-reduction-3.c: Likewise.
* testsuite/libgomp.c-c++-common/task-reduction-7.c: Likewise.
* testsuite/libgomp.c-c++-common/task-reduction-9.c: Likewise.
* testsuite/libgomp.c-c++-common/taskloop-reduction-2.c: Likewise.
* testsuite/libgomp.c-c++-common/teams-nteams-icv-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-16.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: Likewise.
|
|
This adds support for using Cuda Managed Memory with omp_alloc. AMD support
will be added in a future patch.
There is one new predefined allocator, "ompx_gnu_managed_mem_alloc", plus a
corresponding memory space, which can be used to allocate memory in the
"managed" space.
The nvptx plugin is modified to make the necessary Cuda calls, via two new
(optional) plugin interfaces.
gcc/fortran/ChangeLog:
* openmp.cc (is_predefined_allocator): Use GOMP_OMP_PREDEF_ALLOC_MAX
and GOMP_OMPX_PREDEF_ALLOC_MIN/MAX instead of hardcoded values in the
comment.
include/ChangeLog:
* cuda/cuda.h (cuMemAllocManaged): Add declaration and related
CU_MEM_ATTACH_GLOBAL flag.
* gomp-constants.h (GOMP_OMPX_PREDEF_ALLOC_MAX): Update to 201.
(GOMP_OMP_PREDEF_MEMSPACE_MAX): New constant.
(GOMP_OMPX_PREDEF_MEMSPACE_MIN): New constant.
(GOMP_OMPX_PREDEF_MEMSPACE_MAX): New constant.
libgomp/ChangeLog:
* allocator.c (ompx_gnu_max_predefined_alloc): Update to
ompx_gnu_managed_mem_alloc.
(_Static_assert): Fix assertion messages for allocators and add
new assertions for memspace constants.
(omp_max_predefined_mem_space): New define.
(ompx_gnu_min_predefined_mem_space): New define.
(ompx_gnu_max_predefined_mem_space): New define.
(MEMSPACE_ALLOC): Add check for non-standard memspaces.
(MEMSPACE_CALLOC): Likewise.
(MEMSPACE_REALLOC): Likewise.
(MEMSPACE_VALIDATE): Likewise.
(predefined_ompx_gnu_alloc_mapping): Add ompx_gnu_managed_mem_space.
(omp_init_allocator): Add ompx_gnu_managed_mem_space validation.
* config/gcn/allocator.c (gcn_memspace_alloc): Add check for
non-standard memspaces.
(gcn_memspace_calloc): Likewise.
(gcn_memspace_realloc): Likewise.
(gcn_memspace_validate): Update to validate standard vs non-standard
memspaces.
* config/linux/allocator.c (linux_memspace_alloc): Add managed
memory space handling.
(linux_memspace_calloc): Likewise.
(linux_memspace_free): Likewise.
(linux_memspace_realloc): Likewise (returns NULL for fallback).
* config/nvptx/allocator.c (nvptx_memspace_alloc): Add check for
non-standard memspaces.
(nvptx_memspace_calloc): Likewise.
(nvptx_memspace_realloc): Likewise.
(nvptx_memspace_validate): Update to validate standard vs non-standard
memspaces.
* env.c (parse_allocator): Add ompx_gnu_managed_mem_alloc,
ompx_gnu_managed_mem_space, and some static asserts so I don't forget
them again.
* libgomp-plugin.h (GOMP_OFFLOAD_managed_alloc): New declaration.
(GOMP_OFFLOAD_managed_free): New declaration.
* libgomp.h (gomp_managed_alloc): New declaration.
(gomp_managed_free): New declaration.
(struct gomp_device_descr): Add managed_alloc_func and
managed_free_func fields.
* libgomp.texi: Document ompx_gnu_managed_mem_alloc and
ompx_gnu_managed_mem_space, add C++ template documentation, and
describe NVPTX and AMD support.
* omp.h.in: Add ompx_gnu_managed_mem_space and
ompx_gnu_managed_mem_alloc enumerators, and gnu_managed_mem C++
allocator template.
* omp_lib.f90.in: Add Fortran bindings for new allocator and
memory space.
* omp_lib.h.in: Likewise.
* plugin/cuda-lib.def: Add cuMemAllocManaged.
* plugin/plugin-nvptx.c (nvptx_alloc): Add managed parameter to
support cuMemAllocManaged.
(GOMP_OFFLOAD_alloc): Move contents to ...
(cleanup_and_alloc): ... this new function, and add managed support.
(GOMP_OFFLOAD_managed_alloc): New function.
(GOMP_OFFLOAD_managed_free): New function.
* target.c (gomp_managed_alloc): New function.
(gomp_managed_free): New function.
(gomp_load_plugin_for_device): Load optional managed_alloc
and managed_free plugin APIs.
* testsuite/lib/libgomp.exp: Add check_effective_target_omp_managedmem.
* testsuite/libgomp.c++/alloc-managed-1.C: New test.
* testsuite/libgomp.c/alloc-managed-1.c: New test.
* testsuite/libgomp.c/alloc-managed-2.c: New test.
* testsuite/libgomp.c/alloc-managed-3.c: New test.
* testsuite/libgomp.c/alloc-managed-4.c: New test.
* testsuite/libgomp.fortran/alloc-managed-1.f90: New test.
Co-authored-by: Kwok Cheung Yeung <kcyeung@baylibre.com>
Co-authored-by: Thomas Schwinge <tschwinge@baylibre.com>
|
|
omp_default_device [P119677]
Commit r16-5188-g5da963d988e8ea added omp_default_device such that -5
became a conforming device number, but the tests used them to test
for as non-conforming number; now -6 is used.
libgomp/ChangeLog:
PR libgomp/119677
* testsuite/libgomp.c-c++-common/target-is-accessible-1.c: Modify
test as -5 is now a conforming device number.
* testsuite/libgomp.fortran/target-is-accessible-1.f90: Likewise.
|
|
OpenMP TR 14 (OpenMP 6.1) adds omp_default_device < -1 as
named constant alongside omp_initial_device and omp_default_device.
GCC supports it already internally via GOMP_DEVICE_DEFAULT_OMP_61,
but this patch now adds the omp_default_device enum/PARAMETER to
omp.h / omp_lib.
Note that PR119677 requests some cleanups, which still have to be
done.
PR libgomp/119677
gcc/fortran/ChangeLog:
* intrinsic.texi (OpenMP Modules): Add omp_default_device.
* openmp.cc (gfc_resolve_omp_context_selector): Accept
omp_default_device as conforming device number.
libgomp/ChangeLog:
* omp.h.in (omp_default_device): New enum value.
* omp_lib.f90.in: New parameter.
* omp_lib.h.in: Likewise
* target.c (gomp_get_default_device): New. Split off from ...
(resolve_device): ... here; call it.
(omp_target_alloc, omp_target_free, omp_target_is_present,
omp_target_memcpy_check, omp_target_memset, omp_target_memset_async,
omp_target_associate_ptr, omp_get_mapped_ptr,
omp_target_is_accessible, omp_pause_resource,
omp_get_uid_from_device): Handle omp_default_device.
* testsuite/libgomp.c/device_uid.c: Likewise.
* testsuite/libgomp.fortran/device_uid.f90: Likewise.
* testsuite/libgomp.c-c++-common/omp-default-device.c: New test.
* testsuite/libgomp.fortran/omp-default-device.f90: New test.
|
|
Fix the unified-shared memory test,
libgomp.c++/target-std__multimap-concurrent-usm.C
added in commit r16-1010-g83ca283853f195
libgomp: Add testcases for concurrent access to standard C++ containers
on offload targets, a number of USM variants
This tests includes the actual code of target-std__multimap-concurrent.C.
The issue is that multimap.insert allocates memory – which is freed by
the destructor. However, if the memory is allocated on a device
('insert'), it also needs to be freed there ('clear') as in general
freeing device-allocated memory is not possible on the host.
libgomp/ChangeLog:
* testsuite/libgomp.c++/target-std__multimap-concurrent.C: Fix memory
freeing of device allocated memory with USM.
|
|
size [PR122543]
In OpenMP, pointers are implicitly mapped - which means for Fortran that
their pointer target is also mapped. However, for uninitialized memory,
this means that some random pointee with some random amount of memory is
copied - in the good case, size == 0, but if not, odd things can happen.
Solution: Use 'fptr => null()' before the target mapping or - as done here -
declare the pointer inside the region.
libgomp/ChangeLog:
PR libgomp/122543
* testsuite/libgomp.fortran/omp_target_memset.f90: Move fptr inside
the target to avoid implicit mapping of its uninit pointee.
* testsuite/libgomp.fortran/omp_target_memset-2.f90: Likewise.
|
|
... where it currently runs into:
libgomp: cuCtxSynchronize error: an illegal memory access was encountered
... for nvptx, or similarly for GCN:
Memory access fault by GPU node-1 (Agent handle: 0x34d77290) on address 0x7fff3c553000. Reason: Page not present or supervisor privilege.
Fix-up for commit r16-4961-ge2cbcd1b27c0da92bdcd96664064d3d0c1d44e6f
"Fix gimple_copy for OpenMP atomic load/store [PR122281, PR105001]".
PR libgomp/122281
libgomp/
* testsuite/libgomp.c/pr122281.c: Fix for non-USM offloading
execution.
|
|
PR libgomp/122281
PR middle-end/105001
gcc/ChangeLog:
* gimple.cc (gimple_copy): Add missing unshare_expr for
GIMPLE_OMP_ATOMIC_LOAD and GIMPLE_OMP_ATOMIC_STORE.
|
|
This patch introduces a new custom memory allocator for use with pinned
memory (in the case where the Cuda allocator isn't available). In future,
this allocator will also be used for Managed Memory. Both memories are
incompatible with the system malloc because allocated memory cannot share a
page with memory allocated for other purposes.
This means that small allocations will no longer consume an entire page of
pinned memory. Unfortunately, it also means that pinned memory pages will
never be unmapped (although they may be reused). This isn't a technical
limitation; the "free" algorithm could be extended in future, if needed.
The implementation is not perfect; there are various corner cases (especially
related to extending onto new pages) where allocations and reallocations may
be sub-optimal, but it should still be a step forward in support for small
allocations.
I have considered using libmemkind's "fixed" memory but rejected it for three
reasons: 1) libmemkind may not always be present at runtime, 2) there's no
currently documented means to extend a "fixed" kind one page at a time
(although the code appears to have an undocumented function that may do the
job, and/or extending libmemkind to support the MAP_LOCKED mmap flag with its
regular kinds would be straight-forward), 3) Managed Memory benefits from
having the metadata located in different memory and using an external
implementation makes it hard to guarantee this.
libgomp/ChangeLog:
* Makefile.am (libgomp_la_SOURCES): Add simple-allocator.c.
* Makefile.in: Regenerate.
* basic-allocator.c: Mention simple-allocator in the comment.
* config/linux/allocator.c: Include unistd.h.
(pin_ctx): New variable.
(ctxlock): New variable.
(linux_init_pin_ctx): New function.
(linux_memspace_alloc): Use simple-allocator for pinned memory.
(linux_memspace_free): Likewise.
(linux_memspace_realloc): Likewise.
* libgomp.h (gomp_simple_alloc_init_context): New prototype.
(gomp_simple_alloc_register_memory): New prototype.
(gomp_simple_alloc): New prototype.
(gomp_simple_free): New prototype.
(gomp_simple_realloc): New prototype.
* libgomp.texi: Update pinned memory trait documentation.
* testsuite/libgomp.c/alloc-pinned-8.c: New test.
* simple-allocator.c: New file.
|
|
Use Cuda to pin memory, instead of Linux mlock, when available.
There are two advantages: firstly, this gives a significant speed boost for
NVPTX offloading, and secondly, it side-steps the usual OS ulimit/rlimit
setting.
The design adds a device independent plugin API for allocating pinned memory,
and then implements it for NVPTX. At present, the other supported devices do
not have equivalent capabilities (or requirements).
libgomp/ChangeLog:
* config/linux/allocator.c: Include assert.h.
(using_device_for_page_locked): New variable.
(linux_memspace_alloc): Add init0 parameter. Support device pinning.
(linux_memspace_calloc): Set init0 to true.
(linux_memspace_free): Support device pinning.
(linux_memspace_realloc): Support device pinning.
(MEMSPACE_ALLOC): Set init0 to false.
* libgomp-plugin.h
(GOMP_OFFLOAD_page_locked_host_alloc): New prototype.
(GOMP_OFFLOAD_page_locked_host_free): Likewise.
* libgomp.h (gomp_page_locked_host_alloc): Likewise.
(gomp_page_locked_host_free): Likewise.
(struct gomp_device_descr): Add page_locked_host_alloc_func and
page_locked_host_free_func.
* libgomp.texi: Adjust the docs for the pinned trait.
* plugin/plugin-nvptx.c
(GOMP_OFFLOAD_page_locked_host_alloc): New function.
(GOMP_OFFLOAD_page_locked_host_free): Likewise.
* target.c (device_for_page_locked): New variable.
(get_device_for_page_locked): New function.
(gomp_page_locked_host_alloc): Likewise.
(gomp_page_locked_host_free): Likewise.
(gomp_load_plugin_for_device): Add page_locked_host_alloc and
page_locked_host_free.
* testsuite/libgomp.c/alloc-pinned-1.c: Change expectations for NVPTX
devices.
* testsuite/libgomp.c/alloc-pinned-2.c: Likewise.
* testsuite/libgomp.c/alloc-pinned-3.c: Likewise.
* testsuite/libgomp.c/alloc-pinned-4.c: Likewise.
* testsuite/libgomp.c/alloc-pinned-5.c: Likewise.
* testsuite/libgomp.c/alloc-pinned-6.c: Likewise.
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
|
|
This patch addresses the incorrectly placed tests, which fail if the
testsuite is ran and gcc has not been installed yet, as discussed
here:
https://gcc.gnu.org/pipermail/gcc-patches/2025-October/698095.html.
gcc/testsuite/ChangeLog:
* gcc.dg/ipa/ipcp-cb-spec1.c: Moved to libgomp/testsuite/libgomp.c/.
* gcc.dg/ipa/ipcp-cb-spec2.c: Likewise.
* gcc.dg/ipa/ipcp-cb1.c: Likewise.
libgomp/ChangeLog:
* testsuite/libgomp.c/ipcp-cb-spec1.c: Moved from
gcc/testsuite/gcc.dg/ipa/.
* testsuite/libgomp.c/ipcp-cb-spec2.c: Likewise.
* testsuite/libgomp.c/ipcp-cb1.c: Likewise.
Signed-off-by: Josef Melcr <jmelcr02@gmail.com>
|
|
uninitialized reads: Adjust 'libgomp.c++/{target-flex-101.C,target-std__flat_map-concurrent.C,target-std__flat_multimap-concurrent.C}' [PR114457, PR122268, PR120450]
With commit r16-4212-gf256a13f8aed833fe964a2ba541b7b30ad9b4a76
"c++, gimplify: Implement C++26 P2795R5 - Erroneous behavior for uninitialized reads [PR114457]",
we acquired:
{+FAIL: libgomp.c++/target-flex-101.C (internal compiler error: in assign_temp, at function.cc:990)+}
[-PASS:-]{+FAIL:+} libgomp.c++/target-flex-101.C (test for excess errors)
[-PASS:-]{+UNRESOLVED:+} libgomp.c++/target-flex-101.C [-execution test-]{+compilation failed to produce executable+}
... for GCN, nvptx offloading compilation, and on the other hand:
[-XFAIL:-]{+XPASS:+} libgomp.c++/target-std__flat_map-concurrent.C (internal compiler error[-: in assign_temp, at function.cc:990)-]
[-XFAIL:-]{+XPASS:+} libgomp.c++/target-std__flat_map-concurrent.C (test for excess errors)
[-UNRESOLVED:-]{+PASS:+} libgomp.c++/target-std__flat_map-concurrent.C [-compilation failed to produce executable-]{+execution test+}
[-XFAIL:-]{+XPASS:+} libgomp.c++/target-std__flat_multimap-concurrent.C (internal compiler error[-: in assign_temp, at function.cc:990)-]
[-XFAIL:-]{+XPASS:+} libgomp.c++/target-std__flat_multimap-concurrent.C (test for excess errors)
[-UNRESOLVED:-]{+PASS:+} libgomp.c++/target-std__flat_multimap-concurrent.C [-compilation failed to produce executable-]{+execution test+}
... for GCN offloading compilation (already PASSed for nvptx).
Note that these test cases explicitly use '-std=c++23', so don't undergo the
new C++26 P2795R5 functionality. Yet, comparing before vs. after that commit,
in the 'gimple' dumps (that is, early host compilation), there are a lot of
changes where 'gimple_assign <constructor, [...], {CLOBBER(bob)}, NULL, NULL>'s
and relatedly 'gimple_bind's newly appear/no longer appear elsewhere. This
leads to correspondingly different code at the beginning of offloading
compilation. Why/how that now ('libgomp.c++/target-flex-101.C') vs. before
('libgomp.c++/{target-std__flat_map-concurrent.C,target-std__flat_multimap-concurrent.C}')
translates into 'expand' ICEs, I can't tell.
PR c++/114457
PR c++/122268
PR c++/120450
libgomp/
* testsuite/libgomp.c++/target-flex-101.C: XFAIL GCN, nvptx
offloading compilation.
* testsuite/libgomp.c++/target-std__flat_map-concurrent.C:
Un-XFAIL GCN offloading compilation.
* testsuite/libgomp.c++/target-std__flat_multimap-concurrent.C:
Likewise.
|
|
Add missing tests for gfx* context selectors; mark all but the
default-arch declare-variant-4.c with 'dg-excess-errors' to
silence libgomp not-found errors (still passing the
scan-offload-tree-dump check) - or at least causing just
UNRESOLVED errors if the error is
"built without library support ... consider compiling for
the associated generic architecture".
In case the multilib is configured, the result will be
an XPASS.
libgomp/ChangeLog:
* testsuite/libgomp.c/declare-variant-4-gfx10-3-generic.c: Add
dg-excess-errors to handle possible missing libgomp multi lib.
* testsuite/libgomp.c/declare-variant-4-gfx1030.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx1036.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx11-generic.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx1100.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx1103.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx9-4-generic.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx9-generic.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx900.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx906.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx908.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx90a.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx90c.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx942.c: Likewise.
* testsuite/libgomp.c/declare-variant-4-gfx1031.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1032.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1033.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1034.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1035.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1101.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1102.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1150.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1151.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1152.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx1153.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx902.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx904.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx909.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx950.c: New test.
|
|
The definition for gfx942 and gfx950 missed the GFX9_4_GENERIC
family flag.
For OpenMP context selectors: The t-omp-device file missed the
generic selectors.
Additionally, there is now a note in the OpenMP documentation that
there is a one-to-one match for ISA names, ignoring any compatibility.
For instance, for Nvidia GPUs 'isa("sm_70")' is only true when compiling
for 'sm_70', even though sm < 7.0 code also runs on sm_70 hardware.
And, for AMD GPUs, gfx9-4-generic neither matches 'gfx942'
(even though such generic code runs on gfx942) - nor the reverse
(although all gfx9-4-generic code runs on gfx942).
gcc/ChangeLog:
* config/gcn/gcn-devices.def (gfx942, gfx950): Set generic name
to GFX9_4_GENERIC.
* config/gcn/t-omp-device: Include generic names for OpenMP's
ISA trait.
libgomp/ChangeLog:
* libgomp.texi (OpenMP Context Selectors): Add note that there is
currently an exact match between ISA and compilation, ignoring
compatibilities in both ways.
* testsuite/libgomp.c/declare-variant-4.h: Add missing variant
functions for specific and generic AMD GPUs.
* testsuite/libgomp.c/declare-variant-4-gfx10-3-generic.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx11-generic.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx9-4-generic.c: New test.
* testsuite/libgomp.c/declare-variant-4-gfx9-generic.c: New test.
|
|
libatomic/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
* testsuite/Makefile.in: Regenerate.
libcc1/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
libffi/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
* include/Makefile.in: Regenerate.
* man/Makefile.in: Regenerate.
* testsuite/Makefile.in: Regenerate.
libgcobol/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
libgfortran/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
libgm2/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
* libm2cor/Makefile.in: Regenerate.
* libm2iso/Makefile.in: Regenerate.
* libm2log/Makefile.in: Regenerate.
* libm2min/Makefile.in: Regenerate.
* libm2pim/Makefile.in: Regenerate.
libgomp/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
* testsuite/Makefile.in: Regenerate.
libgrust/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
* libformat_parser/Makefile.in: Regenerate.
* libproc_macro_internal/Makefile.in: Regenerate.
libitm/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
* testsuite/Makefile.in: Regenerate.
libobjc/ChangeLog:
* aclocal.m4: Regenerate.
* configure: Regenerate.
libphobos/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
* libdruntime/Makefile.in: Regenerate.
* src/Makefile.in: Regenerate.
* testsuite/Makefile.in: Regenerate.
libquadmath/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
libsanitizer/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* asan/Makefile.in: Regenerate.
* configure: Regenerate.
* hwasan/Makefile.in: Regenerate.
* interception/Makefile.in: Regenerate.
* libbacktrace/Makefile.in: Regenerate.
* lsan/Makefile.in: Regenerate.
* sanitizer_common/Makefile.in: Regenerate.
* tsan/Makefile.in: Regenerate.
* ubsan/Makefile.in: Regenerate.
libssp/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
libstdc++-v3/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
* doc/Makefile.in: Regenerate.
* include/Makefile.in: Regenerate.
* libsupc++/Makefile.in: Regenerate.
* po/Makefile.in: Regenerate.
* python/Makefile.in: Regenerate.
* src/Makefile.in: Regenerate.
* src/c++11/Makefile.in: Regenerate.
* src/c++17/Makefile.in: Regenerate.
* src/c++20/Makefile.in: Regenerate.
* src/c++23/Makefile.in: Regenerate.
* src/c++26/Makefile.in: Regenerate.
* src/c++98/Makefile.in: Regenerate.
* src/experimental/Makefile.in: Regenerate.
* src/filesystem/Makefile.in: Regenerate.
* src/libbacktrace/Makefile.in: Regenerate.
* testsuite/Makefile.in: Regenerate.
libvtv/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
* testsuite/Makefile.in: Regenerate.
lto-plugin/ChangeLog:
* Makefile.in: Regenerate.
* aclocal.m4: Regenerate.
* configure: Regenerate.
|
|
As the testcase shows, a missing unshare_expr caused that the condition
was only evaluated once instead of every time when a 'declare variant'
was resolved.
PR middle-end/121922
gcc/ChangeLog:
* omp-general.cc (omp_dynamic_cond): Use 'unshare_expr' for
the user condition.
libgomp/ChangeLog:
* testsuite/libgomp.c-c++-common/declare-variant-1.c: New test.
Co-authored-by: Sandra Loosemore <sloosemore@baylibre.com>
|
|
acc_delete_finalize_async
OpenACC 2.5 added several functions for C and Fortran; while
acc_{copyout,delete}{,_finalize,_async} exist for both, for some
reasons only the C version of acc_{copyout,delete}_finalize_async
was actually added, even though the documentation (.texi) and
the .map file listed also the auxiliar Fortran functions!
OpenACC 2.5 added the Fortran version with the following odd
interface: 'type, dimension(:[,:]...)'. In OpenACC 2.6, it
was then updated to the Fortran 2018 syntax:
'type(*), dimension(..)', which is also used in openacc.f90
internally.
This commit now also updates the documentation to the newer
syntax - plus fixes a function-name typo: acc_delete_async_finalize
should have the _async at the end not in the middle!
libgomp/ChangeLog:
* libgomp.map (OACC_2.5): Move previously unimplemented
acc_{copyout,delete}_finalize_async_{32,64,array}_h_ to ...
(OACC_2.6.1): ... here.
* libgomp.texi (acc_copyin, acc_present_or_copyin, acc_create,
acc_present_or_create, acc_copyout, acc_update_device,
acc_update_self, acc_is_present): Use 'type(*), dimension(..)'
instead of 'type, dimension(:[,:]...)' for Fortran.
(acc_delete): Likewise; change acc_delete_async_finalize to
acc_delete_finalize_async.
* openacc.f90 (openacc_internal): Add interfaces for
acc_{copyout,delete}_finalize_async_{{32,64,array}_h,_l}.
(openacc): Add generic interfaces for
acc_copyout_finalize_async and acc_delete_finalize_async.
(acc_{copyout,delete}_finalize_async_{32,64,array}_h): New.
* openacc_lib.h: Add generic interfaces for
acc_copyout_finalize_async and acc_delete_finalize_async.
* testsuite/libgomp.oacc-fortran/pr92970-1.f90: New test.
|
|
This adds support for iterators in 'to' and 'from' clauses in the
'target update' OpenMP directive.
gcc/c/
* c-parser.cc (c_parser_omp_clause_from_to): Parse 'iterator' modifier.
* c-typeck.cc (c_finish_omp_clauses): Finish iterators for to/from
clauses.
gcc/cp/
* parser.cc (cp_parser_omp_clause_from_to): Parse 'iterator' modifier.
* semantics.cc (finish_omp_clauses): Finish iterators for to/from
clauses.
gcc/
* gimplify.cc (remove_unused_omp_iterator_vars): Display unused
variable warning for 'to' and 'from' clauses.
(gimplify_scan_omp_clauses): Add argument for iterator loop sequence.
Gimplify the clause decl and size into the iterator loop if iterators
are used.
(gimplify_omp_workshare): Add argument for iterator loops sequence
in call to gimplify_scan_omp_clauses.
(gimplify_omp_target_update): Call remove_unused_omp_iterator_vars and
build_omp_iterators_loops. Add loop sequence as argument when calling
gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses and building
the Gimple statement.
* tree-pretty-print.cc (dump_omp_clause): Call dump_omp_iterators
for to/from clauses with iterators.
* tree.cc (omp_clause_num_ops): Add extra operand for OMP_CLAUSE_FROM
and OMP_CLAUSE_TO.
* tree.h (OMP_CLAUSE_HAS_ITERATORS): Add check for OMP_CLAUSE_TO and
OMP_CLAUSE_FROM.
(OMP_CLAUSE_ITERATORS): Likewise.
gcc/testsuite/
* c-c++-common/gomp/target-update-iterators-1.c: New.
* c-c++-common/gomp/target-update-iterators-2.c: New.
* c-c++-common/gomp/target-update-iterators-3.c: New.
libgomp/
* target.c (gomp_update): Call gomp_merge_iterator_maps. Free
allocated variables.
* testsuite/libgomp.c-c++-common/target-update-iterators-1.c: New.
* testsuite/libgomp.c-c++-common/target-update-iterators-2.c: New.
* testsuite/libgomp.c-c++-common/target-update-iterators-3.c: New.
|
|
This adds preliminary support for iterators in map clauses within OpenMP
'target' constructs (which includes constructs such as 'target enter data').
Iterators with non-constant loop bounds are not currently supported.
gcc/c/
* c-parser.cc (c_parser_omp_variable_list): Use location of the
map expression as the clause location.
(c_parser_omp_clause_map): Parse 'iterator' modifier.
* c-typeck.cc (c_finish_omp_clauses): Finish iterators. Apply
iterators to generated clauses.
gcc/cp/
* parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier.
* semantics.cc (finish_omp_clauses): Finish iterators. Apply
iterators to generated clauses.
gcc/
* gimple-pretty-print.cc (dump_gimple_omp_target): Print expanded
iterator loops.
* gimple.cc (gimple_build_omp_target): Add argument for iterator
loops sequence. Initialize iterator loops field.
* gimple.def (GIMPLE_OMP_TARGET): Set GSS symbol to GSS_OMP_TARGET.
* gimple.h (gomp_target): Set GSS symbol to GSS_OMP_TARGET. Add extra
field for iterator loops.
(gimple_build_omp_target): Add argument for iterator loops sequence.
(gimple_omp_target_iterator_loops): New.
(gimple_omp_target_iterator_loops_ptr): New.
(gimple_omp_target_set_iterator_loops): New.
* gimplify.cc (find_var_decl): New.
(copy_omp_iterator): New.
(remap_omp_iterator_var_1): New.
(remap_omp_iterator_var): New.
(remove_unused_omp_iterator_vars): New.
(struct iterator_loop_info_t): New type.
(iterator_loop_info_map_t): New type.
(build_omp_iterators_loops): New.
(enter_omp_iterator_loop_context_1): New.
(enter_omp_iterator_loop_context): New.
(enter_omp_iterator_loop_context): New.
(exit_omp_iterator_loop_context): New.
(gimplify_adjust_omp_clauses): Add argument for iterator loop
sequence. Gimplify the clause decl and size into the iterator
loop if iterators are used.
(gimplify_omp_workshare): Call remove_unused_omp_iterator_vars and
build_omp_iterators_loops for OpenMP target expressions. Add
loop sequence as argument when calling gimplify_adjust_omp_clauses
and building the Gimple statement.
* gimplify.h (enter_omp_iterator_loop_context): New prototype.
(exit_omp_iterator_loop_context): New prototype.
* gsstruct.def (GSS_OMP_TARGET): New.
* omp-low.cc (lower_omp_map_iterator_expr): New.
(lower_omp_map_iterator_size): New.
(finish_omp_map_iterators): New.
(lower_omp_target): Add sorry if iterators used with deep mapping.
Call lower_omp_map_iterator_expr before assigning to sender ref.
Call lower_omp_map_iterator_size before setting the size. Insert
iterator loop sequence before the statements for the target clause.
* tree-nested.cc (convert_nonlocal_reference_stmt): Walk the iterator
loop sequence of OpenMP target statements.
(convert_local_reference_stmt): Likewise.
(convert_tramp_reference_stmt): Likewise.
* tree-pretty-print.cc (dump_omp_iterators): Dump extra iterator
information if present.
(dump_omp_clause): Call dump_omp_iterators for iterators in map
clauses.
* tree.cc (omp_clause_num_ops): Add operand for OMP_CLAUSE_MAP.
(walk_tree_1): Do not walk last operand of OMP_CLAUSE_MAP.
* tree.h (OMP_CLAUSE_HAS_ITERATORS): New.
(OMP_CLAUSE_ITERATORS): New.
gcc/testsuite/
* c-c++-common/gomp/map-6.c (foo): Amend expected error message.
* c-c++-common/gomp/target-map-iterators-1.c: New.
* c-c++-common/gomp/target-map-iterators-2.c: New.
* c-c++-common/gomp/target-map-iterators-3.c: New.
* c-c++-common/gomp/target-map-iterators-4.c: New.
libgomp/
* target.c (kind_to_name): New.
(gomp_merge_iterator_maps): New.
(gomp_map_vars_internal): Call gomp_merge_iterator_maps. Copy
address of only the first iteration to target vars. Free allocated
variables.
* testsuite/libgomp.c-c++-common/target-map-iterators-1.c: New.
* testsuite/libgomp.c-c++-common/target-map-iterators-2.c: New.
* testsuite/libgomp.c-c++-common/target-map-iterators-3.c: New.
Co-authored-by: Andrew Stubbs <ams@baylibre.com>
|
|
[PR119853, PR119854]
Fix-up for commit aafe942227baf8c2bcd4cac2cb150e49a4b895a9
"GCN, nvptx offloading: Host/device compatibility: Itanium C++ ABI, DSO Object Destruction API [PR119853, PR119854]":
we need to adjust for 'targetm.cxx.use_aeabi_atexit':
gcc/config/arm/arm.cc:#define TARGET_CXX_USE_AEABI_ATEXIT arm_cxx_use_aeabi_atexit
gcc/config/arm/arm.cc:/* The EABI says __aeabi_atexit should be used to register static
gcc/config/arm/arm.cc- destructors. */
gcc/config/arm/arm.cc-
gcc/config/arm/arm.cc-static bool
gcc/config/arm/arm.cc:arm_cxx_use_aeabi_atexit (void)
gcc/config/arm/arm.cc-{
gcc/config/arm/arm.cc- return TARGET_AAPCS_BASED;
gcc/config/arm/arm.cc-}
..., which 'gcc/cp/decl.cc:get_atexit_node' then acts on: call '__aeabi_atexit'
instead of '__cxa_atexit', and swap two arguments.
PR target/119853
PR target/119854
libgomp/
* testsuite/libgomp.c++/target-cdtor-1.C: Adjust for
'targetm.cxx.use_aeabi_atexit'.
* testsuite/libgomp.c++/target-cdtor-2.C: Likewise.
|
|
'dynamic_cast' only for effective-target 'offload_device' [PR119692]
In PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' offloading":
> --- Comment #8 from Rainer Orth <ro at gcc dot gnu.org> ---
> The last commit made things worse on sparc-sun-solaris2.11: since that one
> (dg-timeout 10) I regularly get
>
> WARNING: libgomp.c++/target-exceptions-bad_cast-1.C (test for excess errors)
> program timed out.
> FAIL: libgomp.c++/target-exceptions-bad_cast-1.C (test for excess errors)
> UNRESOLVED: libgomp.c++/target-exceptions-bad_cast-1.C compilation failed to produce executable
> UNRESOLVED: libgomp.c++/target-exceptions-bad_cast-1.C scan-tree-dump-times optimized "gimple_call <__cxa_bad_cast, " 1
>
> Before that, the test had no issue. Compiling the test on an unloaded system
> usually takes less than 1 sec, but when fully loaded, times can go up.
To keep things simple, let's restrict this temporary (yeah...) workaround to
apply only for effective-target 'offload_device', just like the
'dg-xfail-run-if' itself.
PR target/119692
libgomp/
* testsuite/libgomp.c++/pr119692-1-4.C: '{ dg-timeout 10 { target offload_device } }'.
* testsuite/libgomp.c++/pr119692-1-5.C: Likewise.
* testsuite/libgomp.c++/target-exceptions-bad_cast-1.C: Likewise.
* testsuite/libgomp.c++/target-exceptions-bad_cast-2.C: Likewise.
* testsuite/libgomp.oacc-c++/exceptions-bad_cast-1.C: Likewise.
* testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C: Likewise.
|
|
While C/++ support the routines acc_attach{,_async} and
acc_detach{,_finalize}{,_async} routines since a long time, the Fortran
API routines where only added in OpenACC 3.3.
Unfortunately, they cannot directly be implemented in the library as
GCC will introduce a temporary array descriptor in some cases, which
causes the attempted attachment to the this temporary variable instead
of to the original one.
Therefore, those API routines are handled in a special way in the compiler.
gcc/fortran/ChangeLog:
* trans-stmt.cc (gfc_trans_call_acc_attach_detach): New.
(gfc_trans_call): Call it.
libgomp/ChangeLog:
* libgomp.texi (acc_attach, acc_detach): Update for Fortran
version.
* openacc.f90 (acc_attach{,_async}, acc_detach{,_finalize}{,_async}):
Add.
* openacc_lib.h: Likewise.
* testsuite/libgomp.oacc-fortran/acc-attach-detach-1.f90: New test.
* testsuite/libgomp.oacc-fortran/acc-attach-detach-2.f90: New test.
|
|
libstdc++-v3/include/std/ostream contains:
namespace std _GLIBCXX_VISIBILITY(default)
{
...
template<typename _CharT, typename _Traits>
inline basic_ostream<_CharT, _Traits>&
endl(basic_ostream<_CharT, _Traits>& __os)
{ return flush(__os.put(__os.widen('\n'))); }
...
#include <bits/ostream.tcc>
and the latter, libstdc++-v3/include/bits/ostream.tcc, has:
// Inhibit implicit instantiations for required instantiations,
// which are defined via explicit instantiations elsewhere.
#if _GLIBCXX_EXTERN_TEMPLATE
extern template class basic_ostream<char>;
extern template ostream& endl(ostream&);
Before this commit, omp_discover_declare_target_tgt_fn_r marked 'endl'
as (implicitly) declare target - but not the calls in it due to the
'extern' (DECL_EXTERNAL).
Thanks to inlining and as 'endl' is (therefore) not used and, hence,
discarded by the linker; hencet, it works with -O0 and -O1. However,
as the (unused) function still exits, IPA CP (enabled by -O2) will try
to do constant-value propagation and fails as the definition of 'widen'
is not available.
Solution is to still walk 'endl' despite being an 'extern(al)' decl;
this has been restricted for now to DECL_DECLARED_INLINE_P.
gcc/ChangeLog:
* omp-offload.cc (omp_discover_declare_target_tgt_fn_r): Also
walk external functions that are declare inline (and have a
DECL_SAVED_TREE).
libgomp/ChangeLog:
* testsuite/libgomp.c++/declare_target-2.C: New test.
|
|
As gfx942 and gfx950 belong to gfx9-4-generic, the latter two are also added.
Note that there are no specific optimizations for MI300, yet.
For none of the mentioned devices, any multilib is build by default; use
'--with-multilib-list=' when configuring GCC to build them alongside.
gfx942 was added in LLVM (and its mc assembler, used by GCC) in version 18,
generic support in LLVM 19 and gfx950 in LLVM 20.
gcc/ChangeLog:
* config/gcn/gcn-devices.def: Add gfx942, gfx950 and gfx9-4-generic.
* config/gcn/gcn-opts.h (TARGET_CDNA3, TARGET_CDNA3_PLUS,
TARGET_GLC_NAME, TARGET_TARGET_SC_CACHE): Define.
(TARGET_ARCHITECTED_FLAT_SCRATCH): Use also for CDNA3.
* config/gcn/gcn.h (gcn_isa): Add ISA_CDNA3 to the enum.
* config/gcn/gcn.cc (print_operand): Update 'g' to use
TARGET_GLC_NAME; add 'G' to print TARGET_GLC_NAME unconditionally.
* config/gcn/gcn-valu.md (scatter, gather): Use TARGET_GLC_NAME.
* config/gcn/gcn.md: Use %G<num> instead of glc; use 'buffer_inv sc1'
for TARGET_TARGET_SC_CACHE.
* doc/invoke.texi (march): Add gfx942, gfx950 and gfx9-4-generic.
* doc/install.texi (amdgcn*-*-*): Add gfx942, gfx950 and gfx9-4-generic.
* config/gcn/gcn-tables.opt: Regenerate.
libgomp/ChangeLog:
* testsuite/libgomp.c/declare-variant-4.h (gfx942): New variant function.
* testsuite/libgomp.c/declare-variant-4-gfx942.c: New test.
|
|
A mapping clause was missing, causing the code to fail with offloading
when a host pointer was not device accessible.
libgomp/ChangeLog:
PR target/120530
* testsuite/libgomp.c/target-map-zero-sized-3.c (main): Add missing
map clause; remove unused variable.
|
|
The test PASSes for C, but FAILs for C++:
.../libgomp.c-c++-common/omp_target_memset-3.c: In function 'void test_it(void*, int, size_t)':
.../libgomp.c-c++-common/omp_target_memset-3.c:31:7: warning: pointer of type 'void *' used in arithmetic [-Wpointer-arith]
.../libgomp.c-c++-common/omp_target_memset-3.c:33:13: error: invalid conversion from 'void*' to 'int8_t*' {aka 'signed char*'} [-fpermissive]
.../libgomp.c-c++-common/omp_target_memset-3.c:10:19: note: initializing argument 1 of 'void init_val(int8_t*, int, size_t)'
.../libgomp.c-c++-common/omp_target_memset-3.c:37:14: error: invalid conversion from 'void*' to 'int8_t*' {aka 'signed char*'} [-fpermissive]
.../libgomp.c-c++-common/omp_target_memset-3.c:17:20: note: initializing argument 1 of 'void check_val(int8_t*, int, size_t)'
.../libgomp.c-c++-common/omp_target_memset-3.c:38:18: warning: pointer of type 'void *' used in arithmetic [-Wpointer-arith]
.../libgomp.c-c++-common/omp_target_memset-3.c:38:18: error: invalid conversion from 'void*' to 'int8_t*' {aka 'signed char*'} [-fpermissive]
.../libgomp.c-c++-common/omp_target_memset-3.c:17:20: note: initializing argument 1 of 'void check_val(int8_t*, int, size_t)'
.../libgomp.c-c++-common/omp_target_memset-3.c: In function 'int main()':
.../libgomp.c-c++-common/omp_target_memset-3.c:46:7: warning: pointer of type 'void *' used in arithmetic [-Wpointer-arith]
The following two-liner fixes that, tested on x86_64-linux and i686-linux.
2025-06-03 Jakub Jelinek <jakub@redhat.com>
PR libgomp/120444
* testsuite/libgomp.c-c++-common/omp_target_memset-3.c (test_it):
Change ptr argument type from void * to int8_t *.
(main): Change ptr variable type from void * to int8_t * and cast
omp_target_alloc result to the latter type.
|
|
PR libgomp/120444
include/ChangeLog:
* cuda/cuda.h (cuMemsetD8, cuMemsetD8Async): Declare.
libgomp/ChangeLog:
* libgomp-plugin.h (GOMP_OFFLOAD_memset): Declare.
* libgomp.h (struct gomp_device_descr): Add memset_func.
* libgomp.map (GOMP_6.0.1): Add omp_target_memset{,_async}.
* libgomp.texi (Device Memory Routines): Document them.
* omp.h.in (omp_target_memset, omp_target_memset_async): Declare.
* omp_lib.f90.in (omp_target_memset, omp_target_memset_async):
Add interfaces.
* omp_lib.h.in (omp_target_memset, omp_target_memset_async): Likewise.
* plugin/cuda-lib.def: Add cuMemsetD8.
* plugin/plugin-gcn.c (struct hsa_runtime_fn_info): Add
hsa_amd_memory_fill_fn.
(init_hsa_runtime_functions): DLSYM_OPT_FN load it.
(GOMP_OFFLOAD_memset): New.
* plugin/plugin-nvptx.c (GOMP_OFFLOAD_memset): New.
* target.c (omp_target_memset_int, omp_target_memset,
omp_target_memset_async_helper, omp_target_memset_async): New.
(gomp_load_plugin_for_device): Add DLSYM (memset).
* testsuite/libgomp.c-c++-common/omp_target_memset.c: New test.
* testsuite/libgomp.c-c++-common/omp_target_memset-2.c: New test.
* testsuite/libgomp.c-c++-common/omp_target_memset-3.c: New test.
* testsuite/libgomp.fortran/omp_target_memset.f90: New test.
* testsuite/libgomp.fortran/omp_target_memset-2.f90: New test.
|
|
libgomp/
* testsuite/libgomp.c++/target-std__valarray-1.C: New.
* testsuite/libgomp.c++/target-std__valarray-1.output: Likewise.
|
|
offload targets, a number of USM variants
libgomp/
* testsuite/libgomp.c++/target-std__array-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__array-concurrent.C: Adjust.
* testsuite/libgomp.c++/target-std__bitset-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__bitset-concurrent.C: Adjust.
* testsuite/libgomp.c++/target-std__deque-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__deque-concurrent.C: Adjust.
* testsuite/libgomp.c++/target-std__forward_list-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__forward_list-concurrent.C: Adjust.
* testsuite/libgomp.c++/target-std__list-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__list-concurrent.C: Adjust.
* testsuite/libgomp.c++/target-std__map-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__map-concurrent.C: Adjust.
* testsuite/libgomp.c++/target-std__multimap-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__multimap-concurrent.C: Adjust.
* testsuite/libgomp.c++/target-std__multiset-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__multiset-concurrent.C: Adjust.
* testsuite/libgomp.c++/target-std__set-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__set-concurrent.C: Adjust.
* testsuite/libgomp.c++/target-std__span-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__span-concurrent.C: Adjust.
* testsuite/libgomp.c++/target-std__valarray-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__valarray-concurrent.C: Adjust.
* testsuite/libgomp.c++/target-std__vector-concurrent-usm.C: New.
* testsuite/libgomp.c++/target-std__vector-concurrent.C: Adjust.
|
|
offload targets
libgomp/
* testsuite/libgomp.c++/target-std__array-concurrent.C: New.
* testsuite/libgomp.c++/target-std__bitset-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__deque-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__flat_map-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__flat_multimap-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__flat_multiset-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__flat_set-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__forward_list-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__list-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__map-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__multimap-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__multiset-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__set-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__span-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__unordered_map-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__unordered_multimap-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__unordered_multiset-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__unordered_set-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__valarray-concurrent.C: Likewise.
* testsuite/libgomp.c++/target-std__vector-concurrent.C: Likewise.
Co-authored-by: Thomas Schwinge <tschwinge@baylibre.com>
|
|
libgomp/
* testsuite/libgomp.c++/target-std__cmath.C: New.
* testsuite/libgomp.c++/target-std__complex.C: Likewise.
* testsuite/libgomp.c++/target-std__numbers.C: Likewise.
|
|
libgomp/ChangeLog:
* testsuite/libgomp.c++/target-flex-10.C: New test.
* testsuite/libgomp.c++/target-flex-100.C: New test.
* testsuite/libgomp.c++/target-flex-101.C: New test.
* testsuite/libgomp.c++/target-flex-11.C: New test.
* testsuite/libgomp.c++/target-flex-12.C: New test.
* testsuite/libgomp.c++/target-flex-2000.C: New test.
* testsuite/libgomp.c++/target-flex-2001.C: New test.
* testsuite/libgomp.c++/target-flex-2002.C: New test.
* testsuite/libgomp.c++/target-flex-2003.C: New test.
* testsuite/libgomp.c++/target-flex-30.C: New test.
* testsuite/libgomp.c++/target-flex-300.C: New test.
* testsuite/libgomp.c++/target-flex-31.C: New test.
* testsuite/libgomp.c++/target-flex-32.C: New test.
* testsuite/libgomp.c++/target-flex-33.C: New test.
* testsuite/libgomp.c++/target-flex-41.C: New test.
* testsuite/libgomp.c++/target-flex-60.C: New test.
* testsuite/libgomp.c++/target-flex-61.C: New test.
* testsuite/libgomp.c++/target-flex-62.C: New test.
* testsuite/libgomp.c++/target-flex-70.C: New test.
* testsuite/libgomp.c++/target-flex-80.C: New test.
* testsuite/libgomp.c++/target-flex-81.C: New test.
* testsuite/libgomp.c++/target-flex-90.C: New test.
* testsuite/libgomp.c++/target-flex-common.h: New test.
Co-authored-by: Thomas Schwinge <tschwinge@baylibre.com>
|
|
... to avoid running into ICEs per PR119835, until that's resolved properly.
PR middle-end/119835
gcc/
* tree-nrv.cc (pass_nrv::execute): Defuse 'RESULT_DECL' check.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c:
'#pragma GCC optimize "-fno-inline"'.
* testsuite/libgomp.c-c++-common/target-abi-struct-1.c: New.
* testsuite/libgomp.c-c++-common/target-abi-struct-1-O0.c: Adjust.
Co-authored-by: Richard Biener <rguenther@suse.de>
|
|
This patch adds support for "declare mapper" directives (and the "mapper"
modifier on "map" clauses) for C.
gcc/c/ChangeLog:
* c-decl.cc (c_omp_mapper_id, c_omp_mapper_decl, c_omp_mapper_lookup,
c_omp_extract_mapper_directive, c_omp_map_array_section,
c_omp_scan_mapper_bindings_r, c_omp_scan_mapper_bindings): New
functions.
* c-objc-common.h (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES,
LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE,
LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define langhooks for C.
* c-parser.cc (c_parser_omp_clause_map): Add declare_mapper_p
parameter; handle mapper modifier.
(c_parser_omp_all_clauses): Update call to c_parser_omp_clause_map.
(c_parser_omp_target): Instantiate explicit mappers and record bindings
for implicit mappers.
(c_parser_omp_declare_mapper): Parse "declare mapper" directives.
(c_parser_omp_declare): Support "declare mapper".
(c_parser_omp_declare_reduction): Use inform not error_at.
* c-tree.h (c_omp_finish_mapper_clauses, c_omp_mapper_lookup,
c_omp_extract_mapper_directive, c_omp_map_array_section,
c_omp_mapper_id, c_omp_mapper_decl, c_omp_scan_mapper_bindings,
c_omp_instantiate_mappers): Add prototypes.
* c-typeck.cc (c_finish_omp_clauses): Handle GOMP_MAP_PUSH_MAPPER_NAME
and GOMP_MAP_POP_MAPPER_NAME.
(c_omp_finish_mapper_clauses): New function (langhook).
libgomp/ChangeLog:
* testsuite/libgomp.c-c++-common/declare-mapper-9.c: Enable for C.
* testsuite/libgomp.c-c++-common/declare-mapper-10.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-11.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-12.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-13.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-14.c: Likewise.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/declare-mapper-3.c: Enable for C.
* c-c++-common/gomp/declare-mapper-4.c: Likewise.
* c-c++-common/gomp/declare-mapper-5.c: Likewise.
* c-c++-common/gomp/declare-mapper-6.c: Likewise.
* c-c++-common/gomp/declare-mapper-7.c: Likewise.
* c-c++-common/gomp/declare-mapper-8.c: Likewise.
* c-c++-common/gomp/declare-mapper-9.c: Likewise.
* c-c++-common/gomp/declare-mapper-10.c: Likewise.
* c-c++-common/gomp/declare-mapper-12.c: Likewise.
* c-c++-common/gomp/map-6.c: Update dg-error.
* gcc.dg/gomp/udr-3.c: Update for change to dg-note.
* c-c++-common/gomp/declare-mapper-11.c: New.
* gcc.dg/gomp/declare-mapper-10.c: New test.
* gcc.dg/gomp/declare-mapper-11.c: New test.
* gcc.dg/gomp/declare-mapper-13.c: New test.
|
|
This patch adds support for OpenMP 5.0 "declare mapper" functionality
for C++. I've merged it to og13 based on the last version
posted upstream, with some minor changes due to the newly-added
'present' map modifier support. There's also a fix to splay-tree
traversal in gimplify.cc:omp_instantiate_implicit_mappers, and this patch
omits the rearrangement of gimplify.cc:gimplify_{scan,adjust}_omp_clauses
that I separated out into its own patch and applied (to og13) already.
gcc/c-family/
* c-common.h (c_omp_region_type): Add C_ORT_DECLARE_MAPPER and
C_ORT_OMP_DECLARE_MAPPER codes.
(omp_mapper_list): Add forward declaration.
(c_omp_find_nested_mappers, c_omp_instantiate_mappers): Add prototypes.
* c-omp.cc (c_omp_find_nested_mappers): New function.
(remap_mapper_decl_info): New struct.
(remap_mapper_decl_1, omp_instantiate_mapper,
c_omp_instantiate_mappers): New functions.
gcc/cp/
* constexpr.cc (reduced_constant_expression_p): Add OMP_DECLARE_MAPPER
case.
(cxx_eval_constant_expression, potential_constant_expression_1):
Likewise.
* cp-gimplify.cc (cxx_omp_finish_mapper_clauses): New function.
* cp-objcp-common.h (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES,
LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE,
LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define langhooks.
* cp-tree.h (lang_decl_base): Add omp_declare_mapper_p field. Recount
spare bits comment.
(DECL_OMP_DECLARE_MAPPER_P): New macro.
(omp_mapper_id): Add prototype.
(cp_check_omp_declare_mapper): Add prototype.
(omp_instantiate_mappers): Add prototype.
(cxx_omp_finish_mapper_clauses): Add prototype.
(cxx_omp_mapper_lookup): Add prototype.
(cxx_omp_extract_mapper_directive): Add prototype.
(cxx_omp_map_array_section): Add prototype.
* decl.cc (check_initializer): Add OpenMP declare mapper support.
(cp_finish_decl): Set DECL_INITIAL for OpenMP declare mapper var decls
as appropriate.
* decl2.cc (mark_used): Instantiate OpenMP "declare mapper" magic var
decls.
* error.cc (dump_omp_declare_mapper): New function.
(dump_simple_decl): Use above.
* parser.cc (cp_parser_omp_clause_map): Add KIND parameter. Support
"mapper" modifier.
(cp_parser_omp_all_clauses): Add KIND argument to
cp_parser_omp_clause_map call.
(cp_parser_omp_target): Call omp_instantiate_mappers before
finish_omp_clauses.
(cp_parser_omp_declare_mapper): New function.
(cp_parser_omp_declare): Add "declare mapper" support.
* pt.cc (tsubst_decl): Adjust name of "declare mapper" magic var decls
once we know their type.
(tsubst_omp_clauses): Call omp_instantiate_mappers before
finish_omp_clauses, for target regions.
(tsubst_expr): Support OMP_DECLARE_MAPPER nodes.
(instantiate_decl): Instantiate initialiser (i.e definition) for OpenMP
declare mappers.
* semantics.cc (gimplify.h): Include.
(omp_mapper_id, omp_mapper_lookup, omp_extract_mapper_directive,
cxx_omp_map_array_section, cp_check_omp_declare_mapper): New functions.
(finish_omp_clauses): Delete GOMP_MAP_PUSH_MAPPER_NAME and
GOMP_MAP_POP_MAPPER_NAME artificial clauses.
(omp_target_walk_data): Add MAPPERS field.
(finish_omp_target_clauses_r): Scan for uses of struct/union/class type
variables.
(finish_omp_target_clauses): Create artificial mapper binding clauses
for used structs/unions/classes in offload region.
gcc/fortran/
* parse.cc (tree.h, fold-const.h, tree-hash-traits.h): Add includes
(for additions to omp-general.h).
gcc/
* gimplify.cc (gimplify_omp_ctx): Add IMPLICIT_MAPPERS field.
(new_omp_context): Initialise IMPLICIT_MAPPERS hash map.
(delete_omp_context): Delete IMPLICIT_MAPPERS hash map.
(instantiate_mapper_info): New structs.
(remap_mapper_decl_1, omp_mapper_copy_decl, omp_instantiate_mapper,
omp_instantiate_implicit_mappers): New functions.
(gimplify_scan_omp_clauses): Handle MAPPER_BINDING clauses.
(gimplify_adjust_omp_clauses): Instantiate implicit declared mappers.
(gimplify_omp_declare_mapper): New function.
(gimplify_expr): Call above function.
* langhooks-def.h (lhd_omp_mapper_lookup,
lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): Add
prototypes.
(LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES,
LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE,
LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define macros.
(LANG_HOOK_DECLS): Add above macros.
* langhooks.cc (lhd_omp_mapper_lookup,
lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): New
dummy functions.
* langhooks.h (lang_hooks_for_decls): Add OMP_FINISH_MAPPER_CLAUSES,
OMP_MAPPER_LOOKUP, OMP_EXTRACT_MAPPER_DIRECTIVE, OMP_MAP_ARRAY_SECTION
hooks.
* omp-general.h (omp_name_type<T>): Add templatized struct, hash type
traits (for omp_name_type<tree> specialization).
(omp_mapper_list<T>): Add struct.
* tree-core.h (omp_clause_code): Add OMP_CLAUSE__MAPPER_BINDING_.
* tree-pretty-print.cc (dump_omp_clause): Support GOMP_MAP_UNSET,
GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping
clauses. Support OMP_CLAUSE__MAPPER_BINDING_ and OMP_DECLARE_MAPPER.
* tree.cc (omp_clause_num_ops, omp_clause_code_name): Add
OMP_CLAUSE__MAPPER_BINDING_.
* tree.def (OMP_DECLARE_MAPPER): New tree code.
* tree.h (OMP_DECLARE_MAPPER_ID, OMP_DECLARE_MAPPER_DECL,
OMP_DECLARE_MAPPER_CLAUSES): New defines.
(OMP_CLAUSE__MAPPER_BINDING__ID, OMP_CLAUSE__MAPPER_BINDING__DECL,
OMP_CLAUSE__MAPPER_BINDING__MAPPER): New defines.
include/
* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_UNSET,
GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping
clause types.
gcc/testsuite/
* c-c++-common/gomp/map-6.c: Update error scan output.
* c-c++-common/gomp/declare-mapper-3.c: New test (only enabled for C++
for now).
* c-c++-common/gomp/declare-mapper-4.c: Likewise.
* c-c++-common/gomp/declare-mapper-5.c: Likewise.
* c-c++-common/gomp/declare-mapper-6.c: Likewise.
* c-c++-common/gomp/declare-mapper-7.c: Likewise.
* c-c++-common/gomp/declare-mapper-8.c: Likewise.
* c-c++-common/gomp/declare-mapper-9.c: Likewise.
* c-c++-common/gomp/declare-mapper-10.c: Likewise.
* c-c++-common/gomp/declare-mapper-12.c: Likewise.
* g++.dg/gomp/declare-mapper-1.C: New test.
* g++.dg/gomp/declare-mapper-2.C: New test.
* g++.dg/gomp/declare-mapper-3.C: New test.
libgomp/
* testsuite/libgomp.c++/declare-mapper-1.C: New test.
* testsuite/libgomp.c++/declare-mapper-2.C: New test.
* testsuite/libgomp.c++/declare-mapper-3.C: New test.
* testsuite/libgomp.c++/declare-mapper-4.C: New test.
* testsuite/libgomp.c++/declare-mapper-5.C: New test.
* testsuite/libgomp.c++/declare-mapper-6.C: New test.
* testsuite/libgomp.c++/declare-mapper-7.C: New test.
* testsuite/libgomp.c++/declare-mapper-8.C: New test.
* testsuite/libgomp.c-c++-common/declare-mapper-9.c: New test (only
enabled for C++ for now).
* testsuite/libgomp.c-c++-common/declare-mapper-10.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-11.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-12.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-13.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-14.c: Likewise.
Co-authored-by: Tobias Burnus <tburnus@baylibre.com>
|
|
libgomp/ChangeLog:
PR libgomp/93226
* libgomp-plugin.h (GOMP_OFFLOAD_openacc_async_dev2dev): New
prototype.
* libgomp.h (struct acc_dispatch_t): Add dev2dev_func.
(gomp_copy_dev2dev): New prototype.
* libgomp.map (OACC_2.6.1): New; add acc_memcpy_device{,_async}.
* libgomp.texi (acc_memcpy_device): New.
* oacc-mem.c (memcpy_tofrom_device): Change to take from/to
device boolean; use memcpy not memmove; add early return if
size == 0 or same device + same ptr.
(acc_memcpy_to_device, acc_memcpy_to_device_async,
acc_memcpy_from_device, acc_memcpy_from_device_async): Update.
(acc_memcpy_device, acc_memcpy_device_async): New.
* openacc.f90 (acc_memcpy_device, acc_memcpy_device_async):
Add interface.
* openacc_lib.h (acc_memcpy_device, acc_memcpy_device_async):
Likewise.
* openacc.h (acc_memcpy_device, acc_memcpy_device_async): Add
prototype.
* plugin/plugin-gcn.c (GOMP_OFFLOAD_openacc_async_host2dev):
Update comment.
(GOMP_OFFLOAD_openacc_async_dev2host): Update call.
(GOMP_OFFLOAD_openacc_async_dev2dev): New.
* plugin/plugin-nvptx.c (cuda_memcpy_dev_sanity_check): New.
(GOMP_OFFLOAD_dev2dev): Call it.
(GOMP_OFFLOAD_openacc_async_dev2dev): New.
* target.c (gomp_copy_dev2dev): New.
(gomp_load_plugin_for_device): Load dev2dev and async_dev2dev.
* testsuite/libgomp.oacc-c-c++-common/acc_memcpy_device-1.c: New test.
* testsuite/libgomp.oacc-fortran/acc_memcpy_device-1.f90: New test.
|
|
[PR118694]
This should have been part of commit r16-838-gb3d07ec7ac2ccd or
r16-883-g5d6ed6d604ff94 - all showing the same issue:
'!$omp target' followed by a metadirective with 'teams'; if
the metadirective cannot be early resolved, a diagnostic
error is shown about using directives between 'target' and
'teams'.
While the message is misleading, the problem is that the
host invokes 'target' differently when 'teams' is present;
in this case, host fallback + amdgcn offload require the
no-teams case, nvptx offload the teams case such that it
only can be resolved at runtime.
Mark the error as 'dg-bogus + xfail' to silence the FAIL,
when nvptx offloading is compiled for. (If not, the
metadirective can be resolved early during compilation.)
libgomp/ChangeLog:
PR middle-end/118694
* testsuite/libgomp.fortran/metadirective-1.f90: xfail when
compiling (also) for nvptx offloading as an error is then expected.
|
|
[PR118694]
OpenMP's 'target teams' is strictly coupled with 'teams'; if the latter
exists, the kernel is launched in directly with multiple teams. Thus,
the host has to know whether the teams construct exists or not. For
#pragma omp target
#pragma omp metadirective when (device={arch("nvptx")}: teams loop)
it is simple when 'nvptx' offloading is not supported, otherwise it depends
on the default device at runtime as the user code asks for a single team for
host fallback and gcn offload and multiple for nvptx offload.
In any case, this commit ensures that no FAIL is printed, whatever a
future solution might look like. Instead of a dg-bogus combined with an
'xfail offload_target_nvptx', one an also argue that a dg-error for
'target offload_target_nvptx' would be more appropriate.
libgomp/ChangeLog:
PR middle-end/118694
* testsuite/libgomp.c-c++-common/metadirective-1.c: xfail when
compiling (also) for nvptx offloading as an error is then expected.
|
|
We've got 'gcc/stor-layout.cc:finalize_type_size':
/* Handle empty records as per the x86-64 psABI. */
TYPE_EMPTY_P (type) = targetm.calls.empty_record_p (type);
(Indeed x86_64 is still the only target to define 'TARGET_EMPTY_RECORD_P',
calling 'gcc/tree.cc-default_is_empty_record'.)
And so it happens that for an empty struct used in code offloaded from x86_64
host (but not powerpc64le host, for example), we get to see 'TYPE_EMPTY_P' in
offloading compilation (where the offload targets (currently?) don't use it
themselves, and therefore aren't prepared to handle it).
For nvptx offloading compilation, this causes wrong code generation:
'ptxas [...] error : Call has wrong number of parameters', as nvptx code
generation for function definition doesn't pay attention to this flag (say, in
'gcc/config/nvptx/nvptx.cc:pass_in_memory', or whereever else would be
appropriate to handle that), but the generic code 'gcc/calls.cc:expand_call'
via 'gcc/function.cc:aggregate_value_p' does pay attention to it, and we thus
get mismatching function definition vs. function call.
This issue apparently isn't a problem for GCN offloading, but I don't know if
that's by design or by accident.
Richard Biener:
> It looks like TYPE_EMPTY_P is only used during RTL expansion for ABI
> purposes, so computing it during layout_type is premature as shown here.
>
> I would suggest to simply re-compute it at offload stream-in time.
(For avoidance of doubt, the additions to 'gcc.target/nvptx/abi-struct-arg.c',
'gcc.target/nvptx/abi-struct-ret.c' are not dependent on the offload streaming
code changes, but are just to mirror the changes to
'libgomp.oacc-c-c++-common/abi-struct-1.c'.)
PR lto/120308
gcc/
* lto-streamer-out.cc (hash_tree): Don't handle 'TYPE_EMPTY_P' for
'lto_stream_offload_p'.
* tree-streamer-in.cc (unpack_ts_type_common_value_fields):
Likewise.
* tree-streamer-out.cc (pack_ts_type_common_value_fields):
Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c: Add empty
structure testing.
gcc/testsuite/
* gcc.target/nvptx/abi-struct-arg.c: Add empty structure testing.
* gcc.target/nvptx/abi-struct-ret.c: Likewise.
|
|
'libgomp.oacc-c-c++-common/abi-struct-1.c'
libgomp/
* testsuite/libgomp.c-c++-common/target-abi-struct-1-O0.c: New.
* testsuite/libgomp.oacc-c-c++-common/abi-struct-1.c: Likewise.
|
|
libgomp/
* testsuite/libgomp.oacc-fortran/lib-13.f90: End data region after
wait API calls.
|
|
The check whether the location expression in map clause has allocatable
components was failing for some derived-type array expressions such as
map(var%tiles(1))
as the compiler produced
_4 = var.tiles;
MEMREF(_4, _5);
This commit now also handles this case.
gcc/fortran/ChangeLog:
* trans-openmp.cc (gfc_omp_deep_mapping_do): Handle SSA_NAME if
a def_stmt is available.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/alloc-comp-4.f90: New test.
|
|
map(var[:n]), n = 0
For map(ptr[:0]), the used map kind is GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
and it is permitted that 'ptr' does not exist. 'ptr' is set to the device
pointee if it exists or to the host value otherwise.
For map(ptr[:3]), the variable is first mapped and then ptr is updated to point
to the just-mapped device data; the attachment uses GOMP_MAP_ATTACH.
For map(ptr[:n]), generates always a GOMP_MAP_ATTACH, but when n == 0, it
was failing with:
"pointer target not mapped for attach"
The solution is not to fail but first to check whether it was mapped before.
It turned out that for the mapping part, GCC adds a run-time check whether
n == 0 - and uses GOMP_MAP_ZERO_LEN_ARRAY_SECTION for the mapping.
Thus, we just have to check whether there such a mapping for the address
for which the GOMP_MAP_ATTACH. was requested. And, if there was, the
error diagnostic can be skipped.
Unsurprisingly, this issue occurs in real-world code; it was detected in
a code that distributes work via MPI and for some processes, some bounds
ended up to be zero.
libgomp/ChangeLog:
* target.c (gomp_attach_pointer): Return bool; accept additional
bool to optionally silence the fatal pointee-not-found error.
(gomp_map_vars_internal): If the pointee could not be found,
check whether it was mapped as GOMP_MAP_ZERO_LEN_ARRAY_SECTION.
* libgomp.h (gomp_attach_pointer): Update prototype.
* oacc-mem.c (acc_attach_async, goacc_enter_data_internal): Update
calls.
* testsuite/libgomp.c/target-map-zero-sized.c: New test.
* testsuite/libgomp.c/target-map-zero-sized-2.c: New test.
* testsuite/libgomp.c/target-map-zero-sized-3.c: New test.
|
|
'dynamic_cast'" [PR119692]
PR target/119692
libgomp/
* testsuite/libgomp.c++/pr119692-1-4.C: '{ dg-timeout 10 }'.
* testsuite/libgomp.c++/pr119692-1-5.C: Likewise.
* testsuite/libgomp.c++/target-exceptions-bad_cast-1.C: Likewise.
* testsuite/libgomp.c++/target-exceptions-bad_cast-2.C: Likewise.
* testsuite/libgomp.oacc-c++/exceptions-bad_cast-1.C: Likewise.
* testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C: Likewise.
|
|
gcc/
* config/nvptx/nvptx-sm.def: Add '61'.
* config/nvptx/nvptx-gen.h: Regenerate.
* config/nvptx/nvptx-gen.opt: Likewise.
* config/nvptx/nvptx.cc (first_ptx_version_supporting_sm): Adjust.
* config/nvptx/nvptx.opt (-march-map=sm_61, -march-map=sm_62):
Likewise.
* config.gcc: Likewise.
* doc/invoke.texi (Nvidia PTX Options): Document '-march=sm_61'.
* config/nvptx/gen-multilib-matches-tests: Extend.
gcc/testsuite/
* gcc.target/nvptx/march-map=sm_61.c: Adjust.
* gcc.target/nvptx/march-map=sm_62.c: Likewise.
* gcc.target/nvptx/march=sm_61.c: New.
libgomp/
* testsuite/libgomp.c/declare-variant-3-sm61.c: New.
* testsuite/libgomp.c/declare-variant-3.h: Adjust.
|
|
While the tests checked whether the CUDA/HIP runtime is available
before processing them, the execution was then done unconditionally,
leading to FAIL when the default device was the host (or the wrong
offload device).
Now the test is only executed ('run') when the default device is an
Nvidia or AMD GPU (depending on the test case, cf. the test file name).
Otherwise, only a 'link' test is done. (Except when the effective-target
check cannot find the runtime lib - then the test is skipped [as before].)
Note: The cublas/hipblas tests use variant functions and iterate over
all devices, such that the cublas or hipblas, respectively, is only
called when the active device is an AMD or Nvidia device, respectively,
while for the host and other device types the fallback is called.
libgomp/ChangeLog:
* testsuite/libgomp.c/interop-cuda-full.c: Use 'link' instead
of 'run' when the default device is "! offload_device_nvptx".
* testsuite/libgomp.c/interop-cuda-libonly.c: Likewise.
* testsuite/libgomp.c/interop-hip-nvidia-full.c: Likewise.
* testsuite/libgomp.c/interop-hip-nvidia-no-headers.c: Likewise.
* testsuite/libgomp.c/interop-hip-nvidia-no-hip-header.c: Likewise.
* testsuite/libgomp.fortran/interop-hip-nvidia-full.F90: Likewise.
* testsuite/libgomp.fortran/interop-hip-nvidia-no-module.F90: Likewise.
* testsuite/libgomp.c/interop-hip-amd-full.c: Use 'link' instead
of 'run' when the default device is "! offload_device_gcn".
* testsuite/libgomp.c/interop-hip-amd-no-hip-header.c: Likewise.
* testsuite/libgomp.fortran/interop-hip-amd-full.F90: Likewise.
* testsuite/libgomp.fortran/interop-hip-amd-no-module.F90: Likewise.
|
