#include "sp.h"

//---------------------------------------------------------------------
// copy the exact forcing term to the right hand side;  because 
// this forcing term is known, we can store it on the whole grid
// including the boundary                   
//---------------------------------------------------------------------
__kernel void compute_rhs2(__global const double *g_forcing,
                           __global double *g_rhs,
                           int nx2,
                           int ny2,
                           int nz2)
{
  int i, j, k, m;

  k = get_global_id(2);
  j = get_global_id(1);
  i = get_global_id(0);
  if (k > (nz2+1) || j > (ny2+1) || i > (nx2+1)) return;

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

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

