libhsail-rt/rt/workitems.c

/* workitems.c -- The main runtime entry that performs work-item execution in
   various ways and the builtin functions closely related to the
   implementation.

   Copyright (C) 2015-2020 Free Software Foundation, Inc.
   Contributed by Pekka Jaaskelainen <pekka.jaaskelainen@parmance.com>
   for General Processor Tech.

   Permission is hereby granted, free of charge, to any person obtaining a
   copy of this software and associated documentation files
   (the "Software"), to deal in the Software without restriction, including
   without limitation the rights to use, copy, modify, merge, publish,
   distribute, sublicense, and/or sell copies of the Software, and to
   permit persons to whom the Software is furnished to do so, subject to
   the following conditions:

   The above copyright notice and this permission notice shall be included
   in all copies or substantial portions of the Software.

   THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
   OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
   MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
   IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,
   DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
   OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
   USE OR OTHER DEALINGS IN THE SOFTWARE.
*/

/* The fiber based multiple work-item work-group execution uses ucontext
   based user mode threading.  However, if gccbrig is able to optimize the
   kernel to a much faster work-group function that implements the multiple
   WI execution using loops instead of fibers requiring slow context switches,
   the fiber-based implementation won't be called.
 */

#include <stdlib.h>
#include <signal.h>
#include <string.h>

#include "workitems.h"
#include "phsa-rt.h"

#ifdef HAVE_FIBERS
#include "fibers.h"
#endif

#ifdef BENCHMARK_PHSA_RT
#include <stdio.h>
#include <time.h>

static uint64_t wi_count = 0;
static uint64_t wis_skipped = 0;
static uint64_t wi_total = 0;
static clock_t start_time;

#endif

#ifdef DEBUG_PHSA_RT
#include <stdio.h>
#endif

#define PRIVATE_SEGMENT_ALIGN 256
#define FIBER_STACK_SIZE (64*1024)
#define GROUP_SEGMENT_ALIGN 256

/* Preserve this amount of additional space in the alloca stack as we need to
   store the alloca frame pointer to the alloca frame, thus must preserve
   space for it.  This thus supports at most 1024 functions with allocas in
   a call chain.  */
#define ALLOCA_OVERHEAD 1024*4

uint32_t __hsail_workitemabsid (uint32_t dim, PHSAWorkItem *context);

uint32_t __hsail_workitemid (uint32_t dim, PHSAWorkItem *context);

uint32_t __hsail_gridgroups (uint32_t dim, PHSAWorkItem *context);

uint32_t __hsail_currentworkgroupsize (uint32_t dim, PHSAWorkItem *wi);

uint32_t __hsail_workgroupsize (uint32_t dim, PHSAWorkItem *wi);

void
phsa_fatal_error (int code)
{
  exit (code);
}

#ifdef HAVE_FIBERS
/* ucontext-based work-item thread implementation.  Runs all work-items in
   separate fibers.  */

static void
phsa_work_item_thread (int arg0, int arg1)
{
  void *arg = fiber_int_args_to_ptr (arg0, arg1);

  PHSAWorkItem *wi = (PHSAWorkItem *) arg;
  volatile PHSAWorkGroup *wg = wi->wg;
  PHSAKernelLaunchData *l_data = wi->launch_data;

  do
    {
      int retcode
	= fiber_barrier_reach ((fiber_barrier_t *) l_data->wg_start_barrier);

      /* At this point the threads can assume that either more_wgs is 0 or
	 the current_work_group_* is set to point to the WG executed next.  */
      if (!wi->wg->more_wgs)
	break;

      wi->group_x = wg->x;
      wi->group_y = wg->y;
      wi->group_z = wg->z;

      wi->cur_wg_size_x = __hsail_currentworkgroupsize (0, wi);
      wi->cur_wg_size_y = __hsail_currentworkgroupsize (1, wi);
      wi->cur_wg_size_z = __hsail_currentworkgroupsize (2, wi);

#ifdef DEBUG_PHSA_RT
      printf (
	"Running work-item %lu/%lu/%lu for wg %lu/%lu/%lu / %lu/%lu/%lu...\n",
	wi->x, wi->y, wi->z, wi->group_x, wi->group_y, wi->group_z,
	l_data->wg_max_x, l_data->wg_max_y, l_data->wg_max_z);
#endif

      if (wi->x < __hsail_currentworkgroupsize (0, wi)
	  && wi->y < __hsail_currentworkgroupsize (1, wi)
	  && wi->z < __hsail_currentworkgroupsize (2, wi))
	{
	  l_data->kernel (l_data->kernarg_addr, wi, wg->group_base_ptr,
			  wg->initial_group_offset, wg->private_base_ptr);
#ifdef DEBUG_PHSA_RT
	  printf ("done.\n");
#endif
#ifdef BENCHMARK_PHSA_RT
	  wi_count++;
#endif
	}
      else
	{
#ifdef DEBUG_PHSA_RT
	  printf ("skipped (partial WG).\n");
#endif
#ifdef BENCHMARK_PHSA_RT
	  wis_skipped++;
#endif
	}

      retcode
	= fiber_barrier_reach ((fiber_barrier_t *)
			       l_data->wg_completion_barrier);

      /* The first thread updates the WG to execute next etc.  */

      if (retcode == 0)
	{
#ifdef EXECUTE_WGS_BACKWARDS
	  if (wg->x == l_data->wg_min_x)
	    {
	      wg->x = l_data->wg_max_x - 1;
	      if (wg->y == l_data->wg_min_y)
		{
		  wg->y = l_data->wg_max_y - 1;
		  if (wg->z == l_data->wg_min_z)
		    wg->more_wgs = 0;
		  else
		    wg->z--;
		}
	      else
		wg->y--;
	    }
	  else
	    wg->x--;
#else
	  if (wg->x + 1 >= l_data->wg_max_x)
	    {
	      wg->x = l_data->wg_min_x;
	      if (wg->y + 1 >= l_data->wg_max_y)
		{
		  wg->y = l_data->wg_min_y;
		  if (wg->z + 1 >= l_data->wg_max_z)
		    wg->more_wgs = 0;
		  else
		    wg->z++;
		}
	      else
		wg->y++;
	    }
	  else
	    wg->x++;
#endif
	  wi->group_x = wg->x;
	  wi->group_y = wg->y;
	  wi->group_z = wg->z;

	  wi->cur_wg_size_x = __hsail_currentworkgroupsize (0, wi);
	  wi->cur_wg_size_y = __hsail_currentworkgroupsize (1, wi);
	  wi->cur_wg_size_z = __hsail_currentworkgroupsize (2, wi);

	  /* Reinitialize the work-group barrier according to the new WG's
	     size, which might not be the same as the previous ones, due
	     to "partial WGs".  */
	  size_t wg_size = __hsail_currentworkgroupsize (0, wi)
			   * __hsail_currentworkgroupsize (1, wi)
			   * __hsail_currentworkgroupsize (2, wi);

#ifdef DEBUG_PHSA_RT
	  printf ("Reinitializing the WG barrier to %lu.\n", wg_size);
#endif
	  fiber_barrier_init ((fiber_barrier_t *)
			      wi->launch_data->wg_sync_barrier,
			      wg_size);

#ifdef BENCHMARK_PHSA_RT
	  if (wi_count % 1000 == 0)
	    {
	      clock_t spent_time = clock () - start_time;
	      double spent_time_sec = (double) spent_time / CLOCKS_PER_SEC;
	      double wis_per_sec = wi_count / spent_time_sec;
	      uint64_t eta_sec
		= (wi_total - wi_count - wis_skipped) / wis_per_sec;

	      printf ("%lu WIs executed %lu skipped in %lus (%lu WIs/s, ETA in "
		      "%lu s)\n",
		      wi_count, wis_skipped, (uint64_t) spent_time_sec,
		      (uint64_t) wis_per_sec, (uint64_t) eta_sec);
	    }
#endif
	}
    }
  while (1);

  fiber_exit ();
}
#endif

#define MIN(a, b) ((a < b) ? a : b)
#define MAX(a, b) ((a > b) ? a : b)

#ifdef HAVE_FIBERS
/* Spawns a given number of work-items to execute a set of work-groups,
   blocks until their completion.  */

static void
phsa_execute_wi_gang (PHSAKernelLaunchData *context, void *group_base_ptr,
		      uint32_t group_local_offset, size_t wg_size_x,
		      size_t wg_size_y, size_t wg_size_z)
{
  PHSAWorkItem *wi_threads = NULL;
  PHSAWorkGroup wg;
  size_t flat_wi_id = 0, x, y, z, max_x, max_y, max_z;
  uint32_t group_x, group_y, group_z;
  fiber_barrier_t wg_start_barrier;
  fiber_barrier_t wg_completion_barrier;
  fiber_barrier_t wg_sync_barrier;

  max_x = wg_size_x == 0 ? 1 : wg_size_x;
  max_y = wg_size_y == 0 ? 1 : wg_size_y;
  max_z = wg_size_z == 0 ? 1 : wg_size_z;

  size_t wg_size = max_x * max_y * max_z;
  if (wg_size > PHSA_MAX_WG_SIZE)
    phsa_fatal_error (2);

  wg.private_segment_total_size = context->dp->private_segment_size * wg_size;
  if (wg.private_segment_total_size > 0
      && posix_memalign (&wg.private_base_ptr, PRIVATE_SEGMENT_ALIGN,
			 wg.private_segment_total_size)
	   != 0)
    phsa_fatal_error (3);

  wg.alloca_stack_p = wg.private_segment_total_size + ALLOCA_OVERHEAD;
  wg.alloca_frame_p = wg.alloca_stack_p;
  wg.initial_group_offset = group_local_offset;

#ifdef EXECUTE_WGS_BACKWARDS
  group_x = context->wg_max_x - 1;
  group_y = context->wg_max_y - 1;
  group_z = context->wg_max_z - 1;
#else
  group_x = context->wg_min_x;
  group_y = context->wg_min_y;
  group_z = context->wg_min_z;
#endif

  fiber_barrier_init (&wg_sync_barrier, wg_size);
  fiber_barrier_init (&wg_start_barrier, wg_size);
  fiber_barrier_init (&wg_completion_barrier, wg_size);

  context->wg_start_barrier = &wg_start_barrier;
  context->wg_sync_barrier = &wg_sync_barrier;
  context->wg_completion_barrier = &wg_completion_barrier;

  wg.more_wgs = 1;
  wg.group_base_ptr = group_base_ptr;

#ifdef BENCHMARK_PHSA_RT
  wi_count = 0;
  wis_skipped = 0;
  start_time = clock ();
#endif
  wi_threads = malloc (sizeof (PHSAWorkItem) * max_x * max_y * max_z);
  for (x = 0; x < max_x; ++x)
    for (y = 0; y < max_y; ++y)
      for (z = 0; z < max_z; ++z)
	{
	  PHSAWorkItem *wi = &wi_threads[flat_wi_id];
	  wi->launch_data = context;
	  wi->wg = &wg;

	  wg.x = wi->group_x = group_x;
	  wg.y = wi->group_y = group_y;
	  wg.z = wi->group_z = group_z;

	  wi->wg_size_x = context->dp->workgroup_size_x;
	  wi->wg_size_y = context->dp->workgroup_size_y;
	  wi->wg_size_z = context->dp->workgroup_size_z;

	  wi->cur_wg_size_x = __hsail_currentworkgroupsize (0, wi);
	  wi->cur_wg_size_y = __hsail_currentworkgroupsize (1, wi);
	  wi->cur_wg_size_z = __hsail_currentworkgroupsize (2, wi);

	  wi->x = x;
	  wi->y = y;
	  wi->z = z;

	  /* TODO: set the stack size according to the private
		   segment size.  Too big stack consumes huge amount of
		   memory in case of huge number of WIs and a too small stack
		   will fail in mysterious and potentially dangerous ways.  */

	  fiber_init (&wi->fiber, phsa_work_item_thread, wi,
		      FIBER_STACK_SIZE, PRIVATE_SEGMENT_ALIGN);
	  ++flat_wi_id;
	}

  do
    {
      --flat_wi_id;
      fiber_join (&wi_threads[flat_wi_id].fiber);
    }
  while (flat_wi_id > 0);

  if (wg.private_segment_total_size > 0)
    free (wg.private_base_ptr);

  free (wi_threads);
}

/* Spawn the work-item threads to execute work-groups and let
   them execute all the WGs, including a potential partial WG.  */

static void
phsa_spawn_work_items (PHSAKernelLaunchData *context, void *group_base_ptr,
		       uint32_t group_local_offset)
{
  hsa_kernel_dispatch_packet_t *dp = context->dp;
  size_t x, y, z;

  context->group_segment_start_addr = (size_t) group_base_ptr;

  /* HSA seems to allow the WG size to be larger than the grid size.  We need to
     saturate the effective WG size to the grid size to prevent the extra WIs
     from executing.  */
  size_t sat_wg_size_x, sat_wg_size_y, sat_wg_size_z, sat_wg_size;
  sat_wg_size_x = MIN (dp->workgroup_size_x, dp->grid_size_x);
  sat_wg_size_y = MIN (dp->workgroup_size_y, dp->grid_size_y);
  sat_wg_size_z = MIN (dp->workgroup_size_z, dp->grid_size_z);
  sat_wg_size = sat_wg_size_x * sat_wg_size_y * sat_wg_size_z;

#ifdef BENCHMARK_PHSA_RT
  wi_total = (uint64_t) dp->grid_size_x
	     * (dp->grid_size_y > 0 ? dp->grid_size_y : 1)
	     * (dp->grid_size_z > 0 ? dp->grid_size_z : 1);
#endif

  /* For now execute all work groups in a single coarse thread (does not utilize
     multicore/multithread).  */
  context->wg_min_x = context->wg_min_y = context->wg_min_z = 0;

  int dims = dp->setup & 0x3;

  context->wg_max_x = ((uint64_t) dp->grid_size_x + dp->workgroup_size_x - 1)
		      / dp->workgroup_size_x;

  context->wg_max_y
    = dims < 2 ? 1 : ((uint64_t) dp->grid_size_y + dp->workgroup_size_y - 1)
		       / dp->workgroup_size_y;

  context->wg_max_z
    = dims < 3 ? 1 : ((uint64_t) dp->grid_size_z + dp->workgroup_size_z - 1)
		       / dp->workgroup_size_z;

#ifdef DEBUG_PHSA_RT
  printf ("### launching work-groups %lu/%lu/%lu to %lu/%lu/%lu with "
	  "wg size %lu/%lu/%lu grid size %u/%u/%u\n",
	  context->wg_min_x, context->wg_min_y, context->wg_min_z,
	  context->wg_max_x, context->wg_max_y, context->wg_max_z,
	  sat_wg_size_x, sat_wg_size_y, sat_wg_size_z, dp->grid_size_x,
	  dp->grid_size_y, dp->grid_size_z);
#endif

  phsa_execute_wi_gang (context, group_base_ptr, group_local_offset,
			sat_wg_size_x, sat_wg_size_y, sat_wg_size_z);
}
#endif

/* Executes the given work-group function for all work groups in the grid.

   A work-group function is a version of the original kernel which executes
   the kernel for all work-items in a work-group.  It is produced by gccbrig
   if it can handle the kernel's barrier usage and is much faster way to
   execute massive numbers of work-items in a non-SPMD machine than fibers
   (easily 100x faster).  */
static void
phsa_execute_work_groups (PHSAKernelLaunchData *context, void *group_base_ptr,
			  uint32_t group_local_offset)
{
  hsa_kernel_dispatch_packet_t *dp = context->dp;
  size_t x, y, z, wg_x, wg_y, wg_z;

  context->group_segment_start_addr = (size_t) group_base_ptr;

  /* HSA seems to allow the WG size to be larger than the grid size.  We need
     to saturate the effective WG size to the grid size to prevent the extra WIs
     from executing.  */
  size_t sat_wg_size_x, sat_wg_size_y, sat_wg_size_z, sat_wg_size;
  sat_wg_size_x = MIN (dp->workgroup_size_x, dp->grid_size_x);
  sat_wg_size_y = MIN (dp->workgroup_size_y, dp->grid_size_y);
  sat_wg_size_z = MIN (dp->workgroup_size_z, dp->grid_size_z);
  sat_wg_size = sat_wg_size_x * sat_wg_size_y * sat_wg_size_z;

#ifdef BENCHMARK_PHSA_RT
  wi_total = (uint64_t) dp->grid_size_x
	     * (dp->grid_size_y > 0 ? dp->grid_size_y : 1)
	     * (dp->grid_size_z > 0 ? dp->grid_size_z : 1);
#endif

  context->wg_min_x = context->wg_min_y = context->wg_min_z = 0;

  int dims = dp->setup & 0x3;

  context->wg_max_x = ((uint64_t) dp->grid_size_x + dp->workgroup_size_x - 1)
		      / dp->workgroup_size_x;

  context->wg_max_y
    = dims < 2 ? 1 : ((uint64_t) dp->grid_size_y + dp->workgroup_size_y - 1)
		       / dp->workgroup_size_y;

  context->wg_max_z
    = dims < 3 ? 1 : ((uint64_t) dp->grid_size_z + dp->workgroup_size_z - 1)
		       / dp->workgroup_size_z;

#ifdef DEBUG_PHSA_RT
  printf ("### launching work-groups %lu/%lu/%lu to %lu/%lu/%lu with "
	  "wg size %lu/%lu/%lu grid size %u/%u/%u\n",
	  context->wg_min_x, context->wg_min_y, context->wg_min_z,
	  context->wg_max_x, context->wg_max_y, context->wg_max_z,
	  sat_wg_size_x, sat_wg_size_y, sat_wg_size_z, dp->grid_size_x,
	  dp->grid_size_y, dp->grid_size_z);
#endif

  PHSAWorkItem wi;
  PHSAWorkGroup wg;
  wi.wg = &wg;
  wi.x = wi.y = wi.z = 0;
  wi.launch_data = context;

#ifdef BENCHMARK_PHSA_RT
  start_time = clock ();
  uint64_t wg_count = 0;
#endif

  size_t wg_size = __hsail_workgroupsize (0, &wi)
		   * __hsail_workgroupsize (1, &wi)
		   * __hsail_workgroupsize (2, &wi);

  void *private_base_ptr = NULL;
  if (dp->private_segment_size > 0
      && posix_memalign (&private_base_ptr, PRIVATE_SEGMENT_ALIGN,
			 dp->private_segment_size * wg_size)
	   != 0)
    phsa_fatal_error (3);

  wg.alloca_stack_p = dp->private_segment_size * wg_size + ALLOCA_OVERHEAD;
  wg.alloca_frame_p = wg.alloca_stack_p;

  wg.private_base_ptr = private_base_ptr;
  wg.group_base_ptr = group_base_ptr;

#ifdef DEBUG_PHSA_RT
  printf ("priv seg size %u wg_size %lu @ %p\n", dp->private_segment_size,
	  wg_size, private_base_ptr);
#endif

  for (wg_z = context->wg_min_z; wg_z < context->wg_max_z; ++wg_z)
    for (wg_y = context->wg_min_y; wg_y < context->wg_max_y; ++wg_y)
      for (wg_x = context->wg_min_x; wg_x < context->wg_max_x; ++wg_x)
	{
	  wi.group_x = wg_x;
	  wi.group_y = wg_y;
	  wi.group_z = wg_z;

	  wi.wg_size_x = context->dp->workgroup_size_x;
	  wi.wg_size_y = context->dp->workgroup_size_y;
	  wi.wg_size_z = context->dp->workgroup_size_z;

	  wi.cur_wg_size_x = __hsail_currentworkgroupsize (0, &wi);
	  wi.cur_wg_size_y = __hsail_currentworkgroupsize (1, &wi);
	  wi.cur_wg_size_z = __hsail_currentworkgroupsize (2, &wi);

	  context->kernel (context->kernarg_addr, &wi, group_base_ptr,
			   group_local_offset, private_base_ptr);

#if defined (BENCHMARK_PHSA_RT)
	  wg_count++;
	  if (wg_count % 1000000 == 0)
	    {
	      clock_t spent_time = clock () - start_time;
	      uint64_t wi_count = wg_x * sat_wg_size_x + wg_y * sat_wg_size_y
				  + wg_z * sat_wg_size_z;
	      double spent_time_sec = (double) spent_time / CLOCKS_PER_SEC;
	      double wis_per_sec = wi_count / spent_time_sec;
	      uint64_t eta_sec = (wi_total - wi_count) / wis_per_sec;

	      printf ("%lu WIs executed in %lus (%lu WIs/s, ETA in %lu s)\n",
		      wi_count, (uint64_t) spent_time_sec,
		      (uint64_t) wis_per_sec, (uint64_t) eta_sec);
	    }
#endif
	}

#ifdef BENCHMARK_PHSA_RT
  clock_t spent_time = clock () - start_time;
  double spent_time_sec = (double) spent_time / CLOCKS_PER_SEC;
  double wis_per_sec = wi_total / spent_time_sec;

  printf ("### %lu WIs executed in %lu s (%lu WIs / s)\n", wi_total,
	  (uint64_t) spent_time_sec, (uint64_t) wis_per_sec);
#endif
  free (private_base_ptr);
  private_base_ptr = NULL;
}

/* gccbrig generates the following from each HSAIL kernel:

   1) The actual kernel function (a single work-item kernel or a work-group
      function) generated from HSAIL (BRIG).

	 static void _Kernel (void* args, void* context, void* group_base_ptr)
	 {
	   ...
	 }

  2) A public facing kernel function that is called from the PHSA runtime:

   a) A single work-item function (that requires fibers for multi-WI):

      void Kernel (void* context)
      {
	 __launch_launch_kernel (_Kernel, context);
      }

      or

    b) a when gccbrig could generate a work-group function:

      void Kernel (void* context)
      {
		__hsail_launch_wg_function (_Kernel, context);
      }
*/

#ifdef HAVE_FIBERS

void
__hsail_launch_kernel (gccbrigKernelFunc kernel, PHSAKernelLaunchData *context,
		       void *group_base_ptr, uint32_t group_local_offset)
{
  context->kernel = kernel;
  phsa_spawn_work_items (context, group_base_ptr, group_local_offset);
}
#endif

void
__hsail_launch_wg_function (gccbrigKernelFunc kernel,
			    PHSAKernelLaunchData *context, void *group_base_ptr,
			    uint32_t group_local_offset)
{
  context->kernel = kernel;
  phsa_execute_work_groups (context, group_base_ptr, group_local_offset);
}

uint32_t
__hsail_workitemabsid (uint32_t dim, PHSAWorkItem *context)
{
  hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;

  uint32_t id;
  switch (dim)
    {
    default:
    case 0:
      /* Overflow semantics in the case of WG dim > grid dim.  */
      id = ((uint64_t) context->group_x * dp->workgroup_size_x + context->x)
	   % dp->grid_size_x;
      break;
    case 1:
      id = ((uint64_t) context->group_y * dp->workgroup_size_y + context->y)
	   % dp->grid_size_y;
      break;
    case 2:
      id = ((uint64_t) context->group_z * dp->workgroup_size_z + context->z)
	   % dp->grid_size_z;
      break;
    }
  return id;
}

uint64_t
__hsail_workitemabsid_u64 (uint32_t dim, PHSAWorkItem *context)
{
  hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;

  uint64_t id;
  switch (dim)
    {
    default:
    case 0:
      /* Overflow semantics in the case of WG dim > grid dim.  */
      id = ((uint64_t) context->group_x * dp->workgroup_size_x + context->x)
	   % dp->grid_size_x;
      break;
    case 1:
      id = ((uint64_t) context->group_y * dp->workgroup_size_y + context->y)
	   % dp->grid_size_y;
      break;
    case 2:
      id = ((uint64_t) context->group_z * dp->workgroup_size_z + context->z)
	   % dp->grid_size_z;
      break;
    }
  return id;
}


uint32_t
__hsail_workitemid (uint32_t dim, PHSAWorkItem *context)
{
  PHSAWorkItem *c = (PHSAWorkItem *) context;
  hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;

  /* The number of dimensions is in the two least significant bits.  */
  int dims = dp->setup & 0x3;

  uint32_t id;
  switch (dim)
    {
    default:
    case 0:
      id = c->x;
      break;
    case 1:
      id = dims < 2 ? 0 : c->y;
      break;
    case 2:
      id = dims < 3 ? 0 : c->z;
      break;
    }
  return id;
}

uint32_t
__hsail_gridgroups (uint32_t dim, PHSAWorkItem *context)
{
  hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;
  int dims = dp->setup & 0x3;

  uint32_t id;
  switch (dim)
    {
    default:
    case 0:
      id = (dp->grid_size_x + dp->workgroup_size_x - 1) / dp->workgroup_size_x;
      break;
    case 1:
      id = dims < 2 ? 1 : (dp->grid_size_y + dp->workgroup_size_y - 1)
			    / dp->workgroup_size_y;
      break;
    case 2:
      id = dims < 3 ? 1 : (dp->grid_size_z + dp->workgroup_size_z - 1)
			    / dp->workgroup_size_z;
      break;
    }
  return id;
}

uint32_t
__hsail_workitemflatid (PHSAWorkItem *c)
{
  hsa_kernel_dispatch_packet_t *dp = c->launch_data->dp;

  return c->x + c->y * dp->workgroup_size_x
	 + c->z * dp->workgroup_size_x * dp->workgroup_size_y;
}

uint32_t
__hsail_currentworkitemflatid (PHSAWorkItem *c)
{
  hsa_kernel_dispatch_packet_t *dp = c->launch_data->dp;

  return c->x + c->y * __hsail_currentworkgroupsize (0, c)
	 + c->z * __hsail_currentworkgroupsize (0, c)
	     * __hsail_currentworkgroupsize (1, c);
}

void
__hsail_setworkitemid (uint32_t dim, uint32_t id, PHSAWorkItem *context)
{
  switch (dim)
    {
    default:
    case 0:
      context->x = id;
      break;
    case 1:
      context->y = id;
      break;
    case 2:
      context->z = id;
      break;
    }
}

uint64_t
__hsail_workitemflatabsid_u64 (PHSAWorkItem *context)
{
  PHSAWorkItem *c = (PHSAWorkItem *) context;
  hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;

  /* Work-item flattened absolute ID = ID0 + ID1 * max0 + ID2 * max0 * max1.  */
  uint64_t id0 = __hsail_workitemabsid (0, context);
  uint64_t id1 = __hsail_workitemabsid (1, context);
  uint64_t id2 = __hsail_workitemabsid (2, context);

  uint64_t max0 = dp->grid_size_x;
  uint64_t max1 = dp->grid_size_y;
  uint64_t id = id0 + id1 * max0 + id2 * max0 * max1;

  return id;
}

uint32_t
__hsail_workitemflatabsid_u32 (PHSAWorkItem *context)
{
  PHSAWorkItem *c = (PHSAWorkItem *) context;
  hsa_kernel_dispatch_packet_t *dp = context->launch_data->dp;

  /* work-item flattened absolute ID = ID0 + ID1 * max0 + ID2 * max0 * max1.  */
  uint64_t id0 = __hsail_workitemabsid (0, context);
  uint64_t id1 = __hsail_workitemabsid (1, context);
  uint64_t id2 = __hsail_workitemabsid (2, context);

  uint64_t max0 = dp->grid_size_x;
  uint64_t max1 = dp->grid_size_y;
  uint64_t id = id0 + id1 * max0 + id2 * max0 * max1;
  return (uint32_t) id;
}

uint32_t
__hsail_currentworkgroupsize (uint32_t dim, PHSAWorkItem *wi)
{
  hsa_kernel_dispatch_packet_t *dp = wi->launch_data->dp;
  uint32_t wg_size = 0;
  switch (dim)
    {
    default:
    case 0:
      if ((uint64_t) wi->group_x < dp->grid_size_x / dp->workgroup_size_x)
	wg_size = dp->workgroup_size_x; /* Full WG.  */
      else
	wg_size = dp->grid_size_x % dp->workgroup_size_x; /* Partial WG.  */
      break;
    case 1:
      if ((uint64_t) wi->group_y < dp->grid_size_y / dp->workgroup_size_y)
	wg_size = dp->workgroup_size_y; /* Full WG.  */
      else
	wg_size = dp->grid_size_y % dp->workgroup_size_y; /* Partial WG.  */
      break;
    case 2:
      if ((uint64_t) wi->group_z < dp->grid_size_z / dp->workgroup_size_z)
	wg_size = dp->workgroup_size_z; /* Full WG.  */
      else
	wg_size = dp->grid_size_z % dp->workgroup_size_z; /* Partial WG.  */
      break;
    }
  return wg_size;
}

uint32_t
__hsail_workgroupsize (uint32_t dim, PHSAWorkItem *wi)
{
  hsa_kernel_dispatch_packet_t *dp = wi->launch_data->dp;
  switch (dim)
    {
    default:
    case 0:
      return dp->workgroup_size_x;
    case 1:
      return dp->workgroup_size_y;
    case 2:
      return dp->workgroup_size_z;
    }
}

uint32_t
__hsail_gridsize (uint32_t dim, PHSAWorkItem *wi)
{
  hsa_kernel_dispatch_packet_t *dp = wi->launch_data->dp;
  switch (dim)
    {
    default:
    case 0:
      return dp->grid_size_x;
    case 1:
      return dp->grid_size_y;
    case 2:
      return dp->grid_size_z;
    }
}

uint32_t
__hsail_workgroupid (uint32_t dim, PHSAWorkItem *wi)
{
  switch (dim)
    {
    default:
    case 0:
      return wi->group_x;
    case 1:
      return wi->group_y;
    case 2:
      return wi->group_z;
    }
}

uint32_t
__hsail_dim (PHSAWorkItem *wi)
{
  hsa_kernel_dispatch_packet_t *dp = wi->launch_data->dp;
  return dp->setup & 0x3;
}

uint64_t
__hsail_packetid (PHSAWorkItem *wi)
{
  return wi->launch_data->packet_id;
}

uint32_t
__hsail_packetcompletionsig_sig32 (PHSAWorkItem *wi)
{
  return (uint32_t) wi->launch_data->dp->completion_signal.handle;
}

uint64_t
__hsail_packetcompletionsig_sig64 (PHSAWorkItem *wi)
{
  return (uint64_t) (wi->launch_data->dp->completion_signal.handle);
}

#ifdef HAVE_FIBERS
void
__hsail_barrier (PHSAWorkItem *wi)
{
  fiber_barrier_reach ((fiber_barrier_t *) wi->launch_data->wg_sync_barrier);
}
#endif

/* Return a 32b private segment address that points to a dynamically
   allocated chunk of 'size' with 'align'.

   Allocates the space from the end of the private segment allocated
   for the whole work group.  In implementations with separate private
   memories per WI, we will need to have a stack pointer per WI.  But in
   the current implementation, the segment is shared, so we possibly
   save some space in case all WIs do not call the alloca.

   The "alloca frames" are organized as follows:

   wg->alloca_stack_p points to the last allocated data (initially
   outside the private segment)
   wg->alloca_frame_p points to the first address _outside_ the current
   function's allocations (initially to the same as alloca_stack_p)

   The data is allocated downwards from the end of the private segment.

   In the beginning of a new function which has allocas, a new alloca
   frame is pushed which adds the current alloca_frame_p (the current
   function's frame starting point) to the top of the alloca stack and
   alloca_frame_p is set to the current stack position.

   At the exit points of a function with allocas, the alloca frame
   is popped before returning.  This involves popping the alloca_frame_p
   to the one of the previous function in the call stack, and alloca_stack_p
   similarly, to the position of the last word alloca'd by the previous
   function.
 */

uint32_t
__hsail_alloca (uint32_t size, uint32_t align, PHSAWorkItem *wi)
{
  volatile PHSAWorkGroup *wg = wi->wg;
  int64_t new_pos = wg->alloca_stack_p - size;
  while (new_pos % align != 0)
    new_pos--;
  if (new_pos < 0)
    phsa_fatal_error (2);

  wg->alloca_stack_p = new_pos;

#ifdef DEBUG_ALLOCA
  printf ("--- alloca (%u, %u) sp @%u fp @%u\n", size, align,
	  wg->alloca_stack_p, wg->alloca_frame_p);
#endif
  return new_pos;
}

/* Initializes a new "alloca frame" in the private segment.
   This should be called at all the function entry points in case
   the function contains at least one call to alloca.  */

void
__hsail_alloca_push_frame (PHSAWorkItem *wi)
{
  volatile PHSAWorkGroup *wg = wi->wg;

  /* Store the alloca_frame_p without any alignment padding so
     we know exactly where the previous frame ended after popping
     it.  */
#ifdef DEBUG_ALLOCA
  printf ("--- push frame ");
#endif
  uint32_t last_word_offs = __hsail_alloca (4, 1, wi);
  memcpy (wg->private_base_ptr + last_word_offs,
	  (const void *) &wg->alloca_frame_p, 4);
  wg->alloca_frame_p = last_word_offs;

#ifdef DEBUG_ALLOCA
  printf ("--- sp @%u fp @%u\n", wg->alloca_stack_p, wg->alloca_frame_p);
#endif
}

/* Frees the current "alloca frame" and restores the frame
   pointer.
   This should be called at all the function return points in case
   the function contains at least one call to alloca.  Restores the
   alloca stack to the condition it was before pushing the frame
   the last time.  */
void
__hsail_alloca_pop_frame (PHSAWorkItem *wi)
{
  volatile PHSAWorkGroup *wg = wi->wg;

  wg->alloca_stack_p = wg->alloca_frame_p;
  memcpy ((void *) &wg->alloca_frame_p,
	  (const void *) (wg->private_base_ptr + wg->alloca_frame_p), 4);
  /* Now frame_p points to the beginning of the previous function's
     frame and stack_p to its end.  */

  wg->alloca_stack_p += 4;

#ifdef DEBUG_ALLOCA
  printf ("--- pop frame sp @%u fp @%u\n", wg->alloca_stack_p,
	  wg->alloca_frame_p);
#endif
}