Age | Commit message (Collapse) | Author | Files | Lines |
|
|
|
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.
|
|
|
|
One of the buffers that printed the list of set 'omp requires'
requirements missed the 'self' clause addition, being potentially
to short when all device-affecting clauses were passed. Solved it
by moving the sizeof(<string of all permitted values>" into a new
'#define' just above the associated gomp_requires_to_name function.
libgomp/ChangeLog:
* target.c (GOMP_REQUIRES_NAME_BUF_LEN): Define.
(GOMP_offload_register_ver, gomp_target_init): Use it for the
char buffer size.
|
|
libgomp/ChangeLog:
* libgomp.texi (omp_init_allocator): Refer to 'Memory allocation'
for available memory spaces.
(OMP_ALLOCATOR): Move list of traits and predefined memspaces
and allocators to ...
(Memory allocation): ... here. Document omp(x)::allocator::*;
minor wording tweaks, be more explicit about memkind, pinned and
pool_size.
Co-authored-by: waffl3x <waffl3x@baylibre.com>
|
|
|
|
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.
|
|
|
|
By adding omp_get_initial_device and omp_get_num_devices builtins for
C, C++, and Fortran, the following can be achieved:
* By making them pure, multiple calls can be avoiding in some cases.
* Some comparisons can be optimized at compile time.
omp_get_initial_device will be converted to omp_get_num_devices for
consistency; note that OpenMP 6 also permits omp_initial_device (== -1)
as value.
If GCC has not been configure for offloading, either intrinsic will
leads to 0 - and on the offload side, -1 (= omp_initial_device) is
returned for omp_initial_device.
gcc/fortran/ChangeLog:
* f95-lang.cc (ATTR_PURE_NOTHROW_LIST): Define.
* trans-expr.cc (get_builtin_fn): Handle omp_get_num_devices
and omp_get_intrinsic_device.
* gfortran.h (gfc_option_t): Add disable_omp_... for them.
* options.cc (gfc_handle_option): Handle them with
-fno-builtin-.
gcc/ChangeLog:
* gimple-fold.cc (gimple_fold_builtin_omp_get_initial_device,
gimple_fold_builtin_omp_get_num_devices): New.
(gimple_fold_builtin): Call them.
* omp-builtins.def (BUILT_IN_OMP_GET_INITIAL_DEVICE): Add
(BUILT_IN_OMP_GET_NUM_DEVICES): Make uservisible + pure.
libgomp/ChangeLog:
* libgomp.texi (omp_get_num_devices, omp_get_intrinsic_device):
Document builtin handling.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/omp_get_num_devices_initial_device-2.c: New test.
* c-c++-common/gomp/omp_get_num_devices_initial_device.c: New test.
* gfortran.dg/gomp/omp_get_num_devices_initial_device-2.f90: New test.
* gfortran.dg/gomp/omp_get_num_devices_initial_device.f90: New test.
Co-authored-by: Sandra Loosemore <sloosemore@baylibre.com>
|
|
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.
|
|
|
|
GCC uses the 6.0 types - which are unfortunately not quite compatible with
code expecting 5.1/5.2 data types. Therefore, this commit adds a note to
hopefully reduce surprises. Namely:
For C/C++: while OpenMP 5.1 and 5.2 used 'int *ret_code', OpenMP 6.0 uses
'omp_interop_rc_t *ret_code' in omp_interop_{int,ptr,str} and 'int' instead
of 'omp_interop_rc_t ret_code' in omp_get_interop_rc_desc.
Neither C nor C++ like passing the wrong pointer type, albeit for C, GCC < 14
and clang only warn (gcc >= r14-6037-g9715c545d33b3a has an error) and
using -fpermissive turns it into a warning and -Wno-incompatible-pointer-types
silences it for C.
C++ also dislikes passing an int to an enum, albeit -fpermissive turns the
error into a warning with g++ (but not clang++). And, here, using an enum
on the caller side works with both int and enum on the callee side.
libgomp/ChangeLog:
* libgomp.texi (omp_interop_{int,ptr,str,rc_desc}): Add note about
the 'ret_code' type change in OpenMP 6.
Co-authored-by: Sandra Loosemore <sloosemore@baylibre.com>
|
|
|
|
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.
|
|
[PR116792]
In r15-3752-g48261bd26df624 I added a test plugin that overrode the
regular output, instead emitting diagnostics in crude HTML form.
In r15-4760-g0b73e9382ab51c I added support for multiple kinds of
diagnostic output simultaneously, adding
-fdiagnostics-add-output=DIAGNOSTICS-OUTPUT-SPEC
-fdiagnostics-set-output=DIAGNOSTICS-OUTPUT-SPEC
for adding/changing the kind of diagnostics output, supporting
"text" and "sarif" output schemes.
This patch promotes the HTML output code from the test plugins so
that it is available from "-fdiagnostics-add-output=", using a
new "experimental-html" scheme, to allow simultaneous text, sarif
and html output, and to make it easier to experiment with. The
patch adds Python-based testing of the emitted HTML.
The patch does not affect the generated HTML, which is still crude, and
not yet ready for end-users. I hope to improve it in followups.
gcc/ChangeLog:
PR other/116792
* Makefile.in (OBJS-libcommon): Add diagnostic-format-html.o.
* diagnostic-format-html.cc: Move here from
testsuite/gcc.dg/plugin/diagnostic_plugin_xhtml_format.cc.
Simplify includes. Rename "xhtml" to "html" throughout.
(write_escaped_text): Drop.
(class xhtml_stream_output_format): Drop.
(class html_file_output_format): Reimplement using
diagnostic_output_file.
(diagnostic_output_format_init_xhtml): Drop.
(diagnostic_output_format_init_xhtml_stderr): Drop.
(diagnostic_output_format_init_xhtml_file): Drop.
(diagnostic_output_format_open_html_file): New.
(make_html_sink): New.
(xhtml_format_selftests): Convert to...
(diagnostic_format_html_cc_tests): ...this.
(plugin_is_GPL_compatible): Drop.
(plugin_init): Drop.
* diagnostic-format-html.h: New file.
* doc/invoke.texi (-fdiagnostics-add-output=): Add
"experimental-html" scheme.
* opts-diagnostic.cc: Include "diagnostic-format-html.h".
(class html_scheme_handler): New.
(output_factory::output_factory): Add html_scheme_handler.
(html_scheme_handler::make_sink): New.
* selftest-run-tests.cc (selftest::run_tests): Call the new
selftests.
* selftest.h (selftest::diagnostic_format_html_cc_tests): New
decl.
gcc/testsuite/ChangeLog:
PR other/116792
* gcc.dg/plugin/diagnostic_plugin_xhtml_format.cc: Move to
gcc/diagnostic-format-html.cc.
* gcc.dg/html-output/html-output.exp: New support script.
* gcc.dg/html-output/missing-semicolon.c: New test.
* gcc.dg/html-output/missing-semicolon.py: New test script.
* gcc.dg/plugin/diagnostic-test-xhtml-1.c: Deleted test.
* gcc.dg/plugin/plugin.exp (plugin_test_list): Drop moved plugin
and its deleted test.
* lib/gcc-dg.exp (load_lib): Add load_lib of scanhtml.exp.
* lib/htmltest.py: New support script.
* lib/scanhtml.exp: New support script, based on scansarif.exp.
libatomic/ChangeLog:
PR other/116792
* testsuite/lib/libatomic.exp: Add load_lib of scanhtml.exp.
libgomp/ChangeLog:
PR other/116792
* testsuite/lib/libgomp.exp: Add load_lib of scanhtml.exp.
libitm/ChangeLog:
PR other/116792
* testsuite/lib/libitm.exp: Add load_lib of scanhtml.exp.
libphobos/ChangeLog:
PR other/116792
* testsuite/lib/libphobos-dg.exp: Add load_lib of scanhtml.exp.
libvtv/ChangeLog:
PR other/116792
* testsuite/lib/libvtv-dg.exp: Add load_lib of scanhtml.exp.
Signed-off-by: David Malcolm <dmalcolm@redhat.com>
|
|
|
|
When host memory is device accessible - independent whether mapping is done or
not (i.e. self map), the 'vtab' pointer becomes accessible, which stores the
dynamic type's type and size information.
In principle, we want to test: USM available but mapping is still done, but
as there is no simple + reliable not-crashing way to test for this, those
checks are skipped in the (pre)existing test file map-alloc-comp-9.f90.
Or rather: those are only active with self-maps, which is currently only true
for the host.
This commit adds map-alloc-comp-9-usm.f90 which runs the same test with
'omp requires unified_shared_memory'. While OpenMP permits both actual
mapping and self maps with this flag, it in theory covers the missing cases.
However, currently, GCC always uses self maps with USM. Still, having a
device-run self-maps check is better than nothing, even if it misses the
most interesting case.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/map-alloc-comp-9.f90: Process differently
when USE_USM_REQUIREMENT is set.
* testsuite/libgomp.fortran/map-alloc-comp-9-usm.f90: New test.
|
|
|