1 /* A pass for lowering gimple to HSAIL
2 Copyright (C) 2013-2018 Free Software Foundation, Inc.
3 Contributed by Martin Jambor <mjambor@suse.cz> and
4 Martin Liska <mliska@suse.cz>.
6 This file is part of GCC.
8 GCC is free software; you can redistribute it and/or modify
9 it under the terms of the GNU General Public License as published by
10 the Free Software Foundation; either version 3, or (at your option)
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 use B128, while for loading the rest we have to use the
695 input type (??? or maybe also flattened to a equally sized non-vector
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
))
937 segment
= BRIG_SEGMENT_GLOBAL
;
938 allocation
= BRIG_ALLOCATION_PROGRAM
;
940 else if (lookup_attribute ("hsa_global_segment",
941 DECL_ATTRIBUTES (decl
)))
942 segment
= BRIG_SEGMENT_GLOBAL
;
944 segment
= BRIG_SEGMENT_PRIVATE
;
946 sym
= new hsa_symbol (BRIG_TYPE_NONE
, segment
, BRIG_LINKAGE_FUNCTION
,
947 false, allocation
, align
);
948 sym
->fillup_for_decl (decl
);
949 hsa_cfun
->m_private_variables
.safe_push (sym
);
952 sym
->m_name
= hsa_get_declaration_name (decl
);
958 /* For a given HSA function declaration, return a host
959 function declaration. */
962 hsa_get_host_function (tree decl
)
964 hsa_function_summary
*s
965 = hsa_summaries
->get (cgraph_node::get_create (decl
));
966 gcc_assert (s
->m_kind
!= HSA_NONE
);
967 gcc_assert (s
->m_gpu_implementation_p
);
969 return s
->m_bound_function
? s
->m_bound_function
->decl
: NULL
;
972 /* Return true if function DECL has a host equivalent function. */
975 get_brig_function_name (tree decl
)
979 hsa_function_summary
*s
= hsa_summaries
->get (cgraph_node::get_create (d
));
980 if (s
->m_kind
!= HSA_NONE
981 && s
->m_gpu_implementation_p
982 && s
->m_bound_function
)
983 d
= s
->m_bound_function
->decl
;
985 /* IPA split can create a function that has no host equivalent. */
989 char *name
= xstrdup (hsa_get_declaration_name (d
));
990 hsa_sanitize_name (name
);
995 /* Create a spill symbol of type TYPE. */
998 hsa_get_spill_symbol (BrigType16_t type
)
1000 hsa_symbol
*sym
= new hsa_symbol (type
, BRIG_SEGMENT_SPILL
,
1001 BRIG_LINKAGE_FUNCTION
);
1002 hsa_cfun
->m_spill_symbols
.safe_push (sym
);
1006 /* Create a symbol for a read-only string constant. */
1008 hsa_get_string_cst_symbol (tree string_cst
)
1010 gcc_checking_assert (TREE_CODE (string_cst
) == STRING_CST
);
1012 hsa_symbol
**slot
= hsa_cfun
->m_string_constants_map
.get (string_cst
);
1016 hsa_op_immed
*cst
= new hsa_op_immed (string_cst
);
1017 hsa_symbol
*sym
= new hsa_symbol (cst
->m_type
, BRIG_SEGMENT_GLOBAL
,
1018 BRIG_LINKAGE_MODULE
, true,
1019 BRIG_ALLOCATION_AGENT
);
1020 sym
->m_cst_value
= cst
;
1021 sym
->m_dim
= TREE_STRING_LENGTH (string_cst
);
1022 sym
->m_name_number
= hsa_cfun
->m_global_symbols
.length ();
1024 hsa_cfun
->m_global_symbols
.safe_push (sym
);
1025 hsa_cfun
->m_string_constants_map
.put (string_cst
, sym
);
1029 /* Make the type of a MOV instruction larger if mandated by HSAIL rules. */
1032 hsa_fixup_mov_insn_type (hsa_insn_basic
*insn
)
1034 insn
->m_type
= hsa_extend_inttype_to_32bit (insn
->m_type
);
1035 if (insn
->m_type
== BRIG_TYPE_B8
|| insn
->m_type
== BRIG_TYPE_B16
)
1036 insn
->m_type
= BRIG_TYPE_B32
;
1039 /* Constructor of the ancestor of all operands. K is BRIG kind that identified
1040 what the operator is. */
1042 hsa_op_base::hsa_op_base (BrigKind16_t k
)
1043 : m_next (NULL
), m_brig_op_offset (0), m_kind (k
)
1045 hsa_operands
.safe_push (this);
1048 /* Constructor of ancestor of all operands which have a type. K is BRIG kind
1049 that identified what the operator is. T is the type of the operator. */
1051 hsa_op_with_type::hsa_op_with_type (BrigKind16_t k
, BrigType16_t t
)
1052 : hsa_op_base (k
), m_type (t
)
1057 hsa_op_with_type::get_in_type (BrigType16_t dtype
, hsa_bb
*hbb
)
1059 if (m_type
== dtype
)
1064 if (hsa_needs_cvt (dtype
, m_type
))
1066 dest
= new hsa_op_reg (dtype
);
1067 hbb
->append_insn (new hsa_insn_cvt (dest
, this));
1069 else if (is_a
<hsa_op_reg
*> (this))
1071 /* In the end, HSA registers do not really have types, only sizes, so if
1072 the sizes match, we can use the register directly. */
1073 gcc_checking_assert (hsa_type_bit_size (dtype
)
1074 == hsa_type_bit_size (m_type
));
1079 dest
= new hsa_op_reg (m_type
);
1081 hsa_insn_basic
*mov
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
1082 dest
->m_type
, dest
, this);
1083 hsa_fixup_mov_insn_type (mov
);
1084 hbb
->append_insn (mov
);
1085 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
1086 type of the operand must be same as type of the instruction. */
1087 dest
->m_type
= dtype
;
1093 /* If this operand has integer type smaller than 32 bits, extend it to 32 bits,
1094 adding instructions to HBB if needed. */
1097 hsa_op_with_type::extend_int_to_32bit (hsa_bb
*hbb
)
1099 if (m_type
== BRIG_TYPE_U8
|| m_type
== BRIG_TYPE_U16
)
1100 return get_in_type (BRIG_TYPE_U32
, hbb
);
1101 else if (m_type
== BRIG_TYPE_S8
|| m_type
== BRIG_TYPE_S16
)
1102 return get_in_type (BRIG_TYPE_S32
, hbb
);
1107 /* Constructor of class representing HSA immediate values. TREE_VAL is the
1108 tree representation of the immediate value. If min32int is true,
1109 always expand integer types to one that has at least 32 bits. */
1111 hsa_op_immed::hsa_op_immed (tree tree_val
, bool min32int
)
1112 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
,
1113 hsa_type_for_tree_type (TREE_TYPE (tree_val
), NULL
,
1116 if (hsa_seen_error ())
1119 gcc_checking_assert ((is_gimple_min_invariant (tree_val
)
1120 && (!POINTER_TYPE_P (TREE_TYPE (tree_val
))
1121 || TREE_CODE (tree_val
) == INTEGER_CST
))
1122 || TREE_CODE (tree_val
) == CONSTRUCTOR
);
1123 m_tree_value
= tree_val
;
1125 /* Verify that all elements of a constructor are constants. */
1126 if (TREE_CODE (m_tree_value
) == CONSTRUCTOR
)
1127 for (unsigned i
= 0; i
< CONSTRUCTOR_NELTS (m_tree_value
); i
++)
1129 tree v
= CONSTRUCTOR_ELT (m_tree_value
, i
)->value
;
1130 if (!CONSTANT_CLASS_P (v
))
1132 HSA_SORRY_AT (EXPR_LOCATION (tree_val
),
1133 "HSA ctor should have only constants");
1139 /* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
1140 integer representation of the immediate value. TYPE is BRIG type. */
1142 hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value
, BrigType16_t type
)
1143 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
, type
),
1146 gcc_assert (hsa_type_integer_p (type
));
1147 m_int_value
= integer_value
;
1150 hsa_op_immed::hsa_op_immed ()
1151 : hsa_op_with_type (BRIG_KIND_NONE
, BRIG_TYPE_NONE
)
1155 /* New operator to allocate immediate operands from obstack. */
1158 hsa_op_immed::operator new (size_t size
)
1160 return obstack_alloc (&hsa_obstack
, size
);
1165 hsa_op_immed::~hsa_op_immed ()
1169 /* Change type of the immediate value to T. */
1172 hsa_op_immed::set_type (BrigType16_t t
)
1177 /* Constructor of class representing HSA registers and pseudo-registers. T is
1178 the BRIG type of the new register. */
1180 hsa_op_reg::hsa_op_reg (BrigType16_t t
)
1181 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER
, t
), m_gimple_ssa (NULL_TREE
),
1182 m_def_insn (NULL
), m_spill_sym (NULL
), m_order (hsa_cfun
->m_reg_count
++),
1183 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
1187 /* New operator to allocate a register from obstack. */
1190 hsa_op_reg::operator new (size_t size
)
1192 return obstack_alloc (&hsa_obstack
, size
);
1195 /* Verify register operand. */
1198 hsa_op_reg::verify_ssa ()
1200 /* Verify that each HSA register has a definition assigned.
1201 Exceptions are VAR_DECL and PARM_DECL that are a default
1203 gcc_checking_assert (m_def_insn
1204 || (m_gimple_ssa
!= NULL
1205 && (!SSA_NAME_VAR (m_gimple_ssa
)
1206 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa
))
1208 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa
)));
1210 /* Verify that every use of the register is really present
1211 in an instruction. */
1212 for (unsigned i
= 0; i
< m_uses
.length (); i
++)
1214 hsa_insn_basic
*use
= m_uses
[i
];
1216 bool is_visited
= false;
1217 for (unsigned j
= 0; j
< use
->operand_count (); j
++)
1219 hsa_op_base
*u
= use
->get_op (j
);
1220 hsa_op_address
*addr
; addr
= dyn_cast
<hsa_op_address
*> (u
);
1221 if (addr
&& addr
->m_reg
)
1226 bool r
= !addr
&& use
->op_output_p (j
);
1230 error ("HSA SSA name defined by instruction that is supposed "
1232 debug_hsa_operand (this);
1233 debug_hsa_insn (use
);
1234 internal_error ("HSA SSA verification failed");
1243 error ("HSA SSA name not among operands of instruction that is "
1244 "supposed to use it");
1245 debug_hsa_operand (this);
1246 debug_hsa_insn (use
);
1247 internal_error ("HSA SSA verification failed");
1252 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, hsa_op_reg
*r
,
1253 HOST_WIDE_INT offset
)
1254 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (r
),
1255 m_imm_offset (offset
)
1259 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, HOST_WIDE_INT offset
)
1260 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (NULL
),
1261 m_imm_offset (offset
)
1265 hsa_op_address::hsa_op_address (hsa_op_reg
*r
, HOST_WIDE_INT offset
)
1266 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (NULL
), m_reg (r
),
1267 m_imm_offset (offset
)
1271 /* New operator to allocate address operands from obstack. */
1274 hsa_op_address::operator new (size_t size
)
1276 return obstack_alloc (&hsa_obstack
, size
);
1279 /* Constructor of an operand referring to HSAIL code. */
1281 hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF
),
1282 m_directive_offset (0)
1286 /* Constructor of an operand representing a code list. Set it up so that it
1287 can contain ELEMENTS number of elements. */
1289 hsa_op_code_list::hsa_op_code_list (unsigned elements
)
1290 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST
)
1292 m_offsets
.create (1);
1293 m_offsets
.safe_grow_cleared (elements
);
1296 /* New operator to allocate code list operands from obstack. */
1299 hsa_op_code_list::operator new (size_t size
)
1301 return obstack_alloc (&hsa_obstack
, size
);
1304 /* Constructor of an operand representing an operand list.
1305 Set it up so that it can contain ELEMENTS number of elements. */
1307 hsa_op_operand_list::hsa_op_operand_list (unsigned elements
)
1308 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST
)
1310 m_offsets
.create (elements
);
1311 m_offsets
.safe_grow (elements
);
1314 /* New operator to allocate operand list operands from obstack. */
1317 hsa_op_operand_list::operator new (size_t size
)
1319 return obstack_alloc (&hsa_obstack
, size
);
1322 hsa_op_operand_list::~hsa_op_operand_list ()
1324 m_offsets
.release ();
1329 hsa_function_representation::reg_for_gimple_ssa (tree ssa
)
1333 gcc_checking_assert (TREE_CODE (ssa
) == SSA_NAME
);
1334 if (m_ssa_map
[SSA_NAME_VERSION (ssa
)])
1335 return m_ssa_map
[SSA_NAME_VERSION (ssa
)];
1337 hreg
= new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa
),
1339 hreg
->m_gimple_ssa
= ssa
;
1340 m_ssa_map
[SSA_NAME_VERSION (ssa
)] = hreg
;
1346 hsa_op_reg::set_definition (hsa_insn_basic
*insn
)
1348 if (hsa_cfun
->m_in_ssa
)
1350 gcc_checking_assert (!m_def_insn
);
1357 /* Constructor of the class which is the bases of all instructions and directly
1358 represents the most basic ones. NOPS is the number of operands that the
1359 operand vector will contain (and which will be cleared). OP is the opcode
1360 of the instruction. This constructor does not set type. */
1362 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
)
1364 m_next (NULL
), m_bb (NULL
), m_opcode (opc
), m_number (0),
1365 m_type (BRIG_TYPE_NONE
), m_brig_offset (0)
1368 m_operands
.safe_grow_cleared (nops
);
1370 hsa_instructions
.safe_push (this);
1373 /* Make OP the operand number INDEX of operands of this instruction. If OP is a
1374 register or an address containing a register, then either set the definition
1375 of the register to this instruction if it an output operand or add this
1376 instruction to the uses if it is an input one. */
1379 hsa_insn_basic::set_op (int index
, hsa_op_base
*op
)
1381 /* Each address operand is always use. */
1382 hsa_op_address
*addr
= dyn_cast
<hsa_op_address
*> (op
);
1383 if (addr
&& addr
->m_reg
)
1384 addr
->m_reg
->m_uses
.safe_push (this);
1387 hsa_op_reg
*reg
= dyn_cast
<hsa_op_reg
*> (op
);
1390 if (op_output_p (index
))
1391 reg
->set_definition (this);
1393 reg
->m_uses
.safe_push (this);
1397 m_operands
[index
] = op
;
1400 /* Get INDEX-th operand of the instruction. */
1403 hsa_insn_basic::get_op (int index
)
1405 return m_operands
[index
];
1408 /* Get address of INDEX-th operand of the instruction. */
1411 hsa_insn_basic::get_op_addr (int index
)
1413 return &m_operands
[index
];
1416 /* Get number of operands of the instruction. */
1418 hsa_insn_basic::operand_count ()
1420 return m_operands
.length ();
1423 /* Constructor of the class which is the bases of all instructions and directly
1424 represents the most basic ones. NOPS is the number of operands that the
1425 operand vector will contain (and which will be cleared). OPC is the opcode
1426 of the instruction, T is the type of the instruction. */
1428 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
, BrigType16_t t
,
1429 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1430 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1431 : m_prev (NULL
), m_next (NULL
), m_bb (NULL
), m_opcode (opc
),m_number (0),
1432 m_type (t
), m_brig_offset (0)
1435 m_operands
.safe_grow_cleared (nops
);
1439 gcc_checking_assert (nops
>= 1);
1445 gcc_checking_assert (nops
>= 2);
1451 gcc_checking_assert (nops
>= 3);
1457 gcc_checking_assert (nops
>= 4);
1461 hsa_instructions
.safe_push (this);
1464 /* New operator to allocate basic instruction from obstack. */
1467 hsa_insn_basic::operator new (size_t size
)
1469 return obstack_alloc (&hsa_obstack
, size
);
1472 /* Verify the instruction. */
1475 hsa_insn_basic::verify ()
1477 hsa_op_address
*addr
;
1480 /* Iterate all register operands and verify that the instruction
1481 is set in uses of the register. */
1482 for (unsigned i
= 0; i
< operand_count (); i
++)
1484 hsa_op_base
*use
= get_op (i
);
1486 if ((addr
= dyn_cast
<hsa_op_address
*> (use
)) && addr
->m_reg
)
1488 gcc_assert (addr
->m_reg
->m_def_insn
!= this);
1492 if ((reg
= dyn_cast
<hsa_op_reg
*> (use
)) && !op_output_p (i
))
1495 for (j
= 0; j
< reg
->m_uses
.length (); j
++)
1497 if (reg
->m_uses
[j
] == this)
1501 if (j
== reg
->m_uses
.length ())
1503 error ("HSA instruction uses a register but is not among "
1504 "recorded register uses");
1505 debug_hsa_operand (reg
);
1506 debug_hsa_insn (this);
1507 internal_error ("HSA instruction verification failed");
1513 /* Constructor of an instruction representing a PHI node. NOPS is the number
1514 of operands (equal to the number of predecessors). */
1516 hsa_insn_phi::hsa_insn_phi (unsigned nops
, hsa_op_reg
*dst
)
1517 : hsa_insn_basic (nops
, HSA_OPCODE_PHI
), m_dest (dst
)
1519 dst
->set_definition (this);
1522 /* Constructor of class representing instructions for control flow and
1525 hsa_insn_br::hsa_insn_br (unsigned nops
, int opc
, BrigType16_t t
,
1526 BrigWidth8_t width
, hsa_op_base
*arg0
,
1527 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1529 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1534 /* Constructor of class representing instruction for conditional jump, CTRL is
1535 the control register determining whether the jump will be carried out, the
1536 new instruction is automatically added to its uses list. */
1538 hsa_insn_cbr::hsa_insn_cbr (hsa_op_reg
*ctrl
)
1539 : hsa_insn_br (1, BRIG_OPCODE_CBR
, BRIG_TYPE_B1
, BRIG_WIDTH_1
, ctrl
)
1543 /* Constructor of class representing instruction for switch jump, CTRL is
1544 the index register. */
1546 hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg
*index
, unsigned jump_count
)
1547 : hsa_insn_basic (1, BRIG_OPCODE_SBR
, BRIG_TYPE_B1
, index
),
1548 m_width (BRIG_WIDTH_1
), m_jump_table (vNULL
),
1549 m_label_code_list (new hsa_op_code_list (jump_count
))
1553 /* Replace all occurrences of OLD_BB with NEW_BB in the statements
1557 hsa_insn_sbr::replace_all_labels (basic_block old_bb
, basic_block new_bb
)
1559 for (unsigned i
= 0; i
< m_jump_table
.length (); i
++)
1560 if (m_jump_table
[i
] == old_bb
)
1561 m_jump_table
[i
] = new_bb
;
1564 hsa_insn_sbr::~hsa_insn_sbr ()
1566 m_jump_table
.release ();
1569 /* Constructor of comparison instruction. CMP is the comparison operation and T
1570 is the result type. */
1572 hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp
, BrigType16_t t
,
1573 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1575 : hsa_insn_basic (3 , BRIG_OPCODE_CMP
, t
, arg0
, arg1
, arg2
), m_compare (cmp
)
1579 /* Constructor of classes representing memory accesses. OPC is the opcode (must
1580 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
1581 operands are provided as ARG0 and ARG1. */
1583 hsa_insn_mem::hsa_insn_mem (int opc
, BrigType16_t t
, hsa_op_base
*arg0
,
1585 : hsa_insn_basic (2, opc
, t
, arg0
, arg1
),
1586 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1588 gcc_checking_assert (opc
== BRIG_OPCODE_LD
|| opc
== BRIG_OPCODE_ST
);
1591 /* Constructor for descendants allowing different opcodes and number of
1592 operands, it passes its arguments directly to hsa_insn_basic
1593 constructor. The instruction operands are provided as ARG[0-3]. */
1596 hsa_insn_mem::hsa_insn_mem (unsigned nops
, int opc
, BrigType16_t t
,
1597 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1598 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1599 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1600 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1604 /* Constructor of class representing atomic instructions. OPC is the principal
1605 opcode, AOP is the specific atomic operation opcode. T is the type of the
1606 instruction. The instruction operands are provided as ARG[0-3]. */
1608 hsa_insn_atomic::hsa_insn_atomic (int nops
, int opc
,
1609 enum BrigAtomicOperation aop
,
1610 BrigType16_t t
, BrigMemoryOrder memorder
,
1612 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1614 : hsa_insn_mem (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
), m_atomicop (aop
),
1615 m_memoryorder (memorder
),
1616 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM
)
1618 gcc_checking_assert (opc
== BRIG_OPCODE_ATOMICNORET
||
1619 opc
== BRIG_OPCODE_ATOMIC
||
1620 opc
== BRIG_OPCODE_SIGNAL
||
1621 opc
== BRIG_OPCODE_SIGNALNORET
);
1624 /* Constructor of class representing signal instructions. OPC is the prinicpal
1625 opcode, SOP is the specific signal operation opcode. T is the type of the
1626 instruction. The instruction operands are provided as ARG[0-3]. */
1628 hsa_insn_signal::hsa_insn_signal (int nops
, int opc
,
1629 enum BrigAtomicOperation sop
,
1630 BrigType16_t t
, BrigMemoryOrder memorder
,
1631 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1632 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1633 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1634 m_memory_order (memorder
), m_signalop (sop
)
1638 /* Constructor of class representing segment conversion instructions. OPC is
1639 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
1640 and SRCT are destination and source types respectively, SEG is the segment
1641 we are converting to or from. The instruction operands are
1642 provided as ARG0 and ARG1. */
1644 hsa_insn_seg::hsa_insn_seg (int opc
, BrigType16_t dest
, BrigType16_t srct
,
1645 BrigSegment8_t seg
, hsa_op_base
*arg0
,
1647 : hsa_insn_basic (2, opc
, dest
, arg0
, arg1
), m_src_type (srct
),
1650 gcc_checking_assert (opc
== BRIG_OPCODE_STOF
|| opc
== BRIG_OPCODE_FTOS
);
1653 /* Constructor of class representing a call instruction. CALLEE is the tree
1654 representation of the function being called. */
1656 hsa_insn_call::hsa_insn_call (tree callee
)
1657 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (callee
),
1658 m_output_arg (NULL
), m_args_code_list (NULL
), m_result_code_list (NULL
)
1662 hsa_insn_call::hsa_insn_call (hsa_internal_fn
*fn
)
1663 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (NULL
),
1664 m_called_internal_fn (fn
), m_output_arg (NULL
), m_args_code_list (NULL
),
1665 m_result_code_list (NULL
)
1669 hsa_insn_call::~hsa_insn_call ()
1671 for (unsigned i
= 0; i
< m_input_args
.length (); i
++)
1672 delete m_input_args
[i
];
1674 delete m_output_arg
;
1676 m_input_args
.release ();
1677 m_input_arg_insns
.release ();
1680 /* Constructor of class representing the argument block required to invoke
1682 hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind
,
1683 hsa_insn_call
* call
)
1684 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK
), m_kind (brig_kind
),
1689 hsa_insn_comment::hsa_insn_comment (const char *s
)
1690 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT
)
1692 unsigned l
= strlen (s
);
1694 /* Append '// ' to the string. */
1695 char *buf
= XNEWVEC (char, l
+ 4);
1696 sprintf (buf
, "// %s", s
);
1700 hsa_insn_comment::~hsa_insn_comment ()
1702 gcc_checking_assert (m_comment
);
1707 /* Constructor of class representing the queue instruction in HSAIL. */
1709 hsa_insn_queue::hsa_insn_queue (int nops
, int opcode
, BrigSegment segment
,
1710 BrigMemoryOrder memory_order
,
1711 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1712 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1713 : hsa_insn_basic (nops
, opcode
, BRIG_TYPE_U64
, arg0
, arg1
, arg2
, arg3
),
1714 m_segment (segment
), m_memory_order (memory_order
)
1718 /* Constructor of class representing the source type instruction in HSAIL. */
1720 hsa_insn_srctype::hsa_insn_srctype (int nops
, BrigOpcode opcode
,
1721 BrigType16_t destt
, BrigType16_t srct
,
1722 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1723 hsa_op_base
*arg2
= NULL
)
1724 : hsa_insn_basic (nops
, opcode
, destt
, arg0
, arg1
, arg2
),
1725 m_source_type (srct
)
1728 /* Constructor of class representing the packed instruction in HSAIL. */
1730 hsa_insn_packed::hsa_insn_packed (int nops
, BrigOpcode opcode
,
1731 BrigType16_t destt
, BrigType16_t srct
,
1732 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1734 : hsa_insn_srctype (nops
, opcode
, destt
, srct
, arg0
, arg1
, arg2
)
1736 m_operand_list
= new hsa_op_operand_list (nops
- 1);
1739 /* Constructor of class representing the convert instruction in HSAIL. */
1741 hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type
*dest
, hsa_op_with_type
*src
)
1742 : hsa_insn_basic (2, BRIG_OPCODE_CVT
, dest
->m_type
, dest
, src
)
1746 /* Constructor of class representing the alloca in HSAIL. */
1748 hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type
*dest
,
1749 hsa_op_with_type
*size
, unsigned alignment
)
1750 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA
, dest
->m_type
, dest
, size
),
1751 m_align (BRIG_ALIGNMENT_8
)
1753 gcc_assert (dest
->m_type
== BRIG_TYPE_U32
);
1755 m_align
= hsa_alignment_encoding (alignment
);
1758 /* Append an instruction INSN into the basic block. */
1761 hsa_bb::append_insn (hsa_insn_basic
*insn
)
1763 gcc_assert (insn
->m_opcode
!= 0 || insn
->operand_count () == 0);
1764 gcc_assert (!insn
->m_bb
);
1767 insn
->m_prev
= m_last_insn
;
1768 insn
->m_next
= NULL
;
1770 m_last_insn
->m_next
= insn
;
1773 m_first_insn
= insn
;
1777 hsa_bb::append_phi (hsa_insn_phi
*hphi
)
1781 hphi
->m_prev
= m_last_phi
;
1782 hphi
->m_next
= NULL
;
1784 m_last_phi
->m_next
= hphi
;
1790 /* Insert HSA instruction NEW_INSN immediately before an existing instruction
1794 hsa_insert_insn_before (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1796 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1798 if (hbb
->m_first_insn
== old_insn
)
1799 hbb
->m_first_insn
= new_insn
;
1800 new_insn
->m_prev
= old_insn
->m_prev
;
1801 new_insn
->m_next
= old_insn
;
1802 if (old_insn
->m_prev
)
1803 old_insn
->m_prev
->m_next
= new_insn
;
1804 old_insn
->m_prev
= new_insn
;
1807 /* Append HSA instruction NEW_INSN immediately after an existing instruction
1811 hsa_append_insn_after (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1813 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1815 if (hbb
->m_last_insn
== old_insn
)
1816 hbb
->m_last_insn
= new_insn
;
1817 new_insn
->m_prev
= old_insn
;
1818 new_insn
->m_next
= old_insn
->m_next
;
1819 if (old_insn
->m_next
)
1820 old_insn
->m_next
->m_prev
= new_insn
;
1821 old_insn
->m_next
= new_insn
;
1824 /* Return a register containing the calculated value of EXP which must be an
1825 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
1826 integer constants as returned by get_inner_reference.
1827 Newly generated HSA instructions will be appended to HBB.
1828 Perform all calculations in ADDRTYPE. */
1830 static hsa_op_with_type
*
1831 gen_address_calculation (tree exp
, hsa_bb
*hbb
, BrigType16_t addrtype
)
1835 if (TREE_CODE (exp
) == NOP_EXPR
)
1836 exp
= TREE_OPERAND (exp
, 0);
1838 switch (TREE_CODE (exp
))
1841 return hsa_cfun
->reg_for_gimple_ssa (exp
)->get_in_type (addrtype
, hbb
);
1845 hsa_op_immed
*imm
= new hsa_op_immed (exp
);
1846 if (addrtype
!= imm
->m_type
)
1847 imm
->m_type
= addrtype
;
1852 opcode
= BRIG_OPCODE_ADD
;
1856 opcode
= BRIG_OPCODE_MUL
;
1863 hsa_op_reg
*res
= new hsa_op_reg (addrtype
);
1864 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, addrtype
);
1865 insn
->set_op (0, res
);
1867 hsa_op_with_type
*op1
= gen_address_calculation (TREE_OPERAND (exp
, 0), hbb
,
1869 hsa_op_with_type
*op2
= gen_address_calculation (TREE_OPERAND (exp
, 1), hbb
,
1871 insn
->set_op (1, op1
);
1872 insn
->set_op (2, op2
);
1874 hbb
->append_insn (insn
);
1878 /* If R1 is NULL, just return R2, otherwise append an instruction adding them
1879 to HBB and return the register holding the result. */
1882 add_addr_regs_if_needed (hsa_op_reg
*r1
, hsa_op_reg
*r2
, hsa_bb
*hbb
)
1884 gcc_checking_assert (r2
);
1888 hsa_op_reg
*res
= new hsa_op_reg (r1
->m_type
);
1889 gcc_assert (!hsa_needs_cvt (r1
->m_type
, r2
->m_type
));
1890 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_ADD
, res
->m_type
);
1891 insn
->set_op (0, res
);
1892 insn
->set_op (1, r1
);
1893 insn
->set_op (2, r2
);
1894 hbb
->append_insn (insn
);
1898 /* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
1899 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
1902 process_mem_base (tree base
, hsa_symbol
**symbol
, BrigType16_t
*addrtype
,
1903 hsa_op_reg
**reg
, offset_int
*offset
, hsa_bb
*hbb
)
1905 if (TREE_CODE (base
) == SSA_NAME
)
1908 hsa_op_with_type
*ssa
1909 = hsa_cfun
->reg_for_gimple_ssa (base
)->get_in_type (*addrtype
, hbb
);
1910 *reg
= dyn_cast
<hsa_op_reg
*> (ssa
);
1912 else if (TREE_CODE (base
) == ADDR_EXPR
)
1914 tree decl
= TREE_OPERAND (base
, 0);
1916 if (!DECL_P (decl
) || TREE_CODE (decl
) == FUNCTION_DECL
)
1918 HSA_SORRY_AT (EXPR_LOCATION (base
),
1919 "support for HSA does not implement a memory reference "
1920 "to a non-declaration type");
1924 gcc_assert (!*symbol
);
1926 *symbol
= get_symbol_for_decl (decl
);
1927 *addrtype
= hsa_get_segment_addr_type ((*symbol
)->m_segment
);
1929 else if (TREE_CODE (base
) == INTEGER_CST
)
1930 *offset
+= wi::to_offset (base
);
1935 /* Forward declaration of a function. */
1938 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
);
1940 /* Generate HSA address operand for a given tree memory reference REF. If
1941 instructions need to be created to calculate the address, they will be added
1942 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
1943 the function assumes that the caller will handle possible
1944 bit-field references. Otherwise if we reference a bit-field, sorry message
1947 static hsa_op_address
*
1948 gen_hsa_addr (tree ref
, hsa_bb
*hbb
, HOST_WIDE_INT
*output_bitsize
= NULL
,
1949 HOST_WIDE_INT
*output_bitpos
= NULL
)
1951 hsa_symbol
*symbol
= NULL
;
1952 hsa_op_reg
*reg
= NULL
;
1953 offset_int offset
= 0;
1955 tree varoffset
= NULL_TREE
;
1956 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
1957 HOST_WIDE_INT bitsize
= 0, bitpos
= 0;
1958 BrigType16_t flat_addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
1960 if (TREE_CODE (ref
) == STRING_CST
)
1962 symbol
= hsa_get_string_cst_symbol (ref
);
1965 else if (TREE_CODE (ref
) == BIT_FIELD_REF
1966 && (!multiple_p (bit_field_size (ref
), BITS_PER_UNIT
)
1967 || !multiple_p (bit_field_offset (ref
), BITS_PER_UNIT
)))
1969 HSA_SORRY_ATV (EXPR_LOCATION (origref
),
1970 "support for HSA does not implement "
1971 "bit field references such as %E", ref
);
1975 if (handled_component_p (ref
))
1978 int unsignedp
, volatilep
, preversep
;
1979 poly_int64 pbitsize
, pbitpos
;
1982 new_ref
= get_inner_reference (ref
, &pbitsize
, &pbitpos
, &varoffset
,
1983 &mode
, &unsignedp
, &preversep
,
1985 /* When this isn't true, the switch below will report an
1986 appropriate error. */
1987 if (pbitsize
.is_constant () && pbitpos
.is_constant ())
1989 bitsize
= pbitsize
.to_constant ();
1990 bitpos
= pbitpos
.to_constant ();
1993 offset
= wi::rshift (offset
, LOG2_BITS_PER_UNIT
, SIGNED
);
1997 switch (TREE_CODE (ref
))
2001 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
2002 symbol
= hsa_cfun
->create_hsa_temporary (flat_addrtype
);
2003 hsa_op_reg
*r
= new hsa_op_reg (flat_addrtype
);
2004 gen_hsa_addr_insns (ref
, r
, hbb
);
2005 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
2006 r
, new hsa_op_address (symbol
)));
2012 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
2013 hsa_op_with_type
*r
= hsa_cfun
->reg_for_gimple_ssa (ref
);
2014 if (r
->m_type
== BRIG_TYPE_B1
)
2015 r
= r
->get_in_type (BRIG_TYPE_U32
, hbb
);
2016 symbol
= hsa_cfun
->create_hsa_temporary (r
->m_type
);
2018 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
2019 r
, new hsa_op_address (symbol
)));
2027 gcc_assert (!symbol
);
2028 symbol
= get_symbol_for_decl (ref
);
2029 addrtype
= hsa_get_segment_addr_type (symbol
->m_segment
);
2033 process_mem_base (TREE_OPERAND (ref
, 0), &symbol
, &addrtype
, ®
,
2036 if (!integer_zerop (TREE_OPERAND (ref
, 1)))
2037 offset
+= wi::to_offset (TREE_OPERAND (ref
, 1));
2040 case TARGET_MEM_REF
:
2041 process_mem_base (TMR_BASE (ref
), &symbol
, &addrtype
, ®
, &offset
, hbb
);
2042 if (TMR_INDEX (ref
))
2045 hsa_op_base
*idx
= hsa_cfun
->reg_for_gimple_ssa
2046 (TMR_INDEX (ref
))->get_in_type (addrtype
, hbb
);
2047 if (TMR_STEP (ref
) && !integer_onep (TMR_STEP (ref
)))
2049 disp1
= new hsa_op_reg (addrtype
);
2050 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_MUL
,
2053 /* As step must respect addrtype, we overwrite the type
2054 of an immediate value. */
2055 hsa_op_immed
*step
= new hsa_op_immed (TMR_STEP (ref
));
2056 step
->m_type
= addrtype
;
2058 insn
->set_op (0, disp1
);
2059 insn
->set_op (1, idx
);
2060 insn
->set_op (2, step
);
2061 hbb
->append_insn (insn
);
2064 disp1
= as_a
<hsa_op_reg
*> (idx
);
2065 reg
= add_addr_regs_if_needed (reg
, disp1
, hbb
);
2067 if (TMR_INDEX2 (ref
))
2069 if (TREE_CODE (TMR_INDEX2 (ref
)) == SSA_NAME
)
2071 hsa_op_base
*disp2
= hsa_cfun
->reg_for_gimple_ssa
2072 (TMR_INDEX2 (ref
))->get_in_type (addrtype
, hbb
);
2073 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (disp2
),
2076 else if (TREE_CODE (TMR_INDEX2 (ref
)) == INTEGER_CST
)
2077 offset
+= wi::to_offset (TMR_INDEX2 (ref
));
2081 offset
+= wi::to_offset (TMR_OFFSET (ref
));
2084 HSA_SORRY_AT (EXPR_LOCATION (origref
),
2085 "support for HSA does not implement function pointers");
2088 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does "
2089 "not implement memory access to %E", origref
);
2095 if (TREE_CODE (varoffset
) == INTEGER_CST
)
2096 offset
+= wi::to_offset (varoffset
);
2099 hsa_op_base
*off_op
= gen_address_calculation (varoffset
, hbb
,
2101 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (off_op
),
2106 gcc_checking_assert ((symbol
2108 == hsa_get_segment_addr_type (symbol
->m_segment
))
2111 == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
)));
2113 HOST_WIDE_INT hwi_offset
= offset
.to_shwi ();
2115 /* Calculate remaining bitsize offset (if presented). */
2116 bitpos
%= BITS_PER_UNIT
;
2117 /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
2118 is not a reason to think this is a bit-field access. */
2120 && (bitsize
>= BITS_PER_UNIT
)
2121 && !(bitsize
& (bitsize
- 1)))
2124 if ((bitpos
|| bitsize
) && (output_bitpos
== NULL
|| output_bitsize
== NULL
))
2125 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does not "
2126 "implement unhandled bit field reference such as %E", ref
);
2128 if (output_bitsize
!= NULL
&& output_bitpos
!= NULL
)
2130 *output_bitsize
= bitsize
;
2131 *output_bitpos
= bitpos
;
2134 return new hsa_op_address (symbol
, reg
, hwi_offset
);
2137 /* Generate HSA address operand for a given tree memory reference REF. If
2138 instructions need to be created to calculate the address, they will be added
2139 to the end of HBB. OUTPUT_ALIGN is alignment of the created address. */
2141 static hsa_op_address
*
2142 gen_hsa_addr_with_align (tree ref
, hsa_bb
*hbb
, BrigAlignment8_t
*output_align
)
2144 hsa_op_address
*addr
= gen_hsa_addr (ref
, hbb
);
2145 if (addr
->m_reg
|| !addr
->m_symbol
)
2146 *output_align
= hsa_object_alignment (ref
);
2149 /* If the address consists only of a symbol and an offset, we
2150 compute the alignment ourselves to take into account any alignment
2151 promotions we might have done for the HSA symbol representation. */
2152 unsigned align
= hsa_byte_alignment (addr
->m_symbol
->m_align
);
2153 unsigned misalign
= addr
->m_imm_offset
& (align
- 1);
2155 align
= least_bit_hwi (misalign
);
2156 *output_align
= hsa_alignment_encoding (BITS_PER_UNIT
* align
);
2161 /* Generate HSA address for a function call argument of given TYPE.
2162 INDEX is used to generate corresponding name of the arguments.
2163 Special value -1 represents fact that result value is created. */
2165 static hsa_op_address
*
2166 gen_hsa_addr_for_arg (tree tree_type
, int index
)
2168 hsa_symbol
*sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
2170 sym
->m_type
= hsa_type_for_tree_type (tree_type
, &sym
->m_dim
);
2172 if (index
== -1) /* Function result. */
2173 sym
->m_name
= "res";
2174 else /* Function call arguments. */
2177 sym
->m_name_number
= index
;
2180 return new hsa_op_address (sym
);
2183 /* Generate HSA instructions that process all necessary conversions
2184 of an ADDR to flat addressing and place the result into DEST.
2185 Instructions are appended to HBB. */
2188 convert_addr_to_flat_segment (hsa_op_address
*addr
, hsa_op_reg
*dest
,
2191 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_LDA
);
2192 insn
->set_op (1, addr
);
2193 if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
!= BRIG_SEGMENT_GLOBAL
)
2195 /* LDA produces segment-relative address, we need to convert
2196 it to the flat one. */
2198 tmp
= new hsa_op_reg (hsa_get_segment_addr_type
2199 (addr
->m_symbol
->m_segment
));
2201 seg
= new hsa_insn_seg (BRIG_OPCODE_STOF
,
2202 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
2203 tmp
->m_type
, addr
->m_symbol
->m_segment
, dest
,
2206 insn
->set_op (0, tmp
);
2207 insn
->m_type
= tmp
->m_type
;
2208 hbb
->append_insn (insn
);
2209 hbb
->append_insn (seg
);
2213 insn
->set_op (0, dest
);
2214 insn
->m_type
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2215 hbb
->append_insn (insn
);
2219 /* Generate HSA instructions that calculate address of VAL including all
2220 necessary conversions to flat addressing and place the result into DEST.
2221 Instructions are appended to HBB. */
2224 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
)
2226 /* Handle cases like tmp = NULL, where we just emit a move instruction
2228 if (TREE_CODE (val
) == INTEGER_CST
)
2230 hsa_op_immed
*c
= new hsa_op_immed (val
);
2231 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2232 dest
->m_type
, dest
, c
);
2233 hbb
->append_insn (insn
);
2237 hsa_op_address
*addr
;
2239 gcc_assert (dest
->m_type
== hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2240 if (TREE_CODE (val
) == ADDR_EXPR
)
2241 val
= TREE_OPERAND (val
, 0);
2242 addr
= gen_hsa_addr (val
, hbb
);
2244 if (TREE_CODE (val
) == CONST_DECL
2245 && is_gimple_reg_type (TREE_TYPE (val
)))
2247 gcc_assert (addr
->m_symbol
2248 && addr
->m_symbol
->m_segment
== BRIG_SEGMENT_READONLY
);
2249 /* CONST_DECLs are in readonly segment which however does not have
2250 addresses convertible to flat segments. So copy it to a private one
2251 and take address of that. */
2252 BrigType16_t csttype
2253 = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (val
),
2255 hsa_op_reg
*r
= new hsa_op_reg (csttype
);
2256 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD
, csttype
, r
,
2257 new hsa_op_address (addr
->m_symbol
)));
2258 hsa_symbol
*copysym
= hsa_cfun
->create_hsa_temporary (csttype
);
2259 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, csttype
, r
,
2260 new hsa_op_address (copysym
)));
2261 addr
->m_symbol
= copysym
;
2263 else if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
== BRIG_SEGMENT_READONLY
)
2265 HSA_SORRY_ATV (EXPR_LOCATION (val
), "support for HSA does "
2266 "not implement taking addresses of complex "
2267 "CONST_DECLs such as %E", val
);
2272 convert_addr_to_flat_segment (addr
, dest
, hbb
);
2275 /* Return an HSA register or HSA immediate value operand corresponding to
2276 gimple operand OP. */
2278 static hsa_op_with_type
*
2279 hsa_reg_or_immed_for_gimple_op (tree op
, hsa_bb
*hbb
)
2283 if (TREE_CODE (op
) == SSA_NAME
)
2284 tmp
= hsa_cfun
->reg_for_gimple_ssa (op
);
2285 else if (!POINTER_TYPE_P (TREE_TYPE (op
)))
2286 return new hsa_op_immed (op
);
2289 tmp
= new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2290 gen_hsa_addr_insns (op
, tmp
, hbb
);
2295 /* Create a simple movement instruction with register destination DEST and
2296 register or immediate source SRC and append it to the end of HBB. */
2299 hsa_build_append_simple_mov (hsa_op_reg
*dest
, hsa_op_base
*src
, hsa_bb
*hbb
)
2301 /* Moves of packed data between registers need to adhere to the same type
2302 rules like when dealing with memory. */
2303 BrigType16_t tp
= mem_type_for_type (dest
->m_type
);
2304 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, tp
, dest
, src
);
2305 hsa_fixup_mov_insn_type (insn
);
2306 unsigned dest_size
= hsa_type_bit_size (dest
->m_type
);
2307 if (hsa_op_reg
*sreg
= dyn_cast
<hsa_op_reg
*> (src
))
2308 gcc_assert (dest_size
== hsa_type_bit_size (sreg
->m_type
));
2312 = hsa_type_bit_size (as_a
<hsa_op_immed
*> (src
)->m_type
);
2313 gcc_assert ((dest_size
== imm_size
)
2314 /* Eventually < 32bit registers will be promoted to 32bit. */
2315 || (dest_size
< 32 && imm_size
== 32));
2317 hbb
->append_insn (insn
);
2320 /* Generate HSAIL instructions loading a bit field into register DEST.
2321 VALUE_REG is a register of a SSA name that is used in the bit field
2322 reference. To identify a bit field BITPOS is offset to the loaded memory
2323 and BITSIZE is number of bits of the bit field.
2324 Add instructions to HBB. */
2327 gen_hsa_insns_for_bitfield (hsa_op_reg
*dest
, hsa_op_reg
*value_reg
,
2328 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2331 unsigned type_bitsize
2332 = hsa_type_bit_size (hsa_extend_inttype_to_32bit (dest
->m_type
));
2333 unsigned left_shift
= type_bitsize
- (bitsize
+ bitpos
);
2334 unsigned right_shift
= left_shift
+ bitpos
;
2338 hsa_op_reg
*value_reg_2
2339 = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest
->m_type
));
2340 hsa_op_immed
*c
= new hsa_op_immed (left_shift
, BRIG_TYPE_U32
);
2342 hsa_insn_basic
*lshift
2343 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, value_reg_2
->m_type
,
2344 value_reg_2
, value_reg
, c
);
2346 hbb
->append_insn (lshift
);
2348 value_reg
= value_reg_2
;
2353 hsa_op_reg
*value_reg_2
2354 = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest
->m_type
));
2355 hsa_op_immed
*c
= new hsa_op_immed (right_shift
, BRIG_TYPE_U32
);
2357 hsa_insn_basic
*rshift
2358 = new hsa_insn_basic (3, BRIG_OPCODE_SHR
, value_reg_2
->m_type
,
2359 value_reg_2
, value_reg
, c
);
2361 hbb
->append_insn (rshift
);
2363 value_reg
= value_reg_2
;
2366 hsa_insn_basic
*assignment
2367 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, NULL
, value_reg
);
2368 hsa_fixup_mov_insn_type (assignment
);
2369 hbb
->append_insn (assignment
);
2370 assignment
->set_output_in_type (dest
, 0, hbb
);
2374 /* Generate HSAIL instructions loading a bit field into register DEST. ADDR is
2375 prepared memory address which is used to load the bit field. To identify a
2376 bit field BITPOS is offset to the loaded memory and BITSIZE is number of
2377 bits of the bit field. Add instructions to HBB. Load must be performed in
2381 gen_hsa_insns_for_bitfield_load (hsa_op_reg
*dest
, hsa_op_address
*addr
,
2382 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2383 hsa_bb
*hbb
, BrigAlignment8_t align
)
2385 hsa_op_reg
*value_reg
= new hsa_op_reg (dest
->m_type
);
2387 = new hsa_insn_mem (BRIG_OPCODE_LD
,
2388 hsa_extend_inttype_to_32bit (dest
->m_type
),
2390 mem
->set_align (align
);
2391 hbb
->append_insn (mem
);
2392 gen_hsa_insns_for_bitfield (dest
, value_reg
, bitsize
, bitpos
, hbb
);
2395 /* Return the alignment of base memory accesses we issue to perform bit-field
2396 memory access REF. */
2398 static BrigAlignment8_t
2399 hsa_bitmemref_alignment (tree ref
)
2401 unsigned HOST_WIDE_INT bit_offset
= 0;
2405 if (TREE_CODE (ref
) == BIT_FIELD_REF
)
2407 if (!tree_fits_uhwi_p (TREE_OPERAND (ref
, 2)))
2408 return BRIG_ALIGNMENT_1
;
2409 bit_offset
+= tree_to_uhwi (TREE_OPERAND (ref
, 2));
2411 else if (TREE_CODE (ref
) == COMPONENT_REF
2412 && DECL_BIT_FIELD (TREE_OPERAND (ref
, 1)))
2413 bit_offset
+= int_bit_position (TREE_OPERAND (ref
, 1));
2416 ref
= TREE_OPERAND (ref
, 0);
2419 unsigned HOST_WIDE_INT bits
= bit_offset
% BITS_PER_UNIT
;
2420 unsigned HOST_WIDE_INT byte_bits
= bit_offset
- bits
;
2421 BrigAlignment8_t base
= hsa_object_alignment (ref
);
2424 return MIN (base
, hsa_alignment_encoding (least_bit_hwi (byte_bits
)));
2427 /* Generate HSAIL instructions loading something into register DEST. RHS is
2428 tree representation of the loaded data, which are loaded as type TYPE. Add
2429 instructions to HBB. */
2432 gen_hsa_insns_for_load (hsa_op_reg
*dest
, tree rhs
, tree type
, hsa_bb
*hbb
)
2434 /* The destination SSA name will give us the type. */
2435 if (TREE_CODE (rhs
) == VIEW_CONVERT_EXPR
)
2436 rhs
= TREE_OPERAND (rhs
, 0);
2438 if (TREE_CODE (rhs
) == SSA_NAME
)
2440 hsa_op_reg
*src
= hsa_cfun
->reg_for_gimple_ssa (rhs
);
2441 hsa_build_append_simple_mov (dest
, src
, hbb
);
2443 else if (is_gimple_min_invariant (rhs
)
2444 || TREE_CODE (rhs
) == ADDR_EXPR
)
2446 if (POINTER_TYPE_P (TREE_TYPE (rhs
)))
2448 if (dest
->m_type
!= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
))
2450 HSA_SORRY_ATV (EXPR_LOCATION (rhs
),
2451 "support for HSA does not implement conversion "
2452 "of %E to the requested non-pointer type.", rhs
);
2456 gen_hsa_addr_insns (rhs
, dest
, hbb
);
2458 else if (TREE_CODE (rhs
) == COMPLEX_CST
)
2460 hsa_op_immed
*real_part
= new hsa_op_immed (TREE_REALPART (rhs
));
2461 hsa_op_immed
*imag_part
= new hsa_op_immed (TREE_IMAGPART (rhs
));
2463 hsa_op_reg
*real_part_reg
2464 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2466 hsa_op_reg
*imag_part_reg
2467 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2470 hsa_build_append_simple_mov (real_part_reg
, real_part
, hbb
);
2471 hsa_build_append_simple_mov (imag_part_reg
, imag_part
, hbb
);
2473 BrigType16_t src_type
= hsa_bittype_for_type (real_part_reg
->m_type
);
2475 hsa_insn_packed
*insn
2476 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
,
2477 src_type
, dest
, real_part_reg
,
2479 hbb
->append_insn (insn
);
2483 hsa_op_immed
*imm
= new hsa_op_immed (rhs
);
2484 hsa_build_append_simple_mov (dest
, imm
, hbb
);
2487 else if (TREE_CODE (rhs
) == REALPART_EXPR
|| TREE_CODE (rhs
) == IMAGPART_EXPR
)
2489 tree pack_type
= TREE_TYPE (TREE_OPERAND (rhs
, 0));
2491 hsa_op_reg
*packed_reg
2492 = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type
, true));
2494 tree complex_rhs
= TREE_OPERAND (rhs
, 0);
2495 gen_hsa_insns_for_load (packed_reg
, complex_rhs
, TREE_TYPE (complex_rhs
),
2498 hsa_op_reg
*real_reg
2499 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2501 hsa_op_reg
*imag_reg
2502 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2504 BrigKind16_t brig_type
= packed_reg
->m_type
;
2505 hsa_insn_packed
*packed
2506 = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND
,
2507 hsa_bittype_for_type (real_reg
->m_type
),
2508 brig_type
, real_reg
, imag_reg
, packed_reg
);
2510 hbb
->append_insn (packed
);
2512 hsa_op_reg
*source
= TREE_CODE (rhs
) == REALPART_EXPR
?
2513 real_reg
: imag_reg
;
2515 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2516 dest
->m_type
, NULL
, source
);
2517 hsa_fixup_mov_insn_type (insn
);
2518 hbb
->append_insn (insn
);
2519 insn
->set_output_in_type (dest
, 0, hbb
);
2521 else if (TREE_CODE (rhs
) == BIT_FIELD_REF
2522 && TREE_CODE (TREE_OPERAND (rhs
, 0)) == SSA_NAME
)
2524 tree ssa_name
= TREE_OPERAND (rhs
, 0);
2525 HOST_WIDE_INT bitsize
= tree_to_uhwi (TREE_OPERAND (rhs
, 1));
2526 HOST_WIDE_INT bitpos
= tree_to_uhwi (TREE_OPERAND (rhs
, 2));
2528 hsa_op_reg
*imm_value
= hsa_cfun
->reg_for_gimple_ssa (ssa_name
);
2529 gen_hsa_insns_for_bitfield (dest
, imm_value
, bitsize
, bitpos
, hbb
);
2531 else if (DECL_P (rhs
) || TREE_CODE (rhs
) == MEM_REF
2532 || TREE_CODE (rhs
) == TARGET_MEM_REF
2533 || handled_component_p (rhs
))
2535 HOST_WIDE_INT bitsize
, bitpos
;
2537 /* Load from memory. */
2538 hsa_op_address
*addr
;
2539 addr
= gen_hsa_addr (rhs
, hbb
, &bitsize
, &bitpos
);
2541 /* Handle load of a bit field. */
2544 HSA_SORRY_AT (EXPR_LOCATION (rhs
),
2545 "support for HSA does not implement load from a bit "
2546 "field bigger than 64 bits");
2550 if (bitsize
|| bitpos
)
2551 gen_hsa_insns_for_bitfield_load (dest
, addr
, bitsize
, bitpos
, hbb
,
2552 hsa_bitmemref_alignment (rhs
));
2556 /* Not dest->m_type, that's possibly extended. */
2557 mtype
= mem_type_for_type (hsa_type_for_scalar_tree_type (type
,
2559 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
, dest
,
2561 mem
->set_align (hsa_object_alignment (rhs
));
2562 hbb
->append_insn (mem
);
2566 HSA_SORRY_ATV (EXPR_LOCATION (rhs
),
2567 "support for HSA does not implement loading "
2572 /* Return number of bits necessary for representation of a bit field,
2573 starting at BITPOS with size of BITSIZE. */
2576 get_bitfield_size (unsigned bitpos
, unsigned bitsize
)
2578 unsigned s
= bitpos
+ bitsize
;
2579 unsigned sizes
[] = {8, 16, 32, 64};
2581 for (unsigned i
= 0; i
< 4; i
++)
2589 /* Generate HSAIL instructions storing into memory. LHS is the destination of
2590 the store, SRC is the source operand. Add instructions to HBB. */
2593 gen_hsa_insns_for_store (tree lhs
, hsa_op_base
*src
, hsa_bb
*hbb
)
2595 HOST_WIDE_INT bitsize
= 0, bitpos
= 0;
2596 BrigAlignment8_t req_align
;
2598 mtype
= mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
2600 hsa_op_address
*addr
;
2601 addr
= gen_hsa_addr (lhs
, hbb
, &bitsize
, &bitpos
);
2603 /* Handle store to a bit field. */
2606 HSA_SORRY_AT (EXPR_LOCATION (lhs
),
2607 "support for HSA does not implement store to a bit field "
2608 "bigger than 64 bits");
2612 unsigned type_bitsize
= get_bitfield_size (bitpos
, bitsize
);
2614 /* HSAIL does not support MOV insn with 16-bits integers. */
2615 if (type_bitsize
< 32)
2618 if (bitpos
|| (bitsize
&& type_bitsize
!= bitsize
))
2620 unsigned HOST_WIDE_INT mask
= 0;
2621 BrigType16_t mem_type
2622 = get_integer_type_by_bytes (type_bitsize
/ BITS_PER_UNIT
,
2623 !TYPE_UNSIGNED (TREE_TYPE (lhs
)));
2625 for (unsigned i
= 0; i
< type_bitsize
; i
++)
2626 if (i
< bitpos
|| i
>= bitpos
+ bitsize
)
2627 mask
|= ((unsigned HOST_WIDE_INT
)1 << i
);
2629 hsa_op_reg
*value_reg
= new hsa_op_reg (mem_type
);
2631 req_align
= hsa_bitmemref_alignment (lhs
);
2632 /* Load value from memory. */
2633 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mem_type
,
2635 mem
->set_align (req_align
);
2636 hbb
->append_insn (mem
);
2638 /* AND the loaded value with prepared mask. */
2639 hsa_op_reg
*cleared_reg
= new hsa_op_reg (mem_type
);
2642 = get_integer_type_by_bytes (type_bitsize
/ BITS_PER_UNIT
, false);
2643 hsa_op_immed
*c
= new hsa_op_immed (mask
, t
);
2645 hsa_insn_basic
*clearing
2646 = new hsa_insn_basic (3, BRIG_OPCODE_AND
, mem_type
, cleared_reg
,
2648 hbb
->append_insn (clearing
);
2650 /* Shift to left a value that is going to be stored. */
2651 hsa_op_reg
*new_value_reg
= new hsa_op_reg (mem_type
);
2653 hsa_insn_basic
*basic
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, mem_type
,
2654 new_value_reg
, src
);
2655 hsa_fixup_mov_insn_type (basic
);
2656 hbb
->append_insn (basic
);
2660 hsa_op_reg
*shifted_value_reg
= new hsa_op_reg (mem_type
);
2661 c
= new hsa_op_immed (bitpos
, BRIG_TYPE_U32
);
2663 hsa_insn_basic
*basic
2664 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, mem_type
,
2665 shifted_value_reg
, new_value_reg
, c
);
2666 hbb
->append_insn (basic
);
2668 new_value_reg
= shifted_value_reg
;
2671 /* OR the prepared value with prepared chunk loaded from memory. */
2672 hsa_op_reg
*prepared_reg
= new hsa_op_reg (mem_type
);
2673 basic
= new hsa_insn_basic (3, BRIG_OPCODE_OR
, mem_type
, prepared_reg
,
2674 new_value_reg
, cleared_reg
);
2675 hbb
->append_insn (basic
);
2681 req_align
= hsa_object_alignment (lhs
);
2683 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
, addr
);
2684 mem
->set_align (req_align
);
2686 /* The HSAIL verifier has another constraint: if the source is an immediate
2687 then it must match the destination type. If it's a register the low bits
2688 will be used for sub-word stores. We're always allocating new operands so
2689 we can modify the above in place. */
2690 if (hsa_op_immed
*imm
= dyn_cast
<hsa_op_immed
*> (src
))
2692 if (!hsa_type_packed_p (imm
->m_type
))
2693 imm
->m_type
= mem
->m_type
;
2696 /* ...and all vector immediates apparently need to be vectors of
2698 unsigned bs
= hsa_type_bit_size (imm
->m_type
);
2699 gcc_assert (bs
== hsa_type_bit_size (mem
->m_type
));
2703 imm
->m_type
= BRIG_TYPE_U8X4
;
2706 imm
->m_type
= BRIG_TYPE_U8X8
;
2709 imm
->m_type
= BRIG_TYPE_U8X16
;
2717 hbb
->append_insn (mem
);
2720 /* Generate memory copy instructions that are going to be used
2721 for copying a SRC memory to TARGET memory,
2722 represented by pointer in a register. MIN_ALIGN is minimal alignment
2723 of provided HSA addresses. */
2726 gen_hsa_memory_copy (hsa_bb
*hbb
, hsa_op_address
*target
, hsa_op_address
*src
,
2727 unsigned size
, BrigAlignment8_t min_align
)
2729 hsa_op_address
*addr
;
2732 unsigned offset
= 0;
2733 unsigned min_byte_align
= hsa_byte_alignment (min_align
);
2747 if (s
> min_byte_align
)
2750 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2752 hsa_op_reg
*tmp
= new hsa_op_reg (t
);
2753 addr
= new hsa_op_address (src
->m_symbol
, src
->m_reg
,
2754 src
->m_imm_offset
+ offset
);
2755 mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, t
, tmp
, addr
);
2756 hbb
->append_insn (mem
);
2758 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2759 target
->m_imm_offset
+ offset
);
2760 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, tmp
, addr
);
2761 hbb
->append_insn (mem
);
2767 /* Create a memset mask that is created by copying a CONSTANT byte value
2768 to an integer of BYTE_SIZE bytes. */
2770 static unsigned HOST_WIDE_INT
2771 build_memset_value (unsigned HOST_WIDE_INT constant
, unsigned byte_size
)
2776 HOST_WIDE_INT v
= constant
;
2778 for (unsigned i
= 1; i
< byte_size
; i
++)
2779 v
|= constant
<< (8 * i
);
2784 /* Generate memory set instructions that are going to be used
2785 for setting a CONSTANT byte value to TARGET memory of SIZE bytes.
2786 MIN_ALIGN is minimal alignment of provided HSA addresses. */
2789 gen_hsa_memory_set (hsa_bb
*hbb
, hsa_op_address
*target
,
2790 unsigned HOST_WIDE_INT constant
,
2791 unsigned size
, BrigAlignment8_t min_align
)
2793 hsa_op_address
*addr
;
2796 unsigned offset
= 0;
2797 unsigned min_byte_align
= hsa_byte_alignment (min_align
);
2811 if (s
> min_byte_align
)
2814 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2815 target
->m_imm_offset
+ offset
);
2817 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2818 HOST_WIDE_INT c
= build_memset_value (constant
, s
);
2820 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, new hsa_op_immed (c
, t
),
2822 hbb
->append_insn (mem
);
2828 /* Generate HSAIL instructions for a single assignment
2829 of an empty constructor to an ADDR_LHS. Constructor is passed as a
2830 tree RHS and all instructions are appended to HBB. ALIGN is
2831 alignment of the address. */
2834 gen_hsa_ctor_assignment (hsa_op_address
*addr_lhs
, tree rhs
, hsa_bb
*hbb
,
2835 BrigAlignment8_t align
)
2837 if (CONSTRUCTOR_NELTS (rhs
))
2839 HSA_SORRY_AT (EXPR_LOCATION (rhs
),
2840 "support for HSA does not implement load from constructor");
2844 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2845 gen_hsa_memory_set (hbb
, addr_lhs
, 0, size
, align
);
2848 /* Generate HSA instructions for a single assignment of RHS to LHS.
2849 HBB is the basic block they will be appended to. */
2852 gen_hsa_insns_for_single_assignment (tree lhs
, tree rhs
, hsa_bb
*hbb
)
2854 if (TREE_CODE (lhs
) == SSA_NAME
)
2856 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
2857 if (hsa_seen_error ())
2860 gen_hsa_insns_for_load (dest
, rhs
, TREE_TYPE (lhs
), hbb
);
2862 else if (TREE_CODE (rhs
) == SSA_NAME
2863 || (is_gimple_min_invariant (rhs
) && TREE_CODE (rhs
) != STRING_CST
))
2865 /* Store to memory. */
2866 hsa_op_base
*src
= hsa_reg_or_immed_for_gimple_op (rhs
, hbb
);
2867 if (hsa_seen_error ())
2870 gen_hsa_insns_for_store (lhs
, src
, hbb
);
2874 BrigAlignment8_t lhs_align
;
2875 hsa_op_address
*addr_lhs
= gen_hsa_addr_with_align (lhs
, hbb
,
2878 if (TREE_CODE (rhs
) == CONSTRUCTOR
)
2879 gen_hsa_ctor_assignment (addr_lhs
, rhs
, hbb
, lhs_align
);
2882 BrigAlignment8_t rhs_align
;
2883 hsa_op_address
*addr_rhs
= gen_hsa_addr_with_align (rhs
, hbb
,
2886 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2887 gen_hsa_memory_copy (hbb
, addr_lhs
, addr_rhs
, size
,
2888 MIN (lhs_align
, rhs_align
));
2893 /* Prepend before INSN a load from spill symbol of SPILL_REG. Return the
2894 register into which we loaded. If this required another register to convert
2895 from a B1 type, return it in *PTMP2, otherwise store NULL into it. We
2896 assume we are out of SSA so the returned register does not have its
2900 hsa_spill_in (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2902 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2903 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2904 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2906 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, spill_sym
->m_type
,
2908 hsa_insert_insn_before (mem
, insn
);
2911 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2913 hsa_insn_basic
*cvtinsn
;
2915 reg
= new hsa_op_reg (spill_reg
->m_type
);
2917 cvtinsn
= new hsa_insn_cvt (reg
, *ptmp2
);
2918 hsa_insert_insn_before (cvtinsn
, insn
);
2923 /* Append after INSN a store to spill symbol of SPILL_REG. Return the register
2924 from which we stored. If this required another register to convert to a B1
2925 type, return it in *PTMP2, otherwise store NULL into it. We assume we are
2926 out of SSA so the returned register does not have its use updated. */
2929 hsa_spill_out (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2931 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2932 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2933 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2934 hsa_op_reg
*returnreg
;
2938 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2940 hsa_insn_basic
*cvtinsn
;
2941 *ptmp2
= new hsa_op_reg (spill_sym
->m_type
);
2942 reg
->m_type
= spill_reg
->m_type
;
2944 cvtinsn
= new hsa_insn_cvt (*ptmp2
, returnreg
);
2945 hsa_append_insn_after (cvtinsn
, insn
);
2950 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, spill_sym
->m_type
, reg
,
2952 hsa_append_insn_after (mem
, insn
);
2956 /* Generate a comparison instruction that will compare LHS and RHS with
2957 comparison specified by CODE and put result into register DEST. DEST has to
2958 have its type set already but must not have its definition set yet.
2959 Generated instructions will be added to HBB. */
2962 gen_hsa_cmp_insn_from_gimple (enum tree_code code
, tree lhs
, tree rhs
,
2963 hsa_op_reg
*dest
, hsa_bb
*hbb
)
2965 BrigCompareOperation8_t compare
;
2970 compare
= BRIG_COMPARE_LT
;
2973 compare
= BRIG_COMPARE_LE
;
2976 compare
= BRIG_COMPARE_GT
;
2979 compare
= BRIG_COMPARE_GE
;
2982 compare
= BRIG_COMPARE_EQ
;
2985 compare
= BRIG_COMPARE_NE
;
2987 case UNORDERED_EXPR
:
2988 compare
= BRIG_COMPARE_NAN
;
2991 compare
= BRIG_COMPARE_NUM
;
2994 compare
= BRIG_COMPARE_LTU
;
2997 compare
= BRIG_COMPARE_LEU
;
3000 compare
= BRIG_COMPARE_GTU
;
3003 compare
= BRIG_COMPARE_GEU
;
3006 compare
= BRIG_COMPARE_EQU
;
3009 compare
= BRIG_COMPARE_NEU
;
3013 HSA_SORRY_ATV (EXPR_LOCATION (lhs
),
3014 "support for HSA does not implement comparison tree "
3015 "code %s\n", get_tree_code_name (code
));
3019 /* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer)
3020 as a result of comparison. */
3022 BrigType16_t dest_type
= hsa_type_integer_p (dest
->m_type
)
3023 ? (BrigType16_t
) BRIG_TYPE_B1
: dest
->m_type
;
3025 hsa_insn_cmp
*cmp
= new hsa_insn_cmp (compare
, dest_type
);
3026 hsa_op_with_type
*op1
= hsa_reg_or_immed_for_gimple_op (lhs
, hbb
);
3027 cmp
->set_op (1, op1
->extend_int_to_32bit (hbb
));
3028 hsa_op_with_type
*op2
= hsa_reg_or_immed_for_gimple_op (rhs
, hbb
);
3029 cmp
->set_op (2, op2
->extend_int_to_32bit (hbb
));
3031 hbb
->append_insn (cmp
);
3032 cmp
->set_output_in_type (dest
, 0, hbb
);
3035 /* Generate an unary instruction with OPCODE and append it to a basic block
3036 HBB. The instruction uses DEST as a destination and OP1
3037 as a single operand. */
3040 gen_hsa_unary_operation (BrigOpcode opcode
, hsa_op_reg
*dest
,
3041 hsa_op_with_type
*op1
, hsa_bb
*hbb
)
3043 gcc_checking_assert (dest
);
3044 hsa_insn_basic
*insn
;
3046 if (opcode
== BRIG_OPCODE_MOV
&& hsa_needs_cvt (dest
->m_type
, op1
->m_type
))
3048 insn
= new hsa_insn_cvt (dest
, op1
);
3049 hbb
->append_insn (insn
);
3053 op1
= op1
->extend_int_to_32bit (hbb
);
3054 if (opcode
== BRIG_OPCODE_FIRSTBIT
|| opcode
== BRIG_OPCODE_LASTBIT
)
3056 BrigType16_t srctype
= hsa_type_integer_p (op1
->m_type
) ? op1
->m_type
3057 : hsa_unsigned_type_for_type (op1
->m_type
);
3058 insn
= new hsa_insn_srctype (2, opcode
, BRIG_TYPE_U32
, srctype
, NULL
,
3063 BrigType16_t optype
= hsa_extend_inttype_to_32bit (dest
->m_type
);
3064 insn
= new hsa_insn_basic (2, opcode
, optype
, NULL
, op1
);
3066 if (opcode
== BRIG_OPCODE_MOV
)
3067 hsa_fixup_mov_insn_type (insn
);
3068 else if (opcode
== BRIG_OPCODE_ABS
|| opcode
== BRIG_OPCODE_NEG
)
3070 /* ABS and NEG only exist in _s form :-/ */
3071 if (insn
->m_type
== BRIG_TYPE_U32
)
3072 insn
->m_type
= BRIG_TYPE_S32
;
3073 else if (insn
->m_type
== BRIG_TYPE_U64
)
3074 insn
->m_type
= BRIG_TYPE_S64
;
3078 hbb
->append_insn (insn
);
3079 insn
->set_output_in_type (dest
, 0, hbb
);
3082 /* Generate a binary instruction with OPCODE and append it to a basic block
3083 HBB. The instruction uses DEST as a destination and operands OP1
3087 gen_hsa_binary_operation (int opcode
, hsa_op_reg
*dest
,
3088 hsa_op_with_type
*op1
, hsa_op_with_type
*op2
,
3091 gcc_checking_assert (dest
);
3093 BrigType16_t optype
= hsa_extend_inttype_to_32bit (dest
->m_type
);
3094 op1
= op1
->extend_int_to_32bit (hbb
);
3095 op2
= op2
->extend_int_to_32bit (hbb
);
3097 if ((opcode
== BRIG_OPCODE_SHL
|| opcode
== BRIG_OPCODE_SHR
)
3098 && is_a
<hsa_op_immed
*> (op2
))
3100 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3101 i
->set_type (BRIG_TYPE_U32
);
3103 if ((opcode
== BRIG_OPCODE_OR
3104 || opcode
== BRIG_OPCODE_XOR
3105 || opcode
== BRIG_OPCODE_AND
)
3106 && is_a
<hsa_op_immed
*> (op2
))
3108 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3109 i
->set_type (hsa_unsigned_type_for_type (i
->m_type
));
3112 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, optype
, NULL
,
3114 hbb
->append_insn (insn
);
3115 insn
->set_output_in_type (dest
, 0, hbb
);
3118 /* Generate HSA instructions for a single assignment. HBB is the basic block
3119 they will be appended to. */
3122 gen_hsa_insns_for_operation_assignment (gimple
*assign
, hsa_bb
*hbb
)
3124 tree_code code
= gimple_assign_rhs_code (assign
);
3125 gimple_rhs_class rhs_class
= get_gimple_rhs_class (gimple_expr_code (assign
));
3127 tree lhs
= gimple_assign_lhs (assign
);
3128 tree rhs1
= gimple_assign_rhs1 (assign
);
3129 tree rhs2
= gimple_assign_rhs2 (assign
);
3130 tree rhs3
= gimple_assign_rhs3 (assign
);
3138 /* The opcode is changed to BRIG_OPCODE_CVT if BRIG types
3139 needs a conversion. */
3140 opcode
= BRIG_OPCODE_MOV
;
3144 case POINTER_PLUS_EXPR
:
3145 opcode
= BRIG_OPCODE_ADD
;
3148 opcode
= BRIG_OPCODE_SUB
;
3151 opcode
= BRIG_OPCODE_MUL
;
3153 case MULT_HIGHPART_EXPR
:
3154 opcode
= BRIG_OPCODE_MULHI
;
3157 case TRUNC_DIV_EXPR
:
3158 case EXACT_DIV_EXPR
:
3159 opcode
= BRIG_OPCODE_DIV
;
3162 case FLOOR_DIV_EXPR
:
3163 case ROUND_DIV_EXPR
:
3164 HSA_SORRY_AT (gimple_location (assign
),
3165 "support for HSA does not implement CEIL_DIV_EXPR, "
3166 "FLOOR_DIV_EXPR or ROUND_DIV_EXPR");
3168 case TRUNC_MOD_EXPR
:
3169 opcode
= BRIG_OPCODE_REM
;
3172 case FLOOR_MOD_EXPR
:
3173 case ROUND_MOD_EXPR
:
3174 HSA_SORRY_AT (gimple_location (assign
),
3175 "support for HSA does not implement CEIL_MOD_EXPR, "
3176 "FLOOR_MOD_EXPR or ROUND_MOD_EXPR");
3179 opcode
= BRIG_OPCODE_NEG
;
3182 /* There is a native HSA instruction for scalar FMAs but not for vector
3184 if (TREE_CODE (TREE_TYPE (lhs
)) == VECTOR_TYPE
)
3187 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3188 hsa_op_with_type
*op1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3189 hsa_op_with_type
*op2
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3190 hsa_op_with_type
*op3
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3191 hsa_op_reg
*tmp
= new hsa_op_reg (dest
->m_type
);
3192 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp
, op1
, op2
, hbb
);
3193 gen_hsa_binary_operation (BRIG_OPCODE_ADD
, dest
, tmp
, op3
, hbb
);
3196 opcode
= BRIG_OPCODE_MAD
;
3199 opcode
= BRIG_OPCODE_MIN
;
3202 opcode
= BRIG_OPCODE_MAX
;
3205 opcode
= BRIG_OPCODE_ABS
;
3208 opcode
= BRIG_OPCODE_SHL
;
3211 opcode
= BRIG_OPCODE_SHR
;
3216 hsa_insn_basic
*insn
= NULL
;
3217 int code1
= code
== LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3218 int code2
= code
!= LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3219 BrigType16_t btype
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
3222 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3223 hsa_op_reg
*op1
= new hsa_op_reg (btype
);
3224 hsa_op_reg
*op2
= new hsa_op_reg (btype
);
3225 hsa_op_with_type
*shift1
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3227 tree type
= TREE_TYPE (rhs2
);
3228 unsigned HOST_WIDE_INT bitsize
= TREE_INT_CST_LOW (TYPE_SIZE (type
));
3230 hsa_op_with_type
*shift2
= NULL
;
3231 if (TREE_CODE (rhs2
) == INTEGER_CST
)
3232 shift2
= new hsa_op_immed (bitsize
- tree_to_uhwi (rhs2
),
3234 else if (TREE_CODE (rhs2
) == SSA_NAME
)
3236 hsa_op_reg
*s
= hsa_cfun
->reg_for_gimple_ssa (rhs2
);
3237 s
= as_a
<hsa_op_reg
*> (s
->extend_int_to_32bit (hbb
));
3238 hsa_op_reg
*d
= new hsa_op_reg (s
->m_type
);
3239 hsa_op_immed
*size_imm
= new hsa_op_immed (bitsize
, BRIG_TYPE_U32
);
3241 insn
= new hsa_insn_basic (3, BRIG_OPCODE_SUB
, d
->m_type
,
3243 hbb
->append_insn (insn
);
3250 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3251 gen_hsa_binary_operation (code1
, op1
, src
, shift1
, hbb
);
3252 gen_hsa_binary_operation (code2
, op2
, src
, shift2
, hbb
);
3253 gen_hsa_binary_operation (BRIG_OPCODE_OR
, dest
, op1
, op2
, hbb
);
3258 opcode
= BRIG_OPCODE_OR
;
3261 opcode
= BRIG_OPCODE_XOR
;
3264 opcode
= BRIG_OPCODE_AND
;
3267 opcode
= BRIG_OPCODE_NOT
;
3269 case FIX_TRUNC_EXPR
:
3271 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3272 hsa_op_with_type
*v
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3274 if (hsa_needs_cvt (dest
->m_type
, v
->m_type
))
3276 hsa_op_reg
*tmp
= new hsa_op_reg (v
->m_type
);
3278 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3279 tmp
->m_type
, tmp
, v
);
3280 hbb
->append_insn (insn
);
3282 hsa_insn_basic
*cvtinsn
= new hsa_insn_cvt (dest
, tmp
);
3283 hbb
->append_insn (cvtinsn
);
3287 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3288 dest
->m_type
, dest
, v
);
3289 hbb
->append_insn (insn
);
3294 opcode
= BRIG_OPCODE_TRUNC
;
3303 case UNORDERED_EXPR
:
3313 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3315 gen_hsa_cmp_insn_from_gimple (code
, rhs1
, rhs2
, dest
, hbb
);
3321 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3322 hsa_op_with_type
*ctrl
= NULL
;
3325 if (CONSTANT_CLASS_P (cond
) || TREE_CODE (cond
) == SSA_NAME
)
3326 ctrl
= hsa_reg_or_immed_for_gimple_op (cond
, hbb
);
3329 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
3331 gen_hsa_cmp_insn_from_gimple (TREE_CODE (cond
),
3332 TREE_OPERAND (cond
, 0),
3333 TREE_OPERAND (cond
, 1),
3339 hsa_op_with_type
*op2
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3340 hsa_op_with_type
*op3
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3341 op2
= op2
->extend_int_to_32bit (hbb
);
3342 op3
= op3
->extend_int_to_32bit (hbb
);
3344 BrigType16_t type
= hsa_extend_inttype_to_32bit (dest
->m_type
);
3345 BrigType16_t utype
= hsa_unsigned_type_for_type (type
);
3346 if (is_a
<hsa_op_immed
*> (op2
))
3347 op2
->m_type
= utype
;
3348 if (is_a
<hsa_op_immed
*> (op3
))
3349 op3
->m_type
= utype
;
3351 hsa_insn_basic
*insn
3352 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
,
3353 hsa_bittype_for_type (type
),
3354 NULL
, ctrl
, op2
, op3
);
3356 hbb
->append_insn (insn
);
3357 insn
->set_output_in_type (dest
, 0, hbb
);
3363 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3364 hsa_op_with_type
*rhs1_reg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3365 rhs1_reg
= rhs1_reg
->extend_int_to_32bit (hbb
);
3366 hsa_op_with_type
*rhs2_reg
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3367 rhs2_reg
= rhs2_reg
->extend_int_to_32bit (hbb
);
3369 if (hsa_seen_error ())
3372 BrigType16_t src_type
= hsa_bittype_for_type (rhs1_reg
->m_type
);
3373 rhs1_reg
= rhs1_reg
->get_in_type (src_type
, hbb
);
3374 rhs2_reg
= rhs2_reg
->get_in_type (src_type
, hbb
);
3376 hsa_insn_packed
*insn
3377 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
, src_type
,
3378 dest
, rhs1_reg
, rhs2_reg
);
3379 hbb
->append_insn (insn
);
3384 /* Implement others as we come across them. */
3385 HSA_SORRY_ATV (gimple_location (assign
),
3386 "support for HSA does not implement operation %s",
3387 get_tree_code_name (code
));
3392 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3393 hsa_op_with_type
*op1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3394 hsa_op_with_type
*op2
3395 = rhs2
? hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
) : NULL
;
3397 if (hsa_seen_error ())
3402 case GIMPLE_TERNARY_RHS
:
3404 hsa_op_with_type
*op3
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3405 op3
= op3
->extend_int_to_32bit (hbb
);
3406 hsa_insn_basic
*insn
= new hsa_insn_basic (4, opcode
, dest
->m_type
, dest
,
3408 hbb
->append_insn (insn
);
3412 case GIMPLE_BINARY_RHS
:
3413 gen_hsa_binary_operation (opcode
, dest
, op1
, op2
, hbb
);
3416 case GIMPLE_UNARY_RHS
:
3417 gen_hsa_unary_operation (opcode
, dest
, op1
, hbb
);
3424 /* Generate HSA instructions for a given gimple condition statement COND.
3425 Instructions will be appended to HBB, which also needs to be the
3426 corresponding structure to the basic_block of COND. */
3429 gen_hsa_insns_for_cond_stmt (gimple
*cond
, hsa_bb
*hbb
)
3431 hsa_op_reg
*ctrl
= new hsa_op_reg (BRIG_TYPE_B1
);
3434 gen_hsa_cmp_insn_from_gimple (gimple_cond_code (cond
),
3435 gimple_cond_lhs (cond
),
3436 gimple_cond_rhs (cond
),
3439 cbr
= new hsa_insn_cbr (ctrl
);
3440 hbb
->append_insn (cbr
);
3443 /* Maximum number of elements in a jump table for an HSA SBR instruction. */
3445 #define HSA_MAXIMUM_SBR_LABELS 16
3447 /* Return lowest value of a switch S that is handled in a non-default
3451 get_switch_low (gswitch
*s
)
3453 unsigned labels
= gimple_switch_num_labels (s
);
3454 gcc_checking_assert (labels
>= 1);
3456 return CASE_LOW (gimple_switch_label (s
, 1));
3459 /* Return highest value of a switch S that is handled in a non-default
3463 get_switch_high (gswitch
*s
)
3465 unsigned labels
= gimple_switch_num_labels (s
);
3467 /* Compare last label to maximum number of labels. */
3468 tree label
= gimple_switch_label (s
, labels
- 1);
3469 tree low
= CASE_LOW (label
);
3470 tree high
= CASE_HIGH (label
);
3472 return high
!= NULL_TREE
? high
: low
;
3476 get_switch_size (gswitch
*s
)
3478 return int_const_binop (MINUS_EXPR
, get_switch_high (s
), get_switch_low (s
));
3481 /* Generate HSA instructions for a given gimple switch.
3482 Instructions will be appended to HBB. */
3485 gen_hsa_insns_for_switch_stmt (gswitch
*s
, hsa_bb
*hbb
)
3487 gimple_stmt_iterator it
= gsi_for_stmt (s
);
3490 /* Create preambule that verifies that index - lowest_label >= 0. */
3491 edge e
= split_block (hbb
->m_bb
, gsi_stmt (it
));
3492 e
->flags
&= ~EDGE_FALLTHRU
;
3493 e
->flags
|= EDGE_TRUE_VALUE
;
3495 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
3496 tree index_tree
= gimple_switch_index (s
);
3497 tree lowest
= get_switch_low (s
);
3498 tree highest
= get_switch_high (s
);
3500 hsa_op_reg
*index
= hsa_cfun
->reg_for_gimple_ssa (index_tree
);
3501 index
= as_a
<hsa_op_reg
*> (index
->extend_int_to_32bit (hbb
));
3503 hsa_op_reg
*cmp1_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3504 hsa_op_immed
*cmp1_immed
= new hsa_op_immed (lowest
, true);
3505 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_GE
, cmp1_reg
->m_type
,
3506 cmp1_reg
, index
, cmp1_immed
));
3508 hsa_op_reg
*cmp2_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3509 hsa_op_immed
*cmp2_immed
= new hsa_op_immed (highest
, true);
3510 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_LE
, cmp2_reg
->m_type
,
3511 cmp2_reg
, index
, cmp2_immed
));
3513 hsa_op_reg
*cmp_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3514 hbb
->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_AND
, cmp_reg
->m_type
,
3515 cmp_reg
, cmp1_reg
, cmp2_reg
));
3517 hbb
->append_insn (new hsa_insn_cbr (cmp_reg
));
3519 tree default_label
= gimple_switch_default_label (s
);
3520 basic_block default_label_bb
= label_to_block_fn (func
,
3521 CASE_LABEL (default_label
));
3523 if (!gimple_seq_empty_p (phi_nodes (default_label_bb
)))
3525 default_label_bb
= split_edge (find_edge (e
->dest
, default_label_bb
));
3526 hsa_init_new_bb (default_label_bb
);
3529 make_edge (e
->src
, default_label_bb
, EDGE_FALSE_VALUE
);
3531 hsa_cfun
->m_modified_cfg
= true;
3533 /* Basic block with the SBR instruction. */
3534 hbb
= hsa_init_new_bb (e
->dest
);
3536 hsa_op_reg
*sub_index
= new hsa_op_reg (index
->m_type
);
3537 hbb
->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB
, sub_index
->m_type
,
3539 new hsa_op_immed (lowest
, true)));
3541 hsa_op_base
*tmp
= sub_index
->get_in_type (BRIG_TYPE_U64
, hbb
);
3542 sub_index
= as_a
<hsa_op_reg
*> (tmp
);
3543 unsigned labels
= gimple_switch_num_labels (s
);
3544 unsigned HOST_WIDE_INT size
= tree_to_uhwi (get_switch_size (s
));
3546 hsa_insn_sbr
*sbr
= new hsa_insn_sbr (sub_index
, size
+ 1);
3548 /* Prepare array with default label destination. */
3549 for (unsigned HOST_WIDE_INT i
= 0; i
<= size
; i
++)
3550 sbr
->m_jump_table
.safe_push (default_label_bb
);
3552 /* Iterate all labels and fill up the jump table. */
3553 for (unsigned i
= 1; i
< labels
; i
++)
3555 tree label
= gimple_switch_label (s
, i
);
3556 basic_block bb
= label_to_block_fn (func
, CASE_LABEL (label
));
3558 unsigned HOST_WIDE_INT sub_low
3559 = tree_to_uhwi (int_const_binop (MINUS_EXPR
, CASE_LOW (label
), lowest
));
3561 unsigned HOST_WIDE_INT sub_high
= sub_low
;
3562 tree high
= CASE_HIGH (label
);
3564 sub_high
= tree_to_uhwi (int_const_binop (MINUS_EXPR
, high
, lowest
));
3566 for (unsigned HOST_WIDE_INT j
= sub_low
; j
<= sub_high
; j
++)
3567 sbr
->m_jump_table
[j
] = bb
;
3570 hbb
->append_insn (sbr
);
3573 /* Verify that the function DECL can be handled by HSA. */
3576 verify_function_arguments (tree decl
)
3578 tree type
= TREE_TYPE (decl
);
3579 if (DECL_STATIC_CHAIN (decl
))
3581 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3582 "HSA does not support nested functions: %qD", decl
);
3585 else if (!TYPE_ARG_TYPES (type
) || stdarg_p (type
))
3587 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3588 "HSA does not support functions with variadic arguments "
3589 "(or unknown return type): %qD", decl
);
3594 /* Return BRIG type for FORMAL_ARG_TYPE. If the formal argument type is NULL,
3595 return ACTUAL_ARG_TYPE. */
3598 get_format_argument_type (tree formal_arg_type
, BrigType16_t actual_arg_type
)
3600 if (formal_arg_type
== NULL
)
3601 return actual_arg_type
;
3603 BrigType16_t decl_type
3604 = hsa_type_for_scalar_tree_type (formal_arg_type
, false);
3605 return mem_type_for_type (decl_type
);
3608 /* Generate HSA instructions for a direct call instruction.
3609 Instructions will be appended to HBB, which also needs to be the
3610 corresponding structure to the basic_block of STMT.
3611 If ASSIGN_LHS is false, do not copy HSA function result argument into the
3612 corresponding HSA representation of the gimple statement LHS. */
3615 gen_hsa_insns_for_direct_call (gimple
*stmt
, hsa_bb
*hbb
,
3616 bool assign_lhs
= true)
3618 tree decl
= gimple_call_fndecl (stmt
);
3619 verify_function_arguments (decl
);
3620 if (hsa_seen_error ())
3623 hsa_insn_call
*call_insn
= new hsa_insn_call (decl
);
3624 hsa_cfun
->m_called_functions
.safe_push (call_insn
->m_called_function
);
3626 /* Argument block start. */
3627 hsa_insn_arg_block
*arg_start
3628 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3629 hbb
->append_insn (arg_start
);
3631 tree parm_type_chain
= TYPE_ARG_TYPES (gimple_call_fntype (stmt
));
3633 /* Preparation of arguments that will be passed to function. */
3634 const unsigned args
= gimple_call_num_args (stmt
);
3635 for (unsigned i
= 0; i
< args
; ++i
)
3637 tree parm
= gimple_call_arg (stmt
, (int)i
);
3638 tree parm_decl_type
= parm_type_chain
!= NULL_TREE
3639 ? TREE_VALUE (parm_type_chain
) : NULL_TREE
;
3640 hsa_op_address
*addr
;
3642 if (AGGREGATE_TYPE_P (TREE_TYPE (parm
)))
3644 addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3645 BrigAlignment8_t align
;
3646 hsa_op_address
*src
= gen_hsa_addr_with_align (parm
, hbb
, &align
);
3647 gen_hsa_memory_copy (hbb
, addr
, src
,
3648 addr
->m_symbol
->total_byte_size (), align
);
3652 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3654 if (parm_decl_type
!= NULL
&& AGGREGATE_TYPE_P (parm_decl_type
))
3656 HSA_SORRY_AT (gimple_location (stmt
),
3657 "support for HSA does not implement an aggregate "
3658 "formal argument in a function call, while actual "
3659 "argument is not an aggregate");
3663 BrigType16_t formal_arg_type
3664 = get_format_argument_type (parm_decl_type
, src
->m_type
);
3665 if (hsa_seen_error ())
3668 if (src
->m_type
!= formal_arg_type
)
3669 src
= src
->get_in_type (formal_arg_type
, hbb
);
3672 = gen_hsa_addr_for_arg (parm_decl_type
!= NULL_TREE
?
3673 parm_decl_type
: TREE_TYPE (parm
), i
);
3674 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, formal_arg_type
,
3677 hbb
->append_insn (mem
);
3680 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3681 if (parm_type_chain
)
3682 parm_type_chain
= TREE_CHAIN (parm_type_chain
);
3685 call_insn
->m_args_code_list
= new hsa_op_code_list (args
);
3686 hbb
->append_insn (call_insn
);
3688 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
3690 tree result
= gimple_call_lhs (stmt
);
3691 hsa_insn_mem
*result_insn
= NULL
;
3692 if (!VOID_TYPE_P (result_type
))
3694 hsa_op_address
*addr
= gen_hsa_addr_for_arg (result_type
, -1);
3696 /* Even if result of a function call is unused, we have to emit
3697 declaration for the result. */
3698 if (result
&& assign_lhs
)
3700 tree lhs_type
= TREE_TYPE (result
);
3702 if (hsa_seen_error ())
3705 if (AGGREGATE_TYPE_P (lhs_type
))
3707 BrigAlignment8_t align
;
3708 hsa_op_address
*result_addr
3709 = gen_hsa_addr_with_align (result
, hbb
, &align
);
3710 gen_hsa_memory_copy (hbb
, result_addr
, addr
,
3711 addr
->m_symbol
->total_byte_size (), align
);
3716 = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type
,
3719 hsa_op_reg
*dst
= hsa_cfun
->reg_for_gimple_ssa (result
);
3720 result_insn
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
, dst
, addr
);
3721 hbb
->append_insn (result_insn
);
3725 call_insn
->m_output_arg
= addr
->m_symbol
;
3726 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3732 HSA_SORRY_AT (gimple_location (stmt
),
3733 "support for HSA does not implement an assignment of "
3734 "return value from a void function");
3738 call_insn
->m_result_code_list
= new hsa_op_code_list (0);
3741 /* Argument block end. */
3742 hsa_insn_arg_block
*arg_end
3743 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3744 hbb
->append_insn (arg_end
);
3747 /* Generate HSA instructions for a direct call of an internal fn.
3748 Instructions will be appended to HBB, which also needs to be the
3749 corresponding structure to the basic_block of STMT. */
3752 gen_hsa_insns_for_call_of_internal_fn (gimple
*stmt
, hsa_bb
*hbb
)
3754 tree lhs
= gimple_call_lhs (stmt
);
3758 tree lhs_type
= TREE_TYPE (lhs
);
3759 tree rhs1
= gimple_call_arg (stmt
, 0);
3760 tree rhs1_type
= TREE_TYPE (rhs1
);
3761 enum internal_fn fn
= gimple_call_internal_fn (stmt
);
3762 hsa_internal_fn
*ifn
3763 = new hsa_internal_fn (fn
, tree_to_uhwi (TYPE_SIZE (rhs1_type
)));
3764 hsa_insn_call
*call_insn
= new hsa_insn_call (ifn
);
3766 gcc_checking_assert (FLOAT_TYPE_P (rhs1_type
));
3768 if (!hsa_emitted_internal_decls
->find (call_insn
->m_called_internal_fn
))
3769 hsa_cfun
->m_called_internal_fns
.safe_push (call_insn
->m_called_internal_fn
);
3771 hsa_insn_arg_block
*arg_start
3772 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3773 hbb
->append_insn (arg_start
);
3775 unsigned num_args
= gimple_call_num_args (stmt
);
3777 /* Function arguments. */
3778 for (unsigned i
= 0; i
< num_args
; i
++)
3780 tree parm
= gimple_call_arg (stmt
, (int)i
);
3781 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3783 hsa_op_address
*addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3784 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, src
->m_type
,
3787 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3788 hbb
->append_insn (mem
);
3791 call_insn
->m_args_code_list
= new hsa_op_code_list (num_args
);
3792 hbb
->append_insn (call_insn
);
3794 /* Assign returned value. */
3795 hsa_op_address
*addr
= gen_hsa_addr_for_arg (lhs_type
, -1);
3797 call_insn
->m_output_arg
= addr
->m_symbol
;
3798 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3800 /* Argument block end. */
3801 hsa_insn_arg_block
*arg_end
3802 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3803 hbb
->append_insn (arg_end
);
3806 /* Generate HSA instructions for a return value instruction.
3807 Instructions will be appended to HBB, which also needs to be the
3808 corresponding structure to the basic_block of STMT. */
3811 gen_hsa_insns_for_return (greturn
*stmt
, hsa_bb
*hbb
)
3813 tree retval
= gimple_return_retval (stmt
);
3816 hsa_op_address
*addr
= new hsa_op_address (hsa_cfun
->m_output_arg
);
3818 if (AGGREGATE_TYPE_P (TREE_TYPE (retval
)))
3820 BrigAlignment8_t align
;
3821 hsa_op_address
*retval_addr
= gen_hsa_addr_with_align (retval
, hbb
,
3823 gen_hsa_memory_copy (hbb
, addr
, retval_addr
,
3824 hsa_cfun
->m_output_arg
->total_byte_size (),
3829 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (retval
),
3831 BrigType16_t mtype
= mem_type_for_type (t
);
3833 /* Store of return value. */
3834 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (retval
, hbb
);
3835 src
= src
->get_in_type (mtype
, hbb
);
3836 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
,
3838 hbb
->append_insn (mem
);
3842 /* HSAIL return instruction emission. */
3843 hsa_insn_basic
*ret
= new hsa_insn_basic (0, BRIG_OPCODE_RET
);
3844 hbb
->append_insn (ret
);
3847 /* Set OP_INDEX-th operand of the instruction to DEST, as the DEST
3848 can have a different type, conversion instructions are possibly
3852 hsa_insn_basic::set_output_in_type (hsa_op_reg
*dest
, unsigned op_index
,
3855 gcc_checking_assert (op_output_p (op_index
));
3857 if (dest
->m_type
== m_type
)
3859 set_op (op_index
, dest
);
3863 hsa_insn_basic
*insn
;
3865 if (hsa_needs_cvt (dest
->m_type
, m_type
))
3867 tmp
= new hsa_op_reg (m_type
);
3868 insn
= new hsa_insn_cvt (dest
, tmp
);
3870 else if (hsa_type_bit_size (dest
->m_type
) == hsa_type_bit_size (m_type
))
3872 /* When output, HSA registers do not really have types, only sizes, so if
3873 the sizes match, we can use the register directly. */
3874 set_op (op_index
, dest
);
3879 tmp
= new hsa_op_reg (m_type
);
3880 insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
,
3881 dest
, tmp
->get_in_type (dest
->m_type
, hbb
));
3882 hsa_fixup_mov_insn_type (insn
);
3884 set_op (op_index
, tmp
);
3885 hbb
->append_insn (insn
);
3888 /* Generate instruction OPCODE to query a property of HSA grid along the
3889 given DIMENSION. Store result into DEST and append the instruction to
3893 query_hsa_grid_dim (hsa_op_reg
*dest
, int opcode
, hsa_op_immed
*dimension
,
3896 hsa_insn_basic
*insn
= new hsa_insn_basic (2, opcode
, BRIG_TYPE_U32
, NULL
,
3898 hbb
->append_insn (insn
);
3899 insn
->set_output_in_type (dest
, 0, hbb
);
3902 /* Generate instruction OPCODE to query a property of HSA grid along the given
3903 dimension which is an immediate in first argument of STMT. Store result
3904 into the register corresponding to LHS of STMT and append the instruction to
3908 query_hsa_grid_dim (gimple
*stmt
, int opcode
, hsa_bb
*hbb
)
3910 tree lhs
= gimple_call_lhs (dyn_cast
<gcall
*> (stmt
));
3911 if (lhs
== NULL_TREE
)
3914 tree arg
= gimple_call_arg (stmt
, 0);
3915 unsigned HOST_WIDE_INT dim
= 5;
3916 if (tree_fits_uhwi_p (arg
))
3917 dim
= tree_to_uhwi (arg
);
3920 HSA_SORRY_AT (gimple_location (stmt
),
3921 "HSA grid query dimension must be immediate constant 0, 1 "
3926 hsa_op_immed
*hdim
= new hsa_op_immed (dim
, (BrigKind16_t
) BRIG_TYPE_U32
);
3927 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3928 query_hsa_grid_dim (dest
, opcode
, hdim
, hbb
);
3931 /* Generate instruction OPCODE to query a property of HSA grid that is
3932 independent of any dimension. Store result into the register corresponding
3933 to LHS of STMT and append the instruction to HBB. */
3936 query_hsa_grid_nodim (gimple
*stmt
, BrigOpcode16_t opcode
, hsa_bb
*hbb
)
3938 tree lhs
= gimple_call_lhs (dyn_cast
<gcall
*> (stmt
));
3939 if (lhs
== NULL_TREE
)
3941 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3942 BrigType16_t brig_type
= hsa_unsigned_type_for_type (dest
->m_type
);
3943 hsa_insn_basic
*insn
= new hsa_insn_basic (1, opcode
, brig_type
, dest
);
3944 hbb
->append_insn (insn
);
3947 /* Emit instructions that set hsa_num_threads according to provided VALUE.
3948 Instructions are appended to basic block HBB. */
3951 gen_set_num_threads (tree value
, hsa_bb
*hbb
)
3953 hbb
->append_insn (new hsa_insn_comment ("omp_set_num_threads"));
3954 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (value
, hbb
);
3956 src
= src
->get_in_type (hsa_num_threads
->m_type
, hbb
);
3957 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
3959 hsa_insn_basic
*basic
3960 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
, src
, addr
);
3961 hbb
->append_insn (basic
);
3964 /* Return byte offset of a FIELD_NAME in GOMP_hsa_kernel_dispatch which
3965 is defined in plugin-hsa.c. */
3967 static HOST_WIDE_INT
3968 get_hsa_kernel_dispatch_offset (const char *field_name
)
3970 tree
*hsa_kernel_dispatch_type
= hsa_get_kernel_dispatch_type ();
3971 if (*hsa_kernel_dispatch_type
== NULL
)
3973 /* Collection of information needed for a dispatch of a kernel from a
3974 kernel. Keep in sync with libgomp's plugin-hsa.c. */
3976 *hsa_kernel_dispatch_type
= make_node (RECORD_TYPE
);
3977 tree id_f1
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3978 get_identifier ("queue"), ptr_type_node
);
3979 DECL_CHAIN (id_f1
) = NULL_TREE
;
3980 tree id_f2
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3981 get_identifier ("omp_data_memory"),
3983 DECL_CHAIN (id_f2
) = id_f1
;
3984 tree id_f3
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3985 get_identifier ("kernarg_address"),
3987 DECL_CHAIN (id_f3
) = id_f2
;
3988 tree id_f4
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3989 get_identifier ("object"),
3991 DECL_CHAIN (id_f4
) = id_f3
;
3992 tree id_f5
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3993 get_identifier ("signal"),
3995 DECL_CHAIN (id_f5
) = id_f4
;
3996 tree id_f6
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3997 get_identifier ("private_segment_size"),
3999 DECL_CHAIN (id_f6
) = id_f5
;
4000 tree id_f7
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
4001 get_identifier ("group_segment_size"),
4003 DECL_CHAIN (id_f7
) = id_f6
;
4004 tree id_f8
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
4005 get_identifier ("kernel_dispatch_count"),
4007 DECL_CHAIN (id_f8
) = id_f7
;
4008 tree id_f9
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
4009 get_identifier ("debug"),
4011 DECL_CHAIN (id_f9
) = id_f8
;
4012 tree id_f10
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
4013 get_identifier ("omp_level"),
4015 DECL_CHAIN (id_f10
) = id_f9
;
4016 tree id_f11
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
4017 get_identifier ("children_dispatches"),
4019 DECL_CHAIN (id_f11
) = id_f10
;
4020 tree id_f12
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
4021 get_identifier ("omp_num_threads"),
4023 DECL_CHAIN (id_f12
) = id_f11
;
4026 finish_builtin_struct (*hsa_kernel_dispatch_type
, "__hsa_kernel_dispatch",
4028 TYPE_ARTIFICIAL (*hsa_kernel_dispatch_type
) = 1;
4031 for (tree chain
= TYPE_FIELDS (*hsa_kernel_dispatch_type
);
4032 chain
!= NULL_TREE
; chain
= TREE_CHAIN (chain
))
4033 if (id_equal (DECL_NAME (chain
), field_name
))
4034 return int_byte_position (chain
);
4039 /* Return an HSA register that will contain number of threads for
4040 a future dispatched kernel. Instructions are added to HBB. */
4043 gen_num_threads_for_dispatch (hsa_bb
*hbb
)
4045 /* Step 1) Assign to number of threads:
4046 MIN (HSA_DEFAULT_NUM_THREADS, hsa_num_threads). */
4047 hsa_op_reg
*threads
= new hsa_op_reg (hsa_num_threads
->m_type
);
4048 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
4050 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD
, threads
->m_type
,
4053 hsa_op_immed
*limit
= new hsa_op_immed (HSA_DEFAULT_NUM_THREADS
,
4055 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
4057 = new hsa_insn_cmp (BRIG_COMPARE_LT
, r
->m_type
, r
, threads
, limit
);
4058 hbb
->append_insn (cmp
);
4060 BrigType16_t btype
= hsa_bittype_for_type (threads
->m_type
);
4061 hsa_op_reg
*tmp
= new hsa_op_reg (threads
->m_type
);
4063 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp
, r
,
4066 /* Step 2) If the number is equal to zero,
4067 return shadow->omp_num_threads. */
4068 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4070 hsa_op_reg
*shadow_thread_count
= new hsa_op_reg (BRIG_TYPE_U32
);
4072 = new hsa_op_address (shadow_reg_ptr
,
4073 get_hsa_kernel_dispatch_offset ("omp_num_threads"));
4074 hsa_insn_basic
*basic
4075 = new hsa_insn_mem (BRIG_OPCODE_LD
, shadow_thread_count
->m_type
,
4076 shadow_thread_count
, addr
);
4077 hbb
->append_insn (basic
);
4079 hsa_op_reg
*tmp2
= new hsa_op_reg (threads
->m_type
);
4080 r
= new hsa_op_reg (BRIG_TYPE_B1
);
4081 hsa_op_immed
*imm
= new hsa_op_immed (0, shadow_thread_count
->m_type
);
4082 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_EQ
, r
->m_type
, r
, tmp
, imm
));
4083 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp2
, r
,
4084 shadow_thread_count
, tmp
));
4086 hsa_op_base
*dest
= tmp2
->get_in_type (BRIG_TYPE_U16
, hbb
);
4088 return as_a
<hsa_op_reg
*> (dest
);
4091 /* Build OPCODE query for all three hsa dimensions, multiply them and store the
4092 result into DEST. */
4095 multiply_grid_dim_characteristics (hsa_op_reg
*dest
, int opcode
, hsa_bb
*hbb
)
4097 hsa_op_reg
*dimx
= new hsa_op_reg (BRIG_TYPE_U32
);
4098 query_hsa_grid_dim (dimx
, opcode
,
4099 new hsa_op_immed (0, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4100 hsa_op_reg
*dimy
= new hsa_op_reg (BRIG_TYPE_U32
);
4101 query_hsa_grid_dim (dimy
, opcode
,
4102 new hsa_op_immed (1, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4103 hsa_op_reg
*dimz
= new hsa_op_reg (BRIG_TYPE_U32
);
4104 query_hsa_grid_dim (dimz
, opcode
,
4105 new hsa_op_immed (2, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4106 hsa_op_reg
*tmp
= new hsa_op_reg (dest
->m_type
);
4107 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp
,
4108 dimx
->get_in_type (dest
->m_type
, hbb
),
4109 dimy
->get_in_type (dest
->m_type
, hbb
), hbb
);
4110 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, dest
, tmp
,
4111 dimz
->get_in_type (dest
->m_type
, hbb
), hbb
);
4114 /* Emit instructions that assign number of threads to lhs of gimple STMT.
4115 Instructions are appended to basic block HBB. */
4118 gen_get_num_threads (gimple
*stmt
, hsa_bb
*hbb
)
4120 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4123 hbb
->append_insn (new hsa_insn_comment ("omp_get_num_threads"));
4124 tree lhs
= gimple_call_lhs (stmt
);
4125 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4126 multiply_grid_dim_characteristics (dest
, BRIG_OPCODE_CURRENTWORKGROUPSIZE
,
4130 /* Emit instructions that assign number of teams to lhs of gimple STMT.
4131 Instructions are appended to basic block HBB. */
4134 gen_get_num_teams (gimple
*stmt
, hsa_bb
*hbb
)
4136 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4139 hbb
->append_insn (new hsa_insn_comment ("omp_get_num_teams"));
4140 tree lhs
= gimple_call_lhs (stmt
);
4141 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4142 multiply_grid_dim_characteristics (dest
, BRIG_OPCODE_GRIDGROUPS
, hbb
);
4145 /* Emit instructions that assign a team number to lhs of gimple STMT.
4146 Instructions are appended to basic block HBB. */
4149 gen_get_team_num (gimple
*stmt
, hsa_bb
*hbb
)
4151 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4154 hbb
->append_insn (new hsa_insn_comment ("omp_get_team_num"));
4155 tree lhs
= gimple_call_lhs (stmt
);
4156 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4158 hsa_op_reg
*gnum_x
= new hsa_op_reg (BRIG_TYPE_U32
);
4159 query_hsa_grid_dim (gnum_x
, BRIG_OPCODE_GRIDGROUPS
,
4160 new hsa_op_immed (0, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4161 hsa_op_reg
*gnum_y
= new hsa_op_reg (BRIG_TYPE_U32
);
4162 query_hsa_grid_dim (gnum_y
, BRIG_OPCODE_GRIDGROUPS
,
4163 new hsa_op_immed (1, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4165 hsa_op_reg
*gno_z
= new hsa_op_reg (BRIG_TYPE_U32
);
4166 query_hsa_grid_dim (gno_z
, BRIG_OPCODE_WORKGROUPID
,
4167 new hsa_op_immed (2, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4169 hsa_op_reg
*tmp1
= new hsa_op_reg (dest
->m_type
);
4170 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp1
,
4171 gnum_x
->get_in_type (dest
->m_type
, hbb
),
4172 gnum_y
->get_in_type (dest
->m_type
, hbb
), hbb
);
4173 hsa_op_reg
*tmp2
= new hsa_op_reg (dest
->m_type
);
4174 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp2
, tmp1
,
4175 gno_z
->get_in_type (dest
->m_type
, hbb
), hbb
);
4177 hsa_op_reg
*gno_y
= new hsa_op_reg (BRIG_TYPE_U32
);
4178 query_hsa_grid_dim (gno_y
, BRIG_OPCODE_WORKGROUPID
,
4179 new hsa_op_immed (1, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4180 hsa_op_reg
*tmp3
= new hsa_op_reg (dest
->m_type
);
4181 gen_hsa_binary_operation (BRIG_OPCODE_MUL
, tmp3
,
4182 gnum_x
->get_in_type (dest
->m_type
, hbb
),
4183 gno_y
->get_in_type (dest
->m_type
, hbb
), hbb
);
4184 hsa_op_reg
*tmp4
= new hsa_op_reg (dest
->m_type
);
4185 gen_hsa_binary_operation (BRIG_OPCODE_ADD
, tmp4
, tmp3
, tmp2
, hbb
);
4186 hsa_op_reg
*gno_x
= new hsa_op_reg (BRIG_TYPE_U32
);
4187 query_hsa_grid_dim (gno_x
, BRIG_OPCODE_WORKGROUPID
,
4188 new hsa_op_immed (0, (BrigKind16_t
) BRIG_TYPE_U32
), hbb
);
4189 gen_hsa_binary_operation (BRIG_OPCODE_ADD
, dest
, tmp4
,
4190 gno_x
->get_in_type (dest
->m_type
, hbb
), hbb
);
4193 /* Emit instructions that get levels-var ICV to lhs of gimple STMT.
4194 Instructions are appended to basic block HBB. */
4197 gen_get_level (gimple
*stmt
, hsa_bb
*hbb
)
4199 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4202 hbb
->append_insn (new hsa_insn_comment ("omp_get_level"));
4204 tree lhs
= gimple_call_lhs (stmt
);
4205 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4207 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4208 if (shadow_reg_ptr
== NULL
)
4210 HSA_SORRY_AT (gimple_location (stmt
),
4211 "support for HSA does not implement omp_get_level called "
4212 "from a function not being inlined within a kernel");
4216 hsa_op_address
*addr
4217 = new hsa_op_address (shadow_reg_ptr
,
4218 get_hsa_kernel_dispatch_offset ("omp_level"));
4220 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, BRIG_TYPE_U64
,
4221 (hsa_op_base
*) NULL
, addr
);
4222 hbb
->append_insn (mem
);
4223 mem
->set_output_in_type (dest
, 0, hbb
);
4226 /* Emit instruction that implement omp_get_max_threads of gimple STMT. */
4229 gen_get_max_threads (gimple
*stmt
, hsa_bb
*hbb
)
4231 tree lhs
= gimple_call_lhs (stmt
);
4235 hbb
->append_insn (new hsa_insn_comment ("omp_get_max_threads"));
4237 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4238 hsa_op_with_type
*num_theads_reg
= gen_num_threads_for_dispatch (hbb
)
4239 ->get_in_type (dest
->m_type
, hbb
);
4240 hsa_build_append_simple_mov (dest
, num_theads_reg
, hbb
);
4243 /* Emit instructions that implement alloca builtin gimple STMT.
4244 Instructions are appended to basic block HBB. */
4247 gen_hsa_alloca (gcall
*call
, hsa_bb
*hbb
)
4249 tree lhs
= gimple_call_lhs (call
);
4250 if (lhs
== NULL_TREE
)
4253 built_in_function fn
= DECL_FUNCTION_CODE (gimple_call_fndecl (call
));
4255 gcc_checking_assert (ALLOCA_FUNCTION_CODE_P (fn
));
4257 unsigned bit_alignment
= 0;
4259 if (fn
!= BUILT_IN_ALLOCA
)
4261 tree alignment_tree
= gimple_call_arg (call
, 1);
4262 if (TREE_CODE (alignment_tree
) != INTEGER_CST
)
4264 HSA_SORRY_ATV (gimple_location (call
),
4265 "support for HSA does not implement "
4266 "__builtin_alloca_with_align with a non-constant "
4267 "alignment: %E", alignment_tree
);
4270 bit_alignment
= tree_to_uhwi (alignment_tree
);
4273 tree rhs1
= gimple_call_arg (call
, 0);
4274 hsa_op_with_type
*size
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
)
4275 ->get_in_type (BRIG_TYPE_U32
, hbb
);
4276 hsa_op_with_type
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4279 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
));
4280 hsa_insn_alloca
*a
= new hsa_insn_alloca (tmp
, size
, bit_alignment
);
4281 hbb
->append_insn (a
);
4284 = new hsa_insn_seg (BRIG_OPCODE_STOF
,
4285 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
4286 tmp
->m_type
, BRIG_SEGMENT_PRIVATE
, dest
, tmp
);
4287 hbb
->append_insn (seg
);
4290 /* Emit instructions that implement clrsb builtin STMT:
4291 Returns the number of leading redundant sign bits in x, i.e. the number
4292 of bits following the most significant bit that are identical to it.
4293 There are no special cases for 0 or other values.
4294 Instructions are appended to basic block HBB. */
4297 gen_hsa_clrsb (gcall
*call
, hsa_bb
*hbb
)
4299 tree lhs
= gimple_call_lhs (call
);
4300 if (lhs
== NULL_TREE
)
4303 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4304 tree rhs1
= gimple_call_arg (call
, 0);
4305 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4306 arg
->extend_int_to_32bit (hbb
);
4307 BrigType16_t bittype
= hsa_bittype_for_type (arg
->m_type
);
4308 unsigned bitsize
= tree_to_uhwi (TYPE_SIZE (TREE_TYPE (rhs1
)));
4310 /* FIRSTBIT instruction is defined just for 32 and 64-bits wide integers. */
4311 gcc_checking_assert (bitsize
== 32 || bitsize
== 64);
4313 /* Set true to MOST_SIG if the most significant bit is set to one. */
4314 hsa_op_immed
*c
= new hsa_op_immed (1ul << (bitsize
- 1),
4315 hsa_uint_for_bitsize (bitsize
));
4317 hsa_op_reg
*and_reg
= new hsa_op_reg (bittype
);
4318 gen_hsa_binary_operation (BRIG_OPCODE_AND
, and_reg
, arg
, c
, hbb
);
4320 hsa_op_reg
*most_sign
= new hsa_op_reg (BRIG_TYPE_B1
);
4322 = new hsa_insn_cmp (BRIG_COMPARE_EQ
, most_sign
->m_type
, most_sign
,
4324 hbb
->append_insn (cmp
);
4326 /* If the most significant bit is one, negate the input. Otherwise
4327 shift the input value to left by one bit. */
4328 hsa_op_reg
*arg_neg
= new hsa_op_reg (arg
->m_type
);
4329 gen_hsa_unary_operation (BRIG_OPCODE_NEG
, arg_neg
, arg
, hbb
);
4331 hsa_op_reg
*shifted_arg
= new hsa_op_reg (arg
->m_type
);
4332 gen_hsa_binary_operation (BRIG_OPCODE_SHL
, shifted_arg
, arg
,
4333 new hsa_op_immed (1, BRIG_TYPE_U64
), hbb
);
4335 /* Assign the value that can be used for FIRSTBIT instruction according
4336 to the most significant bit. */
4337 hsa_op_reg
*tmp
= new hsa_op_reg (bittype
);
4338 hsa_insn_basic
*cmov
4339 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, bittype
, tmp
, most_sign
,
4340 arg_neg
, shifted_arg
);
4341 hbb
->append_insn (cmov
);
4343 hsa_op_reg
*leading_bits
= new hsa_op_reg (BRIG_TYPE_S32
);
4344 gen_hsa_unary_operation (BRIG_OPCODE_FIRSTBIT
, leading_bits
,
4345 tmp
->get_in_type (hsa_uint_for_bitsize (bitsize
),
4348 /* Set flag if the input value is equal to zero. */
4349 hsa_op_reg
*is_zero
= new hsa_op_reg (BRIG_TYPE_B1
);
4350 cmp
= new hsa_insn_cmp (BRIG_COMPARE_EQ
, is_zero
->m_type
, is_zero
, arg
,
4351 new hsa_op_immed (0, arg
->m_type
));
4352 hbb
->append_insn (cmp
);
4354 /* Return the number of leading bits,
4355 or (bitsize - 1) if the input value is zero. */
4356 cmov
= new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, BRIG_TYPE_B32
, NULL
, is_zero
,
4357 new hsa_op_immed (bitsize
- 1, BRIG_TYPE_U32
),
4358 leading_bits
->get_in_type (BRIG_TYPE_B32
, hbb
));
4359 hbb
->append_insn (cmov
);
4360 cmov
->set_output_in_type (dest
, 0, hbb
);
4363 /* Emit instructions that implement ffs builtin STMT:
4364 Returns one plus the index of the least significant 1-bit of x,
4365 or if x is zero, returns zero.
4366 Instructions are appended to basic block HBB. */
4369 gen_hsa_ffs (gcall
*call
, hsa_bb
*hbb
)
4371 tree lhs
= gimple_call_lhs (call
);
4372 if (lhs
== NULL_TREE
)
4375 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4377 tree rhs1
= gimple_call_arg (call
, 0);
4378 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4379 arg
= arg
->extend_int_to_32bit (hbb
);
4381 hsa_op_reg
*tmp
= new hsa_op_reg (BRIG_TYPE_U32
);
4382 hsa_insn_srctype
*insn
= new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT
,
4383 tmp
->m_type
, arg
->m_type
,
4385 hbb
->append_insn (insn
);
4387 hsa_insn_basic
*addition
4388 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
, NULL
, tmp
,
4389 new hsa_op_immed (1, tmp
->m_type
));
4390 hbb
->append_insn (addition
);
4391 addition
->set_output_in_type (dest
, 0, hbb
);
4395 gen_hsa_popcount_to_dest (hsa_op_reg
*dest
, hsa_op_with_type
*arg
, hsa_bb
*hbb
)
4397 gcc_checking_assert (hsa_type_integer_p (arg
->m_type
));
4399 if (hsa_type_bit_size (arg
->m_type
) < 32)
4400 arg
= arg
->get_in_type (BRIG_TYPE_B32
, hbb
);
4402 BrigType16_t srctype
= hsa_bittype_for_type (arg
->m_type
);
4403 if (!hsa_btype_p (arg
->m_type
))
4404 arg
= arg
->get_in_type (srctype
, hbb
);
4406 hsa_insn_srctype
*popcount
4407 = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT
, BRIG_TYPE_U32
,
4408 srctype
, NULL
, arg
);
4409 hbb
->append_insn (popcount
);
4410 popcount
->set_output_in_type (dest
, 0, hbb
);
4413 /* Emit instructions that implement parity builtin STMT:
4414 Returns the parity of x, i.e. the number of 1-bits in x modulo 2.
4415 Instructions are appended to basic block HBB. */
4418 gen_hsa_parity (gcall
*call
, hsa_bb
*hbb
)
4420 tree lhs
= gimple_call_lhs (call
);
4421 if (lhs
== NULL_TREE
)
4424 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4425 tree rhs1
= gimple_call_arg (call
, 0);
4426 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4428 hsa_op_reg
*popcount
= new hsa_op_reg (BRIG_TYPE_U32
);
4429 gen_hsa_popcount_to_dest (popcount
, arg
, hbb
);
4431 hsa_insn_basic
*insn
4432 = new hsa_insn_basic (3, BRIG_OPCODE_REM
, popcount
->m_type
, NULL
, popcount
,
4433 new hsa_op_immed (2, popcount
->m_type
));
4434 hbb
->append_insn (insn
);
4435 insn
->set_output_in_type (dest
, 0, hbb
);
4438 /* Emit instructions that implement popcount builtin STMT.
4439 Instructions are appended to basic block HBB. */
4442 gen_hsa_popcount (gcall
*call
, hsa_bb
*hbb
)
4444 tree lhs
= gimple_call_lhs (call
);
4445 if (lhs
== NULL_TREE
)
4448 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4449 tree rhs1
= gimple_call_arg (call
, 0);
4450 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4452 gen_hsa_popcount_to_dest (dest
, arg
, hbb
);
4455 /* Emit instructions that implement DIVMOD builtin STMT.
4456 Instructions are appended to basic block HBB. */
4459 gen_hsa_divmod (gcall
*call
, hsa_bb
*hbb
)
4461 tree lhs
= gimple_call_lhs (call
);
4462 if (lhs
== NULL_TREE
)
4465 tree rhs0
= gimple_call_arg (call
, 0);
4466 tree rhs1
= gimple_call_arg (call
, 1);
4468 hsa_op_with_type
*arg0
= hsa_reg_or_immed_for_gimple_op (rhs0
, hbb
);
4469 arg0
= arg0
->extend_int_to_32bit (hbb
);
4470 hsa_op_with_type
*arg1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4471 arg1
= arg1
->extend_int_to_32bit (hbb
);
4473 hsa_op_reg
*dest0
= new hsa_op_reg (arg0
->m_type
);
4474 hsa_op_reg
*dest1
= new hsa_op_reg (arg1
->m_type
);
4476 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_DIV
, dest0
->m_type
,
4478 hbb
->append_insn (insn
);
4479 insn
= new hsa_insn_basic (3, BRIG_OPCODE_REM
, dest1
->m_type
, dest1
, arg0
,
4481 hbb
->append_insn (insn
);
4483 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4484 BrigType16_t dst_type
= hsa_extend_inttype_to_32bit (dest
->m_type
);
4485 BrigType16_t src_type
= hsa_bittype_for_type (dest0
->m_type
);
4487 insn
= new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dst_type
,
4488 src_type
, NULL
, dest0
, dest1
);
4489 hbb
->append_insn (insn
);
4490 insn
->set_output_in_type (dest
, 0, hbb
);
4493 /* Set VALUE to a shadow kernel debug argument and append a new instruction
4494 to HBB basic block. */
4497 set_debug_value (hsa_bb
*hbb
, hsa_op_with_type
*value
)
4499 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4500 if (shadow_reg_ptr
== NULL
)
4503 hsa_op_address
*addr
4504 = new hsa_op_address (shadow_reg_ptr
,
4505 get_hsa_kernel_dispatch_offset ("debug"));
4506 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, BRIG_TYPE_U64
, value
,
4508 hbb
->append_insn (mem
);
4512 omp_simple_builtin::generate (gimple
*stmt
, hsa_bb
*hbb
)
4516 if (m_warning_message
)
4517 HSA_SORRY_AT (gimple_location (stmt
), m_warning_message
);
4519 HSA_SORRY_ATV (gimple_location (stmt
),
4520 "Support for HSA does not implement calls to %s\n",
4523 else if (m_warning_message
!= NULL
)
4524 warning_at (gimple_location (stmt
), OPT_Whsa
, m_warning_message
);
4526 if (m_return_value
!= NULL
)
4528 tree lhs
= gimple_call_lhs (stmt
);
4532 hbb
->append_insn (new hsa_insn_comment (m_name
));
4534 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4535 hsa_op_with_type
*op
= m_return_value
->get_in_type (dest
->m_type
, hbb
);
4536 hsa_build_append_simple_mov (dest
, op
, hbb
);
4540 /* If STMT is a call of a known library function, generate code to perform
4541 it and return true. */
4544 gen_hsa_insns_for_known_library_call (gimple
*stmt
, hsa_bb
*hbb
)
4546 bool handled
= false;
4547 const char *name
= hsa_get_declaration_name (gimple_call_fndecl (stmt
));
4550 size_t len
= strlen (name
);
4551 if (len
> 0 && name
[len
- 1] == '_')
4553 copy
= XNEWVEC (char, len
+ 1);
4554 strcpy (copy
, name
);
4555 copy
[len
- 1] = '\0';
4559 /* Handle omp_* routines. */
4560 if (strstr (name
, "omp_") == name
)
4562 hsa_init_simple_builtins ();
4563 omp_simple_builtin
*builtin
= omp_simple_builtins
->get (name
);
4566 builtin
->generate (stmt
, hbb
);
4571 if (strcmp (name
, "omp_set_num_threads") == 0)
4572 gen_set_num_threads (gimple_call_arg (stmt
, 0), hbb
);
4573 else if (strcmp (name
, "omp_get_thread_num") == 0)
4575 hbb
->append_insn (new hsa_insn_comment (name
));
4576 query_hsa_grid_nodim (stmt
, BRIG_OPCODE_WORKITEMFLATABSID
, hbb
);
4578 else if (strcmp (name
, "omp_get_num_threads") == 0)
4580 hbb
->append_insn (new hsa_insn_comment (name
));
4581 gen_get_num_threads (stmt
, hbb
);
4583 else if (strcmp (name
, "omp_get_num_teams") == 0)
4584 gen_get_num_teams (stmt
, hbb
);
4585 else if (strcmp (name
, "omp_get_team_num") == 0)
4586 gen_get_team_num (stmt
, hbb
);
4587 else if (strcmp (name
, "omp_get_level") == 0)
4588 gen_get_level (stmt
, hbb
);
4589 else if (strcmp (name
, "omp_get_active_level") == 0)
4590 gen_get_level (stmt
, hbb
);
4591 else if (strcmp (name
, "omp_in_parallel") == 0)
4592 gen_get_level (stmt
, hbb
);
4593 else if (strcmp (name
, "omp_get_max_threads") == 0)
4594 gen_get_max_threads (stmt
, hbb
);
4606 if (strcmp (name
, "__hsa_set_debug_value") == 0)
4609 if (hsa_cfun
->has_shadow_reg_p ())
4611 tree rhs1
= gimple_call_arg (stmt
, 0);
4612 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4614 src
= src
->get_in_type (BRIG_TYPE_U64
, hbb
);
4615 set_debug_value (hbb
, src
);
4624 /* Helper functions to create a single unary HSA operations out of calls to
4625 builtins. OPCODE is the HSA operation to be generated. STMT is a gimple
4626 call to a builtin. HBB is the HSA BB to which the instruction should be
4627 added. Note that nothing will be created if STMT does not have a LHS. */
4630 gen_hsa_unaryop_for_builtin (BrigOpcode opcode
, gimple
*stmt
, hsa_bb
*hbb
)
4632 tree lhs
= gimple_call_lhs (stmt
);
4635 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4636 hsa_op_with_type
*op
4637 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4638 gen_hsa_unary_operation (opcode
, dest
, op
, hbb
);
4641 /* Helper functions to create a call to standard library if LHS of the
4642 STMT is used. HBB is the HSA BB to which the instruction should be
4646 gen_hsa_unaryop_builtin_call (gimple
*stmt
, hsa_bb
*hbb
)
4648 tree lhs
= gimple_call_lhs (stmt
);
4652 if (gimple_call_internal_p (stmt
))
4653 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
4655 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4658 /* Helper functions to create a single unary HSA operations out of calls to
4659 builtins (if unsafe math optimizations are enable). Otherwise, create
4660 a call to standard library function.
4661 OPCODE is the HSA operation to be generated. STMT is a gimple
4662 call to a builtin. HBB is the HSA BB to which the instruction should be
4663 added. Note that nothing will be created if STMT does not have a LHS. */
4666 gen_hsa_unaryop_or_call_for_builtin (BrigOpcode opcode
, gimple
*stmt
,
4669 if (flag_unsafe_math_optimizations
)
4670 gen_hsa_unaryop_for_builtin (opcode
, stmt
, hbb
);
4672 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4675 /* Generate HSA address corresponding to a value VAL (as opposed to a memory
4676 reference tree), for example an SSA_NAME or an ADDR_EXPR. HBB is the HSA BB
4677 to which the instruction should be added. */
4679 static hsa_op_address
*
4680 get_address_from_value (tree val
, hsa_bb
*hbb
)
4682 switch (TREE_CODE (val
))
4686 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4688 = hsa_cfun
->reg_for_gimple_ssa (val
)->get_in_type (addrtype
, hbb
);
4689 return new hsa_op_address (NULL
, as_a
<hsa_op_reg
*> (reg
), 0);
4692 return gen_hsa_addr (TREE_OPERAND (val
, 0), hbb
);
4695 if (tree_fits_shwi_p (val
))
4696 return new hsa_op_address (NULL
, NULL
, tree_to_shwi (val
));
4700 HSA_SORRY_ATV (EXPR_LOCATION (val
),
4701 "support for HSA does not implement memory access to %E",
4703 return new hsa_op_address (NULL
, NULL
, 0);
4707 /* Expand assignment of a result of a string BUILTIN to DST.
4708 Size of the operation is N bytes, where instructions
4709 will be append to HBB. */
4712 expand_lhs_of_string_op (gimple
*stmt
,
4713 unsigned HOST_WIDE_INT n
, hsa_bb
*hbb
,
4714 enum built_in_function builtin
)
4716 /* If LHS is expected, we need to emit a PHI instruction. */
4717 tree lhs
= gimple_call_lhs (stmt
);
4721 hsa_op_reg
*lhs_reg
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4723 hsa_op_with_type
*dst_reg
4724 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4725 hsa_op_with_type
*tmp
;
4729 case BUILT_IN_MEMPCPY
:
4731 tmp
= new hsa_op_reg (dst_reg
->m_type
);
4733 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
,
4735 new hsa_op_immed (n
, dst_reg
->m_type
));
4736 hbb
->append_insn (add
);
4739 case BUILT_IN_MEMCPY
:
4740 case BUILT_IN_MEMSET
:
4747 hbb
->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV
, lhs_reg
->m_type
,
4751 #define HSA_MEMORY_BUILTINS_LIMIT 128
4753 /* Expand a string builtin (from a gimple STMT) in a way that
4754 according to MISALIGNED_FLAG we process either direct emission
4755 (a bunch of memory load and store instructions), or we emit a function call
4756 of a library function (for instance 'memcpy'). Actually, a basic block
4757 for direct emission is just prepared, where caller is responsible
4758 for emission of corresponding instructions.
4759 All instruction are appended to HBB. */
4762 expand_string_operation_builtin (gimple
*stmt
, hsa_bb
*hbb
,
4763 hsa_op_reg
*misaligned_flag
)
4765 edge e
= split_block (hbb
->m_bb
, stmt
);
4766 basic_block condition_bb
= e
->src
;
4767 hbb
->append_insn (new hsa_insn_cbr (misaligned_flag
));
4769 /* Prepare the control flow. */
4770 edge condition_edge
= EDGE_SUCC (condition_bb
, 0);
4771 basic_block call_bb
= split_edge (condition_edge
);
4773 basic_block expanded_bb
= split_edge (EDGE_SUCC (call_bb
, 0));
4774 basic_block cont_bb
= EDGE_SUCC (expanded_bb
, 0)->dest
;
4775 basic_block merge_bb
= split_edge (EDGE_PRED (cont_bb
, 0));
4777 condition_edge
->flags
&= ~EDGE_FALLTHRU
;
4778 condition_edge
->flags
|= EDGE_TRUE_VALUE
;
4779 make_edge (condition_bb
, expanded_bb
, EDGE_FALSE_VALUE
);
4781 redirect_edge_succ (EDGE_SUCC (call_bb
, 0), merge_bb
);
4783 hsa_cfun
->m_modified_cfg
= true;
4785 hsa_init_new_bb (expanded_bb
);
4787 /* Slow path: function call. */
4788 gen_hsa_insns_for_direct_call (stmt
, hsa_init_new_bb (call_bb
), false);
4790 return hsa_bb_for_bb (expanded_bb
);
4793 /* Expand a memory copy BUILTIN (BUILT_IN_MEMCPY, BUILT_IN_MEMPCPY) from
4794 a gimple STMT and store all necessary instruction to HBB basic block. */
4797 expand_memory_copy (gimple
*stmt
, hsa_bb
*hbb
, enum built_in_function builtin
)
4799 tree byte_size
= gimple_call_arg (stmt
, 2);
4801 if (!tree_fits_uhwi_p (byte_size
))
4803 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4807 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
4809 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
4811 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4815 tree dst
= gimple_call_arg (stmt
, 0);
4816 tree src
= gimple_call_arg (stmt
, 1);
4818 hsa_op_address
*dst_addr
= get_address_from_value (dst
, hbb
);
4819 hsa_op_address
*src_addr
= get_address_from_value (src
, hbb
);
4821 /* As gen_hsa_memory_copy relies on memory alignment
4822 greater or equal to 8 bytes, we need to verify the alignment. */
4823 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4824 hsa_op_reg
*src_addr_reg
= new hsa_op_reg (addrtype
);
4825 hsa_op_reg
*dst_addr_reg
= new hsa_op_reg (addrtype
);
4827 convert_addr_to_flat_segment (src_addr
, src_addr_reg
, hbb
);
4828 convert_addr_to_flat_segment (dst_addr
, dst_addr_reg
, hbb
);
4830 /* Process BIT OR for source and destination addresses. */
4831 hsa_op_reg
*or_reg
= new hsa_op_reg (addrtype
);
4832 gen_hsa_binary_operation (BRIG_OPCODE_OR
, or_reg
, src_addr_reg
,
4835 /* Process BIT AND with 0x7 to identify the desired alignment
4837 hsa_op_reg
*masked
= new hsa_op_reg (addrtype
);
4839 gen_hsa_binary_operation (BRIG_OPCODE_AND
, masked
, or_reg
,
4840 new hsa_op_immed (7, addrtype
), hbb
);
4842 hsa_op_reg
*misaligned
= new hsa_op_reg (BRIG_TYPE_B1
);
4843 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE
, misaligned
->m_type
,
4845 new hsa_op_immed (0, masked
->m_type
)));
4847 hsa_bb
*native_impl_bb
4848 = expand_string_operation_builtin (stmt
, hbb
, misaligned
);
4850 gen_hsa_memory_copy (native_impl_bb
, dst_addr
, src_addr
, n
, BRIG_ALIGNMENT_8
);
4852 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb
->m_bb
, 0)->dest
);
4853 expand_lhs_of_string_op (stmt
, n
, merge_bb
, builtin
);
4857 /* Expand a memory set BUILTIN (BUILT_IN_MEMSET, BUILT_IN_BZERO) from
4858 a gimple STMT and store all necessary instruction to HBB basic block.
4859 The operation set N bytes with a CONSTANT value. */
4862 expand_memory_set (gimple
*stmt
, unsigned HOST_WIDE_INT n
,
4863 unsigned HOST_WIDE_INT constant
, hsa_bb
*hbb
,
4864 enum built_in_function builtin
)
4866 tree dst
= gimple_call_arg (stmt
, 0);
4867 hsa_op_address
*dst_addr
= get_address_from_value (dst
, hbb
);
4869 /* As gen_hsa_memory_set relies on memory alignment
4870 greater or equal to 8 bytes, we need to verify the alignment. */
4871 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4872 hsa_op_reg
*dst_addr_reg
= new hsa_op_reg (addrtype
);
4873 convert_addr_to_flat_segment (dst_addr
, dst_addr_reg
, hbb
);
4875 /* Process BIT AND with 0x7 to identify the desired alignment
4877 hsa_op_reg
*masked
= new hsa_op_reg (addrtype
);
4879 gen_hsa_binary_operation (BRIG_OPCODE_AND
, masked
, dst_addr_reg
,
4880 new hsa_op_immed (7, addrtype
), hbb
);
4882 hsa_op_reg
*misaligned
= new hsa_op_reg (BRIG_TYPE_B1
);
4883 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE
, misaligned
->m_type
,
4885 new hsa_op_immed (0, masked
->m_type
)));
4887 hsa_bb
*native_impl_bb
4888 = expand_string_operation_builtin (stmt
, hbb
, misaligned
);
4890 gen_hsa_memory_set (native_impl_bb
, dst_addr
, constant
, n
, BRIG_ALIGNMENT_8
);
4892 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb
->m_bb
, 0)->dest
);
4893 expand_lhs_of_string_op (stmt
, n
, merge_bb
, builtin
);
4896 /* Store into MEMORDER the memory order specified by tree T, which must be an
4897 integer constant representing a C++ memory order. If it isn't, issue an HSA
4898 sorry message using LOC and return true, otherwise return false and store
4899 the name of the requested order to *MNAME. */
4902 hsa_memorder_from_tree (tree t
, BrigMemoryOrder
*memorder
, const char **mname
,
4905 if (!tree_fits_uhwi_p (t
))
4907 HSA_SORRY_ATV (loc
, "support for HSA does not implement memory model %E",
4912 unsigned HOST_WIDE_INT mm
= tree_to_uhwi (t
);
4913 switch (mm
& MEMMODEL_BASE_MASK
)
4915 case MEMMODEL_RELAXED
:
4916 *memorder
= BRIG_MEMORY_ORDER_RELAXED
;
4919 case MEMMODEL_CONSUME
:
4920 /* HSA does not have an equivalent, but we can use the slightly stronger
4922 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
4925 case MEMMODEL_ACQUIRE
:
4926 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
4929 case MEMMODEL_RELEASE
:
4930 *memorder
= BRIG_MEMORY_ORDER_SC_RELEASE
;
4933 case MEMMODEL_ACQ_REL
:
4934 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
4937 case MEMMODEL_SEQ_CST
:
4938 /* Callers implementing a simple load or store need to remove the release
4939 or acquire part respectively. */
4940 *memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
4945 HSA_SORRY_AT (loc
, "support for HSA does not implement the specified "
4953 /* Helper function to create an HSA atomic operation instruction out of calls
4954 to atomic builtins. RET_ORIG is true if the built-in is the variant that
4955 return s the value before applying operation, and false if it should return
4956 the value after applying the operation (if it returns value at all). ACODE
4957 is the atomic operation code, STMT is a gimple call to a builtin. HBB is
4958 the HSA BB to which the instruction should be added. If SIGNAL is true, the
4959 created operation will work on HSA signals rather than atomic variables. */
4962 gen_hsa_atomic_for_builtin (bool ret_orig
, enum BrigAtomicOperation acode
,
4963 gimple
*stmt
, hsa_bb
*hbb
, bool signal
)
4965 tree lhs
= gimple_call_lhs (stmt
);
4967 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
4968 BrigType16_t hsa_type
= hsa_type_for_scalar_tree_type (type
, false);
4969 BrigType16_t mtype
= mem_type_for_type (hsa_type
);
4970 BrigMemoryOrder memorder
;
4973 if (hsa_memorder_from_tree (gimple_call_arg (stmt
, 2), &memorder
, &mmname
,
4974 gimple_location (stmt
)))
4977 /* Certain atomic insns must have Bx memory types. */
4980 case BRIG_ATOMIC_LD
:
4981 case BRIG_ATOMIC_ST
:
4982 case BRIG_ATOMIC_AND
:
4983 case BRIG_ATOMIC_OR
:
4984 case BRIG_ATOMIC_XOR
:
4985 case BRIG_ATOMIC_EXCH
:
4986 mtype
= hsa_bittype_for_type (mtype
);
4997 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4999 dest
= new hsa_op_reg (hsa_type
);
5000 opcode
= signal
? BRIG_OPCODE_SIGNAL
: BRIG_OPCODE_ATOMIC
;
5006 opcode
= signal
? BRIG_OPCODE_SIGNALNORET
: BRIG_OPCODE_ATOMICNORET
;
5010 if (acode
== BRIG_ATOMIC_ST
)
5012 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
5013 memorder
= BRIG_MEMORY_ORDER_SC_RELEASE
;
5015 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
5016 && memorder
!= BRIG_MEMORY_ORDER_SC_RELEASE
5017 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
5019 HSA_SORRY_ATV (gimple_location (stmt
),
5020 "support for HSA does not implement memory model for "
5021 "ATOMIC_ST: %s", mmname
);
5026 hsa_insn_basic
*atominsn
;
5030 atominsn
= new hsa_insn_signal (nops
, opcode
, acode
, mtype
, memorder
);
5031 tgt
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
5035 atominsn
= new hsa_insn_atomic (nops
, opcode
, acode
, mtype
, memorder
);
5036 hsa_op_address
*addr
;
5037 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5038 if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
== BRIG_SEGMENT_PRIVATE
)
5040 HSA_SORRY_AT (gimple_location (stmt
),
5041 "HSA does not implement atomic operations in private "
5048 hsa_op_with_type
*op
5049 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1), hbb
);
5052 atominsn
->set_op (0, dest
);
5053 atominsn
->set_op (1, tgt
);
5054 atominsn
->set_op (2, op
);
5058 atominsn
->set_op (0, tgt
);
5059 atominsn
->set_op (1, op
);
5062 hbb
->append_insn (atominsn
);
5064 /* HSA does not natively support the variants that return the modified value,
5065 so re-do the operation again non-atomically if that is what was
5067 if (lhs
&& !ret_orig
)
5072 case BRIG_ATOMIC_ADD
:
5073 arith
= BRIG_OPCODE_ADD
;
5075 case BRIG_ATOMIC_AND
:
5076 arith
= BRIG_OPCODE_AND
;
5078 case BRIG_ATOMIC_OR
:
5079 arith
= BRIG_OPCODE_OR
;
5081 case BRIG_ATOMIC_SUB
:
5082 arith
= BRIG_OPCODE_SUB
;
5084 case BRIG_ATOMIC_XOR
:
5085 arith
= BRIG_OPCODE_XOR
;
5090 hsa_op_reg
*real_dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5091 gen_hsa_binary_operation (arith
, real_dest
, dest
, op
, hbb
);
5095 /* Generate HSA instructions for an internal fn.
5096 Instructions will be appended to HBB, which also needs to be the
5097 corresponding structure to the basic_block of STMT. */
5100 gen_hsa_insn_for_internal_fn_call (gcall
*stmt
, hsa_bb
*hbb
)
5102 gcc_checking_assert (gimple_call_internal_fn (stmt
));
5103 internal_fn fn
= gimple_call_internal_fn (stmt
);
5105 bool is_float_type_p
= false;
5106 if (gimple_call_lhs (stmt
) != NULL
5107 && TREE_TYPE (gimple_call_lhs (stmt
)) == float_type_node
)
5108 is_float_type_p
= true;
5113 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
5117 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
5121 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
5125 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
5129 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_NRSQRT
, stmt
, hbb
);
5133 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
5138 if (is_float_type_p
)
5139 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
5141 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5147 if (is_float_type_p
)
5148 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
5150 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5157 if (is_float_type_p
)
5158 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
5160 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5167 if (is_float_type_p
)
5168 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
5170 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5175 gen_hsa_clrsb (stmt
, hbb
);
5179 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
5183 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
5187 gen_hsa_ffs (stmt
, hbb
);
5191 gen_hsa_parity (stmt
, hbb
);
5195 gen_hsa_popcount (stmt
, hbb
);
5199 gen_hsa_divmod (stmt
, hbb
);
5212 case IFN_SIGNIFICAND
:
5224 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
5228 HSA_SORRY_ATV (gimple_location (stmt
),
5229 "support for HSA does not implement internal function: %s",
5230 internal_fn_name (fn
));
5235 /* Generate HSA instructions for the given call statement STMT. Instructions
5236 will be appended to HBB. */
5239 gen_hsa_insns_for_call (gimple
*stmt
, hsa_bb
*hbb
)
5241 gcall
*call
= as_a
<gcall
*> (stmt
);
5242 tree lhs
= gimple_call_lhs (stmt
);
5245 if (gimple_call_internal_p (stmt
))
5247 gen_hsa_insn_for_internal_fn_call (call
, hbb
);
5251 if (!gimple_call_builtin_p (stmt
, BUILT_IN_NORMAL
))
5253 tree function_decl
= gimple_call_fndecl (stmt
);
5254 /* Prefetch pass can create type-mismatching prefetch builtin calls which
5255 fail the gimple_call_builtin_p test above. Handle them here. */
5256 if (DECL_BUILT_IN_CLASS (function_decl
)
5257 && DECL_FUNCTION_CODE (function_decl
) == BUILT_IN_PREFETCH
)
5260 if (function_decl
== NULL_TREE
)
5262 HSA_SORRY_AT (gimple_location (stmt
),
5263 "support for HSA does not implement indirect calls");
5267 if (hsa_callable_function_p (function_decl
))
5268 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5269 else if (!gen_hsa_insns_for_known_library_call (stmt
, hbb
))
5270 HSA_SORRY_AT (gimple_location (stmt
),
5271 "HSA supports only calls of functions marked with pragma "
5272 "omp declare target");
5276 tree fndecl
= gimple_call_fndecl (stmt
);
5277 enum built_in_function builtin
= DECL_FUNCTION_CODE (fndecl
);
5281 case BUILT_IN_FABSF
:
5282 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_ABS
, stmt
, hbb
);
5286 case BUILT_IN_CEILF
:
5287 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
5290 case BUILT_IN_FLOOR
:
5291 case BUILT_IN_FLOORF
:
5292 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
5296 case BUILT_IN_RINTF
:
5297 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
5301 case BUILT_IN_SQRTF
:
5302 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
5305 case BUILT_IN_TRUNC
:
5306 case BUILT_IN_TRUNCF
:
5307 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
5314 /* HSAIL does not provide an instruction for double argument type. */
5315 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5319 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
5322 case BUILT_IN_EXP2F
:
5323 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
5326 case BUILT_IN_LOG2F
:
5327 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
5331 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
5334 case BUILT_IN_CLRSB
:
5335 case BUILT_IN_CLRSBL
:
5336 case BUILT_IN_CLRSBLL
:
5337 gen_hsa_clrsb (call
, hbb
);
5342 case BUILT_IN_CLZLL
:
5343 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
5348 case BUILT_IN_CTZLL
:
5349 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
5354 case BUILT_IN_FFSLL
:
5355 gen_hsa_ffs (call
, hbb
);
5358 case BUILT_IN_PARITY
:
5359 case BUILT_IN_PARITYL
:
5360 case BUILT_IN_PARITYLL
:
5361 gen_hsa_parity (call
, hbb
);
5364 case BUILT_IN_POPCOUNT
:
5365 case BUILT_IN_POPCOUNTL
:
5366 case BUILT_IN_POPCOUNTLL
:
5367 gen_hsa_popcount (call
, hbb
);
5370 case BUILT_IN_ATOMIC_LOAD_1
:
5371 case BUILT_IN_ATOMIC_LOAD_2
:
5372 case BUILT_IN_ATOMIC_LOAD_4
:
5373 case BUILT_IN_ATOMIC_LOAD_8
:
5374 case BUILT_IN_ATOMIC_LOAD_16
:
5378 src
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5380 BrigMemoryOrder memorder
;
5382 if (hsa_memorder_from_tree (gimple_call_arg (stmt
, 1), &memorder
,
5383 &mmname
, gimple_location (stmt
)))
5386 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
5387 memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
5389 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
5390 && memorder
!= BRIG_MEMORY_ORDER_SC_ACQUIRE
5391 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
5393 HSA_SORRY_ATV (gimple_location (stmt
),
5394 "support for HSA does not implement "
5395 "memory model for atomic loads: %s", mmname
);
5401 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
5403 mtype
= mem_type_for_type (t
);
5404 mtype
= hsa_bittype_for_type (mtype
);
5405 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5409 mtype
= BRIG_TYPE_B64
;
5410 dest
= new hsa_op_reg (mtype
);
5413 hsa_insn_basic
*atominsn
;
5414 atominsn
= new hsa_insn_atomic (2, BRIG_OPCODE_ATOMIC
, BRIG_ATOMIC_LD
,
5415 mtype
, memorder
, dest
, src
);
5417 hbb
->append_insn (atominsn
);
5421 case BUILT_IN_ATOMIC_EXCHANGE_1
:
5422 case BUILT_IN_ATOMIC_EXCHANGE_2
:
5423 case BUILT_IN_ATOMIC_EXCHANGE_4
:
5424 case BUILT_IN_ATOMIC_EXCHANGE_8
:
5425 case BUILT_IN_ATOMIC_EXCHANGE_16
:
5426 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_EXCH
, stmt
, hbb
, false);
5430 case BUILT_IN_ATOMIC_FETCH_ADD_1
:
5431 case BUILT_IN_ATOMIC_FETCH_ADD_2
:
5432 case BUILT_IN_ATOMIC_FETCH_ADD_4
:
5433 case BUILT_IN_ATOMIC_FETCH_ADD_8
:
5434 case BUILT_IN_ATOMIC_FETCH_ADD_16
:
5435 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ADD
, stmt
, hbb
, false);
5439 case BUILT_IN_ATOMIC_FETCH_SUB_1
:
5440 case BUILT_IN_ATOMIC_FETCH_SUB_2
:
5441 case BUILT_IN_ATOMIC_FETCH_SUB_4
:
5442 case BUILT_IN_ATOMIC_FETCH_SUB_8
:
5443 case BUILT_IN_ATOMIC_FETCH_SUB_16
:
5444 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_SUB
, stmt
, hbb
, false);
5448 case BUILT_IN_ATOMIC_FETCH_AND_1
:
5449 case BUILT_IN_ATOMIC_FETCH_AND_2
:
5450 case BUILT_IN_ATOMIC_FETCH_AND_4
:
5451 case BUILT_IN_ATOMIC_FETCH_AND_8
:
5452 case BUILT_IN_ATOMIC_FETCH_AND_16
:
5453 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_AND
, stmt
, hbb
, false);
5457 case BUILT_IN_ATOMIC_FETCH_XOR_1
:
5458 case BUILT_IN_ATOMIC_FETCH_XOR_2
:
5459 case BUILT_IN_ATOMIC_FETCH_XOR_4
:
5460 case BUILT_IN_ATOMIC_FETCH_XOR_8
:
5461 case BUILT_IN_ATOMIC_FETCH_XOR_16
:
5462 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_XOR
, stmt
, hbb
, false);
5466 case BUILT_IN_ATOMIC_FETCH_OR_1
:
5467 case BUILT_IN_ATOMIC_FETCH_OR_2
:
5468 case BUILT_IN_ATOMIC_FETCH_OR_4
:
5469 case BUILT_IN_ATOMIC_FETCH_OR_8
:
5470 case BUILT_IN_ATOMIC_FETCH_OR_16
:
5471 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_OR
, stmt
, hbb
, false);
5475 case BUILT_IN_ATOMIC_STORE_1
:
5476 case BUILT_IN_ATOMIC_STORE_2
:
5477 case BUILT_IN_ATOMIC_STORE_4
:
5478 case BUILT_IN_ATOMIC_STORE_8
:
5479 case BUILT_IN_ATOMIC_STORE_16
:
5480 /* Since there cannot be any LHS, the first parameter is meaningless. */
5481 gen_hsa_atomic_for_builtin (true, BRIG_ATOMIC_ST
, stmt
, hbb
, false);
5485 case BUILT_IN_ATOMIC_ADD_FETCH_1
:
5486 case BUILT_IN_ATOMIC_ADD_FETCH_2
:
5487 case BUILT_IN_ATOMIC_ADD_FETCH_4
:
5488 case BUILT_IN_ATOMIC_ADD_FETCH_8
:
5489 case BUILT_IN_ATOMIC_ADD_FETCH_16
:
5490 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_ADD
, stmt
, hbb
, false);
5493 case BUILT_IN_ATOMIC_SUB_FETCH_1
:
5494 case BUILT_IN_ATOMIC_SUB_FETCH_2
:
5495 case BUILT_IN_ATOMIC_SUB_FETCH_4
:
5496 case BUILT_IN_ATOMIC_SUB_FETCH_8
:
5497 case BUILT_IN_ATOMIC_SUB_FETCH_16
:
5498 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_SUB
, stmt
, hbb
, false);
5501 case BUILT_IN_ATOMIC_AND_FETCH_1
:
5502 case BUILT_IN_ATOMIC_AND_FETCH_2
:
5503 case BUILT_IN_ATOMIC_AND_FETCH_4
:
5504 case BUILT_IN_ATOMIC_AND_FETCH_8
:
5505 case BUILT_IN_ATOMIC_AND_FETCH_16
:
5506 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_AND
, stmt
, hbb
, false);
5509 case BUILT_IN_ATOMIC_XOR_FETCH_1
:
5510 case BUILT_IN_ATOMIC_XOR_FETCH_2
:
5511 case BUILT_IN_ATOMIC_XOR_FETCH_4
:
5512 case BUILT_IN_ATOMIC_XOR_FETCH_8
:
5513 case BUILT_IN_ATOMIC_XOR_FETCH_16
:
5514 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_XOR
, stmt
, hbb
, false);
5517 case BUILT_IN_ATOMIC_OR_FETCH_1
:
5518 case BUILT_IN_ATOMIC_OR_FETCH_2
:
5519 case BUILT_IN_ATOMIC_OR_FETCH_4
:
5520 case BUILT_IN_ATOMIC_OR_FETCH_8
:
5521 case BUILT_IN_ATOMIC_OR_FETCH_16
:
5522 gen_hsa_atomic_for_builtin (false, BRIG_ATOMIC_OR
, stmt
, hbb
, false);
5525 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_1
:
5526 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_2
:
5527 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_4
:
5528 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8
:
5529 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_16
:
5531 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
5533 = hsa_bittype_for_type (hsa_type_for_scalar_tree_type (type
, false));
5534 BrigMemoryOrder memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
5535 hsa_insn_basic
*atominsn
;
5537 atominsn
= new hsa_insn_atomic (4, BRIG_OPCODE_ATOMIC
,
5538 BRIG_ATOMIC_CAS
, atype
, memorder
);
5539 tgt
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5542 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5544 dest
= new hsa_op_reg (atype
);
5546 atominsn
->set_op (0, dest
);
5547 atominsn
->set_op (1, tgt
);
5549 hsa_op_with_type
*op
5550 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1), hbb
);
5551 atominsn
->set_op (2, op
);
5552 op
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 2), hbb
);
5553 atominsn
->set_op (3, op
);
5555 hbb
->append_insn (atominsn
);
5559 case BUILT_IN_HSA_WORKGROUPID
:
5560 query_hsa_grid_dim (stmt
, BRIG_OPCODE_WORKGROUPID
, hbb
);
5562 case BUILT_IN_HSA_WORKITEMID
:
5563 query_hsa_grid_dim (stmt
, BRIG_OPCODE_WORKITEMID
, hbb
);
5565 case BUILT_IN_HSA_WORKITEMABSID
:
5566 query_hsa_grid_dim (stmt
, BRIG_OPCODE_WORKITEMABSID
, hbb
);
5568 case BUILT_IN_HSA_GRIDSIZE
:
5569 query_hsa_grid_dim (stmt
, BRIG_OPCODE_GRIDSIZE
, hbb
);
5571 case BUILT_IN_HSA_CURRENTWORKGROUPSIZE
:
5572 query_hsa_grid_dim (stmt
, BRIG_OPCODE_CURRENTWORKGROUPSIZE
, hbb
);
5575 case BUILT_IN_GOMP_BARRIER
:
5576 hbb
->append_insn (new hsa_insn_br (0, BRIG_OPCODE_BARRIER
, BRIG_TYPE_NONE
,
5579 case BUILT_IN_GOMP_PARALLEL
:
5580 HSA_SORRY_AT (gimple_location (stmt
),
5581 "support for HSA does not implement non-gridified "
5582 "OpenMP parallel constructs.");
5585 case BUILT_IN_OMP_GET_THREAD_NUM
:
5587 query_hsa_grid_nodim (stmt
, BRIG_OPCODE_WORKITEMFLATABSID
, hbb
);
5591 case BUILT_IN_OMP_GET_NUM_THREADS
:
5593 gen_get_num_threads (stmt
, hbb
);
5596 case BUILT_IN_GOMP_TEAMS
:
5598 gen_set_num_threads (gimple_call_arg (stmt
, 1), hbb
);
5601 case BUILT_IN_OMP_GET_NUM_TEAMS
:
5603 gen_get_num_teams (stmt
, hbb
);
5606 case BUILT_IN_OMP_GET_TEAM_NUM
:
5608 gen_get_team_num (stmt
, hbb
);
5611 case BUILT_IN_MEMCPY
:
5612 case BUILT_IN_MEMPCPY
:
5614 expand_memory_copy (stmt
, hbb
, builtin
);
5617 case BUILT_IN_MEMSET
:
5619 tree c
= gimple_call_arg (stmt
, 1);
5621 if (TREE_CODE (c
) != INTEGER_CST
)
5623 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5627 tree byte_size
= gimple_call_arg (stmt
, 2);
5629 if (!tree_fits_uhwi_p (byte_size
))
5631 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5635 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
5637 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5639 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5643 unsigned HOST_WIDE_INT constant
5644 = tree_to_uhwi (fold_convert (unsigned_char_type_node
, c
));
5646 expand_memory_set (stmt
, n
, constant
, hbb
, builtin
);
5650 case BUILT_IN_BZERO
:
5652 tree byte_size
= gimple_call_arg (stmt
, 1);
5654 if (!tree_fits_uhwi_p (byte_size
))
5656 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5660 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
5662 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5664 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5668 expand_memory_set (stmt
, n
, 0, hbb
, builtin
);
5672 CASE_BUILT_IN_ALLOCA
:
5674 gen_hsa_alloca (call
, hbb
);
5677 case BUILT_IN_PREFETCH
:
5681 tree name_tree
= DECL_NAME (fndecl
);
5682 const char *s
= IDENTIFIER_POINTER (name_tree
);
5683 size_t len
= strlen (s
);
5684 if (len
> 4 && (strncmp (s
, "__builtin_GOMP_", 15) == 0))
5685 HSA_SORRY_ATV (gimple_location (stmt
),
5686 "support for HSA does not implement GOMP function %s",
5689 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5695 /* Generate HSA instructions for a given gimple statement. Instructions will be
5699 gen_hsa_insns_for_gimple_stmt (gimple
*stmt
, hsa_bb
*hbb
)
5701 switch (gimple_code (stmt
))
5704 if (gimple_clobber_p (stmt
))
5707 if (gimple_assign_single_p (stmt
))
5709 tree lhs
= gimple_assign_lhs (stmt
);
5710 tree rhs
= gimple_assign_rhs1 (stmt
);
5711 gen_hsa_insns_for_single_assignment (lhs
, rhs
, hbb
);
5714 gen_hsa_insns_for_operation_assignment (stmt
, hbb
);
5717 gen_hsa_insns_for_return (as_a
<greturn
*> (stmt
), hbb
);
5720 gen_hsa_insns_for_cond_stmt (stmt
, hbb
);
5723 gen_hsa_insns_for_call (stmt
, hbb
);
5726 /* ??? HSA supports some debug facilities. */
5730 tree label
= gimple_label_label (as_a
<glabel
*> (stmt
));
5731 if (FORCED_LABEL (label
))
5732 HSA_SORRY_AT (gimple_location (stmt
),
5733 "support for HSA does not implement gimple label with "
5740 hbb
->append_insn (new hsa_insn_basic (0, BRIG_OPCODE_NOP
));
5745 gen_hsa_insns_for_switch_stmt (as_a
<gswitch
*> (stmt
), hbb
);
5749 HSA_SORRY_ATV (gimple_location (stmt
),
5750 "support for HSA does not implement gimple statement %s",
5751 gimple_code_name
[(int) gimple_code (stmt
)]);
5755 /* Generate a HSA PHI from a gimple PHI. */
5758 gen_hsa_phi_from_gimple_phi (gimple
*phi_stmt
, hsa_bb
*hbb
)
5761 unsigned count
= gimple_phi_num_args (phi_stmt
);
5764 = hsa_cfun
->reg_for_gimple_ssa (gimple_phi_result (phi_stmt
));
5765 hphi
= new hsa_insn_phi (count
, dest
);
5766 hphi
->m_bb
= hbb
->m_bb
;
5768 auto_vec
<tree
, 8> aexprs
;
5769 auto_vec
<hsa_op_reg
*, 8> aregs
;
5771 /* Calling split_edge when processing a PHI node messes up with the order of
5772 gimple phi node arguments (it moves the one associated with the edge to
5773 the end). We need to keep the order of edges and arguments of HSA phi
5774 node arguments consistent, so we do all required splitting as the first
5775 step, and in reverse order as to not be affected by the re-orderings. */
5776 for (unsigned j
= count
; j
!= 0; j
--)
5779 tree op
= gimple_phi_arg_def (phi_stmt
, i
);
5780 if (TREE_CODE (op
) != ADDR_EXPR
)
5783 edge e
= gimple_phi_arg_edge (as_a
<gphi
*> (phi_stmt
), i
);
5784 hsa_bb
*hbb_src
= hsa_init_new_bb (split_edge (e
));
5785 hsa_op_address
*addr
= gen_hsa_addr (TREE_OPERAND (op
, 0),
5789 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
5790 hsa_insn_basic
*insn
5791 = new hsa_insn_basic (2, BRIG_OPCODE_LDA
, BRIG_TYPE_U64
,
5793 hbb_src
->append_insn (insn
);
5794 aexprs
.safe_push (op
);
5795 aregs
.safe_push (dest
);
5798 tree lhs
= gimple_phi_result (phi_stmt
);
5799 for (unsigned i
= 0; i
< count
; i
++)
5801 tree op
= gimple_phi_arg_def (phi_stmt
, i
);
5803 if (TREE_CODE (op
) == SSA_NAME
)
5805 hsa_op_reg
*hreg
= hsa_cfun
->reg_for_gimple_ssa (op
);
5806 hphi
->set_op (i
, hreg
);
5810 gcc_assert (is_gimple_min_invariant (op
));
5811 tree t
= TREE_TYPE (op
);
5812 if (!POINTER_TYPE_P (t
)
5813 || (TREE_CODE (op
) == STRING_CST
5814 && TREE_CODE (TREE_TYPE (t
)) == INTEGER_TYPE
))
5815 hphi
->set_op (i
, new hsa_op_immed (op
));
5816 else if (POINTER_TYPE_P (TREE_TYPE (lhs
))
5817 && TREE_CODE (op
) == INTEGER_CST
)
5819 /* Handle assignment of NULL value to a pointer type. */
5820 hphi
->set_op (i
, new hsa_op_immed (op
));
5822 else if (TREE_CODE (op
) == ADDR_EXPR
)
5824 hsa_op_reg
*dest
= NULL
;
5825 for (unsigned a_idx
= 0; a_idx
< aexprs
.length (); a_idx
++)
5826 if (aexprs
[a_idx
] == op
)
5828 dest
= aregs
[a_idx
];
5832 hphi
->set_op (i
, dest
);
5836 HSA_SORRY_AT (gimple_location (phi_stmt
),
5837 "support for HSA does not handle PHI nodes with "
5838 "constant address operands");
5844 hbb
->append_phi (hphi
);
5847 /* Constructor of class containing HSA-specific information about a basic
5848 block. CFG_BB is the CFG BB this HSA BB is associated with. IDX is the new
5849 index of this BB (so that the constructor does not attempt to use
5850 hsa_cfun during its construction). */
5852 hsa_bb::hsa_bb (basic_block cfg_bb
, int idx
)
5853 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5854 m_last_phi (NULL
), m_index (idx
)
5856 gcc_assert (!cfg_bb
->aux
);
5860 /* Constructor of class containing HSA-specific information about a basic
5861 block. CFG_BB is the CFG BB this HSA BB is associated with. */
5863 hsa_bb::hsa_bb (basic_block cfg_bb
)
5864 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5865 m_last_phi (NULL
), m_index (hsa_cfun
->m_hbb_count
++)
5867 gcc_assert (!cfg_bb
->aux
);
5871 /* Create and initialize and return a new hsa_bb structure for a given CFG
5875 hsa_init_new_bb (basic_block bb
)
5877 void *m
= obstack_alloc (&hsa_obstack
, sizeof (hsa_bb
));
5878 return new (m
) hsa_bb (bb
);
5881 /* Initialize OMP in an HSA basic block PROLOGUE. */
5884 init_prologue (void)
5886 if (!hsa_cfun
->m_kern_p
)
5889 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5891 /* Create a magic number that is going to be printed by libgomp. */
5892 unsigned index
= hsa_get_number_decl_kernel_mappings ();
5894 /* Emit store to debug argument. */
5895 if (PARAM_VALUE (PARAM_HSA_GEN_DEBUG_STORES
) > 0)
5896 set_debug_value (prologue
, new hsa_op_immed (1000 + index
, BRIG_TYPE_U64
));
5899 /* Initialize hsa_num_threads to a default value. */
5902 init_hsa_num_threads (void)
5904 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5906 /* Save the default value to private variable hsa_num_threads. */
5907 hsa_insn_basic
*basic
5908 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
,
5909 new hsa_op_immed (0, hsa_num_threads
->m_type
),
5910 new hsa_op_address (hsa_num_threads
));
5911 prologue
->append_insn (basic
);
5914 /* Go over gimple representation and generate our internal HSA one. */
5917 gen_body_from_gimple ()
5921 /* Verify CFG for complex edges we are unable to handle. */
5925 FOR_EACH_BB_FN (bb
, cfun
)
5927 FOR_EACH_EDGE (e
, ei
, bb
->succs
)
5929 /* Verify all unsupported flags for edges that point
5930 to the same basic block. */
5931 if (e
->flags
& EDGE_EH
)
5933 HSA_SORRY_AT (UNKNOWN_LOCATION
,
5934 "support for HSA does not implement exception "
5941 FOR_EACH_BB_FN (bb
, cfun
)
5943 gimple_stmt_iterator gsi
;
5944 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5948 hbb
= hsa_init_new_bb (bb
);
5950 for (gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5952 gen_hsa_insns_for_gimple_stmt (gsi_stmt (gsi
), hbb
);
5953 if (hsa_seen_error ())
5958 FOR_EACH_BB_FN (bb
, cfun
)
5960 gimple_stmt_iterator gsi
;
5961 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5962 gcc_assert (hbb
!= NULL
);
5964 for (gsi
= gsi_start_phis (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5965 if (!virtual_operand_p (gimple_phi_result (gsi_stmt (gsi
))))
5966 gen_hsa_phi_from_gimple_phi (gsi_stmt (gsi
), hbb
);
5969 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
5971 fprintf (dump_file
, "------- Generated SSA form -------\n");
5972 dump_hsa_cfun (dump_file
);
5977 gen_function_decl_parameters (hsa_function_representation
*f
,
5983 for (parm
= TYPE_ARG_TYPES (TREE_TYPE (decl
)), i
= 0;
5985 parm
= TREE_CHAIN (parm
), i
++)
5987 /* Result type if last in the tree list. */
5988 if (TREE_CHAIN (parm
) == NULL
)
5991 tree v
= TREE_VALUE (parm
);
5993 hsa_symbol
*arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5995 arg
->m_type
= hsa_type_for_tree_type (v
, &arg
->m_dim
);
5996 arg
->m_name_number
= i
;
5998 f
->m_input_args
.safe_push (arg
);
6001 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
6002 if (!VOID_TYPE_P (result_type
))
6004 f
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
6006 f
->m_output_arg
->m_type
6007 = hsa_type_for_tree_type (result_type
, &f
->m_output_arg
->m_dim
);
6008 f
->m_output_arg
->m_name
= "res";
6012 /* Generate the vector of parameters of the HSA representation of the current
6013 function. This also includes the output parameter representing the
6017 gen_function_def_parameters ()
6021 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
6023 for (parm
= DECL_ARGUMENTS (cfun
->decl
); parm
;
6024 parm
= DECL_CHAIN (parm
))
6026 struct hsa_symbol
**slot
;
6029 = new hsa_symbol (BRIG_TYPE_NONE
, hsa_cfun
->m_kern_p
6030 ? BRIG_SEGMENT_KERNARG
: BRIG_SEGMENT_ARG
,
6031 BRIG_LINKAGE_FUNCTION
);
6032 arg
->fillup_for_decl (parm
);
6034 hsa_cfun
->m_input_args
.safe_push (arg
);
6036 if (hsa_seen_error ())
6039 arg
->m_name
= hsa_get_declaration_name (parm
);
6041 /* Copy all input arguments and create corresponding private symbols
6043 hsa_symbol
*private_arg
;
6044 hsa_op_address
*parm_addr
= new hsa_op_address (arg
);
6046 if (TREE_ADDRESSABLE (parm
)
6047 || (!is_gimple_reg (parm
) && !TREE_READONLY (parm
)))
6049 private_arg
= hsa_cfun
->create_hsa_temporary (arg
->m_type
);
6050 private_arg
->fillup_for_decl (parm
);
6052 BrigAlignment8_t align
= MIN (arg
->m_align
, private_arg
->m_align
);
6054 hsa_op_address
*private_arg_addr
= new hsa_op_address (private_arg
);
6055 gen_hsa_memory_copy (prologue
, private_arg_addr
, parm_addr
,
6056 arg
->total_byte_size (), align
);
6061 slot
= hsa_cfun
->m_local_symbols
->find_slot (private_arg
, INSERT
);
6062 gcc_assert (!*slot
);
6063 *slot
= private_arg
;
6065 if (is_gimple_reg (parm
))
6067 tree ddef
= ssa_default_def (cfun
, parm
);
6068 if (ddef
&& !has_zero_uses (ddef
))
6070 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (ddef
),
6072 BrigType16_t mtype
= mem_type_for_type (t
);
6073 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (ddef
);
6074 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
,
6076 gcc_assert (!parm_addr
->m_reg
);
6077 prologue
->append_insn (mem
);
6082 if (!VOID_TYPE_P (TREE_TYPE (TREE_TYPE (cfun
->decl
))))
6084 struct hsa_symbol
**slot
;
6086 hsa_cfun
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
6087 BRIG_LINKAGE_FUNCTION
);
6088 hsa_cfun
->m_output_arg
->fillup_for_decl (DECL_RESULT (cfun
->decl
));
6090 if (hsa_seen_error ())
6093 hsa_cfun
->m_output_arg
->m_name
= "res";
6094 slot
= hsa_cfun
->m_local_symbols
->find_slot (hsa_cfun
->m_output_arg
,
6096 gcc_assert (!*slot
);
6097 *slot
= hsa_cfun
->m_output_arg
;
6101 /* Generate function representation that corresponds to
6102 a function declaration. */
6104 hsa_function_representation
*
6105 hsa_generate_function_declaration (tree decl
)
6107 hsa_function_representation
*fun
6108 = new hsa_function_representation (decl
, false, 0);
6110 fun
->m_declaration_p
= true;
6111 fun
->m_name
= get_brig_function_name (decl
);
6112 gen_function_decl_parameters (fun
, decl
);
6118 /* Generate function representation that corresponds to
6121 hsa_function_representation
*
6122 hsa_generate_internal_fn_decl (hsa_internal_fn
*fn
)
6124 hsa_function_representation
*fun
= new hsa_function_representation (fn
);
6126 fun
->m_name
= fn
->name ();
6128 for (unsigned i
= 0; i
< fn
->get_arity (); i
++)
6131 = new hsa_symbol (fn
->get_argument_type (i
), BRIG_SEGMENT_ARG
,
6133 arg
->m_name_number
= i
;
6134 fun
->m_input_args
.safe_push (arg
);
6137 fun
->m_output_arg
= new hsa_symbol (fn
->get_argument_type (-1),
6138 BRIG_SEGMENT_ARG
, BRIG_LINKAGE_NONE
);
6139 fun
->m_output_arg
->m_name
= "res";
6144 /* Return true if switch statement S can be transformed
6145 to a SBR instruction in HSAIL. */
6148 transformable_switch_to_sbr_p (gswitch
*s
)
6150 /* Identify if a switch statement can be transformed to
6151 SBR instruction, like:
6153 sbr_u32 $s1 [@label1, @label2, @label3];
6156 tree size
= get_switch_size (s
);
6157 if (!tree_fits_uhwi_p (size
))
6160 if (tree_to_uhwi (size
) > HSA_MAXIMUM_SBR_LABELS
)
6166 /* Structure hold connection between PHI nodes and immediate
6167 values hold by there nodes. */
6169 struct phi_definition
6171 phi_definition (unsigned phi_i
, unsigned label_i
, tree imm
):
6172 phi_index (phi_i
), label_index (label_i
), phi_value (imm
)
6176 unsigned label_index
;
6180 /* Sum slice of a vector V, starting from index START and ending
6181 at the index END - 1. */
6183 template <typename T
>
6185 T
sum_slice (const auto_vec
<T
> &v
, unsigned start
, unsigned end
,
6190 for (unsigned i
= start
; i
< end
; i
++)
6196 /* Function transforms GIMPLE SWITCH statements to a series of IF statements.
6197 Let's assume following example:
6211 The transformation encompasses following steps:
6212 1) all immediate values used by edges coming from the switch basic block
6214 2) all these edges are removed
6215 3) the switch statement (in L0) is replaced by:
6221 4) newly created basic block Lx' is used for generation of
6223 5) else branch of the last condition goes to LD
6224 6) fix all immediate values in PHI nodes that were propagated though
6225 edges that were removed in step 2
6227 Note: if a case is made by a range C1..C2, then process
6228 following transformation:
6230 switch_cond_op1 = C1 <= index;
6231 switch_cond_op2 = index <= C2;
6232 switch_cond_and = switch_cond_op1 & switch_cond_op2;
6233 if (switch_cond_and != 0)
6241 convert_switch_statements (void)
6243 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
6246 bool modified_cfg
= false;
6248 FOR_EACH_BB_FN (bb
, func
)
6250 gimple_stmt_iterator gsi
= gsi_last_bb (bb
);
6251 if (gsi_end_p (gsi
))
6254 gimple
*stmt
= gsi_stmt (gsi
);
6256 if (gimple_code (stmt
) == GIMPLE_SWITCH
)
6258 gswitch
*s
= as_a
<gswitch
*> (stmt
);
6260 /* If the switch can utilize SBR insn, skip the statement. */
6261 if (transformable_switch_to_sbr_p (s
))
6264 modified_cfg
= true;
6266 unsigned labels
= gimple_switch_num_labels (s
);
6267 tree index
= gimple_switch_index (s
);
6268 tree index_type
= TREE_TYPE (index
);
6269 tree default_label
= gimple_switch_default_label (s
);
6270 basic_block default_label_bb
6271 = label_to_block_fn (func
, CASE_LABEL (default_label
));
6272 basic_block cur_bb
= bb
;
6274 auto_vec
<edge
> new_edges
;
6275 auto_vec
<phi_definition
*> phi_todo_list
;
6276 auto_vec
<profile_count
> edge_counts
;
6277 auto_vec
<profile_probability
> edge_probabilities
;
6279 /* Investigate all labels that and PHI nodes in these edges which
6280 should be fixed after we add new collection of edges. */
6281 for (unsigned i
= 0; i
< labels
; i
++)
6283 tree label
= gimple_switch_label (s
, i
);
6284 basic_block label_bb
= label_to_block_fn (func
, CASE_LABEL (label
));
6285 edge e
= find_edge (bb
, label_bb
);
6286 edge_counts
.safe_push (e
->count ());
6287 edge_probabilities
.safe_push (e
->probability
);
6288 gphi_iterator phi_gsi
;
6290 /* Save PHI definitions that will be destroyed because of an edge
6291 is going to be removed. */
6292 unsigned phi_index
= 0;
6293 for (phi_gsi
= gsi_start_phis (e
->dest
);
6294 !gsi_end_p (phi_gsi
); gsi_next (&phi_gsi
))
6296 gphi
*phi
= phi_gsi
.phi ();
6297 for (unsigned j
= 0; j
< gimple_phi_num_args (phi
); j
++)
6299 if (gimple_phi_arg_edge (phi
, j
) == e
)
6301 tree imm
= gimple_phi_arg_def (phi
, j
);
6302 phi_definition
*p
= new phi_definition (phi_index
, i
,
6304 phi_todo_list
.safe_push (p
);
6312 /* Remove all edges for the current basic block. */
6313 for (int i
= EDGE_COUNT (bb
->succs
) - 1; i
>= 0; i
--)
6315 edge e
= EDGE_SUCC (bb
, i
);
6319 /* Iterate all non-default labels. */
6320 for (unsigned i
= 1; i
< labels
; i
++)
6322 tree label
= gimple_switch_label (s
, i
);
6323 tree low
= CASE_LOW (label
);
6324 tree high
= CASE_HIGH (label
);
6326 if (!useless_type_conversion_p (TREE_TYPE (low
), index_type
))
6327 low
= fold_convert (index_type
, low
);
6329 gimple_stmt_iterator cond_gsi
= gsi_last_bb (cur_bb
);
6333 tree tmp1
= make_temp_ssa_name (boolean_type_node
, NULL
,
6336 gimple
*assign1
= gimple_build_assign (tmp1
, LE_EXPR
, low
,
6339 tree tmp2
= make_temp_ssa_name (boolean_type_node
, NULL
,
6342 if (!useless_type_conversion_p (TREE_TYPE (high
), index_type
))
6343 high
= fold_convert (index_type
, high
);
6344 gimple
*assign2
= gimple_build_assign (tmp2
, LE_EXPR
, index
,
6347 tree tmp3
= make_temp_ssa_name (boolean_type_node
, NULL
,
6349 gimple
*assign3
= gimple_build_assign (tmp3
, BIT_AND_EXPR
, tmp1
,
6352 gsi_insert_before (&cond_gsi
, assign1
, GSI_SAME_STMT
);
6353 gsi_insert_before (&cond_gsi
, assign2
, GSI_SAME_STMT
);
6354 gsi_insert_before (&cond_gsi
, assign3
, GSI_SAME_STMT
);
6356 tree b
= constant_boolean_node (false, boolean_type_node
);
6357 c
= gimple_build_cond (NE_EXPR
, tmp3
, b
, NULL
, NULL
);
6360 c
= gimple_build_cond (EQ_EXPR
, index
, low
, NULL
, NULL
);
6362 gimple_set_location (c
, gimple_location (stmt
));
6364 gsi_insert_before (&cond_gsi
, c
, GSI_SAME_STMT
);
6366 basic_block label_bb
6367 = label_to_block_fn (func
, CASE_LABEL (label
));
6368 edge new_edge
= make_edge (cur_bb
, label_bb
, EDGE_TRUE_VALUE
);
6369 profile_probability prob_sum
= sum_slice
<profile_probability
>
6370 (edge_probabilities
, i
, labels
, profile_probability::never ())
6371 + edge_probabilities
[0];
6373 if (prob_sum
.initialized_p ())
6374 new_edge
->probability
= edge_probabilities
[i
] / prob_sum
;
6376 new_edges
.safe_push (new_edge
);
6380 /* Prepare another basic block that will contain
6382 basic_block next_bb
= create_empty_bb (cur_bb
);
6385 add_bb_to_loop (next_bb
, cur_bb
->loop_father
);
6386 loops_state_set (LOOPS_NEED_FIXUP
);
6389 edge next_edge
= make_edge (cur_bb
, next_bb
, EDGE_FALSE_VALUE
);
6390 next_edge
->probability
= new_edge
->probability
.invert ();
6391 next_bb
->count
= next_edge
->count ();
6394 else /* Link last IF statement and default label
6397 edge e
= make_edge (cur_bb
, default_label_bb
, EDGE_FALSE_VALUE
);
6398 e
->probability
= new_edge
->probability
.invert ();
6399 new_edges
.safe_insert (0, e
);
6403 /* Restore original PHI immediate value. */
6404 for (unsigned i
= 0; i
< phi_todo_list
.length (); i
++)
6406 phi_definition
*phi_def
= phi_todo_list
[i
];
6407 edge new_edge
= new_edges
[phi_def
->label_index
];
6409 gphi_iterator it
= gsi_start_phis (new_edge
->dest
);
6410 for (unsigned i
= 0; i
< phi_def
->phi_index
; i
++)
6413 gphi
*phi
= it
.phi ();
6414 add_phi_arg (phi
, phi_def
->phi_value
, new_edge
, UNKNOWN_LOCATION
);
6418 /* Remove the original GIMPLE switch statement. */
6419 gsi_remove (&gsi
, true);
6424 dump_function_to_file (current_function_decl
, dump_file
, TDF_DETAILS
);
6426 return modified_cfg
;
6429 /* Expand builtins that can't be handled by HSA back-end. */
6434 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
6437 FOR_EACH_BB_FN (bb
, func
)
6439 for (gimple_stmt_iterator gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
);
6442 gimple
*stmt
= gsi_stmt (gsi
);
6444 if (gimple_code (stmt
) != GIMPLE_CALL
)
6447 gcall
*call
= as_a
<gcall
*> (stmt
);
6449 if (!gimple_call_builtin_p (call
, BUILT_IN_NORMAL
))
6452 tree fndecl
= gimple_call_fndecl (stmt
);
6453 enum built_in_function fn
= DECL_FUNCTION_CODE (fndecl
);
6456 case BUILT_IN_CEXPF
:
6457 case BUILT_IN_CEXPIF
:
6458 case BUILT_IN_CEXPI
:
6460 /* Similar to builtins.c (expand_builtin_cexpi), the builtin
6461 can be transformed to: cexp(I * z) = ccos(z) + I * csin(z). */
6462 tree lhs
= gimple_call_lhs (stmt
);
6463 tree rhs
= gimple_call_arg (stmt
, 0);
6464 tree rhs_type
= TREE_TYPE (rhs
);
6465 bool float_type_p
= rhs_type
== float_type_node
;
6466 tree real_part
= make_temp_ssa_name (rhs_type
, NULL
,
6468 tree imag_part
= make_temp_ssa_name (rhs_type
, NULL
,
6472 = mathfn_built_in (rhs_type
, fn
== float_type_p
6473 ? BUILT_IN_COSF
: BUILT_IN_COS
);
6474 gcall
*cos
= gimple_build_call (cos_fndecl
, 1, rhs
);
6475 gimple_call_set_lhs (cos
, real_part
);
6476 gsi_insert_before (&gsi
, cos
, GSI_SAME_STMT
);
6479 = mathfn_built_in (rhs_type
, fn
== float_type_p
6480 ? BUILT_IN_SINF
: BUILT_IN_SIN
);
6481 gcall
*sin
= gimple_build_call (sin_fndecl
, 1, rhs
);
6482 gimple_call_set_lhs (sin
, imag_part
);
6483 gsi_insert_before (&gsi
, sin
, GSI_SAME_STMT
);
6486 gassign
*assign
= gimple_build_assign (lhs
, COMPLEX_EXPR
,
6487 real_part
, imag_part
);
6488 gsi_insert_before (&gsi
, assign
, GSI_SAME_STMT
);
6489 gsi_remove (&gsi
, true);
6500 /* Emit HSA module variables that are global for the entire module. */
6503 emit_hsa_module_variables (void)
6505 hsa_num_threads
= new hsa_symbol (BRIG_TYPE_U32
, BRIG_SEGMENT_PRIVATE
,
6506 BRIG_LINKAGE_MODULE
, true);
6508 hsa_num_threads
->m_name
= "hsa_num_threads";
6510 hsa_brig_emit_omp_symbols ();
6513 /* Generate HSAIL representation of the current function and write into a
6514 special section of the output file. If KERNEL is set, the function will be
6515 considered an HSA kernel callable from the host, otherwise it will be
6516 compiled as an HSA function callable from other HSA code. */
6519 generate_hsa (bool kernel
)
6521 hsa_init_data_for_cfun ();
6523 if (hsa_num_threads
== NULL
)
6524 emit_hsa_module_variables ();
6526 bool modified_cfg
= convert_switch_statements ();
6527 /* Initialize hsa_cfun. */
6528 hsa_cfun
= new hsa_function_representation (cfun
->decl
, kernel
,
6529 SSANAMES (cfun
)->length (),
6531 hsa_cfun
->init_extra_bbs ();
6535 HSA_SORRY_AT (UNKNOWN_LOCATION
,
6536 "support for HSA does not implement transactional memory");
6540 verify_function_arguments (cfun
->decl
);
6541 if (hsa_seen_error ())
6544 hsa_cfun
->m_name
= get_brig_function_name (cfun
->decl
);
6546 gen_function_def_parameters ();
6547 if (hsa_seen_error ())
6552 gen_body_from_gimple ();
6553 if (hsa_seen_error ())
6556 if (hsa_cfun
->m_kernel_dispatch_count
)
6557 init_hsa_num_threads ();
6559 if (hsa_cfun
->m_kern_p
)
6561 hsa_function_summary
*s
6562 = hsa_summaries
->get (cgraph_node::get (hsa_cfun
->m_decl
));
6563 hsa_add_kern_decl_mapping (current_function_decl
, hsa_cfun
->m_name
,
6564 hsa_cfun
->m_maximum_omp_data_size
,
6565 s
->m_gridified_kernel_p
);
6570 for (unsigned i
= 0; i
< hsa_cfun
->m_ssa_map
.length (); i
++)
6571 if (hsa_cfun
->m_ssa_map
[i
])
6572 hsa_cfun
->m_ssa_map
[i
]->verify_ssa ();
6575 FOR_EACH_BB_FN (bb
, cfun
)
6577 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
6579 for (hsa_insn_basic
*insn
= hbb
->m_first_insn
; insn
;
6580 insn
= insn
->m_next
)
6586 hsa_brig_emit_function ();
6589 hsa_deinit_data_for_cfun ();
6594 const pass_data pass_data_gen_hsail
=
6597 "hsagen", /* name */
6598 OPTGROUP_OMP
, /* optinfo_flags */
6599 TV_NONE
, /* tv_id */
6600 PROP_cfg
| PROP_ssa
, /* properties_required */
6601 0, /* properties_provided */
6602 0, /* properties_destroyed */
6603 0, /* todo_flags_start */
6604 0 /* todo_flags_finish */
6607 class pass_gen_hsail
: public gimple_opt_pass
6610 pass_gen_hsail (gcc::context
*ctxt
)
6611 : gimple_opt_pass(pass_data_gen_hsail
, ctxt
)
6614 /* opt_pass methods: */
6615 bool gate (function
*);
6616 unsigned int execute (function
*);
6618 }; // class pass_gen_hsail
6620 /* Determine whether or not to run generation of HSAIL. */
6623 pass_gen_hsail::gate (function
*f
)
6625 return hsa_gen_requested_p ()
6626 && hsa_gpu_implementation_p (f
->decl
);
6630 pass_gen_hsail::execute (function
*)
6632 hsa_function_summary
*s
6633 = hsa_summaries
->get (cgraph_node::get_create (current_function_decl
));
6636 generate_hsa (s
->m_kind
== HSA_KERNEL
);
6637 TREE_ASM_WRITTEN (current_function_decl
) = 1;
6638 return TODO_discard_function
;
6643 /* Create the instance of hsa gen pass. */
6646 make_pass_gen_hsail (gcc::context
*ctxt
)
6648 return new pass_gen_hsail (ctxt
);