Google Git
Sign in
gnu / gcc / refs/heads/releases/gcc-9 / . / gcc / hsa-gen.c
blob: 5d4bc1979c75baacb04b51d5fbf655487651de6a [file] [log] [blame]
/* A pass for lowering gimple to HSAIL
Copyright (C) 2013-2019 Free Software Foundation, Inc.
Contributed by Martin Jambor <mjambor@suse.cz> and
Martin Liska <mliska@suse.cz>.
This file is part of GCC.
GCC is free software; you can redistribute it and/or modify
it under the terms of the GNU General Public License as published by
the Free Software Foundation; either version 3, or (at your option)
any later version.
GCC is distributed in the hope that it will be useful,
but WITHOUT ANY WARRANTY; without even the implied warranty of
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
GNU General Public License for more details.
You should have received a copy of the GNU General Public License
along with GCC; see the file COPYING3. If not see
<http://www.gnu.org/licenses/>. */
#include "config.h"
#include "system.h"
#include "coretypes.h"
#include "memmodel.h"
#include "tm.h"
#include "is-a.h"
#include "hash-table.h"
#include "vec.h"
#include "tree.h"
#include "tree-pass.h"
#include "function.h"
#include "basic-block.h"
#include "cfg.h"
#include "fold-const.h"
#include "gimple.h"
#include "gimple-iterator.h"
#include "bitmap.h"
#include "dumpfile.h"
#include "gimple-pretty-print.h"
#include "diagnostic-core.h"
#include "gimple-ssa.h"
#include "tree-phinodes.h"
#include "stringpool.h"
#include "tree-vrp.h"
#include "tree-ssanames.h"
#include "tree-dfa.h"
#include "ssa-iterators.h"
#include "cgraph.h"
#include "print-tree.h"
#include "symbol-summary.h"
#include "hsa-common.h"
#include "cfghooks.h"
#include "tree-cfg.h"
#include "cfgloop.h"
#include "cfganal.h"
#include "builtins.h"
#include "params.h"
#include "gomp-constants.h"
#include "internal-fn.h"
#include "builtins.h"
#include "stor-layout.h"
#include "stringpool.h"
#include "attribs.h"
/* Print a warning message and set that we have seen an error. */
#define HSA_SORRY_ATV(location, message, ...) \
do \
{ \
hsa_fail_cfun (); \
auto_diagnostic_group d; \
if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
HSA_SORRY_MSG)) \
inform (location, message, __VA_ARGS__); \
} \
while (false)
/* Same as previous, but highlight a location. */
#define HSA_SORRY_AT(location, message) \
do \
{ \
hsa_fail_cfun (); \
auto_diagnostic_group d; \
if (warning_at (EXPR_LOCATION (hsa_cfun->m_decl), OPT_Whsa, \
HSA_SORRY_MSG)) \
inform (location, message); \
} \
while (false)
/* Default number of threads used by kernel dispatch. */
#define HSA_DEFAULT_NUM_THREADS 64
/* Following structures are defined in the final version
of HSA specification. */
/* HSA queue packet is shadow structure, originally provided by AMD. */
struct hsa_queue_packet
{
uint16_t header;
uint16_t setup;
uint16_t workgroup_size_x;
uint16_t workgroup_size_y;
uint16_t workgroup_size_z;
uint16_t reserved0;
uint32_t grid_size_x;
uint32_t grid_size_y;
uint32_t grid_size_z;
uint32_t private_segment_size;
uint32_t group_segment_size;
uint64_t kernel_object;
void *kernarg_address;
uint64_t reserved2;
uint64_t completion_signal;
};
/* HSA queue is shadow structure, originally provided by AMD. */
struct hsa_queue
{
int type;
uint32_t features;
void *base_address;
uint64_t doorbell_signal;
uint32_t size;
uint32_t reserved1;
uint64_t id;
};
static struct obstack hsa_obstack;
/* List of pointers to all instructions that come from an object allocator. */
static vec <hsa_insn_basic *> hsa_instructions;
/* List of pointers to all operands that come from an object allocator. */
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_emitted_to_brig (false)
{
}
hsa_symbol::hsa_symbol (BrigType16_t type, BrigSegment8_t segment,
BrigLinkage8_t linkage, bool global_scope_p,
BrigAllocation allocation, BrigAlignment8_t align)
: 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_emitted_to_brig (false), m_align (align)
{
}
unsigned HOST_WIDE_INT
hsa_symbol::total_byte_size ()
{
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;
if (m_dim)
s *= m_dim;
return s;
}
/* Forward declaration. */
static BrigType16_t
hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p,
bool min32int);
void
hsa_symbol::fillup_for_decl (tree decl)
{
m_decl = decl;
m_type = hsa_type_for_tree_type (TREE_TYPE (decl), &m_dim, false);
if (hsa_seen_error ())
{
m_seen_error = true;
return;
}
m_align = MAX (m_align, hsa_natural_alignment (m_type));
}
/* Constructor of class representing global HSA function/kernel information and
state. FNDECL is function declaration, KERNEL_P is true if the function
is going to become a HSA kernel. If the function has body, SSA_NAMES_COUNT
should be set to number of SSA names used in the function.
MODIFIED_CFG is set to true in case we modified control-flow graph
of the function. */
hsa_function_representation::hsa_function_representation
(tree fdecl, bool kernel_p, unsigned ssa_names_count, bool modified_cfg)
: 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 (),
m_modified_cfg (modified_cfg)
{
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);
m_ssa_map.safe_grow_cleared (ssa_names_count);
}
/* 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 () {}
/* Destructor of class holding function/kernel-wide information and state. */
hsa_function_representation::~hsa_function_representation ()
{
/* Kernel names are deallocated at the end of BRIG output when deallocating
hsa_decl_kernel_mapping. */
if (!m_kern_p || m_seen_error)
free (m_name);
for (unsigned i = 0; i < m_input_args.length (); i++)
delete m_input_args[i];
m_input_args.release ();
delete m_output_arg;
delete m_local_symbols;
for (unsigned i = 0; i < m_spill_symbols.length (); i++)
delete m_spill_symbols[i];
m_spill_symbols.release ();
hsa_symbol *sym;
for (unsigned i = 0; i < m_global_symbols.iterate (i, &sym); i++)
if (sym->m_linkage != BRIG_ALLOCATION_PROGRAM)
delete sym;
m_global_symbols.release ();
for (unsigned i = 0; i < m_private_variables.length (); i++)
delete m_private_variables[i];
m_private_variables.release ();
m_called_functions.release ();
m_ssa_map.release ();
for (unsigned i = 0; i < m_called_internal_fns.length (); i++)
delete m_called_internal_fns[i];
}
hsa_op_reg *
hsa_function_representation::get_shadow_reg ()
{
/* If we compile a function with kernel dispatch and does not set
an optimization level, the function won't be inlined and
we return NULL. */
if (!m_kern_p)
return NULL;
if (m_shadow_reg)
return m_shadow_reg;
/* Append the shadow argument. */
hsa_symbol *shadow = new hsa_symbol (BRIG_TYPE_U64, BRIG_SEGMENT_KERNARG,
BRIG_LINKAGE_FUNCTION);
m_input_args.safe_push (shadow);
shadow->m_name = "hsa_runtime_shadow";
hsa_op_reg *r = new hsa_op_reg (BRIG_TYPE_U64);
hsa_op_address *addr = new hsa_op_address (shadow);
hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, BRIG_TYPE_U64, r, addr);
hsa_bb_for_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun))->append_insn (mem);
m_shadow_reg = r;
return r;
}
bool hsa_function_representation::has_shadow_reg_p ()
{
return m_shadow_reg != NULL;
}
void
hsa_function_representation::init_extra_bbs ()
{
hsa_init_new_bb (ENTRY_BLOCK_PTR_FOR_FN (cfun));
hsa_init_new_bb (EXIT_BLOCK_PTR_FOR_FN (cfun));
}
void
hsa_function_representation::update_dominance ()
{
if (m_modified_cfg)
{
free_dominance_info (CDI_DOMINATORS);
calculate_dominance_info (CDI_DOMINATORS);
}
}
hsa_symbol *
hsa_function_representation::create_hsa_temporary (BrigType16_t type)
{
hsa_symbol *s = new hsa_symbol (type, BRIG_SEGMENT_PRIVATE,
BRIG_LINKAGE_FUNCTION);
s->m_name_number = m_temp_symbol_count++;
hsa_cfun->m_private_variables.safe_push (s);
return s;
}
BrigLinkage8_t
hsa_function_representation::get_linkage ()
{
if (m_internal_fn)
return BRIG_LINKAGE_PROGRAM;
return m_kern_p || TREE_PUBLIC (m_decl) ?
BRIG_LINKAGE_PROGRAM : BRIG_LINKAGE_MODULE;
}
/* Hash map of simple OMP builtins. */
static hash_map <nofree_string_hash, omp_simple_builtin> *omp_simple_builtins
= NULL;
/* Warning messages for OMP builtins. */
#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 feateres"
/* Initialize hash map with simple OMP builtins. */
static void
hsa_init_simple_builtins ()
{
if (omp_simple_builtins != NULL)
return;
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_is_initial_device", NULL, false,
new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_get_dynamic", NULL, false,
new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_set_dynamic", NULL, false, NULL),
omp_simple_builtin ("omp_init_lock", HSA_WARN_LOCK_ROUTINE, true),
omp_simple_builtin ("omp_init_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
true),
omp_simple_builtin ("omp_init_nest_lock_with_hint", HSA_WARN_LOCK_ROUTINE,
true),
omp_simple_builtin ("omp_destroy_lock", HSA_WARN_LOCK_ROUTINE, true),
omp_simple_builtin ("omp_set_lock", HSA_WARN_LOCK_ROUTINE, true),
omp_simple_builtin ("omp_unset_lock", HSA_WARN_LOCK_ROUTINE, true),
omp_simple_builtin ("omp_test_lock", HSA_WARN_LOCK_ROUTINE, true),
omp_simple_builtin ("omp_get_wtime", HSA_WARN_TIMING_ROUTINE, true),
omp_simple_builtin ("omp_get_wtick", HSA_WARN_TIMING_ROUTINE, true),
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_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_get_max_active_levels", NULL, false,
new hsa_op_immed (1, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_in_final", NULL, false,
new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_get_proc_bind", HSA_WARN_AFFINITY, false,
new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_get_num_places", HSA_WARN_AFFINITY, false,
new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_get_place_num_procs", HSA_WARN_AFFINITY, false,
new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_get_place_proc_ids", HSA_WARN_AFFINITY, false,
NULL),
omp_simple_builtin ("omp_get_place_num", HSA_WARN_AFFINITY, false,
new hsa_op_immed (-1, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_get_partition_num_places", HSA_WARN_AFFINITY,
false,
new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_get_partition_place_nums", HSA_WARN_AFFINITY,
false, NULL),
omp_simple_builtin ("omp_set_default_device",
"omp_set_default_device has undefined semantics "
"within target regions, support for HSA ignores it",
false, NULL),
omp_simple_builtin ("omp_get_default_device",
"omp_get_default_device has undefined semantics "
"within target regions, support for HSA ignores it",
false,
new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_get_num_devices",
"omp_get_num_devices has undefined semantics "
"within target regions, support for HSA ignores it",
false,
new hsa_op_immed (0, (BrigType16_t) BRIG_TYPE_S32)),
omp_simple_builtin ("omp_get_num_procs", NULL, true, NULL),
omp_simple_builtin ("omp_get_cancellation", NULL, true, NULL),
omp_simple_builtin ("omp_set_nested", NULL, true, NULL),
omp_simple_builtin ("omp_get_nested", NULL, true, NULL),
omp_simple_builtin ("omp_set_schedule", NULL, true, NULL),
omp_simple_builtin ("omp_get_schedule", NULL, true, NULL),
omp_simple_builtin ("omp_get_thread_limit", NULL, true, NULL),
omp_simple_builtin ("omp_get_team_size", NULL, true, NULL),
omp_simple_builtin ("omp_get_ancestor_thread_num", NULL, true, NULL),
omp_simple_builtin ("omp_get_max_task_priority", NULL, true, NULL)
};
unsigned count = sizeof (omp_builtins) / sizeof (omp_simple_builtin);
for (unsigned i = 0; i < count; i++)
omp_simple_builtins->put (omp_builtins[i].m_name, omp_builtins[i]);
}
/* Allocate HSA structures that we need only while generating with this. */
static void
hsa_init_data_for_cfun ()
{
hsa_init_compilation_unit_data ();
gcc_obstack_init (&hsa_obstack);
}
/* Deinitialize HSA subsystem and free all allocated memory. */
static void
hsa_deinit_data_for_cfun (void)
{
basic_block bb;
FOR_ALL_BB_FN (bb, cfun)
if (bb->aux)
{
hsa_bb *hbb = hsa_bb_for_bb (bb);
hbb->~hsa_bb ();
bb->aux = NULL;
}
for (unsigned int i = 0; i < hsa_operands.length (); i++)
hsa_destroy_operand (hsa_operands[i]);
hsa_operands.release ();
for (unsigned i = 0; i < hsa_instructions.length (); i++)
hsa_destroy_insn (hsa_instructions[i]);
hsa_instructions.release ();
if (omp_simple_builtins != NULL)
{
delete omp_simple_builtins;
omp_simple_builtins = NULL;
}
obstack_free (&hsa_obstack, NULL);
delete hsa_cfun;
}
/* Return the type which holds addresses in the given SEGMENT. */
static BrigType16_t
hsa_get_segment_addr_type (BrigSegment8_t segment)
{
switch (segment)
{
case BRIG_SEGMENT_NONE:
gcc_unreachable ();
case BRIG_SEGMENT_FLAT:
case BRIG_SEGMENT_GLOBAL:
case BRIG_SEGMENT_READONLY:
case BRIG_SEGMENT_KERNARG:
return hsa_machine_large_p () ? BRIG_TYPE_U64 : BRIG_TYPE_U32;
case BRIG_SEGMENT_GROUP:
case BRIG_SEGMENT_PRIVATE:
case BRIG_SEGMENT_SPILL:
case BRIG_SEGMENT_ARG:
return BRIG_TYPE_U32;
}
gcc_unreachable ();
}
/* Return integer brig type according to provided SIZE in bytes. If SIGN
is set to true, return signed integer type. */
static BrigType16_t
get_integer_type_by_bytes (unsigned size, bool sign)
{
if (sign)
switch (size)
{
case 1:
return BRIG_TYPE_S8;
case 2:
return BRIG_TYPE_S16;
case 4:
return BRIG_TYPE_S32;
case 8:
return BRIG_TYPE_S64;
default:
break;
}
else
switch (size)
{
case 1:
return BRIG_TYPE_U8;
case 2:
return BRIG_TYPE_U16;
case 4:
return BRIG_TYPE_U32;
case 8:
return BRIG_TYPE_U64;
default:
break;
}
return 0;
}
/* If T points to an integral type smaller than 32 bits, change it to a 32bit
equivalent and return the result. Otherwise just return the result. */
static BrigType16_t
hsa_extend_inttype_to_32bit (BrigType16_t t)
{
if (t == BRIG_TYPE_U8 || t == BRIG_TYPE_U16)
return BRIG_TYPE_U32;
else if (t == BRIG_TYPE_S8 || t == BRIG_TYPE_S16)
return BRIG_TYPE_S32;
return t;
}
/* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
are assumed to use flat addressing. If min32int is true, always expand
integer types to one that has at least 32 bits. */
static BrigType16_t
hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
{
HOST_WIDE_INT bsize;
const_tree base;
BrigType16_t res = BRIG_TYPE_NONE;
gcc_checking_assert (TYPE_P (type));
gcc_checking_assert (!AGGREGATE_TYPE_P (type));
if (POINTER_TYPE_P (type))
return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
if (TREE_CODE (type) == VECTOR_TYPE)
base = TREE_TYPE (type);
else if (TREE_CODE (type) == COMPLEX_TYPE)
{
base = TREE_TYPE (type);
min32int = true;
}
else
base = type;
if (!tree_fits_uhwi_p (TYPE_SIZE (base)))
{
HSA_SORRY_ATV (EXPR_LOCATION (type),
"support for HSA does not implement huge or "
"variable-sized type %qT", type);
return res;
}
bsize = tree_to_uhwi (TYPE_SIZE (base));
unsigned byte_size = bsize / BITS_PER_UNIT;
if (INTEGRAL_TYPE_P (base))
res = get_integer_type_by_bytes (byte_size, !TYPE_UNSIGNED (base));
else if (SCALAR_FLOAT_TYPE_P (base))
{
switch (bsize)
{
case 16:
res = BRIG_TYPE_F16;
break;
case 32:
res = BRIG_TYPE_F32;
break;
case 64:
res = BRIG_TYPE_F64;
break;
default:
break;
}
}
if (res == BRIG_TYPE_NONE)
{
HSA_SORRY_ATV (EXPR_LOCATION (type),
"support for HSA does not implement type %qT", type);
return res;
}
if (TREE_CODE (type) == VECTOR_TYPE)
{
HOST_WIDE_INT tsize = tree_to_uhwi (TYPE_SIZE (type));
if (bsize == tsize)
{
HSA_SORRY_ATV (EXPR_LOCATION (type),
"support for HSA does not implement a vector type "
"where a type and unit size are equal: %qT", type);
return res;
}
switch (tsize)
{
case 32:
res |= BRIG_TYPE_PACK_32;
break;
case 64:
res |= BRIG_TYPE_PACK_64;
break;
case 128:
res |= BRIG_TYPE_PACK_128;
break;
default:
HSA_SORRY_ATV (EXPR_LOCATION (type),
"support for HSA does not implement type %qT", type);
}
}
if (min32int)
/* Registers/immediate operands can only be 32bit or more except for
f16. */
res = hsa_extend_inttype_to_32bit (res);
if (TREE_CODE (type) == COMPLEX_TYPE)
{
unsigned bsize = 2 * hsa_type_bit_size (res);
res = hsa_bittype_for_bitsize (bsize);
}
return res;
}
/* Returns the BRIG type we need to load/store entities of TYPE. */
static BrigType16_t
mem_type_for_type (BrigType16_t type)
{
/* HSA has non-intuitive constraints on load/store types. If it's
a bit-type it _must_ be B128, if it's not a bit-type it must be
64bit max. So for loading entities of 128 bits (e.g. vectors)
we have to use B128, while for loading the rest we have to use the
input type (??? or maybe also flattened to a equally sized non-vector
unsigned type?). */
if ((type & BRIG_TYPE_PACK_MASK) == BRIG_TYPE_PACK_128)
return BRIG_TYPE_B128;
else if (hsa_btype_p (type) || hsa_type_packed_p (type))
{
unsigned bitsize = hsa_type_bit_size (type);
if (bitsize < 128)
return hsa_uint_for_bitsize (bitsize);
else
return hsa_bittype_for_bitsize (bitsize);
}
return type;
}
/* Return HSA type for tree TYPE. If it cannot fit into BrigType16_t, some
kind of array will be generated, setting DIM appropriately. Otherwise, it
will be set to zero. */
static BrigType16_t
hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p = NULL,
bool min32int = false)
{
gcc_checking_assert (TYPE_P (type));
if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type)))
{
HSA_SORRY_ATV (EXPR_LOCATION (type), "support for HSA does not "
"implement huge or variable-sized type %qT", type);
return BRIG_TYPE_NONE;
}
if (RECORD_OR_UNION_TYPE_P (type))
{
if (dim_p)
*dim_p = tree_to_uhwi (TYPE_SIZE_UNIT (type));
return BRIG_TYPE_U8 | BRIG_TYPE_ARRAY;
}
if (TREE_CODE (type) == ARRAY_TYPE)
{
/* We try to be nice and use the real base-type when this is an array of
scalars and only resort to an array of bytes if the type is more
complex. */
unsigned HOST_WIDE_INT dim = 1;
while (TREE_CODE (type) == ARRAY_TYPE)
{
tree domain = TYPE_DOMAIN (type);
if (!TYPE_MIN_VALUE (domain)
|| !TYPE_MAX_VALUE (domain)
|| !tree_fits_shwi_p (TYPE_MIN_VALUE (domain))
|| !tree_fits_shwi_p (TYPE_MAX_VALUE (domain)))
{
HSA_SORRY_ATV (EXPR_LOCATION (type),
"support for HSA does not implement array "
"%qT with unknown bounds", type);
return BRIG_TYPE_NONE;
}
HOST_WIDE_INT min = tree_to_shwi (TYPE_MIN_VALUE (domain));
HOST_WIDE_INT max = tree_to_shwi (TYPE_MAX_VALUE (domain));
dim = dim * (unsigned HOST_WIDE_INT) (max - min + 1);
type = TREE_TYPE (type);
}
BrigType16_t res;
if (RECORD_OR_UNION_TYPE_P (type))
{
dim = dim * tree_to_uhwi (TYPE_SIZE_UNIT (type));
res = BRIG_TYPE_U8;
}
else
res = hsa_type_for_scalar_tree_type (type, false);
if (dim_p)
*dim_p = dim;
return res | BRIG_TYPE_ARRAY;
}
/* Scalar case: */
if (dim_p)
*dim_p = 0;
return hsa_type_for_scalar_tree_type (type, min32int);
}
/* Returns true if converting from STYPE into DTYPE needs the _CVT
opcode. If false a normal _MOV is enough. */
static bool
hsa_needs_cvt (BrigType16_t dtype, BrigType16_t stype)
{
if (hsa_btype_p (dtype))
return false;
/* float <-> int conversions are real converts. */
if (hsa_type_float_p (dtype) != hsa_type_float_p (stype))
return true;
/* When both types have different size, then we need CVT as well. */
if (hsa_type_bit_size (dtype) != hsa_type_bit_size (stype))
return true;
return false;
}
/* Return declaration name if it exists or create one from UID if it does not.
If DECL is a local variable, make UID part of its name. */
const char *
hsa_get_declaration_name (tree decl)
{
if (!DECL_NAME (decl))
{
char buf[64];
snprintf (buf, 64, "__hsa_anon_%u", DECL_UID (decl));
size_t len = strlen (buf);
char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1);
memcpy (copy, buf, len + 1);
return copy;
}
tree name_tree;
if (TREE_CODE (decl) == FUNCTION_DECL
|| (TREE_CODE (decl) == VAR_DECL && is_global_var (decl)))
name_tree = DECL_ASSEMBLER_NAME (decl);
else
name_tree = DECL_NAME (decl);
const char *name = IDENTIFIER_POINTER (name_tree);
/* User-defined assembly names have prepended asterisk symbol. */
if (name[0] == '*')
name++;
if ((TREE_CODE (decl) == VAR_DECL)
&& decl_function_context (decl))
{
size_t len = strlen (name);
char *buf = (char *) alloca (len + 32);
snprintf (buf, len + 32, "%s_%u", name, DECL_UID (decl));
len = strlen (buf);
char *copy = (char *) obstack_alloc (&hsa_obstack, len + 1);
memcpy (copy, buf, len + 1);
return copy;
}
else
return name;
}
/* Lookup or create the associated hsa_symbol structure with a given VAR_DECL
or lookup the hsa_structure corresponding to a PARM_DECL. */
static hsa_symbol *
get_symbol_for_decl (tree decl)
{
hsa_symbol **slot;
hsa_symbol dummy (BRIG_TYPE_NONE, BRIG_SEGMENT_NONE, BRIG_LINKAGE_NONE);
gcc_assert (TREE_CODE (decl) == PARM_DECL
|| TREE_CODE (decl) == RESULT_DECL
|| TREE_CODE (decl) == VAR_DECL
|| TREE_CODE (decl) == CONST_DECL);
dummy.m_decl = decl;
bool is_in_global_vars = ((TREE_CODE (decl) == VAR_DECL)
&& !decl_function_context (decl));
if (is_in_global_vars)
slot = hsa_global_variable_symbols->find_slot (&dummy, INSERT);
else
slot = hsa_cfun->m_local_symbols->find_slot (&dummy, INSERT);
gcc_checking_assert (slot);
if (*slot)
{
hsa_symbol *sym = (*slot);
/* If the symbol is problematic, mark current function also as
problematic. */
if (sym->m_seen_error)
hsa_fail_cfun ();
/* PR hsa/70234: If a global variable was marked to be emitted,
but HSAIL generation of a function using the variable fails,
we should retry to emit the variable in context of a different
function.
Iterate elements whether a symbol is already in m_global_symbols
of not. */
if (is_in_global_vars && !sym->m_emitted_to_brig)
{
for (unsigned i = 0; i < hsa_cfun->m_global_symbols.length (); i++)
if (hsa_cfun->m_global_symbols[i] == sym)
return *slot;
hsa_cfun->m_global_symbols.safe_push (sym);
}
return *slot;
}
else
{
hsa_symbol *sym;
/* PARM_DECLs and RESULT_DECL should be already in m_local_symbols. */
gcc_assert (TREE_CODE (decl) == VAR_DECL
|| TREE_CODE (decl) == CONST_DECL);
BrigAlignment8_t align = hsa_object_alignment (decl);
if (is_in_global_vars)
{
gcc_checking_assert (TREE_CODE (decl) != CONST_DECL);
sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_GLOBAL,
BRIG_LINKAGE_PROGRAM, true,
BRIG_ALLOCATION_PROGRAM, align);
hsa_cfun->m_global_symbols.safe_push (sym);
sym->fillup_for_decl (decl);
if (sym->m_align > align)
{
sym->m_seen_error = true;
HSA_SORRY_ATV (EXPR_LOCATION (decl),
"HSA specification requires that %E is at least "
"naturally aligned", decl);
}
}
else
{
/* As generation of efficient memory copy instructions relies
on alignment greater or equal to 8 bytes,
we need to increase alignment of all aggregate types.. */
if (AGGREGATE_TYPE_P (TREE_TYPE (decl)))
align = MAX ((BrigAlignment8_t) BRIG_ALIGNMENT_8, align);
BrigAllocation allocation = BRIG_ALLOCATION_AUTOMATIC;
BrigSegment8_t segment;
if (TREE_CODE (decl) == CONST_DECL)
{
segment = BRIG_SEGMENT_READONLY;
allocation = BRIG_ALLOCATION_AGENT;
}
else if (lookup_attribute ("hsa_group_segment",
DECL_ATTRIBUTES (decl)))
segment = BRIG_SEGMENT_GROUP;
else if (TREE_STATIC (decl))
{
segment = BRIG_SEGMENT_GLOBAL;
allocation = BRIG_ALLOCATION_PROGRAM;
}
else if (lookup_attribute ("hsa_global_segment",
DECL_ATTRIBUTES (decl)))
segment = BRIG_SEGMENT_GLOBAL;
else
segment = BRIG_SEGMENT_PRIVATE;
sym = new hsa_symbol (BRIG_TYPE_NONE, segment, BRIG_LINKAGE_FUNCTION,
false, allocation, align);
sym->fillup_for_decl (decl);
hsa_cfun->m_private_variables.safe_push (sym);
}
sym->m_name = hsa_get_declaration_name (decl);
*slot = sym;
return sym;
}
}
/* For a given HSA function declaration, return a host
function declaration. */
tree
hsa_get_host_function (tree decl)
{
hsa_function_summary *s = hsa_summaries->get (cgraph_node::get_create (decl));
gcc_assert (s->m_gpu_implementation_p);
return s->m_bound_function ? s->m_bound_function->decl : NULL;
}
/* Return true if function DECL has a host equivalent function. */
static char *
get_brig_function_name (tree decl)
{
tree d = decl;
hsa_function_summary *s = hsa_summaries->get (cgraph_node::get_create (d));
if (s != NULL
&& s->m_gpu_implementation_p
&& s->m_bound_function)
d = s->m_bound_function->decl;
/* IPA split can create a function that has no host equivalent. */
if (d == NULL)
d = decl;
char *name = xstrdup (hsa_get_declaration_name (d));
hsa_sanitize_name (name);
return name;
}
/* Create a spill symbol of type TYPE. */
hsa_symbol *
hsa_get_spill_symbol (BrigType16_t type)
{
hsa_symbol *sym = new hsa_symbol (type, BRIG_SEGMENT_SPILL,
BRIG_LINKAGE_FUNCTION);
hsa_cfun->m_spill_symbols.safe_push (sym);
return sym;
}
/* Create a symbol for a read-only string constant. */
hsa_symbol *
hsa_get_string_cst_symbol (tree string_cst)
{
gcc_checking_assert (TREE_CODE (string_cst) == STRING_CST);
hsa_symbol **slot = hsa_cfun->m_string_constants_map.get (string_cst);
if (slot)
return *slot;
hsa_op_immed *cst = new hsa_op_immed (string_cst);
hsa_symbol *sym = new hsa_symbol (cst->m_type, BRIG_SEGMENT_GLOBAL,
BRIG_LINKAGE_MODULE, true,
BRIG_ALLOCATION_AGENT);
sym->m_cst_value = cst;
sym->m_dim = TREE_STRING_LENGTH (string_cst);
sym->m_name_number = hsa_cfun->m_global_symbols.length ();
hsa_cfun->m_global_symbols.safe_push (sym);
hsa_cfun->m_string_constants_map.put (string_cst, sym);
return sym;
}
/* Make the type of a MOV instruction larger if mandated by HSAIL rules. */
static void
hsa_fixup_mov_insn_type (hsa_insn_basic *insn)
{
insn->m_type = hsa_extend_inttype_to_32bit (insn->m_type);
if (insn->m_type == BRIG_TYPE_B8 || insn->m_type == BRIG_TYPE_B16)
insn->m_type = BRIG_TYPE_B32;
}
/* 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_operands.safe_push (this);
}
/* Constructor of ancestor of all operands which have a type. K is BRIG kind
that identified what the operator is. T is the type of the operator. */
hsa_op_with_type::hsa_op_with_type (BrigKind16_t k, BrigType16_t t)
: hsa_op_base (k), m_type (t)
{
}
hsa_op_with_type *
hsa_op_with_type::get_in_type (BrigType16_t dtype, hsa_bb *hbb)
{
if (m_type == dtype)
return this;
hsa_op_reg *dest;
if (hsa_needs_cvt (dtype, m_type))
{
dest = new hsa_op_reg (dtype);
hbb->append_insn (new hsa_insn_cvt (dest, this));
}
else if (is_a <hsa_op_reg *> (this))
{
/* In the end, HSA registers do not really have types, only sizes, so if
the sizes match, we can use the register directly. */
gcc_checking_assert (hsa_type_bit_size (dtype)
== hsa_type_bit_size (m_type));
return this;
}
else
{
dest = new hsa_op_reg (m_type);
hsa_insn_basic *mov = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
dest->m_type, dest, this);
hsa_fixup_mov_insn_type (mov);
hbb->append_insn (mov);
/* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
type of the operand must be same as type of the instruction. */
dest->m_type = dtype;
}
return dest;
}
/* If this operand has integer type smaller than 32 bits, extend it to 32 bits,
adding instructions to HBB if needed. */
hsa_op_with_type *
hsa_op_with_type::extend_int_to_32bit (hsa_bb *hbb)
{
if (m_type == BRIG_TYPE_U8 || m_type == BRIG_TYPE_U16)
return get_in_type (BRIG_TYPE_U32, hbb);
else if (m_type == BRIG_TYPE_S8 || m_type == BRIG_TYPE_S16)
return get_in_type (BRIG_TYPE_S32, hbb);
else
return this;
}
/* Constructor of class representing HSA immediate values. TREE_VAL is the
tree representation of the immediate value. If min32int is true,
always expand integer types to one that has at least 32 bits. */
hsa_op_immed::hsa_op_immed (tree tree_val, bool min32int)
: hsa_op_with_type (BRIG_KIND_OPERAND_CONSTANT_BYTES,
hsa_type_for_tree_type (TREE_TYPE (tree_val), NULL,
min32int))
{
if (hsa_seen_error ())
return;
gcc_checking_assert ((is_gimple_min_invariant (tree_val)
&& (!POINTER_TYPE_P (TREE_TYPE (tree_val))
|| TREE_CODE (tree_val) == INTEGER_CST))
|| TREE_CODE (tree_val) == CONSTRUCTOR);
m_tree_value = tree_val;
/* Verify that all elements of a constructor are constants. */
if (TREE_CODE (m_tree_value) == CONSTRUCTOR)
for (unsigned i = 0; i < CONSTRUCTOR_NELTS (m_tree_value); i++)
{
tree v = CONSTRUCTOR_ELT (m_tree_value, i)->value;
if (!CONSTANT_CLASS_P (v))
{
HSA_SORRY_AT (EXPR_LOCATION (tree_val),
"HSA ctor should have only constants");
return;
}
}
}
/* Constructor of class representing HSA immediate values. INTEGER_VALUE is the
integer representation of the immediate value. TYPE is BRIG type. */
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)
{
gcc_assert (hsa_type_integer_p (type));
m_int_value = integer_value;
}
hsa_op_immed::hsa_op_immed ()
: hsa_op_with_type (BRIG_KIND_NONE, BRIG_TYPE_NONE)
{
}
/* New operator to allocate immediate operands from obstack. */
void *
hsa_op_immed::operator new (size_t size)
{
return obstack_alloc (&hsa_obstack, size);
}
/* Destructor. */
hsa_op_immed::~hsa_op_immed ()
{
}
/* Change type of the immediate value to T. */
void
hsa_op_immed::set_type (BrigType16_t t)
{
m_type = t;
}
/* Constructor of class representing HSA registers and pseudo-registers. T is
the BRIG type of the new register. */
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)
{
}
/* New operator to allocate a register from obstack. */
void *
hsa_op_reg::operator new (size_t size)
{
return obstack_alloc (&hsa_obstack, size);
}
/* Verify register operand. */
void
hsa_op_reg::verify_ssa ()
{
/* Verify that each HSA register has a definition assigned.
Exceptions are VAR_DECL and PARM_DECL that are a default
definition. */
gcc_checking_assert (m_def_insn
|| (m_gimple_ssa != NULL
&& (!SSA_NAME_VAR (m_gimple_ssa)
|| (TREE_CODE (SSA_NAME_VAR (m_gimple_ssa))
!= PARM_DECL))
&& SSA_NAME_IS_DEFAULT_DEF (m_gimple_ssa)));
/* Verify that every use of the register is really present
in an instruction. */
for (unsigned i = 0; i < m_uses.length (); i++)
{
hsa_insn_basic *use = m_uses[i];
bool is_visited = false;
for (unsigned j = 0; j < use->operand_count (); j++)
{
hsa_op_base *u = use->get_op (j);
hsa_op_address *addr; addr = dyn_cast <hsa_op_address *> (u);
if (addr && addr->m_reg)
u = addr->m_reg;
if (u == this)
{
bool r = !addr && use->op_output_p (j);
if (r)
{
error ("HSA SSA name defined by instruction that is supposed "
"to be using it");
debug_hsa_operand (this);
debug_hsa_insn (use);
internal_error ("HSA SSA verification failed");
}
is_visited = true;
}
}
if (!is_visited)
{
error ("HSA SSA name not among operands of instruction that is "
"supposed to use it");
debug_hsa_operand (this);
debug_hsa_insn (use);
internal_error ("HSA SSA verification failed");
}
}
}
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)
{
}
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)
{
}
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)
{
}
/* New operator to allocate address operands from obstack. */
void *
hsa_op_address::operator new (size_t size)
{
return obstack_alloc (&hsa_obstack, size);
}
/* Constructor of an operand referring to HSAIL code. */
hsa_op_code_ref::hsa_op_code_ref () : hsa_op_base (BRIG_KIND_OPERAND_CODE_REF),
m_directive_offset (0)
{
}
/* Constructor of an operand representing a code list. Set it up so that it
can contain ELEMENTS number of elements. */
hsa_op_code_list::hsa_op_code_list (unsigned elements)
: hsa_op_base (BRIG_KIND_OPERAND_CODE_LIST)
{
m_offsets.create (1);
m_offsets.safe_grow_cleared (elements);
}
/* New operator to allocate code list operands from obstack. */
void *
hsa_op_code_list::operator new (size_t size)
{
return obstack_alloc (&hsa_obstack, size);
}
/* Constructor of an operand representing an operand list.
Set it up so that it can contain ELEMENTS number of elements. */
hsa_op_operand_list::hsa_op_operand_list (unsigned elements)
: hsa_op_base (BRIG_KIND_OPERAND_OPERAND_LIST)
{
m_offsets.create (elements);
m_offsets.safe_grow (elements);
}
/* New operator to allocate operand list operands from obstack. */
void *
hsa_op_operand_list::operator new (size_t size)
{
return obstack_alloc (&hsa_obstack, size);
}
hsa_op_operand_list::~hsa_op_operand_list ()
{
m_offsets.release ();
}
hsa_op_reg *
hsa_function_representation::reg_for_gimple_ssa (tree ssa)
{
hsa_op_reg *hreg;
gcc_checking_assert (TREE_CODE (ssa) == SSA_NAME);
if (m_ssa_map[SSA_NAME_VERSION (ssa)])
return m_ssa_map[SSA_NAME_VERSION (ssa)];
hreg = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa),
false));
hreg->m_gimple_ssa = ssa;
m_ssa_map[SSA_NAME_VERSION (ssa)] = hreg;
return hreg;
}
void
hsa_op_reg::set_definition (hsa_insn_basic *insn)
{
if (hsa_cfun->m_in_ssa)
{
gcc_checking_assert (!m_def_insn);
m_def_insn = insn;
}
else
m_def_insn = NULL;
}
/* Constructor of the class which is the bases of all instructions and directly
represents the most basic ones. NOPS is the number of operands that the
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)
{
if (nops > 0)
m_operands.safe_grow_cleared (nops);
hsa_instructions.safe_push (this);
}
/* Make OP the operand number INDEX of operands of this instruction. If OP is a
register or an address containing a register, then either set the definition
of the register to this instruction if it an output operand or add this
instruction to the uses if it is an input one. */
void
hsa_insn_basic::set_op (int index, hsa_op_base *op)
{
/* Each address operand is always use. */
hsa_op_address *addr = dyn_cast <hsa_op_address *> (op);
if (addr && addr->m_reg)
addr->m_reg->m_uses.safe_push (this);
else
{
hsa_op_reg *reg = dyn_cast <hsa_op_reg *> (op);
if (reg)
{
if (op_output_p (index))
reg->set_definition (this);
else
reg->m_uses.safe_push (this);
}
}
m_operands[index] = op;
}
/* Get INDEX-th operand of the instruction. */
hsa_op_base *
hsa_insn_basic::get_op (int index)
{
return m_operands[index];
}
/* Get address of INDEX-th operand of the instruction. */
hsa_op_base **
hsa_insn_basic::get_op_addr (int index)
{
return &m_operands[index];
}
/* Get number of operands of the instruction. */
unsigned int
hsa_insn_basic::operand_count ()
{
return m_operands.length ();
}
/* Constructor of the class which is the bases of all instructions and directly
represents the most basic ones. NOPS is the number of operands that the
operand vector will contain (and which will be cleared). OPC is the opcode
of the instruction, T is the type of the instruction. */
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)
{
if (nops > 0)
m_operands.safe_grow_cleared (nops);
if (arg0 != NULL)
{
gcc_checking_assert (nops >= 1);
set_op (0, arg0);
}
if (arg1 != NULL)
{
gcc_checking_assert (nops >= 2);
set_op (1, arg1);
}
if (arg2 != NULL)
{
gcc_checking_assert (nops >= 3);
set_op (2, arg2);
}
if (arg3 != NULL)
{
gcc_checking_assert (nops >= 4);
set_op (3, arg3);
}
hsa_instructions.safe_push (this);
}
/* New operator to allocate basic instruction from obstack. */
void *
hsa_insn_basic::operator new (size_t size)
{
return obstack_alloc (&hsa_obstack, size);
}
/* Verify the instruction. */
void
hsa_insn_basic::verify ()
{
hsa_op_address *addr;
hsa_op_reg *reg;
/* Iterate all register operands and verify that the instruction
is set in uses of the register. */
for (unsigned i = 0; i < operand_count (); i++)
{
hsa_op_base *use = get_op (i);
if ((addr = dyn_cast <hsa_op_address *> (use)) && addr->m_reg)
{
gcc_assert (addr->m_reg->m_def_insn != this);
use = addr->m_reg;
}
if ((reg = dyn_cast <hsa_op_reg *> (use)) && !op_output_p (i))
{
unsigned j;
for (j = 0; j < reg->m_uses.length (); j++)
{
if (reg->m_uses[j] == this)
break;
}
if (j == reg->m_uses.length ())
{
error ("HSA instruction uses a register but is not among "
"recorded register uses");
debug_hsa_operand (reg);
debug_hsa_insn (this);
internal_error ("HSA instruction verification failed");
}
}
}
}
/* Constructor of an instruction representing a PHI node. NOPS is the number
of operands (equal to the number of predecessors). */
hsa_insn_phi::hsa_insn_phi (unsigned nops, hsa_op_reg *dst)
: hsa_insn_basic (nops, HSA_OPCODE_PHI), m_dest (dst)
{
dst->set_definition (this);
}
/* Constructor of class representing instructions for control flow and
sychronization, */
hsa_insn_br::hsa_insn_br (unsigned nops, int opc, BrigType16_t t,
BrigWidth8_t width, 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_width (width)
{
}
/* Constructor of class representing instruction for conditional jump, CTRL is
the control register determining whether the jump will be carried out, the
new instruction is automatically added to its uses list. */
hsa_insn_cbr::hsa_insn_cbr (hsa_op_reg *ctrl)
: hsa_insn_br (1, BRIG_OPCODE_CBR, BRIG_TYPE_B1, BRIG_WIDTH_1, ctrl)
{
}
/* Constructor of class representing instruction for switch jump, CTRL is
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_label_code_list (new hsa_op_code_list (jump_count))
{
}
/* Replace all occurrences of OLD_BB with NEW_BB in the statements
jump table. */
void
hsa_insn_sbr::replace_all_labels (basic_block old_bb, basic_block new_bb)
{
for (unsigned i = 0; i < m_jump_table.length (); i++)
if (m_jump_table[i] == old_bb)
m_jump_table[i] = new_bb;
}
hsa_insn_sbr::~hsa_insn_sbr ()
{
m_jump_table.release ();
}
/* Constructor of comparison instruction. CMP is the comparison operation and T
is the result type. */
hsa_insn_cmp::hsa_insn_cmp (BrigCompareOperation8_t cmp, BrigType16_t t,
hsa_op_base *arg0, hsa_op_base *arg1,
hsa_op_base *arg2)
: hsa_insn_basic (3 , BRIG_OPCODE_CMP, t, arg0, arg1, arg2), m_compare (cmp)
{
}
/* Constructor of classes representing memory accesses. OPC is the opcode (must
be BRIG_OPCODE_ST or BRIG_OPCODE_LD) and T is the type. The instruction
operands are provided as ARG0 and ARG1. */
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)
{
gcc_checking_assert (opc == BRIG_OPCODE_LD || opc == BRIG_OPCODE_ST);
}
/* Constructor for descendants allowing different opcodes and number of
operands, it passes its arguments directly to hsa_insn_basic
constructor. The instruction operands are provided as ARG[0-3]. */
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)
{
}
/* Constructor of class representing atomic instructions. 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]. */
hsa_insn_atomic::hsa_insn_atomic (int nops, int opc,
enum BrigAtomicOperation aop,
BrigType16_t t, BrigMemoryOrder memorder,
hsa_op_base *arg0,
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)
{
gcc_checking_assert (opc == BRIG_OPCODE_ATOMICNORET ||
opc == BRIG_OPCODE_ATOMIC ||
opc == BRIG_OPCODE_SIGNAL ||
opc == BRIG_OPCODE_SIGNALNORET);
}
/* Constructor of class representing signal instructions. OPC is the prinicpal
opcode, SOP is the specific signal operation opcode. T is the type of the
instruction. The instruction operands are provided as ARG[0-3]. */
hsa_insn_signal::hsa_insn_signal (int nops, int opc,
enum BrigAtomicOperation sop,
BrigType16_t t, BrigMemoryOrder memorder,
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_memory_order (memorder), m_signalop (sop)
{
}
/* Constructor of class representing segment conversion instructions. OPC is
the opcode which must be either BRIG_OPCODE_STOF or BRIG_OPCODE_FTOS. DEST
and SRCT are destination and source types respectively, SEG is the segment
we are converting to or from. The instruction operands are
provided as ARG0 and ARG1. */
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)
{
gcc_checking_assert (opc == BRIG_OPCODE_STOF || opc == BRIG_OPCODE_FTOS);
}
/* Constructor of class representing a call instruction. CALLEE is the tree
representation of the function being called. */
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)
{
}
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)
{
}
hsa_insn_call::~hsa_insn_call ()
{
for (unsigned i = 0; i < m_input_args.length (); i++)
delete m_input_args[i];
delete m_output_arg;
m_input_args.release ();
m_input_arg_insns.release ();
}
/* Constructor of class representing the argument block required to invoke
a call in HSAIL. */
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)
{
}
hsa_insn_comment::hsa_insn_comment (const char *s)
: hsa_insn_basic (0, BRIG_KIND_DIRECTIVE_COMMENT)
{
unsigned l = strlen (s);
/* Append '// ' to the string. */
char *buf = XNEWVEC (char, l + 4);
sprintf (buf, "// %s", s);
m_comment = buf;
}
hsa_insn_comment::~hsa_insn_comment ()
{
gcc_checking_assert (m_comment);
free (m_comment);
m_comment = NULL;
}
/* Constructor of class representing the queue instruction in HSAIL. */
hsa_insn_queue::hsa_insn_queue (int nops, int opcode, BrigSegment segment,
BrigMemoryOrder memory_order,
hsa_op_base *arg0, hsa_op_base *arg1,
hsa_op_base *arg2, hsa_op_base *arg3)
: hsa_insn_basic (nops, opcode, BRIG_TYPE_U64, arg0, arg1, arg2, arg3),
m_segment (segment), m_memory_order (memory_order)
{
}
/* Constructor of class representing the source type instruction in HSAIL. */
hsa_insn_srctype::hsa_insn_srctype (int nops, BrigOpcode opcode,
BrigType16_t destt, BrigType16_t srct,
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)
{}
/* Constructor of class representing the packed instruction in HSAIL. */
hsa_insn_packed::hsa_insn_packed (int nops, BrigOpcode opcode,
BrigType16_t destt, BrigType16_t srct,
hsa_op_base *arg0, hsa_op_base *arg1,
hsa_op_base *arg2)
: hsa_insn_srctype (nops, opcode, destt, srct, arg0, arg1, arg2)
{
m_operand_list = new hsa_op_operand_list (nops - 1);
}
/* Constructor of class representing the convert instruction in HSAIL. */
hsa_insn_cvt::hsa_insn_cvt (hsa_op_with_type *dest, hsa_op_with_type *src)
: hsa_insn_basic (2, BRIG_OPCODE_CVT, dest->m_type, dest, src)
{
}
/* Constructor of class representing the alloca in HSAIL. */
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)
{
gcc_assert (dest->m_type == BRIG_TYPE_U32);
if (alignment)
m_align = hsa_alignment_encoding (alignment);
}
/* Append an instruction INSN into the basic block. */
void
hsa_bb::append_insn (hsa_insn_basic *insn)
{
gcc_assert (insn->m_opcode != 0 || insn->operand_count () == 0);
gcc_assert (!insn->m_bb);
insn->m_bb = m_bb;
insn->m_prev = m_last_insn;
insn->m_next = NULL;
if (m_last_insn)
m_last_insn->m_next = insn;
m_last_insn = insn;
if (!m_first_insn)
m_first_insn = insn;
}
void
hsa_bb::append_phi (hsa_insn_phi *hphi)
{
hphi->m_bb = m_bb;
hphi->m_prev = m_last_phi;
hphi->m_next = NULL;
if (m_last_phi)
m_last_phi->m_next = hphi;
m_last_phi = hphi;
if (!m_first_phi)
m_first_phi = hphi;
}
/* Insert HSA instruction NEW_INSN immediately before an existing instruction
OLD_INSN. */
static void
hsa_insert_insn_before (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
{
hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
if (hbb->m_first_insn == old_insn)
hbb->m_first_insn = new_insn;
new_insn->m_prev = old_insn->m_prev;
new_insn->m_next = old_insn;
if (old_insn->m_prev)
old_insn->m_prev->m_next = new_insn;
old_insn->m_prev = new_insn;
}
/* Append HSA instruction NEW_INSN immediately after an existing instruction
OLD_INSN. */
static void
hsa_append_insn_after (hsa_insn_basic *new_insn, hsa_insn_basic *old_insn)
{
hsa_bb *hbb = hsa_bb_for_bb (old_insn->m_bb);
if (hbb->m_last_insn == old_insn)
hbb->m_last_insn = new_insn;
new_insn->m_prev = old_insn;
new_insn->m_next = old_insn->m_next;
if (old_insn->m_next)
old_insn->m_next->m_prev = new_insn;
old_insn->m_next = new_insn;
}
/* Return a register containing the calculated value of EXP which must be an
expression consisting of PLUS_EXPRs, MULT_EXPRs, NOP_EXPRs, SSA_NAMEs and
integer constants as returned by get_inner_reference.
Newly generated HSA instructions will be appended to HBB.
Perform all calculations in ADDRTYPE. */
static hsa_op_with_type *
gen_address_calculation (tree exp, hsa_bb *hbb, BrigType16_t addrtype)
{
int opcode;
if (TREE_CODE (exp) == NOP_EXPR)
exp = TREE_OPERAND (exp, 0);
switch (TREE_CODE (exp))
{
case SSA_NAME:
return hsa_cfun->reg_for_gimple_ssa (exp)->get_in_type (addrtype, hbb);
case INTEGER_CST:
{
hsa_op_immed *imm = new hsa_op_immed (exp);
if (addrtype != imm->m_type)
imm->m_type = addrtype;
return imm;
}
case PLUS_EXPR:
opcode = BRIG_OPCODE_ADD;
break;
case MULT_EXPR:
opcode = BRIG_OPCODE_MUL;
break;
default:
gcc_unreachable ();
}
hsa_op_reg *res = new hsa_op_reg (addrtype);
hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, addrtype);
insn->set_op (0, res);
hsa_op_with_type *op1 = gen_address_calculation (TREE_OPERAND (exp, 0), hbb,
addrtype);
hsa_op_with_type *op2 = gen_address_calculation (TREE_OPERAND (exp, 1), hbb,
addrtype);
insn->set_op (1, op1);
insn->set_op (2, op2);
hbb->append_insn (insn);
return res;
}
/* If R1 is NULL, just return R2, otherwise append an instruction adding them
to HBB and return the register holding the result. */
static hsa_op_reg *
add_addr_regs_if_needed (hsa_op_reg *r1, hsa_op_reg *r2, hsa_bb *hbb)
{
gcc_checking_assert (r2);
if (!r1)
return r2;
hsa_op_reg *res = new hsa_op_reg (r1->m_type);
gcc_assert (!hsa_needs_cvt (r1->m_type, r2->m_type));
hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_ADD, res->m_type);
insn->set_op (0, res);
insn->set_op (1, r1);
insn->set_op (2, r2);
hbb->append_insn (insn);
return res;
}
/* Helper of gen_hsa_addr. Update *SYMBOL, *ADDRTYPE, *REG and *OFFSET to
reflect BASE which is the first operand of a MEM_REF or a TARGET_MEM_REF. */
static void
process_mem_base (tree base, hsa_symbol **symbol, BrigType16_t *addrtype,
hsa_op_reg **reg, offset_int *offset, hsa_bb *hbb)
{
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);
*reg = dyn_cast <hsa_op_reg *> (ssa);
}
else if (TREE_CODE (base) == ADDR_EXPR)
{
tree decl = TREE_OPERAND (base, 0);
if (!DECL_P (decl) || TREE_CODE (decl) == FUNCTION_DECL)
{
HSA_SORRY_AT (EXPR_LOCATION (base),
"support for HSA does not implement a memory reference "
"to a non-declaration type");
return;
}
gcc_assert (!*symbol);
*symbol = get_symbol_for_decl (decl);
*addrtype = hsa_get_segment_addr_type ((*symbol)->m_segment);
}
else if (TREE_CODE (base) == INTEGER_CST)
*offset += wi::to_offset (base);
else
gcc_unreachable ();
}
/* Forward declaration of a function. */
static void
gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb);
/* Generate HSA address operand for a given tree memory reference REF. If
instructions need to be created to calculate the address, they will be added
to the end of HBB. If a caller provider OUTPUT_BITSIZE and OUTPUT_BITPOS,
the function assumes that the caller will handle possible
bit-field references. Otherwise if we reference a bit-field, sorry message
is displayed. */
static hsa_op_address *
gen_hsa_addr (tree ref, hsa_bb *hbb, HOST_WIDE_INT *output_bitsize = NULL,
HOST_WIDE_INT *output_bitpos = NULL)
{
hsa_symbol *symbol = NULL;
hsa_op_reg *reg = NULL;
offset_int offset = 0;
tree origref = ref;
tree varoffset = NULL_TREE;
BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
HOST_WIDE_INT bitsize = 0, bitpos = 0;
BrigType16_t flat_addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
if (TREE_CODE (ref) == STRING_CST)
{
symbol = hsa_get_string_cst_symbol (ref);
goto out;
}
else if (TREE_CODE (ref) == BIT_FIELD_REF
&& (!multiple_p (bit_field_size (ref), BITS_PER_UNIT)
|| !multiple_p (bit_field_offset (ref), BITS_PER_UNIT)))
{
HSA_SORRY_ATV (EXPR_LOCATION (origref),
"support for HSA does not implement "
"bit field references such as %E", ref);
goto out;
}
if (handled_component_p (ref))
{
machine_mode mode;
int unsignedp, volatilep, preversep;
poly_int64 pbitsize, pbitpos;
tree new_ref;
new_ref = get_inner_reference (ref, &pbitsize, &pbitpos, &varoffset,
&mode, &unsignedp, &preversep,
&volatilep);
/* When this isn't true, the switch below will report an
appropriate error. */
if (pbitsize.is_constant () && pbitpos.is_constant ())
{
bitsize = pbitsize.to_constant ();
bitpos = pbitpos.to_constant ();
ref = new_ref;
offset = bitpos;
offset = wi::rshift (offset, LOG2_BITS_PER_UNIT, SIGNED);
}
}
switch (TREE_CODE (ref))
{
case ADDR_EXPR:
{
addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
symbol = hsa_cfun->create_hsa_temporary (flat_addrtype);
hsa_op_reg *r = new hsa_op_reg (flat_addrtype);
gen_hsa_addr_insns (ref, r, hbb);
hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
r, new hsa_op_address (symbol)));
break;
}
case SSA_NAME:
{
addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
hsa_op_with_type *r = hsa_cfun->reg_for_gimple_ssa (ref);
if (r->m_type == BRIG_TYPE_B1)
r = r->get_in_type (BRIG_TYPE_U32, hbb);
symbol = hsa_cfun->create_hsa_temporary (r->m_type);
hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
r, new hsa_op_address (symbol)));
break;
}
case PARM_DECL:
case VAR_DECL:
case RESULT_DECL:
case CONST_DECL:
gcc_assert (!symbol);
symbol = get_symbol_for_decl (ref);
addrtype = hsa_get_segment_addr_type (symbol->m_segment);
break;
case MEM_REF:
process_mem_base (TREE_OPERAND (ref, 0), &symbol, &addrtype, &reg,
&offset, hbb);
if (!integer_zerop (TREE_OPERAND (ref, 1)))
offset += wi::to_offset (TREE_OPERAND (ref, 1));
break;
case TARGET_MEM_REF:
process_mem_base (TMR_BASE (ref), &symbol, &addrtype, &reg, &offset, hbb);
if (TMR_INDEX (ref))
{
hsa_op_reg *disp1;
hsa_op_base *idx = hsa_cfun->reg_for_gimple_ssa
(TMR_INDEX (ref))->get_in_type (addrtype, hbb);
if (TMR_STEP (ref) && !integer_onep (TMR_STEP (ref)))
{
disp1 = new hsa_op_reg (addrtype);
hsa_insn_basic *insn = new hsa_insn_basic (3, BRIG_OPCODE_MUL,
addrtype);
/* As step must respect addrtype, we overwrite the type
of an immediate value. */
hsa_op_immed *step = new hsa_op_immed (TMR_STEP (ref));
step->m_type = addrtype;
insn->set_op (0, disp1);
insn->set_op (1, idx);
insn->set_op (2, step);
hbb->append_insn (insn);
}
else
disp1 = as_a <hsa_op_reg *> (idx);
reg = add_addr_regs_if_needed (reg, disp1, hbb);
}
if (TMR_INDEX2 (ref))
{
if (TREE_CODE (TMR_INDEX2 (ref)) == SSA_NAME)
{
hsa_op_base *disp2 = hsa_cfun->reg_for_gimple_ssa
(TMR_INDEX2 (ref))->get_in_type (addrtype, hbb);
reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (disp2),
hbb);
}
else if (TREE_CODE (TMR_INDEX2 (ref)) == INTEGER_CST)
offset += wi::to_offset (TMR_INDEX2 (ref));
else
gcc_unreachable ();
}
offset += wi::to_offset (TMR_OFFSET (ref));
break;
case FUNCTION_DECL:
HSA_SORRY_AT (EXPR_LOCATION (origref),
"support for HSA does not implement function pointers");
goto out;
default:
HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does "
"not implement memory access to %E", origref);
goto out;
}
if (varoffset)
{
if (TREE_CODE (varoffset) == INTEGER_CST)
offset += wi::to_offset (varoffset);
else
{
hsa_op_base *off_op = gen_address_calculation (varoffset, hbb,
addrtype);
reg = add_addr_regs_if_needed (reg, as_a <hsa_op_reg *> (off_op),
hbb);
}
}
gcc_checking_assert ((symbol
&& addrtype
== hsa_get_segment_addr_type (symbol->m_segment))
|| (!symbol
&& addrtype
== hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT)));
out:
HOST_WIDE_INT hwi_offset = offset.to_shwi ();
/* Calculate remaining bitsize offset (if presented). */
bitpos %= BITS_PER_UNIT;
/* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
is not a reason to think this is a bit-field access. */
if (bitpos == 0
&& (bitsize >= BITS_PER_UNIT)
&& !(bitsize & (bitsize - 1)))
bitsize = 0;
if ((bitpos || bitsize) && (output_bitpos == NULL || output_bitsize == NULL))
HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does not "
"implement unhandled bit field reference such as %E", ref);
if (output_bitsize != NULL && output_bitpos != NULL)
{
*output_bitsize = bitsize;
*output_bitpos = bitpos;
}
return new hsa_op_address (symbol, reg, hwi_offset);
}
/* Generate HSA address operand for a given tree memory reference REF. If
instructions need to be created to calculate the address, they will be added
to the end of HBB. OUTPUT_ALIGN is alignment of the created address. */
static hsa_op_address *
gen_hsa_addr_with_align (tree ref, hsa_bb *hbb, BrigAlignment8_t *output_align)
{
hsa_op_address *addr = gen_hsa_addr (ref, hbb);
if (addr->m_reg || !addr->m_symbol)
*output_align = hsa_object_alignment (ref);
else
{
/* If the address consists only of a symbol and an offset, we
compute the alignment ourselves to take into account any alignment
promotions we might have done for the HSA symbol representation. */
unsigned align = hsa_byte_alignment (addr->m_symbol->m_align);
unsigned misalign = addr->m_imm_offset & (align - 1);
if (misalign)
align = least_bit_hwi (misalign);
*output_align = hsa_alignment_encoding (BITS_PER_UNIT * align);
}
return addr;
}
/* Generate HSA address for a function call argument of given TYPE.
INDEX is used to generate corresponding name of the arguments.
Special value -1 represents fact that result value is created. */
static hsa_op_address *
gen_hsa_addr_for_arg (tree tree_type, int index)
{
hsa_symbol *sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_ARG,
BRIG_LINKAGE_ARG);
sym->m_type = hsa_type_for_tree_type (tree_type, &sym->m_dim);
if (index == -1) /* Function result. */
sym->m_name = "res";
else /* Function call arguments. */
{
sym->m_name = NULL;
sym->m_name_number = index;
}
return new hsa_op_address (sym);
}
/* Generate HSA instructions that process all necessary conversions
of an ADDR to flat addressing and place the result into DEST.
Instructions are appended to HBB. */
static void
convert_addr_to_flat_segment (hsa_op_address *addr, hsa_op_reg *dest,
hsa_bb *hbb)
{
hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_LDA);
insn->set_op (1, addr);
if (addr->m_symbol && addr->m_symbol->m_segment != BRIG_SEGMENT_GLOBAL)
{
/* LDA produces segment-relative address, we need to convert
it to the flat one. */
hsa_op_reg *tmp;
tmp = new hsa_op_reg (hsa_get_segment_addr_type
(addr->m_symbol->m_segment));
hsa_insn_seg *seg;
seg = new hsa_insn_seg (BRIG_OPCODE_STOF,
hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT),
tmp->m_type, addr->m_symbol->m_segment, dest,
tmp);
insn->set_op (0, tmp);
insn->m_type = tmp->m_type;
hbb->append_insn (insn);
hbb->append_insn (seg);
}
else
{
insn->set_op (0, dest);
insn->m_type = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
hbb->append_insn (insn);
}
}
/* Generate HSA instructions that calculate address of VAL including all
necessary conversions to flat addressing and place the result into DEST.
Instructions are appended to HBB. */
static void
gen_hsa_addr_insns (tree val, hsa_op_reg *dest, hsa_bb *hbb)
{
/* Handle cases like tmp = NULL, where we just emit a move instruction
to a register. */
if (TREE_CODE (val) == INTEGER_CST)
{
hsa_op_immed *c = new hsa_op_immed (val);
hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
dest->m_type, dest, c);
hbb->append_insn (insn);
return;
}
hsa_op_address *addr;
gcc_assert (dest->m_type == hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
if (TREE_CODE (val) == ADDR_EXPR)
val = TREE_OPERAND (val, 0);
addr = gen_hsa_addr (val, hbb);
if (TREE_CODE (val) == CONST_DECL
&& is_gimple_reg_type (TREE_TYPE (val)))
{
gcc_assert (addr->m_symbol
&& addr->m_symbol->m_segment == BRIG_SEGMENT_READONLY);
/* CONST_DECLs are in readonly segment which however does not have
addresses convertible to flat segments. So copy it to a private one
and take address of that. */
BrigType16_t csttype
= mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (val),
false));
hsa_op_reg *r = new hsa_op_reg (csttype);
hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_LD, csttype, r,
new hsa_op_address (addr->m_symbol)));
hsa_symbol *copysym = hsa_cfun->create_hsa_temporary (csttype);
hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, csttype, r,
new hsa_op_address (copysym)));
addr->m_symbol = copysym;
}
else if (addr->m_symbol && addr->m_symbol->m_segment == BRIG_SEGMENT_READONLY)
{
HSA_SORRY_ATV (EXPR_LOCATION (val), "support for HSA does "
"not implement taking addresses of complex "
"CONST_DECLs such as %E", val);
return;
}
convert_addr_to_flat_segment (addr, dest, hbb);
}
/* Return an HSA register or HSA immediate value operand corresponding to
gimple operand OP. */
static hsa_op_with_type *
hsa_reg_or_immed_for_gimple_op (tree op, hsa_bb *hbb)
{
hsa_op_reg *tmp;
if (TREE_CODE (op) == SSA_NAME)
tmp = hsa_cfun->reg_for_gimple_ssa (op);
else if (!POINTER_TYPE_P (TREE_TYPE (op)))
return new hsa_op_immed (op);
else
{
tmp = new hsa_op_reg (hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT));
gen_hsa_addr_insns (op, tmp, hbb);
}
return tmp;
}
/* Create a simple movement instruction with register destination DEST and
register or immediate source SRC and append it to the end of HBB. */
void
hsa_build_append_simple_mov (hsa_op_reg *dest, hsa_op_base *src, hsa_bb *hbb)
{
/* Moves of packed data between registers need to adhere to the same type
rules like when dealing with memory. */
BrigType16_t tp = mem_type_for_type (dest->m_type);
hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, tp, dest, src);
hsa_fixup_mov_insn_type (insn);
unsigned dest_size = hsa_type_bit_size (dest->m_type);
if (hsa_op_reg *sreg = dyn_cast <hsa_op_reg *> (src))
gcc_assert (dest_size == hsa_type_bit_size (sreg->m_type));
else
{
unsigned imm_size
= hsa_type_bit_size (as_a <hsa_op_immed *> (src)->m_type);
gcc_assert ((dest_size == imm_size)
/* Eventually < 32bit registers will be promoted to 32bit. */
|| (dest_size < 32 && imm_size == 32));
}
hbb->append_insn (insn);
}
/* Generate HSAIL instructions loading a bit field into register DEST.
VALUE_REG is a register of a SSA name that is used in the bit field
reference. To identify a bit field BITPOS is offset to the loaded memory
and BITSIZE is number of bits of the bit field.
Add instructions to HBB. */
static void
gen_hsa_insns_for_bitfield (hsa_op_reg *dest, hsa_op_reg *value_reg,
HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos,
hsa_bb *hbb)
{
unsigned type_bitsize
= hsa_type_bit_size (hsa_extend_inttype_to_32bit (dest->m_type));
unsigned left_shift = type_bitsize - (bitsize + bitpos);
unsigned right_shift = left_shift + bitpos;
if (left_shift)
{
hsa_op_reg *value_reg_2
= new hsa_op_reg (hsa_extend_inttype_to_32bit (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);
hbb->append_insn (lshift);
value_reg = value_reg_2;
}
if (right_shift)
{
hsa_op_reg *value_reg_2
= new hsa_op_reg (hsa_extend_inttype_to_32bit (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);
hbb->append_insn (rshift);
value_reg = value_reg_2;
}
hsa_insn_basic *assignment
= new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, NULL, value_reg);
hsa_fixup_mov_insn_type (assignment);
hbb->append_insn (assignment);
assignment->set_output_in_type (dest, 0, hbb);
}
/* Generate HSAIL instructions loading a bit field into register DEST. ADDR is
prepared memory address which is used to load the bit field. To identify a
bit field BITPOS is offset to the loaded memory and BITSIZE is number of
bits of the bit field. Add instructions to HBB. Load must be performed in
alignment ALIGN. */
static void
gen_hsa_insns_for_bitfield_load (hsa_op_reg *dest, hsa_op_address *addr,
HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos,
hsa_bb *hbb, BrigAlignment8_t align)
{
hsa_op_reg *value_reg = new hsa_op_reg (dest->m_type);
hsa_insn_mem *mem
= new hsa_insn_mem (BRIG_OPCODE_LD,
hsa_extend_inttype_to_32bit (dest->m_type),
value_reg, addr);
mem->set_align (align);
hbb->append_insn (mem);
gen_hsa_insns_for_bitfield (dest, value_reg, bitsize, bitpos, hbb);
}
/* Return the alignment of base memory accesses we issue to perform bit-field
memory access REF. */
static BrigAlignment8_t
hsa_bitmemref_alignment (tree ref)
{
unsigned HOST_WIDE_INT bit_offset = 0;
while (true)
{
if (TREE_CODE (ref) == BIT_FIELD_REF)
{
if (!tree_fits_uhwi_p (TREE_OPERAND (ref, 2)))
return BRIG_ALIGNMENT_1;
bit_offset += tree_to_uhwi (TREE_OPERAND (ref, 2));
}
else if (TREE_CODE (ref) == COMPONENT_REF
&& DECL_BIT_FIELD (TREE_OPERAND (ref, 1)))
bit_offset += int_bit_position (TREE_OPERAND (ref, 1));
else
break;
ref = TREE_OPERAND (ref, 0);
}
unsigned HOST_WIDE_INT bits = bit_offset % BITS_PER_UNIT;
unsigned HOST_WIDE_INT byte_bits = bit_offset - bits;
BrigAlignment8_t base = hsa_object_alignment (ref);
if (byte_bits == 0)
return base;
return MIN (base, hsa_alignment_encoding (least_bit_hwi (byte_bits)));
}
/* Generate HSAIL instructions loading something into register DEST. RHS is
tree representation of the loaded data, which are loaded as type TYPE. Add
instructions to HBB. */
static void
gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb)
{
/* The destination SSA name will give us the type. */
if (TREE_CODE (rhs) == VIEW_CONVERT_EXPR)
rhs = TREE_OPERAND (rhs, 0);
if (TREE_CODE (rhs) == SSA_NAME)
{
hsa_op_reg *src = hsa_cfun->reg_for_gimple_ssa (rhs);
hsa_build_append_simple_mov (dest, src, hbb);
}
else if (is_gimple_min_invariant (rhs)
|| TREE_CODE (rhs) == ADDR_EXPR)
{
if (POINTER_TYPE_P (TREE_TYPE (rhs)))
{
if (dest->m_type != hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT))
{
HSA_SORRY_ATV (EXPR_LOCATION (rhs),
"support for HSA does not implement conversion "
"of %E to the requested non-pointer type.", rhs);
return;
}
gen_hsa_addr_insns (rhs, dest, hbb);
}
else if (TREE_CODE (rhs) == COMPLEX_CST)
{
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_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);
hbb->append_insn (insn);
}
else
{
hsa_op_immed *imm = new hsa_op_immed (rhs);
hsa_build_append_simple_mov (dest, imm, hbb);
}
}
else if (TREE_CODE (rhs) == REALPART_EXPR || TREE_CODE (rhs) == IMAGPART_EXPR)
{
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));
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 *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),
brig_type, real_reg, imag_reg, packed_reg);
hbb->append_insn (packed);
hsa_op_reg *source = TREE_CODE (rhs) == REALPART_EXPR ?
real_reg : imag_reg;
hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
dest->m_type, NULL, source);
hsa_fixup_mov_insn_type (insn);
hbb->append_insn (insn);
insn->set_output_in_type (dest, 0, hbb);
}
else if (TREE_CODE (rhs) == BIT_FIELD_REF
&& TREE_CODE (TREE_OPERAND (rhs, 0)) == SSA_NAME)
{
tree ssa_name = TREE_OPERAND (rhs, 0);
HOST_WIDE_INT bitsize = tree_to_uhwi (TREE_OPERAND (rhs, 1));
HOST_WIDE_INT bitpos = tree_to_uhwi (TREE_OPERAND (rhs, 2));
hsa_op_reg *imm_value = hsa_cfun->reg_for_gimple_ssa (ssa_name);
gen_hsa_insns_for_bitfield (dest, imm_value, bitsize, bitpos, hbb);
}
else if (DECL_P (rhs) || TREE_CODE (rhs) == MEM_REF
|| TREE_CODE (rhs) == TARGET_MEM_REF
|| handled_component_p (rhs))
{
HOST_WIDE_INT bitsize, bitpos;
/* Load from memory. */
hsa_op_address *addr;
addr = gen_hsa_addr (rhs, hbb, &bitsize, &bitpos);
/* Handle load of a bit field. */
if (bitsize > 64)
{
HSA_SORRY_AT (EXPR_LOCATION (rhs),
"support for HSA does not implement load from a bit "
"field bigger than 64 bits");
return;
}
if (bitsize || bitpos)
gen_hsa_insns_for_bitfield_load (dest, addr, bitsize, bitpos, hbb,
hsa_bitmemref_alignment (rhs));
else
{
BrigType16_t mtype;
/* Not dest->m_type, that's possibly extended. */
mtype = mem_type_for_type (hsa_type_for_scalar_tree_type (type,
false));
hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mtype, dest,
addr);
mem->set_align (hsa_object_alignment (rhs));
hbb->append_insn (mem);
}
}
else
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,
starting at BITPOS with size of BITSIZE. */
static unsigned
get_bitfield_size (unsigned bitpos, unsigned bitsize)
{
unsigned s = bitpos + bitsize;
unsigned sizes[] = {8, 16, 32, 64};
for (unsigned i = 0; i < 4; i++)
if (s <= sizes[i])
return sizes[i];
gcc_unreachable ();
return 0;
}
/* Generate HSAIL instructions storing into memory. LHS is the destination of
the store, SRC is the source operand. Add instructions to HBB. */
static void
gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb)
{
HOST_WIDE_INT bitsize = 0, bitpos = 0;
BrigAlignment8_t req_align;
BrigType16_t mtype;
mtype = mem_type_for_type (hsa_type_for_scalar_tree_type (TREE_TYPE (lhs),
false));
hsa_op_address *addr;
addr = gen_hsa_addr (lhs, hbb, &bitsize, &bitpos);
/* Handle store to a bit field. */
if (bitsize > 64)
{
HSA_SORRY_AT (EXPR_LOCATION (lhs),
"support for HSA does not implement store to a bit field "
"bigger than 64 bits");
return;
}
unsigned type_bitsize = get_bitfield_size (bitpos, bitsize);
/* HSAIL does not support MOV insn with 16-bits integers. */
if (type_bitsize < 32)
type_bitsize = 32;
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)));
for (unsigned i = 0; i < type_bitsize; i++)
if (i < bitpos || i >= bitpos + bitsize)
mask |= ((unsigned HOST_WIDE_INT)1 << i);
hsa_op_reg *value_reg = new hsa_op_reg (mem_type);
req_align = hsa_bitmemref_alignment (lhs);
/* Load value from memory. */
hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, mem_type,
value_reg, addr);
mem->set_align (req_align);
hbb->append_insn (mem);
/* AND the loaded value with prepared mask. */
hsa_op_reg *cleared_reg = new hsa_op_reg (mem_type);
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);
hbb->append_insn (clearing);
/* Shift to left a value that is going to be stored. */
hsa_op_reg *new_value_reg = new hsa_op_reg (mem_type);
hsa_insn_basic *basic = new hsa_insn_basic (2, BRIG_OPCODE_MOV, mem_type,
new_value_reg, src);
hsa_fixup_mov_insn_type (basic);
hbb->append_insn (basic);
if (bitpos)
{
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);
hbb->append_insn (basic);
new_value_reg = shifted_value_reg;
}
/* OR the prepared value with prepared chunk loaded from memory. */
hsa_op_reg *prepared_reg= new hsa_op_reg (mem_type);
basic = new hsa_insn_basic (3, BRIG_OPCODE_OR, mem_type, prepared_reg,
new_value_reg, cleared_reg);
hbb->append_insn (basic);
src = prepared_reg;
mtype = mem_type;
}
else
req_align = hsa_object_alignment (lhs);
hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, mtype, src, addr);
mem->set_align (req_align);
/* 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 (!hsa_type_packed_p (imm->m_type))
imm->m_type = mem->m_type;
else
{
/* ...and all vector immediates apparently need to be vectors of
unsigned bytes. */
unsigned bs = hsa_type_bit_size (imm->m_type);
gcc_assert (bs == hsa_type_bit_size (mem->m_type));
switch (bs)
{
case 32:
imm->m_type = BRIG_TYPE_U8X4;
break;
case 64:
imm->m_type = BRIG_TYPE_U8X8;
break;
case 128:
imm->m_type = BRIG_TYPE_U8X16;
break;
default:
gcc_unreachable ();
}
}
}
hbb->append_insn (mem);
}
/* Generate memory copy instructions that are going to be used
for copying a SRC memory to TARGET memory,
represented by pointer in a register. MIN_ALIGN is minimal alignment
of provided HSA addresses. */
static void
gen_hsa_memory_copy (hsa_bb *hbb, hsa_op_address *target, hsa_op_address *src,
unsigned size, BrigAlignment8_t min_align)
{
hsa_op_address *addr;
hsa_insn_mem *mem;
unsigned offset = 0;
unsigned min_byte_align = hsa_byte_alignment (min_align);
while (size)
{
unsigned s;
if (size >= 8)
s = 8;
else if (size >= 4)
s = 4;
else if (size >= 2)
s = 2;
else
s = 1;
if (s > min_byte_align)
s = min_byte_align;
BrigType16_t t = get_integer_type_by_bytes (s, false);
hsa_op_reg *tmp = new hsa_op_reg (t);
addr = new hsa_op_address (src->m_symbol, src->m_reg,
src->m_imm_offset + offset);
mem = new hsa_insn_mem (BRIG_OPCODE_LD, t, tmp, addr);
hbb->append_insn (mem);
addr = new hsa_op_address (target->m_symbol, target->m_reg,
target->m_imm_offset + offset);
mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, tmp, addr);
hbb->append_insn (mem);
offset += s;
size -= s;
}
}
/* Create a memset mask that is created by copying a CONSTANT byte value
to an integer of BYTE_SIZE bytes. */
static unsigned HOST_WIDE_INT
build_memset_value (unsigned HOST_WIDE_INT constant, unsigned byte_size)
{
if (constant == 0)
return 0;
HOST_WIDE_INT v = constant;
for (unsigned i = 1; i < byte_size; i++)
v |= constant << (8 * i);
return v;
}
/* Generate memory set instructions that are going to be used
for setting a CONSTANT byte value to TARGET memory of SIZE bytes.
MIN_ALIGN is minimal alignment of provided HSA addresses. */
static void
gen_hsa_memory_set (hsa_bb *hbb, hsa_op_address *target,
unsigned HOST_WIDE_INT constant,
unsigned size, BrigAlignment8_t min_align)
{
hsa_op_address *addr;
hsa_insn_mem *mem;
unsigned offset = 0;
unsigned min_byte_align = hsa_byte_alignment (min_align);
while (size)
{
unsigned s;
if (size >= 8)
s = 8;
else if (size >= 4)
s = 4;
else if (size >= 2)
s = 2;
else
s = 1;
if (s > min_byte_align)
s = min_byte_align;
addr = new hsa_op_address (target->m_symbol, target->m_reg,
target->m_imm_offset + offset);
BrigType16_t t = get_integer_type_by_bytes (s, false);
HOST_WIDE_INT c = build_memset_value (constant, s);
mem = new hsa_insn_mem (BRIG_OPCODE_ST, t, new hsa_op_immed (c, t),
addr);
hbb->append_insn (mem);
offset += s;
size -= s;
}
}
/* Generate HSAIL instructions for a single assignment
of an empty constructor to an ADDR_LHS. Constructor is passed as a
tree RHS and all instructions are appended to HBB. ALIGN is
alignment of the address. */
void
gen_hsa_ctor_assignment (hsa_op_address *addr_lhs, tree rhs, hsa_bb *hbb,
BrigAlignment8_t align)
{
if (CONSTRUCTOR_NELTS (rhs))
{
HSA_SORRY_AT (EXPR_LOCATION (rhs),
"support for HSA does not implement load from constructor");
return;
}
unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs)));
gen_hsa_memory_set (hbb, addr_lhs, 0, size, align);
}
/* Generate HSA instructions for a single assignment of RHS to LHS.
HBB is the basic block they will be appended to. */
static void
gen_hsa_insns_for_single_assignment (tree lhs, tree rhs, hsa_bb *hbb)
{
if (TREE_CODE (lhs) == SSA_NAME)
{
hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
if (hsa_seen_error ())
return;
gen_hsa_insns_for_load (dest, rhs, TREE_TYPE (lhs), hbb);
}
else if (TREE_CODE (rhs) == SSA_NAME
|| (is_gimple_min_invariant (rhs) && TREE_CODE (rhs) != STRING_CST))
{
/* Store to memory. */
hsa_op_base *src = hsa_reg_or_immed_for_gimple_op (rhs, hbb);
if (hsa_seen_error ())
return;
gen_hsa_insns_for_store (lhs, src, hbb);
}
else
{
BrigAlignment8_t lhs_align;
hsa_op_address *addr_lhs = gen_hsa_addr_with_align (lhs, hbb,
&lhs_align);
if (TREE_CODE (rhs) == CONSTRUCTOR)
gen_hsa_ctor_assignment (addr_lhs, rhs, hbb, lhs_align);
else
{
BrigAlignment8_t rhs_align;
hsa_op_address *addr_rhs = gen_hsa_addr_with_align (rhs, hbb,
&rhs_align);
unsigned size = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (rhs)));
gen_hsa_memory_copy (hbb, addr_lhs, addr_rhs, size,
MIN (lhs_align, rhs_align));
}
}
}
/* Prepend before INSN a load from spill symbol of SPILL_REG. Return the
register into which we loaded. If this required another register to convert
from a B1 type, return it in *PTMP2, otherwise store NULL into it. We
assume we are out of SSA so the returned register does not have its
definition set. */
hsa_op_reg *
hsa_spill_in (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2)
{
hsa_symbol *spill_sym = spill_reg->m_spill_sym;
hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type);
hsa_op_address *addr = new hsa_op_address (spill_sym);
hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, spill_sym->m_type,
reg, addr);
hsa_insert_insn_before (mem, insn);
*ptmp2 = NULL;
if (spill_reg->m_type == BRIG_TYPE_B1)
{
hsa_insn_basic *cvtinsn;
*ptmp2 = reg;
reg = new hsa_op_reg (spill_reg->m_type);
cvtinsn = new hsa_insn_cvt (reg, *ptmp2);
hsa_insert_insn_before (cvtinsn, insn);
}
return reg;
}
/* Append after INSN a store to spill symbol of SPILL_REG. Return the register
from which we stored. If this required another register to convert to a B1
type, return it in *PTMP2, otherwise store NULL into it. We assume we are
out of SSA so the returned register does not have its use updated. */
hsa_op_reg *
hsa_spill_out (hsa_insn_basic *insn, hsa_op_reg *spill_reg, hsa_op_reg **ptmp2)
{
hsa_symbol *spill_sym = spill_reg->m_spill_sym;
hsa_op_reg *reg = new hsa_op_reg (spill_sym->m_type);
hsa_op_address *addr = new hsa_op_address (spill_sym);
hsa_op_reg *returnreg;
*ptmp2 = NULL;
returnreg = reg;
if (spill_reg->m_type == BRIG_TYPE_B1)
{
hsa_insn_basic *cvtinsn;
*ptmp2 = new hsa_op_reg (spill_sym->m_type);
reg->m_type = spill_reg->m_type;
cvtinsn = new hsa_insn_cvt (*ptmp2, returnreg);
hsa_append_insn_after (cvtinsn, insn);
insn = cvtinsn;
reg = *ptmp2;
}
hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_ST, spill_sym->m_type, reg,
addr);
hsa_append_insn_after (mem, insn);
return returnreg;
}
/* Generate a comparison instruction that will compare LHS and RHS with
comparison specified by CODE and put result into register DEST. DEST has to
have its type set already but must not have its definition set yet.
Generated instructions will be added to HBB. */
static void
gen_hsa_cmp_insn_from_gimple (enum tree_code code, tree lhs, tree rhs,
hsa_op_reg *dest, hsa_bb *hbb)
{
BrigCompareOperation8_t compare;
switch (code)
{
case LT_EXPR:
compare = BRIG_COMPARE_LT;
break;
case LE_EXPR:
compare = BRIG_COMPARE_LE;
break;
case GT_EXPR:
compare = BRIG_COMPARE_GT;
break;
case GE_EXPR:
compare = BRIG_COMPARE_GE;
break;
case EQ_EXPR:
compare = BRIG_COMPARE_EQ;
break;
case NE_EXPR:
compare = BRIG_COMPARE_NE;
break;
case UNORDERED_EXPR:
compare = BRIG_COMPARE_NAN;
break;
case ORDERED_EXPR:
compare = BRIG_COMPARE_NUM;
break;
case UNLT_EXPR:
compare = BRIG_COMPARE_LTU;
break;
case UNLE_EXPR:
compare = BRIG_COMPARE_LEU;
break;
case UNGT_EXPR:
compare = BRIG_COMPARE_GTU;
break;
case UNGE_EXPR:
compare = BRIG_COMPARE_GEU;
break;
case UNEQ_EXPR:
compare = BRIG_COMPARE_EQU;
break;
case LTGT_EXPR:
compare = BRIG_COMPARE_NEU;
break;
default:
HSA_SORRY_ATV (EXPR_LOCATION (lhs),
"support for HSA does not implement comparison tree "
"code %s\n", get_tree_code_name (code));
return;
}
/* CMP instruction returns e.g. 0xffffffff (for a 32-bit with integer)
as a result of comparison. */
BrigType16_t dest_type = hsa_type_integer_p (dest->m_type)
? (BrigType16_t) BRIG_TYPE_B1 : dest->m_type;
hsa_insn_cmp *cmp = new hsa_insn_cmp (compare, dest_type);
hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (lhs, hbb);
cmp->set_op (1, op1->extend_int_to_32bit (hbb));
hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs, hbb);
cmp->set_op (2, op2->extend_int_to_32bit (hbb));
hbb->append_insn (cmp);
cmp->set_output_in_type (dest, 0, hbb);
}
/* Generate an unary instruction with OPCODE and append it to a basic block
HBB. The instruction uses DEST as a destination and OP1
as a single operand. */
static void
gen_hsa_unary_operation (BrigOpcode opcode, hsa_op_reg *dest,
hsa_op_with_type *op1, hsa_bb *hbb)
{
gcc_checking_assert (dest);
hsa_insn_basic *insn;
if (opcode == BRIG_OPCODE_MOV && hsa_needs_cvt (dest->m_type, op1->m_type))
{
insn = new hsa_insn_cvt (dest, op1);
hbb->append_insn (insn);
return;
}
op1 = op1->extend_int_to_32bit (hbb);
if (opcode == BRIG_OPCODE_FIRSTBIT || opcode == BRIG_OPCODE_LASTBIT)
{
BrigType16_t srctype = hsa_type_integer_p (op1->m_type) ? op1->m_type
: hsa_unsigned_type_for_type (op1->m_type);
insn = new hsa_insn_srctype (2, opcode, BRIG_TYPE_U32, srctype, NULL,
op1);
}
else
{
BrigType16_t optype = hsa_extend_inttype_to_32bit (dest->m_type);
insn = new hsa_insn_basic (2, opcode, optype, NULL, op1);
if (opcode == BRIG_OPCODE_MOV)
hsa_fixup_mov_insn_type (insn);
else if (opcode == BRIG_OPCODE_ABS || opcode == BRIG_OPCODE_NEG)
{
/* ABS and NEG only exist in _s form :-/ */
if (insn->m_type == BRIG_TYPE_U32)
insn->m_type = BRIG_TYPE_S32;
else if (insn->m_type == BRIG_TYPE_U64)
insn->m_type = BRIG_TYPE_S64;
}
}
hbb->append_insn (insn);
insn->set_output_in_type (dest, 0, hbb);
}
/* Generate a binary instruction with OPCODE and append it to a basic block
HBB. The instruction uses DEST as a destination and operands OP1
and OP2. */
static void
gen_hsa_binary_operation (int opcode, hsa_op_reg *dest,
hsa_op_with_type *op1, hsa_op_with_type *op2,
hsa_bb *hbb)
{
gcc_checking_assert (dest);
BrigType16_t optype = hsa_extend_inttype_to_32bit (dest->m_type);
op1 = op1->extend_int_to_32bit (hbb);
op2 = op2->extend_int_to_32bit (hbb);
if ((opcode == BRIG_OPCODE_SHL || opcode == BRIG_OPCODE_SHR)
&& is_a <hsa_op_immed *> (op2))
{
hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2);
i->set_type (BRIG_TYPE_U32);
}
if ((opcode == BRIG_OPCODE_OR
|| opcode == BRIG_OPCODE_XOR
|| opcode == BRIG_OPCODE_AND)
&& is_a <hsa_op_immed *> (op2))
{
hsa_op_immed *i = dyn_cast <hsa_op_immed *> (op2);
i->set_type (hsa_unsigned_type_for_type (i->m_type));
}
hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, optype, NULL,
op1, op2);
hbb->append_insn (insn);
insn->set_output_in_type (dest, 0, hbb);
}
/* Generate HSA instructions for a single assignment. HBB is the basic block
they will be appended to. */
static void
gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
{
tree_code code = gimple_assign_rhs_code (assign);
gimple_rhs_class rhs_class = get_gimple_rhs_class (gimple_expr_code (assign));
tree lhs = gimple_assign_lhs (assign);
tree rhs1 = gimple_assign_rhs1 (assign);
tree rhs2 = gimple_assign_rhs2 (assign);
tree rhs3 = gimple_assign_rhs3 (assign);
BrigOpcode opcode;
switch (code)
{
CASE_CONVERT:
case FLOAT_EXPR:
/* The opcode is changed to BRIG_OPCODE_CVT if BRIG types
needs a conversion. */
opcode = BRIG_OPCODE_MOV;
break;
case PLUS_EXPR:
case POINTER_PLUS_EXPR:
opcode = BRIG_OPCODE_ADD;
break;
case MINUS_EXPR:
opcode = BRIG_OPCODE_SUB;
break;
case MULT_EXPR:
opcode = BRIG_OPCODE_MUL;
break;
case MULT_HIGHPART_EXPR: