aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite
AgeCommit message (Collapse)AuthorFilesLines
2019-07-09tree-vect-stmts.c (scan_operand_equal_p): Look through MEM_REF with SSA_NAME ↵Jakub Jelinek2-2/+2
address of POINTER_PLUS_EXPR. * tree-vect-stmts.c (scan_operand_equal_p): Look through MEM_REF with SSA_NAME address of POINTER_PLUS_EXPR. Handle MULT_EXPR and casts in offset when different, both through gimple stmts and through trees. Rewritten using loops to minimize code duplication for each operand. * g++.dg/vect/simd-6.cc: Replace xfail with target x86. * g++.dg/vect/simd-9.cc: Likewise. * testsuite/libgomp.c++/scan-13.C: Replace xfail with target x86. * testsuite/libgomp.c++/scan-16.C: Likewise. From-SVN: r273249
2019-07-06omp-low.c (lower_rec_input_clauses): For lastprivate clauses in ↵Jakub Jelinek2-0/+238
ctx->for_simd_scan_phase simd copy the outer var to... * omp-low.c (lower_rec_input_clauses): For lastprivate clauses in ctx->for_simd_scan_phase simd copy the outer var to the privatized variable(s). For conditional lastprivate look through outer GIMPLE_OMP_SCAN context. (lower_omp_1): For conditional lastprivate look through outer GIMPLE_OMP_SCAN context. * testsuite/libgomp.c/scan-19.c: New test. * testsuite/libgomp.c/scan-20.c: New test. From-SVN: r273169
2019-07-06omp-low.c (struct omp_context): Add for_simd_scan_phase member.Jakub Jelinek16-0/+2126
* omp-low.c (struct omp_context): Add for_simd_scan_phase member. (maybe_lookup_ctx): Add forward declaration. (omp_find_scan): Likewise. Walk into body of simd if composited with worksharing loop. (scan_omp_simd_scan): New function. (scan_omp_1_stmt): Call it. (lower_rec_simd_input_clauses): Don't create rvar nor rvar2 if ctx->for_simd_scan_phase. (lower_rec_input_clauses): Do much less work for inscan reductions in ctx->for_simd_scan_phase is_simd regions. (lower_omp_scan): Set is_simd also on simd constructs composited with worksharing loop, unless ctx->for_simd_scan_phase. Never emit a sorry message. Don't change GIMPLE_OMP_SCAN stmts into nops and emit their body after in simd constructs composited with worksharing loop. (lower_omp_for_scan): Handle worksharing loop composited with simd. * c-c++-common/gomp/scan-4.c: Don't expect sorry message. * testsuite/libgomp.c/scan-11.c: New test. * testsuite/libgomp.c/scan-12.c: New test. * testsuite/libgomp.c/scan-13.c: New test. * testsuite/libgomp.c/scan-14.c: New test. * testsuite/libgomp.c/scan-15.c: New test. * testsuite/libgomp.c/scan-16.c: New test. * testsuite/libgomp.c/scan-17.c: New test. * testsuite/libgomp.c/scan-18.c: New test. * testsuite/libgomp.c++/scan-9.C: New test. * testsuite/libgomp.c++/scan-10.C: New test. * testsuite/libgomp.c++/scan-11.C: New test. * testsuite/libgomp.c++/scan-12.C: New test. * testsuite/libgomp.c++/scan-13.C: New test. * testsuite/libgomp.c++/scan-14.C: New test. * testsuite/libgomp.c++/scan-15.C: New test. * testsuite/libgomp.c++/scan-16.C: New test. From-SVN: r273157
2019-07-04omp-expand.c (expand_omp_for_static_nochunk): Don't emit GOMP_loop_start at ↵Jakub Jelinek2-0/+232
the start of second worksharing loop in a scan. * omp-expand.c (expand_omp_for_static_nochunk): Don't emit GOMP_loop_start at the start of second worksharing loop in a scan. For nowait, don't emit GOMP_loop_end_nowait at the end of first worksharing loop in a scan even if there are conditional lastprivates, and do emit GOMP_loop_end_nowait at the end of second worksharing loop. * testsuite/libgomp.c/scan-9.c: New test. * testsuite/libgomp.c/scan-10.c: New test. From-SVN: r273095
2019-07-03tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__SCANTEMP_ clause.Jakub Jelinek16-0/+2078
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE__SCANTEMP_ clause. * tree.h (OMP_CLAUSE_DECL): Use OMP_CLAUSE__SCANTEMP_ instead of OMP_CLAUSE__CONDTEMP_ as range's upper bound. (OMP_CLAUSE__SCANTEMP__ALLOC, OMP_CLAUSE__SCANTEMP__CONTROL): Define. * tree.c (omp_clause_num_ops, omp_clause_code_name): Add OMP_CLAUSE__SCANTEMP_ entry. (walk_tree_1): Handle OMP_CLAUSE__SCANTEMP_. * tree-pretty-print.c (dump_omp_clause): Likewise. * tree-nested.c (convert_nonlocal_omp_clauses, convert_local_omp_clauses): Likewise. * omp-general.h (struct omp_for_data): Add have_scantemp and have_nonctrl_scantemp members. * omp-general.c (omp_extract_for_data): Initialize them. * omp-low.c (struct omp_context): Add scan_exclusive member. (scan_omp_1_stmt): Don't unnecessarily mask gimple_omp_for_kind result again with GF_OMP_FOR_KIND_MASK. Initialize also ctx->scan_exclusive. (lower_rec_simd_input_clauses): Use ctx->scan_exclusive instead of !ctx->scan_inclusive. (lower_rec_input_clauses): Simplify gimplification of dtors using gimplify_and_add. For non-is_simd test OMP_CLAUSE_REDUCTION_INSCAN rather than rvarp. Handle OMP_CLAUSE_REDUCTION_INSCAN in worksharing loops. Don't add barrier for reduction_omp_orig_ref if ctx->scan_??xclusive. (lower_reduction_clauses): Don't do anything for ctx->scan_??xclusive. (lower_omp_scan): Use ctx->scan_exclusive instead of !ctx->scan_inclusive. Handle worksharing loops with inscan reductions. Use new_vard != new_var instead of repeated omp_is_reference calls. (omp_find_scan, lower_omp_for_scan): New functions. (lower_omp_for): Call lower_omp_for_scan for worksharing loops with inscan reductions. * omp-expand.c (expand_omp_scantemp_alloc): New function. (expand_omp_for_static_nochunk): Handle fd->have_nonctrl_scantemp and fd->have_scantemp. * c-c++-common/gomp/scan-3.c (f1): Don't expect a sorry message. * c-c++-common/gomp/scan-5.c (foo): Likewise. * testsuite/libgomp.c++/scan-1.C: New test. * testsuite/libgomp.c++/scan-2.C: New test. * testsuite/libgomp.c++/scan-3.C: New test. * testsuite/libgomp.c++/scan-4.C: New test. * testsuite/libgomp.c++/scan-5.C: New test. * testsuite/libgomp.c++/scan-6.C: New test. * testsuite/libgomp.c++/scan-7.C: New test. * testsuite/libgomp.c++/scan-8.C: New test. * testsuite/libgomp.c/scan-1.c: New test. * testsuite/libgomp.c/scan-2.c: New test. * testsuite/libgomp.c/scan-3.c: New test. * testsuite/libgomp.c/scan-4.c: New test. * testsuite/libgomp.c/scan-5.c: New test. * testsuite/libgomp.c/scan-6.c: New test. * testsuite/libgomp.c/scan-7.c: New test. * testsuite/libgomp.c/scan-8.c: New test. From-SVN: r272958
2019-06-19Test cases to verify OpenACC 'firstprivate' mappingsThomas Schwinge2-0/+9
gcc/testsuite/ * c-c++-common/goacc/firstprivate-mappings-1.c: New file. * g++.dg/goacc/firstprivate-mappings-1.C: Likewise. libgomp/ * testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C: New file. * testsuite/libgomp.oacc-c-c++-common/firstprivate-mappings-1.c: Likewise. From-SVN: r272451
2019-06-19Add missing results check in 'libgomp.fortran/allocatable3.f90'Thomas Schwinge1-0/+1
libgomp/ * testsuite/libgomp.fortran/allocatable3.f90: Add missing results check. From-SVN: r272449
2019-06-19Add 'libgomp.oacc-fortran/allocatable-array-1.f90'Cesar Philippidis1-0/+27
libgomp/ * testsuite/libgomp.oacc-fortran/allocatable-array-1.f90: New file. From-SVN: r272448
2019-06-19[PR90743] Fortran 'allocatable' with OpenACC data/OpenMP 'target' 'map' clausesThomas Schwinge4-0/+300
Test what OpenMP 5.0 has to say on this topic. And, do the same for OpenACC. libgomp/ PR fortran/90743 * oacc-parallel.c (GOACC_parallel_keyed): Handle NULL mapping case. * testsuite/libgomp.fortran/target-allocatable-1-1.f90: New file. * testsuite/libgomp.fortran/target-allocatable-1-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/allocatable-1-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/allocatable-1-2.f90: Likewise. From-SVN: r272447
2019-06-19[PR90861] Document status quo for OpenACC 'declare' not cleaning up for VLAsThomas Schwinge1-3/+44
gcc/testsuite/ PR testsuite/90861 * c-c++-common/goacc/declare-pr90861.c: New file. libgomp/ PR testsuite/90861 * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Update. From-SVN: r272446
2019-06-19[PR90862] OpenACC 'declare' ICE when nested inside another constructThomas Schwinge1-6/+92
gcc/ PR middle-end/90862 * omp-low.c (check_omp_nesting_restrictions): Handle GF_OMP_TARGET_KIND_OACC_DECLARE. gcc/testsuite/ PR middle-end/90862 * c-c++-common/goacc/declare-1.c: Update. * c-c++-common/goacc/declare-2.c: Likewise. libgomp/ PR middle-end/90862 * testsuite/libgomp.oacc-c-c++-common/declare-1.c: Update. From-SVN: r272444
2019-06-16[openacc, parloops] Fix SIGSEGV in oacc_entry_exit_ok_1Tom de Vries1-0/+15
When compiling the test-case with r268755, we run into a SIGSEGV in oacc_entry_exit_ok_1 when trying to dereference a NULL red: ... struct reduction_info *red; red = reduction_phi (reduction_list, use_stmt); tree val = PHI_RESULT (red->keep_res); ... Fix this by handling ref == NULL. Bootstrapped and reg-tested on x86_64. Build and reg-tested on x86_64 with nvptx accelerator. 2019-06-16 Tom de Vries <tdevries@suse.de> PR tree-optimization/89376 * tree-parloops.c (oacc_entry_exit_ok_1): Handle red == NULL. * testsuite/libgomp.oacc-c-c++-common/pr89376.c: New test. From-SVN: r272338
2019-06-15[nvptx, libgomp] Update pr85381-{2,4}.c test-casesTom de Vries2-23/+2
After the fix for "PR tree-optimization/89713 - Assume loop with an exit is finite" ( r272234 ) empty oacc loops are removed before expand. Update pr85381-{2,4}.c accordingly. 2019-06-15 Tom de Vries <tdevries@suse.de> PR tree-optimization/89713 * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Expect no bar.sync. * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Same. From-SVN: r272324
2019-06-15re PR middle-end/90779 (Fortran array initialization in offload regions)Jakub Jelinek2-0/+30
PR middle-end/90779 * gimplify.c: Include omp-offload.h and context.h. (gimplify_bind_expr): Add "omp declare target" attributes to static block scope variables inside of target region or target functions. * c-c++-common/goacc/routine-5.c (func2): Don't expect error for static block scope variable in #pragma acc routine. * testsuite/libgomp.c/pr90779.c: New test. * testsuite/libgomp.fortran/pr90779.f90: New test. From-SVN: r272322
2019-06-15[openacc] Disable pass_thread_jumps for IFN_UNIQUETom de Vries1-0/+34
If we compile the openacc testcase with -fopenacc -O2, we run into a SIGSEGV or assert. The root cause for this is that pass_thread_jumps breaks the invariant that OACC_FORK and OACC_JOIN mark the start and end of a single-entry-single-exit region. Fix this by bailing out when encountering an IFN_UNIQUE in thread_jumps::profitable_jump_thread_path. Bootstrapped and reg-tested on x86_64. Build and reg-tested libgomp on x86_64 with nvptx accelerator. 2019-06-15 Tom de Vries <tdevries@suse.de> PR tree-optimization/90009 * tree-ssa-threadbackward.c (thread_jumps::profitable_jump_thread_path): Return NULL if bb contains IFN_UNIQUE. * testsuite/libgomp.oacc-c-c++-common/pr90009.c: New test. From-SVN: r272321
2019-06-13PR tree-optimization/89713 - Assume loop with an exit is finiteFeng Xue1-0/+31
gcc/ChangeLog: * doc/invoke.texi (-ffinite-loops): Document new option. * common.opt (-ffinite-loops): New option. * tree-ssa-dce.c (mark_stmt_if_obviously_necessary): Mark IFN_GOACC_LOOP calls as necessary. * tree-ssa-loop-niter.c (finite_loop_p): Assume loop with an exit is finite. * omp-offload.c (oacc_xform_loop): Skip lowering if return value of IFN_GOACC_LOOP call is not used. * opts.c (default_options_table): Enable -ffinite-loops at -O2+. gcc/testsuite/ChangeLog: * g++.dg/tree-ssa/empty-loop.C: New test. * gcc.dg/tree-ssa/dce-2.c: New test. * gcc.dg/const-1.c: Add -fno-finite-loops option. * gcc.dg/graphite/graphite.exp: Likewise. * gcc.dg/loop-unswitch-1.c: Likewise. * gcc.dg/predict-9.c: Likewise. * gcc.dg/pure-2.c: Likewise. * gcc.dg/tree-ssa/20040211-1.c: Likewise. * gcc.dg/tree-ssa/loop-10.c: Likewise. * gcc.dg/tree-ssa/split-path-6.c: Likewise. * gcc.dg/tree-ssa/ssa-thread-12.c: Likewise. libgomp/ChangeLog: * testsuite/libgomp.oacc-c-c++-common/pr84955-1.c: New test. From-SVN: r272234
2019-06-11re PR target/90811 ([nvptx] ptxas error on OpenMP offloaded code)Jakub Jelinek1-0/+29
PR target/90811 * config/nvptx/nvptx.c (nvptx_output_softstack_switch): Use and.b%d instead of and.u%d. * testsuite/libgomp.c/pr90811.c: New test. From-SVN: r272161
2019-06-05omp-low.c (lower_rec_input_clauses): For lastprivate conditional references...Jakub Jelinek2-0/+166
* omp-low.c (lower_rec_input_clauses): For lastprivate conditional references, lookup in in hash map MEM_REF operand instead of the MEM_REF itself. (lower_omp_1): When looking for lastprivate conditional assignments, handle MEM_REFs with REFERENCE_TYPE operands. * testsuite/libgomp.c++/lastprivate-conditional-1.C: New test. * testsuite/libgomp.c++/lastprivate-conditional-2.C: New test. From-SVN: r271948
2019-06-04gimplify.c (gimplify_scan_omp_clauses): Don't sorry_at on lastprivate ↵Jakub Jelinek4-0/+240
conditional on combined for simd. * gimplify.c (gimplify_scan_omp_clauses): Don't sorry_at on lastprivate conditional on combined for simd. * omp-low.c (struct omp_context): Add combined_into_simd_safelen0 member. (lower_rec_input_clauses): For gimple_omp_for_combined_into_p max_vf 1 constructs, don't remove lastprivate_conditional_map, but instead set ctx->combined_into_simd_safelen0 and adjust hash_map, so that it points to parent construct temporaries. (lower_lastprivate_clauses): Handle ctx->combined_into_simd_safelen0 like !ctx->lastprivate_conditional_map. (lower_omp_1) <case GIMPLE_ASSIGN>: If up->combined_into_simd_safelen0, use up->outer context instead of up. * omp-expand.c (expand_omp_for_generic): Perform cond_var bump even if gimple_omp_for_combined_p. (expand_omp_for_static_nochunk): Likewise. (expand_omp_for_static_chunk): Add forgotten cond_var bump that was probably moved over into expand_omp_for_generic rather than being copied there. gcc/cp/ * cp-tree.h (CP_OMP_CLAUSE_INFO): Allow for any clauses up to _condvar_ instead of only up to linear. gcc/testsuite/ * c-c++-common/gomp/lastprivate-conditional-2.c (foo): Don't expect a sorry_at on any of the clauses. libgomp/ * testsuite/libgomp.c-c++-common/lastprivate-conditional-7.c: New test. * testsuite/libgomp.c-c++-common/lastprivate-conditional-8.c: New test. * testsuite/libgomp.c-c++-common/lastprivate-conditional-9.c: New test. * testsuite/libgomp.c-c++-common/lastprivate-conditional-10.c: New test. From-SVN: r271907
2019-05-30Generalize getconf _NPROCESSORS_ONLNRainer Orth2-4/+6
libgomp: * configure.ac: Call AX_COUNT_CPUS. Substitute CPU_COUNT. * testsuite/Makefile.am (check-am): Use CPU_COUNT as processor count fallback. * aclocal.m4: Regenerate. * configure: Regenerate. * Makefile.in, testsuite/Makefile.in: Regenerate. config: * ax_count_cpus.m4: New file. From-SVN: r271769
2019-05-29gimplify.c (struct gimplify_omp_ctx): Add clauses member.Jakub Jelinek3-0/+302
* gimplify.c (struct gimplify_omp_ctx): Add clauses member. (gimplify_scan_omp_clauses): Initialize ctx->clauses. (gimplify_adjust_omp_clauses_1): Transform lastprivate conditional explicit clause on combined parallel into implicit shared clause. (gimplify_adjust_omp_clauses): Move lastprivate conditional clause and firstprivate if the decl has one too from combined parallel to the worksharing construct. gcc/testsuite/ * c-c++-common/gomp/lastprivate-conditional-2.c (foo): Don't expect sorry on lastprivate conditional on parallel for. * c-c++-common/gomp/lastprivate-conditional-3.c (foo): Add tests for lastprivate conditional warnings on parallel for constructs. * c-c++-common/gomp/lastprivate-conditional-4.c: New test. libgomp/ * testsuite/libgomp.c-c++-common/lastprivate_conditional_4.c: Rename to ... * testsuite/libgomp.c-c++-common/lastprivate-conditional-4.c: ... this. * testsuite/libgomp.c-c++-common/lastprivate-conditional-5.c: New test. * testsuite/libgomp.c-c++-common/lastprivate-conditional-6.c: New test. From-SVN: r271733
2019-05-27gimplify.c (gimplify_scan_omp_clauses): Allow lastprivate conditional on ↵Jakub Jelinek1-0/+161
sections construct. * gimplify.c (gimplify_scan_omp_clauses): Allow lastprivate conditional on sections construct. * omp-low.c (lower_lastprivate_conditional_clauses): Handle sections construct. (lower_omp_sections): Handle lastprivate conditional. (lower_omp_1) <case GIMPLE_ASSIGN>: Handle sections construct with lastprivate_conditional_map. * omp-expand.c (expand_omp_sections): Handle lastprivate conditional. libgomp/ * testsuite/libgomp.c-c++-common/lastprivate_conditional_4.c: New test. From-SVN: r271673
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
generated by cgit v1.1 at 2025-09-26 21:50:40 +0000