| blob | f2b6c4bfc74f6bb7a0e30b4385efc175690b8374 |
1 2015-06-29 James Norris <jnorris@codesourcery.com>
3 * omp-low.c (expand_omp_target): Remove unnecessary conversion.
5 2015-06-26 Nathan Sidwell <nathan@codesourcery.com>
7 * config/nvptx/nvptx.md (UNSPEC_ID): Rename to ...
8 (UNSPECV_ID): ... here.
9 (oacc_id): Make volatile.
10 * omp-builtins.def (BUILT_IN_GOACC_ID): Not a constant function.
12 2015-06-24 James Norris <jnorris@codesourcery.com>
14 * c/c-parser.c (finish_oacc_declare): Additional syntax checking
15 and add restriction.
16 * cp/decl.c (finish_oacc_declare): Handle null statement list and
17 add restriction.
18 * cp/parser.c (cp_parser_oacc_declare): Additional syntax checking.
20 2015-06-24 Tom de Vries <tom@codesourcery.com>
22 * tree-ssa-loop-manip.c (get_virtual_phi): Factor out of ...
23 (rewrite_virtuals_into_loop_closed_ssa): ... here.
25 2015-06-24 Tom de Vries <tom@codesourcery.com>
27 * tree-ssa-loop-manip.c (replace_uses_in_dominated_bbs): Factor out
28 of ...
29 (rewrite_virtuals_into_loop_closed_ssa): ... here.
31 2015-06-24 Tom de Vries <tom@codesourcery.com>
33 * dominance.c (bitmap_get_dominated_by): New function.
34 * dominance.h (bitmap_get_dominated_by): Declare.
35 * tree-ssa-loop-manip.c (rewrite_virtuals_into_loop_closed_ssa): Use
36 bitmap_get_dominated_by.
38 2015-06-24 Tom de Vries <tom@codesourcery.com>
40 * tree-parloops.c (replace_uses_in_bbs_by)
41 (rewrite_virtuals_into_loop_closed_ssa): Move to ...
42 * tree-ssa-loop-manip.c: here.
43 * tree-ssa-loop-manip.h (rewrite_virtuals_into_loop_closed_ssa):
44 Declare.
46 2015-06-23 Nathan Sidwell <nathan@codesourcery.com>
48 Port from trunk.
49 * config/nvptx/nvptx.md (sel_true<mode>, sel_false<mode>): New
50 conditional selects.
51 (setcc_int<mode>, setcc_float<mode>): Reformat.
53 2015-06-22 Nathan Sidwell <nathan@codesourcery.com>
55 * omp-low.c (expand_oacc_get_num_threads): Use OACC enum.
56 (expand_oacc_get_thread_num, make_predication_test): Likewise.
57 * builtins.c (expand_oacc_id): Use UINTVAL.
59 * omp-builtins.def (BUILT_IN_GOACC_NTID, BUILTIN_NCTAID): Replace
60 with ...
61 (BUILT_IN_GOACC_NID): ... this.
62 (BUILT_IN_GOACC_TID, BUILTIN_CTAID): Replace with ...
63 (BUILT_IN_GOACC_ID): ... this.
64 * builtins.c: Include omp-low.h.
65 (expand_oacc_buoltin): Replace with ...
66 (expand_oacc_id): ... this.
67 (expand_builtin, is_simple_builtin): Adjust.oo
68 * omp-low.h (enum oacc_loop_levels): New.
69 * omp-low.c (MASK_GANG, MASK_WORKER, MASK_VECTOR): Replace with ...
70 (OACC_LOOP_MASK): ... this.
71 (scan_omp_for, scan_omp_target): Adjust.
72 (expand_oacc_get_num_threads): Adjust and use a loop.
73 (expand_oacc_get_thread_num): Likewise.
74 (oacc_loop_needs_thread_barrier_p, find_omp_for_region_gwv,
75 find_omp_taarget_region_data, required_predication_mask,
76 generate_vector_broadcast, generate_oacc_broadcast): Adjust.
77 (make_predication_test): Adjust and use a loop.
78 (predicate_bb, oacc_broadcast, oacc_init_count_vars): Adjust.
79 * config/nvptx/nvptx.md (UNSPEC_NTID, UNSPEC_TID, UNSPEC_NCTAID,
80 UNSPEC_CTAID): Replace with ...
81 (UNSPEC_NID, UNSPEC_ID): ... these.
82 (*oacc_ntid_insn, oacc_ntid, *oacc_tid_insn, oacc_tid,
83 *oacc_nctaid_insn, oacc_nctaid, *oacc_ctaid_insn,
84 oacc_ctaid): Replace with ...
85 (oacc_nid, oacc_id): ... these.
86 * config/nvptx/nvptx.c (nvptx_print_operand [CASE 'd']): Remove.
88 2015-06-19 Tom de Vries <tom@codesourcery.com>
90 * tree-parloops.c (struct reduction_info): Add reduc_addr field.
91 (create_call_for_reduction_1): Handle case that reduc_addr is non-NULL.
92 (gen_parallel_loop): Init clsn_data for oacc_kernels_p case.
93 (try_create_reduction_list): Add and handle oacc_kernels_p parameter.
94 (parallelize_loops): Add argument to call to try_create_reduction_list.
96 2015-06-18 Tom de Vries <tom@codesourcery.com>
98 * passes.def: Move pass_fre later in pass group pass_oacc_kernels,
99 replacing pass_copy_prop.
101 2015-06-18 Tom de Vries <tom@codesourcery.com>
103 * tree-parloops.c (rewrite_virtuals_into_loop_closed_ssa): New function.
104 (transform_to_exit_first_loop_alt): Use
105 rewrite_virtuals_into_loop_closed_ssa.
107 2015-06-18 Tom de Vries <tom@codesourcery.com>
109 * tree-parloops.c (transform_to_exit_first_loop_alt): Add update_stmt
110 for cond_stmt.
111 (try_get_loop_niter): Declare forward.
112 (try_transform_to_exit_first_loop_alt): Use try_get_loop_niter to get
113 updated number of iterations. Extract alt_bount, and instantiate it.
115 2015-06-17 Cesar Philippidis <cesar@codesourcery.com>
117 * omp-low.c (extract_omp_for_data): Only use static_chunk
118 scheduling for acc gang loops when a the gang clause contains
119 a non-'*' static argument.
121 2015-06-16 Tom de Vries <tom@codesourcery.com>
123 * varpool.c (make_offloadable_1): Declare decl with ATTRIBUTE_UNUSED.
125 2015-06-16 Tom de Vries <tom@codesourcery.com>
127 * omp-low.c (lower_reduction_clauses): Remove unused variable call.
129 2015-06-16 Chung-Lin Tang <cltang@codesourcery.com>
131 * omp-low.c (struct omp_region): Add inside_kernels_p field.
132 (expand_omp_for_generic): Adjust to generate a 'sequential' loop
133 when GOMP builtin arguments are BUILT_IN_NONE.
134 (expand_omp_for): Use expand_omp_for_generic() to generate a
135 non-parallelized loop for OMP_FORs inside OpenACC kernels regions.
136 (expand_omp): Mark inside_kernels_p field true for regions
137 nested inside OpenACC kernels constructs.
139 2015-06-15 Cesar Philippidis <cesar@codesourcery.com>
141 * omp-low.c (expand_omp_for_static_nochunk): Update entry_bb after
142 calling oacc-boradcast.
143 (expand_omp_for_static_chunk): Likewise.
144 (generate_oacc_broadcast): Insert a barrier after reading from
145 shared memory.
146 (set<tree> live_in): New static set.
147 (populate_loop_live_in): New function.
148 (oacc_populate_live_in_1): New function.
149 (oacc_populate_live_in): New function.
150 (oacc_broadcast): Extract loop data from omp_regions instead of
151 omp_for_data, populate live_in set, broadcast worker variables, and
152 return updated entry_bb.
154 2015-06-15 Julian Brown <julian@codesourcery.com>
156 * omp-low.c (build_oacc_threadbarrier): New function.
157 (oacc_loop_needs_threadbarrier_p): New function.
158 (expand_omp_for_static_nochunk, expand_omp_for_static_chunk):
159 Insert threadbarrier after worker loops.
160 (find_omp_for_region_data): Rename to...
161 (find_omp_for_region_gwv): This. Return mask, rather than modifying
162 REGION structure.
163 (build_omp_regions_1): Move modification of REGION structure to
164 here, after calling above function with new name.
165 (generate_oacc_broadcast): Use new build_oacc_threadbarrier
166 function.
167 (make_gimple_omp_edges): Make edges out of OpenACC worker loop exit
168 block abnormal.
169 * tree-ssa-alias.c (ref_maybe_used_by_call_p_1): Add
170 BUILT_IN_GOACC_THREADBARRIER.
172 2015-06-15 Julian Brown <julian@codesourcery.com>
174 * omp-low.c (make_predication_test): Split false edge out of
175 cond_bb, making latter part abnormal.
176 (predicate_bb): Set EDGE_ABNORMAL on edges before broadcast
177 operations.
179 2015-06-15 Julian Brown <julian@codesourcery.com>
181 * builtins.c (expand_oacc_builtin): Return const1_rtx for
182 ntid/nctaid builtins when the associated patterns are not present.
183 * omp-builtins.def (BUILT_IN_GOACC_GET_THREAD_NUM)
184 (BUILT_IN_GOACC_GET_NUM_THREADS): Remove.
185 * omp-low.c (struct omp_for_data): Remove gang, worker, vector
186 fields.
187 (extract_omp_for_data): Don't initialise deleted gang, worker,
188 vector fields.
189 (expand_oacc_get_num_threads, expand_oacc_get_thread_num): New
190 functions.
191 (lower_reduction_clauses): Use above functions.
192 (expand_omp_for_static_nochunk): Likewise.
193 (expand_omp_for_static_chunk): Likewise.
195 2015-06-11 Cesar Philippidis <cesar@codesourcery.com>
197 * omp-low.c (extract_omp_for_data): Adjust the chunk_size
198 based on the static argument of the gang clause.
200 2015-06-10 Cesar Philippidis <cesar@codesourcery.com>
202 * omp-low.c (enclosing_target_ctx): Un-ifdef. Remove null checking
203 assert.
204 (scan_omp_for): Check for incompatible combination of acc loop clauses.
206 2015-06-09 Tom de Vries <tom@codesourcery.com>
208 backport from trunk:
209 2015-06-09 Tom de Vries <tom@codesourcery.com>
211 * gimple-iterator.h (gimple_seq_nondebug_singleton_p): Don't
212 always return false.
214 2015-06-08 James Norris <jnorris@codesourcery.com>
216 * builtin-types.def (BT_FN_VOID_PTR_INT_UINT): New type.
217 * gimple-pretty-print.c (dump_gimple_omp_target): Handle declare
218 directive.
219 * gimple.h (gf_mask): Add enum.
220 (is_gimple_omp_oacc): Add declare directive.
221 * gimplify.c (omp_notice_variable): Handle device_resident.
222 (gimplify_omp_target_update): Handle declare directive.
223 (gimplify_expr): Handle declare directive.
224 * omp-builtins.def (BUILT_IN_GOACC_STATIC, BUILT_IN_GOACC_DECLARE):
225 New types.
226 * omp-low.c (expand_omp_target): Handle declare directive.
227 (build_omp_regions_1): Likewise.
228 (lower_omp_target): Likewise.
229 (make_gimple_omp_edges): Likewise.
230 * varpool.c (gomp-constants.h): Add inclusion.
231 (make_offloadable_1, make_offloadable): New functions.
232 (get_create): Add calls to make_offloadable.
234 2015-06-08 Tom de Vries <tom@codesourcery.com>
236 backport from trunk:
237 2015-06-08 Tom de Vries <tom@codesourcery.com>
239 PR tree-optimization/66442
240 * gimple-iterator.h (gimple_seq_nondebug_singleton_p): Add function.
241 * tree-parloops.c (try_transform_to_exit_first_loop_alt): Return false
242 if the loop latch is not a singleton. Use
243 gimple_seq_nondebug_singleton_p instead of gimple_seq_singleton_p.
245 2015-06-04 Tom de Vries <tom@codesourcery.com>
247 * omp-low.c (gimple_stmt_ssa_operand_references_var_p): Remove function.
248 (gimple_stmt_omp_data_i_init_p): Rewrite without
249 gimple_stmt_ssa_operand_references_var_p.
251 2015-06-04 Tom de Vries <tom@codesourcery.com>
253 * omp-low.c (release_dangling_ssa_names): Factor out of ...
254 (pass_expand_omp_ssa::execute): ... here.
256 2015-06-04 Tom de Vries <tom@codesourcery.com>
258 * omp-low.c (release_first_vuse_in_edge_dest): Add missing comment.
260 2015-06-04 Tom de Vries <tom@codesourcery.com>
262 * cfgloop.h (struct loop): Add in_oacc_kernels_region field.
263 * omp-low.c (mark_loops_in_oacc_kernels_region): New function.
264 (loop_get_oacc_kernels_region_entry): New function.
265 (expand_omp_target): Call mark_loops_in_oacc_kernels_region.
266 (loop_in_oacc_kernels_region_p): Remove function.
267 * omp-low.h (loop_in_oacc_kernels_region_p): Remove declaration.
268 (loop_get_oacc_kernels_region_entry): Declare.
269 * tree-parloops.c (parallelize_loops): Use in_oacc_kernels_region field and
270 loop_get_oacc_kernels_region_entry.
271 * tree-ssa-loop-ch.c (pass_ch_execute): Use in_oacc_kernels_region field.
273 2015-06-04 Tom de Vries <tom@codesourcery.com>
275 backport from trunk:
276 2015-05-28 H.J. Lu <hongjiu.lu@intel.com>
278 * builtins.c (expand_builtin_acc_on_device): Mark parameters
279 with ATTRIBUTE_UNUSED.
281 2015-06-03 Cesar Philippidis <cesar@codesourcery.com>
283 * builtins.c (expand_builtin_oacc_thread_broadcast): Allow
284 constant value operads.
285 * omp-low.c (expand_omp_for_static_nochunk): Broadcast variables
286 to openacc vector threads when necessary.
287 (expand_omp_for_static_chunk): Likewise.
288 (generate_vector_broadcast): Teach how to handle non-integral typed
289 variables.
290 (populate_loop_use): New function.
291 (oacc_broadcast_1): New function.
292 (oacc_broadcast): New function.
294 2015-06-02 Bernd Schmidt <bernds@codesourcery.com>
296 * omp-low.c (oacc_alloc_broadcast_storage): Remove unused parameter
297 CLAUSES and unnecessary code. All callers changed.
298 (expand_omp_target): Only look for entry-end statement if offloaded.
299 (predicate_bb): Extend code to also jump around atomic regions.
301 2015-06-01 Tom de Vries <tom@codesourcery.com>
303 Revert:
304 2015-05-28 Tom de Vries <tom@codesourcery.com>
306 * passes.def: Replace pass_copy_prop with pass_fre. Surround with
307 pass_tree_loop_done and pass_tree_loop_init.
309 2015-06-01 Bernd Schmidt <bernds@codesourcery.com>
311 * gimple.h (struct gimple_statement_omp_parallel_layout): Add a
312 broadcast_array field.
313 (gimple_omp_target_broadcast_array,
314 gimple_omp_target_set_broadcast_array): New inline functions.
315 * omp-low.c (struct omp_region): Add a broadcast_array field.
316 (find_omp_target_region_data): Initialize it. Change STMT arg to
317 be a gomp_target *. All callers changed.
318 (struct omp_context): Add worker_sync_elt field.
319 (alloc_var_ganglocal): Properly handle a NULL underlying_var.
320 (oacc_alloc_broadcast_storage): Likewise.
321 (required_predication_mask): New function.
322 (requires_vector_predicate): Remove. All callers changed to use
323 required_predication_mask.
324 (generate_vector_broadcast): Return a statement suitable as a
325 block splitting point.
326 (generate_oacc_broadcast, make_predication_test): New static
327 functions.
328 (predicate_bb): New arg MASK. All callers changed. Use
329 generate_oacc_braodcast and use the split point it returns.
330 Handle WSVP regions by jumping across entire loops. Use
331 make_predication_test. Correctly handle GIMPLE_OMP_RETURN.
332 (lower_omp_target): Call oacc_alloc_broadcast_storage. Call
333 gimple_omp_target_set_broadcast_array on the stmt.
335 2015-06-01 Tom de Vries <tom@codesourcery.com>
337 * builtin-attrs.def (DOT_DOT_r_r_r): Remove DEF_ATTR_FOR_STRING.
338 (DOT_DOT_DOT_DOT_r_r_r): Add DEF_ATTR_FOR_STRING.
339 (ATTR_FNSPEC_DOT_DOT_r_r_r_NOTHROW_LIST): Remove DEF_ATTR_TREE_LIST.
340 (ATTR_FNSPEC_DOT_DOT_DOT_DOT_r_r_r_NOTHROW_LIST): Add
341 DEF_ATTR_TREE_LIST.
342 * omp-builtins.def (BUILT_IN_GOACC_DATA_START)
343 (BUILT_IN_GOACC_ENTER_EXIT_DATA, BUILT_IN_GOACC_KERNELS_INTERNAL)
344 (BUILT_IN_GOACC_UPDATE): Fix fnspec.
346 2015-06-01 Tom de Vries <tom@codesourcery.com>
348 * omp-low.c (expand_omp_target): Fix GIMPLE_OMP_ENTRY_END handling for
349 ssa.
351 2015-05-29 Bernd Schmidt <bernds@codesourcery.com>
353 * omp-low.c (struct omp_context): Add worker_var and worker_count
354 fields.
355 (oacc_init_count_vars): New function.
356 (lower_omp_target): Call it.
358 * config/nvptx/nvptx.md (UNSPECV_BARSYNC): New constant.
359 (oacc_threadbarrier): New expander.
360 (threadbarrier_insn): New pattern.
361 * config/nvptx/nvptx.c (nvptx_cannot_copy_insn_p):
362 * omp-builtins.def (BUILT_IN_GOACC_THREADBARRIER): Define.
363 * builtins.c (expand_oacc_threadbarrier): Expand it.
364 (expand_builtin): Handle it.
366 * gimple.def (GIMPLE_OMP_ENTRY_END): New code.
367 * gimple.h (gimple_build_omp_entry_end): Declare.
368 (CASE_GIMPLE_OMP): Add GIMPLE_OMP_ENTRY_END.
369 * gimple.c (gimple_build_omp_entry_end): New function.
370 * gimple-low.c (lower_stmt): Handle GIMPLE_OMP_ENTRY_END.
371 * omp-low.c (expand_omp_target): Remove the entry end statement.
372 (predicate_bb): Don't predicate a block ending with
373 GIMPLE_OMP_ENTRY_END.
374 (lower_omp_target): Append a GIMPLE_OMP_ENTRY_END after the
375 ganglocal data allocation.
376 (make_gimple_omp_edges): Handle GIMPLE_OMP_ENTRY_END.
377 * gimple-pretty-print.c (pp_gimple_stmt_1): Likewise.
379 2015-05-28 Tom de Vries <tom@codesourcery.com>
381 * passes.def: Replace pass_copy_prop with pass_fre. Surround with
382 pass_tree_loop_done and pass_tree_loop_init.
383 * tree-parloops.c (gen_parallel_loop): Bail out of parallelization if
384 try_transform_to_exit_first_loop_alt fails.
386 2015-05-28 Tom de Vries <tom@codesourcery.com>
388 PR tree-optimization/65443
389 * tree-parloops.c (replace_imm_uses, replace_uses_in_bb_by)
390 (replace_uses_in_bbs_by, transform_to_exit_first_loop_alt)
391 (try_transform_to_exit_first_loop_alt): New function.
392 (transform_to_exit_first_loop): Use
393 try_transform_to_exit_first_loop_alt.
395 2015-05-28 Tom de Vries <tom@codesourcery.com>
397 PR tree-optimization/65419
398 * cfgexpand.c (pass_data_expand): Add PROP_gimple_lompifn to
399 properties_required field.
400 * gimplify.c (gimplify_omp_workshare): Use IFN_GOACC_DATA_END_WITH_ARG
401 instead of BUILT_IN_GOACC_DATA_END. Clear PROP_gimple_lompifn in
402 curr_properties.
403 (gimplify_function_tree): Tentatively set PROP_gimple_lompifn in
404 curr_properties.
405 * internal-fn.c (expand_GOACC_DATA_END_WITH_ARG): New dummy function.
406 * internal-fn.def (GOACC_DATA_END_WITH_ARG): New DEF_INTERNAL_FN.
407 * omp-low.c (lower_omp_target): Set argument of GOACC_DATA_END_WITH_ARG.
408 (pass_data_late_lower_omp): New pass_data.
409 (pass_late_lower_omp): New pass.
410 (pass_late_lower_omp::gate, pass_late_lower_omp::execute)
411 (make_pass_late_lower_omp): New function.
412 * passes.def: Add pass_late_lower_omp.
413 * tree-inline.c (expand_call_inline): Handle PROP_gimple_lompifn.
414 * tree-pass.h (PROP_gimple_lompifn): Add define.
416 2015-05-28 Julian Brown <julian@codesourcery.com>
418 PR libgomp/65742
420 * builtins.c (expand_builtin_acc_on_device): Don't use open-coded
421 sequence for !ACCEL_COMPILER.
423 2015-05-27 Tom de Vries <tom@codesourcery.com>
425 * tree-parloops.c (parallelize_loops): Remove checks limiting type of
426 loops allowed.
428 2015-05-21 Bernd Schmidt <bernds@codesourcery.com>
430 * omp-low.c (struct omp_region): Add a gwv_this field.
431 (bb_region_map): New variable.
432 (find_omp_for_region_data, find_omp_target_region_data): New static
433 functions.
434 (build_omp_regions_1): Call them. Build the bb_region_map.
435 (enclosing_target_region, requires_vector_predicate,
436 generate_vector_broadcast, predicate_bb, find_predicatable_bbs,
437 predicate_omp_regions): New static functions.
438 (execute_expand_omp): Allocate and free bb_region_map.
440 2015-05-20 Bernd Schmidt <bernds@codesourcery.com>
442 * config/nvptx/nvptx.c: Include "dumpfile,h".
443 (condition_unidirectional_p): New static function.
444 (nvptx_print_operand): Use it for new 'U' handling.
445 (nvptx_reorg): Compute warp_equal_pseudos.
446 * config/nvptx/nvptx.h (struct machine_function): New field
447 warp_equal_pseudos.
448 * config/nvptx/nvptx.md (br_true, br_false): Add %U modifier.
450 2015-05-19 Bernd Schmidt <bernds@codesourcery.com>
452 * omp-builtins.def (GOACC_thread_broadcast,
453 GOACC_thread_broadcast_ll): New builtins.
454 * optabs.def (oacc_thread_broadcast_optab): New optab.
455 * builtins.c (expand_builtin_oacc_thread_broadcast): New function.
456 (expand_builtin): Use it.
457 * config/nvptx/nvptx.c (nvptx_cannot_copy_insn_p): New function.
458 (TARGET_CANNOT_COPY_INSN_P): Define.
459 * config/nvptx/nvptx.md (UNSPECV_WARP_BCAST): New constant.
460 (oacc_thread_broadcastsi): New pattern.
461 (oacc_thread_broadcastdi): New expander.
463 2015-05-19 Tom de Vries <tom@codesourcery.com>
465 * omp-low.c (enclosing_target_ctx): Comment out.
467 2015-05-18 Tom de Vries <tom@codesourcery.com>
469 PR tree-optimization/65637
470 * omp-low.c (expand_omp_for_static_chunk): Handle case that fin_bb has 2
471 predecessors.
473 2015-05-18 Tom de Vries <tom@codesourcery.com>
475 PR tree-optimization/65637
476 * omp-low.c (find_phi_with_arg_on_edge): New function.
477 (expand_omp_for_static_chunk): Fix inner loop phi.
479 2015-05-18 Tom de Vries <tom@codesourcery.com>
481 PR tree-optimization/65637
482 * omp-low.c (expand_omp_for_static_chunk): Fix gcc_assert for the case
483 that head is NULL.
485 2015-05-15 Cesar Philippidis <cesar@codesourcery.com>
487 * omp-low.c (scan_omp_for): Remove check for nested parallel and
488 kernels regions.
490 2015-05-13 Bernd Schmidt <bernds@codesourcery.com>
492 * config/nvptx/nvptx-opts.h: New file.
493 * config/nvptx/nvptx.c (nvptx_file_start): Print the correct .target.
494 * config/nvptx/nvptx.h: Include "nvptx-opts.h".
495 (ASM_SPEC): Define.
496 (TARGET_SM35): New macro.
497 * config/nvptx/nvptx.md (atomic_fetch_<logic><mode>): Enable with the
498 correct predicate.
499 * config/nvptx/nvptx.opt (ptx_isa, sm_30, sm_35): New enum and its
500 values.
501 (misa=): New option.
503 2015-05-13 Cesar Philippidis <cesar@codesourcery.com>
505 * except.c (finish_eh_generation): Don't finalize exeception
506 handlers for functions containing the "oacc function" attribute.
507 * omp-low.c (create_omp_child_function): Add an "oacc function"
508 attribute to acc regions.
510 * langhooks-def.h (lhd_omp_mappable_type): Add bool argument.
511 * langhooks.c (lhd_omp_mappable_type): Likewise, for the
512 parameter.
513 * langhooks.h (struct lang_hooks_for_types): Add bool oacc
514 argument.
515 * gimplify.c (omp_notice_variable): Use it when calling the
516 omp_mappable_type langhook.
518 2015-05-13 Thomas Schwinge <thomas@codesourcery.com>
519 Bernd Schmidt <bernds@codesourcery.com>
520 Cesar Philippidis <cesar@codesourcery.com>
521 Chung-Lin Tang <cltang@codesourcery.com>
522 James Norris <jnorris@codesourcery.com>
523 Joseph Myers <joseph@codesourcery.com>
524 Julian Brown <julian@codesourcery.com>
525 Tom de Vries <tom@codesourcery.com>
527 * builtin-types.def
528 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR):
529 Remove.
530 (BT_FN_INT_INT_INT_INT)
531 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_SIZE_INT_INT_VAR):
532 New function types.
533 * builtins.c (expand_oacc_builtin, expand_oacc_ganglocal_ptr): New
534 functions.
535 (expand_builtin, is_simple_builtin): Use them.
536 * config/nvptx/mkoffload.c (process): Protect
537 GOMP_offload_register prototype from C++ name mangling.
538 * config/nvptx/nvptx.c (nvptx_file_start): Print declaration of
539 sdata.
540 * config/nvptx/nvptx.md (UNSPEC_NCTAID, UNSPEC_CTAID)
541 (UNSPEC_SHARED_DATA): New constants.
542 (oacc_nctaid_insn, oacc_nctaid, oacc_ctaid_insn, oacc_ctaid)
543 (ganglocal_ptr<mode>, ganglocal_ptr): New patterns.
544 * doc/md.texi (oacc_ntid, oacc_tid): Document.
545 * gimple.h (gimple_statement_omp_parallel_layout): Add
546 ganglocal_size member.
547 (gimple_omp_target_ganglocal_size)
548 (gimple_omp_target_set_ganglocal_size): New functions.
549 * gimplify.c (gimplify_omp_var_data): Add GOVD_USE_DEVICE,
550 GOVD_FORCE_MAP, GOVD_GANGLOCAL.
551 (omp_region_type): Add ORT_HOST_DATA.
552 (omp_region_kind, acc_region_kind): New enum types.
553 (gimplify_omp_ctx): Add region_kind, acc_region_kind members.
554 (new_omp_context, omp_add_variable, omp_notice_variable)
555 (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses_1): Update
556 for OpenACC.
557 (gimplify_scan_omp_clauses): Add region_kind formal parameter.
558 Adjust all users.
559 (gimplify_oacc_host_data_1, gimplify_oacc_host_data): New
560 functions.
561 (gimplify_expr): Update handling of OACC_HOST_DATA.
562 * omp-builtins.def (BUILT_IN_GOACC_KERNELS_INTERNAL)
563 (BUILT_IN_GOACC_KERNELS, BUILT_IN_GOACC_PARALLEL): Change type
564 from BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR
565 to
566 BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_INT_INT_INT_SIZE_INT_INT_VAR.
567 Adjust all users.
568 (BUILT_IN_GOACC_GET_THREAD_NUM, BUILT_IN_GOACC_GET_NUM_THREADS):
569 Change type from from BT_FN_INT to BT_FN_INT_INT_INT_INT. Adjust
570 all users.
571 (BUILT_IN_GOACC_NTID, BUILT_IN_GOACC_TID, BUILT_IN_GOACC_NCTAID)
572 (BUILT_IN_GOACC_CTAID, BUILT_IN_GOACC_GET_GANGLOCAL_PTR)
573 (BUILT_IN_GOACC_DEVICEPTR): New builtins.
574 * omp-low.c (omp_context): Add oacc_reduction_set, ganglocal_init,
575 ganglocal_ptr, ganglocal_size, ganglocal_size_host members.
576 (new_omp_context, delete_omp_context): Initialize/deinitialize
577 these, respectively.
578 (omp_for_data): Add gang, worker, vector members.
579 (extract_omp_for_data): Populate these.
580 (oacc_max_threads, oacc_finalize_reduction_data): Rewrite
581 functions.
582 (is_oacc_parallel, oacc_parallel_max_reduction_array_size)
583 (align_and_expand, alloc_var_ganglocal, install_var_ganglocal)
584 (install_array_var_ganglocal)
585 (oacc_outermost_parallel_kernels_context, oacc_inside_routine)
586 (is_oacc_multithreaded, oacc_needs_global_memory)
587 (is_atomic_compatible_reduction, oacc_serial_reduction)
588 (oacc_process_reduction_data_helper): New functions.
589 (build_outer_var_ref, fixup_remapped_decl, scan_sharing_clauses)
590 (check_omp_nesting_restrictions, lower_rec_input_clauses)
591 (lower_reduction_clauses, oacc_initialize_reduction_data)
592 (oacc_process_reduction_data, lower_omp_target)
593 (lower_omp_regimplify_p): Update for OpenACC.
594 * tree-parloops.c (create_parallel_loop): For OpenACC, switch from
595 vector to gang parallelism.
596 * tree-pretty-print.c (dump_omp_clause): Handle
597 GOMP_MAP_FORCE_TO_GANGLOCAL.
599 2015-05-13 Bernd Schmidt <bernds@codesourcery.com>
601 * config/nvptx/mkoffload.c (enum Kind, struct Token, enum Vis)
602 (struct Stmt): Remove.
603 (read_file, tokenize, write_token, write_tokens, alloc_stmt)
604 (alloc_comment, append_stmt, rev_stmts, write_stmt, write_stmts)
605 (parse_insn, parse_list_nosemi, parse_init, parse_file): Remove
606 functions and macros.
607 (decls, vars, fns): Remove variables.
608 (maybe_unlink): Use save_temps rather than debug to keep files.
609 (tool_cleanup): Unlink ptx_cfile_name and ptx_name.
610 (read_file): Accept a pointer to a length and store into it.
611 (process): Don't try to parse the input file, just write it out as a
612 string, but looking for maps. Also write out the length.
613 (main): Don't use -S to compile ptx code. Add -lgomp. Add
614 COLLECT_MKOFFLOAD_OPTIONS. Scan for -fopenacc and produce an empty
615 image if it is not set. Scan for -save-temps.
616 * gcc.c (mkoffload_options): New static variable.
617 (display_help): Mention -Xoffload
618 (driver_handle_option): Handle it.
619 (add_mkoffload_option): New static function.
620 (set_collect_gcc_options): If offloading, set
621 COLLECT_MKOFFLOAD_OPTIONS.
622 * doc/invoke.texi (-Xoffload): Document.
623 * common.opt (Xoffload): New option.
624 * gcc.c (process_command): Use spec_machine rather than
625 spec_host_machine to build tooldir_prefix2.
627 2015-05-11 Thomas Schwinge <thomas@codesourcery.com>
628 Bernd Schmidt <bernds@codesourcery.com>
629 Cesar Philippidis <cesar@codesourcery.com>
630 Chung-Lin Tang <cltang@codesourcery.com>
631 James Norris <jnorris@codesourcery.com>
632 Joseph Myers <joseph@codesourcery.com>
633 Julian Brown <julian@codesourcery.com>
634 Tom de Vries <tom@codesourcery.com>
636 * gimplify.c (gimplify_scan_omp_clauses)
637 (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_TILE.
638 (gimplify_expr): Don't verify OACC_KERNELS_COMBINED, and
639 OACC_PARALLEL_COMBINED.
640 * omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_BIND,
641 OMP_CLAUSE_NOHOST, OMP_CLAUSE_TILE.
642 (check_omp_nesting_restrictions): Support GIMPLE_OMP_ATOMIC_LOAD,
643 GIMPLE_OMP_ATOMIC_STORE inside OpenACC contexts.
644 * tree-core.h (omp_clause_code): Add OMP_CLAUSE_BIND,
645 OMP_CLAUSE_NOHOST, OMP_CLAUSE_TILE, OMP_CLAUSE_DEVICE_TYPE.
646 * tree.c (omp_clause_num_ops, omp_clause_code_name, walk_tree_1):
647 Update for these.
648 * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_TILE.
649 * tree.h (OACC_KERNELS_COMBINED, OACC_PARALLEL_COMBINED): Remove
650 macros.
651 (OMP_CLAUSE_BIND_NAME, OMP_CLAUSE_TILE_LIST)
652 (OMP_CLAUSE_DEVICE_TYPE_DEVICES, OMP_CLAUSE_DEVICE_TYPE_CLAUSES):
653 Add macros.
655 2015-04-21 Tom de Vries <tom@codesourcery.com>
657 * builtin-attrs.def (DOT_DOT_r_r_r): Add DEF_ATTR_FOR_STRING.
658 (ATTR_FNSPEC_DOT_DOT_r_r_r_NOTHROW_LIST): Add DEF_ATTR_TREE_LIST.
659 * omp-builtins.def (BUILT_IN_GOACC_DATA_START)
660 (BUILT_IN_GOACC_ENTER_EXIT_DATA, BUILT_IN_GOACC_UPDATE): Use
661 DEF_GOACC_BUILTIN_FNSPEC instead of DEF_GOACC_BUILTIN.
663 * passes.def: Add pass_fre after pass_ch_oacc_kernels.
665 * passes.def: Add pass_scev_cprop to pass_oacc_kernels.
666 * tree-ssa-loop.c (pass_scev_cprop::clone): New function.
668 * passes.def: Add pass_parallelize_loops_oacc_kernels in pass group
669 pass_oacc_kernels.
670 * tree-parloops.c (create_parallel_loop, gen_parallel_loop): Add
671 function parameters region_entry and bool oacc_kernels_p. Handle
672 oacc_kernels_p.
673 Call create_parallel_loop with additional args.
674 (parallelize_loops): Add function parameter oacc_kernels_p. Calculate
675 dominance info. Skip loops that are not in a kernels region. Call
676 gen_parallel_loop with additional args.
677 (pass_parallelize_loops::execute): Call parallelize_loops with false
678 argument.
679 (pass_data_parallelize_loops_oacc_kernels): New pass_data.
680 (class pass_parallelize_loops_oacc_kernels): New pass.
681 (pass_parallelize_loops_oacc_kernels::execute)
682 (make_pass_parallelize_loops_oacc_kernels): New function.
683 * tree-pass.h (make_pass_parallelize_loops_oacc_kernels): Declare.
685 * passes.def: Add pass_copy_prop to pass group pass_oacc_kernels.
686 * tree-ssa-copy.c (stmt_may_generate_copy): Handle .omp_data_i init
687 conservatively.
689 * passes.def: Add pass_lim in pass group pass_ch_oacc_kernels.
691 * passes.def: Run pass_tree_loop_init and pass_tree_loop_done in pass
692 group pass_oacc_kernels.
693 * tree-ssa-loop.c (pass_tree_loop_init::clone)
694 (pass_tree_loop_done::clone): New function.
696 * omp-low.c (loop_in_oacc_kernels_region_p): New function.
697 * omp-low.h (loop_in_oacc_kernels_region_p): Declare.
698 * passes.def: Add pass_ch_oacc_kernels to pass group pass_oacc_kernels.
699 * tree-pass.h (make_pass_ch_oacc_kernels): Declare
700 * tree-ssa-loop-ch.c: Include omp-low.h.
701 (pass_ch_execute): Declare.
702 (pass_ch::execute): Factor out ...
703 (pass_ch_execute): ... this new function. If handling oacc kernels,
704 skip loops that are not in oacc kernels region.
705 (pass_ch_oacc_kernels::execute):
706 (pass_data_ch_oacc_kernels): New pass_data.
707 (class pass_ch_oacc_kernels): New pass.
708 (pass_ch_oacc_kernels::execute, make_pass_ch_oacc_kernels): New
709 function.
711 * passes.def: Add pass group pass_oacc_kernels.
712 * tree-pass.h (make_pass_oacc_kernels): Declare.
713 * tree-ssa-loop.c (gate_oacc_kernels): New static function.
714 (pass_data_oacc_kernels): New pass_data.
715 (class pass_oacc_kernels): New pass.
716 (make_pass_oacc_kernels): New function.
718 * omp-low.c: Include gimple-pretty-print.h.
719 (release_first_vuse_in_edge_dest): New function.
720 (expand_omp_target): When not in ssa, don't split off oacc kernels
721 region, clear PROP_gimple_eomp in cfun->curr_properties to force later
722 expanssion, and add GOACC_kernels_internal call.
723 When in ssa, split off oacc kernels and convert GOACC_kernels_internal
724 into GOACC_kernels call. Handle ssa-code.
725 (pass_data_expand_omp): Don't set PROP_gimple_eomp unconditionally in
726 properties_provided field.
727 (pass_expand_omp::execute): Set PROP_gimple_eomp in
728 cfun->curr_properties tentatively.
729 (pass_data_expand_omp_ssa): Add TODO_remove_unused_locals to
730 todo_flags_finish field.
731 (pass_expand_omp_ssa::execute): Release dangling SSA_NAMEs after calling
732 execute_expand_omp.
733 (gimple_stmt_ssa_operand_references_var_p)
734 (gimple_stmt_omp_data_i_init_p): New function.
735 * omp-low.h (gimple_stmt_omp_data_i_init_p): Declare.
736 * passes.def: Add pass_expand_omp_ssa after pass_fre. Add
737 pass_expand_omp_ssa after pass_all_early_optimizations.
738 * tree-ssa-ccp.c: Include omp-low.h.
739 (surely_varying_stmt_p, ccp_visit_stmt): Handle .omp_data_i init
740 conservatively.
741 * tree-ssa-forwprop.c: Include omp-low.h.
742 (pass_forwprop::execute): Handle .omp_data_i init conservatively.
743 * tree-ssa-sccvn.c: Include omp-low.h.
744 (visit_use): Handle .omp_data_i init conservatively.
745 * cgraph.c (cgraph_node::release_body): Don't release offloadable
746 functions.
748 * builtin-attrs.def (DOT_DOT_DOT_r_r_r): Add DEF_ATTR_FOR_STRING.
749 (ATTR_FNSPEC_DOT_DOT_DOT_r_r_r_NOTHROW_LIST): Add
750 DEF_ATTR_TREE_LIST.
751 * omp-builtins.def (BUILT_IN_GOACC_KERNELS_INTERNAL): Add
752 DEF_GOACC_BUILTIN_FNSPEC.
754 * builtins.def (DEF_GOACC_BUILTIN_FNSPEC): Define.
756 2015-03-21 Tom de Vries <tom@codesourcery.com>
758 PR tree-optimization/65460
759 * omp-low.c (expand_omp_target): Set parallelized_function on
760 cgraph_node for child_fn.
762 2015-03-21 Tom de Vries <tom@codesourcery.com>
764 backport from trunk:
765 2015-03-21 Tom de Vries <tom@codesourcery.com>
767 PR tree-optimization/65458
768 * cgraph.c (cgraph_node::dump): Handle parallelized_function field.
769 * cgraph.h (cgraph_node): Add parallelized_function field.
770 * lto-cgraph.c (lto_output_node): Write parallelized_function field.
771 (input_overwrite_node): Read parallelized_function field.
772 * omp-low.c (expand_omp_taskreg, finalize_task_copyfn): Set
773 parallelized_function on cgraph_node for child_fn.
774 * tree-parloops.c: Add include of plugin-api.h, ipa-ref.h and cgraph.h.
775 Remove include of gt-tree-parloops.h.
776 (parallelized_functions): Remove static variable.
777 (parallelized_function_p): Rewrite using parallelized_function field of
778 cgraph_node.
779 (create_loop_fn): Remove adding to parallelized_functions.
780 * Makefile.in (GTFILES): Remove tree-parloops.c
782 2015-03-21 Tom de Vries <tom@codesourcery.com>
784 backport from trunk:
785 2015-03-18 Tom de Vries <tom@codesourcery.com>
787 * tree-parloops.c (parallelize_loops): Make static.
788 * tree-parloops.h (parallelize_loops): Remove extern declaration.
790 2015-01-13 Thomas Schwinge <thomas@codesourcery.com>
792 * tree-core.h: Don't include "gomp-constants.h".
793 (struct tree_omp_clause): Change type of map_kind member from enum
794 gomp_map_kind to unsigned char.
795 * tree.h (OMP_CLAUSE_MAP_KIND): Cast it to enum gomp_map_kind.
796 (OMP_CLAUSE_SET_MAP_KIND): New macro.
797 * gimplify.c (gimplify_adjust_omp_clauses_1)
798 (gimplify_adjust_omp_clauses): Use OMP_CLAUSE_SET_MAP_KIND.
799 * omp-low.c (oacc_initialize_reduction_data): Likewise.
800 * tree-nested.c (convert_nonlocal_reference_stmt)
801 (convert_local_reference_stmt, convert_gimple_call): Likewise.
802 * tree-streamer-in.c (unpack_ts_omp_clause_value_fields):
803 Likewise.
805 * lto-streamer-out.c: Include "gomp-constants.h".
806 * tree-streamer-in.c: Likewise.
807 * tree-streamer-out.c: Likewise.
809 2015-01-12 Thomas Schwinge <thomas@codesourcery.com>
811 * tree-core.h (OMP_CLAUSE_MAP_ALLOC, OMP_CLAUSE_MAP_TO)
812 (OMP_CLAUSE_MAP_FROM, OMP_CLAUSE_MAP_TOFROM)
813 (OMP_CLAUSE_MAP_POINTER, OMP_CLAUSE_MAP_TO_PSET)
814 (OMP_CLAUSE_MAP_FORCE_ALLOC, OMP_CLAUSE_MAP_FORCE_TO)
815 (OMP_CLAUSE_MAP_FORCE_FROM, OMP_CLAUSE_MAP_FORCE_TOFROM)
816 (OMP_CLAUSE_MAP_FORCE_PRESENT, OMP_CLAUSE_MAP_FORCE_DEALLOC)
817 (OMP_CLAUSE_MAP_FORCE_DEVICEPTR, OMP_CLAUSE_MAP_LAST): Remove from
818 enum omp_clause_map_kind. Change all users to
819 include/gomp-constants.h's enum gomp_map_kind.
820 (enum omp_clause_map_kind): Remove.
822 2014-12-18 Thomas Schwinge <thomas@codesourcery.com>
824 * builtins.c (expand_builtin_acc_on_device): Fix logic error.
826 * config/i386/intelmic-offload.h: New file.
827 * config/nvptx/offload.h: Likewise.
828 * config.gcc <*-intelmic-*, *-intelmicemul-*, nvptx-*>: Point to
829 them via tm_file.
831 2014-12-18 Thomas Schwinge <thomas@codesourcery.com>
832 Jakub Jelinek <jakub@redhat.com>
834 * builtins.c (expand_builtin_acc_on_device): Use
835 do_compare_rtx_and_jump instead of emit_cmp_and_jump_insns.
837 * builtins.c (expand_builtin_acc_on_device): Make more RTXy.
839 2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
840 Bernd Schmidt <bernds@codesourcery.com>
842 * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_AUTO, and
843 OMP_CLAUSE_SEQ.
844 * tree.c (omp_clause_num_ops, omp_clause_code_name, walk_tree_1):
845 Update accordingly.
846 * gimplify.c (gimplify_scan_omp_clauses)
847 (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_AUTO,
848 OMP_CLAUSE_GANG, OMP_CLAUSE_SEQ, OMP_CLAUSE_VECTOR, and
849 OMP_CLAUSE_WORKER.
850 * tree-pretty-print.c (dump_omp_clause): Likewise.
851 * omp-low.c (scan_sharing_clauses): Likewise.
852 (struct omp_context): Add gwv_below, and gwv_this members.
853 (enclosing_target_ctx, oacc_loop_or_target_p): New functions.
854 (scan_omp_for, scan_omp_target): Use those.
856 2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
858 * tree-nested.c (convert_nonlocal_reference_stmt)
859 (convert_local_reference_stmt, convert_tramp_reference_stmt)
860 (convert_gimple_call) <GIMPLE_OMP_TARGET>: Use
861 is_gimple_omp_offloaded.
863 * omp-low.c (lower_reduction_clauses): Simplify OpenACC handling.
865 * gimple.h: Rename is_gimple_omp_oacc_specifically to
866 is_gimple_omp_oacc. Update all users.
868 * gimplify.c (enum omp_region_type): Restore ORT_TARGET_DATA and
869 ORT_TARGET. Update all users.
871 * doc/generic.texi (OpenMP): Move OpenACC stuff into...
872 (OpenACC): ... this new subsection.
874 * builtin-types.def: Remove BT_FN_VOID_INT_PTR_INT.
876 * oacc-builtins.def: Merge into omp-builtins.def. Update all users.
878 * builtins.c (expand_builtin_acc_on_device): Use
879 include/gomp-constants.h.
880 * omp-low.c (expand_omp_target, oacc_process_reduction_data)
881 (lower_omp_target): Likewise.
882 * tree-core.h (enum omp_clause_map_kind): Likewise.
884 * omp-low.c (build_omp_regions_1, make_gimple_omp_edges): Simplify
885 multi-line if conditions.
887 2014-12-10 Thomas Schwinge <thomas@codesourcery.com>
889 * omp-low.c (scan_omp_target): Remove taskreg_nesting_level and
890 target_nesting_level assertions.
891 (check_omp_nesting_restrictions): Rework OpenACC constructs
892 handling. Update and extend the relevant test cases.
894 2014-12-10 Thomas Schwinge <thomas@codesourcery.com>
895 Bernd Schmidt <bernds@codesourcery.com>
897 * gimple.def (GIMPLE_OACC_KERNELS, GIMPLE_OACC_PARALLEL): Merge
898 into GIMPLE_OMP_TARGET. Update all users.
900 2014-12-10 Thomas Schwinge <thomas@codesourcery.com>
902 * cgraphbuild.c (pass_build_cgraph_edges::execute): Remove
903 handling of GIMPLE_OACC_PARALLEL.
904 * gimple-pretty-print.c (dump_gimple_omp_target): Dump a bit more
905 data, pretty-printing.
907 * omp-low.c (build_omp_regions_1, make_gimple_omp_edges)
908 <GIMPLE_OMP_TARGET>: Handle
909 GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA the same as
910 GF_OMP_TARGET_KIND_OACC_UPDATE.
912 2014-11-13 Cesar Philippidis <cesar@codesourcery.com>
914 * omp-low.c (oacc_get_reduction_array_id): Fix whitespace.
916 2014-11-13 Cesar Philippidis <cesar@codesourcery.com>
918 * omp-low.c (omp_get_id): Rename to...
919 (oacc_get_reduction_array_id): ... this.
920 (lookup_oacc_reduction): ... this.
921 (lookup_reduction): Rename to...
922 (maybe_lookup_reduction): Rename to...
923 (maybe_lookup_oacc_reduction): ... this.
924 (scan_sharing_clauses): Update calls to renamed fns.
925 (lower_reduction_var_helper): Rename to...
926 (oacc_lower_reduction_var_helper): ... this.
927 (lower_reduction_clauses): Rename to...
928 (oacc_lower_reduction_clauses): ... this.
929 (omp_gimple_assign_with_ops): Rename to...
930 (oacc_gimple_assign_with_ops): ... this.
931 (initialize_reduction_data): Rename to ...
932 (oacc_initialize_reduction_data): ... this.
933 (finalize_reduction_data): Rename to...
934 (oacc_finalize_reduction_data): ... this.
935 (process_reduction_data): Rename to...
936 (oacc_process_reduction_data): ... this.
937 (lower_omp_target): Update calls to renamed fns.
939 2014-11-13 Thomas Schwinge <thomas@codesourcery.com>
941 * gimplify.c (gimplify_omp_workshare) <OACC_DATA, OACC_KERNELS,
942 OACC_PARALLEL>: Don't request ORT_TARGET_MAP_FORCE.
943 (enum gimplify_omp_var_data, enum omp_region_type): Remove
944 GOVD_MAP_FORCE, and ORT_TARGET_MAP_FORCE, respectively. Update
945 all users.
947 * omp-low.c (scan_sharing_clauses) <OMP_CLAUSE_MAP>: Revert
948 earlier change.
950 2014-11-12 Cesar Philippidis <cesar@codesourcery.com>
952 * tree.c (omp_clause_code_name): Add missing comma
953 after "_Cilk_for_count_".
955 2014-11-12 Thomas Schwinge <thomas@codesourcery.com>
957 * gimplify.c (gimplify_omp_for): Eliminate govd_private; always
958 use GOVD_PRIVATE.
960 * omp-low.c (scan_oacc_offload, expand_oacc_offload)
961 (lower_oacc_offload): Merge into scan_omp_target,
962 expand_omp_target, lower_omp_target, respectively. Update all
963 users.
965 2014-11-11 Thomas Schwinge <thomas@codesourcery.com>
967 * omp-low.c (scan_sharing_clauses): Remove bogus assertion.
969 * omp-low.c (delete_omp_context): Dispose of reduction_map.
971 * omp-low.c (maybe_lookup_reduction): Don't require an OpenACC
972 context.
973 (lower_oacc_offload): Simplify use of maybe_lookup_reduction.
975 * omp-low.c (lower_omp_target): Restore two gcc_asserts.
977 2014-11-06 Thomas Schwinge <thomas@codesourcery.com>
979 * gimple.h (is_gimple_omp_oacc_specifically): Return true for
980 GIMPLE_OMP_TARGET's GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
982 * omp-low.c (lower_reduction_clauses): Initialize tid.
984 2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
986 * tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE_HOST and
987 OMP_CLAUSE_OACC_DEVICE. Update all users.
989 * gimplify.c (gimplify_oacc_cache): New function.
990 (gimplify_expr): Use it for OACC_CACHE.
991 (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle
992 OMP_CLAUSE__CACHE_.
994 * tree-core.h (enum omp_clause_code): Move OMP_NO_CLAUSE_CACHE
995 next to, and handle it like a data clause. Rename it to
996 OMP_CLAUSE__CACHE_. Update all users.
998 * invoke.texi: Update for OpenACC.
999 * sourcebuild.texi: Likewise.
1001 * tree.def (OACC_WAIT): Remove. Update all users.
1003 * omp-builtins.def (BUILT_IN_OMP_SET_NUM_THREADS): Remove.
1005 2014-11-03 Cesar Philippidis <cesar@codesourcery.com>
1007 * builtins.def (DEF_GOACC_BUILTIN): Revert erroneous checkin.
1009 2014-11-03 Cesar Philippidis <cesar@codesourcery.com>
1011 * oacc-builtins.def (BUILT_IN_GOACC_GET_NUM_THREADS): New built-in
1012 function.
1013 (BUILT_IN_ACC_GET_DEVICE_TYPE): Likewise.
1014 * omp-low.c (oacc_max_threads): New function.
1015 (lower_reduction_clauses): Use the GOACC thread builtin functions to
1016 determine the thread count. Handle multiple reduction variables.
1017 (expand_omp_for_static_nochunk): Likewise.
1018 (expand_omp_for_static_chunk): Likewise.
1019 (finalize_reduction_data): General cleanups.
1020 (process_reduction_data): Use acc_get_device_type to determine nthreads
1021 at runtime.
1023 2014-11-03 Cesar Philippidis <cesar@codesourcery.com>
1025 * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
1026 * gimple-pretty-print.c (dump_gimple_omp_target): Handle it.
1027 * gimplify.c (gimplify_scan_omp_clauses): Remove switch stmt which
1028 declared OMP_CLAUSE_MAP_FORCE_DEALLOC as unimplemented.
1029 (gimplify_omp_target_update): Handle OACC_ENTER_DATA and
1030 OACC_EXIT_DATA.
1031 (gimplify_expr): Shuffle around OACC_ENTER_DATA, OACC_EXIT_DATA and
1032 OACC_WAIT.
1033 * oacc-builtins.def (BUILD_INT_GOACC_ENTER_EXIT_DATA): New built-in
1034 function.
1035 * omp-low.c (expand_omp_target): Handle
1036 GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA. Don't use quick_push when
1037 there is an unknown number of wait args.
1038 (lower_omp_target): Handle GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
1040 2014-10-24 James Norris <jnorris@codesourcery.com>
1042 * builtin-types.def (BT_FN_VOID_INT_INT_VAR,
1043 BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR,
1044 BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR): Add.
1045 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT): Remove.
1046 * c-family/c-common.c (enum c_builtin_type, c_define_builtins): Remove
1047 unused FUNCTION_TYPE_10. Add new FUNCTION_TYPE_VAR_8 and
1048 FUNCTION_TYPE_VAR_12.
1049 * c-family/c-pragma.c (oacc_pragmas): Add entry for wait directive.
1050 * c/c-parser.c (c_parser_omp_clause_name): Add handling of wait and
1051 async clauses.
1052 (c_parser_oacc_wait_list, c_parser_oacc_clause_async,
1053 c_parser_oacc_clause_wait): New.
1054 (c_parser_oacc_all_clauses): Handle async and wait clauses.
1055 (OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK,
1056 OACC_UPDATE_CLAUSE_MASK): Add async and wait clauses.
1057 (OACC_WAIT_CLAUSE_MASK): New.
1058 (c_parser_oacc_wait): New.
1059 (c_parser_omp_construct): Handle wait directive.
1060 * c/c-typeck.c (c_finish_omp_clauses): Handle async and wait clauses.
1061 * fortran/f95-lang.c (gfc_init_builtin_functions): Remove unused
1062 FUNCTION_TYPE_10. Add new FUNCTION_TYPE_VAR_2, FUNCTION_TYPE_VAR_8,
1063 and FUNCTION_TYPE_VAR_12.
1064 * fortran/types.def (BT_FN_VOID_INT_INT_VAR,
1065 BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR,
1066 BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR): Add.
1067 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT): Remove.
1068 * lto/lto-lang.c (enum lto_builtin_type, lto_define_builtins): Remove
1069 unsigned FUNCTION_TYPE_10. Add new FUNCTION_TYPE_VAR_8 and
1070 FUNCTION_TYPE_VAR_12.
1071 * oacc-builtins.def (BUILT_IN_GOACC_KERNELS BUILT_IN_GOACC_PARALLEL,
1072 BUILT_IN_GOACC_UPDATE, BUILT_IN_GOACC_WAIT): Change type.
1073 * omp-low.c (scan_sharing_clauses): Handle async and wait clauses.
1074 (expand_oacc_offload): Lower kernel directive and async and wait
1075 clauses.
1077 2014-10-23 Thomas Schwinge <thomas@codesourcery.com>
1079 * omp-low.c (finalize_reduction_data): Don't skip the reduction
1080 variable's original value.
1082 2014-10-20 Cesar Philippidis <cesar@codesourcery.com>
1084 * gimplify.c (gimplify_scan_omp_clauses): Remove switch stmt which
1085 declared OMP_CLAUSE_MAP_FORCE_DEALLOC as unimplemented.
1086 (gimplify_expr): Remove OACC_WAIT, since it handled directly by the
1087 front ends.
1089 2014-10-15 James Norris <jnorris@codesourcery.com>
1091 * builtin-types.def (BT_FN_VOID_INT_PTR_INT): New type.
1092 * oacc-builtins.def (BUILT_IN_GOACC_WAIT): New builtin.
1094 2014-10-09 Thomas Schwinge <thomas@codesourcery.com>
1096 * omp-low.c (process_reduction_data): Initialize variable inner.
1098 2014-10-08 Cesar Philippidis <cesar@codesourcery.com>
1100 * omp-low.c (lower_reduction_clauses): Clarify comment.
1101 (process_reduction_data): Scan for nonempty bind statements at
1102 the beginning of parallel blocks.
1104 2014-09-23 Thomas Schwinge <thomas@codesourcery.com>
1106 * builtins.def (DEF_GOACC_BUILTIN_COMPILER): New macro.
1107 * oacc-builtins.def (BUILT_IN_GOACC_UPDATE): New builtin.
1108 * builtins.c (expand_builtin_acc_on_device): New function.
1109 (expand_builtin): Use it to handle BUILT_IN_ACC_ON_DEVICE.
1110 (is_inexpensive_builtin): Handle BUILT_IN_ACC_ON_DEVICE.
1112 2014-09-08 Thomas Schwinge <thomas@codesourcery.com>
1114 * configure.ac (offload_targets): Remove.
1115 * configure: Regenerate.
1117 2014-09-05 Bernd Schmidt <bernds@codesourcery.com>
1119 * configure.ac (enable-accelerator): Remove option.
1120 (enable_accelerator): Don't subst.
1121 * configure: Regenerate.
1123 2014-09-03 Thomas Schwinge <thomas@codesourcery.com>
1125 * omp-low.c (expand_oacc_offload): Add child_fn to offload_funcs.
1127 2014-08-19 Ilya Verbin <ilya.verbin@intel.com>
1129 * Makefile.in (GTFILES): Add omp-low.h.
1130 * cgraphunit.c: Include omp-low.h.
1131 (init_offload_var_table): New function.
1132 (ipa_passes): Run ipa_write_summaries for target only when
1133 offload_funcs or offload_vars is not empty.
1134 * gengtype.c (ifiles): Add omp-low.h.
1135 * lto/lto.c (read_cgraph_and_symbols): Call input_offload_tables.
1136 * lto-cgraph.c: Include omp-low.h.
1137 (output_offload_tables): New function.
1138 (input_offload_tables): Likewise.
1139 * lto-section-in.c (lto_section_name): Add offload_table.
1140 * lto-streamer-out.c (lto_output): Call output_offload_tables.
1141 * lto-streamer.h: Declare output_offload_tables
1142 and input_offload_tables.
1143 (lto_section_type): Add LTO_section_offload_table.
1144 * omp-low.c: Define offload_funcs and offload_vars.
1145 (expand_omp_target): Add child_fn to offload_funcs.
1146 (omp_finish_file): Remove obsolete code.
1147 * omp-low.h: Declare offload_funcs and offload_vars.
1149 2014-08-04 Bernd Schmidt <bernds@codesourcery.com>
1151 * Makefile.in (tool_prefix): Remove.
1152 (unlibsubdir): Change depending on enable_as_accelerator.
1153 (GCC_TARGET_INSTALL_NAME): Revert previous change.
1154 (install_driver): Rework offloading changes.
1155 * configure.ac (tool_prefix): Remove.
1156 (enable-as-accelerator-for): Don't require --enable-as-accelerator.
1157 Update the program_transform_name.
1158 (offload_targets): Add 'g' modifier to sed script.
1159 * configure: Regenerate.
1161 2014-07-28 Cesar Philippidis <cesar@codesourcery.com>
1163 * omp-low.c (get_base_type): New function.
1164 (scan_sharing_clauses): Replace TREE_TYPE with get_base_type.
1165 (lower_reduction_clauses): Use it.
1166 (initialize_reduction_data): Likewise.
1167 (finalize_reduction_data): Likewise. Remove new_var.
1169 2014-07-28 Tom de Vries <tom@codesourcery.com>
1171 * omp-low.c (process_reduction_data): Remove unused variables.
1173 2014-07-24 Thomas Schwinge <thomas@codesourcery.com>
1175 * omp-low.c (extract_omp_for_data): Force OpenACC loop to use a
1176 chunk size of one.
1178 * omp-low.c (expand_omp_for_static_chunk): Merge changes
1179 previously applied to expand_omp_for_static_nochunk.
1181 2014-07-14 Cesar Philippidis <cesar@codesourcery.com>
1183 * omp-low.c (extract_omp_for_data): Likewise.
1184 (expand_omp_for_static_nochunk): Likewise.
1186 2014-07-14 Cesar Philippidis <cesar@codesourcery.com>
1188 * omp-low.c (process_reduction_data): Check for new
1189 binding level and skip for-loops which don't have a
1190 reduction clause.
1192 2014-07-08 Cesar Philippidis <cesar@codesourcery.com>
1193 Thomas Schwinge <thomas@codesourcery.com>
1195 * omp-low.c (omp_get_id): New function.
1196 (lookup_reduction): New function.
1197 (maybe_lookup_reduction): New function.
1198 (build_outer_var_ref): Remove openacc assert.
1199 (new_omp_context): Preserve ctx->reduction_map.
1200 (scan_sharing_clauses): Handle OMP_CLAUSE_REDUCTION.
1201 (scan_oacc_offload): Initialize ctx->reduction_map.
1202 (lower_reduction_clauses): Handle OpenACC reductions.
1203 (omp_gimple_assign_with_ops): New function.
1204 (initialize_reduction_data): New function.
1205 (finalize_reduction_data): New function.
1206 (process_reduction_data): New function.
1207 (lower_oacc_offload): Handle reductions.
1208 * gcc/omp-builtins.def (BUILT_IN_OMP_SET_NUM_THREADS): New.
1210 2014-06-13 Thomas Schwinge <thomas@codesourcery.com>
1211 Ilmir Usmanov <i.usmanov@samsung.com>
1213 * tree-pretty-print.c (dump_omp_clause): Instead of
1214 OMP_CLAUSE_DECL, use more specific accessors for some clauses.
1216 2014-06-13 Thomas Schwinge <thomas@codesourcery.com>
1217 Cesar Philippidis <cesar@codesourcery.com>
1219 * omp-low.c (extract_omp_for_data, scan_sharing_clauses): For
1220 OpenACC directives, the collapse(1) clause is supported.
1222 2014-06-13 Cesar Philippidis <cesar@codesourcery.com>
1224 * omp-low.c (scan_sharing_clauses): Preliminary support for
1225 OMP_CLAUSE_PRIVATE for OpenACC.
1227 * omp-low.c (scan_sharing_clauses): For clauses currently not
1228 supported with OpenACC directives, emit a sorry message instead of
1229 aborting.
1231 2014-06-12 Thomas Schwinge <thomas@codesourcery.com>
1232 James Norris <jnorris@codesourcery.com>
1234 * omp-low.c (scan_sharing_clauses): Allow OMP_CLAUSE_IF.
1235 (expand_oacc_offload, expand_omp_target): Handle it.
1237 2014-06-06 Thomas Schwinge <thomas@codesourcery.com>
1238 James Norris <jnorris@codesourcery.com>
1240 * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_UPDATE, and
1241 extend mask size, GF_OMP_TARGET_KIND_MASK.
1242 (is_gimple_omp_oacc_specifically): Handle
1243 GF_OMP_TARGET_KIND_OACC_UPDATE.
1244 * gimplify.c (gimplify_omp_target_update, gimplify_expr):
1245 Likewise.
1246 * gimple-pretty-print.c (dump_gimple_omp_target): Likewise.
1247 * omp-low.c (scan_omp_target, expand_omp_target)
1248 (build_omp_regions_1, lower_omp_target, lower_omp_1)
1249 (make_gimple_omp_edges): Likewise.
1250 * oacc-builtins.def (BUILT_IN_GOACC_UPDATE): New builtin.
1252 2014-06-05 Thomas Schwinge <thomas@codesourcery.com>
1254 * gimplify.c (gimplify_scan_omp_clauses)
1255 (gimplify_adjust_omp_clauses): Handle
1256 OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
1257 * omp-low.c (scan_sharing_clauses, lower_oacc_offload)
1258 (lower_omp_target): Likewise.
1259 * tree-core.h (enum omp_clause_map_kind)
1260 <OMP_CLAUSE_MAP_FORCE_DEVICEPTR>: Update comment.
1262 * gimplify.c (gimplify_scan_omp_clauses) <case OMP_CLAUSE_MAP>:
1263 Don't block OMP_CLAUSE_MAP_FORCE_PRESENT.
1265 2014-06-04 Thomas Schwinge <thomas@codesourcery.com>
1267 * cgraphunit.c (ipa_passes, compile): Handle flag_openacc next to
1268 flag_openmp.
1269 * ipa-inline-analysis.c (inline_generate_summary): Likewise.
1270 * lto-streamer.c (lto_streamer_init, gate_lto_out): Likewise.
1271 * passes.c (ipa_write_summaries): Likewise.
1273 2014-05-12 Bernd Schmidt <bernds@codesourcery.com>
1275 * lto-wrapper.c (ompbegin): New static variable.
1276 (find_ompbeginend): Renamed from find_ompend. Also look for
1277 crtompbegin.o.
1278 (run_gcc): Also use ompbegin.
1280 2014-05-08 Thomas Schwinge <thomas@codesourcery.com>
1282 * omp-low.c (get_offload_symbol_decl): Mark __OPENMP_TARGET__ as
1283 addressable.
1285 2014-04-04 Bernd Schmidt <bernds@codesourcery.com>
1287 * lto-wrapper.c (replace_special_characters): Remove functions and
1288 all calls to it.
1290 * lto-section-names.h (OFFLOAD_VAR_TABLE_SECTION_NAME,
1291 OFFLOAD_FUNC_TABLE_SECTION_NAME): Define.
1292 * lto-wrapper.c (OFFLOAD_FUNC_TABLE_SECTION_NAME): Don't define.
1293 (ompend): New static variable.
1294 (copy_file, find_ompend): New static functions.
1295 (run_gcc): Call find_ompend if we have offload images. Add its
1296 return value to the output.
1297 * omp-low.c: Include "lto-section-names.h".
1298 (omp_finish_file): Initialize section names from macros defined there.
1300 * omp-low.c (offload_symbol_decl): New static variable.
1301 (get_offload_symbol_decl): New static function.
1302 (expand_oacc_offload, expand_omp_target): Use it.
1304 2014-04-02 Thomas Schwinge <thomas@codesourcery.com>
1306 * omp-low.c (omp_finish_file): Use num_vars instead of num_funcs
1307 when recording offload symbols v_vars.
1309 * omp-low.c (expand_oacc_offload): Handle __OPENMP_TARGET__
1310 symbol.
1312 2014-03-20 Bernd Schmidt <bernds@codesourcery.com>
1314 Mostly by Michael Zolotukhin:
1315 * lto-wrapper.c (OFFLOAD_FUNC_TABLE_SECTION_NAME,
1316 OFFLOAD_TARGET_NAMES_ENV): New defines.
1317 (offload_names): New static variable.
1318 (free_array_of_ptrs, parse_env_var, access_check,
1319 prepare_target_image, replace_special_characters,
1320 compile_images_for_openmp_targets): New static functions.
1321 (run_gcc): Determine whether offload sections are present. If so,
1322 run compile_images_for_openmp_targets and print the image names.
1324 * Makefile.in (COLLECT2_OBJS): Add collect-utils.o.
1325 (LTO_WRAPPER_OBJS): New variable.
1326 (lto-wrapper$(exeext)): Use it.
1327 * collect2.c: Include "collect-utils.h".
1328 (vflag, debug): Remove variables. Rename all uses of vflag to
1329 verbose.
1330 (at_file_supplied): No longer static.
1331 (tool_name): New variable.
1332 (do_wait, fork_execute, maybe_unlink): Don't declare.
1333 (tool_cleanup): Renamed from collect_atexit. New arg from_signal.
1334 (collect_atexit): New static function.
1335 (notice): Remove function.
1336 (handler): Just call tool_cleanup.
1337 (maybe_run_lto_and_relink, main, do_dsymutil): Add new arg to
1338 fork_execute calls.
1339 (collect_wait, do_wait, collect_execute): Remove functions.
1340 (maybe_unlink): No longer static.
1341 * collect2.h (vflag, debug): Don't declare.
1342 (at_file_supplied): Declare.
1343 * collect-utils.c (utils_cleanup): New arg from_signal. All callers
1344 changed.
1345 (collect_execute): Replace with implementation from collect2, plus a
1346 new arg use_atfile. All callers changed.
1347 (collect_wait): Replace with implementation from collect2.
1348 (maybe_unlink_file): Remove function.
1349 (fork_execute): Replace with implementation from collect2, plus a
1350 new arg use_atfile. All callers changed.
1351 (do_wait): Add call to utils_cleanup to the error path.
1352 * collect-utils.h (collect_execute, fork_execute, utils_cleanup,
1353 tool_cleanup): Adjust declarations.
1354 * lto-wrapper.c (tool_cleanup): Add unused bool argument.
1355 * tlink.c: Include "collect-utils.h".
1356 (tlink_execute): New arg use_atfile. All callers changed.
1357 (tlink_init, tlink_execute): Remove declarations.
1359 * Makefile.in (ALL_HOST_BACKEND_OBJS): Add collect-utils.o.
1360 (lto-wrapper$(exeext)): Link with collect-utils.o.
1361 * collect-utils.c: New file.
1362 * collect-utils.h: New file.
1363 * lto-wrapper.c: Include "collect-utils.h".
1364 (args_name): Delete variable.
1365 (tool_name): New variable.
1366 (tool_cleanup): New function.
1367 (maybe_unlink): Renamed from maybe_unlink_file. All callers changed.
1368 (lto_wrapper_cleanup, fatal_signal, fatal, fatal_perror,
1369 collect_execute, collect_wait, fork_execute): Remove functions.
1371 Mostly by Michael Zolotukhin:
1372 * omp-low.c: Include "common/common-target.h".
1373 (expand_omp_target): Pass in address of __OPENMP_TARGET__.
1374 (add_decls_addresses_to_decl_constructor, omp_finish_file): New
1375 functions.
1376 * omp-low.h (omp_finish_file): Declare.
1377 * toplev.c: Include "omp-low.h".
1378 (compile_file): Call omp_finish_file.
1379 * target.def (record_offload_symbol): New hook.
1380 * doc/tm.texi.in (TARGET_RECORD_OFFLOAD_SYMBOL): Add.
1381 * doc/tm.texi: Regenerate.
1382 * configure.ac (ENABLE_OFFLOADING): Define if we have offload_targets.
1383 * configure: Regenerate.
1384 * config.in: Regenerate.
1386 * config/darwin.c: Include "lto-section-names.h".
1387 (LTO_SEGMENT_NAME): Don't define.
1388 * config/i386/winnt.c: Include "lto-section-names.h".
1389 * lto-streamer.c: Include "lto-section-names.h".
1390 * lto-streamer.h (LTO_SECTION_NAME_PREFIX, OMP_SECTION_NAME_PREFIX):
1391 Don't define.
1392 (section_name_prefix): Don't declare.
1393 * lto-wrapper.c: Include "lto-section-names.h".
1394 (LTO_SECTION_NAME_PREFIX): Don't define.
1395 * lto-section-names.h: New file.
1396 * cgraphunit.c: Include "lto-section-names.h".
1398 * configure.ac (real_target_noncanonical, tool_prefix,
1399 accel_dir_suffix, offload_targets): Compute new variables.
1400 (--enable-as-accelerator-for, --enable-as-accelerator,
1401 --enable-offload-targets): New options.
1402 * configure: Regenerate.
1403 * config.in: Regenerate.
1404 * Makefile.in (real_target_noncanonical, tool_prefix,
1405 accel_dir_suffix, enable_as_accelerator): New variables substituted
1406 by configure.
1407 (libsubdir, libexecsubdir, GCC_TARGET_INSTALL_NAME): Tweak for the
1408 possibility of being configured as an offload compiler.
1409 (install_driver): Likewise.
1410 (DRIVER_DEFINES): Pass new defines DEFAULT_REAL_TARGET_MACHINE and
1411 ACCEL_DIR_SUFFIX.
1412 * gcc.c (spec_host_machine, accel_dir_suffix): New variables.
1413 (process_command): Tweak path construction for the possibility
1414 of being configured as an offload compiler.
1415 (main): Likewise. Look up specs in just_machine_suffix only if not
1416 ACCEL_COMPILER. Construct OFFLOAD_TARGET_NAMES environment variable
1417 if we have OFFLOAD_TARGETS.
1419 2014-03-20 Thomas Schwinge <thomas@codesourcery.com>
1421 * gimple.h (enum gf_mask): Add GF_OMP_FOR_KIND_OACC_LOOP.
1422 (is_gimple_omp_oacc_specifically): Handle it.
1423 * gimple-pretty-print.c (dump_gimple_omp_for): Likewise.
1424 * gimple.def (GIMPLE_OMP_FOR): Update for OpenACC loop.
1425 * gimple.c (gimple_build_omp_for): Don't explicitly mention some
1426 clauses.
1427 (gimple_copy) <GIMPLE_OMP_FOR>: Handle GF_OMP_FOR_KIND_OACC_LOOP.
1428 * omp-low.c (extract_omp_for_data, scan_sharing_clauses)
1429 (check_omp_nesting_restrictions, lower_rec_input_clauses)
1430 (lower_lastprivate_clauses, lower_reduction_clauses)
1431 (expand_omp_for_generic, expand_omp_for_static_nochunk)
1432 (expand_omp_for_static_chunk, maybe_add_implicit_barrier_cancel)
1433 (lower_omp_for): Likewise.
1434 * tree-inline.c (remap_gimple_stmt): Likewise.
1435 * tree-nested.c (walk_gimple_omp_for)
1436 (convert_nonlocal_reference_stmt, convert_local_reference_stmt)
1437 (convert_gimple_call): Likewise.
1438 * doc/gimple.texi (GIMPLE_OMP_FOR): Don't explicitly mention some
1439 clauses.
1440 * gimplify.c (gimplify_omp_for, gimplify_expr): Handle OACC_LOOP.
1442 * tree.def (OACC_LOOP): Sort after OMP_DISTRIBUTE.
1443 * tree.h (OMP_LOOP_CHECK): Update accordingly.
1444 * gimplify.c (is_gimple_stmt): Sort OACC_LOOP after
1445 OMP_DISTRIBUTE.
1446 * tree-pretty-print.c (dump_generic_node): Likewise.
1447 * doc/generic.texi (OACC_LOOP): Sort after OACC_HOST_DATA.
1448 (OMP_FOR): Fix and extend for OACC_LOOP.
1450 * gimple.h (enum gf_mask): Add and use GF_OMP_FOR_SIMD.
1451 * omp-low.c: Update accordingly.
1453 * gimple.h (enum gf_mask): Rewrite "<< 0" shift expressions used
1454 for GF_OMP_FOR_KIND_MASK, GF_OMP_FOR_KIND_FOR,
1455 GF_OMP_FOR_KIND_DISTRIBUTE, GF_OMP_FOR_KIND_SIMD,
1456 GF_OMP_FOR_KIND_CILKSIMD, GF_OMP_TARGET_KIND_MASK,
1457 GF_OMP_TARGET_KIND_REGION, GF_OMP_TARGET_KIND_DATA,
1458 GF_OMP_TARGET_KIND_UPDATE, GF_OMP_TARGET_KIND_OACC_DATA.
1460 * omp-low.c (check_omp_nesting_restrictions): Allow nesting of
1461 OpenACC constructs inside of OpenACC data constructs.
1463 2014-03-18 Ilmir Usmanov <i.usmanov@samsung.com>
1465 * tree.def (OACC_LOOP): New tree code.
1466 * tree-pretty-print.c (dump_generic_node): Show it.
1467 * tree.h (OACC_KERNELS_COMBINED, OACC_PARALLEL_COMBINED): New macros.
1468 * doc/generic.texi: Document loop directive.
1469 * gimplify.c (is_gimple_stmt, gimplify_expr): Stub gimplification of
1470 loop directive and combined directives.
1472 2014-03-18 Thomas Schwinge <thomas@codesourcery.com>
1474 * builtin-types.def
1475 (BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT): New type.
1476 * oacc-builtins.def (BUILT_IN_GOACC_KERNELS)
1477 (BUILT_IN_GOACC_PARALLEL): Switch to that one.
1478 * gimplify.c (gimplify_scan_omp_clauses)
1479 (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_NUM_GANGS,
1480 OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_VECTOR_LENGTH.
1481 * omp-low.c (scan_sharing_clauses, expand_oacc_offload): Likewise.
1483 2014-03-13 Thomas Schwinge <thomas@codesourcery.com>
1485 * tree.h (OMP_CLAUSE_VECTOR_EXPR): Check for OMP_CLAUSE_VECTOR
1486 instead of OMP_CLAUSE_VECTOR_LENGTH.
1487 (OMP_WAIT_EXPR): Rename to OMP_CLAUSE_WAIT_EXPR. Change
1488 all users.
1490 2014-03-13 Thomas Schwinge <thomas@codesourcery.com>
1492 * omp-low.c (create_omp_child_function): Use
1493 is_gimple_omp_offloaded when looking for offloaded regions.
1495 2014-03-12 Thomas Schwinge <thomas@codesourcery.com>
1497 * omp-low.c (scan_sharing_clauses): Move offloaded logic into...
1498 * gimple.h (is_gimple_omp_offloaded): ... this new static inline
1499 function.
1501 2014-02-28 Thomas Schwinge <thomas@codesourcery.com>
1503 * gimple.def (GIMPLE_OACC_KERNELS): New code.
1504 * doc/gimple.texi: Document it.
1505 * gimple.h (gimple_has_substatements, CASE_GIMPLE_OMP)
1506 (is_gimple_omp_oacc_specifically): Handle it.
1507 (gimple_statement_oacc_kernels): New struct.
1508 (gimple_build_oacc_kernels): New prototype.
1509 (gimple_oacc_kernels_clauses, gimple_oacc_kernels_clauses_ptr)
1510 (gimple_oacc_kernels_set_clauses, gimple_oacc_kernels_child_fn)
1511 (gimple_oacc_kernels_child_fn_ptr)
1512 (gimple_oacc_kernels_set_child_fn, gimple_oacc_kernels_data_arg)
1513 (gimple_oacc_kernels_data_arg_ptr)
1514 (gimple_oacc_kernels_set_data_arg): New inline functions.
1515 * gimple.c (gimple_build_oacc_kernels): New function.
1516 (gimple_copy): Handle GIMPLE_OACC_KERNELS.
1517 * gimple-low.c (lower_stmt): Likewise.
1518 * gimple-walk.c (walk_gimple_op, walk_gimple_stmt): Likewise.
1519 * gimple-pretty-print.c (pp_gimple_stmt_1): Likewise.
1520 (dump_gimple_oacc_parallel): Rename to dump_gimple_oacc_offload.
1521 Also handle GIMPLE_OACC_KERNELS. Update all callers.
1522 * gimplify.c (gimplify_omp_workshare, gimplify_expr): Handle
1523 OACC_KERNELS.
1524 * oacc-builtins.def (BUILT_IN_GOACC_KERNELS): New builtin.
1525 * omp-low.c (scan_oacc_parallel, expand_oacc_parallel)
1526 (lower_oacc_parallel): Rename to scan_oacc_offload,
1527 expand_oacc_offload, and lower_oacc_offload. Also handle
1528 GIMPLE_OACC_KERNELS. Update all callers.
1529 (scan_sharing_clauses, scan_omp_1_stmt, expand_omp, lower_omp_1)
1530 (diagnose_sb_0, diagnose_sb_1, diagnose_sb_2)
1531 (make_gimple_omp_edges): Handle GIMPLE_OACC_KERNELS.
1532 * tree-inline.c (remap_gimple_stmt, estimate_num_insns): Likewise.
1533 * tree-nested.c (convert_nonlocal_reference_stmt)
1534 (convert_local_reference_stmt, convert_tramp_reference_stmt)
1535 (convert_gimple_call): Likewise.
1537 2014-02-27 Thomas Schwinge <thomas@codesourcery.com>
1539 * gimplify.c (gimplify_oacc_parallel): Merge into
1540 gimplify_omp_workshare. Update all callers.
1542 2014-02-27 Ilmir Usmanov <i.usmanov@samsung.com>
1544 Fix OpenACC ASYNC clause: it cannot contain integer-expression-list.
1546 gcc/
1547 * tree.h (OMP_CLAUSE_ASYNC_EXPR): Reorder.
1548 * tree.c (omp_clause_num_ops, omp_clause_code_name): Reorder clauses.
1549 * tree-core.h (enum omp_clause_code): Likewise.
1550 * omp-low.c (scan_sharing_clauses): Likewise.
1551 * gimplify.c (gimplify_scan_omp_clauses): Likewise.
1553 2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
1555 * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DATA.
1556 (is_gimple_omp_oacc_specifically): Handle it.
1557 * gimple-pretty-print.c (dump_gimple_omp_target): Likewise.
1558 * gimplify.c (gimplify_omp_workshare, gimplify_expr): Likewise.
1559 * omp-low.c (scan_sharing_clauses, scan_omp_target)
1560 (expand_omp_target, lower_omp_target, lower_omp_1): Likewise.
1561 * gimple.def (GIMPLE_OMP_TARGET): Update comment.
1562 * gimple.c (gimple_build_omp_target): Likewise.
1563 (gimple_copy): Catch unimplemented case.
1564 * tree-inline.c (remap_gimple_stmt): Likewise.
1565 * tree-nested.c (convert_nonlocal_reference_stmt)
1566 (convert_local_reference_stmt, convert_gimple_call): Likewise.
1567 * oacc-builtins.def (BUILT_IN_GOACC_DATA_START)
1568 (BUILT_IN_GOACC_DATA_END): New builtins.
1570 * omp-low.c (scan_sharing_clauses): Catch unexpected occurrences
1571 of OMP_CLAUSE_TO, OMP_CLAUSE_FROM, OMP_CLAUSE_MAP.
1573 * gimplify.c (enum omp_region_type): Make ORT_TARGET_OFFLOAD a
1574 flag for ORT_TARGET, in its negation replacing ORT_TARGET_DATA.
1575 Update all users.
1577 * omp-low.c (gimple_code_is_oacc): Move to...
1578 * gimple.h (is_gimple_omp_oacc_specifically): ... here. Update
1579 users, and also use it in more places where currently we've only
1580 been checking for GIMPLE_OACC_PARALLEL.
1582 2014-02-18 Thomas Schwinge <thomas@codesourcery.com>
1584 * omp-low.c (diagnose_sb_0, diagnose_sb_1, diagnose_sb_2): Handle
1585 GIMPLE_OACC_PARALLEL.
1586 * gimplify.c (gimplify_case_label_expr): Update comment.
1588 * omp-low.c (diagnose_sb_0): Generalize detection which kind of
1589 structured block we're in.
1591 2014-02-17 Thomas Schwinge <thomas@codesourcery.com>
1593 * tree-core.h (omp_clause_code) <map>: Mention pcopy, pcopyin,
1594 pcopyout, pcreate OpenACC clauses.
1596 2014-02-14 Ilmir Usmanov <i.usmanov@samsung.com>
1597 Dmitry Bocharnikov <dmitry.b@samsung.com>
1598 Evgeny Gavrin <e.gavrin@samsung.com>
1600 Add OpenACC 1.0 support to GENERIC, except loop directive and subarrays.
1602 gcc/
1603 * gimplify.c (is_gimple_stmt, gimplify_scan_omp_clauses,
1604 implify_adjust_omp_clauses, gimplify_expr): Stub OpenACC directives
1605 and clauses.
1606 * omp-low.c (scan_sharing_clauses): Likewise.
1607 (gimple_code_is_oacc): New helper function.
1608 (scan_sharing_clauses): Use it.
1609 * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_HOST,
1610 OMP_CLAUSE_OACC_DEVICE, OMP_CLAUSE_DEVICE_RESIDENT,
1611 OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE_GANG, OMP_CLAUSE_WAIT,
1612 OMP_NO_CLAUSE_CACHE, OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_ASYNC,
1613 OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_NUM_GANGS,
1614 OMP_CLAUSE_NUM_WORKERS and OMP_CLAUSE_VECTOR_LENGTH clauses.
1615 * tree-pretty-print.c (dump_omp_clause): Print OpenACC clause.
1616 (dump_generic_node): Print OpenACC directives and its clauses.
1617 * tree.c (omp_clause_num_ops, omp_clause_code_name,
1618 walk_tree_1): Add OpenACC clauses.
1619 * tree.def (OACC_KERNELS, OACC_DATA, OACC_HOST_DATA, OACC_DECLARE,
1620 OACC_UPDATE, OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_WAIT,
1621 OACC_CACHE): New tree node.
1622 * tree.h (OACC_KERNELS_BODY, OACC_KERNELS_CLAUSES,
1623 OACC_CLAUSE_NUM_GANGS_EXPR, OMP_CLAUSE_NUM_WORKERS_EXPR,
1624 OMP_CLAUSE_VECTOR_LENGTH_EXPR, OMP_CLAUSE_VECTOR_EXPR,
1625 OMP_CLAUSE_WORKER_EXPR, OMP_CLAUSE_GANG_EXPR, OMP_CLAUSE_ASYNC_EXPR,
1626 OMP_WAIT_EXPR, OACC_DATA_BODY, OACC_DATA_CLAUSES, OACC_DECLARE_CLAUSES,
1627 OACC_UPDATE_CLAUSES, OACC_WAIT_CLAUSES, OACC_CACHE_CLAUSES,
1628 OACC_HOST_DATA_BODY, OACC_HOST_DATA_CLAUSES): New macros.
1629 gcc/doc/
1630 * generic.texi: Document OACC_KERNELS, OACC_DATA, OACC_HOST_DATA,
1631 OACC_DECLARE, OACC_UPDATE, OACC_ENTER_DATA, OACC_EXIT_DATA, OACC_WAIT,
1632 OACC_CACHE.
1634 2014-02-07 Ilya Tocar <ilya.tocar@intel.com>
1636 * ipa-inline-analysis.c (inline_free_summary): Always remove
1637 hooks.
1639 2014-01-28 Thomas Schwinge <thomas@codesourcery.com>
1641 * tree-core.h (omp_clause_code): Update description for
1642 OMP_CLAUSE_MAP.
1644 * tree-core.h (omp_clause_map_kind): Add OMP_CLAUSE_MAP_FORCE,
1645 OMP_CLAUSE_MAP_FORCE_ALLOC, OMP_CLAUSE_MAP_FORCE_TO,
1646 OMP_CLAUSE_MAP_FORCE_FROM, OMP_CLAUSE_MAP_FORCE_TOFROM,
1647 OMP_CLAUSE_MAP_FORCE_PRESENT, OMP_CLAUSE_MAP_FORCE_DEALLOC, and
1648 OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
1649 * tree-pretty-print.c (dump_omp_clause): Handle these.
1650 * gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_FORCE.
1651 (omp_region_type): Add ORT_TARGET_MAP_FORCE.
1652 (omp_add_variable, omp_notice_threadprivate_variable)
1653 (omp_notice_variable, gimplify_scan_omp_clauses)
1654 (gimplify_adjust_omp_clauses_1): Extend accordingly.
1655 (gimplify_oacc_parallel): Add ORT_TARGET_MAP_FORCE to ORT_TARGET
1656 usage.
1657 * omp-low.c (install_var_field, scan_sharing_clauses)
1658 (lower_oacc_parallel, lower_omp_target): Extend accordingly.
1660 * tree-core.h (omp_clause_map_kind): Make the identifiers' bit
1661 patterns more obvious. Add comments.
1662 * omp-low.c (lower_oacc_parallel, lower_omp_target): Test for
1663 omp_clause_map_kind flags set instead of for values.
1665 * gimplify.c (gimplify_call_expr, gimplify_modify_expr)
1666 (omp_firstprivatize_variable, omp_notice_threadprivate_variable)
1667 (omp_notice_variable, gimplify_adjust_omp_clauses)
1668 (gimplify_omp_workshare): Treat ORT_TARGET as a flag, not as a
1669 value.
1671 2013-12-18 Thomas Schwinge <thomas@codesourcery.com>
1673 * omp-low.c (lower_oacc_parallel): Switch kinds array to unsigned
1674 short, and shift alignment description to begin at bit 8.
1676 2013-12-04 Thomas Schwinge <thomas@codesourcery.com>
1678 * gimple.h (is_a_helper): Handle GIMPLE_OACC_PARALLEL.
1680 2013-11-30 Balaji V. Iyer <balaji.v.iyer@intel.com>
1682 * config/i386/i386.c (ix86_simd_clone_compute_vecsize_and_simdlen):
1683 Removed a carriage return from the warning string.
1685 2013-11-21 Jakub Jelinek <jakub@redhat.com>
1687 * config/i386/i386.c (ix86_cilkplus_default_vecsize_mangle,
1688 ix86_cilkplus_vecsize_for_mangle): Remove.
1689 (ix86_simd_clone_compute_vecsize_and_simdlen,
1690 ix86_simd_clone_adjust, ix86_simd_clone_usable): New functions.
1691 (TARGET_CILKPLUS_DEFAULT_VECSIZE_MANGLE,
1692 TARGET_CILKPLUS_VECSIZE_FOR_MANGLE): Remove.
1693 (TARGET_SIMD_CLONE_COMPUTE_VECSIZE_AND_SIMDLEN,
1694 TARGET_SIMD_CLONE_ADJUST, TARGET_SIMD_CLONE_USABLE): Define.
1695 * omp-low.c (class argno_map): Remove.
1696 (argno_map::argno_map): Remove.
1697 (simd_clone_vector_of_formal_parm_types): New function.
1698 (simd_clone_struct_alloc): Use struct cgraph_simd_clone
1699 instead of struct simd_clone and struct cgraph_simd_clone_arg
1700 instead of struct simd_clone_arg.
1701 (simd_clone_struct_copy): Likewise.
1702 (simd_clone_clauses_extract): Likewise. Return struct
1703 cgraph_simd_clone *, don't change anything in node.
1704 Use simd_clone_vector_of_formal_parm_types instead of argno_map.
1705 Fail if ignoring linear step.
1706 (vecsize_mangle): Remove.
1707 (simd_clone_compute_base_data_type): Add clone_info argument.
1708 Use simd_clone_vector_of_formal_parm_types instead of argno_map.
1709 (simd_clone_compute_vecsize_and_simdlen): Remove.
1710 (simd_clone_mangle): Remove new_node argument, add clone_info
1711 argument, return the identifier rather than changing anything in
1712 node. If mangled name doesn't already exist among simd clones,
1713 fail.
1714 (simd_clone_create): Handle !old_node->definition case, don't
1715 ICE if cgraph_function_versioning fails.
1716 (simd_clone_adjust_return_type): Handle simdlens larger than
1717 what can be returned or passed in hw register. Handle
1718 !node->definition.
1719 (simd_clone_adjust_argument_types): Likewise. Use
1720 ipa_get_vector_of_formal_parms or
1721 simd_clone_vector_of_formal_parm_types instead of argno_map.
1722 (simd_clone_init_simd_arrays): Handle simdlens larger than
1723 what can be passed in hw register.
1724 (ipa_simd_modify_function_body): Likewise. Call walk_gimple_op
1725 also on GIMPLE_RETURN.
1726 (simd_clone_adjust): Call targetm.simd_clone.adjust. Create
1727 PHI for the iter, rather than waiting for into ssa. Handle
1728 aligned clause on uniform pointer arguments. Handle linear
1729 arguments properly.
1730 (expand_simd_clones): Don't do anything if noclone attribute
1731 is present. Ignore non-prototype C declarations. Rewritten to
1732 only clone if it is possible and desirable to create the simd clone.
1733 Use new target simd_clone hooks to drive the process, allow
1734 creating more than two clones. Handle !node->definition cloning.
1735 (ipa_omp_simd_clone): Use FOR_EACH_FUNCTION instead of
1736 FOR_EACH_DEFINED_FUNCTION.
1737 (pass_omp_simd_clone::gate): Return true also for flag_openmp_simd.
1738 * ipa-cp.c (determine_versionability): Use simd_clones field instead
1739 of has_simd_clones. Also fail if node->simdclone is non-NULL.
1740 (initialize_node_lattices): Set disable = true for simd clones.
1741 * cgraph.c (get_simd_clone): Remove.
1742 * ipa.c (symtab_remove_unreachable_nodes): If node with simd clones
1743 is kept, keep also the simd clones.
1744 * ipa-prop.h (ipa_get_vector_of_formal_parm_types): New prototype.
1745 * ipa-prop.c (get_vector_of_formal_parm_types): Renamed to ...
1746 (ipa_get_vector_of_formal_parm_types): ... this. No longer static.
1747 (ipa_modify_formal_parameters): Adjust caller.
1748 * target.def (TARGET_CILKPLUS_DEFAULT_VECSIZE_MANGLE,
1749 TARGET_CILKPLUS_VECSIZE_FOR_MANGLE): Remove.
1750 (TARGET_SIMD_CLONE_COMPUTE_VECSIZE_AND_SIMDLEN,
1751 TARGET_SIMD_CLONE_ADJUST, TARGET_SIMD_CLONE_USABLE): New target
1752 hooks.
1753 * tree-vect-data-refs.c (vect_analyze_data_refs): Use simd_clones
1754 field instead of has_simd_clones.
1755 * target.h (struct cgraph_node, struct cgraph_simd_node): Declare.
1756 * doc/tm.texi.in (TARGET_CILKPLUS_DEFAULT_VECSIZE_MANGLE,
1757 TARGET_CILKPLUS_VECSIZE_FOR_MANGLE): Remove.
1758 (TARGET_SIMD_CLONE_COMPUTE_VECSIZE_AND_SIMDLEN,
1759 TARGET_SIMD_CLONE_ADJUST, TARGET_SIMD_CLONE_USABLE): Add.
1760 * doc/tm.texi: Regenerated.
1761 * cgraph.h (enum simd_clone_arg_type): Rename to ...
1762 (enum cgraph_simd_clone_arg_type): ... this. Add
1763 SIMD_CLONE_ARG_TYPE_MASK.
1764 (struct simd_clone_arg): Rename to ...
1765 (struct cgraph_simd_clone_arg): ... this. Add orig_type and
1766 vector_type fields. Adjust arg_type field's type.
1767 (struct simd_clone): Rename to ...
1768 (struct cgraph_simd_clone): ... this. Remove hw_vector_size
1769 field. Add vecsize_int, vecsize_float, prev_clone, next_clone
1770 and origin fields. Adjust args field's type.
1771 (struct cgraph_node): Adjust simdclone field's type. Remove
1772 simdclone_of and has_simd_clones fields. Add simd_clones field.
1773 (get_simd_clone): Remove prototype.
1774 * tree-vect-stmts.c (vectorizable_simd_clone_call): Use simd_clones
1775 field instead of has_simd_clones, improve walking of simd clones,
1776 handle arguments and/or return types larger than corresponding
1777 ISA hw vector size, use targetm.simd_clone.usable hook.
1779 2013-11-21 Ilya Tocar <ilya.tocar@intel.com>
1781 * cgraph.h (symtab_node): Add need_dump.
1782 * cgraphunit.c (ipa_passes): Run ipa_write_summaries for omp.
1783 (compile): Intialize streamer for omp.
1784 * ipa-inline-analysis.c (inline_generate_summary): Add flag_openmp.
1785 * lto-cgraph.c (lto_set_symtab_encoder_in_partition): Respect
1786 need_dump flag.
1787 (select_what_to_dump): New.
1788 * lto-streamer.c (section_name_prefix): New.
1789 (lto_get_section_name): Use section_name_prefix.
1790 (lto_streamer_init): Add flag_openmp.
1791 * lto-streamer.h (OMP_SECTION_NAME_PREFIX): New.
1792 (section_name_prefix): Ditto.
1793 (select_what_to_dump): Ditto.
1794 * lto/lto-partition.c (add_symbol_to_partition_1): Set need_dump.
1795 (lto_promote_cross_file_statics): Dump everyhtinh.
1796 * passes.c (ipa_write_summaries): Add parameter,
1797 call select_what_to_dump.
1798 * tree-pass.h (void ipa_write_summaries): Add parameter.
1800 2013-11-15 Jakub Jelinek <jakub@redhat.com>
1802 * tree-vect-stmts.c (vectorizable_simd_clone_call): Avoid
1803 -Wsign-compare warning.
1805 2013-11-12 Aldy Hernandez <aldyh@redhat.com>
1807 * ipa-prop.h (enum ipa_parm_op): Adjust comment to IPA_PARM_OP_NEW
1808 entry.
1810 2013-11-12 Aldy Hernandez <aldyh@redhat.com>
1812 * ipa-prop.c (ipa_modify_expr): Pass expr by address.
1813 (ipa_get_adjustment_candidate): Add another level of indirection
1814 to expr. Adjust comment.
1815 (ipa_get_adjustment_candidate): Add another level of indirection
1816 to first argument.
1818 2013-11-11 Aldy Hernandez <aldyh@redhat.com>
1820 * ipa-prop.c (get_ssa_base_param): New.
1821 * ipa-prop.h (ipa_modify_expr): Rename from ipa_sra_modify_expr.
1822 Remove ipa_sra_modify_function_body.
1823 (ipa_get_adjustment_candidate): Rename from
1824 sra_ipa_get_adjustment_candidate.
1825 * omp-low.c (ipa_simd_modify_stmt_ops): Rename
1826 sra_ipa_get_adjustment_candidate to ipa_get_adjustment_candidate.
1827 * tree-sra.c (get_ssa_base_param): Remove default_def argument.
1828 (create_access): Remove lass argument to get_ssa_base_param.
1829 (disqualify_base_of_expr): Same.
1830 (sra_ipa_get_adjustment_candidate): Rename to
1831 ipa_get_adjustment_candidate and move to ipa-prop.c.
1832 (sra_ipa_modify_expr): Rename to ipa_modify_expr and move to
1833 ipa-prop.c.
1834 (sra_ipa_modify_assign): Rename sra_ipa_modify_expr to
1835 ipa_modify_expr.
1836 (ipa_sra_modify_function_body): Same.
1838 2013-11-11 Jakub Jelinek <jakub@redhat.com>
1840 * tree-vect-data-refs.c (vect_analyze_data_refs): Check loop->safelen
1841 rather than loop->simdlen.
1842 * tree-vect-stmts.c (vectorizable_simd_clone_call): Likewise.
1844 2013-11-08 Aldy Hernandez <aldyh@redhat.com>
1846 * ipa-prop.c (ipa_modify_formal_parameters): Remove
1847 synth_parm_prefix argument.
1848 Use operator enum instead of bit fields.
1849 Add assert for properly handling vector of references.
1850 (ipa_modify_call_arguments): Use operator enum instead of bit
1851 fields.
1852 (ipa_combine_adjustments): Same.
1853 Assert that IPA_PARM_OP_NEW is not used.
1854 (ipa_dump_param_adjustments): Rename reduction to new_decl.
1855 Use operator enum instead of bit fields.
1856 * ipa-prop.h (enum ipa_parm_op): New.
1857 (struct ipa_parm_adjustment): New field op.
1858 Rename reduction to new_decl.
1859 Rename new_arg_prefix to arg_prefix.
1860 Remove new_param, remove_param, copy_param.
1861 (ipa_modify_formal_parameters): Remove argument.
1862 * omp-low.c (simd_clone_adjust_argument_types): Set arg_prefix.
1863 Use operator enum instead of bit fields.
1864 (simd_clone_adjust_argument_types): Use operator enum instead of
1865 bit fields.
1866 Remove last argument to ipa_modify_formal_parameters call.
1867 (simd_clone_init_simd_arrays): Use operator enum.
1868 (ipa_simd_modify_stmt_ops): Rename reduction to new_decl.
1869 (ipa_simd_modify_function_body): Same.
1870 * tree-sra.c (turn_representatives_into_adjustments): Use operator
1871 enum. Set arg_prefix.
1872 (get_adjustment_for_base): Use operator enum.
1873 (sra_ipa_get_adjustment_candidate): Same.
1874 (sra_ipa_modify_expr): Rename reduction to new_decl.
1875 (sra_ipa_reset_debug_stmts): Use operator enum.
1876 (modify_function): Do not pass prefix argument.
1878 2013-11-07 Jakub Jelinek <jakub@redhat.com>
1880 * tree-vectorizer.h (enum stmt_vec_info_type): Add
1881 call_simd_clone_vec_info_type.
1882 * expr.c (store_constructor): Allow CONSTRUCTOR with VECTOR_TYPE
1883 (same sized) elements even if the type of the CONSTRUCTOR has
1884 vector mode and target is a REG.
1885 * tree-vect-data-refs.c: Include cgraph.h.
1886 (vect_analyze_data_refs): Inline by hand find_data_references_in_loop
1887 and find_data_references_in_bb, if find_data_references_in_stmt
1888 fails, still allow calls to #pragma omp declare simd functions
1889 in #pragma omp simd loops unless they contain data references among
1890 the call arguments or in lhs.
1891 * tree-vect-loop.c (vect_determine_vectorization_factor): If a call
1892 doesn't have lhs, set STMT_VINFO_VECTYPE to vector type corresponding
1893 to any of the argument types and exclude it from adjustments of the
1894 vectorization factor.
1895 * tree-vect-stmts.c: Include tree-ssa-loop.h and
1896 tree-scalar-evolution.h.
1897 (vectorizable_function): Don't handle functions with simd clones here.
1898 (vectorizable_call): Nor here. Return early if call doesn't have lhs.
1899 (struct simd_call_arg_info): New type.
1900 (vectorizable_simd_clone_call): New function.
1901 (vect_analyze_stmt, vect_transform_stmt): Call it.
1903 2013-11-07 Thomas Schwinge <thomas@codesourcery.com>
1905 * gimple.def (GIMPLE_OACC_PARALLEL): New code.
1906 * doc/gimple.texi: Document it.
1907 * gimple.h (gimple_build_oacc_parallel): New declaration.
1908 (gimple_oacc_parallel_clauses, gimple_oacc_parallel_clauses_ptr)
1909 (gimple_oacc_parallel_set_clauses, gimple_oacc_parallel_child_fn)
1910 (gimple_oacc_parallel_child_fn_ptr)
1911 (gimple_oacc_parallel_set_child_fn, gimple_oacc_parallel_data_arg)
1912 (gimple_oacc_parallel_data_arg_ptr)
1913 (gimple_oacc_parallel_set_data_arg): New inline functions.
1914 (CASE_GIMPLE_OMP): Add GIMPLE_OACC_PARALLEL.
1915 * gimple.c (gimple_build_oacc_parallel): New function.
1916 (walk_gimple_op, walk_gimple_stmt, gimple_copy): Handle
1917 GIMPLE_OACC_PARALLEL.
1918 * gimplify.c (is_gimple_stmt): Handle GIMPLE_OACC_PARALLEL.
1919 (gimplify_oacc_parallel): New function.
1920 (gimplify_expr): Handle OACC_PARALLEL.
1921 * cgraphbuild.c (build_cgraph_edges): Handle GIMPLE_OACC_PARALLEL.
1922 * gimple-low.c (lower_stmt): Likewise.
1923 * gimple-pretty-print.c (pp_gimple_stmt_1): Likewise.
1924 (dump_gimple_oacc_parallel): New function.
1925 * oacc-builtins.def (BUILT_IN_GOACC_PARALLEL): New macro.
1926 * omp-low.c (scan_oacc_parallel, expand_oacc_parallel)
1927 (lower_oacc_parallel): New functions.
1928 (use_pointer_for_field, build_outer_var_ref, scan_sharing_clauses)
1929 (create_omp_child_function, check_omp_nesting_restrictions)
1930 (scan_omp_1_stmt, lower_rec_simd_input_clauses)
1931 (lower_lastprivate_clauses, lower_reduction_clauses)
1932 (lower_copyprivate_clauses, lower_send_clauses)
1933 (lower_send_shared_vars, expand_omp)
1934 (maybe_add_implicit_barrier_cancel, create_task_copyfn)
1935 (lower_omp_1, make_gimple_omp_edges): Handle GIMPLE_OACC_PARALLEL,
1936 or catch it.
1937 * tree-inline.c (remap_gimple_stmt): Likewise.
1938 * tree-nested.c (convert_nonlocal_reference_stmt)
1939 (convert_local_reference_stmt, convert_tramp_reference_stmt)
1940 (convert_gimple_call): Likewise.
1942 * tree.def (OACC_PARALLEL): New code.
1943 * doc/generic.texi (OpenMP): Document it.
1944 * tree.h (OMP_BODY, OMP_CLAUSES): Include it.
1945 (OACC_PARALLEL_BODY, OACC_PARALLEL_CLAUSES): New macros.
1946 * tree-pretty-print.c (dump_generic_node): Handle OACC_PARALLEL.
1948 * doc/invoke.texi (-fopenacc): Update.
1950 * gimplify.c (gimplify_body): Consider flag_openacc additionally
1951 to flag_openmp.
1952 * omp-low.c (execute_expand_omp, execute_lower_omp)
1953 (gate_diagnose_omp_blocks): Likewise.
1955 * oacc-builtins.def: New file.
1956 * Makefile.in (BUILTINS_DEF): Add oacc-builtins.def.
1957 * builtins.def (DEF_GOACC_BUILTIN): New macro.
1958 Include "oacc-builtins.def".
1960 * gcc.c (LINK_COMMAND_SPEC, GOMP_SELF_SPECS): For -fopenacc, link
1961 to libgomp and its dependencies.
1962 * config/arc/arc.h (LINK_COMMAND_SPEC): Likewise.
1963 * config/darwin.h (LINK_COMMAND_SPEC_A): Likewise.
1964 * config/i386/mingw32.h (GOMP_SELF_SPECS): Likewise.
1965 * config/ia64/hpux.h (LIB_SPEC): Likewise.
1966 * config/pa/pa-hpux11.h (LIB_SPEC): Likewise.
1967 * config/pa/pa64-hpux.h (LIB_SPEC): Likewise.
1968 * doc/invoke.texi (-fopenacc): Update.
1970 * doc/invoke.texi (-fopenacc): Document it.
1971 * doc/sourcebuild.texi (fopenacc): Document it.
1973 Backport from trunk, r204517:
1975 * config/arc/arc.h (LINK_COMMAND_SPEC): For
1976 -ftree-parallelize-loops=*, link to libgomp and its dependencies.
1977 * config/ia64/hpux.h (LIB_SPEC): Likewise.
1978 * config/pa/pa-hpux11.h (LIB_SPEC): Likewise.
1979 * config/pa/pa64-hpux.h (LIB_SPEC): Likewise.
1980 * gcc.c (GOMP_SELF_SPECS): Update comment about libgomp's
1981 dependencies.
1983 2013-11-07 Aldy Hernandez <aldyh@redhat.com>
1985 * omp-low.c (ipa_simd_modify_function_body): Avoid
1986 regimplification of GIMPLE_RETURNs.
1988 2013-11-06 Aldy Hernandez <aldyh@redhat.com>
1990 * ipa-prop.h (sra_ipa_get_adjustment_candidate): Protoize.
1991 * omp-low.c (struct modify_stmt_info): New.
1992 (ipa_simd_modify_function_body_ops_1): Remove.
1993 (ipa_simd_modify_stmt_ops): New.
1994 (ipa_simd_modify_function_body_ops): Remove.
1995 (ipa_simd_modify_function_body): Set new callback info.
1996 Remove special casing. Handle all operators with walk_gimple_op.
1997 * tree-sra.c (get_ssa_base_param): Add new argument. Use it.
1998 (disqualify_base_of_expr): Pass new argument to
1999 get_ssa_base_param.
2000 (sra_ipa_modify_expr): Abstract candidate search into...
2001 (sra_ipa_get_adjustment_candidate): ...here.
2003 2013-11-06 Aldy Hernandez <aldyh@redhat.com>
2005 * omp-low.c (simd_clone_mangle): Linear step of 1 is mangled as
2006 'l'.
2008 2013-11-04 Aldy Hernandez <aldyh@redhat.com>
2010 * omp-low.c (ipa_simd_modify_function_body): Adjust tree operands
2011 on the LHS of an assign.
2012 (ipa_simd_modify_function_body_ops_1): New.
2013 (ipa_simd_modify_function_body_ops): New.
2015 2013-11-02 Aldy Hernandez <aldyh@redhat.com>
2017 * omp-low.c (ipa_simd_modify_function_body): Handle empty returns.
2019 2013-11-01 Jakub Jelinek <jakub@redhat.com>
2021 * cgraph.h (enum linear_stride_type): Remove.
2022 (enum simd_clone_arg_type): New.
2023 (struct simd_clone_arg): Remove linear_stride, linear_stride_num
2024 and uniform fields. Add arg_type and linear_step.
2025 * omp-low.c (simd_clone_struct_copy): Formatting.
2026 (simd_clone_struct_alloc): Likewise. Use size_t.
2027 (simd_clone_clauses_extract, simd_clone_compute_base_data_type,
2028 simd_clone_adjust_argument_types): Adjust for struct simd_clone_arg
2029 changes.
2030 (simd_clone_mangle): Likewise. Handle negative linear step.
2032 2013-11-01 Aldy Hernandez <aldyh@redhat.com>
2034 * tree-vect-stmts.c: Include cgraph.h.
2035 (vectorizable_call): Allow > 3 arguments when a SIMD clone may be
2036 available.
2037 (vectorizable_function): Use SIMD clone if available.
2038 * ipa-cp.c (determine_versionability): Nodes with SIMD clones are
2039 not versionable.
2040 * ggc.h (ggc_alloc_cleared_simd_clone_stat): New.
2041 * cgraph.h (enum linear_stride_type): New.
2042 (struct simd_clone_arg): New.
2043 (struct simd_clone): New.
2044 (struct cgraph_node): Add simdclone and simdclone_of fields.
2045 (get_simd_clone): Protoize.
2046 * cgraph.c (get_simd_clone): New.
2047 Add `has_simd_clones' field.
2048 * ipa-cp.c (determine_versionability): Disallow functions with
2049 simd clones.
2050 * ipa-prop.h (ipa_sra_modify_function_body): Protoize.
2051 (sra_ipa_modify_expr): Same.
2052 (struct ipa_parm_adjustment): Add new_arg_prefix and new_param
2053 fields. Document their use.
2054 * ipa-prop.c (ipa_modify_formal_parameters): Handle creating brand
2055 new parameters and minor cleanups.
2056 * omp-low.c: Add new pass_omp_simd_clone support code.
2057 (make_pass_omp_simd_clone): New.
2058 (pass_data_omp_simd_clone): Declare.
2059 (class pass_omp_simd_clone): Declare.
2060 (vecsize_mangle): New.
2061 (ipa_omp_simd_clone): New.
2062 (simd_clone_clauses_extract): New.
2063 (simd_clone_compute_base_data_type): New.
2064 (simd_clone_compute_vecsize_and_simdlen): New.
2065 (simd_clone_create): New.
2066 (simd_clone_adjust_return_type): New.
2067 (simd_clone_adjust_return_types): New.
2068 (simd_clone_adjust): New.
2069 (simd_clone_init_simd_arrays): New.
2070 (ipa_simd_modify_function_body): New.
2071 (simd_clone_mangle): New.
2072 (simd_clone_struct_alloc): New.
2073 (simd_clone_struct_copy): New.
2074 (class argno_map): New.
2075 (argno_map::argno_map(tree)): New.
2076 (argno_map::~argno_map): New.
2077 (argno_map::operator []): New.
2078 (argno_map::length): New.
2079 (expand_simd_clones): New.
2080 (create_tmp_simd_array): New.
2081 * tree.h (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE): New.
2082 * tree-core.h (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE): Document.
2083 * tree-pass.h (make_pass_omp_simd_clone): New.
2084 * passes.def (pass_omp_simd_clone): New.
2085 * target.def: Define new hook prefix "TARGET_CILKPLUS_".
2086 (default_vecsize_mangle): New.
2087 (vecsize_for_mangle): New.
2088 * doc/tm.texi.in: Add placeholder for
2089 TARGET_CILKPLUS_DEFAULT_VECSIZE_MANGLE and
2090 TARGET_CILKPLUS_VECSIZE_FOR_MANGLE.
2091 * tree-sra.c (sra_ipa_modify_expr): Remove static modifier.
2092 (ipa_sra_modify_function_body): Same.
2093 * tree.h (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE): Define.
2094 * doc/tm.texi: Regenerate.
2095 * config/i386/i386.c (ix86_cilkplus_default_vecsize_mangle): New.
2096 (ix86_cilkplus_vecsize_for_mangle): New.
2097 (TARGET_CILKPLUS_DEFAULT_VECSIZE_MANGLE): New.
2098 (TARGET_CILKPLUS_VECSIZE_FOR_MANGLE): New.
2100 2013-10-10 Jakub Jelinek <jakub@redhat.com>
2102 * tree.c (omp_remove_redundant_declare_simd_attrs): Adjust
2103 for the clauses being stored into TREE_VALUE of an extra TREE_LIST
2104 pointed out by TREE_VALUE of the attribute.
2106 2013-10-04 Jakub Jelinek <jakub@redhat.com>
2108 * omp-low.c (expand_omp_target): When handling IF clause on
2109 #pragma omp target, split new_bb rather than entry_bb. If
2110 not GF_OMP_TARGET_KIND_REGION, split new_bb right before
2111 the GOMP_TARGET stmt, rather than after labels.
2113 2013-09-26 Jakub Jelinek <jakub@redhat.com>
2115 * omp-low.c (expand_task_call): If there are depend clauses,
2116 pass bit 8 in 7th argument and pass pointer to depend array
2117 as 8th argument.
2118 (lower_depend_clauses): New function.
2119 (lower_omp_taskreg): Handle depend clauses.
2120 * omp-builtins.def (BUILT_IN_GOMP_TASK): Use
2121 BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR
2122 instead of BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT.
2123 * builtin-types.def
2124 (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT): Remove.
2125 (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR): New.
2127 2013-09-25 Jakub Jelinek <jakub@redhat.com>
2129 * omp-low.c (lower_omp_sections, lower_omp_for, lower_omp_taskreg):
2130 Emit ctx->cancel_label before destructors.
2132 * gimple-pretty-print.c (dump_gimple_omp_block,
2133 pp_gimple_stmt_1): Handle GIMPLE_OMP_TASKGROUP.
2134 * tree-nested.c (convert_nonlocal_reference_stmt,
2135 convert_local_reference_stmt, convert_gimple_call): Likewise.
2136 * tree-cfg.c (make_edges): Likewise.
2137 * gimple.h (gimple_build_omp_taskgroup): New prototype.
2138 (gimple_has_substatement): Handle GIMPLE_OMP_TASKGROUP.
2139 (CASE_GIMPLE_OMP): Likewise.
2140 * gimplify.c (is_gimple_stmt, gimplify_expr): Handle OMP_TASKGROUP.
2141 * omp-low.c (check_omp_nesting_restrictions): Warn if #pragma omp
2142 cancel is used in nowait loop or sections construct.
2143 (scan_omp_1_stmt, expand_omp_synch, expand_omp, lower_omp_1): Handle
2144 GIMPLE_OMP_TASKGROUP.
2145 (diagnose_sb_1, diagnose_sb_2): Likewise. Handle GIMPLE_OMP_TARGET
2146 and GIMPLE_OMP_TEAMS.
2147 (lower_omp_taskgroup): New function.
2148 * tree-inline.c (remap_gimple_stmt, estimate_num_insns): Handle
2149 GIMPLE_OMP_TASKGROUP.
2150 * gimple-low.c (lower_stmt): Likewise.
2151 * tree.h (OMP_TASKGROUP_BODY): Define.
2152 * tree.def (OMP_TASKGROUP): New tree.
2153 * tree-pretty-print.c (dump_generic_node): Handle OMP_TASKGROUP.
2154 * gimple.c (gimple_build_omp_taskgroup): New function.
2155 (walk_gimple_stmt, gimple_copy): Handle GIMPLE_OMP_TASKGROUP.
2156 * gimple.def (GIMPLE_OMP_TASKGROUP): New GIMPLE code.
2158 2013-09-19 Jakub Jelinek <jakub@redhat.com>
2160 * omp-builtins.def (BUILT_IN_GOMP_TARGET_DATA,
2161 BUILT_IN_GOMP_TARGET_UPDATE): Use BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR
2162 rather than BT_FN_VOID_INT_SIZE_PTR_PTR_PTR.
2163 * builtin-types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR): Remove.
2164 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR): New.
2165 * omp-low.c (expand_omp_target): Remove fnname argument from
2166 GOMP_target, add openmp_target argument to GOMP_target,
2167 GOMP_target_data and GOMP_target_update calls.
2169 2013-09-18 Jakub Jelinek <jakub@redhat.com>
2171 * tree.h (OMP_CLAUSE_REDUCTION_OMP_ORIG_REF): Define.
2172 * tree-core.h (OMP_CLAUSE_REDUCTION_OMP_ORIG_REF): Document.
2173 * omp-low.c (lower_rec_simd_input_clauses): Don't set DECL_VALUE_EXPR
2174 on new_var if it is not a DECL_P.
2175 (lower_rec_input_clauses): Don't force max_vf = 1
2176 if OMP_CLAUSE_REDUCTION_PLACEHOLDER. Add barrier also if any
2177 OMP_CLAUSE_REDUCTION_OMP_ORIG_REF is seen. For OMP_CLAUSE_PRIVATE
2178 in simd, fix last argument to omp_clause_default_ctor langhook.
2179 Handle OMP_CLAUSE_REDUCTION_PLACEHOLDER in simd loops, if
2180 OMP_CLAUSE_REDUCTION_GIMPLE_INIT is NULL, emit omp_clause_default_ctor
2181 if any and emit omp_clause_dtor if any. Handle C++ references in
2182 OMP_CLAUSE_REDUCTION clauses.
2183 (lower_reduction_clauses): Adjust comment for UDRs. Handle
2184 C++ references in OMP_CLAUSE_REDUCTION clauses.
2185 (lower_omp_taskreg): Emit reduction merges before destructors.
2186 * tree-pretty-print.c (dump_omp_clause): Don't emit any reduction
2187 operator name if OMP_CLAUSE_REDUCTION_CODE is ERROR_MARK.
2188 * gimplify.c (omp_add_variable): Ignore GOVD_LOCAL decls for which
2189 privatize_by_reference returns true.
2191 * tree-core.h (OMP_CLAUSE_DEPEND_KIND, OMP_CLAUSE_MAP_KIND,
2192 OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION, OMP_CLAUSE_PROC_BIND_KIND):
2193 Move definitions...
2194 * tree.h (OMP_CLAUSE_DEPEND_KIND, OMP_CLAUSE_MAP_KIND,
2195 OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION, OMP_CLAUSE_PROC_BIND_KIND):
2196 ... here.
2198 2013-09-16 Jakub Jelinek <jakub@redhat.com>
2200 * omp-low.c (install_var_field): Use (mask & 4) to request double
2201 indirection.
2202 (scan_sharing_clauses): For OMP_CLAUSE_MAP_POINTER arrays pass
2203 7 instead of 3 to install_var_field.
2204 (lower_omp_target): For OMP_CLAUSE_MAP_POINTER arrays add extra
2205 indirection.
2207 2013-09-06 Jakub Jelinek <jakub@redhat.com>
2209 * omp-low.c (scan_sharing_clauses): Handle VLAs in
2210 OMP_CLAUSE_{MAP,TO,FROM}. Set DECL_ALIGN (field) before
2211 calling insert_field_into_struct.
2212 (scan_omp_target): Reverse TYPE_FIELDS, verify that
2213 all field alignments are the same.
2214 (lower_omp_target): Use maybe_lookup_field instead of
2215 lookup_sfield to check if field is present. Handle VLAs.
2216 * tree-pretty-print.c (dump_omp_clause): Only check
2217 OMP_CLAUSE_MAP_KIND on OMP_CLAUSE_MAP clauses.
2218 * gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_TO_ONLY.
2219 (omp_firstprivatize_variable, omp_add_variable,
2220 gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Handle
2221 VLAs in OMP_CLAUSE_{MAP,TO,FROM}.
2223 2013-09-05 Jakub Jelinek <jakub@redhat.com>
2225 * gimplify.c (gimplify_call_expr): Don't call fold_stmt
2226 inside of #pragma omp target construct.
2227 (gimplify_modify_expr): Likewise.
2228 * omp-low.c (target_nesting_level): New variable.
2229 (lower_omp_target): Increase/restore target_nesting_level
2230 around lowering #pragma omp target body. Use TYPE_SIZE_UNIT
2231 instead of TYPE_SIZE if OMP_CLAUSE_SIZE is missing.
2232 Or log2 of needed alignment into high 5 bits of kind.
2233 (lower_omp): Call fold_stmt on all stmts inside of
2234 #pragma omp target construct.
2236 * tree-cfg.c (make_edges): For GIMPLE_OMP_TARGET
2237 with GF_OMP_TARGET_KIND_UPDATE, don't look for
2238 GIMPLE_OMP_RETURN and immediately restore previous
2239 region.
2240 * langhooks.c (lhd_omp_mappable_type): New function.
2241 * omp-low.c (scan_sharing_clauses): Ignore OMP_CLAUSE_SHARED
2242 in GIMPLE_OMP_TEAMS constructs. Handle OMP_CLAUSE_NUM_TEAMS,
2243 OMP_CLAUSE_THREAD_LIMIT, OMP_CLAUSE_DEVICE, OMP_CLAUSE_MAP,
2244 OMP_CLAUSE_TO and OMP_CLAUSE_FROM.
2245 (create_omp_child_function): If current function has
2246 "omp declare target" attribute or if current region
2247 is OMP_TARGET or lexically nested in it, add that
2248 attribute to the omp child function.
2249 (scan_omp_target, scan_omp_teams): New functions.
2250 (check_omp_nesting_restrictions): Fix a typo in TEAMS nesting
2251 check.
2252 (scan_omp_1_stmt): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
2253 (lower_rec_input_clauses): Ignore OMP_CLAUSE_SHARED
2254 in GIMPLE_OMP_TEAMS constructs.
2255 (expand_omp_synch): Handle GIMPLE_OMP_TEAMS.
2256 (expand_omp_target): New function.
2257 (expand_omp): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
2258 (build_omp_regions_1): For GIMPLE_OMP_TARGET with
2259 GF_OMP_TARGET_KIND_UPDATE, don't look for GIMPLE_OMP_RETURN and
2260 immediately restore previous region.
2261 (lower_omp_single): Emit a CLOBBER stmt after GIMPLE_OMP_RETURN.
2262 (lower_omp_taskreg): Likewise.
2263 (lower_omp_target, lower_omp_teams): New functions.
2264 (lower_omp_1): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
2265 * tree.h (enum omp_clause_map_kind): Improve description of
2266 OMP_CLAUSE_MAP_POINTER.
2267 (OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION): Define.
2268 * gimple.def (GIMPLE_OMP_TARGET): Use GSS_OMP_PARALLEL instead
2269 of GSS_OMP_SINGLE.
2270 * langhooks.h (struct lang_hooks_for_types): Add
2271 omp_mappable_type hook.
2272 * langhooks-def.h (lhd_omp_mappable_type): New prototype.
2273 (LANG_HOOKS_OMP_MAPPABLE_TYPE): Define.
2274 (LANG_HOOKS_FOR_TYPES_INITIALIZER): Use it.
2275 * gimple.h (gimple_omp_target_clauses, gimple_omp_target_clauses_ptr,
2276 gimple_omp_target_set_clauses): Use gimple_omp_parallel
2277 instead of gimple_omp_single.
2278 (gimple_omp_target_child_fn, gimple_omp_target_child_fn_ptr,
2279 gimple_omp_target_set_child_fn, gimple_omp_target_data_arg,
2280 gimple_omp_target_data_arg_ptr, gimple_omp_target_set_data_arg): New
2281 inlines.
2282 * omp-builtins.def (BUILT_IN_GOMP_TARGET, BUILT_IN_GOMP_TARGET_DATA,
2283 BUILT_IN_GOMP_TARGET_END_DATA, BUILT_IN_GOMP_TARGET_UPDATE,
2284 BUILT_IN_GOMP_TEAMS): New builtins.
2285 * gimple-pretty-print.c (dump_gimple_omp_target): Print child_fn
2286 for #pragma omp target.
2287 * gimplify.c (omp_notice_variable): Diagnose if implicitly mapped
2288 decl doesn't have mappable type.
2289 (gimplify_scan_omp_clauses): For OMP_CLAUSE_MAP don't clear
2290 notice_outer.
2291 (gimplify_adjust_omp_clauses_1): Add OMP_CLAUSE_MAP even if the same
2292 decl is already referenced in target data construct surrounding it.
2293 (gimplify_adjust_omp_clauses): Likewise. Handle
2294 OMP_CLAUSE_THREAD_LIMIT.
2295 (gimplify_omp_workshare): Fix up gimplification of target or target
2296 data construct body. For target data add GOMP_target_end_data
2297 call in a try/finally cleanup.
2298 * builtin-types.def (BT_FN_VOID_UINT_UINT,
2299 BT_FN_VOID_INT_SIZE_PTR_PTR_PTR,
2300 BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR): New.
2302 2013-07-14 Jakub Jelinek <jakub@redhat.com>
2304 * tree.h (omp_declare_simd_clauses_equal): Remove prototype.
2305 (omp_remove_redundant_declare_simd_attrs): New prototype.
2306 * tree.c (omp_declare_simd_clauses_equal): Make static.
2307 (omp_remove_redundant_declare_simd_attrs): New function.
2309 2013-07-09 Jakub Jelinek <jakub@redhat.com>
2311 * gimple-pretty-print.c (dump_gimple_omp_return): Print
2312 gimple_omp_return_lhs if non-NULL.
2313 * tree-ssa-alias.c (ref_maybe_used_by_call_p_1,
2314 call_may_clobber_ref_p_1): Handle BUILT_IN_GOMP_BARRIER_CANCEL,
2315 BUILT_IN_GOMP_LOOP_END_CANCEL, BUILT_IN_GOMP_SECTIONS_END_CANCEL.
2316 * gimple.h (gimple_omp_return_set_lhs, gimple_omp_return_lhs,
2317 gimple_omp_return_lhs_ptr): New inlines.
2318 * gimple.def (GIMPLE_OMP_RETURN): Use GSS_OMP_ATOMIC_STORE
2319 instead of GSS_BASE.
2320 * gimple.c (walk_gimple_op) <case GIMPLE_OMP_RETURN>: Walk lhs.
2321 * builtin-types.def (BT_FN_BOOL_INT, BT_FN_BOOL_INT_BOOL): New.
2322 * omp-builtins.def (BUILT_IN_GOMP_CANCELLATION_POINT): Use
2323 ATTR_NOTHROW_LEAF_LIST instead of ATTR_NULL. Return type is now
2324 bool.
2325 (BUILT_IN_GOMP_CANCEL): Likewise. Add second argument with bool type.
2326 (BUILT_IN_BARRIER_CANCEL, BUILT_IN_GOMP_LOOP_END_CANCEL,
2327 BUILT_IN_GOMP_SECTIONS_END_CANCEL): New builtins.
2328 * omp-low.c (struct omp_context): Add cancel_label and cancellable
2329 fields.
2330 (extract_omp_for_data): Set have_nowait even for simd implicitly.
2331 (check_omp_nesting_restrictions): Verify nesting restrictions for
2332 #pragma omp cancel and #pragma omp cancellation point.
2333 Set ctx->cancellable for regions that can be cancelled or also
2334 for any task region that contains #pragma omp cancellation point.
2335 (scan_omp_1_stmt): Check nesting restrictions even if ctx == NULL.
2336 (build_omp_barrier): Return gimple instead of tree, add lhs argument,
2337 if non-NULL, build GOMP_barrier_cancel builtin instead and set its
2338 call lhs to lhs.
2339 (lower_rec_input_clauses): Adjust build_omp_barrier caller.
2340 (expand_omp_for_static_nochunk, expand_omp_for_static_chunk,
2341 expand_omp_single): Likewise. If OMP_RETURN has lhs, pass it to
2342 build_omp_barrier.
2343 (expand_omp_for_generic): If OMP_RETURN has lhs, use
2344 GOMP_loop_end_cancel libcall instead of GOMP_loop_end and set its
2345 lhs from OMP_RETURN's lhs.
2346 (expand_omp_sections): If OMP_RETURN has lhs, use
2347 GOMP_sections_end_cancel libcall instead of GOMP_sections_end and set
2348 its lhs from OMP_RETURN's lhs.
2349 (maybe_add_implicit_barrier_cancel): New function.
2350 (lower_omp_sections): If ctx->cancellable, emit cancel_label before
2351 OMP_RETURN. Call maybe_add_implicit_barrier_cancel.
2352 (lower_omp_for): Likewise.
2353 (lower_omp_single): Call maybe_add_implicit_barrier_cancel.
2354 (lower_omp_taskreg): If ctx->cancellable, emit cancel_label before
2355 OMP_RETURN.
2356 (lower_omp_1): If ctx->cancellable, create ctx->cancel_label.
2357 Adjust GOMP_barrier libcalls to GOMP_barrier_cancel plus conditional
2358 branch in cancellable regions, adjust GOMP_cancel and
2359 GOMP_cancellation_point in cancellable regions or remove
2360 GOMP_cancellation_point in non-cancellable regions.
2362 2013-07-03 Jakub Jelinek <jakub@redhat.com>
2364 * omp-low.c (expand_omp_single): Don't force barrier for
2365 copyprivate.
2367 2013-06-28 Jakub Jelinek <jakub@redhat.com>
2368 Aldy Hernandez <aldyh@redhat.com>
2370 * internal-fn.def (GOMP_SIMD_LANE, GOMP_SIMD_VF,
2371 GOMP_SIMD_LAST_LANE): New internal functions.
2372 * omp-low.c (omp_max_vf, lower_rec_simd_input_clauses): New
2373 functions.
2374 (lower_rec_input_clauses): Add fd argument. Enforce max_vf = 1
2375 if any data sharing clauses mention VLAs or for array reductions.
2376 Handle OMP_CLAUSE__LOOPTEMP_ clause. For
2377 OMP_CLAUSE_{{FIRST,LAST,}PRIVATE,LINEAR,REDUCTION} on SIMD
2378 constructs use "omp simd array" temporaries. For OMP_CLAUSE_LINEAR
2379 adjust initial value in combined constructs. Don't emit any
2380 barriers for #pragma omp distribute. If max_vf is lower than
2381 current safelen, prepend an OMP_CLAUSE_SAFELEN clause.
2382 (lower_lastprivate_clauses): Handle "omp simd array" temporaries.
2383 (lower_reduction_clauses): Exit early for #pragma omp simd.
2384 (expand_omp_simd): Set loop->simduid from OMP_CLAUSE__SIMDUID_
2385 and cfun->has_simduid_loops if set.
2386 If OMP_CLAUSE_SAFELEN (1) is present, don't set loop->safelen
2387 nor loop->force_vect.
2388 (lower_omp_sections, lower_omp_single, lower_omp_taskreg): Adjust
2389 lower_rec_input_clauses callers.
2390 (lower_omp_for_lastprivate): Unshare vinit.
2391 (lower_omp_for): Add OMP_CLAUSE__LOOPTEMP_ clauses before calling
2392 lower_rec_input_clauses. Adjust lower_rec_input_clauses caller.
2393 Always call lower_omp_for_lastprivate at the same place, even for
2394 #pragma omp simd.
2395 * tree.h (enum clause_code): Add OMP_CLAUSE__SIMDUID_.
2396 (OMP_CLAUSE__SIMDUID__DECL): Define.
2397 * tree-vectorizer.c: Include hash-table.h and tree-ssa-propagate.h.
2398 (simduid_to_vf, decl_to_simduid): New classes.
2399 (simduid_to_vf::hash, simduid_to_vf::equal, decl_to_simduid::hash,
2400 decl_to_simduid::equal): New methods.
2401 (note_simd_array_uses_struct): New struct.
2402 (adjust_simduid_builtins, note_simd_array_uses_cb,
2403 note_simd_array_uses): New functions.
2404 (vectorize_loops): Adjust "omp simd array" temporary array sizes
2405 and fold GOMP_SIMD_{LANE,VF,LAST_LANE} builtins.
2406 * tree-vectorizer.h (struct _stmt_vec_info): Add simd_lane_access_p
2407 field.
2408 (STMT_VINFO_SIMD_LANE_ACCESS_P): Define.
2409 * tree-data-ref.c (get_references_in_stmt): Allow GOMP_SIMD_LANE
2410 builtins in their own loops.
2411 * tree-inline.c (copy_cfg_body): Propagate has_force_vect_loops
2412 and has_simduid_loops.
2413 * function.h (struct function): Add has_simduid_loops field.
2414 * tree-ssa-ccp.c (likely_value): For GOMP_SIMD_{LANE,LAST_LANE,VF}
2415 builtins ignore the undefined magic argument.
2416 * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE__SIMDUID_
2417 clause.
2418 * cfgloop.h (struct loop): Add simduid field.
2419 * Makefile.in (tree-vectorizer.o): Depend on $(HASH_TABLE_H)
2420 and tree-ssa-propagate.h.
2421 * tree-vect-data-refs.c (vect_analyze_data_refs): Check for SIMD
2422 lane access.
2423 * gimplify.c (omp_add_variable): Handle combination of aligned
2424 clause and some data sharing clause for the same decl.
2425 (gimplify_omp_for): For collapse (2) and above simd loops
2426 predetermine loop iteration vars as lastprivate instead of
2427 linear.
2428 * tree.c (omp_clause_num_ops, omp_clause_code_name): Add
2429 entries for OMP_CLAUSE__SIMDUID_.
2430 (walk_tree_1): Handle OMP_CLAUSE__SIMDUID_.
2431 * tree-vect-loop.c (vectorizable_live_operation): Handle live
2432 GOMP_SIMD_LANE result.
2433 * tree-vect-stmts.c (vectorizable_call): Vectorize GOMP_SIMD_LANE
2434 builtin.
2435 (vectorizable_store, vectorizable_load): Handle
2436 STMT_VINFO_SIMD_LANE_ACCESS_P.
2437 * internal-fn.c (expand_GOMP_SIMD_LANE, expand_GOMP_SIMD_VF,
2438 expand_GOMP_SIMD_LAST_LANE): New functions.
2440 2013-06-21 Jakub Jelinek <jakub@redhat.com>
2442 * gimple.h (enum gf_mask): Adjust GF_OMP_FOR_COMBINED
2443 value representation, add GF_OMP_FOR_COMBINED_INTO.
2444 (gimple_omp_for_combined_into_p,
2445 gimple_omp_for_set_combined_into_p): New inlines.
2446 * gimplify.c (enum omp_region_type): Remove outdated
2447 ORT_SIMD comment.
2448 (struct gimplify_omp_ctx): Add combined_loop field.
2449 (gimplify_omp_for): Call gimple_omp_for_set_combined_into_p
2450 for inner for/simd constructs combined with an outer
2451 loop construct (for or distribute).
2452 * tree.c (omp_clause_num_ops): Add OMP_CLAUSE__LOOPTEMP_
2453 entry.
2454 (omp_clause_code_name): Likewise.
2455 (walk_tree_1): Handle OMP_CLAUSE__LOOPTEMP_.
2456 * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE__LOOPTEMP_.
2457 * tree.h (enum omp_clause_code): Add OMP_CLAUSE__LOOPTEMP_.
2458 (OMP_CLAUSE_DECL): Allow also on OMP_CLAUSE__LOOPTEMP_.
2459 * omp-low.c (extract_omp_for_data): Rename non_ws to simd. Don't set
2460 fd->chunk_size for non-chunk OMP_CLAUSE_SCHEDULE_STATIC, unless
2461 fd->have_ordered. For OMP_CLAUSE_SCHEDULE_STATIC non-ordered loops
2462 compute fd->iter_type the same as for simd.
2463 (get_ws_args_for): Add par_stmt argument, if
2464 gimple_omp_for_combined_into_p, use first two _looptemp_ clauses
2465 temporaries instead of fd->loop.n{1,2}.
2466 (determine_parallel_type): Adjust caller.
2467 (scan_sharing_clauses): Handle OMP_CLAUSE__LOOPTEMP_.
2468 (find_combined_for): New function.
2469 (scan_omp_parallel): If gimple_omp_parallel_combined_p and
2470 it is combined with gimple_omp_for_combined_into_p OMP_FOR,
2471 add OMP_CLAUSE__LOOPTEMP_ clauses to the parallel.
2472 (check_omp_nesting_restrictions): Don't insist that the only construct
2473 nested in OMP_DISTRIBUTE must be OMP_PARALLEL.
2474 (lower_rec_input_clauses, lower_send_clauses): Handle
2475 OMP_CLAUSE__LOOPTEMP_.
2476 (expand_omp_for_init_counts, expand_omp_for_init_vars,
2477 extract_omp_for_update_vars): New functions.
2478 (expand_omp_for_generic): Add inner_stmt argument. Use
2479 expand_omp_for_{init,update}* helper functions. Handle combined loop
2480 constructs.
2481 (expand_omp_for_static_nochunk, expand_omp_for_static_chunk):
2482 Likewise. Handle fd->collapse > 1 and broken_loop cases.
2483 (expand_omp_simd): Use expand_omp_for_init* helper functions. Handle
2484 combined loop constructs.
2485 (expand_omp_for): Add inner_stmt argument. Pass it through to
2486 expand_omp_for_{generic,static_{,no}chunk}. Use
2487 expand_omp_for_static* even for fd->collapse > 1 and/or broken_loop
2488 cases, just not when fd->have_ordered.
2489 (expand_omp): Adjust expand_omp_for caller.
2490 (lower_omp_for): If gimple_omp_parallel_combined_p, add
2491 OMP_CLAUSE__LOOPTEMP_ clauses to the GIMPLE_FOR stmt.
2493 2013-06-14 Jakub Jelinek <jakub@redhat.com>
2495 * gimple-pretty-print.c (dump_gimple_omp_for): Don't handle
2496 GF_OMP_FOR_KIND_FOR_SIMD.
2497 * gimple.h (GF_OMP_FOR_KIND_FOR_SIMD): Remove.
2498 (GF_OMP_FOR_COMBINED): New.
2499 (gimple_omp_for_combined_p, gimple_omp_for_set_combined_p): New
2500 inline functions.
2501 * gimplify.c (is_gimple_stmt): Don't handle OMP_FOR_SIMD.
2502 (find_combined_omp_for): New function.
2503 (gimplify_omp_for): Handle combined OMP_DISTRIBUTE and OMP_FOR
2504 loops.
2505 * Makefile.in (c-family/c-omp.o): Depend on $(C_PRAGMA_H).
2506 * omp-low.c (build_outer_var_ref): Fix up simd handling.
2507 (check_omp_nesting_restrictions): Don't handle
2508 GF_OMP_FOR_KIND_FOR_SIMD.
2509 * tree.def (OMP_FOR_SIMD): Remove.
2510 * tree-pretty-print.c (dump_generic_node): Don't handle OMP_FOR_SIMD.
2511 Handle NULL OMP_FOR_INIT.
2513 2013-06-12 Jakub Jelinek <jakub@redhat.com>
2515 * gimplify.c (gimplify_scan_omp_clauses): Handle
2516 OMP_CLAUSE_THREAD_LIMIT.
2517 * tree-pretty-print.c (dump_omp_clause): Likewise.
2518 * tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries for
2519 OMP_CLAUSE_THREAD_LIMIT.
2520 * tree.h (enum omp_clause_code): Add OMP_CLAUSE_THREAD_LIMIT.
2521 (OMP_CLAUSE_THREAD_LIMIT_EXPR): Define.
2523 2013-06-04 Jakub Jelinek <jakub@redhat.com>
2525 * gimplify.c (gimplify_scan_omp_clauses): Handle array
2526 sections on OMP_CLAUSE_{MAP,TO,FROM} clauses, handle
2527 OMP_CLAUSE_DEPEND clause.
2528 (gimplify_adjust_omp_clauses): Handle array sections on
2529 OMP_CLAUSE_MAP, handle OMP_CLAUSE_DEPEND clause.
2530 * tree.c (omp_clause_num_ops): OMP_CLAUSE_{MAP,TO,FROM}
2531 now have 2 arguments, move OMP_CLAUSE_UNIFORM before these
2532 3.
2533 (omp_clause_code_name): Adjust for OMP_CLAUSE_UNIFORM movement.
2534 (walk_tree_1): Adjust to handle 2 arguments of
2535 OMP_CLAUSE_{MAP,TO,FROM}.
2536 * tree-pretty-print.c (dump_omp_clause): For OMP_CLAUSE_{MAP,TO,FROM}
2537 print OMP_CLAUSE_SIZE, and for OMP_CLAUSE_MAP handle
2538 OMP_CLAUSE_MAP_POINTER.
2539 * tree.h (enum omp_clause_code): Move OMP_CLAUSE_UNIFORM before
2540 OMP_CLAUSE_{MAP,TO,FROM}.
2541 (OMP_CLAUSE_SIZE): Define.
2542 (enum omp_clause_map_kind): Add OMP_CLAUSE_MAP_POINTER.
2543 * omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_DEPEND.
2545 2013-05-29 Jakub Jelinek <jakub@redhat.com>
2547 * omp-builtins.def (BUILT_IN_OMP_GET_TEAM_NUM,
2548 BUILT_IN_OMP_GET_NUM_TEAMS): New built-ins.
2549 * omp-low.c (extract_omp_for_data, expand_omp_for_static_nochunk,
2550 expand_omp_for_static_chunk): Handle #pragma omp distribute.
2551 (expand_omp_for): Add assertion for non-finished distribute collapse
2552 > 1 support.
2553 (check_omp_nesting_restrictions): Allow orphaned distribute construct.
2555 * omp-low.c (check_omp_nesting_restrictions): Add some
2556 accelerator related nesting restrictions.
2557 (scan_omp_1_stmt): Call check_omp_nesting_restrictions
2558 even for GOMP_taskgroup_{start,end}.
2559 * gimplify.c (omp_notice_threadprivate_variable): Fix a typo.
2560 (gimplify_body): For functions with "omp declare target" attribute
2561 add ORT_TARGET region around the body.
2563 2013-05-27 Jakub Jelinek <jakub@redhat.com>
2565 * tree.def (OMP_TEAMS, OMP_TARGET_DATA, OMP_TARGET,
2566 OMP_TARGET_UPDATE): New tree codes.
2567 * tree-cfg.c (make_edges): Handle GIMPLE_OMP_TARGET
2568 and GIMPLE_OMP_TEAMS.
2569 * omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_DIST_SCHEDULE.
2570 * gimple-low.c (lower_stmt): Handle GIMPLE_OMP_TARGET
2571 and GIMPLE_OMP_TEAMS.
2572 * tree.h (OMP_TEAMS_BODY, OMP_TEAMS_CLAUSES, OMP_TARGET_DATA_BODY,
2573 OMP_TARGET_DATA_CLAUSES, OMP_TARGET_BODY, OMP_TARGET_CLAUSES,
2574 OMP_TARGET_UPDATE_CLAUSES): Define.
2575 * tree-nested.c (convert_nonlocal_reference_stmt,
2576 convert_local_reference_stmt, convert_gimple_call): Handle
2577 GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
2578 * tree-inline.c (estimate_num_insns): Likewise.
2579 (remap_gimple_stmt): Likewise. Adjust gimple_build_omp_for
2580 caller.
2581 * gimple.def: Adjust comments describing OMP_CLAUSEs.
2582 (GIMPLE_OMP_TARGET, GIMPLE_OMP_TEAMS): New GIMPLE stmts.
2583 * tree-parloops.c (create_parallel_loop): Adjust gimple_build_omp_for
2584 caller.
2585 * tree-pretty-print.c (dump_generic_node): Handle OMP_TEAMS,
2586 OMP_TARGET, OMP_TARGET_DATA and OMP_TARGET_UPDATE.
2587 * gimple.h (GF_OMP_TARGET_KIND_MASK, GF_OMP_TARGET_KIND_REGION,
2588 GF_OMP_TARGET_KIND_DATA, GF_OMP_TARGET_KIND_UPDATE): New.
2589 (gimple_build_omp_for): Add kind argument to prototype.
2590 (gimple_build_omp_target, gimple_build_omp_teams): New prototypes.
2591 (gimple_has_substatements): Handle GIMPLE_OMP_TARGET and
2592 GIMPLE_OMP_TEAMS.
2593 (gimple_omp_subcode): Change GIMPLE_OMP_SINGLE to GIMPLE_OMP_TEAMS.
2594 (gimple_omp_target_clauses, gimple_omp_target_clauses_ptr,
2595 gimple_omp_target_set_clauses, gimple_omp_target_kind,
2596 gimple_omp_target_set_kind, gimple_omp_teams_clauses,
2597 gimple_omp_teams_clauses_ptr, gimple_omp_teams_set_clauses): New
2598 inline functions.
2599 (gimple_return_set_retval): Handle GIMPLE_OMP_TARGET and
2600 GIMPLE_OMP_TEAMS.
2601 * gimple.c (gimple_build_omp_for): Add kind argument, call
2602 gimple_omp_for_set_kind.
2603 (gimple_build_omp_target, gimple_build_omp_teams): New functions.
2604 (walk_gimple_op, walk_gimple_stmt, gimple_copy): Handle
2605 GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
2606 * gimple-pretty-print.c (dump_gimple_omp_target,
2607 dump_gimple_omp_teams): New functions.
2608 (pp_gimple_stmt_1): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
2609 * gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP.
2610 (enum omp_region_type): Add ORT_TEAMS, ORT_TARGET and ORT_TARGET_DATA.
2611 (omp_add_variable): Add temporary assertions.
2612 (omp_notice_threadprivate_variable): Complain if threadprivate vars
2613 appear in target region.
2614 (omp_notice_variable): ORT_TARGET, ORT_TARGET_DATA and ORT_TEAMS
2615 handling.
2616 (omp_check_private): Ignore ORT_TARGET and ORT_TARGET_DATA regions.
2617 (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_MAP, OMP_CLAUSE_TO,
2618 OMP_CLAUSE_FROM, OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_DIST_SCHEDULE
2619 and OMP_CLAUSE_DEVICE.
2620 (gimplify_adjust_omp_clauses): Likewise.
2621 (gimplify_adjust_omp_clauses_1): Handle GOVD_MAP. Fix up
2622 check for privatization by also testing for GOVD_LINEAR.
2623 (gimplify_omp_for): Adjust gimple_build_omp_for caller.
2624 Clear *expr_p.
2625 (gimplify_omp_workshare): Handle also OMP_TARGET, OMP_TARGET_DATA
2626 and OMP_TEAMS. Clear *expr_p.
2627 (gimplify_omp_target_update): New function.
2628 (gimplify_expr): Handle OMP_TARGET, OMP_TARGET_DATA, OMP_TARGET_UPDATE
2629 and OMP_TEAMS.
2631 2013-05-20 Jakub Jelinek <jakub@redhat.com>
2633 * omp-low.c (expand_omp_simd): For collapse > 1 loops,
2634 if some loop condition might be not true initially, add runtime
2635 test and skip the whole loop.
2637 2013-05-14 Jakub Jelinek <jakub@redhat.com>
2639 * cfgloop.h (struct loop): Add safelen and force_vect fields.
2640 * function.h (struct function): Add has_force_vect_loops field.
2641 * omp-low.c (expand_omp_simd): If !broken_loop, create loop for
2642 the simd region and set safelen and force_vect fields in it.
2643 * tree-vectorizer.c (vectorize_loops): If loop has force_vect set,
2644 vectorize it even if flag_vectorize isn't set. Clear loop->force_vect
2645 after vectorization.
2646 * tree-ssa-loop.c (gate_tree_vectorize): Return true even
2647 cfun->has_force_vect_loops.
2648 * tree-ssa-loop-ivcanon.c (tree_unroll_loops_completely_1): Don't
2649 unroll loops with loop->force_vect.
2650 * tree-vect-data-refs.c (vect_analyze_data_ref_dependence): For
2651 unknown or bad data dependency, if loop->safelen is non-zero, just
2652 decrease *max_vf to loop->safelen if needed and return false.
2653 * tree-if-conv.c (main_tree_if_conversion): If-convert also loops with
2654 loop->force_vect.
2655 (gate_tree_if_conversion): Return true even if
2656 cfun->has_force_vect_loops.
2658 2013-05-09 Jakub Jelinek <jakub@redhat.com>
2660 * tree.c (omp_declare_simd_clauses_equal): New function.
2661 (attribute_value_equal): Call it for -fopenmp if
2662 TREE_VALUE of the attributes are both OMP_CLAUSEs.
2663 * tree.h (omp_declare_simd_clauses_equal): Declare.
2665 2013-04-30 Jakub Jelinek <jakub@redhat.com>
2667 * gimple-pretty-print.c (dump_gimple_omp_atomic_load,
2668 dump_gimple_omp_atomic_store): Handle gimple_omp_atomic_seq_cst_p.
2669 * gimple.h (enum gf_mask): Add GF_OMP_ATOMIC_SEQ_CST.
2670 (gimple_omp_atomic_set_seq_cst, gimple_omp_atomic_seq_cst_p): New
2671 inline functions.
2672 * omp-low.c (expand_omp_atomic_load, expand_omp_atomic_store,
2673 expand_omp_atomic_fetch_op): If gimple_omp_atomic_seq_cst_p,
2674 pass MEMMODEL_SEQ_CST instead of MEMMODEL_RELAXED to the builtin.
2675 * gimplify.c (gimplify_omp_atomic): Handle OMP_ATOMIC_SEQ_CST.
2676 * tree-pretty-print.c (dump_generic_node): Handle OMP_ATOMIC_SEQ_CST.
2677 * tree.def (OMP_ATOMIC): Add comment that OMP_ATOMIC* must stay
2678 consecutive.
2679 * tree.h (OMP_ATOMIC_SEQ_CST): Define.
2681 * omp-low.c (check_omp_nesting_restrictions): Diagnose
2682 OpenMP constructs nested inside simd region. Don't treat
2683 #pragma omp simd as work-sharing region. Disallow work-sharing
2684 constructs inside of critical region. Complain if ordered
2685 region is nested inside of parallel region without loop
2686 region in between.
2687 (scan_omp_1_stmt): Call check_omp_nesting_restrictions even
2688 for GOMP_{cancel{,lation_point},taskyield,taskwait} calls.
2690 2013-04-23 Jakub Jelinek <jakub@redhat.com>
2692 * Makefile.in (omp-low.o): Depend on $(TARGET_H).
2693 * gimplify.c (gimplify_adjust_omp_clauses): For linear clauses
2694 if outer_context is non-NULL, but not ORT_COMBINED_PARALLEL,
2695 call omp_notice_variable. Remove aligned clauses that can't
2696 be handled yet.
2697 * omp-low.c: Include target.h.
2698 (scan_sharing_clauses): For aligned clauses with global arrays
2699 register local replacement.
2700 (omp_clause_aligned_alignment): New function.
2701 (lower_rec_input_clauses): For aligned clauses for global
2702 arrays or automatic pointers emit __builtin_assume_aligned
2703 before the loop if possible.
2704 (expand_omp_regimplify_p, expand_omp_build_assign): New functions.
2705 (expand_omp_simd): Use them. Handle pointer iterators and broken
2706 loops.
2707 (lower_omp_for): Call lower_omp on gimple_omp_body_ptr after
2708 calling lower_rec_input_clauses, not before it.
2710 2013-04-19 Jakub Jelinek <jakub@redhat.com>
2712 * tree.h (OMP_CLAUSE_LINEAR_NO_COPYIN,
2713 OMP_CLAUSE_LINEAR_NO_COPYOUT): Define.
2714 * omp-low.c (extract_omp_for_data): Handle #pragma omp simd.
2715 (build_outer_var_ref): For #pragma omp simd allow linear etc.
2716 clauses to bind even to private vars.
2717 (scan_sharing_clauses): Handle OMP_CLAUSE_LINEAR, OMP_CLAUSE_ALIGNED
2718 and OMP_CLAUSE_SAFELEN.
2719 (lower_rec_input_clauses): Handle OMP_CLAUSE_LINEAR. Don't emit
2720 a GOMP_barrier call for firstprivate/lastprivate in #pragma omp simd.
2721 (lower_lastprivate_clauses): Handle also OMP_CLAUSE_LINEAR.
2722 (expand_omp_simd): New function.
2723 (expand_omp_for): Handle #pragma omp simd.
2724 * gimplify.c (enum gimplify_omp_var_data): Add GOVD_LINEAR and
2725 GOVD_ALIGNED, add GOVD_LINEAR into GOVD_DATA_SHARE_CLASS.
2726 (enum omp_region_type): Add ORT_SIMD.
2727 (gimple_add_tmp_var, gimplify_var_or_parm_decl, omp_check_private,
2728 omp_firstprivatize_variable, omp_notice_variable): Handle ORT_SIMD
2729 like ORT_WORKSHARE.
2730 (omp_is_private): Likewise. Add SIMD argument, tweak diagnostics
2731 and add extra errors in simd constructs.
2732 (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle
2733 OMP_CLAUSE_LINEAR, OMP_CLAUSE_ALIGNED and OMP_CLAUSE_SAFELEN.
2734 (gimplify_adjust_omp_clauses_1): Handle GOVD_LASTPRIVATE and
2735 GOVD_ALIGNED.
2736 (gimplify_omp_for): Handle #pragma omp simd.
2738 2013-04-10 Jakub Jelinek <jakub@redhat.com>
2740 * builtin-types.def (DEF_FUNCTION_TYPE_8): Document.
2741 (BT_FN_VOID_OMPFN_PTR_UINT, BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG,
2742 BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG): Remove.
2743 (BT_FN_VOID_OMPFN_PTR_UINT_UINT_UINT,
2744 BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_UINT,
2745 BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT): New.
2746 * gimplify.c (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses):
2747 Handle OMP_CLAUSE_PROC_BIND.
2748 * omp-builtins.def (BUILT_IN_GOMP_TASKGROUP_START,
2749 BUILT_IN_GOMP_TASKGROUP_END, BUILT_IN_GOMP_PARALLEL_LOOP_STATIC,
2750 BUILT_IN_GOMP_PARALLEL_LOOP_DYNAMIC,
2751 BUILT_IN_GOMP_PARALLEL_LOOP_GUIDED,
2752 BUILT_IN_GOMP_PARALLEL_LOOP_RUNTIME, BUILT_IN_GOMP_PARALLEL,
2753 BUILT_IN_GOMP_PARALLEL_SECTIONS): New built-ins.
2754 (BUILT_IN_GOMP_PARALLEL_LOOP_STATIC_START,
2755 BUILT_IN_GOMP_PARALLEL_LOOP_DYNAMIC_START,
2756 BUILT_IN_GOMP_PARALLEL_LOOP_GUIDED_START,
2757 BUILT_IN_GOMP_PARALLEL_LOOP_RUNTIME_START,
2758 BUILT_IN_GOMP_PARALLEL_START, BUILT_IN_GOMP_PARALLEL_END,
2759 BUILT_IN_GOMP_PARALLEL_SECTIONS_START): Remove.
2760 * omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_PROC_BIND.
2761 (expand_parallel_call): Expand #pragma omp parallel* as
2762 calls to the new GOMP_parallel_* APIs without _start at the end,
2763 instead of GOMP_parallel_*_start followed by fn.omp_fn.N call,
2764 followed by GOMP_parallel_end. Handle OMP_CLAUSE_PROC_BIND.
2765 * tree-ssa-alias.c (ref_maybe_used_by_call_p_1,
2766 call_may_clobber_ref_p_1): Handle BUILT_IN_GOMP_TASKGROUP_END
2767 instead of BUILT_IN_GOMP_PARALLEL_END.
2769 2013-03-27 Jakub Jelinek <jakub@redhat.com>
2771 * gimple-pretty-print.c (dump_gimple_omp_for): Handle different
2772 GIMPLE_OMP_FOR kinds.
2773 * tree.def (OMP_SIMD, OMP_FOR_SIMD, OMP_DISTRIBUTE): New tree codes.
2774 * gimple.h (enum gf_mask): Add GF_OMP_FOR_KIND_MASK,
2775 GF_OMP_FOR_KIND_FOR, GF_OMP_FOR_KIND_SIMD, GF_OMP_FOR_KIND_FOR_SIMD
2776 and GF_OMP_FOR_KIND_DISTRIBUTE.
2777 (gimple_omp_for_kind, gimple_omp_for_set_kind): New inline functions.
2778 * gimplify.c (is_gimple_stmt, gimplify_omp_for, gimplify_expr): Handle
2779 OMP_SIMD, OMP_FOR_SIMD and OMP_DISTRIBUTE.
2780 * tree.c (omp_clause_num_ops, omp_clause_code_name, walk_tree_1):
2781 Handle new OpenMP 4.0 clauses.
2782 * tree-pretty-print.c (dump_omp_clause): Likewise.
2783 (dump_generic_node): Handle OMP_SIMD, OMP_FOR_SIMD and OMP_DISTRIBUTE.
2784 * tree.h (enum omp_clause_code): Add OMP_CLAUSE_LINEAR,
2785 OMP_CLAUSE_ALIGNED, OMP_CLAUSE_DEPEND, OMP_CLAUSE_FROM, OMP_CLAUSE_TO,
2786 OMP_CLAUSE_UNIFORM, OMP_CLAUSE_MAP, OMP_CLAUSE_DEVICE,
2787 OMP_CLAUSE_DIST_SCHEDULE, OMP_CLAUSE_INBRANCH, OMP_CLAUSE_NOTINBRANCH,
2788 OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_PROC_BIND, OMP_CLAUSE_SAFELEN,
2789 OMP_CLAUSE_SIMDLEN, OMP_CLAUSE_FOR, OMP_CLAUSE_PARALLEL,
2790 OMP_CLAUSE_SECTIONS and OMP_CLAUSE_TASKGROUP.
2791 (OMP_LOOP_CHECK): Define.
2792 (OMP_FOR_BODY, OMP_FOR_CLAUSES, OMP_FOR_INIT, OMP_FOR_COND,
2793 OMP_FOR_INCR, OMP_FOR_PRE_BODY): Use OMP_LOOP_CHECK instead of
2794 OMP_FOR_CHECK.
2795 (OMP_CLAUSE_DECL): Extend check range up to OMP_CLAUSE_MAP.
2796 (OMP_CLAUSE_LINEAR_STEP, OMP_CLAUSE_ALIGNED_ALIGNMENT,
2797 OMP_CLAUSE_NUM_TEAMS_EXPR, OMP_CLAUSE_DEVICE_ID,
2798 OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR, OMP_CLAUSE_SAFELEN_EXPR,
2799 OMP_CLAUSE_SIMDLEN_EXPR): Define.
2800 (enum omp_clause_depend_kind, enum omp_clause_map_kind,
2801 enum omp_clause_proc_bind_kind): New enums.
2802 (OMP_CLAUSE_DEPEND_KIND, OMP_CLAUSE_MAP_KIND,
2803 OMP_CLAUSE_PROC_BIND_KIND): Define.
2804 (struct tree_omp_clause): Add subcode.depend_kind, subcode.map_kind
2805 and subcode.proc_bind_kind.
2806 (find_omp_clause): New prototype.
2807 * omp-builtins.def (BUILT_IN_GOMP_CANCEL,
2808 BUILT_IN_GOMP_CANCELLATION_POINT): New built-ins.
2809 * tree-flow.h (find_omp_clause): Remove prototype.
2810 \f
2811 Copyright (C) 2013-2015 Free Software Foundation, Inc.
2813 Copying and distribution of this file, with or without modification,
2814 are permitted in any medium without royalty provided the copyright
2815 notice and this notice are preserved.
2817 Local Variables:
2818 mode: change-log
2819 change-log-default-name: "ChangeLog.gomp"
2820 End: