// KernelAnalyzer: avg=8 max=255

float distanceSq(float2 this_, float4 pt) { // 5
	float2 p=(float2)(pt.x-this_.x, pt.y-this_.y);
	//return p.x*p.x + p.y*p.y;
	p=p*p;
	return p.x+p.y;
}

float4 invLengthSq(float4 this_, float4 pt1, float4 pt2) { // 5
	float4 p1=pt1-this_;
	float4 p2=pt2-this_;
	//return p.x*p.x + p.y*p.y;
	p1=p1*p1;
	p2=p2*p2;
	return native_recip((float4)(p1.x+p1.y, p1.z+p1.w, p2.x+p2.y, p2.z+p2.w));
}


__kernel void templateKernel(
__global float *a,
__constant float4 *points /*__attribute__((max_constant_size (8192)))*/,
const unsigned int points_length,
const float step_0,
const float step_1,
const unsigned int widthPerKernel,
const unsigned int width,
const float top_0,
const float top_1) { // 64.000*FullHD~132.710.400.000/0,3~1/2 TFlOp/s
	//float2 point;
	float4 point4;
	//point.y=get_global_id(0)*step_1;
	point4.y=point4.w=top_0+get_global_id(0)*step_0;

	for(int x=0; x<widthPerKernel; x++) {
		//point.x=(get_global_id(1)+x)*step_0;
		point4.x=point4.z=top_1+(get_global_id(1)+x)*step_1;

		float s=0;
		// points_length*(1+1+5=7)~28.000
		/*for(int i=0; i<points_length; i++) { 
			s+=native_recip(distanceSq(point, points[i]));
		}*/
		float4 s4=0;
		for(int i=0; i<points_length; i+=4) {
			/*s4.s0+=native_recip(distanceSq(point, points[i  ]));
			s4.s1+=native_recip(distanceSq(point, points[i+1]));
			s4.s2+=native_recip(distanceSq(point, points[i+2]));
			s4.s3+=native_recip(distanceSq(point, points[i+3]));*/
			s4+=invLengthSq(point4, (float4)(points[i].x,points[i].y,points[i+1].x,points[i+1].y), (float4)(points[i+2].x,points[i+2].y,points[i+3].x,points[i+3].y));
			//s+=s4.s0+s4.s1+s4.s2+s4.s3;
		}
		s+=s4.s0+s4.s1+s4.s2+s4.s3;
		// points_length*(7+1+1=9)~36.000
		/*float w=0;
		for(int i=0; i<points_length; i++) { 
			w+=( native_divide(native_recip(distanceSq(point, points[i])), s) )*points[i].z;//0.43836
			//w=mad( ( native_divide(native_recip(distanceSq(point, points[i])), s) ), points[i].z, w );//0.438259
		}*/
		float4 w4=0;
		for(int i=0; i<points_length; i+=4) {
			/*w4.s0+=( native_divide(native_recip(distanceSq(point, points[i  ])), s) )*points[i  ].z;
			w4.s1+=( native_divide(native_recip(distanceSq(point, points[i+1])), s) )*points[i+1].z;
			w4.s2+=( native_divide(native_recip(distanceSq(point, points[i+2])), s) )*points[i+2].z;
			w4.s3+=( native_divide(native_recip(distanceSq(point, points[i+3])), s) )*points[i+3].z;*/
			w4+=( native_divide(
								(float4)invLengthSq(point4, (float4)(points[i].x,points[i].y,points[i+1].x,points[i+1].y),
											(float4)(points[i+2].x,points[i+2].y,points[i+3].x,points[i+3].y)
								)
							, (float4)s) 
				)*(float4)(points[i].z,points[i+1].z,points[i+2].z,points[i+3].z);
			//w+=w4.s0+w4.s1+w4.s2+w4.s3;
		}
		//w+=w4.s0+w4.s1+w4.s2+w4.s3;

		a[get_global_id(0)*/*get_global_size(1)*/width+get_global_id(1)+x]=w4.s0+w4.s1+w4.s2+w4.s3;
	}
}