1 /* A pass for lowering gimple to HSAIL
2 Copyright (C) 2013-2016 Free Software Foundation, Inc.
3 Contributed by Martin Jambor <mjambor@suse.cz> and
4 Martin Liska <mliska@suse.cz>.
6 This file is part of GCC.
8 GCC is free software; you can redistribute it and/or modify
9 it under the terms of the GNU General Public License as published by
10 the Free Software Foundation; either version 3, or (at your option)
13 GCC is distributed in the hope that it will be useful,
14 but WITHOUT ANY WARRANTY; without even the implied warranty of
15 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
16 GNU General Public License for more details.
18 You should have received a copy of the GNU General Public License
19 along with GCC; see the file COPYING3. If not see
20 <http://www.gnu.org/licenses/>. */
24 #include "coretypes.h"
27 #include "hash-table.h"
30 #include "tree-pass.h"
33 #include "basic-block.h"
34 #include "fold-const.h"
36 #include "gimple-iterator.h"
39 #include "gimple-pretty-print.h"
40 #include "diagnostic-core.h"
41 #include "alloc-pool.h"
42 #include "gimple-ssa.h"
43 #include "tree-phinodes.h"
44 #include "stringpool.h"
45 #include "tree-ssanames.h"
47 #include "ssa-iterators.h"
49 #include "print-tree.h"
50 #include "symbol-summary.h"
58 #include "gomp-constants.h"
59 #include "internal-fn.h"
61 #include "stor-layout.h"
63 /* Print a warning message and set that we have seen an error. */
65 #define HSA_SORRY_ATV(location, message, ...) \
69 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
71 inform (location, message, __VA_ARGS__); \
75 /* Same as previous, but highlight a location. */
77 #define HSA_SORRY_AT(location, message) \
81 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
83 inform (location, message); \
87 /* Default number of threads used by kernel dispatch. */
89 #define HSA_DEFAULT_NUM_THREADS 64
91 /* Following structures are defined in the final version
92 of HSA specification. */
94 /* HSA queue packet is shadow structure, originally provided by AMD. */
96 struct hsa_queue_packet
100 uint16_t workgroup_size_x
;
101 uint16_t workgroup_size_y
;
102 uint16_t workgroup_size_z
;
104 uint32_t grid_size_x
;
105 uint32_t grid_size_y
;
106 uint32_t grid_size_z
;
107 uint32_t private_segment_size
;
108 uint32_t group_segment_size
;
109 uint64_t kernel_object
;
110 void *kernarg_address
;
112 uint64_t completion_signal
;
115 /* HSA queue is shadow structure, originally provided by AMD. */
122 uint64_t doorbell_signal
;
128 /* Alloc pools for allocating basic hsa structures such as operands,
129 instructions and other basic entities. */
130 static object_allocator
<hsa_op_address
> *hsa_allocp_operand_address
;
131 static object_allocator
<hsa_op_immed
> *hsa_allocp_operand_immed
;
132 static object_allocator
<hsa_op_reg
> *hsa_allocp_operand_reg
;
133 static object_allocator
<hsa_op_code_list
> *hsa_allocp_operand_code_list
;
134 static object_allocator
<hsa_op_operand_list
> *hsa_allocp_operand_operand_list
;
135 static object_allocator
<hsa_insn_basic
> *hsa_allocp_inst_basic
;
136 static object_allocator
<hsa_insn_phi
> *hsa_allocp_inst_phi
;
137 static object_allocator
<hsa_insn_mem
> *hsa_allocp_inst_mem
;
138 static object_allocator
<hsa_insn_atomic
> *hsa_allocp_inst_atomic
;
139 static object_allocator
<hsa_insn_signal
> *hsa_allocp_inst_signal
;
140 static object_allocator
<hsa_insn_seg
> *hsa_allocp_inst_seg
;
141 static object_allocator
<hsa_insn_cmp
> *hsa_allocp_inst_cmp
;
142 static object_allocator
<hsa_insn_br
> *hsa_allocp_inst_br
;
143 static object_allocator
<hsa_insn_sbr
> *hsa_allocp_inst_sbr
;
144 static object_allocator
<hsa_insn_call
> *hsa_allocp_inst_call
;
145 static object_allocator
<hsa_insn_arg_block
> *hsa_allocp_inst_arg_block
;
146 static object_allocator
<hsa_insn_comment
> *hsa_allocp_inst_comment
;
147 static object_allocator
<hsa_insn_queue
> *hsa_allocp_inst_queue
;
148 static object_allocator
<hsa_insn_srctype
> *hsa_allocp_inst_srctype
;
149 static object_allocator
<hsa_insn_packed
> *hsa_allocp_inst_packed
;
150 static object_allocator
<hsa_insn_cvt
> *hsa_allocp_inst_cvt
;
151 static object_allocator
<hsa_insn_alloca
> *hsa_allocp_inst_alloca
;
152 static object_allocator
<hsa_bb
> *hsa_allocp_bb
;
154 /* List of pointers to all instructions that come from an object allocator. */
155 static vec
<hsa_insn_basic
*> hsa_instructions
;
157 /* List of pointers to all operands that come from an object allocator. */
158 static vec
<hsa_op_base
*> hsa_operands
;
160 hsa_symbol::hsa_symbol ()
161 : m_decl (NULL_TREE
), m_name (NULL
), m_name_number (0),
162 m_directive_offset (0), m_type (BRIG_TYPE_NONE
),
163 m_segment (BRIG_SEGMENT_NONE
), m_linkage (BRIG_LINKAGE_NONE
), m_dim (0),
164 m_cst_value (NULL
), m_global_scope_p (false), m_seen_error (false),
165 m_allocation (BRIG_ALLOCATION_AUTOMATIC
), m_emitted_to_brig (false)
170 hsa_symbol::hsa_symbol (BrigType16_t type
, BrigSegment8_t segment
,
171 BrigLinkage8_t linkage
, bool global_scope_p
,
172 BrigAllocation allocation
, BrigAlignment8_t align
)
173 : m_decl (NULL_TREE
), m_name (NULL
), m_name_number (0),
174 m_directive_offset (0), m_type (type
), m_segment (segment
),
175 m_linkage (linkage
), m_dim (0), m_cst_value (NULL
),
176 m_global_scope_p (global_scope_p
), m_seen_error (false),
177 m_allocation (allocation
), m_emitted_to_brig (false), m_align (align
)
181 unsigned HOST_WIDE_INT
182 hsa_symbol::total_byte_size ()
184 unsigned HOST_WIDE_INT s
185 = hsa_type_bit_size (~BRIG_TYPE_ARRAY_MASK
& m_type
);
186 gcc_assert (s
% BITS_PER_UNIT
== 0);
195 /* Forward declaration. */
198 hsa_type_for_tree_type (const_tree type
, unsigned HOST_WIDE_INT
*dim_p
,
202 hsa_symbol::fillup_for_decl (tree decl
)
205 m_type
= hsa_type_for_tree_type (TREE_TYPE (decl
), &m_dim
, false);
206 if (hsa_seen_error ())
212 m_align
= MAX (m_align
, hsa_natural_alignment (m_type
));
215 /* Constructor of class representing global HSA function/kernel information and
216 state. FNDECL is function declaration, KERNEL_P is true if the function
217 is going to become a HSA kernel. If the function has body, SSA_NAMES_COUNT
218 should be set to number of SSA names used in the function.
219 MODIFIED_CFG is set to true in case we modified control-flow graph
222 hsa_function_representation::hsa_function_representation
223 (tree fdecl
, bool kernel_p
, unsigned ssa_names_count
, bool modified_cfg
)
225 m_reg_count (0), m_input_args (vNULL
),
226 m_output_arg (NULL
), m_spill_symbols (vNULL
), m_global_symbols (vNULL
),
227 m_private_variables (vNULL
), m_called_functions (vNULL
),
228 m_called_internal_fns (vNULL
), m_hbb_count (0),
229 m_in_ssa (true), m_kern_p (kernel_p
), m_declaration_p (false),
230 m_decl (fdecl
), m_internal_fn (NULL
), m_shadow_reg (NULL
),
231 m_kernel_dispatch_count (0), m_maximum_omp_data_size (0),
232 m_seen_error (false), m_temp_symbol_count (0), m_ssa_map (),
233 m_modified_cfg (modified_cfg
)
235 int sym_init_len
= (vec_safe_length (cfun
->local_decls
) / 2) + 1;;
236 m_local_symbols
= new hash_table
<hsa_noop_symbol_hasher
> (sym_init_len
);
237 m_ssa_map
.safe_grow_cleared (ssa_names_count
);
240 /* Constructor of class representing HSA function information that
241 is derived for an internal function. */
242 hsa_function_representation::hsa_function_representation (hsa_internal_fn
*fn
)
243 : m_reg_count (0), m_input_args (vNULL
),
244 m_output_arg (NULL
), m_local_symbols (NULL
),
245 m_spill_symbols (vNULL
), m_global_symbols (vNULL
),
246 m_private_variables (vNULL
), m_called_functions (vNULL
),
247 m_called_internal_fns (vNULL
), m_hbb_count (0),
248 m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL
),
249 m_internal_fn (fn
), m_shadow_reg (NULL
), m_kernel_dispatch_count (0),
250 m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
253 /* Destructor of class holding function/kernel-wide information and state. */
255 hsa_function_representation::~hsa_function_representation ()
257 /* Kernel names are deallocated at the end of BRIG output when deallocating
258 hsa_decl_kernel_mapping. */
259 if (!m_kern_p
|| m_seen_error
)
262 for (unsigned i
= 0; i
< m_input_args
.length (); i
++)
263 delete m_input_args
[i
];
264 m_input_args
.release ();
267 delete m_local_symbols
;
269 for (unsigned i
= 0; i
< m_spill_symbols
.length (); i
++)
270 delete m_spill_symbols
[i
];
271 m_spill_symbols
.release ();
274 for (unsigned i
= 0; i
< m_global_symbols
.iterate (i
, &sym
); i
++)
275 if (sym
->m_linkage
!= BRIG_ALLOCATION_PROGRAM
)
277 m_global_symbols
.release ();
279 for (unsigned i
= 0; i
< m_private_variables
.length (); i
++)
280 delete m_private_variables
[i
];
281 m_private_variables
.release ();
282 m_called_functions
.release ();
283 m_ssa_map
.release ();
285 for (unsigned i
= 0; i
< m_called_internal_fns
.length (); i
++)
286 delete m_called_internal_fns
[i
];
290 hsa_function_representation::get_shadow_reg ()
292 /* If we compile a function with kernel dispatch and does not set
293 an optimization level, the function won't be inlined and
301 /* Append the shadow argument. */
302 hsa_symbol
*shadow
= new hsa_symbol (BRIG_TYPE_U64
, BRIG_SEGMENT_KERNARG
,
303 BRIG_LINKAGE_FUNCTION
);
304 m_input_args
.safe_push (shadow
);
305 shadow
->m_name
= "hsa_runtime_shadow";
307 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_U64
);
308 hsa_op_address
*addr
= new hsa_op_address (shadow
);
310 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, BRIG_TYPE_U64
, r
, addr
);
311 hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
))->append_insn (mem
);
317 bool hsa_function_representation::has_shadow_reg_p ()
319 return m_shadow_reg
!= NULL
;
323 hsa_function_representation::init_extra_bbs ()
325 hsa_init_new_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
326 hsa_init_new_bb (EXIT_BLOCK_PTR_FOR_FN (cfun
));
330 hsa_function_representation::update_dominance ()
334 free_dominance_info (CDI_DOMINATORS
);
335 calculate_dominance_info (CDI_DOMINATORS
);
340 hsa_function_representation::create_hsa_temporary (BrigType16_t type
)
342 hsa_symbol
*s
= new hsa_symbol (type
, BRIG_SEGMENT_PRIVATE
,
343 BRIG_LINKAGE_FUNCTION
);
344 s
->m_name_number
= m_temp_symbol_count
++;
346 hsa_cfun
->m_private_variables
.safe_push (s
);
351 hsa_function_representation::get_linkage ()
354 return BRIG_LINKAGE_PROGRAM
;
356 return m_kern_p
|| TREE_PUBLIC (m_decl
) ?
357 BRIG_LINKAGE_PROGRAM
: BRIG_LINKAGE_MODULE
;
360 /* Hash map of simple OMP builtins. */
361 static hash_map
<nofree_string_hash
, omp_simple_builtin
> *omp_simple_builtins
364 /* Warning messages for OMP builtins. */
366 #define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP " \
368 #define HSA_WARN_TIMING_ROUTINE "support for HSA does not implement OpenMP " \
370 #define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \
371 "undefined semantics within target regions, support for HSA ignores them"
372 #define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP " \
375 /* Initialize hash map with simple OMP builtins. */
378 hsa_init_simple_builtins ()
380 if (omp_simple_builtins
!= NULL
)
384 = new hash_map
<nofree_string_hash
, omp_simple_builtin
> ();
386 omp_simple_builtin omp_builtins
[] =
388 omp_simple_builtin ("omp_get_initial_device", NULL
, false,
389 new hsa_op_immed (GOMP_DEVICE_HOST
,
390 (BrigType16_t
) BRIG_TYPE_S32
)),
391 omp_simple_builtin ("omp_is_initial_device", NULL
, false,
392 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
393 omp_simple_builtin ("omp_get_dynamic", NULL
, false,
394 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
395 omp_simple_builtin ("omp_set_dynamic", NULL
, false, NULL
),
396 omp_simple_builtin ("omp_init_lock", HSA_WARN_LOCK_ROUTINE
, true),
397 omp_simple_builtin ("omp_init_lock_with_hint", HSA_WARN_LOCK_ROUTINE
,
399 omp_simple_builtin ("omp_init_nest_lock_with_hint", HSA_WARN_LOCK_ROUTINE
,
401 omp_simple_builtin ("omp_destroy_lock", HSA_WARN_LOCK_ROUTINE
, true),
402 omp_simple_builtin ("omp_set_lock", HSA_WARN_LOCK_ROUTINE
, true),
403 omp_simple_builtin ("omp_unset_lock", HSA_WARN_LOCK_ROUTINE
, true),
404 omp_simple_builtin ("omp_test_lock", HSA_WARN_LOCK_ROUTINE
, true),
405 omp_simple_builtin ("omp_get_wtime", HSA_WARN_TIMING_ROUTINE
, true),
406 omp_simple_builtin ("omp_get_wtick", HSA_WARN_TIMING_ROUTINE
, true),
407 omp_simple_builtin ("omp_target_alloc", HSA_WARN_MEMORY_ROUTINE
, false,
408 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_U64
)),
409 omp_simple_builtin ("omp_target_free", HSA_WARN_MEMORY_ROUTINE
, false),
410 omp_simple_builtin ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE
,
412 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
413 omp_simple_builtin ("omp_target_memcpy", HSA_WARN_MEMORY_ROUTINE
, false,
414 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
415 omp_simple_builtin ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE
,
417 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
418 omp_simple_builtin ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE
,
420 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
421 omp_simple_builtin ("omp_target_disassociate_ptr",
422 HSA_WARN_MEMORY_ROUTINE
,
424 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
425 omp_simple_builtin ("omp_set_max_active_levels",
426 "Support for HSA only allows only one active level, "
427 "call to omp_set_max_active_levels will be ignored "
428 "in the generated HSAIL",
430 omp_simple_builtin ("omp_get_max_active_levels", NULL
, false,
431 new hsa_op_immed (1, (BrigType16_t
) BRIG_TYPE_S32
)),
432 omp_simple_builtin ("omp_in_final", NULL
, false,
433 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
434 omp_simple_builtin ("omp_get_proc_bind", HSA_WARN_AFFINITY
, false,
435 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
436 omp_simple_builtin ("omp_get_num_places", HSA_WARN_AFFINITY
, false,
437 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
438 omp_simple_builtin ("omp_get_place_num_procs", HSA_WARN_AFFINITY
, false,
439 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
440 omp_simple_builtin ("omp_get_place_proc_ids", HSA_WARN_AFFINITY
, false,
442 omp_simple_builtin ("omp_get_place_num", HSA_WARN_AFFINITY
, false,
443 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
444 omp_simple_builtin ("omp_get_partition_num_places", HSA_WARN_AFFINITY
,
446 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
447 omp_simple_builtin ("omp_get_partition_place_nums", HSA_WARN_AFFINITY
,
449 omp_simple_builtin ("omp_set_default_device",
450 "omp_set_default_device has undefined semantics "
451 "within target regions, support for HSA ignores it",
453 omp_simple_builtin ("omp_get_default_device",
454 "omp_get_default_device has undefined semantics "
455 "within target regions, support for HSA ignores it",
457 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
458 omp_simple_builtin ("omp_get_num_devices",
459 "omp_get_num_devices has undefined semantics "
460 "within target regions, support for HSA ignores it",
462 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
463 omp_simple_builtin ("omp_get_num_procs", NULL
, true, NULL
),
464 omp_simple_builtin ("omp_get_cancellation", NULL
, true, NULL
),
465 omp_simple_builtin ("omp_set_nested", NULL
, true, NULL
),
466 omp_simple_builtin ("omp_get_nested", NULL
, true, NULL
),
467 omp_simple_builtin ("omp_set_schedule", NULL
, true, NULL
),
468 omp_simple_builtin ("omp_get_schedule", NULL
, true, NULL
),
469 omp_simple_builtin ("omp_get_thread_limit", NULL
, true, NULL
),
470 omp_simple_builtin ("omp_get_team_size", NULL
, true, NULL
),
471 omp_simple_builtin ("omp_get_ancestor_thread_num", NULL
, true, NULL
),
472 omp_simple_builtin ("omp_get_max_task_priority", NULL
, true, NULL
)
475 unsigned count
= sizeof (omp_builtins
) / sizeof (omp_simple_builtin
);
477 for (unsigned i
= 0; i
< count
; i
++)
478 omp_simple_builtins
->put (omp_builtins
[i
].m_name
, omp_builtins
[i
]);
481 /* Allocate HSA structures that we need only while generating with this. */
484 hsa_init_data_for_cfun ()
486 hsa_init_compilation_unit_data ();
487 hsa_allocp_operand_address
488 = new object_allocator
<hsa_op_address
> ("HSA address operands");
489 hsa_allocp_operand_immed
490 = new object_allocator
<hsa_op_immed
> ("HSA immediate operands");
491 hsa_allocp_operand_reg
492 = new object_allocator
<hsa_op_reg
> ("HSA register operands");
493 hsa_allocp_operand_code_list
494 = new object_allocator
<hsa_op_code_list
> ("HSA code list operands");
495 hsa_allocp_operand_operand_list
496 = new object_allocator
<hsa_op_operand_list
> ("HSA operand list operands");
497 hsa_allocp_inst_basic
498 = new object_allocator
<hsa_insn_basic
> ("HSA basic instructions");
500 = new object_allocator
<hsa_insn_phi
> ("HSA phi operands");
502 = new object_allocator
<hsa_insn_mem
> ("HSA memory instructions");
503 hsa_allocp_inst_atomic
504 = new object_allocator
<hsa_insn_atomic
> ("HSA atomic instructions");
505 hsa_allocp_inst_signal
506 = new object_allocator
<hsa_insn_signal
> ("HSA signal instructions");
508 = new object_allocator
<hsa_insn_seg
> ("HSA segment conversion "
511 = new object_allocator
<hsa_insn_cmp
> ("HSA comparison instructions");
513 = new object_allocator
<hsa_insn_br
> ("HSA branching instructions");
515 = new object_allocator
<hsa_insn_sbr
> ("HSA switch branching instructions");
517 = new object_allocator
<hsa_insn_call
> ("HSA call instructions");
518 hsa_allocp_inst_arg_block
519 = new object_allocator
<hsa_insn_arg_block
> ("HSA arg block instructions");
520 hsa_allocp_inst_comment
521 = new object_allocator
<hsa_insn_comment
> ("HSA comment instructions");
522 hsa_allocp_inst_queue
523 = new object_allocator
<hsa_insn_queue
> ("HSA queue instructions");
524 hsa_allocp_inst_srctype
525 = new object_allocator
<hsa_insn_srctype
> ("HSA source type instructions");
526 hsa_allocp_inst_packed
527 = new object_allocator
<hsa_insn_packed
> ("HSA packed instructions");
529 = new object_allocator
<hsa_insn_cvt
> ("HSA convert instructions");
530 hsa_allocp_inst_alloca
531 = new object_allocator
<hsa_insn_alloca
> ("HSA alloca instructions");
532 hsa_allocp_bb
= new object_allocator
<hsa_bb
> ("HSA basic blocks");
535 /* Deinitialize HSA subsystem and free all allocated memory. */
538 hsa_deinit_data_for_cfun (void)
542 FOR_ALL_BB_FN (bb
, cfun
)
545 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
550 for (unsigned int i
= 0; i
< hsa_operands
.length (); i
++)
551 hsa_destroy_operand (hsa_operands
[i
]);
553 hsa_operands
.release ();
555 for (unsigned i
= 0; i
< hsa_instructions
.length (); i
++)
556 hsa_destroy_insn (hsa_instructions
[i
]);
558 hsa_instructions
.release ();
560 if (omp_simple_builtins
!= NULL
)
562 delete omp_simple_builtins
;
563 omp_simple_builtins
= NULL
;
566 delete hsa_allocp_operand_address
;
567 delete hsa_allocp_operand_immed
;
568 delete hsa_allocp_operand_reg
;
569 delete hsa_allocp_operand_code_list
;
570 delete hsa_allocp_operand_operand_list
;
571 delete hsa_allocp_inst_basic
;
572 delete hsa_allocp_inst_phi
;
573 delete hsa_allocp_inst_atomic
;
574 delete hsa_allocp_inst_mem
;
575 delete hsa_allocp_inst_signal
;
576 delete hsa_allocp_inst_seg
;
577 delete hsa_allocp_inst_cmp
;
578 delete hsa_allocp_inst_br
;
579 delete hsa_allocp_inst_sbr
;
580 delete hsa_allocp_inst_call
;
581 delete hsa_allocp_inst_arg_block
;
582 delete hsa_allocp_inst_comment
;
583 delete hsa_allocp_inst_queue
;
584 delete hsa_allocp_inst_srctype
;
585 delete hsa_allocp_inst_packed
;
586 delete hsa_allocp_inst_cvt
;
587 delete hsa_allocp_inst_alloca
;
588 delete hsa_allocp_bb
;
592 /* Return the type which holds addresses in the given SEGMENT. */
595 hsa_get_segment_addr_type (BrigSegment8_t segment
)
599 case BRIG_SEGMENT_NONE
:
602 case BRIG_SEGMENT_FLAT
:
603 case BRIG_SEGMENT_GLOBAL
:
604 case BRIG_SEGMENT_READONLY
:
605 case BRIG_SEGMENT_KERNARG
:
606 return hsa_machine_large_p () ? BRIG_TYPE_U64
: BRIG_TYPE_U32
;
608 case BRIG_SEGMENT_GROUP
:
609 case BRIG_SEGMENT_PRIVATE
:
610 case BRIG_SEGMENT_SPILL
:
611 case BRIG_SEGMENT_ARG
:
612 return BRIG_TYPE_U32
;
617 /* Return integer brig type according to provided SIZE in bytes. If SIGN
618 is set to true, return signed integer type. */
621 get_integer_type_by_bytes (unsigned size
, bool sign
)
629 return BRIG_TYPE_S16
;
631 return BRIG_TYPE_S32
;
633 return BRIG_TYPE_S64
;
643 return BRIG_TYPE_U16
;
645 return BRIG_TYPE_U32
;
647 return BRIG_TYPE_U64
;
655 /* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
656 are assumed to use flat addressing. If min32int is true, always expand
657 integer types to one that has at least 32 bits. */
660 hsa_type_for_scalar_tree_type (const_tree type
, bool min32int
)
664 BrigType16_t res
= BRIG_TYPE_NONE
;
666 gcc_checking_assert (TYPE_P (type
));
667 gcc_checking_assert (!AGGREGATE_TYPE_P (type
));
668 if (POINTER_TYPE_P (type
))
669 return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
671 if (TREE_CODE (type
) == VECTOR_TYPE
|| TREE_CODE (type
) == COMPLEX_TYPE
)
672 base
= TREE_TYPE (type
);
676 if (!tree_fits_uhwi_p (TYPE_SIZE (base
)))
678 HSA_SORRY_ATV (EXPR_LOCATION (type
),
679 "support for HSA does not implement huge or "
680 "variable-sized type %T", type
);
684 bsize
= tree_to_uhwi (TYPE_SIZE (base
));
685 unsigned byte_size
= bsize
/ BITS_PER_UNIT
;
686 if (INTEGRAL_TYPE_P (base
))
687 res
= get_integer_type_by_bytes (byte_size
, !TYPE_UNSIGNED (base
));
688 else if (SCALAR_FLOAT_TYPE_P (base
))
706 if (res
== BRIG_TYPE_NONE
)
708 HSA_SORRY_ATV (EXPR_LOCATION (type
),
709 "support for HSA does not implement type %T", type
);
713 if (TREE_CODE (type
) == VECTOR_TYPE
)
715 HOST_WIDE_INT tsize
= tree_to_uhwi (TYPE_SIZE (type
));
719 HSA_SORRY_ATV (EXPR_LOCATION (type
),
720 "support for HSA does not implement a vector type "
721 "where a type and unit size are equal: %T", type
);
728 res
|= BRIG_TYPE_PACK_32
;
731 res
|= BRIG_TYPE_PACK_64
;
734 res
|= BRIG_TYPE_PACK_128
;
737 HSA_SORRY_ATV (EXPR_LOCATION (type
),
738 "support for HSA does not implement type %T", type
);
744 /* Registers/immediate operands can only be 32bit or more except for
746 if (res
== BRIG_TYPE_U8
|| res
== BRIG_TYPE_U16
)
748 else if (res
== BRIG_TYPE_S8
|| res
== BRIG_TYPE_S16
)
752 if (TREE_CODE (type
) == COMPLEX_TYPE
)
754 unsigned bsize
= 2 * hsa_type_bit_size (res
);
755 res
= hsa_bittype_for_bitsize (bsize
);
761 /* Returns the BRIG type we need to load/store entities of TYPE. */
764 mem_type_for_type (BrigType16_t type
)
766 /* HSA has non-intuitive constraints on load/store types. If it's
767 a bit-type it _must_ be B128, if it's not a bit-type it must be
768 64bit max. So for loading entities of 128 bits (e.g. vectors)
769 we have to to B128, while for loading the rest we have to use the
770 input type (??? or maybe also flattened to a equally sized non-vector
772 if ((type
& BRIG_TYPE_PACK_MASK
) == BRIG_TYPE_PACK_128
)
773 return BRIG_TYPE_B128
;
774 else if (hsa_btype_p (type
) || hsa_type_packed_p (type
))
776 unsigned bitsize
= hsa_type_bit_size (type
);
778 return hsa_uint_for_bitsize (bitsize
);
780 return hsa_bittype_for_bitsize (bitsize
);
785 /* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some
786 kind of array will be generated, setting DIM appropriately. Otherwise, it
787 will be set to zero. */
790 hsa_type_for_tree_type (const_tree type
, unsigned HOST_WIDE_INT
*dim_p
= NULL
,
791 bool min32int
= false)
793 gcc_checking_assert (TYPE_P (type
));
794 if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type
)))
796 HSA_SORRY_ATV (EXPR_LOCATION (type
), "support for HSA does not "
797 "implement huge or variable-sized type %T", type
);
798 return BRIG_TYPE_NONE
;
801 if (RECORD_OR_UNION_TYPE_P (type
))
804 *dim_p
= tree_to_uhwi (TYPE_SIZE_UNIT (type
));
805 return BRIG_TYPE_U8
| BRIG_TYPE_ARRAY
;
808 if (TREE_CODE (type
) == ARRAY_TYPE
)
810 /* We try to be nice and use the real base-type when this is an array of
811 scalars and only resort to an array of bytes if the type is more
814 unsigned HOST_WIDE_INT dim
= 1;
816 while (TREE_CODE (type
) == ARRAY_TYPE
)
818 tree domain
= TYPE_DOMAIN (type
);
819 if (!TYPE_MIN_VALUE (domain
)
820 || !TYPE_MAX_VALUE (domain
)
821 || !tree_fits_shwi_p (TYPE_MIN_VALUE (domain
))
822 || !tree_fits_shwi_p (TYPE_MAX_VALUE (domain
)))
824 HSA_SORRY_ATV (EXPR_LOCATION (type
),
825 "support for HSA does not implement array %T with "
826 "unknown bounds", type
);
827 return BRIG_TYPE_NONE
;
829 HOST_WIDE_INT min
= tree_to_shwi (TYPE_MIN_VALUE (domain
));
830 HOST_WIDE_INT max
= tree_to_shwi (TYPE_MAX_VALUE (domain
));
831 dim
= dim
* (unsigned HOST_WIDE_INT
) (max
- min
+ 1);
832 type
= TREE_TYPE (type
);
836 if (RECORD_OR_UNION_TYPE_P (type
))
838 dim
= dim
* tree_to_uhwi (TYPE_SIZE_UNIT (type
));
842 res
= hsa_type_for_scalar_tree_type (type
, false);
846 return res
| BRIG_TYPE_ARRAY
;
853 return hsa_type_for_scalar_tree_type (type
, min32int
);
856 /* Returns true if converting from STYPE into DTYPE needs the _CVT
857 opcode. If false a normal _MOV is enough. */
860 hsa_needs_cvt (BrigType16_t dtype
, BrigType16_t stype
)
862 if (hsa_btype_p (dtype
))
865 /* float <-> int conversions are real converts. */
866 if (hsa_type_float_p (dtype
) != hsa_type_float_p (stype
))
868 /* When both types have different size, then we need CVT as well. */
869 if (hsa_type_bit_size (dtype
) != hsa_type_bit_size (stype
))
874 /* Lookup or create the associated hsa_symbol structure with a given VAR_DECL
875 or lookup the hsa_structure corresponding to a PARM_DECL. */
878 get_symbol_for_decl (tree decl
)
881 hsa_symbol
dummy (BRIG_TYPE_NONE
, BRIG_SEGMENT_NONE
, BRIG_LINKAGE_NONE
);
883 gcc_assert (TREE_CODE (decl
) == PARM_DECL
884 || TREE_CODE (decl
) == RESULT_DECL
885 || TREE_CODE (decl
) == VAR_DECL
);
889 bool is_in_global_vars
890 = TREE_CODE (decl
) == VAR_DECL
&& is_global_var (decl
);
892 if (is_in_global_vars
)
893 slot
= hsa_global_variable_symbols
->find_slot (&dummy
, INSERT
);
895 slot
= hsa_cfun
->m_local_symbols
->find_slot (&dummy
, INSERT
);
897 gcc_checking_assert (slot
);
900 hsa_symbol
*sym
= (*slot
);
902 /* If the symbol is problematic, mark current function also as
904 if (sym
->m_seen_error
)
907 /* PR hsa/70234: If a global variable was marked to be emitted,
908 but HSAIL generation of a function using the variable fails,
909 we should retry to emit the variable in context of a different
912 Iterate elements whether a symbol is already in m_global_symbols
914 if (is_in_global_vars
&& !sym
->m_emitted_to_brig
)
916 for (unsigned i
= 0; i
< hsa_cfun
->m_global_symbols
.length (); i
++)
917 if (hsa_cfun
->m_global_symbols
[i
] == sym
)
919 hsa_cfun
->m_global_symbols
.safe_push (sym
);
927 gcc_assert (TREE_CODE (decl
) == VAR_DECL
);
928 BrigAlignment8_t align
= hsa_object_alignment (decl
);
930 if (is_in_global_vars
)
932 sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_GLOBAL
,
933 BRIG_LINKAGE_PROGRAM
, true,
934 BRIG_ALLOCATION_PROGRAM
, align
);
935 hsa_cfun
->m_global_symbols
.safe_push (sym
);
936 sym
->fillup_for_decl (decl
);
937 if (sym
->m_align
> align
)
939 sym
->m_seen_error
= true;
940 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
941 "HSA specification requires that %E is at least "
942 "naturally aligned", decl
);
947 /* As generation of efficient memory copy instructions relies
948 on alignment greater or equal to 8 bytes,
949 we need to increase alignment of all aggregate types.. */
950 if (AGGREGATE_TYPE_P (TREE_TYPE (decl
)))
951 align
= MAX ((BrigAlignment8_t
) BRIG_ALIGNMENT_8
, align
);
953 /* PARM_DECL and RESULT_DECL should be already in m_local_symbols. */
954 gcc_assert (TREE_CODE (decl
) == VAR_DECL
);
956 sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_PRIVATE
,
957 BRIG_LINKAGE_FUNCTION
);
958 sym
->m_align
= align
;
959 sym
->fillup_for_decl (decl
);
960 hsa_cfun
->m_private_variables
.safe_push (sym
);
963 sym
->m_name
= hsa_get_declaration_name (decl
);
969 /* For a given HSA function declaration, return a host
970 function declaration. */
973 hsa_get_host_function (tree decl
)
975 hsa_function_summary
*s
976 = hsa_summaries
->get (cgraph_node::get_create (decl
));
977 gcc_assert (s
->m_kind
!= HSA_NONE
);
978 gcc_assert (s
->m_gpu_implementation_p
);
980 return s
->m_binded_function
->decl
;
983 /* Return true if function DECL has a host equivalent function. */
986 get_brig_function_name (tree decl
)
990 hsa_function_summary
*s
= hsa_summaries
->get (cgraph_node::get_create (d
));
991 if (s
->m_kind
!= HSA_NONE
&& s
->m_gpu_implementation_p
)
992 d
= s
->m_binded_function
->decl
;
994 /* IPA split can create a function that has no host equivalent. */
998 char *name
= xstrdup (hsa_get_declaration_name (d
));
999 hsa_sanitize_name (name
);
1004 /* Create a spill symbol of type TYPE. */
1007 hsa_get_spill_symbol (BrigType16_t type
)
1009 hsa_symbol
*sym
= new hsa_symbol (type
, BRIG_SEGMENT_SPILL
,
1010 BRIG_LINKAGE_FUNCTION
);
1011 hsa_cfun
->m_spill_symbols
.safe_push (sym
);
1015 /* Create a symbol for a read-only string constant. */
1017 hsa_get_string_cst_symbol (tree string_cst
)
1019 gcc_checking_assert (TREE_CODE (string_cst
) == STRING_CST
);
1021 hsa_symbol
**slot
= hsa_cfun
->m_string_constants_map
.get (string_cst
);
1025 hsa_op_immed
*cst
= new hsa_op_immed (string_cst
);
1026 hsa_symbol
*sym
= new hsa_symbol (cst
->m_type
, BRIG_SEGMENT_GLOBAL
,
1027 BRIG_LINKAGE_MODULE
, true,
1028 BRIG_ALLOCATION_AGENT
);
1029 sym
->m_cst_value
= cst
;
1030 sym
->m_dim
= TREE_STRING_LENGTH (string_cst
);
1031 sym
->m_name_number
= hsa_cfun
->m_global_symbols
.length ();
1033 hsa_cfun
->m_global_symbols
.safe_push (sym
);
1034 hsa_cfun
->m_string_constants_map
.put (string_cst
, sym
);
1038 /* Constructor of the ancestor of all operands. K is BRIG kind that identified
1039 what the operator is. */
1041 hsa_op_base::hsa_op_base (BrigKind16_t k
)
1042 : m_next (NULL
), m_brig_op_offset (0), m_kind (k
)
1044 hsa_operands
.safe_push (this);
1047 /* Constructor of ancestor of all operands which have a type. K is BRIG kind
1048 that identified what the operator is. T is the type of the operator. */
1050 hsa_op_with_type::hsa_op_with_type (BrigKind16_t k
, BrigType16_t t
)
1051 : hsa_op_base (k
), m_type (t
)
1056 hsa_op_with_type::get_in_type (BrigType16_t dtype
, hsa_bb
*hbb
)
1058 if (m_type
== dtype
)
1063 if (hsa_needs_cvt (dtype
, m_type
))
1065 dest
= new hsa_op_reg (dtype
);
1066 hbb
->append_insn (new hsa_insn_cvt (dest
, this));
1070 dest
= new hsa_op_reg (m_type
);
1071 hbb
->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
1072 dest
->m_type
, dest
, this));
1074 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
1075 type of the operand must be same as type of the instruction. */
1076 dest
->m_type
= dtype
;
1082 /* Constructor of class representing HSA immediate values. TREE_VAL is the
1083 tree representation of the immediate value. If min32int is true,
1084 always expand integer types to one that has at least 32 bits. */
1086 hsa_op_immed::hsa_op_immed (tree tree_val
, bool min32int
)
1087 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
,
1088 hsa_type_for_tree_type (TREE_TYPE (tree_val
), NULL
,
1091 if (hsa_seen_error ())
1094 gcc_checking_assert ((is_gimple_min_invariant (tree_val
)
1095 && (!POINTER_TYPE_P (TREE_TYPE (tree_val
))
1096 || TREE_CODE (tree_val
) == INTEGER_CST
))
1097 || TREE_CODE (tree_val
) == CONSTRUCTOR
);
1098 m_tree_value
= tree_val
;
1100 /* Verify that all elements of a constructor are constants. */
1101 if (TREE_CODE (m_tree_value
) == CONSTRUCTOR
)
1102 for (unsigned i
= 0;
1103 i
< vec_safe_length (CONSTRUCTOR_ELTS (m_tree_value
)); i
++)
1105 tree v
= CONSTRUCTOR_ELT (m_tree_value
, i
)->value
;
1106 if (!CONSTANT_CLASS_P (v
))
1108 HSA_SORRY_AT (EXPR_LOCATION (tree_val
),
1109 "HSA ctor should have only constants");
1115 /* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
1116 integer representation of the immediate value. TYPE is BRIG type. */
1118 hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value
, BrigType16_t type
)
1119 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
, type
),
1122 gcc_assert (hsa_type_integer_p (type
));
1123 m_int_value
= integer_value
;
1126 hsa_op_immed::hsa_op_immed ()
1127 : hsa_op_with_type (BRIG_KIND_NONE
, BRIG_TYPE_NONE
)
1131 /* New operator to allocate immediate operands from pool alloc. */
1134 hsa_op_immed::operator new (size_t)
1136 return hsa_allocp_operand_immed
->allocate_raw ();
1141 hsa_op_immed::~hsa_op_immed ()
1145 /* Change type of the immediate value to T. */
1148 hsa_op_immed::set_type (BrigType16_t t
)
1153 /* Constructor of class representing HSA registers and pseudo-registers. T is
1154 the BRIG type of the new register. */
1156 hsa_op_reg::hsa_op_reg (BrigType16_t t
)
1157 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER
, t
), m_gimple_ssa (NULL_TREE
),
1158 m_def_insn (NULL
), m_spill_sym (NULL
), m_order (hsa_cfun
->m_reg_count
++),
1159 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
1163 /* New operator to allocate a register from pool alloc. */
1166 hsa_op_reg::operator new (size_t)
1168 return hsa_allocp_operand_reg
->allocate_raw ();
1171 /* Verify register operand. */
1174 hsa_op_reg::verify_ssa ()
1176 /* Verify that each HSA register has a definition assigned.
1177 Exceptions are VAR_DECL and PARM_DECL that are a default
1179 gcc_checking_assert (m_def_insn
1180 || (m_gimple_ssa
!= NULL
1181 && (!SSA_NAME_VAR (m_gimple_ssa
)
1182 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa
))
1184 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa
)));
1186 /* Verify that every use of the register is really present
1187 in an instruction. */
1188 for (unsigned i
= 0; i
< m_uses
.length (); i
++)
1190 hsa_insn_basic
*use
= m_uses
[i
];
1192 bool is_visited
= false;
1193 for (unsigned j
= 0; j
< use
->operand_count (); j
++)
1195 hsa_op_base
*u
= use
->get_op (j
);
1196 hsa_op_address
*addr
; addr
= dyn_cast
<hsa_op_address
*> (u
);
1197 if (addr
&& addr
->m_reg
)
1202 bool r
= !addr
&& use
->op_output_p (j
);
1206 error ("HSA SSA name defined by instruction that is supposed "
1208 debug_hsa_operand (this);
1209 debug_hsa_insn (use
);
1210 internal_error ("HSA SSA verification failed");
1219 error ("HSA SSA name not among operands of instruction that is "
1220 "supposed to use it");
1221 debug_hsa_operand (this);
1222 debug_hsa_insn (use
);
1223 internal_error ("HSA SSA verification failed");
1228 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, hsa_op_reg
*r
,
1229 HOST_WIDE_INT offset
)
1230 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (r
),
1231 m_imm_offset (offset
)
1235 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, HOST_WIDE_INT offset
)
1236 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (NULL
),
1237 m_imm_offset (offset
)
1241 hsa_op_address::hsa_op_address (hsa_op_reg
*r
, HOST_WIDE_INT offset
)
1242 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (NULL
), m_reg (r
),
1243 m_imm_offset (offset
)
1247 /* New operator to allocate address operands from pool alloc. */
1250 hsa_op_address::operator new (size_t)
1252 return hsa_allocp_operand_address
->allocate_raw ();
1255 /* Constructor of an operand referring to HSAIL code. */
1257 hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF
),
1258 m_directive_offset (0)
1262 /* Constructor of an operand representing a code list. Set it up so that it
1263 can contain ELEMENTS number of elements. */
1265 hsa_op_code_list::hsa_op_code_list (unsigned elements
)
1266 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST
)
1268 m_offsets
.create (1);
1269 m_offsets
.safe_grow_cleared (elements
);
1272 /* New operator to allocate code list operands from pool alloc. */
1275 hsa_op_code_list::operator new (size_t)
1277 return hsa_allocp_operand_code_list
->allocate_raw ();
1280 /* Constructor of an operand representing an operand list.
1281 Set it up so that it can contain ELEMENTS number of elements. */
1283 hsa_op_operand_list::hsa_op_operand_list (unsigned elements
)
1284 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST
)
1286 m_offsets
.create (elements
);
1287 m_offsets
.safe_grow (elements
);
1290 /* New operator to allocate operand list operands from pool alloc. */
1293 hsa_op_operand_list::operator new (size_t)
1295 return hsa_allocp_operand_operand_list
->allocate_raw ();
1298 hsa_op_operand_list::~hsa_op_operand_list ()
1300 m_offsets
.release ();
1305 hsa_function_representation::reg_for_gimple_ssa (tree ssa
)
1309 gcc_checking_assert (TREE_CODE (ssa
) == SSA_NAME
);
1310 if (m_ssa_map
[SSA_NAME_VERSION (ssa
)])
1311 return m_ssa_map
[SSA_NAME_VERSION (ssa
)];
1313 hreg
= new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa
),
1315 hreg
->m_gimple_ssa
= ssa
;
1316 m_ssa_map
[SSA_NAME_VERSION (ssa
)] = hreg
;
1322 hsa_op_reg::set_definition (hsa_insn_basic
*insn
)
1324 if (hsa_cfun
->m_in_ssa
)
1326 gcc_checking_assert (!m_def_insn
);
1333 /* Constructor of the class which is the bases of all instructions and directly
1334 represents the most basic ones. NOPS is the number of operands that the
1335 operand vector will contain (and which will be cleared). OP is the opcode
1336 of the instruction. This constructor does not set type. */
1338 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
)
1340 m_next (NULL
), m_bb (NULL
), m_opcode (opc
), m_number (0),
1341 m_type (BRIG_TYPE_NONE
), m_brig_offset (0)
1344 m_operands
.safe_grow_cleared (nops
);
1346 hsa_instructions
.safe_push (this);
1349 /* Make OP the operand number INDEX of operands of this instruction. If OP is a
1350 register or an address containing a register, then either set the definition
1351 of the register to this instruction if it an output operand or add this
1352 instruction to the uses if it is an input one. */
1355 hsa_insn_basic::set_op (int index
, hsa_op_base
*op
)
1357 /* Each address operand is always use. */
1358 hsa_op_address
*addr
= dyn_cast
<hsa_op_address
*> (op
);
1359 if (addr
&& addr
->m_reg
)
1360 addr
->m_reg
->m_uses
.safe_push (this);
1363 hsa_op_reg
*reg
= dyn_cast
<hsa_op_reg
*> (op
);
1366 if (op_output_p (index
))
1367 reg
->set_definition (this);
1369 reg
->m_uses
.safe_push (this);
1373 m_operands
[index
] = op
;
1376 /* Get INDEX-th operand of the instruction. */
1379 hsa_insn_basic::get_op (int index
)
1381 return m_operands
[index
];
1384 /* Get address of INDEX-th operand of the instruction. */
1387 hsa_insn_basic::get_op_addr (int index
)
1389 return &m_operands
[index
];
1392 /* Get number of operands of the instruction. */
1394 hsa_insn_basic::operand_count ()
1396 return m_operands
.length ();
1399 /* Constructor of the class which is the bases of all instructions and directly
1400 represents the most basic ones. NOPS is the number of operands that the
1401 operand vector will contain (and which will be cleared). OPC is the opcode
1402 of the instruction, T is the type of the instruction. */
1404 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
, BrigType16_t t
,
1405 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1406 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1407 : m_prev (NULL
), m_next (NULL
), m_bb (NULL
), m_opcode (opc
),m_number (0),
1408 m_type (t
), m_brig_offset (0)
1411 m_operands
.safe_grow_cleared (nops
);
1415 gcc_checking_assert (nops
>= 1);
1421 gcc_checking_assert (nops
>= 2);
1427 gcc_checking_assert (nops
>= 3);
1433 gcc_checking_assert (nops
>= 4);
1437 hsa_instructions
.safe_push (this);
1440 /* New operator to allocate basic instruction from pool alloc. */
1443 hsa_insn_basic::operator new (size_t)
1445 return hsa_allocp_inst_basic
->allocate_raw ();
1448 /* Verify the instruction. */
1451 hsa_insn_basic::verify ()
1453 hsa_op_address
*addr
;
1456 /* Iterate all register operands and verify that the instruction
1457 is set in uses of the register. */
1458 for (unsigned i
= 0; i
< operand_count (); i
++)
1460 hsa_op_base
*use
= get_op (i
);
1462 if ((addr
= dyn_cast
<hsa_op_address
*> (use
)) && addr
->m_reg
)
1464 gcc_assert (addr
->m_reg
->m_def_insn
!= this);
1468 if ((reg
= dyn_cast
<hsa_op_reg
*> (use
)) && !op_output_p (i
))
1471 for (j
= 0; j
< reg
->m_uses
.length (); j
++)
1473 if (reg
->m_uses
[j
] == this)
1477 if (j
== reg
->m_uses
.length ())
1479 error ("HSA instruction uses a register but is not among "
1480 "recorded register uses");
1481 debug_hsa_operand (reg
);
1482 debug_hsa_insn (this);
1483 internal_error ("HSA instruction verification failed");
1489 /* Constructor of an instruction representing a PHI node. NOPS is the number
1490 of operands (equal to the number of predecessors). */
1492 hsa_insn_phi::hsa_insn_phi (unsigned nops
, hsa_op_reg
*dst
)
1493 : hsa_insn_basic (nops
, HSA_OPCODE_PHI
), m_dest (dst
)
1495 dst
->set_definition (this);
1498 /* New operator to allocate PHI instruction from pool alloc. */
1501 hsa_insn_phi::operator new (size_t)
1503 return hsa_allocp_inst_phi
->allocate_raw ();
1506 /* Constructor of class representing instruction for conditional jump, CTRL is
1507 the control register determining whether the jump will be carried out, the
1508 new instruction is automatically added to its uses list. */
1510 hsa_insn_br::hsa_insn_br (hsa_op_reg
*ctrl
)
1511 : hsa_insn_basic (1, BRIG_OPCODE_CBR
, BRIG_TYPE_B1
, ctrl
),
1512 m_width (BRIG_WIDTH_1
)
1516 /* New operator to allocate branch instruction from pool alloc. */
1519 hsa_insn_br::operator new (size_t)
1521 return hsa_allocp_inst_br
->allocate_raw ();
1524 /* Constructor of class representing instruction for switch jump, CTRL is
1525 the index register. */
1527 hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg
*index
, unsigned jump_count
)
1528 : hsa_insn_basic (1, BRIG_OPCODE_SBR
, BRIG_TYPE_B1
, index
),
1529 m_width (BRIG_WIDTH_1
), m_jump_table (vNULL
),
1530 m_label_code_list (new hsa_op_code_list (jump_count
))
1534 /* New operator to allocate switch branch instruction from pool alloc. */
1537 hsa_insn_sbr::operator new (size_t)
1539 return hsa_allocp_inst_sbr
->allocate_raw ();
1542 /* Replace all occurrences of OLD_BB with NEW_BB in the statements
1546 hsa_insn_sbr::replace_all_labels (basic_block old_bb
, basic_block new_bb
)
1548 for (unsigned i
= 0; i
< m_jump_table
.length (); i
++)
1549 if (m_jump_table
[i
] == old_bb
)
1550 m_jump_table
[i
] = new_bb
;
1553 hsa_insn_sbr::~hsa_insn_sbr ()
1555 m_jump_table
.release ();
1558 /* Constructor of comparison instruction. CMP is the comparison operation and T
1559 is the result type. */
1561 hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp
, BrigType16_t t
,
1562 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1564 : hsa_insn_basic (3 , BRIG_OPCODE_CMP
, t
, arg0
, arg1
, arg2
), m_compare (cmp
)
1568 /* New operator to allocate compare instruction from pool alloc. */
1571 hsa_insn_cmp::operator new (size_t)
1573 return hsa_allocp_inst_cmp
->allocate_raw ();
1576 /* Constructor of classes representing memory accesses. OPC is the opcode (must
1577 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
1578 operands are provided as ARG0 and ARG1. */
1580 hsa_insn_mem::hsa_insn_mem (int opc
, BrigType16_t t
, hsa_op_base
*arg0
,
1582 : hsa_insn_basic (2, opc
, t
, arg0
, arg1
),
1583 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1585 gcc_checking_assert (opc
== BRIG_OPCODE_LD
|| opc
== BRIG_OPCODE_ST
);
1588 /* Constructor for descendants allowing different opcodes and number of
1589 operands, it passes its arguments directly to hsa_insn_basic
1590 constructor. The instruction operands are provided as ARG[0-3]. */
1593 hsa_insn_mem::hsa_insn_mem (unsigned nops
, int opc
, BrigType16_t t
,
1594 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1595 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1596 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1597 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1601 /* New operator to allocate memory instruction from pool alloc. */
1604 hsa_insn_mem::operator new (size_t)
1606 return hsa_allocp_inst_mem
->allocate_raw ();
1609 /* Constructor of class representing atomic instructions and signals. OPC is
1610 the principal opcode, aop is the specific atomic operation opcode. T is the
1611 type of the instruction. The instruction operands
1612 are provided as ARG[0-3]. */
1614 hsa_insn_atomic::hsa_insn_atomic (int nops
, int opc
,
1615 enum BrigAtomicOperation aop
,
1616 BrigType16_t t
, BrigMemoryOrder memorder
,
1618 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1620 : hsa_insn_mem (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
), m_atomicop (aop
),
1621 m_memoryorder (memorder
),
1622 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM
)
1624 gcc_checking_assert (opc
== BRIG_OPCODE_ATOMICNORET
||
1625 opc
== BRIG_OPCODE_ATOMIC
||
1626 opc
== BRIG_OPCODE_SIGNAL
||
1627 opc
== BRIG_OPCODE_SIGNALNORET
);
1630 /* New operator to allocate signal instruction from pool alloc. */
1633 hsa_insn_atomic::operator new (size_t)
1635 return hsa_allocp_inst_atomic
->allocate_raw ();
1638 /* Constructor of class representing signal instructions. OPC is the prinicpal
1639 opcode, sop is the specific signal operation opcode. T is the type of the
1640 instruction. The instruction operands are provided as ARG[0-3]. */
1642 hsa_insn_signal::hsa_insn_signal (int nops
, int opc
,
1643 enum BrigAtomicOperation sop
,
1644 BrigType16_t t
, hsa_op_base
*arg0
,
1645 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1647 : hsa_insn_atomic (nops
, opc
, sop
, t
, BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
,
1648 arg0
, arg1
, arg2
, arg3
)
1652 /* New operator to allocate signal instruction from pool alloc. */
1655 hsa_insn_signal::operator new (size_t)
1657 return hsa_allocp_inst_signal
->allocate_raw ();
1660 /* Constructor of class representing segment conversion instructions. OPC is
1661 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
1662 and SRCT are destination and source types respectively, SEG is the segment
1663 we are converting to or from. The instruction operands are
1664 provided as ARG0 and ARG1. */
1666 hsa_insn_seg::hsa_insn_seg (int opc
, BrigType16_t dest
, BrigType16_t srct
,
1667 BrigSegment8_t seg
, hsa_op_base
*arg0
,
1669 : hsa_insn_basic (2, opc
, dest
, arg0
, arg1
), m_src_type (srct
),
1672 gcc_checking_assert (opc
== BRIG_OPCODE_STOF
|| opc
== BRIG_OPCODE_FTOS
);
1675 /* New operator to allocate address conversion instruction from pool alloc. */
1678 hsa_insn_seg::operator new (size_t)
1680 return hsa_allocp_inst_seg
->allocate_raw ();
1683 /* Constructor of class representing a call instruction. CALLEE is the tree
1684 representation of the function being called. */
1686 hsa_insn_call::hsa_insn_call (tree callee
)
1687 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (callee
),
1688 m_output_arg (NULL
), m_args_code_list (NULL
), m_result_code_list (NULL
)
1692 hsa_insn_call::hsa_insn_call (hsa_internal_fn
*fn
)
1693 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (NULL
),
1694 m_called_internal_fn (fn
), m_output_arg (NULL
), m_args_code_list (NULL
),
1695 m_result_code_list (NULL
)
1699 /* New operator to allocate call instruction from pool alloc. */
1702 hsa_insn_call::operator new (size_t)
1704 return hsa_allocp_inst_call
->allocate_raw ();
1707 hsa_insn_call::~hsa_insn_call ()
1709 for (unsigned i
= 0; i
< m_input_args
.length (); i
++)
1710 delete m_input_args
[i
];
1712 delete m_output_arg
;
1714 m_input_args
.release ();
1715 m_input_arg_insns
.release ();
1718 /* Constructor of class representing the argument block required to invoke
1720 hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind
,
1721 hsa_insn_call
* call
)
1722 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK
), m_kind (brig_kind
),
1727 /* New operator to allocate argument block instruction from pool alloc. */
1730 hsa_insn_arg_block::operator new (size_t)
1732 return hsa_allocp_inst_arg_block
->allocate_raw ();
1735 hsa_insn_comment::hsa_insn_comment (const char *s
)
1736 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT
)
1738 unsigned l
= strlen (s
);
1740 /* Append '// ' to the string. */
1741 char *buf
= XNEWVEC (char, l
+ 4);
1742 sprintf (buf
, "// %s", s
);
1746 /* New operator to allocate comment instruction from pool alloc. */
1749 hsa_insn_comment::operator new (size_t)
1751 return hsa_allocp_inst_comment
->allocate_raw ();
1754 hsa_insn_comment::~hsa_insn_comment ()
1756 gcc_checking_assert (m_comment
);
1761 /* Constructor of class representing the queue instruction in HSAIL. */
1762 hsa_insn_queue::hsa_insn_queue (int nops
, BrigOpcode opcode
)
1763 : hsa_insn_basic (nops
, opcode
, BRIG_TYPE_U64
)
1767 /* New operator to allocate source type instruction from pool alloc. */
1770 hsa_insn_srctype::operator new (size_t)
1772 return hsa_allocp_inst_srctype
->allocate_raw ();
1775 /* Constructor of class representing the source type instruction in HSAIL. */
1777 hsa_insn_srctype::hsa_insn_srctype (int nops
, BrigOpcode opcode
,
1778 BrigType16_t destt
, BrigType16_t srct
,
1779 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1780 hsa_op_base
*arg2
= NULL
)
1781 : hsa_insn_basic (nops
, opcode
, destt
, arg0
, arg1
, arg2
),
1782 m_source_type (srct
)
1785 /* New operator to allocate packed instruction from pool alloc. */
1788 hsa_insn_packed::operator new (size_t)
1790 return hsa_allocp_inst_packed
->allocate_raw ();
1793 /* Constructor of class representing the packed instruction in HSAIL. */
1795 hsa_insn_packed::hsa_insn_packed (int nops
, BrigOpcode opcode
,
1796 BrigType16_t destt
, BrigType16_t srct
,
1797 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1799 : hsa_insn_srctype (nops
, opcode
, destt
, srct
, arg0
, arg1
, arg2
)
1801 m_operand_list
= new hsa_op_operand_list (nops
- 1);
1804 /* New operator to allocate convert instruction from pool alloc. */
1807 hsa_insn_cvt::operator new (size_t)
1809 return hsa_allocp_inst_cvt
->allocate_raw ();
1812 /* Constructor of class representing the convert instruction in HSAIL. */
1814 hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type
*dest
, hsa_op_with_type
*src
)
1815 : hsa_insn_basic (2, BRIG_OPCODE_CVT
, dest
->m_type
, dest
, src
)
1819 /* New operator to allocate alloca from pool alloc. */
1822 hsa_insn_alloca::operator new (size_t)
1824 return hsa_allocp_inst_alloca
->allocate_raw ();
1827 /* Constructor of class representing the alloca in HSAIL. */
1829 hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type
*dest
,
1830 hsa_op_with_type
*size
, unsigned alignment
)
1831 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA
, dest
->m_type
, dest
, size
),
1832 m_align (BRIG_ALIGNMENT_8
)
1834 gcc_assert (dest
->m_type
== BRIG_TYPE_U32
);
1836 m_align
= hsa_alignment_encoding (alignment
);
1839 /* Append an instruction INSN into the basic block. */
1842 hsa_bb::append_insn (hsa_insn_basic
*insn
)
1844 gcc_assert (insn
->m_opcode
!= 0 || insn
->operand_count () == 0);
1845 gcc_assert (!insn
->m_bb
);
1848 insn
->m_prev
= m_last_insn
;
1849 insn
->m_next
= NULL
;
1851 m_last_insn
->m_next
= insn
;
1854 m_first_insn
= insn
;
1857 /* Insert HSA instruction NEW_INSN immediately before an existing instruction
1861 hsa_insert_insn_before (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1863 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1865 if (hbb
->m_first_insn
== old_insn
)
1866 hbb
->m_first_insn
= new_insn
;
1867 new_insn
->m_prev
= old_insn
->m_prev
;
1868 new_insn
->m_next
= old_insn
;
1869 if (old_insn
->m_prev
)
1870 old_insn
->m_prev
->m_next
= new_insn
;
1871 old_insn
->m_prev
= new_insn
;
1874 /* Append HSA instruction NEW_INSN immediately after an existing instruction
1878 hsa_append_insn_after (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1880 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1882 if (hbb
->m_last_insn
== old_insn
)
1883 hbb
->m_last_insn
= new_insn
;
1884 new_insn
->m_prev
= old_insn
;
1885 new_insn
->m_next
= old_insn
->m_next
;
1886 if (old_insn
->m_next
)
1887 old_insn
->m_next
->m_prev
= new_insn
;
1888 old_insn
->m_next
= new_insn
;
1891 /* Return a register containing the calculated value of EXP which must be an
1892 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
1893 integer constants as returned by get_inner_reference.
1894 Newly generated HSA instructions will be appended to HBB.
1895 Perform all calculations in ADDRTYPE. */
1897 static hsa_op_with_type
*
1898 gen_address_calculation (tree exp
, hsa_bb
*hbb
, BrigType16_t addrtype
)
1902 if (TREE_CODE (exp
) == NOP_EXPR
)
1903 exp
= TREE_OPERAND (exp
, 0);
1905 switch (TREE_CODE (exp
))
1908 return hsa_cfun
->reg_for_gimple_ssa (exp
)->get_in_type (addrtype
, hbb
);
1912 hsa_op_immed
*imm
= new hsa_op_immed (exp
);
1913 if (addrtype
!= imm
->m_type
)
1914 imm
->m_type
= addrtype
;
1919 opcode
= BRIG_OPCODE_ADD
;
1923 opcode
= BRIG_OPCODE_MUL
;
1930 hsa_op_reg
*res
= new hsa_op_reg (addrtype
);
1931 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, addrtype
);
1932 insn
->set_op (0, res
);
1934 hsa_op_with_type
*op1
= gen_address_calculation (TREE_OPERAND (exp
, 0), hbb
,
1936 hsa_op_with_type
*op2
= gen_address_calculation (TREE_OPERAND (exp
, 1), hbb
,
1938 insn
->set_op (1, op1
);
1939 insn
->set_op (2, op2
);
1941 hbb
->append_insn (insn
);
1945 /* If R1 is NULL, just return R2, otherwise append an instruction adding them
1946 to HBB and return the register holding the result. */
1949 add_addr_regs_if_needed (hsa_op_reg
*r1
, hsa_op_reg
*r2
, hsa_bb
*hbb
)
1951 gcc_checking_assert (r2
);
1955 hsa_op_reg
*res
= new hsa_op_reg (r1
->m_type
);
1956 gcc_assert (!hsa_needs_cvt (r1
->m_type
, r2
->m_type
));
1957 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_ADD
, res
->m_type
);
1958 insn
->set_op (0, res
);
1959 insn
->set_op (1, r1
);
1960 insn
->set_op (2, r2
);
1961 hbb
->append_insn (insn
);
1965 /* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
1966 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
1969 process_mem_base (tree base
, hsa_symbol
**symbol
, BrigType16_t
*addrtype
,
1970 hsa_op_reg
**reg
, offset_int
*offset
, hsa_bb
*hbb
)
1972 if (TREE_CODE (base
) == SSA_NAME
)
1975 hsa_op_with_type
*ssa
1976 = hsa_cfun
->reg_for_gimple_ssa (base
)->get_in_type (*addrtype
, hbb
);
1977 *reg
= dyn_cast
<hsa_op_reg
*> (ssa
);
1979 else if (TREE_CODE (base
) == ADDR_EXPR
)
1981 tree decl
= TREE_OPERAND (base
, 0);
1983 if (!DECL_P (decl
) || TREE_CODE (decl
) == FUNCTION_DECL
)
1985 HSA_SORRY_AT (EXPR_LOCATION (base
),
1986 "support for HSA does not implement a memory reference "
1987 "to a non-declaration type");
1991 gcc_assert (!*symbol
);
1993 *symbol
= get_symbol_for_decl (decl
);
1994 *addrtype
= hsa_get_segment_addr_type ((*symbol
)->m_segment
);
1996 else if (TREE_CODE (base
) == INTEGER_CST
)
1997 *offset
+= wi::to_offset (base
);
2002 /* Forward declaration of a function. */
2005 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
);
2007 /* Generate HSA address operand for a given tree memory reference REF. If
2008 instructions need to be created to calculate the address, they will be added
2009 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
2010 the function assumes that the caller will handle possible
2011 bit-field references. Otherwise if we reference a bit-field, sorry message
2014 static hsa_op_address
*
2015 gen_hsa_addr (tree ref
, hsa_bb
*hbb
, HOST_WIDE_INT
*output_bitsize
= NULL
,
2016 HOST_WIDE_INT
*output_bitpos
= NULL
)
2018 hsa_symbol
*symbol
= NULL
;
2019 hsa_op_reg
*reg
= NULL
;
2020 offset_int offset
= 0;
2022 tree varoffset
= NULL_TREE
;
2023 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2024 HOST_WIDE_INT bitsize
= 0, bitpos
= 0;
2025 BrigType16_t flat_addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2027 if (TREE_CODE (ref
) == STRING_CST
)
2029 symbol
= hsa_get_string_cst_symbol (ref
);
2032 else if (TREE_CODE (ref
) == BIT_FIELD_REF
2033 && ((tree_to_uhwi (TREE_OPERAND (ref
, 1)) % BITS_PER_UNIT
) != 0
2034 || (tree_to_uhwi (TREE_OPERAND (ref
, 2)) % BITS_PER_UNIT
) != 0))
2036 HSA_SORRY_ATV (EXPR_LOCATION (origref
),
2037 "support for HSA does not implement "
2038 "bit field references such as %E", ref
);
2042 if (handled_component_p (ref
))
2044 enum machine_mode mode
;
2045 int unsignedp
, volatilep
, preversep
;
2047 ref
= get_inner_reference (ref
, &bitsize
, &bitpos
, &varoffset
, &mode
,
2048 &unsignedp
, &preversep
, &volatilep
, false);
2051 offset
= wi::rshift (offset
, LOG2_BITS_PER_UNIT
, SIGNED
);
2054 switch (TREE_CODE (ref
))
2058 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
2059 symbol
= hsa_cfun
->create_hsa_temporary (flat_addrtype
);
2060 hsa_op_reg
*r
= new hsa_op_reg (flat_addrtype
);
2061 gen_hsa_addr_insns (ref
, r
, hbb
);
2062 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
2063 r
, new hsa_op_address (symbol
)));
2069 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
2070 symbol
= hsa_cfun
->create_hsa_temporary (flat_addrtype
);
2071 hsa_op_reg
*r
= hsa_cfun
->reg_for_gimple_ssa (ref
);
2073 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
2074 r
, new hsa_op_address (symbol
)));
2081 gcc_assert (!symbol
);
2082 symbol
= get_symbol_for_decl (ref
);
2083 addrtype
= hsa_get_segment_addr_type (symbol
->m_segment
);
2087 process_mem_base (TREE_OPERAND (ref
, 0), &symbol
, &addrtype
, ®
,
2090 if (!integer_zerop (TREE_OPERAND (ref
, 1)))
2091 offset
+= wi::to_offset (TREE_OPERAND (ref
, 1));
2094 case TARGET_MEM_REF
:
2095 process_mem_base (TMR_BASE (ref
), &symbol
, &addrtype
, ®
, &offset
, hbb
);
2096 if (TMR_INDEX (ref
))
2099 hsa_op_base
*idx
= hsa_cfun
->reg_for_gimple_ssa
2100 (TMR_INDEX (ref
))->get_in_type (addrtype
, hbb
);
2101 if (TMR_STEP (ref
) && !integer_onep (TMR_STEP (ref
)))
2103 disp1
= new hsa_op_reg (addrtype
);
2104 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_MUL
,
2107 /* As step must respect addrtype, we overwrite the type
2108 of an immediate value. */
2109 hsa_op_immed
*step
= new hsa_op_immed (TMR_STEP (ref
));
2110 step
->m_type
= addrtype
;
2112 insn
->set_op (0, disp1
);
2113 insn
->set_op (1, idx
);
2114 insn
->set_op (2, step
);
2115 hbb
->append_insn (insn
);
2118 disp1
= as_a
<hsa_op_reg
*> (idx
);
2119 reg
= add_addr_regs_if_needed (reg
, disp1
, hbb
);
2121 if (TMR_INDEX2 (ref
))
2123 if (TREE_CODE (TMR_INDEX2 (ref
)) == SSA_NAME
)
2125 hsa_op_base
*disp2
= hsa_cfun
->reg_for_gimple_ssa
2126 (TMR_INDEX2 (ref
))->get_in_type (addrtype
, hbb
);
2127 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (disp2
),
2130 else if (TREE_CODE (TMR_INDEX2 (ref
)) == INTEGER_CST
)
2131 offset
+= wi::to_offset (TMR_INDEX2 (ref
));
2135 offset
+= wi::to_offset (TMR_OFFSET (ref
));
2138 HSA_SORRY_AT (EXPR_LOCATION (origref
),
2139 "support for HSA does not implement function pointers");
2142 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does "
2143 "not implement memory access to %E", origref
);
2149 if (TREE_CODE (varoffset
) == INTEGER_CST
)
2150 offset
+= wi::to_offset (varoffset
);
2153 hsa_op_base
*off_op
= gen_address_calculation (varoffset
, hbb
,
2155 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (off_op
),
2160 gcc_checking_assert ((symbol
2162 == hsa_get_segment_addr_type (symbol
->m_segment
))
2165 == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
)));
2167 HOST_WIDE_INT hwi_offset
= offset
.to_shwi ();
2169 /* Calculate remaining bitsize offset (if presented). */
2170 bitpos
%= BITS_PER_UNIT
;
2171 /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
2172 is not a reason to think this is a bit-field access. */
2174 && (bitsize
>= BITS_PER_UNIT
)
2175 && !(bitsize
& (bitsize
- 1)))
2178 if ((bitpos
|| bitsize
) && (output_bitpos
== NULL
|| output_bitsize
== NULL
))
2179 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does not "
2180 "implement unhandled bit field reference such as %E", ref
);
2182 if (output_bitsize
!= NULL
&& output_bitpos
!= NULL
)
2184 *output_bitsize
= bitsize
;
2185 *output_bitpos
= bitpos
;
2188 return new hsa_op_address (symbol
, reg
, hwi_offset
);
2191 /* Generate HSA address operand for a given tree memory reference REF. If
2192 instructions need to be created to calculate the address, they will be added
2193 to the end of HBB. OUTPUT_ALIGN is alignment of the created address. */
2195 static hsa_op_address
*
2196 gen_hsa_addr_with_align (tree ref
, hsa_bb
*hbb
, BrigAlignment8_t
*output_align
)
2198 hsa_op_address
*addr
= gen_hsa_addr (ref
, hbb
);
2199 if (addr
->m_reg
|| !addr
->m_symbol
)
2200 *output_align
= hsa_object_alignment (ref
);
2203 /* If the address consists only of a symbol and an offset, we
2204 compute the alignment ourselves to take into account any alignment
2205 promotions we might have done for the HSA symbol representation. */
2206 unsigned align
= hsa_byte_alignment (addr
->m_symbol
->m_align
);
2207 unsigned misalign
= addr
->m_imm_offset
& (align
- 1);
2209 align
= (misalign
& -misalign
);
2210 *output_align
= hsa_alignment_encoding (BITS_PER_UNIT
* align
);
2215 /* Generate HSA address for a function call argument of given TYPE.
2216 INDEX is used to generate corresponding name of the arguments.
2217 Special value -1 represents fact that result value is created. */
2219 static hsa_op_address
*
2220 gen_hsa_addr_for_arg (tree tree_type
, int index
)
2222 hsa_symbol
*sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
2224 sym
->m_type
= hsa_type_for_tree_type (tree_type
, &sym
->m_dim
);
2226 if (index
== -1) /* Function result. */
2227 sym
->m_name
= "res";
2228 else /* Function call arguments. */
2231 sym
->m_name_number
= index
;
2234 return new hsa_op_address (sym
);
2237 /* Generate HSA instructions that process all necessary conversions
2238 of an ADDR to flat addressing and place the result into DEST.
2239 Instructions are appended to HBB. */
2242 convert_addr_to_flat_segment (hsa_op_address
*addr
, hsa_op_reg
*dest
,
2245 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_LDA
);
2246 insn
->set_op (1, addr
);
2247 if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
!= BRIG_SEGMENT_GLOBAL
)
2249 /* LDA produces segment-relative address, we need to convert
2250 it to the flat one. */
2252 tmp
= new hsa_op_reg (hsa_get_segment_addr_type
2253 (addr
->m_symbol
->m_segment
));
2255 seg
= new hsa_insn_seg (BRIG_OPCODE_STOF
,
2256 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
2257 tmp
->m_type
, addr
->m_symbol
->m_segment
, dest
,
2260 insn
->set_op (0, tmp
);
2261 insn
->m_type
= tmp
->m_type
;
2262 hbb
->append_insn (insn
);
2263 hbb
->append_insn (seg
);
2267 insn
->set_op (0, dest
);
2268 insn
->m_type
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2269 hbb
->append_insn (insn
);
2273 /* Generate HSA instructions that calculate address of VAL including all
2274 necessary conversions to flat addressing and place the result into DEST.
2275 Instructions are appended to HBB. */
2278 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
)
2280 /* Handle cases like tmp = NULL, where we just emit a move instruction
2282 if (TREE_CODE (val
) == INTEGER_CST
)
2284 hsa_op_immed
*c
= new hsa_op_immed (val
);
2285 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2286 dest
->m_type
, dest
, c
);
2287 hbb
->append_insn (insn
);
2291 hsa_op_address
*addr
;
2293 gcc_assert (dest
->m_type
== hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2294 if (TREE_CODE (val
) == ADDR_EXPR
)
2295 val
= TREE_OPERAND (val
, 0);
2296 addr
= gen_hsa_addr (val
, hbb
);
2298 convert_addr_to_flat_segment (addr
, dest
, hbb
);
2301 /* Return an HSA register or HSA immediate value operand corresponding to
2302 gimple operand OP. */
2304 static hsa_op_with_type
*
2305 hsa_reg_or_immed_for_gimple_op (tree op
, hsa_bb
*hbb
)
2309 if (TREE_CODE (op
) == SSA_NAME
)
2310 tmp
= hsa_cfun
->reg_for_gimple_ssa (op
);
2311 else if (!POINTER_TYPE_P (TREE_TYPE (op
)))
2312 return new hsa_op_immed (op
);
2315 tmp
= new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2316 gen_hsa_addr_insns (op
, tmp
, hbb
);
2321 /* Create a simple movement instruction with register destination DEST and
2322 register or immediate source SRC and append it to the end of HBB. */
2325 hsa_build_append_simple_mov (hsa_op_reg
*dest
, hsa_op_base
*src
, hsa_bb
*hbb
)
2327 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
,
2329 if (hsa_op_reg
*sreg
= dyn_cast
<hsa_op_reg
*> (src
))
2330 gcc_assert (hsa_type_bit_size (dest
->m_type
)
2331 == hsa_type_bit_size (sreg
->m_type
));
2333 gcc_assert (hsa_type_bit_size (dest
->m_type
)
2334 == hsa_type_bit_size (as_a
<hsa_op_immed
*> (src
)->m_type
));
2336 hbb
->append_insn (insn
);
2339 /* Generate HSAIL instructions loading a bit field into register DEST.
2340 VALUE_REG is a register of a SSA name that is used in the bit field
2341 reference. To identify a bit field BITPOS is offset to the loaded memory
2342 and BITSIZE is number of bits of the bit field.
2343 Add instructions to HBB. */
2346 gen_hsa_insns_for_bitfield (hsa_op_reg
*dest
, hsa_op_reg
*value_reg
,
2347 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2350 unsigned type_bitsize
= hsa_type_bit_size (dest
->m_type
);
2351 unsigned left_shift
= type_bitsize
- (bitsize
+ bitpos
);
2352 unsigned right_shift
= left_shift
+ bitpos
;
2356 hsa_op_reg
*value_reg_2
= new hsa_op_reg (dest
->m_type
);
2357 hsa_op_immed
*c
= new hsa_op_immed (left_shift
, BRIG_TYPE_U32
);
2359 hsa_insn_basic
*lshift
2360 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, value_reg_2
->m_type
,
2361 value_reg_2
, value_reg
, c
);
2363 hbb
->append_insn (lshift
);
2365 value_reg
= value_reg_2
;
2370 hsa_op_reg
*value_reg_2
= new hsa_op_reg (dest
->m_type
);
2371 hsa_op_immed
*c
= new hsa_op_immed (right_shift
, BRIG_TYPE_U32
);
2373 hsa_insn_basic
*rshift
2374 = new hsa_insn_basic (3, BRIG_OPCODE_SHR
, value_reg_2
->m_type
,
2375 value_reg_2
, value_reg
, c
);
2377 hbb
->append_insn (rshift
);
2379 value_reg
= value_reg_2
;
2382 hsa_insn_basic
*assignment
2383 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, dest
, value_reg
);
2384 hbb
->append_insn (assignment
);
2388 /* Generate HSAIL instructions loading a bit field into register DEST. ADDR is
2389 prepared memory address which is used to load the bit field. To identify a
2390 bit field BITPOS is offset to the loaded memory and BITSIZE is number of
2391 bits of the bit field. Add instructions to HBB. Load must be performed in
2395 gen_hsa_insns_for_bitfield_load (hsa_op_reg
*dest
, hsa_op_address
*addr
,
2396 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2397 hsa_bb
*hbb
, BrigAlignment8_t align
)
2399 hsa_op_reg
*value_reg
= new hsa_op_reg (dest
->m_type
);
2400 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, dest
->m_type
, value_reg
,
2402 mem
->set_align (align
);
2403 hbb
->append_insn (mem
);
2404 gen_hsa_insns_for_bitfield (dest
, value_reg
, bitsize
, bitpos
, hbb
);
2407 /* Return the alignment of base memory accesses we issue to perform bit-field
2408 memory access REF. */
2410 static BrigAlignment8_t
2411 hsa_bitmemref_alignment (tree ref
)
2413 unsigned HOST_WIDE_INT bit_offset
= 0;
2417 if (TREE_CODE (ref
) == BIT_FIELD_REF
)
2419 if (!tree_fits_uhwi_p (TREE_OPERAND (ref
, 2)))
2420 return BRIG_ALIGNMENT_1
;
2421 bit_offset
+= tree_to_uhwi (TREE_OPERAND (ref
, 2));
2423 else if (TREE_CODE (ref
) == COMPONENT_REF
2424 && DECL_BIT_FIELD (TREE_OPERAND (ref
, 1)))
2425 bit_offset
+= int_bit_position (TREE_OPERAND (ref
, 1));
2428 ref
= TREE_OPERAND (ref
, 0);
2431 unsigned HOST_WIDE_INT bits
= bit_offset
% BITS_PER_UNIT
;
2432 unsigned HOST_WIDE_INT byte_bits
= bit_offset
- bits
;
2433 BrigAlignment8_t base
= hsa_object_alignment (ref
);
2436 return MIN (base
, hsa_alignment_encoding (byte_bits
& -byte_bits
));
2439 /* Generate HSAIL instructions loading something into register DEST. RHS is
2440 tree representation of the loaded data, which are loaded as type TYPE. Add
2441 instructions to HBB. */
2444 gen_hsa_insns_for_load (hsa_op_reg
*dest
, tree rhs
, tree type
, hsa_bb
*hbb
)
2446 /* The destination SSA name will give us the type. */
2447 if (TREE_CODE (rhs
) == VIEW_CONVERT_EXPR
)
2448 rhs
= TREE_OPERAND (rhs
, 0);
2450 if (TREE_CODE (rhs
) == SSA_NAME
)
2452 hsa_op_reg
*src
= hsa_cfun
->reg_for_gimple_ssa (rhs
);
2453 hsa_build_append_simple_mov (dest
, src
, hbb
);
2455 else if (is_gimple_min_invariant (rhs
)
2456 || TREE_CODE (rhs
) == ADDR_EXPR
)
2458 if (POINTER_TYPE_P (TREE_TYPE (rhs
)))
2460 if (dest
->m_type
!= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
))
2462 HSA_SORRY_ATV (EXPR_LOCATION (rhs
),
2463 "support for HSA does not implement conversion "
2464 "of %E to the requested non-pointer type.", rhs
);
2468 gen_hsa_addr_insns (rhs
, dest
, hbb
);
2470 else if (TREE_CODE (rhs
) == COMPLEX_CST
)
2472 hsa_op_immed
*real_part
= new hsa_op_immed (TREE_REALPART (rhs
));
2473 hsa_op_immed
*imag_part
= new hsa_op_immed (TREE_IMAGPART (rhs
));
2475 hsa_op_reg
*real_part_reg
2476 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2478 hsa_op_reg
*imag_part_reg
2479 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2482 hsa_build_append_simple_mov (real_part_reg
, real_part
, hbb
);
2483 hsa_build_append_simple_mov (imag_part_reg
, imag_part
, hbb
);
2485 BrigType16_t src_type
= hsa_bittype_for_type (real_part_reg
->m_type
);
2487 hsa_insn_packed
*insn
2488 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
,
2489 src_type
, dest
, real_part_reg
,
2491 hbb
->append_insn (insn
);
2495 hsa_op_immed
*imm
= new hsa_op_immed (rhs
);
2496 hsa_build_append_simple_mov (dest
, imm
, hbb
);
2499 else if (TREE_CODE (rhs
) == REALPART_EXPR
|| TREE_CODE (rhs
) == IMAGPART_EXPR
)
2501 tree pack_type
= TREE_TYPE (TREE_OPERAND (rhs
, 0));
2503 hsa_op_reg
*packed_reg
2504 = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type
, true));
2506 tree complex_rhs
= TREE_OPERAND (rhs
, 0);
2507 gen_hsa_insns_for_load (packed_reg
, complex_rhs
, TREE_TYPE (complex_rhs
),
2510 hsa_op_reg
*real_reg
2511 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2513 hsa_op_reg
*imag_reg
2514 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2516 BrigKind16_t brig_type
= packed_reg
->m_type
;
2517 hsa_insn_packed
*packed
2518 = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND
,
2519 hsa_bittype_for_type (real_reg
->m_type
),
2520 brig_type
, real_reg
, imag_reg
, packed_reg
);
2522 hbb
->append_insn (packed
);
2524 hsa_op_reg
*source
= TREE_CODE (rhs
) == REALPART_EXPR
?
2525 real_reg
: imag_reg
;
2527 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2528 dest
->m_type
, dest
, source
);
2530 hbb
->append_insn (insn
);
2532 else if (TREE_CODE (rhs
) == BIT_FIELD_REF
2533 && TREE_CODE (TREE_OPERAND (rhs
, 0)) == SSA_NAME
)
2535 tree ssa_name
= TREE_OPERAND (rhs
, 0);
2536 HOST_WIDE_INT bitsize
= tree_to_uhwi (TREE_OPERAND (rhs
, 1));
2537 HOST_WIDE_INT bitpos
= tree_to_uhwi (TREE_OPERAND (rhs
, 2));
2539 hsa_op_reg
*imm_value
= hsa_cfun
->reg_for_gimple_ssa (ssa_name
);
2540 gen_hsa_insns_for_bitfield (dest
, imm_value
, bitsize
, bitpos
, hbb
);
2542 else if (DECL_P (rhs
) || TREE_CODE (rhs
) == MEM_REF
2543 || TREE_CODE (rhs
) == TARGET_MEM_REF
2544 || handled_component_p (rhs
))
2546 HOST_WIDE_INT bitsize
, bitpos
;
2548 /* Load from memory. */
2549 hsa_op_address
*addr
;
2550 addr
= gen_hsa_addr (rhs
, hbb
, &bitsize
, &bitpos
);
2552 /* Handle load of a bit field. */
2555 HSA_SORRY_AT (EXPR_LOCATION (rhs
),
2556 "support for HSA does not implement load from a bit "
2557 "field bigger than 64 bits");
2561 if (bitsize
|| bitpos
)
2562 gen_hsa_insns_for_bitfield_load (dest
, addr
, bitsize
, bitpos
, hbb
,
2563 hsa_bitmemref_alignment (rhs
));
2567 /* Not dest->m_type, that's possibly extended. */
2568 mtype
= mem_type_for_type (hsa_type_for_scalar_tree_type (type
,
2570 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
, dest
,
2572 mem
->set_align (hsa_object_alignment (rhs
));
2573 hbb
->append_insn (mem
);
2577 HSA_SORRY_ATV (EXPR_LOCATION (rhs
),
2578 "support for HSA does not implement loading "
2583 /* Return number of bits necessary for representation of a bit field,
2584 starting at BITPOS with size of BITSIZE. */
2587 get_bitfield_size (unsigned bitpos
, unsigned bitsize
)
2589 unsigned s
= bitpos
+ bitsize
;
2590 unsigned sizes
[] = {8, 16, 32, 64};
2592 for (unsigned i
= 0; i
< 4; i
++)
2600 /* Generate HSAIL instructions storing into memory. LHS is the destination of
2601 the store, SRC is the source operand. Add instructions to HBB. */
2604 gen_hsa_insns_for_store (tree lhs
, hsa_op_base
*src
, hsa_bb
*hbb
)
2606 HOST_WIDE_INT bitsize
= 0, bitpos
= 0;
2607 BrigAlignment8_t req_align
;
2609 mtype
= mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
2611 hsa_op_address
*addr
;
2612 addr
= gen_hsa_addr (lhs
, hbb
, &bitsize
, &bitpos
);
2614 /* Handle store to a bit field. */
2617 HSA_SORRY_AT (EXPR_LOCATION (lhs
),
2618 "support for HSA does not implement store to a bit field "
2619 "bigger than 64 bits");
2623 unsigned type_bitsize
= get_bitfield_size (bitpos
, bitsize
);
2625 /* HSAIL does not support MOV insn with 16-bits integers. */
2626 if (type_bitsize
< 32)
2629 if (bitpos
|| (bitsize
&& type_bitsize
!= bitsize
))
2631 unsigned HOST_WIDE_INT mask
= 0;
2632 BrigType16_t mem_type
2633 = get_integer_type_by_bytes (type_bitsize
/ BITS_PER_UNIT
,
2634 !TYPE_UNSIGNED (TREE_TYPE (lhs
)));
2636 for (unsigned i
= 0; i
< type_bitsize
; i
++)
2637 if (i
< bitpos
|| i
>= bitpos
+ bitsize
)
2638 mask
|= ((unsigned HOST_WIDE_INT
)1 << i
);
2640 hsa_op_reg
*value_reg
= new hsa_op_reg (mem_type
);
2642 req_align
= hsa_bitmemref_alignment (lhs
);
2643 /* Load value from memory. */
2644 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mem_type
,
2646 mem
->set_align (req_align
);
2647 hbb
->append_insn (mem
);
2649 /* AND the loaded value with prepared mask. */
2650 hsa_op_reg
*cleared_reg
= new hsa_op_reg (mem_type
);
2653 = get_integer_type_by_bytes (type_bitsize
/ BITS_PER_UNIT
, false);
2654 hsa_op_immed
*c
= new hsa_op_immed (mask
, t
);
2656 hsa_insn_basic
*clearing
2657 = new hsa_insn_basic (3, BRIG_OPCODE_AND
, mem_type
, cleared_reg
,
2659 hbb
->append_insn (clearing
);
2661 /* Shift to left a value that is going to be stored. */
2662 hsa_op_reg
*new_value_reg
= new hsa_op_reg (mem_type
);
2664 hsa_insn_basic
*basic
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, mem_type
,
2665 new_value_reg
, src
);
2666 hbb
->append_insn (basic
);
2670 hsa_op_reg
*shifted_value_reg
= new hsa_op_reg (mem_type
);
2671 c
= new hsa_op_immed (bitpos
, BRIG_TYPE_U32
);
2673 hsa_insn_basic
*basic
2674 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, mem_type
,
2675 shifted_value_reg
, new_value_reg
, c
);
2676 hbb
->append_insn (basic
);
2678 new_value_reg
= shifted_value_reg
;
2681 /* OR the prepared value with prepared chunk loaded from memory. */
2682 hsa_op_reg
*prepared_reg
= new hsa_op_reg (mem_type
);
2683 basic
= new hsa_insn_basic (3, BRIG_OPCODE_OR
, mem_type
, prepared_reg
,
2684 new_value_reg
, cleared_reg
);
2685 hbb
->append_insn (basic
);
2691 req_align
= hsa_object_alignment (lhs
);
2693 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
, addr
);
2694 mem
->set_align (req_align
);
2696 /* The HSAIL verifier has another constraint: if the source is an immediate
2697 then it must match the destination type. If it's a register the low bits
2698 will be used for sub-word stores. We're always allocating new operands so
2699 we can modify the above in place. */
2700 if (hsa_op_immed
*imm
= dyn_cast
<hsa_op_immed
*> (src
))
2702 if (!hsa_type_packed_p (imm
->m_type
))
2703 imm
->m_type
= mem
->m_type
;
2706 /* ...and all vector immediates apparently need to be vectors of
2708 unsigned bs
= hsa_type_bit_size (imm
->m_type
);
2709 gcc_assert (bs
== hsa_type_bit_size (mem
->m_type
));
2713 imm
->m_type
= BRIG_TYPE_U8X4
;
2716 imm
->m_type
= BRIG_TYPE_U8X8
;
2719 imm
->m_type
= BRIG_TYPE_U8X16
;
2727 hbb
->append_insn (mem
);
2730 /* Generate memory copy instructions that are going to be used
2731 for copying a SRC memory to TARGET memory,
2732 represented by pointer in a register. MIN_ALIGN is minimal alignment
2733 of provided HSA addresses. */
2736 gen_hsa_memory_copy (hsa_bb
*hbb
, hsa_op_address
*target
, hsa_op_address
*src
,
2737 unsigned size
, BrigAlignment8_t min_align
)
2739 hsa_op_address
*addr
;
2742 unsigned offset
= 0;
2743 unsigned min_byte_align
= hsa_byte_alignment (min_align
);
2757 if (s
> min_byte_align
)
2760 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2762 hsa_op_reg
*tmp
= new hsa_op_reg (t
);
2763 addr
= new hsa_op_address (src
->m_symbol
, src
->m_reg
,
2764 src
->m_imm_offset
+ offset
);
2765 mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, t
, tmp
, addr
);
2766 hbb
->append_insn (mem
);
2768 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2769 target
->m_imm_offset
+ offset
);
2770 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, tmp
, addr
);
2771 hbb
->append_insn (mem
);
2777 /* Create a memset mask that is created by copying a CONSTANT byte value
2778 to an integer of BYTE_SIZE bytes. */
2780 static unsigned HOST_WIDE_INT
2781 build_memset_value (unsigned HOST_WIDE_INT constant
, unsigned byte_size
)
2786 HOST_WIDE_INT v
= constant
;
2788 for (unsigned i
= 1; i
< byte_size
; i
++)
2789 v
|= constant
<< (8 * i
);
2794 /* Generate memory set instructions that are going to be used
2795 for setting a CONSTANT byte value to TARGET memory of SIZE bytes.
2796 MIN_ALIGN is minimal alignment of provided HSA addresses. */
2799 gen_hsa_memory_set (hsa_bb
*hbb
, hsa_op_address
*target
,
2800 unsigned HOST_WIDE_INT constant
,
2801 unsigned size
, BrigAlignment8_t min_align
)
2803 hsa_op_address
*addr
;
2806 unsigned offset
= 0;
2807 unsigned min_byte_align
= hsa_byte_alignment (min_align
);
2821 if (s
> min_byte_align
)
2824 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2825 target
->m_imm_offset
+ offset
);
2827 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2828 HOST_WIDE_INT c
= build_memset_value (constant
, s
);
2830 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, new hsa_op_immed (c
, t
),
2832 hbb
->append_insn (mem
);
2838 /* Generate HSAIL instructions for a single assignment
2839 of an empty constructor to an ADDR_LHS. Constructor is passed as a
2840 tree RHS and all instructions are appended to HBB. ALIGN is
2841 alignment of the address. */
2844 gen_hsa_ctor_assignment (hsa_op_address
*addr_lhs
, tree rhs
, hsa_bb
*hbb
,
2845 BrigAlignment8_t align
)
2847 if (vec_safe_length (CONSTRUCTOR_ELTS (rhs
)))
2849 HSA_SORRY_AT (EXPR_LOCATION (rhs
),
2850 "support for HSA does not implement load from constructor");
2854 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2855 gen_hsa_memory_set (hbb
, addr_lhs
, 0, size
, align
);
2858 /* Generate HSA instructions for a single assignment of RHS to LHS.
2859 HBB is the basic block they will be appended to. */
2862 gen_hsa_insns_for_single_assignment (tree lhs
, tree rhs
, hsa_bb
*hbb
)
2864 if (TREE_CODE (lhs
) == SSA_NAME
)
2866 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
2867 if (hsa_seen_error ())
2870 gen_hsa_insns_for_load (dest
, rhs
, TREE_TYPE (lhs
), hbb
);
2872 else if (TREE_CODE (rhs
) == SSA_NAME
2873 || (is_gimple_min_invariant (rhs
) && TREE_CODE (rhs
) != STRING_CST
))
2875 /* Store to memory. */
2876 hsa_op_base
*src
= hsa_reg_or_immed_for_gimple_op (rhs
, hbb
);
2877 if (hsa_seen_error ())
2880 gen_hsa_insns_for_store (lhs
, src
, hbb
);
2884 BrigAlignment8_t lhs_align
;
2885 hsa_op_address
*addr_lhs
= gen_hsa_addr_with_align (lhs
, hbb
,
2888 if (TREE_CODE (rhs
) == CONSTRUCTOR
)
2889 gen_hsa_ctor_assignment (addr_lhs
, rhs
, hbb
, lhs_align
);
2892 BrigAlignment8_t rhs_align
;
2893 hsa_op_address
*addr_rhs
= gen_hsa_addr_with_align (rhs
, hbb
,
2896 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2897 gen_hsa_memory_copy (hbb
, addr_lhs
, addr_rhs
, size
,
2898 MIN (lhs_align
, rhs_align
));
2903 /* Prepend before INSN a load from spill symbol of SPILL_REG. Return the
2904 register into which we loaded. If this required another register to convert
2905 from a B1 type, return it in *PTMP2, otherwise store NULL into it. We
2906 assume we are out of SSA so the returned register does not have its
2910 hsa_spill_in (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2912 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2913 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2914 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2916 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, spill_sym
->m_type
,
2918 hsa_insert_insn_before (mem
, insn
);
2921 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2923 hsa_insn_basic
*cvtinsn
;
2925 reg
= new hsa_op_reg (spill_reg
->m_type
);
2927 cvtinsn
= new hsa_insn_cvt (reg
, *ptmp2
);
2928 hsa_insert_insn_before (cvtinsn
, insn
);
2933 /* Append after INSN a store to spill symbol of SPILL_REG. Return the register
2934 from which we stored. If this required another register to convert to a B1
2935 type, return it in *PTMP2, otherwise store NULL into it. We assume we are
2936 out of SSA so the returned register does not have its use updated. */
2939 hsa_spill_out (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2941 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2942 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2943 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2944 hsa_op_reg
*returnreg
;
2948 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2950 hsa_insn_basic
*cvtinsn
;
2951 *ptmp2
= new hsa_op_reg (spill_sym
->m_type
);
2952 reg
->m_type
= spill_reg
->m_type
;
2954 cvtinsn
= new hsa_insn_cvt (*ptmp2
, returnreg
);
2955 hsa_append_insn_after (cvtinsn
, insn
);
2960 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, spill_sym
->m_type
, reg
,
2962 hsa_append_insn_after (mem
, insn
);
2966 /* Generate a comparison instruction that will compare LHS and RHS with
2967 comparison specified by CODE and put result into register DEST. DEST has to
2968 have its type set already but must not have its definition set yet.
2969 Generated instructions will be added to HBB. */
2972 gen_hsa_cmp_insn_from_gimple (enum tree_code code
, tree lhs
, tree rhs
,
2973 hsa_op_reg
*dest
, hsa_bb
*hbb
)
2975 BrigCompareOperation8_t compare
;
2980 compare
= BRIG_COMPARE_LT
;
2983 compare
= BRIG_COMPARE_LE
;
2986 compare
= BRIG_COMPARE_GT
;
2989 compare
= BRIG_COMPARE_GE
;
2992 compare
= BRIG_COMPARE_EQ
;
2995 compare
= BRIG_COMPARE_NE
;
2997 case UNORDERED_EXPR
:
2998 compare
= BRIG_COMPARE_NAN
;
3001 compare
= BRIG_COMPARE_NUM
;
3004 compare
= BRIG_COMPARE_LTU
;
3007 compare
= BRIG_COMPARE_LEU
;
3010 compare
= BRIG_COMPARE_GTU
;
3013 compare
= BRIG_COMPARE_GEU
;
3016 compare
= BRIG_COMPARE_EQU
;
3019 compare
= BRIG_COMPARE_NEU
;
3023 HSA_SORRY_ATV (EXPR_LOCATION (lhs
),
3024 "support for HSA does not implement comparison tree "
3025 "code %s\n", get_tree_code_name (code
));
3029 /* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer)
3030 as a result of comparison. */
3032 BrigType16_t dest_type
= hsa_type_integer_p (dest
->m_type
)
3033 ? (BrigType16_t
) BRIG_TYPE_B1
: dest
->m_type
;
3035 hsa_insn_cmp
*cmp
= new hsa_insn_cmp (compare
, dest_type
);
3036 cmp
->set_op (1, hsa_reg_or_immed_for_gimple_op (lhs
, hbb
));
3037 cmp
->set_op (2, hsa_reg_or_immed_for_gimple_op (rhs
, hbb
));
3039 hbb
->append_insn (cmp
);
3040 cmp
->set_output_in_type (dest
, 0, hbb
);
3043 /* Generate an unary instruction with OPCODE and append it to a basic block
3044 HBB. The instruction uses DEST as a destination and OP1
3045 as a single operand. */
3048 gen_hsa_unary_operation (BrigOpcode opcode
, hsa_op_reg
*dest
,
3049 hsa_op_with_type
*op1
, hsa_bb
*hbb
)
3051 gcc_checking_assert (dest
);
3052 hsa_insn_basic
*insn
;
3054 if (opcode
== BRIG_OPCODE_MOV
&& hsa_needs_cvt (dest
->m_type
, op1
->m_type
))
3055 insn
= new hsa_insn_cvt (dest
, op1
);
3056 else if (opcode
== BRIG_OPCODE_FIRSTBIT
|| opcode
== BRIG_OPCODE_LASTBIT
)
3057 insn
= new hsa_insn_srctype (2, opcode
, BRIG_TYPE_U32
, op1
->m_type
, NULL
,
3061 insn
= new hsa_insn_basic (2, opcode
, dest
->m_type
, dest
, op1
);
3063 if (opcode
== BRIG_OPCODE_ABS
|| opcode
== BRIG_OPCODE_NEG
)
3065 /* ABS and NEG only exist in _s form :-/ */
3066 if (insn
->m_type
== BRIG_TYPE_U32
)
3067 insn
->m_type
= BRIG_TYPE_S32
;
3068 else if (insn
->m_type
== BRIG_TYPE_U64
)
3069 insn
->m_type
= BRIG_TYPE_S64
;
3073 hbb
->append_insn (insn
);
3075 if (opcode
== BRIG_OPCODE_FIRSTBIT
|| opcode
== BRIG_OPCODE_LASTBIT
)
3076 insn
->set_output_in_type (dest
, 0, hbb
);
3079 /* Generate a binary instruction with OPCODE and append it to a basic block
3080 HBB. The instruction uses DEST as a destination and operands OP1
3084 gen_hsa_binary_operation (int opcode
, hsa_op_reg
*dest
,
3085 hsa_op_base
*op1
, hsa_op_base
*op2
, hsa_bb
*hbb
)
3087 gcc_checking_assert (dest
);
3089 if ((opcode
== BRIG_OPCODE_SHL
|| opcode
== BRIG_OPCODE_SHR
)
3090 && is_a
<hsa_op_immed
*> (op2
))
3092 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3093 i
->set_type (BRIG_TYPE_U32
);
3095 if ((opcode
== BRIG_OPCODE_OR
3096 || opcode
== BRIG_OPCODE_XOR
3097 || opcode
== BRIG_OPCODE_AND
)
3098 && is_a
<hsa_op_immed
*> (op2
))
3100 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3101 i
->set_type (hsa_unsigned_type_for_type (i
->m_type
));
3104 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, dest
->m_type
, dest
,
3106 hbb
->append_insn (insn
);
3109 /* Generate HSA instructions for a single assignment. HBB is the basic block
3110 they will be appended to. */
3113 gen_hsa_insns_for_operation_assignment (gimple
*assign
, hsa_bb
*hbb
)
3115 tree_code code
= gimple_assign_rhs_code (assign
);
3116 gimple_rhs_class rhs_class
= get_gimple_rhs_class (gimple_expr_code (assign
));
3118 tree lhs
= gimple_assign_lhs (assign
);
3119 tree rhs1
= gimple_assign_rhs1 (assign
);
3120 tree rhs2
= gimple_assign_rhs2 (assign
);
3121 tree rhs3
= gimple_assign_rhs3 (assign
);
3129 /* The opcode is changed to BRIG_OPCODE_CVT if BRIG types
3130 needs a conversion. */
3131 opcode
= BRIG_OPCODE_MOV
;
3135 case POINTER_PLUS_EXPR
:
3136 opcode
= BRIG_OPCODE_ADD
;
3139 opcode
= BRIG_OPCODE_SUB
;
3142 opcode
= BRIG_OPCODE_MUL
;
3144 case MULT_HIGHPART_EXPR
:
3145 opcode
= BRIG_OPCODE_MULHI
;
3148 case TRUNC_DIV_EXPR
:
3149 case EXACT_DIV_EXPR
:
3150 opcode
= BRIG_OPCODE_DIV
;
3153 case FLOOR_DIV_EXPR
:
3154 case ROUND_DIV_EXPR
:
3155 HSA_SORRY_AT (gimple_location (assign
),
3156 "support for HSA does not implement CEIL_DIV_EXPR, "
3157 "FLOOR_DIV_EXPR or ROUND_DIV_EXPR");
3159 case TRUNC_MOD_EXPR
:
3160 opcode
= BRIG_OPCODE_REM
;
3163 case FLOOR_MOD_EXPR
:
3164 case ROUND_MOD_EXPR
:
3165 HSA_SORRY_AT (gimple_location (assign
),
3166 "support for HSA does not implement CEIL_MOD_EXPR, "
3167 "FLOOR_MOD_EXPR or ROUND_MOD_EXPR");
3170 opcode
= BRIG_OPCODE_NEG
;
3173 opcode
= BRIG_OPCODE_MIN
;
3176 opcode
= BRIG_OPCODE_MAX
;
3179 opcode
= BRIG_OPCODE_ABS
;
3182 opcode
= BRIG_OPCODE_SHL
;
3185 opcode
= BRIG_OPCODE_SHR
;
3190 hsa_insn_basic
*insn
= NULL
;
3191 int code1
= code
== LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3192 int code2
= code
!= LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3193 BrigType16_t btype
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
3196 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3197 hsa_op_reg
*op1
= new hsa_op_reg (btype
);
3198 hsa_op_reg
*op2
= new hsa_op_reg (btype
);
3199 hsa_op_with_type
*shift1
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3201 tree type
= TREE_TYPE (rhs2
);
3202 unsigned HOST_WIDE_INT bitsize
= TREE_INT_CST_LOW (TYPE_SIZE (type
));
3204 hsa_op_with_type
*shift2
= NULL
;
3205 if (TREE_CODE (rhs2
) == INTEGER_CST
)
3206 shift2
= new hsa_op_immed (bitsize
- tree_to_uhwi (rhs2
),
3208 else if (TREE_CODE (rhs2
) == SSA_NAME
)
3210 hsa_op_reg
*s
= hsa_cfun
->reg_for_gimple_ssa (rhs2
);
3211 hsa_op_reg
*d
= new hsa_op_reg (s
->m_type
);
3212 hsa_op_immed
*size_imm
= new hsa_op_immed (bitsize
, BRIG_TYPE_U32
);
3214 insn
= new hsa_insn_basic (3, BRIG_OPCODE_SUB
, d
->m_type
,
3216 hbb
->append_insn (insn
);
3223 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3224 gen_hsa_binary_operation (code1
, op1
, src
, shift1
, hbb
);
3225 gen_hsa_binary_operation (code2
, op2
, src
, shift2
, hbb
);
3226 gen_hsa_binary_operation (BRIG_OPCODE_OR
, dest
, op1
, op2
, hbb
);
3231 opcode
= BRIG_OPCODE_OR
;
3234 opcode
= BRIG_OPCODE_XOR
;
3237 opcode
= BRIG_OPCODE_AND
;
3240 opcode
= BRIG_OPCODE_NOT
;
3242 case FIX_TRUNC_EXPR
:
3244 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3245 hsa_op_with_type
*v
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3247 if (hsa_needs_cvt (dest
->m_type
, v
->m_type
))
3249 hsa_op_reg
*tmp
= new hsa_op_reg (v
->m_type
);
3251 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3252 tmp
->m_type
, tmp
, v
);
3253 hbb
->append_insn (insn
);
3255 hsa_insn_basic
*cvtinsn
= new hsa_insn_cvt (dest
, tmp
);
3256 hbb
->append_insn (cvtinsn
);
3260 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3261 dest
->m_type
, dest
, v
);
3262 hbb
->append_insn (insn
);
3267 opcode
= BRIG_OPCODE_TRUNC
;
3276 case UNORDERED_EXPR
:
3286 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3288 gen_hsa_cmp_insn_from_gimple (code
, rhs1
, rhs2
, dest
, hbb
);
3294 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3295 hsa_op_with_type
*ctrl
= NULL
;
3298 if (CONSTANT_CLASS_P (cond
) || TREE_CODE (cond
) == SSA_NAME
)
3299 ctrl
= hsa_reg_or_immed_for_gimple_op (cond
, hbb
);
3302 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
3304 gen_hsa_cmp_insn_from_gimple (TREE_CODE (cond
),
3305 TREE_OPERAND (cond
, 0),
3306 TREE_OPERAND (cond
, 1),
3312 hsa_op_with_type
*op2
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3313 hsa_op_with_type
*op3
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3315 BrigType16_t utype
= hsa_unsigned_type_for_type (dest
->m_type
);
3316 if (is_a
<hsa_op_immed
*> (op2
))
3317 op2
->m_type
= utype
;
3318 if (is_a
<hsa_op_immed
*> (op3
))
3319 op3
->m_type
= utype
;
3321 hsa_insn_basic
*insn
3322 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
,
3323 hsa_bittype_for_type (dest
->m_type
),
3324 dest
, ctrl
, op2
, op3
);
3326 hbb
->append_insn (insn
);
3332 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3333 hsa_op_with_type
*rhs1_reg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3334 hsa_op_with_type
*rhs2_reg
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3336 if (hsa_seen_error ())
3339 BrigType16_t src_type
= hsa_bittype_for_type (rhs1_reg
->m_type
);
3340 rhs1_reg
= rhs1_reg
->get_in_type (src_type
, hbb
);
3341 rhs2_reg
= rhs2_reg
->get_in_type (src_type
, hbb
);
3343 hsa_insn_packed
*insn
3344 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
, src_type
,
3345 dest
, rhs1_reg
, rhs2_reg
);
3346 hbb
->append_insn (insn
);
3351 /* Implement others as we come across them. */
3352 HSA_SORRY_ATV (gimple_location (assign
),
3353 "support for HSA does not implement operation %s",
3354 get_tree_code_name (code
));
3359 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3361 hsa_op_with_type
*op1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3362 hsa_op_with_type
*op2
= rhs2
!= NULL_TREE
?
3363 hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
) : NULL
;
3365 if (hsa_seen_error ())
3370 case GIMPLE_TERNARY_RHS
:
3375 case GIMPLE_BINARY_RHS
:
3376 gen_hsa_binary_operation (opcode
, dest
, op1
, op2
, hbb
);
3379 case GIMPLE_UNARY_RHS
:
3380 gen_hsa_unary_operation (opcode
, dest
, op1
, hbb
);
3387 /* Generate HSA instructions for a given gimple condition statement COND.
3388 Instructions will be appended to HBB, which also needs to be the
3389 corresponding structure to the basic_block of COND. */
3392 gen_hsa_insns_for_cond_stmt (gimple
*cond
, hsa_bb
*hbb
)
3394 hsa_op_reg
*ctrl
= new hsa_op_reg (BRIG_TYPE_B1
);
3397 gen_hsa_cmp_insn_from_gimple (gimple_cond_code (cond
),
3398 gimple_cond_lhs (cond
),
3399 gimple_cond_rhs (cond
),
3402 cbr
= new hsa_insn_br (ctrl
);
3403 hbb
->append_insn (cbr
);
3406 /* Maximum number of elements in a jump table for an HSA SBR instruction. */
3408 #define HSA_MAXIMUM_SBR_LABELS 16
3410 /* Return lowest value of a switch S that is handled in a non-default
3414 get_switch_low (gswitch
*s
)
3416 unsigned labels
= gimple_switch_num_labels (s
);
3417 gcc_checking_assert (labels
>= 1);
3419 return CASE_LOW (gimple_switch_label (s
, 1));
3422 /* Return highest value of a switch S that is handled in a non-default
3426 get_switch_high (gswitch
*s
)
3428 unsigned labels
= gimple_switch_num_labels (s
);
3430 /* Compare last label to maximum number of labels. */
3431 tree label
= gimple_switch_label (s
, labels
- 1);
3432 tree low
= CASE_LOW (label
);
3433 tree high
= CASE_HIGH (label
);
3435 return high
!= NULL_TREE
? high
: low
;
3439 get_switch_size (gswitch
*s
)
3441 return int_const_binop (MINUS_EXPR
, get_switch_high (s
), get_switch_low (s
));
3444 /* Generate HSA instructions for a given gimple switch.
3445 Instructions will be appended to HBB. */
3448 gen_hsa_insns_for_switch_stmt (gswitch
*s
, hsa_bb
*hbb
)
3450 gimple_stmt_iterator it
= gsi_for_stmt (s
);
3453 /* Create preambule that verifies that index - lowest_label >= 0. */
3454 edge e
= split_block (hbb
->m_bb
, gsi_stmt (it
));
3455 e
->flags
&= ~EDGE_FALLTHRU
;
3456 e
->flags
|= EDGE_TRUE_VALUE
;
3458 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
3459 tree index_tree
= gimple_switch_index (s
);
3460 tree lowest
= get_switch_low (s
);
3461 tree highest
= get_switch_high (s
);
3463 hsa_op_reg
*index
= hsa_cfun
->reg_for_gimple_ssa (index_tree
);
3465 hsa_op_reg
*cmp1_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3466 hsa_op_immed
*cmp1_immed
= new hsa_op_immed (lowest
);
3467 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_GE
, cmp1_reg
->m_type
,
3468 cmp1_reg
, index
, cmp1_immed
));
3470 hsa_op_reg
*cmp2_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3471 hsa_op_immed
*cmp2_immed
= new hsa_op_immed (highest
);
3472 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_LE
, cmp2_reg
->m_type
,
3473 cmp2_reg
, index
, cmp2_immed
));
3475 hsa_op_reg
*cmp_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3476 hbb
->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_AND
, cmp_reg
->m_type
,
3477 cmp_reg
, cmp1_reg
, cmp2_reg
));
3479 hbb
->append_insn (new hsa_insn_br (cmp_reg
));
3481 tree default_label
= gimple_switch_default_label (s
);
3482 basic_block default_label_bb
= label_to_block_fn (func
,
3483 CASE_LABEL (default_label
));
3485 make_edge (e
->src
, default_label_bb
, EDGE_FALSE_VALUE
);
3487 hsa_cfun
->m_modified_cfg
= true;
3489 /* Basic block with the SBR instruction. */
3490 hbb
= hsa_init_new_bb (e
->dest
);
3492 hsa_op_reg
*sub_index
= new hsa_op_reg (index
->m_type
);
3493 hbb
->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB
, sub_index
->m_type
,
3495 new hsa_op_immed (lowest
)));
3497 hsa_op_base
*tmp
= sub_index
->get_in_type (BRIG_TYPE_U64
, hbb
);
3498 sub_index
= as_a
<hsa_op_reg
*> (tmp
);
3499 unsigned labels
= gimple_switch_num_labels (s
);
3500 unsigned HOST_WIDE_INT size
= tree_to_uhwi (get_switch_size (s
));
3502 hsa_insn_sbr
*sbr
= new hsa_insn_sbr (sub_index
, size
+ 1);
3504 /* Prepare array with default label destination. */
3505 for (unsigned HOST_WIDE_INT i
= 0; i
<= size
; i
++)
3506 sbr
->m_jump_table
.safe_push (default_label_bb
);
3508 /* Iterate all labels and fill up the jump table. */
3509 for (unsigned i
= 1; i
< labels
; i
++)
3511 tree label
= gimple_switch_label (s
, i
);
3512 basic_block bb
= label_to_block_fn (func
, CASE_LABEL (label
));
3514 unsigned HOST_WIDE_INT sub_low
3515 = tree_to_uhwi (int_const_binop (MINUS_EXPR
, CASE_LOW (label
), lowest
));
3517 unsigned HOST_WIDE_INT sub_high
= sub_low
;
3518 tree high
= CASE_HIGH (label
);
3520 sub_high
= tree_to_uhwi (int_const_binop (MINUS_EXPR
, high
, lowest
));
3522 for (unsigned HOST_WIDE_INT j
= sub_low
; j
<= sub_high
; j
++)
3523 sbr
->m_jump_table
[j
] = bb
;
3526 hbb
->append_insn (sbr
);
3529 /* Verify that the function DECL can be handled by HSA. */
3532 verify_function_arguments (tree decl
)
3534 if (DECL_STATIC_CHAIN (decl
))
3536 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3537 "HSA does not support nested functions: %D", decl
);
3540 else if (!TYPE_ARG_TYPES (TREE_TYPE (decl
)))
3542 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3543 "HSA does not support functions with variadic arguments "
3544 "(or unknown return type): %D", decl
);
3549 /* Return BRIG type for FORMAL_ARG_TYPE. If the formal argument type is NULL,
3550 return ACTUAL_ARG_TYPE. */
3553 get_format_argument_type (tree formal_arg_type
, BrigType16_t actual_arg_type
)
3555 if (formal_arg_type
== NULL
)
3556 return actual_arg_type
;
3558 BrigType16_t decl_type
3559 = hsa_type_for_scalar_tree_type (formal_arg_type
, false);
3560 return mem_type_for_type (decl_type
);
3563 /* Generate HSA instructions for a direct call instruction.
3564 Instructions will be appended to HBB, which also needs to be the
3565 corresponding structure to the basic_block of STMT.
3566 If ASSIGN_LHS is false, do not copy HSA function result argument into the
3567 corresponding HSA representation of the gimple statement LHS. */
3570 gen_hsa_insns_for_direct_call (gimple
*stmt
, hsa_bb
*hbb
,
3571 bool assign_lhs
= true)
3573 tree decl
= gimple_call_fndecl (stmt
);
3574 verify_function_arguments (decl
);
3575 if (hsa_seen_error ())
3578 hsa_insn_call
*call_insn
= new hsa_insn_call (decl
);
3579 hsa_cfun
->m_called_functions
.safe_push (call_insn
->m_called_function
);
3581 /* Argument block start. */
3582 hsa_insn_arg_block
*arg_start
3583 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3584 hbb
->append_insn (arg_start
);
3586 tree parm_type_chain
= TYPE_ARG_TYPES (gimple_call_fntype (stmt
));
3588 /* Preparation of arguments that will be passed to function. */
3589 const unsigned args
= gimple_call_num_args (stmt
);
3590 for (unsigned i
= 0; i
< args
; ++i
)
3592 tree parm
= gimple_call_arg (stmt
, (int)i
);
3593 tree parm_decl_type
= parm_type_chain
!= NULL_TREE
3594 ? TREE_VALUE (parm_type_chain
) : NULL_TREE
;
3595 hsa_op_address
*addr
;
3597 if (AGGREGATE_TYPE_P (TREE_TYPE (parm
)))
3599 addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3600 BrigAlignment8_t align
;
3601 hsa_op_address
*src
= gen_hsa_addr_with_align (parm
, hbb
, &align
);
3602 gen_hsa_memory_copy (hbb
, addr
, src
,
3603 addr
->m_symbol
->total_byte_size (), align
);
3607 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3609 if (parm_decl_type
!= NULL
&& AGGREGATE_TYPE_P (parm_decl_type
))
3611 HSA_SORRY_AT (gimple_location (stmt
),
3612 "support for HSA does not implement an aggregate "
3613 "formal argument in a function call, while actual "
3614 "argument is not an aggregate");
3618 BrigType16_t formal_arg_type
3619 = get_format_argument_type (parm_decl_type
, src
->m_type
);
3620 if (hsa_seen_error ())
3623 if (src
->m_type
!= formal_arg_type
)
3624 src
= src
->get_in_type (formal_arg_type
, hbb
);
3627 = gen_hsa_addr_for_arg (parm_decl_type
!= NULL_TREE
?
3628 parm_decl_type
: TREE_TYPE (parm
), i
);
3629 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, formal_arg_type
,
3632 hbb
->append_insn (mem
);
3635 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3636 if (parm_type_chain
)
3637 parm_type_chain
= TREE_CHAIN (parm_type_chain
);
3640 call_insn
->m_args_code_list
= new hsa_op_code_list (args
);
3641 hbb
->append_insn (call_insn
);
3643 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
3645 tree result
= gimple_call_lhs (stmt
);
3646 hsa_insn_mem
*result_insn
= NULL
;
3647 if (!VOID_TYPE_P (result_type
))
3649 hsa_op_address
*addr
= gen_hsa_addr_for_arg (result_type
, -1);
3651 /* Even if result of a function call is unused, we have to emit
3652 declaration for the result. */
3653 if (result
&& assign_lhs
)
3655 tree lhs_type
= TREE_TYPE (result
);
3657 if (hsa_seen_error ())
3660 if (AGGREGATE_TYPE_P (lhs_type
))
3662 BrigAlignment8_t align
;
3663 hsa_op_address
*result_addr
3664 = gen_hsa_addr_with_align (result
, hbb
, &align
);
3665 gen_hsa_memory_copy (hbb
, result_addr
, addr
,
3666 addr
->m_symbol
->total_byte_size (), align
);
3671 = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type
,
3674 hsa_op_reg
*dst
= hsa_cfun
->reg_for_gimple_ssa (result
);
3675 result_insn
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
, dst
, addr
);
3676 hbb
->append_insn (result_insn
);
3680 call_insn
->m_output_arg
= addr
->m_symbol
;
3681 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3687 HSA_SORRY_AT (gimple_location (stmt
),
3688 "support for HSA does not implement an assignment of "
3689 "return value from a void function");
3693 call_insn
->m_result_code_list
= new hsa_op_code_list (0);
3696 /* Argument block end. */
3697 hsa_insn_arg_block
*arg_end
3698 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3699 hbb
->append_insn (arg_end
);
3702 /* Generate HSA instructions for a direct call of an internal fn.
3703 Instructions will be appended to HBB, which also needs to be the
3704 corresponding structure to the basic_block of STMT. */
3707 gen_hsa_insns_for_call_of_internal_fn (gimple
*stmt
, hsa_bb
*hbb
)
3709 tree lhs
= gimple_call_lhs (stmt
);
3713 tree lhs_type
= TREE_TYPE (lhs
);
3714 tree rhs1
= gimple_call_arg (stmt
, 0);
3715 tree rhs1_type
= TREE_TYPE (rhs1
);
3716 enum internal_fn fn
= gimple_call_internal_fn (stmt
);
3717 hsa_internal_fn
*ifn
3718 = new hsa_internal_fn (fn
, tree_to_uhwi (TYPE_SIZE (rhs1_type
)));
3719 hsa_insn_call
*call_insn
= new hsa_insn_call (ifn
);
3721 gcc_checking_assert (FLOAT_TYPE_P (rhs1_type
));
3723 if (!hsa_emitted_internal_decls
->find (call_insn
->m_called_internal_fn
))
3724 hsa_cfun
->m_called_internal_fns
.safe_push (call_insn
->m_called_internal_fn
);
3726 hsa_insn_arg_block
*arg_start
3727 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3728 hbb
->append_insn (arg_start
);
3730 unsigned num_args
= gimple_call_num_args (stmt
);
3732 /* Function arguments. */
3733 for (unsigned i
= 0; i
< num_args
; i
++)
3735 tree parm
= gimple_call_arg (stmt
, (int)i
);
3736 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3738 hsa_op_address
*addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3739 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, src
->m_type
,
3742 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3743 hbb
->append_insn (mem
);
3746 call_insn
->m_args_code_list
= new hsa_op_code_list (num_args
);
3747 hbb
->append_insn (call_insn
);
3749 /* Assign returned value. */
3750 hsa_op_address
*addr
= gen_hsa_addr_for_arg (lhs_type
, -1);
3752 call_insn
->m_output_arg
= addr
->m_symbol
;
3753 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3755 /* Argument block end. */
3756 hsa_insn_arg_block
*arg_end
3757 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3758 hbb
->append_insn (arg_end
);
3761 /* Generate HSA instructions for a return value instruction.
3762 Instructions will be appended to HBB, which also needs to be the
3763 corresponding structure to the basic_block of STMT. */
3766 gen_hsa_insns_for_return (greturn
*stmt
, hsa_bb
*hbb
)
3768 tree retval
= gimple_return_retval (stmt
);
3771 hsa_op_address
*addr
= new hsa_op_address (hsa_cfun
->m_output_arg
);
3773 if (AGGREGATE_TYPE_P (TREE_TYPE (retval
)))
3775 BrigAlignment8_t align
;
3776 hsa_op_address
*retval_addr
= gen_hsa_addr_with_align (retval
, hbb
,
3778 gen_hsa_memory_copy (hbb
, addr
, retval_addr
,
3779 hsa_cfun
->m_output_arg
->total_byte_size (),
3784 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (retval
),
3786 BrigType16_t mtype
= mem_type_for_type (t
);
3788 /* Store of return value. */
3789 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (retval
, hbb
);
3790 src
= src
->get_in_type (mtype
, hbb
);
3791 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
,
3793 hbb
->append_insn (mem
);
3797 /* HSAIL return instruction emission. */
3798 hsa_insn_basic
*ret
= new hsa_insn_basic (0, BRIG_OPCODE_RET
);
3799 hbb
->append_insn (ret
);
3802 /* Set OP_INDEX-th operand of the instruction to DEST, as the DEST
3803 can have a different type, conversion instructions are possibly
3807 hsa_insn_basic::set_output_in_type (hsa_op_reg
*dest
, unsigned op_index
,
3810 hsa_insn_basic
*insn
;
3811 gcc_checking_assert (op_output_p (op_index
));
3813 if (dest
->m_type
== m_type
)
3815 set_op (op_index
, dest
);
3819 hsa_op_reg
*tmp
= new hsa_op_reg (m_type
);
3820 set_op (op_index
, tmp
);
3822 if (hsa_needs_cvt (dest
->m_type
, m_type
))
3823 insn
= new hsa_insn_cvt (dest
, tmp
);
3825 insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
,
3826 dest
, tmp
->get_in_type (dest
->m_type
, hbb
));
3828 hbb
->append_insn (insn
);
3831 /* Generate instruction OPCODE to query a property of HSA grid along the
3832 given DIMENSION. Store result into DEST and append the instruction to
3836 query_hsa_grid (hsa_op_reg
*dest
, BrigType16_t opcode
, int dimension
,
3839 /* We're using just one-dimensional kernels, so hard-coded
3842 = new hsa_op_immed (dimension
, (BrigKind16_t
) BRIG_TYPE_U32
);
3843 hsa_insn_basic
*insn
= new hsa_insn_basic (2, opcode
, BRIG_TYPE_U32
, NULL
,
3845 hbb
->append_insn (insn
);
3846 insn
->set_output_in_type (dest
, 0, hbb
);
3849 /* Generate a special HSA-related instruction for gimple STMT.
3850 Instructions are appended to basic block HBB. */
3853 query_hsa_grid (gimple
*stmt
, BrigOpcode16_t opcode
, int dimension
,
3856 tree lhs
= gimple_call_lhs (dyn_cast
<gcall
*> (stmt
));
3857 if (lhs
== NULL_TREE
)
3860 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3862 query_hsa_grid (dest
, opcode
, dimension
, hbb
);
3865 /* Emit instructions that set hsa_num_threads according to provided VALUE.
3866 Instructions are appended to basic block HBB. */
3869 gen_set_num_threads (tree value
, hsa_bb
*hbb
)
3871 hbb
->append_insn (new hsa_insn_comment ("omp_set_num_threads"));
3872 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (value
, hbb
);
3874 src
= src
->get_in_type (hsa_num_threads
->m_type
, hbb
);
3875 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
3877 hsa_insn_basic
*basic
3878 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
, src
, addr
);
3879 hbb
->append_insn (basic
);
3882 /* Return byte offset of a FIELD_NAME in GOMP_hsa_kernel_dispatch which
3883 is defined in plugin-hsa.c. */
3885 static HOST_WIDE_INT
3886 get_hsa_kernel_dispatch_offset (const char *field_name
)
3888 tree
*hsa_kernel_dispatch_type
= hsa_get_kernel_dispatch_type ();
3889 if (*hsa_kernel_dispatch_type
== NULL
)
3891 /* Collection of information needed for a dispatch of a kernel from a
3892 kernel. Keep in sync with libgomp's plugin-hsa.c. */
3894 *hsa_kernel_dispatch_type
= make_node (RECORD_TYPE
);
3895 tree id_f1
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3896 get_identifier ("queue"), ptr_type_node
);
3897 DECL_CHAIN (id_f1
) = NULL_TREE
;
3898 tree id_f2
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3899 get_identifier ("omp_data_memory"),
3901 DECL_CHAIN (id_f2
) = id_f1
;
3902 tree id_f3
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3903 get_identifier ("kernarg_address"),
3905 DECL_CHAIN (id_f3
) = id_f2
;
3906 tree id_f4
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3907 get_identifier ("object"),
3909 DECL_CHAIN (id_f4
) = id_f3
;
3910 tree id_f5
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3911 get_identifier ("signal"),
3913 DECL_CHAIN (id_f5
) = id_f4
;
3914 tree id_f6
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3915 get_identifier ("private_segment_size"),
3917 DECL_CHAIN (id_f6
) = id_f5
;
3918 tree id_f7
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3919 get_identifier ("group_segment_size"),
3921 DECL_CHAIN (id_f7
) = id_f6
;
3922 tree id_f8
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3923 get_identifier ("kernel_dispatch_count"),
3925 DECL_CHAIN (id_f8
) = id_f7
;
3926 tree id_f9
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3927 get_identifier ("debug"),
3929 DECL_CHAIN (id_f9
) = id_f8
;
3930 tree id_f10
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3931 get_identifier ("omp_level"),
3933 DECL_CHAIN (id_f10
) = id_f9
;
3934 tree id_f11
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3935 get_identifier ("children_dispatches"),
3937 DECL_CHAIN (id_f11
) = id_f10
;
3938 tree id_f12
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3939 get_identifier ("omp_num_threads"),
3941 DECL_CHAIN (id_f12
) = id_f11
;
3944 finish_builtin_struct (*hsa_kernel_dispatch_type
, "__hsa_kernel_dispatch",
3946 TYPE_ARTIFICIAL (*hsa_kernel_dispatch_type
) = 1;
3949 for (tree chain
= TYPE_FIELDS (*hsa_kernel_dispatch_type
);
3950 chain
!= NULL_TREE
; chain
= TREE_CHAIN (chain
))
3951 if (strcmp (field_name
, IDENTIFIER_POINTER (DECL_NAME (chain
))) == 0)
3952 return int_byte_position (chain
);
3957 /* Return an HSA register that will contain number of threads for
3958 a future dispatched kernel. Instructions are added to HBB. */
3961 gen_num_threads_for_dispatch (hsa_bb
*hbb
)
3963 /* Step 1) Assign to number of threads:
3964 MIN (HSA_DEFAULT_NUM_THREADS, hsa_num_threads). */
3965 hsa_op_reg
*threads
= new hsa_op_reg (hsa_num_threads
->m_type
);
3966 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
3968 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD
, threads
->m_type
,
3971 hsa_op_immed
*limit
= new hsa_op_immed (HSA_DEFAULT_NUM_THREADS
,
3973 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
3975 = new hsa_insn_cmp (BRIG_COMPARE_LT
, r
->m_type
, r
, threads
, limit
);
3976 hbb
->append_insn (cmp
);
3978 BrigType16_t btype
= hsa_bittype_for_type (threads
->m_type
);
3979 hsa_op_reg
*tmp
= new hsa_op_reg (threads
->m_type
);
3981 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp
, r
,
3984 /* Step 2) If the number is equal to zero,
3985 return shadow->omp_num_threads. */
3986 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
3988 hsa_op_reg
*shadow_thread_count
= new hsa_op_reg (BRIG_TYPE_U32
);
3990 = new hsa_op_address (shadow_reg_ptr
,
3991 get_hsa_kernel_dispatch_offset ("omp_num_threads"));
3992 hsa_insn_basic
*basic
3993 = new hsa_insn_mem (BRIG_OPCODE_LD
, shadow_thread_count
->m_type
,
3994 shadow_thread_count
, addr
);
3995 hbb
->append_insn (basic
);
3997 hsa_op_reg
*tmp2
= new hsa_op_reg (threads
->m_type
);
3998 r
= new hsa_op_reg (BRIG_TYPE_B1
);
3999 hsa_op_immed
*imm
= new hsa_op_immed (0, shadow_thread_count
->m_type
);
4000 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_EQ
, r
->m_type
, r
, tmp
, imm
));
4001 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp2
, r
,
4002 shadow_thread_count
, tmp
));
4004 hsa_op_base
*dest
= tmp2
->get_in_type (BRIG_TYPE_U16
, hbb
);
4006 return as_a
<hsa_op_reg
*> (dest
);
4010 /* Emit instructions that assign number of teams to lhs of gimple STMT.
4011 Instructions are appended to basic block HBB. */
4014 gen_get_num_teams (gimple
*stmt
, hsa_bb
*hbb
)
4016 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4019 hbb
->append_insn (new hsa_insn_comment ("omp_get_num_teams"));
4021 tree lhs
= gimple_call_lhs (stmt
);
4022 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4023 hsa_op_immed
*one
= new hsa_op_immed (1, dest
->m_type
);
4025 hsa_insn_basic
*basic
4026 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, dest
, one
);
4028 hbb
->append_insn (basic
);
4031 /* Emit instructions that assign a team number to lhs of gimple STMT.
4032 Instructions are appended to basic block HBB. */
4035 gen_get_team_num (gimple
*stmt
, hsa_bb
*hbb
)
4037 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4040 hbb
->append_insn (new hsa_insn_comment ("omp_get_team_num"));
4042 tree lhs
= gimple_call_lhs (stmt
);
4043 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4044 hsa_op_immed
*zero
= new hsa_op_immed (0, dest
->m_type
);
4046 hsa_insn_basic
*basic
4047 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, dest
, zero
);
4049 hbb
->append_insn (basic
);
4052 /* Emit instructions that get levels-var ICV to lhs of gimple STMT.
4053 Instructions are appended to basic block HBB. */
4056 gen_get_level (gimple
*stmt
, hsa_bb
*hbb
)
4058 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4061 hbb
->append_insn (new hsa_insn_comment ("omp_get_level"));
4063 tree lhs
= gimple_call_lhs (stmt
);
4064 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4066 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4067 if (shadow_reg_ptr
== NULL
)
4069 HSA_SORRY_AT (gimple_location (stmt
),
4070 "support for HSA does not implement omp_get_level called "
4071 "from a function not being inlined within a kernel");
4075 hsa_op_address
*addr
4076 = new hsa_op_address (shadow_reg_ptr
,
4077 get_hsa_kernel_dispatch_offset ("omp_level"));
4079 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, BRIG_TYPE_U64
,
4080 (hsa_op_base
*) NULL
, addr
);
4081 hbb
->append_insn (mem
);
4082 mem
->set_output_in_type (dest
, 0, hbb
);
4085 /* Emit instruction that implement omp_get_max_threads of gimple STMT. */
4088 gen_get_max_threads (gimple
*stmt
, hsa_bb
*hbb
)
4090 tree lhs
= gimple_call_lhs (stmt
);
4094 hbb
->append_insn (new hsa_insn_comment ("omp_get_max_threads"));
4096 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4097 hsa_op_with_type
*num_theads_reg
= gen_num_threads_for_dispatch (hbb
)
4098 ->get_in_type (dest
->m_type
, hbb
);
4099 hsa_build_append_simple_mov (dest
, num_theads_reg
, hbb
);
4102 /* Emit instructions that implement alloca builtin gimple STMT.
4103 Instructions are appended to basic block HBB. */
4106 gen_hsa_alloca (gcall
*call
, hsa_bb
*hbb
)
4108 tree lhs
= gimple_call_lhs (call
);
4109 if (lhs
== NULL_TREE
)
4112 built_in_function fn
= DECL_FUNCTION_CODE (gimple_call_fndecl (call
));
4114 gcc_checking_assert (fn
== BUILT_IN_ALLOCA
4115 || fn
== BUILT_IN_ALLOCA_WITH_ALIGN
);
4117 unsigned bit_alignment
= 0;
4119 if (fn
== BUILT_IN_ALLOCA_WITH_ALIGN
)
4121 tree alignment_tree
= gimple_call_arg (call
, 1);
4122 if (TREE_CODE (alignment_tree
) != INTEGER_CST
)
4124 HSA_SORRY_ATV (gimple_location (call
),
4125 "support for HSA does not implement "
4126 "__builtin_alloca_with_align with a non-constant "
4127 "alignment: %E", alignment_tree
);
4130 bit_alignment
= tree_to_uhwi (alignment_tree
);
4133 tree rhs1
= gimple_call_arg (call
, 0);
4134 hsa_op_with_type
*size
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
)
4135 ->get_in_type (BRIG_TYPE_U32
, hbb
);
4136 hsa_op_with_type
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4139 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
));
4140 hsa_insn_alloca
*a
= new hsa_insn_alloca (tmp
, size
, bit_alignment
);
4141 hbb
->append_insn (a
);
4144 = new hsa_insn_seg (BRIG_OPCODE_STOF
,
4145 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
4146 tmp
->m_type
, BRIG_SEGMENT_PRIVATE
, dest
, tmp
);
4147 hbb
->append_insn (seg
);
4150 /* Emit instructions that implement clrsb builtin STMT:
4151 Returns the number of leading redundant sign bits in x, i.e. the number
4152 of bits following the most significant bit that are identical to it.
4153 There are no special cases for 0 or other values.
4154 Instructions are appended to basic block HBB. */
4157 gen_hsa_clrsb (gcall
*call
, hsa_bb
*hbb
)
4159 tree lhs
= gimple_call_lhs (call
);
4160 if (lhs
== NULL_TREE
)
4163 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4164 tree rhs1
= gimple_call_arg (call
, 0);
4165 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4166 BrigType16_t bittype
= hsa_bittype_for_type (arg
->m_type
);
4167 unsigned bitsize
= tree_to_uhwi (TYPE_SIZE (TREE_TYPE (rhs1
)));
4169 /* FIRSTBIT instruction is defined just for 32 and 64-bits wide integers. */
4170 gcc_checking_assert (bitsize
== 32 || bitsize
== 64);
4172 /* Set true to MOST_SIG if the most significant bit is set to one. */
4173 hsa_op_immed
*c
= new hsa_op_immed (1ul << (bitsize
- 1),
4174 hsa_uint_for_bitsize (bitsize
));
4176 hsa_op_reg
*and_reg
= new hsa_op_reg (bittype
);
4177 gen_hsa_binary_operation (BRIG_OPCODE_AND
, and_reg
, arg
, c
, hbb
);
4179 hsa_op_reg
*most_sign
= new hsa_op_reg (BRIG_TYPE_B1
);
4181 = new hsa_insn_cmp (BRIG_COMPARE_EQ
, most_sign
->m_type
, most_sign
,
4183 hbb
->append_insn (cmp
);
4185 /* If the most significant bit is one, negate the input. Otherwise
4186 shift the input value to left by one bit. */
4187 hsa_op_reg
*arg_neg
= new hsa_op_reg (arg
->m_type
);
4188 gen_hsa_unary_operation (BRIG_OPCODE_NEG
, arg_neg
, arg
, hbb
);
4190 hsa_op_reg
*shifted_arg
= new hsa_op_reg (arg
->m_type
);
4191 gen_hsa_binary_operation (BRIG_OPCODE_SHL
, shifted_arg
, arg
,
4192 new hsa_op_immed (1, BRIG_TYPE_U64
), hbb
);
4194 /* Assign the value that can be used for FIRSTBIT instruction according
4195 to the most significant bit. */
4196 hsa_op_reg
*tmp
= new hsa_op_reg (bittype
);
4197 hsa_insn_basic
*cmov
4198 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, bittype
, tmp
, most_sign
,
4199 arg_neg
, shifted_arg
);
4200 hbb
->append_insn (cmov
);
4202 hsa_op_reg
*leading_bits
= new hsa_op_reg (BRIG_TYPE_S32
);
4203 gen_hsa_unary_operation (BRIG_OPCODE_FIRSTBIT
, leading_bits
,
4204 tmp
->get_in_type (hsa_uint_for_bitsize (bitsize
),
4207 /* Set flag if the input value is equal to zero. */
4208 hsa_op_reg
*is_zero
= new hsa_op_reg (BRIG_TYPE_B1
);
4209 cmp
= new hsa_insn_cmp (BRIG_COMPARE_EQ
, is_zero
->m_type
, is_zero
, arg
,
4210 new hsa_op_immed (0, arg
->m_type
));
4211 hbb
->append_insn (cmp
);
4213 /* Return the number of leading bits,
4214 or (bitsize - 1) if the input value is zero. */
4215 cmov
= new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, BRIG_TYPE_B32
, NULL
, is_zero
,
4216 new hsa_op_immed (bitsize
- 1, BRIG_TYPE_U32
),
4217 leading_bits
->get_in_type (BRIG_TYPE_B32
, hbb
));
4218 hbb
->append_insn (cmov
);
4219 cmov
->set_output_in_type (dest
, 0, hbb
);
4222 /* Emit instructions that implement ffs builtin STMT:
4223 Returns one plus the index of the least significant 1-bit of x,
4224 or if x is zero, returns zero.
4225 Instructions are appended to basic block HBB. */
4228 gen_hsa_ffs (gcall
*call
, hsa_bb
*hbb
)
4230 tree lhs
= gimple_call_lhs (call
);
4231 if (lhs
== NULL_TREE
)
4234 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4236 tree rhs1
= gimple_call_arg (call
, 0);
4237 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4239 hsa_op_reg
*tmp
= new hsa_op_reg (BRIG_TYPE_U32
);
4240 hsa_insn_srctype
*insn
= new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT
,
4241 tmp
->m_type
, arg
->m_type
,
4243 hbb
->append_insn (insn
);
4245 hsa_insn_basic
*addition
4246 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
, NULL
, tmp
,
4247 new hsa_op_immed (1, tmp
->m_type
));
4248 hbb
->append_insn (addition
);
4249 addition
->set_output_in_type (dest
, 0, hbb
);
4253 gen_hsa_popcount_to_dest (hsa_op_reg
*dest
, hsa_op_with_type
*arg
, hsa_bb
*hbb
)
4255 gcc_checking_assert (hsa_type_integer_p (arg
->m_type
));
4257 if (hsa_type_bit_size (arg
->m_type
) < 32)
4258 arg
= arg
->get_in_type (BRIG_TYPE_B32
, hbb
);
4260 if (!hsa_btype_p (arg
->m_type
))
4261 arg
= arg
->get_in_type (hsa_bittype_for_type (arg
->m_type
), hbb
);
4263 hsa_insn_srctype
*popcount
4264 = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT
, BRIG_TYPE_U32
,
4265 arg
->m_type
, NULL
, arg
);
4266 hbb
->append_insn (popcount
);
4267 popcount
->set_output_in_type (dest
, 0, hbb
);
4270 /* Emit instructions that implement parity builtin STMT:
4271 Returns the parity of x, i.e. the number of 1-bits in x modulo 2.
4272 Instructions are appended to basic block HBB. */
4275 gen_hsa_parity (gcall
*call
, hsa_bb
*hbb
)
4277 tree lhs
= gimple_call_lhs (call
);
4278 if (lhs
== NULL_TREE
)
4281 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4282 tree rhs1
= gimple_call_arg (call
, 0);
4283 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4285 hsa_op_reg
*popcount
= new hsa_op_reg (BRIG_TYPE_U32
);
4286 gen_hsa_popcount_to_dest (popcount
, arg
, hbb
);
4288 hsa_insn_basic
*insn
4289 = new hsa_insn_basic (3, BRIG_OPCODE_REM
, popcount
->m_type
, NULL
, popcount
,
4290 new hsa_op_immed (2, popcount
->m_type
));
4291 hbb
->append_insn (insn
);
4292 insn
->set_output_in_type (dest
, 0, hbb
);
4295 /* Emit instructions that implement popcount builtin STMT.
4296 Instructions are appended to basic block HBB. */
4299 gen_hsa_popcount (gcall
*call
, hsa_bb
*hbb
)
4301 tree lhs
= gimple_call_lhs (call
);
4302 if (lhs
== NULL_TREE
)
4305 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4306 tree rhs1
= gimple_call_arg (call
, 0);
4307 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4309 gen_hsa_popcount_to_dest (dest
, arg
, hbb
);
4312 /* Set VALUE to a shadow kernel debug argument and append a new instruction
4313 to HBB basic block. */
4316 set_debug_value (hsa_bb
*hbb
, hsa_op_with_type
*value
)
4318 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4319 if (shadow_reg_ptr
== NULL
)
4322 hsa_op_address
*addr
4323 = new hsa_op_address (shadow_reg_ptr
,
4324 get_hsa_kernel_dispatch_offset ("debug"));
4325 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, BRIG_TYPE_U64
, value
,
4327 hbb
->append_insn (mem
);
4331 omp_simple_builtin::generate (gimple
*stmt
, hsa_bb
*hbb
)
4335 if (m_warning_message
)
4336 HSA_SORRY_AT (gimple_location (stmt
), m_warning_message
)
4338 HSA_SORRY_ATV (gimple_location (stmt
),
4339 "Support for HSA does not implement calls to %s\n",
4342 else if (m_warning_message
!= NULL
)
4343 warning_at (gimple_location (stmt
), OPT_Whsa
, m_warning_message
);
4345 if (m_return_value
!= NULL
)
4347 tree lhs
= gimple_call_lhs (stmt
);
4351 hbb
->append_insn (new hsa_insn_comment (m_name
));
4353 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4354 hsa_op_with_type
*op
= m_return_value
->get_in_type (dest
->m_type
, hbb
);
4355 hsa_build_append_simple_mov (dest
, op
, hbb
);
4359 /* If STMT is a call of a known library function, generate code to perform
4360 it and return true. */
4363 gen_hsa_insns_for_known_library_call (gimple
*stmt
, hsa_bb
*hbb
)
4365 bool handled
= false;
4366 const char *name
= hsa_get_declaration_name (gimple_call_fndecl (stmt
));
4369 size_t len
= strlen (name
);
4370 if (len
> 0 && name
[len
- 1] == '_')
4372 copy
= XNEWVEC (char, len
+ 1);
4373 strcpy (copy
, name
);
4374 copy
[len
- 1] = '\0';
4378 /* Handle omp_* routines. */
4379 if (strstr (name
, "omp_") == name
)
4381 hsa_init_simple_builtins ();
4382 omp_simple_builtin
*builtin
= omp_simple_builtins
->get (name
);
4385 builtin
->generate (stmt
, hbb
);
4390 if (strcmp (name
, "omp_set_num_threads") == 0)
4391 gen_set_num_threads (gimple_call_arg (stmt
, 0), hbb
);
4392 else if (strcmp (name
, "omp_get_thread_num") == 0)
4394 hbb
->append_insn (new hsa_insn_comment (name
));
4395 query_hsa_grid (stmt
, BRIG_OPCODE_WORKITEMABSID
, 0, hbb
);
4397 else if (strcmp (name
, "omp_get_num_threads") == 0)
4399 hbb
->append_insn (new hsa_insn_comment (name
));
4400 query_hsa_grid (stmt
, BRIG_OPCODE_GRIDSIZE
, 0, hbb
);
4402 else if (strcmp (name
, "omp_get_num_teams") == 0)
4403 gen_get_num_teams (stmt
, hbb
);
4404 else if (strcmp (name
, "omp_get_team_num") == 0)
4405 gen_get_team_num (stmt
, hbb
);
4406 else if (strcmp (name
, "omp_get_level") == 0)
4407 gen_get_level (stmt
, hbb
);
4408 else if (strcmp (name
, "omp_get_active_level") == 0)
4409 gen_get_level (stmt
, hbb
);
4410 else if (strcmp (name
, "omp_in_parallel") == 0)
4411 gen_get_level (stmt
, hbb
);
4412 else if (strcmp (name
, "omp_get_max_threads") == 0)
4413 gen_get_max_threads (stmt
, hbb
);
4425 if (strcmp (name
, "__hsa_set_debug_value") == 0)
4428 if (hsa_cfun
->has_shadow_reg_p ())
4430 tree rhs1
= gimple_call_arg (stmt
, 0);
4431 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4433 src
= src
->get_in_type (BRIG_TYPE_U64
, hbb
);
4434 set_debug_value (hbb
, src
);
4443 /* Helper functions to create a single unary HSA operations out of calls to
4444 builtins. OPCODE is the HSA operation to be generated. STMT is a gimple
4445 call to a builtin. HBB is the HSA BB to which the instruction should be
4446 added. Note that nothing will be created if STMT does not have a LHS. */
4449 gen_hsa_unaryop_for_builtin (BrigOpcode opcode
, gimple
*stmt
, hsa_bb
*hbb
)
4451 tree lhs
= gimple_call_lhs (stmt
);
4454 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4455 hsa_op_with_type
*op
4456 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4457 gen_hsa_unary_operation (opcode
, dest
, op
, hbb
);
4460 /* Helper functions to create a call to standard library if LHS of the
4461 STMT is used. HBB is the HSA BB to which the instruction should be
4465 gen_hsa_unaryop_builtin_call (gimple
*stmt
, hsa_bb
*hbb
)
4467 tree lhs
= gimple_call_lhs (stmt
);
4471 if (gimple_call_internal_p (stmt
))
4472 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
4474 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4477 /* Helper functions to create a single unary HSA operations out of calls to
4478 builtins (if unsafe math optimizations are enable). Otherwise, create
4479 a call to standard library function.
4480 OPCODE is the HSA operation to be generated. STMT is a gimple
4481 call to a builtin. HBB is the HSA BB to which the instruction should be
4482 added. Note that nothing will be created if STMT does not have a LHS. */
4485 gen_hsa_unaryop_or_call_for_builtin (BrigOpcode opcode
, gimple
*stmt
,
4488 if (flag_unsafe_math_optimizations
)
4489 gen_hsa_unaryop_for_builtin (opcode
, stmt
, hbb
);
4491 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4494 /* Generate HSA address corresponding to a value VAL (as opposed to a memory
4495 reference tree), for example an SSA_NAME or an ADDR_EXPR. HBB is the HSA BB
4496 to which the instruction should be added. */
4498 static hsa_op_address
*
4499 get_address_from_value (tree val
, hsa_bb
*hbb
)
4501 switch (TREE_CODE (val
))
4505 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4507 = hsa_cfun
->reg_for_gimple_ssa (val
)->get_in_type (addrtype
, hbb
);
4508 return new hsa_op_address (NULL
, as_a
<hsa_op_reg
*> (reg
), 0);
4511 return gen_hsa_addr (TREE_OPERAND (val
, 0), hbb
);
4514 if (tree_fits_shwi_p (val
))
4515 return new hsa_op_address (NULL
, NULL
, tree_to_shwi (val
));
4516 /* Otherwise fall-through */
4519 HSA_SORRY_ATV (EXPR_LOCATION (val
),
4520 "support for HSA does not implement memory access to %E",
4522 return new hsa_op_address (NULL
, NULL
, 0);
4526 /* Expand assignment of a result of a string BUILTIN to DST.
4527 Size of the operation is N bytes, where instructions
4528 will be append to HBB. */
4531 expand_lhs_of_string_op (gimple
*stmt
,
4532 unsigned HOST_WIDE_INT n
, hsa_bb
*hbb
,
4533 enum built_in_function builtin
)
4535 /* If LHS is expected, we need to emit a PHI instruction. */
4536 tree lhs
= gimple_call_lhs (stmt
);
4540 hsa_op_reg
*lhs_reg
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4542 hsa_op_with_type
*dst_reg
4543 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4544 hsa_op_with_type
*tmp
;
4548 case BUILT_IN_MEMPCPY
:
4550 tmp
= new hsa_op_reg (dst_reg
->m_type
);
4552 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
,
4554 new hsa_op_immed (n
, dst_reg
->m_type
));
4555 hbb
->append_insn (add
);
4558 case BUILT_IN_MEMCPY
:
4559 case BUILT_IN_MEMSET
:
4566 hbb
->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV
, lhs_reg
->m_type
,
4570 #define HSA_MEMORY_BUILTINS_LIMIT 128
4572 /* Expand a string builtin (from a gimple STMT) in a way that
4573 according to MISALIGNED_FLAG we process either direct emission
4574 (a bunch of memory load and store instructions), or we emit a function call
4575 of a library function (for instance 'memcpy'). Actually, a basic block
4576 for direct emission is just prepared, where caller is responsible
4577 for emission of corresponding instructions.
4578 All instruction are appended to HBB. */
4581 expand_string_operation_builtin (gimple
*stmt
, hsa_bb
*hbb
,
4582 hsa_op_reg
*misaligned_flag
)
4584 edge e
= split_block (hbb
->m_bb
, stmt
);
4585 basic_block condition_bb
= e
->src
;
4586 hbb
->append_insn (new hsa_insn_br (misaligned_flag
));
4588 /* Prepare the control flow. */
4589 edge condition_edge
= EDGE_SUCC (condition_bb
, 0);
4590 basic_block call_bb
= split_edge (condition_edge
);
4592 basic_block expanded_bb
= split_edge (EDGE_SUCC (call_bb
, 0));
4593 basic_block cont_bb
= EDGE_SUCC (expanded_bb
, 0)->dest
;
4594 basic_block merge_bb
= split_edge (EDGE_PRED (cont_bb
, 0));
4596 condition_edge
->flags
&= ~EDGE_FALLTHRU
;
4597 condition_edge
->flags
|= EDGE_TRUE_VALUE
;
4598 make_edge (condition_bb
, expanded_bb
, EDGE_FALSE_VALUE
);
4600 redirect_edge_succ (EDGE_SUCC (call_bb
, 0), merge_bb
);
4602 hsa_cfun
->m_modified_cfg
= true;
4604 hsa_init_new_bb (expanded_bb
);
4606 /* Slow path: function call. */
4607 gen_hsa_insns_for_direct_call (stmt
, hsa_init_new_bb (call_bb
), false);
4609 return hsa_bb_for_bb (expanded_bb
);
4612 /* Expand a memory copy BUILTIN (BUILT_IN_MEMCPY, BUILT_IN_MEMPCPY) from
4613 a gimple STMT and store all necessary instruction to HBB basic block. */
4616 expand_memory_copy (gimple
*stmt
, hsa_bb
*hbb
, enum built_in_function builtin
)
4618 tree byte_size
= gimple_call_arg (stmt
, 2);
4620 if (!tree_fits_uhwi_p (byte_size
))
4622 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4626 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
4628 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
4630 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4634 tree dst
= gimple_call_arg (stmt
, 0);
4635 tree src
= gimple_call_arg (stmt
, 1);
4637 hsa_op_address
*dst_addr
= get_address_from_value (dst
, hbb
);
4638 hsa_op_address
*src_addr
= get_address_from_value (src
, hbb
);
4640 /* As gen_hsa_memory_copy relies on memory alignment
4641 greater or equal to 8 bytes, we need to verify the alignment. */
4642 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4643 hsa_op_reg
*src_addr_reg
= new hsa_op_reg (addrtype
);
4644 hsa_op_reg
*dst_addr_reg
= new hsa_op_reg (addrtype
);
4646 convert_addr_to_flat_segment (src_addr
, src_addr_reg
, hbb
);
4647 convert_addr_to_flat_segment (dst_addr
, dst_addr_reg
, hbb
);
4649 /* Process BIT OR for source and destination addresses. */
4650 hsa_op_reg
*or_reg
= new hsa_op_reg (addrtype
);
4651 gen_hsa_binary_operation (BRIG_OPCODE_OR
, or_reg
, src_addr_reg
,
4654 /* Process BIT AND with 0x7 to identify the desired alignment
4656 hsa_op_reg
*masked
= new hsa_op_reg (addrtype
);
4658 gen_hsa_binary_operation (BRIG_OPCODE_AND
, masked
, or_reg
,
4659 new hsa_op_immed (7, addrtype
), hbb
);
4661 hsa_op_reg
*misaligned
= new hsa_op_reg (BRIG_TYPE_B1
);
4662 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE
, misaligned
->m_type
,
4664 new hsa_op_immed (0, masked
->m_type
)));
4666 hsa_bb
*native_impl_bb
4667 = expand_string_operation_builtin (stmt
, hbb
, misaligned
);
4669 gen_hsa_memory_copy (native_impl_bb
, dst_addr
, src_addr
, n
, BRIG_ALIGNMENT_8
);
4671 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb
->m_bb
, 0)->dest
);
4672 expand_lhs_of_string_op (stmt
, n
, merge_bb
, builtin
);
4676 /* Expand a memory set BUILTIN (BUILT_IN_MEMSET, BUILT_IN_BZERO) from
4677 a gimple STMT and store all necessary instruction to HBB basic block.
4678 The operation set N bytes with a CONSTANT value. */
4681 expand_memory_set (gimple
*stmt
, unsigned HOST_WIDE_INT n
,
4682 unsigned HOST_WIDE_INT constant
, hsa_bb
*hbb
,
4683 enum built_in_function builtin
)
4685 tree dst
= gimple_call_arg (stmt
, 0);
4686 hsa_op_address
*dst_addr
= get_address_from_value (dst
, hbb
);
4688 /* As gen_hsa_memory_set relies on memory alignment
4689 greater or equal to 8 bytes, we need to verify the alignment. */
4690 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4691 hsa_op_reg
*dst_addr_reg
= new hsa_op_reg (addrtype
);
4692 convert_addr_to_flat_segment (dst_addr
, dst_addr_reg
, hbb
);
4694 /* Process BIT AND with 0x7 to identify the desired alignment
4696 hsa_op_reg
*masked
= new hsa_op_reg (addrtype
);
4698 gen_hsa_binary_operation (BRIG_OPCODE_AND
, masked
, dst_addr_reg
,
4699 new hsa_op_immed (7, addrtype
), hbb
);
4701 hsa_op_reg
*misaligned
= new hsa_op_reg (BRIG_TYPE_B1
);
4702 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE
, misaligned
->m_type
,
4704 new hsa_op_immed (0, masked
->m_type
)));
4706 hsa_bb
*native_impl_bb
4707 = expand_string_operation_builtin (stmt
, hbb
, misaligned
);
4709 gen_hsa_memory_set (native_impl_bb
, dst_addr
, constant
, n
, BRIG_ALIGNMENT_8
);
4711 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb
->m_bb
, 0)->dest
);
4712 expand_lhs_of_string_op (stmt
, n
, merge_bb
, builtin
);
4715 /* Return string for MEMMODEL. */
4718 get_memory_order_name (unsigned memmodel
)
4720 switch (memmodel
& MEMMODEL_BASE_MASK
)
4722 case MEMMODEL_RELAXED
:
4724 case MEMMODEL_CONSUME
:
4726 case MEMMODEL_ACQUIRE
:
4728 case MEMMODEL_RELEASE
:
4730 case MEMMODEL_ACQ_REL
:
4732 case MEMMODEL_SEQ_CST
:
4739 /* Return memory order according to predefined __atomic memory model
4740 constants. LOCATION is provided to locate the problematic statement. */
4742 static BrigMemoryOrder
4743 get_memory_order (unsigned memmodel
, location_t location
)
4745 switch (memmodel
& MEMMODEL_BASE_MASK
)
4747 case MEMMODEL_RELAXED
:
4748 return BRIG_MEMORY_ORDER_RELAXED
;
4749 case MEMMODEL_CONSUME
:
4750 /* HSA does not have an equivalent, but we can use the slightly stronger
4752 case MEMMODEL_ACQUIRE
:
4753 return BRIG_MEMORY_ORDER_SC_ACQUIRE
;
4754 case MEMMODEL_RELEASE
:
4755 return BRIG_MEMORY_ORDER_SC_RELEASE
;
4756 case MEMMODEL_ACQ_REL
:
4757 case MEMMODEL_SEQ_CST
:
4758 /* Callers implementing a simple load or store need to remove the release
4759 or acquire part respectively. */
4760 return BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
4763 const char *mmname
= get_memory_order_name (memmodel
);
4764 HSA_SORRY_ATV (location
,
4765 "support for HSA does not implement the specified "
4766 " memory model%s %s",
4767 mmname
? ": " : "", mmname
? mmname
: "");
4768 return BRIG_MEMORY_ORDER_NONE
;
4773 /* Helper function to create an HSA atomic binary operation instruction out of
4774 calls to atomic builtins. RET_ORIG is true if the built-in is the variant
4775 that return s the value before applying operation, and false if it should
4776 return the value after applying the operation (if it returns value at all).
4777 ACODE is the atomic operation code, STMT is a gimple call to a builtin. HBB
4778 is the HSA BB to which the instruction should be added. */
4781 gen_hsa_ternary_atomic_for_builtin (bool ret_orig
,
4782 enum BrigAtomicOperation acode
,
4786 tree lhs
= gimple_call_lhs (stmt
);
4788 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
4789 BrigType16_t hsa_type
= hsa_type_for_scalar_tree_type (type
, false);
4790 BrigType16_t mtype
= mem_type_for_type (hsa_type
);
4791 tree model
= gimple_call_arg (stmt
, 2);
4793 if (!tree_fits_uhwi_p (model
))
4795 HSA_SORRY_ATV (gimple_location (stmt
),
4796 "support for HSA does not implement memory model %E",
4801 unsigned HOST_WIDE_INT mmodel
= tree_to_uhwi (model
);
4803 BrigMemoryOrder memorder
= get_memory_order (mmodel
, gimple_location (stmt
));
4805 /* Certain atomic insns must have Bx memory types. */
4808 case BRIG_ATOMIC_LD
:
4809 case BRIG_ATOMIC_ST
:
4810 case BRIG_ATOMIC_AND
:
4811 case BRIG_ATOMIC_OR
:
4812 case BRIG_ATOMIC_XOR
:
4813 case BRIG_ATOMIC_EXCH
:
4814 mtype
= hsa_bittype_for_type (mtype
);
4825 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4827 dest
= new hsa_op_reg (hsa_type
);
4828 opcode
= BRIG_OPCODE_ATOMIC
;
4834 opcode
= BRIG_OPCODE_ATOMICNORET
;
4838 if (acode
== BRIG_ATOMIC_ST
)
4840 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
4841 memorder
= BRIG_MEMORY_ORDER_SC_RELEASE
;
4843 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
4844 && memorder
!= BRIG_MEMORY_ORDER_SC_RELEASE
4845 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
4847 HSA_SORRY_ATV (gimple_location (stmt
),
4848 "support for HSA does not implement memory model for "
4849 "ATOMIC_ST: %s", get_memory_order_name (mmodel
));
4854 hsa_insn_atomic
*atominsn
= new hsa_insn_atomic (nops
, opcode
, acode
, mtype
,
4857 hsa_op_address
*addr
;
4858 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
4859 if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
== BRIG_SEGMENT_PRIVATE
)
4861 HSA_SORRY_AT (gimple_location (stmt
),
4862 "HSA does not implement atomic operations in private "
4866 hsa_op_base
*op
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1),
4871 atominsn
->set_op (0, dest
);
4872 atominsn
->set_op (1, addr
);
4873 atominsn
->set_op (2, op
);
4877 atominsn
->set_op (0, addr
);
4878 atominsn
->set_op (1, op
);
4881 hbb
->append_insn (atominsn
);
4883 /* HSA does not natively support the variants that return the modified value,
4884 so re-do the operation again non-atomically if that is what was
4886 if (lhs
&& !ret_orig
)
4891 case BRIG_ATOMIC_ADD
:
4892 arith
= BRIG_OPCODE_ADD
;
4894 case BRIG_ATOMIC_AND
:
4895 arith
= BRIG_OPCODE_AND
;
4897 case BRIG_ATOMIC_OR
:
4898 arith
= BRIG_OPCODE_OR
;
4900 case BRIG_ATOMIC_SUB
:
4901 arith
= BRIG_OPCODE_SUB
;
4903 case BRIG_ATOMIC_XOR
:
4904 arith
= BRIG_OPCODE_XOR
;
4909 hsa_op_reg
*real_dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4910 gen_hsa_binary_operation (arith
, real_dest
, dest
, op
, hbb
);
4914 /* Generate HSA instructions for an internal fn.
4915 Instructions will be appended to HBB, which also needs to be the
4916 corresponding structure to the basic_block of STMT. */
4919 gen_hsa_insn_for_internal_fn_call (gcall
*stmt
, hsa_bb
*hbb
)
4921 gcc_checking_assert (gimple_call_internal_fn (stmt
));
4922 internal_fn fn
= gimple_call_internal_fn (stmt
);
4924 bool is_float_type_p
= false;
4925 if (gimple_call_lhs (stmt
) != NULL
4926 && TREE_TYPE (gimple_call_lhs (stmt
)) == float_type_node
)
4927 is_float_type_p
= true;
4932 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
4936 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
4940 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
4944 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
4948 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
4953 if (is_float_type_p
)
4954 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
4956 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4962 if (is_float_type_p
)
4963 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
4965 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4972 if (is_float_type_p
)
4973 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
4975 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4982 if (is_float_type_p
)
4983 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
4985 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4990 gen_hsa_clrsb (stmt
, hbb
);
4994 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
4998 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
5002 gen_hsa_ffs (stmt
, hbb
);
5006 gen_hsa_parity (stmt
, hbb
);
5010 gen_hsa_popcount (stmt
, hbb
);
5023 case IFN_SIGNIFICAND
:
5035 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
5038 HSA_SORRY_ATV (gimple_location (stmt
),
5039 "support for HSA does not implement internal function: %s",
5040 internal_fn_name (fn
));
5045 /* Generate HSA instructions for the given call statement STMT. Instructions
5046 will be appended to HBB. */
5049 gen_hsa_insns_for_call (gimple
*stmt
, hsa_bb
*hbb
)
5051 gcall
*call
= as_a
<gcall
*> (stmt
);
5052 tree lhs
= gimple_call_lhs (stmt
);
5055 if (gimple_call_internal_p (stmt
))
5057 gen_hsa_insn_for_internal_fn_call (call
, hbb
);
5061 if (!gimple_call_builtin_p (stmt
, BUILT_IN_NORMAL
))
5063 tree function_decl
= gimple_call_fndecl (stmt
);
5064 if (function_decl
== NULL_TREE
)
5066 HSA_SORRY_AT (gimple_location (stmt
),
5067 "support for HSA does not implement indirect calls");
5071 if (hsa_callable_function_p (function_decl
))
5072 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5073 else if (!gen_hsa_insns_for_known_library_call (stmt
, hbb
))
5074 HSA_SORRY_AT (gimple_location (stmt
),
5075 "HSA supports only calls of functions marked with pragma "
5076 "omp declare target");
5080 tree fndecl
= gimple_call_fndecl (stmt
);
5081 enum built_in_function builtin
= DECL_FUNCTION_CODE (fndecl
);
5085 case BUILT_IN_FABSF
:
5086 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_ABS
, stmt
, hbb
);
5090 case BUILT_IN_CEILF
:
5091 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
5094 case BUILT_IN_FLOOR
:
5095 case BUILT_IN_FLOORF
:
5096 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
5100 case BUILT_IN_RINTF
:
5101 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
5105 case BUILT_IN_SQRTF
:
5106 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
5109 case BUILT_IN_TRUNC
:
5110 case BUILT_IN_TRUNCF
:
5111 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
5118 /* HSAIL does not provide an instruction for double argument type. */
5119 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5123 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
5126 case BUILT_IN_EXP2F
:
5127 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
5130 case BUILT_IN_LOG2F
:
5131 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
5135 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
5138 case BUILT_IN_CLRSB
:
5139 case BUILT_IN_CLRSBL
:
5140 case BUILT_IN_CLRSBLL
:
5141 gen_hsa_clrsb (call
, hbb
);
5146 case BUILT_IN_CLZLL
:
5147 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
5152 case BUILT_IN_CTZLL
:
5153 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
5158 case BUILT_IN_FFSLL
:
5159 gen_hsa_ffs (call
, hbb
);
5162 case BUILT_IN_PARITY
:
5163 case BUILT_IN_PARITYL
:
5164 case BUILT_IN_PARITYLL
:
5165 gen_hsa_parity (call
, hbb
);
5168 case BUILT_IN_POPCOUNT
:
5169 case BUILT_IN_POPCOUNTL
:
5170 case BUILT_IN_POPCOUNTLL
:
5171 gen_hsa_popcount (call
, hbb
);
5174 case BUILT_IN_ATOMIC_LOAD_1
:
5175 case BUILT_IN_ATOMIC_LOAD_2
:
5176 case BUILT_IN_ATOMIC_LOAD_4
:
5177 case BUILT_IN_ATOMIC_LOAD_8
:
5178 case BUILT_IN_ATOMIC_LOAD_16
:
5181 hsa_op_address
*addr
;
5182 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5183 tree model
= gimple_call_arg (stmt
, 1);
5184 if (!tree_fits_uhwi_p (model
))
5186 HSA_SORRY_ATV (gimple_location (stmt
),
5187 "support for HSA does not implement "
5193 unsigned HOST_WIDE_INT mmodel
= tree_to_uhwi (model
);
5194 BrigMemoryOrder memorder
= get_memory_order (mmodel
,
5195 gimple_location (stmt
));
5197 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
5198 memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
5200 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
5201 && memorder
!= BRIG_MEMORY_ORDER_SC_ACQUIRE
5202 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
5204 HSA_SORRY_ATV (gimple_location (stmt
),
5205 "support for HSA does not implement "
5206 "memory model for ATOMIC_LD: %s",
5207 get_memory_order_name (mmodel
));
5213 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
5215 mtype
= mem_type_for_type (t
);
5216 mtype
= hsa_bittype_for_type (mtype
);
5217 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5221 mtype
= BRIG_TYPE_B64
;
5222 dest
= new hsa_op_reg (mtype
);
5225 hsa_insn_atomic
*atominsn
5226 = new hsa_insn_atomic (2, BRIG_OPCODE_ATOMIC
, BRIG_ATOMIC_LD
, mtype
,
5227 memorder
, dest
, addr
);
5229 hbb
->append_insn (atominsn
);
5233 case BUILT_IN_ATOMIC_EXCHANGE_1
:
5234 case BUILT_IN_ATOMIC_EXCHANGE_2
:
5235 case BUILT_IN_ATOMIC_EXCHANGE_4
:
5236 case BUILT_IN_ATOMIC_EXCHANGE_8
:
5237 case BUILT_IN_ATOMIC_EXCHANGE_16
:
5238 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_EXCH
, stmt
, hbb
);
5241 case BUILT_IN_ATOMIC_FETCH_ADD_1
:
5242 case BUILT_IN_ATOMIC_FETCH_ADD_2
:
5243 case BUILT_IN_ATOMIC_FETCH_ADD_4
:
5244 case BUILT_IN_ATOMIC_FETCH_ADD_8
:
5245 case BUILT_IN_ATOMIC_FETCH_ADD_16
:
5246 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_ADD
, stmt
, hbb
);
5249 case BUILT_IN_ATOMIC_FETCH_SUB_1
:
5250 case BUILT_IN_ATOMIC_FETCH_SUB_2
:
5251 case BUILT_IN_ATOMIC_FETCH_SUB_4
:
5252 case BUILT_IN_ATOMIC_FETCH_SUB_8
:
5253 case BUILT_IN_ATOMIC_FETCH_SUB_16
:
5254 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_SUB
, stmt
, hbb
);
5257 case BUILT_IN_ATOMIC_FETCH_AND_1
:
5258 case BUILT_IN_ATOMIC_FETCH_AND_2
:
5259 case BUILT_IN_ATOMIC_FETCH_AND_4
:
5260 case BUILT_IN_ATOMIC_FETCH_AND_8
:
5261 case BUILT_IN_ATOMIC_FETCH_AND_16
:
5262 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_AND
, stmt
, hbb
);
5265 case BUILT_IN_ATOMIC_FETCH_XOR_1
:
5266 case BUILT_IN_ATOMIC_FETCH_XOR_2
:
5267 case BUILT_IN_ATOMIC_FETCH_XOR_4
:
5268 case BUILT_IN_ATOMIC_FETCH_XOR_8
:
5269 case BUILT_IN_ATOMIC_FETCH_XOR_16
:
5270 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_XOR
, stmt
, hbb
);
5273 case BUILT_IN_ATOMIC_FETCH_OR_1
:
5274 case BUILT_IN_ATOMIC_FETCH_OR_2
:
5275 case BUILT_IN_ATOMIC_FETCH_OR_4
:
5276 case BUILT_IN_ATOMIC_FETCH_OR_8
:
5277 case BUILT_IN_ATOMIC_FETCH_OR_16
:
5278 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_OR
, stmt
, hbb
);
5281 case BUILT_IN_ATOMIC_STORE_1
:
5282 case BUILT_IN_ATOMIC_STORE_2
:
5283 case BUILT_IN_ATOMIC_STORE_4
:
5284 case BUILT_IN_ATOMIC_STORE_8
:
5285 case BUILT_IN_ATOMIC_STORE_16
:
5286 /* Since there cannot be any LHS, the first parameter is meaningless. */
5287 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_ST
, stmt
, hbb
);
5290 case BUILT_IN_ATOMIC_ADD_FETCH_1
:
5291 case BUILT_IN_ATOMIC_ADD_FETCH_2
:
5292 case BUILT_IN_ATOMIC_ADD_FETCH_4
:
5293 case BUILT_IN_ATOMIC_ADD_FETCH_8
:
5294 case BUILT_IN_ATOMIC_ADD_FETCH_16
:
5295 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_ADD
, stmt
, hbb
);
5298 case BUILT_IN_ATOMIC_SUB_FETCH_1
:
5299 case BUILT_IN_ATOMIC_SUB_FETCH_2
:
5300 case BUILT_IN_ATOMIC_SUB_FETCH_4
:
5301 case BUILT_IN_ATOMIC_SUB_FETCH_8
:
5302 case BUILT_IN_ATOMIC_SUB_FETCH_16
:
5303 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_SUB
, stmt
, hbb
);
5306 case BUILT_IN_ATOMIC_AND_FETCH_1
:
5307 case BUILT_IN_ATOMIC_AND_FETCH_2
:
5308 case BUILT_IN_ATOMIC_AND_FETCH_4
:
5309 case BUILT_IN_ATOMIC_AND_FETCH_8
:
5310 case BUILT_IN_ATOMIC_AND_FETCH_16
:
5311 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_AND
, stmt
, hbb
);
5314 case BUILT_IN_ATOMIC_XOR_FETCH_1
:
5315 case BUILT_IN_ATOMIC_XOR_FETCH_2
:
5316 case BUILT_IN_ATOMIC_XOR_FETCH_4
:
5317 case BUILT_IN_ATOMIC_XOR_FETCH_8
:
5318 case BUILT_IN_ATOMIC_XOR_FETCH_16
:
5319 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_XOR
, stmt
, hbb
);
5322 case BUILT_IN_ATOMIC_OR_FETCH_1
:
5323 case BUILT_IN_ATOMIC_OR_FETCH_2
:
5324 case BUILT_IN_ATOMIC_OR_FETCH_4
:
5325 case BUILT_IN_ATOMIC_OR_FETCH_8
:
5326 case BUILT_IN_ATOMIC_OR_FETCH_16
:
5327 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_OR
, stmt
, hbb
);
5330 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_1
:
5331 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_2
:
5332 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_4
:
5333 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8
:
5334 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_16
:
5336 /* TODO: Use the appropriate memory model for now. */
5337 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
5340 = hsa_bittype_for_type (hsa_type_for_scalar_tree_type (type
, false));
5342 hsa_insn_atomic
*atominsn
5343 = new hsa_insn_atomic (4, BRIG_OPCODE_ATOMIC
, BRIG_ATOMIC_CAS
, atype
,
5344 BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
);
5345 hsa_op_address
*addr
;
5346 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5349 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5351 dest
= new hsa_op_reg (atype
);
5353 /* Should check what the memory scope is. */
5354 atominsn
->m_memoryscope
= BRIG_MEMORY_SCOPE_WORKGROUP
;
5355 atominsn
->set_op (0, dest
);
5356 atominsn
->set_op (1, addr
);
5358 hsa_op_with_type
*op
5359 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1), hbb
);
5360 atominsn
->set_op (2, op
);
5361 op
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 2), hbb
);
5362 atominsn
->set_op (3, op
);
5364 hbb
->append_insn (atominsn
);
5367 case BUILT_IN_GOMP_PARALLEL
:
5368 HSA_SORRY_AT (gimple_location (stmt
),
5369 "support for HSA does not implement non-gridified "
5370 "OpenMP parallel constructs.");
5372 case BUILT_IN_OMP_GET_THREAD_NUM
:
5374 query_hsa_grid (stmt
, BRIG_OPCODE_WORKITEMABSID
, 0, hbb
);
5378 case BUILT_IN_OMP_GET_NUM_THREADS
:
5380 query_hsa_grid (stmt
, BRIG_OPCODE_GRIDSIZE
, 0, hbb
);
5383 case BUILT_IN_GOMP_TEAMS
:
5385 gen_set_num_threads (gimple_call_arg (stmt
, 1), hbb
);
5388 case BUILT_IN_OMP_GET_NUM_TEAMS
:
5390 gen_get_num_teams (stmt
, hbb
);
5393 case BUILT_IN_OMP_GET_TEAM_NUM
:
5395 gen_get_team_num (stmt
, hbb
);
5398 case BUILT_IN_MEMCPY
:
5399 case BUILT_IN_MEMPCPY
:
5401 expand_memory_copy (stmt
, hbb
, builtin
);
5404 case BUILT_IN_MEMSET
:
5406 tree c
= gimple_call_arg (stmt
, 1);
5408 if (TREE_CODE (c
) != INTEGER_CST
)
5410 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5414 tree byte_size
= gimple_call_arg (stmt
, 2);
5416 if (!tree_fits_uhwi_p (byte_size
))
5418 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5422 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
5424 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5426 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5430 unsigned HOST_WIDE_INT constant
5431 = tree_to_uhwi (fold_convert (unsigned_char_type_node
, c
));
5433 expand_memory_set (stmt
, n
, constant
, hbb
, builtin
);
5437 case BUILT_IN_BZERO
:
5439 tree byte_size
= gimple_call_arg (stmt
, 1);
5441 if (!tree_fits_uhwi_p (byte_size
))
5443 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5447 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
5449 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5451 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5455 expand_memory_set (stmt
, n
, 0, hbb
, builtin
);
5459 case BUILT_IN_ALLOCA
:
5460 case BUILT_IN_ALLOCA_WITH_ALIGN
:
5462 gen_hsa_alloca (call
, hbb
);
5467 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5473 /* Generate HSA instructions for a given gimple statement. Instructions will be
5477 gen_hsa_insns_for_gimple_stmt (gimple
*stmt
, hsa_bb
*hbb
)
5479 switch (gimple_code (stmt
))
5482 if (gimple_clobber_p (stmt
))
5485 if (gimple_assign_single_p (stmt
))
5487 tree lhs
= gimple_assign_lhs (stmt
);
5488 tree rhs
= gimple_assign_rhs1 (stmt
);
5489 gen_hsa_insns_for_single_assignment (lhs
, rhs
, hbb
);
5492 gen_hsa_insns_for_operation_assignment (stmt
, hbb
);
5495 gen_hsa_insns_for_return (as_a
<greturn
*> (stmt
), hbb
);
5498 gen_hsa_insns_for_cond_stmt (stmt
, hbb
);
5501 gen_hsa_insns_for_call (stmt
, hbb
);
5504 /* ??? HSA supports some debug facilities. */
5508 tree label
= gimple_label_label (as_a
<glabel
*> (stmt
));
5509 if (FORCED_LABEL (label
))
5510 HSA_SORRY_AT (gimple_location (stmt
),
5511 "support for HSA does not implement gimple label with "
5518 hbb
->append_insn (new hsa_insn_basic (0, BRIG_OPCODE_NOP
));
5523 gen_hsa_insns_for_switch_stmt (as_a
<gswitch
*> (stmt
), hbb
);
5527 HSA_SORRY_ATV (gimple_location (stmt
),
5528 "support for HSA does not implement gimple statement %s",
5529 gimple_code_name
[(int) gimple_code (stmt
)]);
5533 /* Generate a HSA PHI from a gimple PHI. */
5536 gen_hsa_phi_from_gimple_phi (gimple
*phi_stmt
, hsa_bb
*hbb
)
5539 unsigned count
= gimple_phi_num_args (phi_stmt
);
5542 = hsa_cfun
->reg_for_gimple_ssa (gimple_phi_result (phi_stmt
));
5543 hphi
= new hsa_insn_phi (count
, dest
);
5544 hphi
->m_bb
= hbb
->m_bb
;
5546 tree lhs
= gimple_phi_result (phi_stmt
);
5548 for (unsigned i
= 0; i
< count
; i
++)
5550 tree op
= gimple_phi_arg_def (phi_stmt
, i
);
5552 if (TREE_CODE (op
) == SSA_NAME
)
5554 hsa_op_reg
*hreg
= hsa_cfun
->reg_for_gimple_ssa (op
);
5555 hphi
->set_op (i
, hreg
);
5559 gcc_assert (is_gimple_min_invariant (op
));
5560 tree t
= TREE_TYPE (op
);
5561 if (!POINTER_TYPE_P (t
)
5562 || (TREE_CODE (op
) == STRING_CST
5563 && TREE_CODE (TREE_TYPE (t
)) == INTEGER_TYPE
))
5564 hphi
->set_op (i
, new hsa_op_immed (op
));
5565 else if (POINTER_TYPE_P (TREE_TYPE (lhs
))
5566 && TREE_CODE (op
) == INTEGER_CST
)
5568 /* Handle assignment of NULL value to a pointer type. */
5569 hphi
->set_op (i
, new hsa_op_immed (op
));
5571 else if (TREE_CODE (op
) == ADDR_EXPR
)
5573 edge e
= gimple_phi_arg_edge (as_a
<gphi
*> (phi_stmt
), i
);
5574 hsa_bb
*hbb_src
= hsa_init_new_bb (split_edge (e
));
5575 hsa_op_address
*addr
= gen_hsa_addr (TREE_OPERAND (op
, 0),
5579 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
5580 hsa_insn_basic
*insn
5581 = new hsa_insn_basic (2, BRIG_OPCODE_LDA
, BRIG_TYPE_U64
,
5583 hbb_src
->append_insn (insn
);
5585 hphi
->set_op (i
, dest
);
5589 HSA_SORRY_AT (gimple_location (phi_stmt
),
5590 "support for HSA does not handle PHI nodes with "
5591 "constant address operands");
5597 hphi
->m_prev
= hbb
->m_last_phi
;
5598 hphi
->m_next
= NULL
;
5599 if (hbb
->m_last_phi
)
5600 hbb
->m_last_phi
->m_next
= hphi
;
5601 hbb
->m_last_phi
= hphi
;
5602 if (!hbb
->m_first_phi
)
5603 hbb
->m_first_phi
= hphi
;
5606 /* Constructor of class containing HSA-specific information about a basic
5607 block. CFG_BB is the CFG BB this HSA BB is associated with. IDX is the new
5608 index of this BB (so that the constructor does not attempt to use
5609 hsa_cfun during its construction). */
5611 hsa_bb::hsa_bb (basic_block cfg_bb
, int idx
)
5612 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5613 m_last_phi (NULL
), m_index (idx
), m_liveout (BITMAP_ALLOC (NULL
)),
5614 m_livein (BITMAP_ALLOC (NULL
))
5616 gcc_assert (!cfg_bb
->aux
);
5620 /* Constructor of class containing HSA-specific information about a basic
5621 block. CFG_BB is the CFG BB this HSA BB is associated with. */
5623 hsa_bb::hsa_bb (basic_block cfg_bb
)
5624 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5625 m_last_phi (NULL
), m_index (hsa_cfun
->m_hbb_count
++),
5626 m_liveout (BITMAP_ALLOC (NULL
)), m_livein (BITMAP_ALLOC (NULL
))
5628 gcc_assert (!cfg_bb
->aux
);
5632 /* Destructor of class representing HSA BB. */
5636 BITMAP_FREE (m_livein
);
5637 BITMAP_FREE (m_liveout
);
5640 /* Create and initialize and return a new hsa_bb structure for a given CFG
5644 hsa_init_new_bb (basic_block bb
)
5646 return new (*hsa_allocp_bb
) hsa_bb (bb
);
5649 /* Initialize OMP in an HSA basic block PROLOGUE. */
5652 init_prologue (void)
5654 if (!hsa_cfun
->m_kern_p
)
5657 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5659 /* Create a magic number that is going to be printed by libgomp. */
5660 unsigned index
= hsa_get_number_decl_kernel_mappings ();
5662 /* Emit store to debug argument. */
5663 if (PARAM_VALUE (PARAM_HSA_GEN_DEBUG_STORES
) > 0)
5664 set_debug_value (prologue
, new hsa_op_immed (1000 + index
, BRIG_TYPE_U64
));
5667 /* Initialize hsa_num_threads to a default value. */
5670 init_hsa_num_threads (void)
5672 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5674 /* Save the default value to private variable hsa_num_threads. */
5675 hsa_insn_basic
*basic
5676 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
,
5677 new hsa_op_immed (0, hsa_num_threads
->m_type
),
5678 new hsa_op_address (hsa_num_threads
));
5679 prologue
->append_insn (basic
);
5682 /* Go over gimple representation and generate our internal HSA one. */
5685 gen_body_from_gimple ()
5689 /* Verify CFG for complex edges we are unable to handle. */
5693 FOR_EACH_BB_FN (bb
, cfun
)
5695 FOR_EACH_EDGE (e
, ei
, bb
->succs
)
5697 /* Verify all unsupported flags for edges that point
5698 to the same basic block. */
5699 if (e
->flags
& EDGE_EH
)
5701 HSA_SORRY_AT (UNKNOWN_LOCATION
,
5702 "support for HSA does not implement exception "
5709 FOR_EACH_BB_FN (bb
, cfun
)
5711 gimple_stmt_iterator gsi
;
5712 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5716 hbb
= hsa_init_new_bb (bb
);
5718 for (gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5720 gen_hsa_insns_for_gimple_stmt (gsi_stmt (gsi
), hbb
);
5721 if (hsa_seen_error ())
5726 FOR_EACH_BB_FN (bb
, cfun
)
5728 gimple_stmt_iterator gsi
;
5729 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5730 gcc_assert (hbb
!= NULL
);
5732 for (gsi
= gsi_start_phis (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5733 if (!virtual_operand_p (gimple_phi_result (gsi_stmt (gsi
))))
5734 gen_hsa_phi_from_gimple_phi (gsi_stmt (gsi
), hbb
);
5737 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
5739 fprintf (dump_file
, "------- Generated SSA form -------\n");
5740 dump_hsa_cfun (dump_file
);
5745 gen_function_decl_parameters (hsa_function_representation
*f
,
5751 for (parm
= TYPE_ARG_TYPES (TREE_TYPE (decl
)), i
= 0;
5753 parm
= TREE_CHAIN (parm
), i
++)
5755 /* Result type if last in the tree list. */
5756 if (TREE_CHAIN (parm
) == NULL
)
5759 tree v
= TREE_VALUE (parm
);
5761 hsa_symbol
*arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5763 arg
->m_type
= hsa_type_for_tree_type (v
, &arg
->m_dim
);
5764 arg
->m_name_number
= i
;
5766 f
->m_input_args
.safe_push (arg
);
5769 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
5770 if (!VOID_TYPE_P (result_type
))
5772 f
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5774 f
->m_output_arg
->m_type
5775 = hsa_type_for_tree_type (result_type
, &f
->m_output_arg
->m_dim
);
5776 f
->m_output_arg
->m_name
= "res";
5780 /* Generate the vector of parameters of the HSA representation of the current
5781 function. This also includes the output parameter representing the
5785 gen_function_def_parameters ()
5789 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5791 for (parm
= DECL_ARGUMENTS (cfun
->decl
); parm
;
5792 parm
= DECL_CHAIN (parm
))
5794 struct hsa_symbol
**slot
;
5797 = new hsa_symbol (BRIG_TYPE_NONE
, hsa_cfun
->m_kern_p
5798 ? BRIG_SEGMENT_KERNARG
: BRIG_SEGMENT_ARG
,
5799 BRIG_LINKAGE_FUNCTION
);
5800 arg
->fillup_for_decl (parm
);
5802 hsa_cfun
->m_input_args
.safe_push (arg
);
5804 if (hsa_seen_error ())
5807 arg
->m_name
= hsa_get_declaration_name (parm
);
5809 /* Copy all input arguments and create corresponding private symbols
5811 hsa_symbol
*private_arg
;
5812 hsa_op_address
*parm_addr
= new hsa_op_address (arg
);
5814 if (TREE_ADDRESSABLE (parm
)
5815 || (!is_gimple_reg (parm
) && !TREE_READONLY (parm
)))
5817 private_arg
= hsa_cfun
->create_hsa_temporary (arg
->m_type
);
5818 private_arg
->fillup_for_decl (parm
);
5820 BrigAlignment8_t align
= MIN (arg
->m_align
, private_arg
->m_align
);
5822 hsa_op_address
*private_arg_addr
= new hsa_op_address (private_arg
);
5823 gen_hsa_memory_copy (prologue
, private_arg_addr
, parm_addr
,
5824 arg
->total_byte_size (), align
);
5829 slot
= hsa_cfun
->m_local_symbols
->find_slot (private_arg
, INSERT
);
5830 gcc_assert (!*slot
);
5831 *slot
= private_arg
;
5833 if (is_gimple_reg (parm
))
5835 tree ddef
= ssa_default_def (cfun
, parm
);
5836 if (ddef
&& !has_zero_uses (ddef
))
5838 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (ddef
),
5840 BrigType16_t mtype
= mem_type_for_type (t
);
5841 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (ddef
);
5842 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
,
5844 gcc_assert (!parm_addr
->m_reg
);
5845 prologue
->append_insn (mem
);
5850 if (!VOID_TYPE_P (TREE_TYPE (TREE_TYPE (cfun
->decl
))))
5852 struct hsa_symbol
**slot
;
5854 hsa_cfun
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5855 BRIG_LINKAGE_FUNCTION
);
5856 hsa_cfun
->m_output_arg
->fillup_for_decl (DECL_RESULT (cfun
->decl
));
5858 if (hsa_seen_error ())
5861 hsa_cfun
->m_output_arg
->m_name
= "res";
5862 slot
= hsa_cfun
->m_local_symbols
->find_slot (hsa_cfun
->m_output_arg
,
5864 gcc_assert (!*slot
);
5865 *slot
= hsa_cfun
->m_output_arg
;
5869 /* Generate function representation that corresponds to
5870 a function declaration. */
5872 hsa_function_representation
*
5873 hsa_generate_function_declaration (tree decl
)
5875 hsa_function_representation
*fun
5876 = new hsa_function_representation (decl
, false, 0);
5878 fun
->m_declaration_p
= true;
5879 fun
->m_name
= get_brig_function_name (decl
);
5880 gen_function_decl_parameters (fun
, decl
);
5886 /* Generate function representation that corresponds to
5889 hsa_function_representation
*
5890 hsa_generate_internal_fn_decl (hsa_internal_fn
*fn
)
5892 hsa_function_representation
*fun
= new hsa_function_representation (fn
);
5894 fun
->m_name
= fn
->name ();
5896 for (unsigned i
= 0; i
< fn
->get_arity (); i
++)
5899 = new hsa_symbol (fn
->get_argument_type (i
), BRIG_SEGMENT_ARG
,
5901 arg
->m_name_number
= i
;
5902 fun
->m_input_args
.safe_push (arg
);
5905 fun
->m_output_arg
= new hsa_symbol (fn
->get_argument_type (-1),
5906 BRIG_SEGMENT_ARG
, BRIG_LINKAGE_NONE
);
5907 fun
->m_output_arg
->m_name
= "res";
5912 /* Return true if switch statement S can be transformed
5913 to a SBR instruction in HSAIL. */
5916 transformable_switch_to_sbr_p (gswitch
*s
)
5918 /* Identify if a switch statement can be transformed to
5919 SBR instruction, like:
5921 sbr_u32 $s1 [@label1, @label2, @label3];
5924 tree size
= get_switch_size (s
);
5925 if (!tree_fits_uhwi_p (size
))
5928 if (tree_to_uhwi (size
) > HSA_MAXIMUM_SBR_LABELS
)
5934 /* Structure hold connection between PHI nodes and immediate
5935 values hold by there nodes. */
5937 struct phi_definition
5939 phi_definition (unsigned phi_i
, unsigned label_i
, tree imm
):
5940 phi_index (phi_i
), label_index (label_i
), phi_value (imm
)
5944 unsigned label_index
;
5948 /* Sum slice of a vector V, starting from index START and ending
5949 at the index END - 1. */
5951 template <typename T
>
5953 T
sum_slice (const auto_vec
<T
> &v
, unsigned start
, unsigned end
)
5957 for (unsigned i
= start
; i
< end
; i
++)
5963 /* Function transforms GIMPLE SWITCH statements to a series of IF statements.
5964 Let's assume following example:
5978 The transformation encompasses following steps:
5979 1) all immediate values used by edges coming from the switch basic block
5981 2) all these edges are removed
5982 3) the switch statement (in L0) is replaced by:
5988 4) newly created basic block Lx' is used for generation of
5990 5) else branch of the last condition goes to LD
5991 6) fix all immediate values in PHI nodes that were propagated though
5992 edges that were removed in step 2
5994 Note: if a case is made by a range C1..C2, then process
5995 following transformation:
5997 switch_cond_op1 = C1 <= index;
5998 switch_cond_op2 = index <= C2;
5999 switch_cond_and = switch_cond_op1 & switch_cond_op2;
6000 if (switch_cond_and != 0)
6008 convert_switch_statements (void)
6010 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
6013 bool modified_cfg
= false;
6015 FOR_EACH_BB_FN (bb
, func
)
6017 gimple_stmt_iterator gsi
= gsi_last_bb (bb
);
6018 if (gsi_end_p (gsi
))
6021 gimple
*stmt
= gsi_stmt (gsi
);
6023 if (gimple_code (stmt
) == GIMPLE_SWITCH
)
6025 gswitch
*s
= as_a
<gswitch
*> (stmt
);
6027 /* If the switch can utilize SBR insn, skip the statement. */
6028 if (transformable_switch_to_sbr_p (s
))
6031 modified_cfg
= true;
6033 unsigned labels
= gimple_switch_num_labels (s
);
6034 tree index
= gimple_switch_index (s
);
6035 tree index_type
= TREE_TYPE (index
);
6036 tree default_label
= gimple_switch_default_label (s
);
6037 basic_block default_label_bb
6038 = label_to_block_fn (func
, CASE_LABEL (default_label
));
6039 basic_block cur_bb
= bb
;
6041 auto_vec
<edge
> new_edges
;
6042 auto_vec
<phi_definition
*> phi_todo_list
;
6043 auto_vec
<gcov_type
> edge_counts
;
6044 auto_vec
<int> edge_probabilities
;
6046 /* Investigate all labels that and PHI nodes in these edges which
6047 should be fixed after we add new collection of edges. */
6048 for (unsigned i
= 0; i
< labels
; i
++)
6050 tree label
= gimple_switch_label (s
, i
);
6051 basic_block label_bb
= label_to_block_fn (func
, CASE_LABEL (label
));
6052 edge e
= find_edge (bb
, label_bb
);
6053 edge_counts
.safe_push (e
->count
);
6054 edge_probabilities
.safe_push (e
->probability
);
6055 gphi_iterator phi_gsi
;
6057 /* Save PHI definitions that will be destroyed because of an edge
6058 is going to be removed. */
6059 unsigned phi_index
= 0;
6060 for (phi_gsi
= gsi_start_phis (e
->dest
);
6061 !gsi_end_p (phi_gsi
); gsi_next (&phi_gsi
))
6063 gphi
*phi
= phi_gsi
.phi ();
6064 for (unsigned j
= 0; j
< gimple_phi_num_args (phi
); j
++)
6066 if (gimple_phi_arg_edge (phi
, j
) == e
)
6068 tree imm
= gimple_phi_arg_def (phi
, j
);
6069 phi_definition
*p
= new phi_definition (phi_index
, i
,
6071 phi_todo_list
.safe_push (p
);
6079 /* Remove all edges for the current basic block. */
6080 for (int i
= EDGE_COUNT (bb
->succs
) - 1; i
>= 0; i
--)
6082 edge e
= EDGE_SUCC (bb
, i
);
6086 /* Iterate all non-default labels. */
6087 for (unsigned i
= 1; i
< labels
; i
++)
6089 tree label
= gimple_switch_label (s
, i
);
6090 tree low
= CASE_LOW (label
);
6091 tree high
= CASE_HIGH (label
);
6093 if (!useless_type_conversion_p (TREE_TYPE (low
), index_type
))
6094 low
= fold_convert (index_type
, low
);
6096 gimple_stmt_iterator cond_gsi
= gsi_last_bb (cur_bb
);
6100 tree tmp1
= make_temp_ssa_name (boolean_type_node
, NULL
,
6103 gimple
*assign1
= gimple_build_assign (tmp1
, LE_EXPR
, low
,
6106 tree tmp2
= make_temp_ssa_name (boolean_type_node
, NULL
,
6109 if (!useless_type_conversion_p (TREE_TYPE (high
), index_type
))
6110 high
= fold_convert (index_type
, high
);
6111 gimple
*assign2
= gimple_build_assign (tmp2
, LE_EXPR
, index
,
6114 tree tmp3
= make_temp_ssa_name (boolean_type_node
, NULL
,
6116 gimple
*assign3
= gimple_build_assign (tmp3
, BIT_AND_EXPR
, tmp1
,
6119 gsi_insert_before (&cond_gsi
, assign1
, GSI_SAME_STMT
);
6120 gsi_insert_before (&cond_gsi
, assign2
, GSI_SAME_STMT
);
6121 gsi_insert_before (&cond_gsi
, assign3
, GSI_SAME_STMT
);
6123 tree b
= constant_boolean_node (false, boolean_type_node
);
6124 c
= gimple_build_cond (NE_EXPR
, tmp3
, b
, NULL
, NULL
);
6127 c
= gimple_build_cond (EQ_EXPR
, index
, low
, NULL
, NULL
);
6129 gimple_set_location (c
, gimple_location (stmt
));
6131 gsi_insert_before (&cond_gsi
, c
, GSI_SAME_STMT
);
6133 basic_block label_bb
6134 = label_to_block_fn (func
, CASE_LABEL (label
));
6135 edge new_edge
= make_edge (cur_bb
, label_bb
, EDGE_TRUE_VALUE
);
6136 int prob_sum
= sum_slice
<int> (edge_probabilities
, i
, labels
) +
6137 edge_probabilities
[0];
6140 new_edge
->probability
6141 = RDIV (REG_BR_PROB_BASE
* edge_probabilities
[i
], prob_sum
);
6143 new_edge
->count
= edge_counts
[i
];
6144 new_edges
.safe_push (new_edge
);
6148 /* Prepare another basic block that will contain
6150 basic_block next_bb
= create_empty_bb (cur_bb
);
6153 add_bb_to_loop (next_bb
, cur_bb
->loop_father
);
6154 loops_state_set (LOOPS_NEED_FIXUP
);
6157 edge next_edge
= make_edge (cur_bb
, next_bb
, EDGE_FALSE_VALUE
);
6158 next_edge
->probability
6159 = inverse_probability (new_edge
->probability
);
6160 next_edge
->count
= edge_counts
[0]
6161 + sum_slice
<gcov_type
> (edge_counts
, i
, labels
);
6162 next_bb
->frequency
= EDGE_FREQUENCY (next_edge
);
6165 else /* Link last IF statement and default label
6168 edge e
= make_edge (cur_bb
, default_label_bb
, EDGE_FALSE_VALUE
);
6169 e
->probability
= inverse_probability (new_edge
->probability
);
6170 e
->count
= edge_counts
[0];
6171 new_edges
.safe_insert (0, e
);
6175 /* Restore original PHI immediate value. */
6176 for (unsigned i
= 0; i
< phi_todo_list
.length (); i
++)
6178 phi_definition
*phi_def
= phi_todo_list
[i
];
6179 edge new_edge
= new_edges
[phi_def
->label_index
];
6181 gphi_iterator it
= gsi_start_phis (new_edge
->dest
);
6182 for (unsigned i
= 0; i
< phi_def
->phi_index
; i
++)
6185 gphi
*phi
= it
.phi ();
6186 add_phi_arg (phi
, phi_def
->phi_value
, new_edge
, UNKNOWN_LOCATION
);
6190 /* Remove the original GIMPLE switch statement. */
6191 gsi_remove (&gsi
, true);
6196 dump_function_to_file (current_function_decl
, dump_file
, TDF_DETAILS
);
6198 return modified_cfg
;
6201 /* Expand builtins that can't be handled by HSA back-end. */
6206 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
6209 FOR_EACH_BB_FN (bb
, func
)
6211 for (gimple_stmt_iterator gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
);
6214 gimple
*stmt
= gsi_stmt (gsi
);
6216 if (gimple_code (stmt
) != GIMPLE_CALL
)
6219 gcall
*call
= as_a
<gcall
*> (stmt
);
6221 if (!gimple_call_builtin_p (call
, BUILT_IN_NORMAL
))
6224 tree fndecl
= gimple_call_fndecl (stmt
);
6225 enum built_in_function fn
= DECL_FUNCTION_CODE (fndecl
);
6228 case BUILT_IN_CEXPF
:
6229 case BUILT_IN_CEXPIF
:
6230 case BUILT_IN_CEXPI
:
6232 /* Similar to builtins.c (expand_builtin_cexpi), the builtin
6233 can be transformed to: cexp(I * z) = ccos(z) + I * csin(z). */
6234 tree lhs
= gimple_call_lhs (stmt
);
6235 tree rhs
= gimple_call_arg (stmt
, 0);
6236 tree rhs_type
= TREE_TYPE (rhs
);
6237 bool float_type_p
= rhs_type
== float_type_node
;
6238 tree real_part
= make_temp_ssa_name (rhs_type
, NULL
,
6240 tree imag_part
= make_temp_ssa_name (rhs_type
, NULL
,
6244 = mathfn_built_in (rhs_type
, fn
== float_type_p
6245 ? BUILT_IN_COSF
: BUILT_IN_COS
);
6246 gcall
*cos
= gimple_build_call (cos_fndecl
, 1, rhs
);
6247 gimple_call_set_lhs (cos
, real_part
);
6248 gsi_insert_before (&gsi
, cos
, GSI_SAME_STMT
);
6251 = mathfn_built_in (rhs_type
, fn
== float_type_p
6252 ? BUILT_IN_SINF
: BUILT_IN_SIN
);
6253 gcall
*sin
= gimple_build_call (sin_fndecl
, 1, rhs
);
6254 gimple_call_set_lhs (sin
, imag_part
);
6255 gsi_insert_before (&gsi
, sin
, GSI_SAME_STMT
);
6258 gassign
*assign
= gimple_build_assign (lhs
, COMPLEX_EXPR
,
6259 real_part
, imag_part
);
6260 gsi_insert_before (&gsi
, assign
, GSI_SAME_STMT
);
6261 gsi_remove (&gsi
, true);
6272 /* Emit HSA module variables that are global for the entire module. */
6275 emit_hsa_module_variables (void)
6277 hsa_num_threads
= new hsa_symbol (BRIG_TYPE_U32
, BRIG_SEGMENT_PRIVATE
,
6278 BRIG_LINKAGE_MODULE
, true);
6280 hsa_num_threads
->m_name
= "hsa_num_threads";
6282 hsa_brig_emit_omp_symbols ();
6285 /* Generate HSAIL representation of the current function and write into a
6286 special section of the output file. If KERNEL is set, the function will be
6287 considered an HSA kernel callable from the host, otherwise it will be
6288 compiled as an HSA function callable from other HSA code. */
6291 generate_hsa (bool kernel
)
6293 hsa_init_data_for_cfun ();
6295 if (hsa_num_threads
== NULL
)
6296 emit_hsa_module_variables ();
6298 bool modified_cfg
= convert_switch_statements ();
6299 /* Initialize hsa_cfun. */
6300 hsa_cfun
= new hsa_function_representation (cfun
->decl
, kernel
,
6301 SSANAMES (cfun
)->length (),
6303 hsa_cfun
->init_extra_bbs ();
6307 HSA_SORRY_AT (UNKNOWN_LOCATION
,
6308 "support for HSA does not implement transactional memory");
6312 verify_function_arguments (cfun
->decl
);
6313 if (hsa_seen_error ())
6316 hsa_cfun
->m_name
= get_brig_function_name (cfun
->decl
);
6318 gen_function_def_parameters ();
6319 if (hsa_seen_error ())
6324 gen_body_from_gimple ();
6325 if (hsa_seen_error ())
6328 if (hsa_cfun
->m_kernel_dispatch_count
)
6329 init_hsa_num_threads ();
6331 if (hsa_cfun
->m_kern_p
)
6333 hsa_function_summary
*s
6334 = hsa_summaries
->get (cgraph_node::get (hsa_cfun
->m_decl
));
6335 hsa_add_kern_decl_mapping (current_function_decl
, hsa_cfun
->m_name
,
6336 hsa_cfun
->m_maximum_omp_data_size
,
6337 s
->m_gridified_kernel_p
);
6342 for (unsigned i
= 0; i
< hsa_cfun
->m_ssa_map
.length (); i
++)
6343 if (hsa_cfun
->m_ssa_map
[i
])
6344 hsa_cfun
->m_ssa_map
[i
]->verify_ssa ();
6347 FOR_EACH_BB_FN (bb
, cfun
)
6349 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
6351 for (hsa_insn_basic
*insn
= hbb
->m_first_insn
; insn
;
6352 insn
= insn
->m_next
)
6358 hsa_brig_emit_function ();
6361 hsa_deinit_data_for_cfun ();
6366 const pass_data pass_data_gen_hsail
=
6369 "hsagen", /* name */
6370 OPTGROUP_NONE
, /* optinfo_flags */
6371 TV_NONE
, /* tv_id */
6372 PROP_cfg
| PROP_ssa
, /* properties_required */
6373 0, /* properties_provided */
6374 0, /* properties_destroyed */
6375 0, /* todo_flags_start */
6376 0 /* todo_flags_finish */
6379 class pass_gen_hsail
: public gimple_opt_pass
6382 pass_gen_hsail (gcc::context
*ctxt
)
6383 : gimple_opt_pass(pass_data_gen_hsail
, ctxt
)
6386 /* opt_pass methods: */
6387 bool gate (function
*);
6388 unsigned int execute (function
*);
6390 }; // class pass_gen_hsail
6392 /* Determine whether or not to run generation of HSAIL. */
6395 pass_gen_hsail::gate (function
*f
)
6397 return hsa_gen_requested_p ()
6398 && hsa_gpu_implementation_p (f
->decl
);
6402 pass_gen_hsail::execute (function
*)
6404 hsa_function_summary
*s
6405 = hsa_summaries
->get (cgraph_node::get_create (current_function_decl
));
6408 generate_hsa (s
->m_kind
== HSA_KERNEL
);
6409 TREE_ASM_WRITTEN (current_function_decl
) = 1;
6410 return TODO_discard_function
;
6415 /* Create the instance of hsa gen pass. */
6418 make_pass_gen_hsail (gcc::context
*ctxt
)
6420 return new pass_gen_hsail (ctxt
);