aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite/libgomp.c
AgeCommit message (Collapse)AuthorFilesLines
2024-03-06Revert "Set num_threads to 50 on 32-bit hppa in two libgomp loop tests"John David Anglin1-7/+1
This reverts commit b14209715e659f6d3ca0f9eef9a4851e7bd6e373.
2024-02-12libgomp: testsuite: Don't XPASS libgomp.c/alloc-pinned-1.c etc. on non-Linux ↵Rainer Orth2-2/+8
targets [PR113448] Two libgomp tests XPASS on Solaris (any non-Linux target actually) since their introduction: XPASS: libgomp.c/alloc-pinned-1.c execution test XPASS: libgomp.c/alloc-pinned-2.c execution test The problem is that the test just prints OS unsupported and exits successfully, while the test is XFAILed: /* { dg-xfail-run-if "Pinning not implemented on this host" { ! *-*-linux-gnu } } */ Fixed by aborting immediately after the message above in the non-Linux case. Tested on i386-pc-solaris2.11 and i686-pc-linux-gnu. 2024-02-02 Rainer Orth <ro@CeBiTec.Uni-Bielefeld.DE> libgomp: PR testsuite/113448 * testsuite/libgomp.c/alloc-pinned-1.c [!__linux__] (CHECK_SIZE): Call abort. * testsuite/libgomp.c/alloc-pinned-2.c [!__linux__] (CHECK_SIZE): Likewise.
2024-02-01Set num_threads to 50 on 32-bit hppa in two libgomp loop testsJohn David Anglin1-1/+7
We support a maximum of 50 threads on 32-bit hppa. 2024-02-01 John David Anglin <danglin@gcc.gnu.org> libgomp/ChangeLog: * testsuite/libgomp.c++/loop-3.C: Set num_threads to 50 on 32-bit hppa. * testsuite/libgomp.c/omp-loop03.c: Likewise.
2024-01-29libgomp.c/declare-variant-4.h: Fix used variant function for gfx1030/gfx1100Tobias Burnus1-2/+2
libgomp/ChangeLog: * testsuite/libgomp.c/declare-variant-4.h: Use gfx1100/gfx1030 function not gfx90a for gfx1100/gfx1030 context selector. Signed-off-by: Tobias Burnus <tburnus@baylibre.com>
2024-01-26amdgcn: config.gcc - enable gfx1030 and gfx1100 multilib; add them to the docsTobias Burnus3-0/+32
gcc/ChangeLog: * config.gcc (amdgcn-*-*): Add gfx1030 and gfx1100 to TM_MULTILIB_CONFIG. * doc/install.texi (Configuration amdgcn-*-*): Mention gfx1030/gfx1100. * doc/invoke.texi (AMD GCN Options): Add gfx1030 and gfx1100 to -march/-mtune. libgomp/ChangeLog: * testsuite/libgomp.c/declare-variant-4.h: Add variant functions for gfx1030 and gfx1100. * testsuite/libgomp.c/declare-variant-4-gfx1030.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx1100.c: New test. Signed-off-by: Tobias Burnus <tburnus@baylibre.com>
2024-01-22xfail libgomp.c/declare-variant-4-{fiji,gfx803}.cTobias Burnus2-0/+4
Since r14-4734-g56ed1055b2f40ac162ae8d382280ac07a33f789f, GCC no longer builds the Fiji (alias gfx803) libraries by default as support for it was removed in ROCm 4.0 and will be removed in LLVM 18. Thus, unless gfx803 is explicitly enabled, the following testcases will fail to link as libgomp is not available for Fiji. Hence, this commit xfails those testcases. libgomp/ChangeLog: * testsuite/libgomp.c/declare-variant-4-fiji.c: Xfail as fiji support is no longer enabled by default. * testsuite/libgomp.c/declare-variant-4-gfx803.c: Likewise. Signed-off-by: Tobias Burnus <tburnus@baylibre.com>
2024-01-20Don't run libgomp.c/simd-math-1.c on hppa*-*-hpux*John David Anglin1-1/+1
hppa*-*-hpux* lacks necessary math functions. 2024-01-20 John David Anglin <danglin@gcc.gnu.org> libgomp/ChangeLog: * testsuite/libgomp.c/simd-math-1.c: Don't run on hppa*-*-hpux*.
2024-01-17openmp: Add OpenMP _BitInt support [PR113409]Jakub Jelinek1-0/+65
The following patch adds support for _BitInt iterators of OpenMP canonical loops (with the preexisting limitation that when not using compile time static scheduling the iterators in the library are at most unsigned long long or signed long, so one can't in the runtime/dynamic/guided etc. cases iterate more than what those types can represent, like is the case of e.g. __int128 iterators too) and the testcase also covers linear/reduction clauses for them. 2024-01-17 Jakub Jelinek <jakub@redhat.com> PR middle-end/113409 * omp-general.cc (omp_adjust_for_condition): Handle BITINT_TYPE like INTEGER_TYPE. (omp_extract_for_data): Use build_bitint_type rather than build_nonstandard_integer_type if either iter_type or loop->v type is BITINT_TYPE. * omp-expand.cc (expand_omp_for_generic, expand_omp_taskloop_for_outer, expand_omp_taskloop_for_inner): Handle BITINT_TYPE like INTEGER_TYPE. * testsuite/libgomp.c/bitint-1.c: New test.
2024-01-03Update copyright years.Jakub Jelinek2-2/+2
2023-12-18libgomp: Make libgomp.c/declare-variant-1.c test x86 specificJakub Jelinek1-9/+8
As written earlier, this test was written with the x86 specifics in mind and adding dg-final directives for it for other arches makes it unreadable. If a declare variant call can be resolved in gimple already as in the aarch64 or gcn cases, it can be done in gcc.dg/gomp/ and I believe we have tests like that already, the point of the test is that it is not known during gimplification time which exact call should be chosen as it depends on which declare simd clone it will be in. 2023-12-18 Jakub Jelinek <jakub@redhat.com> * testsuite/libgomp.c/declare-variant-1.c: Restrict the test to x86, drop because of that unneeded target selector from other directives and remove the aarch64 specific ones.
2023-12-15Fix tests for gompAndre Vieira1-6/+7
This is to fix testisms initially introduced by: commit f5fc001a84a7dbb942a6252b3162dd38b4aae311 Author: Andre Vieira <andre.simoesdiasvieira@arm.com> Date: Mon Dec 11 14:24:41 2023 +0000 aarch64: enable mixed-types for aarch64 simdclones gcc/testsuite/ChangeLog: * gcc.dg/gomp/pr87887-1.c: Fixed test. * gcc.dg/gomp/pr89246-1.c: Likewise. * gcc.dg/gomp/simd-clones-2.c: Likewise. libgomp/ChangeLog: * testsuite/libgomp.c/declare-variant-1.c: Fixed test. * testsuite/libgomp.fortran/declare-simd-1.f90: Likewise.
2023-12-13libgomp: basic pinned memory on LinuxAndrew Stubbs4-0/+541
Implement the OpenMP pinned memory trait on Linux hosts using the mlock syscall. Pinned allocations are performed using mmap, not malloc, to ensure that they can be unpinned safely when freed. This implementation will work OK for page-scale allocations, and finer-grained allocations will be implemented in a future patch. libgomp/ChangeLog: * allocator.c (MEMSPACE_ALLOC): Add PIN. (MEMSPACE_CALLOC): Add PIN. (MEMSPACE_REALLOC): Add PIN. (MEMSPACE_FREE): Add PIN. (MEMSPACE_VALIDATE): Add PIN. (omp_init_allocator): Use MEMSPACE_VALIDATE to check pinning. (omp_aligned_alloc): Add pinning to all MEMSPACE_* calls. (omp_aligned_calloc): Likewise. (omp_realloc): Likewise. (omp_free): Likewise. * config/linux/allocator.c: New file. * config/nvptx/allocator.c (MEMSPACE_ALLOC): Add PIN. (MEMSPACE_CALLOC): Add PIN. (MEMSPACE_REALLOC): Add PIN. (MEMSPACE_FREE): Add PIN. (MEMSPACE_VALIDATE): Add PIN. * config/gcn/allocator.c (MEMSPACE_ALLOC): Add PIN. (MEMSPACE_CALLOC): Add PIN. (MEMSPACE_REALLOC): Add PIN. (MEMSPACE_FREE): Add PIN. * libgomp.texi: Switch pinned trait to supported. (MEMSPACE_VALIDATE): Add PIN. * testsuite/libgomp.c/alloc-pinned-1.c: New test. * testsuite/libgomp.c/alloc-pinned-2.c: New test. * testsuite/libgomp.c/alloc-pinned-3.c: New test. * testsuite/libgomp.c/alloc-pinned-4.c: New test. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
2023-12-11aarch64: enable mixed-types for aarch64 simdclonesAndre Vieira1-2/+4
This patch enables the use of mixed-types for simd clones for AArch64, adds aarch64 as a target_vect_simd_clones and corrects the way the simdlen is chosen for non-specified simdlen clauses according to the 'Vector Function Application Binary Interface Specification for AArch64'. Additionally this patch also restricts combinations of simdlen and return/argument types that map to vectors larger than 128 bits as we currently do not have a way to represent these types in a way that is consistent internally and externally. gcc/ChangeLog: * config/aarch64/aarch64.cc (lane_size): New function. (aarch64_simd_clone_compute_vecsize_and_simdlen): Determine simdlen according to NDS rule and reject combination of simdlen and types that lead to vectors larger than 128bits. gcc/testsuite/ChangeLog: * lib/target-supports.exp: Add aarch64 targets to vect_simd_clones. * c-c++-common/gomp/declare-variant-14.c: Adapt test for aarch64. * c-c++-common/gomp/pr60823-1.c: Likewise. * c-c++-common/gomp/pr60823-2.c: Likewise. * c-c++-common/gomp/pr60823-3.c: Likewise. * g++.dg/gomp/attrs-10.C: Likewise. * g++.dg/gomp/declare-simd-1.C: Likewise. * g++.dg/gomp/declare-simd-3.C: Likewise. * g++.dg/gomp/declare-simd-4.C: Likewise. * g++.dg/gomp/declare-simd-7.C: Likewise. * g++.dg/gomp/declare-simd-8.C: Likewise. * g++.dg/gomp/pr88182.C: Likewise. * gcc.dg/declare-simd.c: Likewise. * gcc.dg/gomp/declare-simd-1.c: Likewise. * gcc.dg/gomp/declare-simd-3.c: Likewise. * gcc.dg/gomp/pr87887-1.c: Likewise. * gcc.dg/gomp/pr87895-1.c: Likewise. * gcc.dg/gomp/pr89246-1.c: Likewise. * gcc.dg/gomp/pr99542.c: Likewise. * gcc.dg/gomp/simd-clones-2.c: Likewise. * gcc.dg/vect/vect-simd-clone-1.c: Likewise. * gcc.dg/vect/vect-simd-clone-2.c: Likewise. * gcc.dg/vect/vect-simd-clone-4.c: Likewise. * gcc.dg/vect/vect-simd-clone-5.c: Likewise. * gcc.dg/vect/vect-simd-clone-6.c: Likewise. * gcc.dg/vect/vect-simd-clone-7.c: Likewise. * gcc.dg/vect/vect-simd-clone-8.c: Likewise. * gfortran.dg/gomp/declare-simd-2.f90: Likewise. * gfortran.dg/gomp/declare-simd-coarray-lib.f90: Likewise. * gfortran.dg/gomp/declare-variant-14.f90: Likewise. * gfortran.dg/gomp/pr79154-1.f90: Likewise. * gfortran.dg/gomp/pr83977.f90: Likewise. libgomp/ChangeLog: * testsuite/libgomp.c/declare-variant-1.c: Adapt test for aarch64. * testsuite/libgomp.fortran/declare-simd-1.f90: Likewise.
2023-12-06amdgcn, libgomp: low-latency allocatorAndrew Stubbs1-1/+1
This implements the OpenMP low-latency memory allocator for AMD GCN using the small per-team LDS memory (Local Data Store). Since addresses can now refer to LDS space, the "Global" address space is no-longer compatible. This patch therefore switches the backend to use entirely "Flat" addressing (which supports both memories). A future patch will re-enable "global" instructions for cases where it is known to be safe to do so. gcc/ChangeLog: * config/gcn/gcn-builtins.def (DISPATCH_PTR): New built-in. * config/gcn/gcn.cc (gcn_init_machine_status): Disable global addressing. (gcn_expand_builtin_1): Implement GCN_BUILTIN_DISPATCH_PTR. libgomp/ChangeLog: * config/gcn/libgomp-gcn.h (TEAM_ARENA_START): Move to here. (TEAM_ARENA_FREE): Likewise. (TEAM_ARENA_END): Likewise. (GCN_LOWLAT_HEAP): New. * config/gcn/team.c (LITTLEENDIAN_CPU): New, and import hsa.h. (__gcn_lowlat_init): New prototype. (gomp_gcn_enter_kernel): Initialize the low-latency heap. * libgomp.h (TEAM_ARENA_START): Move to libgomp.h. (TEAM_ARENA_FREE): Likewise. (TEAM_ARENA_END): Likewise. * plugin/plugin-gcn.c (lowlat_size): New variable. (print_kernel_dispatch): Label the group_segment_size purpose. (init_environment_variables): Read GOMP_GCN_LOWLAT_POOL. (create_kernel_dispatch): Pass low-latency head allocation to kernel. (run_kernel): Use shadow; don't assume values. * testsuite/libgomp.c/omp_alloc-traits.c: Enable for amdgcn. * config/gcn/allocator.c: New file. * libgomp.texi: Document low-latency implementation details.
2023-12-06openmp, nvptx: low-lat memory access traitsAndrew Stubbs7-6/+107
The NVPTX low latency memory is not accessible outside the team that allocates it, and therefore should be unavailable for allocators with the access trait "all". This change means that the omp_low_lat_mem_alloc predefined allocator no longer works (but omp_cgroup_mem_alloc still does). libgomp/ChangeLog: * allocator.c (MEMSPACE_VALIDATE): New macro. (omp_init_allocator): Use MEMSPACE_VALIDATE. (omp_aligned_alloc): Use OMP_LOW_LAT_MEM_ALLOC_INVALID. (omp_aligned_calloc): Likewise. (omp_realloc): Likewise. * config/nvptx/allocator.c (nvptx_memspace_validate): New function. (MEMSPACE_VALIDATE): New macro. (OMP_LOW_LAT_MEM_ALLOC_INVALID): New define. * libgomp.texi: Document low-latency implementation details. * testsuite/libgomp.c/omp_alloc-1.c (main): Add gnu_lowlat. * testsuite/libgomp.c/omp_alloc-2.c (main): Add gnu_lowlat. * testsuite/libgomp.c/omp_alloc-3.c (main): Add gnu_lowlat. * testsuite/libgomp.c/omp_alloc-4.c (main): Add access trait. * testsuite/libgomp.c/omp_alloc-5.c (main): Add gnu_lowlat. * testsuite/libgomp.c/omp_alloc-6.c (main): Add access trait. * testsuite/libgomp.c/omp_alloc-traits.c: New test.
2023-12-06libgomp, nvptx: low-latency memory allocatorAndrew Stubbs6-0/+544
This patch adds support for allocating low-latency ".shared" memory on NVPTX GPU device, via the omp_low_lat_mem_space and omp_alloc. The memory can be allocated, reallocated, and freed using a basic but fast algorithm, is thread safe and the size of the low-latency heap can be configured using the GOMP_NVPTX_LOWLAT_POOL environment variable. The use of the PTX dynamic_smem_size feature means that low-latency allocator will not work with the PTX 3.1 multilib. For now, the omp_low_lat_mem_alloc allocator also works, but that will change when I implement the access traits. libgomp/ChangeLog: * allocator.c (MEMSPACE_ALLOC): New macro. (MEMSPACE_CALLOC): New macro. (MEMSPACE_REALLOC): New macro. (MEMSPACE_FREE): New macro. (predefined_alloc_mapping): New array. Add _Static_assert to match. (ARRAY_SIZE): New macro. (omp_aligned_alloc): Use MEMSPACE_ALLOC. Implement fall-backs for predefined allocators. Simplify existing fall-backs. (omp_free): Use MEMSPACE_FREE. (omp_calloc): Use MEMSPACE_CALLOC. Implement fall-backs for predefined allocators. Simplify existing fall-backs. (omp_realloc): Use MEMSPACE_REALLOC, MEMSPACE_ALLOC, and MEMSPACE_FREE. Implement fall-backs for predefined allocators. Simplify existing fall-backs. * config/nvptx/team.c (__nvptx_lowlat_pool): New asm variable. (__nvptx_lowlat_init): New prototype. (gomp_nvptx_main): Call __nvptx_lowlat_init. * libgomp.texi: Update memory space table. * plugin/plugin-nvptx.c (lowlat_pool_size): New variable. (GOMP_OFFLOAD_init_device): Read the GOMP_NVPTX_LOWLAT_POOL envvar. (GOMP_OFFLOAD_run): Apply lowlat_pool_size. * basic-allocator.c: New file. * config/nvptx/allocator.c: New file. * testsuite/libgomp.c/omp_alloc-1.c: New test. * testsuite/libgomp.c/omp_alloc-2.c: New test. * testsuite/libgomp.c/omp_alloc-3.c: New test. * testsuite/libgomp.c/omp_alloc-4.c: New test. * testsuite/libgomp.c/omp_alloc-5.c: New test. * testsuite/libgomp.c/omp_alloc-6.c: New test. Co-authored-by: Kwok Cheung Yeung <kcy@codesourcery.com> Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
2023-11-30Fix 'libgomp.c/declare-variant-4-*.c', add 'libgomp.c/declare-variant-4.c'Thomas Schwinge8-12/+31
These test cases being 'dg-skip-if [...] { ! amdgcn-*-* }' meant they were only ever considered for 'amdgcn-*-*' target -- that is, GCN *target*, not GCN *offload target*, what is intended to be tested here, but for which they thus always were UNSUPPORTED. Use the same style of 'dg-[...]' directives as for the nvptx offloading test cases, 'libgomp.c/declare-variant-3*.c'. Fix-up for commit 1fd508744eccda9ad9c6d6fcce5b2ea9c568818d "amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors". libgomp/ * testsuite/libgomp.c/declare-variant-4-fiji.c: Adjust. * testsuite/libgomp.c/declare-variant-4-gfx803.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.h: Likewise. * testsuite/libgomp.c/declare-variant-4.c: New.
2023-11-30Spin 'dg-do run' part of 'libgomp.c/declare-variant-3-sm30.c' off into new ↵Thomas Schwinge3-1/+14
'libgomp.c/declare-variant-3.c' Having nvptx offloading configured doesn't imply being able to run nvptx offloading test cases on the test host. Also, make 'libgomp.c/declare-variant-3.c' work for all non-offloading and offloading cases. Fix-up for commit 59b8ade88774b4dcf1691a8f650cdbb86cc30862 "[libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c". libgomp/ * testsuite/libgomp.c/declare-variant-3-sm30.c: Turn 'dg-do run' into 'dg-do link'. * testsuite/libgomp.c/declare-variant-3.c: New. * testsuite/libgomp.c/declare-variant-3.h: Extend.
2023-11-30In 'libgomp.c/declare-variant-{3,4}-*.c', restrict 'scan-offload-tree-dump's ↵Thomas Schwinge12-12/+12
to 'only_for_offload_target [...]' ... to care for the case where not just one but both of GCN and nvptx offloading are enabled. In that case, we currently get: UNRESOLVED: libgomp.c/declare-variant-3-sm30.c scan-amdgcn-amdhsa-offload-tree-dump optimized "= f30 \\(\\);" ... in addition to: PASS: libgomp.c/declare-variant-3-sm30.c scan-nvptx-none-offload-tree-dump optimized "= f30 \\(\\);" Etc. Fix-up for commit 59b8ade88774b4dcf1691a8f650cdbb86cc30862 "[libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c", and commit 1fd508744eccda9ad9c6d6fcce5b2ea9c568818d "amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectors". libgomp/ * testsuite/libgomp.c/declare-variant-3-sm30.c: Restrict 'scan-offload-tree-dump' to 'only_for_offload_target nvptx-none'. * testsuite/libgomp.c/declare-variant-3-sm35.c: Likewise. * testsuite/libgomp.c/declare-variant-3-sm53.c: Likewise. * testsuite/libgomp.c/declare-variant-3-sm70.c: Likewise. * testsuite/libgomp.c/declare-variant-3-sm75.c: Likewise. * testsuite/libgomp.c/declare-variant-3-sm80.c: Likewise. * testsuite/libgomp.c/declare-variant-4-fiji.c: Restrict 'scan-offload-tree-dump' to 'only_for_offload_target amdgcn-amdhsa'. * testsuite/libgomp.c/declare-variant-4-gfx803.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.
2023-11-30Fix 'libgomp.c/declare-variant-3-*.c' compilation for configurations where ↵Thomas Schwinge6-0/+6
GCN offloading is enabled in addition to nvptx The GCN offloading compiler doesn't like '-misa=sm_30' etc.; restrict to '-foffload=nvptx-none' compilation only. Fix-up for commit 59b8ade88774b4dcf1691a8f650cdbb86cc30862 "[libgomp, testsuite, nvptx] Add libgomp.c/declare-variant-3-sm*.c". libgomp/ * testsuite/libgomp.c/declare-variant-3-sm30.c: 'dg-additional-options -foffload=nvptx-none'. * testsuite/libgomp.c/declare-variant-3-sm35.c: Likewise. * testsuite/libgomp.c/declare-variant-3-sm53.c: Likewise. * testsuite/libgomp.c/declare-variant-3-sm70.c: Likewise. * testsuite/libgomp.c/declare-variant-3-sm75.c: Likewise. * testsuite/libgomp.c/declare-variant-3-sm80.c: Likewise.
2023-11-29In 'libgomp.c/target-simd-clone-{1,2,3}.c', restrict ↵Thomas Schwinge3-5/+5
'scan-offload-ipa-dump's to 'only_for_offload_target amdgcn-amdhsa' This gets rid of UNRESOLVEDs if nvptx offloading compilation is enabled in addition to GCN: PASS: libgomp.c/target-simd-clone-1.c (test for excess errors) PASS: libgomp.c/target-simd-clone-1.c scan-amdgcn-amdhsa-offload-ipa-dump simdclone "Generated local clone _ZGV.*N.*_addit" -UNRESOLVED: libgomp.c/target-simd-clone-1.c scan-nvptx-none-offload-ipa-dump simdclone "Generated local clone _ZGV.*N.*_addit" PASS: libgomp.c/target-simd-clone-1.c scan-amdgcn-amdhsa-offload-ipa-dump simdclone "Generated local clone _ZGV.*M.*_addit" -UNRESOLVED: libgomp.c/target-simd-clone-1.c scan-nvptx-none-offload-ipa-dump simdclone "Generated local clone _ZGV.*M.*_addit" PASS: libgomp.c/target-simd-clone-2.c (test for excess errors) PASS: libgomp.c/target-simd-clone-2.c scan-amdgcn-amdhsa-offload-ipa-dump-not simdclone "Generated .* clone" -UNRESOLVED: libgomp.c/target-simd-clone-2.c scan-nvptx-none-offload-ipa-dump-not simdclone "Generated .* clone" PASS: libgomp.c/target-simd-clone-3.c (test for excess errors) PASS: libgomp.c/target-simd-clone-3.c scan-amdgcn-amdhsa-offload-ipa-dump simdclone "device doesn't match" -UNRESOLVED: libgomp.c/target-simd-clone-3.c scan-nvptx-none-offload-ipa-dump simdclone "device doesn't match" PASS: libgomp.c/target-simd-clone-3.c scan-amdgcn-amdhsa-offload-ipa-dump-not simdclone "Generated .* clone" -UNRESOLVED: libgomp.c/target-simd-clone-3.c scan-nvptx-none-offload-ipa-dump-not simdclone "Generated .* clone" Minor fix-up for commit 309e2d95e3b930c6f15c8a5346b913158404c76d 'OpenMP: Generate SIMD clones for functions with "declare target"'. libgomp/ * testsuite/libgomp.c/target-simd-clone-1.c: Restrict 'scan-offload-ipa-dump's to 'only_for_offload_target amdgcn-amdhsa'. * testsuite/libgomp.c/target-simd-clone-2.c: Likewise. * testsuite/libgomp.c/target-simd-clone-3.c: Likewise.
2023-11-22Adjust 'libgomp.c/declare-variant-{3,4}-[...]' for inter-procedural value ↵Thomas Schwinge2-0/+15
range propagation ..., that is, commit 53ba8d669550d3a1f809048428b97ca607f95cf5 "inter-procedural value range propagation", after which we see: [-PASS:-]{+FAIL:+} libgomp.c/declare-variant-3-sm30.c scan-nvptx-none-offload-tree-dump optimized "= f30 \\(\\);" Etc. That's due to: @@ -144,13 +144,11 @@ __attribute__((omp target entrypoint, noclone)) void main._omp_fn.0 (const struct .omp_data_t.3 & restrict .omp_data_i) { - int _3; int * _5; <bb 2> [local count: 1073741824]: - _3 = f30 (); _5 = *.omp_data_i_4(D).v; - *_5 = _3; + *_5 = 30; return; It's nice to see this optimization work here, too, but it does interfere with how we're currently testing OpenMP 'declare variant'. libgomp/ * testsuite/libgomp.c/declare-variant-3.h (f30, f35, f53, f70) (f75, f80, f): Add '__attribute__ ((noipa))'. * testsuite/libgomp.c/declare-variant-4.h (gfx803, gfx900, gfx906) (gfx908, gfx90a, f): Likewise.
2023-09-20OpenMP: Add ME support for 'omp allocate' stack variablesTobias Burnus3-0/+529
Call GOMP_alloc/free for 'omp allocate' allocated variables. This is for C only as C++ and Fortran show a sorry already in the FE. Note that this only applies to stack variables as the C FE shows a sorry for static variables. gcc/ChangeLog: * gimplify.cc (gimplify_bind_expr): Call GOMP_alloc/free for 'omp allocate' variables; move stack cleanup after other cleanup. (omp_notice_variable): Process original decl when decl of the value-expression for a 'omp allocate' variable is passed. * omp-low.cc (scan_omp_1_op): Handle 'omp allocate' variables libgomp/ChangeLog: * libgomp.texi (OpenMP 5.1 Impl.): Mark 'omp allocate' as implemented for C only. * testsuite/libgomp.c/allocate-4.c: New test. * testsuite/libgomp.c/allocate-5.c: New test. * testsuite/libgomp.c/allocate-6.c: New test. gcc/testsuite/ChangeLog: * c-c++-common/gomp/allocate-11.c: Remove C-only dg-message for 'sorry, unimplemented'. * c-c++-common/gomp/allocate-12.c: Likewise. * c-c++-common/gomp/allocate-15.c: Likewise. * c-c++-common/gomp/allocate-9.c: Likewise. * c-c++-common/gomp/allocate-10.c: New test. * c-c++-common/gomp/allocate-17.c: New test.
2023-08-23libgomp, testsuite: Do not call nonstandard functionsFrancois-Xavier Coudert1-0/+9
The following functions are not standard, and not always available (e.g., on darwin). They should not be called unless available: gamma, gammaf, scalb, scalbf, significand, and significandf. libgomp/ChangeLog: * testsuite/lib/libgomp.exp: Add effective target. * testsuite/libgomp.c/simd-math-1.c: Avoid calling nonstandard functions.
2023-07-26OpenMP: Call cuMemcpy2D/cuMemcpy3D for nvptx for omp_target_memcpy_rectTobias Burnus1-3/+3
When copying a 2D or 3D rectangular memmory block, the performance is better when using CUDA's cuMemcpy2D/cuMemcpy3D instead of copying the data one by one. That's what this commit does. Additionally, it permits device-to-device copies, if neccessary using a temporary variable on the host. include/ChangeLog: * cuda/cuda.h (CUlimit): Add CUDA_ERROR_NOT_INITIALIZED, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_INVALID_HANDLE. (CUarray, CUmemorytype, CUDA_MEMCPY2D, CUDA_MEMCPY3D, CUDA_MEMCPY3D_PEER): New typdefs. (cuMemcpy2D, cuMemcpy2DAsync, cuMemcpy2DUnaligned, cuMemcpy3D, cuMemcpy3DAsync, cuMemcpy3DPeer, cuMemcpy3DPeerAsync): New prototypes. libgomp/ChangeLog: * libgomp-plugin.h (GOMP_OFFLOAD_memcpy2d, GOMP_OFFLOAD_memcpy3d): New prototypes. * libgomp.h (struct gomp_device_descr): Add memcpy2d_func and memcpy3d_func. * libgomp.texi (nvtpx): Document when cuMemcpy2D/cuMemcpy3D is used. * oacc-host.c (memcpy2d_func, .memcpy3d_func): Init with NULL. * plugin/cuda-lib.def (cuMemcpy2D, cuMemcpy2DUnaligned, cuMemcpy3D): Invoke via CUDA_ONE_CALL. * plugin/plugin-nvptx.c (GOMP_OFFLOAD_memcpy2d, GOMP_OFFLOAD_memcpy3d): New. * target.c (omp_target_memcpy_rect_worker): (omp_target_memcpy_rect_check, omp_target_memcpy_rect_copy): Permit all device-to-device copyies; invoke new plugins for 2D and 3D copying when available. (gomp_load_plugin_for_device): DLSYM the new plugin functions. * testsuite/libgomp.c/target-12.c: Fix dimension bug. * testsuite/libgomp.fortran/target-12.f90: Likewise. * testsuite/libgomp.fortran/target-memcpy-rect-1.f90: New test.
2023-06-19Fix DejaGnu directive syntax error in 'libgomp.c/target-51.c'Thomas Schwinge1-1/+1
ERROR: libgomp.c/target-51.c: unknown dg option: \} for "}" Fix-up for recent commit 01fe115ba7eafebcf97bbac9e157038a003d0c85 "libgomp.c/target-51.c: Accept more error-msg variants in dg-output". libgomp/ * testsuite/libgomp.c/target-51.c: Fix DejaGnu directive syntax error.
2023-06-19libgomp.c/target-51.c: Accept more error-msg variants in dg-outputTobias Burnus1-2/+1
Depending on the details, the testcase can fail with different but related messages; all of the following all could be observed for this testcase: libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY, but device cannot be used for offloading libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY, but device not found libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY, but only the host device is available Before, the last two were tested for with 'target offload_device' and '! offload_device', respectively. Now, all three are accepted by matching '.*' already after 'but' and without distinguishing whether the effective target is an offload_device or not. (For completeness, there is a fourth error that follows this pattern: 'OMP_TARGET_OFFLOAD is set to MANDATORY, but device is finalized'.) libgomp/ * testsuite/libgomp.c/target-51.c: Accept more error msg variants as expected dg-output.
2023-06-19OpenMP (C/C++): Keep pointer value of unmapped ptr with default mapping ↵Tobias Burnus1-6/+15
[PR110270] For C/C++ pointers, default implicit mapping firstprivatizes the pointer but if the memory it points to is mapped, the it is updated to point to the device memory (by attaching a zero sized array section of the pointed-to storage). However, if the pointed-to storage wasn't mapped, the pointer was set to NULL on the device side (OpenMP 5.0/5.1 semantic). With this commit, the pointer retains the on-host address in that case (OpenMP 5.2 semantic). The new semantic avoids an explicit map/firstprivate/is_device_ptr in the following sensible cases: Special values (e.g. pointer or 0x1, 0x2 etc.), explicitly device allocated memory (e.g. omp_target_alloc), and with (unified) shared memory. (Note: With (U)SM, mappings still must be tracked, at least when omp_target_associate_ptr does not fail when passing in two destinct pointers.) libgomp/ PR middle-end/110270 * target.c (gomp_map_vars_internal): Copy host value instead of NULL for GOMP_MAP_ZERO_LEN_ARRAY_SECTION if not mapped. * libgomp.texi (OpenMP 5.2 Impl.): Mark as 'Y'. * testsuite/libgomp.c/target-19.c: Update expected value. * testsuite/libgomp.c++/target-18.C: Likewise. * testsuite/libgomp.c++/target-19.C: Likewise. * testsuite/libgomp.c-c++-common/requires-unified-addr-2.c: New test. * testsuite/libgomp.c-c++-common/target-implicit-map-3.c: New test. * testsuite/libgomp.c-c++-common/target-implicit-map-4.c: New test.
2023-06-16libgomp: Fix OMP_TARGET_OFFLOAD=mandatoryTobias Burnus2-0/+43
It turned out that gomp_init_targets_once() was not run when directly calling 'omp target' or 'omp target (enter/exit) data' causing an abort with OMP_TARGET_OFFLOAD=mandatory wrongly claiming that no device is available. It was called a tiny bit later but few lines too late for updating the default-device-var. libgomp/ChangeLog: * target.c (resolve_device): Call gomp_get_num_devices early to ensure gomp_init_targets_once was called before using default-device-var. * testsuite/libgomp.c/target-55.c: New test. * testsuite/libgomp.c/target-55a.c: New test.
2023-06-15libgomp: Extend OMP_ALLOCATOR, add affinity env var docTobias Burnus6-0/+104
Support OpenMP 5.1's syntax for OMP_ALLOCATOR as well, which permits besides predefined allocators also predefined memspaces optionally followed by traits. Additionally, this commit adds the previously lacking documentation for OMP_ALLOCATOR, OMP_AFFINITY_FORMAT and OMP_DISPLAY_AFFINITY. libgomp/ChangeLog: * env.c (gomp_def_allocator_envvar): New var. (parse_allocator): Handle OpenMP 5.1 syntax. (cleanup_env): New. (omp_display_env): Output gomp_def_allocator_envvar for an allocator with traits. * libgomp.texi (OMP_ALLOCATOR, OMP_AFFINITY_FORMAT, OMP_DISPLAY_AFFINITY): New. * testsuite/libgomp.c/allocator-1.c: New test. * testsuite/libgomp.c/allocator-2.c: New test. * testsuite/libgomp.c/allocator-3.c: New test. * testsuite/libgomp.c/allocator-4.c: New test. * testsuite/libgomp.c/allocator-5.c: New test. * testsuite/libgomp.c/allocator-6.c: New test.
2023-06-14Align a 'OMP_TARGET_OFFLOAD=mandatory' diagnostic with othersThomas Schwinge1-1/+1
On 2023-06-14T11:42:22+0200, Tobias Burnus <tobias@codesourcery.com> wrote: > On 14.06.23 10:09, Thomas Schwinge wrote: >> Let me know if I should also adjust the new 'target { ! offload_device }' >> diagnostic "[...] MANDATORY but only the host device is available" to >> include a comma before 'but', for consistency with the other existing >> diagnostics (cited above)? > > I think it makes sense to be consistent. Thus: Yes, please add the commas. Fix-up for recent commit 18c8b56c7d67a9e37acf28822587786f0fc0efbc "OpenMP: Set default-device-var with OMP_TARGET_OFFLOAD=mandatory". libgomp/ * target.c (resolve_device): Align a 'OMP_TARGET_OFFLOAD=mandatory' diagnostic with others. * testsuite/libgomp.c/target-51.c: Adjust.
2023-06-14driver: Forward '-lgfortran', '-lm' to offloading compilationThomas Schwinge1-1/+0
..., so that users don't manually need to specify '-foffload-options=-lgfortran', '-foffload-options=-lm' in addition to '-lgfortran', '-lm' (specified manually, or implicitly by the driver). gcc/ * gcc.cc (driver_handle_option): Forward host '-lgfortran', '-lm' to offloading compilation. * config/gcn/mkoffload.cc (main): Adjust. * config/nvptx/mkoffload.cc (main): Likewise. * doc/invoke.texi (foffload-options): Update example. libgomp/ * testsuite/libgomp.fortran/fortran.exp (lang_link_flags): Don't set. * testsuite/libgomp.oacc-fortran/fortran.exp (lang_link_flags): Likewise. * testsuite/libgomp.c/simd-math-1.c: Remove '-foffload-options=-lm'. * testsuite/libgomp.fortran/fortran-torture_execute_math.f90: Likewise. * testsuite/libgomp.oacc-fortran/fortran-torture_execute_math.f90: Likewise.
2023-06-14Fix typo in 'libgomp.c/target-51.c'Thomas Schwinge1-1/+1
..., and therefore, given 'target offload_device': PASS: libgomp.c/target-51.c (test for excess errors) PASS: libgomp.c/target-51.c execution test [-FAIL:-]{+PASS:+} libgomp.c/target-51.c output pattern test Fix-up for recent commit 18c8b56c7d67a9e37acf28822587786f0fc0efbc "OpenMP: Set default-device-var with OMP_TARGET_OFFLOAD=mandatory". libgomp/ * testsuite/libgomp.c/target-51.c: Fix typo.
2023-06-14OpenMP: Set default-device-var with OMP_TARGET_OFFLOAD=mandatoryTobias Burnus8-0/+210
OMP_TARGET_OFFLOAD=mandatory handling was before inconsistent. Hence, in OpenMP 5.2 it was clarified/extended by having implications on the default-device-var; additionally, omp_initial_device and omp_invalid_device enum values/PARAMETERs were added; support for it was added in r13-1066-g1158fe43407568 including aborting for omp_invalid_device and non-conforming device numbers. Only the mandatory handling was missing. Namely, while the default-device-var is usually initialized to value 0, with 'mandatory' it must have the value 'omp_invalid_device' if and only if zero non-host devices are available. (The OMP_DEFAULT_DEVICE env var overrides this as it comes semantically after the initialization.) To achieve this, default-device-var is now initialized to MIN_INT. If there is no 'mandatory', it is set to 0 directly after env var parsing. Otherwise, it is updated in gomp_target_init to either 0 or omp_invalid_device. To ensure INT_MIN is never seen by the user, both the omp_get_default_device API routine and omp_display_env (user call and OMP_DISPLAY_ENV env var) call gomp_init_targets_once() in that case. libgomp/ChangeLog: * env.c (gomp_default_icv_values): Init default_device_var to an nonconforming value - INT_MIN. (initialize_env): After env-var parsing, set default_device_var to device 0 unless OMP_TARGET_OFFLOAD=mandatory. (omp_display_env): If default_device_var is INT_MIN, call gomp_init_targets_once. * icv-device.c (omp_get_default_device): Likewise. * libgomp.texi (OMP_DEFAULT_DEVICE): Update init description. (OpenMP 5.2 Impl. Status): Mark OMP_TARGET_OFFLOAD=mandatory as 'Y'. * target.c (resolve_device): Improve error message device-num < 0 with 'mandatory' and no no-host devices available. (gomp_target_init): Set default-device-var if INT_MIN. * testsuite/libgomp.c/target-48.c: New test. * testsuite/libgomp.c/target-49.c: New test. * testsuite/libgomp.c/target-50.c: New test. * testsuite/libgomp.c/target-50a.c: New test. * testsuite/libgomp.c/target-51.c: New test. * testsuite/libgomp.c/target-52.c: New test. * testsuite/libgomp.c/target-53.c: New test. * testsuite/libgomp.c/target-54.c: New test.
2023-05-15libgomp testsuite: Have each '*.exp' file specify the compiler to use [PR91884]Thomas Schwinge1-0/+2
..., which is still 'GCC_UNDER_TEST' for all of them; no change in behavior. PR testsuite/91884 libgomp/ * testsuite/lib/libgomp.exp (libgomp_target_compile): Don't specify compiler. * testsuite/libgomp.c++/c++.exp (ALWAYS_CFLAGS): Specify compiler. * testsuite/libgomp.c/c.exp (ALWAYS_CFLAGS): Likewise. * testsuite/libgomp.fortran/fortran.exp (ALWAYS_CFLAGS): Likewise. * testsuite/libgomp.graphite/graphite.exp (ALWAYS_CFLAGS): Likewise. * testsuite/libgomp.oacc-c++/c++.exp (ALWAYS_CFLAGS): Likewise. * testsuite/libgomp.oacc-c/c.exp (ALWAYS_CFLAGS): Likewise. * testsuite/libgomp.oacc-fortran/fortran.exp (ALWAYS_CFLAGS): Likewise.
2023-05-12libgomp testsuite: Localize 'lang_[...]' etc.Thomas Schwinge1-11/+0
..., instead of letting them bleed into the next '*.exp' file, requiring clean-up there. libgomp/ * testsuite/libgomp.c++/c++.exp: Localize 'lang_[...]' etc. * testsuite/libgomp.c/c.exp: Likewise. * testsuite/libgomp.fortran/fortran.exp: Likewise. * testsuite/libgomp.graphite/graphite.exp: Likewise. * testsuite/libgomp.oacc-c++/c++.exp: Likewise. * testsuite/libgomp.oacc-c/c.exp: Likewise. * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
2023-05-09libgomp testsuite: Use 'lang_test_file_found' instead of 'lang_test_file'Thomas Schwinge1-2/+2
The value of 'lang_test_file' isn't actually used anywhere. libgomp/ * testsuite/libgomp.c++/c++.exp: Don't set 'lang_test_file'. * testsuite/libgomp.fortran/fortran.exp: Likewise. * testsuite/libgomp.oacc-c++/c++.exp: Likewise. * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. * testsuite/libgomp.c/c.exp: Unset 'lang_test_file_found' instead of 'lang_test_file'. * testsuite/libgomp.oacc-c/c.exp: Likewise. * testsuite/libgomp.graphite/graphite.exp: Likewise. * testsuite/lib/libgomp.exp (libgomp_target_compile): Look for 'lang_test_file_found' instead of 'lang_test_file'.
2023-04-25'omp scan' struct block seq update for OpenMP 5.xTobias Burnus1-0/+121
While OpenMP 5.0 required a single structured block before and after the 'omp scan' directive, OpenMP 5.1 changed this to a 'structured block sequence, denoting 2 or more executable statements in OpenMP 5.1 (whoops!) and zero or more in OpenMP 5.2. This commit updates C/C++ to accept zero statements (but till requires the '{' ... '}' for the final-loop-body) and updates Fortran to accept zero or more than one statements. If there is no preceeding or succeeding executable statement, a warning is shown. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_scan_loop_body): Handle zero exec statements before/after 'omp scan'. gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_scan_loop_body): Handle zero exec statements before/after 'omp scan'. gcc/fortran/ChangeLog: * openmp.cc (gfc_resolve_omp_do_blocks): Handle zero or more than one exec statements before/after 'omp scan'. * trans-openmp.cc (gfc_trans_omp_do): Likewise. libgomp/ChangeLog: * testsuite/libgomp.c-c++-common/scan-1.c: New test. * testsuite/libgomp.c/scan-23.c: New test. * testsuite/libgomp.fortran/scan-2.f90: New test. gcc/testsuite/ChangeLog: * g++.dg/gomp/attrs-7.C: Update dg-error/dg-warning. * gfortran.dg/gomp/loop-2.f90: Likewise. * gfortran.dg/gomp/reduction5.f90: Likewise. * gfortran.dg/gomp/reduction6.f90: Likewise. * gfortran.dg/gomp/scan-1.f90: Likewise. * gfortran.dg/gomp/taskloop-2.f90: Likewise. * c-c++-common/gomp/scan-6.c: New test. * gfortran.dg/gomp/scan-8.f90: New test.
2023-03-02amdgcn: Enable SIMD vectorization of math functionsKwok Cheung Yeung1-0/+217
Calls to vectorized versions of routines in the math library will now be inserted when vectorizing code containing supported math functions. 2023-03-02 Kwok Cheung Yeung <kcy@codesourcery.com> Paul-Antoine Arras <pa@codesourcery.com> gcc/ * builtins.cc (mathfn_built_in_explicit): New. * config/gcn/gcn.cc: Include case-cfn-macros.h. (mathfn_built_in_explicit): Add prototype. (gcn_vectorize_builtin_vectorized_function): New. (gcn_libc_has_function): New. (TARGET_LIBC_HAS_FUNCTION): Define. (TARGET_VECTORIZE_BUILTIN_VECTORIZED_FUNCTION): Define. gcc/testsuite/ * gcc.target/gcn/simd-math-1.c: New testcase. * gcc.target/gcn/simd-math-2.c: New testcase. libgomp/ * testsuite/libgomp.c/simd-math-1.c: New testcase.
2023-01-19openmp: Fix up OpenMP expansion of non-rectangular loops [PR108459]Jakub Jelinek1-0/+41
expand_omp_for_init_counts was using for the case where collapse(2) inner loop has init expression dependent on non-constant multiple of the outer iterator and the condition upper bound expression doesn't depend on the outer iterator fold_unary (NEGATE_EXPR, ...). This will just return NULL if it can't be folded, we need fold_build1 instead. 2023-01-19 Jakub Jelinek <jakub@redhat.com> PR middle-end/108459 * omp-expand.cc (expand_omp_for_init_counts): Use fold_build1 rather than fold_unary for NEGATE_EXPR. * testsuite/libgomp.c/pr108459.c: New test.
2023-01-16Update copyright years.Jakub Jelinek2-2/+2
2022-11-30amdgcn: Support AMD-specific 'isa' traits in OpenMP context selectorsPaul-Antoine Arras7-0/+106
Add support for gfx803 as an alias for fiji. Add test cases for all supported 'isa' values. gcc/ChangeLog: * config/gcn/gcn.cc (gcn_omp_device_kind_arch_isa): Add gfx803. * config/gcn/t-omp-device: Add gfx803. libgomp/ChangeLog: * testsuite/libgomp.c/declare-variant-4-fiji.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx803.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx900.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx906.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx908.c: New test. * testsuite/libgomp.c/declare-variant-4-gfx90a.c: New test. * testsuite/libgomp.c/declare-variant-4.h: New header file.
2022-11-25OpenMP: Generate SIMD clones for functions with "declare target"Sandra Loosemore3-0/+122
This patch causes the IPA simdclone pass to generate clones for functions with the "omp declare target" attribute as if they had "omp declare simd", provided the function appears to be suitable for SIMD execution. The filter is conservative, rejecting functions that write memory or that call other functions not known to be safe. A new option -fopenmp-target-simd-clone is added to control this transformation; it's enabled for offload processing at -O2 and higher. gcc/ChangeLog: * common.opt (fopenmp-target-simd-clone): New option. (target_simd_clone_device): New enum to go with it. * doc/invoke.texi (-fopenmp-target-simd-clone): Document. * flag-types.h (enum omp_target_simd_clone_device_kind): New. * omp-simd-clone.cc (auto_simd_fail): New function. (auto_simd_check_stmt): New function. (plausible_type_for_simd_clone): New function. (ok_for_auto_simd_clone): New function. (simd_clone_create): Add force_local argument, make the symbol have internal linkage if it is true. (expand_simd_clones): Also check for cloneable functions with "omp declare target". Pass explicit_p argument to simd_clone.compute_vecsize_and_simdlen target hook. * opts.cc (default_options_table): Add -fopenmp-target-simd-clone. * target.def (TARGET_SIMD_CLONE_COMPUTE_VECSIZE_AND_SIMDLEN): Add bool explicit_p argument. * doc/tm.texi: Regenerated. * config/aarch64/aarch64.cc (aarch64_simd_clone_compute_vecsize_and_simdlen): Update. * config/gcn/gcn.cc (gcn_simd_clone_compute_vecsize_and_simdlen): Update. * config/i386/i386.cc (ix86_simd_clone_compute_vecsize_and_simdlen): Update. gcc/testsuite/ChangeLog: * g++.dg/gomp/target-simd-clone-1.C: New. * g++.dg/gomp/target-simd-clone-2.C: New. * gcc.dg/gomp/target-simd-clone-1.c: New. * gcc.dg/gomp/target-simd-clone-2.c: New. * gcc.dg/gomp/target-simd-clone-3.c: New. * gcc.dg/gomp/target-simd-clone-4.c: New. * gcc.dg/gomp/target-simd-clone-5.c: New. * gcc.dg/gomp/target-simd-clone-6.c: New. * gcc.dg/gomp/target-simd-clone-7.c: New. * gcc.dg/gomp/target-simd-clone-8.c: New. * lib/scanoffloadipa.exp: New. libgomp/ChangeLog: * testsuite/lib/libgomp.exp: Load scanoffloadipa.exp library. * testsuite/libgomp.c/target-simd-clone-1.c: New. * testsuite/libgomp.c/target-simd-clone-2.c: New. * testsuite/libgomp.c/target-simd-clone-3.c: New.
2022-10-17Fix nvptx-specific '-foffload-options' syntax in ↵Thomas Schwinge1-1/+1
'libgomp.c/reverse-offload-sm30.c' That is, '-mptx=_' is only valid in '-foffload-options=nvptx-none', too. Fix test case added in recent commit r13-2625-g6b43f556f392a7165582aca36a19fe7389d995b2 "nvptx/mkoffload.cc: Warn instead of error when reverse offload is not possible". libgomp/ * testsuite/libgomp.c/reverse-offload-sm30.c: Fix nvptx-specific '-foffload-options' syntax.
2022-10-12libgomp: Fix up creation of artificial teamsJakub Jelinek2-0/+40
When not in explicit parallel/target/teams construct, we in some cases create an artificial parallel with a single thread (either to handle target nowait or for task reduction purposes). In those cases, it handled again artificially created implicit task (created by gomp_new_icv for cases where we needed to write to some ICVs), but as the testcases show, didn't take into account possibility of this being done from explicit task(s). The code would destroy/free the previous task and replace it with the new implicit task. If task is an explicit task (when teams is NULL, all explicit tasks behave like if (0)), it is a pointer to a local stack variable, so freeing it doesn't work, and additionally we shouldn't lose the explicit tasks - the new implicit task should instead replace the ancestor task which is the first implicit one. 2022-10-12 Jakub Jelinek <jakub@redhat.com> * task.c (gomp_create_artificial_team): Fix up handling of invocations from within explicit task. * target.c (GOMP_target_ext): Likewise. * testsuite/libgomp.c/task-7.c: New test. * testsuite/libgomp.c/task-8.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-17.c: New test. * testsuite/libgomp.c-c++-common/task-reduction-18.c: New test.
2022-09-12nvptx/mkoffload.cc: Warn instead of error when reverse offload is not possibleTobias Burnus1-0/+15
Reverse offload requests at least -misa=sm_35; with this patch, a warning instead of an error is shown, still permitting reverse offload for all other configured device types. This is achieved by not calling GOMP_offload_register_ver (and stopping generating pointless 'static const char' variables, once known.) The tool_name as progname changes adds "nvptx " and "gcn " to the "mkoffload: warning/error:" diagnostic. gcc/ChangeLog: * config/nvptx/mkoffload.cc (process): Replace a fatal_error by a warning + not enabling offloading if -misa=sm_30 prevents reverse offload. (main): Use tool_name as progname for diagnostic. * config/gcn/mkoffload.cc (main): Likewise. libgomp/ChangeLog: * libgomp.texi (Offload-Target Specifics: nvptx): Document that reverse offload requires >= -march=sm_35. * testsuite/libgomp.c-c++-common/requires-4.c: Build for nvptx with -misa=sm_35. * testsuite/libgomp.c-c++-common/requires-5.c: Likewise. * testsuite/libgomp.c-c++-common/requires-6.c: Likewise. * testsuite/libgomp.c-c++-common/reverse-offload-1.c: Likewise. * testsuite/libgomp.fortran/reverse-offload-1.f90: Likewise. * testsuite/libgomp.c/reverse-offload-sm30.c: New test.
2022-09-08openmp: Implement doacross(sink: omp_cur_iteration - 1)Jakub Jelinek4-0/+888
This patch implements doacross(sink: omp_cur_iteration - 1) that the previous patchset emitted a sorry on during omp expansion. It can be implemented with existing library functions. To recap, depend(source)/doacross(source:)/doacross(source:omp_cur_iteration) is implemented calling GOMP_doacross_post or GOMP_doacross_ull_post, called with an array of long or unsigned long long elements, one for all collapsed loops together and one for each further ordered loop if any. We initialize that array in each thread when grabbing further set of iterations and update it at the end of loops, so that it represents the current iteration (as 0 based counters). When the worksharing loop is created, we tell the library through another similar array the counts (the loop needs to be rectangular) in each dimension, first element is count of all logical iterations in the collapsed loops. depend(sink:v1 op N1, v2 op N2, ...) is then implemented by conditionally calling GOMP_doacross_wait/GOMP_doacross_ull_wait. For N? of 0 there is no check, otherwise if it wants to wait in a particular dimension for a previous iteration, we check that the corresponding iterator isn't the first one (or first few), where the previous iterator in that dimension would be out of range, and similarly for checking of next iteration in a dimension that it isn't the last one (or last few) where it would be similarly out of bounds. Then the collapsed loop counters are folded into a single 0 based counter (first argument) and then other 0 based iterations counters on what iteration it should wait for. Now, doacross(sink: omp_cur_iteration - 1) is supposed to wait for the previous logical iteration in the combined iteration space of all ordered loops. For the very first iteration in that combined iteration space it does nothing, there is no previous iteration. And similarly it does nothing if there are more ordered loops than collapsed loop and it isn't the first logical iteration of the combined loops inside of the collapsed loops, because as implemented we know the previous iteration in that case is always executed by the same thread as the current one. In the implementation, we use the same value as is stored in the first element of the array for GOMP_doacross_post/GOMP_doacross_ull_post, if that value is 0, we do nothing. The rest is different based on if ordered argument is equal to collapse or not. If it is, then we otherwise call GOMP_doacross_wait/GOMP_doacross_ull_wait with a single argument, one less than that counter we compare against 0. If ordered argument is bigger than collapse, we add a per-thread boolean variable .first.N, which we set to true at the start of the outermost ordered loop inside of the collapsed set of loops and set to false at the end of the innermost ordered loop. If .first.N is false, we don't do anything (we know the previous iteration was handled by the current thread and by my reading of the spec we don't need to emit even a memory barrier in that case, because it is just synchronization with the same thread), otherwise we call GOMP_doacross_wait/GOMP_doacross_ull_wait with the first argument one less than the counter we compare against 0, and then one less than 2nd and following counts if iterations we pass to the workshare initialization. If say .counts.N passed to the workshare initialization is { 256, 13, 5, 2 } for collapse(3) ordered(6) loop, then GOMP_doacross_post/GOMP_doacross_ull_post is called with arguments equal to .ordereda.N[0] - 1, 12, 4, 1. 2022-09-08 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-expand.cc (expand_omp_ordered_sink): Add CONT_BB argument. Add doacross(sink:omp_cur_iteration-1) support. (expand_omp_ordered_source_sink): Clear counts[fd->ordered + 1]. Adjust expand_omp_ordered_sink caller. (expand_omp_for_ordered_loops): If counts[fd->ordered + 1] is non-NULL, set that variable to true at the start of outermost non-collapsed loop and set it to false at the end of innermost ordered loop. (expand_omp_for_generic): If fd->ordered, allocate 1 + (fd->ordered - fd->collapse) further elements in counts array. Copy to counts + 2 + fd->ordered the counts of fd->collapse .. fd->ordered - 1 loop if any. gcc/testsuite/ * c-c++-common/gomp/doacross-7.c: New test. libgomp/ * libgomp.texi (OpenMP 5.2): Mention that omp_cur_iteration is now fully supported. * testsuite/libgomp.c/doacross-4.c: New test. * testsuite/libgomp.c/doacross-5.c: New test. * testsuite/libgomp.c/doacross-6.c: New test. * testsuite/libgomp.c/doacross-7.c: New test.
2022-08-17OpenMP: Fix var replacement with 'simd' and linear-step vars [PR106548]Tobias Burnus1-0/+254
gcc/ChangeLog: PR middle-end/106548 * omp-low.cc (lower_rec_input_clauses): Use build_outer_var_ref for 'simd' linear-step values that are variable. libgomp/ChangeLog: PR middle-end/106548 * testsuite/libgomp.c/linear-2.c: New test.
2022-06-21libgomp: Fix up target-31.c test [PR106045]Jakub Jelinek1-1/+1
The i variable is used inside of the parallel in: #pragma omp simd safelen(32) private (v) for (i = 0; i < 64; i++) { v = 3 * i; ll[i] = u1 + v * u2[0] + u2[1] + x + y[0] + y[1] + v + h[0] + u3[i]; } where i is predetermined linear (so while inside of the body it is safe, private per SIMD lane var) the final value is written to the shared variable, and in: for (i = 0; i < 64; i++) if (ll[i] != u1 + 3 * i * u2[0] + u2[1] + x + y[0] + y[1] + 3 * i + 13 + 14 + i) #pragma omp atomic write err = 1; which is a normal loop and so it isn't in any way privatized there. So we have a data race, fixed by adding private (i) clause to the parallel. 2022-06-21 Jakub Jelinek <jakub@redhat.com> Paul Iannetta <piannetta@kalrayinc.com> PR libgomp/106045 * testsuite/libgomp.c/target-31.c: Add private (i) clause.
2022-06-13openmp: Conforming device numbers and omp_{initial,invalid}_deviceJakub Jelinek4-3/+63
OpenMP 5.2 changed once more what device numbers are allowed. In 5.1, valid device numbers were [0, omp_get_num_devices()]. 5.2 makes also -1 valid (calls it omp_initial_device), which is equivalent in behavior to omp_get_num_devices() number but has the advantage that it is a constant. And it also introduces omp_invalid_device which is also a constant with implementation defined value < -1. That value should act like sNaN, any time any device construct (GOMP_target*) or OpenMP runtime API routine is asked for such a device, the program is terminated. And if OMP_TARGET_OFFLOAD=mandatory, all non-conforming device numbers (which is all but [-1, omp_get_num_devices()] other than omp_invalid_device) must be treated like omp_invalid_device. For device constructs, we have a compatibility problem, we've historically used 2 magic negative values to mean something special. GOMP_DEVICE_ICV (-1) means device clause wasn't present, pick the omp_get_default_device () number GOMP_DEVICE_FALLBACK (-2) means the host device (this is used e.g. for #pragma omp target if (cond) where if cond is false, we pass -2 But 5.2 requires that omp_initial_device is -1 (there were discussions about it, advantage of -1 is that one can say iterate over the [-1, omp_get_num_devices()-1] range to get all devices starting with the host/initial one. And also, if user passes -2, unless it is omp_invalid_device, we need to treat it like non-conforming with OMP_TARGET_OFFLOAD=mandatory. So, the patch does on the compiler side some number remapping, user_device_num >= -2U ? user_device_num - 1 : user_device_num. This remapping is done at compile time if device clause has constant argument, otherwise at runtime, and means that for user -1 (omp_initial_device) we pass -2 to GOMP_* in the runtime library where it treats it like host fallback, while -2 is remapped to -3 (one of the non-conforming device numbers, for those it doesn't matter which one is which). omp_invalid_device is then -4. For the OpenMP device runtime APIs, no remapping is done. This patch doesn't deal with the initial default-device-var for OMP_TARGET_OFFLOAD=mandatory , the spec says that the inital ICV value for that should in that case depend on whether there are any offloading devices or not (if not, should be omp_invalid_device), but that means we can't determine the number of devices lazily (and let libraries have the possibility to register their offloading data etc.). 2022-06-13 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-expand.cc (expand_omp_target): Remap user provided device clause arguments, -1 to -2 and -2 to -3, either at compile time if constant, or at runtime. include/ * gomp-constants.h (GOMP_DEVICE_INVALID): Define. libgomp/ * omp.h.in (omp_initial_device, omp_invalid_device): New enumerators. * omp_lib.f90.in (omp_initial_device, omp_invalid_device): New parameters. * omp_lib.h.in (omp_initial_device, omp_invalid_device): Likewise. * target.c (resolve_device): Add remapped argument, handle GOMP_DEVICE_ICV only if remapped is true (and clear remapped), for negative values, treat GOMP_DEVICE_FALLBACK as fallback only if remapped, otherwise treat omp_initial_device that way. For omp_invalid_device, always emit gomp_fatal, even when OMP_TARGET_OFFLOAD isn't mandatory. (GOMP_target, GOMP_target_ext, GOMP_target_data, GOMP_target_data_ext, GOMP_target_update, GOMP_target_update_ext, GOMP_target_enter_exit_data): Pass true as remapped argument to resolve_device. (omp_target_alloc, omp_target_free, omp_target_is_present, omp_target_memcpy_check, omp_target_associate_ptr, omp_target_disassociate_ptr, omp_get_mapped_ptr, omp_target_is_accessible): Pass false as remapped argument to resolve_device. Treat omp_initial_device the same as gomp_get_num_devices (). Don't bypass resolve_device calls if device_num is negative. (omp_pause_resource): Treat omp_initial_device the same as gomp_get_num_devices (). Call resolve_device. * icv-device.c (omp_set_default_device): Always set to device_num even when it is negative. * libgomp.texi: Document that Conforming device numbers, omp_initial_device and omp_invalid_device is implemented. * testsuite/libgomp.c/target-41.c (main): Add test with omp_initial_device. * testsuite/libgomp.c/target-45.c: New test. * testsuite/libgomp.c/target-46.c: New test. * testsuite/libgomp.c/target-47.c: New test. * testsuite/libgomp.c-c++-common/target-is-accessible-1.c (main): Add test with omp_initial_device. Use -5 instead of -1 for negative value test. * testsuite/libgomp.fortran/target-is-accessible-1.f90 (main): Likewise. Reorder stop numbers.
generated by cgit v1.1 at 2025-05-23 05:28:04 +0000