2016-05-04 Thomas Preud'homme <thomas.preudhomme@arm.com>
[official-gcc.git] / gcc / hsa-gen.c
blob5baf6073e3b1bd7179e5c3fca5f409d3770c9266
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)
11 any later version.
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/>. */
22 #include "config.h"
23 #include "system.h"
24 #include "coretypes.h"
25 #include "tm.h"
26 #include "is-a.h"
27 #include "hash-table.h"
28 #include "vec.h"
29 #include "tree.h"
30 #include "tree-pass.h"
31 #include "cfg.h"
32 #include "function.h"
33 #include "basic-block.h"
34 #include "fold-const.h"
35 #include "gimple.h"
36 #include "gimple-iterator.h"
37 #include "bitmap.h"
38 #include "dumpfile.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"
46 #include "tree-dfa.h"
47 #include "ssa-iterators.h"
48 #include "cgraph.h"
49 #include "print-tree.h"
50 #include "symbol-summary.h"
51 #include "hsa.h"
52 #include "cfghooks.h"
53 #include "tree-cfg.h"
54 #include "cfgloop.h"
55 #include "cfganal.h"
56 #include "builtins.h"
57 #include "params.h"
58 #include "gomp-constants.h"
59 #include "internal-fn.h"
60 #include "builtins.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, ...) \
66 do \
67 { \
68 hsa_fail_cfun (); \
69 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
70 HSA_SORRY_MSG)) \
71 inform (location, message, __VA_ARGS__); \
72 } \
73 while (false);
75 /* Same as previous, but highlight a location. */
77 #define HSA_SORRY_AT(location, message) \
78 do \
79 { \
80 hsa_fail_cfun (); \
81 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
82 HSA_SORRY_MSG)) \
83 inform (location, message); \
84 } \
85 while (false);
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
98 uint16_t header;
99 uint16_t setup;
100 uint16_t workgroup_size_x;
101 uint16_t workgroup_size_y;
102 uint16_t workgroup_size_z;
103 uint16_t reserved0;
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;
111 uint64_t reserved2;
112 uint64_t completion_signal;
115 /* HSA queue is shadow structure, originally provided by AMD. */
117 struct hsa_queue
119 int type;
120 uint32_t features;
121 void *base_address;
122 uint64_t doorbell_signal;
123 uint32_t size;
124 uint32_t reserved1;
125 uint64_t id;
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);
187 s /= BITS_PER_UNIT;
189 if (m_dim)
190 s *= m_dim;
192 return s;
195 /* Forward declaration. */
197 static BrigType16_t
198 hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p,
199 bool min32int);
201 void
202 hsa_symbol::fillup_for_decl (tree decl)
204 m_decl = decl;
205 m_type = hsa_type_for_tree_type (TREE_TYPE (decl), &m_dim, false);
207 if (hsa_seen_error ())
208 m_seen_error = true;
211 /* Constructor of class representing global HSA function/kernel information and
212 state. FNDECL is function declaration, KERNEL_P is true if the function
213 is going to become a HSA kernel. If the function has body, SSA_NAMES_COUNT
214 should be set to number of SSA names used in the function.
215 MODIFIED_CFG is set to true in case we modified control-flow graph
216 of the function. */
218 hsa_function_representation::hsa_function_representation
219 (tree fdecl, bool kernel_p, unsigned ssa_names_count, bool modified_cfg)
220 : m_name (NULL),
221 m_reg_count (0), m_input_args (vNULL),
222 m_output_arg (NULL), m_spill_symbols (vNULL), m_global_symbols (vNULL),
223 m_private_variables (vNULL), m_called_functions (vNULL),
224 m_called_internal_fns (vNULL), m_hbb_count (0),
225 m_in_ssa (true), m_kern_p (kernel_p), m_declaration_p (false),
226 m_decl (fdecl), m_internal_fn (NULL), m_shadow_reg (NULL),
227 m_kernel_dispatch_count (0), m_maximum_omp_data_size (0),
228 m_seen_error (false), m_temp_symbol_count (0), m_ssa_map (),
229 m_modified_cfg (modified_cfg)
231 int sym_init_len = (vec_safe_length (cfun->local_decls) / 2) + 1;;
232 m_local_symbols = new hash_table <hsa_noop_symbol_hasher> (sym_init_len);
233 m_ssa_map.safe_grow_cleared (ssa_names_count);
236 /* Constructor of class representing HSA function information that
237 is derived for an internal function. */
238 hsa_function_representation::hsa_function_representation (hsa_internal_fn *fn)
239 : m_reg_count (0), m_input_args (vNULL),
240 m_output_arg (NULL), m_local_symbols (NULL),
241 m_spill_symbols (vNULL), m_global_symbols (vNULL),
242 m_private_variables (vNULL), m_called_functions (vNULL),
243 m_called_internal_fns (vNULL), m_hbb_count (0),
244 m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL),
245 m_internal_fn (fn), m_shadow_reg (NULL), m_kernel_dispatch_count (0),
246 m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
247 m_ssa_map () {}
249 /* Destructor of class holding function/kernel-wide information and state. */
251 hsa_function_representation::~hsa_function_representation ()
253 /* Kernel names are deallocated at the end of BRIG output when deallocating
254 hsa_decl_kernel_mapping. */
255 if (!m_kern_p || m_seen_error)
256 free (m_name);
258 for (unsigned i = 0; i < m_input_args.length (); i++)
259 delete m_input_args[i];
260 m_input_args.release ();
262 delete m_output_arg;
263 delete m_local_symbols;
265 for (unsigned i = 0; i < m_spill_symbols.length (); i++)
266 delete m_spill_symbols[i];
267 m_spill_symbols.release ();
269 hsa_symbol *sym;
270 for (unsigned i = 0; i < m_global_symbols.iterate (i, &sym); i++)
271 if (sym->m_linkage != BRIG_ALLOCATION_PROGRAM)
272 delete sym;
273 m_global_symbols.release ();
275 for (unsigned i = 0; i < m_private_variables.length (); i++)
276 delete m_private_variables[i];
277 m_private_variables.release ();
278 m_called_functions.release ();
279 m_ssa_map.release ();
281 for (unsigned i = 0; i < m_called_internal_fns.length (); i++)
282 delete m_called_internal_fns[i];
285 hsa_op_reg *
286 hsa_function_representation::get_shadow_reg ()
288 /* If we compile a function with kernel dispatch and does not set
289 an optimization level, the function won't be inlined and
290 we return NULL. */
291 if (!m_kern_p)
292 return NULL;
294 if (m_shadow_reg)
295 return m_shadow_reg;
297 /* Append the shadow argument. */
298 hsa_symbol *shadow = new hsa_symbol (BRIG_TYPE_U64, BRIG_SEGMENT_KERNARG,
299 BRIG_LINKAGE_FUNCTION);
300 m_input_args.safe_push (shadow);
301 shadow->m_name = "hsa_runtime_shadow";
303 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_U64);
304 hsa_op_address *addr = new hsa_op_address (shadow);
306 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, r, addr);
307 hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun))->append_insn (mem);
308 m_shadow_reg = r;
310 return r;
313 bool hsa_function_representation::has_shadow_reg_p ()
315 return m_shadow_reg != NULL;
318 void
319 hsa_function_representation::init_extra_bbs ()
321 hsa_init_new_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
322 hsa_init_new_bb (EXIT_BLOCK_PTR_FOR_FN (cfun));
325 void
326 hsa_function_representation::update_dominance ()
328 if (m_modified_cfg)
330 free_dominance_info (CDI_DOMINATORS);
331 calculate_dominance_info (CDI_DOMINATORS);
335 hsa_symbol *
336 hsa_function_representation::create_hsa_temporary (BrigType16_t type)
338 hsa_symbol *s = new hsa_symbol (type, BRIG_SEGMENT_PRIVATE,
339 BRIG_LINKAGE_FUNCTION);
340 s->m_name_number = m_temp_symbol_count++;
342 hsa_cfun->m_private_variables.safe_push (s);
343 return s;
346 BrigLinkage8_t
347 hsa_function_representation::get_linkage ()
349 if (m_internal_fn)
350 return BRIG_LINKAGE_PROGRAM;
352 return m_kern_p || TREE_PUBLIC (m_decl) ?
353 BRIG_LINKAGE_PROGRAM : BRIG_LINKAGE_MODULE;
356 /* Hash map of simple OMP builtins. */
357 static hash_map <nofree_string_hash, omp_simple_builtin> *omp_simple_builtins
358 = NULL;
360 /* Warning messages for OMP builtins. */
362 #define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP " \
363 "lock routines"
364 #define HSA_WARN_TIMING_ROUTINE "support for HSA does not implement OpenMP " \
365 "timing routines"
366 #define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \
367 "undefined semantics within target regions, support for HSA ignores them"
368 #define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP " \
369 "affinity feateres"
371 /* Initialize hash map with simple OMP builtins. */
373 static void
374 hsa_init_simple_builtins ()
376 if (omp_simple_builtins != NULL)
377 return;
379 omp_simple_builtins
380 = new hash_map <nofree_string_hash, omp_simple_builtin> ();
382 omp_simple_builtin omp_builtins[] =
384 omp_simple_builtin ("omp_get_initial_device", NULL, false,
385 new hsa_op_immed (GOMP_DEVICE_HOST,
386 (BrigType16_t) BRIG_TYPE_S32)),
387 omp_simple_builtin ("omp_is_initial_device", NULL, false,
388 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
389 omp_simple_builtin ("omp_get_dynamic", NULL, false,
390 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
391 omp_simple_builtin ("omp_set_dynamic", NULL, false, NULL),
392 omp_simple_builtin ("omp_init_lock", HSA_WARN_LOCK_ROUTINE, true),
393 omp_simple_builtin ("omp_init_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
394 true),
395 omp_simple_builtin ("omp_init_nest_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
396 true),
397 omp_simple_builtin ("omp_destroy_lock", HSA_WARN_LOCK_ROUTINE, true),
398 omp_simple_builtin ("omp_set_lock", HSA_WARN_LOCK_ROUTINE, true),
399 omp_simple_builtin ("omp_unset_lock", HSA_WARN_LOCK_ROUTINE, true),
400 omp_simple_builtin ("omp_test_lock", HSA_WARN_LOCK_ROUTINE, true),
401 omp_simple_builtin ("omp_get_wtime", HSA_WARN_TIMING_ROUTINE, true),
402 omp_simple_builtin ("omp_get_wtick", HSA_WARN_TIMING_ROUTINE, true),
403 omp_simple_builtin ("omp_target_alloc", HSA_WARN_MEMORY_ROUTINE, false,
404 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_U64)),
405 omp_simple_builtin ("omp_target_free", HSA_WARN_MEMORY_ROUTINE, false),
406 omp_simple_builtin ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE,
407 false,
408 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
409 omp_simple_builtin ("omp_target_memcpy", HSA_WARN_MEMORY_ROUTINE, false,
410 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
411 omp_simple_builtin ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE,
412 false,
413 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
414 omp_simple_builtin ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE,
415 false,
416 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
417 omp_simple_builtin ("omp_target_disassociate_ptr",
418 HSA_WARN_MEMORY_ROUTINE,
419 false,
420 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
421 omp_simple_builtin ("omp_set_max_active_levels",
422 "Support for HSA only allows only one active level, "
423 "call to omp_set_max_active_levels will be ignored "
424 "in the generated HSAIL",
425 false, NULL),
426 omp_simple_builtin ("omp_get_max_active_levels", NULL, false,
427 new hsa_op_immed (1, (BrigType16_t) BRIG_TYPE_S32)),
428 omp_simple_builtin ("omp_in_final", NULL, false,
429 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
430 omp_simple_builtin ("omp_get_proc_bind", HSA_WARN_AFFINITY, false,
431 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
432 omp_simple_builtin ("omp_get_num_places", HSA_WARN_AFFINITY, false,
433 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
434 omp_simple_builtin ("omp_get_place_num_procs", HSA_WARN_AFFINITY, false,
435 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
436 omp_simple_builtin ("omp_get_place_proc_ids", HSA_WARN_AFFINITY, false,
437 NULL),
438 omp_simple_builtin ("omp_get_place_num", HSA_WARN_AFFINITY, false,
439 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
440 omp_simple_builtin ("omp_get_partition_num_places", HSA_WARN_AFFINITY,
441 false,
442 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
443 omp_simple_builtin ("omp_get_partition_place_nums", HSA_WARN_AFFINITY,
444 false, NULL),
445 omp_simple_builtin ("omp_set_default_device",
446 "omp_set_default_device has undefined semantics "
447 "within target regions, support for HSA ignores it",
448 false, NULL),
449 omp_simple_builtin ("omp_get_default_device",
450 "omp_get_default_device has undefined semantics "
451 "within target regions, support for HSA ignores it",
452 false,
453 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
454 omp_simple_builtin ("omp_get_num_devices",
455 "omp_get_num_devices has undefined semantics "
456 "within target regions, support for HSA ignores it",
457 false,
458 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
459 omp_simple_builtin ("omp_get_num_procs", NULL, true, NULL),
460 omp_simple_builtin ("omp_get_cancellation", NULL, true, NULL),
461 omp_simple_builtin ("omp_set_nested", NULL, true, NULL),
462 omp_simple_builtin ("omp_get_nested", NULL, true, NULL),
463 omp_simple_builtin ("omp_set_schedule", NULL, true, NULL),
464 omp_simple_builtin ("omp_get_schedule", NULL, true, NULL),
465 omp_simple_builtin ("omp_get_thread_limit", NULL, true, NULL),
466 omp_simple_builtin ("omp_get_team_size", NULL, true, NULL),
467 omp_simple_builtin ("omp_get_ancestor_thread_num", NULL, true, NULL),
468 omp_simple_builtin ("omp_get_max_task_priority", NULL, true, NULL)
471 unsigned count = sizeof (omp_builtins) / sizeof (omp_simple_builtin);
473 for (unsigned i = 0; i < count; i++)
474 omp_simple_builtins->put (omp_builtins[i].m_name, omp_builtins[i]);
477 /* Allocate HSA structures that we need only while generating with this. */
479 static void
480 hsa_init_data_for_cfun ()
482 hsa_init_compilation_unit_data ();
483 hsa_allocp_operand_address
484 = new object_allocator<hsa_op_address> ("HSA address operands");
485 hsa_allocp_operand_immed
486 = new object_allocator<hsa_op_immed> ("HSA immediate operands");
487 hsa_allocp_operand_reg
488 = new object_allocator<hsa_op_reg> ("HSA register operands");
489 hsa_allocp_operand_code_list
490 = new object_allocator<hsa_op_code_list> ("HSA code list operands");
491 hsa_allocp_operand_operand_list
492 = new object_allocator<hsa_op_operand_list> ("HSA operand list operands");
493 hsa_allocp_inst_basic
494 = new object_allocator<hsa_insn_basic> ("HSA basic instructions");
495 hsa_allocp_inst_phi
496 = new object_allocator<hsa_insn_phi> ("HSA phi operands");
497 hsa_allocp_inst_mem
498 = new object_allocator<hsa_insn_mem> ("HSA memory instructions");
499 hsa_allocp_inst_atomic
500 = new object_allocator<hsa_insn_atomic> ("HSA atomic instructions");
501 hsa_allocp_inst_signal
502 = new object_allocator<hsa_insn_signal> ("HSA signal instructions");
503 hsa_allocp_inst_seg
504 = new object_allocator<hsa_insn_seg> ("HSA segment conversion "
505 "instructions");
506 hsa_allocp_inst_cmp
507 = new object_allocator<hsa_insn_cmp> ("HSA comparison instructions");
508 hsa_allocp_inst_br
509 = new object_allocator<hsa_insn_br> ("HSA branching instructions");
510 hsa_allocp_inst_sbr
511 = new object_allocator<hsa_insn_sbr> ("HSA switch branching instructions");
512 hsa_allocp_inst_call
513 = new object_allocator<hsa_insn_call> ("HSA call instructions");
514 hsa_allocp_inst_arg_block
515 = new object_allocator<hsa_insn_arg_block> ("HSA arg block instructions");
516 hsa_allocp_inst_comment
517 = new object_allocator<hsa_insn_comment> ("HSA comment instructions");
518 hsa_allocp_inst_queue
519 = new object_allocator<hsa_insn_queue> ("HSA queue instructions");
520 hsa_allocp_inst_srctype
521 = new object_allocator<hsa_insn_srctype> ("HSA source type instructions");
522 hsa_allocp_inst_packed
523 = new object_allocator<hsa_insn_packed> ("HSA packed instructions");
524 hsa_allocp_inst_cvt
525 = new object_allocator<hsa_insn_cvt> ("HSA convert instructions");
526 hsa_allocp_inst_alloca
527 = new object_allocator<hsa_insn_alloca> ("HSA alloca instructions");
528 hsa_allocp_bb = new object_allocator<hsa_bb> ("HSA basic blocks");
531 /* Deinitialize HSA subsystem and free all allocated memory. */
533 static void
534 hsa_deinit_data_for_cfun (void)
536 basic_block bb;
538 FOR_ALL_BB_FN (bb, cfun)
539 if (bb->aux)
541 hsa_bb *hbb = hsa_bb_for_bb (bb);
542 hbb->~hsa_bb ();
543 bb->aux = NULL;
546 for (unsigned int i = 0; i < hsa_operands.length (); i++)
547 hsa_destroy_operand (hsa_operands[i]);
549 hsa_operands.release ();
551 for (unsigned i = 0; i < hsa_instructions.length (); i++)
552 hsa_destroy_insn (hsa_instructions[i]);
554 hsa_instructions.release ();
556 if (omp_simple_builtins != NULL)
558 delete omp_simple_builtins;
559 omp_simple_builtins = NULL;
562 delete hsa_allocp_operand_address;
563 delete hsa_allocp_operand_immed;
564 delete hsa_allocp_operand_reg;
565 delete hsa_allocp_operand_code_list;
566 delete hsa_allocp_operand_operand_list;
567 delete hsa_allocp_inst_basic;
568 delete hsa_allocp_inst_phi;
569 delete hsa_allocp_inst_atomic;
570 delete hsa_allocp_inst_mem;
571 delete hsa_allocp_inst_signal;
572 delete hsa_allocp_inst_seg;
573 delete hsa_allocp_inst_cmp;
574 delete hsa_allocp_inst_br;
575 delete hsa_allocp_inst_sbr;
576 delete hsa_allocp_inst_call;
577 delete hsa_allocp_inst_arg_block;
578 delete hsa_allocp_inst_comment;
579 delete hsa_allocp_inst_queue;
580 delete hsa_allocp_inst_srctype;
581 delete hsa_allocp_inst_packed;
582 delete hsa_allocp_inst_cvt;
583 delete hsa_allocp_inst_alloca;
584 delete hsa_allocp_bb;
585 delete hsa_cfun;
588 /* Return the type which holds addresses in the given SEGMENT. */
590 static BrigType16_t
591 hsa_get_segment_addr_type (BrigSegment8_t segment)
593 switch (segment)
595 case BRIG_SEGMENT_NONE:
596 gcc_unreachable ();
598 case BRIG_SEGMENT_FLAT:
599 case BRIG_SEGMENT_GLOBAL:
600 case BRIG_SEGMENT_READONLY:
601 case BRIG_SEGMENT_KERNARG:
602 return hsa_machine_large_p () ? BRIG_TYPE_U64 : BRIG_TYPE_U32;
604 case BRIG_SEGMENT_GROUP:
605 case BRIG_SEGMENT_PRIVATE:
606 case BRIG_SEGMENT_SPILL:
607 case BRIG_SEGMENT_ARG:
608 return BRIG_TYPE_U32;
610 gcc_unreachable ();
613 /* Return integer brig type according to provided SIZE in bytes. If SIGN
614 is set to true, return signed integer type. */
616 static BrigType16_t
617 get_integer_type_by_bytes (unsigned size, bool sign)
619 if (sign)
620 switch (size)
622 case 1:
623 return BRIG_TYPE_S8;
624 case 2:
625 return BRIG_TYPE_S16;
626 case 4:
627 return BRIG_TYPE_S32;
628 case 8:
629 return BRIG_TYPE_S64;
630 default:
631 break;
633 else
634 switch (size)
636 case 1:
637 return BRIG_TYPE_U8;
638 case 2:
639 return BRIG_TYPE_U16;
640 case 4:
641 return BRIG_TYPE_U32;
642 case 8:
643 return BRIG_TYPE_U64;
644 default:
645 break;
648 return 0;
651 /* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
652 are assumed to use flat addressing. If min32int is true, always expand
653 integer types to one that has at least 32 bits. */
655 static BrigType16_t
656 hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
658 HOST_WIDE_INT bsize;
659 const_tree base;
660 BrigType16_t res = BRIG_TYPE_NONE;
662 gcc_checking_assert (TYPE_P (type));
663 gcc_checking_assert (!AGGREGATE_TYPE_P (type));
664 if (POINTER_TYPE_P (type))
665 return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
667 if (TREE_CODE (type) == VECTOR_TYPE || TREE_CODE (type) == COMPLEX_TYPE)
668 base = TREE_TYPE (type);
669 else
670 base = type;
672 if (!tree_fits_uhwi_p (TYPE_SIZE (base)))
674 HSA_SORRY_ATV (EXPR_LOCATION (type),
675 "support for HSA does not implement huge or "
676 "variable-sized type %T", type);
677 return res;
680 bsize = tree_to_uhwi (TYPE_SIZE (base));
681 unsigned byte_size = bsize / BITS_PER_UNIT;
682 if (INTEGRAL_TYPE_P (base))
683 res = get_integer_type_by_bytes (byte_size, !TYPE_UNSIGNED (base));
684 else if (SCALAR_FLOAT_TYPE_P (base))
686 switch (bsize)
688 case 16:
689 res = BRIG_TYPE_F16;
690 break;
691 case 32:
692 res = BRIG_TYPE_F32;
693 break;
694 case 64:
695 res = BRIG_TYPE_F64;
696 break;
697 default:
698 break;
702 if (res == BRIG_TYPE_NONE)
704 HSA_SORRY_ATV (EXPR_LOCATION (type),
705 "support for HSA does not implement type %T", type);
706 return res;
709 if (TREE_CODE (type) == VECTOR_TYPE)
711 HOST_WIDE_INT tsize = tree_to_uhwi (TYPE_SIZE (type));
713 if (bsize == tsize)
715 HSA_SORRY_ATV (EXPR_LOCATION (type),
716 "support for HSA does not implement a vector type "
717 "where a type and unit size are equal: %T", type);
718 return res;
721 switch (tsize)
723 case 32:
724 res |= BRIG_TYPE_PACK_32;
725 break;
726 case 64:
727 res |= BRIG_TYPE_PACK_64;
728 break;
729 case 128:
730 res |= BRIG_TYPE_PACK_128;
731 break;
732 default:
733 HSA_SORRY_ATV (EXPR_LOCATION (type),
734 "support for HSA does not implement type %T", type);
738 if (min32int)
740 /* Registers/immediate operands can only be 32bit or more except for
741 f16. */
742 if (res == BRIG_TYPE_U8 || res == BRIG_TYPE_U16)
743 res = BRIG_TYPE_U32;
744 else if (res == BRIG_TYPE_S8 || res == BRIG_TYPE_S16)
745 res = BRIG_TYPE_S32;
748 if (TREE_CODE (type) == COMPLEX_TYPE)
750 unsigned bsize = 2 * hsa_type_bit_size (res);
751 res = hsa_bittype_for_bitsize (bsize);
754 return res;
757 /* Returns the BRIG type we need to load/store entities of TYPE. */
759 static BrigType16_t
760 mem_type_for_type (BrigType16_t type)
762 /* HSA has non-intuitive constraints on load/store types. If it's
763 a bit-type it _must_ be B128, if it's not a bit-type it must be
764 64bit max. So for loading entities of 128 bits (e.g. vectors)
765 we have to to B128, while for loading the rest we have to use the
766 input type (??? or maybe also flattened to a equally sized non-vector
767 unsigned type?). */
768 if ((type & BRIG_TYPE_PACK_MASK) == BRIG_TYPE_PACK_128)
769 return BRIG_TYPE_B128;
770 else if (hsa_btype_p (type) || hsa_type_packed_p (type))
772 unsigned bitsize = hsa_type_bit_size (type);
773 if (bitsize < 128)
774 return hsa_uint_for_bitsize (bitsize);
775 else
776 return hsa_bittype_for_bitsize (bitsize);
778 return type;
781 /* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some
782 kind of array will be generated, setting DIM appropriately. Otherwise, it
783 will be set to zero. */
785 static BrigType16_t
786 hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p = NULL,
787 bool min32int = false)
789 gcc_checking_assert (TYPE_P (type));
790 if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type)))
792 HSA_SORRY_ATV (EXPR_LOCATION (type), "support for HSA does not "
793 "implement huge or variable-sized type %T", type);
794 return BRIG_TYPE_NONE;
797 if (RECORD_OR_UNION_TYPE_P (type))
799 if (dim_p)
800 *dim_p = tree_to_uhwi (TYPE_SIZE_UNIT (type));
801 return BRIG_TYPE_U8 | BRIG_TYPE_ARRAY;
804 if (TREE_CODE (type) == ARRAY_TYPE)
806 /* We try to be nice and use the real base-type when this is an array of
807 scalars and only resort to an array of bytes if the type is more
808 complex. */
810 unsigned HOST_WIDE_INT dim = 1;
812 while (TREE_CODE (type) == ARRAY_TYPE)
814 tree domain = TYPE_DOMAIN (type);
815 if (!TYPE_MIN_VALUE (domain)
816 || !TYPE_MAX_VALUE (domain)
817 || !tree_fits_shwi_p (TYPE_MIN_VALUE (domain))
818 || !tree_fits_shwi_p (TYPE_MAX_VALUE (domain)))
820 HSA_SORRY_ATV (EXPR_LOCATION (type),
821 "support for HSA does not implement array %T with "
822 "unknown bounds", type);
823 return BRIG_TYPE_NONE;
825 HOST_WIDE_INT min = tree_to_shwi (TYPE_MIN_VALUE (domain));
826 HOST_WIDE_INT max = tree_to_shwi (TYPE_MAX_VALUE (domain));
827 dim = dim * (unsigned HOST_WIDE_INT) (max - min + 1);
828 type = TREE_TYPE (type);
831 BrigType16_t res;
832 if (RECORD_OR_UNION_TYPE_P (type))
834 dim = dim * tree_to_uhwi (TYPE_SIZE_UNIT (type));
835 res = BRIG_TYPE_U8;
837 else
838 res = hsa_type_for_scalar_tree_type (type, false);
840 if (dim_p)
841 *dim_p = dim;
842 return res | BRIG_TYPE_ARRAY;
845 /* Scalar case: */
846 if (dim_p)
847 *dim_p = 0;
849 return hsa_type_for_scalar_tree_type (type, min32int);
852 /* Returns true if converting from STYPE into DTYPE needs the _CVT
853 opcode. If false a normal _MOV is enough. */
855 static bool
856 hsa_needs_cvt (BrigType16_t dtype, BrigType16_t stype)
858 if (hsa_btype_p (dtype))
859 return false;
861 /* float <-> int conversions are real converts. */
862 if (hsa_type_float_p (dtype) != hsa_type_float_p (stype))
863 return true;
864 /* When both types have different size, then we need CVT as well. */
865 if (hsa_type_bit_size (dtype) != hsa_type_bit_size (stype))
866 return true;
867 return false;
870 /* Lookup or create the associated hsa_symbol structure with a given VAR_DECL
871 or lookup the hsa_structure corresponding to a PARM_DECL. */
873 static hsa_symbol *
874 get_symbol_for_decl (tree decl)
876 hsa_symbol **slot;
877 hsa_symbol dummy (BRIG_TYPE_NONE, BRIG_SEGMENT_NONE, BRIG_LINKAGE_NONE);
879 gcc_assert (TREE_CODE (decl) == PARM_DECL
880 || TREE_CODE (decl) == RESULT_DECL
881 || TREE_CODE (decl) == VAR_DECL);
883 dummy.m_decl = decl;
885 bool is_in_global_vars
886 = TREE_CODE (decl) == VAR_DECL && is_global_var (decl);
888 if (is_in_global_vars)
889 slot = hsa_global_variable_symbols->find_slot (&dummy, INSERT);
890 else
891 slot = hsa_cfun->m_local_symbols->find_slot (&dummy, INSERT);
893 gcc_checking_assert (slot);
894 if (*slot)
896 hsa_symbol *sym = (*slot);
898 /* If the symbol is problematic, mark current function also as
899 problematic. */
900 if (sym->m_seen_error)
901 hsa_fail_cfun ();
903 /* PR hsa/70234: If a global variable was marked to be emitted,
904 but HSAIL generation of a function using the variable fails,
905 we should retry to emit the variable in context of a different
906 function.
908 Iterate elements whether a symbol is already in m_global_symbols
909 of not. */
910 if (is_in_global_vars && !sym->m_emitted_to_brig)
912 for (unsigned i = 0; i < hsa_cfun->m_global_symbols.length (); i++)
913 if (hsa_cfun->m_global_symbols[i] == sym)
914 return *slot;
915 hsa_cfun->m_global_symbols.safe_push (sym);
918 return *slot;
920 else
922 hsa_symbol *sym;
923 gcc_assert (TREE_CODE (decl) == VAR_DECL);
924 BrigAlignment8_t align = hsa_object_alignment (decl);
926 if (is_in_global_vars)
928 sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_GLOBAL,
929 BRIG_LINKAGE_PROGRAM, true,
930 BRIG_ALLOCATION_PROGRAM, align);
931 hsa_cfun->m_global_symbols.safe_push (sym);
933 else
935 /* As generation of efficient memory copy instructions relies
936 on alignment greater or equal to 8 bytes,
937 we need to increase alignment of all aggregate types.. */
938 if (AGGREGATE_TYPE_P (TREE_TYPE (decl)))
939 align = MAX ((BrigAlignment8_t) BRIG_ALIGNMENT_8, align);
941 /* PARM_DECL and RESULT_DECL should be already in m_local_symbols. */
942 gcc_assert (TREE_CODE (decl) == VAR_DECL);
944 sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_PRIVATE,
945 BRIG_LINKAGE_FUNCTION);
946 sym->m_align = align;
947 hsa_cfun->m_private_variables.safe_push (sym);
950 sym->fillup_for_decl (decl);
951 sym->m_name = hsa_get_declaration_name (decl);
953 *slot = sym;
954 return sym;
958 /* For a given HSA function declaration, return a host
959 function declaration. */
961 tree
962 hsa_get_host_function (tree decl)
964 hsa_function_summary *s
965 = hsa_summaries->get (cgraph_node::get_create (decl));
966 gcc_assert (s->m_kind != HSA_NONE);
967 gcc_assert (s->m_gpu_implementation_p);
969 return s->m_binded_function->decl;
972 /* Return true if function DECL has a host equivalent function. */
974 static char *
975 get_brig_function_name (tree decl)
977 tree d = decl;
979 hsa_function_summary *s = hsa_summaries->get (cgraph_node::get_create (d));
980 if (s->m_kind != HSA_NONE && s->m_gpu_implementation_p)
981 d = s->m_binded_function->decl;
983 /* IPA split can create a function that has no host equivalent. */
984 if (d == NULL)
985 d = decl;
987 char *name = xstrdup (hsa_get_declaration_name (d));
988 hsa_sanitize_name (name);
990 return name;
993 /* Create a spill symbol of type TYPE. */
995 hsa_symbol *
996 hsa_get_spill_symbol (BrigType16_t type)
998 hsa_symbol *sym = new hsa_symbol (type, BRIG_SEGMENT_SPILL,
999 BRIG_LINKAGE_FUNCTION);
1000 hsa_cfun->m_spill_symbols.safe_push (sym);
1001 return sym;
1004 /* Create a symbol for a read-only string constant. */
1005 hsa_symbol *
1006 hsa_get_string_cst_symbol (tree string_cst)
1008 gcc_checking_assert (TREE_CODE (string_cst) == STRING_CST);
1010 hsa_symbol **slot = hsa_cfun->m_string_constants_map.get (string_cst);
1011 if (slot)
1012 return *slot;
1014 hsa_op_immed *cst = new hsa_op_immed (string_cst);
1015 hsa_symbol *sym = new hsa_symbol (cst->m_type, BRIG_SEGMENT_GLOBAL,
1016 BRIG_LINKAGE_MODULE, true,
1017 BRIG_ALLOCATION_AGENT);
1018 sym->m_cst_value = cst;
1019 sym->m_dim = TREE_STRING_LENGTH (string_cst);
1020 sym->m_name_number = hsa_cfun->m_global_symbols.length ();
1022 hsa_cfun->m_global_symbols.safe_push (sym);
1023 hsa_cfun->m_string_constants_map.put (string_cst, sym);
1024 return sym;
1027 /* Constructor of the ancestor of all operands. K is BRIG kind that identified
1028 what the operator is. */
1030 hsa_op_base::hsa_op_base (BrigKind16_t k)
1031 : m_next (NULL), m_brig_op_offset (0), m_kind (k)
1033 hsa_operands.safe_push (this);
1036 /* Constructor of ancestor of all operands which have a type. K is BRIG kind
1037 that identified what the operator is. T is the type of the operator. */
1039 hsa_op_with_type::hsa_op_with_type (BrigKind16_t k, BrigType16_t t)
1040 : hsa_op_base (k), m_type (t)
1044 hsa_op_with_type *
1045 hsa_op_with_type::get_in_type (BrigType16_t dtype, hsa_bb *hbb)
1047 if (m_type == dtype)
1048 return this;
1050 hsa_op_reg *dest;
1052 if (hsa_needs_cvt (dtype, m_type))
1054 dest = new hsa_op_reg (dtype);
1055 hbb->append_insn (new hsa_insn_cvt (dest, this));
1057 else
1059 dest = new hsa_op_reg (m_type);
1060 hbb->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV,
1061 dest->m_type, dest, this));
1063 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
1064 type of the operand must be same as type of the instruction. */
1065 dest->m_type = dtype;
1068 return dest;
1071 /* Constructor of class representing HSA immediate values. TREE_VAL is the
1072 tree representation of the immediate value. If min32int is true,
1073 always expand integer types to one that has at least 32 bits. */
1075 hsa_op_immed::hsa_op_immed (tree tree_val, bool min32int)
1076 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES,
1077 hsa_type_for_tree_type (TREE_TYPE (tree_val), NULL,
1078 min32int))
1080 if (hsa_seen_error ())
1081 return;
1083 gcc_checking_assert ((is_gimple_min_invariant (tree_val)
1084 && (!POINTER_TYPE_P (TREE_TYPE (tree_val))
1085 || TREE_CODE (tree_val) == INTEGER_CST))
1086 || TREE_CODE (tree_val) == CONSTRUCTOR);
1087 m_tree_value = tree_val;
1089 /* Verify that all elements of a constructor are constants. */
1090 if (TREE_CODE (m_tree_value) == CONSTRUCTOR)
1091 for (unsigned i = 0;
1092 i < vec_safe_length (CONSTRUCTOR_ELTS (m_tree_value)); i++)
1094 tree v = CONSTRUCTOR_ELT (m_tree_value, i)->value;
1095 if (!CONSTANT_CLASS_P (v))
1097 HSA_SORRY_AT (EXPR_LOCATION (tree_val),
1098 "HSA ctor should have only constants");
1099 return;
1104 /* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
1105 integer representation of the immediate value. TYPE is BRIG type. */
1107 hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value, BrigType16_t type)
1108 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES, type),
1109 m_tree_value (NULL)
1111 gcc_assert (hsa_type_integer_p (type));
1112 m_int_value = integer_value;
1115 hsa_op_immed::hsa_op_immed ()
1116 : hsa_op_with_type (BRIG_KIND_NONE, BRIG_TYPE_NONE)
1120 /* New operator to allocate immediate operands from pool alloc. */
1122 void *
1123 hsa_op_immed::operator new (size_t)
1125 return hsa_allocp_operand_immed->allocate_raw ();
1128 /* Destructor. */
1130 hsa_op_immed::~hsa_op_immed ()
1134 /* Change type of the immediate value to T. */
1136 void
1137 hsa_op_immed::set_type (BrigType16_t t)
1139 m_type = t;
1142 /* Constructor of class representing HSA registers and pseudo-registers. T is
1143 the BRIG type of the new register. */
1145 hsa_op_reg::hsa_op_reg (BrigType16_t t)
1146 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER, t), m_gimple_ssa (NULL_TREE),
1147 m_def_insn (NULL), m_spill_sym (NULL), m_order (hsa_cfun->m_reg_count++),
1148 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
1152 /* New operator to allocate a register from pool alloc. */
1154 void *
1155 hsa_op_reg::operator new (size_t)
1157 return hsa_allocp_operand_reg->allocate_raw ();
1160 /* Verify register operand. */
1162 void
1163 hsa_op_reg::verify_ssa ()
1165 /* Verify that each HSA register has a definition assigned.
1166 Exceptions are VAR_DECL and PARM_DECL that are a default
1167 definition. */
1168 gcc_checking_assert (m_def_insn
1169 || (m_gimple_ssa != NULL
1170 && (!SSA_NAME_VAR (m_gimple_ssa)
1171 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa))
1172 != PARM_DECL))
1173 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa)));
1175 /* Verify that every use of the register is really present
1176 in an instruction. */
1177 for (unsigned i = 0; i < m_uses.length (); i++)
1179 hsa_insn_basic *use = m_uses[i];
1181 bool is_visited = false;
1182 for (unsigned j = 0; j < use->operand_count (); j++)
1184 hsa_op_base *u = use->get_op (j);
1185 hsa_op_address *addr; addr = dyn_cast <hsa_op_address *> (u);
1186 if (addr && addr->m_reg)
1187 u = addr->m_reg;
1189 if (u == this)
1191 bool r = !addr && use->op_output_p (j);
1193 if (r)
1195 error ("HSA SSA name defined by instruction that is supposed "
1196 "to be using it");
1197 debug_hsa_operand (this);
1198 debug_hsa_insn (use);
1199 internal_error ("HSA SSA verification failed");
1202 is_visited = true;
1206 if (!is_visited)
1208 error ("HSA SSA name not among operands of instruction that is "
1209 "supposed to use it");
1210 debug_hsa_operand (this);
1211 debug_hsa_insn (use);
1212 internal_error ("HSA SSA verification failed");
1217 hsa_op_address::hsa_op_address (hsa_symbol *sym, hsa_op_reg *r,
1218 HOST_WIDE_INT offset)
1219 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (r),
1220 m_imm_offset (offset)
1224 hsa_op_address::hsa_op_address (hsa_symbol *sym, HOST_WIDE_INT offset)
1225 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (NULL),
1226 m_imm_offset (offset)
1230 hsa_op_address::hsa_op_address (hsa_op_reg *r, HOST_WIDE_INT offset)
1231 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (NULL), m_reg (r),
1232 m_imm_offset (offset)
1236 /* New operator to allocate address operands from pool alloc. */
1238 void *
1239 hsa_op_address::operator new (size_t)
1241 return hsa_allocp_operand_address->allocate_raw ();
1244 /* Constructor of an operand referring to HSAIL code. */
1246 hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF),
1247 m_directive_offset (0)
1251 /* Constructor of an operand representing a code list. Set it up so that it
1252 can contain ELEMENTS number of elements. */
1254 hsa_op_code_list::hsa_op_code_list (unsigned elements)
1255 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST)
1257 m_offsets.create (1);
1258 m_offsets.safe_grow_cleared (elements);
1261 /* New operator to allocate code list operands from pool alloc. */
1263 void *
1264 hsa_op_code_list::operator new (size_t)
1266 return hsa_allocp_operand_code_list->allocate_raw ();
1269 /* Constructor of an operand representing an operand list.
1270 Set it up so that it can contain ELEMENTS number of elements. */
1272 hsa_op_operand_list::hsa_op_operand_list (unsigned elements)
1273 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST)
1275 m_offsets.create (elements);
1276 m_offsets.safe_grow (elements);
1279 /* New operator to allocate operand list operands from pool alloc. */
1281 void *
1282 hsa_op_operand_list::operator new (size_t)
1284 return hsa_allocp_operand_operand_list->allocate_raw ();
1287 hsa_op_operand_list::~hsa_op_operand_list ()
1289 m_offsets.release ();
1293 hsa_op_reg *
1294 hsa_function_representation::reg_for_gimple_ssa (tree ssa)
1296 hsa_op_reg *hreg;
1298 gcc_checking_assert (TREE_CODE (ssa) == SSA_NAME);
1299 if (m_ssa_map[SSA_NAME_VERSION (ssa)])
1300 return m_ssa_map[SSA_NAME_VERSION (ssa)];
1302 hreg = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa),
1303 true));
1304 hreg->m_gimple_ssa = ssa;
1305 m_ssa_map[SSA_NAME_VERSION (ssa)] = hreg;
1307 return hreg;
1310 void
1311 hsa_op_reg::set_definition (hsa_insn_basic *insn)
1313 if (hsa_cfun->m_in_ssa)
1315 gcc_checking_assert (!m_def_insn);
1316 m_def_insn = insn;
1318 else
1319 m_def_insn = NULL;
1322 /* Constructor of the class which is the bases of all instructions and directly
1323 represents the most basic ones. NOPS is the number of operands that the
1324 operand vector will contain (and which will be cleared). OP is the opcode
1325 of the instruction. This constructor does not set type. */
1327 hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc)
1328 : m_prev (NULL),
1329 m_next (NULL), m_bb (NULL), m_opcode (opc), m_number (0),
1330 m_type (BRIG_TYPE_NONE), m_brig_offset (0)
1332 if (nops > 0)
1333 m_operands.safe_grow_cleared (nops);
1335 hsa_instructions.safe_push (this);
1338 /* Make OP the operand number INDEX of operands of this instruction. If OP is a
1339 register or an address containing a register, then either set the definition
1340 of the register to this instruction if it an output operand or add this
1341 instruction to the uses if it is an input one. */
1343 void
1344 hsa_insn_basic::set_op (int index, hsa_op_base *op)
1346 /* Each address operand is always use. */
1347 hsa_op_address *addr = dyn_cast <hsa_op_address *> (op);
1348 if (addr && addr->m_reg)
1349 addr->m_reg->m_uses.safe_push (this);
1350 else
1352 hsa_op_reg *reg = dyn_cast <hsa_op_reg *> (op);
1353 if (reg)
1355 if (op_output_p (index))
1356 reg->set_definition (this);
1357 else
1358 reg->m_uses.safe_push (this);
1362 m_operands[index] = op;
1365 /* Get INDEX-th operand of the instruction. */
1367 hsa_op_base *
1368 hsa_insn_basic::get_op (int index)
1370 return m_operands[index];
1373 /* Get address of INDEX-th operand of the instruction. */
1375 hsa_op_base **
1376 hsa_insn_basic::get_op_addr (int index)
1378 return &m_operands[index];
1381 /* Get number of operands of the instruction. */
1382 unsigned int
1383 hsa_insn_basic::operand_count ()
1385 return m_operands.length ();
1388 /* Constructor of the class which is the bases of all instructions and directly
1389 represents the most basic ones. NOPS is the number of operands that the
1390 operand vector will contain (and which will be cleared). OPC is the opcode
1391 of the instruction, T is the type of the instruction. */
1393 hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc, BrigType16_t t,
1394 hsa_op_base *arg0, hsa_op_base *arg1,
1395 hsa_op_base *arg2, hsa_op_base *arg3)
1396 : m_prev (NULL), m_next (NULL), m_bb (NULL), m_opcode (opc),m_number (0),
1397 m_type (t), m_brig_offset (0)
1399 if (nops > 0)
1400 m_operands.safe_grow_cleared (nops);
1402 if (arg0 != NULL)
1404 gcc_checking_assert (nops >= 1);
1405 set_op (0, arg0);
1408 if (arg1 != NULL)
1410 gcc_checking_assert (nops >= 2);
1411 set_op (1, arg1);
1414 if (arg2 != NULL)
1416 gcc_checking_assert (nops >= 3);
1417 set_op (2, arg2);
1420 if (arg3 != NULL)
1422 gcc_checking_assert (nops >= 4);
1423 set_op (3, arg3);
1426 hsa_instructions.safe_push (this);
1429 /* New operator to allocate basic instruction from pool alloc. */
1431 void *
1432 hsa_insn_basic::operator new (size_t)
1434 return hsa_allocp_inst_basic->allocate_raw ();
1437 /* Verify the instruction. */
1439 void
1440 hsa_insn_basic::verify ()
1442 hsa_op_address *addr;
1443 hsa_op_reg *reg;
1445 /* Iterate all register operands and verify that the instruction
1446 is set in uses of the register. */
1447 for (unsigned i = 0; i < operand_count (); i++)
1449 hsa_op_base *use = get_op (i);
1451 if ((addr = dyn_cast <hsa_op_address *> (use)) && addr->m_reg)
1453 gcc_assert (addr->m_reg->m_def_insn != this);
1454 use = addr->m_reg;
1457 if ((reg = dyn_cast <hsa_op_reg *> (use)) && !op_output_p (i))
1459 unsigned j;
1460 for (j = 0; j < reg->m_uses.length (); j++)
1462 if (reg->m_uses[j] == this)
1463 break;
1466 if (j == reg->m_uses.length ())
1468 error ("HSA instruction uses a register but is not among "
1469 "recorded register uses");
1470 debug_hsa_operand (reg);
1471 debug_hsa_insn (this);
1472 internal_error ("HSA instruction verification failed");
1478 /* Constructor of an instruction representing a PHI node. NOPS is the number
1479 of operands (equal to the number of predecessors). */
1481 hsa_insn_phi::hsa_insn_phi (unsigned nops, hsa_op_reg *dst)
1482 : hsa_insn_basic (nops, HSA_OPCODE_PHI), m_dest (dst)
1484 dst->set_definition (this);
1487 /* New operator to allocate PHI instruction from pool alloc. */
1489 void *
1490 hsa_insn_phi::operator new (size_t)
1492 return hsa_allocp_inst_phi->allocate_raw ();
1495 /* Constructor of class representing instruction for conditional jump, CTRL is
1496 the control register determining whether the jump will be carried out, the
1497 new instruction is automatically added to its uses list. */
1499 hsa_insn_br::hsa_insn_br (hsa_op_reg *ctrl)
1500 : hsa_insn_basic (1, BRIG_OPCODE_CBR, BRIG_TYPE_B1, ctrl),
1501 m_width (BRIG_WIDTH_1)
1505 /* New operator to allocate branch instruction from pool alloc. */
1507 void *
1508 hsa_insn_br::operator new (size_t)
1510 return hsa_allocp_inst_br->allocate_raw ();
1513 /* Constructor of class representing instruction for switch jump, CTRL is
1514 the index register. */
1516 hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg *index, unsigned jump_count)
1517 : hsa_insn_basic (1, BRIG_OPCODE_SBR, BRIG_TYPE_B1, index),
1518 m_width (BRIG_WIDTH_1), m_jump_table (vNULL),
1519 m_label_code_list (new hsa_op_code_list (jump_count))
1523 /* New operator to allocate switch branch instruction from pool alloc. */
1525 void *
1526 hsa_insn_sbr::operator new (size_t)
1528 return hsa_allocp_inst_sbr->allocate_raw ();
1531 /* Replace all occurrences of OLD_BB with NEW_BB in the statements
1532 jump table. */
1534 void
1535 hsa_insn_sbr::replace_all_labels (basic_block old_bb, basic_block new_bb)
1537 for (unsigned i = 0; i < m_jump_table.length (); i++)
1538 if (m_jump_table[i] == old_bb)
1539 m_jump_table[i] = new_bb;
1542 hsa_insn_sbr::~hsa_insn_sbr ()
1544 m_jump_table.release ();
1547 /* Constructor of comparison instruction. CMP is the comparison operation and T
1548 is the result type. */
1550 hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp, BrigType16_t t,
1551 hsa_op_base *arg0, hsa_op_base *arg1,
1552 hsa_op_base *arg2)
1553 : hsa_insn_basic (3 , BRIG_OPCODE_CMP, t, arg0, arg1, arg2), m_compare (cmp)
1557 /* New operator to allocate compare instruction from pool alloc. */
1559 void *
1560 hsa_insn_cmp::operator new (size_t)
1562 return hsa_allocp_inst_cmp->allocate_raw ();
1565 /* Constructor of classes representing memory accesses. OPC is the opcode (must
1566 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
1567 operands are provided as ARG0 and ARG1. */
1569 hsa_insn_mem::hsa_insn_mem (int opc, BrigType16_t t, hsa_op_base *arg0,
1570 hsa_op_base *arg1)
1571 : hsa_insn_basic (2, opc, t, arg0, arg1),
1572 m_align (hsa_natural_alignment (t)), m_equiv_class (0)
1574 gcc_checking_assert (opc == BRIG_OPCODE_LD || opc == BRIG_OPCODE_ST);
1577 /* Constructor for descendants allowing different opcodes and number of
1578 operands, it passes its arguments directly to hsa_insn_basic
1579 constructor. The instruction operands are provided as ARG[0-3]. */
1582 hsa_insn_mem::hsa_insn_mem (unsigned nops, int opc, BrigType16_t t,
1583 hsa_op_base *arg0, hsa_op_base *arg1,
1584 hsa_op_base *arg2, hsa_op_base *arg3)
1585 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
1586 m_align (hsa_natural_alignment (t)), m_equiv_class (0)
1590 /* New operator to allocate memory instruction from pool alloc. */
1592 void *
1593 hsa_insn_mem::operator new (size_t)
1595 return hsa_allocp_inst_mem->allocate_raw ();
1598 /* Constructor of class representing atomic instructions and signals. OPC is
1599 the principal opcode, aop is the specific atomic operation opcode. T is the
1600 type of the instruction. The instruction operands
1601 are provided as ARG[0-3]. */
1603 hsa_insn_atomic::hsa_insn_atomic (int nops, int opc,
1604 enum BrigAtomicOperation aop,
1605 BrigType16_t t, BrigMemoryOrder memorder,
1606 hsa_op_base *arg0,
1607 hsa_op_base *arg1, hsa_op_base *arg2,
1608 hsa_op_base *arg3)
1609 : hsa_insn_mem (nops, opc, t, arg0, arg1, arg2, arg3), m_atomicop (aop),
1610 m_memoryorder (memorder),
1611 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM)
1613 gcc_checking_assert (opc == BRIG_OPCODE_ATOMICNORET ||
1614 opc == BRIG_OPCODE_ATOMIC ||
1615 opc == BRIG_OPCODE_SIGNAL ||
1616 opc == BRIG_OPCODE_SIGNALNORET);
1619 /* New operator to allocate signal instruction from pool alloc. */
1621 void *
1622 hsa_insn_atomic::operator new (size_t)
1624 return hsa_allocp_inst_atomic->allocate_raw ();
1627 /* Constructor of class representing signal instructions. OPC is the prinicpal
1628 opcode, sop is the specific signal operation opcode. T is the type of the
1629 instruction. The instruction operands are provided as ARG[0-3]. */
1631 hsa_insn_signal::hsa_insn_signal (int nops, int opc,
1632 enum BrigAtomicOperation sop,
1633 BrigType16_t t, hsa_op_base *arg0,
1634 hsa_op_base *arg1, hsa_op_base *arg2,
1635 hsa_op_base *arg3)
1636 : hsa_insn_atomic (nops, opc, sop, t, BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE,
1637 arg0, arg1, arg2, arg3)
1641 /* New operator to allocate signal instruction from pool alloc. */
1643 void *
1644 hsa_insn_signal::operator new (size_t)
1646 return hsa_allocp_inst_signal->allocate_raw ();
1649 /* Constructor of class representing segment conversion instructions. OPC is
1650 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
1651 and SRCT are destination and source types respectively, SEG is the segment
1652 we are converting to or from. The instruction operands are
1653 provided as ARG0 and ARG1. */
1655 hsa_insn_seg::hsa_insn_seg (int opc, BrigType16_t dest, BrigType16_t srct,
1656 BrigSegment8_t seg, hsa_op_base *arg0,
1657 hsa_op_base *arg1)
1658 : hsa_insn_basic (2, opc, dest, arg0, arg1), m_src_type (srct),
1659 m_segment (seg)
1661 gcc_checking_assert (opc == BRIG_OPCODE_STOF || opc == BRIG_OPCODE_FTOS);
1664 /* New operator to allocate address conversion instruction from pool alloc. */
1666 void *
1667 hsa_insn_seg::operator new (size_t)
1669 return hsa_allocp_inst_seg->allocate_raw ();
1672 /* Constructor of class representing a call instruction. CALLEE is the tree
1673 representation of the function being called. */
1675 hsa_insn_call::hsa_insn_call (tree callee)
1676 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (callee),
1677 m_output_arg (NULL), m_args_code_list (NULL), m_result_code_list (NULL)
1681 hsa_insn_call::hsa_insn_call (hsa_internal_fn *fn)
1682 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (NULL),
1683 m_called_internal_fn (fn), m_output_arg (NULL), m_args_code_list (NULL),
1684 m_result_code_list (NULL)
1688 /* New operator to allocate call instruction from pool alloc. */
1690 void *
1691 hsa_insn_call::operator new (size_t)
1693 return hsa_allocp_inst_call->allocate_raw ();
1696 hsa_insn_call::~hsa_insn_call ()
1698 for (unsigned i = 0; i < m_input_args.length (); i++)
1699 delete m_input_args[i];
1701 delete m_output_arg;
1703 m_input_args.release ();
1704 m_input_arg_insns.release ();
1707 /* Constructor of class representing the argument block required to invoke
1708 a call in HSAIL. */
1709 hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind,
1710 hsa_insn_call * call)
1711 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK), m_kind (brig_kind),
1712 m_call_insn (call)
1716 /* New operator to allocate argument block instruction from pool alloc. */
1718 void *
1719 hsa_insn_arg_block::operator new (size_t)
1721 return hsa_allocp_inst_arg_block->allocate_raw ();
1724 hsa_insn_comment::hsa_insn_comment (const char *s)
1725 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT)
1727 unsigned l = strlen (s);
1729 /* Append '// ' to the string. */
1730 char *buf = XNEWVEC (char, l + 4);
1731 sprintf (buf, "// %s", s);
1732 m_comment = buf;
1735 /* New operator to allocate comment instruction from pool alloc. */
1737 void *
1738 hsa_insn_comment::operator new (size_t)
1740 return hsa_allocp_inst_comment->allocate_raw ();
1743 hsa_insn_comment::~hsa_insn_comment ()
1745 gcc_checking_assert (m_comment);
1746 free (m_comment);
1747 m_comment = NULL;
1750 /* Constructor of class representing the queue instruction in HSAIL. */
1751 hsa_insn_queue::hsa_insn_queue (int nops, BrigOpcode opcode)
1752 : hsa_insn_basic (nops, opcode, BRIG_TYPE_U64)
1756 /* New operator to allocate source type instruction from pool alloc. */
1758 void *
1759 hsa_insn_srctype::operator new (size_t)
1761 return hsa_allocp_inst_srctype->allocate_raw ();
1764 /* Constructor of class representing the source type instruction in HSAIL. */
1766 hsa_insn_srctype::hsa_insn_srctype (int nops, BrigOpcode opcode,
1767 BrigType16_t destt, BrigType16_t srct,
1768 hsa_op_base *arg0, hsa_op_base *arg1,
1769 hsa_op_base *arg2 = NULL)
1770 : hsa_insn_basic (nops, opcode, destt, arg0, arg1, arg2),
1771 m_source_type (srct)
1774 /* New operator to allocate packed instruction from pool alloc. */
1776 void *
1777 hsa_insn_packed::operator new (size_t)
1779 return hsa_allocp_inst_packed->allocate_raw ();
1782 /* Constructor of class representing the packed instruction in HSAIL. */
1784 hsa_insn_packed::hsa_insn_packed (int nops, BrigOpcode opcode,
1785 BrigType16_t destt, BrigType16_t srct,
1786 hsa_op_base *arg0, hsa_op_base *arg1,
1787 hsa_op_base *arg2)
1788 : hsa_insn_srctype (nops, opcode, destt, srct, arg0, arg1, arg2)
1790 m_operand_list = new hsa_op_operand_list (nops - 1);
1793 /* New operator to allocate convert instruction from pool alloc. */
1795 void *
1796 hsa_insn_cvt::operator new (size_t)
1798 return hsa_allocp_inst_cvt->allocate_raw ();
1801 /* Constructor of class representing the convert instruction in HSAIL. */
1803 hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type *dest, hsa_op_with_type *src)
1804 : hsa_insn_basic (2, BRIG_OPCODE_CVT, dest->m_type, dest, src)
1808 /* New operator to allocate alloca from pool alloc. */
1810 void *
1811 hsa_insn_alloca::operator new (size_t)
1813 return hsa_allocp_inst_alloca->allocate_raw ();
1816 /* Constructor of class representing the alloca in HSAIL. */
1818 hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type *dest,
1819 hsa_op_with_type *size, unsigned alignment)
1820 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA, dest->m_type, dest, size),
1821 m_align (BRIG_ALIGNMENT_8)
1823 gcc_assert (dest->m_type == BRIG_TYPE_U32);
1824 if (alignment)
1825 m_align = hsa_alignment_encoding (alignment);
1828 /* Append an instruction INSN into the basic block. */
1830 void
1831 hsa_bb::append_insn (hsa_insn_basic *insn)
1833 gcc_assert (insn->m_opcode != 0 || insn->operand_count () == 0);
1834 gcc_assert (!insn->m_bb);
1836 insn->m_bb = m_bb;
1837 insn->m_prev = m_last_insn;
1838 insn->m_next = NULL;
1839 if (m_last_insn)
1840 m_last_insn->m_next = insn;
1841 m_last_insn = insn;
1842 if (!m_first_insn)
1843 m_first_insn = insn;
1846 /* Insert HSA instruction NEW_INSN immediately before an existing instruction
1847 OLD_INSN. */
1849 static void
1850 hsa_insert_insn_before (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
1852 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
1854 if (hbb->m_first_insn == old_insn)
1855 hbb->m_first_insn = new_insn;
1856 new_insn->m_prev = old_insn->m_prev;
1857 new_insn->m_next = old_insn;
1858 if (old_insn->m_prev)
1859 old_insn->m_prev->m_next = new_insn;
1860 old_insn->m_prev = new_insn;
1863 /* Append HSA instruction NEW_INSN immediately after an existing instruction
1864 OLD_INSN. */
1866 static void
1867 hsa_append_insn_after (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
1869 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
1871 if (hbb->m_last_insn == old_insn)
1872 hbb->m_last_insn = new_insn;
1873 new_insn->m_prev = old_insn;
1874 new_insn->m_next = old_insn->m_next;
1875 if (old_insn->m_next)
1876 old_insn->m_next->m_prev = new_insn;
1877 old_insn->m_next = new_insn;
1880 /* Return a register containing the calculated value of EXP which must be an
1881 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
1882 integer constants as returned by get_inner_reference.
1883 Newly generated HSA instructions will be appended to HBB.
1884 Perform all calculations in ADDRTYPE. */
1886 static hsa_op_with_type *
1887 gen_address_calculation (tree exp, hsa_bb *hbb, BrigType16_t addrtype)
1889 int opcode;
1891 if (TREE_CODE (exp) == NOP_EXPR)
1892 exp = TREE_OPERAND (exp, 0);
1894 switch (TREE_CODE (exp))
1896 case SSA_NAME:
1897 return hsa_cfun->reg_for_gimple_ssa (exp)->get_in_type (addrtype, hbb);
1899 case INTEGER_CST:
1901 hsa_op_immed *imm = new hsa_op_immed (exp);
1902 if (addrtype != imm->m_type)
1903 imm->m_type = addrtype;
1904 return imm;
1907 case PLUS_EXPR:
1908 opcode = BRIG_OPCODE_ADD;
1909 break;
1911 case MULT_EXPR:
1912 opcode = BRIG_OPCODE_MUL;
1913 break;
1915 default:
1916 gcc_unreachable ();
1919 hsa_op_reg *res = new hsa_op_reg (addrtype);
1920 hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, addrtype);
1921 insn->set_op (0, res);
1923 hsa_op_with_type *op1 = gen_address_calculation (TREE_OPERAND (exp, 0), hbb,
1924 addrtype);
1925 hsa_op_with_type *op2 = gen_address_calculation (TREE_OPERAND (exp, 1), hbb,
1926 addrtype);
1927 insn->set_op (1, op1);
1928 insn->set_op (2, op2);
1930 hbb->append_insn (insn);
1931 return res;
1934 /* If R1 is NULL, just return R2, otherwise append an instruction adding them
1935 to HBB and return the register holding the result. */
1937 static hsa_op_reg *
1938 add_addr_regs_if_needed (hsa_op_reg *r1, hsa_op_reg *r2, hsa_bb *hbb)
1940 gcc_checking_assert (r2);
1941 if (!r1)
1942 return r2;
1944 hsa_op_reg *res = new hsa_op_reg (r1->m_type);
1945 gcc_assert (!hsa_needs_cvt (r1->m_type, r2->m_type));
1946 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_ADD, res->m_type);
1947 insn->set_op (0, res);
1948 insn->set_op (1, r1);
1949 insn->set_op (2, r2);
1950 hbb->append_insn (insn);
1951 return res;
1954 /* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
1955 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
1957 static void
1958 process_mem_base (tree base, hsa_symbol **symbol, BrigType16_t *addrtype,
1959 hsa_op_reg **reg, offset_int *offset, hsa_bb *hbb)
1961 if (TREE_CODE (base) == SSA_NAME)
1963 gcc_assert (!*reg);
1964 hsa_op_with_type *ssa
1965 = hsa_cfun->reg_for_gimple_ssa (base)->get_in_type (*addrtype, hbb);
1966 *reg = dyn_cast <hsa_op_reg *> (ssa);
1968 else if (TREE_CODE (base) == ADDR_EXPR)
1970 tree decl = TREE_OPERAND (base, 0);
1972 if (!DECL_P (decl) || TREE_CODE (decl) == FUNCTION_DECL)
1974 HSA_SORRY_AT (EXPR_LOCATION (base),
1975 "support for HSA does not implement a memory reference "
1976 "to a non-declaration type");
1977 return;
1980 gcc_assert (!*symbol);
1982 *symbol = get_symbol_for_decl (decl);
1983 *addrtype = hsa_get_segment_addr_type ((*symbol)->m_segment);
1985 else if (TREE_CODE (base) == INTEGER_CST)
1986 *offset += wi::to_offset (base);
1987 else
1988 gcc_unreachable ();
1991 /* Forward declaration of a function. */
1993 static void
1994 gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb);
1996 /* Generate HSA address operand for a given tree memory reference REF. If
1997 instructions need to be created to calculate the address, they will be added
1998 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
1999 the function assumes that the caller will handle possible
2000 bit-field references. Otherwise if we reference a bit-field, sorry message
2001 is displayed. */
2003 static hsa_op_address *
2004 gen_hsa_addr (tree ref, hsa_bb *hbb, HOST_WIDE_INT *output_bitsize = NULL,
2005 HOST_WIDE_INT *output_bitpos = NULL)
2007 hsa_symbol *symbol = NULL;
2008 hsa_op_reg *reg = NULL;
2009 offset_int offset = 0;
2010 tree origref = ref;
2011 tree varoffset = NULL_TREE;
2012 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
2013 HOST_WIDE_INT bitsize = 0, bitpos = 0;
2014 BrigType16_t flat_addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
2016 if (TREE_CODE (ref) == STRING_CST)
2018 symbol = hsa_get_string_cst_symbol (ref);
2019 goto out;
2021 else if (TREE_CODE (ref) == BIT_FIELD_REF
2022 && ((tree_to_uhwi (TREE_OPERAND (ref, 1)) % BITS_PER_UNIT) != 0
2023 || (tree_to_uhwi (TREE_OPERAND (ref, 2)) % BITS_PER_UNIT) != 0))
2025 HSA_SORRY_ATV (EXPR_LOCATION (origref),
2026 "support for HSA does not implement "
2027 "bit field references such as %E", ref);
2028 goto out;
2031 if (handled_component_p (ref))
2033 enum machine_mode mode;
2034 int unsignedp, volatilep, preversep;
2036 ref = get_inner_reference (ref, &bitsize, &bitpos, &varoffset, &mode,
2037 &unsignedp, &preversep, &volatilep, false);
2039 offset = bitpos;
2040 offset = wi::rshift (offset, LOG2_BITS_PER_UNIT, SIGNED);
2043 switch (TREE_CODE (ref))
2045 case ADDR_EXPR:
2047 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
2048 symbol = hsa_cfun->create_hsa_temporary (flat_addrtype);
2049 hsa_op_reg *r = new hsa_op_reg (flat_addrtype);
2050 gen_hsa_addr_insns (ref, r, hbb);
2051 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
2052 r, new hsa_op_address (symbol)));
2054 break;
2056 case SSA_NAME:
2058 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
2059 symbol = hsa_cfun->create_hsa_temporary (flat_addrtype);
2060 hsa_op_reg *r = hsa_cfun->reg_for_gimple_ssa (ref);
2062 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
2063 r, new hsa_op_address (symbol)));
2065 break;
2067 case PARM_DECL:
2068 case VAR_DECL:
2069 case RESULT_DECL:
2070 gcc_assert (!symbol);
2071 symbol = get_symbol_for_decl (ref);
2072 addrtype = hsa_get_segment_addr_type (symbol->m_segment);
2073 break;
2075 case MEM_REF:
2076 process_mem_base (TREE_OPERAND (ref, 0), &symbol, &addrtype, &reg,
2077 &offset, hbb);
2079 if (!integer_zerop (TREE_OPERAND (ref, 1)))
2080 offset += wi::to_offset (TREE_OPERAND (ref, 1));
2081 break;
2083 case TARGET_MEM_REF:
2084 process_mem_base (TMR_BASE (ref), &symbol, &addrtype, &reg, &offset, hbb);
2085 if (TMR_INDEX (ref))
2087 hsa_op_reg *disp1;
2088 hsa_op_base *idx = hsa_cfun->reg_for_gimple_ssa
2089 (TMR_INDEX (ref))->get_in_type (addrtype, hbb);
2090 if (TMR_STEP (ref) && !integer_onep (TMR_STEP (ref)))
2092 disp1 = new hsa_op_reg (addrtype);
2093 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_MUL,
2094 addrtype);
2096 /* As step must respect addrtype, we overwrite the type
2097 of an immediate value. */
2098 hsa_op_immed *step = new hsa_op_immed (TMR_STEP (ref));
2099 step->m_type = addrtype;
2101 insn->set_op (0, disp1);
2102 insn->set_op (1, idx);
2103 insn->set_op (2, step);
2104 hbb->append_insn (insn);
2106 else
2107 disp1 = as_a <hsa_op_reg *> (idx);
2108 reg = add_addr_regs_if_needed (reg, disp1, hbb);
2110 if (TMR_INDEX2 (ref))
2112 if (TREE_CODE (TMR_INDEX2 (ref)) == SSA_NAME)
2114 hsa_op_base *disp2 = hsa_cfun->reg_for_gimple_ssa
2115 (TMR_INDEX2 (ref))->get_in_type (addrtype, hbb);
2116 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (disp2),
2117 hbb);
2119 else if (TREE_CODE (TMR_INDEX2 (ref)) == INTEGER_CST)
2120 offset += wi::to_offset (TMR_INDEX2 (ref));
2121 else
2122 gcc_unreachable ();
2124 offset += wi::to_offset (TMR_OFFSET (ref));
2125 break;
2126 case FUNCTION_DECL:
2127 HSA_SORRY_AT (EXPR_LOCATION (origref),
2128 "support for HSA does not implement function pointers");
2129 goto out;
2130 default:
2131 HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does "
2132 "not implement memory access to %E", origref);
2133 goto out;
2136 if (varoffset)
2138 if (TREE_CODE (varoffset) == INTEGER_CST)
2139 offset += wi::to_offset (varoffset);
2140 else
2142 hsa_op_base *off_op = gen_address_calculation (varoffset, hbb,
2143 addrtype);
2144 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (off_op),
2145 hbb);
2149 gcc_checking_assert ((symbol
2150 && addrtype
2151 == hsa_get_segment_addr_type (symbol->m_segment))
2152 || (!symbol
2153 && addrtype
2154 == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT)));
2155 out:
2156 HOST_WIDE_INT hwi_offset = offset.to_shwi ();
2158 /* Calculate remaining bitsize offset (if presented). */
2159 bitpos %= BITS_PER_UNIT;
2160 /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
2161 is not a reason to think this is a bit-field access. */
2162 if (bitpos == 0
2163 && (bitsize >= BITS_PER_UNIT)
2164 && !(bitsize & (bitsize - 1)))
2165 bitsize = 0;
2167 if ((bitpos || bitsize) && (output_bitpos == NULL || output_bitsize == NULL))
2168 HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does not "
2169 "implement unhandled bit field reference such as %E", ref);
2171 if (output_bitsize != NULL && output_bitpos != NULL)
2173 *output_bitsize = bitsize;
2174 *output_bitpos = bitpos;
2177 return new hsa_op_address (symbol, reg, hwi_offset);
2180 /* Generate HSA address operand for a given tree memory reference REF. If
2181 instructions need to be created to calculate the address, they will be added
2182 to the end of HBB. OUTPUT_ALIGN is alignment of the created address. */
2184 static hsa_op_address *
2185 gen_hsa_addr_with_align (tree ref, hsa_bb *hbb, BrigAlignment8_t *output_align)
2187 hsa_op_address *addr = gen_hsa_addr (ref, hbb);
2188 if (addr->m_reg || !addr->m_symbol)
2189 *output_align = hsa_object_alignment (ref);
2190 else
2192 /* If the address consists only of a symbol and an offset, we
2193 compute the alignment ourselves to take into account any alignment
2194 promotions we might have done for the HSA symbol representation. */
2195 unsigned align = hsa_byte_alignment (addr->m_symbol->m_align);
2196 unsigned misalign = addr->m_imm_offset & (align - 1);
2197 if (misalign)
2198 align = (misalign & -misalign);
2199 *output_align = hsa_alignment_encoding (BITS_PER_UNIT * align);
2201 return addr;
2204 /* Generate HSA address for a function call argument of given TYPE.
2205 INDEX is used to generate corresponding name of the arguments.
2206 Special value -1 represents fact that result value is created. */
2208 static hsa_op_address *
2209 gen_hsa_addr_for_arg (tree tree_type, int index)
2211 hsa_symbol *sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
2212 BRIG_LINKAGE_ARG);
2213 sym->m_type = hsa_type_for_tree_type (tree_type, &sym->m_dim);
2215 if (index == -1) /* Function result. */
2216 sym->m_name = "res";
2217 else /* Function call arguments. */
2219 sym->m_name = NULL;
2220 sym->m_name_number = index;
2223 return new hsa_op_address (sym);
2226 /* Generate HSA instructions that process all necessary conversions
2227 of an ADDR to flat addressing and place the result into DEST.
2228 Instructions are appended to HBB. */
2230 static void
2231 convert_addr_to_flat_segment (hsa_op_address *addr, hsa_op_reg *dest,
2232 hsa_bb *hbb)
2234 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_LDA);
2235 insn->set_op (1, addr);
2236 if (addr->m_symbol && addr->m_symbol->m_segment != BRIG_SEGMENT_GLOBAL)
2238 /* LDA produces segment-relative address, we need to convert
2239 it to the flat one. */
2240 hsa_op_reg *tmp;
2241 tmp = new hsa_op_reg (hsa_get_segment_addr_type
2242 (addr->m_symbol->m_segment));
2243 hsa_insn_seg *seg;
2244 seg = new hsa_insn_seg (BRIG_OPCODE_STOF,
2245 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT),
2246 tmp->m_type, addr->m_symbol->m_segment, dest,
2247 tmp);
2249 insn->set_op (0, tmp);
2250 insn->m_type = tmp->m_type;
2251 hbb->append_insn (insn);
2252 hbb->append_insn (seg);
2254 else
2256 insn->set_op (0, dest);
2257 insn->m_type = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
2258 hbb->append_insn (insn);
2262 /* Generate HSA instructions that calculate address of VAL including all
2263 necessary conversions to flat addressing and place the result into DEST.
2264 Instructions are appended to HBB. */
2266 static void
2267 gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb)
2269 /* Handle cases like tmp = NULL, where we just emit a move instruction
2270 to a register. */
2271 if (TREE_CODE (val) == INTEGER_CST)
2273 hsa_op_immed *c = new hsa_op_immed (val);
2274 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
2275 dest->m_type, dest, c);
2276 hbb->append_insn (insn);
2277 return;
2280 hsa_op_address *addr;
2282 gcc_assert (dest->m_type == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
2283 if (TREE_CODE (val) == ADDR_EXPR)
2284 val = TREE_OPERAND (val, 0);
2285 addr = gen_hsa_addr (val, hbb);
2287 convert_addr_to_flat_segment (addr, dest, hbb);
2290 /* Return an HSA register or HSA immediate value operand corresponding to
2291 gimple operand OP. */
2293 static hsa_op_with_type *
2294 hsa_reg_or_immed_for_gimple_op (tree op, hsa_bb *hbb)
2296 hsa_op_reg *tmp;
2298 if (TREE_CODE (op) == SSA_NAME)
2299 tmp = hsa_cfun->reg_for_gimple_ssa (op);
2300 else if (!POINTER_TYPE_P (TREE_TYPE (op)))
2301 return new hsa_op_immed (op);
2302 else
2304 tmp = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
2305 gen_hsa_addr_insns (op, tmp, hbb);
2307 return tmp;
2310 /* Create a simple movement instruction with register destination DEST and
2311 register or immediate source SRC and append it to the end of HBB. */
2313 void
2314 hsa_build_append_simple_mov (hsa_op_reg *dest, hsa_op_base *src, hsa_bb *hbb)
2316 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type,
2317 dest, src);
2318 if (hsa_op_reg *sreg = dyn_cast <hsa_op_reg *> (src))
2319 gcc_assert (hsa_type_bit_size (dest->m_type)
2320 == hsa_type_bit_size (sreg->m_type));
2321 else
2322 gcc_assert (hsa_type_bit_size (dest->m_type)
2323 == hsa_type_bit_size (as_a <hsa_op_immed *> (src)->m_type));
2325 hbb->append_insn (insn);
2328 /* Generate HSAIL instructions loading a bit field into register DEST.
2329 VALUE_REG is a register of a SSA name that is used in the bit field
2330 reference. To identify a bit field BITPOS is offset to the loaded memory
2331 and BITSIZE is number of bits of the bit field.
2332 Add instructions to HBB. */
2334 static void
2335 gen_hsa_insns_for_bitfield (hsa_op_reg *dest, hsa_op_reg *value_reg,
2336 HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos,
2337 hsa_bb *hbb)
2339 unsigned type_bitsize = hsa_type_bit_size (dest->m_type);
2340 unsigned left_shift = type_bitsize - (bitsize + bitpos);
2341 unsigned right_shift = left_shift + bitpos;
2343 if (left_shift)
2345 hsa_op_reg *value_reg_2 = new hsa_op_reg (dest->m_type);
2346 hsa_op_immed *c = new hsa_op_immed (left_shift, BRIG_TYPE_U32);
2348 hsa_insn_basic *lshift
2349 = new hsa_insn_basic (3, BRIG_OPCODE_SHL, value_reg_2->m_type,
2350 value_reg_2, value_reg, c);
2352 hbb->append_insn (lshift);
2354 value_reg = value_reg_2;
2357 if (right_shift)
2359 hsa_op_reg *value_reg_2 = new hsa_op_reg (dest->m_type);
2360 hsa_op_immed *c = new hsa_op_immed (right_shift, BRIG_TYPE_U32);
2362 hsa_insn_basic *rshift
2363 = new hsa_insn_basic (3, BRIG_OPCODE_SHR, value_reg_2->m_type,
2364 value_reg_2, value_reg, c);
2366 hbb->append_insn (rshift);
2368 value_reg = value_reg_2;
2371 hsa_insn_basic *assignment
2372 = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, dest, value_reg);
2373 hbb->append_insn (assignment);
2377 /* Generate HSAIL instructions loading a bit field into register DEST. ADDR is
2378 prepared memory address which is used to load the bit field. To identify a
2379 bit field BITPOS is offset to the loaded memory and BITSIZE is number of
2380 bits of the bit field. Add instructions to HBB. Load must be performed in
2381 alignment ALIGN. */
2383 static void
2384 gen_hsa_insns_for_bitfield_load (hsa_op_reg *dest, hsa_op_address *addr,
2385 HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos,
2386 hsa_bb *hbb, BrigAlignment8_t align)
2388 hsa_op_reg *value_reg = new hsa_op_reg (dest->m_type);
2389 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, dest->m_type, value_reg,
2390 addr);
2391 mem->set_align (align);
2392 hbb->append_insn (mem);
2393 gen_hsa_insns_for_bitfield (dest, value_reg, bitsize, bitpos, hbb);
2396 /* Return the alignment of base memory accesses we issue to perform bit-field
2397 memory access REF. */
2399 static BrigAlignment8_t
2400 hsa_bitmemref_alignment (tree ref)
2402 unsigned HOST_WIDE_INT bit_offset = 0;
2404 while (true)
2406 if (TREE_CODE (ref) == BIT_FIELD_REF)
2408 if (!tree_fits_uhwi_p (TREE_OPERAND (ref, 2)))
2409 return BRIG_ALIGNMENT_1;
2410 bit_offset += tree_to_uhwi (TREE_OPERAND (ref, 2));
2412 else if (TREE_CODE (ref) == COMPONENT_REF
2413 && DECL_BIT_FIELD (TREE_OPERAND (ref, 1)))
2414 bit_offset += int_bit_position (TREE_OPERAND (ref, 1));
2415 else
2416 break;
2417 ref = TREE_OPERAND (ref, 0);
2420 unsigned HOST_WIDE_INT bits = bit_offset % BITS_PER_UNIT;
2421 unsigned HOST_WIDE_INT byte_bits = bit_offset - bits;
2422 BrigAlignment8_t base = hsa_object_alignment (ref);
2423 if (byte_bits == 0)
2424 return base;
2425 return MIN (base, hsa_alignment_encoding (byte_bits & -byte_bits));
2428 /* Generate HSAIL instructions loading something into register DEST. RHS is
2429 tree representation of the loaded data, which are loaded as type TYPE. Add
2430 instructions to HBB. */
2432 static void
2433 gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb)
2435 /* The destination SSA name will give us the type. */
2436 if (TREE_CODE (rhs) == VIEW_CONVERT_EXPR)
2437 rhs = TREE_OPERAND (rhs, 0);
2439 if (TREE_CODE (rhs) == SSA_NAME)
2441 hsa_op_reg *src = hsa_cfun->reg_for_gimple_ssa (rhs);
2442 hsa_build_append_simple_mov (dest, src, hbb);
2444 else if (is_gimple_min_invariant (rhs)
2445 || TREE_CODE (rhs) == ADDR_EXPR)
2447 if (POINTER_TYPE_P (TREE_TYPE (rhs)))
2449 if (dest->m_type != hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT))
2451 HSA_SORRY_ATV (EXPR_LOCATION (rhs),
2452 "support for HSA does not implement conversion "
2453 "of %E to the requested non-pointer type.", rhs);
2454 return;
2457 gen_hsa_addr_insns (rhs, dest, hbb);
2459 else if (TREE_CODE (rhs) == COMPLEX_CST)
2461 hsa_op_immed *real_part = new hsa_op_immed (TREE_REALPART (rhs));
2462 hsa_op_immed *imag_part = new hsa_op_immed (TREE_IMAGPART (rhs));
2464 hsa_op_reg *real_part_reg
2465 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type),
2466 true));
2467 hsa_op_reg *imag_part_reg
2468 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type),
2469 true));
2471 hsa_build_append_simple_mov (real_part_reg, real_part, hbb);
2472 hsa_build_append_simple_mov (imag_part_reg, imag_part, hbb);
2474 BrigType16_t src_type = hsa_bittype_for_type (real_part_reg->m_type);
2476 hsa_insn_packed *insn
2477 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type,
2478 src_type, dest, real_part_reg,
2479 imag_part_reg);
2480 hbb->append_insn (insn);
2482 else
2484 hsa_op_immed *imm = new hsa_op_immed (rhs);
2485 hsa_build_append_simple_mov (dest, imm, hbb);
2488 else if (TREE_CODE (rhs) == REALPART_EXPR || TREE_CODE (rhs) == IMAGPART_EXPR)
2490 tree pack_type = TREE_TYPE (TREE_OPERAND (rhs, 0));
2492 hsa_op_reg *packed_reg
2493 = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type, true));
2495 tree complex_rhs = TREE_OPERAND (rhs, 0);
2496 gen_hsa_insns_for_load (packed_reg, complex_rhs, TREE_TYPE (complex_rhs),
2497 hbb);
2499 hsa_op_reg *real_reg
2500 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true));
2502 hsa_op_reg *imag_reg
2503 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true));
2505 BrigKind16_t brig_type = packed_reg->m_type;
2506 hsa_insn_packed *packed
2507 = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND,
2508 hsa_bittype_for_type (real_reg->m_type),
2509 brig_type, real_reg, imag_reg, packed_reg);
2511 hbb->append_insn (packed);
2513 hsa_op_reg *source = TREE_CODE (rhs) == REALPART_EXPR ?
2514 real_reg : imag_reg;
2516 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
2517 dest->m_type, dest, source);
2519 hbb->append_insn (insn);
2521 else if (TREE_CODE (rhs) == BIT_FIELD_REF
2522 && TREE_CODE (TREE_OPERAND (rhs, 0)) == SSA_NAME)
2524 tree ssa_name = TREE_OPERAND (rhs, 0);
2525 HOST_WIDE_INT bitsize = tree_to_uhwi (TREE_OPERAND (rhs, 1));
2526 HOST_WIDE_INT bitpos = tree_to_uhwi (TREE_OPERAND (rhs, 2));
2528 hsa_op_reg *imm_value = hsa_cfun->reg_for_gimple_ssa (ssa_name);
2529 gen_hsa_insns_for_bitfield (dest, imm_value, bitsize, bitpos, hbb);
2531 else if (DECL_P (rhs) || TREE_CODE (rhs) == MEM_REF
2532 || TREE_CODE (rhs) == TARGET_MEM_REF
2533 || handled_component_p (rhs))
2535 HOST_WIDE_INT bitsize, bitpos;
2537 /* Load from memory. */
2538 hsa_op_address *addr;
2539 addr = gen_hsa_addr (rhs, hbb, &bitsize, &bitpos);
2541 /* Handle load of a bit field. */
2542 if (bitsize > 64)
2544 HSA_SORRY_AT (EXPR_LOCATION (rhs),
2545 "support for HSA does not implement load from a bit "
2546 "field bigger than 64 bits");
2547 return;
2550 if (bitsize || bitpos)
2551 gen_hsa_insns_for_bitfield_load (dest, addr, bitsize, bitpos, hbb,
2552 hsa_bitmemref_alignment (rhs));
2553 else
2555 BrigType16_t mtype;
2556 /* Not dest->m_type, that's possibly extended. */
2557 mtype = mem_type_for_type (hsa_type_for_scalar_tree_type (type,
2558 false));
2559 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, dest,
2560 addr);
2561 mem->set_align (hsa_object_alignment (rhs));
2562 hbb->append_insn (mem);
2565 else
2566 HSA_SORRY_ATV (EXPR_LOCATION (rhs),
2567 "support for HSA does not implement loading "
2568 "of expression %E",
2569 rhs);
2572 /* Return number of bits necessary for representation of a bit field,
2573 starting at BITPOS with size of BITSIZE. */
2575 static unsigned
2576 get_bitfield_size (unsigned bitpos, unsigned bitsize)
2578 unsigned s = bitpos + bitsize;
2579 unsigned sizes[] = {8, 16, 32, 64};
2581 for (unsigned i = 0; i < 4; i++)
2582 if (s <= sizes[i])
2583 return sizes[i];
2585 gcc_unreachable ();
2586 return 0;
2589 /* Generate HSAIL instructions storing into memory. LHS is the destination of
2590 the store, SRC is the source operand. Add instructions to HBB. */
2592 static void
2593 gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb)
2595 HOST_WIDE_INT bitsize = 0, bitpos = 0;
2596 BrigAlignment8_t req_align;
2597 BrigType16_t mtype;
2598 mtype = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
2599 false));
2600 hsa_op_address *addr;
2601 addr = gen_hsa_addr (lhs, hbb, &bitsize, &bitpos);
2603 /* Handle store to a bit field. */
2604 if (bitsize > 64)
2606 HSA_SORRY_AT (EXPR_LOCATION (lhs),
2607 "support for HSA does not implement store to a bit field "
2608 "bigger than 64 bits");
2609 return;
2612 unsigned type_bitsize = get_bitfield_size (bitpos, bitsize);
2614 /* HSAIL does not support MOV insn with 16-bits integers. */
2615 if (type_bitsize < 32)
2616 type_bitsize = 32;
2618 if (bitpos || (bitsize && type_bitsize != bitsize))
2620 unsigned HOST_WIDE_INT mask = 0;
2621 BrigType16_t mem_type
2622 = get_integer_type_by_bytes (type_bitsize / BITS_PER_UNIT,
2623 !TYPE_UNSIGNED (TREE_TYPE (lhs)));
2625 for (unsigned i = 0; i < type_bitsize; i++)
2626 if (i < bitpos || i >= bitpos + bitsize)
2627 mask |= ((unsigned HOST_WIDE_INT)1 << i);
2629 hsa_op_reg *value_reg = new hsa_op_reg (mem_type);
2631 req_align = hsa_bitmemref_alignment (lhs);
2632 /* Load value from memory. */
2633 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mem_type,
2634 value_reg, addr);
2635 mem->set_align (req_align);
2636 hbb->append_insn (mem);
2638 /* AND the loaded value with prepared mask. */
2639 hsa_op_reg *cleared_reg = new hsa_op_reg (mem_type);
2641 BrigType16_t t
2642 = get_integer_type_by_bytes (type_bitsize / BITS_PER_UNIT, false);
2643 hsa_op_immed *c = new hsa_op_immed (mask, t);
2645 hsa_insn_basic *clearing
2646 = new hsa_insn_basic (3, BRIG_OPCODE_AND, mem_type, cleared_reg,
2647 value_reg, c);
2648 hbb->append_insn (clearing);
2650 /* Shift to left a value that is going to be stored. */
2651 hsa_op_reg *new_value_reg = new hsa_op_reg (mem_type);
2653 hsa_insn_basic *basic = new hsa_insn_basic (2, BRIG_OPCODE_MOV, mem_type,
2654 new_value_reg, src);
2655 hbb->append_insn (basic);
2657 if (bitpos)
2659 hsa_op_reg *shifted_value_reg = new hsa_op_reg (mem_type);
2660 c = new hsa_op_immed (bitpos, BRIG_TYPE_U32);
2662 hsa_insn_basic *basic
2663 = new hsa_insn_basic (3, BRIG_OPCODE_SHL, mem_type,
2664 shifted_value_reg, new_value_reg, c);
2665 hbb->append_insn (basic);
2667 new_value_reg = shifted_value_reg;
2670 /* OR the prepared value with prepared chunk loaded from memory. */
2671 hsa_op_reg *prepared_reg= new hsa_op_reg (mem_type);
2672 basic = new hsa_insn_basic (3, BRIG_OPCODE_OR, mem_type, prepared_reg,
2673 new_value_reg, cleared_reg);
2674 hbb->append_insn (basic);
2676 src = prepared_reg;
2677 mtype = mem_type;
2679 else
2680 req_align = hsa_object_alignment (lhs);
2682 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, mtype, src, addr);
2683 mem->set_align (req_align);
2685 /* The HSAIL verifier has another constraint: if the source is an immediate
2686 then it must match the destination type. If it's a register the low bits
2687 will be used for sub-word stores. We're always allocating new operands so
2688 we can modify the above in place. */
2689 if (hsa_op_immed *imm = dyn_cast <hsa_op_immed *> (src))
2691 if (!hsa_type_packed_p (imm->m_type))
2692 imm->m_type = mem->m_type;
2693 else
2695 /* ...and all vector immediates apparently need to be vectors of
2696 unsigned bytes. */
2697 unsigned bs = hsa_type_bit_size (imm->m_type);
2698 gcc_assert (bs == hsa_type_bit_size (mem->m_type));
2699 switch (bs)
2701 case 32:
2702 imm->m_type = BRIG_TYPE_U8X4;
2703 break;
2704 case 64:
2705 imm->m_type = BRIG_TYPE_U8X8;
2706 break;
2707 case 128:
2708 imm->m_type = BRIG_TYPE_U8X16;
2709 break;
2710 default:
2711 gcc_unreachable ();
2716 hbb->append_insn (mem);
2719 /* Generate memory copy instructions that are going to be used
2720 for copying a SRC memory to TARGET memory,
2721 represented by pointer in a register. MIN_ALIGN is minimal alignment
2722 of provided HSA addresses. */
2724 static void
2725 gen_hsa_memory_copy (hsa_bb *hbb, hsa_op_address *target, hsa_op_address *src,
2726 unsigned size, BrigAlignment8_t min_align)
2728 hsa_op_address *addr;
2729 hsa_insn_mem *mem;
2731 unsigned offset = 0;
2732 unsigned min_byte_align = hsa_byte_alignment (min_align);
2734 while (size)
2736 unsigned s;
2737 if (size >= 8)
2738 s = 8;
2739 else if (size >= 4)
2740 s = 4;
2741 else if (size >= 2)
2742 s = 2;
2743 else
2744 s = 1;
2746 if (s > min_byte_align)
2747 s = min_byte_align;
2749 BrigType16_t t = get_integer_type_by_bytes (s, false);
2751 hsa_op_reg *tmp = new hsa_op_reg (t);
2752 addr = new hsa_op_address (src->m_symbol, src->m_reg,
2753 src->m_imm_offset + offset);
2754 mem = new hsa_insn_mem (BRIG_OPCODE_LD, t, tmp, addr);
2755 hbb->append_insn (mem);
2757 addr = new hsa_op_address (target->m_symbol, target->m_reg,
2758 target->m_imm_offset + offset);
2759 mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, tmp, addr);
2760 hbb->append_insn (mem);
2761 offset += s;
2762 size -= s;
2766 /* Create a memset mask that is created by copying a CONSTANT byte value
2767 to an integer of BYTE_SIZE bytes. */
2769 static unsigned HOST_WIDE_INT
2770 build_memset_value (unsigned HOST_WIDE_INT constant, unsigned byte_size)
2772 if (constant == 0)
2773 return 0;
2775 HOST_WIDE_INT v = constant;
2777 for (unsigned i = 1; i < byte_size; i++)
2778 v |= constant << (8 * i);
2780 return v;
2783 /* Generate memory set instructions that are going to be used
2784 for setting a CONSTANT byte value to TARGET memory of SIZE bytes.
2785 MIN_ALIGN is minimal alignment of provided HSA addresses. */
2787 static void
2788 gen_hsa_memory_set (hsa_bb *hbb, hsa_op_address *target,
2789 unsigned HOST_WIDE_INT constant,
2790 unsigned size, BrigAlignment8_t min_align)
2792 hsa_op_address *addr;
2793 hsa_insn_mem *mem;
2795 unsigned offset = 0;
2796 unsigned min_byte_align = hsa_byte_alignment (min_align);
2798 while (size)
2800 unsigned s;
2801 if (size >= 8)
2802 s = 8;
2803 else if (size >= 4)
2804 s = 4;
2805 else if (size >= 2)
2806 s = 2;
2807 else
2808 s = 1;
2810 if (s > min_byte_align)
2811 s = min_byte_align;
2813 addr = new hsa_op_address (target->m_symbol, target->m_reg,
2814 target->m_imm_offset + offset);
2816 BrigType16_t t = get_integer_type_by_bytes (s, false);
2817 HOST_WIDE_INT c = build_memset_value (constant, s);
2819 mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, new hsa_op_immed (c, t),
2820 addr);
2821 hbb->append_insn (mem);
2822 offset += s;
2823 size -= s;
2827 /* Generate HSAIL instructions for a single assignment
2828 of an empty constructor to an ADDR_LHS. Constructor is passed as a
2829 tree RHS and all instructions are appended to HBB. ALIGN is
2830 alignment of the address. */
2832 void
2833 gen_hsa_ctor_assignment (hsa_op_address *addr_lhs, tree rhs, hsa_bb *hbb,
2834 BrigAlignment8_t align)
2836 if (vec_safe_length (CONSTRUCTOR_ELTS (rhs)))
2838 HSA_SORRY_AT (EXPR_LOCATION (rhs),
2839 "support for HSA does not implement load from constructor");
2840 return;
2843 unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs)));
2844 gen_hsa_memory_set (hbb, addr_lhs, 0, size, align);
2847 /* Generate HSA instructions for a single assignment of RHS to LHS.
2848 HBB is the basic block they will be appended to. */
2850 static void
2851 gen_hsa_insns_for_single_assignment (tree lhs, tree rhs, hsa_bb *hbb)
2853 if (TREE_CODE (lhs) == SSA_NAME)
2855 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
2856 if (hsa_seen_error ())
2857 return;
2859 gen_hsa_insns_for_load (dest, rhs, TREE_TYPE (lhs), hbb);
2861 else if (TREE_CODE (rhs) == SSA_NAME
2862 || (is_gimple_min_invariant (rhs) && TREE_CODE (rhs) != STRING_CST))
2864 /* Store to memory. */
2865 hsa_op_base *src = hsa_reg_or_immed_for_gimple_op (rhs, hbb);
2866 if (hsa_seen_error ())
2867 return;
2869 gen_hsa_insns_for_store (lhs, src, hbb);
2871 else
2873 BrigAlignment8_t lhs_align;
2874 hsa_op_address *addr_lhs = gen_hsa_addr_with_align (lhs, hbb,
2875 &lhs_align);
2877 if (TREE_CODE (rhs) == CONSTRUCTOR)
2878 gen_hsa_ctor_assignment (addr_lhs, rhs, hbb, lhs_align);
2879 else
2881 BrigAlignment8_t rhs_align;
2882 hsa_op_address *addr_rhs = gen_hsa_addr_with_align (rhs, hbb,
2883 &rhs_align);
2885 unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs)));
2886 gen_hsa_memory_copy (hbb, addr_lhs, addr_rhs, size,
2887 MIN (lhs_align, rhs_align));
2892 /* Prepend before INSN a load from spill symbol of SPILL_REG. Return the
2893 register into which we loaded. If this required another register to convert
2894 from a B1 type, return it in *PTMP2, otherwise store NULL into it. We
2895 assume we are out of SSA so the returned register does not have its
2896 definition set. */
2898 hsa_op_reg *
2899 hsa_spill_in (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2)
2901 hsa_symbol *spill_sym = spill_reg->m_spill_sym;
2902 hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type);
2903 hsa_op_address *addr = new hsa_op_address (spill_sym);
2905 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, spill_sym->m_type,
2906 reg, addr);
2907 hsa_insert_insn_before (mem, insn);
2909 *ptmp2 = NULL;
2910 if (spill_reg->m_type == BRIG_TYPE_B1)
2912 hsa_insn_basic *cvtinsn;
2913 *ptmp2 = reg;
2914 reg = new hsa_op_reg (spill_reg->m_type);
2916 cvtinsn = new hsa_insn_cvt (reg, *ptmp2);
2917 hsa_insert_insn_before (cvtinsn, insn);
2919 return reg;
2922 /* Append after INSN a store to spill symbol of SPILL_REG. Return the register
2923 from which we stored. If this required another register to convert to a B1
2924 type, return it in *PTMP2, otherwise store NULL into it. We assume we are
2925 out of SSA so the returned register does not have its use updated. */
2927 hsa_op_reg *
2928 hsa_spill_out (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2)
2930 hsa_symbol *spill_sym = spill_reg->m_spill_sym;
2931 hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type);
2932 hsa_op_address *addr = new hsa_op_address (spill_sym);
2933 hsa_op_reg *returnreg;
2935 *ptmp2 = NULL;
2936 returnreg = reg;
2937 if (spill_reg->m_type == BRIG_TYPE_B1)
2939 hsa_insn_basic *cvtinsn;
2940 *ptmp2 = new hsa_op_reg (spill_sym->m_type);
2941 reg->m_type = spill_reg->m_type;
2943 cvtinsn = new hsa_insn_cvt (*ptmp2, returnreg);
2944 hsa_append_insn_after (cvtinsn, insn);
2945 insn = cvtinsn;
2946 reg = *ptmp2;
2949 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, spill_sym->m_type, reg,
2950 addr);
2951 hsa_append_insn_after (mem, insn);
2952 return returnreg;
2955 /* Generate a comparison instruction that will compare LHS and RHS with
2956 comparison specified by CODE and put result into register DEST. DEST has to
2957 have its type set already but must not have its definition set yet.
2958 Generated instructions will be added to HBB. */
2960 static void
2961 gen_hsa_cmp_insn_from_gimple (enum tree_code code, tree lhs, tree rhs,
2962 hsa_op_reg *dest, hsa_bb *hbb)
2964 BrigCompareOperation8_t compare;
2966 switch (code)
2968 case LT_EXPR:
2969 compare = BRIG_COMPARE_LT;
2970 break;
2971 case LE_EXPR:
2972 compare = BRIG_COMPARE_LE;
2973 break;
2974 case GT_EXPR:
2975 compare = BRIG_COMPARE_GT;
2976 break;
2977 case GE_EXPR:
2978 compare = BRIG_COMPARE_GE;
2979 break;
2980 case EQ_EXPR:
2981 compare = BRIG_COMPARE_EQ;
2982 break;
2983 case NE_EXPR:
2984 compare = BRIG_COMPARE_NE;
2985 break;
2986 case UNORDERED_EXPR:
2987 compare = BRIG_COMPARE_NAN;
2988 break;
2989 case ORDERED_EXPR:
2990 compare = BRIG_COMPARE_NUM;
2991 break;
2992 case UNLT_EXPR:
2993 compare = BRIG_COMPARE_LTU;
2994 break;
2995 case UNLE_EXPR:
2996 compare = BRIG_COMPARE_LEU;
2997 break;
2998 case UNGT_EXPR:
2999 compare = BRIG_COMPARE_GTU;
3000 break;
3001 case UNGE_EXPR:
3002 compare = BRIG_COMPARE_GEU;
3003 break;
3004 case UNEQ_EXPR:
3005 compare = BRIG_COMPARE_EQU;
3006 break;
3007 case LTGT_EXPR:
3008 compare = BRIG_COMPARE_NEU;
3009 break;
3011 default:
3012 HSA_SORRY_ATV (EXPR_LOCATION (lhs),
3013 "support for HSA does not implement comparison tree "
3014 "code %s\n", get_tree_code_name (code));
3015 return;
3018 /* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer)
3019 as a result of comparison. */
3021 BrigType16_t dest_type = hsa_type_integer_p (dest->m_type)
3022 ? (BrigType16_t) BRIG_TYPE_B1 : dest->m_type;
3024 hsa_insn_cmp *cmp = new hsa_insn_cmp (compare, dest_type);
3025 cmp->set_op (1, hsa_reg_or_immed_for_gimple_op (lhs, hbb));
3026 cmp->set_op (2, hsa_reg_or_immed_for_gimple_op (rhs, hbb));
3028 hbb->append_insn (cmp);
3029 cmp->set_output_in_type (dest, 0, hbb);
3032 /* Generate an unary instruction with OPCODE and append it to a basic block
3033 HBB. The instruction uses DEST as a destination and OP1
3034 as a single operand. */
3036 static void
3037 gen_hsa_unary_operation (BrigOpcode opcode, hsa_op_reg *dest,
3038 hsa_op_with_type *op1, hsa_bb *hbb)
3040 gcc_checking_assert (dest);
3041 hsa_insn_basic *insn;
3043 if (opcode == BRIG_OPCODE_MOV && hsa_needs_cvt (dest->m_type, op1->m_type))
3044 insn = new hsa_insn_cvt (dest, op1);
3045 else if (opcode == BRIG_OPCODE_FIRSTBIT || opcode == BRIG_OPCODE_LASTBIT)
3046 insn = new hsa_insn_srctype (2, opcode, BRIG_TYPE_U32, op1->m_type, NULL,
3047 op1);
3048 else
3050 insn = new hsa_insn_basic (2, opcode, dest->m_type, dest, op1);
3052 if (opcode == BRIG_OPCODE_ABS || opcode == BRIG_OPCODE_NEG)
3054 /* ABS and NEG only exist in _s form :-/ */
3055 if (insn->m_type == BRIG_TYPE_U32)
3056 insn->m_type = BRIG_TYPE_S32;
3057 else if (insn->m_type == BRIG_TYPE_U64)
3058 insn->m_type = BRIG_TYPE_S64;
3062 hbb->append_insn (insn);
3064 if (opcode == BRIG_OPCODE_FIRSTBIT || opcode == BRIG_OPCODE_LASTBIT)
3065 insn->set_output_in_type (dest, 0, hbb);
3068 /* Generate a binary instruction with OPCODE and append it to a basic block
3069 HBB. The instruction uses DEST as a destination and operands OP1
3070 and OP2. */
3072 static void
3073 gen_hsa_binary_operation (int opcode, hsa_op_reg *dest,
3074 hsa_op_base *op1, hsa_op_base *op2, hsa_bb *hbb)
3076 gcc_checking_assert (dest);
3078 if ((opcode == BRIG_OPCODE_SHL || opcode == BRIG_OPCODE_SHR)
3079 && is_a <hsa_op_immed *> (op2))
3081 hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2);
3082 i->set_type (BRIG_TYPE_U32);
3084 if ((opcode == BRIG_OPCODE_OR
3085 || opcode == BRIG_OPCODE_XOR
3086 || opcode == BRIG_OPCODE_AND)
3087 && is_a <hsa_op_immed *> (op2))
3089 hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2);
3090 i->set_type (hsa_unsigned_type_for_type (i->m_type));
3093 hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, dest->m_type, dest,
3094 op1, op2);
3095 hbb->append_insn (insn);
3098 /* Generate HSA instructions for a single assignment. HBB is the basic block
3099 they will be appended to. */
3101 static void
3102 gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
3104 tree_code code = gimple_assign_rhs_code (assign);
3105 gimple_rhs_class rhs_class = get_gimple_rhs_class (gimple_expr_code (assign));
3107 tree lhs = gimple_assign_lhs (assign);
3108 tree rhs1 = gimple_assign_rhs1 (assign);
3109 tree rhs2 = gimple_assign_rhs2 (assign);
3110 tree rhs3 = gimple_assign_rhs3 (assign);
3112 BrigOpcode opcode;
3114 switch (code)
3116 CASE_CONVERT:
3117 case FLOAT_EXPR:
3118 /* The opcode is changed to BRIG_OPCODE_CVT if BRIG types
3119 needs a conversion. */
3120 opcode = BRIG_OPCODE_MOV;
3121 break;
3123 case PLUS_EXPR:
3124 case POINTER_PLUS_EXPR:
3125 opcode = BRIG_OPCODE_ADD;
3126 break;
3127 case MINUS_EXPR:
3128 opcode = BRIG_OPCODE_SUB;
3129 break;
3130 case MULT_EXPR:
3131 opcode = BRIG_OPCODE_MUL;
3132 break;
3133 case MULT_HIGHPART_EXPR:
3134 opcode = BRIG_OPCODE_MULHI;
3135 break;
3136 case RDIV_EXPR:
3137 case TRUNC_DIV_EXPR:
3138 case EXACT_DIV_EXPR:
3139 opcode = BRIG_OPCODE_DIV;
3140 break;
3141 case CEIL_DIV_EXPR:
3142 case FLOOR_DIV_EXPR:
3143 case ROUND_DIV_EXPR:
3144 HSA_SORRY_AT (gimple_location (assign),
3145 "support for HSA does not implement CEIL_DIV_EXPR, "
3146 "FLOOR_DIV_EXPR or ROUND_DIV_EXPR");
3147 return;
3148 case TRUNC_MOD_EXPR:
3149 opcode = BRIG_OPCODE_REM;
3150 break;
3151 case CEIL_MOD_EXPR:
3152 case FLOOR_MOD_EXPR:
3153 case ROUND_MOD_EXPR:
3154 HSA_SORRY_AT (gimple_location (assign),
3155 "support for HSA does not implement CEIL_MOD_EXPR, "
3156 "FLOOR_MOD_EXPR or ROUND_MOD_EXPR");
3157 return;
3158 case NEGATE_EXPR:
3159 opcode = BRIG_OPCODE_NEG;
3160 break;
3161 case MIN_EXPR:
3162 opcode = BRIG_OPCODE_MIN;
3163 break;
3164 case MAX_EXPR:
3165 opcode = BRIG_OPCODE_MAX;
3166 break;
3167 case ABS_EXPR:
3168 opcode = BRIG_OPCODE_ABS;
3169 break;
3170 case LSHIFT_EXPR:
3171 opcode = BRIG_OPCODE_SHL;
3172 break;
3173 case RSHIFT_EXPR:
3174 opcode = BRIG_OPCODE_SHR;
3175 break;
3176 case LROTATE_EXPR:
3177 case RROTATE_EXPR:
3179 hsa_insn_basic *insn = NULL;
3180 int code1 = code == LROTATE_EXPR ? BRIG_OPCODE_SHL : BRIG_OPCODE_SHR;
3181 int code2 = code != LROTATE_EXPR ? BRIG_OPCODE_SHL : BRIG_OPCODE_SHR;
3182 BrigType16_t btype = hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
3183 true);
3185 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3186 hsa_op_reg *op1 = new hsa_op_reg (btype);
3187 hsa_op_reg *op2 = new hsa_op_reg (btype);
3188 hsa_op_with_type *shift1 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3190 tree type = TREE_TYPE (rhs2);
3191 unsigned HOST_WIDE_INT bitsize = TREE_INT_CST_LOW (TYPE_SIZE (type));
3193 hsa_op_with_type *shift2 = NULL;
3194 if (TREE_CODE (rhs2) == INTEGER_CST)
3195 shift2 = new hsa_op_immed (bitsize - tree_to_uhwi (rhs2),
3196 BRIG_TYPE_U32);
3197 else if (TREE_CODE (rhs2) == SSA_NAME)
3199 hsa_op_reg *s = hsa_cfun->reg_for_gimple_ssa (rhs2);
3200 hsa_op_reg *d = new hsa_op_reg (s->m_type);
3201 hsa_op_immed *size_imm = new hsa_op_immed (bitsize, BRIG_TYPE_U32);
3203 insn = new hsa_insn_basic (3, BRIG_OPCODE_SUB, d->m_type,
3204 d, s, size_imm);
3205 hbb->append_insn (insn);
3207 shift2 = d;
3209 else
3210 gcc_unreachable ();
3212 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3213 gen_hsa_binary_operation (code1, op1, src, shift1, hbb);
3214 gen_hsa_binary_operation (code2, op2, src, shift2, hbb);
3215 gen_hsa_binary_operation (BRIG_OPCODE_OR, dest, op1, op2, hbb);
3217 return;
3219 case BIT_IOR_EXPR:
3220 opcode = BRIG_OPCODE_OR;
3221 break;
3222 case BIT_XOR_EXPR:
3223 opcode = BRIG_OPCODE_XOR;
3224 break;
3225 case BIT_AND_EXPR:
3226 opcode = BRIG_OPCODE_AND;
3227 break;
3228 case BIT_NOT_EXPR:
3229 opcode = BRIG_OPCODE_NOT;
3230 break;
3231 case FIX_TRUNC_EXPR:
3233 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3234 hsa_op_with_type *v = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3236 if (hsa_needs_cvt (dest->m_type, v->m_type))
3238 hsa_op_reg *tmp = new hsa_op_reg (v->m_type);
3240 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_TRUNC,
3241 tmp->m_type, tmp, v);
3242 hbb->append_insn (insn);
3244 hsa_insn_basic *cvtinsn = new hsa_insn_cvt (dest, tmp);
3245 hbb->append_insn (cvtinsn);
3247 else
3249 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_TRUNC,
3250 dest->m_type, dest, v);
3251 hbb->append_insn (insn);
3254 return;
3256 opcode = BRIG_OPCODE_TRUNC;
3257 break;
3259 case LT_EXPR:
3260 case LE_EXPR:
3261 case GT_EXPR:
3262 case GE_EXPR:
3263 case EQ_EXPR:
3264 case NE_EXPR:
3265 case UNORDERED_EXPR:
3266 case ORDERED_EXPR:
3267 case UNLT_EXPR:
3268 case UNLE_EXPR:
3269 case UNGT_EXPR:
3270 case UNGE_EXPR:
3271 case UNEQ_EXPR:
3272 case LTGT_EXPR:
3274 hsa_op_reg *dest
3275 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3277 gen_hsa_cmp_insn_from_gimple (code, rhs1, rhs2, dest, hbb);
3278 return;
3280 case COND_EXPR:
3282 hsa_op_reg *dest
3283 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3284 hsa_op_with_type *ctrl = NULL;
3285 tree cond = rhs1;
3287 if (CONSTANT_CLASS_P (cond) || TREE_CODE (cond) == SSA_NAME)
3288 ctrl = hsa_reg_or_immed_for_gimple_op (cond, hbb);
3289 else
3291 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_B1);
3293 gen_hsa_cmp_insn_from_gimple (TREE_CODE (cond),
3294 TREE_OPERAND (cond, 0),
3295 TREE_OPERAND (cond, 1),
3296 r, hbb);
3298 ctrl = r;
3301 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3302 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
3304 BrigType16_t utype = hsa_unsigned_type_for_type (dest->m_type);
3305 if (is_a <hsa_op_immed *> (op2))
3306 op2->m_type = utype;
3307 if (is_a <hsa_op_immed *> (op3))
3308 op3->m_type = utype;
3310 hsa_insn_basic *insn
3311 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV,
3312 hsa_bittype_for_type (dest->m_type),
3313 dest, ctrl, op2, op3);
3315 hbb->append_insn (insn);
3316 return;
3318 case COMPLEX_EXPR:
3320 hsa_op_reg *dest
3321 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3322 hsa_op_with_type *rhs1_reg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3323 hsa_op_with_type *rhs2_reg = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3325 if (hsa_seen_error ())
3326 return;
3328 BrigType16_t src_type = hsa_bittype_for_type (rhs1_reg->m_type);
3329 rhs1_reg = rhs1_reg->get_in_type (src_type, hbb);
3330 rhs2_reg = rhs2_reg->get_in_type (src_type, hbb);
3332 hsa_insn_packed *insn
3333 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type, src_type,
3334 dest, rhs1_reg, rhs2_reg);
3335 hbb->append_insn (insn);
3337 return;
3339 default:
3340 /* Implement others as we come across them. */
3341 HSA_SORRY_ATV (gimple_location (assign),
3342 "support for HSA does not implement operation %s",
3343 get_tree_code_name (code));
3344 return;
3348 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3350 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3351 hsa_op_with_type *op2 = rhs2 != NULL_TREE ?
3352 hsa_reg_or_immed_for_gimple_op (rhs2, hbb) : NULL;
3354 if (hsa_seen_error ())
3355 return;
3357 switch (rhs_class)
3359 case GIMPLE_TERNARY_RHS:
3360 gcc_unreachable ();
3361 return;
3363 /* Fall through */
3364 case GIMPLE_BINARY_RHS:
3365 gen_hsa_binary_operation (opcode, dest, op1, op2, hbb);
3366 break;
3367 /* Fall through */
3368 case GIMPLE_UNARY_RHS:
3369 gen_hsa_unary_operation (opcode, dest, op1, hbb);
3370 break;
3371 default:
3372 gcc_unreachable ();
3376 /* Generate HSA instructions for a given gimple condition statement COND.
3377 Instructions will be appended to HBB, which also needs to be the
3378 corresponding structure to the basic_block of COND. */
3380 static void
3381 gen_hsa_insns_for_cond_stmt (gimple *cond, hsa_bb *hbb)
3383 hsa_op_reg *ctrl = new hsa_op_reg (BRIG_TYPE_B1);
3384 hsa_insn_br *cbr;
3386 gen_hsa_cmp_insn_from_gimple (gimple_cond_code (cond),
3387 gimple_cond_lhs (cond),
3388 gimple_cond_rhs (cond),
3389 ctrl, hbb);
3391 cbr = new hsa_insn_br (ctrl);
3392 hbb->append_insn (cbr);
3395 /* Maximum number of elements in a jump table for an HSA SBR instruction. */
3397 #define HSA_MAXIMUM_SBR_LABELS 16
3399 /* Return lowest value of a switch S that is handled in a non-default
3400 label. */
3402 static tree
3403 get_switch_low (gswitch *s)
3405 unsigned labels = gimple_switch_num_labels (s);
3406 gcc_checking_assert (labels >= 1);
3408 return CASE_LOW (gimple_switch_label (s, 1));
3411 /* Return highest value of a switch S that is handled in a non-default
3412 label. */
3414 static tree
3415 get_switch_high (gswitch *s)
3417 unsigned labels = gimple_switch_num_labels (s);
3419 /* Compare last label to maximum number of labels. */
3420 tree label = gimple_switch_label (s, labels - 1);
3421 tree low = CASE_LOW (label);
3422 tree high = CASE_HIGH (label);
3424 return high != NULL_TREE ? high : low;
3427 static tree
3428 get_switch_size (gswitch *s)
3430 return int_const_binop (MINUS_EXPR, get_switch_high (s), get_switch_low (s));
3433 /* Generate HSA instructions for a given gimple switch.
3434 Instructions will be appended to HBB. */
3436 static void
3437 gen_hsa_insns_for_switch_stmt (gswitch *s, hsa_bb *hbb)
3439 gimple_stmt_iterator it = gsi_for_stmt (s);
3440 gsi_prev (&it);
3442 /* Create preambule that verifies that index - lowest_label >= 0. */
3443 edge e = split_block (hbb->m_bb, gsi_stmt (it));
3444 e->flags &= ~EDGE_FALLTHRU;
3445 e->flags |= EDGE_TRUE_VALUE;
3447 function *func = DECL_STRUCT_FUNCTION (current_function_decl);
3448 tree index_tree = gimple_switch_index (s);
3449 tree lowest = get_switch_low (s);
3450 tree highest = get_switch_high (s);
3452 hsa_op_reg *index = hsa_cfun->reg_for_gimple_ssa (index_tree);
3454 hsa_op_reg *cmp1_reg = new hsa_op_reg (BRIG_TYPE_B1);
3455 hsa_op_immed *cmp1_immed = new hsa_op_immed (lowest);
3456 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_GE, cmp1_reg->m_type,
3457 cmp1_reg, index, cmp1_immed));
3459 hsa_op_reg *cmp2_reg = new hsa_op_reg (BRIG_TYPE_B1);
3460 hsa_op_immed *cmp2_immed = new hsa_op_immed (highest);
3461 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_LE, cmp2_reg->m_type,
3462 cmp2_reg, index, cmp2_immed));
3464 hsa_op_reg *cmp_reg = new hsa_op_reg (BRIG_TYPE_B1);
3465 hbb->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_AND, cmp_reg->m_type,
3466 cmp_reg, cmp1_reg, cmp2_reg));
3468 hbb->append_insn (new hsa_insn_br (cmp_reg));
3470 tree default_label = gimple_switch_default_label (s);
3471 basic_block default_label_bb = label_to_block_fn (func,
3472 CASE_LABEL (default_label));
3474 make_edge (e->src, default_label_bb, EDGE_FALSE_VALUE);
3476 hsa_cfun->m_modified_cfg = true;
3478 /* Basic block with the SBR instruction. */
3479 hbb = hsa_init_new_bb (e->dest);
3481 hsa_op_reg *sub_index = new hsa_op_reg (index->m_type);
3482 hbb->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB, sub_index->m_type,
3483 sub_index, index,
3484 new hsa_op_immed (lowest)));
3486 hsa_op_base *tmp = sub_index->get_in_type (BRIG_TYPE_U64, hbb);
3487 sub_index = as_a <hsa_op_reg *> (tmp);
3488 unsigned labels = gimple_switch_num_labels (s);
3489 unsigned HOST_WIDE_INT size = tree_to_uhwi (get_switch_size (s));
3491 hsa_insn_sbr *sbr = new hsa_insn_sbr (sub_index, size + 1);
3493 /* Prepare array with default label destination. */
3494 for (unsigned HOST_WIDE_INT i = 0; i <= size; i++)
3495 sbr->m_jump_table.safe_push (default_label_bb);
3497 /* Iterate all labels and fill up the jump table. */
3498 for (unsigned i = 1; i < labels; i++)
3500 tree label = gimple_switch_label (s, i);
3501 basic_block bb = label_to_block_fn (func, CASE_LABEL (label));
3503 unsigned HOST_WIDE_INT sub_low
3504 = tree_to_uhwi (int_const_binop (MINUS_EXPR, CASE_LOW (label), lowest));
3506 unsigned HOST_WIDE_INT sub_high = sub_low;
3507 tree high = CASE_HIGH (label);
3508 if (high != NULL)
3509 sub_high = tree_to_uhwi (int_const_binop (MINUS_EXPR, high, lowest));
3511 for (unsigned HOST_WIDE_INT j = sub_low; j <= sub_high; j++)
3512 sbr->m_jump_table[j] = bb;
3515 hbb->append_insn (sbr);
3518 /* Verify that the function DECL can be handled by HSA. */
3520 static void
3521 verify_function_arguments (tree decl)
3523 if (DECL_STATIC_CHAIN (decl))
3525 HSA_SORRY_ATV (EXPR_LOCATION (decl),
3526 "HSA does not support nested functions: %D", decl);
3527 return;
3529 else if (!TYPE_ARG_TYPES (TREE_TYPE (decl)))
3531 HSA_SORRY_ATV (EXPR_LOCATION (decl),
3532 "HSA does not support functions with variadic arguments "
3533 "(or unknown return type): %D", decl);
3534 return;
3538 /* Return BRIG type for FORMAL_ARG_TYPE. If the formal argument type is NULL,
3539 return ACTUAL_ARG_TYPE. */
3541 static BrigType16_t
3542 get_format_argument_type (tree formal_arg_type, BrigType16_t actual_arg_type)
3544 if (formal_arg_type == NULL)
3545 return actual_arg_type;
3547 BrigType16_t decl_type
3548 = hsa_type_for_scalar_tree_type (formal_arg_type, false);
3549 return mem_type_for_type (decl_type);
3552 /* Generate HSA instructions for a direct call instruction.
3553 Instructions will be appended to HBB, which also needs to be the
3554 corresponding structure to the basic_block of STMT.
3555 If ASSIGN_LHS is false, do not copy HSA function result argument into the
3556 corresponding HSA representation of the gimple statement LHS. */
3558 static void
3559 gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb,
3560 bool assign_lhs = true)
3562 tree decl = gimple_call_fndecl (stmt);
3563 verify_function_arguments (decl);
3564 if (hsa_seen_error ())
3565 return;
3567 hsa_insn_call *call_insn = new hsa_insn_call (decl);
3568 hsa_cfun->m_called_functions.safe_push (call_insn->m_called_function);
3570 /* Argument block start. */
3571 hsa_insn_arg_block *arg_start
3572 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn);
3573 hbb->append_insn (arg_start);
3575 tree parm_type_chain = TYPE_ARG_TYPES (gimple_call_fntype (stmt));
3577 /* Preparation of arguments that will be passed to function. */
3578 const unsigned args = gimple_call_num_args (stmt);
3579 for (unsigned i = 0; i < args; ++i)
3581 tree parm = gimple_call_arg (stmt, (int)i);
3582 tree parm_decl_type = parm_type_chain != NULL_TREE
3583 ? TREE_VALUE (parm_type_chain) : NULL_TREE;
3584 hsa_op_address *addr;
3586 if (AGGREGATE_TYPE_P (TREE_TYPE (parm)))
3588 addr = gen_hsa_addr_for_arg (TREE_TYPE (parm), i);
3589 BrigAlignment8_t align;
3590 hsa_op_address *src = gen_hsa_addr_with_align (parm, hbb, &align);
3591 gen_hsa_memory_copy (hbb, addr, src,
3592 addr->m_symbol->total_byte_size (), align);
3594 else
3596 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (parm, hbb);
3598 if (parm_decl_type != NULL && AGGREGATE_TYPE_P (parm_decl_type))
3600 HSA_SORRY_AT (gimple_location (stmt),
3601 "support for HSA does not implement an aggregate "
3602 "formal argument in a function call, while actual "
3603 "argument is not an aggregate");
3604 return;
3607 BrigType16_t formal_arg_type
3608 = get_format_argument_type (parm_decl_type, src->m_type);
3609 if (hsa_seen_error ())
3610 return;
3612 if (src->m_type != formal_arg_type)
3613 src = src->get_in_type (formal_arg_type, hbb);
3615 addr
3616 = gen_hsa_addr_for_arg (parm_decl_type != NULL_TREE ?
3617 parm_decl_type: TREE_TYPE (parm), i);
3618 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, formal_arg_type,
3619 src, addr);
3621 hbb->append_insn (mem);
3624 call_insn->m_input_args.safe_push (addr->m_symbol);
3625 if (parm_type_chain)
3626 parm_type_chain = TREE_CHAIN (parm_type_chain);
3629 call_insn->m_args_code_list = new hsa_op_code_list (args);
3630 hbb->append_insn (call_insn);
3632 tree result_type = TREE_TYPE (TREE_TYPE (decl));
3634 tree result = gimple_call_lhs (stmt);
3635 hsa_insn_mem *result_insn = NULL;
3636 if (!VOID_TYPE_P (result_type))
3638 hsa_op_address *addr = gen_hsa_addr_for_arg (result_type, -1);
3640 /* Even if result of a function call is unused, we have to emit
3641 declaration for the result. */
3642 if (result && assign_lhs)
3644 tree lhs_type = TREE_TYPE (result);
3646 if (hsa_seen_error ())
3647 return;
3649 if (AGGREGATE_TYPE_P (lhs_type))
3651 BrigAlignment8_t align;
3652 hsa_op_address *result_addr
3653 = gen_hsa_addr_with_align (result, hbb, &align);
3654 gen_hsa_memory_copy (hbb, result_addr, addr,
3655 addr->m_symbol->total_byte_size (), align);
3657 else
3659 BrigType16_t mtype
3660 = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type,
3661 false));
3663 hsa_op_reg *dst = hsa_cfun->reg_for_gimple_ssa (result);
3664 result_insn = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, dst, addr);
3665 hbb->append_insn (result_insn);
3669 call_insn->m_output_arg = addr->m_symbol;
3670 call_insn->m_result_code_list = new hsa_op_code_list (1);
3672 else
3674 if (result)
3676 HSA_SORRY_AT (gimple_location (stmt),
3677 "support for HSA does not implement an assignment of "
3678 "return value from a void function");
3679 return;
3682 call_insn->m_result_code_list = new hsa_op_code_list (0);
3685 /* Argument block end. */
3686 hsa_insn_arg_block *arg_end
3687 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn);
3688 hbb->append_insn (arg_end);
3691 /* Generate HSA instructions for a direct call of an internal fn.
3692 Instructions will be appended to HBB, which also needs to be the
3693 corresponding structure to the basic_block of STMT. */
3695 static void
3696 gen_hsa_insns_for_call_of_internal_fn (gimple *stmt, hsa_bb *hbb)
3698 tree lhs = gimple_call_lhs (stmt);
3699 if (!lhs)
3700 return;
3702 tree lhs_type = TREE_TYPE (lhs);
3703 tree rhs1 = gimple_call_arg (stmt, 0);
3704 tree rhs1_type = TREE_TYPE (rhs1);
3705 enum internal_fn fn = gimple_call_internal_fn (stmt);
3706 hsa_internal_fn *ifn
3707 = new hsa_internal_fn (fn, tree_to_uhwi (TYPE_SIZE (rhs1_type)));
3708 hsa_insn_call *call_insn = new hsa_insn_call (ifn);
3710 gcc_checking_assert (FLOAT_TYPE_P (rhs1_type));
3712 if (!hsa_emitted_internal_decls->find (call_insn->m_called_internal_fn))
3713 hsa_cfun->m_called_internal_fns.safe_push (call_insn->m_called_internal_fn);
3715 hsa_insn_arg_block *arg_start
3716 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn);
3717 hbb->append_insn (arg_start);
3719 unsigned num_args = gimple_call_num_args (stmt);
3721 /* Function arguments. */
3722 for (unsigned i = 0; i < num_args; i++)
3724 tree parm = gimple_call_arg (stmt, (int)i);
3725 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (parm, hbb);
3727 hsa_op_address *addr = gen_hsa_addr_for_arg (TREE_TYPE (parm), i);
3728 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, src->m_type,
3729 src, addr);
3731 call_insn->m_input_args.safe_push (addr->m_symbol);
3732 hbb->append_insn (mem);
3735 call_insn->m_args_code_list = new hsa_op_code_list (num_args);
3736 hbb->append_insn (call_insn);
3738 /* Assign returned value. */
3739 hsa_op_address *addr = gen_hsa_addr_for_arg (lhs_type, -1);
3741 call_insn->m_output_arg = addr->m_symbol;
3742 call_insn->m_result_code_list = new hsa_op_code_list (1);
3744 /* Argument block end. */
3745 hsa_insn_arg_block *arg_end
3746 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn);
3747 hbb->append_insn (arg_end);
3750 /* Generate HSA instructions for a return value instruction.
3751 Instructions will be appended to HBB, which also needs to be the
3752 corresponding structure to the basic_block of STMT. */
3754 static void
3755 gen_hsa_insns_for_return (greturn *stmt, hsa_bb *hbb)
3757 tree retval = gimple_return_retval (stmt);
3758 if (retval)
3760 hsa_op_address *addr = new hsa_op_address (hsa_cfun->m_output_arg);
3762 if (AGGREGATE_TYPE_P (TREE_TYPE (retval)))
3764 BrigAlignment8_t align;
3765 hsa_op_address *retval_addr = gen_hsa_addr_with_align (retval, hbb,
3766 &align);
3767 gen_hsa_memory_copy (hbb, addr, retval_addr,
3768 hsa_cfun->m_output_arg->total_byte_size (),
3769 align);
3771 else
3773 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (retval),
3774 false);
3775 BrigType16_t mtype = mem_type_for_type (t);
3777 /* Store of return value. */
3778 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (retval, hbb);
3779 src = src->get_in_type (mtype, hbb);
3780 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, mtype, src,
3781 addr);
3782 hbb->append_insn (mem);
3786 /* HSAIL return instruction emission. */
3787 hsa_insn_basic *ret = new hsa_insn_basic (0, BRIG_OPCODE_RET);
3788 hbb->append_insn (ret);
3791 /* Set OP_INDEX-th operand of the instruction to DEST, as the DEST
3792 can have a different type, conversion instructions are possibly
3793 appended to HBB. */
3795 void
3796 hsa_insn_basic::set_output_in_type (hsa_op_reg *dest, unsigned op_index,
3797 hsa_bb *hbb)
3799 hsa_insn_basic *insn;
3800 gcc_checking_assert (op_output_p (op_index));
3802 if (dest->m_type == m_type)
3804 set_op (op_index, dest);
3805 return;
3808 hsa_op_reg *tmp = new hsa_op_reg (m_type);
3809 set_op (op_index, tmp);
3811 if (hsa_needs_cvt (dest->m_type, m_type))
3812 insn = new hsa_insn_cvt (dest, tmp);
3813 else
3814 insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type,
3815 dest, tmp->get_in_type (dest->m_type, hbb));
3817 hbb->append_insn (insn);
3820 /* Generate instruction OPCODE to query a property of HSA grid along the
3821 given DIMENSION. Store result into DEST and append the instruction to
3822 HBB. */
3824 static void
3825 query_hsa_grid (hsa_op_reg *dest, BrigType16_t opcode, int dimension,
3826 hsa_bb *hbb)
3828 /* We're using just one-dimensional kernels, so hard-coded
3829 dimension X. */
3830 hsa_op_immed *imm
3831 = new hsa_op_immed (dimension, (BrigKind16_t) BRIG_TYPE_U32);
3832 hsa_insn_basic *insn = new hsa_insn_basic (2, opcode, BRIG_TYPE_U32, NULL,
3833 imm);
3834 hbb->append_insn (insn);
3835 insn->set_output_in_type (dest, 0, hbb);
3838 /* Generate a special HSA-related instruction for gimple STMT.
3839 Instructions are appended to basic block HBB. */
3841 static void
3842 query_hsa_grid (gimple *stmt, BrigOpcode16_t opcode, int dimension,
3843 hsa_bb *hbb)
3845 tree lhs = gimple_call_lhs (dyn_cast <gcall *> (stmt));
3846 if (lhs == NULL_TREE)
3847 return;
3849 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3851 query_hsa_grid (dest, opcode, dimension, hbb);
3854 /* Emit instructions that set hsa_num_threads according to provided VALUE.
3855 Instructions are appended to basic block HBB. */
3857 static void
3858 gen_set_num_threads (tree value, hsa_bb *hbb)
3860 hbb->append_insn (new hsa_insn_comment ("omp_set_num_threads"));
3861 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (value, hbb);
3863 src = src->get_in_type (hsa_num_threads->m_type, hbb);
3864 hsa_op_address *addr = new hsa_op_address (hsa_num_threads);
3866 hsa_insn_basic *basic
3867 = new hsa_insn_mem (BRIG_OPCODE_ST, hsa_num_threads->m_type, src, addr);
3868 hbb->append_insn (basic);
3871 /* Return byte offset of a FIELD_NAME in GOMP_hsa_kernel_dispatch which
3872 is defined in plugin-hsa.c. */
3874 static HOST_WIDE_INT
3875 get_hsa_kernel_dispatch_offset (const char *field_name)
3877 tree *hsa_kernel_dispatch_type = hsa_get_kernel_dispatch_type ();
3878 if (*hsa_kernel_dispatch_type == NULL)
3880 /* Collection of information needed for a dispatch of a kernel from a
3881 kernel. Keep in sync with libgomp's plugin-hsa.c. */
3883 *hsa_kernel_dispatch_type = make_node (RECORD_TYPE);
3884 tree id_f1 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3885 get_identifier ("queue"), ptr_type_node);
3886 DECL_CHAIN (id_f1) = NULL_TREE;
3887 tree id_f2 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3888 get_identifier ("omp_data_memory"),
3889 ptr_type_node);
3890 DECL_CHAIN (id_f2) = id_f1;
3891 tree id_f3 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3892 get_identifier ("kernarg_address"),
3893 ptr_type_node);
3894 DECL_CHAIN (id_f3) = id_f2;
3895 tree id_f4 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3896 get_identifier ("object"),
3897 uint64_type_node);
3898 DECL_CHAIN (id_f4) = id_f3;
3899 tree id_f5 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3900 get_identifier ("signal"),
3901 uint64_type_node);
3902 DECL_CHAIN (id_f5) = id_f4;
3903 tree id_f6 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3904 get_identifier ("private_segment_size"),
3905 uint32_type_node);
3906 DECL_CHAIN (id_f6) = id_f5;
3907 tree id_f7 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3908 get_identifier ("group_segment_size"),
3909 uint32_type_node);
3910 DECL_CHAIN (id_f7) = id_f6;
3911 tree id_f8 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3912 get_identifier ("kernel_dispatch_count"),
3913 uint64_type_node);
3914 DECL_CHAIN (id_f8) = id_f7;
3915 tree id_f9 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3916 get_identifier ("debug"),
3917 uint64_type_node);
3918 DECL_CHAIN (id_f9) = id_f8;
3919 tree id_f10 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3920 get_identifier ("omp_level"),
3921 uint64_type_node);
3922 DECL_CHAIN (id_f10) = id_f9;
3923 tree id_f11 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3924 get_identifier ("children_dispatches"),
3925 ptr_type_node);
3926 DECL_CHAIN (id_f11) = id_f10;
3927 tree id_f12 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3928 get_identifier ("omp_num_threads"),
3929 uint32_type_node);
3930 DECL_CHAIN (id_f12) = id_f11;
3933 finish_builtin_struct (*hsa_kernel_dispatch_type, "__hsa_kernel_dispatch",
3934 id_f12, NULL_TREE);
3935 TYPE_ARTIFICIAL (*hsa_kernel_dispatch_type) = 1;
3938 for (tree chain = TYPE_FIELDS (*hsa_kernel_dispatch_type);
3939 chain != NULL_TREE; chain = TREE_CHAIN (chain))
3940 if (strcmp (field_name, IDENTIFIER_POINTER (DECL_NAME (chain))) == 0)
3941 return int_byte_position (chain);
3943 gcc_unreachable ();
3946 /* Return an HSA register that will contain number of threads for
3947 a future dispatched kernel. Instructions are added to HBB. */
3949 static hsa_op_reg *
3950 gen_num_threads_for_dispatch (hsa_bb *hbb)
3952 /* Step 1) Assign to number of threads:
3953 MIN (HSA_DEFAULT_NUM_THREADS, hsa_num_threads). */
3954 hsa_op_reg *threads = new hsa_op_reg (hsa_num_threads->m_type);
3955 hsa_op_address *addr = new hsa_op_address (hsa_num_threads);
3957 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD, threads->m_type,
3958 threads, addr));
3960 hsa_op_immed *limit = new hsa_op_immed (HSA_DEFAULT_NUM_THREADS,
3961 BRIG_TYPE_U32);
3962 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_B1);
3963 hsa_insn_cmp * cmp
3964 = new hsa_insn_cmp (BRIG_COMPARE_LT, r->m_type, r, threads, limit);
3965 hbb->append_insn (cmp);
3967 BrigType16_t btype = hsa_bittype_for_type (threads->m_type);
3968 hsa_op_reg *tmp = new hsa_op_reg (threads->m_type);
3970 hbb->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp, r,
3971 threads, limit));
3973 /* Step 2) If the number is equal to zero,
3974 return shadow->omp_num_threads. */
3975 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg ();
3977 hsa_op_reg *shadow_thread_count = new hsa_op_reg (BRIG_TYPE_U32);
3978 addr
3979 = new hsa_op_address (shadow_reg_ptr,
3980 get_hsa_kernel_dispatch_offset ("omp_num_threads"));
3981 hsa_insn_basic *basic
3982 = new hsa_insn_mem (BRIG_OPCODE_LD, shadow_thread_count->m_type,
3983 shadow_thread_count, addr);
3984 hbb->append_insn (basic);
3986 hsa_op_reg *tmp2 = new hsa_op_reg (threads->m_type);
3987 r = new hsa_op_reg (BRIG_TYPE_B1);
3988 hsa_op_immed *imm = new hsa_op_immed (0, shadow_thread_count->m_type);
3989 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_EQ, r->m_type, r, tmp, imm));
3990 hbb->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp2, r,
3991 shadow_thread_count, tmp));
3993 hsa_op_base *dest = tmp2->get_in_type (BRIG_TYPE_U16, hbb);
3995 return as_a <hsa_op_reg *> (dest);
3999 /* Emit instructions that assign number of teams to lhs of gimple STMT.
4000 Instructions are appended to basic block HBB. */
4002 static void
4003 gen_get_num_teams (gimple *stmt, hsa_bb *hbb)
4005 if (gimple_call_lhs (stmt) == NULL_TREE)
4006 return;
4008 hbb->append_insn (new hsa_insn_comment ("omp_get_num_teams"));
4010 tree lhs = gimple_call_lhs (stmt);
4011 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4012 hsa_op_immed *one = new hsa_op_immed (1, dest->m_type);
4014 hsa_insn_basic *basic
4015 = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, dest, one);
4017 hbb->append_insn (basic);
4020 /* Emit instructions that assign a team number to lhs of gimple STMT.
4021 Instructions are appended to basic block HBB. */
4023 static void
4024 gen_get_team_num (gimple *stmt, hsa_bb *hbb)
4026 if (gimple_call_lhs (stmt) == NULL_TREE)
4027 return;
4029 hbb->append_insn (new hsa_insn_comment ("omp_get_team_num"));
4031 tree lhs = gimple_call_lhs (stmt);
4032 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4033 hsa_op_immed *zero = new hsa_op_immed (0, dest->m_type);
4035 hsa_insn_basic *basic
4036 = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, dest, zero);
4038 hbb->append_insn (basic);
4041 /* Emit instructions that get levels-var ICV to lhs of gimple STMT.
4042 Instructions are appended to basic block HBB. */
4044 static void
4045 gen_get_level (gimple *stmt, hsa_bb *hbb)
4047 if (gimple_call_lhs (stmt) == NULL_TREE)
4048 return;
4050 hbb->append_insn (new hsa_insn_comment ("omp_get_level"));
4052 tree lhs = gimple_call_lhs (stmt);
4053 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4055 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg ();
4056 if (shadow_reg_ptr == NULL)
4058 HSA_SORRY_AT (gimple_location (stmt),
4059 "support for HSA does not implement omp_get_level called "
4060 "from a function not being inlined within a kernel");
4061 return;
4064 hsa_op_address *addr
4065 = new hsa_op_address (shadow_reg_ptr,
4066 get_hsa_kernel_dispatch_offset ("omp_level"));
4068 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64,
4069 (hsa_op_base *) NULL, addr);
4070 hbb->append_insn (mem);
4071 mem->set_output_in_type (dest, 0, hbb);
4074 /* Emit instruction that implement omp_get_max_threads of gimple STMT. */
4076 static void
4077 gen_get_max_threads (gimple *stmt, hsa_bb *hbb)
4079 tree lhs = gimple_call_lhs (stmt);
4080 if (!lhs)
4081 return;
4083 hbb->append_insn (new hsa_insn_comment ("omp_get_max_threads"));
4085 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4086 hsa_op_with_type *num_theads_reg = gen_num_threads_for_dispatch (hbb)
4087 ->get_in_type (dest->m_type, hbb);
4088 hsa_build_append_simple_mov (dest, num_theads_reg, hbb);
4091 /* Emit instructions that implement alloca builtin gimple STMT.
4092 Instructions are appended to basic block HBB. */
4094 static void
4095 gen_hsa_alloca (gcall *call, hsa_bb *hbb)
4097 tree lhs = gimple_call_lhs (call);
4098 if (lhs == NULL_TREE)
4099 return;
4101 built_in_function fn = DECL_FUNCTION_CODE (gimple_call_fndecl (call));
4103 gcc_checking_assert (fn == BUILT_IN_ALLOCA
4104 || fn == BUILT_IN_ALLOCA_WITH_ALIGN);
4106 unsigned bit_alignment = 0;
4108 if (fn == BUILT_IN_ALLOCA_WITH_ALIGN)
4110 tree alignment_tree = gimple_call_arg (call, 1);
4111 if (TREE_CODE (alignment_tree) != INTEGER_CST)
4113 HSA_SORRY_ATV (gimple_location (call),
4114 "support for HSA does not implement "
4115 "__builtin_alloca_with_align with a non-constant "
4116 "alignment: %E", alignment_tree);
4119 bit_alignment = tree_to_uhwi (alignment_tree);
4122 tree rhs1 = gimple_call_arg (call, 0);
4123 hsa_op_with_type *size = hsa_reg_or_immed_for_gimple_op (rhs1, hbb)
4124 ->get_in_type (BRIG_TYPE_U32, hbb);
4125 hsa_op_with_type *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4127 hsa_op_reg *tmp
4128 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE));
4129 hsa_insn_alloca *a = new hsa_insn_alloca (tmp, size, bit_alignment);
4130 hbb->append_insn (a);
4132 hsa_insn_seg *seg
4133 = new hsa_insn_seg (BRIG_OPCODE_STOF,
4134 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT),
4135 tmp->m_type, BRIG_SEGMENT_PRIVATE, dest, tmp);
4136 hbb->append_insn (seg);
4139 /* Emit instructions that implement clrsb builtin STMT:
4140 Returns the number of leading redundant sign bits in x, i.e. the number
4141 of bits following the most significant bit that are identical to it.
4142 There are no special cases for 0 or other values.
4143 Instructions are appended to basic block HBB. */
4145 static void
4146 gen_hsa_clrsb (gcall *call, hsa_bb *hbb)
4148 tree lhs = gimple_call_lhs (call);
4149 if (lhs == NULL_TREE)
4150 return;
4152 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4153 tree rhs1 = gimple_call_arg (call, 0);
4154 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4155 BrigType16_t bittype = hsa_bittype_for_type (arg->m_type);
4156 unsigned bitsize = tree_to_uhwi (TYPE_SIZE (TREE_TYPE (rhs1)));
4158 /* FIRSTBIT instruction is defined just for 32 and 64-bits wide integers. */
4159 gcc_checking_assert (bitsize == 32 || bitsize == 64);
4161 /* Set true to MOST_SIG if the most significant bit is set to one. */
4162 hsa_op_immed *c = new hsa_op_immed (1ul << (bitsize - 1),
4163 hsa_uint_for_bitsize (bitsize));
4165 hsa_op_reg *and_reg = new hsa_op_reg (bittype);
4166 gen_hsa_binary_operation (BRIG_OPCODE_AND, and_reg, arg, c, hbb);
4168 hsa_op_reg *most_sign = new hsa_op_reg (BRIG_TYPE_B1);
4169 hsa_insn_cmp *cmp
4170 = new hsa_insn_cmp (BRIG_COMPARE_EQ, most_sign->m_type, most_sign,
4171 and_reg, c);
4172 hbb->append_insn (cmp);
4174 /* If the most significant bit is one, negate the input. Otherwise
4175 shift the input value to left by one bit. */
4176 hsa_op_reg *arg_neg = new hsa_op_reg (arg->m_type);
4177 gen_hsa_unary_operation (BRIG_OPCODE_NEG, arg_neg, arg, hbb);
4179 hsa_op_reg *shifted_arg = new hsa_op_reg (arg->m_type);
4180 gen_hsa_binary_operation (BRIG_OPCODE_SHL, shifted_arg, arg,
4181 new hsa_op_immed (1, BRIG_TYPE_U64), hbb);
4183 /* Assign the value that can be used for FIRSTBIT instruction according
4184 to the most significant bit. */
4185 hsa_op_reg *tmp = new hsa_op_reg (bittype);
4186 hsa_insn_basic *cmov
4187 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, bittype, tmp, most_sign,
4188 arg_neg, shifted_arg);
4189 hbb->append_insn (cmov);
4191 hsa_op_reg *leading_bits = new hsa_op_reg (BRIG_TYPE_S32);
4192 gen_hsa_unary_operation (BRIG_OPCODE_FIRSTBIT, leading_bits,
4193 tmp->get_in_type (hsa_uint_for_bitsize (bitsize),
4194 hbb), hbb);
4196 /* Set flag if the input value is equal to zero. */
4197 hsa_op_reg *is_zero = new hsa_op_reg (BRIG_TYPE_B1);
4198 cmp = new hsa_insn_cmp (BRIG_COMPARE_EQ, is_zero->m_type, is_zero, arg,
4199 new hsa_op_immed (0, arg->m_type));
4200 hbb->append_insn (cmp);
4202 /* Return the number of leading bits,
4203 or (bitsize - 1) if the input value is zero. */
4204 cmov = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, BRIG_TYPE_B32, NULL, is_zero,
4205 new hsa_op_immed (bitsize - 1, BRIG_TYPE_U32),
4206 leading_bits->get_in_type (BRIG_TYPE_B32, hbb));
4207 hbb->append_insn (cmov);
4208 cmov->set_output_in_type (dest, 0, hbb);
4211 /* Emit instructions that implement ffs builtin STMT:
4212 Returns one plus the index of the least significant 1-bit of x,
4213 or if x is zero, returns zero.
4214 Instructions are appended to basic block HBB. */
4216 static void
4217 gen_hsa_ffs (gcall *call, hsa_bb *hbb)
4219 tree lhs = gimple_call_lhs (call);
4220 if (lhs == NULL_TREE)
4221 return;
4223 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4225 tree rhs1 = gimple_call_arg (call, 0);
4226 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4228 hsa_op_reg *tmp = new hsa_op_reg (BRIG_TYPE_U32);
4229 hsa_insn_srctype *insn = new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT,
4230 tmp->m_type, arg->m_type,
4231 tmp, arg);
4232 hbb->append_insn (insn);
4234 hsa_insn_basic *addition
4235 = new hsa_insn_basic (3, BRIG_OPCODE_ADD, tmp->m_type, NULL, tmp,
4236 new hsa_op_immed (1, tmp->m_type));
4237 hbb->append_insn (addition);
4238 addition->set_output_in_type (dest, 0, hbb);
4241 static void
4242 gen_hsa_popcount_to_dest (hsa_op_reg *dest, hsa_op_with_type *arg, hsa_bb *hbb)
4244 gcc_checking_assert (hsa_type_integer_p (arg->m_type));
4246 if (hsa_type_bit_size (arg->m_type) < 32)
4247 arg = arg->get_in_type (BRIG_TYPE_B32, hbb);
4249 if (!hsa_btype_p (arg->m_type))
4250 arg = arg->get_in_type (hsa_bittype_for_type (arg->m_type), hbb);
4252 hsa_insn_srctype *popcount
4253 = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT, BRIG_TYPE_U32,
4254 arg->m_type, NULL, arg);
4255 hbb->append_insn (popcount);
4256 popcount->set_output_in_type (dest, 0, hbb);
4259 /* Emit instructions that implement parity builtin STMT:
4260 Returns the parity of x, i.e. the number of 1-bits in x modulo 2.
4261 Instructions are appended to basic block HBB. */
4263 static void
4264 gen_hsa_parity (gcall *call, hsa_bb *hbb)
4266 tree lhs = gimple_call_lhs (call);
4267 if (lhs == NULL_TREE)
4268 return;
4270 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4271 tree rhs1 = gimple_call_arg (call, 0);
4272 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4274 hsa_op_reg *popcount = new hsa_op_reg (BRIG_TYPE_U32);
4275 gen_hsa_popcount_to_dest (popcount, arg, hbb);
4277 hsa_insn_basic *insn
4278 = new hsa_insn_basic (3, BRIG_OPCODE_REM, popcount->m_type, NULL, popcount,
4279 new hsa_op_immed (2, popcount->m_type));
4280 hbb->append_insn (insn);
4281 insn->set_output_in_type (dest, 0, hbb);
4284 /* Emit instructions that implement popcount builtin STMT.
4285 Instructions are appended to basic block HBB. */
4287 static void
4288 gen_hsa_popcount (gcall *call, hsa_bb *hbb)
4290 tree lhs = gimple_call_lhs (call);
4291 if (lhs == NULL_TREE)
4292 return;
4294 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4295 tree rhs1 = gimple_call_arg (call, 0);
4296 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4298 gen_hsa_popcount_to_dest (dest, arg, hbb);
4301 /* Set VALUE to a shadow kernel debug argument and append a new instruction
4302 to HBB basic block. */
4304 static void
4305 set_debug_value (hsa_bb *hbb, hsa_op_with_type *value)
4307 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg ();
4308 if (shadow_reg_ptr == NULL)
4309 return;
4311 hsa_op_address *addr
4312 = new hsa_op_address (shadow_reg_ptr,
4313 get_hsa_kernel_dispatch_offset ("debug"));
4314 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, BRIG_TYPE_U64, value,
4315 addr);
4316 hbb->append_insn (mem);
4319 void
4320 omp_simple_builtin::generate (gimple *stmt, hsa_bb *hbb)
4322 if (m_sorry)
4324 if (m_warning_message)
4325 HSA_SORRY_AT (gimple_location (stmt), m_warning_message)
4326 else
4327 HSA_SORRY_ATV (gimple_location (stmt),
4328 "Support for HSA does not implement calls to %s\n",
4329 m_name)
4331 else if (m_warning_message != NULL)
4332 warning_at (gimple_location (stmt), OPT_Whsa, m_warning_message);
4334 if (m_return_value != NULL)
4336 tree lhs = gimple_call_lhs (stmt);
4337 if (!lhs)
4338 return;
4340 hbb->append_insn (new hsa_insn_comment (m_name));
4342 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4343 hsa_op_with_type *op = m_return_value->get_in_type (dest->m_type, hbb);
4344 hsa_build_append_simple_mov (dest, op, hbb);
4348 /* If STMT is a call of a known library function, generate code to perform
4349 it and return true. */
4351 static bool
4352 gen_hsa_insns_for_known_library_call (gimple *stmt, hsa_bb *hbb)
4354 bool handled = false;
4355 const char *name = hsa_get_declaration_name (gimple_call_fndecl (stmt));
4357 char *copy = NULL;
4358 size_t len = strlen (name);
4359 if (len > 0 && name[len - 1] == '_')
4361 copy = XNEWVEC (char, len + 1);
4362 strcpy (copy, name);
4363 copy[len - 1] = '\0';
4364 name = copy;
4367 /* Handle omp_* routines. */
4368 if (strstr (name, "omp_") == name)
4370 hsa_init_simple_builtins ();
4371 omp_simple_builtin *builtin = omp_simple_builtins->get (name);
4372 if (builtin)
4374 builtin->generate (stmt, hbb);
4375 return true;
4378 handled = true;
4379 if (strcmp (name, "omp_set_num_threads") == 0)
4380 gen_set_num_threads (gimple_call_arg (stmt, 0), hbb);
4381 else if (strcmp (name, "omp_get_thread_num") == 0)
4383 hbb->append_insn (new hsa_insn_comment (name));
4384 query_hsa_grid (stmt, BRIG_OPCODE_WORKITEMABSID, 0, hbb);
4386 else if (strcmp (name, "omp_get_num_threads") == 0)
4388 hbb->append_insn (new hsa_insn_comment (name));
4389 query_hsa_grid (stmt, BRIG_OPCODE_GRIDSIZE, 0, hbb);
4391 else if (strcmp (name, "omp_get_num_teams") == 0)
4392 gen_get_num_teams (stmt, hbb);
4393 else if (strcmp (name, "omp_get_team_num") == 0)
4394 gen_get_team_num (stmt, hbb);
4395 else if (strcmp (name, "omp_get_level") == 0)
4396 gen_get_level (stmt, hbb);
4397 else if (strcmp (name, "omp_get_active_level") == 0)
4398 gen_get_level (stmt, hbb);
4399 else if (strcmp (name, "omp_in_parallel") == 0)
4400 gen_get_level (stmt, hbb);
4401 else if (strcmp (name, "omp_get_max_threads") == 0)
4402 gen_get_max_threads (stmt, hbb);
4403 else
4404 handled = false;
4406 if (handled)
4408 if (copy)
4409 free (copy);
4410 return true;
4414 if (strcmp (name, "__hsa_set_debug_value") == 0)
4416 handled = true;
4417 if (hsa_cfun->has_shadow_reg_p ())
4419 tree rhs1 = gimple_call_arg (stmt, 0);
4420 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4422 src = src->get_in_type (BRIG_TYPE_U64, hbb);
4423 set_debug_value (hbb, src);
4427 if (copy)
4428 free (copy);
4429 return handled;
4432 /* Helper functions to create a single unary HSA operations out of calls to
4433 builtins. OPCODE is the HSA operation to be generated. STMT is a gimple
4434 call to a builtin. HBB is the HSA BB to which the instruction should be
4435 added. Note that nothing will be created if STMT does not have a LHS. */
4437 static void
4438 gen_hsa_unaryop_for_builtin (BrigOpcode opcode, gimple *stmt, hsa_bb *hbb)
4440 tree lhs = gimple_call_lhs (stmt);
4441 if (!lhs)
4442 return;
4443 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4444 hsa_op_with_type *op
4445 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb);
4446 gen_hsa_unary_operation (opcode, dest, op, hbb);
4449 /* Helper functions to create a call to standard library if LHS of the
4450 STMT is used. HBB is the HSA BB to which the instruction should be
4451 added. */
4453 static void
4454 gen_hsa_unaryop_builtin_call (gimple *stmt, hsa_bb *hbb)
4456 tree lhs = gimple_call_lhs (stmt);
4457 if (!lhs)
4458 return;
4460 if (gimple_call_internal_p (stmt))
4461 gen_hsa_insns_for_call_of_internal_fn (stmt, hbb);
4462 else
4463 gen_hsa_insns_for_direct_call (stmt, hbb);
4466 /* Helper functions to create a single unary HSA operations out of calls to
4467 builtins (if unsafe math optimizations are enable). Otherwise, create
4468 a call to standard library function.
4469 OPCODE is the HSA operation to be generated. STMT is a gimple
4470 call to a builtin. HBB is the HSA BB to which the instruction should be
4471 added. Note that nothing will be created if STMT does not have a LHS. */
4473 static void
4474 gen_hsa_unaryop_or_call_for_builtin (BrigOpcode opcode, gimple *stmt,
4475 hsa_bb *hbb)
4477 if (flag_unsafe_math_optimizations)
4478 gen_hsa_unaryop_for_builtin (opcode, stmt, hbb);
4479 else
4480 gen_hsa_unaryop_builtin_call (stmt, hbb);
4483 /* Generate HSA address corresponding to a value VAL (as opposed to a memory
4484 reference tree), for example an SSA_NAME or an ADDR_EXPR. HBB is the HSA BB
4485 to which the instruction should be added. */
4487 static hsa_op_address *
4488 get_address_from_value (tree val, hsa_bb *hbb)
4490 switch (TREE_CODE (val))
4492 case SSA_NAME:
4494 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
4495 hsa_op_base *reg
4496 = hsa_cfun->reg_for_gimple_ssa (val)->get_in_type (addrtype, hbb);
4497 return new hsa_op_address (NULL, as_a <hsa_op_reg *> (reg), 0);
4499 case ADDR_EXPR:
4500 return gen_hsa_addr (TREE_OPERAND (val, 0), hbb);
4502 case INTEGER_CST:
4503 if (tree_fits_shwi_p (val))
4504 return new hsa_op_address (NULL, NULL, tree_to_shwi (val));
4505 /* Otherwise fall-through */
4507 default:
4508 HSA_SORRY_ATV (EXPR_LOCATION (val),
4509 "support for HSA does not implement memory access to %E",
4510 val);
4511 return new hsa_op_address (NULL, NULL, 0);
4515 /* Expand assignment of a result of a string BUILTIN to DST.
4516 Size of the operation is N bytes, where instructions
4517 will be append to HBB. */
4519 static void
4520 expand_lhs_of_string_op (gimple *stmt,
4521 unsigned HOST_WIDE_INT n, hsa_bb *hbb,
4522 enum built_in_function builtin)
4524 /* If LHS is expected, we need to emit a PHI instruction. */
4525 tree lhs = gimple_call_lhs (stmt);
4526 if (!lhs)
4527 return;
4529 hsa_op_reg *lhs_reg = hsa_cfun->reg_for_gimple_ssa (lhs);
4531 hsa_op_with_type *dst_reg
4532 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb);
4533 hsa_op_with_type *tmp;
4535 switch (builtin)
4537 case BUILT_IN_MEMPCPY:
4539 tmp = new hsa_op_reg (dst_reg->m_type);
4540 hsa_insn_basic *add
4541 = new hsa_insn_basic (3, BRIG_OPCODE_ADD, tmp->m_type,
4542 tmp, dst_reg,
4543 new hsa_op_immed (n, dst_reg->m_type));
4544 hbb->append_insn (add);
4545 break;
4547 case BUILT_IN_MEMCPY:
4548 case BUILT_IN_MEMSET:
4549 tmp = dst_reg;
4550 break;
4551 default:
4552 gcc_unreachable ();
4555 hbb->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV, lhs_reg->m_type,
4556 lhs_reg, tmp));
4559 #define HSA_MEMORY_BUILTINS_LIMIT 128
4561 /* Expand a string builtin (from a gimple STMT) in a way that
4562 according to MISALIGNED_FLAG we process either direct emission
4563 (a bunch of memory load and store instructions), or we emit a function call
4564 of a library function (for instance 'memcpy'). Actually, a basic block
4565 for direct emission is just prepared, where caller is responsible
4566 for emission of corresponding instructions.
4567 All instruction are appended to HBB. */
4569 hsa_bb *
4570 expand_string_operation_builtin (gimple *stmt, hsa_bb *hbb,
4571 hsa_op_reg *misaligned_flag)
4573 edge e = split_block (hbb->m_bb, stmt);
4574 basic_block condition_bb = e->src;
4575 hbb->append_insn (new hsa_insn_br (misaligned_flag));
4577 /* Prepare the control flow. */
4578 edge condition_edge = EDGE_SUCC (condition_bb, 0);
4579 basic_block call_bb = split_edge (condition_edge);
4581 basic_block expanded_bb = split_edge (EDGE_SUCC (call_bb, 0));
4582 basic_block cont_bb = EDGE_SUCC (expanded_bb, 0)->dest;
4583 basic_block merge_bb = split_edge (EDGE_PRED (cont_bb, 0));
4585 condition_edge->flags &= ~EDGE_FALLTHRU;
4586 condition_edge->flags |= EDGE_TRUE_VALUE;
4587 make_edge (condition_bb, expanded_bb, EDGE_FALSE_VALUE);
4589 redirect_edge_succ (EDGE_SUCC (call_bb, 0), merge_bb);
4591 hsa_cfun->m_modified_cfg = true;
4593 hsa_init_new_bb (expanded_bb);
4595 /* Slow path: function call. */
4596 gen_hsa_insns_for_direct_call (stmt, hsa_init_new_bb (call_bb), false);
4598 return hsa_bb_for_bb (expanded_bb);
4601 /* Expand a memory copy BUILTIN (BUILT_IN_MEMCPY, BUILT_IN_MEMPCPY) from
4602 a gimple STMT and store all necessary instruction to HBB basic block. */
4604 static void
4605 expand_memory_copy (gimple *stmt, hsa_bb *hbb, enum built_in_function builtin)
4607 tree byte_size = gimple_call_arg (stmt, 2);
4609 if (!tree_fits_uhwi_p (byte_size))
4611 gen_hsa_insns_for_direct_call (stmt, hbb);
4612 return;
4615 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size);
4617 if (n > HSA_MEMORY_BUILTINS_LIMIT)
4619 gen_hsa_insns_for_direct_call (stmt, hbb);
4620 return;
4623 tree dst = gimple_call_arg (stmt, 0);
4624 tree src = gimple_call_arg (stmt, 1);
4626 hsa_op_address *dst_addr = get_address_from_value (dst, hbb);
4627 hsa_op_address *src_addr = get_address_from_value (src, hbb);
4629 /* As gen_hsa_memory_copy relies on memory alignment
4630 greater or equal to 8 bytes, we need to verify the alignment. */
4631 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
4632 hsa_op_reg *src_addr_reg = new hsa_op_reg (addrtype);
4633 hsa_op_reg *dst_addr_reg = new hsa_op_reg (addrtype);
4635 convert_addr_to_flat_segment (src_addr, src_addr_reg, hbb);
4636 convert_addr_to_flat_segment (dst_addr, dst_addr_reg, hbb);
4638 /* Process BIT OR for source and destination addresses. */
4639 hsa_op_reg *or_reg = new hsa_op_reg (addrtype);
4640 gen_hsa_binary_operation (BRIG_OPCODE_OR, or_reg, src_addr_reg,
4641 dst_addr_reg, hbb);
4643 /* Process BIT AND with 0x7 to identify the desired alignment
4644 of 8 bytes. */
4645 hsa_op_reg *masked = new hsa_op_reg (addrtype);
4647 gen_hsa_binary_operation (BRIG_OPCODE_AND, masked, or_reg,
4648 new hsa_op_immed (7, addrtype), hbb);
4650 hsa_op_reg *misaligned = new hsa_op_reg (BRIG_TYPE_B1);
4651 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE, misaligned->m_type,
4652 misaligned, masked,
4653 new hsa_op_immed (0, masked->m_type)));
4655 hsa_bb *native_impl_bb
4656 = expand_string_operation_builtin (stmt, hbb, misaligned);
4658 gen_hsa_memory_copy (native_impl_bb, dst_addr, src_addr, n, BRIG_ALIGNMENT_8);
4659 hsa_bb *merge_bb
4660 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb->m_bb, 0)->dest);
4661 expand_lhs_of_string_op (stmt, n, merge_bb, builtin);
4665 /* Expand a memory set BUILTIN (BUILT_IN_MEMSET, BUILT_IN_BZERO) from
4666 a gimple STMT and store all necessary instruction to HBB basic block.
4667 The operation set N bytes with a CONSTANT value. */
4669 static void
4670 expand_memory_set (gimple *stmt, unsigned HOST_WIDE_INT n,
4671 unsigned HOST_WIDE_INT constant, hsa_bb *hbb,
4672 enum built_in_function builtin)
4674 tree dst = gimple_call_arg (stmt, 0);
4675 hsa_op_address *dst_addr = get_address_from_value (dst, hbb);
4677 /* As gen_hsa_memory_set relies on memory alignment
4678 greater or equal to 8 bytes, we need to verify the alignment. */
4679 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
4680 hsa_op_reg *dst_addr_reg = new hsa_op_reg (addrtype);
4681 convert_addr_to_flat_segment (dst_addr, dst_addr_reg, hbb);
4683 /* Process BIT AND with 0x7 to identify the desired alignment
4684 of 8 bytes. */
4685 hsa_op_reg *masked = new hsa_op_reg (addrtype);
4687 gen_hsa_binary_operation (BRIG_OPCODE_AND, masked, dst_addr_reg,
4688 new hsa_op_immed (7, addrtype), hbb);
4690 hsa_op_reg *misaligned = new hsa_op_reg (BRIG_TYPE_B1);
4691 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE, misaligned->m_type,
4692 misaligned, masked,
4693 new hsa_op_immed (0, masked->m_type)));
4695 hsa_bb *native_impl_bb
4696 = expand_string_operation_builtin (stmt, hbb, misaligned);
4698 gen_hsa_memory_set (native_impl_bb, dst_addr, constant, n, BRIG_ALIGNMENT_8);
4699 hsa_bb *merge_bb
4700 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb->m_bb, 0)->dest);
4701 expand_lhs_of_string_op (stmt, n, merge_bb, builtin);
4704 /* Return string for MEMMODEL. */
4706 static const char *
4707 get_memory_order_name (unsigned memmodel)
4709 switch (memmodel & MEMMODEL_BASE_MASK)
4711 case MEMMODEL_RELAXED:
4712 return "relaxed";
4713 case MEMMODEL_CONSUME:
4714 return "consume";
4715 case MEMMODEL_ACQUIRE:
4716 return "acquire";
4717 case MEMMODEL_RELEASE:
4718 return "release";
4719 case MEMMODEL_ACQ_REL:
4720 return "acq_rel";
4721 case MEMMODEL_SEQ_CST:
4722 return "seq_cst";
4723 default:
4724 return NULL;
4728 /* Return memory order according to predefined __atomic memory model
4729 constants. LOCATION is provided to locate the problematic statement. */
4731 static BrigMemoryOrder
4732 get_memory_order (unsigned memmodel, location_t location)
4734 switch (memmodel & MEMMODEL_BASE_MASK)
4736 case MEMMODEL_RELAXED:
4737 return BRIG_MEMORY_ORDER_RELAXED;
4738 case MEMMODEL_CONSUME:
4739 /* HSA does not have an equivalent, but we can use the slightly stronger
4740 ACQUIRE. */
4741 case MEMMODEL_ACQUIRE:
4742 return BRIG_MEMORY_ORDER_SC_ACQUIRE;
4743 case MEMMODEL_RELEASE:
4744 return BRIG_MEMORY_ORDER_SC_RELEASE;
4745 case MEMMODEL_ACQ_REL:
4746 case MEMMODEL_SEQ_CST:
4747 /* Callers implementing a simple load or store need to remove the release
4748 or acquire part respectively. */
4749 return BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE;
4750 default:
4752 const char *mmname = get_memory_order_name (memmodel);
4753 HSA_SORRY_ATV (location,
4754 "support for HSA does not implement the specified "
4755 " memory model%s %s",
4756 mmname ? ": " : "", mmname ? mmname : "");
4757 return BRIG_MEMORY_ORDER_NONE;
4762 /* Helper function to create an HSA atomic binary operation instruction out of
4763 calls to atomic builtins. RET_ORIG is true if the built-in is the variant
4764 that return s the value before applying operation, and false if it should
4765 return the value after applying the operation (if it returns value at all).
4766 ACODE is the atomic operation code, STMT is a gimple call to a builtin. HBB
4767 is the HSA BB to which the instruction should be added. */
4769 static void
4770 gen_hsa_ternary_atomic_for_builtin (bool ret_orig,
4771 enum BrigAtomicOperation acode,
4772 gimple *stmt,
4773 hsa_bb *hbb)
4775 tree lhs = gimple_call_lhs (stmt);
4777 tree type = TREE_TYPE (gimple_call_arg (stmt, 1));
4778 BrigType16_t hsa_type = hsa_type_for_scalar_tree_type (type, false);
4779 BrigType16_t mtype = mem_type_for_type (hsa_type);
4780 tree model = gimple_call_arg (stmt, 2);
4782 if (!tree_fits_uhwi_p (model))
4784 HSA_SORRY_ATV (gimple_location (stmt),
4785 "support for HSA does not implement memory model %E",
4786 model);
4787 return;
4790 unsigned HOST_WIDE_INT mmodel = tree_to_uhwi (model);
4792 BrigMemoryOrder memorder = get_memory_order (mmodel, gimple_location (stmt));
4794 /* Certain atomic insns must have Bx memory types. */
4795 switch (acode)
4797 case BRIG_ATOMIC_LD:
4798 case BRIG_ATOMIC_ST:
4799 case BRIG_ATOMIC_AND:
4800 case BRIG_ATOMIC_OR:
4801 case BRIG_ATOMIC_XOR:
4802 case BRIG_ATOMIC_EXCH:
4803 mtype = hsa_bittype_for_type (mtype);
4804 break;
4805 default:
4806 break;
4809 hsa_op_reg *dest;
4810 int nops, opcode;
4811 if (lhs)
4813 if (ret_orig)
4814 dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4815 else
4816 dest = new hsa_op_reg (hsa_type);
4817 opcode = BRIG_OPCODE_ATOMIC;
4818 nops = 3;
4820 else
4822 dest = NULL;
4823 opcode = BRIG_OPCODE_ATOMICNORET;
4824 nops = 2;
4827 if (acode == BRIG_ATOMIC_ST)
4829 if (memorder == BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE)
4830 memorder = BRIG_MEMORY_ORDER_SC_RELEASE;
4832 if (memorder != BRIG_MEMORY_ORDER_RELAXED
4833 && memorder != BRIG_MEMORY_ORDER_SC_RELEASE
4834 && memorder != BRIG_MEMORY_ORDER_NONE)
4836 HSA_SORRY_ATV (gimple_location (stmt),
4837 "support for HSA does not implement memory model for "
4838 "ATOMIC_ST: %s", get_memory_order_name (mmodel));
4839 return;
4843 hsa_insn_atomic *atominsn = new hsa_insn_atomic (nops, opcode, acode, mtype,
4844 memorder);
4846 hsa_op_address *addr;
4847 addr = get_address_from_value (gimple_call_arg (stmt, 0), hbb);
4848 if (addr->m_symbol && addr->m_symbol->m_segment == BRIG_SEGMENT_PRIVATE)
4850 HSA_SORRY_AT (gimple_location (stmt),
4851 "HSA does not implement atomic operations in private "
4852 "segment");
4853 return;
4855 hsa_op_base *op = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1),
4856 hbb);
4858 if (lhs)
4860 atominsn->set_op (0, dest);
4861 atominsn->set_op (1, addr);
4862 atominsn->set_op (2, op);
4864 else
4866 atominsn->set_op (0, addr);
4867 atominsn->set_op (1, op);
4870 hbb->append_insn (atominsn);
4872 /* HSA does not natively support the variants that return the modified value,
4873 so re-do the operation again non-atomically if that is what was
4874 requested. */
4875 if (lhs && !ret_orig)
4877 int arith;
4878 switch (acode)
4880 case BRIG_ATOMIC_ADD:
4881 arith = BRIG_OPCODE_ADD;
4882 break;
4883 case BRIG_ATOMIC_AND:
4884 arith = BRIG_OPCODE_AND;
4885 break;
4886 case BRIG_ATOMIC_OR:
4887 arith = BRIG_OPCODE_OR;
4888 break;
4889 case BRIG_ATOMIC_SUB:
4890 arith = BRIG_OPCODE_SUB;
4891 break;
4892 case BRIG_ATOMIC_XOR:
4893 arith = BRIG_OPCODE_XOR;
4894 break;
4895 default:
4896 gcc_unreachable ();
4898 hsa_op_reg *real_dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4899 gen_hsa_binary_operation (arith, real_dest, dest, op, hbb);
4903 /* Generate HSA instructions for an internal fn.
4904 Instructions will be appended to HBB, which also needs to be the
4905 corresponding structure to the basic_block of STMT. */
4907 static void
4908 gen_hsa_insn_for_internal_fn_call (gcall *stmt, hsa_bb *hbb)
4910 gcc_checking_assert (gimple_call_internal_fn (stmt));
4911 internal_fn fn = gimple_call_internal_fn (stmt);
4913 bool is_float_type_p = false;
4914 if (gimple_call_lhs (stmt) != NULL
4915 && TREE_TYPE (gimple_call_lhs (stmt)) == float_type_node)
4916 is_float_type_p = true;
4918 switch (fn)
4920 case IFN_CEIL:
4921 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL, stmt, hbb);
4922 break;
4924 case IFN_FLOOR:
4925 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR, stmt, hbb);
4926 break;
4928 case IFN_RINT:
4929 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT, stmt, hbb);
4930 break;
4932 case IFN_SQRT:
4933 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT, stmt, hbb);
4934 break;
4936 case IFN_TRUNC:
4937 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC, stmt, hbb);
4938 break;
4940 case IFN_COS:
4942 if (is_float_type_p)
4943 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS, stmt, hbb);
4944 else
4945 gen_hsa_unaryop_builtin_call (stmt, hbb);
4947 break;
4949 case IFN_EXP2:
4951 if (is_float_type_p)
4952 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2, stmt, hbb);
4953 else
4954 gen_hsa_unaryop_builtin_call (stmt, hbb);
4956 break;
4959 case IFN_LOG2:
4961 if (is_float_type_p)
4962 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2, stmt, hbb);
4963 else
4964 gen_hsa_unaryop_builtin_call (stmt, hbb);
4966 break;
4969 case IFN_SIN:
4971 if (is_float_type_p)
4972 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN, stmt, hbb);
4973 else
4974 gen_hsa_unaryop_builtin_call (stmt, hbb);
4975 break;
4978 case IFN_CLRSB:
4979 gen_hsa_clrsb (stmt, hbb);
4980 break;
4982 case IFN_CLZ:
4983 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT, stmt, hbb);
4984 break;
4986 case IFN_CTZ:
4987 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT, stmt, hbb);
4988 break;
4990 case IFN_FFS:
4991 gen_hsa_ffs (stmt, hbb);
4992 break;
4994 case IFN_PARITY:
4995 gen_hsa_parity (stmt, hbb);
4996 break;
4998 case IFN_POPCOUNT:
4999 gen_hsa_popcount (stmt, hbb);
5000 break;
5002 case IFN_ACOS:
5003 case IFN_ASIN:
5004 case IFN_ATAN:
5005 case IFN_EXP:
5006 case IFN_EXP10:
5007 case IFN_EXPM1:
5008 case IFN_LOG:
5009 case IFN_LOG10:
5010 case IFN_LOG1P:
5011 case IFN_LOGB:
5012 case IFN_SIGNIFICAND:
5013 case IFN_TAN:
5014 case IFN_NEARBYINT:
5015 case IFN_ROUND:
5016 case IFN_ATAN2:
5017 case IFN_COPYSIGN:
5018 case IFN_FMOD:
5019 case IFN_POW:
5020 case IFN_REMAINDER:
5021 case IFN_SCALB:
5022 case IFN_FMIN:
5023 case IFN_FMAX:
5024 gen_hsa_insns_for_call_of_internal_fn (stmt, hbb);
5026 default:
5027 HSA_SORRY_ATV (gimple_location (stmt),
5028 "support for HSA does not implement internal function: %s",
5029 internal_fn_name (fn));
5030 break;
5034 /* Generate HSA instructions for the given call statement STMT. Instructions
5035 will be appended to HBB. */
5037 static void
5038 gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb)
5040 gcall *call = as_a <gcall *> (stmt);
5041 tree lhs = gimple_call_lhs (stmt);
5042 hsa_op_reg *dest;
5044 if (gimple_call_internal_p (stmt))
5046 gen_hsa_insn_for_internal_fn_call (call, hbb);
5047 return;
5050 if (!gimple_call_builtin_p (stmt, BUILT_IN_NORMAL))
5052 tree function_decl = gimple_call_fndecl (stmt);
5053 if (function_decl == NULL_TREE)
5055 HSA_SORRY_AT (gimple_location (stmt),
5056 "support for HSA does not implement indirect calls");
5057 return;
5060 if (hsa_callable_function_p (function_decl))
5061 gen_hsa_insns_for_direct_call (stmt, hbb);
5062 else if (!gen_hsa_insns_for_known_library_call (stmt, hbb))
5063 HSA_SORRY_AT (gimple_location (stmt),
5064 "HSA supports only calls of functions marked with pragma "
5065 "omp declare target");
5066 return;
5069 tree fndecl = gimple_call_fndecl (stmt);
5070 enum built_in_function builtin = DECL_FUNCTION_CODE (fndecl);
5071 switch (builtin)
5073 case BUILT_IN_FABS:
5074 case BUILT_IN_FABSF:
5075 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_ABS, stmt, hbb);
5076 break;
5078 case BUILT_IN_CEIL:
5079 case BUILT_IN_CEILF:
5080 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL, stmt, hbb);
5081 break;
5083 case BUILT_IN_FLOOR:
5084 case BUILT_IN_FLOORF:
5085 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR, stmt, hbb);
5086 break;
5088 case BUILT_IN_RINT:
5089 case BUILT_IN_RINTF:
5090 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT, stmt, hbb);
5091 break;
5093 case BUILT_IN_SQRT:
5094 case BUILT_IN_SQRTF:
5095 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT, stmt, hbb);
5096 break;
5098 case BUILT_IN_TRUNC:
5099 case BUILT_IN_TRUNCF:
5100 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC, stmt, hbb);
5101 break;
5103 case BUILT_IN_COS:
5104 case BUILT_IN_SIN:
5105 case BUILT_IN_EXP2:
5106 case BUILT_IN_LOG2:
5107 /* HSAIL does not provide an instruction for double argument type. */
5108 gen_hsa_unaryop_builtin_call (stmt, hbb);
5109 break;
5111 case BUILT_IN_COSF:
5112 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS, stmt, hbb);
5113 break;
5115 case BUILT_IN_EXP2F:
5116 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2, stmt, hbb);
5117 break;
5119 case BUILT_IN_LOG2F:
5120 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2, stmt, hbb);
5121 break;
5123 case BUILT_IN_SINF:
5124 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN, stmt, hbb);
5125 break;
5127 case BUILT_IN_CLRSB:
5128 case BUILT_IN_CLRSBL:
5129 case BUILT_IN_CLRSBLL:
5130 gen_hsa_clrsb (call, hbb);
5131 break;
5133 case BUILT_IN_CLZ:
5134 case BUILT_IN_CLZL:
5135 case BUILT_IN_CLZLL:
5136 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT, stmt, hbb);
5137 break;
5139 case BUILT_IN_CTZ:
5140 case BUILT_IN_CTZL:
5141 case BUILT_IN_CTZLL:
5142 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT, stmt, hbb);
5143 break;
5145 case BUILT_IN_FFS:
5146 case BUILT_IN_FFSL:
5147 case BUILT_IN_FFSLL:
5148 gen_hsa_ffs (call, hbb);
5149 break;
5151 case BUILT_IN_PARITY:
5152 case BUILT_IN_PARITYL:
5153 case BUILT_IN_PARITYLL:
5154 gen_hsa_parity (call, hbb);
5155 break;
5157 case BUILT_IN_POPCOUNT:
5158 case BUILT_IN_POPCOUNTL:
5159 case BUILT_IN_POPCOUNTLL:
5160 gen_hsa_popcount (call, hbb);
5161 break;
5163 case BUILT_IN_ATOMIC_LOAD_1:
5164 case BUILT_IN_ATOMIC_LOAD_2:
5165 case BUILT_IN_ATOMIC_LOAD_4:
5166 case BUILT_IN_ATOMIC_LOAD_8:
5167 case BUILT_IN_ATOMIC_LOAD_16:
5169 BrigType16_t mtype;
5170 hsa_op_address *addr;
5171 addr = get_address_from_value (gimple_call_arg (stmt, 0), hbb);
5172 tree model = gimple_call_arg (stmt, 1);
5173 if (!tree_fits_uhwi_p (model))
5175 HSA_SORRY_ATV (gimple_location (stmt),
5176 "support for HSA does not implement "
5177 "memory model: %E",
5178 model);
5179 return;
5182 unsigned HOST_WIDE_INT mmodel = tree_to_uhwi (model);
5183 BrigMemoryOrder memorder = get_memory_order (mmodel,
5184 gimple_location (stmt));
5186 if (memorder == BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE)
5187 memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE;
5189 if (memorder != BRIG_MEMORY_ORDER_RELAXED
5190 && memorder != BRIG_MEMORY_ORDER_SC_ACQUIRE
5191 && memorder != BRIG_MEMORY_ORDER_NONE)
5193 HSA_SORRY_ATV (gimple_location (stmt),
5194 "support for HSA does not implement "
5195 "memory model for ATOMIC_LD: %s",
5196 get_memory_order_name (mmodel));
5197 return;
5200 if (lhs)
5202 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
5203 false);
5204 mtype = mem_type_for_type (t);
5205 mtype = hsa_bittype_for_type (mtype);
5206 dest = hsa_cfun->reg_for_gimple_ssa (lhs);
5208 else
5210 mtype = BRIG_TYPE_B64;
5211 dest = new hsa_op_reg (mtype);
5214 hsa_insn_atomic *atominsn
5215 = new hsa_insn_atomic (2, BRIG_OPCODE_ATOMIC, BRIG_ATOMIC_LD, mtype,
5216 memorder, dest, addr);
5218 hbb->append_insn (atominsn);
5219 break;
5222 case BUILT_IN_ATOMIC_EXCHANGE_1:
5223 case BUILT_IN_ATOMIC_EXCHANGE_2:
5224 case BUILT_IN_ATOMIC_EXCHANGE_4:
5225 case BUILT_IN_ATOMIC_EXCHANGE_8:
5226 case BUILT_IN_ATOMIC_EXCHANGE_16:
5227 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_EXCH, stmt, hbb);
5228 break;
5230 case BUILT_IN_ATOMIC_FETCH_ADD_1:
5231 case BUILT_IN_ATOMIC_FETCH_ADD_2:
5232 case BUILT_IN_ATOMIC_FETCH_ADD_4:
5233 case BUILT_IN_ATOMIC_FETCH_ADD_8:
5234 case BUILT_IN_ATOMIC_FETCH_ADD_16:
5235 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_ADD, stmt, hbb);
5236 break;
5238 case BUILT_IN_ATOMIC_FETCH_SUB_1:
5239 case BUILT_IN_ATOMIC_FETCH_SUB_2:
5240 case BUILT_IN_ATOMIC_FETCH_SUB_4:
5241 case BUILT_IN_ATOMIC_FETCH_SUB_8:
5242 case BUILT_IN_ATOMIC_FETCH_SUB_16:
5243 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_SUB, stmt, hbb);
5244 break;
5246 case BUILT_IN_ATOMIC_FETCH_AND_1:
5247 case BUILT_IN_ATOMIC_FETCH_AND_2:
5248 case BUILT_IN_ATOMIC_FETCH_AND_4:
5249 case BUILT_IN_ATOMIC_FETCH_AND_8:
5250 case BUILT_IN_ATOMIC_FETCH_AND_16:
5251 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_AND, stmt, hbb);
5252 break;
5254 case BUILT_IN_ATOMIC_FETCH_XOR_1:
5255 case BUILT_IN_ATOMIC_FETCH_XOR_2:
5256 case BUILT_IN_ATOMIC_FETCH_XOR_4:
5257 case BUILT_IN_ATOMIC_FETCH_XOR_8:
5258 case BUILT_IN_ATOMIC_FETCH_XOR_16:
5259 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_XOR, stmt, hbb);
5260 break;
5262 case BUILT_IN_ATOMIC_FETCH_OR_1:
5263 case BUILT_IN_ATOMIC_FETCH_OR_2:
5264 case BUILT_IN_ATOMIC_FETCH_OR_4:
5265 case BUILT_IN_ATOMIC_FETCH_OR_8:
5266 case BUILT_IN_ATOMIC_FETCH_OR_16:
5267 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_OR, stmt, hbb);
5268 break;
5270 case BUILT_IN_ATOMIC_STORE_1:
5271 case BUILT_IN_ATOMIC_STORE_2:
5272 case BUILT_IN_ATOMIC_STORE_4:
5273 case BUILT_IN_ATOMIC_STORE_8:
5274 case BUILT_IN_ATOMIC_STORE_16:
5275 /* Since there cannot be any LHS, the first parameter is meaningless. */
5276 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_ST, stmt, hbb);
5277 break;
5279 case BUILT_IN_ATOMIC_ADD_FETCH_1:
5280 case BUILT_IN_ATOMIC_ADD_FETCH_2:
5281 case BUILT_IN_ATOMIC_ADD_FETCH_4:
5282 case BUILT_IN_ATOMIC_ADD_FETCH_8:
5283 case BUILT_IN_ATOMIC_ADD_FETCH_16:
5284 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_ADD, stmt, hbb);
5285 break;
5287 case BUILT_IN_ATOMIC_SUB_FETCH_1:
5288 case BUILT_IN_ATOMIC_SUB_FETCH_2:
5289 case BUILT_IN_ATOMIC_SUB_FETCH_4:
5290 case BUILT_IN_ATOMIC_SUB_FETCH_8:
5291 case BUILT_IN_ATOMIC_SUB_FETCH_16:
5292 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_SUB, stmt, hbb);
5293 break;
5295 case BUILT_IN_ATOMIC_AND_FETCH_1:
5296 case BUILT_IN_ATOMIC_AND_FETCH_2:
5297 case BUILT_IN_ATOMIC_AND_FETCH_4:
5298 case BUILT_IN_ATOMIC_AND_FETCH_8:
5299 case BUILT_IN_ATOMIC_AND_FETCH_16:
5300 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_AND, stmt, hbb);
5301 break;
5303 case BUILT_IN_ATOMIC_XOR_FETCH_1:
5304 case BUILT_IN_ATOMIC_XOR_FETCH_2:
5305 case BUILT_IN_ATOMIC_XOR_FETCH_4:
5306 case BUILT_IN_ATOMIC_XOR_FETCH_8:
5307 case BUILT_IN_ATOMIC_XOR_FETCH_16:
5308 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_XOR, stmt, hbb);
5309 break;
5311 case BUILT_IN_ATOMIC_OR_FETCH_1:
5312 case BUILT_IN_ATOMIC_OR_FETCH_2:
5313 case BUILT_IN_ATOMIC_OR_FETCH_4:
5314 case BUILT_IN_ATOMIC_OR_FETCH_8:
5315 case BUILT_IN_ATOMIC_OR_FETCH_16:
5316 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_OR, stmt, hbb);
5317 break;
5319 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_1:
5320 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_2:
5321 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_4:
5322 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8:
5323 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_16:
5325 /* TODO: Use the appropriate memory model for now. */
5326 tree type = TREE_TYPE (gimple_call_arg (stmt, 1));
5328 BrigType16_t atype
5329 = hsa_bittype_for_type (hsa_type_for_scalar_tree_type (type, false));
5331 hsa_insn_atomic *atominsn
5332 = new hsa_insn_atomic (4, BRIG_OPCODE_ATOMIC, BRIG_ATOMIC_CAS, atype,
5333 BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE);
5334 hsa_op_address *addr;
5335 addr = get_address_from_value (gimple_call_arg (stmt, 0), hbb);
5337 if (lhs != NULL)
5338 dest = hsa_cfun->reg_for_gimple_ssa (lhs);
5339 else
5340 dest = new hsa_op_reg (atype);
5342 /* Should check what the memory scope is. */
5343 atominsn->m_memoryscope = BRIG_MEMORY_SCOPE_WORKGROUP;
5344 atominsn->set_op (0, dest);
5345 atominsn->set_op (1, addr);
5347 hsa_op_with_type *op
5348 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1), hbb);
5349 atominsn->set_op (2, op);
5350 op = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 2), hbb);
5351 atominsn->set_op (3, op);
5353 hbb->append_insn (atominsn);
5354 break;
5356 case BUILT_IN_GOMP_PARALLEL:
5357 HSA_SORRY_AT (gimple_location (stmt),
5358 "support for HSA does not implement non-gridified "
5359 "OpenMP parallel constructs.");
5360 break;
5361 case BUILT_IN_OMP_GET_THREAD_NUM:
5363 query_hsa_grid (stmt, BRIG_OPCODE_WORKITEMABSID, 0, hbb);
5364 break;
5367 case BUILT_IN_OMP_GET_NUM_THREADS:
5369 query_hsa_grid (stmt, BRIG_OPCODE_GRIDSIZE, 0, hbb);
5370 break;
5372 case BUILT_IN_GOMP_TEAMS:
5374 gen_set_num_threads (gimple_call_arg (stmt, 1), hbb);
5375 break;
5377 case BUILT_IN_OMP_GET_NUM_TEAMS:
5379 gen_get_num_teams (stmt, hbb);
5380 break;
5382 case BUILT_IN_OMP_GET_TEAM_NUM:
5384 gen_get_team_num (stmt, hbb);
5385 break;
5387 case BUILT_IN_MEMCPY:
5388 case BUILT_IN_MEMPCPY:
5390 expand_memory_copy (stmt, hbb, builtin);
5391 break;
5393 case BUILT_IN_MEMSET:
5395 tree c = gimple_call_arg (stmt, 1);
5397 if (TREE_CODE (c) != INTEGER_CST)
5399 gen_hsa_insns_for_direct_call (stmt, hbb);
5400 return;
5403 tree byte_size = gimple_call_arg (stmt, 2);
5405 if (!tree_fits_uhwi_p (byte_size))
5407 gen_hsa_insns_for_direct_call (stmt, hbb);
5408 return;
5411 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size);
5413 if (n > HSA_MEMORY_BUILTINS_LIMIT)
5415 gen_hsa_insns_for_direct_call (stmt, hbb);
5416 return;
5419 unsigned HOST_WIDE_INT constant
5420 = tree_to_uhwi (fold_convert (unsigned_char_type_node, c));
5422 expand_memory_set (stmt, n, constant, hbb, builtin);
5424 break;
5426 case BUILT_IN_BZERO:
5428 tree byte_size = gimple_call_arg (stmt, 1);
5430 if (!tree_fits_uhwi_p (byte_size))
5432 gen_hsa_insns_for_direct_call (stmt, hbb);
5433 return;
5436 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size);
5438 if (n > HSA_MEMORY_BUILTINS_LIMIT)
5440 gen_hsa_insns_for_direct_call (stmt, hbb);
5441 return;
5444 expand_memory_set (stmt, n, 0, hbb, builtin);
5446 break;
5448 case BUILT_IN_ALLOCA:
5449 case BUILT_IN_ALLOCA_WITH_ALIGN:
5451 gen_hsa_alloca (call, hbb);
5452 break;
5454 default:
5456 gen_hsa_insns_for_direct_call (stmt, hbb);
5457 return;
5462 /* Generate HSA instructions for a given gimple statement. Instructions will be
5463 appended to HBB. */
5465 static void
5466 gen_hsa_insns_for_gimple_stmt (gimple *stmt, hsa_bb *hbb)
5468 switch (gimple_code (stmt))
5470 case GIMPLE_ASSIGN:
5471 if (gimple_clobber_p (stmt))
5472 break;
5474 if (gimple_assign_single_p (stmt))
5476 tree lhs = gimple_assign_lhs (stmt);
5477 tree rhs = gimple_assign_rhs1 (stmt);
5478 gen_hsa_insns_for_single_assignment (lhs, rhs, hbb);
5480 else
5481 gen_hsa_insns_for_operation_assignment (stmt, hbb);
5482 break;
5483 case GIMPLE_RETURN:
5484 gen_hsa_insns_for_return (as_a <greturn *> (stmt), hbb);
5485 break;
5486 case GIMPLE_COND:
5487 gen_hsa_insns_for_cond_stmt (stmt, hbb);
5488 break;
5489 case GIMPLE_CALL:
5490 gen_hsa_insns_for_call (stmt, hbb);
5491 break;
5492 case GIMPLE_DEBUG:
5493 /* ??? HSA supports some debug facilities. */
5494 break;
5495 case GIMPLE_LABEL:
5497 tree label = gimple_label_label (as_a <glabel *> (stmt));
5498 if (FORCED_LABEL (label))
5499 HSA_SORRY_AT (gimple_location (stmt),
5500 "support for HSA does not implement gimple label with "
5501 "address taken");
5503 break;
5505 case GIMPLE_NOP:
5507 hbb->append_insn (new hsa_insn_basic (0, BRIG_OPCODE_NOP));
5508 break;
5510 case GIMPLE_SWITCH:
5512 gen_hsa_insns_for_switch_stmt (as_a <gswitch *> (stmt), hbb);
5513 break;
5515 default:
5516 HSA_SORRY_ATV (gimple_location (stmt),
5517 "support for HSA does not implement gimple statement %s",
5518 gimple_code_name[(int) gimple_code (stmt)]);
5522 /* Generate a HSA PHI from a gimple PHI. */
5524 static void
5525 gen_hsa_phi_from_gimple_phi (gimple *phi_stmt, hsa_bb *hbb)
5527 hsa_insn_phi *hphi;
5528 unsigned count = gimple_phi_num_args (phi_stmt);
5530 hsa_op_reg *dest
5531 = hsa_cfun->reg_for_gimple_ssa (gimple_phi_result (phi_stmt));
5532 hphi = new hsa_insn_phi (count, dest);
5533 hphi->m_bb = hbb->m_bb;
5535 tree lhs = gimple_phi_result (phi_stmt);
5537 for (unsigned i = 0; i < count; i++)
5539 tree op = gimple_phi_arg_def (phi_stmt, i);
5541 if (TREE_CODE (op) == SSA_NAME)
5543 hsa_op_reg *hreg = hsa_cfun->reg_for_gimple_ssa (op);
5544 hphi->set_op (i, hreg);
5546 else
5548 gcc_assert (is_gimple_min_invariant (op));
5549 tree t = TREE_TYPE (op);
5550 if (!POINTER_TYPE_P (t)
5551 || (TREE_CODE (op) == STRING_CST
5552 && TREE_CODE (TREE_TYPE (t)) == INTEGER_TYPE))
5553 hphi->set_op (i, new hsa_op_immed (op));
5554 else if (POINTER_TYPE_P (TREE_TYPE (lhs))
5555 && TREE_CODE (op) == INTEGER_CST)
5557 /* Handle assignment of NULL value to a pointer type. */
5558 hphi->set_op (i, new hsa_op_immed (op));
5560 else if (TREE_CODE (op) == ADDR_EXPR)
5562 edge e = gimple_phi_arg_edge (as_a <gphi *> (phi_stmt), i);
5563 hsa_bb *hbb_src = hsa_init_new_bb (split_edge (e));
5564 hsa_op_address *addr = gen_hsa_addr (TREE_OPERAND (op, 0),
5565 hbb_src);
5567 hsa_op_reg *dest
5568 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
5569 hsa_insn_basic *insn
5570 = new hsa_insn_basic (2, BRIG_OPCODE_LDA, BRIG_TYPE_U64,
5571 dest, addr);
5572 hbb_src->append_insn (insn);
5574 hphi->set_op (i, dest);
5576 else
5578 HSA_SORRY_AT (gimple_location (phi_stmt),
5579 "support for HSA does not handle PHI nodes with "
5580 "constant address operands");
5581 return;
5586 hphi->m_prev = hbb->m_last_phi;
5587 hphi->m_next = NULL;
5588 if (hbb->m_last_phi)
5589 hbb->m_last_phi->m_next = hphi;
5590 hbb->m_last_phi = hphi;
5591 if (!hbb->m_first_phi)
5592 hbb->m_first_phi = hphi;
5595 /* Constructor of class containing HSA-specific information about a basic
5596 block. CFG_BB is the CFG BB this HSA BB is associated with. IDX is the new
5597 index of this BB (so that the constructor does not attempt to use
5598 hsa_cfun during its construction). */
5600 hsa_bb::hsa_bb (basic_block cfg_bb, int idx)
5601 : m_bb (cfg_bb), m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL),
5602 m_last_phi (NULL), m_index (idx), m_liveout (BITMAP_ALLOC (NULL)),
5603 m_livein (BITMAP_ALLOC (NULL))
5605 gcc_assert (!cfg_bb->aux);
5606 cfg_bb->aux = this;
5609 /* Constructor of class containing HSA-specific information about a basic
5610 block. CFG_BB is the CFG BB this HSA BB is associated with. */
5612 hsa_bb::hsa_bb (basic_block cfg_bb)
5613 : m_bb (cfg_bb), m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL),
5614 m_last_phi (NULL), m_index (hsa_cfun->m_hbb_count++),
5615 m_liveout (BITMAP_ALLOC (NULL)), m_livein (BITMAP_ALLOC (NULL))
5617 gcc_assert (!cfg_bb->aux);
5618 cfg_bb->aux = this;
5621 /* Destructor of class representing HSA BB. */
5623 hsa_bb::~hsa_bb ()
5625 BITMAP_FREE (m_livein);
5626 BITMAP_FREE (m_liveout);
5629 /* Create and initialize and return a new hsa_bb structure for a given CFG
5630 basic block BB. */
5632 hsa_bb *
5633 hsa_init_new_bb (basic_block bb)
5635 return new (*hsa_allocp_bb) hsa_bb (bb);
5638 /* Initialize OMP in an HSA basic block PROLOGUE. */
5640 static void
5641 init_prologue (void)
5643 if (!hsa_cfun->m_kern_p)
5644 return;
5646 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
5648 /* Create a magic number that is going to be printed by libgomp. */
5649 unsigned index = hsa_get_number_decl_kernel_mappings ();
5651 /* Emit store to debug argument. */
5652 if (PARAM_VALUE (PARAM_HSA_GEN_DEBUG_STORES) > 0)
5653 set_debug_value (prologue, new hsa_op_immed (1000 + index, BRIG_TYPE_U64));
5656 /* Initialize hsa_num_threads to a default value. */
5658 static void
5659 init_hsa_num_threads (void)
5661 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
5663 /* Save the default value to private variable hsa_num_threads. */
5664 hsa_insn_basic *basic
5665 = new hsa_insn_mem (BRIG_OPCODE_ST, hsa_num_threads->m_type,
5666 new hsa_op_immed (0, hsa_num_threads->m_type),
5667 new hsa_op_address (hsa_num_threads));
5668 prologue->append_insn (basic);
5671 /* Go over gimple representation and generate our internal HSA one. */
5673 static void
5674 gen_body_from_gimple ()
5676 basic_block bb;
5678 /* Verify CFG for complex edges we are unable to handle. */
5679 edge_iterator ei;
5680 edge e;
5682 FOR_EACH_BB_FN (bb, cfun)
5684 FOR_EACH_EDGE (e, ei, bb->succs)
5686 /* Verify all unsupported flags for edges that point
5687 to the same basic block. */
5688 if (e->flags & EDGE_EH)
5690 HSA_SORRY_AT (UNKNOWN_LOCATION,
5691 "support for HSA does not implement exception "
5692 "handling");
5693 return;
5698 FOR_EACH_BB_FN (bb, cfun)
5700 gimple_stmt_iterator gsi;
5701 hsa_bb *hbb = hsa_bb_for_bb (bb);
5702 if (hbb)
5703 continue;
5705 hbb = hsa_init_new_bb (bb);
5707 for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
5709 gen_hsa_insns_for_gimple_stmt (gsi_stmt (gsi), hbb);
5710 if (hsa_seen_error ())
5711 return;
5715 FOR_EACH_BB_FN (bb, cfun)
5717 gimple_stmt_iterator gsi;
5718 hsa_bb *hbb = hsa_bb_for_bb (bb);
5719 gcc_assert (hbb != NULL);
5721 for (gsi = gsi_start_phis (bb); !gsi_end_p (gsi); gsi_next (&gsi))
5722 if (!virtual_operand_p (gimple_phi_result (gsi_stmt (gsi))))
5723 gen_hsa_phi_from_gimple_phi (gsi_stmt (gsi), hbb);
5726 if (dump_file && (dump_flags & TDF_DETAILS))
5728 fprintf (dump_file, "------- Generated SSA form -------\n");
5729 dump_hsa_cfun (dump_file);
5733 static void
5734 gen_function_decl_parameters (hsa_function_representation *f,
5735 tree decl)
5737 tree parm;
5738 unsigned i;
5740 for (parm = TYPE_ARG_TYPES (TREE_TYPE (decl)), i = 0;
5741 parm;
5742 parm = TREE_CHAIN (parm), i++)
5744 /* Result type if last in the tree list. */
5745 if (TREE_CHAIN (parm) == NULL)
5746 break;
5748 tree v = TREE_VALUE (parm);
5750 hsa_symbol *arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
5751 BRIG_LINKAGE_NONE);
5752 arg->m_type = hsa_type_for_tree_type (v, &arg->m_dim);
5753 arg->m_name_number = i;
5755 f->m_input_args.safe_push (arg);
5758 tree result_type = TREE_TYPE (TREE_TYPE (decl));
5759 if (!VOID_TYPE_P (result_type))
5761 f->m_output_arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
5762 BRIG_LINKAGE_NONE);
5763 f->m_output_arg->m_type
5764 = hsa_type_for_tree_type (result_type, &f->m_output_arg->m_dim);
5765 f->m_output_arg->m_name = "res";
5769 /* Generate the vector of parameters of the HSA representation of the current
5770 function. This also includes the output parameter representing the
5771 result. */
5773 static void
5774 gen_function_def_parameters ()
5776 tree parm;
5778 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
5780 for (parm = DECL_ARGUMENTS (cfun->decl); parm;
5781 parm = DECL_CHAIN (parm))
5783 struct hsa_symbol **slot;
5785 hsa_symbol *arg
5786 = new hsa_symbol (BRIG_TYPE_NONE, hsa_cfun->m_kern_p
5787 ? BRIG_SEGMENT_KERNARG : BRIG_SEGMENT_ARG,
5788 BRIG_LINKAGE_FUNCTION);
5789 arg->fillup_for_decl (parm);
5791 hsa_cfun->m_input_args.safe_push (arg);
5793 if (hsa_seen_error ())
5794 return;
5796 arg->m_name = hsa_get_declaration_name (parm);
5798 /* Copy all input arguments and create corresponding private symbols
5799 for them. */
5800 hsa_symbol *private_arg;
5801 hsa_op_address *parm_addr = new hsa_op_address (arg);
5803 if (TREE_ADDRESSABLE (parm)
5804 || (!is_gimple_reg (parm) && !TREE_READONLY (parm)))
5806 private_arg = hsa_cfun->create_hsa_temporary (arg->m_type);
5807 private_arg->fillup_for_decl (parm);
5809 BrigAlignment8_t align = MIN (arg->m_align, private_arg->m_align);
5811 hsa_op_address *private_arg_addr = new hsa_op_address (private_arg);
5812 gen_hsa_memory_copy (prologue, private_arg_addr, parm_addr,
5813 arg->total_byte_size (), align);
5815 else
5816 private_arg = arg;
5818 slot = hsa_cfun->m_local_symbols->find_slot (private_arg, INSERT);
5819 gcc_assert (!*slot);
5820 *slot = private_arg;
5822 if (is_gimple_reg (parm))
5824 tree ddef = ssa_default_def (cfun, parm);
5825 if (ddef && !has_zero_uses (ddef))
5827 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (ddef),
5828 false);
5829 BrigType16_t mtype = mem_type_for_type (t);
5830 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (ddef);
5831 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mtype,
5832 dest, parm_addr);
5833 gcc_assert (!parm_addr->m_reg);
5834 prologue->append_insn (mem);
5839 if (!VOID_TYPE_P (TREE_TYPE (TREE_TYPE (cfun->decl))))
5841 struct hsa_symbol **slot;
5843 hsa_cfun->m_output_arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
5844 BRIG_LINKAGE_FUNCTION);
5845 hsa_cfun->m_output_arg->fillup_for_decl (DECL_RESULT (cfun->decl));
5847 if (hsa_seen_error ())
5848 return;
5850 hsa_cfun->m_output_arg->m_name = "res";
5851 slot = hsa_cfun->m_local_symbols->find_slot (hsa_cfun->m_output_arg,
5852 INSERT);
5853 gcc_assert (!*slot);
5854 *slot = hsa_cfun->m_output_arg;
5858 /* Generate function representation that corresponds to
5859 a function declaration. */
5861 hsa_function_representation *
5862 hsa_generate_function_declaration (tree decl)
5864 hsa_function_representation *fun
5865 = new hsa_function_representation (decl, false, 0);
5867 fun->m_declaration_p = true;
5868 fun->m_name = get_brig_function_name (decl);
5869 gen_function_decl_parameters (fun, decl);
5871 return fun;
5875 /* Generate function representation that corresponds to
5876 an internal FN. */
5878 hsa_function_representation *
5879 hsa_generate_internal_fn_decl (hsa_internal_fn *fn)
5881 hsa_function_representation *fun = new hsa_function_representation (fn);
5883 fun->m_name = fn->name ();
5885 for (unsigned i = 0; i < fn->get_arity (); i++)
5887 hsa_symbol *arg
5888 = new hsa_symbol (fn->get_argument_type (i), BRIG_SEGMENT_ARG,
5889 BRIG_LINKAGE_NONE);
5890 arg->m_name_number = i;
5891 fun->m_input_args.safe_push (arg);
5894 fun->m_output_arg = new hsa_symbol (fn->get_argument_type (-1),
5895 BRIG_SEGMENT_ARG, BRIG_LINKAGE_NONE);
5896 fun->m_output_arg->m_name = "res";
5898 return fun;
5901 /* Return true if switch statement S can be transformed
5902 to a SBR instruction in HSAIL. */
5904 static bool
5905 transformable_switch_to_sbr_p (gswitch *s)
5907 /* Identify if a switch statement can be transformed to
5908 SBR instruction, like:
5910 sbr_u32 $s1 [@label1, @label2, @label3];
5913 tree size = get_switch_size (s);
5914 if (!tree_fits_uhwi_p (size))
5915 return false;
5917 if (tree_to_uhwi (size) > HSA_MAXIMUM_SBR_LABELS)
5918 return false;
5920 return true;
5923 /* Structure hold connection between PHI nodes and immediate
5924 values hold by there nodes. */
5926 struct phi_definition
5928 phi_definition (unsigned phi_i, unsigned label_i, tree imm):
5929 phi_index (phi_i), label_index (label_i), phi_value (imm)
5932 unsigned phi_index;
5933 unsigned label_index;
5934 tree phi_value;
5937 /* Sum slice of a vector V, starting from index START and ending
5938 at the index END - 1. */
5940 template <typename T>
5941 static
5942 T sum_slice (const auto_vec <T> &v, unsigned start, unsigned end)
5944 T s = 0;
5946 for (unsigned i = start; i < end; i++)
5947 s += v[i];
5949 return s;
5952 /* Function transforms GIMPLE SWITCH statements to a series of IF statements.
5953 Let's assume following example:
5956 switch (index)
5957 case C1:
5958 L1: hard_work_1 ();
5959 break;
5960 case C2..C3:
5961 L2: hard_work_2 ();
5962 break;
5963 default:
5964 LD: hard_work_3 ();
5965 break;
5967 The transformation encompasses following steps:
5968 1) all immediate values used by edges coming from the switch basic block
5969 are saved
5970 2) all these edges are removed
5971 3) the switch statement (in L0) is replaced by:
5972 if (index == C1)
5973 goto L1;
5974 else
5975 goto L1';
5977 4) newly created basic block Lx' is used for generation of
5978 a next condition
5979 5) else branch of the last condition goes to LD
5980 6) fix all immediate values in PHI nodes that were propagated though
5981 edges that were removed in step 2
5983 Note: if a case is made by a range C1..C2, then process
5984 following transformation:
5986 switch_cond_op1 = C1 <= index;
5987 switch_cond_op2 = index <= C2;
5988 switch_cond_and = switch_cond_op1 & switch_cond_op2;
5989 if (switch_cond_and != 0)
5990 goto Lx;
5991 else
5992 goto Ly;
5996 static bool
5997 convert_switch_statements (void)
5999 function *func = DECL_STRUCT_FUNCTION (current_function_decl);
6000 basic_block bb;
6002 bool modified_cfg = false;
6004 FOR_EACH_BB_FN (bb, func)
6006 gimple_stmt_iterator gsi = gsi_last_bb (bb);
6007 if (gsi_end_p (gsi))
6008 continue;
6010 gimple *stmt = gsi_stmt (gsi);
6012 if (gimple_code (stmt) == GIMPLE_SWITCH)
6014 gswitch *s = as_a <gswitch *> (stmt);
6016 /* If the switch can utilize SBR insn, skip the statement. */
6017 if (transformable_switch_to_sbr_p (s))
6018 continue;
6020 modified_cfg = true;
6022 unsigned labels = gimple_switch_num_labels (s);
6023 tree index = gimple_switch_index (s);
6024 tree index_type = TREE_TYPE (index);
6025 tree default_label = gimple_switch_default_label (s);
6026 basic_block default_label_bb
6027 = label_to_block_fn (func, CASE_LABEL (default_label));
6028 basic_block cur_bb = bb;
6030 auto_vec <edge> new_edges;
6031 auto_vec <phi_definition *> phi_todo_list;
6032 auto_vec <gcov_type> edge_counts;
6033 auto_vec <int> edge_probabilities;
6035 /* Investigate all labels that and PHI nodes in these edges which
6036 should be fixed after we add new collection of edges. */
6037 for (unsigned i = 0; i < labels; i++)
6039 tree label = gimple_switch_label (s, i);
6040 basic_block label_bb = label_to_block_fn (func, CASE_LABEL (label));
6041 edge e = find_edge (bb, label_bb);
6042 edge_counts.safe_push (e->count);
6043 edge_probabilities.safe_push (e->probability);
6044 gphi_iterator phi_gsi;
6046 /* Save PHI definitions that will be destroyed because of an edge
6047 is going to be removed. */
6048 unsigned phi_index = 0;
6049 for (phi_gsi = gsi_start_phis (e->dest);
6050 !gsi_end_p (phi_gsi); gsi_next (&phi_gsi))
6052 gphi *phi = phi_gsi.phi ();
6053 for (unsigned j = 0; j < gimple_phi_num_args (phi); j++)
6055 if (gimple_phi_arg_edge (phi, j) == e)
6057 tree imm = gimple_phi_arg_def (phi, j);
6058 phi_definition *p = new phi_definition (phi_index, i,
6059 imm);
6060 phi_todo_list.safe_push (p);
6061 break;
6064 phi_index++;
6068 /* Remove all edges for the current basic block. */
6069 for (int i = EDGE_COUNT (bb->succs) - 1; i >= 0; i--)
6071 edge e = EDGE_SUCC (bb, i);
6072 remove_edge (e);
6075 /* Iterate all non-default labels. */
6076 for (unsigned i = 1; i < labels; i++)
6078 tree label = gimple_switch_label (s, i);
6079 tree low = CASE_LOW (label);
6080 tree high = CASE_HIGH (label);
6082 if (!useless_type_conversion_p (TREE_TYPE (low), index_type))
6083 low = fold_convert (index_type, low);
6085 gimple_stmt_iterator cond_gsi = gsi_last_bb (cur_bb);
6086 gimple *c = NULL;
6087 if (high)
6089 tree tmp1 = make_temp_ssa_name (boolean_type_node, NULL,
6090 "switch_cond_op1");
6092 gimple *assign1 = gimple_build_assign (tmp1, LE_EXPR, low,
6093 index);
6095 tree tmp2 = make_temp_ssa_name (boolean_type_node, NULL,
6096 "switch_cond_op2");
6098 if (!useless_type_conversion_p (TREE_TYPE (high), index_type))
6099 high = fold_convert (index_type, high);
6100 gimple *assign2 = gimple_build_assign (tmp2, LE_EXPR, index,
6101 high);
6103 tree tmp3 = make_temp_ssa_name (boolean_type_node, NULL,
6104 "switch_cond_and");
6105 gimple *assign3 = gimple_build_assign (tmp3, BIT_AND_EXPR, tmp1,
6106 tmp2);
6108 gsi_insert_before (&cond_gsi, assign1, GSI_SAME_STMT);
6109 gsi_insert_before (&cond_gsi, assign2, GSI_SAME_STMT);
6110 gsi_insert_before (&cond_gsi, assign3, GSI_SAME_STMT);
6112 tree b = constant_boolean_node (false, boolean_type_node);
6113 c = gimple_build_cond (NE_EXPR, tmp3, b, NULL, NULL);
6115 else
6116 c = gimple_build_cond (EQ_EXPR, index, low, NULL, NULL);
6118 gimple_set_location (c, gimple_location (stmt));
6120 gsi_insert_before (&cond_gsi, c, GSI_SAME_STMT);
6122 basic_block label_bb
6123 = label_to_block_fn (func, CASE_LABEL (label));
6124 edge new_edge = make_edge (cur_bb, label_bb, EDGE_TRUE_VALUE);
6125 int prob_sum = sum_slice <int> (edge_probabilities, i, labels) +
6126 edge_probabilities[0];
6128 if (prob_sum)
6129 new_edge->probability
6130 = RDIV (REG_BR_PROB_BASE * edge_probabilities[i], prob_sum);
6132 new_edge->count = edge_counts[i];
6133 new_edges.safe_push (new_edge);
6135 if (i < labels - 1)
6137 /* Prepare another basic block that will contain
6138 next condition. */
6139 basic_block next_bb = create_empty_bb (cur_bb);
6140 if (current_loops)
6142 add_bb_to_loop (next_bb, cur_bb->loop_father);
6143 loops_state_set (LOOPS_NEED_FIXUP);
6146 edge next_edge = make_edge (cur_bb, next_bb, EDGE_FALSE_VALUE);
6147 next_edge->probability
6148 = inverse_probability (new_edge->probability);
6149 next_edge->count = edge_counts[0]
6150 + sum_slice <gcov_type> (edge_counts, i, labels);
6151 next_bb->frequency = EDGE_FREQUENCY (next_edge);
6152 cur_bb = next_bb;
6154 else /* Link last IF statement and default label
6155 of the switch. */
6157 edge e = make_edge (cur_bb, default_label_bb, EDGE_FALSE_VALUE);
6158 e->probability = inverse_probability (new_edge->probability);
6159 e->count = edge_counts[0];
6160 new_edges.safe_insert (0, e);
6164 /* Restore original PHI immediate value. */
6165 for (unsigned i = 0; i < phi_todo_list.length (); i++)
6167 phi_definition *phi_def = phi_todo_list[i];
6168 edge new_edge = new_edges[phi_def->label_index];
6170 gphi_iterator it = gsi_start_phis (new_edge->dest);
6171 for (unsigned i = 0; i < phi_def->phi_index; i++)
6172 gsi_next (&it);
6174 gphi *phi = it.phi ();
6175 add_phi_arg (phi, phi_def->phi_value, new_edge, UNKNOWN_LOCATION);
6176 delete phi_def;
6179 /* Remove the original GIMPLE switch statement. */
6180 gsi_remove (&gsi, true);
6184 if (dump_file)
6185 dump_function_to_file (current_function_decl, dump_file, TDF_DETAILS);
6187 return modified_cfg;
6190 /* Expand builtins that can't be handled by HSA back-end. */
6192 static void
6193 expand_builtins ()
6195 function *func = DECL_STRUCT_FUNCTION (current_function_decl);
6196 basic_block bb;
6198 FOR_EACH_BB_FN (bb, func)
6200 for (gimple_stmt_iterator gsi = gsi_start_bb (bb); !gsi_end_p (gsi);
6201 gsi_next (&gsi))
6203 gimple *stmt = gsi_stmt (gsi);
6205 if (gimple_code (stmt) != GIMPLE_CALL)
6206 continue;
6208 gcall *call = as_a <gcall *> (stmt);
6210 if (!gimple_call_builtin_p (call, BUILT_IN_NORMAL))
6211 continue;
6213 tree fndecl = gimple_call_fndecl (stmt);
6214 enum built_in_function fn = DECL_FUNCTION_CODE (fndecl);
6215 switch (fn)
6217 case BUILT_IN_CEXPF:
6218 case BUILT_IN_CEXPIF:
6219 case BUILT_IN_CEXPI:
6221 /* Similar to builtins.c (expand_builtin_cexpi), the builtin
6222 can be transformed to: cexp(I * z) = ccos(z) + I * csin(z). */
6223 tree lhs = gimple_call_lhs (stmt);
6224 tree rhs = gimple_call_arg (stmt, 0);
6225 tree rhs_type = TREE_TYPE (rhs);
6226 bool float_type_p = rhs_type == float_type_node;
6227 tree real_part = make_temp_ssa_name (rhs_type, NULL,
6228 "cexp_real_part");
6229 tree imag_part = make_temp_ssa_name (rhs_type, NULL,
6230 "cexp_imag_part");
6232 tree cos_fndecl
6233 = mathfn_built_in (rhs_type, fn == float_type_p
6234 ? BUILT_IN_COSF : BUILT_IN_COS);
6235 gcall *cos = gimple_build_call (cos_fndecl, 1, rhs);
6236 gimple_call_set_lhs (cos, real_part);
6237 gsi_insert_before (&gsi, cos, GSI_SAME_STMT);
6239 tree sin_fndecl
6240 = mathfn_built_in (rhs_type, fn == float_type_p
6241 ? BUILT_IN_SINF : BUILT_IN_SIN);
6242 gcall *sin = gimple_build_call (sin_fndecl, 1, rhs);
6243 gimple_call_set_lhs (sin, imag_part);
6244 gsi_insert_before (&gsi, sin, GSI_SAME_STMT);
6247 gassign *assign = gimple_build_assign (lhs, COMPLEX_EXPR,
6248 real_part, imag_part);
6249 gsi_insert_before (&gsi, assign, GSI_SAME_STMT);
6250 gsi_remove (&gsi, true);
6252 break;
6254 default:
6255 break;
6261 /* Emit HSA module variables that are global for the entire module. */
6263 static void
6264 emit_hsa_module_variables (void)
6266 hsa_num_threads = new hsa_symbol (BRIG_TYPE_U32, BRIG_SEGMENT_PRIVATE,
6267 BRIG_LINKAGE_MODULE, true);
6269 hsa_num_threads->m_name = "hsa_num_threads";
6271 hsa_brig_emit_omp_symbols ();
6274 /* Generate HSAIL representation of the current function and write into a
6275 special section of the output file. If KERNEL is set, the function will be
6276 considered an HSA kernel callable from the host, otherwise it will be
6277 compiled as an HSA function callable from other HSA code. */
6279 static void
6280 generate_hsa (bool kernel)
6282 hsa_init_data_for_cfun ();
6284 if (hsa_num_threads == NULL)
6285 emit_hsa_module_variables ();
6287 bool modified_cfg = convert_switch_statements ();
6288 /* Initialize hsa_cfun. */
6289 hsa_cfun = new hsa_function_representation (cfun->decl, kernel,
6290 SSANAMES (cfun)->length (),
6291 modified_cfg);
6292 hsa_cfun->init_extra_bbs ();
6294 if (flag_tm)
6296 HSA_SORRY_AT (UNKNOWN_LOCATION,
6297 "support for HSA does not implement transactional memory");
6298 goto fail;
6301 verify_function_arguments (cfun->decl);
6302 if (hsa_seen_error ())
6303 goto fail;
6305 hsa_cfun->m_name = get_brig_function_name (cfun->decl);
6307 gen_function_def_parameters ();
6308 if (hsa_seen_error ())
6309 goto fail;
6311 init_prologue ();
6313 gen_body_from_gimple ();
6314 if (hsa_seen_error ())
6315 goto fail;
6317 if (hsa_cfun->m_kernel_dispatch_count)
6318 init_hsa_num_threads ();
6320 if (hsa_cfun->m_kern_p)
6322 hsa_function_summary *s
6323 = hsa_summaries->get (cgraph_node::get (hsa_cfun->m_decl));
6324 hsa_add_kern_decl_mapping (current_function_decl, hsa_cfun->m_name,
6325 hsa_cfun->m_maximum_omp_data_size,
6326 s->m_gridified_kernel_p);
6329 if (flag_checking)
6331 for (unsigned i = 0; i < hsa_cfun->m_ssa_map.length (); i++)
6332 if (hsa_cfun->m_ssa_map[i])
6333 hsa_cfun->m_ssa_map[i]->verify_ssa ();
6335 basic_block bb;
6336 FOR_EACH_BB_FN (bb, cfun)
6338 hsa_bb *hbb = hsa_bb_for_bb (bb);
6340 for (hsa_insn_basic *insn = hbb->m_first_insn; insn;
6341 insn = insn->m_next)
6342 insn->verify ();
6346 hsa_regalloc ();
6347 hsa_brig_emit_function ();
6349 fail:
6350 hsa_deinit_data_for_cfun ();
6353 namespace {
6355 const pass_data pass_data_gen_hsail =
6357 GIMPLE_PASS,
6358 "hsagen", /* name */
6359 OPTGROUP_NONE, /* optinfo_flags */
6360 TV_NONE, /* tv_id */
6361 PROP_cfg | PROP_ssa, /* properties_required */
6362 0, /* properties_provided */
6363 0, /* properties_destroyed */
6364 0, /* todo_flags_start */
6365 0 /* todo_flags_finish */
6368 class pass_gen_hsail : public gimple_opt_pass
6370 public:
6371 pass_gen_hsail (gcc::context *ctxt)
6372 : gimple_opt_pass(pass_data_gen_hsail, ctxt)
6375 /* opt_pass methods: */
6376 bool gate (function *);
6377 unsigned int execute (function *);
6379 }; // class pass_gen_hsail
6381 /* Determine whether or not to run generation of HSAIL. */
6383 bool
6384 pass_gen_hsail::gate (function *f)
6386 return hsa_gen_requested_p ()
6387 && hsa_gpu_implementation_p (f->decl);
6390 unsigned int
6391 pass_gen_hsail::execute (function *)
6393 hsa_function_summary *s
6394 = hsa_summaries->get (cgraph_node::get_create (current_function_decl));
6396 expand_builtins ();
6397 generate_hsa (s->m_kind == HSA_KERNEL);
6398 TREE_ASM_WRITTEN (current_function_decl) = 1;
6399 return TODO_discard_function;
6402 } // anon namespace
6404 /* Create the instance of hsa gen pass. */
6406 gimple_opt_pass *
6407 make_pass_gen_hsail (gcc::context *ctxt)
6409 return new pass_gen_hsail (ctxt);