1/* A pass for lowering gimple to HSAIL
2 Copyright (C) 2013-2017 Free Software Foundation, Inc.
3 Contributed by Martin Jambor <mjambor@suse.cz> and
4 Martin Liska <mliska@suse.cz>.
5
6This file is part of GCC.
7
8GCC is free software; you can redistribute it and/or modify
9it under the terms of the GNU General Public License as published by
10the Free Software Foundation; either version 3, or (at your option)
11any later version.
12
13GCC is distributed in the hope that it will be useful,
14but WITHOUT ANY WARRANTY; without even the implied warranty of
15MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
16GNU General Public License for more details.
17
18You should have received a copy of the GNU General Public License
19along with GCC; see the file COPYING3. If not see
20<http://www.gnu.org/licenses/>. */
21
22#include "config.h"
23#include "system.h"
24#include "coretypes.h"
25#include "memmodel.h"
26#include "tm.h"
27#include "is-a.h"
28#include "hash-table.h"
29#include "vec.h"
30#include "tree.h"
31#include "tree-pass.h"
32#include "function.h"
33#include "basic-block.h"
34#include "cfg.h"
35#include "fold-const.h"
36#include "gimple.h"
37#include "gimple-iterator.h"
38#include "bitmap.h"
39#include "dumpfile.h"
40#include "gimple-pretty-print.h"
41#include "diagnostic-core.h"
42#include "gimple-ssa.h"
43#include "tree-phinodes.h"
44#include "stringpool.h"
45#include "tree-vrp.h"
46#include "tree-ssanames.h"
47#include "tree-dfa.h"
48#include "ssa-iterators.h"
49#include "cgraph.h"
50#include "print-tree.h"
51#include "symbol-summary.h"
52#include "hsa-common.h"
53#include "cfghooks.h"
54#include "tree-cfg.h"
55#include "cfgloop.h"
56#include "cfganal.h"
57#include "builtins.h"
58#include "params.h"
59#include "gomp-constants.h"
60#include "internal-fn.h"
61#include "builtins.h"
62#include "stor-layout.h"
63#include "stringpool.h"
64#include "attribs.h"
65
66/* Print a warning message and set that we have seen an error. */
67
68#define HSA_SORRY_ATV(location, message, ...) \
69 do \
70 { \
71 hsa_fail_cfun (); \
72 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
73 HSA_SORRY_MSG)) \
74 inform (location, message, __VA_ARGS__); \
75 } \
76 while (false)
77
78/* Same as previous, but highlight a location. */
79
80#define HSA_SORRY_AT(location, message) \
81 do \
82 { \
83 hsa_fail_cfun (); \
84 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
85 HSA_SORRY_MSG)) \
86 inform (location, message); \
87 } \
88 while (false)
89
90/* Default number of threads used by kernel dispatch. */
91
92#define HSA_DEFAULT_NUM_THREADS 64
93
94/* Following structures are defined in the final version
95 of HSA specification. */
96
97/* HSA queue packet is shadow structure, originally provided by AMD. */
98
99struct hsa_queue_packet
100{
101 uint16_t header;
102 uint16_t setup;
103 uint16_t workgroup_size_x;
104 uint16_t workgroup_size_y;
105 uint16_t workgroup_size_z;
106 uint16_t reserved0;
107 uint32_t grid_size_x;
108 uint32_t grid_size_y;
109 uint32_t grid_size_z;
110 uint32_t private_segment_size;
111 uint32_t group_segment_size;
112 uint64_t kernel_object;
113 void *kernarg_address;
114 uint64_t reserved2;
115 uint64_t completion_signal;
116};
117
118/* HSA queue is shadow structure, originally provided by AMD. */
119
120struct hsa_queue
121{
122 int type;
123 uint32_t features;
124 void *base_address;
125 uint64_t doorbell_signal;
126 uint32_t size;
127 uint32_t reserved1;
128 uint64_t id;
129};
130
131static struct obstack hsa_obstack;
132
133/* List of pointers to all instructions that come from an object allocator. */
134static vec <hsa_insn_basic *> hsa_instructions;
135
136/* List of pointers to all operands that come from an object allocator. */
137static vec <hsa_op_base *> hsa_operands;
138
139hsa_symbol::hsa_symbol ()
140 : m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
141 m_directive_offset (0), m_type (BRIG_TYPE_NONE),
142 m_segment (BRIG_SEGMENT_NONE), m_linkage (BRIG_LINKAGE_NONE), m_dim (0),
143 m_cst_value (NULL), m_global_scope_p (false), m_seen_error (false),
144 m_allocation (BRIG_ALLOCATION_AUTOMATIC), m_emitted_to_brig (false)
145{
146}
147
148
149hsa_symbol::hsa_symbol (BrigType16_t type, BrigSegment8_t segment,
150 BrigLinkage8_t linkage, bool global_scope_p,
151 BrigAllocation allocation, BrigAlignment8_t align)
152 : m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
153 m_directive_offset (0), m_type (type), m_segment (segment),
154 m_linkage (linkage), m_dim (0), m_cst_value (NULL),
155 m_global_scope_p (global_scope_p), m_seen_error (false),
156 m_allocation (allocation), m_emitted_to_brig (false), m_align (align)
157{
158}
159
160unsigned HOST_WIDE_INT
161hsa_symbol::total_byte_size ()
162{
163 unsigned HOST_WIDE_INT s
164 = hsa_type_bit_size (~BRIG_TYPE_ARRAY_MASK & m_type);
165 gcc_assert (s % BITS_PER_UNIT == 0);
166 s /= BITS_PER_UNIT;
167
168 if (m_dim)
169 s *= m_dim;
170
171 return s;
172}
173
174/* Forward declaration. */
175
176static BrigType16_t
177hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p,
178 bool min32int);
179
180void
181hsa_symbol::fillup_for_decl (tree decl)
182{
183 m_decl = decl;
184 m_type = hsa_type_for_tree_type (TREE_TYPE (decl), &m_dim, false);
185 if (hsa_seen_error ())
186 {
187 m_seen_error = true;
188 return;
189 }
190
191 m_align = MAX (m_align, hsa_natural_alignment (m_type));
192}
193
194/* Constructor of class representing global HSA function/kernel information and
195 state. FNDECL is function declaration, KERNEL_P is true if the function
196 is going to become a HSA kernel. If the function has body, SSA_NAMES_COUNT
197 should be set to number of SSA names used in the function.
198 MODIFIED_CFG is set to true in case we modified control-flow graph
199 of the function. */
200
201hsa_function_representation::hsa_function_representation
202 (tree fdecl, bool kernel_p, unsigned ssa_names_count, bool modified_cfg)
203 : m_name (NULL),
204 m_reg_count (0), m_input_args (vNULL),
205 m_output_arg (NULL), m_spill_symbols (vNULL), m_global_symbols (vNULL),
206 m_private_variables (vNULL), m_called_functions (vNULL),
207 m_called_internal_fns (vNULL), m_hbb_count (0),
208 m_in_ssa (true), m_kern_p (kernel_p), m_declaration_p (false),
209 m_decl (fdecl), m_internal_fn (NULL), m_shadow_reg (NULL),
210 m_kernel_dispatch_count (0), m_maximum_omp_data_size (0),
211 m_seen_error (false), m_temp_symbol_count (0), m_ssa_map (),
212 m_modified_cfg (modified_cfg)
213{
214 int sym_init_len = (vec_safe_length (cfun->local_decls) / 2) + 1;
215 m_local_symbols = new hash_table <hsa_noop_symbol_hasher> (sym_init_len);
216 m_ssa_map.safe_grow_cleared (ssa_names_count);
217}
218
219/* Constructor of class representing HSA function information that
220 is derived for an internal function. */
221hsa_function_representation::hsa_function_representation (hsa_internal_fn *fn)
222 : m_reg_count (0), m_input_args (vNULL),
223 m_output_arg (NULL), m_local_symbols (NULL),
224 m_spill_symbols (vNULL), m_global_symbols (vNULL),
225 m_private_variables (vNULL), m_called_functions (vNULL),
226 m_called_internal_fns (vNULL), m_hbb_count (0),
227 m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL),
228 m_internal_fn (fn), m_shadow_reg (NULL), m_kernel_dispatch_count (0),
229 m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
230 m_ssa_map () {}
231
232/* Destructor of class holding function/kernel-wide information and state. */
233
234hsa_function_representation::~hsa_function_representation ()
235{
236 /* Kernel names are deallocated at the end of BRIG output when deallocating
237 hsa_decl_kernel_mapping. */
238 if (!m_kern_p || m_seen_error)
239 free (m_name);
240
241 for (unsigned i = 0; i < m_input_args.length (); i++)
242 delete m_input_args[i];
243 m_input_args.release ();
244
245 delete m_output_arg;
246 delete m_local_symbols;
247
248 for (unsigned i = 0; i < m_spill_symbols.length (); i++)
249 delete m_spill_symbols[i];
250 m_spill_symbols.release ();
251
252 hsa_symbol *sym;
253 for (unsigned i = 0; i < m_global_symbols.iterate (i, &sym); i++)
254 if (sym->m_linkage != BRIG_ALLOCATION_PROGRAM)
255 delete sym;
256 m_global_symbols.release ();
257
258 for (unsigned i = 0; i < m_private_variables.length (); i++)
259 delete m_private_variables[i];
260 m_private_variables.release ();
261 m_called_functions.release ();
262 m_ssa_map.release ();
263
264 for (unsigned i = 0; i < m_called_internal_fns.length (); i++)
265 delete m_called_internal_fns[i];
266}
267
268hsa_op_reg *
269hsa_function_representation::get_shadow_reg ()
270{
271 /* If we compile a function with kernel dispatch and does not set
272 an optimization level, the function won't be inlined and
273 we return NULL. */
274 if (!m_kern_p)
275 return NULL;
276
277 if (m_shadow_reg)
278 return m_shadow_reg;
279
280 /* Append the shadow argument. */
281 hsa_symbol *shadow = new hsa_symbol (BRIG_TYPE_U64, BRIG_SEGMENT_KERNARG,
282 BRIG_LINKAGE_FUNCTION);
283 m_input_args.safe_push (shadow);
284 shadow->m_name = "hsa_runtime_shadow";
285
286 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_U64);
287 hsa_op_address *addr = new hsa_op_address (shadow);
288
289 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, r, addr);
290 hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun))->append_insn (mem);
291 m_shadow_reg = r;
292
293 return r;
294}
295
296bool hsa_function_representation::has_shadow_reg_p ()
297{
298 return m_shadow_reg != NULL;
299}
300
301void
302hsa_function_representation::init_extra_bbs ()
303{
304 hsa_init_new_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
305 hsa_init_new_bb (EXIT_BLOCK_PTR_FOR_FN (cfun));
306}
307
308void
309hsa_function_representation::update_dominance ()
310{
311 if (m_modified_cfg)
312 {
313 free_dominance_info (CDI_DOMINATORS);
314 calculate_dominance_info (CDI_DOMINATORS);
315 }
316}
317
318hsa_symbol *
319hsa_function_representation::create_hsa_temporary (BrigType16_t type)
320{
321 hsa_symbol *s = new hsa_symbol (type, BRIG_SEGMENT_PRIVATE,
322 BRIG_LINKAGE_FUNCTION);
323 s->m_name_number = m_temp_symbol_count++;
324
325 hsa_cfun->m_private_variables.safe_push (s);
326 return s;
327}
328
329BrigLinkage8_t
330hsa_function_representation::get_linkage ()
331{
332 if (m_internal_fn)
333 return BRIG_LINKAGE_PROGRAM;
334
335 return m_kern_p || TREE_PUBLIC (m_decl) ?
336 BRIG_LINKAGE_PROGRAM : BRIG_LINKAGE_MODULE;
337}
338
339/* Hash map of simple OMP builtins. */
340static hash_map <nofree_string_hash, omp_simple_builtin> *omp_simple_builtins
341 = NULL;
342
343/* Warning messages for OMP builtins. */
344
345#define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP " \
346 "lock routines"
347#define HSA_WARN_TIMING_ROUTINE "support for HSA does not implement OpenMP " \
348 "timing routines"
349#define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \
350 "undefined semantics within target regions, support for HSA ignores them"
351#define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP " \
352 "affinity feateres"
353
354/* Initialize hash map with simple OMP builtins. */
355
356static void
357hsa_init_simple_builtins ()
358{
359 if (omp_simple_builtins != NULL)
360 return;
361
362 omp_simple_builtins
363 = new hash_map <nofree_string_hash, omp_simple_builtin> ();
364
365 omp_simple_builtin omp_builtins[] =
366 {
367 omp_simple_builtin ("omp_get_initial_device", NULL, false,
368 new hsa_op_immed (GOMP_DEVICE_HOST,
369 (BrigType16_t) BRIG_TYPE_S32)),
370 omp_simple_builtin ("omp_is_initial_device", NULL, false,
371 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
372 omp_simple_builtin ("omp_get_dynamic", NULL, false,
373 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
374 omp_simple_builtin ("omp_set_dynamic", NULL, false, NULL),
375 omp_simple_builtin ("omp_init_lock", HSA_WARN_LOCK_ROUTINE, true),
376 omp_simple_builtin ("omp_init_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
377 true),
378 omp_simple_builtin ("omp_init_nest_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
379 true),
380 omp_simple_builtin ("omp_destroy_lock", HSA_WARN_LOCK_ROUTINE, true),
381 omp_simple_builtin ("omp_set_lock", HSA_WARN_LOCK_ROUTINE, true),
382 omp_simple_builtin ("omp_unset_lock", HSA_WARN_LOCK_ROUTINE, true),
383 omp_simple_builtin ("omp_test_lock", HSA_WARN_LOCK_ROUTINE, true),
384 omp_simple_builtin ("omp_get_wtime", HSA_WARN_TIMING_ROUTINE, true),
385 omp_simple_builtin ("omp_get_wtick", HSA_WARN_TIMING_ROUTINE, true),
386 omp_simple_builtin ("omp_target_alloc", HSA_WARN_MEMORY_ROUTINE, false,
387 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_U64)),
388 omp_simple_builtin ("omp_target_free", HSA_WARN_MEMORY_ROUTINE, false),
389 omp_simple_builtin ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE,
390 false,
391 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
392 omp_simple_builtin ("omp_target_memcpy", HSA_WARN_MEMORY_ROUTINE, false,
393 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
394 omp_simple_builtin ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE,
395 false,
396 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
397 omp_simple_builtin ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE,
398 false,
399 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
400 omp_simple_builtin ("omp_target_disassociate_ptr",
401 HSA_WARN_MEMORY_ROUTINE,
402 false,
403 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
404 omp_simple_builtin ("omp_set_max_active_levels",
405 "Support for HSA only allows only one active level, "
406 "call to omp_set_max_active_levels will be ignored "
407 "in the generated HSAIL",
408 false, NULL),
409 omp_simple_builtin ("omp_get_max_active_levels", NULL, false,
410 new hsa_op_immed (1, (BrigType16_t) BRIG_TYPE_S32)),
411 omp_simple_builtin ("omp_in_final", NULL, false,
412 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
413 omp_simple_builtin ("omp_get_proc_bind", HSA_WARN_AFFINITY, false,
414 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
415 omp_simple_builtin ("omp_get_num_places", HSA_WARN_AFFINITY, false,
416 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
417 omp_simple_builtin ("omp_get_place_num_procs", HSA_WARN_AFFINITY, false,
418 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
419 omp_simple_builtin ("omp_get_place_proc_ids", HSA_WARN_AFFINITY, false,
420 NULL),
421 omp_simple_builtin ("omp_get_place_num", HSA_WARN_AFFINITY, false,
422 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
423 omp_simple_builtin ("omp_get_partition_num_places", HSA_WARN_AFFINITY,
424 false,
425 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
426 omp_simple_builtin ("omp_get_partition_place_nums", HSA_WARN_AFFINITY,
427 false, NULL),
428 omp_simple_builtin ("omp_set_default_device",
429 "omp_set_default_device has undefined semantics "
430 "within target regions, support for HSA ignores it",
431 false, NULL),
432 omp_simple_builtin ("omp_get_default_device",
433 "omp_get_default_device has undefined semantics "
434 "within target regions, support for HSA ignores it",
435 false,
436 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
437 omp_simple_builtin ("omp_get_num_devices",
438 "omp_get_num_devices has undefined semantics "
439 "within target regions, support for HSA ignores it",
440 false,
441 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
442 omp_simple_builtin ("omp_get_num_procs", NULL, true, NULL),
443 omp_simple_builtin ("omp_get_cancellation", NULL, true, NULL),
444 omp_simple_builtin ("omp_set_nested", NULL, true, NULL),
445 omp_simple_builtin ("omp_get_nested", NULL, true, NULL),
446 omp_simple_builtin ("omp_set_schedule", NULL, true, NULL),
447 omp_simple_builtin ("omp_get_schedule", NULL, true, NULL),
448 omp_simple_builtin ("omp_get_thread_limit", NULL, true, NULL),
449 omp_simple_builtin ("omp_get_team_size", NULL, true, NULL),
450 omp_simple_builtin ("omp_get_ancestor_thread_num", NULL, true, NULL),
451 omp_simple_builtin ("omp_get_max_task_priority", NULL, true, NULL)
452 };
453
454 unsigned count = sizeof (omp_builtins) / sizeof (omp_simple_builtin);
455
456 for (unsigned i = 0; i < count; i++)
457 omp_simple_builtins->put (omp_builtins[i].m_name, omp_builtins[i]);
458}
459
460/* Allocate HSA structures that we need only while generating with this. */
461
462static void
463hsa_init_data_for_cfun ()
464{
465 hsa_init_compilation_unit_data ();
466 gcc_obstack_init (&hsa_obstack);
467}
468
469/* Deinitialize HSA subsystem and free all allocated memory. */
470
471static void
472hsa_deinit_data_for_cfun (void)
473{
474 basic_block bb;
475
476 FOR_ALL_BB_FN (bb, cfun)
477 if (bb->aux)
478 {
479 hsa_bb *hbb = hsa_bb_for_bb (bb);
480 hbb->~hsa_bb ();
481 bb->aux = NULL;
482 }
483
484 for (unsigned int i = 0; i < hsa_operands.length (); i++)
485 hsa_destroy_operand (hsa_operands[i]);
486
487 hsa_operands.release ();
488
489 for (unsigned i = 0; i < hsa_instructions.length (); i++)
490 hsa_destroy_insn (hsa_instructions[i]);
491
492 hsa_instructions.release ();
493
494 if (omp_simple_builtins != NULL)
495 {
496 delete omp_simple_builtins;
497 omp_simple_builtins = NULL;
498 }
499
500 obstack_free (&hsa_obstack, NULL);
501 delete hsa_cfun;
502}
503
504/* Return the type which holds addresses in the given SEGMENT. */
505
506static BrigType16_t
507hsa_get_segment_addr_type (BrigSegment8_t segment)
508{
509 switch (segment)
510 {
511 case BRIG_SEGMENT_NONE:
512 gcc_unreachable ();
513
514 case BRIG_SEGMENT_FLAT:
515 case BRIG_SEGMENT_GLOBAL:
516 case BRIG_SEGMENT_READONLY:
517 case BRIG_SEGMENT_KERNARG:
518 return hsa_machine_large_p () ? BRIG_TYPE_U64 : BRIG_TYPE_U32;
519
520 case BRIG_SEGMENT_GROUP:
521 case BRIG_SEGMENT_PRIVATE:
522 case BRIG_SEGMENT_SPILL:
523 case BRIG_SEGMENT_ARG:
524 return BRIG_TYPE_U32;
525 }
526 gcc_unreachable ();
527}
528
529/* Return integer brig type according to provided SIZE in bytes. If SIGN
530 is set to true, return signed integer type. */
531
532static BrigType16_t
533get_integer_type_by_bytes (unsigned size, bool sign)
534{
535 if (sign)
536 switch (size)
537 {
538 case 1:
539 return BRIG_TYPE_S8;
540 case 2:
541 return BRIG_TYPE_S16;
542 case 4:
543 return BRIG_TYPE_S32;
544 case 8:
545 return BRIG_TYPE_S64;
546 default:
547 break;
548 }
549 else
550 switch (size)
551 {
552 case 1:
553 return BRIG_TYPE_U8;
554 case 2:
555 return BRIG_TYPE_U16;
556 case 4:
557 return BRIG_TYPE_U32;
558 case 8:
559 return BRIG_TYPE_U64;
560 default:
561 break;
562 }
563
564 return 0;
565}
566
567/* If T points to an integral type smaller than 32 bits, change it to a 32bit
568 equivalent and return the result. Otherwise just return the result. */
569
570static BrigType16_t
571hsa_extend_inttype_to_32bit (BrigType16_t t)
572{
573 if (t == BRIG_TYPE_U8 || t == BRIG_TYPE_U16)
574 return BRIG_TYPE_U32;
575 else if (t == BRIG_TYPE_S8 || t == BRIG_TYPE_S16)
576 return BRIG_TYPE_S32;
577 return t;
578}
579
580/* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
581 are assumed to use flat addressing. If min32int is true, always expand
582 integer types to one that has at least 32 bits. */
583
584static BrigType16_t
585hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
586{
587 HOST_WIDE_INT bsize;
588 const_tree base;
589 BrigType16_t res = BRIG_TYPE_NONE;
590
591 gcc_checking_assert (TYPE_P (type));
592 gcc_checking_assert (!AGGREGATE_TYPE_P (type));
593 if (POINTER_TYPE_P (type))
594 return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
595
596 if (TREE_CODE (type) == VECTOR_TYPE)
597 base = TREE_TYPE (type);
598 else if (TREE_CODE (type) == COMPLEX_TYPE)
599 {
600 base = TREE_TYPE (type);
601 min32int = true;
602 }
603 else
604 base = type;
605
606 if (!tree_fits_uhwi_p (TYPE_SIZE (base)))
607 {
608 HSA_SORRY_ATV (EXPR_LOCATION (type),
609 "support for HSA does not implement huge or "
610 "variable-sized type %qT", type);
611 return res;
612 }
613
614 bsize = tree_to_uhwi (TYPE_SIZE (base));
615 unsigned byte_size = bsize / BITS_PER_UNIT;
616 if (INTEGRAL_TYPE_P (base))
617 res = get_integer_type_by_bytes (byte_size, !TYPE_UNSIGNED (base));
618 else if (SCALAR_FLOAT_TYPE_P (base))
619 {
620 switch (bsize)
621 {
622 case 16:
623 res = BRIG_TYPE_F16;
624 break;
625 case 32:
626 res = BRIG_TYPE_F32;
627 break;
628 case 64:
629 res = BRIG_TYPE_F64;
630 break;
631 default:
632 break;
633 }
634 }
635
636 if (res == BRIG_TYPE_NONE)
637 {
638 HSA_SORRY_ATV (EXPR_LOCATION (type),
639 "support for HSA does not implement type %qT", type);
640 return res;
641 }
642
643 if (TREE_CODE (type) == VECTOR_TYPE)
644 {
645 HOST_WIDE_INT tsize = tree_to_uhwi (TYPE_SIZE (type));
646
647 if (bsize == tsize)
648 {
649 HSA_SORRY_ATV (EXPR_LOCATION (type),
650 "support for HSA does not implement a vector type "
651 "where a type and unit size are equal: %qT", type);
652 return res;
653 }
654
655 switch (tsize)
656 {
657 case 32:
658 res |= BRIG_TYPE_PACK_32;
659 break;
660 case 64:
661 res |= BRIG_TYPE_PACK_64;
662 break;
663 case 128:
664 res |= BRIG_TYPE_PACK_128;
665 break;
666 default:
667 HSA_SORRY_ATV (EXPR_LOCATION (type),
668 "support for HSA does not implement type %qT", type);
669 }
670 }
671
672 if (min32int)
673 /* Registers/immediate operands can only be 32bit or more except for
674 f16. */
675 res = hsa_extend_inttype_to_32bit (res);
676
677 if (TREE_CODE (type) == COMPLEX_TYPE)
678 {
679 unsigned bsize = 2 * hsa_type_bit_size (res);
680 res = hsa_bittype_for_bitsize (bsize);
681 }
682
683 return res;
684}
685
686/* Returns the BRIG type we need to load/store entities of TYPE. */
687
688static BrigType16_t
689mem_type_for_type (BrigType16_t type)
690{
691 /* HSA has non-intuitive constraints on load/store types. If it's
692 a bit-type it _must_ be B128, if it's not a bit-type it must be
693 64bit max. So for loading entities of 128 bits (e.g. vectors)
694 we have to to B128, while for loading the rest we have to use the
695 input type (??? or maybe also flattened to a equally sized non-vector
696 unsigned type?). */
697 if ((type & BRIG_TYPE_PACK_MASK) == BRIG_TYPE_PACK_128)
698 return BRIG_TYPE_B128;
699 else if (hsa_btype_p (type) || hsa_type_packed_p (type))
700 {
701 unsigned bitsize = hsa_type_bit_size (type);
702 if (bitsize < 128)
703 return hsa_uint_for_bitsize (bitsize);
704 else
705 return hsa_bittype_for_bitsize (bitsize);
706 }
707 return type;
708}
709
710/* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some
711 kind of array will be generated, setting DIM appropriately. Otherwise, it
712 will be set to zero. */
713
714static BrigType16_t
715hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p = NULL,
716 bool min32int = false)
717{
718 gcc_checking_assert (TYPE_P (type));
719 if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type)))
720 {
721 HSA_SORRY_ATV (EXPR_LOCATION (type), "support for HSA does not "
722 "implement huge or variable-sized type %qT", type);
723 return BRIG_TYPE_NONE;
724 }
725
726 if (RECORD_OR_UNION_TYPE_P (type))
727 {
728 if (dim_p)
729 *dim_p = tree_to_uhwi (TYPE_SIZE_UNIT (type));
730 return BRIG_TYPE_U8 | BRIG_TYPE_ARRAY;
731 }
732
733 if (TREE_CODE (type) == ARRAY_TYPE)
734 {
735 /* We try to be nice and use the real base-type when this is an array of
736 scalars and only resort to an array of bytes if the type is more
737 complex. */
738
739 unsigned HOST_WIDE_INT dim = 1;
740
741 while (TREE_CODE (type) == ARRAY_TYPE)
742 {
743 tree domain = TYPE_DOMAIN (type);
744 if (!TYPE_MIN_VALUE (domain)
745 || !TYPE_MAX_VALUE (domain)
746 || !tree_fits_shwi_p (TYPE_MIN_VALUE (domain))
747 || !tree_fits_shwi_p (TYPE_MAX_VALUE (domain)))
748 {
749 HSA_SORRY_ATV (EXPR_LOCATION (type),
750 "support for HSA does not implement array "
751 "%qT with unknown bounds", type);
752 return BRIG_TYPE_NONE;
753 }
754 HOST_WIDE_INT min = tree_to_shwi (TYPE_MIN_VALUE (domain));
755 HOST_WIDE_INT max = tree_to_shwi (TYPE_MAX_VALUE (domain));
756 dim = dim * (unsigned HOST_WIDE_INT) (max - min + 1);
757 type = TREE_TYPE (type);
758 }
759
760 BrigType16_t res;
761 if (RECORD_OR_UNION_TYPE_P (type))
762 {
763 dim = dim * tree_to_uhwi (TYPE_SIZE_UNIT (type));
764 res = BRIG_TYPE_U8;
765 }
766 else
767 res = hsa_type_for_scalar_tree_type (type, false);
768
769 if (dim_p)
770 *dim_p = dim;
771 return res | BRIG_TYPE_ARRAY;
772 }
773
774 /* Scalar case: */
775 if (dim_p)
776 *dim_p = 0;
777
778 return hsa_type_for_scalar_tree_type (type, min32int);
779}
780
781/* Returns true if converting from STYPE into DTYPE needs the _CVT
782 opcode. If false a normal _MOV is enough. */
783
784static bool
785hsa_needs_cvt (BrigType16_t dtype, BrigType16_t stype)
786{
787 if (hsa_btype_p (dtype))
788 return false;
789
790 /* float <-> int conversions are real converts. */
791 if (hsa_type_float_p (dtype) != hsa_type_float_p (stype))
792 return true;
793 /* When both types have different size, then we need CVT as well. */
794 if (hsa_type_bit_size (dtype) != hsa_type_bit_size (stype))
795 return true;
796 return false;
797}
798
799/* Return declaration name if it exists or create one from UID if it does not.
800 If DECL is a local variable, make UID part of its name. */
801
802const char *
803hsa_get_declaration_name (tree decl)
804{
805 if (!DECL_NAME (decl))
806 {
807 char buf[64];
808 snprintf (buf, 64, "__hsa_anon_%u", DECL_UID (decl));
809 size_t len = strlen (buf);
810 char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1);
811 memcpy (copy, buf, len + 1);
812 return copy;
813 }
814
815 tree name_tree;
816 if (TREE_CODE (decl) == FUNCTION_DECL
817 || (TREE_CODE (decl) == VAR_DECL && is_global_var (decl)))
818 name_tree = DECL_ASSEMBLER_NAME (decl);
819 else
820 name_tree = DECL_NAME (decl);
821
822 const char *name = IDENTIFIER_POINTER (name_tree);
823 /* User-defined assembly names have prepended asterisk symbol. */
824 if (name[0] == '*')
825 name++;
826
827 if ((TREE_CODE (decl) == VAR_DECL)
828 && decl_function_context (decl))
829 {
830 size_t len = strlen (name);
831 char *buf = (char *) alloca (len + 32);
832 snprintf (buf, len + 32, "%s_%u", name, DECL_UID (decl));
833 len = strlen (buf);
834 char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1);
835 memcpy (copy, buf, len + 1);
836 return copy;
837 }
838 else
839 return name;
840}
841
842/* Lookup or create the associated hsa_symbol structure with a given VAR_DECL
843 or lookup the hsa_structure corresponding to a PARM_DECL. */
844
845static hsa_symbol *
846get_symbol_for_decl (tree decl)
847{
848 hsa_symbol **slot;
849 hsa_symbol dummy (BRIG_TYPE_NONE, BRIG_SEGMENT_NONE, BRIG_LINKAGE_NONE);
850
851 gcc_assert (TREE_CODE (decl) == PARM_DECL
852 || TREE_CODE (decl) == RESULT_DECL
853 || TREE_CODE (decl) == VAR_DECL
854 || TREE_CODE (decl) == CONST_DECL);
855
856 dummy.m_decl = decl;
857
858 bool is_in_global_vars = ((TREE_CODE (decl) == VAR_DECL)
859 && !decl_function_context (decl));
860
861 if (is_in_global_vars)
862 slot = hsa_global_variable_symbols->find_slot (&dummy, INSERT);
863 else
864 slot = hsa_cfun->m_local_symbols->find_slot (&dummy, INSERT);
865
866 gcc_checking_assert (slot);
867 if (*slot)
868 {
869 hsa_symbol *sym = (*slot);
870
871 /* If the symbol is problematic, mark current function also as
872 problematic. */
873 if (sym->m_seen_error)
874 hsa_fail_cfun ();
875
876 /* PR hsa/70234: If a global variable was marked to be emitted,
877 but HSAIL generation of a function using the variable fails,
878 we should retry to emit the variable in context of a different
879 function.
880
881 Iterate elements whether a symbol is already in m_global_symbols
882 of not. */
883 if (is_in_global_vars && !sym->m_emitted_to_brig)
884 {
885 for (unsigned i = 0; i < hsa_cfun->m_global_symbols.length (); i++)
886 if (hsa_cfun->m_global_symbols[i] == sym)
887 return *slot;
888 hsa_cfun->m_global_symbols.safe_push (sym);
889 }
890
891 return *slot;
892 }
893 else
894 {
895 hsa_symbol *sym;
896 /* PARM_DECLs and RESULT_DECL should be already in m_local_symbols. */
897 gcc_assert (TREE_CODE (decl) == VAR_DECL
898 || TREE_CODE (decl) == CONST_DECL);
899 BrigAlignment8_t align = hsa_object_alignment (decl);
900
901 if (is_in_global_vars)
902 {
903 gcc_checking_assert (TREE_CODE (decl) != CONST_DECL);
904 sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_GLOBAL,
905 BRIG_LINKAGE_PROGRAM, true,
906 BRIG_ALLOCATION_PROGRAM, align);
907 hsa_cfun->m_global_symbols.safe_push (sym);
908 sym->fillup_for_decl (decl);
909 if (sym->m_align > align)
910 {
911 sym->m_seen_error = true;
912 HSA_SORRY_ATV (EXPR_LOCATION (decl),
913 "HSA specification requires that %E is at least "
914 "naturally aligned", decl);
915 }
916 }
917 else
918 {
919 /* As generation of efficient memory copy instructions relies
920 on alignment greater or equal to 8 bytes,
921 we need to increase alignment of all aggregate types.. */
922 if (AGGREGATE_TYPE_P (TREE_TYPE (decl)))
923 align = MAX ((BrigAlignment8_t) BRIG_ALIGNMENT_8, align);
924
925 BrigAllocation allocation = BRIG_ALLOCATION_AUTOMATIC;
926 BrigSegment8_t segment;
927 if (TREE_CODE (decl) == CONST_DECL)
928 {
929 segment = BRIG_SEGMENT_READONLY;
930 allocation = BRIG_ALLOCATION_AGENT;
931 }
932 else if (lookup_attribute ("hsa_group_segment",
933 DECL_ATTRIBUTES (decl)))
934 segment = BRIG_SEGMENT_GROUP;
935 else if (TREE_STATIC (decl)
936 || lookup_attribute ("hsa_global_segment",
937 DECL_ATTRIBUTES (decl)))
938 segment = BRIG_SEGMENT_GLOBAL;
939 else
940 segment = BRIG_SEGMENT_PRIVATE;
941
942 sym = new hsa_symbol (BRIG_TYPE_NONE, segment, BRIG_LINKAGE_FUNCTION,
943 false, allocation, align);
944 sym->fillup_for_decl (decl);
945 hsa_cfun->m_private_variables.safe_push (sym);
946 }
947
948 sym->m_name = hsa_get_declaration_name (decl);
949 *slot = sym;
950 return sym;
951 }
952}
953
954/* For a given HSA function declaration, return a host
955 function declaration. */
956
957tree
958hsa_get_host_function (tree decl)
959{
960 hsa_function_summary *s
961 = hsa_summaries->get (cgraph_node::get_create (decl));
962 gcc_assert (s->m_kind != HSA_NONE);
963 gcc_assert (s->m_gpu_implementation_p);
964
965 return s->m_bound_function ? s->m_bound_function->decl : NULL;
966}
967
968/* Return true if function DECL has a host equivalent function. */
969
970static char *
971get_brig_function_name (tree decl)
972{
973 tree d = decl;
974
975 hsa_function_summary *s = hsa_summaries->get (cgraph_node::get_create (d));
976 if (s->m_kind != HSA_NONE
977 && s->m_gpu_implementation_p
978 && s->m_bound_function)
979 d = s->m_bound_function->decl;
980
981 /* IPA split can create a function that has no host equivalent. */
982 if (d == NULL)
983 d = decl;
984
985 char *name = xstrdup (hsa_get_declaration_name (d));
986 hsa_sanitize_name (name);
987
988 return name;
989}
990
991/* Create a spill symbol of type TYPE. */
992
993hsa_symbol *
994hsa_get_spill_symbol (BrigType16_t type)
995{
996 hsa_symbol *sym = new hsa_symbol (type, BRIG_SEGMENT_SPILL,
997 BRIG_LINKAGE_FUNCTION);
998 hsa_cfun->m_spill_symbols.safe_push (sym);
999 return sym;
1000}
1001
1002/* Create a symbol for a read-only string constant. */
1003hsa_symbol *
1004hsa_get_string_cst_symbol (tree string_cst)
1005{
1006 gcc_checking_assert (TREE_CODE (string_cst) == STRING_CST);
1007
1008 hsa_symbol **slot = hsa_cfun->m_string_constants_map.get (string_cst);
1009 if (slot)
1010 return *slot;
1011
1012 hsa_op_immed *cst = new hsa_op_immed (string_cst);
1013 hsa_symbol *sym = new hsa_symbol (cst->m_type, BRIG_SEGMENT_GLOBAL,
1014 BRIG_LINKAGE_MODULE, true,
1015 BRIG_ALLOCATION_AGENT);
1016 sym->m_cst_value = cst;
1017 sym->m_dim = TREE_STRING_LENGTH (string_cst);
1018 sym->m_name_number = hsa_cfun->m_global_symbols.length ();
1019
1020 hsa_cfun->m_global_symbols.safe_push (sym);
1021 hsa_cfun->m_string_constants_map.put (string_cst, sym);
1022 return sym;
1023}
1024
1025/* Make the type of a MOV instruction larger if mandated by HSAIL rules. */
1026
1027static void
1028hsa_fixup_mov_insn_type (hsa_insn_basic *insn)
1029{
1030 insn->m_type = hsa_extend_inttype_to_32bit (insn->m_type);
1031 if (insn->m_type == BRIG_TYPE_B8 || insn->m_type == BRIG_TYPE_B16)
1032 insn->m_type = BRIG_TYPE_B32;
1033}
1034
1035/* Constructor of the ancestor of all operands. K is BRIG kind that identified
1036 what the operator is. */
1037
1038hsa_op_base::hsa_op_base (BrigKind16_t k)
1039 : m_next (NULL), m_brig_op_offset (0), m_kind (k)
1040{
1041 hsa_operands.safe_push (this);
1042}
1043
1044/* Constructor of ancestor of all operands which have a type. K is BRIG kind
1045 that identified what the operator is. T is the type of the operator. */
1046
1047hsa_op_with_type::hsa_op_with_type (BrigKind16_t k, BrigType16_t t)
1048 : hsa_op_base (k), m_type (t)
1049{
1050}
1051
1052hsa_op_with_type *
1053hsa_op_with_type::get_in_type (BrigType16_t dtype, hsa_bb *hbb)
1054{
1055 if (m_type == dtype)
1056 return this;
1057
1058 hsa_op_reg *dest;
1059
1060 if (hsa_needs_cvt (dtype, m_type))
1061 {
1062 dest = new hsa_op_reg (dtype);
1063 hbb->append_insn (new hsa_insn_cvt (dest, this));
1064 }
1065 else if (is_a <hsa_op_reg *> (this))
1066 {
1067 /* In the end, HSA registers do not really have types, only sizes, so if
1068 the sizes match, we can use the register directly. */
1069 gcc_checking_assert (hsa_type_bit_size (dtype)
1070 == hsa_type_bit_size (m_type));
1071 return this;
1072 }
1073 else
1074 {
1075 dest = new hsa_op_reg (m_type);
1076
1077 hsa_insn_basic *mov = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
1078 dest->m_type, dest, this);
1079 hsa_fixup_mov_insn_type (mov);
1080 hbb->append_insn (mov);
1081 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
1082 type of the operand must be same as type of the instruction. */
1083 dest->m_type = dtype;
1084 }
1085
1086 return dest;
1087}
1088
1089/* If this operand has integer type smaller than 32 bits, extend it to 32 bits,
1090 adding instructions to HBB if needed. */
1091
1092hsa_op_with_type *
1093hsa_op_with_type::extend_int_to_32bit (hsa_bb *hbb)
1094{
1095 if (m_type == BRIG_TYPE_U8 || m_type == BRIG_TYPE_U16)
1096 return get_in_type (BRIG_TYPE_U32, hbb);
1097 else if (m_type == BRIG_TYPE_S8 || m_type == BRIG_TYPE_S16)
1098 return get_in_type (BRIG_TYPE_S32, hbb);
1099 else
1100 return this;
1101}
1102
1103/* Constructor of class representing HSA immediate values. TREE_VAL is the
1104 tree representation of the immediate value. If min32int is true,
1105 always expand integer types to one that has at least 32 bits. */
1106
1107hsa_op_immed::hsa_op_immed (tree tree_val, bool min32int)
1108 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES,
1109 hsa_type_for_tree_type (TREE_TYPE (tree_val), NULL,
1110 min32int))
1111{
1112 if (hsa_seen_error ())
1113 return;
1114
1115 gcc_checking_assert ((is_gimple_min_invariant (tree_val)
1116 && (!POINTER_TYPE_P (TREE_TYPE (tree_val))
1117 || TREE_CODE (tree_val) == INTEGER_CST))
1118 || TREE_CODE (tree_val) == CONSTRUCTOR);
1119 m_tree_value = tree_val;
1120
1121 /* Verify that all elements of a constructor are constants. */
1122 if (TREE_CODE (m_tree_value) == CONSTRUCTOR)
1123 for (unsigned i = 0; i < CONSTRUCTOR_NELTS (m_tree_value); i++)
1124 {
1125 tree v = CONSTRUCTOR_ELT (m_tree_value, i)->value;
1126 if (!CONSTANT_CLASS_P (v))
1127 {
1128 HSA_SORRY_AT (EXPR_LOCATION (tree_val),
1129 "HSA ctor should have only constants");
1130 return;
1131 }
1132 }
1133}
1134
1135/* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
1136 integer representation of the immediate value. TYPE is BRIG type. */
1137
1138hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value, BrigType16_t type)
1139 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES, type),
1140 m_tree_value (NULL)
1141{
1142 gcc_assert (hsa_type_integer_p (type));
1143 m_int_value = integer_value;
1144}
1145
1146hsa_op_immed::hsa_op_immed ()
1147 : hsa_op_with_type (BRIG_KIND_NONE, BRIG_TYPE_NONE)
1148{
1149}
1150
1151/* New operator to allocate immediate operands from obstack. */
1152
1153void *
1154hsa_op_immed::operator new (size_t size)
1155{
1156 return obstack_alloc (&hsa_obstack, size);
1157}
1158
1159/* Destructor. */
1160
1161hsa_op_immed::~hsa_op_immed ()
1162{
1163}
1164
1165/* Change type of the immediate value to T. */
1166
1167void
1168hsa_op_immed::set_type (BrigType16_t t)
1169{
1170 m_type = t;
1171}
1172
1173/* Constructor of class representing HSA registers and pseudo-registers. T is
1174 the BRIG type of the new register. */
1175
1176hsa_op_reg::hsa_op_reg (BrigType16_t t)
1177 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER, t), m_gimple_ssa (NULL_TREE),
1178 m_def_insn (NULL), m_spill_sym (NULL), m_order (hsa_cfun->m_reg_count++),
1179 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
1180{
1181}
1182
1183/* New operator to allocate a register from obstack. */
1184
1185void *
1186hsa_op_reg::operator new (size_t size)
1187{
1188 return obstack_alloc (&hsa_obstack, size);
1189}
1190
1191/* Verify register operand. */
1192
1193void
1194hsa_op_reg::verify_ssa ()
1195{
1196 /* Verify that each HSA register has a definition assigned.
1197 Exceptions are VAR_DECL and PARM_DECL that are a default
1198 definition. */
1199 gcc_checking_assert (m_def_insn
1200 || (m_gimple_ssa != NULL
1201 && (!SSA_NAME_VAR (m_gimple_ssa)
1202 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa))
1203 != PARM_DECL))
1204 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa)));
1205
1206 /* Verify that every use of the register is really present
1207 in an instruction. */
1208 for (unsigned i = 0; i < m_uses.length (); i++)
1209 {
1210 hsa_insn_basic *use = m_uses[i];
1211
1212 bool is_visited = false;
1213 for (unsigned j = 0; j < use->operand_count (); j++)
1214 {
1215 hsa_op_base *u = use->get_op (j);
1216 hsa_op_address *addr; addr = dyn_cast <hsa_op_address *> (u);
1217 if (addr && addr->m_reg)
1218 u = addr->m_reg;
1219
1220 if (u == this)
1221 {
1222 bool r = !addr && use->op_output_p (j);
1223
1224 if (r)
1225 {
1226 error ("HSA SSA name defined by instruction that is supposed "
1227 "to be using it");
1228 debug_hsa_operand (this);
1229 debug_hsa_insn (use);
1230 internal_error ("HSA SSA verification failed");
1231 }
1232
1233 is_visited = true;
1234 }
1235 }
1236
1237 if (!is_visited)
1238 {
1239 error ("HSA SSA name not among operands of instruction that is "
1240 "supposed to use it");
1241 debug_hsa_operand (this);
1242 debug_hsa_insn (use);
1243 internal_error ("HSA SSA verification failed");
1244 }
1245 }
1246}
1247
1248hsa_op_address::hsa_op_address (hsa_symbol *sym, hsa_op_reg *r,
1249 HOST_WIDE_INT offset)
1250 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (r),
1251 m_imm_offset (offset)
1252{
1253}
1254
1255hsa_op_address::hsa_op_address (hsa_symbol *sym, HOST_WIDE_INT offset)
1256 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (NULL),
1257 m_imm_offset (offset)
1258{
1259}
1260
1261hsa_op_address::hsa_op_address (hsa_op_reg *r, HOST_WIDE_INT offset)
1262 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (NULL), m_reg (r),
1263 m_imm_offset (offset)
1264{
1265}
1266
1267/* New operator to allocate address operands from obstack. */
1268
1269void *
1270hsa_op_address::operator new (size_t size)
1271{
1272 return obstack_alloc (&hsa_obstack, size);
1273}
1274
1275/* Constructor of an operand referring to HSAIL code. */
1276
1277hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF),
1278 m_directive_offset (0)
1279{
1280}
1281
1282/* Constructor of an operand representing a code list. Set it up so that it
1283 can contain ELEMENTS number of elements. */
1284
1285hsa_op_code_list::hsa_op_code_list (unsigned elements)
1286 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST)
1287{
1288 m_offsets.create (1);
1289 m_offsets.safe_grow_cleared (elements);
1290}
1291
1292/* New operator to allocate code list operands from obstack. */
1293
1294void *
1295hsa_op_code_list::operator new (size_t size)
1296{
1297 return obstack_alloc (&hsa_obstack, size);
1298}
1299
1300/* Constructor of an operand representing an operand list.
1301 Set it up so that it can contain ELEMENTS number of elements. */
1302
1303hsa_op_operand_list::hsa_op_operand_list (unsigned elements)
1304 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST)
1305{
1306 m_offsets.create (elements);
1307 m_offsets.safe_grow (elements);
1308}
1309
1310/* New operator to allocate operand list operands from obstack. */
1311
1312void *
1313hsa_op_operand_list::operator new (size_t size)
1314{
1315 return obstack_alloc (&hsa_obstack, size);
1316}
1317
1318hsa_op_operand_list::~hsa_op_operand_list ()
1319{
1320 m_offsets.release ();
1321}
1322
1323
1324hsa_op_reg *
1325hsa_function_representation::reg_for_gimple_ssa (tree ssa)
1326{
1327 hsa_op_reg *hreg;
1328
1329 gcc_checking_assert (TREE_CODE (ssa) == SSA_NAME);
1330 if (m_ssa_map[SSA_NAME_VERSION (ssa)])
1331 return m_ssa_map[SSA_NAME_VERSION (ssa)];
1332
1333 hreg = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa),
1334 false));
1335 hreg->m_gimple_ssa = ssa;
1336 m_ssa_map[SSA_NAME_VERSION (ssa)] = hreg;
1337
1338 return hreg;
1339}
1340
1341void
1342hsa_op_reg::set_definition (hsa_insn_basic *insn)
1343{
1344 if (hsa_cfun->m_in_ssa)
1345 {
1346 gcc_checking_assert (!m_def_insn);
1347 m_def_insn = insn;
1348 }
1349 else
1350 m_def_insn = NULL;
1351}
1352
1353/* Constructor of the class which is the bases of all instructions and directly
1354 represents the most basic ones. NOPS is the number of operands that the
1355 operand vector will contain (and which will be cleared). OP is the opcode
1356 of the instruction. This constructor does not set type. */
1357
1358hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc)
1359 : m_prev (NULL),
1360 m_next (NULL), m_bb (NULL), m_opcode (opc), m_number (0),
1361 m_type (BRIG_TYPE_NONE), m_brig_offset (0)
1362{
1363 if (nops > 0)
1364 m_operands.safe_grow_cleared (nops);
1365
1366 hsa_instructions.safe_push (this);
1367}
1368
1369/* Make OP the operand number INDEX of operands of this instruction. If OP is a
1370 register or an address containing a register, then either set the definition
1371 of the register to this instruction if it an output operand or add this
1372 instruction to the uses if it is an input one. */
1373
1374void
1375hsa_insn_basic::set_op (int index, hsa_op_base *op)
1376{
1377 /* Each address operand is always use. */
1378 hsa_op_address *addr = dyn_cast <hsa_op_address *> (op);
1379 if (addr && addr->m_reg)
1380 addr->m_reg->m_uses.safe_push (this);
1381 else
1382 {
1383 hsa_op_reg *reg = dyn_cast <hsa_op_reg *> (op);
1384 if (reg)
1385 {
1386 if (op_output_p (index))
1387 reg->set_definition (this);
1388 else
1389 reg->m_uses.safe_push (this);
1390 }
1391 }
1392
1393 m_operands[index] = op;
1394}
1395
1396/* Get INDEX-th operand of the instruction. */
1397
1398hsa_op_base *
1399hsa_insn_basic::get_op (int index)
1400{
1401 return m_operands[index];
1402}
1403
1404/* Get address of INDEX-th operand of the instruction. */
1405
1406hsa_op_base **
1407hsa_insn_basic::get_op_addr (int index)
1408{
1409 return &m_operands[index];
1410}
1411
1412/* Get number of operands of the instruction. */
1413unsigned int
1414hsa_insn_basic::operand_count ()
1415{
1416 return m_operands.length ();
1417}
1418
1419/* Constructor of the class which is the bases of all instructions and directly
1420 represents the most basic ones. NOPS is the number of operands that the
1421 operand vector will contain (and which will be cleared). OPC is the opcode
1422 of the instruction, T is the type of the instruction. */
1423
1424hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc, BrigType16_t t,
1425 hsa_op_base *arg0, hsa_op_base *arg1,
1426 hsa_op_base *arg2, hsa_op_base *arg3)
1427 : m_prev (NULL), m_next (NULL), m_bb (NULL), m_opcode (opc),m_number (0),
1428 m_type (t), m_brig_offset (0)
1429{
1430 if (nops > 0)
1431 m_operands.safe_grow_cleared (nops);
1432
1433 if (arg0 != NULL)
1434 {
1435 gcc_checking_assert (nops >= 1);
1436 set_op (0, arg0);
1437 }
1438
1439 if (arg1 != NULL)
1440 {
1441 gcc_checking_assert (nops >= 2);
1442 set_op (1, arg1);
1443 }
1444
1445 if (arg2 != NULL)
1446 {
1447 gcc_checking_assert (nops >= 3);
1448 set_op (2, arg2);
1449 }
1450
1451 if (arg3 != NULL)
1452 {
1453 gcc_checking_assert (nops >= 4);
1454 set_op (3, arg3);
1455 }
1456
1457 hsa_instructions.safe_push (this);
1458}
1459
1460/* New operator to allocate basic instruction from obstack. */
1461
1462void *
1463hsa_insn_basic::operator new (size_t size)
1464{
1465 return obstack_alloc (&hsa_obstack, size);
1466}
1467
1468/* Verify the instruction. */
1469
1470void
1471hsa_insn_basic::verify ()
1472{
1473 hsa_op_address *addr;
1474 hsa_op_reg *reg;
1475
1476 /* Iterate all register operands and verify that the instruction
1477 is set in uses of the register. */
1478 for (unsigned i = 0; i < operand_count (); i++)
1479 {
1480 hsa_op_base *use = get_op (i);
1481
1482 if ((addr = dyn_cast <hsa_op_address *> (use)) && addr->m_reg)
1483 {
1484 gcc_assert (addr->m_reg->m_def_insn != this);
1485 use = addr->m_reg;
1486 }
1487
1488 if ((reg = dyn_cast <hsa_op_reg *> (use)) && !op_output_p (i))
1489 {
1490 unsigned j;
1491 for (j = 0; j < reg->m_uses.length (); j++)
1492 {
1493 if (reg->m_uses[j] == this)
1494 break;
1495 }
1496
1497 if (j == reg->m_uses.length ())
1498 {
1499 error ("HSA instruction uses a register but is not among "
1500 "recorded register uses");
1501 debug_hsa_operand (reg);
1502 debug_hsa_insn (this);
1503 internal_error ("HSA instruction verification failed");
1504 }
1505 }
1506 }
1507}
1508
1509/* Constructor of an instruction representing a PHI node. NOPS is the number
1510 of operands (equal to the number of predecessors). */
1511
1512hsa_insn_phi::hsa_insn_phi (unsigned nops, hsa_op_reg *dst)
1513 : hsa_insn_basic (nops, HSA_OPCODE_PHI), m_dest (dst)
1514{
1515 dst->set_definition (this);
1516}
1517
1518/* Constructor of class representing instructions for control flow and
1519 sychronization, */
1520
1521hsa_insn_br::hsa_insn_br (unsigned nops, int opc, BrigType16_t t,
1522 BrigWidth8_t width, hsa_op_base *arg0,
1523 hsa_op_base *arg1, hsa_op_base *arg2,
1524 hsa_op_base *arg3)
1525 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
1526 m_width (width)
1527{
1528}
1529
1530/* Constructor of class representing instruction for conditional jump, CTRL is
1531 the control register determining whether the jump will be carried out, the
1532 new instruction is automatically added to its uses list. */
1533
1534hsa_insn_cbr::hsa_insn_cbr (hsa_op_reg *ctrl)
1535 : hsa_insn_br (1, BRIG_OPCODE_CBR, BRIG_TYPE_B1, BRIG_WIDTH_1, ctrl)
1536{
1537}
1538
1539/* Constructor of class representing instruction for switch jump, CTRL is
1540 the index register. */
1541
1542hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg *index, unsigned jump_count)
1543 : hsa_insn_basic (1, BRIG_OPCODE_SBR, BRIG_TYPE_B1, index),
1544 m_width (BRIG_WIDTH_1), m_jump_table (vNULL),
1545 m_label_code_list (new hsa_op_code_list (jump_count))
1546{
1547}
1548
1549/* Replace all occurrences of OLD_BB with NEW_BB in the statements
1550 jump table. */
1551
1552void
1553hsa_insn_sbr::replace_all_labels (basic_block old_bb, basic_block new_bb)
1554{
1555 for (unsigned i = 0; i < m_jump_table.length (); i++)
1556 if (m_jump_table[i] == old_bb)
1557 m_jump_table[i] = new_bb;
1558}
1559
1560hsa_insn_sbr::~hsa_insn_sbr ()
1561{
1562 m_jump_table.release ();
1563}
1564
1565/* Constructor of comparison instruction. CMP is the comparison operation and T
1566 is the result type. */
1567
1568hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp, BrigType16_t t,
1569 hsa_op_base *arg0, hsa_op_base *arg1,
1570 hsa_op_base *arg2)
1571 : hsa_insn_basic (3 , BRIG_OPCODE_CMP, t, arg0, arg1, arg2), m_compare (cmp)
1572{
1573}
1574
1575/* Constructor of classes representing memory accesses. OPC is the opcode (must
1576 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
1577 operands are provided as ARG0 and ARG1. */
1578
1579hsa_insn_mem::hsa_insn_mem (int opc, BrigType16_t t, hsa_op_base *arg0,
1580 hsa_op_base *arg1)
1581 : hsa_insn_basic (2, opc, t, arg0, arg1),
1582 m_align (hsa_natural_alignment (t)), m_equiv_class (0)
1583{
1584 gcc_checking_assert (opc == BRIG_OPCODE_LD || opc == BRIG_OPCODE_ST);
1585}
1586
1587/* Constructor for descendants allowing different opcodes and number of
1588 operands, it passes its arguments directly to hsa_insn_basic
1589 constructor. The instruction operands are provided as ARG[0-3]. */
1590
1591
1592hsa_insn_mem::hsa_insn_mem (unsigned nops, int opc, BrigType16_t t,
1593 hsa_op_base *arg0, hsa_op_base *arg1,
1594 hsa_op_base *arg2, hsa_op_base *arg3)
1595 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
1596 m_align (hsa_natural_alignment (t)), m_equiv_class (0)
1597{
1598}
1599
1600/* Constructor of class representing atomic instructions. OPC is the principal
1601 opcode, AOP is the specific atomic operation opcode. T is the type of the
1602 instruction. The instruction operands are provided as ARG[0-3]. */
1603
1604hsa_insn_atomic::hsa_insn_atomic (int nops, int opc,
1605 enum BrigAtomicOperation aop,
1606 BrigType16_t t, BrigMemoryOrder memorder,
1607 hsa_op_base *arg0,
1608 hsa_op_base *arg1, hsa_op_base *arg2,
1609 hsa_op_base *arg3)
1610 : hsa_insn_mem (nops, opc, t, arg0, arg1, arg2, arg3), m_atomicop (aop),
1611 m_memoryorder (memorder),
1612 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM)
1613{
1614 gcc_checking_assert (opc == BRIG_OPCODE_ATOMICNORET ||
1615 opc == BRIG_OPCODE_ATOMIC ||
1616 opc == BRIG_OPCODE_SIGNAL ||
1617 opc == BRIG_OPCODE_SIGNALNORET);
1618}
1619
1620/* Constructor of class representing signal instructions. OPC is the prinicpal
1621 opcode, SOP is the specific signal operation opcode. T is the type of the
1622 instruction. The instruction operands are provided as ARG[0-3]. */
1623
1624hsa_insn_signal::hsa_insn_signal (int nops, int opc,
1625 enum BrigAtomicOperation sop,
1626 BrigType16_t t, BrigMemoryOrder memorder,
1627 hsa_op_base *arg0, hsa_op_base *arg1,
1628 hsa_op_base *arg2, hsa_op_base *arg3)
1629 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
1630 m_memory_order (memorder), m_signalop (sop)
1631{
1632}
1633
1634/* Constructor of class representing segment conversion instructions. OPC is
1635 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
1636 and SRCT are destination and source types respectively, SEG is the segment
1637 we are converting to or from. The instruction operands are
1638 provided as ARG0 and ARG1. */
1639
1640hsa_insn_seg::hsa_insn_seg (int opc, BrigType16_t dest, BrigType16_t srct,
1641 BrigSegment8_t seg, hsa_op_base *arg0,
1642 hsa_op_base *arg1)
1643 : hsa_insn_basic (2, opc, dest, arg0, arg1), m_src_type (srct),
1644 m_segment (seg)
1645{
1646 gcc_checking_assert (opc == BRIG_OPCODE_STOF || opc == BRIG_OPCODE_FTOS);
1647}
1648
1649/* Constructor of class representing a call instruction. CALLEE is the tree
1650 representation of the function being called. */
1651
1652hsa_insn_call::hsa_insn_call (tree callee)
1653 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (callee),
1654 m_output_arg (NULL), m_args_code_list (NULL), m_result_code_list (NULL)
1655{
1656}
1657
1658hsa_insn_call::hsa_insn_call (hsa_internal_fn *fn)
1659 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (NULL),
1660 m_called_internal_fn (fn), m_output_arg (NULL), m_args_code_list (NULL),
1661 m_result_code_list (NULL)
1662{
1663}
1664
1665hsa_insn_call::~hsa_insn_call ()
1666{
1667 for (unsigned i = 0; i < m_input_args.length (); i++)
1668 delete m_input_args[i];
1669
1670 delete m_output_arg;
1671
1672 m_input_args.release ();
1673 m_input_arg_insns.release ();
1674}
1675
1676/* Constructor of class representing the argument block required to invoke
1677 a call in HSAIL. */
1678hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind,
1679 hsa_insn_call * call)
1680 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK), m_kind (brig_kind),
1681 m_call_insn (call)
1682{
1683}
1684
1685hsa_insn_comment::hsa_insn_comment (const char *s)
1686 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT)
1687{
1688 unsigned l = strlen (s);
1689
1690 /* Append '// ' to the string. */
1691 char *buf = XNEWVEC (char, l + 4);
1692 sprintf (buf, "// %s", s);
1693 m_comment = buf;
1694}
1695
1696hsa_insn_comment::~hsa_insn_comment ()
1697{
1698 gcc_checking_assert (m_comment);
1699 free (m_comment);
1700 m_comment = NULL;
1701}
1702
1703/* Constructor of class representing the queue instruction in HSAIL. */
1704
1705hsa_insn_queue::hsa_insn_queue (int nops, int opcode, BrigSegment segment,
1706 BrigMemoryOrder memory_order,
1707 hsa_op_base *arg0, hsa_op_base *arg1,
1708 hsa_op_base *arg2, hsa_op_base *arg3)
1709 : hsa_insn_basic (nops, opcode, BRIG_TYPE_U64, arg0, arg1, arg2, arg3),
1710 m_segment (segment), m_memory_order (memory_order)
1711{
1712}
1713
1714/* Constructor of class representing the source type instruction in HSAIL. */
1715
1716hsa_insn_srctype::hsa_insn_srctype (int nops, BrigOpcode opcode,
1717 BrigType16_t destt, BrigType16_t srct,
1718 hsa_op_base *arg0, hsa_op_base *arg1,
1719 hsa_op_base *arg2 = NULL)
1720 : hsa_insn_basic (nops, opcode, destt, arg0, arg1, arg2),
1721 m_source_type (srct)
1722{}
1723
1724/* Constructor of class representing the packed instruction in HSAIL. */
1725
1726hsa_insn_packed::hsa_insn_packed (int nops, BrigOpcode opcode,
1727 BrigType16_t destt, BrigType16_t srct,
1728 hsa_op_base *arg0, hsa_op_base *arg1,
1729 hsa_op_base *arg2)
1730 : hsa_insn_srctype (nops, opcode, destt, srct, arg0, arg1, arg2)
1731{
1732 m_operand_list = new hsa_op_operand_list (nops - 1);
1733}
1734
1735/* Constructor of class representing the convert instruction in HSAIL. */
1736
1737hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type *dest, hsa_op_with_type *src)
1738 : hsa_insn_basic (2, BRIG_OPCODE_CVT, dest->m_type, dest, src)
1739{
1740}
1741
1742/* Constructor of class representing the alloca in HSAIL. */
1743
1744hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type *dest,
1745 hsa_op_with_type *size, unsigned alignment)
1746 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA, dest->m_type, dest, size),
1747 m_align (BRIG_ALIGNMENT_8)
1748{
1749 gcc_assert (dest->m_type == BRIG_TYPE_U32);
1750 if (alignment)
1751 m_align = hsa_alignment_encoding (alignment);
1752}
1753
1754/* Append an instruction INSN into the basic block. */
1755
1756void
1757hsa_bb::append_insn (hsa_insn_basic *insn)
1758{
1759 gcc_assert (insn->m_opcode != 0 || insn->operand_count () == 0);
1760 gcc_assert (!insn->m_bb);
1761
1762 insn->m_bb = m_bb;
1763 insn->m_prev = m_last_insn;
1764 insn->m_next = NULL;
1765 if (m_last_insn)
1766 m_last_insn->m_next = insn;
1767 m_last_insn = insn;
1768 if (!m_first_insn)
1769 m_first_insn = insn;
1770}
1771
1772void
1773hsa_bb::append_phi (hsa_insn_phi *hphi)
1774{
1775 hphi->m_bb = m_bb;
1776
1777 hphi->m_prev = m_last_phi;
1778 hphi->m_next = NULL;
1779 if (m_last_phi)
1780 m_last_phi->m_next = hphi;
1781 m_last_phi = hphi;
1782 if (!m_first_phi)
1783 m_first_phi = hphi;
1784}
1785
1786/* Insert HSA instruction NEW_INSN immediately before an existing instruction
1787 OLD_INSN. */
1788
1789static void
1790hsa_insert_insn_before (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
1791{
1792 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
1793
1794 if (hbb->m_first_insn == old_insn)
1795 hbb->m_first_insn = new_insn;
1796 new_insn->m_prev = old_insn->m_prev;
1797 new_insn->m_next = old_insn;
1798 if (old_insn->m_prev)
1799 old_insn->m_prev->m_next = new_insn;
1800 old_insn->m_prev = new_insn;
1801}
1802
1803/* Append HSA instruction NEW_INSN immediately after an existing instruction
1804 OLD_INSN. */
1805
1806static void
1807hsa_append_insn_after (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
1808{
1809 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
1810
1811 if (hbb->m_last_insn == old_insn)
1812 hbb->m_last_insn = new_insn;
1813 new_insn->m_prev = old_insn;
1814 new_insn->m_next = old_insn->m_next;
1815 if (old_insn->m_next)
1816 old_insn->m_next->m_prev = new_insn;
1817 old_insn->m_next = new_insn;
1818}
1819
1820/* Return a register containing the calculated value of EXP which must be an
1821 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
1822 integer constants as returned by get_inner_reference.
1823 Newly generated HSA instructions will be appended to HBB.
1824 Perform all calculations in ADDRTYPE. */
1825
1826static hsa_op_with_type *
1827gen_address_calculation (tree exp, hsa_bb *hbb, BrigType16_t addrtype)
1828{
1829 int opcode;
1830
1831 if (TREE_CODE (exp) == NOP_EXPR)
1832 exp = TREE_OPERAND (exp, 0);
1833
1834 switch (TREE_CODE (exp))
1835 {
1836 case SSA_NAME:
1837 return hsa_cfun->reg_for_gimple_ssa (exp)->get_in_type (addrtype, hbb);
1838
1839 case INTEGER_CST:
1840 {
1841 hsa_op_immed *imm = new hsa_op_immed (exp);
1842 if (addrtype != imm->m_type)
1843 imm->m_type = addrtype;
1844 return imm;
1845 }
1846
1847 case PLUS_EXPR:
1848 opcode = BRIG_OPCODE_ADD;
1849 break;
1850
1851 case MULT_EXPR:
1852 opcode = BRIG_OPCODE_MUL;
1853 break;
1854
1855 default:
1856 gcc_unreachable ();
1857 }
1858
1859 hsa_op_reg *res = new hsa_op_reg (addrtype);
1860 hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, addrtype);
1861 insn->set_op (0, res);
1862
1863 hsa_op_with_type *op1 = gen_address_calculation (TREE_OPERAND (exp, 0), hbb,
1864 addrtype);
1865 hsa_op_with_type *op2 = gen_address_calculation (TREE_OPERAND (exp, 1), hbb,
1866 addrtype);
1867 insn->set_op (1, op1);
1868 insn->set_op (2, op2);
1869
1870 hbb->append_insn (insn);
1871 return res;
1872}
1873
1874/* If R1 is NULL, just return R2, otherwise append an instruction adding them
1875 to HBB and return the register holding the result. */
1876
1877static hsa_op_reg *
1878add_addr_regs_if_needed (hsa_op_reg *r1, hsa_op_reg *r2, hsa_bb *hbb)
1879{
1880 gcc_checking_assert (r2);
1881 if (!r1)
1882 return r2;
1883
1884 hsa_op_reg *res = new hsa_op_reg (r1->m_type);
1885 gcc_assert (!hsa_needs_cvt (r1->m_type, r2->m_type));
1886 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_ADD, res->m_type);
1887 insn->set_op (0, res);
1888 insn->set_op (1, r1);
1889 insn->set_op (2, r2);
1890 hbb->append_insn (insn);
1891 return res;
1892}
1893
1894/* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
1895 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
1896
1897static void
1898process_mem_base (tree base, hsa_symbol **symbol, BrigType16_t *addrtype,
1899 hsa_op_reg **reg, offset_int *offset, hsa_bb *hbb)
1900{
1901 if (TREE_CODE (base) == SSA_NAME)
1902 {
1903 gcc_assert (!*reg);
1904 hsa_op_with_type *ssa
1905 = hsa_cfun->reg_for_gimple_ssa (base)->get_in_type (*addrtype, hbb);
1906 *reg = dyn_cast <hsa_op_reg *> (ssa);
1907 }
1908 else if (TREE_CODE (base) == ADDR_EXPR)
1909 {
1910 tree decl = TREE_OPERAND (base, 0);
1911
1912 if (!DECL_P (decl) || TREE_CODE (decl) == FUNCTION_DECL)
1913 {
1914 HSA_SORRY_AT (EXPR_LOCATION (base),
1915 "support for HSA does not implement a memory reference "
1916 "to a non-declaration type");
1917 return;
1918 }
1919
1920 gcc_assert (!*symbol);
1921
1922 *symbol = get_symbol_for_decl (decl);
1923 *addrtype = hsa_get_segment_addr_type ((*symbol)->m_segment);
1924 }
1925 else if (TREE_CODE (base) == INTEGER_CST)
1926 *offset += wi::to_offset (base);
1927 else
1928 gcc_unreachable ();
1929}
1930
1931/* Forward declaration of a function. */
1932
1933static void
1934gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb);
1935
1936/* Generate HSA address operand for a given tree memory reference REF. If
1937 instructions need to be created to calculate the address, they will be added
1938 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
1939 the function assumes that the caller will handle possible
1940 bit-field references. Otherwise if we reference a bit-field, sorry message
1941 is displayed. */
1942
1943static hsa_op_address *
1944gen_hsa_addr (tree ref, hsa_bb *hbb, HOST_WIDE_INT *output_bitsize = NULL,
1945 HOST_WIDE_INT *output_bitpos = NULL)
1946{
1947 hsa_symbol *symbol = NULL;
1948 hsa_op_reg *reg = NULL;
1949 offset_int offset = 0;
1950 tree origref = ref;
1951 tree varoffset = NULL_TREE;
1952 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
1953 HOST_WIDE_INT bitsize = 0, bitpos = 0;
1954 BrigType16_t flat_addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
1955
1956 if (TREE_CODE (ref) == STRING_CST)
1957 {
1958 symbol = hsa_get_string_cst_symbol (ref);
1959 goto out;
1960 }
1961 else if (TREE_CODE (ref) == BIT_FIELD_REF
1962 && ((tree_to_uhwi (TREE_OPERAND (ref, 1)) % BITS_PER_UNIT) != 0
1963 || (tree_to_uhwi (TREE_OPERAND (ref, 2)) % BITS_PER_UNIT) != 0))
1964 {
1965 HSA_SORRY_ATV (EXPR_LOCATION (origref),
1966 "support for HSA does not implement "
1967 "bit field references such as %E", ref);
1968 goto out;
1969 }
1970
1971 if (handled_component_p (ref))
1972 {
1973 machine_mode mode;
1974 int unsignedp, volatilep, preversep;
1975
1976 ref = get_inner_reference (ref, &bitsize, &bitpos, &varoffset, &mode,
1977 &unsignedp, &preversep, &volatilep);
1978
1979 offset = bitpos;
1980 offset = wi::rshift (offset, LOG2_BITS_PER_UNIT, SIGNED);
1981 }
1982
1983 switch (TREE_CODE (ref))
1984 {
1985 case ADDR_EXPR:
1986 {
1987 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
1988 symbol = hsa_cfun->create_hsa_temporary (flat_addrtype);
1989 hsa_op_reg *r = new hsa_op_reg (flat_addrtype);
1990 gen_hsa_addr_insns (ref, r, hbb);
1991 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
1992 r, new hsa_op_address (symbol)));
1993
1994 break;
1995 }
1996 case SSA_NAME:
1997 {
1998 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
1999 hsa_op_with_type *r = hsa_cfun->reg_for_gimple_ssa (ref);
2000 if (r->m_type == BRIG_TYPE_B1)
2001 r = r->get_in_type (BRIG_TYPE_U32, hbb);
2002 symbol = hsa_cfun->create_hsa_temporary (r->m_type);
2003
2004 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
2005 r, new hsa_op_address (symbol)));
2006
2007 break;
2008 }
2009 case PARM_DECL:
2010 case VAR_DECL:
2011 case RESULT_DECL:
2012 case CONST_DECL:
2013 gcc_assert (!symbol);
2014 symbol = get_symbol_for_decl (ref);
2015 addrtype = hsa_get_segment_addr_type (symbol->m_segment);
2016 break;
2017
2018 case MEM_REF:
2019 process_mem_base (TREE_OPERAND (ref, 0), &symbol, &addrtype, &reg,
2020 &offset, hbb);
2021
2022 if (!integer_zerop (TREE_OPERAND (ref, 1)))
2023 offset += wi::to_offset (TREE_OPERAND (ref, 1));
2024 break;
2025
2026 case TARGET_MEM_REF:
2027 process_mem_base (TMR_BASE (ref), &symbol, &addrtype, &reg, &offset, hbb);
2028 if (TMR_INDEX (ref))
2029 {
2030 hsa_op_reg *disp1;
2031 hsa_op_base *idx = hsa_cfun->reg_for_gimple_ssa
2032 (TMR_INDEX (ref))->get_in_type (addrtype, hbb);
2033 if (TMR_STEP (ref) && !integer_onep (TMR_STEP (ref)))
2034 {
2035 disp1 = new hsa_op_reg (addrtype);
2036 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_MUL,
2037 addrtype);
2038
2039 /* As step must respect addrtype, we overwrite the type
2040 of an immediate value. */
2041 hsa_op_immed *step = new hsa_op_immed (TMR_STEP (ref));
2042 step->m_type = addrtype;
2043
2044 insn->set_op (0, disp1);
2045 insn->set_op (1, idx);
2046 insn->set_op (2, step);
2047 hbb->append_insn (insn);
2048 }
2049 else
2050 disp1 = as_a <hsa_op_reg *> (idx);
2051 reg = add_addr_regs_if_needed (reg, disp1, hbb);
2052 }
2053 if (TMR_INDEX2 (ref))
2054 {
2055 if (TREE_CODE (TMR_INDEX2 (ref)) == SSA_NAME)
2056 {
2057 hsa_op_base *disp2 = hsa_cfun->reg_for_gimple_ssa
2058 (TMR_INDEX2 (ref))->get_in_type (addrtype, hbb);
2059 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (disp2),
2060 hbb);
2061 }
2062 else if (TREE_CODE (TMR_INDEX2 (ref)) == INTEGER_CST)
2063 offset += wi::to_offset (TMR_INDEX2 (ref));
2064 else
2065 gcc_unreachable ();
2066 }
2067 offset += wi::to_offset (TMR_OFFSET (ref));
2068 break;
2069 case FUNCTION_DECL:
2070 HSA_SORRY_AT (EXPR_LOCATION (origref),
2071 "support for HSA does not implement function pointers");
2072 goto out;
2073 default:
2074 HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does "
2075 "not implement memory access to %E", origref);
2076 goto out;
2077 }
2078
2079 if (varoffset)
2080 {
2081 if (TREE_CODE (varoffset) == INTEGER_CST)
2082 offset += wi::to_offset (varoffset);
2083 else
2084 {
2085 hsa_op_base *off_op = gen_address_calculation (varoffset, hbb,
2086 addrtype);
2087 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (off_op),
2088 hbb);
2089 }
2090 }
2091
2092 gcc_checking_assert ((symbol
2093 && addrtype
2094 == hsa_get_segment_addr_type (symbol->m_segment))
2095 || (!symbol
2096 && addrtype
2097 == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT)));
2098out:
2099 HOST_WIDE_INT hwi_offset = offset.to_shwi ();
2100
2101 /* Calculate remaining bitsize offset (if presented). */
2102 bitpos %= BITS_PER_UNIT;
2103 /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
2104 is not a reason to think this is a bit-field access. */
2105 if (bitpos == 0
2106 && (bitsize >= BITS_PER_UNIT)
2107 && !(bitsize & (bitsize - 1)))
2108 bitsize = 0;
2109
2110 if ((bitpos || bitsize) && (output_bitpos == NULL || output_bitsize == NULL))
2111 HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does not "
2112 "implement unhandled bit field reference such as %E", ref);
2113
2114 if (output_bitsize != NULL && output_bitpos != NULL)
2115 {
2116 *output_bitsize = bitsize;
2117 *output_bitpos = bitpos;
2118 }
2119
2120 return new hsa_op_address (symbol, reg, hwi_offset);
2121}
2122
2123/* Generate HSA address operand for a given tree memory reference REF. If
2124 instructions need to be created to calculate the address, they will be added
2125 to the end of HBB. OUTPUT_ALIGN is alignment of the created address. */
2126
2127static hsa_op_address *
2128gen_hsa_addr_with_align (tree ref, hsa_bb *hbb, BrigAlignment8_t *output_align)
2129{
2130 hsa_op_address *addr = gen_hsa_addr (ref, hbb);
2131 if (addr->m_reg || !addr->m_symbol)
2132 *output_align = hsa_object_alignment (ref);
2133 else
2134 {
2135 /* If the address consists only of a symbol and an offset, we
2136 compute the alignment ourselves to take into account any alignment
2137 promotions we might have done for the HSA symbol representation. */
2138 unsigned align = hsa_byte_alignment (addr->m_symbol->m_align);
2139 unsigned misalign = addr->m_imm_offset & (align - 1);
2140 if (misalign)
2141 align = least_bit_hwi (misalign);
2142 *output_align = hsa_alignment_encoding (BITS_PER_UNIT * align);
2143 }
2144 return addr;
2145}
2146
2147/* Generate HSA address for a function call argument of given TYPE.
2148 INDEX is used to generate corresponding name of the arguments.
2149 Special value -1 represents fact that result value is created. */
2150
2151static hsa_op_address *
2152gen_hsa_addr_for_arg (tree tree_type, int index)
2153{
2154 hsa_symbol *sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
2155 BRIG_LINKAGE_ARG);
2156 sym->m_type = hsa_type_for_tree_type (tree_type, &sym->m_dim);
2157
2158 if (index == -1) /* Function result. */
2159 sym->m_name = "res";
2160 else /* Function call arguments. */
2161 {
2162 sym->m_name = NULL;
2163 sym->m_name_number = index;
2164 }
2165
2166 return new hsa_op_address (sym);
2167}
2168
2169/* Generate HSA instructions that process all necessary conversions
2170 of an ADDR to flat addressing and place the result into DEST.
2171 Instructions are appended to HBB. */
2172
2173static void
2174convert_addr_to_flat_segment (hsa_op_address *addr, hsa_op_reg *dest,
2175 hsa_bb *hbb)
2176{
2177 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_LDA);
2178 insn->set_op (1, addr);
2179 if (addr->m_symbol && addr->m_symbol->m_segment != BRIG_SEGMENT_GLOBAL)
2180 {
2181 /* LDA produces segment-relative address, we need to convert
2182 it to the flat one. */
2183 hsa_op_reg *tmp;
2184 tmp = new hsa_op_reg (hsa_get_segment_addr_type
2185 (addr->m_symbol->m_segment));
2186 hsa_insn_seg *seg;
2187 seg = new hsa_insn_seg (BRIG_OPCODE_STOF,
2188 hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT),
2189 tmp->m_type, addr->m_symbol->m_segment, dest,
2190 tmp);
2191
2192 insn->set_op (0, tmp);
2193 insn->m_type = tmp->m_type;
2194 hbb->append_insn (insn);
2195 hbb->append_insn (seg);
2196 }
2197 else
2198 {
2199 insn->set_op (0, dest);
2200 insn->m_type = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
2201 hbb->append_insn (insn);
2202 }
2203}
2204
2205/* Generate HSA instructions that calculate address of VAL including all
2206 necessary conversions to flat addressing and place the result into DEST.
2207 Instructions are appended to HBB. */
2208
2209static void
2210gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb)
2211{
2212 /* Handle cases like tmp = NULL, where we just emit a move instruction
2213 to a register. */
2214 if (TREE_CODE (val) == INTEGER_CST)
2215 {
2216 hsa_op_immed *c = new hsa_op_immed (val);
2217 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
2218 dest->m_type, dest, c);
2219 hbb->append_insn (insn);
2220 return;
2221 }
2222
2223 hsa_op_address *addr;
2224
2225 gcc_assert (dest->m_type == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
2226 if (TREE_CODE (val) == ADDR_EXPR)
2227 val = TREE_OPERAND (val, 0);
2228 addr = gen_hsa_addr (val, hbb);
2229
2230 if (TREE_CODE (val) == CONST_DECL
2231 && is_gimple_reg_type (TREE_TYPE (val)))
2232 {
2233 gcc_assert (addr->m_symbol
2234 && addr->m_symbol->m_segment == BRIG_SEGMENT_READONLY);
2235 /* CONST_DECLs are in readonly segment which however does not have
2236 addresses convertible to flat segments. So copy it to a private one
2237 and take address of that. */
2238 BrigType16_t csttype
2239 = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (val),
2240 false));
2241 hsa_op_reg *r = new hsa_op_reg (csttype);
2242 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD, csttype, r,
2243 new hsa_op_address (addr->m_symbol)));
2244 hsa_symbol *copysym = hsa_cfun->create_hsa_temporary (csttype);
2245 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, csttype, r,
2246 new hsa_op_address (copysym)));
2247 addr->m_symbol = copysym;
2248 }
2249 else if (addr->m_symbol && addr->m_symbol->m_segment == BRIG_SEGMENT_READONLY)
2250 {
2251 HSA_SORRY_ATV (EXPR_LOCATION (val), "support for HSA does "
2252 "not implement taking addresses of complex "
2253 "CONST_DECLs such as %E", val);
2254 return;
2255 }
2256
2257
2258 convert_addr_to_flat_segment (addr, dest, hbb);
2259}
2260
2261/* Return an HSA register or HSA immediate value operand corresponding to
2262 gimple operand OP. */
2263
2264static hsa_op_with_type *
2265hsa_reg_or_immed_for_gimple_op (tree op, hsa_bb *hbb)
2266{
2267 hsa_op_reg *tmp;
2268
2269 if (TREE_CODE (op) == SSA_NAME)
2270 tmp = hsa_cfun->reg_for_gimple_ssa (op);
2271 else if (!POINTER_TYPE_P (TREE_TYPE (op)))
2272 return new hsa_op_immed (op);
2273 else
2274 {
2275 tmp = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
2276 gen_hsa_addr_insns (op, tmp, hbb);
2277 }
2278 return tmp;
2279}
2280
2281/* Create a simple movement instruction with register destination DEST and
2282 register or immediate source SRC and append it to the end of HBB. */
2283
2284void
2285hsa_build_append_simple_mov (hsa_op_reg *dest, hsa_op_base *src, hsa_bb *hbb)
2286{
2287 /* Moves of packed data between registers need to adhere to the same type
2288 rules like when dealing with memory. */
2289 BrigType16_t tp = mem_type_for_type (dest->m_type);
2290 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, tp, dest, src);
2291 hsa_fixup_mov_insn_type (insn);
2292 unsigned dest_size = hsa_type_bit_size (dest->m_type);
2293 if (hsa_op_reg *sreg = dyn_cast <hsa_op_reg *> (src))
2294 gcc_assert (dest_size == hsa_type_bit_size (sreg->m_type));
2295 else
2296 {
2297 unsigned imm_size
2298 = hsa_type_bit_size (as_a <hsa_op_immed *> (src)->m_type);
2299 gcc_assert ((dest_size == imm_size)
2300 /* Eventually < 32bit registers will be promoted to 32bit. */
2301 || (dest_size < 32 && imm_size == 32));
2302 }
2303 hbb->append_insn (insn);
2304}
2305
2306/* Generate HSAIL instructions loading a bit field into register DEST.
2307 VALUE_REG is a register of a SSA name that is used in the bit field
2308 reference. To identify a bit field BITPOS is offset to the loaded memory
2309 and BITSIZE is number of bits of the bit field.
2310 Add instructions to HBB. */
2311
2312static void
2313gen_hsa_insns_for_bitfield (hsa_op_reg *dest, hsa_op_reg *value_reg,
2314 HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos,
2315 hsa_bb *hbb)
2316{
2317 unsigned type_bitsize
2318 = hsa_type_bit_size (hsa_extend_inttype_to_32bit (dest->m_type));
2319 unsigned left_shift = type_bitsize - (bitsize + bitpos);
2320 unsigned right_shift = left_shift + bitpos;
2321
2322 if (left_shift)
2323 {
2324 hsa_op_reg *value_reg_2
2325 = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest->m_type));
2326 hsa_op_immed *c = new hsa_op_immed (left_shift, BRIG_TYPE_U32);
2327
2328 hsa_insn_basic *lshift
2329 = new hsa_insn_basic (3, BRIG_OPCODE_SHL, value_reg_2->m_type,
2330 value_reg_2, value_reg, c);
2331
2332 hbb->append_insn (lshift);
2333
2334 value_reg = value_reg_2;
2335 }
2336
2337 if (right_shift)
2338 {
2339 hsa_op_reg *value_reg_2
2340 = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest->m_type));
2341 hsa_op_immed *c = new hsa_op_immed (right_shift, BRIG_TYPE_U32);
2342
2343 hsa_insn_basic *rshift
2344 = new hsa_insn_basic (3, BRIG_OPCODE_SHR, value_reg_2->m_type,
2345 value_reg_2, value_reg, c);
2346
2347 hbb->append_insn (rshift);
2348
2349 value_reg = value_reg_2;
2350 }
2351
2352 hsa_insn_basic *assignment
2353 = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, NULL, value_reg);
2354 hsa_fixup_mov_insn_type (assignment);
2355 hbb->append_insn (assignment);
2356 assignment->set_output_in_type (dest, 0, hbb);
2357}
2358
2359
2360/* Generate HSAIL instructions loading a bit field into register DEST. ADDR is
2361 prepared memory address which is used to load the bit field. To identify a
2362 bit field BITPOS is offset to the loaded memory and BITSIZE is number of
2363 bits of the bit field. Add instructions to HBB. Load must be performed in
2364 alignment ALIGN. */
2365
2366static void
2367gen_hsa_insns_for_bitfield_load (hsa_op_reg *dest, hsa_op_address *addr,
2368 HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos,
2369 hsa_bb *hbb, BrigAlignment8_t align)
2370{
2371 hsa_op_reg *value_reg = new hsa_op_reg (dest->m_type);
2372 hsa_insn_mem *mem
2373 = new hsa_insn_mem (BRIG_OPCODE_LD,
2374 hsa_extend_inttype_to_32bit (dest->m_type),
2375 value_reg, addr);
2376 mem->set_align (align);
2377 hbb->append_insn (mem);
2378 gen_hsa_insns_for_bitfield (dest, value_reg, bitsize, bitpos, hbb);
2379}
2380
2381/* Return the alignment of base memory accesses we issue to perform bit-field
2382 memory access REF. */
2383
2384static BrigAlignment8_t
2385hsa_bitmemref_alignment (tree ref)
2386{
2387 unsigned HOST_WIDE_INT bit_offset = 0;
2388
2389 while (true)
2390 {
2391 if (TREE_CODE (ref) == BIT_FIELD_REF)
2392 {
2393 if (!tree_fits_uhwi_p (TREE_OPERAND (ref, 2)))
2394 return BRIG_ALIGNMENT_1;
2395 bit_offset += tree_to_uhwi (TREE_OPERAND (ref, 2));
2396 }
2397 else if (TREE_CODE (ref) == COMPONENT_REF
2398 && DECL_BIT_FIELD (TREE_OPERAND (ref, 1)))
2399 bit_offset += int_bit_position (TREE_OPERAND (ref, 1));
2400 else
2401 break;
2402 ref = TREE_OPERAND (ref, 0);
2403 }
2404
2405 unsigned HOST_WIDE_INT bits = bit_offset % BITS_PER_UNIT;
2406 unsigned HOST_WIDE_INT byte_bits = bit_offset - bits;
2407 BrigAlignment8_t base = hsa_object_alignment (ref);
2408 if (byte_bits == 0)
2409 return base;
2410 return MIN (base, hsa_alignment_encoding (least_bit_hwi (byte_bits)));
2411}
2412
2413/* Generate HSAIL instructions loading something into register DEST. RHS is
2414 tree representation of the loaded data, which are loaded as type TYPE. Add
2415 instructions to HBB. */
2416
2417static void
2418gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb)
2419{
2420 /* The destination SSA name will give us the type. */
2421 if (TREE_CODE (rhs) == VIEW_CONVERT_EXPR)
2422 rhs = TREE_OPERAND (rhs, 0);
2423
2424 if (TREE_CODE (rhs) == SSA_NAME)
2425 {
2426 hsa_op_reg *src = hsa_cfun->reg_for_gimple_ssa (rhs);
2427 hsa_build_append_simple_mov (dest, src, hbb);
2428 }
2429 else if (is_gimple_min_invariant (rhs)
2430 || TREE_CODE (rhs) == ADDR_EXPR)
2431 {
2432 if (POINTER_TYPE_P (TREE_TYPE (rhs)))
2433 {
2434 if (dest->m_type != hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT))
2435 {
2436 HSA_SORRY_ATV (EXPR_LOCATION (rhs),
2437 "support for HSA does not implement conversion "
2438 "of %E to the requested non-pointer type.", rhs);
2439 return;
2440 }
2441
2442 gen_hsa_addr_insns (rhs, dest, hbb);
2443 }
2444 else if (TREE_CODE (rhs) == COMPLEX_CST)
2445 {
2446 hsa_op_immed *real_part = new hsa_op_immed (TREE_REALPART (rhs));
2447 hsa_op_immed *imag_part = new hsa_op_immed (TREE_IMAGPART (rhs));
2448
2449 hsa_op_reg *real_part_reg
2450 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type),
2451 true));
2452 hsa_op_reg *imag_part_reg
2453 = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type),
2454 true));
2455
2456 hsa_build_append_simple_mov (real_part_reg, real_part, hbb);
2457 hsa_build_append_simple_mov (imag_part_reg, imag_part, hbb);
2458
2459 BrigType16_t src_type = hsa_bittype_for_type (real_part_reg->m_type);
2460
2461 hsa_insn_packed *insn
2462 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type,
2463 src_type, dest, real_part_reg,
2464 imag_part_reg);
2465 hbb->append_insn (insn);
2466 }
2467 else
2468 {
2469 hsa_op_immed *imm = new hsa_op_immed (rhs);
2470 hsa_build_append_simple_mov (dest, imm, hbb);
2471 }
2472 }
2473 else if (TREE_CODE (rhs) == REALPART_EXPR || TREE_CODE (rhs) == IMAGPART_EXPR)
2474 {
2475 tree pack_type = TREE_TYPE (TREE_OPERAND (rhs, 0));
2476
2477 hsa_op_reg *packed_reg
2478 = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type, true));
2479
2480 tree complex_rhs = TREE_OPERAND (rhs, 0);
2481 gen_hsa_insns_for_load (packed_reg, complex_rhs, TREE_TYPE (complex_rhs),
2482 hbb);
2483
2484 hsa_op_reg *real_reg
2485 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true));
2486
2487 hsa_op_reg *imag_reg
2488 = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true));
2489
2490 BrigKind16_t brig_type = packed_reg->m_type;
2491 hsa_insn_packed *packed
2492 = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND,
2493 hsa_bittype_for_type (real_reg->m_type),
2494 brig_type, real_reg, imag_reg, packed_reg);
2495
2496 hbb->append_insn (packed);
2497
2498 hsa_op_reg *source = TREE_CODE (rhs) == REALPART_EXPR ?
2499 real_reg : imag_reg;
2500
2501 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
2502 dest->m_type, NULL, source);
2503 hsa_fixup_mov_insn_type (insn);
2504 hbb->append_insn (insn);
2505 insn->set_output_in_type (dest, 0, hbb);
2506 }
2507 else if (TREE_CODE (rhs) == BIT_FIELD_REF
2508 && TREE_CODE (TREE_OPERAND (rhs, 0)) == SSA_NAME)
2509 {
2510 tree ssa_name = TREE_OPERAND (rhs, 0);
2511 HOST_WIDE_INT bitsize = tree_to_uhwi (TREE_OPERAND (rhs, 1));
2512 HOST_WIDE_INT bitpos = tree_to_uhwi (TREE_OPERAND (rhs, 2));
2513
2514 hsa_op_reg *imm_value = hsa_cfun->reg_for_gimple_ssa (ssa_name);
2515 gen_hsa_insns_for_bitfield (dest, imm_value, bitsize, bitpos, hbb);
2516 }
2517 else if (DECL_P (rhs) || TREE_CODE (rhs) == MEM_REF
2518 || TREE_CODE (rhs) == TARGET_MEM_REF
2519 || handled_component_p (rhs))
2520 {
2521 HOST_WIDE_INT bitsize, bitpos;
2522
2523 /* Load from memory. */
2524 hsa_op_address *addr;
2525 addr = gen_hsa_addr (rhs, hbb, &bitsize, &bitpos);
2526
2527 /* Handle load of a bit field. */
2528 if (bitsize > 64)
2529 {
2530 HSA_SORRY_AT (EXPR_LOCATION (rhs),
2531 "support for HSA does not implement load from a bit "
2532 "field bigger than 64 bits");
2533 return;
2534 }
2535
2536 if (bitsize || bitpos)
2537 gen_hsa_insns_for_bitfield_load (dest, addr, bitsize, bitpos, hbb,
2538 hsa_bitmemref_alignment (rhs));
2539 else
2540 {
2541 BrigType16_t mtype;
2542 /* Not dest->m_type, that's possibly extended. */
2543 mtype = mem_type_for_type (hsa_type_for_scalar_tree_type (type,
2544 false));
2545 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, dest,
2546 addr);
2547 mem->set_align (hsa_object_alignment (rhs));
2548 hbb->append_insn (mem);
2549 }
2550 }
2551 else
2552 HSA_SORRY_ATV (EXPR_LOCATION (rhs),
2553 "support for HSA does not implement loading "
2554 "of expression %E",
2555 rhs);
2556}
2557
2558/* Return number of bits necessary for representation of a bit field,
2559 starting at BITPOS with size of BITSIZE. */
2560
2561static unsigned
2562get_bitfield_size (unsigned bitpos, unsigned bitsize)
2563{
2564 unsigned s = bitpos + bitsize;
2565 unsigned sizes[] = {8, 16, 32, 64};
2566
2567 for (unsigned i = 0; i < 4; i++)
2568 if (s <= sizes[i])
2569 return sizes[i];
2570
2571 gcc_unreachable ();
2572 return 0;
2573}
2574
2575/* Generate HSAIL instructions storing into memory. LHS is the destination of
2576 the store, SRC is the source operand. Add instructions to HBB. */
2577
2578static void
2579gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb)
2580{
2581 HOST_WIDE_INT bitsize = 0, bitpos = 0;
2582 BrigAlignment8_t req_align;
2583 BrigType16_t mtype;
2584 mtype = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
2585 false));
2586 hsa_op_address *addr;
2587 addr = gen_hsa_addr (lhs, hbb, &bitsize, &bitpos);
2588
2589 /* Handle store to a bit field. */
2590 if (bitsize > 64)
2591 {
2592 HSA_SORRY_AT (EXPR_LOCATION (lhs),
2593 "support for HSA does not implement store to a bit field "
2594 "bigger than 64 bits");
2595 return;
2596 }
2597
2598 unsigned type_bitsize = get_bitfield_size (bitpos, bitsize);
2599
2600 /* HSAIL does not support MOV insn with 16-bits integers. */
2601 if (type_bitsize < 32)
2602 type_bitsize = 32;
2603
2604 if (bitpos || (bitsize && type_bitsize != bitsize))
2605 {
2606 unsigned HOST_WIDE_INT mask = 0;
2607 BrigType16_t mem_type
2608 = get_integer_type_by_bytes (type_bitsize / BITS_PER_UNIT,
2609 !TYPE_UNSIGNED (TREE_TYPE (lhs)));
2610
2611 for (unsigned i = 0; i < type_bitsize; i++)
2612 if (i < bitpos || i >= bitpos + bitsize)
2613 mask |= ((unsigned HOST_WIDE_INT)1 << i);
2614
2615 hsa_op_reg *value_reg = new hsa_op_reg (mem_type);
2616
2617 req_align = hsa_bitmemref_alignment (lhs);
2618 /* Load value from memory. */
2619 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mem_type,
2620 value_reg, addr);
2621 mem->set_align (req_align);
2622 hbb->append_insn (mem);
2623
2624 /* AND the loaded value with prepared mask. */
2625 hsa_op_reg *cleared_reg = new hsa_op_reg (mem_type);
2626
2627 BrigType16_t t
2628 = get_integer_type_by_bytes (type_bitsize / BITS_PER_UNIT, false);
2629 hsa_op_immed *c = new hsa_op_immed (mask, t);
2630
2631 hsa_insn_basic *clearing
2632 = new hsa_insn_basic (3, BRIG_OPCODE_AND, mem_type, cleared_reg,
2633 value_reg, c);
2634 hbb->append_insn (clearing);
2635
2636 /* Shift to left a value that is going to be stored. */
2637 hsa_op_reg *new_value_reg = new hsa_op_reg (mem_type);
2638
2639 hsa_insn_basic *basic = new hsa_insn_basic (2, BRIG_OPCODE_MOV, mem_type,
2640 new_value_reg, src);
2641 hsa_fixup_mov_insn_type (basic);
2642 hbb->append_insn (basic);
2643
2644 if (bitpos)
2645 {
2646 hsa_op_reg *shifted_value_reg = new hsa_op_reg (mem_type);
2647 c = new hsa_op_immed (bitpos, BRIG_TYPE_U32);
2648
2649 hsa_insn_basic *basic
2650 = new hsa_insn_basic (3, BRIG_OPCODE_SHL, mem_type,
2651 shifted_value_reg, new_value_reg, c);
2652 hbb->append_insn (basic);
2653
2654 new_value_reg = shifted_value_reg;
2655 }
2656
2657 /* OR the prepared value with prepared chunk loaded from memory. */
2658 hsa_op_reg *prepared_reg= new hsa_op_reg (mem_type);
2659 basic = new hsa_insn_basic (3, BRIG_OPCODE_OR, mem_type, prepared_reg,
2660 new_value_reg, cleared_reg);
2661 hbb->append_insn (basic);
2662
2663 src = prepared_reg;
2664 mtype = mem_type;
2665 }
2666 else
2667 req_align = hsa_object_alignment (lhs);
2668
2669 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, mtype, src, addr);
2670 mem->set_align (req_align);
2671
2672 /* The HSAIL verifier has another constraint: if the source is an immediate
2673 then it must match the destination type. If it's a register the low bits
2674 will be used for sub-word stores. We're always allocating new operands so
2675 we can modify the above in place. */
2676 if (hsa_op_immed *imm = dyn_cast <hsa_op_immed *> (src))
2677 {
2678 if (!hsa_type_packed_p (imm->m_type))
2679 imm->m_type = mem->m_type;
2680 else
2681 {
2682 /* ...and all vector immediates apparently need to be vectors of
2683 unsigned bytes. */
2684 unsigned bs = hsa_type_bit_size (imm->m_type);
2685 gcc_assert (bs == hsa_type_bit_size (mem->m_type));
2686 switch (bs)
2687 {
2688 case 32:
2689 imm->m_type = BRIG_TYPE_U8X4;
2690 break;
2691 case 64:
2692 imm->m_type = BRIG_TYPE_U8X8;
2693 break;
2694 case 128:
2695 imm->m_type = BRIG_TYPE_U8X16;
2696 break;
2697 default:
2698 gcc_unreachable ();
2699 }
2700 }
2701 }
2702
2703 hbb->append_insn (mem);
2704}
2705
2706/* Generate memory copy instructions that are going to be used
2707 for copying a SRC memory to TARGET memory,
2708 represented by pointer in a register. MIN_ALIGN is minimal alignment
2709 of provided HSA addresses. */
2710
2711static void
2712gen_hsa_memory_copy (hsa_bb *hbb, hsa_op_address *target, hsa_op_address *src,
2713 unsigned size, BrigAlignment8_t min_align)
2714{
2715 hsa_op_address *addr;
2716 hsa_insn_mem *mem;
2717
2718 unsigned offset = 0;
2719 unsigned min_byte_align = hsa_byte_alignment (min_align);
2720
2721 while (size)
2722 {
2723 unsigned s;
2724 if (size >= 8)
2725 s = 8;
2726 else if (size >= 4)
2727 s = 4;
2728 else if (size >= 2)
2729 s = 2;
2730 else
2731 s = 1;
2732
2733 if (s > min_byte_align)
2734 s = min_byte_align;
2735
2736 BrigType16_t t = get_integer_type_by_bytes (s, false);
2737
2738 hsa_op_reg *tmp = new hsa_op_reg (t);
2739 addr = new hsa_op_address (src->m_symbol, src->m_reg,
2740 src->m_imm_offset + offset);
2741 mem = new hsa_insn_mem (BRIG_OPCODE_LD, t, tmp, addr);
2742 hbb->append_insn (mem);
2743
2744 addr = new hsa_op_address (target->m_symbol, target->m_reg,
2745 target->m_imm_offset + offset);
2746 mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, tmp, addr);
2747 hbb->append_insn (mem);
2748 offset += s;
2749 size -= s;
2750 }
2751}
2752
2753/* Create a memset mask that is created by copying a CONSTANT byte value
2754 to an integer of BYTE_SIZE bytes. */
2755
2756static unsigned HOST_WIDE_INT
2757build_memset_value (unsigned HOST_WIDE_INT constant, unsigned byte_size)
2758{
2759 if (constant == 0)
2760 return 0;
2761
2762 HOST_WIDE_INT v = constant;
2763
2764 for (unsigned i = 1; i < byte_size; i++)
2765 v |= constant << (8 * i);
2766
2767 return v;
2768}
2769
2770/* Generate memory set instructions that are going to be used
2771 for setting a CONSTANT byte value to TARGET memory of SIZE bytes.
2772 MIN_ALIGN is minimal alignment of provided HSA addresses. */
2773
2774static void
2775gen_hsa_memory_set (hsa_bb *hbb, hsa_op_address *target,
2776 unsigned HOST_WIDE_INT constant,
2777 unsigned size, BrigAlignment8_t min_align)
2778{
2779 hsa_op_address *addr;
2780 hsa_insn_mem *mem;
2781
2782 unsigned offset = 0;
2783 unsigned min_byte_align = hsa_byte_alignment (min_align);
2784
2785 while (size)
2786 {
2787 unsigned s;
2788 if (size >= 8)
2789 s = 8;
2790 else if (size >= 4)
2791 s = 4;
2792 else if (size >= 2)
2793 s = 2;
2794 else
2795 s = 1;
2796
2797 if (s > min_byte_align)
2798 s = min_byte_align;
2799
2800 addr = new hsa_op_address (target->m_symbol, target->m_reg,
2801 target->m_imm_offset + offset);
2802
2803 BrigType16_t t = get_integer_type_by_bytes (s, false);
2804 HOST_WIDE_INT c = build_memset_value (constant, s);
2805
2806 mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, new hsa_op_immed (c, t),
2807 addr);
2808 hbb->append_insn (mem);
2809 offset += s;
2810 size -= s;
2811 }
2812}
2813
2814/* Generate HSAIL instructions for a single assignment
2815 of an empty constructor to an ADDR_LHS. Constructor is passed as a
2816 tree RHS and all instructions are appended to HBB. ALIGN is
2817 alignment of the address. */
2818
2819void
2820gen_hsa_ctor_assignment (hsa_op_address *addr_lhs, tree rhs, hsa_bb *hbb,
2821 BrigAlignment8_t align)
2822{
2823 if (CONSTRUCTOR_NELTS (rhs))
2824 {
2825 HSA_SORRY_AT (EXPR_LOCATION (rhs),
2826 "support for HSA does not implement load from constructor");
2827 return;
2828 }
2829
2830 unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs)));
2831 gen_hsa_memory_set (hbb, addr_lhs, 0, size, align);
2832}
2833
2834/* Generate HSA instructions for a single assignment of RHS to LHS.
2835 HBB is the basic block they will be appended to. */
2836
2837static void
2838gen_hsa_insns_for_single_assignment (tree lhs, tree rhs, hsa_bb *hbb)
2839{
2840 if (TREE_CODE (lhs) == SSA_NAME)
2841 {
2842 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
2843 if (hsa_seen_error ())
2844 return;
2845
2846 gen_hsa_insns_for_load (dest, rhs, TREE_TYPE (lhs), hbb);
2847 }
2848 else if (TREE_CODE (rhs) == SSA_NAME
2849 || (is_gimple_min_invariant (rhs) && TREE_CODE (rhs) != STRING_CST))
2850 {
2851 /* Store to memory. */
2852 hsa_op_base *src = hsa_reg_or_immed_for_gimple_op (rhs, hbb);
2853 if (hsa_seen_error ())
2854 return;
2855
2856 gen_hsa_insns_for_store (lhs, src, hbb);
2857 }
2858 else
2859 {
2860 BrigAlignment8_t lhs_align;
2861 hsa_op_address *addr_lhs = gen_hsa_addr_with_align (lhs, hbb,
2862 &lhs_align);
2863
2864 if (TREE_CODE (rhs) == CONSTRUCTOR)
2865 gen_hsa_ctor_assignment (addr_lhs, rhs, hbb, lhs_align);
2866 else
2867 {
2868 BrigAlignment8_t rhs_align;
2869 hsa_op_address *addr_rhs = gen_hsa_addr_with_align (rhs, hbb,
2870 &rhs_align);
2871
2872 unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs)));
2873 gen_hsa_memory_copy (hbb, addr_lhs, addr_rhs, size,
2874 MIN (lhs_align, rhs_align));
2875 }
2876 }
2877}
2878
2879/* Prepend before INSN a load from spill symbol of SPILL_REG. Return the
2880 register into which we loaded. If this required another register to convert
2881 from a B1 type, return it in *PTMP2, otherwise store NULL into it. We
2882 assume we are out of SSA so the returned register does not have its
2883 definition set. */
2884
2885hsa_op_reg *
2886hsa_spill_in (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2)
2887{
2888 hsa_symbol *spill_sym = spill_reg->m_spill_sym;
2889 hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type);
2890 hsa_op_address *addr = new hsa_op_address (spill_sym);
2891
2892 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, spill_sym->m_type,
2893 reg, addr);
2894 hsa_insert_insn_before (mem, insn);
2895
2896 *ptmp2 = NULL;
2897 if (spill_reg->m_type == BRIG_TYPE_B1)
2898 {
2899 hsa_insn_basic *cvtinsn;
2900 *ptmp2 = reg;
2901 reg = new hsa_op_reg (spill_reg->m_type);
2902
2903 cvtinsn = new hsa_insn_cvt (reg, *ptmp2);
2904 hsa_insert_insn_before (cvtinsn, insn);
2905 }
2906 return reg;
2907}
2908
2909/* Append after INSN a store to spill symbol of SPILL_REG. Return the register
2910 from which we stored. If this required another register to convert to a B1
2911 type, return it in *PTMP2, otherwise store NULL into it. We assume we are
2912 out of SSA so the returned register does not have its use updated. */
2913
2914hsa_op_reg *
2915hsa_spill_out (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2)
2916{
2917 hsa_symbol *spill_sym = spill_reg->m_spill_sym;
2918 hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type);
2919 hsa_op_address *addr = new hsa_op_address (spill_sym);
2920 hsa_op_reg *returnreg;
2921
2922 *ptmp2 = NULL;
2923 returnreg = reg;
2924 if (spill_reg->m_type == BRIG_TYPE_B1)
2925 {
2926 hsa_insn_basic *cvtinsn;
2927 *ptmp2 = new hsa_op_reg (spill_sym->m_type);
2928 reg->m_type = spill_reg->m_type;
2929
2930 cvtinsn = new hsa_insn_cvt (*ptmp2, returnreg);
2931 hsa_append_insn_after (cvtinsn, insn);
2932 insn = cvtinsn;
2933 reg = *ptmp2;
2934 }
2935
2936 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, spill_sym->m_type, reg,
2937 addr);
2938 hsa_append_insn_after (mem, insn);
2939 return returnreg;
2940}
2941
2942/* Generate a comparison instruction that will compare LHS and RHS with
2943 comparison specified by CODE and put result into register DEST. DEST has to
2944 have its type set already but must not have its definition set yet.
2945 Generated instructions will be added to HBB. */
2946
2947static void
2948gen_hsa_cmp_insn_from_gimple (enum tree_code code, tree lhs, tree rhs,
2949 hsa_op_reg *dest, hsa_bb *hbb)
2950{
2951 BrigCompareOperation8_t compare;
2952
2953 switch (code)
2954 {
2955 case LT_EXPR:
2956 compare = BRIG_COMPARE_LT;
2957 break;
2958 case LE_EXPR:
2959 compare = BRIG_COMPARE_LE;
2960 break;
2961 case GT_EXPR:
2962 compare = BRIG_COMPARE_GT;
2963 break;
2964 case GE_EXPR:
2965 compare = BRIG_COMPARE_GE;
2966 break;
2967 case EQ_EXPR:
2968 compare = BRIG_COMPARE_EQ;
2969 break;
2970 case NE_EXPR:
2971 compare = BRIG_COMPARE_NE;
2972 break;
2973 case UNORDERED_EXPR:
2974 compare = BRIG_COMPARE_NAN;
2975 break;
2976 case ORDERED_EXPR:
2977 compare = BRIG_COMPARE_NUM;
2978 break;
2979 case UNLT_EXPR:
2980 compare = BRIG_COMPARE_LTU;
2981 break;
2982 case UNLE_EXPR:
2983 compare = BRIG_COMPARE_LEU;
2984 break;
2985 case UNGT_EXPR:
2986 compare = BRIG_COMPARE_GTU;
2987 break;
2988 case UNGE_EXPR:
2989 compare = BRIG_COMPARE_GEU;
2990 break;
2991 case UNEQ_EXPR:
2992 compare = BRIG_COMPARE_EQU;
2993 break;
2994 case LTGT_EXPR:
2995 compare = BRIG_COMPARE_NEU;
2996 break;
2997
2998 default:
2999 HSA_SORRY_ATV (EXPR_LOCATION (lhs),
3000 "support for HSA does not implement comparison tree "
3001 "code %s\n", get_tree_code_name (code));
3002 return;
3003 }
3004
3005 /* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer)
3006 as a result of comparison. */
3007
3008 BrigType16_t dest_type = hsa_type_integer_p (dest->m_type)
3009 ? (BrigType16_t) BRIG_TYPE_B1 : dest->m_type;
3010
3011 hsa_insn_cmp *cmp = new hsa_insn_cmp (compare, dest_type);
3012 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (lhs, hbb);
3013 cmp->set_op (1, op1->extend_int_to_32bit (hbb));
3014 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs, hbb);
3015 cmp->set_op (2, op2->extend_int_to_32bit (hbb));
3016
3017 hbb->append_insn (cmp);
3018 cmp->set_output_in_type (dest, 0, hbb);
3019}
3020
3021/* Generate an unary instruction with OPCODE and append it to a basic block
3022 HBB. The instruction uses DEST as a destination and OP1
3023 as a single operand. */
3024
3025static void
3026gen_hsa_unary_operation (BrigOpcode opcode, hsa_op_reg *dest,
3027 hsa_op_with_type *op1, hsa_bb *hbb)
3028{
3029 gcc_checking_assert (dest);
3030 hsa_insn_basic *insn;
3031
3032 if (opcode == BRIG_OPCODE_MOV && hsa_needs_cvt (dest->m_type, op1->m_type))
3033 {
3034 insn = new hsa_insn_cvt (dest, op1);
3035 hbb->append_insn (insn);
3036 return;
3037 }
3038
3039 op1 = op1->extend_int_to_32bit (hbb);
3040 if (opcode == BRIG_OPCODE_FIRSTBIT || opcode == BRIG_OPCODE_LASTBIT)
3041 {
3042 BrigType16_t srctype = hsa_type_integer_p (op1->m_type) ? op1->m_type
3043 : hsa_unsigned_type_for_type (op1->m_type);
3044 insn = new hsa_insn_srctype (2, opcode, BRIG_TYPE_U32, srctype, NULL,
3045 op1);
3046 }
3047 else
3048 {
3049 BrigType16_t optype = hsa_extend_inttype_to_32bit (dest->m_type);
3050 insn = new hsa_insn_basic (2, opcode, optype, NULL, op1);
3051
3052 if (opcode == BRIG_OPCODE_MOV)
3053 hsa_fixup_mov_insn_type (insn);
3054 else if (opcode == BRIG_OPCODE_ABS || opcode == BRIG_OPCODE_NEG)
3055 {
3056 /* ABS and NEG only exist in _s form :-/ */
3057 if (insn->m_type == BRIG_TYPE_U32)
3058 insn->m_type = BRIG_TYPE_S32;
3059 else if (insn->m_type == BRIG_TYPE_U64)
3060 insn->m_type = BRIG_TYPE_S64;
3061 }
3062 }
3063
3064 hbb->append_insn (insn);
3065 insn->set_output_in_type (dest, 0, hbb);
3066}
3067
3068/* Generate a binary instruction with OPCODE and append it to a basic block
3069 HBB. The instruction uses DEST as a destination and operands OP1
3070 and OP2. */
3071
3072static void
3073gen_hsa_binary_operation (int opcode, hsa_op_reg *dest,
3074 hsa_op_with_type *op1, hsa_op_with_type *op2,
3075 hsa_bb *hbb)
3076{
3077 gcc_checking_assert (dest);
3078
3079 BrigType16_t optype = hsa_extend_inttype_to_32bit (dest->m_type);
3080 op1 = op1->extend_int_to_32bit (hbb);
3081 op2 = op2->extend_int_to_32bit (hbb);
3082
3083 if ((opcode == BRIG_OPCODE_SHL || opcode == BRIG_OPCODE_SHR)
3084 && is_a <hsa_op_immed *> (op2))
3085 {
3086 hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2);
3087 i->set_type (BRIG_TYPE_U32);
3088 }
3089 if ((opcode == BRIG_OPCODE_OR
3090 || opcode == BRIG_OPCODE_XOR
3091 || opcode == BRIG_OPCODE_AND)
3092 && is_a <hsa_op_immed *> (op2))
3093 {
3094 hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2);
3095 i->set_type (hsa_unsigned_type_for_type (i->m_type));
3096 }
3097
3098 hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, optype, NULL,
3099 op1, op2);
3100 hbb->append_insn (insn);
3101 insn->set_output_in_type (dest, 0, hbb);
3102}
3103
3104/* Generate HSA instructions for a single assignment. HBB is the basic block
3105 they will be appended to. */
3106
3107static void
3108gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
3109{
3110 tree_code code = gimple_assign_rhs_code (assign);
3111 gimple_rhs_class rhs_class = get_gimple_rhs_class (gimple_expr_code (assign));
3112
3113 tree lhs = gimple_assign_lhs (assign);
3114 tree rhs1 = gimple_assign_rhs1 (assign);
3115 tree rhs2 = gimple_assign_rhs2 (assign);
3116 tree rhs3 = gimple_assign_rhs3 (assign);
3117
3118 BrigOpcode opcode;
3119
3120 switch (code)
3121 {
3122 CASE_CONVERT:
3123 case FLOAT_EXPR:
3124 /* The opcode is changed to BRIG_OPCODE_CVT if BRIG types
3125 needs a conversion. */
3126 opcode = BRIG_OPCODE_MOV;
3127 break;
3128
3129 case PLUS_EXPR:
3130 case POINTER_PLUS_EXPR:
3131 opcode = BRIG_OPCODE_ADD;
3132 break;
3133 case MINUS_EXPR:
3134 opcode = BRIG_OPCODE_SUB;
3135 break;
3136 case MULT_EXPR:
3137 opcode = BRIG_OPCODE_MUL;
3138 break;
3139 case MULT_HIGHPART_EXPR:
3140 opcode = BRIG_OPCODE_MULHI;
3141 break;
3142 case RDIV_EXPR:
3143 case TRUNC_DIV_EXPR:
3144 case EXACT_DIV_EXPR:
3145 opcode = BRIG_OPCODE_DIV;
3146 break;
3147 case CEIL_DIV_EXPR:
3148 case FLOOR_DIV_EXPR:
3149 case ROUND_DIV_EXPR:
3150 HSA_SORRY_AT (gimple_location (assign),
3151 "support for HSA does not implement CEIL_DIV_EXPR, "
3152 "FLOOR_DIV_EXPR or ROUND_DIV_EXPR");
3153 return;
3154 case TRUNC_MOD_EXPR:
3155 opcode = BRIG_OPCODE_REM;
3156 break;
3157 case CEIL_MOD_EXPR:
3158 case FLOOR_MOD_EXPR:
3159 case ROUND_MOD_EXPR:
3160 HSA_SORRY_AT (gimple_location (assign),
3161 "support for HSA does not implement CEIL_MOD_EXPR, "
3162 "FLOOR_MOD_EXPR or ROUND_MOD_EXPR");
3163 return;
3164 case NEGATE_EXPR:
3165 opcode = BRIG_OPCODE_NEG;
3166 break;
3167 case FMA_EXPR:
3168 /* There is a native HSA instruction for scalar FMAs but not for vector
3169 ones. */
3170 if (TREE_CODE (TREE_TYPE (lhs)) == VECTOR_TYPE)
3171 {
3172 hsa_op_reg *dest
3173 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3174 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3175 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3176 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
3177 hsa_op_reg *tmp = new hsa_op_reg (dest->m_type);
3178 gen_hsa_binary_operation (BRIG_OPCODE_MUL, tmp, op1, op2, hbb);
3179 gen_hsa_binary_operation (BRIG_OPCODE_ADD, dest, tmp, op3, hbb);
3180 return;
3181 }
3182 opcode = BRIG_OPCODE_MAD;
3183 break;
3184 case MIN_EXPR:
3185 opcode = BRIG_OPCODE_MIN;
3186 break;
3187 case MAX_EXPR:
3188 opcode = BRIG_OPCODE_MAX;
3189 break;
3190 case ABS_EXPR:
3191 opcode = BRIG_OPCODE_ABS;
3192 break;
3193 case LSHIFT_EXPR:
3194 opcode = BRIG_OPCODE_SHL;
3195 break;
3196 case RSHIFT_EXPR:
3197 opcode = BRIG_OPCODE_SHR;
3198 break;
3199 case LROTATE_EXPR:
3200 case RROTATE_EXPR:
3201 {
3202 hsa_insn_basic *insn = NULL;
3203 int code1 = code == LROTATE_EXPR ? BRIG_OPCODE_SHL : BRIG_OPCODE_SHR;
3204 int code2 = code != LROTATE_EXPR ? BRIG_OPCODE_SHL : BRIG_OPCODE_SHR;
3205 BrigType16_t btype = hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
3206 true);
3207
3208 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3209 hsa_op_reg *op1 = new hsa_op_reg (btype);
3210 hsa_op_reg *op2 = new hsa_op_reg (btype);
3211 hsa_op_with_type *shift1 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3212
3213 tree type = TREE_TYPE (rhs2);
3214 unsigned HOST_WIDE_INT bitsize = TREE_INT_CST_LOW (TYPE_SIZE (type));
3215
3216 hsa_op_with_type *shift2 = NULL;
3217 if (TREE_CODE (rhs2) == INTEGER_CST)
3218 shift2 = new hsa_op_immed (bitsize - tree_to_uhwi (rhs2),
3219 BRIG_TYPE_U32);
3220 else if (TREE_CODE (rhs2) == SSA_NAME)
3221 {
3222 hsa_op_reg *s = hsa_cfun->reg_for_gimple_ssa (rhs2);
3223 s = as_a <hsa_op_reg *> (s->extend_int_to_32bit (hbb));
3224 hsa_op_reg *d = new hsa_op_reg (s->m_type);
3225 hsa_op_immed *size_imm = new hsa_op_immed (bitsize, BRIG_TYPE_U32);
3226
3227 insn = new hsa_insn_basic (3, BRIG_OPCODE_SUB, d->m_type,
3228 d, s, size_imm);
3229 hbb->append_insn (insn);
3230
3231 shift2 = d;
3232 }
3233 else
3234 gcc_unreachable ();
3235
3236 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3237 gen_hsa_binary_operation (code1, op1, src, shift1, hbb);
3238 gen_hsa_binary_operation (code2, op2, src, shift2, hbb);
3239 gen_hsa_binary_operation (BRIG_OPCODE_OR, dest, op1, op2, hbb);
3240
3241 return;
3242 }
3243 case BIT_IOR_EXPR:
3244 opcode = BRIG_OPCODE_OR;
3245 break;
3246 case BIT_XOR_EXPR:
3247 opcode = BRIG_OPCODE_XOR;
3248 break;
3249 case BIT_AND_EXPR:
3250 opcode = BRIG_OPCODE_AND;
3251 break;
3252 case BIT_NOT_EXPR:
3253 opcode = BRIG_OPCODE_NOT;
3254 break;
3255 case FIX_TRUNC_EXPR:
3256 {
3257 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3258 hsa_op_with_type *v = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3259
3260 if (hsa_needs_cvt (dest->m_type, v->m_type))
3261 {
3262 hsa_op_reg *tmp = new hsa_op_reg (v->m_type);
3263
3264 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_TRUNC,
3265 tmp->m_type, tmp, v);
3266 hbb->append_insn (insn);
3267
3268 hsa_insn_basic *cvtinsn = new hsa_insn_cvt (dest, tmp);
3269 hbb->append_insn (cvtinsn);
3270 }
3271 else
3272 {
3273 hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_TRUNC,
3274 dest->m_type, dest, v);
3275 hbb->append_insn (insn);
3276 }
3277
3278 return;
3279 }
3280 opcode = BRIG_OPCODE_TRUNC;
3281 break;
3282
3283 case LT_EXPR:
3284 case LE_EXPR:
3285 case GT_EXPR:
3286 case GE_EXPR:
3287 case EQ_EXPR:
3288 case NE_EXPR:
3289 case UNORDERED_EXPR:
3290 case ORDERED_EXPR:
3291 case UNLT_EXPR:
3292 case UNLE_EXPR:
3293 case UNGT_EXPR:
3294 case UNGE_EXPR:
3295 case UNEQ_EXPR:
3296 case LTGT_EXPR:
3297 {
3298 hsa_op_reg *dest
3299 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3300
3301 gen_hsa_cmp_insn_from_gimple (code, rhs1, rhs2, dest, hbb);
3302 return;
3303 }
3304 case COND_EXPR:
3305 {
3306 hsa_op_reg *dest
3307 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3308 hsa_op_with_type *ctrl = NULL;
3309 tree cond = rhs1;
3310
3311 if (CONSTANT_CLASS_P (cond) || TREE_CODE (cond) == SSA_NAME)
3312 ctrl = hsa_reg_or_immed_for_gimple_op (cond, hbb);
3313 else
3314 {
3315 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_B1);
3316
3317 gen_hsa_cmp_insn_from_gimple (TREE_CODE (cond),
3318 TREE_OPERAND (cond, 0),
3319 TREE_OPERAND (cond, 1),
3320 r, hbb);
3321
3322 ctrl = r;
3323 }
3324
3325 hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3326 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
3327 op2 = op2->extend_int_to_32bit (hbb);
3328 op3 = op3->extend_int_to_32bit (hbb);
3329
3330 BrigType16_t type = hsa_extend_inttype_to_32bit (dest->m_type);
3331 BrigType16_t utype = hsa_unsigned_type_for_type (type);
3332 if (is_a <hsa_op_immed *> (op2))
3333 op2->m_type = utype;
3334 if (is_a <hsa_op_immed *> (op3))
3335 op3->m_type = utype;
3336
3337 hsa_insn_basic *insn
3338 = new hsa_insn_basic (4, BRIG_OPCODE_CMOV,
3339 hsa_bittype_for_type (type),
3340 NULL, ctrl, op2, op3);
3341
3342 hbb->append_insn (insn);
3343 insn->set_output_in_type (dest, 0, hbb);
3344 return;
3345 }
3346 case COMPLEX_EXPR:
3347 {
3348 hsa_op_reg *dest
3349 = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
3350 hsa_op_with_type *rhs1_reg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3351 rhs1_reg = rhs1_reg->extend_int_to_32bit (hbb);
3352 hsa_op_with_type *rhs2_reg = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
3353 rhs2_reg = rhs2_reg->extend_int_to_32bit (hbb);
3354
3355 if (hsa_seen_error ())
3356 return;
3357
3358 BrigType16_t src_type = hsa_bittype_for_type (rhs1_reg->m_type);
3359 rhs1_reg = rhs1_reg->get_in_type (src_type, hbb);
3360 rhs2_reg = rhs2_reg->get_in_type (src_type, hbb);
3361
3362 hsa_insn_packed *insn
3363 = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type, src_type,
3364 dest, rhs1_reg, rhs2_reg);
3365 hbb->append_insn (insn);
3366
3367 return;
3368 }
3369 default:
3370 /* Implement others as we come across them. */
3371 HSA_SORRY_ATV (gimple_location (assign),
3372 "support for HSA does not implement operation %s",
3373 get_tree_code_name (code));
3374 return;
3375 }
3376
3377
3378 hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
3379 hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
3380 hsa_op_with_type *op2
3381 = rhs2 ? hsa_reg_or_immed_for_gimple_op (rhs2, hbb) : NULL;
3382
3383 if (hsa_seen_error ())
3384 return;
3385
3386 switch (rhs_class)
3387 {
3388 case GIMPLE_TERNARY_RHS:
3389 {
3390 hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
3391 op3 = op3->extend_int_to_32bit (hbb);
3392 hsa_insn_basic *insn = new hsa_insn_basic (4, opcode, dest->m_type, dest,
3393 op1, op2, op3);
3394 hbb->append_insn (insn);
3395 }
3396 return;
3397
3398 case GIMPLE_BINARY_RHS:
3399 gen_hsa_binary_operation (opcode, dest, op1, op2, hbb);
3400 break;
3401
3402 case GIMPLE_UNARY_RHS:
3403 gen_hsa_unary_operation (opcode, dest, op1, hbb);
3404 break;
3405 default:
3406 gcc_unreachable ();
3407 }
3408}
3409
3410/* Generate HSA instructions for a given gimple condition statement COND.
3411 Instructions will be appended to HBB, which also needs to be the
3412 corresponding structure to the basic_block of COND. */
3413
3414static void
3415gen_hsa_insns_for_cond_stmt (gimple *cond, hsa_bb *hbb)
3416{
3417 hsa_op_reg *ctrl = new hsa_op_reg (BRIG_TYPE_B1);
3418 hsa_insn_cbr *cbr;
3419
3420 gen_hsa_cmp_insn_from_gimple (gimple_cond_code (cond),
3421 gimple_cond_lhs (cond),
3422 gimple_cond_rhs (cond),
3423 ctrl, hbb);
3424
3425 cbr = new hsa_insn_cbr (ctrl);
3426 hbb->append_insn (cbr);
3427}
3428
3429/* Maximum number of elements in a jump table for an HSA SBR instruction. */
3430
3431#define HSA_MAXIMUM_SBR_LABELS 16
3432
3433/* Return lowest value of a switch S that is handled in a non-default
3434 label. */
3435
3436static tree
3437get_switch_low (gswitch *s)
3438{
3439 unsigned labels = gimple_switch_num_labels (s);
3440 gcc_checking_assert (labels >= 1);
3441
3442 return CASE_LOW (gimple_switch_label (s, 1));
3443}
3444
3445/* Return highest value of a switch S that is handled in a non-default
3446 label. */
3447
3448static tree
3449get_switch_high (gswitch *s)
3450{
3451 unsigned labels = gimple_switch_num_labels (s);
3452
3453 /* Compare last label to maximum number of labels. */
3454 tree label = gimple_switch_label (s, labels - 1);
3455 tree low = CASE_LOW (label);
3456 tree high = CASE_HIGH (label);
3457
3458 return high != NULL_TREE ? high : low;
3459}
3460
3461static tree
3462get_switch_size (gswitch *s)
3463{
3464 return int_const_binop (MINUS_EXPR, get_switch_high (s), get_switch_low (s));
3465}
3466
3467/* Generate HSA instructions for a given gimple switch.
3468 Instructions will be appended to HBB. */
3469
3470static void
3471gen_hsa_insns_for_switch_stmt (gswitch *s, hsa_bb *hbb)
3472{
3473 gimple_stmt_iterator it = gsi_for_stmt (s);
3474 gsi_prev (&it);
3475
3476 /* Create preambule that verifies that index - lowest_label >= 0. */
3477 edge e = split_block (hbb->m_bb, gsi_stmt (it));
3478 e->flags &= ~EDGE_FALLTHRU;
3479 e->flags |= EDGE_TRUE_VALUE;
3480
3481 function *func = DECL_STRUCT_FUNCTION (current_function_decl);
3482 tree index_tree = gimple_switch_index (s);
3483 tree lowest = get_switch_low (s);
3484 tree highest = get_switch_high (s);
3485
3486 hsa_op_reg *index = hsa_cfun->reg_for_gimple_ssa (index_tree);
3487 index = as_a <hsa_op_reg *> (index->extend_int_to_32bit (hbb));
3488
3489 hsa_op_reg *cmp1_reg = new hsa_op_reg (BRIG_TYPE_B1);
3490 hsa_op_immed *cmp1_immed = new hsa_op_immed (lowest, true);
3491 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_GE, cmp1_reg->m_type,
3492 cmp1_reg, index, cmp1_immed));
3493
3494 hsa_op_reg *cmp2_reg = new hsa_op_reg (BRIG_TYPE_B1);
3495 hsa_op_immed *cmp2_immed = new hsa_op_immed (highest, true);
3496 hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_LE, cmp2_reg->m_type,
3497 cmp2_reg, index, cmp2_immed));
3498
3499 hsa_op_reg *cmp_reg = new hsa_op_reg (BRIG_TYPE_B1);
3500 hbb->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_AND, cmp_reg->m_type,
3501 cmp_reg, cmp1_reg, cmp2_reg));
3502
3503 hbb->append_insn (new hsa_insn_cbr (cmp_reg));
3504
3505 tree default_label = gimple_switch_default_label (s);
3506 basic_block default_label_bb = label_to_block_fn (func,
3507 CASE_LABEL (default_label));
3508
3509 if (!gimple_seq_empty_p (phi_nodes (default_label_bb)))
3510 {
3511 default_label_bb = split_edge (find_edge (e->dest, default_label_bb));
3512 hsa_init_new_bb (default_label_bb);
3513 }
3514
3515 make_edge (e->src, default_label_bb, EDGE_FALSE_VALUE);
3516
3517 hsa_cfun->m_modified_cfg = true;
3518
3519 /* Basic block with the SBR instruction. */
3520 hbb = hsa_init_new_bb (e->dest);
3521
3522 hsa_op_reg *sub_index = new hsa_op_reg (index->m_type);
3523 hbb->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB, sub_index->m_type,
3524 sub_index, index,
3525 new hsa_op_immed (lowest, true)));
3526
3527 hsa_op_base *tmp = sub_index->get_in_type (BRIG_TYPE_U64, hbb);
3528 sub_index = as_a <hsa_op_reg *> (tmp);
3529 unsigned labels = gimple_switch_num_labels (s);
3530 unsigned HOST_WIDE_INT size = tree_to_uhwi (get_switch_size (s));
3531
3532 hsa_insn_sbr *sbr = new hsa_insn_sbr (sub_index, size + 1);
3533
3534 /* Prepare array with default label destination. */
3535 for (unsigned HOST_WIDE_INT i = 0; i <= size; i++)
3536 sbr->m_jump_table.safe_push (default_label_bb);
3537
3538 /* Iterate all labels and fill up the jump table. */
3539 for (unsigned i = 1; i < labels; i++)
3540 {
3541 tree label = gimple_switch_label (s, i);
3542 basic_block bb = label_to_block_fn (func, CASE_LABEL (label));
3543
3544 unsigned HOST_WIDE_INT sub_low
3545 = tree_to_uhwi (int_const_binop (MINUS_EXPR, CASE_LOW (label), lowest));
3546
3547 unsigned HOST_WIDE_INT sub_high = sub_low;
3548 tree high = CASE_HIGH (label);
3549 if (high != NULL)
3550 sub_high = tree_to_uhwi (int_const_binop (MINUS_EXPR, high, lowest));
3551
3552 for (unsigned HOST_WIDE_INT j = sub_low; j <= sub_high; j++)
3553 sbr->m_jump_table[j] = bb;
3554 }
3555
3556 hbb->append_insn (sbr);
3557}
3558
3559/* Verify that the function DECL can be handled by HSA. */
3560
3561static void
3562verify_function_arguments (tree decl)
3563{
3564 tree type = TREE_TYPE (decl);
3565 if (DECL_STATIC_CHAIN (decl))
3566 {
3567 HSA_SORRY_ATV (EXPR_LOCATION (decl),
3568 "HSA does not support nested functions: %qD", decl);
3569 return;
3570 }
3571 else if (!TYPE_ARG_TYPES (type) || stdarg_p (type))
3572 {
3573 HSA_SORRY_ATV (EXPR_LOCATION (decl),
3574 "HSA does not support functions with variadic arguments "
3575 "(or unknown return type): %qD", decl);
3576 return;
3577 }
3578}
3579
3580/* Return BRIG type for FORMAL_ARG_TYPE. If the formal argument type is NULL,
3581 return ACTUAL_ARG_TYPE. */
3582
3583static BrigType16_t
3584get_format_argument_type (tree formal_arg_type, BrigType16_t actual_arg_type)
3585{
3586 if (formal_arg_type == NULL)
3587 return actual_arg_type;
3588
3589 BrigType16_t decl_type
3590 = hsa_type_for_scalar_tree_type (formal_arg_type, false);
3591 return mem_type_for_type (decl_type);
3592}
3593
3594/* Generate HSA instructions for a direct call instruction.
3595 Instructions will be appended to HBB, which also needs to be the
3596 corresponding structure to the basic_block of STMT.
3597 If ASSIGN_LHS is false, do not copy HSA function result argument into the
3598 corresponding HSA representation of the gimple statement LHS. */
3599
3600static void
3601gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb,
3602 bool assign_lhs = true)
3603{
3604 tree decl = gimple_call_fndecl (stmt);
3605 verify_function_arguments (decl);
3606 if (hsa_seen_error ())
3607 return;
3608
3609 hsa_insn_call *call_insn = new hsa_insn_call (decl);
3610 hsa_cfun->m_called_functions.safe_push (call_insn->m_called_function);
3611
3612 /* Argument block start. */
3613 hsa_insn_arg_block *arg_start
3614 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn);
3615 hbb->append_insn (arg_start);
3616
3617 tree parm_type_chain = TYPE_ARG_TYPES (gimple_call_fntype (stmt));
3618
3619 /* Preparation of arguments that will be passed to function. */
3620 const unsigned args = gimple_call_num_args (stmt);
3621 for (unsigned i = 0; i < args; ++i)
3622 {
3623 tree parm = gimple_call_arg (stmt, (int)i);
3624 tree parm_decl_type = parm_type_chain != NULL_TREE
3625 ? TREE_VALUE (parm_type_chain) : NULL_TREE;
3626 hsa_op_address *addr;
3627
3628 if (AGGREGATE_TYPE_P (TREE_TYPE (parm)))
3629 {
3630 addr = gen_hsa_addr_for_arg (TREE_TYPE (parm), i);
3631 BrigAlignment8_t align;
3632 hsa_op_address *src = gen_hsa_addr_with_align (parm, hbb, &align);
3633 gen_hsa_memory_copy (hbb, addr, src,
3634 addr->m_symbol->total_byte_size (), align);
3635 }
3636 else
3637 {
3638 hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (parm, hbb);
3639
3640 if (parm_decl_type != NULL && AGGREGATE_TYPE_P (parm_decl_type))
3641 {
3642 HSA_SORRY_AT (gimple_location (stmt),
3643 "support for HSA does not implement an aggregate "
3644 "formal argument in a function call, while actual "
3645 "argument is not an aggregate");
3646 return;
3647 }
3648
3649 BrigType16_t formal_arg_type
3650 = get_format_argument_type (parm_decl_type, src->m_type);
3651 if (hsa_seen_error ())
3652 return;
3653
3654 if (src->m_type != formal_arg_type)
3655 src = src->get_in_type (formal_arg_type, hbb);
3656
3657 addr
3658 = gen_hsa_addr_for_arg (parm_decl_type != NULL_TREE ?
3659 parm_decl_type: TREE_TYPE (parm), i);
3660 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, formal_arg_type,
3661 src, addr);
3662
3663 hbb->append_insn (mem);
3664 }
3665
3666 call_insn->m_input_args.safe_push (addr->m_symbol);
3667 if (parm_type_chain)
3668 parm_type_chain = TREE_CHAIN (parm_type_chain);
3669 }
3670
3671 call_insn->m_args_code_list = new hsa_op_code_list (args);
3672 hbb->append_insn (call_insn);
3673
3674 tree result_type = TREE_TYPE (TREE_TYPE (decl));
3675
3676 tree result = gimple_call_lhs (stmt);
3677 hsa_insn_mem *result_insn = NULL;
3678 if (!VOID_TYPE_P (result_type))
3679 {
3680 hsa_op_address *addr = gen_hsa_addr_for_arg (result_type, -1);
3681
3682 /* Even if result of a function call is unused, we have to emit
3683 declaration for the result. */
3684 if (result && assign_lhs)
3685 {
3686 tree lhs_type = TREE_TYPE (result);
3687
3688 if (hsa_seen_error ())
3689 return;
3690
3691 if (AGGREGATE_TYPE_P (lhs_type))
3692 {
3693 BrigAlignment8_t align;
3694 hsa_op_address *result_addr
3695 = gen_hsa_addr_with_align (result, hbb, &align);
3696 gen_hsa_memory_copy (hbb, result_addr, addr,
3697 addr->m_symbol->total_byte_size (), align);
3698 }
3699 else
3700 {
3701 BrigType16_t mtype
3702 = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type,
3703 false));
3704
3705 hsa_op_reg *dst = hsa_cfun->reg_for_gimple_ssa (result);
3706 result_insn = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, dst, addr);
3707 hbb->append_insn (result_insn);
3708 }
3709 }
3710
3711 call_insn->m_output_arg = addr->m_symbol;
3712 call_insn->m_result_code_list = new hsa_op_code_list (1);
3713 }
3714 else
3715 {
3716 if (result)
3717 {
3718 HSA_SORRY_AT (gimple_location (stmt),
3719 "support for HSA does not implement an assignment of "
3720 "return value from a void function");
3721 return;
3722 }
3723
3724 call_insn->m_result_code_list = new hsa_op_code_list (0);
3725 }
3726
3727 /* Argument block end. */
3728 hsa_insn_arg_block *arg_end
3729 = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn);
3730 hbb->append_insn (arg_end);
3731}
3732
3733/* Generate HSA instructions for a direct call of an internal fn.
3734 Instructions will be appended to HBB, which also needs to be the
3735 corresponding structure to the ba