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"
63 #include "stringpool.h"
66 /* Print a warning message and set that we have seen an error. */
68 #define HSA_SORRY_ATV(location, message, ...) \
72 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
74 inform (location, message, __VA_ARGS__); \
78 /* Same as previous, but highlight a location. */
80 #define HSA_SORRY_AT(location, message) \
84 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
86 inform (location, message); \
90 /* Default number of threads used by kernel dispatch. */
92 #define HSA_DEFAULT_NUM_THREADS 64
94 /* Following structures are defined in the final version
95 of HSA specification. */
97 /* HSA queue packet is shadow structure, originally provided by AMD. */
99 struct hsa_queue_packet
103 uint16_t workgroup_size_x
;
104 uint16_t workgroup_size_y
;
105 uint16_t workgroup_size_z
;
107 uint32_t grid_size_x
;
108 uint32_t grid_size_y
;
109 uint32_t grid_size_z
;
110 uint32_t private_segment_size
;
111 uint32_t group_segment_size
;
112 uint64_t kernel_object
;
113 void *kernarg_address
;
115 uint64_t completion_signal
;
118 /* HSA queue is shadow structure, originally provided by AMD. */
125 uint64_t doorbell_signal
;
131 static struct obstack hsa_obstack
;
133 /* List of pointers to all instructions that come from an object allocator. */
134 static vec
<hsa_insn_basic
*> hsa_instructions
;
136 /* List of pointers to all operands that come from an object allocator. */
137 static vec
<hsa_op_base
*> hsa_operands
;
139 hsa_symbol::hsa_symbol ()
140 : m_decl (NULL_TREE
), m_name (NULL
), m_name_number (0),
141 m_directive_offset (0), m_type (BRIG_TYPE_NONE
),
142 m_segment (BRIG_SEGMENT_NONE
), m_linkage (BRIG_LINKAGE_NONE
), m_dim (0),
143 m_cst_value (NULL
), m_global_scope_p (false), m_seen_error (false),
144 m_allocation (BRIG_ALLOCATION_AUTOMATIC
), m_emitted_to_brig (false)
149 hsa_symbol::hsa_symbol (BrigType16_t type
, BrigSegment8_t segment
,
150 BrigLinkage8_t linkage
, bool global_scope_p
,
151 BrigAllocation allocation
, BrigAlignment8_t align
)
152 : m_decl (NULL_TREE
), m_name (NULL
), m_name_number (0),
153 m_directive_offset (0), m_type (type
), m_segment (segment
),
154 m_linkage (linkage
), m_dim (0), m_cst_value (NULL
),
155 m_global_scope_p (global_scope_p
), m_seen_error (false),
156 m_allocation (allocation
), m_emitted_to_brig (false), m_align (align
)
160 unsigned HOST_WIDE_INT
161 hsa_symbol::total_byte_size ()
163 unsigned HOST_WIDE_INT s
164 = hsa_type_bit_size (~BRIG_TYPE_ARRAY_MASK
& m_type
);
165 gcc_assert (s
% BITS_PER_UNIT
== 0);
174 /* Forward declaration. */
177 hsa_type_for_tree_type (const_tree type
, unsigned HOST_WIDE_INT
*dim_p
,
181 hsa_symbol::fillup_for_decl (tree decl
)
184 m_type
= hsa_type_for_tree_type (TREE_TYPE (decl
), &m_dim
, false);
185 if (hsa_seen_error ())
191 m_align
= MAX (m_align
, hsa_natural_alignment (m_type
));
194 /* Constructor of class representing global HSA function/kernel information and
195 state. FNDECL is function declaration, KERNEL_P is true if the function
196 is going to become a HSA kernel. If the function has body, SSA_NAMES_COUNT
197 should be set to number of SSA names used in the function.
198 MODIFIED_CFG is set to true in case we modified control-flow graph
201 hsa_function_representation::hsa_function_representation
202 (tree fdecl
, bool kernel_p
, unsigned ssa_names_count
, bool modified_cfg
)
204 m_reg_count (0), m_input_args (vNULL
),
205 m_output_arg (NULL
), m_spill_symbols (vNULL
), m_global_symbols (vNULL
),
206 m_private_variables (vNULL
), m_called_functions (vNULL
),
207 m_called_internal_fns (vNULL
), m_hbb_count (0),
208 m_in_ssa (true), m_kern_p (kernel_p
), m_declaration_p (false),
209 m_decl (fdecl
), m_internal_fn (NULL
), m_shadow_reg (NULL
),
210 m_kernel_dispatch_count (0), m_maximum_omp_data_size (0),
211 m_seen_error (false), m_temp_symbol_count (0), m_ssa_map (),
212 m_modified_cfg (modified_cfg
)
214 int sym_init_len
= (vec_safe_length (cfun
->local_decls
) / 2) + 1;;
215 m_local_symbols
= new hash_table
<hsa_noop_symbol_hasher
> (sym_init_len
);
216 m_ssa_map
.safe_grow_cleared (ssa_names_count
);
219 /* Constructor of class representing HSA function information that
220 is derived for an internal function. */
221 hsa_function_representation::hsa_function_representation (hsa_internal_fn
*fn
)
222 : m_reg_count (0), m_input_args (vNULL
),
223 m_output_arg (NULL
), m_local_symbols (NULL
),
224 m_spill_symbols (vNULL
), m_global_symbols (vNULL
),
225 m_private_variables (vNULL
), m_called_functions (vNULL
),
226 m_called_internal_fns (vNULL
), m_hbb_count (0),
227 m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL
),
228 m_internal_fn (fn
), m_shadow_reg (NULL
), m_kernel_dispatch_count (0),
229 m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
232 /* Destructor of class holding function/kernel-wide information and state. */
234 hsa_function_representation::~hsa_function_representation ()
236 /* Kernel names are deallocated at the end of BRIG output when deallocating
237 hsa_decl_kernel_mapping. */
238 if (!m_kern_p
|| m_seen_error
)
241 for (unsigned i
= 0; i
< m_input_args
.length (); i
++)
242 delete m_input_args
[i
];
243 m_input_args
.release ();
246 delete m_local_symbols
;
248 for (unsigned i
= 0; i
< m_spill_symbols
.length (); i
++)
249 delete m_spill_symbols
[i
];
250 m_spill_symbols
.release ();
253 for (unsigned i
= 0; i
< m_global_symbols
.iterate (i
, &sym
); i
++)
254 if (sym
->m_linkage
!= BRIG_ALLOCATION_PROGRAM
)
256 m_global_symbols
.release ();
258 for (unsigned i
= 0; i
< m_private_variables
.length (); i
++)
259 delete m_private_variables
[i
];
260 m_private_variables
.release ();
261 m_called_functions
.release ();
262 m_ssa_map
.release ();
264 for (unsigned i
= 0; i
< m_called_internal_fns
.length (); i
++)
265 delete m_called_internal_fns
[i
];
269 hsa_function_representation::get_shadow_reg ()
271 /* If we compile a function with kernel dispatch and does not set
272 an optimization level, the function won't be inlined and
280 /* Append the shadow argument. */
281 hsa_symbol
*shadow
= new hsa_symbol (BRIG_TYPE_U64
, BRIG_SEGMENT_KERNARG
,
282 BRIG_LINKAGE_FUNCTION
);
283 m_input_args
.safe_push (shadow
);
284 shadow
->m_name
= "hsa_runtime_shadow";
286 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_U64
);
287 hsa_op_address
*addr
= new hsa_op_address (shadow
);
289 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, BRIG_TYPE_U64
, r
, addr
);
290 hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
))->append_insn (mem
);
296 bool hsa_function_representation::has_shadow_reg_p ()
298 return m_shadow_reg
!= NULL
;
302 hsa_function_representation::init_extra_bbs ()
304 hsa_init_new_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
305 hsa_init_new_bb (EXIT_BLOCK_PTR_FOR_FN (cfun
));
309 hsa_function_representation::update_dominance ()
313 free_dominance_info (CDI_DOMINATORS
);
314 calculate_dominance_info (CDI_DOMINATORS
);
319 hsa_function_representation::create_hsa_temporary (BrigType16_t type
)
321 hsa_symbol
*s
= new hsa_symbol (type
, BRIG_SEGMENT_PRIVATE
,
322 BRIG_LINKAGE_FUNCTION
);
323 s
->m_name_number
= m_temp_symbol_count
++;
325 hsa_cfun
->m_private_variables
.safe_push (s
);
330 hsa_function_representation::get_linkage ()
333 return BRIG_LINKAGE_PROGRAM
;
335 return m_kern_p
|| TREE_PUBLIC (m_decl
) ?
336 BRIG_LINKAGE_PROGRAM
: BRIG_LINKAGE_MODULE
;
339 /* Hash map of simple OMP builtins. */
340 static hash_map
<nofree_string_hash
, omp_simple_builtin
> *omp_simple_builtins
343 /* Warning messages for OMP builtins. */
345 #define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP " \
347 #define HSA_WARN_TIMING_ROUTINE "support for HSA does not implement OpenMP " \
349 #define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \
350 "undefined semantics within target regions, support for HSA ignores them"
351 #define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP " \
354 /* Initialize hash map with simple OMP builtins. */
357 hsa_init_simple_builtins ()
359 if (omp_simple_builtins
!= NULL
)
363 = new hash_map
<nofree_string_hash
, omp_simple_builtin
> ();
365 omp_simple_builtin omp_builtins
[] =
367 omp_simple_builtin ("omp_get_initial_device", NULL
, false,
368 new hsa_op_immed (GOMP_DEVICE_HOST
,
369 (BrigType16_t
) BRIG_TYPE_S32
)),
370 omp_simple_builtin ("omp_is_initial_device", NULL
, false,
371 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
372 omp_simple_builtin ("omp_get_dynamic", NULL
, false,
373 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
374 omp_simple_builtin ("omp_set_dynamic", NULL
, false, NULL
),
375 omp_simple_builtin ("omp_init_lock", HSA_WARN_LOCK_ROUTINE
, true),
376 omp_simple_builtin ("omp_init_lock_with_hint", HSA_WARN_LOCK_ROUTINE
,
378 omp_simple_builtin ("omp_init_nest_lock_with_hint", HSA_WARN_LOCK_ROUTINE
,
380 omp_simple_builtin ("omp_destroy_lock", HSA_WARN_LOCK_ROUTINE
, true),
381 omp_simple_builtin ("omp_set_lock", HSA_WARN_LOCK_ROUTINE
, true),
382 omp_simple_builtin ("omp_unset_lock", HSA_WARN_LOCK_ROUTINE
, true),
383 omp_simple_builtin ("omp_test_lock", HSA_WARN_LOCK_ROUTINE
, true),
384 omp_simple_builtin ("omp_get_wtime", HSA_WARN_TIMING_ROUTINE
, true),
385 omp_simple_builtin ("omp_get_wtick", HSA_WARN_TIMING_ROUTINE
, true),
386 omp_simple_builtin ("omp_target_alloc", HSA_WARN_MEMORY_ROUTINE
, false,
387 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_U64
)),
388 omp_simple_builtin ("omp_target_free", HSA_WARN_MEMORY_ROUTINE
, false),
389 omp_simple_builtin ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE
,
391 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
392 omp_simple_builtin ("omp_target_memcpy", HSA_WARN_MEMORY_ROUTINE
, false,
393 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
394 omp_simple_builtin ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE
,
396 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
397 omp_simple_builtin ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE
,
399 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
400 omp_simple_builtin ("omp_target_disassociate_ptr",
401 HSA_WARN_MEMORY_ROUTINE
,
403 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
404 omp_simple_builtin ("omp_set_max_active_levels",
405 "Support for HSA only allows only one active level, "
406 "call to omp_set_max_active_levels will be ignored "
407 "in the generated HSAIL",
409 omp_simple_builtin ("omp_get_max_active_levels", NULL
, false,
410 new hsa_op_immed (1, (BrigType16_t
) BRIG_TYPE_S32
)),
411 omp_simple_builtin ("omp_in_final", NULL
, false,
412 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
413 omp_simple_builtin ("omp_get_proc_bind", HSA_WARN_AFFINITY
, false,
414 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
415 omp_simple_builtin ("omp_get_num_places", HSA_WARN_AFFINITY
, false,
416 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
417 omp_simple_builtin ("omp_get_place_num_procs", HSA_WARN_AFFINITY
, false,
418 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
419 omp_simple_builtin ("omp_get_place_proc_ids", HSA_WARN_AFFINITY
, false,
421 omp_simple_builtin ("omp_get_place_num", HSA_WARN_AFFINITY
, false,
422 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
423 omp_simple_builtin ("omp_get_partition_num_places", HSA_WARN_AFFINITY
,
425 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
426 omp_simple_builtin ("omp_get_partition_place_nums", HSA_WARN_AFFINITY
,
428 omp_simple_builtin ("omp_set_default_device",
429 "omp_set_default_device has undefined semantics "
430 "within target regions, support for HSA ignores it",
432 omp_simple_builtin ("omp_get_default_device",
433 "omp_get_default_device has undefined semantics "
434 "within target regions, support for HSA ignores it",
436 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
437 omp_simple_builtin ("omp_get_num_devices",
438 "omp_get_num_devices has undefined semantics "
439 "within target regions, support for HSA ignores it",
441 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
442 omp_simple_builtin ("omp_get_num_procs", NULL
, true, NULL
),
443 omp_simple_builtin ("omp_get_cancellation", NULL
, true, NULL
),
444 omp_simple_builtin ("omp_set_nested", NULL
, true, NULL
),
445 omp_simple_builtin ("omp_get_nested", NULL
, true, NULL
),
446 omp_simple_builtin ("omp_set_schedule", NULL
, true, NULL
),
447 omp_simple_builtin ("omp_get_schedule", NULL
, true, NULL
),
448 omp_simple_builtin ("omp_get_thread_limit", NULL
, true, NULL
),
449 omp_simple_builtin ("omp_get_team_size", NULL
, true, NULL
),
450 omp_simple_builtin ("omp_get_ancestor_thread_num", NULL
, true, NULL
),
451 omp_simple_builtin ("omp_get_max_task_priority", NULL
, true, NULL
)
454 unsigned count
= sizeof (omp_builtins
) / sizeof (omp_simple_builtin
);
456 for (unsigned i
= 0; i
< count
; i
++)
457 omp_simple_builtins
->put (omp_builtins
[i
].m_name
, omp_builtins
[i
]);
460 /* Allocate HSA structures that we need only while generating with this. */
463 hsa_init_data_for_cfun ()
465 hsa_init_compilation_unit_data ();
466 gcc_obstack_init (&hsa_obstack
);
469 /* Deinitialize HSA subsystem and free all allocated memory. */
472 hsa_deinit_data_for_cfun (void)
476 FOR_ALL_BB_FN (bb
, cfun
)
479 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
484 for (unsigned int i
= 0; i
< hsa_operands
.length (); i
++)
485 hsa_destroy_operand (hsa_operands
[i
]);
487 hsa_operands
.release ();
489 for (unsigned i
= 0; i
< hsa_instructions
.length (); i
++)
490 hsa_destroy_insn (hsa_instructions
[i
]);
492 hsa_instructions
.release ();
494 if (omp_simple_builtins
!= NULL
)
496 delete omp_simple_builtins
;
497 omp_simple_builtins
= NULL
;
500 obstack_free (&hsa_obstack
, NULL
);
504 /* Return the type which holds addresses in the given SEGMENT. */
507 hsa_get_segment_addr_type (BrigSegment8_t segment
)
511 case BRIG_SEGMENT_NONE
:
514 case BRIG_SEGMENT_FLAT
:
515 case BRIG_SEGMENT_GLOBAL
:
516 case BRIG_SEGMENT_READONLY
:
517 case BRIG_SEGMENT_KERNARG
:
518 return hsa_machine_large_p () ? BRIG_TYPE_U64
: BRIG_TYPE_U32
;
520 case BRIG_SEGMENT_GROUP
:
521 case BRIG_SEGMENT_PRIVATE
:
522 case BRIG_SEGMENT_SPILL
:
523 case BRIG_SEGMENT_ARG
:
524 return BRIG_TYPE_U32
;
529 /* Return integer brig type according to provided SIZE in bytes. If SIGN
530 is set to true, return signed integer type. */
533 get_integer_type_by_bytes (unsigned size
, bool sign
)
541 return BRIG_TYPE_S16
;
543 return BRIG_TYPE_S32
;
545 return BRIG_TYPE_S64
;
555 return BRIG_TYPE_U16
;
557 return BRIG_TYPE_U32
;
559 return BRIG_TYPE_U64
;
567 /* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
568 are assumed to use flat addressing. If min32int is true, always expand
569 integer types to one that has at least 32 bits. */
572 hsa_type_for_scalar_tree_type (const_tree type
, bool min32int
)
576 BrigType16_t res
= BRIG_TYPE_NONE
;
578 gcc_checking_assert (TYPE_P (type
));
579 gcc_checking_assert (!AGGREGATE_TYPE_P (type
));
580 if (POINTER_TYPE_P (type
))
581 return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
583 if (TREE_CODE (type
) == VECTOR_TYPE
|| TREE_CODE (type
) == COMPLEX_TYPE
)
584 base
= TREE_TYPE (type
);
588 if (!tree_fits_uhwi_p (TYPE_SIZE (base
)))
590 HSA_SORRY_ATV (EXPR_LOCATION (type
),
591 "support for HSA does not implement huge or "
592 "variable-sized type %qT", type
);
596 bsize
= tree_to_uhwi (TYPE_SIZE (base
));
597 unsigned byte_size
= bsize
/ BITS_PER_UNIT
;
598 if (INTEGRAL_TYPE_P (base
))
599 res
= get_integer_type_by_bytes (byte_size
, !TYPE_UNSIGNED (base
));
600 else if (SCALAR_FLOAT_TYPE_P (base
))
618 if (res
== BRIG_TYPE_NONE
)
620 HSA_SORRY_ATV (EXPR_LOCATION (type
),
621 "support for HSA does not implement type %qT", type
);
625 if (TREE_CODE (type
) == VECTOR_TYPE
)
627 HOST_WIDE_INT tsize
= tree_to_uhwi (TYPE_SIZE (type
));
631 HSA_SORRY_ATV (EXPR_LOCATION (type
),
632 "support for HSA does not implement a vector type "
633 "where a type and unit size are equal: %qT", type
);
640 res
|= BRIG_TYPE_PACK_32
;
643 res
|= BRIG_TYPE_PACK_64
;
646 res
|= BRIG_TYPE_PACK_128
;
649 HSA_SORRY_ATV (EXPR_LOCATION (type
),
650 "support for HSA does not implement type %qT", type
);
656 /* Registers/immediate operands can only be 32bit or more except for
658 if (res
== BRIG_TYPE_U8
|| res
== BRIG_TYPE_U16
)
660 else if (res
== BRIG_TYPE_S8
|| res
== BRIG_TYPE_S16
)
664 if (TREE_CODE (type
) == COMPLEX_TYPE
)
666 unsigned bsize
= 2 * hsa_type_bit_size (res
);
667 res
= hsa_bittype_for_bitsize (bsize
);
673 /* Returns the BRIG type we need to load/store entities of TYPE. */
676 mem_type_for_type (BrigType16_t type
)
678 /* HSA has non-intuitive constraints on load/store types. If it's
679 a bit-type it _must_ be B128, if it's not a bit-type it must be
680 64bit max. So for loading entities of 128 bits (e.g. vectors)
681 we have to to B128, while for loading the rest we have to use the
682 input type (??? or maybe also flattened to a equally sized non-vector
684 if ((type
& BRIG_TYPE_PACK_MASK
) == BRIG_TYPE_PACK_128
)
685 return BRIG_TYPE_B128
;
686 else if (hsa_btype_p (type
) || hsa_type_packed_p (type
))
688 unsigned bitsize
= hsa_type_bit_size (type
);
690 return hsa_uint_for_bitsize (bitsize
);
692 return hsa_bittype_for_bitsize (bitsize
);
697 /* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some
698 kind of array will be generated, setting DIM appropriately. Otherwise, it
699 will be set to zero. */
702 hsa_type_for_tree_type (const_tree type
, unsigned HOST_WIDE_INT
*dim_p
= NULL
,
703 bool min32int
= false)
705 gcc_checking_assert (TYPE_P (type
));
706 if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type
)))
708 HSA_SORRY_ATV (EXPR_LOCATION (type
), "support for HSA does not "
709 "implement huge or variable-sized type %qT", type
);
710 return BRIG_TYPE_NONE
;
713 if (RECORD_OR_UNION_TYPE_P (type
))
716 *dim_p
= tree_to_uhwi (TYPE_SIZE_UNIT (type
));
717 return BRIG_TYPE_U8
| BRIG_TYPE_ARRAY
;
720 if (TREE_CODE (type
) == ARRAY_TYPE
)
722 /* We try to be nice and use the real base-type when this is an array of
723 scalars and only resort to an array of bytes if the type is more
726 unsigned HOST_WIDE_INT dim
= 1;
728 while (TREE_CODE (type
) == ARRAY_TYPE
)
730 tree domain
= TYPE_DOMAIN (type
);
731 if (!TYPE_MIN_VALUE (domain
)
732 || !TYPE_MAX_VALUE (domain
)
733 || !tree_fits_shwi_p (TYPE_MIN_VALUE (domain
))
734 || !tree_fits_shwi_p (TYPE_MAX_VALUE (domain
)))
736 HSA_SORRY_ATV (EXPR_LOCATION (type
),
737 "support for HSA does not implement array "
738 "%qT with unknown bounds", type
);
739 return BRIG_TYPE_NONE
;
741 HOST_WIDE_INT min
= tree_to_shwi (TYPE_MIN_VALUE (domain
));
742 HOST_WIDE_INT max
= tree_to_shwi (TYPE_MAX_VALUE (domain
));
743 dim
= dim
* (unsigned HOST_WIDE_INT
) (max
- min
+ 1);
744 type
= TREE_TYPE (type
);
748 if (RECORD_OR_UNION_TYPE_P (type
))
750 dim
= dim
* tree_to_uhwi (TYPE_SIZE_UNIT (type
));
754 res
= hsa_type_for_scalar_tree_type (type
, false);
758 return res
| BRIG_TYPE_ARRAY
;
765 return hsa_type_for_scalar_tree_type (type
, min32int
);
768 /* Returns true if converting from STYPE into DTYPE needs the _CVT
769 opcode. If false a normal _MOV is enough. */
772 hsa_needs_cvt (BrigType16_t dtype
, BrigType16_t stype
)
774 if (hsa_btype_p (dtype
))
777 /* float <-> int conversions are real converts. */
778 if (hsa_type_float_p (dtype
) != hsa_type_float_p (stype
))
780 /* When both types have different size, then we need CVT as well. */
781 if (hsa_type_bit_size (dtype
) != hsa_type_bit_size (stype
))
786 /* Return declaration name if it exists or create one from UID if it does not.
787 If DECL is a local variable, make UID part of its name. */
790 hsa_get_declaration_name (tree decl
)
792 if (!DECL_NAME (decl
))
795 snprintf (buf
, 64, "__hsa_anon_%u", DECL_UID (decl
));
796 size_t len
= strlen (buf
);
797 char *copy
= (char *) obstack_alloc (&hsa_obstack
, len
+ 1);
798 memcpy (copy
, buf
, len
+ 1);
803 if (TREE_CODE (decl
) == FUNCTION_DECL
804 || (TREE_CODE (decl
) == VAR_DECL
&& is_global_var (decl
)))
805 name_tree
= DECL_ASSEMBLER_NAME (decl
);
807 name_tree
= DECL_NAME (decl
);
809 const char *name
= IDENTIFIER_POINTER (name_tree
);
810 /* User-defined assembly names have prepended asterisk symbol. */
814 if ((TREE_CODE (decl
) == VAR_DECL
)
815 && decl_function_context (decl
))
817 size_t len
= strlen (name
);
818 char *buf
= (char *) alloca (len
+ 32);
819 snprintf (buf
, len
+ 32, "%s_%u", name
, DECL_UID (decl
));
821 char *copy
= (char *) obstack_alloc (&hsa_obstack
, len
+ 1);
822 memcpy (copy
, buf
, len
+ 1);
829 /* Lookup or create the associated hsa_symbol structure with a given VAR_DECL
830 or lookup the hsa_structure corresponding to a PARM_DECL. */
833 get_symbol_for_decl (tree decl
)
836 hsa_symbol
dummy (BRIG_TYPE_NONE
, BRIG_SEGMENT_NONE
, BRIG_LINKAGE_NONE
);
838 gcc_assert (TREE_CODE (decl
) == PARM_DECL
839 || TREE_CODE (decl
) == RESULT_DECL
840 || TREE_CODE (decl
) == VAR_DECL
841 || TREE_CODE (decl
) == CONST_DECL
);
845 bool is_in_global_vars
= ((TREE_CODE (decl
) == VAR_DECL
)
846 && !decl_function_context (decl
));
848 if (is_in_global_vars
)
849 slot
= hsa_global_variable_symbols
->find_slot (&dummy
, INSERT
);
851 slot
= hsa_cfun
->m_local_symbols
->find_slot (&dummy
, INSERT
);
853 gcc_checking_assert (slot
);
856 hsa_symbol
*sym
= (*slot
);
858 /* If the symbol is problematic, mark current function also as
860 if (sym
->m_seen_error
)
863 /* PR hsa/70234: If a global variable was marked to be emitted,
864 but HSAIL generation of a function using the variable fails,
865 we should retry to emit the variable in context of a different
868 Iterate elements whether a symbol is already in m_global_symbols
870 if (is_in_global_vars
&& !sym
->m_emitted_to_brig
)
872 for (unsigned i
= 0; i
< hsa_cfun
->m_global_symbols
.length (); i
++)
873 if (hsa_cfun
->m_global_symbols
[i
] == sym
)
875 hsa_cfun
->m_global_symbols
.safe_push (sym
);
883 /* PARM_DECLs and RESULT_DECL should be already in m_local_symbols. */
884 gcc_assert (TREE_CODE (decl
) == VAR_DECL
885 || TREE_CODE (decl
) == CONST_DECL
);
886 BrigAlignment8_t align
= hsa_object_alignment (decl
);
888 if (is_in_global_vars
)
890 gcc_checking_assert (TREE_CODE (decl
) != CONST_DECL
);
891 sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_GLOBAL
,
892 BRIG_LINKAGE_PROGRAM
, true,
893 BRIG_ALLOCATION_PROGRAM
, align
);
894 hsa_cfun
->m_global_symbols
.safe_push (sym
);
895 sym
->fillup_for_decl (decl
);
896 if (sym
->m_align
> align
)
898 sym
->m_seen_error
= true;
899 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
900 "HSA specification requires that %E is at least "
901 "naturally aligned", decl
);
906 /* As generation of efficient memory copy instructions relies
907 on alignment greater or equal to 8 bytes,
908 we need to increase alignment of all aggregate types.. */
909 if (AGGREGATE_TYPE_P (TREE_TYPE (decl
)))
910 align
= MAX ((BrigAlignment8_t
) BRIG_ALIGNMENT_8
, align
);
912 BrigAllocation allocation
= BRIG_ALLOCATION_AUTOMATIC
;
913 BrigSegment8_t segment
;
914 if (TREE_CODE (decl
) == CONST_DECL
)
916 segment
= BRIG_SEGMENT_READONLY
;
917 allocation
= BRIG_ALLOCATION_AGENT
;
919 else if (lookup_attribute ("hsa_group_segment",
920 DECL_ATTRIBUTES (decl
)))
921 segment
= BRIG_SEGMENT_GROUP
;
922 else if (TREE_STATIC (decl
)
923 || lookup_attribute ("hsa_global_segment",
924 DECL_ATTRIBUTES (decl
)))
925 segment
= BRIG_SEGMENT_GLOBAL
;
927 segment
= BRIG_SEGMENT_PRIVATE
;
929 sym
= new hsa_symbol (BRIG_TYPE_NONE
, segment
, BRIG_LINKAGE_FUNCTION
,
930 false, allocation
, align
);
931 sym
->fillup_for_decl (decl
);
932 hsa_cfun
->m_private_variables
.safe_push (sym
);
935 sym
->m_name
= hsa_get_declaration_name (decl
);
941 /* For a given HSA function declaration, return a host
942 function declaration. */
945 hsa_get_host_function (tree decl
)
947 hsa_function_summary
*s
948 = hsa_summaries
->get (cgraph_node::get_create (decl
));
949 gcc_assert (s
->m_kind
!= HSA_NONE
);
950 gcc_assert (s
->m_gpu_implementation_p
);
952 return s
->m_bound_function
? s
->m_bound_function
->decl
: NULL
;
955 /* Return true if function DECL has a host equivalent function. */
958 get_brig_function_name (tree decl
)
962 hsa_function_summary
*s
= hsa_summaries
->get (cgraph_node::get_create (d
));
963 if (s
->m_kind
!= HSA_NONE
964 && s
->m_gpu_implementation_p
965 && s
->m_bound_function
)
966 d
= s
->m_bound_function
->decl
;
968 /* IPA split can create a function that has no host equivalent. */
972 char *name
= xstrdup (hsa_get_declaration_name (d
));
973 hsa_sanitize_name (name
);
978 /* Create a spill symbol of type TYPE. */
981 hsa_get_spill_symbol (BrigType16_t type
)
983 hsa_symbol
*sym
= new hsa_symbol (type
, BRIG_SEGMENT_SPILL
,
984 BRIG_LINKAGE_FUNCTION
);
985 hsa_cfun
->m_spill_symbols
.safe_push (sym
);
989 /* Create a symbol for a read-only string constant. */
991 hsa_get_string_cst_symbol (tree string_cst
)
993 gcc_checking_assert (TREE_CODE (string_cst
) == STRING_CST
);
995 hsa_symbol
**slot
= hsa_cfun
->m_string_constants_map
.get (string_cst
);
999 hsa_op_immed
*cst
= new hsa_op_immed (string_cst
);
1000 hsa_symbol
*sym
= new hsa_symbol (cst
->m_type
, BRIG_SEGMENT_GLOBAL
,
1001 BRIG_LINKAGE_MODULE
, true,
1002 BRIG_ALLOCATION_AGENT
);
1003 sym
->m_cst_value
= cst
;
1004 sym
->m_dim
= TREE_STRING_LENGTH (string_cst
);
1005 sym
->m_name_number
= hsa_cfun
->m_global_symbols
.length ();
1007 hsa_cfun
->m_global_symbols
.safe_push (sym
);
1008 hsa_cfun
->m_string_constants_map
.put (string_cst
, sym
);
1012 /* Constructor of the ancestor of all operands. K is BRIG kind that identified
1013 what the operator is. */
1015 hsa_op_base::hsa_op_base (BrigKind16_t k
)
1016 : m_next (NULL
), m_brig_op_offset (0), m_kind (k
)
1018 hsa_operands
.safe_push (this);
1021 /* Constructor of ancestor of all operands which have a type. K is BRIG kind
1022 that identified what the operator is. T is the type of the operator. */
1024 hsa_op_with_type::hsa_op_with_type (BrigKind16_t k
, BrigType16_t t
)
1025 : hsa_op_base (k
), m_type (t
)
1030 hsa_op_with_type::get_in_type (BrigType16_t dtype
, hsa_bb
*hbb
)
1032 if (m_type
== dtype
)
1037 if (hsa_needs_cvt (dtype
, m_type
))
1039 dest
= new hsa_op_reg (dtype
);
1040 hbb
->append_insn (new hsa_insn_cvt (dest
, this));
1042 else if (is_a
<hsa_op_reg
*> (this))
1044 /* In the end, HSA registers do not really have types, only sizes, so if
1045 the sizes match, we can use the register directly. */
1046 gcc_checking_assert (hsa_type_bit_size (dtype
)
1047 == hsa_type_bit_size (m_type
));
1052 dest
= new hsa_op_reg (m_type
);
1053 hbb
->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
1054 dest
->m_type
, dest
, this));
1056 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
1057 type of the operand must be same as type of the instruction. */
1058 dest
->m_type
= dtype
;
1064 /* Constructor of class representing HSA immediate values. TREE_VAL is the
1065 tree representation of the immediate value. If min32int is true,
1066 always expand integer types to one that has at least 32 bits. */
1068 hsa_op_immed::hsa_op_immed (tree tree_val
, bool min32int
)
1069 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
,
1070 hsa_type_for_tree_type (TREE_TYPE (tree_val
), NULL
,
1073 if (hsa_seen_error ())
1076 gcc_checking_assert ((is_gimple_min_invariant (tree_val
)
1077 && (!POINTER_TYPE_P (TREE_TYPE (tree_val
))
1078 || TREE_CODE (tree_val
) == INTEGER_CST
))
1079 || TREE_CODE (tree_val
) == CONSTRUCTOR
);
1080 m_tree_value
= tree_val
;
1082 /* Verify that all elements of a constructor are constants. */
1083 if (TREE_CODE (m_tree_value
) == CONSTRUCTOR
)
1084 for (unsigned i
= 0; i
< CONSTRUCTOR_NELTS (m_tree_value
); i
++)
1086 tree v
= CONSTRUCTOR_ELT (m_tree_value
, i
)->value
;
1087 if (!CONSTANT_CLASS_P (v
))
1089 HSA_SORRY_AT (EXPR_LOCATION (tree_val
),
1090 "HSA ctor should have only constants");
1096 /* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
1097 integer representation of the immediate value. TYPE is BRIG type. */
1099 hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value
, BrigType16_t type
)
1100 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
, type
),
1103 gcc_assert (hsa_type_integer_p (type
));
1104 m_int_value
= integer_value
;
1107 hsa_op_immed::hsa_op_immed ()
1108 : hsa_op_with_type (BRIG_KIND_NONE
, BRIG_TYPE_NONE
)
1112 /* New operator to allocate immediate operands from obstack. */
1115 hsa_op_immed::operator new (size_t size
)
1117 return obstack_alloc (&hsa_obstack
, size
);
1122 hsa_op_immed::~hsa_op_immed ()
1126 /* Change type of the immediate value to T. */
1129 hsa_op_immed::set_type (BrigType16_t t
)
1134 /* Constructor of class representing HSA registers and pseudo-registers. T is
1135 the BRIG type of the new register. */
1137 hsa_op_reg::hsa_op_reg (BrigType16_t t
)
1138 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER
, t
), m_gimple_ssa (NULL_TREE
),
1139 m_def_insn (NULL
), m_spill_sym (NULL
), m_order (hsa_cfun
->m_reg_count
++),
1140 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
1144 /* New operator to allocate a register from obstack. */
1147 hsa_op_reg::operator new (size_t size
)
1149 return obstack_alloc (&hsa_obstack
, size
);
1152 /* Verify register operand. */
1155 hsa_op_reg::verify_ssa ()
1157 /* Verify that each HSA register has a definition assigned.
1158 Exceptions are VAR_DECL and PARM_DECL that are a default
1160 gcc_checking_assert (m_def_insn
1161 || (m_gimple_ssa
!= NULL
1162 && (!SSA_NAME_VAR (m_gimple_ssa
)
1163 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa
))
1165 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa
)));
1167 /* Verify that every use of the register is really present
1168 in an instruction. */
1169 for (unsigned i
= 0; i
< m_uses
.length (); i
++)
1171 hsa_insn_basic
*use
= m_uses
[i
];
1173 bool is_visited
= false;
1174 for (unsigned j
= 0; j
< use
->operand_count (); j
++)
1176 hsa_op_base
*u
= use
->get_op (j
);
1177 hsa_op_address
*addr
; addr
= dyn_cast
<hsa_op_address
*> (u
);
1178 if (addr
&& addr
->m_reg
)
1183 bool r
= !addr
&& use
->op_output_p (j
);
1187 error ("HSA SSA name defined by instruction that is supposed "
1189 debug_hsa_operand (this);
1190 debug_hsa_insn (use
);
1191 internal_error ("HSA SSA verification failed");
1200 error ("HSA SSA name not among operands of instruction that is "
1201 "supposed to use it");
1202 debug_hsa_operand (this);
1203 debug_hsa_insn (use
);
1204 internal_error ("HSA SSA verification failed");
1209 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, hsa_op_reg
*r
,
1210 HOST_WIDE_INT offset
)
1211 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (r
),
1212 m_imm_offset (offset
)
1216 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, HOST_WIDE_INT offset
)
1217 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (NULL
),
1218 m_imm_offset (offset
)
1222 hsa_op_address::hsa_op_address (hsa_op_reg
*r
, HOST_WIDE_INT offset
)
1223 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (NULL
), m_reg (r
),
1224 m_imm_offset (offset
)
1228 /* New operator to allocate address operands from obstack. */
1231 hsa_op_address::operator new (size_t size
)
1233 return obstack_alloc (&hsa_obstack
, size
);
1236 /* Constructor of an operand referring to HSAIL code. */
1238 hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF
),
1239 m_directive_offset (0)
1243 /* Constructor of an operand representing a code list. Set it up so that it
1244 can contain ELEMENTS number of elements. */
1246 hsa_op_code_list::hsa_op_code_list (unsigned elements
)
1247 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST
)
1249 m_offsets
.create (1);
1250 m_offsets
.safe_grow_cleared (elements
);
1253 /* New operator to allocate code list operands from obstack. */
1256 hsa_op_code_list::operator new (size_t size
)
1258 return obstack_alloc (&hsa_obstack
, size
);
1261 /* Constructor of an operand representing an operand list.
1262 Set it up so that it can contain ELEMENTS number of elements. */
1264 hsa_op_operand_list::hsa_op_operand_list (unsigned elements
)
1265 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST
)
1267 m_offsets
.create (elements
);
1268 m_offsets
.safe_grow (elements
);
1271 /* New operator to allocate operand list operands from obstack. */
1274 hsa_op_operand_list::operator new (size_t size
)
1276 return obstack_alloc (&hsa_obstack
, size
);
1279 hsa_op_operand_list::~hsa_op_operand_list ()
1281 m_offsets
.release ();
1286 hsa_function_representation::reg_for_gimple_ssa (tree ssa
)
1290 gcc_checking_assert (TREE_CODE (ssa
) == SSA_NAME
);
1291 if (m_ssa_map
[SSA_NAME_VERSION (ssa
)])
1292 return m_ssa_map
[SSA_NAME_VERSION (ssa
)];
1294 hreg
= new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa
),
1296 hreg
->m_gimple_ssa
= ssa
;
1297 m_ssa_map
[SSA_NAME_VERSION (ssa
)] = hreg
;
1303 hsa_op_reg::set_definition (hsa_insn_basic
*insn
)
1305 if (hsa_cfun
->m_in_ssa
)
1307 gcc_checking_assert (!m_def_insn
);
1314 /* Constructor of the class which is the bases of all instructions and directly
1315 represents the most basic ones. NOPS is the number of operands that the
1316 operand vector will contain (and which will be cleared). OP is the opcode
1317 of the instruction. This constructor does not set type. */
1319 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
)
1321 m_next (NULL
), m_bb (NULL
), m_opcode (opc
), m_number (0),
1322 m_type (BRIG_TYPE_NONE
), m_brig_offset (0)
1325 m_operands
.safe_grow_cleared (nops
);
1327 hsa_instructions
.safe_push (this);
1330 /* Make OP the operand number INDEX of operands of this instruction. If OP is a
1331 register or an address containing a register, then either set the definition
1332 of the register to this instruction if it an output operand or add this
1333 instruction to the uses if it is an input one. */
1336 hsa_insn_basic::set_op (int index
, hsa_op_base
*op
)
1338 /* Each address operand is always use. */
1339 hsa_op_address
*addr
= dyn_cast
<hsa_op_address
*> (op
);
1340 if (addr
&& addr
->m_reg
)
1341 addr
->m_reg
->m_uses
.safe_push (this);
1344 hsa_op_reg
*reg
= dyn_cast
<hsa_op_reg
*> (op
);
1347 if (op_output_p (index
))
1348 reg
->set_definition (this);
1350 reg
->m_uses
.safe_push (this);
1354 m_operands
[index
] = op
;
1357 /* Get INDEX-th operand of the instruction. */
1360 hsa_insn_basic::get_op (int index
)
1362 return m_operands
[index
];
1365 /* Get address of INDEX-th operand of the instruction. */
1368 hsa_insn_basic::get_op_addr (int index
)
1370 return &m_operands
[index
];
1373 /* Get number of operands of the instruction. */
1375 hsa_insn_basic::operand_count ()
1377 return m_operands
.length ();
1380 /* Constructor of the class which is the bases of all instructions and directly
1381 represents the most basic ones. NOPS is the number of operands that the
1382 operand vector will contain (and which will be cleared). OPC is the opcode
1383 of the instruction, T is the type of the instruction. */
1385 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
, BrigType16_t t
,
1386 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1387 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1388 : m_prev (NULL
), m_next (NULL
), m_bb (NULL
), m_opcode (opc
),m_number (0),
1389 m_type (t
), m_brig_offset (0)
1392 m_operands
.safe_grow_cleared (nops
);
1396 gcc_checking_assert (nops
>= 1);
1402 gcc_checking_assert (nops
>= 2);
1408 gcc_checking_assert (nops
>= 3);
1414 gcc_checking_assert (nops
>= 4);
1418 hsa_instructions
.safe_push (this);
1421 /* New operator to allocate basic instruction from obstack. */
1424 hsa_insn_basic::operator new (size_t size
)
1426 return obstack_alloc (&hsa_obstack
, size
);
1429 /* Verify the instruction. */
1432 hsa_insn_basic::verify ()
1434 hsa_op_address
*addr
;
1437 /* Iterate all register operands and verify that the instruction
1438 is set in uses of the register. */
1439 for (unsigned i
= 0; i
< operand_count (); i
++)
1441 hsa_op_base
*use
= get_op (i
);
1443 if ((addr
= dyn_cast
<hsa_op_address
*> (use
)) && addr
->m_reg
)
1445 gcc_assert (addr
->m_reg
->m_def_insn
!= this);
1449 if ((reg
= dyn_cast
<hsa_op_reg
*> (use
)) && !op_output_p (i
))
1452 for (j
= 0; j
< reg
->m_uses
.length (); j
++)
1454 if (reg
->m_uses
[j
] == this)
1458 if (j
== reg
->m_uses
.length ())
1460 error ("HSA instruction uses a register but is not among "
1461 "recorded register uses");
1462 debug_hsa_operand (reg
);
1463 debug_hsa_insn (this);
1464 internal_error ("HSA instruction verification failed");
1470 /* Constructor of an instruction representing a PHI node. NOPS is the number
1471 of operands (equal to the number of predecessors). */
1473 hsa_insn_phi::hsa_insn_phi (unsigned nops
, hsa_op_reg
*dst
)
1474 : hsa_insn_basic (nops
, HSA_OPCODE_PHI
), m_dest (dst
)
1476 dst
->set_definition (this);
1479 /* Constructor of class representing instructions for control flow and
1482 hsa_insn_br::hsa_insn_br (unsigned nops
, int opc
, BrigType16_t t
,
1483 BrigWidth8_t width
, hsa_op_base
*arg0
,
1484 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1486 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1491 /* Constructor of class representing instruction for conditional jump, CTRL is
1492 the control register determining whether the jump will be carried out, the
1493 new instruction is automatically added to its uses list. */
1495 hsa_insn_cbr::hsa_insn_cbr (hsa_op_reg
*ctrl
)
1496 : hsa_insn_br (1, BRIG_OPCODE_CBR
, BRIG_TYPE_B1
, BRIG_WIDTH_1
, ctrl
)
1500 /* Constructor of class representing instruction for switch jump, CTRL is
1501 the index register. */
1503 hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg
*index
, unsigned jump_count
)
1504 : hsa_insn_basic (1, BRIG_OPCODE_SBR
, BRIG_TYPE_B1
, index
),
1505 m_width (BRIG_WIDTH_1
), m_jump_table (vNULL
),
1506 m_label_code_list (new hsa_op_code_list (jump_count
))
1510 /* Replace all occurrences of OLD_BB with NEW_BB in the statements
1514 hsa_insn_sbr::replace_all_labels (basic_block old_bb
, basic_block new_bb
)
1516 for (unsigned i
= 0; i
< m_jump_table
.length (); i
++)
1517 if (m_jump_table
[i
] == old_bb
)
1518 m_jump_table
[i
] = new_bb
;
1521 hsa_insn_sbr::~hsa_insn_sbr ()
1523 m_jump_table
.release ();
1526 /* Constructor of comparison instruction. CMP is the comparison operation and T
1527 is the result type. */
1529 hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp
, BrigType16_t t
,
1530 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1532 : hsa_insn_basic (3 , BRIG_OPCODE_CMP
, t
, arg0
, arg1
, arg2
), m_compare (cmp
)
1536 /* Constructor of classes representing memory accesses. OPC is the opcode (must
1537 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
1538 operands are provided as ARG0 and ARG1. */
1540 hsa_insn_mem::hsa_insn_mem (int opc
, BrigType16_t t
, hsa_op_base
*arg0
,
1542 : hsa_insn_basic (2, opc
, t
, arg0
, arg1
),
1543 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1545 gcc_checking_assert (opc
== BRIG_OPCODE_LD
|| opc
== BRIG_OPCODE_ST
);
1548 /* Constructor for descendants allowing different opcodes and number of
1549 operands, it passes its arguments directly to hsa_insn_basic
1550 constructor. The instruction operands are provided as ARG[0-3]. */
1553 hsa_insn_mem::hsa_insn_mem (unsigned nops
, int opc
, BrigType16_t t
,
1554 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1555 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1556 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1557 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1561 /* Constructor of class representing atomic instructions. OPC is the principal
1562 opcode, AOP is the specific atomic operation opcode. T is the type of the
1563 instruction. The instruction operands are provided as ARG[0-3]. */
1565 hsa_insn_atomic::hsa_insn_atomic (int nops
, int opc
,
1566 enum BrigAtomicOperation aop
,
1567 BrigType16_t t
, BrigMemoryOrder memorder
,
1569 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1571 : hsa_insn_mem (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
), m_atomicop (aop
),
1572 m_memoryorder (memorder
),
1573 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM
)
1575 gcc_checking_assert (opc
== BRIG_OPCODE_ATOMICNORET
||
1576 opc
== BRIG_OPCODE_ATOMIC
||
1577 opc
== BRIG_OPCODE_SIGNAL
||
1578 opc
== BRIG_OPCODE_SIGNALNORET
);
1581 /* Constructor of class representing signal instructions. OPC is the prinicpal
1582 opcode, SOP is the specific signal operation opcode. T is the type of the
1583 instruction. The instruction operands are provided as ARG[0-3]. */
1585 hsa_insn_signal::hsa_insn_signal (int nops
, int opc
,
1586 enum BrigAtomicOperation sop
,
1587 BrigType16_t t
, BrigMemoryOrder memorder
,
1588 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1589 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1590 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1591 m_memory_order (memorder
), m_signalop (sop
)
1595 /* Constructor of class representing segment conversion instructions. OPC is
1596 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
1597 and SRCT are destination and source types respectively, SEG is the segment
1598 we are converting to or from. The instruction operands are
1599 provided as ARG0 and ARG1. */
1601 hsa_insn_seg::hsa_insn_seg (int opc
, BrigType16_t dest
, BrigType16_t srct
,
1602 BrigSegment8_t seg
, hsa_op_base
*arg0
,
1604 : hsa_insn_basic (2, opc
, dest
, arg0
, arg1
), m_src_type (srct
),
1607 gcc_checking_assert (opc
== BRIG_OPCODE_STOF
|| opc
== BRIG_OPCODE_FTOS
);
1610 /* Constructor of class representing a call instruction. CALLEE is the tree
1611 representation of the function being called. */
1613 hsa_insn_call::hsa_insn_call (tree callee
)
1614 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (callee
),
1615 m_output_arg (NULL
), m_args_code_list (NULL
), m_result_code_list (NULL
)
1619 hsa_insn_call::hsa_insn_call (hsa_internal_fn
*fn
)
1620 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (NULL
),
1621 m_called_internal_fn (fn
), m_output_arg (NULL
), m_args_code_list (NULL
),
1622 m_result_code_list (NULL
)
1626 hsa_insn_call::~hsa_insn_call ()
1628 for (unsigned i
= 0; i
< m_input_args
.length (); i
++)
1629 delete m_input_args
[i
];
1631 delete m_output_arg
;
1633 m_input_args
.release ();
1634 m_input_arg_insns
.release ();
1637 /* Constructor of class representing the argument block required to invoke
1639 hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind
,
1640 hsa_insn_call
* call
)
1641 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK
), m_kind (brig_kind
),
1646 hsa_insn_comment::hsa_insn_comment (const char *s
)
1647 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT
)
1649 unsigned l
= strlen (s
);
1651 /* Append '// ' to the string. */
1652 char *buf
= XNEWVEC (char, l
+ 4);
1653 sprintf (buf
, "// %s", s
);
1657 hsa_insn_comment::~hsa_insn_comment ()
1659 gcc_checking_assert (m_comment
);
1664 /* Constructor of class representing the queue instruction in HSAIL. */
1666 hsa_insn_queue::hsa_insn_queue (int nops
, int opcode
, BrigSegment segment
,
1667 BrigMemoryOrder memory_order
,
1668 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1669 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1670 : hsa_insn_basic (nops
, opcode
, BRIG_TYPE_U64
, arg0
, arg1
, arg2
, arg3
),
1671 m_segment (segment
), m_memory_order (memory_order
)
1675 /* Constructor of class representing the source type instruction in HSAIL. */
1677 hsa_insn_srctype::hsa_insn_srctype (int nops
, BrigOpcode opcode
,
1678 BrigType16_t destt
, BrigType16_t srct
,
1679 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1680 hsa_op_base
*arg2
= NULL
)
1681 : hsa_insn_basic (nops
, opcode
, destt
, arg0
, arg1
, arg2
),
1682 m_source_type (srct
)
1685 /* Constructor of class representing the packed instruction in HSAIL. */
1687 hsa_insn_packed::hsa_insn_packed (int nops
, BrigOpcode opcode
,
1688 BrigType16_t destt
, BrigType16_t srct
,
1689 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1691 : hsa_insn_srctype (nops
, opcode
, destt
, srct
, arg0
, arg1
, arg2
)
1693 m_operand_list
= new hsa_op_operand_list (nops
- 1);
1696 /* Constructor of class representing the convert instruction in HSAIL. */
1698 hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type
*dest
, hsa_op_with_type
*src
)
1699 : hsa_insn_basic (2, BRIG_OPCODE_CVT
, dest
->m_type
, dest
, src
)
1703 /* Constructor of class representing the alloca in HSAIL. */
1705 hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type
*dest
,
1706 hsa_op_with_type
*size
, unsigned alignment
)
1707 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA
, dest
->m_type
, dest
, size
),
1708 m_align (BRIG_ALIGNMENT_8
)
1710 gcc_assert (dest
->m_type
== BRIG_TYPE_U32
);
1712 m_align
= hsa_alignment_encoding (alignment
);
1715 /* Append an instruction INSN into the basic block. */
1718 hsa_bb::append_insn (hsa_insn_basic
*insn
)
1720 gcc_assert (insn
->m_opcode
!= 0 || insn
->operand_count () == 0);
1721 gcc_assert (!insn
->m_bb
);
1724 insn
->m_prev
= m_last_insn
;
1725 insn
->m_next
= NULL
;
1727 m_last_insn
->m_next
= insn
;
1730 m_first_insn
= insn
;
1734 hsa_bb::append_phi (hsa_insn_phi
*hphi
)
1738 hphi
->m_prev
= m_last_phi
;
1739 hphi
->m_next
= NULL
;
1741 m_last_phi
->m_next
= hphi
;
1747 /* Insert HSA instruction NEW_INSN immediately before an existing instruction
1751 hsa_insert_insn_before (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1753 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1755 if (hbb
->m_first_insn
== old_insn
)
1756 hbb
->m_first_insn
= new_insn
;
1757 new_insn
->m_prev
= old_insn
->m_prev
;
1758 new_insn
->m_next
= old_insn
;
1759 if (old_insn
->m_prev
)
1760 old_insn
->m_prev
->m_next
= new_insn
;
1761 old_insn
->m_prev
= new_insn
;
1764 /* Append HSA instruction NEW_INSN immediately after an existing instruction
1768 hsa_append_insn_after (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1770 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1772 if (hbb
->m_last_insn
== old_insn
)
1773 hbb
->m_last_insn
= new_insn
;
1774 new_insn
->m_prev
= old_insn
;
1775 new_insn
->m_next
= old_insn
->m_next
;
1776 if (old_insn
->m_next
)
1777 old_insn
->m_next
->m_prev
= new_insn
;
1778 old_insn
->m_next
= new_insn
;
1781 /* Return a register containing the calculated value of EXP which must be an
1782 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
1783 integer constants as returned by get_inner_reference.
1784 Newly generated HSA instructions will be appended to HBB.
1785 Perform all calculations in ADDRTYPE. */
1787 static hsa_op_with_type
*
1788 gen_address_calculation (tree exp
, hsa_bb
*hbb
, BrigType16_t addrtype
)
1792 if (TREE_CODE (exp
) == NOP_EXPR
)
1793 exp
= TREE_OPERAND (exp
, 0);
1795 switch (TREE_CODE (exp
))
1798 return hsa_cfun
->reg_for_gimple_ssa (exp
)->get_in_type (addrtype
, hbb
);
1802 hsa_op_immed
*imm
= new hsa_op_immed (exp
);
1803 if (addrtype
!= imm
->m_type
)
1804 imm
->m_type
= addrtype
;
1809 opcode
= BRIG_OPCODE_ADD
;
1813 opcode
= BRIG_OPCODE_MUL
;
1820 hsa_op_reg
*res
= new hsa_op_reg (addrtype
);
1821 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, addrtype
);
1822 insn
->set_op (0, res
);
1824 hsa_op_with_type
*op1
= gen_address_calculation (TREE_OPERAND (exp
, 0), hbb
,
1826 hsa_op_with_type
*op2
= gen_address_calculation (TREE_OPERAND (exp
, 1), hbb
,
1828 insn
->set_op (1, op1
);
1829 insn
->set_op (2, op2
);
1831 hbb
->append_insn (insn
);
1835 /* If R1 is NULL, just return R2, otherwise append an instruction adding them
1836 to HBB and return the register holding the result. */
1839 add_addr_regs_if_needed (hsa_op_reg
*r1
, hsa_op_reg
*r2
, hsa_bb
*hbb
)
1841 gcc_checking_assert (r2
);
1845 hsa_op_reg
*res
= new hsa_op_reg (r1
->m_type
);
1846 gcc_assert (!hsa_needs_cvt (r1
->m_type
, r2
->m_type
));
1847 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_ADD
, res
->m_type
);
1848 insn
->set_op (0, res
);
1849 insn
->set_op (1, r1
);
1850 insn
->set_op (2, r2
);
1851 hbb
->append_insn (insn
);
1855 /* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
1856 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
1859 process_mem_base (tree base
, hsa_symbol
**symbol
, BrigType16_t
*addrtype
,
1860 hsa_op_reg
**reg
, offset_int
*offset
, hsa_bb
*hbb
)
1862 if (TREE_CODE (base
) == SSA_NAME
)
1865 hsa_op_with_type
*ssa
1866 = hsa_cfun
->reg_for_gimple_ssa (base
)->get_in_type (*addrtype
, hbb
);
1867 *reg
= dyn_cast
<hsa_op_reg
*> (ssa
);
1869 else if (TREE_CODE (base
) == ADDR_EXPR
)
1871 tree decl
= TREE_OPERAND (base
, 0);
1873 if (!DECL_P (decl
) || TREE_CODE (decl
) == FUNCTION_DECL
)
1875 HSA_SORRY_AT (EXPR_LOCATION (base
),
1876 "support for HSA does not implement a memory reference "
1877 "to a non-declaration type");
1881 gcc_assert (!*symbol
);
1883 *symbol
= get_symbol_for_decl (decl
);
1884 *addrtype
= hsa_get_segment_addr_type ((*symbol
)->m_segment
);
1886 else if (TREE_CODE (base
) == INTEGER_CST
)
1887 *offset
+= wi::to_offset (base
);
1892 /* Forward declaration of a function. */
1895 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
);
1897 /* Generate HSA address operand for a given tree memory reference REF. If
1898 instructions need to be created to calculate the address, they will be added
1899 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
1900 the function assumes that the caller will handle possible
1901 bit-field references. Otherwise if we reference a bit-field, sorry message
1904 static hsa_op_address
*
1905 gen_hsa_addr (tree ref
, hsa_bb
*hbb
, HOST_WIDE_INT
*output_bitsize
= NULL
,
1906 HOST_WIDE_INT
*output_bitpos
= NULL
)
1908 hsa_symbol
*symbol
= NULL
;
1909 hsa_op_reg
*reg
= NULL
;
1910 offset_int offset
= 0;
1912 tree varoffset
= NULL_TREE
;
1913 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
1914 HOST_WIDE_INT bitsize
= 0, bitpos
= 0;
1915 BrigType16_t flat_addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
1917 if (TREE_CODE (ref
) == STRING_CST
)
1919 symbol
= hsa_get_string_cst_symbol (ref
);
1922 else if (TREE_CODE (ref
) == BIT_FIELD_REF
1923 && ((tree_to_uhwi (TREE_OPERAND (ref
, 1)) % BITS_PER_UNIT
) != 0
1924 || (tree_to_uhwi (TREE_OPERAND (ref
, 2)) % BITS_PER_UNIT
) != 0))
1926 HSA_SORRY_ATV (EXPR_LOCATION (origref
),
1927 "support for HSA does not implement "
1928 "bit field references such as %E", ref
);
1932 if (handled_component_p (ref
))
1935 int unsignedp
, volatilep
, preversep
;
1937 ref
= get_inner_reference (ref
, &bitsize
, &bitpos
, &varoffset
, &mode
,
1938 &unsignedp
, &preversep
, &volatilep
);
1941 offset
= wi::rshift (offset
, LOG2_BITS_PER_UNIT
, SIGNED
);
1944 switch (TREE_CODE (ref
))
1948 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
1949 symbol
= hsa_cfun
->create_hsa_temporary (flat_addrtype
);
1950 hsa_op_reg
*r
= new hsa_op_reg (flat_addrtype
);
1951 gen_hsa_addr_insns (ref
, r
, hbb
);
1952 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
1953 r
, new hsa_op_address (symbol
)));
1959 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
1960 symbol
= hsa_cfun
->create_hsa_temporary (flat_addrtype
);
1961 hsa_op_reg
*r
= hsa_cfun
->reg_for_gimple_ssa (ref
);
1963 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
1964 r
, new hsa_op_address (symbol
)));
1972 gcc_assert (!symbol
);
1973 symbol
= get_symbol_for_decl (ref
);
1974 addrtype
= hsa_get_segment_addr_type (symbol
->m_segment
);
1978 process_mem_base (TREE_OPERAND (ref
, 0), &symbol
, &addrtype
, ®
,
1981 if (!integer_zerop (TREE_OPERAND (ref
, 1)))
1982 offset
+= wi::to_offset (TREE_OPERAND (ref
, 1));
1985 case TARGET_MEM_REF
:
1986 process_mem_base (TMR_BASE (ref
), &symbol
, &addrtype
, ®
, &offset
, hbb
);
1987 if (TMR_INDEX (ref
))
1990 hsa_op_base
*idx
= hsa_cfun
->reg_for_gimple_ssa
1991 (TMR_INDEX (ref
))->get_in_type (addrtype
, hbb
);
1992 if (TMR_STEP (ref
) && !integer_onep (TMR_STEP (ref
)))
1994 disp1
= new hsa_op_reg (addrtype
);
1995 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_MUL
,
1998 /* As step must respect addrtype, we overwrite the type
1999 of an immediate value. */
2000 hsa_op_immed
*step
= new hsa_op_immed (TMR_STEP (ref
));
2001 step
->m_type
= addrtype
;
2003 insn
->set_op (0, disp1
);
2004 insn
->set_op (1, idx
);
2005 insn
->set_op (2, step
);
2006 hbb
->append_insn (insn
);
2009 disp1
= as_a
<hsa_op_reg
*> (idx
);
2010 reg
= add_addr_regs_if_needed (reg
, disp1
, hbb
);
2012 if (TMR_INDEX2 (ref
))
2014 if (TREE_CODE (TMR_INDEX2 (ref
)) == SSA_NAME
)
2016 hsa_op_base
*disp2
= hsa_cfun
->reg_for_gimple_ssa
2017 (TMR_INDEX2 (ref
))->get_in_type (addrtype
, hbb
);
2018 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (disp2
),
2021 else if (TREE_CODE (TMR_INDEX2 (ref
)) == INTEGER_CST
)
2022 offset
+= wi::to_offset (TMR_INDEX2 (ref
));
2026 offset
+= wi::to_offset (TMR_OFFSET (ref
));
2029 HSA_SORRY_AT (EXPR_LOCATION (origref
),
2030 "support for HSA does not implement function pointers");
2033 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does "
2034 "not implement memory access to %E", origref
);
2040 if (TREE_CODE (varoffset
) == INTEGER_CST
)
2041 offset
+= wi::to_offset (varoffset
);
2044 hsa_op_base
*off_op
= gen_address_calculation (varoffset
, hbb
,
2046 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (off_op
),
2051 gcc_checking_assert ((symbol
2053 == hsa_get_segment_addr_type (symbol
->m_segment
))
2056 == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
)));
2058 HOST_WIDE_INT hwi_offset
= offset
.to_shwi ();
2060 /* Calculate remaining bitsize offset (if presented). */
2061 bitpos
%= BITS_PER_UNIT
;
2062 /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
2063 is not a reason to think this is a bit-field access. */
2065 && (bitsize
>= BITS_PER_UNIT
)
2066 && !(bitsize
& (bitsize
- 1)))
2069 if ((bitpos
|| bitsize
) && (output_bitpos
== NULL
|| output_bitsize
== NULL
))
2070 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does not "
2071 "implement unhandled bit field reference such as %E", ref
);
2073 if (output_bitsize
!= NULL
&& output_bitpos
!= NULL
)
2075 *output_bitsize
= bitsize
;
2076 *output_bitpos
= bitpos
;
2079 return new hsa_op_address (symbol
, reg
, hwi_offset
);
2082 /* Generate HSA address operand for a given tree memory reference REF. If
2083 instructions need to be created to calculate the address, they will be added
2084 to the end of HBB. OUTPUT_ALIGN is alignment of the created address. */
2086 static hsa_op_address
*
2087 gen_hsa_addr_with_align (tree ref
, hsa_bb
*hbb
, BrigAlignment8_t
*output_align
)
2089 hsa_op_address
*addr
= gen_hsa_addr (ref
, hbb
);
2090 if (addr
->m_reg
|| !addr
->m_symbol
)
2091 *output_align
= hsa_object_alignment (ref
);
2094 /* If the address consists only of a symbol and an offset, we
2095 compute the alignment ourselves to take into account any alignment
2096 promotions we might have done for the HSA symbol representation. */
2097 unsigned align
= hsa_byte_alignment (addr
->m_symbol
->m_align
);
2098 unsigned misalign
= addr
->m_imm_offset
& (align
- 1);
2100 align
= least_bit_hwi (misalign
);
2101 *output_align
= hsa_alignment_encoding (BITS_PER_UNIT
* align
);
2106 /* Generate HSA address for a function call argument of given TYPE.
2107 INDEX is used to generate corresponding name of the arguments.
2108 Special value -1 represents fact that result value is created. */
2110 static hsa_op_address
*
2111 gen_hsa_addr_for_arg (tree tree_type
, int index
)
2113 hsa_symbol
*sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
2115 sym
->m_type
= hsa_type_for_tree_type (tree_type
, &sym
->m_dim
);
2117 if (index
== -1) /* Function result. */
2118 sym
->m_name
= "res";
2119 else /* Function call arguments. */
2122 sym
->m_name_number
= index
;
2125 return new hsa_op_address (sym
);
2128 /* Generate HSA instructions that process all necessary conversions
2129 of an ADDR to flat addressing and place the result into DEST.
2130 Instructions are appended to HBB. */
2133 convert_addr_to_flat_segment (hsa_op_address
*addr
, hsa_op_reg
*dest
,
2136 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_LDA
);
2137 insn
->set_op (1, addr
);
2138 if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
!= BRIG_SEGMENT_GLOBAL
)
2140 /* LDA produces segment-relative address, we need to convert
2141 it to the flat one. */
2143 tmp
= new hsa_op_reg (hsa_get_segment_addr_type
2144 (addr
->m_symbol
->m_segment
));
2146 seg
= new hsa_insn_seg (BRIG_OPCODE_STOF
,
2147 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
2148 tmp
->m_type
, addr
->m_symbol
->m_segment
, dest
,
2151 insn
->set_op (0, tmp
);
2152 insn
->m_type
= tmp
->m_type
;
2153 hbb
->append_insn (insn
);
2154 hbb
->append_insn (seg
);
2158 insn
->set_op (0, dest
);
2159 insn
->m_type
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2160 hbb
->append_insn (insn
);
2164 /* Generate HSA instructions that calculate address of VAL including all
2165 necessary conversions to flat addressing and place the result into DEST.
2166 Instructions are appended to HBB. */
2169 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
)
2171 /* Handle cases like tmp = NULL, where we just emit a move instruction
2173 if (TREE_CODE (val
) == INTEGER_CST
)
2175 hsa_op_immed
*c
= new hsa_op_immed (val
);
2176 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2177 dest
->m_type
, dest
, c
);
2178 hbb
->append_insn (insn
);
2182 hsa_op_address
*addr
;
2184 gcc_assert (dest
->m_type
== hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2185 if (TREE_CODE (val
) == ADDR_EXPR
)
2186 val
= TREE_OPERAND (val
, 0);
2187 addr
= gen_hsa_addr (val
, hbb
);
2189 if (TREE_CODE (val
) == CONST_DECL
2190 && is_gimple_reg_type (TREE_TYPE (val
)))
2192 gcc_assert (addr
->m_symbol
2193 && addr
->m_symbol
->m_segment
== BRIG_SEGMENT_READONLY
);
2194 /* CONST_DECLs are in readonly segment which however does not have
2195 addresses convertible to flat segments. So copy it to a private one
2196 and take address of that. */
2197 BrigType16_t csttype
2198 = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (val
),
2200 hsa_op_reg
*r
= new hsa_op_reg (csttype
);
2201 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD
, csttype
, r
,
2202 new hsa_op_address (addr
->m_symbol
)));
2203 hsa_symbol
*copysym
= hsa_cfun
->create_hsa_temporary (csttype
);
2204 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, csttype
, r
,
2205 new hsa_op_address (copysym
)));
2206 addr
->m_symbol
= copysym
;
2208 else if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
== BRIG_SEGMENT_READONLY
)
2210 HSA_SORRY_ATV (EXPR_LOCATION (val
), "support for HSA does "
2211 "not implement taking addresses of complex "
2212 "CONST_DECLs such as %E", val
);
2217 convert_addr_to_flat_segment (addr
, dest
, hbb
);
2220 /* Return an HSA register or HSA immediate value operand corresponding to
2221 gimple operand OP. */
2223 static hsa_op_with_type
*
2224 hsa_reg_or_immed_for_gimple_op (tree op
, hsa_bb
*hbb
)
2228 if (TREE_CODE (op
) == SSA_NAME
)
2229 tmp
= hsa_cfun
->reg_for_gimple_ssa (op
);
2230 else if (!POINTER_TYPE_P (TREE_TYPE (op
)))
2231 return new hsa_op_immed (op
);
2234 tmp
= new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2235 gen_hsa_addr_insns (op
, tmp
, hbb
);
2240 /* Create a simple movement instruction with register destination DEST and
2241 register or immediate source SRC and append it to the end of HBB. */
2244 hsa_build_append_simple_mov (hsa_op_reg
*dest
, hsa_op_base
*src
, hsa_bb
*hbb
)
2246 /* Moves of packed data between registers need to adhere to the same type
2247 rules like when dealing with memory. */
2248 BrigType16_t tp
= mem_type_for_type (dest
->m_type
);
2249 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, tp
, dest
, src
);
2250 if (hsa_op_reg
*sreg
= dyn_cast
<hsa_op_reg
*> (src
))
2251 gcc_assert (hsa_type_bit_size (dest
->m_type
)
2252 == hsa_type_bit_size (sreg
->m_type
));
2254 gcc_assert (hsa_type_bit_size (dest
->m_type
)
2255 == hsa_type_bit_size (as_a
<hsa_op_immed
*> (src
)->m_type
));
2257 hbb
->append_insn (insn
);
2260 /* Generate HSAIL instructions loading a bit field into register DEST.
2261 VALUE_REG is a register of a SSA name that is used in the bit field
2262 reference. To identify a bit field BITPOS is offset to the loaded memory
2263 and BITSIZE is number of bits of the bit field.
2264 Add instructions to HBB. */
2267 gen_hsa_insns_for_bitfield (hsa_op_reg
*dest
, hsa_op_reg
*value_reg
,
2268 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2271 unsigned type_bitsize
= hsa_type_bit_size (dest
->m_type
);
2272 unsigned left_shift
= type_bitsize
- (bitsize
+ bitpos
);
2273 unsigned right_shift
= left_shift
+ bitpos
;
2277 hsa_op_reg
*value_reg_2
= new hsa_op_reg (dest
->m_type
);
2278 hsa_op_immed
*c
= new hsa_op_immed (left_shift
, BRIG_TYPE_U32
);
2280 hsa_insn_basic
*lshift
2281 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, value_reg_2
->m_type
,
2282 value_reg_2
, value_reg
, c
);
2284 hbb
->append_insn (lshift
);
2286 value_reg
= value_reg_2
;
2291 hsa_op_reg
*value_reg_2
= new hsa_op_reg (dest
->m_type
);
2292 hsa_op_immed
*c
= new hsa_op_immed (right_shift
, BRIG_TYPE_U32
);
2294 hsa_insn_basic
*rshift
2295 = new hsa_insn_basic (3, BRIG_OPCODE_SHR
, value_reg_2
->m_type
,
2296 value_reg_2
, value_reg
, c
);
2298 hbb
->append_insn (rshift
);
2300 value_reg
= value_reg_2
;
2303 hsa_insn_basic
*assignment
2304 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, dest
, value_reg
);
2305 hbb
->append_insn (assignment
);
2309 /* Generate HSAIL instructions loading a bit field into register DEST. ADDR is
2310 prepared memory address which is used to load the bit field. To identify a
2311 bit field BITPOS is offset to the loaded memory and BITSIZE is number of
2312 bits of the bit field. Add instructions to HBB. Load must be performed in
2316 gen_hsa_insns_for_bitfield_load (hsa_op_reg
*dest
, hsa_op_address
*addr
,
2317 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2318 hsa_bb
*hbb
, BrigAlignment8_t align
)
2320 hsa_op_reg
*value_reg
= new hsa_op_reg (dest
->m_type
);
2321 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, dest
->m_type
, value_reg
,
2323 mem
->set_align (align
);
2324 hbb
->append_insn (mem
);
2325 gen_hsa_insns_for_bitfield (dest
, value_reg
, bitsize
, bitpos
, hbb
);
2328 /* Return the alignment of base memory accesses we issue to perform bit-field
2329 memory access REF. */
2331 static BrigAlignment8_t
2332 hsa_bitmemref_alignment (tree ref
)
2334 unsigned HOST_WIDE_INT bit_offset
= 0;
2338 if (TREE_CODE (ref
) == BIT_FIELD_REF
)
2340 if (!tree_fits_uhwi_p (TREE_OPERAND (ref
, 2)))
2341 return BRIG_ALIGNMENT_1
;
2342 bit_offset
+= tree_to_uhwi (TREE_OPERAND (ref
, 2));
2344 else if (TREE_CODE (ref
) == COMPONENT_REF
2345 && DECL_BIT_FIELD (TREE_OPERAND (ref
, 1)))
2346 bit_offset
+= int_bit_position (TREE_OPERAND (ref
, 1));
2349 ref
= TREE_OPERAND (ref
, 0);
2352 unsigned HOST_WIDE_INT bits
= bit_offset
% BITS_PER_UNIT
;
2353 unsigned HOST_WIDE_INT byte_bits
= bit_offset
- bits
;
2354 BrigAlignment8_t base
= hsa_object_alignment (ref
);
2357 return MIN (base
, hsa_alignment_encoding (least_bit_hwi (byte_bits
)));
2360 /* Generate HSAIL instructions loading something into register DEST. RHS is
2361 tree representation of the loaded data, which are loaded as type TYPE. Add
2362 instructions to HBB. */
2365 gen_hsa_insns_for_load (hsa_op_reg
*dest
, tree rhs
, tree type
, hsa_bb
*hbb
)
2367 /* The destination SSA name will give us the type. */
2368 if (TREE_CODE (rhs
) == VIEW_CONVERT_EXPR
)
2369 rhs
= TREE_OPERAND (rhs
, 0);
2371 if (TREE_CODE (rhs
) == SSA_NAME
)
2373 hsa_op_reg
*src
= hsa_cfun
->reg_for_gimple_ssa (rhs
);
2374 hsa_build_append_simple_mov (dest
, src
, hbb
);
2376 else if (is_gimple_min_invariant (rhs
)
2377 || TREE_CODE (rhs
) == ADDR_EXPR
)
2379 if (POINTER_TYPE_P (TREE_TYPE (rhs
)))
2381 if (dest
->m_type
!= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
))
2383 HSA_SORRY_ATV (EXPR_LOCATION (rhs
),
2384 "support for HSA does not implement conversion "
2385 "of %E to the requested non-pointer type.", rhs
);
2389 gen_hsa_addr_insns (rhs
, dest
, hbb
);
2391 else if (TREE_CODE (rhs
) == COMPLEX_CST
)
2393 hsa_op_immed
*real_part
= new hsa_op_immed (TREE_REALPART (rhs
));
2394 hsa_op_immed
*imag_part
= new hsa_op_immed (TREE_IMAGPART (rhs
));
2396 hsa_op_reg
*real_part_reg
2397 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2399 hsa_op_reg
*imag_part_reg
2400 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2403 hsa_build_append_simple_mov (real_part_reg
, real_part
, hbb
);
2404 hsa_build_append_simple_mov (imag_part_reg
, imag_part
, hbb
);
2406 BrigType16_t src_type
= hsa_bittype_for_type (real_part_reg
->m_type
);
2408 hsa_insn_packed
*insn
2409 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
,
2410 src_type
, dest
, real_part_reg
,
2412 hbb
->append_insn (insn
);
2416 hsa_op_immed
*imm
= new hsa_op_immed (rhs
);
2417 hsa_build_append_simple_mov (dest
, imm
, hbb
);
2420 else if (TREE_CODE (rhs
) == REALPART_EXPR
|| TREE_CODE (rhs
) == IMAGPART_EXPR
)
2422 tree pack_type
= TREE_TYPE (TREE_OPERAND (rhs
, 0));
2424 hsa_op_reg
*packed_reg
2425 = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type
, true));
2427 tree complex_rhs
= TREE_OPERAND (rhs
, 0);
2428 gen_hsa_insns_for_load (packed_reg
, complex_rhs
, TREE_TYPE (complex_rhs
),
2431 hsa_op_reg
*real_reg
2432 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2434 hsa_op_reg
*imag_reg
2435 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2437 BrigKind16_t brig_type
= packed_reg
->m_type
;
2438 hsa_insn_packed
*packed
2439 = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND
,
2440 hsa_bittype_for_type (real_reg
->m_type
),
2441 brig_type
, real_reg
, imag_reg
, packed_reg
);
2443 hbb
->append_insn (packed
);
2445 hsa_op_reg
*source
= TREE_CODE (rhs
) == REALPART_EXPR
?
2446 real_reg
: imag_reg
;
2448 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2449 dest
->m_type
, dest
, source
);
2451 hbb
->append_insn (insn
);
2453 else if (TREE_CODE (rhs
) == BIT_FIELD_REF
2454 && TREE_CODE (TREE_OPERAND (rhs
, 0)) == SSA_NAME
)
2456 tree ssa_name
= TREE_OPERAND (rhs
, 0);
2457 HOST_WIDE_INT bitsize
= tree_to_uhwi (TREE_OPERAND (rhs
, 1));
2458 HOST_WIDE_INT bitpos
= tree_to_uhwi (TREE_OPERAND (rhs
, 2));
2460 hsa_op_reg
*imm_value
= hsa_cfun
->reg_for_gimple_ssa (ssa_name
);
2461 gen_hsa_insns_for_bitfield (dest
, imm_value
, bitsize
, bitpos
, hbb
);
2463 else if (DECL_P (rhs
) || TREE_CODE (rhs
) == MEM_REF
2464 || TREE_CODE (rhs
) == TARGET_MEM_REF
2465 || handled_component_p (rhs
))
2467 HOST_WIDE_INT bitsize
, bitpos
;
2469 /* Load from memory. */
2470 hsa_op_address
*addr
;
2471 addr
= gen_hsa_addr (rhs
, hbb
, &bitsize
, &bitpos
);
2473 /* Handle load of a bit field. */
2476 HSA_SORRY_AT (EXPR_LOCATION (rhs
),
2477 "support for HSA does not implement load from a bit "
2478 "field bigger than 64 bits");
2482 if (bitsize
|| bitpos
)
2483 gen_hsa_insns_for_bitfield_load (dest
, addr
, bitsize
, bitpos
, hbb
,
2484 hsa_bitmemref_alignment (rhs
));
2488 /* Not dest->m_type, that's possibly extended. */
2489 mtype
= mem_type_for_type (hsa_type_for_scalar_tree_type (type
,
2491 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
, dest
,
2493 mem
->set_align (hsa_object_alignment (rhs
));
2494 hbb
->append_insn (mem
);
2498 HSA_SORRY_ATV (EXPR_LOCATION (rhs
),
2499 "support for HSA does not implement loading "
2504 /* Return number of bits necessary for representation of a bit field,
2505 starting at BITPOS with size of BITSIZE. */
2508 get_bitfield_size (unsigned bitpos
, unsigned bitsize
)
2510 unsigned s
= bitpos
+ bitsize
;
2511 unsigned sizes
[] = {8, 16, 32, 64};
2513 for (unsigned i
= 0; i
< 4; i
++)
2521 /* Generate HSAIL instructions storing into memory. LHS is the destination of
2522 the store, SRC is the source operand. Add instructions to HBB. */
2525 gen_hsa_insns_for_store (tree lhs
, hsa_op_base
*src
, hsa_bb
*hbb
)
2527 HOST_WIDE_INT bitsize
= 0, bitpos
= 0;
2528 BrigAlignment8_t req_align
;
2530 mtype
= mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
2532 hsa_op_address
*addr
;
2533 addr
= gen_hsa_addr (lhs
, hbb
, &bitsize
, &bitpos
);
2535 /* Handle store to a bit field. */
2538 HSA_SORRY_AT (EXPR_LOCATION (lhs
),
2539 "support for HSA does not implement store to a bit field "
2540 "bigger than 64 bits");
2544 unsigned type_bitsize
= get_bitfield_size (bitpos
, bitsize
);
2546 /* HSAIL does not support MOV insn with 16-bits integers. */
2547 if (type_bitsize
< 32)
2550 if (bitpos
|| (bitsize
&& type_bitsize
!= bitsize
))
2552 unsigned HOST_WIDE_INT mask
= 0;
2553 BrigType16_t mem_type
2554 = get_integer_type_by_bytes (type_bitsize
/ BITS_PER_UNIT
,
2555 !TYPE_UNSIGNED (TREE_TYPE (lhs
)));
2557 for (unsigned i
= 0; i
< type_bitsize
; i
++)
2558 if (i
< bitpos
|| i
>= bitpos
+ bitsize
)
2559 mask
|= ((unsigned HOST_WIDE_INT
)1 << i
);
2561 hsa_op_reg
*value_reg
= new hsa_op_reg (mem_type
);
2563 req_align
= hsa_bitmemref_alignment (lhs
);
2564 /* Load value from memory. */
2565 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mem_type
,
2567 mem
->set_align (req_align
);
2568 hbb
->append_insn (mem
);
2570 /* AND the loaded value with prepared mask. */
2571 hsa_op_reg
*cleared_reg
= new hsa_op_reg (mem_type
);
2574 = get_integer_type_by_bytes (type_bitsize
/ BITS_PER_UNIT
, false);
2575 hsa_op_immed
*c
= new hsa_op_immed (mask
, t
);
2577 hsa_insn_basic
*clearing
2578 = new hsa_insn_basic (3, BRIG_OPCODE_AND
, mem_type
, cleared_reg
,
2580 hbb
->append_insn (clearing
);
2582 /* Shift to left a value that is going to be stored. */
2583 hsa_op_reg
*new_value_reg
= new hsa_op_reg (mem_type
);
2585 hsa_insn_basic
*basic
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, mem_type
,
2586 new_value_reg
, src
);
2587 hbb
->append_insn (basic
);
2591 hsa_op_reg
*shifted_value_reg
= new hsa_op_reg (mem_type
);
2592 c
= new hsa_op_immed (bitpos
, BRIG_TYPE_U32
);
2594 hsa_insn_basic
*basic
2595 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, mem_type
,
2596 shifted_value_reg
, new_value_reg
, c
);
2597 hbb
->append_insn (basic
);
2599 new_value_reg
= shifted_value_reg
;
2602 /* OR the prepared value with prepared chunk loaded from memory. */
2603 hsa_op_reg
*prepared_reg
= new hsa_op_reg (mem_type
);
2604 basic
= new hsa_insn_basic (3, BRIG_OPCODE_OR
, mem_type
, prepared_reg
,
2605 new_value_reg
, cleared_reg
);
2606 hbb
->append_insn (basic
);
2612 req_align
= hsa_object_alignment (lhs
);
2614 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
, addr
);
2615 mem
->set_align (req_align
);
2617 /* The HSAIL verifier has another constraint: if the source is an immediate
2618 then it must match the destination type. If it's a register the low bits
2619 will be used for sub-word stores. We're always allocating new operands so
2620 we can modify the above in place. */
2621 if (hsa_op_immed
*imm
= dyn_cast
<hsa_op_immed
*> (src
))
2623 if (!hsa_type_packed_p (imm
->m_type
))
2624 imm
->m_type
= mem
->m_type
;
2627 /* ...and all vector immediates apparently need to be vectors of
2629 unsigned bs
= hsa_type_bit_size (imm
->m_type
);
2630 gcc_assert (bs
== hsa_type_bit_size (mem
->m_type
));
2634 imm
->m_type
= BRIG_TYPE_U8X4
;
2637 imm
->m_type
= BRIG_TYPE_U8X8
;
2640 imm
->m_type
= BRIG_TYPE_U8X16
;
2648 hbb
->append_insn (mem
);
2651 /* Generate memory copy instructions that are going to be used
2652 for copying a SRC memory to TARGET memory,
2653 represented by pointer in a register. MIN_ALIGN is minimal alignment
2654 of provided HSA addresses. */
2657 gen_hsa_memory_copy (hsa_bb
*hbb
, hsa_op_address
*target
, hsa_op_address
*src
,
2658 unsigned size
, BrigAlignment8_t min_align
)
2660 hsa_op_address
*addr
;
2663 unsigned offset
= 0;
2664 unsigned min_byte_align
= hsa_byte_alignment (min_align
);
2678 if (s
> min_byte_align
)
2681 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2683 hsa_op_reg
*tmp
= new hsa_op_reg (t
);
2684 addr
= new hsa_op_address (src
->m_symbol
, src
->m_reg
,
2685 src
->m_imm_offset
+ offset
);
2686 mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, t
, tmp
, addr
);
2687 hbb
->append_insn (mem
);
2689 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2690 target
->m_imm_offset
+ offset
);
2691 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, tmp
, addr
);
2692 hbb
->append_insn (mem
);
2698 /* Create a memset mask that is created by copying a CONSTANT byte value
2699 to an integer of BYTE_SIZE bytes. */
2701 static unsigned HOST_WIDE_INT
2702 build_memset_value (unsigned HOST_WIDE_INT constant
, unsigned byte_size
)
2707 HOST_WIDE_INT v
= constant
;
2709 for (unsigned i
= 1; i
< byte_size
; i
++)
2710 v
|= constant
<< (8 * i
);
2715 /* Generate memory set instructions that are going to be used
2716 for setting a CONSTANT byte value to TARGET memory of SIZE bytes.
2717 MIN_ALIGN is minimal alignment of provided HSA addresses. */
2720 gen_hsa_memory_set (hsa_bb
*hbb
, hsa_op_address
*target
,
2721 unsigned HOST_WIDE_INT constant
,
2722 unsigned size
, BrigAlignment8_t min_align
)
2724 hsa_op_address
*addr
;
2727 unsigned offset
= 0;
2728 unsigned min_byte_align
= hsa_byte_alignment (min_align
);
2742 if (s
> min_byte_align
)
2745 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2746 target
->m_imm_offset
+ offset
);
2748 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2749 HOST_WIDE_INT c
= build_memset_value (constant
, s
);
2751 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, new hsa_op_immed (c
, t
),
2753 hbb
->append_insn (mem
);
2759 /* Generate HSAIL instructions for a single assignment
2760 of an empty constructor to an ADDR_LHS. Constructor is passed as a
2761 tree RHS and all instructions are appended to HBB. ALIGN is
2762 alignment of the address. */
2765 gen_hsa_ctor_assignment (hsa_op_address
*addr_lhs
, tree rhs
, hsa_bb
*hbb
,
2766 BrigAlignment8_t align
)
2768 if (CONSTRUCTOR_NELTS (rhs
))
2770 HSA_SORRY_AT (EXPR_LOCATION (rhs
),
2771 "support for HSA does not implement load from constructor");
2775 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2776 gen_hsa_memory_set (hbb
, addr_lhs
, 0, size
, align
);
2779 /* Generate HSA instructions for a single assignment of RHS to LHS.
2780 HBB is the basic block they will be appended to. */
2783 gen_hsa_insns_for_single_assignment (tree lhs
, tree rhs
, hsa_bb
*hbb
)
2785 if (TREE_CODE (lhs
) == SSA_NAME
)
2787 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
2788 if (hsa_seen_error ())
2791 gen_hsa_insns_for_load (dest
, rhs
, TREE_TYPE (lhs
), hbb
);
2793 else if (TREE_CODE (rhs
) == SSA_NAME
2794 || (is_gimple_min_invariant (rhs
) && TREE_CODE (rhs
) != STRING_CST
))
2796 /* Store to memory. */
2797 hsa_op_base
*src
= hsa_reg_or_immed_for_gimple_op (rhs
, hbb
);
2798 if (hsa_seen_error ())
2801 gen_hsa_insns_for_store (lhs
, src
, hbb
);
2805 BrigAlignment8_t lhs_align
;
2806 hsa_op_address
*addr_lhs
= gen_hsa_addr_with_align (lhs
, hbb
,
2809 if (TREE_CODE (rhs
) == CONSTRUCTOR
)
2810 gen_hsa_ctor_assignment (addr_lhs
, rhs
, hbb
, lhs_align
);
2813 BrigAlignment8_t rhs_align
;
2814 hsa_op_address
*addr_rhs
= gen_hsa_addr_with_align (rhs
, hbb
,
2817 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2818 gen_hsa_memory_copy (hbb
, addr_lhs
, addr_rhs
, size
,
2819 MIN (lhs_align
, rhs_align
));
2824 /* Prepend before INSN a load from spill symbol of SPILL_REG. Return the
2825 register into which we loaded. If this required another register to convert
2826 from a B1 type, return it in *PTMP2, otherwise store NULL into it. We
2827 assume we are out of SSA so the returned register does not have its
2831 hsa_spill_in (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2833 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2834 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2835 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2837 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, spill_sym
->m_type
,
2839 hsa_insert_insn_before (mem
, insn
);
2842 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2844 hsa_insn_basic
*cvtinsn
;
2846 reg
= new hsa_op_reg (spill_reg
->m_type
);
2848 cvtinsn
= new hsa_insn_cvt (reg
, *ptmp2
);
2849 hsa_insert_insn_before (cvtinsn
, insn
);
2854 /* Append after INSN a store to spill symbol of SPILL_REG. Return the register
2855 from which we stored. If this required another register to convert to a B1
2856 type, return it in *PTMP2, otherwise store NULL into it. We assume we are
2857 out of SSA so the returned register does not have its use updated. */
2860 hsa_spill_out (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2862 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2863 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2864 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2865 hsa_op_reg
*returnreg
;
2869 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2871 hsa_insn_basic
*cvtinsn
;
2872 *ptmp2
= new hsa_op_reg (spill_sym
->m_type
);
2873 reg
->m_type
= spill_reg
->m_type
;
2875 cvtinsn
= new hsa_insn_cvt (*ptmp2
, returnreg
);
2876 hsa_append_insn_after (cvtinsn
, insn
);
2881 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, spill_sym
->m_type
, reg
,
2883 hsa_append_insn_after (mem
, insn
);
2887 /* Generate a comparison instruction that will compare LHS and RHS with
2888 comparison specified by CODE and put result into register DEST. DEST has to
2889 have its type set already but must not have its definition set yet.
2890 Generated instructions will be added to HBB. */
2893 gen_hsa_cmp_insn_from_gimple (enum tree_code code
, tree lhs
, tree rhs
,
2894 hsa_op_reg
*dest
, hsa_bb
*hbb
)
2896 BrigCompareOperation8_t compare
;
2901 compare
= BRIG_COMPARE_LT
;
2904 compare
= BRIG_COMPARE_LE
;
2907 compare
= BRIG_COMPARE_GT
;
2910 compare
= BRIG_COMPARE_GE
;
2913 compare
= BRIG_COMPARE_EQ
;
2916 compare
= BRIG_COMPARE_NE
;
2918 case UNORDERED_EXPR
:
2919 compare
= BRIG_COMPARE_NAN
;
2922 compare
= BRIG_COMPARE_NUM
;
2925 compare
= BRIG_COMPARE_LTU
;
2928 compare
= BRIG_COMPARE_LEU
;
2931 compare
= BRIG_COMPARE_GTU
;
2934 compare
= BRIG_COMPARE_GEU
;
2937 compare
= BRIG_COMPARE_EQU
;
2940 compare
= BRIG_COMPARE_NEU
;
2944 HSA_SORRY_ATV (EXPR_LOCATION (lhs
),
2945 "support for HSA does not implement comparison tree "
2946 "code %s\n", get_tree_code_name (code
));
2950 /* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer)
2951 as a result of comparison. */
2953 BrigType16_t dest_type
= hsa_type_integer_p (dest
->m_type
)
2954 ? (BrigType16_t
) BRIG_TYPE_B1
: dest
->m_type
;
2956 hsa_insn_cmp
*cmp
= new hsa_insn_cmp (compare
, dest_type
);
2957 cmp
->set_op (1, hsa_reg_or_immed_for_gimple_op (lhs
, hbb
));
2958 cmp
->set_op (2, hsa_reg_or_immed_for_gimple_op (rhs
, hbb
));
2960 hbb
->append_insn (cmp
);
2961 cmp
->set_output_in_type (dest
, 0, hbb
);
2964 /* Generate an unary instruction with OPCODE and append it to a basic block
2965 HBB. The instruction uses DEST as a destination and OP1
2966 as a single operand. */
2969 gen_hsa_unary_operation (BrigOpcode opcode
, hsa_op_reg
*dest
,
2970 hsa_op_with_type
*op1
, hsa_bb
*hbb
)
2972 gcc_checking_assert (dest
);
2973 hsa_insn_basic
*insn
;
2975 if (opcode
== BRIG_OPCODE_MOV
&& hsa_needs_cvt (dest
->m_type
, op1
->m_type
))
2976 insn
= new hsa_insn_cvt (dest
, op1
);
2977 else if (opcode
== BRIG_OPCODE_FIRSTBIT
|| opcode
== BRIG_OPCODE_LASTBIT
)
2979 BrigType16_t srctype
= hsa_type_integer_p (op1
->m_type
) ? op1
->m_type
2980 : hsa_unsigned_type_for_type (op1
->m_type
);
2981 insn
= new hsa_insn_srctype (2, opcode
, BRIG_TYPE_U32
, srctype
, NULL
,
2986 insn
= new hsa_insn_basic (2, opcode
, dest
->m_type
, dest
, op1
);
2988 if (opcode
== BRIG_OPCODE_ABS
|| opcode
== BRIG_OPCODE_NEG
)
2990 /* ABS and NEG only exist in _s form :-/ */
2991 if (insn
->m_type
== BRIG_TYPE_U32
)
2992 insn
->m_type
= BRIG_TYPE_S32
;
2993 else if (insn
->m_type
== BRIG_TYPE_U64
)
2994 insn
->m_type
= BRIG_TYPE_S64
;
2998 hbb
->append_insn (insn
);
3000 if (opcode
== BRIG_OPCODE_FIRSTBIT
|| opcode
== BRIG_OPCODE_LASTBIT
)
3001 insn
->set_output_in_type (dest
, 0, hbb
);
3004 /* Generate a binary instruction with OPCODE and append it to a basic block
3005 HBB. The instruction uses DEST as a destination and operands OP1
3009 gen_hsa_binary_operation (int opcode
, hsa_op_reg
*dest
,
3010 hsa_op_base
*op1
, hsa_op_base
*op2
, hsa_bb
*hbb
)
3012 gcc_checking_assert (dest
);
3014 if ((opcode
== BRIG_OPCODE_SHL
|| opcode
== BRIG_OPCODE_SHR
)
3015 && is_a
<hsa_op_immed
*> (op2
))
3017 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3018 i
->set_type (BRIG_TYPE_U32
);
3020 if ((opcode
== BRIG_OPCODE_OR
3021 || opcode
== BRIG_OPCODE_XOR
3022 || opcode
== BRIG_OPCODE_AND
)
3023 && is_a
<hsa_op_immed
*> (op2
))
3025 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3026 i
->set_type (hsa_unsigned_type_for_type (i
->m_type
));
3029 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, dest
->m_type
, dest
,
3031 hbb
->append_insn (insn
);
3034 /* Generate HSA instructions for a single assignment. HBB is the basic block
3035 they will be appended to. */
3038 gen_hsa_insns_for_operation_assignment (gimple
*assign
, hsa_bb
*hbb
)
3040 tree_code code
= gimple_assign_rhs_code (assign
);
3041 gimple_rhs_class rhs_class
= get_gimple_rhs_class (gimple_expr_code (assign
));
3043 tree lhs
= gimple_assign_lhs (assign
);
3044 tree rhs1
= gimple_assign_rhs1 (assign
);
3045 tree rhs2
= gimple_assign_rhs2 (assign
);
3046 tree rhs3
= gimple_assign_rhs3 (assign
);
3054 /* The opcode is changed to BRIG_OPCODE_CVT if BRIG types
3055 needs a conversion. */
3056 opcode
= BRIG_OPCODE_MOV
;
3060 case POINTER_PLUS_EXPR
:
3061 opcode
= BRIG_OPCODE_ADD
;
3064 opcode
= BRIG_OPCODE_SUB
;
3067 opcode
= BRIG_OPCODE_MUL
;
3069 case MULT_HIGHPART_EXPR
:
3070 opcode
= BRIG_OPCODE_MULHI
;
3073 case TRUNC_DIV_EXPR
:
3074 case EXACT_DIV_EXPR
:
3075 opcode
= BRIG_OPCODE_DIV
;
3078 case FLOOR_DIV_EXPR
:
3079 case ROUND_DIV_EXPR
:
3080 HSA_SORRY_AT (gimple_location (assign
),
3081 "support for HSA does not implement CEIL_DIV_EXPR, "
3082 "FLOOR_DIV_EXPR or ROUND_DIV_EXPR");
3084 case TRUNC_MOD_EXPR
:
3085 opcode
= BRIG_OPCODE_REM
;
3088 case FLOOR_MOD_EXPR
:
3089 case ROUND_MOD_EXPR
:
3090 HSA_SORRY_AT (gimple_location (assign
),
3091 "support for HSA does not implement CEIL_MOD_EXPR, "
3092 "FLOOR_MOD_EXPR or ROUND_MOD_EXPR");
3095 opcode
= BRIG_OPCODE_NEG
;
3098 /* There is a native HSA instruction for scalar FMAs but not for vector
3100 if (TREE_CODE (TREE_TYPE (lhs
)) == VECTOR_TYPE
)
3103 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3104 hsa_op_with_type
*op1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3105 hsa_op_with_type
*op2
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3106 hsa_op_with_type
*op3
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3107 hsa_op_reg
*tmp
= new hsa_op_reg (dest
->m_type
);
3108 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp
, op1
, op2
, hbb
);
3109 gen_hsa_binary_operation (BRIG_OPCODE_ADD
, dest
, tmp
, op3
, hbb
);
3112 opcode
= BRIG_OPCODE_MAD
;
3115 opcode
= BRIG_OPCODE_MIN
;
3118 opcode
= BRIG_OPCODE_MAX
;
3121 opcode
= BRIG_OPCODE_ABS
;
3124 opcode
= BRIG_OPCODE_SHL
;
3127 opcode
= BRIG_OPCODE_SHR
;
3132 hsa_insn_basic
*insn
= NULL
;
3133 int code1
= code
== LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3134 int code2
= code
!= LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3135 BrigType16_t btype
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
3138 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3139 hsa_op_reg
*op1
= new hsa_op_reg (btype
);
3140 hsa_op_reg
*op2
= new hsa_op_reg (btype
);
3141 hsa_op_with_type
*shift1
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3143 tree type
= TREE_TYPE (rhs2
);
3144 unsigned HOST_WIDE_INT bitsize
= TREE_INT_CST_LOW (TYPE_SIZE (type
));
3146 hsa_op_with_type
*shift2
= NULL
;
3147 if (TREE_CODE (rhs2
) == INTEGER_CST
)
3148 shift2
= new hsa_op_immed (bitsize
- tree_to_uhwi (rhs2
),
3150 else if (TREE_CODE (rhs2
) == SSA_NAME
)
3152 hsa_op_reg
*s
= hsa_cfun
->reg_for_gimple_ssa (rhs2
);
3153 hsa_op_reg
*d
= new hsa_op_reg (s
->m_type
);
3154 hsa_op_immed
*size_imm
= new hsa_op_immed (bitsize
, BRIG_TYPE_U32
);
3156 insn
= new hsa_insn_basic (3, BRIG_OPCODE_SUB
, d
->m_type
,
3158 hbb
->append_insn (insn
);
3165 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3166 gen_hsa_binary_operation (code1
, op1
, src
, shift1
, hbb
);
3167 gen_hsa_binary_operation (code2
, op2
, src
, shift2
, hbb
);
3168 gen_hsa_binary_operation (BRIG_OPCODE_OR
, dest
, op1
, op2
, hbb
);
3173 opcode
= BRIG_OPCODE_OR
;
3176 opcode
= BRIG_OPCODE_XOR
;
3179 opcode
= BRIG_OPCODE_AND
;
3182 opcode
= BRIG_OPCODE_NOT
;
3184 case FIX_TRUNC_EXPR
:
3186 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3187 hsa_op_with_type
*v
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3189 if (hsa_needs_cvt (dest
->m_type
, v
->m_type
))
3191 hsa_op_reg
*tmp
= new hsa_op_reg (v
->m_type
);
3193 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3194 tmp
->m_type
, tmp
, v
);
3195 hbb
->append_insn (insn
);
3197 hsa_insn_basic
*cvtinsn
= new hsa_insn_cvt (dest
, tmp
);
3198 hbb
->append_insn (cvtinsn
);
3202 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3203 dest
->m_type
, dest
, v
);
3204 hbb
->append_insn (insn
);
3209 opcode
= BRIG_OPCODE_TRUNC
;
3218 case UNORDERED_EXPR
:
3228 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3230 gen_hsa_cmp_insn_from_gimple (code
, rhs1
, rhs2
, dest
, hbb
);
3236 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3237 hsa_op_with_type
*ctrl
= NULL
;
3240 if (CONSTANT_CLASS_P (cond
) || TREE_CODE (cond
) == SSA_NAME
)
3241 ctrl
= hsa_reg_or_immed_for_gimple_op (cond
, hbb
);
3244 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
3246 gen_hsa_cmp_insn_from_gimple (TREE_CODE (cond
),
3247 TREE_OPERAND (cond
, 0),
3248 TREE_OPERAND (cond
, 1),
3254 hsa_op_with_type
*op2
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3255 hsa_op_with_type
*op3
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3257 BrigType16_t utype
= hsa_unsigned_type_for_type (dest
->m_type
);
3258 if (is_a
<hsa_op_immed
*> (op2
))
3259 op2
->m_type
= utype
;
3260 if (is_a
<hsa_op_immed
*> (op3
))
3261 op3
->m_type
= utype
;
3263 hsa_insn_basic
*insn
3264 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
,
3265 hsa_bittype_for_type (dest
->m_type
),
3266 dest
, ctrl
, op2
, op3
);
3268 hbb
->append_insn (insn
);
3274 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3275 hsa_op_with_type
*rhs1_reg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3276 hsa_op_with_type
*rhs2_reg
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3278 if (hsa_seen_error ())
3281 BrigType16_t src_type
= hsa_bittype_for_type (rhs1_reg
->m_type
);
3282 rhs1_reg
= rhs1_reg
->get_in_type (src_type
, hbb
);
3283 rhs2_reg
= rhs2_reg
->get_in_type (src_type
, hbb
);
3285 hsa_insn_packed
*insn
3286 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
, src_type
,
3287 dest
, rhs1_reg
, rhs2_reg
);
3288 hbb
->append_insn (insn
);
3293 /* Implement others as we come across them. */
3294 HSA_SORRY_ATV (gimple_location (assign
),
3295 "support for HSA does not implement operation %s",
3296 get_tree_code_name (code
));
3301 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3303 hsa_op_with_type
*op1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3304 hsa_op_with_type
*op2
= rhs2
!= NULL_TREE
?
3305 hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
) : NULL
;
3307 if (hsa_seen_error ())
3312 case GIMPLE_TERNARY_RHS
:
3314 hsa_op_with_type
*op3
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3315 hsa_insn_basic
*insn
= new hsa_insn_basic (4, opcode
, dest
->m_type
, dest
,
3317 hbb
->append_insn (insn
);
3321 case GIMPLE_BINARY_RHS
:
3322 gen_hsa_binary_operation (opcode
, dest
, op1
, op2
, hbb
);
3325 case GIMPLE_UNARY_RHS
:
3326 gen_hsa_unary_operation (opcode
, dest
, op1
, hbb
);
3333 /* Generate HSA instructions for a given gimple condition statement COND.
3334 Instructions will be appended to HBB, which also needs to be the
3335 corresponding structure to the basic_block of COND. */
3338 gen_hsa_insns_for_cond_stmt (gimple
*cond
, hsa_bb
*hbb
)
3340 hsa_op_reg
*ctrl
= new hsa_op_reg (BRIG_TYPE_B1
);
3343 gen_hsa_cmp_insn_from_gimple (gimple_cond_code (cond
),
3344 gimple_cond_lhs (cond
),
3345 gimple_cond_rhs (cond
),
3348 cbr
= new hsa_insn_cbr (ctrl
);
3349 hbb
->append_insn (cbr
);
3352 /* Maximum number of elements in a jump table for an HSA SBR instruction. */
3354 #define HSA_MAXIMUM_SBR_LABELS 16
3356 /* Return lowest value of a switch S that is handled in a non-default
3360 get_switch_low (gswitch
*s
)
3362 unsigned labels
= gimple_switch_num_labels (s
);
3363 gcc_checking_assert (labels
>= 1);
3365 return CASE_LOW (gimple_switch_label (s
, 1));
3368 /* Return highest value of a switch S that is handled in a non-default
3372 get_switch_high (gswitch
*s
)
3374 unsigned labels
= gimple_switch_num_labels (s
);
3376 /* Compare last label to maximum number of labels. */
3377 tree label
= gimple_switch_label (s
, labels
- 1);
3378 tree low
= CASE_LOW (label
);
3379 tree high
= CASE_HIGH (label
);
3381 return high
!= NULL_TREE
? high
: low
;
3385 get_switch_size (gswitch
*s
)
3387 return int_const_binop (MINUS_EXPR
, get_switch_high (s
), get_switch_low (s
));
3390 /* Generate HSA instructions for a given gimple switch.
3391 Instructions will be appended to HBB. */
3394 gen_hsa_insns_for_switch_stmt (gswitch
*s
, hsa_bb
*hbb
)
3396 gimple_stmt_iterator it
= gsi_for_stmt (s
);
3399 /* Create preambule that verifies that index - lowest_label >= 0. */
3400 edge e
= split_block (hbb
->m_bb
, gsi_stmt (it
));
3401 e
->flags
&= ~EDGE_FALLTHRU
;
3402 e
->flags
|= EDGE_TRUE_VALUE
;
3404 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
3405 tree index_tree
= gimple_switch_index (s
);
3406 tree lowest
= get_switch_low (s
);
3407 tree highest
= get_switch_high (s
);
3409 hsa_op_reg
*index
= hsa_cfun
->reg_for_gimple_ssa (index_tree
);
3411 hsa_op_reg
*cmp1_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3412 hsa_op_immed
*cmp1_immed
= new hsa_op_immed (lowest
);
3413 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_GE
, cmp1_reg
->m_type
,
3414 cmp1_reg
, index
, cmp1_immed
));
3416 hsa_op_reg
*cmp2_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3417 hsa_op_immed
*cmp2_immed
= new hsa_op_immed (highest
);
3418 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_LE
, cmp2_reg
->m_type
,
3419 cmp2_reg
, index
, cmp2_immed
));
3421 hsa_op_reg
*cmp_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3422 hbb
->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_AND
, cmp_reg
->m_type
,
3423 cmp_reg
, cmp1_reg
, cmp2_reg
));
3425 hbb
->append_insn (new hsa_insn_cbr (cmp_reg
));
3427 tree default_label
= gimple_switch_default_label (s
);
3428 basic_block default_label_bb
= label_to_block_fn (func
,
3429 CASE_LABEL (default_label
));
3431 if (!gimple_seq_empty_p (phi_nodes (default_label_bb
)))
3433 default_label_bb
= split_edge (find_edge (e
->dest
, default_label_bb
));
3434 hsa_init_new_bb (default_label_bb
);
3437 make_edge (e
->src
, default_label_bb
, EDGE_FALSE_VALUE
);
3439 hsa_cfun
->m_modified_cfg
= true;
3441 /* Basic block with the SBR instruction. */
3442 hbb
= hsa_init_new_bb (e
->dest
);
3444 hsa_op_reg
*sub_index
= new hsa_op_reg (index
->m_type
);
3445 hbb
->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB
, sub_index
->m_type
,
3447 new hsa_op_immed (lowest
)));
3449 hsa_op_base
*tmp
= sub_index
->get_in_type (BRIG_TYPE_U64
, hbb
);
3450 sub_index
= as_a
<hsa_op_reg
*> (tmp
);
3451 unsigned labels
= gimple_switch_num_labels (s
);
3452 unsigned HOST_WIDE_INT size
= tree_to_uhwi (get_switch_size (s
));
3454 hsa_insn_sbr
*sbr
= new hsa_insn_sbr (sub_index
, size
+ 1);
3456 /* Prepare array with default label destination. */
3457 for (unsigned HOST_WIDE_INT i
= 0; i
<= size
; i
++)
3458 sbr
->m_jump_table
.safe_push (default_label_bb
);
3460 /* Iterate all labels and fill up the jump table. */
3461 for (unsigned i
= 1; i
< labels
; i
++)
3463 tree label
= gimple_switch_label (s
, i
);
3464 basic_block bb
= label_to_block_fn (func
, CASE_LABEL (label
));
3466 unsigned HOST_WIDE_INT sub_low
3467 = tree_to_uhwi (int_const_binop (MINUS_EXPR
, CASE_LOW (label
), lowest
));
3469 unsigned HOST_WIDE_INT sub_high
= sub_low
;
3470 tree high
= CASE_HIGH (label
);
3472 sub_high
= tree_to_uhwi (int_const_binop (MINUS_EXPR
, high
, lowest
));
3474 for (unsigned HOST_WIDE_INT j
= sub_low
; j
<= sub_high
; j
++)
3475 sbr
->m_jump_table
[j
] = bb
;
3478 hbb
->append_insn (sbr
);
3481 /* Verify that the function DECL can be handled by HSA. */
3484 verify_function_arguments (tree decl
)
3486 tree type
= TREE_TYPE (decl
);
3487 if (DECL_STATIC_CHAIN (decl
))
3489 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3490 "HSA does not support nested functions: %qD", decl
);
3493 else if (!TYPE_ARG_TYPES (type
) || stdarg_p (type
))
3495 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3496 "HSA does not support functions with variadic arguments "
3497 "(or unknown return type): %qD", decl
);
3502 /* Return BRIG type for FORMAL_ARG_TYPE. If the formal argument type is NULL,
3503 return ACTUAL_ARG_TYPE. */
3506 get_format_argument_type (tree formal_arg_type
, BrigType16_t actual_arg_type
)
3508 if (formal_arg_type
== NULL
)
3509 return actual_arg_type
;
3511 BrigType16_t decl_type
3512 = hsa_type_for_scalar_tree_type (formal_arg_type
, false);
3513 return mem_type_for_type (decl_type
);
3516 /* Generate HSA instructions for a direct call instruction.
3517 Instructions will be appended to HBB, which also needs to be the
3518 corresponding structure to the basic_block of STMT.
3519 If ASSIGN_LHS is false, do not copy HSA function result argument into the
3520 corresponding HSA representation of the gimple statement LHS. */
3523 gen_hsa_insns_for_direct_call (gimple
*stmt
, hsa_bb
*hbb
,
3524 bool assign_lhs
= true)
3526 tree decl
= gimple_call_fndecl (stmt
);
3527 verify_function_arguments (decl
);
3528 if (hsa_seen_error ())
3531 hsa_insn_call
*call_insn
= new hsa_insn_call (decl
);
3532 hsa_cfun
->m_called_functions
.safe_push (call_insn
->m_called_function
);
3534 /* Argument block start. */
3535 hsa_insn_arg_block
*arg_start
3536 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3537 hbb
->append_insn (arg_start
);
3539 tree parm_type_chain
= TYPE_ARG_TYPES (gimple_call_fntype (stmt
));
3541 /* Preparation of arguments that will be passed to function. */
3542 const unsigned args
= gimple_call_num_args (stmt
);
3543 for (unsigned i
= 0; i
< args
; ++i
)
3545 tree parm
= gimple_call_arg (stmt
, (int)i
);
3546 tree parm_decl_type
= parm_type_chain
!= NULL_TREE
3547 ? TREE_VALUE (parm_type_chain
) : NULL_TREE
;
3548 hsa_op_address
*addr
;
3550 if (AGGREGATE_TYPE_P (TREE_TYPE (parm
)))
3552 addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3553 BrigAlignment8_t align
;
3554 hsa_op_address
*src
= gen_hsa_addr_with_align (parm
, hbb
, &align
);
3555 gen_hsa_memory_copy (hbb
, addr
, src
,
3556 addr
->m_symbol
->total_byte_size (), align
);
3560 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3562 if (parm_decl_type
!= NULL
&& AGGREGATE_TYPE_P (parm_decl_type
))
3564 HSA_SORRY_AT (gimple_location (stmt
),
3565 "support for HSA does not implement an aggregate "
3566 "formal argument in a function call, while actual "
3567 "argument is not an aggregate");
3571 BrigType16_t formal_arg_type
3572 = get_format_argument_type (parm_decl_type
, src
->m_type
);
3573 if (hsa_seen_error ())
3576 if (src
->m_type
!= formal_arg_type
)
3577 src
= src
->get_in_type (formal_arg_type
, hbb
);
3580 = gen_hsa_addr_for_arg (parm_decl_type
!= NULL_TREE
?
3581 parm_decl_type
: TREE_TYPE (parm
), i
);
3582 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, formal_arg_type
,
3585 hbb
->append_insn (mem
);
3588 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3589 if (parm_type_chain
)
3590 parm_type_chain
= TREE_CHAIN (parm_type_chain
);
3593 call_insn
->m_args_code_list
= new hsa_op_code_list (args
);
3594 hbb
->append_insn (call_insn
);
3596 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
3598 tree result
= gimple_call_lhs (stmt
);
3599 hsa_insn_mem
*result_insn
= NULL
;
3600 if (!VOID_TYPE_P (result_type
))
3602 hsa_op_address
*addr
= gen_hsa_addr_for_arg (result_type
, -1);
3604 /* Even if result of a function call is unused, we have to emit
3605 declaration for the result. */
3606 if (result
&& assign_lhs
)
3608 tree lhs_type
= TREE_TYPE (result
);
3610 if (hsa_seen_error ())
3613 if (AGGREGATE_TYPE_P (lhs_type
))
3615 BrigAlignment8_t align
;
3616 hsa_op_address
*result_addr
3617 = gen_hsa_addr_with_align (result
, hbb
, &align
);
3618 gen_hsa_memory_copy (hbb
, result_addr
, addr
,
3619 addr
->m_symbol
->total_byte_size (), align
);
3624 = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type
,
3627 hsa_op_reg
*dst
= hsa_cfun
->reg_for_gimple_ssa (result
);
3628 result_insn
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
, dst
, addr
);
3629 hbb
->append_insn (result_insn
);
3633 call_insn
->m_output_arg
= addr
->m_symbol
;
3634 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3640 HSA_SORRY_AT (gimple_location (stmt
),
3641 "support for HSA does not implement an assignment of "
3642 "return value from a void function");
3646 call_insn
->m_result_code_list
= new hsa_op_code_list (0);
3649 /* Argument block end. */
3650 hsa_insn_arg_block
*arg_end
3651 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3652 hbb
->append_insn (arg_end
);
3655 /* Generate HSA instructions for a direct call of an internal fn.
3656 Instructions will be appended to HBB, which also needs to be the
3657 corresponding structure to the basic_block of STMT. */
3660 gen_hsa_insns_for_call_of_internal_fn (gimple
*stmt
, hsa_bb
*hbb
)
3662 tree lhs
= gimple_call_lhs (stmt
);
3666 tree lhs_type
= TREE_TYPE (lhs
);
3667 tree rhs1
= gimple_call_arg (stmt
, 0);
3668 tree rhs1_type
= TREE_TYPE (rhs1
);
3669 enum internal_fn fn
= gimple_call_internal_fn (stmt
);
3670 hsa_internal_fn
*ifn
3671 = new hsa_internal_fn (fn
, tree_to_uhwi (TYPE_SIZE (rhs1_type
)));
3672 hsa_insn_call
*call_insn
= new hsa_insn_call (ifn
);
3674 gcc_checking_assert (FLOAT_TYPE_P (rhs1_type
));
3676 if (!hsa_emitted_internal_decls
->find (call_insn
->m_called_internal_fn
))
3677 hsa_cfun
->m_called_internal_fns
.safe_push (call_insn
->m_called_internal_fn
);
3679 hsa_insn_arg_block
*arg_start
3680 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3681 hbb
->append_insn (arg_start
);
3683 unsigned num_args
= gimple_call_num_args (stmt
);
3685 /* Function arguments. */
3686 for (unsigned i
= 0; i
< num_args
; i
++)
3688 tree parm
= gimple_call_arg (stmt
, (int)i
);
3689 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3691 hsa_op_address
*addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3692 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, src
->m_type
,
3695 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3696 hbb
->append_insn (mem
);
3699 call_insn
->m_args_code_list
= new hsa_op_code_list (num_args
);
3700 hbb
->append_insn (call_insn
);
3702 /* Assign returned value. */
3703 hsa_op_address
*addr
= gen_hsa_addr_for_arg (lhs_type
, -1);
3705 call_insn
->m_output_arg
= addr
->m_symbol
;
3706 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3708 /* Argument block end. */
3709 hsa_insn_arg_block
*arg_end
3710 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3711 hbb
->append_insn (arg_end
);
3714 /* Generate HSA instructions for a return value instruction.
3715 Instructions will be appended to HBB, which also needs to be the
3716 corresponding structure to the basic_block of STMT. */
3719 gen_hsa_insns_for_return (greturn
*stmt
, hsa_bb
*hbb
)
3721 tree retval
= gimple_return_retval (stmt
);
3724 hsa_op_address
*addr
= new hsa_op_address (hsa_cfun
->m_output_arg
);
3726 if (AGGREGATE_TYPE_P (TREE_TYPE (retval
)))
3728 BrigAlignment8_t align
;
3729 hsa_op_address
*retval_addr
= gen_hsa_addr_with_align (retval
, hbb
,
3731 gen_hsa_memory_copy (hbb
, addr
, retval_addr
,
3732 hsa_cfun
->m_output_arg
->total_byte_size (),
3737 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (retval
),
3739 BrigType16_t mtype
= mem_type_for_type (t
);
3741 /* Store of return value. */
3742 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (retval
, hbb
);
3743 src
= src
->get_in_type (mtype
, hbb
);
3744 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
,
3746 hbb
->append_insn (mem
);
3750 /* HSAIL return instruction emission. */
3751 hsa_insn_basic
*ret
= new hsa_insn_basic (0, BRIG_OPCODE_RET
);
3752 hbb
->append_insn (ret
);
3755 /* Set OP_INDEX-th operand of the instruction to DEST, as the DEST
3756 can have a different type, conversion instructions are possibly
3760 hsa_insn_basic::set_output_in_type (hsa_op_reg
*dest
, unsigned op_index
,
3763 hsa_insn_basic
*insn
;
3764 gcc_checking_assert (op_output_p (op_index
));
3766 if (dest
->m_type
== m_type
)
3768 set_op (op_index
, dest
);
3772 hsa_op_reg
*tmp
= new hsa_op_reg (m_type
);
3773 set_op (op_index
, tmp
);
3775 if (hsa_needs_cvt (dest
->m_type
, m_type
))
3776 insn
= new hsa_insn_cvt (dest
, tmp
);
3778 insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
,
3779 dest
, tmp
->get_in_type (dest
->m_type
, hbb
));
3781 hbb
->append_insn (insn
);
3784 /* Generate instruction OPCODE to query a property of HSA grid along the
3785 given DIMENSION. Store result into DEST and append the instruction to
3789 query_hsa_grid_dim (hsa_op_reg
*dest
, int opcode
, hsa_op_immed
*dimension
,
3792 hsa_insn_basic
*insn
= new hsa_insn_basic (2, opcode
, BRIG_TYPE_U32
, NULL
,
3794 hbb
->append_insn (insn
);
3795 insn
->set_output_in_type (dest
, 0, hbb
);
3798 /* Generate instruction OPCODE to query a property of HSA grid along the given
3799 dimension which is an immediate in first argument of STMT. Store result
3800 into the register corresponding to LHS of STMT and append the instruction to
3804 query_hsa_grid_dim (gimple
*stmt
, int opcode
, hsa_bb
*hbb
)
3806 tree lhs
= gimple_call_lhs (dyn_cast
<gcall
*> (stmt
));
3807 if (lhs
== NULL_TREE
)
3810 tree arg
= gimple_call_arg (stmt
, 0);
3811 unsigned HOST_WIDE_INT dim
= 5;
3812 if (tree_fits_uhwi_p (arg
))
3813 dim
= tree_to_uhwi (arg
);
3816 HSA_SORRY_AT (gimple_location (stmt
),
3817 "HSA grid query dimension must be immediate constant 0, 1 "
3822 hsa_op_immed
*hdim
= new hsa_op_immed (dim
, (BrigKind16_t
) BRIG_TYPE_U32
);
3823 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3824 query_hsa_grid_dim (dest
, opcode
, hdim
, hbb
);
3827 /* Generate instruction OPCODE to query a property of HSA grid that is
3828 independent of any dimension. Store result into the register corresponding
3829 to LHS of STMT and append the instruction to HBB. */
3832 query_hsa_grid_nodim (gimple
*stmt
, BrigOpcode16_t opcode
, hsa_bb
*hbb
)
3834 tree lhs
= gimple_call_lhs (dyn_cast
<gcall
*> (stmt
));
3835 if (lhs
== NULL_TREE
)
3837 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3838 BrigType16_t brig_type
= hsa_unsigned_type_for_type (dest
->m_type
);
3839 hsa_insn_basic
*insn
= new hsa_insn_basic (1, opcode
, brig_type
, dest
);
3840 hbb
->append_insn (insn
);
3843 /* Emit instructions that set hsa_num_threads according to provided VALUE.
3844 Instructions are appended to basic block HBB. */
3847 gen_set_num_threads (tree value
, hsa_bb
*hbb
)
3849 hbb
->append_insn (new hsa_insn_comment ("omp_set_num_threads"));
3850 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (value
, hbb
);
3852 src
= src
->get_in_type (hsa_num_threads
->m_type
, hbb
);
3853 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
3855 hsa_insn_basic
*basic
3856 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
, src
, addr
);
3857 hbb
->append_insn (basic
);
3860 /* Return byte offset of a FIELD_NAME in GOMP_hsa_kernel_dispatch which
3861 is defined in plugin-hsa.c. */
3863 static HOST_WIDE_INT
3864 get_hsa_kernel_dispatch_offset (const char *field_name
)
3866 tree
*hsa_kernel_dispatch_type
= hsa_get_kernel_dispatch_type ();
3867 if (*hsa_kernel_dispatch_type
== NULL
)
3869 /* Collection of information needed for a dispatch of a kernel from a
3870 kernel. Keep in sync with libgomp's plugin-hsa.c. */
3872 *hsa_kernel_dispatch_type
= make_node (RECORD_TYPE
);
3873 tree id_f1
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3874 get_identifier ("queue"), ptr_type_node
);
3875 DECL_CHAIN (id_f1
) = NULL_TREE
;
3876 tree id_f2
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3877 get_identifier ("omp_data_memory"),
3879 DECL_CHAIN (id_f2
) = id_f1
;
3880 tree id_f3
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3881 get_identifier ("kernarg_address"),
3883 DECL_CHAIN (id_f3
) = id_f2
;
3884 tree id_f4
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3885 get_identifier ("object"),
3887 DECL_CHAIN (id_f4
) = id_f3
;
3888 tree id_f5
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3889 get_identifier ("signal"),
3891 DECL_CHAIN (id_f5
) = id_f4
;
3892 tree id_f6
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3893 get_identifier ("private_segment_size"),
3895 DECL_CHAIN (id_f6
) = id_f5
;
3896 tree id_f7
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3897 get_identifier ("group_segment_size"),
3899 DECL_CHAIN (id_f7
) = id_f6
;
3900 tree id_f8
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3901 get_identifier ("kernel_dispatch_count"),
3903 DECL_CHAIN (id_f8
) = id_f7
;
3904 tree id_f9
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3905 get_identifier ("debug"),
3907 DECL_CHAIN (id_f9
) = id_f8
;
3908 tree id_f10
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3909 get_identifier ("omp_level"),
3911 DECL_CHAIN (id_f10
) = id_f9
;
3912 tree id_f11
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3913 get_identifier ("children_dispatches"),
3915 DECL_CHAIN (id_f11
) = id_f10
;
3916 tree id_f12
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3917 get_identifier ("omp_num_threads"),
3919 DECL_CHAIN (id_f12
) = id_f11
;
3922 finish_builtin_struct (*hsa_kernel_dispatch_type
, "__hsa_kernel_dispatch",
3924 TYPE_ARTIFICIAL (*hsa_kernel_dispatch_type
) = 1;
3927 for (tree chain
= TYPE_FIELDS (*hsa_kernel_dispatch_type
);
3928 chain
!= NULL_TREE
; chain
= TREE_CHAIN (chain
))
3929 if (id_equal (DECL_NAME (chain
), field_name
))
3930 return int_byte_position (chain
);
3935 /* Return an HSA register that will contain number of threads for
3936 a future dispatched kernel. Instructions are added to HBB. */
3939 gen_num_threads_for_dispatch (hsa_bb
*hbb
)
3941 /* Step 1) Assign to number of threads:
3942 MIN (HSA_DEFAULT_NUM_THREADS, hsa_num_threads). */
3943 hsa_op_reg
*threads
= new hsa_op_reg (hsa_num_threads
->m_type
);
3944 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
3946 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD
, threads
->m_type
,
3949 hsa_op_immed
*limit
= new hsa_op_immed (HSA_DEFAULT_NUM_THREADS
,
3951 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
3953 = new hsa_insn_cmp (BRIG_COMPARE_LT
, r
->m_type
, r
, threads
, limit
);
3954 hbb
->append_insn (cmp
);
3956 BrigType16_t btype
= hsa_bittype_for_type (threads
->m_type
);
3957 hsa_op_reg
*tmp
= new hsa_op_reg (threads
->m_type
);
3959 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp
, r
,
3962 /* Step 2) If the number is equal to zero,
3963 return shadow->omp_num_threads. */
3964 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
3966 hsa_op_reg
*shadow_thread_count
= new hsa_op_reg (BRIG_TYPE_U32
);
3968 = new hsa_op_address (shadow_reg_ptr
,
3969 get_hsa_kernel_dispatch_offset ("omp_num_threads"));
3970 hsa_insn_basic
*basic
3971 = new hsa_insn_mem (BRIG_OPCODE_LD
, shadow_thread_count
->m_type
,
3972 shadow_thread_count
, addr
);
3973 hbb
->append_insn (basic
);
3975 hsa_op_reg
*tmp2
= new hsa_op_reg (threads
->m_type
);
3976 r
= new hsa_op_reg (BRIG_TYPE_B1
);
3977 hsa_op_immed
*imm
= new hsa_op_immed (0, shadow_thread_count
->m_type
);
3978 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_EQ
, r
->m_type
, r
, tmp
, imm
));
3979 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp2
, r
,
3980 shadow_thread_count
, tmp
));
3982 hsa_op_base
*dest
= tmp2
->get_in_type (BRIG_TYPE_U16
, hbb
);
3984 return as_a
<hsa_op_reg
*> (dest
);
3987 /* Build OPCODE query for all three hsa dimensions, multiply them and store the
3988 result into DEST. */
3991 multiply_grid_dim_characteristics (hsa_op_reg
*dest
, int opcode
, hsa_bb
*hbb
)
3993 hsa_op_reg
*dimx
= new hsa_op_reg (BRIG_TYPE_U32
);
3994 query_hsa_grid_dim (dimx
, opcode
,
3995 new hsa_op_immed (0, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
3996 hsa_op_reg
*dimy
= new hsa_op_reg (BRIG_TYPE_U32
);
3997 query_hsa_grid_dim (dimy
, opcode
,
3998 new hsa_op_immed (1, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
3999 hsa_op_reg
*dimz
= new hsa_op_reg (BRIG_TYPE_U32
);
4000 query_hsa_grid_dim (dimz
, opcode
,
4001 new hsa_op_immed (2, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4002 hsa_op_reg
*tmp
= new hsa_op_reg (dest
->m_type
);
4003 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp
,
4004 dimx
->get_in_type (dest
->m_type
, hbb
),
4005 dimy
->get_in_type (dest
->m_type
, hbb
), hbb
);
4006 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, dest
, tmp
,
4007 dimz
->get_in_type (dest
->m_type
, hbb
), hbb
);
4010 /* Emit instructions that assign number of threads to lhs of gimple STMT.
4011 Instructions are appended to basic block HBB. */
4014 gen_get_num_threads (gimple
*stmt
, hsa_bb
*hbb
)
4016 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4019 hbb
->append_insn (new hsa_insn_comment ("omp_get_num_threads"));
4020 tree lhs
= gimple_call_lhs (stmt
);
4021 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4022 multiply_grid_dim_characteristics (dest
, BRIG_OPCODE_CURRENTWORKGROUPSIZE
,
4026 /* Emit instructions that assign number of teams to lhs of gimple STMT.
4027 Instructions are appended to basic block HBB. */
4030 gen_get_num_teams (gimple
*stmt
, hsa_bb
*hbb
)
4032 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4035 hbb
->append_insn (new hsa_insn_comment ("omp_get_num_teams"));
4036 tree lhs
= gimple_call_lhs (stmt
);
4037 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4038 multiply_grid_dim_characteristics (dest
, BRIG_OPCODE_GRIDGROUPS
, hbb
);
4041 /* Emit instructions that assign a team number to lhs of gimple STMT.
4042 Instructions are appended to basic block HBB. */
4045 gen_get_team_num (gimple
*stmt
, hsa_bb
*hbb
)
4047 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4050 hbb
->append_insn (new hsa_insn_comment ("omp_get_team_num"));
4051 tree lhs
= gimple_call_lhs (stmt
);
4052 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4054 hsa_op_reg
*gnum_x
= new hsa_op_reg (BRIG_TYPE_U32
);
4055 query_hsa_grid_dim (gnum_x
, BRIG_OPCODE_GRIDGROUPS
,
4056 new hsa_op_immed (0, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4057 hsa_op_reg
*gnum_y
= new hsa_op_reg (BRIG_TYPE_U32
);
4058 query_hsa_grid_dim (gnum_y
, BRIG_OPCODE_GRIDGROUPS
,
4059 new hsa_op_immed (1, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4061 hsa_op_reg
*gno_z
= new hsa_op_reg (BRIG_TYPE_U32
);
4062 query_hsa_grid_dim (gno_z
, BRIG_OPCODE_WORKGROUPID
,
4063 new hsa_op_immed (2, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4065 hsa_op_reg
*tmp1
= new hsa_op_reg (dest
->m_type
);
4066 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp1
,
4067 gnum_x
->get_in_type (dest
->m_type
, hbb
),
4068 gnum_y
->get_in_type (dest
->m_type
, hbb
), hbb
);
4069 hsa_op_reg
*tmp2
= new hsa_op_reg (dest
->m_type
);
4070 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp2
, tmp1
,
4071 gno_z
->get_in_type (dest
->m_type
, hbb
), hbb
);
4073 hsa_op_reg
*gno_y
= new hsa_op_reg (BRIG_TYPE_U32
);
4074 query_hsa_grid_dim (gno_y
, BRIG_OPCODE_WORKGROUPID
,
4075 new hsa_op_immed (1, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4076 hsa_op_reg
*tmp3
= new hsa_op_reg (dest
->m_type
);
4077 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp3
,
4078 gnum_x
->get_in_type (dest
->m_type
, hbb
),
4079 gno_y
->get_in_type (dest
->m_type
, hbb
), hbb
);
4080 hsa_op_reg
*tmp4
= new hsa_op_reg (dest
->m_type
);
4081 gen_hsa_binary_operation (BRIG_OPCODE_ADD
, tmp4
, tmp3
, tmp2
, hbb
);
4082 hsa_op_reg
*gno_x
= new hsa_op_reg (BRIG_TYPE_U32
);
4083 query_hsa_grid_dim (gno_x
, BRIG_OPCODE_WORKGROUPID
,
4084 new hsa_op_immed (0, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4085 gen_hsa_binary_operation (BRIG_OPCODE_ADD
, dest
, tmp4
,
4086 gno_x
->get_in_type (dest
->m_type
, hbb
), hbb
);
4089 /* Emit instructions that get levels-var ICV to lhs of gimple STMT.
4090 Instructions are appended to basic block HBB. */
4093 gen_get_level (gimple
*stmt
, hsa_bb
*hbb
)
4095 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4098 hbb
->append_insn (new hsa_insn_comment ("omp_get_level"));
4100 tree lhs
= gimple_call_lhs (stmt
);
4101 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4103 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4104 if (shadow_reg_ptr
== NULL
)
4106 HSA_SORRY_AT (gimple_location (stmt
),
4107 "support for HSA does not implement omp_get_level called "
4108 "from a function not being inlined within a kernel");
4112 hsa_op_address
*addr
4113 = new hsa_op_address (shadow_reg_ptr
,
4114 get_hsa_kernel_dispatch_offset ("omp_level"));
4116 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, BRIG_TYPE_U64
,
4117 (hsa_op_base
*) NULL
, addr
);
4118 hbb
->append_insn (mem
);
4119 mem
->set_output_in_type (dest
, 0, hbb
);
4122 /* Emit instruction that implement omp_get_max_threads of gimple STMT. */
4125 gen_get_max_threads (gimple
*stmt
, hsa_bb
*hbb
)
4127 tree lhs
= gimple_call_lhs (stmt
);
4131 hbb
->append_insn (new hsa_insn_comment ("omp_get_max_threads"));
4133 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4134 hsa_op_with_type
*num_theads_reg
= gen_num_threads_for_dispatch (hbb
)
4135 ->get_in_type (dest
->m_type
, hbb
);
4136 hsa_build_append_simple_mov (dest
, num_theads_reg
, hbb
);
4139 /* Emit instructions that implement alloca builtin gimple STMT.
4140 Instructions are appended to basic block HBB. */
4143 gen_hsa_alloca (gcall
*call
, hsa_bb
*hbb
)
4145 tree lhs
= gimple_call_lhs (call
);
4146 if (lhs
== NULL_TREE
)
4149 built_in_function fn
= DECL_FUNCTION_CODE (gimple_call_fndecl (call
));
4151 gcc_checking_assert (fn
== BUILT_IN_ALLOCA
4152 || fn
== BUILT_IN_ALLOCA_WITH_ALIGN
);
4154 unsigned bit_alignment
= 0;
4156 if (fn
== BUILT_IN_ALLOCA_WITH_ALIGN
)
4158 tree alignment_tree
= gimple_call_arg (call
, 1);
4159 if (TREE_CODE (alignment_tree
) != INTEGER_CST
)
4161 HSA_SORRY_ATV (gimple_location (call
),
4162 "support for HSA does not implement "
4163 "__builtin_alloca_with_align with a non-constant "
4164 "alignment: %E", alignment_tree
);
4167 bit_alignment
= tree_to_uhwi (alignment_tree
);
4170 tree rhs1
= gimple_call_arg (call
, 0);
4171 hsa_op_with_type
*size
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
)
4172 ->get_in_type (BRIG_TYPE_U32
, hbb
);
4173 hsa_op_with_type
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4176 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
));
4177 hsa_insn_alloca
*a
= new hsa_insn_alloca (tmp
, size
, bit_alignment
);
4178 hbb
->append_insn (a
);
4181 = new hsa_insn_seg (BRIG_OPCODE_STOF
,
4182 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
4183 tmp
->m_type
, BRIG_SEGMENT_PRIVATE
, dest
, tmp
);
4184 hbb
->append_insn (seg
);
4187 /* Emit instructions that implement clrsb builtin STMT:
4188 Returns the number of leading redundant sign bits in x, i.e. the number
4189 of bits following the most significant bit that are identical to it.
4190 There are no special cases for 0 or other values.
4191 Instructions are appended to basic block HBB. */
4194 gen_hsa_clrsb (gcall
*call
, hsa_bb
*hbb
)
4196 tree lhs
= gimple_call_lhs (call
);
4197 if (lhs
== NULL_TREE
)
4200 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4201 tree rhs1
= gimple_call_arg (call
, 0);
4202 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4203 BrigType16_t bittype
= hsa_bittype_for_type (arg
->m_type
);
4204 unsigned bitsize
= tree_to_uhwi (TYPE_SIZE (TREE_TYPE (rhs1
)));
4206 /* FIRSTBIT instruction is defined just for 32 and 64-bits wide integers. */
4207 gcc_checking_assert (bitsize
== 32 || bitsize
== 64);
4209 /* Set true to MOST_SIG if the most significant bit is set to one. */
4210 hsa_op_immed
*c
= new hsa_op_immed (1ul << (bitsize
- 1),
4211 hsa_uint_for_bitsize (bitsize
));
4213 hsa_op_reg
*and_reg
= new hsa_op_reg (bittype
);
4214 gen_hsa_binary_operation (BRIG_OPCODE_AND
, and_reg
, arg
, c
, hbb
);
4216 hsa_op_reg
*most_sign
= new hsa_op_reg (BRIG_TYPE_B1
);
4218 = new hsa_insn_cmp (BRIG_COMPARE_EQ
, most_sign
->m_type
, most_sign
,
4220 hbb
->append_insn (cmp
);
4222 /* If the most significant bit is one, negate the input. Otherwise
4223 shift the input value to left by one bit. */
4224 hsa_op_reg
*arg_neg
= new hsa_op_reg (arg
->m_type
);
4225 gen_hsa_unary_operation (BRIG_OPCODE_NEG
, arg_neg
, arg
, hbb
);
4227 hsa_op_reg
*shifted_arg
= new hsa_op_reg (arg
->m_type
);
4228 gen_hsa_binary_operation (BRIG_OPCODE_SHL
, shifted_arg
, arg
,
4229 new hsa_op_immed (1, BRIG_TYPE_U64
), hbb
);
4231 /* Assign the value that can be used for FIRSTBIT instruction according
4232 to the most significant bit. */
4233 hsa_op_reg
*tmp
= new hsa_op_reg (bittype
);
4234 hsa_insn_basic
*cmov
4235 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, bittype
, tmp
, most_sign
,
4236 arg_neg
, shifted_arg
);
4237 hbb
->append_insn (cmov
);
4239 hsa_op_reg
*leading_bits
= new hsa_op_reg (BRIG_TYPE_S32
);
4240 gen_hsa_unary_operation (BRIG_OPCODE_FIRSTBIT
, leading_bits
,
4241 tmp
->get_in_type (hsa_uint_for_bitsize (bitsize
),
4244 /* Set flag if the input value is equal to zero. */
4245 hsa_op_reg
*is_zero
= new hsa_op_reg (BRIG_TYPE_B1
);
4246 cmp
= new hsa_insn_cmp (BRIG_COMPARE_EQ
, is_zero
->m_type
, is_zero
, arg
,
4247 new hsa_op_immed (0, arg
->m_type
));
4248 hbb
->append_insn (cmp
);
4250 /* Return the number of leading bits,
4251 or (bitsize - 1) if the input value is zero. */
4252 cmov
= new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, BRIG_TYPE_B32
, NULL
, is_zero
,
4253 new hsa_op_immed (bitsize
- 1, BRIG_TYPE_U32
),
4254 leading_bits
->get_in_type (BRIG_TYPE_B32
, hbb
));
4255 hbb
->append_insn (cmov
);
4256 cmov
->set_output_in_type (dest
, 0, hbb
);
4259 /* Emit instructions that implement ffs builtin STMT:
4260 Returns one plus the index of the least significant 1-bit of x,
4261 or if x is zero, returns zero.
4262 Instructions are appended to basic block HBB. */
4265 gen_hsa_ffs (gcall
*call
, hsa_bb
*hbb
)
4267 tree lhs
= gimple_call_lhs (call
);
4268 if (lhs
== NULL_TREE
)
4271 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4273 tree rhs1
= gimple_call_arg (call
, 0);
4274 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4276 hsa_op_reg
*tmp
= new hsa_op_reg (BRIG_TYPE_U32
);
4277 hsa_insn_srctype
*insn
= new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT
,
4278 tmp
->m_type
, arg
->m_type
,
4280 hbb
->append_insn (insn
);
4282 hsa_insn_basic
*addition
4283 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
, NULL
, tmp
,
4284 new hsa_op_immed (1, tmp
->m_type
));
4285 hbb
->append_insn (addition
);
4286 addition
->set_output_in_type (dest
, 0, hbb
);
4290 gen_hsa_popcount_to_dest (hsa_op_reg
*dest
, hsa_op_with_type
*arg
, hsa_bb
*hbb
)
4292 gcc_checking_assert (hsa_type_integer_p (arg
->m_type
));
4294 if (hsa_type_bit_size (arg
->m_type
) < 32)
4295 arg
= arg
->get_in_type (BRIG_TYPE_B32
, hbb
);
4297 BrigType16_t srctype
= hsa_bittype_for_type (arg
->m_type
);
4298 if (!hsa_btype_p (arg
->m_type
))
4299 arg
= arg
->get_in_type (srctype
, hbb
);
4301 hsa_insn_srctype
*popcount
4302 = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT
, BRIG_TYPE_U32
,
4303 srctype
, NULL
, arg
);
4304 hbb
->append_insn (popcount
);
4305 popcount
->set_output_in_type (dest
, 0, hbb
);
4308 /* Emit instructions that implement parity builtin STMT:
4309 Returns the parity of x, i.e. the number of 1-bits in x modulo 2.
4310 Instructions are appended to basic block HBB. */
4313 gen_hsa_parity (gcall
*call
, hsa_bb
*hbb
)
4315 tree lhs
= gimple_call_lhs (call
);
4316 if (lhs
== NULL_TREE
)
4319 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4320 tree rhs1
= gimple_call_arg (call
, 0);
4321 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4323 hsa_op_reg
*popcount
= new hsa_op_reg (BRIG_TYPE_U32
);
4324 gen_hsa_popcount_to_dest (popcount
, arg
, hbb
);
4326 hsa_insn_basic
*insn
4327 = new hsa_insn_basic (3, BRIG_OPCODE_REM
, popcount
->m_type
, NULL
, popcount
,
4328 new hsa_op_immed (2, popcount
->m_type
));
4329 hbb
->append_insn (insn
);
4330 insn
->set_output_in_type (dest
, 0, hbb
);
4333 /* Emit instructions that implement popcount builtin STMT.
4334 Instructions are appended to basic block HBB. */
4337 gen_hsa_popcount (gcall
*call
, hsa_bb
*hbb
)
4339 tree lhs
= gimple_call_lhs (call
);
4340 if (lhs
== NULL_TREE
)
4343 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4344 tree rhs1
= gimple_call_arg (call
, 0);
4345 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4347 gen_hsa_popcount_to_dest (dest
, arg
, hbb
);
4350 /* Emit instructions that implement DIVMOD builtin STMT.
4351 Instructions are appended to basic block HBB. */
4354 gen_hsa_divmod (gcall
*call
, hsa_bb
*hbb
)
4356 tree lhs
= gimple_call_lhs (call
);
4357 if (lhs
== NULL_TREE
)
4360 tree rhs0
= gimple_call_arg (call
, 0);
4361 tree rhs1
= gimple_call_arg (call
, 1);
4363 hsa_op_with_type
*arg0
= hsa_reg_or_immed_for_gimple_op (rhs0
, hbb
);
4364 hsa_op_with_type
*arg1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4366 hsa_op_reg
*dest0
= new hsa_op_reg (arg0
->m_type
);
4367 hsa_op_reg
*dest1
= new hsa_op_reg (arg1
->m_type
);
4369 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_DIV
, dest0
->m_type
,
4371 hbb
->append_insn (insn
);
4372 insn
= new hsa_insn_basic (3, BRIG_OPCODE_REM
, dest1
->m_type
, dest1
, arg0
,
4374 hbb
->append_insn (insn
);
4376 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4377 BrigType16_t src_type
= hsa_bittype_for_type (dest0
->m_type
);
4379 insn
= new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
,
4380 src_type
, dest
, dest0
, dest1
);
4381 hbb
->append_insn (insn
);
4384 /* Set VALUE to a shadow kernel debug argument and append a new instruction
4385 to HBB basic block. */
4388 set_debug_value (hsa_bb
*hbb
, hsa_op_with_type
*value
)
4390 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4391 if (shadow_reg_ptr
== NULL
)
4394 hsa_op_address
*addr
4395 = new hsa_op_address (shadow_reg_ptr
,
4396 get_hsa_kernel_dispatch_offset ("debug"));
4397 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, BRIG_TYPE_U64
, value
,
4399 hbb
->append_insn (mem
);
4403 omp_simple_builtin::generate (gimple
*stmt
, hsa_bb
*hbb
)
4407 if (m_warning_message
)
4408 HSA_SORRY_AT (gimple_location (stmt
), m_warning_message
);
4410 HSA_SORRY_ATV (gimple_location (stmt
),
4411 "Support for HSA does not implement calls to %s\n",
4414 else if (m_warning_message
!= NULL
)
4415 warning_at (gimple_location (stmt
), OPT_Whsa
, m_warning_message
);
4417 if (m_return_value
!= NULL
)
4419 tree lhs
= gimple_call_lhs (stmt
);
4423 hbb
->append_insn (new hsa_insn_comment (m_name
));
4425 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4426 hsa_op_with_type
*op
= m_return_value
->get_in_type (dest
->m_type
, hbb
);
4427 hsa_build_append_simple_mov (dest
, op
, hbb
);
4431 /* If STMT is a call of a known library function, generate code to perform
4432 it and return true. */
4435 gen_hsa_insns_for_known_library_call (gimple
*stmt
, hsa_bb
*hbb
)
4437 bool handled
= false;
4438 const char *name
= hsa_get_declaration_name (gimple_call_fndecl (stmt
));
4441 size_t len
= strlen (name
);
4442 if (len
> 0 && name
[len
- 1] == '_')
4444 copy
= XNEWVEC (char, len
+ 1);
4445 strcpy (copy
, name
);
4446 copy
[len
- 1] = '\0';
4450 /* Handle omp_* routines. */
4451 if (strstr (name
, "omp_") == name
)
4453 hsa_init_simple_builtins ();
4454 omp_simple_builtin
*builtin
= omp_simple_builtins
->get (name
);
4457 builtin
->generate (stmt
, hbb
);
4462 if (strcmp (name
, "omp_set_num_threads") == 0)
4463 gen_set_num_threads (gimple_call_arg (stmt
, 0), hbb
);
4464 else if (strcmp (name
, "omp_get_thread_num") == 0)
4466 hbb
->append_insn (new hsa_insn_comment (name
));
4467 query_hsa_grid_nodim (stmt
, BRIG_OPCODE_WORKITEMFLATABSID
, hbb
);
4469 else if (strcmp (name
, "omp_get_num_threads") == 0)
4471 hbb
->append_insn (new hsa_insn_comment (name
));
4472 gen_get_num_threads (stmt
, hbb
);
4474 else if (strcmp (name
, "omp_get_num_teams") == 0)
4475 gen_get_num_teams (stmt
, hbb
);
4476 else if (strcmp (name
, "omp_get_team_num") == 0)
4477 gen_get_team_num (stmt
, hbb
);
4478 else if (strcmp (name
, "omp_get_level") == 0)
4479 gen_get_level (stmt
, hbb
);
4480 else if (strcmp (name
, "omp_get_active_level") == 0)
4481 gen_get_level (stmt
, hbb
);
4482 else if (strcmp (name
, "omp_in_parallel") == 0)
4483 gen_get_level (stmt
, hbb
);
4484 else if (strcmp (name
, "omp_get_max_threads") == 0)
4485 gen_get_max_threads (stmt
, hbb
);
4497 if (strcmp (name
, "__hsa_set_debug_value") == 0)
4500 if (hsa_cfun
->has_shadow_reg_p ())
4502 tree rhs1
= gimple_call_arg (stmt
, 0);
4503 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4505 src
= src
->get_in_type (BRIG_TYPE_U64
, hbb
);
4506 set_debug_value (hbb
, src
);
4515 /* Helper functions to create a single unary HSA operations out of calls to
4516 builtins. OPCODE is the HSA operation to be generated. STMT is a gimple
4517 call to a builtin. HBB is the HSA BB to which the instruction should be
4518 added. Note that nothing will be created if STMT does not have a LHS. */
4521 gen_hsa_unaryop_for_builtin (BrigOpcode opcode
, gimple
*stmt
, hsa_bb
*hbb
)
4523 tree lhs
= gimple_call_lhs (stmt
);
4526 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4527 hsa_op_with_type
*op
4528 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4529 gen_hsa_unary_operation (opcode
, dest
, op
, hbb
);
4532 /* Helper functions to create a call to standard library if LHS of the
4533 STMT is used. HBB is the HSA BB to which the instruction should be
4537 gen_hsa_unaryop_builtin_call (gimple
*stmt
, hsa_bb
*hbb
)
4539 tree lhs
= gimple_call_lhs (stmt
);
4543 if (gimple_call_internal_p (stmt
))
4544 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
4546 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4549 /* Helper functions to create a single unary HSA operations out of calls to
4550 builtins (if unsafe math optimizations are enable). Otherwise, create
4551 a call to standard library function.
4552 OPCODE is the HSA operation to be generated. STMT is a gimple
4553 call to a builtin. HBB is the HSA BB to which the instruction should be
4554 added. Note that nothing will be created if STMT does not have a LHS. */
4557 gen_hsa_unaryop_or_call_for_builtin (BrigOpcode opcode
, gimple
*stmt
,
4560 if (flag_unsafe_math_optimizations
)
4561 gen_hsa_unaryop_for_builtin (opcode
, stmt
, hbb
);
4563 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4566 /* Generate HSA address corresponding to a value VAL (as opposed to a memory
4567 reference tree), for example an SSA_NAME or an ADDR_EXPR. HBB is the HSA BB
4568 to which the instruction should be added. */
4570 static hsa_op_address
*
4571 get_address_from_value (tree val
, hsa_bb
*hbb
)
4573 switch (TREE_CODE (val
))
4577 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4579 = hsa_cfun
->reg_for_gimple_ssa (val
)->get_in_type (addrtype
, hbb
);
4580 return new hsa_op_address (NULL
, as_a
<hsa_op_reg
*> (reg
), 0);
4583 return gen_hsa_addr (TREE_OPERAND (val
, 0), hbb
);
4586 if (tree_fits_shwi_p (val
))
4587 return new hsa_op_address (NULL
, NULL
, tree_to_shwi (val
));
4591 HSA_SORRY_ATV (EXPR_LOCATION (val
),
4592 "support for HSA does not implement memory access to %E",
4594 return new hsa_op_address (NULL
, NULL
, 0);
4598 /* Expand assignment of a result of a string BUILTIN to DST.
4599 Size of the operation is N bytes, where instructions
4600 will be append to HBB. */
4603 expand_lhs_of_string_op (gimple
*stmt
,
4604 unsigned HOST_WIDE_INT n
, hsa_bb
*hbb
,
4605 enum built_in_function builtin
)
4607 /* If LHS is expected, we need to emit a PHI instruction. */
4608 tree lhs
= gimple_call_lhs (stmt
);
4612 hsa_op_reg
*lhs_reg
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4614 hsa_op_with_type
*dst_reg
4615 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4616 hsa_op_with_type
*tmp
;
4620 case BUILT_IN_MEMPCPY
:
4622 tmp
= new hsa_op_reg (dst_reg
->m_type
);
4624 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
,
4626 new hsa_op_immed (n
, dst_reg
->m_type
));
4627 hbb
->append_insn (add
);
4630 case BUILT_IN_MEMCPY
:
4631 case BUILT_IN_MEMSET
:
4638 hbb
->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV
, lhs_reg
->m_type
,
4642 #define HSA_MEMORY_BUILTINS_LIMIT 128
4644 /* Expand a string builtin (from a gimple STMT) in a way that
4645 according to MISALIGNED_FLAG we process either direct emission
4646 (a bunch of memory load and store instructions), or we emit a function call
4647 of a library function (for instance 'memcpy'). Actually, a basic block
4648 for direct emission is just prepared, where caller is responsible
4649 for emission of corresponding instructions.
4650 All instruction are appended to HBB. */
4653 expand_string_operation_builtin (gimple
*stmt
, hsa_bb
*hbb
,
4654 hsa_op_reg
*misaligned_flag
)
4656 edge e
= split_block (hbb
->m_bb
, stmt
);
4657 basic_block condition_bb
= e
->src
;
4658 hbb
->append_insn (new hsa_insn_cbr (misaligned_flag
));
4660 /* Prepare the control flow. */
4661 edge condition_edge
= EDGE_SUCC (condition_bb
, 0);
4662 basic_block call_bb
= split_edge (condition_edge
);
4664 basic_block expanded_bb
= split_edge (EDGE_SUCC (call_bb
, 0));
4665 basic_block cont_bb
= EDGE_SUCC (expanded_bb
, 0)->dest
;
4666 basic_block merge_bb
= split_edge (EDGE_PRED (cont_bb
, 0));
4668 condition_edge
->flags
&= ~EDGE_FALLTHRU
;
4669 condition_edge
->flags
|= EDGE_TRUE_VALUE
;
4670 make_edge (condition_bb
, expanded_bb
, EDGE_FALSE_VALUE
);
4672 redirect_edge_succ (EDGE_SUCC (call_bb
, 0), merge_bb
);
4674 hsa_cfun
->m_modified_cfg
= true;
4676 hsa_init_new_bb (expanded_bb
);
4678 /* Slow path: function call. */
4679 gen_hsa_insns_for_direct_call (stmt
, hsa_init_new_bb (call_bb
), false);
4681 return hsa_bb_for_bb (expanded_bb
);
4684 /* Expand a memory copy BUILTIN (BUILT_IN_MEMCPY, BUILT_IN_MEMPCPY) from
4685 a gimple STMT and store all necessary instruction to HBB basic block. */
4688 expand_memory_copy (gimple
*stmt
, hsa_bb
*hbb
, enum built_in_function builtin
)
4690 tree byte_size
= gimple_call_arg (stmt
, 2);
4692 if (!tree_fits_uhwi_p (byte_size
))
4694 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4698 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
4700 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
4702 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4706 tree dst
= gimple_call_arg (stmt
, 0);
4707 tree src
= gimple_call_arg (stmt
, 1);
4709 hsa_op_address
*dst_addr
= get_address_from_value (dst
, hbb
);
4710 hsa_op_address
*src_addr
= get_address_from_value (src
, hbb
);
4712 /* As gen_hsa_memory_copy relies on memory alignment
4713 greater or equal to 8 bytes, we need to verify the alignment. */
4714 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4715 hsa_op_reg
*src_addr_reg
= new hsa_op_reg (addrtype
);
4716 hsa_op_reg
*dst_addr_reg
= new hsa_op_reg (addrtype
);
4718 convert_addr_to_flat_segment (src_addr
, src_addr_reg
, hbb
);
4719 convert_addr_to_flat_segment (dst_addr
, dst_addr_reg
, hbb
);
4721 /* Process BIT OR for source and destination addresses. */
4722 hsa_op_reg
*or_reg
= new hsa_op_reg (addrtype
);
4723 gen_hsa_binary_operation (BRIG_OPCODE_OR
, or_reg
, src_addr_reg
,
4726 /* Process BIT AND with 0x7 to identify the desired alignment
4728 hsa_op_reg
*masked
= new hsa_op_reg (addrtype
);
4730 gen_hsa_binary_operation (BRIG_OPCODE_AND
, masked
, or_reg
,
4731 new hsa_op_immed (7, addrtype
), hbb
);
4733 hsa_op_reg
*misaligned
= new hsa_op_reg (BRIG_TYPE_B1
);
4734 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE
, misaligned
->m_type
,
4736 new hsa_op_immed (0, masked
->m_type
)));
4738 hsa_bb
*native_impl_bb
4739 = expand_string_operation_builtin (stmt
, hbb
, misaligned
);
4741 gen_hsa_memory_copy (native_impl_bb
, dst_addr
, src_addr
, n
, BRIG_ALIGNMENT_8
);
4743 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb
->m_bb
, 0)->dest
);
4744 expand_lhs_of_string_op (stmt
, n
, merge_bb
, builtin
);
4748 /* Expand a memory set BUILTIN (BUILT_IN_MEMSET, BUILT_IN_BZERO) from
4749 a gimple STMT and store all necessary instruction to HBB basic block.
4750 The operation set N bytes with a CONSTANT value. */
4753 expand_memory_set (gimple
*stmt
, unsigned HOST_WIDE_INT n
,
4754 unsigned HOST_WIDE_INT constant
, hsa_bb
*hbb
,
4755 enum built_in_function builtin
)
4757 tree dst
= gimple_call_arg (stmt
, 0);
4758 hsa_op_address
*dst_addr
= get_address_from_value (dst
, hbb
);
4760 /* As gen_hsa_memory_set relies on memory alignment
4761 greater or equal to 8 bytes, we need to verify the alignment. */
4762 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4763 hsa_op_reg
*dst_addr_reg
= new hsa_op_reg (addrtype
);
4764 convert_addr_to_flat_segment (dst_addr
, dst_addr_reg
, hbb
);
4766 /* Process BIT AND with 0x7 to identify the desired alignment
4768 hsa_op_reg
*masked
= new hsa_op_reg (addrtype
);
4770 gen_hsa_binary_operation (BRIG_OPCODE_AND
, masked
, dst_addr_reg
,
4771 new hsa_op_immed (7, addrtype
), hbb
);
4773 hsa_op_reg
*misaligned
= new hsa_op_reg (BRIG_TYPE_B1
);
4774 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE
, misaligned
->m_type
,
4776 new hsa_op_immed (0, masked
->m_type
)));
4778 hsa_bb
*native_impl_bb
4779 = expand_string_operation_builtin (stmt
, hbb
, misaligned
);
4781 gen_hsa_memory_set (native_impl_bb
, dst_addr
, constant
, n
, BRIG_ALIGNMENT_8
);
4783 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb
->m_bb
, 0)->dest
);
4784 expand_lhs_of_string_op (stmt
, n
, merge_bb
, builtin
);
4787 /* Store into MEMORDER the memory order specified by tree T, which must be an
4788 integer constant representing a C++ memory order. If it isn't, issue an HSA
4789 sorry message using LOC and return true, otherwise return false and store
4790 the name of the requested order to *MNAME. */
4793 hsa_memorder_from_tree (tree t
, BrigMemoryOrder
*memorder
, const char **mname
,
4796 if (!tree_fits_uhwi_p (t
))
4798 HSA_SORRY_ATV (loc
, "support for HSA does not implement memory model %E",
4803 unsigned HOST_WIDE_INT mm
= tree_to_uhwi (t
);
4804 switch (mm
& MEMMODEL_BASE_MASK
)
4806 case MEMMODEL_RELAXED
:
4807 *memorder
= BRIG_MEMORY_ORDER_RELAXED
;
4810 case MEMMODEL_CONSUME
:
4811 /* HSA does not have an equivalent, but we can use the slightly stronger
4813 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
4816 case MEMMODEL_ACQUIRE
:
4817 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
4820 case MEMMODEL_RELEASE
:
4821 *memorder
= BRIG_MEMORY_ORDER_SC_RELEASE
;
4824 case MEMMODEL_ACQ_REL
:
4825 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
4828 case MEMMODEL_SEQ_CST
:
4829 /* Callers implementing a simple load or store need to remove the release
4830 or acquire part respectively. */
4831 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
4836 HSA_SORRY_AT (loc
, "support for HSA does not implement the specified "
4844 /* Helper function to create an HSA atomic operation instruction out of calls
4845 to atomic builtins. RET_ORIG is true if the built-in is the variant that
4846 return s the value before applying operation, and false if it should return
4847 the value after applying the operation (if it returns value at all). ACODE
4848 is the atomic operation code, STMT is a gimple call to a builtin. HBB is
4849 the HSA BB to which the instruction should be added. If SIGNAL is true, the
4850 created operation will work on HSA signals rather than atomic variables. */
4853 gen_hsa_atomic_for_builtin (bool ret_orig
, enum BrigAtomicOperation acode
,
4854 gimple
*stmt
, hsa_bb
*hbb
, bool signal
)
4856 tree lhs
= gimple_call_lhs (stmt
);
4858 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
4859 BrigType16_t hsa_type
= hsa_type_for_scalar_tree_type (type
, false);
4860 BrigType16_t mtype
= mem_type_for_type (hsa_type
);
4861 BrigMemoryOrder memorder
;
4864 if (hsa_memorder_from_tree (gimple_call_arg (stmt
, 2), &memorder
, &mmname
,
4865 gimple_location (stmt
)))
4868 /* Certain atomic insns must have Bx memory types. */
4871 case BRIG_ATOMIC_LD
:
4872 case BRIG_ATOMIC_ST
:
4873 case BRIG_ATOMIC_AND
:
4874 case BRIG_ATOMIC_OR
:
4875 case BRIG_ATOMIC_XOR
:
4876 case BRIG_ATOMIC_EXCH
:
4877 mtype
= hsa_bittype_for_type (mtype
);
4888 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4890 dest
= new hsa_op_reg (hsa_type
);
4891 opcode
= signal
? BRIG_OPCODE_SIGNAL
: BRIG_OPCODE_ATOMIC
;
4897 opcode
= signal
? BRIG_OPCODE_SIGNALNORET
: BRIG_OPCODE_ATOMICNORET
;
4901 if (acode
== BRIG_ATOMIC_ST
)
4903 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
4904 memorder
= BRIG_MEMORY_ORDER_SC_RELEASE
;
4906 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
4907 && memorder
!= BRIG_MEMORY_ORDER_SC_RELEASE
4908 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
4910 HSA_SORRY_ATV (gimple_location (stmt
),
4911 "support for HSA does not implement memory model for "
4912 "ATOMIC_ST: %s", mmname
);
4917 hsa_insn_basic
*atominsn
;
4921 atominsn
= new hsa_insn_signal (nops
, opcode
, acode
, mtype
, memorder
);
4922 tgt
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4926 atominsn
= new hsa_insn_atomic (nops
, opcode
, acode
, mtype
, memorder
);
4927 hsa_op_address
*addr
;
4928 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
4929 if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
== BRIG_SEGMENT_PRIVATE
)
4931 HSA_SORRY_AT (gimple_location (stmt
),
4932 "HSA does not implement atomic operations in private "
4939 hsa_op_base
*op
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1),
4943 atominsn
->set_op (0, dest
);
4944 atominsn
->set_op (1, tgt
);
4945 atominsn
->set_op (2, op
);
4949 atominsn
->set_op (0, tgt
);
4950 atominsn
->set_op (1, op
);
4953 hbb
->append_insn (atominsn
);
4955 /* HSA does not natively support the variants that return the modified value,
4956 so re-do the operation again non-atomically if that is what was
4958 if (lhs
&& !ret_orig
)
4963 case BRIG_ATOMIC_ADD
:
4964 arith
= BRIG_OPCODE_ADD
;
4966 case BRIG_ATOMIC_AND
:
4967 arith
= BRIG_OPCODE_AND
;
4969 case BRIG_ATOMIC_OR
:
4970 arith
= BRIG_OPCODE_OR
;
4972 case BRIG_ATOMIC_SUB
:
4973 arith
= BRIG_OPCODE_SUB
;
4975 case BRIG_ATOMIC_XOR
:
4976 arith
= BRIG_OPCODE_XOR
;
4981 hsa_op_reg
*real_dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4982 gen_hsa_binary_operation (arith
, real_dest
, dest
, op
, hbb
);
4986 /* Generate HSA instructions for an internal fn.
4987 Instructions will be appended to HBB, which also needs to be the
4988 corresponding structure to the basic_block of STMT. */
4991 gen_hsa_insn_for_internal_fn_call (gcall
*stmt
, hsa_bb
*hbb
)
4993 gcc_checking_assert (gimple_call_internal_fn (stmt
));
4994 internal_fn fn
= gimple_call_internal_fn (stmt
);
4996 bool is_float_type_p
= false;
4997 if (gimple_call_lhs (stmt
) != NULL
4998 && TREE_TYPE (gimple_call_lhs (stmt
)) == float_type_node
)
4999 is_float_type_p
= true;
5004 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
5008 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
5012 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
5016 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
5020 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_NRSQRT
, stmt
, hbb
);
5024 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
5029 if (is_float_type_p
)
5030 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
5032 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5038 if (is_float_type_p
)
5039 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
5041 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5048 if (is_float_type_p
)
5049 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
5051 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5058 if (is_float_type_p
)
5059 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
5061 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5066 gen_hsa_clrsb (stmt
, hbb
);
5070 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
5074 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
5078 gen_hsa_ffs (stmt
, hbb
);
5082 gen_hsa_parity (stmt
, hbb
);
5086 gen_hsa_popcount (stmt
, hbb
);
5090 gen_hsa_divmod (stmt
, hbb
);
5103 case IFN_SIGNIFICAND
:
5115 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
5119 HSA_SORRY_ATV (gimple_location (stmt
),
5120 "support for HSA does not implement internal function: %s",
5121 internal_fn_name (fn
));
5126 /* Generate HSA instructions for the given call statement STMT. Instructions
5127 will be appended to HBB. */
5130 gen_hsa_insns_for_call (gimple
*stmt
, hsa_bb
*hbb
)
5132 gcall
*call
= as_a
<gcall
*> (stmt
);
5133 tree lhs
= gimple_call_lhs (stmt
);
5136 if (gimple_call_internal_p (stmt
))
5138 gen_hsa_insn_for_internal_fn_call (call
, hbb
);
5142 if (!gimple_call_builtin_p (stmt
, BUILT_IN_NORMAL
))
5144 tree function_decl
= gimple_call_fndecl (stmt
);
5145 /* Prefetch pass can create type-mismatching prefetch builtin calls which
5146 fail the gimple_call_builtin_p test above. Handle them here. */
5147 if (DECL_BUILT_IN_CLASS (function_decl
)
5148 && DECL_FUNCTION_CODE (function_decl
) == BUILT_IN_PREFETCH
)
5151 if (function_decl
== NULL_TREE
)
5153 HSA_SORRY_AT (gimple_location (stmt
),
5154 "support for HSA does not implement indirect calls");
5158 if (hsa_callable_function_p (function_decl
))
5159 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5160 else if (!gen_hsa_insns_for_known_library_call (stmt
, hbb
))
5161 HSA_SORRY_AT (gimple_location (stmt
),
5162 "HSA supports only calls of functions marked with pragma "
5163 "omp declare target");
5167 tree fndecl
= gimple_call_fndecl (stmt
);
5168 enum built_in_function builtin
= DECL_FUNCTION_CODE (fndecl
);
5172 case BUILT_IN_FABSF
:
5173 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_ABS
, stmt
, hbb
);
5177 case BUILT_IN_CEILF
:
5178 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
5181 case BUILT_IN_FLOOR
:
5182 case BUILT_IN_FLOORF
:
5183 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
5187 case BUILT_IN_RINTF
:
5188 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
5192 case BUILT_IN_SQRTF
:
5193 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
5196 case BUILT_IN_TRUNC
:
5197 case BUILT_IN_TRUNCF
:
5198 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
5205 /* HSAIL does not provide an instruction for double argument type. */
5206 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5210 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
5213 case BUILT_IN_EXP2F
:
5214 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
5217 case BUILT_IN_LOG2F
:
5218 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
5222 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
5225 case BUILT_IN_CLRSB
:
5226 case BUILT_IN_CLRSBL
:
5227 case BUILT_IN_CLRSBLL
:
5228 gen_hsa_clrsb (call
, hbb
);
5233 case BUILT_IN_CLZLL
:
5234 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
5239 case BUILT_IN_CTZLL
:
5240 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
5245 case BUILT_IN_FFSLL
:
5246 gen_hsa_ffs (call
, hbb
);
5249 case BUILT_IN_PARITY
:
5250 case BUILT_IN_PARITYL
:
5251 case BUILT_IN_PARITYLL
:
5252 gen_hsa_parity (call
, hbb
);
5255 case BUILT_IN_POPCOUNT
:
5256 case BUILT_IN_POPCOUNTL
:
5257 case BUILT_IN_POPCOUNTLL
:
5258 gen_hsa_popcount (call
, hbb
);
5261 case BUILT_IN_ATOMIC_LOAD_1
:
5262 case BUILT_IN_ATOMIC_LOAD_2
:
5263 case BUILT_IN_ATOMIC_LOAD_4
:
5264 case BUILT_IN_ATOMIC_LOAD_8
:
5265 case BUILT_IN_ATOMIC_LOAD_16
:
5269 src
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5271 BrigMemoryOrder memorder
;
5273 if (hsa_memorder_from_tree (gimple_call_arg (stmt
, 1), &memorder
,
5274 &mmname
, gimple_location (stmt
)))
5277 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
5278 memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
5280 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
5281 && memorder
!= BRIG_MEMORY_ORDER_SC_ACQUIRE
5282 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
5284 HSA_SORRY_ATV (gimple_location (stmt
),
5285 "support for HSA does not implement "
5286 "memory model for atomic loads: %s", mmname
);
5292 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
5294 mtype
= mem_type_for_type (t
);
5295 mtype
= hsa_bittype_for_type (mtype
);
5296 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5300 mtype
= BRIG_TYPE_B64
;
5301 dest
= new hsa_op_reg (mtype
);
5304 hsa_insn_basic
*atominsn
;
5305 atominsn
= new hsa_insn_atomic (2, BRIG_OPCODE_ATOMIC
, BRIG_ATOMIC_LD
,
5306 mtype
, memorder
, dest
, src
);
5308 hbb
->append_insn (atominsn
);
5312 case BUILT_IN_ATOMIC_EXCHANGE_1
:
5313 case BUILT_IN_ATOMIC_EXCHANGE_2
:
5314 case BUILT_IN_ATOMIC_EXCHANGE_4
:
5315 case BUILT_IN_ATOMIC_EXCHANGE_8
:
5316 case BUILT_IN_ATOMIC_EXCHANGE_16
:
5317 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_EXCH
, stmt
, hbb
, false);
5321 case BUILT_IN_ATOMIC_FETCH_ADD_1
:
5322 case BUILT_IN_ATOMIC_FETCH_ADD_2
:
5323 case BUILT_IN_ATOMIC_FETCH_ADD_4
:
5324 case BUILT_IN_ATOMIC_FETCH_ADD_8
:
5325 case BUILT_IN_ATOMIC_FETCH_ADD_16
:
5326 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ADD
, stmt
, hbb
, false);
5330 case BUILT_IN_ATOMIC_FETCH_SUB_1
:
5331 case BUILT_IN_ATOMIC_FETCH_SUB_2
:
5332 case BUILT_IN_ATOMIC_FETCH_SUB_4
:
5333 case BUILT_IN_ATOMIC_FETCH_SUB_8
:
5334 case BUILT_IN_ATOMIC_FETCH_SUB_16
:
5335 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_SUB
, stmt
, hbb
, false);
5339 case BUILT_IN_ATOMIC_FETCH_AND_1
:
5340 case BUILT_IN_ATOMIC_FETCH_AND_2
:
5341 case BUILT_IN_ATOMIC_FETCH_AND_4
:
5342 case BUILT_IN_ATOMIC_FETCH_AND_8
:
5343 case BUILT_IN_ATOMIC_FETCH_AND_16
:
5344 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_AND
, stmt
, hbb
, false);
5348 case BUILT_IN_ATOMIC_FETCH_XOR_1
:
5349 case BUILT_IN_ATOMIC_FETCH_XOR_2
:
5350 case BUILT_IN_ATOMIC_FETCH_XOR_4
:
5351 case BUILT_IN_ATOMIC_FETCH_XOR_8
:
5352 case BUILT_IN_ATOMIC_FETCH_XOR_16
:
5353 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_XOR
, stmt
, hbb
, false);
5357 case BUILT_IN_ATOMIC_FETCH_OR_1
:
5358 case BUILT_IN_ATOMIC_FETCH_OR_2
:
5359 case BUILT_IN_ATOMIC_FETCH_OR_4
:
5360 case BUILT_IN_ATOMIC_FETCH_OR_8
:
5361 case BUILT_IN_ATOMIC_FETCH_OR_16
:
5362 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_OR
, stmt
, hbb
, false);
5366 case BUILT_IN_ATOMIC_STORE_1
:
5367 case BUILT_IN_ATOMIC_STORE_2
:
5368 case BUILT_IN_ATOMIC_STORE_4
:
5369 case BUILT_IN_ATOMIC_STORE_8
:
5370 case BUILT_IN_ATOMIC_STORE_16
:
5371 /* Since there cannot be any LHS, the first parameter is meaningless. */
5372 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ST
, stmt
, hbb
, false);
5376 case BUILT_IN_ATOMIC_ADD_FETCH_1
:
5377 case BUILT_IN_ATOMIC_ADD_FETCH_2
:
5378 case BUILT_IN_ATOMIC_ADD_FETCH_4
:
5379 case BUILT_IN_ATOMIC_ADD_FETCH_8
:
5380 case BUILT_IN_ATOMIC_ADD_FETCH_16
:
5381 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_ADD
, stmt
, hbb
, false);
5384 case BUILT_IN_ATOMIC_SUB_FETCH_1
:
5385 case BUILT_IN_ATOMIC_SUB_FETCH_2
:
5386 case BUILT_IN_ATOMIC_SUB_FETCH_4
:
5387 case BUILT_IN_ATOMIC_SUB_FETCH_8
:
5388 case BUILT_IN_ATOMIC_SUB_FETCH_16
:
5389 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_SUB
, stmt
, hbb
, false);
5392 case BUILT_IN_ATOMIC_AND_FETCH_1
:
5393 case BUILT_IN_ATOMIC_AND_FETCH_2
:
5394 case BUILT_IN_ATOMIC_AND_FETCH_4
:
5395 case BUILT_IN_ATOMIC_AND_FETCH_8
:
5396 case BUILT_IN_ATOMIC_AND_FETCH_16
:
5397 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_AND
, stmt
, hbb
, false);
5400 case BUILT_IN_ATOMIC_XOR_FETCH_1
:
5401 case BUILT_IN_ATOMIC_XOR_FETCH_2
:
5402 case BUILT_IN_ATOMIC_XOR_FETCH_4
:
5403 case BUILT_IN_ATOMIC_XOR_FETCH_8
:
5404 case BUILT_IN_ATOMIC_XOR_FETCH_16
:
5405 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_XOR
, stmt
, hbb
, false);
5408 case BUILT_IN_ATOMIC_OR_FETCH_1
:
5409 case BUILT_IN_ATOMIC_OR_FETCH_2
:
5410 case BUILT_IN_ATOMIC_OR_FETCH_4
:
5411 case BUILT_IN_ATOMIC_OR_FETCH_8
:
5412 case BUILT_IN_ATOMIC_OR_FETCH_16
:
5413 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_OR
, stmt
, hbb
, false);
5416 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_1
:
5417 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_2
:
5418 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_4
:
5419 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8
:
5420 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_16
:
5422 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
5424 = hsa_bittype_for_type (hsa_type_for_scalar_tree_type (type
, false));
5425 BrigMemoryOrder memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
5426 hsa_insn_basic
*atominsn
;
5428 atominsn
= new hsa_insn_atomic (4, BRIG_OPCODE_ATOMIC
,
5429 BRIG_ATOMIC_CAS
, atype
, memorder
);
5430 tgt
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5433 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5435 dest
= new hsa_op_reg (atype
);
5437 atominsn
->set_op (0, dest
);
5438 atominsn
->set_op (1, tgt
);
5440 hsa_op_with_type
*op
5441 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1), hbb
);
5442 atominsn
->set_op (2, op
);
5443 op
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 2), hbb
);
5444 atominsn
->set_op (3, op
);
5446 hbb
->append_insn (atominsn
);
5450 case BUILT_IN_HSA_WORKGROUPID
:
5451 query_hsa_grid_dim (stmt
, BRIG_OPCODE_WORKGROUPID
, hbb
);
5453 case BUILT_IN_HSA_WORKITEMID
:
5454 query_hsa_grid_dim (stmt
, BRIG_OPCODE_WORKITEMID
, hbb
);
5456 case BUILT_IN_HSA_WORKITEMABSID
:
5457 query_hsa_grid_dim (stmt
, BRIG_OPCODE_WORKITEMABSID
, hbb
);
5459 case BUILT_IN_HSA_GRIDSIZE
:
5460 query_hsa_grid_dim (stmt
, BRIG_OPCODE_GRIDSIZE
, hbb
);
5462 case BUILT_IN_HSA_CURRENTWORKGROUPSIZE
:
5463 query_hsa_grid_dim (stmt
, BRIG_OPCODE_CURRENTWORKGROUPSIZE
, hbb
);
5466 case BUILT_IN_GOMP_BARRIER
:
5467 hbb
->append_insn (new hsa_insn_br (0, BRIG_OPCODE_BARRIER
, BRIG_TYPE_NONE
,
5470 case BUILT_IN_GOMP_PARALLEL
:
5471 HSA_SORRY_AT (gimple_location (stmt
),
5472 "support for HSA does not implement non-gridified "
5473 "OpenMP parallel constructs.");
5476 case BUILT_IN_OMP_GET_THREAD_NUM
:
5478 query_hsa_grid_nodim (stmt
, BRIG_OPCODE_WORKITEMFLATABSID
, hbb
);
5482 case BUILT_IN_OMP_GET_NUM_THREADS
:
5484 gen_get_num_threads (stmt
, hbb
);
5487 case BUILT_IN_GOMP_TEAMS
:
5489 gen_set_num_threads (gimple_call_arg (stmt
, 1), hbb
);
5492 case BUILT_IN_OMP_GET_NUM_TEAMS
:
5494 gen_get_num_teams (stmt
, hbb
);
5497 case BUILT_IN_OMP_GET_TEAM_NUM
:
5499 gen_get_team_num (stmt
, hbb
);
5502 case BUILT_IN_MEMCPY
:
5503 case BUILT_IN_MEMPCPY
:
5505 expand_memory_copy (stmt
, hbb
, builtin
);
5508 case BUILT_IN_MEMSET
:
5510 tree c
= gimple_call_arg (stmt
, 1);
5512 if (TREE_CODE (c
) != INTEGER_CST
)
5514 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5518 tree byte_size
= gimple_call_arg (stmt
, 2);
5520 if (!tree_fits_uhwi_p (byte_size
))
5522 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5526 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
5528 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5530 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5534 unsigned HOST_WIDE_INT constant
5535 = tree_to_uhwi (fold_convert (unsigned_char_type_node
, c
));
5537 expand_memory_set (stmt
, n
, constant
, hbb
, builtin
);
5541 case BUILT_IN_BZERO
:
5543 tree byte_size
= gimple_call_arg (stmt
, 1);
5545 if (!tree_fits_uhwi_p (byte_size
))
5547 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5551 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
5553 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5555 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5559 expand_memory_set (stmt
, n
, 0, hbb
, builtin
);
5563 case BUILT_IN_ALLOCA
:
5564 case BUILT_IN_ALLOCA_WITH_ALIGN
:
5566 gen_hsa_alloca (call
, hbb
);
5569 case BUILT_IN_PREFETCH
:
5573 tree name_tree
= DECL_NAME (fndecl
);
5574 const char *s
= IDENTIFIER_POINTER (name_tree
);
5575 size_t len
= strlen (s
);
5576 if (len
> 4 && (strncmp (s
, "__builtin_GOMP_", 15) == 0))
5577 HSA_SORRY_ATV (gimple_location (stmt
),
5578 "support for HSA does not implement GOMP function %s",
5581 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5587 /* Generate HSA instructions for a given gimple statement. Instructions will be
5591 gen_hsa_insns_for_gimple_stmt (gimple
*stmt
, hsa_bb
*hbb
)
5593 switch (gimple_code (stmt
))
5596 if (gimple_clobber_p (stmt
))
5599 if (gimple_assign_single_p (stmt
))
5601 tree lhs
= gimple_assign_lhs (stmt
);
5602 tree rhs
= gimple_assign_rhs1 (stmt
);
5603 gen_hsa_insns_for_single_assignment (lhs
, rhs
, hbb
);
5606 gen_hsa_insns_for_operation_assignment (stmt
, hbb
);
5609 gen_hsa_insns_for_return (as_a
<greturn
*> (stmt
), hbb
);
5612 gen_hsa_insns_for_cond_stmt (stmt
, hbb
);
5615 gen_hsa_insns_for_call (stmt
, hbb
);
5618 /* ??? HSA supports some debug facilities. */
5622 tree label
= gimple_label_label (as_a
<glabel
*> (stmt
));
5623 if (FORCED_LABEL (label
))
5624 HSA_SORRY_AT (gimple_location (stmt
),
5625 "support for HSA does not implement gimple label with "
5632 hbb
->append_insn (new hsa_insn_basic (0, BRIG_OPCODE_NOP
));
5637 gen_hsa_insns_for_switch_stmt (as_a
<gswitch
*> (stmt
), hbb
);
5641 HSA_SORRY_ATV (gimple_location (stmt
),
5642 "support for HSA does not implement gimple statement %s",
5643 gimple_code_name
[(int) gimple_code (stmt
)]);
5647 /* Generate a HSA PHI from a gimple PHI. */
5650 gen_hsa_phi_from_gimple_phi (gimple
*phi_stmt
, hsa_bb
*hbb
)
5653 unsigned count
= gimple_phi_num_args (phi_stmt
);
5656 = hsa_cfun
->reg_for_gimple_ssa (gimple_phi_result (phi_stmt
));
5657 hphi
= new hsa_insn_phi (count
, dest
);
5658 hphi
->m_bb
= hbb
->m_bb
;
5660 auto_vec
<tree
, 8> aexprs
;
5661 auto_vec
<hsa_op_reg
*, 8> aregs
;
5663 /* Calling split_edge when processing a PHI node messes up with the order of
5664 gimple phi node arguments (it moves the one associated with the edge to
5665 the end). We need to keep the order of edges and arguments of HSA phi
5666 node arguments consistent, so we do all required splitting as the first
5667 step, and in reverse order as to not be affected by the re-orderings. */
5668 for (unsigned j
= count
; j
!= 0; j
--)
5671 tree op
= gimple_phi_arg_def (phi_stmt
, i
);
5672 if (TREE_CODE (op
) != ADDR_EXPR
)
5675 edge e
= gimple_phi_arg_edge (as_a
<gphi
*> (phi_stmt
), i
);
5676 hsa_bb
*hbb_src
= hsa_init_new_bb (split_edge (e
));
5677 hsa_op_address
*addr
= gen_hsa_addr (TREE_OPERAND (op
, 0),
5681 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
5682 hsa_insn_basic
*insn
5683 = new hsa_insn_basic (2, BRIG_OPCODE_LDA
, BRIG_TYPE_U64
,
5685 hbb_src
->append_insn (insn
);
5686 aexprs
.safe_push (op
);
5687 aregs
.safe_push (dest
);
5690 tree lhs
= gimple_phi_result (phi_stmt
);
5691 for (unsigned i
= 0; i
< count
; i
++)
5693 tree op
= gimple_phi_arg_def (phi_stmt
, i
);
5695 if (TREE_CODE (op
) == SSA_NAME
)
5697 hsa_op_reg
*hreg
= hsa_cfun
->reg_for_gimple_ssa (op
);
5698 hphi
->set_op (i
, hreg
);
5702 gcc_assert (is_gimple_min_invariant (op
));
5703 tree t
= TREE_TYPE (op
);
5704 if (!POINTER_TYPE_P (t
)
5705 || (TREE_CODE (op
) == STRING_CST
5706 && TREE_CODE (TREE_TYPE (t
)) == INTEGER_TYPE
))
5707 hphi
->set_op (i
, new hsa_op_immed (op
));
5708 else if (POINTER_TYPE_P (TREE_TYPE (lhs
))
5709 && TREE_CODE (op
) == INTEGER_CST
)
5711 /* Handle assignment of NULL value to a pointer type. */
5712 hphi
->set_op (i
, new hsa_op_immed (op
));
5714 else if (TREE_CODE (op
) == ADDR_EXPR
)
5716 hsa_op_reg
*dest
= NULL
;
5717 for (unsigned a_idx
= 0; a_idx
< aexprs
.length (); a_idx
++)
5718 if (aexprs
[a_idx
] == op
)
5720 dest
= aregs
[a_idx
];
5724 hphi
->set_op (i
, dest
);
5728 HSA_SORRY_AT (gimple_location (phi_stmt
),
5729 "support for HSA does not handle PHI nodes with "
5730 "constant address operands");
5736 hbb
->append_phi (hphi
);
5739 /* Constructor of class containing HSA-specific information about a basic
5740 block. CFG_BB is the CFG BB this HSA BB is associated with. IDX is the new
5741 index of this BB (so that the constructor does not attempt to use
5742 hsa_cfun during its construction). */
5744 hsa_bb::hsa_bb (basic_block cfg_bb
, int idx
)
5745 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5746 m_last_phi (NULL
), m_index (idx
)
5748 gcc_assert (!cfg_bb
->aux
);
5752 /* Constructor of class containing HSA-specific information about a basic
5753 block. CFG_BB is the CFG BB this HSA BB is associated with. */
5755 hsa_bb::hsa_bb (basic_block cfg_bb
)
5756 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5757 m_last_phi (NULL
), m_index (hsa_cfun
->m_hbb_count
++)
5759 gcc_assert (!cfg_bb
->aux
);
5763 /* Create and initialize and return a new hsa_bb structure for a given CFG
5767 hsa_init_new_bb (basic_block bb
)
5769 void *m
= obstack_alloc (&hsa_obstack
, sizeof (hsa_bb
));
5770 return new (m
) hsa_bb (bb
);
5773 /* Initialize OMP in an HSA basic block PROLOGUE. */
5776 init_prologue (void)
5778 if (!hsa_cfun
->m_kern_p
)
5781 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5783 /* Create a magic number that is going to be printed by libgomp. */
5784 unsigned index
= hsa_get_number_decl_kernel_mappings ();
5786 /* Emit store to debug argument. */
5787 if (PARAM_VALUE (PARAM_HSA_GEN_DEBUG_STORES
) > 0)
5788 set_debug_value (prologue
, new hsa_op_immed (1000 + index
, BRIG_TYPE_U64
));
5791 /* Initialize hsa_num_threads to a default value. */
5794 init_hsa_num_threads (void)
5796 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5798 /* Save the default value to private variable hsa_num_threads. */
5799 hsa_insn_basic
*basic
5800 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
,
5801 new hsa_op_immed (0, hsa_num_threads
->m_type
),
5802 new hsa_op_address (hsa_num_threads
));
5803 prologue
->append_insn (basic
);
5806 /* Go over gimple representation and generate our internal HSA one. */
5809 gen_body_from_gimple ()
5813 /* Verify CFG for complex edges we are unable to handle. */
5817 FOR_EACH_BB_FN (bb
, cfun
)
5819 FOR_EACH_EDGE (e
, ei
, bb
->succs
)
5821 /* Verify all unsupported flags for edges that point
5822 to the same basic block. */
5823 if (e
->flags
& EDGE_EH
)
5825 HSA_SORRY_AT (UNKNOWN_LOCATION
,
5826 "support for HSA does not implement exception "
5833 FOR_EACH_BB_FN (bb
, cfun
)
5835 gimple_stmt_iterator gsi
;
5836 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5840 hbb
= hsa_init_new_bb (bb
);
5842 for (gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5844 gen_hsa_insns_for_gimple_stmt (gsi_stmt (gsi
), hbb
);
5845 if (hsa_seen_error ())
5850 FOR_EACH_BB_FN (bb
, cfun
)
5852 gimple_stmt_iterator gsi
;
5853 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5854 gcc_assert (hbb
!= NULL
);
5856 for (gsi
= gsi_start_phis (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5857 if (!virtual_operand_p (gimple_phi_result (gsi_stmt (gsi
))))
5858 gen_hsa_phi_from_gimple_phi (gsi_stmt (gsi
), hbb
);
5861 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
5863 fprintf (dump_file
, "------- Generated SSA form -------\n");
5864 dump_hsa_cfun (dump_file
);
5869 gen_function_decl_parameters (hsa_function_representation
*f
,
5875 for (parm
= TYPE_ARG_TYPES (TREE_TYPE (decl
)), i
= 0;
5877 parm
= TREE_CHAIN (parm
), i
++)
5879 /* Result type if last in the tree list. */
5880 if (TREE_CHAIN (parm
) == NULL
)
5883 tree v
= TREE_VALUE (parm
);
5885 hsa_symbol
*arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5887 arg
->m_type
= hsa_type_for_tree_type (v
, &arg
->m_dim
);
5888 arg
->m_name_number
= i
;
5890 f
->m_input_args
.safe_push (arg
);
5893 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
5894 if (!VOID_TYPE_P (result_type
))
5896 f
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5898 f
->m_output_arg
->m_type
5899 = hsa_type_for_tree_type (result_type
, &f
->m_output_arg
->m_dim
);
5900 f
->m_output_arg
->m_name
= "res";
5904 /* Generate the vector of parameters of the HSA representation of the current
5905 function. This also includes the output parameter representing the
5909 gen_function_def_parameters ()
5913 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5915 for (parm
= DECL_ARGUMENTS (cfun
->decl
); parm
;
5916 parm
= DECL_CHAIN (parm
))
5918 struct hsa_symbol
**slot
;
5921 = new hsa_symbol (BRIG_TYPE_NONE
, hsa_cfun
->m_kern_p
5922 ? BRIG_SEGMENT_KERNARG
: BRIG_SEGMENT_ARG
,
5923 BRIG_LINKAGE_FUNCTION
);
5924 arg
->fillup_for_decl (parm
);
5926 hsa_cfun
->m_input_args
.safe_push (arg
);
5928 if (hsa_seen_error ())
5931 arg
->m_name
= hsa_get_declaration_name (parm
);
5933 /* Copy all input arguments and create corresponding private symbols
5935 hsa_symbol
*private_arg
;
5936 hsa_op_address
*parm_addr
= new hsa_op_address (arg
);
5938 if (TREE_ADDRESSABLE (parm
)
5939 || (!is_gimple_reg (parm
) && !TREE_READONLY (parm
)))
5941 private_arg
= hsa_cfun
->create_hsa_temporary (arg
->m_type
);
5942 private_arg
->fillup_for_decl (parm
);
5944 BrigAlignment8_t align
= MIN (arg
->m_align
, private_arg
->m_align
);
5946 hsa_op_address
*private_arg_addr
= new hsa_op_address (private_arg
);
5947 gen_hsa_memory_copy (prologue
, private_arg_addr
, parm_addr
,
5948 arg
->total_byte_size (), align
);
5953 slot
= hsa_cfun
->m_local_symbols
->find_slot (private_arg
, INSERT
);
5954 gcc_assert (!*slot
);
5955 *slot
= private_arg
;
5957 if (is_gimple_reg (parm
))
5959 tree ddef
= ssa_default_def (cfun
, parm
);
5960 if (ddef
&& !has_zero_uses (ddef
))
5962 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (ddef
),
5964 BrigType16_t mtype
= mem_type_for_type (t
);
5965 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (ddef
);
5966 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
,
5968 gcc_assert (!parm_addr
->m_reg
);
5969 prologue
->append_insn (mem
);
5974 if (!VOID_TYPE_P (TREE_TYPE (TREE_TYPE (cfun
->decl
))))
5976 struct hsa_symbol
**slot
;
5978 hsa_cfun
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5979 BRIG_LINKAGE_FUNCTION
);
5980 hsa_cfun
->m_output_arg
->fillup_for_decl (DECL_RESULT (cfun
->decl
));
5982 if (hsa_seen_error ())
5985 hsa_cfun
->m_output_arg
->m_name
= "res";
5986 slot
= hsa_cfun
->m_local_symbols
->find_slot (hsa_cfun
->m_output_arg
,
5988 gcc_assert (!*slot
);
5989 *slot
= hsa_cfun
->m_output_arg
;
5993 /* Generate function representation that corresponds to
5994 a function declaration. */
5996 hsa_function_representation
*
5997 hsa_generate_function_declaration (tree decl
)
5999 hsa_function_representation
*fun
6000 = new hsa_function_representation (decl
, false, 0);
6002 fun
->m_declaration_p
= true;
6003 fun
->m_name
= get_brig_function_name (decl
);
6004 gen_function_decl_parameters (fun
, decl
);
6010 /* Generate function representation that corresponds to
6013 hsa_function_representation
*
6014 hsa_generate_internal_fn_decl (hsa_internal_fn
*fn
)
6016 hsa_function_representation
*fun
= new hsa_function_representation (fn
);
6018 fun
->m_name
= fn
->name ();
6020 for (unsigned i
= 0; i
< fn
->get_arity (); i
++)
6023 = new hsa_symbol (fn
->get_argument_type (i
), BRIG_SEGMENT_ARG
,
6025 arg
->m_name_number
= i
;
6026 fun
->m_input_args
.safe_push (arg
);
6029 fun
->m_output_arg
= new hsa_symbol (fn
->get_argument_type (-1),
6030 BRIG_SEGMENT_ARG
, BRIG_LINKAGE_NONE
);
6031 fun
->m_output_arg
->m_name
= "res";
6036 /* Return true if switch statement S can be transformed
6037 to a SBR instruction in HSAIL. */
6040 transformable_switch_to_sbr_p (gswitch
*s
)
6042 /* Identify if a switch statement can be transformed to
6043 SBR instruction, like:
6045 sbr_u32 $s1 [@label1, @label2, @label3];
6048 tree size
= get_switch_size (s
);
6049 if (!tree_fits_uhwi_p (size
))
6052 if (tree_to_uhwi (size
) > HSA_MAXIMUM_SBR_LABELS
)
6058 /* Structure hold connection between PHI nodes and immediate
6059 values hold by there nodes. */
6061 struct phi_definition
6063 phi_definition (unsigned phi_i
, unsigned label_i
, tree imm
):
6064 phi_index (phi_i
), label_index (label_i
), phi_value (imm
)
6068 unsigned label_index
;
6072 /* Sum slice of a vector V, starting from index START and ending
6073 at the index END - 1. */
6075 template <typename T
>
6077 T
sum_slice (const auto_vec
<T
> &v
, unsigned start
, unsigned end
,
6082 for (unsigned i
= start
; i
< end
; i
++)
6088 /* Function transforms GIMPLE SWITCH statements to a series of IF statements.
6089 Let's assume following example:
6103 The transformation encompasses following steps:
6104 1) all immediate values used by edges coming from the switch basic block
6106 2) all these edges are removed
6107 3) the switch statement (in L0) is replaced by:
6113 4) newly created basic block Lx' is used for generation of
6115 5) else branch of the last condition goes to LD
6116 6) fix all immediate values in PHI nodes that were propagated though
6117 edges that were removed in step 2
6119 Note: if a case is made by a range C1..C2, then process
6120 following transformation:
6122 switch_cond_op1 = C1 <= index;
6123 switch_cond_op2 = index <= C2;
6124 switch_cond_and = switch_cond_op1 & switch_cond_op2;
6125 if (switch_cond_and != 0)
6133 convert_switch_statements (void)
6135 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
6138 bool modified_cfg
= false;
6140 FOR_EACH_BB_FN (bb
, func
)
6142 gimple_stmt_iterator gsi
= gsi_last_bb (bb
);
6143 if (gsi_end_p (gsi
))
6146 gimple
*stmt
= gsi_stmt (gsi
);
6148 if (gimple_code (stmt
) == GIMPLE_SWITCH
)
6150 gswitch
*s
= as_a
<gswitch
*> (stmt
);
6152 /* If the switch can utilize SBR insn, skip the statement. */
6153 if (transformable_switch_to_sbr_p (s
))
6156 modified_cfg
= true;
6158 unsigned labels
= gimple_switch_num_labels (s
);
6159 tree index
= gimple_switch_index (s
);
6160 tree index_type
= TREE_TYPE (index
);
6161 tree default_label
= gimple_switch_default_label (s
);
6162 basic_block default_label_bb
6163 = label_to_block_fn (func
, CASE_LABEL (default_label
));
6164 basic_block cur_bb
= bb
;
6166 auto_vec
<edge
> new_edges
;
6167 auto_vec
<phi_definition
*> phi_todo_list
;
6168 auto_vec
<profile_count
> edge_counts
;
6169 auto_vec
<profile_probability
> edge_probabilities
;
6171 /* Investigate all labels that and PHI nodes in these edges which
6172 should be fixed after we add new collection of edges. */
6173 for (unsigned i
= 0; i
< labels
; i
++)
6175 tree label
= gimple_switch_label (s
, i
);
6176 basic_block label_bb
= label_to_block_fn (func
, CASE_LABEL (label
));
6177 edge e
= find_edge (bb
, label_bb
);
6178 edge_counts
.safe_push (e
->count
);
6179 edge_probabilities
.safe_push (e
->probability
);
6180 gphi_iterator phi_gsi
;
6182 /* Save PHI definitions that will be destroyed because of an edge
6183 is going to be removed. */
6184 unsigned phi_index
= 0;
6185 for (phi_gsi
= gsi_start_phis (e
->dest
);
6186 !gsi_end_p (phi_gsi
); gsi_next (&phi_gsi
))
6188 gphi
*phi
= phi_gsi
.phi ();
6189 for (unsigned j
= 0; j
< gimple_phi_num_args (phi
); j
++)
6191 if (gimple_phi_arg_edge (phi
, j
) == e
)
6193 tree imm
= gimple_phi_arg_def (phi
, j
);
6194 phi_definition
*p
= new phi_definition (phi_index
, i
,
6196 phi_todo_list
.safe_push (p
);
6204 /* Remove all edges for the current basic block. */
6205 for (int i
= EDGE_COUNT (bb
->succs
) - 1; i
>= 0; i
--)
6207 edge e
= EDGE_SUCC (bb
, i
);
6211 /* Iterate all non-default labels. */
6212 for (unsigned i
= 1; i
< labels
; i
++)
6214 tree label
= gimple_switch_label (s
, i
);
6215 tree low
= CASE_LOW (label
);
6216 tree high
= CASE_HIGH (label
);
6218 if (!useless_type_conversion_p (TREE_TYPE (low
), index_type
))
6219 low
= fold_convert (index_type
, low
);
6221 gimple_stmt_iterator cond_gsi
= gsi_last_bb (cur_bb
);
6225 tree tmp1
= make_temp_ssa_name (boolean_type_node
, NULL
,
6228 gimple
*assign1
= gimple_build_assign (tmp1
, LE_EXPR
, low
,
6231 tree tmp2
= make_temp_ssa_name (boolean_type_node
, NULL
,
6234 if (!useless_type_conversion_p (TREE_TYPE (high
), index_type
))
6235 high
= fold_convert (index_type
, high
);
6236 gimple
*assign2
= gimple_build_assign (tmp2
, LE_EXPR
, index
,
6239 tree tmp3
= make_temp_ssa_name (boolean_type_node
, NULL
,
6241 gimple
*assign3
= gimple_build_assign (tmp3
, BIT_AND_EXPR
, tmp1
,
6244 gsi_insert_before (&cond_gsi
, assign1
, GSI_SAME_STMT
);
6245 gsi_insert_before (&cond_gsi
, assign2
, GSI_SAME_STMT
);
6246 gsi_insert_before (&cond_gsi
, assign3
, GSI_SAME_STMT
);
6248 tree b
= constant_boolean_node (false, boolean_type_node
);
6249 c
= gimple_build_cond (NE_EXPR
, tmp3
, b
, NULL
, NULL
);
6252 c
= gimple_build_cond (EQ_EXPR
, index
, low
, NULL
, NULL
);
6254 gimple_set_location (c
, gimple_location (stmt
));
6256 gsi_insert_before (&cond_gsi
, c
, GSI_SAME_STMT
);
6258 basic_block label_bb
6259 = label_to_block_fn (func
, CASE_LABEL (label
));
6260 edge new_edge
= make_edge (cur_bb
, label_bb
, EDGE_TRUE_VALUE
);
6261 profile_probability prob_sum
= sum_slice
<profile_probability
>
6262 (edge_probabilities
, i
, labels
, profile_probability::never ())
6263 + edge_probabilities
[0];
6265 if (prob_sum
.initialized_p ())
6266 new_edge
->probability
= edge_probabilities
[i
] / prob_sum
;
6268 new_edge
->count
= edge_counts
[i
];
6269 new_edges
.safe_push (new_edge
);
6273 /* Prepare another basic block that will contain
6275 basic_block next_bb
= create_empty_bb (cur_bb
);
6278 add_bb_to_loop (next_bb
, cur_bb
->loop_father
);
6279 loops_state_set (LOOPS_NEED_FIXUP
);
6282 edge next_edge
= make_edge (cur_bb
, next_bb
, EDGE_FALSE_VALUE
);
6283 next_edge
->probability
= new_edge
->probability
.invert ();
6284 next_edge
->count
= edge_counts
[0]
6285 + sum_slice
<profile_count
> (edge_counts
, i
, labels
,
6286 profile_count::zero ());
6287 next_bb
->frequency
= EDGE_FREQUENCY (next_edge
);
6290 else /* Link last IF statement and default label
6293 edge e
= make_edge (cur_bb
, default_label_bb
, EDGE_FALSE_VALUE
);
6294 e
->probability
= new_edge
->probability
.invert ();
6295 e
->count
= edge_counts
[0];
6296 new_edges
.safe_insert (0, e
);
6300 /* Restore original PHI immediate value. */
6301 for (unsigned i
= 0; i
< phi_todo_list
.length (); i
++)
6303 phi_definition
*phi_def
= phi_todo_list
[i
];
6304 edge new_edge
= new_edges
[phi_def
->label_index
];
6306 gphi_iterator it
= gsi_start_phis (new_edge
->dest
);
6307 for (unsigned i
= 0; i
< phi_def
->phi_index
; i
++)
6310 gphi
*phi
= it
.phi ();
6311 add_phi_arg (phi
, phi_def
->phi_value
, new_edge
, UNKNOWN_LOCATION
);
6315 /* Remove the original GIMPLE switch statement. */
6316 gsi_remove (&gsi
, true);
6321 dump_function_to_file (current_function_decl
, dump_file
, TDF_DETAILS
);
6323 return modified_cfg
;
6326 /* Expand builtins that can't be handled by HSA back-end. */
6331 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
6334 FOR_EACH_BB_FN (bb
, func
)
6336 for (gimple_stmt_iterator gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
);
6339 gimple
*stmt
= gsi_stmt (gsi
);
6341 if (gimple_code (stmt
) != GIMPLE_CALL
)
6344 gcall
*call
= as_a
<gcall
*> (stmt
);
6346 if (!gimple_call_builtin_p (call
, BUILT_IN_NORMAL
))
6349 tree fndecl
= gimple_call_fndecl (stmt
);
6350 enum built_in_function fn
= DECL_FUNCTION_CODE (fndecl
);
6353 case BUILT_IN_CEXPF
:
6354 case BUILT_IN_CEXPIF
:
6355 case BUILT_IN_CEXPI
:
6357 /* Similar to builtins.c (expand_builtin_cexpi), the builtin
6358 can be transformed to: cexp(I * z) = ccos(z) + I * csin(z). */
6359 tree lhs
= gimple_call_lhs (stmt
);
6360 tree rhs
= gimple_call_arg (stmt
, 0);
6361 tree rhs_type
= TREE_TYPE (rhs
);
6362 bool float_type_p
= rhs_type
== float_type_node
;
6363 tree real_part
= make_temp_ssa_name (rhs_type
, NULL
,
6365 tree imag_part
= make_temp_ssa_name (rhs_type
, NULL
,
6369 = mathfn_built_in (rhs_type
, fn
== float_type_p
6370 ? BUILT_IN_COSF
: BUILT_IN_COS
);
6371 gcall
*cos
= gimple_build_call (cos_fndecl
, 1, rhs
);
6372 gimple_call_set_lhs (cos
, real_part
);
6373 gsi_insert_before (&gsi
, cos
, GSI_SAME_STMT
);
6376 = mathfn_built_in (rhs_type
, fn
== float_type_p
6377 ? BUILT_IN_SINF
: BUILT_IN_SIN
);
6378 gcall
*sin
= gimple_build_call (sin_fndecl
, 1, rhs
);
6379 gimple_call_set_lhs (sin
, imag_part
);
6380 gsi_insert_before (&gsi
, sin
, GSI_SAME_STMT
);
6383 gassign
*assign
= gimple_build_assign (lhs
, COMPLEX_EXPR
,
6384 real_part
, imag_part
);
6385 gsi_insert_before (&gsi
, assign
, GSI_SAME_STMT
);
6386 gsi_remove (&gsi
, true);
6397 /* Emit HSA module variables that are global for the entire module. */
6400 emit_hsa_module_variables (void)
6402 hsa_num_threads
= new hsa_symbol (BRIG_TYPE_U32
, BRIG_SEGMENT_PRIVATE
,
6403 BRIG_LINKAGE_MODULE
, true);
6405 hsa_num_threads
->m_name
= "hsa_num_threads";
6407 hsa_brig_emit_omp_symbols ();
6410 /* Generate HSAIL representation of the current function and write into a
6411 special section of the output file. If KERNEL is set, the function will be
6412 considered an HSA kernel callable from the host, otherwise it will be
6413 compiled as an HSA function callable from other HSA code. */
6416 generate_hsa (bool kernel
)
6418 hsa_init_data_for_cfun ();
6420 if (hsa_num_threads
== NULL
)
6421 emit_hsa_module_variables ();
6423 bool modified_cfg
= convert_switch_statements ();
6424 /* Initialize hsa_cfun. */
6425 hsa_cfun
= new hsa_function_representation (cfun
->decl
, kernel
,
6426 SSANAMES (cfun
)->length (),
6428 hsa_cfun
->init_extra_bbs ();
6432 HSA_SORRY_AT (UNKNOWN_LOCATION
,
6433 "support for HSA does not implement transactional memory");
6437 verify_function_arguments (cfun
->decl
);
6438 if (hsa_seen_error ())
6441 hsa_cfun
->m_name
= get_brig_function_name (cfun
->decl
);
6443 gen_function_def_parameters ();
6444 if (hsa_seen_error ())
6449 gen_body_from_gimple ();
6450 if (hsa_seen_error ())
6453 if (hsa_cfun
->m_kernel_dispatch_count
)
6454 init_hsa_num_threads ();
6456 if (hsa_cfun
->m_kern_p
)
6458 hsa_function_summary
*s
6459 = hsa_summaries
->get (cgraph_node::get (hsa_cfun
->m_decl
));
6460 hsa_add_kern_decl_mapping (current_function_decl
, hsa_cfun
->m_name
,
6461 hsa_cfun
->m_maximum_omp_data_size
,
6462 s
->m_gridified_kernel_p
);
6467 for (unsigned i
= 0; i
< hsa_cfun
->m_ssa_map
.length (); i
++)
6468 if (hsa_cfun
->m_ssa_map
[i
])
6469 hsa_cfun
->m_ssa_map
[i
]->verify_ssa ();
6472 FOR_EACH_BB_FN (bb
, cfun
)
6474 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
6476 for (hsa_insn_basic
*insn
= hbb
->m_first_insn
; insn
;
6477 insn
= insn
->m_next
)
6483 hsa_brig_emit_function ();
6486 hsa_deinit_data_for_cfun ();
6491 const pass_data pass_data_gen_hsail
=
6494 "hsagen", /* name */
6495 OPTGROUP_OMP
, /* optinfo_flags */
6496 TV_NONE
, /* tv_id */
6497 PROP_cfg
| PROP_ssa
, /* properties_required */
6498 0, /* properties_provided */
6499 0, /* properties_destroyed */
6500 0, /* todo_flags_start */
6501 0 /* todo_flags_finish */
6504 class pass_gen_hsail
: public gimple_opt_pass
6507 pass_gen_hsail (gcc::context
*ctxt
)
6508 : gimple_opt_pass(pass_data_gen_hsail
, ctxt
)
6511 /* opt_pass methods: */
6512 bool gate (function
*);
6513 unsigned int execute (function
*);
6515 }; // class pass_gen_hsail
6517 /* Determine whether or not to run generation of HSAIL. */
6520 pass_gen_hsail::gate (function
*f
)
6522 return hsa_gen_requested_p ()
6523 && hsa_gpu_implementation_p (f
->decl
);
6527 pass_gen_hsail::execute (function
*)
6529 hsa_function_summary
*s
6530 = hsa_summaries
->get (cgraph_node::get_create (current_function_decl
));
6533 generate_hsa (s
->m_kind
== HSA_KERNEL
);
6534 TREE_ASM_WRITTEN (current_function_decl
) = 1;
6535 return TODO_discard_function
;
6540 /* Create the instance of hsa gen pass. */
6543 make_pass_gen_hsail (gcc::context
*ctxt
)
6545 return new pass_gen_hsail (ctxt
);