#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

#define mm 10

__kernel void kernel_zran3_2(__global double* base_oz,
							__global double *g_ten,
							__global int *g_j1,
							__global int *g_j2,
							__global int *g_j3,
							int n1, int n2, int n3,
							int offset
							)
{
	int i1, i2, i3;
	int i;
	__global double *oz = base_oz + offset;
	i3 = get_global_id(0);
	if (i3 >= n3) return;

	__global double (*ten)[mm][2] = (__global double(*)[mm][2])g_ten;
	__global int (*j1)[mm][2] = (__global int(*)[mm][2])g_j1;
	__global int (*j2)[mm][2] = (__global int(*)[mm][2])g_j2;
	__global int (*j3)[mm][2] = (__global int(*)[mm][2])g_j3;

	  for (i = 0; i < mm; i++) {
	    ten[i3][i][1] = 0.0;
	    j1[i3][i][1] = 0;
	    j2[i3][i][1] = 0;
	    j3[i3][i][1] = 0;
	    ten[i3][i][0] = 1.0;
	    j1[i3][i][0] = 0;
	    j2[i3][i][0] = 0;
	    j3[i3][i][0] = 0;
	  }

	  i3 = i3+1;
  for (i2 = 1; i2 < n2-1; i2++)
	for (i1 = 1; i1 < n1-1; i1++)
	if (i3 >= 1 && i3 < n3-1)
	{
		__global double *zi3 = oz + i3*n2*n1;

		if (zi3[i2*n1+i1] > ten[i3-1][0][1]) {
		  ten[i3-1][0][1] = zi3[i2*n1+i1];
		  j1[i3-1][0][1] = i1;
		  j2[i3-1][0][1] = i2;
		  j3[i3-1][0][1] = i3;
		  bubble(ten[i3-1], j1[i3-1], j2[i3-1], j3[i3-1], mm, 1);
		}
		if (zi3[i2*n1+i1] < ten[i3-1][0][0]) {
		  ten[i3-1][0][0] = zi3[i2*n1+i1];
		  j1[i3-1][0][0] = i1;
		  j2[i3-1][0][0] = i2;
		  j3[i3-1][0][0] = i3;
		  bubble(ten[i3-1], j1[i3-1], j2[i3-1], j3[i3-1], mm, 0);
		}
	}

}

