public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [HSA] support global variables
@ 2015-11-11 15:21 Martin Liška
  0 siblings, 0 replies; only message in thread
From: Martin Liška @ 2015-11-11 15:21 UTC (permalink / raw)
  To: GCC Patches

[-- Attachment #1: Type: text/plain, Size: 214 bytes --]

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

[-- Attachment #2: 0001-HSA-support-global-variables.patch --]
[-- Type: text/x-patch, Size: 22215 bytes --]

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(-)

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


^ permalink raw reply	[flat|nested] only message in thread

only message in thread, other threads:[~2015-11-11 15:21 UTC | newest]

Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-11-11 15:21 [HSA] support global variables Martin Liška

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).