__kernel void initScalar(__global float* values,
			 __private const float4 min,
			 __private const float4 size)
{
  __private uint ind,ix,iy,iz,nbx,nby,nbz;
  __private float px,py,pz,s;

  ix = get_global_id(0);
  iy = get_global_id(1);
  iz = get_global_id(2);
  nbx = get_global_size(0);
  nby = get_global_size(1);
  nbz = get_global_size(2);

  px = min.x + (float)(ix)*size.x;
  py = min.y + (float)(iy)*size.y;
  pz = min.z + (float)(iz)*size.z;

  if ((px < 0.5f) && (py < 0.5f) && (pz < 0.5f))
    s = 1.0f;
  else
    s = 0.0f;
  // Write
  ind = iz*nby*nbx + iy*nbx + ix;
  values[ind] = s;
}

// velocity field
__kernel void initVelocity(__global float* values_x,
			   __global float* values_y,
			   __global float* values_z,
			   __private const float4 min,
			   __private const float4 size
			   )
{
  __private uint ix,iy,iz,nbx,nby,nbz, ind;

  ix = get_global_id(0);
  iy = get_global_id(1);
  iz = get_global_id(2);
  nbx = get_global_size(0);
  nby = get_global_size(1);
  nbz = get_global_size(2);


  // Write x component
  ind = iz*nby*nbx + iy*nbx + ix;
  values_x[ind] = 1.0f;
  // Write y component
  values_y[ind] = 1.0f;
  // Write z component
  values_z[ind] = 1.0f;

}