[PR c++/84968] reject stmt-exprs in noexcept constexprs
[official-gcc.git] / gcc / hsa-gen.c
blob7974fffe360f77eb4e35a8c4a588e4b2865d5de3
1 /* A pass for lowering gimple to HSAIL
2 Copyright (C) 2013-2018 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 "memmodel.h"
26 #include "tm.h"
27 #include "is-a.h"
28 #include "hash-table.h"
29 #include "vec.h"
30 #include "tree.h"
31 #include "tree-pass.h"
32 #include "function.h"
33 #include "basic-block.h"
34 #include "cfg.h"
35 #include "fold-const.h"
36 #include "gimple.h"
37 #include "gimple-iterator.h"
38 #include "bitmap.h"
39 #include "dumpfile.h"
40 #include "gimple-pretty-print.h"
41 #include "diagnostic-core.h"
42 #include "gimple-ssa.h"
43 #include "tree-phinodes.h"
44 #include "stringpool.h"
45 #include "tree-vrp.h"
46 #include "tree-ssanames.h"
47 #include "tree-dfa.h"
48 #include "ssa-iterators.h"
49 #include "cgraph.h"
50 #include "print-tree.h"
51 #include "symbol-summary.h"
52 #include "hsa-common.h"
53 #include "cfghooks.h"
54 #include "tree-cfg.h"
55 #include "cfgloop.h"
56 #include "cfganal.h"
57 #include "builtins.h"
58 #include "params.h"
59 #include "gomp-constants.h"
60 #include "internal-fn.h"
61 #include "builtins.h"
62 #include "stor-layout.h"
63 #include "stringpool.h"
64 #include "attribs.h"
66 /* Print a warning message and set that we have seen an error. */
68 #define HSA_SORRY_ATV(location, message, ...) \
69 do \
70 { \
71 hsa_fail_cfun (); \
72 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
73 HSA_SORRY_MSG)) \
74 inform (location, message, __VA_ARGS__); \
75 } \
76 while (false)
78 /* Same as previous, but highlight a location. */
80 #define HSA_SORRY_AT(location, message) \
81 do \
82 { \
83 hsa_fail_cfun (); \
84 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
85 HSA_SORRY_MSG)) \
86 inform (location, message); \
87 } \
88 while (false)
90 /* Default number of threads used by kernel dispatch. */
92 #define HSA_DEFAULT_NUM_THREADS 64
94 /* Following structures are defined in the final version
95 of HSA specification. */
97 /* HSA queue packet is shadow structure, originally provided by AMD. */
99 struct hsa_queue_packet
101 uint16_t header;
102 uint16_t setup;
103 uint16_t workgroup_size_x;
104 uint16_t workgroup_size_y;
105 uint16_t workgroup_size_z;
106 uint16_t reserved0;
107 uint32_t grid_size_x;
108 uint32_t grid_size_y;
109 uint32_t grid_size_z;
110 uint32_t private_segment_size;
111 uint32_t group_segment_size;
112 uint64_t kernel_object;
113 void *kernarg_address;
114 uint64_t reserved2;
115 uint64_t completion_signal;
118 /* HSA queue is shadow structure, originally provided by AMD. */
120 struct hsa_queue
122 int type;
123 uint32_t features;
124 void *base_address;
125 uint64_t doorbell_signal;
126 uint32_t size;
127 uint32_t reserved1;
128 uint64_t id;
131 static struct obstack hsa_obstack;
133 /* List of pointers to all instructions that come from an object allocator. */
134 static vec <hsa_insn_basic *> hsa_instructions;
136 /* List of pointers to all operands that come from an object allocator. */
137 static vec <hsa_op_base *> hsa_operands;
139 hsa_symbol::hsa_symbol ()
140 : m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
141 m_directive_offset (0), m_type (BRIG_TYPE_NONE),
142 m_segment (BRIG_SEGMENT_NONE), m_linkage (BRIG_LINKAGE_NONE), m_dim (0),
143 m_cst_value (NULL), m_global_scope_p (false), m_seen_error (false),
144 m_allocation (BRIG_ALLOCATION_AUTOMATIC), m_emitted_to_brig (false)
149 hsa_symbol::hsa_symbol (BrigType16_t type, BrigSegment8_t segment,
150 BrigLinkage8_t linkage, bool global_scope_p,
151 BrigAllocation allocation, BrigAlignment8_t align)
152 : m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
153 m_directive_offset (0), m_type (type), m_segment (segment),
154 m_linkage (linkage), m_dim (0), m_cst_value (NULL),
155 m_global_scope_p (global_scope_p), m_seen_error (false),
156 m_allocation (allocation), m_emitted_to_brig (false), m_align (align)
160 unsigned HOST_WIDE_INT
161 hsa_symbol::total_byte_size ()
163 unsigned HOST_WIDE_INT s
164 = hsa_type_bit_size (~BRIG_TYPE_ARRAY_MASK & m_type);
165 gcc_assert (s % BITS_PER_UNIT == 0);
166 s /= BITS_PER_UNIT;
168 if (m_dim)
169 s *= m_dim;
171 return s;
174 /* Forward declaration. */
176 static BrigType16_t
177 hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p,
178 bool min32int);
180 void
181 hsa_symbol::fillup_for_decl (tree decl)
183 m_decl = decl;
184 m_type = hsa_type_for_tree_type (TREE_TYPE (decl), &m_dim, false);
185 if (hsa_seen_error ())
187 m_seen_error = true;
188 return;
191 m_align = MAX (m_align, hsa_natural_alignment (m_type));
194 /* Constructor of class representing global HSA function/kernel information and
195 state. FNDECL is function declaration, KERNEL_P is true if the function
196 is going to become a HSA kernel. If the function has body, SSA_NAMES_COUNT
197 should be set to number of SSA names used in the function.
198 MODIFIED_CFG is set to true in case we modified control-flow graph
199 of the function. */
201 hsa_function_representation::hsa_function_representation
202 (tree fdecl, bool kernel_p, unsigned ssa_names_count, bool modified_cfg)
203 : m_name (NULL),
204 m_reg_count (0), m_input_args (vNULL),
205 m_output_arg (NULL), m_spill_symbols (vNULL), m_global_symbols (vNULL),
206 m_private_variables (vNULL), m_called_functions (vNULL),
207 m_called_internal_fns (vNULL), m_hbb_count (0),
208 m_in_ssa (true), m_kern_p (kernel_p), m_declaration_p (false),
209 m_decl (fdecl), m_internal_fn (NULL), m_shadow_reg (NULL),
210 m_kernel_dispatch_count (0), m_maximum_omp_data_size (0),
211 m_seen_error (false), m_temp_symbol_count (0), m_ssa_map (),
212 m_modified_cfg (modified_cfg)
214 int sym_init_len = (vec_safe_length (cfun->local_decls) / 2) + 1;
215 m_local_symbols = new hash_table <hsa_noop_symbol_hasher> (sym_init_len);
216 m_ssa_map.safe_grow_cleared (ssa_names_count);
219 /* Constructor of class representing HSA function information that
220 is derived for an internal function. */
221 hsa_function_representation::hsa_function_representation (hsa_internal_fn *fn)
222 : m_reg_count (0), m_input_args (vNULL),
223 m_output_arg (NULL), m_local_symbols (NULL),
224 m_spill_symbols (vNULL), m_global_symbols (vNULL),
225 m_private_variables (vNULL), m_called_functions (vNULL),
226 m_called_internal_fns (vNULL), m_hbb_count (0),
227 m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL),
228 m_internal_fn (fn), m_shadow_reg (NULL), m_kernel_dispatch_count (0),
229 m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
230 m_ssa_map () {}
232 /* Destructor of class holding function/kernel-wide information and state. */
234 hsa_function_representation::~hsa_function_representation ()
236 /* Kernel names are deallocated at the end of BRIG output when deallocating
237 hsa_decl_kernel_mapping. */
238 if (!m_kern_p || m_seen_error)
239 free (m_name);
241 for (unsigned i = 0; i < m_input_args.length (); i++)
242 delete m_input_args[i];
243 m_input_args.release ();
245 delete m_output_arg;
246 delete m_local_symbols;
248 for (unsigned i = 0; i < m_spill_symbols.length (); i++)
249 delete m_spill_symbols[i];
250 m_spill_symbols.release ();
252 hsa_symbol *sym;
253 for (unsigned i = 0; i < m_global_symbols.iterate (i, &sym); i++)
254 if (sym->m_linkage != BRIG_ALLOCATION_PROGRAM)
255 delete sym;
256 m_global_symbols.release ();
258 for (unsigned i = 0; i < m_private_variables.length (); i++)
259 delete m_private_variables[i];
260 m_private_variables.release ();
261 m_called_functions.release ();
262 m_ssa_map.release ();
264 for (unsigned i = 0; i < m_called_internal_fns.length (); i++)
265 delete m_called_internal_fns[i];
268 hsa_op_reg *
269 hsa_function_representation::get_shadow_reg ()
271 /* If we compile a function with kernel dispatch and does not set
272 an optimization level, the function won't be inlined and
273 we return NULL. */
274 if (!m_kern_p)
275 return NULL;
277 if (m_shadow_reg)
278 return m_shadow_reg;
280 /* Append the shadow argument. */
281 hsa_symbol *shadow = new hsa_symbol (BRIG_TYPE_U64, BRIG_SEGMENT_KERNARG,
282 BRIG_LINKAGE_FUNCTION);
283 m_input_args.safe_push (shadow);
284 shadow->m_name = "hsa_runtime_shadow";
286 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_U64);
287 hsa_op_address *addr = new hsa_op_address (shadow);
289 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, r, addr);
290 hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun))->append_insn (mem);
291 m_shadow_reg = r;
293 return r;
296 bool hsa_function_representation::has_shadow_reg_p ()
298 return m_shadow_reg != NULL;
301 void
302 hsa_function_representation::init_extra_bbs ()
304 hsa_init_new_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
305 hsa_init_new_bb (EXIT_BLOCK_PTR_FOR_FN (cfun));
308 void
309 hsa_function_representation::update_dominance ()
311 if (m_modified_cfg)
313 free_dominance_info (CDI_DOMINATORS);
314 calculate_dominance_info (CDI_DOMINATORS);
318 hsa_symbol *
319 hsa_function_representation::create_hsa_temporary (BrigType16_t type)
321 hsa_symbol *s = new hsa_symbol (type, BRIG_SEGMENT_PRIVATE,
322 BRIG_LINKAGE_FUNCTION);
323 s->m_name_number = m_temp_symbol_count++;
325 hsa_cfun->m_private_variables.safe_push (s);
326 return s;
329 BrigLinkage8_t
330 hsa_function_representation::get_linkage ()
332 if (m_internal_fn)
333 return BRIG_LINKAGE_PROGRAM;
335 return m_kern_p || TREE_PUBLIC (m_decl) ?
336 BRIG_LINKAGE_PROGRAM : BRIG_LINKAGE_MODULE;
339 /* Hash map of simple OMP builtins. */
340 static hash_map <nofree_string_hash, omp_simple_builtin> *omp_simple_builtins
341 = NULL;
343 /* Warning messages for OMP builtins. */
345 #define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP " \
346 "lock routines"
347 #define HSA_WARN_TIMING_ROUTINE "support for HSA does not implement OpenMP " \
348 "timing routines"
349 #define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \
350 "undefined semantics within target regions, support for HSA ignores them"
351 #define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP " \
352 "affinity feateres"
354 /* Initialize hash map with simple OMP builtins. */
356 static void
357 hsa_init_simple_builtins ()
359 if (omp_simple_builtins != NULL)
360 return;
362 omp_simple_builtins
363 = new hash_map <nofree_string_hash, omp_simple_builtin> ();
365 omp_simple_builtin omp_builtins[] =
367 omp_simple_builtin ("omp_get_initial_device", NULL, false,
368 new hsa_op_immed (GOMP_DEVICE_HOST,
369 (BrigType16_t) BRIG_TYPE_S32)),
370 omp_simple_builtin ("omp_is_initial_device", NULL, false,
371 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
372 omp_simple_builtin ("omp_get_dynamic", NULL, false,
373 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
374 omp_simple_builtin ("omp_set_dynamic", NULL, false, NULL),
375 omp_simple_builtin ("omp_init_lock", HSA_WARN_LOCK_ROUTINE, true),
376 omp_simple_builtin ("omp_init_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
377 true),
378 omp_simple_builtin ("omp_init_nest_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
379 true),
380 omp_simple_builtin ("omp_destroy_lock", HSA_WARN_LOCK_ROUTINE, true),
381 omp_simple_builtin ("omp_set_lock", HSA_WARN_LOCK_ROUTINE, true),
382 omp_simple_builtin ("omp_unset_lock", HSA_WARN_LOCK_ROUTINE, true),
383 omp_simple_builtin ("omp_test_lock", HSA_WARN_LOCK_ROUTINE, true),
384 omp_simple_builtin ("omp_get_wtime", HSA_WARN_TIMING_ROUTINE, true),
385 omp_simple_builtin ("omp_get_wtick", HSA_WARN_TIMING_ROUTINE, true),
386 omp_simple_builtin ("omp_target_alloc", HSA_WARN_MEMORY_ROUTINE, false,
387 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_U64)),
388 omp_simple_builtin ("omp_target_free", HSA_WARN_MEMORY_ROUTINE, false),
389 omp_simple_builtin ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE,
390 false,
391 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
392 omp_simple_builtin ("omp_target_memcpy", HSA_WARN_MEMORY_ROUTINE, false,
393 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
394 omp_simple_builtin ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE,
395 false,
396 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
397 omp_simple_builtin ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE,
398 false,
399 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
400 omp_simple_builtin ("omp_target_disassociate_ptr",
401 HSA_WARN_MEMORY_ROUTINE,
402 false,
403 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
404 omp_simple_builtin ("omp_set_max_active_levels",
405 "Support for HSA only allows only one active level, "
406 "call to omp_set_max_active_levels will be ignored "
407 "in the generated HSAIL",
408 false, NULL),
409 omp_simple_builtin ("omp_get_max_active_levels", NULL, false,
410 new hsa_op_immed (1, (BrigType16_t) BRIG_TYPE_S32)),
411 omp_simple_builtin ("omp_in_final", NULL, false,
412 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
413 omp_simple_builtin ("omp_get_proc_bind", HSA_WARN_AFFINITY, false,
414 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
415 omp_simple_builtin ("omp_get_num_places", HSA_WARN_AFFINITY, false,
416 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
417 omp_simple_builtin ("omp_get_place_num_procs", HSA_WARN_AFFINITY, false,
418 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
419 omp_simple_builtin ("omp_get_place_proc_ids", HSA_WARN_AFFINITY, false,
420 NULL),
421 omp_simple_builtin ("omp_get_place_num", HSA_WARN_AFFINITY, false,
422 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
423 omp_simple_builtin ("omp_get_partition_num_places", HSA_WARN_AFFINITY,
424 false,
425 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
426 omp_simple_builtin ("omp_get_partition_place_nums", HSA_WARN_AFFINITY,
427 false, NULL),
428 omp_simple_builtin ("omp_set_default_device",
429 "omp_set_default_device has undefined semantics "
430 "within target regions, support for HSA ignores it",
431 false, NULL),
432 omp_simple_builtin ("omp_get_default_device",
433 "omp_get_default_device has undefined semantics "
434 "within target regions, support for HSA ignores it",
435 false,
436 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
437 omp_simple_builtin ("omp_get_num_devices",
438 "omp_get_num_devices has undefined semantics "
439 "within target regions, support for HSA ignores it",
440 false,
441 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
442 omp_simple_builtin ("omp_get_num_procs", NULL, true, NULL),
443 omp_simple_builtin ("omp_get_cancellation", NULL, true, NULL),
444 omp_simple_builtin ("omp_set_nested", NULL, true, NULL),
445 omp_simple_builtin ("omp_get_nested", NULL, true, NULL),
446 omp_simple_builtin ("omp_set_schedule", NULL, true, NULL),
447 omp_simple_builtin ("omp_get_schedule", NULL, true, NULL),
448 omp_simple_builtin ("omp_get_thread_limit", NULL, true, NULL),
449 omp_simple_builtin ("omp_get_team_size", NULL, true, NULL),
450 omp_simple_builtin ("omp_get_ancestor_thread_num", NULL, true, NULL),
451 omp_simple_builtin ("omp_get_max_task_priority", NULL, true, NULL)
454 unsigned count = sizeof (omp_builtins) / sizeof (omp_simple_builtin);
456 for (unsigned i = 0; i < count; i++)
457 omp_simple_builtins->put (omp_builtins[i].m_name, omp_builtins[i]);
460 /* Allocate HSA structures that we need only while generating with this. */
462 static void
463 hsa_init_data_for_cfun ()
465 hsa_init_compilation_unit_data ();
466 gcc_obstack_init (&hsa_obstack);
469 /* Deinitialize HSA subsystem and free all allocated memory. */
471 static void
472 hsa_deinit_data_for_cfun (void)
474 basic_block bb;
476 FOR_ALL_BB_FN (bb, cfun)
477 if (bb->aux)
479 hsa_bb *hbb = hsa_bb_for_bb (bb);
480 hbb->~hsa_bb ();
481 bb->aux = NULL;
484 for (unsigned int i = 0; i < hsa_operands.length (); i++)
485 hsa_destroy_operand (hsa_operands[i]);
487 hsa_operands.release ();
489 for (unsigned i = 0; i < hsa_instructions.length (); i++)
490 hsa_destroy_insn (hsa_instructions[i]);
492 hsa_instructions.release ();
494 if (omp_simple_builtins != NULL)
496 delete omp_simple_builtins;
497 omp_simple_builtins = NULL;
500 obstack_free (&hsa_obstack, NULL);
501 delete hsa_cfun;
504 /* Return the type which holds addresses in the given SEGMENT. */
506 static BrigType16_t
507 hsa_get_segment_addr_type (BrigSegment8_t segment)
509 switch (segment)
511 case BRIG_SEGMENT_NONE:
512 gcc_unreachable ();
514 case BRIG_SEGMENT_FLAT:
515 case BRIG_SEGMENT_GLOBAL:
516 case BRIG_SEGMENT_READONLY:
517 case BRIG_SEGMENT_KERNARG:
518 return hsa_machine_large_p () ? BRIG_TYPE_U64 : BRIG_TYPE_U32;
520 case BRIG_SEGMENT_GROUP:
521 case BRIG_SEGMENT_PRIVATE:
522 case BRIG_SEGMENT_SPILL:
523 case BRIG_SEGMENT_ARG:
524 return BRIG_TYPE_U32;
526 gcc_unreachable ();
529 /* Return integer brig type according to provided SIZE in bytes. If SIGN
530 is set to true, return signed integer type. */
532 static BrigType16_t
533 get_integer_type_by_bytes (unsigned size, bool sign)
535 if (sign)
536 switch (size)
538 case 1:
539 return BRIG_TYPE_S8;
540 case 2:
541 return BRIG_TYPE_S16;
542 case 4:
543 return BRIG_TYPE_S32;
544 case 8:
545 return BRIG_TYPE_S64;
546 default:
547 break;
549 else
550 switch (size)
552 case 1:
553 return BRIG_TYPE_U8;
554 case 2:
555 return BRIG_TYPE_U16;
556 case 4:
557 return BRIG_TYPE_U32;
558 case 8:
559 return BRIG_TYPE_U64;
560 default:
561 break;
564 return 0;
567 /* If T points to an integral type smaller than 32 bits, change it to a 32bit
568 equivalent and return the result. Otherwise just return the result. */
570 static BrigType16_t
571 hsa_extend_inttype_to_32bit (BrigType16_t t)
573 if (t == BRIG_TYPE_U8 || t == BRIG_TYPE_U16)
574 return BRIG_TYPE_U32;
575 else if (t == BRIG_TYPE_S8 || t == BRIG_TYPE_S16)
576 return BRIG_TYPE_S32;
577 return t;
580 /* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
581 are assumed to use flat addressing. If min32int is true, always expand
582 integer types to one that has at least 32 bits. */
584 static BrigType16_t
585 hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
587 HOST_WIDE_INT bsize;
588 const_tree base;
589 BrigType16_t res = BRIG_TYPE_NONE;
591 gcc_checking_assert (TYPE_P (type));
592 gcc_checking_assert (!AGGREGATE_TYPE_P (type));
593 if (POINTER_TYPE_P (type))
594 return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
596 if (TREE_CODE (type) == VECTOR_TYPE)
597 base = TREE_TYPE (type);
598 else if (TREE_CODE (type) == COMPLEX_TYPE)
600 base = TREE_TYPE (type);
601 min32int = true;
603 else
604 base = type;
606 if (!tree_fits_uhwi_p (TYPE_SIZE (base)))
608 HSA_SORRY_ATV (EXPR_LOCATION (type),
609 "support for HSA does not implement huge or "
610 "variable-sized type %qT", type);
611 return res;
614 bsize = tree_to_uhwi (TYPE_SIZE (base));
615 unsigned byte_size = bsize / BITS_PER_UNIT;
616 if (INTEGRAL_TYPE_P (base))
617 res = get_integer_type_by_bytes (byte_size, !TYPE_UNSIGNED (base));
618 else if (SCALAR_FLOAT_TYPE_P (base))
620 switch (bsize)
622 case 16:
623 res = BRIG_TYPE_F16;
624 break;
625 case 32:
626 res = BRIG_TYPE_F32;
627 break;
628 case 64:
629 res = BRIG_TYPE_F64;
630 break;
631 default:
632 break;
636 if (res == BRIG_TYPE_NONE)
638 HSA_SORRY_ATV (EXPR_LOCATION (type),
639 "support for HSA does not implement type %qT", type);
640 return res;
643 if (TREE_CODE (type) == VECTOR_TYPE)
645 HOST_WIDE_INT tsize = tree_to_uhwi (TYPE_SIZE (type));
647 if (bsize == tsize)
649 HSA_SORRY_ATV (EXPR_LOCATION (type),
650 "support for HSA does not implement a vector type "
651 "where a type and unit size are equal: %qT", type);
652 return res;
655 switch (tsize)
657 case 32:
658 res |= BRIG_TYPE_PACK_32;
659 break;
660 case 64:
661 res |= BRIG_TYPE_PACK_64;
662 break;
663 case 128:
664 res |= BRIG_TYPE_PACK_128;
665 break;
666 default:
667 HSA_SORRY_ATV (EXPR_LOCATION (type),
668 "support for HSA does not implement type %qT", type);
672 if (min32int)
673 /* Registers/immediate operands can only be 32bit or more except for
674 f16. */
675 res = hsa_extend_inttype_to_32bit (res);
677 if (TREE_CODE (type) == COMPLEX_TYPE)
679 unsigned bsize = 2 * hsa_type_bit_size (res);
680 res = hsa_bittype_for_bitsize (bsize);
683 return res;
686 /* Returns the BRIG type we need to load/store entities of TYPE. */
688 static BrigType16_t
689 mem_type_for_type (BrigType16_t type)
691 /* HSA has non-intuitive constraints on load/store types. If it's
692 a bit-type it _must_ be B128, if it's not a bit-type it must be
693 64bit max. So for loading entities of 128 bits (e.g. vectors)
694 we have to use B128, while for loading the rest we have to use the
695 input type (??? or maybe also flattened to a equally sized non-vector
696 unsigned type?). */
697 if ((type & BRIG_TYPE_PACK_MASK) == BRIG_TYPE_PACK_128)
698 return BRIG_TYPE_B128;
699 else if (hsa_btype_p (type) || hsa_type_packed_p (type))
701 unsigned bitsize = hsa_type_bit_size (type);
702 if (bitsize < 128)
703 return hsa_uint_for_bitsize (bitsize);
704 else
705 return hsa_bittype_for_bitsize (bitsize);
707 return type;
710 /* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some
711 kind of array will be generated, setting DIM appropriately. Otherwise, it
712 will be set to zero. */
714 static BrigType16_t
715 hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p = NULL,
716 bool min32int = false)
718 gcc_checking_assert (TYPE_P (type));
719 if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type)))
721 HSA_SORRY_ATV (EXPR_LOCATION (type), "support for HSA does not "
722 "implement huge or variable-sized type %qT", type);
723 return BRIG_TYPE_NONE;
726 if (RECORD_OR_UNION_TYPE_P (type))
728 if (dim_p)
729 *dim_p = tree_to_uhwi (TYPE_SIZE_UNIT (type));
730 return BRIG_TYPE_U8 | BRIG_TYPE_ARRAY;
733 if (TREE_CODE (type) == ARRAY_TYPE)
735 /* We try to be nice and use the real base-type when this is an array of
736 scalars and only resort to an array of bytes if the type is more
737 complex. */
739 unsigned HOST_WIDE_INT dim = 1;
741 while (TREE_CODE (type) == ARRAY_TYPE)
743 tree domain = TYPE_DOMAIN (type);
744 if (!TYPE_MIN_VALUE (domain)
745 || !TYPE_MAX_VALUE (domain)
746 || !tree_fits_shwi_p (TYPE_MIN_VALUE (domain))
747 || !tree_fits_shwi_p (TYPE_MAX_VALUE (domain)))
749 HSA_SORRY_ATV (EXPR_LOCATION (type),
750 "support for HSA does not implement array "
751 "%qT with unknown bounds", type);
752 return BRIG_TYPE_NONE;
754 HOST_WIDE_INT min = tree_to_shwi (TYPE_MIN_VALUE (domain));
755 HOST_WIDE_INT max = tree_to_shwi (TYPE_MAX_VALUE (domain));
756 dim = dim * (unsigned HOST_WIDE_INT) (max - min + 1);
757 type = TREE_TYPE (type);
760 BrigType16_t res;
761 if (RECORD_OR_UNION_TYPE_P (type))
763 dim = dim * tree_to_uhwi (TYPE_SIZE_UNIT (type));
764 res = BRIG_TYPE_U8;
766 else
767 res = hsa_type_for_scalar_tree_type (type, false);
769 if (dim_p)
770 *dim_p = dim;
771 return res | BRIG_TYPE_ARRAY;
774 /* Scalar case: */
775 if (dim_p)
776 *dim_p = 0;
778 return hsa_type_for_scalar_tree_type (type, min32int);
781 /* Returns true if converting from STYPE into DTYPE needs the _CVT
782 opcode. If false a normal _MOV is enough. */
784 static bool
785 hsa_needs_cvt (BrigType16_t dtype, BrigType16_t stype)
787 if (hsa_btype_p (dtype))
788 return false;
790 /* float <-> int conversions are real converts. */
791 if (hsa_type_float_p (dtype) != hsa_type_float_p (stype))
792 return true;
793 /* When both types have different size, then we need CVT as well. */
794 if (hsa_type_bit_size (dtype) != hsa_type_bit_size (stype))
795 return true;
796 return false;
799 /* Return declaration name if it exists or create one from UID if it does not.
800 If DECL is a local variable, make UID part of its name. */
802 const char *
803 hsa_get_declaration_name (tree decl)
805 if (!DECL_NAME (decl))
807 char buf[64];
808 snprintf (buf, 64, "__hsa_anon_%u", DECL_UID (decl));
809 size_t len = strlen (buf);
810 char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1);
811 memcpy (copy, buf, len + 1);
812 return copy;
815 tree name_tree;
816 if (TREE_CODE (decl) == FUNCTION_DECL
817 || (TREE_CODE (decl) == VAR_DECL && is_global_var (decl)))
818 name_tree = DECL_ASSEMBLER_NAME (decl);
819 else
820 name_tree = DECL_NAME (decl);
822 const char *name = IDENTIFIER_POINTER (name_tree);
823 /* User-defined assembly names have prepended asterisk symbol. */
824 if (name[0] == '*')
825 name++;
827 if ((TREE_CODE (decl) == VAR_DECL)
828 && decl_function_context (decl))
830 size_t len = strlen (name);
831 char *buf = (char *) alloca (len + 32);
832 snprintf (buf, len + 32, "%s_%u", name, DECL_UID (decl));
833 len = strlen (buf);
834 char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1);
835 memcpy (copy, buf, len + 1);
836 return copy;
838 else
839 return name;
842 /* Lookup or create the associated hsa_symbol structure with a given VAR_DECL
843 or lookup the hsa_structure corresponding to a PARM_DECL. */
845 static hsa_symbol *
846 get_symbol_for_decl (tree decl)
848 hsa_symbol **slot;
849 hsa_symbol dummy (BRIG_TYPE_NONE, BRIG_SEGMENT_NONE, BRIG_LINKAGE_NONE);
851 gcc_assert (TREE_CODE (decl) == PARM_DECL
852 || TREE_CODE (decl) == RESULT_DECL
853 || TREE_CODE (decl) == VAR_DECL
854 || TREE_CODE (decl) == CONST_DECL);
856 dummy.m_decl = decl;
858 bool is_in_global_vars = ((TREE_CODE (decl) == VAR_DECL)
859 && !decl_function_context (decl));
861 if (is_in_global_vars)
862 slot = hsa_global_variable_symbols->find_slot (&dummy, INSERT);
863 else
864 slot = hsa_cfun->m_local_symbols->find_slot (&dummy, INSERT);
866 gcc_checking_assert (slot);
867 if (*slot)
869 hsa_symbol *sym = (*slot);
871 /* If the symbol is problematic, mark current function also as
872 problematic. */
873 if (sym->m_seen_error)
874 hsa_fail_cfun ();
876 /* PR hsa/70234: If a global variable was marked to be emitted,
877 but HSAIL generation of a function using the variable fails,
878 we should retry to emit the variable in context of a different
879 function.
881 Iterate elements whether a symbol is already in m_global_symbols
882 of not. */
883 if (is_in_global_vars && !sym->m_emitted_to_brig)
885 for (unsigned i = 0; i < hsa_cfun->m_global_symbols.length (); i++)
886 if (hsa_cfun->m_global_symbols[i] == sym)
887 return *slot;
888 hsa_cfun->m_global_symbols.safe_push (sym);
891 return *slot;
893 else
895 hsa_symbol *sym;
896 /* PARM_DECLs and RESULT_DECL should be already in m_local_symbols. */
897 gcc_assert (TREE_CODE (decl) == VAR_DECL
898 || TREE_CODE (decl) == CONST_DECL);
899 BrigAlignment8_t align = hsa_object_alignment (decl);
901 if (is_in_global_vars)
903 gcc_checking_assert (TREE_CODE (decl) != CONST_DECL);
904 sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_GLOBAL,
905 BRIG_LINKAGE_PROGRAM, true,
906 BRIG_ALLOCATION_PROGRAM, align);
907 hsa_cfun->m_global_symbols.safe_push (sym);
908 sym->fillup_for_decl (decl);
909 if (sym->m_align > align)
911 sym->m_seen_error = true;
912 HSA_SORRY_ATV (EXPR_LOCATION (decl),
913 "HSA specification requires that %E is at least "
914 "naturally aligned", decl);
917 else
919 /* As generation of efficient memory copy instructions relies
920 on alignment greater or equal to 8 bytes,
921 we need to increase alignment of all aggregate types.. */
922 if (AGGREGATE_TYPE_P (TREE_TYPE (decl)))
923 align = MAX ((BrigAlignment8_t) BRIG_ALIGNMENT_8, align);
925 BrigAllocation allocation = BRIG_ALLOCATION_AUTOMATIC;
926 BrigSegment8_t segment;
927 if (TREE_CODE (decl) == CONST_DECL)
929 segment = BRIG_SEGMENT_READONLY;
930 allocation = BRIG_ALLOCATION_AGENT;
932 else if (lookup_attribute ("hsa_group_segment",
933 DECL_ATTRIBUTES (decl)))
934 segment = BRIG_SEGMENT_GROUP;
935 else if (TREE_STATIC (decl))
937 segment = BRIG_SEGMENT_GLOBAL;
938 allocation = BRIG_ALLOCATION_PROGRAM;
940 else if (lookup_attribute ("hsa_global_segment",
941 DECL_ATTRIBUTES (decl)))
942 segment = BRIG_SEGMENT_GLOBAL;
943 else
944 segment = BRIG_SEGMENT_PRIVATE;
946 sym = new hsa_symbol (BRIG_TYPE_NONE, segment, BRIG_LINKAGE_FUNCTION,
947 false, allocation, align);
948 sym->fillup_for_decl (decl);
949 hsa_cfun->m_private_variables.safe_push (sym);
952 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_bound_function ? s->m_bound_function->decl : NULL;
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
981 && s->m_gpu_implementation_p
982 && s->m_bound_function)
983 d = s->m_bound_function->decl;
985 /* IPA split can create a function that has no host equivalent. */
986 if (d == NULL)
987 d = decl;
989 char *name = xstrdup (hsa_get_declaration_name (d));
990 hsa_sanitize_name (name);
992 return name;
995 /* Create a spill symbol of type TYPE. */
997 hsa_symbol *
998 hsa_get_spill_symbol (BrigType16_t type)
1000 hsa_symbol *sym = new hsa_symbol (type, BRIG_SEGMENT_SPILL,
1001 BRIG_LINKAGE_FUNCTION);
1002 hsa_cfun->m_spill_symbols.safe_push (sym);
1003 return sym;
1006 /* Create a symbol for a read-only string constant. */
1007 hsa_symbol *
1008 hsa_get_string_cst_symbol (tree string_cst)
1010 gcc_checking_assert (TREE_CODE (string_cst) == STRING_CST);
1012 hsa_symbol **slot = hsa_cfun->m_string_constants_map.get (string_cst);
1013 if (slot)
1014 return *slot;
1016 hsa_op_immed *cst = new hsa_op_immed (string_cst);
1017 hsa_symbol *sym = new hsa_symbol (cst->m_type, BRIG_SEGMENT_GLOBAL,
1018 BRIG_LINKAGE_MODULE, true,
1019 BRIG_ALLOCATION_AGENT);
1020 sym->m_cst_value = cst;
1021 sym->m_dim = TREE_STRING_LENGTH (string_cst);
1022 sym->m_name_number = hsa_cfun->m_global_symbols.length ();
1024 hsa_cfun->m_global_symbols.safe_push (sym);
1025 hsa_cfun->m_string_constants_map.put (string_cst, sym);
1026 return sym;
1029 /* Make the type of a MOV instruction larger if mandated by HSAIL rules. */
1031 static void
1032 hsa_fixup_mov_insn_type (hsa_insn_basic *insn)
1034 insn->m_type = hsa_extend_inttype_to_32bit (insn->m_type);
1035 if (insn->m_type == BRIG_TYPE_B8 || insn->m_type == BRIG_TYPE_B16)
1036 insn->m_type = BRIG_TYPE_B32;
1039 /* Constructor of the ancestor of all operands. K is BRIG kind that identified
1040 what the operator is. */
1042 hsa_op_base::hsa_op_base (BrigKind16_t k)
1043 : m_next (NULL), m_brig_op_offset (0), m_kind (k)
1045 hsa_operands.safe_push (this);
1048 /* Constructor of ancestor of all operands which have a type. K is BRIG kind
1049 that identified what the operator is. T is the type of the operator. */
1051 hsa_op_with_type::hsa_op_with_type (BrigKind16_t k, BrigType16_t t)
1052 : hsa_op_base (k), m_type (t)
1056 hsa_op_with_type *
1057 hsa_op_with_type::get_in_type (BrigType16_t dtype, hsa_bb *hbb)
1059 if (m_type == dtype)
1060 return this;
1062 hsa_op_reg *dest;
1064 if (hsa_needs_cvt (dtype, m_type))
1066 dest = new hsa_op_reg (dtype);
1067 hbb->append_insn (new hsa_insn_cvt (dest, this));
1069 else if (is_a <hsa_op_reg *> (this))
1071 /* In the end, HSA registers do not really have types, only sizes, so if
1072 the sizes match, we can use the register directly. */
1073 gcc_checking_assert (hsa_type_bit_size (dtype)
1074 == hsa_type_bit_size (m_type));
1075 return this;
1077 else
1079 dest = new hsa_op_reg (m_type);
1081 hsa_insn_basic *mov = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
1082 dest->m_type, dest, this);
1083 hsa_fixup_mov_insn_type (mov);
1084 hbb->append_insn (mov);
1085 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
1086 type of the operand must be same as type of the instruction. */
1087 dest->m_type = dtype;
1090 return dest;
1093 /* If this operand has integer type smaller than 32 bits, extend it to 32 bits,
1094 adding instructions to HBB if needed. */
1096 hsa_op_with_type *
1097 hsa_op_with_type::extend_int_to_32bit (hsa_bb *hbb)
1099 if (m_type == BRIG_TYPE_U8 || m_type == BRIG_TYPE_U16)
1100 return get_in_type (BRIG_TYPE_U32, hbb);
1101 else if (m_type == BRIG_TYPE_S8 || m_type == BRIG_TYPE_S16)
1102 return get_in_type (BRIG_TYPE_S32, hbb);
1103 else
1104 return this;
1107 /* Constructor of class representing HSA immediate values. TREE_VAL is the
1108 tree representation of the immediate value. If min32int is true,
1109 always expand integer types to one that has at least 32 bits. */
1111 hsa_op_immed::hsa_op_immed (tree tree_val, bool min32int)
1112 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES,
1113 hsa_type_for_tree_type (TREE_TYPE (tree_val), NULL,
1114 min32int))
1116 if (hsa_seen_error ())
1117 return;
1119 gcc_checking_assert ((is_gimple_min_invariant (tree_val)
1120 && (!POINTER_TYPE_P (TREE_TYPE (tree_val))
1121 || TREE_CODE (tree_val) == INTEGER_CST))
1122 || TREE_CODE (tree_val) == CONSTRUCTOR);
1123 m_tree_value = tree_val;
1125 /* Verify that all elements of a constructor are constants. */
1126 if (TREE_CODE (m_tree_value) == CONSTRUCTOR)
1127 for (unsigned i = 0; i < CONSTRUCTOR_NELTS (m_tree_value); i++)
1129 tree v = CONSTRUCTOR_ELT (m_tree_value, i)->value;
1130 if (!CONSTANT_CLASS_P (v))
1132 HSA_SORRY_AT (EXPR_LOCATION (tree_val),
1133 "HSA ctor should have only constants");
1134 return;
1139 /* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
1140 integer representation of the immediate value. TYPE is BRIG type. */
1142 hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value, BrigType16_t type)
1143 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES, type),
1144 m_tree_value (NULL)
1146 gcc_assert (hsa_type_integer_p (type));
1147 m_int_value = integer_value;
1150 hsa_op_immed::hsa_op_immed ()
1151 : hsa_op_with_type (BRIG_KIND_NONE, BRIG_TYPE_NONE)
1155 /* New operator to allocate immediate operands from obstack. */
1157 void *
1158 hsa_op_immed::operator new (size_t size)
1160 return obstack_alloc (&hsa_obstack, size);
1163 /* Destructor. */
1165 hsa_op_immed::~hsa_op_immed ()
1169 /* Change type of the immediate value to T. */
1171 void
1172 hsa_op_immed::set_type (BrigType16_t t)
1174 m_type = t;
1177 /* Constructor of class representing HSA registers and pseudo-registers. T is
1178 the BRIG type of the new register. */
1180 hsa_op_reg::hsa_op_reg (BrigType16_t t)
1181 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER, t), m_gimple_ssa (NULL_TREE),
1182 m_def_insn (NULL), m_spill_sym (NULL), m_order (hsa_cfun->m_reg_count++),
1183 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
1187 /* New operator to allocate a register from obstack. */
1189 void *
1190 hsa_op_reg::operator new (size_t size)
1192 return obstack_alloc (&hsa_obstack, size);
1195 /* Verify register operand. */
1197 void
1198 hsa_op_reg::verify_ssa ()
1200 /* Verify that each HSA register has a definition assigned.
1201 Exceptions are VAR_DECL and PARM_DECL that are a default
1202 definition. */
1203 gcc_checking_assert (m_def_insn
1204 || (m_gimple_ssa != NULL
1205 && (!SSA_NAME_VAR (m_gimple_ssa)
1206 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa))
1207 != PARM_DECL))
1208 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa)));
1210 /* Verify that every use of the register is really present
1211 in an instruction. */
1212 for (unsigned i = 0; i < m_uses.length (); i++)
1214 hsa_insn_basic *use = m_uses[i];
1216 bool is_visited = false;
1217 for (unsigned j = 0; j < use->operand_count (); j++)
1219 hsa_op_base *u = use->get_op (j);
1220 hsa_op_address *addr; addr = dyn_cast <hsa_op_address *> (u);
1221 if (addr && addr->m_reg)
1222 u = addr->m_reg;
1224 if (u == this)
1226 bool r = !addr && use->op_output_p (j);
1228 if (r)
1230 error ("HSA SSA name defined by instruction that is supposed "
1231 "to be using it");
1232 debug_hsa_operand (this);
1233 debug_hsa_insn (use);
1234 internal_error ("HSA SSA verification failed");
1237 is_visited = true;
1241 if (!is_visited)
1243 error ("HSA SSA name not among operands of instruction that is "
1244 "supposed to use it");
1245 debug_hsa_operand (this);
1246 debug_hsa_insn (use);
1247 internal_error ("HSA SSA verification failed");
1252 hsa_op_address::hsa_op_address (hsa_symbol *sym, hsa_op_reg *r,
1253 HOST_WIDE_INT offset)
1254 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (r),
1255 m_imm_offset (offset)
1259 hsa_op_address::hsa_op_address (hsa_symbol *sym, HOST_WIDE_INT offset)
1260 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (NULL),
1261 m_imm_offset (offset)
1265 hsa_op_address::hsa_op_address (hsa_op_reg *r, HOST_WIDE_INT offset)
1266 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (NULL), m_reg (r),
1267 m_imm_offset (offset)
1271 /* New operator to allocate address operands from obstack. */
1273 void *
1274 hsa_op_address::operator new (size_t size)
1276 return obstack_alloc (&hsa_obstack, size);
1279 /* Constructor of an operand referring to HSAIL code. */
1281 hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF),
1282 m_directive_offset (0)
1286 /* Constructor of an operand representing a code list. Set it up so that it
1287 can contain ELEMENTS number of elements. */
1289 hsa_op_code_list::hsa_op_code_list (unsigned elements)
1290 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST)
1292 m_offsets.create (1);
1293 m_offsets.safe_grow_cleared (elements);
1296 /* New operator to allocate code list operands from obstack. */
1298 void *
1299 hsa_op_code_list::operator new (size_t size)
1301 return obstack_alloc (&hsa_obstack, size);
1304 /* Constructor of an operand representing an operand list.
1305 Set it up so that it can contain ELEMENTS number of elements. */
1307 hsa_op_operand_list::hsa_op_operand_list (unsigned elements)
1308 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST)
1310 m_offsets.create (elements);
1311 m_offsets.safe_grow (elements);
1314 /* New operator to allocate operand list operands from obstack. */
1316 void *
1317 hsa_op_operand_list::operator new (size_t size)
1319 return obstack_alloc (&hsa_obstack, size);
1322 hsa_op_operand_list::~hsa_op_operand_list ()
1324 m_offsets.release ();
1328 hsa_op_reg *
1329 hsa_function_representation::reg_for_gimple_ssa (tree ssa)
1331 hsa_op_reg *hreg;
1333 gcc_checking_assert (TREE_CODE (ssa) == SSA_NAME);
1334 if (m_ssa_map[SSA_NAME_VERSION (ssa)])
1335 return m_ssa_map[SSA_NAME_VERSION (ssa)];
1337 hreg = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa),
1338 false));
1339 hreg->m_gimple_ssa = ssa;
1340 m_ssa_map[SSA_NAME_VERSION (ssa)] = hreg;
1342 return hreg;
1345 void
1346 hsa_op_reg::set_definition (hsa_insn_basic *insn)
1348 if (hsa_cfun->m_in_ssa)
1350 gcc_checking_assert (!m_def_insn);
1351 m_def_insn = insn;
1353 else
1354 m_def_insn = NULL;
1357 /* Constructor of the class which is the bases of all instructions and directly
1358 represents the most basic ones. NOPS is the number of operands that the
1359 operand vector will contain (and which will be cleared). OP is the opcode
1360 of the instruction. This constructor does not set type. */
1362 hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc)
1363 : m_prev (NULL),
1364 m_next (NULL), m_bb (NULL), m_opcode (opc), m_number (0),
1365 m_type (BRIG_TYPE_NONE), m_brig_offset (0)
1367 if (nops > 0)
1368 m_operands.safe_grow_cleared (nops);
1370 hsa_instructions.safe_push (this);
1373 /* Make OP the operand number INDEX of operands of this instruction. If OP is a
1374 register or an address containing a register, then either set the definition
1375 of the register to this instruction if it an output operand or add this
1376 instruction to the uses if it is an input one. */
1378 void
1379 hsa_insn_basic::set_op (int index, hsa_op_base *op)
1381 /* Each address operand is always use. */
1382 hsa_op_address *addr = dyn_cast <hsa_op_address *> (op);
1383 if (addr && addr->m_reg)
1384 addr->m_reg->m_uses.safe_push (this);
1385 else
1387 hsa_op_reg *reg = dyn_cast <hsa_op_reg *> (op);
1388 if (reg)
1390 if (op_output_p (index))
1391 reg->set_definition (this);
1392 else
1393 reg->m_uses.safe_push (this);
1397 m_operands[index] = op;
1400 /* Get INDEX-th operand of the instruction. */
1402 hsa_op_base *
1403 hsa_insn_basic::get_op (int index)
1405 return m_operands[index];
1408 /* Get address of INDEX-th operand of the instruction. */
1410 hsa_op_base **
1411 hsa_insn_basic::get_op_addr (int index)
1413 return &m_operands[index];
1416 /* Get number of operands of the instruction. */
1417 unsigned int
1418 hsa_insn_basic::operand_count ()
1420 return m_operands.length ();
1423 /* Constructor of the class which is the bases of all instructions and directly
1424 represents the most basic ones. NOPS is the number of operands that the
1425 operand vector will contain (and which will be cleared). OPC is the opcode
1426 of the instruction, T is the type of the instruction. */
1428 hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc, BrigType16_t t,
1429 hsa_op_base *arg0, hsa_op_base *arg1,
1430 hsa_op_base *arg2, hsa_op_base *arg3)
1431 : m_prev (NULL), m_next (NULL), m_bb (NULL), m_opcode (opc),m_number (0),
1432 m_type (t), m_brig_offset (0)
1434 if (nops > 0)
1435 m_operands.safe_grow_cleared (nops);
1437 if (arg0 != NULL)
1439 gcc_checking_assert (nops >= 1);
1440 set_op (0, arg0);
1443 if (arg1 != NULL)
1445 gcc_checking_assert (nops >= 2);
1446 set_op (1, arg1);
1449 if (arg2 != NULL)
1451 gcc_checking_assert (nops >= 3);
1452 set_op (2, arg2);
1455 if (arg3 != NULL)
1457 gcc_checking_assert (nops >= 4);
1458 set_op (3, arg3);
1461 hsa_instructions.safe_push (this);
1464 /* New operator to allocate basic instruction from obstack. */
1466 void *
1467 hsa_insn_basic::operator new (size_t size)
1469 return obstack_alloc (&hsa_obstack, size);
1472 /* Verify the instruction. */
1474 void
1475 hsa_insn_basic::verify ()
1477 hsa_op_address *addr;
1478 hsa_op_reg *reg;
1480 /* Iterate all register operands and verify that the instruction
1481 is set in uses of the register. */
1482 for (unsigned i = 0; i < operand_count (); i++)
1484 hsa_op_base *use = get_op (i);
1486 if ((addr = dyn_cast <hsa_op_address *> (use)) && addr->m_reg)
1488 gcc_assert (addr->m_reg->m_def_insn != this);
1489 use = addr->m_reg;
1492 if ((reg = dyn_cast <hsa_op_reg *> (use)) && !op_output_p (i))
1494 unsigned j;
1495 for (j = 0; j < reg->m_uses.length (); j++)
1497 if (reg->m_uses[j] == this)
1498 break;
1501 if (j == reg->m_uses.length ())
1503 error ("HSA instruction uses a register but is not among "
1504 "recorded register uses");
1505 debug_hsa_operand (reg);
1506 debug_hsa_insn (this);
1507 internal_error ("HSA instruction verification failed");
1513 /* Constructor of an instruction representing a PHI node. NOPS is the number
1514 of operands (equal to the number of predecessors). */
1516 hsa_insn_phi::hsa_insn_phi (unsigned nops, hsa_op_reg *dst)
1517 : hsa_insn_basic (nops, HSA_OPCODE_PHI), m_dest (dst)
1519 dst->set_definition (this);
1522 /* Constructor of class representing instructions for control flow and
1523 sychronization, */
1525 hsa_insn_br::hsa_insn_br (unsigned nops, int opc, BrigType16_t t,
1526 BrigWidth8_t width, hsa_op_base *arg0,
1527 hsa_op_base *arg1, hsa_op_base *arg2,
1528 hsa_op_base *arg3)
1529 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
1530 m_width (width)
1534 /* Constructor of class representing instruction for conditional jump, CTRL is
1535 the control register determining whether the jump will be carried out, the
1536 new instruction is automatically added to its uses list. */
1538 hsa_insn_cbr::hsa_insn_cbr (hsa_op_reg *ctrl)
1539 : hsa_insn_br (1, BRIG_OPCODE_CBR, BRIG_TYPE_B1, BRIG_WIDTH_1, ctrl)
1543 /* Constructor of class representing instruction for switch jump, CTRL is
1544 the index register. */
1546 hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg *index, unsigned jump_count)
1547 : hsa_insn_basic (1, BRIG_OPCODE_SBR, BRIG_TYPE_B1, index),
1548 m_width (BRIG_WIDTH_1), m_jump_table (vNULL),
1549 m_label_code_list (new hsa_op_code_list (jump_count))
1553 /* Replace all occurrences of OLD_BB with NEW_BB in the statements
1554 jump table. */
1556 void
1557 hsa_insn_sbr::replace_all_labels (basic_block old_bb, basic_block new_bb)
1559 for (unsigned i = 0; i < m_jump_table.length (); i++)
1560 if (m_jump_table[i] == old_bb)
1561 m_jump_table[i] = new_bb;
1564 hsa_insn_sbr::~hsa_insn_sbr ()
1566 m_jump_table.release ();
1569 /* Constructor of comparison instruction. CMP is the comparison operation and T
1570 is the result type. */
1572 hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp, BrigType16_t t,
1573 hsa_op_base *arg0, hsa_op_base *arg1,
1574 hsa_op_base *arg2)
1575 : hsa_insn_basic (3 , BRIG_OPCODE_CMP, t, arg0, arg1, arg2), m_compare (cmp)
1579 /* Constructor of classes representing memory accesses. OPC is the opcode (must
1580 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
1581 operands are provided as ARG0 and ARG1. */
1583 hsa_insn_mem::hsa_insn_mem (int opc, BrigType16_t t, hsa_op_base *arg0,
1584 hsa_op_base *arg1)
1585 : hsa_insn_basic (2, opc, t, arg0, arg1),
1586 m_align (hsa_natural_alignment (t)), m_equiv_class (0)
1588 gcc_checking_assert (opc == BRIG_OPCODE_LD || opc == BRIG_OPCODE_ST);
1591 /* Constructor for descendants allowing different opcodes and number of
1592 operands, it passes its arguments directly to hsa_insn_basic
1593 constructor. The instruction operands are provided as ARG[0-3]. */
1596 hsa_insn_mem::hsa_insn_mem (unsigned nops, int opc, BrigType16_t t,
1597 hsa_op_base *arg0, hsa_op_base *arg1,
1598 hsa_op_base *arg2, hsa_op_base *arg3)
1599 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
1600 m_align (hsa_natural_alignment (t)), m_equiv_class (0)
1604 /* Constructor of class representing atomic instructions. OPC is the principal
1605 opcode, AOP is the specific atomic operation opcode. T is the type of the
1606 instruction. The instruction operands are provided as ARG[0-3]. */
1608 hsa_insn_atomic::hsa_insn_atomic (int nops, int opc,
1609 enum BrigAtomicOperation aop,
1610 BrigType16_t t, BrigMemoryOrder memorder,
1611 hsa_op_base *arg0,
1612 hsa_op_base *arg1, hsa_op_base *arg2,
1613 hsa_op_base *arg3)
1614 : hsa_insn_mem (nops, opc, t, arg0, arg1, arg2, arg3), m_atomicop (aop),
1615 m_memoryorder (memorder),
1616 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM)
1618 gcc_checking_assert (opc == BRIG_OPCODE_ATOMICNORET ||
1619 opc == BRIG_OPCODE_ATOMIC ||
1620 opc == BRIG_OPCODE_SIGNAL ||
1621 opc == BRIG_OPCODE_SIGNALNORET);
1624 /* Constructor of class representing signal instructions. OPC is the prinicpal
1625 opcode, SOP is the specific signal operation opcode. T is the type of the
1626 instruction. The instruction operands are provided as ARG[0-3]. */
1628 hsa_insn_signal::hsa_insn_signal (int nops, int opc,
1629 enum BrigAtomicOperation sop,
1630 BrigType16_t t, BrigMemoryOrder memorder,
1631 hsa_op_base *arg0, hsa_op_base *arg1,
1632 hsa_op_base *arg2, hsa_op_base *arg3)
1633 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
1634 m_memory_order (memorder), m_signalop (sop)
1638 /* Constructor of class representing segment conversion instructions. OPC is
1639 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
1640 and SRCT are destination and source types respectively, SEG is the segment
1641 we are converting to or from. The instruction operands are
1642 provided as ARG0 and ARG1. */
1644 hsa_insn_seg::hsa_insn_seg (int opc, BrigType16_t dest, BrigType16_t srct,
1645 BrigSegment8_t seg, hsa_op_base *arg0,
1646 hsa_op_base *arg1)
1647 : hsa_insn_basic (2, opc, dest, arg0, arg1), m_src_type (srct),
1648 m_segment (seg)
1650 gcc_checking_assert (opc == BRIG_OPCODE_STOF || opc == BRIG_OPCODE_FTOS);
1653 /* Constructor of class representing a call instruction. CALLEE is the tree
1654 representation of the function being called. */
1656 hsa_insn_call::hsa_insn_call (tree callee)
1657 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (callee),
1658 m_output_arg (NULL), m_args_code_list (NULL), m_result_code_list (NULL)
1662 hsa_insn_call::hsa_insn_call (hsa_internal_fn *fn)
1663 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (NULL),
1664 m_called_internal_fn (fn), m_output_arg (NULL), m_args_code_list (NULL),
1665 m_result_code_list (NULL)
1669 hsa_insn_call::~hsa_insn_call ()
1671 for (unsigned i = 0; i < m_input_args.length (); i++)
1672 delete m_input_args[i];
1674 delete m_output_arg;
1676 m_input_args.release ();
1677 m_input_arg_insns.release ();
1680 /* Constructor of class representing the argument block required to invoke
1681 a call in HSAIL. */
1682 hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind,
1683 hsa_insn_call * call)
1684 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK), m_kind (brig_kind),
1685 m_call_insn (call)
1689 hsa_insn_comment::hsa_insn_comment (const char *s)
1690 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT)
1692 unsigned l = strlen (s);
1694 /* Append '// ' to the string. */
1695 char *buf = XNEWVEC (char, l + 4);
1696 sprintf (buf, "// %s", s);
1697 m_comment = buf;
1700 hsa_insn_comment::~hsa_insn_comment ()
1702 gcc_checking_assert (m_comment);
1703 free (m_comment);
1704 m_comment = NULL;
1707 /* Constructor of class representing the queue instruction in HSAIL. */
1709 hsa_insn_queue::hsa_insn_queue (int nops, int opcode, BrigSegment segment,
1710 BrigMemoryOrder memory_order,
1711 hsa_op_base *arg0, hsa_op_base *arg1,
1712 hsa_op_base *arg2, hsa_op_base *arg3)
1713 : hsa_insn_basic (nops, opcode, BRIG_TYPE_U64, arg0, arg1, arg2, arg3),
1714 m_segment (segment), m_memory_order (memory_order)
1718 /* Constructor of class representing the source type instruction in HSAIL. */
1720 hsa_insn_srctype::hsa_insn_srctype (int nops, BrigOpcode opcode,
1721 BrigType16_t destt, BrigType16_t srct,
1722 hsa_op_base *arg0, hsa_op_base *arg1,
1723 hsa_op_base *arg2 = NULL)
1724 : hsa_insn_basic (nops, opcode, destt, arg0, arg1, arg2),
1725 m_source_type (srct)
1728 /* Constructor of class representing the packed instruction in HSAIL. */
1730 hsa_insn_packed::hsa_insn_packed (int nops, BrigOpcode opcode,
1731 BrigType16_t destt, BrigType16_t srct,
1732 hsa_op_base *arg0, hsa_op_base *arg1,
1733 hsa_op_base *arg2)
1734 : hsa_insn_srctype (nops, opcode, destt, srct, arg0, arg1, arg2)
1736 m_operand_list = new hsa_op_operand_list (nops - 1);
1739 /* Constructor of class representing the convert instruction in HSAIL. */
1741 hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type *dest, hsa_op_with_type *src)
1742 : hsa_insn_basic (2, BRIG_OPCODE_CVT, dest->m_type, dest, src)
1746 /* Constructor of class representing the alloca in HSAIL. */
1748 hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type *dest,
1749 hsa_op_with_type *size, unsigned alignment)
1750 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA, dest->m_type, dest, size),
1751 m_align (BRIG_ALIGNMENT_8)
1753 gcc_assert (dest->m_type == BRIG_TYPE_U32);
1754 if (alignment)
1755 m_align = hsa_alignment_encoding (alignment);
1758 /* Append an instruction INSN into the basic block. */
1760 void
1761 hsa_bb::append_insn (hsa_insn_basic *insn)
1763 gcc_assert (insn->m_opcode != 0 || insn->operand_count () == 0);
1764 gcc_assert (!insn->m_bb);
1766 insn->m_bb = m_bb;
1767 insn->m_prev = m_last_insn;
1768 insn->m_next = NULL;
1769 if (m_last_insn)
1770 m_last_insn->m_next = insn;
1771 m_last_insn = insn;
1772 if (!m_first_insn)
1773 m_first_insn = insn;
1776 void
1777 hsa_bb::append_phi (hsa_insn_phi *hphi)
1779 hphi->m_bb = m_bb;
1781 hphi->m_prev = m_last_phi;
1782 hphi->m_next = NULL;
1783 if (m_last_phi)
1784 m_last_phi->m_next = hphi;
1785 m_last_phi = hphi;
1786 if (!m_first_phi)
1787 m_first_phi = hphi;
1790 /* Insert HSA instruction NEW_INSN immediately before an existing instruction
1791 OLD_INSN. */
1793 static void
1794 hsa_insert_insn_before (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
1796 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
1798 if (hbb->m_first_insn == old_insn)
1799 hbb->m_first_insn = new_insn;
1800 new_insn->m_prev = old_insn->m_prev;
1801 new_insn->m_next = old_insn;
1802 if (old_insn->m_prev)
1803 old_insn->m_prev->m_next = new_insn;
1804 old_insn->m_prev = new_insn;
1807 /* Append HSA instruction NEW_INSN immediately after an existing instruction
1808 OLD_INSN. */
1810 static void
1811 hsa_append_insn_after (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
1813 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
1815 if (hbb->m_last_insn == old_insn)
1816 hbb->m_last_insn = new_insn;
1817 new_insn->m_prev = old_insn;
1818 new_insn->m_next = old_insn->m_next;
1819 if (old_insn->m_next)
1820 old_insn->m_next->m_prev = new_insn;
1821 old_insn->m_next = new_insn;
1824 /* Return a register containing the calculated value of EXP which must be an
1825 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
1826 integer constants as returned by get_inner_reference.
1827 Newly generated HSA instructions will be appended to HBB.
1828 Perform all calculations in ADDRTYPE. */
1830 static hsa_op_with_type *
1831 gen_address_calculation (tree exp, hsa_bb *hbb, BrigType16_t addrtype)
1833 int opcode;
1835 if (TREE_CODE (exp) == NOP_EXPR)
1836 exp = TREE_OPERAND (exp, 0);
1838 switch (TREE_CODE (exp))
1840 case SSA_NAME:
1841 return hsa_cfun->reg_for_gimple_ssa (exp)->get_in_type (addrtype, hbb);
1843 case INTEGER_CST:
1845 hsa_op_immed *imm = new hsa_op_immed (exp);
1846 if (addrtype != imm->m_type)
1847 imm->m_type = addrtype;
1848 return imm;
1851 case PLUS_EXPR:
1852 opcode = BRIG_OPCODE_ADD;
1853 break;
1855 case MULT_EXPR:
1856 opcode = BRIG_OPCODE_MUL;
1857 break;
1859 default:
1860 gcc_unreachable ();
1863 hsa_op_reg *res = new hsa_op_reg (addrtype);
1864 hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, addrtype);
1865 insn->set_op (0, res);
1867 hsa_op_with_type *op1 = gen_address_calculation (TREE_OPERAND (exp, 0), hbb,
1868 addrtype);
1869 hsa_op_with_type *op2 = gen_address_calculation (TREE_OPERAND (exp, 1), hbb,
1870 addrtype);
1871 insn->set_op (1, op1);
1872 insn->set_op (2, op2);
1874 hbb->append_insn (insn);
1875 return res;
1878 /* If R1 is NULL, just return R2, otherwise append an instruction adding them
1879 to HBB and return the register holding the result. */
1881 static hsa_op_reg *
1882 add_addr_regs_if_needed (hsa_op_reg *r1, hsa_op_reg *r2, hsa_bb *hbb)
1884 gcc_checking_assert (r2);
1885 if (!r1)
1886 return r2;
1888 hsa_op_reg *res = new hsa_op_reg (r1->m_type);
1889 gcc_assert (!hsa_needs_cvt (r1->m_type, r2->m_type));
1890 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_ADD, res->m_type);
1891 insn->set_op (0, res);
1892 insn->set_op (1, r1);
1893 insn->set_op (2, r2);
1894 hbb->append_insn (insn);
1895 return res;
1898 /* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
1899 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
1901 static void
1902 process_mem_base (tree base, hsa_symbol **symbol, BrigType16_t *addrtype,
1903 hsa_op_reg **reg, offset_int *offset, hsa_bb *hbb)
1905 if (TREE_CODE (base) == SSA_NAME)
1907 gcc_assert (!*reg);
1908 hsa_op_with_type *ssa
1909 = hsa_cfun->reg_for_gimple_ssa (base)->get_in_type (*addrtype, hbb);
1910 *reg = dyn_cast <hsa_op_reg *> (ssa);
1912 else if (TREE_CODE (base) == ADDR_EXPR)
1914 tree decl = TREE_OPERAND (base, 0);
1916 if (!DECL_P (decl) || TREE_CODE (decl) == FUNCTION_DECL)
1918 HSA_SORRY_AT (EXPR_LOCATION (base),
1919 "support for HSA does not implement a memory reference "
1920 "to a non-declaration type");
1921 return;
1924 gcc_assert (!*symbol);
1926 *symbol = get_symbol_for_decl (decl);
1927 *addrtype = hsa_get_segment_addr_type ((*symbol)->m_segment);
1929 else if (TREE_CODE (base) == INTEGER_CST)
1930 *offset += wi::to_offset (base);
1931 else
1932 gcc_unreachable ();
1935 /* Forward declaration of a function. */
1937 static void
1938 gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb);
1940 /* Generate HSA address operand for a given tree memory reference REF. If
1941 instructions need to be created to calculate the address, they will be added
1942 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
1943 the function assumes that the caller will handle possible
1944 bit-field references. Otherwise if we reference a bit-field, sorry message
1945 is displayed. */
1947 static hsa_op_address *
1948 gen_hsa_addr (tree ref, hsa_bb *hbb, HOST_WIDE_INT *output_bitsize = NULL,
1949 HOST_WIDE_INT *output_bitpos = NULL)
1951 hsa_symbol *symbol = NULL;
1952 hsa_op_reg *reg = NULL;
1953 offset_int offset = 0;
1954 tree origref = ref;
1955 tree varoffset = NULL_TREE;
1956 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
1957 HOST_WIDE_INT bitsize = 0, bitpos = 0;
1958 BrigType16_t flat_addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
1960 if (TREE_CODE (ref) == STRING_CST)
1962 symbol = hsa_get_string_cst_symbol (ref);
1963 goto out;
1965 else if (TREE_CODE (ref) == BIT_FIELD_REF
1966 && (!multiple_p (bit_field_size (ref), BITS_PER_UNIT)
1967 || !multiple_p (bit_field_offset (ref), BITS_PER_UNIT)))
1969 HSA_SORRY_ATV (EXPR_LOCATION (origref),
1970 "support for HSA does not implement "
1971 "bit field references such as %E", ref);
1972 goto out;
1975 if (handled_component_p (ref))
1977 machine_mode mode;
1978 int unsignedp, volatilep, preversep;
1979 poly_int64 pbitsize, pbitpos;
1980 tree new_ref;
1982 new_ref = get_inner_reference (ref, &pbitsize, &pbitpos, &varoffset,
1983 &mode, &unsignedp, &preversep,
1984 &volatilep);
1985 /* When this isn't true, the switch below will report an
1986 appropriate error. */
1987 if (pbitsize.is_constant () && pbitpos.is_constant ())
1989 bitsize = pbitsize.to_constant ();
1990 bitpos = pbitpos.to_constant ();
1991 ref = new_ref;
1992 offset = bitpos;
1993 offset = wi::rshift (offset, LOG2_BITS_PER_UNIT, SIGNED);
1997 switch (TREE_CODE (ref))
1999 case ADDR_EXPR:
2001 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
2002 symbol = hsa_cfun->create_hsa_temporary (flat_addrtype);
2003 hsa_op_reg *r = new hsa_op_reg (flat_addrtype);
2004 gen_hsa_addr_insns (ref, r, hbb);
2005 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
2006 r, new hsa_op_address (symbol)));
2008 break;
2010 case SSA_NAME:
2012 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
2013 hsa_op_with_type *r = hsa_cfun->reg_for_gimple_ssa (ref);
2014 if (r->m_type == BRIG_TYPE_B1)
2015 r = r->get_in_type (BRIG_TYPE_U32, hbb);
2016 symbol = hsa_cfun->create_hsa_temporary (r->m_type);
2018 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
2019 r, new hsa_op_address (symbol)));
2021 break;
2023 case PARM_DECL:
2024 case VAR_DECL:
2025 case RESULT_DECL:
2026 case CONST_DECL:
2027 gcc_assert (!symbol);
2028 symbol = get_symbol_for_decl (ref);
2029 addrtype = hsa_get_segment_addr_type (symbol->m_segment);
2030 break;
2032 case MEM_REF:
2033 process_mem_base (TREE_OPERAND (ref, 0), &symbol, &addrtype, &reg,
2034 &offset, hbb);
2036 if (!integer_zerop (TREE_OPERAND (ref, 1)))
2037 offset += wi::to_offset (TREE_OPERAND (ref, 1));
2038 break;
2040 case TARGET_MEM_REF:
2041 process_mem_base (TMR_BASE (ref), &symbol, &addrtype, &reg, &offset, hbb);
2042 if (TMR_INDEX (ref))
2044 hsa_op_reg *disp1;
2045 hsa_op_base *idx = hsa_cfun->reg_for_gimple_ssa
2046 (TMR_INDEX (ref))->get_in_type (addrtype, hbb);
2047 if (TMR_STEP (ref) && !integer_onep (TMR_STEP (ref)))
2049 disp1 = new hsa_op_reg (addrtype);
2050 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_MUL,
2051 addrtype);
2053 /* As step must respect addrtype, we overwrite the type
2054 of an immediate value. */
2055 hsa_op_immed *step = new hsa_op_immed (TMR_STEP (ref));
2056 step->m_type = addrtype;
2058 insn->set_op (0, disp1);
2059 insn->set_op (1, idx);
2060 insn->set_op (2, step);
2061 hbb->append_insn (insn);
2063 else
2064 disp1 = as_a <hsa_op_reg *> (idx);
2065 reg = add_addr_regs_if_needed (reg, disp1, hbb);
2067 if (TMR_INDEX2 (ref))
2069 if (TREE_CODE (TMR_INDEX2 (ref)) == SSA_NAME)
2071 hsa_op_base *disp2 = hsa_cfun->reg_for_gimple_ssa
2072 (TMR_INDEX2 (ref))->get_in_type (addrtype, hbb);
2073 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (disp2),
2074 hbb);
2076 else if (TREE_CODE (TMR_INDEX2 (ref)) == INTEGER_CST)
2077 offset += wi::to_offset (TMR_INDEX2 (ref));
2078 else
2079 gcc_unreachable ();
2081 offset += wi::to_offset (TMR_OFFSET (ref));
2082 break;
2083 case FUNCTION_DECL:
2084 HSA_SORRY_AT (EXPR_LOCATION (origref),
2085 "support for HSA does not implement function pointers");
2086 goto out;
2087 default:
2088 HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does "
2089 "not implement memory access to %E", origref);
2090 goto out;
2093 if (varoffset)
2095 if (TREE_CODE (varoffset) == INTEGER_CST)
2096 offset += wi::to_offset (varoffset);
2097 else
2099 hsa_op_base *off_op = gen_address_calculation (varoffset, hbb,
2100 addrtype);
2101 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (off_op),
2102 hbb);
2106 gcc_checking_assert ((symbol
2107 && addrtype
2108 == hsa_get_segment_addr_type (symbol->m_segment))
2109 || (!symbol
2110 && addrtype
2111 == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT)));
2112 out:
2113 HOST_WIDE_INT hwi_offset = offset.to_shwi ();
2115 /* Calculate remaining bitsize offset (if presented). */
2116 bitpos %= BITS_PER_UNIT;
2117 /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
2118 is not a reason to think this is a bit-field access. */
2119 if (bitpos == 0
2120 && (bitsize >= BITS_PER_UNIT)
2121 && !(bitsize & (bitsize - 1)))
2122 bitsize = 0;
2124 if ((bitpos || bitsize) && (output_bitpos == NULL || output_bitsize == NULL))
2125 HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does not "
2126 "implement unhandled bit field reference such as %E", ref);
2128 if (output_bitsize != NULL && output_bitpos != NULL)
2130 *output_bitsize = bitsize;
2131 *output_bitpos = bitpos;
2134 return new hsa_op_address (symbol, reg, hwi_offset);
2137 /* Generate HSA address operand for a given tree memory reference REF. If
2138 instructions need to be created to calculate the address, they will be added
2139 to the end of HBB. OUTPUT_ALIGN is alignment of the created address. */
2141 static hsa_op_address *
2142 gen_hsa_addr_with_align (tree ref, hsa_bb *hbb, BrigAlignment8_t *output_align)
2144 hsa_op_address *addr = gen_hsa_addr (ref, hbb);
2145 if (addr->m_reg || !addr->m_symbol)
2146 *output_align = hsa_object_alignment (ref);
2147 else
2149 /* If the address consists only of a symbol and an offset, we
2150 compute the alignment ourselves to take into account any alignment
2151 promotions we might have done for the HSA symbol representation. */
2152 unsigned align = hsa_byte_alignment (addr->m_symbol->m_align);
2153 unsigned misalign = addr->m_imm_offset & (align - 1);
2154 if (misalign)
2155 align = least_bit_hwi (misalign);
2156 *output_align = hsa_alignment_encoding (BITS_PER_UNIT * align);
2158 return addr;
2161 /* Generate HSA address for a function call argument of given TYPE.
2162 INDEX is used to generate corresponding name of the arguments.
2163 Special value -1 represents fact that result value is created. */
2165 static hsa_op_address *
2166 gen_hsa_addr_for_arg (tree tree_type, int index)
2168 hsa_symbol *sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
2169 BRIG_LINKAGE_ARG);
2170 sym->m_type = hsa_type_for_tree_type (tree_type, &sym->m_dim);
2172 if (index == -1) /* Function result. */
2173 sym->m_name = "res";
2174 else /* Function call arguments. */
2176 sym->m_name = NULL;
2177 sym->m_name_number = index;
2180 return new hsa_op_address (sym);
2183 /* Generate HSA instructions that process all necessary conversions
2184 of an ADDR to flat addressing and place the result into DEST.
2185 Instructions are appended to HBB. */
2187 static void
2188 convert_addr_to_flat_segment (hsa_op_address *addr, hsa_op_reg *dest,
2189 hsa_bb *hbb)
2191 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_LDA);
2192 insn->set_op (1, addr);
2193 if (addr->m_symbol && addr->m_symbol->m_segment != BRIG_SEGMENT_GLOBAL)
2195 /* LDA produces segment-relative address, we need to convert
2196 it to the flat one. */
2197 hsa_op_reg *tmp;
2198 tmp = new hsa_op_reg (hsa_get_segment_addr_type
2199 (addr->m_symbol->m_segment));
2200 hsa_insn_seg *seg;
2201 seg = new hsa_insn_seg (BRIG_OPCODE_STOF,
2202 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT),
2203 tmp->m_type, addr->m_symbol->m_segment, dest,
2204 tmp);
2206 insn->set_op (0, tmp);
2207 insn->m_type = tmp->m_type;
2208 hbb->append_insn (insn);
2209 hbb->append_insn (seg);
2211 else
2213 insn->set_op (0, dest);
2214 insn->m_type = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
2215 hbb->append_insn (insn);
2219 /* Generate HSA instructions that calculate address of VAL including all
2220 necessary conversions to flat addressing and place the result into DEST.
2221 Instructions are appended to HBB. */
2223 static void
2224 gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb)
2226 /* Handle cases like tmp = NULL, where we just emit a move instruction
2227 to a register. */
2228 if (TREE_CODE (val) == INTEGER_CST)
2230 hsa_op_immed *c = new hsa_op_immed (val);
2231 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
2232 dest->m_type, dest, c);
2233 hbb->append_insn (insn);
2234 return;
2237 hsa_op_address *addr;
2239 gcc_assert (dest->m_type == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
2240 if (TREE_CODE (val) == ADDR_EXPR)
2241 val = TREE_OPERAND (val, 0);
2242 addr = gen_hsa_addr (val, hbb);
2244 if (TREE_CODE (val) == CONST_DECL
2245 && is_gimple_reg_type (TREE_TYPE (val)))
2247 gcc_assert (addr->m_symbol
2248 && addr->m_symbol->m_segment == BRIG_SEGMENT_READONLY);
2249 /* CONST_DECLs are in readonly segment which however does not have
2250 addresses convertible to flat segments. So copy it to a private one
2251 and take address of that. */
2252 BrigType16_t csttype
2253 = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (val),
2254 false));
2255 hsa_op_reg *r = new hsa_op_reg (csttype);
2256 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD, csttype, r,
2257 new hsa_op_address (addr->m_symbol)));
2258 hsa_symbol *copysym = hsa_cfun->create_hsa_temporary (csttype);
2259 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, csttype, r,
2260 new hsa_op_address (copysym)));
2261 addr->m_symbol = copysym;
2263 else if (addr->m_symbol && addr->m_symbol->m_segment == BRIG_SEGMENT_READONLY)
2265 HSA_SORRY_ATV (EXPR_LOCATION (val), "support for HSA does "
2266 "not implement taking addresses of complex "
2267 "CONST_DECLs such as %E", val);
2268 return;
2272 convert_addr_to_flat_segment (addr, dest, hbb);
2275 /* Return an HSA register or HSA immediate value operand corresponding to
2276 gimple operand OP. */
2278 static hsa_op_with_type *
2279 hsa_reg_or_immed_for_gimple_op (tree op, hsa_bb *hbb)
2281 hsa_op_reg *tmp;
2283 if (TREE_CODE (op) == SSA_NAME)
2284 tmp = hsa_cfun->reg_for_gimple_ssa (op);
2285 else if (!POINTER_TYPE_P (TREE_TYPE (op)))
2286 return new hsa_op_immed (op);
2287 else
2289 tmp = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
2290 gen_hsa_addr_insns (op, tmp, hbb);
2292 return tmp;
2295 /* Create a simple movement instruction with register destination DEST and
2296 register or immediate source SRC and append it to the end of HBB. */
2298 void
2299 hsa_build_append_simple_mov (hsa_op_reg *dest, hsa_op_base *src, hsa_bb *hbb)
2301 /* Moves of packed data between registers need to adhere to the same type
2302 rules like when dealing with memory. */
2303 BrigType16_t tp = mem_type_for_type (dest->m_type);
2304 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, tp, dest, src);
2305 hsa_fixup_mov_insn_type (insn);
2306 unsigned dest_size = hsa_type_bit_size (dest->m_type);
2307 if (hsa_op_reg *sreg = dyn_cast <hsa_op_reg *> (src))
2308 gcc_assert (dest_size == hsa_type_bit_size (sreg->m_type));
2309 else
2311 unsigned imm_size
2312 = hsa_type_bit_size (as_a <hsa_op_immed *> (src)->m_type);
2313 gcc_assert ((dest_size == imm_size)
2314 /* Eventually < 32bit registers will be promoted to 32bit. */
2315 || (dest_size < 32 && imm_size == 32));
2317 hbb->append_insn (insn);
2320 /* Generate HSAIL instructions loading a bit field into register DEST.
2321 VALUE_REG is a register of a SSA name that is used in the bit field
2322 reference. To identify a bit field BITPOS is offset to the loaded memory
2323 and BITSIZE is number of bits of the bit field.
2324 Add instructions to HBB. */
2326 static void
2327 gen_hsa_insns_for_bitfield (hsa_op_reg *dest, hsa_op_reg *value_reg,
2328 HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos,
2329 hsa_bb *hbb)
2331 unsigned type_bitsize
2332 = hsa_type_bit_size (hsa_extend_inttype_to_32bit (dest->m_type));
2333 unsigned left_shift = type_bitsize - (bitsize + bitpos);
2334 unsigned right_shift = left_shift + bitpos;
2336 if (left_shift)
2338 hsa_op_reg *value_reg_2
2339 = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest->m_type));
2340 hsa_op_immed *c = new hsa_op_immed (left_shift, BRIG_TYPE_U32);
2342 hsa_insn_basic *lshift
2343 = new hsa_insn_basic (3, BRIG_OPCODE_SHL, value_reg_2->m_type,
2344 value_reg_2, value_reg, c);
2346 hbb->append_insn (lshift);
2348 value_reg = value_reg_2;
2351 if (right_shift)
2353 hsa_op_reg *value_reg_2
2354 = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest->m_type));
2355 hsa_op_immed *c = new hsa_op_immed (right_shift, BRIG_TYPE_U32);
2357 hsa_insn_basic *rshift
2358 = new hsa_insn_basic (3, BRIG_OPCODE_SHR, value_reg_2->m_type,
2359 value_reg_2, value_reg, c);
2361 hbb->append_insn (rshift);
2363 value_reg = value_reg_2;
2366 hsa_insn_basic *assignment
2367 = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, NULL, value_reg);
2368 hsa_fixup_mov_insn_type (assignment);
2369 hbb->append_insn (assignment);
2370 assignment->set_output_in_type (dest, 0, hbb);
2374 /* Generate HSAIL instructions loading a bit field into register DEST. ADDR is
2375 prepared memory address which is used to load the bit field. To identify a
2376 bit field BITPOS is offset to the loaded memory and BITSIZE is number of
2377 bits of the bit field. Add instructions to HBB. Load must be performed in
2378 alignment ALIGN. */
2380 static void
2381 gen_hsa_insns_for_bitfield_load (hsa_op_reg *dest, hsa_op_address *addr,
2382 HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos,
2383 hsa_bb *hbb, BrigAlignment8_t align)
2385 hsa_op_reg *value_reg = new hsa_op_reg (dest->m_type);
2386 hsa_insn_mem *mem
2387 = new hsa_insn_mem (BRIG_OPCODE_LD,
2388 hsa_extend_inttype_to_32bit (dest->m_type),
2389 value_reg, addr);
2390 mem->set_align (align);
2391 hbb->append_insn (mem);
2392 gen_hsa_insns_for_bitfield (dest, value_reg, bitsize, bitpos, hbb);
2395 /* Return the alignment of base memory accesses we issue to perform bit-field
2396 memory access REF. */
2398 static BrigAlignment8_t
2399 hsa_bitmemref_alignment (tree ref)
2401 unsigned HOST_WIDE_INT bit_offset = 0;
2403 while (true)
2405 if (TREE_CODE (ref) == BIT_FIELD_REF)
2407 if (!tree_fits_uhwi_p (TREE_OPERAND (ref, 2)))
2408 return BRIG_ALIGNMENT_1;
2409 bit_offset += tree_to_uhwi (TREE_OPERAND (ref, 2));
2411 else if (TREE_CODE (ref) == COMPONENT_REF
2412 && DECL_BIT_FIELD (TREE_OPERAND (ref, 1)))
2413 bit_offset += int_bit_position (TREE_OPERAND (ref, 1));
2414 else
2415 break;
2416 ref = TREE_OPERAND (ref, 0);
2419 unsigned HOST_WIDE_INT bits = bit_offset % BITS_PER_UNIT;
2420 unsigned HOST_WIDE_INT byte_bits = bit_offset - bits;
2421 BrigAlignment8_t base = hsa_object_alignment (ref);
2422 if (byte_bits == 0)
2423 return base;
2424 return MIN (base, hsa_alignment_encoding (least_bit_hwi (byte_bits)));
2427 /* Generate HSAIL instructions loading something into register DEST. RHS is
2428 tree representation of the loaded data, which are loaded as type TYPE. Add
2429 instructions to HBB. */
2431 static void
2432 gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb)
2434 /* The destination SSA name will give us the type. */
2435 if (TREE_CODE (rhs) == VIEW_CONVERT_EXPR)
2436 rhs = TREE_OPERAND (rhs, 0);
2438 if (TREE_CODE (rhs) == SSA_NAME)
2440 hsa_op_reg *src = hsa_cfun->reg_for_gimple_ssa (rhs);
2441 hsa_build_append_simple_mov (dest, src, hbb);
2443 else if (is_gimple_min_invariant (rhs)
2444 || TREE_CODE (rhs) == ADDR_EXPR)
2446 if (POINTER_TYPE_P (TREE_TYPE (rhs)))
2448 if (dest->m_type != hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT))
2450 HSA_SORRY_ATV (EXPR_LOCATION (rhs),
2451 "support for HSA does not implement conversion "
2452 "of %E to the requested non-pointer type.", rhs);
2453 return;
2456 gen_hsa_addr_insns (rhs, dest, hbb);
2458 else if (TREE_CODE (rhs) == COMPLEX_CST)
2460 hsa_op_immed *real_part = new hsa_op_immed (TREE_REALPART (rhs));
2461 hsa_op_immed *imag_part = new hsa_op_immed (TREE_IMAGPART (rhs));
2463 hsa_op_reg *real_part_reg
2464 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type),
2465 true));
2466 hsa_op_reg *imag_part_reg
2467 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type),
2468 true));
2470 hsa_build_append_simple_mov (real_part_reg, real_part, hbb);
2471 hsa_build_append_simple_mov (imag_part_reg, imag_part, hbb);
2473 BrigType16_t src_type = hsa_bittype_for_type (real_part_reg->m_type);
2475 hsa_insn_packed *insn
2476 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type,
2477 src_type, dest, real_part_reg,
2478 imag_part_reg);
2479 hbb->append_insn (insn);
2481 else
2483 hsa_op_immed *imm = new hsa_op_immed (rhs);
2484 hsa_build_append_simple_mov (dest, imm, hbb);
2487 else if (TREE_CODE (rhs) == REALPART_EXPR || TREE_CODE (rhs) == IMAGPART_EXPR)
2489 tree pack_type = TREE_TYPE (TREE_OPERAND (rhs, 0));
2491 hsa_op_reg *packed_reg
2492 = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type, true));
2494 tree complex_rhs = TREE_OPERAND (rhs, 0);
2495 gen_hsa_insns_for_load (packed_reg, complex_rhs, TREE_TYPE (complex_rhs),
2496 hbb);
2498 hsa_op_reg *real_reg
2499 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true));
2501 hsa_op_reg *imag_reg
2502 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true));
2504 BrigKind16_t brig_type = packed_reg->m_type;
2505 hsa_insn_packed *packed
2506 = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND,
2507 hsa_bittype_for_type (real_reg->m_type),
2508 brig_type, real_reg, imag_reg, packed_reg);
2510 hbb->append_insn (packed);
2512 hsa_op_reg *source = TREE_CODE (rhs) == REALPART_EXPR ?
2513 real_reg : imag_reg;
2515 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
2516 dest->m_type, NULL, source);
2517 hsa_fixup_mov_insn_type (insn);
2518 hbb->append_insn (insn);
2519 insn->set_output_in_type (dest, 0, hbb);
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 hsa_fixup_mov_insn_type (basic);
2656 hbb->append_insn (basic);
2658 if (bitpos)
2660 hsa_op_reg *shifted_value_reg = new hsa_op_reg (mem_type);
2661 c = new hsa_op_immed (bitpos, BRIG_TYPE_U32);
2663 hsa_insn_basic *basic
2664 = new hsa_insn_basic (3, BRIG_OPCODE_SHL, mem_type,
2665 shifted_value_reg, new_value_reg, c);
2666 hbb->append_insn (basic);
2668 new_value_reg = shifted_value_reg;
2671 /* OR the prepared value with prepared chunk loaded from memory. */
2672 hsa_op_reg *prepared_reg= new hsa_op_reg (mem_type);
2673 basic = new hsa_insn_basic (3, BRIG_OPCODE_OR, mem_type, prepared_reg,
2674 new_value_reg, cleared_reg);
2675 hbb->append_insn (basic);
2677 src = prepared_reg;
2678 mtype = mem_type;
2680 else
2681 req_align = hsa_object_alignment (lhs);
2683 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, mtype, src, addr);
2684 mem->set_align (req_align);
2686 /* The HSAIL verifier has another constraint: if the source is an immediate
2687 then it must match the destination type. If it's a register the low bits
2688 will be used for sub-word stores. We're always allocating new operands so
2689 we can modify the above in place. */
2690 if (hsa_op_immed *imm = dyn_cast <hsa_op_immed *> (src))
2692 if (!hsa_type_packed_p (imm->m_type))
2693 imm->m_type = mem->m_type;
2694 else
2696 /* ...and all vector immediates apparently need to be vectors of
2697 unsigned bytes. */
2698 unsigned bs = hsa_type_bit_size (imm->m_type);
2699 gcc_assert (bs == hsa_type_bit_size (mem->m_type));
2700 switch (bs)
2702 case 32:
2703 imm->m_type = BRIG_TYPE_U8X4;
2704 break;
2705 case 64:
2706 imm->m_type = BRIG_TYPE_U8X8;
2707 break;
2708 case 128:
2709 imm->m_type = BRIG_TYPE_U8X16;
2710 break;
2711 default:
2712 gcc_unreachable ();
2717 hbb->append_insn (mem);
2720 /* Generate memory copy instructions that are going to be used
2721 for copying a SRC memory to TARGET memory,
2722 represented by pointer in a register. MIN_ALIGN is minimal alignment
2723 of provided HSA addresses. */
2725 static void
2726 gen_hsa_memory_copy (hsa_bb *hbb, hsa_op_address *target, hsa_op_address *src,
2727 unsigned size, BrigAlignment8_t min_align)
2729 hsa_op_address *addr;
2730 hsa_insn_mem *mem;
2732 unsigned offset = 0;
2733 unsigned min_byte_align = hsa_byte_alignment (min_align);
2735 while (size)
2737 unsigned s;
2738 if (size >= 8)
2739 s = 8;
2740 else if (size >= 4)
2741 s = 4;
2742 else if (size >= 2)
2743 s = 2;
2744 else
2745 s = 1;
2747 if (s > min_byte_align)
2748 s = min_byte_align;
2750 BrigType16_t t = get_integer_type_by_bytes (s, false);
2752 hsa_op_reg *tmp = new hsa_op_reg (t);
2753 addr = new hsa_op_address (src->m_symbol, src->m_reg,
2754 src->m_imm_offset + offset);
2755 mem = new hsa_insn_mem (BRIG_OPCODE_LD, t, tmp, addr);
2756 hbb->append_insn (mem);
2758 addr = new hsa_op_address (target->m_symbol, target->m_reg,
2759 target->m_imm_offset + offset);
2760 mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, tmp, addr);
2761 hbb->append_insn (mem);
2762 offset += s;
2763 size -= s;
2767 /* Create a memset mask that is created by copying a CONSTANT byte value
2768 to an integer of BYTE_SIZE bytes. */
2770 static unsigned HOST_WIDE_INT
2771 build_memset_value (unsigned HOST_WIDE_INT constant, unsigned byte_size)
2773 if (constant == 0)
2774 return 0;
2776 HOST_WIDE_INT v = constant;
2778 for (unsigned i = 1; i < byte_size; i++)
2779 v |= constant << (8 * i);
2781 return v;
2784 /* Generate memory set instructions that are going to be used
2785 for setting a CONSTANT byte value to TARGET memory of SIZE bytes.
2786 MIN_ALIGN is minimal alignment of provided HSA addresses. */
2788 static void
2789 gen_hsa_memory_set (hsa_bb *hbb, hsa_op_address *target,
2790 unsigned HOST_WIDE_INT constant,
2791 unsigned size, BrigAlignment8_t min_align)
2793 hsa_op_address *addr;
2794 hsa_insn_mem *mem;
2796 unsigned offset = 0;
2797 unsigned min_byte_align = hsa_byte_alignment (min_align);
2799 while (size)
2801 unsigned s;
2802 if (size >= 8)
2803 s = 8;
2804 else if (size >= 4)
2805 s = 4;
2806 else if (size >= 2)
2807 s = 2;
2808 else
2809 s = 1;
2811 if (s > min_byte_align)
2812 s = min_byte_align;
2814 addr = new hsa_op_address (target->m_symbol, target->m_reg,
2815 target->m_imm_offset + offset);
2817 BrigType16_t t = get_integer_type_by_bytes (s, false);
2818 HOST_WIDE_INT c = build_memset_value (constant, s);
2820 mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, new hsa_op_immed (c, t),
2821 addr);
2822 hbb->append_insn (mem);
2823 offset += s;
2824 size -= s;
2828 /* Generate HSAIL instructions for a single assignment
2829 of an empty constructor to an ADDR_LHS. Constructor is passed as a
2830 tree RHS and all instructions are appended to HBB. ALIGN is
2831 alignment of the address. */
2833 void
2834 gen_hsa_ctor_assignment (hsa_op_address *addr_lhs, tree rhs, hsa_bb *hbb,
2835 BrigAlignment8_t align)
2837 if (CONSTRUCTOR_NELTS (rhs))
2839 HSA_SORRY_AT (EXPR_LOCATION (rhs),
2840 "support for HSA does not implement load from constructor");
2841 return;
2844 unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs)));
2845 gen_hsa_memory_set (hbb, addr_lhs, 0, size, align);
2848 /* Generate HSA instructions for a single assignment of RHS to LHS.
2849 HBB is the basic block they will be appended to. */
2851 static void
2852 gen_hsa_insns_for_single_assignment (tree lhs, tree rhs, hsa_bb *hbb)
2854 if (TREE_CODE (lhs) == SSA_NAME)
2856 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
2857 if (hsa_seen_error ())
2858 return;
2860 gen_hsa_insns_for_load (dest, rhs, TREE_TYPE (lhs), hbb);
2862 else if (TREE_CODE (rhs) == SSA_NAME
2863 || (is_gimple_min_invariant (rhs) && TREE_CODE (rhs) != STRING_CST))
2865 /* Store to memory. */
2866 hsa_op_base *src = hsa_reg_or_immed_for_gimple_op (rhs, hbb);
2867 if (hsa_seen_error ())
2868 return;
2870 gen_hsa_insns_for_store (lhs, src, hbb);
2872 else
2874 BrigAlignment8_t lhs_align;
2875 hsa_op_address *addr_lhs = gen_hsa_addr_with_align (lhs, hbb,
2876 &lhs_align);
2878 if (TREE_CODE (rhs) == CONSTRUCTOR)
2879 gen_hsa_ctor_assignment (addr_lhs, rhs, hbb, lhs_align);
2880 else
2882 BrigAlignment8_t rhs_align;
2883 hsa_op_address *addr_rhs = gen_hsa_addr_with_align (rhs, hbb,
2884 &rhs_align);
2886 unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs)));
2887 gen_hsa_memory_copy (hbb, addr_lhs, addr_rhs, size,
2888 MIN (lhs_align, rhs_align));
2893 /* Prepend before INSN a load from spill symbol of SPILL_REG. Return the
2894 register into which we loaded. If this required another register to convert
2895 from a B1 type, return it in *PTMP2, otherwise store NULL into it. We
2896 assume we are out of SSA so the returned register does not have its
2897 definition set. */
2899 hsa_op_reg *
2900 hsa_spill_in (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2)
2902 hsa_symbol *spill_sym = spill_reg->m_spill_sym;
2903 hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type);
2904 hsa_op_address *addr = new hsa_op_address (spill_sym);
2906 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, spill_sym->m_type,
2907 reg, addr);
2908 hsa_insert_insn_before (mem, insn);
2910 *ptmp2 = NULL;
2911 if (spill_reg->m_type == BRIG_TYPE_B1)
2913 hsa_insn_basic *cvtinsn;
2914 *ptmp2 = reg;
2915 reg = new hsa_op_reg (spill_reg->m_type);
2917 cvtinsn = new hsa_insn_cvt (reg, *ptmp2);
2918 hsa_insert_insn_before (cvtinsn, insn);
2920 return reg;
2923 /* Append after INSN a store to spill symbol of SPILL_REG. Return the register
2924 from which we stored. If this required another register to convert to a B1
2925 type, return it in *PTMP2, otherwise store NULL into it. We assume we are
2926 out of SSA so the returned register does not have its use updated. */
2928 hsa_op_reg *
2929 hsa_spill_out (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2)
2931 hsa_symbol *spill_sym = spill_reg->m_spill_sym;
2932 hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type);
2933 hsa_op_address *addr = new hsa_op_address (spill_sym);
2934 hsa_op_reg *returnreg;
2936 *ptmp2 = NULL;
2937 returnreg = reg;
2938 if (spill_reg->m_type == BRIG_TYPE_B1)
2940 hsa_insn_basic *cvtinsn;
2941 *ptmp2 = new hsa_op_reg (spill_sym->m_type);
2942 reg->m_type = spill_reg->m_type;
2944 cvtinsn = new hsa_insn_cvt (*ptmp2, returnreg);
2945 hsa_append_insn_after (cvtinsn, insn);
2946 insn = cvtinsn;
2947 reg = *ptmp2;
2950 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, spill_sym->m_type, reg,
2951 addr);
2952 hsa_append_insn_after (mem, insn);
2953 return returnreg;
2956 /* Generate a comparison instruction that will compare LHS and RHS with
2957 comparison specified by CODE and put result into register DEST. DEST has to
2958 have its type set already but must not have its definition set yet.
2959 Generated instructions will be added to HBB. */
2961 static void
2962 gen_hsa_cmp_insn_from_gimple (enum tree_code code, tree lhs, tree rhs,
2963 hsa_op_reg *dest, hsa_bb *hbb)
2965 BrigCompareOperation8_t compare;
2967 switch (code)
2969 case LT_EXPR:
2970 compare = BRIG_COMPARE_LT;
2971 break;
2972 case LE_EXPR:
2973 compare = BRIG_COMPARE_LE;
2974 break;
2975 case GT_EXPR:
2976 compare = BRIG_COMPARE_GT;
2977 break;
2978 case GE_EXPR:
2979 compare = BRIG_COMPARE_GE;
2980 break;
2981 case EQ_EXPR:
2982 compare = BRIG_COMPARE_EQ;
2983 break;
2984 case NE_EXPR:
2985 compare = BRIG_COMPARE_NE;
2986 break;
2987 case UNORDERED_EXPR:
2988 compare = BRIG_COMPARE_NAN;
2989 break;
2990 case ORDERED_EXPR:
2991 compare = BRIG_COMPARE_NUM;
2992 break;
2993 case UNLT_EXPR:
2994 compare = BRIG_COMPARE_LTU;
2995 break;
2996 case UNLE_EXPR:
2997 compare = BRIG_COMPARE_LEU;
2998 break;
2999 case UNGT_EXPR:
3000 compare = BRIG_COMPARE_GTU;
3001 break;
3002 case UNGE_EXPR:
3003 compare = BRIG_COMPARE_GEU;
3004 break;
3005 case UNEQ_EXPR:
3006 compare = BRIG_COMPARE_EQU;
3007 break;
3008 case LTGT_EXPR:
3009 compare = BRIG_COMPARE_NEU;
3010 break;
3012 default:
3013 HSA_SORRY_ATV (EXPR_LOCATION (lhs),
3014 "support for HSA does not implement comparison tree "
3015 "code %s\n", get_tree_code_name (code));
3016 return;
3019 /* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer)
3020 as a result of comparison. */
3022 BrigType16_t dest_type = hsa_type_integer_p (dest->m_type)
3023 ? (BrigType16_t) BRIG_TYPE_B1 : dest->m_type;
3025 hsa_insn_cmp *cmp = new hsa_insn_cmp (compare, dest_type);
3026 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (lhs, hbb);
3027 cmp->set_op (1, op1->extend_int_to_32bit (hbb));
3028 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs, hbb);
3029 cmp->set_op (2, op2->extend_int_to_32bit (hbb));
3031 hbb->append_insn (cmp);
3032 cmp->set_output_in_type (dest, 0, hbb);
3035 /* Generate an unary instruction with OPCODE and append it to a basic block
3036 HBB. The instruction uses DEST as a destination and OP1
3037 as a single operand. */
3039 static void
3040 gen_hsa_unary_operation (BrigOpcode opcode, hsa_op_reg *dest,
3041 hsa_op_with_type *op1, hsa_bb *hbb)
3043 gcc_checking_assert (dest);
3044 hsa_insn_basic *insn;
3046 if (opcode == BRIG_OPCODE_MOV && hsa_needs_cvt (dest->m_type, op1->m_type))
3048 insn = new hsa_insn_cvt (dest, op1);
3049 hbb->append_insn (insn);
3050 return;
3053 op1 = op1->extend_int_to_32bit (hbb);
3054 if (opcode == BRIG_OPCODE_FIRSTBIT || opcode == BRIG_OPCODE_LASTBIT)
3056 BrigType16_t srctype = hsa_type_integer_p (op1->m_type) ? op1->m_type
3057 : hsa_unsigned_type_for_type (op1->m_type);
3058 insn = new hsa_insn_srctype (2, opcode, BRIG_TYPE_U32, srctype, NULL,
3059 op1);
3061 else
3063 BrigType16_t optype = hsa_extend_inttype_to_32bit (dest->m_type);
3064 insn = new hsa_insn_basic (2, opcode, optype, NULL, op1);
3066 if (opcode == BRIG_OPCODE_MOV)
3067 hsa_fixup_mov_insn_type (insn);
3068 else if (opcode == BRIG_OPCODE_ABS || opcode == BRIG_OPCODE_NEG)
3070 /* ABS and NEG only exist in _s form :-/ */
3071 if (insn->m_type == BRIG_TYPE_U32)
3072 insn->m_type = BRIG_TYPE_S32;
3073 else if (insn->m_type == BRIG_TYPE_U64)
3074 insn->m_type = BRIG_TYPE_S64;
3078 hbb->append_insn (insn);
3079 insn->set_output_in_type (dest, 0, hbb);
3082 /* Generate a binary instruction with OPCODE and append it to a basic block
3083 HBB. The instruction uses DEST as a destination and operands OP1
3084 and OP2. */
3086 static void
3087 gen_hsa_binary_operation (int opcode, hsa_op_reg *dest,
3088 hsa_op_with_type *op1, hsa_op_with_type *op2,
3089 hsa_bb *hbb)
3091 gcc_checking_assert (dest);
3093 BrigType16_t optype = hsa_extend_inttype_to_32bit (dest->m_type);
3094 op1 = op1->extend_int_to_32bit (hbb);
3095 op2 = op2->extend_int_to_32bit (hbb);
3097 if ((opcode == BRIG_OPCODE_SHL || opcode == BRIG_OPCODE_SHR)
3098 && is_a <hsa_op_immed *> (op2))
3100 hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2);
3101 i->set_type (BRIG_TYPE_U32);
3103 if ((opcode == BRIG_OPCODE_OR
3104 || opcode == BRIG_OPCODE_XOR
3105 || opcode == BRIG_OPCODE_AND)
3106 && is_a <hsa_op_immed *> (op2))
3108 hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2);
3109 i->set_type (hsa_unsigned_type_for_type (i->m_type));
3112 hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, optype, NULL,
3113 op1, op2);
3114 hbb->append_insn (insn);
3115 insn->set_output_in_type (dest, 0, hbb);
3118 /* Generate HSA instructions for a single assignment. HBB is the basic block
3119 they will be appended to. */
3121 static void
3122 gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
3124 tree_code code = gimple_assign_rhs_code (assign);
3125 gimple_rhs_class rhs_class = get_gimple_rhs_class (gimple_expr_code (assign));
3127 tree lhs = gimple_assign_lhs (assign);
3128 tree rhs1 = gimple_assign_rhs1 (assign);
3129 tree rhs2 = gimple_assign_rhs2 (assign);
3130 tree rhs3 = gimple_assign_rhs3 (assign);
3132 BrigOpcode opcode;
3134 switch (code)
3136 CASE_CONVERT:
3137 case FLOAT_EXPR:
3138 /* The opcode is changed to BRIG_OPCODE_CVT if BRIG types
3139 needs a conversion. */
3140 opcode = BRIG_OPCODE_MOV;
3141 break;
3143 case PLUS_EXPR:
3144 case POINTER_PLUS_EXPR:
3145 opcode = BRIG_OPCODE_ADD;
3146 break;
3147 case MINUS_EXPR:
3148 opcode = BRIG_OPCODE_SUB;
3149 break;
3150 case MULT_EXPR:
3151 opcode = BRIG_OPCODE_MUL;
3152 break;
3153 case MULT_HIGHPART_EXPR:
3154 opcode = BRIG_OPCODE_MULHI;
3155 break;
3156 case RDIV_EXPR:
3157 case TRUNC_DIV_EXPR:
3158 case EXACT_DIV_EXPR:
3159 opcode = BRIG_OPCODE_DIV;
3160 break;
3161 case CEIL_DIV_EXPR:
3162 case FLOOR_DIV_EXPR:
3163 case ROUND_DIV_EXPR:
3164 HSA_SORRY_AT (gimple_location (assign),
3165 "support for HSA does not implement CEIL_DIV_EXPR, "
3166 "FLOOR_DIV_EXPR or ROUND_DIV_EXPR");
3167 return;
3168 case TRUNC_MOD_EXPR:
3169 opcode = BRIG_OPCODE_REM;
3170 break;
3171 case CEIL_MOD_EXPR:
3172 case FLOOR_MOD_EXPR:
3173 case ROUND_MOD_EXPR:
3174 HSA_SORRY_AT (gimple_location (assign),
3175 "support for HSA does not implement CEIL_MOD_EXPR, "
3176 "FLOOR_MOD_EXPR or ROUND_MOD_EXPR");
3177 return;
3178 case NEGATE_EXPR:
3179 opcode = BRIG_OPCODE_NEG;
3180 break;
3181 case FMA_EXPR:
3182 /* There is a native HSA instruction for scalar FMAs but not for vector
3183 ones. */
3184 if (TREE_CODE (TREE_TYPE (lhs)) == VECTOR_TYPE)
3186 hsa_op_reg *dest
3187 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3188 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3189 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3190 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
3191 hsa_op_reg *tmp = new hsa_op_reg (dest->m_type);
3192 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp, op1, op2, hbb);
3193 gen_hsa_binary_operation (BRIG_OPCODE_ADD, dest, tmp, op3, hbb);
3194 return;
3196 opcode = BRIG_OPCODE_MAD;
3197 break;
3198 case MIN_EXPR:
3199 opcode = BRIG_OPCODE_MIN;
3200 break;
3201 case MAX_EXPR:
3202 opcode = BRIG_OPCODE_MAX;
3203 break;
3204 case ABS_EXPR:
3205 opcode = BRIG_OPCODE_ABS;
3206 break;
3207 case LSHIFT_EXPR:
3208 opcode = BRIG_OPCODE_SHL;
3209 break;
3210 case RSHIFT_EXPR:
3211 opcode = BRIG_OPCODE_SHR;
3212 break;
3213 case LROTATE_EXPR:
3214 case RROTATE_EXPR:
3216 hsa_insn_basic *insn = NULL;
3217 int code1 = code == LROTATE_EXPR ? BRIG_OPCODE_SHL : BRIG_OPCODE_SHR;
3218 int code2 = code != LROTATE_EXPR ? BRIG_OPCODE_SHL : BRIG_OPCODE_SHR;
3219 BrigType16_t btype = hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
3220 true);
3222 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3223 hsa_op_reg *op1 = new hsa_op_reg (btype);
3224 hsa_op_reg *op2 = new hsa_op_reg (btype);
3225 hsa_op_with_type *shift1 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3227 tree type = TREE_TYPE (rhs2);
3228 unsigned HOST_WIDE_INT bitsize = TREE_INT_CST_LOW (TYPE_SIZE (type));
3230 hsa_op_with_type *shift2 = NULL;
3231 if (TREE_CODE (rhs2) == INTEGER_CST)
3232 shift2 = new hsa_op_immed (bitsize - tree_to_uhwi (rhs2),
3233 BRIG_TYPE_U32);
3234 else if (TREE_CODE (rhs2) == SSA_NAME)
3236 hsa_op_reg *s = hsa_cfun->reg_for_gimple_ssa (rhs2);
3237 s = as_a <hsa_op_reg *> (s->extend_int_to_32bit (hbb));
3238 hsa_op_reg *d = new hsa_op_reg (s->m_type);
3239 hsa_op_immed *size_imm = new hsa_op_immed (bitsize, BRIG_TYPE_U32);
3241 insn = new hsa_insn_basic (3, BRIG_OPCODE_SUB, d->m_type,
3242 d, s, size_imm);
3243 hbb->append_insn (insn);
3245 shift2 = d;
3247 else
3248 gcc_unreachable ();
3250 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3251 gen_hsa_binary_operation (code1, op1, src, shift1, hbb);
3252 gen_hsa_binary_operation (code2, op2, src, shift2, hbb);
3253 gen_hsa_binary_operation (BRIG_OPCODE_OR, dest, op1, op2, hbb);
3255 return;
3257 case BIT_IOR_EXPR:
3258 opcode = BRIG_OPCODE_OR;
3259 break;
3260 case BIT_XOR_EXPR:
3261 opcode = BRIG_OPCODE_XOR;
3262 break;
3263 case BIT_AND_EXPR:
3264 opcode = BRIG_OPCODE_AND;
3265 break;
3266 case BIT_NOT_EXPR:
3267 opcode = BRIG_OPCODE_NOT;
3268 break;
3269 case FIX_TRUNC_EXPR:
3271 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3272 hsa_op_with_type *v = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3274 if (hsa_needs_cvt (dest->m_type, v->m_type))
3276 hsa_op_reg *tmp = new hsa_op_reg (v->m_type);
3278 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_TRUNC,
3279 tmp->m_type, tmp, v);
3280 hbb->append_insn (insn);
3282 hsa_insn_basic *cvtinsn = new hsa_insn_cvt (dest, tmp);
3283 hbb->append_insn (cvtinsn);
3285 else
3287 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_TRUNC,
3288 dest->m_type, dest, v);
3289 hbb->append_insn (insn);
3292 return;
3294 opcode = BRIG_OPCODE_TRUNC;
3295 break;
3297 case LT_EXPR:
3298 case LE_EXPR:
3299 case GT_EXPR:
3300 case GE_EXPR:
3301 case EQ_EXPR:
3302 case NE_EXPR:
3303 case UNORDERED_EXPR:
3304 case ORDERED_EXPR:
3305 case UNLT_EXPR:
3306 case UNLE_EXPR:
3307 case UNGT_EXPR:
3308 case UNGE_EXPR:
3309 case UNEQ_EXPR:
3310 case LTGT_EXPR:
3312 hsa_op_reg *dest
3313 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3315 gen_hsa_cmp_insn_from_gimple (code, rhs1, rhs2, dest, hbb);
3316 return;
3318 case COND_EXPR:
3320 hsa_op_reg *dest
3321 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3322 hsa_op_with_type *ctrl = NULL;
3323 tree cond = rhs1;
3325 if (CONSTANT_CLASS_P (cond) || TREE_CODE (cond) == SSA_NAME)
3326 ctrl = hsa_reg_or_immed_for_gimple_op (cond, hbb);
3327 else
3329 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_B1);
3331 gen_hsa_cmp_insn_from_gimple (TREE_CODE (cond),
3332 TREE_OPERAND (cond, 0),
3333 TREE_OPERAND (cond, 1),
3334 r, hbb);
3336 ctrl = r;
3339 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3340 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
3341 op2 = op2->extend_int_to_32bit (hbb);
3342 op3 = op3->extend_int_to_32bit (hbb);
3344 BrigType16_t type = hsa_extend_inttype_to_32bit (dest->m_type);
3345 BrigType16_t utype = hsa_unsigned_type_for_type (type);
3346 if (is_a <hsa_op_immed *> (op2))
3347 op2->m_type = utype;
3348 if (is_a <hsa_op_immed *> (op3))
3349 op3->m_type = utype;
3351 hsa_insn_basic *insn
3352 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV,
3353 hsa_bittype_for_type (type),
3354 NULL, ctrl, op2, op3);
3356 hbb->append_insn (insn);
3357 insn->set_output_in_type (dest, 0, hbb);
3358 return;
3360 case COMPLEX_EXPR:
3362 hsa_op_reg *dest
3363 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3364 hsa_op_with_type *rhs1_reg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3365 rhs1_reg = rhs1_reg->extend_int_to_32bit (hbb);
3366 hsa_op_with_type *rhs2_reg = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3367 rhs2_reg = rhs2_reg->extend_int_to_32bit (hbb);
3369 if (hsa_seen_error ())
3370 return;
3372 BrigType16_t src_type = hsa_bittype_for_type (rhs1_reg->m_type);
3373 rhs1_reg = rhs1_reg->get_in_type (src_type, hbb);
3374 rhs2_reg = rhs2_reg->get_in_type (src_type, hbb);
3376 hsa_insn_packed *insn
3377 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type, src_type,
3378 dest, rhs1_reg, rhs2_reg);
3379 hbb->append_insn (insn);
3381 return;
3383 default:
3384 /* Implement others as we come across them. */
3385 HSA_SORRY_ATV (gimple_location (assign),
3386 "support for HSA does not implement operation %s",
3387 get_tree_code_name (code));
3388 return;
3392 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3393 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3394 hsa_op_with_type *op2
3395 = rhs2 ? hsa_reg_or_immed_for_gimple_op (rhs2, hbb) : NULL;
3397 if (hsa_seen_error ())
3398 return;
3400 switch (rhs_class)
3402 case GIMPLE_TERNARY_RHS:
3404 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
3405 op3 = op3->extend_int_to_32bit (hbb);
3406 hsa_insn_basic *insn = new hsa_insn_basic (4, opcode, dest->m_type, dest,
3407 op1, op2, op3);
3408 hbb->append_insn (insn);
3410 return;
3412 case GIMPLE_BINARY_RHS:
3413 gen_hsa_binary_operation (opcode, dest, op1, op2, hbb);
3414 break;
3416 case GIMPLE_UNARY_RHS:
3417 gen_hsa_unary_operation (opcode, dest, op1, hbb);
3418 break;
3419 default:
3420 gcc_unreachable ();
3424 /* Generate HSA instructions for a given gimple condition statement COND.
3425 Instructions will be appended to HBB, which also needs to be the
3426 corresponding structure to the basic_block of COND. */
3428 static void
3429 gen_hsa_insns_for_cond_stmt (gimple *cond, hsa_bb *hbb)
3431 hsa_op_reg *ctrl = new hsa_op_reg (BRIG_TYPE_B1);
3432 hsa_insn_cbr *cbr;
3434 gen_hsa_cmp_insn_from_gimple (gimple_cond_code (cond),
3435 gimple_cond_lhs (cond),
3436 gimple_cond_rhs (cond),
3437 ctrl, hbb);
3439 cbr = new hsa_insn_cbr (ctrl);
3440 hbb->append_insn (cbr);
3443 /* Maximum number of elements in a jump table for an HSA SBR instruction. */
3445 #define HSA_MAXIMUM_SBR_LABELS 16
3447 /* Return lowest value of a switch S that is handled in a non-default
3448 label. */
3450 static tree
3451 get_switch_low (gswitch *s)
3453 unsigned labels = gimple_switch_num_labels (s);
3454 gcc_checking_assert (labels >= 1);
3456 return CASE_LOW (gimple_switch_label (s, 1));
3459 /* Return highest value of a switch S that is handled in a non-default
3460 label. */
3462 static tree
3463 get_switch_high (gswitch *s)
3465 unsigned labels = gimple_switch_num_labels (s);
3467 /* Compare last label to maximum number of labels. */
3468 tree label = gimple_switch_label (s, labels - 1);
3469 tree low = CASE_LOW (label);
3470 tree high = CASE_HIGH (label);
3472 return high != NULL_TREE ? high : low;
3475 static tree
3476 get_switch_size (gswitch *s)
3478 return int_const_binop (MINUS_EXPR, get_switch_high (s), get_switch_low (s));
3481 /* Generate HSA instructions for a given gimple switch.
3482 Instructions will be appended to HBB. */
3484 static void
3485 gen_hsa_insns_for_switch_stmt (gswitch *s, hsa_bb *hbb)
3487 gimple_stmt_iterator it = gsi_for_stmt (s);
3488 gsi_prev (&it);
3490 /* Create preambule that verifies that index - lowest_label >= 0. */
3491 edge e = split_block (hbb->m_bb, gsi_stmt (it));
3492 e->flags &= ~EDGE_FALLTHRU;
3493 e->flags |= EDGE_TRUE_VALUE;
3495 function *func = DECL_STRUCT_FUNCTION (current_function_decl);
3496 tree index_tree = gimple_switch_index (s);
3497 tree lowest = get_switch_low (s);
3498 tree highest = get_switch_high (s);
3500 hsa_op_reg *index = hsa_cfun->reg_for_gimple_ssa (index_tree);
3501 index = as_a <hsa_op_reg *> (index->extend_int_to_32bit (hbb));
3503 hsa_op_reg *cmp1_reg = new hsa_op_reg (BRIG_TYPE_B1);
3504 hsa_op_immed *cmp1_immed = new hsa_op_immed (lowest, true);
3505 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_GE, cmp1_reg->m_type,
3506 cmp1_reg, index, cmp1_immed));
3508 hsa_op_reg *cmp2_reg = new hsa_op_reg (BRIG_TYPE_B1);
3509 hsa_op_immed *cmp2_immed = new hsa_op_immed (highest, true);
3510 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_LE, cmp2_reg->m_type,
3511 cmp2_reg, index, cmp2_immed));
3513 hsa_op_reg *cmp_reg = new hsa_op_reg (BRIG_TYPE_B1);
3514 hbb->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_AND, cmp_reg->m_type,
3515 cmp_reg, cmp1_reg, cmp2_reg));
3517 hbb->append_insn (new hsa_insn_cbr (cmp_reg));
3519 tree default_label = gimple_switch_default_label (s);
3520 basic_block default_label_bb = label_to_block_fn (func,
3521 CASE_LABEL (default_label));
3523 if (!gimple_seq_empty_p (phi_nodes (default_label_bb)))
3525 default_label_bb = split_edge (find_edge (e->dest, default_label_bb));
3526 hsa_init_new_bb (default_label_bb);
3529 make_edge (e->src, default_label_bb, EDGE_FALSE_VALUE);
3531 hsa_cfun->m_modified_cfg = true;
3533 /* Basic block with the SBR instruction. */
3534 hbb = hsa_init_new_bb (e->dest);
3536 hsa_op_reg *sub_index = new hsa_op_reg (index->m_type);
3537 hbb->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB, sub_index->m_type,
3538 sub_index, index,
3539 new hsa_op_immed (lowest, true)));
3541 hsa_op_base *tmp = sub_index->get_in_type (BRIG_TYPE_U64, hbb);
3542 sub_index = as_a <hsa_op_reg *> (tmp);
3543 unsigned labels = gimple_switch_num_labels (s);
3544 unsigned HOST_WIDE_INT size = tree_to_uhwi (get_switch_size (s));
3546 hsa_insn_sbr *sbr = new hsa_insn_sbr (sub_index, size + 1);
3548 /* Prepare array with default label destination. */
3549 for (unsigned HOST_WIDE_INT i = 0; i <= size; i++)
3550 sbr->m_jump_table.safe_push (default_label_bb);
3552 /* Iterate all labels and fill up the jump table. */
3553 for (unsigned i = 1; i < labels; i++)
3555 tree label = gimple_switch_label (s, i);
3556 basic_block bb = label_to_block_fn (func, CASE_LABEL (label));
3558 unsigned HOST_WIDE_INT sub_low
3559 = tree_to_uhwi (int_const_binop (MINUS_EXPR, CASE_LOW (label), lowest));
3561 unsigned HOST_WIDE_INT sub_high = sub_low;
3562 tree high = CASE_HIGH (label);
3563 if (high != NULL)
3564 sub_high = tree_to_uhwi (int_const_binop (MINUS_EXPR, high, lowest));
3566 for (unsigned HOST_WIDE_INT j = sub_low; j <= sub_high; j++)
3567 sbr->m_jump_table[j] = bb;
3570 hbb->append_insn (sbr);
3573 /* Verify that the function DECL can be handled by HSA. */
3575 static void
3576 verify_function_arguments (tree decl)
3578 tree type = TREE_TYPE (decl);
3579 if (DECL_STATIC_CHAIN (decl))
3581 HSA_SORRY_ATV (EXPR_LOCATION (decl),
3582 "HSA does not support nested functions: %qD", decl);
3583 return;
3585 else if (!TYPE_ARG_TYPES (type) || stdarg_p (type))
3587 HSA_SORRY_ATV (EXPR_LOCATION (decl),
3588 "HSA does not support functions with variadic arguments "
3589 "(or unknown return type): %qD", decl);
3590 return;
3594 /* Return BRIG type for FORMAL_ARG_TYPE. If the formal argument type is NULL,
3595 return ACTUAL_ARG_TYPE. */
3597 static BrigType16_t
3598 get_format_argument_type (tree formal_arg_type, BrigType16_t actual_arg_type)
3600 if (formal_arg_type == NULL)
3601 return actual_arg_type;
3603 BrigType16_t decl_type
3604 = hsa_type_for_scalar_tree_type (formal_arg_type, false);
3605 return mem_type_for_type (decl_type);
3608 /* Generate HSA instructions for a direct call instruction.
3609 Instructions will be appended to HBB, which also needs to be the
3610 corresponding structure to the basic_block of STMT.
3611 If ASSIGN_LHS is false, do not copy HSA function result argument into the
3612 corresponding HSA representation of the gimple statement LHS. */
3614 static void
3615 gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb,
3616 bool assign_lhs = true)
3618 tree decl = gimple_call_fndecl (stmt);
3619 verify_function_arguments (decl);
3620 if (hsa_seen_error ())
3621 return;
3623 hsa_insn_call *call_insn = new hsa_insn_call (decl);
3624 hsa_cfun->m_called_functions.safe_push (call_insn->m_called_function);
3626 /* Argument block start. */
3627 hsa_insn_arg_block *arg_start
3628 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn);
3629 hbb->append_insn (arg_start);
3631 tree parm_type_chain = TYPE_ARG_TYPES (gimple_call_fntype (stmt));
3633 /* Preparation of arguments that will be passed to function. */
3634 const unsigned args = gimple_call_num_args (stmt);
3635 for (unsigned i = 0; i < args; ++i)
3637 tree parm = gimple_call_arg (stmt, (int)i);
3638 tree parm_decl_type = parm_type_chain != NULL_TREE
3639 ? TREE_VALUE (parm_type_chain) : NULL_TREE;
3640 hsa_op_address *addr;
3642 if (AGGREGATE_TYPE_P (TREE_TYPE (parm)))
3644 addr = gen_hsa_addr_for_arg (TREE_TYPE (parm), i);
3645 BrigAlignment8_t align;
3646 hsa_op_address *src = gen_hsa_addr_with_align (parm, hbb, &align);
3647 gen_hsa_memory_copy (hbb, addr, src,
3648 addr->m_symbol->total_byte_size (), align);
3650 else
3652 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (parm, hbb);
3654 if (parm_decl_type != NULL && AGGREGATE_TYPE_P (parm_decl_type))
3656 HSA_SORRY_AT (gimple_location (stmt),
3657 "support for HSA does not implement an aggregate "
3658 "formal argument in a function call, while actual "
3659 "argument is not an aggregate");
3660 return;
3663 BrigType16_t formal_arg_type
3664 = get_format_argument_type (parm_decl_type, src->m_type);
3665 if (hsa_seen_error ())
3666 return;
3668 if (src->m_type != formal_arg_type)
3669 src = src->get_in_type (formal_arg_type, hbb);
3671 addr
3672 = gen_hsa_addr_for_arg (parm_decl_type != NULL_TREE ?
3673 parm_decl_type: TREE_TYPE (parm), i);
3674 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, formal_arg_type,
3675 src, addr);
3677 hbb->append_insn (mem);
3680 call_insn->m_input_args.safe_push (addr->m_symbol);
3681 if (parm_type_chain)
3682 parm_type_chain = TREE_CHAIN (parm_type_chain);
3685 call_insn->m_args_code_list = new hsa_op_code_list (args);
3686 hbb->append_insn (call_insn);
3688 tree result_type = TREE_TYPE (TREE_TYPE (decl));
3690 tree result = gimple_call_lhs (stmt);
3691 hsa_insn_mem *result_insn = NULL;
3692 if (!VOID_TYPE_P (result_type))
3694 hsa_op_address *addr = gen_hsa_addr_for_arg (result_type, -1);
3696 /* Even if result of a function call is unused, we have to emit
3697 declaration for the result. */
3698 if (result && assign_lhs)
3700 tree lhs_type = TREE_TYPE (result);
3702 if (hsa_seen_error ())
3703 return;
3705 if (AGGREGATE_TYPE_P (lhs_type))
3707 BrigAlignment8_t align;
3708 hsa_op_address *result_addr
3709 = gen_hsa_addr_with_align (result, hbb, &align);
3710 gen_hsa_memory_copy (hbb, result_addr, addr,
3711 addr->m_symbol->total_byte_size (), align);
3713 else
3715 BrigType16_t mtype
3716 = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type,
3717 false));
3719 hsa_op_reg *dst = hsa_cfun->reg_for_gimple_ssa (result);
3720 result_insn = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, dst, addr);
3721 hbb->append_insn (result_insn);
3725 call_insn->m_output_arg = addr->m_symbol;
3726 call_insn->m_result_code_list = new hsa_op_code_list (1);
3728 else
3730 if (result)
3732 HSA_SORRY_AT (gimple_location (stmt),
3733 "support for HSA does not implement an assignment of "
3734 "return value from a void function");
3735 return;
3738 call_insn->m_result_code_list = new hsa_op_code_list (0);
3741 /* Argument block end. */
3742 hsa_insn_arg_block *arg_end
3743 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn);
3744 hbb->append_insn (arg_end);
3747 /* Generate HSA instructions for a direct call of an internal fn.
3748 Instructions will be appended to HBB, which also needs to be the
3749 corresponding structure to the basic_block of STMT. */
3751 static void
3752 gen_hsa_insns_for_call_of_internal_fn (gimple *stmt, hsa_bb *hbb)
3754 tree lhs = gimple_call_lhs (stmt);
3755 if (!lhs)
3756 return;
3758 tree lhs_type = TREE_TYPE (lhs);
3759 tree rhs1 = gimple_call_arg (stmt, 0);
3760 tree rhs1_type = TREE_TYPE (rhs1);
3761 enum internal_fn fn = gimple_call_internal_fn (stmt);
3762 hsa_internal_fn *ifn
3763 = new hsa_internal_fn (fn, tree_to_uhwi (TYPE_SIZE (rhs1_type)));
3764 hsa_insn_call *call_insn = new hsa_insn_call (ifn);
3766 gcc_checking_assert (FLOAT_TYPE_P (rhs1_type));
3768 if (!hsa_emitted_internal_decls->find (call_insn->m_called_internal_fn))
3769 hsa_cfun->m_called_internal_fns.safe_push (call_insn->m_called_internal_fn);
3771 hsa_insn_arg_block *arg_start
3772 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn);
3773 hbb->append_insn (arg_start);
3775 unsigned num_args = gimple_call_num_args (stmt);
3777 /* Function arguments. */
3778 for (unsigned i = 0; i < num_args; i++)
3780 tree parm = gimple_call_arg (stmt, (int)i);
3781 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (parm, hbb);
3783 hsa_op_address *addr = gen_hsa_addr_for_arg (TREE_TYPE (parm), i);
3784 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, src->m_type,
3785 src, addr);
3787 call_insn->m_input_args.safe_push (addr->m_symbol);
3788 hbb->append_insn (mem);
3791 call_insn->m_args_code_list = new hsa_op_code_list (num_args);
3792 hbb->append_insn (call_insn);
3794 /* Assign returned value. */
3795 hsa_op_address *addr = gen_hsa_addr_for_arg (lhs_type, -1);
3797 call_insn->m_output_arg = addr->m_symbol;
3798 call_insn->m_result_code_list = new hsa_op_code_list (1);
3800 /* Argument block end. */
3801 hsa_insn_arg_block *arg_end
3802 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn);
3803 hbb->append_insn (arg_end);
3806 /* Generate HSA instructions for a return value instruction.
3807 Instructions will be appended to HBB, which also needs to be the
3808 corresponding structure to the basic_block of STMT. */
3810 static void
3811 gen_hsa_insns_for_return (greturn *stmt, hsa_bb *hbb)
3813 tree retval = gimple_return_retval (stmt);
3814 if (retval)
3816 hsa_op_address *addr = new hsa_op_address (hsa_cfun->m_output_arg);
3818 if (AGGREGATE_TYPE_P (TREE_TYPE (retval)))
3820 BrigAlignment8_t align;
3821 hsa_op_address *retval_addr = gen_hsa_addr_with_align (retval, hbb,
3822 &align);
3823 gen_hsa_memory_copy (hbb, addr, retval_addr,
3824 hsa_cfun->m_output_arg->total_byte_size (),
3825 align);
3827 else
3829 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (retval),
3830 false);
3831 BrigType16_t mtype = mem_type_for_type (t);
3833 /* Store of return value. */
3834 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (retval, hbb);
3835 src = src->get_in_type (mtype, hbb);
3836 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, mtype, src,
3837 addr);
3838 hbb->append_insn (mem);
3842 /* HSAIL return instruction emission. */
3843 hsa_insn_basic *ret = new hsa_insn_basic (0, BRIG_OPCODE_RET);
3844 hbb->append_insn (ret);
3847 /* Set OP_INDEX-th operand of the instruction to DEST, as the DEST
3848 can have a different type, conversion instructions are possibly
3849 appended to HBB. */
3851 void
3852 hsa_insn_basic::set_output_in_type (hsa_op_reg *dest, unsigned op_index,
3853 hsa_bb *hbb)
3855 gcc_checking_assert (op_output_p (op_index));
3857 if (dest->m_type == m_type)
3859 set_op (op_index, dest);
3860 return;
3863 hsa_insn_basic *insn;
3864 hsa_op_reg *tmp;
3865 if (hsa_needs_cvt (dest->m_type, m_type))
3867 tmp = new hsa_op_reg (m_type);
3868 insn = new hsa_insn_cvt (dest, tmp);
3870 else if (hsa_type_bit_size (dest->m_type) == hsa_type_bit_size (m_type))
3872 /* When output, HSA registers do not really have types, only sizes, so if
3873 the sizes match, we can use the register directly. */
3874 set_op (op_index, dest);
3875 return;
3877 else
3879 tmp = new hsa_op_reg (m_type);
3880 insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type,
3881 dest, tmp->get_in_type (dest->m_type, hbb));
3882 hsa_fixup_mov_insn_type (insn);
3884 set_op (op_index, tmp);
3885 hbb->append_insn (insn);
3888 /* Generate instruction OPCODE to query a property of HSA grid along the
3889 given DIMENSION. Store result into DEST and append the instruction to
3890 HBB. */
3892 static void
3893 query_hsa_grid_dim (hsa_op_reg *dest, int opcode, hsa_op_immed *dimension,
3894 hsa_bb *hbb)
3896 hsa_insn_basic *insn = new hsa_insn_basic (2, opcode, BRIG_TYPE_U32, NULL,
3897 dimension);
3898 hbb->append_insn (insn);
3899 insn->set_output_in_type (dest, 0, hbb);
3902 /* Generate instruction OPCODE to query a property of HSA grid along the given
3903 dimension which is an immediate in first argument of STMT. Store result
3904 into the register corresponding to LHS of STMT and append the instruction to
3905 HBB. */
3907 static void
3908 query_hsa_grid_dim (gimple *stmt, int opcode, hsa_bb *hbb)
3910 tree lhs = gimple_call_lhs (dyn_cast <gcall *> (stmt));
3911 if (lhs == NULL_TREE)
3912 return;
3914 tree arg = gimple_call_arg (stmt, 0);
3915 unsigned HOST_WIDE_INT dim = 5;
3916 if (tree_fits_uhwi_p (arg))
3917 dim = tree_to_uhwi (arg);
3918 if (dim > 2)
3920 HSA_SORRY_AT (gimple_location (stmt),
3921 "HSA grid query dimension must be immediate constant 0, 1 "
3922 "or 2");
3923 return;
3926 hsa_op_immed *hdim = new hsa_op_immed (dim, (BrigKind16_t) BRIG_TYPE_U32);
3927 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3928 query_hsa_grid_dim (dest, opcode, hdim, hbb);
3931 /* Generate instruction OPCODE to query a property of HSA grid that is
3932 independent of any dimension. Store result into the register corresponding
3933 to LHS of STMT and append the instruction to HBB. */
3935 static void
3936 query_hsa_grid_nodim (gimple *stmt, BrigOpcode16_t opcode, hsa_bb *hbb)
3938 tree lhs = gimple_call_lhs (dyn_cast <gcall *> (stmt));
3939 if (lhs == NULL_TREE)
3940 return;
3941 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3942 BrigType16_t brig_type = hsa_unsigned_type_for_type (dest->m_type);
3943 hsa_insn_basic *insn = new hsa_insn_basic (1, opcode, brig_type, dest);
3944 hbb->append_insn (insn);
3947 /* Emit instructions that set hsa_num_threads according to provided VALUE.
3948 Instructions are appended to basic block HBB. */
3950 static void
3951 gen_set_num_threads (tree value, hsa_bb *hbb)
3953 hbb->append_insn (new hsa_insn_comment ("omp_set_num_threads"));
3954 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (value, hbb);
3956 src = src->get_in_type (hsa_num_threads->m_type, hbb);
3957 hsa_op_address *addr = new hsa_op_address (hsa_num_threads);
3959 hsa_insn_basic *basic
3960 = new hsa_insn_mem (BRIG_OPCODE_ST, hsa_num_threads->m_type, src, addr);
3961 hbb->append_insn (basic);
3964 /* Return byte offset of a FIELD_NAME in GOMP_hsa_kernel_dispatch which
3965 is defined in plugin-hsa.c. */
3967 static HOST_WIDE_INT
3968 get_hsa_kernel_dispatch_offset (const char *field_name)
3970 tree *hsa_kernel_dispatch_type = hsa_get_kernel_dispatch_type ();
3971 if (*hsa_kernel_dispatch_type == NULL)
3973 /* Collection of information needed for a dispatch of a kernel from a
3974 kernel. Keep in sync with libgomp's plugin-hsa.c. */
3976 *hsa_kernel_dispatch_type = make_node (RECORD_TYPE);
3977 tree id_f1 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3978 get_identifier ("queue"), ptr_type_node);
3979 DECL_CHAIN (id_f1) = NULL_TREE;
3980 tree id_f2 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3981 get_identifier ("omp_data_memory"),
3982 ptr_type_node);
3983 DECL_CHAIN (id_f2) = id_f1;
3984 tree id_f3 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3985 get_identifier ("kernarg_address"),
3986 ptr_type_node);
3987 DECL_CHAIN (id_f3) = id_f2;
3988 tree id_f4 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3989 get_identifier ("object"),
3990 uint64_type_node);
3991 DECL_CHAIN (id_f4) = id_f3;
3992 tree id_f5 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3993 get_identifier ("signal"),
3994 uint64_type_node);
3995 DECL_CHAIN (id_f5) = id_f4;
3996 tree id_f6 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3997 get_identifier ("private_segment_size"),
3998 uint32_type_node);
3999 DECL_CHAIN (id_f6) = id_f5;
4000 tree id_f7 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
4001 get_identifier ("group_segment_size"),
4002 uint32_type_node);
4003 DECL_CHAIN (id_f7) = id_f6;
4004 tree id_f8 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
4005 get_identifier ("kernel_dispatch_count"),
4006 uint64_type_node);
4007 DECL_CHAIN (id_f8) = id_f7;
4008 tree id_f9 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
4009 get_identifier ("debug"),
4010 uint64_type_node);
4011 DECL_CHAIN (id_f9) = id_f8;
4012 tree id_f10 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
4013 get_identifier ("omp_level"),
4014 uint64_type_node);
4015 DECL_CHAIN (id_f10) = id_f9;
4016 tree id_f11 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
4017 get_identifier ("children_dispatches"),
4018 ptr_type_node);
4019 DECL_CHAIN (id_f11) = id_f10;
4020 tree id_f12 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
4021 get_identifier ("omp_num_threads"),
4022 uint32_type_node);
4023 DECL_CHAIN (id_f12) = id_f11;
4026 finish_builtin_struct (*hsa_kernel_dispatch_type, "__hsa_kernel_dispatch",
4027 id_f12, NULL_TREE);
4028 TYPE_ARTIFICIAL (*hsa_kernel_dispatch_type) = 1;
4031 for (tree chain = TYPE_FIELDS (*hsa_kernel_dispatch_type);
4032 chain != NULL_TREE; chain = TREE_CHAIN (chain))
4033 if (id_equal (DECL_NAME (chain), field_name))
4034 return int_byte_position (chain);
4036 gcc_unreachable ();
4039 /* Return an HSA register that will contain number of threads for
4040 a future dispatched kernel. Instructions are added to HBB. */
4042 static hsa_op_reg *
4043 gen_num_threads_for_dispatch (hsa_bb *hbb)
4045 /* Step 1) Assign to number of threads:
4046 MIN (HSA_DEFAULT_NUM_THREADS, hsa_num_threads). */
4047 hsa_op_reg *threads = new hsa_op_reg (hsa_num_threads->m_type);
4048 hsa_op_address *addr = new hsa_op_address (hsa_num_threads);
4050 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD, threads->m_type,
4051 threads, addr));
4053 hsa_op_immed *limit = new hsa_op_immed (HSA_DEFAULT_NUM_THREADS,
4054 BRIG_TYPE_U32);
4055 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_B1);
4056 hsa_insn_cmp * cmp
4057 = new hsa_insn_cmp (BRIG_COMPARE_LT, r->m_type, r, threads, limit);
4058 hbb->append_insn (cmp);
4060 BrigType16_t btype = hsa_bittype_for_type (threads->m_type);
4061 hsa_op_reg *tmp = new hsa_op_reg (threads->m_type);
4063 hbb->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp, r,
4064 threads, limit));
4066 /* Step 2) If the number is equal to zero,
4067 return shadow->omp_num_threads. */
4068 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg ();
4070 hsa_op_reg *shadow_thread_count = new hsa_op_reg (BRIG_TYPE_U32);
4071 addr
4072 = new hsa_op_address (shadow_reg_ptr,
4073 get_hsa_kernel_dispatch_offset ("omp_num_threads"));
4074 hsa_insn_basic *basic
4075 = new hsa_insn_mem (BRIG_OPCODE_LD, shadow_thread_count->m_type,
4076 shadow_thread_count, addr);
4077 hbb->append_insn (basic);
4079 hsa_op_reg *tmp2 = new hsa_op_reg (threads->m_type);
4080 r = new hsa_op_reg (BRIG_TYPE_B1);
4081 hsa_op_immed *imm = new hsa_op_immed (0, shadow_thread_count->m_type);
4082 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_EQ, r->m_type, r, tmp, imm));
4083 hbb->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp2, r,
4084 shadow_thread_count, tmp));
4086 hsa_op_base *dest = tmp2->get_in_type (BRIG_TYPE_U16, hbb);
4088 return as_a <hsa_op_reg *> (dest);
4091 /* Build OPCODE query for all three hsa dimensions, multiply them and store the
4092 result into DEST. */
4094 static void
4095 multiply_grid_dim_characteristics (hsa_op_reg *dest, int opcode, hsa_bb *hbb)
4097 hsa_op_reg *dimx = new hsa_op_reg (BRIG_TYPE_U32);
4098 query_hsa_grid_dim (dimx, opcode,
4099 new hsa_op_immed (0, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4100 hsa_op_reg *dimy = new hsa_op_reg (BRIG_TYPE_U32);
4101 query_hsa_grid_dim (dimy, opcode,
4102 new hsa_op_immed (1, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4103 hsa_op_reg *dimz = new hsa_op_reg (BRIG_TYPE_U32);
4104 query_hsa_grid_dim (dimz, opcode,
4105 new hsa_op_immed (2, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4106 hsa_op_reg *tmp = new hsa_op_reg (dest->m_type);
4107 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp,
4108 dimx->get_in_type (dest->m_type, hbb),
4109 dimy->get_in_type (dest->m_type, hbb), hbb);
4110 gen_hsa_binary_operation (BRIG_OPCODE_MUL, dest, tmp,
4111 dimz->get_in_type (dest->m_type, hbb), hbb);
4114 /* Emit instructions that assign number of threads to lhs of gimple STMT.
4115 Instructions are appended to basic block HBB. */
4117 static void
4118 gen_get_num_threads (gimple *stmt, hsa_bb *hbb)
4120 if (gimple_call_lhs (stmt) == NULL_TREE)
4121 return;
4123 hbb->append_insn (new hsa_insn_comment ("omp_get_num_threads"));
4124 tree lhs = gimple_call_lhs (stmt);
4125 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4126 multiply_grid_dim_characteristics (dest, BRIG_OPCODE_CURRENTWORKGROUPSIZE,
4127 hbb);
4130 /* Emit instructions that assign number of teams to lhs of gimple STMT.
4131 Instructions are appended to basic block HBB. */
4133 static void
4134 gen_get_num_teams (gimple *stmt, hsa_bb *hbb)
4136 if (gimple_call_lhs (stmt) == NULL_TREE)
4137 return;
4139 hbb->append_insn (new hsa_insn_comment ("omp_get_num_teams"));
4140 tree lhs = gimple_call_lhs (stmt);
4141 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4142 multiply_grid_dim_characteristics (dest, BRIG_OPCODE_GRIDGROUPS, hbb);
4145 /* Emit instructions that assign a team number to lhs of gimple STMT.
4146 Instructions are appended to basic block HBB. */
4148 static void
4149 gen_get_team_num (gimple *stmt, hsa_bb *hbb)
4151 if (gimple_call_lhs (stmt) == NULL_TREE)
4152 return;
4154 hbb->append_insn (new hsa_insn_comment ("omp_get_team_num"));
4155 tree lhs = gimple_call_lhs (stmt);
4156 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4158 hsa_op_reg *gnum_x = new hsa_op_reg (BRIG_TYPE_U32);
4159 query_hsa_grid_dim (gnum_x, BRIG_OPCODE_GRIDGROUPS,
4160 new hsa_op_immed (0, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4161 hsa_op_reg *gnum_y = new hsa_op_reg (BRIG_TYPE_U32);
4162 query_hsa_grid_dim (gnum_y, BRIG_OPCODE_GRIDGROUPS,
4163 new hsa_op_immed (1, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4165 hsa_op_reg *gno_z = new hsa_op_reg (BRIG_TYPE_U32);
4166 query_hsa_grid_dim (gno_z, BRIG_OPCODE_WORKGROUPID,
4167 new hsa_op_immed (2, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4169 hsa_op_reg *tmp1 = new hsa_op_reg (dest->m_type);
4170 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp1,
4171 gnum_x->get_in_type (dest->m_type, hbb),
4172 gnum_y->get_in_type (dest->m_type, hbb), hbb);
4173 hsa_op_reg *tmp2 = new hsa_op_reg (dest->m_type);
4174 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp2, tmp1,
4175 gno_z->get_in_type (dest->m_type, hbb), hbb);
4177 hsa_op_reg *gno_y = new hsa_op_reg (BRIG_TYPE_U32);
4178 query_hsa_grid_dim (gno_y, BRIG_OPCODE_WORKGROUPID,
4179 new hsa_op_immed (1, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4180 hsa_op_reg *tmp3 = new hsa_op_reg (dest->m_type);
4181 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp3,
4182 gnum_x->get_in_type (dest->m_type, hbb),
4183 gno_y->get_in_type (dest->m_type, hbb), hbb);
4184 hsa_op_reg *tmp4 = new hsa_op_reg (dest->m_type);
4185 gen_hsa_binary_operation (BRIG_OPCODE_ADD, tmp4, tmp3, tmp2, hbb);
4186 hsa_op_reg *gno_x = new hsa_op_reg (BRIG_TYPE_U32);
4187 query_hsa_grid_dim (gno_x, BRIG_OPCODE_WORKGROUPID,
4188 new hsa_op_immed (0, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4189 gen_hsa_binary_operation (BRIG_OPCODE_ADD, dest, tmp4,
4190 gno_x->get_in_type (dest->m_type, hbb), hbb);
4193 /* Emit instructions that get levels-var ICV to lhs of gimple STMT.
4194 Instructions are appended to basic block HBB. */
4196 static void
4197 gen_get_level (gimple *stmt, hsa_bb *hbb)
4199 if (gimple_call_lhs (stmt) == NULL_TREE)
4200 return;
4202 hbb->append_insn (new hsa_insn_comment ("omp_get_level"));
4204 tree lhs = gimple_call_lhs (stmt);
4205 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4207 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg ();
4208 if (shadow_reg_ptr == NULL)
4210 HSA_SORRY_AT (gimple_location (stmt),
4211 "support for HSA does not implement omp_get_level called "
4212 "from a function not being inlined within a kernel");
4213 return;
4216 hsa_op_address *addr
4217 = new hsa_op_address (shadow_reg_ptr,
4218 get_hsa_kernel_dispatch_offset ("omp_level"));
4220 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64,
4221 (hsa_op_base *) NULL, addr);
4222 hbb->append_insn (mem);
4223 mem->set_output_in_type (dest, 0, hbb);
4226 /* Emit instruction that implement omp_get_max_threads of gimple STMT. */
4228 static void
4229 gen_get_max_threads (gimple *stmt, hsa_bb *hbb)
4231 tree lhs = gimple_call_lhs (stmt);
4232 if (!lhs)
4233 return;
4235 hbb->append_insn (new hsa_insn_comment ("omp_get_max_threads"));
4237 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4238 hsa_op_with_type *num_theads_reg = gen_num_threads_for_dispatch (hbb)
4239 ->get_in_type (dest->m_type, hbb);
4240 hsa_build_append_simple_mov (dest, num_theads_reg, hbb);
4243 /* Emit instructions that implement alloca builtin gimple STMT.
4244 Instructions are appended to basic block HBB. */
4246 static void
4247 gen_hsa_alloca (gcall *call, hsa_bb *hbb)
4249 tree lhs = gimple_call_lhs (call);
4250 if (lhs == NULL_TREE)
4251 return;
4253 built_in_function fn = DECL_FUNCTION_CODE (gimple_call_fndecl (call));
4255 gcc_checking_assert (ALLOCA_FUNCTION_CODE_P (fn));
4257 unsigned bit_alignment = 0;
4259 if (fn != BUILT_IN_ALLOCA)
4261 tree alignment_tree = gimple_call_arg (call, 1);
4262 if (TREE_CODE (alignment_tree) != INTEGER_CST)
4264 HSA_SORRY_ATV (gimple_location (call),
4265 "support for HSA does not implement "
4266 "__builtin_alloca_with_align with a non-constant "
4267 "alignment: %E", alignment_tree);
4270 bit_alignment = tree_to_uhwi (alignment_tree);
4273 tree rhs1 = gimple_call_arg (call, 0);
4274 hsa_op_with_type *size = hsa_reg_or_immed_for_gimple_op (rhs1, hbb)
4275 ->get_in_type (BRIG_TYPE_U32, hbb);
4276 hsa_op_with_type *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4278 hsa_op_reg *tmp
4279 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE));
4280 hsa_insn_alloca *a = new hsa_insn_alloca (tmp, size, bit_alignment);
4281 hbb->append_insn (a);
4283 hsa_insn_seg *seg
4284 = new hsa_insn_seg (BRIG_OPCODE_STOF,
4285 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT),
4286 tmp->m_type, BRIG_SEGMENT_PRIVATE, dest, tmp);
4287 hbb->append_insn (seg);
4290 /* Emit instructions that implement clrsb builtin STMT:
4291 Returns the number of leading redundant sign bits in x, i.e. the number
4292 of bits following the most significant bit that are identical to it.
4293 There are no special cases for 0 or other values.
4294 Instructions are appended to basic block HBB. */
4296 static void
4297 gen_hsa_clrsb (gcall *call, hsa_bb *hbb)
4299 tree lhs = gimple_call_lhs (call);
4300 if (lhs == NULL_TREE)
4301 return;
4303 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4304 tree rhs1 = gimple_call_arg (call, 0);
4305 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4306 arg->extend_int_to_32bit (hbb);
4307 BrigType16_t bittype = hsa_bittype_for_type (arg->m_type);
4308 unsigned bitsize = tree_to_uhwi (TYPE_SIZE (TREE_TYPE (rhs1)));
4310 /* FIRSTBIT instruction is defined just for 32 and 64-bits wide integers. */
4311 gcc_checking_assert (bitsize == 32 || bitsize == 64);
4313 /* Set true to MOST_SIG if the most significant bit is set to one. */
4314 hsa_op_immed *c = new hsa_op_immed (1ul << (bitsize - 1),
4315 hsa_uint_for_bitsize (bitsize));
4317 hsa_op_reg *and_reg = new hsa_op_reg (bittype);
4318 gen_hsa_binary_operation (BRIG_OPCODE_AND, and_reg, arg, c, hbb);
4320 hsa_op_reg *most_sign = new hsa_op_reg (BRIG_TYPE_B1);
4321 hsa_insn_cmp *cmp
4322 = new hsa_insn_cmp (BRIG_COMPARE_EQ, most_sign->m_type, most_sign,
4323 and_reg, c);
4324 hbb->append_insn (cmp);
4326 /* If the most significant bit is one, negate the input. Otherwise
4327 shift the input value to left by one bit. */
4328 hsa_op_reg *arg_neg = new hsa_op_reg (arg->m_type);
4329 gen_hsa_unary_operation (BRIG_OPCODE_NEG, arg_neg, arg, hbb);
4331 hsa_op_reg *shifted_arg = new hsa_op_reg (arg->m_type);
4332 gen_hsa_binary_operation (BRIG_OPCODE_SHL, shifted_arg, arg,
4333 new hsa_op_immed (1, BRIG_TYPE_U64), hbb);
4335 /* Assign the value that can be used for FIRSTBIT instruction according
4336 to the most significant bit. */
4337 hsa_op_reg *tmp = new hsa_op_reg (bittype);
4338 hsa_insn_basic *cmov
4339 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, bittype, tmp, most_sign,
4340 arg_neg, shifted_arg);
4341 hbb->append_insn (cmov);
4343 hsa_op_reg *leading_bits = new hsa_op_reg (BRIG_TYPE_S32);
4344 gen_hsa_unary_operation (BRIG_OPCODE_FIRSTBIT, leading_bits,
4345 tmp->get_in_type (hsa_uint_for_bitsize (bitsize),
4346 hbb), hbb);
4348 /* Set flag if the input value is equal to zero. */
4349 hsa_op_reg *is_zero = new hsa_op_reg (BRIG_TYPE_B1);
4350 cmp = new hsa_insn_cmp (BRIG_COMPARE_EQ, is_zero->m_type, is_zero, arg,
4351 new hsa_op_immed (0, arg->m_type));
4352 hbb->append_insn (cmp);
4354 /* Return the number of leading bits,
4355 or (bitsize - 1) if the input value is zero. */
4356 cmov = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, BRIG_TYPE_B32, NULL, is_zero,
4357 new hsa_op_immed (bitsize - 1, BRIG_TYPE_U32),
4358 leading_bits->get_in_type (BRIG_TYPE_B32, hbb));
4359 hbb->append_insn (cmov);
4360 cmov->set_output_in_type (dest, 0, hbb);
4363 /* Emit instructions that implement ffs builtin STMT:
4364 Returns one plus the index of the least significant 1-bit of x,
4365 or if x is zero, returns zero.
4366 Instructions are appended to basic block HBB. */
4368 static void
4369 gen_hsa_ffs (gcall *call, hsa_bb *hbb)
4371 tree lhs = gimple_call_lhs (call);
4372 if (lhs == NULL_TREE)
4373 return;
4375 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4377 tree rhs1 = gimple_call_arg (call, 0);
4378 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4379 arg = arg->extend_int_to_32bit (hbb);
4381 hsa_op_reg *tmp = new hsa_op_reg (BRIG_TYPE_U32);
4382 hsa_insn_srctype *insn = new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT,
4383 tmp->m_type, arg->m_type,
4384 tmp, arg);
4385 hbb->append_insn (insn);
4387 hsa_insn_basic *addition
4388 = new hsa_insn_basic (3, BRIG_OPCODE_ADD, tmp->m_type, NULL, tmp,
4389 new hsa_op_immed (1, tmp->m_type));
4390 hbb->append_insn (addition);
4391 addition->set_output_in_type (dest, 0, hbb);
4394 static void
4395 gen_hsa_popcount_to_dest (hsa_op_reg *dest, hsa_op_with_type *arg, hsa_bb *hbb)
4397 gcc_checking_assert (hsa_type_integer_p (arg->m_type));
4399 if (hsa_type_bit_size (arg->m_type) < 32)
4400 arg = arg->get_in_type (BRIG_TYPE_B32, hbb);
4402 BrigType16_t srctype = hsa_bittype_for_type (arg->m_type);
4403 if (!hsa_btype_p (arg->m_type))
4404 arg = arg->get_in_type (srctype, hbb);
4406 hsa_insn_srctype *popcount
4407 = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT, BRIG_TYPE_U32,
4408 srctype, NULL, arg);
4409 hbb->append_insn (popcount);
4410 popcount->set_output_in_type (dest, 0, hbb);
4413 /* Emit instructions that implement parity builtin STMT:
4414 Returns the parity of x, i.e. the number of 1-bits in x modulo 2.
4415 Instructions are appended to basic block HBB. */
4417 static void
4418 gen_hsa_parity (gcall *call, hsa_bb *hbb)
4420 tree lhs = gimple_call_lhs (call);
4421 if (lhs == NULL_TREE)
4422 return;
4424 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4425 tree rhs1 = gimple_call_arg (call, 0);
4426 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4428 hsa_op_reg *popcount = new hsa_op_reg (BRIG_TYPE_U32);
4429 gen_hsa_popcount_to_dest (popcount, arg, hbb);
4431 hsa_insn_basic *insn
4432 = new hsa_insn_basic (3, BRIG_OPCODE_REM, popcount->m_type, NULL, popcount,
4433 new hsa_op_immed (2, popcount->m_type));
4434 hbb->append_insn (insn);
4435 insn->set_output_in_type (dest, 0, hbb);
4438 /* Emit instructions that implement popcount builtin STMT.
4439 Instructions are appended to basic block HBB. */
4441 static void
4442 gen_hsa_popcount (gcall *call, hsa_bb *hbb)
4444 tree lhs = gimple_call_lhs (call);
4445 if (lhs == NULL_TREE)
4446 return;
4448 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4449 tree rhs1 = gimple_call_arg (call, 0);
4450 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4452 gen_hsa_popcount_to_dest (dest, arg, hbb);
4455 /* Emit instructions that implement DIVMOD builtin STMT.
4456 Instructions are appended to basic block HBB. */
4458 static void
4459 gen_hsa_divmod (gcall *call, hsa_bb *hbb)
4461 tree lhs = gimple_call_lhs (call);
4462 if (lhs == NULL_TREE)
4463 return;
4465 tree rhs0 = gimple_call_arg (call, 0);
4466 tree rhs1 = gimple_call_arg (call, 1);
4468 hsa_op_with_type *arg0 = hsa_reg_or_immed_for_gimple_op (rhs0, hbb);
4469 arg0 = arg0->extend_int_to_32bit (hbb);
4470 hsa_op_with_type *arg1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4471 arg1 = arg1->extend_int_to_32bit (hbb);
4473 hsa_op_reg *dest0 = new hsa_op_reg (arg0->m_type);
4474 hsa_op_reg *dest1 = new hsa_op_reg (arg1->m_type);
4476 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_DIV, dest0->m_type,
4477 dest0, arg0, arg1);
4478 hbb->append_insn (insn);
4479 insn = new hsa_insn_basic (3, BRIG_OPCODE_REM, dest1->m_type, dest1, arg0,
4480 arg1);
4481 hbb->append_insn (insn);
4483 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4484 BrigType16_t dst_type = hsa_extend_inttype_to_32bit (dest->m_type);
4485 BrigType16_t src_type = hsa_bittype_for_type (dest0->m_type);
4487 insn = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dst_type,
4488 src_type, NULL, dest0, dest1);
4489 hbb->append_insn (insn);
4490 insn->set_output_in_type (dest, 0, hbb);
4493 /* Set VALUE to a shadow kernel debug argument and append a new instruction
4494 to HBB basic block. */
4496 static void
4497 set_debug_value (hsa_bb *hbb, hsa_op_with_type *value)
4499 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg ();
4500 if (shadow_reg_ptr == NULL)
4501 return;
4503 hsa_op_address *addr
4504 = new hsa_op_address (shadow_reg_ptr,
4505 get_hsa_kernel_dispatch_offset ("debug"));
4506 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, BRIG_TYPE_U64, value,
4507 addr);
4508 hbb->append_insn (mem);
4511 void
4512 omp_simple_builtin::generate (gimple *stmt, hsa_bb *hbb)
4514 if (m_sorry)
4516 if (m_warning_message)
4517 HSA_SORRY_AT (gimple_location (stmt), m_warning_message);
4518 else
4519 HSA_SORRY_ATV (gimple_location (stmt),
4520 "Support for HSA does not implement calls to %s\n",
4521 m_name);
4523 else if (m_warning_message != NULL)
4524 warning_at (gimple_location (stmt), OPT_Whsa, m_warning_message);
4526 if (m_return_value != NULL)
4528 tree lhs = gimple_call_lhs (stmt);
4529 if (!lhs)
4530 return;
4532 hbb->append_insn (new hsa_insn_comment (m_name));
4534 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4535 hsa_op_with_type *op = m_return_value->get_in_type (dest->m_type, hbb);
4536 hsa_build_append_simple_mov (dest, op, hbb);
4540 /* If STMT is a call of a known library function, generate code to perform
4541 it and return true. */
4543 static bool
4544 gen_hsa_insns_for_known_library_call (gimple *stmt, hsa_bb *hbb)
4546 bool handled = false;
4547 const char *name = hsa_get_declaration_name (gimple_call_fndecl (stmt));
4549 char *copy = NULL;
4550 size_t len = strlen (name);
4551 if (len > 0 && name[len - 1] == '_')
4553 copy = XNEWVEC (char, len + 1);
4554 strcpy (copy, name);
4555 copy[len - 1] = '\0';
4556 name = copy;
4559 /* Handle omp_* routines. */
4560 if (strstr (name, "omp_") == name)
4562 hsa_init_simple_builtins ();
4563 omp_simple_builtin *builtin = omp_simple_builtins->get (name);
4564 if (builtin)
4566 builtin->generate (stmt, hbb);
4567 return true;
4570 handled = true;
4571 if (strcmp (name, "omp_set_num_threads") == 0)
4572 gen_set_num_threads (gimple_call_arg (stmt, 0), hbb);
4573 else if (strcmp (name, "omp_get_thread_num") == 0)
4575 hbb->append_insn (new hsa_insn_comment (name));
4576 query_hsa_grid_nodim (stmt, BRIG_OPCODE_WORKITEMFLATABSID, hbb);
4578 else if (strcmp (name, "omp_get_num_threads") == 0)
4580 hbb->append_insn (new hsa_insn_comment (name));
4581 gen_get_num_threads (stmt, hbb);
4583 else if (strcmp (name, "omp_get_num_teams") == 0)
4584 gen_get_num_teams (stmt, hbb);
4585 else if (strcmp (name, "omp_get_team_num") == 0)
4586 gen_get_team_num (stmt, hbb);
4587 else if (strcmp (name, "omp_get_level") == 0)
4588 gen_get_level (stmt, hbb);
4589 else if (strcmp (name, "omp_get_active_level") == 0)
4590 gen_get_level (stmt, hbb);
4591 else if (strcmp (name, "omp_in_parallel") == 0)
4592 gen_get_level (stmt, hbb);
4593 else if (strcmp (name, "omp_get_max_threads") == 0)
4594 gen_get_max_threads (stmt, hbb);
4595 else
4596 handled = false;
4598 if (handled)
4600 if (copy)
4601 free (copy);
4602 return true;
4606 if (strcmp (name, "__hsa_set_debug_value") == 0)
4608 handled = true;
4609 if (hsa_cfun->has_shadow_reg_p ())
4611 tree rhs1 = gimple_call_arg (stmt, 0);
4612 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4614 src = src->get_in_type (BRIG_TYPE_U64, hbb);
4615 set_debug_value (hbb, src);
4619 if (copy)
4620 free (copy);
4621 return handled;
4624 /* Helper functions to create a single unary HSA operations out of calls to
4625 builtins. OPCODE is the HSA operation to be generated. STMT is a gimple
4626 call to a builtin. HBB is the HSA BB to which the instruction should be
4627 added. Note that nothing will be created if STMT does not have a LHS. */
4629 static void
4630 gen_hsa_unaryop_for_builtin (BrigOpcode opcode, gimple *stmt, hsa_bb *hbb)
4632 tree lhs = gimple_call_lhs (stmt);
4633 if (!lhs)
4634 return;
4635 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4636 hsa_op_with_type *op
4637 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb);
4638 gen_hsa_unary_operation (opcode, dest, op, hbb);
4641 /* Helper functions to create a call to standard library if LHS of the
4642 STMT is used. HBB is the HSA BB to which the instruction should be
4643 added. */
4645 static void
4646 gen_hsa_unaryop_builtin_call (gimple *stmt, hsa_bb *hbb)
4648 tree lhs = gimple_call_lhs (stmt);
4649 if (!lhs)
4650 return;
4652 if (gimple_call_internal_p (stmt))
4653 gen_hsa_insns_for_call_of_internal_fn (stmt, hbb);
4654 else
4655 gen_hsa_insns_for_direct_call (stmt, hbb);
4658 /* Helper functions to create a single unary HSA operations out of calls to
4659 builtins (if unsafe math optimizations are enable). Otherwise, create
4660 a call to standard library function.
4661 OPCODE is the HSA operation to be generated. STMT is a gimple
4662 call to a builtin. HBB is the HSA BB to which the instruction should be
4663 added. Note that nothing will be created if STMT does not have a LHS. */
4665 static void
4666 gen_hsa_unaryop_or_call_for_builtin (BrigOpcode opcode, gimple *stmt,
4667 hsa_bb *hbb)
4669 if (flag_unsafe_math_optimizations)
4670 gen_hsa_unaryop_for_builtin (opcode, stmt, hbb);
4671 else
4672 gen_hsa_unaryop_builtin_call (stmt, hbb);
4675 /* Generate HSA address corresponding to a value VAL (as opposed to a memory
4676 reference tree), for example an SSA_NAME or an ADDR_EXPR. HBB is the HSA BB
4677 to which the instruction should be added. */
4679 static hsa_op_address *
4680 get_address_from_value (tree val, hsa_bb *hbb)
4682 switch (TREE_CODE (val))
4684 case SSA_NAME:
4686 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
4687 hsa_op_base *reg
4688 = hsa_cfun->reg_for_gimple_ssa (val)->get_in_type (addrtype, hbb);
4689 return new hsa_op_address (NULL, as_a <hsa_op_reg *> (reg), 0);
4691 case ADDR_EXPR:
4692 return gen_hsa_addr (TREE_OPERAND (val, 0), hbb);
4694 case INTEGER_CST:
4695 if (tree_fits_shwi_p (val))
4696 return new hsa_op_address (NULL, NULL, tree_to_shwi (val));
4697 /* fall-through */
4699 default:
4700 HSA_SORRY_ATV (EXPR_LOCATION (val),
4701 "support for HSA does not implement memory access to %E",
4702 val);
4703 return new hsa_op_address (NULL, NULL, 0);
4707 /* Expand assignment of a result of a string BUILTIN to DST.
4708 Size of the operation is N bytes, where instructions
4709 will be append to HBB. */
4711 static void
4712 expand_lhs_of_string_op (gimple *stmt,
4713 unsigned HOST_WIDE_INT n, hsa_bb *hbb,
4714 enum built_in_function builtin)
4716 /* If LHS is expected, we need to emit a PHI instruction. */
4717 tree lhs = gimple_call_lhs (stmt);
4718 if (!lhs)
4719 return;
4721 hsa_op_reg *lhs_reg = hsa_cfun->reg_for_gimple_ssa (lhs);
4723 hsa_op_with_type *dst_reg
4724 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb);
4725 hsa_op_with_type *tmp;
4727 switch (builtin)
4729 case BUILT_IN_MEMPCPY:
4731 tmp = new hsa_op_reg (dst_reg->m_type);
4732 hsa_insn_basic *add
4733 = new hsa_insn_basic (3, BRIG_OPCODE_ADD, tmp->m_type,
4734 tmp, dst_reg,
4735 new hsa_op_immed (n, dst_reg->m_type));
4736 hbb->append_insn (add);
4737 break;
4739 case BUILT_IN_MEMCPY:
4740 case BUILT_IN_MEMSET:
4741 tmp = dst_reg;
4742 break;
4743 default:
4744 gcc_unreachable ();
4747 hbb->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV, lhs_reg->m_type,
4748 lhs_reg, tmp));
4751 #define HSA_MEMORY_BUILTINS_LIMIT 128
4753 /* Expand a string builtin (from a gimple STMT) in a way that
4754 according to MISALIGNED_FLAG we process either direct emission
4755 (a bunch of memory load and store instructions), or we emit a function call
4756 of a library function (for instance 'memcpy'). Actually, a basic block
4757 for direct emission is just prepared, where caller is responsible
4758 for emission of corresponding instructions.
4759 All instruction are appended to HBB. */
4761 hsa_bb *
4762 expand_string_operation_builtin (gimple *stmt, hsa_bb *hbb,
4763 hsa_op_reg *misaligned_flag)
4765 edge e = split_block (hbb->m_bb, stmt);
4766 basic_block condition_bb = e->src;
4767 hbb->append_insn (new hsa_insn_cbr (misaligned_flag));
4769 /* Prepare the control flow. */
4770 edge condition_edge = EDGE_SUCC (condition_bb, 0);
4771 basic_block call_bb = split_edge (condition_edge);
4773 basic_block expanded_bb = split_edge (EDGE_SUCC (call_bb, 0));
4774 basic_block cont_bb = EDGE_SUCC (expanded_bb, 0)->dest;
4775 basic_block merge_bb = split_edge (EDGE_PRED (cont_bb, 0));
4777 condition_edge->flags &= ~EDGE_FALLTHRU;
4778 condition_edge->flags |= EDGE_TRUE_VALUE;
4779 make_edge (condition_bb, expanded_bb, EDGE_FALSE_VALUE);
4781 redirect_edge_succ (EDGE_SUCC (call_bb, 0), merge_bb);
4783 hsa_cfun->m_modified_cfg = true;
4785 hsa_init_new_bb (expanded_bb);
4787 /* Slow path: function call. */
4788 gen_hsa_insns_for_direct_call (stmt, hsa_init_new_bb (call_bb), false);
4790 return hsa_bb_for_bb (expanded_bb);
4793 /* Expand a memory copy BUILTIN (BUILT_IN_MEMCPY, BUILT_IN_MEMPCPY) from
4794 a gimple STMT and store all necessary instruction to HBB basic block. */
4796 static void
4797 expand_memory_copy (gimple *stmt, hsa_bb *hbb, enum built_in_function builtin)
4799 tree byte_size = gimple_call_arg (stmt, 2);
4801 if (!tree_fits_uhwi_p (byte_size))
4803 gen_hsa_insns_for_direct_call (stmt, hbb);
4804 return;
4807 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size);
4809 if (n > HSA_MEMORY_BUILTINS_LIMIT)
4811 gen_hsa_insns_for_direct_call (stmt, hbb);
4812 return;
4815 tree dst = gimple_call_arg (stmt, 0);
4816 tree src = gimple_call_arg (stmt, 1);
4818 hsa_op_address *dst_addr = get_address_from_value (dst, hbb);
4819 hsa_op_address *src_addr = get_address_from_value (src, hbb);
4821 /* As gen_hsa_memory_copy relies on memory alignment
4822 greater or equal to 8 bytes, we need to verify the alignment. */
4823 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
4824 hsa_op_reg *src_addr_reg = new hsa_op_reg (addrtype);
4825 hsa_op_reg *dst_addr_reg = new hsa_op_reg (addrtype);
4827 convert_addr_to_flat_segment (src_addr, src_addr_reg, hbb);
4828 convert_addr_to_flat_segment (dst_addr, dst_addr_reg, hbb);
4830 /* Process BIT OR for source and destination addresses. */
4831 hsa_op_reg *or_reg = new hsa_op_reg (addrtype);
4832 gen_hsa_binary_operation (BRIG_OPCODE_OR, or_reg, src_addr_reg,
4833 dst_addr_reg, hbb);
4835 /* Process BIT AND with 0x7 to identify the desired alignment
4836 of 8 bytes. */
4837 hsa_op_reg *masked = new hsa_op_reg (addrtype);
4839 gen_hsa_binary_operation (BRIG_OPCODE_AND, masked, or_reg,
4840 new hsa_op_immed (7, addrtype), hbb);
4842 hsa_op_reg *misaligned = new hsa_op_reg (BRIG_TYPE_B1);
4843 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE, misaligned->m_type,
4844 misaligned, masked,
4845 new hsa_op_immed (0, masked->m_type)));
4847 hsa_bb *native_impl_bb
4848 = expand_string_operation_builtin (stmt, hbb, misaligned);
4850 gen_hsa_memory_copy (native_impl_bb, dst_addr, src_addr, n, BRIG_ALIGNMENT_8);
4851 hsa_bb *merge_bb
4852 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb->m_bb, 0)->dest);
4853 expand_lhs_of_string_op (stmt, n, merge_bb, builtin);
4857 /* Expand a memory set BUILTIN (BUILT_IN_MEMSET, BUILT_IN_BZERO) from
4858 a gimple STMT and store all necessary instruction to HBB basic block.
4859 The operation set N bytes with a CONSTANT value. */
4861 static void
4862 expand_memory_set (gimple *stmt, unsigned HOST_WIDE_INT n,
4863 unsigned HOST_WIDE_INT constant, hsa_bb *hbb,
4864 enum built_in_function builtin)
4866 tree dst = gimple_call_arg (stmt, 0);
4867 hsa_op_address *dst_addr = get_address_from_value (dst, hbb);
4869 /* As gen_hsa_memory_set relies on memory alignment
4870 greater or equal to 8 bytes, we need to verify the alignment. */
4871 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
4872 hsa_op_reg *dst_addr_reg = new hsa_op_reg (addrtype);
4873 convert_addr_to_flat_segment (dst_addr, dst_addr_reg, hbb);
4875 /* Process BIT AND with 0x7 to identify the desired alignment
4876 of 8 bytes. */
4877 hsa_op_reg *masked = new hsa_op_reg (addrtype);
4879 gen_hsa_binary_operation (BRIG_OPCODE_AND, masked, dst_addr_reg,
4880 new hsa_op_immed (7, addrtype), hbb);
4882 hsa_op_reg *misaligned = new hsa_op_reg (BRIG_TYPE_B1);
4883 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE, misaligned->m_type,
4884 misaligned, masked,
4885 new hsa_op_immed (0, masked->m_type)));
4887 hsa_bb *native_impl_bb
4888 = expand_string_operation_builtin (stmt, hbb, misaligned);
4890 gen_hsa_memory_set (native_impl_bb, dst_addr, constant, n, BRIG_ALIGNMENT_8);
4891 hsa_bb *merge_bb
4892 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb->m_bb, 0)->dest);
4893 expand_lhs_of_string_op (stmt, n, merge_bb, builtin);
4896 /* Store into MEMORDER the memory order specified by tree T, which must be an
4897 integer constant representing a C++ memory order. If it isn't, issue an HSA
4898 sorry message using LOC and return true, otherwise return false and store
4899 the name of the requested order to *MNAME. */
4901 static bool
4902 hsa_memorder_from_tree (tree t, BrigMemoryOrder *memorder, const char **mname,
4903 location_t loc)
4905 if (!tree_fits_uhwi_p (t))
4907 HSA_SORRY_ATV (loc, "support for HSA does not implement memory model %E",
4909 return true;
4912 unsigned HOST_WIDE_INT mm = tree_to_uhwi (t);
4913 switch (mm & MEMMODEL_BASE_MASK)
4915 case MEMMODEL_RELAXED:
4916 *memorder = BRIG_MEMORY_ORDER_RELAXED;
4917 *mname = "relaxed";
4918 break;
4919 case MEMMODEL_CONSUME:
4920 /* HSA does not have an equivalent, but we can use the slightly stronger
4921 ACQUIRE. */
4922 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE;
4923 *mname = "consume";
4924 break;
4925 case MEMMODEL_ACQUIRE:
4926 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE;
4927 *mname = "acquire";
4928 break;
4929 case MEMMODEL_RELEASE:
4930 *memorder = BRIG_MEMORY_ORDER_SC_RELEASE;
4931 *mname = "release";
4932 break;
4933 case MEMMODEL_ACQ_REL:
4934 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE;
4935 *mname = "acq_rel";
4936 break;
4937 case MEMMODEL_SEQ_CST:
4938 /* Callers implementing a simple load or store need to remove the release
4939 or acquire part respectively. */
4940 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE;
4941 *mname = "seq_cst";
4942 break;
4943 default:
4945 HSA_SORRY_AT (loc, "support for HSA does not implement the specified "
4946 "memory model");
4947 return true;
4950 return false;
4953 /* Helper function to create an HSA atomic operation instruction out of calls
4954 to atomic builtins. RET_ORIG is true if the built-in is the variant that
4955 return s the value before applying operation, and false if it should return
4956 the value after applying the operation (if it returns value at all). ACODE
4957 is the atomic operation code, STMT is a gimple call to a builtin. HBB is
4958 the HSA BB to which the instruction should be added. If SIGNAL is true, the
4959 created operation will work on HSA signals rather than atomic variables. */
4961 static void
4962 gen_hsa_atomic_for_builtin (bool ret_orig, enum BrigAtomicOperation acode,
4963 gimple *stmt, hsa_bb *hbb, bool signal)
4965 tree lhs = gimple_call_lhs (stmt);
4967 tree type = TREE_TYPE (gimple_call_arg (stmt, 1));
4968 BrigType16_t hsa_type = hsa_type_for_scalar_tree_type (type, false);
4969 BrigType16_t mtype = mem_type_for_type (hsa_type);
4970 BrigMemoryOrder memorder;
4971 const char *mmname;
4973 if (hsa_memorder_from_tree (gimple_call_arg (stmt, 2), &memorder, &mmname,
4974 gimple_location (stmt)))
4975 return;
4977 /* Certain atomic insns must have Bx memory types. */
4978 switch (acode)
4980 case BRIG_ATOMIC_LD:
4981 case BRIG_ATOMIC_ST:
4982 case BRIG_ATOMIC_AND:
4983 case BRIG_ATOMIC_OR:
4984 case BRIG_ATOMIC_XOR:
4985 case BRIG_ATOMIC_EXCH:
4986 mtype = hsa_bittype_for_type (mtype);
4987 break;
4988 default:
4989 break;
4992 hsa_op_reg *dest;
4993 int nops, opcode;
4994 if (lhs)
4996 if (ret_orig)
4997 dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4998 else
4999 dest = new hsa_op_reg (hsa_type);
5000 opcode = signal ? BRIG_OPCODE_SIGNAL : BRIG_OPCODE_ATOMIC;
5001 nops = 3;
5003 else
5005 dest = NULL;
5006 opcode = signal ? BRIG_OPCODE_SIGNALNORET : BRIG_OPCODE_ATOMICNORET;
5007 nops = 2;
5010 if (acode == BRIG_ATOMIC_ST)
5012 if (memorder == BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE)
5013 memorder = BRIG_MEMORY_ORDER_SC_RELEASE;
5015 if (memorder != BRIG_MEMORY_ORDER_RELAXED
5016 && memorder != BRIG_MEMORY_ORDER_SC_RELEASE
5017 && memorder != BRIG_MEMORY_ORDER_NONE)
5019 HSA_SORRY_ATV (gimple_location (stmt),
5020 "support for HSA does not implement memory model for "
5021 "ATOMIC_ST: %s", mmname);
5022 return;
5026 hsa_insn_basic *atominsn;
5027 hsa_op_base *tgt;
5028 if (signal)
5030 atominsn = new hsa_insn_signal (nops, opcode, acode, mtype, memorder);
5031 tgt = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb);
5033 else
5035 atominsn = new hsa_insn_atomic (nops, opcode, acode, mtype, memorder);
5036 hsa_op_address *addr;
5037 addr = get_address_from_value (gimple_call_arg (stmt, 0), hbb);
5038 if (addr->m_symbol && addr->m_symbol->m_segment == BRIG_SEGMENT_PRIVATE)
5040 HSA_SORRY_AT (gimple_location (stmt),
5041 "HSA does not implement atomic operations in private "
5042 "segment");
5043 return;
5045 tgt = addr;
5048 hsa_op_with_type *op
5049 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1), hbb);
5050 if (lhs)
5052 atominsn->set_op (0, dest);
5053 atominsn->set_op (1, tgt);
5054 atominsn->set_op (2, op);
5056 else
5058 atominsn->set_op (0, tgt);
5059 atominsn->set_op (1, op);
5062 hbb->append_insn (atominsn);
5064 /* HSA does not natively support the variants that return the modified value,
5065 so re-do the operation again non-atomically if that is what was
5066 requested. */
5067 if (lhs && !ret_orig)
5069 int arith;
5070 switch (acode)
5072 case BRIG_ATOMIC_ADD:
5073 arith = BRIG_OPCODE_ADD;
5074 break;
5075 case BRIG_ATOMIC_AND:
5076 arith = BRIG_OPCODE_AND;
5077 break;
5078 case BRIG_ATOMIC_OR:
5079 arith = BRIG_OPCODE_OR;
5080 break;
5081 case BRIG_ATOMIC_SUB:
5082 arith = BRIG_OPCODE_SUB;
5083 break;
5084 case BRIG_ATOMIC_XOR:
5085 arith = BRIG_OPCODE_XOR;
5086 break;
5087 default:
5088 gcc_unreachable ();
5090 hsa_op_reg *real_dest = hsa_cfun->reg_for_gimple_ssa (lhs);
5091 gen_hsa_binary_operation (arith, real_dest, dest, op, hbb);
5095 /* Generate HSA instructions for an internal fn.
5096 Instructions will be appended to HBB, which also needs to be the
5097 corresponding structure to the basic_block of STMT. */
5099 static void
5100 gen_hsa_insn_for_internal_fn_call (gcall *stmt, hsa_bb *hbb)
5102 gcc_checking_assert (gimple_call_internal_fn (stmt));
5103 internal_fn fn = gimple_call_internal_fn (stmt);
5105 bool is_float_type_p = false;
5106 if (gimple_call_lhs (stmt) != NULL
5107 && TREE_TYPE (gimple_call_lhs (stmt)) == float_type_node)
5108 is_float_type_p = true;
5110 switch (fn)
5112 case IFN_CEIL:
5113 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL, stmt, hbb);
5114 break;
5116 case IFN_FLOOR:
5117 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR, stmt, hbb);
5118 break;
5120 case IFN_RINT:
5121 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT, stmt, hbb);
5122 break;
5124 case IFN_SQRT:
5125 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT, stmt, hbb);
5126 break;
5128 case IFN_RSQRT:
5129 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_NRSQRT, stmt, hbb);
5130 break;
5132 case IFN_TRUNC:
5133 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC, stmt, hbb);
5134 break;
5136 case IFN_COS:
5138 if (is_float_type_p)
5139 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS, stmt, hbb);
5140 else
5141 gen_hsa_unaryop_builtin_call (stmt, hbb);
5143 break;
5145 case IFN_EXP2:
5147 if (is_float_type_p)
5148 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2, stmt, hbb);
5149 else
5150 gen_hsa_unaryop_builtin_call (stmt, hbb);
5152 break;
5155 case IFN_LOG2:
5157 if (is_float_type_p)
5158 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2, stmt, hbb);
5159 else
5160 gen_hsa_unaryop_builtin_call (stmt, hbb);
5162 break;
5165 case IFN_SIN:
5167 if (is_float_type_p)
5168 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN, stmt, hbb);
5169 else
5170 gen_hsa_unaryop_builtin_call (stmt, hbb);
5171 break;
5174 case IFN_CLRSB:
5175 gen_hsa_clrsb (stmt, hbb);
5176 break;
5178 case IFN_CLZ:
5179 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT, stmt, hbb);
5180 break;
5182 case IFN_CTZ:
5183 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT, stmt, hbb);
5184 break;
5186 case IFN_FFS:
5187 gen_hsa_ffs (stmt, hbb);
5188 break;
5190 case IFN_PARITY:
5191 gen_hsa_parity (stmt, hbb);
5192 break;
5194 case IFN_POPCOUNT:
5195 gen_hsa_popcount (stmt, hbb);
5196 break;
5198 case IFN_DIVMOD:
5199 gen_hsa_divmod (stmt, hbb);
5200 break;
5202 case IFN_ACOS:
5203 case IFN_ASIN:
5204 case IFN_ATAN:
5205 case IFN_EXP:
5206 case IFN_EXP10:
5207 case IFN_EXPM1:
5208 case IFN_LOG:
5209 case IFN_LOG10:
5210 case IFN_LOG1P:
5211 case IFN_LOGB:
5212 case IFN_SIGNIFICAND:
5213 case IFN_TAN:
5214 case IFN_NEARBYINT:
5215 case IFN_ROUND:
5216 case IFN_ATAN2:
5217 case IFN_COPYSIGN:
5218 case IFN_FMOD:
5219 case IFN_POW:
5220 case IFN_REMAINDER:
5221 case IFN_SCALB:
5222 case IFN_FMIN:
5223 case IFN_FMAX:
5224 gen_hsa_insns_for_call_of_internal_fn (stmt, hbb);
5225 break;
5227 default:
5228 HSA_SORRY_ATV (gimple_location (stmt),
5229 "support for HSA does not implement internal function: %s",
5230 internal_fn_name (fn));
5231 break;
5235 /* Generate HSA instructions for the given call statement STMT. Instructions
5236 will be appended to HBB. */
5238 static void
5239 gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb)
5241 gcall *call = as_a <gcall *> (stmt);
5242 tree lhs = gimple_call_lhs (stmt);
5243 hsa_op_reg *dest;
5245 if (gimple_call_internal_p (stmt))
5247 gen_hsa_insn_for_internal_fn_call (call, hbb);
5248 return;
5251 if (!gimple_call_builtin_p (stmt, BUILT_IN_NORMAL))
5253 tree function_decl = gimple_call_fndecl (stmt);
5254 /* Prefetch pass can create type-mismatching prefetch builtin calls which
5255 fail the gimple_call_builtin_p test above. Handle them here. */
5256 if (DECL_BUILT_IN_CLASS (function_decl)
5257 && DECL_FUNCTION_CODE (function_decl) == BUILT_IN_PREFETCH)
5258 return;
5260 if (function_decl == NULL_TREE)
5262 HSA_SORRY_AT (gimple_location (stmt),
5263 "support for HSA does not implement indirect calls");
5264 return;
5267 if (hsa_callable_function_p (function_decl))
5268 gen_hsa_insns_for_direct_call (stmt, hbb);
5269 else if (!gen_hsa_insns_for_known_library_call (stmt, hbb))
5270 HSA_SORRY_AT (gimple_location (stmt),
5271 "HSA supports only calls of functions marked with pragma "
5272 "omp declare target");
5273 return;
5276 tree fndecl = gimple_call_fndecl (stmt);
5277 enum built_in_function builtin = DECL_FUNCTION_CODE (fndecl);
5278 switch (builtin)
5280 case BUILT_IN_FABS:
5281 case BUILT_IN_FABSF:
5282 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_ABS, stmt, hbb);
5283 break;
5285 case BUILT_IN_CEIL:
5286 case BUILT_IN_CEILF:
5287 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL, stmt, hbb);
5288 break;
5290 case BUILT_IN_FLOOR:
5291 case BUILT_IN_FLOORF:
5292 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR, stmt, hbb);
5293 break;
5295 case BUILT_IN_RINT:
5296 case BUILT_IN_RINTF:
5297 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT, stmt, hbb);
5298 break;
5300 case BUILT_IN_SQRT:
5301 case BUILT_IN_SQRTF:
5302 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT, stmt, hbb);
5303 break;
5305 case BUILT_IN_TRUNC:
5306 case BUILT_IN_TRUNCF:
5307 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC, stmt, hbb);
5308 break;
5310 case BUILT_IN_COS:
5311 case BUILT_IN_SIN:
5312 case BUILT_IN_EXP2:
5313 case BUILT_IN_LOG2:
5314 /* HSAIL does not provide an instruction for double argument type. */
5315 gen_hsa_unaryop_builtin_call (stmt, hbb);
5316 break;
5318 case BUILT_IN_COSF:
5319 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS, stmt, hbb);
5320 break;
5322 case BUILT_IN_EXP2F:
5323 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2, stmt, hbb);
5324 break;
5326 case BUILT_IN_LOG2F:
5327 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2, stmt, hbb);
5328 break;
5330 case BUILT_IN_SINF:
5331 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN, stmt, hbb);
5332 break;
5334 case BUILT_IN_CLRSB:
5335 case BUILT_IN_CLRSBL:
5336 case BUILT_IN_CLRSBLL:
5337 gen_hsa_clrsb (call, hbb);
5338 break;
5340 case BUILT_IN_CLZ:
5341 case BUILT_IN_CLZL:
5342 case BUILT_IN_CLZLL:
5343 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT, stmt, hbb);
5344 break;
5346 case BUILT_IN_CTZ:
5347 case BUILT_IN_CTZL:
5348 case BUILT_IN_CTZLL:
5349 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT, stmt, hbb);
5350 break;
5352 case BUILT_IN_FFS:
5353 case BUILT_IN_FFSL:
5354 case BUILT_IN_FFSLL:
5355 gen_hsa_ffs (call, hbb);
5356 break;
5358 case BUILT_IN_PARITY:
5359 case BUILT_IN_PARITYL:
5360 case BUILT_IN_PARITYLL:
5361 gen_hsa_parity (call, hbb);
5362 break;
5364 case BUILT_IN_POPCOUNT:
5365 case BUILT_IN_POPCOUNTL:
5366 case BUILT_IN_POPCOUNTLL:
5367 gen_hsa_popcount (call, hbb);
5368 break;
5370 case BUILT_IN_ATOMIC_LOAD_1:
5371 case BUILT_IN_ATOMIC_LOAD_2:
5372 case BUILT_IN_ATOMIC_LOAD_4:
5373 case BUILT_IN_ATOMIC_LOAD_8:
5374 case BUILT_IN_ATOMIC_LOAD_16:
5376 BrigType16_t mtype;
5377 hsa_op_base *src;
5378 src = get_address_from_value (gimple_call_arg (stmt, 0), hbb);
5380 BrigMemoryOrder memorder;
5381 const char *mmname;
5382 if (hsa_memorder_from_tree (gimple_call_arg (stmt, 1), &memorder,
5383 &mmname, gimple_location (stmt)))
5384 return;
5386 if (memorder == BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE)
5387 memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE;
5389 if (memorder != BRIG_MEMORY_ORDER_RELAXED
5390 && memorder != BRIG_MEMORY_ORDER_SC_ACQUIRE
5391 && memorder != BRIG_MEMORY_ORDER_NONE)
5393 HSA_SORRY_ATV (gimple_location (stmt),
5394 "support for HSA does not implement "
5395 "memory model for atomic loads: %s", mmname);
5396 return;
5399 if (lhs)
5401 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
5402 false);
5403 mtype = mem_type_for_type (t);
5404 mtype = hsa_bittype_for_type (mtype);
5405 dest = hsa_cfun->reg_for_gimple_ssa (lhs);
5407 else
5409 mtype = BRIG_TYPE_B64;
5410 dest = new hsa_op_reg (mtype);
5413 hsa_insn_basic *atominsn;
5414 atominsn = new hsa_insn_atomic (2, BRIG_OPCODE_ATOMIC, BRIG_ATOMIC_LD,
5415 mtype, memorder, dest, src);
5417 hbb->append_insn (atominsn);
5418 break;
5421 case BUILT_IN_ATOMIC_EXCHANGE_1:
5422 case BUILT_IN_ATOMIC_EXCHANGE_2:
5423 case BUILT_IN_ATOMIC_EXCHANGE_4:
5424 case BUILT_IN_ATOMIC_EXCHANGE_8:
5425 case BUILT_IN_ATOMIC_EXCHANGE_16:
5426 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_EXCH, stmt, hbb, false);
5427 break;
5428 break;
5430 case BUILT_IN_ATOMIC_FETCH_ADD_1:
5431 case BUILT_IN_ATOMIC_FETCH_ADD_2:
5432 case BUILT_IN_ATOMIC_FETCH_ADD_4:
5433 case BUILT_IN_ATOMIC_FETCH_ADD_8:
5434 case BUILT_IN_ATOMIC_FETCH_ADD_16:
5435 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ADD, stmt, hbb, false);
5436 break;
5437 break;
5439 case BUILT_IN_ATOMIC_FETCH_SUB_1:
5440 case BUILT_IN_ATOMIC_FETCH_SUB_2:
5441 case BUILT_IN_ATOMIC_FETCH_SUB_4:
5442 case BUILT_IN_ATOMIC_FETCH_SUB_8:
5443 case BUILT_IN_ATOMIC_FETCH_SUB_16:
5444 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_SUB, stmt, hbb, false);
5445 break;
5446 break;
5448 case BUILT_IN_ATOMIC_FETCH_AND_1:
5449 case BUILT_IN_ATOMIC_FETCH_AND_2:
5450 case BUILT_IN_ATOMIC_FETCH_AND_4:
5451 case BUILT_IN_ATOMIC_FETCH_AND_8:
5452 case BUILT_IN_ATOMIC_FETCH_AND_16:
5453 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_AND, stmt, hbb, false);
5454 break;
5455 break;
5457 case BUILT_IN_ATOMIC_FETCH_XOR_1:
5458 case BUILT_IN_ATOMIC_FETCH_XOR_2:
5459 case BUILT_IN_ATOMIC_FETCH_XOR_4:
5460 case BUILT_IN_ATOMIC_FETCH_XOR_8:
5461 case BUILT_IN_ATOMIC_FETCH_XOR_16:
5462 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_XOR, stmt, hbb, false);
5463 break;
5464 break;
5466 case BUILT_IN_ATOMIC_FETCH_OR_1:
5467 case BUILT_IN_ATOMIC_FETCH_OR_2:
5468 case BUILT_IN_ATOMIC_FETCH_OR_4:
5469 case BUILT_IN_ATOMIC_FETCH_OR_8:
5470 case BUILT_IN_ATOMIC_FETCH_OR_16:
5471 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_OR, stmt, hbb, false);
5472 break;
5473 break;
5475 case BUILT_IN_ATOMIC_STORE_1:
5476 case BUILT_IN_ATOMIC_STORE_2:
5477 case BUILT_IN_ATOMIC_STORE_4:
5478 case BUILT_IN_ATOMIC_STORE_8:
5479 case BUILT_IN_ATOMIC_STORE_16:
5480 /* Since there cannot be any LHS, the first parameter is meaningless. */
5481 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ST, stmt, hbb, false);
5482 break;
5483 break;
5485 case BUILT_IN_ATOMIC_ADD_FETCH_1:
5486 case BUILT_IN_ATOMIC_ADD_FETCH_2:
5487 case BUILT_IN_ATOMIC_ADD_FETCH_4:
5488 case BUILT_IN_ATOMIC_ADD_FETCH_8:
5489 case BUILT_IN_ATOMIC_ADD_FETCH_16:
5490 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_ADD, stmt, hbb, false);
5491 break;
5493 case BUILT_IN_ATOMIC_SUB_FETCH_1:
5494 case BUILT_IN_ATOMIC_SUB_FETCH_2:
5495 case BUILT_IN_ATOMIC_SUB_FETCH_4:
5496 case BUILT_IN_ATOMIC_SUB_FETCH_8:
5497 case BUILT_IN_ATOMIC_SUB_FETCH_16:
5498 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_SUB, stmt, hbb, false);
5499 break;
5501 case BUILT_IN_ATOMIC_AND_FETCH_1:
5502 case BUILT_IN_ATOMIC_AND_FETCH_2:
5503 case BUILT_IN_ATOMIC_AND_FETCH_4:
5504 case BUILT_IN_ATOMIC_AND_FETCH_8:
5505 case BUILT_IN_ATOMIC_AND_FETCH_16:
5506 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_AND, stmt, hbb, false);
5507 break;
5509 case BUILT_IN_ATOMIC_XOR_FETCH_1:
5510 case BUILT_IN_ATOMIC_XOR_FETCH_2:
5511 case BUILT_IN_ATOMIC_XOR_FETCH_4:
5512 case BUILT_IN_ATOMIC_XOR_FETCH_8:
5513 case BUILT_IN_ATOMIC_XOR_FETCH_16:
5514 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_XOR, stmt, hbb, false);
5515 break;
5517 case BUILT_IN_ATOMIC_OR_FETCH_1:
5518 case BUILT_IN_ATOMIC_OR_FETCH_2:
5519 case BUILT_IN_ATOMIC_OR_FETCH_4:
5520 case BUILT_IN_ATOMIC_OR_FETCH_8:
5521 case BUILT_IN_ATOMIC_OR_FETCH_16:
5522 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_OR, stmt, hbb, false);
5523 break;
5525 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_1:
5526 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_2:
5527 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_4:
5528 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8:
5529 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_16:
5531 tree type = TREE_TYPE (gimple_call_arg (stmt, 1));
5532 BrigType16_t atype
5533 = hsa_bittype_for_type (hsa_type_for_scalar_tree_type (type, false));
5534 BrigMemoryOrder memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE;
5535 hsa_insn_basic *atominsn;
5536 hsa_op_base *tgt;
5537 atominsn = new hsa_insn_atomic (4, BRIG_OPCODE_ATOMIC,
5538 BRIG_ATOMIC_CAS, atype, memorder);
5539 tgt = get_address_from_value (gimple_call_arg (stmt, 0), hbb);
5541 if (lhs != NULL)
5542 dest = hsa_cfun->reg_for_gimple_ssa (lhs);
5543 else
5544 dest = new hsa_op_reg (atype);
5546 atominsn->set_op (0, dest);
5547 atominsn->set_op (1, tgt);
5549 hsa_op_with_type *op
5550 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1), hbb);
5551 atominsn->set_op (2, op);
5552 op = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 2), hbb);
5553 atominsn->set_op (3, op);
5555 hbb->append_insn (atominsn);
5556 break;
5559 case BUILT_IN_HSA_WORKGROUPID:
5560 query_hsa_grid_dim (stmt, BRIG_OPCODE_WORKGROUPID, hbb);
5561 break;
5562 case BUILT_IN_HSA_WORKITEMID:
5563 query_hsa_grid_dim (stmt, BRIG_OPCODE_WORKITEMID, hbb);
5564 break;
5565 case BUILT_IN_HSA_WORKITEMABSID:
5566 query_hsa_grid_dim (stmt, BRIG_OPCODE_WORKITEMABSID, hbb);
5567 break;
5568 case BUILT_IN_HSA_GRIDSIZE:
5569 query_hsa_grid_dim (stmt, BRIG_OPCODE_GRIDSIZE, hbb);
5570 break;
5571 case BUILT_IN_HSA_CURRENTWORKGROUPSIZE:
5572 query_hsa_grid_dim (stmt, BRIG_OPCODE_CURRENTWORKGROUPSIZE, hbb);
5573 break;
5575 case BUILT_IN_GOMP_BARRIER:
5576 hbb->append_insn (new hsa_insn_br (0, BRIG_OPCODE_BARRIER, BRIG_TYPE_NONE,
5577 BRIG_WIDTH_ALL));
5578 break;
5579 case BUILT_IN_GOMP_PARALLEL:
5580 HSA_SORRY_AT (gimple_location (stmt),
5581 "support for HSA does not implement non-gridified "
5582 "OpenMP parallel constructs.");
5583 break;
5585 case BUILT_IN_OMP_GET_THREAD_NUM:
5587 query_hsa_grid_nodim (stmt, BRIG_OPCODE_WORKITEMFLATABSID, hbb);
5588 break;
5591 case BUILT_IN_OMP_GET_NUM_THREADS:
5593 gen_get_num_threads (stmt, hbb);
5594 break;
5596 case BUILT_IN_GOMP_TEAMS:
5598 gen_set_num_threads (gimple_call_arg (stmt, 1), hbb);
5599 break;
5601 case BUILT_IN_OMP_GET_NUM_TEAMS:
5603 gen_get_num_teams (stmt, hbb);
5604 break;
5606 case BUILT_IN_OMP_GET_TEAM_NUM:
5608 gen_get_team_num (stmt, hbb);
5609 break;
5611 case BUILT_IN_MEMCPY:
5612 case BUILT_IN_MEMPCPY:
5614 expand_memory_copy (stmt, hbb, builtin);
5615 break;
5617 case BUILT_IN_MEMSET:
5619 tree c = gimple_call_arg (stmt, 1);
5621 if (TREE_CODE (c) != INTEGER_CST)
5623 gen_hsa_insns_for_direct_call (stmt, hbb);
5624 return;
5627 tree byte_size = gimple_call_arg (stmt, 2);
5629 if (!tree_fits_uhwi_p (byte_size))
5631 gen_hsa_insns_for_direct_call (stmt, hbb);
5632 return;
5635 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size);
5637 if (n > HSA_MEMORY_BUILTINS_LIMIT)
5639 gen_hsa_insns_for_direct_call (stmt, hbb);
5640 return;
5643 unsigned HOST_WIDE_INT constant
5644 = tree_to_uhwi (fold_convert (unsigned_char_type_node, c));
5646 expand_memory_set (stmt, n, constant, hbb, builtin);
5648 break;
5650 case BUILT_IN_BZERO:
5652 tree byte_size = gimple_call_arg (stmt, 1);
5654 if (!tree_fits_uhwi_p (byte_size))
5656 gen_hsa_insns_for_direct_call (stmt, hbb);
5657 return;
5660 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size);
5662 if (n > HSA_MEMORY_BUILTINS_LIMIT)
5664 gen_hsa_insns_for_direct_call (stmt, hbb);
5665 return;
5668 expand_memory_set (stmt, n, 0, hbb, builtin);
5670 break;
5672 CASE_BUILT_IN_ALLOCA:
5674 gen_hsa_alloca (call, hbb);
5675 break;
5677 case BUILT_IN_PREFETCH:
5678 break;
5679 default:
5681 tree name_tree = DECL_NAME (fndecl);
5682 const char *s = IDENTIFIER_POINTER (name_tree);
5683 size_t len = strlen (s);
5684 if (len > 4 && (strncmp (s, "__builtin_GOMP_", 15) == 0))
5685 HSA_SORRY_ATV (gimple_location (stmt),
5686 "support for HSA does not implement GOMP function %s",
5688 else
5689 gen_hsa_insns_for_direct_call (stmt, hbb);
5690 return;
5695 /* Generate HSA instructions for a given gimple statement. Instructions will be
5696 appended to HBB. */
5698 static void
5699 gen_hsa_insns_for_gimple_stmt (gimple *stmt, hsa_bb *hbb)
5701 switch (gimple_code (stmt))
5703 case GIMPLE_ASSIGN:
5704 if (gimple_clobber_p (stmt))
5705 break;
5707 if (gimple_assign_single_p (stmt))
5709 tree lhs = gimple_assign_lhs (stmt);
5710 tree rhs = gimple_assign_rhs1 (stmt);
5711 gen_hsa_insns_for_single_assignment (lhs, rhs, hbb);
5713 else
5714 gen_hsa_insns_for_operation_assignment (stmt, hbb);
5715 break;
5716 case GIMPLE_RETURN:
5717 gen_hsa_insns_for_return (as_a <greturn *> (stmt), hbb);
5718 break;
5719 case GIMPLE_COND:
5720 gen_hsa_insns_for_cond_stmt (stmt, hbb);
5721 break;
5722 case GIMPLE_CALL:
5723 gen_hsa_insns_for_call (stmt, hbb);
5724 break;
5725 case GIMPLE_DEBUG:
5726 /* ??? HSA supports some debug facilities. */
5727 break;
5728 case GIMPLE_LABEL:
5730 tree label = gimple_label_label (as_a <glabel *> (stmt));
5731 if (FORCED_LABEL (label))
5732 HSA_SORRY_AT (gimple_location (stmt),
5733 "support for HSA does not implement gimple label with "
5734 "address taken");
5736 break;
5738 case GIMPLE_NOP:
5740 hbb->append_insn (new hsa_insn_basic (0, BRIG_OPCODE_NOP));
5741 break;
5743 case GIMPLE_SWITCH:
5745 gen_hsa_insns_for_switch_stmt (as_a <gswitch *> (stmt), hbb);
5746 break;
5748 default:
5749 HSA_SORRY_ATV (gimple_location (stmt),
5750 "support for HSA does not implement gimple statement %s",
5751 gimple_code_name[(int) gimple_code (stmt)]);
5755 /* Generate a HSA PHI from a gimple PHI. */
5757 static void
5758 gen_hsa_phi_from_gimple_phi (gimple *phi_stmt, hsa_bb *hbb)
5760 hsa_insn_phi *hphi;
5761 unsigned count = gimple_phi_num_args (phi_stmt);
5763 hsa_op_reg *dest
5764 = hsa_cfun->reg_for_gimple_ssa (gimple_phi_result (phi_stmt));
5765 hphi = new hsa_insn_phi (count, dest);
5766 hphi->m_bb = hbb->m_bb;
5768 auto_vec <tree, 8> aexprs;
5769 auto_vec <hsa_op_reg *, 8> aregs;
5771 /* Calling split_edge when processing a PHI node messes up with the order of
5772 gimple phi node arguments (it moves the one associated with the edge to
5773 the end). We need to keep the order of edges and arguments of HSA phi
5774 node arguments consistent, so we do all required splitting as the first
5775 step, and in reverse order as to not be affected by the re-orderings. */
5776 for (unsigned j = count; j != 0; j--)
5778 unsigned i = j - 1;
5779 tree op = gimple_phi_arg_def (phi_stmt, i);
5780 if (TREE_CODE (op) != ADDR_EXPR)
5781 continue;
5783 edge e = gimple_phi_arg_edge (as_a <gphi *> (phi_stmt), i);
5784 hsa_bb *hbb_src = hsa_init_new_bb (split_edge (e));
5785 hsa_op_address *addr = gen_hsa_addr (TREE_OPERAND (op, 0),
5786 hbb_src);
5788 hsa_op_reg *dest
5789 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
5790 hsa_insn_basic *insn
5791 = new hsa_insn_basic (2, BRIG_OPCODE_LDA, BRIG_TYPE_U64,
5792 dest, addr);
5793 hbb_src->append_insn (insn);
5794 aexprs.safe_push (op);
5795 aregs.safe_push (dest);
5798 tree lhs = gimple_phi_result (phi_stmt);
5799 for (unsigned i = 0; i < count; i++)
5801 tree op = gimple_phi_arg_def (phi_stmt, i);
5803 if (TREE_CODE (op) == SSA_NAME)
5805 hsa_op_reg *hreg = hsa_cfun->reg_for_gimple_ssa (op);
5806 hphi->set_op (i, hreg);
5808 else
5810 gcc_assert (is_gimple_min_invariant (op));
5811 tree t = TREE_TYPE (op);
5812 if (!POINTER_TYPE_P (t)
5813 || (TREE_CODE (op) == STRING_CST
5814 && TREE_CODE (TREE_TYPE (t)) == INTEGER_TYPE))
5815 hphi->set_op (i, new hsa_op_immed (op));
5816 else if (POINTER_TYPE_P (TREE_TYPE (lhs))
5817 && TREE_CODE (op) == INTEGER_CST)
5819 /* Handle assignment of NULL value to a pointer type. */
5820 hphi->set_op (i, new hsa_op_immed (op));
5822 else if (TREE_CODE (op) == ADDR_EXPR)
5824 hsa_op_reg *dest = NULL;
5825 for (unsigned a_idx = 0; a_idx < aexprs.length (); a_idx++)
5826 if (aexprs[a_idx] == op)
5828 dest = aregs[a_idx];
5829 break;
5831 gcc_assert (dest);
5832 hphi->set_op (i, dest);
5834 else
5836 HSA_SORRY_AT (gimple_location (phi_stmt),
5837 "support for HSA does not handle PHI nodes with "
5838 "constant address operands");
5839 return;
5844 hbb->append_phi (hphi);
5847 /* Constructor of class containing HSA-specific information about a basic
5848 block. CFG_BB is the CFG BB this HSA BB is associated with. IDX is the new
5849 index of this BB (so that the constructor does not attempt to use
5850 hsa_cfun during its construction). */
5852 hsa_bb::hsa_bb (basic_block cfg_bb, int idx)
5853 : m_bb (cfg_bb), m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL),
5854 m_last_phi (NULL), m_index (idx)
5856 gcc_assert (!cfg_bb->aux);
5857 cfg_bb->aux = this;
5860 /* Constructor of class containing HSA-specific information about a basic
5861 block. CFG_BB is the CFG BB this HSA BB is associated with. */
5863 hsa_bb::hsa_bb (basic_block cfg_bb)
5864 : m_bb (cfg_bb), m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL),
5865 m_last_phi (NULL), m_index (hsa_cfun->m_hbb_count++)
5867 gcc_assert (!cfg_bb->aux);
5868 cfg_bb->aux = this;
5871 /* Create and initialize and return a new hsa_bb structure for a given CFG
5872 basic block BB. */
5874 hsa_bb *
5875 hsa_init_new_bb (basic_block bb)
5877 void *m = obstack_alloc (&hsa_obstack, sizeof (hsa_bb));
5878 return new (m) hsa_bb (bb);
5881 /* Initialize OMP in an HSA basic block PROLOGUE. */
5883 static void
5884 init_prologue (void)
5886 if (!hsa_cfun->m_kern_p)
5887 return;
5889 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
5891 /* Create a magic number that is going to be printed by libgomp. */
5892 unsigned index = hsa_get_number_decl_kernel_mappings ();
5894 /* Emit store to debug argument. */
5895 if (PARAM_VALUE (PARAM_HSA_GEN_DEBUG_STORES) > 0)
5896 set_debug_value (prologue, new hsa_op_immed (1000 + index, BRIG_TYPE_U64));
5899 /* Initialize hsa_num_threads to a default value. */
5901 static void
5902 init_hsa_num_threads (void)
5904 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
5906 /* Save the default value to private variable hsa_num_threads. */
5907 hsa_insn_basic *basic
5908 = new hsa_insn_mem (BRIG_OPCODE_ST, hsa_num_threads->m_type,
5909 new hsa_op_immed (0, hsa_num_threads->m_type),
5910 new hsa_op_address (hsa_num_threads));
5911 prologue->append_insn (basic);
5914 /* Go over gimple representation and generate our internal HSA one. */
5916 static void
5917 gen_body_from_gimple ()
5919 basic_block bb;
5921 /* Verify CFG for complex edges we are unable to handle. */
5922 edge_iterator ei;
5923 edge e;
5925 FOR_EACH_BB_FN (bb, cfun)
5927 FOR_EACH_EDGE (e, ei, bb->succs)
5929 /* Verify all unsupported flags for edges that point
5930 to the same basic block. */
5931 if (e->flags & EDGE_EH)
5933 HSA_SORRY_AT (UNKNOWN_LOCATION,
5934 "support for HSA does not implement exception "
5935 "handling");
5936 return;
5941 FOR_EACH_BB_FN (bb, cfun)
5943 gimple_stmt_iterator gsi;
5944 hsa_bb *hbb = hsa_bb_for_bb (bb);
5945 if (hbb)
5946 continue;
5948 hbb = hsa_init_new_bb (bb);
5950 for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
5952 gen_hsa_insns_for_gimple_stmt (gsi_stmt (gsi), hbb);
5953 if (hsa_seen_error ())
5954 return;
5958 FOR_EACH_BB_FN (bb, cfun)
5960 gimple_stmt_iterator gsi;
5961 hsa_bb *hbb = hsa_bb_for_bb (bb);
5962 gcc_assert (hbb != NULL);
5964 for (gsi = gsi_start_phis (bb); !gsi_end_p (gsi); gsi_next (&gsi))
5965 if (!virtual_operand_p (gimple_phi_result (gsi_stmt (gsi))))
5966 gen_hsa_phi_from_gimple_phi (gsi_stmt (gsi), hbb);
5969 if (dump_file && (dump_flags & TDF_DETAILS))
5971 fprintf (dump_file, "------- Generated SSA form -------\n");
5972 dump_hsa_cfun (dump_file);
5976 static void
5977 gen_function_decl_parameters (hsa_function_representation *f,
5978 tree decl)
5980 tree parm;
5981 unsigned i;
5983 for (parm = TYPE_ARG_TYPES (TREE_TYPE (decl)), i = 0;
5984 parm;
5985 parm = TREE_CHAIN (parm), i++)
5987 /* Result type if last in the tree list. */
5988 if (TREE_CHAIN (parm) == NULL)
5989 break;
5991 tree v = TREE_VALUE (parm);
5993 hsa_symbol *arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
5994 BRIG_LINKAGE_NONE);
5995 arg->m_type = hsa_type_for_tree_type (v, &arg->m_dim);
5996 arg->m_name_number = i;
5998 f->m_input_args.safe_push (arg);
6001 tree result_type = TREE_TYPE (TREE_TYPE (decl));
6002 if (!VOID_TYPE_P (result_type))
6004 f->m_output_arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
6005 BRIG_LINKAGE_NONE);
6006 f->m_output_arg->m_type
6007 = hsa_type_for_tree_type (result_type, &f->m_output_arg->m_dim);
6008 f->m_output_arg->m_name = "res";
6012 /* Generate the vector of parameters of the HSA representation of the current
6013 function. This also includes the output parameter representing the
6014 result. */
6016 static void
6017 gen_function_def_parameters ()
6019 tree parm;
6021 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
6023 for (parm = DECL_ARGUMENTS (cfun->decl); parm;
6024 parm = DECL_CHAIN (parm))
6026 struct hsa_symbol **slot;
6028 hsa_symbol *arg
6029 = new hsa_symbol (BRIG_TYPE_NONE, hsa_cfun->m_kern_p
6030 ? BRIG_SEGMENT_KERNARG : BRIG_SEGMENT_ARG,
6031 BRIG_LINKAGE_FUNCTION);
6032 arg->fillup_for_decl (parm);
6034 hsa_cfun->m_input_args.safe_push (arg);
6036 if (hsa_seen_error ())
6037 return;
6039 arg->m_name = hsa_get_declaration_name (parm);
6041 /* Copy all input arguments and create corresponding private symbols
6042 for them. */
6043 hsa_symbol *private_arg;
6044 hsa_op_address *parm_addr = new hsa_op_address (arg);
6046 if (TREE_ADDRESSABLE (parm)
6047 || (!is_gimple_reg (parm) && !TREE_READONLY (parm)))
6049 private_arg = hsa_cfun->create_hsa_temporary (arg->m_type);
6050 private_arg->fillup_for_decl (parm);
6052 BrigAlignment8_t align = MIN (arg->m_align, private_arg->m_align);
6054 hsa_op_address *private_arg_addr = new hsa_op_address (private_arg);
6055 gen_hsa_memory_copy (prologue, private_arg_addr, parm_addr,
6056 arg->total_byte_size (), align);
6058 else
6059 private_arg = arg;
6061 slot = hsa_cfun->m_local_symbols->find_slot (private_arg, INSERT);
6062 gcc_assert (!*slot);
6063 *slot = private_arg;
6065 if (is_gimple_reg (parm))
6067 tree ddef = ssa_default_def (cfun, parm);
6068 if (ddef && !has_zero_uses (ddef))
6070 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (ddef),
6071 false);
6072 BrigType16_t mtype = mem_type_for_type (t);
6073 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (ddef);
6074 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mtype,
6075 dest, parm_addr);
6076 gcc_assert (!parm_addr->m_reg);
6077 prologue->append_insn (mem);
6082 if (!VOID_TYPE_P (TREE_TYPE (TREE_TYPE (cfun->decl))))
6084 struct hsa_symbol **slot;
6086 hsa_cfun->m_output_arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
6087 BRIG_LINKAGE_FUNCTION);
6088 hsa_cfun->m_output_arg->fillup_for_decl (DECL_RESULT (cfun->decl));
6090 if (hsa_seen_error ())
6091 return;
6093 hsa_cfun->m_output_arg->m_name = "res";
6094 slot = hsa_cfun->m_local_symbols->find_slot (hsa_cfun->m_output_arg,
6095 INSERT);
6096 gcc_assert (!*slot);
6097 *slot = hsa_cfun->m_output_arg;
6101 /* Generate function representation that corresponds to
6102 a function declaration. */
6104 hsa_function_representation *
6105 hsa_generate_function_declaration (tree decl)
6107 hsa_function_representation *fun
6108 = new hsa_function_representation (decl, false, 0);
6110 fun->m_declaration_p = true;
6111 fun->m_name = get_brig_function_name (decl);
6112 gen_function_decl_parameters (fun, decl);
6114 return fun;
6118 /* Generate function representation that corresponds to
6119 an internal FN. */
6121 hsa_function_representation *
6122 hsa_generate_internal_fn_decl (hsa_internal_fn *fn)
6124 hsa_function_representation *fun = new hsa_function_representation (fn);
6126 fun->m_name = fn->name ();
6128 for (unsigned i = 0; i < fn->get_arity (); i++)
6130 hsa_symbol *arg
6131 = new hsa_symbol (fn->get_argument_type (i), BRIG_SEGMENT_ARG,
6132 BRIG_LINKAGE_NONE);
6133 arg->m_name_number = i;
6134 fun->m_input_args.safe_push (arg);
6137 fun->m_output_arg = new hsa_symbol (fn->get_argument_type (-1),
6138 BRIG_SEGMENT_ARG, BRIG_LINKAGE_NONE);
6139 fun->m_output_arg->m_name = "res";
6141 return fun;
6144 /* Return true if switch statement S can be transformed
6145 to a SBR instruction in HSAIL. */
6147 static bool
6148 transformable_switch_to_sbr_p (gswitch *s)
6150 /* Identify if a switch statement can be transformed to
6151 SBR instruction, like:
6153 sbr_u32 $s1 [@label1, @label2, @label3];
6156 tree size = get_switch_size (s);
6157 if (!tree_fits_uhwi_p (size))
6158 return false;
6160 if (tree_to_uhwi (size) > HSA_MAXIMUM_SBR_LABELS)
6161 return false;
6163 return true;
6166 /* Structure hold connection between PHI nodes and immediate
6167 values hold by there nodes. */
6169 struct phi_definition
6171 phi_definition (unsigned phi_i, unsigned label_i, tree imm):
6172 phi_index (phi_i), label_index (label_i), phi_value (imm)
6175 unsigned phi_index;
6176 unsigned label_index;
6177 tree phi_value;
6180 /* Sum slice of a vector V, starting from index START and ending
6181 at the index END - 1. */
6183 template <typename T>
6184 static
6185 T sum_slice (const auto_vec <T> &v, unsigned start, unsigned end,
6186 T zero)
6188 T s = zero;
6190 for (unsigned i = start; i < end; i++)
6191 s += v[i];
6193 return s;
6196 /* Function transforms GIMPLE SWITCH statements to a series of IF statements.
6197 Let's assume following example:
6200 switch (index)
6201 case C1:
6202 L1: hard_work_1 ();
6203 break;
6204 case C2..C3:
6205 L2: hard_work_2 ();
6206 break;
6207 default:
6208 LD: hard_work_3 ();
6209 break;
6211 The transformation encompasses following steps:
6212 1) all immediate values used by edges coming from the switch basic block
6213 are saved
6214 2) all these edges are removed
6215 3) the switch statement (in L0) is replaced by:
6216 if (index == C1)
6217 goto L1;
6218 else
6219 goto L1';
6221 4) newly created basic block Lx' is used for generation of
6222 a next condition
6223 5) else branch of the last condition goes to LD
6224 6) fix all immediate values in PHI nodes that were propagated though
6225 edges that were removed in step 2
6227 Note: if a case is made by a range C1..C2, then process
6228 following transformation:
6230 switch_cond_op1 = C1 <= index;
6231 switch_cond_op2 = index <= C2;
6232 switch_cond_and = switch_cond_op1 & switch_cond_op2;
6233 if (switch_cond_and != 0)
6234 goto Lx;
6235 else
6236 goto Ly;
6240 static bool
6241 convert_switch_statements (void)
6243 function *func = DECL_STRUCT_FUNCTION (current_function_decl);
6244 basic_block bb;
6246 bool modified_cfg = false;
6248 FOR_EACH_BB_FN (bb, func)
6250 gimple_stmt_iterator gsi = gsi_last_bb (bb);
6251 if (gsi_end_p (gsi))
6252 continue;
6254 gimple *stmt = gsi_stmt (gsi);
6256 if (gimple_code (stmt) == GIMPLE_SWITCH)
6258 gswitch *s = as_a <gswitch *> (stmt);
6260 /* If the switch can utilize SBR insn, skip the statement. */
6261 if (transformable_switch_to_sbr_p (s))
6262 continue;
6264 modified_cfg = true;
6266 unsigned labels = gimple_switch_num_labels (s);
6267 tree index = gimple_switch_index (s);
6268 tree index_type = TREE_TYPE (index);
6269 tree default_label = gimple_switch_default_label (s);
6270 basic_block default_label_bb
6271 = label_to_block_fn (func, CASE_LABEL (default_label));
6272 basic_block cur_bb = bb;
6274 auto_vec <edge> new_edges;
6275 auto_vec <phi_definition *> phi_todo_list;
6276 auto_vec <profile_count> edge_counts;
6277 auto_vec <profile_probability> edge_probabilities;
6279 /* Investigate all labels that and PHI nodes in these edges which
6280 should be fixed after we add new collection of edges. */
6281 for (unsigned i = 0; i < labels; i++)
6283 tree label = gimple_switch_label (s, i);
6284 basic_block label_bb = label_to_block_fn (func, CASE_LABEL (label));
6285 edge e = find_edge (bb, label_bb);
6286 edge_counts.safe_push (e->count ());
6287 edge_probabilities.safe_push (e->probability);
6288 gphi_iterator phi_gsi;
6290 /* Save PHI definitions that will be destroyed because of an edge
6291 is going to be removed. */
6292 unsigned phi_index = 0;
6293 for (phi_gsi = gsi_start_phis (e->dest);
6294 !gsi_end_p (phi_gsi); gsi_next (&phi_gsi))
6296 gphi *phi = phi_gsi.phi ();
6297 for (unsigned j = 0; j < gimple_phi_num_args (phi); j++)
6299 if (gimple_phi_arg_edge (phi, j) == e)
6301 tree imm = gimple_phi_arg_def (phi, j);
6302 phi_definition *p = new phi_definition (phi_index, i,
6303 imm);
6304 phi_todo_list.safe_push (p);
6305 break;
6308 phi_index++;
6312 /* Remove all edges for the current basic block. */
6313 for (int i = EDGE_COUNT (bb->succs) - 1; i >= 0; i--)
6315 edge e = EDGE_SUCC (bb, i);
6316 remove_edge (e);
6319 /* Iterate all non-default labels. */
6320 for (unsigned i = 1; i < labels; i++)
6322 tree label = gimple_switch_label (s, i);
6323 tree low = CASE_LOW (label);
6324 tree high = CASE_HIGH (label);
6326 if (!useless_type_conversion_p (TREE_TYPE (low), index_type))
6327 low = fold_convert (index_type, low);
6329 gimple_stmt_iterator cond_gsi = gsi_last_bb (cur_bb);
6330 gimple *c = NULL;
6331 if (high)
6333 tree tmp1 = make_temp_ssa_name (boolean_type_node, NULL,
6334 "switch_cond_op1");
6336 gimple *assign1 = gimple_build_assign (tmp1, LE_EXPR, low,
6337 index);
6339 tree tmp2 = make_temp_ssa_name (boolean_type_node, NULL,
6340 "switch_cond_op2");
6342 if (!useless_type_conversion_p (TREE_TYPE (high), index_type))
6343 high = fold_convert (index_type, high);
6344 gimple *assign2 = gimple_build_assign (tmp2, LE_EXPR, index,
6345 high);
6347 tree tmp3 = make_temp_ssa_name (boolean_type_node, NULL,
6348 "switch_cond_and");
6349 gimple *assign3 = gimple_build_assign (tmp3, BIT_AND_EXPR, tmp1,
6350 tmp2);
6352 gsi_insert_before (&cond_gsi, assign1, GSI_SAME_STMT);
6353 gsi_insert_before (&cond_gsi, assign2, GSI_SAME_STMT);
6354 gsi_insert_before (&cond_gsi, assign3, GSI_SAME_STMT);
6356 tree b = constant_boolean_node (false, boolean_type_node);
6357 c = gimple_build_cond (NE_EXPR, tmp3, b, NULL, NULL);
6359 else
6360 c = gimple_build_cond (EQ_EXPR, index, low, NULL, NULL);
6362 gimple_set_location (c, gimple_location (stmt));
6364 gsi_insert_before (&cond_gsi, c, GSI_SAME_STMT);
6366 basic_block label_bb
6367 = label_to_block_fn (func, CASE_LABEL (label));
6368 edge new_edge = make_edge (cur_bb, label_bb, EDGE_TRUE_VALUE);
6369 profile_probability prob_sum = sum_slice <profile_probability>
6370 (edge_probabilities, i, labels, profile_probability::never ())
6371 + edge_probabilities[0];
6373 if (prob_sum.initialized_p ())
6374 new_edge->probability = edge_probabilities[i] / prob_sum;
6376 new_edges.safe_push (new_edge);
6378 if (i < labels - 1)
6380 /* Prepare another basic block that will contain
6381 next condition. */
6382 basic_block next_bb = create_empty_bb (cur_bb);
6383 if (current_loops)
6385 add_bb_to_loop (next_bb, cur_bb->loop_father);
6386 loops_state_set (LOOPS_NEED_FIXUP);
6389 edge next_edge = make_edge (cur_bb, next_bb, EDGE_FALSE_VALUE);
6390 next_edge->probability = new_edge->probability.invert ();
6391 next_bb->count = next_edge->count ();
6392 cur_bb = next_bb;
6394 else /* Link last IF statement and default label
6395 of the switch. */
6397 edge e = make_edge (cur_bb, default_label_bb, EDGE_FALSE_VALUE);
6398 e->probability = new_edge->probability.invert ();
6399 new_edges.safe_insert (0, e);
6403 /* Restore original PHI immediate value. */
6404 for (unsigned i = 0; i < phi_todo_list.length (); i++)
6406 phi_definition *phi_def = phi_todo_list[i];
6407 edge new_edge = new_edges[phi_def->label_index];
6409 gphi_iterator it = gsi_start_phis (new_edge->dest);
6410 for (unsigned i = 0; i < phi_def->phi_index; i++)
6411 gsi_next (&it);
6413 gphi *phi = it.phi ();
6414 add_phi_arg (phi, phi_def->phi_value, new_edge, UNKNOWN_LOCATION);
6415 delete phi_def;
6418 /* Remove the original GIMPLE switch statement. */
6419 gsi_remove (&gsi, true);
6423 if (dump_file)
6424 dump_function_to_file (current_function_decl, dump_file, TDF_DETAILS);
6426 return modified_cfg;
6429 /* Expand builtins that can't be handled by HSA back-end. */
6431 static void
6432 expand_builtins ()
6434 function *func = DECL_STRUCT_FUNCTION (current_function_decl);
6435 basic_block bb;
6437 FOR_EACH_BB_FN (bb, func)
6439 for (gimple_stmt_iterator gsi = gsi_start_bb (bb); !gsi_end_p (gsi);
6440 gsi_next (&gsi))
6442 gimple *stmt = gsi_stmt (gsi);
6444 if (gimple_code (stmt) != GIMPLE_CALL)
6445 continue;
6447 gcall *call = as_a <gcall *> (stmt);
6449 if (!gimple_call_builtin_p (call, BUILT_IN_NORMAL))
6450 continue;
6452 tree fndecl = gimple_call_fndecl (stmt);
6453 enum built_in_function fn = DECL_FUNCTION_CODE (fndecl);
6454 switch (fn)
6456 case BUILT_IN_CEXPF:
6457 case BUILT_IN_CEXPIF:
6458 case BUILT_IN_CEXPI:
6460 /* Similar to builtins.c (expand_builtin_cexpi), the builtin
6461 can be transformed to: cexp(I * z) = ccos(z) + I * csin(z). */
6462 tree lhs = gimple_call_lhs (stmt);
6463 tree rhs = gimple_call_arg (stmt, 0);
6464 tree rhs_type = TREE_TYPE (rhs);
6465 bool float_type_p = rhs_type == float_type_node;
6466 tree real_part = make_temp_ssa_name (rhs_type, NULL,
6467 "cexp_real_part");
6468 tree imag_part = make_temp_ssa_name (rhs_type, NULL,
6469 "cexp_imag_part");
6471 tree cos_fndecl
6472 = mathfn_built_in (rhs_type, fn == float_type_p
6473 ? BUILT_IN_COSF : BUILT_IN_COS);
6474 gcall *cos = gimple_build_call (cos_fndecl, 1, rhs);
6475 gimple_call_set_lhs (cos, real_part);
6476 gsi_insert_before (&gsi, cos, GSI_SAME_STMT);
6478 tree sin_fndecl
6479 = mathfn_built_in (rhs_type, fn == float_type_p
6480 ? BUILT_IN_SINF : BUILT_IN_SIN);
6481 gcall *sin = gimple_build_call (sin_fndecl, 1, rhs);
6482 gimple_call_set_lhs (sin, imag_part);
6483 gsi_insert_before (&gsi, sin, GSI_SAME_STMT);
6486 gassign *assign = gimple_build_assign (lhs, COMPLEX_EXPR,
6487 real_part, imag_part);
6488 gsi_insert_before (&gsi, assign, GSI_SAME_STMT);
6489 gsi_remove (&gsi, true);
6491 break;
6493 default:
6494 break;
6500 /* Emit HSA module variables that are global for the entire module. */
6502 static void
6503 emit_hsa_module_variables (void)
6505 hsa_num_threads = new hsa_symbol (BRIG_TYPE_U32, BRIG_SEGMENT_PRIVATE,
6506 BRIG_LINKAGE_MODULE, true);
6508 hsa_num_threads->m_name = "hsa_num_threads";
6510 hsa_brig_emit_omp_symbols ();
6513 /* Generate HSAIL representation of the current function and write into a
6514 special section of the output file. If KERNEL is set, the function will be
6515 considered an HSA kernel callable from the host, otherwise it will be
6516 compiled as an HSA function callable from other HSA code. */
6518 static void
6519 generate_hsa (bool kernel)
6521 hsa_init_data_for_cfun ();
6523 if (hsa_num_threads == NULL)
6524 emit_hsa_module_variables ();
6526 bool modified_cfg = convert_switch_statements ();
6527 /* Initialize hsa_cfun. */
6528 hsa_cfun = new hsa_function_representation (cfun->decl, kernel,
6529 SSANAMES (cfun)->length (),
6530 modified_cfg);
6531 hsa_cfun->init_extra_bbs ();
6533 if (flag_tm)
6535 HSA_SORRY_AT (UNKNOWN_LOCATION,
6536 "support for HSA does not implement transactional memory");
6537 goto fail;
6540 verify_function_arguments (cfun->decl);
6541 if (hsa_seen_error ())
6542 goto fail;
6544 hsa_cfun->m_name = get_brig_function_name (cfun->decl);
6546 gen_function_def_parameters ();
6547 if (hsa_seen_error ())
6548 goto fail;
6550 init_prologue ();
6552 gen_body_from_gimple ();
6553 if (hsa_seen_error ())
6554 goto fail;
6556 if (hsa_cfun->m_kernel_dispatch_count)
6557 init_hsa_num_threads ();
6559 if (hsa_cfun->m_kern_p)
6561 hsa_function_summary *s
6562 = hsa_summaries->get (cgraph_node::get (hsa_cfun->m_decl));
6563 hsa_add_kern_decl_mapping (current_function_decl, hsa_cfun->m_name,
6564 hsa_cfun->m_maximum_omp_data_size,
6565 s->m_gridified_kernel_p);
6568 if (flag_checking)
6570 for (unsigned i = 0; i < hsa_cfun->m_ssa_map.length (); i++)
6571 if (hsa_cfun->m_ssa_map[i])
6572 hsa_cfun->m_ssa_map[i]->verify_ssa ();
6574 basic_block bb;
6575 FOR_EACH_BB_FN (bb, cfun)
6577 hsa_bb *hbb = hsa_bb_for_bb (bb);
6579 for (hsa_insn_basic *insn = hbb->m_first_insn; insn;
6580 insn = insn->m_next)
6581 insn->verify ();
6585 hsa_regalloc ();
6586 hsa_brig_emit_function ();
6588 fail:
6589 hsa_deinit_data_for_cfun ();
6592 namespace {
6594 const pass_data pass_data_gen_hsail =
6596 GIMPLE_PASS,
6597 "hsagen", /* name */
6598 OPTGROUP_OMP, /* optinfo_flags */
6599 TV_NONE, /* tv_id */
6600 PROP_cfg | PROP_ssa, /* properties_required */
6601 0, /* properties_provided */
6602 0, /* properties_destroyed */
6603 0, /* todo_flags_start */
6604 0 /* todo_flags_finish */
6607 class pass_gen_hsail : public gimple_opt_pass
6609 public:
6610 pass_gen_hsail (gcc::context *ctxt)
6611 : gimple_opt_pass(pass_data_gen_hsail, ctxt)
6614 /* opt_pass methods: */
6615 bool gate (function *);
6616 unsigned int execute (function *);
6618 }; // class pass_gen_hsail
6620 /* Determine whether or not to run generation of HSAIL. */
6622 bool
6623 pass_gen_hsail::gate (function *f)
6625 return hsa_gen_requested_p ()
6626 && hsa_gpu_implementation_p (f->decl);
6629 unsigned int
6630 pass_gen_hsail::execute (function *)
6632 hsa_function_summary *s
6633 = hsa_summaries->get (cgraph_node::get_create (current_function_decl));
6635 expand_builtins ();
6636 generate_hsa (s->m_kind == HSA_KERNEL);
6637 TREE_ASM_WRITTEN (current_function_decl) = 1;
6638 return TODO_discard_function;
6641 } // anon namespace
6643 /* Create the instance of hsa gen pass. */
6645 gimple_opt_pass *
6646 make_pass_gen_hsail (gcc::context *ctxt)
6648 return new pass_gen_hsail (ctxt);