aboutsummaryrefslogtreecommitdiff
path: root/gcc/omp-low.c
AgeCommit message (Collapse)AuthorFilesLines
2016-10-13Move MEMMODEL_* from coretypes.h to memmodel.hThomas Preud'homme1-0/+1
2016-10-13 Thomas Preud'homme <thomas.preudhomme@arm.com> gcc/ * coretypes.h: Move MEMMODEL_* macros and enum memmodel definition into ... * memmodel.h: This file. * alias.c, asan.c, auto-inc-dec.c, bb-reorder.c, bt-load.c, caller-save.c, calls.c, ccmp.c, cfgbuild.c, cfgcleanup.c, cfgexpand.c, cfgloopanal.c, cfgrtl.c, cilk-common.c, combine.c, combine-stack-adj.c, common/config/aarch64/aarch64-common.c, common/config/arm/arm-common.c, common/config/bfin/bfin-common.c, common/config/c6x/c6x-common.c, common/config/i386/i386-common.c, common/config/ia64/ia64-common.c, common/config/nvptx/nvptx-common.c, compare-elim.c, config/aarch64/aarch64-builtins.c, config/aarch64/aarch64-c.c, config/aarch64/cortex-a57-fma-steering.c, config/arc/arc.c, config/arc/arc-c.c, config/arm/arm-builtins.c, config/arm/arm-c.c, config/avr/avr.c, config/avr/avr-c.c, config/avr/avr-log.c, config/bfin/bfin.c, config/c6x/c6x.c, config/cr16/cr16.c, config/cris/cris.c, config/darwin-c.c, config/darwin.c, config/epiphany/epiphany.c, config/epiphany/mode-switch-use.c, config/epiphany/resolve-sw-modes.c, config/fr30/fr30.c, config/frv/frv.c, config/ft32/ft32.c, config/h8300/h8300.c, config/i386/i386-c.c, config/i386/winnt.c, config/iq2000/iq2000.c, config/lm32/lm32.c, config/m32c/m32c.c, config/m32r/m32r.c, config/m68k/m68k.c, config/mcore/mcore.c, config/microblaze/microblaze.c, config/mmix/mmix.c, config/mn10300/mn10300.c, config/moxie/moxie.c, config/msp430/msp430.c, config/nds32/nds32-cost.c, config/nds32/nds32-intrinsic.c, config/nds32/nds32-md-auxiliary.c, config/nds32/nds32-memory-manipulation.c, config/nds32/nds32-predicates.c, config/nds32/nds32.c, config/nios2/nios2.c, config/nvptx/nvptx.c, config/pa/pa.c, config/pdp11/pdp11.c, config/rl78/rl78.c, config/rs6000/rs6000-c.c, config/rx/rx.c, config/s390/s390-c.c, config/s390/s390.c, config/sh/sh.c, config/sh/sh-c.c, config/sh/sh-mem.cc, config/sh/sh_treg_combine.cc, config/sol2.c, config/spu/spu.c, config/stormy16/stormy16.c, config/tilegx/tilegx.c, config/tilepro/tilepro.c, config/v850/v850.c, config/vax/vax.c, config/visium/visium.c, config/vms/vms-c.c, config/xtensa/xtensa.c, coverage.c, cppbuiltin.c, cprop.c, cse.c, cselib.c, dbxout.c, dce.c, df-core.c, df-problems.c, df-scan.c, dojump.c, dse.c, dwarf2asm.c, dwarf2cfi.c, dwarf2out.c, emit-rtl.c, except.c, explow.c, expmed.c, expr.c, final.c, fold-const.c, function.c, fwprop.c, gcse.c, ggc-page.c, haifa-sched.c, hsa-brig.c, hsa-gen.c, hw-doloop.c, ifcvt.c, init-regs.c, internal-fn.c, ira-build.c, ira-color.c, ira-conflicts.c, ira-costs.c, ira-emit.c, ira-lives.c, ira.c, jump.c, loop-doloop.c, loop-invariant.c, loop-iv.c, loop-unroll.c, lower-subreg.c, lra.c, lra-assigns.c, lra-coalesce.c, lra-constraints.c, lra-eliminations.c, lra-lives.c, lra-remat.c, lra-spills.c, mode-switching.c, modulo-sched.c, omp-low.c, passes.c, postreload-gcse.c, postreload.c, predict.c, print-rtl-function.c, recog.c, ree.c, reg-stack.c, regcprop.c, reginfo.c, regrename.c, reload.c, reload1.c, reorg.c, resource.c, rtl-chkp.c, rtl-tests.c, rtlanal.c, rtlhooks.c, sched-deps.c, sched-rgn.c, sdbout.c, sel-sched-ir.c, sel-sched.c, shrink-wrap.c, simplify-rtx.c, stack-ptr-mod.c, stmt.c, stor-layout.c, target-globals.c, targhooks.c, toplev.c, tree-nested.c, tree-outof-ssa.c, tree-profile.c, tree-ssa-coalesce.c, tree-ssa-ifcombine.c, tree-ssa-loop-ivopts.c, tree-ssa-loop.c, tree-ssa-reassoc.c, tree-ssa-sccvn.c, tree-vect-data-refs.c, ubsan.c, valtrack.c, var-tracking.c, varasm.c: Include memmodel.h. * genattrtab.c (write_header): Include memmodel.h in generated file. * genautomata.c (main): Likewise. * gengtype.c (open_base_files): Likewise. * genopinit.c (main): Likewise. * genconditions.c (write_header): Include memmodel.h earlier in generated file. * genemit.c (main): Likewise. * genoutput.c (output_prologue): Likewise. * genpeep.c (main): Likewise. * genpreds.c (write_insn_preds_c): Likewise. * genrecog.c (write_header): Likewise. * Makefile.in (PLUGIN_HEADERS): Include memmodel.h gcc/ada/ * gcc-interface/utils2.c: Include memmodel.h. gcc/c-family/ * c-cppbuiltin.c: Include memmodel.h. * c-opts.c: Likewise. * c-pragma.c: Likewise. * c-warn.c: Likewise. gcc/c/ * c-typeck.c: Include memmodel.h. gcc/cp/ * decl2.c: Include memmodel.h. * rtti.c: Likewise. gcc/fortran/ * trans-intrinsic.c: Include memmodel.h. gcc/go/ * go-backend.c: Include memmodel.h. libgcc/ * libgcov-profiler.c: Replace MEMMODEL_* macros by their __ATOMIC_* equivalent. * config/tilepro/atomic.c: Likewise and stop casting model to enum memmodel. From-SVN: r241121
2016-10-09tree-ssa.c (target_for_debug_bind, [...]): Use VAR_P and/or ↵Jakub Jelinek1-12/+8
VAR_OR_FUNCTION_DECL_P macros. * tree-ssa.c (target_for_debug_bind, verify_phi_args, ssa_undefined_value_p, maybe_optimize_var): Use VAR_P and/or VAR_OR_FUNCTION_DECL_P macros. * tree-chkp.c (chkp_register_var_initializer, chkp_make_static_bounds, chkp_get_bounds_for_decl_addr, chkp_parse_array_and_component_ref, chkp_find_bounds_1): Likewise. * ipa-polymorphic-call.c (decl_maybe_in_construction_p): Likewise. * hsa-gen.c (get_symbol_for_decl): Likewise. * cgraphunit.c (check_global_declaration, analyze_functions, handle_alias_pairs, thunk_adjust, cgraph_node::expand_thunk): Likewise. * gimple-fold.c (can_refer_decl_in_current_unit_p, canonicalize_constructor_val, gimple_get_virt_method_for_vtable): Likewise. * tree.c (set_decl_section_name, copy_node_stat, need_assembler_name_p, free_lang_data_in_decl, find_decls_types_r, merge_dllimport_decl_attributes, handle_dll_attribute, decl_init_priority_insert, auto_var_in_fn_p, array_at_struct_end_p, verify_type): Likewise. * gimple-ssa-isolate-paths.c (find_implicit_erroneous_behavior, find_explicit_erroneous_behavior): Likewise. * sdbout.c (sdbout_toplevel_data, sdbout_late_global_decl): Likewise. * ipa.c (process_references): Likewise. * tree-chkp-opt.c (chkp_get_check_result): Likewise. * varasm.c (get_block_for_decl, use_blocks_for_decl_p, make_decl_rtl, notice_global_symbol, assemble_variable, mark_decl_referenced, build_constant_desc, output_constant_def_contents, do_assemble_alias, make_decl_one_only, default_section_type_flags, categorize_decl_for_section, default_encode_section_info): Likewise. * trans-mem.c (requires_barrier): Likewise. * gimple-expr.c (mark_addressable): Likewise. * cfgexpand.c (add_scope_conflicts_1, expand_one_var, expand_used_vars_for_block, clear_tree_used, stack_protect_decl_p, expand_debug_expr): Likewise. * tree-dump.c (dequeue_and_dump): Likewise. * ubsan.c (instrument_bool_enum_load): Likewise. * tree-pretty-print.c (print_declaration): Likewise. * simplify-rtx.c (delegitimize_mem_from_attrs): Likewise. * tree-ssa-uninit.c (warn_uninitialized_vars): Likewise. * asan.c (asan_protect_global, instrument_derefs): Likewise. * tree-into-ssa.c (rewrite_stmt, maybe_register_def, pass_build_ssa::execute): Likewise. * var-tracking.c (var_debug_decl, track_expr_p): Likewise. * tree-ssa-loop-ivopts.c (force_expr_to_var_cost, split_address_cost): Likewise. * ipa-split.c (test_nonssa_use, consider_split, mark_nonssa_use): Likewise. * tree-inline.c (insert_debug_decl_map, remap_ssa_name, can_be_nonlocal, remap_decls, copy_debug_stmt, initialize_inlined_parameters, add_local_variables, reset_debug_binding, replace_locals_op): Likewise. * dse.c (can_escape): Likewise. * ipa-devirt.c (compare_virtual_tables, referenced_from_vtable_p): Likewise. * tree-diagnostic.c (default_tree_printer): Likewise. * tree-streamer-in.c (unpack_ts_decl_common_value_fields, unpack_ts_decl_with_vis_value_fields, lto_input_ts_decl_common_tree_pointers): Likewise. * builtins.c (builtin_save_expr, fold_builtin_expect, readonly_data_expr): Likewise. * tree-ssa-structalias.c (new_var_info, get_constraint_for_ssa_var, create_variable_info_for, set_uids_in_ptset, visit_loadstore): Likewise. * gimple-streamer-out.c (output_gimple_stmt): Likewise. * gimplify.c (force_constant_size, gimplify_bind_expr, gimplify_decl_expr, gimplify_var_or_parm_decl, gimplify_compound_lval, gimplify_init_constructor, gimplify_modify_expr, gimplify_asm_expr, gimplify_oacc_declare, gimplify_type_sizes): Likewise. * cgraphbuild.c (record_reference, record_type_list, mark_address, mark_load, mark_store, pass_build_cgraph_edges::execute): Likewise. * tree-ssa-live.c (mark_all_vars_used_1, remove_unused_scope_block_p, remove_unused_locals): Likewise. * tree-ssa-alias.c (ptr_deref_may_alias_decl_p, ptrs_compare_unequal, ref_maybe_used_by_call_p_1, call_may_clobber_ref_p_1): Likewise. * function.c (instantiate_expr, instantiate_decls_1, setjmp_vars_warning, add_local_decl): Likewise. * alias.c (ao_ref_from_mem, get_alias_set, compare_base_symbol_refs): Likewise. * tree-stdarg.c (find_va_list_reference, va_list_counter_struct_op, va_list_ptr_read, va_list_ptr_write, check_all_va_list_escapes, optimize_va_list_gpr_fpr_size): Likewise. * tree-nrv.c (pass_nrv::execute): Likewise. * tsan.c (instrument_expr): Likewise. * tree-ssa-dce.c (remove_dead_stmt): Likewise. * vtable-verify.c (verify_bb_vtables): Likewise. * tree-dfa.c (ssa_default_def, set_ssa_default_def, get_ref_base_and_extent): Likewise. * toplev.c (wrapup_global_declaration_1, wrapup_global_declaration_2): Likewise. * tree-sra.c (static bool constant_decl_p, find_var_candidates, analyze_all_variable_accesses): Likewise. * tree-nested.c (get_nonlocal_debug_decl, convert_nonlocal_omp_clauses, note_nonlocal_vla_type, note_nonlocal_block_vlas, convert_nonlocal_reference_stmt, get_local_debug_decl, convert_local_omp_clauses, convert_local_reference_stmt, nesting_copy_decl, remap_vla_decls): Likewise. * tree-vect-data-refs.c (vect_can_force_dr_alignment_p): Likewise. * stmt.c (decl_overlaps_hard_reg_set_p): Likewise. * dbxout.c (dbxout_late_global_decl, dbxout_type_fields, dbxout_symbol, dbxout_common_check): Likewise. * expr.c (expand_assignment, expand_expr_real_2, expand_expr_real_1, string_constant): Likewise. * hsa.c (hsa_get_declaration_name): Likewise. * passes.c (rest_of_decl_compilation): Likewise. * tree-ssanames.c (make_ssa_name_fn): Likewise. * tree-streamer-out.c (pack_ts_decl_common_value_fields, pack_ts_decl_with_vis_value_fields, write_ts_decl_common_tree_pointers): Likewise. * stor-layout.c (place_field): Likewise. * symtab.c (symtab_node::maybe_create_reference, symtab_node::verify_base, symtab_node::make_decl_local, symtab_node::copy_visibility_from, symtab_node::can_increase_alignment_p): Likewise. * dwarf2out.c (add_var_loc_to_decl, tls_mem_loc_descriptor, decl_by_reference_p, reference_to_unused, rtl_for_decl_location, fortran_common, add_location_or_const_value_attribute, add_scalar_info, add_linkage_name, set_block_abstract_flags, local_function_static, gen_variable_die, dwarf2out_late_global_decl, optimize_one_addr_into_implicit_ptr, optimize_location_into_implicit_ptr): Likewise. * gimple-low.c (record_vars_into): Likewise. * ipa-visibility.c (update_vtable_references): Likewise. * tree-ssa-address.c (fixed_address_object_p, copy_ref_info): Likewise. * lto-streamer-out.c (tree_is_indexable, get_symbol_initial_value, DFS::DFS_write_tree_body, write_symbol): Likewise. * langhooks.c (lhd_warn_unused_global_decl, lhd_set_decl_assembler_name): Likewise. * attribs.c (decl_attributes): Likewise. * except.c (output_ttype): Likewise. * varpool.c (varpool_node::get_create, ctor_for_folding, varpool_node::assemble_decl, varpool_node::create_alias): Likewise. * fold-const.c (fold_unary_loc): Likewise. * ipa-prop.c (ipa_compute_jump_functions_for_edge, ipa_find_agg_cst_from_init): Likewise. * omp-low.c (expand_omp_regimplify_p, expand_omp_taskreg, expand_omp_target, lower_omp_regimplify_p, grid_reg_assignment_to_local_var_p, grid_remap_prebody_decls, find_link_var_op): Likewise. * tree-chrec.c (chrec_contains_symbols): Likewise. * tree-cfg.c (verify_address, verify_expr, verify_expr_location_1, gimple_duplicate_bb, move_stmt_op, replace_block_vars_by_duplicates, execute_fixup_cfg): Likewise. From-SVN: r240900
2016-09-26ipa-inline-analysis.c (find_foldable_builtin_expect): Use ↵Marek Polacek1-10/+4
gimple_call_internal_p. * ipa-inline-analysis.c (find_foldable_builtin_expect): Use gimple_call_internal_p. * ipa-split.c (find_return_bb): Likewise. (execute_split_functions): Likewise. * omp-low.c (dump_oacc_loop_part): Likewise. (oacc_loop_xform_head_tail): Likewise. * predict.c (predict_loops): Likewise. * sanopt.c (pass_sanopt::execute): Likewise. * tree-cfg.c (get_abnormal_succ_dispatcher): Likewise. * tree-parloops.c (oacc_entry_exit_ok_1): Likewise. * tree-stdarg.c (gimple_call_ifn_va_arg_p): Remove function. (expand_ifn_va_arg_1): Use gimple_call_internal_p. (expand_ifn_va_arg): Likewise. * tree-vect-loop.c (vect_determine_vectorization_factor): Likewise. (optimize_mask_stores): Likewise. * tree-vect-stmts.c (vect_simd_lane_linear): Likewise. (vect_transform_stmt): Likewise. * tree-vectorizer.c (vect_loop_vectorized_call): Likewise. * tsan.c (instrument_memory_accesses): Likewise. From-SVN: r240498
2016-09-16Add inline functions for various bitwise operations.Jason Merrill1-2/+2
* hwint.h (least_bit_hwi, pow2_or_zerop, pow2p_hwi, ctz_or_zero): New. * hwint.c (exact_log2): Use pow2p_hwi. (ctz_hwi, ffs_hwi): Use least_bit_hwi. * alias.c (memrefs_conflict_p): Use pow2_or_zerop. * builtins.c (get_object_alignment_2, get_object_alignment) (get_pointer_alignment, fold_builtin_atomic_always_lock_free): Use least_bit_hwi. * calls.c (compute_argument_addresses, store_one_arg): Use least_bit_hwi. * cfgexpand.c (expand_one_stack_var_at): Use least_bit_hwi. * combine.c (force_to_mode): Use least_bit_hwi. * emit-rtl.c (set_mem_attributes_minus_bitpos, adjust_address_1): Use least_bit_hwi. * expmed.c (synth_mult, expand_divmod): Use ctz_or_zero, ctz_hwi. (init_expmed_one_conv): Use pow2p_hwi. * fold-const.c (round_up_loc, round_down_loc): Use pow2_or_zerop. (fold_binary_loc): Use pow2p_hwi. * function.c (assign_parm_find_stack_rtl): Use least_bit_hwi. * gimple-fold.c (gimple_fold_builtin_memory_op): Use pow2p_hwi. * gimple-ssa-strength-reduction.c (replace_ref): Use least_bit_hwi. * hsa-gen.c (gen_hsa_addr_with_align, hsa_bitmemref_alignment): Use least_bit_hwi. * ipa-cp.c (ipcp_alignment_lattice::meet_with_1): Use least_bit_hwi. * ipa-prop.c (ipa_modify_call_arguments): Use least_bit_hwi. * omp-low.c (oacc_loop_fixed_partitions) (oacc_loop_auto_partitions): Use least_bit_hwi. * rtlanal.c (nonzero_bits1): Use ctz_or_zero. * stor-layout.c (place_field): Use least_bit_hwi. * tree-pretty-print.c (dump_generic_node): Use pow2p_hwi. * tree-sra.c (build_ref_for_offset): Use least_bit_hwi. * tree-ssa-ccp.c (ccp_finalize): Use least_bit_hwi. * tree-ssa-math-opts.c (bswap_replace): Use least_bit_hwi. * tree-ssa-strlen.c (handle_builtin_memcmp): Use pow2p_hwi. * tree-vect-data-refs.c (vect_analyze_group_access_1) (vect_grouped_store_supported, vect_grouped_load_supported) (vect_permute_load_chain, vect_shift_permute_load_chain) (vect_transform_grouped_load): Use pow2p_hwi. * tree-vect-generic.c (expand_vector_divmod): Use ctz_or_zero. * tree-vect-patterns.c (vect_recog_divmod_pattern): Use ctz_or_zero. * tree-vect-stmts.c (vectorizable_mask_load_store): Use least_bit_hwi. * tsan.c (instrument_expr): Use least_bit_hwi. * var-tracking.c (negative_power_of_two_p): Use pow2_or_zerop. From-SVN: r240194
2016-09-08re PR fortran/77516 (ICE in is_gimple_min_invariant, at gimple-expr.c:706)Jakub Jelinek1-1/+3
PR fortran/77516 * omp-low.c (lower_rec_simd_input_clauses): Use max_vf for non-positive OMP_CLAUSE_SAFELEN_EXPR. * gfortran.dg/gomp/pr77516.f90: New test. From-SVN: r240037
2016-09-02re PR c/65467 ([libgomp] sorry, unimplemented: '_Atomic' with OpenMP)Jakub Jelinek1-1/+2
PR c/65467 * gimplify.c (gimplify_adjust_omp_clauses_1): Diagnose implicit map and firstprivate clauses on target construct for _Atomic qualified decls. (gimplify_adjust_omp_clauses): Diagnose explicit firstprivate clauses on target construct for _Atomic qualified decls. * omp-low.c (use_pointer_for_field): Return true for _Atomic qualified decls. * omp-simd-clone.c (simd_clone_clauses_extract): Warn and give up for _Atomic qualified arguments not mentioned in uniform clause. c/ * c-parser.c (c_parser_declspecs): Don't sorry about _Atomic if flag_openmp. (c_parser_omp_variable_list): Use convert_lvalue_to_rvalue instead of mark_exp_read on low_bound/length expression. (c_parser_omp_clause_num_gangs, c_parser_omp_clause_num_threads, c_parser_omp_clause_num_tasks, c_parser_omp_clause_grainsize, c_parser_omp_clause_priority, c_parser_omp_clause_hint, c_parser_omp_clause_num_workers, c_parser_oacc_shape_clause, c_parser_oacc_clause_tile, c_parser_omp_clause_schedule, c_parser_omp_clause_vector_length, c_parser_omp_clause_num_teams, c_parser_omp_clause_thread_limit, c_parser_omp_clause_aligned, c_parser_omp_clause_linear, c_parser_omp_clause_safelen, c_parser_omp_clause_simdlen, c_parser_omp_clause_device, c_parser_omp_clause_dist_schedule): Use convert_lvalue_to_rvalue instead of mark_expr_read. (c_parser_omp_declare_reduction): Reject _Atomic qualified types. * c-objc-common.h (LANG_HOOKS_OMP_CLAUSE_COPY_CTOR, LANG_HOOKS_OMP_CLAUSE_ASSIGN_OP): Redefine. * c-tree.h (c_omp_clause_copy_ctor): New prototype. * c-typeck.c (handle_omp_array_sections_1): Diagnose _Atomic qualified array section bases outside of depend clause, for depend clause use convert_lvalue_to_rvalue on the base. (c_finish_omp_clauses): Reject _Atomic qualified vars in reduction, linear, aligned, map, to and from clauses. (c_omp_clause_copy_ctor): New function. c-family/ * c-omp.c (c_finish_omp_atomic): Reject _Atomic qualified expressions. (c_finish_omp_for): Reject _Atomic qualified iterators. testsuite/ * gcc.dg/gomp/_Atomic-1.c: New test. * gcc.dg/gomp/_Atomic-2.c: New test. * gcc.dg/gomp/_Atomic-3.c: New test. * gcc.dg/gomp/_Atomic-4.c: New test. * gcc.dg/gomp/_Atomic-5.c: New test. From-SVN: r239964
2016-08-17omp-low.c (lower_oacc_reductions): Adjust variable lookup to use ↵Chung-Lin Tang1-3/+12
maybe_lookup_decl... 2016-08-17 Chung-Lin Tang <cltang@codesourcery.com> * omp-low.c (lower_oacc_reductions): Adjust variable lookup to use maybe_lookup_decl, to handle nested acc loop directives. testsuite/ * c-c++-common/goacc/reduction-6.c: New testcase. From-SVN: r239530
2016-08-08re PR middle-end/72781 (-Wuninitialized false positives in OpenMP code)Jakub Jelinek1-1/+9
PR middle-end/72781 * omp-low.c (lower_lastprivate_clauses): Set TREE_NO_WARNING on the private vars for lastprivate and for linear iterator. * gcc.dg/gomp/pr72781.c: New test. From-SVN: r239249
2016-07-21function-tests.c (build_trivial_generic_function): Set BLOCK_SUPERCONTEXT of ↵Richard Biener1-0/+2
DECL_INITIAL. 2016-07-21 Richard Biener <rguenther@suse.de> * function-tests.c (build_trivial_generic_function): Set BLOCK_SUPERCONTEXT of DECL_INITIAL. * omp-low.c (create_omp_child_function): Likewise. (grid_expand_target_grid_body): Likewise. * cgraphunit.c (init_lowered_empty_function): Likewise. (cgraph_node::expand_thunk): Likewise. * tree-parloops.c (create_loop_fn): Likewise. * ipa.c (cgraph_build_static_cdtor_1): Likewise. cp/ * vtable-class-hierarchy.c (vtv_generate_init_routine): Set DECL_IGNORED_P. java/ * jcf-parse.c (java_emit_static_constructor): Set BLOCK_SUPERCONTEXT of DECL_INITIAL. From-SVN: r238589
2016-07-15c-parser.c (c_parser_oacc_declare): Don't scan for GOMP_MAP_POINTER.Cesar Philippidis1-0/+4
gcc/c/ * c-parser.c (c_parser_oacc_declare): Don't scan for GOMP_MAP_POINTER. * c-typeck.c (handle_omp_array_sections): Mark data clauses with GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} as potentially having zero-length subarrays. gcc/cp/ * parser.c (cp_parser_oacc_declare): Don't scan for GOMP_MAP_POINTER. * semantics.c (handle_omp_array_sections): Mark data clauses with GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} as potentially having zero-length subarrays. gcc/ * omp-low.c (lower_omp_target): Mark data clauses with GOMP_MAP_FORCE_{PRESENT,TO,FROM,TOFROM} as potentially having zero-length subarrays. libgomp/ * testsuite/libgomp.oacc-c-c++-common/zero_length_subarrays.c: New test. From-SVN: r238376
2016-07-11re PR middle-end/71758 (ICE in verify_gimple_in_cfg, at tree-cfg.c:5212)Jakub Jelinek1-1/+7
PR middle-end/71758 * omp-low.c (expand_omp_target): Gimplify device. * c-c++-common/gomp/pr71758.c: New test. * gfortran.dg/gomp/pr71758.f90: New test. From-SVN: r238231
2016-06-24re PR tree-optimization/71647 (aligned(x:32) in #pragma omp simd does not work)Jakub Jelinek1-4/+6
PR tree-optimization/71647 * omp-low.c (lower_rec_input_clauses): Convert omp_clause_aligned_alignment (c) to size_type_node for the last argument of __builtin_assume_aligned. * gcc.target/i386/pr71647.c: New test. From-SVN: r237769
2016-06-10[PR middle-end/71373] Handle more OMP_CLAUSE_* in nested function decompositionThomas Schwinge1-2/+2
gcc/ * gimplify.c (gimplify_adjust_omp_clauses): Discard OMP_CLAUSE_TILE. * omp-low.c (scan_sharing_clauses): Don't expect OMP_CLAUSE_TILE. gcc/testsuite/ * c-c++-common/goacc/combined-directives.c: XFAIL tree scanning for OpenACC tile clauses. * gfortran.dg/goacc/combined-directives.f90: Likewise. gcc/ PR middle-end/71373 * tree-nested.c (convert_nonlocal_omp_clauses) (convert_local_omp_clauses): Handle OMP_CLAUSE_ASYNC, OMP_CLAUSE_WAIT, OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_AUTO, OMP_CLAUSE__CACHE_, OMP_CLAUSE_TILE. gcc/testsuite/ PR middle-end/71373 * gcc.dg/goacc/nested-function-1.c: New file. * gcc.dg/goacc/nested-function-2.c: Likewise. * gcc.dg/goacc/pr71373.c: Likewise. * gfortran.dg/goacc/cray-2.f95: Likewise. * gfortran.dg/goacc/loop-1-2.f95: Likewise. * gfortran.dg/goacc/loop-3-2.f95: Likewise. * gfortran.dg/goacc/cray.f95: Update. * gfortran.dg/goacc/loop-1.f95: Likewise. * gfortran.dg/goacc/loop-3.f95: Likewise. * gfortran.dg/goacc/subroutines.f90: Update, and rename to... * gfortran.dg/goacc/nested-function-1.f90: ... this new file. libgomp/testsuite/ PR middle-end/71373 * libgomp.oacc-c/nested-function-1.c: New file. * libgomp.oacc-c/nested-function-2.c: Likewise. * libgomp.oacc-fortran/nested-function-1.f90: Likewise. * libgomp.oacc-fortran/nested-function-2.f90: Likewise. * libgomp.oacc-fortran/nested-function-3.f90: Likewise. Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com> From-SVN: r237291
2016-06-10[PR c/71381] C/C++ OpenACC cache directive rejects valid syntaxThomas Schwinge1-6/+0
gcc/c/ PR c/71381 * c-parser.c (c_parser_omp_variable_list) <OMP_CLAUSE__CACHE_>: Loosen checking. gcc/cp/ PR c/71381 * parser.c (cp_parser_omp_var_list_no_open) <OMP_CLAUSE__CACHE_>: Loosen checking. gcc/fortran/ PR c/71381 * openmp.c (gfc_match_oacc_cache): Add comment. gcc/testsuite/ PR c/71381 * c-c++-common/goacc/cache-1.c: Update. Move invalid usage tests to... * c-c++-common/goacc/cache-2.c: ... this new file. * gfortran.dg/goacc/cache-1.f95: Move invalid usage tests to... * gfortran.dg/goacc/cache-2.f95: ... this new file. * gfortran.dg/goacc/coarray.f95: Update OpenACC cache directive usage. * gfortran.dg/goacc/cray.f95: Likewise. * gfortran.dg/goacc/loop-1.f95: Likewise. libgomp/ PR c/71381 * testsuite/libgomp.oacc-c-c++-common/cache-1.c: #include "../../../gcc/testsuite/c-c++-common/goacc/cache-1.c". * testsuite/libgomp.oacc-fortran/cache-1.f95: New file. gcc/ * omp-low.c (scan_sharing_clauses): Don't expect OMP_CLAUSE__CACHE_. From-SVN: r237290
2016-06-01Remove the unused OMP_CLAUSE_DEVICE_RESIDENTThomas Schwinge1-2/+0
gcc/ * tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE_DEVICE_RESIDENT. Adjust all users. From-SVN: r236985
2016-05-16[PR 70857] Copy RESULT_DECL of HSA outlined kernel functionMartin Jambor1-0/+3
2016-05-16 Martin Jambor <mjambor@suse.cz> PR hsa/70857 * omp-low.c (grid_expand_target_grid_body): Copy RESULT_DECL of the outlined kernel function. From-SVN: r236291
2016-05-03gimplify.h (get_initialized_tmp_var): Add allow_ssa parameter default true.Richard Biener1-0/+2
2016-05-03 Richard Biener <rguenther@suse.de> * gimplify.h (get_initialized_tmp_var): Add allow_ssa parameter default true. (gimplify_arg): Likewise. * gimplify.c (gimplify_expr): Add overload with allow_ssa parameter, re-writing the result to a decl if required. (internal_get_tmp_var): Add allow_ssa parameter and override into_ssa with it. (get_formal_tmp_var): Adjust. (get_initialized_tmp_var): Add allow_ssa parameter. (gimplify_arg): Add allow_ssa parameter and avoid generating SSA names for the result false. (gimplify_call_expr): If the call may return twice do not gimplify parameters into SSA. (prepare_gimple_addressable): Do not allow an SSA name as temporary. (gimplify_modify_expr): Adjust assert. For noreturn calls with a SSA name LHS adjust its def. (gimplify_save_expr): Do not allow an SSA name as save-expr result. (gimplify_one_sizepos): Do not allow an SSA name as a sizepos. (gimplify_body): Init GIMPLE SSA data structures and gimplify into-SSA. (gimplify_scan_omp_clauses): Make sure OMP_CLAUSE_SIZE is not an SSA name. Likewise for OMP_CLAUSE_REDUCTION operands. (gimplify_omp_for): Likewise for OMP_CLAUSE_DECL. Likewise for OMP_FOR_COND, OMP_FOR_INCR and OMP_CLAUSE_LINEAR_STEP. (optimize_target_teams): Do not allow SSA names for clause operands. (gimplify_expr): Likewise for where we mark the result addressable. * passes.def (pass_init_datastructures): Remove. * tree-into-ssa.c (mark_def_sites): Ignore existing SSA names. (rewrite_stmt): Likewise. * tree-inline.c (initialize_cfun): Properly transfer SSA state. (replace_locals_op): Replace SSA names. (copy_gimple_seq_and_replace_locals): Init src_cfun. * gimple-low.c (lower_builtin_setjmp): Deal with SSA. * cgraph.c (release_function_body): Free CFG annotations only when we have a CFG. Simplify. * gimple-fold.c (gimplify_and_update_call_from_tree): Use force_gimple_operand instead of get_initialized_tmp_var. * tree-pass.h (make_pass_init_datastructures): Remove. * tree-ssa.c (execute_init_datastructures): Remove. (pass_data_init_datastructures): Likewise. (class pass_init_datastructures): Likewise. (make_pass_init_datastructures): Likewise. * omp-low.c (create_omp_child_function): Init SSA data structures. (grid_expand_target_grid_body): Likewise. * tree-cfg.c (move_block_to_fn): Double-check the DEF is an SSA name before adding it to names_to_release. (remove_bb): Always release SSA defs. * tree-ssa-ccp.c (get_default_value): Check SSA_NAME_VAR before dereferencing it. * cgraphunit.c (init_lowered_empty_function): Always int SSA data structures. * tree-ssanames.c (release_defs): Remove assert that we are in SSA form. * trans-mem.c (diagnose_tm_1): Handle SSA name function. c-family/ * cilk.c (cilk_gimplify_call_params_in_spawned_fn): Do not allow call args to gimplify to SSA names. * gcc.dg/pr30172-1.c: Adjust. * gcc.dg/pr63743.c: Likewise. * gcc.dg/tm/pr51696.c: Likewise. * c-c++-common/tm/safe-1.c: Likewise. * gcc.dg/tree-prof/val-prof-3.c: Likewise. * gcc.dg/plugin/self-assign-test-1.c: XFAIL case that needs CSE. * g++.dg/plugin/self-assign-test-1.C: Likewise. * g++.dg/plugin/self-assign-test-2.C: Likewise. From-SVN: r235817
2016-05-02omp-low.c (lower_oacc_head_tail): Assert there is at least one marker.Nathan Sidwell1-6/+2
* omp-low.c (lower_oacc_head_tail): Assert there is at least one marker. (oacc_loop_process): Check mask for loop termination. From-SVN: r235775
2016-05-02omp-low.c (struct oacc_loop): Add 'inner' field.Nathan Sidwell1-14/+38
gcc/ * omp-low.c (struct oacc_loop): Add 'inner' field. (new_oacc_loop_raw): Initialize it to zero. (oacc_loop_fixed_partitions): Initialize it. (oacc_loop_auto_partitions): Partition outermost loop to outermost available partitioning. gcc/testsuite/ * c-c++-common/goacc/loop-auto-1.c: Adjust expected warnings. libgomp/ * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Adjust expected partitioning. From-SVN: r235756
2016-04-18tree.h (TYPE_ALIGN, DECL_ALIGN): Return shifted amount.Michael Matz1-12/+12
* tree.h (TYPE_ALIGN, DECL_ALIGN): Return shifted amount. (SET_TYPE_ALIGN, SET_DECL_ALIGN): New. * tree-core.h (tree_type_common.align): Use bit-field. (tree_type_common.spare): New. (tree_decl_common.off_align): Make smaller. (tree_decl_common.align): Use bit-field. * expr.c (expand_expr_addr_expr_1): Use SET_TYPE_ALIGN. * omp-low.c (install_var_field): Use SET_DECL_ALIGN. (scan_sharing_clauses): Ditto. (finish_taskreg_scan): Use SET_DECL_ALIGN and SET_TYPE_ALIGN. (omp_finish_file): Ditto. * stor-layout.c (do_type_align): Use SET_DECL_ALIGN. (layout_decl): Ditto. (relayout_decl): Ditto. (finalize_record_size): Use SET_TYPE_ALIGN. (finalize_type_size): Ditto. (finish_builtin_struct): Ditto. (layout_type): Ditto. (initialize_sizetypes): Ditto. * targhooks.c (std_gimplify_va_arg_expr): Use SET_TYPE_ALIGN. * tree-nested.c (insert_field_into_struct): Use SET_TYPE_ALIGN. (lookup_field_for_decl): Use SET_DECL_ALIGN. (get_chain_field): Ditto. (get_trampoline_type): Ditto. (get_nl_goto_field): Ditto. * tree-streamer-in.c (unpack_ts_decl_common_value_fields): Use SET_DECL_ALIGN. (unpack_ts_type_common_value_fields): Use SET_TYPE_ALIGN. * gimple-expr.c (copy_var_decl): Use SET_DECL_ALIGN. * tree.c (make_node_stat): Use SET_DECL_ALIGN and SET_TYPE_ALIGN. (build_qualified_type): Use SET_TYPE_ALIGN. (build_aligned_type, build_range_type_1): Ditto. (build_atomic_base): Ditto. (build_common_tree_nodes): Ditto. * cfgexpand.c (align_local_variable): Use SET_DECL_ALIGN. (expand_one_stack_var_at): Ditto. * coverage.c (build_var): Use SET_DECL_ALIGN. * except.c (init_eh): Ditto. * function.c (assign_parm_setup_block): Ditto. * symtab.c (increase_alignment_1): Ditto. * tree-ssa-ccp.c (fold_builtin_alloca_with_align): Ditto. * tree-vect-stmts.c (ensure_base_align): Ditto. * varasm.c (align_variable): Ditto. (assemble_variable): Ditto. (build_constant_desc): Ditto. (output_constant_def_contents): Ditto. * config/arm/arm.c (arm_relayout_function): Use SET_DECL_ALIGN. * config/avr/avr.c (avr_adjust_type_node): Use SET_TYPE_ALIGN. * config/mips/mips.c (mips_std_gimplify_va_arg_expr): Ditto. * config/msp430/msp430.c (msp430_gimplify_va_arg_expr): Ditto. * config/spu/spu.c (spu_build_builtin_va_list): Use SET_DECL_ALIGN. ada/ * gcc-interface/decl.c (gnat_to_gnu_entity): Use SET_TYPE_ALIGN. (gnat_to_gnu_field): Ditto. (components_to_record): Ditto. (create_variant_part_from): Ditto. (copy_and_substitute_in_size): Ditto. (substitute_in_type): Ditto. * gcc-interface/utils.c (make_aligning_type): Use SET_TYPE_ALIGN. (make_packable_type): Ditto. (maybe_pad_type): Ditto. (finish_fat_pointer_type): Ditto. (finish_record_type): Ditto and use SET_DECL_ALIGN. (rest_of_record_type_compilation): Use SET_TYPE_ALIGN. (create_field_decl): Use SET_DECL_ALIGN. c-family/ * c-common.c (handle_aligned_attribute): Use SET_TYPE_ALIGN and SET_DECL_ALIGN. c/ * c-decl.c (merge_decls): Use SET_DECL_ALIGN and SET_TYPE_ALIGN. (grokdeclarator, parser_xref_tag, finish_enum): Use SET_TYPE_ALIGN. cp/ * class.c (build_vtable): Use SET_DECL_ALIGN and SET_TYPE_ALIGN. (layout_class_type): Ditto. (build_base_field): Use SET_DECL_ALIGN. (fixup_attribute_variants): Use SET_TYPE_ALIGN. * decl.c (duplicate_decls): Use SET_DECL_ALIGN. (record_unknown_type): Use SET_TYPE_ALIGN. (cxx_init_decl_processing): Ditto. (copy_type_enum): Ditto. (grokfndecl): Use SET_DECL_ALIGN. (copy_type_enum): Use SET_TYPE_ALIGN. * pt.c (instantiate_class_template_1): Use SET_TYPE_ALIGN. (tsubst): Ditto. * tree.c (cp_build_qualified_type_real): Use SET_TYPE_ALIGN. * lambda.c (maybe_add_lambda_conv_op): Use SET_DECL_ALIGN. * method.c (implicitly_declare_fn): Use SET_DECL_ALIGN. * rtti.c (emit_tinfo_decl): Ditto. fortran/ * trans-io.c (gfc_build_io_library_fndecls): Use SET_TYPE_ALIGN. * trans-common.c (build_common_decl): Use SET_DECL_ALIGN. * trans-types.c (gfc_add_field_to_struct): Use SET_DECL_ALIGN. go/ * go-gcc.cc (Gcc_backend::implicit_variable): Use SET_DECL_ALIGN. java/ * class.c (add_method_1): Use SET_DECL_ALIGN. (make_class_data): Ditto. (emit_register_classes_in_jcr_section): Ditto. * typeck.c (build_java_array_type): Ditto. objc/ * objc-act.c (objc_build_struct): Use SET_DECL_ALIGN. libcc1/ * plugin.cc (plugin_finish_record_or_union): Use SET_TYPE_ALIGN. From-SVN: r235172
2016-04-15Split out OMP constructs' SIMD clone supporting codeThomas Schwinge1-1606/+0
gcc/ * omp-low.c (simd_clone_struct_alloc, simd_clone_struct_copy) (simd_clone_vector_of_formal_parm_types) (simd_clone_clauses_extract, simd_clone_compute_base_data_type) (simd_clone_mangle, simd_clone_create) (simd_clone_adjust_return_type, create_tmp_simd_array) (simd_clone_adjust_argument_types, simd_clone_init_simd_arrays) (struct modify_stmt_info, ipa_simd_modify_stmt_ops) (ipa_simd_modify_function_body, simd_clone_linear_addend) (simd_clone_adjust, expand_simd_clones, ipa_omp_simd_clone) (pass_data_omp_simd_clone, class pass_omp_simd_clone) (pass_omp_simd_clone::gate, make_pass_omp_simd_clone): Move into... * omp-simd-clone.c: ... this new file. (simd_clone_vector_of_formal_parm_types): Make it static. * Makefile.in (OBJS): Add omp-simd-clone.o. From-SVN: r235017
2016-04-14re PR middle-end/70643 (broken openacc reduction inside a fortran module)Cesar Philippidis1-1/+1
PR middle-end/70643 gcc/ * omp-low.c (lower_oacc_reductions): Check for TREE_CONSTANT when building a mem ref for the incoming reduction variable. libgomp/ * testsuite/libgomp.oacc-fortran/pr70643.f90: New test. From-SVN: r234973
2016-04-12omp-low.c (lower_omp_target): Use GOMP_MAP_FIRSTPRIVATE_INT regardless ↵Jakub Jelinek1-69/+1
whether there are depend clauses or not. * omp-low.c (lower_omp_target): Use GOMP_MAP_FIRSTPRIVATE_INT regardless whether there are depend clauses or not. * libgomp.h (struct gomp_target_task): Remove firstprivate_copies field. * target.c (gomp_target_fallback_firstprivate, gomp_target_unshare_firstprivate): Removed. (GOMP_target_ext): Copy firstprivate vars into gomp_allocaed memory before waiting for dependencies. (gomp_target_task_fn): Don't copy firstprivate vars here. * task.c (GOMP_PLUGIN_target_task_completion): Don't free firstprivate_copies here. (gomp_create_target_task): Don't initialize firstprivate_copies field. * testsuite/libgomp.c/target-25.c (main): Use map (to:) instead of explicit/implicit firstprivate. From-SVN: r234894
2016-04-08re PR lto/70289 ([openacc] ICE in input_varpool_node)Cesar Philippidis1-17/+65
gcc/ PR lto/70289 PR ipa/70348 PR tree-optimization/70373 PR middle-end/70533 PR middle-end/70534 PR middle-end/70535 * gimplify.c (gimplify_adjust_omp_clauses): Add or adjust data clauses for acc parallel reductions as necessary. Error on those that are private. * omp-low.c (scan_sharing_clauses): Don't install variables which are used in acc parallel reductions. (lower_rec_input_clauses): Remove dead code. (lower_oacc_reductions): Add support for reference reductions. (lower_reduction_clauses): Remove dead code. (lower_omp_target): Don't remap variables appearing in acc parallel reductions. * tree.h (OMP_CLAUSE_MAP_IN_REDUCTION): New macro. gcc/testsuite/ * c-c++-common/goacc/reduction-5.c: New test. * c-c++-common/goacc/reduction-promotions.c: New test. * gfortran.dg/goacc/reduction-3.f95: New test. * gfortran.dg/goacc/reduction-promotions.f90: New test. libgomp/ * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gang-np-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gw-np-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-reduction-gwv-np-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-reduction-vector-p-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-reduction-worker-p-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/loop-reduction-wv-p-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c: Add test coverage. * testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr70289.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr70373.c: New test. * testsuite/libgomp.oacc-c-c++-common/reduction-1.c: Add test coverage. * testsuite/libgomp.oacc-c-c++-common/reduction-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-6.c: New test. * testsuite/libgomp.oacc-c-c++-common/reduction.h: New test. * testsuite/libgomp.oacc-fortran/parallel-reduction.f90: New test. * testsuite/libgomp.oacc-fortran/pr70289.f90: New test. * testsuite/libgomp.oacc-fortran/reduction-1.f90: Add test coverage. * testsuite/libgomp.oacc-fortran/reduction-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-4.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-5.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-6.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-7.f90: New test. From-SVN: r234840
2016-04-08Remove incorrect warning for parallel implicit firstprivate clauseTom de Vries1-1/+6
2016-04-08 Tom de Vries <tom@codesourcery.com> * omp-low.c (lower_omp_target): Set TREE_NO_WARNING for oacc implicit firstprivate clause. * c-c++-common/goacc/uninit-firstprivate-clause.c: New test. * gfortran.dg/goacc/uninit-firstprivate-clause.f95: New test. From-SVN: r234826
2016-04-07cgraph.h (struct cgraph_simd_clone): Add mask_mode field.Jakub Jelinek1-41/+119
* cgraph.h (struct cgraph_simd_clone): Add mask_mode field. * omp-low.c (simd_clone_init_simd_arrays, simd_clone_adjust): Handle node->simdclone->mask_mode != VOIDmode masks. (simd_clone_adjust_argument_types): Likewise. Move sc var definition earlier, use it instead of node->simdclone. * config/i386/i386.c (ix86_simd_clone_compute_vecsize_and_simdlen): Set clonei->mask_mode. * c-c++-common/attr-simd.c: Add scan-assembler* directives for AVX512F clones. * c-c++-common/attr-simd-2.c: Likewise. * c-c++-common/attr-simd-4.c: Likewise. * gcc.dg/gomp/simd-clones-2.c: Likewise. * gcc.dg/gomp/simd-clones-3.c: Likewise. From-SVN: r234816
2016-04-06Fix new -Wparentheses warnings encountered during bootstrapPatrick Palka1-13/+15
gcc/ChangeLog: PR c/70436 * gimplify.c (gimplify_omp_ordered): Add explicit braces to resolve a future -Wparentheses warning. * omp-low.c (scan_sharing_clauses): Likewise. * tree-parloops.c (eliminate_local_variables): Likewise. gcc/cp/ChangeLog: PR c/70436 * cp-tree.h (FOR_EACH_CLONE): Restructure macro to avoid potentially generating a future -Wparentheses warning in its callers. gcc/fortran/ChangeLog: PR c/70436 * openmp.c (gfc_find_omp_udr): Add explicit braces to resolve a future -Wparentheses warning. gcc/testsuite/ChangeLog: PR c/70436 * g++.dg/plugin/pragma_plugin.c (handle_pragma_sayhello): Add explicit braces to resolve a future -Wparentheses warning. From-SVN: r234801
2016-04-06re PR middle-end/70550 (-Wuninitialized false positives in OpenMP code)Jakub Jelinek1-2/+22
PR middle-end/70550 * tree.h (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT): Define. * gimplify.c (gimplify_adjust_omp_clauses_1): Set it for implicit firstprivate clauses. * omp-low.c (lower_send_clauses): Set TREE_NO_WARNING for OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT !by_ref vars in task contexts. (lower_omp_target): Set TREE_NO_WARNING for non-addressable possibly uninitialized vars which are copied into addressable temporaries or copied for GOMP_MAP_FIRSTPRIVATE_INT. * c-c++-common/gomp/pr70550-1.c: New test. * c-c++-common/gomp/pr70550-2.c: New test. From-SVN: r234779
2016-03-17Rename GOMP_MAP_FORCE_DEALLOC to GOMP_MAP_DELETEThomas Schwinge1-1/+1
Also rename the Fortran OMP_MAP_FORCE_DEALLOC to OMP_MAP_DELETE. include/ * gomp-constants.h (enum gomp_map_kind): Rename GOMP_MAP_FORCE_DEALLOC to GOMP_MAP_DELETE. Adjust all users. gcc/fortran/ * gfortran.h (enum gfc_omp_map_op): Rename OMP_MAP_FORCE_DEALLOC to OMP_MAP_DELETE. Adjust all users. From-SVN: r234294
2016-03-07re PR middle-end/69916 ([openacc] ICE in single_succ_edge called from ↵Nathan Sidwell1-40/+62
oacc_loop_xform_loop) gcc/ PR middle-end/69916 * omp-low.c (struct oacc_loop): Add ifns. (new_oacc_loop_raw): Initialize it. (finish_oacc_loop): Clear mask & flags if no ifns. (oacc_loop_discover_walk): Count IFN_GOACC_LOOP calls. (oacc_loop_xform_loop): Add ifns arg & adjust. (oacc_loop_process): Adjust oacc_loop_xform_loop call. gcc/testsuite/ PR middle-end/69916 * c-c-++-common/goacc/pr69916.c: New. From-SVN: r234026
2016-03-05Handle oacc region in oacc routineTom de Vries1-13/+28
2016-03-05 Tom de Vries <tom@codesourcery.com> * omp-low.c (check_omp_nesting_restrictions): Check for non-oacc construct in oacc routine. Check for oacc region in oacc routine. * c-c++-common/goacc/nesting-fail-1.c (f_acc_routine): New function. * c-c++-common/goacc-gomp/nesting-fail-1.c (f_acc_routine): New function. From-SVN: r233998
2016-03-02re PR libgomp/69555 (libgomp.c++/target-6.C fails because of undefined ↵Jakub Jelinek1-8/+24
behaviour) PR libgomp/69555 * gimplify.c (gimplify_decl_expr): For decls with REFERENCE_TYPE, also gimplify_type_sizes the type they refer to. (omp_notice_variable): Handle reference vars to VLAs. * omp-low.c (lower_omp_target): Emit setup of OMP_CLAUSE_PRIVATE reference to VLA decls in the second pass instead of first pass. * testsuite/libgomp.c++/pr69555-1.C: New test. * testsuite/libgomp.c++/pr69555-2.C: New test. From-SVN: r233913
2016-02-26[omp, hsa] Do not gridify simd constructsMartin Jambor1-4/+22
2016-02-26 Martin Jambor <mjambor@suse.cz> * omp-low.c (grid_find_ungridifiable_statement): Store problematic statements to wi->info. Also disallow omp simd constructs. (grid_target_follows_gridifiable_pattern): Use wi.info to dump reason for not gridifying. Dump special string for omp_for. From-SVN: r233746
2016-02-16Don't call call_cgraph_insertion_hooks in simd_clone_createTom de Vries1-1/+0
2016-02-16 Tom de Vries <tom@codesourcery.com> PR lto/67709 * omp-low.c (simd_clone_create): Remove call to symtab->call_cgraph_insertion_hooks. * testsuite/libgomp.fortran/declare-simd-4.f90: New test. From-SVN: r233447
2016-02-02Merge BUILT_IN_GOACC_HOST_DATA into BUILT_IN_GOACC_DATA_STARTThomas Schwinge1-4/+1
gcc/ * omp-builtins.def (BUILT_IN_GOACC_HOST_DATA): Remove. * omp-low.c (expand_omp_target): Use BUILT_IN_GOACC_DATA_START instead. libgomp/ * libgomp.map (GOACC_2.0): Remove GOACC_host_data. * oacc-parallel.c (GOACC_host_data): Remove function definition. From-SVN: r233074
2016-02-01omp-low.c (oacc_parse_default_dims): Avoid -Wsign-compare warning...Jakub Jelinek1-2/+2
* omp-low.c (oacc_parse_default_dims): Avoid -Wsign-compare warning, make sure value fits into int rather than just unsigned int. From-SVN: r233044
2016-02-01nvptx.c (PTX_GANG_DEFAULT): New.Nathan Sidwell1-24/+128
gcc/ * config/nvptx/nvptx.c (PTX_GANG_DEFAULT): New. (nvptx_goacc_validate_dims): Extend to handle global defaults. * target.def (OACC_VALIDATE_DIMS): Extend documentation. * doc/tm.texti: Rebuilt. * doc/invoke.texi (fopenacc-dim): Document. * lto-wrapper.c (merge_and_complain): Add OPT_fopenacc_dim_ case. (append_compiler_options): Likewise. * omp-low.c (oacc_default_dims, oacc_min_dims): New. (oacc_parse_default_dims): New. (oacc_validate_dims): Add USED arg. Select non-unity default when possible. (oacc_loop_fixed_partitions): Return mask of used partitions. (oacc_loop_auto_partitions): Emit dump info. (oacc_loop_partition): Return mask of used partitions. (execute_oacc_device_lower): Parse default dimension arg. Adjust loop partitioning and validation calls. gcc/c-family/ * c.opt (fopenacc-dim=): New option. gcc/fortran/ * lang.opt (fopenacc-dim=): New option. libgomp/ * testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c: New. * testsuite/libgomp.oacc-fortran/routine-7.f90: Serialize loop. From-SVN: r233041
2016-01-26omp-low.h (oacc_fn_attrib_kernels_p): Declare.Nathan Sidwell1-30/+76
* omp-low.h (oacc_fn_attrib_kernels_p): Declare. (set_oacc_fn_attrib): Add IS_KERNEL arg. * omp-low.c (set_oacc_fn_attrib): Add IS_KERNEL arg. (oacc_fn_attrib_kernels_p, oacc_fn_attrib_level): New. (expand_omp_target): Pass is_kernel to set_oacc_fn_attrib. (oacc_validate_dims): Add LEVEL arg, don't return level. (new_oacc_loop_routine): Use oacc_fn_attrib_level, not oacc_validate_dims. (execute_oacc_device_lower): Adjust, add more dump output. * tree-ssa-loop.c (gate_oacc_kernels): Use oacc_fn_attrib_kernels_p. * tree-parloops.c (create_parallel_loop): Adjust set_oacc_fn_attrib call. From-SVN: r232830
2016-01-25omp-low.c (lower_omp_target): Set DECL_VALUE_EXPR of new_var even for the ↵Jakub Jelinek1-0/+9
non-array case. * omp-low.c (lower_omp_target) <case USE_DEVICE_PTR>: Set DECL_VALUE_EXPR of new_var even for the non-array case. Look through DECL_VALUE_EXPR for expansion. * c-c++-common/goacc/use_device-1.c: New test. From-SVN: r232804
2016-01-21omp-low.c (expand_omp_target): Avoid -Wmaybe-uninitialized warning.Jakub Jelinek1-6/+5
* omp-low.c (expand_omp_target): Avoid -Wmaybe-uninitialized warning. Fix up formatting. From-SVN: r232641
2016-01-19Merge of HSAMartin Jambor1-95/+1361
2016-01-19 Martin Jambor <mjambor@suse.cz> Martin Liska <mliska@suse.cz> Michael Matz <matz@suse.de> libgomp/ * plugin/Makefrag.am: Add HSA plugin requirements. * plugin/configfrag.ac (HSA_RUNTIME_INCLUDE): New variable. (HSA_RUNTIME_LIB): Likewise. (HSA_RUNTIME_CPPFLAGS): Likewise. (HSA_RUNTIME_INCLUDE): New substitution. (HSA_RUNTIME_LIB): Likewise. (HSA_RUNTIME_LDFLAGS): Likewise. (hsa-runtime): New configure option. (hsa-runtime-include): Likewise. (hsa-runtime-lib): Likewise. (PLUGIN_HSA): New substitution variable. Fill HSA_RUNTIME_INCLUDE and HSA_RUNTIME_LIB according to the new configure options. (PLUGIN_HSA_CPPFLAGS): Likewise. (PLUGIN_HSA_LDFLAGS): Likewise. (PLUGIN_HSA_LIBS): Likewise. Check that we have access to HSA run-time. * libgomp-plugin.h (offload_target_type): New element OFFLOAD_TARGET_TYPE_HSA. * libgomp.h (gomp_target_task): New fields firstprivate_copies and args. (bool gomp_create_target_task): Updated. (gomp_device_descr): Extra parameter of run_func and async_run_func, new field can_run_func. * libgomp_g.h (GOMP_target_ext): Update prototype. * oacc-host.c (host_run): Added a new parameter args. * target.c (calculate_firstprivate_requirements): New function. (copy_firstprivate_data): Likewise. (gomp_target_fallback_firstprivate): Use them. (gomp_target_unshare_firstprivate): New function. (gomp_get_target_fn_addr): Allow returning NULL for shared memory devices. (GOMP_target): Do host fallback for all shared memory devices. Do not pass any args to plugins. (GOMP_target_ext): Introduce device-specific argument parameter args. Allow host fallback if device shares memory. Do not remap data if device has shared memory. (gomp_target_task_fn): Likewise. Also treat shared memory devices like host fallback for mappings. (GOMP_target_data): Treat shared memory devices like host fallback. (GOMP_target_data_ext): Likewise. (GOMP_target_update): Likewise. (GOMP_target_update_ext): Likewise. Also pass NULL as args to gomp_create_target_task. (GOMP_target_enter_exit_data): Likewise. (omp_target_alloc): Treat shared memory devices like host fallback. (omp_target_free): Likewise. (omp_target_is_present): Likewise. (omp_target_memcpy): Likewise. (omp_target_memcpy_rect): Likewise. (omp_target_associate_ptr): Likewise. (gomp_load_plugin_for_device): Also load can_run. * task.c (GOMP_PLUGIN_target_task_completion): Free firstprivate_copies. (gomp_create_target_task): Accept new argument args and store it to ttask. * plugin/plugin-hsa.c: New file. gcc/ * Makefile.in (OBJS): Add new source files. (GTFILES): Add hsa.c. * common.opt (disable_hsa): New variable. (-Whsa): New warning. * config.in (ENABLE_HSA): New. * configure.ac: Treat hsa differently from other accelerators. (OFFLOAD_TARGETS): Define ENABLE_OFFLOADING according to $enable_offloading. (ENABLE_HSA): Define ENABLE_HSA according to $enable_hsa. * doc/install.texi (Configuration): Document --with-hsa-runtime, --with-hsa-runtime-include, --with-hsa-runtime-lib and --with-hsa-kmt-lib. * doc/invoke.texi (-Whsa): Document. (hsa-gen-debug-stores): Likewise. * lto-wrapper.c (compile_images_for_offload_targets): Do not attempt to invoke offload compiler for hsa acclerator. * opts.c (common_handle_option): Determine whether HSA offloading should be performed. * params.def (PARAM_HSA_GEN_DEBUG_STORES): New parameter. * builtin-types.def (BT_FN_VOID_UINT_PTR_INT_PTR): New. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT): Removed. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR): New. * gimple-low.c (lower_stmt): Also handle GIMPLE_OMP_GRID_BODY. * gimple-pretty-print.c (dump_gimple_omp_for): Also handle GF_OMP_FOR_KIND_GRID_LOOP. (dump_gimple_omp_block): Also handle GIMPLE_OMP_GRID_BODY. (pp_gimple_stmt_1): Likewise. * gimple-walk.c (walk_gimple_stmt): Likewise. * gimple.c (gimple_build_omp_grid_body): New function. (gimple_copy): Also handle GIMPLE_OMP_GRID_BODY. * gimple.def (GIMPLE_OMP_GRID_BODY): New. * gimple.h (enum gf_mask): Added GF_OMP_PARALLEL_GRID_PHONY, GF_OMP_FOR_KIND_GRID_LOOP, GF_OMP_FOR_GRID_PHONY and GF_OMP_TEAMS_GRID_PHONY. (gimple_statement_omp_single_layout): Updated comments. (gimple_build_omp_grid_body): New function. (gimple_has_substatements): Also handle GIMPLE_OMP_GRID_BODY. (gimple_omp_for_grid_phony): New function. (gimple_omp_for_set_grid_phony): Likewise. (gimple_omp_parallel_grid_phony): Likewise. (gimple_omp_parallel_set_grid_phony): Likewise. (gimple_omp_teams_grid_phony): Likewise. (gimple_omp_teams_set_grid_phony): Likewise. (gimple_return_set_retbnd): Also handle GIMPLE_OMP_GRID_BODY. * omp-builtins.def (BUILT_IN_GOMP_OFFLOAD_REGISTER): New. (BUILT_IN_GOMP_OFFLOAD_UNREGISTER): Likewise. (BUILT_IN_GOMP_TARGET): Updated type. * omp-low.c: Include symbol-summary.h, hsa.h and params.h. (adjust_for_condition): New function. (get_omp_for_step_from_incr): Likewise. (extract_omp_for_data): Moved parts to adjust_for_condition and get_omp_for_step_from_incr. (build_outer_var_ref): Handle GIMPLE_OMP_GRID_BODY. (fixup_child_record_type): Bail out if receiver_decl is NULL. (scan_sharing_clauses): Handle OMP_CLAUSE__GRIDDIM_. (scan_omp_parallel): Do not create child functions for phony constructs. (check_omp_nesting_restrictions): Handle GIMPLE_OMP_GRID_BODY. (scan_omp_1_op): Checking assert we are not remapping to ERROR_MARK. Also also handle GIMPLE_OMP_GRID_BODY. (parallel_needs_hsa_kernel_p): New function. (expand_parallel_call): Register apprpriate parallel child functions as HSA kernels. (grid_launch_attributes_trees): New type. (grid_attr_trees): New variable. (grid_create_kernel_launch_attr_types): New function. (grid_insert_store_range_dim): Likewise. (grid_get_kernel_launch_attributes): Likewise. (get_target_argument_identifier_1): Likewise. (get_target_argument_identifier): Likewise. (get_target_argument_value): Likewise. (push_target_argument_according_to_value): Likewise. (get_target_arguments): Likewise. (expand_omp_target): Call get_target_arguments instead of looking up for teams and thread limit. (grid_expand_omp_for_loop): New function. (grid_arg_decl_map): New type. (grid_remap_kernel_arg_accesses): New function. (grid_expand_target_kernel_body): New function. (expand_omp): Call it. (lower_omp_for): Do not emit phony constructs. (lower_omp_taskreg): Do not emit phony constructs but create for them a temporary variable receiver_decl. (lower_omp_taskreg): Do not emit phony constructs. (lower_omp_teams): Likewise. (lower_omp_grid_body): New function. (lower_omp_1): Call it. (grid_reg_assignment_to_local_var_p): New function. (grid_seq_only_contains_local_assignments): Likewise. (grid_find_single_omp_among_assignments_1): Likewise. (grid_find_single_omp_among_assignments): Likewise. (grid_find_ungridifiable_statement): Likewise. (grid_target_follows_gridifiable_pattern): Likewise. (grid_remap_prebody_decls): Likewise. (grid_copy_leading_local_assignments): Likewise. (grid_process_kernel_body_copy): Likewise. (grid_attempt_target_gridification): Likewise. (grid_gridify_all_targets_stmt): Likewise. (grid_gridify_all_targets): Likewise. (execute_lower_omp): Call grid_gridify_all_targets. (make_gimple_omp_edges): Handle GIMPLE_OMP_GRID_BODY. * tree-core.h (omp_clause_code): Added OMP_CLAUSE__GRIDDIM_. (tree_omp_clause): Added union field dimension. * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE__GRIDDIM_. * tree.c (omp_clause_num_ops): Added number of arguments of OMP_CLAUSE__GRIDDIM_. (omp_clause_code_name): Added name of OMP_CLAUSE__GRIDDIM_. (walk_tree_1): Handle OMP_CLAUSE__GRIDDIM_. * tree.h (OMP_CLAUSE_GRIDDIM_DIMENSION): New. (OMP_CLAUSE_SET_GRIDDIM_DIMENSION): Likewise. (OMP_CLAUSE_GRIDDIM_SIZE): Likewise. (OMP_CLAUSE_GRIDDIM_GROUP): Likewise. * passes.def: Schedule pass_ipa_hsa and pass_gen_hsail. * tree-pass.h (make_pass_gen_hsail): Declare. (make_pass_ipa_hsa): Likewise. * ipa-hsa.c: New file. * lto-section-in.c (lto_section_name): Add hsa section name. * lto-streamer.h (lto_section_type): Add hsa section. * timevar.def (TV_IPA_HSA): New. * hsa-brig-format.h: New file. * hsa-brig.c: New file. * hsa-dump.c: Likewise. * hsa-gen.c: Likewise. * hsa.c: Likewise. * hsa.h: Likewise. * toplev.c (compile_file): Call hsa_output_brig. * hsa-regalloc.c: New file. gcc/fortran/ * types.def (BT_FN_VOID_UINT_PTR_INT_PTR): New. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT): Removed. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR): New. gcc/lto/ * lto-partition.c: Include "hsa.h" (add_symbol_to_partition_1): Put hsa implementations into the same partition as host implementations. liboffloadmic/ * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_async_run): New unused parameter. (GOMP_OFFLOAD_run): Likewise. include/ * gomp-constants.h (GOMP_DEVICE_HSA): New macro. (GOMP_VERSION_HSA): Likewise. (GOMP_TARGET_ARG_DEVICE_MASK): Likewise. (GOMP_TARGET_ARG_DEVICE_ALL): Likewise. (GOMP_TARGET_ARG_SUBSEQUENT_PARAM): Likewise. (GOMP_TARGET_ARG_ID_MASK): Likewise. (GOMP_TARGET_ARG_NUM_TEAMS): Likewise. (GOMP_TARGET_ARG_THREAD_LIMIT): Likewise. (GOMP_TARGET_ARG_VALUE_SHIFT): Likewise. (GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES): Likewise. From-SVN: r232549
2016-01-18Add oacc_kernels_p argument to pass_parallelize_loopsTom de Vries1-1/+1
2016-01-18 Tom de Vries <tom@codesourcery.com> * omp-low.c (set_oacc_fn_attrib): Make extern. * omp-low.h (set_oacc_fn_attrib): Declare. * tree-parloops.c (struct reduction_info): Add reduc_addr field. (create_call_for_reduction_1): Handle case that reduc_addr is non-NULL. (create_parallel_loop, gen_parallel_loop, try_create_reduction_list): Add and handle function parameter oacc_kernels_p. (find_reduc_addr, get_omp_data_i_param): New function. (ref_conflicts_with_region, oacc_entry_exit_ok_1) (oacc_entry_exit_single_gang, oacc_entry_exit_ok): New function. (parallelize_loops): Add and handle function parameter oacc_kernels_p. Calculate dominance info. Skip loops that are not in a kernels region in oacc_kernels_p mode. Skip inner loops of parallelized loops. (pass_parallelize_loops::execute): Call parallelize_loops with oacc_kernels_p argument. (pass_parallelize_loops::clone, pass_parallelize_loops::set_pass_param): New member function. (pass_parallelize_loops::bool oacc_kernels_p): New member var. * passes.def: Add argument to pass_parallelize_loops instantation. From-SVN: r232512
2016-01-17omp-low.c (mark_loops_in_oacc_kernels_region): Work around ↵Jakub Jelinek1-1/+1
-Wmaybe-uninitialized warning. * omp-low.c (mark_loops_in_oacc_kernels_region): Work around -Wmaybe-uninitialized warning. From-SVN: r232484
2016-01-16Release_defs in expand_omp_atomic_fetch_opTom de Vries1-1/+5
2016-01-16 Tom de Vries <tom@codesourcery.com> * omp-low.c (expand_omp_atomic_fetch_op): Release defs of update stmt. From-SVN: r232472
2016-01-14Mark symbols in offload tables with force_output in read_offload_tablesTom de Vries1-5/+0
2016-01-14 Tom de Vries <tom@codesourcery.com> PR tree-optimization/68773 * c-parser.c (c_parser_oacc_declare, c_parser_omp_declare_target): Don't set force_output. * parser.c (cp_parser_oacc_declare, cp_parser_omp_declare_target): Don't set force_output. * omp-low.c (expand_omp_target): Don't set force_output. * varpool.c (varpool_node::get_create): Same. * lto-cgraph.c (input_offload_tables): Mark entries in offload_vars and offload_funcs with force_output. From-SVN: r232384
2016-01-04Update copyright years.Jakub Jelinek1-1/+1
From-SVN: r232055
2015-12-23Merge OMP_CLAUSE_USE_DEVICE into OMP_CLAUSE_USE_DEVICE_PTRThomas Schwinge1-9/+2
gcc/c/ * c-parser.c (c_parser_oacc_clause_use_device): Merge function into... (c_parser_omp_clause_use_device_ptr): ... this function. Adjust all users. gcc/ * tree-core.h (enum omp_clause_code): Merge OMP_CLAUSE_USE_DEVICE into OMP_CLAUSE_USE_DEVICE_PTR. Adjust all users. From-SVN: r231926
2015-12-15c-common.c (c_common_attribute_table): Handle "omp declare target link" ↵Ilya Verbin1-8/+123
attribute. gcc/c-family/ * c-common.c (c_common_attribute_table): Handle "omp declare target link" attribute. gcc/ * cgraphunit.c (output_in_order): Do not assemble "omp declare target link" variables in ACCEL_COMPILER. * gimplify.c (gimplify_adjust_omp_clauses): Do not remove mapping of "omp declare target link" variables. * omp-low.c (scan_sharing_clauses): Do not remove mapping of "omp declare target link" variables. (add_decls_addresses_to_decl_constructor): For "omp declare target link" variables output address of the artificial pointer instead of address of the variable. Set most significant bit of the size to mark them. (pass_data_omp_target_link): New pass_data. (pass_omp_target_link): New class. (find_link_var_op): New static function. (make_pass_omp_target_link): New function. * passes.def: Add pass_omp_target_link. * tree-pass.h (make_pass_omp_target_link): Declare. * varpool.c (symbol_table::output_variables): Do not assemble "omp declare target link" variables in ACCEL_COMPILER. gcc/lto/ * lto.c: Include stringpool.h and fold-const.h. (offload_handle_link_vars): New static function. (lto_main): Call offload_handle_link_vars. libgomp/ * libgomp.h (REFCOUNT_LINK): Define. (struct splay_tree_key_s): Add link_key. * target.c (gomp_map_vars): Treat REFCOUNT_LINK objects as not mapped. Replace target address of the pointer with target address of newly mapped object in the splay tree. Set link pointer on target to the device address of the mapped object. (gomp_unmap_vars): Restore target address of the pointer in the splay tree for REFCOUNT_LINK objects after unmapping. (gomp_load_image_to_device): Set refcount to REFCOUNT_LINK for "omp declare target link" objects. (gomp_unload_image_from_device): Replace j with i. Force unmap of all "omp declare target link" objects, which were mapped for the image. (gomp_exit_data): Restore target address of the pointer in the splay tree for REFCOUNT_LINK objects after unmapping. * testsuite/libgomp.c/target-link-1.c: New file. From-SVN: r231655
2015-12-08Clear restrict in install_var_fieldTom de Vries1-0/+7
2015-12-08 Tom de Vries <tom@codesourcery.com> PR tree-optimization/68640 * omp-low.c (install_var_field): Clear the restrict qualifier on the var type. From-SVN: r231411
2015-12-02Mark pointers to allocated target vars as restricted, if possibleTom de Vries1-7/+86
2015-12-02 Tom de Vries <tom@codesourcery.com> * omp-low.c (install_var_field, scan_sharing_clauses): Add and handle parameter base_pointers_restrict. (omp_target_base_pointers_restrict_p): New function. (scan_omp_target): Call scan_sharing_clauses with base_pointers_restrict arg. * c-c++-common/goacc/kernels-alias-2.c: New test. * c-c++-common/goacc/kernels-alias-3.c: New test. * c-c++-common/goacc/kernels-alias-4.c: New test. * c-c++-common/goacc/kernels-alias-5.c: New test. * c-c++-common/goacc/kernels-alias-6.c: New test. * c-c++-common/goacc/kernels-alias-7.c: New test. * c-c++-common/goacc/kernels-alias-8.c: New test. * c-c++-common/goacc/kernels-alias.c: New test. From-SVN: r231182
generated by cgit v1.1 at 2025-08-04 03:39:39 +0000