#include "mg.h"
#include "mg_dim.h"

#define PSINV_DIM       PSINV_DIM_GPU
#define RESID_DIM       RESID_DIM_GPU
#define RPRJ3_DIM       RPRJ3_DIM_GPU
#define INTERP_1_DIM    INTERP_1_DIM_GPU
#define NORM2U3_DIM     NORM2U3_DIM_GPU
#define COMM3_1_DIM     COMM3_1_DIM_GPU
#define COMM3_2_DIM     COMM3_2_DIM_GPU
#define COMM3_3_DIM     COMM3_3_DIM_GPU
#define ZERO3_DIM       ZERO3_DIM_GPU

__kernel void kernel_comm3_2(__global double *u,
                             int n1, int n2, int n3, int offset)
{
  int i3 = get_global_id(1) + 1;
  int i1 = get_global_id(0);
  if (i1 >= n1) return;

  u[i3*n2*n1+0*n1+i1+offset]      = u[i3*n2*n1+(n2-2)*n1+i1+offset];
  u[i3*n2*n1+(n2-1)*n1+i1+offset] = u[i3*n2*n1+1*n1+i1+offset];
}

