1 /* A pass for lowering gimple to HSAIL
2 Copyright (C) 2013-2016 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"
27 #include "hash-table.h"
30 #include "tree-pass.h"
33 #include "basic-block.h"
34 #include "fold-const.h"
36 #include "gimple-iterator.h"
39 #include "gimple-pretty-print.h"
40 #include "diagnostic-core.h"
41 #include "alloc-pool.h"
42 #include "gimple-ssa.h"
43 #include "tree-phinodes.h"
44 #include "stringpool.h"
45 #include "tree-ssanames.h"
47 #include "ssa-iterators.h"
49 #include "print-tree.h"
50 #include "symbol-summary.h"
58 #include "gomp-constants.h"
59 #include "internal-fn.h"
61 #include "stor-layout.h"
63 /* Print a warning message and set that we have seen an error. */
65 #define HSA_SORRY_ATV(location, message, ...) \
69 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
71 inform (location, message, __VA_ARGS__); \
75 /* Same as previous, but highlight a location. */
77 #define HSA_SORRY_AT(location, message) \
81 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
83 inform (location, message); \
87 /* Default number of threads used by kernel dispatch. */
89 #define HSA_DEFAULT_NUM_THREADS 64
91 /* Following structures are defined in the final version
92 of HSA specification. */
94 /* HSA queue packet is shadow structure, originally provided by AMD. */
96 struct hsa_queue_packet
100 uint16_t workgroup_size_x
;
101 uint16_t workgroup_size_y
;
102 uint16_t workgroup_size_z
;
104 uint32_t grid_size_x
;
105 uint32_t grid_size_y
;
106 uint32_t grid_size_z
;
107 uint32_t private_segment_size
;
108 uint32_t group_segment_size
;
109 uint64_t kernel_object
;
110 void *kernarg_address
;
112 uint64_t completion_signal
;
115 /* HSA queue is shadow structure, originally provided by AMD. */
122 uint64_t doorbell_signal
;
128 /* Alloc pools for allocating basic hsa structures such as operands,
129 instructions and other basic entities. */
130 static object_allocator
<hsa_op_address
> *hsa_allocp_operand_address
;
131 static object_allocator
<hsa_op_immed
> *hsa_allocp_operand_immed
;
132 static object_allocator
<hsa_op_reg
> *hsa_allocp_operand_reg
;
133 static object_allocator
<hsa_op_code_list
> *hsa_allocp_operand_code_list
;
134 static object_allocator
<hsa_op_operand_list
> *hsa_allocp_operand_operand_list
;
135 static object_allocator
<hsa_insn_basic
> *hsa_allocp_inst_basic
;
136 static object_allocator
<hsa_insn_phi
> *hsa_allocp_inst_phi
;
137 static object_allocator
<hsa_insn_mem
> *hsa_allocp_inst_mem
;
138 static object_allocator
<hsa_insn_atomic
> *hsa_allocp_inst_atomic
;
139 static object_allocator
<hsa_insn_signal
> *hsa_allocp_inst_signal
;
140 static object_allocator
<hsa_insn_seg
> *hsa_allocp_inst_seg
;
141 static object_allocator
<hsa_insn_cmp
> *hsa_allocp_inst_cmp
;
142 static object_allocator
<hsa_insn_br
> *hsa_allocp_inst_br
;
143 static object_allocator
<hsa_insn_sbr
> *hsa_allocp_inst_sbr
;
144 static object_allocator
<hsa_insn_call
> *hsa_allocp_inst_call
;
145 static object_allocator
<hsa_insn_arg_block
> *hsa_allocp_inst_arg_block
;
146 static object_allocator
<hsa_insn_comment
> *hsa_allocp_inst_comment
;
147 static object_allocator
<hsa_insn_queue
> *hsa_allocp_inst_queue
;
148 static object_allocator
<hsa_insn_srctype
> *hsa_allocp_inst_srctype
;
149 static object_allocator
<hsa_insn_packed
> *hsa_allocp_inst_packed
;
150 static object_allocator
<hsa_insn_cvt
> *hsa_allocp_inst_cvt
;
151 static object_allocator
<hsa_insn_alloca
> *hsa_allocp_inst_alloca
;
152 static object_allocator
<hsa_bb
> *hsa_allocp_bb
;
154 /* List of pointers to all instructions that come from an object allocator. */
155 static vec
<hsa_insn_basic
*> hsa_instructions
;
157 /* List of pointers to all operands that come from an object allocator. */
158 static vec
<hsa_op_base
*> hsa_operands
;
160 hsa_symbol::hsa_symbol ()
161 : m_decl (NULL_TREE
), m_name (NULL
), m_name_number (0),
162 m_directive_offset (0), m_type (BRIG_TYPE_NONE
),
163 m_segment (BRIG_SEGMENT_NONE
), m_linkage (BRIG_LINKAGE_NONE
), m_dim (0),
164 m_cst_value (NULL
), m_global_scope_p (false), m_seen_error (false),
165 m_allocation (BRIG_ALLOCATION_AUTOMATIC
), m_emitted_to_brig (false)
170 hsa_symbol::hsa_symbol (BrigType16_t type
, BrigSegment8_t segment
,
171 BrigLinkage8_t linkage
, bool global_scope_p
,
172 BrigAllocation allocation
, BrigAlignment8_t align
)
173 : m_decl (NULL_TREE
), m_name (NULL
), m_name_number (0),
174 m_directive_offset (0), m_type (type
), m_segment (segment
),
175 m_linkage (linkage
), m_dim (0), m_cst_value (NULL
),
176 m_global_scope_p (global_scope_p
), m_seen_error (false),
177 m_allocation (allocation
), m_emitted_to_brig (false), m_align (align
)
181 unsigned HOST_WIDE_INT
182 hsa_symbol::total_byte_size ()
184 unsigned HOST_WIDE_INT s
185 = hsa_type_bit_size (~BRIG_TYPE_ARRAY_MASK
& m_type
);
186 gcc_assert (s
% BITS_PER_UNIT
== 0);
195 /* Forward declaration. */
198 hsa_type_for_tree_type (const_tree type
, unsigned HOST_WIDE_INT
*dim_p
,
202 hsa_symbol::fillup_for_decl (tree decl
)
205 m_type
= hsa_type_for_tree_type (TREE_TYPE (decl
), &m_dim
, false);
207 if (hsa_seen_error ())
211 /* Constructor of class representing global HSA function/kernel information and
212 state. FNDECL is function declaration, KERNEL_P is true if the function
213 is going to become a HSA kernel. If the function has body, SSA_NAMES_COUNT
214 should be set to number of SSA names used in the function.
215 MODIFIED_CFG is set to true in case we modified control-flow graph
218 hsa_function_representation::hsa_function_representation
219 (tree fdecl
, bool kernel_p
, unsigned ssa_names_count
, bool modified_cfg
)
221 m_reg_count (0), m_input_args (vNULL
),
222 m_output_arg (NULL
), m_spill_symbols (vNULL
), m_global_symbols (vNULL
),
223 m_private_variables (vNULL
), m_called_functions (vNULL
),
224 m_called_internal_fns (vNULL
), m_hbb_count (0),
225 m_in_ssa (true), m_kern_p (kernel_p
), m_declaration_p (false),
226 m_decl (fdecl
), m_internal_fn (NULL
), m_shadow_reg (NULL
),
227 m_kernel_dispatch_count (0), m_maximum_omp_data_size (0),
228 m_seen_error (false), m_temp_symbol_count (0), m_ssa_map (),
229 m_modified_cfg (modified_cfg
)
231 int sym_init_len
= (vec_safe_length (cfun
->local_decls
) / 2) + 1;;
232 m_local_symbols
= new hash_table
<hsa_noop_symbol_hasher
> (sym_init_len
);
233 m_ssa_map
.safe_grow_cleared (ssa_names_count
);
236 /* Constructor of class representing HSA function information that
237 is derived for an internal function. */
238 hsa_function_representation::hsa_function_representation (hsa_internal_fn
*fn
)
239 : m_reg_count (0), m_input_args (vNULL
),
240 m_output_arg (NULL
), m_local_symbols (NULL
),
241 m_spill_symbols (vNULL
), m_global_symbols (vNULL
),
242 m_private_variables (vNULL
), m_called_functions (vNULL
),
243 m_called_internal_fns (vNULL
), m_hbb_count (0),
244 m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL
),
245 m_internal_fn (fn
), m_shadow_reg (NULL
), m_kernel_dispatch_count (0),
246 m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
249 /* Destructor of class holding function/kernel-wide information and state. */
251 hsa_function_representation::~hsa_function_representation ()
253 /* Kernel names are deallocated at the end of BRIG output when deallocating
254 hsa_decl_kernel_mapping. */
255 if (!m_kern_p
|| m_seen_error
)
258 for (unsigned i
= 0; i
< m_input_args
.length (); i
++)
259 delete m_input_args
[i
];
260 m_input_args
.release ();
263 delete m_local_symbols
;
265 for (unsigned i
= 0; i
< m_spill_symbols
.length (); i
++)
266 delete m_spill_symbols
[i
];
267 m_spill_symbols
.release ();
270 for (unsigned i
= 0; i
< m_global_symbols
.iterate (i
, &sym
); i
++)
271 if (sym
->m_linkage
!= BRIG_ALLOCATION_PROGRAM
)
273 m_global_symbols
.release ();
275 for (unsigned i
= 0; i
< m_private_variables
.length (); i
++)
276 delete m_private_variables
[i
];
277 m_private_variables
.release ();
278 m_called_functions
.release ();
279 m_ssa_map
.release ();
281 for (unsigned i
= 0; i
< m_called_internal_fns
.length (); i
++)
282 delete m_called_internal_fns
[i
];
286 hsa_function_representation::get_shadow_reg ()
288 /* If we compile a function with kernel dispatch and does not set
289 an optimization level, the function won't be inlined and
297 /* Append the shadow argument. */
298 hsa_symbol
*shadow
= new hsa_symbol (BRIG_TYPE_U64
, BRIG_SEGMENT_KERNARG
,
299 BRIG_LINKAGE_FUNCTION
);
300 m_input_args
.safe_push (shadow
);
301 shadow
->m_name
= "hsa_runtime_shadow";
303 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_U64
);
304 hsa_op_address
*addr
= new hsa_op_address (shadow
);
306 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, BRIG_TYPE_U64
, r
, addr
);
307 hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
))->append_insn (mem
);
313 bool hsa_function_representation::has_shadow_reg_p ()
315 return m_shadow_reg
!= NULL
;
319 hsa_function_representation::init_extra_bbs ()
321 hsa_init_new_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
322 hsa_init_new_bb (EXIT_BLOCK_PTR_FOR_FN (cfun
));
326 hsa_function_representation::update_dominance ()
330 free_dominance_info (CDI_DOMINATORS
);
331 calculate_dominance_info (CDI_DOMINATORS
);
336 hsa_function_representation::create_hsa_temporary (BrigType16_t type
)
338 hsa_symbol
*s
= new hsa_symbol (type
, BRIG_SEGMENT_PRIVATE
,
339 BRIG_LINKAGE_FUNCTION
);
340 s
->m_name_number
= m_temp_symbol_count
++;
342 hsa_cfun
->m_private_variables
.safe_push (s
);
347 hsa_function_representation::get_linkage ()
350 return BRIG_LINKAGE_PROGRAM
;
352 return m_kern_p
|| TREE_PUBLIC (m_decl
) ?
353 BRIG_LINKAGE_PROGRAM
: BRIG_LINKAGE_MODULE
;
356 /* Hash map of simple OMP builtins. */
357 static hash_map
<nofree_string_hash
, omp_simple_builtin
> *omp_simple_builtins
360 /* Warning messages for OMP builtins. */
362 #define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP " \
364 #define HSA_WARN_TIMING_ROUTINE "support for HSA does not implement OpenMP " \
366 #define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \
367 "undefined semantics within target regions, support for HSA ignores them"
368 #define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP " \
371 /* Initialize hash map with simple OMP builtins. */
374 hsa_init_simple_builtins ()
376 if (omp_simple_builtins
!= NULL
)
380 = new hash_map
<nofree_string_hash
, omp_simple_builtin
> ();
382 omp_simple_builtin omp_builtins
[] =
384 omp_simple_builtin ("omp_get_initial_device", NULL
, false,
385 new hsa_op_immed (GOMP_DEVICE_HOST
,
386 (BrigType16_t
) BRIG_TYPE_S32
)),
387 omp_simple_builtin ("omp_is_initial_device", NULL
, false,
388 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
389 omp_simple_builtin ("omp_get_dynamic", NULL
, false,
390 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
391 omp_simple_builtin ("omp_set_dynamic", NULL
, false, NULL
),
392 omp_simple_builtin ("omp_init_lock", HSA_WARN_LOCK_ROUTINE
, true),
393 omp_simple_builtin ("omp_init_lock_with_hint", HSA_WARN_LOCK_ROUTINE
,
395 omp_simple_builtin ("omp_init_nest_lock_with_hint", HSA_WARN_LOCK_ROUTINE
,
397 omp_simple_builtin ("omp_destroy_lock", HSA_WARN_LOCK_ROUTINE
, true),
398 omp_simple_builtin ("omp_set_lock", HSA_WARN_LOCK_ROUTINE
, true),
399 omp_simple_builtin ("omp_unset_lock", HSA_WARN_LOCK_ROUTINE
, true),
400 omp_simple_builtin ("omp_test_lock", HSA_WARN_LOCK_ROUTINE
, true),
401 omp_simple_builtin ("omp_get_wtime", HSA_WARN_TIMING_ROUTINE
, true),
402 omp_simple_builtin ("omp_get_wtick", HSA_WARN_TIMING_ROUTINE
, true),
403 omp_simple_builtin ("omp_target_alloc", HSA_WARN_MEMORY_ROUTINE
, false,
404 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_U64
)),
405 omp_simple_builtin ("omp_target_free", HSA_WARN_MEMORY_ROUTINE
, false),
406 omp_simple_builtin ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE
,
408 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
409 omp_simple_builtin ("omp_target_memcpy", HSA_WARN_MEMORY_ROUTINE
, false,
410 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
411 omp_simple_builtin ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE
,
413 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
414 omp_simple_builtin ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE
,
416 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
417 omp_simple_builtin ("omp_target_disassociate_ptr",
418 HSA_WARN_MEMORY_ROUTINE
,
420 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
421 omp_simple_builtin ("omp_set_max_active_levels",
422 "Support for HSA only allows only one active level, "
423 "call to omp_set_max_active_levels will be ignored "
424 "in the generated HSAIL",
426 omp_simple_builtin ("omp_get_max_active_levels", NULL
, false,
427 new hsa_op_immed (1, (BrigType16_t
) BRIG_TYPE_S32
)),
428 omp_simple_builtin ("omp_in_final", NULL
, false,
429 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
430 omp_simple_builtin ("omp_get_proc_bind", HSA_WARN_AFFINITY
, false,
431 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
432 omp_simple_builtin ("omp_get_num_places", HSA_WARN_AFFINITY
, false,
433 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
434 omp_simple_builtin ("omp_get_place_num_procs", HSA_WARN_AFFINITY
, false,
435 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
436 omp_simple_builtin ("omp_get_place_proc_ids", HSA_WARN_AFFINITY
, false,
438 omp_simple_builtin ("omp_get_place_num", HSA_WARN_AFFINITY
, false,
439 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
440 omp_simple_builtin ("omp_get_partition_num_places", HSA_WARN_AFFINITY
,
442 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
443 omp_simple_builtin ("omp_get_partition_place_nums", HSA_WARN_AFFINITY
,
445 omp_simple_builtin ("omp_set_default_device",
446 "omp_set_default_device has undefined semantics "
447 "within target regions, support for HSA ignores it",
449 omp_simple_builtin ("omp_get_default_device",
450 "omp_get_default_device has undefined semantics "
451 "within target regions, support for HSA ignores it",
453 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
454 omp_simple_builtin ("omp_get_num_devices",
455 "omp_get_num_devices has undefined semantics "
456 "within target regions, support for HSA ignores it",
458 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
459 omp_simple_builtin ("omp_get_num_procs", NULL
, true, NULL
),
460 omp_simple_builtin ("omp_get_cancellation", NULL
, true, NULL
),
461 omp_simple_builtin ("omp_set_nested", NULL
, true, NULL
),
462 omp_simple_builtin ("omp_get_nested", NULL
, true, NULL
),
463 omp_simple_builtin ("omp_set_schedule", NULL
, true, NULL
),
464 omp_simple_builtin ("omp_get_schedule", NULL
, true, NULL
),
465 omp_simple_builtin ("omp_get_thread_limit", NULL
, true, NULL
),
466 omp_simple_builtin ("omp_get_team_size", NULL
, true, NULL
),
467 omp_simple_builtin ("omp_get_ancestor_thread_num", NULL
, true, NULL
),
468 omp_simple_builtin ("omp_get_max_task_priority", NULL
, true, NULL
)
471 unsigned count
= sizeof (omp_builtins
) / sizeof (omp_simple_builtin
);
473 for (unsigned i
= 0; i
< count
; i
++)
474 omp_simple_builtins
->put (omp_builtins
[i
].m_name
, omp_builtins
[i
]);
477 /* Allocate HSA structures that we need only while generating with this. */
480 hsa_init_data_for_cfun ()
482 hsa_init_compilation_unit_data ();
483 hsa_allocp_operand_address
484 = new object_allocator
<hsa_op_address
> ("HSA address operands");
485 hsa_allocp_operand_immed
486 = new object_allocator
<hsa_op_immed
> ("HSA immediate operands");
487 hsa_allocp_operand_reg
488 = new object_allocator
<hsa_op_reg
> ("HSA register operands");
489 hsa_allocp_operand_code_list
490 = new object_allocator
<hsa_op_code_list
> ("HSA code list operands");
491 hsa_allocp_operand_operand_list
492 = new object_allocator
<hsa_op_operand_list
> ("HSA operand list operands");
493 hsa_allocp_inst_basic
494 = new object_allocator
<hsa_insn_basic
> ("HSA basic instructions");
496 = new object_allocator
<hsa_insn_phi
> ("HSA phi operands");
498 = new object_allocator
<hsa_insn_mem
> ("HSA memory instructions");
499 hsa_allocp_inst_atomic
500 = new object_allocator
<hsa_insn_atomic
> ("HSA atomic instructions");
501 hsa_allocp_inst_signal
502 = new object_allocator
<hsa_insn_signal
> ("HSA signal instructions");
504 = new object_allocator
<hsa_insn_seg
> ("HSA segment conversion "
507 = new object_allocator
<hsa_insn_cmp
> ("HSA comparison instructions");
509 = new object_allocator
<hsa_insn_br
> ("HSA branching instructions");
511 = new object_allocator
<hsa_insn_sbr
> ("HSA switch branching instructions");
513 = new object_allocator
<hsa_insn_call
> ("HSA call instructions");
514 hsa_allocp_inst_arg_block
515 = new object_allocator
<hsa_insn_arg_block
> ("HSA arg block instructions");
516 hsa_allocp_inst_comment
517 = new object_allocator
<hsa_insn_comment
> ("HSA comment instructions");
518 hsa_allocp_inst_queue
519 = new object_allocator
<hsa_insn_queue
> ("HSA queue instructions");
520 hsa_allocp_inst_srctype
521 = new object_allocator
<hsa_insn_srctype
> ("HSA source type instructions");
522 hsa_allocp_inst_packed
523 = new object_allocator
<hsa_insn_packed
> ("HSA packed instructions");
525 = new object_allocator
<hsa_insn_cvt
> ("HSA convert instructions");
526 hsa_allocp_inst_alloca
527 = new object_allocator
<hsa_insn_alloca
> ("HSA alloca instructions");
528 hsa_allocp_bb
= new object_allocator
<hsa_bb
> ("HSA basic blocks");
531 /* Deinitialize HSA subsystem and free all allocated memory. */
534 hsa_deinit_data_for_cfun (void)
538 FOR_ALL_BB_FN (bb
, cfun
)
541 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
546 for (unsigned int i
= 0; i
< hsa_operands
.length (); i
++)
547 hsa_destroy_operand (hsa_operands
[i
]);
549 hsa_operands
.release ();
551 for (unsigned i
= 0; i
< hsa_instructions
.length (); i
++)
552 hsa_destroy_insn (hsa_instructions
[i
]);
554 hsa_instructions
.release ();
556 if (omp_simple_builtins
!= NULL
)
558 delete omp_simple_builtins
;
559 omp_simple_builtins
= NULL
;
562 delete hsa_allocp_operand_address
;
563 delete hsa_allocp_operand_immed
;
564 delete hsa_allocp_operand_reg
;
565 delete hsa_allocp_operand_code_list
;
566 delete hsa_allocp_operand_operand_list
;
567 delete hsa_allocp_inst_basic
;
568 delete hsa_allocp_inst_phi
;
569 delete hsa_allocp_inst_atomic
;
570 delete hsa_allocp_inst_mem
;
571 delete hsa_allocp_inst_signal
;
572 delete hsa_allocp_inst_seg
;
573 delete hsa_allocp_inst_cmp
;
574 delete hsa_allocp_inst_br
;
575 delete hsa_allocp_inst_sbr
;
576 delete hsa_allocp_inst_call
;
577 delete hsa_allocp_inst_arg_block
;
578 delete hsa_allocp_inst_comment
;
579 delete hsa_allocp_inst_queue
;
580 delete hsa_allocp_inst_srctype
;
581 delete hsa_allocp_inst_packed
;
582 delete hsa_allocp_inst_cvt
;
583 delete hsa_allocp_inst_alloca
;
584 delete hsa_allocp_bb
;
588 /* Return the type which holds addresses in the given SEGMENT. */
591 hsa_get_segment_addr_type (BrigSegment8_t segment
)
595 case BRIG_SEGMENT_NONE
:
598 case BRIG_SEGMENT_FLAT
:
599 case BRIG_SEGMENT_GLOBAL
:
600 case BRIG_SEGMENT_READONLY
:
601 case BRIG_SEGMENT_KERNARG
:
602 return hsa_machine_large_p () ? BRIG_TYPE_U64
: BRIG_TYPE_U32
;
604 case BRIG_SEGMENT_GROUP
:
605 case BRIG_SEGMENT_PRIVATE
:
606 case BRIG_SEGMENT_SPILL
:
607 case BRIG_SEGMENT_ARG
:
608 return BRIG_TYPE_U32
;
613 /* Return integer brig type according to provided SIZE in bytes. If SIGN
614 is set to true, return signed integer type. */
617 get_integer_type_by_bytes (unsigned size
, bool sign
)
625 return BRIG_TYPE_S16
;
627 return BRIG_TYPE_S32
;
629 return BRIG_TYPE_S64
;
639 return BRIG_TYPE_U16
;
641 return BRIG_TYPE_U32
;
643 return BRIG_TYPE_U64
;
651 /* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
652 are assumed to use flat addressing. If min32int is true, always expand
653 integer types to one that has at least 32 bits. */
656 hsa_type_for_scalar_tree_type (const_tree type
, bool min32int
)
660 BrigType16_t res
= BRIG_TYPE_NONE
;
662 gcc_checking_assert (TYPE_P (type
));
663 gcc_checking_assert (!AGGREGATE_TYPE_P (type
));
664 if (POINTER_TYPE_P (type
))
665 return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
667 if (TREE_CODE (type
) == VECTOR_TYPE
|| TREE_CODE (type
) == COMPLEX_TYPE
)
668 base
= TREE_TYPE (type
);
672 if (!tree_fits_uhwi_p (TYPE_SIZE (base
)))
674 HSA_SORRY_ATV (EXPR_LOCATION (type
),
675 "support for HSA does not implement huge or "
676 "variable-sized type %T", type
);
680 bsize
= tree_to_uhwi (TYPE_SIZE (base
));
681 unsigned byte_size
= bsize
/ BITS_PER_UNIT
;
682 if (INTEGRAL_TYPE_P (base
))
683 res
= get_integer_type_by_bytes (byte_size
, !TYPE_UNSIGNED (base
));
684 else if (SCALAR_FLOAT_TYPE_P (base
))
702 if (res
== BRIG_TYPE_NONE
)
704 HSA_SORRY_ATV (EXPR_LOCATION (type
),
705 "support for HSA does not implement type %T", type
);
709 if (TREE_CODE (type
) == VECTOR_TYPE
)
711 HOST_WIDE_INT tsize
= tree_to_uhwi (TYPE_SIZE (type
));
715 HSA_SORRY_ATV (EXPR_LOCATION (type
),
716 "support for HSA does not implement a vector type "
717 "where a type and unit size are equal: %T", type
);
724 res
|= BRIG_TYPE_PACK_32
;
727 res
|= BRIG_TYPE_PACK_64
;
730 res
|= BRIG_TYPE_PACK_128
;
733 HSA_SORRY_ATV (EXPR_LOCATION (type
),
734 "support for HSA does not implement type %T", type
);
740 /* Registers/immediate operands can only be 32bit or more except for
742 if (res
== BRIG_TYPE_U8
|| res
== BRIG_TYPE_U16
)
744 else if (res
== BRIG_TYPE_S8
|| res
== BRIG_TYPE_S16
)
748 if (TREE_CODE (type
) == COMPLEX_TYPE
)
750 unsigned bsize
= 2 * hsa_type_bit_size (res
);
751 res
= hsa_bittype_for_bitsize (bsize
);
757 /* Returns the BRIG type we need to load/store entities of TYPE. */
760 mem_type_for_type (BrigType16_t type
)
762 /* HSA has non-intuitive constraints on load/store types. If it's
763 a bit-type it _must_ be B128, if it's not a bit-type it must be
764 64bit max. So for loading entities of 128 bits (e.g. vectors)
765 we have to to B128, while for loading the rest we have to use the
766 input type (??? or maybe also flattened to a equally sized non-vector
768 if ((type
& BRIG_TYPE_PACK_MASK
) == BRIG_TYPE_PACK_128
)
769 return BRIG_TYPE_B128
;
770 else if (hsa_btype_p (type
) || hsa_type_packed_p (type
))
772 unsigned bitsize
= hsa_type_bit_size (type
);
774 return hsa_uint_for_bitsize (bitsize
);
776 return hsa_bittype_for_bitsize (bitsize
);
781 /* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some
782 kind of array will be generated, setting DIM appropriately. Otherwise, it
783 will be set to zero. */
786 hsa_type_for_tree_type (const_tree type
, unsigned HOST_WIDE_INT
*dim_p
= NULL
,
787 bool min32int
= false)
789 gcc_checking_assert (TYPE_P (type
));
790 if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type
)))
792 HSA_SORRY_ATV (EXPR_LOCATION (type
), "support for HSA does not "
793 "implement huge or variable-sized type %T", type
);
794 return BRIG_TYPE_NONE
;
797 if (RECORD_OR_UNION_TYPE_P (type
))
800 *dim_p
= tree_to_uhwi (TYPE_SIZE_UNIT (type
));
801 return BRIG_TYPE_U8
| BRIG_TYPE_ARRAY
;
804 if (TREE_CODE (type
) == ARRAY_TYPE
)
806 /* We try to be nice and use the real base-type when this is an array of
807 scalars and only resort to an array of bytes if the type is more
810 unsigned HOST_WIDE_INT dim
= 1;
812 while (TREE_CODE (type
) == ARRAY_TYPE
)
814 tree domain
= TYPE_DOMAIN (type
);
815 if (!TYPE_MIN_VALUE (domain
)
816 || !TYPE_MAX_VALUE (domain
)
817 || !tree_fits_shwi_p (TYPE_MIN_VALUE (domain
))
818 || !tree_fits_shwi_p (TYPE_MAX_VALUE (domain
)))
820 HSA_SORRY_ATV (EXPR_LOCATION (type
),
821 "support for HSA does not implement array %T with "
822 "unknown bounds", type
);
823 return BRIG_TYPE_NONE
;
825 HOST_WIDE_INT min
= tree_to_shwi (TYPE_MIN_VALUE (domain
));
826 HOST_WIDE_INT max
= tree_to_shwi (TYPE_MAX_VALUE (domain
));
827 dim
= dim
* (unsigned HOST_WIDE_INT
) (max
- min
+ 1);
828 type
= TREE_TYPE (type
);
832 if (RECORD_OR_UNION_TYPE_P (type
))
834 dim
= dim
* tree_to_uhwi (TYPE_SIZE_UNIT (type
));
838 res
= hsa_type_for_scalar_tree_type (type
, false);
842 return res
| BRIG_TYPE_ARRAY
;
849 return hsa_type_for_scalar_tree_type (type
, min32int
);
852 /* Returns true if converting from STYPE into DTYPE needs the _CVT
853 opcode. If false a normal _MOV is enough. */
856 hsa_needs_cvt (BrigType16_t dtype
, BrigType16_t stype
)
858 if (hsa_btype_p (dtype
))
861 /* float <-> int conversions are real converts. */
862 if (hsa_type_float_p (dtype
) != hsa_type_float_p (stype
))
864 /* When both types have different size, then we need CVT as well. */
865 if (hsa_type_bit_size (dtype
) != hsa_type_bit_size (stype
))
870 /* Lookup or create the associated hsa_symbol structure with a given VAR_DECL
871 or lookup the hsa_structure corresponding to a PARM_DECL. */
874 get_symbol_for_decl (tree decl
)
877 hsa_symbol
dummy (BRIG_TYPE_NONE
, BRIG_SEGMENT_NONE
, BRIG_LINKAGE_NONE
);
879 gcc_assert (TREE_CODE (decl
) == PARM_DECL
880 || TREE_CODE (decl
) == RESULT_DECL
881 || TREE_CODE (decl
) == VAR_DECL
);
885 bool is_in_global_vars
886 = TREE_CODE (decl
) == VAR_DECL
&& is_global_var (decl
);
888 if (is_in_global_vars
)
889 slot
= hsa_global_variable_symbols
->find_slot (&dummy
, INSERT
);
891 slot
= hsa_cfun
->m_local_symbols
->find_slot (&dummy
, INSERT
);
893 gcc_checking_assert (slot
);
896 hsa_symbol
*sym
= (*slot
);
898 /* If the symbol is problematic, mark current function also as
900 if (sym
->m_seen_error
)
903 /* PR hsa/70234: If a global variable was marked to be emitted,
904 but HSAIL generation of a function using the variable fails,
905 we should retry to emit the variable in context of a different
908 Iterate elements whether a symbol is already in m_global_symbols
910 if (is_in_global_vars
&& !sym
->m_emitted_to_brig
)
912 for (unsigned i
= 0; i
< hsa_cfun
->m_global_symbols
.length (); i
++)
913 if (hsa_cfun
->m_global_symbols
[i
] == sym
)
915 hsa_cfun
->m_global_symbols
.safe_push (sym
);
923 gcc_assert (TREE_CODE (decl
) == VAR_DECL
);
924 BrigAlignment8_t align
= hsa_object_alignment (decl
);
926 if (is_in_global_vars
)
928 sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_GLOBAL
,
929 BRIG_LINKAGE_PROGRAM
, true,
930 BRIG_ALLOCATION_PROGRAM
, align
);
931 hsa_cfun
->m_global_symbols
.safe_push (sym
);
935 /* As generation of efficient memory copy instructions relies
936 on alignment greater or equal to 8 bytes,
937 we need to increase alignment of all aggregate types.. */
938 if (AGGREGATE_TYPE_P (TREE_TYPE (decl
)))
939 align
= MAX ((BrigAlignment8_t
) BRIG_ALIGNMENT_8
, align
);
941 /* PARM_DECL and RESULT_DECL should be already in m_local_symbols. */
942 gcc_assert (TREE_CODE (decl
) == VAR_DECL
);
944 sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_PRIVATE
,
945 BRIG_LINKAGE_FUNCTION
);
946 sym
->m_align
= align
;
947 hsa_cfun
->m_private_variables
.safe_push (sym
);
950 sym
->fillup_for_decl (decl
);
951 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_binded_function
->decl
;
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
&& s
->m_gpu_implementation_p
)
981 d
= s
->m_binded_function
->decl
;
983 /* IPA split can create a function that has no host equivalent. */
987 char *name
= xstrdup (hsa_get_declaration_name (d
));
988 hsa_sanitize_name (name
);
993 /* Create a spill symbol of type TYPE. */
996 hsa_get_spill_symbol (BrigType16_t type
)
998 hsa_symbol
*sym
= new hsa_symbol (type
, BRIG_SEGMENT_SPILL
,
999 BRIG_LINKAGE_FUNCTION
);
1000 hsa_cfun
->m_spill_symbols
.safe_push (sym
);
1004 /* Create a symbol for a read-only string constant. */
1006 hsa_get_string_cst_symbol (tree string_cst
)
1008 gcc_checking_assert (TREE_CODE (string_cst
) == STRING_CST
);
1010 hsa_symbol
**slot
= hsa_cfun
->m_string_constants_map
.get (string_cst
);
1014 hsa_op_immed
*cst
= new hsa_op_immed (string_cst
);
1015 hsa_symbol
*sym
= new hsa_symbol (cst
->m_type
, BRIG_SEGMENT_GLOBAL
,
1016 BRIG_LINKAGE_MODULE
, true,
1017 BRIG_ALLOCATION_AGENT
);
1018 sym
->m_cst_value
= cst
;
1019 sym
->m_dim
= TREE_STRING_LENGTH (string_cst
);
1020 sym
->m_name_number
= hsa_cfun
->m_global_symbols
.length ();
1022 hsa_cfun
->m_global_symbols
.safe_push (sym
);
1023 hsa_cfun
->m_string_constants_map
.put (string_cst
, sym
);
1027 /* Constructor of the ancestor of all operands. K is BRIG kind that identified
1028 what the operator is. */
1030 hsa_op_base::hsa_op_base (BrigKind16_t k
)
1031 : m_next (NULL
), m_brig_op_offset (0), m_kind (k
)
1033 hsa_operands
.safe_push (this);
1036 /* Constructor of ancestor of all operands which have a type. K is BRIG kind
1037 that identified what the operator is. T is the type of the operator. */
1039 hsa_op_with_type::hsa_op_with_type (BrigKind16_t k
, BrigType16_t t
)
1040 : hsa_op_base (k
), m_type (t
)
1045 hsa_op_with_type::get_in_type (BrigType16_t dtype
, hsa_bb
*hbb
)
1047 if (m_type
== dtype
)
1052 if (hsa_needs_cvt (dtype
, m_type
))
1054 dest
= new hsa_op_reg (dtype
);
1055 hbb
->append_insn (new hsa_insn_cvt (dest
, this));
1059 dest
= new hsa_op_reg (m_type
);
1060 hbb
->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
1061 dest
->m_type
, dest
, this));
1063 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
1064 type of the operand must be same as type of the instruction. */
1065 dest
->m_type
= dtype
;
1071 /* Constructor of class representing HSA immediate values. TREE_VAL is the
1072 tree representation of the immediate value. If min32int is true,
1073 always expand integer types to one that has at least 32 bits. */
1075 hsa_op_immed::hsa_op_immed (tree tree_val
, bool min32int
)
1076 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
,
1077 hsa_type_for_tree_type (TREE_TYPE (tree_val
), NULL
,
1080 if (hsa_seen_error ())
1083 gcc_checking_assert ((is_gimple_min_invariant (tree_val
)
1084 && (!POINTER_TYPE_P (TREE_TYPE (tree_val
))
1085 || TREE_CODE (tree_val
) == INTEGER_CST
))
1086 || TREE_CODE (tree_val
) == CONSTRUCTOR
);
1087 m_tree_value
= tree_val
;
1089 /* Verify that all elements of a constructor are constants. */
1090 if (TREE_CODE (m_tree_value
) == CONSTRUCTOR
)
1091 for (unsigned i
= 0;
1092 i
< vec_safe_length (CONSTRUCTOR_ELTS (m_tree_value
)); i
++)
1094 tree v
= CONSTRUCTOR_ELT (m_tree_value
, i
)->value
;
1095 if (!CONSTANT_CLASS_P (v
))
1097 HSA_SORRY_AT (EXPR_LOCATION (tree_val
),
1098 "HSA ctor should have only constants");
1104 /* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
1105 integer representation of the immediate value. TYPE is BRIG type. */
1107 hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value
, BrigType16_t type
)
1108 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
, type
),
1111 gcc_assert (hsa_type_integer_p (type
));
1112 m_int_value
= integer_value
;
1115 hsa_op_immed::hsa_op_immed ()
1116 : hsa_op_with_type (BRIG_KIND_NONE
, BRIG_TYPE_NONE
)
1120 /* New operator to allocate immediate operands from pool alloc. */
1123 hsa_op_immed::operator new (size_t)
1125 return hsa_allocp_operand_immed
->allocate_raw ();
1130 hsa_op_immed::~hsa_op_immed ()
1134 /* Change type of the immediate value to T. */
1137 hsa_op_immed::set_type (BrigType16_t t
)
1142 /* Constructor of class representing HSA registers and pseudo-registers. T is
1143 the BRIG type of the new register. */
1145 hsa_op_reg::hsa_op_reg (BrigType16_t t
)
1146 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER
, t
), m_gimple_ssa (NULL_TREE
),
1147 m_def_insn (NULL
), m_spill_sym (NULL
), m_order (hsa_cfun
->m_reg_count
++),
1148 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
1152 /* New operator to allocate a register from pool alloc. */
1155 hsa_op_reg::operator new (size_t)
1157 return hsa_allocp_operand_reg
->allocate_raw ();
1160 /* Verify register operand. */
1163 hsa_op_reg::verify_ssa ()
1165 /* Verify that each HSA register has a definition assigned.
1166 Exceptions are VAR_DECL and PARM_DECL that are a default
1168 gcc_checking_assert (m_def_insn
1169 || (m_gimple_ssa
!= NULL
1170 && (!SSA_NAME_VAR (m_gimple_ssa
)
1171 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa
))
1173 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa
)));
1175 /* Verify that every use of the register is really present
1176 in an instruction. */
1177 for (unsigned i
= 0; i
< m_uses
.length (); i
++)
1179 hsa_insn_basic
*use
= m_uses
[i
];
1181 bool is_visited
= false;
1182 for (unsigned j
= 0; j
< use
->operand_count (); j
++)
1184 hsa_op_base
*u
= use
->get_op (j
);
1185 hsa_op_address
*addr
; addr
= dyn_cast
<hsa_op_address
*> (u
);
1186 if (addr
&& addr
->m_reg
)
1191 bool r
= !addr
&& use
->op_output_p (j
);
1195 error ("HSA SSA name defined by instruction that is supposed "
1197 debug_hsa_operand (this);
1198 debug_hsa_insn (use
);
1199 internal_error ("HSA SSA verification failed");
1208 error ("HSA SSA name not among operands of instruction that is "
1209 "supposed to use it");
1210 debug_hsa_operand (this);
1211 debug_hsa_insn (use
);
1212 internal_error ("HSA SSA verification failed");
1217 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, hsa_op_reg
*r
,
1218 HOST_WIDE_INT offset
)
1219 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (r
),
1220 m_imm_offset (offset
)
1224 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, HOST_WIDE_INT offset
)
1225 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (NULL
),
1226 m_imm_offset (offset
)
1230 hsa_op_address::hsa_op_address (hsa_op_reg
*r
, HOST_WIDE_INT offset
)
1231 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (NULL
), m_reg (r
),
1232 m_imm_offset (offset
)
1236 /* New operator to allocate address operands from pool alloc. */
1239 hsa_op_address::operator new (size_t)
1241 return hsa_allocp_operand_address
->allocate_raw ();
1244 /* Constructor of an operand referring to HSAIL code. */
1246 hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF
),
1247 m_directive_offset (0)
1251 /* Constructor of an operand representing a code list. Set it up so that it
1252 can contain ELEMENTS number of elements. */
1254 hsa_op_code_list::hsa_op_code_list (unsigned elements
)
1255 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST
)
1257 m_offsets
.create (1);
1258 m_offsets
.safe_grow_cleared (elements
);
1261 /* New operator to allocate code list operands from pool alloc. */
1264 hsa_op_code_list::operator new (size_t)
1266 return hsa_allocp_operand_code_list
->allocate_raw ();
1269 /* Constructor of an operand representing an operand list.
1270 Set it up so that it can contain ELEMENTS number of elements. */
1272 hsa_op_operand_list::hsa_op_operand_list (unsigned elements
)
1273 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST
)
1275 m_offsets
.create (elements
);
1276 m_offsets
.safe_grow (elements
);
1279 /* New operator to allocate operand list operands from pool alloc. */
1282 hsa_op_operand_list::operator new (size_t)
1284 return hsa_allocp_operand_operand_list
->allocate_raw ();
1287 hsa_op_operand_list::~hsa_op_operand_list ()
1289 m_offsets
.release ();
1294 hsa_function_representation::reg_for_gimple_ssa (tree ssa
)
1298 gcc_checking_assert (TREE_CODE (ssa
) == SSA_NAME
);
1299 if (m_ssa_map
[SSA_NAME_VERSION (ssa
)])
1300 return m_ssa_map
[SSA_NAME_VERSION (ssa
)];
1302 hreg
= new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa
),
1304 hreg
->m_gimple_ssa
= ssa
;
1305 m_ssa_map
[SSA_NAME_VERSION (ssa
)] = hreg
;
1311 hsa_op_reg::set_definition (hsa_insn_basic
*insn
)
1313 if (hsa_cfun
->m_in_ssa
)
1315 gcc_checking_assert (!m_def_insn
);
1322 /* Constructor of the class which is the bases of all instructions and directly
1323 represents the most basic ones. NOPS is the number of operands that the
1324 operand vector will contain (and which will be cleared). OP is the opcode
1325 of the instruction. This constructor does not set type. */
1327 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
)
1329 m_next (NULL
), m_bb (NULL
), m_opcode (opc
), m_number (0),
1330 m_type (BRIG_TYPE_NONE
), m_brig_offset (0)
1333 m_operands
.safe_grow_cleared (nops
);
1335 hsa_instructions
.safe_push (this);
1338 /* Make OP the operand number INDEX of operands of this instruction. If OP is a
1339 register or an address containing a register, then either set the definition
1340 of the register to this instruction if it an output operand or add this
1341 instruction to the uses if it is an input one. */
1344 hsa_insn_basic::set_op (int index
, hsa_op_base
*op
)
1346 /* Each address operand is always use. */
1347 hsa_op_address
*addr
= dyn_cast
<hsa_op_address
*> (op
);
1348 if (addr
&& addr
->m_reg
)
1349 addr
->m_reg
->m_uses
.safe_push (this);
1352 hsa_op_reg
*reg
= dyn_cast
<hsa_op_reg
*> (op
);
1355 if (op_output_p (index
))
1356 reg
->set_definition (this);
1358 reg
->m_uses
.safe_push (this);
1362 m_operands
[index
] = op
;
1365 /* Get INDEX-th operand of the instruction. */
1368 hsa_insn_basic::get_op (int index
)
1370 return m_operands
[index
];
1373 /* Get address of INDEX-th operand of the instruction. */
1376 hsa_insn_basic::get_op_addr (int index
)
1378 return &m_operands
[index
];
1381 /* Get number of operands of the instruction. */
1383 hsa_insn_basic::operand_count ()
1385 return m_operands
.length ();
1388 /* Constructor of the class which is the bases of all instructions and directly
1389 represents the most basic ones. NOPS is the number of operands that the
1390 operand vector will contain (and which will be cleared). OPC is the opcode
1391 of the instruction, T is the type of the instruction. */
1393 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
, BrigType16_t t
,
1394 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1395 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1396 : m_prev (NULL
), m_next (NULL
), m_bb (NULL
), m_opcode (opc
),m_number (0),
1397 m_type (t
), m_brig_offset (0)
1400 m_operands
.safe_grow_cleared (nops
);
1404 gcc_checking_assert (nops
>= 1);
1410 gcc_checking_assert (nops
>= 2);
1416 gcc_checking_assert (nops
>= 3);
1422 gcc_checking_assert (nops
>= 4);
1426 hsa_instructions
.safe_push (this);
1429 /* New operator to allocate basic instruction from pool alloc. */
1432 hsa_insn_basic::operator new (size_t)
1434 return hsa_allocp_inst_basic
->allocate_raw ();
1437 /* Verify the instruction. */
1440 hsa_insn_basic::verify ()
1442 hsa_op_address
*addr
;
1445 /* Iterate all register operands and verify that the instruction
1446 is set in uses of the register. */
1447 for (unsigned i
= 0; i
< operand_count (); i
++)
1449 hsa_op_base
*use
= get_op (i
);
1451 if ((addr
= dyn_cast
<hsa_op_address
*> (use
)) && addr
->m_reg
)
1453 gcc_assert (addr
->m_reg
->m_def_insn
!= this);
1457 if ((reg
= dyn_cast
<hsa_op_reg
*> (use
)) && !op_output_p (i
))
1460 for (j
= 0; j
< reg
->m_uses
.length (); j
++)
1462 if (reg
->m_uses
[j
] == this)
1466 if (j
== reg
->m_uses
.length ())
1468 error ("HSA instruction uses a register but is not among "
1469 "recorded register uses");
1470 debug_hsa_operand (reg
);
1471 debug_hsa_insn (this);
1472 internal_error ("HSA instruction verification failed");
1478 /* Constructor of an instruction representing a PHI node. NOPS is the number
1479 of operands (equal to the number of predecessors). */
1481 hsa_insn_phi::hsa_insn_phi (unsigned nops
, hsa_op_reg
*dst
)
1482 : hsa_insn_basic (nops
, HSA_OPCODE_PHI
), m_dest (dst
)
1484 dst
->set_definition (this);
1487 /* New operator to allocate PHI instruction from pool alloc. */
1490 hsa_insn_phi::operator new (size_t)
1492 return hsa_allocp_inst_phi
->allocate_raw ();
1495 /* Constructor of class representing instruction for conditional jump, CTRL is
1496 the control register determining whether the jump will be carried out, the
1497 new instruction is automatically added to its uses list. */
1499 hsa_insn_br::hsa_insn_br (hsa_op_reg
*ctrl
)
1500 : hsa_insn_basic (1, BRIG_OPCODE_CBR
, BRIG_TYPE_B1
, ctrl
),
1501 m_width (BRIG_WIDTH_1
)
1505 /* New operator to allocate branch instruction from pool alloc. */
1508 hsa_insn_br::operator new (size_t)
1510 return hsa_allocp_inst_br
->allocate_raw ();
1513 /* Constructor of class representing instruction for switch jump, CTRL is
1514 the index register. */
1516 hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg
*index
, unsigned jump_count
)
1517 : hsa_insn_basic (1, BRIG_OPCODE_SBR
, BRIG_TYPE_B1
, index
),
1518 m_width (BRIG_WIDTH_1
), m_jump_table (vNULL
),
1519 m_label_code_list (new hsa_op_code_list (jump_count
))
1523 /* New operator to allocate switch branch instruction from pool alloc. */
1526 hsa_insn_sbr::operator new (size_t)
1528 return hsa_allocp_inst_sbr
->allocate_raw ();
1531 /* Replace all occurrences of OLD_BB with NEW_BB in the statements
1535 hsa_insn_sbr::replace_all_labels (basic_block old_bb
, basic_block new_bb
)
1537 for (unsigned i
= 0; i
< m_jump_table
.length (); i
++)
1538 if (m_jump_table
[i
] == old_bb
)
1539 m_jump_table
[i
] = new_bb
;
1542 hsa_insn_sbr::~hsa_insn_sbr ()
1544 m_jump_table
.release ();
1547 /* Constructor of comparison instruction. CMP is the comparison operation and T
1548 is the result type. */
1550 hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp
, BrigType16_t t
,
1551 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1553 : hsa_insn_basic (3 , BRIG_OPCODE_CMP
, t
, arg0
, arg1
, arg2
), m_compare (cmp
)
1557 /* New operator to allocate compare instruction from pool alloc. */
1560 hsa_insn_cmp::operator new (size_t)
1562 return hsa_allocp_inst_cmp
->allocate_raw ();
1565 /* Constructor of classes representing memory accesses. OPC is the opcode (must
1566 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
1567 operands are provided as ARG0 and ARG1. */
1569 hsa_insn_mem::hsa_insn_mem (int opc
, BrigType16_t t
, hsa_op_base
*arg0
,
1571 : hsa_insn_basic (2, opc
, t
, arg0
, arg1
),
1572 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1574 gcc_checking_assert (opc
== BRIG_OPCODE_LD
|| opc
== BRIG_OPCODE_ST
);
1577 /* Constructor for descendants allowing different opcodes and number of
1578 operands, it passes its arguments directly to hsa_insn_basic
1579 constructor. The instruction operands are provided as ARG[0-3]. */
1582 hsa_insn_mem::hsa_insn_mem (unsigned nops
, int opc
, BrigType16_t t
,
1583 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1584 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1585 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1586 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1590 /* New operator to allocate memory instruction from pool alloc. */
1593 hsa_insn_mem::operator new (size_t)
1595 return hsa_allocp_inst_mem
->allocate_raw ();
1598 /* Constructor of class representing atomic instructions and signals. OPC is
1599 the principal opcode, aop is the specific atomic operation opcode. T is the
1600 type of the instruction. The instruction operands
1601 are provided as ARG[0-3]. */
1603 hsa_insn_atomic::hsa_insn_atomic (int nops
, int opc
,
1604 enum BrigAtomicOperation aop
,
1605 BrigType16_t t
, BrigMemoryOrder memorder
,
1607 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1609 : hsa_insn_mem (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
), m_atomicop (aop
),
1610 m_memoryorder (memorder
),
1611 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM
)
1613 gcc_checking_assert (opc
== BRIG_OPCODE_ATOMICNORET
||
1614 opc
== BRIG_OPCODE_ATOMIC
||
1615 opc
== BRIG_OPCODE_SIGNAL
||
1616 opc
== BRIG_OPCODE_SIGNALNORET
);
1619 /* New operator to allocate signal instruction from pool alloc. */
1622 hsa_insn_atomic::operator new (size_t)
1624 return hsa_allocp_inst_atomic
->allocate_raw ();
1627 /* Constructor of class representing signal instructions. OPC is the prinicpal
1628 opcode, sop is the specific signal operation opcode. T is the type of the
1629 instruction. The instruction operands are provided as ARG[0-3]. */
1631 hsa_insn_signal::hsa_insn_signal (int nops
, int opc
,
1632 enum BrigAtomicOperation sop
,
1633 BrigType16_t t
, hsa_op_base
*arg0
,
1634 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1636 : hsa_insn_atomic (nops
, opc
, sop
, t
, BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
,
1637 arg0
, arg1
, arg2
, arg3
)
1641 /* New operator to allocate signal instruction from pool alloc. */
1644 hsa_insn_signal::operator new (size_t)
1646 return hsa_allocp_inst_signal
->allocate_raw ();
1649 /* Constructor of class representing segment conversion instructions. OPC is
1650 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
1651 and SRCT are destination and source types respectively, SEG is the segment
1652 we are converting to or from. The instruction operands are
1653 provided as ARG0 and ARG1. */
1655 hsa_insn_seg::hsa_insn_seg (int opc
, BrigType16_t dest
, BrigType16_t srct
,
1656 BrigSegment8_t seg
, hsa_op_base
*arg0
,
1658 : hsa_insn_basic (2, opc
, dest
, arg0
, arg1
), m_src_type (srct
),
1661 gcc_checking_assert (opc
== BRIG_OPCODE_STOF
|| opc
== BRIG_OPCODE_FTOS
);
1664 /* New operator to allocate address conversion instruction from pool alloc. */
1667 hsa_insn_seg::operator new (size_t)
1669 return hsa_allocp_inst_seg
->allocate_raw ();
1672 /* Constructor of class representing a call instruction. CALLEE is the tree
1673 representation of the function being called. */
1675 hsa_insn_call::hsa_insn_call (tree callee
)
1676 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (callee
),
1677 m_output_arg (NULL
), m_args_code_list (NULL
), m_result_code_list (NULL
)
1681 hsa_insn_call::hsa_insn_call (hsa_internal_fn
*fn
)
1682 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (NULL
),
1683 m_called_internal_fn (fn
), m_output_arg (NULL
), m_args_code_list (NULL
),
1684 m_result_code_list (NULL
)
1688 /* New operator to allocate call instruction from pool alloc. */
1691 hsa_insn_call::operator new (size_t)
1693 return hsa_allocp_inst_call
->allocate_raw ();
1696 hsa_insn_call::~hsa_insn_call ()
1698 for (unsigned i
= 0; i
< m_input_args
.length (); i
++)
1699 delete m_input_args
[i
];
1701 delete m_output_arg
;
1703 m_input_args
.release ();
1704 m_input_arg_insns
.release ();
1707 /* Constructor of class representing the argument block required to invoke
1709 hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind
,
1710 hsa_insn_call
* call
)
1711 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK
), m_kind (brig_kind
),
1716 /* New operator to allocate argument block instruction from pool alloc. */
1719 hsa_insn_arg_block::operator new (size_t)
1721 return hsa_allocp_inst_arg_block
->allocate_raw ();
1724 hsa_insn_comment::hsa_insn_comment (const char *s
)
1725 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT
)
1727 unsigned l
= strlen (s
);
1729 /* Append '// ' to the string. */
1730 char *buf
= XNEWVEC (char, l
+ 4);
1731 sprintf (buf
, "// %s", s
);
1735 /* New operator to allocate comment instruction from pool alloc. */
1738 hsa_insn_comment::operator new (size_t)
1740 return hsa_allocp_inst_comment
->allocate_raw ();
1743 hsa_insn_comment::~hsa_insn_comment ()
1745 gcc_checking_assert (m_comment
);
1750 /* Constructor of class representing the queue instruction in HSAIL. */
1751 hsa_insn_queue::hsa_insn_queue (int nops
, BrigOpcode opcode
)
1752 : hsa_insn_basic (nops
, opcode
, BRIG_TYPE_U64
)
1756 /* New operator to allocate source type instruction from pool alloc. */
1759 hsa_insn_srctype::operator new (size_t)
1761 return hsa_allocp_inst_srctype
->allocate_raw ();
1764 /* Constructor of class representing the source type instruction in HSAIL. */
1766 hsa_insn_srctype::hsa_insn_srctype (int nops
, BrigOpcode opcode
,
1767 BrigType16_t destt
, BrigType16_t srct
,
1768 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1769 hsa_op_base
*arg2
= NULL
)
1770 : hsa_insn_basic (nops
, opcode
, destt
, arg0
, arg1
, arg2
),
1771 m_source_type (srct
)
1774 /* New operator to allocate packed instruction from pool alloc. */
1777 hsa_insn_packed::operator new (size_t)
1779 return hsa_allocp_inst_packed
->allocate_raw ();
1782 /* Constructor of class representing the packed instruction in HSAIL. */
1784 hsa_insn_packed::hsa_insn_packed (int nops
, BrigOpcode opcode
,
1785 BrigType16_t destt
, BrigType16_t srct
,
1786 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1788 : hsa_insn_srctype (nops
, opcode
, destt
, srct
, arg0
, arg1
, arg2
)
1790 m_operand_list
= new hsa_op_operand_list (nops
- 1);
1793 /* New operator to allocate convert instruction from pool alloc. */
1796 hsa_insn_cvt::operator new (size_t)
1798 return hsa_allocp_inst_cvt
->allocate_raw ();
1801 /* Constructor of class representing the convert instruction in HSAIL. */
1803 hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type
*dest
, hsa_op_with_type
*src
)
1804 : hsa_insn_basic (2, BRIG_OPCODE_CVT
, dest
->m_type
, dest
, src
)
1808 /* New operator to allocate alloca from pool alloc. */
1811 hsa_insn_alloca::operator new (size_t)
1813 return hsa_allocp_inst_alloca
->allocate_raw ();
1816 /* Constructor of class representing the alloca in HSAIL. */
1818 hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type
*dest
,
1819 hsa_op_with_type
*size
, unsigned alignment
)
1820 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA
, dest
->m_type
, dest
, size
),
1821 m_align (BRIG_ALIGNMENT_8
)
1823 gcc_assert (dest
->m_type
== BRIG_TYPE_U32
);
1825 m_align
= hsa_alignment_encoding (alignment
);
1828 /* Append an instruction INSN into the basic block. */
1831 hsa_bb::append_insn (hsa_insn_basic
*insn
)
1833 gcc_assert (insn
->m_opcode
!= 0 || insn
->operand_count () == 0);
1834 gcc_assert (!insn
->m_bb
);
1837 insn
->m_prev
= m_last_insn
;
1838 insn
->m_next
= NULL
;
1840 m_last_insn
->m_next
= insn
;
1843 m_first_insn
= insn
;
1846 /* Insert HSA instruction NEW_INSN immediately before an existing instruction
1850 hsa_insert_insn_before (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1852 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1854 if (hbb
->m_first_insn
== old_insn
)
1855 hbb
->m_first_insn
= new_insn
;
1856 new_insn
->m_prev
= old_insn
->m_prev
;
1857 new_insn
->m_next
= old_insn
;
1858 if (old_insn
->m_prev
)
1859 old_insn
->m_prev
->m_next
= new_insn
;
1860 old_insn
->m_prev
= new_insn
;
1863 /* Append HSA instruction NEW_INSN immediately after an existing instruction
1867 hsa_append_insn_after (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1869 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1871 if (hbb
->m_last_insn
== old_insn
)
1872 hbb
->m_last_insn
= new_insn
;
1873 new_insn
->m_prev
= old_insn
;
1874 new_insn
->m_next
= old_insn
->m_next
;
1875 if (old_insn
->m_next
)
1876 old_insn
->m_next
->m_prev
= new_insn
;
1877 old_insn
->m_next
= new_insn
;
1880 /* Return a register containing the calculated value of EXP which must be an
1881 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
1882 integer constants as returned by get_inner_reference.
1883 Newly generated HSA instructions will be appended to HBB.
1884 Perform all calculations in ADDRTYPE. */
1886 static hsa_op_with_type
*
1887 gen_address_calculation (tree exp
, hsa_bb
*hbb
, BrigType16_t addrtype
)
1891 if (TREE_CODE (exp
) == NOP_EXPR
)
1892 exp
= TREE_OPERAND (exp
, 0);
1894 switch (TREE_CODE (exp
))
1897 return hsa_cfun
->reg_for_gimple_ssa (exp
)->get_in_type (addrtype
, hbb
);
1901 hsa_op_immed
*imm
= new hsa_op_immed (exp
);
1902 if (addrtype
!= imm
->m_type
)
1903 imm
->m_type
= addrtype
;
1908 opcode
= BRIG_OPCODE_ADD
;
1912 opcode
= BRIG_OPCODE_MUL
;
1919 hsa_op_reg
*res
= new hsa_op_reg (addrtype
);
1920 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, addrtype
);
1921 insn
->set_op (0, res
);
1923 hsa_op_with_type
*op1
= gen_address_calculation (TREE_OPERAND (exp
, 0), hbb
,
1925 hsa_op_with_type
*op2
= gen_address_calculation (TREE_OPERAND (exp
, 1), hbb
,
1927 insn
->set_op (1, op1
);
1928 insn
->set_op (2, op2
);
1930 hbb
->append_insn (insn
);
1934 /* If R1 is NULL, just return R2, otherwise append an instruction adding them
1935 to HBB and return the register holding the result. */
1938 add_addr_regs_if_needed (hsa_op_reg
*r1
, hsa_op_reg
*r2
, hsa_bb
*hbb
)
1940 gcc_checking_assert (r2
);
1944 hsa_op_reg
*res
= new hsa_op_reg (r1
->m_type
);
1945 gcc_assert (!hsa_needs_cvt (r1
->m_type
, r2
->m_type
));
1946 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_ADD
, res
->m_type
);
1947 insn
->set_op (0, res
);
1948 insn
->set_op (1, r1
);
1949 insn
->set_op (2, r2
);
1950 hbb
->append_insn (insn
);
1954 /* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
1955 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
1958 process_mem_base (tree base
, hsa_symbol
**symbol
, BrigType16_t
*addrtype
,
1959 hsa_op_reg
**reg
, offset_int
*offset
, hsa_bb
*hbb
)
1961 if (TREE_CODE (base
) == SSA_NAME
)
1964 hsa_op_with_type
*ssa
1965 = hsa_cfun
->reg_for_gimple_ssa (base
)->get_in_type (*addrtype
, hbb
);
1966 *reg
= dyn_cast
<hsa_op_reg
*> (ssa
);
1968 else if (TREE_CODE (base
) == ADDR_EXPR
)
1970 tree decl
= TREE_OPERAND (base
, 0);
1972 if (!DECL_P (decl
) || TREE_CODE (decl
) == FUNCTION_DECL
)
1974 HSA_SORRY_AT (EXPR_LOCATION (base
),
1975 "support for HSA does not implement a memory reference "
1976 "to a non-declaration type");
1980 gcc_assert (!*symbol
);
1982 *symbol
= get_symbol_for_decl (decl
);
1983 *addrtype
= hsa_get_segment_addr_type ((*symbol
)->m_segment
);
1985 else if (TREE_CODE (base
) == INTEGER_CST
)
1986 *offset
+= wi::to_offset (base
);
1991 /* Forward declaration of a function. */
1994 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
);
1996 /* Generate HSA address operand for a given tree memory reference REF. If
1997 instructions need to be created to calculate the address, they will be added
1998 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
1999 the function assumes that the caller will handle possible
2000 bit-field references. Otherwise if we reference a bit-field, sorry message
2003 static hsa_op_address
*
2004 gen_hsa_addr (tree ref
, hsa_bb
*hbb
, HOST_WIDE_INT
*output_bitsize
= NULL
,
2005 HOST_WIDE_INT
*output_bitpos
= NULL
)
2007 hsa_symbol
*symbol
= NULL
;
2008 hsa_op_reg
*reg
= NULL
;
2009 offset_int offset
= 0;
2011 tree varoffset
= NULL_TREE
;
2012 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2013 HOST_WIDE_INT bitsize
= 0, bitpos
= 0;
2014 BrigType16_t flat_addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2016 if (TREE_CODE (ref
) == STRING_CST
)
2018 symbol
= hsa_get_string_cst_symbol (ref
);
2021 else if (TREE_CODE (ref
) == BIT_FIELD_REF
2022 && ((tree_to_uhwi (TREE_OPERAND (ref
, 1)) % BITS_PER_UNIT
) != 0
2023 || (tree_to_uhwi (TREE_OPERAND (ref
, 2)) % BITS_PER_UNIT
) != 0))
2025 HSA_SORRY_ATV (EXPR_LOCATION (origref
),
2026 "support for HSA does not implement "
2027 "bit field references such as %E", ref
);
2031 if (handled_component_p (ref
))
2033 enum machine_mode mode
;
2034 int unsignedp
, volatilep
, preversep
;
2036 ref
= get_inner_reference (ref
, &bitsize
, &bitpos
, &varoffset
, &mode
,
2037 &unsignedp
, &preversep
, &volatilep
, false);
2040 offset
= wi::rshift (offset
, LOG2_BITS_PER_UNIT
, SIGNED
);
2043 switch (TREE_CODE (ref
))
2047 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
2048 symbol
= hsa_cfun
->create_hsa_temporary (flat_addrtype
);
2049 hsa_op_reg
*r
= new hsa_op_reg (flat_addrtype
);
2050 gen_hsa_addr_insns (ref
, r
, hbb
);
2051 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
2052 r
, new hsa_op_address (symbol
)));
2058 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
2059 symbol
= hsa_cfun
->create_hsa_temporary (flat_addrtype
);
2060 hsa_op_reg
*r
= hsa_cfun
->reg_for_gimple_ssa (ref
);
2062 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
2063 r
, new hsa_op_address (symbol
)));
2070 gcc_assert (!symbol
);
2071 symbol
= get_symbol_for_decl (ref
);
2072 addrtype
= hsa_get_segment_addr_type (symbol
->m_segment
);
2076 process_mem_base (TREE_OPERAND (ref
, 0), &symbol
, &addrtype
, ®
,
2079 if (!integer_zerop (TREE_OPERAND (ref
, 1)))
2080 offset
+= wi::to_offset (TREE_OPERAND (ref
, 1));
2083 case TARGET_MEM_REF
:
2084 process_mem_base (TMR_BASE (ref
), &symbol
, &addrtype
, ®
, &offset
, hbb
);
2085 if (TMR_INDEX (ref
))
2088 hsa_op_base
*idx
= hsa_cfun
->reg_for_gimple_ssa
2089 (TMR_INDEX (ref
))->get_in_type (addrtype
, hbb
);
2090 if (TMR_STEP (ref
) && !integer_onep (TMR_STEP (ref
)))
2092 disp1
= new hsa_op_reg (addrtype
);
2093 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_MUL
,
2096 /* As step must respect addrtype, we overwrite the type
2097 of an immediate value. */
2098 hsa_op_immed
*step
= new hsa_op_immed (TMR_STEP (ref
));
2099 step
->m_type
= addrtype
;
2101 insn
->set_op (0, disp1
);
2102 insn
->set_op (1, idx
);
2103 insn
->set_op (2, step
);
2104 hbb
->append_insn (insn
);
2107 disp1
= as_a
<hsa_op_reg
*> (idx
);
2108 reg
= add_addr_regs_if_needed (reg
, disp1
, hbb
);
2110 if (TMR_INDEX2 (ref
))
2112 if (TREE_CODE (TMR_INDEX2 (ref
)) == SSA_NAME
)
2114 hsa_op_base
*disp2
= hsa_cfun
->reg_for_gimple_ssa
2115 (TMR_INDEX2 (ref
))->get_in_type (addrtype
, hbb
);
2116 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (disp2
),
2119 else if (TREE_CODE (TMR_INDEX2 (ref
)) == INTEGER_CST
)
2120 offset
+= wi::to_offset (TMR_INDEX2 (ref
));
2124 offset
+= wi::to_offset (TMR_OFFSET (ref
));
2127 HSA_SORRY_AT (EXPR_LOCATION (origref
),
2128 "support for HSA does not implement function pointers");
2131 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does "
2132 "not implement memory access to %E", origref
);
2138 if (TREE_CODE (varoffset
) == INTEGER_CST
)
2139 offset
+= wi::to_offset (varoffset
);
2142 hsa_op_base
*off_op
= gen_address_calculation (varoffset
, hbb
,
2144 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (off_op
),
2149 gcc_checking_assert ((symbol
2151 == hsa_get_segment_addr_type (symbol
->m_segment
))
2154 == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
)));
2156 HOST_WIDE_INT hwi_offset
= offset
.to_shwi ();
2158 /* Calculate remaining bitsize offset (if presented). */
2159 bitpos
%= BITS_PER_UNIT
;
2160 /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
2161 is not a reason to think this is a bit-field access. */
2163 && (bitsize
>= BITS_PER_UNIT
)
2164 && !(bitsize
& (bitsize
- 1)))
2167 if ((bitpos
|| bitsize
) && (output_bitpos
== NULL
|| output_bitsize
== NULL
))
2168 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does not "
2169 "implement unhandled bit field reference such as %E", ref
);
2171 if (output_bitsize
!= NULL
&& output_bitpos
!= NULL
)
2173 *output_bitsize
= bitsize
;
2174 *output_bitpos
= bitpos
;
2177 return new hsa_op_address (symbol
, reg
, hwi_offset
);
2180 /* Generate HSA address operand for a given tree memory reference REF. If
2181 instructions need to be created to calculate the address, they will be added
2182 to the end of HBB. OUTPUT_ALIGN is alignment of the created address. */
2184 static hsa_op_address
*
2185 gen_hsa_addr_with_align (tree ref
, hsa_bb
*hbb
, BrigAlignment8_t
*output_align
)
2187 hsa_op_address
*addr
= gen_hsa_addr (ref
, hbb
);
2188 if (addr
->m_reg
|| !addr
->m_symbol
)
2189 *output_align
= hsa_object_alignment (ref
);
2192 /* If the address consists only of a symbol and an offset, we
2193 compute the alignment ourselves to take into account any alignment
2194 promotions we might have done for the HSA symbol representation. */
2195 unsigned align
= hsa_byte_alignment (addr
->m_symbol
->m_align
);
2196 unsigned misalign
= addr
->m_imm_offset
& (align
- 1);
2198 align
= (misalign
& -misalign
);
2199 *output_align
= hsa_alignment_encoding (BITS_PER_UNIT
* align
);
2204 /* Generate HSA address for a function call argument of given TYPE.
2205 INDEX is used to generate corresponding name of the arguments.
2206 Special value -1 represents fact that result value is created. */
2208 static hsa_op_address
*
2209 gen_hsa_addr_for_arg (tree tree_type
, int index
)
2211 hsa_symbol
*sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
2213 sym
->m_type
= hsa_type_for_tree_type (tree_type
, &sym
->m_dim
);
2215 if (index
== -1) /* Function result. */
2216 sym
->m_name
= "res";
2217 else /* Function call arguments. */
2220 sym
->m_name_number
= index
;
2223 return new hsa_op_address (sym
);
2226 /* Generate HSA instructions that process all necessary conversions
2227 of an ADDR to flat addressing and place the result into DEST.
2228 Instructions are appended to HBB. */
2231 convert_addr_to_flat_segment (hsa_op_address
*addr
, hsa_op_reg
*dest
,
2234 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_LDA
);
2235 insn
->set_op (1, addr
);
2236 if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
!= BRIG_SEGMENT_GLOBAL
)
2238 /* LDA produces segment-relative address, we need to convert
2239 it to the flat one. */
2241 tmp
= new hsa_op_reg (hsa_get_segment_addr_type
2242 (addr
->m_symbol
->m_segment
));
2244 seg
= new hsa_insn_seg (BRIG_OPCODE_STOF
,
2245 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
2246 tmp
->m_type
, addr
->m_symbol
->m_segment
, dest
,
2249 insn
->set_op (0, tmp
);
2250 insn
->m_type
= tmp
->m_type
;
2251 hbb
->append_insn (insn
);
2252 hbb
->append_insn (seg
);
2256 insn
->set_op (0, dest
);
2257 insn
->m_type
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2258 hbb
->append_insn (insn
);
2262 /* Generate HSA instructions that calculate address of VAL including all
2263 necessary conversions to flat addressing and place the result into DEST.
2264 Instructions are appended to HBB. */
2267 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
)
2269 /* Handle cases like tmp = NULL, where we just emit a move instruction
2271 if (TREE_CODE (val
) == INTEGER_CST
)
2273 hsa_op_immed
*c
= new hsa_op_immed (val
);
2274 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2275 dest
->m_type
, dest
, c
);
2276 hbb
->append_insn (insn
);
2280 hsa_op_address
*addr
;
2282 gcc_assert (dest
->m_type
== hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2283 if (TREE_CODE (val
) == ADDR_EXPR
)
2284 val
= TREE_OPERAND (val
, 0);
2285 addr
= gen_hsa_addr (val
, hbb
);
2287 convert_addr_to_flat_segment (addr
, dest
, hbb
);
2290 /* Return an HSA register or HSA immediate value operand corresponding to
2291 gimple operand OP. */
2293 static hsa_op_with_type
*
2294 hsa_reg_or_immed_for_gimple_op (tree op
, hsa_bb
*hbb
)
2298 if (TREE_CODE (op
) == SSA_NAME
)
2299 tmp
= hsa_cfun
->reg_for_gimple_ssa (op
);
2300 else if (!POINTER_TYPE_P (TREE_TYPE (op
)))
2301 return new hsa_op_immed (op
);
2304 tmp
= new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2305 gen_hsa_addr_insns (op
, tmp
, hbb
);
2310 /* Create a simple movement instruction with register destination DEST and
2311 register or immediate source SRC and append it to the end of HBB. */
2314 hsa_build_append_simple_mov (hsa_op_reg
*dest
, hsa_op_base
*src
, hsa_bb
*hbb
)
2316 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
,
2318 if (hsa_op_reg
*sreg
= dyn_cast
<hsa_op_reg
*> (src
))
2319 gcc_assert (hsa_type_bit_size (dest
->m_type
)
2320 == hsa_type_bit_size (sreg
->m_type
));
2322 gcc_assert (hsa_type_bit_size (dest
->m_type
)
2323 == hsa_type_bit_size (as_a
<hsa_op_immed
*> (src
)->m_type
));
2325 hbb
->append_insn (insn
);
2328 /* Generate HSAIL instructions loading a bit field into register DEST.
2329 VALUE_REG is a register of a SSA name that is used in the bit field
2330 reference. To identify a bit field BITPOS is offset to the loaded memory
2331 and BITSIZE is number of bits of the bit field.
2332 Add instructions to HBB. */
2335 gen_hsa_insns_for_bitfield (hsa_op_reg
*dest
, hsa_op_reg
*value_reg
,
2336 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2339 unsigned type_bitsize
= hsa_type_bit_size (dest
->m_type
);
2340 unsigned left_shift
= type_bitsize
- (bitsize
+ bitpos
);
2341 unsigned right_shift
= left_shift
+ bitpos
;
2345 hsa_op_reg
*value_reg_2
= new hsa_op_reg (dest
->m_type
);
2346 hsa_op_immed
*c
= new hsa_op_immed (left_shift
, BRIG_TYPE_U32
);
2348 hsa_insn_basic
*lshift
2349 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, value_reg_2
->m_type
,
2350 value_reg_2
, value_reg
, c
);
2352 hbb
->append_insn (lshift
);
2354 value_reg
= value_reg_2
;
2359 hsa_op_reg
*value_reg_2
= new hsa_op_reg (dest
->m_type
);
2360 hsa_op_immed
*c
= new hsa_op_immed (right_shift
, BRIG_TYPE_U32
);
2362 hsa_insn_basic
*rshift
2363 = new hsa_insn_basic (3, BRIG_OPCODE_SHR
, value_reg_2
->m_type
,
2364 value_reg_2
, value_reg
, c
);
2366 hbb
->append_insn (rshift
);
2368 value_reg
= value_reg_2
;
2371 hsa_insn_basic
*assignment
2372 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, dest
, value_reg
);
2373 hbb
->append_insn (assignment
);
2377 /* Generate HSAIL instructions loading a bit field into register DEST. ADDR is
2378 prepared memory address which is used to load the bit field. To identify a
2379 bit field BITPOS is offset to the loaded memory and BITSIZE is number of
2380 bits of the bit field. Add instructions to HBB. Load must be performed in
2384 gen_hsa_insns_for_bitfield_load (hsa_op_reg
*dest
, hsa_op_address
*addr
,
2385 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2386 hsa_bb
*hbb
, BrigAlignment8_t align
)
2388 hsa_op_reg
*value_reg
= new hsa_op_reg (dest
->m_type
);
2389 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, dest
->m_type
, value_reg
,
2391 mem
->set_align (align
);
2392 hbb
->append_insn (mem
);
2393 gen_hsa_insns_for_bitfield (dest
, value_reg
, bitsize
, bitpos
, hbb
);
2396 /* Return the alignment of base memory accesses we issue to perform bit-field
2397 memory access REF. */
2399 static BrigAlignment8_t
2400 hsa_bitmemref_alignment (tree ref
)
2402 unsigned HOST_WIDE_INT bit_offset
= 0;
2406 if (TREE_CODE (ref
) == BIT_FIELD_REF
)
2408 if (!tree_fits_uhwi_p (TREE_OPERAND (ref
, 2)))
2409 return BRIG_ALIGNMENT_1
;
2410 bit_offset
+= tree_to_uhwi (TREE_OPERAND (ref
, 2));
2412 else if (TREE_CODE (ref
) == COMPONENT_REF
2413 && DECL_BIT_FIELD (TREE_OPERAND (ref
, 1)))
2414 bit_offset
+= int_bit_position (TREE_OPERAND (ref
, 1));
2417 ref
= TREE_OPERAND (ref
, 0);
2420 unsigned HOST_WIDE_INT bits
= bit_offset
% BITS_PER_UNIT
;
2421 unsigned HOST_WIDE_INT byte_bits
= bit_offset
- bits
;
2422 BrigAlignment8_t base
= hsa_object_alignment (ref
);
2425 return MIN (base
, hsa_alignment_encoding (byte_bits
& -byte_bits
));
2428 /* Generate HSAIL instructions loading something into register DEST. RHS is
2429 tree representation of the loaded data, which are loaded as type TYPE. Add
2430 instructions to HBB. */
2433 gen_hsa_insns_for_load (hsa_op_reg
*dest
, tree rhs
, tree type
, hsa_bb
*hbb
)
2435 /* The destination SSA name will give us the type. */
2436 if (TREE_CODE (rhs
) == VIEW_CONVERT_EXPR
)
2437 rhs
= TREE_OPERAND (rhs
, 0);
2439 if (TREE_CODE (rhs
) == SSA_NAME
)
2441 hsa_op_reg
*src
= hsa_cfun
->reg_for_gimple_ssa (rhs
);
2442 hsa_build_append_simple_mov (dest
, src
, hbb
);
2444 else if (is_gimple_min_invariant (rhs
)
2445 || TREE_CODE (rhs
) == ADDR_EXPR
)
2447 if (POINTER_TYPE_P (TREE_TYPE (rhs
)))
2449 if (dest
->m_type
!= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
))
2451 HSA_SORRY_ATV (EXPR_LOCATION (rhs
),
2452 "support for HSA does not implement conversion "
2453 "of %E to the requested non-pointer type.", rhs
);
2457 gen_hsa_addr_insns (rhs
, dest
, hbb
);
2459 else if (TREE_CODE (rhs
) == COMPLEX_CST
)
2461 hsa_op_immed
*real_part
= new hsa_op_immed (TREE_REALPART (rhs
));
2462 hsa_op_immed
*imag_part
= new hsa_op_immed (TREE_IMAGPART (rhs
));
2464 hsa_op_reg
*real_part_reg
2465 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2467 hsa_op_reg
*imag_part_reg
2468 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2471 hsa_build_append_simple_mov (real_part_reg
, real_part
, hbb
);
2472 hsa_build_append_simple_mov (imag_part_reg
, imag_part
, hbb
);
2474 BrigType16_t src_type
= hsa_bittype_for_type (real_part_reg
->m_type
);
2476 hsa_insn_packed
*insn
2477 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
,
2478 src_type
, dest
, real_part_reg
,
2480 hbb
->append_insn (insn
);
2484 hsa_op_immed
*imm
= new hsa_op_immed (rhs
);
2485 hsa_build_append_simple_mov (dest
, imm
, hbb
);
2488 else if (TREE_CODE (rhs
) == REALPART_EXPR
|| TREE_CODE (rhs
) == IMAGPART_EXPR
)
2490 tree pack_type
= TREE_TYPE (TREE_OPERAND (rhs
, 0));
2492 hsa_op_reg
*packed_reg
2493 = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type
, true));
2495 tree complex_rhs
= TREE_OPERAND (rhs
, 0);
2496 gen_hsa_insns_for_load (packed_reg
, complex_rhs
, TREE_TYPE (complex_rhs
),
2499 hsa_op_reg
*real_reg
2500 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2502 hsa_op_reg
*imag_reg
2503 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2505 BrigKind16_t brig_type
= packed_reg
->m_type
;
2506 hsa_insn_packed
*packed
2507 = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND
,
2508 hsa_bittype_for_type (real_reg
->m_type
),
2509 brig_type
, real_reg
, imag_reg
, packed_reg
);
2511 hbb
->append_insn (packed
);
2513 hsa_op_reg
*source
= TREE_CODE (rhs
) == REALPART_EXPR
?
2514 real_reg
: imag_reg
;
2516 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2517 dest
->m_type
, dest
, source
);
2519 hbb
->append_insn (insn
);
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 hbb
->append_insn (basic
);
2659 hsa_op_reg
*shifted_value_reg
= new hsa_op_reg (mem_type
);
2660 c
= new hsa_op_immed (bitpos
, BRIG_TYPE_U32
);
2662 hsa_insn_basic
*basic
2663 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, mem_type
,
2664 shifted_value_reg
, new_value_reg
, c
);
2665 hbb
->append_insn (basic
);
2667 new_value_reg
= shifted_value_reg
;
2670 /* OR the prepared value with prepared chunk loaded from memory. */
2671 hsa_op_reg
*prepared_reg
= new hsa_op_reg (mem_type
);
2672 basic
= new hsa_insn_basic (3, BRIG_OPCODE_OR
, mem_type
, prepared_reg
,
2673 new_value_reg
, cleared_reg
);
2674 hbb
->append_insn (basic
);
2680 req_align
= hsa_object_alignment (lhs
);
2682 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
, addr
);
2683 mem
->set_align (req_align
);
2685 /* The HSAIL verifier has another constraint: if the source is an immediate
2686 then it must match the destination type. If it's a register the low bits
2687 will be used for sub-word stores. We're always allocating new operands so
2688 we can modify the above in place. */
2689 if (hsa_op_immed
*imm
= dyn_cast
<hsa_op_immed
*> (src
))
2691 if (!hsa_type_packed_p (imm
->m_type
))
2692 imm
->m_type
= mem
->m_type
;
2695 /* ...and all vector immediates apparently need to be vectors of
2697 unsigned bs
= hsa_type_bit_size (imm
->m_type
);
2698 gcc_assert (bs
== hsa_type_bit_size (mem
->m_type
));
2702 imm
->m_type
= BRIG_TYPE_U8X4
;
2705 imm
->m_type
= BRIG_TYPE_U8X8
;
2708 imm
->m_type
= BRIG_TYPE_U8X16
;
2716 hbb
->append_insn (mem
);
2719 /* Generate memory copy instructions that are going to be used
2720 for copying a SRC memory to TARGET memory,
2721 represented by pointer in a register. MIN_ALIGN is minimal alignment
2722 of provided HSA addresses. */
2725 gen_hsa_memory_copy (hsa_bb
*hbb
, hsa_op_address
*target
, hsa_op_address
*src
,
2726 unsigned size
, BrigAlignment8_t min_align
)
2728 hsa_op_address
*addr
;
2731 unsigned offset
= 0;
2732 unsigned min_byte_align
= hsa_byte_alignment (min_align
);
2746 if (s
> min_byte_align
)
2749 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2751 hsa_op_reg
*tmp
= new hsa_op_reg (t
);
2752 addr
= new hsa_op_address (src
->m_symbol
, src
->m_reg
,
2753 src
->m_imm_offset
+ offset
);
2754 mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, t
, tmp
, addr
);
2755 hbb
->append_insn (mem
);
2757 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2758 target
->m_imm_offset
+ offset
);
2759 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, tmp
, addr
);
2760 hbb
->append_insn (mem
);
2766 /* Create a memset mask that is created by copying a CONSTANT byte value
2767 to an integer of BYTE_SIZE bytes. */
2769 static unsigned HOST_WIDE_INT
2770 build_memset_value (unsigned HOST_WIDE_INT constant
, unsigned byte_size
)
2775 HOST_WIDE_INT v
= constant
;
2777 for (unsigned i
= 1; i
< byte_size
; i
++)
2778 v
|= constant
<< (8 * i
);
2783 /* Generate memory set instructions that are going to be used
2784 for setting a CONSTANT byte value to TARGET memory of SIZE bytes.
2785 MIN_ALIGN is minimal alignment of provided HSA addresses. */
2788 gen_hsa_memory_set (hsa_bb
*hbb
, hsa_op_address
*target
,
2789 unsigned HOST_WIDE_INT constant
,
2790 unsigned size
, BrigAlignment8_t min_align
)
2792 hsa_op_address
*addr
;
2795 unsigned offset
= 0;
2796 unsigned min_byte_align
= hsa_byte_alignment (min_align
);
2810 if (s
> min_byte_align
)
2813 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2814 target
->m_imm_offset
+ offset
);
2816 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2817 HOST_WIDE_INT c
= build_memset_value (constant
, s
);
2819 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, new hsa_op_immed (c
, t
),
2821 hbb
->append_insn (mem
);
2827 /* Generate HSAIL instructions for a single assignment
2828 of an empty constructor to an ADDR_LHS. Constructor is passed as a
2829 tree RHS and all instructions are appended to HBB. ALIGN is
2830 alignment of the address. */
2833 gen_hsa_ctor_assignment (hsa_op_address
*addr_lhs
, tree rhs
, hsa_bb
*hbb
,
2834 BrigAlignment8_t align
)
2836 if (vec_safe_length (CONSTRUCTOR_ELTS (rhs
)))
2838 HSA_SORRY_AT (EXPR_LOCATION (rhs
),
2839 "support for HSA does not implement load from constructor");
2843 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2844 gen_hsa_memory_set (hbb
, addr_lhs
, 0, size
, align
);
2847 /* Generate HSA instructions for a single assignment of RHS to LHS.
2848 HBB is the basic block they will be appended to. */
2851 gen_hsa_insns_for_single_assignment (tree lhs
, tree rhs
, hsa_bb
*hbb
)
2853 if (TREE_CODE (lhs
) == SSA_NAME
)
2855 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
2856 if (hsa_seen_error ())
2859 gen_hsa_insns_for_load (dest
, rhs
, TREE_TYPE (lhs
), hbb
);
2861 else if (TREE_CODE (rhs
) == SSA_NAME
2862 || (is_gimple_min_invariant (rhs
) && TREE_CODE (rhs
) != STRING_CST
))
2864 /* Store to memory. */
2865 hsa_op_base
*src
= hsa_reg_or_immed_for_gimple_op (rhs
, hbb
);
2866 if (hsa_seen_error ())
2869 gen_hsa_insns_for_store (lhs
, src
, hbb
);
2873 BrigAlignment8_t lhs_align
;
2874 hsa_op_address
*addr_lhs
= gen_hsa_addr_with_align (lhs
, hbb
,
2877 if (TREE_CODE (rhs
) == CONSTRUCTOR
)
2878 gen_hsa_ctor_assignment (addr_lhs
, rhs
, hbb
, lhs_align
);
2881 BrigAlignment8_t rhs_align
;
2882 hsa_op_address
*addr_rhs
= gen_hsa_addr_with_align (rhs
, hbb
,
2885 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2886 gen_hsa_memory_copy (hbb
, addr_lhs
, addr_rhs
, size
,
2887 MIN (lhs_align
, rhs_align
));
2892 /* Prepend before INSN a load from spill symbol of SPILL_REG. Return the
2893 register into which we loaded. If this required another register to convert
2894 from a B1 type, return it in *PTMP2, otherwise store NULL into it. We
2895 assume we are out of SSA so the returned register does not have its
2899 hsa_spill_in (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2901 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2902 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2903 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2905 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, spill_sym
->m_type
,
2907 hsa_insert_insn_before (mem
, insn
);
2910 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2912 hsa_insn_basic
*cvtinsn
;
2914 reg
= new hsa_op_reg (spill_reg
->m_type
);
2916 cvtinsn
= new hsa_insn_cvt (reg
, *ptmp2
);
2917 hsa_insert_insn_before (cvtinsn
, insn
);
2922 /* Append after INSN a store to spill symbol of SPILL_REG. Return the register
2923 from which we stored. If this required another register to convert to a B1
2924 type, return it in *PTMP2, otherwise store NULL into it. We assume we are
2925 out of SSA so the returned register does not have its use updated. */
2928 hsa_spill_out (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2930 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2931 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2932 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2933 hsa_op_reg
*returnreg
;
2937 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2939 hsa_insn_basic
*cvtinsn
;
2940 *ptmp2
= new hsa_op_reg (spill_sym
->m_type
);
2941 reg
->m_type
= spill_reg
->m_type
;
2943 cvtinsn
= new hsa_insn_cvt (*ptmp2
, returnreg
);
2944 hsa_append_insn_after (cvtinsn
, insn
);
2949 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, spill_sym
->m_type
, reg
,
2951 hsa_append_insn_after (mem
, insn
);
2955 /* Generate a comparison instruction that will compare LHS and RHS with
2956 comparison specified by CODE and put result into register DEST. DEST has to
2957 have its type set already but must not have its definition set yet.
2958 Generated instructions will be added to HBB. */
2961 gen_hsa_cmp_insn_from_gimple (enum tree_code code
, tree lhs
, tree rhs
,
2962 hsa_op_reg
*dest
, hsa_bb
*hbb
)
2964 BrigCompareOperation8_t compare
;
2969 compare
= BRIG_COMPARE_LT
;
2972 compare
= BRIG_COMPARE_LE
;
2975 compare
= BRIG_COMPARE_GT
;
2978 compare
= BRIG_COMPARE_GE
;
2981 compare
= BRIG_COMPARE_EQ
;
2984 compare
= BRIG_COMPARE_NE
;
2986 case UNORDERED_EXPR
:
2987 compare
= BRIG_COMPARE_NAN
;
2990 compare
= BRIG_COMPARE_NUM
;
2993 compare
= BRIG_COMPARE_LTU
;
2996 compare
= BRIG_COMPARE_LEU
;
2999 compare
= BRIG_COMPARE_GTU
;
3002 compare
= BRIG_COMPARE_GEU
;
3005 compare
= BRIG_COMPARE_EQU
;
3008 compare
= BRIG_COMPARE_NEU
;
3012 HSA_SORRY_ATV (EXPR_LOCATION (lhs
),
3013 "support for HSA does not implement comparison tree "
3014 "code %s\n", get_tree_code_name (code
));
3018 /* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer)
3019 as a result of comparison. */
3021 BrigType16_t dest_type
= hsa_type_integer_p (dest
->m_type
)
3022 ? (BrigType16_t
) BRIG_TYPE_B1
: dest
->m_type
;
3024 hsa_insn_cmp
*cmp
= new hsa_insn_cmp (compare
, dest_type
);
3025 cmp
->set_op (1, hsa_reg_or_immed_for_gimple_op (lhs
, hbb
));
3026 cmp
->set_op (2, hsa_reg_or_immed_for_gimple_op (rhs
, hbb
));
3028 hbb
->append_insn (cmp
);
3029 cmp
->set_output_in_type (dest
, 0, hbb
);
3032 /* Generate an unary instruction with OPCODE and append it to a basic block
3033 HBB. The instruction uses DEST as a destination and OP1
3034 as a single operand. */
3037 gen_hsa_unary_operation (BrigOpcode opcode
, hsa_op_reg
*dest
,
3038 hsa_op_with_type
*op1
, hsa_bb
*hbb
)
3040 gcc_checking_assert (dest
);
3041 hsa_insn_basic
*insn
;
3043 if (opcode
== BRIG_OPCODE_MOV
&& hsa_needs_cvt (dest
->m_type
, op1
->m_type
))
3044 insn
= new hsa_insn_cvt (dest
, op1
);
3045 else if (opcode
== BRIG_OPCODE_FIRSTBIT
|| opcode
== BRIG_OPCODE_LASTBIT
)
3046 insn
= new hsa_insn_srctype (2, opcode
, BRIG_TYPE_U32
, op1
->m_type
, NULL
,
3050 insn
= new hsa_insn_basic (2, opcode
, dest
->m_type
, dest
, op1
);
3052 if (opcode
== BRIG_OPCODE_ABS
|| opcode
== BRIG_OPCODE_NEG
)
3054 /* ABS and NEG only exist in _s form :-/ */
3055 if (insn
->m_type
== BRIG_TYPE_U32
)
3056 insn
->m_type
= BRIG_TYPE_S32
;
3057 else if (insn
->m_type
== BRIG_TYPE_U64
)
3058 insn
->m_type
= BRIG_TYPE_S64
;
3062 hbb
->append_insn (insn
);
3064 if (opcode
== BRIG_OPCODE_FIRSTBIT
|| opcode
== BRIG_OPCODE_LASTBIT
)
3065 insn
->set_output_in_type (dest
, 0, hbb
);
3068 /* Generate a binary instruction with OPCODE and append it to a basic block
3069 HBB. The instruction uses DEST as a destination and operands OP1
3073 gen_hsa_binary_operation (int opcode
, hsa_op_reg
*dest
,
3074 hsa_op_base
*op1
, hsa_op_base
*op2
, hsa_bb
*hbb
)
3076 gcc_checking_assert (dest
);
3078 if ((opcode
== BRIG_OPCODE_SHL
|| opcode
== BRIG_OPCODE_SHR
)
3079 && is_a
<hsa_op_immed
*> (op2
))
3081 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3082 i
->set_type (BRIG_TYPE_U32
);
3084 if ((opcode
== BRIG_OPCODE_OR
3085 || opcode
== BRIG_OPCODE_XOR
3086 || opcode
== BRIG_OPCODE_AND
)
3087 && is_a
<hsa_op_immed
*> (op2
))
3089 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3090 i
->set_type (hsa_unsigned_type_for_type (i
->m_type
));
3093 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, dest
->m_type
, dest
,
3095 hbb
->append_insn (insn
);
3098 /* Generate HSA instructions for a single assignment. HBB is the basic block
3099 they will be appended to. */
3102 gen_hsa_insns_for_operation_assignment (gimple
*assign
, hsa_bb
*hbb
)
3104 tree_code code
= gimple_assign_rhs_code (assign
);
3105 gimple_rhs_class rhs_class
= get_gimple_rhs_class (gimple_expr_code (assign
));
3107 tree lhs
= gimple_assign_lhs (assign
);
3108 tree rhs1
= gimple_assign_rhs1 (assign
);
3109 tree rhs2
= gimple_assign_rhs2 (assign
);
3110 tree rhs3
= gimple_assign_rhs3 (assign
);
3118 /* The opcode is changed to BRIG_OPCODE_CVT if BRIG types
3119 needs a conversion. */
3120 opcode
= BRIG_OPCODE_MOV
;
3124 case POINTER_PLUS_EXPR
:
3125 opcode
= BRIG_OPCODE_ADD
;
3128 opcode
= BRIG_OPCODE_SUB
;
3131 opcode
= BRIG_OPCODE_MUL
;
3133 case MULT_HIGHPART_EXPR
:
3134 opcode
= BRIG_OPCODE_MULHI
;
3137 case TRUNC_DIV_EXPR
:
3138 case EXACT_DIV_EXPR
:
3139 opcode
= BRIG_OPCODE_DIV
;
3142 case FLOOR_DIV_EXPR
:
3143 case ROUND_DIV_EXPR
:
3144 HSA_SORRY_AT (gimple_location (assign
),
3145 "support for HSA does not implement CEIL_DIV_EXPR, "
3146 "FLOOR_DIV_EXPR or ROUND_DIV_EXPR");
3148 case TRUNC_MOD_EXPR
:
3149 opcode
= BRIG_OPCODE_REM
;
3152 case FLOOR_MOD_EXPR
:
3153 case ROUND_MOD_EXPR
:
3154 HSA_SORRY_AT (gimple_location (assign
),
3155 "support for HSA does not implement CEIL_MOD_EXPR, "
3156 "FLOOR_MOD_EXPR or ROUND_MOD_EXPR");
3159 opcode
= BRIG_OPCODE_NEG
;
3162 opcode
= BRIG_OPCODE_MIN
;
3165 opcode
= BRIG_OPCODE_MAX
;
3168 opcode
= BRIG_OPCODE_ABS
;
3171 opcode
= BRIG_OPCODE_SHL
;
3174 opcode
= BRIG_OPCODE_SHR
;
3179 hsa_insn_basic
*insn
= NULL
;
3180 int code1
= code
== LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3181 int code2
= code
!= LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3182 BrigType16_t btype
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
3185 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3186 hsa_op_reg
*op1
= new hsa_op_reg (btype
);
3187 hsa_op_reg
*op2
= new hsa_op_reg (btype
);
3188 hsa_op_with_type
*shift1
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3190 tree type
= TREE_TYPE (rhs2
);
3191 unsigned HOST_WIDE_INT bitsize
= TREE_INT_CST_LOW (TYPE_SIZE (type
));
3193 hsa_op_with_type
*shift2
= NULL
;
3194 if (TREE_CODE (rhs2
) == INTEGER_CST
)
3195 shift2
= new hsa_op_immed (bitsize
- tree_to_uhwi (rhs2
),
3197 else if (TREE_CODE (rhs2
) == SSA_NAME
)
3199 hsa_op_reg
*s
= hsa_cfun
->reg_for_gimple_ssa (rhs2
);
3200 hsa_op_reg
*d
= new hsa_op_reg (s
->m_type
);
3201 hsa_op_immed
*size_imm
= new hsa_op_immed (bitsize
, BRIG_TYPE_U32
);
3203 insn
= new hsa_insn_basic (3, BRIG_OPCODE_SUB
, d
->m_type
,
3205 hbb
->append_insn (insn
);
3212 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3213 gen_hsa_binary_operation (code1
, op1
, src
, shift1
, hbb
);
3214 gen_hsa_binary_operation (code2
, op2
, src
, shift2
, hbb
);
3215 gen_hsa_binary_operation (BRIG_OPCODE_OR
, dest
, op1
, op2
, hbb
);
3220 opcode
= BRIG_OPCODE_OR
;
3223 opcode
= BRIG_OPCODE_XOR
;
3226 opcode
= BRIG_OPCODE_AND
;
3229 opcode
= BRIG_OPCODE_NOT
;
3231 case FIX_TRUNC_EXPR
:
3233 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3234 hsa_op_with_type
*v
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3236 if (hsa_needs_cvt (dest
->m_type
, v
->m_type
))
3238 hsa_op_reg
*tmp
= new hsa_op_reg (v
->m_type
);
3240 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3241 tmp
->m_type
, tmp
, v
);
3242 hbb
->append_insn (insn
);
3244 hsa_insn_basic
*cvtinsn
= new hsa_insn_cvt (dest
, tmp
);
3245 hbb
->append_insn (cvtinsn
);
3249 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3250 dest
->m_type
, dest
, v
);
3251 hbb
->append_insn (insn
);
3256 opcode
= BRIG_OPCODE_TRUNC
;
3265 case UNORDERED_EXPR
:
3275 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3277 gen_hsa_cmp_insn_from_gimple (code
, rhs1
, rhs2
, dest
, hbb
);
3283 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3284 hsa_op_with_type
*ctrl
= NULL
;
3287 if (CONSTANT_CLASS_P (cond
) || TREE_CODE (cond
) == SSA_NAME
)
3288 ctrl
= hsa_reg_or_immed_for_gimple_op (cond
, hbb
);
3291 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
3293 gen_hsa_cmp_insn_from_gimple (TREE_CODE (cond
),
3294 TREE_OPERAND (cond
, 0),
3295 TREE_OPERAND (cond
, 1),
3301 hsa_op_with_type
*op2
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3302 hsa_op_with_type
*op3
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3304 BrigType16_t utype
= hsa_unsigned_type_for_type (dest
->m_type
);
3305 if (is_a
<hsa_op_immed
*> (op2
))
3306 op2
->m_type
= utype
;
3307 if (is_a
<hsa_op_immed
*> (op3
))
3308 op3
->m_type
= utype
;
3310 hsa_insn_basic
*insn
3311 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
,
3312 hsa_bittype_for_type (dest
->m_type
),
3313 dest
, ctrl
, op2
, op3
);
3315 hbb
->append_insn (insn
);
3321 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3322 hsa_op_with_type
*rhs1_reg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3323 hsa_op_with_type
*rhs2_reg
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3325 if (hsa_seen_error ())
3328 BrigType16_t src_type
= hsa_bittype_for_type (rhs1_reg
->m_type
);
3329 rhs1_reg
= rhs1_reg
->get_in_type (src_type
, hbb
);
3330 rhs2_reg
= rhs2_reg
->get_in_type (src_type
, hbb
);
3332 hsa_insn_packed
*insn
3333 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
, src_type
,
3334 dest
, rhs1_reg
, rhs2_reg
);
3335 hbb
->append_insn (insn
);
3340 /* Implement others as we come across them. */
3341 HSA_SORRY_ATV (gimple_location (assign
),
3342 "support for HSA does not implement operation %s",
3343 get_tree_code_name (code
));
3348 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3350 hsa_op_with_type
*op1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3351 hsa_op_with_type
*op2
= rhs2
!= NULL_TREE
?
3352 hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
) : NULL
;
3354 if (hsa_seen_error ())
3359 case GIMPLE_TERNARY_RHS
:
3364 case GIMPLE_BINARY_RHS
:
3365 gen_hsa_binary_operation (opcode
, dest
, op1
, op2
, hbb
);
3368 case GIMPLE_UNARY_RHS
:
3369 gen_hsa_unary_operation (opcode
, dest
, op1
, hbb
);
3376 /* Generate HSA instructions for a given gimple condition statement COND.
3377 Instructions will be appended to HBB, which also needs to be the
3378 corresponding structure to the basic_block of COND. */
3381 gen_hsa_insns_for_cond_stmt (gimple
*cond
, hsa_bb
*hbb
)
3383 hsa_op_reg
*ctrl
= new hsa_op_reg (BRIG_TYPE_B1
);
3386 gen_hsa_cmp_insn_from_gimple (gimple_cond_code (cond
),
3387 gimple_cond_lhs (cond
),
3388 gimple_cond_rhs (cond
),
3391 cbr
= new hsa_insn_br (ctrl
);
3392 hbb
->append_insn (cbr
);
3395 /* Maximum number of elements in a jump table for an HSA SBR instruction. */
3397 #define HSA_MAXIMUM_SBR_LABELS 16
3399 /* Return lowest value of a switch S that is handled in a non-default
3403 get_switch_low (gswitch
*s
)
3405 unsigned labels
= gimple_switch_num_labels (s
);
3406 gcc_checking_assert (labels
>= 1);
3408 return CASE_LOW (gimple_switch_label (s
, 1));
3411 /* Return highest value of a switch S that is handled in a non-default
3415 get_switch_high (gswitch
*s
)
3417 unsigned labels
= gimple_switch_num_labels (s
);
3419 /* Compare last label to maximum number of labels. */
3420 tree label
= gimple_switch_label (s
, labels
- 1);
3421 tree low
= CASE_LOW (label
);
3422 tree high
= CASE_HIGH (label
);
3424 return high
!= NULL_TREE
? high
: low
;
3428 get_switch_size (gswitch
*s
)
3430 return int_const_binop (MINUS_EXPR
, get_switch_high (s
), get_switch_low (s
));
3433 /* Generate HSA instructions for a given gimple switch.
3434 Instructions will be appended to HBB. */
3437 gen_hsa_insns_for_switch_stmt (gswitch
*s
, hsa_bb
*hbb
)
3439 gimple_stmt_iterator it
= gsi_for_stmt (s
);
3442 /* Create preambule that verifies that index - lowest_label >= 0. */
3443 edge e
= split_block (hbb
->m_bb
, gsi_stmt (it
));
3444 e
->flags
&= ~EDGE_FALLTHRU
;
3445 e
->flags
|= EDGE_TRUE_VALUE
;
3447 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
3448 tree index_tree
= gimple_switch_index (s
);
3449 tree lowest
= get_switch_low (s
);
3450 tree highest
= get_switch_high (s
);
3452 hsa_op_reg
*index
= hsa_cfun
->reg_for_gimple_ssa (index_tree
);
3454 hsa_op_reg
*cmp1_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3455 hsa_op_immed
*cmp1_immed
= new hsa_op_immed (lowest
);
3456 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_GE
, cmp1_reg
->m_type
,
3457 cmp1_reg
, index
, cmp1_immed
));
3459 hsa_op_reg
*cmp2_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3460 hsa_op_immed
*cmp2_immed
= new hsa_op_immed (highest
);
3461 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_LE
, cmp2_reg
->m_type
,
3462 cmp2_reg
, index
, cmp2_immed
));
3464 hsa_op_reg
*cmp_reg
= new hsa_op_reg (BRIG_TYPE_B1
);
3465 hbb
->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_AND
, cmp_reg
->m_type
,
3466 cmp_reg
, cmp1_reg
, cmp2_reg
));
3468 hbb
->append_insn (new hsa_insn_br (cmp_reg
));
3470 tree default_label
= gimple_switch_default_label (s
);
3471 basic_block default_label_bb
= label_to_block_fn (func
,
3472 CASE_LABEL (default_label
));
3474 make_edge (e
->src
, default_label_bb
, EDGE_FALSE_VALUE
);
3476 hsa_cfun
->m_modified_cfg
= true;
3478 /* Basic block with the SBR instruction. */
3479 hbb
= hsa_init_new_bb (e
->dest
);
3481 hsa_op_reg
*sub_index
= new hsa_op_reg (index
->m_type
);
3482 hbb
->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB
, sub_index
->m_type
,
3484 new hsa_op_immed (lowest
)));
3486 hsa_op_base
*tmp
= sub_index
->get_in_type (BRIG_TYPE_U64
, hbb
);
3487 sub_index
= as_a
<hsa_op_reg
*> (tmp
);
3488 unsigned labels
= gimple_switch_num_labels (s
);
3489 unsigned HOST_WIDE_INT size
= tree_to_uhwi (get_switch_size (s
));
3491 hsa_insn_sbr
*sbr
= new hsa_insn_sbr (sub_index
, size
+ 1);
3493 /* Prepare array with default label destination. */
3494 for (unsigned HOST_WIDE_INT i
= 0; i
<= size
; i
++)
3495 sbr
->m_jump_table
.safe_push (default_label_bb
);
3497 /* Iterate all labels and fill up the jump table. */
3498 for (unsigned i
= 1; i
< labels
; i
++)
3500 tree label
= gimple_switch_label (s
, i
);
3501 basic_block bb
= label_to_block_fn (func
, CASE_LABEL (label
));
3503 unsigned HOST_WIDE_INT sub_low
3504 = tree_to_uhwi (int_const_binop (MINUS_EXPR
, CASE_LOW (label
), lowest
));
3506 unsigned HOST_WIDE_INT sub_high
= sub_low
;
3507 tree high
= CASE_HIGH (label
);
3509 sub_high
= tree_to_uhwi (int_const_binop (MINUS_EXPR
, high
, lowest
));
3511 for (unsigned HOST_WIDE_INT j
= sub_low
; j
<= sub_high
; j
++)
3512 sbr
->m_jump_table
[j
] = bb
;
3515 hbb
->append_insn (sbr
);
3518 /* Verify that the function DECL can be handled by HSA. */
3521 verify_function_arguments (tree decl
)
3523 if (DECL_STATIC_CHAIN (decl
))
3525 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3526 "HSA does not support nested functions: %D", decl
);
3529 else if (!TYPE_ARG_TYPES (TREE_TYPE (decl
)))
3531 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3532 "HSA does not support functions with variadic arguments "
3533 "(or unknown return type): %D", decl
);
3538 /* Return BRIG type for FORMAL_ARG_TYPE. If the formal argument type is NULL,
3539 return ACTUAL_ARG_TYPE. */
3542 get_format_argument_type (tree formal_arg_type
, BrigType16_t actual_arg_type
)
3544 if (formal_arg_type
== NULL
)
3545 return actual_arg_type
;
3547 BrigType16_t decl_type
3548 = hsa_type_for_scalar_tree_type (formal_arg_type
, false);
3549 return mem_type_for_type (decl_type
);
3552 /* Generate HSA instructions for a direct call instruction.
3553 Instructions will be appended to HBB, which also needs to be the
3554 corresponding structure to the basic_block of STMT.
3555 If ASSIGN_LHS is false, do not copy HSA function result argument into the
3556 corresponding HSA representation of the gimple statement LHS. */
3559 gen_hsa_insns_for_direct_call (gimple
*stmt
, hsa_bb
*hbb
,
3560 bool assign_lhs
= true)
3562 tree decl
= gimple_call_fndecl (stmt
);
3563 verify_function_arguments (decl
);
3564 if (hsa_seen_error ())
3567 hsa_insn_call
*call_insn
= new hsa_insn_call (decl
);
3568 hsa_cfun
->m_called_functions
.safe_push (call_insn
->m_called_function
);
3570 /* Argument block start. */
3571 hsa_insn_arg_block
*arg_start
3572 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3573 hbb
->append_insn (arg_start
);
3575 tree parm_type_chain
= TYPE_ARG_TYPES (gimple_call_fntype (stmt
));
3577 /* Preparation of arguments that will be passed to function. */
3578 const unsigned args
= gimple_call_num_args (stmt
);
3579 for (unsigned i
= 0; i
< args
; ++i
)
3581 tree parm
= gimple_call_arg (stmt
, (int)i
);
3582 tree parm_decl_type
= parm_type_chain
!= NULL_TREE
3583 ? TREE_VALUE (parm_type_chain
) : NULL_TREE
;
3584 hsa_op_address
*addr
;
3586 if (AGGREGATE_TYPE_P (TREE_TYPE (parm
)))
3588 addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3589 BrigAlignment8_t align
;
3590 hsa_op_address
*src
= gen_hsa_addr_with_align (parm
, hbb
, &align
);
3591 gen_hsa_memory_copy (hbb
, addr
, src
,
3592 addr
->m_symbol
->total_byte_size (), align
);
3596 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3598 if (parm_decl_type
!= NULL
&& AGGREGATE_TYPE_P (parm_decl_type
))
3600 HSA_SORRY_AT (gimple_location (stmt
),
3601 "support for HSA does not implement an aggregate "
3602 "formal argument in a function call, while actual "
3603 "argument is not an aggregate");
3607 BrigType16_t formal_arg_type
3608 = get_format_argument_type (parm_decl_type
, src
->m_type
);
3609 if (hsa_seen_error ())
3612 if (src
->m_type
!= formal_arg_type
)
3613 src
= src
->get_in_type (formal_arg_type
, hbb
);
3616 = gen_hsa_addr_for_arg (parm_decl_type
!= NULL_TREE
?
3617 parm_decl_type
: TREE_TYPE (parm
), i
);
3618 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, formal_arg_type
,
3621 hbb
->append_insn (mem
);
3624 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3625 if (parm_type_chain
)
3626 parm_type_chain
= TREE_CHAIN (parm_type_chain
);
3629 call_insn
->m_args_code_list
= new hsa_op_code_list (args
);
3630 hbb
->append_insn (call_insn
);
3632 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
3634 tree result
= gimple_call_lhs (stmt
);
3635 hsa_insn_mem
*result_insn
= NULL
;
3636 if (!VOID_TYPE_P (result_type
))
3638 hsa_op_address
*addr
= gen_hsa_addr_for_arg (result_type
, -1);
3640 /* Even if result of a function call is unused, we have to emit
3641 declaration for the result. */
3642 if (result
&& assign_lhs
)
3644 tree lhs_type
= TREE_TYPE (result
);
3646 if (hsa_seen_error ())
3649 if (AGGREGATE_TYPE_P (lhs_type
))
3651 BrigAlignment8_t align
;
3652 hsa_op_address
*result_addr
3653 = gen_hsa_addr_with_align (result
, hbb
, &align
);
3654 gen_hsa_memory_copy (hbb
, result_addr
, addr
,
3655 addr
->m_symbol
->total_byte_size (), align
);
3660 = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type
,
3663 hsa_op_reg
*dst
= hsa_cfun
->reg_for_gimple_ssa (result
);
3664 result_insn
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
, dst
, addr
);
3665 hbb
->append_insn (result_insn
);
3669 call_insn
->m_output_arg
= addr
->m_symbol
;
3670 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3676 HSA_SORRY_AT (gimple_location (stmt
),
3677 "support for HSA does not implement an assignment of "
3678 "return value from a void function");
3682 call_insn
->m_result_code_list
= new hsa_op_code_list (0);
3685 /* Argument block end. */
3686 hsa_insn_arg_block
*arg_end
3687 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3688 hbb
->append_insn (arg_end
);
3691 /* Generate HSA instructions for a direct call of an internal fn.
3692 Instructions will be appended to HBB, which also needs to be the
3693 corresponding structure to the basic_block of STMT. */
3696 gen_hsa_insns_for_call_of_internal_fn (gimple
*stmt
, hsa_bb
*hbb
)
3698 tree lhs
= gimple_call_lhs (stmt
);
3702 tree lhs_type
= TREE_TYPE (lhs
);
3703 tree rhs1
= gimple_call_arg (stmt
, 0);
3704 tree rhs1_type
= TREE_TYPE (rhs1
);
3705 enum internal_fn fn
= gimple_call_internal_fn (stmt
);
3706 hsa_internal_fn
*ifn
3707 = new hsa_internal_fn (fn
, tree_to_uhwi (TYPE_SIZE (rhs1_type
)));
3708 hsa_insn_call
*call_insn
= new hsa_insn_call (ifn
);
3710 gcc_checking_assert (FLOAT_TYPE_P (rhs1_type
));
3712 if (!hsa_emitted_internal_decls
->find (call_insn
->m_called_internal_fn
))
3713 hsa_cfun
->m_called_internal_fns
.safe_push (call_insn
->m_called_internal_fn
);
3715 hsa_insn_arg_block
*arg_start
3716 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3717 hbb
->append_insn (arg_start
);
3719 unsigned num_args
= gimple_call_num_args (stmt
);
3721 /* Function arguments. */
3722 for (unsigned i
= 0; i
< num_args
; i
++)
3724 tree parm
= gimple_call_arg (stmt
, (int)i
);
3725 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3727 hsa_op_address
*addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3728 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, src
->m_type
,
3731 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3732 hbb
->append_insn (mem
);
3735 call_insn
->m_args_code_list
= new hsa_op_code_list (num_args
);
3736 hbb
->append_insn (call_insn
);
3738 /* Assign returned value. */
3739 hsa_op_address
*addr
= gen_hsa_addr_for_arg (lhs_type
, -1);
3741 call_insn
->m_output_arg
= addr
->m_symbol
;
3742 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3744 /* Argument block end. */
3745 hsa_insn_arg_block
*arg_end
3746 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3747 hbb
->append_insn (arg_end
);
3750 /* Generate HSA instructions for a return value instruction.
3751 Instructions will be appended to HBB, which also needs to be the
3752 corresponding structure to the basic_block of STMT. */
3755 gen_hsa_insns_for_return (greturn
*stmt
, hsa_bb
*hbb
)
3757 tree retval
= gimple_return_retval (stmt
);
3760 hsa_op_address
*addr
= new hsa_op_address (hsa_cfun
->m_output_arg
);
3762 if (AGGREGATE_TYPE_P (TREE_TYPE (retval
)))
3764 BrigAlignment8_t align
;
3765 hsa_op_address
*retval_addr
= gen_hsa_addr_with_align (retval
, hbb
,
3767 gen_hsa_memory_copy (hbb
, addr
, retval_addr
,
3768 hsa_cfun
->m_output_arg
->total_byte_size (),
3773 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (retval
),
3775 BrigType16_t mtype
= mem_type_for_type (t
);
3777 /* Store of return value. */
3778 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (retval
, hbb
);
3779 src
= src
->get_in_type (mtype
, hbb
);
3780 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
,
3782 hbb
->append_insn (mem
);
3786 /* HSAIL return instruction emission. */
3787 hsa_insn_basic
*ret
= new hsa_insn_basic (0, BRIG_OPCODE_RET
);
3788 hbb
->append_insn (ret
);
3791 /* Set OP_INDEX-th operand of the instruction to DEST, as the DEST
3792 can have a different type, conversion instructions are possibly
3796 hsa_insn_basic::set_output_in_type (hsa_op_reg
*dest
, unsigned op_index
,
3799 hsa_insn_basic
*insn
;
3800 gcc_checking_assert (op_output_p (op_index
));
3802 if (dest
->m_type
== m_type
)
3804 set_op (op_index
, dest
);
3808 hsa_op_reg
*tmp
= new hsa_op_reg (m_type
);
3809 set_op (op_index
, tmp
);
3811 if (hsa_needs_cvt (dest
->m_type
, m_type
))
3812 insn
= new hsa_insn_cvt (dest
, tmp
);
3814 insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
,
3815 dest
, tmp
->get_in_type (dest
->m_type
, hbb
));
3817 hbb
->append_insn (insn
);
3820 /* Generate instruction OPCODE to query a property of HSA grid along the
3821 given DIMENSION. Store result into DEST and append the instruction to
3825 query_hsa_grid (hsa_op_reg
*dest
, BrigType16_t opcode
, int dimension
,
3828 /* We're using just one-dimensional kernels, so hard-coded
3831 = new hsa_op_immed (dimension
, (BrigKind16_t
) BRIG_TYPE_U32
);
3832 hsa_insn_basic
*insn
= new hsa_insn_basic (2, opcode
, BRIG_TYPE_U32
, NULL
,
3834 hbb
->append_insn (insn
);
3835 insn
->set_output_in_type (dest
, 0, hbb
);
3838 /* Generate a special HSA-related instruction for gimple STMT.
3839 Instructions are appended to basic block HBB. */
3842 query_hsa_grid (gimple
*stmt
, BrigOpcode16_t opcode
, int dimension
,
3845 tree lhs
= gimple_call_lhs (dyn_cast
<gcall
*> (stmt
));
3846 if (lhs
== NULL_TREE
)
3849 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3851 query_hsa_grid (dest
, opcode
, dimension
, hbb
);
3854 /* Emit instructions that set hsa_num_threads according to provided VALUE.
3855 Instructions are appended to basic block HBB. */
3858 gen_set_num_threads (tree value
, hsa_bb
*hbb
)
3860 hbb
->append_insn (new hsa_insn_comment ("omp_set_num_threads"));
3861 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (value
, hbb
);
3863 src
= src
->get_in_type (hsa_num_threads
->m_type
, hbb
);
3864 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
3866 hsa_insn_basic
*basic
3867 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
, src
, addr
);
3868 hbb
->append_insn (basic
);
3871 /* Return byte offset of a FIELD_NAME in GOMP_hsa_kernel_dispatch which
3872 is defined in plugin-hsa.c. */
3874 static HOST_WIDE_INT
3875 get_hsa_kernel_dispatch_offset (const char *field_name
)
3877 tree
*hsa_kernel_dispatch_type
= hsa_get_kernel_dispatch_type ();
3878 if (*hsa_kernel_dispatch_type
== NULL
)
3880 /* Collection of information needed for a dispatch of a kernel from a
3881 kernel. Keep in sync with libgomp's plugin-hsa.c. */
3883 *hsa_kernel_dispatch_type
= make_node (RECORD_TYPE
);
3884 tree id_f1
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3885 get_identifier ("queue"), ptr_type_node
);
3886 DECL_CHAIN (id_f1
) = NULL_TREE
;
3887 tree id_f2
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3888 get_identifier ("omp_data_memory"),
3890 DECL_CHAIN (id_f2
) = id_f1
;
3891 tree id_f3
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3892 get_identifier ("kernarg_address"),
3894 DECL_CHAIN (id_f3
) = id_f2
;
3895 tree id_f4
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3896 get_identifier ("object"),
3898 DECL_CHAIN (id_f4
) = id_f3
;
3899 tree id_f5
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3900 get_identifier ("signal"),
3902 DECL_CHAIN (id_f5
) = id_f4
;
3903 tree id_f6
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3904 get_identifier ("private_segment_size"),
3906 DECL_CHAIN (id_f6
) = id_f5
;
3907 tree id_f7
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3908 get_identifier ("group_segment_size"),
3910 DECL_CHAIN (id_f7
) = id_f6
;
3911 tree id_f8
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3912 get_identifier ("kernel_dispatch_count"),
3914 DECL_CHAIN (id_f8
) = id_f7
;
3915 tree id_f9
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3916 get_identifier ("debug"),
3918 DECL_CHAIN (id_f9
) = id_f8
;
3919 tree id_f10
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3920 get_identifier ("omp_level"),
3922 DECL_CHAIN (id_f10
) = id_f9
;
3923 tree id_f11
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3924 get_identifier ("children_dispatches"),
3926 DECL_CHAIN (id_f11
) = id_f10
;
3927 tree id_f12
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3928 get_identifier ("omp_num_threads"),
3930 DECL_CHAIN (id_f12
) = id_f11
;
3933 finish_builtin_struct (*hsa_kernel_dispatch_type
, "__hsa_kernel_dispatch",
3935 TYPE_ARTIFICIAL (*hsa_kernel_dispatch_type
) = 1;
3938 for (tree chain
= TYPE_FIELDS (*hsa_kernel_dispatch_type
);
3939 chain
!= NULL_TREE
; chain
= TREE_CHAIN (chain
))
3940 if (strcmp (field_name
, IDENTIFIER_POINTER (DECL_NAME (chain
))) == 0)
3941 return int_byte_position (chain
);
3946 /* Return an HSA register that will contain number of threads for
3947 a future dispatched kernel. Instructions are added to HBB. */
3950 gen_num_threads_for_dispatch (hsa_bb
*hbb
)
3952 /* Step 1) Assign to number of threads:
3953 MIN (HSA_DEFAULT_NUM_THREADS, hsa_num_threads). */
3954 hsa_op_reg
*threads
= new hsa_op_reg (hsa_num_threads
->m_type
);
3955 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
3957 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD
, threads
->m_type
,
3960 hsa_op_immed
*limit
= new hsa_op_immed (HSA_DEFAULT_NUM_THREADS
,
3962 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
3964 = new hsa_insn_cmp (BRIG_COMPARE_LT
, r
->m_type
, r
, threads
, limit
);
3965 hbb
->append_insn (cmp
);
3967 BrigType16_t btype
= hsa_bittype_for_type (threads
->m_type
);
3968 hsa_op_reg
*tmp
= new hsa_op_reg (threads
->m_type
);
3970 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp
, r
,
3973 /* Step 2) If the number is equal to zero,
3974 return shadow->omp_num_threads. */
3975 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
3977 hsa_op_reg
*shadow_thread_count
= new hsa_op_reg (BRIG_TYPE_U32
);
3979 = new hsa_op_address (shadow_reg_ptr
,
3980 get_hsa_kernel_dispatch_offset ("omp_num_threads"));
3981 hsa_insn_basic
*basic
3982 = new hsa_insn_mem (BRIG_OPCODE_LD
, shadow_thread_count
->m_type
,
3983 shadow_thread_count
, addr
);
3984 hbb
->append_insn (basic
);
3986 hsa_op_reg
*tmp2
= new hsa_op_reg (threads
->m_type
);
3987 r
= new hsa_op_reg (BRIG_TYPE_B1
);
3988 hsa_op_immed
*imm
= new hsa_op_immed (0, shadow_thread_count
->m_type
);
3989 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_EQ
, r
->m_type
, r
, tmp
, imm
));
3990 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp2
, r
,
3991 shadow_thread_count
, tmp
));
3993 hsa_op_base
*dest
= tmp2
->get_in_type (BRIG_TYPE_U16
, hbb
);
3995 return as_a
<hsa_op_reg
*> (dest
);
3999 /* Emit instructions that assign number of teams to lhs of gimple STMT.
4000 Instructions are appended to basic block HBB. */
4003 gen_get_num_teams (gimple
*stmt
, hsa_bb
*hbb
)
4005 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4008 hbb
->append_insn (new hsa_insn_comment ("omp_get_num_teams"));
4010 tree lhs
= gimple_call_lhs (stmt
);
4011 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4012 hsa_op_immed
*one
= new hsa_op_immed (1, dest
->m_type
);
4014 hsa_insn_basic
*basic
4015 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, dest
, one
);
4017 hbb
->append_insn (basic
);
4020 /* Emit instructions that assign a team number to lhs of gimple STMT.
4021 Instructions are appended to basic block HBB. */
4024 gen_get_team_num (gimple
*stmt
, hsa_bb
*hbb
)
4026 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4029 hbb
->append_insn (new hsa_insn_comment ("omp_get_team_num"));
4031 tree lhs
= gimple_call_lhs (stmt
);
4032 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4033 hsa_op_immed
*zero
= new hsa_op_immed (0, dest
->m_type
);
4035 hsa_insn_basic
*basic
4036 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, dest
, zero
);
4038 hbb
->append_insn (basic
);
4041 /* Emit instructions that get levels-var ICV to lhs of gimple STMT.
4042 Instructions are appended to basic block HBB. */
4045 gen_get_level (gimple
*stmt
, hsa_bb
*hbb
)
4047 if (gimple_call_lhs (stmt
) == NULL_TREE
)
4050 hbb
->append_insn (new hsa_insn_comment ("omp_get_level"));
4052 tree lhs
= gimple_call_lhs (stmt
);
4053 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4055 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4056 if (shadow_reg_ptr
== NULL
)
4058 HSA_SORRY_AT (gimple_location (stmt
),
4059 "support for HSA does not implement omp_get_level called "
4060 "from a function not being inlined within a kernel");
4064 hsa_op_address
*addr
4065 = new hsa_op_address (shadow_reg_ptr
,
4066 get_hsa_kernel_dispatch_offset ("omp_level"));
4068 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, BRIG_TYPE_U64
,
4069 (hsa_op_base
*) NULL
, addr
);
4070 hbb
->append_insn (mem
);
4071 mem
->set_output_in_type (dest
, 0, hbb
);
4074 /* Emit instruction that implement omp_get_max_threads of gimple STMT. */
4077 gen_get_max_threads (gimple
*stmt
, hsa_bb
*hbb
)
4079 tree lhs
= gimple_call_lhs (stmt
);
4083 hbb
->append_insn (new hsa_insn_comment ("omp_get_max_threads"));
4085 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4086 hsa_op_with_type
*num_theads_reg
= gen_num_threads_for_dispatch (hbb
)
4087 ->get_in_type (dest
->m_type
, hbb
);
4088 hsa_build_append_simple_mov (dest
, num_theads_reg
, hbb
);
4091 /* Emit instructions that implement alloca builtin gimple STMT.
4092 Instructions are appended to basic block HBB. */
4095 gen_hsa_alloca (gcall
*call
, hsa_bb
*hbb
)
4097 tree lhs
= gimple_call_lhs (call
);
4098 if (lhs
== NULL_TREE
)
4101 built_in_function fn
= DECL_FUNCTION_CODE (gimple_call_fndecl (call
));
4103 gcc_checking_assert (fn
== BUILT_IN_ALLOCA
4104 || fn
== BUILT_IN_ALLOCA_WITH_ALIGN
);
4106 unsigned bit_alignment
= 0;
4108 if (fn
== BUILT_IN_ALLOCA_WITH_ALIGN
)
4110 tree alignment_tree
= gimple_call_arg (call
, 1);
4111 if (TREE_CODE (alignment_tree
) != INTEGER_CST
)
4113 HSA_SORRY_ATV (gimple_location (call
),
4114 "support for HSA does not implement "
4115 "__builtin_alloca_with_align with a non-constant "
4116 "alignment: %E", alignment_tree
);
4119 bit_alignment
= tree_to_uhwi (alignment_tree
);
4122 tree rhs1
= gimple_call_arg (call
, 0);
4123 hsa_op_with_type
*size
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
)
4124 ->get_in_type (BRIG_TYPE_U32
, hbb
);
4125 hsa_op_with_type
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4128 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
));
4129 hsa_insn_alloca
*a
= new hsa_insn_alloca (tmp
, size
, bit_alignment
);
4130 hbb
->append_insn (a
);
4133 = new hsa_insn_seg (BRIG_OPCODE_STOF
,
4134 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
4135 tmp
->m_type
, BRIG_SEGMENT_PRIVATE
, dest
, tmp
);
4136 hbb
->append_insn (seg
);
4139 /* Emit instructions that implement clrsb builtin STMT:
4140 Returns the number of leading redundant sign bits in x, i.e. the number
4141 of bits following the most significant bit that are identical to it.
4142 There are no special cases for 0 or other values.
4143 Instructions are appended to basic block HBB. */
4146 gen_hsa_clrsb (gcall
*call
, hsa_bb
*hbb
)
4148 tree lhs
= gimple_call_lhs (call
);
4149 if (lhs
== NULL_TREE
)
4152 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4153 tree rhs1
= gimple_call_arg (call
, 0);
4154 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4155 BrigType16_t bittype
= hsa_bittype_for_type (arg
->m_type
);
4156 unsigned bitsize
= tree_to_uhwi (TYPE_SIZE (TREE_TYPE (rhs1
)));
4158 /* FIRSTBIT instruction is defined just for 32 and 64-bits wide integers. */
4159 gcc_checking_assert (bitsize
== 32 || bitsize
== 64);
4161 /* Set true to MOST_SIG if the most significant bit is set to one. */
4162 hsa_op_immed
*c
= new hsa_op_immed (1ul << (bitsize
- 1),
4163 hsa_uint_for_bitsize (bitsize
));
4165 hsa_op_reg
*and_reg
= new hsa_op_reg (bittype
);
4166 gen_hsa_binary_operation (BRIG_OPCODE_AND
, and_reg
, arg
, c
, hbb
);
4168 hsa_op_reg
*most_sign
= new hsa_op_reg (BRIG_TYPE_B1
);
4170 = new hsa_insn_cmp (BRIG_COMPARE_EQ
, most_sign
->m_type
, most_sign
,
4172 hbb
->append_insn (cmp
);
4174 /* If the most significant bit is one, negate the input. Otherwise
4175 shift the input value to left by one bit. */
4176 hsa_op_reg
*arg_neg
= new hsa_op_reg (arg
->m_type
);
4177 gen_hsa_unary_operation (BRIG_OPCODE_NEG
, arg_neg
, arg
, hbb
);
4179 hsa_op_reg
*shifted_arg
= new hsa_op_reg (arg
->m_type
);
4180 gen_hsa_binary_operation (BRIG_OPCODE_SHL
, shifted_arg
, arg
,
4181 new hsa_op_immed (1, BRIG_TYPE_U64
), hbb
);
4183 /* Assign the value that can be used for FIRSTBIT instruction according
4184 to the most significant bit. */
4185 hsa_op_reg
*tmp
= new hsa_op_reg (bittype
);
4186 hsa_insn_basic
*cmov
4187 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, bittype
, tmp
, most_sign
,
4188 arg_neg
, shifted_arg
);
4189 hbb
->append_insn (cmov
);
4191 hsa_op_reg
*leading_bits
= new hsa_op_reg (BRIG_TYPE_S32
);
4192 gen_hsa_unary_operation (BRIG_OPCODE_FIRSTBIT
, leading_bits
,
4193 tmp
->get_in_type (hsa_uint_for_bitsize (bitsize
),
4196 /* Set flag if the input value is equal to zero. */
4197 hsa_op_reg
*is_zero
= new hsa_op_reg (BRIG_TYPE_B1
);
4198 cmp
= new hsa_insn_cmp (BRIG_COMPARE_EQ
, is_zero
->m_type
, is_zero
, arg
,
4199 new hsa_op_immed (0, arg
->m_type
));
4200 hbb
->append_insn (cmp
);
4202 /* Return the number of leading bits,
4203 or (bitsize - 1) if the input value is zero. */
4204 cmov
= new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, BRIG_TYPE_B32
, NULL
, is_zero
,
4205 new hsa_op_immed (bitsize
- 1, BRIG_TYPE_U32
),
4206 leading_bits
->get_in_type (BRIG_TYPE_B32
, hbb
));
4207 hbb
->append_insn (cmov
);
4208 cmov
->set_output_in_type (dest
, 0, hbb
);
4211 /* Emit instructions that implement ffs builtin STMT:
4212 Returns one plus the index of the least significant 1-bit of x,
4213 or if x is zero, returns zero.
4214 Instructions are appended to basic block HBB. */
4217 gen_hsa_ffs (gcall
*call
, hsa_bb
*hbb
)
4219 tree lhs
= gimple_call_lhs (call
);
4220 if (lhs
== NULL_TREE
)
4223 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4225 tree rhs1
= gimple_call_arg (call
, 0);
4226 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4228 hsa_op_reg
*tmp
= new hsa_op_reg (BRIG_TYPE_U32
);
4229 hsa_insn_srctype
*insn
= new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT
,
4230 tmp
->m_type
, arg
->m_type
,
4232 hbb
->append_insn (insn
);
4234 hsa_insn_basic
*addition
4235 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
, NULL
, tmp
,
4236 new hsa_op_immed (1, tmp
->m_type
));
4237 hbb
->append_insn (addition
);
4238 addition
->set_output_in_type (dest
, 0, hbb
);
4242 gen_hsa_popcount_to_dest (hsa_op_reg
*dest
, hsa_op_with_type
*arg
, hsa_bb
*hbb
)
4244 gcc_checking_assert (hsa_type_integer_p (arg
->m_type
));
4246 if (hsa_type_bit_size (arg
->m_type
) < 32)
4247 arg
= arg
->get_in_type (BRIG_TYPE_B32
, hbb
);
4249 if (!hsa_btype_p (arg
->m_type
))
4250 arg
= arg
->get_in_type (hsa_bittype_for_type (arg
->m_type
), hbb
);
4252 hsa_insn_srctype
*popcount
4253 = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT
, BRIG_TYPE_U32
,
4254 arg
->m_type
, NULL
, arg
);
4255 hbb
->append_insn (popcount
);
4256 popcount
->set_output_in_type (dest
, 0, hbb
);
4259 /* Emit instructions that implement parity builtin STMT:
4260 Returns the parity of x, i.e. the number of 1-bits in x modulo 2.
4261 Instructions are appended to basic block HBB. */
4264 gen_hsa_parity (gcall
*call
, hsa_bb
*hbb
)
4266 tree lhs
= gimple_call_lhs (call
);
4267 if (lhs
== NULL_TREE
)
4270 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4271 tree rhs1
= gimple_call_arg (call
, 0);
4272 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4274 hsa_op_reg
*popcount
= new hsa_op_reg (BRIG_TYPE_U32
);
4275 gen_hsa_popcount_to_dest (popcount
, arg
, hbb
);
4277 hsa_insn_basic
*insn
4278 = new hsa_insn_basic (3, BRIG_OPCODE_REM
, popcount
->m_type
, NULL
, popcount
,
4279 new hsa_op_immed (2, popcount
->m_type
));
4280 hbb
->append_insn (insn
);
4281 insn
->set_output_in_type (dest
, 0, hbb
);
4284 /* Emit instructions that implement popcount builtin STMT.
4285 Instructions are appended to basic block HBB. */
4288 gen_hsa_popcount (gcall
*call
, hsa_bb
*hbb
)
4290 tree lhs
= gimple_call_lhs (call
);
4291 if (lhs
== NULL_TREE
)
4294 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4295 tree rhs1
= gimple_call_arg (call
, 0);
4296 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4298 gen_hsa_popcount_to_dest (dest
, arg
, hbb
);
4301 /* Set VALUE to a shadow kernel debug argument and append a new instruction
4302 to HBB basic block. */
4305 set_debug_value (hsa_bb
*hbb
, hsa_op_with_type
*value
)
4307 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4308 if (shadow_reg_ptr
== NULL
)
4311 hsa_op_address
*addr
4312 = new hsa_op_address (shadow_reg_ptr
,
4313 get_hsa_kernel_dispatch_offset ("debug"));
4314 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, BRIG_TYPE_U64
, value
,
4316 hbb
->append_insn (mem
);
4320 omp_simple_builtin::generate (gimple
*stmt
, hsa_bb
*hbb
)
4324 if (m_warning_message
)
4325 HSA_SORRY_AT (gimple_location (stmt
), m_warning_message
)
4327 HSA_SORRY_ATV (gimple_location (stmt
),
4328 "Support for HSA does not implement calls to %s\n",
4331 else if (m_warning_message
!= NULL
)
4332 warning_at (gimple_location (stmt
), OPT_Whsa
, m_warning_message
);
4334 if (m_return_value
!= NULL
)
4336 tree lhs
= gimple_call_lhs (stmt
);
4340 hbb
->append_insn (new hsa_insn_comment (m_name
));
4342 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4343 hsa_op_with_type
*op
= m_return_value
->get_in_type (dest
->m_type
, hbb
);
4344 hsa_build_append_simple_mov (dest
, op
, hbb
);
4348 /* If STMT is a call of a known library function, generate code to perform
4349 it and return true. */
4352 gen_hsa_insns_for_known_library_call (gimple
*stmt
, hsa_bb
*hbb
)
4354 bool handled
= false;
4355 const char *name
= hsa_get_declaration_name (gimple_call_fndecl (stmt
));
4358 size_t len
= strlen (name
);
4359 if (len
> 0 && name
[len
- 1] == '_')
4361 copy
= XNEWVEC (char, len
+ 1);
4362 strcpy (copy
, name
);
4363 copy
[len
- 1] = '\0';
4367 /* Handle omp_* routines. */
4368 if (strstr (name
, "omp_") == name
)
4370 hsa_init_simple_builtins ();
4371 omp_simple_builtin
*builtin
= omp_simple_builtins
->get (name
);
4374 builtin
->generate (stmt
, hbb
);
4379 if (strcmp (name
, "omp_set_num_threads") == 0)
4380 gen_set_num_threads (gimple_call_arg (stmt
, 0), hbb
);
4381 else if (strcmp (name
, "omp_get_thread_num") == 0)
4383 hbb
->append_insn (new hsa_insn_comment (name
));
4384 query_hsa_grid (stmt
, BRIG_OPCODE_WORKITEMABSID
, 0, hbb
);
4386 else if (strcmp (name
, "omp_get_num_threads") == 0)
4388 hbb
->append_insn (new hsa_insn_comment (name
));
4389 query_hsa_grid (stmt
, BRIG_OPCODE_GRIDSIZE
, 0, hbb
);
4391 else if (strcmp (name
, "omp_get_num_teams") == 0)
4392 gen_get_num_teams (stmt
, hbb
);
4393 else if (strcmp (name
, "omp_get_team_num") == 0)
4394 gen_get_team_num (stmt
, hbb
);
4395 else if (strcmp (name
, "omp_get_level") == 0)
4396 gen_get_level (stmt
, hbb
);
4397 else if (strcmp (name
, "omp_get_active_level") == 0)
4398 gen_get_level (stmt
, hbb
);
4399 else if (strcmp (name
, "omp_in_parallel") == 0)
4400 gen_get_level (stmt
, hbb
);
4401 else if (strcmp (name
, "omp_get_max_threads") == 0)
4402 gen_get_max_threads (stmt
, hbb
);
4414 if (strcmp (name
, "__hsa_set_debug_value") == 0)
4417 if (hsa_cfun
->has_shadow_reg_p ())
4419 tree rhs1
= gimple_call_arg (stmt
, 0);
4420 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4422 src
= src
->get_in_type (BRIG_TYPE_U64
, hbb
);
4423 set_debug_value (hbb
, src
);
4432 /* Helper functions to create a single unary HSA operations out of calls to
4433 builtins. OPCODE is the HSA operation to be generated. STMT is a gimple
4434 call to a builtin. HBB is the HSA BB to which the instruction should be
4435 added. Note that nothing will be created if STMT does not have a LHS. */
4438 gen_hsa_unaryop_for_builtin (BrigOpcode opcode
, gimple
*stmt
, hsa_bb
*hbb
)
4440 tree lhs
= gimple_call_lhs (stmt
);
4443 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4444 hsa_op_with_type
*op
4445 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4446 gen_hsa_unary_operation (opcode
, dest
, op
, hbb
);
4449 /* Helper functions to create a call to standard library if LHS of the
4450 STMT is used. HBB is the HSA BB to which the instruction should be
4454 gen_hsa_unaryop_builtin_call (gimple
*stmt
, hsa_bb
*hbb
)
4456 tree lhs
= gimple_call_lhs (stmt
);
4460 if (gimple_call_internal_p (stmt
))
4461 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
4463 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4466 /* Helper functions to create a single unary HSA operations out of calls to
4467 builtins (if unsafe math optimizations are enable). Otherwise, create
4468 a call to standard library function.
4469 OPCODE is the HSA operation to be generated. STMT is a gimple
4470 call to a builtin. HBB is the HSA BB to which the instruction should be
4471 added. Note that nothing will be created if STMT does not have a LHS. */
4474 gen_hsa_unaryop_or_call_for_builtin (BrigOpcode opcode
, gimple
*stmt
,
4477 if (flag_unsafe_math_optimizations
)
4478 gen_hsa_unaryop_for_builtin (opcode
, stmt
, hbb
);
4480 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4483 /* Generate HSA address corresponding to a value VAL (as opposed to a memory
4484 reference tree), for example an SSA_NAME or an ADDR_EXPR. HBB is the HSA BB
4485 to which the instruction should be added. */
4487 static hsa_op_address
*
4488 get_address_from_value (tree val
, hsa_bb
*hbb
)
4490 switch (TREE_CODE (val
))
4494 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4496 = hsa_cfun
->reg_for_gimple_ssa (val
)->get_in_type (addrtype
, hbb
);
4497 return new hsa_op_address (NULL
, as_a
<hsa_op_reg
*> (reg
), 0);
4500 return gen_hsa_addr (TREE_OPERAND (val
, 0), hbb
);
4503 if (tree_fits_shwi_p (val
))
4504 return new hsa_op_address (NULL
, NULL
, tree_to_shwi (val
));
4505 /* Otherwise fall-through */
4508 HSA_SORRY_ATV (EXPR_LOCATION (val
),
4509 "support for HSA does not implement memory access to %E",
4511 return new hsa_op_address (NULL
, NULL
, 0);
4515 /* Expand assignment of a result of a string BUILTIN to DST.
4516 Size of the operation is N bytes, where instructions
4517 will be append to HBB. */
4520 expand_lhs_of_string_op (gimple
*stmt
,
4521 unsigned HOST_WIDE_INT n
, hsa_bb
*hbb
,
4522 enum built_in_function builtin
)
4524 /* If LHS is expected, we need to emit a PHI instruction. */
4525 tree lhs
= gimple_call_lhs (stmt
);
4529 hsa_op_reg
*lhs_reg
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4531 hsa_op_with_type
*dst_reg
4532 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4533 hsa_op_with_type
*tmp
;
4537 case BUILT_IN_MEMPCPY
:
4539 tmp
= new hsa_op_reg (dst_reg
->m_type
);
4541 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
,
4543 new hsa_op_immed (n
, dst_reg
->m_type
));
4544 hbb
->append_insn (add
);
4547 case BUILT_IN_MEMCPY
:
4548 case BUILT_IN_MEMSET
:
4555 hbb
->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV
, lhs_reg
->m_type
,
4559 #define HSA_MEMORY_BUILTINS_LIMIT 128
4561 /* Expand a string builtin (from a gimple STMT) in a way that
4562 according to MISALIGNED_FLAG we process either direct emission
4563 (a bunch of memory load and store instructions), or we emit a function call
4564 of a library function (for instance 'memcpy'). Actually, a basic block
4565 for direct emission is just prepared, where caller is responsible
4566 for emission of corresponding instructions.
4567 All instruction are appended to HBB. */
4570 expand_string_operation_builtin (gimple
*stmt
, hsa_bb
*hbb
,
4571 hsa_op_reg
*misaligned_flag
)
4573 edge e
= split_block (hbb
->m_bb
, stmt
);
4574 basic_block condition_bb
= e
->src
;
4575 hbb
->append_insn (new hsa_insn_br (misaligned_flag
));
4577 /* Prepare the control flow. */
4578 edge condition_edge
= EDGE_SUCC (condition_bb
, 0);
4579 basic_block call_bb
= split_edge (condition_edge
);
4581 basic_block expanded_bb
= split_edge (EDGE_SUCC (call_bb
, 0));
4582 basic_block cont_bb
= EDGE_SUCC (expanded_bb
, 0)->dest
;
4583 basic_block merge_bb
= split_edge (EDGE_PRED (cont_bb
, 0));
4585 condition_edge
->flags
&= ~EDGE_FALLTHRU
;
4586 condition_edge
->flags
|= EDGE_TRUE_VALUE
;
4587 make_edge (condition_bb
, expanded_bb
, EDGE_FALSE_VALUE
);
4589 redirect_edge_succ (EDGE_SUCC (call_bb
, 0), merge_bb
);
4591 hsa_cfun
->m_modified_cfg
= true;
4593 hsa_init_new_bb (expanded_bb
);
4595 /* Slow path: function call. */
4596 gen_hsa_insns_for_direct_call (stmt
, hsa_init_new_bb (call_bb
), false);
4598 return hsa_bb_for_bb (expanded_bb
);
4601 /* Expand a memory copy BUILTIN (BUILT_IN_MEMCPY, BUILT_IN_MEMPCPY) from
4602 a gimple STMT and store all necessary instruction to HBB basic block. */
4605 expand_memory_copy (gimple
*stmt
, hsa_bb
*hbb
, enum built_in_function builtin
)
4607 tree byte_size
= gimple_call_arg (stmt
, 2);
4609 if (!tree_fits_uhwi_p (byte_size
))
4611 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4615 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
4617 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
4619 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4623 tree dst
= gimple_call_arg (stmt
, 0);
4624 tree src
= gimple_call_arg (stmt
, 1);
4626 hsa_op_address
*dst_addr
= get_address_from_value (dst
, hbb
);
4627 hsa_op_address
*src_addr
= get_address_from_value (src
, hbb
);
4629 /* As gen_hsa_memory_copy relies on memory alignment
4630 greater or equal to 8 bytes, we need to verify the alignment. */
4631 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4632 hsa_op_reg
*src_addr_reg
= new hsa_op_reg (addrtype
);
4633 hsa_op_reg
*dst_addr_reg
= new hsa_op_reg (addrtype
);
4635 convert_addr_to_flat_segment (src_addr
, src_addr_reg
, hbb
);
4636 convert_addr_to_flat_segment (dst_addr
, dst_addr_reg
, hbb
);
4638 /* Process BIT OR for source and destination addresses. */
4639 hsa_op_reg
*or_reg
= new hsa_op_reg (addrtype
);
4640 gen_hsa_binary_operation (BRIG_OPCODE_OR
, or_reg
, src_addr_reg
,
4643 /* Process BIT AND with 0x7 to identify the desired alignment
4645 hsa_op_reg
*masked
= new hsa_op_reg (addrtype
);
4647 gen_hsa_binary_operation (BRIG_OPCODE_AND
, masked
, or_reg
,
4648 new hsa_op_immed (7, addrtype
), hbb
);
4650 hsa_op_reg
*misaligned
= new hsa_op_reg (BRIG_TYPE_B1
);
4651 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE
, misaligned
->m_type
,
4653 new hsa_op_immed (0, masked
->m_type
)));
4655 hsa_bb
*native_impl_bb
4656 = expand_string_operation_builtin (stmt
, hbb
, misaligned
);
4658 gen_hsa_memory_copy (native_impl_bb
, dst_addr
, src_addr
, n
, BRIG_ALIGNMENT_8
);
4660 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb
->m_bb
, 0)->dest
);
4661 expand_lhs_of_string_op (stmt
, n
, merge_bb
, builtin
);
4665 /* Expand a memory set BUILTIN (BUILT_IN_MEMSET, BUILT_IN_BZERO) from
4666 a gimple STMT and store all necessary instruction to HBB basic block.
4667 The operation set N bytes with a CONSTANT value. */
4670 expand_memory_set (gimple
*stmt
, unsigned HOST_WIDE_INT n
,
4671 unsigned HOST_WIDE_INT constant
, hsa_bb
*hbb
,
4672 enum built_in_function builtin
)
4674 tree dst
= gimple_call_arg (stmt
, 0);
4675 hsa_op_address
*dst_addr
= get_address_from_value (dst
, hbb
);
4677 /* As gen_hsa_memory_set relies on memory alignment
4678 greater or equal to 8 bytes, we need to verify the alignment. */
4679 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4680 hsa_op_reg
*dst_addr_reg
= new hsa_op_reg (addrtype
);
4681 convert_addr_to_flat_segment (dst_addr
, dst_addr_reg
, hbb
);
4683 /* Process BIT AND with 0x7 to identify the desired alignment
4685 hsa_op_reg
*masked
= new hsa_op_reg (addrtype
);
4687 gen_hsa_binary_operation (BRIG_OPCODE_AND
, masked
, dst_addr_reg
,
4688 new hsa_op_immed (7, addrtype
), hbb
);
4690 hsa_op_reg
*misaligned
= new hsa_op_reg (BRIG_TYPE_B1
);
4691 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_NE
, misaligned
->m_type
,
4693 new hsa_op_immed (0, masked
->m_type
)));
4695 hsa_bb
*native_impl_bb
4696 = expand_string_operation_builtin (stmt
, hbb
, misaligned
);
4698 gen_hsa_memory_set (native_impl_bb
, dst_addr
, constant
, n
, BRIG_ALIGNMENT_8
);
4700 = hsa_init_new_bb (EDGE_SUCC (native_impl_bb
->m_bb
, 0)->dest
);
4701 expand_lhs_of_string_op (stmt
, n
, merge_bb
, builtin
);
4704 /* Return string for MEMMODEL. */
4707 get_memory_order_name (unsigned memmodel
)
4709 switch (memmodel
& MEMMODEL_BASE_MASK
)
4711 case MEMMODEL_RELAXED
:
4713 case MEMMODEL_CONSUME
:
4715 case MEMMODEL_ACQUIRE
:
4717 case MEMMODEL_RELEASE
:
4719 case MEMMODEL_ACQ_REL
:
4721 case MEMMODEL_SEQ_CST
:
4728 /* Return memory order according to predefined __atomic memory model
4729 constants. LOCATION is provided to locate the problematic statement. */
4731 static BrigMemoryOrder
4732 get_memory_order (unsigned memmodel
, location_t location
)
4734 switch (memmodel
& MEMMODEL_BASE_MASK
)
4736 case MEMMODEL_RELAXED
:
4737 return BRIG_MEMORY_ORDER_RELAXED
;
4738 case MEMMODEL_CONSUME
:
4739 /* HSA does not have an equivalent, but we can use the slightly stronger
4741 case MEMMODEL_ACQUIRE
:
4742 return BRIG_MEMORY_ORDER_SC_ACQUIRE
;
4743 case MEMMODEL_RELEASE
:
4744 return BRIG_MEMORY_ORDER_SC_RELEASE
;
4745 case MEMMODEL_ACQ_REL
:
4746 case MEMMODEL_SEQ_CST
:
4747 /* Callers implementing a simple load or store need to remove the release
4748 or acquire part respectively. */
4749 return BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
4752 const char *mmname
= get_memory_order_name (memmodel
);
4753 HSA_SORRY_ATV (location
,
4754 "support for HSA does not implement the specified "
4755 " memory model%s %s",
4756 mmname
? ": " : "", mmname
? mmname
: "");
4757 return BRIG_MEMORY_ORDER_NONE
;
4762 /* Helper function to create an HSA atomic binary operation instruction out of
4763 calls to atomic builtins. RET_ORIG is true if the built-in is the variant
4764 that return s the value before applying operation, and false if it should
4765 return the value after applying the operation (if it returns value at all).
4766 ACODE is the atomic operation code, STMT is a gimple call to a builtin. HBB
4767 is the HSA BB to which the instruction should be added. */
4770 gen_hsa_ternary_atomic_for_builtin (bool ret_orig
,
4771 enum BrigAtomicOperation acode
,
4775 tree lhs
= gimple_call_lhs (stmt
);
4777 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
4778 BrigType16_t hsa_type
= hsa_type_for_scalar_tree_type (type
, false);
4779 BrigType16_t mtype
= mem_type_for_type (hsa_type
);
4780 tree model
= gimple_call_arg (stmt
, 2);
4782 if (!tree_fits_uhwi_p (model
))
4784 HSA_SORRY_ATV (gimple_location (stmt
),
4785 "support for HSA does not implement memory model %E",
4790 unsigned HOST_WIDE_INT mmodel
= tree_to_uhwi (model
);
4792 BrigMemoryOrder memorder
= get_memory_order (mmodel
, gimple_location (stmt
));
4794 /* Certain atomic insns must have Bx memory types. */
4797 case BRIG_ATOMIC_LD
:
4798 case BRIG_ATOMIC_ST
:
4799 case BRIG_ATOMIC_AND
:
4800 case BRIG_ATOMIC_OR
:
4801 case BRIG_ATOMIC_XOR
:
4802 case BRIG_ATOMIC_EXCH
:
4803 mtype
= hsa_bittype_for_type (mtype
);
4814 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4816 dest
= new hsa_op_reg (hsa_type
);
4817 opcode
= BRIG_OPCODE_ATOMIC
;
4823 opcode
= BRIG_OPCODE_ATOMICNORET
;
4827 if (acode
== BRIG_ATOMIC_ST
)
4829 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
4830 memorder
= BRIG_MEMORY_ORDER_SC_RELEASE
;
4832 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
4833 && memorder
!= BRIG_MEMORY_ORDER_SC_RELEASE
4834 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
4836 HSA_SORRY_ATV (gimple_location (stmt
),
4837 "support for HSA does not implement memory model for "
4838 "ATOMIC_ST: %s", get_memory_order_name (mmodel
));
4843 hsa_insn_atomic
*atominsn
= new hsa_insn_atomic (nops
, opcode
, acode
, mtype
,
4846 hsa_op_address
*addr
;
4847 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
4848 if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
== BRIG_SEGMENT_PRIVATE
)
4850 HSA_SORRY_AT (gimple_location (stmt
),
4851 "HSA does not implement atomic operations in private "
4855 hsa_op_base
*op
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1),
4860 atominsn
->set_op (0, dest
);
4861 atominsn
->set_op (1, addr
);
4862 atominsn
->set_op (2, op
);
4866 atominsn
->set_op (0, addr
);
4867 atominsn
->set_op (1, op
);
4870 hbb
->append_insn (atominsn
);
4872 /* HSA does not natively support the variants that return the modified value,
4873 so re-do the operation again non-atomically if that is what was
4875 if (lhs
&& !ret_orig
)
4880 case BRIG_ATOMIC_ADD
:
4881 arith
= BRIG_OPCODE_ADD
;
4883 case BRIG_ATOMIC_AND
:
4884 arith
= BRIG_OPCODE_AND
;
4886 case BRIG_ATOMIC_OR
:
4887 arith
= BRIG_OPCODE_OR
;
4889 case BRIG_ATOMIC_SUB
:
4890 arith
= BRIG_OPCODE_SUB
;
4892 case BRIG_ATOMIC_XOR
:
4893 arith
= BRIG_OPCODE_XOR
;
4898 hsa_op_reg
*real_dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4899 gen_hsa_binary_operation (arith
, real_dest
, dest
, op
, hbb
);
4903 /* Generate HSA instructions for an internal fn.
4904 Instructions will be appended to HBB, which also needs to be the
4905 corresponding structure to the basic_block of STMT. */
4908 gen_hsa_insn_for_internal_fn_call (gcall
*stmt
, hsa_bb
*hbb
)
4910 gcc_checking_assert (gimple_call_internal_fn (stmt
));
4911 internal_fn fn
= gimple_call_internal_fn (stmt
);
4913 bool is_float_type_p
= false;
4914 if (gimple_call_lhs (stmt
) != NULL
4915 && TREE_TYPE (gimple_call_lhs (stmt
)) == float_type_node
)
4916 is_float_type_p
= true;
4921 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
4925 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
4929 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
4933 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
4937 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
4942 if (is_float_type_p
)
4943 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
4945 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4951 if (is_float_type_p
)
4952 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
4954 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4961 if (is_float_type_p
)
4962 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
4964 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4971 if (is_float_type_p
)
4972 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
4974 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4979 gen_hsa_clrsb (stmt
, hbb
);
4983 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
4987 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
4991 gen_hsa_ffs (stmt
, hbb
);
4995 gen_hsa_parity (stmt
, hbb
);
4999 gen_hsa_popcount (stmt
, hbb
);
5012 case IFN_SIGNIFICAND
:
5024 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
5027 HSA_SORRY_ATV (gimple_location (stmt
),
5028 "support for HSA does not implement internal function: %s",
5029 internal_fn_name (fn
));
5034 /* Generate HSA instructions for the given call statement STMT. Instructions
5035 will be appended to HBB. */
5038 gen_hsa_insns_for_call (gimple
*stmt
, hsa_bb
*hbb
)
5040 gcall
*call
= as_a
<gcall
*> (stmt
);
5041 tree lhs
= gimple_call_lhs (stmt
);
5044 if (gimple_call_internal_p (stmt
))
5046 gen_hsa_insn_for_internal_fn_call (call
, hbb
);
5050 if (!gimple_call_builtin_p (stmt
, BUILT_IN_NORMAL
))
5052 tree function_decl
= gimple_call_fndecl (stmt
);
5053 if (function_decl
== NULL_TREE
)
5055 HSA_SORRY_AT (gimple_location (stmt
),
5056 "support for HSA does not implement indirect calls");
5060 if (hsa_callable_function_p (function_decl
))
5061 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5062 else if (!gen_hsa_insns_for_known_library_call (stmt
, hbb
))
5063 HSA_SORRY_AT (gimple_location (stmt
),
5064 "HSA supports only calls of functions marked with pragma "
5065 "omp declare target");
5069 tree fndecl
= gimple_call_fndecl (stmt
);
5070 enum built_in_function builtin
= DECL_FUNCTION_CODE (fndecl
);
5074 case BUILT_IN_FABSF
:
5075 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_ABS
, stmt
, hbb
);
5079 case BUILT_IN_CEILF
:
5080 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
5083 case BUILT_IN_FLOOR
:
5084 case BUILT_IN_FLOORF
:
5085 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
5089 case BUILT_IN_RINTF
:
5090 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
5094 case BUILT_IN_SQRTF
:
5095 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
5098 case BUILT_IN_TRUNC
:
5099 case BUILT_IN_TRUNCF
:
5100 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
5107 /* HSAIL does not provide an instruction for double argument type. */
5108 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
5112 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
5115 case BUILT_IN_EXP2F
:
5116 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
5119 case BUILT_IN_LOG2F
:
5120 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
5124 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
5127 case BUILT_IN_CLRSB
:
5128 case BUILT_IN_CLRSBL
:
5129 case BUILT_IN_CLRSBLL
:
5130 gen_hsa_clrsb (call
, hbb
);
5135 case BUILT_IN_CLZLL
:
5136 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
5141 case BUILT_IN_CTZLL
:
5142 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
5147 case BUILT_IN_FFSLL
:
5148 gen_hsa_ffs (call
, hbb
);
5151 case BUILT_IN_PARITY
:
5152 case BUILT_IN_PARITYL
:
5153 case BUILT_IN_PARITYLL
:
5154 gen_hsa_parity (call
, hbb
);
5157 case BUILT_IN_POPCOUNT
:
5158 case BUILT_IN_POPCOUNTL
:
5159 case BUILT_IN_POPCOUNTLL
:
5160 gen_hsa_popcount (call
, hbb
);
5163 case BUILT_IN_ATOMIC_LOAD_1
:
5164 case BUILT_IN_ATOMIC_LOAD_2
:
5165 case BUILT_IN_ATOMIC_LOAD_4
:
5166 case BUILT_IN_ATOMIC_LOAD_8
:
5167 case BUILT_IN_ATOMIC_LOAD_16
:
5170 hsa_op_address
*addr
;
5171 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5172 tree model
= gimple_call_arg (stmt
, 1);
5173 if (!tree_fits_uhwi_p (model
))
5175 HSA_SORRY_ATV (gimple_location (stmt
),
5176 "support for HSA does not implement "
5182 unsigned HOST_WIDE_INT mmodel
= tree_to_uhwi (model
);
5183 BrigMemoryOrder memorder
= get_memory_order (mmodel
,
5184 gimple_location (stmt
));
5186 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
5187 memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
5189 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
5190 && memorder
!= BRIG_MEMORY_ORDER_SC_ACQUIRE
5191 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
5193 HSA_SORRY_ATV (gimple_location (stmt
),
5194 "support for HSA does not implement "
5195 "memory model for ATOMIC_LD: %s",
5196 get_memory_order_name (mmodel
));
5202 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
5204 mtype
= mem_type_for_type (t
);
5205 mtype
= hsa_bittype_for_type (mtype
);
5206 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5210 mtype
= BRIG_TYPE_B64
;
5211 dest
= new hsa_op_reg (mtype
);
5214 hsa_insn_atomic
*atominsn
5215 = new hsa_insn_atomic (2, BRIG_OPCODE_ATOMIC
, BRIG_ATOMIC_LD
, mtype
,
5216 memorder
, dest
, addr
);
5218 hbb
->append_insn (atominsn
);
5222 case BUILT_IN_ATOMIC_EXCHANGE_1
:
5223 case BUILT_IN_ATOMIC_EXCHANGE_2
:
5224 case BUILT_IN_ATOMIC_EXCHANGE_4
:
5225 case BUILT_IN_ATOMIC_EXCHANGE_8
:
5226 case BUILT_IN_ATOMIC_EXCHANGE_16
:
5227 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_EXCH
, stmt
, hbb
);
5230 case BUILT_IN_ATOMIC_FETCH_ADD_1
:
5231 case BUILT_IN_ATOMIC_FETCH_ADD_2
:
5232 case BUILT_IN_ATOMIC_FETCH_ADD_4
:
5233 case BUILT_IN_ATOMIC_FETCH_ADD_8
:
5234 case BUILT_IN_ATOMIC_FETCH_ADD_16
:
5235 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_ADD
, stmt
, hbb
);
5238 case BUILT_IN_ATOMIC_FETCH_SUB_1
:
5239 case BUILT_IN_ATOMIC_FETCH_SUB_2
:
5240 case BUILT_IN_ATOMIC_FETCH_SUB_4
:
5241 case BUILT_IN_ATOMIC_FETCH_SUB_8
:
5242 case BUILT_IN_ATOMIC_FETCH_SUB_16
:
5243 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_SUB
, stmt
, hbb
);
5246 case BUILT_IN_ATOMIC_FETCH_AND_1
:
5247 case BUILT_IN_ATOMIC_FETCH_AND_2
:
5248 case BUILT_IN_ATOMIC_FETCH_AND_4
:
5249 case BUILT_IN_ATOMIC_FETCH_AND_8
:
5250 case BUILT_IN_ATOMIC_FETCH_AND_16
:
5251 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_AND
, stmt
, hbb
);
5254 case BUILT_IN_ATOMIC_FETCH_XOR_1
:
5255 case BUILT_IN_ATOMIC_FETCH_XOR_2
:
5256 case BUILT_IN_ATOMIC_FETCH_XOR_4
:
5257 case BUILT_IN_ATOMIC_FETCH_XOR_8
:
5258 case BUILT_IN_ATOMIC_FETCH_XOR_16
:
5259 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_XOR
, stmt
, hbb
);
5262 case BUILT_IN_ATOMIC_FETCH_OR_1
:
5263 case BUILT_IN_ATOMIC_FETCH_OR_2
:
5264 case BUILT_IN_ATOMIC_FETCH_OR_4
:
5265 case BUILT_IN_ATOMIC_FETCH_OR_8
:
5266 case BUILT_IN_ATOMIC_FETCH_OR_16
:
5267 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_OR
, stmt
, hbb
);
5270 case BUILT_IN_ATOMIC_STORE_1
:
5271 case BUILT_IN_ATOMIC_STORE_2
:
5272 case BUILT_IN_ATOMIC_STORE_4
:
5273 case BUILT_IN_ATOMIC_STORE_8
:
5274 case BUILT_IN_ATOMIC_STORE_16
:
5275 /* Since there cannot be any LHS, the first parameter is meaningless. */
5276 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_ST
, stmt
, hbb
);
5279 case BUILT_IN_ATOMIC_ADD_FETCH_1
:
5280 case BUILT_IN_ATOMIC_ADD_FETCH_2
:
5281 case BUILT_IN_ATOMIC_ADD_FETCH_4
:
5282 case BUILT_IN_ATOMIC_ADD_FETCH_8
:
5283 case BUILT_IN_ATOMIC_ADD_FETCH_16
:
5284 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_ADD
, stmt
, hbb
);
5287 case BUILT_IN_ATOMIC_SUB_FETCH_1
:
5288 case BUILT_IN_ATOMIC_SUB_FETCH_2
:
5289 case BUILT_IN_ATOMIC_SUB_FETCH_4
:
5290 case BUILT_IN_ATOMIC_SUB_FETCH_8
:
5291 case BUILT_IN_ATOMIC_SUB_FETCH_16
:
5292 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_SUB
, stmt
, hbb
);
5295 case BUILT_IN_ATOMIC_AND_FETCH_1
:
5296 case BUILT_IN_ATOMIC_AND_FETCH_2
:
5297 case BUILT_IN_ATOMIC_AND_FETCH_4
:
5298 case BUILT_IN_ATOMIC_AND_FETCH_8
:
5299 case BUILT_IN_ATOMIC_AND_FETCH_16
:
5300 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_AND
, stmt
, hbb
);
5303 case BUILT_IN_ATOMIC_XOR_FETCH_1
:
5304 case BUILT_IN_ATOMIC_XOR_FETCH_2
:
5305 case BUILT_IN_ATOMIC_XOR_FETCH_4
:
5306 case BUILT_IN_ATOMIC_XOR_FETCH_8
:
5307 case BUILT_IN_ATOMIC_XOR_FETCH_16
:
5308 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_XOR
, stmt
, hbb
);
5311 case BUILT_IN_ATOMIC_OR_FETCH_1
:
5312 case BUILT_IN_ATOMIC_OR_FETCH_2
:
5313 case BUILT_IN_ATOMIC_OR_FETCH_4
:
5314 case BUILT_IN_ATOMIC_OR_FETCH_8
:
5315 case BUILT_IN_ATOMIC_OR_FETCH_16
:
5316 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_OR
, stmt
, hbb
);
5319 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_1
:
5320 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_2
:
5321 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_4
:
5322 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8
:
5323 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_16
:
5325 /* TODO: Use the appropriate memory model for now. */
5326 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
5329 = hsa_bittype_for_type (hsa_type_for_scalar_tree_type (type
, false));
5331 hsa_insn_atomic
*atominsn
5332 = new hsa_insn_atomic (4, BRIG_OPCODE_ATOMIC
, BRIG_ATOMIC_CAS
, atype
,
5333 BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
);
5334 hsa_op_address
*addr
;
5335 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5338 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5340 dest
= new hsa_op_reg (atype
);
5342 /* Should check what the memory scope is. */
5343 atominsn
->m_memoryscope
= BRIG_MEMORY_SCOPE_WORKGROUP
;
5344 atominsn
->set_op (0, dest
);
5345 atominsn
->set_op (1, addr
);
5347 hsa_op_with_type
*op
5348 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1), hbb
);
5349 atominsn
->set_op (2, op
);
5350 op
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 2), hbb
);
5351 atominsn
->set_op (3, op
);
5353 hbb
->append_insn (atominsn
);
5356 case BUILT_IN_GOMP_PARALLEL
:
5357 HSA_SORRY_AT (gimple_location (stmt
),
5358 "support for HSA does not implement non-gridified "
5359 "OpenMP parallel constructs.");
5361 case BUILT_IN_OMP_GET_THREAD_NUM
:
5363 query_hsa_grid (stmt
, BRIG_OPCODE_WORKITEMABSID
, 0, hbb
);
5367 case BUILT_IN_OMP_GET_NUM_THREADS
:
5369 query_hsa_grid (stmt
, BRIG_OPCODE_GRIDSIZE
, 0, hbb
);
5372 case BUILT_IN_GOMP_TEAMS
:
5374 gen_set_num_threads (gimple_call_arg (stmt
, 1), hbb
);
5377 case BUILT_IN_OMP_GET_NUM_TEAMS
:
5379 gen_get_num_teams (stmt
, hbb
);
5382 case BUILT_IN_OMP_GET_TEAM_NUM
:
5384 gen_get_team_num (stmt
, hbb
);
5387 case BUILT_IN_MEMCPY
:
5388 case BUILT_IN_MEMPCPY
:
5390 expand_memory_copy (stmt
, hbb
, builtin
);
5393 case BUILT_IN_MEMSET
:
5395 tree c
= gimple_call_arg (stmt
, 1);
5397 if (TREE_CODE (c
) != INTEGER_CST
)
5399 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5403 tree byte_size
= gimple_call_arg (stmt
, 2);
5405 if (!tree_fits_uhwi_p (byte_size
))
5407 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5411 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
5413 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5415 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5419 unsigned HOST_WIDE_INT constant
5420 = tree_to_uhwi (fold_convert (unsigned_char_type_node
, c
));
5422 expand_memory_set (stmt
, n
, constant
, hbb
, builtin
);
5426 case BUILT_IN_BZERO
:
5428 tree byte_size
= gimple_call_arg (stmt
, 1);
5430 if (!tree_fits_uhwi_p (byte_size
))
5432 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5436 unsigned HOST_WIDE_INT n
= tree_to_uhwi (byte_size
);
5438 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5440 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5444 expand_memory_set (stmt
, n
, 0, hbb
, builtin
);
5448 case BUILT_IN_ALLOCA
:
5449 case BUILT_IN_ALLOCA_WITH_ALIGN
:
5451 gen_hsa_alloca (call
, hbb
);
5456 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5462 /* Generate HSA instructions for a given gimple statement. Instructions will be
5466 gen_hsa_insns_for_gimple_stmt (gimple
*stmt
, hsa_bb
*hbb
)
5468 switch (gimple_code (stmt
))
5471 if (gimple_clobber_p (stmt
))
5474 if (gimple_assign_single_p (stmt
))
5476 tree lhs
= gimple_assign_lhs (stmt
);
5477 tree rhs
= gimple_assign_rhs1 (stmt
);
5478 gen_hsa_insns_for_single_assignment (lhs
, rhs
, hbb
);
5481 gen_hsa_insns_for_operation_assignment (stmt
, hbb
);
5484 gen_hsa_insns_for_return (as_a
<greturn
*> (stmt
), hbb
);
5487 gen_hsa_insns_for_cond_stmt (stmt
, hbb
);
5490 gen_hsa_insns_for_call (stmt
, hbb
);
5493 /* ??? HSA supports some debug facilities. */
5497 tree label
= gimple_label_label (as_a
<glabel
*> (stmt
));
5498 if (FORCED_LABEL (label
))
5499 HSA_SORRY_AT (gimple_location (stmt
),
5500 "support for HSA does not implement gimple label with "
5507 hbb
->append_insn (new hsa_insn_basic (0, BRIG_OPCODE_NOP
));
5512 gen_hsa_insns_for_switch_stmt (as_a
<gswitch
*> (stmt
), hbb
);
5516 HSA_SORRY_ATV (gimple_location (stmt
),
5517 "support for HSA does not implement gimple statement %s",
5518 gimple_code_name
[(int) gimple_code (stmt
)]);
5522 /* Generate a HSA PHI from a gimple PHI. */
5525 gen_hsa_phi_from_gimple_phi (gimple
*phi_stmt
, hsa_bb
*hbb
)
5528 unsigned count
= gimple_phi_num_args (phi_stmt
);
5531 = hsa_cfun
->reg_for_gimple_ssa (gimple_phi_result (phi_stmt
));
5532 hphi
= new hsa_insn_phi (count
, dest
);
5533 hphi
->m_bb
= hbb
->m_bb
;
5535 tree lhs
= gimple_phi_result (phi_stmt
);
5537 for (unsigned i
= 0; i
< count
; i
++)
5539 tree op
= gimple_phi_arg_def (phi_stmt
, i
);
5541 if (TREE_CODE (op
) == SSA_NAME
)
5543 hsa_op_reg
*hreg
= hsa_cfun
->reg_for_gimple_ssa (op
);
5544 hphi
->set_op (i
, hreg
);
5548 gcc_assert (is_gimple_min_invariant (op
));
5549 tree t
= TREE_TYPE (op
);
5550 if (!POINTER_TYPE_P (t
)
5551 || (TREE_CODE (op
) == STRING_CST
5552 && TREE_CODE (TREE_TYPE (t
)) == INTEGER_TYPE
))
5553 hphi
->set_op (i
, new hsa_op_immed (op
));
5554 else if (POINTER_TYPE_P (TREE_TYPE (lhs
))
5555 && TREE_CODE (op
) == INTEGER_CST
)
5557 /* Handle assignment of NULL value to a pointer type. */
5558 hphi
->set_op (i
, new hsa_op_immed (op
));
5560 else if (TREE_CODE (op
) == ADDR_EXPR
)
5562 edge e
= gimple_phi_arg_edge (as_a
<gphi
*> (phi_stmt
), i
);
5563 hsa_bb
*hbb_src
= hsa_init_new_bb (split_edge (e
));
5564 hsa_op_address
*addr
= gen_hsa_addr (TREE_OPERAND (op
, 0),
5568 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
5569 hsa_insn_basic
*insn
5570 = new hsa_insn_basic (2, BRIG_OPCODE_LDA
, BRIG_TYPE_U64
,
5572 hbb_src
->append_insn (insn
);
5574 hphi
->set_op (i
, dest
);
5578 HSA_SORRY_AT (gimple_location (phi_stmt
),
5579 "support for HSA does not handle PHI nodes with "
5580 "constant address operands");
5586 hphi
->m_prev
= hbb
->m_last_phi
;
5587 hphi
->m_next
= NULL
;
5588 if (hbb
->m_last_phi
)
5589 hbb
->m_last_phi
->m_next
= hphi
;
5590 hbb
->m_last_phi
= hphi
;
5591 if (!hbb
->m_first_phi
)
5592 hbb
->m_first_phi
= hphi
;
5595 /* Constructor of class containing HSA-specific information about a basic
5596 block. CFG_BB is the CFG BB this HSA BB is associated with. IDX is the new
5597 index of this BB (so that the constructor does not attempt to use
5598 hsa_cfun during its construction). */
5600 hsa_bb::hsa_bb (basic_block cfg_bb
, int idx
)
5601 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5602 m_last_phi (NULL
), m_index (idx
), m_liveout (BITMAP_ALLOC (NULL
)),
5603 m_livein (BITMAP_ALLOC (NULL
))
5605 gcc_assert (!cfg_bb
->aux
);
5609 /* Constructor of class containing HSA-specific information about a basic
5610 block. CFG_BB is the CFG BB this HSA BB is associated with. */
5612 hsa_bb::hsa_bb (basic_block cfg_bb
)
5613 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5614 m_last_phi (NULL
), m_index (hsa_cfun
->m_hbb_count
++),
5615 m_liveout (BITMAP_ALLOC (NULL
)), m_livein (BITMAP_ALLOC (NULL
))
5617 gcc_assert (!cfg_bb
->aux
);
5621 /* Destructor of class representing HSA BB. */
5625 BITMAP_FREE (m_livein
);
5626 BITMAP_FREE (m_liveout
);
5629 /* Create and initialize and return a new hsa_bb structure for a given CFG
5633 hsa_init_new_bb (basic_block bb
)
5635 return new (*hsa_allocp_bb
) hsa_bb (bb
);
5638 /* Initialize OMP in an HSA basic block PROLOGUE. */
5641 init_prologue (void)
5643 if (!hsa_cfun
->m_kern_p
)
5646 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5648 /* Create a magic number that is going to be printed by libgomp. */
5649 unsigned index
= hsa_get_number_decl_kernel_mappings ();
5651 /* Emit store to debug argument. */
5652 if (PARAM_VALUE (PARAM_HSA_GEN_DEBUG_STORES
) > 0)
5653 set_debug_value (prologue
, new hsa_op_immed (1000 + index
, BRIG_TYPE_U64
));
5656 /* Initialize hsa_num_threads to a default value. */
5659 init_hsa_num_threads (void)
5661 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5663 /* Save the default value to private variable hsa_num_threads. */
5664 hsa_insn_basic
*basic
5665 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
,
5666 new hsa_op_immed (0, hsa_num_threads
->m_type
),
5667 new hsa_op_address (hsa_num_threads
));
5668 prologue
->append_insn (basic
);
5671 /* Go over gimple representation and generate our internal HSA one. */
5674 gen_body_from_gimple ()
5678 /* Verify CFG for complex edges we are unable to handle. */
5682 FOR_EACH_BB_FN (bb
, cfun
)
5684 FOR_EACH_EDGE (e
, ei
, bb
->succs
)
5686 /* Verify all unsupported flags for edges that point
5687 to the same basic block. */
5688 if (e
->flags
& EDGE_EH
)
5690 HSA_SORRY_AT (UNKNOWN_LOCATION
,
5691 "support for HSA does not implement exception "
5698 FOR_EACH_BB_FN (bb
, cfun
)
5700 gimple_stmt_iterator gsi
;
5701 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5705 hbb
= hsa_init_new_bb (bb
);
5707 for (gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5709 gen_hsa_insns_for_gimple_stmt (gsi_stmt (gsi
), hbb
);
5710 if (hsa_seen_error ())
5715 FOR_EACH_BB_FN (bb
, cfun
)
5717 gimple_stmt_iterator gsi
;
5718 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5719 gcc_assert (hbb
!= NULL
);
5721 for (gsi
= gsi_start_phis (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5722 if (!virtual_operand_p (gimple_phi_result (gsi_stmt (gsi
))))
5723 gen_hsa_phi_from_gimple_phi (gsi_stmt (gsi
), hbb
);
5726 if (dump_file
&& (dump_flags
& TDF_DETAILS
))
5728 fprintf (dump_file
, "------- Generated SSA form -------\n");
5729 dump_hsa_cfun (dump_file
);
5734 gen_function_decl_parameters (hsa_function_representation
*f
,
5740 for (parm
= TYPE_ARG_TYPES (TREE_TYPE (decl
)), i
= 0;
5742 parm
= TREE_CHAIN (parm
), i
++)
5744 /* Result type if last in the tree list. */
5745 if (TREE_CHAIN (parm
) == NULL
)
5748 tree v
= TREE_VALUE (parm
);
5750 hsa_symbol
*arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5752 arg
->m_type
= hsa_type_for_tree_type (v
, &arg
->m_dim
);
5753 arg
->m_name_number
= i
;
5755 f
->m_input_args
.safe_push (arg
);
5758 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
5759 if (!VOID_TYPE_P (result_type
))
5761 f
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5763 f
->m_output_arg
->m_type
5764 = hsa_type_for_tree_type (result_type
, &f
->m_output_arg
->m_dim
);
5765 f
->m_output_arg
->m_name
= "res";
5769 /* Generate the vector of parameters of the HSA representation of the current
5770 function. This also includes the output parameter representing the
5774 gen_function_def_parameters ()
5778 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5780 for (parm
= DECL_ARGUMENTS (cfun
->decl
); parm
;
5781 parm
= DECL_CHAIN (parm
))
5783 struct hsa_symbol
**slot
;
5786 = new hsa_symbol (BRIG_TYPE_NONE
, hsa_cfun
->m_kern_p
5787 ? BRIG_SEGMENT_KERNARG
: BRIG_SEGMENT_ARG
,
5788 BRIG_LINKAGE_FUNCTION
);
5789 arg
->fillup_for_decl (parm
);
5791 hsa_cfun
->m_input_args
.safe_push (arg
);
5793 if (hsa_seen_error ())
5796 arg
->m_name
= hsa_get_declaration_name (parm
);
5798 /* Copy all input arguments and create corresponding private symbols
5800 hsa_symbol
*private_arg
;
5801 hsa_op_address
*parm_addr
= new hsa_op_address (arg
);
5803 if (TREE_ADDRESSABLE (parm
)
5804 || (!is_gimple_reg (parm
) && !TREE_READONLY (parm
)))
5806 private_arg
= hsa_cfun
->create_hsa_temporary (arg
->m_type
);
5807 private_arg
->fillup_for_decl (parm
);
5809 BrigAlignment8_t align
= MIN (arg
->m_align
, private_arg
->m_align
);
5811 hsa_op_address
*private_arg_addr
= new hsa_op_address (private_arg
);
5812 gen_hsa_memory_copy (prologue
, private_arg_addr
, parm_addr
,
5813 arg
->total_byte_size (), align
);
5818 slot
= hsa_cfun
->m_local_symbols
->find_slot (private_arg
, INSERT
);
5819 gcc_assert (!*slot
);
5820 *slot
= private_arg
;
5822 if (is_gimple_reg (parm
))
5824 tree ddef
= ssa_default_def (cfun
, parm
);
5825 if (ddef
&& !has_zero_uses (ddef
))
5827 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (ddef
),
5829 BrigType16_t mtype
= mem_type_for_type (t
);
5830 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (ddef
);
5831 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
,
5833 gcc_assert (!parm_addr
->m_reg
);
5834 prologue
->append_insn (mem
);
5839 if (!VOID_TYPE_P (TREE_TYPE (TREE_TYPE (cfun
->decl
))))
5841 struct hsa_symbol
**slot
;
5843 hsa_cfun
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5844 BRIG_LINKAGE_FUNCTION
);
5845 hsa_cfun
->m_output_arg
->fillup_for_decl (DECL_RESULT (cfun
->decl
));
5847 if (hsa_seen_error ())
5850 hsa_cfun
->m_output_arg
->m_name
= "res";
5851 slot
= hsa_cfun
->m_local_symbols
->find_slot (hsa_cfun
->m_output_arg
,
5853 gcc_assert (!*slot
);
5854 *slot
= hsa_cfun
->m_output_arg
;
5858 /* Generate function representation that corresponds to
5859 a function declaration. */
5861 hsa_function_representation
*
5862 hsa_generate_function_declaration (tree decl
)
5864 hsa_function_representation
*fun
5865 = new hsa_function_representation (decl
, false, 0);
5867 fun
->m_declaration_p
= true;
5868 fun
->m_name
= get_brig_function_name (decl
);
5869 gen_function_decl_parameters (fun
, decl
);
5875 /* Generate function representation that corresponds to
5878 hsa_function_representation
*
5879 hsa_generate_internal_fn_decl (hsa_internal_fn
*fn
)
5881 hsa_function_representation
*fun
= new hsa_function_representation (fn
);
5883 fun
->m_name
= fn
->name ();
5885 for (unsigned i
= 0; i
< fn
->get_arity (); i
++)
5888 = new hsa_symbol (fn
->get_argument_type (i
), BRIG_SEGMENT_ARG
,
5890 arg
->m_name_number
= i
;
5891 fun
->m_input_args
.safe_push (arg
);
5894 fun
->m_output_arg
= new hsa_symbol (fn
->get_argument_type (-1),
5895 BRIG_SEGMENT_ARG
, BRIG_LINKAGE_NONE
);
5896 fun
->m_output_arg
->m_name
= "res";
5901 /* Return true if switch statement S can be transformed
5902 to a SBR instruction in HSAIL. */
5905 transformable_switch_to_sbr_p (gswitch
*s
)
5907 /* Identify if a switch statement can be transformed to
5908 SBR instruction, like:
5910 sbr_u32 $s1 [@label1, @label2, @label3];
5913 tree size
= get_switch_size (s
);
5914 if (!tree_fits_uhwi_p (size
))
5917 if (tree_to_uhwi (size
) > HSA_MAXIMUM_SBR_LABELS
)
5923 /* Structure hold connection between PHI nodes and immediate
5924 values hold by there nodes. */
5926 struct phi_definition
5928 phi_definition (unsigned phi_i
, unsigned label_i
, tree imm
):
5929 phi_index (phi_i
), label_index (label_i
), phi_value (imm
)
5933 unsigned label_index
;
5937 /* Sum slice of a vector V, starting from index START and ending
5938 at the index END - 1. */
5940 template <typename T
>
5942 T
sum_slice (const auto_vec
<T
> &v
, unsigned start
, unsigned end
)
5946 for (unsigned i
= start
; i
< end
; i
++)
5952 /* Function transforms GIMPLE SWITCH statements to a series of IF statements.
5953 Let's assume following example:
5967 The transformation encompasses following steps:
5968 1) all immediate values used by edges coming from the switch basic block
5970 2) all these edges are removed
5971 3) the switch statement (in L0) is replaced by:
5977 4) newly created basic block Lx' is used for generation of
5979 5) else branch of the last condition goes to LD
5980 6) fix all immediate values in PHI nodes that were propagated though
5981 edges that were removed in step 2
5983 Note: if a case is made by a range C1..C2, then process
5984 following transformation:
5986 switch_cond_op1 = C1 <= index;
5987 switch_cond_op2 = index <= C2;
5988 switch_cond_and = switch_cond_op1 & switch_cond_op2;
5989 if (switch_cond_and != 0)
5997 convert_switch_statements (void)
5999 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
6002 bool modified_cfg
= false;
6004 FOR_EACH_BB_FN (bb
, func
)
6006 gimple_stmt_iterator gsi
= gsi_last_bb (bb
);
6007 if (gsi_end_p (gsi
))
6010 gimple
*stmt
= gsi_stmt (gsi
);
6012 if (gimple_code (stmt
) == GIMPLE_SWITCH
)
6014 gswitch
*s
= as_a
<gswitch
*> (stmt
);
6016 /* If the switch can utilize SBR insn, skip the statement. */
6017 if (transformable_switch_to_sbr_p (s
))
6020 modified_cfg
= true;
6022 unsigned labels
= gimple_switch_num_labels (s
);
6023 tree index
= gimple_switch_index (s
);
6024 tree index_type
= TREE_TYPE (index
);
6025 tree default_label
= gimple_switch_default_label (s
);
6026 basic_block default_label_bb
6027 = label_to_block_fn (func
, CASE_LABEL (default_label
));
6028 basic_block cur_bb
= bb
;
6030 auto_vec
<edge
> new_edges
;
6031 auto_vec
<phi_definition
*> phi_todo_list
;
6032 auto_vec
<gcov_type
> edge_counts
;
6033 auto_vec
<int> edge_probabilities
;
6035 /* Investigate all labels that and PHI nodes in these edges which
6036 should be fixed after we add new collection of edges. */
6037 for (unsigned i
= 0; i
< labels
; i
++)
6039 tree label
= gimple_switch_label (s
, i
);
6040 basic_block label_bb
= label_to_block_fn (func
, CASE_LABEL (label
));
6041 edge e
= find_edge (bb
, label_bb
);
6042 edge_counts
.safe_push (e
->count
);
6043 edge_probabilities
.safe_push (e
->probability
);
6044 gphi_iterator phi_gsi
;
6046 /* Save PHI definitions that will be destroyed because of an edge
6047 is going to be removed. */
6048 unsigned phi_index
= 0;
6049 for (phi_gsi
= gsi_start_phis (e
->dest
);
6050 !gsi_end_p (phi_gsi
); gsi_next (&phi_gsi
))
6052 gphi
*phi
= phi_gsi
.phi ();
6053 for (unsigned j
= 0; j
< gimple_phi_num_args (phi
); j
++)
6055 if (gimple_phi_arg_edge (phi
, j
) == e
)
6057 tree imm
= gimple_phi_arg_def (phi
, j
);
6058 phi_definition
*p
= new phi_definition (phi_index
, i
,
6060 phi_todo_list
.safe_push (p
);
6068 /* Remove all edges for the current basic block. */
6069 for (int i
= EDGE_COUNT (bb
->succs
) - 1; i
>= 0; i
--)
6071 edge e
= EDGE_SUCC (bb
, i
);
6075 /* Iterate all non-default labels. */
6076 for (unsigned i
= 1; i
< labels
; i
++)
6078 tree label
= gimple_switch_label (s
, i
);
6079 tree low
= CASE_LOW (label
);
6080 tree high
= CASE_HIGH (label
);
6082 if (!useless_type_conversion_p (TREE_TYPE (low
), index_type
))
6083 low
= fold_convert (index_type
, low
);
6085 gimple_stmt_iterator cond_gsi
= gsi_last_bb (cur_bb
);
6089 tree tmp1
= make_temp_ssa_name (boolean_type_node
, NULL
,
6092 gimple
*assign1
= gimple_build_assign (tmp1
, LE_EXPR
, low
,
6095 tree tmp2
= make_temp_ssa_name (boolean_type_node
, NULL
,
6098 if (!useless_type_conversion_p (TREE_TYPE (high
), index_type
))
6099 high
= fold_convert (index_type
, high
);
6100 gimple
*assign2
= gimple_build_assign (tmp2
, LE_EXPR
, index
,
6103 tree tmp3
= make_temp_ssa_name (boolean_type_node
, NULL
,
6105 gimple
*assign3
= gimple_build_assign (tmp3
, BIT_AND_EXPR
, tmp1
,
6108 gsi_insert_before (&cond_gsi
, assign1
, GSI_SAME_STMT
);
6109 gsi_insert_before (&cond_gsi
, assign2
, GSI_SAME_STMT
);
6110 gsi_insert_before (&cond_gsi
, assign3
, GSI_SAME_STMT
);
6112 tree b
= constant_boolean_node (false, boolean_type_node
);
6113 c
= gimple_build_cond (NE_EXPR
, tmp3
, b
, NULL
, NULL
);
6116 c
= gimple_build_cond (EQ_EXPR
, index
, low
, NULL
, NULL
);
6118 gimple_set_location (c
, gimple_location (stmt
));
6120 gsi_insert_before (&cond_gsi
, c
, GSI_SAME_STMT
);
6122 basic_block label_bb
6123 = label_to_block_fn (func
, CASE_LABEL (label
));
6124 edge new_edge
= make_edge (cur_bb
, label_bb
, EDGE_TRUE_VALUE
);
6125 int prob_sum
= sum_slice
<int> (edge_probabilities
, i
, labels
) +
6126 edge_probabilities
[0];
6129 new_edge
->probability
6130 = RDIV (REG_BR_PROB_BASE
* edge_probabilities
[i
], prob_sum
);
6132 new_edge
->count
= edge_counts
[i
];
6133 new_edges
.safe_push (new_edge
);
6137 /* Prepare another basic block that will contain
6139 basic_block next_bb
= create_empty_bb (cur_bb
);
6142 add_bb_to_loop (next_bb
, cur_bb
->loop_father
);
6143 loops_state_set (LOOPS_NEED_FIXUP
);
6146 edge next_edge
= make_edge (cur_bb
, next_bb
, EDGE_FALSE_VALUE
);
6147 next_edge
->probability
6148 = inverse_probability (new_edge
->probability
);
6149 next_edge
->count
= edge_counts
[0]
6150 + sum_slice
<gcov_type
> (edge_counts
, i
, labels
);
6151 next_bb
->frequency
= EDGE_FREQUENCY (next_edge
);
6154 else /* Link last IF statement and default label
6157 edge e
= make_edge (cur_bb
, default_label_bb
, EDGE_FALSE_VALUE
);
6158 e
->probability
= inverse_probability (new_edge
->probability
);
6159 e
->count
= edge_counts
[0];
6160 new_edges
.safe_insert (0, e
);
6164 /* Restore original PHI immediate value. */
6165 for (unsigned i
= 0; i
< phi_todo_list
.length (); i
++)
6167 phi_definition
*phi_def
= phi_todo_list
[i
];
6168 edge new_edge
= new_edges
[phi_def
->label_index
];
6170 gphi_iterator it
= gsi_start_phis (new_edge
->dest
);
6171 for (unsigned i
= 0; i
< phi_def
->phi_index
; i
++)
6174 gphi
*phi
= it
.phi ();
6175 add_phi_arg (phi
, phi_def
->phi_value
, new_edge
, UNKNOWN_LOCATION
);
6179 /* Remove the original GIMPLE switch statement. */
6180 gsi_remove (&gsi
, true);
6185 dump_function_to_file (current_function_decl
, dump_file
, TDF_DETAILS
);
6187 return modified_cfg
;
6190 /* Expand builtins that can't be handled by HSA back-end. */
6195 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
6198 FOR_EACH_BB_FN (bb
, func
)
6200 for (gimple_stmt_iterator gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
);
6203 gimple
*stmt
= gsi_stmt (gsi
);
6205 if (gimple_code (stmt
) != GIMPLE_CALL
)
6208 gcall
*call
= as_a
<gcall
*> (stmt
);
6210 if (!gimple_call_builtin_p (call
, BUILT_IN_NORMAL
))
6213 tree fndecl
= gimple_call_fndecl (stmt
);
6214 enum built_in_function fn
= DECL_FUNCTION_CODE (fndecl
);
6217 case BUILT_IN_CEXPF
:
6218 case BUILT_IN_CEXPIF
:
6219 case BUILT_IN_CEXPI
:
6221 /* Similar to builtins.c (expand_builtin_cexpi), the builtin
6222 can be transformed to: cexp(I * z) = ccos(z) + I * csin(z). */
6223 tree lhs
= gimple_call_lhs (stmt
);
6224 tree rhs
= gimple_call_arg (stmt
, 0);
6225 tree rhs_type
= TREE_TYPE (rhs
);
6226 bool float_type_p
= rhs_type
== float_type_node
;
6227 tree real_part
= make_temp_ssa_name (rhs_type
, NULL
,
6229 tree imag_part
= make_temp_ssa_name (rhs_type
, NULL
,
6233 = mathfn_built_in (rhs_type
, fn
== float_type_p
6234 ? BUILT_IN_COSF
: BUILT_IN_COS
);
6235 gcall
*cos
= gimple_build_call (cos_fndecl
, 1, rhs
);
6236 gimple_call_set_lhs (cos
, real_part
);
6237 gsi_insert_before (&gsi
, cos
, GSI_SAME_STMT
);
6240 = mathfn_built_in (rhs_type
, fn
== float_type_p
6241 ? BUILT_IN_SINF
: BUILT_IN_SIN
);
6242 gcall
*sin
= gimple_build_call (sin_fndecl
, 1, rhs
);
6243 gimple_call_set_lhs (sin
, imag_part
);
6244 gsi_insert_before (&gsi
, sin
, GSI_SAME_STMT
);
6247 gassign
*assign
= gimple_build_assign (lhs
, COMPLEX_EXPR
,
6248 real_part
, imag_part
);
6249 gsi_insert_before (&gsi
, assign
, GSI_SAME_STMT
);
6250 gsi_remove (&gsi
, true);
6261 /* Emit HSA module variables that are global for the entire module. */
6264 emit_hsa_module_variables (void)
6266 hsa_num_threads
= new hsa_symbol (BRIG_TYPE_U32
, BRIG_SEGMENT_PRIVATE
,
6267 BRIG_LINKAGE_MODULE
, true);
6269 hsa_num_threads
->m_name
= "hsa_num_threads";
6271 hsa_brig_emit_omp_symbols ();
6274 /* Generate HSAIL representation of the current function and write into a
6275 special section of the output file. If KERNEL is set, the function will be
6276 considered an HSA kernel callable from the host, otherwise it will be
6277 compiled as an HSA function callable from other HSA code. */
6280 generate_hsa (bool kernel
)
6282 hsa_init_data_for_cfun ();
6284 if (hsa_num_threads
== NULL
)
6285 emit_hsa_module_variables ();
6287 bool modified_cfg
= convert_switch_statements ();
6288 /* Initialize hsa_cfun. */
6289 hsa_cfun
= new hsa_function_representation (cfun
->decl
, kernel
,
6290 SSANAMES (cfun
)->length (),
6292 hsa_cfun
->init_extra_bbs ();
6296 HSA_SORRY_AT (UNKNOWN_LOCATION
,
6297 "support for HSA does not implement transactional memory");
6301 verify_function_arguments (cfun
->decl
);
6302 if (hsa_seen_error ())
6305 hsa_cfun
->m_name
= get_brig_function_name (cfun
->decl
);
6307 gen_function_def_parameters ();
6308 if (hsa_seen_error ())
6313 gen_body_from_gimple ();
6314 if (hsa_seen_error ())
6317 if (hsa_cfun
->m_kernel_dispatch_count
)
6318 init_hsa_num_threads ();
6320 if (hsa_cfun
->m_kern_p
)
6322 hsa_function_summary
*s
6323 = hsa_summaries
->get (cgraph_node::get (hsa_cfun
->m_decl
));
6324 hsa_add_kern_decl_mapping (current_function_decl
, hsa_cfun
->m_name
,
6325 hsa_cfun
->m_maximum_omp_data_size
,
6326 s
->m_gridified_kernel_p
);
6331 for (unsigned i
= 0; i
< hsa_cfun
->m_ssa_map
.length (); i
++)
6332 if (hsa_cfun
->m_ssa_map
[i
])
6333 hsa_cfun
->m_ssa_map
[i
]->verify_ssa ();
6336 FOR_EACH_BB_FN (bb
, cfun
)
6338 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
6340 for (hsa_insn_basic
*insn
= hbb
->m_first_insn
; insn
;
6341 insn
= insn
->m_next
)
6347 hsa_brig_emit_function ();
6350 hsa_deinit_data_for_cfun ();
6355 const pass_data pass_data_gen_hsail
=
6358 "hsagen", /* name */
6359 OPTGROUP_NONE
, /* optinfo_flags */
6360 TV_NONE
, /* tv_id */
6361 PROP_cfg
| PROP_ssa
, /* properties_required */
6362 0, /* properties_provided */
6363 0, /* properties_destroyed */
6364 0, /* todo_flags_start */
6365 0 /* todo_flags_finish */
6368 class pass_gen_hsail
: public gimple_opt_pass
6371 pass_gen_hsail (gcc::context
*ctxt
)
6372 : gimple_opt_pass(pass_data_gen_hsail
, ctxt
)
6375 /* opt_pass methods: */
6376 bool gate (function
*);
6377 unsigned int execute (function
*);
6379 }; // class pass_gen_hsail
6381 /* Determine whether or not to run generation of HSAIL. */
6384 pass_gen_hsail::gate (function
*f
)
6386 return hsa_gen_requested_p ()
6387 && hsa_gpu_implementation_p (f
->decl
);
6391 pass_gen_hsail::execute (function
*)
6393 hsa_function_summary
*s
6394 = hsa_summaries
->get (cgraph_node::get_create (current_function_decl
));
6397 generate_hsa (s
->m_kind
== HSA_KERNEL
);
6398 TREE_ASM_WRITTEN (current_function_decl
) = 1;
6399 return TODO_discard_function
;
6404 /* Create the instance of hsa gen pass. */
6407 make_pass_gen_hsail (gcc::context
*ctxt
)
6409 return new pass_gen_hsail (ctxt
);