1 /* A pass for lowering gimple to HSAIL
2 Copyright (C) 2013-2017 Free Software Foundation, Inc.
3 Contributed by Martin Jambor <mjambor@suse.cz> and
4 Martin Liska <mliska@suse.cz>.
6 This file is part of GCC.
8 GCC is free software; you can redistribute it and/or modify
9 it under the terms of the GNU General Public License as published by
10 the Free Software Foundation; either version 3, or (at your option)
13 GCC is distributed in the hope that it will be useful,
14 but WITHOUT ANY WARRANTY; without even the implied warranty of
15 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
16 GNU General Public License for more details.
18 You should have received a copy of the GNU General Public License
19 along with GCC; see the file COPYING3. If not see
20 <http://www.gnu.org/licenses/>. */
24 #include "coretypes.h"
28 #include "hash-table.h"
31 #include "tree-pass.h"
33 #include "basic-block.h"
35 #include "fold-const.h"
37 #include "gimple-iterator.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"
46 #include "tree-ssanames.h"
48 #include "ssa-iterators.h"
50 #include "print-tree.h"
51 #include "symbol-summary.h"
52 #include "hsa-common.h"
59 #include "gomp-constants.h"
60 #include "internal-fn.h"
62 #include "stor-layout.h"
64 /* Print a warning message and set that we have seen an error. */
66 #define HSA_SORRY_ATV(location, message, ...) \
70 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
72 inform (location, message, __VA_ARGS__); \
76 /* Same as previous, but highlight a location. */
78 #define HSA_SORRY_AT(location, message) \
82 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
84 inform (location, message); \
88 /* Default number of threads used by kernel dispatch. */
90 #define HSA_DEFAULT_NUM_THREADS 64
92 /* Following structures are defined in the final version
93 of HSA specification. */
95 /* HSA queue packet is shadow structure, originally provided by AMD. */
97 struct hsa_queue_packet
101 uint16_t workgroup_size_x;
102 uint16_t workgroup_size_y;
103 uint16_t workgroup_size_z;
105 uint32_t grid_size_x;
106 uint32_t grid_size_y;
107 uint32_t grid_size_z;
108 uint32_t private_segment_size;
109 uint32_t group_segment_size;
110 uint64_t kernel_object;
111 void *kernarg_address;
113 uint64_t completion_signal;
116 /* HSA queue is shadow structure, originally provided by AMD. */
123 uint64_t doorbell_signal;
129 static struct obstack hsa_obstack;
131 /* List of pointers to all instructions that come from an object allocator. */
132 static vec <hsa_insn_basic *> hsa_instructions;
134 /* List of pointers to all operands that come from an object allocator. */
135 static vec <hsa_op_base *> hsa_operands;
137 hsa_symbol::hsa_symbol ()
138 : m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
139 m_directive_offset (0), m_type (BRIG_TYPE_NONE),
140 m_segment (BRIG_SEGMENT_NONE), m_linkage (BRIG_LINKAGE_NONE), m_dim (0),
141 m_cst_value (NULL), m_global_scope_p (false), m_seen_error (false),
142 m_allocation (BRIG_ALLOCATION_AUTOMATIC), m_emitted_to_brig (false)
147 hsa_symbol::hsa_symbol (BrigType16_t type, BrigSegment8_t segment,
148 BrigLinkage8_t linkage, bool global_scope_p,
149 BrigAllocation allocation, BrigAlignment8_t align)
150 : m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
151 m_directive_offset (0), m_type (type), m_segment (segment),
152 m_linkage (linkage), m_dim (0), m_cst_value (NULL),
153 m_global_scope_p (global_scope_p), m_seen_error (false),
154 m_allocation (allocation), m_emitted_to_brig (false), m_align (align)
158 unsigned HOST_WIDE_INT
159 hsa_symbol::total_byte_size ()
161 unsigned HOST_WIDE_INT s
162 = hsa_type_bit_size (~BRIG_TYPE_ARRAY_MASK & m_type);
163 gcc_assert (s % BITS_PER_UNIT == 0);
172 /* Forward declaration. */
175 hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p,
179 hsa_symbol::fillup_for_decl (tree decl)
182 m_type = hsa_type_for_tree_type (TREE_TYPE (decl), &m_dim, false);
183 if (hsa_seen_error ())
189 m_align = MAX (m_align, hsa_natural_alignment (m_type));
192 /* Constructor of class representing global HSA function/kernel information and
193 state. FNDECL is function declaration, KERNEL_P is true if the function
194 is going to become a HSA kernel. If the function has body, SSA_NAMES_COUNT
195 should be set to number of SSA names used in the function.
196 MODIFIED_CFG is set to true in case we modified control-flow graph
199 hsa_function_representation::hsa_function_representation
200 (tree fdecl, bool kernel_p, unsigned ssa_names_count, bool modified_cfg)
202 m_reg_count (0), m_input_args (vNULL),
203 m_output_arg (NULL), m_spill_symbols (vNULL), m_global_symbols (vNULL),
204 m_private_variables (vNULL), m_called_functions (vNULL),
205 m_called_internal_fns (vNULL), m_hbb_count (0),
206 m_in_ssa (true), m_kern_p (kernel_p), m_declaration_p (false),
207 m_decl (fdecl), m_internal_fn (NULL), m_shadow_reg (NULL),
208 m_kernel_dispatch_count (0), m_maximum_omp_data_size (0),
209 m_seen_error (false), m_temp_symbol_count (0), m_ssa_map (),
210 m_modified_cfg (modified_cfg)
212 int sym_init_len = (vec_safe_length (cfun->local_decls) / 2) + 1;;
213 m_local_symbols = new hash_table <hsa_noop_symbol_hasher> (sym_init_len);
214 m_ssa_map.safe_grow_cleared (ssa_names_count);
217 /* Constructor of class representing HSA function information that
218 is derived for an internal function. */
219 hsa_function_representation::hsa_function_representation (hsa_internal_fn *fn)
220 : m_reg_count (0), m_input_args (vNULL),
221 m_output_arg (NULL), m_local_symbols (NULL),
222 m_spill_symbols (vNULL), m_global_symbols (vNULL),
223 m_private_variables (vNULL), m_called_functions (vNULL),
224 m_called_internal_fns (vNULL), m_hbb_count (0),
225 m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL),
226 m_internal_fn (fn), m_shadow_reg (NULL), m_kernel_dispatch_count (0),
227 m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
230 /* Destructor of class holding function/kernel-wide information and state. */
232 hsa_function_representation::~hsa_function_representation ()
234 /* Kernel names are deallocated at the end of BRIG output when deallocating
235 hsa_decl_kernel_mapping. */
236 if (!m_kern_p || m_seen_error)
239 for (unsigned i = 0; i < m_input_args.length (); i++)
240 delete m_input_args[i];
241 m_input_args.release ();
244 delete m_local_symbols;
246 for (unsigned i = 0; i < m_spill_symbols.length (); i++)
247 delete m_spill_symbols[i];
248 m_spill_symbols.release ();
251 for (unsigned i = 0; i < m_global_symbols.iterate (i, &sym); i++)
252 if (sym->m_linkage != BRIG_ALLOCATION_PROGRAM)
254 m_global_symbols.release ();
256 for (unsigned i = 0; i < m_private_variables.length (); i++)
257 delete m_private_variables[i];
258 m_private_variables.release ();
259 m_called_functions.release ();
260 m_ssa_map.release ();
262 for (unsigned i = 0; i < m_called_internal_fns.length (); i++)
263 delete m_called_internal_fns[i];
267 hsa_function_representation::get_shadow_reg ()
269 /* If we compile a function with kernel dispatch and does not set
270 an optimization level, the function won't be inlined and
278 /* Append the shadow argument. */
279 hsa_symbol *shadow = new hsa_symbol (BRIG_TYPE_U64, BRIG_SEGMENT_KERNARG,
280 BRIG_LINKAGE_FUNCTION);
281 m_input_args.safe_push (shadow);
282 shadow->m_name = "hsa_runtime_shadow";
284 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_U64);
285 hsa_op_address *addr = new hsa_op_address (shadow);
287 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, r, addr);
288 hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun))->append_insn (mem);
294 bool hsa_function_representation::has_shadow_reg_p ()
296 return m_shadow_reg != NULL;
300 hsa_function_representation::init_extra_bbs ()
302 hsa_init_new_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
303 hsa_init_new_bb (EXIT_BLOCK_PTR_FOR_FN (cfun));
307 hsa_function_representation::update_dominance ()
311 free_dominance_info (CDI_DOMINATORS);
312 calculate_dominance_info (CDI_DOMINATORS);
317 hsa_function_representation::create_hsa_temporary (BrigType16_t type)
319 hsa_symbol *s = new hsa_symbol (type, BRIG_SEGMENT_PRIVATE,
320 BRIG_LINKAGE_FUNCTION);
321 s->m_name_number = m_temp_symbol_count++;
323 hsa_cfun->m_private_variables.safe_push (s);
328 hsa_function_representation::get_linkage ()
331 return BRIG_LINKAGE_PROGRAM;
333 return m_kern_p || TREE_PUBLIC (m_decl) ?
334 BRIG_LINKAGE_PROGRAM : BRIG_LINKAGE_MODULE;
337 /* Hash map of simple OMP builtins. */
338 static hash_map <nofree_string_hash, omp_simple_builtin> *omp_simple_builtins
341 /* Warning messages for OMP builtins. */
343 #define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP " \
345 #define HSA_WARN_TIMING_ROUTINE "support for HSA does not implement OpenMP " \
347 #define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \
348 "undefined semantics within target regions, support for HSA ignores them"
349 #define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP " \
352 /* Initialize hash map with simple OMP builtins. */
355 hsa_init_simple_builtins ()
357 if (omp_simple_builtins != NULL)
361 = new hash_map <nofree_string_hash, omp_simple_builtin> ();
363 omp_simple_builtin omp_builtins[] =
365 omp_simple_builtin ("omp_get_initial_device", NULL, false,
366 new hsa_op_immed (GOMP_DEVICE_HOST,
367 (BrigType16_t) BRIG_TYPE_S32)),
368 omp_simple_builtin ("omp_is_initial_device", NULL, false,
369 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
370 omp_simple_builtin ("omp_get_dynamic", NULL, false,
371 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
372 omp_simple_builtin ("omp_set_dynamic", NULL, false, NULL),
373 omp_simple_builtin ("omp_init_lock", HSA_WARN_LOCK_ROUTINE, true),
374 omp_simple_builtin ("omp_init_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
376 omp_simple_builtin ("omp_init_nest_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
378 omp_simple_builtin ("omp_destroy_lock", HSA_WARN_LOCK_ROUTINE, true),
379 omp_simple_builtin ("omp_set_lock", HSA_WARN_LOCK_ROUTINE, true),
380 omp_simple_builtin ("omp_unset_lock", HSA_WARN_LOCK_ROUTINE, true),
381 omp_simple_builtin ("omp_test_lock", HSA_WARN_LOCK_ROUTINE, true),
382 omp_simple_builtin ("omp_get_wtime", HSA_WARN_TIMING_ROUTINE, true),
383 omp_simple_builtin ("omp_get_wtick", HSA_WARN_TIMING_ROUTINE, true),
384 omp_simple_builtin ("omp_target_alloc", HSA_WARN_MEMORY_ROUTINE, false,
385 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_U64)),
386 omp_simple_builtin ("omp_target_free", HSA_WARN_MEMORY_ROUTINE, false),
387 omp_simple_builtin ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE,
389 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
390 omp_simple_builtin ("omp_target_memcpy", HSA_WARN_MEMORY_ROUTINE, false,
391 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
392 omp_simple_builtin ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE,
394 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
395 omp_simple_builtin ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE,
397 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
398 omp_simple_builtin ("omp_target_disassociate_ptr",
399 HSA_WARN_MEMORY_ROUTINE,
401 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
402 omp_simple_builtin ("omp_set_max_active_levels",
403 "Support for HSA only allows only one active level, "
404 "call to omp_set_max_active_levels will be ignored "
405 "in the generated HSAIL",
407 omp_simple_builtin ("omp_get_max_active_levels", NULL, false,
408 new hsa_op_immed (1, (BrigType16_t) BRIG_TYPE_S32)),
409 omp_simple_builtin ("omp_in_final", NULL, false,
410 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
411 omp_simple_builtin ("omp_get_proc_bind", HSA_WARN_AFFINITY, false,
412 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
413 omp_simple_builtin ("omp_get_num_places", HSA_WARN_AFFINITY, false,
414 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
415 omp_simple_builtin ("omp_get_place_num_procs", HSA_WARN_AFFINITY, false,
416 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
417 omp_simple_builtin ("omp_get_place_proc_ids", HSA_WARN_AFFINITY, false,
419 omp_simple_builtin ("omp_get_place_num", HSA_WARN_AFFINITY, false,
420 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
421 omp_simple_builtin ("omp_get_partition_num_places", HSA_WARN_AFFINITY,
423 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
424 omp_simple_builtin ("omp_get_partition_place_nums", HSA_WARN_AFFINITY,
426 omp_simple_builtin ("omp_set_default_device",
427 "omp_set_default_device has undefined semantics "
428 "within target regions, support for HSA ignores it",
430 omp_simple_builtin ("omp_get_default_device",
431 "omp_get_default_device has undefined semantics "
432 "within target regions, support for HSA ignores it",
434 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
435 omp_simple_builtin ("omp_get_num_devices",
436 "omp_get_num_devices has undefined semantics "
437 "within target regions, support for HSA ignores it",
439 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
440 omp_simple_builtin ("omp_get_num_procs", NULL, true, NULL),
441 omp_simple_builtin ("omp_get_cancellation", NULL, true, NULL),
442 omp_simple_builtin ("omp_set_nested", NULL, true, NULL),
443 omp_simple_builtin ("omp_get_nested", NULL, true, NULL),
444 omp_simple_builtin ("omp_set_schedule", NULL, true, NULL),
445 omp_simple_builtin ("omp_get_schedule", NULL, true, NULL),
446 omp_simple_builtin ("omp_get_thread_limit", NULL, true, NULL),
447 omp_simple_builtin ("omp_get_team_size", NULL, true, NULL),
448 omp_simple_builtin ("omp_get_ancestor_thread_num", NULL, true, NULL),
449 omp_simple_builtin ("omp_get_max_task_priority", NULL, true, NULL)
452 unsigned count = sizeof (omp_builtins) / sizeof (omp_simple_builtin);
454 for (unsigned i = 0; i < count; i++)
455 omp_simple_builtins->put (omp_builtins[i].m_name, omp_builtins[i]);
458 /* Allocate HSA structures that we need only while generating with this. */
461 hsa_init_data_for_cfun ()
463 hsa_init_compilation_unit_data ();
464 gcc_obstack_init (&hsa_obstack);
467 /* Deinitialize HSA subsystem and free all allocated memory. */
470 hsa_deinit_data_for_cfun (void)
474 FOR_ALL_BB_FN (bb, cfun)
477 hsa_bb *hbb = hsa_bb_for_bb (bb);
482 for (unsigned int i = 0; i < hsa_operands.length (); i++)
483 hsa_destroy_operand (hsa_operands[i]);
485 hsa_operands.release ();
487 for (unsigned i = 0; i < hsa_instructions.length (); i++)
488 hsa_destroy_insn (hsa_instructions[i]);
490 hsa_instructions.release ();
492 if (omp_simple_builtins != NULL)
494 delete omp_simple_builtins;
495 omp_simple_builtins = NULL;
498 obstack_free (&hsa_obstack, NULL);
502 /* Return the type which holds addresses in the given SEGMENT. */
505 hsa_get_segment_addr_type (BrigSegment8_t segment)
509 case BRIG_SEGMENT_NONE:
512 case BRIG_SEGMENT_FLAT:
513 case BRIG_SEGMENT_GLOBAL:
514 case BRIG_SEGMENT_READONLY:
515 case BRIG_SEGMENT_KERNARG:
516 return hsa_machine_large_p () ? BRIG_TYPE_U64 : BRIG_TYPE_U32;
518 case BRIG_SEGMENT_GROUP:
519 case BRIG_SEGMENT_PRIVATE:
520 case BRIG_SEGMENT_SPILL:
521 case BRIG_SEGMENT_ARG:
522 return BRIG_TYPE_U32;
527 /* Return integer brig type according to provided SIZE in bytes. If SIGN
528 is set to true, return signed integer type. */
531 get_integer_type_by_bytes (unsigned size, bool sign)
539 return BRIG_TYPE_S16;
541 return BRIG_TYPE_S32;
543 return BRIG_TYPE_S64;
553 return BRIG_TYPE_U16;
555 return BRIG_TYPE_U32;
557 return BRIG_TYPE_U64;
565 /* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
566 are assumed to use flat addressing. If min32int is true, always expand
567 integer types to one that has at least 32 bits. */
570 hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
574 BrigType16_t res = BRIG_TYPE_NONE;
576 gcc_checking_assert (TYPE_P (type));
577 gcc_checking_assert (!AGGREGATE_TYPE_P (type));
578 if (POINTER_TYPE_P (type))
579 return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
581 if (TREE_CODE (type) == VECTOR_TYPE || TREE_CODE (type) == COMPLEX_TYPE)
582 base = TREE_TYPE (type);
586 if (!tree_fits_uhwi_p (TYPE_SIZE (base)))
588 HSA_SORRY_ATV (EXPR_LOCATION (type),
589 "support for HSA does not implement huge or "
590 "variable-sized type %qT", type);
594 bsize = tree_to_uhwi (TYPE_SIZE (base));
595 unsigned byte_size = bsize / BITS_PER_UNIT;
596 if (INTEGRAL_TYPE_P (base))
597 res = get_integer_type_by_bytes (byte_size, !TYPE_UNSIGNED (base));
598 else if (SCALAR_FLOAT_TYPE_P (base))
616 if (res == BRIG_TYPE_NONE)
618 HSA_SORRY_ATV (EXPR_LOCATION (type),
619 "support for HSA does not implement type %qT", type);
623 if (TREE_CODE (type) == VECTOR_TYPE)
625 HOST_WIDE_INT tsize = tree_to_uhwi (TYPE_SIZE (type));
629 HSA_SORRY_ATV (EXPR_LOCATION (type),
630 "support for HSA does not implement a vector type "
631 "where a type and unit size are equal: %qT", type);
638 res |= BRIG_TYPE_PACK_32;
641 res |= BRIG_TYPE_PACK_64;
644 res |= BRIG_TYPE_PACK_128;
647 HSA_SORRY_ATV (EXPR_LOCATION (type),
648 "support for HSA does not implement type %qT", type);
654 /* Registers/immediate operands can only be 32bit or more except for
656 if (res == BRIG_TYPE_U8 || res == BRIG_TYPE_U16)
658 else if (res == BRIG_TYPE_S8 || res == BRIG_TYPE_S16)
662 if (TREE_CODE (type) == COMPLEX_TYPE)
664 unsigned bsize = 2 * hsa_type_bit_size (res);
665 res = hsa_bittype_for_bitsize (bsize);
671 /* Returns the BRIG type we need to load/store entities of TYPE. */
674 mem_type_for_type (BrigType16_t type)
676 /* HSA has non-intuitive constraints on load/store types. If it's
677 a bit-type it _must_ be B128, if it's not a bit-type it must be
678 64bit max. So for loading entities of 128 bits (e.g. vectors)
679 we have to to B128, while for loading the rest we have to use the
680 input type (??? or maybe also flattened to a equally sized non-vector
682 if ((type & BRIG_TYPE_PACK_MASK) == BRIG_TYPE_PACK_128)
683 return BRIG_TYPE_B128;
684 else if (hsa_btype_p (type) || hsa_type_packed_p (type))
686 unsigned bitsize = hsa_type_bit_size (type);
688 return hsa_uint_for_bitsize (bitsize);
690 return hsa_bittype_for_bitsize (bitsize);
695 /* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some
696 kind of array will be generated, setting DIM appropriately. Otherwise, it
697 will be set to zero. */
700 hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p = NULL,
701 bool min32int = false)
703 gcc_checking_assert (TYPE_P (type));
704 if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type)))
706 HSA_SORRY_ATV (EXPR_LOCATION (type), "support for HSA does not "
707 "implement huge or variable-sized type %qT", type);
708 return BRIG_TYPE_NONE;
711 if (RECORD_OR_UNION_TYPE_P (type))
714 *dim_p = tree_to_uhwi (TYPE_SIZE_UNIT (type));
715 return BRIG_TYPE_U8 | BRIG_TYPE_ARRAY;
718 if (TREE_CODE (type) == ARRAY_TYPE)
720 /* We try to be nice and use the real base-type when this is an array of
721 scalars and only resort to an array of bytes if the type is more
724 unsigned HOST_WIDE_INT dim = 1;
726 while (TREE_CODE (type) == ARRAY_TYPE)
728 tree domain = TYPE_DOMAIN (type);
729 if (!TYPE_MIN_VALUE (domain)
730 || !TYPE_MAX_VALUE (domain)
731 || !tree_fits_shwi_p (TYPE_MIN_VALUE (domain))
732 || !tree_fits_shwi_p (TYPE_MAX_VALUE (domain)))
734 HSA_SORRY_ATV (EXPR_LOCATION (type),
735 "support for HSA does not implement array "
736 "%qT with unknown bounds", type);
737 return BRIG_TYPE_NONE;
739 HOST_WIDE_INT min = tree_to_shwi (TYPE_MIN_VALUE (domain));
740 HOST_WIDE_INT max = tree_to_shwi (TYPE_MAX_VALUE (domain));
741 dim = dim * (unsigned HOST_WIDE_INT) (max - min + 1);
742 type = TREE_TYPE (type);
746 if (RECORD_OR_UNION_TYPE_P (type))
748 dim = dim * tree_to_uhwi (TYPE_SIZE_UNIT (type));
752 res = hsa_type_for_scalar_tree_type (type, false);
756 return res | BRIG_TYPE_ARRAY;
763 return hsa_type_for_scalar_tree_type (type, min32int);
766 /* Returns true if converting from STYPE into DTYPE needs the _CVT
767 opcode. If false a normal _MOV is enough. */
770 hsa_needs_cvt (BrigType16_t dtype, BrigType16_t stype)
772 if (hsa_btype_p (dtype))
775 /* float <-> int conversions are real converts. */
776 if (hsa_type_float_p (dtype) != hsa_type_float_p (stype))
778 /* When both types have different size, then we need CVT as well. */
779 if (hsa_type_bit_size (dtype) != hsa_type_bit_size (stype))
784 /* Return declaration name if it exists or create one from UID if it does not.
785 If DECL is a local variable, make UID part of its name. */
788 hsa_get_declaration_name (tree decl)
790 if (!DECL_NAME (decl))
793 snprintf (buf, 64, "__hsa_anon_%u", DECL_UID (decl));
794 size_t len = strlen (buf);
795 char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1);
796 memcpy (copy, buf, len + 1);
801 if (TREE_CODE (decl) == FUNCTION_DECL
802 || (TREE_CODE (decl) == VAR_DECL && is_global_var (decl)))
803 name_tree = DECL_ASSEMBLER_NAME (decl);
805 name_tree = DECL_NAME (decl);
807 const char *name = IDENTIFIER_POINTER (name_tree);
808 /* User-defined assembly names have prepended asterisk symbol. */
812 if ((TREE_CODE (decl) == VAR_DECL)
813 && decl_function_context (decl))
815 size_t len = strlen (name);
816 char *buf = (char *) alloca (len + 32);
817 snprintf (buf, len + 32, "%s_%u", name, DECL_UID (decl));
819 char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1);
820 memcpy (copy, buf, len + 1);
827 /* Lookup or create the associated hsa_symbol structure with a given VAR_DECL
828 or lookup the hsa_structure corresponding to a PARM_DECL. */
831 get_symbol_for_decl (tree decl)
834 hsa_symbol dummy (BRIG_TYPE_NONE, BRIG_SEGMENT_NONE, BRIG_LINKAGE_NONE);
836 gcc_assert (TREE_CODE (decl) == PARM_DECL
837 || TREE_CODE (decl) == RESULT_DECL
838 || TREE_CODE (decl) == VAR_DECL
839 || TREE_CODE (decl) == CONST_DECL);
843 bool is_in_global_vars = ((TREE_CODE (decl) == VAR_DECL)
844 && !decl_function_context (decl));
846 if (is_in_global_vars)
847 slot = hsa_global_variable_symbols->find_slot (&dummy, INSERT);
849 slot = hsa_cfun->m_local_symbols->find_slot (&dummy, INSERT);
851 gcc_checking_assert (slot);
854 hsa_symbol *sym = (*slot);
856 /* If the symbol is problematic, mark current function also as
858 if (sym->m_seen_error)
861 /* PR hsa/70234: If a global variable was marked to be emitted,
862 but HSAIL generation of a function using the variable fails,
863 we should retry to emit the variable in context of a different
866 Iterate elements whether a symbol is already in m_global_symbols
868 if (is_in_global_vars && !sym->m_emitted_to_brig)
870 for (unsigned i = 0; i < hsa_cfun->m_global_symbols.length (); i++)
871 if (hsa_cfun->m_global_symbols[i] == sym)
873 hsa_cfun->m_global_symbols.safe_push (sym);
881 /* PARM_DECLs and RESULT_DECL should be already in m_local_symbols. */
882 gcc_assert (TREE_CODE (decl) == VAR_DECL
883 || TREE_CODE (decl) == CONST_DECL);
884 BrigAlignment8_t align = hsa_object_alignment (decl);
886 if (is_in_global_vars)
888 gcc_checking_assert (TREE_CODE (decl) != CONST_DECL);
889 sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_GLOBAL,
890 BRIG_LINKAGE_PROGRAM, true,
891 BRIG_ALLOCATION_PROGRAM, align);
892 hsa_cfun->m_global_symbols.safe_push (sym);
893 sym->fillup_for_decl (decl);
894 if (sym->m_align > align)
896 sym->m_seen_error = true;
897 HSA_SORRY_ATV (EXPR_LOCATION (decl),
898 "HSA specification requires that %E is at least "
899 "naturally aligned", decl);
904 /* As generation of efficient memory copy instructions relies
905 on alignment greater or equal to 8 bytes,
906 we need to increase alignment of all aggregate types.. */
907 if (AGGREGATE_TYPE_P (TREE_TYPE (decl)))
908 align = MAX ((BrigAlignment8_t) BRIG_ALIGNMENT_8, align);
910 BrigAllocation allocation = BRIG_ALLOCATION_AUTOMATIC;
911 BrigSegment8_t segment;
912 if (TREE_CODE (decl) == CONST_DECL)
914 segment = BRIG_SEGMENT_READONLY;
915 allocation = BRIG_ALLOCATION_AGENT;
917 else if (lookup_attribute ("hsa_group_segment",
918 DECL_ATTRIBUTES (decl)))
919 segment = BRIG_SEGMENT_GROUP;
920 else if (TREE_STATIC (decl)
921 || lookup_attribute ("hsa_global_segment",
922 DECL_ATTRIBUTES (decl)))
923 segment = BRIG_SEGMENT_GLOBAL;
925 segment = BRIG_SEGMENT_PRIVATE;
927 sym = new hsa_symbol (BRIG_TYPE_NONE, segment, BRIG_LINKAGE_FUNCTION,
928 false, allocation, align);
929 sym->fillup_for_decl (decl);
930 hsa_cfun->m_private_variables.safe_push (sym);
933 sym->m_name = hsa_get_declaration_name (decl);
939 /* For a given HSA function declaration, return a host
940 function declaration. */
943 hsa_get_host_function (tree decl)
945 hsa_function_summary *s
946 = hsa_summaries->get (cgraph_node::get_create (decl));
947 gcc_assert (s->m_kind != HSA_NONE);
948 gcc_assert (s->m_gpu_implementation_p);
950 return s->m_bound_function ? s->m_bound_function->decl : NULL;
953 /* Return true if function DECL has a host equivalent function. */
956 get_brig_function_name (tree decl)
960 hsa_function_summary *s = hsa_summaries->get (cgraph_node::get_create (d));
961 if (s->m_kind != HSA_NONE
962 && s->m_gpu_implementation_p
963 && s->m_bound_function)
964 d = s->m_bound_function->decl;
966 /* IPA split can create a function that has no host equivalent. */
970 char *name = xstrdup (hsa_get_declaration_name (d));
971 hsa_sanitize_name (name);
976 /* Create a spill symbol of type TYPE. */
979 hsa_get_spill_symbol (BrigType16_t type)
981 hsa_symbol *sym = new hsa_symbol (type, BRIG_SEGMENT_SPILL,
982 BRIG_LINKAGE_FUNCTION);
983 hsa_cfun->m_spill_symbols.safe_push (sym);
987 /* Create a symbol for a read-only string constant. */
989 hsa_get_string_cst_symbol (tree string_cst)
991 gcc_checking_assert (TREE_CODE (string_cst) == STRING_CST);
993 hsa_symbol **slot = hsa_cfun->m_string_constants_map.get (string_cst);
997 hsa_op_immed *cst = new hsa_op_immed (string_cst);
998 hsa_symbol *sym = new hsa_symbol (cst->m_type, BRIG_SEGMENT_GLOBAL,
999 BRIG_LINKAGE_MODULE, true,
1000 BRIG_ALLOCATION_AGENT);
1001 sym->m_cst_value = cst;
1002 sym->m_dim = TREE_STRING_LENGTH (string_cst);
1003 sym->m_name_number = hsa_cfun->m_global_symbols.length ();
1005 hsa_cfun->m_global_symbols.safe_push (sym);
1006 hsa_cfun->m_string_constants_map.put (string_cst, sym);
1010 /* Constructor of the ancestor of all operands. K is BRIG kind that identified
1011 what the operator is. */
1013 hsa_op_base::hsa_op_base (BrigKind16_t k)
1014 : m_next (NULL), m_brig_op_offset (0), m_kind (k)
1016 hsa_operands.safe_push (this);
1019 /* Constructor of ancestor of all operands which have a type. K is BRIG kind
1020 that identified what the operator is. T is the type of the operator. */
1022 hsa_op_with_type::hsa_op_with_type (BrigKind16_t k, BrigType16_t t)
1023 : hsa_op_base (k), m_type (t)
1028 hsa_op_with_type::get_in_type (BrigType16_t dtype, hsa_bb *hbb)
1030 if (m_type == dtype)
1035 if (hsa_needs_cvt (dtype, m_type))
1037 dest = new hsa_op_reg (dtype);
1038 hbb->append_insn (new hsa_insn_cvt (dest, this));
1040 else if (is_a <hsa_op_reg *> (this))
1042 /* In the end, HSA registers do not really have types, only sizes, so if
1043 the sizes match, we can use the register directly. */
1044 gcc_checking_assert (hsa_type_bit_size (dtype)
1045 == hsa_type_bit_size (m_type));
1050 dest = new hsa_op_reg (m_type);
1051 hbb->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV,
1052 dest->m_type, dest, this));
1054 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
1055 type of the operand must be same as type of the instruction. */
1056 dest->m_type = dtype;
1062 /* Constructor of class representing HSA immediate values. TREE_VAL is the
1063 tree representation of the immediate value. If min32int is true,
1064 always expand integer types to one that has at least 32 bits. */
1066 hsa_op_immed::hsa_op_immed (tree tree_val, bool min32int)
1067 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES,
1068 hsa_type_for_tree_type (TREE_TYPE (tree_val), NULL,
1071 if (hsa_seen_error ())
1074 gcc_checking_assert ((is_gimple_min_invariant (tree_val)
1075 && (!POINTER_TYPE_P (TREE_TYPE (tree_val))
1076 || TREE_CODE (tree_val) == INTEGER_CST))
1077 || TREE_CODE (tree_val) == CONSTRUCTOR);
1078 m_tree_value = tree_val;
1080 /* Verify that all elements of a constructor are constants. */
1081 if (TREE_CODE (m_tree_value) == CONSTRUCTOR)
1082 for (unsigned i = 0; i < CONSTRUCTOR_NELTS (m_tree_value); i++)
1084 tree v = CONSTRUCTOR_ELT (m_tree_value, i)->value;
1085 if (!CONSTANT_CLASS_P (v))
1087 HSA_SORRY_AT (EXPR_LOCATION (tree_val),
1088 "HSA ctor should have only constants");
1094 /* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
1095 integer representation of the immediate value. TYPE is BRIG type. */
1097 hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value, BrigType16_t type)
1098 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES, type),
1101 gcc_assert (hsa_type_integer_p (type));
1102 m_int_value = integer_value;
1105 hsa_op_immed::hsa_op_immed ()
1106 : hsa_op_with_type (BRIG_KIND_NONE, BRIG_TYPE_NONE)
1110 /* New operator to allocate immediate operands from obstack. */
1113 hsa_op_immed::operator new (size_t size)
1115 return obstack_alloc (&hsa_obstack, size);
1120 hsa_op_immed::~hsa_op_immed ()
1124 /* Change type of the immediate value to T. */
1127 hsa_op_immed::set_type (BrigType16_t t)
1132 /* Constructor of class representing HSA registers and pseudo-registers. T is
1133 the BRIG type of the new register. */
1135 hsa_op_reg::hsa_op_reg (BrigType16_t t)
1136 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER, t), m_gimple_ssa (NULL_TREE),
1137 m_def_insn (NULL), m_spill_sym (NULL), m_order (hsa_cfun->m_reg_count++),
1138 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
1142 /* New operator to allocate a register from obstack. */
1145 hsa_op_reg::operator new (size_t size)
1147 return obstack_alloc (&hsa_obstack, size);
1150 /* Verify register operand. */
1153 hsa_op_reg::verify_ssa ()
1155 /* Verify that each HSA register has a definition assigned.
1156 Exceptions are VAR_DECL and PARM_DECL that are a default
1158 gcc_checking_assert (m_def_insn
1159 || (m_gimple_ssa != NULL
1160 && (!SSA_NAME_VAR (m_gimple_ssa)
1161 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa))
1163 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa)));
1165 /* Verify that every use of the register is really present
1166 in an instruction. */
1167 for (unsigned i = 0; i < m_uses.length (); i++)
1169 hsa_insn_basic *use = m_uses[i];
1171 bool is_visited = false;
1172 for (unsigned j = 0; j < use->operand_count (); j++)
1174 hsa_op_base *u = use->get_op (j);
1175 hsa_op_address *addr; addr = dyn_cast <hsa_op_address *> (u);
1176 if (addr && addr->m_reg)
1181 bool r = !addr && use->op_output_p (j);
1185 error ("HSA SSA name defined by instruction that is supposed "
1187 debug_hsa_operand (this);
1188 debug_hsa_insn (use);
1189 internal_error ("HSA SSA verification failed");
1198 error ("HSA SSA name not among operands of instruction that is "
1199 "supposed to use it");
1200 debug_hsa_operand (this);
1201 debug_hsa_insn (use);
1202 internal_error ("HSA SSA verification failed");
1207 hsa_op_address::hsa_op_address (hsa_symbol *sym, hsa_op_reg *r,
1208 HOST_WIDE_INT offset)
1209 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (r),
1210 m_imm_offset (offset)
1214 hsa_op_address::hsa_op_address (hsa_symbol *sym, HOST_WIDE_INT offset)
1215 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (NULL),
1216 m_imm_offset (offset)
1220 hsa_op_address::hsa_op_address (hsa_op_reg *r, HOST_WIDE_INT offset)
1221 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (NULL), m_reg (r),
1222 m_imm_offset (offset)
1226 /* New operator to allocate address operands from obstack. */
1229 hsa_op_address::operator new (size_t size)
1231 return obstack_alloc (&hsa_obstack, size);
1234 /* Constructor of an operand referring to HSAIL code. */
1236 hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF),
1237 m_directive_offset (0)
1241 /* Constructor of an operand representing a code list. Set it up so that it
1242 can contain ELEMENTS number of elements. */
1244 hsa_op_code_list::hsa_op_code_list (unsigned elements)
1245 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST)
1247 m_offsets.create (1);
1248 m_offsets.safe_grow_cleared (elements);
1251 /* New operator to allocate code list operands from obstack. */
1254 hsa_op_code_list::operator new (size_t size)
1256 return obstack_alloc (&hsa_obstack, size);
1259 /* Constructor of an operand representing an operand list.
1260 Set it up so that it can contain ELEMENTS number of elements. */
1262 hsa_op_operand_list::hsa_op_operand_list (unsigned elements)
1263 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST)
1265 m_offsets.create (elements);
1266 m_offsets.safe_grow (elements);
1269 /* New operator to allocate operand list operands from obstack. */
1272 hsa_op_operand_list::operator new (size_t size)
1274 return obstack_alloc (&hsa_obstack, size);
1277 hsa_op_operand_list::~hsa_op_operand_list ()
1279 m_offsets.release ();
1284 hsa_function_representation::reg_for_gimple_ssa (tree ssa)
1288 gcc_checking_assert (TREE_CODE (ssa) == SSA_NAME);
1289 if (m_ssa_map[SSA_NAME_VERSION (ssa)])
1290 return m_ssa_map[SSA_NAME_VERSION (ssa)];
1292 hreg = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa),
1294 hreg->m_gimple_ssa = ssa;
1295 m_ssa_map[SSA_NAME_VERSION (ssa)] = hreg;
1301 hsa_op_reg::set_definition (hsa_insn_basic *insn)
1303 if (hsa_cfun->m_in_ssa)
1305 gcc_checking_assert (!m_def_insn);
1312 /* Constructor of the class which is the bases of all instructions and directly
1313 represents the most basic ones. NOPS is the number of operands that the
1314 operand vector will contain (and which will be cleared). OP is the opcode
1315 of the instruction. This constructor does not set type. */
1317 hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc)
1319 m_next (NULL), m_bb (NULL), m_opcode (opc), m_number (0),
1320 m_type (BRIG_TYPE_NONE), m_brig_offset (0)
1323 m_operands.safe_grow_cleared (nops);
1325 hsa_instructions.safe_push (this);
1328 /* Make OP the operand number INDEX of operands of this instruction. If OP is a
1329 register or an address containing a register, then either set the definition
1330 of the register to this instruction if it an output operand or add this
1331 instruction to the uses if it is an input one. */
1334 hsa_insn_basic::set_op (int index, hsa_op_base *op)
1336 /* Each address operand is always use. */
1337 hsa_op_address *addr = dyn_cast <hsa_op_address *> (op);
1338 if (addr && addr->m_reg)
1339 addr->m_reg->m_uses.safe_push (this);
1342 hsa_op_reg *reg = dyn_cast <hsa_op_reg *> (op);
1345 if (op_output_p (index))
1346 reg->set_definition (this);
1348 reg->m_uses.safe_push (this);
1352 m_operands[index] = op;
1355 /* Get INDEX-th operand of the instruction. */
1358 hsa_insn_basic::get_op (int index)
1360 return m_operands[index];
1363 /* Get address of INDEX-th operand of the instruction. */
1366 hsa_insn_basic::get_op_addr (int index)
1368 return &m_operands[index];
1371 /* Get number of operands of the instruction. */
1373 hsa_insn_basic::operand_count ()
1375 return m_operands.length ();
1378 /* Constructor of the class which is the bases of all instructions and directly
1379 represents the most basic ones. NOPS is the number of operands that the
1380 operand vector will contain (and which will be cleared). OPC is the opcode
1381 of the instruction, T is the type of the instruction. */
1383 hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc, BrigType16_t t,
1384 hsa_op_base *arg0, hsa_op_base *arg1,
1385 hsa_op_base *arg2, hsa_op_base *arg3)
1386 : m_prev (NULL), m_next (NULL), m_bb (NULL), m_opcode (opc),m_number (0),
1387 m_type (t), m_brig_offset (0)
1390 m_operands.safe_grow_cleared (nops);
1394 gcc_checking_assert (nops >= 1);
1400 gcc_checking_assert (nops >= 2);
1406 gcc_checking_assert (nops >= 3);
1412 gcc_checking_assert (nops >= 4);
1416 hsa_instructions.safe_push (this);
1419 /* New operator to allocate basic instruction from obstack. */
1422 hsa_insn_basic::operator new (size_t size)
1424 return obstack_alloc (&hsa_obstack, size);
1427 /* Verify the instruction. */
1430 hsa_insn_basic::verify ()
1432 hsa_op_address *addr;
1435 /* Iterate all register operands and verify that the instruction
1436 is set in uses of the register. */
1437 for (unsigned i = 0; i < operand_count (); i++)
1439 hsa_op_base *use = get_op (i);
1441 if ((addr = dyn_cast <hsa_op_address *> (use)) && addr->m_reg)
1443 gcc_assert (addr->m_reg->m_def_insn != this);
1447 if ((reg = dyn_cast <hsa_op_reg *> (use)) && !op_output_p (i))
1450 for (j = 0; j < reg->m_uses.length (); j++)
1452 if (reg->m_uses[j] == this)
1456 if (j == reg->m_uses.length ())
1458 error ("HSA instruction uses a register but is not among "
1459 "recorded register uses");
1460 debug_hsa_operand (reg);
1461 debug_hsa_insn (this);
1462 internal_error ("HSA instruction verification failed");
1468 /* Constructor of an instruction representing a PHI node. NOPS is the number
1469 of operands (equal to the number of predecessors). */
1471 hsa_insn_phi::hsa_insn_phi (unsigned nops, hsa_op_reg *dst)
1472 : hsa_insn_basic (nops, HSA_OPCODE_PHI), m_dest (dst)
1474 dst->set_definition (this);
1477 /* Constructor of class representing instructions for control flow and
1480 hsa_insn_br::hsa_insn_br (unsigned nops, int opc, BrigType16_t t,
1481 BrigWidth8_t width, hsa_op_base *arg0,
1482 hsa_op_base *arg1, hsa_op_base *arg2,
1484 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
1489 /* Constructor of class representing instruction for conditional jump, CTRL is
1490 the control register determining whether the jump will be carried out, the
1491 new instruction is automatically added to its uses list. */
1493 hsa_insn_cbr::hsa_insn_cbr (hsa_op_reg *ctrl)
1494 : hsa_insn_br (1, BRIG_OPCODE_CBR, BRIG_TYPE_B1, BRIG_WIDTH_1, ctrl)
1498 /* Constructor of class representing instruction for switch jump, CTRL is
1499 the index register. */
1501 hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg *index, unsigned jump_count)
1502 : hsa_insn_basic (1, BRIG_OPCODE_SBR, BRIG_TYPE_B1, index),
1503 m_width (BRIG_WIDTH_1), m_jump_table (vNULL),
1504 m_label_code_list (new hsa_op_code_list (jump_count))
1508 /* Replace all occurrences of OLD_BB with NEW_BB in the statements
1512 hsa_insn_sbr::replace_all_labels (basic_block old_bb, basic_block new_bb)
1514 for (unsigned i = 0; i < m_jump_table.length (); i++)
1515 if (m_jump_table[i] == old_bb)
1516 m_jump_table[i] = new_bb;
1519 hsa_insn_sbr::~hsa_insn_sbr ()
1521 m_jump_table.release ();
1524 /* Constructor of comparison instruction. CMP is the comparison operation and T
1525 is the result type. */
1527 hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp, BrigType16_t t,
1528 hsa_op_base *arg0, hsa_op_base *arg1,
1530 : hsa_insn_basic (3 , BRIG_OPCODE_CMP, t, arg0, arg1, arg2), m_compare (cmp)
1534 /* Constructor of classes representing memory accesses. OPC is the opcode (must
1535 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
1536 operands are provided as ARG0 and ARG1. */
1538 hsa_insn_mem::hsa_insn_mem (int opc, BrigType16_t t, hsa_op_base *arg0,
1540 : hsa_insn_basic (2, opc, t, arg0, arg1),
1541 m_align (hsa_natural_alignment (t)), m_equiv_class (0)
1543 gcc_checking_assert (opc == BRIG_OPCODE_LD || opc == BRIG_OPCODE_ST);
1546 /* Constructor for descendants allowing different opcodes and number of
1547 operands, it passes its arguments directly to hsa_insn_basic
1548 constructor. The instruction operands are provided as ARG[0-3]. */
1551 hsa_insn_mem::hsa_insn_mem (unsigned nops, int opc, BrigType16_t t,
1552 hsa_op_base *arg0, hsa_op_base *arg1,
1553 hsa_op_base *arg2, hsa_op_base *arg3)
1554 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
1555 m_align (hsa_natural_alignment (t)), m_equiv_class (0)
1559 /* Constructor of class representing atomic instructions. OPC is the principal
1560 opcode, AOP is the specific atomic operation opcode. T is the type of the
1561 instruction. The instruction operands are provided as ARG[0-3]. */
1563 hsa_insn_atomic::hsa_insn_atomic (int nops, int opc,
1564 enum BrigAtomicOperation aop,
1565 BrigType16_t t, BrigMemoryOrder memorder,
1567 hsa_op_base *arg1, hsa_op_base *arg2,
1569 : hsa_insn_mem (nops, opc, t, arg0, arg1, arg2, arg3), m_atomicop (aop),
1570 m_memoryorder (memorder),
1571 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM)
1573 gcc_checking_assert (opc == BRIG_OPCODE_ATOMICNORET ||
1574 opc == BRIG_OPCODE_ATOMIC ||
1575 opc == BRIG_OPCODE_SIGNAL ||
1576 opc == BRIG_OPCODE_SIGNALNORET);
1579 /* Constructor of class representing signal instructions. OPC is the prinicpal
1580 opcode, SOP is the specific signal operation opcode. T is the type of the
1581 instruction. The instruction operands are provided as ARG[0-3]. */
1583 hsa_insn_signal::hsa_insn_signal (int nops, int opc,
1584 enum BrigAtomicOperation sop,
1585 BrigType16_t t, BrigMemoryOrder memorder,
1586 hsa_op_base *arg0, hsa_op_base *arg1,
1587 hsa_op_base *arg2, hsa_op_base *arg3)
1588 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
1589 m_memory_order (memorder), m_signalop (sop)
1593 /* Constructor of class representing segment conversion instructions. OPC is
1594 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
1595 and SRCT are destination and source types respectively, SEG is the segment
1596 we are converting to or from. The instruction operands are
1597 provided as ARG0 and ARG1. */
1599 hsa_insn_seg::hsa_insn_seg (int opc, BrigType16_t dest, BrigType16_t srct,
1600 BrigSegment8_t seg, hsa_op_base *arg0,
1602 : hsa_insn_basic (2, opc, dest, arg0, arg1), m_src_type (srct),
1605 gcc_checking_assert (opc == BRIG_OPCODE_STOF || opc == BRIG_OPCODE_FTOS);
1608 /* Constructor of class representing a call instruction. CALLEE is the tree
1609 representation of the function being called. */
1611 hsa_insn_call::hsa_insn_call (tree callee)
1612 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (callee),
1613 m_output_arg (NULL), m_args_code_list (NULL), m_result_code_list (NULL)
1617 hsa_insn_call::hsa_insn_call (hsa_internal_fn *fn)
1618 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (NULL),
1619 m_called_internal_fn (fn), m_output_arg (NULL), m_args_code_list (NULL),
1620 m_result_code_list (NULL)
1624 hsa_insn_call::~hsa_insn_call ()
1626 for (unsigned i = 0; i < m_input_args.length (); i++)
1627 delete m_input_args[i];
1629 delete m_output_arg;
1631 m_input_args.release ();
1632 m_input_arg_insns.release ();
1635 /* Constructor of class representing the argument block required to invoke
1637 hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind,
1638 hsa_insn_call * call)
1639 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK), m_kind (brig_kind),
1644 hsa_insn_comment::hsa_insn_comment (const char *s)
1645 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT)
1647 unsigned l = strlen (s);
1649 /* Append '// ' to the string. */
1650 char *buf = XNEWVEC (char, l + 4);
1651 sprintf (buf, "// %s", s);
1655 hsa_insn_comment::~hsa_insn_comment ()
1657 gcc_checking_assert (m_comment);
1662 /* Constructor of class representing the queue instruction in HSAIL. */
1664 hsa_insn_queue::hsa_insn_queue (int nops, int opcode, BrigSegment segment,
1665 BrigMemoryOrder memory_order,
1666 hsa_op_base *arg0, hsa_op_base *arg1,
1667 hsa_op_base *arg2, hsa_op_base *arg3)
1668 : hsa_insn_basic (nops, opcode, BRIG_TYPE_U64, arg0, arg1, arg2, arg3),
1669 m_segment (segment), m_memory_order (memory_order)
1673 /* Constructor of class representing the source type instruction in HSAIL. */
1675 hsa_insn_srctype::hsa_insn_srctype (int nops, BrigOpcode opcode,
1676 BrigType16_t destt, BrigType16_t srct,
1677 hsa_op_base *arg0, hsa_op_base *arg1,
1678 hsa_op_base *arg2 = NULL)
1679 : hsa_insn_basic (nops, opcode, destt, arg0, arg1, arg2),
1680 m_source_type (srct)
1683 /* Constructor of class representing the packed instruction in HSAIL. */
1685 hsa_insn_packed::hsa_insn_packed (int nops, BrigOpcode opcode,
1686 BrigType16_t destt, BrigType16_t srct,
1687 hsa_op_base *arg0, hsa_op_base *arg1,
1689 : hsa_insn_srctype (nops, opcode, destt, srct, arg0, arg1, arg2)
1691 m_operand_list = new hsa_op_operand_list (nops - 1);
1694 /* Constructor of class representing the convert instruction in HSAIL. */
1696 hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type *dest, hsa_op_with_type *src)
1697 : hsa_insn_basic (2, BRIG_OPCODE_CVT, dest->m_type, dest, src)
1701 /* Constructor of class representing the alloca in HSAIL. */
1703 hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type *dest,
1704 hsa_op_with_type *size, unsigned alignment)
1705 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA, dest->m_type, dest, size),
1706 m_align (BRIG_ALIGNMENT_8)
1708 gcc_assert (dest->m_type == BRIG_TYPE_U32);
1710 m_align = hsa_alignment_encoding (alignment);
1713 /* Append an instruction INSN into the basic block. */
1716 hsa_bb::append_insn (hsa_insn_basic *insn)
1718 gcc_assert (insn->m_opcode != 0 || insn->operand_count () == 0);
1719 gcc_assert (!insn->m_bb);
1722 insn->m_prev = m_last_insn;
1723 insn->m_next = NULL;
1725 m_last_insn->m_next = insn;
1728 m_first_insn = insn;
1732 hsa_bb::append_phi (hsa_insn_phi *hphi)
1736 hphi->m_prev = m_last_phi;
1737 hphi->m_next = NULL;
1739 m_last_phi->m_next = hphi;
1745 /* Insert HSA instruction NEW_INSN immediately before an existing instruction
1749 hsa_insert_insn_before (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
1751 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
1753 if (hbb->m_first_insn == old_insn)
1754 hbb->m_first_insn = new_insn;
1755 new_insn->m_prev = old_insn->m_prev;
1756 new_insn->m_next = old_insn;
1757 if (old_insn->m_prev)
1758 old_insn->m_prev->m_next = new_insn;
1759 old_insn->m_prev = new_insn;
1762 /* Append HSA instruction NEW_INSN immediately after an existing instruction
1766 hsa_append_insn_after (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
1768 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
1770 if (hbb->m_last_insn == old_insn)
1771 hbb->m_last_insn = new_insn;
1772 new_insn->m_prev = old_insn;
1773 new_insn->m_next = old_insn->m_next;
1774 if (old_insn->m_next)
1775 old_insn->m_next->m_prev = new_insn;
1776 old_insn->m_next = new_insn;
1779 /* Return a register containing the calculated value of EXP which must be an
1780 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
1781 integer constants as returned by get_inner_reference.
1782 Newly generated HSA instructions will be appended to HBB.
1783 Perform all calculations in ADDRTYPE. */
1785 static hsa_op_with_type *
1786 gen_address_calculation (tree exp, hsa_bb *hbb, BrigType16_t addrtype)
1790 if (TREE_CODE (exp) == NOP_EXPR)
1791 exp = TREE_OPERAND (exp, 0);
1793 switch (TREE_CODE (exp))
1796 return hsa_cfun->reg_for_gimple_ssa (exp)->get_in_type (addrtype, hbb);
1800 hsa_op_immed *imm = new hsa_op_immed (exp);
1801 if (addrtype != imm->m_type)
1802 imm->m_type = addrtype;
1807 opcode = BRIG_OPCODE_ADD;
1811 opcode = BRIG_OPCODE_MUL;
1818 hsa_op_reg *res = new hsa_op_reg (addrtype);
1819 hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, addrtype);
1820 insn->set_op (0, res);
1822 hsa_op_with_type *op1 = gen_address_calculation (TREE_OPERAND (exp, 0), hbb,
1824 hsa_op_with_type *op2 = gen_address_calculation (TREE_OPERAND (exp, 1), hbb,
1826 insn->set_op (1, op1);
1827 insn->set_op (2, op2);
1829 hbb->append_insn (insn);
1833 /* If R1 is NULL, just return R2, otherwise append an instruction adding them
1834 to HBB and return the register holding the result. */
1837 add_addr_regs_if_needed (hsa_op_reg *r1, hsa_op_reg *r2, hsa_bb *hbb)
1839 gcc_checking_assert (r2);
1843 hsa_op_reg *res = new hsa_op_reg (r1->m_type);
1844 gcc_assert (!hsa_needs_cvt (r1->m_type, r2->m_type));
1845 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_ADD, res->m_type);
1846 insn->set_op (0, res);
1847 insn->set_op (1, r1);
1848 insn->set_op (2, r2);
1849 hbb->append_insn (insn);
1853 /* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
1854 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
1857 process_mem_base (tree base, hsa_symbol **symbol, BrigType16_t *addrtype,
1858 hsa_op_reg **reg, offset_int *offset, hsa_bb *hbb)
1860 if (TREE_CODE (base) == SSA_NAME)
1863 hsa_op_with_type *ssa
1864 = hsa_cfun->reg_for_gimple_ssa (base)->get_in_type (*addrtype, hbb);
1865 *reg = dyn_cast <hsa_op_reg *> (ssa);
1867 else if (TREE_CODE (base) == ADDR_EXPR)
1869 tree decl = TREE_OPERAND (base, 0);
1871 if (!DECL_P (decl) || TREE_CODE (decl) == FUNCTION_DECL)
1873 HSA_SORRY_AT (EXPR_LOCATION (base),
1874 "support for HSA does not implement a memory reference "
1875 "to a non-declaration type");
1879 gcc_assert (!*symbol);
1881 *symbol = get_symbol_for_decl (decl);
1882 *addrtype = hsa_get_segment_addr_type ((*symbol)->m_segment);
1884 else if (TREE_CODE (base) == INTEGER_CST)
1885 *offset += wi::to_offset (base);
1890 /* Forward declaration of a function. */
1893 gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb);
1895 /* Generate HSA address operand for a given tree memory reference REF. If
1896 instructions need to be created to calculate the address, they will be added
1897 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
1898 the function assumes that the caller will handle possible
1899 bit-field references. Otherwise if we reference a bit-field, sorry message
1902 static hsa_op_address *
1903 gen_hsa_addr (tree ref, hsa_bb *hbb, HOST_WIDE_INT *output_bitsize = NULL,
1904 HOST_WIDE_INT *output_bitpos = NULL)
1906 hsa_symbol *symbol = NULL;
1907 hsa_op_reg *reg = NULL;
1908 offset_int offset = 0;
1910 tree varoffset = NULL_TREE;
1911 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
1912 HOST_WIDE_INT bitsize = 0, bitpos = 0;
1913 BrigType16_t flat_addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
1915 if (TREE_CODE (ref) == STRING_CST)
1917 symbol = hsa_get_string_cst_symbol (ref);
1920 else if (TREE_CODE (ref) == BIT_FIELD_REF
1921 && ((tree_to_uhwi (TREE_OPERAND (ref, 1)) % BITS_PER_UNIT) != 0
1922 || (tree_to_uhwi (TREE_OPERAND (ref, 2)) % BITS_PER_UNIT) != 0))
1924 HSA_SORRY_ATV (EXPR_LOCATION (origref),
1925 "support for HSA does not implement "
1926 "bit field references such as %E", ref);
1930 if (handled_component_p (ref))
1933 int unsignedp, volatilep, preversep;
1935 ref = get_inner_reference (ref, &bitsize, &bitpos, &varoffset, &mode,
1936 &unsignedp, &preversep, &volatilep);
1939 offset = wi::rshift (offset, LOG2_BITS_PER_UNIT, SIGNED);
1942 switch (TREE_CODE (ref))
1946 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
1947 symbol = hsa_cfun->create_hsa_temporary (flat_addrtype);
1948 hsa_op_reg *r = new hsa_op_reg (flat_addrtype);
1949 gen_hsa_addr_insns (ref, r, hbb);
1950 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
1951 r, new hsa_op_address (symbol)));
1957 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
1958 symbol = hsa_cfun->create_hsa_temporary (flat_addrtype);
1959 hsa_op_reg *r = hsa_cfun->reg_for_gimple_ssa (ref);
1961 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
1962 r, new hsa_op_address (symbol)));
1970 gcc_assert (!symbol);
1971 symbol = get_symbol_for_decl (ref);
1972 addrtype = hsa_get_segment_addr_type (symbol->m_segment);
1976 process_mem_base (TREE_OPERAND (ref, 0), &symbol, &addrtype, ®,
1979 if (!integer_zerop (TREE_OPERAND (ref, 1)))
1980 offset += wi::to_offset (TREE_OPERAND (ref, 1));
1983 case TARGET_MEM_REF:
1984 process_mem_base (TMR_BASE (ref), &symbol, &addrtype, ®, &offset, hbb);
1985 if (TMR_INDEX (ref))
1988 hsa_op_base *idx = hsa_cfun->reg_for_gimple_ssa
1989 (TMR_INDEX (ref))->get_in_type (addrtype, hbb);
1990 if (TMR_STEP (ref) && !integer_onep (TMR_STEP (ref)))
1992 disp1 = new hsa_op_reg (addrtype);
1993 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_MUL,
1996 /* As step must respect addrtype, we overwrite the type
1997 of an immediate value. */
1998 hsa_op_immed *step = new hsa_op_immed (TMR_STEP (ref));
1999 step->m_type = addrtype;
2001 insn->set_op (0, disp1);
2002 insn->set_op (1, idx);
2003 insn->set_op (2, step);
2004 hbb->append_insn (insn);
2007 disp1 = as_a <hsa_op_reg *> (idx);
2008 reg = add_addr_regs_if_needed (reg, disp1, hbb);
2010 if (TMR_INDEX2 (ref))
2012 if (TREE_CODE (TMR_INDEX2 (ref)) == SSA_NAME)
2014 hsa_op_base *disp2 = hsa_cfun->reg_for_gimple_ssa
2015 (TMR_INDEX2 (ref))->get_in_type (addrtype, hbb);
2016 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (disp2),
2019 else if (TREE_CODE (TMR_INDEX2 (ref)) == INTEGER_CST)
2020 offset += wi::to_offset (TMR_INDEX2 (ref));
2024 offset += wi::to_offset (TMR_OFFSET (ref));
2027 HSA_SORRY_AT (EXPR_LOCATION (origref),
2028 "support for HSA does not implement function pointers");
2031 HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does "
2032 "not implement memory access to %E", origref);
2038 if (TREE_CODE (varoffset) == INTEGER_CST)
2039 offset += wi::to_offset (varoffset);
2042 hsa_op_base *off_op = gen_address_calculation (varoffset, hbb,
2044 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (off_op),
2049 gcc_checking_assert ((symbol
2051 == hsa_get_segment_addr_type (symbol->m_segment))
2054 == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT)));
2056 HOST_WIDE_INT hwi_offset = offset.to_shwi ();
2058 /* Calculate remaining bitsize offset (if presented). */
2059 bitpos %= BITS_PER_UNIT;
2060 /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
2061 is not a reason to think this is a bit-field access. */
2063 && (bitsize >= BITS_PER_UNIT)
2064 && !(bitsize & (bitsize - 1)))
2067 if ((bitpos || bitsize) && (output_bitpos == NULL || output_bitsize == NULL))
2068 HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does not "
2069 "implement unhandled bit field reference such as %E", ref);
2071 if (output_bitsize != NULL && output_bitpos != NULL)
2073 *output_bitsize = bitsize;
2074 *output_bitpos = bitpos;
2077 return new hsa_op_address (symbol, reg, hwi_offset);
2080 /* Generate HSA address operand for a given tree memory reference REF. If
2081 instructions need to be created to calculate the address, they will be added
2082 to the end of HBB. OUTPUT_ALIGN is alignment of the created address. */
2084 static hsa_op_address *
2085 gen_hsa_addr_with_align (tree ref, hsa_bb *hbb, BrigAlignment8_t *output_align)
2087 hsa_op_address *addr = gen_hsa_addr (ref, hbb);
2088 if (addr->m_reg || !addr->m_symbol)
2089 *output_align = hsa_object_alignment (ref);
2092 /* If the address consists only of a symbol and an offset, we
2093 compute the alignment ourselves to take into account any alignment
2094 promotions we might have done for the HSA symbol representation. */
2095 unsigned align = hsa_byte_alignment (addr->m_symbol->m_align);
2096 unsigned misalign = addr->m_imm_offset & (align - 1);
2098 align = least_bit_hwi (misalign);
2099 *output_align = hsa_alignment_encoding (BITS_PER_UNIT * align);
2104 /* Generate HSA address for a function call argument of given TYPE.
2105 INDEX is used to generate corresponding name of the arguments.
2106 Special value -1 represents fact that result value is created. */
2108 static hsa_op_address *
2109 gen_hsa_addr_for_arg (tree tree_type, int index)
2111 hsa_symbol *sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
2113 sym->m_type = hsa_type_for_tree_type (tree_type, &sym->m_dim);
2115 if (index == -1) /* Function result. */
2116 sym->m_name = "res";
2117 else /* Function call arguments. */
2120 sym->m_name_number = index;
2123 return new hsa_op_address (sym);
2126 /* Generate HSA instructions that process all necessary conversions
2127 of an ADDR to flat addressing and place the result into DEST.
2128 Instructions are appended to HBB. */
2131 convert_addr_to_flat_segment (hsa_op_address *addr, hsa_op_reg *dest,
2134 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_LDA);
2135 insn->set_op (1, addr);
2136 if (addr->m_symbol && addr->m_symbol->m_segment != BRIG_SEGMENT_GLOBAL)
2138 /* LDA produces segment-relative address, we need to convert
2139 it to the flat one. */
2141 tmp = new hsa_op_reg (hsa_get_segment_addr_type
2142 (addr->m_symbol->m_segment));
2144 seg = new hsa_insn_seg (BRIG_OPCODE_STOF,
2145 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT),
2146 tmp->m_type, addr->m_symbol->m_segment, dest,
2149 insn->set_op (0, tmp);
2150 insn->m_type = tmp->m_type;
2151 hbb->append_insn (insn);
2152 hbb->append_insn (seg);
2156 insn->set_op (0, dest);
2157 insn->m_type = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
2158 hbb->append_insn (insn);
2162 /* Generate HSA instructions that calculate address of VAL including all
2163 necessary conversions to flat addressing and place the result into DEST.
2164 Instructions are appended to HBB. */
2167 gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb)
2169 /* Handle cases like tmp = NULL, where we just emit a move instruction
2171 if (TREE_CODE (val) == INTEGER_CST)
2173 hsa_op_immed *c = new hsa_op_immed (val);
2174 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
2175 dest->m_type, dest, c);
2176 hbb->append_insn (insn);
2180 hsa_op_address *addr;
2182 gcc_assert (dest->m_type == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
2183 if (TREE_CODE (val) == ADDR_EXPR)
2184 val = TREE_OPERAND (val, 0);
2185 addr = gen_hsa_addr (val, hbb);
2187 if (TREE_CODE (val) == CONST_DECL
2188 && is_gimple_reg_type (TREE_TYPE (val)))
2190 gcc_assert (addr->m_symbol
2191 && addr->m_symbol->m_segment == BRIG_SEGMENT_READONLY);
2192 /* CONST_DECLs are in readonly segment which however does not have
2193 addresses convertible to flat segments. So copy it to a private one
2194 and take address of that. */
2195 BrigType16_t csttype
2196 = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (val),
2198 hsa_op_reg *r = new hsa_op_reg (csttype);
2199 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD, csttype, r,
2200 new hsa_op_address (addr->m_symbol)));
2201 hsa_symbol *copysym = hsa_cfun->create_hsa_temporary (csttype);
2202 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, csttype, r,
2203 new hsa_op_address (copysym)));
2204 addr->m_symbol = copysym;
2206 else if (addr->m_symbol && addr->m_symbol->m_segment == BRIG_SEGMENT_READONLY)
2208 HSA_SORRY_ATV (EXPR_LOCATION (val), "support for HSA does "
2209 "not implement taking addresses of complex "
2210 "CONST_DECLs such as %E", val);
2215 convert_addr_to_flat_segment (addr, dest, hbb);
2218 /* Return an HSA register or HSA immediate value operand corresponding to
2219 gimple operand OP. */
2221 static hsa_op_with_type *
2222 hsa_reg_or_immed_for_gimple_op (tree op, hsa_bb *hbb)
2226 if (TREE_CODE (op) == SSA_NAME)
2227 tmp = hsa_cfun->reg_for_gimple_ssa (op);
2228 else if (!POINTER_TYPE_P (TREE_TYPE (op)))
2229 return new hsa_op_immed (op);
2232 tmp = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
2233 gen_hsa_addr_insns (op, tmp, hbb);
2238 /* Create a simple movement instruction with register destination DEST and
2239 register or immediate source SRC and append it to the end of HBB. */
2242 hsa_build_append_simple_mov (hsa_op_reg *dest, hsa_op_base *src, hsa_bb *hbb)
2244 /* Moves of packed data between registers need to adhere to the same type
2245 rules like when dealing with memory. */
2246 BrigType16_t tp = mem_type_for_type (dest->m_type);
2247 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, tp, dest, src);
2248 if (hsa_op_reg *sreg = dyn_cast <hsa_op_reg *> (src))
2249 gcc_assert (hsa_type_bit_size (dest->m_type)
2250 == hsa_type_bit_size (sreg->m_type));
2252 gcc_assert (hsa_type_bit_size (dest->m_type)
2253 == hsa_type_bit_size (as_a <hsa_op_immed *> (src)->m_type));
2255 hbb->append_insn (insn);
2258 /* Generate HSAIL instructions loading a bit field into register DEST.
2259 VALUE_REG is a register of a SSA name that is used in the bit field
2260 reference. To identify a bit field BITPOS is offset to the loaded memory
2261 and BITSIZE is number of bits of the bit field.
2262 Add instructions to HBB. */
2265 gen_hsa_insns_for_bitfield (hsa_op_reg *dest, hsa_op_reg *value_reg,
2266 HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos,
2269 unsigned type_bitsize = hsa_type_bit_size (dest->m_type);
2270 unsigned left_shift = type_bitsize - (bitsize + bitpos);
2271 unsigned right_shift = left_shift + bitpos;
2275 hsa_op_reg *value_reg_2 = new hsa_op_reg (dest->m_type);
2276 hsa_op_immed *c = new hsa_op_immed (left_shift, BRIG_TYPE_U32);
2278 hsa_insn_basic *lshift
2279 = new hsa_insn_basic (3, BRIG_OPCODE_SHL, value_reg_2->m_type,
2280 value_reg_2, value_reg, c);
2282 hbb->append_insn (lshift);
2284 value_reg = value_reg_2;
2289 hsa_op_reg *value_reg_2 = new hsa_op_reg (dest->m_type);
2290 hsa_op_immed *c = new hsa_op_immed (right_shift, BRIG_TYPE_U32);
2292 hsa_insn_basic *rshift
2293 = new hsa_insn_basic (3, BRIG_OPCODE_SHR, value_reg_2->m_type,
2294 value_reg_2, value_reg, c);
2296 hbb->append_insn (rshift);
2298 value_reg = value_reg_2;
2301 hsa_insn_basic *assignment
2302 = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, dest, value_reg);
2303 hbb->append_insn (assignment);
2307 /* Generate HSAIL instructions loading a bit field into register DEST. ADDR is
2308 prepared memory address which is used to load the bit field. To identify a
2309 bit field BITPOS is offset to the loaded memory and BITSIZE is number of
2310 bits of the bit field. Add instructions to HBB. Load must be performed in
2314 gen_hsa_insns_for_bitfield_load (hsa_op_reg *dest, hsa_op_address *addr,
2315 HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos,
2316 hsa_bb *hbb, BrigAlignment8_t align)
2318 hsa_op_reg *value_reg = new hsa_op_reg (dest->m_type);
2319 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, dest->m_type, value_reg,
2321 mem->set_align (align);
2322 hbb->append_insn (mem);
2323 gen_hsa_insns_for_bitfield (dest, value_reg, bitsize, bitpos, hbb);
2326 /* Return the alignment of base memory accesses we issue to perform bit-field
2327 memory access REF. */
2329 static BrigAlignment8_t
2330 hsa_bitmemref_alignment (tree ref)
2332 unsigned HOST_WIDE_INT bit_offset = 0;
2336 if (TREE_CODE (ref) == BIT_FIELD_REF)
2338 if (!tree_fits_uhwi_p (TREE_OPERAND (ref, 2)))
2339 return BRIG_ALIGNMENT_1;
2340 bit_offset += tree_to_uhwi (TREE_OPERAND (ref, 2));
2342 else if (TREE_CODE (ref) == COMPONENT_REF
2343 && DECL_BIT_FIELD (TREE_OPERAND (ref, 1)))
2344 bit_offset += int_bit_position (TREE_OPERAND (ref, 1));
2347 ref = TREE_OPERAND (ref, 0);
2350 unsigned HOST_WIDE_INT bits = bit_offset % BITS_PER_UNIT;
2351 unsigned HOST_WIDE_INT byte_bits = bit_offset - bits;
2352 BrigAlignment8_t base = hsa_object_alignment (ref);
2355 return MIN (base, hsa_alignment_encoding (least_bit_hwi (byte_bits)));
2358 /* Generate HSAIL instructions loading something into register DEST. RHS is
2359 tree representation of the loaded data, which are loaded as type TYPE. Add
2360 instructions to HBB. */
2363 gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb)
2365 /* The destination SSA name will give us the type. */
2366 if (TREE_CODE (rhs) == VIEW_CONVERT_EXPR)
2367 rhs = TREE_OPERAND (rhs, 0);
2369 if (TREE_CODE (rhs) == SSA_NAME)
2371 hsa_op_reg *src = hsa_cfun->reg_for_gimple_ssa (rhs);
2372 hsa_build_append_simple_mov (dest, src, hbb);
2374 else if (is_gimple_min_invariant (rhs)
2375 || TREE_CODE (rhs) == ADDR_EXPR)
2377 if (POINTER_TYPE_P (TREE_TYPE (rhs)))
2379 if (dest->m_type != hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT))
2381 HSA_SORRY_ATV (EXPR_LOCATION (rhs),
2382 "support for HSA does not implement conversion "
2383 "of %E to the requested non-pointer type.", rhs);
2387 gen_hsa_addr_insns (rhs, dest, hbb);
2389 else if (TREE_CODE (rhs) == COMPLEX_CST)
2391 hsa_op_immed *real_part = new hsa_op_immed (TREE_REALPART (rhs));
2392 hsa_op_immed *imag_part = new hsa_op_immed (TREE_IMAGPART (rhs));
2394 hsa_op_reg *real_part_reg
2395 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type),
2397 hsa_op_reg *imag_part_reg
2398 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type),
2401 hsa_build_append_simple_mov (real_part_reg, real_part, hbb);
2402 hsa_build_append_simple_mov (imag_part_reg, imag_part, hbb);
2404 BrigType16_t src_type = hsa_bittype_for_type (real_part_reg->m_type);
2406 hsa_insn_packed *insn
2407 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type,
2408 src_type, dest, real_part_reg,
2410 hbb->append_insn (insn);
2414 hsa_op_immed *imm = new hsa_op_immed (rhs);
2415 hsa_build_append_simple_mov (dest, imm, hbb);
2418 else if (TREE_CODE (rhs) == REALPART_EXPR || TREE_CODE (rhs) == IMAGPART_EXPR)
2420 tree pack_type = TREE_TYPE (TREE_OPERAND (rhs, 0));
2422 hsa_op_reg *packed_reg
2423 = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type, true));
2425 tree complex_rhs = TREE_OPERAND (rhs, 0);
2426 gen_hsa_insns_for_load (packed_reg, complex_rhs, TREE_TYPE (complex_rhs),
2429 hsa_op_reg *real_reg
2430 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true));
2432 hsa_op_reg *imag_reg
2433 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true));
2435 BrigKind16_t brig_type = packed_reg->m_type;
2436 hsa_insn_packed *packed
2437 = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND,
2438 hsa_bittype_for_type (real_reg->m_type),
2439 brig_type, real_reg, imag_reg, packed_reg);
2441 hbb->append_insn (packed);
2443 hsa_op_reg *source = TREE_CODE (rhs) == REALPART_EXPR ?
2444 real_reg : imag_reg;
2446 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
2447 dest->m_type, dest, source);
2449 hbb->append_insn (insn);
2451 else if (TREE_CODE (rhs) == BIT_FIELD_REF
2452 && TREE_CODE (TREE_OPERAND (rhs, 0)) == SSA_NAME)
2454 tree ssa_name = TREE_OPERAND (rhs, 0);
2455 HOST_WIDE_INT bitsize = tree_to_uhwi (TREE_OPERAND (rhs, 1));
2456 HOST_WIDE_INT bitpos = tree_to_uhwi (TREE_OPERAND (rhs, 2));
2458 hsa_op_reg *imm_value = hsa_cfun->reg_for_gimple_ssa (ssa_name);
2459 gen_hsa_insns_for_bitfield (dest, imm_value, bitsize, bitpos, hbb);
2461 else if (DECL_P (rhs) || TREE_CODE (rhs) == MEM_REF
2462 || TREE_CODE (rhs) == TARGET_MEM_REF
2463 || handled_component_p (rhs))
2465 HOST_WIDE_INT bitsize, bitpos;
2467 /* Load from memory. */
2468 hsa_op_address *addr;
2469 addr = gen_hsa_addr (rhs, hbb, &bitsize, &bitpos);
2471 /* Handle load of a bit field. */
2474 HSA_SORRY_AT (EXPR_LOCATION (rhs),
2475 "support for HSA does not implement load from a bit "
2476 "field bigger than 64 bits");
2480 if (bitsize || bitpos)
2481 gen_hsa_insns_for_bitfield_load (dest, addr, bitsize, bitpos, hbb,
2482 hsa_bitmemref_alignment (rhs));
2486 /* Not dest->m_type, that's possibly extended. */
2487 mtype = mem_type_for_type (hsa_type_for_scalar_tree_type (type,
2489 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, dest,
2491 mem->set_align (hsa_object_alignment (rhs));
2492 hbb->append_insn (mem);
2496 HSA_SORRY_ATV (EXPR_LOCATION (rhs),
2497 "support for HSA does not implement loading "
2502 /* Return number of bits necessary for representation of a bit field,
2503 starting at BITPOS with size of BITSIZE. */
2506 get_bitfield_size (unsigned bitpos, unsigned bitsize)
2508 unsigned s = bitpos + bitsize;
2509 unsigned sizes[] = {8, 16, 32, 64};
2511 for (unsigned i = 0; i < 4; i++)
2519 /* Generate HSAIL instructions storing into memory. LHS is the destination of
2520 the store, SRC is the source operand. Add instructions to HBB. */
2523 gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb)
2525 HOST_WIDE_INT bitsize = 0, bitpos = 0;
2526 BrigAlignment8_t req_align;
2528 mtype = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
2530 hsa_op_address *addr;
2531 addr = gen_hsa_addr (lhs, hbb, &bitsize, &bitpos);
2533 /* Handle store to a bit field. */
2536 HSA_SORRY_AT (EXPR_LOCATION (lhs),
2537 "support for HSA does not implement store to a bit field "
2538 "bigger than 64 bits");
2542 unsigned type_bitsize = get_bitfield_size (bitpos, bitsize);
2544 /* HSAIL does not support MOV insn with 16-bits integers. */
2545 if (type_bitsize < 32)
2548 if (bitpos || (bitsize && type_bitsize != bitsize))
2550 unsigned HOST_WIDE_INT mask = 0;
2551 BrigType16_t mem_type
2552 = get_integer_type_by_bytes (type_bitsize / BITS_PER_UNIT,
2553 !TYPE_UNSIGNED (TREE_TYPE (lhs)));
2555 for (unsigned i = 0; i < type_bitsize; i++)
2556 if (i < bitpos || i >= bitpos + bitsize)
2557 mask |= ((unsigned HOST_WIDE_INT)1 << i);
2559 hsa_op_reg *value_reg = new hsa_op_reg (mem_type);
2561 req_align = hsa_bitmemref_alignment (lhs);
2562 /* Load value from memory. */
2563 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mem_type,
2565 mem->set_align (req_align);
2566 hbb->append_insn (mem);
2568 /* AND the loaded value with prepared mask. */
2569 hsa_op_reg *cleared_reg = new hsa_op_reg (mem_type);
2572 = get_integer_type_by_bytes (type_bitsize / BITS_PER_UNIT, false);
2573 hsa_op_immed *c = new hsa_op_immed (mask, t);
2575 hsa_insn_basic *clearing
2576 = new hsa_insn_basic (3, BRIG_OPCODE_AND, mem_type, cleared_reg,
2578 hbb->append_insn (clearing);
2580 /* Shift to left a value that is going to be stored. */
2581 hsa_op_reg *new_value_reg = new hsa_op_reg (mem_type);
2583 hsa_insn_basic *basic = new hsa_insn_basic (2, BRIG_OPCODE_MOV, mem_type,
2584 new_value_reg, src);
2585 hbb->append_insn (basic);
2589 hsa_op_reg *shifted_value_reg = new hsa_op_reg (mem_type);
2590 c = new hsa_op_immed (bitpos, BRIG_TYPE_U32);
2592 hsa_insn_basic *basic
2593 = new hsa_insn_basic (3, BRIG_OPCODE_SHL, mem_type,
2594 shifted_value_reg, new_value_reg, c);
2595 hbb->append_insn (basic);
2597 new_value_reg = shifted_value_reg;
2600 /* OR the prepared value with prepared chunk loaded from memory. */
2601 hsa_op_reg *prepared_reg= new hsa_op_reg (mem_type);
2602 basic = new hsa_insn_basic (3, BRIG_OPCODE_OR, mem_type, prepared_reg,
2603 new_value_reg, cleared_reg);
2604 hbb->append_insn (basic);
2610 req_align = hsa_object_alignment (lhs);
2612 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, mtype, src, addr);
2613 mem->set_align (req_align);
2615 /* The HSAIL verifier has another constraint: if the source is an immediate
2616 then it must match the destination type. If it's a register the low bits
2617 will be used for sub-word stores. We're always allocating new operands so
2618 we can modify the above in place. */
2619 if (hsa_op_immed *imm = dyn_cast <hsa_op_immed *> (src))
2621 if (!hsa_type_packed_p (imm->m_type))
2622 imm->m_type = mem->m_type;
2625 /* ...and all vector immediates apparently need to be vectors of
2627 unsigned bs = hsa_type_bit_size (imm->m_type);
2628 gcc_assert (bs == hsa_type_bit_size (mem->m_type));
2632 imm->m_type = BRIG_TYPE_U8X4;
2635 imm->m_type = BRIG_TYPE_U8X8;
2638 imm->m_type = BRIG_TYPE_U8X16;
2646 hbb->append_insn (mem);
2649 /* Generate memory copy instructions that are going to be used
2650 for copying a SRC memory to TARGET memory,
2651 represented by pointer in a register. MIN_ALIGN is minimal alignment
2652 of provided HSA addresses. */
2655 gen_hsa_memory_copy (hsa_bb *hbb, hsa_op_address *target, hsa_op_address *src,
2656 unsigned size, BrigAlignment8_t min_align)
2658 hsa_op_address *addr;
2661 unsigned offset = 0;
2662 unsigned min_byte_align = hsa_byte_alignment (min_align);
2676 if (s > min_byte_align)
2679 BrigType16_t t = get_integer_type_by_bytes (s, false);
2681 hsa_op_reg *tmp = new hsa_op_reg (t);
2682 addr = new hsa_op_address (src->m_symbol, src->m_reg,
2683 src->m_imm_offset + offset);
2684 mem = new hsa_insn_mem (BRIG_OPCODE_LD, t, tmp, addr);
2685 hbb->append_insn (mem);
2687 addr = new hsa_op_address (target->m_symbol, target->m_reg,
2688 target->m_imm_offset + offset);
2689 mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, tmp, addr);
2690 hbb->append_insn (mem);
2696 /* Create a memset mask that is created by copying a CONSTANT byte value
2697 to an integer of BYTE_SIZE bytes. */
2699 static unsigned HOST_WIDE_INT
2700 build_memset_value (unsigned HOST_WIDE_INT constant, unsigned byte_size)
2705 HOST_WIDE_INT v = constant;
2707 for (unsigned i = 1; i < byte_size; i++)
2708 v |= constant << (8 * i);
2713 /* Generate memory set instructions that are going to be used
2714 for setting a CONSTANT byte value to TARGET memory of SIZE bytes.
2715 MIN_ALIGN is minimal alignment of provided HSA addresses. */
2718 gen_hsa_memory_set (hsa_bb *hbb, hsa_op_address *target,
2719 unsigned HOST_WIDE_INT constant,
2720 unsigned size, BrigAlignment8_t min_align)
2722 hsa_op_address *addr;
2725 unsigned offset = 0;
2726 unsigned min_byte_align = hsa_byte_alignment (min_align);
2740 if (s > min_byte_align)
2743 addr = new hsa_op_address (target->m_symbol, target->m_reg,
2744 target->m_imm_offset + offset);
2746 BrigType16_t t = get_integer_type_by_bytes (s, false);
2747 HOST_WIDE_INT c = build_memset_value (constant, s);
2749 mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, new hsa_op_immed (c, t),
2751 hbb->append_insn (mem);
2757 /* Generate HSAIL instructions for a single assignment
2758 of an empty constructor to an ADDR_LHS. Constructor is passed as a
2759 tree RHS and all instructions are appended to HBB. ALIGN is
2760 alignment of the address. */
2763 gen_hsa_ctor_assignment (hsa_op_address *addr_lhs, tree rhs, hsa_bb *hbb,
2764 BrigAlignment8_t align)
2766 if (CONSTRUCTOR_NELTS (rhs))
2768 HSA_SORRY_AT (EXPR_LOCATION (rhs),
2769 "support for HSA does not implement load from constructor");
2773 unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs)));
2774 gen_hsa_memory_set (hbb, addr_lhs, 0, size, align);
2777 /* Generate HSA instructions for a single assignment of RHS to LHS.
2778 HBB is the basic block they will be appended to. */
2781 gen_hsa_insns_for_single_assignment (tree lhs, tree rhs, hsa_bb *hbb)
2783 if (TREE_CODE (lhs) == SSA_NAME)
2785 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
2786 if (hsa_seen_error ())
2789 gen_hsa_insns_for_load (dest, rhs, TREE_TYPE (lhs), hbb);
2791 else if (TREE_CODE (rhs) == SSA_NAME
2792 || (is_gimple_min_invariant (rhs) && TREE_CODE (rhs) != STRING_CST))
2794 /* Store to memory. */
2795 hsa_op_base *src = hsa_reg_or_immed_for_gimple_op (rhs, hbb);
2796 if (hsa_seen_error ())
2799 gen_hsa_insns_for_store (lhs, src, hbb);
2803 BrigAlignment8_t lhs_align;
2804 hsa_op_address *addr_lhs = gen_hsa_addr_with_align (lhs, hbb,
2807 if (TREE_CODE (rhs) == CONSTRUCTOR)
2808 gen_hsa_ctor_assignment (addr_lhs, rhs, hbb, lhs_align);
2811 BrigAlignment8_t rhs_align;
2812 hsa_op_address *addr_rhs = gen_hsa_addr_with_align (rhs, hbb,
2815 unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs)));
2816 gen_hsa_memory_copy (hbb, addr_lhs, addr_rhs, size,
2817 MIN (lhs_align, rhs_align));
2822 /* Prepend before INSN a load from spill symbol of SPILL_REG. Return the
2823 register into which we loaded. If this required another register to convert
2824 from a B1 type, return it in *PTMP2, otherwise store NULL into it. We
2825 assume we are out of SSA so the returned register does not have its
2829 hsa_spill_in (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2)
2831 hsa_symbol *spill_sym = spill_reg->m_spill_sym;
2832 hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type);
2833 hsa_op_address *addr = new hsa_op_address (spill_sym);
2835 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, spill_sym->m_type,
2837 hsa_insert_insn_before (mem, insn);
2840 if (spill_reg->m_type == BRIG_TYPE_B1)
2842 hsa_insn_basic *cvtinsn;
2844 reg = new hsa_op_reg (spill_reg->m_type);
2846 cvtinsn = new hsa_insn_cvt (reg, *ptmp2);
2847 hsa_insert_insn_before (cvtinsn, insn);
2852 /* Append after INSN a store to spill symbol of SPILL_REG. Return the register
2853 from which we stored. If this required another register to convert to a B1
2854 type, return it in *PTMP2, otherwise store NULL into it. We assume we are
2855 out of SSA so the returned register does not have its use updated. */
2858 hsa_spill_out (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2)
2860 hsa_symbol *spill_sym = spill_reg->m_spill_sym;
2861 hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type);
2862 hsa_op_address *addr = new hsa_op_address (spill_sym);
2863 hsa_op_reg *returnreg;
2867 if (spill_reg->m_type == BRIG_TYPE_B1)
2869 hsa_insn_basic *cvtinsn;
2870 *ptmp2 = new hsa_op_reg (spill_sym->m_type);
2871 reg->m_type = spill_reg->m_type;
2873 cvtinsn = new hsa_insn_cvt (*ptmp2, returnreg);
2874 hsa_append_insn_after (cvtinsn, insn);
2879 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, spill_sym->m_type, reg,
2881 hsa_append_insn_after (mem, insn);
2885 /* Generate a comparison instruction that will compare LHS and RHS with
2886 comparison specified by CODE and put result into register DEST. DEST has to
2887 have its type set already but must not have its definition set yet.
2888 Generated instructions will be added to HBB. */
2891 gen_hsa_cmp_insn_from_gimple (enum tree_code code, tree lhs, tree rhs,
2892 hsa_op_reg *dest, hsa_bb *hbb)
2894 BrigCompareOperation8_t compare;
2899 compare = BRIG_COMPARE_LT;
2902 compare = BRIG_COMPARE_LE;
2905 compare = BRIG_COMPARE_GT;
2908 compare = BRIG_COMPARE_GE;
2911 compare = BRIG_COMPARE_EQ;
2914 compare = BRIG_COMPARE_NE;
2916 case UNORDERED_EXPR:
2917 compare = BRIG_COMPARE_NAN;
2920 compare = BRIG_COMPARE_NUM;
2923 compare = BRIG_COMPARE_LTU;
2926 compare = BRIG_COMPARE_LEU;
2929 compare = BRIG_COMPARE_GTU;
2932 compare = BRIG_COMPARE_GEU;
2935 compare = BRIG_COMPARE_EQU;
2938 compare = BRIG_COMPARE_NEU;
2942 HSA_SORRY_ATV (EXPR_LOCATION (lhs),
2943 "support for HSA does not implement comparison tree "
2944 "code %s\n", get_tree_code_name (code));
2948 /* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer)
2949 as a result of comparison. */
2951 BrigType16_t dest_type = hsa_type_integer_p (dest->m_type)
2952 ? (BrigType16_t) BRIG_TYPE_B1 : dest->m_type;
2954 hsa_insn_cmp *cmp = new hsa_insn_cmp (compare, dest_type);
2955 cmp->set_op (1, hsa_reg_or_immed_for_gimple_op (lhs, hbb));
2956 cmp->set_op (2, hsa_reg_or_immed_for_gimple_op (rhs, hbb));
2958 hbb->append_insn (cmp);
2959 cmp->set_output_in_type (dest, 0, hbb);
2962 /* Generate an unary instruction with OPCODE and append it to a basic block
2963 HBB. The instruction uses DEST as a destination and OP1
2964 as a single operand. */
2967 gen_hsa_unary_operation (BrigOpcode opcode, hsa_op_reg *dest,
2968 hsa_op_with_type *op1, hsa_bb *hbb)
2970 gcc_checking_assert (dest);
2971 hsa_insn_basic *insn;
2973 if (opcode == BRIG_OPCODE_MOV && hsa_needs_cvt (dest->m_type, op1->m_type))
2974 insn = new hsa_insn_cvt (dest, op1);
2975 else if (opcode == BRIG_OPCODE_FIRSTBIT || opcode == BRIG_OPCODE_LASTBIT)
2977 BrigType16_t srctype = hsa_type_integer_p (op1->m_type) ? op1->m_type
2978 : hsa_unsigned_type_for_type (op1->m_type);
2979 insn = new hsa_insn_srctype (2, opcode, BRIG_TYPE_U32, srctype, NULL,
2984 insn = new hsa_insn_basic (2, opcode, dest->m_type, dest, op1);
2986 if (opcode == BRIG_OPCODE_ABS || opcode == BRIG_OPCODE_NEG)
2988 /* ABS and NEG only exist in _s form :-/ */
2989 if (insn->m_type == BRIG_TYPE_U32)
2990 insn->m_type = BRIG_TYPE_S32;
2991 else if (insn->m_type == BRIG_TYPE_U64)
2992 insn->m_type = BRIG_TYPE_S64;
2996 hbb->append_insn (insn);
2998 if (opcode == BRIG_OPCODE_FIRSTBIT || opcode == BRIG_OPCODE_LASTBIT)
2999 insn->set_output_in_type (dest, 0, hbb);
3002 /* Generate a binary instruction with OPCODE and append it to a basic block
3003 HBB. The instruction uses DEST as a destination and operands OP1
3007 gen_hsa_binary_operation (int opcode, hsa_op_reg *dest,
3008 hsa_op_base *op1, hsa_op_base *op2, hsa_bb *hbb)
3010 gcc_checking_assert (dest);
3012 if ((opcode == BRIG_OPCODE_SHL || opcode == BRIG_OPCODE_SHR)
3013 && is_a <hsa_op_immed *> (op2))
3015 hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2);
3016 i->set_type (BRIG_TYPE_U32);
3018 if ((opcode == BRIG_OPCODE_OR
3019 || opcode == BRIG_OPCODE_XOR
3020 || opcode == BRIG_OPCODE_AND)
3021 && is_a <hsa_op_immed *> (op2))
3023 hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2);
3024 i->set_type (hsa_unsigned_type_for_type (i->m_type));
3027 hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, dest->m_type, dest,
3029 hbb->append_insn (insn);
3032 /* Generate HSA instructions for a single assignment. HBB is the basic block
3033 they will be appended to. */
3036 gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
3038 tree_code code = gimple_assign_rhs_code (assign);
3039 gimple_rhs_class rhs_class = get_gimple_rhs_class (gimple_expr_code (assign));
3041 tree lhs = gimple_assign_lhs (assign);
3042 tree rhs1 = gimple_assign_rhs1 (assign);
3043 tree rhs2 = gimple_assign_rhs2 (assign);
3044 tree rhs3 = gimple_assign_rhs3 (assign);
3052 /* The opcode is changed to BRIG_OPCODE_CVT if BRIG types
3053 needs a conversion. */
3054 opcode = BRIG_OPCODE_MOV;
3058 case POINTER_PLUS_EXPR:
3059 opcode = BRIG_OPCODE_ADD;
3062 opcode = BRIG_OPCODE_SUB;
3065 opcode = BRIG_OPCODE_MUL;
3067 case MULT_HIGHPART_EXPR:
3068 opcode = BRIG_OPCODE_MULHI;
3071 case TRUNC_DIV_EXPR:
3072 case EXACT_DIV_EXPR:
3073 opcode = BRIG_OPCODE_DIV;
3076 case FLOOR_DIV_EXPR:
3077 case ROUND_DIV_EXPR:
3078 HSA_SORRY_AT (gimple_location (assign),
3079 "support for HSA does not implement CEIL_DIV_EXPR, "
3080 "FLOOR_DIV_EXPR or ROUND_DIV_EXPR");
3082 case TRUNC_MOD_EXPR:
3083 opcode = BRIG_OPCODE_REM;
3086 case FLOOR_MOD_EXPR:
3087 case ROUND_MOD_EXPR:
3088 HSA_SORRY_AT (gimple_location (assign),
3089 "support for HSA does not implement CEIL_MOD_EXPR, "
3090 "FLOOR_MOD_EXPR or ROUND_MOD_EXPR");
3093 opcode = BRIG_OPCODE_NEG;
3096 /* There is a native HSA instruction for scalar FMAs but not for vector
3098 if (TREE_CODE (TREE_TYPE (lhs)) == VECTOR_TYPE)
3101 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3102 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3103 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3104 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
3105 hsa_op_reg *tmp = new hsa_op_reg (dest->m_type);
3106 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp, op1, op2, hbb);
3107 gen_hsa_binary_operation (BRIG_OPCODE_ADD, dest, tmp, op3, hbb);
3110 opcode = BRIG_OPCODE_MAD;
3113 opcode = BRIG_OPCODE_MIN;
3116 opcode = BRIG_OPCODE_MAX;
3119 opcode = BRIG_OPCODE_ABS;
3122 opcode = BRIG_OPCODE_SHL;
3125 opcode = BRIG_OPCODE_SHR;
3130 hsa_insn_basic *insn = NULL;
3131 int code1 = code == LROTATE_EXPR ? BRIG_OPCODE_SHL : BRIG_OPCODE_SHR;
3132 int code2 = code != LROTATE_EXPR ? BRIG_OPCODE_SHL : BRIG_OPCODE_SHR;
3133 BrigType16_t btype = hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
3136 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3137 hsa_op_reg *op1 = new hsa_op_reg (btype);
3138 hsa_op_reg *op2 = new hsa_op_reg (btype);
3139 hsa_op_with_type *shift1 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3141 tree type = TREE_TYPE (rhs2);
3142 unsigned HOST_WIDE_INT bitsize = TREE_INT_CST_LOW (TYPE_SIZE (type));
3144 hsa_op_with_type *shift2 = NULL;
3145 if (TREE_CODE (rhs2) == INTEGER_CST)
3146 shift2 = new hsa_op_immed (bitsize - tree_to_uhwi (rhs2),
3148 else if (TREE_CODE (rhs2) == SSA_NAME)
3150 hsa_op_reg *s = hsa_cfun->reg_for_gimple_ssa (rhs2);
3151 hsa_op_reg *d = new hsa_op_reg (s->m_type);
3152 hsa_op_immed *size_imm = new hsa_op_immed (bitsize, BRIG_TYPE_U32);
3154 insn = new hsa_insn_basic (3, BRIG_OPCODE_SUB, d->m_type,
3156 hbb->append_insn (insn);
3163 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3164 gen_hsa_binary_operation (code1, op1, src, shift1, hbb);
3165 gen_hsa_binary_operation (code2, op2, src, shift2, hbb);
3166 gen_hsa_binary_operation (BRIG_OPCODE_OR, dest, op1, op2, hbb);
3171 opcode = BRIG_OPCODE_OR;
3174 opcode = BRIG_OPCODE_XOR;
3177 opcode = BRIG_OPCODE_AND;
3180 opcode = BRIG_OPCODE_NOT;
3182 case FIX_TRUNC_EXPR:
3184 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3185 hsa_op_with_type *v = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3187 if (hsa_needs_cvt (dest->m_type, v->m_type))
3189 hsa_op_reg *tmp = new hsa_op_reg (v->m_type);
3191 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_TRUNC,
3192 tmp->m_type, tmp, v);
3193 hbb->append_insn (insn);
3195 hsa_insn_basic *cvtinsn = new hsa_insn_cvt (dest, tmp);
3196 hbb->append_insn (cvtinsn);
3200 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_TRUNC,
3201 dest->m_type, dest, v);
3202 hbb->append_insn (insn);
3207 opcode = BRIG_OPCODE_TRUNC;
3216 case UNORDERED_EXPR:
3226 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3228 gen_hsa_cmp_insn_from_gimple (code, rhs1, rhs2, dest, hbb);
3234 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3235 hsa_op_with_type *ctrl = NULL;
3238 if (CONSTANT_CLASS_P (cond) || TREE_CODE (cond) == SSA_NAME)
3239 ctrl = hsa_reg_or_immed_for_gimple_op (cond, hbb);
3242 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_B1);
3244 gen_hsa_cmp_insn_from_gimple (TREE_CODE (cond),
3245 TREE_OPERAND (cond, 0),
3246 TREE_OPERAND (cond, 1),
3252 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3253 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
3255 BrigType16_t utype = hsa_unsigned_type_for_type (dest->m_type);
3256 if (is_a <hsa_op_immed *> (op2))
3257 op2->m_type = utype;
3258 if (is_a <hsa_op_immed *> (op3))
3259 op3->m_type = utype;
3261 hsa_insn_basic *insn
3262 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV,
3263 hsa_bittype_for_type (dest->m_type),
3264 dest, ctrl, op2, op3);
3266 hbb->append_insn (insn);
3272 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3273 hsa_op_with_type *rhs1_reg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3274 hsa_op_with_type *rhs2_reg = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3276 if (hsa_seen_error ())
3279 BrigType16_t src_type = hsa_bittype_for_type (rhs1_reg->m_type);
3280 rhs1_reg = rhs1_reg->get_in_type (src_type, hbb);
3281 rhs2_reg = rhs2_reg->get_in_type (src_type, hbb);
3283 hsa_insn_packed *insn
3284 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type, src_type,
3285 dest, rhs1_reg, rhs2_reg);
3286 hbb->append_insn (insn);
3291 /* Implement others as we come across them. */
3292 HSA_SORRY_ATV (gimple_location (assign),
3293 "support for HSA does not implement operation %s",
3294 get_tree_code_name (code));
3299 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3301 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3302 hsa_op_with_type *op2 = rhs2 != NULL_TREE ?
3303 hsa_reg_or_immed_for_gimple_op (rhs2, hbb) : NULL;
3305 if (hsa_seen_error ())
3310 case GIMPLE_TERNARY_RHS:
3312 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
3313 hsa_insn_basic *insn = new hsa_insn_basic (4, opcode, dest->m_type, dest,
3315 hbb->append_insn (insn);
3319 case GIMPLE_BINARY_RHS:
3320 gen_hsa_binary_operation (opcode, dest, op1, op2, hbb);
3323 case GIMPLE_UNARY_RHS:
3324 gen_hsa_unary_operation (opcode, dest, op1, hbb);
3331 /* Generate HSA instructions for a given gimple condition statement COND.
3332 Instructions will be appended to HBB, which also needs to be the
3333 corresponding structure to the basic_block of COND. */
3336 gen_hsa_insns_for_cond_stmt (gimple *cond, hsa_bb *hbb)
3338 hsa_op_reg *ctrl = new hsa_op_reg (BRIG_TYPE_B1);
3341 gen_hsa_cmp_insn_from_gimple (gimple_cond_code (cond),
3342 gimple_cond_lhs (cond),
3343 gimple_cond_rhs (cond),
3346 cbr = new hsa_insn_cbr (ctrl);
3347 hbb->append_insn (cbr);
3350 /* Maximum number of elements in a jump table for an HSA SBR instruction. */
3352 #define HSA_MAXIMUM_SBR_LABELS 16
3354 /* Return lowest value of a switch S that is handled in a non-default
3358 get_switch_low (gswitch *s)
3360 unsigned labels = gimple_switch_num_labels (s);
3361 gcc_checking_assert (labels >= 1);
3363 return CASE_LOW (gimple_switch_label (s, 1));
3366 /* Return highest value of a switch S that is handled in a non-default
3370 get_switch_high (gswitch *s)
3372 unsigned labels = gimple_switch_num_labels (s);
3374 /* Compare last label to maximum number of labels. */
3375 tree label = gimple_switch_label (s, labels - 1);
3376 tree low = CASE_LOW (label);
3377 tree high = CASE_HIGH (label);
3379 return high != NULL_TREE ? high : low;
3383 get_switch_size (gswitch *s)
3385 return int_const_binop (MINUS_EXPR, get_switch_high (s), get_switch_low (s));
3388 /* Generate HSA instructions for a given gimple switch.
3389 Instructions will be appended to HBB. */
3392 gen_hsa_insns_for_switch_stmt (gswitch *s, hsa_bb *hbb)
3394 gimple_stmt_iterator it = gsi_for_stmt (s);
3397 /* Create preambule that verifies that index - lowest_label >= 0. */
3398 edge e = split_block (hbb->m_bb, gsi_stmt (it));
3399 e->flags &= ~EDGE_FALLTHRU;
3400 e->flags |= EDGE_TRUE_VALUE;
3402 function *func = DECL_STRUCT_FUNCTION (current_function_decl);
3403 tree index_tree = gimple_switch_index (s);
3404 tree lowest = get_switch_low (s);
3405 tree highest = get_switch_high (s);
3407 hsa_op_reg *index = hsa_cfun->reg_for_gimple_ssa (index_tree);
3409 hsa_op_reg *cmp1_reg = new hsa_op_reg (BRIG_TYPE_B1);
3410 hsa_op_immed *cmp1_immed = new hsa_op_immed (lowest);
3411 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_GE, cmp1_reg->m_type,
3412 cmp1_reg, index, cmp1_immed));
3414 hsa_op_reg *cmp2_reg = new hsa_op_reg (BRIG_TYPE_B1);
3415 hsa_op_immed *cmp2_immed = new hsa_op_immed (highest);
3416 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_LE, cmp2_reg->m_type,
3417 cmp2_reg, index, cmp2_immed));
3419 hsa_op_reg *cmp_reg = new hsa_op_reg (BRIG_TYPE_B1);
3420 hbb->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_AND, cmp_reg->m_type,
3421 cmp_reg, cmp1_reg, cmp2_reg));
3423 hbb->append_insn (new hsa_insn_cbr (cmp_reg));
3425 tree default_label = gimple_switch_default_label (s);
3426 basic_block default_label_bb = label_to_block_fn (func,
3427 CASE_LABEL (default_label));
3429 if (!gimple_seq_empty_p (phi_nodes (default_label_bb)))
3431 default_label_bb = split_edge (find_edge (e->dest, default_label_bb));
3432 hsa_init_new_bb (default_label_bb);
3435 make_edge (e->src, default_label_bb, EDGE_FALSE_VALUE);
3437 hsa_cfun->m_modified_cfg = true;
3439 /* Basic block with the SBR instruction. */
3440 hbb = hsa_init_new_bb (e->dest);
3442 hsa_op_reg *sub_index = new hsa_op_reg (index->m_type);
3443 hbb->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB, sub_index->m_type,
3445 new hsa_op_immed (lowest)));
3447 hsa_op_base *tmp = sub_index->get_in_type (BRIG_TYPE_U64, hbb);
3448 sub_index = as_a <hsa_op_reg *> (tmp);
3449 unsigned labels = gimple_switch_num_labels (s);
3450 unsigned HOST_WIDE_INT size = tree_to_uhwi (get_switch_size (s));
3452 hsa_insn_sbr *sbr = new hsa_insn_sbr (sub_index, size + 1);
3454 /* Prepare array with default label destination. */
3455 for (unsigned HOST_WIDE_INT i = 0; i <= size; i++)
3456 sbr->m_jump_table.safe_push (default_label_bb);
3458 /* Iterate all labels and fill up the jump table. */
3459 for (unsigned i = 1; i < labels; i++)
3461 tree label = gimple_switch_label (s, i);
3462 basic_block bb = label_to_block_fn (func, CASE_LABEL (label));
3464 unsigned HOST_WIDE_INT sub_low
3465 = tree_to_uhwi (int_const_binop (MINUS_EXPR, CASE_LOW (label), lowest));
3467 unsigned HOST_WIDE_INT sub_high = sub_low;
3468 tree high = CASE_HIGH (label);
3470 sub_high = tree_to_uhwi (int_const_binop (MINUS_EXPR, high, lowest));
3472 for (unsigned HOST_WIDE_INT j = sub_low; j <= sub_high; j++)
3473 sbr->m_jump_table[j] = bb;
3476 hbb->append_insn (sbr);
3479 /* Verify that the function DECL can be handled by HSA. */
3482 verify_function_arguments (tree decl)
3484 tree type = TREE_TYPE (decl);
3485 if (DECL_STATIC_CHAIN (decl))
3487 HSA_SORRY_ATV (EXPR_LOCATION (decl),
3488 "HSA does not support nested functions: %qD", decl);
3491 else if (!TYPE_ARG_TYPES (type) || stdarg_p (type))
3493 HSA_SORRY_ATV (EXPR_LOCATION (decl),
3494 "HSA does not support functions with variadic arguments "
3495 "(or unknown return type): %qD", decl);
3500 /* Return BRIG type for FORMAL_ARG_TYPE. If the formal argument type is NULL,
3501 return ACTUAL_ARG_TYPE. */
3504 get_format_argument_type (tree formal_arg_type, BrigType16_t actual_arg_type)
3506 if (formal_arg_type == NULL)
3507 return actual_arg_type;
3509 BrigType16_t decl_type
3510 = hsa_type_for_scalar_tree_type (formal_arg_type, false);
3511 return mem_type_for_type (decl_type);
3514 /* Generate HSA instructions for a direct call instruction.
3515 Instructions will be appended to HBB, which also needs to be the
3516 corresponding structure to the basic_block of STMT.
3517 If ASSIGN_LHS is false, do not copy HSA function result argument into the
3518 corresponding HSA representation of the gimple statement LHS. */
3521 gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb,
3522 bool assign_lhs = true)
3524 tree decl = gimple_call_fndecl (stmt);
3525 verify_function_arguments (decl);
3526 if (hsa_seen_error ())
3529 hsa_insn_call *call_insn = new hsa_insn_call (decl);
3530 hsa_cfun->m_called_functions.safe_push (call_insn->m_called_function);
3532 /* Argument block start. */
3533 hsa_insn_arg_block *arg_start
3534 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn);
3535 hbb->append_insn (arg_start);
3537 tree parm_type_chain = TYPE_ARG_TYPES (gimple_call_fntype (stmt));
3539 /* Preparation of arguments that will be passed to function. */
3540 const unsigned args = gimple_call_num_args (stmt);
3541 for (unsigned i = 0; i < args; ++i)
3543 tree parm = gimple_call_arg (stmt, (int)i);
3544 tree parm_decl_type = parm_type_chain != NULL_TREE
3545 ? TREE_VALUE (parm_type_chain) : NULL_TREE;
3546 hsa_op_address *addr;
3548 if (AGGREGATE_TYPE_P (TREE_TYPE (parm)))
3550 addr = gen_hsa_addr_for_arg (TREE_TYPE (parm), i);
3551 BrigAlignment8_t align;
3552 hsa_op_address *src = gen_hsa_addr_with_align (parm, hbb, &align);
3553 gen_hsa_memory_copy (hbb, addr, src,
3554 addr->m_symbol->total_byte_size (), align);
3558 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (parm, hbb);
3560 if (parm_decl_type != NULL && AGGREGATE_TYPE_P (parm_decl_type))
3562 HSA_SORRY_AT (gimple_location (stmt),
3563 "support for HSA does not implement an aggregate "
3564 "formal argument in a function call, while actual "
3565 "argument is not an aggregate");
3569 BrigType16_t formal_arg_type
3570 = get_format_argument_type (parm_decl_type, src->m_type);
3571 if (hsa_seen_error ())
3574 if (src->m_type != formal_arg_type)
3575 src = src->get_in_type (formal_arg_type, hbb);
3578 = gen_hsa_addr_for_arg (parm_decl_type != NULL_TREE ?
3579 parm_decl_type: TREE_TYPE (parm), i);
3580 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, formal_arg_type,
3583 hbb->append_insn (mem);
3586 call_insn->m_input_args.safe_push (addr->m_symbol);
3587 if (parm_type_chain)
3588 parm_type_chain = TREE_CHAIN (parm_type_chain);
3591 call_insn->m_args_code_list = new hsa_op_code_list (args);
3592 hbb->append_insn (call_insn);
3594 tree result_type = TREE_TYPE (TREE_TYPE (decl));
3596 tree result = gimple_call_lhs (stmt);
3597 hsa_insn_mem *result_insn = NULL;
3598 if (!VOID_TYPE_P (result_type))
3600 hsa_op_address *addr = gen_hsa_addr_for_arg (result_type, -1);
3602 /* Even if result of a function call is unused, we have to emit
3603 declaration for the result. */
3604 if (result && assign_lhs)
3606 tree lhs_type = TREE_TYPE (result);
3608 if (hsa_seen_error ())
3611 if (AGGREGATE_TYPE_P (lhs_type))
3613 BrigAlignment8_t align;
3614 hsa_op_address *result_addr
3615 = gen_hsa_addr_with_align (result, hbb, &align);
3616 gen_hsa_memory_copy (hbb, result_addr, addr,
3617 addr->m_symbol->total_byte_size (), align);
3622 = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type,
3625 hsa_op_reg *dst = hsa_cfun->reg_for_gimple_ssa (result);
3626 result_insn = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, dst, addr);
3627 hbb->append_insn (result_insn);
3631 call_insn->m_output_arg = addr->m_symbol;
3632 call_insn->m_result_code_list = new hsa_op_code_list (1);
3638 HSA_SORRY_AT (gimple_location (stmt),
3639 "support for HSA does not implement an assignment of "
3640 "return value from a void function");
3644 call_insn->m_result_code_list = new hsa_op_code_list (0);
3647 /* Argument block end. */
3648 hsa_insn_arg_block *arg_end
3649 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn);
3650 hbb->append_insn (arg_end);
3653 /* Generate HSA instructions for a direct call of an internal fn.
3654 Instructions will be appended to HBB, which also needs to be the
3655 corresponding structure to the basic_block of STMT. */
3658 gen_hsa_insns_for_call_of_internal_fn (gimple *stmt, hsa_bb *hbb)
3660 tree lhs = gimple_call_lhs (stmt);
3664 tree lhs_type = TREE_TYPE (lhs);
3665 tree rhs1 = gimple_call_arg (stmt, 0);
3666 tree rhs1_type = TREE_TYPE (rhs1);
3667 enum internal_fn fn = gimple_call_internal_fn (stmt);
3668 hsa_internal_fn *ifn
3669 = new hsa_internal_fn (fn, tree_to_uhwi (TYPE_SIZE (rhs1_type)));
3670 hsa_insn_call *call_insn = new hsa_insn_call (ifn);
3672 gcc_checking_assert (FLOAT_TYPE_P (rhs1_type));
3674 if (!hsa_emitted_internal_decls->find (call_insn->m_called_internal_fn))
3675 hsa_cfun->m_called_internal_fns.safe_push (call_insn->m_called_internal_fn);
3677 hsa_insn_arg_block *arg_start
3678 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn);
3679 hbb->append_insn (arg_start);
3681 unsigned num_args = gimple_call_num_args (stmt);
3683 /* Function arguments. */
3684 for (unsigned i = 0; i < num_args; i++)
3686 tree parm = gimple_call_arg (stmt, (int)i);
3687 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (parm, hbb);
3689 hsa_op_address *addr = gen_hsa_addr_for_arg (TREE_TYPE (parm), i);
3690 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, src->m_type,
3693 call_insn->m_input_args.safe_push (addr->m_symbol);
3694 hbb->append_insn (mem);
3697 call_insn->m_args_code_list = new hsa_op_code_list (num_args);
3698 hbb->append_insn (call_insn);
3700 /* Assign returned value. */
3701 hsa_op_address *addr = gen_hsa_addr_for_arg (lhs_type, -1);
3703 call_insn->m_output_arg = addr->m_symbol;
3704 call_insn->m_result_code_list = new hsa_op_code_list (1);
3706 /* Argument block end. */
3707 hsa_insn_arg_block *arg_end
3708 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn);
3709 hbb->append_insn (arg_end);
3712 /* Generate HSA instructions for a return value instruction.
3713 Instructions will be appended to HBB, which also needs to be the
3714 corresponding structure to the basic_block of STMT. */
3717 gen_hsa_insns_for_return (greturn *stmt, hsa_bb *hbb)
3719 tree retval = gimple_return_retval (stmt);
3722 hsa_op_address *addr = new hsa_op_address (hsa_cfun->m_output_arg);
3724 if (AGGREGATE_TYPE_P (TREE_TYPE (retval)))
3726 BrigAlignment8_t align;
3727 hsa_op_address *retval_addr = gen_hsa_addr_with_align (retval, hbb,
3729 gen_hsa_memory_copy (hbb, addr, retval_addr,
3730 hsa_cfun->m_output_arg->total_byte_size (),
3735 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (retval),
3737 BrigType16_t mtype = mem_type_for_type (t);
3739 /* Store of return value. */
3740 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (retval, hbb);
3741 src = src->get_in_type (mtype, hbb);
3742 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, mtype, src,
3744 hbb->append_insn (mem);
3748 /* HSAIL return instruction emission. */
3749 hsa_insn_basic *ret = new hsa_insn_basic (0, BRIG_OPCODE_RET);
3750 hbb->append_insn (ret);
3753 /* Set OP_INDEX-th operand of the instruction to DEST, as the DEST
3754 can have a different type, conversion instructions are possibly
3758 hsa_insn_basic::set_output_in_type (hsa_op_reg *dest, unsigned op_index,
3761 hsa_insn_basic *insn;
3762 gcc_checking_assert (op_output_p (op_index));
3764 if (dest->m_type == m_type)
3766 set_op (op_index, dest);
3770 hsa_op_reg *tmp = new hsa_op_reg (m_type);
3771 set_op (op_index, tmp);
3773 if (hsa_needs_cvt (dest->m_type, m_type))
3774 insn = new hsa_insn_cvt (dest, tmp);
3776 insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type,
3777 dest, tmp->get_in_type (dest->m_type, hbb));
3779 hbb->append_insn (insn);
3782 /* Generate instruction OPCODE to query a property of HSA grid along the
3783 given DIMENSION. Store result into DEST and append the instruction to
3787 query_hsa_grid_dim (hsa_op_reg *dest, int opcode, hsa_op_immed *dimension,
3790 hsa_insn_basic *insn = new hsa_insn_basic (2, opcode, BRIG_TYPE_U32, NULL,
3792 hbb->append_insn (insn);
3793 insn->set_output_in_type (dest, 0, hbb);
3796 /* Generate instruction OPCODE to query a property of HSA grid along the given
3797 dimension which is an immediate in first argument of STMT. Store result
3798 into the register corresponding to LHS of STMT and append the instruction to
3802 query_hsa_grid_dim (gimple *stmt, int opcode, hsa_bb *hbb)
3804 tree lhs = gimple_call_lhs (dyn_cast <gcall *> (stmt));
3805 if (lhs == NULL_TREE)
3808 tree arg = gimple_call_arg (stmt, 0);
3809 unsigned HOST_WIDE_INT dim = 5;
3810 if (tree_fits_uhwi_p (arg))
3811 dim = tree_to_uhwi (arg);
3814 HSA_SORRY_AT (gimple_location (stmt),
3815 "HSA grid query dimension must be immediate constant 0, 1 "
3820 hsa_op_immed *hdim = new hsa_op_immed (dim, (BrigKind16_t) BRIG_TYPE_U32);
3821 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3822 query_hsa_grid_dim (dest, opcode, hdim, hbb);
3825 /* Generate instruction OPCODE to query a property of HSA grid that is
3826 independent of any dimension. Store result into the register corresponding
3827 to LHS of STMT and append the instruction to HBB. */
3830 query_hsa_grid_nodim (gimple *stmt, BrigOpcode16_t opcode, hsa_bb *hbb)
3832 tree lhs = gimple_call_lhs (dyn_cast <gcall *> (stmt));
3833 if (lhs == NULL_TREE)
3835 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3836 BrigType16_t brig_type = hsa_unsigned_type_for_type (dest->m_type);
3837 hsa_insn_basic *insn = new hsa_insn_basic (1, opcode, brig_type, dest);
3838 hbb->append_insn (insn);
3841 /* Emit instructions that set hsa_num_threads according to provided VALUE.
3842 Instructions are appended to basic block HBB. */
3845 gen_set_num_threads (tree value, hsa_bb *hbb)
3847 hbb->append_insn (new hsa_insn_comment ("omp_set_num_threads"));
3848 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (value, hbb);
3850 src = src->get_in_type (hsa_num_threads->m_type, hbb);
3851 hsa_op_address *addr = new hsa_op_address (hsa_num_threads);
3853 hsa_insn_basic *basic
3854 = new hsa_insn_mem (BRIG_OPCODE_ST, hsa_num_threads->m_type, src, addr);
3855 hbb->append_insn (basic);
3858 /* Return byte offset of a FIELD_NAME in GOMP_hsa_kernel_dispatch which
3859 is defined in plugin-hsa.c. */
3861 static HOST_WIDE_INT
3862 get_hsa_kernel_dispatch_offset (const char *field_name)
3864 tree *hsa_kernel_dispatch_type = hsa_get_kernel_dispatch_type ();
3865 if (*hsa_kernel_dispatch_type == NULL)
3867 /* Collection of information needed for a dispatch of a kernel from a
3868 kernel. Keep in sync with libgomp's plugin-hsa.c. */
3870 *hsa_kernel_dispatch_type = make_node (RECORD_TYPE);
3871 tree id_f1 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3872 get_identifier ("queue"), ptr_type_node);
3873 DECL_CHAIN (id_f1) = NULL_TREE;
3874 tree id_f2 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3875 get_identifier ("omp_data_memory"),
3877 DECL_CHAIN (id_f2) = id_f1;
3878 tree id_f3 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3879 get_identifier ("kernarg_address"),
3881 DECL_CHAIN (id_f3) = id_f2;
3882 tree id_f4 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3883 get_identifier ("object"),
3885 DECL_CHAIN (id_f4) = id_f3;
3886 tree id_f5 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3887 get_identifier ("signal"),
3889 DECL_CHAIN (id_f5) = id_f4;
3890 tree id_f6 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3891 get_identifier ("private_segment_size"),
3893 DECL_CHAIN (id_f6) = id_f5;
3894 tree id_f7 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3895 get_identifier ("group_segment_size"),
3897 DECL_CHAIN (id_f7) = id_f6;
3898 tree id_f8 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3899 get_identifier ("kernel_dispatch_count"),
3901 DECL_CHAIN (id_f8) = id_f7;
3902 tree id_f9 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3903 get_identifier ("debug"),
3905 DECL_CHAIN (id_f9) = id_f8;
3906 tree id_f10 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3907 get_identifier ("omp_level"),
3909 DECL_CHAIN (id_f10) = id_f9;
3910 tree id_f11 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3911 get_identifier ("children_dispatches"),
3913 DECL_CHAIN (id_f11) = id_f10;
3914 tree id_f12 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
3915 get_identifier ("omp_num_threads"),
3917 DECL_CHAIN (id_f12) = id_f11;
3920 finish_builtin_struct (*hsa_kernel_dispatch_type, "__hsa_kernel_dispatch",
3922 TYPE_ARTIFICIAL (*hsa_kernel_dispatch_type) = 1;
3925 for (tree chain = TYPE_FIELDS (*hsa_kernel_dispatch_type);
3926 chain != NULL_TREE; chain = TREE_CHAIN (chain))
3927 if (id_equal (DECL_NAME (chain), field_name))
3928 return int_byte_position (chain);
3933 /* Return an HSA register that will contain number of threads for
3934 a future dispatched kernel. Instructions are added to HBB. */
3937 gen_num_threads_for_dispatch (hsa_bb *hbb)
3939 /* Step 1) Assign to number of threads:
3940 MIN (HSA_DEFAULT_NUM_THREADS, hsa_num_threads). */
3941 hsa_op_reg *threads = new hsa_op_reg (hsa_num_threads->m_type);
3942 hsa_op_address *addr = new hsa_op_address (hsa_num_threads);
3944 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD, threads->m_type,
3947 hsa_op_immed *limit = new hsa_op_immed (HSA_DEFAULT_NUM_THREADS,
3949 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_B1);
3951 = new hsa_insn_cmp (BRIG_COMPARE_LT, r->m_type, r, threads, limit);
3952 hbb->append_insn (cmp);
3954 BrigType16_t btype = hsa_bittype_for_type (threads->m_type);
3955 hsa_op_reg *tmp = new hsa_op_reg (threads->m_type);
3957 hbb->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp, r,
3960 /* Step 2) If the number is equal to zero,
3961 return shadow->omp_num_threads. */
3962 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg ();
3964 hsa_op_reg *shadow_thread_count = new hsa_op_reg (BRIG_TYPE_U32);
3966 = new hsa_op_address (shadow_reg_ptr,
3967 get_hsa_kernel_dispatch_offset ("omp_num_threads"));
3968 hsa_insn_basic *basic
3969 = new hsa_insn_mem (BRIG_OPCODE_LD, shadow_thread_count->m_type,
3970 shadow_thread_count, addr);
3971 hbb->append_insn (basic);
3973 hsa_op_reg *tmp2 = new hsa_op_reg (threads->m_type);
3974 r = new hsa_op_reg (BRIG_TYPE_B1);
3975 hsa_op_immed *imm = new hsa_op_immed (0, shadow_thread_count->m_type);
3976 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_EQ, r->m_type, r, tmp, imm));
3977 hbb->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp2, r,
3978 shadow_thread_count, tmp));
3980 hsa_op_base *dest = tmp2->get_in_type (BRIG_TYPE_U16, hbb);
3982 return as_a <hsa_op_reg *> (dest);
3985 /* Build OPCODE query for all three hsa dimensions, multiply them and store the
3986 result into DEST. */
3989 multiply_grid_dim_characteristics (hsa_op_reg *dest, int opcode, hsa_bb *hbb)
3991 hsa_op_reg *dimx = new hsa_op_reg (BRIG_TYPE_U32);
3992 query_hsa_grid_dim (dimx, opcode,
3993 new hsa_op_immed (0, (BrigKind16_t) BRIG_TYPE_U32), hbb);
3994 hsa_op_reg *dimy = new hsa_op_reg (BRIG_TYPE_U32);
3995 query_hsa_grid_dim (dimy, opcode,
3996 new hsa_op_immed (1, (BrigKind16_t) BRIG_TYPE_U32), hbb);
3997 hsa_op_reg *dimz = new hsa_op_reg (BRIG_TYPE_U32);
3998 query_hsa_grid_dim (dimz, opcode,
3999 new hsa_op_immed (2, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4000 hsa_op_reg *tmp = new hsa_op_reg (dest->m_type);
4001 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp,
4002 dimx->get_in_type (dest->m_type, hbb),
4003 dimy->get_in_type (dest->m_type, hbb), hbb);
4004 gen_hsa_binary_operation (BRIG_OPCODE_MUL, dest, tmp,
4005 dimz->get_in_type (dest->m_type, hbb), hbb);
4008 /* Emit instructions that assign number of threads to lhs of gimple STMT.
4009 Instructions are appended to basic block HBB. */
4012 gen_get_num_threads (gimple *stmt, hsa_bb *hbb)
4014 if (gimple_call_lhs (stmt) == NULL_TREE)
4017 hbb->append_insn (new hsa_insn_comment ("omp_get_num_threads"));
4018 tree lhs = gimple_call_lhs (stmt);
4019 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4020 multiply_grid_dim_characteristics (dest, BRIG_OPCODE_CURRENTWORKGROUPSIZE,
4024 /* Emit instructions that assign number of teams to lhs of gimple STMT.
4025 Instructions are appended to basic block HBB. */
4028 gen_get_num_teams (gimple *stmt, hsa_bb *hbb)
4030 if (gimple_call_lhs (stmt) == NULL_TREE)
4033 hbb->append_insn (new hsa_insn_comment ("omp_get_num_teams"));
4034 tree lhs = gimple_call_lhs (stmt);
4035 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4036 multiply_grid_dim_characteristics (dest, BRIG_OPCODE_GRIDGROUPS, hbb);
4039 /* Emit instructions that assign a team number to lhs of gimple STMT.
4040 Instructions are appended to basic block HBB. */
4043 gen_get_team_num (gimple *stmt, hsa_bb *hbb)
4045 if (gimple_call_lhs (stmt) == NULL_TREE)
4048 hbb->append_insn (new hsa_insn_comment ("omp_get_team_num"));
4049 tree lhs = gimple_call_lhs (stmt);
4050 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4052 hsa_op_reg *gnum_x = new hsa_op_reg (BRIG_TYPE_U32);
4053 query_hsa_grid_dim (gnum_x, BRIG_OPCODE_GRIDGROUPS,
4054 new hsa_op_immed (0, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4055 hsa_op_reg *gnum_y = new hsa_op_reg (BRIG_TYPE_U32);
4056 query_hsa_grid_dim (gnum_y, BRIG_OPCODE_GRIDGROUPS,
4057 new hsa_op_immed (1, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4059 hsa_op_reg *gno_z = new hsa_op_reg (BRIG_TYPE_U32);
4060 query_hsa_grid_dim (gno_z, BRIG_OPCODE_WORKGROUPID,
4061 new hsa_op_immed (2, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4063 hsa_op_reg *tmp1 = new hsa_op_reg (dest->m_type);
4064 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp1,
4065 gnum_x->get_in_type (dest->m_type, hbb),
4066 gnum_y->get_in_type (dest->m_type, hbb), hbb);
4067 hsa_op_reg *tmp2 = new hsa_op_reg (dest->m_type);
4068 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp2, tmp1,
4069 gno_z->get_in_type (dest->m_type, hbb), hbb);
4071 hsa_op_reg *gno_y = new hsa_op_reg (BRIG_TYPE_U32);
4072 query_hsa_grid_dim (gno_y, BRIG_OPCODE_WORKGROUPID,
4073 new hsa_op_immed (1, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4074 hsa_op_reg *tmp3 = new hsa_op_reg (dest->m_type);
4075 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp3,
4076 gnum_x->get_in_type (dest->m_type, hbb),
4077 gno_y->get_in_type (dest->m_type, hbb), hbb);
4078 hsa_op_reg *tmp4 = new hsa_op_reg (dest->m_type);
4079 gen_hsa_binary_operation (BRIG_OPCODE_ADD, tmp4, tmp3, tmp2, hbb);
4080 hsa_op_reg *gno_x = new hsa_op_reg (BRIG_TYPE_U32);
4081 query_hsa_grid_dim (gno_x, BRIG_OPCODE_WORKGROUPID,
4082 new hsa_op_immed (0, (BrigKind16_t) BRIG_TYPE_U32), hbb);
4083 gen_hsa_binary_operation (BRIG_OPCODE_ADD, dest, tmp4,
4084 gno_x->get_in_type (dest->m_type, hbb), hbb);
4087 /* Emit instructions that get levels-var ICV to lhs of gimple STMT.
4088 Instructions are appended to basic block HBB. */
4091 gen_get_level (gimple *stmt, hsa_bb *hbb)
4093 if (gimple_call_lhs (stmt) == NULL_TREE)
4096 hbb->append_insn (new hsa_insn_comment ("omp_get_level"));
4098 tree lhs = gimple_call_lhs (stmt);
4099 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4101 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg ();
4102 if (shadow_reg_ptr == NULL)
4104 HSA_SORRY_AT (gimple_location (stmt),
4105 "support for HSA does not implement omp_get_level called "
4106 "from a function not being inlined within a kernel");
4110 hsa_op_address *addr
4111 = new hsa_op_address (shadow_reg_ptr,
4112 get_hsa_kernel_dispatch_offset ("omp_level"));
4114 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64,
4115 (hsa_op_base *) NULL, addr);
4116 hbb->append_insn (mem);
4117 mem->set_output_in_type (dest, 0, hbb);
4120 /* Emit instruction that implement omp_get_max_threads of gimple STMT. */
4123 gen_get_max_threads (gimple *stmt, hsa_bb *hbb)
4125 tree lhs = gimple_call_lhs (stmt);
4129 hbb->append_insn (new hsa_insn_comment ("omp_get_max_threads"));
4131 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4132 hsa_op_with_type *num_theads_reg = gen_num_threads_for_dispatch (hbb)
4133 ->get_in_type (dest->m_type, hbb);
4134 hsa_build_append_simple_mov (dest, num_theads_reg, hbb);
4137 /* Emit instructions that implement alloca builtin gimple STMT.
4138 Instructions are appended to basic block HBB. */
4141 gen_hsa_alloca (gcall *call, hsa_bb *hbb)
4143 tree lhs = gimple_call_lhs (call);
4144 if (lhs == NULL_TREE)
4147 built_in_function fn = DECL_FUNCTION_CODE (gimple_call_fndecl (call));
4149 gcc_checking_assert (fn == BUILT_IN_ALLOCA
4150 || fn == BUILT_IN_ALLOCA_WITH_ALIGN);
4152 unsigned bit_alignment = 0;
4154 if (fn == BUILT_IN_ALLOCA_WITH_ALIGN)
4156 tree alignment_tree = gimple_call_arg (call, 1);
4157 if (TREE_CODE (alignment_tree) != INTEGER_CST)
4159 HSA_SORRY_ATV (gimple_location (call),
4160 "support for HSA does not implement "
4161 "__builtin_alloca_with_align with a non-constant "
4162 "alignment: %E", alignment_tree);
4165 bit_alignment = tree_to_uhwi (alignment_tree);
4168 tree rhs1 = gimple_call_arg (call, 0);
4169 hsa_op_with_type *size = hsa_reg_or_immed_for_gimple_op (rhs1, hbb)
4170 ->get_in_type (BRIG_TYPE_U32, hbb);
4171 hsa_op_with_type *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4174 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE));
4175 hsa_insn_alloca *a = new hsa_insn_alloca (tmp, size, bit_alignment);
4176 hbb->append_insn (a);
4179 = new hsa_insn_seg (BRIG_OPCODE_STOF,
4180 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT),
4181 tmp->m_type, BRIG_SEGMENT_PRIVATE, dest, tmp);
4182 hbb->append_insn (seg);
4185 /* Emit instructions that implement clrsb builtin STMT:
4186 Returns the number of leading redundant sign bits in x, i.e. the number
4187 of bits following the most significant bit that are identical to it.
4188 There are no special cases for 0 or other values.
4189 Instructions are appended to basic block HBB. */
4192 gen_hsa_clrsb (gcall *call, hsa_bb *hbb)
4194 tree lhs = gimple_call_lhs (call);
4195 if (lhs == NULL_TREE)
4198 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4199 tree rhs1 = gimple_call_arg (call, 0);
4200 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4201 BrigType16_t bittype = hsa_bittype_for_type (arg->m_type);
4202 unsigned bitsize = tree_to_uhwi (TYPE_SIZE (TREE_TYPE (rhs1)));
4204 /* FIRSTBIT instruction is defined just for 32 and 64-bits wide integers. */
4205 gcc_checking_assert (bitsize == 32 || bitsize == 64);
4207 /* Set true to MOST_SIG if the most significant bit is set to one. */
4208 hsa_op_immed *c = new hsa_op_immed (1ul << (bitsize - 1),
4209 hsa_uint_for_bitsize (bitsize));
4211 hsa_op_reg *and_reg = new hsa_op_reg (bittype);
4212 gen_hsa_binary_operation (BRIG_OPCODE_AND, and_reg, arg, c, hbb);
4214 hsa_op_reg *most_sign = new hsa_op_reg (BRIG_TYPE_B1);
4216 = new hsa_insn_cmp (BRIG_COMPARE_EQ, most_sign->m_type, most_sign,
4218 hbb->append_insn (cmp);
4220 /* If the most significant bit is one, negate the input. Otherwise
4221 shift the input value to left by one bit. */
4222 hsa_op_reg *arg_neg = new hsa_op_reg (arg->m_type);
4223 gen_hsa_unary_operation (BRIG_OPCODE_NEG, arg_neg, arg, hbb);
4225 hsa_op_reg *shifted_arg = new hsa_op_reg (arg->m_type);
4226 gen_hsa_binary_operation (BRIG_OPCODE_SHL, shifted_arg, arg,
4227 new hsa_op_immed (1, BRIG_TYPE_U64), hbb);
4229 /* Assign the value that can be used for FIRSTBIT instruction according
4230 to the most significant bit. */
4231 hsa_op_reg *tmp = new hsa_op_reg (bittype);
4232 hsa_insn_basic *cmov
4233 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, bittype, tmp, most_sign,
4234 arg_neg, shifted_arg);
4235 hbb->append_insn (cmov);
4237 hsa_op_reg *leading_bits = new hsa_op_reg (BRIG_TYPE_S32);
4238 gen_hsa_unary_operation (BRIG_OPCODE_FIRSTBIT, leading_bits,
4239 tmp->get_in_type (hsa_uint_for_bitsize (bitsize),
4242 /* Set flag if the input value is equal to zero. */
4243 hsa_op_reg *is_zero = new hsa_op_reg (BRIG_TYPE_B1);
4244 cmp = new hsa_insn_cmp (BRIG_COMPARE_EQ, is_zero->m_type, is_zero, arg,
4245 new hsa_op_immed (0, arg->m_type));
4246 hbb->append_insn (cmp);
4248 /* Return the number of leading bits,
4249 or (bitsize - 1) if the input value is zero. */
4250 cmov = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, BRIG_TYPE_B32, NULL, is_zero,
4251 new hsa_op_immed (bitsize - 1, BRIG_TYPE_U32),
4252 leading_bits->get_in_type (BRIG_TYPE_B32, hbb));
4253 hbb->append_insn (cmov);
4254 cmov->set_output_in_type (dest, 0, hbb);
4257 /* Emit instructions that implement ffs builtin STMT:
4258 Returns one plus the index of the least significant 1-bit of x,
4259 or if x is zero, returns zero.
4260 Instructions are appended to basic block HBB. */
4263 gen_hsa_ffs (gcall *call, hsa_bb *hbb)
4265 tree lhs = gimple_call_lhs (call);
4266 if (lhs == NULL_TREE)
4269 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4271 tree rhs1 = gimple_call_arg (call, 0);
4272 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4274 hsa_op_reg *tmp = new hsa_op_reg (BRIG_TYPE_U32);
4275 hsa_insn_srctype *insn = new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT,
4276 tmp->m_type, arg->m_type,
4278 hbb->append_insn (insn);
4280 hsa_insn_basic *addition
4281 = new hsa_insn_basic (3, BRIG_OPCODE_ADD, tmp->m_type, NULL, tmp,
4282 new hsa_op_immed (1, tmp->m_type));
4283 hbb->append_insn (addition);
4284 addition->set_output_in_type (dest, 0, hbb);
4288 gen_hsa_popcount_to_dest (hsa_op_reg *dest, hsa_op_with_type *arg, hsa_bb *hbb)
4290 gcc_checking_assert (hsa_type_integer_p (arg->m_type));
4292 if (hsa_type_bit_size (arg->m_type) < 32)
4293 arg = arg->get_in_type (BRIG_TYPE_B32, hbb);
4295 BrigType16_t srctype = hsa_bittype_for_type (arg->m_type);
4296 if (!hsa_btype_p (arg->m_type))
4297 arg = arg->get_in_type (srctype, hbb);
4299 hsa_insn_srctype *popcount
4300 = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT, BRIG_TYPE_U32,
4301 srctype, NULL, arg);
4302 hbb->append_insn (popcount);
4303 popcount->set_output_in_type (dest, 0, hbb);
4306 /* Emit instructions that implement parity builtin STMT:
4307 Returns the parity of x, i.e. the number of 1-bits in x modulo 2.
4308 Instructions are appended to basic block HBB. */
4311 gen_hsa_parity (gcall *call, hsa_bb *hbb)
4313 tree lhs = gimple_call_lhs (call);
4314 if (lhs == NULL_TREE)
4317 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4318 tree rhs1 = gimple_call_arg (call, 0);
4319 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4321 hsa_op_reg *popcount = new hsa_op_reg (BRIG_TYPE_U32);
4322 gen_hsa_popcount_to_dest (popcount, arg, hbb);
4324 hsa_insn_basic *insn
4325 = new hsa_insn_basic (3, BRIG_OPCODE_REM, popcount->m_type, NULL, popcount,
4326 new hsa_op_immed (2, popcount->m_type));
4327 hbb->append_insn (insn);
4328 insn->set_output_in_type (dest, 0, hbb);
4331 /* Emit instructions that implement popcount builtin STMT.
4332 Instructions are appended to basic block HBB. */
4335 gen_hsa_popcount (gcall *call, hsa_bb *hbb)
4337 tree lhs = gimple_call_lhs (call);
4338 if (lhs == NULL_TREE)
4341 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4342 tree rhs1 = gimple_call_arg (call, 0);
4343 hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4345 gen_hsa_popcount_to_dest (dest, arg, hbb);
4348 /* Emit instructions that implement DIVMOD builtin STMT.
4349 Instructions are appended to basic block HBB. */
4352 gen_hsa_divmod (gcall *call, hsa_bb *hbb)
4354 tree lhs = gimple_call_lhs (call);
4355 if (lhs == NULL_TREE)
4358 tree rhs0 = gimple_call_arg (call, 0);
4359 tree rhs1 = gimple_call_arg (call, 1);
4361 hsa_op_with_type *arg0 = hsa_reg_or_immed_for_gimple_op (rhs0, hbb);
4362 hsa_op_with_type *arg1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4364 hsa_op_reg *dest0 = new hsa_op_reg (arg0->m_type);
4365 hsa_op_reg *dest1 = new hsa_op_reg (arg1->m_type);
4367 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_DIV, dest0->m_type,
4369 hbb->append_insn (insn);
4370 insn = new hsa_insn_basic (3, BRIG_OPCODE_REM, dest1->m_type, dest1, arg0,
4372 hbb->append_insn (insn);
4374 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4375 BrigType16_t src_type = hsa_bittype_for_type (dest0->m_type);
4377 insn = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type,
4378 src_type, dest, dest0, dest1);
4379 hbb->append_insn (insn);
4382 /* Set VALUE to a shadow kernel debug argument and append a new instruction
4383 to HBB basic block. */
4386 set_debug_value (hsa_bb *hbb, hsa_op_with_type *value)
4388 hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg ();
4389 if (shadow_reg_ptr == NULL)
4392 hsa_op_address *addr
4393 = new hsa_op_address (shadow_reg_ptr,
4394 get_hsa_kernel_dispatch_offset ("debug"));
4395 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, BRIG_TYPE_U64, value,
4397 hbb->append_insn (mem);
4401 omp_simple_builtin::generate (gimple *stmt, hsa_bb *hbb)
4405 if (m_warning_message)
4406 HSA_SORRY_AT (gimple_location (stmt), m_warning_message);
4408 HSA_SORRY_ATV (gimple_location (stmt),
4409 "Support for HSA does not implement calls to %s\n",
4412 else if (m_warning_message != NULL)
4413 warning_at (gimple_location (stmt), OPT_Whsa, m_warning_message);
4415 if (m_return_value != NULL)
4417 tree lhs = gimple_call_lhs (stmt);
4421 hbb->append_insn (new hsa_insn_comment (m_name));
4423 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4424 hsa_op_with_type *op = m_return_value->get_in_type (dest->m_type, hbb);
4425 hsa_build_append_simple_mov (dest, op, hbb);
4429 /* If STMT is a call of a known library function, generate code to perform
4430 it and return true. */
4433 gen_hsa_insns_for_known_library_call (gimple *stmt, hsa_bb *hbb)
4435 bool handled = false;
4436 const char *name = hsa_get_declaration_name (gimple_call_fndecl (stmt));
4439 size_t len = strlen (name);
4440 if (len > 0 && name[len - 1] == '_')
4442 copy = XNEWVEC (char, len + 1);
4443 strcpy (copy, name);
4444 copy[len - 1] = '\0';
4448 /* Handle omp_* routines. */
4449 if (strstr (name, "omp_") == name)
4451 hsa_init_simple_builtins ();
4452 omp_simple_builtin *builtin = omp_simple_builtins->get (name);
4455 builtin->generate (stmt, hbb);
4460 if (strcmp (name, "omp_set_num_threads") == 0)
4461 gen_set_num_threads (gimple_call_arg (stmt, 0), hbb);
4462 else if (strcmp (name, "omp_get_thread_num") == 0)
4464 hbb->append_insn (new hsa_insn_comment (name));
4465 query_hsa_grid_nodim (stmt, BRIG_OPCODE_WORKITEMFLATABSID, hbb);
4467 else if (strcmp (name, "omp_get_num_threads") == 0)
4469 hbb->append_insn (new hsa_insn_comment (name));
4470 gen_get_num_threads (stmt, hbb);
4472 else if (strcmp (name, "omp_get_num_teams") == 0)
4473 gen_get_num_teams (stmt, hbb);
4474 else if (strcmp (name, "omp_get_team_num") == 0)
4475 gen_get_team_num (stmt, hbb);
4476 else if (strcmp (name, "omp_get_level") == 0)
4477 gen_get_level (stmt, hbb);
4478 else if (strcmp (name, "omp_get_active_level") == 0)
4479 gen_get_level (stmt, hbb);
4480 else if (strcmp (name, "omp_in_parallel") == 0)
4481 gen_get_level (stmt, hbb);
4482 else if (strcmp (name, "omp_get_max_threads") == 0)
4483 gen_get_max_threads (stmt, hbb);
4495 if (strcmp (name, "__hsa_set_debug_value") == 0)
4498 if (hsa_cfun->has_shadow_reg_p ())
4500 tree rhs1 = gimple_call_arg (stmt, 0);
4501 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
4503 src = src->get_in_type (BRIG_TYPE_U64, hbb);
4504 set_debug_value (hbb, src);
4513 /* Helper functions to create a single unary HSA operations out of calls to
4514 builtins. OPCODE is the HSA operation to be generated. STMT is a gimple
4515 call to a builtin. HBB is the HSA BB to which the instruction should be
4516 added. Note that nothing will be created if STMT does not have a LHS. */
4519 gen_hsa_unaryop_for_builtin (BrigOpcode opcode, gimple *stmt, hsa_bb *hbb)
4521 tree lhs = gimple_call_lhs (stmt);
4524 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4525 hsa_op_with_type *op
4526 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb);
4527 gen_hsa_unary_operation (opcode, dest, op, hbb);
4530 /* Helper functions to create a call to standard library if LHS of the
4531 STMT is used. HBB is the HSA BB to which the instruction should be
4535 gen_hsa_unaryop_builtin_call (gimple *stmt, hsa_bb *hbb)
4537 tree lhs = gimple_call_lhs (stmt);
4541 if (gimple_call_internal_p (stmt))
4542 gen_hsa_insns_for_call_of_internal_fn (stmt, hbb);
4544 gen_hsa_insns_for_direct_call (stmt, hbb);
4547 /* Helper functions to create a single unary HSA operations out of calls to
4548 builtins (if unsafe math optimizations are enable). Otherwise, create
4549 a call to standard library function.
4550 OPCODE is the HSA operation to be generated. STMT is a gimple
4551 call to a builtin. HBB is the HSA BB to which the instruction should be
4552 added. Note that nothing will be created if STMT does not have a LHS. */
4555 gen_hsa_unaryop_or_call_for_builtin (BrigOpcode opcode, gimple *stmt,
4558 if (flag_unsafe_math_optimizations)
4559 gen_hsa_unaryop_for_builtin (opcode, stmt, hbb);
4561 gen_hsa_unaryop_builtin_call (stmt, hbb);
4564 /* Generate HSA address corresponding to a value VAL (as opposed to a memory
4565 reference tree), for example an SSA_NAME or an ADDR_EXPR. HBB is the HSA BB
4566 to which the instruction should be added. */
4568 static hsa_op_address *
4569 get_address_from_value (tree val, hsa_bb *hbb)
4571 switch (TREE_CODE (val))
4575 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
4577 = hsa_cfun->reg_for_gimple_ssa (val)->get_in_type (addrtype, hbb);
4578 return new hsa_op_address (NULL, as_a <hsa_op_reg *> (reg), 0);
4581 return gen_hsa_addr (TREE_OPERAND (val, 0), hbb);
4584 if (tree_fits_shwi_p (val))
4585 return new hsa_op_address (NULL, NULL, tree_to_shwi (val));
4589 HSA_SORRY_ATV (EXPR_LOCATION (val),
4590 "support for HSA does not implement memory access to %E",
4592 return new hsa_op_address (NULL, NULL, 0);
4596 /* Expand assignment of a result of a string BUILTIN to DST.
4597 Size of the operation is N bytes, where instructions
4598 will be append to HBB. */
4601 expand_lhs_of_string_op (gimple *stmt,
4602 unsigned HOST_WIDE_INT n, hsa_bb *hbb,
4603 enum built_in_function builtin)
4605 /* If LHS is expected, we need to emit a PHI instruction. */
4606 tree lhs = gimple_call_lhs (stmt);
4610 hsa_op_reg *lhs_reg = hsa_cfun->reg_for_gimple_ssa (lhs);
4612 hsa_op_with_type *dst_reg
4613 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb);
4614 hsa_op_with_type *tmp;
4618 case BUILT_IN_MEMPCPY:
4620 tmp = new hsa_op_reg (dst_reg->m_type);
4622 = new hsa_insn_basic (3, BRIG_OPCODE_ADD, tmp->m_type,
4624 new hsa_op_immed (n, dst_reg->m_type));
4625 hbb->append_insn (add);
4628 case BUILT_IN_MEMCPY:
4629 case BUILT_IN_MEMSET:
4636 hbb->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV, lhs_reg->m_type,
4640 #define HSA_MEMORY_BUILTINS_LIMIT 128
4642 /* Expand a string builtin (from a gimple STMT) in a way that
4643 according to MISALIGNED_FLAG we process either direct emission
4644 (a bunch of memory load and store instructions), or we emit a function call
4645 of a library function (for instance 'memcpy'). Actually, a basic block
4646 for direct emission is just prepared, where caller is responsible
4647 for emission of corresponding instructions.
4648 All instruction are appended to HBB. */
4651 expand_string_operation_builtin (gimple *stmt, hsa_bb *hbb,
4652 hsa_op_reg *misaligned_flag)
4654 edge e = split_block (hbb->m_bb, stmt);
4655 basic_block condition_bb = e->src;
4656 hbb->append_insn (new hsa_insn_cbr (misaligned_flag));
4658 /* Prepare the control flow. */
4659 edge condition_edge = EDGE_SUCC (condition_bb, 0);
4660 basic_block call_bb = split_edge (condition_edge);
4662 basic_block expanded_bb = split_edge (EDGE_SUCC (call_bb, 0));
4663 basic_block cont_bb = EDGE_SUCC (expanded_bb, 0)->dest;
4664 basic_block merge_bb = split_edge (EDGE_PRED (cont_bb, 0));
4666 condition_edge->flags &= ~EDGE_FALLTHRU;
4667 condition_edge->flags |= EDGE_TRUE_VALUE;
4668 make_edge (condition_bb, expanded_bb, EDGE_FALSE_VALUE);
4670 redirect_edge_succ (EDGE_SUCC (call_bb, 0), merge_bb);
4672 hsa_cfun->m_modified_cfg = true;
4674 hsa_init_new_bb (expanded_bb);
4676 /* Slow path: function call. */
4677 gen_hsa_insns_for_direct_call (stmt, hsa_init_new_bb (call_bb), false);
4679 return hsa_bb_for_bb (expanded_bb);
4682 /* Expand a memory copy BUILTIN (BUILT_IN_MEMCPY, BUILT_IN_MEMPCPY) from
4683 a gimple STMT and store all necessary instruction to HBB basic block. */
4686 expand_memory_copy (gimple *stmt, hsa_bb *hbb, enum built_in_function builtin)
4688 tree byte_size = gimple_call_arg (stmt, 2);
4690 if (!tree_fits_uhwi_p (byte_size))
4692 gen_hsa_insns_for_direct_call (stmt, hbb);
4696 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size);
4698 if (n > HSA_MEMORY_BUILTINS_LIMIT)
4700 gen_hsa_insns_for_direct_call (stmt, hbb);
4704 tree dst = gimple_call_arg (stmt, 0);
4705 tree src = gimple_call_arg (stmt, 1);
4707 hsa_op_address *dst_addr = get_address_from_value (dst, hbb);
4708 hsa_op_address *src_addr = get_address_from_value (src, hbb);
4710 /* As gen_hsa_memory_copy relies on memory alignment
4711 greater or equal to 8 bytes, we need to verify the alignment. */
4712 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
4713 hsa_op_reg *src_addr_reg = new hsa_op_reg (addrtype);
4714 hsa_op_reg *dst_addr_reg = new hsa_op_reg (addrtype);
4716 convert_addr_to_flat_segment (src_addr, src_addr_reg, hbb);
4717 convert_addr_to_flat_segment (dst_addr, dst_addr_reg, hbb);
4719 /* Process BIT OR for source and destination addresses. */
4720 hsa_op_reg *or_reg = new hsa_op_reg (addrtype);
4721 gen_hsa_binary_operation (BRIG_OPCODE_OR, or_reg, src_addr_reg,
4724 /* Process BIT AND with 0x7 to identify the desired alignment
4726 hsa_op_reg *masked = new hsa_op_reg (addrtype);
4728 gen_hsa_binary_operation (BRIG_OPCODE_AND, masked, or_reg,
4729 new hsa_op_immed (7, addrtype), hbb);
4731 hsa_op_reg *misaligned = new hsa_op_reg (BRIG_TYPE_B1);
4732 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE, misaligned->m_type,
4734 new hsa_op_immed (0, masked->m_type)));
4736 hsa_bb *native_impl_bb
4737 = expand_string_operation_builtin (stmt, hbb, misaligned);
4739 gen_hsa_memory_copy (native_impl_bb, dst_addr, src_addr, n, BRIG_ALIGNMENT_8);
4741 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb->m_bb, 0)->dest);
4742 expand_lhs_of_string_op (stmt, n, merge_bb, builtin);
4746 /* Expand a memory set BUILTIN (BUILT_IN_MEMSET, BUILT_IN_BZERO) from
4747 a gimple STMT and store all necessary instruction to HBB basic block.
4748 The operation set N bytes with a CONSTANT value. */
4751 expand_memory_set (gimple *stmt, unsigned HOST_WIDE_INT n,
4752 unsigned HOST_WIDE_INT constant, hsa_bb *hbb,
4753 enum built_in_function builtin)
4755 tree dst = gimple_call_arg (stmt, 0);
4756 hsa_op_address *dst_addr = get_address_from_value (dst, hbb);
4758 /* As gen_hsa_memory_set relies on memory alignment
4759 greater or equal to 8 bytes, we need to verify the alignment. */
4760 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
4761 hsa_op_reg *dst_addr_reg = new hsa_op_reg (addrtype);
4762 convert_addr_to_flat_segment (dst_addr, dst_addr_reg, hbb);
4764 /* Process BIT AND with 0x7 to identify the desired alignment
4766 hsa_op_reg *masked = new hsa_op_reg (addrtype);
4768 gen_hsa_binary_operation (BRIG_OPCODE_AND, masked, dst_addr_reg,
4769 new hsa_op_immed (7, addrtype), hbb);
4771 hsa_op_reg *misaligned = new hsa_op_reg (BRIG_TYPE_B1);
4772 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE, misaligned->m_type,
4774 new hsa_op_immed (0, masked->m_type)));
4776 hsa_bb *native_impl_bb
4777 = expand_string_operation_builtin (stmt, hbb, misaligned);
4779 gen_hsa_memory_set (native_impl_bb, dst_addr, constant, n, BRIG_ALIGNMENT_8);
4781 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb->m_bb, 0)->dest);
4782 expand_lhs_of_string_op (stmt, n, merge_bb, builtin);
4785 /* Store into MEMORDER the memory order specified by tree T, which must be an
4786 integer constant representing a C++ memory order. If it isn't, issue an HSA
4787 sorry message using LOC and return true, otherwise return false and store
4788 the name of the requested order to *MNAME. */
4791 hsa_memorder_from_tree (tree t, BrigMemoryOrder *memorder, const char **mname,
4794 if (!tree_fits_uhwi_p (t))
4796 HSA_SORRY_ATV (loc, "support for HSA does not implement memory model %E",
4801 unsigned HOST_WIDE_INT mm = tree_to_uhwi (t);
4802 switch (mm & MEMMODEL_BASE_MASK)
4804 case MEMMODEL_RELAXED:
4805 *memorder = BRIG_MEMORY_ORDER_RELAXED;
4808 case MEMMODEL_CONSUME:
4809 /* HSA does not have an equivalent, but we can use the slightly stronger
4811 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE;
4814 case MEMMODEL_ACQUIRE:
4815 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE;
4818 case MEMMODEL_RELEASE:
4819 *memorder = BRIG_MEMORY_ORDER_SC_RELEASE;
4822 case MEMMODEL_ACQ_REL:
4823 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE;
4826 case MEMMODEL_SEQ_CST:
4827 /* Callers implementing a simple load or store need to remove the release
4828 or acquire part respectively. */
4829 *memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE;
4834 HSA_SORRY_AT (loc, "support for HSA does not implement the specified "
4842 /* Helper function to create an HSA atomic operation instruction out of calls
4843 to atomic builtins. RET_ORIG is true if the built-in is the variant that
4844 return s the value before applying operation, and false if it should return
4845 the value after applying the operation (if it returns value at all). ACODE
4846 is the atomic operation code, STMT is a gimple call to a builtin. HBB is
4847 the HSA BB to which the instruction should be added. If SIGNAL is true, the
4848 created operation will work on HSA signals rather than atomic variables. */
4851 gen_hsa_atomic_for_builtin (bool ret_orig, enum BrigAtomicOperation acode,
4852 gimple *stmt, hsa_bb *hbb, bool signal)
4854 tree lhs = gimple_call_lhs (stmt);
4856 tree type = TREE_TYPE (gimple_call_arg (stmt, 1));
4857 BrigType16_t hsa_type = hsa_type_for_scalar_tree_type (type, false);
4858 BrigType16_t mtype = mem_type_for_type (hsa_type);
4859 BrigMemoryOrder memorder;
4862 if (hsa_memorder_from_tree (gimple_call_arg (stmt, 2), &memorder, &mmname,
4863 gimple_location (stmt)))
4866 /* Certain atomic insns must have Bx memory types. */
4869 case BRIG_ATOMIC_LD:
4870 case BRIG_ATOMIC_ST:
4871 case BRIG_ATOMIC_AND:
4872 case BRIG_ATOMIC_OR:
4873 case BRIG_ATOMIC_XOR:
4874 case BRIG_ATOMIC_EXCH:
4875 mtype = hsa_bittype_for_type (mtype);
4886 dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4888 dest = new hsa_op_reg (hsa_type);
4889 opcode = signal ? BRIG_OPCODE_SIGNAL : BRIG_OPCODE_ATOMIC;
4895 opcode = signal ? BRIG_OPCODE_SIGNALNORET : BRIG_OPCODE_ATOMICNORET;
4899 if (acode == BRIG_ATOMIC_ST)
4901 if (memorder == BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE)
4902 memorder = BRIG_MEMORY_ORDER_SC_RELEASE;
4904 if (memorder != BRIG_MEMORY_ORDER_RELAXED
4905 && memorder != BRIG_MEMORY_ORDER_SC_RELEASE
4906 && memorder != BRIG_MEMORY_ORDER_NONE)
4908 HSA_SORRY_ATV (gimple_location (stmt),
4909 "support for HSA does not implement memory model for "
4910 "ATOMIC_ST: %s", mmname);
4915 hsa_insn_basic *atominsn;
4919 atominsn = new hsa_insn_signal (nops, opcode, acode, mtype, memorder);
4920 tgt = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb);
4924 atominsn = new hsa_insn_atomic (nops, opcode, acode, mtype, memorder);
4925 hsa_op_address *addr;
4926 addr = get_address_from_value (gimple_call_arg (stmt, 0), hbb);
4927 if (addr->m_symbol && addr->m_symbol->m_segment == BRIG_SEGMENT_PRIVATE)
4929 HSA_SORRY_AT (gimple_location (stmt),
4930 "HSA does not implement atomic operations in private "
4937 hsa_op_base *op = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1),
4941 atominsn->set_op (0, dest);
4942 atominsn->set_op (1, tgt);
4943 atominsn->set_op (2, op);
4947 atominsn->set_op (0, tgt);
4948 atominsn->set_op (1, op);
4951 hbb->append_insn (atominsn);
4953 /* HSA does not natively support the variants that return the modified value,
4954 so re-do the operation again non-atomically if that is what was
4956 if (lhs && !ret_orig)
4961 case BRIG_ATOMIC_ADD:
4962 arith = BRIG_OPCODE_ADD;
4964 case BRIG_ATOMIC_AND:
4965 arith = BRIG_OPCODE_AND;
4967 case BRIG_ATOMIC_OR:
4968 arith = BRIG_OPCODE_OR;
4970 case BRIG_ATOMIC_SUB:
4971 arith = BRIG_OPCODE_SUB;
4973 case BRIG_ATOMIC_XOR:
4974 arith = BRIG_OPCODE_XOR;
4979 hsa_op_reg *real_dest = hsa_cfun->reg_for_gimple_ssa (lhs);
4980 gen_hsa_binary_operation (arith, real_dest, dest, op, hbb);
4984 /* Generate HSA instructions for an internal fn.
4985 Instructions will be appended to HBB, which also needs to be the
4986 corresponding structure to the basic_block of STMT. */
4989 gen_hsa_insn_for_internal_fn_call (gcall *stmt, hsa_bb *hbb)
4991 gcc_checking_assert (gimple_call_internal_fn (stmt));
4992 internal_fn fn = gimple_call_internal_fn (stmt);
4994 bool is_float_type_p = false;
4995 if (gimple_call_lhs (stmt) != NULL
4996 && TREE_TYPE (gimple_call_lhs (stmt)) == float_type_node)
4997 is_float_type_p = true;
5002 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL, stmt, hbb);
5006 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR, stmt, hbb);
5010 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT, stmt, hbb);
5014 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT, stmt, hbb);
5018 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_NRSQRT, stmt, hbb);
5022 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC, stmt, hbb);
5027 if (is_float_type_p)
5028 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS, stmt, hbb);
5030 gen_hsa_unaryop_builtin_call (stmt, hbb);
5036 if (is_float_type_p)
5037 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2, stmt, hbb);
5039 gen_hsa_unaryop_builtin_call (stmt, hbb);
5046 if (is_float_type_p)
5047 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2, stmt, hbb);
5049 gen_hsa_unaryop_builtin_call (stmt, hbb);
5056 if (is_float_type_p)
5057 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN, stmt, hbb);
5059 gen_hsa_unaryop_builtin_call (stmt, hbb);
5064 gen_hsa_clrsb (stmt, hbb);
5068 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT, stmt, hbb);
5072 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT, stmt, hbb);
5076 gen_hsa_ffs (stmt, hbb);
5080 gen_hsa_parity (stmt, hbb);
5084 gen_hsa_popcount (stmt, hbb);
5088 gen_hsa_divmod (stmt, hbb);
5101 case IFN_SIGNIFICAND:
5113 gen_hsa_insns_for_call_of_internal_fn (stmt, hbb);
5117 HSA_SORRY_ATV (gimple_location (stmt),
5118 "support for HSA does not implement internal function: %s",
5119 internal_fn_name (fn));
5124 /* Generate HSA instructions for the given call statement STMT. Instructions
5125 will be appended to HBB. */
5128 gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb)
5130 gcall *call = as_a <gcall *> (stmt);
5131 tree lhs = gimple_call_lhs (stmt);
5134 if (gimple_call_internal_p (stmt))
5136 gen_hsa_insn_for_internal_fn_call (call, hbb);
5140 if (!gimple_call_builtin_p (stmt, BUILT_IN_NORMAL))
5142 tree function_decl = gimple_call_fndecl (stmt);
5143 /* Prefetch pass can create type-mismatching prefetch builtin calls which
5144 fail the gimple_call_builtin_p test above. Handle them here. */
5145 if (DECL_BUILT_IN_CLASS (function_decl)
5146 && DECL_FUNCTION_CODE (function_decl) == BUILT_IN_PREFETCH)
5149 if (function_decl == NULL_TREE)
5151 HSA_SORRY_AT (gimple_location (stmt),
5152 "support for HSA does not implement indirect calls");
5156 if (hsa_callable_function_p (function_decl))
5157 gen_hsa_insns_for_direct_call (stmt, hbb);
5158 else if (!gen_hsa_insns_for_known_library_call (stmt, hbb))
5159 HSA_SORRY_AT (gimple_location (stmt),
5160 "HSA supports only calls of functions marked with pragma "
5161 "omp declare target");
5165 tree fndecl = gimple_call_fndecl (stmt);
5166 enum built_in_function builtin = DECL_FUNCTION_CODE (fndecl);
5170 case BUILT_IN_FABSF:
5171 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_ABS, stmt, hbb);
5175 case BUILT_IN_CEILF:
5176 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL, stmt, hbb);
5179 case BUILT_IN_FLOOR:
5180 case BUILT_IN_FLOORF:
5181 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR, stmt, hbb);
5185 case BUILT_IN_RINTF:
5186 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT, stmt, hbb);
5190 case BUILT_IN_SQRTF:
5191 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT, stmt, hbb);
5194 case BUILT_IN_TRUNC:
5195 case BUILT_IN_TRUNCF:
5196 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC, stmt, hbb);
5203 /* HSAIL does not provide an instruction for double argument type. */
5204 gen_hsa_unaryop_builtin_call (stmt, hbb);
5208 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS, stmt, hbb);
5211 case BUILT_IN_EXP2F:
5212 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2, stmt, hbb);
5215 case BUILT_IN_LOG2F:
5216 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2, stmt, hbb);
5220 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN, stmt, hbb);
5223 case BUILT_IN_CLRSB:
5224 case BUILT_IN_CLRSBL:
5225 case BUILT_IN_CLRSBLL:
5226 gen_hsa_clrsb (call, hbb);
5231 case BUILT_IN_CLZLL:
5232 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT, stmt, hbb);
5237 case BUILT_IN_CTZLL:
5238 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT, stmt, hbb);
5243 case BUILT_IN_FFSLL:
5244 gen_hsa_ffs (call, hbb);
5247 case BUILT_IN_PARITY:
5248 case BUILT_IN_PARITYL:
5249 case BUILT_IN_PARITYLL:
5250 gen_hsa_parity (call, hbb);
5253 case BUILT_IN_POPCOUNT:
5254 case BUILT_IN_POPCOUNTL:
5255 case BUILT_IN_POPCOUNTLL:
5256 gen_hsa_popcount (call, hbb);
5259 case BUILT_IN_ATOMIC_LOAD_1:
5260 case BUILT_IN_ATOMIC_LOAD_2:
5261 case BUILT_IN_ATOMIC_LOAD_4:
5262 case BUILT_IN_ATOMIC_LOAD_8:
5263 case BUILT_IN_ATOMIC_LOAD_16:
5267 src = get_address_from_value (gimple_call_arg (stmt, 0), hbb);
5269 BrigMemoryOrder memorder;
5271 if (hsa_memorder_from_tree (gimple_call_arg (stmt, 1), &memorder,
5272 &mmname, gimple_location (stmt)))
5275 if (memorder == BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE)
5276 memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE;
5278 if (memorder != BRIG_MEMORY_ORDER_RELAXED
5279 && memorder != BRIG_MEMORY_ORDER_SC_ACQUIRE
5280 && memorder != BRIG_MEMORY_ORDER_NONE)
5282 HSA_SORRY_ATV (gimple_location (stmt),
5283 "support for HSA does not implement "
5284 "memory model for atomic loads: %s", mmname);
5290 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
5292 mtype = mem_type_for_type (t);
5293 mtype = hsa_bittype_for_type (mtype);
5294 dest = hsa_cfun->reg_for_gimple_ssa (lhs);
5298 mtype = BRIG_TYPE_B64;
5299 dest = new hsa_op_reg (mtype);
5302 hsa_insn_basic *atominsn;
5303 atominsn = new hsa_insn_atomic (2, BRIG_OPCODE_ATOMIC, BRIG_ATOMIC_LD,
5304 mtype, memorder, dest, src);
5306 hbb->append_insn (atominsn);
5310 case BUILT_IN_ATOMIC_EXCHANGE_1:
5311 case BUILT_IN_ATOMIC_EXCHANGE_2:
5312 case BUILT_IN_ATOMIC_EXCHANGE_4:
5313 case BUILT_IN_ATOMIC_EXCHANGE_8:
5314 case BUILT_IN_ATOMIC_EXCHANGE_16:
5315 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_EXCH, stmt, hbb, false);
5319 case BUILT_IN_ATOMIC_FETCH_ADD_1:
5320 case BUILT_IN_ATOMIC_FETCH_ADD_2:
5321 case BUILT_IN_ATOMIC_FETCH_ADD_4:
5322 case BUILT_IN_ATOMIC_FETCH_ADD_8:
5323 case BUILT_IN_ATOMIC_FETCH_ADD_16:
5324 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ADD, stmt, hbb, false);
5328 case BUILT_IN_ATOMIC_FETCH_SUB_1:
5329 case BUILT_IN_ATOMIC_FETCH_SUB_2:
5330 case BUILT_IN_ATOMIC_FETCH_SUB_4:
5331 case BUILT_IN_ATOMIC_FETCH_SUB_8:
5332 case BUILT_IN_ATOMIC_FETCH_SUB_16:
5333 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_SUB, stmt, hbb, false);
5337 case BUILT_IN_ATOMIC_FETCH_AND_1:
5338 case BUILT_IN_ATOMIC_FETCH_AND_2:
5339 case BUILT_IN_ATOMIC_FETCH_AND_4:
5340 case BUILT_IN_ATOMIC_FETCH_AND_8:
5341 case BUILT_IN_ATOMIC_FETCH_AND_16:
5342 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_AND, stmt, hbb, false);
5346 case BUILT_IN_ATOMIC_FETCH_XOR_1:
5347 case BUILT_IN_ATOMIC_FETCH_XOR_2:
5348 case BUILT_IN_ATOMIC_FETCH_XOR_4:
5349 case BUILT_IN_ATOMIC_FETCH_XOR_8:
5350 case BUILT_IN_ATOMIC_FETCH_XOR_16:
5351 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_XOR, stmt, hbb, false);
5355 case BUILT_IN_ATOMIC_FETCH_OR_1:
5356 case BUILT_IN_ATOMIC_FETCH_OR_2:
5357 case BUILT_IN_ATOMIC_FETCH_OR_4:
5358 case BUILT_IN_ATOMIC_FETCH_OR_8:
5359 case BUILT_IN_ATOMIC_FETCH_OR_16:
5360 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_OR, stmt, hbb, false);
5364 case BUILT_IN_ATOMIC_STORE_1:
5365 case BUILT_IN_ATOMIC_STORE_2:
5366 case BUILT_IN_ATOMIC_STORE_4:
5367 case BUILT_IN_ATOMIC_STORE_8:
5368 case BUILT_IN_ATOMIC_STORE_16:
5369 /* Since there cannot be any LHS, the first parameter is meaningless. */
5370 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ST, stmt, hbb, false);
5374 case BUILT_IN_ATOMIC_ADD_FETCH_1:
5375 case BUILT_IN_ATOMIC_ADD_FETCH_2:
5376 case BUILT_IN_ATOMIC_ADD_FETCH_4:
5377 case BUILT_IN_ATOMIC_ADD_FETCH_8:
5378 case BUILT_IN_ATOMIC_ADD_FETCH_16:
5379 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_ADD, stmt, hbb, false);
5382 case BUILT_IN_ATOMIC_SUB_FETCH_1:
5383 case BUILT_IN_ATOMIC_SUB_FETCH_2:
5384 case BUILT_IN_ATOMIC_SUB_FETCH_4:
5385 case BUILT_IN_ATOMIC_SUB_FETCH_8:
5386 case BUILT_IN_ATOMIC_SUB_FETCH_16:
5387 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_SUB, stmt, hbb, false);
5390 case BUILT_IN_ATOMIC_AND_FETCH_1:
5391 case BUILT_IN_ATOMIC_AND_FETCH_2:
5392 case BUILT_IN_ATOMIC_AND_FETCH_4:
5393 case BUILT_IN_ATOMIC_AND_FETCH_8:
5394 case BUILT_IN_ATOMIC_AND_FETCH_16:
5395 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_AND, stmt, hbb, false);
5398 case BUILT_IN_ATOMIC_XOR_FETCH_1:
5399 case BUILT_IN_ATOMIC_XOR_FETCH_2:
5400 case BUILT_IN_ATOMIC_XOR_FETCH_4:
5401 case BUILT_IN_ATOMIC_XOR_FETCH_8:
5402 case BUILT_IN_ATOMIC_XOR_FETCH_16:
5403 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_XOR, stmt, hbb, false);
5406 case BUILT_IN_ATOMIC_OR_FETCH_1:
5407 case BUILT_IN_ATOMIC_OR_FETCH_2:
5408 case BUILT_IN_ATOMIC_OR_FETCH_4:
5409 case BUILT_IN_ATOMIC_OR_FETCH_8:
5410 case BUILT_IN_ATOMIC_OR_FETCH_16:
5411 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_OR, stmt, hbb, false);
5414 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_1:
5415 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_2:
5416 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_4:
5417 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8:
5418 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_16:
5420 tree type = TREE_TYPE (gimple_call_arg (stmt, 1));
5422 = hsa_bittype_for_type (hsa_type_for_scalar_tree_type (type, false));
5423 BrigMemoryOrder memorder = BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE;
5424 hsa_insn_basic *atominsn;
5426 atominsn = new hsa_insn_atomic (4, BRIG_OPCODE_ATOMIC,
5427 BRIG_ATOMIC_CAS, atype, memorder);
5428 tgt = get_address_from_value (gimple_call_arg (stmt, 0), hbb);
5431 dest = hsa_cfun->reg_for_gimple_ssa (lhs);
5433 dest = new hsa_op_reg (atype);
5435 atominsn->set_op (0, dest);
5436 atominsn->set_op (1, tgt);
5438 hsa_op_with_type *op
5439 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1), hbb);
5440 atominsn->set_op (2, op);
5441 op = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 2), hbb);
5442 atominsn->set_op (3, op);
5444 hbb->append_insn (atominsn);
5448 case BUILT_IN_HSA_WORKGROUPID:
5449 query_hsa_grid_dim (stmt, BRIG_OPCODE_WORKGROUPID, hbb);
5451 case BUILT_IN_HSA_WORKITEMID:
5452 query_hsa_grid_dim (stmt, BRIG_OPCODE_WORKITEMID, hbb);
5454 case BUILT_IN_HSA_WORKITEMABSID:
5455 query_hsa_grid_dim (stmt, BRIG_OPCODE_WORKITEMABSID, hbb);
5457 case BUILT_IN_HSA_GRIDSIZE:
5458 query_hsa_grid_dim (stmt, BRIG_OPCODE_GRIDSIZE, hbb);
5460 case BUILT_IN_HSA_CURRENTWORKGROUPSIZE:
5461 query_hsa_grid_dim (stmt, BRIG_OPCODE_CURRENTWORKGROUPSIZE, hbb);
5464 case BUILT_IN_GOMP_BARRIER:
5465 hbb->append_insn (new hsa_insn_br (0, BRIG_OPCODE_BARRIER, BRIG_TYPE_NONE,
5468 case BUILT_IN_GOMP_PARALLEL:
5469 HSA_SORRY_AT (gimple_location (stmt),
5470 "support for HSA does not implement non-gridified "
5471 "OpenMP parallel constructs.");
5474 case BUILT_IN_OMP_GET_THREAD_NUM:
5476 query_hsa_grid_nodim (stmt, BRIG_OPCODE_WORKITEMFLATABSID, hbb);
5480 case BUILT_IN_OMP_GET_NUM_THREADS:
5482 gen_get_num_threads (stmt, hbb);
5485 case BUILT_IN_GOMP_TEAMS:
5487 gen_set_num_threads (gimple_call_arg (stmt, 1), hbb);
5490 case BUILT_IN_OMP_GET_NUM_TEAMS:
5492 gen_get_num_teams (stmt, hbb);
5495 case BUILT_IN_OMP_GET_TEAM_NUM:
5497 gen_get_team_num (stmt, hbb);
5500 case BUILT_IN_MEMCPY:
5501 case BUILT_IN_MEMPCPY:
5503 expand_memory_copy (stmt, hbb, builtin);
5506 case BUILT_IN_MEMSET:
5508 tree c = gimple_call_arg (stmt, 1);
5510 if (TREE_CODE (c) != INTEGER_CST)
5512 gen_hsa_insns_for_direct_call (stmt, hbb);
5516 tree byte_size = gimple_call_arg (stmt, 2);
5518 if (!tree_fits_uhwi_p (byte_size))
5520 gen_hsa_insns_for_direct_call (stmt, hbb);
5524 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size);
5526 if (n > HSA_MEMORY_BUILTINS_LIMIT)
5528 gen_hsa_insns_for_direct_call (stmt, hbb);
5532 unsigned HOST_WIDE_INT constant
5533 = tree_to_uhwi (fold_convert (unsigned_char_type_node, c));
5535 expand_memory_set (stmt, n, constant, hbb, builtin);
5539 case BUILT_IN_BZERO:
5541 tree byte_size = gimple_call_arg (stmt, 1);
5543 if (!tree_fits_uhwi_p (byte_size))
5545 gen_hsa_insns_for_direct_call (stmt, hbb);
5549 unsigned HOST_WIDE_INT n = tree_to_uhwi (byte_size);
5551 if (n > HSA_MEMORY_BUILTINS_LIMIT)
5553 gen_hsa_insns_for_direct_call (stmt, hbb);
5557 expand_memory_set (stmt, n, 0, hbb, builtin);
5561 case BUILT_IN_ALLOCA:
5562 case BUILT_IN_ALLOCA_WITH_ALIGN:
5564 gen_hsa_alloca (call, hbb);
5567 case BUILT_IN_PREFETCH:
5571 tree name_tree = DECL_NAME (fndecl);
5572 const char *s = IDENTIFIER_POINTER (name_tree);
5573 size_t len = strlen (s);
5574 if (len > 4 && (strncmp (s, "__builtin_GOMP_", 15) == 0))
5575 HSA_SORRY_ATV (gimple_location (stmt),
5576 "support for HSA does not implement GOMP function %s",
5579 gen_hsa_insns_for_direct_call (stmt, hbb);
5585 /* Generate HSA instructions for a given gimple statement. Instructions will be
5589 gen_hsa_insns_for_gimple_stmt (gimple *stmt, hsa_bb *hbb)
5591 switch (gimple_code (stmt))
5594 if (gimple_clobber_p (stmt))
5597 if (gimple_assign_single_p (stmt))
5599 tree lhs = gimple_assign_lhs (stmt);
5600 tree rhs = gimple_assign_rhs1 (stmt);
5601 gen_hsa_insns_for_single_assignment (lhs, rhs, hbb);
5604 gen_hsa_insns_for_operation_assignment (stmt, hbb);
5607 gen_hsa_insns_for_return (as_a <greturn *> (stmt), hbb);
5610 gen_hsa_insns_for_cond_stmt (stmt, hbb);
5613 gen_hsa_insns_for_call (stmt, hbb);
5616 /* ??? HSA supports some debug facilities. */
5620 tree label = gimple_label_label (as_a <glabel *> (stmt));
5621 if (FORCED_LABEL (label))
5622 HSA_SORRY_AT (gimple_location (stmt),
5623 "support for HSA does not implement gimple label with "
5630 hbb->append_insn (new hsa_insn_basic (0, BRIG_OPCODE_NOP));
5635 gen_hsa_insns_for_switch_stmt (as_a <gswitch *> (stmt), hbb);
5639 HSA_SORRY_ATV (gimple_location (stmt),
5640 "support for HSA does not implement gimple statement %s",
5641 gimple_code_name[(int) gimple_code (stmt)]);
5645 /* Generate a HSA PHI from a gimple PHI. */
5648 gen_hsa_phi_from_gimple_phi (gimple *phi_stmt, hsa_bb *hbb)
5651 unsigned count = gimple_phi_num_args (phi_stmt);
5654 = hsa_cfun->reg_for_gimple_ssa (gimple_phi_result (phi_stmt));
5655 hphi = new hsa_insn_phi (count, dest);
5656 hphi->m_bb = hbb->m_bb;
5658 tree lhs = gimple_phi_result (phi_stmt);
5660 for (unsigned i = 0; i < count; i++)
5662 tree op = gimple_phi_arg_def (phi_stmt, i);
5664 if (TREE_CODE (op) == SSA_NAME)
5666 hsa_op_reg *hreg = hsa_cfun->reg_for_gimple_ssa (op);
5667 hphi->set_op (i, hreg);
5671 gcc_assert (is_gimple_min_invariant (op));
5672 tree t = TREE_TYPE (op);
5673 if (!POINTER_TYPE_P (t)
5674 || (TREE_CODE (op) == STRING_CST
5675 && TREE_CODE (TREE_TYPE (t)) == INTEGER_TYPE))
5676 hphi->set_op (i, new hsa_op_immed (op));
5677 else if (POINTER_TYPE_P (TREE_TYPE (lhs))
5678 && TREE_CODE (op) == INTEGER_CST)
5680 /* Handle assignment of NULL value to a pointer type. */
5681 hphi->set_op (i, new hsa_op_immed (op));
5683 else if (TREE_CODE (op) == ADDR_EXPR)
5685 edge e = gimple_phi_arg_edge (as_a <gphi *> (phi_stmt), i);
5686 hsa_bb *hbb_src = hsa_init_new_bb (split_edge (e));
5687 hsa_op_address *addr = gen_hsa_addr (TREE_OPERAND (op, 0),
5691 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
5692 hsa_insn_basic *insn
5693 = new hsa_insn_basic (2, BRIG_OPCODE_LDA, BRIG_TYPE_U64,
5695 hbb_src->append_insn (insn);
5697 hphi->set_op (i, dest);
5701 HSA_SORRY_AT (gimple_location (phi_stmt),
5702 "support for HSA does not handle PHI nodes with "
5703 "constant address operands");
5709 hbb->append_phi (hphi);
5712 /* Constructor of class containing HSA-specific information about a basic
5713 block. CFG_BB is the CFG BB this HSA BB is associated with. IDX is the new
5714 index of this BB (so that the constructor does not attempt to use
5715 hsa_cfun during its construction). */
5717 hsa_bb::hsa_bb (basic_block cfg_bb, int idx)
5718 : m_bb (cfg_bb), m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL),
5719 m_last_phi (NULL), m_index (idx)
5721 gcc_assert (!cfg_bb->aux);
5725 /* Constructor of class containing HSA-specific information about a basic
5726 block. CFG_BB is the CFG BB this HSA BB is associated with. */
5728 hsa_bb::hsa_bb (basic_block cfg_bb)
5729 : m_bb (cfg_bb), m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL),
5730 m_last_phi (NULL), m_index (hsa_cfun->m_hbb_count++)
5732 gcc_assert (!cfg_bb->aux);
5736 /* Create and initialize and return a new hsa_bb structure for a given CFG
5740 hsa_init_new_bb (basic_block bb)
5742 void *m = obstack_alloc (&hsa_obstack, sizeof (hsa_bb));
5743 return new (m) hsa_bb (bb);
5746 /* Initialize OMP in an HSA basic block PROLOGUE. */
5749 init_prologue (void)
5751 if (!hsa_cfun->m_kern_p)
5754 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
5756 /* Create a magic number that is going to be printed by libgomp. */
5757 unsigned index = hsa_get_number_decl_kernel_mappings ();
5759 /* Emit store to debug argument. */
5760 if (PARAM_VALUE (PARAM_HSA_GEN_DEBUG_STORES) > 0)
5761 set_debug_value (prologue, new hsa_op_immed (1000 + index, BRIG_TYPE_U64));
5764 /* Initialize hsa_num_threads to a default value. */
5767 init_hsa_num_threads (void)
5769 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
5771 /* Save the default value to private variable hsa_num_threads. */
5772 hsa_insn_basic *basic
5773 = new hsa_insn_mem (BRIG_OPCODE_ST, hsa_num_threads->m_type,
5774 new hsa_op_immed (0, hsa_num_threads->m_type),
5775 new hsa_op_address (hsa_num_threads));
5776 prologue->append_insn (basic);
5779 /* Go over gimple representation and generate our internal HSA one. */
5782 gen_body_from_gimple ()
5786 /* Verify CFG for complex edges we are unable to handle. */
5790 FOR_EACH_BB_FN (bb, cfun)
5792 FOR_EACH_EDGE (e, ei, bb->succs)
5794 /* Verify all unsupported flags for edges that point
5795 to the same basic block. */
5796 if (e->flags & EDGE_EH)
5798 HSA_SORRY_AT (UNKNOWN_LOCATION,
5799 "support for HSA does not implement exception "
5806 FOR_EACH_BB_FN (bb, cfun)
5808 gimple_stmt_iterator gsi;
5809 hsa_bb *hbb = hsa_bb_for_bb (bb);
5813 hbb = hsa_init_new_bb (bb);
5815 for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
5817 gen_hsa_insns_for_gimple_stmt (gsi_stmt (gsi), hbb);
5818 if (hsa_seen_error ())
5823 FOR_EACH_BB_FN (bb, cfun)
5825 gimple_stmt_iterator gsi;
5826 hsa_bb *hbb = hsa_bb_for_bb (bb);
5827 gcc_assert (hbb != NULL);
5829 for (gsi = gsi_start_phis (bb); !gsi_end_p (gsi); gsi_next (&gsi))
5830 if (!virtual_operand_p (gimple_phi_result (gsi_stmt (gsi))))
5831 gen_hsa_phi_from_gimple_phi (gsi_stmt (gsi), hbb);
5834 if (dump_file && (dump_flags & TDF_DETAILS))
5836 fprintf (dump_file, "------- Generated SSA form -------\n");
5837 dump_hsa_cfun (dump_file);
5842 gen_function_decl_parameters (hsa_function_representation *f,
5848 for (parm = TYPE_ARG_TYPES (TREE_TYPE (decl)), i = 0;
5850 parm = TREE_CHAIN (parm), i++)
5852 /* Result type if last in the tree list. */
5853 if (TREE_CHAIN (parm) == NULL)
5856 tree v = TREE_VALUE (parm);
5858 hsa_symbol *arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
5860 arg->m_type = hsa_type_for_tree_type (v, &arg->m_dim);
5861 arg->m_name_number = i;
5863 f->m_input_args.safe_push (arg);
5866 tree result_type = TREE_TYPE (TREE_TYPE (decl));
5867 if (!VOID_TYPE_P (result_type))
5869 f->m_output_arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
5871 f->m_output_arg->m_type
5872 = hsa_type_for_tree_type (result_type, &f->m_output_arg->m_dim);
5873 f->m_output_arg->m_name = "res";
5877 /* Generate the vector of parameters of the HSA representation of the current
5878 function. This also includes the output parameter representing the
5882 gen_function_def_parameters ()
5886 hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
5888 for (parm = DECL_ARGUMENTS (cfun->decl); parm;
5889 parm = DECL_CHAIN (parm))
5891 struct hsa_symbol **slot;
5894 = new hsa_symbol (BRIG_TYPE_NONE, hsa_cfun->m_kern_p
5895 ? BRIG_SEGMENT_KERNARG : BRIG_SEGMENT_ARG,
5896 BRIG_LINKAGE_FUNCTION);
5897 arg->fillup_for_decl (parm);
5899 hsa_cfun->m_input_args.safe_push (arg);
5901 if (hsa_seen_error ())
5904 arg->m_name = hsa_get_declaration_name (parm);
5906 /* Copy all input arguments and create corresponding private symbols
5908 hsa_symbol *private_arg;
5909 hsa_op_address *parm_addr = new hsa_op_address (arg);
5911 if (TREE_ADDRESSABLE (parm)
5912 || (!is_gimple_reg (parm) && !TREE_READONLY (parm)))
5914 private_arg = hsa_cfun->create_hsa_temporary (arg->m_type);
5915 private_arg->fillup_for_decl (parm);
5917 BrigAlignment8_t align = MIN (arg->m_align, private_arg->m_align);
5919 hsa_op_address *private_arg_addr = new hsa_op_address (private_arg);
5920 gen_hsa_memory_copy (prologue, private_arg_addr, parm_addr,
5921 arg->total_byte_size (), align);
5926 slot = hsa_cfun->m_local_symbols->find_slot (private_arg, INSERT);
5927 gcc_assert (!*slot);
5928 *slot = private_arg;
5930 if (is_gimple_reg (parm))
5932 tree ddef = ssa_default_def (cfun, parm);
5933 if (ddef && !has_zero_uses (ddef))
5935 BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (ddef),
5937 BrigType16_t mtype = mem_type_for_type (t);
5938 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (ddef);
5939 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mtype,
5941 gcc_assert (!parm_addr->m_reg);
5942 prologue->append_insn (mem);
5947 if (!VOID_TYPE_P (TREE_TYPE (TREE_TYPE (cfun->decl))))
5949 struct hsa_symbol **slot;
5951 hsa_cfun->m_output_arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
5952 BRIG_LINKAGE_FUNCTION);
5953 hsa_cfun->m_output_arg->fillup_for_decl (DECL_RESULT (cfun->decl));
5955 if (hsa_seen_error ())
5958 hsa_cfun->m_output_arg->m_name = "res";
5959 slot = hsa_cfun->m_local_symbols->find_slot (hsa_cfun->m_output_arg,
5961 gcc_assert (!*slot);
5962 *slot = hsa_cfun->m_output_arg;
5966 /* Generate function representation that corresponds to
5967 a function declaration. */
5969 hsa_function_representation *
5970 hsa_generate_function_declaration (tree decl)
5972 hsa_function_representation *fun
5973 = new hsa_function_representation (decl, false, 0);
5975 fun->m_declaration_p = true;
5976 fun->m_name = get_brig_function_name (decl);
5977 gen_function_decl_parameters (fun, decl);
5983 /* Generate function representation that corresponds to
5986 hsa_function_representation *
5987 hsa_generate_internal_fn_decl (hsa_internal_fn *fn)
5989 hsa_function_representation *fun = new hsa_function_representation (fn);
5991 fun->m_name = fn->name ();
5993 for (unsigned i = 0; i < fn->get_arity (); i++)
5996 = new hsa_symbol (fn->get_argument_type (i), BRIG_SEGMENT_ARG,
5998 arg->m_name_number = i;
5999 fun->m_input_args.safe_push (arg);
6002 fun->m_output_arg = new hsa_symbol (fn->get_argument_type (-1),
6003 BRIG_SEGMENT_ARG, BRIG_LINKAGE_NONE);
6004 fun->m_output_arg->m_name = "res";
6009 /* Return true if switch statement S can be transformed
6010 to a SBR instruction in HSAIL. */
6013 transformable_switch_to_sbr_p (gswitch *s)
6015 /* Identify if a switch statement can be transformed to
6016 SBR instruction, like:
6018 sbr_u32 $s1 [@label1, @label2, @label3];
6021 tree size = get_switch_size (s);
6022 if (!tree_fits_uhwi_p (size))
6025 if (tree_to_uhwi (size) > HSA_MAXIMUM_SBR_LABELS)
6031 /* Structure hold connection between PHI nodes and immediate
6032 values hold by there nodes. */
6034 struct phi_definition
6036 phi_definition (unsigned phi_i, unsigned label_i, tree imm):
6037 phi_index (phi_i), label_index (label_i), phi_value (imm)
6041 unsigned label_index;
6045 /* Sum slice of a vector V, starting from index START and ending
6046 at the index END - 1. */
6048 template <typename T>
6050 T sum_slice (const auto_vec <T> &v, unsigned start, unsigned end,
6055 for (unsigned i = start; i < end; i++)
6061 /* Function transforms GIMPLE SWITCH statements to a series of IF statements.
6062 Let's assume following example:
6076 The transformation encompasses following steps:
6077 1) all immediate values used by edges coming from the switch basic block
6079 2) all these edges are removed
6080 3) the switch statement (in L0) is replaced by:
6086 4) newly created basic block Lx' is used for generation of
6088 5) else branch of the last condition goes to LD
6089 6) fix all immediate values in PHI nodes that were propagated though
6090 edges that were removed in step 2
6092 Note: if a case is made by a range C1..C2, then process
6093 following transformation:
6095 switch_cond_op1 = C1 <= index;
6096 switch_cond_op2 = index <= C2;
6097 switch_cond_and = switch_cond_op1 & switch_cond_op2;
6098 if (switch_cond_and != 0)
6106 convert_switch_statements (void)
6108 function *func = DECL_STRUCT_FUNCTION (current_function_decl);
6111 bool modified_cfg = false;
6113 FOR_EACH_BB_FN (bb, func)
6115 gimple_stmt_iterator gsi = gsi_last_bb (bb);
6116 if (gsi_end_p (gsi))
6119 gimple *stmt = gsi_stmt (gsi);
6121 if (gimple_code (stmt) == GIMPLE_SWITCH)
6123 gswitch *s = as_a <gswitch *> (stmt);
6125 /* If the switch can utilize SBR insn, skip the statement. */
6126 if (transformable_switch_to_sbr_p (s))
6129 modified_cfg = true;
6131 unsigned labels = gimple_switch_num_labels (s);
6132 tree index = gimple_switch_index (s);
6133 tree index_type = TREE_TYPE (index);
6134 tree default_label = gimple_switch_default_label (s);
6135 basic_block default_label_bb
6136 = label_to_block_fn (func, CASE_LABEL (default_label));
6137 basic_block cur_bb = bb;
6139 auto_vec <edge> new_edges;
6140 auto_vec <phi_definition *> phi_todo_list;
6141 auto_vec <profile_count> edge_counts;
6142 auto_vec <profile_probability> edge_probabilities;
6144 /* Investigate all labels that and PHI nodes in these edges which
6145 should be fixed after we add new collection of edges. */
6146 for (unsigned i = 0; i < labels; i++)
6148 tree label = gimple_switch_label (s, i);
6149 basic_block label_bb = label_to_block_fn (func, CASE_LABEL (label));
6150 edge e = find_edge (bb, label_bb);
6151 edge_counts.safe_push (e->count);
6152 edge_probabilities.safe_push (e->probability);
6153 gphi_iterator phi_gsi;
6155 /* Save PHI definitions that will be destroyed because of an edge
6156 is going to be removed. */
6157 unsigned phi_index = 0;
6158 for (phi_gsi = gsi_start_phis (e->dest);
6159 !gsi_end_p (phi_gsi); gsi_next (&phi_gsi))
6161 gphi *phi = phi_gsi.phi ();
6162 for (unsigned j = 0; j < gimple_phi_num_args (phi); j++)
6164 if (gimple_phi_arg_edge (phi, j) == e)
6166 tree imm = gimple_phi_arg_def (phi, j);
6167 phi_definition *p = new phi_definition (phi_index, i,
6169 phi_todo_list.safe_push (p);
6177 /* Remove all edges for the current basic block. */
6178 for (int i = EDGE_COUNT (bb->succs) - 1; i >= 0; i--)
6180 edge e = EDGE_SUCC (bb, i);
6184 /* Iterate all non-default labels. */
6185 for (unsigned i = 1; i < labels; i++)
6187 tree label = gimple_switch_label (s, i);
6188 tree low = CASE_LOW (label);
6189 tree high = CASE_HIGH (label);
6191 if (!useless_type_conversion_p (TREE_TYPE (low), index_type))
6192 low = fold_convert (index_type, low);
6194 gimple_stmt_iterator cond_gsi = gsi_last_bb (cur_bb);
6198 tree tmp1 = make_temp_ssa_name (boolean_type_node, NULL,
6201 gimple *assign1 = gimple_build_assign (tmp1, LE_EXPR, low,
6204 tree tmp2 = make_temp_ssa_name (boolean_type_node, NULL,
6207 if (!useless_type_conversion_p (TREE_TYPE (high), index_type))
6208 high = fold_convert (index_type, high);
6209 gimple *assign2 = gimple_build_assign (tmp2, LE_EXPR, index,
6212 tree tmp3 = make_temp_ssa_name (boolean_type_node, NULL,
6214 gimple *assign3 = gimple_build_assign (tmp3, BIT_AND_EXPR, tmp1,
6217 gsi_insert_before (&cond_gsi, assign1, GSI_SAME_STMT);
6218 gsi_insert_before (&cond_gsi, assign2, GSI_SAME_STMT);
6219 gsi_insert_before (&cond_gsi, assign3, GSI_SAME_STMT);
6221 tree b = constant_boolean_node (false, boolean_type_node);
6222 c = gimple_build_cond (NE_EXPR, tmp3, b, NULL, NULL);
6225 c = gimple_build_cond (EQ_EXPR, index, low, NULL, NULL);
6227 gimple_set_location (c, gimple_location (stmt));
6229 gsi_insert_before (&cond_gsi, c, GSI_SAME_STMT);
6231 basic_block label_bb
6232 = label_to_block_fn (func, CASE_LABEL (label));
6233 edge new_edge = make_edge (cur_bb, label_bb, EDGE_TRUE_VALUE);
6234 profile_probability prob_sum = sum_slice <profile_probability>
6235 (edge_probabilities, i, labels, profile_probability::never ())
6236 + edge_probabilities[0];
6238 if (prob_sum.initialized_p ())
6239 new_edge->probability = edge_probabilities[i] / prob_sum;
6241 new_edge->count = edge_counts[i];
6242 new_edges.safe_push (new_edge);
6246 /* Prepare another basic block that will contain
6248 basic_block next_bb = create_empty_bb (cur_bb);
6251 add_bb_to_loop (next_bb, cur_bb->loop_father);
6252 loops_state_set (LOOPS_NEED_FIXUP);
6255 edge next_edge = make_edge (cur_bb, next_bb, EDGE_FALSE_VALUE);
6256 next_edge->probability = new_edge->probability.invert ();
6257 next_edge->count = edge_counts[0]
6258 + sum_slice <profile_count> (edge_counts, i, labels,
6259 profile_count::zero ());
6260 next_bb->frequency = EDGE_FREQUENCY (next_edge);
6263 else /* Link last IF statement and default label
6266 edge e = make_edge (cur_bb, default_label_bb, EDGE_FALSE_VALUE);
6267 e->probability = new_edge->probability.invert ();
6268 e->count = edge_counts[0];
6269 new_edges.safe_insert (0, e);
6273 /* Restore original PHI immediate value. */
6274 for (unsigned i = 0; i < phi_todo_list.length (); i++)
6276 phi_definition *phi_def = phi_todo_list[i];
6277 edge new_edge = new_edges[phi_def->label_index];
6279 gphi_iterator it = gsi_start_phis (new_edge->dest);
6280 for (unsigned i = 0; i < phi_def->phi_index; i++)
6283 gphi *phi = it.phi ();
6284 add_phi_arg (phi, phi_def->phi_value, new_edge, UNKNOWN_LOCATION);
6288 /* Remove the original GIMPLE switch statement. */
6289 gsi_remove (&gsi, true);
6294 dump_function_to_file (current_function_decl, dump_file, TDF_DETAILS);
6296 return modified_cfg;
6299 /* Expand builtins that can't be handled by HSA back-end. */
6304 function *func = DECL_STRUCT_FUNCTION (current_function_decl);
6307 FOR_EACH_BB_FN (bb, func)
6309 for (gimple_stmt_iterator gsi = gsi_start_bb (bb); !gsi_end_p (gsi);
6312 gimple *stmt = gsi_stmt (gsi);
6314 if (gimple_code (stmt) != GIMPLE_CALL)
6317 gcall *call = as_a <gcall *> (stmt);
6319 if (!gimple_call_builtin_p (call, BUILT_IN_NORMAL))
6322 tree fndecl = gimple_call_fndecl (stmt);
6323 enum built_in_function fn = DECL_FUNCTION_CODE (fndecl);
6326 case BUILT_IN_CEXPF:
6327 case BUILT_IN_CEXPIF:
6328 case BUILT_IN_CEXPI:
6330 /* Similar to builtins.c (expand_builtin_cexpi), the builtin
6331 can be transformed to: cexp(I * z) = ccos(z) + I * csin(z). */
6332 tree lhs = gimple_call_lhs (stmt);
6333 tree rhs = gimple_call_arg (stmt, 0);
6334 tree rhs_type = TREE_TYPE (rhs);
6335 bool float_type_p = rhs_type == float_type_node;
6336 tree real_part = make_temp_ssa_name (rhs_type, NULL,
6338 tree imag_part = make_temp_ssa_name (rhs_type, NULL,
6342 = mathfn_built_in (rhs_type, fn == float_type_p
6343 ? BUILT_IN_COSF : BUILT_IN_COS);
6344 gcall *cos = gimple_build_call (cos_fndecl, 1, rhs);
6345 gimple_call_set_lhs (cos, real_part);
6346 gsi_insert_before (&gsi, cos, GSI_SAME_STMT);
6349 = mathfn_built_in (rhs_type, fn == float_type_p
6350 ? BUILT_IN_SINF : BUILT_IN_SIN);
6351 gcall *sin = gimple_build_call (sin_fndecl, 1, rhs);
6352 gimple_call_set_lhs (sin, imag_part);
6353 gsi_insert_before (&gsi, sin, GSI_SAME_STMT);
6356 gassign *assign = gimple_build_assign (lhs, COMPLEX_EXPR,
6357 real_part, imag_part);
6358 gsi_insert_before (&gsi, assign, GSI_SAME_STMT);
6359 gsi_remove (&gsi, true);
6370 /* Emit HSA module variables that are global for the entire module. */
6373 emit_hsa_module_variables (void)
6375 hsa_num_threads = new hsa_symbol (BRIG_TYPE_U32, BRIG_SEGMENT_PRIVATE,
6376 BRIG_LINKAGE_MODULE, true);
6378 hsa_num_threads->m_name = "hsa_num_threads";
6380 hsa_brig_emit_omp_symbols ();
6383 /* Generate HSAIL representation of the current function and write into a
6384 special section of the output file. If KERNEL is set, the function will be
6385 considered an HSA kernel callable from the host, otherwise it will be
6386 compiled as an HSA function callable from other HSA code. */
6389 generate_hsa (bool kernel)
6391 hsa_init_data_for_cfun ();
6393 if (hsa_num_threads == NULL)
6394 emit_hsa_module_variables ();
6396 bool modified_cfg = convert_switch_statements ();
6397 /* Initialize hsa_cfun. */
6398 hsa_cfun = new hsa_function_representation (cfun->decl, kernel,
6399 SSANAMES (cfun)->length (),
6401 hsa_cfun->init_extra_bbs ();
6405 HSA_SORRY_AT (UNKNOWN_LOCATION,
6406 "support for HSA does not implement transactional memory");
6410 verify_function_arguments (cfun->decl);
6411 if (hsa_seen_error ())
6414 hsa_cfun->m_name = get_brig_function_name (cfun->decl);
6416 gen_function_def_parameters ();
6417 if (hsa_seen_error ())
6422 gen_body_from_gimple ();
6423 if (hsa_seen_error ())
6426 if (hsa_cfun->m_kernel_dispatch_count)
6427 init_hsa_num_threads ();
6429 if (hsa_cfun->m_kern_p)
6431 hsa_function_summary *s
6432 = hsa_summaries->get (cgraph_node::get (hsa_cfun->m_decl));
6433 hsa_add_kern_decl_mapping (current_function_decl, hsa_cfun->m_name,
6434 hsa_cfun->m_maximum_omp_data_size,
6435 s->m_gridified_kernel_p);
6440 for (unsigned i = 0; i < hsa_cfun->m_ssa_map.length (); i++)
6441 if (hsa_cfun->m_ssa_map[i])
6442 hsa_cfun->m_ssa_map[i]->verify_ssa ();
6445 FOR_EACH_BB_FN (bb, cfun)
6447 hsa_bb *hbb = hsa_bb_for_bb (bb);
6449 for (hsa_insn_basic *insn = hbb->m_first_insn; insn;
6450 insn = insn->m_next)
6456 hsa_brig_emit_function ();
6459 hsa_deinit_data_for_cfun ();
6464 const pass_data pass_data_gen_hsail =
6467 "hsagen", /* name */
6468 OPTGROUP_OMP, /* optinfo_flags */
6469 TV_NONE, /* tv_id */
6470 PROP_cfg | PROP_ssa, /* properties_required */
6471 0, /* properties_provided */
6472 0, /* properties_destroyed */
6473 0, /* todo_flags_start */
6474 0 /* todo_flags_finish */
6477 class pass_gen_hsail : public gimple_opt_pass
6480 pass_gen_hsail (gcc::context *ctxt)
6481 : gimple_opt_pass(pass_data_gen_hsail, ctxt)
6484 /* opt_pass methods: */
6485 bool gate (function *);
6486 unsigned int execute (function *);
6488 }; // class pass_gen_hsail
6490 /* Determine whether or not to run generation of HSAIL. */
6493 pass_gen_hsail::gate (function *f)
6495 return hsa_gen_requested_p ()
6496 && hsa_gpu_implementation_p (f->decl);
6500 pass_gen_hsail::execute (function *)
6502 hsa_function_summary *s
6503 = hsa_summaries->get (cgraph_node::get_create (current_function_decl));
6506 generate_hsa (s->m_kind == HSA_KERNEL);
6507 TREE_ASM_WRITTEN (current_function_decl) = 1;
6508 return TODO_discard_function;
6513 /* Create the instance of hsa gen pass. */
6516 make_pass_gen_hsail (gcc::context *ctxt)
6518 return new pass_gen_hsail (ctxt);