This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [hsa merge 09/10] Majority of the HSA back-end
- From: Martin Jambor <mjambor at suse dot cz>
- To: Jakub Jelinek <jakub at redhat dot com>
- Cc: GCC Patches <gcc-patches at gcc dot gnu dot org>, Martin Liska <mliska at suse dot cz>
- Date: Fri, 15 Jan 2016 16:08:14 +0100
- Subject: Re: [hsa merge 09/10] Majority of the HSA back-end
- Authentication-results: sourceware.org; auth=none
- References: <20160113173925 dot 220029649 at virgil dot suse dot cz> <20160113173925 dot 914372077 at virgil dot suse dot cz> <20160114140533 dot GG3017 at tucnak dot redhat dot com>
Hi,
thanks Jakub. Below you'll find a patch, which is mostly work of
Martin Liska, that should address all the review comments. We have
then also went over the "XXX" marks (my bad that I forgot that Michael
uses this mark), removed half of them and turned the rest into TODOs.
Let me just quickly answer two comments as well:
On Thu, Jan 14, 2016 at 03:05:33PM +0100, Jakub Jelinek wrote:
> On Wed, Jan 13, 2016 at 06:39:34PM +0100, Martin Jambor wrote:
>
...
> > +#define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \
> > + "undefined semantics within target regions, support for HSA ignores them"
>
> Well, if you don't support them in HSA target regions, you'd better punt and
> not error on them.
We don't error, apart from issuing a warning we basically ignore them.
I believe we can do it even in the long term and that it is in fact
useful because the standard says that the "effect" if these routines
is "unspecified" if they get called from a target region.
Perhaps this is even something we should warn about earlier in omp
lowering/expansion.
...
> > +unsigned
> > +hsa_internal_fn::get_arity ()
> > +{
> > + switch (m_fn)
> > + {
> > + case IFN_ACOS:
> > + case IFN_ASIN:
> > + case IFN_ATAN:
> > + case IFN_COS:
> > + case IFN_EXP:
> > + case IFN_EXP10:
> > + case IFN_EXP2:
> > + case IFN_EXPM1:
> > + case IFN_LOG:
> > + case IFN_LOG10:
> > + case IFN_LOG1P:
> > + case IFN_LOG2:
> > + case IFN_LOGB:
> > + case IFN_SIGNIFICAND:
> > + case IFN_SIN:
> > + case IFN_SQRT:
> > + case IFN_TAN:
> > + case IFN_CEIL:
> > + case IFN_FLOOR:
> > + case IFN_NEARBYINT:
> > + case IFN_RINT:
> > + case IFN_ROUND:
> > + case IFN_TRUNC:
> > + return 1;
> > + case IFN_ATAN2:
> > + case IFN_COPYSIGN:
> > + case IFN_FMOD:
> > + case IFN_POW:
> > + case IFN_REMAINDER:
> > + case IFN_SCALB:
> > + case IFN_LDEXP:
> > + return 2;
> > + break;
> > + case IFN_CLRSB:
> > + case IFN_CLZ:
> > + case IFN_CTZ:
> > + case IFN_FFS:
> > + case IFN_PARITY:
> > + case IFN_POPCOUNT:
> > + default:
> > + gcc_unreachable ();
>
> There are various other IFNs (e.g. for __builtin_{add,sub,mul}_overflow,
> lots of others). How do you ensure you don't ICE on those?
Martin added a comment explaining this. This can only be reached when
we already know we are processing a known builtin, filtered by
gen_hsa_insn_for_internal_fn_call.
Thanks for looking at the code,
Martin
2016-01-15 Martin Liska <mliska@suse.cz>
Martin Jambor <mjambor@suse.cz>
* hsa-brig.c (struct function_linkage_pair): Fix GNU coding style
and replace sprintf with snprintf.
(hsa_brig_section::init): Likewise.
(hsa_brig_section::output): Likewise.
(hsa_brig_section::get_ptr_by_offset): Likewise.
(brig_string_slot_hasher::hash): Likewise.
(brig_string_slot_hasher::equal): Likewise.
(brig_string_slot_hasher::remove): Likewise.
(brig_emit_string): Likewise.
(brig_init): Likewise.
(emit_directive_variable): Likewise.
(emit_function_directives): Likewise.
(emit_bb_label_directive): Likewise.
(emit_immediate_scalar_to_buffer): Likewise.
(hsa_op_immed::emit_to_buffer): Likewise.
(emit_immediate_operand): Likewise.
(emit_address_operand): Likewise.
(emit_memory_insn): Likewise.
(emit_alloca_insn): Likewise.
(emit_cmp_insn): Likewise.
(emit_branch_insn): Likewise.
(emit_switch_insn): Likewise.
(emit_call_insn): Likewise.
(emit_arg_block_insn): Likewise.
(emit_packed_insn): Likewise.
(emit_basic_insn): Likewise.
(hsa_brig_emit_function): Likewise.
(hsa_output_global_variables): Likewise.
(hsa_output_kernels): Likewise.
(hsa_output_libgomp_mapping): Likewise.
(hsa_output_brig): Likewise.
* hsa-dump.c (dump_hsa_immed): Likewise.
(dump_hsa_insn_1): Likewise.
* hsa-gen.c (hsa_symbol::total_byte_size): Likewise.
(hsa_init_simple_builtins): Likewise.
(hsa_init_data_for_cfun): Likewise.
(hsa_type_for_scalar_tree_type): Likewise.
(get_symbol_for_decl): Likewise.
(hsa_get_host_function): Likewise.
(hsa_op_immed::hsa_op_immed): Likewise.
(hsa_insn_mem::hsa_insn_mem): Likewise.
(hsa_insn_atomic::hsa_insn_atomic): Likewise.
(hsa_insn_seg::hsa_insn_seg): Likewise.
(hsa_insn_srctype::hsa_insn_srctype): Likewise.
(process_mem_base): Likewise.
(gen_hsa_insns_for_bitfield): Likewise.
(gen_hsa_insns_for_load): Likewise.
(gen_hsa_insns_for_store): Likewise.
(gen_hsa_insns_for_operation_assignment): Likewise.
(gen_hsa_insns_for_switch_stmt): Likewise.
(get_format_argument_type): Likewise.
(gen_hsa_insns_for_direct_call): Likewise.
(gen_hsa_insns_for_call_of_internal_fn): Likewise.
(gen_hsa_insns_for_return): Likewise.
(query_hsa_grid): Likewise.
(gen_set_num_threads): Likewise.
(gen_num_threads_for_dispatch): Likewise.
(gen_get_num_teams): Likewise.
(gen_get_team_num): Likewise.
(gen_get_level): Likewise.
(gen_hsa_alloca): Likewise.
(gen_hsa_clrsb): Likewise.
(gen_hsa_ffs): Likewise.
(gen_hsa_popcount_to_dest): Likewise.
(gen_hsa_parity): Likewise.
(set_debug_value): Likewise.
(omp_simple_builtin::generate): Likewise.
(gen_hsa_insns_for_kernel_call): Likewise.
(gen_hsa_unaryop_for_builtin): Likewise.
(get_address_from_value): Likewise.
(gen_hsa_ternary_atomic_for_builtin): Likewise.
(gen_hsa_insn_for_internal_fn_call): Likewise.
(gen_hsa_insns_for_call): Likewise.
(gen_hsa_phi_from_gimple_phi): Likewise.
(init_hsa_num_threads): Likewise.
(gen_body_from_gimple): Likewise.
(gen_function_decl_parameters): Likewise.
(gen_function_def_parameters): Likewise.
(hsa_generate_function_declaration): Likewise.
(hsa_generate_internal_fn_decl): Likewise.
(convert_switch_statements): Likewise.
(expand_builtins): Likewise.
(generate_hsa): Likewise.
(pass_gen_hsail::execute): Likewise.
* hsa.c (hsa_deinit_compilation_unit_data): Likewise.
(hsa_get_declaration_name): Likewise.
(hsa_internal_fn::get_arity): Likewise.
(hsa_internal_fn::get_argument_type): Likewise.
* hsa.h (struct hsa_symbol): Likewise.
---
gcc/hsa-brig.c | 292 ++++++++++----------
gcc/hsa-dump.c | 11 +-
gcc/hsa-gen.c | 832 ++++++++++++++++++++++++++++++---------------------------
gcc/hsa.c | 23 +-
gcc/hsa.h | 20 +-
5 files changed, 615 insertions(+), 563 deletions(-)
diff --git a/gcc/hsa-brig.c b/gcc/hsa-brig.c
index 17e3398..7e71f6f 100644
--- a/gcc/hsa-brig.c
+++ b/gcc/hsa-brig.c
@@ -52,6 +52,10 @@ along with GCC; see the file COPYING3. If not see
#define BRIG_CHUNK_MAX_SIZE (64 * 1024)
+/* Required HSA section alignment. */
+
+#define HSA_SECTION_ALIGNMENT 16
+
/* Chunks of BRIG binary data. */
struct hsa_brig_data_chunk
@@ -72,7 +76,7 @@ public:
const char *section_name;
/* Size in bytes of all data stored in the section. */
unsigned total_size;
- /* The size of the header of the section including padding. */
+ /* The size of the header of the section including padding. */
unsigned header_byte_count;
/* The size of the header of the section without any padding. */
unsigned header_byte_delta;
@@ -80,7 +84,7 @@ public:
/* Buffers of binary data, each containing BRIG_CHUNK_MAX_SIZE bytes. */
vec <struct hsa_brig_data_chunk> chunks;
- /* More convenient access to the last chunk from the vector above. */
+ /* More convenient access to the last chunk from the vector above. */
struct hsa_brig_data_chunk *cur_chunk;
void allocate_new_chunk ();
@@ -110,8 +114,8 @@ static vec <hsa_insn_sbr *> *switch_instructions;
struct function_linkage_pair
{
- function_linkage_pair (tree decl, unsigned int off):
- function_decl (decl), offset (off) {}
+ function_linkage_pair (tree decl, unsigned int off)
+ : function_decl (decl), offset (off) {}
/* Declaration of called function. */
tree function_decl;
@@ -145,7 +149,7 @@ hsa_brig_section::init (const char *name)
certainly wasn't to have the first character of name and padding, which
are a part of sizeof (BrigSectionHeader), included in the first addend,
this is what the disassembler expects. */
- total_size = sizeof (BrigSectionHeader) + strlen(section_name);
+ total_size = sizeof (BrigSectionHeader) + strlen (section_name);
chunks.create (1);
allocate_new_chunk ();
header_byte_delta = total_size;
@@ -175,8 +179,8 @@ hsa_brig_section::output ()
section_header.byteCount = htole64 (total_size);
section_header.headerByteCount = htole32 (header_byte_count);
- section_header.nameLength = htole32 (strlen(section_name));
- assemble_string ((const char*) §ion_header, 16);
+ section_header.nameLength = htole32 (strlen (section_name));
+ assemble_string ((const char *) §ion_header, 16);
assemble_string (section_name, (section_header.nameLength));
memset (&padding, 0, sizeof (padding));
/* This is also a consequence of the wrong header size computation described
@@ -229,14 +233,13 @@ hsa_brig_section::round_size_up (int factor)
/* Return pointer to data by global OFFSET in the section. */
-void*
+void *
hsa_brig_section::get_ptr_by_offset (unsigned int offset)
{
gcc_assert (offset < total_size);
-
offset -= header_byte_delta;
- unsigned int i;
+ unsigned i;
for (i = 0; offset >= chunks[i].size; i++)
offset -= chunks[i].size;
@@ -272,8 +275,8 @@ brig_string_slot_hasher::hash (const value_type ds)
int i;
for (i = 0; i < ds->len; i++)
- r = r * 67 + (unsigned)ds->s[i] - 113;
- r = r * 67 + (unsigned)ds->prefix - 113;
+ r = r * 67 + (unsigned) ds->s[i] - 113;
+ r = r * 67 + (unsigned) ds->prefix - 113;
return r;
}
@@ -283,7 +286,8 @@ inline bool
brig_string_slot_hasher::equal (const value_type ds1, const compare_type ds2)
{
if (ds1->len == ds2->len)
- return ds1->prefix == ds2->prefix && memcmp (ds1->s, ds2->s, ds1->len) == 0;
+ return ds1->prefix == ds2->prefix
+ && memcmp (ds1->s, ds2->s, ds1->len) == 0;
return 0;
}
@@ -293,7 +297,7 @@ brig_string_slot_hasher::equal (const value_type ds1, const compare_type ds2)
inline void
brig_string_slot_hasher::remove (value_type ds)
{
- free (const_cast<char*> (ds->s));
+ free (const_cast<char *> (ds->s));
free (ds);
}
@@ -315,7 +319,6 @@ brig_emit_string (const char *str, char prefix = 0, bool sanitize = true)
brig_string_slot **slot;
char *str2;
- /* XXX Sanitize the names without all the strdup. */
str2 = xstrdup (str);
if (sanitize)
@@ -331,15 +334,16 @@ brig_emit_string (const char *str, char prefix = 0, bool sanitize = true)
brig_string_slot *new_slot = XCNEW (brig_string_slot);
/* In theory we should fill in BrigData but that would mean copying
- the string to a buffer for no reason, so we just emulate it. */
+ the string to a buffer for no reason, so we just emulate it. */
offset = brig_data.add (&hdr_len, sizeof (hdr_len));
if (prefix)
- brig_data.add (&prefix, 1);
+ brig_data.add (&prefix, 1);
brig_data.add (str2, slen);
brig_data.round_size_up (4);
- /* XXX could use the string we just copied into brig_string->cur_chunk */
+ /* TODO: could use the string we just copied into
+ brig_string->cur_chunk */
new_slot->s = str2;
new_slot->len = slen;
new_slot->prefix = prefix;
@@ -395,8 +399,8 @@ brig_init (void)
part = main_input_filename;
else
part++;
- asprintf (&modname, "&__hsa_module_%s", part);
- char* extension = strchr (modname, '.');
+ modname = concat ("&__hsa_module_", part, NULL);
+ char *extension = strchr (modname, '.');
if (extension)
*extension = '\0';
@@ -421,7 +425,7 @@ brig_init (void)
else
moddir.name = brig_emit_string ("__hsa_module_unnamed", '&');
moddir.base.kind = htole16 (BRIG_KIND_DIRECTIVE_MODULE);
- moddir.hsailMajor = htole32 (BRIG_VERSION_HSAIL_MAJOR) ;
+ moddir.hsailMajor = htole32 (BRIG_VERSION_HSAIL_MAJOR);
moddir.hsailMinor = htole32 (BRIG_VERSION_HSAIL_MINOR);
moddir.profile = hsa_full_profile_p () ? BRIG_PROFILE_FULL: BRIG_PROFILE_BASE;
if (hsa_machine_large_p ())
@@ -513,8 +517,8 @@ emit_directive_variable (struct hsa_symbol *symbol)
else
{
char buf[64];
- sprintf (buf, "__%s_%i", hsa_seg_name (symbol->m_segment),
- symbol->m_name_number);
+ snprintf (buf, 64, "__%s_%i", hsa_seg_name (symbol->m_segment),
+ symbol->m_name_number);
name_offset = brig_emit_string (buf, prefix);
}
@@ -527,8 +531,8 @@ emit_directive_variable (struct hsa_symbol *symbol)
dirvar.align = MAX (hsa_natural_alignment (dirvar.type),
(BrigAlignment8_t) BRIG_ALIGNMENT_4);
dirvar.linkage = symbol->m_linkage;
- dirvar.dim.lo = (uint32_t) symbol->m_dim;
- dirvar.dim.hi = (uint32_t) ((unsigned long long) symbol->m_dim >> 32);
+ dirvar.dim.lo = symbol->m_dim;
+ dirvar.dim.hi = symbol->m_dim >> 32;
/* Global variables are just declared and linked via HSA runtime. */
if (symbol->m_linkage != BRIG_ALLOCATION_PROGRAM)
@@ -565,7 +569,7 @@ emit_function_directives (hsa_function_representation *f, bool is_declaration)
}
name_offset = brig_emit_string (f->m_name, '&');
- inarg_off = brig_code.total_size + sizeof(fndir)
+ inarg_off = brig_code.total_size + sizeof (fndir)
+ (f->m_output_arg ? sizeof (struct BrigDirectiveVariable) : 0);
scoped_off = inarg_off
+ f->m_input_args.length () * sizeof (struct BrigDirectiveVariable);
@@ -603,15 +607,15 @@ emit_function_directives (hsa_function_representation *f, bool is_declaration)
else
{
/* Internal function. */
- hsa_internal_fn **slot = hsa_emitted_internal_decls->find_slot
- (f->m_internal_fn, INSERT);
+ hsa_internal_fn **slot
+ = hsa_emitted_internal_decls->find_slot (f->m_internal_fn, INSERT);
hsa_internal_fn *int_fn = new hsa_internal_fn (f->m_internal_fn);
int_fn->m_offset = brig_code.total_size;
*slot = int_fn;
}
brig_code.add (&fndir, sizeof (fndir));
- /* XXX terrible hack: we need to set instCount after we emit all
+ /* terrible hack: we need to set instCount after we emit all
insns, but we need to emit directive in order, and we emit directives
during insn emitting. So we need to emit the FUNCTION directive
early, then the insns, and then we need to set instCount, so remember
@@ -619,8 +623,8 @@ emit_function_directives (hsa_function_representation *f, bool is_declaration)
directly to after fndir here. */
ptr_to_fndir
= (BrigDirectiveExecutable *)(brig_code.cur_chunk->data
- + brig_code.cur_chunk->size
- - sizeof (fndir));
+ + brig_code.cur_chunk->size
+ - sizeof (fndir));
if (f->m_output_arg)
emit_directive_variable (f->m_output_arg);
@@ -651,14 +655,16 @@ static void
emit_bb_label_directive (hsa_bb *hbb)
{
struct BrigDirectiveLabel lbldir;
- char buf[32];
lbldir.base.byteCount = htole16 (sizeof (lbldir));
lbldir.base.kind = htole16 (BRIG_KIND_DIRECTIVE_LABEL);
- sprintf (buf, "BB_%u_%i", DECL_UID (current_function_decl), hbb->m_index);
+ char buf[32];
+ snprintf (buf, 32, "BB_%u_%i", DECL_UID (current_function_decl),
+ hbb->m_index);
lbldir.name = htole32 (brig_emit_string (buf, '@'));
- hbb->m_label_ref.m_directive_offset = brig_code.add (&lbldir, sizeof (lbldir));
+ hbb->m_label_ref.m_directive_offset = brig_code.add (&lbldir,
+ sizeof (lbldir));
brig_insn_count++;
}
@@ -788,7 +794,7 @@ emit_immediate_scalar_to_buffer (tree value, char *data, unsigned need_len)
tree type = TREE_TYPE (value);
gcc_checking_assert (TREE_CODE (type) != VECTOR_TYPE);
- unsigned data_len = tree_to_uhwi (TYPE_SIZE (type))/BITS_PER_UNIT;
+ unsigned data_len = tree_to_uhwi (TYPE_SIZE (type)) / BITS_PER_UNIT;
if (INTEGRAL_TYPE_P (type)
|| (POINTER_TYPE_P (type) && TREE_CODE (value) == INTEGER_CST))
switch (data_len)
@@ -803,7 +809,7 @@ emit_immediate_scalar_to_buffer (tree value, char *data, unsigned need_len)
bytes.b32 = (uint32_t) TREE_INT_CST_LOW (value);
break;
case 8:
- bytes.b64 = (uint64_t) int_cst_value (value);
+ bytes.b64 = (uint64_t) TREE_INT_CST_LOW (value);
break;
default:
gcc_unreachable ();
@@ -861,8 +867,8 @@ hsa_op_immed::emit_to_buffer (tree value)
for (i = 0; i < num; i++)
{
unsigned actual;
- actual = emit_immediate_scalar_to_buffer
- (VECTOR_CST_ELT (value, i), p, 0);
+ actual
+ = emit_immediate_scalar_to_buffer (VECTOR_CST_ELT (value, i), p, 0);
total_len -= actual;
p += actual;
}
@@ -876,14 +882,16 @@ hsa_op_immed::emit_to_buffer (tree value)
{
gcc_assert (total_len % 2 == 0);
unsigned actual;
- actual = emit_immediate_scalar_to_buffer
- (TREE_REALPART (value), p, total_len / 2);
+ actual
+ = emit_immediate_scalar_to_buffer (TREE_REALPART (value), p,
+ total_len / 2);
gcc_assert (actual == total_len / 2);
p += actual;
- actual = emit_immediate_scalar_to_buffer
- (TREE_IMAGPART (value), p, total_len / 2);
+ actual
+ = emit_immediate_scalar_to_buffer (TREE_IMAGPART (value), p,
+ total_len / 2);
gcc_assert (actual == total_len / 2);
}
else if (TREE_CODE (value) == CONSTRUCTOR)
@@ -891,8 +899,8 @@ hsa_op_immed::emit_to_buffer (tree value)
unsigned len = vec_safe_length (CONSTRUCTOR_ELTS (value));
for (unsigned i = 0; i < len; i++)
{
- unsigned actual = emit_immediate_scalar_to_buffer
- (CONSTRUCTOR_ELT (value, i)->value, p, 0);
+ tree v = CONSTRUCTOR_ELT (value, i)->value;
+ unsigned actual = emit_immediate_scalar_to_buffer (v, p, 0);
total_len -= actual;
p += actual;
}
@@ -918,7 +926,7 @@ emit_immediate_operand (hsa_op_immed *imm)
uint32_t byteCount = htole32 (imm->m_brig_repr_size);
out.type = htole16 (imm->m_type);
out.bytes = htole32 (brig_data.add (&byteCount, sizeof (byteCount)));
- brig_operand.add (&out, sizeof(out));
+ brig_operand.add (&out, sizeof (out));
brig_data.add (imm->m_brig_repr, imm->m_brig_repr_size);
brig_data.round_size_up (4);
}
@@ -970,8 +978,8 @@ emit_address_operand (hsa_op_address *addr)
if (sizeof (addr->m_imm_offset) == 8)
{
- out.offset.lo = htole32 ((uint32_t)addr->m_imm_offset);
- out.offset.hi = htole32 (((long long) addr->m_imm_offset) >> 32);
+ out.offset.lo = htole32 (addr->m_imm_offset);
+ out.offset.hi = htole32 (addr->m_imm_offset >> 32);
}
else
{
@@ -1097,7 +1105,6 @@ emit_internal_fn_decl (hsa_internal_fn *fn)
return e;
}
-
/* Enqueue all operands of INSN and return offset to BRIG data section
to list of operand offsets. */
@@ -1160,7 +1167,7 @@ emit_operands (hsa_op_base *op0, hsa_op_base *op1 = NULL,
}
/* Emit an HSA memory instruction and all necessary directives, schedule
- necessary operands for writing . */
+ necessary operands for writing. */
static void
emit_memory_insn (hsa_insn_mem *mem)
@@ -1185,7 +1192,7 @@ emit_memory_insn (hsa_insn_mem *mem)
repr.segment = addr->m_symbol->m_segment;
else
repr.segment = BRIG_SEGMENT_FLAT;
- repr.modifier.allBits = 0 ;
+ repr.modifier.allBits = 0;
repr.equivClass = mem->m_equiv_class;
repr.align = mem->m_align;
if (mem->m_opcode == BRIG_OPCODE_LD)
@@ -1225,7 +1232,7 @@ emit_signal_insn (hsa_insn_signal *mem)
}
/* Emit an HSA atomic memory instruction and all necessary directives, schedule
- necessary operands for writing . */
+ necessary operands for writing. */
static void
emit_atomic_insn (hsa_insn_atomic *mem)
@@ -1263,7 +1270,7 @@ emit_atomic_insn (hsa_insn_atomic *mem)
}
/* Emit an HSA LDA instruction and all necessary directives, schedule
- necessary operands for writing . */
+ necessary operands for writing. */
static void
emit_addr_insn (hsa_insn_basic *insn)
@@ -1289,7 +1296,7 @@ emit_addr_insn (hsa_insn_basic *insn)
}
/* Emit an HSA segment conversion instruction and all necessary directives,
- schedule necessary operands for writing . */
+ schedule necessary operands for writing. */
static void
emit_segment_insn (hsa_insn_seg *seg)
@@ -1311,7 +1318,7 @@ emit_segment_insn (hsa_insn_seg *seg)
}
/* Emit an HSA alloca instruction and all necessary directives,
- schedule necessary operands for writing . */
+ schedule necessary operands for writing. */
static void
emit_alloca_insn (hsa_insn_alloca *alloca)
@@ -1330,7 +1337,7 @@ emit_alloca_insn (hsa_insn_alloca *alloca)
repr.base.type = htole16 (alloca->m_type);
repr.base.operands = htole32 (emit_insn_operands (alloca));
repr.segment = BRIG_SEGMENT_PRIVATE;
- repr.modifier.allBits = 0 ;
+ repr.modifier.allBits = 0;
repr.equivClass = 0;
repr.align = alloca->m_align;
repr.width = BRIG_WIDTH_NONE;
@@ -1340,7 +1347,7 @@ emit_alloca_insn (hsa_insn_alloca *alloca)
}
/* Emit an HSA comparison instruction and all necessary directives,
- schedule necessary operands for writing . */
+ schedule necessary operands for writing. */
static void
emit_cmp_insn (hsa_insn_cmp *cmp)
@@ -1357,7 +1364,8 @@ emit_cmp_insn (hsa_insn_cmp *cmp)
if (is_a <hsa_op_reg *> (cmp->get_op (1)))
repr.sourceType = htole16 (as_a <hsa_op_reg *> (cmp->get_op (1))->m_type);
else
- repr.sourceType = htole16 (as_a <hsa_op_immed *> (cmp->get_op (1))->m_type);
+ repr.sourceType
+ = htole16 (as_a <hsa_op_immed *> (cmp->get_op (1))->m_type);
repr.modifier.allBits = 0;
repr.compare = cmp->m_compare;
repr.pack = 0;
@@ -1367,7 +1375,7 @@ emit_cmp_insn (hsa_insn_cmp *cmp)
}
/* Emit an HSA branching instruction and all necessary directives, schedule
- necessary operands for writing . */
+ necessary operands for writing. */
static void
emit_branch_insn (hsa_insn_br *br)
@@ -1395,8 +1403,9 @@ emit_branch_insn (hsa_insn_br *br)
}
gcc_assert (target);
- repr.base.operands = htole32
- (emit_operands (br->get_op (0), &hsa_bb_for_bb (target)->m_label_ref));
+ repr.base.operands
+ = htole32 (emit_operands (br->get_op (0),
+ &hsa_bb_for_bb (target)->m_label_ref));
memset (&repr.reserved, 0, sizeof (repr.reserved));
brig_code.add (&repr, sizeof (repr));
@@ -1440,8 +1449,8 @@ emit_switch_insn (hsa_insn_sbr *sbr)
/* For Conditional jumps the type is always B1. */
hsa_op_reg *index = as_a <hsa_op_reg *> (sbr->get_op (0));
repr.base.type = htole16 (index->m_type);
- repr.base.operands = htole32
- (emit_operands (sbr->get_op (0), sbr->m_label_code_list));
+ repr.base.operands
+ = htole32 (emit_operands (sbr->get_op (0), sbr->m_label_code_list));
memset (&repr.reserved, 0, sizeof (repr.reserved));
brig_code.add (&repr, sizeof (repr));
@@ -1502,19 +1511,21 @@ emit_call_insn (hsa_insn_call *call)
repr.base.opcode = htole16 (BRIG_OPCODE_CALL);
repr.base.type = htole16 (BRIG_TYPE_NONE);
- repr.base.operands = htole32
- (emit_operands (call->m_result_code_list, &call->m_func,
- call->m_args_code_list));
+ repr.base.operands
+ = htole32 (emit_operands (call->m_result_code_list, &call->m_func,
+ call->m_args_code_list));
/* Internal functions have not set m_called_function. */
if (call->m_called_function)
- function_call_linkage.safe_push
- (function_linkage_pair (call->m_called_function,
- call->m_func.m_brig_op_offset));
+ {
+ function_linkage_pair pair (call->m_called_function,
+ call->m_func.m_brig_op_offset);
+ function_call_linkage.safe_push (pair);
+ }
else
{
- hsa_internal_fn *slot = hsa_emitted_internal_decls->find
- (call->m_called_internal_fn);
+ hsa_internal_fn *slot
+ = hsa_emitted_internal_decls->find (call->m_called_internal_fn);
gcc_assert (slot);
gcc_assert (slot->m_offset > 0);
call->m_func.m_directive_offset = slot->m_offset;
@@ -1543,15 +1554,17 @@ emit_arg_block_insn (hsa_insn_arg_block *insn)
for (unsigned i = 0; i < insn->m_call_insn->m_input_args.length (); i++)
{
- insn->m_call_insn->m_args_code_list->m_offsets[i] = htole32
- (emit_directive_variable (insn->m_call_insn->m_input_args[i]));
+ insn->m_call_insn->m_args_code_list->m_offsets[i]
+ = htole32 (emit_directive_variable
+ (insn->m_call_insn->m_input_args[i]));
brig_insn_count++;
}
if (insn->m_call_insn->m_output_arg)
{
- insn->m_call_insn->m_result_code_list->m_offsets[0] = htole32
- (emit_directive_variable (insn->m_call_insn->m_output_arg));
+ insn->m_call_insn->m_result_code_list->m_offsets[0]
+ = htole32 (emit_directive_variable
+ (insn->m_call_insn->m_output_arg));
brig_insn_count++;
}
@@ -1652,8 +1665,8 @@ emit_packed_insn (hsa_insn_packed *insn)
for (unsigned i = 1; i < operand_count; i++)
{
gcc_checking_assert (insn->get_op (i));
- insn->m_operand_list->m_offsets[i - 1] = htole32
- (enqueue_op (insn->get_op (i)));
+ insn->m_operand_list->m_offsets[i - 1]
+ = htole32 (enqueue_op (insn->get_op (i)));
}
repr.base.operands = htole32 (emit_operands (insn->get_op (0),
@@ -1665,13 +1678,13 @@ emit_packed_insn (hsa_insn_packed *insn)
for (unsigned i = 0; i < operand_count - 1; i++)
{
gcc_checking_assert (insn->get_op (i));
- insn->m_operand_list->m_offsets[i] = htole32
- (enqueue_op (insn->get_op (i)));
+ insn->m_operand_list->m_offsets[i]
+ = htole32 (enqueue_op (insn->get_op (i)));
}
- repr.base.operands = htole32
- (emit_operands (insn->m_operand_list,
- insn->get_op (insn->operand_count () - 1)));
+ unsigned ops = emit_operands (insn->m_operand_list,
+ insn->get_op (insn->operand_count () - 1));
+ repr.base.operands = htole32 (ops);
}
@@ -1696,7 +1709,7 @@ emit_basic_insn (hsa_insn_basic *insn)
switch (insn->m_opcode)
{
/* And the bit-logical operations need bit types and whine about
- arithmetic types :-/ */
+ arithmetic types :-/ */
case BRIG_OPCODE_AND:
case BRIG_OPCODE_OR:
case BRIG_OPCODE_XOR:
@@ -1717,8 +1730,7 @@ emit_basic_insn (hsa_insn_basic *insn)
repr.round = BRIG_ROUND_FLOAT_NEAR_EVEN;
else
repr.round = 0;
- /* We assume that destination and sources agree in packing
- layout. */
+ /* We assume that destination and sources agree in packing layout. */
if (insn->num_used_ops () >= 2)
repr.pack = BRIG_PACK_PP;
else
@@ -1734,7 +1746,7 @@ emit_basic_insn (hsa_insn_basic *insn)
}
/* Emit an HSA instruction and all necessary directives, schedule necessary
- operands for writing . */
+ operands for writing. */
static void
emit_insn (hsa_insn_basic *insn)
@@ -1886,8 +1898,8 @@ hsa_brig_emit_function (void)
for (unsigned j = 0; j < sbr->m_jump_table.length (); j++)
{
hsa_bb *hbb = hsa_bb_for_bb (sbr->m_jump_table[j]);
- sbr->m_label_code_list->m_offsets[j] =
- hbb->m_label_ref.m_directive_offset;
+ sbr->m_label_code_list->m_offsets[j]
+ = hbb->m_label_ref.m_directive_offset;
}
}
@@ -1914,10 +1926,7 @@ hsa_brig_emit_omp_symbols (void)
emit_directive_variable (hsa_num_threads);
}
-/* Unit constructor and destructor statements. */
-
-static GTY(()) tree hsa_ctor_statements;
-static GTY(()) tree hsa_dtor_statements;
+static GTY(()) tree hsa_cdtor_statements[2];
/* Create and return __hsa_global_variables symbol that contains
all informations consumed by libgomp to link global variables
@@ -1961,8 +1970,8 @@ hsa_output_global_variables ()
hsa_sanitize_name (copy);
tree var_name = build_string (len, copy);
- TREE_TYPE (var_name) = build_array_type
- (char_type_node, build_index_type (size_int (len)));
+ TREE_TYPE (var_name)
+ = build_array_type (char_type_node, build_index_type (size_int (len)));
free (copy);
vec<constructor_elt, va_gc> *variable_info_vec = NULL;
@@ -2021,9 +2030,8 @@ hsa_output_kernels (tree *host_func_table, tree *kernels)
for (unsigned i = 0; i < map_count; ++i)
{
tree decl = hsa_get_decl_kernel_mapping_decl (i);
- CONSTRUCTOR_APPEND_ELT
- (host_functions_vec, NULL_TREE,
- build_fold_addr_expr (hsa_get_host_function (decl)));
+ tree host_fn = build_fold_addr_expr (hsa_get_host_function (decl));
+ CONSTRUCTOR_APPEND_ELT (host_functions_vec, NULL_TREE, host_fn);
}
tree host_functions_ctor = build_constructor (host_functions_array_type,
host_functions_vec);
@@ -2070,8 +2078,9 @@ hsa_output_kernels (tree *host_func_table, tree *kernels)
NULL_TREE);
int_num_of_kernels = build_int_cstu (uint32_type_node, map_count);
- tree kernel_info_vector_type = build_array_type
- (kernel_info_type, build_index_type (int_num_of_kernels));
+ tree kernel_info_vector_type
+ = build_array_type (kernel_info_type,
+ build_index_type (int_num_of_kernels));
TYPE_ARTIFICIAL (kernel_info_vector_type) = 1;
vec<constructor_elt, va_gc> *kernel_info_vector_vec = NULL;
@@ -2089,8 +2098,8 @@ hsa_output_kernels (tree *host_func_table, tree *kernels)
len++;
tree kern_name = build_string (len, copy);
- TREE_TYPE (kern_name) = build_array_type
- (char_type_node, build_index_type (size_int (len)));
+ TREE_TYPE (kern_name)
+ = build_array_type (char_type_node, build_index_type (size_int (len)));
free (copy);
unsigned omp_size = hsa_get_decl_kernel_mapping_omp_size (i);
@@ -2100,9 +2109,9 @@ hsa_output_kernels (tree *host_func_table, tree *kernels)
gridified_kernel_p);
unsigned count = 0;
- kernel_dependencies_vector_type = build_array_type
- (build_pointer_type (char_type_node),
- build_index_type (size_int (0)));
+ kernel_dependencies_vector_type
+ = build_array_type (build_pointer_type (char_type_node),
+ build_index_type (size_int (0)));
vec<constructor_elt, va_gc> *kernel_dependencies_vec = NULL;
if (hsa_decl_kernel_dependencies)
@@ -2114,9 +2123,9 @@ hsa_output_kernels (tree *host_func_table, tree *kernels)
vec <const char *> *dependencies = *slot;
count = dependencies->length ();
- kernel_dependencies_vector_type = build_array_type
- (build_pointer_type (char_type_node),
- build_index_type (size_int (count)));
+ kernel_dependencies_vector_type
+ = build_array_type (build_pointer_type (char_type_node),
+ build_index_type (size_int (count)));
TYPE_ARTIFICIAL (kernel_dependencies_vector_type) = 1;
for (unsigned j = 0; j < count; j++)
@@ -2124,8 +2133,9 @@ hsa_output_kernels (tree *host_func_table, tree *kernels)
const char *d = (*dependencies)[j];
len = strlen (d);
tree dependency_name = build_string (len, d);
- TREE_TYPE (dependency_name) = build_array_type
- (char_type_node, build_index_type (size_int (len)));
+ TREE_TYPE (dependency_name)
+ = build_array_type (char_type_node,
+ build_index_type (size_int (len)));
CONSTRUCTOR_APPEND_ELT
(kernel_dependencies_vec, NULL_TREE,
@@ -2163,8 +2173,9 @@ hsa_output_kernels (tree *host_func_table, tree *kernels)
DECL_IGNORED_P (dependencies_list) = 1;
DECL_EXTERNAL (dependencies_list) = 0;
TREE_CONSTANT (dependencies_list) = 1;
- DECL_INITIAL (dependencies_list) = build_constructor
- (kernel_dependencies_vector_type, kernel_dependencies_vec);
+ DECL_INITIAL (dependencies_list)
+ = build_constructor (kernel_dependencies_vector_type,
+ kernel_dependencies_vec);
varpool_node::finalize_decl (dependencies_list);
CONSTRUCTOR_APPEND_ELT (kernel_info_vec, NULL_TREE,
@@ -2275,18 +2286,16 @@ hsa_output_libgomp_mapping (tree brig_decl)
DECL_INITIAL (hsa_img_descriptor) = img_desc_ctor;
varpool_node::finalize_decl (hsa_img_descriptor);
- /* Construct the "host_table" libgomp expects. */
- tree libgomp_host_table_type = build_array_type (ptr_type_node,
- build_index_type
- (build_int_cst
- (integer_type_node, 4)));
+ /* Construct the "host_table" libgomp expects. */
+ tree index_type = build_index_type (build_int_cst (integer_type_node, 4));
+ tree libgomp_host_table_type = build_array_type (ptr_type_node, index_type);
TYPE_ARTIFICIAL (libgomp_host_table_type) = 1;
vec<constructor_elt, va_gc> *libgomp_host_table_vec = NULL;
tree host_func_table_addr = build_fold_addr_expr (host_func_table);
CONSTRUCTOR_APPEND_ELT (libgomp_host_table_vec, NULL_TREE,
host_func_table_addr);
- offset_int func_table_size = wi::to_offset (TYPE_SIZE_UNIT (ptr_type_node))
- * kernel_count;
+ offset_int func_table_size
+ = wi::to_offset (TYPE_SIZE_UNIT (ptr_type_node)) * kernel_count;
CONSTRUCTOR_APPEND_ELT (libgomp_host_table_vec, NULL_TREE,
fold_build2 (POINTER_PLUS_EXPR,
TREE_TYPE (host_func_table_addr),
@@ -2315,8 +2324,8 @@ hsa_output_libgomp_mapping (tree brig_decl)
/* Generate an initializer with a call to the registration routine. */
- tree offload_register = builtin_decl_explicit
- (BUILT_IN_GOMP_OFFLOAD_REGISTER);
+ tree offload_register
+ = builtin_decl_explicit (BUILT_IN_GOMP_OFFLOAD_REGISTER);
gcc_checking_assert (offload_register);
append_to_statement_list
@@ -2327,12 +2336,13 @@ hsa_output_libgomp_mapping (tree brig_decl)
build_fold_addr_expr (hsa_libgomp_host_table),
build_int_cst (integer_type_node, GOMP_DEVICE_HSA),
build_fold_addr_expr (hsa_img_descriptor)),
- &hsa_ctor_statements);
+ &hsa_cdtor_statements[0]);
- cgraph_build_static_cdtor ('I', hsa_ctor_statements, DEFAULT_INIT_PRIORITY);
+ cgraph_build_static_cdtor ('I', hsa_cdtor_statements[0],
+ DEFAULT_INIT_PRIORITY);
- tree offload_unregister = builtin_decl_explicit
- (BUILT_IN_GOMP_OFFLOAD_UNREGISTER);
+ tree offload_unregister
+ = builtin_decl_explicit (BUILT_IN_GOMP_OFFLOAD_UNREGISTER);
gcc_checking_assert (offload_unregister);
append_to_statement_list
@@ -2343,14 +2353,11 @@ hsa_output_libgomp_mapping (tree brig_decl)
build_fold_addr_expr (hsa_libgomp_host_table),
build_int_cst (integer_type_node, GOMP_DEVICE_HSA),
build_fold_addr_expr (hsa_img_descriptor)),
- &hsa_dtor_statements);
- cgraph_build_static_cdtor ('D', hsa_dtor_statements, DEFAULT_INIT_PRIORITY);
+ &hsa_cdtor_statements[1]);
+ cgraph_build_static_cdtor ('D', hsa_cdtor_statements[1],
+ DEFAULT_INIT_PRIORITY);
}
-/* Required HSA section alignment. */
-
-#define HSA_SECTION_ALIGNMENT 16
-
/* Emit the brig module we have compiled to a section in the final assembly and
also create a compile unit static constructor that will register the brig
module with libgomp. */
@@ -2369,8 +2376,8 @@ hsa_output_brig (void)
BrigCodeOffset32_t *func_offset = function_offsets->get (p.function_decl);
gcc_assert (*func_offset);
- BrigOperandCodeRef *code_ref = (BrigOperandCodeRef *)
- (brig_operand.get_ptr_by_offset (p.offset));
+ BrigOperandCodeRef *code_ref
+ = (BrigOperandCodeRef *) (brig_operand.get_ptr_by_offset (p.offset));
gcc_assert (code_ref->base.kind == BRIG_KIND_OPERAND_CODE_REF);
code_ref->ref = htole32 (*func_offset);
}
@@ -2380,8 +2387,9 @@ hsa_output_brig (void)
then change the linkage to program linkage. Doing so, we will emit
a valid BRIG image. */
if (hsa_failed_functions != NULL && emitted_declarations != NULL)
- for (hash_map <tree, BrigDirectiveExecutable *>::iterator it =
- emitted_declarations->begin (); it != emitted_declarations->end ();
+ for (hash_map <tree, BrigDirectiveExecutable *>::iterator it
+ = emitted_declarations->begin ();
+ it != emitted_declarations->end ();
++it)
{
if (hsa_failed_functions->contains ((*it).first))
@@ -2410,7 +2418,7 @@ hsa_output_brig (void)
BrigModuleHeader module_header;
memcpy (&module_header.identification, "HSA BRIG",
- sizeof(module_header.identification));
+ sizeof (module_header.identification));
module_header.brigMajor = htole32 (BRIG_VERSION_BRIG_MAJOR);
module_header.brigMinor = htole32 (BRIG_VERSION_BRIG_MINOR);
uint64_t section_index[3];
@@ -2423,17 +2431,21 @@ hsa_output_brig (void)
operand_padding = HSA_SECTION_ALIGNMENT
- brig_operand.total_size % HSA_SECTION_ALIGNMENT;
- uint64_t module_size = sizeof (module_header) + sizeof (section_index)
- + brig_data.total_size + data_padding
- + brig_code.total_size + code_padding
- + brig_operand.total_size + operand_padding;
+ uint64_t module_size = sizeof (module_header)
+ + sizeof (section_index)
+ + brig_data.total_size
+ + data_padding
+ + brig_code.total_size
+ + code_padding
+ + brig_operand.total_size
+ + operand_padding;
gcc_assert ((module_size % 16) == 0);
module_header.byteCount = htole64 (module_size);
memset (&module_header.hash, 0, sizeof (module_header.hash));
module_header.reserved = 0;
module_header.sectionCount = htole32 (3);
module_header.sectionIndex = htole64 (sizeof (module_header));
- assemble_string ((const char *) &module_header, sizeof(module_header));
+ assemble_string ((const char *) &module_header, sizeof (module_header));
uint64_t off = sizeof (module_header) + sizeof (section_index);
section_index[0] = htole64 (off);
off += brig_data.total_size + data_padding;
@@ -2443,7 +2455,7 @@ hsa_output_brig (void)
assemble_string ((const char *) §ion_index, sizeof (section_index));
char padding[HSA_SECTION_ALIGNMENT];
- memset (padding, 0, sizeof(padding));
+ memset (padding, 0, sizeof (padding));
brig_data.output ();
assemble_string (padding, data_padding);
diff --git a/gcc/hsa-dump.c b/gcc/hsa-dump.c
index aee525e..af79bcb 100644
--- a/gcc/hsa-dump.c
+++ b/gcc/hsa-dump.c
@@ -543,7 +543,7 @@ hsa_memsem_name (enum BrigMemoryOrder mo)
}
}
-/* Return textual name for memory scope. */
+/* Return textual name for memory scope. */
static const char *
hsa_memscope_name (enum BrigMemoryScope scope)
@@ -649,8 +649,9 @@ dump_hsa_reg (FILE *f, hsa_op_reg *reg, bool dump_type = false)
static void
dump_hsa_immed (FILE *f, hsa_op_immed *imm)
{
- bool unsigned_int_type = (BRIG_TYPE_U8 | BRIG_TYPE_U16 | BRIG_TYPE_U32
- | BRIG_TYPE_U64) & imm->m_type;
+ bool unsigned_int_type
+ = (BRIG_TYPE_U8 | BRIG_TYPE_U16 | BRIG_TYPE_U32 | BRIG_TYPE_U64)
+ & imm->m_type;
if (imm->m_tree_value)
print_generic_expr (f, imm->m_tree_value, 0);
@@ -662,7 +663,7 @@ dump_hsa_immed (FILE *f, hsa_op_immed *imm)
fprintf (f, HOST_WIDE_INT_PRINT_DEC, imm->m_int_value);
else
fprintf (f, HOST_WIDE_INT_PRINT_UNSIGNED,
- (unsigned HOST_WIDE_INT)imm->m_int_value);
+ (unsigned HOST_WIDE_INT) imm->m_int_value);
}
fprintf (f, " (%s)", hsa_type_name (imm->m_type));
@@ -967,7 +968,7 @@ dump_hsa_insn_1 (FILE *f, hsa_insn_basic *insn, int *indent)
fprintf (f, "(");
for (unsigned i = 0; i < call->m_input_args.length (); i++)
- {
+ {
fprintf (f, "%%__arg_%u", i);
if (i != call->m_input_args.length () - 1)
diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index 2628fbc..6d1cc98 100644
--- a/gcc/hsa-gen.c
+++ b/gcc/hsa-gen.c
@@ -126,7 +126,7 @@ struct hsa_queue
};
/* Alloc pools for allocating basic hsa structures such as operands,
- instructions and other basic entities.s */
+ instructions and other basic entities. */
static object_allocator<hsa_op_address> *hsa_allocp_operand_address;
static object_allocator<hsa_op_immed> *hsa_allocp_operand_immed;
static object_allocator<hsa_op_reg> *hsa_allocp_operand_reg;
@@ -158,11 +158,11 @@ static vec <hsa_insn_basic *> hsa_instructions;
static vec <hsa_op_base *> hsa_operands;
hsa_symbol::hsa_symbol ()
-: m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
- m_directive_offset (0), m_type (BRIG_TYPE_NONE),
- m_segment (BRIG_SEGMENT_NONE), m_linkage (BRIG_LINKAGE_NONE), m_dim (0),
- m_cst_value (NULL), m_global_scope_p (false), m_seen_error (false),
- m_allocation (BRIG_ALLOCATION_AUTOMATIC)
+ : m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
+ m_directive_offset (0), m_type (BRIG_TYPE_NONE),
+ m_segment (BRIG_SEGMENT_NONE), m_linkage (BRIG_LINKAGE_NONE), m_dim (0),
+ m_cst_value (NULL), m_global_scope_p (false), m_seen_error (false),
+ m_allocation (BRIG_ALLOCATION_AUTOMATIC)
{
}
@@ -170,18 +170,19 @@ hsa_symbol::hsa_symbol ()
hsa_symbol::hsa_symbol (BrigType16_t type, BrigSegment8_t segment,
BrigLinkage8_t linkage, bool global_scope_p,
BrigAllocation allocation)
-: m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
- m_directive_offset (0), m_type (type), m_segment (segment),
- m_linkage (linkage), m_dim (0), m_cst_value (NULL),
- m_global_scope_p (global_scope_p), m_seen_error (false),
- m_allocation (allocation)
+ : m_decl (NULL_TREE), m_name (NULL), m_name_number (0),
+ m_directive_offset (0), m_type (type), m_segment (segment),
+ m_linkage (linkage), m_dim (0), m_cst_value (NULL),
+ m_global_scope_p (global_scope_p), m_seen_error (false),
+ m_allocation (allocation)
{
}
unsigned HOST_WIDE_INT
hsa_symbol::total_byte_size ()
{
- unsigned HOST_WIDE_INT s = hsa_type_bit_size (~BRIG_TYPE_ARRAY_MASK & m_type);
+ unsigned HOST_WIDE_INT s
+ = hsa_type_bit_size (~BRIG_TYPE_ARRAY_MASK & m_type);
gcc_assert (s % BITS_PER_UNIT == 0);
s /= BITS_PER_UNIT;
@@ -213,15 +214,16 @@ hsa_symbol::fillup_for_decl (tree decl)
should be set to number of SSA names used in the function. */
hsa_function_representation::hsa_function_representation
- (tree fdecl, bool kernel_p, unsigned ssa_names_count): m_name (NULL),
- m_reg_count (0), m_input_args (vNULL),
- m_output_arg (NULL), m_spill_symbols (vNULL), m_global_symbols (vNULL),
- m_private_variables (vNULL), m_called_functions (vNULL),
- m_called_internal_fns (vNULL), m_hbb_count (0),
- m_in_ssa (true), m_kern_p (kernel_p), m_declaration_p (false), m_decl (fdecl),
- m_internal_fn (NULL), m_shadow_reg (NULL), m_kernel_dispatch_count (0),
- m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
- m_ssa_map ()
+ (tree fdecl, bool kernel_p, unsigned ssa_names_count)
+ : m_name (NULL),
+ m_reg_count (0), m_input_args (vNULL),
+ m_output_arg (NULL), m_spill_symbols (vNULL), m_global_symbols (vNULL),
+ m_private_variables (vNULL), m_called_functions (vNULL),
+ m_called_internal_fns (vNULL), m_hbb_count (0),
+ m_in_ssa (true), m_kern_p (kernel_p), m_declaration_p (false),
+ m_decl (fdecl), m_internal_fn (NULL), m_shadow_reg (NULL),
+ m_kernel_dispatch_count (0), m_maximum_omp_data_size (0),
+ m_seen_error (false), m_temp_symbol_count (0), m_ssa_map ()
{
int sym_init_len = (vec_safe_length (cfun->local_decls) / 2) + 1;;
m_local_symbols = new hash_table <hsa_noop_symbol_hasher> (sym_init_len);
@@ -230,16 +232,16 @@ hsa_function_representation::hsa_function_representation
/* Constructor of class representing HSA function information that
is derived for an internal function. */
-hsa_function_representation::hsa_function_representation (hsa_internal_fn *fn):
- m_reg_count (0), m_input_args (vNULL),
- m_output_arg (NULL), m_local_symbols (NULL),
- m_spill_symbols (vNULL), m_global_symbols (vNULL),
- m_private_variables (vNULL), m_called_functions (vNULL),
- m_called_internal_fns (vNULL), m_hbb_count (0),
- m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL),
- m_internal_fn (fn), m_shadow_reg (NULL), m_kernel_dispatch_count (0),
- m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
- m_ssa_map () {}
+hsa_function_representation::hsa_function_representation (hsa_internal_fn *fn)
+ : m_reg_count (0), m_input_args (vNULL),
+ m_output_arg (NULL), m_local_symbols (NULL),
+ m_spill_symbols (vNULL), m_global_symbols (vNULL),
+ m_private_variables (vNULL), m_called_functions (vNULL),
+ m_called_internal_fns (vNULL), m_hbb_count (0),
+ m_in_ssa (true), m_kern_p (false), m_declaration_p (true), m_decl (NULL),
+ m_internal_fn (fn), m_shadow_reg (NULL), m_kernel_dispatch_count (0),
+ m_maximum_omp_data_size (0), m_seen_error (false), m_temp_symbol_count (0),
+ m_ssa_map () {}
/* Destructor of class holding function/kernel-wide information and state. */
@@ -344,14 +346,14 @@ static hash_map <nofree_string_hash, omp_simple_builtin> *omp_simple_builtins
/* Warning messages for OMP builtins. */
-#define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP lock " \
- "routines"
+#define HSA_WARN_LOCK_ROUTINE "support for HSA does not implement OpenMP " \
+ "lock routines"
#define HSA_WARN_TIMING_ROUTINE "support for HSA does not implement OpenMP " \
"timing routines"
#define HSA_WARN_MEMORY_ROUTINE "OpenMP device memory library routines have " \
"undefined semantics within target regions, support for HSA ignores them"
-#define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP affinity " \
- "featerues"
+#define HSA_WARN_AFFINITY "Support for HSA does not implement OpenMP " \
+ "affinity feateres"
/* Initialize hash map with simple OMP builtins. */
@@ -361,14 +363,14 @@ hsa_init_simple_builtins ()
if (omp_simple_builtins != NULL)
return;
- omp_simple_builtins = new hash_map <nofree_string_hash, omp_simple_builtin>
- ();
+ omp_simple_builtins
+ = new hash_map <nofree_string_hash, omp_simple_builtin> ();
omp_simple_builtin omp_builtins[] =
{
- omp_simple_builtin
- ("omp_get_initial_device", NULL, false,
- new hsa_op_immed (GOMP_DEVICE_HOST, (BrigType16_t) BRIG_TYPE_S32)),
+ omp_simple_builtin ("omp_get_initial_device", NULL, false,
+ new hsa_op_immed (GOMP_DEVICE_HOST,
+ (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_is_initial_device", NULL, false,
new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_get_dynamic", NULL, false,
@@ -388,25 +390,26 @@ hsa_init_simple_builtins ()
omp_simple_builtin ("omp_target_alloc", HSA_WARN_MEMORY_ROUTINE, false,
new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_U64)),
omp_simple_builtin ("omp_target_free", HSA_WARN_MEMORY_ROUTINE, false),
- omp_simple_builtin
- ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE,
- false, new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
+ omp_simple_builtin ("omp_target_is_present", HSA_WARN_MEMORY_ROUTINE,
+ false,
+ new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_target_memcpy", HSA_WARN_MEMORY_ROUTINE, false,
new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
- omp_simple_builtin
- ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE,
- false, new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
- omp_simple_builtin
- ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE, false,
- new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
- omp_simple_builtin
- ("omp_target_disassociate_ptr", HSA_WARN_MEMORY_ROUTINE,
- false, new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
- omp_simple_builtin
- ("omp_set_max_active_levels",
- "Support for HSA only allows only one active level, call to "
- "omp_set_max_active_levels will be ignored in the generated HSAIL",
- false, NULL),
+ omp_simple_builtin ("omp_target_memcpy_rect", HSA_WARN_MEMORY_ROUTINE,
+ false,
+ new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
+ omp_simple_builtin ("omp_target_associate_ptr", HSA_WARN_MEMORY_ROUTINE,
+ false,
+ new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
+ omp_simple_builtin ("omp_target_disassociate_ptr",
+ HSA_WARN_MEMORY_ROUTINE,
+ false,
+ new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
+ omp_simple_builtin ("omp_set_max_active_levels",
+ "Support for HSA only allows only one active level, "
+ "call to omp_set_max_active_levels will be ignored "
+ "in the generated HSAIL",
+ false, NULL),
omp_simple_builtin ("omp_get_max_active_levels", NULL, false,
new hsa_op_immed (1, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_in_final", NULL, false,
@@ -485,7 +488,8 @@ hsa_init_data_for_cfun ()
hsa_allocp_inst_signal
= new object_allocator<hsa_insn_signal> ("HSA signal instructions");
hsa_allocp_inst_seg
- = new object_allocator<hsa_insn_seg> ("HSA segment conversion instructions");
+ = new object_allocator<hsa_insn_seg> ("HSA segment conversion "
+ "instructions");
hsa_allocp_inst_cmp
= new object_allocator<hsa_insn_cmp> ("HSA comparison instructions");
hsa_allocp_inst_br
@@ -721,7 +725,7 @@ hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
if (min32int)
{
/* Registers/immediate operands can only be 32bit or more except for
- f16. */
+ f16. */
if (res == BRIG_TYPE_U8 || res == BRIG_TYPE_U16)
res = BRIG_TYPE_U32;
else if (res == BRIG_TYPE_S8 || res == BRIG_TYPE_S16)
@@ -863,7 +867,8 @@ get_symbol_for_decl (tree decl)
dummy.m_decl = decl;
- bool is_in_global_vars = TREE_CODE (decl) == VAR_DECL && is_global_var (decl);
+ bool is_in_global_vars
+ = TREE_CODE (decl) == VAR_DECL && is_global_var (decl);
if (is_in_global_vars)
slot = hsa_global_variable_symbols->find_slot (&dummy, INSERT);
@@ -916,7 +921,8 @@ get_symbol_for_decl (tree decl)
tree
hsa_get_host_function (tree decl)
{
- hsa_function_summary *s = hsa_summaries->get (cgraph_node::get_create (decl));
+ hsa_function_summary *s
+ = hsa_summaries->get (cgraph_node::get_create (decl));
gcc_assert (s->m_kind != HSA_NONE);
gcc_assert (s->m_gpu_implementation_p);
@@ -981,8 +987,8 @@ hsa_get_string_cst_symbol (tree string_cst)
/* Constructor of the ancestor of all operands. K is BRIG kind that identified
what the operator is. */
-hsa_op_base::hsa_op_base (BrigKind16_t k): m_next (NULL), m_brig_op_offset (0),
- m_kind (k)
+hsa_op_base::hsa_op_base (BrigKind16_t k)
+ : m_next (NULL), m_brig_op_offset (0), m_kind (k)
{
hsa_operands.safe_push (this);
}
@@ -1046,8 +1052,8 @@ hsa_op_immed::hsa_op_immed (tree tree_val, bool min32int)
m_brig_repr_size = TREE_STRING_LENGTH (m_tree_value);
else if (TREE_CODE (m_tree_value) == CONSTRUCTOR)
{
- m_brig_repr_size = tree_to_uhwi
- (TYPE_SIZE_UNIT (TREE_TYPE (m_tree_value)));
+ m_brig_repr_size
+ = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (m_tree_value)));
/* Verify that all elements of a constructor are constants. */
for (unsigned i = 0;
@@ -1071,7 +1077,7 @@ hsa_op_immed::hsa_op_immed (tree tree_val, bool min32int)
hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value, BrigType16_t type)
: hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES, type),
- m_tree_value (NULL), m_brig_repr (NULL)
+ m_tree_value (NULL), m_brig_repr (NULL)
{
gcc_assert (hsa_type_integer_p (type));
m_int_value = integer_value;
@@ -1101,8 +1107,8 @@ hsa_op_immed::hsa_op_immed (HOST_WIDE_INT integer_value, BrigType16_t type)
memcpy (m_brig_repr, &bytes, m_brig_repr_size);
}
-hsa_op_immed::hsa_op_immed ():
- hsa_op_with_type (BRIG_KIND_NONE, BRIG_TYPE_NONE), m_brig_repr (NULL)
+hsa_op_immed::hsa_op_immed ()
+ : hsa_op_with_type (BRIG_KIND_NONE, BRIG_TYPE_NONE), m_brig_repr (NULL)
{
}
@@ -1134,8 +1140,8 @@ hsa_op_immed::set_type (BrigType16_t t)
hsa_op_reg::hsa_op_reg (BrigType16_t t)
: hsa_op_with_type (BRIG_KIND_OPERAND_REGISTER, t), m_gimple_ssa (NULL_TREE),
- m_def_insn (NULL), m_spill_sym (NULL), m_order (hsa_cfun->m_reg_count++),
- m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
+ m_def_insn (NULL), m_spill_sym (NULL), m_order (hsa_cfun->m_reg_count++),
+ m_lr_begin (0), m_lr_end (0), m_reg_class (0), m_hard_num (0)
{
}
@@ -1207,19 +1213,19 @@ hsa_op_reg::verify_ssa ()
hsa_op_address::hsa_op_address (hsa_symbol *sym, hsa_op_reg *r,
HOST_WIDE_INT offset)
: hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (r),
- m_imm_offset (offset)
+ m_imm_offset (offset)
{
}
hsa_op_address::hsa_op_address (hsa_symbol *sym, HOST_WIDE_INT offset)
: hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (sym), m_reg (NULL),
- m_imm_offset (offset)
+ m_imm_offset (offset)
{
}
hsa_op_address::hsa_op_address (hsa_op_reg *r, HOST_WIDE_INT offset)
: hsa_op_base (BRIG_KIND_OPERAND_ADDRESS), m_symbol (NULL), m_reg (r),
- m_imm_offset (offset)
+ m_imm_offset (offset)
{
}
@@ -1314,9 +1320,10 @@ hsa_op_reg::set_definition (hsa_insn_basic *insn)
operand vector will contain (and which will be cleared). OP is the opcode
of the instruction. This constructor does not set type. */
-hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc): m_prev (NULL),
- m_next (NULL), m_bb (NULL), m_opcode (opc), m_number (0),
- m_type (BRIG_TYPE_NONE), m_brig_offset (0)
+hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc)
+ : m_prev (NULL),
+ m_next (NULL), m_bb (NULL), m_opcode (opc), m_number (0),
+ m_type (BRIG_TYPE_NONE), m_brig_offset (0)
{
if (nops > 0)
m_operands.safe_grow_cleared (nops);
@@ -1381,9 +1388,9 @@ hsa_insn_basic::operand_count ()
hsa_insn_basic::hsa_insn_basic (unsigned nops, int opc, BrigType16_t t,
hsa_op_base *arg0, hsa_op_base *arg1,
- hsa_op_base *arg2, hsa_op_base *arg3):
- m_prev (NULL), m_next (NULL), m_bb (NULL), m_opcode (opc),m_number (0),
- m_type (t), m_brig_offset (0)
+ hsa_op_base *arg2, hsa_op_base *arg3)
+ : m_prev (NULL), m_next (NULL), m_bb (NULL), m_opcode (opc),m_number (0),
+ m_type (t), m_brig_offset (0)
{
if (nops > 0)
m_operands.safe_grow_cleared (nops);
@@ -1486,8 +1493,8 @@ hsa_insn_phi::operator new (size_t)
new instruction is automatically added to its uses list. */
hsa_insn_br::hsa_insn_br (hsa_op_reg *ctrl)
-: hsa_insn_basic (1, BRIG_OPCODE_CBR, BRIG_TYPE_B1, ctrl),
- m_width (BRIG_WIDTH_1)
+ : hsa_insn_basic (1, BRIG_OPCODE_CBR, BRIG_TYPE_B1, ctrl),
+ m_width (BRIG_WIDTH_1)
{
}
@@ -1503,9 +1510,9 @@ hsa_insn_br::operator new (size_t)
the index register. */
hsa_insn_sbr::hsa_insn_sbr (hsa_op_reg *index, unsigned jump_count)
-: hsa_insn_basic (1, BRIG_OPCODE_SBR, BRIG_TYPE_B1, index),
- m_width (BRIG_WIDTH_1), m_jump_table (vNULL), m_default_bb (NULL),
- m_label_code_list (new hsa_op_code_list (jump_count))
+ : hsa_insn_basic (1, BRIG_OPCODE_SBR, BRIG_TYPE_B1, index),
+ m_width (BRIG_WIDTH_1), m_jump_table (vNULL), m_default_bb (NULL),
+ m_label_code_list (new hsa_op_code_list (jump_count))
{
}
@@ -1558,7 +1565,7 @@ hsa_insn_cmp::operator new (size_t)
hsa_insn_mem::hsa_insn_mem (int opc, BrigType16_t t, hsa_op_base *arg0,
hsa_op_base *arg1)
: hsa_insn_basic (2, opc, t, arg0, arg1),
- m_align (hsa_natural_alignment (t)), m_equiv_class (0)
+ m_align (hsa_natural_alignment (t)), m_equiv_class (0)
{
gcc_checking_assert (opc == BRIG_OPCODE_LD || opc == BRIG_OPCODE_ST);
}
@@ -1572,7 +1579,7 @@ hsa_insn_mem::hsa_insn_mem (unsigned nops, int opc, BrigType16_t t,
hsa_op_base *arg0, hsa_op_base *arg1,
hsa_op_base *arg2, hsa_op_base *arg3)
: hsa_insn_basic (nops, opc, t, arg0, arg1, arg2, arg3),
- m_align (hsa_natural_alignment (t)), m_equiv_class (0)
+ m_align (hsa_natural_alignment (t)), m_equiv_class (0)
{
}
@@ -1584,7 +1591,7 @@ hsa_insn_mem::operator new (size_t)
return hsa_allocp_inst_mem->allocate_raw ();
}
-/* Constructor of class representing atomic instructions and signals. OPC is
+/* Constructor of class representing atomic instructions and signals. OPC is
the principal opcode, aop is the specific atomic operation opcode. T is the
type of the instruction. The instruction operands
are provided as ARG[0-3]. */
@@ -1596,8 +1603,8 @@ hsa_insn_atomic::hsa_insn_atomic (int nops, int opc,
hsa_op_base *arg1, hsa_op_base *arg2,
hsa_op_base *arg3)
: hsa_insn_mem (nops, opc, t, arg0, arg1, arg2, arg3), m_atomicop (aop),
- m_memoryorder (memorder),
- m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM)
+ m_memoryorder (memorder),
+ m_memoryscope (BRIG_MEMORY_SCOPE_SYSTEM)
{
gcc_checking_assert (opc == BRIG_OPCODE_ATOMICNORET ||
opc == BRIG_OPCODE_ATOMIC ||
@@ -1645,7 +1652,7 @@ hsa_insn_seg::hsa_insn_seg (int opc, BrigType16_t dest, BrigType16_t srct,
BrigSegment8_t seg, hsa_op_base *arg0,
hsa_op_base *arg1)
: hsa_insn_basic (2, opc, dest, arg0, arg1), m_src_type (srct),
- m_segment (seg)
+ m_segment (seg)
{
gcc_checking_assert (opc == BRIG_OPCODE_STOF || opc == BRIG_OPCODE_FTOS);
}
@@ -1663,14 +1670,14 @@ hsa_insn_seg::operator new (size_t)
hsa_insn_call::hsa_insn_call (tree callee)
: hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (callee),
- m_output_arg (NULL), m_args_code_list (NULL), m_result_code_list (NULL)
+ m_output_arg (NULL), m_args_code_list (NULL), m_result_code_list (NULL)
{
}
hsa_insn_call::hsa_insn_call (hsa_internal_fn *fn)
: hsa_insn_basic (0, BRIG_OPCODE_CALL), m_called_function (NULL),
- m_called_internal_fn (fn), m_output_arg (NULL), m_args_code_list (NULL),
- m_result_code_list (NULL)
+ m_called_internal_fn (fn), m_output_arg (NULL), m_args_code_list (NULL),
+ m_result_code_list (NULL)
{
}
@@ -1698,7 +1705,7 @@ hsa_insn_call::~hsa_insn_call ()
hsa_insn_arg_block::hsa_insn_arg_block (BrigKind brig_kind,
hsa_insn_call * call)
: hsa_insn_basic (0, HSA_OPCODE_ARG_BLOCK), m_kind (brig_kind),
- m_call_insn (call)
+ m_call_insn (call)
{
}
@@ -1757,7 +1764,7 @@ hsa_insn_srctype::hsa_insn_srctype (int nops, BrigOpcode opcode,
hsa_op_base *arg0, hsa_op_base *arg1,
hsa_op_base *arg2 = NULL)
: hsa_insn_basic (nops, opcode, destt, arg0, arg1, arg2),
- m_source_type (srct)
+ m_source_type (srct)
{}
/* New operator to allocate packed instruction from pool alloc. */
@@ -1807,7 +1814,7 @@ hsa_insn_alloca::operator new (size_t)
hsa_insn_alloca::hsa_insn_alloca (hsa_op_with_type *dest,
hsa_op_with_type *size, unsigned alignment)
: hsa_insn_basic (2, BRIG_OPCODE_ALLOCA, dest->m_type, dest, size),
- m_align (BRIG_ALIGNMENT_8)
+ m_align (BRIG_ALIGNMENT_8)
{
gcc_assert (dest->m_type == BRIG_TYPE_U32);
if (alignment)
@@ -1950,8 +1957,8 @@ process_mem_base (tree base, hsa_symbol **symbol, BrigType16_t *addrtype,
if (TREE_CODE (base) == SSA_NAME)
{
gcc_assert (!*reg);
- hsa_op_with_type *ssa = hsa_cfun->reg_for_gimple_ssa (base)->get_in_type
- (*addrtype, hbb);
+ hsa_op_with_type *ssa
+ = hsa_cfun->reg_for_gimple_ssa (base)->get_in_type (*addrtype, hbb);
*reg = dyn_cast <hsa_op_reg *> (ssa);
}
else if (TREE_CODE (base) == ADDR_EXPR)
@@ -2290,8 +2297,9 @@ gen_hsa_insns_for_bitfield (hsa_op_reg *dest, hsa_op_reg *value_reg,
hsa_op_reg *value_reg_2 = new hsa_op_reg (dest->m_type);
hsa_op_immed *c = new hsa_op_immed (left_shift, BRIG_TYPE_U32);
- hsa_insn_basic *lshift = new hsa_insn_basic
- (3, BRIG_OPCODE_SHL, value_reg_2->m_type, value_reg_2, value_reg, c);
+ hsa_insn_basic *lshift
+ = new hsa_insn_basic (3, BRIG_OPCODE_SHL, value_reg_2->m_type,
+ value_reg_2, value_reg, c);
hbb->append_insn (lshift);
@@ -2303,16 +2311,17 @@ gen_hsa_insns_for_bitfield (hsa_op_reg *dest, hsa_op_reg *value_reg,
hsa_op_reg *value_reg_2 = new hsa_op_reg (dest->m_type);
hsa_op_immed *c = new hsa_op_immed (right_shift, BRIG_TYPE_U32);
- hsa_insn_basic *rshift = new hsa_insn_basic
- (3, BRIG_OPCODE_SHR, value_reg_2->m_type, value_reg_2, value_reg, c);
+ hsa_insn_basic *rshift
+ = new hsa_insn_basic (3, BRIG_OPCODE_SHR, value_reg_2->m_type,
+ value_reg_2, value_reg, c);
hbb->append_insn (rshift);
value_reg = value_reg_2;
}
- hsa_insn_basic *assignment = new hsa_insn_basic
- (2, BRIG_OPCODE_MOV, dest->m_type, dest, value_reg);
+ hsa_insn_basic *assignment
+ = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, dest, value_reg);
hbb->append_insn (assignment);
}
@@ -2404,19 +2413,22 @@ gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb)
hsa_op_immed *real_part = new hsa_op_immed (TREE_REALPART (rhs));
hsa_op_immed *imag_part = new hsa_op_immed (TREE_IMAGPART (rhs));
- hsa_op_reg *real_part_reg = new hsa_op_reg
- (hsa_type_for_scalar_tree_type (TREE_TYPE (type), true));
- hsa_op_reg *imag_part_reg = new hsa_op_reg
- (hsa_type_for_scalar_tree_type (TREE_TYPE (type), true));
+ hsa_op_reg *real_part_reg
+ = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type),
+ true));
+ hsa_op_reg *imag_part_reg
+ = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (type),
+ true));
hsa_build_append_simple_mov (real_part_reg, real_part, hbb);
hsa_build_append_simple_mov (imag_part_reg, imag_part, hbb);
BrigType16_t src_type = hsa_bittype_for_type (real_part_reg->m_type);
- hsa_insn_packed *insn = new hsa_insn_packed
- (3, BRIG_OPCODE_COMBINE, dest->m_type, src_type, dest,
- real_part_reg, imag_part_reg);
+ hsa_insn_packed *insn
+ = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type,
+ src_type, dest, real_part_reg,
+ imag_part_reg);
hbb->append_insn (insn);
}
else
@@ -2429,22 +2441,23 @@ gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb)
{
tree pack_type = TREE_TYPE (TREE_OPERAND (rhs, 0));
- hsa_op_reg *packed_reg = new hsa_op_reg
- (hsa_type_for_scalar_tree_type (pack_type, true));
+ hsa_op_reg *packed_reg
+ = new hsa_op_reg (hsa_type_for_scalar_tree_type (pack_type, true));
tree complex_rhs = TREE_OPERAND (rhs, 0);
gen_hsa_insns_for_load (packed_reg, complex_rhs, TREE_TYPE (complex_rhs),
hbb);
- hsa_op_reg *real_reg = new hsa_op_reg
- (hsa_type_for_scalar_tree_type (type, true));
+ hsa_op_reg *real_reg
+ = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true));
- hsa_op_reg *imag_reg = new hsa_op_reg
- (hsa_type_for_scalar_tree_type (type, true));
+ hsa_op_reg *imag_reg
+ = new hsa_op_reg (hsa_type_for_scalar_tree_type (type, true));
BrigKind16_t brig_type = packed_reg->m_type;
- hsa_insn_packed *packed = new hsa_insn_packed
- (3, BRIG_OPCODE_EXPAND, hsa_bittype_for_type (real_reg->m_type),
+ hsa_insn_packed *packed
+ = new hsa_insn_packed (3, BRIG_OPCODE_EXPAND,
+ hsa_bittype_for_type (real_reg->m_type),
brig_type, real_reg, imag_reg, packed_reg);
hbb->append_insn (packed);
@@ -2502,9 +2515,10 @@ gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb)
}
}
else
- HSA_SORRY_ATV
- (EXPR_LOCATION (rhs),
- "support for HSA does not implement loading of expression %E", rhs);
+ HSA_SORRY_ATV (EXPR_LOCATION (rhs),
+ "support for HSA does not implement loading "
+ "of expression %E",
+ rhs);
}
/* Return number of bits necessary for representation of a bit field,
@@ -2556,8 +2570,9 @@ gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb)
if (bitpos || (bitsize && type_bitsize != bitsize))
{
unsigned HOST_WIDE_INT mask = 0;
- BrigType16_t mem_type = get_integer_type_by_bytes
- (type_bitsize / BITS_PER_UNIT, !TYPE_UNSIGNED (TREE_TYPE (lhs)));
+ BrigType16_t mem_type
+ = get_integer_type_by_bytes (type_bitsize / BITS_PER_UNIT,
+ !TYPE_UNSIGNED (TREE_TYPE (lhs)));
for (unsigned i = 0; i < type_bitsize; i++)
if (i < bitpos || i >= bitpos + bitsize)
@@ -2575,11 +2590,13 @@ gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb)
/* AND the loaded value with prepared mask. */
hsa_op_reg *cleared_reg = new hsa_op_reg (mem_type);
- hsa_op_immed *c = new hsa_op_immed
- (mask, get_integer_type_by_bytes (type_bitsize / BITS_PER_UNIT, false));
+ BrigType16_t t
+ = get_integer_type_by_bytes (type_bitsize / BITS_PER_UNIT, false);
+ hsa_op_immed *c = new hsa_op_immed (mask, t);
- hsa_insn_basic *clearing = new hsa_insn_basic
- (3, BRIG_OPCODE_AND, mem_type, cleared_reg, value_reg, c);
+ hsa_insn_basic *clearing
+ = new hsa_insn_basic (3, BRIG_OPCODE_AND, mem_type, cleared_reg,
+ value_reg, c);
hbb->append_insn (clearing);
/* Shift to left a value that is going to be stored. */
@@ -2594,8 +2611,9 @@ gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb)
hsa_op_reg *shifted_value_reg = new hsa_op_reg (mem_type);
c = new hsa_op_immed (bitpos, BRIG_TYPE_U32);
- hsa_insn_basic *basic = new hsa_insn_basic
- (3, BRIG_OPCODE_SHL, mem_type, shifted_value_reg, new_value_reg, c);
+ hsa_insn_basic *basic
+ = new hsa_insn_basic (3, BRIG_OPCODE_SHL, mem_type,
+ shifted_value_reg, new_value_reg, c);
hbb->append_insn (basic);
new_value_reg = shifted_value_reg;
@@ -2616,11 +2634,10 @@ gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb)
hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, mtype, src, addr);
mem->set_align (req_align);
- /* XXX The HSAIL verifier has another constraint: if the source
- is an immediate then it must match the destination type. If
- it's a register the low bits will be used for sub-word stores.
- We're always allocating new operands so we can modify the above
- in place. */
+ /* The HSAIL verifier has another constraint: if the source is an immediate
+ then it must match the destination type. If it's a register the low bits
+ will be used for sub-word stores. We're always allocating new operands so
+ we can modify the above in place. */
if (hsa_op_immed *imm = dyn_cast <hsa_op_immed *> (src))
{
if ((imm->m_type & BRIG_TYPE_PACK_MASK) == BRIG_TYPE_PACK_NONE)
@@ -2628,7 +2645,7 @@ gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb)
else
{
/* ...and all vector immediates apparently need to be vectors of
- unsigned bytes. */
+ unsigned bytes. */
unsigned bs = hsa_type_bit_size (imm->m_type);
gcc_assert (bs == hsa_type_bit_size (mem->m_type));
switch (bs)
@@ -3189,16 +3206,16 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
case UNEQ_EXPR:
case LTGT_EXPR:
{
- hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa
- (gimple_assign_lhs (assign));
+ hsa_op_reg *dest
+ = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
gen_hsa_cmp_insn_from_gimple (code, rhs1, rhs2, dest, hbb);
return;
}
case COND_EXPR:
{
- hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa
- (gimple_assign_lhs (assign));
+ hsa_op_reg *dest
+ = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
hsa_op_with_type *ctrl = NULL;
tree cond = rhs1;
@@ -3216,10 +3233,8 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
ctrl = r;
}
- hsa_op_with_type *rhs2_reg = hsa_reg_or_immed_for_gimple_op
- (rhs2, hbb);
- hsa_op_with_type *rhs3_reg = hsa_reg_or_immed_for_gimple_op
- (rhs3, hbb);
+ hsa_op_with_type *rhs2_reg = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
+ hsa_op_with_type *rhs3_reg = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
BrigType16_t btype = hsa_bittype_for_type (dest->m_type);
hsa_op_reg *tmp = new hsa_op_reg (btype);
@@ -3227,8 +3242,9 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
rhs2_reg->m_type = btype;
rhs3_reg->m_type = btype;
- hsa_insn_basic *insn = new hsa_insn_basic
- (4, BRIG_OPCODE_CMOV, tmp->m_type, tmp, ctrl, rhs2_reg, rhs3_reg);
+ hsa_insn_basic *insn
+ = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, tmp->m_type, tmp, ctrl,
+ rhs2_reg, rhs3_reg);
hbb->append_insn (insn);
@@ -3242,8 +3258,8 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
}
case COMPLEX_EXPR:
{
- hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa
- (gimple_assign_lhs (assign));
+ hsa_op_reg *dest
+ = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
hsa_op_with_type *rhs1_reg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
hsa_op_with_type *rhs2_reg = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
@@ -3254,9 +3270,9 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
rhs1_reg = rhs1_reg->get_in_type (src_type, hbb);
rhs2_reg = rhs2_reg->get_in_type (src_type, hbb);
- hsa_insn_packed *insn = new hsa_insn_packed
- (3, BRIG_OPCODE_COMBINE, dest->m_type, src_type, dest,
- rhs1_reg, rhs2_reg);
+ hsa_insn_packed *insn
+ = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type, src_type,
+ dest, rhs1_reg, rhs2_reg);
hbb->append_insn (insn);
return;
@@ -3378,8 +3394,8 @@ gen_hsa_insns_for_switch_stmt (gswitch *s, hsa_bb *hbb)
hsa_insn_sbr *sbr = new hsa_insn_sbr (sub_index, size + 1);
tree default_label = gimple_switch_default_label (s);
- basic_block default_label_bb = label_to_block_fn
- (func, CASE_LABEL (default_label));
+ basic_block default_label_bb = label_to_block_fn (func,
+ CASE_LABEL (default_label));
sbr->m_default_bb = default_label_bb;
@@ -3393,8 +3409,8 @@ gen_hsa_insns_for_switch_stmt (gswitch *s, hsa_bb *hbb)
tree label = gimple_switch_label (s, i);
basic_block bb = label_to_block_fn (func, CASE_LABEL (label));
- unsigned HOST_WIDE_INT sub_low = tree_to_uhwi
- (int_const_binop (MINUS_EXPR, CASE_LOW (label), lowest));
+ unsigned HOST_WIDE_INT sub_low
+ = tree_to_uhwi (int_const_binop (MINUS_EXPR, CASE_LOW (label), lowest));
unsigned HOST_WIDE_INT sub_high = sub_low;
tree high = CASE_HIGH (label);
@@ -3437,8 +3453,8 @@ get_format_argument_type (tree formal_arg_type, BrigType16_t actual_arg_type)
if (formal_arg_type == NULL)
return actual_arg_type;
- BrigType16_t decl_type = hsa_type_for_scalar_tree_type
- (formal_arg_type, false);
+ BrigType16_t decl_type
+ = hsa_type_for_scalar_tree_type (formal_arg_type, false);
return mem_type_for_type (decl_type);
}
@@ -3458,8 +3474,8 @@ gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb)
hsa_cfun->m_called_functions.safe_push (call_insn->m_called_function);
/* Argument block start. */
- hsa_insn_arg_block *arg_start = new hsa_insn_arg_block
- (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn);
+ hsa_insn_arg_block *arg_start
+ = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn);
hbb->append_insn (arg_start);
tree parm_type_chain = TYPE_ARG_TYPES (gimple_call_fntype (stmt));
@@ -3493,16 +3509,17 @@ gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb)
return;
}
- BrigType16_t formal_arg_type = get_format_argument_type
- (parm_decl_type, src->m_type);
+ BrigType16_t formal_arg_type
+ = get_format_argument_type (parm_decl_type, src->m_type);
if (hsa_seen_error ())
return;
if (src->m_type != formal_arg_type)
src = src->get_in_type (formal_arg_type, hbb);
- addr = gen_hsa_addr_for_arg
- (parm_decl_type != NULL_TREE ? parm_decl_type: TREE_TYPE (parm), i);
+ addr
+ = gen_hsa_addr_for_arg (parm_decl_type != NULL_TREE ?
+ parm_decl_type: TREE_TYPE (parm), i);
hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, formal_arg_type,
src, addr);
@@ -3542,8 +3559,9 @@ gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb)
}
else
{
- BrigType16_t mtype = mem_type_for_type
- (hsa_type_for_scalar_tree_type (lhs_type, false));
+ BrigType16_t mtype
+ = mem_type_for_type (hsa_type_for_scalar_tree_type (lhs_type,
+ false));
hsa_op_reg *dst = hsa_cfun->reg_for_gimple_ssa (result);
result_insn = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, dst, addr);
@@ -3568,8 +3586,8 @@ gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb)
}
/* Argument block end. */
- hsa_insn_arg_block *arg_end = new hsa_insn_arg_block
- (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn);
+ hsa_insn_arg_block *arg_end
+ = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn);
hbb->append_insn (arg_end);
}
@@ -3588,16 +3606,17 @@ gen_hsa_insns_for_call_of_internal_fn (gimple *stmt, hsa_bb *hbb)
tree rhs1 = gimple_call_arg (stmt, 0);
tree rhs1_type = TREE_TYPE (rhs1);
enum internal_fn fn = gimple_call_internal_fn (stmt);
- hsa_insn_call *call_insn = new hsa_insn_call
- (new hsa_internal_fn (fn, tree_to_uhwi (TYPE_SIZE (rhs1_type))));
+ hsa_internal_fn *ifn
+ = new hsa_internal_fn (fn, tree_to_uhwi (TYPE_SIZE (rhs1_type)));
+ hsa_insn_call *call_insn = new hsa_insn_call (ifn);
gcc_checking_assert (FLOAT_TYPE_P (rhs1_type));
if (!hsa_emitted_internal_decls->find (call_insn->m_called_internal_fn))
hsa_cfun->m_called_internal_fns.safe_push (call_insn->m_called_internal_fn);
- hsa_insn_arg_block *arg_start = new hsa_insn_arg_block
- (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn);
+ hsa_insn_arg_block *arg_start
+ = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_START, call_insn);
hbb->append_insn (arg_start);
unsigned num_args = gimple_call_num_args (stmt);
@@ -3626,8 +3645,8 @@ gen_hsa_insns_for_call_of_internal_fn (gimple *stmt, hsa_bb *hbb)
call_insn->m_result_code_list = new hsa_op_code_list (1);
/* Argument block end. */
- hsa_insn_arg_block *arg_end = new hsa_insn_arg_block
- (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn);
+ hsa_insn_arg_block *arg_end
+ = new hsa_insn_arg_block (BRIG_KIND_DIRECTIVE_ARG_BLOCK_END, call_insn);
hbb->append_insn (arg_end);
}
@@ -3651,8 +3670,9 @@ gen_hsa_insns_for_return (greturn *stmt, hsa_bb *hbb)
}
else
{
- BrigType16_t mtype = mem_type_for_type
- (hsa_type_for_scalar_tree_type (TREE_TYPE (retval), false));
+ BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (retval),
+ false);
+ BrigType16_t mtype = mem_type_for_type (t);
/* Store of return value. */
hsa_op_with_type *src = hsa_reg_or_immed_for_gimple_op (retval, hbb);
@@ -3707,8 +3727,8 @@ query_hsa_grid (hsa_op_reg *dest, BrigType16_t opcode, int dimension,
{
/* We're using just one-dimensional kernels, so hard-coded
dimension X. */
- hsa_op_immed *imm = new hsa_op_immed (dimension,
- (BrigKind16_t) BRIG_TYPE_U32);
+ hsa_op_immed *imm
+ = new hsa_op_immed (dimension, (BrigKind16_t) BRIG_TYPE_U32);
hsa_insn_basic *insn = new hsa_insn_basic (2, opcode, BRIG_TYPE_U32, NULL,
imm);
hbb->append_insn (insn);
@@ -3743,8 +3763,8 @@ gen_set_num_threads (tree value, hsa_bb *hbb)
src = src->get_in_type (hsa_num_threads->m_type, hbb);
hsa_op_address *addr = new hsa_op_address (hsa_num_threads);
- hsa_insn_basic *basic = new hsa_insn_mem
- (BRIG_OPCODE_ST, hsa_num_threads->m_type, src, addr);
+ hsa_insn_basic *basic
+ = new hsa_insn_mem (BRIG_OPCODE_ST, hsa_num_threads->m_type, src, addr);
hbb->append_insn (basic);
}
@@ -3841,35 +3861,35 @@ gen_num_threads_for_dispatch (hsa_bb *hbb)
hsa_op_immed *limit = new hsa_op_immed (HSA_DEFAULT_NUM_THREADS,
BRIG_TYPE_U32);
hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_B1);
- hbb->append_insn
- (new hsa_insn_cmp (BRIG_COMPARE_LT, r->m_type, r, threads, limit));
+ hsa_insn_cmp * cmp
+ = new hsa_insn_cmp (BRIG_COMPARE_LT, r->m_type, r, threads, limit);
+ hbb->append_insn (cmp);
BrigType16_t btype = hsa_bittype_for_type (threads->m_type);
hsa_op_reg *tmp = new hsa_op_reg (threads->m_type);
- hbb->append_insn
- (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp, r,
- threads, limit));
+ hbb->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp, r,
+ threads, limit));
/* Step 2) If the number is equal to zero,
return shadow->omp_num_threads. */
hsa_op_reg *shadow_reg_ptr = hsa_cfun->get_shadow_reg ();
hsa_op_reg *shadow_thread_count = new hsa_op_reg (BRIG_TYPE_U32);
- addr = new hsa_op_address
- (shadow_reg_ptr, get_hsa_kernel_dispatch_offset ("omp_num_threads"));
- hsa_insn_basic *basic = new hsa_insn_mem
- (BRIG_OPCODE_LD, shadow_thread_count->m_type, shadow_thread_count, addr);
+ addr
+ = new hsa_op_address (shadow_reg_ptr,
+ get_hsa_kernel_dispatch_offset ("omp_num_threads"));
+ hsa_insn_basic *basic
+ = new hsa_insn_mem (BRIG_OPCODE_LD, shadow_thread_count->m_type,
+ shadow_thread_count, addr);
hbb->append_insn (basic);
hsa_op_reg *tmp2 = new hsa_op_reg (threads->m_type);
r = new hsa_op_reg (BRIG_TYPE_B1);
- hbb->append_insn
- (new hsa_insn_cmp (BRIG_COMPARE_EQ, r->m_type, r, tmp,
- new hsa_op_immed (0, shadow_thread_count->m_type)));
- hbb->append_insn
- (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp2, r,
- shadow_thread_count, tmp));
+ hsa_op_immed *imm = new hsa_op_immed (0, shadow_thread_count->m_type);
+ hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_EQ, r->m_type, r, tmp, imm));
+ hbb->append_insn (new hsa_insn_basic (4, BRIG_OPCODE_CMOV, btype, tmp2, r,
+ shadow_thread_count, tmp));
hsa_op_base *dest = tmp2->get_in_type (BRIG_TYPE_U16, hbb);
@@ -3892,8 +3912,8 @@ gen_get_num_teams (gimple *stmt, hsa_bb *hbb)
hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
hsa_op_immed *one = new hsa_op_immed (1, dest->m_type);
- hsa_insn_basic *basic = new hsa_insn_basic
- (2, BRIG_OPCODE_MOV, dest->m_type, dest, one);
+ hsa_insn_basic *basic
+ = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, dest, one);
hbb->append_insn (basic);
}
@@ -3913,8 +3933,8 @@ gen_get_team_num (gimple *stmt, hsa_bb *hbb)
hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
hsa_op_immed *zero = new hsa_op_immed (0, dest->m_type);
- hsa_insn_basic *basic = new hsa_insn_basic
- (2, BRIG_OPCODE_MOV, dest->m_type, dest, zero);
+ hsa_insn_basic *basic
+ = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, dest, zero);
hbb->append_insn (basic);
}
@@ -3942,8 +3962,9 @@ gen_get_level (gimple *stmt, hsa_bb *hbb)
return;
}
- hsa_op_address *addr = new hsa_op_address
- (shadow_reg_ptr, get_hsa_kernel_dispatch_offset ("omp_level"));
+ hsa_op_address *addr
+ = new hsa_op_address (shadow_reg_ptr,
+ get_hsa_kernel_dispatch_offset ("omp_level"));
hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, NULL,
addr);
@@ -3990,10 +4011,10 @@ gen_hsa_alloca (gcall *call, hsa_bb *hbb)
tree alignment_tree = gimple_call_arg (call, 1);
if (TREE_CODE (alignment_tree) != INTEGER_CST)
{
- HSA_SORRY_ATV
- (gimple_location (call), "support for HSA does not implement "
- "__builtin_alloca_with_align with a non-constant "
- "alignment: %E", alignment_tree);
+ HSA_SORRY_ATV (gimple_location (call),
+ "support for HSA does not implement "
+ "__builtin_alloca_with_align with a non-constant "
+ "alignment: %E", alignment_tree);
}
bit_alignment = tree_to_uhwi (alignment_tree);
@@ -4004,14 +4025,15 @@ gen_hsa_alloca (gcall *call, hsa_bb *hbb)
->get_in_type (BRIG_TYPE_U32, hbb);
hsa_op_with_type *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
- hsa_op_reg *tmp = new hsa_op_reg
- (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE));
+ hsa_op_reg *tmp
+ = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE));
hsa_insn_alloca *a = new hsa_insn_alloca (tmp, size, bit_alignment);
hbb->append_insn (a);
- hsa_insn_seg *seg = new hsa_insn_seg
- (BRIG_OPCODE_STOF, hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT),
- tmp->m_type, BRIG_SEGMENT_PRIVATE, dest, tmp);
+ hsa_insn_seg *seg
+ = new hsa_insn_seg (BRIG_OPCODE_STOF,
+ hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT),
+ tmp->m_type, BRIG_SEGMENT_PRIVATE, dest, tmp);
hbb->append_insn (seg);
}
@@ -4043,8 +4065,9 @@ gen_hsa_clrsb (gcall *call, hsa_bb *hbb)
gen_hsa_binary_operation (BRIG_OPCODE_AND, and_reg, arg, c, hbb);
hsa_op_reg *most_sign = new hsa_op_reg (BRIG_TYPE_B1);
- hsa_insn_cmp *cmp = new hsa_insn_cmp
- (BRIG_COMPARE_EQ, most_sign->m_type, most_sign, and_reg, c);
+ hsa_insn_cmp *cmp
+ = new hsa_insn_cmp (BRIG_COMPARE_EQ, most_sign->m_type, most_sign,
+ and_reg, c);
hbb->append_insn (cmp);
/* If the most significant bit is one, negate the input. Otherwise
@@ -4059,8 +4082,9 @@ gen_hsa_clrsb (gcall *call, hsa_bb *hbb)
/* Assign the value that can be used for FIRSTBIT instruction according
to the most significant bit. */
hsa_op_reg *tmp = new hsa_op_reg (bittype);
- hsa_insn_basic *cmov = new hsa_insn_basic
- (4, BRIG_OPCODE_CMOV, bittype, tmp, most_sign, arg_neg, shifted_arg);
+ hsa_insn_basic *cmov
+ = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, bittype, tmp, most_sign,
+ arg_neg, shifted_arg);
hbb->append_insn (cmov);
hsa_op_reg *leading_bits = new hsa_op_reg (BRIG_TYPE_S32);
@@ -4070,16 +4094,14 @@ gen_hsa_clrsb (gcall *call, hsa_bb *hbb)
/* Set flag if the input value is equal to zero. */
hsa_op_reg *is_zero = new hsa_op_reg (BRIG_TYPE_B1);
- cmp = new hsa_insn_cmp
- (BRIG_COMPARE_EQ, is_zero->m_type, is_zero, arg,
- new hsa_op_immed (0, arg->m_type));
+ cmp = new hsa_insn_cmp (BRIG_COMPARE_EQ, is_zero->m_type, is_zero, arg,
+ new hsa_op_immed (0, arg->m_type));
hbb->append_insn (cmp);
/* Return the number of leading bits, or 31 if the input value is zero. */
- cmov = new hsa_insn_basic
- (4, BRIG_OPCODE_CMOV, BRIG_TYPE_B32, NULL, is_zero,
- new hsa_op_immed (31, BRIG_TYPE_U32),
- leading_bits->get_in_type (BRIG_TYPE_B32, hbb));
+ cmov = new hsa_insn_basic (4, BRIG_OPCODE_CMOV, BRIG_TYPE_B32, NULL, is_zero,
+ new hsa_op_immed (31, BRIG_TYPE_U32),
+ leading_bits->get_in_type (BRIG_TYPE_B32, hbb));
hbb->append_insn (cmov);
cmov->set_output_in_type (dest, 0, hbb);
}
@@ -4102,13 +4124,14 @@ gen_hsa_ffs (gcall *call, hsa_bb *hbb)
hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
hsa_op_reg *tmp = new hsa_op_reg (BRIG_TYPE_U32);
- hsa_insn_srctype *insn = new hsa_insn_srctype
- (2, BRIG_OPCODE_LASTBIT, tmp->m_type, arg->m_type, tmp, arg);
+ hsa_insn_srctype *insn = new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT,
+ tmp->m_type, arg->m_type,
+ tmp, arg);
hbb->append_insn (insn);
- hsa_insn_basic *addition = new hsa_insn_basic
- (3, BRIG_OPCODE_ADD, tmp->m_type, NULL, tmp,
- new hsa_op_immed (1, tmp->m_type));
+ hsa_insn_basic *addition
+ = new hsa_insn_basic (3, BRIG_OPCODE_ADD, tmp->m_type, NULL, tmp,
+ new hsa_op_immed (1, tmp->m_type));
hbb->append_insn (addition);
addition->set_output_in_type (dest, 0, hbb);
}
@@ -4124,8 +4147,9 @@ gen_hsa_popcount_to_dest (hsa_op_reg *dest, hsa_op_with_type *arg, hsa_bb *hbb)
if (!hsa_btype_p (arg->m_type))
arg = arg->get_in_type (hsa_bittype_for_type (arg->m_type), hbb);
- hsa_insn_srctype *popcount = new hsa_insn_srctype
- (2, BRIG_OPCODE_POPCOUNT, BRIG_TYPE_U32, arg->m_type, NULL, arg);
+ hsa_insn_srctype *popcount
+ = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT, BRIG_TYPE_U32,
+ arg->m_type, NULL, arg);
hbb->append_insn (popcount);
popcount->set_output_in_type (dest, 0, hbb);
}
@@ -4148,9 +4172,9 @@ gen_hsa_parity (gcall *call, hsa_bb *hbb)
hsa_op_reg *popcount = new hsa_op_reg (BRIG_TYPE_U32);
gen_hsa_popcount_to_dest (popcount, arg, hbb);
- hsa_insn_basic *insn = new hsa_insn_basic
- (3, BRIG_OPCODE_REM, popcount->m_type, NULL, popcount,
- new hsa_op_immed (2, popcount->m_type));
+ hsa_insn_basic *insn
+ = new hsa_insn_basic (3, BRIG_OPCODE_REM, popcount->m_type, NULL, popcount,
+ new hsa_op_immed (2, popcount->m_type));
hbb->append_insn (insn);
insn->set_output_in_type (dest, 0, hbb);
}
@@ -4182,8 +4206,9 @@ set_debug_value (hsa_bb *hbb, hsa_op_with_type *value)
if (shadow_reg_ptr == NULL)
return;
- hsa_op_address *addr = new hsa_op_address
- (shadow_reg_ptr, get_hsa_kernel_dispatch_offset ("debug"));
+ hsa_op_address *addr
+ = new hsa_op_address (shadow_reg_ptr,
+ get_hsa_kernel_dispatch_offset ("debug"));
hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, BRIG_TYPE_U64, value,
addr);
hbb->append_insn (mem);
@@ -4213,8 +4238,8 @@ omp_simple_builtin::generate (gimple *stmt, hsa_bb *hbb)
hbb->append_insn (new hsa_insn_comment (m_name));
hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
- hsa_build_append_simple_mov
- (dest, m_return_value->get_in_type (dest->m_type, hbb), hbb);
+ hsa_op_with_type *op = m_return_value->get_in_type (dest->m_type, hbb);
+ hsa_build_append_simple_mov (dest, op, hbb);
}
}
@@ -4326,8 +4351,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
/* Get my kernel dispatch argument. */
hbb->append_insn (new hsa_insn_comment ("get kernel dispatch structure"));
- addr = new hsa_op_address
- (shadow_reg_ptr, get_hsa_kernel_dispatch_offset ("children_dispatches"));
+ HOST_WIDE_INT o = get_hsa_kernel_dispatch_offset ("children_dispatches");
+ addr = new hsa_op_address (shadow_reg_ptr, o);
hsa_op_reg *shadow_reg_base_ptr = new hsa_op_reg (BRIG_TYPE_U64);
mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, shadow_reg_base_ptr,
@@ -4344,12 +4369,12 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
hbb->append_insn (mem);
/* Load an address of the command queue to a register. */
- hbb->append_insn (new hsa_insn_comment
- ("load base address of command queue"));
+ hbb->append_insn (new hsa_insn_comment ("load base address of command "
+ "queue"));
hsa_op_reg *queue_reg = new hsa_op_reg (BRIG_TYPE_U64);
- addr = new hsa_op_address (shadow_reg,
- get_hsa_kernel_dispatch_offset ("queue"));
+ o = get_hsa_kernel_dispatch_offset ("queue");
+ addr = new hsa_op_address (shadow_reg, o);
mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, queue_reg, addr);
@@ -4358,9 +4383,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
/* Load an address of prepared memory for a kernel arguments. */
hbb->append_insn (new hsa_insn_comment ("get a kernarg address"));
hsa_op_reg *kernarg_reg = new hsa_op_reg (BRIG_TYPE_U64);
-
- addr = new hsa_op_address
- (shadow_reg, get_hsa_kernel_dispatch_offset ("kernarg_address"));
+ o = get_hsa_kernel_dispatch_offset ("kernarg_address");
+ addr = new hsa_op_address (shadow_reg, o);
mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, kernarg_reg, addr);
hbb->append_insn (mem);
@@ -4368,9 +4392,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
/* Load an kernel object we want to call. */
hbb->append_insn (new hsa_insn_comment ("get a kernel object"));
hsa_op_reg *object_reg = new hsa_op_reg (BRIG_TYPE_U64);
-
- addr = new hsa_op_address
- (shadow_reg, get_hsa_kernel_dispatch_offset ("object"));
+ o = get_hsa_kernel_dispatch_offset ("object");
+ addr = new hsa_op_address (shadow_reg, o);
mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, object_reg, addr);
hbb->append_insn (mem);
@@ -4379,8 +4402,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
hbb->append_insn (new hsa_insn_comment ("get a signal by kernel call index"));
hsa_op_reg *signal_reg = new hsa_op_reg (BRIG_TYPE_U64);
- addr = new hsa_op_address
- (shadow_reg, get_hsa_kernel_dispatch_offset ("signal"));
+ o = get_hsa_kernel_dispatch_offset ("signal");
+ addr = new hsa_op_address (shadow_reg, o);
mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, signal_reg, addr);
hbb->append_insn (mem);
@@ -4399,22 +4422,22 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
/* Get private segment size. */
hsa_op_reg *private_seg_reg = new hsa_op_reg (BRIG_TYPE_U32);
- hbb->append_insn (new hsa_insn_comment
- ("get a kernel private segment size by kernel call index"));
+ hbb->append_insn (new hsa_insn_comment ("get a kernel private segment size "
+ "by kernel call index"));
- addr = new hsa_op_address
- (shadow_reg, get_hsa_kernel_dispatch_offset ("private_segment_size"));
+ o = get_hsa_kernel_dispatch_offset ("private_segment_size");
+ addr = new hsa_op_address (shadow_reg, o);
mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U32, private_seg_reg, addr);
hbb->append_insn (mem);
/* Get group segment size. */
hsa_op_reg *group_seg_reg = new hsa_op_reg (BRIG_TYPE_U32);
- hbb->append_insn (new hsa_insn_comment
- ("get a kernel group segment size by kernel call index"));
+ hbb->append_insn (new hsa_insn_comment ("get a kernel group segment size "
+ "by kernel call index"));
- addr = new hsa_op_address
- (shadow_reg, get_hsa_kernel_dispatch_offset ("group_segment_size"));
+ o = get_hsa_kernel_dispatch_offset ("group_segment_size");
+ addr = new hsa_op_address (shadow_reg, o);
mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U32, group_seg_reg, addr);
hbb->append_insn (mem);
@@ -4438,13 +4461,14 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
hsa_op_reg *queue_addr_reg = new hsa_op_reg (BRIG_TYPE_U64);
c = new hsa_op_immed (addr_offset, BRIG_TYPE_U64);
- hsa_insn_basic *insn = new hsa_insn_basic
- (3, BRIG_OPCODE_ADD, BRIG_TYPE_U64, queue_addr_reg, queue_reg, c);
+ hsa_insn_basic *insn
+ = new hsa_insn_basic (3, BRIG_OPCODE_ADD, BRIG_TYPE_U64, queue_addr_reg,
+ queue_reg, c);
hbb->append_insn (insn);
- hbb->append_insn (new hsa_insn_comment
- ("get base address of prepared packet"));
+ hbb->append_insn (new hsa_insn_comment ("get base address of prepared "
+ "packet"));
hsa_op_reg *queue_addr_value_reg = new hsa_op_reg (BRIG_TYPE_U64);
addr = new hsa_op_address (queue_addr_reg);
@@ -4454,16 +4478,15 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
c = new hsa_op_immed (sizeof (hsa_queue_packet), BRIG_TYPE_U64);
hsa_op_reg *queue_packet_offset_reg = new hsa_op_reg (BRIG_TYPE_U64);
- insn = new hsa_insn_basic
- (3, BRIG_OPCODE_MUL, BRIG_TYPE_U64, queue_packet_offset_reg,
- queue_index_reg, c);
+ insn = new hsa_insn_basic (3, BRIG_OPCODE_MUL, BRIG_TYPE_U64,
+ queue_packet_offset_reg, queue_index_reg, c);
hbb->append_insn (insn);
hsa_op_reg *queue_packet_reg = new hsa_op_reg (BRIG_TYPE_U64);
- insn = new hsa_insn_basic
- (3, BRIG_OPCODE_ADD, BRIG_TYPE_U64, queue_packet_reg, queue_addr_value_reg,
- queue_packet_offset_reg);
+ insn = new hsa_insn_basic (3, BRIG_OPCODE_ADD, BRIG_TYPE_U64,
+ queue_packet_reg, queue_addr_value_reg,
+ queue_packet_offset_reg);
hbb->append_insn (insn);
@@ -4478,8 +4501,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
addr);
hbb->append_insn (mem);
- hsa_op_with_type *packet_setup_u32 = packet_setup_reg->get_in_type
- (BRIG_TYPE_U32, hbb);
+ hsa_op_with_type *packet_setup_u32
+ = packet_setup_reg->get_in_type (BRIG_TYPE_U32, hbb);
hsa_op_reg *packet_setup_u32_2 = new hsa_op_reg (BRIG_TYPE_U32);
c = new hsa_op_immed (1, BRIG_TYPE_U32);
@@ -4488,8 +4511,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
hbb->append_insn (insn);
- hsa_op_with_type *packet_setup_reg_2 = packet_setup_u32_2->get_in_type
- (BRIG_TYPE_U16, hbb);
+ hsa_op_with_type *packet_setup_reg_2
+ = packet_setup_u32_2->get_in_type (BRIG_TYPE_U16, hbb);
addr = new hsa_op_address (queue_packet_reg,
offsetof (hsa_queue_packet, setup));
@@ -4501,8 +4524,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
emit passed grid_size. */
hsa_op_reg *threads_reg = gen_num_threads_for_dispatch (hbb);
- hbb->append_insn (new hsa_insn_comment
- ("set packet->grid_size_x = hsa_num_threads"));
+ hbb->append_insn (new hsa_insn_comment ("set packet->grid_size_x = "
+ "hsa_num_threads"));
addr = new hsa_op_address (queue_packet_reg,
offsetof (hsa_queue_packet, grid_size_x));
@@ -4511,20 +4534,20 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
hbb->append_insn (mem);
/* Write to shadow_reg->omp_num_threads = hsa_num_threads. */
- hbb->append_insn (new hsa_insn_comment
- ("set shadow_reg->omp_num_threads = hsa_num_threads"));
+ hbb->append_insn (new hsa_insn_comment ("set shadow_reg->omp_num_threads = "
+ "hsa_num_threads"));
- addr = new hsa_op_address
- (shadow_reg, get_hsa_kernel_dispatch_offset ("omp_num_threads"));
- hbb->append_insn
- (new hsa_insn_mem (BRIG_OPCODE_ST, threads_reg->m_type, threads_reg, addr));
+ o = get_hsa_kernel_dispatch_offset ("omp_num_threads");
+ addr = new hsa_op_address (shadow_reg, o);
+ hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, threads_reg->m_type,
+ threads_reg, addr));
/* Write to packet->workgroup_size_x. */
- hbb->append_insn (new hsa_insn_comment
- ("set packet->workgroup_size_x = hsa_num_threads"));
+ hbb->append_insn (new hsa_insn_comment ("set packet->workgroup_size_x = "
+ "hsa_num_threads"));
- addr = new hsa_op_address
- (queue_packet_reg, offsetof (hsa_queue_packet, workgroup_size_x));
+ addr = new hsa_op_address (queue_packet_reg,
+ offsetof (hsa_queue_packet, workgroup_size_x));
mem = new hsa_insn_mem (BRIG_OPCODE_ST, BRIG_TYPE_U16, threads_reg,
addr);
hbb->append_insn (mem);
@@ -4568,8 +4591,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
/* Write to packet->private_segment_size. */
hbb->append_insn (new hsa_insn_comment ("set packet->private_segment_size"));
- hsa_op_with_type *private_seg_reg_u16 = private_seg_reg->get_in_type
- (BRIG_TYPE_U16, hbb);
+ hsa_op_with_type *private_seg_reg_u16
+ = private_seg_reg->get_in_type (BRIG_TYPE_U16, hbb);
addr = new hsa_op_address (queue_packet_reg,
offsetof (hsa_queue_packet, private_segment_size));
@@ -4580,8 +4603,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
/* Write to packet->group_segment_size. */
hbb->append_insn (new hsa_insn_comment ("set packet->group_segment_size"));
- hsa_op_with_type *group_seg_reg_u16 = group_seg_reg->get_in_type
- (BRIG_TYPE_U16, hbb);
+ hsa_op_with_type *group_seg_reg_u16
+ = group_seg_reg->get_in_type (BRIG_TYPE_U16, hbb);
addr = new hsa_op_address (queue_packet_reg,
offsetof (hsa_queue_packet, group_segment_size));
@@ -4602,9 +4625,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
hsa_op_reg *omp_data_memory_reg = new hsa_op_reg (BRIG_TYPE_U64);
- addr = new hsa_op_address
- (shadow_reg, get_hsa_kernel_dispatch_offset ("omp_data_memory"));
-
+ o = get_hsa_kernel_dispatch_offset ("omp_data_memory");
+ addr = new hsa_op_address (shadow_reg, o);
mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, omp_data_memory_reg,
addr);
hbb->append_insn (mem);
@@ -4638,8 +4660,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
else
gcc_unreachable ();
- hbb->append_insn (new hsa_insn_comment
- ("write memory pointer to packet->kernarg_address"));
+ hbb->append_insn (new hsa_insn_comment ("write memory pointer to "
+ "packet->kernarg_address"));
addr = new hsa_op_address (queue_packet_reg,
offsetof (hsa_queue_packet, kernarg_address));
@@ -4647,8 +4669,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
hbb->append_insn (mem);
/* Write to packet->kernarg_address. */
- hbb->append_insn (new hsa_insn_comment
- ("write argument0 to *packet->kernarg_address"));
+ hbb->append_insn (new hsa_insn_comment ("write argument0 to "
+ "*packet->kernarg_address"));
addr = new hsa_op_address (kernarg_reg);
@@ -4657,8 +4679,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
hbb->append_insn (mem);
/* Pass shadow argument to another dispatched kernel module. */
- hbb->append_insn (new hsa_insn_comment
- ("write argument1 to *packet->kernarg_address"));
+ hbb->append_insn (new hsa_insn_comment ("write argument1 to "
+ "*packet->kernarg_address"));
addr = new hsa_op_address (kernarg_reg, 8);
mem = new hsa_insn_mem (BRIG_OPCODE_ST, BRIG_TYPE_U64, shadow_reg, addr);
@@ -4673,8 +4695,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
hbb->append_insn (mem);
/* Atomically write to packer->header. */
- hbb->append_insn
- (new hsa_insn_comment ("store atomically to packet->header"));
+ hbb->append_insn (new hsa_insn_comment ("store atomically to "
+ "packet->header"));
addr = new hsa_op_address (queue_packet_reg,
offsetof (hsa_queue_packet, header));
@@ -4682,9 +4704,10 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
/* Store 5122 << 16 + 1 to packet->header. */
c = new hsa_op_immed (70658, BRIG_TYPE_U32);
- hsa_insn_atomic *atomic = new hsa_insn_atomic
- (2, BRIG_OPCODE_ATOMICNORET, BRIG_ATOMIC_ST, BRIG_TYPE_B32,
- BRIG_MEMORY_ORDER_SC_RELEASE, addr, c);
+ hsa_insn_atomic *atomic
+ = new hsa_insn_atomic (2, BRIG_OPCODE_ATOMICNORET, BRIG_ATOMIC_ST,
+ BRIG_TYPE_B32, BRIG_MEMORY_ORDER_SC_RELEASE, addr,
+ c);
atomic->m_memoryscope = BRIG_MEMORY_SCOPE_SYSTEM;
hbb->append_insn (atomic);
@@ -4737,9 +4760,9 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
new_hbb->append_insn (signal);
hsa_op_reg *ctrl = new hsa_op_reg (BRIG_TYPE_B1);
- hsa_insn_cmp *cmp = new hsa_insn_cmp
- (BRIG_COMPARE_EQ, ctrl->m_type, ctrl, signal_result_reg,
- new hsa_op_immed (1, signal_result_reg->m_type));
+ hsa_op_immed *imm = new hsa_op_immed (1, signal_result_reg->m_type);
+ hsa_insn_cmp *cmp = new hsa_insn_cmp (BRIG_COMPARE_EQ, ctrl->m_type, ctrl,
+ signal_result_reg, imm);
new_hbb->append_insn (cmp);
new_hbb->append_insn (new hsa_insn_br (ctrl));
@@ -4747,8 +4770,8 @@ gen_hsa_insns_for_kernel_call (hsa_bb *hbb, gcall *call)
if (TREE_CODE (argument) == ADDR_EXPR)
{
/* Emit instructions that copy back OMP arguments to a caller kernel. */
- memcpy_hbb->append_insn
- (new hsa_insn_comment ("OMP arg memcpy back instructions"));
+ memcpy_hbb->append_insn (new hsa_insn_comment ("OMP arg memcpy back "
+ "instructions"));
hsa_op_address *src_addr = new hsa_op_address (omp_var_decl);
gen_hsa_memory_copy (memcpy_hbb, src_addr, dst_addr, omp_var_decl->m_dim);
@@ -4769,8 +4792,8 @@ gen_hsa_unaryop_for_builtin (BrigOpcode opcode, gimple *stmt, hsa_bb *hbb)
if (!lhs)
return;
hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
- hsa_op_with_type *op = hsa_reg_or_immed_for_gimple_op
- (gimple_call_arg (stmt, 0), hbb);
+ hsa_op_with_type *op
+ = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 0), hbb);
gen_hsa_unary_operation (opcode, dest, op, hbb);
}
@@ -4820,8 +4843,8 @@ get_address_from_value (tree val, hsa_bb *hbb)
case SSA_NAME:
{
BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
- hsa_op_base *reg = hsa_cfun->reg_for_gimple_ssa (val)->get_in_type
- (addrtype, hbb);
+ hsa_op_base *reg
+ = hsa_cfun->reg_for_gimple_ssa (val)->get_in_type (addrtype, hbb);
return new hsa_op_address (NULL, as_a <hsa_op_reg *> (reg), 0);
}
case ADDR_EXPR:
@@ -4897,7 +4920,8 @@ get_memory_order (unsigned memmodel, location_t location)
static void
gen_hsa_ternary_atomic_for_builtin (bool ret_orig,
- enum BrigAtomicOperation acode, gimple *stmt,
+ enum BrigAtomicOperation acode,
+ gimple *stmt,
hsa_bb *hbb)
{
tree lhs = gimple_call_lhs (stmt);
@@ -4909,16 +4933,15 @@ gen_hsa_ternary_atomic_for_builtin (bool ret_orig,
if (!tree_fits_uhwi_p (model))
{
- HSA_SORRY_ATV
- (gimple_location (stmt),
- "support for HSA does not implement memory model %E", model);
+ HSA_SORRY_ATV (gimple_location (stmt),
+ "support for HSA does not implement memory model %E",
+ model);
return;
}
unsigned HOST_WIDE_INT mmodel = tree_to_uhwi (model);
- BrigMemoryOrder memorder = get_memory_order
- (mmodel, gimple_location (stmt));
+ BrigMemoryOrder memorder = get_memory_order (mmodel, gimple_location (stmt));
/* Certain atomic insns must have Bx memory types. */
switch (acode)
@@ -5089,7 +5112,7 @@ gen_hsa_insn_for_internal_fn_call (gcall *stmt, hsa_bb *hbb)
gen_hsa_unaryop_or_call_for_builtin (BRIG_OPCODE_NSIN, stmt, hbb);
else
gen_hsa_unaryop_builtin_call (stmt, hbb);
- break;
+ break;
}
case IFN_CLRSB:
@@ -5291,9 +5314,10 @@ gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb)
tree model = gimple_call_arg (stmt, 1);
if (!tree_fits_uhwi_p (model))
{
- HSA_SORRY_ATV
- (gimple_location (stmt),
- "support for HSA does not implement memory model: %E", model);
+ HSA_SORRY_ATV (gimple_location (stmt),
+ "support for HSA does not implement "
+ "memory model: %E",
+ model);
return;
}
@@ -5304,17 +5328,18 @@ gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb)
if (memorder != BRIG_MEMORY_ORDER_RELAXED
&& memorder != BRIG_MEMORY_ORDER_SC_RELEASE)
{
- HSA_SORRY_ATV
- (gimple_location (stmt),
- "support for HSA does not implement memory model for "
- "ATOMIC_LD: %s", get_memory_order_name (mmodel));
+ HSA_SORRY_ATV (gimple_location (stmt),
+ "support for HSA does not implement "
+ "memory model for ATOMIC_LD: %s",
+ get_memory_order_name (mmodel));
return;
}
if (lhs)
{
- mtype = mem_type_for_type
- (hsa_type_for_scalar_tree_type (TREE_TYPE (lhs), false));
+ BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
+ false);
+ mtype = mem_type_for_type (t);
mtype = hsa_bittype_for_type (mtype);
dest = hsa_cfun->reg_for_gimple_ssa (lhs);
}
@@ -5435,15 +5460,15 @@ gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb)
case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_8:
case BUILT_IN_SYNC_VAL_COMPARE_AND_SWAP_16:
{
- /* XXX Ignore mem model for now. */
+ /* TODO: Use the appropriate memory model for now. */
tree type = TREE_TYPE (gimple_call_arg (stmt, 1));
- BrigType16_t atype = hsa_bittype_for_type
- (hsa_type_for_scalar_tree_type (type, false));
+ BrigType16_t atype
+ = hsa_bittype_for_type (hsa_type_for_scalar_tree_type (type, false));
- hsa_insn_atomic *atominsn = new hsa_insn_atomic
- (4, BRIG_OPCODE_ATOMIC, BRIG_ATOMIC_CAS, atype,
- BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE);
+ hsa_insn_atomic *atominsn
+ = new hsa_insn_atomic (4, BRIG_OPCODE_ATOMIC, BRIG_ATOMIC_CAS, atype,
+ BRIG_MEMORY_ORDER_SC_ACQUIRE_RELEASE);
hsa_op_address *addr;
addr = get_address_from_value (gimple_call_arg (stmt, 0), hbb);
@@ -5452,14 +5477,16 @@ gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb)
else
dest = new hsa_op_reg (atype);
- /* Should check what the memory scope is */
+ /* Should check what the memory scope is. */
atominsn->m_memoryscope = BRIG_MEMORY_SCOPE_WORKGROUP;
atominsn->set_op (0, dest);
atominsn->set_op (1, addr);
- atominsn->set_op
- (2, hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1), hbb));
- atominsn->set_op
- (3, hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 2), hbb));
+
+ hsa_op_with_type *op
+ = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1), hbb);
+ atominsn->set_op (2, op);
+ op = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 2), hbb);
+ atominsn->set_op (3, op);
hbb->append_insn (atominsn);
break;
@@ -5472,9 +5499,9 @@ gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb)
called = TREE_OPERAND (called, 0);
gcc_checking_assert (TREE_CODE (called) == FUNCTION_DECL);
- hsa_add_kernel_dependency
- (hsa_cfun->m_decl,
- hsa_brig_function_name (hsa_get_declaration_name (called)));
+ const char *name
+ = hsa_brig_function_name (hsa_get_declaration_name (called));
+ hsa_add_kernel_dependency (hsa_cfun->m_decl, name);
gen_hsa_insns_for_kernel_call (hbb, as_a <gcall *> (stmt));
break;
@@ -5543,9 +5570,10 @@ gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb)
if (builtin == BUILT_IN_MEMPCPY)
{
tmp = new hsa_op_reg (dst_reg->m_type);
- hsa_insn_basic *add = new hsa_insn_basic
- (3, BRIG_OPCODE_ADD, tmp->m_type,
- tmp, dst_reg, new hsa_op_immed (n, dst_reg->m_type));
+ hsa_insn_basic *add
+ = new hsa_insn_basic (3, BRIG_OPCODE_ADD, tmp->m_type,
+ tmp, dst_reg,
+ new hsa_op_immed (n, dst_reg->m_type));
hbb->append_insn (add);
}
else
@@ -5585,8 +5613,8 @@ gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb)
hsa_op_address *dst_addr;
dst_addr = get_address_from_value (dst, hbb);
- unsigned HOST_WIDE_INT constant = tree_to_uhwi
- (fold_convert (unsigned_char_type_node, c));
+ unsigned HOST_WIDE_INT constant
+ = tree_to_uhwi (fold_convert (unsigned_char_type_node, c));
gen_hsa_memory_set (hbb, dst_addr, constant, n);
@@ -5704,8 +5732,8 @@ gen_hsa_phi_from_gimple_phi (gimple *phi_stmt, hsa_bb *hbb)
hsa_insn_phi *hphi;
unsigned count = gimple_phi_num_args (phi_stmt);
- hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa
- (gimple_phi_result (phi_stmt));
+ hsa_op_reg *dest
+ = hsa_cfun->reg_for_gimple_ssa (gimple_phi_result (phi_stmt));
hphi = new hsa_insn_phi (count, dest);
hphi->m_bb = hbb->m_bb;
@@ -5742,8 +5770,9 @@ gen_hsa_phi_from_gimple_phi (gimple *phi_stmt, hsa_bb *hbb)
hbb_src);
hsa_op_reg *dest = new hsa_op_reg (BRIG_TYPE_U64);
- hsa_insn_basic *insn = new hsa_insn_basic
- (2, BRIG_OPCODE_LDA, BRIG_TYPE_U64, dest, addr);
+ hsa_insn_basic *insn
+ = new hsa_insn_basic (2, BRIG_OPCODE_LDA, BRIG_TYPE_U64,
+ dest, addr);
hbb_src->append_insn (insn);
hphi->set_op (i, dest);
@@ -5772,10 +5801,10 @@ gen_hsa_phi_from_gimple_phi (gimple *phi_stmt, hsa_bb *hbb)
index of this BB (so that the constructor does not attempt to use
hsa_cfun during its construction). */
-hsa_bb::hsa_bb (basic_block cfg_bb, int idx): m_bb (cfg_bb),
- m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL),
- m_last_phi (NULL), m_index (idx), m_liveout (BITMAP_ALLOC (NULL)),
- m_livein (BITMAP_ALLOC (NULL))
+hsa_bb::hsa_bb (basic_block cfg_bb, int idx)
+ : m_bb (cfg_bb), m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL),
+ m_last_phi (NULL), m_index (idx), m_liveout (BITMAP_ALLOC (NULL)),
+ m_livein (BITMAP_ALLOC (NULL))
{
gcc_assert (!cfg_bb->aux);
cfg_bb->aux = this;
@@ -5784,10 +5813,10 @@ hsa_bb::hsa_bb (basic_block cfg_bb, int idx): m_bb (cfg_bb),
/* Constructor of class containing HSA-specific information about a basic
block. CFG_BB is the CFG BB this HSA BB is associated with. */
-hsa_bb::hsa_bb (basic_block cfg_bb): m_bb (cfg_bb),
- m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL),
- m_last_phi (NULL), m_index (hsa_cfun->m_hbb_count++),
- m_liveout (BITMAP_ALLOC (NULL)), m_livein (BITMAP_ALLOC (NULL))
+hsa_bb::hsa_bb (basic_block cfg_bb)
+ : m_bb (cfg_bb), m_first_insn (NULL), m_last_insn (NULL), m_first_phi (NULL),
+ m_last_phi (NULL), m_index (hsa_cfun->m_hbb_count++),
+ m_liveout (BITMAP_ALLOC (NULL)), m_livein (BITMAP_ALLOC (NULL))
{
gcc_assert (!cfg_bb->aux);
cfg_bb->aux = this;
@@ -5836,10 +5865,10 @@ init_hsa_num_threads (void)
hsa_bb *prologue = hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
/* Save the default value to private variable hsa_num_threads. */
- hsa_insn_basic *basic = new hsa_insn_mem
- (BRIG_OPCODE_ST, hsa_num_threads->m_type,
- new hsa_op_immed (0, hsa_num_threads->m_type),
- new hsa_op_address (hsa_num_threads));
+ hsa_insn_basic *basic
+ = new hsa_insn_mem (BRIG_OPCODE_ST, hsa_num_threads->m_type,
+ new hsa_op_immed (0, hsa_num_threads->m_type),
+ new hsa_op_address (hsa_num_threads));
prologue->append_insn (basic);
}
@@ -5862,9 +5891,9 @@ gen_body_from_gimple ()
to the same basic block. */
if (e->flags & EDGE_EH)
{
- HSA_SORRY_AT
- (UNKNOWN_LOCATION,
- "support for HSA does not implement exception handling");
+ HSA_SORRY_AT (UNKNOWN_LOCATION,
+ "support for HSA does not implement exception "
+ "handling");
return;
}
}
@@ -5935,8 +5964,8 @@ gen_function_decl_parameters (hsa_function_representation *f,
{
f->m_output_arg = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
BRIG_LINKAGE_NONE);
- f->m_output_arg->m_type = hsa_type_for_tree_type
- (result_type, &f->m_output_arg->m_dim);
+ f->m_output_arg->m_type
+ = hsa_type_for_tree_type (result_type, &f->m_output_arg->m_dim);
f->m_output_arg->m_name = "res";
}
}
@@ -5957,10 +5986,10 @@ gen_function_def_parameters ()
{
struct hsa_symbol **slot;
- hsa_symbol *arg = new hsa_symbol
- (BRIG_TYPE_NONE,
- hsa_cfun->m_kern_p ? BRIG_SEGMENT_KERNARG : BRIG_SEGMENT_ARG,
- BRIG_LINKAGE_FUNCTION);
+ hsa_symbol *arg
+ = new hsa_symbol (BRIG_TYPE_NONE, hsa_cfun->m_kern_p
+ ? BRIG_SEGMENT_KERNARG : BRIG_SEGMENT_ARG,
+ BRIG_LINKAGE_FUNCTION);
arg->fillup_for_decl (parm);
hsa_cfun->m_input_args.safe_push (arg);
@@ -5997,8 +6026,9 @@ gen_function_def_parameters ()
tree ddef = ssa_default_def (cfun, parm);
if (ddef && !has_zero_uses (ddef))
{
- BrigType16_t mtype = mem_type_for_type
- (hsa_type_for_scalar_tree_type (TREE_TYPE (ddef), false));
+ BrigType16_t t = hsa_type_for_scalar_tree_type (TREE_TYPE (ddef),
+ false);
+ BrigType16_t mtype = mem_type_for_type (t);
hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (ddef);
hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mtype,
dest, parm_addr);
@@ -6033,8 +6063,8 @@ gen_function_def_parameters ()
hsa_function_representation *
hsa_generate_function_declaration (tree decl)
{
- hsa_function_representation *fun = new hsa_function_representation
- (decl, false, 0);
+ hsa_function_representation *fun
+ = new hsa_function_representation (decl, false, 0);
fun->m_declaration_p = true;
fun->m_name = get_brig_function_name (decl);
@@ -6056,14 +6086,15 @@ hsa_generate_internal_fn_decl (hsa_internal_fn *fn)
for (unsigned i = 0; i < fn->get_arity (); i++)
{
- hsa_symbol *arg = new hsa_symbol
- (fn->get_argument_type (i), BRIG_SEGMENT_ARG, BRIG_LINKAGE_NONE);
+ hsa_symbol *arg
+ = new hsa_symbol (fn->get_argument_type (i), BRIG_SEGMENT_ARG,
+ BRIG_LINKAGE_NONE);
arg->m_name_number = i;
fun->m_input_args.safe_push (arg);
}
- fun->m_output_arg = new hsa_symbol
- (fn->get_argument_type (-1), BRIG_SEGMENT_ARG, BRIG_LINKAGE_NONE);
+ fun->m_output_arg = new hsa_symbol (fn->get_argument_type (-1),
+ BRIG_SEGMENT_ARG, BRIG_LINKAGE_NONE);
fun->m_output_arg->m_name = "res";
return fun;
@@ -6194,8 +6225,8 @@ convert_switch_statements ()
tree index = gimple_switch_index (s);
tree index_type = TREE_TYPE (index);
tree default_label = gimple_switch_default_label (s);
- basic_block default_label_bb = label_to_block_fn
- (func, CASE_LABEL (default_label));
+ basic_block default_label_bb
+ = label_to_block_fn (func, CASE_LABEL (default_label));
basic_block cur_bb = bb;
auto_vec <edge> new_edges;
@@ -6226,8 +6257,9 @@ convert_switch_statements ()
if (gimple_phi_arg_edge (phi, j) == e)
{
tree imm = gimple_phi_arg_def (phi, j);
- phi_todo_list.safe_push
- (new phi_definition (phi_index, i, imm));
+ phi_definition *p = new phi_definition (phi_index, i,
+ imm);
+ phi_todo_list.safe_push (p);
break;
}
}
@@ -6279,8 +6311,8 @@ convert_switch_statements ()
gsi_insert_before (&cond_gsi, assign2, GSI_SAME_STMT);
gsi_insert_before (&cond_gsi, assign3, GSI_SAME_STMT);
- c = gimple_build_cond (NE_EXPR, tmp3, constant_boolean_node
- (false, boolean_type_node), NULL, NULL);
+ tree b = constant_boolean_node (false, boolean_type_node);
+ c = gimple_build_cond (NE_EXPR, tmp3, b, NULL, NULL);
}
else
c = gimple_build_cond (EQ_EXPR, index, low, NULL, NULL);
@@ -6289,15 +6321,15 @@ convert_switch_statements ()
gsi_insert_before (&cond_gsi, c, GSI_SAME_STMT);
- basic_block label_bb = label_to_block_fn
- (func, CASE_LABEL (label));
+ basic_block label_bb
+ = label_to_block_fn (func, CASE_LABEL (label));
edge new_edge = make_edge (cur_bb, label_bb, EDGE_TRUE_VALUE);
int prob_sum = sum_slice <int> (edge_probabilities, i, labels) +
edge_probabilities[0];
if (prob_sum)
- new_edge->probability = RDIV
- (REG_BR_PROB_BASE * edge_probabilities[i], prob_sum);
+ new_edge->probability
+ = RDIV (REG_BR_PROB_BASE * edge_probabilities[i], prob_sum);
new_edge->count = edge_counts[i];
new_edges.safe_push (new_edge);
@@ -6314,8 +6346,8 @@ convert_switch_statements ()
}
edge next_edge = make_edge (cur_bb, next_bb, EDGE_FALSE_VALUE);
- next_edge->probability = inverse_probability
- (new_edge->probability);
+ next_edge->probability
+ = inverse_probability (new_edge->probability);
next_edge->count = edge_counts[0]
+ sum_slice <gcov_type> (edge_counts, i, labels);
next_bb->frequency = EDGE_FREQUENCY (next_edge);
@@ -6403,14 +6435,16 @@ expand_builtins ()
tree imag_part = make_temp_ssa_name (rhs_type, NULL,
"cexp_imag_part");
- tree cos_fndecl = mathfn_built_in
- (rhs_type, fn == float_type_p ? BUILT_IN_COSF : BUILT_IN_COS);
+ tree cos_fndecl
+ = mathfn_built_in (rhs_type, fn == float_type_p
+ ? BUILT_IN_COSF : BUILT_IN_COS);
gcall *cos = gimple_build_call (cos_fndecl, 1, rhs);
gimple_call_set_lhs (cos, real_part);
gsi_insert_before (&gsi, cos, GSI_SAME_STMT);
- tree sin_fndecl = mathfn_built_in
- (rhs_type, fn == float_type_p ? BUILT_IN_SINF : BUILT_IN_SIN);
+ tree sin_fndecl
+ = mathfn_built_in (rhs_type, fn == float_type_p
+ ? BUILT_IN_SINF : BUILT_IN_SIN);
gcall *sin = gimple_build_call (sin_fndecl, 1, rhs);
gimple_call_set_lhs (sin, imag_part);
gsi_insert_before (&gsi, sin, GSI_SAME_STMT);
@@ -6489,8 +6523,8 @@ generate_hsa (bool kernel)
if (hsa_cfun->m_kern_p)
{
- hsa_function_summary *s = hsa_summaries->get
- (cgraph_node::get (hsa_cfun->m_decl));
+ hsa_function_summary *s
+ = hsa_summaries->get (cgraph_node::get (hsa_cfun->m_decl));
hsa_add_kern_decl_mapping (current_function_decl, hsa_cfun->m_name,
hsa_cfun->m_maximum_omp_data_size,
s->m_gridified_kernel_p);
@@ -6525,9 +6559,9 @@ const pass_data pass_data_gen_hsail =
{
GIMPLE_PASS,
"hsagen", /* name */
- OPTGROUP_NONE, /* optinfo_flags */
+ OPTGROUP_NONE, /* optinfo_flags */
TV_NONE, /* tv_id */
- PROP_cfg | PROP_ssa, /* properties_required */
+ PROP_cfg | PROP_ssa, /* properties_required */
0, /* properties_provided */
0, /* properties_destroyed */
0, /* todo_flags_start */
@@ -6559,8 +6593,8 @@ pass_gen_hsail::gate (function *f)
unsigned int
pass_gen_hsail::execute (function *)
{
- hsa_function_summary *s = hsa_summaries->get
- (cgraph_node::get_create (current_function_decl));
+ hsa_function_summary *s
+ = hsa_summaries->get (cgraph_node::get_create (current_function_decl));
convert_switch_statements ();
expand_builtins ();
diff --git a/gcc/hsa.c b/gcc/hsa.c
index 0dd6184..ec23f81 100644
--- a/gcc/hsa.c
+++ b/gcc/hsa.c
@@ -41,7 +41,7 @@ along with GCC; see the file COPYING3. If not see
#include "ctype.h"
/* Structure containing intermediate HSA representation of the generated
- function. */
+ function. */
class hsa_function_representation *hsa_cfun;
/* Element of the mapping vector between a host decl and an HSA kernel. */
@@ -61,7 +61,8 @@ struct GTY(()) hsa_decl_kernel_map_element
/* Mapping between decls and corresponding HSA kernels in this compilation
unit. */
-static GTY (()) vec<hsa_decl_kernel_map_element, va_gc> *hsa_decl_kernel_mapping;
+static GTY (()) vec<hsa_decl_kernel_map_element, va_gc>
+ *hsa_decl_kernel_mapping;
/* Mapping between decls and corresponding HSA kernels
called by the function. */
@@ -118,9 +119,10 @@ hsa_deinit_compilation_unit_data (void)
delete hsa_failed_functions;
delete hsa_emitted_internal_decls;
- for (hash_table <hsa_noop_symbol_hasher>::iterator it =
- hsa_global_variable_symbols->begin ();
- it != hsa_global_variable_symbols->end (); ++it)
+ for (hash_table <hsa_noop_symbol_hasher>::iterator it
+ = hsa_global_variable_symbols->begin ();
+ it != hsa_global_variable_symbols->end ();
+ ++it)
{
hsa_symbol *sym = *it;
delete sym;
@@ -727,10 +729,9 @@ hsa_get_declaration_name (tree decl)
{
if (!DECL_NAME (decl))
{
- char *b = XNEWVEC (char, 64);
- sprintf (b, "__hsa_anonymous_%i", DECL_UID (decl));
- const char *ggc_str = ggc_alloc_string (b, strlen (b) + 1);
- free (b);
+ char buf[64];
+ snprintf (buf, 64, "__hsa_anonymous_%i", DECL_UID (decl));
+ const char *ggc_str = ggc_strdup (buf);
return ggc_str;
}
@@ -888,6 +889,8 @@ hsa_internal_fn::get_arity ()
case IFN_PARITY:
case IFN_POPCOUNT:
default:
+ /* As we produce sorry message for unknown internal functions,
+ reaching this label is definitely a bug. */
gcc_unreachable ();
}
}
@@ -935,6 +938,8 @@ hsa_internal_fn::get_argument_type (int n)
return BRIG_TYPE_S32;
}
default:
+ /* As we produce sorry message for unknown internal functions,
+ reaching this label is definitely a bug. */
gcc_unreachable ();
}
}
diff --git a/gcc/hsa.h b/gcc/hsa.h
index 29b29e2..f0436f3 100644
--- a/gcc/hsa.h
+++ b/gcc/hsa.h
@@ -40,7 +40,7 @@ hsa_gen_requested_p (void)
return !flag_disable_hsa;
}
-/* Standard warning message if we failed to generate HSAIL for a function */
+/* Standard warning message if we failed to generate HSAIL for a function. */
#define HSA_SORRY_MSG "could not emit HSAIL for the function"
@@ -75,7 +75,7 @@ struct hsa_symbol
tree m_decl;
/* Name of the symbol, that will be written into output and dumps. Can be
- NULL, see name_number below.*/
+ NULL, see name_number below. */
const char *m_name;
/* If name is NULL, artificial name will be formed from the segment name and
@@ -115,7 +115,7 @@ private:
hsa_symbol ();
};
-/* Abstract class for HSA instruction operands. */
+/* Abstract class for HSA instruction operands. */
class hsa_op_base
{
@@ -231,7 +231,7 @@ public:
be 'c', 's', 'd' or 'q'. */
char m_reg_class;
/* If allocated, the number of the HW register (within its HSA register
- class). */
+ class). */
char m_hard_num;
private:
@@ -318,7 +318,7 @@ is_a_helper <hsa_op_address *>::test (hsa_op_base *p)
return p->m_kind == BRIG_KIND_OPERAND_ADDRESS;
}
-/* A reference to code HSA operand. It can be either reference
+/* A reference to code HSA operand. It can be either reference
to a start of a BB or a start of a function. */
class hsa_op_code_ref : public hsa_op_base
@@ -624,10 +624,10 @@ public:
/* The segment is of the memory access is either the segment of the symbol in
the address operand or flat address is there is no symbol there. */
- /* Required alignment of the memory operation. */
+ /* Required alignment of the memory operation. */
BrigAlignment8_t m_align;
- /* HSA equiv class, basically an alias set number. */
+ /* HSA equiv class, basically an alias set number. */
uint8_t m_equiv_class;
/* TODO: Add width modifier, perhaps also other things. */
@@ -672,7 +672,7 @@ public:
/* Things like acquire/release/aligned. */
enum BrigMemoryOrder m_memoryorder;
- /* Scope of the atomic operation. */
+ /* Scope of the atomic operation. */
enum BrigMemoryScope m_memoryscope;
private:
@@ -801,7 +801,7 @@ public:
void *operator new (size_t);
- /* Called function */
+ /* Called function. */
tree m_called_function;
/* Called internal function. */
@@ -1013,7 +1013,7 @@ public:
hsa_insn_alloca (hsa_op_with_type *dest, hsa_op_with_type *size,
unsigned alignment = 0);
- /* Required alignment of the allocation. */
+ /* Required alignment of the allocation. */
BrigAlignment8_t m_align;
/* Pool allocator. */
--
2.6.4