Re: [hsa 3/10] HSA libgomp plugin [part 1/2]

["Martin Liška" <mliska@suse.cz>]

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

On 12/09/2015 01:15 PM, Jakub Jelinek wrote:
> On Mon, Dec 07, 2015 at 12:20:49PM +0100, Martin Jambor wrote:
>> +#include <stdio.h>
>> +#include <stdlib.h>
>> +#include <string.h>
>> +#include <pthread.h>
>> +#include "libgomp-plugin.h"
>> +#include "gomp-constants.h"
>> +#include "hsa.h"
>> +#include "hsa_ext_finalize.h"
> 
> If these 2 headers are coming from outside of gcc, better use <> for them
> instead of "", and better place them above the libgomp specific ones.
> 
>> +#include "dlfcn.h"
> 
> Ditto.
> 
>> +/* Flag to decide whether print to stderr information about what is going on.
>> +   Set in init_debug depending on environment variables.  */
>> +
>> +static bool debug;
>> +
>> +/* Flag to decide if the runtime should suppress a possible fallback to host
>> +   execution.  */
>> +
>> +static bool suppress_host_fallback;
>> +
>> +/* Initialize debug and suppress_host_fallback according to the environment.  */
>> +
>> +static void
>> +init_enviroment_variables (void)
>> +{
>> +  if (getenv ("HSA_DEBUG"))
>> +    debug = true;
>> +  else
>> +    debug = false;
>> +
>> +  if (getenv ("HSA_SUPPRESS_HOST_FALLBACK"))
>> +    suppress_host_fallback = true;
>> +  else
>> +    suppress_host_fallback = false;
> 
> Wonder whether we want these custom env vars in suid programs, allowing
> users to change behavior might be undesirable.  So perhaps use secure_getenv
> instead of getenv if found by configure?
>> +
>> +  const char* hsa_error;
> 
> Space before * instead of after it (multiple spots).
> 
>> +  hsa_status_string (status, &hsa_error);
>> +
>> +  unsigned l = strlen (hsa_error);
>> +
>> +  char *err = GOMP_PLUGIN_malloc (sizeof (char) * l);
> 
> sizeof (char) is 1, please don't multiply by it, that is just confusing.
> It makes sense in macros where you don't know what the type really is, but
> not here.
> 
>> +  memcpy (err, hsa_error, l - 1);
>> +  err[l] = '\0';
>> +
>> +  fprintf (stderr, "HSA warning: %s (%s)\n", str, err);
> 
> Why do you allocate err at all, if you free it right away?  Can't you use
> hsa_error directly instead?
>> +
>> +  free (err);
> 
>> +static void
>> +hsa_fatal (const char *str, hsa_status_t status)
>> +{
>> +  const char* hsa_error;
>> +  hsa_status_string (status, &hsa_error);
>> +  GOMP_PLUGIN_fatal ("HSA fatal error: %s (%s)", str, hsa_error);
> 
> Like you do e.g. here?
> 
>> +/* Callback of dispatch queues to report errors.  */
>> +
>> +static void
>> +queue_callback(hsa_status_t status, hsa_queue_t* queue __attribute__ ((unused)),
> 
> Missing space before (.
>> +/* Callback of hsa_agent_iterate_regions.  Determine if a memory REGION can be
>> +   used for kernarg allocations and if so write it to the memory pointed to by
>> +   DATA and break the query.  */
>> +
>> +static hsa_status_t get_kernarg_memory_region (hsa_region_t region, void* data)
> 
> Missing newline between return type and function name.
> 
>> +      hsa_region_t* ret = (hsa_region_t*) data;
> 
> Two spots with wrong formatting above.
>> +{
>> +  if (agent->first_module)
>> +      agent->first_module->prev = module;
> 
> Wrong indentation.
> 
>> +  unsigned size = end - start;
>> +  char *buf = (char *) malloc (size);
>> +  memcpy (buf, start, size);
> 
> malloc could return NULL and you crash.  Should this use GOMP_PLUGIN_malloc
> instead?
> 
>> +  struct GOMP_hsa_kernel_dispatch *shadow = GOMP_PLUGIN_malloc_cleared
>> +    (sizeof (struct GOMP_hsa_kernel_dispatch));
> 
> Formatting, = should go on the next line, and if even that doesn't help,
> maybe use sizeof (*shadow)?
>> +
>> +  shadow->queue = agent->command_q;
>> +  shadow->omp_data_memory = omp_data_size > 0
>> +    ? GOMP_PLUGIN_malloc (omp_data_size) : NULL;
> 
> = should go on the next line.
> 
>> +  unsigned dispatch_count = kernel->dependencies_count;
>> +  shadow->kernel_dispatch_count = dispatch_count;
>> +
>> +  shadow->children_dispatches = GOMP_PLUGIN_malloc
>> +    (dispatch_count * sizeof (struct GOMP_hsa_kernel_dispatch *));
> 
> Likewise.
>> +indent_stream (FILE *f, unsigned indent)
>> +{
>> +  for (int i = 0; i < indent; i++)
>> +    fputc (' ', f);
> 
> Perhaps fprintf (f, "%*s", indent, "");
> instead?
> 
>> +  struct GOMP_hsa_kernel_dispatch *shadow = create_single_kernel_dispatch
>> +    (kernel, omp_data_size);
> 
> Formatting issues again.
> 
>> +  shadow->omp_num_threads = 64;
>> +  shadow->debug = 0;
>> +  shadow->omp_level = kernel->gridified_kernel_p ? 1 : 0;
>> +
>> +  /* Create kernel dispatch data structures.  We do not allow to have
>> +     a kernel dispatch with depth bigger than one.  */
>> +  for (unsigned i = 0; i < kernel->dependencies_count; i++)
>> +    {
>> +      struct kernel_info *dependency = get_kernel_for_agent
>> +	(kernel->agent, kernel->dependencies[i]);
>> +      shadow->children_dispatches[i] = create_single_kernel_dispatch
>> +	(dependency, omp_data_size);
>> +      shadow->children_dispatches[i]->queue =
> 
> Never put = at the end of line.
>> +	kernel->agent->kernel_dispatch_command_q;
> 
> 	Jakub
> 

Hello.

There's a new version of the patch that changes many GNU coding style issues,
trimming of last character in hsa_warn is removed.

Martin

[-- Attachment #2: 0001-Clean-up-plugin-hsa.c-file.patch --]
[-- Type: text/x-patch, Size: 15289 bytes --]

From a289f04ce17d8f3a0c30ce87b029854736f82a16 Mon Sep 17 00:00:00 2001
From: marxin <mliska@suse.cz>
Date: Wed, 9 Dec 2015 14:06:55 +0100
Subject: [PATCH 1/2] Clean-up plugin-hsa.c file

libgomp/ChangeLog:

2015-12-09  Martin Liska  <mliska@suse.cz>

	* plugin/plugin-hsa.c (hsa_warn): Change output string format.
	(hsa_fatal): Dtto.
	(struct agent_info): Fix GNU coding style.
	(suitable_hsa_agent_p): Likewise.
	(get_kernarg_memory_region): Likewise.
	(add_module_to_agent): Likewise.
	(add_shared_library): Use GOMP_PLUGIN_malloc instead
	of malloc.
	(create_and_finalize_hsa_program): Fix GNU coding style.
	(create_single_kernel_dispatch): Dtto.
	(release_kernel_dispatch): Dtto.
	(init_single_kernel): Dtto.
	(indent_stream): Use printf.
	(print_kernel_dispatch): Fix GNU coding style.
	(create_kernel_dispatch): Dtto.
	(GOMP_OFFLOAD_run): Dtto.
---
 libgomp/plugin/plugin-hsa.c | 144 ++++++++++++++++++++++----------------------
 1 file changed, 71 insertions(+), 73 deletions(-)

diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c
index b132954..46b9c03 100644
--- a/libgomp/plugin/plugin-hsa.c
+++ b/libgomp/plugin/plugin-hsa.c
@@ -31,11 +31,11 @@
 #include <stdlib.h>
 #include <string.h>
 #include <pthread.h>
+#include <hsa.h>
+#include <hsa_ext_finalize.h>
+#include <dlfcn.h>
 #include "libgomp-plugin.h"
 #include "gomp-constants.h"
-#include "hsa.h"
-#include "hsa_ext_finalize.h"
-#include "dlfcn.h"
 
 /* Part of the libgomp plugin interface.  Return the name of the accelerator,
    which is "hsa".  */
@@ -128,18 +128,10 @@ hsa_warn (const char *str, hsa_status_t status)
   if (!debug)
     return;
 
-  const char* hsa_error;
+  const char *hsa_error;
   hsa_status_string (status, &hsa_error);
 
-  unsigned l = strlen (hsa_error);
-
-  char *err = GOMP_PLUGIN_malloc (sizeof (char) * l);
-  memcpy (err, hsa_error, l - 1);
-  err[l] = '\0';
-
-  fprintf (stderr, "HSA warning: %s (%s)\n", str, err);
-
-  free (err);
+  fprintf (stderr, "HSA warning: %s\nRuntime message: %s", str, hsa_error);
 }
 
 /* Report a fatal error STR together with the HSA error corresponding to STATUS
@@ -148,9 +140,10 @@ hsa_warn (const char *str, hsa_status_t status)
 static void
 hsa_fatal (const char *str, hsa_status_t status)
 {
-  const char* hsa_error;
+  const char *hsa_error;
   hsa_status_string (status, &hsa_error);
-  GOMP_PLUGIN_fatal ("HSA fatal error: %s (%s)", str, hsa_error);
+  GOMP_PLUGIN_fatal ("HSA fatal error: %s\nRuntime message: %s", str,
+		     hsa_error);
 }
 
 struct hsa_kernel_description
@@ -258,9 +251,9 @@ struct agent_info
   /* The HSA ISA of this agent.  */
   hsa_isa_t isa;
   /* Command queue of the agent.  */
-  hsa_queue_t* command_q;
+  hsa_queue_t *command_q;
   /* Kernel from kernel dispatch command queue.  */
-  hsa_queue_t* kernel_dispatch_command_q;
+  hsa_queue_t *kernel_dispatch_command_q;
   /* The HSA memory region from which to allocate kernel arguments.  */
   hsa_region_t kernarg_region;
 
@@ -332,8 +325,8 @@ static bool
 suitable_hsa_agent_p (hsa_agent_t agent)
 {
   hsa_device_type_t device_type;
-  hsa_status_t status = hsa_agent_get_info (agent, HSA_AGENT_INFO_DEVICE,
-					  &device_type);
+  hsa_status_t status
+    = hsa_agent_get_info (agent, HSA_AGENT_INFO_DEVICE, &device_type);
   if (status != HSA_STATUS_SUCCESS || device_type != HSA_DEVICE_TYPE_GPU)
     return false;
 
@@ -410,8 +403,9 @@ init_hsa_context (void)
 /* Callback of dispatch queues to report errors.  */
 
 static void
-queue_callback(hsa_status_t status, hsa_queue_t* queue __attribute__ ((unused)),
-	       void* data __attribute__ ((unused)))
+queue_callback (hsa_status_t status,
+		hsa_queue_t *queue __attribute__ ((unused)),
+		void *data __attribute__ ((unused)))
 {
   hsa_fatal ("Asynchronous queue error", status);
 }
@@ -420,7 +414,8 @@ queue_callback(hsa_status_t status, hsa_queue_t* queue __attribute__ ((unused)),
    used for kernarg allocations and if so write it to the memory pointed to by
    DATA and break the query.  */
 
-static hsa_status_t get_kernarg_memory_region (hsa_region_t region, void* data)
+static hsa_status_t
+get_kernarg_memory_region (hsa_region_t region, void *data)
 {
   hsa_status_t status;
   hsa_region_segment_t segment;
@@ -437,7 +432,7 @@ static hsa_status_t get_kernarg_memory_region (hsa_region_t region, void* data)
     return status;
   if (flags & HSA_REGION_GLOBAL_FLAG_KERNARG)
     {
-      hsa_region_t* ret = (hsa_region_t*) data;
+      hsa_region_t *ret = (hsa_region_t *) data;
       *ret = region;
       return HSA_STATUS_INFO_BREAK;
     }
@@ -528,7 +523,7 @@ static void
 add_module_to_agent (struct agent_info *agent, struct module_info *module)
 {
   if (agent->first_module)
-      agent->first_module->prev = module;
+    agent->first_module->prev = module;
   module->next = agent->first_module;
   module->prev = NULL;
   agent->first_module = module;
@@ -653,12 +648,12 @@ add_shared_library (const char *file_name, struct agent_info *agent)
     return NULL;
 
   unsigned size = end - start;
-  char *buf = (char *) malloc (size);
+  char *buf = (char *) GOMP_PLUGIN_malloc (size);
   memcpy (buf, start, size);
 
   library = GOMP_PLUGIN_malloc (sizeof (struct agent_info));
   library->file_name = (char *) GOMP_PLUGIN_malloc
-    ((strlen (file_name) + 1) * sizeof (char));
+    ((strlen (file_name) + 1));
   strcpy (library->file_name, file_name);
   library->image = (hsa_ext_module_t) buf;
 
@@ -746,11 +741,11 @@ create_and_finalize_hsa_program (struct agent_info *agent)
   hsa_ext_control_directives_t control_directives;
   memset (&control_directives, 0, sizeof (control_directives));
   hsa_code_object_t code_object;
-  status = hsa_ext_program_finalize(prog_handle, agent->isa,
-				    HSA_EXT_FINALIZER_CALL_CONVENTION_AUTO,
-				    control_directives, "",
-				    HSA_CODE_OBJECT_TYPE_PROGRAM,
-				    &code_object);
+  status = hsa_ext_program_finalize (prog_handle, agent->isa,
+				     HSA_EXT_FINALIZER_CALL_CONVENTION_AUTO,
+				     control_directives, "",
+				     HSA_CODE_OBJECT_TYPE_PROGRAM,
+				     &code_object);
   if (status != HSA_STATUS_SUCCESS)
     {
       hsa_warn ("Finalization of the HSA program failed", status);
@@ -760,8 +755,9 @@ create_and_finalize_hsa_program (struct agent_info *agent)
   HSA_DEBUG ("Finalization done\n");
   hsa_ext_program_destroy (prog_handle);
 
-  status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN,
-				 "", &agent->executable);
+  status
+    = hsa_executable_create (HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN,
+			     "", &agent->executable);
   if (status != HSA_STATUS_SUCCESS)
     hsa_fatal ("Could not create HSA executable", status);
 
@@ -773,8 +769,9 @@ create_and_finalize_hsa_program (struct agent_info *agent)
 	{
 	  struct global_var_info *var;
 	  var = &module->image_desc->global_variables[i];
-	  status = hsa_executable_global_variable_define
-	    (agent->executable, var->name, var->address);
+	  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);
@@ -787,11 +784,11 @@ create_and_finalize_hsa_program (struct agent_info *agent)
       module = module->next;
     }
 
-  status = hsa_executable_load_code_object(agent->executable, agent->id,
-					   code_object, "");
+  status = hsa_executable_load_code_object (agent->executable, agent->id,
+					    code_object, "");
   if (status != HSA_STATUS_SUCCESS)
     hsa_fatal ("Could not add a code object to the HSA executable", status);
-  status = hsa_executable_freeze(agent->executable, "");
+  status = hsa_executable_freeze (agent->executable, "");
   if (status != HSA_STATUS_SUCCESS)
     hsa_fatal ("Could not freeze the HSA executable", status);
 
@@ -818,17 +815,17 @@ create_single_kernel_dispatch (struct kernel_info *kernel,
 			       unsigned omp_data_size)
 {
   struct agent_info *agent = kernel->agent;
-  struct GOMP_hsa_kernel_dispatch *shadow = GOMP_PLUGIN_malloc_cleared
-    (sizeof (struct GOMP_hsa_kernel_dispatch));
+  struct GOMP_hsa_kernel_dispatch *shadow
+    = GOMP_PLUGIN_malloc_cleared (sizeof (struct GOMP_hsa_kernel_dispatch));
 
   shadow->queue = agent->command_q;
-  shadow->omp_data_memory = omp_data_size > 0
-    ? GOMP_PLUGIN_malloc (omp_data_size) : NULL;
+  shadow->omp_data_memory
+    = omp_data_size > 0 ? GOMP_PLUGIN_malloc (omp_data_size) : NULL;
   unsigned dispatch_count = kernel->dependencies_count;
   shadow->kernel_dispatch_count = dispatch_count;
 
-  shadow->children_dispatches = GOMP_PLUGIN_malloc
-    (dispatch_count * sizeof (struct GOMP_hsa_kernel_dispatch *));
+  shadow->children_dispatches
+    = GOMP_PLUGIN_malloc (dispatch_count * sizeof (shadow));
 
   shadow->object = kernel->object;
 
@@ -841,9 +838,9 @@ create_single_kernel_dispatch (struct kernel_info *kernel,
   shadow->private_segment_size = kernel->private_segment_size;
   shadow->group_segment_size = kernel->group_segment_size;
 
-  status = hsa_memory_allocate
-    (agent->kernarg_region, kernel->kernarg_segment_size,
-     &shadow->kernarg_address);
+  status
+    = hsa_memory_allocate (agent->kernarg_region, kernel->kernarg_segment_size,
+			   &shadow->kernarg_address);
   if (status != HSA_STATUS_SUCCESS)
     hsa_fatal ("Could not allocate memory for HSA kernel arguments", status);
 
@@ -855,8 +852,8 @@ create_single_kernel_dispatch (struct kernel_info *kernel,
 static void
 release_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *shadow)
 {
-  HSA_DEBUG ("Released kernel dispatch: %p has value: %lu (%p)\n",
-	     shadow, shadow->debug, (void *)shadow->debug);
+  HSA_DEBUG ("Released kernel dispatch: %p has value: %lu (%p)\n", shadow,
+	     shadow->debug, (void *) shadow->debug);
 
   hsa_memory_free (shadow->kernarg_address);
 
@@ -890,8 +887,10 @@ init_single_kernel (struct kernel_info *kernel, unsigned *max_omp_data_size)
       goto failure;
     }
   HSA_DEBUG ("Located kernel %s\n", kernel->name);
-  status = hsa_executable_symbol_get_info
-    (kernel_symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel->object);
+  status
+    = hsa_executable_symbol_get_info (kernel_symbol,
+				      HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
+				      &kernel->object);
   if (status != HSA_STATUS_SUCCESS)
     hsa_fatal ("Could not extract a kernel object from its symbol", status);
   status = hsa_executable_symbol_get_info
@@ -927,8 +926,8 @@ init_single_kernel (struct kernel_info *kernel, unsigned *max_omp_data_size)
 
   for (unsigned i = 0; i < kernel->dependencies_count; i++)
     {
-      struct kernel_info *dependency = get_kernel_for_agent
-	(agent, kernel->dependencies[i]);
+      struct kernel_info *dependency
+	= get_kernel_for_agent (agent, kernel->dependencies[i]);
 
       if (dependency == NULL)
 	{
@@ -959,8 +958,7 @@ failure:
 static void
 indent_stream (FILE *f, unsigned indent)
 {
-  for (int i = 0; i < indent; i++)
-    fputc (' ', f);
+  fprintf (f, "%*s", indent, "");
 }
 
 /* Dump kernel DISPATCH data structure and indent it by INDENT spaces.  */
@@ -995,7 +993,7 @@ print_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *dispatch, unsigned inden
   fprintf (stderr, "\n");
 
   for (unsigned i = 0; i < dispatch->kernel_dispatch_count; i++)
-      print_kernel_dispatch (dispatch->children_dispatches[i], indent + 2);
+    print_kernel_dispatch (dispatch->children_dispatches[i], indent + 2);
 }
 
 /* Create kernel dispatch data structure for a KERNEL and all its
@@ -1004,8 +1002,8 @@ print_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *dispatch, unsigned inden
 static struct GOMP_hsa_kernel_dispatch *
 create_kernel_dispatch (struct kernel_info *kernel, unsigned omp_data_size)
 {
-  struct GOMP_hsa_kernel_dispatch *shadow = create_single_kernel_dispatch
-    (kernel, omp_data_size);
+  struct GOMP_hsa_kernel_dispatch *shadow
+    = create_single_kernel_dispatch (kernel, omp_data_size);
   shadow->omp_num_threads = 64;
   shadow->debug = 0;
   shadow->omp_level = kernel->gridified_kernel_p ? 1 : 0;
@@ -1014,12 +1012,12 @@ create_kernel_dispatch (struct kernel_info *kernel, unsigned omp_data_size)
      a kernel dispatch with depth bigger than one.  */
   for (unsigned i = 0; i < kernel->dependencies_count; i++)
     {
-      struct kernel_info *dependency = get_kernel_for_agent
-	(kernel->agent, kernel->dependencies[i]);
-      shadow->children_dispatches[i] = create_single_kernel_dispatch
-	(dependency, omp_data_size);
-      shadow->children_dispatches[i]->queue =
-	kernel->agent->kernel_dispatch_command_q;
+      struct kernel_info *dependency
+	= get_kernel_for_agent (kernel->agent, kernel->dependencies[i]);
+      shadow->children_dispatches[i]
+	= create_single_kernel_dispatch (dependency, omp_data_size);
+      shadow->children_dispatches[i]->queue
+	= kernel->agent->kernel_dispatch_command_q;
       shadow->children_dispatches[i]->omp_level = 1;
     }
 
@@ -1146,7 +1144,7 @@ failure:
    FN_PTR which must point to a kernel_info structure.  */
 
 void
-GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void** args)
+GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void **args)
 {
   struct kernel_info *kernel = (struct kernel_info *) fn_ptr;
   struct agent_info *agent = kernel->agent;
@@ -1166,8 +1164,8 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void** args)
   if (!kernel->initialized)
     GOMP_PLUGIN_fatal ("Called kernel must be initialized");
 
-  struct GOMP_hsa_kernel_dispatch *shadow = create_kernel_dispatch
-    (kernel, kernel->max_omp_data_size);
+  struct GOMP_hsa_kernel_dispatch *shadow
+    = create_kernel_dispatch (kernel, kernel->max_omp_data_size);
 
   if (debug)
     {
@@ -1179,16 +1177,16 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void** args)
   HSA_DEBUG ("Got AQL index %llu\n", (long long int) index);
 
   /* Wait until the queue is not full before writing the packet.   */
-  while (index - hsa_queue_load_read_index_acquire(agent->command_q)
+  while (index - hsa_queue_load_read_index_acquire (agent->command_q)
 	 >= agent->command_q->size)
     ;
 
   hsa_kernel_dispatch_packet_t *packet;
-  packet = ((hsa_kernel_dispatch_packet_t*) agent->command_q->base_address)
-    + index % agent->command_q->size;
+  packet = ((hsa_kernel_dispatch_packet_t *) agent->command_q->base_address)
+	   + index % agent->command_q->size;
 
-  memset (((uint8_t *)packet) + 4, 0, sizeof (*packet) - 4);
-  packet->setup  |= (uint16_t) 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
+  memset (((uint8_t *) packet) + 4, 0, sizeof (*packet) - 4);
+  packet->setup |= (uint16_t) 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
   packet->grid_size_x = kla->gdims[0];
   uint32_t wgs = kla->wdims[0];
   if (wgs == 0)
@@ -1223,7 +1221,7 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void** args)
 
   HSA_DEBUG ("Going to dispatch kernel %s\n", kernel->name);
 
-  __atomic_store_n ((uint16_t*)(&packet->header), header, __ATOMIC_RELEASE);
+  __atomic_store_n ((uint16_t *) (&packet->header), header, __ATOMIC_RELEASE);
   hsa_signal_store_release (agent->command_q->doorbell_signal, index);
 
   /* TODO: GPU agents in Carrizo APUs cannot properly update L2 cache for
-- 
2.6.3