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
)
170 hsa_symbol::hsa_symbol (BrigType16_t type
, BrigSegment8_t segment
,
171 BrigLinkage8_t linkage
, bool global_scope_p
,
172 BrigAllocation allocation
)
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
)
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. */
216 hsa_function_representation::hsa_function_representation
217 (tree fdecl
, bool kernel_p
, unsigned ssa_names_count
)
219 m_reg_count (0), m_input_args (vNULL
),
220 m_output_arg (NULL
), m_spill_symbols (vNULL
), m_global_symbols (vNULL
),
221 m_private_variables (vNULL
), m_called_functions (vNULL
),
222 m_called_internal_fns (vNULL
), m_hbb_count (0),
223 m_in_ssa (true), m_kern_p (kernel_p
), m_declaration_p (false),
224 m_decl (fdecl
), m_internal_fn (NULL
), m_shadow_reg (NULL
),
225 m_kernel_dispatch_count (0), m_maximum_omp_data_size (0),
226 m_seen_error (false), m_temp_symbol_count (0), m_ssa_map ()
228 int sym_init_len
= (vec_safe_length (cfun
->local_decls
) / 2) + 1;;
229 m_local_symbols
= new hash_table
<hsa_noop_symbol_hasher
> (sym_init_len
);
230 m_ssa_map
.safe_grow_cleared (ssa_names_count
);
233 /* Constructor of class representing HSA function information that
234 is derived for an internal function. */
235 hsa_function_representation::hsa_function_representation (hsa_internal_fn
*fn
)
236 : m_reg_count (0), m_input_args (vNULL
),
237 m_output_arg (NULL
), m_local_symbols (NULL
),
238 m_spill_symbols (vNULL
), m_global_symbols (vNULL
),
239 m_private_variables (vNULL
), m_called_functions (vNULL
),
240 m_called_internal_fns (vNULL
), m_hbb_count (0),
241 m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL
),
242 m_internal_fn (fn
), m_shadow_reg (NULL
), m_kernel_dispatch_count (0),
243 m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
246 /* Destructor of class holding function/kernel-wide information and state. */
248 hsa_function_representation::~hsa_function_representation ()
250 /* Kernel names are deallocated at the end of BRIG output when deallocating
251 hsa_decl_kernel_mapping. */
252 if (!m_kern_p
|| m_seen_error
)
255 for (unsigned i
= 0; i
< m_input_args
.length (); i
++)
256 delete m_input_args
[i
];
257 m_input_args
.release ();
260 delete m_local_symbols
;
262 for (unsigned i
= 0; i
< m_spill_symbols
.length (); i
++)
263 delete m_spill_symbols
[i
];
264 m_spill_symbols
.release ();
267 for (unsigned i
= 0; i
< m_global_symbols
.iterate (i
, &sym
); i
++)
268 if (sym
->m_linkage
!= BRIG_ALLOCATION_PROGRAM
)
270 m_global_symbols
.release ();
272 for (unsigned i
= 0; i
< m_private_variables
.length (); i
++)
273 delete m_private_variables
[i
];
274 m_private_variables
.release ();
275 m_called_functions
.release ();
276 m_ssa_map
.release ();
278 for (unsigned i
= 0; i
< m_called_internal_fns
.length (); i
++)
279 delete m_called_internal_fns
[i
];
283 hsa_function_representation::get_shadow_reg ()
285 /* If we compile a function with kernel dispatch and does not set
286 an optimization level, the function won't be inlined and
294 /* Append the shadow argument. */
295 hsa_symbol
*shadow
= new hsa_symbol (BRIG_TYPE_U64
, BRIG_SEGMENT_KERNARG
,
296 BRIG_LINKAGE_FUNCTION
);
297 m_input_args
.safe_push (shadow
);
298 shadow
->m_name
= "hsa_runtime_shadow";
300 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_U64
);
301 hsa_op_address
*addr
= new hsa_op_address (shadow
);
303 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, BRIG_TYPE_U64
, r
, addr
);
304 hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
))->append_insn (mem
);
310 bool hsa_function_representation::has_shadow_reg_p ()
312 return m_shadow_reg
!= NULL
;
316 hsa_function_representation::init_extra_bbs ()
318 hsa_init_new_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
319 hsa_init_new_bb (EXIT_BLOCK_PTR_FOR_FN (cfun
));
323 hsa_function_representation::create_hsa_temporary (BrigType16_t type
)
325 hsa_symbol
*s
= new hsa_symbol (type
, BRIG_SEGMENT_PRIVATE
,
326 BRIG_LINKAGE_FUNCTION
);
327 s
->m_name_number
= m_temp_symbol_count
++;
329 hsa_cfun
->m_private_variables
.safe_push (s
);
334 hsa_function_representation::get_linkage ()
337 return BRIG_LINKAGE_PROGRAM
;
339 return m_kern_p
|| TREE_PUBLIC (m_decl
) ?
340 BRIG_LINKAGE_PROGRAM
: BRIG_LINKAGE_MODULE
;
343 /* Hash map of simple OMP builtins. */
344 static hash_map
<nofree_string_hash
, omp_simple_builtin
> *omp_simple_builtins
347 /* Warning messages for OMP builtins. */
349 #define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP " \
351 #define HSA_WARN_TIMING_ROUTINE "support for HSA does not implement OpenMP " \
353 #define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \
354 "undefined semantics within target regions, support for HSA ignores them"
355 #define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP " \
358 /* Initialize hash map with simple OMP builtins. */
361 hsa_init_simple_builtins ()
363 if (omp_simple_builtins
!= NULL
)
367 = new hash_map
<nofree_string_hash
, omp_simple_builtin
> ();
369 omp_simple_builtin omp_builtins
[] =
371 omp_simple_builtin ("omp_get_initial_device", NULL
, false,
372 new hsa_op_immed (GOMP_DEVICE_HOST
,
373 (BrigType16_t
) BRIG_TYPE_S32
)),
374 omp_simple_builtin ("omp_is_initial_device", NULL
, false,
375 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
376 omp_simple_builtin ("omp_get_dynamic", NULL
, false,
377 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
378 omp_simple_builtin ("omp_set_dynamic", NULL
, false, NULL
),
379 omp_simple_builtin ("omp_init_lock", HSA_WARN_LOCK_ROUTINE
, true),
380 omp_simple_builtin ("omp_init_lock_with_hint", HSA_WARN_LOCK_ROUTINE
,
382 omp_simple_builtin ("omp_init_nest_lock_with_hint", HSA_WARN_LOCK_ROUTINE
,
384 omp_simple_builtin ("omp_destroy_lock", HSA_WARN_LOCK_ROUTINE
, true),
385 omp_simple_builtin ("omp_set_lock", HSA_WARN_LOCK_ROUTINE
, true),
386 omp_simple_builtin ("omp_unset_lock", HSA_WARN_LOCK_ROUTINE
, true),
387 omp_simple_builtin ("omp_test_lock", HSA_WARN_LOCK_ROUTINE
, true),
388 omp_simple_builtin ("omp_get_wtime", HSA_WARN_TIMING_ROUTINE
, true),
389 omp_simple_builtin ("omp_get_wtick", HSA_WARN_TIMING_ROUTINE
, true),
390 omp_simple_builtin ("omp_target_alloc", HSA_WARN_MEMORY_ROUTINE
, false,
391 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_U64
)),
392 omp_simple_builtin ("omp_target_free", HSA_WARN_MEMORY_ROUTINE
, false),
393 omp_simple_builtin ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE
,
395 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
396 omp_simple_builtin ("omp_target_memcpy", HSA_WARN_MEMORY_ROUTINE
, false,
397 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
398 omp_simple_builtin ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE
,
400 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
401 omp_simple_builtin ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE
,
403 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
404 omp_simple_builtin ("omp_target_disassociate_ptr",
405 HSA_WARN_MEMORY_ROUTINE
,
407 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
408 omp_simple_builtin ("omp_set_max_active_levels",
409 "Support for HSA only allows only one active level, "
410 "call to omp_set_max_active_levels will be ignored "
411 "in the generated HSAIL",
413 omp_simple_builtin ("omp_get_max_active_levels", NULL
, false,
414 new hsa_op_immed (1, (BrigType16_t
) BRIG_TYPE_S32
)),
415 omp_simple_builtin ("omp_in_final", NULL
, false,
416 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
417 omp_simple_builtin ("omp_get_proc_bind", HSA_WARN_AFFINITY
, false,
418 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
419 omp_simple_builtin ("omp_get_num_places", HSA_WARN_AFFINITY
, false,
420 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
421 omp_simple_builtin ("omp_get_place_num_procs", HSA_WARN_AFFINITY
, false,
422 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
423 omp_simple_builtin ("omp_get_place_proc_ids", HSA_WARN_AFFINITY
, false,
425 omp_simple_builtin ("omp_get_place_num", HSA_WARN_AFFINITY
, false,
426 new hsa_op_immed (-1, (BrigType16_t
) BRIG_TYPE_S32
)),
427 omp_simple_builtin ("omp_get_partition_num_places", HSA_WARN_AFFINITY
,
429 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
430 omp_simple_builtin ("omp_get_partition_place_nums", HSA_WARN_AFFINITY
,
432 omp_simple_builtin ("omp_set_default_device",
433 "omp_set_default_device has undefined semantics "
434 "within target regions, support for HSA ignores it",
436 omp_simple_builtin ("omp_get_default_device",
437 "omp_get_default_device has undefined semantics "
438 "within target regions, support for HSA ignores it",
440 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
441 omp_simple_builtin ("omp_get_num_devices",
442 "omp_get_num_devices has undefined semantics "
443 "within target regions, support for HSA ignores it",
445 new hsa_op_immed (0, (BrigType16_t
) BRIG_TYPE_S32
)),
446 omp_simple_builtin ("omp_get_num_procs", NULL
, true, NULL
),
447 omp_simple_builtin ("omp_get_cancellation", NULL
, true, NULL
),
448 omp_simple_builtin ("omp_set_nested", NULL
, true, NULL
),
449 omp_simple_builtin ("omp_get_nested", NULL
, true, NULL
),
450 omp_simple_builtin ("omp_set_schedule", NULL
, true, NULL
),
451 omp_simple_builtin ("omp_get_schedule", NULL
, true, NULL
),
452 omp_simple_builtin ("omp_get_thread_limit", NULL
, true, NULL
),
453 omp_simple_builtin ("omp_get_team_size", NULL
, true, NULL
),
454 omp_simple_builtin ("omp_get_ancestor_thread_num", NULL
, true, NULL
),
455 omp_simple_builtin ("omp_get_max_task_priority", NULL
, true, NULL
)
458 unsigned count
= sizeof (omp_builtins
) / sizeof (omp_simple_builtin
);
460 for (unsigned i
= 0; i
< count
; i
++)
461 omp_simple_builtins
->put (omp_builtins
[i
].m_name
, omp_builtins
[i
]);
464 /* Allocate HSA structures that we need only while generating with this. */
467 hsa_init_data_for_cfun ()
469 hsa_init_compilation_unit_data ();
470 hsa_allocp_operand_address
471 = new object_allocator
<hsa_op_address
> ("HSA address operands");
472 hsa_allocp_operand_immed
473 = new object_allocator
<hsa_op_immed
> ("HSA immediate operands");
474 hsa_allocp_operand_reg
475 = new object_allocator
<hsa_op_reg
> ("HSA register operands");
476 hsa_allocp_operand_code_list
477 = new object_allocator
<hsa_op_code_list
> ("HSA code list operands");
478 hsa_allocp_operand_operand_list
479 = new object_allocator
<hsa_op_operand_list
> ("HSA operand list operands");
480 hsa_allocp_inst_basic
481 = new object_allocator
<hsa_insn_basic
> ("HSA basic instructions");
483 = new object_allocator
<hsa_insn_phi
> ("HSA phi operands");
485 = new object_allocator
<hsa_insn_mem
> ("HSA memory instructions");
486 hsa_allocp_inst_atomic
487 = new object_allocator
<hsa_insn_atomic
> ("HSA atomic instructions");
488 hsa_allocp_inst_signal
489 = new object_allocator
<hsa_insn_signal
> ("HSA signal instructions");
491 = new object_allocator
<hsa_insn_seg
> ("HSA segment conversion "
494 = new object_allocator
<hsa_insn_cmp
> ("HSA comparison instructions");
496 = new object_allocator
<hsa_insn_br
> ("HSA branching instructions");
498 = new object_allocator
<hsa_insn_sbr
> ("HSA switch branching instructions");
500 = new object_allocator
<hsa_insn_call
> ("HSA call instructions");
501 hsa_allocp_inst_arg_block
502 = new object_allocator
<hsa_insn_arg_block
> ("HSA arg block instructions");
503 hsa_allocp_inst_comment
504 = new object_allocator
<hsa_insn_comment
> ("HSA comment instructions");
505 hsa_allocp_inst_queue
506 = new object_allocator
<hsa_insn_queue
> ("HSA queue instructions");
507 hsa_allocp_inst_srctype
508 = new object_allocator
<hsa_insn_srctype
> ("HSA source type instructions");
509 hsa_allocp_inst_packed
510 = new object_allocator
<hsa_insn_packed
> ("HSA packed instructions");
512 = new object_allocator
<hsa_insn_cvt
> ("HSA convert instructions");
513 hsa_allocp_inst_alloca
514 = new object_allocator
<hsa_insn_alloca
> ("HSA alloca instructions");
515 hsa_allocp_bb
= new object_allocator
<hsa_bb
> ("HSA basic blocks");
518 /* Deinitialize HSA subsystem and free all allocated memory. */
521 hsa_deinit_data_for_cfun (void)
525 FOR_ALL_BB_FN (bb
, cfun
)
528 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
533 for (unsigned int i
= 0; i
< hsa_operands
.length (); i
++)
534 hsa_destroy_operand (hsa_operands
[i
]);
536 hsa_operands
.release ();
538 for (unsigned i
= 0; i
< hsa_instructions
.length (); i
++)
539 hsa_destroy_insn (hsa_instructions
[i
]);
541 hsa_instructions
.release ();
543 if (omp_simple_builtins
!= NULL
)
545 delete omp_simple_builtins
;
546 omp_simple_builtins
= NULL
;
549 delete hsa_allocp_operand_address
;
550 delete hsa_allocp_operand_immed
;
551 delete hsa_allocp_operand_reg
;
552 delete hsa_allocp_operand_code_list
;
553 delete hsa_allocp_operand_operand_list
;
554 delete hsa_allocp_inst_basic
;
555 delete hsa_allocp_inst_phi
;
556 delete hsa_allocp_inst_atomic
;
557 delete hsa_allocp_inst_mem
;
558 delete hsa_allocp_inst_signal
;
559 delete hsa_allocp_inst_seg
;
560 delete hsa_allocp_inst_cmp
;
561 delete hsa_allocp_inst_br
;
562 delete hsa_allocp_inst_sbr
;
563 delete hsa_allocp_inst_call
;
564 delete hsa_allocp_inst_arg_block
;
565 delete hsa_allocp_inst_comment
;
566 delete hsa_allocp_inst_queue
;
567 delete hsa_allocp_inst_srctype
;
568 delete hsa_allocp_inst_packed
;
569 delete hsa_allocp_inst_cvt
;
570 delete hsa_allocp_inst_alloca
;
571 delete hsa_allocp_bb
;
575 /* Return the type which holds addresses in the given SEGMENT. */
578 hsa_get_segment_addr_type (BrigSegment8_t segment
)
582 case BRIG_SEGMENT_NONE
:
585 case BRIG_SEGMENT_FLAT
:
586 case BRIG_SEGMENT_GLOBAL
:
587 case BRIG_SEGMENT_READONLY
:
588 case BRIG_SEGMENT_KERNARG
:
589 return hsa_machine_large_p () ? BRIG_TYPE_U64
: BRIG_TYPE_U32
;
591 case BRIG_SEGMENT_GROUP
:
592 case BRIG_SEGMENT_PRIVATE
:
593 case BRIG_SEGMENT_SPILL
:
594 case BRIG_SEGMENT_ARG
:
595 return BRIG_TYPE_U32
;
600 /* Return integer brig type according to provided SIZE in bytes. If SIGN
601 is set to true, return signed integer type. */
604 get_integer_type_by_bytes (unsigned size
, bool sign
)
612 return BRIG_TYPE_S16
;
614 return BRIG_TYPE_S32
;
616 return BRIG_TYPE_S64
;
626 return BRIG_TYPE_U16
;
628 return BRIG_TYPE_U32
;
630 return BRIG_TYPE_U64
;
638 /* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
639 are assumed to use flat addressing. If min32int is true, always expand
640 integer types to one that has at least 32 bits. */
643 hsa_type_for_scalar_tree_type (const_tree type
, bool min32int
)
647 BrigType16_t res
= BRIG_TYPE_NONE
;
649 gcc_checking_assert (TYPE_P (type
));
650 gcc_checking_assert (!AGGREGATE_TYPE_P (type
));
651 if (POINTER_TYPE_P (type
))
652 return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
654 if (TREE_CODE (type
) == VECTOR_TYPE
|| TREE_CODE (type
) == COMPLEX_TYPE
)
655 base
= TREE_TYPE (type
);
659 if (!tree_fits_uhwi_p (TYPE_SIZE (base
)))
661 HSA_SORRY_ATV (EXPR_LOCATION (type
),
662 "support for HSA does not implement huge or "
663 "variable-sized type %T", type
);
667 bsize
= tree_to_uhwi (TYPE_SIZE (base
));
668 unsigned byte_size
= bsize
/ BITS_PER_UNIT
;
669 if (INTEGRAL_TYPE_P (base
))
670 res
= get_integer_type_by_bytes (byte_size
, !TYPE_UNSIGNED (base
));
671 else if (SCALAR_FLOAT_TYPE_P (base
))
689 if (res
== BRIG_TYPE_NONE
)
691 HSA_SORRY_ATV (EXPR_LOCATION (type
),
692 "support for HSA does not implement type %T", type
);
696 if (TREE_CODE (type
) == VECTOR_TYPE
)
698 HOST_WIDE_INT tsize
= tree_to_uhwi (TYPE_SIZE (type
));
702 HSA_SORRY_ATV (EXPR_LOCATION (type
),
703 "support for HSA does not implement a vector type "
704 "where a type and unit size are equal: %T", type
);
711 res
|= BRIG_TYPE_PACK_32
;
714 res
|= BRIG_TYPE_PACK_64
;
717 res
|= BRIG_TYPE_PACK_128
;
720 HSA_SORRY_ATV (EXPR_LOCATION (type
),
721 "support for HSA does not implement type %T", type
);
727 /* Registers/immediate operands can only be 32bit or more except for
729 if (res
== BRIG_TYPE_U8
|| res
== BRIG_TYPE_U16
)
731 else if (res
== BRIG_TYPE_S8
|| res
== BRIG_TYPE_S16
)
735 if (TREE_CODE (type
) == COMPLEX_TYPE
)
737 unsigned bsize
= 2 * hsa_type_bit_size (res
);
738 res
= hsa_bittype_for_bitsize (bsize
);
744 /* Returns the BRIG type we need to load/store entities of TYPE. */
747 mem_type_for_type (BrigType16_t type
)
749 /* HSA has non-intuitive constraints on load/store types. If it's
750 a bit-type it _must_ be B128, if it's not a bit-type it must be
751 64bit max. So for loading entities of 128 bits (e.g. vectors)
752 we have to to B128, while for loading the rest we have to use the
753 input type (??? or maybe also flattened to a equally sized non-vector
755 if ((type
& BRIG_TYPE_PACK_MASK
) == BRIG_TYPE_PACK_128
)
756 return BRIG_TYPE_B128
;
757 else if (hsa_btype_p (type
))
759 unsigned bitsize
= hsa_type_bit_size (type
);
761 return hsa_uint_for_bitsize (bitsize
);
766 /* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some
767 kind of array will be generated, setting DIM appropriately. Otherwise, it
768 will be set to zero. */
771 hsa_type_for_tree_type (const_tree type
, unsigned HOST_WIDE_INT
*dim_p
= NULL
,
772 bool min32int
= false)
774 gcc_checking_assert (TYPE_P (type
));
775 if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type
)))
777 HSA_SORRY_ATV (EXPR_LOCATION (type
), "support for HSA does not "
778 "implement huge or variable-sized type %T", type
);
779 return BRIG_TYPE_NONE
;
782 if (RECORD_OR_UNION_TYPE_P (type
))
785 *dim_p
= tree_to_uhwi (TYPE_SIZE_UNIT (type
));
786 return BRIG_TYPE_U8
| BRIG_TYPE_ARRAY
;
789 if (TREE_CODE (type
) == ARRAY_TYPE
)
791 /* We try to be nice and use the real base-type when this is an array of
792 scalars and only resort to an array of bytes if the type is more
795 unsigned HOST_WIDE_INT dim
= 1;
797 while (TREE_CODE (type
) == ARRAY_TYPE
)
799 tree domain
= TYPE_DOMAIN (type
);
800 if (!TYPE_MIN_VALUE (domain
)
801 || !TYPE_MAX_VALUE (domain
)
802 || !tree_fits_shwi_p (TYPE_MIN_VALUE (domain
))
803 || !tree_fits_shwi_p (TYPE_MAX_VALUE (domain
)))
805 HSA_SORRY_ATV (EXPR_LOCATION (type
),
806 "support for HSA does not implement array %T with "
807 "unknown bounds", type
);
808 return BRIG_TYPE_NONE
;
810 HOST_WIDE_INT min
= tree_to_shwi (TYPE_MIN_VALUE (domain
));
811 HOST_WIDE_INT max
= tree_to_shwi (TYPE_MAX_VALUE (domain
));
812 dim
= dim
* (unsigned HOST_WIDE_INT
) (max
- min
+ 1);
813 type
= TREE_TYPE (type
);
817 if (RECORD_OR_UNION_TYPE_P (type
))
819 dim
= dim
* tree_to_uhwi (TYPE_SIZE_UNIT (type
));
823 res
= hsa_type_for_scalar_tree_type (type
, false);
827 return res
| BRIG_TYPE_ARRAY
;
834 return hsa_type_for_scalar_tree_type (type
, min32int
);
837 /* Returns true if converting from STYPE into DTYPE needs the _CVT
838 opcode. If false a normal _MOV is enough. */
841 hsa_needs_cvt (BrigType16_t dtype
, BrigType16_t stype
)
843 if (hsa_btype_p (dtype
))
846 /* float <-> int conversions are real converts. */
847 if (hsa_type_float_p (dtype
) != hsa_type_float_p (stype
))
849 /* When both types have different size, then we need CVT as well. */
850 if (hsa_type_bit_size (dtype
) != hsa_type_bit_size (stype
))
855 /* Lookup or create the associated hsa_symbol structure with a given VAR_DECL
856 or lookup the hsa_structure corresponding to a PARM_DECL. */
859 get_symbol_for_decl (tree decl
)
862 hsa_symbol
dummy (BRIG_TYPE_NONE
, BRIG_SEGMENT_NONE
, BRIG_LINKAGE_NONE
);
864 gcc_assert (TREE_CODE (decl
) == PARM_DECL
865 || TREE_CODE (decl
) == RESULT_DECL
866 || TREE_CODE (decl
) == VAR_DECL
);
870 bool is_in_global_vars
871 = TREE_CODE (decl
) == VAR_DECL
&& is_global_var (decl
);
873 if (is_in_global_vars
)
874 slot
= hsa_global_variable_symbols
->find_slot (&dummy
, INSERT
);
876 slot
= hsa_cfun
->m_local_symbols
->find_slot (&dummy
, INSERT
);
878 gcc_checking_assert (slot
);
881 /* If the symbol is problematic, mark current function also as
883 if ((*slot
)->m_seen_error
)
891 gcc_assert (TREE_CODE (decl
) == VAR_DECL
);
893 if (is_in_global_vars
)
895 sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_GLOBAL
,
896 BRIG_LINKAGE_PROGRAM
, true,
897 BRIG_ALLOCATION_PROGRAM
);
898 hsa_cfun
->m_global_symbols
.safe_push (sym
);
902 /* PARM_DECL and RESULT_DECL should be already in m_local_symbols. */
903 gcc_assert (TREE_CODE (decl
) == VAR_DECL
);
905 sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_PRIVATE
,
906 BRIG_LINKAGE_FUNCTION
);
907 hsa_cfun
->m_private_variables
.safe_push (sym
);
910 sym
->fillup_for_decl (decl
);
911 sym
->m_name
= hsa_get_declaration_name (decl
);
918 /* For a given HSA function declaration, return a host
919 function declaration. */
922 hsa_get_host_function (tree decl
)
924 hsa_function_summary
*s
925 = hsa_summaries
->get (cgraph_node::get_create (decl
));
926 gcc_assert (s
->m_kind
!= HSA_NONE
);
927 gcc_assert (s
->m_gpu_implementation_p
);
929 return s
->m_binded_function
->decl
;
932 /* Return true if function DECL has a host equivalent function. */
935 get_brig_function_name (tree decl
)
939 hsa_function_summary
*s
= hsa_summaries
->get (cgraph_node::get_create (d
));
940 if (s
->m_kind
!= HSA_NONE
&& s
->m_gpu_implementation_p
)
941 d
= s
->m_binded_function
->decl
;
943 /* IPA split can create a function that has no host equivalent. */
947 char *name
= xstrdup (hsa_get_declaration_name (d
));
948 hsa_sanitize_name (name
);
953 /* Create a spill symbol of type TYPE. */
956 hsa_get_spill_symbol (BrigType16_t type
)
958 hsa_symbol
*sym
= new hsa_symbol (type
, BRIG_SEGMENT_SPILL
,
959 BRIG_LINKAGE_FUNCTION
);
960 hsa_cfun
->m_spill_symbols
.safe_push (sym
);
964 /* Create a symbol for a read-only string constant. */
966 hsa_get_string_cst_symbol (tree string_cst
)
968 gcc_checking_assert (TREE_CODE (string_cst
) == STRING_CST
);
970 hsa_symbol
**slot
= hsa_cfun
->m_string_constants_map
.get (string_cst
);
974 hsa_op_immed
*cst
= new hsa_op_immed (string_cst
);
975 hsa_symbol
*sym
= new hsa_symbol (cst
->m_type
, BRIG_SEGMENT_GLOBAL
,
976 BRIG_LINKAGE_MODULE
, true,
977 BRIG_ALLOCATION_AGENT
);
978 sym
->m_cst_value
= cst
;
979 sym
->m_dim
= TREE_STRING_LENGTH (string_cst
);
980 sym
->m_name_number
= hsa_cfun
->m_global_symbols
.length ();
982 hsa_cfun
->m_global_symbols
.safe_push (sym
);
983 hsa_cfun
->m_string_constants_map
.put (string_cst
, sym
);
987 /* Constructor of the ancestor of all operands. K is BRIG kind that identified
988 what the operator is. */
990 hsa_op_base::hsa_op_base (BrigKind16_t k
)
991 : m_next (NULL
), m_brig_op_offset (0), m_kind (k
)
993 hsa_operands
.safe_push (this);
996 /* Constructor of ancestor of all operands which have a type. K is BRIG kind
997 that identified what the operator is. T is the type of the operator. */
999 hsa_op_with_type::hsa_op_with_type (BrigKind16_t k
, BrigType16_t t
)
1000 : hsa_op_base (k
), m_type (t
)
1005 hsa_op_with_type::get_in_type (BrigType16_t dtype
, hsa_bb
*hbb
)
1007 if (m_type
== dtype
)
1012 if (hsa_needs_cvt (dtype
, m_type
))
1014 dest
= new hsa_op_reg (dtype
);
1015 hbb
->append_insn (new hsa_insn_cvt (dest
, this));
1019 dest
= new hsa_op_reg (m_type
);
1020 hbb
->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
1021 dest
->m_type
, dest
, this));
1023 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
1024 type of the operand must be same as type of the instruction. */
1025 dest
->m_type
= dtype
;
1031 /* Constructor of class representing HSA immediate values. TREE_VAL is the
1032 tree representation of the immediate value. If min32int is true,
1033 always expand integer types to one that has at least 32 bits. */
1035 hsa_op_immed::hsa_op_immed (tree tree_val
, bool min32int
)
1036 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
,
1037 hsa_type_for_tree_type (TREE_TYPE (tree_val
), NULL
,
1041 if (hsa_seen_error ())
1044 gcc_checking_assert ((is_gimple_min_invariant (tree_val
)
1045 && (!POINTER_TYPE_P (TREE_TYPE (tree_val
))
1046 || TREE_CODE (tree_val
) == INTEGER_CST
))
1047 || TREE_CODE (tree_val
) == CONSTRUCTOR
);
1048 m_tree_value
= tree_val
;
1049 m_brig_repr_size
= hsa_get_imm_brig_type_len (m_type
);
1051 if (TREE_CODE (m_tree_value
) == STRING_CST
)
1052 m_brig_repr_size
= TREE_STRING_LENGTH (m_tree_value
);
1053 else if (TREE_CODE (m_tree_value
) == CONSTRUCTOR
)
1056 = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (m_tree_value
)));
1058 /* Verify that all elements of a constructor are constants. */
1059 for (unsigned i
= 0;
1060 i
< vec_safe_length (CONSTRUCTOR_ELTS (m_tree_value
)); i
++)
1062 tree v
= CONSTRUCTOR_ELT (m_tree_value
, i
)->value
;
1063 if (!CONSTANT_CLASS_P (v
))
1065 HSA_SORRY_AT (EXPR_LOCATION (tree_val
),
1066 "HSA ctor should have only constants");
1072 emit_to_buffer (m_tree_value
);
1075 /* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
1076 integer representation of the immediate value. TYPE is BRIG type. */
1078 hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value
, BrigType16_t type
)
1079 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES
, type
),
1080 m_tree_value (NULL
), m_brig_repr (NULL
)
1082 gcc_assert (hsa_type_integer_p (type
));
1083 m_int_value
= integer_value
;
1084 m_brig_repr_size
= hsa_type_bit_size (type
) / BITS_PER_UNIT
;
1088 switch (m_brig_repr_size
)
1091 bytes
.b8
= (uint8_t) m_int_value
;
1094 bytes
.b16
= (uint16_t) m_int_value
;
1097 bytes
.b32
= (uint32_t) m_int_value
;
1100 bytes
.b64
= (uint64_t) m_int_value
;
1106 m_brig_repr
= XNEWVEC (char, m_brig_repr_size
);
1107 memcpy (m_brig_repr
, &bytes
, m_brig_repr_size
);
1110 hsa_op_immed::hsa_op_immed ()
1111 : hsa_op_with_type (BRIG_KIND_NONE
, BRIG_TYPE_NONE
), m_brig_repr (NULL
)
1115 /* New operator to allocate immediate operands from pool alloc. */
1118 hsa_op_immed::operator new (size_t)
1120 return hsa_allocp_operand_immed
->allocate_raw ();
1125 hsa_op_immed::~hsa_op_immed ()
1130 /* Change type of the immediate value to T. */
1133 hsa_op_immed::set_type (BrigType16_t t
)
1138 /* Constructor of class representing HSA registers and pseudo-registers. T is
1139 the BRIG type of the new register. */
1141 hsa_op_reg::hsa_op_reg (BrigType16_t t
)
1142 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER
, t
), m_gimple_ssa (NULL_TREE
),
1143 m_def_insn (NULL
), m_spill_sym (NULL
), m_order (hsa_cfun
->m_reg_count
++),
1144 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
1148 /* New operator to allocate a register from pool alloc. */
1151 hsa_op_reg::operator new (size_t)
1153 return hsa_allocp_operand_reg
->allocate_raw ();
1156 /* Verify register operand. */
1159 hsa_op_reg::verify_ssa ()
1161 /* Verify that each HSA register has a definition assigned.
1162 Exceptions are VAR_DECL and PARM_DECL that are a default
1164 gcc_checking_assert (m_def_insn
1165 || (m_gimple_ssa
!= NULL
1166 && (!SSA_NAME_VAR (m_gimple_ssa
)
1167 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa
))
1169 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa
)));
1171 /* Verify that every use of the register is really present
1172 in an instruction. */
1173 for (unsigned i
= 0; i
< m_uses
.length (); i
++)
1175 hsa_insn_basic
*use
= m_uses
[i
];
1177 bool is_visited
= false;
1178 for (unsigned j
= 0; j
< use
->operand_count (); j
++)
1180 hsa_op_base
*u
= use
->get_op (j
);
1181 hsa_op_address
*addr
; addr
= dyn_cast
<hsa_op_address
*> (u
);
1182 if (addr
&& addr
->m_reg
)
1187 bool r
= !addr
&& use
->op_output_p (j
);
1191 error ("HSA SSA name defined by instruction that is supposed "
1193 debug_hsa_operand (this);
1194 debug_hsa_insn (use
);
1195 internal_error ("HSA SSA verification failed");
1204 error ("HSA SSA name not among operands of instruction that is "
1205 "supposed to use it");
1206 debug_hsa_operand (this);
1207 debug_hsa_insn (use
);
1208 internal_error ("HSA SSA verification failed");
1213 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, hsa_op_reg
*r
,
1214 HOST_WIDE_INT offset
)
1215 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (r
),
1216 m_imm_offset (offset
)
1220 hsa_op_address::hsa_op_address (hsa_symbol
*sym
, HOST_WIDE_INT offset
)
1221 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (sym
), m_reg (NULL
),
1222 m_imm_offset (offset
)
1226 hsa_op_address::hsa_op_address (hsa_op_reg
*r
, HOST_WIDE_INT offset
)
1227 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS
), m_symbol (NULL
), m_reg (r
),
1228 m_imm_offset (offset
)
1232 /* New operator to allocate address operands from pool alloc. */
1235 hsa_op_address::operator new (size_t)
1237 return hsa_allocp_operand_address
->allocate_raw ();
1240 /* Constructor of an operand referring to HSAIL code. */
1242 hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF
),
1243 m_directive_offset (0)
1247 /* Constructor of an operand representing a code list. Set it up so that it
1248 can contain ELEMENTS number of elements. */
1250 hsa_op_code_list::hsa_op_code_list (unsigned elements
)
1251 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST
)
1253 m_offsets
.create (1);
1254 m_offsets
.safe_grow_cleared (elements
);
1257 /* New operator to allocate code list operands from pool alloc. */
1260 hsa_op_code_list::operator new (size_t)
1262 return hsa_allocp_operand_code_list
->allocate_raw ();
1265 /* Constructor of an operand representing an operand list.
1266 Set it up so that it can contain ELEMENTS number of elements. */
1268 hsa_op_operand_list::hsa_op_operand_list (unsigned elements
)
1269 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST
)
1271 m_offsets
.create (elements
);
1272 m_offsets
.safe_grow (elements
);
1275 /* New operator to allocate operand list operands from pool alloc. */
1278 hsa_op_operand_list::operator new (size_t)
1280 return hsa_allocp_operand_operand_list
->allocate_raw ();
1283 hsa_op_operand_list::~hsa_op_operand_list ()
1285 m_offsets
.release ();
1290 hsa_function_representation::reg_for_gimple_ssa (tree ssa
)
1294 gcc_checking_assert (TREE_CODE (ssa
) == SSA_NAME
);
1295 if (m_ssa_map
[SSA_NAME_VERSION (ssa
)])
1296 return m_ssa_map
[SSA_NAME_VERSION (ssa
)];
1298 hreg
= new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa
),
1300 hreg
->m_gimple_ssa
= ssa
;
1301 m_ssa_map
[SSA_NAME_VERSION (ssa
)] = hreg
;
1307 hsa_op_reg::set_definition (hsa_insn_basic
*insn
)
1309 if (hsa_cfun
->m_in_ssa
)
1311 gcc_checking_assert (!m_def_insn
);
1318 /* Constructor of the class which is the bases of all instructions and directly
1319 represents the most basic ones. NOPS is the number of operands that the
1320 operand vector will contain (and which will be cleared). OP is the opcode
1321 of the instruction. This constructor does not set type. */
1323 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
)
1325 m_next (NULL
), m_bb (NULL
), m_opcode (opc
), m_number (0),
1326 m_type (BRIG_TYPE_NONE
), m_brig_offset (0)
1329 m_operands
.safe_grow_cleared (nops
);
1331 hsa_instructions
.safe_push (this);
1334 /* Make OP the operand number INDEX of operands of this instruction. If OP is a
1335 register or an address containing a register, then either set the definition
1336 of the register to this instruction if it an output operand or add this
1337 instruction to the uses if it is an input one. */
1340 hsa_insn_basic::set_op (int index
, hsa_op_base
*op
)
1342 /* Each address operand is always use. */
1343 hsa_op_address
*addr
= dyn_cast
<hsa_op_address
*> (op
);
1344 if (addr
&& addr
->m_reg
)
1345 addr
->m_reg
->m_uses
.safe_push (this);
1348 hsa_op_reg
*reg
= dyn_cast
<hsa_op_reg
*> (op
);
1351 if (op_output_p (index
))
1352 reg
->set_definition (this);
1354 reg
->m_uses
.safe_push (this);
1358 m_operands
[index
] = op
;
1361 /* Get INDEX-th operand of the instruction. */
1364 hsa_insn_basic::get_op (int index
)
1366 return m_operands
[index
];
1369 /* Get address of INDEX-th operand of the instruction. */
1372 hsa_insn_basic::get_op_addr (int index
)
1374 return &m_operands
[index
];
1377 /* Get number of operands of the instruction. */
1379 hsa_insn_basic::operand_count ()
1381 return m_operands
.length ();
1384 /* Constructor of the class which is the bases of all instructions and directly
1385 represents the most basic ones. NOPS is the number of operands that the
1386 operand vector will contain (and which will be cleared). OPC is the opcode
1387 of the instruction, T is the type of the instruction. */
1389 hsa_insn_basic::hsa_insn_basic (unsigned nops
, int opc
, BrigType16_t t
,
1390 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1391 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1392 : m_prev (NULL
), m_next (NULL
), m_bb (NULL
), m_opcode (opc
),m_number (0),
1393 m_type (t
), m_brig_offset (0)
1396 m_operands
.safe_grow_cleared (nops
);
1400 gcc_checking_assert (nops
>= 1);
1406 gcc_checking_assert (nops
>= 2);
1412 gcc_checking_assert (nops
>= 3);
1418 gcc_checking_assert (nops
>= 4);
1422 hsa_instructions
.safe_push (this);
1425 /* New operator to allocate basic instruction from pool alloc. */
1428 hsa_insn_basic::operator new (size_t)
1430 return hsa_allocp_inst_basic
->allocate_raw ();
1433 /* Verify the instruction. */
1436 hsa_insn_basic::verify ()
1438 hsa_op_address
*addr
;
1441 /* Iterate all register operands and verify that the instruction
1442 is set in uses of the register. */
1443 for (unsigned i
= 0; i
< operand_count (); i
++)
1445 hsa_op_base
*use
= get_op (i
);
1447 if ((addr
= dyn_cast
<hsa_op_address
*> (use
)) && addr
->m_reg
)
1449 gcc_assert (addr
->m_reg
->m_def_insn
!= this);
1453 if ((reg
= dyn_cast
<hsa_op_reg
*> (use
)) && !op_output_p (i
))
1456 for (j
= 0; j
< reg
->m_uses
.length (); j
++)
1458 if (reg
->m_uses
[j
] == this)
1462 if (j
== reg
->m_uses
.length ())
1464 error ("HSA instruction uses a register but is not among "
1465 "recorded register uses");
1466 debug_hsa_operand (reg
);
1467 debug_hsa_insn (this);
1468 internal_error ("HSA instruction verification failed");
1474 /* Constructor of an instruction representing a PHI node. NOPS is the number
1475 of operands (equal to the number of predecessors). */
1477 hsa_insn_phi::hsa_insn_phi (unsigned nops
, hsa_op_reg
*dst
)
1478 : hsa_insn_basic (nops
, HSA_OPCODE_PHI
), m_dest (dst
)
1480 dst
->set_definition (this);
1483 /* New operator to allocate PHI instruction from pool alloc. */
1486 hsa_insn_phi::operator new (size_t)
1488 return hsa_allocp_inst_phi
->allocate_raw ();
1491 /* Constructor of class representing instruction for conditional jump, CTRL is
1492 the control register determining whether the jump will be carried out, the
1493 new instruction is automatically added to its uses list. */
1495 hsa_insn_br::hsa_insn_br (hsa_op_reg
*ctrl
)
1496 : hsa_insn_basic (1, BRIG_OPCODE_CBR
, BRIG_TYPE_B1
, ctrl
),
1497 m_width (BRIG_WIDTH_1
)
1501 /* New operator to allocate branch instruction from pool alloc. */
1504 hsa_insn_br::operator new (size_t)
1506 return hsa_allocp_inst_br
->allocate_raw ();
1509 /* Constructor of class representing instruction for switch jump, CTRL is
1510 the index register. */
1512 hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg
*index
, unsigned jump_count
)
1513 : hsa_insn_basic (1, BRIG_OPCODE_SBR
, BRIG_TYPE_B1
, index
),
1514 m_width (BRIG_WIDTH_1
), m_jump_table (vNULL
), m_default_bb (NULL
),
1515 m_label_code_list (new hsa_op_code_list (jump_count
))
1519 /* New operator to allocate switch branch instruction from pool alloc. */
1522 hsa_insn_sbr::operator new (size_t)
1524 return hsa_allocp_inst_sbr
->allocate_raw ();
1527 /* Replace all occurrences of OLD_BB with NEW_BB in the statements
1531 hsa_insn_sbr::replace_all_labels (basic_block old_bb
, basic_block new_bb
)
1533 for (unsigned i
= 0; i
< m_jump_table
.length (); i
++)
1534 if (m_jump_table
[i
] == old_bb
)
1535 m_jump_table
[i
] = new_bb
;
1538 hsa_insn_sbr::~hsa_insn_sbr ()
1540 m_jump_table
.release ();
1543 /* Constructor of comparison instruction. CMP is the comparison operation and T
1544 is the result type. */
1546 hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp
, BrigType16_t t
,
1547 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1549 : hsa_insn_basic (3 , BRIG_OPCODE_CMP
, t
, arg0
, arg1
, arg2
), m_compare (cmp
)
1553 /* New operator to allocate compare instruction from pool alloc. */
1556 hsa_insn_cmp::operator new (size_t)
1558 return hsa_allocp_inst_cmp
->allocate_raw ();
1561 /* Constructor of classes representing memory accesses. OPC is the opcode (must
1562 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
1563 operands are provided as ARG0 and ARG1. */
1565 hsa_insn_mem::hsa_insn_mem (int opc
, BrigType16_t t
, hsa_op_base
*arg0
,
1567 : hsa_insn_basic (2, opc
, t
, arg0
, arg1
),
1568 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1570 gcc_checking_assert (opc
== BRIG_OPCODE_LD
|| opc
== BRIG_OPCODE_ST
);
1573 /* Constructor for descendants allowing different opcodes and number of
1574 operands, it passes its arguments directly to hsa_insn_basic
1575 constructor. The instruction operands are provided as ARG[0-3]. */
1578 hsa_insn_mem::hsa_insn_mem (unsigned nops
, int opc
, BrigType16_t t
,
1579 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1580 hsa_op_base
*arg2
, hsa_op_base
*arg3
)
1581 : hsa_insn_basic (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
),
1582 m_align (hsa_natural_alignment (t
)), m_equiv_class (0)
1586 /* New operator to allocate memory instruction from pool alloc. */
1589 hsa_insn_mem::operator new (size_t)
1591 return hsa_allocp_inst_mem
->allocate_raw ();
1594 /* Constructor of class representing atomic instructions and signals. OPC is
1595 the principal opcode, aop is the specific atomic operation opcode. T is the
1596 type of the instruction. The instruction operands
1597 are provided as ARG[0-3]. */
1599 hsa_insn_atomic::hsa_insn_atomic (int nops
, int opc
,
1600 enum BrigAtomicOperation aop
,
1601 BrigType16_t t
, BrigMemoryOrder memorder
,
1603 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1605 : hsa_insn_mem (nops
, opc
, t
, arg0
, arg1
, arg2
, arg3
), m_atomicop (aop
),
1606 m_memoryorder (memorder
),
1607 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM
)
1609 gcc_checking_assert (opc
== BRIG_OPCODE_ATOMICNORET
||
1610 opc
== BRIG_OPCODE_ATOMIC
||
1611 opc
== BRIG_OPCODE_SIGNAL
||
1612 opc
== BRIG_OPCODE_SIGNALNORET
);
1615 /* New operator to allocate signal instruction from pool alloc. */
1618 hsa_insn_atomic::operator new (size_t)
1620 return hsa_allocp_inst_atomic
->allocate_raw ();
1623 /* Constructor of class representing signal instructions. OPC is the prinicpal
1624 opcode, sop is the specific signal operation opcode. T is the type of the
1625 instruction. The instruction operands are provided as ARG[0-3]. */
1627 hsa_insn_signal::hsa_insn_signal (int nops
, int opc
,
1628 enum BrigAtomicOperation sop
,
1629 BrigType16_t t
, hsa_op_base
*arg0
,
1630 hsa_op_base
*arg1
, hsa_op_base
*arg2
,
1632 : hsa_insn_atomic (nops
, opc
, sop
, t
, BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
,
1633 arg0
, arg1
, arg2
, arg3
)
1637 /* New operator to allocate signal instruction from pool alloc. */
1640 hsa_insn_signal::operator new (size_t)
1642 return hsa_allocp_inst_signal
->allocate_raw ();
1645 /* Constructor of class representing segment conversion instructions. OPC is
1646 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
1647 and SRCT are destination and source types respectively, SEG is the segment
1648 we are converting to or from. The instruction operands are
1649 provided as ARG0 and ARG1. */
1651 hsa_insn_seg::hsa_insn_seg (int opc
, BrigType16_t dest
, BrigType16_t srct
,
1652 BrigSegment8_t seg
, hsa_op_base
*arg0
,
1654 : hsa_insn_basic (2, opc
, dest
, arg0
, arg1
), m_src_type (srct
),
1657 gcc_checking_assert (opc
== BRIG_OPCODE_STOF
|| opc
== BRIG_OPCODE_FTOS
);
1660 /* New operator to allocate address conversion instruction from pool alloc. */
1663 hsa_insn_seg::operator new (size_t)
1665 return hsa_allocp_inst_seg
->allocate_raw ();
1668 /* Constructor of class representing a call instruction. CALLEE is the tree
1669 representation of the function being called. */
1671 hsa_insn_call::hsa_insn_call (tree callee
)
1672 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (callee
),
1673 m_output_arg (NULL
), m_args_code_list (NULL
), m_result_code_list (NULL
)
1677 hsa_insn_call::hsa_insn_call (hsa_internal_fn
*fn
)
1678 : hsa_insn_basic (0, BRIG_OPCODE_CALL
), m_called_function (NULL
),
1679 m_called_internal_fn (fn
), m_output_arg (NULL
), m_args_code_list (NULL
),
1680 m_result_code_list (NULL
)
1684 /* New operator to allocate call instruction from pool alloc. */
1687 hsa_insn_call::operator new (size_t)
1689 return hsa_allocp_inst_call
->allocate_raw ();
1692 hsa_insn_call::~hsa_insn_call ()
1694 for (unsigned i
= 0; i
< m_input_args
.length (); i
++)
1695 delete m_input_args
[i
];
1697 delete m_output_arg
;
1699 m_input_args
.release ();
1700 m_input_arg_insns
.release ();
1703 /* Constructor of class representing the argument block required to invoke
1705 hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind
,
1706 hsa_insn_call
* call
)
1707 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK
), m_kind (brig_kind
),
1712 /* New operator to allocate argument block instruction from pool alloc. */
1715 hsa_insn_arg_block::operator new (size_t)
1717 return hsa_allocp_inst_arg_block
->allocate_raw ();
1720 hsa_insn_comment::hsa_insn_comment (const char *s
)
1721 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT
)
1723 unsigned l
= strlen (s
);
1725 /* Append '// ' to the string. */
1726 char *buf
= XNEWVEC (char, l
+ 4);
1727 sprintf (buf
, "// %s", s
);
1731 /* New operator to allocate comment instruction from pool alloc. */
1734 hsa_insn_comment::operator new (size_t)
1736 return hsa_allocp_inst_comment
->allocate_raw ();
1739 hsa_insn_comment::~hsa_insn_comment ()
1741 gcc_checking_assert (m_comment
);
1746 /* Constructor of class representing the queue instruction in HSAIL. */
1747 hsa_insn_queue::hsa_insn_queue (int nops
, BrigOpcode opcode
)
1748 : hsa_insn_basic (nops
, opcode
, BRIG_TYPE_U64
)
1752 /* New operator to allocate source type instruction from pool alloc. */
1755 hsa_insn_srctype::operator new (size_t)
1757 return hsa_allocp_inst_srctype
->allocate_raw ();
1760 /* Constructor of class representing the source type instruction in HSAIL. */
1762 hsa_insn_srctype::hsa_insn_srctype (int nops
, BrigOpcode opcode
,
1763 BrigType16_t destt
, BrigType16_t srct
,
1764 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1765 hsa_op_base
*arg2
= NULL
)
1766 : hsa_insn_basic (nops
, opcode
, destt
, arg0
, arg1
, arg2
),
1767 m_source_type (srct
)
1770 /* New operator to allocate packed instruction from pool alloc. */
1773 hsa_insn_packed::operator new (size_t)
1775 return hsa_allocp_inst_packed
->allocate_raw ();
1778 /* Constructor of class representing the packed instruction in HSAIL. */
1780 hsa_insn_packed::hsa_insn_packed (int nops
, BrigOpcode opcode
,
1781 BrigType16_t destt
, BrigType16_t srct
,
1782 hsa_op_base
*arg0
, hsa_op_base
*arg1
,
1784 : hsa_insn_srctype (nops
, opcode
, destt
, srct
, arg0
, arg1
, arg2
)
1786 m_operand_list
= new hsa_op_operand_list (nops
- 1);
1789 /* New operator to allocate convert instruction from pool alloc. */
1792 hsa_insn_cvt::operator new (size_t)
1794 return hsa_allocp_inst_cvt
->allocate_raw ();
1797 /* Constructor of class representing the convert instruction in HSAIL. */
1799 hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type
*dest
, hsa_op_with_type
*src
)
1800 : hsa_insn_basic (2, BRIG_OPCODE_CVT
, dest
->m_type
, dest
, src
)
1804 /* New operator to allocate alloca from pool alloc. */
1807 hsa_insn_alloca::operator new (size_t)
1809 return hsa_allocp_inst_alloca
->allocate_raw ();
1812 /* Constructor of class representing the alloca in HSAIL. */
1814 hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type
*dest
,
1815 hsa_op_with_type
*size
, unsigned alignment
)
1816 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA
, dest
->m_type
, dest
, size
),
1817 m_align (BRIG_ALIGNMENT_8
)
1819 gcc_assert (dest
->m_type
== BRIG_TYPE_U32
);
1821 m_align
= hsa_alignment_encoding (alignment
);
1824 /* Append an instruction INSN into the basic block. */
1827 hsa_bb::append_insn (hsa_insn_basic
*insn
)
1829 gcc_assert (insn
->m_opcode
!= 0 || insn
->operand_count () == 0);
1830 gcc_assert (!insn
->m_bb
);
1833 insn
->m_prev
= m_last_insn
;
1834 insn
->m_next
= NULL
;
1836 m_last_insn
->m_next
= insn
;
1839 m_first_insn
= insn
;
1842 /* Insert HSA instruction NEW_INSN immediately before an existing instruction
1846 hsa_insert_insn_before (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1848 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1850 if (hbb
->m_first_insn
== old_insn
)
1851 hbb
->m_first_insn
= new_insn
;
1852 new_insn
->m_prev
= old_insn
->m_prev
;
1853 new_insn
->m_next
= old_insn
;
1854 if (old_insn
->m_prev
)
1855 old_insn
->m_prev
->m_next
= new_insn
;
1856 old_insn
->m_prev
= new_insn
;
1859 /* Append HSA instruction NEW_INSN immediately after an existing instruction
1863 hsa_append_insn_after (hsa_insn_basic
*new_insn
, hsa_insn_basic
*old_insn
)
1865 hsa_bb
*hbb
= hsa_bb_for_bb (old_insn
->m_bb
);
1867 if (hbb
->m_last_insn
== old_insn
)
1868 hbb
->m_last_insn
= new_insn
;
1869 new_insn
->m_prev
= old_insn
;
1870 new_insn
->m_next
= old_insn
->m_next
;
1871 if (old_insn
->m_next
)
1872 old_insn
->m_next
->m_prev
= new_insn
;
1873 old_insn
->m_next
= new_insn
;
1876 /* Return a register containing the calculated value of EXP which must be an
1877 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
1878 integer constants as returned by get_inner_reference.
1879 Newly generated HSA instructions will be appended to HBB.
1880 Perform all calculations in ADDRTYPE. */
1882 static hsa_op_with_type
*
1883 gen_address_calculation (tree exp
, hsa_bb
*hbb
, BrigType16_t addrtype
)
1887 if (TREE_CODE (exp
) == NOP_EXPR
)
1888 exp
= TREE_OPERAND (exp
, 0);
1890 switch (TREE_CODE (exp
))
1893 return hsa_cfun
->reg_for_gimple_ssa (exp
)->get_in_type (addrtype
, hbb
);
1897 hsa_op_immed
*imm
= new hsa_op_immed (exp
);
1898 if (addrtype
!= imm
->m_type
)
1899 imm
->m_type
= addrtype
;
1904 opcode
= BRIG_OPCODE_ADD
;
1908 opcode
= BRIG_OPCODE_MUL
;
1915 hsa_op_reg
*res
= new hsa_op_reg (addrtype
);
1916 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, addrtype
);
1917 insn
->set_op (0, res
);
1919 hsa_op_with_type
*op1
= gen_address_calculation (TREE_OPERAND (exp
, 0), hbb
,
1921 hsa_op_with_type
*op2
= gen_address_calculation (TREE_OPERAND (exp
, 1), hbb
,
1923 insn
->set_op (1, op1
);
1924 insn
->set_op (2, op2
);
1926 hbb
->append_insn (insn
);
1930 /* If R1 is NULL, just return R2, otherwise append an instruction adding them
1931 to HBB and return the register holding the result. */
1934 add_addr_regs_if_needed (hsa_op_reg
*r1
, hsa_op_reg
*r2
, hsa_bb
*hbb
)
1936 gcc_checking_assert (r2
);
1940 hsa_op_reg
*res
= new hsa_op_reg (r1
->m_type
);
1941 gcc_assert (!hsa_needs_cvt (r1
->m_type
, r2
->m_type
));
1942 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_ADD
, res
->m_type
);
1943 insn
->set_op (0, res
);
1944 insn
->set_op (1, r1
);
1945 insn
->set_op (2, r2
);
1946 hbb
->append_insn (insn
);
1950 /* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
1951 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
1954 process_mem_base (tree base
, hsa_symbol
**symbol
, BrigType16_t
*addrtype
,
1955 hsa_op_reg
**reg
, offset_int
*offset
, hsa_bb
*hbb
)
1957 if (TREE_CODE (base
) == SSA_NAME
)
1960 hsa_op_with_type
*ssa
1961 = hsa_cfun
->reg_for_gimple_ssa (base
)->get_in_type (*addrtype
, hbb
);
1962 *reg
= dyn_cast
<hsa_op_reg
*> (ssa
);
1964 else if (TREE_CODE (base
) == ADDR_EXPR
)
1966 tree decl
= TREE_OPERAND (base
, 0);
1968 if (!DECL_P (decl
) || TREE_CODE (decl
) == FUNCTION_DECL
)
1970 HSA_SORRY_AT (EXPR_LOCATION (base
),
1971 "support for HSA does not implement a memory reference "
1972 "to a non-declaration type");
1976 gcc_assert (!*symbol
);
1978 *symbol
= get_symbol_for_decl (decl
);
1979 *addrtype
= hsa_get_segment_addr_type ((*symbol
)->m_segment
);
1981 else if (TREE_CODE (base
) == INTEGER_CST
)
1982 *offset
+= wi::to_offset (base
);
1987 /* Forward declaration of a function. */
1990 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
);
1992 /* Generate HSA address operand for a given tree memory reference REF. If
1993 instructions need to be created to calculate the address, they will be added
1994 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
1995 the function assumes that the caller will handle possible
1996 bit-field references. Otherwise if we reference a bit-field, sorry message
1999 static hsa_op_address
*
2000 gen_hsa_addr (tree ref
, hsa_bb
*hbb
, HOST_WIDE_INT
*output_bitsize
= NULL
,
2001 HOST_WIDE_INT
*output_bitpos
= NULL
)
2003 hsa_symbol
*symbol
= NULL
;
2004 hsa_op_reg
*reg
= NULL
;
2005 offset_int offset
= 0;
2007 tree varoffset
= NULL_TREE
;
2008 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2009 HOST_WIDE_INT bitsize
= 0, bitpos
= 0;
2010 BrigType16_t flat_addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2012 if (TREE_CODE (ref
) == STRING_CST
)
2014 symbol
= hsa_get_string_cst_symbol (ref
);
2017 else if (TREE_CODE (ref
) == BIT_FIELD_REF
2018 && ((tree_to_uhwi (TREE_OPERAND (ref
, 1)) % BITS_PER_UNIT
) != 0
2019 || (tree_to_uhwi (TREE_OPERAND (ref
, 2)) % BITS_PER_UNIT
) != 0))
2021 HSA_SORRY_ATV (EXPR_LOCATION (origref
),
2022 "support for HSA does not implement "
2023 "bit field references such as %E", ref
);
2027 if (handled_component_p (ref
))
2029 enum machine_mode mode
;
2030 int unsignedp
, volatilep
, preversep
;
2032 ref
= get_inner_reference (ref
, &bitsize
, &bitpos
, &varoffset
, &mode
,
2033 &unsignedp
, &preversep
, &volatilep
, false);
2036 offset
= wi::rshift (offset
, LOG2_BITS_PER_UNIT
, SIGNED
);
2039 switch (TREE_CODE (ref
))
2043 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
2044 symbol
= hsa_cfun
->create_hsa_temporary (flat_addrtype
);
2045 hsa_op_reg
*r
= new hsa_op_reg (flat_addrtype
);
2046 gen_hsa_addr_insns (ref
, r
, hbb
);
2047 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
2048 r
, new hsa_op_address (symbol
)));
2054 addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
);
2055 symbol
= hsa_cfun
->create_hsa_temporary (flat_addrtype
);
2056 hsa_op_reg
*r
= hsa_cfun
->reg_for_gimple_ssa (ref
);
2058 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST
, r
->m_type
,
2059 r
, new hsa_op_address (symbol
)));
2066 gcc_assert (!symbol
);
2067 symbol
= get_symbol_for_decl (ref
);
2068 addrtype
= hsa_get_segment_addr_type (symbol
->m_segment
);
2072 process_mem_base (TREE_OPERAND (ref
, 0), &symbol
, &addrtype
, ®
,
2075 if (!integer_zerop (TREE_OPERAND (ref
, 1)))
2076 offset
+= wi::to_offset (TREE_OPERAND (ref
, 1));
2079 case TARGET_MEM_REF
:
2080 process_mem_base (TMR_BASE (ref
), &symbol
, &addrtype
, ®
, &offset
, hbb
);
2081 if (TMR_INDEX (ref
))
2084 hsa_op_base
*idx
= hsa_cfun
->reg_for_gimple_ssa
2085 (TMR_INDEX (ref
))->get_in_type (addrtype
, hbb
);
2086 if (TMR_STEP (ref
) && !integer_onep (TMR_STEP (ref
)))
2088 disp1
= new hsa_op_reg (addrtype
);
2089 hsa_insn_basic
*insn
= new hsa_insn_basic (3, BRIG_OPCODE_MUL
,
2092 /* As step must respect addrtype, we overwrite the type
2093 of an immediate value. */
2094 hsa_op_immed
*step
= new hsa_op_immed (TMR_STEP (ref
));
2095 step
->m_type
= addrtype
;
2097 insn
->set_op (0, disp1
);
2098 insn
->set_op (1, idx
);
2099 insn
->set_op (2, step
);
2100 hbb
->append_insn (insn
);
2103 disp1
= as_a
<hsa_op_reg
*> (idx
);
2104 reg
= add_addr_regs_if_needed (reg
, disp1
, hbb
);
2106 if (TMR_INDEX2 (ref
))
2108 hsa_op_base
*disp2
= hsa_cfun
->reg_for_gimple_ssa
2109 (TMR_INDEX2 (ref
))->get_in_type (addrtype
, hbb
);
2110 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (disp2
), hbb
);
2112 offset
+= wi::to_offset (TMR_OFFSET (ref
));
2115 HSA_SORRY_AT (EXPR_LOCATION (origref
),
2116 "support for HSA does not implement function pointers");
2119 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does "
2120 "not implement memory access to %E", origref
);
2126 if (TREE_CODE (varoffset
) == INTEGER_CST
)
2127 offset
+= wi::to_offset (varoffset
);
2130 hsa_op_base
*off_op
= gen_address_calculation (varoffset
, hbb
,
2132 reg
= add_addr_regs_if_needed (reg
, as_a
<hsa_op_reg
*> (off_op
),
2137 gcc_checking_assert ((symbol
2139 == hsa_get_segment_addr_type (symbol
->m_segment
))
2142 == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
)));
2144 HOST_WIDE_INT hwi_offset
= offset
.to_shwi ();
2146 /* Calculate remaining bitsize offset (if presented). */
2147 bitpos
%= BITS_PER_UNIT
;
2148 /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
2149 is not a reason to think this is a bit-field access. */
2151 && (bitsize
>= BITS_PER_UNIT
)
2152 && !(bitsize
& (bitsize
- 1)))
2155 if ((bitpos
|| bitsize
) && (output_bitpos
== NULL
|| output_bitsize
== NULL
))
2156 HSA_SORRY_ATV (EXPR_LOCATION (origref
), "support for HSA does not "
2157 "implement unhandled bit field reference such as %E", ref
);
2159 if (output_bitsize
!= NULL
&& output_bitpos
!= NULL
)
2161 *output_bitsize
= bitsize
;
2162 *output_bitpos
= bitpos
;
2165 return new hsa_op_address (symbol
, reg
, hwi_offset
);
2168 /* Generate HSA address for a function call argument of given TYPE.
2169 INDEX is used to generate corresponding name of the arguments.
2170 Special value -1 represents fact that result value is created. */
2172 static hsa_op_address
*
2173 gen_hsa_addr_for_arg (tree tree_type
, int index
)
2175 hsa_symbol
*sym
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
2177 sym
->m_type
= hsa_type_for_tree_type (tree_type
, &sym
->m_dim
);
2179 if (index
== -1) /* Function result. */
2180 sym
->m_name
= "res";
2181 else /* Function call arguments. */
2184 sym
->m_name_number
= index
;
2187 return new hsa_op_address (sym
);
2190 /* Generate HSA instructions that calculate address of VAL including all
2191 necessary conversions to flat addressing and place the result into DEST.
2192 Instructions are appended to HBB. */
2195 gen_hsa_addr_insns (tree val
, hsa_op_reg
*dest
, hsa_bb
*hbb
)
2197 /* Handle cases like tmp = NULL, where we just emit a move instruction
2199 if (TREE_CODE (val
) == INTEGER_CST
)
2201 hsa_op_immed
*c
= new hsa_op_immed (val
);
2202 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2203 dest
->m_type
, dest
, c
);
2204 hbb
->append_insn (insn
);
2208 hsa_op_address
*addr
;
2210 gcc_assert (dest
->m_type
== hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2211 if (TREE_CODE (val
) == ADDR_EXPR
)
2212 val
= TREE_OPERAND (val
, 0);
2213 addr
= gen_hsa_addr (val
, hbb
);
2214 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_LDA
);
2215 insn
->set_op (1, addr
);
2216 if (addr
->m_symbol
&& addr
->m_symbol
->m_segment
!= BRIG_SEGMENT_GLOBAL
)
2218 /* LDA produces segment-relative address, we need to convert
2219 it to the flat one. */
2221 tmp
= new hsa_op_reg (hsa_get_segment_addr_type
2222 (addr
->m_symbol
->m_segment
));
2224 seg
= new hsa_insn_seg (BRIG_OPCODE_STOF
,
2225 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
2226 tmp
->m_type
, addr
->m_symbol
->m_segment
, dest
,
2229 insn
->set_op (0, tmp
);
2230 insn
->m_type
= tmp
->m_type
;
2231 hbb
->append_insn (insn
);
2232 hbb
->append_insn (seg
);
2236 insn
->set_op (0, dest
);
2237 insn
->m_type
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
2238 hbb
->append_insn (insn
);
2242 /* Return an HSA register or HSA immediate value operand corresponding to
2243 gimple operand OP. */
2245 static hsa_op_with_type
*
2246 hsa_reg_or_immed_for_gimple_op (tree op
, hsa_bb
*hbb
)
2250 if (TREE_CODE (op
) == SSA_NAME
)
2251 tmp
= hsa_cfun
->reg_for_gimple_ssa (op
);
2252 else if (!POINTER_TYPE_P (TREE_TYPE (op
)))
2253 return new hsa_op_immed (op
);
2256 tmp
= new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
));
2257 gen_hsa_addr_insns (op
, tmp
, hbb
);
2262 /* Create a simple movement instruction with register destination DEST and
2263 register or immediate source SRC and append it to the end of HBB. */
2266 hsa_build_append_simple_mov (hsa_op_reg
*dest
, hsa_op_base
*src
, hsa_bb
*hbb
)
2268 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
,
2270 if (hsa_op_reg
*sreg
= dyn_cast
<hsa_op_reg
*> (src
))
2271 gcc_assert (hsa_type_bit_size (dest
->m_type
)
2272 == hsa_type_bit_size (sreg
->m_type
));
2274 gcc_assert (hsa_type_bit_size (dest
->m_type
)
2275 == hsa_type_bit_size (as_a
<hsa_op_immed
*> (src
)->m_type
));
2277 hbb
->append_insn (insn
);
2280 /* Generate HSAIL instructions loading a bit field into register DEST.
2281 VALUE_REG is a register of a SSA name that is used in the bit field
2282 reference. To identify a bit field BITPOS is offset to the loaded memory
2283 and BITSIZE is number of bits of the bit field.
2284 Add instructions to HBB. */
2287 gen_hsa_insns_for_bitfield (hsa_op_reg
*dest
, hsa_op_reg
*value_reg
,
2288 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2291 unsigned type_bitsize
= hsa_type_bit_size (dest
->m_type
);
2292 unsigned left_shift
= type_bitsize
- (bitsize
+ bitpos
);
2293 unsigned right_shift
= left_shift
+ bitpos
;
2297 hsa_op_reg
*value_reg_2
= new hsa_op_reg (dest
->m_type
);
2298 hsa_op_immed
*c
= new hsa_op_immed (left_shift
, BRIG_TYPE_U32
);
2300 hsa_insn_basic
*lshift
2301 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, value_reg_2
->m_type
,
2302 value_reg_2
, value_reg
, c
);
2304 hbb
->append_insn (lshift
);
2306 value_reg
= value_reg_2
;
2311 hsa_op_reg
*value_reg_2
= new hsa_op_reg (dest
->m_type
);
2312 hsa_op_immed
*c
= new hsa_op_immed (right_shift
, BRIG_TYPE_U32
);
2314 hsa_insn_basic
*rshift
2315 = new hsa_insn_basic (3, BRIG_OPCODE_SHR
, value_reg_2
->m_type
,
2316 value_reg_2
, value_reg
, c
);
2318 hbb
->append_insn (rshift
);
2320 value_reg
= value_reg_2
;
2323 hsa_insn_basic
*assignment
2324 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, dest
, value_reg
);
2325 hbb
->append_insn (assignment
);
2329 /* Generate HSAIL instructions loading a bit field into register DEST. ADDR is
2330 prepared memory address which is used to load the bit field. To identify a
2331 bit field BITPOS is offset to the loaded memory and BITSIZE is number of
2332 bits of the bit field. Add instructions to HBB. Load must be performed in
2336 gen_hsa_insns_for_bitfield_load (hsa_op_reg
*dest
, hsa_op_address
*addr
,
2337 HOST_WIDE_INT bitsize
, HOST_WIDE_INT bitpos
,
2338 hsa_bb
*hbb
, BrigAlignment8_t align
)
2340 hsa_op_reg
*value_reg
= new hsa_op_reg (dest
->m_type
);
2341 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, dest
->m_type
, value_reg
,
2343 mem
->set_align (align
);
2344 hbb
->append_insn (mem
);
2345 gen_hsa_insns_for_bitfield (dest
, value_reg
, bitsize
, bitpos
, hbb
);
2348 /* Return the alignment of base memory accesses we issue to perform bit-field
2349 memory access REF. */
2351 static BrigAlignment8_t
2352 hsa_bitmemref_alignment (tree ref
)
2354 unsigned HOST_WIDE_INT bit_offset
= 0;
2358 if (TREE_CODE (ref
) == BIT_FIELD_REF
)
2360 if (!tree_fits_uhwi_p (TREE_OPERAND (ref
, 2)))
2361 return BRIG_ALIGNMENT_1
;
2362 bit_offset
+= tree_to_uhwi (TREE_OPERAND (ref
, 2));
2364 else if (TREE_CODE (ref
) == COMPONENT_REF
2365 && DECL_BIT_FIELD (TREE_OPERAND (ref
, 1)))
2366 bit_offset
+= int_bit_position (TREE_OPERAND (ref
, 1));
2369 ref
= TREE_OPERAND (ref
, 0);
2372 unsigned HOST_WIDE_INT bits
= bit_offset
% BITS_PER_UNIT
;
2373 unsigned HOST_WIDE_INT byte_bits
= bit_offset
- bits
;
2374 BrigAlignment8_t base
= hsa_alignment_encoding (get_object_alignment (ref
));
2377 return MIN (base
, hsa_alignment_encoding (byte_bits
& -byte_bits
));
2380 /* Generate HSAIL instructions loading something into register DEST. RHS is
2381 tree representation of the loaded data, which are loaded as type TYPE. Add
2382 instructions to HBB. */
2385 gen_hsa_insns_for_load (hsa_op_reg
*dest
, tree rhs
, tree type
, hsa_bb
*hbb
)
2387 /* The destination SSA name will give us the type. */
2388 if (TREE_CODE (rhs
) == VIEW_CONVERT_EXPR
)
2389 rhs
= TREE_OPERAND (rhs
, 0);
2391 if (TREE_CODE (rhs
) == SSA_NAME
)
2393 hsa_op_reg
*src
= hsa_cfun
->reg_for_gimple_ssa (rhs
);
2394 hsa_build_append_simple_mov (dest
, src
, hbb
);
2396 else if (is_gimple_min_invariant (rhs
)
2397 || TREE_CODE (rhs
) == ADDR_EXPR
)
2399 if (POINTER_TYPE_P (TREE_TYPE (rhs
)))
2401 if (dest
->m_type
!= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
))
2403 HSA_SORRY_ATV (EXPR_LOCATION (rhs
),
2404 "support for HSA does not implement conversion "
2405 "of %E to the requested non-pointer type.", rhs
);
2409 gen_hsa_addr_insns (rhs
, dest
, hbb
);
2411 else if (TREE_CODE (rhs
) == COMPLEX_CST
)
2413 hsa_op_immed
*real_part
= new hsa_op_immed (TREE_REALPART (rhs
));
2414 hsa_op_immed
*imag_part
= new hsa_op_immed (TREE_IMAGPART (rhs
));
2416 hsa_op_reg
*real_part_reg
2417 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2419 hsa_op_reg
*imag_part_reg
2420 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type
),
2423 hsa_build_append_simple_mov (real_part_reg
, real_part
, hbb
);
2424 hsa_build_append_simple_mov (imag_part_reg
, imag_part
, hbb
);
2426 BrigType16_t src_type
= hsa_bittype_for_type (real_part_reg
->m_type
);
2428 hsa_insn_packed
*insn
2429 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
,
2430 src_type
, dest
, real_part_reg
,
2432 hbb
->append_insn (insn
);
2436 hsa_op_immed
*imm
= new hsa_op_immed (rhs
);
2437 hsa_build_append_simple_mov (dest
, imm
, hbb
);
2440 else if (TREE_CODE (rhs
) == REALPART_EXPR
|| TREE_CODE (rhs
) == IMAGPART_EXPR
)
2442 tree pack_type
= TREE_TYPE (TREE_OPERAND (rhs
, 0));
2444 hsa_op_reg
*packed_reg
2445 = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type
, true));
2447 tree complex_rhs
= TREE_OPERAND (rhs
, 0);
2448 gen_hsa_insns_for_load (packed_reg
, complex_rhs
, TREE_TYPE (complex_rhs
),
2451 hsa_op_reg
*real_reg
2452 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2454 hsa_op_reg
*imag_reg
2455 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type
, true));
2457 BrigKind16_t brig_type
= packed_reg
->m_type
;
2458 hsa_insn_packed
*packed
2459 = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND
,
2460 hsa_bittype_for_type (real_reg
->m_type
),
2461 brig_type
, real_reg
, imag_reg
, packed_reg
);
2463 hbb
->append_insn (packed
);
2465 hsa_op_reg
*source
= TREE_CODE (rhs
) == REALPART_EXPR
?
2466 real_reg
: imag_reg
;
2468 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
2469 dest
->m_type
, dest
, source
);
2471 hbb
->append_insn (insn
);
2473 else if (TREE_CODE (rhs
) == BIT_FIELD_REF
2474 && TREE_CODE (TREE_OPERAND (rhs
, 0)) == SSA_NAME
)
2476 tree ssa_name
= TREE_OPERAND (rhs
, 0);
2477 HOST_WIDE_INT bitsize
= tree_to_uhwi (TREE_OPERAND (rhs
, 1));
2478 HOST_WIDE_INT bitpos
= tree_to_uhwi (TREE_OPERAND (rhs
, 2));
2480 hsa_op_reg
*imm_value
= hsa_cfun
->reg_for_gimple_ssa (ssa_name
);
2481 gen_hsa_insns_for_bitfield (dest
, imm_value
, bitsize
, bitpos
, hbb
);
2483 else if (DECL_P (rhs
) || TREE_CODE (rhs
) == MEM_REF
2484 || TREE_CODE (rhs
) == TARGET_MEM_REF
2485 || handled_component_p (rhs
))
2487 HOST_WIDE_INT bitsize
, bitpos
;
2489 /* Load from memory. */
2490 hsa_op_address
*addr
;
2491 addr
= gen_hsa_addr (rhs
, hbb
, &bitsize
, &bitpos
);
2493 /* Handle load of a bit field. */
2496 HSA_SORRY_AT (EXPR_LOCATION (rhs
),
2497 "support for HSA does not implement load from a bit "
2498 "field bigger than 64 bits");
2502 if (bitsize
|| bitpos
)
2503 gen_hsa_insns_for_bitfield_load (dest
, addr
, bitsize
, bitpos
, hbb
,
2504 hsa_bitmemref_alignment (rhs
));
2508 /* Not dest->m_type, that's possibly extended. */
2509 mtype
= mem_type_for_type (hsa_type_for_scalar_tree_type (type
,
2511 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
, dest
,
2513 mem
->set_align (hsa_alignment_encoding (get_object_alignment (rhs
)));
2514 hbb
->append_insn (mem
);
2518 HSA_SORRY_ATV (EXPR_LOCATION (rhs
),
2519 "support for HSA does not implement loading "
2524 /* Return number of bits necessary for representation of a bit field,
2525 starting at BITPOS with size of BITSIZE. */
2528 get_bitfield_size (unsigned bitpos
, unsigned bitsize
)
2530 unsigned s
= bitpos
+ bitsize
;
2531 unsigned sizes
[] = {8, 16, 32, 64};
2533 for (unsigned i
= 0; i
< 4; i
++)
2541 /* Generate HSAIL instructions storing into memory. LHS is the destination of
2542 the store, SRC is the source operand. Add instructions to HBB. */
2545 gen_hsa_insns_for_store (tree lhs
, hsa_op_base
*src
, hsa_bb
*hbb
)
2547 HOST_WIDE_INT bitsize
= 0, bitpos
= 0;
2548 BrigAlignment8_t req_align
;
2550 mtype
= mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
2552 hsa_op_address
*addr
;
2553 addr
= gen_hsa_addr (lhs
, hbb
, &bitsize
, &bitpos
);
2555 /* Handle store to a bit field. */
2558 HSA_SORRY_AT (EXPR_LOCATION (lhs
),
2559 "support for HSA does not implement store to a bit field "
2560 "bigger than 64 bits");
2564 unsigned type_bitsize
= get_bitfield_size (bitpos
, bitsize
);
2566 /* HSAIL does not support MOV insn with 16-bits integers. */
2567 if (type_bitsize
< 32)
2570 if (bitpos
|| (bitsize
&& type_bitsize
!= bitsize
))
2572 unsigned HOST_WIDE_INT mask
= 0;
2573 BrigType16_t mem_type
2574 = get_integer_type_by_bytes (type_bitsize
/ BITS_PER_UNIT
,
2575 !TYPE_UNSIGNED (TREE_TYPE (lhs
)));
2577 for (unsigned i
= 0; i
< type_bitsize
; i
++)
2578 if (i
< bitpos
|| i
>= bitpos
+ bitsize
)
2579 mask
|= ((unsigned HOST_WIDE_INT
)1 << i
);
2581 hsa_op_reg
*value_reg
= new hsa_op_reg (mem_type
);
2583 req_align
= hsa_bitmemref_alignment (lhs
);
2584 /* Load value from memory. */
2585 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mem_type
,
2587 mem
->set_align (req_align
);
2588 hbb
->append_insn (mem
);
2590 /* AND the loaded value with prepared mask. */
2591 hsa_op_reg
*cleared_reg
= new hsa_op_reg (mem_type
);
2594 = get_integer_type_by_bytes (type_bitsize
/ BITS_PER_UNIT
, false);
2595 hsa_op_immed
*c
= new hsa_op_immed (mask
, t
);
2597 hsa_insn_basic
*clearing
2598 = new hsa_insn_basic (3, BRIG_OPCODE_AND
, mem_type
, cleared_reg
,
2600 hbb
->append_insn (clearing
);
2602 /* Shift to left a value that is going to be stored. */
2603 hsa_op_reg
*new_value_reg
= new hsa_op_reg (mem_type
);
2605 hsa_insn_basic
*basic
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, mem_type
,
2606 new_value_reg
, src
);
2607 hbb
->append_insn (basic
);
2611 hsa_op_reg
*shifted_value_reg
= new hsa_op_reg (mem_type
);
2612 c
= new hsa_op_immed (bitpos
, BRIG_TYPE_U32
);
2614 hsa_insn_basic
*basic
2615 = new hsa_insn_basic (3, BRIG_OPCODE_SHL
, mem_type
,
2616 shifted_value_reg
, new_value_reg
, c
);
2617 hbb
->append_insn (basic
);
2619 new_value_reg
= shifted_value_reg
;
2622 /* OR the prepared value with prepared chunk loaded from memory. */
2623 hsa_op_reg
*prepared_reg
= new hsa_op_reg (mem_type
);
2624 basic
= new hsa_insn_basic (3, BRIG_OPCODE_OR
, mem_type
, prepared_reg
,
2625 new_value_reg
, cleared_reg
);
2626 hbb
->append_insn (basic
);
2632 req_align
= hsa_alignment_encoding (get_object_alignment (lhs
));
2634 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
, addr
);
2635 mem
->set_align (req_align
);
2637 /* The HSAIL verifier has another constraint: if the source is an immediate
2638 then it must match the destination type. If it's a register the low bits
2639 will be used for sub-word stores. We're always allocating new operands so
2640 we can modify the above in place. */
2641 if (hsa_op_immed
*imm
= dyn_cast
<hsa_op_immed
*> (src
))
2643 if ((imm
->m_type
& BRIG_TYPE_PACK_MASK
) == BRIG_TYPE_PACK_NONE
)
2644 imm
->m_type
= mem
->m_type
;
2647 /* ...and all vector immediates apparently need to be vectors of
2649 unsigned bs
= hsa_type_bit_size (imm
->m_type
);
2650 gcc_assert (bs
== hsa_type_bit_size (mem
->m_type
));
2654 imm
->m_type
= BRIG_TYPE_U8X4
;
2657 imm
->m_type
= BRIG_TYPE_U8X8
;
2660 imm
->m_type
= BRIG_TYPE_U8X16
;
2668 hbb
->append_insn (mem
);
2671 /* Generate memory copy instructions that are going to be used
2672 for copying a HSA symbol SRC_SYMBOL (or SRC_REG) to TARGET memory,
2673 represented by pointer in a register. */
2676 gen_hsa_memory_copy (hsa_bb
*hbb
, hsa_op_address
*target
, hsa_op_address
*src
,
2679 hsa_op_address
*addr
;
2682 unsigned offset
= 0;
2696 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2698 hsa_op_reg
*tmp
= new hsa_op_reg (t
);
2699 addr
= new hsa_op_address (src
->m_symbol
, src
->m_reg
,
2700 src
->m_imm_offset
+ offset
);
2701 mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, t
, tmp
, addr
);
2702 hbb
->append_insn (mem
);
2704 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2705 target
->m_imm_offset
+ offset
);
2706 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, tmp
, addr
);
2707 hbb
->append_insn (mem
);
2713 /* Create a memset mask that is created by copying a CONSTANT byte value
2714 to an integer of BYTE_SIZE bytes. */
2716 static unsigned HOST_WIDE_INT
2717 build_memset_value (unsigned HOST_WIDE_INT constant
, unsigned byte_size
)
2722 HOST_WIDE_INT v
= constant
;
2724 for (unsigned i
= 1; i
< byte_size
; i
++)
2725 v
|= constant
<< (8 * i
);
2730 /* Generate memory set instructions that are going to be used
2731 for setting a CONSTANT byte value to TARGET memory of SIZE bytes. */
2734 gen_hsa_memory_set (hsa_bb
*hbb
, hsa_op_address
*target
,
2735 unsigned HOST_WIDE_INT constant
,
2738 hsa_op_address
*addr
;
2741 unsigned offset
= 0;
2755 addr
= new hsa_op_address (target
->m_symbol
, target
->m_reg
,
2756 target
->m_imm_offset
+ offset
);
2758 BrigType16_t t
= get_integer_type_by_bytes (s
, false);
2759 HOST_WIDE_INT c
= build_memset_value (constant
, s
);
2761 mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, t
, new hsa_op_immed (c
, t
),
2763 hbb
->append_insn (mem
);
2769 /* Generate HSAIL instructions for a single assignment
2770 of an empty constructor to an ADDR_LHS. Constructor is passed as a
2771 tree RHS and all instructions are appended to HBB. */
2774 gen_hsa_ctor_assignment (hsa_op_address
*addr_lhs
, tree rhs
, hsa_bb
*hbb
)
2776 if (vec_safe_length (CONSTRUCTOR_ELTS (rhs
)))
2778 HSA_SORRY_AT (EXPR_LOCATION (rhs
),
2779 "support for HSA does not implement load from constructor");
2783 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2784 gen_hsa_memory_set (hbb
, addr_lhs
, 0, size
);
2787 /* Generate HSA instructions for a single assignment of RHS to LHS.
2788 HBB is the basic block they will be appended to. */
2791 gen_hsa_insns_for_single_assignment (tree lhs
, tree rhs
, hsa_bb
*hbb
)
2793 if (TREE_CODE (lhs
) == SSA_NAME
)
2795 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
2796 if (hsa_seen_error ())
2799 gen_hsa_insns_for_load (dest
, rhs
, TREE_TYPE (lhs
), hbb
);
2801 else if (TREE_CODE (rhs
) == SSA_NAME
2802 || (is_gimple_min_invariant (rhs
) && TREE_CODE (rhs
) != STRING_CST
))
2804 /* Store to memory. */
2805 hsa_op_base
*src
= hsa_reg_or_immed_for_gimple_op (rhs
, hbb
);
2806 if (hsa_seen_error ())
2809 gen_hsa_insns_for_store (lhs
, src
, hbb
);
2813 hsa_op_address
*addr_lhs
= gen_hsa_addr (lhs
, hbb
);
2815 if (TREE_CODE (rhs
) == CONSTRUCTOR
)
2816 gen_hsa_ctor_assignment (addr_lhs
, rhs
, hbb
);
2819 hsa_op_address
*addr_rhs
= gen_hsa_addr (rhs
, hbb
);
2821 unsigned size
= tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs
)));
2822 gen_hsa_memory_copy (hbb
, addr_lhs
, addr_rhs
, size
);
2827 /* Prepend before INSN a load from spill symbol of SPILL_REG. Return the
2828 register into which we loaded. If this required another register to convert
2829 from a B1 type, return it in *PTMP2, otherwise store NULL into it. We
2830 assume we are out of SSA so the returned register does not have its
2834 hsa_spill_in (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2836 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2837 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2838 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2840 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, spill_sym
->m_type
,
2842 hsa_insert_insn_before (mem
, insn
);
2845 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2847 hsa_insn_basic
*cvtinsn
;
2849 reg
= new hsa_op_reg (spill_reg
->m_type
);
2851 cvtinsn
= new hsa_insn_cvt (reg
, *ptmp2
);
2852 hsa_insert_insn_before (cvtinsn
, insn
);
2857 /* Append after INSN a store to spill symbol of SPILL_REG. Return the register
2858 from which we stored. If this required another register to convert to a B1
2859 type, return it in *PTMP2, otherwise store NULL into it. We assume we are
2860 out of SSA so the returned register does not have its use updated. */
2863 hsa_spill_out (hsa_insn_basic
*insn
, hsa_op_reg
*spill_reg
, hsa_op_reg
**ptmp2
)
2865 hsa_symbol
*spill_sym
= spill_reg
->m_spill_sym
;
2866 hsa_op_reg
*reg
= new hsa_op_reg (spill_sym
->m_type
);
2867 hsa_op_address
*addr
= new hsa_op_address (spill_sym
);
2868 hsa_op_reg
*returnreg
;
2872 if (spill_reg
->m_type
== BRIG_TYPE_B1
)
2874 hsa_insn_basic
*cvtinsn
;
2875 *ptmp2
= new hsa_op_reg (spill_sym
->m_type
);
2876 reg
->m_type
= spill_reg
->m_type
;
2878 cvtinsn
= new hsa_insn_cvt (*ptmp2
, returnreg
);
2879 hsa_append_insn_after (cvtinsn
, insn
);
2884 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, spill_sym
->m_type
, reg
,
2886 hsa_append_insn_after (mem
, insn
);
2890 /* Generate a comparison instruction that will compare LHS and RHS with
2891 comparison specified by CODE and put result into register DEST. DEST has to
2892 have its type set already but must not have its definition set yet.
2893 Generated instructions will be added to HBB. */
2896 gen_hsa_cmp_insn_from_gimple (enum tree_code code
, tree lhs
, tree rhs
,
2897 hsa_op_reg
*dest
, hsa_bb
*hbb
)
2899 BrigCompareOperation8_t compare
;
2904 compare
= BRIG_COMPARE_LT
;
2907 compare
= BRIG_COMPARE_LE
;
2910 compare
= BRIG_COMPARE_GT
;
2913 compare
= BRIG_COMPARE_GE
;
2916 compare
= BRIG_COMPARE_EQ
;
2919 compare
= BRIG_COMPARE_NE
;
2921 case UNORDERED_EXPR
:
2922 compare
= BRIG_COMPARE_NAN
;
2925 compare
= BRIG_COMPARE_NUM
;
2928 compare
= BRIG_COMPARE_LTU
;
2931 compare
= BRIG_COMPARE_LEU
;
2934 compare
= BRIG_COMPARE_GTU
;
2937 compare
= BRIG_COMPARE_GEU
;
2940 compare
= BRIG_COMPARE_EQU
;
2943 compare
= BRIG_COMPARE_NEU
;
2947 HSA_SORRY_ATV (EXPR_LOCATION (lhs
),
2948 "support for HSA does not implement comparison tree "
2949 "code %s\n", get_tree_code_name (code
));
2953 /* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer)
2954 as a result of comparison. */
2956 BrigType16_t dest_type
= hsa_type_integer_p (dest
->m_type
)
2957 ? (BrigType16_t
) BRIG_TYPE_B1
: dest
->m_type
;
2959 hsa_insn_cmp
*cmp
= new hsa_insn_cmp (compare
, dest_type
);
2960 cmp
->set_op (1, hsa_reg_or_immed_for_gimple_op (lhs
, hbb
));
2961 cmp
->set_op (2, hsa_reg_or_immed_for_gimple_op (rhs
, hbb
));
2963 hbb
->append_insn (cmp
);
2964 cmp
->set_output_in_type (dest
, 0, hbb
);
2967 /* Generate an unary instruction with OPCODE and append it to a basic block
2968 HBB. The instruction uses DEST as a destination and OP1
2969 as a single operand. */
2972 gen_hsa_unary_operation (BrigOpcode opcode
, hsa_op_reg
*dest
,
2973 hsa_op_with_type
*op1
, hsa_bb
*hbb
)
2975 gcc_checking_assert (dest
);
2976 hsa_insn_basic
*insn
;
2978 if (opcode
== BRIG_OPCODE_MOV
&& hsa_needs_cvt (dest
->m_type
, op1
->m_type
))
2979 insn
= new hsa_insn_cvt (dest
, op1
);
2980 else if (opcode
== BRIG_OPCODE_FIRSTBIT
|| opcode
== BRIG_OPCODE_LASTBIT
)
2981 insn
= new hsa_insn_srctype (2, opcode
, BRIG_TYPE_U32
, op1
->m_type
, NULL
,
2985 insn
= new hsa_insn_basic (2, opcode
, dest
->m_type
, dest
, op1
);
2987 if (opcode
== BRIG_OPCODE_ABS
|| opcode
== BRIG_OPCODE_NEG
)
2989 /* ABS and NEG only exist in _s form :-/ */
2990 if (insn
->m_type
== BRIG_TYPE_U32
)
2991 insn
->m_type
= BRIG_TYPE_S32
;
2992 else if (insn
->m_type
== BRIG_TYPE_U64
)
2993 insn
->m_type
= BRIG_TYPE_S64
;
2997 hbb
->append_insn (insn
);
2999 if (opcode
== BRIG_OPCODE_FIRSTBIT
|| opcode
== BRIG_OPCODE_LASTBIT
)
3000 insn
->set_output_in_type (dest
, 0, hbb
);
3003 /* Generate a binary instruction with OPCODE and append it to a basic block
3004 HBB. The instruction uses DEST as a destination and operands OP1
3008 gen_hsa_binary_operation (int opcode
, hsa_op_reg
*dest
,
3009 hsa_op_base
*op1
, hsa_op_base
*op2
, hsa_bb
*hbb
)
3011 gcc_checking_assert (dest
);
3013 if ((opcode
== BRIG_OPCODE_SHL
|| opcode
== BRIG_OPCODE_SHR
)
3014 && is_a
<hsa_op_immed
*> (op2
))
3016 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3017 i
->set_type (BRIG_TYPE_U32
);
3019 if ((opcode
== BRIG_OPCODE_OR
3020 || opcode
== BRIG_OPCODE_XOR
3021 || opcode
== BRIG_OPCODE_AND
)
3022 && is_a
<hsa_op_immed
*> (op2
))
3024 hsa_op_immed
*i
= dyn_cast
<hsa_op_immed
*> (op2
);
3025 i
->set_type (hsa_uint_for_bitsize (hsa_type_bit_size (i
->m_type
)));
3028 hsa_insn_basic
*insn
= new hsa_insn_basic (3, opcode
, dest
->m_type
, dest
,
3030 hbb
->append_insn (insn
);
3033 /* Generate HSA instructions for a single assignment. HBB is the basic block
3034 they will be appended to. */
3037 gen_hsa_insns_for_operation_assignment (gimple
*assign
, hsa_bb
*hbb
)
3039 tree_code code
= gimple_assign_rhs_code (assign
);
3040 gimple_rhs_class rhs_class
= get_gimple_rhs_class (gimple_expr_code (assign
));
3042 tree lhs
= gimple_assign_lhs (assign
);
3043 tree rhs1
= gimple_assign_rhs1 (assign
);
3044 tree rhs2
= gimple_assign_rhs2 (assign
);
3045 tree rhs3
= gimple_assign_rhs3 (assign
);
3053 /* The opcode is changed to BRIG_OPCODE_CVT if BRIG types
3054 needs a conversion. */
3055 opcode
= BRIG_OPCODE_MOV
;
3059 case POINTER_PLUS_EXPR
:
3060 opcode
= BRIG_OPCODE_ADD
;
3063 opcode
= BRIG_OPCODE_SUB
;
3066 opcode
= BRIG_OPCODE_MUL
;
3068 case MULT_HIGHPART_EXPR
:
3069 opcode
= BRIG_OPCODE_MULHI
;
3072 case TRUNC_DIV_EXPR
:
3073 case EXACT_DIV_EXPR
:
3074 opcode
= BRIG_OPCODE_DIV
;
3077 case FLOOR_DIV_EXPR
:
3078 case ROUND_DIV_EXPR
:
3079 HSA_SORRY_AT (gimple_location (assign
),
3080 "support for HSA does not implement CEIL_DIV_EXPR, "
3081 "FLOOR_DIV_EXPR or ROUND_DIV_EXPR");
3083 case TRUNC_MOD_EXPR
:
3084 opcode
= BRIG_OPCODE_REM
;
3087 case FLOOR_MOD_EXPR
:
3088 case ROUND_MOD_EXPR
:
3089 HSA_SORRY_AT (gimple_location (assign
),
3090 "support for HSA does not implement CEIL_MOD_EXPR, "
3091 "FLOOR_MOD_EXPR or ROUND_MOD_EXPR");
3094 opcode
= BRIG_OPCODE_NEG
;
3097 opcode
= BRIG_OPCODE_MIN
;
3100 opcode
= BRIG_OPCODE_MAX
;
3103 opcode
= BRIG_OPCODE_ABS
;
3106 opcode
= BRIG_OPCODE_SHL
;
3109 opcode
= BRIG_OPCODE_SHR
;
3114 hsa_insn_basic
*insn
= NULL
;
3115 int code1
= code
== LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3116 int code2
= code
!= LROTATE_EXPR
? BRIG_OPCODE_SHL
: BRIG_OPCODE_SHR
;
3117 BrigType16_t btype
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
3120 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3121 hsa_op_reg
*op1
= new hsa_op_reg (btype
);
3122 hsa_op_reg
*op2
= new hsa_op_reg (btype
);
3123 hsa_op_with_type
*shift1
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3125 tree type
= TREE_TYPE (rhs2
);
3126 unsigned HOST_WIDE_INT bitsize
= TREE_INT_CST_LOW (TYPE_SIZE (type
));
3128 hsa_op_with_type
*shift2
= NULL
;
3129 if (TREE_CODE (rhs2
) == INTEGER_CST
)
3130 shift2
= new hsa_op_immed (bitsize
- tree_to_uhwi (rhs2
),
3132 else if (TREE_CODE (rhs2
) == SSA_NAME
)
3134 hsa_op_reg
*s
= hsa_cfun
->reg_for_gimple_ssa (rhs2
);
3135 hsa_op_reg
*d
= new hsa_op_reg (s
->m_type
);
3136 hsa_op_immed
*size_imm
= new hsa_op_immed (bitsize
, BRIG_TYPE_U32
);
3138 insn
= new hsa_insn_basic (3, BRIG_OPCODE_SUB
, d
->m_type
,
3140 hbb
->append_insn (insn
);
3147 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3148 gen_hsa_binary_operation (code1
, op1
, src
, shift1
, hbb
);
3149 gen_hsa_binary_operation (code2
, op2
, src
, shift2
, hbb
);
3150 gen_hsa_binary_operation (BRIG_OPCODE_OR
, dest
, op1
, op2
, hbb
);
3155 opcode
= BRIG_OPCODE_OR
;
3158 opcode
= BRIG_OPCODE_XOR
;
3161 opcode
= BRIG_OPCODE_AND
;
3164 opcode
= BRIG_OPCODE_NOT
;
3166 case FIX_TRUNC_EXPR
:
3168 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3169 hsa_op_with_type
*v
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3171 if (hsa_needs_cvt (dest
->m_type
, v
->m_type
))
3173 hsa_op_reg
*tmp
= new hsa_op_reg (v
->m_type
);
3175 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3176 tmp
->m_type
, tmp
, v
);
3177 hbb
->append_insn (insn
);
3179 hsa_insn_basic
*cvtinsn
= new hsa_insn_cvt (dest
, tmp
);
3180 hbb
->append_insn (cvtinsn
);
3184 hsa_insn_basic
*insn
= new hsa_insn_basic (2, BRIG_OPCODE_TRUNC
,
3185 dest
->m_type
, dest
, v
);
3186 hbb
->append_insn (insn
);
3191 opcode
= BRIG_OPCODE_TRUNC
;
3200 case UNORDERED_EXPR
:
3210 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3212 gen_hsa_cmp_insn_from_gimple (code
, rhs1
, rhs2
, dest
, hbb
);
3218 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3219 hsa_op_with_type
*ctrl
= NULL
;
3222 if (CONSTANT_CLASS_P (cond
) || TREE_CODE (cond
) == SSA_NAME
)
3223 ctrl
= hsa_reg_or_immed_for_gimple_op (cond
, hbb
);
3226 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
3228 gen_hsa_cmp_insn_from_gimple (TREE_CODE (cond
),
3229 TREE_OPERAND (cond
, 0),
3230 TREE_OPERAND (cond
, 1),
3236 hsa_op_with_type
*rhs2_reg
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3237 hsa_op_with_type
*rhs3_reg
= hsa_reg_or_immed_for_gimple_op (rhs3
, hbb
);
3239 BrigType16_t btype
= hsa_bittype_for_type (dest
->m_type
);
3240 hsa_op_reg
*tmp
= new hsa_op_reg (btype
);
3242 rhs2_reg
->m_type
= btype
;
3243 rhs3_reg
->m_type
= btype
;
3245 hsa_insn_basic
*insn
3246 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, tmp
->m_type
, tmp
, ctrl
,
3247 rhs2_reg
, rhs3_reg
);
3249 hbb
->append_insn (insn
);
3251 /* As operands of a CMOV insn must be Bx types, we have to emit
3252 a conversion insn. */
3253 hsa_insn_basic
*mov
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
,
3254 dest
->m_type
, dest
, tmp
);
3255 hbb
->append_insn (mov
);
3262 = hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3263 hsa_op_with_type
*rhs1_reg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3264 hsa_op_with_type
*rhs2_reg
= hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
);
3266 if (hsa_seen_error ())
3269 BrigType16_t src_type
= hsa_bittype_for_type (rhs1_reg
->m_type
);
3270 rhs1_reg
= rhs1_reg
->get_in_type (src_type
, hbb
);
3271 rhs2_reg
= rhs2_reg
->get_in_type (src_type
, hbb
);
3273 hsa_insn_packed
*insn
3274 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE
, dest
->m_type
, src_type
,
3275 dest
, rhs1_reg
, rhs2_reg
);
3276 hbb
->append_insn (insn
);
3281 /* Implement others as we come across them. */
3282 HSA_SORRY_ATV (gimple_location (assign
),
3283 "support for HSA does not implement operation %s",
3284 get_tree_code_name (code
));
3289 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (gimple_assign_lhs (assign
));
3291 hsa_op_with_type
*op1
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
3292 hsa_op_with_type
*op2
= rhs2
!= NULL_TREE
?
3293 hsa_reg_or_immed_for_gimple_op (rhs2
, hbb
) : NULL
;
3295 if (hsa_seen_error ())
3300 case GIMPLE_TERNARY_RHS
:
3305 case GIMPLE_BINARY_RHS
:
3306 gen_hsa_binary_operation (opcode
, dest
, op1
, op2
, hbb
);
3309 case GIMPLE_UNARY_RHS
:
3310 gen_hsa_unary_operation (opcode
, dest
, op1
, hbb
);
3317 /* Generate HSA instructions for a given gimple condition statement COND.
3318 Instructions will be appended to HBB, which also needs to be the
3319 corresponding structure to the basic_block of COND. */
3322 gen_hsa_insns_for_cond_stmt (gimple
*cond
, hsa_bb
*hbb
)
3324 hsa_op_reg
*ctrl
= new hsa_op_reg (BRIG_TYPE_B1
);
3327 gen_hsa_cmp_insn_from_gimple (gimple_cond_code (cond
),
3328 gimple_cond_lhs (cond
),
3329 gimple_cond_rhs (cond
),
3332 cbr
= new hsa_insn_br (ctrl
);
3333 hbb
->append_insn (cbr
);
3336 /* Maximum number of elements in a jump table for an HSA SBR instruction. */
3338 #define HSA_MAXIMUM_SBR_LABELS 16
3340 /* Return lowest value of a switch S that is handled in a non-default
3344 get_switch_low (gswitch
*s
)
3346 unsigned labels
= gimple_switch_num_labels (s
);
3347 gcc_checking_assert (labels
>= 1);
3349 return CASE_LOW (gimple_switch_label (s
, 1));
3352 /* Return highest value of a switch S that is handled in a non-default
3356 get_switch_high (gswitch
*s
)
3358 unsigned labels
= gimple_switch_num_labels (s
);
3360 /* Compare last label to maximum number of labels. */
3361 tree label
= gimple_switch_label (s
, labels
- 1);
3362 tree low
= CASE_LOW (label
);
3363 tree high
= CASE_HIGH (label
);
3365 return high
!= NULL_TREE
? high
: low
;
3369 get_switch_size (gswitch
*s
)
3371 return int_const_binop (MINUS_EXPR
, get_switch_high (s
), get_switch_low (s
));
3374 /* Generate HSA instructions for a given gimple switch.
3375 Instructions will be appended to HBB. */
3378 gen_hsa_insns_for_switch_stmt (gswitch
*s
, hsa_bb
*hbb
)
3380 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
3381 tree index_tree
= gimple_switch_index (s
);
3382 tree lowest
= get_switch_low (s
);
3384 hsa_op_reg
*index
= hsa_cfun
->reg_for_gimple_ssa (index_tree
);
3385 hsa_op_reg
*sub_index
= new hsa_op_reg (index
->m_type
);
3386 hbb
->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB
, sub_index
->m_type
,
3388 new hsa_op_immed (lowest
)));
3390 hsa_op_base
*tmp
= sub_index
->get_in_type (BRIG_TYPE_U64
, hbb
);
3391 sub_index
= as_a
<hsa_op_reg
*> (tmp
);
3392 unsigned labels
= gimple_switch_num_labels (s
);
3393 unsigned HOST_WIDE_INT size
= tree_to_uhwi (get_switch_size (s
));
3395 hsa_insn_sbr
*sbr
= new hsa_insn_sbr (sub_index
, size
+ 1);
3396 tree default_label
= gimple_switch_default_label (s
);
3397 basic_block default_label_bb
= label_to_block_fn (func
,
3398 CASE_LABEL (default_label
));
3400 sbr
->m_default_bb
= default_label_bb
;
3402 /* Prepare array with default label destination. */
3403 for (unsigned HOST_WIDE_INT i
= 0; i
<= size
; i
++)
3404 sbr
->m_jump_table
.safe_push (default_label_bb
);
3406 /* Iterate all labels and fill up the jump table. */
3407 for (unsigned i
= 1; i
< labels
; i
++)
3409 tree label
= gimple_switch_label (s
, i
);
3410 basic_block bb
= label_to_block_fn (func
, CASE_LABEL (label
));
3412 unsigned HOST_WIDE_INT sub_low
3413 = tree_to_uhwi (int_const_binop (MINUS_EXPR
, CASE_LOW (label
), lowest
));
3415 unsigned HOST_WIDE_INT sub_high
= sub_low
;
3416 tree high
= CASE_HIGH (label
);
3418 sub_high
= tree_to_uhwi (int_const_binop (MINUS_EXPR
, high
, lowest
));
3420 for (unsigned HOST_WIDE_INT j
= sub_low
; j
<= sub_high
; j
++)
3421 sbr
->m_jump_table
[j
] = bb
;
3424 hbb
->append_insn (sbr
);
3427 /* Verify that the function DECL can be handled by HSA. */
3430 verify_function_arguments (tree decl
)
3432 if (DECL_STATIC_CHAIN (decl
))
3434 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3435 "HSA does not support nested functions: %D", decl
);
3438 else if (!TYPE_ARG_TYPES (TREE_TYPE (decl
)))
3440 HSA_SORRY_ATV (EXPR_LOCATION (decl
),
3441 "HSA does not support functions with variadic arguments "
3442 "(or unknown return type): %D", decl
);
3447 /* Return BRIG type for FORMAL_ARG_TYPE. If the formal argument type is NULL,
3448 return ACTUAL_ARG_TYPE. */
3451 get_format_argument_type (tree formal_arg_type
, BrigType16_t actual_arg_type
)
3453 if (formal_arg_type
== NULL
)
3454 return actual_arg_type
;
3456 BrigType16_t decl_type
3457 = hsa_type_for_scalar_tree_type (formal_arg_type
, false);
3458 return mem_type_for_type (decl_type
);
3461 /* Generate HSA instructions for a direct call instruction.
3462 Instructions will be appended to HBB, which also needs to be the
3463 corresponding structure to the basic_block of STMT. */
3466 gen_hsa_insns_for_direct_call (gimple
*stmt
, hsa_bb
*hbb
)
3468 tree decl
= gimple_call_fndecl (stmt
);
3469 verify_function_arguments (decl
);
3470 if (hsa_seen_error ())
3473 hsa_insn_call
*call_insn
= new hsa_insn_call (decl
);
3474 hsa_cfun
->m_called_functions
.safe_push (call_insn
->m_called_function
);
3476 /* Argument block start. */
3477 hsa_insn_arg_block
*arg_start
3478 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3479 hbb
->append_insn (arg_start
);
3481 tree parm_type_chain
= TYPE_ARG_TYPES (gimple_call_fntype (stmt
));
3483 /* Preparation of arguments that will be passed to function. */
3484 const unsigned args
= gimple_call_num_args (stmt
);
3485 for (unsigned i
= 0; i
< args
; ++i
)
3487 tree parm
= gimple_call_arg (stmt
, (int)i
);
3488 tree parm_decl_type
= parm_type_chain
!= NULL_TREE
3489 ? TREE_VALUE (parm_type_chain
) : NULL_TREE
;
3490 hsa_op_address
*addr
;
3492 if (AGGREGATE_TYPE_P (TREE_TYPE (parm
)))
3494 addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3495 hsa_op_address
*src
= gen_hsa_addr (parm
, hbb
);
3496 gen_hsa_memory_copy (hbb
, addr
, src
,
3497 addr
->m_symbol
->total_byte_size ());
3501 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3503 if (parm_decl_type
!= NULL
&& AGGREGATE_TYPE_P (parm_decl_type
))
3505 HSA_SORRY_AT (gimple_location (stmt
),
3506 "support for HSA does not implement an aggregate "
3507 "formal argument in a function call, while actual "
3508 "argument is not an aggregate");
3512 BrigType16_t formal_arg_type
3513 = get_format_argument_type (parm_decl_type
, src
->m_type
);
3514 if (hsa_seen_error ())
3517 if (src
->m_type
!= formal_arg_type
)
3518 src
= src
->get_in_type (formal_arg_type
, hbb
);
3521 = gen_hsa_addr_for_arg (parm_decl_type
!= NULL_TREE
?
3522 parm_decl_type
: TREE_TYPE (parm
), i
);
3523 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, formal_arg_type
,
3526 hbb
->append_insn (mem
);
3529 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3530 if (parm_type_chain
)
3531 parm_type_chain
= TREE_CHAIN (parm_type_chain
);
3534 call_insn
->m_args_code_list
= new hsa_op_code_list (args
);
3535 hbb
->append_insn (call_insn
);
3537 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
3539 tree result
= gimple_call_lhs (stmt
);
3540 hsa_insn_mem
*result_insn
= NULL
;
3541 if (!VOID_TYPE_P (result_type
))
3543 hsa_op_address
*addr
= gen_hsa_addr_for_arg (result_type
, -1);
3545 /* Even if result of a function call is unused, we have to emit
3546 declaration for the result. */
3549 tree lhs_type
= TREE_TYPE (result
);
3551 if (hsa_seen_error ())
3554 if (AGGREGATE_TYPE_P (lhs_type
))
3556 hsa_op_address
*result_addr
= gen_hsa_addr (result
, hbb
);
3557 gen_hsa_memory_copy (hbb
, result_addr
, addr
,
3558 addr
->m_symbol
->total_byte_size ());
3563 = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type
,
3566 hsa_op_reg
*dst
= hsa_cfun
->reg_for_gimple_ssa (result
);
3567 result_insn
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
, dst
, addr
);
3568 hbb
->append_insn (result_insn
);
3572 call_insn
->m_output_arg
= addr
->m_symbol
;
3573 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3579 HSA_SORRY_AT (gimple_location (stmt
),
3580 "support for HSA does not implement an assignment of "
3581 "return value from a void function");
3585 call_insn
->m_result_code_list
= new hsa_op_code_list (0);
3588 /* Argument block end. */
3589 hsa_insn_arg_block
*arg_end
3590 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3591 hbb
->append_insn (arg_end
);
3594 /* Generate HSA instructions for a direct call of an internal fn.
3595 Instructions will be appended to HBB, which also needs to be the
3596 corresponding structure to the basic_block of STMT. */
3599 gen_hsa_insns_for_call_of_internal_fn (gimple
*stmt
, hsa_bb
*hbb
)
3601 tree lhs
= gimple_call_lhs (stmt
);
3605 tree lhs_type
= TREE_TYPE (lhs
);
3606 tree rhs1
= gimple_call_arg (stmt
, 0);
3607 tree rhs1_type
= TREE_TYPE (rhs1
);
3608 enum internal_fn fn
= gimple_call_internal_fn (stmt
);
3609 hsa_internal_fn
*ifn
3610 = new hsa_internal_fn (fn
, tree_to_uhwi (TYPE_SIZE (rhs1_type
)));
3611 hsa_insn_call
*call_insn
= new hsa_insn_call (ifn
);
3613 gcc_checking_assert (FLOAT_TYPE_P (rhs1_type
));
3615 if (!hsa_emitted_internal_decls
->find (call_insn
->m_called_internal_fn
))
3616 hsa_cfun
->m_called_internal_fns
.safe_push (call_insn
->m_called_internal_fn
);
3618 hsa_insn_arg_block
*arg_start
3619 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START
, call_insn
);
3620 hbb
->append_insn (arg_start
);
3622 unsigned num_args
= gimple_call_num_args (stmt
);
3624 /* Function arguments. */
3625 for (unsigned i
= 0; i
< num_args
; i
++)
3627 tree parm
= gimple_call_arg (stmt
, (int)i
);
3628 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (parm
, hbb
);
3630 hsa_op_address
*addr
= gen_hsa_addr_for_arg (TREE_TYPE (parm
), i
);
3631 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, src
->m_type
,
3634 call_insn
->m_input_args
.safe_push (addr
->m_symbol
);
3635 hbb
->append_insn (mem
);
3638 call_insn
->m_args_code_list
= new hsa_op_code_list (num_args
);
3639 hbb
->append_insn (call_insn
);
3641 /* Assign returned value. */
3642 hsa_op_address
*addr
= gen_hsa_addr_for_arg (lhs_type
, -1);
3644 call_insn
->m_output_arg
= addr
->m_symbol
;
3645 call_insn
->m_result_code_list
= new hsa_op_code_list (1);
3647 /* Argument block end. */
3648 hsa_insn_arg_block
*arg_end
3649 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END
, call_insn
);
3650 hbb
->append_insn (arg_end
);
3653 /* Generate HSA instructions for a return value instruction.
3654 Instructions will be appended to HBB, which also needs to be the
3655 corresponding structure to the basic_block of STMT. */
3658 gen_hsa_insns_for_return (greturn
*stmt
, hsa_bb
*hbb
)
3660 tree retval
= gimple_return_retval (stmt
);
3663 hsa_op_address
*addr
= new hsa_op_address (hsa_cfun
->m_output_arg
);
3665 if (AGGREGATE_TYPE_P (TREE_TYPE (retval
)))
3667 hsa_op_address
*retval_addr
= gen_hsa_addr (retval
, hbb
);
3668 gen_hsa_memory_copy (hbb
, addr
, retval_addr
,
3669 hsa_cfun
->m_output_arg
->total_byte_size ());
3673 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (retval
),
3675 BrigType16_t mtype
= mem_type_for_type (t
);
3677 /* Store of return value. */
3678 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (retval
, hbb
);
3679 src
= src
->get_in_type (mtype
, hbb
);
3680 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, mtype
, src
,
3682 hbb
->append_insn (mem
);
3686 /* HSAIL return instruction emission. */
3687 hsa_insn_basic
*ret
= new hsa_insn_basic (0, BRIG_OPCODE_RET
);
3688 hbb
->append_insn (ret
);
3691 /* Set OP_INDEX-th operand of the instruction to DEST, as the DEST
3692 can have a different type, conversion instructions are possibly
3696 hsa_insn_basic::set_output_in_type (hsa_op_reg
*dest
, unsigned op_index
,
3699 hsa_insn_basic
*insn
;
3700 gcc_checking_assert (op_output_p (op_index
));
3702 if (dest
->m_type
== m_type
)
3704 set_op (op_index
, dest
);
3708 hsa_op_reg
*tmp
= new hsa_op_reg (m_type
);
3709 set_op (op_index
, tmp
);
3711 if (hsa_needs_cvt (dest
->m_type
, m_type
))
3712 insn
= new hsa_insn_cvt (dest
, tmp
);
3714 insn
= new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
,
3715 dest
, tmp
->get_in_type (dest
->m_type
, hbb
));
3717 hbb
->append_insn (insn
);
3720 /* Generate instruction OPCODE to query a property of HSA grid along the
3721 given DIMENSION. Store result into DEST and append the instruction to
3725 query_hsa_grid (hsa_op_reg
*dest
, BrigType16_t opcode
, int dimension
,
3728 /* We're using just one-dimensional kernels, so hard-coded
3731 = new hsa_op_immed (dimension
, (BrigKind16_t
) BRIG_TYPE_U32
);
3732 hsa_insn_basic
*insn
= new hsa_insn_basic (2, opcode
, BRIG_TYPE_U32
, NULL
,
3734 hbb
->append_insn (insn
);
3735 insn
->set_output_in_type (dest
, 0, hbb
);
3738 /* Generate a special HSA-related instruction for gimple STMT.
3739 Instructions are appended to basic block HBB. */
3742 query_hsa_grid (gimple
*stmt
, BrigOpcode16_t opcode
, int dimension
,
3745 tree lhs
= gimple_call_lhs (dyn_cast
<gcall
*> (stmt
));
3746 if (lhs
== NULL_TREE
)
3749 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3751 query_hsa_grid (dest
, opcode
, dimension
, hbb
);
3754 /* Emit instructions that set hsa_num_threads according to provided VALUE.
3755 Instructions are appended to basic block HBB. */
3758 gen_set_num_threads (tree value
, hsa_bb
*hbb
)
3760 hbb
->append_insn (new hsa_insn_comment ("omp_set_num_threads"));
3761 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (value
, hbb
);
3763 src
= src
->get_in_type (hsa_num_threads
->m_type
, hbb
);
3764 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
3766 hsa_insn_basic
*basic
3767 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
, src
, addr
);
3768 hbb
->append_insn (basic
);
3771 static GTY (()) tree hsa_kernel_dispatch_type
= NULL
;
3773 /* Return byte offset of a FIELD_NAME in GOMP_hsa_kernel_dispatch which
3774 is defined in plugin-hsa.c. */
3776 static HOST_WIDE_INT
3777 get_hsa_kernel_dispatch_offset (const char *field_name
)
3779 if (hsa_kernel_dispatch_type
== NULL
)
3781 /* Collection of information needed for a dispatch of a kernel from a
3782 kernel. Keep in sync with libgomp's plugin-hsa.c. */
3784 hsa_kernel_dispatch_type
= make_node (RECORD_TYPE
);
3785 tree id_f1
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3786 get_identifier ("queue"), ptr_type_node
);
3787 DECL_CHAIN (id_f1
) = NULL_TREE
;
3788 tree id_f2
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3789 get_identifier ("omp_data_memory"),
3791 DECL_CHAIN (id_f2
) = id_f1
;
3792 tree id_f3
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3793 get_identifier ("kernarg_address"),
3795 DECL_CHAIN (id_f3
) = id_f2
;
3796 tree id_f4
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3797 get_identifier ("object"),
3799 DECL_CHAIN (id_f4
) = id_f3
;
3800 tree id_f5
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3801 get_identifier ("signal"),
3803 DECL_CHAIN (id_f5
) = id_f4
;
3804 tree id_f6
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3805 get_identifier ("private_segment_size"),
3807 DECL_CHAIN (id_f6
) = id_f5
;
3808 tree id_f7
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3809 get_identifier ("group_segment_size"),
3811 DECL_CHAIN (id_f7
) = id_f6
;
3812 tree id_f8
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3813 get_identifier ("kernel_dispatch_count"),
3815 DECL_CHAIN (id_f8
) = id_f7
;
3816 tree id_f9
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3817 get_identifier ("debug"),
3819 DECL_CHAIN (id_f9
) = id_f8
;
3820 tree id_f10
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3821 get_identifier ("omp_level"),
3823 DECL_CHAIN (id_f10
) = id_f9
;
3824 tree id_f11
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3825 get_identifier ("children_dispatches"),
3827 DECL_CHAIN (id_f11
) = id_f10
;
3828 tree id_f12
= build_decl (BUILTINS_LOCATION
, FIELD_DECL
,
3829 get_identifier ("omp_num_threads"),
3831 DECL_CHAIN (id_f12
) = id_f11
;
3834 finish_builtin_struct (hsa_kernel_dispatch_type
, "__hsa_kernel_dispatch",
3836 TYPE_ARTIFICIAL (hsa_kernel_dispatch_type
) = 1;
3839 for (tree chain
= TYPE_FIELDS (hsa_kernel_dispatch_type
);
3840 chain
!= NULL_TREE
; chain
= TREE_CHAIN (chain
))
3841 if (strcmp (field_name
, IDENTIFIER_POINTER (DECL_NAME (chain
))) == 0)
3842 return int_byte_position (chain
);
3847 /* Return an HSA register that will contain number of threads for
3848 a future dispatched kernel. Instructions are added to HBB. */
3851 gen_num_threads_for_dispatch (hsa_bb
*hbb
)
3853 /* Step 1) Assign to number of threads:
3854 MIN (HSA_DEFAULT_NUM_THREADS, hsa_num_threads). */
3855 hsa_op_reg
*threads
= new hsa_op_reg (hsa_num_threads
->m_type
);
3856 hsa_op_address
*addr
= new hsa_op_address (hsa_num_threads
);
3858 hbb
->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD
, threads
->m_type
,
3861 hsa_op_immed
*limit
= new hsa_op_immed (HSA_DEFAULT_NUM_THREADS
,
3863 hsa_op_reg
*r
= new hsa_op_reg (BRIG_TYPE_B1
);
3865 = new hsa_insn_cmp (BRIG_COMPARE_LT
, r
->m_type
, r
, threads
, limit
);
3866 hbb
->append_insn (cmp
);
3868 BrigType16_t btype
= hsa_bittype_for_type (threads
->m_type
);
3869 hsa_op_reg
*tmp
= new hsa_op_reg (threads
->m_type
);
3871 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp
, r
,
3874 /* Step 2) If the number is equal to zero,
3875 return shadow->omp_num_threads. */
3876 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
3878 hsa_op_reg
*shadow_thread_count
= new hsa_op_reg (BRIG_TYPE_U32
);
3880 = new hsa_op_address (shadow_reg_ptr
,
3881 get_hsa_kernel_dispatch_offset ("omp_num_threads"));
3882 hsa_insn_basic
*basic
3883 = new hsa_insn_mem (BRIG_OPCODE_LD
, shadow_thread_count
->m_type
,
3884 shadow_thread_count
, addr
);
3885 hbb
->append_insn (basic
);
3887 hsa_op_reg
*tmp2
= new hsa_op_reg (threads
->m_type
);
3888 r
= new hsa_op_reg (BRIG_TYPE_B1
);
3889 hsa_op_immed
*imm
= new hsa_op_immed (0, shadow_thread_count
->m_type
);
3890 hbb
->append_insn (new hsa_insn_cmp (BRIG_COMPARE_EQ
, r
->m_type
, r
, tmp
, imm
));
3891 hbb
->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, btype
, tmp2
, r
,
3892 shadow_thread_count
, tmp
));
3894 hsa_op_base
*dest
= tmp2
->get_in_type (BRIG_TYPE_U16
, hbb
);
3896 return as_a
<hsa_op_reg
*> (dest
);
3900 /* Emit instructions that assign number of teams to lhs of gimple STMT.
3901 Instructions are appended to basic block HBB. */
3904 gen_get_num_teams (gimple
*stmt
, hsa_bb
*hbb
)
3906 if (gimple_call_lhs (stmt
) == NULL_TREE
)
3909 hbb
->append_insn (new hsa_insn_comment ("omp_get_num_teams"));
3911 tree lhs
= gimple_call_lhs (stmt
);
3912 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3913 hsa_op_immed
*one
= new hsa_op_immed (1, dest
->m_type
);
3915 hsa_insn_basic
*basic
3916 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, dest
, one
);
3918 hbb
->append_insn (basic
);
3921 /* Emit instructions that assign a team number to lhs of gimple STMT.
3922 Instructions are appended to basic block HBB. */
3925 gen_get_team_num (gimple
*stmt
, hsa_bb
*hbb
)
3927 if (gimple_call_lhs (stmt
) == NULL_TREE
)
3930 hbb
->append_insn (new hsa_insn_comment ("omp_get_team_num"));
3932 tree lhs
= gimple_call_lhs (stmt
);
3933 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3934 hsa_op_immed
*zero
= new hsa_op_immed (0, dest
->m_type
);
3936 hsa_insn_basic
*basic
3937 = new hsa_insn_basic (2, BRIG_OPCODE_MOV
, dest
->m_type
, dest
, zero
);
3939 hbb
->append_insn (basic
);
3942 /* Emit instructions that get levels-var ICV to lhs of gimple STMT.
3943 Instructions are appended to basic block HBB. */
3946 gen_get_level (gimple
*stmt
, hsa_bb
*hbb
)
3948 if (gimple_call_lhs (stmt
) == NULL_TREE
)
3951 hbb
->append_insn (new hsa_insn_comment ("omp_get_level"));
3953 tree lhs
= gimple_call_lhs (stmt
);
3954 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3956 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
3957 if (shadow_reg_ptr
== NULL
)
3959 HSA_SORRY_AT (gimple_location (stmt
),
3960 "support for HSA does not implement omp_get_level called "
3961 "from a function not being inlined within a kernel");
3965 hsa_op_address
*addr
3966 = new hsa_op_address (shadow_reg_ptr
,
3967 get_hsa_kernel_dispatch_offset ("omp_level"));
3969 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, BRIG_TYPE_U64
,
3970 (hsa_op_base
*) NULL
, addr
);
3971 hbb
->append_insn (mem
);
3972 mem
->set_output_in_type (dest
, 0, hbb
);
3975 /* Emit instruction that implement omp_get_max_threads of gimple STMT. */
3978 gen_get_max_threads (gimple
*stmt
, hsa_bb
*hbb
)
3980 tree lhs
= gimple_call_lhs (stmt
);
3984 hbb
->append_insn (new hsa_insn_comment ("omp_get_max_threads"));
3986 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
3987 hsa_op_with_type
*num_theads_reg
= gen_num_threads_for_dispatch (hbb
)
3988 ->get_in_type (dest
->m_type
, hbb
);
3989 hsa_build_append_simple_mov (dest
, num_theads_reg
, hbb
);
3992 /* Emit instructions that implement alloca builtin gimple STMT.
3993 Instructions are appended to basic block HBB. */
3996 gen_hsa_alloca (gcall
*call
, hsa_bb
*hbb
)
3998 tree lhs
= gimple_call_lhs (call
);
3999 if (lhs
== NULL_TREE
)
4002 built_in_function fn
= DECL_FUNCTION_CODE (gimple_call_fndecl (call
));
4004 gcc_checking_assert (fn
== BUILT_IN_ALLOCA
4005 || fn
== BUILT_IN_ALLOCA_WITH_ALIGN
);
4007 unsigned bit_alignment
= 0;
4009 if (fn
== BUILT_IN_ALLOCA_WITH_ALIGN
)
4011 tree alignment_tree
= gimple_call_arg (call
, 1);
4012 if (TREE_CODE (alignment_tree
) != INTEGER_CST
)
4014 HSA_SORRY_ATV (gimple_location (call
),
4015 "support for HSA does not implement "
4016 "__builtin_alloca_with_align with a non-constant "
4017 "alignment: %E", alignment_tree
);
4020 bit_alignment
= tree_to_uhwi (alignment_tree
);
4023 tree rhs1
= gimple_call_arg (call
, 0);
4024 hsa_op_with_type
*size
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
)
4025 ->get_in_type (BRIG_TYPE_U32
, hbb
);
4026 hsa_op_with_type
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4029 = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE
));
4030 hsa_insn_alloca
*a
= new hsa_insn_alloca (tmp
, size
, bit_alignment
);
4031 hbb
->append_insn (a
);
4034 = new hsa_insn_seg (BRIG_OPCODE_STOF
,
4035 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
),
4036 tmp
->m_type
, BRIG_SEGMENT_PRIVATE
, dest
, tmp
);
4037 hbb
->append_insn (seg
);
4040 /* Emit instructions that implement clrsb builtin STMT:
4041 Returns the number of leading redundant sign bits in x, i.e. the number
4042 of bits following the most significant bit that are identical to it.
4043 There are no special cases for 0 or other values.
4044 Instructions are appended to basic block HBB. */
4047 gen_hsa_clrsb (gcall
*call
, hsa_bb
*hbb
)
4049 tree lhs
= gimple_call_lhs (call
);
4050 if (lhs
== NULL_TREE
)
4053 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4054 tree rhs1
= gimple_call_arg (call
, 0);
4055 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4056 BrigType16_t bittype
= hsa_bittype_for_type (arg
->m_type
);
4057 unsigned bitsize
= tree_to_uhwi (TYPE_SIZE (TREE_TYPE (rhs1
)));
4059 /* FIRSTBIT instruction is defined just for 32 and 64-bits wide integers. */
4060 gcc_checking_assert (bitsize
== 32 || bitsize
== 64);
4062 /* Set true to MOST_SIG if the most significant bit is set to one. */
4063 hsa_op_immed
*c
= new hsa_op_immed (1ul << (bitsize
- 1),
4064 hsa_uint_for_bitsize (bitsize
));
4066 hsa_op_reg
*and_reg
= new hsa_op_reg (bittype
);
4067 gen_hsa_binary_operation (BRIG_OPCODE_AND
, and_reg
, arg
, c
, hbb
);
4069 hsa_op_reg
*most_sign
= new hsa_op_reg (BRIG_TYPE_B1
);
4071 = new hsa_insn_cmp (BRIG_COMPARE_EQ
, most_sign
->m_type
, most_sign
,
4073 hbb
->append_insn (cmp
);
4075 /* If the most significant bit is one, negate the input. Otherwise
4076 shift the input value to left by one bit. */
4077 hsa_op_reg
*arg_neg
= new hsa_op_reg (arg
->m_type
);
4078 gen_hsa_unary_operation (BRIG_OPCODE_NEG
, arg_neg
, arg
, hbb
);
4080 hsa_op_reg
*shifted_arg
= new hsa_op_reg (arg
->m_type
);
4081 gen_hsa_binary_operation (BRIG_OPCODE_SHL
, shifted_arg
, arg
,
4082 new hsa_op_immed (1, BRIG_TYPE_U64
), hbb
);
4084 /* Assign the value that can be used for FIRSTBIT instruction according
4085 to the most significant bit. */
4086 hsa_op_reg
*tmp
= new hsa_op_reg (bittype
);
4087 hsa_insn_basic
*cmov
4088 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, bittype
, tmp
, most_sign
,
4089 arg_neg
, shifted_arg
);
4090 hbb
->append_insn (cmov
);
4092 hsa_op_reg
*leading_bits
= new hsa_op_reg (BRIG_TYPE_S32
);
4093 gen_hsa_unary_operation (BRIG_OPCODE_FIRSTBIT
, leading_bits
,
4094 tmp
->get_in_type (hsa_uint_for_bitsize (bitsize
),
4097 /* Set flag if the input value is equal to zero. */
4098 hsa_op_reg
*is_zero
= new hsa_op_reg (BRIG_TYPE_B1
);
4099 cmp
= new hsa_insn_cmp (BRIG_COMPARE_EQ
, is_zero
->m_type
, is_zero
, arg
,
4100 new hsa_op_immed (0, arg
->m_type
));
4101 hbb
->append_insn (cmp
);
4103 /* Return the number of leading bits,
4104 or (bitsize - 1) if the input value is zero. */
4105 cmov
= new hsa_insn_basic (4, BRIG_OPCODE_CMOV
, BRIG_TYPE_B32
, NULL
, is_zero
,
4106 new hsa_op_immed (bitsize
- 1, BRIG_TYPE_U32
),
4107 leading_bits
->get_in_type (BRIG_TYPE_B32
, hbb
));
4108 hbb
->append_insn (cmov
);
4109 cmov
->set_output_in_type (dest
, 0, hbb
);
4112 /* Emit instructions that implement ffs builtin STMT:
4113 Returns one plus the index of the least significant 1-bit of x,
4114 or if x is zero, returns zero.
4115 Instructions are appended to basic block HBB. */
4118 gen_hsa_ffs (gcall
*call
, hsa_bb
*hbb
)
4120 tree lhs
= gimple_call_lhs (call
);
4121 if (lhs
== NULL_TREE
)
4124 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4126 tree rhs1
= gimple_call_arg (call
, 0);
4127 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4129 hsa_op_reg
*tmp
= new hsa_op_reg (BRIG_TYPE_U32
);
4130 hsa_insn_srctype
*insn
= new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT
,
4131 tmp
->m_type
, arg
->m_type
,
4133 hbb
->append_insn (insn
);
4135 hsa_insn_basic
*addition
4136 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
, NULL
, tmp
,
4137 new hsa_op_immed (1, tmp
->m_type
));
4138 hbb
->append_insn (addition
);
4139 addition
->set_output_in_type (dest
, 0, hbb
);
4143 gen_hsa_popcount_to_dest (hsa_op_reg
*dest
, hsa_op_with_type
*arg
, hsa_bb
*hbb
)
4145 gcc_checking_assert (hsa_type_integer_p (arg
->m_type
));
4147 if (hsa_type_bit_size (arg
->m_type
) < 32)
4148 arg
= arg
->get_in_type (BRIG_TYPE_B32
, hbb
);
4150 if (!hsa_btype_p (arg
->m_type
))
4151 arg
= arg
->get_in_type (hsa_bittype_for_type (arg
->m_type
), hbb
);
4153 hsa_insn_srctype
*popcount
4154 = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT
, BRIG_TYPE_U32
,
4155 arg
->m_type
, NULL
, arg
);
4156 hbb
->append_insn (popcount
);
4157 popcount
->set_output_in_type (dest
, 0, hbb
);
4160 /* Emit instructions that implement parity builtin STMT:
4161 Returns the parity of x, i.e. the number of 1-bits in x modulo 2.
4162 Instructions are appended to basic block HBB. */
4165 gen_hsa_parity (gcall
*call
, hsa_bb
*hbb
)
4167 tree lhs
= gimple_call_lhs (call
);
4168 if (lhs
== NULL_TREE
)
4171 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4172 tree rhs1
= gimple_call_arg (call
, 0);
4173 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4175 hsa_op_reg
*popcount
= new hsa_op_reg (BRIG_TYPE_U32
);
4176 gen_hsa_popcount_to_dest (popcount
, arg
, hbb
);
4178 hsa_insn_basic
*insn
4179 = new hsa_insn_basic (3, BRIG_OPCODE_REM
, popcount
->m_type
, NULL
, popcount
,
4180 new hsa_op_immed (2, popcount
->m_type
));
4181 hbb
->append_insn (insn
);
4182 insn
->set_output_in_type (dest
, 0, hbb
);
4185 /* Emit instructions that implement popcount builtin STMT.
4186 Instructions are appended to basic block HBB. */
4189 gen_hsa_popcount (gcall
*call
, hsa_bb
*hbb
)
4191 tree lhs
= gimple_call_lhs (call
);
4192 if (lhs
== NULL_TREE
)
4195 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4196 tree rhs1
= gimple_call_arg (call
, 0);
4197 hsa_op_with_type
*arg
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4199 gen_hsa_popcount_to_dest (dest
, arg
, hbb
);
4202 /* Set VALUE to a shadow kernel debug argument and append a new instruction
4203 to HBB basic block. */
4206 set_debug_value (hsa_bb
*hbb
, hsa_op_with_type
*value
)
4208 hsa_op_reg
*shadow_reg_ptr
= hsa_cfun
->get_shadow_reg ();
4209 if (shadow_reg_ptr
== NULL
)
4212 hsa_op_address
*addr
4213 = new hsa_op_address (shadow_reg_ptr
,
4214 get_hsa_kernel_dispatch_offset ("debug"));
4215 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_ST
, BRIG_TYPE_U64
, value
,
4217 hbb
->append_insn (mem
);
4221 omp_simple_builtin::generate (gimple
*stmt
, hsa_bb
*hbb
)
4225 if (m_warning_message
)
4226 HSA_SORRY_AT (gimple_location (stmt
), m_warning_message
)
4228 HSA_SORRY_ATV (gimple_location (stmt
),
4229 "Support for HSA does not implement calls to %s\n",
4232 else if (m_warning_message
!= NULL
)
4233 warning_at (gimple_location (stmt
), OPT_Whsa
, m_warning_message
);
4235 if (m_return_value
!= NULL
)
4237 tree lhs
= gimple_call_lhs (stmt
);
4241 hbb
->append_insn (new hsa_insn_comment (m_name
));
4243 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4244 hsa_op_with_type
*op
= m_return_value
->get_in_type (dest
->m_type
, hbb
);
4245 hsa_build_append_simple_mov (dest
, op
, hbb
);
4249 /* If STMT is a call of a known library function, generate code to perform
4250 it and return true. */
4253 gen_hsa_insns_for_known_library_call (gimple
*stmt
, hsa_bb
*hbb
)
4255 bool handled
= false;
4256 const char *name
= hsa_get_declaration_name (gimple_call_fndecl (stmt
));
4259 size_t len
= strlen (name
);
4260 if (len
> 0 && name
[len
- 1] == '_')
4262 copy
= XNEWVEC (char, len
+ 1);
4263 strcpy (copy
, name
);
4264 copy
[len
- 1] = '\0';
4268 /* Handle omp_* routines. */
4269 if (strstr (name
, "omp_") == name
)
4271 hsa_init_simple_builtins ();
4272 omp_simple_builtin
*builtin
= omp_simple_builtins
->get (name
);
4275 builtin
->generate (stmt
, hbb
);
4280 if (strcmp (name
, "omp_set_num_threads") == 0)
4281 gen_set_num_threads (gimple_call_arg (stmt
, 0), hbb
);
4282 else if (strcmp (name
, "omp_get_thread_num") == 0)
4284 hbb
->append_insn (new hsa_insn_comment (name
));
4285 query_hsa_grid (stmt
, BRIG_OPCODE_WORKITEMABSID
, 0, hbb
);
4287 else if (strcmp (name
, "omp_get_num_threads") == 0)
4289 hbb
->append_insn (new hsa_insn_comment (name
));
4290 query_hsa_grid (stmt
, BRIG_OPCODE_GRIDSIZE
, 0, hbb
);
4292 else if (strcmp (name
, "omp_get_num_teams") == 0)
4293 gen_get_num_teams (stmt
, hbb
);
4294 else if (strcmp (name
, "omp_get_team_num") == 0)
4295 gen_get_team_num (stmt
, hbb
);
4296 else if (strcmp (name
, "omp_get_level") == 0)
4297 gen_get_level (stmt
, hbb
);
4298 else if (strcmp (name
, "omp_get_active_level") == 0)
4299 gen_get_level (stmt
, hbb
);
4300 else if (strcmp (name
, "omp_in_parallel") == 0)
4301 gen_get_level (stmt
, hbb
);
4302 else if (strcmp (name
, "omp_get_max_threads") == 0)
4303 gen_get_max_threads (stmt
, hbb
);
4315 if (strcmp (name
, "__hsa_set_debug_value") == 0)
4318 if (hsa_cfun
->has_shadow_reg_p ())
4320 tree rhs1
= gimple_call_arg (stmt
, 0);
4321 hsa_op_with_type
*src
= hsa_reg_or_immed_for_gimple_op (rhs1
, hbb
);
4323 src
= src
->get_in_type (BRIG_TYPE_U64
, hbb
);
4324 set_debug_value (hbb
, src
);
4333 /* Helper functions to create a single unary HSA operations out of calls to
4334 builtins. OPCODE is the HSA operation to be generated. STMT is a gimple
4335 call to a builtin. HBB is the HSA BB to which the instruction should be
4336 added. Note that nothing will be created if STMT does not have a LHS. */
4339 gen_hsa_unaryop_for_builtin (BrigOpcode opcode
, gimple
*stmt
, hsa_bb
*hbb
)
4341 tree lhs
= gimple_call_lhs (stmt
);
4344 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4345 hsa_op_with_type
*op
4346 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 0), hbb
);
4347 gen_hsa_unary_operation (opcode
, dest
, op
, hbb
);
4350 /* Helper functions to create a call to standard library if LHS of the
4351 STMT is used. HBB is the HSA BB to which the instruction should be
4355 gen_hsa_unaryop_builtin_call (gimple
*stmt
, hsa_bb
*hbb
)
4357 tree lhs
= gimple_call_lhs (stmt
);
4361 if (gimple_call_internal_p (stmt
))
4362 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
4364 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4367 /* Helper functions to create a single unary HSA operations out of calls to
4368 builtins (if unsafe math optimizations are enable). Otherwise, create
4369 a call to standard library function.
4370 OPCODE is the HSA operation to be generated. STMT is a gimple
4371 call to a builtin. HBB is the HSA BB to which the instruction should be
4372 added. Note that nothing will be created if STMT does not have a LHS. */
4375 gen_hsa_unaryop_or_call_for_builtin (BrigOpcode opcode
, gimple
*stmt
,
4378 if (flag_unsafe_math_optimizations
)
4379 gen_hsa_unaryop_for_builtin (opcode
, stmt
, hbb
);
4381 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4384 /* Generate HSA address corresponding to a value VAL (as opposed to a memory
4385 reference tree), for example an SSA_NAME or an ADDR_EXPR. HBB is the HSA BB
4386 to which the instruction should be added. */
4388 static hsa_op_address
*
4389 get_address_from_value (tree val
, hsa_bb
*hbb
)
4391 switch (TREE_CODE (val
))
4395 BrigType16_t addrtype
= hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT
);
4397 = hsa_cfun
->reg_for_gimple_ssa (val
)->get_in_type (addrtype
, hbb
);
4398 return new hsa_op_address (NULL
, as_a
<hsa_op_reg
*> (reg
), 0);
4401 return gen_hsa_addr (TREE_OPERAND (val
, 0), hbb
);
4404 if (tree_fits_shwi_p (val
))
4405 return new hsa_op_address (NULL
, NULL
, tree_to_shwi (val
));
4406 /* Otherwise fall-through */
4409 HSA_SORRY_ATV (EXPR_LOCATION (val
),
4410 "support for HSA does not implement memory access to %E",
4412 return new hsa_op_address (NULL
, NULL
, 0);
4416 /* Return string for MEMMODEL. */
4419 get_memory_order_name (unsigned memmodel
)
4421 switch (memmodel
& MEMMODEL_BASE_MASK
)
4423 case MEMMODEL_RELAXED
:
4425 case MEMMODEL_CONSUME
:
4427 case MEMMODEL_ACQUIRE
:
4429 case MEMMODEL_RELEASE
:
4431 case MEMMODEL_ACQ_REL
:
4433 case MEMMODEL_SEQ_CST
:
4440 /* Return memory order according to predefined __atomic memory model
4441 constants. LOCATION is provided to locate the problematic statement. */
4443 static BrigMemoryOrder
4444 get_memory_order (unsigned memmodel
, location_t location
)
4446 switch (memmodel
& MEMMODEL_BASE_MASK
)
4448 case MEMMODEL_RELAXED
:
4449 return BRIG_MEMORY_ORDER_RELAXED
;
4450 case MEMMODEL_CONSUME
:
4451 /* HSA does not have an equivalent, but we can use the slightly stronger
4453 case MEMMODEL_ACQUIRE
:
4454 return BRIG_MEMORY_ORDER_SC_ACQUIRE
;
4455 case MEMMODEL_RELEASE
:
4456 return BRIG_MEMORY_ORDER_SC_RELEASE
;
4457 case MEMMODEL_ACQ_REL
:
4458 case MEMMODEL_SEQ_CST
:
4459 /* Callers implementing a simple load or store need to remove the release
4460 or acquire part respectively. */
4461 return BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
;
4464 const char *mmname
= get_memory_order_name (memmodel
);
4465 HSA_SORRY_ATV (location
,
4466 "support for HSA does not implement the specified "
4467 " memory model%s %s",
4468 mmname
? ": " : "", mmname
? mmname
: "");
4469 return BRIG_MEMORY_ORDER_NONE
;
4474 /* Helper function to create an HSA atomic binary operation instruction out of
4475 calls to atomic builtins. RET_ORIG is true if the built-in is the variant
4476 that return s the value before applying operation, and false if it should
4477 return the value after applying the operation (if it returns value at all).
4478 ACODE is the atomic operation code, STMT is a gimple call to a builtin. HBB
4479 is the HSA BB to which the instruction should be added. */
4482 gen_hsa_ternary_atomic_for_builtin (bool ret_orig
,
4483 enum BrigAtomicOperation acode
,
4487 tree lhs
= gimple_call_lhs (stmt
);
4489 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
4490 BrigType16_t hsa_type
= hsa_type_for_scalar_tree_type (type
, false);
4491 BrigType16_t mtype
= mem_type_for_type (hsa_type
);
4492 tree model
= gimple_call_arg (stmt
, 2);
4494 if (!tree_fits_uhwi_p (model
))
4496 HSA_SORRY_ATV (gimple_location (stmt
),
4497 "support for HSA does not implement memory model %E",
4502 unsigned HOST_WIDE_INT mmodel
= tree_to_uhwi (model
);
4504 BrigMemoryOrder memorder
= get_memory_order (mmodel
, gimple_location (stmt
));
4506 /* Certain atomic insns must have Bx memory types. */
4509 case BRIG_ATOMIC_LD
:
4510 case BRIG_ATOMIC_ST
:
4511 case BRIG_ATOMIC_AND
:
4512 case BRIG_ATOMIC_OR
:
4513 case BRIG_ATOMIC_XOR
:
4514 case BRIG_ATOMIC_EXCH
:
4515 mtype
= hsa_bittype_for_type (mtype
);
4526 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4528 dest
= new hsa_op_reg (hsa_type
);
4529 opcode
= BRIG_OPCODE_ATOMIC
;
4535 opcode
= BRIG_OPCODE_ATOMICNORET
;
4539 if (acode
== BRIG_ATOMIC_ST
)
4541 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
4542 memorder
= BRIG_MEMORY_ORDER_SC_RELEASE
;
4544 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
4545 && memorder
!= BRIG_MEMORY_ORDER_SC_RELEASE
4546 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
4548 HSA_SORRY_ATV (gimple_location (stmt
),
4549 "support for HSA does not implement memory model for "
4550 "ATOMIC_ST: %s", get_memory_order_name (mmodel
));
4555 hsa_insn_atomic
*atominsn
= new hsa_insn_atomic (nops
, opcode
, acode
, mtype
,
4558 hsa_op_address
*addr
;
4559 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
4560 /* TODO: Warn if addr has private segment, because the finalizer will not
4561 accept that (and it does not make much sense). */
4562 hsa_op_base
*op
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1),
4567 atominsn
->set_op (0, dest
);
4568 atominsn
->set_op (1, addr
);
4569 atominsn
->set_op (2, op
);
4573 atominsn
->set_op (0, addr
);
4574 atominsn
->set_op (1, op
);
4577 hbb
->append_insn (atominsn
);
4579 /* HSA does not natively support the variants that return the modified value,
4580 so re-do the operation again non-atomically if that is what was
4582 if (lhs
&& !ret_orig
)
4587 case BRIG_ATOMIC_ADD
:
4588 arith
= BRIG_OPCODE_ADD
;
4590 case BRIG_ATOMIC_AND
:
4591 arith
= BRIG_OPCODE_AND
;
4593 case BRIG_ATOMIC_OR
:
4594 arith
= BRIG_OPCODE_OR
;
4596 case BRIG_ATOMIC_SUB
:
4597 arith
= BRIG_OPCODE_SUB
;
4599 case BRIG_ATOMIC_XOR
:
4600 arith
= BRIG_OPCODE_XOR
;
4605 hsa_op_reg
*real_dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4606 gen_hsa_binary_operation (arith
, real_dest
, dest
, op
, hbb
);
4610 /* Generate HSA instructions for an internal fn.
4611 Instructions will be appended to HBB, which also needs to be the
4612 corresponding structure to the basic_block of STMT. */
4615 gen_hsa_insn_for_internal_fn_call (gcall
*stmt
, hsa_bb
*hbb
)
4617 gcc_checking_assert (gimple_call_internal_fn (stmt
));
4618 internal_fn fn
= gimple_call_internal_fn (stmt
);
4620 bool is_float_type_p
= false;
4621 if (gimple_call_lhs (stmt
) != NULL
4622 && TREE_TYPE (gimple_call_lhs (stmt
)) == float_type_node
)
4623 is_float_type_p
= true;
4628 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
4632 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
4636 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
4640 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
4644 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
4649 if (is_float_type_p
)
4650 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
4652 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4658 if (is_float_type_p
)
4659 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
4661 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4668 if (is_float_type_p
)
4669 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
4671 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4678 if (is_float_type_p
)
4679 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
4681 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4686 gen_hsa_clrsb (stmt
, hbb
);
4690 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
4694 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
4698 gen_hsa_ffs (stmt
, hbb
);
4702 gen_hsa_parity (stmt
, hbb
);
4706 gen_hsa_popcount (stmt
, hbb
);
4719 case IFN_SIGNIFICAND
:
4731 gen_hsa_insns_for_call_of_internal_fn (stmt
, hbb
);
4734 HSA_SORRY_ATV (gimple_location (stmt
),
4735 "support for HSA does not implement internal function: %s",
4736 internal_fn_name (fn
));
4741 #define HSA_MEMORY_BUILTINS_LIMIT 128
4743 /* Generate HSA instructions for the given call statement STMT. Instructions
4744 will be appended to HBB. */
4747 gen_hsa_insns_for_call (gimple
*stmt
, hsa_bb
*hbb
)
4749 gcall
*call
= as_a
<gcall
*> (stmt
);
4750 tree lhs
= gimple_call_lhs (stmt
);
4753 if (gimple_call_internal_p (stmt
))
4755 gen_hsa_insn_for_internal_fn_call (call
, hbb
);
4759 if (!gimple_call_builtin_p (stmt
, BUILT_IN_NORMAL
))
4761 tree function_decl
= gimple_call_fndecl (stmt
);
4762 if (function_decl
== NULL_TREE
)
4764 HSA_SORRY_AT (gimple_location (stmt
),
4765 "support for HSA does not implement indirect calls");
4769 if (hsa_callable_function_p (function_decl
))
4770 gen_hsa_insns_for_direct_call (stmt
, hbb
);
4771 else if (!gen_hsa_insns_for_known_library_call (stmt
, hbb
))
4772 HSA_SORRY_AT (gimple_location (stmt
),
4773 "HSA supports only calls of functions marked with pragma "
4774 "omp declare target");
4778 tree fndecl
= gimple_call_fndecl (stmt
);
4779 enum built_in_function builtin
= DECL_FUNCTION_CODE (fndecl
);
4783 case BUILT_IN_FABSF
:
4784 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_ABS
, stmt
, hbb
);
4788 case BUILT_IN_CEILF
:
4789 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_CEIL
, stmt
, hbb
);
4792 case BUILT_IN_FLOOR
:
4793 case BUILT_IN_FLOORF
:
4794 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FLOOR
, stmt
, hbb
);
4798 case BUILT_IN_RINTF
:
4799 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_RINT
, stmt
, hbb
);
4803 case BUILT_IN_SQRTF
:
4804 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_SQRT
, stmt
, hbb
);
4807 case BUILT_IN_TRUNC
:
4808 case BUILT_IN_TRUNCF
:
4809 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_TRUNC
, stmt
, hbb
);
4816 /* HSAIL does not provide an instruction for double argument type. */
4817 gen_hsa_unaryop_builtin_call (stmt
, hbb
);
4821 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NCOS
, stmt
, hbb
);
4824 case BUILT_IN_EXP2F
:
4825 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NEXP2
, stmt
, hbb
);
4828 case BUILT_IN_LOG2F
:
4829 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NLOG2
, stmt
, hbb
);
4833 gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN
, stmt
, hbb
);
4836 case BUILT_IN_CLRSB
:
4837 case BUILT_IN_CLRSBL
:
4838 case BUILT_IN_CLRSBLL
:
4839 gen_hsa_clrsb (call
, hbb
);
4844 case BUILT_IN_CLZLL
:
4845 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_FIRSTBIT
, stmt
, hbb
);
4850 case BUILT_IN_CTZLL
:
4851 gen_hsa_unaryop_for_builtin (BRIG_OPCODE_LASTBIT
, stmt
, hbb
);
4856 case BUILT_IN_FFSLL
:
4857 gen_hsa_ffs (call
, hbb
);
4860 case BUILT_IN_PARITY
:
4861 case BUILT_IN_PARITYL
:
4862 case BUILT_IN_PARITYLL
:
4863 gen_hsa_parity (call
, hbb
);
4866 case BUILT_IN_POPCOUNT
:
4867 case BUILT_IN_POPCOUNTL
:
4868 case BUILT_IN_POPCOUNTLL
:
4869 gen_hsa_popcount (call
, hbb
);
4872 case BUILT_IN_ATOMIC_LOAD_1
:
4873 case BUILT_IN_ATOMIC_LOAD_2
:
4874 case BUILT_IN_ATOMIC_LOAD_4
:
4875 case BUILT_IN_ATOMIC_LOAD_8
:
4876 case BUILT_IN_ATOMIC_LOAD_16
:
4879 hsa_op_address
*addr
;
4880 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
4881 tree model
= gimple_call_arg (stmt
, 1);
4882 if (!tree_fits_uhwi_p (model
))
4884 HSA_SORRY_ATV (gimple_location (stmt
),
4885 "support for HSA does not implement "
4891 unsigned HOST_WIDE_INT mmodel
= tree_to_uhwi (model
);
4892 BrigMemoryOrder memorder
= get_memory_order (mmodel
,
4893 gimple_location (stmt
));
4895 if (memorder
== BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
)
4896 memorder
= BRIG_MEMORY_ORDER_SC_ACQUIRE
;
4898 if (memorder
!= BRIG_MEMORY_ORDER_RELAXED
4899 && memorder
!= BRIG_MEMORY_ORDER_SC_ACQUIRE
4900 && memorder
!= BRIG_MEMORY_ORDER_NONE
)
4902 HSA_SORRY_ATV (gimple_location (stmt
),
4903 "support for HSA does not implement "
4904 "memory model for ATOMIC_LD: %s",
4905 get_memory_order_name (mmodel
));
4911 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (lhs
),
4913 mtype
= mem_type_for_type (t
);
4914 mtype
= hsa_bittype_for_type (mtype
);
4915 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
4919 mtype
= BRIG_TYPE_B64
;
4920 dest
= new hsa_op_reg (mtype
);
4923 hsa_insn_atomic
*atominsn
4924 = new hsa_insn_atomic (2, BRIG_OPCODE_ATOMIC
, BRIG_ATOMIC_LD
, mtype
,
4925 memorder
, dest
, addr
);
4927 hbb
->append_insn (atominsn
);
4931 case BUILT_IN_ATOMIC_EXCHANGE_1
:
4932 case BUILT_IN_ATOMIC_EXCHANGE_2
:
4933 case BUILT_IN_ATOMIC_EXCHANGE_4
:
4934 case BUILT_IN_ATOMIC_EXCHANGE_8
:
4935 case BUILT_IN_ATOMIC_EXCHANGE_16
:
4936 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_EXCH
, stmt
, hbb
);
4939 case BUILT_IN_ATOMIC_FETCH_ADD_1
:
4940 case BUILT_IN_ATOMIC_FETCH_ADD_2
:
4941 case BUILT_IN_ATOMIC_FETCH_ADD_4
:
4942 case BUILT_IN_ATOMIC_FETCH_ADD_8
:
4943 case BUILT_IN_ATOMIC_FETCH_ADD_16
:
4944 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_ADD
, stmt
, hbb
);
4947 case BUILT_IN_ATOMIC_FETCH_SUB_1
:
4948 case BUILT_IN_ATOMIC_FETCH_SUB_2
:
4949 case BUILT_IN_ATOMIC_FETCH_SUB_4
:
4950 case BUILT_IN_ATOMIC_FETCH_SUB_8
:
4951 case BUILT_IN_ATOMIC_FETCH_SUB_16
:
4952 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_SUB
, stmt
, hbb
);
4955 case BUILT_IN_ATOMIC_FETCH_AND_1
:
4956 case BUILT_IN_ATOMIC_FETCH_AND_2
:
4957 case BUILT_IN_ATOMIC_FETCH_AND_4
:
4958 case BUILT_IN_ATOMIC_FETCH_AND_8
:
4959 case BUILT_IN_ATOMIC_FETCH_AND_16
:
4960 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_AND
, stmt
, hbb
);
4963 case BUILT_IN_ATOMIC_FETCH_XOR_1
:
4964 case BUILT_IN_ATOMIC_FETCH_XOR_2
:
4965 case BUILT_IN_ATOMIC_FETCH_XOR_4
:
4966 case BUILT_IN_ATOMIC_FETCH_XOR_8
:
4967 case BUILT_IN_ATOMIC_FETCH_XOR_16
:
4968 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_XOR
, stmt
, hbb
);
4971 case BUILT_IN_ATOMIC_FETCH_OR_1
:
4972 case BUILT_IN_ATOMIC_FETCH_OR_2
:
4973 case BUILT_IN_ATOMIC_FETCH_OR_4
:
4974 case BUILT_IN_ATOMIC_FETCH_OR_8
:
4975 case BUILT_IN_ATOMIC_FETCH_OR_16
:
4976 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_OR
, stmt
, hbb
);
4979 case BUILT_IN_ATOMIC_STORE_1
:
4980 case BUILT_IN_ATOMIC_STORE_2
:
4981 case BUILT_IN_ATOMIC_STORE_4
:
4982 case BUILT_IN_ATOMIC_STORE_8
:
4983 case BUILT_IN_ATOMIC_STORE_16
:
4984 /* Since there cannot be any LHS, the first parameter is meaningless. */
4985 gen_hsa_ternary_atomic_for_builtin (true, BRIG_ATOMIC_ST
, stmt
, hbb
);
4988 case BUILT_IN_ATOMIC_ADD_FETCH_1
:
4989 case BUILT_IN_ATOMIC_ADD_FETCH_2
:
4990 case BUILT_IN_ATOMIC_ADD_FETCH_4
:
4991 case BUILT_IN_ATOMIC_ADD_FETCH_8
:
4992 case BUILT_IN_ATOMIC_ADD_FETCH_16
:
4993 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_ADD
, stmt
, hbb
);
4996 case BUILT_IN_ATOMIC_SUB_FETCH_1
:
4997 case BUILT_IN_ATOMIC_SUB_FETCH_2
:
4998 case BUILT_IN_ATOMIC_SUB_FETCH_4
:
4999 case BUILT_IN_ATOMIC_SUB_FETCH_8
:
5000 case BUILT_IN_ATOMIC_SUB_FETCH_16
:
5001 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_SUB
, stmt
, hbb
);
5004 case BUILT_IN_ATOMIC_AND_FETCH_1
:
5005 case BUILT_IN_ATOMIC_AND_FETCH_2
:
5006 case BUILT_IN_ATOMIC_AND_FETCH_4
:
5007 case BUILT_IN_ATOMIC_AND_FETCH_8
:
5008 case BUILT_IN_ATOMIC_AND_FETCH_16
:
5009 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_AND
, stmt
, hbb
);
5012 case BUILT_IN_ATOMIC_XOR_FETCH_1
:
5013 case BUILT_IN_ATOMIC_XOR_FETCH_2
:
5014 case BUILT_IN_ATOMIC_XOR_FETCH_4
:
5015 case BUILT_IN_ATOMIC_XOR_FETCH_8
:
5016 case BUILT_IN_ATOMIC_XOR_FETCH_16
:
5017 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_XOR
, stmt
, hbb
);
5020 case BUILT_IN_ATOMIC_OR_FETCH_1
:
5021 case BUILT_IN_ATOMIC_OR_FETCH_2
:
5022 case BUILT_IN_ATOMIC_OR_FETCH_4
:
5023 case BUILT_IN_ATOMIC_OR_FETCH_8
:
5024 case BUILT_IN_ATOMIC_OR_FETCH_16
:
5025 gen_hsa_ternary_atomic_for_builtin (false, BRIG_ATOMIC_OR
, stmt
, hbb
);
5028 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_1
:
5029 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_2
:
5030 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_4
:
5031 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8
:
5032 case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_16
:
5034 /* TODO: Use the appropriate memory model for now. */
5035 tree type
= TREE_TYPE (gimple_call_arg (stmt
, 1));
5038 = hsa_bittype_for_type (hsa_type_for_scalar_tree_type (type
, false));
5040 hsa_insn_atomic
*atominsn
5041 = new hsa_insn_atomic (4, BRIG_OPCODE_ATOMIC
, BRIG_ATOMIC_CAS
, atype
,
5042 BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE
);
5043 hsa_op_address
*addr
;
5044 addr
= get_address_from_value (gimple_call_arg (stmt
, 0), hbb
);
5047 dest
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5049 dest
= new hsa_op_reg (atype
);
5051 /* Should check what the memory scope is. */
5052 atominsn
->m_memoryscope
= BRIG_MEMORY_SCOPE_WORKGROUP
;
5053 atominsn
->set_op (0, dest
);
5054 atominsn
->set_op (1, addr
);
5056 hsa_op_with_type
*op
5057 = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 1), hbb
);
5058 atominsn
->set_op (2, op
);
5059 op
= hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt
, 2), hbb
);
5060 atominsn
->set_op (3, op
);
5062 hbb
->append_insn (atominsn
);
5065 case BUILT_IN_GOMP_PARALLEL
:
5066 HSA_SORRY_AT (gimple_location (stmt
),
5067 "support for HSA does not implement non-gridified "
5068 "OpenMP parallel constructs.");
5070 case BUILT_IN_OMP_GET_THREAD_NUM
:
5072 query_hsa_grid (stmt
, BRIG_OPCODE_WORKITEMABSID
, 0, hbb
);
5076 case BUILT_IN_OMP_GET_NUM_THREADS
:
5078 query_hsa_grid (stmt
, BRIG_OPCODE_GRIDSIZE
, 0, hbb
);
5081 case BUILT_IN_GOMP_TEAMS
:
5083 gen_set_num_threads (gimple_call_arg (stmt
, 1), hbb
);
5086 case BUILT_IN_OMP_GET_NUM_TEAMS
:
5088 gen_get_num_teams (stmt
, hbb
);
5091 case BUILT_IN_OMP_GET_TEAM_NUM
:
5093 gen_get_team_num (stmt
, hbb
);
5096 case BUILT_IN_MEMCPY
:
5097 case BUILT_IN_MEMPCPY
:
5099 tree byte_size
= gimple_call_arg (stmt
, 2);
5101 if (!tree_fits_uhwi_p (byte_size
))
5103 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5107 unsigned n
= tree_to_uhwi (byte_size
);
5109 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5111 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5115 tree dst
= gimple_call_arg (stmt
, 0);
5116 tree src
= gimple_call_arg (stmt
, 1);
5118 hsa_op_address
*dst_addr
= get_address_from_value (dst
, hbb
);
5119 hsa_op_address
*src_addr
= get_address_from_value (src
, hbb
);
5121 gen_hsa_memory_copy (hbb
, dst_addr
, src_addr
, n
);
5123 tree lhs
= gimple_call_lhs (stmt
);
5126 hsa_op_reg
*lhs_reg
= hsa_cfun
->reg_for_gimple_ssa (lhs
);
5127 hsa_op_with_type
*dst_reg
= hsa_reg_or_immed_for_gimple_op (dst
,
5129 hsa_op_with_type
*tmp
;
5131 if (builtin
== BUILT_IN_MEMPCPY
)
5133 tmp
= new hsa_op_reg (dst_reg
->m_type
);
5135 = new hsa_insn_basic (3, BRIG_OPCODE_ADD
, tmp
->m_type
,
5137 new hsa_op_immed (n
, dst_reg
->m_type
));
5138 hbb
->append_insn (add
);
5143 hsa_build_append_simple_mov (lhs_reg
, tmp
, hbb
);
5148 case BUILT_IN_MEMSET
:
5150 tree dst
= gimple_call_arg (stmt
, 0);
5151 tree c
= gimple_call_arg (stmt
, 1);
5153 if (TREE_CODE (c
) != INTEGER_CST
)
5155 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5159 tree byte_size
= gimple_call_arg (stmt
, 2);
5161 if (!tree_fits_uhwi_p (byte_size
))
5163 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5167 unsigned n
= tree_to_uhwi (byte_size
);
5169 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5171 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5175 hsa_op_address
*dst_addr
;
5176 dst_addr
= get_address_from_value (dst
, hbb
);
5177 unsigned HOST_WIDE_INT constant
5178 = tree_to_uhwi (fold_convert (unsigned_char_type_node
, c
));
5180 gen_hsa_memory_set (hbb
, dst_addr
, constant
, n
);
5182 tree lhs
= gimple_call_lhs (stmt
);
5184 gen_hsa_insns_for_single_assignment (lhs
, dst
, hbb
);
5188 case BUILT_IN_BZERO
:
5190 tree dst
= gimple_call_arg (stmt
, 0);
5191 tree byte_size
= gimple_call_arg (stmt
, 1);
5193 if (!tree_fits_uhwi_p (byte_size
))
5195 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5199 unsigned n
= tree_to_uhwi (byte_size
);
5201 if (n
> HSA_MEMORY_BUILTINS_LIMIT
)
5203 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5207 hsa_op_address
*dst_addr
;
5208 dst_addr
= get_address_from_value (dst
, hbb
);
5210 gen_hsa_memory_set (hbb
, dst_addr
, 0, n
);
5214 case BUILT_IN_ALLOCA
:
5215 case BUILT_IN_ALLOCA_WITH_ALIGN
:
5217 gen_hsa_alloca (call
, hbb
);
5222 gen_hsa_insns_for_direct_call (stmt
, hbb
);
5228 /* Generate HSA instructions for a given gimple statement. Instructions will be
5232 gen_hsa_insns_for_gimple_stmt (gimple
*stmt
, hsa_bb
*hbb
)
5234 switch (gimple_code (stmt
))
5237 if (gimple_clobber_p (stmt
))
5240 if (gimple_assign_single_p (stmt
))
5242 tree lhs
= gimple_assign_lhs (stmt
);
5243 tree rhs
= gimple_assign_rhs1 (stmt
);
5244 gen_hsa_insns_for_single_assignment (lhs
, rhs
, hbb
);
5247 gen_hsa_insns_for_operation_assignment (stmt
, hbb
);
5250 gen_hsa_insns_for_return (as_a
<greturn
*> (stmt
), hbb
);
5253 gen_hsa_insns_for_cond_stmt (stmt
, hbb
);
5256 gen_hsa_insns_for_call (stmt
, hbb
);
5259 /* ??? HSA supports some debug facilities. */
5263 tree label
= gimple_label_label (as_a
<glabel
*> (stmt
));
5264 if (FORCED_LABEL (label
))
5265 HSA_SORRY_AT (gimple_location (stmt
),
5266 "support for HSA does not implement gimple label with "
5273 hbb
->append_insn (new hsa_insn_basic (0, BRIG_OPCODE_NOP
));
5278 gen_hsa_insns_for_switch_stmt (as_a
<gswitch
*> (stmt
), hbb
);
5282 HSA_SORRY_ATV (gimple_location (stmt
),
5283 "support for HSA does not implement gimple statement %s",
5284 gimple_code_name
[(int) gimple_code (stmt
)]);
5288 /* Generate a HSA PHI from a gimple PHI. */
5291 gen_hsa_phi_from_gimple_phi (gimple
*phi_stmt
, hsa_bb
*hbb
)
5294 unsigned count
= gimple_phi_num_args (phi_stmt
);
5297 = hsa_cfun
->reg_for_gimple_ssa (gimple_phi_result (phi_stmt
));
5298 hphi
= new hsa_insn_phi (count
, dest
);
5299 hphi
->m_bb
= hbb
->m_bb
;
5301 tree lhs
= gimple_phi_result (phi_stmt
);
5303 for (unsigned i
= 0; i
< count
; i
++)
5305 tree op
= gimple_phi_arg_def (phi_stmt
, i
);
5307 if (TREE_CODE (op
) == SSA_NAME
)
5309 hsa_op_reg
*hreg
= hsa_cfun
->reg_for_gimple_ssa (op
);
5310 hphi
->set_op (i
, hreg
);
5314 gcc_assert (is_gimple_min_invariant (op
));
5315 tree t
= TREE_TYPE (op
);
5316 if (!POINTER_TYPE_P (t
)
5317 || (TREE_CODE (op
) == STRING_CST
5318 && TREE_CODE (TREE_TYPE (t
)) == INTEGER_TYPE
))
5319 hphi
->set_op (i
, new hsa_op_immed (op
));
5320 else if (POINTER_TYPE_P (TREE_TYPE (lhs
))
5321 && TREE_CODE (op
) == INTEGER_CST
)
5323 /* Handle assignment of NULL value to a pointer type. */
5324 hphi
->set_op (i
, new hsa_op_immed (op
));
5326 else if (TREE_CODE (op
) == ADDR_EXPR
)
5328 edge e
= gimple_phi_arg_edge (as_a
<gphi
*> (phi_stmt
), i
);
5329 hsa_bb
*hbb_src
= hsa_init_new_bb (split_edge (e
));
5330 hsa_op_address
*addr
= gen_hsa_addr (TREE_OPERAND (op
, 0),
5333 hsa_op_reg
*dest
= new hsa_op_reg (BRIG_TYPE_U64
);
5334 hsa_insn_basic
*insn
5335 = new hsa_insn_basic (2, BRIG_OPCODE_LDA
, BRIG_TYPE_U64
,
5337 hbb_src
->append_insn (insn
);
5339 hphi
->set_op (i
, dest
);
5343 HSA_SORRY_AT (gimple_location (phi_stmt
),
5344 "support for HSA does not handle PHI nodes with "
5345 "constant address operands");
5351 hphi
->m_prev
= hbb
->m_last_phi
;
5352 hphi
->m_next
= NULL
;
5353 if (hbb
->m_last_phi
)
5354 hbb
->m_last_phi
->m_next
= hphi
;
5355 hbb
->m_last_phi
= hphi
;
5356 if (!hbb
->m_first_phi
)
5357 hbb
->m_first_phi
= hphi
;
5360 /* Constructor of class containing HSA-specific information about a basic
5361 block. CFG_BB is the CFG BB this HSA BB is associated with. IDX is the new
5362 index of this BB (so that the constructor does not attempt to use
5363 hsa_cfun during its construction). */
5365 hsa_bb::hsa_bb (basic_block cfg_bb
, int idx
)
5366 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5367 m_last_phi (NULL
), m_index (idx
), m_liveout (BITMAP_ALLOC (NULL
)),
5368 m_livein (BITMAP_ALLOC (NULL
))
5370 gcc_assert (!cfg_bb
->aux
);
5374 /* Constructor of class containing HSA-specific information about a basic
5375 block. CFG_BB is the CFG BB this HSA BB is associated with. */
5377 hsa_bb::hsa_bb (basic_block cfg_bb
)
5378 : m_bb (cfg_bb
), m_first_insn (NULL
), m_last_insn (NULL
), m_first_phi (NULL
),
5379 m_last_phi (NULL
), m_index (hsa_cfun
->m_hbb_count
++),
5380 m_liveout (BITMAP_ALLOC (NULL
)), m_livein (BITMAP_ALLOC (NULL
))
5382 gcc_assert (!cfg_bb
->aux
);
5386 /* Destructor of class representing HSA BB. */
5390 BITMAP_FREE (m_livein
);
5391 BITMAP_FREE (m_liveout
);
5394 /* Create and initialize and return a new hsa_bb structure for a given CFG
5398 hsa_init_new_bb (basic_block bb
)
5400 return new (*hsa_allocp_bb
) hsa_bb (bb
);
5403 /* Initialize OMP in an HSA basic block PROLOGUE. */
5406 init_prologue (void)
5408 if (!hsa_cfun
->m_kern_p
)
5411 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5413 /* Create a magic number that is going to be printed by libgomp. */
5414 unsigned index
= hsa_get_number_decl_kernel_mappings ();
5416 /* Emit store to debug argument. */
5417 if (PARAM_VALUE (PARAM_HSA_GEN_DEBUG_STORES
) > 0)
5418 set_debug_value (prologue
, new hsa_op_immed (1000 + index
, BRIG_TYPE_U64
));
5421 /* Initialize hsa_num_threads to a default value. */
5424 init_hsa_num_threads (void)
5426 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5428 /* Save the default value to private variable hsa_num_threads. */
5429 hsa_insn_basic
*basic
5430 = new hsa_insn_mem (BRIG_OPCODE_ST
, hsa_num_threads
->m_type
,
5431 new hsa_op_immed (0, hsa_num_threads
->m_type
),
5432 new hsa_op_address (hsa_num_threads
));
5433 prologue
->append_insn (basic
);
5436 /* Go over gimple representation and generate our internal HSA one. */
5439 gen_body_from_gimple ()
5443 /* Verify CFG for complex edges we are unable to handle. */
5447 FOR_EACH_BB_FN (bb
, cfun
)
5449 FOR_EACH_EDGE (e
, ei
, bb
->succs
)
5451 /* Verify all unsupported flags for edges that point
5452 to the same basic block. */
5453 if (e
->flags
& EDGE_EH
)
5455 HSA_SORRY_AT (UNKNOWN_LOCATION
,
5456 "support for HSA does not implement exception "
5463 FOR_EACH_BB_FN (bb
, cfun
)
5465 gimple_stmt_iterator gsi
;
5466 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5470 hbb
= hsa_init_new_bb (bb
);
5472 for (gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5474 gen_hsa_insns_for_gimple_stmt (gsi_stmt (gsi
), hbb
);
5475 if (hsa_seen_error ())
5480 FOR_EACH_BB_FN (bb
, cfun
)
5482 gimple_stmt_iterator gsi
;
5483 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
5484 gcc_assert (hbb
!= NULL
);
5486 for (gsi
= gsi_start_phis (bb
); !gsi_end_p (gsi
); gsi_next (&gsi
))
5487 if (!virtual_operand_p (gimple_phi_result (gsi_stmt (gsi
))))
5488 gen_hsa_phi_from_gimple_phi (gsi_stmt (gsi
), hbb
);
5493 fprintf (dump_file
, "------- Generated SSA form -------\n");
5494 dump_hsa_cfun (dump_file
);
5499 gen_function_decl_parameters (hsa_function_representation
*f
,
5505 for (parm
= TYPE_ARG_TYPES (TREE_TYPE (decl
)), i
= 0;
5507 parm
= TREE_CHAIN (parm
), i
++)
5509 /* Result type if last in the tree list. */
5510 if (TREE_CHAIN (parm
) == NULL
)
5513 tree v
= TREE_VALUE (parm
);
5515 hsa_symbol
*arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5517 arg
->m_type
= hsa_type_for_tree_type (v
, &arg
->m_dim
);
5518 arg
->m_name_number
= i
;
5520 f
->m_input_args
.safe_push (arg
);
5523 tree result_type
= TREE_TYPE (TREE_TYPE (decl
));
5524 if (!VOID_TYPE_P (result_type
))
5526 f
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5528 f
->m_output_arg
->m_type
5529 = hsa_type_for_tree_type (result_type
, &f
->m_output_arg
->m_dim
);
5530 f
->m_output_arg
->m_name
= "res";
5534 /* Generate the vector of parameters of the HSA representation of the current
5535 function. This also includes the output parameter representing the
5539 gen_function_def_parameters ()
5543 hsa_bb
*prologue
= hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun
));
5545 for (parm
= DECL_ARGUMENTS (cfun
->decl
); parm
;
5546 parm
= DECL_CHAIN (parm
))
5548 struct hsa_symbol
**slot
;
5551 = new hsa_symbol (BRIG_TYPE_NONE
, hsa_cfun
->m_kern_p
5552 ? BRIG_SEGMENT_KERNARG
: BRIG_SEGMENT_ARG
,
5553 BRIG_LINKAGE_FUNCTION
);
5554 arg
->fillup_for_decl (parm
);
5556 hsa_cfun
->m_input_args
.safe_push (arg
);
5558 if (hsa_seen_error ())
5561 arg
->m_name
= hsa_get_declaration_name (parm
);
5563 /* Copy all input arguments and create corresponding private symbols
5565 hsa_symbol
*private_arg
;
5566 hsa_op_address
*parm_addr
= new hsa_op_address (arg
);
5568 if (TREE_ADDRESSABLE (parm
)
5569 || (!is_gimple_reg (parm
) && !TREE_READONLY (parm
)))
5571 private_arg
= hsa_cfun
->create_hsa_temporary (arg
->m_type
);
5572 private_arg
->fillup_for_decl (parm
);
5574 hsa_op_address
*private_arg_addr
= new hsa_op_address (private_arg
);
5575 gen_hsa_memory_copy (prologue
, private_arg_addr
, parm_addr
,
5576 arg
->total_byte_size ());
5581 slot
= hsa_cfun
->m_local_symbols
->find_slot (private_arg
, INSERT
);
5582 gcc_assert (!*slot
);
5583 *slot
= private_arg
;
5585 if (is_gimple_reg (parm
))
5587 tree ddef
= ssa_default_def (cfun
, parm
);
5588 if (ddef
&& !has_zero_uses (ddef
))
5590 BrigType16_t t
= hsa_type_for_scalar_tree_type (TREE_TYPE (ddef
),
5592 BrigType16_t mtype
= mem_type_for_type (t
);
5593 hsa_op_reg
*dest
= hsa_cfun
->reg_for_gimple_ssa (ddef
);
5594 hsa_insn_mem
*mem
= new hsa_insn_mem (BRIG_OPCODE_LD
, mtype
,
5596 gcc_assert (!parm_addr
->m_reg
);
5597 prologue
->append_insn (mem
);
5602 if (!VOID_TYPE_P (TREE_TYPE (TREE_TYPE (cfun
->decl
))))
5604 struct hsa_symbol
**slot
;
5606 hsa_cfun
->m_output_arg
= new hsa_symbol (BRIG_TYPE_NONE
, BRIG_SEGMENT_ARG
,
5607 BRIG_LINKAGE_FUNCTION
);
5608 hsa_cfun
->m_output_arg
->fillup_for_decl (DECL_RESULT (cfun
->decl
));
5610 if (hsa_seen_error ())
5613 hsa_cfun
->m_output_arg
->m_name
= "res";
5614 slot
= hsa_cfun
->m_local_symbols
->find_slot (hsa_cfun
->m_output_arg
,
5616 gcc_assert (!*slot
);
5617 *slot
= hsa_cfun
->m_output_arg
;
5621 /* Generate function representation that corresponds to
5622 a function declaration. */
5624 hsa_function_representation
*
5625 hsa_generate_function_declaration (tree decl
)
5627 hsa_function_representation
*fun
5628 = new hsa_function_representation (decl
, false, 0);
5630 fun
->m_declaration_p
= true;
5631 fun
->m_name
= get_brig_function_name (decl
);
5632 gen_function_decl_parameters (fun
, decl
);
5638 /* Generate function representation that corresponds to
5641 hsa_function_representation
*
5642 hsa_generate_internal_fn_decl (hsa_internal_fn
*fn
)
5644 hsa_function_representation
*fun
= new hsa_function_representation (fn
);
5646 fun
->m_name
= fn
->name ();
5648 for (unsigned i
= 0; i
< fn
->get_arity (); i
++)
5651 = new hsa_symbol (fn
->get_argument_type (i
), BRIG_SEGMENT_ARG
,
5653 arg
->m_name_number
= i
;
5654 fun
->m_input_args
.safe_push (arg
);
5657 fun
->m_output_arg
= new hsa_symbol (fn
->get_argument_type (-1),
5658 BRIG_SEGMENT_ARG
, BRIG_LINKAGE_NONE
);
5659 fun
->m_output_arg
->m_name
= "res";
5664 /* Return true if switch statement S can be transformed
5665 to a SBR instruction in HSAIL. */
5668 transformable_switch_to_sbr_p (gswitch
*s
)
5670 /* Identify if a switch statement can be transformed to
5671 SBR instruction, like:
5673 sbr_u32 $s1 [@label1, @label2, @label3];
5676 tree size
= get_switch_size (s
);
5677 if (!tree_fits_uhwi_p (size
))
5680 if (tree_to_uhwi (size
) > HSA_MAXIMUM_SBR_LABELS
)
5686 /* Structure hold connection between PHI nodes and immediate
5687 values hold by there nodes. */
5689 struct phi_definition
5691 phi_definition (unsigned phi_i
, unsigned label_i
, tree imm
):
5692 phi_index (phi_i
), label_index (label_i
), phi_value (imm
)
5696 unsigned label_index
;
5700 /* Sum slice of a vector V, starting from index START and ending
5701 at the index END - 1. */
5703 template <typename T
>
5705 T
sum_slice (const auto_vec
<T
> &v
, unsigned start
, unsigned end
)
5709 for (unsigned i
= start
; i
< end
; i
++)
5715 /* Function transforms GIMPLE SWITCH statements to a series of IF statements.
5716 Let's assume following example:
5730 The transformation encompasses following steps:
5731 1) all immediate values used by edges coming from the switch basic block
5733 2) all these edges are removed
5734 3) the switch statement (in L0) is replaced by:
5740 4) newly created basic block Lx' is used for generation of
5742 5) else branch of the last condition goes to LD
5743 6) fix all immediate values in PHI nodes that were propagated though
5744 edges that were removed in step 2
5746 Note: if a case is made by a range C1..C2, then process
5747 following transformation:
5749 switch_cond_op1 = C1 <= index;
5750 switch_cond_op2 = index <= C2;
5751 switch_cond_and = switch_cond_op1 & switch_cond_op2;
5752 if (switch_cond_and != 0)
5760 convert_switch_statements ()
5762 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
5765 bool need_update
= false;
5767 FOR_EACH_BB_FN (bb
, func
)
5769 gimple_stmt_iterator gsi
= gsi_last_bb (bb
);
5770 if (gsi_end_p (gsi
))
5773 gimple
*stmt
= gsi_stmt (gsi
);
5775 if (gimple_code (stmt
) == GIMPLE_SWITCH
)
5777 gswitch
*s
= as_a
<gswitch
*> (stmt
);
5779 /* If the switch can utilize SBR insn, skip the statement. */
5780 if (transformable_switch_to_sbr_p (s
))
5785 unsigned labels
= gimple_switch_num_labels (s
);
5786 tree index
= gimple_switch_index (s
);
5787 tree index_type
= TREE_TYPE (index
);
5788 tree default_label
= gimple_switch_default_label (s
);
5789 basic_block default_label_bb
5790 = label_to_block_fn (func
, CASE_LABEL (default_label
));
5791 basic_block cur_bb
= bb
;
5793 auto_vec
<edge
> new_edges
;
5794 auto_vec
<phi_definition
*> phi_todo_list
;
5795 auto_vec
<gcov_type
> edge_counts
;
5796 auto_vec
<int> edge_probabilities
;
5798 /* Investigate all labels that and PHI nodes in these edges which
5799 should be fixed after we add new collection of edges. */
5800 for (unsigned i
= 0; i
< labels
; i
++)
5802 tree label
= gimple_switch_label (s
, i
);
5803 basic_block label_bb
= label_to_block_fn (func
, CASE_LABEL (label
));
5804 edge e
= find_edge (bb
, label_bb
);
5805 edge_counts
.safe_push (e
->count
);
5806 edge_probabilities
.safe_push (e
->probability
);
5807 gphi_iterator phi_gsi
;
5809 /* Save PHI definitions that will be destroyed because of an edge
5810 is going to be removed. */
5811 unsigned phi_index
= 0;
5812 for (phi_gsi
= gsi_start_phis (e
->dest
);
5813 !gsi_end_p (phi_gsi
); gsi_next (&phi_gsi
))
5815 gphi
*phi
= phi_gsi
.phi ();
5816 for (unsigned j
= 0; j
< gimple_phi_num_args (phi
); j
++)
5818 if (gimple_phi_arg_edge (phi
, j
) == e
)
5820 tree imm
= gimple_phi_arg_def (phi
, j
);
5821 phi_definition
*p
= new phi_definition (phi_index
, i
,
5823 phi_todo_list
.safe_push (p
);
5831 /* Remove all edges for the current basic block. */
5832 for (int i
= EDGE_COUNT (bb
->succs
) - 1; i
>= 0; i
--)
5834 edge e
= EDGE_SUCC (bb
, i
);
5838 /* Iterate all non-default labels. */
5839 for (unsigned i
= 1; i
< labels
; i
++)
5841 tree label
= gimple_switch_label (s
, i
);
5842 tree low
= CASE_LOW (label
);
5843 tree high
= CASE_HIGH (label
);
5845 if (!useless_type_conversion_p (TREE_TYPE (low
), index_type
))
5846 low
= fold_convert (index_type
, low
);
5848 gimple_stmt_iterator cond_gsi
= gsi_last_bb (cur_bb
);
5852 tree tmp1
= make_temp_ssa_name (boolean_type_node
, NULL
,
5855 gimple
*assign1
= gimple_build_assign (tmp1
, LE_EXPR
, low
,
5858 tree tmp2
= make_temp_ssa_name (boolean_type_node
, NULL
,
5861 if (!useless_type_conversion_p (TREE_TYPE (high
), index_type
))
5862 high
= fold_convert (index_type
, high
);
5863 gimple
*assign2
= gimple_build_assign (tmp2
, LE_EXPR
, index
,
5866 tree tmp3
= make_temp_ssa_name (boolean_type_node
, NULL
,
5868 gimple
*assign3
= gimple_build_assign (tmp3
, BIT_AND_EXPR
, tmp1
,
5871 gsi_insert_before (&cond_gsi
, assign1
, GSI_SAME_STMT
);
5872 gsi_insert_before (&cond_gsi
, assign2
, GSI_SAME_STMT
);
5873 gsi_insert_before (&cond_gsi
, assign3
, GSI_SAME_STMT
);
5875 tree b
= constant_boolean_node (false, boolean_type_node
);
5876 c
= gimple_build_cond (NE_EXPR
, tmp3
, b
, NULL
, NULL
);
5879 c
= gimple_build_cond (EQ_EXPR
, index
, low
, NULL
, NULL
);
5881 gimple_set_location (c
, gimple_location (stmt
));
5883 gsi_insert_before (&cond_gsi
, c
, GSI_SAME_STMT
);
5885 basic_block label_bb
5886 = label_to_block_fn (func
, CASE_LABEL (label
));
5887 edge new_edge
= make_edge (cur_bb
, label_bb
, EDGE_TRUE_VALUE
);
5888 int prob_sum
= sum_slice
<int> (edge_probabilities
, i
, labels
) +
5889 edge_probabilities
[0];
5892 new_edge
->probability
5893 = RDIV (REG_BR_PROB_BASE
* edge_probabilities
[i
], prob_sum
);
5895 new_edge
->count
= edge_counts
[i
];
5896 new_edges
.safe_push (new_edge
);
5900 /* Prepare another basic block that will contain
5902 basic_block next_bb
= create_empty_bb (cur_bb
);
5905 add_bb_to_loop (next_bb
, cur_bb
->loop_father
);
5906 loops_state_set (LOOPS_NEED_FIXUP
);
5909 edge next_edge
= make_edge (cur_bb
, next_bb
, EDGE_FALSE_VALUE
);
5910 next_edge
->probability
5911 = inverse_probability (new_edge
->probability
);
5912 next_edge
->count
= edge_counts
[0]
5913 + sum_slice
<gcov_type
> (edge_counts
, i
, labels
);
5914 next_bb
->frequency
= EDGE_FREQUENCY (next_edge
);
5917 else /* Link last IF statement and default label
5920 edge e
= make_edge (cur_bb
, default_label_bb
, EDGE_FALSE_VALUE
);
5921 e
->probability
= inverse_probability (new_edge
->probability
);
5922 e
->count
= edge_counts
[0];
5923 new_edges
.safe_insert (0, e
);
5927 /* Restore original PHI immediate value. */
5928 for (unsigned i
= 0; i
< phi_todo_list
.length (); i
++)
5930 phi_definition
*phi_def
= phi_todo_list
[i
];
5931 edge new_edge
= new_edges
[phi_def
->label_index
];
5933 gphi_iterator it
= gsi_start_phis (new_edge
->dest
);
5934 for (unsigned i
= 0; i
< phi_def
->phi_index
; i
++)
5937 gphi
*phi
= it
.phi ();
5938 add_phi_arg (phi
, phi_def
->phi_value
, new_edge
, UNKNOWN_LOCATION
);
5942 /* Remove the original GIMPLE switch statement. */
5943 gsi_remove (&gsi
, true);
5948 dump_function_to_file (current_function_decl
, dump_file
, TDF_DETAILS
);
5952 free_dominance_info (CDI_DOMINATORS
);
5953 calculate_dominance_info (CDI_DOMINATORS
);
5957 /* Expand builtins that can't be handled by HSA back-end. */
5962 function
*func
= DECL_STRUCT_FUNCTION (current_function_decl
);
5965 FOR_EACH_BB_FN (bb
, func
)
5967 for (gimple_stmt_iterator gsi
= gsi_start_bb (bb
); !gsi_end_p (gsi
);
5970 gimple
*stmt
= gsi_stmt (gsi
);
5972 if (gimple_code (stmt
) != GIMPLE_CALL
)
5975 gcall
*call
= as_a
<gcall
*> (stmt
);
5977 if (!gimple_call_builtin_p (call
, BUILT_IN_NORMAL
))
5980 tree fndecl
= gimple_call_fndecl (stmt
);
5981 enum built_in_function fn
= DECL_FUNCTION_CODE (fndecl
);
5984 case BUILT_IN_CEXPF
:
5985 case BUILT_IN_CEXPIF
:
5986 case BUILT_IN_CEXPI
:
5988 /* Similar to builtins.c (expand_builtin_cexpi), the builtin
5989 can be transformed to: cexp(I * z) = ccos(z) + I * csin(z). */
5990 tree lhs
= gimple_call_lhs (stmt
);
5991 tree rhs
= gimple_call_arg (stmt
, 0);
5992 tree rhs_type
= TREE_TYPE (rhs
);
5993 bool float_type_p
= rhs_type
== float_type_node
;
5994 tree real_part
= make_temp_ssa_name (rhs_type
, NULL
,
5996 tree imag_part
= make_temp_ssa_name (rhs_type
, NULL
,
6000 = mathfn_built_in (rhs_type
, fn
== float_type_p
6001 ? BUILT_IN_COSF
: BUILT_IN_COS
);
6002 gcall
*cos
= gimple_build_call (cos_fndecl
, 1, rhs
);
6003 gimple_call_set_lhs (cos
, real_part
);
6004 gsi_insert_before (&gsi
, cos
, GSI_SAME_STMT
);
6007 = mathfn_built_in (rhs_type
, fn
== float_type_p
6008 ? BUILT_IN_SINF
: BUILT_IN_SIN
);
6009 gcall
*sin
= gimple_build_call (sin_fndecl
, 1, rhs
);
6010 gimple_call_set_lhs (sin
, imag_part
);
6011 gsi_insert_before (&gsi
, sin
, GSI_SAME_STMT
);
6014 gassign
*assign
= gimple_build_assign (lhs
, COMPLEX_EXPR
,
6015 real_part
, imag_part
);
6016 gsi_insert_before (&gsi
, assign
, GSI_SAME_STMT
);
6017 gsi_remove (&gsi
, true);
6028 /* Emit HSA module variables that are global for the entire module. */
6031 emit_hsa_module_variables (void)
6033 hsa_num_threads
= new hsa_symbol (BRIG_TYPE_U32
, BRIG_SEGMENT_PRIVATE
,
6034 BRIG_LINKAGE_MODULE
, true);
6036 hsa_num_threads
->m_name
= "hsa_num_threads";
6038 hsa_brig_emit_omp_symbols ();
6041 /* Generate HSAIL representation of the current function and write into a
6042 special section of the output file. If KERNEL is set, the function will be
6043 considered an HSA kernel callable from the host, otherwise it will be
6044 compiled as an HSA function callable from other HSA code. */
6047 generate_hsa (bool kernel
)
6049 hsa_init_data_for_cfun ();
6051 if (hsa_num_threads
== NULL
)
6052 emit_hsa_module_variables ();
6054 /* Initialize hsa_cfun. */
6055 hsa_cfun
= new hsa_function_representation (cfun
->decl
, kernel
,
6056 SSANAMES (cfun
)->length ());
6057 hsa_cfun
->init_extra_bbs ();
6061 HSA_SORRY_AT (UNKNOWN_LOCATION
,
6062 "support for HSA does not implement transactional memory");
6066 verify_function_arguments (cfun
->decl
);
6067 if (hsa_seen_error ())
6070 hsa_cfun
->m_name
= get_brig_function_name (cfun
->decl
);
6072 gen_function_def_parameters ();
6073 if (hsa_seen_error ())
6078 gen_body_from_gimple ();
6079 if (hsa_seen_error ())
6082 if (hsa_cfun
->m_kernel_dispatch_count
)
6083 init_hsa_num_threads ();
6085 if (hsa_cfun
->m_kern_p
)
6087 hsa_function_summary
*s
6088 = hsa_summaries
->get (cgraph_node::get (hsa_cfun
->m_decl
));
6089 hsa_add_kern_decl_mapping (current_function_decl
, hsa_cfun
->m_name
,
6090 hsa_cfun
->m_maximum_omp_data_size
,
6091 s
->m_gridified_kernel_p
);
6094 #ifdef ENABLE_CHECKING
6095 for (unsigned i
= 0; i
< hsa_cfun
->m_ssa_map
.length (); i
++)
6096 if (hsa_cfun
->m_ssa_map
[i
])
6097 hsa_cfun
->m_ssa_map
[i
]->verify_ssa ();
6100 FOR_EACH_BB_FN (bb
, cfun
)
6102 hsa_bb
*hbb
= hsa_bb_for_bb (bb
);
6104 for (hsa_insn_basic
*insn
= hbb
->m_first_insn
; insn
; insn
= insn
->m_next
)
6111 hsa_brig_emit_function ();
6114 hsa_deinit_data_for_cfun ();
6119 const pass_data pass_data_gen_hsail
=
6122 "hsagen", /* name */
6123 OPTGROUP_NONE
, /* optinfo_flags */
6124 TV_NONE
, /* tv_id */
6125 PROP_cfg
| PROP_ssa
, /* properties_required */
6126 0, /* properties_provided */
6127 0, /* properties_destroyed */
6128 0, /* todo_flags_start */
6129 0 /* todo_flags_finish */
6132 class pass_gen_hsail
: public gimple_opt_pass
6135 pass_gen_hsail (gcc::context
*ctxt
)
6136 : gimple_opt_pass(pass_data_gen_hsail
, ctxt
)
6139 /* opt_pass methods: */
6140 bool gate (function
*);
6141 unsigned int execute (function
*);
6143 }; // class pass_gen_hsail
6145 /* Determine whether or not to run generation of HSAIL. */
6148 pass_gen_hsail::gate (function
*f
)
6150 return hsa_gen_requested_p ()
6151 && hsa_gpu_implementation_p (f
->decl
);
6155 pass_gen_hsail::execute (function
*)
6157 hsa_function_summary
*s
6158 = hsa_summaries
->get (cgraph_node::get_create (current_function_decl
));
6160 convert_switch_statements ();
6162 generate_hsa (s
->m_kind
== HSA_KERNEL
);
6163 TREE_ASM_WRITTEN (current_function_decl
) = 1;
6164 return TODO_discard_function
;
6169 /* Create the instance of hsa gen pass. */
6172 make_pass_gen_hsail (gcc::context
*ctxt
)
6174 return new pass_gen_hsail (ctxt
);