Hi. Following patch adds support for global variables seen by an HSAIL executable. The HSA runtime can link a name of a global variable with pointer to the variable used by host.
Installed to HSA branch. Martin
>From de58711a6ddbb1e4558a9454d7aeb6d2b33861de Mon Sep 17 00:00:00 2001 From: marxin <mli...@suse.cz> Date: Thu, 5 Nov 2015 11:36:23 +0100 Subject: [PATCH] HSA: support global variables gcc/ChangeLog: 2015-11-05 Martin Liska <mli...@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 <mli...@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(-) diff --git a/gcc/hsa-brig.c b/gcc/hsa-brig.c index d2882fc..f47e9c3 100644 --- a/gcc/hsa-brig.c +++ b/gcc/hsa-brig.c @@ -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 (); diff --git a/gcc/hsa-dump.c b/gcc/hsa-dump.c index 4007eff..1391f7b 100644 --- a/gcc/hsa-dump.c +++ b/gcc/hsa-dump.c @@ -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++) diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c index a00fc10..300bee6 100644 --- a/gcc/hsa-gen.c +++ b/gcc/hsa-gen.c @@ -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; } diff --git a/gcc/hsa.c b/gcc/hsa.c index ab05a1d..4add232 100644 --- a/gcc/hsa.c +++ b/gcc/hsa.c @@ -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) { diff --git a/gcc/hsa.h b/gcc/hsa.h index 025de67..d1f6ee0 100644 --- a/gcc/hsa.h +++ b/gcc/hsa.h @@ -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); diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c index 470b892..6558413 100644 --- a/libgomp/plugin/plugin-hsa.c +++ b/libgomp/plugin/plugin-hsa.c @@ -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