#pragma OPENCL EXTENSION cl_khr_fp64 : enable

#define PROBLEM_SIZE   24
#define IMAX      PROBLEM_SIZE
#define JMAX      PROBLEM_SIZE
#define KMAX      PROBLEM_SIZE
#define IMAXP     IMAX/2*2
#define JMAXP     JMAX/2*2

__kernel void add(__global double *g_u,
                  __global double *g_rhs,
                  int gp0,
                  int gp1,
                  int gp2)
{
  int i, j, k, m;

  k = get_global_id(2) + 1;
  j = get_global_id(1) + 1;
  i = get_global_id(0) + 1;
  if (k > (gp2-2) || j > (gp1-2) || i > (gp0-2)) return;

  __global double (*u)[JMAXP+1][IMAXP+1][5] =
    (__global double (*)[JMAXP+1][IMAXP+1][5])g_u;
  __global double (*rhs)[JMAXP+1][IMAXP+1][5] =
    (__global double (*)[JMAXP+1][IMAXP+1][5])g_rhs;

  for (m = 0; m < 5; m++) {
    u[k][j][i][m] = u[k][j][i][m] + rhs[k][j][i][m];
  }
}
