aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite
AgeCommit message (Collapse)AuthorFilesLines
2019-05-27omp-low.c (lower_omp_1): Look through ordered...Jakub Jelinek1-0/+57
* omp-low.c (lower_omp_1) <case GIMPLE_ASSIGN>: Look through ordered, critical, taskgroup and section regions when looking for a region with non-NULL lastprivate_conditional_map. * testsuite/libgomp.c-c++-common/lastprivate-conditional-3.c: New test. From-SVN: r271672
2019-05-24tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__CONDTEMP_.Jakub Jelinek2-0/+315
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__CONDTEMP_. * tree.h (OMP_CLAUSE_DECL): Use OMP_CLAUSE__CONDTEMP_ instead of OMP_CLAUSE__REDUCTEMP_. * tree.c (omp_clause_num_ops, omp_clause_code_name): Add OMP_CLAUSE__CONDTEMP_. (walk_tree_1): Handle OMP_CLAUSE__CONDTEMP_. * tree-pretty-print.c (dump_omp_clause): Likewise. * tree-nested.c (convert_nonlocal_omp_clauses, convert_local_omp_clauses): Likewise. * gimplify.c (enum gimplify_omp_var_data): Use hexadecimal constants instead of decimal. Add GOVD_LASTPRIVATE_CONDITIONAL. (gimplify_scan_omp_clauses): Don't reject lastprivate conditional on OMP_FOR. (gimplify_omp_for): Warn and disable conditional modifier from lastprivate on loop iterators. * omp-general.h (struct omp_for_data): Add lastprivate_conditional member. * omp-general.c (omp_extract_for_data): Initialize it. * omp-low.c (struct omp_context): Add lastprivate_conditional_map member. (delete_omp_context): Delete it. (lower_lastprivate_conditional_clauses): New function. (lower_lastprivate_clauses): Add BODY_P and CSTMT_LIST arguments, handle lastprivate conditional clauses. (lower_reduction_clauses): Add CLIST argument, emit it into the critical section if any. (lower_omp_sections): Adjust lower_lastprivate_clauses and lower_reduction_clauses callers. (lower_omp_for_lastprivate): Add CLIST argument, pass it through to lower_lastprivate_clauses. (lower_omp_for): Call lower_lastprivate_conditional_clauses, adjust lower_omp_for_lastprivate and lower_reduction_clauses callers, emit clist into a critical section if not emitted there already by lower_reduction_clauses. (lower_omp_taskreg, lower_omp_teams): Adjust lower_reduction_clauses callers. (lower_omp_1): Handle GIMPLE_ASSIGNs storing into lastprivate conditional variables. * omp-expand.c (determine_parallel_type): Punt if OMP_CLAUSE__CONDTEMP_ clause is present. (expand_omp_for_generic, expand_omp_for_static_nochunk, expand_omp_for_static_chunk): Handle lastprivate conditional. (expand_omp_for): Handle fd.lastprivate_conditional like fd.have_reductemp. gcc/testsuite/ * c-c++-common/gomp/lastprivate-conditional-2.c (foo): Don't expect sorry for omp for. * c-c++-common/gomp/lastprivate-conditional-3.c: New test. libgomp/ * testsuite/libgomp.c-c++-common/lastprivate-conditional-1.c: New test. * testsuite/libgomp.c-c++-common/lastprivate-conditional-2.c: New test. From-SVN: r271610
2019-05-17OpenACC Profiling Interface (incomplete)Thomas Schwinge6-0/+1919
libgomp/ * acc_prof.h: New file. * oacc-profiling.c: Likewise. * Makefile.am (nodist_libsubinclude_HEADERS, libgomp_la_SOURCES): Add these, respectively. * Makefile.in: Regenerate. * env.c (initialize_env): Call goacc_profiling_initialize. * oacc-plugin.c (GOMP_PLUGIN_goacc_thread) (GOMP_PLUGIN_goacc_profiling_dispatch): New functions. * oacc-plugin.h (GOMP_PLUGIN_goacc_thread) (GOMP_PLUGIN_goacc_profiling_dispatch): Declare. * libgomp.map (OACC_2.5.1): Add acc_prof_lookup, acc_prof_register, acc_prof_unregister, and acc_register_library. (GOMP_PLUGIN_1.3): Add GOMP_PLUGIN_goacc_profiling_dispatch, and GOMP_PLUGIN_goacc_thread. * oacc-int.h (struct goacc_thread): Add prof_info, api_info, prof_callbacks_enabled members. (goacc_prof_enabled, goacc_profiling_initialize) (_goacc_profiling_dispatch_p, _goacc_profiling_setup_p) (goacc_profiling_dispatch): Declare. (GOACC_PROF_ENABLED, GOACC_PROFILING_DISPATCH_P) (GOACC_PROFILING_SETUP_P): Define. * oacc-async.c (acc_async_test, acc_async_test_all, acc_wait) (acc_wait_async, acc_wait_all, acc_wait_all_async): Update for OpenACC Profiling Interface. * oacc-cuda.c (acc_get_current_cuda_device) (acc_get_current_cuda_context, acc_get_cuda_stream) (acc_set_cuda_stream): Likewise. * oacc-init.c (acc_init_1, goacc_attach_host_thread_to_device) (acc_init, acc_set_device_type, acc_get_device_type) (acc_get_device_num, goacc_lazy_initialize): Likewise. * oacc-mem.c (acc_malloc, acc_free, memcpy_tofrom_device) (acc_deviceptr, acc_hostptr, acc_is_present, acc_map_data) (acc_unmap_data, present_create_copy, delete_copyout) (update_dev_host): Likewise. * oacc-parallel.c (GOACC_parallel_keyed, GOACC_data_start) (GOACC_data_end, GOACC_enter_exit_data, GOACC_update, GOACC_wait): Likewise. * plugin/plugin-nvptx.c (nvptx_exec, nvptx_alloc, nvptx_free) (GOMP_OFFLOAD_openacc_exec, GOMP_OFFLOAD_openacc_async_exec): Likewise. * libgomp.texi: Update. * testsuite/libgomp.oacc-c-c++-common/acc_prof-dispatch-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_prof-valid_bytes-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_prof-version-1.c: Likewise. From-SVN: r271346
2019-05-08Address compiler diagnostics in libgomp.oacc-c-c++-common/pr87835.cThomas Schwinge1-3/+2
source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c: In function 'main': source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c:45: warning: ignoring #pragma loop gang [-Wunknown-pragmas] 45 | #pragma loop gang | source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr87835.c:19:7: warning: unused variable 'b' [-Wunused-variable] 19 | int b[n]; | ^ libgomp/ PR target/87835 * testsuite/libgomp.oacc-c-c++-common/pr87835.c: Update. From-SVN: r271004
2019-02-22[libgomp] In OpenACC testing, by default only build for the offload target ↵Thomas Schwinge7-20/+23
that we're actually going to test ... to avoid compilation overhead, and to keep simple '-foffload=[...]' handling in test cases. libgomp/ * testsuite/libgomp.oacc-c++/c++.exp: Specify "-foffload=$offload_target". * testsuite/libgomp.oacc-c/c.exp: Likewise. * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. * testsuite/lib/libgomp.exp (check_effective_target_openacc_nvidia_accel_configured): Remove, as (conceptually) merged into check_effective_target_openacc_nvidia_accel_selected. Adjust all users. From-SVN: r269109
2019-02-22[libgomp] In OpenACC testing, cycle though all offload targetsThomas Schwinge6-46/+60
... instead of through offload plugins. libgomp/ * plugin/configfrag.ac: Populate and AC_SUBST offload_targets. * testsuite/libgomp-test-support.exp.in: Adjust. * testsuite/lib/libgomp.exp: Likewise. Don't populate openacc_device_types_s. (offload_target_to_openacc_device_type): New proc. * testsuite/libgomp.oacc-c++/c++.exp: Adjust. * testsuite/libgomp.oacc-c/c.exp: Likewise. * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. * Makefile.in: Regenerate. * configure: Likewise. * testsuite/Makefile.in: Likewise. From-SVN: r269108
2019-02-22[libgomp] Clarify difference between offload target, offload plugin, and ↵Thomas Schwinge6-57/+52
OpenACC device type libgomp/ * plugin/configfrag.ac: Populate and AC_SUBST offload_plugins instead of offload_targets, and AC_DEFINE_UNQUOTED OFFLOAD_PLUGINS instead of OFFLOAD_TARGETS. * target.c (gomp_target_init): Adjust. * testsuite/libgomp-test-support.exp.in: Likewise. * testsuite/lib/libgomp.exp: Likewise. Populate openacc_device_types_s instead of offload_targets_s_openacc. (check_effective_target_openacc_nvidia_accel_selected) (check_effective_target_openacc_host_selected): Adjust. * testsuite/libgomp.oacc-c++/c++.exp: Likewise. * testsuite/libgomp.oacc-c/c.exp: Likewise. * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. * Makefile.in: Regenerate. * config.h.in: Likewise. * configure: Likewise. * testsuite/Makefile.in: Likewise. From-SVN: r269107
2019-02-22[libgomp] In OpenACC offloading testing, be more explicit in what is ↵Thomas Schwinge4-8/+11
supported, and what is not, or why not libgomp/ * testsuite/lib/libgomp.exp: Error out for unknown offload target. * testsuite/libgomp.oacc-c++/c++.exp: Likewise. Report if "offloading: supported, but hardware not accessible". * testsuite/libgomp.oacc-c/c.exp: Likewise. * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. From-SVN: r269106
2019-01-31re PR c++/88988 (ICE: Segmentation fault (in lookup_name_real_1))Jakub Jelinek1-0/+28
PR c++/88988 * lambda.c (is_capture_proxy): Don't return true for DECL_OMP_PRIVATIZED_MEMBER artificial vars. * testsuite/libgomp.c++/pr88988.C: New test. From-SVN: r268407
2019-01-28re PR middle-end/89002 (ICE in scan_omp_1_op, at omp-low.c:3166)Jakub Jelinek1-43/+0
PR middle-end/89002 * gimplify.c (gimplify_omp_for): When adding OMP_CLAUSE_*_GIMPLE_SEQ for lastprivate/linear IV, push gimplify context around gimplify_assign and, if it needed any temporaries, pop it into a gimple bind around the sequence. * testsuite/libgomp.c/pr89002.c: New test. From-SVN: r268346
2019-01-28re PR middle-end/89002 (ICE in scan_omp_1_op, at omp-low.c:3166)Jakub Jelinek1-0/+86
PR middle-end/89002 * gimplify.c (gimplify_omp_for): When adding OMP_CLAUSE_*_GIMPLE_SEQ for lastprivate/linear IV, push gimplify context around gimplify_assign and, if it needed any temporaries, pop it into a gimple bind around the sequence. * testsuite/libgomp.c/pr89002.c: New test. From-SVN: r268345
2019-01-28re PR testsuite/89064 (libgomp.graphite/force-parallel-5.c fails starting ↵Richard Biener1-1/+1
with r268257) 2019-01-28 Richard Biener <rguenther@suse.de> PR testsuite/89064 PR tree-optimization/86865 * testsuite/libgomp.graphite/force-parallel-5.c: XFAIL. From-SVN: r268333
2019-01-23[nvptx, libgomp] Fix cuMemAlloc with size zeroTom de Vries1-0/+15
Consider test-case: ... int main (void) { #pragma acc parallel async ; #pragma acc parallel async ; #pragma acc wait return 0; } ... This fails with: ... libgomp: cuMemAlloc error: invalid argument Segmentation fault (core dumped) ... The cuMemAlloc error is due to the fact that we're try to allocate 0 bytes. Fix this by preventing calling map_push with size zero argument in nvptx_exec. This also has the consequence that for the abort-1.c test-case, we end up calling cuMemFree during map_fini for the struct cuda_map allocated in map_init, which fails because an abort happened. Fix this by calling cuMemFree with CUDA_CALL_NOCHECK in cuda_map_destroy. 2019-01-23 Tom de Vries <tdevries@suse.de> PR target/PR88946 * plugin/plugin-nvptx.c (cuda_map_destroy): Use CUDA_CALL_NOCHECK for cuMemFree. (nvptx_exec): Don't call map_push if mapnum == 0. * testsuite/libgomp.oacc-c-c++-common/pr88946.c: New test. From-SVN: r268178
2019-01-23[nvptx, libgomp] Fix assert (!s->map->active) in map_finiTom de Vries1-0/+15
There are currently two situations where this assert triggers: ... libgomp/plugin/plugin-nvptx.c: map_fini: Assertion `!s->map->active' failed. ... First, in abort-1.c, a parallel region triggering an abort: ... int main (void) { #pragma acc parallel abort (); return 0; } ... The abort is detected in nvptx_exec as the CUDA_ERROR_ILLEGAL_INSTRUCTION return status of the cuStreamSynchronize call after kernel launch, which is then handled by calling non-returning function GOMP_PLUGIN_fatal. Consequently, the map_pop in nvptx_exec that in case of cuStreamSynchronize success would remove or inactive the element added by the map_push earlier in nvptx_exec, does not trigger. With the element no longer active, but still marked active and a member of s->map, we run into the assert during GOMP_OFFLOAD_fini_device, which is triggered from atexit handler gomp_target_fini (which is triggered by the GOMP_PLUGIN_fatal mentioned above calling exit). Second, in pr88941.c, an async parallel region without wait: ... int main (void) { #pragma acc parallel async ; /* no #pragma acc wait */ return 0; } ... Because nvptx_exec is handling an async region, it does not call map_pop for the element added by map_push, but schedules an kernel execution completion event to call map_pop. Again, we run into the assert during GOMP_OFFLOAD_fini_device, which is triggered from atexit handler gomp_target_fini, but the exit in this case is triggered by returning from main. So either the kernel is still running, or the kernel has completed but the corresponding event that is supposed to call map_pop is stuck in the event queue, waiting for an event_gc. Fix this by removing the assert, and skipping the freeing of device memory if the map is still marked active (though in the async case, this is more a workaround than an fix). 2019-01-23 Tom de Vries <tdevries@suse.de> PR target/88941 PR target/88939 * plugin/plugin-nvptx.c (cuda_map_destroy): Handle map->active case. (map_fini): Remove "assert (!s->map->active)". * testsuite/libgomp.oacc-c-c++-common/pr88941.c: New test. From-SVN: r268177
2019-01-23[nvptx, libgomp] Fix map_pushTom de Vries1-0/+62
The map field of a struct ptx_stream is a FIFO. The FIFO is implemented as a single linked list, with pop-from-the-front semantics. The function map_pop pops an element, either by: - deallocating the element, if there is more than one element - or marking the element inactive, if there's only one element The responsibility of map_push is to push an element to the back, as well as selecting the element to push, by: - allocating an element, or - reusing the element at the front if inactive and big enough, or - dropping the element at the front if inactive and not big enough, and allocating one that's big enough The current implemention gets at least the first and most basic scenario wrong: > map = cuda_map_create (size); We create an element, and assign it to map. > for (t = s->map; t->next != NULL; t = t->next) > ; We determine the last element in the fifo. > t->next = map; We append the new element. > s->map = map; But here, we throw away the rest of the FIFO, and declare the FIFO to be just the new element. This problem causes the test-case asyncwait-1.c to fail intermittently on some systems. The pr87835.c test-case added here is a a minimized and modified version of asyncwait-1.c (avoiding the kernel construct) that is more likely to fail. Fix this by rewriting map_pop more robustly, by: - seperating the function in two phases: select element, push element - when reusing or dropping an element, making sure that the element is cleanly popped from the queue - rewriting the push element part in such a way that it can handle all cases without needing if statements, such that each line is exercised for each of the three cases. 2019-01-23 Tom de Vries <tdevries@suse.de> PR target/87835 * plugin/plugin-nvptx.c (map_push): Fix adding of allocated element. * testsuite/libgomp.oacc-c-c++-common/pr87835.c: New test. From-SVN: r268176
2019-01-15[nvptx] Handle assignment to gang-level reduction variableTom de Vries1-0/+16
2019-01-15 Tom de Vries <tdevries@suse.de> PR target/80547 * config/nvptx/nvptx.c (nvptx_goacc_reduction_init): Handle lhs == NULL_TREE for gang-level reduction. * testsuite/libgomp.oacc-c-c++-common/gang-reduction-var-assignment.c: New test. From-SVN: r267934
2019-01-12[nvptx] Enable setting vector length using -fopenacc-dim -- testcasesTom de Vries4-0/+212
Add some test-cases that set vector length using -fopenacc-dim. 2019-01-12 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c: New test. * testsuite/libgomp.oacc-fortran/gemm-2.f90: New test. From-SVN: r267897
2019-01-12[nvptx] Add vector_length 64 test-casesTom de Vries3-0/+55
Add some test-cases using vector_length 64. 2019-01-12 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vector-length-64-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-64-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-64-3.c: New test. From-SVN: r267895
2019-01-12[nvptx] Force vl32 if calling vector-partitionable routines -- test-casesTom de Vries2-0/+105
Add test-cases for "[nvptx] Force vl32 if calling vector-partitionable routines". 2019-01-12 Tom de Vries <tdevries@suse.de> PR target/85486 * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr85486.c: New test. From-SVN: r267894
2019-01-12[nvptx] Don't emit barriers for empty loops -- test-casesTom de Vries2-0/+42
Add test-cases for PR85381. 2019-01-12 Tom de Vries <tdevries@suse.de> PR target/85381 * testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr85381.c: New test. From-SVN: r267893
2019-01-12[nvptx] Enable large vectors -- reduction testcasesTom de Vries3-0/+173
Add various reduction test-cases with vector length 128. 2019-01-12 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vred2d-128.c: New test. * testsuite/libgomp.oacc-fortran/gemm.f90: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-10.c: New test. From-SVN: r267892
2019-01-12[nvptx] Enable large vectors -- test-casesTom de Vries3-0/+121
Add various test-cases with vector length 128. 2019-01-12 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c: New test. From-SVN: r267891
2019-01-12[nvptx] Enable large vectorsTom de Vries2-5/+4
Allow vector_length clauses to accept values larger than warp size. Note that this does not enable setting vector_length to values larger than warp size using -fopenacc-dim. 2019-01-12 Tom de Vries <tdevries@suse.de> * config/nvptx/nvptx.c (nvptx_goacc_validate_dims): Take larger vector lengths into account. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Expect vector length to be 128. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Expect vector length 2097152 to be reduced to 1024 instead of 32. From-SVN: r267889
2019-01-11[libgomp, testsuite, openacc] Remove -foffload=-w in reduction-[1-5].cTom de Vries5-15/+0
Before the commit "[libgomp, testsuite, openacc] Don't use const int for dimensions", the "const int" construct was used to set launch dimensions in reductions-[1-5].c. In the case of -xc -O0, the const int is implemented as a variable by the C front-end. Consequently, the nvptx back-end generated warnings that vector_length was overridden to be hard-coded, rather than left to be set at runtime. The test-cases silenced these warnings by switching off all warnings in the accelerator compiler using "-foffload=-w". Given that no warnings occur anymore, remove the "-foffload=-w" setting. 2019-01-11 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/reduction-1.c: Remove -foffload=-w. * testsuite/libgomp.oacc-c-c++-common/reduction-2.c: Same. * testsuite/libgomp.oacc-c-c++-common/reduction-3.c: Same. * testsuite/libgomp.oacc-c-c++-common/reduction-4.c: Same. * testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Same. From-SVN: r267836
2019-01-11[nvptx, testsuite, openacc, libgomp] Add insufficient-resources.cTom de Vries1-0/+21
Add a test-case that tests the "insufficient resources" fatal in the nvptx libgomp plugin. 2019-01-11 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/insufficient-resources.c: New test. From-SVN: r267835
2019-01-10Add testcase from PR71959Nathan Sidwell2-0/+66
libgomp/ PR lto/71959 * testsuite/libgomp.oacc-c++/pr71959-aux.cc: New. * testsuite/libgomp.oacc-c++/pr71959.C: New. Co-Authored-By: Julian Brown <julian@codesourcery.com> From-SVN: r267806
2019-01-09[libgomp, testsuite, openacc] Don't use const int for dimensionsTom de Vries5-15/+15
Const int is handled differently at -O0 for -xc and -xc++, which can cause noise in testsuite/libgomp.oacc-c-c++-common test-cases (which are both run for c and c++) if const int is used for launch dimensions. Fix this by using #defines instead. 2019-01-09 Tom de Vries <tdevries@suse.de> PR target/88756 * testsuite/libgomp.oacc-c-c++-common/reduction-1.c (ng, nw, vl): Use #define instead of "const int". * testsuite/libgomp.oacc-c-c++-common/reduction-2.c (ng, nw, vl): Same. * testsuite/libgomp.oacc-c-c++-common/reduction-3.c (ng, nw, vl): Same. * testsuite/libgomp.oacc-c-c++-common/reduction-4.c (ng, nw, vl): Same. * testsuite/libgomp.oacc-c-c++-common/reduction-5.c (ng, nw, vl): Same. From-SVN: r267747
2019-01-07[nvptx] Fix libgomp.oacc-c-c++-common/vector-length-128-3.cTom de Vries1-1/+1
The vector-length-128-3.c test-case uses GOMP_OPENACC_DIM=-:-:128, but '-' is not yet supported on trunk. Use GOMP_OPENACC_DIM=::128 instead. 2019-01-07 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: Fix GOMP_OPENACC_DIM argument. From-SVN: r267624
2019-01-03[nvptx] Add vector_length 128 testcasesTom de Vries2-0/+81
Add a couple of test-cases using vector length 128, while checking that we override to vector length 32. 2019-01-03 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: New test. From-SVN: r267559
2019-01-01Update copyright years.Jakub Jelinek4-4/+4
From-SVN: r267494
2018-12-27aligned1.f03: Fix invalid code that now causes an error after r267415.Steven G. Kargl1-2/+2
2018-12-27 Steven G. Kargl <kargl@gcc.gnu.org> * libgomp.fortran/aligned1.f03: Fix invalid code that now causes an error after r267415. From-SVN: r267436
2018-12-19[nvptx] Commit passing pr85381-*.c test-casesTom de Vries3-0/+98
Add pr85381*.c test-cases that are already passing without the fix for PR85381. Build and reg-tested on x86_64 with nvptx accelerator. 2018-12-19 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: New test. From-SVN: r267268
2018-12-19[nvptx, libgomp] Move rtl-dump test-cases to libgompTom de Vries3-0/+66
The goacc.exp test-cases nvptx-merged-loop.c and nvptx-sese-1.c are failing during linking due to missing libgomp.spec. Move them to the libgomp testsuite. Build and reg-tested on x86_64 with nvptx accelerator. 2018-12-19 Tom de Vries <tdevries@suse.de> * gcc.dg/goacc/nvptx-merged-loop.c: Move to libgomp/testsuite/libgomp.oacc-c-c++-common. * gcc.dg/goacc/nvptx-sese-1.c: Same. * testsuite/lib/libgomp.exp: Add load_lib of scanoffloadrtl.exp. * testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c: Move from gcc/testsuite/gcc.dg/goacc. * testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c: Same. From-SVN: r267267
2018-12-14Missing changes from "Adjust copy/copyin/copyout/create for OpenACC 2.5"Thomas Schwinge16-271/+0
Most of that patch's changes were already committed as part of r261813 "Update OpenACC data clause semantics to the 2.5 behavior", but not all of them. libgomp/ * oacc-mem.c (acc_present_or_create): Remove definition and change to alias of acc_create. (acc_present_or_copyin): Remove definition and change to alias of acc_copyin. * oacc-parallel.c (GOACC_enter_exit_data): Call acc_create instead of acc_present_or_create. * testsuite/libgomp.oacc-c-c++-common/data-already-1.c: Remove. * testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise. * testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise. Co-Authored-By: Chung-Lin Tang <cltang@codesourcery.com> From-SVN: r267153
2018-12-14[PR88495] An OpenACC async queue is always synchronized with itselfThomas Schwinge2-138/+0
An OpenACC async queue is always synchronized with itself, so invocations like "#pragma acc wait(0) async(0)", or "acc_wait_async (0, 0)" don't make a lot of sense, but are still valid. libgomp/ PR libgomp/88495 * plugin/plugin-nvptx.c (nvptx_wait_async): Don't refuse "identical parameters". * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/lib-80.c: Remove. From-SVN: r267152
2018-12-14[PR88484] OpenACC wait directive without wait argument but with async clauseThomas Schwinge1-0/+78
We don't correctly handle "#pragma acc wait async (a)" for "a >= 0", handling as a no-op whereas it should enqueue the appropriate wait operations on "async (a)". libgomp/ PR libgomp/88484 * oacc-parallel.c (GOACC_wait): Correct handling for "async >= 0". * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: New file. From-SVN: r267151
2018-12-14[PR88407] [OpenACC] Correctly handle unseen async-argumentsThomas Schwinge9-260/+74
... which turn the operation into a no-op. libgomp/ PR libgomp/88407 * plugin/plugin-nvptx.c (nvptx_async_test, nvptx_wait) (nvptx_wait_async): Unseen async-argument is a no-op. * testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-79.c: Likewise. * testsuite/libgomp.oacc-fortran/lib-12.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-71.c: Merge into... * testsuite/libgomp.oacc-c-c++-common/lib-69.c: ... this. Update. * testsuite/libgomp.oacc-c-c++-common/lib-77.c: Merge into... * testsuite/libgomp.oacc-c-c++-common/lib-74.c: ... this. Update From-SVN: r267150
2018-12-14Revise libgomp.oacc-c-c++-common/data-2-lib.c, ↵Thomas Schwinge2-157/+120
libgomp.oacc-c-c++-common/data-2.c These are meant to be functionally equivalent (but no longer are), just using different means. Also, use the OpenACC "*_async" functions recently added. libgomp/ * testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Revise. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise. From-SVN: r267149
2018-12-14Correctly describe OpenACC async/wait dependenciesChung-Lin Tang3-3/+3
libgomp/ * testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Adjust. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com> From-SVN: r267148
2018-12-14[PR88370] acc_get_cuda_stream/acc_set_cuda_stream: acc_async_sync, ↵Thomas Schwinge4-10/+187
acc_async_noval Per my reading of the OpenACC specification (and as supported by secondary documentation, such as code examples, or presentations), it's valid to call "acc_get_cuda_stream"/"acc_set_cuda_stream" also with "acc_async_sync", "acc_async_noval" arguments, not just with the nonnegative values as currently implemented. libgomp/ PR libgomp/88370 * libgomp.texi (acc_get_current_cuda_context, acc_get_cuda_stream) (acc_set_cuda_stream): Clarify. * oacc-cuda.c (acc_get_cuda_stream, acc_set_cuda_stream): Use "async_valid_p". * plugin/plugin-nvptx.c (nvptx_set_cuda_stream): Refuse "async == acc_async_sync". * testsuite/libgomp.oacc-c-c++-common/acc_set_cuda_stream-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-84.c: Update. * testsuite/libgomp.oacc-c-c++-common/lib-85.c: Likewise. From-SVN: r267147
2018-12-14[offloading] Error on missing symbolsTom de Vries5-0/+82
When compiling an OpenMP or OpenACC program containing a reference in the offloaded code to a symbol that has not been included in the offloaded code, the offloading compiler may ICE in lto1. Fix this by erroring out instead, mentioning the problematic symbol: ... error: variable 'var' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code lto1: fatal error: errors during merging of translation units compilation terminated. ... Build x86_64 with nvptx accelerator and reg-tested libgomp. Build x86_64 and reg-tested libgomp. 2018-12-14 Tom de Vries <tdevries@suse.de> * lto-cgraph.c (verify_node_partition): New function. (input_overwrite_node, input_varpool_node): Use verify_node_partition. * testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c: New test. * testsuite/libgomp.c-c++-common/function-not-offloaded.c: New test. * testsuite/libgomp.c-c++-common/variable-not-offloaded.c: New test. * testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c: New test. * testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c: New test. From-SVN: r267134
2018-12-13re PR libgomp/88460 ([nvptx] FAIL: libgomp.c++/for-24.C (internal compiler ↵Jakub Jelinek1-20/+18
error)) PR libgomp/88460 * testsuite/libgomp.c++/for-24.C (results): Include it in omp declare target region. (main): Use map (always, tofrom: results) instead of map (tofrom: results). From-SVN: r267093
2018-12-12re PR fortran/88463 (Rejects conforming source, OpenMP Parallel region ↵Jakub Jelinek2-0/+28
Default(None) reference to module parameter array, separate source) PR fortran/88463 * trans-openmp.c (gfc_omp_predetermined_sharing): Handle TREE_READONLY VAR_DECLs with DECL_EXTERNAL like those with TREE_STATIC. * testsuite/libgomp.fortran/pr88463-1.f90: New test. * testsuite/libgomp.fortran/pr88463-2.f90: New test. From-SVN: r267069
2018-12-12omp-builtins.def (BUILT_IN_GOMP_LOOP_NONMONOTONIC_RUNTIME_START, [...]): Fix ↵Jakub Jelinek1-0/+114
up function types - remove one argument. * omp-builtins.def (BUILT_IN_GOMP_LOOP_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_LOOP_MAYBE_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_LOOP_ULL_MAYBE_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_PARALLEL_LOOP_NONMONOTONIC_RUNTIME, BUILT_IN_GOMP_PARALLEL_LOOP_MAYBE_NONMONOTONIC_RUNTIME): Fix up function types - remove one argument. * testsuite/libgomp.c-c++-common/for-16.c: New test. From-SVN: r267067
2018-12-09re PR libfortran/88411 (Random crashes for ASYNCHRONOUS writes (bad locking?))Thomas Koenig1-0/+30
2018-12-09 Thomas Koenig <tkoenig@gcc.gnu.org> PR fortran/88411 * io/transfer.c (dta_transfer_init): Do not treat as an asynchronous statement unless the statement has ASYNCHRONOUS="YES". (st_write_done): Likewise. (st_read_done): Do not perform async_wait for synchronous I/O on an async unit. (st_read_done): Likewise. 2018-12-09 Thomas Koenig <tkoenig@gcc.gnu.org> PR fortran/88411 * testsuite/libgomp.fortran/async_io_8.f90: New test. From-SVN: r266929
2018-12-08re PR libgomp/87995 (libgomp.c/../libgomp.c-c++-common/cancel-taskgroup-3.c ↵Jakub Jelinek1-4/+9
fails consistently after r265930) PR libgomp/87995 * testsuite/libgomp.c-c++-common/cancel-taskgroup-3.c: Require tls_runtime effective target. (t): New threadprivate variable. (main): Set t in threads which execute iterations of the worksharing loop. Propagate that to the task after the loop and don't abort if the current taskgroup hasn't been cancelled. From-SVN: r266904
2018-12-02tree-nested.c (convert_nonlocal_omp_clauses, [...]): Handle ↵Jakub Jelinek1-0/+60
OMP_CLAUSE_IN_REDUCTION... * tree-nested.c (convert_nonlocal_omp_clauses, convert_local_omp_clauses): Handle OMP_CLAUSE_IN_REDUCTION, OMP_CLAUSE_TASK_REDUCTION and OMP_CLAUSE__SIMT_ clauses. (convert_nonlocal_reference_stmt, convert_local_reference_stmt): Convert clauses for GIMPLE_OMP_TASKGROUP. * testsuite/libgomp.c/task-reduction-3.c: New test. From-SVN: r266723
2018-12-02omp-low.c (check_omp_nesting_restrictions): Allow cancel or cancellation ↵Jakub Jelinek1-0/+32
point with taskgroup clause inside of taskloop. * omp-low.c (check_omp_nesting_restrictions): Allow cancel or cancellation point with taskgroup clause inside of taskloop. Consider a taskloop construct without nogroup clause as implicit taskgroup for diagnostics if cancel/cancellation point with taskgroup clause is closely nested inside of taskgroup region. * c-c++-common/gomp/cancel-1.c (f2): Add various taskloop related tests. * testsuite/libgomp.c-c++-common/cancel-taskgroup-4.c: New test. From-SVN: r266722
2018-11-30[PR88288, OpenACC, libgomp] Adjust offsets for present data clausesCesar Philippidis1-0/+41
Make libgomp respect the on device offset of subarrays which may arise in present data clauses. libgomp/ PR libgomp/88288 * oacc-parallel.c (GOACC_parallel_keyed): Add offset to devaddrs. * testsuite/libgomp.oacc-c-c++-common/pr88288.c: New test. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com> From-SVN: r266688
2018-11-30Add libgomp.oacc-fortran/lib-16-2.f90Thomas Schwinge2-0/+59
This is a copy of libgomp.oacc-fortran/lib-16.f90, but does 'include "openacc_lib.h"' instead of 'use openacc'. libgomp/ * testsuite/libgomp.oacc-fortran/lib-16-2.f90: New file. From-SVN: r266683
generated by cgit v1.1 at 2025-09-27 06:03:35 +0000