annotate gcc/hsa-gen.c @ 131:84e7813d76e9

gcc-8.2
author mir3636
date Thu, 25 Oct 2018 07:37:49 +0900
parents 04ced10e8804
children 1830386684a0
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
111
kono
parents:
diff changeset
1 /* A pass for lowering gimple to HSAIL
131
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
2 Copyright (C) 2013-2018 Free Software Foundation, Inc.
111
kono
parents:
diff changeset
3 Contributed by Martin Jambor <mjambor@suse.cz> and
kono
parents:
diff changeset
4 Martin Liska <mliska@suse.cz>.
kono
parents:
diff changeset
5
kono
parents:
diff changeset
6 This file is part of GCC.
kono
parents:
diff changeset
7
kono
parents:
diff changeset
8 GCC is free software; you can redistribute it and/or modify
kono
parents:
diff changeset
9 it under the terms of the GNU General Public License as published by
kono
parents:
diff changeset
10 the Free Software Foundation; either version 3, or (at your option)
kono
parents:
diff changeset
11 any later version.
kono
parents:
diff changeset
12
kono
parents:
diff changeset
13 GCC is distributed in the hope that it will be useful,
kono
parents:
diff changeset
14 but WITHOUT ANY WARRANTY; without even the implied warranty of
kono
parents:
diff changeset
15 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
kono
parents:
diff changeset
16 GNU General Public License for more details.
kono
parents:
diff changeset
17
kono
parents:
diff changeset
18 You should have received a copy of the GNU General Public License
kono
parents:
diff changeset
19 along with GCC; see the file COPYING3. If not see
kono
parents:
diff changeset
20 <http://www.gnu.org/licenses/>. */
kono
parents:
diff changeset
21
kono
parents:
diff changeset
22 #include "config.h"
kono
parents:
diff changeset
23 #include "system.h"
kono
parents:
diff changeset
24 #include "coretypes.h"
kono
parents:
diff changeset
25 #include "memmodel.h"
kono
parents:
diff changeset
26 #include "tm.h"
kono
parents:
diff changeset
27 #include "is-a.h"
kono
parents:
diff changeset
28 #include "hash-table.h"
kono
parents:
diff changeset
29 #include "vec.h"
kono
parents:
diff changeset
30 #include "tree.h"
kono
parents:
diff changeset
31 #include "tree-pass.h"
kono
parents:
diff changeset
32 #include "function.h"
kono
parents:
diff changeset
33 #include "basic-block.h"
kono
parents:
diff changeset
34 #include "cfg.h"
kono
parents:
diff changeset
35 #include "fold-const.h"
kono
parents:
diff changeset
36 #include "gimple.h"
kono
parents:
diff changeset
37 #include "gimple-iterator.h"
kono
parents:
diff changeset
38 #include "bitmap.h"
kono
parents:
diff changeset
39 #include "dumpfile.h"
kono
parents:
diff changeset
40 #include "gimple-pretty-print.h"
kono
parents:
diff changeset
41 #include "diagnostic-core.h"
kono
parents:
diff changeset
42 #include "gimple-ssa.h"
kono
parents:
diff changeset
43 #include "tree-phinodes.h"
kono
parents:
diff changeset
44 #include "stringpool.h"
kono
parents:
diff changeset
45 #include "tree-vrp.h"
kono
parents:
diff changeset
46 #include "tree-ssanames.h"
kono
parents:
diff changeset
47 #include "tree-dfa.h"
kono
parents:
diff changeset
48 #include "ssa-iterators.h"
kono
parents:
diff changeset
49 #include "cgraph.h"
kono
parents:
diff changeset
50 #include "print-tree.h"
kono
parents:
diff changeset
51 #include "symbol-summary.h"
kono
parents:
diff changeset
52 #include "hsa-common.h"
kono
parents:
diff changeset
53 #include "cfghooks.h"
kono
parents:
diff changeset
54 #include "tree-cfg.h"
kono
parents:
diff changeset
55 #include "cfgloop.h"
kono
parents:
diff changeset
56 #include "cfganal.h"
kono
parents:
diff changeset
57 #include "builtins.h"
kono
parents:
diff changeset
58 #include "params.h"
kono
parents:
diff changeset
59 #include "gomp-constants.h"
kono
parents:
diff changeset
60 #include "internal-fn.h"
kono
parents:
diff changeset
61 #include "builtins.h"
kono
parents:
diff changeset
62 #include "stor-layout.h"
kono
parents:
diff changeset
63 #include "stringpool.h"
kono
parents:
diff changeset
64 #include "attribs.h"
kono
parents:
diff changeset
65
kono
parents:
diff changeset
66 /* Print a warning message and set that we have seen an error. */
kono
parents:
diff changeset
67
kono
parents:
diff changeset
68 #define HSA_SORRY_ATV(location, message, ...) \
kono
parents:
diff changeset
69 do \
kono
parents:
diff changeset
70 { \
kono
parents:
diff changeset
71 hsa_fail_cfun (); \
131
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
72 auto_diagnostic_group d; \
111
kono
parents:
diff changeset
73 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
kono
parents:
diff changeset
74 HSA_SORRY_MSG)) \
kono
parents:
diff changeset
75 inform (location, message, __VA_ARGS__); \
kono
parents:
diff changeset
76 } \
kono
parents:
diff changeset
77 while (false)
kono
parents:
diff changeset
78
kono
parents:
diff changeset
79 /* Same as previous, but highlight a location. */
kono
parents:
diff changeset
80
kono
parents:
diff changeset
81 #define HSA_SORRY_AT(location, message) \
kono
parents:
diff changeset
82 do \
kono
parents:
diff changeset
83 { \
kono
parents:
diff changeset
84 hsa_fail_cfun (); \
131
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
85 auto_diagnostic_group d; \
111
kono
parents:
diff changeset
86 if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
kono
parents:
diff changeset
87 HSA_SORRY_MSG)) \
kono
parents:
diff changeset
88 inform (location, message); \
kono
parents:
diff changeset
89 } \
kono
parents:
diff changeset
90 while (false)
kono
parents:
diff changeset
91
kono
parents:
diff changeset
92 /* Default number of threads used by kernel dispatch. */
kono
parents:
diff changeset
93
kono
parents:
diff changeset
94 #define HSA_DEFAULT_NUM_THREADS 64
kono
parents:
diff changeset
95
kono
parents:
diff changeset
96 /* Following structures are defined in the final version
kono
parents:
diff changeset
97 of HSA specification. */
kono
parents:
diff changeset
98
kono
parents:
diff changeset
99 /* HSA queue packet is shadow structure, originally provided by AMD. */
kono
parents:
diff changeset
100
kono
parents:
diff changeset
101 struct hsa_queue_packet
kono
parents:
diff changeset
102 {
kono
parents:
diff changeset
103 uint16_t header;
kono
parents:
diff changeset
104 uint16_t setup;
kono
parents:
diff changeset
105 uint16_t workgroup_size_x;
kono
parents:
diff changeset
106 uint16_t workgroup_size_y;
kono
parents:
diff changeset
107 uint16_t workgroup_size_z;
kono
parents:
diff changeset
108 uint16_t reserved0;
kono
parents:
diff changeset
109 uint32_t grid_size_x;
kono
parents:
diff changeset
110 uint32_t grid_size_y;
kono
parents:
diff changeset
111 uint32_t grid_size_z;
kono
parents:
diff changeset
112 uint32_t private_segment_size;
kono
parents:
diff changeset
113 uint32_t group_segment_size;
kono
parents:
diff changeset
114 uint64_t kernel_object;
kono
parents:
diff changeset
115 void *kernarg_address;
kono
parents:
diff changeset
116 uint64_t reserved2;
kono
parents:
diff changeset
117 uint64_t completion_signal;
kono
parents:
diff changeset
118 };
kono
parents:
diff changeset
119
kono
parents:
diff changeset
120 /* HSA queue is shadow structure, originally provided by AMD. */
kono
parents:
diff changeset
121
kono
parents:
diff changeset
122 struct hsa_queue
kono
parents:
diff changeset
123 {
kono
parents:
diff changeset
124 int type;
kono
parents:
diff changeset
125 uint32_t features;
kono
parents:
diff changeset
126 void *base_address;
kono
parents:
diff changeset
127 uint64_t doorbell_signal;
kono
parents:
diff changeset
128 uint32_t size;
kono
parents:
diff changeset
129 uint32_t reserved1;
kono
parents:
diff changeset
130 uint64_t id;
kono
parents:
diff changeset
131 };
kono
parents:
diff changeset
132
kono
parents:
diff changeset
133 static struct obstack hsa_obstack;
kono
parents:
diff changeset
134
kono
parents:
diff changeset
135 /* List of pointers to all instructions that come from an object allocator. */
kono
parents:
diff changeset
136 static vec <hsa_insn_basic *> hsa_instructions;
kono
parents:
diff changeset
137
kono
parents:
diff changeset
138 /* List of pointers to all operands that come from an object allocator. */
kono
parents:
diff changeset
139 static vec <hsa_op_base *> hsa_operands;
kono
parents:
diff changeset
140
kono
parents:
diff changeset
141 hsa_symbol::hsa_symbol ()
kono
parents:
diff changeset
142 : m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
kono
parents:
diff changeset
143 m_directive_offset (0), m_type (BRIG_TYPE_NONE),
kono
parents:
diff changeset
144 m_segment (BRIG_SEGMENT_NONE), m_linkage (BRIG_LINKAGE_NONE), m_dim (0),
kono
parents:
diff changeset
145 m_cst_value (NULL), m_global_scope_p (false), m_seen_error (false),
kono
parents:
diff changeset
146 m_allocation (BRIG_ALLOCATION_AUTOMATIC), m_emitted_to_brig (false)
kono
parents:
diff changeset
147 {
kono
parents:
diff changeset
148 }
kono
parents:
diff changeset
149
kono
parents:
diff changeset
150
kono
parents:
diff changeset
151 hsa_symbol::hsa_symbol (BrigType16_t type, BrigSegment8_t segment,
kono
parents:
diff changeset
152 BrigLinkage8_t linkage, bool global_scope_p,
kono
parents:
diff changeset
153 BrigAllocation allocation, BrigAlignment8_t align)
kono
parents:
diff changeset
154 : m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
kono
parents:
diff changeset
155 m_directive_offset (0), m_type (type), m_segment (segment),
kono
parents:
diff changeset
156 m_linkage (linkage), m_dim (0), m_cst_value (NULL),
kono
parents:
diff changeset
157 m_global_scope_p (global_scope_p), m_seen_error (false),
kono
parents:
diff changeset
158 m_allocation (allocation), m_emitted_to_brig (false), m_align (align)
kono
parents:
diff changeset
159 {
kono
parents:
diff changeset
160 }
kono
parents:
diff changeset
161
kono
parents:
diff changeset
162 unsigned HOST_WIDE_INT
kono
parents:
diff changeset
163 hsa_symbol::total_byte_size ()
kono
parents:
diff changeset
164 {
kono
parents:
diff changeset
165 unsigned HOST_WIDE_INT s
kono
parents:
diff changeset
166 = hsa_type_bit_size (~BRIG_TYPE_ARRAY_MASK & m_type);
kono
parents:
diff changeset
167 gcc_assert (s % BITS_PER_UNIT == 0);
kono
parents:
diff changeset
168 s /= BITS_PER_UNIT;
kono
parents:
diff changeset
169
kono
parents:
diff changeset
170 if (m_dim)
kono
parents:
diff changeset
171 s *= m_dim;
kono
parents:
diff changeset
172
kono
parents:
diff changeset
173 return s;
kono
parents:
diff changeset
174 }
kono
parents:
diff changeset
175
kono
parents:
diff changeset
176 /* Forward declaration. */
kono
parents:
diff changeset
177
kono
parents:
diff changeset
178 static BrigType16_t
kono
parents:
diff changeset
179 hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p,
kono
parents:
diff changeset
180 bool min32int);
kono
parents:
diff changeset
181
kono
parents:
diff changeset
182 void
kono
parents:
diff changeset
183 hsa_symbol::fillup_for_decl (tree decl)
kono
parents:
diff changeset
184 {
kono
parents:
diff changeset
185 m_decl = decl;
kono
parents:
diff changeset
186 m_type = hsa_type_for_tree_type (TREE_TYPE (decl), &m_dim, false);
kono
parents:
diff changeset
187 if (hsa_seen_error ())
kono
parents:
diff changeset
188 {
kono
parents:
diff changeset
189 m_seen_error = true;
kono
parents:
diff changeset
190 return;
kono
parents:
diff changeset
191 }
kono
parents:
diff changeset
192
kono
parents:
diff changeset
193 m_align = MAX (m_align, hsa_natural_alignment (m_type));
kono
parents:
diff changeset
194 }
kono
parents:
diff changeset
195
kono
parents:
diff changeset
196 /* Constructor of class representing global HSA function/kernel information and
kono
parents:
diff changeset
197 state. FNDECL is function declaration, KERNEL_P is true if the function
kono
parents:
diff changeset
198 is going to become a HSA kernel. If the function has body, SSA_NAMES_COUNT
kono
parents:
diff changeset
199 should be set to number of SSA names used in the function.
kono
parents:
diff changeset
200 MODIFIED_CFG is set to true in case we modified control-flow graph
kono
parents:
diff changeset
201 of the function. */
kono
parents:
diff changeset
202
kono
parents:
diff changeset
203 hsa_function_representation::hsa_function_representation
kono
parents:
diff changeset
204 (tree fdecl, bool kernel_p, unsigned ssa_names_count, bool modified_cfg)
kono
parents:
diff changeset
205 : m_name (NULL),
kono
parents:
diff changeset
206 m_reg_count (0), m_input_args (vNULL),
kono
parents:
diff changeset
207 m_output_arg (NULL), m_spill_symbols (vNULL), m_global_symbols (vNULL),
kono
parents:
diff changeset
208 m_private_variables (vNULL), m_called_functions (vNULL),
kono
parents:
diff changeset
209 m_called_internal_fns (vNULL), m_hbb_count (0),
kono
parents:
diff changeset
210 m_in_ssa (true), m_kern_p (kernel_p), m_declaration_p (false),
kono
parents:
diff changeset
211 m_decl (fdecl), m_internal_fn (NULL), m_shadow_reg (NULL),
kono
parents:
diff changeset
212 m_kernel_dispatch_count (0), m_maximum_omp_data_size (0),
kono
parents:
diff changeset
213 m_seen_error (false), m_temp_symbol_count (0), m_ssa_map (),
kono
parents:
diff changeset
214 m_modified_cfg (modified_cfg)
kono
parents:
diff changeset
215 {
131
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
216 int sym_init_len = (vec_safe_length (cfun->local_decls) / 2) + 1;
111
kono
parents:
diff changeset
217 m_local_symbols = new hash_table <hsa_noop_symbol_hasher> (sym_init_len);
kono
parents:
diff changeset
218 m_ssa_map.safe_grow_cleared (ssa_names_count);
kono
parents:
diff changeset
219 }
kono
parents:
diff changeset
220
kono
parents:
diff changeset
221 /* Constructor of class representing HSA function information that
kono
parents:
diff changeset
222 is derived for an internal function. */
kono
parents:
diff changeset
223 hsa_function_representation::hsa_function_representation (hsa_internal_fn *fn)
kono
parents:
diff changeset
224 : m_reg_count (0), m_input_args (vNULL),
kono
parents:
diff changeset
225 m_output_arg (NULL), m_local_symbols (NULL),
kono
parents:
diff changeset
226 m_spill_symbols (vNULL), m_global_symbols (vNULL),
kono
parents:
diff changeset
227 m_private_variables (vNULL), m_called_functions (vNULL),
kono
parents:
diff changeset
228 m_called_internal_fns (vNULL), m_hbb_count (0),
kono
parents:
diff changeset
229 m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL),
kono
parents:
diff changeset
230 m_internal_fn (fn), m_shadow_reg (NULL), m_kernel_dispatch_count (0),
kono
parents:
diff changeset
231 m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
kono
parents:
diff changeset
232 m_ssa_map () {}
kono
parents:
diff changeset
233
kono
parents:
diff changeset
234 /* Destructor of class holding function/kernel-wide information and state. */
kono
parents:
diff changeset
235
kono
parents:
diff changeset
236 hsa_function_representation::~hsa_function_representation ()
kono
parents:
diff changeset
237 {
kono
parents:
diff changeset
238 /* Kernel names are deallocated at the end of BRIG output when deallocating
kono
parents:
diff changeset
239 hsa_decl_kernel_mapping. */
kono
parents:
diff changeset
240 if (!m_kern_p || m_seen_error)
kono
parents:
diff changeset
241 free (m_name);
kono
parents:
diff changeset
242
kono
parents:
diff changeset
243 for (unsigned i = 0; i < m_input_args.length (); i++)
kono
parents:
diff changeset
244 delete m_input_args[i];
kono
parents:
diff changeset
245 m_input_args.release ();
kono
parents:
diff changeset
246
kono
parents:
diff changeset
247 delete m_output_arg;
kono
parents:
diff changeset
248 delete m_local_symbols;
kono
parents:
diff changeset
249
kono
parents:
diff changeset
250 for (unsigned i = 0; i < m_spill_symbols.length (); i++)
kono
parents:
diff changeset
251 delete m_spill_symbols[i];
kono
parents:
diff changeset
252 m_spill_symbols.release ();
kono
parents:
diff changeset
253
kono
parents:
diff changeset
254 hsa_symbol *sym;
kono
parents:
diff changeset
255 for (unsigned i = 0; i < m_global_symbols.iterate (i, &sym); i++)
kono
parents:
diff changeset
256 if (sym->m_linkage != BRIG_ALLOCATION_PROGRAM)
kono
parents:
diff changeset
257 delete sym;
kono
parents:
diff changeset
258 m_global_symbols.release ();
kono
parents:
diff changeset
259
kono
parents:
diff changeset
260 for (unsigned i = 0; i < m_private_variables.length (); i++)
kono
parents:
diff changeset
261 delete m_private_variables[i];
kono
parents:
diff changeset
262 m_private_variables.release ();
kono
parents:
diff changeset
263 m_called_functions.release ();
kono
parents:
diff changeset
264 m_ssa_map.release ();
kono
parents:
diff changeset
265
kono
parents:
diff changeset
266 for (unsigned i = 0; i < m_called_internal_fns.length (); i++)
kono
parents:
diff changeset
267 delete m_called_internal_fns[i];
kono
parents:
diff changeset
268 }
kono
parents:
diff changeset
269
kono
parents:
diff changeset
270 hsa_op_reg *
kono
parents:
diff changeset
271 hsa_function_representation::get_shadow_reg ()
kono
parents:
diff changeset
272 {
kono
parents:
diff changeset
273 /* If we compile a function with kernel dispatch and does not set
kono
parents:
diff changeset
274 an optimization level, the function won't be inlined and
kono
parents:
diff changeset
275 we return NULL. */
kono
parents:
diff changeset
276 if (!m_kern_p)
kono
parents:
diff changeset
277 return NULL;
kono
parents:
diff changeset
278
kono
parents:
diff changeset
279 if (m_shadow_reg)
kono
parents:
diff changeset
280 return m_shadow_reg;
kono
parents:
diff changeset
281
kono
parents:
diff changeset
282 /* Append the shadow argument. */
kono
parents:
diff changeset
283 hsa_symbol *shadow = new hsa_symbol (BRIG_TYPE_U64, BRIG_SEGMENT_KERNARG,
kono
parents:
diff changeset
284 BRIG_LINKAGE_FUNCTION);
kono
parents:
diff changeset
285 m_input_args.safe_push (shadow);
kono
parents:
diff changeset
286 shadow->m_name = "hsa_runtime_shadow";
kono
parents:
diff changeset
287
kono
parents:
diff changeset
288 hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_U64);
kono
parents:
diff changeset
289 hsa_op_address *addr = new hsa_op_address (shadow);
kono
parents:
diff changeset
290
kono
parents:
diff changeset
291 hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, r, addr);
kono
parents:
diff changeset
292 hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun))->append_insn (mem);
kono
parents:
diff changeset
293 m_shadow_reg = r;
kono
parents:
diff changeset
294
kono
parents:
diff changeset
295 return r;
kono
parents:
diff changeset
296 }
kono
parents:
diff changeset
297
kono
parents:
diff changeset
298 bool hsa_function_representation::has_shadow_reg_p ()
kono
parents:
diff changeset
299 {
kono
parents:
diff changeset
300 return m_shadow_reg != NULL;
kono
parents:
diff changeset
301 }
kono
parents:
diff changeset
302
kono
parents:
diff changeset
303 void
kono
parents:
diff changeset
304 hsa_function_representation::init_extra_bbs ()
kono
parents:
diff changeset
305 {
kono
parents:
diff changeset
306 hsa_init_new_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
kono
parents:
diff changeset
307 hsa_init_new_bb (EXIT_BLOCK_PTR_FOR_FN (cfun));
kono
parents:
diff changeset
308 }
kono
parents:
diff changeset
309
kono
parents:
diff changeset
310 void
kono
parents:
diff changeset
311 hsa_function_representation::update_dominance ()
kono
parents:
diff changeset
312 {
kono
parents:
diff changeset
313 if (m_modified_cfg)
kono
parents:
diff changeset
314 {
kono
parents:
diff changeset
315 free_dominance_info (CDI_DOMINATORS);
kono
parents:
diff changeset
316 calculate_dominance_info (CDI_DOMINATORS);
kono
parents:
diff changeset
317 }
kono
parents:
diff changeset
318 }
kono
parents:
diff changeset
319
kono
parents:
diff changeset
320 hsa_symbol *
kono
parents:
diff changeset
321 hsa_function_representation::create_hsa_temporary (BrigType16_t type)
kono
parents:
diff changeset
322 {
kono
parents:
diff changeset
323 hsa_symbol *s = new hsa_symbol (type, BRIG_SEGMENT_PRIVATE,
kono
parents:
diff changeset
324 BRIG_LINKAGE_FUNCTION);
kono
parents:
diff changeset
325 s->m_name_number = m_temp_symbol_count++;
kono
parents:
diff changeset
326
kono
parents:
diff changeset
327 hsa_cfun->m_private_variables.safe_push (s);
kono
parents:
diff changeset
328 return s;
kono
parents:
diff changeset
329 }
kono
parents:
diff changeset
330
kono
parents:
diff changeset
331 BrigLinkage8_t
kono
parents:
diff changeset
332 hsa_function_representation::get_linkage ()
kono
parents:
diff changeset
333 {
kono
parents:
diff changeset
334 if (m_internal_fn)
kono
parents:
diff changeset
335 return BRIG_LINKAGE_PROGRAM;
kono
parents:
diff changeset
336
kono
parents:
diff changeset
337 return m_kern_p || TREE_PUBLIC (m_decl) ?
kono
parents:
diff changeset
338 BRIG_LINKAGE_PROGRAM : BRIG_LINKAGE_MODULE;
kono
parents:
diff changeset
339 }
kono
parents:
diff changeset
340
kono
parents:
diff changeset
341 /* Hash map of simple OMP builtins. */
kono
parents:
diff changeset
342 static hash_map <nofree_string_hash, omp_simple_builtin> *omp_simple_builtins
kono
parents:
diff changeset
343 = NULL;
kono
parents:
diff changeset
344
kono
parents:
diff changeset
345 /* Warning messages for OMP builtins. */
kono
parents:
diff changeset
346
kono
parents:
diff changeset
347 #define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP " \
kono
parents:
diff changeset
348 "lock routines"
kono
parents:
diff changeset
349 #define HSA_WARN_TIMING_ROUTINE "support for HSA does not implement OpenMP " \
kono
parents:
diff changeset
350 "timing routines"
kono
parents:
diff changeset
351 #define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \
kono
parents:
diff changeset
352 "undefined semantics within target regions, support for HSA ignores them"
kono
parents:
diff changeset
353 #define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP " \
kono
parents:
diff changeset
354 "affinity feateres"
kono
parents:
diff changeset
355
kono
parents:
diff changeset
356 /* Initialize hash map with simple OMP builtins. */
kono
parents:
diff changeset
357
kono
parents:
diff changeset
358 static void
kono
parents:
diff changeset
359 hsa_init_simple_builtins ()
kono
parents:
diff changeset
360 {
kono
parents:
diff changeset
361 if (omp_simple_builtins != NULL)
kono
parents:
diff changeset
362 return;
kono
parents:
diff changeset
363
kono
parents:
diff changeset
364 omp_simple_builtins
kono
parents:
diff changeset
365 = new hash_map <nofree_string_hash, omp_simple_builtin> ();
kono
parents:
diff changeset
366
kono
parents:
diff changeset
367 omp_simple_builtin omp_builtins[] =
kono
parents:
diff changeset
368 {
kono
parents:
diff changeset
369 omp_simple_builtin ("omp_get_initial_device", NULL, false,
kono
parents:
diff changeset
370 new hsa_op_immed (GOMP_DEVICE_HOST,
kono
parents:
diff changeset
371 (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
372 omp_simple_builtin ("omp_is_initial_device", NULL, false,
kono
parents:
diff changeset
373 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
374 omp_simple_builtin ("omp_get_dynamic", NULL, false,
kono
parents:
diff changeset
375 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
376 omp_simple_builtin ("omp_set_dynamic", NULL, false, NULL),
kono
parents:
diff changeset
377 omp_simple_builtin ("omp_init_lock", HSA_WARN_LOCK_ROUTINE, true),
kono
parents:
diff changeset
378 omp_simple_builtin ("omp_init_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
kono
parents:
diff changeset
379 true),
kono
parents:
diff changeset
380 omp_simple_builtin ("omp_init_nest_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
kono
parents:
diff changeset
381 true),
kono
parents:
diff changeset
382 omp_simple_builtin ("omp_destroy_lock", HSA_WARN_LOCK_ROUTINE, true),
kono
parents:
diff changeset
383 omp_simple_builtin ("omp_set_lock", HSA_WARN_LOCK_ROUTINE, true),
kono
parents:
diff changeset
384 omp_simple_builtin ("omp_unset_lock", HSA_WARN_LOCK_ROUTINE, true),
kono
parents:
diff changeset
385 omp_simple_builtin ("omp_test_lock", HSA_WARN_LOCK_ROUTINE, true),
kono
parents:
diff changeset
386 omp_simple_builtin ("omp_get_wtime", HSA_WARN_TIMING_ROUTINE, true),
kono
parents:
diff changeset
387 omp_simple_builtin ("omp_get_wtick", HSA_WARN_TIMING_ROUTINE, true),
kono
parents:
diff changeset
388 omp_simple_builtin ("omp_target_alloc", HSA_WARN_MEMORY_ROUTINE, false,
kono
parents:
diff changeset
389 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_U64)),
kono
parents:
diff changeset
390 omp_simple_builtin ("omp_target_free", HSA_WARN_MEMORY_ROUTINE, false),
kono
parents:
diff changeset
391 omp_simple_builtin ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE,
kono
parents:
diff changeset
392 false,
kono
parents:
diff changeset
393 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
394 omp_simple_builtin ("omp_target_memcpy", HSA_WARN_MEMORY_ROUTINE, false,
kono
parents:
diff changeset
395 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
396 omp_simple_builtin ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE,
kono
parents:
diff changeset
397 false,
kono
parents:
diff changeset
398 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
399 omp_simple_builtin ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE,
kono
parents:
diff changeset
400 false,
kono
parents:
diff changeset
401 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
402 omp_simple_builtin ("omp_target_disassociate_ptr",
kono
parents:
diff changeset
403 HSA_WARN_MEMORY_ROUTINE,
kono
parents:
diff changeset
404 false,
kono
parents:
diff changeset
405 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
406 omp_simple_builtin ("omp_set_max_active_levels",
kono
parents:
diff changeset
407 "Support for HSA only allows only one active level, "
kono
parents:
diff changeset
408 "call to omp_set_max_active_levels will be ignored "
kono
parents:
diff changeset
409 "in the generated HSAIL",
kono
parents:
diff changeset
410 false, NULL),
kono
parents:
diff changeset
411 omp_simple_builtin ("omp_get_max_active_levels", NULL, false,
kono
parents:
diff changeset
412 new hsa_op_immed (1, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
413 omp_simple_builtin ("omp_in_final", NULL, false,
kono
parents:
diff changeset
414 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
415 omp_simple_builtin ("omp_get_proc_bind", HSA_WARN_AFFINITY, false,
kono
parents:
diff changeset
416 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
417 omp_simple_builtin ("omp_get_num_places", HSA_WARN_AFFINITY, false,
kono
parents:
diff changeset
418 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
419 omp_simple_builtin ("omp_get_place_num_procs", HSA_WARN_AFFINITY, false,
kono
parents:
diff changeset
420 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
421 omp_simple_builtin ("omp_get_place_proc_ids", HSA_WARN_AFFINITY, false,
kono
parents:
diff changeset
422 NULL),
kono
parents:
diff changeset
423 omp_simple_builtin ("omp_get_place_num", HSA_WARN_AFFINITY, false,
kono
parents:
diff changeset
424 new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
425 omp_simple_builtin ("omp_get_partition_num_places", HSA_WARN_AFFINITY,
kono
parents:
diff changeset
426 false,
kono
parents:
diff changeset
427 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
428 omp_simple_builtin ("omp_get_partition_place_nums", HSA_WARN_AFFINITY,
kono
parents:
diff changeset
429 false, NULL),
kono
parents:
diff changeset
430 omp_simple_builtin ("omp_set_default_device",
kono
parents:
diff changeset
431 "omp_set_default_device has undefined semantics "
kono
parents:
diff changeset
432 "within target regions, support for HSA ignores it",
kono
parents:
diff changeset
433 false, NULL),
kono
parents:
diff changeset
434 omp_simple_builtin ("omp_get_default_device",
kono
parents:
diff changeset
435 "omp_get_default_device has undefined semantics "
kono
parents:
diff changeset
436 "within target regions, support for HSA ignores it",
kono
parents:
diff changeset
437 false,
kono
parents:
diff changeset
438 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
439 omp_simple_builtin ("omp_get_num_devices",
kono
parents:
diff changeset
440 "omp_get_num_devices has undefined semantics "
kono
parents:
diff changeset
441 "within target regions, support for HSA ignores it",
kono
parents:
diff changeset
442 false,
kono
parents:
diff changeset
443 new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
kono
parents:
diff changeset
444 omp_simple_builtin ("omp_get_num_procs", NULL, true, NULL),
kono
parents:
diff changeset
445 omp_simple_builtin ("omp_get_cancellation", NULL, true, NULL),
kono
parents:
diff changeset
446 omp_simple_builtin ("omp_set_nested", NULL, true, NULL),
kono
parents:
diff changeset
447 omp_simple_builtin ("omp_get_nested", NULL, true, NULL),
kono
parents:
diff changeset
448 omp_simple_builtin ("omp_set_schedule", NULL, true, NULL),
kono
parents:
diff changeset
449 omp_simple_builtin ("omp_get_schedule", NULL, true, NULL),
kono
parents:
diff changeset
450 omp_simple_builtin ("omp_get_thread_limit", NULL, true, NULL),
kono
parents:
diff changeset
451 omp_simple_builtin ("omp_get_team_size", NULL, true, NULL),
kono
parents:
diff changeset
452 omp_simple_builtin ("omp_get_ancestor_thread_num", NULL, true, NULL),
kono
parents:
diff changeset
453 omp_simple_builtin ("omp_get_max_task_priority", NULL, true, NULL)
kono
parents:
diff changeset
454 };
kono
parents:
diff changeset
455
kono
parents:
diff changeset
456 unsigned count = sizeof (omp_builtins) / sizeof (omp_simple_builtin);
kono
parents:
diff changeset
457
kono
parents:
diff changeset
458 for (unsigned i = 0; i < count; i++)
kono
parents:
diff changeset
459 omp_simple_builtins->put (omp_builtins[i].m_name, omp_builtins[i]);
kono
parents:
diff changeset
460 }
kono
parents:
diff changeset
461
kono
parents:
diff changeset
462 /* Allocate HSA structures that we need only while generating with this. */
kono
parents:
diff changeset
463
kono
parents:
diff changeset
464 static void
kono
parents:
diff changeset
465 hsa_init_data_for_cfun ()
kono
parents:
diff changeset
466 {
kono
parents:
diff changeset
467 hsa_init_compilation_unit_data ();
kono
parents:
diff changeset
468 gcc_obstack_init (&hsa_obstack);
kono
parents:
diff changeset
469 }
kono
parents:
diff changeset
470
kono
parents:
diff changeset
471 /* Deinitialize HSA subsystem and free all allocated memory. */
kono
parents:
diff changeset
472
kono
parents:
diff changeset
473 static void
kono
parents:
diff changeset
474 hsa_deinit_data_for_cfun (void)
kono
parents:
diff changeset
475 {
kono
parents:
diff changeset
476 basic_block bb;
kono
parents:
diff changeset
477
kono
parents:
diff changeset
478 FOR_ALL_BB_FN (bb, cfun)
kono
parents:
diff changeset
479 if (bb->aux)
kono
parents:
diff changeset
480 {
kono
parents:
diff changeset
481 hsa_bb *hbb = hsa_bb_for_bb (bb);
kono
parents:
diff changeset
482 hbb->~hsa_bb ();
kono
parents:
diff changeset
483 bb->aux = NULL;
kono
parents:
diff changeset
484 }
kono
parents:
diff changeset
485
kono
parents:
diff changeset
486 for (unsigned int i = 0; i < hsa_operands.length (); i++)
kono
parents:
diff changeset
487 hsa_destroy_operand (hsa_operands[i]);
kono
parents:
diff changeset
488
kono
parents:
diff changeset
489 hsa_operands.release ();
kono
parents:
diff changeset
490
kono
parents:
diff changeset
491 for (unsigned i = 0; i < hsa_instructions.length (); i++)
kono
parents:
diff changeset
492 hsa_destroy_insn (hsa_instructions[i]);
kono
parents:
diff changeset
493
kono
parents:
diff changeset
494 hsa_instructions.release ();
kono
parents:
diff changeset
495
kono
parents:
diff changeset
496 if (omp_simple_builtins != NULL)
kono
parents:
diff changeset
497 {
kono
parents:
diff changeset
498 delete omp_simple_builtins;
kono
parents:
diff changeset
499 omp_simple_builtins = NULL;
kono
parents:
diff changeset
500 }
kono
parents:
diff changeset
501
kono
parents:
diff changeset
502 obstack_free (&hsa_obstack, NULL);
kono
parents:
diff changeset
503 delete hsa_cfun;
kono
parents:
diff changeset
504 }
kono
parents:
diff changeset
505
kono
parents:
diff changeset
506 /* Return the type which holds addresses in the given SEGMENT. */
kono
parents:
diff changeset
507
kono
parents:
diff changeset
508 static BrigType16_t
kono
parents:
diff changeset
509 hsa_get_segment_addr_type (BrigSegment8_t segment)
kono
parents:
diff changeset
510 {
kono
parents:
diff changeset
511 switch (segment)
kono
parents:
diff changeset
512 {
kono
parents:
diff changeset
513 case BRIG_SEGMENT_NONE:
kono
parents:
diff changeset
514 gcc_unreachable ();
kono
parents:
diff changeset
515
kono
parents:
diff changeset
516 case BRIG_SEGMENT_FLAT:
kono
parents:
diff changeset
517 case BRIG_SEGMENT_GLOBAL:
kono
parents:
diff changeset
518 case BRIG_SEGMENT_READONLY:
kono
parents:
diff changeset
519 case BRIG_SEGMENT_KERNARG:
kono
parents:
diff changeset
520 return hsa_machine_large_p () ? BRIG_TYPE_U64 : BRIG_TYPE_U32;
kono
parents:
diff changeset
521
kono
parents:
diff changeset
522 case BRIG_SEGMENT_GROUP:
kono
parents:
diff changeset
523 case BRIG_SEGMENT_PRIVATE:
kono
parents:
diff changeset
524 case BRIG_SEGMENT_SPILL:
kono
parents:
diff changeset
525 case BRIG_SEGMENT_ARG:
kono
parents:
diff changeset
526 return BRIG_TYPE_U32;
kono
parents:
diff changeset
527 }
kono
parents:
diff changeset
528 gcc_unreachable ();
kono
parents:
diff changeset
529 }
kono
parents:
diff changeset
530
kono
parents:
diff changeset
531 /* Return integer brig type according to provided SIZE in bytes. If SIGN
kono
parents:
diff changeset
532 is set to true, return signed integer type. */
kono
parents:
diff changeset
533
kono
parents:
diff changeset
534 static BrigType16_t
kono
parents:
diff changeset
535 get_integer_type_by_bytes (unsigned size, bool sign)
kono
parents:
diff changeset
536 {
kono
parents:
diff changeset
537 if (sign)
kono
parents:
diff changeset
538 switch (size)
kono
parents:
diff changeset
539 {
kono
parents:
diff changeset
540 case 1:
kono
parents:
diff changeset
541 return BRIG_TYPE_S8;
kono
parents:
diff changeset
542 case 2:
kono
parents:
diff changeset
543 return BRIG_TYPE_S16;
kono
parents:
diff changeset
544 case 4:
kono
parents:
diff changeset
545 return BRIG_TYPE_S32;
kono
parents:
diff changeset
546 case 8:
kono
parents:
diff changeset
547 return BRIG_TYPE_S64;
kono
parents:
diff changeset
548 default:
kono
parents:
diff changeset
549 break;
kono
parents:
diff changeset
550 }
kono
parents:
diff changeset
551 else
kono
parents:
diff changeset
552 switch (size)
kono
parents:
diff changeset
553 {
kono
parents:
diff changeset
554 case 1:
kono
parents:
diff changeset
555 return BRIG_TYPE_U8;
kono
parents:
diff changeset
556 case 2:
kono
parents:
diff changeset
557 return BRIG_TYPE_U16;
kono
parents:
diff changeset
558 case 4:
kono
parents:
diff changeset
559 return BRIG_TYPE_U32;
kono
parents:
diff changeset
560 case 8:
kono
parents:
diff changeset
561 return BRIG_TYPE_U64;
kono
parents:
diff changeset
562 default:
kono
parents:
diff changeset
563 break;
kono
parents:
diff changeset
564 }
kono
parents:
diff changeset
565
kono
parents:
diff changeset
566 return 0;
kono
parents:
diff changeset
567 }
kono
parents:
diff changeset
568
kono
parents:
diff changeset
569 /* If T points to an integral type smaller than 32 bits, change it to a 32bit
kono
parents:
diff changeset
570 equivalent and return the result. Otherwise just return the result. */
kono
parents:
diff changeset
571
kono
parents:
diff changeset
572 static BrigType16_t
kono
parents:
diff changeset
573 hsa_extend_inttype_to_32bit (BrigType16_t t)
kono
parents:
diff changeset
574 {
kono
parents:
diff changeset
575 if (t == BRIG_TYPE_U8 || t == BRIG_TYPE_U16)
kono
parents:
diff changeset
576 return BRIG_TYPE_U32;
kono
parents:
diff changeset
577 else if (t == BRIG_TYPE_S8 || t == BRIG_TYPE_S16)
kono
parents:
diff changeset
578 return BRIG_TYPE_S32;
kono
parents:
diff changeset
579 return t;
kono
parents:
diff changeset
580 }
kono
parents:
diff changeset
581
kono
parents:
diff changeset
582 /* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
kono
parents:
diff changeset
583 are assumed to use flat addressing. If min32int is true, always expand
kono
parents:
diff changeset
584 integer types to one that has at least 32 bits. */
kono
parents:
diff changeset
585
kono
parents:
diff changeset
586 static BrigType16_t
kono
parents:
diff changeset
587 hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
kono
parents:
diff changeset
588 {
kono
parents:
diff changeset
589 HOST_WIDE_INT bsize;
kono
parents:
diff changeset
590 const_tree base;
kono
parents:
diff changeset
591 BrigType16_t res = BRIG_TYPE_NONE;
kono
parents:
diff changeset
592
kono
parents:
diff changeset
593 gcc_checking_assert (TYPE_P (type));
kono
parents:
diff changeset
594 gcc_checking_assert (!AGGREGATE_TYPE_P (type));
kono
parents:
diff changeset
595 if (POINTER_TYPE_P (type))
kono
parents:
diff changeset
596 return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
kono
parents:
diff changeset
597
kono
parents:
diff changeset
598 if (TREE_CODE (type) == VECTOR_TYPE)
kono
parents:
diff changeset
599 base = TREE_TYPE (type);
kono
parents:
diff changeset
600 else if (TREE_CODE (type) == COMPLEX_TYPE)
kono
parents:
diff changeset
601 {
kono
parents:
diff changeset
602 base = TREE_TYPE (type);
kono
parents:
diff changeset
603 min32int = true;
kono
parents:
diff changeset
604 }
kono
parents:
diff changeset
605 else
kono
parents:
diff changeset
606 base = type;
kono
parents:
diff changeset
607
kono
parents:
diff changeset
608 if (!tree_fits_uhwi_p (TYPE_SIZE (base)))
kono
parents:
diff changeset
609 {
kono
parents:
diff changeset
610 HSA_SORRY_ATV (EXPR_LOCATION (type),
kono
parents:
diff changeset
611 "support for HSA does not implement huge or "
kono
parents:
diff changeset
612 "variable-sized type %qT", type);
kono
parents:
diff changeset
613 return res;
kono
parents:
diff changeset
614 }
kono
parents:
diff changeset
615
kono
parents:
diff changeset
616 bsize = tree_to_uhwi (TYPE_SIZE (base));
kono
parents:
diff changeset
617 unsigned byte_size = bsize / BITS_PER_UNIT;
kono
parents:
diff changeset
618 if (INTEGRAL_TYPE_P (base))
kono
parents:
diff changeset
619 res = get_integer_type_by_bytes (byte_size, !TYPE_UNSIGNED (base));
kono
parents:
diff changeset
620 else if (SCALAR_FLOAT_TYPE_P (base))
kono
parents:
diff changeset
621 {
kono
parents:
diff changeset
622 switch (bsize)
kono
parents:
diff changeset
623 {
kono
parents:
diff changeset
624 case 16:
kono
parents:
diff changeset
625 res = BRIG_TYPE_F16;
kono
parents:
diff changeset
626 break;
kono
parents:
diff changeset
627 case 32:
kono
parents:
diff changeset
628 res = BRIG_TYPE_F32;
kono
parents:
diff changeset
629 break;
kono
parents:
diff changeset
630 case 64:
kono
parents:
diff changeset
631 res = BRIG_TYPE_F64;
kono
parents:
diff changeset
632 break;
kono
parents:
diff changeset
633 default:
kono
parents:
diff changeset
634 break;
kono
parents:
diff changeset
635 }
kono
parents:
diff changeset
636 }
kono
parents:
diff changeset
637
kono
parents:
diff changeset
638 if (res == BRIG_TYPE_NONE)
kono
parents:
diff changeset
639 {
kono
parents:
diff changeset
640 HSA_SORRY_ATV (EXPR_LOCATION (type),
kono
parents:
diff changeset
641 "support for HSA does not implement type %qT", type);
kono
parents:
diff changeset
642 return res;
kono
parents:
diff changeset
643 }
kono
parents:
diff changeset
644
kono
parents:
diff changeset
645 if (TREE_CODE (type) == VECTOR_TYPE)
kono
parents:
diff changeset
646 {
kono
parents:
diff changeset
647 HOST_WIDE_INT tsize = tree_to_uhwi (TYPE_SIZE (type));
kono
parents:
diff changeset
648
kono
parents:
diff changeset
649 if (bsize == tsize)
kono
parents:
diff changeset
650 {
kono
parents:
diff changeset
651 HSA_SORRY_ATV (EXPR_LOCATION (type),
kono
parents:
diff changeset
652 "support for HSA does not implement a vector type "
kono
parents:
diff changeset
653 "where a type and unit size are equal: %qT", type);
kono
parents:
diff changeset
654 return res;
kono
parents:
diff changeset
655 }
kono
parents:
diff changeset
656
kono
parents:
diff changeset
657 switch (tsize)
kono
parents:
diff changeset
658 {
kono
parents:
diff changeset
659 case 32:
kono
parents:
diff changeset
660 res |= BRIG_TYPE_PACK_32;
kono
parents:
diff changeset
661 break;
kono
parents:
diff changeset
662 case 64:
kono
parents:
diff changeset
663 res |= BRIG_TYPE_PACK_64;
kono
parents:
diff changeset
664 break;
kono
parents:
diff changeset
665 case 128:
kono
parents:
diff changeset
666 res |= BRIG_TYPE_PACK_128;
kono
parents:
diff changeset
667 break;
kono
parents:
diff changeset
668 default:
kono
parents:
diff changeset
669 HSA_SORRY_ATV (EXPR_LOCATION (type),
kono
parents:
diff changeset
670 "support for HSA does not implement type %qT", type);
kono
parents:
diff changeset
671 }
kono
parents:
diff changeset
672 }
kono
parents:
diff changeset
673
kono
parents:
diff changeset
674 if (min32int)
kono
parents:
diff changeset
675 /* Registers/immediate operands can only be 32bit or more except for
kono
parents:
diff changeset
676 f16. */
kono
parents:
diff changeset
677 res = hsa_extend_inttype_to_32bit (res);
kono
parents:
diff changeset
678
kono
parents:
diff changeset
679 if (TREE_CODE (type) == COMPLEX_TYPE)
kono
parents:
diff changeset
680 {
kono
parents:
diff changeset
681 unsigned bsize = 2 * hsa_type_bit_size (res);
kono
parents:
diff changeset
682 res = hsa_bittype_for_bitsize (bsize);
kono
parents:
diff changeset
683 }
kono
parents:
diff changeset
684
kono
parents:
diff changeset
685 return res;
kono
parents:
diff changeset
686 }
kono
parents:
diff changeset
687
kono
parents:
diff changeset
688 /* Returns the BRIG type we need to load/store entities of TYPE. */
kono
parents:
diff changeset
689
kono
parents:
diff changeset
690 static BrigType16_t
kono
parents:
diff changeset
691 mem_type_for_type (BrigType16_t type)
kono
parents:
diff changeset
692 {
kono
parents:
diff changeset
693 /* HSA has non-intuitive constraints on load/store types. If it's
kono
parents:
diff changeset
694 a bit-type it _must_ be B128, if it's not a bit-type it must be
kono
parents:
diff changeset
695 64bit max. So for loading entities of 128 bits (e.g. vectors)
131
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
696 we have to use B128, while for loading the rest we have to use the
111
kono
parents:
diff changeset
697 input type (??? or maybe also flattened to a equally sized non-vector
kono
parents:
diff changeset
698 unsigned type?). */
kono
parents:
diff changeset
699 if ((type & BRIG_TYPE_PACK_MASK) == BRIG_TYPE_PACK_128)
kono
parents:
diff changeset
700 return BRIG_TYPE_B128;
kono
parents:
diff changeset
701 else if (hsa_btype_p (type) || hsa_type_packed_p (type))
kono
parents:
diff changeset
702 {
kono
parents:
diff changeset
703 unsigned bitsize = hsa_type_bit_size (type);
kono
parents:
diff changeset
704 if (bitsize < 128)
kono
parents:
diff changeset
705 return hsa_uint_for_bitsize (bitsize);
kono
parents:
diff changeset
706 else
kono
parents:
diff changeset
707 return hsa_bittype_for_bitsize (bitsize);
kono
parents:
diff changeset
708 }
kono
parents:
diff changeset
709 return type;
kono
parents:
diff changeset
710 }
kono
parents:
diff changeset
711
kono
parents:
diff changeset
712 /* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some
kono
parents:
diff changeset
713 kind of array will be generated, setting DIM appropriately. Otherwise, it
kono
parents:
diff changeset
714 will be set to zero. */
kono
parents:
diff changeset
715
kono
parents:
diff changeset
716 static BrigType16_t
kono
parents:
diff changeset
717 hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p = NULL,
kono
parents:
diff changeset
718 bool min32int = false)
kono
parents:
diff changeset
719 {
kono
parents:
diff changeset
720 gcc_checking_assert (TYPE_P (type));
kono
parents:
diff changeset
721 if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type)))
kono
parents:
diff changeset
722 {
kono
parents:
diff changeset
723 HSA_SORRY_ATV (EXPR_LOCATION (type), "support for HSA does not "
kono
parents:
diff changeset
724 "implement huge or variable-sized type %qT", type);
kono
parents:
diff changeset
725 return BRIG_TYPE_NONE;
kono
parents:
diff changeset
726 }
kono
parents:
diff changeset
727
kono
parents:
diff changeset
728 if (RECORD_OR_UNION_TYPE_P (type))
kono
parents:
diff changeset
729 {
kono
parents:
diff changeset
730 if (dim_p)
kono
parents:
diff changeset
731 *dim_p = tree_to_uhwi (TYPE_SIZE_UNIT (type));
kono
parents:
diff changeset
732 return BRIG_TYPE_U8 | BRIG_TYPE_ARRAY;
kono
parents:
diff changeset
733 }
kono
parents:
diff changeset
734
kono
parents:
diff changeset
735 if (TREE_CODE (type) == ARRAY_TYPE)
kono
parents:
diff changeset
736 {
kono
parents:
diff changeset
737 /* We try to be nice and use the real base-type when this is an array of
kono
parents:
diff changeset
738 scalars and only resort to an array of bytes if the type is more
kono
parents:
diff changeset
739 complex. */
kono
parents:
diff changeset
740
kono
parents:
diff changeset
741 unsigned HOST_WIDE_INT dim = 1;
kono
parents:
diff changeset
742
kono
parents:
diff changeset
743 while (TREE_CODE (type) == ARRAY_TYPE)
kono
parents:
diff changeset
744 {
kono
parents:
diff changeset
745 tree domain = TYPE_DOMAIN (type);
kono
parents:
diff changeset
746 if (!TYPE_MIN_VALUE (domain)
kono
parents:
diff changeset
747 || !TYPE_MAX_VALUE (domain)
kono
parents:
diff changeset
748 || !tree_fits_shwi_p (TYPE_MIN_VALUE (domain))
kono
parents:
diff changeset
749 || !tree_fits_shwi_p (TYPE_MAX_VALUE (domain)))
kono
parents:
diff changeset
750 {
kono
parents:
diff changeset
751 HSA_SORRY_ATV (EXPR_LOCATION (type),
kono
parents:
diff changeset
752 "support for HSA does not implement array "
kono
parents:
diff changeset
753 "%qT with unknown bounds", type);
kono
parents:
diff changeset
754 return BRIG_TYPE_NONE;
kono
parents:
diff changeset
755 }
kono
parents:
diff changeset
756 HOST_WIDE_INT min = tree_to_shwi (TYPE_MIN_VALUE (domain));
kono
parents:
diff changeset
757 HOST_WIDE_INT max = tree_to_shwi (TYPE_MAX_VALUE (domain));
kono
parents:
diff changeset
758 dim = dim * (unsigned HOST_WIDE_INT) (max - min + 1);
kono
parents:
diff changeset
759 type = TREE_TYPE (type);
kono
parents:
diff changeset
760 }
kono
parents:
diff changeset
761
kono
parents:
diff changeset
762 BrigType16_t res;
kono
parents:
diff changeset
763 if (RECORD_OR_UNION_TYPE_P (type))
kono
parents:
diff changeset
764 {
kono
parents:
diff changeset
765 dim = dim * tree_to_uhwi (TYPE_SIZE_UNIT (type));
kono
parents:
diff changeset
766 res = BRIG_TYPE_U8;
kono
parents:
diff changeset
767 }
kono
parents:
diff changeset
768 else
kono
parents:
diff changeset
769 res = hsa_type_for_scalar_tree_type (type, false);
kono
parents:
diff changeset
770
kono
parents:
diff changeset
771 if (dim_p)
kono
parents:
diff changeset
772 *dim_p = dim;
kono
parents:
diff changeset
773 return res | BRIG_TYPE_ARRAY;
kono
parents:
diff changeset
774 }
kono
parents:
diff changeset
775
kono
parents:
diff changeset
776 /* Scalar case: */
kono
parents:
diff changeset
777 if (dim_p)
kono
parents:
diff changeset
778 *dim_p = 0;
kono
parents:
diff changeset
779
kono
parents:
diff changeset
780 return hsa_type_for_scalar_tree_type (type, min32int);
kono
parents:
diff changeset
781 }
kono
parents:
diff changeset
782
kono
parents:
diff changeset
783 /* Returns true if converting from STYPE into DTYPE needs the _CVT
kono
parents:
diff changeset
784 opcode. If false a normal _MOV is enough. */
kono
parents:
diff changeset
785
kono
parents:
diff changeset
786 static bool
kono
parents:
diff changeset
787 hsa_needs_cvt (BrigType16_t dtype, BrigType16_t stype)
kono
parents:
diff changeset
788 {
kono
parents:
diff changeset
789 if (hsa_btype_p (dtype))
kono
parents:
diff changeset
790 return false;
kono
parents:
diff changeset
791
kono
parents:
diff changeset
792 /* float <-> int conversions are real converts. */
kono
parents:
diff changeset
793 if (hsa_type_float_p (dtype) != hsa_type_float_p (stype))
kono
parents:
diff changeset
794 return true;
kono
parents:
diff changeset
795 /* When both types have different size, then we need CVT as well. */
kono
parents:
diff changeset
796 if (hsa_type_bit_size (dtype) != hsa_type_bit_size (stype))
kono
parents:
diff changeset
797 return true;
kono
parents:
diff changeset
798 return false;
kono
parents:
diff changeset
799 }
kono
parents:
diff changeset
800
kono
parents:
diff changeset
801 /* Return declaration name if it exists or create one from UID if it does not.
kono
parents:
diff changeset
802 If DECL is a local variable, make UID part of its name. */
kono
parents:
diff changeset
803
kono
parents:
diff changeset
804 const char *
kono
parents:
diff changeset
805 hsa_get_declaration_name (tree decl)
kono
parents:
diff changeset
806 {
kono
parents:
diff changeset
807 if (!DECL_NAME (decl))
kono
parents:
diff changeset
808 {
kono
parents:
diff changeset
809 char buf[64];
kono
parents:
diff changeset
810 snprintf (buf, 64, "__hsa_anon_%u", DECL_UID (decl));
kono
parents:
diff changeset
811 size_t len = strlen (buf);
kono
parents:
diff changeset
812 char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1);
kono
parents:
diff changeset
813 memcpy (copy, buf, len + 1);
kono
parents:
diff changeset
814 return copy;
kono
parents:
diff changeset
815 }
kono
parents:
diff changeset
816
kono
parents:
diff changeset
817 tree name_tree;
kono
parents:
diff changeset
818 if (TREE_CODE (decl) == FUNCTION_DECL
kono
parents:
diff changeset
819 || (TREE_CODE (decl) == VAR_DECL && is_global_var (decl)))
kono
parents:
diff changeset
820 name_tree = DECL_ASSEMBLER_NAME (decl);
kono
parents:
diff changeset
821 else
kono
parents:
diff changeset
822 name_tree = DECL_NAME (decl);
kono
parents:
diff changeset
823
kono
parents:
diff changeset
824 const char *name = IDENTIFIER_POINTER (name_tree);
kono
parents:
diff changeset
825 /* User-defined assembly names have prepended asterisk symbol. */
kono
parents:
diff changeset
826 if (name[0] == '*')
kono
parents:
diff changeset
827 name++;
kono
parents:
diff changeset
828
kono
parents:
diff changeset
829 if ((TREE_CODE (decl) == VAR_DECL)
kono
parents:
diff changeset
830 && decl_function_context (decl))
kono
parents:
diff changeset
831 {
kono
parents:
diff changeset
832 size_t len = strlen (name);
kono
parents:
diff changeset
833 char *buf = (char *) alloca (len + 32);
kono
parents:
diff changeset
834 snprintf (buf, len + 32, "%s_%u", name, DECL_UID (decl));
kono
parents:
diff changeset
835 len = strlen (buf);
kono
parents:
diff changeset
836 char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1);
kono
parents:
diff changeset
837 memcpy (copy, buf, len + 1);
kono
parents:
diff changeset
838 return copy;
kono
parents:
diff changeset
839 }
kono
parents:
diff changeset
840 else
kono
parents:
diff changeset
841 return name;
kono
parents:
diff changeset
842 }
kono
parents:
diff changeset
843
kono
parents:
diff changeset
844 /* Lookup or create the associated hsa_symbol structure with a given VAR_DECL
kono
parents:
diff changeset
845 or lookup the hsa_structure corresponding to a PARM_DECL. */
kono
parents:
diff changeset
846
kono
parents:
diff changeset
847 static hsa_symbol *
kono
parents:
diff changeset
848 get_symbol_for_decl (tree decl)
kono
parents:
diff changeset
849 {
kono
parents:
diff changeset
850 hsa_symbol **slot;
kono
parents:
diff changeset
851 hsa_symbol dummy (BRIG_TYPE_NONE, BRIG_SEGMENT_NONE, BRIG_LINKAGE_NONE);
kono
parents:
diff changeset
852
kono
parents:
diff changeset
853 gcc_assert (TREE_CODE (decl) == PARM_DECL
kono
parents:
diff changeset
854 || TREE_CODE (decl) == RESULT_DECL
kono
parents:
diff changeset
855 || TREE_CODE (decl) == VAR_DECL
kono
parents:
diff changeset
856 || TREE_CODE (decl) == CONST_DECL);
kono
parents:
diff changeset
857
kono
parents:
diff changeset
858 dummy.m_decl = decl;
kono
parents:
diff changeset
859
kono
parents:
diff changeset
860 bool is_in_global_vars = ((TREE_CODE (decl) == VAR_DECL)
kono
parents:
diff changeset
861 && !decl_function_context (decl));
kono
parents:
diff changeset
862
kono
parents:
diff changeset
863 if (is_in_global_vars)
kono
parents:
diff changeset
864 slot = hsa_global_variable_symbols->find_slot (&dummy, INSERT);
kono
parents:
diff changeset
865 else
kono
parents:
diff changeset
866 slot = hsa_cfun->m_local_symbols->find_slot (&dummy, INSERT);
kono
parents:
diff changeset
867
kono
parents:
diff changeset
868 gcc_checking_assert (slot);
kono
parents:
diff changeset
869 if (*slot)
kono
parents:
diff changeset
870 {
kono
parents:
diff changeset
871 hsa_symbol *sym = (*slot);
kono
parents:
diff changeset
872
kono
parents:
diff changeset
873 /* If the symbol is problematic, mark current function also as
kono
parents:
diff changeset
874 problematic. */
kono
parents:
diff changeset
875 if (sym->m_seen_error)
kono
parents:
diff changeset
876 hsa_fail_cfun ();
kono
parents:
diff changeset
877
kono
parents:
diff changeset
878 /* PR hsa/70234: If a global variable was marked to be emitted,
kono
parents:
diff changeset
879 but HSAIL generation of a function using the variable fails,
kono
parents:
diff changeset
880 we should retry to emit the variable in context of a different
kono
parents:
diff changeset
881 function.
kono
parents:
diff changeset
882
kono
parents:
diff changeset
883 Iterate elements whether a symbol is already in m_global_symbols
kono
parents:
diff changeset
884 of not. */
kono
parents:
diff changeset
885 if (is_in_global_vars && !sym->m_emitted_to_brig)
kono
parents:
diff changeset
886 {
kono
parents:
diff changeset
887 for (unsigned i = 0; i < hsa_cfun->m_global_symbols.length (); i++)
kono
parents:
diff changeset
888 if (hsa_cfun->m_global_symbols[i] == sym)
kono
parents:
diff changeset
889 return *slot;
kono
parents:
diff changeset
890 hsa_cfun->m_global_symbols.safe_push (sym);
kono
parents:
diff changeset
891 }
kono
parents:
diff changeset
892
kono
parents:
diff changeset
893 return *slot;
kono
parents:
diff changeset
894 }
kono
parents:
diff changeset
895 else
kono
parents:
diff changeset
896 {
kono
parents:
diff changeset
897 hsa_symbol *sym;
kono
parents:
diff changeset
898 /* PARM_DECLs and RESULT_DECL should be already in m_local_symbols. */
kono
parents:
diff changeset
899 gcc_assert (TREE_CODE (decl) == VAR_DECL
kono
parents:
diff changeset
900 || TREE_CODE (decl) == CONST_DECL);
kono
parents:
diff changeset
901 BrigAlignment8_t align = hsa_object_alignment (decl);
kono
parents:
diff changeset
902
kono
parents:
diff changeset
903 if (is_in_global_vars)
kono
parents:
diff changeset
904 {
kono
parents:
diff changeset
905 gcc_checking_assert (TREE_CODE (decl) != CONST_DECL);
kono
parents:
diff changeset
906 sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_GLOBAL,
kono
parents:
diff changeset
907 BRIG_LINKAGE_PROGRAM, true,
kono
parents:
diff changeset
908 BRIG_ALLOCATION_PROGRAM, align);
kono
parents:
diff changeset
909 hsa_cfun->m_global_symbols.safe_push (sym);
kono
parents:
diff changeset
910 sym->fillup_for_decl (decl);
kono
parents:
diff changeset
911 if (sym->m_align > align)
kono
parents:
diff changeset
912 {
kono
parents:
diff changeset
913 sym->m_seen_error = true;
kono
parents:
diff changeset
914 HSA_SORRY_ATV (EXPR_LOCATION (decl),
kono
parents:
diff changeset
915 "HSA specification requires that %E is at least "
kono
parents:
diff changeset
916 "naturally aligned", decl);
kono
parents:
diff changeset
917 }
kono
parents:
diff changeset
918 }
kono
parents:
diff changeset
919 else
kono
parents:
diff changeset
920 {
kono
parents:
diff changeset
921 /* As generation of efficient memory copy instructions relies
kono
parents:
diff changeset
922 on alignment greater or equal to 8 bytes,
kono
parents:
diff changeset
923 we need to increase alignment of all aggregate types.. */
kono
parents:
diff changeset
924 if (AGGREGATE_TYPE_P (TREE_TYPE (decl)))
kono
parents:
diff changeset
925 align = MAX ((BrigAlignment8_t) BRIG_ALIGNMENT_8, align);
kono
parents:
diff changeset
926
kono
parents:
diff changeset
927 BrigAllocation allocation = BRIG_ALLOCATION_AUTOMATIC;
kono
parents:
diff changeset
928 BrigSegment8_t segment;
kono
parents:
diff changeset
929 if (TREE_CODE (decl) == CONST_DECL)
kono
parents:
diff changeset
930 {
kono
parents:
diff changeset
931 segment = BRIG_SEGMENT_READONLY;
kono
parents:
diff changeset
932 allocation = BRIG_ALLOCATION_AGENT;
kono
parents:
diff changeset
933 }
kono
parents:
diff changeset
934 else if (lookup_attribute ("hsa_group_segment",
kono
parents:
diff changeset
935 DECL_ATTRIBUTES (decl)))
kono
parents:
diff changeset
936 segment = BRIG_SEGMENT_GROUP;
131
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
937 else if (TREE_STATIC (decl))
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
938 {
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
939 segment = BRIG_SEGMENT_GLOBAL;
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
940 allocation = BRIG_ALLOCATION_PROGRAM;
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
941 }
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
942 else if (lookup_attribute ("hsa_global_segment",
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
943 DECL_ATTRIBUTES (decl)))
111
kono
parents:
diff changeset
944 segment = BRIG_SEGMENT_GLOBAL;
kono
parents:
diff changeset
945 else
kono
parents:
diff changeset
946 segment = BRIG_SEGMENT_PRIVATE;
kono
parents:
diff changeset
947
kono
parents:
diff changeset
948 sym = new hsa_symbol (BRIG_TYPE_NONE, segment, BRIG_LINKAGE_FUNCTION,
kono
parents:
diff changeset
949 false, allocation, align);
kono
parents:
diff changeset
950 sym->fillup_for_decl (decl);
kono
parents:
diff changeset
951 hsa_cfun->m_private_variables.safe_push (sym);
kono
parents:
diff changeset
952 }
kono
parents:
diff changeset
953
kono
parents:
diff changeset
954 sym->m_name = hsa_get_declaration_name (decl);
kono
parents:
diff changeset
955 *slot = sym;
kono
parents:
diff changeset
956 return sym;
kono
parents:
diff changeset
957 }
kono
parents:
diff changeset
958 }
kono
parents:
diff changeset
959
kono
parents:
diff changeset
960 /* For a given HSA function declaration, return a host
kono
parents:
diff changeset
961 function declaration. */
kono
parents:
diff changeset
962
kono
parents:
diff changeset
963 tree
kono
parents:
diff changeset
964 hsa_get_host_function (tree decl)
kono
parents:
diff changeset
965 {
131
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
966 hsa_function_summary *s = hsa_summaries->get (cgraph_node::get_create (decl));
111
kono
parents:
diff changeset
967 gcc_assert (s->m_gpu_implementation_p);
kono
parents:
diff changeset
968
kono
parents:
diff changeset
969 return s->m_bound_function ? s->m_bound_function->decl : NULL;
kono
parents:
diff changeset
970 }
kono
parents:
diff changeset
971
kono
parents:
diff changeset
972 /* Return true if function DECL has a host equivalent function. */
kono
parents:
diff changeset
973
kono
parents:
diff changeset
974 static char *
kono
parents:
diff changeset
975 get_brig_function_name (tree decl)
kono
parents:
diff changeset
976 {
kono
parents:
diff changeset
977 tree d = decl;
kono
parents:
diff changeset
978
kono
parents:
diff changeset
979 hsa_function_summary *s = hsa_summaries->get (cgraph_node::get_create (d));
131
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
980 if (s != NULL
111
kono
parents:
diff changeset
981 && s->m_gpu_implementation_p
kono
parents:
diff changeset
982 && s->m_bound_function)
kono
parents:
diff changeset
983 d = s->m_bound_function->decl;
kono
parents:
diff changeset
984
kono
parents:
diff changeset
985 /* IPA split can create a function that has no host equivalent. */
kono
parents:
diff changeset
986 if (d == NULL)
kono
parents:
diff changeset
987 d = decl;
kono
parents:
diff changeset
988
kono
parents:
diff changeset
989 char *name = xstrdup (hsa_get_declaration_name (d));
kono
parents:
diff changeset
990 hsa_sanitize_name (name);
kono
parents:
diff changeset
991
kono
parents:
diff changeset
992 return name;
kono
parents:
diff changeset
993 }
kono
parents:
diff changeset
994
kono
parents:
diff changeset
995 /* Create a spill symbol of type TYPE. */
kono
parents:
diff changeset
996
kono
parents:
diff changeset
997 hsa_symbol *
kono
parents:
diff changeset
998 hsa_get_spill_symbol (BrigType16_t type)
kono
parents:
diff changeset
999 {
kono
parents:
diff changeset
1000 hsa_symbol *sym = new hsa_symbol (type, BRIG_SEGMENT_SPILL,
kono
parents:
diff changeset
1001 BRIG_LINKAGE_FUNCTION);
kono
parents:
diff changeset
1002 hsa_cfun->m_spill_symbols.safe_push (sym);
kono
parents:
diff changeset
1003 return sym;
kono
parents:
diff changeset
1004 }
kono
parents:
diff changeset
1005
kono
parents:
diff changeset
1006 /* Create a symbol for a read-only string constant. */
kono
parents:
diff changeset
1007 hsa_symbol *
kono
parents:
diff changeset
1008 hsa_get_string_cst_symbol (tree string_cst)
kono
parents:
diff changeset
1009 {
kono
parents:
diff changeset
1010 gcc_checking_assert (TREE_CODE (string_cst) == STRING_CST);
kono
parents:
diff changeset
1011
kono
parents:
diff changeset
1012 hsa_symbol **slot = hsa_cfun->m_string_constants_map.get (string_cst);
kono
parents:
diff changeset
1013 if (slot)
kono
parents:
diff changeset
1014 return *slot;
kono
parents:
diff changeset
1015
kono
parents:
diff changeset
1016 hsa_op_immed *cst = new hsa_op_immed (string_cst);
kono
parents:
diff changeset
1017 hsa_symbol *sym = new hsa_symbol (cst->m_type, BRIG_SEGMENT_GLOBAL,
kono
parents:
diff changeset
1018 BRIG_LINKAGE_MODULE, true,
kono
parents:
diff changeset
1019 BRIG_ALLOCATION_AGENT);
kono
parents:
diff changeset
1020 sym->m_cst_value = cst;
kono
parents:
diff changeset
1021 sym->m_dim = TREE_STRING_LENGTH (string_cst);
kono
parents:
diff changeset
1022 sym->m_name_number = hsa_cfun->m_global_symbols.length ();
kono
parents:
diff changeset
1023
kono
parents:
diff changeset
1024 hsa_cfun->m_global_symbols.safe_push (sym);
kono
parents:
diff changeset
1025 hsa_cfun->m_string_constants_map.put (string_cst, sym);
kono
parents:
diff changeset
1026 return sym;
kono
parents:
diff changeset
1027 }
kono
parents:
diff changeset
1028
kono
parents:
diff changeset
1029 /* Make the type of a MOV instruction larger if mandated by HSAIL rules. */
kono
parents:
diff changeset
1030
kono
parents:
diff changeset
1031 static void
kono
parents:
diff changeset
1032 hsa_fixup_mov_insn_type (hsa_insn_basic *insn)
kono
parents:
diff changeset
1033 {
kono
parents:
diff changeset
1034 insn->m_type = hsa_extend_inttype_to_32bit (insn->m_type);
kono
parents:
diff changeset
1035 if (insn->m_type == BRIG_TYPE_B8 || insn->m_type == BRIG_TYPE_B16)
kono
parents:
diff changeset
1036 insn->m_type = BRIG_TYPE_B32;
kono
parents:
diff changeset
1037 }
kono
parents:
diff changeset
1038
kono
parents:
diff changeset
1039 /* Constructor of the ancestor of all operands. K is BRIG kind that identified
kono
parents:
diff changeset
1040 what the operator is. */
kono
parents:
diff changeset
1041
kono
parents:
diff changeset
1042 hsa_op_base::hsa_op_base (BrigKind16_t k)
kono
parents:
diff changeset
1043 : m_next (NULL), m_brig_op_offset (0), m_kind (k)
kono
parents:
diff changeset
1044 {
kono
parents:
diff changeset
1045 hsa_operands.safe_push (this);
kono
parents:
diff changeset
1046 }
kono
parents:
diff changeset
1047
kono
parents:
diff changeset
1048 /* Constructor of ancestor of all operands which have a type. K is BRIG kind
kono
parents:
diff changeset
1049 that identified what the operator is. T is the type of the operator. */
kono
parents:
diff changeset
1050
kono
parents:
diff changeset
1051 hsa_op_with_type::hsa_op_with_type (BrigKind16_t k, BrigType16_t t)
kono
parents:
diff changeset
1052 : hsa_op_base (k), m_type (t)
kono
parents:
diff changeset
1053 {
kono
parents:
diff changeset
1054 }
kono
parents:
diff changeset
1055
kono
parents:
diff changeset
1056 hsa_op_with_type *
kono
parents:
diff changeset
1057 hsa_op_with_type::get_in_type (BrigType16_t dtype, hsa_bb *hbb)
kono
parents:
diff changeset
1058 {
kono
parents:
diff changeset
1059 if (m_type == dtype)
kono
parents:
diff changeset
1060 return this;
kono
parents:
diff changeset
1061
kono
parents:
diff changeset
1062 hsa_op_reg *dest;
kono
parents:
diff changeset
1063
kono
parents:
diff changeset
1064 if (hsa_needs_cvt (dtype, m_type))
kono
parents:
diff changeset
1065 {
kono
parents:
diff changeset
1066 dest = new hsa_op_reg (dtype);
kono
parents:
diff changeset
1067 hbb->append_insn (new hsa_insn_cvt (dest, this));
kono
parents:
diff changeset
1068 }
kono
parents:
diff changeset
1069 else if (is_a <hsa_op_reg *> (this))
kono
parents:
diff changeset
1070 {
kono
parents:
diff changeset
1071 /* In the end, HSA registers do not really have types, only sizes, so if
kono
parents:
diff changeset
1072 the sizes match, we can use the register directly. */
kono
parents:
diff changeset
1073 gcc_checking_assert (hsa_type_bit_size (dtype)
kono
parents:
diff changeset
1074 == hsa_type_bit_size (m_type));
kono
parents:
diff changeset
1075 return this;
kono
parents:
diff changeset
1076 }
kono
parents:
diff changeset
1077 else
kono
parents:
diff changeset
1078 {
kono
parents:
diff changeset
1079 dest = new hsa_op_reg (m_type);
kono
parents:
diff changeset
1080
kono
parents:
diff changeset
1081 hsa_insn_basic *mov = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
kono
parents:
diff changeset
1082 dest->m_type, dest, this);
kono
parents:
diff changeset
1083 hsa_fixup_mov_insn_type (mov);
kono
parents:
diff changeset
1084 hbb->append_insn (mov);
kono
parents:
diff changeset
1085 /* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
kono
parents:
diff changeset
1086 type of the operand must be same as type of the instruction. */
kono
parents:
diff changeset
1087 dest->m_type = dtype;
kono
parents:
diff changeset
1088 }
kono
parents:
diff changeset
1089
kono
parents:
diff changeset
1090 return dest;
kono
parents:
diff changeset
1091 }
kono
parents:
diff changeset
1092
kono
parents:
diff changeset
1093 /* If this operand has integer type smaller than 32 bits, extend it to 32 bits,
kono
parents:
diff changeset
1094 adding instructions to HBB if needed. */
kono
parents:
diff changeset
1095
kono
parents:
diff changeset
1096 hsa_op_with_type *
kono
parents:
diff changeset
1097 hsa_op_with_type::extend_int_to_32bit (hsa_bb *hbb)
kono
parents:
diff changeset
1098 {
kono
parents:
diff changeset
1099 if (m_type == BRIG_TYPE_U8 || m_type == BRIG_TYPE_U16)
kono
parents:
diff changeset
1100 return get_in_type (BRIG_TYPE_U32, hbb);
kono
parents:
diff changeset
1101 else if (m_type == BRIG_TYPE_S8 || m_type == BRIG_TYPE_S16)
kono
parents:
diff changeset
1102 return get_in_type (BRIG_TYPE_S32, hbb);
kono
parents:
diff changeset
1103 else
kono
parents:
diff changeset
1104 return this;
kono
parents:
diff changeset
1105 }
kono
parents:
diff changeset
1106
kono
parents:
diff changeset
1107 /* Constructor of class representing HSA immediate values. TREE_VAL is the
kono
parents:
diff changeset
1108 tree representation of the immediate value. If min32int is true,
kono
parents:
diff changeset
1109 always expand integer types to one that has at least 32 bits. */
kono
parents:
diff changeset
1110
kono
parents:
diff changeset
1111 hsa_op_immed::hsa_op_immed (tree tree_val, bool min32int)
kono
parents:
diff changeset
1112 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES,
kono
parents:
diff changeset
1113 hsa_type_for_tree_type (TREE_TYPE (tree_val), NULL,
kono
parents:
diff changeset
1114 min32int))
kono
parents:
diff changeset
1115 {
kono
parents:
diff changeset
1116 if (hsa_seen_error ())
kono
parents:
diff changeset
1117 return;
kono
parents:
diff changeset
1118
kono
parents:
diff changeset
1119 gcc_checking_assert ((is_gimple_min_invariant (tree_val)
kono
parents:
diff changeset
1120 && (!POINTER_TYPE_P (TREE_TYPE (tree_val))
kono
parents:
diff changeset
1121 || TREE_CODE (tree_val) == INTEGER_CST))
kono
parents:
diff changeset
1122 || TREE_CODE (tree_val) == CONSTRUCTOR);
kono
parents:
diff changeset
1123 m_tree_value = tree_val;
kono
parents:
diff changeset
1124
kono
parents:
diff changeset
1125 /* Verify that all elements of a constructor are constants. */
kono
parents:
diff changeset
1126 if (TREE_CODE (m_tree_value) == CONSTRUCTOR)
kono
parents:
diff changeset
1127 for (unsigned i = 0; i < CONSTRUCTOR_NELTS (m_tree_value); i++)
kono
parents:
diff changeset
1128 {
kono
parents:
diff changeset
1129 tree v = CONSTRUCTOR_ELT (m_tree_value, i)->value;
kono
parents:
diff changeset
1130 if (!CONSTANT_CLASS_P (v))
kono
parents:
diff changeset
1131 {
kono
parents:
diff changeset
1132 HSA_SORRY_AT (EXPR_LOCATION (tree_val),
kono
parents:
diff changeset
1133 "HSA ctor should have only constants");
kono
parents:
diff changeset
1134 return;
kono
parents:
diff changeset
1135 }
kono
parents:
diff changeset
1136 }
kono
parents:
diff changeset
1137 }
kono
parents:
diff changeset
1138
kono
parents:
diff changeset
1139 /* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
kono
parents:
diff changeset
1140 integer representation of the immediate value. TYPE is BRIG type. */
kono
parents:
diff changeset
1141
kono
parents:
diff changeset
1142 hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value, BrigType16_t type)
kono
parents:
diff changeset
1143 : hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES, type),
kono
parents:
diff changeset
1144 m_tree_value (NULL)
kono
parents:
diff changeset
1145 {
kono
parents:
diff changeset
1146 gcc_assert (hsa_type_integer_p (type));
kono
parents:
diff changeset
1147 m_int_value = integer_value;
kono
parents:
diff changeset
1148 }
kono
parents:
diff changeset
1149
kono
parents:
diff changeset
1150 hsa_op_immed::hsa_op_immed ()
kono
parents:
diff changeset
1151 : hsa_op_with_type (BRIG_KIND_NONE, BRIG_TYPE_NONE)
kono
parents:
diff changeset
1152 {
kono
parents:
diff changeset
1153 }
kono
parents:
diff changeset
1154
kono
parents:
diff changeset
1155 /* New operator to allocate immediate operands from obstack. */
kono
parents:
diff changeset
1156
kono
parents:
diff changeset
1157 void *
kono
parents:
diff changeset
1158 hsa_op_immed::operator new (size_t size)
kono
parents:
diff changeset
1159 {
kono
parents:
diff changeset
1160 return obstack_alloc (&hsa_obstack, size);
kono
parents:
diff changeset
1161 }
kono
parents:
diff changeset
1162
kono
parents:
diff changeset
1163 /* Destructor. */
kono
parents:
diff changeset
1164
kono
parents:
diff changeset
1165 hsa_op_immed::~hsa_op_immed ()
kono
parents:
diff changeset
1166 {
kono
parents:
diff changeset
1167 }
kono
parents:
diff changeset
1168
kono
parents:
diff changeset
1169 /* Change type of the immediate value to T. */
kono
parents:
diff changeset
1170
kono
parents:
diff changeset
1171 void
kono
parents:
diff changeset
1172 hsa_op_immed::set_type (BrigType16_t t)
kono
parents:
diff changeset
1173 {
kono
parents:
diff changeset
1174 m_type = t;
kono
parents:
diff changeset
1175 }
kono
parents:
diff changeset
1176
kono
parents:
diff changeset
1177 /* Constructor of class representing HSA registers and pseudo-registers. T is
kono
parents:
diff changeset
1178 the BRIG type of the new register. */
kono
parents:
diff changeset
1179
kono
parents:
diff changeset
1180 hsa_op_reg::hsa_op_reg (BrigType16_t t)
kono
parents:
diff changeset
1181 : hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER, t), m_gimple_ssa (NULL_TREE),
kono
parents:
diff changeset
1182 m_def_insn (NULL), m_spill_sym (NULL), m_order (hsa_cfun->m_reg_count++),
kono
parents:
diff changeset
1183 m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
kono
parents:
diff changeset
1184 {
kono
parents:
diff changeset
1185 }
kono
parents:
diff changeset
1186
kono
parents:
diff changeset
1187 /* New operator to allocate a register from obstack. */
kono
parents:
diff changeset
1188
kono
parents:
diff changeset
1189 void *
kono
parents:
diff changeset
1190 hsa_op_reg::operator new (size_t size)
kono
parents:
diff changeset
1191 {
kono
parents:
diff changeset
1192 return obstack_alloc (&hsa_obstack, size);
kono
parents:
diff changeset
1193 }
kono
parents:
diff changeset
1194
kono
parents:
diff changeset
1195 /* Verify register operand. */
kono
parents:
diff changeset
1196
kono
parents:
diff changeset
1197 void
kono
parents:
diff changeset
1198 hsa_op_reg::verify_ssa ()
kono
parents:
diff changeset
1199 {
kono
parents:
diff changeset
1200 /* Verify that each HSA register has a definition assigned.
kono
parents:
diff changeset
1201 Exceptions are VAR_DECL and PARM_DECL that are a default
kono
parents:
diff changeset
1202 definition. */
kono
parents:
diff changeset
1203 gcc_checking_assert (m_def_insn
kono
parents:
diff changeset
1204 || (m_gimple_ssa != NULL
kono
parents:
diff changeset
1205 && (!SSA_NAME_VAR (m_gimple_ssa)
kono
parents:
diff changeset
1206 || (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa))
kono
parents:
diff changeset
1207 != PARM_DECL))
kono
parents:
diff changeset
1208 && SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa)));
kono
parents:
diff changeset
1209
kono
parents:
diff changeset
1210 /* Verify that every use of the register is really present
kono
parents:
diff changeset
1211 in an instruction. */
kono
parents:
diff changeset
1212 for (unsigned i = 0; i < m_uses.length (); i++)
kono
parents:
diff changeset
1213 {
kono
parents:
diff changeset
1214 hsa_insn_basic *use = m_uses[i];
kono
parents:
diff changeset
1215
kono
parents:
diff changeset
1216 bool is_visited = false;
kono
parents:
diff changeset
1217 for (unsigned j = 0; j < use->operand_count (); j++)
kono
parents:
diff changeset
1218 {
kono
parents:
diff changeset
1219 hsa_op_base *u = use->get_op (j);
kono
parents:
diff changeset
1220 hsa_op_address *addr; addr = dyn_cast <hsa_op_address *> (u);
kono
parents:
diff changeset
1221 if (addr && addr->m_reg)
kono
parents:
diff changeset
1222 u = addr->m_reg;
kono
parents:
diff changeset
1223
kono
parents:
diff changeset
1224 if (u == this)
kono
parents:
diff changeset
1225 {
kono
parents:
diff changeset
1226 bool r = !addr && use->op_output_p (j);
kono
parents:
diff changeset
1227
kono
parents:
diff changeset
1228 if (r)
kono
parents:
diff changeset
1229 {
kono
parents:
diff changeset
1230 error ("HSA SSA name defined by instruction that is supposed "
kono
parents:
diff changeset
1231 "to be using it");
kono
parents:
diff changeset
1232 debug_hsa_operand (this);
kono
parents:
diff changeset
1233 debug_hsa_insn (use);
kono
parents:
diff changeset
1234 internal_error ("HSA SSA verification failed");
kono
parents:
diff changeset
1235 }
kono
parents:
diff changeset
1236
kono
parents:
diff changeset
1237 is_visited = true;
kono
parents:
diff changeset
1238 }
kono
parents:
diff changeset
1239 }
kono
parents:
diff changeset
1240
kono
parents:
diff changeset
1241 if (!is_visited)
kono
parents:
diff changeset
1242 {
kono
parents:
diff changeset
1243 error ("HSA SSA name not among operands of instruction that is "
kono
parents:
diff changeset
1244 "supposed to use it");
kono
parents:
diff changeset
1245 debug_hsa_operand (this);
kono
parents:
diff changeset
1246 debug_hsa_insn (use);
kono
parents:
diff changeset
1247 internal_error ("HSA SSA verification failed");
kono
parents:
diff changeset
1248 }
kono
parents:
diff changeset
1249 }
kono
parents:
diff changeset
1250 }
kono
parents:
diff changeset
1251
kono
parents:
diff changeset
1252 hsa_op_address::hsa_op_address (hsa_symbol *sym, hsa_op_reg *r,
kono
parents:
diff changeset
1253 HOST_WIDE_INT offset)
kono
parents:
diff changeset
1254 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (r),
kono
parents:
diff changeset
1255 m_imm_offset (offset)
kono
parents:
diff changeset
1256 {
kono
parents:
diff changeset
1257 }
kono
parents:
diff changeset
1258
kono
parents:
diff changeset
1259 hsa_op_address::hsa_op_address (hsa_symbol *sym, HOST_WIDE_INT offset)
kono
parents:
diff changeset
1260 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (NULL),
kono
parents:
diff changeset
1261 m_imm_offset (offset)
kono
parents:
diff changeset
1262 {
kono
parents:
diff changeset
1263 }
kono
parents:
diff changeset
1264
kono
parents:
diff changeset
1265 hsa_op_address::hsa_op_address (hsa_op_reg *r, HOST_WIDE_INT offset)
kono
parents:
diff changeset
1266 : hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (NULL), m_reg (r),
kono
parents:
diff changeset
1267 m_imm_offset (offset)
kono
parents:
diff changeset
1268 {
kono
parents:
diff changeset
1269 }
kono
parents:
diff changeset
1270
kono
parents:
diff changeset
1271 /* New operator to allocate address operands from obstack. */
kono
parents:
diff changeset
1272
kono
parents:
diff changeset
1273 void *
kono
parents:
diff changeset
1274 hsa_op_address::operator new (size_t size)
kono
parents:
diff changeset
1275 {
kono
parents:
diff changeset
1276 return obstack_alloc (&hsa_obstack, size);
kono
parents:
diff changeset
1277 }
kono
parents:
diff changeset
1278
kono
parents:
diff changeset
1279 /* Constructor of an operand referring to HSAIL code. */
kono
parents:
diff changeset
1280
kono
parents:
diff changeset
1281 hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF),
kono
parents:
diff changeset
1282 m_directive_offset (0)
kono
parents:
diff changeset
1283 {
kono
parents:
diff changeset
1284 }
kono
parents:
diff changeset
1285
kono
parents:
diff changeset
1286 /* Constructor of an operand representing a code list. Set it up so that it
kono
parents:
diff changeset
1287 can contain ELEMENTS number of elements. */
kono
parents:
diff changeset
1288
kono
parents:
diff changeset
1289 hsa_op_code_list::hsa_op_code_list (unsigned elements)
kono
parents:
diff changeset
1290 : hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST)
kono
parents:
diff changeset
1291 {
kono
parents:
diff changeset
1292 m_offsets.create (1);
kono
parents:
diff changeset
1293 m_offsets.safe_grow_cleared (elements);
kono
parents:
diff changeset
1294 }
kono
parents:
diff changeset
1295
kono
parents:
diff changeset
1296 /* New operator to allocate code list operands from obstack. */
kono
parents:
diff changeset
1297
kono
parents:
diff changeset
1298 void *
kono
parents:
diff changeset
1299 hsa_op_code_list::operator new (size_t size)
kono
parents:
diff changeset
1300 {
kono
parents:
diff changeset
1301 return obstack_alloc (&hsa_obstack, size);
kono
parents:
diff changeset
1302 }
kono
parents:
diff changeset
1303
kono
parents:
diff changeset
1304 /* Constructor of an operand representing an operand list.
kono
parents:
diff changeset
1305 Set it up so that it can contain ELEMENTS number of elements. */
kono
parents:
diff changeset
1306
kono
parents:
diff changeset
1307 hsa_op_operand_list::hsa_op_operand_list (unsigned elements)
kono
parents:
diff changeset
1308 : hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST)
kono
parents:
diff changeset
1309 {
kono
parents:
diff changeset
1310 m_offsets.create (elements);
kono
parents:
diff changeset
1311 m_offsets.safe_grow (elements);
kono
parents:
diff changeset
1312 }
kono
parents:
diff changeset
1313
kono
parents:
diff changeset
1314 /* New operator to allocate operand list operands from obstack. */
kono
parents:
diff changeset
1315
kono
parents:
diff changeset
1316 void *
kono
parents:
diff changeset
1317 hsa_op_operand_list::operator new (size_t size)
kono
parents:
diff changeset
1318 {
kono
parents:
diff changeset
1319 return obstack_alloc (&hsa_obstack, size);
kono
parents:
diff changeset
1320 }
kono
parents:
diff changeset
1321
kono
parents:
diff changeset
1322 hsa_op_operand_list::~hsa_op_operand_list ()
kono
parents:
diff changeset
1323 {
kono
parents:
diff changeset
1324 m_offsets.release ();
kono
parents:
diff changeset
1325 }
kono
parents:
diff changeset
1326
kono
parents:
diff changeset
1327
kono
parents:
diff changeset
1328 hsa_op_reg *
kono
parents:
diff changeset
1329 hsa_function_representation::reg_for_gimple_ssa (tree ssa)
kono
parents:
diff changeset
1330 {
kono
parents:
diff changeset
1331 hsa_op_reg *hreg;
kono
parents:
diff changeset
1332
kono
parents:
diff changeset
1333 gcc_checking_assert (TREE_CODE (ssa) == SSA_NAME);
kono
parents:
diff changeset
1334 if (m_ssa_map[SSA_NAME_VERSION (ssa)])
kono
parents:
diff changeset
1335 return m_ssa_map[SSA_NAME_VERSION (ssa)];
kono
parents:
diff changeset
1336
kono
parents:
diff changeset
1337 hreg = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa),
kono
parents:
diff changeset
1338 false));
kono
parents:
diff changeset
1339 hreg->m_gimple_ssa = ssa;
kono
parents:
diff changeset
1340 m_ssa_map[SSA_NAME_VERSION (ssa)] = hreg;
kono
parents:
diff changeset
1341
kono
parents:
diff changeset
1342 return hreg;
kono
parents:
diff changeset
1343 }
kono
parents:
diff changeset
1344
kono
parents:
diff changeset
1345 void
kono
parents:
diff changeset
1346 hsa_op_reg::set_definition (hsa_insn_basic *insn)
kono
parents:
diff changeset
1347 {
kono
parents:
diff changeset
1348 if (hsa_cfun->m_in_ssa)
kono
parents:
diff changeset
1349 {
kono
parents:
diff changeset
1350 gcc_checking_assert (!m_def_insn);
kono
parents:
diff changeset
1351 m_def_insn = insn;
kono
parents:
diff changeset
1352 }
kono
parents:
diff changeset
1353 else
kono
parents:
diff changeset
1354 m_def_insn = NULL;
kono
parents:
diff changeset
1355 }
kono
parents:
diff changeset
1356
kono
parents:
diff changeset
1357 /* Constructor of the class which is the bases of all instructions and directly
kono
parents:
diff changeset
1358 represents the most basic ones. NOPS is the number of operands that the
kono
parents:
diff changeset
1359 operand vector will contain (and which will be cleared). OP is the opcode
kono
parents:
diff changeset
1360 of the instruction. This constructor does not set type. */
kono
parents:
diff changeset
1361
kono
parents:
diff changeset
1362 hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc)
kono
parents:
diff changeset
1363 : m_prev (NULL),
kono
parents:
diff changeset
1364 m_next (NULL), m_bb (NULL), m_opcode (opc), m_number (0),
kono
parents:
diff changeset
1365 m_type (BRIG_TYPE_NONE), m_brig_offset (0)
kono
parents:
diff changeset
1366 {
kono
parents:
diff changeset
1367 if (nops > 0)
kono
parents:
diff changeset
1368 m_operands.safe_grow_cleared (nops);
kono
parents:
diff changeset
1369
kono
parents:
diff changeset
1370 hsa_instructions.safe_push (this);
kono
parents:
diff changeset
1371 }
kono
parents:
diff changeset
1372
kono
parents:
diff changeset
1373 /* Make OP the operand number INDEX of operands of this instruction. If OP is a
kono
parents:
diff changeset
1374 register or an address containing a register, then either set the definition
kono
parents:
diff changeset
1375 of the register to this instruction if it an output operand or add this
kono
parents:
diff changeset
1376 instruction to the uses if it is an input one. */
kono
parents:
diff changeset
1377
kono
parents:
diff changeset
1378 void
kono
parents:
diff changeset
1379 hsa_insn_basic::set_op (int index, hsa_op_base *op)
kono
parents:
diff changeset
1380 {
kono
parents:
diff changeset
1381 /* Each address operand is always use. */
kono
parents:
diff changeset
1382 hsa_op_address *addr = dyn_cast <hsa_op_address *> (op);
kono
parents:
diff changeset
1383 if (addr && addr->m_reg)
kono
parents:
diff changeset
1384 addr->m_reg->m_uses.safe_push (this);
kono
parents:
diff changeset
1385 else
kono
parents:
diff changeset
1386 {
kono
parents:
diff changeset
1387 hsa_op_reg *reg = dyn_cast <hsa_op_reg *> (op);
kono
parents:
diff changeset
1388 if (reg)
kono
parents:
diff changeset
1389 {
kono
parents:
diff changeset
1390 if (op_output_p (index))
kono
parents:
diff changeset
1391 reg->set_definition (this);
kono
parents:
diff changeset
1392 else
kono
parents:
diff changeset
1393 reg->m_uses.safe_push (this);
kono
parents:
diff changeset
1394 }
kono
parents:
diff changeset
1395 }
kono
parents:
diff changeset
1396
kono
parents:
diff changeset
1397 m_operands[index] = op;
kono
parents:
diff changeset
1398 }
kono
parents:
diff changeset
1399
kono
parents:
diff changeset
1400 /* Get INDEX-th operand of the instruction. */
kono
parents:
diff changeset
1401
kono
parents:
diff changeset
1402 hsa_op_base *
kono
parents:
diff changeset
1403 hsa_insn_basic::get_op (int index)
kono
parents:
diff changeset
1404 {
kono
parents:
diff changeset
1405 return m_operands[index];
kono
parents:
diff changeset
1406 }
kono
parents:
diff changeset
1407
kono
parents:
diff changeset
1408 /* Get address of INDEX-th operand of the instruction. */
kono
parents:
diff changeset
1409
kono
parents:
diff changeset
1410 hsa_op_base **
kono
parents:
diff changeset
1411 hsa_insn_basic::get_op_addr (int index)
kono
parents:
diff changeset
1412 {
kono
parents:
diff changeset
1413 return &m_operands[index];
kono
parents:
diff changeset
1414 }
kono
parents:
diff changeset
1415
kono
parents:
diff changeset
1416 /* Get number of operands of the instruction. */
kono
parents:
diff changeset
1417 unsigned int
kono
parents:
diff changeset
1418 hsa_insn_basic::operand_count ()
kono
parents:
diff changeset
1419 {
kono
parents:
diff changeset
1420 return m_operands.length ();
kono
parents:
diff changeset
1421 }
kono
parents:
diff changeset
1422
kono
parents:
diff changeset
1423 /* Constructor of the class which is the bases of all instructions and directly
kono
parents:
diff changeset
1424 represents the most basic ones. NOPS is the number of operands that the
kono
parents:
diff changeset
1425 operand vector will contain (and which will be cleared). OPC is the opcode
kono
parents:
diff changeset
1426 of the instruction, T is the type of the instruction. */
kono
parents:
diff changeset
1427
kono
parents:
diff changeset
1428 hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc, BrigType16_t t,
kono
parents:
diff changeset
1429 hsa_op_base *arg0, hsa_op_base *arg1,
kono
parents:
diff changeset
1430 hsa_op_base *arg2, hsa_op_base *arg3)
kono
parents:
diff changeset
1431 : m_prev (NULL), m_next (NULL), m_bb (NULL), m_opcode (opc),m_number (0),
kono
parents:
diff changeset
1432 m_type (t), m_brig_offset (0)
kono
parents:
diff changeset
1433 {
kono
parents:
diff changeset
1434 if (nops > 0)
kono
parents:
diff changeset
1435 m_operands.safe_grow_cleared (nops);
kono
parents:
diff changeset
1436
kono
parents:
diff changeset
1437 if (arg0 != NULL)
kono
parents:
diff changeset
1438 {
kono
parents:
diff changeset
1439 gcc_checking_assert (nops >= 1);
kono
parents:
diff changeset
1440 set_op (0, arg0);
kono
parents:
diff changeset
1441 }
kono
parents:
diff changeset
1442
kono
parents:
diff changeset
1443 if (arg1 != NULL)
kono
parents:
diff changeset
1444 {
kono
parents:
diff changeset
1445 gcc_checking_assert (nops >= 2);
kono
parents:
diff changeset
1446 set_op (1, arg1);
kono
parents:
diff changeset
1447 }
kono
parents:
diff changeset
1448
kono
parents:
diff changeset
1449 if (arg2 != NULL)
kono
parents:
diff changeset
1450 {
kono
parents:
diff changeset
1451 gcc_checking_assert (nops >= 3);
kono
parents:
diff changeset
1452 set_op (2, arg2);
kono
parents:
diff changeset
1453 }
kono
parents:
diff changeset
1454
kono
parents:
diff changeset
1455 if (arg3 != NULL)
kono
parents:
diff changeset
1456 {
kono
parents:
diff changeset
1457 gcc_checking_assert (nops >= 4);
kono
parents:
diff changeset
1458 set_op (3, arg3);
kono
parents:
diff changeset
1459 }
kono
parents:
diff changeset
1460
kono
parents:
diff changeset
1461 hsa_instructions.safe_push (this);
kono
parents:
diff changeset
1462 }
kono
parents:
diff changeset
1463
kono
parents:
diff changeset
1464 /* New operator to allocate basic instruction from obstack. */
kono
parents:
diff changeset
1465
kono
parents:
diff changeset
1466 void *
kono
parents:
diff changeset
1467 hsa_insn_basic::operator new (size_t size)
kono
parents:
diff changeset
1468 {
kono
parents:
diff changeset
1469 return obstack_alloc (&hsa_obstack, size);
kono
parents:
diff changeset
1470 }
kono
parents:
diff changeset
1471
kono
parents:
diff changeset
1472 /* Verify the instruction. */
kono
parents:
diff changeset
1473
kono
parents:
diff changeset
1474 void
kono
parents:
diff changeset
1475 hsa_insn_basic::verify ()
kono
parents:
diff changeset
1476 {
kono
parents:
diff changeset
1477 hsa_op_address *addr;
kono
parents:
diff changeset
1478 hsa_op_reg *reg;
kono
parents:
diff changeset
1479
kono
parents:
diff changeset
1480 /* Iterate all register operands and verify that the instruction
kono
parents:
diff changeset
1481 is set in uses of the register. */
kono
parents:
diff changeset
1482 for (unsigned i = 0; i < operand_count (); i++)
kono
parents:
diff changeset
1483 {
kono
parents:
diff changeset
1484 hsa_op_base *use = get_op (i);
kono
parents:
diff changeset
1485
kono
parents:
diff changeset
1486 if ((addr = dyn_cast <hsa_op_address *> (use)) && addr->m_reg)
kono
parents:
diff changeset
1487 {
kono
parents:
diff changeset
1488 gcc_assert (addr->m_reg->m_def_insn != this);
kono
parents:
diff changeset
1489 use = addr->m_reg;
kono
parents:
diff changeset
1490 }
kono
parents:
diff changeset
1491
kono
parents:
diff changeset
1492 if ((reg = dyn_cast <hsa_op_reg *> (use)) && !op_output_p (i))
kono
parents:
diff changeset
1493 {
kono
parents:
diff changeset
1494 unsigned j;
kono
parents:
diff changeset
1495 for (j = 0; j < reg->m_uses.length (); j++)
kono
parents:
diff changeset
1496 {
kono
parents:
diff changeset
1497 if (reg->m_uses[j] == this)
kono
parents:
diff changeset
1498 break;
kono
parents:
diff changeset
1499 }
kono
parents:
diff changeset
1500
kono
parents:
diff changeset
1501 if (j == reg->m_uses.length ())
kono
parents:
diff changeset
1502 {
kono
parents:
diff changeset
1503 error ("HSA instruction uses a register but is not among "
kono
parents:
diff changeset
1504 "recorded register uses");
kono
parents:
diff changeset
1505 debug_hsa_operand (reg);
kono
parents:
diff changeset
1506 debug_hsa_insn (this);
kono
parents:
diff changeset
1507 internal_error ("HSA instruction verification failed");
kono
parents:
diff changeset
1508 }
kono
parents:
diff changeset
1509 }
kono
parents:
diff changeset
1510 }
kono
parents:
diff changeset
1511 }
kono
parents:
diff changeset
1512
kono
parents:
diff changeset
1513 /* Constructor of an instruction representing a PHI node. NOPS is the number
kono
parents:
diff changeset
1514 of operands (equal to the number of predecessors). */
kono
parents:
diff changeset
1515
kono
parents:
diff changeset
1516 hsa_insn_phi::hsa_insn_phi (unsigned nops, hsa_op_reg *dst)
kono
parents:
diff changeset
1517 : hsa_insn_basic (nops, HSA_OPCODE_PHI), m_dest (dst)
kono
parents:
diff changeset
1518 {
kono
parents:
diff changeset
1519 dst->set_definition (this);
kono
parents:
diff changeset
1520 }
kono
parents:
diff changeset
1521
kono
parents:
diff changeset
1522 /* Constructor of class representing instructions for control flow and
kono
parents:
diff changeset
1523 sychronization, */
kono
parents:
diff changeset
1524
kono
parents:
diff changeset
1525 hsa_insn_br::hsa_insn_br (unsigned nops, int opc, BrigType16_t t,
kono
parents:
diff changeset
1526 BrigWidth8_t width, hsa_op_base *arg0,
kono
parents:
diff changeset
1527 hsa_op_base *arg1, hsa_op_base *arg2,
kono
parents:
diff changeset
1528 hsa_op_base *arg3)
kono
parents:
diff changeset
1529 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
kono
parents:
diff changeset
1530 m_width (width)
kono
parents:
diff changeset
1531 {
kono
parents:
diff changeset
1532 }
kono
parents:
diff changeset
1533
kono
parents:
diff changeset
1534 /* Constructor of class representing instruction for conditional jump, CTRL is
kono
parents:
diff changeset
1535 the control register determining whether the jump will be carried out, the
kono
parents:
diff changeset
1536 new instruction is automatically added to its uses list. */
kono
parents:
diff changeset
1537
kono
parents:
diff changeset
1538 hsa_insn_cbr::hsa_insn_cbr (hsa_op_reg *ctrl)
kono
parents:
diff changeset
1539 : hsa_insn_br (1, BRIG_OPCODE_CBR, BRIG_TYPE_B1, BRIG_WIDTH_1, ctrl)
kono
parents:
diff changeset
1540 {
kono
parents:
diff changeset
1541 }
kono
parents:
diff changeset
1542
kono
parents:
diff changeset
1543 /* Constructor of class representing instruction for switch jump, CTRL is
kono
parents:
diff changeset
1544 the index register. */
kono
parents:
diff changeset
1545
kono
parents:
diff changeset
1546 hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg *index, unsigned jump_count)
kono
parents:
diff changeset
1547 : hsa_insn_basic (1, BRIG_OPCODE_SBR, BRIG_TYPE_B1, index),
kono
parents:
diff changeset
1548 m_width (BRIG_WIDTH_1), m_jump_table (vNULL),
kono
parents:
diff changeset
1549 m_label_code_list (new hsa_op_code_list (jump_count))
kono
parents:
diff changeset
1550 {
kono
parents:
diff changeset
1551 }
kono
parents:
diff changeset
1552
kono
parents:
diff changeset
1553 /* Replace all occurrences of OLD_BB with NEW_BB in the statements
kono
parents:
diff changeset
1554 jump table. */
kono
parents:
diff changeset
1555
kono
parents:
diff changeset
1556 void
kono
parents:
diff changeset
1557 hsa_insn_sbr::replace_all_labels (basic_block old_bb, basic_block new_bb)
kono
parents:
diff changeset
1558 {
kono
parents:
diff changeset
1559 for (unsigned i = 0; i < m_jump_table.length (); i++)
kono
parents:
diff changeset
1560 if (m_jump_table[i] == old_bb)
kono
parents:
diff changeset
1561 m_jump_table[i] = new_bb;
kono
parents:
diff changeset
1562 }
kono
parents:
diff changeset
1563
kono
parents:
diff changeset
1564 hsa_insn_sbr::~hsa_insn_sbr ()
kono
parents:
diff changeset
1565 {
kono
parents:
diff changeset
1566 m_jump_table.release ();
kono
parents:
diff changeset
1567 }
kono
parents:
diff changeset
1568
kono
parents:
diff changeset
1569 /* Constructor of comparison instruction. CMP is the comparison operation and T
kono
parents:
diff changeset
1570 is the result type. */
kono
parents:
diff changeset
1571
kono
parents:
diff changeset
1572 hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp, BrigType16_t t,
kono
parents:
diff changeset
1573 hsa_op_base *arg0, hsa_op_base *arg1,
kono
parents:
diff changeset
1574 hsa_op_base *arg2)
kono
parents:
diff changeset
1575 : hsa_insn_basic (3 , BRIG_OPCODE_CMP, t, arg0, arg1, arg2), m_compare (cmp)
kono
parents:
diff changeset
1576 {
kono
parents:
diff changeset
1577 }
kono
parents:
diff changeset
1578
kono
parents:
diff changeset
1579 /* Constructor of classes representing memory accesses. OPC is the opcode (must
kono
parents:
diff changeset
1580 be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
kono
parents:
diff changeset
1581 operands are provided as ARG0 and ARG1. */
kono
parents:
diff changeset
1582
kono
parents:
diff changeset
1583 hsa_insn_mem::hsa_insn_mem (int opc, BrigType16_t t, hsa_op_base *arg0,
kono
parents:
diff changeset
1584 hsa_op_base *arg1)
kono
parents:
diff changeset
1585 : hsa_insn_basic (2, opc, t, arg0, arg1),
kono
parents:
diff changeset
1586 m_align (hsa_natural_alignment (t)), m_equiv_class (0)
kono
parents:
diff changeset
1587 {
kono
parents:
diff changeset
1588 gcc_checking_assert (opc == BRIG_OPCODE_LD || opc == BRIG_OPCODE_ST);
kono
parents:
diff changeset
1589 }
kono
parents:
diff changeset
1590
kono
parents:
diff changeset
1591 /* Constructor for descendants allowing different opcodes and number of
kono
parents:
diff changeset
1592 operands, it passes its arguments directly to hsa_insn_basic
kono
parents:
diff changeset
1593 constructor. The instruction operands are provided as ARG[0-3]. */
kono
parents:
diff changeset
1594
kono
parents:
diff changeset
1595
kono
parents:
diff changeset
1596 hsa_insn_mem::hsa_insn_mem (unsigned nops, int opc, BrigType16_t t,
kono
parents:
diff changeset
1597 hsa_op_base *arg0, hsa_op_base *arg1,
kono
parents:
diff changeset
1598 hsa_op_base *arg2, hsa_op_base *arg3)
kono
parents:
diff changeset
1599 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
kono
parents:
diff changeset
1600 m_align (hsa_natural_alignment (t)), m_equiv_class (0)
kono
parents:
diff changeset
1601 {
kono
parents:
diff changeset
1602 }
kono
parents:
diff changeset
1603
kono
parents:
diff changeset
1604 /* Constructor of class representing atomic instructions. OPC is the principal
kono
parents:
diff changeset
1605 opcode, AOP is the specific atomic operation opcode. T is the type of the
kono
parents:
diff changeset
1606 instruction. The instruction operands are provided as ARG[0-3]. */
kono
parents:
diff changeset
1607
kono
parents:
diff changeset
1608 hsa_insn_atomic::hsa_insn_atomic (int nops, int opc,
kono
parents:
diff changeset
1609 enum BrigAtomicOperation aop,
kono
parents:
diff changeset
1610 BrigType16_t t, BrigMemoryOrder memorder,
kono
parents:
diff changeset
1611 hsa_op_base *arg0,
kono
parents:
diff changeset
1612 hsa_op_base *arg1, hsa_op_base *arg2,
kono
parents:
diff changeset
1613 hsa_op_base *arg3)
kono
parents:
diff changeset
1614 : hsa_insn_mem (nops, opc, t, arg0, arg1, arg2, arg3), m_atomicop (aop),
kono
parents:
diff changeset
1615 m_memoryorder (memorder),
kono
parents:
diff changeset
1616 m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM)
kono
parents:
diff changeset
1617 {
kono
parents:
diff changeset
1618 gcc_checking_assert (opc == BRIG_OPCODE_ATOMICNORET ||
kono
parents:
diff changeset
1619 opc == BRIG_OPCODE_ATOMIC ||
kono
parents:
diff changeset
1620 opc == BRIG_OPCODE_SIGNAL ||
kono
parents:
diff changeset
1621 opc == BRIG_OPCODE_SIGNALNORET);
kono
parents:
diff changeset
1622 }
kono
parents:
diff changeset
1623
kono
parents:
diff changeset
1624 /* Constructor of class representing signal instructions. OPC is the prinicpal
kono
parents:
diff changeset
1625 opcode, SOP is the specific signal operation opcode. T is the type of the
kono
parents:
diff changeset
1626 instruction. The instruction operands are provided as ARG[0-3]. */
kono
parents:
diff changeset
1627
kono
parents:
diff changeset
1628 hsa_insn_signal::hsa_insn_signal (int nops, int opc,
kono
parents:
diff changeset
1629 enum BrigAtomicOperation sop,
kono
parents:
diff changeset
1630 BrigType16_t t, BrigMemoryOrder memorder,
kono
parents:
diff changeset
1631 hsa_op_base *arg0, hsa_op_base *arg1,
kono
parents:
diff changeset
1632 hsa_op_base *arg2, hsa_op_base *arg3)
kono
parents:
diff changeset
1633 : hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
kono
parents:
diff changeset
1634 m_memory_order (memorder), m_signalop (sop)
kono
parents:
diff changeset
1635 {
kono
parents:
diff changeset
1636 }
kono
parents:
diff changeset
1637
kono
parents:
diff changeset
1638 /* Constructor of class representing segment conversion instructions. OPC is
kono
parents:
diff changeset
1639 the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
kono
parents:
diff changeset
1640 and SRCT are destination and source types respectively, SEG is the segment
kono
parents:
diff changeset
1641 we are converting to or from. The instruction operands are
kono
parents:
diff changeset
1642 provided as ARG0 and ARG1. */
kono
parents:
diff changeset
1643
kono
parents:
diff changeset
1644 hsa_insn_seg::hsa_insn_seg (int opc, BrigType16_t dest, BrigType16_t srct,
kono
parents:
diff changeset
1645 BrigSegment8_t seg, hsa_op_base *arg0,
kono
parents:
diff changeset
1646 hsa_op_base *arg1)
kono
parents:
diff changeset
1647 : hsa_insn_basic (2, opc, dest, arg0, arg1), m_src_type (srct),
kono
parents:
diff changeset
1648 m_segment (seg)
kono
parents:
diff changeset
1649 {
kono
parents:
diff changeset
1650 gcc_checking_assert (opc == BRIG_OPCODE_STOF || opc == BRIG_OPCODE_FTOS);
kono
parents:
diff changeset
1651 }
kono
parents:
diff changeset
1652
kono
parents:
diff changeset
1653 /* Constructor of class representing a call instruction. CALLEE is the tree
kono
parents:
diff changeset
1654 representation of the function being called. */
kono
parents:
diff changeset
1655
kono
parents:
diff changeset
1656 hsa_insn_call::hsa_insn_call (tree callee)
kono
parents:
diff changeset
1657 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (callee),
kono
parents:
diff changeset
1658 m_output_arg (NULL), m_args_code_list (NULL), m_result_code_list (NULL)
kono
parents:
diff changeset
1659 {
kono
parents:
diff changeset
1660 }
kono
parents:
diff changeset
1661
kono
parents:
diff changeset
1662 hsa_insn_call::hsa_insn_call (hsa_internal_fn *fn)
kono
parents:
diff changeset
1663 : hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (NULL),
kono
parents:
diff changeset
1664 m_called_internal_fn (fn), m_output_arg (NULL), m_args_code_list (NULL),
kono
parents:
diff changeset
1665 m_result_code_list (NULL)
kono
parents:
diff changeset
1666 {
kono
parents:
diff changeset
1667 }
kono
parents:
diff changeset
1668
kono
parents:
diff changeset
1669 hsa_insn_call::~hsa_insn_call ()
kono
parents:
diff changeset
1670 {
kono
parents:
diff changeset
1671 for (unsigned i = 0; i < m_input_args.length (); i++)
kono
parents:
diff changeset
1672 delete m_input_args[i];
kono
parents:
diff changeset
1673
kono
parents:
diff changeset
1674 delete m_output_arg;
kono
parents:
diff changeset
1675
kono
parents:
diff changeset
1676 m_input_args.release ();
kono
parents:
diff changeset
1677 m_input_arg_insns.release ();
kono
parents:
diff changeset
1678 }
kono
parents:
diff changeset
1679
kono
parents:
diff changeset
1680 /* Constructor of class representing the argument block required to invoke
kono
parents:
diff changeset
1681 a call in HSAIL. */
kono
parents:
diff changeset
1682 hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind,
kono
parents:
diff changeset
1683 hsa_insn_call * call)
kono
parents:
diff changeset
1684 : hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK), m_kind (brig_kind),
kono
parents:
diff changeset
1685 m_call_insn (call)
kono
parents:
diff changeset
1686 {
kono
parents:
diff changeset
1687 }
kono
parents:
diff changeset
1688
kono
parents:
diff changeset
1689 hsa_insn_comment::hsa_insn_comment (const char *s)
kono
parents:
diff changeset
1690 : hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT)
kono
parents:
diff changeset
1691 {
kono
parents:
diff changeset
1692 unsigned l = strlen (s);
kono
parents:
diff changeset
1693
kono
parents:
diff changeset
1694 /* Append '// ' to the string. */
kono
parents:
diff changeset
1695 char *buf = XNEWVEC (char, l + 4);
kono
parents:
diff changeset
1696 sprintf (buf, "// %s", s);
kono
parents:
diff changeset
1697 m_comment = buf;
kono
parents:
diff changeset
1698 }
kono
parents:
diff changeset
1699
kono
parents:
diff changeset
1700 hsa_insn_comment::~hsa_insn_comment ()
kono
parents:
diff changeset
1701 {
kono
parents:
diff changeset
1702 gcc_checking_assert (m_comment);
kono
parents:
diff changeset
1703 free (m_comment);
kono
parents:
diff changeset
1704 m_comment = NULL;
kono
parents:
diff changeset
1705 }
kono
parents:
diff changeset
1706
kono
parents:
diff changeset
1707 /* Constructor of class representing the queue instruction in HSAIL. */
kono
parents:
diff changeset
1708
kono
parents:
diff changeset
1709 hsa_insn_queue::hsa_insn_queue (int nops, int opcode, BrigSegment segment,
kono
parents:
diff changeset
1710 BrigMemoryOrder memory_order,
kono
parents:
diff changeset
1711 hsa_op_base *arg0, hsa_op_base *arg1,
kono
parents:
diff changeset
1712 hsa_op_base *arg2, hsa_op_base *arg3)
kono
parents:
diff changeset
1713 : hsa_insn_basic (nops, opcode, BRIG_TYPE_U64, arg0, arg1, arg2, arg3),
kono
parents:
diff changeset
1714 m_segment (segment), m_memory_order (memory_order)
kono
parents:
diff changeset
1715 {
kono
parents:
diff changeset
1716 }
kono
parents:
diff changeset
1717
kono
parents:
diff changeset
1718 /* Constructor of class representing the source type instruction in HSAIL. */
kono
parents:
diff changeset
1719
kono
parents:
diff changeset
1720 hsa_insn_srctype::hsa_insn_srctype (int nops, BrigOpcode opcode,
kono
parents:
diff changeset
1721 BrigType16_t destt, BrigType16_t srct,
kono
parents:
diff changeset
1722 hsa_op_base *arg0, hsa_op_base *arg1,
kono
parents:
diff changeset
1723 hsa_op_base *arg2 = NULL)
kono
parents:
diff changeset
1724 : hsa_insn_basic (nops, opcode, destt, arg0, arg1, arg2),
kono
parents:
diff changeset
1725 m_source_type (srct)
kono
parents:
diff changeset
1726 {}
kono
parents:
diff changeset
1727
kono
parents:
diff changeset
1728 /* Constructor of class representing the packed instruction in HSAIL. */
kono
parents:
diff changeset
1729
kono
parents:
diff changeset
1730 hsa_insn_packed::hsa_insn_packed (int nops, BrigOpcode opcode,
kono
parents:
diff changeset
1731 BrigType16_t destt, BrigType16_t srct,
kono
parents:
diff changeset
1732 hsa_op_base *arg0, hsa_op_base *arg1,
kono
parents:
diff changeset
1733 hsa_op_base *arg2)
kono
parents:
diff changeset
1734 : hsa_insn_srctype (nops, opcode, destt, srct, arg0, arg1, arg2)
kono
parents:
diff changeset
1735 {
kono
parents:
diff changeset
1736 m_operand_list = new hsa_op_operand_list (nops - 1);
kono
parents:
diff changeset
1737 }
kono
parents:
diff changeset
1738
kono
parents:
diff changeset
1739 /* Constructor of class representing the convert instruction in HSAIL. */
kono
parents:
diff changeset
1740
kono
parents:
diff changeset
1741 hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type *dest, hsa_op_with_type *src)
kono
parents:
diff changeset
1742 : hsa_insn_basic (2, BRIG_OPCODE_CVT, dest->m_type, dest, src)
kono
parents:
diff changeset
1743 {
kono
parents:
diff changeset
1744 }
kono
parents:
diff changeset
1745
kono
parents:
diff changeset
1746 /* Constructor of class representing the alloca in HSAIL. */
kono
parents:
diff changeset
1747
kono
parents:
diff changeset
1748 hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type *dest,
kono
parents:
diff changeset
1749 hsa_op_with_type *size, unsigned alignment)
kono
parents:
diff changeset
1750 : hsa_insn_basic (2, BRIG_OPCODE_ALLOCA, dest->m_type, dest, size),
kono
parents:
diff changeset
1751 m_align (BRIG_ALIGNMENT_8)
kono
parents:
diff changeset
1752 {
kono
parents:
diff changeset
1753 gcc_assert (dest->m_type == BRIG_TYPE_U32);
kono
parents:
diff changeset
1754 if (alignment)
kono
parents:
diff changeset
1755 m_align = hsa_alignment_encoding (alignment);
kono
parents:
diff changeset
1756 }
kono
parents:
diff changeset
1757
kono
parents:
diff changeset
1758 /* Append an instruction INSN into the basic block. */
kono
parents:
diff changeset
1759
kono
parents:
diff changeset
1760 void
kono
parents:
diff changeset
1761 hsa_bb::append_insn (hsa_insn_basic *insn)
kono
parents:
diff changeset
1762 {
kono
parents:
diff changeset
1763 gcc_assert (insn->m_opcode != 0 || insn->operand_count () == 0);
kono
parents:
diff changeset
1764 gcc_assert (!insn->m_bb);
kono
parents:
diff changeset
1765
kono
parents:
diff changeset
1766 insn->m_bb = m_bb;
kono
parents:
diff changeset
1767 insn->m_prev = m_last_insn;
kono
parents:
diff changeset
1768 insn->m_next = NULL;
kono
parents:
diff changeset
1769 if (m_last_insn)
kono
parents:
diff changeset
1770 m_last_insn->m_next = insn;
kono
parents:
diff changeset
1771 m_last_insn = insn;
kono
parents:
diff changeset
1772 if (!m_first_insn)
kono
parents:
diff changeset
1773 m_first_insn = insn;
kono
parents:
diff changeset
1774 }
kono
parents:
diff changeset
1775
kono
parents:
diff changeset
1776 void
kono
parents:
diff changeset
1777 hsa_bb::append_phi (hsa_insn_phi *hphi)
kono
parents:
diff changeset
1778 {
kono
parents:
diff changeset
1779 hphi->m_bb = m_bb;
kono
parents:
diff changeset
1780
kono
parents:
diff changeset
1781 hphi->m_prev = m_last_phi;
kono
parents:
diff changeset
1782 hphi->m_next = NULL;
kono
parents:
diff changeset
1783 if (m_last_phi)
kono
parents:
diff changeset
1784 m_last_phi->m_next = hphi;
kono
parents:
diff changeset
1785 m_last_phi = hphi;
kono
parents:
diff changeset
1786 if (!m_first_phi)
kono
parents:
diff changeset
1787 m_first_phi = hphi;
kono
parents:
diff changeset
1788 }
kono
parents:
diff changeset
1789
kono
parents:
diff changeset
1790 /* Insert HSA instruction NEW_INSN immediately before an existing instruction
kono
parents:
diff changeset
1791 OLD_INSN. */
kono
parents:
diff changeset
1792
kono
parents:
diff changeset
1793 static void
kono
parents:
diff changeset
1794 hsa_insert_insn_before (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
kono
parents:
diff changeset
1795 {
kono
parents:
diff changeset
1796 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
kono
parents:
diff changeset
1797
kono
parents:
diff changeset
1798 if (hbb->m_first_insn == old_insn)
kono
parents:
diff changeset
1799 hbb->m_first_insn = new_insn;
kono
parents:
diff changeset
1800 new_insn->m_prev = old_insn->m_prev;
kono
parents:
diff changeset
1801 new_insn->m_next = old_insn;
kono
parents:
diff changeset
1802 if (old_insn->m_prev)
kono
parents:
diff changeset
1803 old_insn->m_prev->m_next = new_insn;
kono
parents:
diff changeset
1804 old_insn->m_prev = new_insn;
kono
parents:
diff changeset
1805 }
kono
parents:
diff changeset
1806
kono
parents:
diff changeset
1807 /* Append HSA instruction NEW_INSN immediately after an existing instruction
kono
parents:
diff changeset
1808 OLD_INSN. */
kono
parents:
diff changeset
1809
kono
parents:
diff changeset
1810 static void
kono
parents:
diff changeset
1811 hsa_append_insn_after (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
kono
parents:
diff changeset
1812 {
kono
parents:
diff changeset
1813 hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
kono
parents:
diff changeset
1814
kono
parents:
diff changeset
1815 if (hbb->m_last_insn == old_insn)
kono
parents:
diff changeset
1816 hbb->m_last_insn = new_insn;
kono
parents:
diff changeset
1817 new_insn->m_prev = old_insn;
kono
parents:
diff changeset
1818 new_insn->m_next = old_insn->m_next;
kono
parents:
diff changeset
1819 if (old_insn->m_next)
kono
parents:
diff changeset
1820 old_insn->m_next->m_prev = new_insn;
kono
parents:
diff changeset
1821 old_insn->m_next = new_insn;
kono
parents:
diff changeset
1822 }
kono
parents:
diff changeset
1823
kono
parents:
diff changeset
1824 /* Return a register containing the calculated value of EXP which must be an
kono
parents:
diff changeset
1825 expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
kono
parents:
diff changeset
1826 integer constants as returned by get_inner_reference.
kono
parents:
diff changeset
1827 Newly generated HSA instructions will be appended to HBB.
kono
parents:
diff changeset
1828 Perform all calculations in ADDRTYPE. */
kono
parents:
diff changeset
1829
kono
parents:
diff changeset
1830 static hsa_op_with_type *
kono
parents:
diff changeset
1831 gen_address_calculation (tree exp, hsa_bb *hbb, BrigType16_t addrtype)
kono
parents:
diff changeset
1832 {
kono
parents:
diff changeset
1833 int opcode;
kono
parents:
diff changeset
1834
kono
parents:
diff changeset
1835 if (TREE_CODE (exp) == NOP_EXPR)
kono
parents:
diff changeset
1836 exp = TREE_OPERAND (exp, 0);
kono
parents:
diff changeset
1837
kono
parents:
diff changeset
1838 switch (TREE_CODE (exp))
kono
parents:
diff changeset
1839 {
kono
parents:
diff changeset
1840 case SSA_NAME:
kono
parents:
diff changeset
1841 return hsa_cfun->reg_for_gimple_ssa (exp)->get_in_type (addrtype, hbb);
kono
parents:
diff changeset
1842
kono
parents:
diff changeset
1843 case INTEGER_CST:
kono
parents:
diff changeset
1844 {
kono
parents:
diff changeset
1845 hsa_op_immed *imm = new hsa_op_immed (exp);
kono
parents:
diff changeset
1846 if (addrtype != imm->m_type)
kono
parents:
diff changeset
1847 imm->m_type = addrtype;
kono
parents:
diff changeset
1848 return imm;
kono
parents:
diff changeset
1849 }
kono
parents:
diff changeset
1850
kono
parents:
diff changeset
1851 case PLUS_EXPR:
kono
parents:
diff changeset
1852 opcode = BRIG_OPCODE_ADD;
kono
parents:
diff changeset
1853 break;
kono
parents:
diff changeset
1854
kono
parents:
diff changeset
1855 case MULT_EXPR:
kono
parents:
diff changeset
1856 opcode = BRIG_OPCODE_MUL;
kono
parents:
diff changeset
1857 break;
kono
parents:
diff changeset
1858
kono
parents:
diff changeset
1859 default:
kono
parents:
diff changeset
1860 gcc_unreachable ();
kono
parents:
diff changeset
1861 }
kono
parents:
diff changeset
1862
kono
parents:
diff changeset
1863 hsa_op_reg *res = new hsa_op_reg (addrtype);
kono
parents:
diff changeset
1864 hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, addrtype);
kono
parents:
diff changeset
1865 insn->set_op (0, res);
kono
parents:
diff changeset
1866
kono
parents:
diff changeset
1867 hsa_op_with_type *op1 = gen_address_calculation (TREE_OPERAND (exp, 0), hbb,
kono
parents:
diff changeset
1868 addrtype);
kono
parents:
diff changeset
1869 hsa_op_with_type *op2 = gen_address_calculation (TREE_OPERAND (exp, 1), hbb,
kono
parents:
diff changeset
1870 addrtype);
kono
parents:
diff changeset
1871 insn->set_op (1, op1);
kono
parents:
diff changeset
1872 insn->set_op (2, op2);
kono
parents:
diff changeset
1873
kono
parents:
diff changeset
1874 hbb->append_insn (insn);
kono
parents:
diff changeset
1875 return res;
kono
parents:
diff changeset
1876 }
kono
parents:
diff changeset
1877
kono
parents:
diff changeset
1878 /* If R1 is NULL, just return R2, otherwise append an instruction adding them
kono
parents:
diff changeset
1879 to HBB and return the register holding the result. */
kono
parents:
diff changeset
1880
kono
parents:
diff changeset
1881 static hsa_op_reg *
kono
parents:
diff changeset
1882 add_addr_regs_if_needed (hsa_op_reg *r1, hsa_op_reg *r2, hsa_bb *hbb)
kono
parents:
diff changeset
1883 {
kono
parents:
diff changeset
1884 gcc_checking_assert (r2);
kono
parents:
diff changeset
1885 if (!r1)
kono
parents:
diff changeset
1886 return r2;
kono
parents:
diff changeset
1887
kono
parents:
diff changeset
1888 hsa_op_reg *res = new hsa_op_reg (r1->m_type);
kono
parents:
diff changeset
1889 gcc_assert (!hsa_needs_cvt (r1->m_type, r2->m_type));
kono
parents:
diff changeset
1890 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_ADD, res->m_type);
kono
parents:
diff changeset
1891 insn->set_op (0, res);
kono
parents:
diff changeset
1892 insn->set_op (1, r1);
kono
parents:
diff changeset
1893 insn->set_op (2, r2);
kono
parents:
diff changeset
1894 hbb->append_insn (insn);
kono
parents:
diff changeset
1895 return res;
kono
parents:
diff changeset
1896 }
kono
parents:
diff changeset
1897
kono
parents:
diff changeset
1898 /* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
kono
parents:
diff changeset
1899 reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
kono
parents:
diff changeset
1900
kono
parents:
diff changeset
1901 static void
kono
parents:
diff changeset
1902 process_mem_base (tree base, hsa_symbol **symbol, BrigType16_t *addrtype,
kono
parents:
diff changeset
1903 hsa_op_reg **reg, offset_int *offset, hsa_bb *hbb)
kono
parents:
diff changeset
1904 {
kono
parents:
diff changeset
1905 if (TREE_CODE (base) == SSA_NAME)
kono
parents:
diff changeset
1906 {
kono
parents:
diff changeset
1907 gcc_assert (!*reg);
kono
parents:
diff changeset
1908 hsa_op_with_type *ssa
kono
parents:
diff changeset
1909 = hsa_cfun->reg_for_gimple_ssa (base)->get_in_type (*addrtype, hbb);
kono
parents:
diff changeset
1910 *reg = dyn_cast <hsa_op_reg *> (ssa);
kono
parents:
diff changeset
1911 }
kono
parents:
diff changeset
1912 else if (TREE_CODE (base) == ADDR_EXPR)
kono
parents:
diff changeset
1913 {
kono
parents:
diff changeset
1914 tree decl = TREE_OPERAND (base, 0);
kono
parents:
diff changeset
1915
kono
parents:
diff changeset
1916 if (!DECL_P (decl) || TREE_CODE (decl) == FUNCTION_DECL)
kono
parents:
diff changeset
1917 {
kono
parents:
diff changeset
1918 HSA_SORRY_AT (EXPR_LOCATION (base),
kono
parents:
diff changeset
1919 "support for HSA does not implement a memory reference "
kono
parents:
diff changeset
1920 "to a non-declaration type");
kono
parents:
diff changeset
1921 return;
kono
parents:
diff changeset
1922 }
kono
parents:
diff changeset
1923
kono
parents:
diff changeset
1924 gcc_assert (!*symbol);
kono
parents:
diff changeset
1925
kono
parents:
diff changeset
1926 *symbol = get_symbol_for_decl (decl);
kono
parents:
diff changeset
1927 *addrtype = hsa_get_segment_addr_type ((*symbol)->m_segment);
kono
parents:
diff changeset
1928 }
kono
parents:
diff changeset
1929 else if (TREE_CODE (base) == INTEGER_CST)
kono
parents:
diff changeset
1930 *offset += wi::to_offset (base);
kono
parents:
diff changeset
1931 else
kono
parents:
diff changeset
1932 gcc_unreachable ();
kono
parents:
diff changeset
1933 }
kono
parents:
diff changeset
1934
kono
parents:
diff changeset
1935 /* Forward declaration of a function. */
kono
parents:
diff changeset
1936
kono
parents:
diff changeset
1937 static void
kono
parents:
diff changeset
1938 gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb);
kono
parents:
diff changeset
1939
kono
parents:
diff changeset
1940 /* Generate HSA address operand for a given tree memory reference REF. If
kono
parents:
diff changeset
1941 instructions need to be created to calculate the address, they will be added
kono
parents:
diff changeset
1942 to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
kono
parents:
diff changeset
1943 the function assumes that the caller will handle possible
kono
parents:
diff changeset
1944 bit-field references. Otherwise if we reference a bit-field, sorry message
kono
parents:
diff changeset
1945 is displayed. */
kono
parents:
diff changeset
1946
kono
parents:
diff changeset
1947 static hsa_op_address *
kono
parents:
diff changeset
1948 gen_hsa_addr (tree ref, hsa_bb *hbb, HOST_WIDE_INT *output_bitsize = NULL,
kono
parents:
diff changeset
1949 HOST_WIDE_INT *output_bitpos = NULL)
kono
parents:
diff changeset
1950 {
kono
parents:
diff changeset
1951 hsa_symbol *symbol = NULL;
kono
parents:
diff changeset
1952 hsa_op_reg *reg = NULL;
kono
parents:
diff changeset
1953 offset_int offset = 0;
kono
parents:
diff changeset
1954 tree origref = ref;
kono
parents:
diff changeset
1955 tree varoffset = NULL_TREE;
kono
parents:
diff changeset
1956 BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
kono
parents:
diff changeset
1957 HOST_WIDE_INT bitsize = 0, bitpos = 0;
kono
parents:
diff changeset
1958 BrigType16_t flat_addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
kono
parents:
diff changeset
1959
kono
parents:
diff changeset
1960 if (TREE_CODE (ref) == STRING_CST)
kono
parents:
diff changeset
1961 {
kono
parents:
diff changeset
1962 symbol = hsa_get_string_cst_symbol (ref);
kono
parents:
diff changeset
1963 goto out;
kono
parents:
diff changeset
1964 }
kono
parents:
diff changeset
1965 else if (TREE_CODE (ref) == BIT_FIELD_REF
131
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1966 && (!multiple_p (bit_field_size (ref), BITS_PER_UNIT)
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1967 || !multiple_p (bit_field_offset (ref), BITS_PER_UNIT)))
111
kono
parents:
diff changeset
1968 {
kono
parents:
diff changeset
1969 HSA_SORRY_ATV (EXPR_LOCATION (origref),
kono
parents:
diff changeset
1970 "support for HSA does not implement "
kono
parents:
diff changeset
1971 "bit field references such as %E", ref);
kono
parents:
diff changeset
1972 goto out;
kono
parents:
diff changeset
1973 }
kono
parents:
diff changeset
1974
kono
parents:
diff changeset
1975 if (handled_component_p (ref))
kono
parents:
diff changeset
1976 {
kono
parents:
diff changeset
1977 machine_mode mode;
kono
parents:
diff changeset
1978 int unsignedp, volatilep, preversep;
131
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1979 poly_int64 pbitsize, pbitpos;
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1980 tree new_ref;
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1981
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1982 new_ref = get_inner_reference (ref, &pbitsize, &pbitpos, &varoffset,
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1983 &mode, &unsignedp, &preversep,
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1984 &volatilep);
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1985 /* When this isn't true, the switch below will report an
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1986 appropriate error. */
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1987 if (pbitsize.is_constant () && pbitpos.is_constant ())
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1988 {
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1989 bitsize = pbitsize.to_constant ();
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1990 bitpos = pbitpos.to_constant ();
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1991 ref = new_ref;
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1992 offset = bitpos;
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1993 offset = wi::rshift (offset, LOG2_BITS_PER_UNIT, SIGNED);
84e7813d76e9 gcc-8.2
mir3636
parents: 111
diff changeset
1994 }
111
kono
parents:
diff changeset
1995 }
kono
parents:
diff changeset
1996
kono
parents:
diff changeset
1997 switch (TREE_CODE (ref))
kono
parents:
diff changeset
1998 {
kono
parents:
diff changeset
1999 case ADDR_EXPR:
kono
parents:
diff changeset
2000 {
kono
parents:
diff changeset
2001 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
kono
parents:
diff changeset
2002 symbol = hsa_cfun->create_hsa_temporary (flat_addrtype);
kono
parents:
diff changeset
2003 hsa_op_reg *r = new hsa_op_reg (flat_addrtype);
kono
parents:
diff changeset
2004 gen_hsa_addr_insns (ref, r, hbb);
kono
parents:
diff changeset
2005 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
kono
parents:
diff changeset
2006 r, new hsa_op_address (symbol)));
kono
parents:
diff changeset
2007
kono
parents:
diff changeset
2008 break;
kono
parents:
diff changeset
2009 }
kono
parents:
diff changeset
2010 case SSA_NAME:
kono
parents:
diff changeset
2011 {
kono
parents:
diff changeset
2012 addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
kono
parents:
diff changeset
2013 hsa_op_with_type *r = hsa_cfun->reg_for_gimple_ssa (ref);
kono
parents:
diff changeset
2014 if (r->m_type == BRIG_TYPE_B1)
kono
parents:
diff changeset
2015 r = r->get_in_type (BRIG_TYPE_U32, hbb);
kono
parents:
diff changeset
2016 symbol = hsa_cfun->create_hsa_temporary (r->m_type);
kono
parents:
diff changeset
2017
kono
parents:
diff changeset
2018 hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
kono
parents:
diff changeset
2019 r, new hsa_op_address (symbol)));
kono
parents:
diff changeset
2020
kono
parents:
diff changeset
2021 break;
kono
parents:
diff changeset
2022 }
kono
parents:
diff changeset
2023 case PARM_DECL:
kono
parents:
diff changeset
2024 case VAR_DECL:
kono
parents:
diff changeset
2025 case RESULT_DECL:
kono
parents:
diff changeset
2026 case CONST_DECL:
kono
parents:
diff changeset
2027 gcc_assert (!symbol);
kono
parents:
diff changeset
2028 symbol = get_symbol_for_decl (ref);
kono
parents:
diff changeset
2029 addrtype = hsa_get_segment_addr_type (symbol->m_segment);
kono
parents:
diff changeset
2030 break;
kono
parents:
diff changeset
2031
kono
parents:
diff changeset
2032 case MEM_REF:
kono
parents:
diff changeset
2033 process_mem_base (TREE_OPERAND (ref, 0), &symbol, &addrtype, &reg,
kono
parents:
diff changeset
2034 &offset, hbb);
kono
parents:
diff changeset
2035
kono
parents:
diff changeset
2036 if (!integer_zerop (TREE_OPERAND (ref, 1)))
kono
parents:
diff changeset
2037 offset += wi::to_offset (TREE_OPERAND (ref, 1));
kono
parents:
diff changeset
2038 break;
kono
parents:
diff changeset
2039
kono
parents:
diff changeset
2040 case TARGET_MEM_REF:
kono
parents:
diff changeset
2041 process_mem_base (TMR_BASE (ref), &symbol, &addrtype, &reg, &offset, hbb);
kono
parents:
diff changeset
2042 if (TMR_INDEX (ref))
kono
parents:
diff changeset
2043 {
kono
parents:
diff changeset
2044 hsa_op_reg *disp1;
kono
parents:
diff changeset
2045 hsa_op_base *idx = hsa_cfun->reg_for_gimple_ssa
kono
parents:
diff changeset
2046 (TMR_INDEX (ref))->get_in_type (addrtype, hbb);
kono
parents:
diff changeset
2047 if (TMR_STEP (ref) && !integer_onep (TMR_STEP (ref)))
kono
parents:
diff changeset
2048 {
kono
parents:
diff changeset
2049 disp1 = new hsa_op_reg (addrtype);
kono
parents:
diff changeset
2050 hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_MUL,
kono
parents:
diff changeset
2051 addrtype);
kono
parents:
diff changeset
2052
kono
parents:
diff changeset
2053 /* As step must respect addrtype, we overwrite the type
kono
parents:
diff changeset
2054 of an immediate value. */
kono
parents:
diff changeset
2055 hsa_op_immed *step = new hsa_op_immed (TMR_STEP (ref));
kono
parents:
diff changeset
2056 step->m_type = addrtype;
kono
parents:
diff changeset
2057
kono
parents:
diff changeset
2058 insn->set_op (0, disp1);
kono
parents:
diff changeset
2059 insn->set_op (1, idx);
kono
parents:
diff changeset
2060 insn->set_op (2, step);
kono
parents:
diff changeset
2061 hbb->append_insn (insn);
kono
parents:
diff changeset
2062 }
kono
parents:
diff changeset
2063 else
kono
parents:
diff changeset
2064 disp1 = as_a <hsa_op_reg *> (idx);
kono
parents:
diff changeset
2065 reg = add_addr_regs_if_needed (reg, disp1, hbb);
kono
parents:
diff changeset
2066 }
kono
parents:
diff changeset
2067 if (TMR_INDEX2 (ref))
kono
parents:
diff changeset
2068 {
kono
parents:
diff changeset
2069 if (TREE_CODE (TMR_INDEX2 (ref)) == SSA_NAME)
kono
parents:
diff changeset
2070 {
kono
parents:
diff changeset
2071 hsa_op_base *disp2 = hsa_cfun->reg_for_gimple_ssa
kono
parents:
diff changeset
2072 (TMR_INDEX2 (ref))->get_in_type (addrtype, hbb);
kono
parents:
diff changeset
2073 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (disp2),
kono
parents:
diff changeset
2074 hbb);
kono
parents:
diff changeset
2075 }
kono
parents:
diff changeset
2076 else if (TREE_CODE (TMR_INDEX2 (ref)) == INTEGER_CST)
kono
parents:
diff changeset
2077 offset += wi::to_offset (TMR_INDEX2 (ref));
kono
parents:
diff changeset
2078 else
kono
parents:
diff changeset
2079 gcc_unreachable ();
kono
parents:
diff changeset
2080 }
kono
parents:
diff changeset
2081 offset += wi::to_offset (TMR_OFFSET (ref));
kono
parents:
diff changeset
2082 break;
kono
parents:
diff changeset
2083 case FUNCTION_DECL:
kono
parents:
diff changeset
2084 HSA_SORRY_AT (EXPR_LOCATION (origref),
kono
parents:
diff changeset
2085 "support for HSA does not implement function pointers");
kono
parents:
diff changeset
2086 goto out;
kono
parents:
diff changeset
2087 default:
kono
parents:
diff changeset
2088 HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does "
kono
parents:
diff changeset
2089 "not implement memory access to %E", origref);
kono
parents:
diff changeset
2090 goto out;
kono
parents:
diff changeset
2091 }
kono
parents:
diff changeset
2092
kono
parents:
diff changeset
2093 if (varoffset)
kono
parents:
diff changeset
2094 {
kono
parents:
diff changeset
2095 if (TREE_CODE (varoffset) == INTEGER_CST)
kono
parents:
diff changeset
2096 offset += wi::to_offset (varoffset);
kono
parents:
diff changeset
2097 else
kono
parents:
diff changeset
2098 {
kono
parents:
diff changeset
2099 hsa_op_base *off_op = gen_address_calculation (varoffset, hbb,
kono
parents:
diff changeset
2100 addrtype);
kono
parents:
diff changeset
2101 reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (off_op),
kono
parents:
diff changeset
2102 hbb);
kono
parents:
diff changeset
2103 }
kono
parents:
diff changeset
2104 }
kono
parents:
diff changeset
2105
kono
parents:
diff changeset
2106 gcc_checking_assert ((symbol