__kernel void raycast (__global float * img,  
					   __global float * h1,
					   __global float * h2, 
					   __global float * out,
					    float4  orig, 
					    float4  eyeVector, 
					    float4  cVec, 
					    float4  rVec, 
					    float16  transformMatrix, 
					    float stepSize, 
					    int length,
					    int width)
{  
	int row = get_global_id(0), col = get_global_id(1);

	if(row < 200 &&  col < 200) 
	{ 	  	
		int i;
		float t,x,y,z,probedVal, maxValue = -INFINITY; 
		float4 t_i, t_j, t_k, value, f, imgPt, pt, v; 
		int4 n; 
		
		
		float4 d = (float4) (h2[0],h1[0],h1[0],h2[0]); 
		float4 c = (float4) (h2[1],h1[1],h1[1],h2[1]); 
		float4 b = (float4) (h2[2],h1[2],h1[2],h2[2]); 
		float4 a = (float4) (h2[3],h1[3],h1[3],h2[3]); 
		
		float4 pos = orig + (float)row * rVec + (float)col * cVec; 
		float4 dir =  (pos - eyeVector) / fabs(pos - eyeVector); 
		
		pt.w = 1.0f; 
		pos.w = 1.0f; 
		dir.w = 1.0f; 
		value.w = 1.0f; 
			
		for(t = 0.0;  t < 20; t+= stepSize) 
		{ 
		
			pos = pos + stepSize * dir; 

			// Begin Probe Operation	
			
			// Transform the value to image space position.  
			imgPt = (float4) (dot(pos,transformMatrix.s0123),
							  dot(pos,transformMatrix.s4567),
							  dot(pos,transformMatrix.s89ab),
							  dot(pos,transformMatrix.scdef)); 
			

			f.xyzw = (float4) (modf(imgPt.x,&x),modf(imgPt.y,&y),modf(imgPt.z,&z),1.0f); 
			n = (int4)((int)x,(int)y,(int)z,1); 
		

			// the t value for h(fx - i) 
			t_i = (float4) (-1.0-f.x, -f.x,f.x - 1.0, f.x - 2.0); 
			
			// the t value for h(fy - j) 
			t_j = (float4) (-1.0-f.y, -f.y, f.y - 1.0,f.x - 2.0); 
			
			// the t value for h(fx - k) 
			t_k = (float4) (-1.0-f.z, -f.z, f.z - 1.0,f.z - 2.0); 
			

			value = ( (d + t_i * (c + t_i * (b + t_i * a)))  *  // h(fx - i) *
					  (d + t_j * (c + t_j * (b + t_j * a)))  *  // h(fy - j) * 
					  (d + t_k * (c + t_k * (b + t_k * a))));   // h(fz - k) 

	
			n.x = abs((n.x-1) * length * width + n.y-1 * width + n.z-1); 
			n.y = abs(n.x * length * width +  n.y * width + n.z); 
			n.z = abs((n.x+1) * length * width +  n.y+1 * width + n.z+1); 
			n.w = abs((n.x+2) * length * width + n.y+2 * width + n.z+2); 
			
			
		/*	if(n.x >= 2912 || n.y >= 2912  || n.z >= 2912  || n.w >= 2912)
			{ 
				printf("Row:%d, Col:%d\n",row,col); 
			} 
			else
			v = (float4)(img[n.x], 
						 img[n.y],
						 img[n.z],
						 img[n.w]);
	
			*/
					  
		   probedVal  = dot(v,value);  // V(n + <i,j,k>) * summations of h(x) components
			
		
		   // End Probe Operation  
		  if(maxValue < probedVal) 
				maxValue = probedVal; 
		} 
	   				
		 out[row * 200 + col] = maxValue; 
	}
}