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 /* If T points to an integral type smaller than 32 bits, change it to a 32bit
568 equivalent and return the result. Otherwise just return the result. */
571 hsa_extend_inttype_to_32bit (BrigType16_t t
)
573 if (t
== BRIG_TYPE_U8
|| t
== BRIG_TYPE_U16
)
574 return BRIG_TYPE_U32
;
575 else if (t
== BRIG_TYPE_S8
|| t
== BRIG_TYPE_S16
)
576 return BRIG_TYPE_S32
;
580 /* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
581 are assumed to use flat addressing. If min32int is true, always expand
582 integer types to one that has at least 32 bits. */
585 hsa_type_for_scalar_tree_type (const_tree type
, bool min32int
)
589 BrigType16_t res
= BRIG_TYPE_NONE
;
591 gcc_checking_assert (TYPE_P (type
));
592 gcc_checking_assert (!AGGREGATE_TYPE_P (type
));
593 if (POINTER_TYPE_P (type
))
594 return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
596 if (TREE_CODE (type
) == VECTOR_TYPE
)
597 base
= TREE_TYPE (type
);
598 else if (TREE_CODE (type
) == COMPLEX_TYPE
)
600 base
= TREE_TYPE (type
);
606 if (!tree_fits_uhwi_p (TYPE_SIZE (base
)))
608 HSA_SORRY_ATV (EXPR_LOCATION (type
),
609 "support for HSA does not implement huge or "
610 "variable-sized type %qT", type
);
614 bsize
= tree_to_uhwi (TYPE_SIZE (base
));
615 unsigned byte_size
= bsize
/ BITS_PER_UNIT
;
616 if (INTEGRAL_TYPE_P (base
))
617 res
= get_integer_type_by_bytes (byte_size
, !TYPE_UNSIGNED (base
));
618 else if (SCALAR_FLOAT_TYPE_P (base
))
636 if (res
== BRIG_TYPE_NONE
)
638 HSA_SORRY_ATV (EXPR_LOCATION (type
),
639 "support for HSA does not implement type %qT", type
);
643 if (TREE_CODE (type
) == VECTOR_TYPE
)
645 HOST_WIDE_INT tsize
= tree_to_uhwi (TYPE_SIZE (type
));
649 HSA_SORRY_ATV (EXPR_LOCATION (type
),
650 "support for HSA does not implement a vector type "
651 "where a type and unit size are equal: %qT", type
);
658 res
|= BRIG_TYPE_PACK_32
;
661 res
|= BRIG_TYPE_PACK_64
;
664 res
|= BRIG_TYPE_PACK_128
;
667 HSA_SORRY_ATV (EXPR_LOCATION (type
),
668 "support for HSA does not implement type %qT", type
);
673 /* Registers/immediate operands can only be 32bit or more except for
675 res
= hsa_extend_inttype_to_32bit (res
);
677 if (TREE_CODE (type
) == COMPLEX_TYPE
)
679 unsigned bsize
= 2 * hsa_type_bit_size (res
);
680 res
= hsa_bittype_for_bitsize (bsize
);
686 /* Returns the BRIG type we need to load/store entities of TYPE. */
689 mem_type_for_type (BrigType16_t type
)
691 /* HSA has non-intuitive constraints on load/store types. If it's
692 a bit-type it _must_ be B128, if it's not a bit-type it must be
693 64bit max. So for loading entities of 128 bits (e.g. vectors)
694 we have to to B128, while for loading the rest we have to use the
695 input type (??? or maybe also flattened to a equally sized non-vector
697 if ((type
& BRIG_TYPE_PACK_MASK
) == BRIG_TYPE_PACK_128
)
698 return BRIG_TYPE_B128
;
699 else if (hsa_btype_p (type
) || hsa_type_packed_p (type
))
701 unsigned bitsize
= hsa_type_bit_size (type
);
703 return hsa_uint_for_bitsize (bitsize
);
705 return hsa_bittype_for_bitsize (bitsize
);
710 /* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some
711 kind of array will be generated, setting DIM appropriately. Otherwise, it
712 will be set to zero. */
715 hsa_type_for_tree_type (const_tree type
, unsigned HOST_WIDE_INT
*dim_p
= NULL
,
716 bool min32int
= false)
718 gcc_checking_assert (TYPE_P (type
));
719 if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type
)))
721 HSA_SORRY_ATV (EXPR_LOCATION (type
), "support for HSA does not "
722 "implement huge or variable-sized type %qT", type
);
723 return BRIG_TYPE_NONE
;
726 if (RECORD_OR_UNION_TYPE_P (type
))
729 *dim_p
= tree_to_uhwi (TYPE_SIZE_UNIT (type
));
730 return BRIG_TYPE_U8
| BRIG_TYPE_ARRAY
;
733 if (TREE_CODE (type
) == ARRAY_TYPE
)
735 /* We try to be nice and use the real base-type when this is an array of
736 scalars and only resort to an array of bytes if the type is more
739 unsigned HOST_WIDE_INT dim
= 1;
741 while (TREE_CODE (type
) == ARRAY_TYPE
)
743 tree domain
= TYPE_DOMAIN (type
);
744 if (!TYPE_MIN_VALUE (domain
)
745 || !TYPE_MAX_VALUE (domain
)
746 || !tree_fits_shwi_p (TYPE_MIN_VALUE (domain
))
747 || !tree_fits_shwi_p (TYPE_MAX_VALUE (domain
)))
749 HSA_SORRY_ATV (EXPR_LOCATION (type
),
750 "support for HSA does not implement array "
751 "%qT with unknown bounds", type
);
752 return BRIG_TYPE_NONE
;
754 HOST_WIDE_INT min
= tree_to_shwi (TYPE_MIN_VALUE (domain
));
755 HOST_WIDE_INT max
= tree_to_shwi (TYPE_MAX_VALUE (domain
));
756 dim
= dim
* (unsigned HOST_WIDE_INT
) (max
- min
+ 1);
757 type
= TREE_TYPE (type
);
761 if (RECORD_OR_UNION_TYPE_P (type
))
763 dim
= dim
* tree_to_uhwi (TYPE_SIZE_UNIT (type
));
767 res
= hsa_type_for_scalar_tree_type (type
, false);
771 return res
| BRIG_TYPE_ARRAY
;
778 return hsa_type_for_scalar_tree_type (type
, min32int
);
781 /* Returns true if converting from STYPE into DTYPE needs the _CVT
782 opcode. If false a normal _MOV is enough. */
785 hsa_needs_cvt (BrigType16_t dtype
, BrigType16_t stype
)
787 if (hsa_btype_p (dtype
))
790 /* float <-> int conversions are real converts. */
791 if (hsa_type_float_p (dtype
) != hsa_type_float_p (stype
))
793 /* When both types have different size, then we need CVT as well. */
794 if (hsa_type_bit_size (dtype
) != hsa_type_bit_size (stype
))
799 /* Return declaration name if it exists or create one from UID if it does not.
800 If DECL is a local variable, make UID part of its name. */
803 hsa_get_declaration_name (tree decl
)
805 if (!DECL_NAME (decl
))
808 snprintf (buf
, 64, "__hsa_anon_%u", DECL_UID (decl
));
809 size_t len
= strlen (buf
);
810 char *copy
= (char *) obstack_alloc (&hsa_obstack
, len
+ 1);
811 memcpy (copy
, buf
, len
+ 1);
816 if (TREE_CODE (decl
) == FUNCTION_DECL
817 || (TREE_CODE (decl
) == VAR_DECL
&& is_global_var (decl
)))
818 name_tree
= DECL_ASSEMBLER_NAME (decl
);
820 name_tree
= DECL_NAME (decl
);
822 const char *name
= IDENTIFIER_POINTER (name_tree
);
823 /* User-defined assembly names have prepended asterisk symbol. */
827 if ((TREE_CODE (decl
) == VAR_DECL
)
828 && decl_function_context (decl
))
830 size_t len
= strlen (name
);
831 char *buf
= (char *) alloca (len
+ 32);
832 snprintf (buf
, len
+ 32, "%s_%u", name
, DECL_UID (decl
));
834 char *copy
= (char *) obstack_alloc (&hsa_obstack
, len
+ 1);
835 memcpy (copy
, buf
, len
+ 1);
842 /* Lookup or create the associated hsa_symbol structure with a given VAR_DECL
843 or lookup the hsa_structure corresponding to a PARM_DECL. */
846 get_symbol_for_decl (tree decl
)
849 hsa_symbol
dummy (BRIG_TYPE_NONE
, BRIG_SEGMENT_NONE
, BRIG_LINKAGE_NONE
);
851 gcc_assert (TREE_CODE (decl
) == PARM_DECL
852 || TREE_CODE (decl
) == RESULT_DECL
853 || TREE_CODE (decl
) == VAR_DECL
854 || TREE_CODE (decl
) == CONST_DECL
);
858 bool is_in_global_vars
= ((TREE_CODE (decl
) == VAR_DECL
)
859 && !decl_function_context (decl
));
861 if (is_in_global_vars
)
862 slot
= hsa_global_variable_symbols
->find_slot (&dummy
, INSERT
);
864 slot
= hsa_cfun
->m_local_symbols
->find_slot (&dummy
, INSERT
);
866 gcc_checking_assert (slot
);
869 hsa_symbol
*sym
= (*slot
);
871 /* If the symbol is problematic, mark current function also as
873 if (sym
->m_seen_error
)
876 /* PR hsa/70234: If a global variable was marked to be emitted,
877 but HSAIL generation of a function using the variable fails,
878 we should retry to emit the variable in context of a different
881 Iterate elements whether a symbol is already in m_global_symbols
883 if (is_in_global_vars
&& !sym
->m_emitted_to_brig
)
885 for (unsigned i
= 0; i
< hsa_cfun
->m_global_symbols
.length (); i
++)
886 if (hsa_cfun
->m_global_symbols
[i
] == sym
)
888 hsa_cfun
->m_global_symbols
.safe_push (sym
);
896 /* PARM_DECLs and RESULT_DECL should be already in m_local_symbols. */
897 gcc_assert (TREE_CODE (decl
) == VAR_DECL
898 || TREE_CODE (decl
) == CONST_DECL
);
899 BrigAlignment8_t align
= hsa_object_alignment (decl
);
901 if (is_in_global_vars
)
903 gcc_checking_assert (TREE_CODE (decl
) != CONST_DECL
);
904 sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_GLOBAL
,
905 BRIG_LINKAGE_PROGRAM
, true,
906 BRIG_ALLOCATION_PROGRAM
, align
);
907 hsa_cfun
->m_global_symbols
.safe_push (sym
);
908 sym
->fillup_for_decl (decl
);
909 if (sym
->m_align
> align
)
911 sym
->m_seen_error
= true;
912 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
913 "HSA specification requires that %E is at least "
914 "naturally aligned", decl
);
919 /* As generation of efficient memory copy instructions relies
920 on alignment greater or equal to 8 bytes,
921 we need to increase alignment of all aggregate types.. */
922 if (AGGREGATE_TYPE_P (TREE_TYPE (decl
)))
923 align
= MAX ((BrigAlignment8_t
) BRIG_ALIGNMENT_8
, align
);
925 BrigAllocation allocation
= BRIG_ALLOCATION_AUTOMATIC
;
926 BrigSegment8_t segment
;
927 if (TREE_CODE (decl
) == CONST_DECL
)
929 segment
= BRIG_SEGMENT_READONLY
;
930 allocation
= BRIG_ALLOCATION_AGENT
;
932 else if (lookup_attribute ("hsa_group_segment",
933 DECL_ATTRIBUTES (decl
)))
934 segment
= BRIG_SEGMENT_GROUP
;
935 else if (TREE_STATIC (decl
)
936 || lookup_attribute ("hsa_global_segment",
937 DECL_ATTRIBUTES (decl
)))
938 segment
= BRIG_SEGMENT_GLOBAL
;
940 segment
= BRIG_SEGMENT_PRIVATE
;
942 sym
= new hsa_symbol (BRIG_TYPE_NONE
, segment
, BRIG_LINKAGE_FUNCTION
,
943 false, allocation
, align
);
944 sym
->fillup_for_decl (decl
);
945 hsa_cfun
->m_private_variables
.safe_push (sym
);
948 sym
->m_name
= hsa_get_declaration_name (decl
);
954 /* For a given HSA function declaration, return a host
955 function declaration. */
958 hsa_get_host_function (tree decl
)
960 hsa_function_summary
*s
961 = hsa_summaries
->get (cgraph_node::get_create (decl
));
962 gcc_assert (s
->m_kind
!= HSA_NONE
);
963 gcc_assert (s
->m_gpu_implementation_p
);
965 return s
->m_bound_function
? s
->m_bound_function
->decl
: NULL
;
968 /* Return true if function DECL has a host equivalent function. */
971 get_brig_function_name (tree decl
)
975 hsa_function_summary
*s
= hsa_summaries
->get (cgraph_node::get_create (d
));
976 if (s
->m_kind
!= HSA_NONE
977 && s
->m_gpu_implementation_p
978 && s
->m_bound_function
)
979 d
= s
->m_bound_function
->decl
;
981 /* IPA split can create a function that has no host equivalent. */
985 char *name
= xstrdup (hsa_get_declaration_name (d
));
986 hsa_sanitize_name (name
);
991 /* Create a spill symbol of type TYPE. */
994 hsa_get_spill_symbol (BrigType16_t type
)
996 hsa_symbol
*sym
= new hsa_symbol (type
, BRIG_SEGMENT_SPILL
,
997 BRIG_LINKAGE_FUNCTION
);
998 hsa_cfun
->m_spill_symbols
.safe_push (sym
);
1002 /* Create a symbol for a read-only string constant. */
1004 hsa_get_string_cst_symbol (tree string_cst
)
1006 gcc_checking_assert (TREE_CODE (string_cst
) == STRING_CST
);
1008 hsa_symbol
**slot
= hsa_cfun
->m_string_constants_map
.get (string_cst
);
1012 hsa_op_immed
*cst
= new hsa_op_immed (string_cst
);
1013 hsa_symbol
*sym
= new hsa_symbol (cst
->m_type
, BRIG_SEGMENT_GLOBAL
,
1014 BRIG_LINKAGE_MODULE
, true,
1015 BRIG_ALLOCATION_AGENT
);
1016 sym
->m_cst_value
= cst
;
1017 sym
->m_dim
= TREE_STRING_LENGTH (string_cst
);
1018 sym
->m_name_number
= hsa_cfun
->m_global_symbols
.length ();
1020 hsa_cfun
->m_global_symbols
.safe_push (sym
);
1021 hsa_cfun
->m_string_constants_map
.put (string_cst
, sym
);
1025 /* Make the type of a MOV instruction larger if mandated by HSAIL rules. */
1028 hsa_fixup_mov_insn_type (hsa_insn_basic
*insn
)
1030 insn
->m_type
= hsa_extend_inttype_to_32bit (insn
->m_type
);
1031 if (insn
->m_type
== BRIG_TYPE_B8
|| insn
->m_type
== BRIG_TYPE_B16
)
1032 insn
->m_type
= BRIG_TYPE_B32
;
1035 /* Constructor of the ancestor of all operands. K is BRIG kind that identified
1036 what the operator is. */
1038 hsa_op_base::hsa_op_base (BrigKind16_t k
)
1039 : m_next (NULL
), m_brig_op_offset (0), m_kind (k
)
1041 hsa_operands
.safe_push (this);
1044 /* Constructor of ancestor of all operands which have a type. K is BRIG kind
1045 that identified what the operator is. T is the type of the operator. */
1047 hsa_op_with_type::hsa_op_with_type (BrigKind16_t k
, BrigType16_t t
)
1048 : hsa_op_base (k
), m_type (t
)
1053 hsa_op_with_type::get_in_type (BrigType16_t dtype
, hsa_bb
*hbb
)
1055 if (m_type
== dtype
)
1060 if (hsa_needs_cvt (dtype
, m_type
))
1062 dest
= new hsa_op_reg (dtype
);
1063 hbb
->append_insn (new hsa_insn_cvt (dest
, this));
1065 else if (is_a
<hsa_op_reg
*> (this))
1067 /* In the end, HSA registers do not really have types, only sizes, so if
1068 the sizes match, we can use the register directly. */
1069 gcc_checking_assert (hsa_type_bit_size (dtype
)
1070 == hsa_type_bit_size (m_type
));
1075 dest
= new hsa_op_reg (m_type
);
1077 hsa_insn_basic
*mov
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
1078 dest
->m_type
, dest
, this);
1079 hsa_fixup_mov_insn_type (mov
);
1080 hbb
->append_insn (mov
);
1081 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
1082 type of the operand must be same as type of the instruction. */
1083 dest
->m_type
= dtype
;
1089 /* If this operand has integer type smaller than 32 bits, extend it to 32 bits,
1090 adding instructions to HBB if needed. */
1093 hsa_op_with_type::extend_int_to_32bit (hsa_bb
*hbb
)
1095 if (m_type
== BRIG_TYPE_U8
|| m_type
== BRIG_TYPE_U16
)
1096 return get_in_type (BRIG_TYPE_U32
, hbb
);
1097 else if (m_type
== BRIG_TYPE_S8
|| m_type
== BRIG_TYPE_S16
)
1098 return get_in_type (BRIG_TYPE_S32
, hbb
);
1103 /* Constructor of class representing HSA immediate values. TREE_VAL is the
1104 tree representation of the immediate value. If min32int is true,
1105 always expand integer types to one that has at least 32 bits. */
1107 hsa_op_immed::hsa_op_immed (tree tree_val
, bool min32int
)
1108 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
,
1109 hsa_type_for_tree_type (TREE_TYPE (tree_val
), NULL
,
1112 if (hsa_seen_error ())
1115 gcc_checking_assert ((is_gimple_min_invariant (tree_val
)
1116 && (!POINTER_TYPE_P (TREE_TYPE (tree_val
))
1117 || TREE_CODE (tree_val
) == INTEGER_CST
))
1118 || TREE_CODE (tree_val
) == CONSTRUCTOR
);
1119 m_tree_value
= tree_val
;
1121 /* Verify that all elements of a constructor are constants. */
1122 if (TREE_CODE (m_tree_value
) == CONSTRUCTOR
)
1123 for (unsigned i
= 0; i
< CONSTRUCTOR_NELTS (m_tree_value
); i
++)
1125 tree v
= CONSTRUCTOR_ELT (m_tree_value
, i
)->value
;
1126 if (!CONSTANT_CLASS_P (v
))
1128 HSA_SORRY_AT (EXPR_LOCATION (tree_val
),
1129 "HSA ctor should have only constants");
1135 /* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
1136 integer representation of the immediate value. TYPE is BRIG type. */
1138 hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value
, BrigType16_t type
)
1139 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
, type
),
1142 gcc_assert (hsa_type_integer_p (type
));
1143 m_int_value
= integer_value
;
1146 hsa_op_immed::hsa_op_immed ()
1147 : hsa_op_with_type (BRIG_KIND_NONE
, BRIG_TYPE_NONE
)
1151 /* New operator to allocate immediate operands from obstack. */
1154 hsa_op_immed::operator new (size_t size
)
1156 return obstack_alloc (&hsa_obstack
, size
);
1161 hsa_op_immed::~hsa_op_immed ()
1165 /* Change type of the immediate value to T. */
1168 hsa_op_immed::set_type (BrigType16_t t
)
1173 /* Constructor of class representing HSA registers and pseudo-registers. T is
1174 the BRIG type of the new register. */
1176 hsa_op_reg::hsa_op_reg (BrigType16_t t
)
1177 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER
, t
), m_gimple_ssa (NULL_TREE
),
1178 m_def_insn (NULL
), m_spill_sym (NULL
), m_order (hsa_cfun
->m_reg_count
++),
1179 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
1183 /* New operator to allocate a register from obstack. */
1186 hsa_op_reg::operator new (size_t size
)
1188 return obstack_alloc (&hsa_obstack
, size
);
1191 /* Verify register operand. */
1194 hsa_op_reg::verify_ssa ()
1196 /* Verify that each HSA register has a definition assigned.
1197 Exceptions are VAR_DECL and PARM_DECL that are a default
1199 gcc_checking_assert (m_def_insn
1200 || (m_gimple_ssa
!= NULL
1201 && (!SSA_NAME_VAR (m_gimple_ssa
)
1202 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa
))
1204 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa
)));
1206 /* Verify that every use of the register is really present
1207 in an instruction. */
1208 for (unsigned i
= 0; i
< m_uses
.length (); i
++)
1210 hsa_insn_basic
*use
= m_uses
[i
];
1212 bool is_visited
= false;
1213 for (unsigned j
= 0; j
< use
->operand_count (); j
++)
1215 hsa_op_base
*u
= use
->get_op (j
);
1216 hsa_op_address
*addr
; addr
= dyn_cast
<hsa_op_address
*> (u
);
1217 if (addr
&& addr
->m_reg
)
1222 bool r
= !addr
&& use
->op_output_p (j
);
1226 error ("HSA SSA name defined by instruction that is supposed "
1228 debug_hsa_operand (this);
1229 debug_hsa_insn (use
);
1230 internal_error ("HSA SSA verification failed");
1239 error ("HSA SSA name not among operands of instruction that is "
1240 "supposed to use it");
1241 debug_hsa_operand (this);
1242 debug_hsa_insn (use
);
1243 internal_error ("HSA SSA verification failed");
1248 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, hsa_op_reg
*r
,
1249 HOST_WIDE_INT offset
)
1250 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (r
),
1251 m_imm_offset (offset
)
1255 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, HOST_WIDE_INT offset
)
1256 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (NULL
),
1257 m_imm_offset (offset
)
1261 hsa_op_address::hsa_op_address (hsa_op_reg
*r
, HOST_WIDE_INT offset
)
1262 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (NULL
), m_reg (r
),
1263 m_imm_offset (offset
)
1267 /* New operator to allocate address operands from obstack. */
1270 hsa_op_address::operator new (size_t size
)
1272 return obstack_alloc (&hsa_obstack
, size
);
1275 /* Constructor of an operand referring to HSAIL code. */
1277 hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF
),
1278 m_directive_offset (0)
1282 /* Constructor of an operand representing a code list. Set it up so that it
1283 can contain ELEMENTS number of elements. */
1285 hsa_op_code_list::hsa_op_code_list (unsigned elements
)
1286 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST
)
1288 m_offsets
.create (1);
1289 m_offsets
.safe_grow_cleared (elements
);
1292 /* New operator to allocate code list operands from obstack. */
1295 hsa_op_code_list::operator new (size_t size
)
1297 return obstack_alloc (&hsa_obstack
, size
);
1300 /* Constructor of an operand representing an operand list.
1301 Set it up so that it can contain ELEMENTS number of elements. */
1303 hsa_op_operand_list::hsa_op_operand_list (unsigned elements
)
1304 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST
)
1306 m_offsets
.create (elements
);
1307 m_offsets
.safe_grow (elements
);
1310 /* New operator to allocate operand list operands from obstack. */
1313 hsa_op_operand_list::operator new (size_t size
)
1315 return obstack_alloc (&hsa_obstack
, size
);
1318 hsa_op_operand_list::~hsa_op_operand_list ()
1320 m_offsets
.release ();
1325 hsa_function_representation::reg_for_gimple_ssa (tree ssa
)
1329 gcc_checking_assert (TREE_CODE (ssa
) == SSA_NAME
);
1330 if (m_ssa_map
[SSA_NAME_VERSION (ssa
)])
1331 return m_ssa_map
[SSA_NAME_VERSION (ssa
)];
1333 hreg
= new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa
),
1335 hreg
->m_gimple_ssa
= ssa
;
1336 m_ssa_map
[SSA_NAME_VERSION (ssa
)] = hreg
;
1342 hsa_op_reg::set_definition (hsa_insn_basic
*insn
)
1344 if (hsa_cfun
->m_in_ssa
)
1346 gcc_checking_assert (!m_def_insn
);
1353 /* Constructor of the class which is the bases of all instructions and directly
1354 represents the most basic ones. NOPS is the number of operands that the
1355 operand vector will contain (and which will be cleared). OP is the opcode
1356 of the instruction. This constructor does not set type. */
1358 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
)
1360 m_next (NULL
), m_bb (NULL
), m_opcode (opc
), m_number (0),
1361 m_type (BRIG_TYPE_NONE
), m_brig_offset (0)
1364 m_operands
.safe_grow_cleared (nops
);
1366 hsa_instructions
.safe_push (this);
1369 /* Make OP the operand number INDEX of operands of this instruction. If OP is a
1370 register or an address containing a register, then either set the definition
1371 of the register to this instruction if it an output operand or add this
1372 instruction to the uses if it is an input one. */
1375 hsa_insn_basic::set_op (int index
, hsa_op_base
*op
)
1377 /* Each address operand is always use. */
1378 hsa_op_address
*addr
= dyn_cast
<hsa_op_address
*> (op
);
1379 if (addr
&& addr
->m_reg
)
1380 addr
->m_reg
->m_uses
.safe_push (this);
1383 hsa_op_reg
*reg
= dyn_cast
<hsa_op_reg
*> (op
);
1386 if (op_output_p (index
))
1387 reg
->set_definition (this);
1389 reg
->m_uses
.safe_push (this);
1393 m_operands
[index
] = op
;
1396 /* Get INDEX-th operand of the instruction. */
1399 hsa_insn_basic::get_op (int index
)
1401 return m_operands
[index
];
1404 /* Get address of INDEX-th operand of the instruction. */
1407 hsa_insn_basic::get_op_addr (int index
)
1409 return &m_operands
[index
];
1412 /* Get number of operands of the instruction. */
1414 hsa_insn_basic::operand_count ()
1416 return m_operands
.length ();
1419 /* Constructor of the class which is the bases of all instructions and directly
1420 represents the most basic ones. NOPS is the number of operands that the
1421 operand vector will contain (and which will be cleared). OPC is the opcode
1422 of the instruction, T is the type of the instruction. */
1424 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
, BrigType16_t t
,
1425 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1426 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1427 : m_prev (NULL
), m_next (NULL
), m_bb (NULL
), m_opcode (opc
),m_number (0),
1428 m_type (t
), m_brig_offset (0)
1431 m_operands
.safe_grow_cleared (nops
);
1435 gcc_checking_assert (nops
>= 1);
1441 gcc_checking_assert (nops
>= 2);
1447 gcc_checking_assert (nops
>= 3);
1453 gcc_checking_assert (nops
>= 4);
1457 hsa_instructions
.safe_push (this);
1460 /* New operator to allocate basic instruction from obstack. */
1463 hsa_insn_basic::operator new (size_t size
)
1465 return obstack_alloc (&hsa_obstack
, size
);
1468 /* Verify the instruction. */
1471 hsa_insn_basic::verify ()
1473 hsa_op_address
*addr
;
1476 /* Iterate all register operands and verify that the instruction
1477 is set in uses of the register. */
1478 for (unsigned i
= 0; i
< operand_count (); i
++)
1480 hsa_op_base
*use
= get_op (i
);
1482 if ((addr
= dyn_cast
<hsa_op_address
*> (use
)) && addr
->m_reg
)
1484 gcc_assert (addr
->m_reg
->m_def_insn
!= this);
1488 if ((reg
= dyn_cast
<hsa_op_reg
*> (use
)) && !op_output_p (i
))
1491 for (j
= 0; j
< reg
->m_uses
.length (); j
++)
1493 if (reg
->m_uses
[j
] == this)
1497 if (j
== reg
->m_uses
.length ())
1499 error ("HSA instruction uses a register but is not among "
1500 "recorded register uses");
1501 debug_hsa_operand (reg
);
1502 debug_hsa_insn (this);
1503 internal_error ("HSA instruction verification failed");
1509 /* Constructor of an instruction representing a PHI node. NOPS is the number
1510 of operands (equal to the number of predecessors). */
1512 hsa_insn_phi::hsa_insn_phi (unsigned nops
, hsa_op_reg
*dst
)
1513 : hsa_insn_basic (nops
, HSA_OPCODE_PHI
), m_dest (dst
)
1515 dst
->set_definition (this);
1518 /* Constructor of class representing instructions for control flow and
1521 hsa_insn_br::hsa_insn_br (unsigned nops
, int opc
, BrigType16_t t
,
1522 BrigWidth8_t width
, hsa_op_base
*arg0
,
1523 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1525 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1530 /* Constructor of class representing instruction for conditional jump, CTRL is
1531 the control register determining whether the jump will be carried out, the
1532 new instruction is automatically added to its uses list. */
1534 hsa_insn_cbr::hsa_insn_cbr (hsa_op_reg
*ctrl
)
1535 : hsa_insn_br (1, BRIG_OPCODE_CBR
, BRIG_TYPE_B1
, BRIG_WIDTH_1
, ctrl
)
1539 /* Constructor of class representing instruction for switch jump, CTRL is
1540 the index register. */
1542 hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg
*index
, unsigned jump_count
)
1543 : hsa_insn_basic (1, BRIG_OPCODE_SBR
, BRIG_TYPE_B1
, index
),
1544 m_width (BRIG_WIDTH_1
), m_jump_table (vNULL
),
1545 m_label_code_list (new hsa_op_code_list (jump_count
))
1549 /* Replace all occurrences of OLD_BB with NEW_BB in the statements
1553 hsa_insn_sbr::replace_all_labels (basic_block old_bb
, basic_block new_bb
)
1555 for (unsigned i
= 0; i
< m_jump_table
.length (); i
++)
1556 if (m_jump_table
[i
] == old_bb
)
1557 m_jump_table
[i
] = new_bb
;
1560 hsa_insn_sbr::~hsa_insn_sbr ()
1562 m_jump_table
.release ();
1565 /* Constructor of comparison instruction. CMP is the comparison operation and T
1566 is the result type. */
1568 hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp
, BrigType16_t t
,
1569 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1571 : hsa_insn_basic (3 , BRIG_OPCODE_CMP
, t
, arg0
, arg1
, arg2
), m_compare (cmp
)
1575 /* Constructor of classes representing memory accesses. OPC is the opcode (must
1576 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
1577 operands are provided as ARG0 and ARG1. */
1579 hsa_insn_mem::hsa_insn_mem (int opc
, BrigType16_t t
, hsa_op_base
*arg0
,
1581 : hsa_insn_basic (2, opc
, t
, arg0
, arg1
),
1582 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1584 gcc_checking_assert (opc
== BRIG_OPCODE_LD
|| opc
== BRIG_OPCODE_ST
);
1587 /* Constructor for descendants allowing different opcodes and number of
1588 operands, it passes its arguments directly to hsa_insn_basic
1589 constructor. The instruction operands are provided as ARG[0-3]. */
1592 hsa_insn_mem::hsa_insn_mem (unsigned nops
, int opc
, BrigType16_t t
,
1593 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1594 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1595 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1596 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1600 /* Constructor of class representing atomic instructions. OPC is the principal
1601 opcode, AOP is the specific atomic operation opcode. T is the type of the
1602 instruction. The instruction operands are provided as ARG[0-3]. */
1604 hsa_insn_atomic::hsa_insn_atomic (int nops
, int opc
,
1605 enum BrigAtomicOperation aop
,
1606 BrigType16_t t
, BrigMemoryOrder memorder
,
1608 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1610 : hsa_insn_mem (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
), m_atomicop (aop
),
1611 m_memoryorder (memorder
),
1612 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM
)
1614 gcc_checking_assert (opc
== BRIG_OPCODE_ATOMICNORET
||
1615 opc
== BRIG_OPCODE_ATOMIC
||
1616 opc
== BRIG_OPCODE_SIGNAL
||
1617 opc
== BRIG_OPCODE_SIGNALNORET
);
1620 /* Constructor of class representing signal instructions. OPC is the prinicpal
1621 opcode, SOP is the specific signal operation opcode. T is the type of the
1622 instruction. The instruction operands are provided as ARG[0-3]. */
1624 hsa_insn_signal::hsa_insn_signal (int nops
, int opc
,
1625 enum BrigAtomicOperation sop
,
1626 BrigType16_t t
, BrigMemoryOrder memorder
,
1627 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1628 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1629 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1630 m_memory_order (memorder
), m_signalop (sop
)
1634 /* Constructor of class representing segment conversion instructions. OPC is
1635 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
1636 and SRCT are destination and source types respectively, SEG is the segment
1637 we are converting to or from. The instruction operands are
1638 provided as ARG0 and ARG1. */
1640 hsa_insn_seg::hsa_insn_seg (int opc
, BrigType16_t dest
, BrigType16_t srct
,
1641 BrigSegment8_t seg
, hsa_op_base
*arg0
,
1643 : hsa_insn_basic (2, opc
, dest
, arg0
, arg1
), m_src_type (srct
),
1646 gcc_checking_assert (opc
== BRIG_OPCODE_STOF
|| opc
== BRIG_OPCODE_FTOS
);
1649 /* Constructor of class representing a call instruction. CALLEE is the tree
1650 representation of the function being called. */
1652 hsa_insn_call::hsa_insn_call (tree callee
)
1653 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (callee
),
1654 m_output_arg (NULL
), m_args_code_list (NULL
), m_result_code_list (NULL
)
1658 hsa_insn_call::hsa_insn_call (hsa_internal_fn
*fn
)
1659 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (NULL
),
1660 m_called_internal_fn (fn
), m_output_arg (NULL
), m_args_code_list (NULL
),
1661 m_result_code_list (NULL
)
1665 hsa_insn_call::~hsa_insn_call ()
1667 for (unsigned i
= 0; i
< m_input_args
.length (); i
++)
1668 delete m_input_args
[i
];
1670 delete m_output_arg
;
1672 m_input_args
.release ();
1673 m_input_arg_insns
.release ();
1676 /* Constructor of class representing the argument block required to invoke
1678 hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind
,
1679 hsa_insn_call
* call
)
1680 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK
), m_kind (brig_kind
),
1685 hsa_insn_comment::hsa_insn_comment (const char *s
)
1686 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT
)
1688 unsigned l
= strlen (s
);
1690 /* Append '// ' to the string. */
1691 char *buf
= XNEWVEC (char, l
+ 4);
1692 sprintf (buf
, "// %s", s
);
1696 hsa_insn_comment::~hsa_insn_comment ()
1698 gcc_checking_assert (m_comment
);
1703 /* Constructor of class representing the queue instruction in HSAIL. */
1705 hsa_insn_queue::hsa_insn_queue (int nops
, int opcode
, BrigSegment segment
,
1706 BrigMemoryOrder memory_order
,
1707 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1708 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1709 : hsa_insn_basic (nops
, opcode
, BRIG_TYPE_U64
, arg0
, arg1
, arg2
, arg3
),
1710 m_segment (segment
), m_memory_order (memory_order
)
1714 /* Constructor of class representing the source type instruction in HSAIL. */
1716 hsa_insn_srctype::hsa_insn_srctype (int nops
, BrigOpcode opcode
,
1717 BrigType16_t destt
, BrigType16_t srct
,
1718 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1719 hsa_op_base
*arg2
= NULL
)
1720 : hsa_insn_basic (nops
, opcode
, destt
, arg0
, arg1
, arg2
),
1721 m_source_type (srct
)
1724 /* Constructor of class representing the packed instruction in HSAIL. */
1726 hsa_insn_packed::hsa_insn_packed (int nops
, BrigOpcode opcode
,
1727 BrigType16_t destt
, BrigType16_t srct
,
1728 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1730 : hsa_insn_srctype (nops
, opcode
, destt
, srct
, arg0
, arg1
, arg2
)
1732 m_operand_list
= new hsa_op_operand_list (nops
- 1);
1735 /* Constructor of class representing the convert instruction in HSAIL. */
1737 hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type
*dest
, hsa_op_with_type
*src
)
1738 : hsa_insn_basic (2, BRIG_OPCODE_CVT
, dest
->m_type
, dest
, src
)
1742 /* Constructor of class representing the alloca in HSAIL. */
1744 hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type
*dest
,
1745 hsa_op_with_type
*size
, unsigned alignment
)
1746 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA
, dest
->m_type
, dest
, size
),
1747 m_align (BRIG_ALIGNMENT_8
)
1749 gcc_assert (dest
->m_type
== BRIG_TYPE_U32
);
1751 m_align
= hsa_alignment_encoding (alignment
);
1754 /* Append an instruction INSN into the basic block. */
1757 hsa_bb::append_insn (hsa_insn_basic
*insn
)
1759 gcc_assert (insn
->m_opcode
!= 0 || insn
->operand_count () == 0);
1760 gcc_assert (!insn
->m_bb
);
1763 insn
->m_prev
= m_last_insn
;
1764 insn
->m_next
= NULL
;
1766 m_last_insn
->m_next
= insn
;
1769 m_first_insn
= insn
;
1773 hsa_bb::append_phi (hsa_insn_phi
*hphi
)
1777 hphi
->m_prev
= m_last_phi
;
1778 hphi
->m_next
= NULL
;
1780 m_last_phi
->m_next
= hphi
;
1786 /* Insert HSA instruction NEW_INSN immediately before an existing instruction
1790 hsa_insert_insn_before (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1792 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1794 if (hbb
->m_first_insn
== old_insn
)
1795 hbb
->m_first_insn
= new_insn
;
1796 new_insn
->m_prev
= old_insn
->m_prev
;
1797 new_insn
->m_next
= old_insn
;
1798 if (old_insn
->m_prev
)
1799 old_insn
->m_prev
->m_next
= new_insn
;
1800 old_insn
->m_prev
= new_insn
;
1803 /* Append HSA instruction NEW_INSN immediately after an existing instruction
1807 hsa_append_insn_after (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1809 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1811 if (hbb
->m_last_insn
== old_insn
)
1812 hbb
->m_last_insn
= new_insn
;
1813 new_insn
->m_prev
= old_insn
;
1814 new_insn
->m_next
= old_insn
->m_next
;
1815 if (old_insn
->m_next
)
1816 old_insn
->m_next
->m_prev
= new_insn
;
1817 old_insn
->m_next
= new_insn
;
1820 /* Return a register containing the calculated value of EXP which must be an
1821 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
1822 integer constants as returned by get_inner_reference.
1823 Newly generated HSA instructions will be appended to HBB.
1824 Perform all calculations in ADDRTYPE. */
1826 static hsa_op_with_type
*
1827 gen_address_calculation (tree exp
, hsa_bb
*hbb
, BrigType16_t addrtype
)
1831 if (TREE_CODE (exp
) == NOP_EXPR
)
1832 exp
= TREE_OPERAND (exp
, 0);
1834 switch (TREE_CODE (exp
))
1837 return hsa_cfun
->reg_for_gimple_ssa (exp
)->get_in_type (addrtype
, hbb
);
1841 hsa_op_immed
*imm
= new hsa_op_immed (exp
);
1842 if (addrtype
!= imm
->m_type
)
1843 imm
->m_type
= addrtype
;
1848 opcode
= BRIG_OPCODE_ADD
;
1852 opcode
= BRIG_OPCODE_MUL
;
1859 hsa_op_reg
*res
= new hsa_op_reg (addrtype
);
1860 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, addrtype
);
1861 insn
->set_op (0, res
);
1863 hsa_op_with_type
*op1
= gen_address_calculation (TREE_OPERAND (exp
, 0), hbb
,
1865 hsa_op_with_type
*op2
= gen_address_calculation (TREE_OPERAND (exp
, 1), hbb
,
1867 insn
->set_op (1, op1
);
1868 insn
->set_op (2, op2
);
1870 hbb
->append_insn (insn
);
1874 /* If R1 is NULL, just return R2, otherwise append an instruction adding them
1875 to HBB and return the register holding the result. */
1878 add_addr_regs_if_needed (hsa_op_reg
*r1
, hsa_op_reg
*r2
, hsa_bb
*hbb
)
1880 gcc_checking_assert (r2
);
1884 hsa_op_reg
*res
= new hsa_op_reg (r1
->m_type
);
1885 gcc_assert (!hsa_needs_cvt (r1
->m_type
, r2
->m_type
));
1886 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_ADD
, res
->m_type
);
1887 insn
->set_op (0, res
);
1888 insn
->set_op (1, r1
);
1889 insn
->set_op (2, r2
);
1890 hbb
->append_insn (insn
);
1894 /* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
1895 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
1898 process_mem_base (tree base
, hsa_symbol
**symbol
, BrigType16_t
*addrtype
,
1899 hsa_op_reg
**reg
, offset_int
*offset
, hsa_bb
*hbb
)
1901 if (TREE_CODE (base
) == SSA_NAME
)
1904 hsa_op_with_type
*ssa
1905 = hsa_cfun
->reg_for_gimple_ssa (base
)->get_in_type (*addrtype
, hbb
);
1906 *reg
= dyn_cast
<hsa_op_reg
*> (ssa
);
1908 else if (TREE_CODE (base
) == ADDR_EXPR
)
1910 tree decl
= TREE_OPERAND (base
, 0);
1912 if (!DECL_P (decl
) || TREE_CODE (decl
) == FUNCTION_DECL
)
1914 HSA_SORRY_AT (EXPR_LOCATION (base
),
1915 "support for HSA does not implement a memory reference "
1916 "to a non-declaration type");
1920 gcc_assert (!*symbol
);
1922 *symbol
= get_symbol_for_decl (decl
);
1923 *addrtype
= hsa_get_segment_addr_type ((*symbol
)->m_segment
);
1925 else if (TREE_CODE (base
) == INTEGER_CST
)
1926 *offset
+= wi::to_offset (base
);
1931 /* Forward declaration of a function. */
1934 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
);
1936 /* Generate HSA address operand for a given tree memory reference REF. If
1937 instructions need to be created to calculate the address, they will be added
1938 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
1939 the function assumes that the caller will handle possible
1940 bit-field references. Otherwise if we reference a bit-field, sorry message
1943 static hsa_op_address
*
1944 gen_hsa_addr (tree ref
, hsa_bb
*hbb
, HOST_WIDE_INT
*output_bitsize
= NULL
,
1945 HOST_WIDE_INT
*output_bitpos
= NULL
)
1947 hsa_symbol
*symbol
= NULL
;
1948 hsa_op_reg
*reg
= NULL
;
1949 offset_int offset
= 0;
1951 tree varoffset
= NULL_TREE
;
1952 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
1953 HOST_WIDE_INT bitsize
= 0, bitpos
= 0;
1954 BrigType16_t flat_addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
1956 if (TREE_CODE (ref
) == STRING_CST
)
1958 symbol
= hsa_get_string_cst_symbol (ref
);
1961 else if (TREE_CODE (ref
) == BIT_FIELD_REF
1962 && ((tree_to_uhwi (TREE_OPERAND (ref
, 1)) % BITS_PER_UNIT
) != 0
1963 || (tree_to_uhwi (TREE_OPERAND (ref
, 2)) % BITS_PER_UNIT
) != 0))
1965 HSA_SORRY_ATV (EXPR_LOCATION (origref
),
1966 "support for HSA does not implement "
1967 "bit field references such as %E", ref
);
1971 if (handled_component_p (ref
))
1974 int unsignedp
, volatilep
, preversep
;
1976 ref
= get_inner_reference (ref
, &bitsize
, &bitpos
, &varoffset
, &mode
,
1977 &unsignedp
, &preversep
, &volatilep
);
1980 offset
= wi::rshift (offset
, LOG2_BITS_PER_UNIT
, SIGNED
);
1983 switch (TREE_CODE (ref
))
1987 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
1988 symbol
= hsa_cfun
->create_hsa_temporary (flat_addrtype
);
1989 hsa_op_reg
*r
= new hsa_op_reg (flat_addrtype
);
1990 gen_hsa_addr_insns (ref
, r
, hbb
);
1991 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
1992 r
, new hsa_op_address (symbol
)));
1998 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
1999 hsa_op_with_type
*r
= hsa_cfun
->reg_for_gimple_ssa (ref
);
2000 if (r
->m_type
== BRIG_TYPE_B1
)
2001 r
= r
->get_in_type (BRIG_TYPE_U32
, hbb
);
2002 symbol
= hsa_cfun
->create_hsa_temporary (r
->m_type
);
2004 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
2005 r
, new hsa_op_address (symbol
)));
2013 gcc_assert (!symbol
);
2014 symbol
= get_symbol_for_decl (ref
);
2015 addrtype
= hsa_get_segment_addr_type (symbol
->m_segment
);
2019 process_mem_base (TREE_OPERAND (ref
, 0), &symbol
, &addrtype
, ®
,
2022 if (!integer_zerop (TREE_OPERAND (ref
, 1)))
2023 offset
+= wi::to_offset (TREE_OPERAND (ref
, 1));
2026 case TARGET_MEM_REF
:
2027 process_mem_base (TMR_BASE (ref
), &symbol
, &addrtype
, ®
, &offset
, hbb
);
2028 if (TMR_INDEX (ref
))
2031 hsa_op_base
*idx
= hsa_cfun
->reg_for_gimple_ssa
2032 (TMR_INDEX (ref
))->get_in_type (addrtype
, hbb
);
2033 if (TMR_STEP (ref
) && !integer_onep (TMR_STEP (ref
)))
2035 disp1
= new hsa_op_reg (addrtype
);
2036 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_MUL
,
2039 /* As step must respect addrtype, we overwrite the type
2040 of an immediate value. */
2041 hsa_op_immed
*step
= new hsa_op_immed (TMR_STEP (ref
));
2042 step
->m_type
= addrtype
;
2044 insn
->set_op (0, disp1
);
2045 insn
->set_op (1, idx
);
2046 insn
->set_op (2, step
);
2047 hbb
->append_insn (insn
);
2050 disp1
= as_a
<hsa_op_reg
*> (idx
);
2051 reg
= add_addr_regs_if_needed (reg
, disp1
, hbb
);
2053 if (TMR_INDEX2 (ref
))
2055 if (TREE_CODE (TMR_INDEX2 (ref
)) == SSA_NAME
)
2057 hsa_op_base
*disp2
= hsa_cfun
->reg_for_gimple_ssa
2058 (TMR_INDEX2 (ref
))->get_in_type (addrtype
, hbb
);
2059 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (disp2
),
2062 else if (TREE_CODE (TMR_INDEX2 (ref
)) == INTEGER_CST
)
2063 offset
+= wi::to_offset (TMR_INDEX2 (ref
));
2067 offset
+= wi::to_offset (TMR_OFFSET (ref
));
2070 HSA_SORRY_AT (EXPR_LOCATION (origref
),
2071 "support for HSA does not implement function pointers");
2074 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does "
2075 "not implement memory access to %E", origref
);
2081 if (TREE_CODE (varoffset
) == INTEGER_CST
)
2082 offset
+= wi::to_offset (varoffset
);
2085 hsa_op_base
*off_op
= gen_address_calculation (varoffset
, hbb
,
2087 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (off_op
),
2092 gcc_checking_assert ((symbol
2094 == hsa_get_segment_addr_type (symbol
->m_segment
))
2097 == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
)));
2099 HOST_WIDE_INT hwi_offset
= offset
.to_shwi ();
2101 /* Calculate remaining bitsize offset (if presented). */
2102 bitpos
%= BITS_PER_UNIT
;
2103 /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
2104 is not a reason to think this is a bit-field access. */
2106 && (bitsize
>= BITS_PER_UNIT
)
2107 && !(bitsize
& (bitsize
- 1)))
2110 if ((bitpos
|| bitsize
) && (output_bitpos
== NULL
|| output_bitsize
== NULL
))
2111 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does not "
2112 "implement unhandled bit field reference such as %E", ref
);
2114 if (output_bitsize
!= NULL
&& output_bitpos
!= NULL
)
2116 *output_bitsize
= bitsize
;
2117 *output_bitpos
= bitpos
;
2120 return new hsa_op_address (symbol
, reg
, hwi_offset
);
2123 /* Generate HSA address operand for a given tree memory reference REF. If
2124 instructions need to be created to calculate the address, they will be added
2125 to the end of HBB. OUTPUT_ALIGN is alignment of the created address. */
2127 static hsa_op_address
*
2128 gen_hsa_addr_with_align (tree ref
, hsa_bb
*hbb
, BrigAlignment8_t
*output_align
)
2130 hsa_op_address
*addr
= gen_hsa_addr (ref
, hbb
);
2131 if (addr
->m_reg
|| !addr
->m_symbol
)
2132 *output_align
= hsa_object_alignment (ref
);
2135 /* If the address consists only of a symbol and an offset, we
2136 compute the alignment ourselves to take into account any alignment
2137 promotions we might have done for the HSA symbol representation. */
2138 unsigned align
= hsa_byte_alignment (addr
->m_symbol
->m_align
);
2139 unsigned misalign
= addr
->m_imm_offset
& (align
- 1);
2141 align
= least_bit_hwi (misalign
);
2142 *output_align
= hsa_alignment_encoding (BITS_PER_UNIT
* align
);
2147 /* Generate HSA address for a function call argument of given TYPE.
2148 INDEX is used to generate corresponding name of the arguments.
2149 Special value -1 represents fact that result value is created. */
2151 static hsa_op_address
*
2152 gen_hsa_addr_for_arg (tree tree_type
, int index
)
2154 hsa_symbol
*sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
2156 sym
->m_type
= hsa_type_for_tree_type (tree_type
, &sym
->m_dim
);
2158 if (index
== -1) /* Function result. */
2159 sym
->m_name
= "res";
2160 else /* Function call arguments. */
2163 sym
->m_name_number
= index
;
2166 return new hsa_op_address (sym
);
2169 /* Generate HSA instructions that process all necessary conversions
2170 of an ADDR to flat addressing and place the result into DEST.
2171 Instructions are appended to HBB. */
2174 convert_addr_to_flat_segment (hsa_op_address
*addr
, hsa_op_reg
*dest
,
2177 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_LDA
);
2178 insn
->set_op (1, addr
);
2179 if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
!= BRIG_SEGMENT_GLOBAL
)
2181 /* LDA produces segment-relative address, we need to convert
2182 it to the flat one. */
2184 tmp
= new hsa_op_reg (hsa_get_segment_addr_type
2185 (addr
->m_symbol
->m_segment
));
2187 seg
= new hsa_insn_seg (BRIG_OPCODE_STOF
,
2188 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
2189 tmp
->m_type
, addr
->m_symbol
->m_segment
, dest
,
2192 insn
->set_op (0, tmp
);
2193 insn
->m_type
= tmp
->m_type
;
2194 hbb
->append_insn (insn
);
2195 hbb
->append_insn (seg
);
2199 insn
->set_op (0, dest
);
2200 insn
->m_type
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2201 hbb
->append_insn (insn
);
2205 /* Generate HSA instructions that calculate address of VAL including all
2206 necessary conversions to flat addressing and place the result into DEST.
2207 Instructions are appended to HBB. */
2210 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
)
2212 /* Handle cases like tmp = NULL, where we just emit a move instruction
2214 if (TREE_CODE (val
) == INTEGER_CST
)
2216 hsa_op_immed
*c
= new hsa_op_immed (val
);
2217 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2218 dest
->m_type
, dest
, c
);
2219 hbb
->append_insn (insn
);
2223 hsa_op_address
*addr
;
2225 gcc_assert (dest
->m_type
== hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2226 if (TREE_CODE (val
) == ADDR_EXPR
)
2227 val
= TREE_OPERAND (val
, 0);
2228 addr
= gen_hsa_addr (val
, hbb
);
2230 if (TREE_CODE (val
) == CONST_DECL
2231 && is_gimple_reg_type (TREE_TYPE (val
)))
2233 gcc_assert (addr
->m_symbol
2234 && addr
->m_symbol
->m_segment
== BRIG_SEGMENT_READONLY
);
2235 /* CONST_DECLs are in readonly segment which however does not have
2236 addresses convertible to flat segments. So copy it to a private one
2237 and take address of that. */
2238 BrigType16_t csttype
2239 = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (val
),
2241 hsa_op_reg
*r
= new hsa_op_reg (csttype
);
2242 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD
, csttype
, r
,
2243 new hsa_op_address (addr
->m_symbol
)));
2244 hsa_symbol
*copysym
= hsa_cfun
->create_hsa_temporary (csttype
);
2245 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, csttype
, r
,
2246 new hsa_op_address (copysym
)));
2247 addr
->m_symbol
= copysym
;
2249 else if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
== BRIG_SEGMENT_READONLY
)
2251 HSA_SORRY_ATV (EXPR_LOCATION (val
), "support for HSA does "
2252 "not implement taking addresses of complex "
2253 "CONST_DECLs such as %E", val
);
2258 convert_addr_to_flat_segment (addr
, dest
, hbb
);
2261 /* Return an HSA register or HSA immediate value operand corresponding to
2262 gimple operand OP. */
2264 static hsa_op_with_type
*
2265 hsa_reg_or_immed_for_gimple_op (tree op
, hsa_bb
*hbb
)
2269 if (TREE_CODE (op
) == SSA_NAME
)
2270 tmp
= hsa_cfun
->reg_for_gimple_ssa (op
);
2271 else if (!POINTER_TYPE_P (TREE_TYPE (op
)))
2272 return new hsa_op_immed (op
);
2275 tmp
= new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2276 gen_hsa_addr_insns (op
, tmp
, hbb
);
2281 /* Create a simple movement instruction with register destination DEST and
2282 register or immediate source SRC and append it to the end of HBB. */
2285 hsa_build_append_simple_mov (hsa_op_reg
*dest
, hsa_op_base
*src
, hsa_bb
*hbb
)
2287 /* Moves of packed data between registers need to adhere to the same type
2288 rules like when dealing with memory. */
2289 BrigType16_t tp
= mem_type_for_type (dest
->m_type
);
2290 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, tp
, dest
, src
);
2291 hsa_fixup_mov_insn_type (insn
);
2292 unsigned dest_size
= hsa_type_bit_size (dest
->m_type
);
2293 if (hsa_op_reg
*sreg
= dyn_cast
<hsa_op_reg
*> (src
))
2294 gcc_assert (dest_size
== hsa_type_bit_size (sreg
->m_type
));
2298 = hsa_type_bit_size (as_a
<hsa_op_immed
*> (src
)->m_type
);
2299 gcc_assert ((dest_size
== imm_size
)
2300 /* Eventually < 32bit registers will be promoted to 32bit. */
2301 || (dest_size
< 32 && imm_size
== 32));
2303 hbb
->append_insn (insn
);
2306 /* Generate HSAIL instructions loading a bit field into register DEST.
2307 VALUE_REG is a register of a SSA name that is used in the bit field
2308 reference. To identify a bit field BITPOS is offset to the loaded memory
2309 and BITSIZE is number of bits of the bit field.
2310 Add instructions to HBB. */
2313 gen_hsa_insns_for_bitfield (hsa_op_reg
*dest
, hsa_op_reg
*value_reg
,
2314 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2317 unsigned type_bitsize
2318 = hsa_type_bit_size (hsa_extend_inttype_to_32bit (dest
->m_type
));
2319 unsigned left_shift
= type_bitsize
- (bitsize
+ bitpos
);
2320 unsigned right_shift
= left_shift
+ bitpos
;
2324 hsa_op_reg
*value_reg_2
2325 = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest
->m_type
));
2326 hsa_op_immed
*c
= new hsa_op_immed (left_shift
, BRIG_TYPE_U32
);
2328 hsa_insn_basic
*lshift
2329 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, value_reg_2
->m_type
,
2330 value_reg_2
, value_reg
, c
);
2332 hbb
->append_insn (lshift
);
2334 value_reg
= value_reg_2
;
2339 hsa_op_reg
*value_reg_2
2340 = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest
->m_type
));
2341 hsa_op_immed
*c
= new hsa_op_immed (right_shift
, BRIG_TYPE_U32
);
2343 hsa_insn_basic
*rshift
2344 = new hsa_insn_basic (3, BRIG_OPCODE_SHR
, value_reg_2
->m_type
,
2345 value_reg_2
, value_reg
, c
);
2347 hbb
->append_insn (rshift
);
2349 value_reg
= value_reg_2
;
2352 hsa_insn_basic
*assignment
2353 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, NULL
, value_reg
);
2354 hsa_fixup_mov_insn_type (assignment
);
2355 hbb
->append_insn (assignment
);
2356 assignment
->set_output_in_type (dest
, 0, hbb
);
2360 /* Generate HSAIL instructions loading a bit field into register DEST. ADDR is
2361 prepared memory address which is used to load the bit field. To identify a
2362 bit field BITPOS is offset to the loaded memory and BITSIZE is number of
2363 bits of the bit field. Add instructions to HBB. Load must be performed in
2367 gen_hsa_insns_for_bitfield_load (hsa_op_reg
*dest
, hsa_op_address
*addr
,
2368 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2369 hsa_bb
*hbb
, BrigAlignment8_t align
)
2371 hsa_op_reg
*value_reg
= new hsa_op_reg (dest
->m_type
);
2373 = new hsa_insn_mem (BRIG_OPCODE_LD
,
2374 hsa_extend_inttype_to_32bit (dest
->m_type
),
2376 mem
->set_align (align
);
2377 hbb
->append_insn (mem
);
2378 gen_hsa_insns_for_bitfield (dest
, value_reg
, bitsize
, bitpos
, hbb
);
2381 /* Return the alignment of base memory accesses we issue to perform bit-field
2382 memory access REF. */
2384 static BrigAlignment8_t
2385 hsa_bitmemref_alignment (tree ref
)
2387 unsigned HOST_WIDE_INT bit_offset
= 0;
2391 if (TREE_CODE (ref
) == BIT_FIELD_REF
)
2393 if (!tree_fits_uhwi_p (TREE_OPERAND (ref
, 2)))
2394 return BRIG_ALIGNMENT_1
;
2395 bit_offset
+= tree_to_uhwi (TREE_OPERAND (ref
, 2));
2397 else if (TREE_CODE (ref
) == COMPONENT_REF
2398 && DECL_BIT_FIELD (TREE_OPERAND (ref
, 1)))
2399 bit_offset
+= int_bit_position (TREE_OPERAND (ref
, 1));
2402 ref
= TREE_OPERAND (ref
, 0);
2405 unsigned HOST_WIDE_INT bits
= bit_offset
% BITS_PER_UNIT
;
2406 unsigned HOST_WIDE_INT byte_bits
= bit_offset
- bits
;
2407 BrigAlignment8_t base
= hsa_object_alignment (ref
);
2410 return MIN (base
, hsa_alignment_encoding (least_bit_hwi (byte_bits
)));
2413 /* Generate HSAIL instructions loading something into register DEST. RHS is
2414 tree representation of the loaded data, which are loaded as type TYPE. Add
2415 instructions to HBB. */
2418 gen_hsa_insns_for_load (hsa_op_reg
*dest
, tree rhs
, tree type
, hsa_bb
*hbb
)
2420 /* The destination SSA name will give us the type. */
2421 if (TREE_CODE (rhs
) == VIEW_CONVERT_EXPR
)
2422 rhs
= TREE_OPERAND (rhs
, 0);
2424 if (TREE_CODE (rhs
) == SSA_NAME
)
2426 hsa_op_reg
*src
= hsa_cfun
->reg_for_gimple_ssa (rhs
);
2427 hsa_build_append_simple_mov (dest
, src
, hbb
);
2429 else if (is_gimple_min_invariant (rhs
)
2430 || TREE_CODE (rhs
) == ADDR_EXPR
)
2432 if (POINTER_TYPE_P (TREE_TYPE (rhs
)))
2434 if (dest
->m_type
!= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
))
2436 HSA_SORRY_ATV (EXPR_LOCATION (rhs
),
2437 "support for HSA does not implement conversion "
2438 "of %E to the requested non-pointer type.", rhs
);
2442 gen_hsa_addr_insns (rhs
, dest
, hbb
);
2444 else if (TREE_CODE (rhs
) == COMPLEX_CST
)
2446 hsa_op_immed
*real_part
= new hsa_op_immed (TREE_REALPART (rhs
));
2447 hsa_op_immed
*imag_part
= new hsa_op_immed (TREE_IMAGPART (rhs
));
2449 hsa_op_reg
*real_part_reg
2450 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2452 hsa_op_reg
*imag_part_reg
2453 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2456 hsa_build_append_simple_mov (real_part_reg
, real_part
, hbb
);
2457 hsa_build_append_simple_mov (imag_part_reg
, imag_part
, hbb
);
2459 BrigType16_t src_type
= hsa_bittype_for_type (real_part_reg
->m_type
);
2461 hsa_insn_packed
*insn
2462 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
,
2463 src_type
, dest
, real_part_reg
,
2465 hbb
->append_insn (insn
);
2469 hsa_op_immed
*imm
= new hsa_op_immed (rhs
);
2470 hsa_build_append_simple_mov (dest
, imm
, hbb
);
2473 else if (TREE_CODE (rhs
) == REALPART_EXPR
|| TREE_CODE (rhs
) == IMAGPART_EXPR
)
2475 tree pack_type
= TREE_TYPE (TREE_OPERAND (rhs
, 0));
2477 hsa_op_reg
*packed_reg
2478 = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type
, true));
2480 tree complex_rhs
= TREE_OPERAND (rhs
, 0);
2481 gen_hsa_insns_for_load (packed_reg
, complex_rhs
, TREE_TYPE (complex_rhs
),
2484 hsa_op_reg
*real_reg
2485 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2487 hsa_op_reg
*imag_reg
2488 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2490 BrigKind16_t brig_type
= packed_reg
->m_type
;
2491 hsa_insn_packed
*packed
2492 = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND
,
2493 hsa_bittype_for_type (real_reg
->m_type
),
2494 brig_type
, real_reg
, imag_reg
, packed_reg
);
2496 hbb
->append_insn (packed
);
2498 hsa_op_reg
*source
= TREE_CODE (rhs
) == REALPART_EXPR
?
2499 real_reg
: imag_reg
;
2501 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2502 dest
->m_type
, NULL
, source
);
2503 hsa_fixup_mov_insn_type (insn
);
2504 hbb
->append_insn (insn
);
2505 insn
->set_output_in_type (dest
, 0, hbb
);
2507 else if (TREE_CODE (rhs
) == BIT_FIELD_REF
2508 && TREE_CODE (TREE_OPERAND (rhs
, 0)) == SSA_NAME
)
2510 tree ssa_name
= TREE_OPERAND (rhs
, 0);
2511 HOST_WIDE_INT bitsize
= tree_to_uhwi (TREE_OPERAND (rhs
, 1));
2512 HOST_WIDE_INT bitpos
= tree_to_uhwi (TREE_OPERAND (rhs
, 2));
2514 hsa_op_reg
*imm_value
= hsa_cfun
->reg_for_gimple_ssa (ssa_name
);
2515 gen_hsa_insns_for_bitfield (dest
, imm_value
, bitsize
, bitpos
, hbb
);
2517 else if (DECL_P (rhs
) || TREE_CODE (rhs
) == MEM_REF
2518 || TREE_CODE (rhs
) == TARGET_MEM_REF
2519 || handled_component_p (rhs
))
2521 HOST_WIDE_INT bitsize
, bitpos
;
2523 /* Load from memory. */
2524 hsa_op_address
*addr
;
2525 addr
= gen_hsa_addr (rhs
, hbb
, &bitsize
, &bitpos
);
2527 /* Handle load of a bit field. */
2530 HSA_SORRY_AT (EXPR_LOCATION (rhs
),
2531 "support for HSA does not implement load from a bit "
2532 "field bigger than 64 bits");
2536 if (bitsize
|| bitpos
)
2537 gen_hsa_insns_for_bitfield_load (dest
, addr
, bitsize
, bitpos
, hbb
,
2538 hsa_bitmemref_alignment (rhs
));
2542 /* Not dest->m_type, that's possibly extended. */
2543 mtype
= mem_type_for_type (hsa_type_for_scalar_tree_type (type
,
2545 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
, dest
,
2547 mem
->set_align (hsa_object_alignment (rhs
));
2548 hbb
->append_insn (mem
);
2552 HSA_SORRY_ATV (EXPR_LOCATION (rhs
),
2553 "support for HSA does not implement loading "
2558 /* Return number of bits necessary for representation of a bit field,
2559 starting at BITPOS with size of BITSIZE. */
2562 get_bitfield_size (unsigned bitpos
, unsigned bitsize
)
2564 unsigned s
= bitpos
+ bitsize
;
2565 unsigned sizes
[] = {8, 16, 32, 64};
2567 for (unsigned i
= 0; i
< 4; i
++)
2575 /* Generate HSAIL instructions storing into memory. LHS is the destination of
2576 the store, SRC is the source operand. Add instructions to HBB. */
2579 gen_hsa_insns_for_store (tree lhs
, hsa_op_base
*src
, hsa_bb
*hbb
)
2581 HOST_WIDE_INT bitsize
= 0, bitpos
= 0;
2582 BrigAlignment8_t req_align
;
2584 mtype
= mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
2586 hsa_op_address
*addr
;
2587 addr
= gen_hsa_addr (lhs
, hbb
, &bitsize
, &bitpos
);
2589 /* Handle store to a bit field. */
2592 HSA_SORRY_AT (EXPR_LOCATION (lhs
),
2593 "support for HSA does not implement store to a bit field "
2594 "bigger than 64 bits");
2598 unsigned type_bitsize
= get_bitfield_size (bitpos
, bitsize
);
2600 /* HSAIL does not support MOV insn with 16-bits integers. */
2601 if (type_bitsize
< 32)
2604 if (bitpos
|| (bitsize
&& type_bitsize
!= bitsize
))
2606 unsigned HOST_WIDE_INT mask
= 0;
2607 BrigType16_t mem_type
2608 = get_integer_type_by_bytes (type_bitsize
/ BITS_PER_UNIT
,
2609 !TYPE_UNSIGNED (TREE_TYPE (lhs
)));
2611 for (unsigned i
= 0; i
< type_bitsize
; i
++)
2612 if (i
< bitpos
|| i
>= bitpos
+ bitsize
)
2613 mask
|= ((unsigned HOST_WIDE_INT
)1 << i
);
2615 hsa_op_reg
*value_reg
= new hsa_op_reg (mem_type
);
2617 req_align
= hsa_bitmemref_alignment (lhs
);
2618 /* Load value from memory. */
2619 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mem_type
,
2621 mem
->set_align (req_align
);
2622 hbb
->append_insn (mem
);
2624 /* AND the loaded value with prepared mask. */
2625 hsa_op_reg
*cleared_reg
= new hsa_op_reg (mem_type
);
2628 = get_integer_type_by_bytes (type_bitsize
/ BITS_PER_UNIT
, false);
2629 hsa_op_immed
*c
= new hsa_op_immed (mask
, t
);
2631 hsa_insn_basic
*clearing
2632 = new hsa_insn_basic (3, BRIG_OPCODE_AND
, mem_type
, cleared_reg
,
2634 hbb
->append_insn (clearing
);
2636 /* Shift to left a value that is going to be stored. */
2637 hsa_op_reg
*new_value_reg
= new hsa_op_reg (mem_type
);
2639 hsa_insn_basic
*basic
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, mem_type
,
2640 new_value_reg
, src
);
2641 hsa_fixup_mov_insn_type (basic
);
2642 hbb
->append_insn (basic
);
2646 hsa_op_reg
*shifted_value_reg
= new hsa_op_reg (mem_type
);
2647 c
= new hsa_op_immed (bitpos
, BRIG_TYPE_U32
);
2649 hsa_insn_basic
*basic
2650 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, mem_type
,
2651 shifted_value_reg
, new_value_reg
, c
);
2652 hbb
->append_insn (basic
);
2654 new_value_reg
= shifted_value_reg
;
2657 /* OR the prepared value with prepared chunk loaded from memory. */
2658 hsa_op_reg
*prepared_reg
= new hsa_op_reg (mem_type
);
2659 basic
= new hsa_insn_basic (3, BRIG_OPCODE_OR
, mem_type
, prepared_reg
,
2660 new_value_reg
, cleared_reg
);
2661 hbb
->append_insn (basic
);
2667 req_align
= hsa_object_alignment (lhs
);
2669 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
, addr
);
2670 mem
->set_align (req_align
);
2672 /* The HSAIL verifier has another constraint: if the source is an immediate
2673 then it must match the destination type. If it's a register the low bits
2674 will be used for sub-word stores. We're always allocating new operands so
2675 we can modify the above in place. */
2676 if (hsa_op_immed
*imm
= dyn_cast
<hsa_op_immed
*> (src
))
2678 if (!hsa_type_packed_p (imm
->m_type
))
2679 imm
->m_type
= mem
->m_type
;
2682 /* ...and all vector immediates apparently need to be vectors of
2684 unsigned bs
= hsa_type_bit_size (imm
->m_type
);
2685 gcc_assert (bs
== hsa_type_bit_size (mem
->m_type
));
2689 imm
->m_type
= BRIG_TYPE_U8X4
;
2692 imm
->m_type
= BRIG_TYPE_U8X8
;
2695 imm
->m_type
= BRIG_TYPE_U8X16
;
2703 hbb
->append_insn (mem
);
2706 /* Generate memory copy instructions that are going to be used
2707 for copying a SRC memory to TARGET memory,
2708 represented by pointer in a register. MIN_ALIGN is minimal alignment
2709 of provided HSA addresses. */
2712 gen_hsa_memory_copy (hsa_bb
*hbb
, hsa_op_address
*target
, hsa_op_address
*src
,
2713 unsigned size
, BrigAlignment8_t min_align
)
2715 hsa_op_address
*addr
;
2718 unsigned offset
= 0;
2719 unsigned min_byte_align
= hsa_byte_alignment (min_align
);
2733 if (s
> min_byte_align
)
2736 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2738 hsa_op_reg
*tmp
= new hsa_op_reg (t
);
2739 addr
= new hsa_op_address (src
->m_symbol
, src
->m_reg
,
2740 src
->m_imm_offset
+ offset
);
2741 mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, t
, tmp
, addr
);
2742 hbb
->append_insn (mem
);
2744 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2745 target
->m_imm_offset
+ offset
);
2746 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, tmp
, addr
);
2747 hbb
->append_insn (mem
);
2753 /* Create a memset mask that is created by copying a CONSTANT byte value
2754 to an integer of BYTE_SIZE bytes. */
2756 static unsigned HOST_WIDE_INT
2757 build_memset_value (unsigned HOST_WIDE_INT constant
, unsigned byte_size
)
2762 HOST_WIDE_INT v
= constant
;
2764 for (unsigned i
= 1; i
< byte_size
; i
++)
2765 v
|= constant
<< (8 * i
);
2770 /* Generate memory set instructions that are going to be used
2771 for setting a CONSTANT byte value to TARGET memory of SIZE bytes.
2772 MIN_ALIGN is minimal alignment of provided HSA addresses. */
2775 gen_hsa_memory_set (hsa_bb
*hbb
, hsa_op_address
*target
,
2776 unsigned HOST_WIDE_INT constant
,
2777 unsigned size
, BrigAlignment8_t min_align
)
2779 hsa_op_address
*addr
;
2782 unsigned offset
= 0;
2783 unsigned min_byte_align
= hsa_byte_alignment (min_align
);
2797 if (s
> min_byte_align
)
2800 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2801 target
->m_imm_offset
+ offset
);
2803 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2804 HOST_WIDE_INT c
= build_memset_value (constant
, s
);
2806 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, new hsa_op_immed (c
, t
),
2808 hbb
->append_insn (mem
);
2814 /* Generate HSAIL instructions for a single assignment
2815 of an empty constructor to an ADDR_LHS. Constructor is passed as a
2816 tree RHS and all instructions are appended to HBB. ALIGN is
2817 alignment of the address. */
2820 gen_hsa_ctor_assignment (hsa_op_address
*addr_lhs
, tree rhs
, hsa_bb
*hbb
,
2821 BrigAlignment8_t align
)
2823 if (CONSTRUCTOR_NELTS (rhs
))
2825 HSA_SORRY_AT (EXPR_LOCATION (rhs
),
2826 "support for HSA does not implement load from constructor");
2830 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2831 gen_hsa_memory_set (hbb
, addr_lhs
, 0, size
, align
);
2834 /* Generate HSA instructions for a single assignment of RHS to LHS.
2835 HBB is the basic block they will be appended to. */
2838 gen_hsa_insns_for_single_assignment (tree lhs
, tree rhs
, hsa_bb
*hbb
)
2840 if (TREE_CODE (lhs
) == SSA_NAME
)
2842 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
2843 if (hsa_seen_error ())
2846 gen_hsa_insns_for_load (dest
, rhs
, TREE_TYPE (lhs
), hbb
);
2848 else if (TREE_CODE (rhs
) == SSA_NAME
2849 || (is_gimple_min_invariant (rhs
) && TREE_CODE (rhs
) != STRING_CST
))
2851 /* Store to memory. */
2852 hsa_op_base
*src
= hsa_reg_or_immed_for_gimple_op (rhs
, hbb
);
2853 if (hsa_seen_error ())
2856 gen_hsa_insns_for_store (lhs
, src
, hbb
);
2860 BrigAlignment8_t lhs_align
;
2861 hsa_op_address
*addr_lhs
= gen_hsa_addr_with_align (lhs
, hbb
,
2864 if (TREE_CODE (rhs
) == CONSTRUCTOR
)
2865 gen_hsa_ctor_assignment (addr_lhs
, rhs
, hbb
, lhs_align
);
2868 BrigAlignment8_t rhs_align
;
2869 hsa_op_address
*addr_rhs
= gen_hsa_addr_with_align (rhs
, hbb
,
2872 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2873 gen_hsa_memory_copy (hbb
, addr_lhs
, addr_rhs
, size
,
2874 MIN (lhs_align
, rhs_align
));
2879 /* Prepend before INSN a load from spill symbol of SPILL_REG. Return the
2880 register into which we loaded. If this required another register to convert
2881 from a B1 type, return it in *PTMP2, otherwise store NULL into it. We
2882 assume we are out of SSA so the returned register does not have its
2886 hsa_spill_in (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2888 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2889 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2890 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2892 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, spill_sym
->m_type
,
2894 hsa_insert_insn_before (mem
, insn
);
2897 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2899 hsa_insn_basic
*cvtinsn
;
2901 reg
= new hsa_op_reg (spill_reg
->m_type
);
2903 cvtinsn
= new hsa_insn_cvt (reg
, *ptmp2
);
2904 hsa_insert_insn_before (cvtinsn
, insn
);
2909 /* Append after INSN a store to spill symbol of SPILL_REG. Return the register
2910 from which we stored. If this required another register to convert to a B1
2911 type, return it in *PTMP2, otherwise store NULL into it. We assume we are
2912 out of SSA so the returned register does not have its use updated. */
2915 hsa_spill_out (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2917 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2918 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2919 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2920 hsa_op_reg
*returnreg
;
2924 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2926 hsa_insn_basic
*cvtinsn
;
2927 *ptmp2
= new hsa_op_reg (spill_sym
->m_type
);
2928 reg
->m_type
= spill_reg
->m_type
;
2930 cvtinsn
= new hsa_insn_cvt (*ptmp2
, returnreg
);
2931 hsa_append_insn_after (cvtinsn
, insn
);
2936 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, spill_sym
->m_type
, reg
,
2938 hsa_append_insn_after (mem
, insn
);
2942 /* Generate a comparison instruction that will compare LHS and RHS with
2943 comparison specified by CODE and put result into register DEST. DEST has to
2944 have its type set already but must not have its definition set yet.
2945 Generated instructions will be added to HBB. */
2948 gen_hsa_cmp_insn_from_gimple (enum tree_code code
, tree lhs
, tree rhs
,
2949 hsa_op_reg
*dest
, hsa_bb
*hbb
)
2951 BrigCompareOperation8_t compare
;
2956 compare
= BRIG_COMPARE_LT
;
2959 compare
= BRIG_COMPARE_LE
;
2962 compare
= BRIG_COMPARE_GT
;
2965 compare
= BRIG_COMPARE_GE
;
2968 compare
= BRIG_COMPARE_EQ
;
2971 compare
= BRIG_COMPARE_NE
;
2973 case UNORDERED_EXPR
:
2974 compare
= BRIG_COMPARE_NAN
;
2977 compare
= BRIG_COMPARE_NUM
;
2980 compare
= BRIG_COMPARE_LTU
;
2983 compare
= BRIG_COMPARE_LEU
;
2986 compare
= BRIG_COMPARE_GTU
;
2989 compare
= BRIG_COMPARE_GEU
;
2992 compare
= BRIG_COMPARE_EQU
;
2995 compare
= BRIG_COMPARE_NEU
;
2999 HSA_SORRY_ATV (EXPR_LOCATION (lhs
),
3000 "support for HSA does not implement comparison tree "
3001 "code %s\n", get_tree_code_name (code
));
3005 /* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer)
3006 as a result of comparison. */
3008 BrigType16_t dest_type
= hsa_type_integer_p (dest
->m_type
)
3009 ? (BrigType16_t
) BRIG_TYPE_B1
: dest
->m_type
;
3011 hsa_insn_cmp
*cmp
= new hsa_insn_cmp (compare
, dest_type
);
3012 hsa_op_with_type
*op1
= hsa_reg_or_immed_for_gimple_op (lhs
, hbb
);
3013 cmp
->set_op (1, op1
->extend_int_to_32bit (hbb
));
3014 hsa_op_with_type
*op2
= hsa_reg_or_immed_for_gimple_op (rhs
, hbb
);
3015 cmp
->set_op (2, op2
->extend_int_to_32bit (hbb
));
3017 hbb
->append_insn (cmp
);
3018 cmp
->set_output_in_type (dest
, 0, hbb
);
3021 /* Generate an unary instruction with OPCODE and append it to a basic block
3022 HBB. The instruction uses DEST as a destination and OP1
3023 as a single operand. */
3026 gen_hsa_unary_operation (BrigOpcode opcode
, hsa_op_reg
*dest
,
3027 hsa_op_with_type
*op1
, hsa_bb
*hbb
)
3029 gcc_checking_assert (dest
);
3030 hsa_insn_basic
*insn
;
3032 if (opcode
== BRIG_OPCODE_MOV
&& hsa_needs_cvt (dest
->m_type
, op1
->m_type
))
3034 insn
= new hsa_insn_cvt (dest
, op1
);
3035 hbb
->append_insn (insn
);
3039 op1
= op1
->extend_int_to_32bit (hbb
);
3040 if (opcode
== BRIG_OPCODE_FIRSTBIT
|| opcode
== BRIG_OPCODE_LASTBIT
)
3042 BrigType16_t srctype
= hsa_type_integer_p (op1
->m_type
) ? op1
->m_type
3043 : hsa_unsigned_type_for_type (op1
->m_type
);
3044 insn
= new hsa_insn_srctype (2, opcode
, BRIG_TYPE_U32
, srctype
, NULL
,
3049 BrigType16_t optype
= hsa_extend_inttype_to_32bit (dest
->m_type
);
3050 insn
= new hsa_insn_basic (2, opcode
, optype
, NULL
, op1
);
3052 if (opcode
== BRIG_OPCODE_MOV
)
3053 hsa_fixup_mov_insn_type (insn
);
3054 else if (opcode
== BRIG_OPCODE_ABS
|| opcode
== BRIG_OPCODE_NEG
)
3056 /* ABS and NEG only exist in _s form :-/ */
3057 if (insn
->m_type
== BRIG_TYPE_U32
)
3058 insn
->m_type
= BRIG_TYPE_S32
;
3059 else if (insn
->m_type
== BRIG_TYPE_U64
)
3060 insn
->m_type
= BRIG_TYPE_S64
;
3064 hbb
->append_insn (insn
);
3065 insn
->set_output_in_type (dest
, 0, hbb
);
3068 /* Generate a binary instruction with OPCODE and append it to a basic block
3069 HBB. The instruction uses DEST as a destination and operands OP1
3073 gen_hsa_binary_operation (int opcode
, hsa_op_reg
*dest
,
3074 hsa_op_with_type
*op1
, hsa_op_with_type
*op2
,
3077 gcc_checking_assert (dest
);
3079 BrigType16_t optype
= hsa_extend_inttype_to_32bit (dest
->m_type
);
3080 op1
= op1
->extend_int_to_32bit (hbb
);
3081 op2
= op2
->extend_int_to_32bit (hbb
);
3083 if ((opcode
== BRIG_OPCODE_SHL
|| opcode
== BRIG_OPCODE_SHR
)
3084 && is_a
<hsa_op_immed
*> (op2
))
3086 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3087 i
->set_type (BRIG_TYPE_U32
);
3089 if ((opcode
== BRIG_OPCODE_OR
3090 || opcode
== BRIG_OPCODE_XOR
3091 || opcode
== BRIG_OPCODE_AND
)
3092 && is_a
<hsa_op_immed
*> (op2
))
3094 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3095 i
->set_type (hsa_unsigned_type_for_type (i
->m_type
));
3098 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, optype
, NULL
,
3100 hbb
->append_insn (insn
);
3101 insn
->set_output_in_type (dest
, 0, hbb
);
3104 /* Generate HSA instructions for a single assignment. HBB is the basic block
3105 they will be appended to. */
3108 gen_hsa_insns_for_operation_assignment (gimple
*assign
, hsa_bb
*hbb
)
3110 tree_code code
= gimple_assign_rhs_code (assign
);
3111 gimple_rhs_class rhs_class
= get_gimple_rhs_class (gimple_expr_code (assign
));
3113 tree lhs
= gimple_assign_lhs (assign
);
3114 tree rhs1
= gimple_assign_rhs1 (assign
);
3115 tree rhs2
= gimple_assign_rhs2 (assign
);
3116 tree rhs3
= gimple_assign_rhs3 (assign
);
3124 /* The opcode is changed to BRIG_OPCODE_CVT if BRIG types
3125 needs a conversion. */
3126 opcode
= BRIG_OPCODE_MOV
;
3130 case POINTER_PLUS_EXPR
:
3131 opcode
= BRIG_OPCODE_ADD
;
3134 opcode
= BRIG_OPCODE_SUB
;
3137 opcode
= BRIG_OPCODE_MUL
;
3139 case MULT_HIGHPART_EXPR
:
3140 opcode
= BRIG_OPCODE_MULHI
;
3143 case TRUNC_DIV_EXPR
:
3144 case EXACT_DIV_EXPR
:
3145 opcode
= BRIG_OPCODE_DIV
;
3148 case FLOOR_DIV_EXPR
:
3149 case ROUND_DIV_EXPR
:
3150 HSA_SORRY_AT (gimple_location (assign
),
3151 "support for HSA does not implement CEIL_DIV_EXPR, "
3152 "FLOOR_DIV_EXPR or ROUND_DIV_EXPR");
3154 case TRUNC_MOD_EXPR
:
3155 opcode
= BRIG_OPCODE_REM
;
3158 case FLOOR_MOD_EXPR
:
3159 case ROUND_MOD_EXPR
:
3160 HSA_SORRY_AT (gimple_location (assign
),
3161 "support for HSA does not implement CEIL_MOD_EXPR, "
3162 "FLOOR_MOD_EXPR or ROUND_MOD_EXPR");
3165 opcode
= BRIG_OPCODE_NEG
;
3168 /* There is a native HSA instruction for scalar FMAs but not for vector
3170 if (TREE_CODE (TREE_TYPE (lhs
)) == VECTOR_TYPE
)
3173 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3174 hsa_op_with_type
*op1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3175 hsa_op_with_type
*op2
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3176 hsa_op_with_type
*op3
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3177 hsa_op_reg
*tmp
= new hsa_op_reg (dest
->m_type
);
3178 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp
, op1
, op2
, hbb
);
3179 gen_hsa_binary_operation (BRIG_OPCODE_ADD
, dest
, tmp
, op3
, hbb
);
3182 opcode
= BRIG_OPCODE_MAD
;
3185 opcode
= BRIG_OPCODE_MIN
;
3188 opcode
= BRIG_OPCODE_MAX
;
3191 opcode
= BRIG_OPCODE_ABS
;
3194 opcode
= BRIG_OPCODE_SHL
;
3197 opcode
= BRIG_OPCODE_SHR
;
3202 hsa_insn_basic
*insn
= NULL
;
3203 int code1
= code
== LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3204 int code2
= code
!= LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3205 BrigType16_t btype
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
3208 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3209 hsa_op_reg
*op1
= new hsa_op_reg (btype
);
3210 hsa_op_reg
*op2
= new hsa_op_reg (btype
);
3211 hsa_op_with_type
*shift1
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3213 tree type
= TREE_TYPE (rhs2
);
3214 unsigned HOST_WIDE_INT bitsize
= TREE_INT_CST_LOW (TYPE_SIZE (type
));
3216 hsa_op_with_type
*shift2
= NULL
;
3217 if (TREE_CODE (rhs2
) == INTEGER_CST
)
3218 shift2
= new hsa_op_immed (bitsize
- tree_to_uhwi (rhs2
),
3220 else if (TREE_CODE (rhs2
) == SSA_NAME
)
3222 hsa_op_reg
*s
= hsa_cfun
->reg_for_gimple_ssa (rhs2
);
3223 s
= as_a
<hsa_op_reg
*> (s
->extend_int_to_32bit (hbb
));
3224 hsa_op_reg
*d
= new hsa_op_reg (s
->m_type
);
3225 hsa_op_immed
*size_imm
= new hsa_op_immed (bitsize
, BRIG_TYPE_U32
);
3227 insn
= new hsa_insn_basic (3, BRIG_OPCODE_SUB
, d
->m_type
,
3229 hbb
->append_insn (insn
);
3236 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3237 gen_hsa_binary_operation (code1
, op1
, src
, shift1
, hbb
);
3238 gen_hsa_binary_operation (code2
, op2
, src
, shift2
, hbb
);
3239 gen_hsa_binary_operation (BRIG_OPCODE_OR
, dest
, op1
, op2
, hbb
);
3244 opcode
= BRIG_OPCODE_OR
;
3247 opcode
= BRIG_OPCODE_XOR
;
3250 opcode
= BRIG_OPCODE_AND
;
3253 opcode
= BRIG_OPCODE_NOT
;
3255 case FIX_TRUNC_EXPR
:
3257 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3258 hsa_op_with_type
*v
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3260 if (hsa_needs_cvt (dest
->m_type
, v
->m_type
))
3262 hsa_op_reg
*tmp
= new hsa_op_reg (v
->m_type
);
3264 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3265 tmp
->m_type
, tmp
, v
);
3266 hbb
->append_insn (insn
);
3268 hsa_insn_basic
*cvtinsn
= new hsa_insn_cvt (dest
, tmp
);
3269 hbb
->append_insn (cvtinsn
);
3273 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3274 dest
->m_type
, dest
, v
);
3275 hbb
->append_insn (insn
);
3280 opcode
= BRIG_OPCODE_TRUNC
;
3289 case UNORDERED_EXPR
:
3299 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3301 gen_hsa_cmp_insn_from_gimple (code
, rhs1
, rhs2
, dest
, hbb
);
3307 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3308 hsa_op_with_type
*ctrl
= NULL
;
3311 if (CONSTANT_CLASS_P (cond
) || TREE_CODE (cond
) == SSA_NAME
)
3312 ctrl
= hsa_reg_or_immed_for_gimple_op (cond
, hbb
);
3315 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
3317 gen_hsa_cmp_insn_from_gimple (TREE_CODE (cond
),
3318 TREE_OPERAND (cond
, 0),
3319 TREE_OPERAND (cond
, 1),
3325 hsa_op_with_type
*op2
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3326 hsa_op_with_type
*op3
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3327 op2
= op2
->extend_int_to_32bit (hbb
);
3328 op3
= op3
->extend_int_to_32bit (hbb
);
3330 BrigType16_t type
= hsa_extend_inttype_to_32bit (dest
->m_type
);
3331 BrigType16_t utype
= hsa_unsigned_type_for_type (type
);
3332 if (is_a
<hsa_op_immed
*> (op2
))
3333 op2
->m_type
= utype
;
3334 if (is_a
<hsa_op_immed
*> (op3
))
3335 op3
->m_type
= utype
;
3337 hsa_insn_basic
*insn
3338 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
,
3339 hsa_bittype_for_type (type
),
3340 NULL
, ctrl
, op2
, op3
);
3342 hbb
->append_insn (insn
);
3343 insn
->set_output_in_type (dest
, 0, hbb
);
3349 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3350 hsa_op_with_type
*rhs1_reg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3351 rhs1_reg
= rhs1_reg
->extend_int_to_32bit (hbb
);
3352 hsa_op_with_type
*rhs2_reg
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3353 rhs2_reg
= rhs2_reg
->extend_int_to_32bit (hbb
);
3355 if (hsa_seen_error ())
3358 BrigType16_t src_type
= hsa_bittype_for_type (rhs1_reg
->m_type
);
3359 rhs1_reg
= rhs1_reg
->get_in_type (src_type
, hbb
);
3360 rhs2_reg
= rhs2_reg
->get_in_type (src_type
, hbb
);
3362 hsa_insn_packed
*insn
3363 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
, src_type
,
3364 dest
, rhs1_reg
, rhs2_reg
);
3365 hbb
->append_insn (insn
);
3370 /* Implement others as we come across them. */
3371 HSA_SORRY_ATV (gimple_location (assign
),
3372 "support for HSA does not implement operation %s",
3373 get_tree_code_name (code
));
3378 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3379 hsa_op_with_type
*op1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3380 hsa_op_with_type
*op2
3381 = rhs2
? hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
) : NULL
;
3383 if (hsa_seen_error ())
3388 case GIMPLE_TERNARY_RHS
:
3390 hsa_op_with_type
*op3
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3391 op3
= op3
->extend_int_to_32bit (hbb
);
3392 hsa_insn_basic
*insn
= new hsa_insn_basic (4, opcode
, dest
->m_type
, dest
,
3394 hbb
->append_insn (insn
);
3398 case GIMPLE_BINARY_RHS
:
3399 gen_hsa_binary_operation (opcode
, dest
, op1
, op2
, hbb
);
3402 case GIMPLE_UNARY_RHS
:
3403 gen_hsa_unary_operation (opcode
, dest
, op1
, hbb
);
3410 /* Generate HSA instructions for a given gimple condition statement COND.
3411 Instructions will be appended to HBB, which also needs to be the
3412 corresponding structure to the basic_block of COND. */
3415 gen_hsa_insns_for_cond_stmt (gimple
*cond
, hsa_bb
*hbb
)
3417 hsa_op_reg
*ctrl
= new hsa_op_reg (BRIG_TYPE_B1
);
3420 gen_hsa_cmp_insn_from_gimple (gimple_cond_code (cond
),
3421 gimple_cond_lhs (cond
),
3422 gimple_cond_rhs (cond
),
3425 cbr
= new hsa_insn_cbr (ctrl
);
3426 hbb
->append_insn (cbr
);
3429 /* Maximum number of elements in a jump table for an HSA SBR instruction. */
3431 #define HSA_MAXIMUM_SBR_LABELS 16
3433 /* Return lowest value of a switch S that is handled in a non-default
3437 get_switch_low (gswitch
*s
)
3439 unsigned labels
= gimple_switch_num_labels (s
);
3440 gcc_checking_assert (labels
>= 1);
3442 return CASE_LOW (gimple_switch_label (s
, 1));
3445 /* Return highest value of a switch S that is handled in a non-default
3449 get_switch_high (gswitch
*s
)
3451 unsigned labels
= gimple_switch_num_labels (s
);
3453 /* Compare last label to maximum number of labels. */
3454 tree label
= gimple_switch_label (s
, labels
- 1);
3455 tree low
= CASE_LOW (label
);
3456 tree high
= CASE_HIGH (label
);
3458 return high
!= NULL_TREE
? high
: low
;
3462 get_switch_size (gswitch
*s
)
3464 return int_const_binop (MINUS_EXPR
, get_switch_high (s
), get_switch_low (s
));
3467 /* Generate HSA instructions for a given gimple switch.
3468 Instructions will be appended to HBB. */
3471 gen_hsa_insns_for_switch_stmt (gswitch
*s
, hsa_bb
*hbb
)
3473 gimple_stmt_iterator it
= gsi_for_stmt (s
);
3476 /* Create preambule that verifies that index - lowest_label >= 0. */
3477 edge e
= split_block (hbb
->m_bb
, gsi_stmt (it
));
3478 e
->flags
&= ~EDGE_FALLTHRU
;
3479 e
->flags
|= EDGE_TRUE_VALUE
;
3481 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
3482 tree index_tree
= gimple_switch_index (s
);
3483 tree lowest
= get_switch_low (s
);
3484 tree highest
= get_switch_high (s
);
3486 hsa_op_reg
*index
= hsa_cfun
->reg_for_gimple_ssa (index_tree
);
3487 index
= as_a
<hsa_op_reg
*> (index
->extend_int_to_32bit (hbb
));
3489 hsa_op_reg
*cmp1_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3490 hsa_op_immed
*cmp1_immed
= new hsa_op_immed (lowest
, true);
3491 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_GE
, cmp1_reg
->m_type
,
3492 cmp1_reg
, index
, cmp1_immed
));
3494 hsa_op_reg
*cmp2_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3495 hsa_op_immed
*cmp2_immed
= new hsa_op_immed (highest
, true);
3496 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_LE
, cmp2_reg
->m_type
,
3497 cmp2_reg
, index
, cmp2_immed
));
3499 hsa_op_reg
*cmp_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3500 hbb
->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_AND
, cmp_reg
->m_type
,
3501 cmp_reg
, cmp1_reg
, cmp2_reg
));
3503 hbb
->append_insn (new hsa_insn_cbr (cmp_reg
));
3505 tree default_label
= gimple_switch_default_label (s
);
3506 basic_block default_label_bb
= label_to_block_fn (func
,
3507 CASE_LABEL (default_label
));
3509 if (!gimple_seq_empty_p (phi_nodes (default_label_bb
)))
3511 default_label_bb
= split_edge (find_edge (e
->dest
, default_label_bb
));
3512 hsa_init_new_bb (default_label_bb
);
3515 make_edge (e
->src
, default_label_bb
, EDGE_FALSE_VALUE
);
3517 hsa_cfun
->m_modified_cfg
= true;
3519 /* Basic block with the SBR instruction. */
3520 hbb
= hsa_init_new_bb (e
->dest
);
3522 hsa_op_reg
*sub_index
= new hsa_op_reg (index
->m_type
);
3523 hbb
->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB
, sub_index
->m_type
,
3525 new hsa_op_immed (lowest
, true)));
3527 hsa_op_base
*tmp
= sub_index
->get_in_type (BRIG_TYPE_U64
, hbb
);
3528 sub_index
= as_a
<hsa_op_reg
*> (tmp
);
3529 unsigned labels
= gimple_switch_num_labels (s
);
3530 unsigned HOST_WIDE_INT size
= tree_to_uhwi (get_switch_size (s
));
3532 hsa_insn_sbr
*sbr
= new hsa_insn_sbr (sub_index
, size
+ 1);
3534 /* Prepare array with default label destination. */
3535 for (unsigned HOST_WIDE_INT i
= 0; i
<= size
; i
++)
3536 sbr
->m_jump_table
.safe_push (default_label_bb
);
3538 /* Iterate all labels and fill up the jump table. */
3539 for (unsigned i
= 1; i
< labels
; i
++)
3541 tree label
= gimple_switch_label (s
, i
);
3542 basic_block bb
= label_to_block_fn (func
, CASE_LABEL (label
));
3544 unsigned HOST_WIDE_INT sub_low
3545 = tree_to_uhwi (int_const_binop (MINUS_EXPR
, CASE_LOW (label
), lowest
));
3547 unsigned HOST_WIDE_INT sub_high
= sub_low
;
3548 tree high
= CASE_HIGH (label
);
3550 sub_high
= tree_to_uhwi (int_const_binop (MINUS_EXPR
, high
, lowest
));
3552 for (unsigned HOST_WIDE_INT j
= sub_low
; j
<= sub_high
; j
++)
3553 sbr
->m_jump_table
[j
] = bb
;
3556 hbb
->append_insn (sbr
);
3559 /* Verify that the function DECL can be handled by HSA. */
3562 verify_function_arguments (tree decl
)
3564 tree type
= TREE_TYPE (decl
);
3565 if (DECL_STATIC_CHAIN (decl
))
3567 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3568 "HSA does not support nested functions: %qD", decl
);
3571 else if (!TYPE_ARG_TYPES (type
) || stdarg_p (type
))
3573 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3574 "HSA does not support functions with variadic arguments "
3575 "(or unknown return type): %qD", decl
);
3580 /* Return BRIG type for FORMAL_ARG_TYPE. If the formal argument type is NULL,
3581 return ACTUAL_ARG_TYPE. */
3584 get_format_argument_type (tree formal_arg_type
, BrigType16_t actual_arg_type
)
3586 if (formal_arg_type
== NULL
)
3587 return actual_arg_type
;
3589 BrigType16_t decl_type
3590 = hsa_type_for_scalar_tree_type (formal_arg_type
, false);
3591 return mem_type_for_type (decl_type
);
3594 /* Generate HSA instructions for a direct call instruction.
3595 Instructions will be appended to HBB, which also needs to be the
3596 corresponding structure to the basic_block of STMT.
3597 If ASSIGN_LHS is false, do not copy HSA function result argument into the
3598 corresponding HSA representation of the gimple statement LHS. */
3601 gen_hsa_insns_for_direct_call (gimple
*stmt
, hsa_bb
*hbb
,
3602 bool assign_lhs
= true)
3604 tree decl
= gimple_call_fndecl (stmt
);
3605 verify_function_arguments (decl
);
3606 if (hsa_seen_error ())
3609 hsa_insn_call
*call_insn
= new hsa_insn_call (decl
);
3610 hsa_cfun
->m_called_functions
.safe_push (call_insn
->m_called_function
);
3612 /* Argument block start. */
3613 hsa_insn_arg_block
*arg_start
3614 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3615 hbb
->append_insn (arg_start
);
3617 tree parm_type_chain
= TYPE_ARG_TYPES (gimple_call_fntype (stmt
));
3619 /* Preparation of arguments that will be passed to function. */
3620 const unsigned args
= gimple_call_num_args (stmt
);
3621 for (unsigned i
= 0; i
< args
; ++i
)
3623 tree parm
= gimple_call_arg (stmt
, (int)i
);
3624 tree parm_decl_type
= parm_type_chain
!= NULL_TREE
3625 ? TREE_VALUE (parm_type_chain
) : NULL_TREE
;
3626 hsa_op_address
*addr
;
3628 if (AGGREGATE_TYPE_P (TREE_TYPE (parm
)))
3630 addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3631 BrigAlignment8_t align
;
3632 hsa_op_address
*src
= gen_hsa_addr_with_align (parm
, hbb
, &align
);
3633 gen_hsa_memory_copy (hbb
, addr
, src
,
3634 addr
->m_symbol
->total_byte_size (), align
);
3638 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3640 if (parm_decl_type
!= NULL
&& AGGREGATE_TYPE_P (parm_decl_type
))
3642 HSA_SORRY_AT (gimple_location (stmt
),
3643 "support for HSA does not implement an aggregate "
3644 "formal argument in a function call, while actual "
3645 "argument is not an aggregate");
3649 BrigType16_t formal_arg_type
3650 = get_format_argument_type (parm_decl_type
, src
->m_type
);
3651 if (hsa_seen_error ())
3654 if (src
->m_type
!= formal_arg_type
)
3655 src
= src
->get_in_type (formal_arg_type
, hbb
);
3658 = gen_hsa_addr_for_arg (parm_decl_type
!= NULL_TREE
?
3659 parm_decl_type
: TREE_TYPE (parm
), i
);
3660 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, formal_arg_type
,
3663 hbb
->append_insn (mem
);
3666 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3667 if (parm_type_chain
)
3668 parm_type_chain
= TREE_CHAIN (parm_type_chain
);
3671 call_insn
->m_args_code_list
= new hsa_op_code_list (args
);
3672 hbb
->append_insn (call_insn
);
3674 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
3676 tree result
= gimple_call_lhs (stmt
);
3677 hsa_insn_mem
*result_insn
= NULL
;
3678 if (!VOID_TYPE_P (result_type
))
3680 hsa_op_address
*addr
= gen_hsa_addr_for_arg (result_type
, -1);
3682 /* Even if result of a function call is unused, we have to emit
3683 declaration for the result. */
3684 if (result
&& assign_lhs
)
3686 tree lhs_type
= TREE_TYPE (result
);
3688 if (hsa_seen_error ())
3691 if (AGGREGATE_TYPE_P (lhs_type
))
3693 BrigAlignment8_t align
;
3694 hsa_op_address
*result_addr
3695 = gen_hsa_addr_with_align (result
, hbb
, &align
);
3696 gen_hsa_memory_copy (hbb
, result_addr
, addr
,
3697 addr
->m_symbol
->total_byte_size (), align
);
3702 = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type
,
3705 hsa_op_reg
*dst
= hsa_cfun
->reg_for_gimple_ssa (result
);
3706 result_insn
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
, dst
, addr
);
3707 hbb
->append_insn (result_insn
);
3711 call_insn
->m_output_arg
= addr
->m_symbol
;
3712 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3718 HSA_SORRY_AT (gimple_location (stmt
),
3719 "support for HSA does not implement an assignment of "
3720 "return value from a void function");
3724 call_insn
->m_result_code_list
= new hsa_op_code_list (0);
3727 /* Argument block end. */
3728 hsa_insn_arg_block
*arg_end
3729 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3730 hbb
->append_insn (arg_end
);
3733 /* Generate HSA instructions for a direct call of an internal fn.
3734 Instructions will be appended to HBB, which also needs to be the
3735 corresponding structure to the basic_block of STMT. */
3738 gen_hsa_insns_for_call_of_internal_fn (gimple
*stmt
, hsa_bb
*hbb
)
3740 tree lhs
= gimple_call_lhs (stmt
);
3744 tree lhs_type
= TREE_TYPE (lhs
);
3745 tree rhs1
= gimple_call_arg (stmt
, 0);
3746 tree rhs1_type
= TREE_TYPE (rhs1
);
3747 enum internal_fn fn
= gimple_call_internal_fn (stmt
);
3748 hsa_internal_fn
*ifn
3749 = new hsa_internal_fn (fn
, tree_to_uhwi (TYPE_SIZE (rhs1_type
)));
3750 hsa_insn_call
*call_insn
= new hsa_insn_call (ifn
);
3752 gcc_checking_assert (FLOAT_TYPE_P (rhs1_type
));
3754 if (!hsa_emitted_internal_decls
->find (call_insn
->m_called_internal_fn
))
3755 hsa_cfun
->m_called_internal_fns
.safe_push (call_insn
->m_called_internal_fn
);
3757 hsa_insn_arg_block
*arg_start
3758 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3759 hbb
->append_insn (arg_start
);
3761 unsigned num_args
= gimple_call_num_args (stmt
);
3763 /* Function arguments. */
3764 for (unsigned i
= 0; i
< num_args
; i
++)
3766 tree parm
= gimple_call_arg (stmt
, (int)i
);
3767 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3769 hsa_op_address
*addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3770 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, src
->m_type
,
3773 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3774 hbb
->append_insn (mem
);
3777 call_insn
->m_args_code_list
= new hsa_op_code_list (num_args
);
3778 hbb
->append_insn (call_insn
);
3780 /* Assign returned value. */
3781 hsa_op_address
*addr
= gen_hsa_addr_for_arg (lhs_type
, -1);
3783 call_insn
->m_output_arg
= addr
->m_symbol
;
3784 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3786 /* Argument block end. */
3787 hsa_insn_arg_block
*arg_end
3788 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3789 hbb
->append_insn (arg_end
);
3792 /* Generate HSA instructions for a return value instruction.
3793 Instructions will be appended to HBB, which also needs to be the
3794 corresponding structure to the basic_block of STMT. */
3797 gen_hsa_insns_for_return (greturn
*stmt
, hsa_bb
*hbb
)
3799 tree retval
= gimple_return_retval (stmt
);
3802 hsa_op_address
*addr
= new hsa_op_address (hsa_cfun
->m_output_arg
);
3804 if (AGGREGATE_TYPE_P (TREE_TYPE (retval
)))
3806 BrigAlignment8_t align
;
3807 hsa_op_address
*retval_addr
= gen_hsa_addr_with_align (retval
, hbb
,
3809 gen_hsa_memory_copy (hbb
, addr
, retval_addr
,
3810 hsa_cfun
->m_output_arg
->total_byte_size (),
3815 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (retval
),
3817 BrigType16_t mtype
= mem_type_for_type (t
);
3819 /* Store of return value. */
3820 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (retval
, hbb
);
3821 src
= src
->get_in_type (mtype
, hbb
);
3822 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
,
3824 hbb
->append_insn (mem
);
3828 /* HSAIL return instruction emission. */
3829 hsa_insn_basic
*ret
= new hsa_insn_basic (0, BRIG_OPCODE_RET
);
3830 hbb
->append_insn (ret
);
3833 /* Set OP_INDEX-th operand of the instruction to DEST, as the DEST
3834 can have a different type, conversion instructions are possibly
3838 hsa_insn_basic::set_output_in_type (hsa_op_reg
*dest
, unsigned op_index
,
3841 gcc_checking_assert (op_output_p (op_index
));
3843 if (dest
->m_type
== m_type
)
3845 set_op (op_index
, dest
);
3849 hsa_insn_basic
*insn
;
3851 if (hsa_needs_cvt (dest
->m_type
, m_type
))
3853 tmp
= new hsa_op_reg (m_type
);
3854 insn
= new hsa_insn_cvt (dest
, tmp
);
3856 else if (hsa_type_bit_size (dest
->m_type
) == hsa_type_bit_size (m_type
))
3858 /* When output, HSA registers do not really have types, only sizes, so if
3859 the sizes match, we can use the register directly. */
3860 set_op (op_index
, dest
);
3865 tmp
= new hsa_op_reg (m_type
);
3866 insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
,
3867 dest
, tmp
->get_in_type (dest
->m_type
, hbb
));
3868 hsa_fixup_mov_insn_type (insn
);
3870 set_op (op_index
, tmp
);
3871 hbb
->append_insn (insn
);
3874 /* Generate instruction OPCODE to query a property of HSA grid along the
3875 given DIMENSION. Store result into DEST and append the instruction to
3879 query_hsa_grid_dim (hsa_op_reg
*dest
, int opcode
, hsa_op_immed
*dimension
,
3882 hsa_insn_basic
*insn
= new hsa_insn_basic (2, opcode
, BRIG_TYPE_U32
, NULL
,
3884 hbb
->append_insn (insn
);
3885 insn
->set_output_in_type (dest
, 0, hbb
);
3888 /* Generate instruction OPCODE to query a property of HSA grid along the given
3889 dimension which is an immediate in first argument of STMT. Store result
3890 into the register corresponding to LHS of STMT and append the instruction to
3894 query_hsa_grid_dim (gimple
*stmt
, int opcode
, hsa_bb
*hbb
)
3896 tree lhs
= gimple_call_lhs (dyn_cast
<gcall
*> (stmt
));
3897 if (lhs
== NULL_TREE
)
3900 tree arg
= gimple_call_arg (stmt
, 0);
3901 unsigned HOST_WIDE_INT dim
= 5;
3902 if (tree_fits_uhwi_p (arg
))
3903 dim
= tree_to_uhwi (arg
);
3906 HSA_SORRY_AT (gimple_location (stmt
),
3907 "HSA grid query dimension must be immediate constant 0, 1 "
3912 hsa_op_immed
*hdim
= new hsa_op_immed (dim
, (BrigKind16_t
) BRIG_TYPE_U32
);
3913 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3914 query_hsa_grid_dim (dest
, opcode
, hdim
, hbb
);
3917 /* Generate instruction OPCODE to query a property of HSA grid that is
3918 independent of any dimension. Store result into the register corresponding
3919 to LHS of STMT and append the instruction to HBB. */
3922 query_hsa_grid_nodim (gimple
*stmt
, BrigOpcode16_t opcode
, hsa_bb
*hbb
)
3924 tree lhs
= gimple_call_lhs (dyn_cast
<gcall
*> (stmt
));
3925 if (lhs
== NULL_TREE
)
3927 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3928 BrigType16_t brig_type
= hsa_unsigned_type_for_type (dest
->m_type
);
3929 hsa_insn_basic
*insn
= new hsa_insn_basic (1, opcode
, brig_type
, dest
);
3930 hbb
->append_insn (insn
);
3933 /* Emit instructions that set hsa_num_threads according to provided VALUE.
3934 Instructions are appended to basic block HBB. */
3937 gen_set_num_threads (tree value
, hsa_bb
*hbb
)
3939 hbb
->append_insn (new hsa_insn_comment ("omp_set_num_threads"));
3940 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (value
, hbb
);
3942 src
= src
->get_in_type (hsa_num_threads
->m_type
, hbb
);
3943 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
3945 hsa_insn_basic
*basic
3946 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
, src
, addr
);
3947 hbb
->append_insn (basic
);
3950 /* Return byte offset of a FIELD_NAME in GOMP_hsa_kernel_dispatch which
3951 is defined in plugin-hsa.c. */
3953 static HOST_WIDE_INT
3954 get_hsa_kernel_dispatch_offset (const char *field_name
)
3956 tree
*hsa_kernel_dispatch_type
= hsa_get_kernel_dispatch_type ();
3957 if (*hsa_kernel_dispatch_type
== NULL
)
3959 /* Collection of information needed for a dispatch of a kernel from a
3960 kernel. Keep in sync with libgomp's plugin-hsa.c. */
3962 *hsa_kernel_dispatch_type
= make_node (RECORD_TYPE
);
3963 tree id_f1
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3964 get_identifier ("queue"), ptr_type_node
);
3965 DECL_CHAIN (id_f1
) = NULL_TREE
;
3966 tree id_f2
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3967 get_identifier ("omp_data_memory"),
3969 DECL_CHAIN (id_f2
) = id_f1
;
3970 tree id_f3
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3971 get_identifier ("kernarg_address"),
3973 DECL_CHAIN (id_f3
) = id_f2
;
3974 tree id_f4
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3975 get_identifier ("object"),
3977 DECL_CHAIN (id_f4
) = id_f3
;
3978 tree id_f5
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3979 get_identifier ("signal"),
3981 DECL_CHAIN (id_f5
) = id_f4
;
3982 tree id_f6
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3983 get_identifier ("private_segment_size"),
3985 DECL_CHAIN (id_f6
) = id_f5
;
3986 tree id_f7
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3987 get_identifier ("group_segment_size"),
3989 DECL_CHAIN (id_f7
) = id_f6
;
3990 tree id_f8
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3991 get_identifier ("kernel_dispatch_count"),
3993 DECL_CHAIN (id_f8
) = id_f7
;
3994 tree id_f9
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3995 get_identifier ("debug"),
3997 DECL_CHAIN (id_f9
) = id_f8
;
3998 tree id_f10
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3999 get_identifier ("omp_level"),
4001 DECL_CHAIN (id_f10
) = id_f9
;
4002 tree id_f11
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
4003 get_identifier ("children_dispatches"),
4005 DECL_CHAIN (id_f11
) = id_f10
;
4006 tree id_f12
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
4007 get_identifier ("omp_num_threads"),
4009 DECL_CHAIN (id_f12
) = id_f11
;
4012 finish_builtin_struct (*hsa_kernel_dispatch_type
, "__hsa_kernel_dispatch",
4014 TYPE_ARTIFICIAL (*hsa_kernel_dispatch_type
) = 1;
4017 for (tree chain
= TYPE_FIELDS (*hsa_kernel_dispatch_type
);
4018 chain
!= NULL_TREE
; chain
= TREE_CHAIN (chain
))
4019 if (id_equal (DECL_NAME (chain
), field_name
))
4020 return int_byte_position (chain
);
4025 /* Return an HSA register that will contain number of threads for
4026 a future dispatched kernel. Instructions are added to HBB. */
4029 gen_num_threads_for_dispatch (hsa_bb
*hbb
)
4031 /* Step 1) Assign to number of threads:
4032 MIN (HSA_DEFAULT_NUM_THREADS, hsa_num_threads). */
4033 hsa_op_reg
*threads
= new hsa_op_reg (hsa_num_threads
->m_type
);
4034 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
4036 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD
, threads
->m_type
,
4039 hsa_op_immed
*limit
= new hsa_op_immed (HSA_DEFAULT_NUM_THREADS
,
4041 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
4043 = new hsa_insn_cmp (BRIG_COMPARE_LT
, r
->m_type
, r
, threads
, limit
);
4044 hbb
->append_insn (cmp
);
4046 BrigType16_t btype
= hsa_bittype_for_type (threads
->m_type
);
4047 hsa_op_reg
*tmp
= new hsa_op_reg (threads
->m_type
);
4049 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp
, r
,
4052 /* Step 2) If the number is equal to zero,
4053 return shadow->omp_num_threads. */
4054 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4056 hsa_op_reg
*shadow_thread_count
= new hsa_op_reg (BRIG_TYPE_U32
);
4058 = new hsa_op_address (shadow_reg_ptr
,
4059 get_hsa_kernel_dispatch_offset ("omp_num_threads"));
4060 hsa_insn_basic
*basic
4061 = new hsa_insn_mem (BRIG_OPCODE_LD
, shadow_thread_count
->m_type
,
4062 shadow_thread_count
, addr
);
4063 hbb
->append_insn (basic
);
4065 hsa_op_reg
*tmp2
= new hsa_op_reg (threads
->m_type
);
4066 r
= new hsa_op_reg (BRIG_TYPE_B1
);
4067 hsa_op_immed
*imm
= new hsa_op_immed (0, shadow_thread_count
->m_type
);
4068 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_EQ
, r
->m_type
, r
, tmp
, imm
));
4069 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp2
, r
,
4070 shadow_thread_count
, tmp
));
4072 hsa_op_base
*dest
= tmp2
->get_in_type (BRIG_TYPE_U16
, hbb
);
4074 return as_a
<hsa_op_reg
*> (dest
);
4077 /* Build OPCODE query for all three hsa dimensions, multiply them and store the
4078 result into DEST. */
4081 multiply_grid_dim_characteristics (hsa_op_reg
*dest
, int opcode
, hsa_bb
*hbb
)
4083 hsa_op_reg
*dimx
= new hsa_op_reg (BRIG_TYPE_U32
);
4084 query_hsa_grid_dim (dimx
, opcode
,
4085 new hsa_op_immed (0, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4086 hsa_op_reg
*dimy
= new hsa_op_reg (BRIG_TYPE_U32
);
4087 query_hsa_grid_dim (dimy
, opcode
,
4088 new hsa_op_immed (1, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4089 hsa_op_reg
*dimz
= new hsa_op_reg (BRIG_TYPE_U32
);
4090 query_hsa_grid_dim (dimz
, opcode
,
4091 new hsa_op_immed (2, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4092 hsa_op_reg
*tmp
= new hsa_op_reg (dest
->m_type
);
4093 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp
,
4094 dimx
->get_in_type (dest
->m_type
, hbb
),
4095 dimy
->get_in_type (dest
->m_type
, hbb
), hbb
);
4096 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, dest
, tmp
,
4097 dimz
->get_in_type (dest
->m_type
, hbb
), hbb
);
4100 /* Emit instructions that assign number of threads to lhs of gimple STMT.
4101 Instructions are appended to basic block HBB. */
4104 gen_get_num_threads (gimple
*stmt
, hsa_bb
*hbb
)
4106 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4109 hbb
->append_insn (new hsa_insn_comment ("omp_get_num_threads"));
4110 tree lhs
= gimple_call_lhs (stmt
);
4111 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4112 multiply_grid_dim_characteristics (dest
, BRIG_OPCODE_CURRENTWORKGROUPSIZE
,
4116 /* Emit instructions that assign number of teams to lhs of gimple STMT.
4117 Instructions are appended to basic block HBB. */
4120 gen_get_num_teams (gimple
*stmt
, hsa_bb
*hbb
)
4122 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4125 hbb
->append_insn (new hsa_insn_comment ("omp_get_num_teams"));
4126 tree lhs
= gimple_call_lhs (stmt
);
4127 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4128 multiply_grid_dim_characteristics (dest
, BRIG_OPCODE_GRIDGROUPS
, hbb
);
4131 /* Emit instructions that assign a team number to lhs of gimple STMT.
4132 Instructions are appended to basic block HBB. */
4135 gen_get_team_num (gimple
*stmt
, hsa_bb
*hbb
)
4137 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4140 hbb
->append_insn (new hsa_insn_comment ("omp_get_team_num"));
4141 tree lhs
= gimple_call_lhs (stmt
);
4142 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4144 hsa_op_reg
*gnum_x
= new hsa_op_reg (BRIG_TYPE_U32
);
4145 query_hsa_grid_dim (gnum_x
, BRIG_OPCODE_GRIDGROUPS
,
4146 new hsa_op_immed (0, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4147 hsa_op_reg
*gnum_y
= new hsa_op_reg (BRIG_TYPE_U32
);
4148 query_hsa_grid_dim (gnum_y
, BRIG_OPCODE_GRIDGROUPS
,
4149 new hsa_op_immed (1, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4151 hsa_op_reg
*gno_z
= new hsa_op_reg (BRIG_TYPE_U32
);
4152 query_hsa_grid_dim (gno_z
, BRIG_OPCODE_WORKGROUPID
,
4153 new hsa_op_immed (2, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4155 hsa_op_reg
*tmp1
= new hsa_op_reg (dest
->m_type
);
4156 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp1
,
4157 gnum_x
->get_in_type (dest
->m_type
, hbb
),
4158 gnum_y
->get_in_type (dest
->m_type
, hbb
), hbb
);
4159 hsa_op_reg
*tmp2
= new hsa_op_reg (dest
->m_type
);
4160 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp2
, tmp1
,
4161 gno_z
->get_in_type (dest
->m_type
, hbb
), hbb
);
4163 hsa_op_reg
*gno_y
= new hsa_op_reg (BRIG_TYPE_U32
);
4164 query_hsa_grid_dim (gno_y
, BRIG_OPCODE_WORKGROUPID
,
4165 new hsa_op_immed (1, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4166 hsa_op_reg
*tmp3
= new hsa_op_reg (dest
->m_type
);
4167 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp3
,
4168 gnum_x
->get_in_type (dest
->m_type
, hbb
),
4169 gno_y
->get_in_type (dest
->m_type
, hbb
), hbb
);
4170 hsa_op_reg
*tmp4
= new hsa_op_reg (dest
->m_type
);
4171 gen_hsa_binary_operation (BRIG_OPCODE_ADD
, tmp4
, tmp3
, tmp2
, hbb
);
4172 hsa_op_reg
*gno_x
= new hsa_op_reg (BRIG_TYPE_U32
);
4173 query_hsa_grid_dim (gno_x
, BRIG_OPCODE_WORKGROUPID
,
4174 new hsa_op_immed (0, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4175 gen_hsa_binary_operation (BRIG_OPCODE_ADD
, dest
, tmp4
,
4176 gno_x
->get_in_type (dest
->m_type
, hbb
), hbb
);
4179 /* Emit instructions that get levels-var ICV to lhs of gimple STMT.
4180 Instructions are appended to basic block HBB. */
4183 gen_get_level (gimple
*stmt
, hsa_bb
*hbb
)
4185 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4188 hbb
->append_insn (new hsa_insn_comment ("omp_get_level"));
4190 tree lhs
= gimple_call_lhs (stmt
);
4191 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4193 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4194 if (shadow_reg_ptr
== NULL
)
4196 HSA_SORRY_AT (gimple_location (stmt
),
4197 "support for HSA does not implement omp_get_level called "
4198 "from a function not being inlined within a kernel");
4202 hsa_op_address
*addr
4203 = new hsa_op_address (shadow_reg_ptr
,
4204 get_hsa_kernel_dispatch_offset ("omp_level"));
4206 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, BRIG_TYPE_U64
,
4207 (hsa_op_base
*) NULL
, addr
);
4208 hbb
->append_insn (mem
);
4209 mem
->set_output_in_type (dest
, 0, hbb
);
4212 /* Emit instruction that implement omp_get_max_threads of gimple STMT. */
4215 gen_get_max_threads (gimple
*stmt
, hsa_bb
*hbb
)
4217 tree lhs
= gimple_call_lhs (stmt
);
4221 hbb
->append_insn (new hsa_insn_comment ("omp_get_max_threads"));
4223 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4224 hsa_op_with_type
*num_theads_reg
= gen_num_threads_for_dispatch (hbb
)
4225 ->get_in_type (dest
->m_type
, hbb
);
4226 hsa_build_append_simple_mov (dest
, num_theads_reg
, hbb
);
4229 /* Emit instructions that implement alloca builtin gimple STMT.
4230 Instructions are appended to basic block HBB. */
4233 gen_hsa_alloca (gcall
*call
, hsa_bb
*hbb
)
4235 tree lhs
= gimple_call_lhs (call
);
4236 if (lhs
== NULL_TREE
)
4239 built_in_function fn
= DECL_FUNCTION_CODE (gimple_call_fndecl (call
));
4241 gcc_checking_assert (ALLOCA_FUNCTION_CODE_P (fn
));
4243 unsigned bit_alignment
= 0;
4245 if (fn
!= BUILT_IN_ALLOCA
)
4247 tree alignment_tree
= gimple_call_arg (call
, 1);
4248 if (TREE_CODE (alignment_tree
) != INTEGER_CST
)
4250 HSA_SORRY_ATV (gimple_location (call
),
4251 "support for HSA does not implement "
4252 "__builtin_alloca_with_align with a non-constant "
4253 "alignment: %E", alignment_tree
);
4256 bit_alignment
= tree_to_uhwi (alignment_tree
);
4259 tree rhs1
= gimple_call_arg (call
, 0);
4260 hsa_op_with_type
*size
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
)
4261 ->get_in_type (BRIG_TYPE_U32
, hbb
);
4262 hsa_op_with_type
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4265 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
));
4266 hsa_insn_alloca
*a
= new hsa_insn_alloca (tmp
, size
, bit_alignment
);
4267 hbb
->append_insn (a
);
4270 = new hsa_insn_seg (BRIG_OPCODE_STOF
,
4271 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
4272 tmp
->m_type
, BRIG_SEGMENT_PRIVATE
, dest
, tmp
);
4273 hbb
->append_insn (seg
);
4276 /* Emit instructions that implement clrsb builtin STMT:
4277 Returns the number of leading redundant sign bits in x, i.e. the number
4278 of bits following the most significant bit that are identical to it.
4279 There are no special cases for 0 or other values.
4280 Instructions are appended to basic block HBB. */
4283 gen_hsa_clrsb (gcall
*call
, hsa_bb
*hbb
)
4285 tree lhs
= gimple_call_lhs (call
);
4286 if (lhs
== NULL_TREE
)
4289 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4290 tree rhs1
= gimple_call_arg (call
, 0);
4291 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4292 arg
->extend_int_to_32bit (hbb
);
4293 BrigType16_t bittype
= hsa_bittype_for_type (arg
->m_type
);
4294 unsigned bitsize
= tree_to_uhwi (TYPE_SIZE (TREE_TYPE (rhs1
)));
4296 /* FIRSTBIT instruction is defined just for 32 and 64-bits wide integers. */
4297 gcc_checking_assert (bitsize
== 32 || bitsize
== 64);
4299 /* Set true to MOST_SIG if the most significant bit is set to one. */
4300 hsa_op_immed
*c
= new hsa_op_immed (1ul << (bitsize
- 1),
4301 hsa_uint_for_bitsize (bitsize
));
4303 hsa_op_reg
*and_reg
= new hsa_op_reg (bittype
);
4304 gen_hsa_binary_operation (BRIG_OPCODE_AND
, and_reg
, arg
, c
, hbb
);
4306 hsa_op_reg
*most_sign
= new hsa_op_reg (BRIG_TYPE_B1
);
4308 = new hsa_insn_cmp (BRIG_COMPARE_EQ
, most_sign
->m_type
, most_sign
,
4310 hbb
->append_insn (cmp
);
4312 /* If the most significant bit is one, negate the input. Otherwise
4313 shift the input value to left by one bit. */
4314 hsa_op_reg
*arg_neg
= new hsa_op_reg (arg
->m_type
);
4315 gen_hsa_unary_operation (BRIG_OPCODE_NEG
, arg_neg
, arg
, hbb
);
4317 hsa_op_reg
*shifted_arg
= new hsa_op_reg (arg
->m_type
);
4318 gen_hsa_binary_operation (BRIG_OPCODE_SHL
, shifted_arg
, arg
,
4319 new hsa_op_immed (1, BRIG_TYPE_U64
), hbb
);
4321 /* Assign the value that can be used for FIRSTBIT instruction according
4322 to the most significant bit. */
4323 hsa_op_reg
*tmp
= new hsa_op_reg (bittype
);
4324 hsa_insn_basic
*cmov
4325 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, bittype
, tmp
, most_sign
,
4326 arg_neg
, shifted_arg
);
4327 hbb
->append_insn (cmov
);
4329 hsa_op_reg
*leading_bits
= new hsa_op_reg (BRIG_TYPE_S32
);
4330 gen_hsa_unary_operation (BRIG_OPCODE_FIRSTBIT
, leading_bits
,
4331 tmp
->get_in_type (hsa_uint_for_bitsize (bitsize
),
4334 /* Set flag if the input value is equal to zero. */
4335 hsa_op_reg
*is_zero
= new hsa_op_reg (BRIG_TYPE_B1
);
4336 cmp
= new hsa_insn_cmp (BRIG_COMPARE_EQ
, is_zero
->m_type
, is_zero
, arg
,
4337 new hsa_op_immed (0, arg
->m_type
));
4338 hbb
->append_insn (cmp
);
4340 /* Return the number of leading bits,
4341 or (bitsize - 1) if the input value is zero. */
4342 cmov
= new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, BRIG_TYPE_B32
, NULL
, is_zero
,
4343 new hsa_op_immed (bitsize
- 1, BRIG_TYPE_U32
),
4344 leading_bits
->get_in_type (BRIG_TYPE_B32
, hbb
));
4345 hbb
->append_insn (cmov
);
4346 cmov
->set_output_in_type (dest
, 0, hbb
);
4349 /* Emit instructions that implement ffs builtin STMT:
4350 Returns one plus the index of the least significant 1-bit of x,
4351 or if x is zero, returns zero.
4352 Instructions are appended to basic block HBB. */
4355 gen_hsa_ffs (gcall
*call
, hsa_bb
*hbb
)
4357 tree lhs
= gimple_call_lhs (call
);
4358 if (lhs
== NULL_TREE
)
4361 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4363 tree rhs1
= gimple_call_arg (call
, 0);
4364 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4365 arg
= arg
->extend_int_to_32bit (hbb
);
4367 hsa_op_reg
*tmp
= new hsa_op_reg (BRIG_TYPE_U32
);
4368 hsa_insn_srctype
*insn
= new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT
,
4369 tmp
->m_type
, arg
->m_type
,
4371 hbb
->append_insn (insn
);
4373 hsa_insn_basic
*addition
4374 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
, NULL
, tmp
,
4375 new hsa_op_immed (1, tmp
->m_type
));
4376 hbb
->append_insn (addition
);
4377 addition
->set_output_in_type (dest
, 0, hbb
);
4381 gen_hsa_popcount_to_dest (hsa_op_reg
*dest
, hsa_op_with_type
*arg
, hsa_bb
*hbb
)
4383 gcc_checking_assert (hsa_type_integer_p (arg
->m_type
));
4385 if (hsa_type_bit_size (arg
->m_type
) < 32)
4386 arg
= arg
->get_in_type (BRIG_TYPE_B32
, hbb
);
4388 BrigType16_t srctype
= hsa_bittype_for_type (arg
->m_type
);
4389 if (!hsa_btype_p (arg
->m_type
))
4390 arg
= arg
->get_in_type (srctype
, hbb
);
4392 hsa_insn_srctype
*popcount
4393 = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT
, BRIG_TYPE_U32
,
4394 srctype
, NULL
, arg
);
4395 hbb
->append_insn (popcount
);
4396 popcount
->set_output_in_type (dest
, 0, hbb
);
4399 /* Emit instructions that implement parity builtin STMT:
4400 Returns the parity of x, i.e. the number of 1-bits in x modulo 2.
4401 Instructions are appended to basic block HBB. */
4404 gen_hsa_parity (gcall
*call
, hsa_bb
*hbb
)
4406 tree lhs
= gimple_call_lhs (call
);
4407 if (lhs
== NULL_TREE
)
4410 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4411 tree rhs1
= gimple_call_arg (call
, 0);
4412 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4414 hsa_op_reg
*popcount
= new hsa_op_reg (BRIG_TYPE_U32
);
4415 gen_hsa_popcount_to_dest (popcount
, arg
, hbb
);
4417 hsa_insn_basic
*insn
4418 = new hsa_insn_basic (3, BRIG_OPCODE_REM
, popcount
->m_type
, NULL
, popcount
,
4419 new hsa_op_immed (2, popcount
->m_type
));
4420 hbb
->append_insn (insn
);
4421 insn
->set_output_in_type (dest
, 0, hbb
);
4424 /* Emit instructions that implement popcount builtin STMT.
4425 Instructions are appended to basic block HBB. */
4428 gen_hsa_popcount (gcall
*call
, hsa_bb
*hbb
)
4430 tree lhs
= gimple_call_lhs (call
);
4431 if (lhs
== NULL_TREE
)
4434 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4435 tree rhs1
= gimple_call_arg (call
, 0);
4436 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4438 gen_hsa_popcount_to_dest (dest
, arg
, hbb
);
4441 /* Emit instructions that implement DIVMOD builtin STMT.
4442 Instructions are appended to basic block HBB. */
4445 gen_hsa_divmod (gcall
*call
, hsa_bb
*hbb
)
4447 tree lhs
= gimple_call_lhs (call
);
4448 if (lhs
== NULL_TREE
)
4451 tree rhs0
= gimple_call_arg (call
, 0);
4452 tree rhs1
= gimple_call_arg (call
, 1);
4454 hsa_op_with_type
*arg0
= hsa_reg_or_immed_for_gimple_op (rhs0
, hbb
);
4455 arg0
= arg0
->extend_int_to_32bit (hbb
);
4456 hsa_op_with_type
*arg1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4457 arg1
= arg1
->extend_int_to_32bit (hbb
);
4459 hsa_op_reg
*dest0
= new hsa_op_reg (arg0
->m_type
);
4460 hsa_op_reg
*dest1
= new hsa_op_reg (arg1
->m_type
);
4462 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_DIV
, dest0
->m_type
,
4464 hbb
->append_insn (insn
);
4465 insn
= new hsa_insn_basic (3, BRIG_OPCODE_REM
, dest1
->m_type
, dest1
, arg0
,
4467 hbb
->append_insn (insn
);
4469 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4470 BrigType16_t dst_type
= hsa_extend_inttype_to_32bit (dest
->m_type
);
4471 BrigType16_t src_type
= hsa_bittype_for_type (dest0
->m_type
);
4473 insn
= new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dst_type
,
4474 src_type
, NULL
, dest0
, dest1
);
4475 hbb
->append_insn (insn
);
4476 insn
->set_output_in_type (dest
, 0, hbb
);
4479 /* Set VALUE to a shadow kernel debug argument and append a new instruction
4480 to HBB basic block. */
4483 set_debug_value (hsa_bb
*hbb
, hsa_op_with_type
*value
)
4485 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4486 if (shadow_reg_ptr
== NULL
)
4489 hsa_op_address
*addr
4490 = new hsa_op_address (shadow_reg_ptr
,
4491 get_hsa_kernel_dispatch_offset ("debug"));
4492 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, BRIG_TYPE_U64
, value
,
4494 hbb
->append_insn (mem
);
4498 omp_simple_builtin::generate (gimple
*stmt
, hsa_bb
*hbb
)
4502 if (m_warning_message
)
4503 HSA_SORRY_AT (gimple_location (stmt
), m_warning_message
);
4505 HSA_SORRY_ATV (gimple_location (stmt
),
4506 "Support for HSA does not implement calls to %s\n",
4509 else if (m_warning_message
!= NULL
)
4510 warning_at (gimple_location (stmt
), OPT_Whsa
, m_warning_message
);
4512 if (m_return_value
!= NULL
)
4514 tree lhs
= gimple_call_lhs (stmt
);
4518 hbb
->append_insn (new hsa_insn_comment (m_name
));
4520 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4521 hsa_op_with_type
*op
= m_return_value
->get_in_type (dest
->m_type
, hbb
);
4522 hsa_build_append_simple_mov (dest
, op
, hbb
);
4526 /* If STMT is a call of a known library function, generate code to perform
4527 it and return true. */
4530 gen_hsa_insns_for_known_library_call (gimple
*stmt
, hsa_bb
*hbb
)
4532 bool handled
= false;
4533 const char *name
= hsa_get_declaration_name (gimple_call_fndecl (stmt
));
4536 size_t len
= strlen (name
);
4537 if (len
> 0 && name
[len
- 1] == '_')
4539 copy
= XNEWVEC (char, len
+ 1);
4540 strcpy (copy
, name
);
4541 copy
[len
- 1] = '\0';
4545 /* Handle omp_* routines. */
4546 if (strstr (name
, "omp_") == name
)
4548 hsa_init_simple_builtins ();
4549 omp_simple_builtin
*builtin
= omp_simple_builtins
->get (name
);
4552 builtin
->generate (stmt
, hbb
);
4557 if (strcmp (name
, "omp_set_num_threads") == 0)
4558 gen_set_num_threads (gimple_call_arg (stmt
, 0), hbb
);
4559 else if (strcmp (name
, "omp_get_thread_num") == 0)
4561 hbb
->append_insn (new hsa_insn_comment (name
));
4562 query_hsa_grid_nodim (stmt
, BRIG_OPCODE_WORKITEMFLATABSID
, hbb
);
4564 else if (strcmp (name
, "omp_get_num_threads") == 0)
4566 hbb
->append_insn (new hsa_insn_comment (name
));
4567 gen_get_num_threads (stmt
, hbb
);
4569 else if (strcmp (name
, "omp_get_num_teams") == 0)
4570 gen_get_num_teams (stmt
, hbb
);
4571 else if (strcmp (name
, "omp_get_team_num") == 0)
4572 gen_get_team_num (stmt
, hbb
);
4573 else if (strcmp (name
, "omp_get_level") == 0)
4574 gen_get_level (stmt
, hbb
);
4575 else if (strcmp (name
, "omp_get_active_level") == 0)
4576 gen_get_level (stmt
, hbb
);
4577 else if (strcmp (name
, "omp_in_parallel") == 0)
4578 gen_get_level (stmt
, hbb
);
4579 else if (strcmp (name
, "omp_get_max_threads") == 0)
4580 gen_get_max_threads (stmt
, hbb
);
4592 if (strcmp (name
, "__hsa_set_debug_value") == 0)
4595 if (hsa_cfun
->has_shadow_reg_p ())
4597 tree rhs1
= gimple_call_arg (stmt
, 0);
4598 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4600 src
= src
->get_in_type (BRIG_TYPE_U64
, hbb
);
4601 set_debug_value (hbb
, src
);
4610 /* Helper functions to create a single unary HSA operations out of calls to
4611 builtins. OPCODE is the HSA operation to be generated. STMT is a gimple
4612 call to a builtin. HBB is the HSA BB to which the instruction should be
4613 added. Note that nothing will be created if STMT does not have a LHS. */
4616 gen_hsa_unaryop_for_builtin (BrigOpcode opcode
, gimple
*stmt
, hsa_bb
*hbb
)
4618 tree lhs
= gimple_call_lhs (stmt
);
4621 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4622 hsa_op_with_type
*op
4623 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4624 gen_hsa_unary_operation (opcode
, dest
, op
, hbb
);
4627 /* Helper functions to create a call to standard library if LHS of the
4628 STMT is used. HBB is the HSA BB to which the instruction should be
4632 gen_hsa_unaryop_builtin_call (gimple
*stmt
, hsa_bb
*hbb
)
4634 tree lhs
= gimple_call_lhs (stmt
);
4638 if (gimple_call_internal_p (stmt
))
4639 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
4641 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4644 /* Helper functions to create a single unary HSA operations out of calls to
4645 builtins (if unsafe math optimizations are enable). Otherwise, create
4646 a call to standard library function.
4647 OPCODE is the HSA operation to be generated. STMT is a gimple
4648 call to a builtin. HBB is the HSA BB to which the instruction should be
4649 added. Note that nothing will be created if STMT does not have a LHS. */
4652 gen_hsa_unaryop_or_call_for_builtin (BrigOpcode opcode
, gimple
*stmt
,
4655 if (flag_unsafe_math_optimizations
)
4656 gen_hsa_unaryop_for_builtin (opcode
, stmt
, hbb
);
4658 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4661 /* Generate HSA address corresponding to a value VAL (as opposed to a memory
4662 reference tree), for example an SSA_NAME or an ADDR_EXPR. HBB is the HSA BB
4663 to which the instruction should be added. */
4665 static hsa_op_address
*
4666 get_address_from_value (tree val
, hsa_bb
*hbb
)
4668 switch (TREE_CODE (val
))
4672 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4674 = hsa_cfun
->reg_for_gimple_ssa (val
)->get_in_type (addrtype
, hbb
);
4675 return new hsa_op_address (NULL
, as_a
<hsa_op_reg
*> (reg
), 0);
4678 return gen_hsa_addr (TREE_OPERAND (val
, 0), hbb
);
4681 if (tree_fits_shwi_p (val
))
4682 return new hsa_op_address (NULL
, NULL
, tree_to_shwi (val
));
4686 HSA_SORRY_ATV (EXPR_LOCATION (val
),
4687 "support for HSA does not implement memory access to %E",
4689 return new hsa_op_address (NULL
, NULL
, 0);
4693 /* Expand assignment of a result of a string BUILTIN to DST.
4694 Size of the operation is N bytes, where instructions
4695 will be append to HBB. */
4698 expand_lhs_of_string_op (gimple
*stmt
,
4699 unsigned HOST_WIDE_INT n
, hsa_bb
*hbb
,
4700 enum built_in_function builtin
)
4702 /* If LHS is expected, we need to emit a PHI instruction. */
4703 tree lhs
= gimple_call_lhs (stmt
);
4707 hsa_op_reg
*lhs_reg
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4709 hsa_op_with_type
*dst_reg
4710 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4711 hsa_op_with_type
*tmp
;
4715 case BUILT_IN_MEMPCPY
:
4717 tmp
= new hsa_op_reg (dst_reg
->m_type
);
4719 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
,
4721 new hsa_op_immed (n
, dst_reg
->m_type
));
4722 hbb
->append_insn (add
);
4725 case BUILT_IN_MEMCPY
:
4726 case BUILT_IN_MEMSET
:
4733 hbb
->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV
, lhs_reg
->m_type
,
4737 #define HSA_MEMORY_BUILTINS_LIMIT 128
4739 /* Expand a string builtin (from a gimple STMT) in a way that
4740 according to MISALIGNED_FLAG we process either direct emission
4741 (a bunch of memory load and store instructions), or we emit a function call
4742 of a library function (for instance 'memcpy'). Actually, a basic block
4743 for direct emission is just prepared, where caller is responsible
4744 for emission of corresponding instructions.
4745 All instruction are appended to HBB. */
4748 expand_string_operation_builtin (gimple
*stmt
, hsa_bb
*hbb
,
4749 hsa_op_reg
*misaligned_flag
)
4751 edge e
= split_block (hbb
->m_bb
, stmt
);
4752 basic_block condition_bb
= e
->src
;
4753 hbb
->append_insn (new hsa_insn_cbr (misaligned_flag
));
4755 /* Prepare the control flow. */
4756 edge condition_edge
= EDGE_SUCC (condition_bb
, 0);
4757 basic_block call_bb
= split_edge (condition_edge
);
4759 basic_block expanded_bb
= split_edge (EDGE_SUCC (call_bb
, 0));
4760 basic_block cont_bb
= EDGE_SUCC (expanded_bb
, 0)->dest
;
4761 basic_block merge_bb
= split_edge (EDGE_PRED (cont_bb
, 0));
4763 condition_edge
->flags
&= ~EDGE_FALLTHRU
;
4764 condition_edge
->flags
|= EDGE_TRUE_VALUE
;
4765 make_edge (condition_bb
, expanded_bb
, EDGE_FALSE_VALUE
);
4767 redirect_edge_succ (EDGE_SUCC (call_bb
, 0), merge_bb
);
4769 hsa_cfun
->m_modified_cfg
= true;
4771 hsa_init_new_bb (expanded_bb
);
4773 /* Slow path: function call. */
4774 gen_hsa_insns_for_direct_call (stmt
, hsa_init_new_bb (call_bb
), false);
4776 return hsa_bb_for_bb (expanded_bb
);
4779 /* Expand a memory copy BUILTIN (BUILT_IN_MEMCPY, BUILT_IN_MEMPCPY) from
4780 a gimple STMT and store all necessary instruction to HBB basic block. */
4783 expand_memory_copy (gimple
*stmt
, hsa_bb
*hbb
, enum built_in_function builtin
)
4785 tree byte_size
= gimple_call_arg (stmt
, 2);
4787 if (!tree_fits_uhwi_p (byte_size
))
4789 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4793 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
4795 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
4797 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4801 tree dst
= gimple_call_arg (stmt
, 0);
4802 tree src
= gimple_call_arg (stmt
, 1);
4804 hsa_op_address
*dst_addr
= get_address_from_value (dst
, hbb
);
4805 hsa_op_address
*src_addr
= get_address_from_value (src
, hbb
);
4807 /* As gen_hsa_memory_copy relies on memory alignment
4808 greater or equal to 8 bytes, we need to verify the alignment. */
4809 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4810 hsa_op_reg
*src_addr_reg
= new hsa_op_reg (addrtype
);
4811 hsa_op_reg
*dst_addr_reg
= new hsa_op_reg (addrtype
);
4813 convert_addr_to_flat_segment (src_addr
, src_addr_reg
, hbb
);
4814 convert_addr_to_flat_segment (dst_addr
, dst_addr_reg
, hbb
);
4816 /* Process BIT OR for source and destination addresses. */
4817 hsa_op_reg
*or_reg
= new hsa_op_reg (addrtype
);
4818 gen_hsa_binary_operation (BRIG_OPCODE_OR
, or_reg
, src_addr_reg
,
4821 /* Process BIT AND with 0x7 to identify the desired alignment
4823 hsa_op_reg
*masked
= new hsa_op_reg (addrtype
);
4825 gen_hsa_binary_operation (BRIG_OPCODE_AND
, masked
, or_reg
,
4826 new hsa_op_immed (7, addrtype
), hbb
);
4828 hsa_op_reg
*misaligned
= new hsa_op_reg (BRIG_TYPE_B1
);
4829 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE
, misaligned
->m_type
,
4831 new hsa_op_immed (0, masked
->m_type
)));
4833 hsa_bb
*native_impl_bb
4834 = expand_string_operation_builtin (stmt
, hbb
, misaligned
);
4836 gen_hsa_memory_copy (native_impl_bb
, dst_addr
, src_addr
, n
, BRIG_ALIGNMENT_8
);
4838 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb
->m_bb
, 0)->dest
);
4839 expand_lhs_of_string_op (stmt
, n
, merge_bb
, builtin
);
4843 /* Expand a memory set BUILTIN (BUILT_IN_MEMSET, BUILT_IN_BZERO) from
4844 a gimple STMT and store all necessary instruction to HBB basic block.
4845 The operation set N bytes with a CONSTANT value. */
4848 expand_memory_set (gimple
*stmt
, unsigned HOST_WIDE_INT n
,
4849 unsigned HOST_WIDE_INT constant
, hsa_bb
*hbb
,
4850 enum built_in_function builtin
)
4852 tree dst
= gimple_call_arg (stmt
, 0);
4853 hsa_op_address
*dst_addr
= get_address_from_value (dst
, hbb
);
4855 /* As gen_hsa_memory_set relies on memory alignment
4856 greater or equal to 8 bytes, we need to verify the alignment. */
4857 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4858 hsa_op_reg
*dst_addr_reg
= new hsa_op_reg (addrtype
);
4859 convert_addr_to_flat_segment (dst_addr
, dst_addr_reg
, hbb
);
4861 /* Process BIT AND with 0x7 to identify the desired alignment
4863 hsa_op_reg
*masked
= new hsa_op_reg (addrtype
);
4865 gen_hsa_binary_operation (BRIG_OPCODE_AND
, masked
, dst_addr_reg
,
4866 new hsa_op_immed (7, addrtype
), hbb
);
4868 hsa_op_reg
*misaligned
= new hsa_op_reg (BRIG_TYPE_B1
);
4869 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE
, misaligned
->m_type
,
4871 new hsa_op_immed (0, masked
->m_type
)));
4873 hsa_bb
*native_impl_bb
4874 = expand_string_operation_builtin (stmt
, hbb
, misaligned
);
4876 gen_hsa_memory_set (native_impl_bb
, dst_addr
, constant
, n
, BRIG_ALIGNMENT_8
);
4878 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb
->m_bb
, 0)->dest
);
4879 expand_lhs_of_string_op (stmt
, n
, merge_bb
, builtin
);
4882 /* Store into MEMORDER the memory order specified by tree T, which must be an
4883 integer constant representing a C++ memory order. If it isn't, issue an HSA
4884 sorry message using LOC and return true, otherwise return false and store
4885 the name of the requested order to *MNAME. */
4888 hsa_memorder_from_tree (tree t
, BrigMemoryOrder
*memorder
, const char **mname
,
4891 if (!tree_fits_uhwi_p (t
))
4893 HSA_SORRY_ATV (loc
, "support for HSA does not implement memory model %E",
4898 unsigned HOST_WIDE_INT mm
= tree_to_uhwi (t
);
4899 switch (mm
& MEMMODEL_BASE_MASK
)
4901 case MEMMODEL_RELAXED
:
4902 *memorder
= BRIG_MEMORY_ORDER_RELAXED
;
4905 case MEMMODEL_CONSUME
:
4906 /* HSA does not have an equivalent, but we can use the slightly stronger
4908 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
4911 case MEMMODEL_ACQUIRE
:
4912 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
4915 case MEMMODEL_RELEASE
:
4916 *memorder
= BRIG_MEMORY_ORDER_SC_RELEASE
;
4919 case MEMMODEL_ACQ_REL
:
4920 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
4923 case MEMMODEL_SEQ_CST
:
4924 /* Callers implementing a simple load or store need to remove the release
4925 or acquire part respectively. */
4926 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
4931 HSA_SORRY_AT (loc
, "support for HSA does not implement the specified "
4939 /* Helper function to create an HSA atomic operation instruction out of calls
4940 to atomic builtins. RET_ORIG is true if the built-in is the variant that
4941 return s the value before applying operation, and false if it should return
4942 the value after applying the operation (if it returns value at all). ACODE
4943 is the atomic operation code, STMT is a gimple call to a builtin. HBB is
4944 the HSA BB to which the instruction should be added. If SIGNAL is true, the
4945 created operation will work on HSA signals rather than atomic variables. */
4948 gen_hsa_atomic_for_builtin (bool ret_orig
, enum BrigAtomicOperation acode
,
4949 gimple
*stmt
, hsa_bb
*hbb
, bool signal
)
4951 tree lhs
= gimple_call_lhs (stmt
);
4953 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
4954 BrigType16_t hsa_type
= hsa_type_for_scalar_tree_type (type
, false);
4955 BrigType16_t mtype
= mem_type_for_type (hsa_type
);
4956 BrigMemoryOrder memorder
;
4959 if (hsa_memorder_from_tree (gimple_call_arg (stmt
, 2), &memorder
, &mmname
,
4960 gimple_location (stmt
)))
4963 /* Certain atomic insns must have Bx memory types. */
4966 case BRIG_ATOMIC_LD
:
4967 case BRIG_ATOMIC_ST
:
4968 case BRIG_ATOMIC_AND
:
4969 case BRIG_ATOMIC_OR
:
4970 case BRIG_ATOMIC_XOR
:
4971 case BRIG_ATOMIC_EXCH
:
4972 mtype
= hsa_bittype_for_type (mtype
);
4983 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4985 dest
= new hsa_op_reg (hsa_type
);
4986 opcode
= signal
? BRIG_OPCODE_SIGNAL
: BRIG_OPCODE_ATOMIC
;
4992 opcode
= signal
? BRIG_OPCODE_SIGNALNORET
: BRIG_OPCODE_ATOMICNORET
;
4996 if (acode
== BRIG_ATOMIC_ST
)
4998 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
4999 memorder
= BRIG_MEMORY_ORDER_SC_RELEASE
;
5001 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
5002 && memorder
!= BRIG_MEMORY_ORDER_SC_RELEASE
5003 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
5005 HSA_SORRY_ATV (gimple_location (stmt
),
5006 "support for HSA does not implement memory model for "
5007 "ATOMIC_ST: %s", mmname
);
5012 hsa_insn_basic
*atominsn
;
5016 atominsn
= new hsa_insn_signal (nops
, opcode
, acode
, mtype
, memorder
);
5017 tgt
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
5021 atominsn
= new hsa_insn_atomic (nops
, opcode
, acode
, mtype
, memorder
);
5022 hsa_op_address
*addr
;
5023 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5024 if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
== BRIG_SEGMENT_PRIVATE
)
5026 HSA_SORRY_AT (gimple_location (stmt
),
5027 "HSA does not implement atomic operations in private "
5034 hsa_op_with_type
*op
5035 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1), hbb
);
5038 atominsn
->set_op (0, dest
);
5039 atominsn
->set_op (1, tgt
);
5040 atominsn
->set_op (2, op
);
5044 atominsn
->set_op (0, tgt
);
5045 atominsn
->set_op (1, op
);
5048 hbb
->append_insn (atominsn
);
5050 /* HSA does not natively support the variants that return the modified value,
5051 so re-do the operation again non-atomically if that is what was
5053 if (lhs
&& !ret_orig
)
5058 case BRIG_ATOMIC_ADD
:
5059 arith
= BRIG_OPCODE_ADD
;
5061 case BRIG_ATOMIC_AND
:
5062 arith
= BRIG_OPCODE_AND
;
5064 case BRIG_ATOMIC_OR
:
5065 arith
= BRIG_OPCODE_OR
;
5067 case BRIG_ATOMIC_SUB
:
5068 arith
= BRIG_OPCODE_SUB
;
5070 case BRIG_ATOMIC_XOR
:
5071 arith
= BRIG_OPCODE_XOR
;
5076 hsa_op_reg
*real_dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5077 gen_hsa_binary_operation (arith
, real_dest
, dest
, op
, hbb
);
5081 /* Generate HSA instructions for an internal fn.
5082 Instructions will be appended to HBB, which also needs to be the
5083 corresponding structure to the basic_block of STMT. */
5086 gen_hsa_insn_for_internal_fn_call (gcall
*stmt
, hsa_bb
*hbb
)
5088 gcc_checking_assert (gimple_call_internal_fn (stmt
));
5089 internal_fn fn
= gimple_call_internal_fn (stmt
);
5091 bool is_float_type_p
= false;
5092 if (gimple_call_lhs (stmt
) != NULL
5093 && TREE_TYPE (gimple_call_lhs (stmt
)) == float_type_node
)
5094 is_float_type_p
= true;
5099 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
5103 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
5107 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
5111 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
5115 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_NRSQRT
, stmt
, hbb
);
5119 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
5124 if (is_float_type_p
)
5125 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
5127 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5133 if (is_float_type_p
)
5134 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
5136 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5143 if (is_float_type_p
)
5144 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
5146 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5153 if (is_float_type_p
)
5154 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
5156 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5161 gen_hsa_clrsb (stmt
, hbb
);
5165 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
5169 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
5173 gen_hsa_ffs (stmt
, hbb
);
5177 gen_hsa_parity (stmt
, hbb
);
5181 gen_hsa_popcount (stmt
, hbb
);
5185 gen_hsa_divmod (stmt
, hbb
);
5198 case IFN_SIGNIFICAND
:
5210 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
5214 HSA_SORRY_ATV (gimple_location (stmt
),
5215 "support for HSA does not implement internal function: %s",
5216 internal_fn_name (fn
));
5221 /* Generate HSA instructions for the given call statement STMT. Instructions
5222 will be appended to HBB. */
5225 gen_hsa_insns_for_call (gimple
*stmt
, hsa_bb
*hbb
)
5227 gcall
*call
= as_a
<gcall
*> (stmt
);
5228 tree lhs
= gimple_call_lhs (stmt
);
5231 if (gimple_call_internal_p (stmt
))
5233 gen_hsa_insn_for_internal_fn_call (call
, hbb
);
5237 if (!gimple_call_builtin_p (stmt
, BUILT_IN_NORMAL
))
5239 tree function_decl
= gimple_call_fndecl (stmt
);
5240 /* Prefetch pass can create type-mismatching prefetch builtin calls which
5241 fail the gimple_call_builtin_p test above. Handle them here. */
5242 if (DECL_BUILT_IN_CLASS (function_decl
)
5243 && DECL_FUNCTION_CODE (function_decl
) == BUILT_IN_PREFETCH
)
5246 if (function_decl
== NULL_TREE
)
5248 HSA_SORRY_AT (gimple_location (stmt
),
5249 "support for HSA does not implement indirect calls");
5253 if (hsa_callable_function_p (function_decl
))
5254 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5255 else if (!gen_hsa_insns_for_known_library_call (stmt
, hbb
))
5256 HSA_SORRY_AT (gimple_location (stmt
),
5257 "HSA supports only calls of functions marked with pragma "
5258 "omp declare target");
5262 tree fndecl
= gimple_call_fndecl (stmt
);
5263 enum built_in_function builtin
= DECL_FUNCTION_CODE (fndecl
);
5267 case BUILT_IN_FABSF
:
5268 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_ABS
, stmt
, hbb
);
5272 case BUILT_IN_CEILF
:
5273 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
5276 case BUILT_IN_FLOOR
:
5277 case BUILT_IN_FLOORF
:
5278 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
5282 case BUILT_IN_RINTF
:
5283 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
5287 case BUILT_IN_SQRTF
:
5288 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
5291 case BUILT_IN_TRUNC
:
5292 case BUILT_IN_TRUNCF
:
5293 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
5300 /* HSAIL does not provide an instruction for double argument type. */
5301 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5305 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
5308 case BUILT_IN_EXP2F
:
5309 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
5312 case BUILT_IN_LOG2F
:
5313 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
5317 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
5320 case BUILT_IN_CLRSB
:
5321 case BUILT_IN_CLRSBL
:
5322 case BUILT_IN_CLRSBLL
:
5323 gen_hsa_clrsb (call
, hbb
);
5328 case BUILT_IN_CLZLL
:
5329 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
5334 case BUILT_IN_CTZLL
:
5335 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
5340 case BUILT_IN_FFSLL
:
5341 gen_hsa_ffs (call
, hbb
);
5344 case BUILT_IN_PARITY
:
5345 case BUILT_IN_PARITYL
:
5346 case BUILT_IN_PARITYLL
:
5347 gen_hsa_parity (call
, hbb
);
5350 case BUILT_IN_POPCOUNT
:
5351 case BUILT_IN_POPCOUNTL
:
5352 case BUILT_IN_POPCOUNTLL
:
5353 gen_hsa_popcount (call
, hbb
);
5356 case BUILT_IN_ATOMIC_LOAD_1
:
5357 case BUILT_IN_ATOMIC_LOAD_2
:
5358 case BUILT_IN_ATOMIC_LOAD_4
:
5359 case BUILT_IN_ATOMIC_LOAD_8
:
5360 case BUILT_IN_ATOMIC_LOAD_16
:
5364 src
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5366 BrigMemoryOrder memorder
;
5368 if (hsa_memorder_from_tree (gimple_call_arg (stmt
, 1), &memorder
,
5369 &mmname
, gimple_location (stmt
)))
5372 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
5373 memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
5375 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
5376 && memorder
!= BRIG_MEMORY_ORDER_SC_ACQUIRE
5377 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
5379 HSA_SORRY_ATV (gimple_location (stmt
),
5380 "support for HSA does not implement "
5381 "memory model for atomic loads: %s", mmname
);
5387 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
5389 mtype
= mem_type_for_type (t
);
5390 mtype
= hsa_bittype_for_type (mtype
);
5391 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5395 mtype
= BRIG_TYPE_B64
;
5396 dest
= new hsa_op_reg (mtype
);
5399 hsa_insn_basic
*atominsn
;
5400 atominsn
= new hsa_insn_atomic (2, BRIG_OPCODE_ATOMIC
, BRIG_ATOMIC_LD
,
5401 mtype
, memorder
, dest
, src
);
5403 hbb
->append_insn (atominsn
);
5407 case BUILT_IN_ATOMIC_EXCHANGE_1
:
5408 case BUILT_IN_ATOMIC_EXCHANGE_2
:
5409 case BUILT_IN_ATOMIC_EXCHANGE_4
:
5410 case BUILT_IN_ATOMIC_EXCHANGE_8
:
5411 case BUILT_IN_ATOMIC_EXCHANGE_16
:
5412 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_EXCH
, stmt
, hbb
, false);
5416 case BUILT_IN_ATOMIC_FETCH_ADD_1
:
5417 case BUILT_IN_ATOMIC_FETCH_ADD_2
:
5418 case BUILT_IN_ATOMIC_FETCH_ADD_4
:
5419 case BUILT_IN_ATOMIC_FETCH_ADD_8
:
5420 case BUILT_IN_ATOMIC_FETCH_ADD_16
:
5421 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ADD
, stmt
, hbb
, false);
5425 case BUILT_IN_ATOMIC_FETCH_SUB_1
:
5426 case BUILT_IN_ATOMIC_FETCH_SUB_2
:
5427 case BUILT_IN_ATOMIC_FETCH_SUB_4
:
5428 case BUILT_IN_ATOMIC_FETCH_SUB_8
:
5429 case BUILT_IN_ATOMIC_FETCH_SUB_16
:
5430 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_SUB
, stmt
, hbb
, false);
5434 case BUILT_IN_ATOMIC_FETCH_AND_1
:
5435 case BUILT_IN_ATOMIC_FETCH_AND_2
:
5436 case BUILT_IN_ATOMIC_FETCH_AND_4
:
5437 case BUILT_IN_ATOMIC_FETCH_AND_8
:
5438 case BUILT_IN_ATOMIC_FETCH_AND_16
:
5439 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_AND
, stmt
, hbb
, false);
5443 case BUILT_IN_ATOMIC_FETCH_XOR_1
:
5444 case BUILT_IN_ATOMIC_FETCH_XOR_2
:
5445 case BUILT_IN_ATOMIC_FETCH_XOR_4
:
5446 case BUILT_IN_ATOMIC_FETCH_XOR_8
:
5447 case BUILT_IN_ATOMIC_FETCH_XOR_16
:
5448 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_XOR
, stmt
, hbb
, false);
5452 case BUILT_IN_ATOMIC_FETCH_OR_1
:
5453 case BUILT_IN_ATOMIC_FETCH_OR_2
:
5454 case BUILT_IN_ATOMIC_FETCH_OR_4
:
5455 case BUILT_IN_ATOMIC_FETCH_OR_8
:
5456 case BUILT_IN_ATOMIC_FETCH_OR_16
:
5457 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_OR
, stmt
, hbb
, false);
5461 case BUILT_IN_ATOMIC_STORE_1
:
5462 case BUILT_IN_ATOMIC_STORE_2
:
5463 case BUILT_IN_ATOMIC_STORE_4
:
5464 case BUILT_IN_ATOMIC_STORE_8
:
5465 case BUILT_IN_ATOMIC_STORE_16
:
5466 /* Since there cannot be any LHS, the first parameter is meaningless. */
5467 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ST
, stmt
, hbb
, false);
5471 case BUILT_IN_ATOMIC_ADD_FETCH_1
:
5472 case BUILT_IN_ATOMIC_ADD_FETCH_2
:
5473 case BUILT_IN_ATOMIC_ADD_FETCH_4
:
5474 case BUILT_IN_ATOMIC_ADD_FETCH_8
:
5475 case BUILT_IN_ATOMIC_ADD_FETCH_16
:
5476 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_ADD
, stmt
, hbb
, false);
5479 case BUILT_IN_ATOMIC_SUB_FETCH_1
:
5480 case BUILT_IN_ATOMIC_SUB_FETCH_2
:
5481 case BUILT_IN_ATOMIC_SUB_FETCH_4
:
5482 case BUILT_IN_ATOMIC_SUB_FETCH_8
:
5483 case BUILT_IN_ATOMIC_SUB_FETCH_16
:
5484 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_SUB
, stmt
, hbb
, false);
5487 case BUILT_IN_ATOMIC_AND_FETCH_1
:
5488 case BUILT_IN_ATOMIC_AND_FETCH_2
:
5489 case BUILT_IN_ATOMIC_AND_FETCH_4
:
5490 case BUILT_IN_ATOMIC_AND_FETCH_8
:
5491 case BUILT_IN_ATOMIC_AND_FETCH_16
:
5492 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_AND
, stmt
, hbb
, false);
5495 case BUILT_IN_ATOMIC_XOR_FETCH_1
:
5496 case BUILT_IN_ATOMIC_XOR_FETCH_2
:
5497 case BUILT_IN_ATOMIC_XOR_FETCH_4
:
5498 case BUILT_IN_ATOMIC_XOR_FETCH_8
:
5499 case BUILT_IN_ATOMIC_XOR_FETCH_16
:
5500 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_XOR
, stmt
, hbb
, false);
5503 case BUILT_IN_ATOMIC_OR_FETCH_1
:
5504 case BUILT_IN_ATOMIC_OR_FETCH_2
:
5505 case BUILT_IN_ATOMIC_OR_FETCH_4
:
5506 case BUILT_IN_ATOMIC_OR_FETCH_8
:
5507 case BUILT_IN_ATOMIC_OR_FETCH_16
:
5508 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_OR
, stmt
, hbb
, false);
5511 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_1
:
5512 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_2
:
5513 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_4
:
5514 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8
:
5515 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_16
:
5517 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
5519 = hsa_bittype_for_type (hsa_type_for_scalar_tree_type (type
, false));
5520 BrigMemoryOrder memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
5521 hsa_insn_basic
*atominsn
;
5523 atominsn
= new hsa_insn_atomic (4, BRIG_OPCODE_ATOMIC
,
5524 BRIG_ATOMIC_CAS
, atype
, memorder
);
5525 tgt
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5528 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5530 dest
= new hsa_op_reg (atype
);
5532 atominsn
->set_op (0, dest
);
5533 atominsn
->set_op (1, tgt
);
5535 hsa_op_with_type
*op
5536 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1), hbb
);
5537 atominsn
->set_op (2, op
);
5538 op
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 2), hbb
);
5539 atominsn
->set_op (3, op
);
5541 hbb
->append_insn (atominsn
);
5545 case BUILT_IN_HSA_WORKGROUPID
:
5546 query_hsa_grid_dim (stmt
, BRIG_OPCODE_WORKGROUPID
, hbb
);
5548 case BUILT_IN_HSA_WORKITEMID
:
5549 query_hsa_grid_dim (stmt
, BRIG_OPCODE_WORKITEMID
, hbb
);
5551 case BUILT_IN_HSA_WORKITEMABSID
:
5552 query_hsa_grid_dim (stmt
, BRIG_OPCODE_WORKITEMABSID
, hbb
);
5554 case BUILT_IN_HSA_GRIDSIZE
:
5555 query_hsa_grid_dim (stmt
, BRIG_OPCODE_GRIDSIZE
, hbb
);
5557 case BUILT_IN_HSA_CURRENTWORKGROUPSIZE
:
5558 query_hsa_grid_dim (stmt
, BRIG_OPCODE_CURRENTWORKGROUPSIZE
, hbb
);
5561 case BUILT_IN_GOMP_BARRIER
:
5562 hbb
->append_insn (new hsa_insn_br (0, BRIG_OPCODE_BARRIER
, BRIG_TYPE_NONE
,
5565 case BUILT_IN_GOMP_PARALLEL
:
5566 HSA_SORRY_AT (gimple_location (stmt
),
5567 "support for HSA does not implement non-gridified "
5568 "OpenMP parallel constructs.");
5571 case BUILT_IN_OMP_GET_THREAD_NUM
:
5573 query_hsa_grid_nodim (stmt
, BRIG_OPCODE_WORKITEMFLATABSID
, hbb
);
5577 case BUILT_IN_OMP_GET_NUM_THREADS
:
5579 gen_get_num_threads (stmt
, hbb
);
5582 case BUILT_IN_GOMP_TEAMS
:
5584 gen_set_num_threads (gimple_call_arg (stmt
, 1), hbb
);
5587 case BUILT_IN_OMP_GET_NUM_TEAMS
:
5589 gen_get_num_teams (stmt
, hbb
);
5592 case BUILT_IN_OMP_GET_TEAM_NUM
:
5594 gen_get_team_num (stmt
, hbb
);
5597 case BUILT_IN_MEMCPY
:
5598 case BUILT_IN_MEMPCPY
:
5600 expand_memory_copy (stmt
, hbb
, builtin
);
5603 case BUILT_IN_MEMSET
:
5605 tree c
= gimple_call_arg (stmt
, 1);
5607 if (TREE_CODE (c
) != INTEGER_CST
)
5609 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5613 tree byte_size
= gimple_call_arg (stmt
, 2);
5615 if (!tree_fits_uhwi_p (byte_size
))
5617 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5621 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
5623 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5625 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5629 unsigned HOST_WIDE_INT constant
5630 = tree_to_uhwi (fold_convert (unsigned_char_type_node
, c
));
5632 expand_memory_set (stmt
, n
, constant
, hbb
, builtin
);
5636 case BUILT_IN_BZERO
:
5638 tree byte_size
= gimple_call_arg (stmt
, 1);
5640 if (!tree_fits_uhwi_p (byte_size
))
5642 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5646 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
5648 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5650 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5654 expand_memory_set (stmt
, n
, 0, hbb
, builtin
);
5658 CASE_BUILT_IN_ALLOCA
:
5660 gen_hsa_alloca (call
, hbb
);
5663 case BUILT_IN_PREFETCH
:
5667 tree name_tree
= DECL_NAME (fndecl
);
5668 const char *s
= IDENTIFIER_POINTER (name_tree
);
5669 size_t len
= strlen (s
);
5670 if (len
> 4 && (strncmp (s
, "__builtin_GOMP_", 15) == 0))
5671 HSA_SORRY_ATV (gimple_location (stmt
),
5672 "support for HSA does not implement GOMP function %s",
5675 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5681 /* Generate HSA instructions for a given gimple statement. Instructions will be
5685 gen_hsa_insns_for_gimple_stmt (gimple
*stmt
, hsa_bb
*hbb
)
5687 switch (gimple_code (stmt
))
5690 if (gimple_clobber_p (stmt
))
5693 if (gimple_assign_single_p (stmt
))
5695 tree lhs
= gimple_assign_lhs (stmt
);
5696 tree rhs
= gimple_assign_rhs1 (stmt
);
5697 gen_hsa_insns_for_single_assignment (lhs
, rhs
, hbb
);
5700 gen_hsa_insns_for_operation_assignment (stmt
, hbb
);
5703 gen_hsa_insns_for_return (as_a
<greturn
*> (stmt
), hbb
);
5706 gen_hsa_insns_for_cond_stmt (stmt
, hbb
);
5709 gen_hsa_insns_for_call (stmt
, hbb
);
5712 /* ??? HSA supports some debug facilities. */
5716 tree label
= gimple_label_label (as_a
<glabel
*> (stmt
));
5717 if (FORCED_LABEL (label
))
5718 HSA_SORRY_AT (gimple_location (stmt
),
5719 "support for HSA does not implement gimple label with "
5726 hbb
->append_insn (new hsa_insn_basic (0, BRIG_OPCODE_NOP
));
5731 gen_hsa_insns_for_switch_stmt (as_a
<gswitch
*> (stmt
), hbb
);
5735 HSA_SORRY_ATV (gimple_location (stmt
),
5736 "support for HSA does not implement gimple statement %s",
5737 gimple_code_name
[(int) gimple_code (stmt
)]);
5741 /* Generate a HSA PHI from a gimple PHI. */
5744 gen_hsa_phi_from_gimple_phi (gimple
*phi_stmt
, hsa_bb
*hbb
)
5747 unsigned count
= gimple_phi_num_args (phi_stmt
);
5750 = hsa_cfun
->reg_for_gimple_ssa (gimple_phi_result (phi_stmt
));
5751 hphi
= new hsa_insn_phi (count
, dest
);
5752 hphi
->m_bb
= hbb
->m_bb
;
5754 auto_vec
<tree
, 8> aexprs
;
5755 auto_vec
<hsa_op_reg
*, 8> aregs
;
5757 /* Calling split_edge when processing a PHI node messes up with the order of
5758 gimple phi node arguments (it moves the one associated with the edge to
5759 the end). We need to keep the order of edges and arguments of HSA phi
5760 node arguments consistent, so we do all required splitting as the first
5761 step, and in reverse order as to not be affected by the re-orderings. */
5762 for (unsigned j
= count
; j
!= 0; j
--)
5765 tree op
= gimple_phi_arg_def (phi_stmt
, i
);
5766 if (TREE_CODE (op
) != ADDR_EXPR
)
5769 edge e
= gimple_phi_arg_edge (as_a
<gphi
*> (phi_stmt
), i
);
5770 hsa_bb
*hbb_src
= hsa_init_new_bb (split_edge (e
));
5771 hsa_op_address
*addr
= gen_hsa_addr (TREE_OPERAND (op
, 0),
5775 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
5776 hsa_insn_basic
*insn
5777 = new hsa_insn_basic (2, BRIG_OPCODE_LDA
, BRIG_TYPE_U64
,
5779 hbb_src
->append_insn (insn
);
5780 aexprs
.safe_push (op
);
5781 aregs
.safe_push (dest
);
5784 tree lhs
= gimple_phi_result (phi_stmt
);
5785 for (unsigned i
= 0; i
< count
; i
++)
5787 tree op
= gimple_phi_arg_def (phi_stmt
, i
);
5789 if (TREE_CODE (op
) == SSA_NAME
)
5791 hsa_op_reg
*hreg
= hsa_cfun
->reg_for_gimple_ssa (op
);
5792 hphi
->set_op (i
, hreg
);
5796 gcc_assert (is_gimple_min_invariant (op
));
5797 tree t
= TREE_TYPE (op
);
5798 if (!POINTER_TYPE_P (t
)
5799 || (TREE_CODE (op
) == STRING_CST
5800 && TREE_CODE (TREE_TYPE (t
)) == INTEGER_TYPE
))
5801 hphi
->set_op (i
, new hsa_op_immed (op
));
5802 else if (POINTER_TYPE_P (TREE_TYPE (lhs
))
5803 && TREE_CODE (op
) == INTEGER_CST
)
5805 /* Handle assignment of NULL value to a pointer type. */
5806 hphi
->set_op (i
, new hsa_op_immed (op
));
5808 else if (TREE_CODE (op
) == ADDR_EXPR
)
5810 hsa_op_reg
*dest
= NULL
;
5811 for (unsigned a_idx
= 0; a_idx
< aexprs
.length (); a_idx
++)
5812 if (aexprs
[a_idx
] == op
)
5814 dest
= aregs
[a_idx
];
5818 hphi
->set_op (i
, dest
);
5822 HSA_SORRY_AT (gimple_location (phi_stmt
),
5823 "support for HSA does not handle PHI nodes with "
5824 "constant address operands");
5830 hbb
->append_phi (hphi
);
5833 /* Constructor of class containing HSA-specific information about a basic
5834 block. CFG_BB is the CFG BB this HSA BB is associated with. IDX is the new
5835 index of this BB (so that the constructor does not attempt to use
5836 hsa_cfun during its construction). */
5838 hsa_bb::hsa_bb (basic_block cfg_bb
, int idx
)
5839 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5840 m_last_phi (NULL
), m_index (idx
)
5842 gcc_assert (!cfg_bb
->aux
);
5846 /* Constructor of class containing HSA-specific information about a basic
5847 block. CFG_BB is the CFG BB this HSA BB is associated with. */
5849 hsa_bb::hsa_bb (basic_block cfg_bb
)
5850 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5851 m_last_phi (NULL
), m_index (hsa_cfun
->m_hbb_count
++)
5853 gcc_assert (!cfg_bb
->aux
);
5857 /* Create and initialize and return a new hsa_bb structure for a given CFG
5861 hsa_init_new_bb (basic_block bb
)
5863 void *m
= obstack_alloc (&hsa_obstack
, sizeof (hsa_bb
));
5864 return new (m
) hsa_bb (bb
);
5867 /* Initialize OMP in an HSA basic block PROLOGUE. */
5870 init_prologue (void)
5872 if (!hsa_cfun
->m_kern_p
)
5875 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5877 /* Create a magic number that is going to be printed by libgomp. */
5878 unsigned index
= hsa_get_number_decl_kernel_mappings ();
5880 /* Emit store to debug argument. */
5881 if (PARAM_VALUE (PARAM_HSA_GEN_DEBUG_STORES
) > 0)
5882 set_debug_value (prologue
, new hsa_op_immed (1000 + index
, BRIG_TYPE_U64
));
5885 /* Initialize hsa_num_threads to a default value. */
5888 init_hsa_num_threads (void)
5890 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5892 /* Save the default value to private variable hsa_num_threads. */
5893 hsa_insn_basic
*basic
5894 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
,
5895 new hsa_op_immed (0, hsa_num_threads
->m_type
),
5896 new hsa_op_address (hsa_num_threads
));
5897 prologue
->append_insn (basic
);
5900 /* Go over gimple representation and generate our internal HSA one. */
5903 gen_body_from_gimple ()
5907 /* Verify CFG for complex edges we are unable to handle. */
5911 FOR_EACH_BB_FN (bb
, cfun
)
5913 FOR_EACH_EDGE (e
, ei
, bb
->succs
)
5915 /* Verify all unsupported flags for edges that point
5916 to the same basic block. */
5917 if (e
->flags
& EDGE_EH
)
5919 HSA_SORRY_AT (UNKNOWN_LOCATION
,
5920 "support for HSA does not implement exception "
5927 FOR_EACH_BB_FN (bb
, cfun
)
5929 gimple_stmt_iterator gsi
;
5930 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5934 hbb
= hsa_init_new_bb (bb
);
5936 for (gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5938 gen_hsa_insns_for_gimple_stmt (gsi_stmt (gsi
), hbb
);
5939 if (hsa_seen_error ())
5944 FOR_EACH_BB_FN (bb
, cfun
)
5946 gimple_stmt_iterator gsi
;
5947 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5948 gcc_assert (hbb
!= NULL
);
5950 for (gsi
= gsi_start_phis (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5951 if (!virtual_operand_p (gimple_phi_result (gsi_stmt (gsi
))))
5952 gen_hsa_phi_from_gimple_phi (gsi_stmt (gsi
), hbb
);
5955 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
5957 fprintf (dump_file
, "------- Generated SSA form -------\n");
5958 dump_hsa_cfun (dump_file
);
5963 gen_function_decl_parameters (hsa_function_representation
*f
,
5969 for (parm
= TYPE_ARG_TYPES (TREE_TYPE (decl
)), i
= 0;
5971 parm
= TREE_CHAIN (parm
), i
++)
5973 /* Result type if last in the tree list. */
5974 if (TREE_CHAIN (parm
) == NULL
)
5977 tree v
= TREE_VALUE (parm
);
5979 hsa_symbol
*arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5981 arg
->m_type
= hsa_type_for_tree_type (v
, &arg
->m_dim
);
5982 arg
->m_name_number
= i
;
5984 f
->m_input_args
.safe_push (arg
);
5987 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
5988 if (!VOID_TYPE_P (result_type
))
5990 f
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5992 f
->m_output_arg
->m_type
5993 = hsa_type_for_tree_type (result_type
, &f
->m_output_arg
->m_dim
);
5994 f
->m_output_arg
->m_name
= "res";
5998 /* Generate the vector of parameters of the HSA representation of the current
5999 function. This also includes the output parameter representing the
6003 gen_function_def_parameters ()
6007 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
6009 for (parm
= DECL_ARGUMENTS (cfun
->decl
); parm
;
6010 parm
= DECL_CHAIN (parm
))
6012 struct hsa_symbol
**slot
;
6015 = new hsa_symbol (BRIG_TYPE_NONE
, hsa_cfun
->m_kern_p
6016 ? BRIG_SEGMENT_KERNARG
: BRIG_SEGMENT_ARG
,
6017 BRIG_LINKAGE_FUNCTION
);
6018 arg
->fillup_for_decl (parm
);
6020 hsa_cfun
->m_input_args
.safe_push (arg
);
6022 if (hsa_seen_error ())
6025 arg
->m_name
= hsa_get_declaration_name (parm
);
6027 /* Copy all input arguments and create corresponding private symbols
6029 hsa_symbol
*private_arg
;
6030 hsa_op_address
*parm_addr
= new hsa_op_address (arg
);
6032 if (TREE_ADDRESSABLE (parm
)
6033 || (!is_gimple_reg (parm
) && !TREE_READONLY (parm
)))
6035 private_arg
= hsa_cfun
->create_hsa_temporary (arg
->m_type
);
6036 private_arg
->fillup_for_decl (parm
);
6038 BrigAlignment8_t align
= MIN (arg
->m_align
, private_arg
->m_align
);
6040 hsa_op_address
*private_arg_addr
= new hsa_op_address (private_arg
);
6041 gen_hsa_memory_copy (prologue
, private_arg_addr
, parm_addr
,
6042 arg
->total_byte_size (), align
);
6047 slot
= hsa_cfun
->m_local_symbols
->find_slot (private_arg
, INSERT
);
6048 gcc_assert (!*slot
);
6049 *slot
= private_arg
;
6051 if (is_gimple_reg (parm
))
6053 tree ddef
= ssa_default_def (cfun
, parm
);
6054 if (ddef
&& !has_zero_uses (ddef
))
6056 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (ddef
),
6058 BrigType16_t mtype
= mem_type_for_type (t
);
6059 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (ddef
);
6060 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
,
6062 gcc_assert (!parm_addr
->m_reg
);
6063 prologue
->append_insn (mem
);
6068 if (!VOID_TYPE_P (TREE_TYPE (TREE_TYPE (cfun
->decl
))))
6070 struct hsa_symbol
**slot
;
6072 hsa_cfun
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
6073 BRIG_LINKAGE_FUNCTION
);
6074 hsa_cfun
->m_output_arg
->fillup_for_decl (DECL_RESULT (cfun
->decl
));
6076 if (hsa_seen_error ())
6079 hsa_cfun
->m_output_arg
->m_name
= "res";
6080 slot
= hsa_cfun
->m_local_symbols
->find_slot (hsa_cfun
->m_output_arg
,
6082 gcc_assert (!*slot
);
6083 *slot
= hsa_cfun
->m_output_arg
;
6087 /* Generate function representation that corresponds to
6088 a function declaration. */
6090 hsa_function_representation
*
6091 hsa_generate_function_declaration (tree decl
)
6093 hsa_function_representation
*fun
6094 = new hsa_function_representation (decl
, false, 0);
6096 fun
->m_declaration_p
= true;
6097 fun
->m_name
= get_brig_function_name (decl
);
6098 gen_function_decl_parameters (fun
, decl
);
6104 /* Generate function representation that corresponds to
6107 hsa_function_representation
*
6108 hsa_generate_internal_fn_decl (hsa_internal_fn
*fn
)
6110 hsa_function_representation
*fun
= new hsa_function_representation (fn
);
6112 fun
->m_name
= fn
->name ();
6114 for (unsigned i
= 0; i
< fn
->get_arity (); i
++)
6117 = new hsa_symbol (fn
->get_argument_type (i
), BRIG_SEGMENT_ARG
,
6119 arg
->m_name_number
= i
;
6120 fun
->m_input_args
.safe_push (arg
);
6123 fun
->m_output_arg
= new hsa_symbol (fn
->get_argument_type (-1),
6124 BRIG_SEGMENT_ARG
, BRIG_LINKAGE_NONE
);
6125 fun
->m_output_arg
->m_name
= "res";
6130 /* Return true if switch statement S can be transformed
6131 to a SBR instruction in HSAIL. */
6134 transformable_switch_to_sbr_p (gswitch
*s
)
6136 /* Identify if a switch statement can be transformed to
6137 SBR instruction, like:
6139 sbr_u32 $s1 [@label1, @label2, @label3];
6142 tree size
= get_switch_size (s
);
6143 if (!tree_fits_uhwi_p (size
))
6146 if (tree_to_uhwi (size
) > HSA_MAXIMUM_SBR_LABELS
)
6152 /* Structure hold connection between PHI nodes and immediate
6153 values hold by there nodes. */
6155 struct phi_definition
6157 phi_definition (unsigned phi_i
, unsigned label_i
, tree imm
):
6158 phi_index (phi_i
), label_index (label_i
), phi_value (imm
)
6162 unsigned label_index
;
6166 /* Sum slice of a vector V, starting from index START and ending
6167 at the index END - 1. */
6169 template <typename T
>
6171 T
sum_slice (const auto_vec
<T
> &v
, unsigned start
, unsigned end
,
6176 for (unsigned i
= start
; i
< end
; i
++)
6182 /* Function transforms GIMPLE SWITCH statements to a series of IF statements.
6183 Let's assume following example:
6197 The transformation encompasses following steps:
6198 1) all immediate values used by edges coming from the switch basic block
6200 2) all these edges are removed
6201 3) the switch statement (in L0) is replaced by:
6207 4) newly created basic block Lx' is used for generation of
6209 5) else branch of the last condition goes to LD
6210 6) fix all immediate values in PHI nodes that were propagated though
6211 edges that were removed in step 2
6213 Note: if a case is made by a range C1..C2, then process
6214 following transformation:
6216 switch_cond_op1 = C1 <= index;
6217 switch_cond_op2 = index <= C2;
6218 switch_cond_and = switch_cond_op1 & switch_cond_op2;
6219 if (switch_cond_and != 0)
6227 convert_switch_statements (void)
6229 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
6232 bool modified_cfg
= false;
6234 FOR_EACH_BB_FN (bb
, func
)
6236 gimple_stmt_iterator gsi
= gsi_last_bb (bb
);
6237 if (gsi_end_p (gsi
))
6240 gimple
*stmt
= gsi_stmt (gsi
);
6242 if (gimple_code (stmt
) == GIMPLE_SWITCH
)
6244 gswitch
*s
= as_a
<gswitch
*> (stmt
);
6246 /* If the switch can utilize SBR insn, skip the statement. */
6247 if (transformable_switch_to_sbr_p (s
))
6250 modified_cfg
= true;
6252 unsigned labels
= gimple_switch_num_labels (s
);
6253 tree index
= gimple_switch_index (s
);
6254 tree index_type
= TREE_TYPE (index
);
6255 tree default_label
= gimple_switch_default_label (s
);
6256 basic_block default_label_bb
6257 = label_to_block_fn (func
, CASE_LABEL (default_label
));
6258 basic_block cur_bb
= bb
;
6260 auto_vec
<edge
> new_edges
;
6261 auto_vec
<phi_definition
*> phi_todo_list
;
6262 auto_vec
<profile_count
> edge_counts
;
6263 auto_vec
<profile_probability
> edge_probabilities
;
6265 /* Investigate all labels that and PHI nodes in these edges which
6266 should be fixed after we add new collection of edges. */
6267 for (unsigned i
= 0; i
< labels
; i
++)
6269 tree label
= gimple_switch_label (s
, i
);
6270 basic_block label_bb
= label_to_block_fn (func
, CASE_LABEL (label
));
6271 edge e
= find_edge (bb
, label_bb
);
6272 edge_counts
.safe_push (e
->count ());
6273 edge_probabilities
.safe_push (e
->probability
);
6274 gphi_iterator phi_gsi
;
6276 /* Save PHI definitions that will be destroyed because of an edge
6277 is going to be removed. */
6278 unsigned phi_index
= 0;
6279 for (phi_gsi
= gsi_start_phis (e
->dest
);
6280 !gsi_end_p (phi_gsi
); gsi_next (&phi_gsi
))
6282 gphi
*phi
= phi_gsi
.phi ();
6283 for (unsigned j
= 0; j
< gimple_phi_num_args (phi
); j
++)
6285 if (gimple_phi_arg_edge (phi
, j
) == e
)
6287 tree imm
= gimple_phi_arg_def (phi
, j
);
6288 phi_definition
*p
= new phi_definition (phi_index
, i
,
6290 phi_todo_list
.safe_push (p
);
6298 /* Remove all edges for the current basic block. */
6299 for (int i
= EDGE_COUNT (bb
->succs
) - 1; i
>= 0; i
--)
6301 edge e
= EDGE_SUCC (bb
, i
);
6305 /* Iterate all non-default labels. */
6306 for (unsigned i
= 1; i
< labels
; i
++)
6308 tree label
= gimple_switch_label (s
, i
);
6309 tree low
= CASE_LOW (label
);
6310 tree high
= CASE_HIGH (label
);
6312 if (!useless_type_conversion_p (TREE_TYPE (low
), index_type
))
6313 low
= fold_convert (index_type
, low
);
6315 gimple_stmt_iterator cond_gsi
= gsi_last_bb (cur_bb
);
6319 tree tmp1
= make_temp_ssa_name (boolean_type_node
, NULL
,
6322 gimple
*assign1
= gimple_build_assign (tmp1
, LE_EXPR
, low
,
6325 tree tmp2
= make_temp_ssa_name (boolean_type_node
, NULL
,
6328 if (!useless_type_conversion_p (TREE_TYPE (high
), index_type
))
6329 high
= fold_convert (index_type
, high
);
6330 gimple
*assign2
= gimple_build_assign (tmp2
, LE_EXPR
, index
,
6333 tree tmp3
= make_temp_ssa_name (boolean_type_node
, NULL
,
6335 gimple
*assign3
= gimple_build_assign (tmp3
, BIT_AND_EXPR
, tmp1
,
6338 gsi_insert_before (&cond_gsi
, assign1
, GSI_SAME_STMT
);
6339 gsi_insert_before (&cond_gsi
, assign2
, GSI_SAME_STMT
);
6340 gsi_insert_before (&cond_gsi
, assign3
, GSI_SAME_STMT
);
6342 tree b
= constant_boolean_node (false, boolean_type_node
);
6343 c
= gimple_build_cond (NE_EXPR
, tmp3
, b
, NULL
, NULL
);
6346 c
= gimple_build_cond (EQ_EXPR
, index
, low
, NULL
, NULL
);
6348 gimple_set_location (c
, gimple_location (stmt
));
6350 gsi_insert_before (&cond_gsi
, c
, GSI_SAME_STMT
);
6352 basic_block label_bb
6353 = label_to_block_fn (func
, CASE_LABEL (label
));
6354 edge new_edge
= make_edge (cur_bb
, label_bb
, EDGE_TRUE_VALUE
);
6355 profile_probability prob_sum
= sum_slice
<profile_probability
>
6356 (edge_probabilities
, i
, labels
, profile_probability::never ())
6357 + edge_probabilities
[0];
6359 if (prob_sum
.initialized_p ())
6360 new_edge
->probability
= edge_probabilities
[i
] / prob_sum
;
6362 new_edges
.safe_push (new_edge
);
6366 /* Prepare another basic block that will contain
6368 basic_block next_bb
= create_empty_bb (cur_bb
);
6371 add_bb_to_loop (next_bb
, cur_bb
->loop_father
);
6372 loops_state_set (LOOPS_NEED_FIXUP
);
6375 edge next_edge
= make_edge (cur_bb
, next_bb
, EDGE_FALSE_VALUE
);
6376 next_edge
->probability
= new_edge
->probability
.invert ();
6377 next_bb
->count
= next_edge
->count ();
6380 else /* Link last IF statement and default label
6383 edge e
= make_edge (cur_bb
, default_label_bb
, EDGE_FALSE_VALUE
);
6384 e
->probability
= new_edge
->probability
.invert ();
6385 new_edges
.safe_insert (0, e
);
6389 /* Restore original PHI immediate value. */
6390 for (unsigned i
= 0; i
< phi_todo_list
.length (); i
++)
6392 phi_definition
*phi_def
= phi_todo_list
[i
];
6393 edge new_edge
= new_edges
[phi_def
->label_index
];
6395 gphi_iterator it
= gsi_start_phis (new_edge
->dest
);
6396 for (unsigned i
= 0; i
< phi_def
->phi_index
; i
++)
6399 gphi
*phi
= it
.phi ();
6400 add_phi_arg (phi
, phi_def
->phi_value
, new_edge
, UNKNOWN_LOCATION
);
6404 /* Remove the original GIMPLE switch statement. */
6405 gsi_remove (&gsi
, true);
6410 dump_function_to_file (current_function_decl
, dump_file
, TDF_DETAILS
);
6412 return modified_cfg
;
6415 /* Expand builtins that can't be handled by HSA back-end. */
6420 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
6423 FOR_EACH_BB_FN (bb
, func
)
6425 for (gimple_stmt_iterator gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
);
6428 gimple
*stmt
= gsi_stmt (gsi
);
6430 if (gimple_code (stmt
) != GIMPLE_CALL
)
6433 gcall
*call
= as_a
<gcall
*> (stmt
);
6435 if (!gimple_call_builtin_p (call
, BUILT_IN_NORMAL
))
6438 tree fndecl
= gimple_call_fndecl (stmt
);
6439 enum built_in_function fn
= DECL_FUNCTION_CODE (fndecl
);
6442 case BUILT_IN_CEXPF
:
6443 case BUILT_IN_CEXPIF
:
6444 case BUILT_IN_CEXPI
:
6446 /* Similar to builtins.c (expand_builtin_cexpi), the builtin
6447 can be transformed to: cexp(I * z) = ccos(z) + I * csin(z). */
6448 tree lhs
= gimple_call_lhs (stmt
);
6449 tree rhs
= gimple_call_arg (stmt
, 0);
6450 tree rhs_type
= TREE_TYPE (rhs
);
6451 bool float_type_p
= rhs_type
== float_type_node
;
6452 tree real_part
= make_temp_ssa_name (rhs_type
, NULL
,
6454 tree imag_part
= make_temp_ssa_name (rhs_type
, NULL
,
6458 = mathfn_built_in (rhs_type
, fn
== float_type_p
6459 ? BUILT_IN_COSF
: BUILT_IN_COS
);
6460 gcall
*cos
= gimple_build_call (cos_fndecl
, 1, rhs
);
6461 gimple_call_set_lhs (cos
, real_part
);
6462 gsi_insert_before (&gsi
, cos
, GSI_SAME_STMT
);
6465 = mathfn_built_in (rhs_type
, fn
== float_type_p
6466 ? BUILT_IN_SINF
: BUILT_IN_SIN
);
6467 gcall
*sin
= gimple_build_call (sin_fndecl
, 1, rhs
);
6468 gimple_call_set_lhs (sin
, imag_part
);
6469 gsi_insert_before (&gsi
, sin
, GSI_SAME_STMT
);
6472 gassign
*assign
= gimple_build_assign (lhs
, COMPLEX_EXPR
,
6473 real_part
, imag_part
);
6474 gsi_insert_before (&gsi
, assign
, GSI_SAME_STMT
);
6475 gsi_remove (&gsi
, true);
6486 /* Emit HSA module variables that are global for the entire module. */
6489 emit_hsa_module_variables (void)
6491 hsa_num_threads
= new hsa_symbol (BRIG_TYPE_U32
, BRIG_SEGMENT_PRIVATE
,
6492 BRIG_LINKAGE_MODULE
, true);
6494 hsa_num_threads
->m_name
= "hsa_num_threads";
6496 hsa_brig_emit_omp_symbols ();
6499 /* Generate HSAIL representation of the current function and write into a
6500 special section of the output file. If KERNEL is set, the function will be
6501 considered an HSA kernel callable from the host, otherwise it will be
6502 compiled as an HSA function callable from other HSA code. */
6505 generate_hsa (bool kernel
)
6507 hsa_init_data_for_cfun ();
6509 if (hsa_num_threads
== NULL
)
6510 emit_hsa_module_variables ();
6512 bool modified_cfg
= convert_switch_statements ();
6513 /* Initialize hsa_cfun. */
6514 hsa_cfun
= new hsa_function_representation (cfun
->decl
, kernel
,
6515 SSANAMES (cfun
)->length (),
6517 hsa_cfun
->init_extra_bbs ();
6521 HSA_SORRY_AT (UNKNOWN_LOCATION
,
6522 "support for HSA does not implement transactional memory");
6526 verify_function_arguments (cfun
->decl
);
6527 if (hsa_seen_error ())
6530 hsa_cfun
->m_name
= get_brig_function_name (cfun
->decl
);
6532 gen_function_def_parameters ();
6533 if (hsa_seen_error ())
6538 gen_body_from_gimple ();
6539 if (hsa_seen_error ())
6542 if (hsa_cfun
->m_kernel_dispatch_count
)
6543 init_hsa_num_threads ();
6545 if (hsa_cfun
->m_kern_p
)
6547 hsa_function_summary
*s
6548 = hsa_summaries
->get (cgraph_node::get (hsa_cfun
->m_decl
));
6549 hsa_add_kern_decl_mapping (current_function_decl
, hsa_cfun
->m_name
,
6550 hsa_cfun
->m_maximum_omp_data_size
,
6551 s
->m_gridified_kernel_p
);
6556 for (unsigned i
= 0; i
< hsa_cfun
->m_ssa_map
.length (); i
++)
6557 if (hsa_cfun
->m_ssa_map
[i
])
6558 hsa_cfun
->m_ssa_map
[i
]->verify_ssa ();
6561 FOR_EACH_BB_FN (bb
, cfun
)
6563 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
6565 for (hsa_insn_basic
*insn
= hbb
->m_first_insn
; insn
;
6566 insn
= insn
->m_next
)
6572 hsa_brig_emit_function ();
6575 hsa_deinit_data_for_cfun ();
6580 const pass_data pass_data_gen_hsail
=
6583 "hsagen", /* name */
6584 OPTGROUP_OMP
, /* optinfo_flags */
6585 TV_NONE
, /* tv_id */
6586 PROP_cfg
| PROP_ssa
, /* properties_required */
6587 0, /* properties_provided */
6588 0, /* properties_destroyed */
6589 0, /* todo_flags_start */
6590 0 /* todo_flags_finish */
6593 class pass_gen_hsail
: public gimple_opt_pass
6596 pass_gen_hsail (gcc::context
*ctxt
)
6597 : gimple_opt_pass(pass_data_gen_hsail
, ctxt
)
6600 /* opt_pass methods: */
6601 bool gate (function
*);
6602 unsigned int execute (function
*);
6604 }; // class pass_gen_hsail
6606 /* Determine whether or not to run generation of HSAIL. */
6609 pass_gen_hsail::gate (function
*f
)
6611 return hsa_gen_requested_p ()
6612 && hsa_gpu_implementation_p (f
->decl
);
6616 pass_gen_hsail::execute (function
*)
6618 hsa_function_summary
*s
6619 = hsa_summaries
->get (cgraph_node::get_create (current_function_decl
));
6622 generate_hsa (s
->m_kind
== HSA_KERNEL
);
6623 TREE_ASM_WRITTEN (current_function_decl
) = 1;
6624 return TODO_discard_function
;
6629 /* Create the instance of hsa gen pass. */
6632 make_pass_gen_hsail (gcc::context
*ctxt
)
6634 return new pass_gen_hsail (ctxt
);