From de58711a6ddbb1e4558a9454d7aeb6d2b33861de Mon Sep 17 00:00:00 2001
From: marxin <mliska@suse.cz>
Date: Thu, 5 Nov 2015 11:36:23 +0100
Subject: [PATCH] HSA: support global variables
gcc/ChangeLog:
2015-11-05 Martin Liska <mliska@suse.cz>
* hsa-brig.c (emit_directive_variable): Do not display warning
for global variables.
(emit_function_directives): Iterate m_global_symbols instead
of m_readonly_variables.
(hsa_output_global_variables): New function.
(hsa_output_kernel_mapping): Remove.
(hsa_output_libgomp_mapping): New function.
(hsa_output_kernels): Likewise.
(hsa_output_brig): Use new functions.
* hsa-dump.c (dump_hsa_cfun): Dump all global symbols.
* hsa-gen.c (hsa_symbol::global_var_p): New predicate.
(hsa_function_representation::~hsa_function_representation):
Release memory.
(get_symbol_for_decl): Simplify logic to just two types
of variables: local and global.
(hsa_get_string_cst_symbol): Use m_global_symbols instead
of m_readonly_variables.
* hsa.c (hsa_init_compilation_unit_data): Initialize
hsa_global_variable_symbols.
(hsa_deinit_compilation_unit_data): Release it.
* hsa.h (struct hsa_symbol): Remove m_readonly_variables and
replace it with m_global_symbols.
(struct hsa_free_symbol_hasher): Remove.
(hsa_free_symbol_hasher::hash): Likewise.
(hsa_free_symbol_hasher::equal): Likewise.
libgomp/ChangeLog:
2015-11-05 Martin Liska <mliska@suse.cz>
* plugin/plugin-hsa.c (struct global_var_info): New structure.
(struct brig_image_desc): Add global variables.
(create_and_finalize_hsa_program): Define all global variables
used in a BRIG module.
---
gcc/hsa-brig.c | 157 +++++++++++++++++++++++++++++++++++++-------
gcc/hsa-dump.c | 10 +++
gcc/hsa-gen.c | 80 ++++++++++++----------
gcc/hsa.c | 18 ++++-
gcc/hsa.h | 35 +++-------
libgomp/plugin/plugin-hsa.c | 30 +++++++++
6 files changed, 242 insertions(+), 88 deletions(-)
@@ -506,12 +506,7 @@ emit_directive_variable (struct hsa_symbol *symbol)
prefix = '&';
if (!symbol->m_cst_value)
- {
- dirvar.allocation = BRIG_ALLOCATION_PROGRAM;
- if (TREE_CODE (symbol->m_decl) == VAR_DECL)
- warning (0, "referring to global symbol %q+D by name from HSA code "
- "won't work", symbol->m_decl);
- }
+ dirvar.allocation = BRIG_ALLOCATION_PROGRAM;
}
else if (symbol->m_global_scope_p)
prefix = '&';
@@ -545,7 +540,10 @@ emit_directive_variable (struct hsa_symbol *symbol)
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.modifier.allBits |= BRIG_VARIABLE_DEFINITION;
+
+ /* Global variables are just declared and linked via HSA runtime. */
+ if (!symbol->global_var_p ())
+ dirvar.modifier.allBits |= BRIG_VARIABLE_DEFINITION;
dirvar.reserved = 0;
if (symbol->m_cst_value)
@@ -571,7 +569,7 @@ emit_function_directives (hsa_function_representation *f, bool is_declaration)
hsa_symbol *sym;
if (!f->m_declaration_p)
- for (int i = 0; f->m_readonly_variables.iterate (i, &sym); i++)
+ for (int i = 0; f->m_global_symbols.iterate (i, &sym); i++)
{
emit_directive_variable (sym);
brig_insn_count++;
@@ -1832,11 +1830,93 @@ hsa_brig_emit_omp_symbols (void)
static GTY(()) tree hsa_ctor_statements;
static GTY(()) tree hsa_dtor_statements;
-/* Create a static constructor that will register out brig stuff with
- libgomp. */
+/* Create and return __hsa_global_variables symbol that contains
+ all informations consumed by libgomp to link global variables
+ with their string names used by an HSA kernel. */
+
+static tree
+hsa_output_global_variables ()
+{
+ unsigned l = hsa_global_variable_symbols->elements ();
+
+ tree variable_info_type = make_node (RECORD_TYPE);
+ tree id_f1 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
+ get_identifier ("name"), ptr_type_node);
+ DECL_CHAIN (id_f1) = NULL_TREE;
+ tree id_f2 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
+ get_identifier ("omp_data_size"),
+ ptr_type_node);
+ DECL_CHAIN (id_f2) = id_f1;
+ finish_builtin_struct (variable_info_type, "__hsa_variable_info", id_f2,
+ NULL_TREE);
+
+ tree int_num_of_global_vars;
+ int_num_of_global_vars = build_int_cst (uint32_type_node, l);
+ tree global_vars_num_index_type = build_index_type (int_num_of_global_vars);
+ tree global_vars_array_type = build_array_type (variable_info_type,
+ global_vars_num_index_type);
+
+ vec<constructor_elt, va_gc> *global_vars_vec = NULL;
+
+ for (hash_table <hsa_noop_symbol_hasher>::iterator it
+ = hsa_global_variable_symbols->begin ();
+ it != hsa_global_variable_symbols->end (); ++it)
+ {
+ unsigned len = strlen ((*it)->m_name);
+ char *copy = XNEWVEC (char, len + 2);
+ copy[0] = '&';
+ memcpy (copy + 1, (*it)->m_name, len);
+ copy[len + 1] = '\0';
+ len++;
+ 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)));
+ free (copy);
+
+ vec<constructor_elt, va_gc> *variable_info_vec = NULL;
+ CONSTRUCTOR_APPEND_ELT (variable_info_vec, NULL_TREE,
+ build1 (ADDR_EXPR,
+ build_pointer_type (TREE_TYPE (var_name)),
+ var_name));
+ CONSTRUCTOR_APPEND_ELT (variable_info_vec, NULL_TREE,
+ build_fold_addr_expr ((*it)->m_decl));
+
+ tree variable_info_ctor = build_constructor (variable_info_type,
+ variable_info_vec);
+
+ CONSTRUCTOR_APPEND_ELT (global_vars_vec, NULL_TREE,
+ variable_info_ctor);
+ }
+
+ tree global_vars_ctor = build_constructor (global_vars_array_type,
+ global_vars_vec);
+
+ char tmp_name[64];
+ ASM_GENERATE_INTERNAL_LABEL (tmp_name, "__hsa_global_variables", 1);
+ tree global_vars_table = build_decl (UNKNOWN_LOCATION, VAR_DECL,
+ get_identifier (tmp_name),
+ global_vars_array_type);
+ TREE_STATIC (global_vars_table) = 1;
+ TREE_READONLY (global_vars_table) = 1;
+ TREE_PUBLIC (global_vars_table) = 0;
+ DECL_ARTIFICIAL (global_vars_table) = 1;
+ DECL_IGNORED_P (global_vars_table) = 1;
+ DECL_EXTERNAL (global_vars_table) = 0;
+ TREE_CONSTANT (global_vars_table) = 1;
+ DECL_INITIAL (global_vars_table) = global_vars_ctor;
+ varpool_node::finalize_decl (global_vars_table);
+
+ return global_vars_table;
+}
+
+/* Create __hsa_host_functions and __hsa_kernels that contain
+ all informations consumed by libgomp to register all kernels
+ in the BRIG binary. */
static void
-hsa_output_kernel_mapping (tree brig_decl)
+hsa_output_kernels (tree *host_func_table, tree *kernels)
{
unsigned map_count = hsa_get_number_decl_kernel_mappings ();
@@ -1870,6 +1950,7 @@ hsa_output_kernel_mapping (tree brig_decl)
TREE_CONSTANT (hsa_host_func_table) = 1;
DECL_INITIAL (hsa_host_func_table) = host_functions_ctor;
varpool_node::finalize_decl (hsa_host_func_table);
+ *host_func_table = hsa_host_func_table;
/* Following code emits list of kernel_info structures. */
@@ -2015,36 +2096,68 @@ hsa_output_kernel_mapping (tree brig_decl)
DECL_INITIAL (hsa_kernels) = build_constructor (kernel_info_vector_type,
kernel_info_vector_vec);
varpool_node::finalize_decl (hsa_kernels);
+ *kernels = hsa_kernels;
+}
+
+/* Create a static constructor that will register out brig stuff with
+ libgomp. */
+
+static void
+hsa_output_libgomp_mapping (tree brig_decl)
+{
+ unsigned kernel_count = hsa_get_number_decl_kernel_mappings ();
+ unsigned global_variable_count = hsa_global_variable_symbols->elements ();
+
+ tree kernels;
+ tree host_func_table;
+
+ hsa_output_kernels (&host_func_table, &kernels);
+ tree global_vars = hsa_output_global_variables ();
tree hsa_image_desc_type = make_node (RECORD_TYPE);
- id_f1 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
+ tree id_f1 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
get_identifier ("brig_module"), ptr_type_node);
DECL_CHAIN (id_f1) = NULL_TREE;
- id_f2 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
+ tree id_f2 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
get_identifier ("kernel_count"),
unsigned_type_node);
DECL_CHAIN (id_f2) = id_f1;
- id_f3 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
+ tree id_f3 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
get_identifier ("hsa_kernel_infos"),
ptr_type_node);
DECL_CHAIN (id_f3) = id_f2;
- finish_builtin_struct (hsa_image_desc_type, "__hsa_image_desc", id_f3,
+ tree id_f4 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
+ get_identifier ("global_variable_count"),
+ unsigned_type_node);
+ DECL_CHAIN (id_f4) = id_f3;
+ tree id_f5 = build_decl (BUILTINS_LOCATION, FIELD_DECL,
+ get_identifier ("hsa_global_variable_infos"),
+ ptr_type_node);
+ DECL_CHAIN (id_f5) = id_f4;
+ finish_builtin_struct (hsa_image_desc_type, "__hsa_image_desc", id_f5,
NULL_TREE);
vec<constructor_elt, va_gc> *img_desc_vec = NULL;
CONSTRUCTOR_APPEND_ELT (img_desc_vec, NULL_TREE,
build_fold_addr_expr (brig_decl));
CONSTRUCTOR_APPEND_ELT (img_desc_vec, NULL_TREE,
- build_int_cstu (unsigned_type_node, map_count));
+ build_int_cstu (unsigned_type_node, kernel_count));
CONSTRUCTOR_APPEND_ELT (img_desc_vec, NULL_TREE,
build1 (ADDR_EXPR,
- build_pointer_type (TREE_TYPE
- (hsa_kernels)),
- hsa_kernels));
+ build_pointer_type (TREE_TYPE (kernels)),
+ kernels));
+ CONSTRUCTOR_APPEND_ELT (img_desc_vec, NULL_TREE,
+ build_int_cstu (unsigned_type_node,
+ global_variable_count));
+ CONSTRUCTOR_APPEND_ELT (img_desc_vec, NULL_TREE,
+ build1 (ADDR_EXPR,
+ build_pointer_type (TREE_TYPE (global_vars)),
+ global_vars));
tree img_desc_ctor = build_constructor (hsa_image_desc_type, img_desc_vec);
+ char tmp_name[64];
ASM_GENERATE_INTERNAL_LABEL (tmp_name, "__hsa_img_descriptor", 1);
tree hsa_img_descriptor = build_decl (UNKNOWN_LOCATION, VAR_DECL,
get_identifier (tmp_name),
@@ -2065,11 +2178,11 @@ hsa_output_kernel_mapping (tree brig_decl)
(build_int_cst
(integer_type_node, 4)));
vec<constructor_elt, va_gc> *libgomp_host_table_vec = NULL;
- tree host_func_table_addr = build_fold_addr_expr (hsa_host_func_table);
+ 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))
- * map_count;
+ * kernel_count;
CONSTRUCTOR_APPEND_ELT (libgomp_host_table_vec, NULL_TREE,
fold_build2 (POINTER_PLUS_EXPR,
TREE_TYPE (host_func_table_addr),
@@ -2234,7 +2347,7 @@ hsa_output_brig (void)
if (saved_section)
switch_to_section (saved_section);
- hsa_output_kernel_mapping (brig_decl);
+ hsa_output_libgomp_mapping (brig_decl);
hsa_free_decl_kernel_mapping ();
brig_release_data ();
@@ -1079,6 +1079,16 @@ dump_hsa_cfun (FILE *f)
{
basic_block bb;
+ if (hsa_cfun->m_global_symbols.length () > 0)
+ fprintf (f, "\nHSAIL in global scope\n");
+
+ for (unsigned i = 0; i < hsa_cfun->m_global_symbols.length (); i++)
+ {
+ fprintf (f, " ");
+ dump_hsa_symbol (f, hsa_cfun->m_global_symbols[i]);
+ fprintf (f, "\n");
+ }
+
fprintf (f, "\nHSAIL IL for %s\n", hsa_cfun->m_name);
for (unsigned i = 0; i < hsa_cfun->m_private_variables.length (); i++)
@@ -199,6 +199,12 @@ hsa_symbol::fillup_for_decl (tree decl)
m_seen_error = true;
}
+bool
+hsa_symbol::global_var_p ()
+{
+ return m_decl && is_global_var (m_decl);
+}
+
/* 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
@@ -207,7 +213,7 @@ hsa_symbol::fillup_for_decl (tree decl)
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_readonly_variables (vNULL),
+ m_output_arg (NULL), m_spill_symbols (vNULL), m_global_symbols (vNULL),
m_private_variables (vNULL), m_called_functions (vNULL), m_hbb_count (0),
m_in_ssa (true), m_kern_p (kernel_p), m_declaration_p (false), m_decl (fdecl),
m_shadow_reg (NULL), m_kernel_dispatch_count (0), m_maximum_omp_data_size (0),
@@ -238,9 +244,11 @@ hsa_function_representation::~hsa_function_representation ()
delete m_spill_symbols[i];
m_spill_symbols.release ();
- for (unsigned i = 0; i < m_readonly_variables.length (); i++)
- delete m_readonly_variables[i];
- m_readonly_variables.release ();
+ hsa_symbol *sym;
+ for (unsigned i = 0; i < m_global_symbols.iterate (i, &sym); i++)
+ if (!sym->global_var_p ())
+ delete sym;
+ m_global_symbols.release ();
for (unsigned i = 0; i < m_private_variables.length (); i++)
delete m_private_variables[i];
@@ -684,7 +692,7 @@ hsa_needs_cvt (BrigType16_t dtype, BrigType16_t stype)
static hsa_symbol *
get_symbol_for_decl (tree decl)
{
- hsa_symbol **slot, *sym;
+ hsa_symbol **slot;
hsa_symbol dummy (BRIG_TYPE_NONE, BRIG_SEGMENT_NONE, BRIG_LINKAGE_NONE);
gcc_assert (TREE_CODE (decl) == PARM_DECL
@@ -693,50 +701,50 @@ get_symbol_for_decl (tree decl)
dummy.m_decl = decl;
- slot = hsa_cfun->m_local_symbols->find_slot (&dummy, INSERT);
+ 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);
+ else
+ slot = hsa_cfun->m_local_symbols->find_slot (&dummy, INSERT);
+
gcc_checking_assert (slot);
if (*slot)
{
- sym = *slot;
-
/* If the symbol is problematic, mark current function also as
problematic. */
- if (sym->m_seen_error)
+ if ((*slot)->m_seen_error)
hsa_fail_cfun ();
- return sym;
+ return *slot;
}
-
- if (TREE_CODE (decl) == VAR_DECL && is_global_var (decl))
+ else
{
- sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_READONLY,
- BRIG_LINKAGE_MODULE);
+ hsa_symbol *sym;
+ gcc_assert (TREE_CODE (decl) == VAR_DECL);
- /* Following type of global variables can be handled. */
- if (TREE_READONLY (decl) && !TREE_ADDRESSABLE (decl)
- && DECL_INITIAL (decl) && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
- && TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) == INTEGER_TYPE)
+ if (is_in_global_vars)
{
- sym->m_cst_value = new hsa_op_immed (DECL_INITIAL (decl), false);
+ sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_GLOBAL,
+ BRIG_LINKAGE_PROGRAM);
+ hsa_cfun->m_global_symbols.safe_push (sym);
}
else
- HSA_SORRY_ATV (EXPR_LOCATION (decl), "referring to global symbol "
- "%q+D by name from HSA code won't work", decl);
+ {
+ /* PARM_DECL and RESULT_DECL should be already in m_local_symbols. */
+ gcc_assert (TREE_CODE (decl) == VAR_DECL);
- hsa_cfun->m_readonly_variables.safe_push (sym);
- }
- else
- {
- gcc_assert (TREE_CODE (decl) == VAR_DECL);
- sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_PRIVATE,
- BRIG_LINKAGE_FUNCTION);
- hsa_cfun->m_private_variables.safe_push (sym);
- }
+ sym = new hsa_symbol (BRIG_TYPE_NONE, BRIG_SEGMENT_PRIVATE,
+ BRIG_LINKAGE_FUNCTION);
+ hsa_cfun->m_private_variables.safe_push (sym);
+ }
- sym->fillup_for_decl (decl);
- sym->m_name = hsa_get_declaration_name (decl);
- *slot = sym;
- return sym;
+ sym->fillup_for_decl (decl);
+ sym->m_name = hsa_get_declaration_name (decl);
+
+ *slot = sym;
+ return sym;
+ }
}
/* For a given HSA function declaration, return a host
@@ -799,10 +807,10 @@ hsa_get_string_cst_symbol (tree string_cst)
BRIG_SEGMENT_GLOBAL, BRIG_LINKAGE_MODULE);
sym->m_cst_value = cst;
sym->m_dim = TREE_STRING_LENGTH (string_cst);
- sym->m_name_number = hsa_cfun->m_readonly_variables.length ();
+ sym->m_name_number = hsa_cfun->m_global_symbols.length ();
sym->m_global_scope_p = true;
- hsa_cfun->m_readonly_variables.safe_push (sym);
+ hsa_cfun->m_global_symbols.safe_push (sym);
hsa_cfun->m_string_constants_map.put (string_cst, sym);
return sym;
}
@@ -63,7 +63,7 @@ static GTY (()) vec<hsa_decl_kernel_map_element, va_gc> *hsa_decl_kernel_mapping
hash_map <tree, vec <const char *> *> *hsa_decl_kernel_dependencies;
/* Hash function to lookup a symbol for a decl. */
-hash_table <hsa_free_symbol_hasher> *hsa_global_variable_symbols;
+hash_table <hsa_noop_symbol_hasher> *hsa_global_variable_symbols;
/* HSA summaries. */
hsa_summary_t *hsa_summaries = NULL;
@@ -96,6 +96,7 @@ hsa_init_compilation_unit_data (void)
compilation_unit_data_initialized = true;
+ hsa_global_variable_symbols = new hash_table <hsa_noop_symbol_hasher> (8);
hsa_failed_functions = new hash_set <tree> ();
}
@@ -105,8 +106,19 @@ hsa_init_compilation_unit_data (void)
void
hsa_deinit_compilation_unit_data (void)
{
- if (hsa_failed_functions)
- delete hsa_failed_functions;
+ gcc_assert (compilation_unit_data_initialized);
+
+ delete hsa_failed_functions;
+
+ 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;
+ }
+
+ delete hsa_global_variable_symbols;
if (hsa_num_threads)
{
@@ -69,6 +69,10 @@ struct hsa_symbol
or a variable, local or global. */
void fillup_for_decl (tree decl);
+ /* Return true if the symbol is a global variable that should be preserved
+ after a function is emitted to BRIG. */
+ bool global_var_p ();
+
/* Pointer to the original tree, which is PARM_DECL for input parameters and
RESULT_DECL for the output parameters. */
tree m_decl;
@@ -1005,30 +1009,6 @@ hsa_noop_symbol_hasher::equal (const value_type a, const compare_type b)
return (DECL_UID (a->m_decl) == DECL_UID (b->m_decl));
}
-/* Class for hashing global hsa_symbols. */
-
-struct hsa_free_symbol_hasher : free_ptr_hash <hsa_symbol>
-{
- static inline hashval_t hash (const value_type);
- static inline bool equal (const value_type, const compare_type);
-};
-
-/* Hash hsa_symbol. */
-
-inline hashval_t
-hsa_free_symbol_hasher::hash (const value_type item)
-{
- return DECL_UID (item->m_decl);
-}
-
-/* Return true if the DECL_UIDs of decls both symbols refer to are equal. */
-
-inline bool
-hsa_free_symbol_hasher::equal (const value_type a, const compare_type b)
-{
- return (DECL_UID (a->m_decl) == DECL_UID (b->m_decl));
-}
-
/* Structure that encapsulates intermediate representation of a HSA
function. */
@@ -1078,9 +1058,9 @@ public:
/* Vector of pointers to spill symbols. */
vec <struct hsa_symbol *> m_spill_symbols;
- /* Vector of pointers to symbols (string constants and global,
- non-addressable variables with a constructor). */
- vec <struct hsa_symbol *> m_readonly_variables;
+ /* Vector of pointers to global variables and transformed string constants
+ that are used by the function. */
+ vec <struct hsa_symbol *> m_global_symbols;
/* Private function artificial variables. */
vec <struct hsa_symbol *> m_private_variables;
@@ -1189,6 +1169,7 @@ extern hsa_summary_t *hsa_summaries;
extern hsa_symbol *hsa_num_threads;
extern unsigned hsa_kernel_calls_counter;
extern hash_set <tree> *hsa_failed_functions;
+extern hash_table <hsa_noop_symbol_hasher> *hsa_global_variable_symbols;
bool hsa_callable_function_p (tree fndecl);
void hsa_init_compilation_unit_data (void);
@@ -161,6 +161,12 @@ struct hsa_kernel_description
const char **kernel_dependencies;
};
+struct global_var_info
+{
+ const char *name;
+ void *address;
+};
+
/* Data passed by the static initializer of a compilation unit containing BRIG
to GOMP_offload_register. */
@@ -169,6 +175,8 @@ struct brig_image_desc
hsa_ext_module_t brig_module;
const unsigned kernel_count;
struct hsa_kernel_description *kernel_infos;
+ const unsigned global_variable_count;
+ struct global_var_info *global_variables;
};
struct agent_info;
@@ -750,6 +758,28 @@ create_and_finalize_hsa_program (struct agent_info *agent)
if (status != HSA_STATUS_SUCCESS)
hsa_fatal ("Could not create HSA executable", status);
+ module = agent->first_module;
+ while (module)
+ {
+ /* Initialize all global variables declared in the module. */
+ for (unsigned i = 0; i < module->image_desc->global_variable_count; i++)
+ {
+ struct global_var_info *var;
+ var = &module->image_desc->global_variables[i];
+ status = hsa_executable_global_variable_define
+ (agent->executable, var->name, var->address);
+
+ HSA_DEBUG ("Defining global variable: %s, address: %p\n", var->name,
+ var->address);
+
+ if (status != HSA_STATUS_SUCCESS)
+ hsa_fatal ("Could not define a global variable in the HSA program",
+ status);
+ }
+
+ module = module->next;
+ }
+
status = hsa_executable_load_code_object(agent->executable, agent->id,
code_object, "");
if (status != HSA_STATUS_SUCCESS)
--
2.6.2