

typedef struct{
	uint	A;
	uint	B;
	uint	C;
	//uint	idM;
}cl_triangle;

typedef struct{
	float3	position;
	float3	normal;
	float3	uv;		
}cl_vertex;

typedef struct{
	float3	origin;
	float3	direction;
}Ray;
typedef struct{
	float3	A;
	float3	B;
	float3	C;
}TriangleFace;

typedef struct{
	float3	position;		
	float4	diffuse;
	//float16 matrix;
	float3 direction;
	float3 attenuation;
	float3	AEI;//angle, exponent intenzity
	uint3	type;
}cl_light;
typedef struct{
	float16 matrix;
	uint16 indexes;// s0 = start vertical s1 = count vertical, s2 = start horizontal, s3 = count horizontal, s4 = start intenzity , s5 = count intenzity
}cl_IESProfile;
typedef struct{
	float3	ambient;
	float4	diffuse;
	float3	specular;
	float3	emission;
	float3	RRS; //reflection refraction shinnes		
	uint4	TexWHOT; //id texture
	//cl_uint4 idTex; // 0 = idtex
}cl_material;
/*
typedef struct{
cl_uint4 TexWHOT; // O = offset texture(id), W = width and H= height texture if width or heigth ==0 then material havent texture, T = Type texture
}cl_texture;
*/
typedef struct{
	float2	A;
	float2	B;
	float2	C;
}UVface;

typedef struct
{
	float3	max;// max of aabb box
	float3	min;// min of aabb box	
}cl_AABB;

typedef struct{
	cl_AABB		aabb;
	int4		IPLR;// i = Index, P = Parent, L = index of Left child, P = index of Right child
	uint4		kdRMMA; //S = Start, C = count, M = id materialu
}cl_bvh;
typedef struct{
	int kdroot;
	int material;
	float dist;
}cl_bvhCache;
typedef struct{
	//cl_bvh bvh;
	//bool LR;///first = true left second = false rigth third del
	uint3	LR; //left right index
}cl_stackBVH;

typedef struct{
	uint4	SCA;//S = start index in array  C = count triangles A = axis S = 0 && C = 0 => no leaf
	int4	IPLR;// I = Index, P = Parent, L = index of Left child, P = index of Right child
	float4 splitPlanePos;
}cl_kdtree;
typedef struct{
	Ray ray;
	float3 color;
	uint depthR;
	float intensity;
	float2 intz;
}cl_RaytracerStack;
typedef struct{
	float	A[3];
	float	B[3];
	//float *	A;
	//float *	B;
	//float3 A;
	//float3 B;
	uint4	node;	
}cl_traversalInfo;
typedef struct
{
	uint2 wh;
	uint index;
}cl_textureIndex;
typedef struct
{
	float4 color;
	float dist;
}cl_colorDistance;
typedef struct
{
	int2 idTM; //id triangle , id material
	float3 tuv;
}cl_zBuffer;

#define MAXZBUFFER 100
#define fovX 60.0f
#define MISS (float3)( -1.0f, 0.0f, 0.0f )
#define LPOS (float3)( 0.0f, 0.0f, -20.0f)
#define SHADOW_OFFSET 0.005f
#define COL (float3)( 0.5f, 0.2f, 0.7f)
#define MAXSTACKKD 100
#define MAXSTACKBVH 100
#define IOR_AIR 1.0f
#define MAXRAYTRACERSTACK 1024
#define BCOLOR (float4)(0.4f, 0.3f, 0.6f, 1.0f)
#define BCOLOR3 (float3)(0.4f, 0.3f, 0.6f)

float4 readImage(__global float4 *dbtextures, uint2 wh, uint offset, uint type, float2 coordinate)
{
	float4 rColor = (float4)(0.0f, 0.0f, 0.0f, 1.0f);	
	float2 fxy =  convert_float2(wh) * coordinate;//(float2)(wh.x * coordinate.x, wh.y * coordinate.y);

	int2 uixy1 = convert_int2(fxy);// (uint2)(fxy.x, fxy.y);
	uixy1 = uixy1 % convert_int2(wh);//cpu//(uint2)(uixy1.x % wh.x, uixy1.y % wh.y);

	int tmp1 = uixy1.y * wh.x;
	rColor = dbtextures[ tmp1 + uixy1.x + offset];

	if (type == 1)// interpolation
	{
		//uint2 uixy2 = uixy1 % wh;//(uint2)(uixy1.x % wh.x, uixy1.y % wh.y);

		float2 ab = (fxy - convert_float2(uixy1));// (float2)(fxy.x - uixy1.x, fxy.y - uixy1.y);


		int tmp2 = (( uixy1.y + 1) % wh.y ) * wh.x;
		int tmp3 = ( uixy1.x + 1 ) % wh.x;

		float4 Q11, Q12, Q21, Q22;
		Q11 = rColor; //dbtextures[tmp1 + uixy1.x + offset];
		Q21 = dbtextures[tmp1 + tmp3 + offset];
		Q12	= dbtextures[tmp2 + uixy1.x + offset];
		Q22 = dbtextures[tmp2 + tmp3 + offset];

		rColor = ab.x * ( ab.y * Q22 + (1.0f - ab.y) * Q21) + (1.0f - ab.x) * ( ab.y * Q12 + ( 1.0f - ab.y) * Q11);		
	}

	return rColor;
}

float3 IntersectionTriangle( Ray R, float3 v1, float3 v2, float3 v3 , int8 typeR)
{
	float3 e1 = v2 - v1;
	float3 e2 = v3 - v1;
	//float3 v1 = T.v1;

	float u;
	float v;
	float t;
	float maxDist = 100000;
	const float threshold = 1e-6f;

	float3 p = cross(R.direction, e2);		
	float det = dot(e1, p);		
	float3 w0 = (R.origin - v1);
	if (det > threshold)
	{
		u = dot(w0,p);
		if (u < 0.0f || u > det)
			//if (u < 0.0f && u > 0.2 || u > det && u < det - 0.2)
			return MISS;


		float3 q = cross(w0, e1);
		v = dot(R.direction, q);
		if (v < 0.0f || u + v > det)
			//if (v < 0.0f && v > 0.2 || u + v > det && u + v < det - 0.2)
			return MISS;

		t = dot(e2, q) / det;

	}
	else
	{
		if (det < -threshold)
		{

			u = dot(w0,p);
			if (u > 0.0f || u < det)
				//if (u < 0.0f && u > 0.2f || u > det && u < det - 0.2f)
				return MISS;


			float3 q = cross(w0, e1);

			v = dot(R.direction,q);
			if (v > 0.0f || u + v < det)
				//if (v < 0.0f && v > 0.2f || u + v > det && u + v < det - 0.2)
				return MISS;


			t = dot(e2,q) / det;

		} else
			return MISS;
	}

	u = u / det;
	v = v / det;
	if (typeR.s0 == 0)
		return (float3)(t, u, v);

	float3 V2 = v1 + e1;
	float3 V3 = v1 + e2;
	float3 A = R.origin + R.direction * t; // point

	float3 MN1 = e1;
	float3 MA1 = A - v1;
	float dist1 = length( cross(MN1, MA1)) / length(MN1);

	float3 MN2 = e2;
	float3 MA2 =  A - v1;
	float dist2 = length( cross(MN2, MA2)) / length(MN2);


	float3 MN3 = V3 - V2;
	float3 MA3 = A - V2;
	float dist3 = length( cross(MN3, MA3)) / length(MN3);

	float maxD = 0.01f;
	if (dist1 <= maxD || dist2 <= maxD || dist3 <= maxD)
		return (float3)(t, u , v);
	else
		return MISS;
	//float distE1 = length(e1);
	//float distE2 = length(e2);

	//float tmp1 = 1.0 / det;

	if (/*(u + v < 1.0 && u + v > 1 - det) ||*/ (u > 0.0f && u < 1.0f && v <= 0.05f) || (v > 0.0f && v < 1.0f && u <= 0.05f))
		return (float3)(t, u , v);
	else
		return MISS;
}

float3 IntersectionTriangle2( Ray R, float3 v1, float3 v2, float3 v3 , int8 typeR )
{
	float3    u, v, n;             // triangle vectors
	float3    w0, w, I;          // ray vectors
	float     r, a, b;             // params to calc ray-plane intersect


	// get triangle edge vectors and plane normal
	u = v2 - v1;
	v = v3 - v1;

	n =  normalize( cross( u, v ));// cross product
	//return n;
	/*
	if ( n.x == 0.0f && n.y == 0.0f && n.z == 0.0f )
	return MISS;	
	*/
	// ray direction vector
	w0 = R.origin - v1;
	a = - dot( n, w0 );
	b = dot( n, R.direction );
	/*
	if ( fabs(b) < SMALL_NUM ) {     // ray is parallel to triangle plane

	if (a == 0)
	{                // ray lies in triangle plane
	//return 2;				//ray lies in triangle plane
	//printf("miss");
	return MISS;
	}
	else j  return MISS;             // ray disjoint from plane
	}
	*/
	// get intersect point of ray with triangle plane
	r = a / b;
	//return (float3)(r, 1.0f, 1.0f);
	if ( r < 0.0f )                   // ray goes away from triangle
		return MISS;                  // => no intersect
	// for a segment, also test if (r > 1.0) => no intersect
	//*I = R.origine_ + r * dir; //wtf??
	I = R.origin + r * R.direction;           // intersect point of ray and plane

	// is I inside T?
	float    uu, uv, vv, wu, wv, D;
	//uu = u.DotProduct(u);
	uu = dot( u,u );
	//uv = u.DotProduct(v);
	uv = dot( u, v );
	//vv = v.DotProduct(v);
	vv = dot( v, v );
	w = I - v1;
	//wu = w.DotProduct(u);
	wu = dot( w ,u );
	//wv = w.DotProduct(v);
	wv = dot( w, v );
	D = uv * uv - uu * vv;

	// get and test parametric coords
	float s, t;
	s = (uv * wv - vv * wu) / D;

	if (s < 0.0f || s > 1.0f)        // I is outside T
		return MISS;
	t = (uv * wu - uu * wv) / D;

	if ( t < 0.0f || ( s + t ) > 1.0f )  // I is outside T
		return MISS;

	//return (float3)( 1.0f, 0.5, 0.5 );
	return (float3)( r, s, t );                      // I is in T
}

float3 TransformCoordHom(float3 v, __global float *mat)
{
	float3 ret;
	ret.x = v.x * mat[0] + v.y * mat[4] + v.z * mat[8] + mat[12];
	ret.y = v.x * mat[1] + v.y * mat[5] + v.z * mat[9] + mat[13];
	ret.z = v.x * mat[2] + v.y * mat[6] + v.z * mat[10] + mat[14];
	float w = v.x * mat[3] + v.y * mat[7] + v.z * mat[11] + mat[15];
	return ret/w;
}

float3 TransformNormal16(float3 v, float16 mat)
{
	float3 ret;
	ret.x = v.x * mat.s0 + v.y * mat.s4 + v.z * mat.s8;
	ret.y = v.x * mat.s1 + v.y * mat.s5 + v.z * mat.s9;
	ret.z = v.x * mat.s2 + v.y * mat.s6 + v.z * mat.sa;
	return ret;
}

float3 TransformNormal(float3 v, __global float *mat)
{
	float3 ret;
	ret.x = v.x * mat[0] + v.y * mat[4] + v.z * mat[8];
	ret.y = v.x * mat[1] + v.y * mat[5] + v.z * mat[9];
	ret.z = v.x * mat[2] + v.y * mat[6] + v.z * mat[10];
	return ret;
}

Ray GenerateRay(float x_p, float y_p, uint width, uint height, __global float* cam, __global float *pcam)
{
	float px = 2.0f * ((x_p / width) - 0.5f);
	//float py = 2.0f * ((-y_p / height) + 0.5f);
	float py = -2.0f * ((-y_p / height) + 0.5f);/// for invert opengl texdata
	float3 dir = (float3)( px, py, 1.0f ) ;


	float3 direction_trans = TransformNormal(TransformCoordHom(dir, pcam), cam);

	Ray r;

	r.origin.x = cam[12];
	r.origin.y = cam[13];
	r.origin.z = cam[14];
	r.direction = normalize(direction_trans);
	return r;
}

//http://www.codeproject.com/KB/graphics/Simple_Ray_TracinginC6.aspx
float3 GetPhongNormal( TriangleFace trr, float3 tuv )
{
	return ( 1.0f - ( tuv.y + tuv.z )) * trr.A + trr.B * tuv.y + trr.C * tuv.z;
}

float2 GetUV( UVface trr, float3 tuv )
{
	//float3 edge1 = trr.b - trr.a;
	//float3 edge2 = trr.c - trr.a;
	//return clamp(( 1.0f - ( tuv.y + tuv.z )) * trr.A + trr.B * tuv.y + trr.C * tuv.z, (float2)(0.0f,0.0f), (float2)(1.0f, 1.0f));
	//return ( 1.0f - ( tuv.y + tuv.z )) * trr.A + trr.B * tuv.y + trr.C * tuv.z;
	/*
	float tmp1 = ((tuv.y * (trr.A.x - trr.B.x) + tuv.z * (trr.C.x - trr.B.x) + trr.B.x) );
	float tmp2 = ((tuv.y * (trr.A.y - trr.B.y) + tuv.z * (trr.C.y - trr.B.y) + trr.B.y) );
	return (float2)(tmp1, tmp2);
	*/
	float2 tmp = ( 1.0f - ( tuv.y + tuv.z )) * trr.A + trr.B * tuv.y + trr.C * tuv.z;
	tmp -= (float2)(10.0f, 10.0f);
	float2 tmp2 =(float2)(0, 0);
	tmp = fract(tmp,&tmp2);

	return tmp;
}

/*
bool checkShadowRay(Ray ray, __global cl_vertex *verts, __global cl_triangle *triIndex, uint countT,int8 typeR)
{
//TriangleFace tf;
float3 testI;
for (uint i = 0; i < countT; i++)
{	
//tf.v1 = verts[triIndex[i].v1].position;
//tf.e1 = verts[triIndex[i].e1].position;
//tf.e2 = verts[triIndex[i].e2].position;		

//testI = IntersectionTriangle( ray, tf );
testI = IntersectionTriangle( ray, verts[triIndex[i].A].position,  verts[triIndex[i].B].position, verts[triIndex[i].C].position, typeR);
if (testI.x > 0.0f )
{
return false; //is in shadow
}

}
return true; //no shadow
}
*/

//blood force
int getIndexTriangle(float3 *info, Ray ray, uint start, uint count, __global cl_vertex *verts, __global cl_triangle *triIndex, int8 typeR)
{	
	float nearI = FLT_MAX;
	int indexI = -1;
	float3 testI;	

	int end = start + count;
	for (uint i = start; i < end; i++)
	{			
		testI = IntersectionTriangle( ray, verts[triIndex[i].A].position,  verts[triIndex[i].B].position, verts[triIndex[i].C].position, typeR);
		if (testI.x > 0.0f && testI.x < nearI)
		{
			indexI = i;

			nearI = testI.x;
			*info = testI;			
		}		
	}

	return indexI;
}

bool IntersectionAABB(Ray ray, cl_AABB aabb, float *distance)
{
	float3 t1 = (aabb.min - ray.origin) / ray.direction;
	float3 t2 = (aabb.max - ray.origin) / ray.direction;

	float3 tx1 = (float3)(fmin(t1.x, t2.x), fmin(t1.y, t2.y), fmin(t1.z, t2.z));
	float3 tx2 = (float3)(fmax(t1.x, t2.x), fmax(t1.y, t2.y), fmax(t1.z, t2.z));
	float mit = fmax(tx1.x, fmax(tx1.y, tx1.z));
	float mat = fmin(tx2.x, fmin(tx2.y, tx2.z));

	if (mit <= mat && mit < *distance)
	{
		//*distance = mit;
		return true;
	}

	//return true;
	return false;
}

bool IntersectionAABB2(Ray r, cl_AABB aabb, float *t1)
{
	float t0 = 0.0f;
	//*t1 = 1e9f;

	float tmin = -FLT_MAX;
	float tmax = FLT_MAX;
	if (r.direction.x != 0.0f)
	{
		float invDirX = 1.0f/r.direction.x;
		if (invDirX < 0.0f)
		{
			tmin = (aabb.max.x - r.origin.x) * invDirX;
			tmax = (aabb.min.x - r.origin.x) * invDirX;
		}
		else
		{
			tmin = (aabb.min.x - r.origin.x) * invDirX;
			tmax = (aabb.max.x - r.origin.x) * invDirX;
		}
	}
	else if (r.origin.x < aabb.min.x || r.origin.x > aabb.max.x)
		return false;

	if (r.direction.y != 0.0f)
	{
		float tymin, tymax;
		float invDirY = 1.0f/r.direction.y;
		if (invDirY < 0.0f)
		{
			tymin = (aabb.max.y - r.origin.y) * invDirY;
			tymax = (aabb.min.y - r.origin.y) * invDirY;
		}
		else
		{
			tymin = (aabb.min.y - r.origin.y) * invDirY;
			tymax = (aabb.max.y - r.origin.y) * invDirY;
		}
		if ( (tmin > tymax) || (tymin > tmax) )
			return false;
		if (tymin > tmin)
			tmin = tymin;
		if (tymax < tmax)
			tmax = tymax;
	}
	else if (r.origin.y < aabb.min.y || r.origin.y > aabb.max.y)
		return false;

	if (r.direction.z != 0.0f)
	{
		float tzmin, tzmax;
		float invDirZ = 1.0f/r.direction.z;
		if (invDirZ < 0.0f)
		{
			tzmin = (aabb.max.z - r.origin.z) * invDirZ;
			tzmax = (aabb.min.z - r.origin.z) * invDirZ;
		}
		else
		{
			tzmin = (aabb.min.z - r.origin.z) * invDirZ;
			tzmax = (aabb.max.z - r.origin.z) * invDirZ;
		}
		if ( (tmin > tzmax) || (tzmin > tmax) )
			return false;
		if (tzmin > tmin)
			tmin = tzmin;
		if (tzmax < tmax)
			tmax = tzmax;
	}
	else if (r.origin.z < aabb.min.z || r.origin.z > aabb.max.z)
		return false;

	bool ret = ( (tmin < *t1) && (tmax > t0) );
	/*
	if (ret && dist)
	*dist = tmin;
	*/
	return ret;
}

bool IntersectionAABB3(Ray r, cl_AABB *aabb, float *t0, float *t1)
{
	float tmin = -FLT_MAX;
	float tmax = FLT_MAX;
	if (r.direction.x != 0.0f)
	{
		float invDirX = 1.0f/r.direction.x;
		if (invDirX < 0.0f)
		{
			tmin = (aabb->max.x - r.origin.x) * invDirX;
			tmax = (aabb->min.x - r.origin.x) * invDirX;
		}
		else
		{
			tmin = (aabb->min.x - r.origin.x) * invDirX;
			tmax = (aabb->max.x - r.origin.x) * invDirX;
		}
	}
	else if (r.origin.x < aabb->min.x || r.origin.x > aabb->max.x)
		return false;

	if (r.direction.y != 0.0f)
	{
		float tymin, tymax;
		float invDirY = 1.0f/r.direction.y;
		if (invDirY < 0.0f)
		{
			tymin = (aabb->max.y - r.origin.y) * invDirY;
			tymax = (aabb->min.y - r.origin.y) * invDirY;
		}
		else
		{
			tymin = (aabb->min.y - r.origin.y) * invDirY;
			tymax = (aabb->max.y - r.origin.y) * invDirY;
		}
		if ( (tmin > tymax) || (tymin > tmax) )
			return false;
		if (tymin > tmin)
			tmin = tymin;
		if (tymax < tmax)
			tmax = tymax;
	}
	else if (r.origin.y < aabb->min.y || r.origin.y > aabb->max.y)
		return false;

	if (r.direction.z != 0.0f)
	{
		float tzmin, tzmax;
		float invDirZ = 1.0f/r.direction.z;
		if (invDirZ < 0.0f)
		{
			tzmin = (aabb->max.z - r.origin.z) * invDirZ;
			tzmax = (aabb->min.z - r.origin.z) * invDirZ;
		}
		else
		{
			tzmin = (aabb->min.z - r.origin.z) * invDirZ;
			tzmax = (aabb->max.z - r.origin.z) * invDirZ;
		}
		if ( (tmin > tzmax) || (tzmin > tmax) )
			return false;
		if (tzmin > tmin)
			tmin = tzmin;
		if (tzmax < tmax)
			tmax = tzmax;
	}
	else if (r.origin.z < aabb->min.z || r.origin.z > aabb->max.z)
		return false;

	if (*t0<tmin)
		*t0=tmin;
	if (*t1>tmax)
		*t1=tmax;
	return (*t0 <= *t1);
}

bool PointIsIn(float3 v, cl_AABB aabb)
{
	return v.x >= aabb.min.x && v.x <= aabb.max.x
		&& v.y >= aabb.min.y && v.y <= aabb.max.y
		&& v.z >= aabb.min.z && v.z <= aabb.max.z;
}

bool findEntryAndExitPoints(float3 *A, float3 *B, cl_AABB aabb, Ray ray)
{
	float tmin = -FLT_MAX;//-0.0000000000001;//-1e9f;
	float tmax = FLT_MAX;//0.0000000000001;//1e9f;
	if (ray.direction.x != 0.0f)
	{
		float invDirX = 1.0f / ray.direction.x;
		if (invDirX < 0.0f)
		{
			tmin = (aabb.max.x - ray.origin.x) * invDirX;
			tmax = (aabb.min.x - ray.origin.x) * invDirX;
		}
		else
		{
			tmin = (aabb.min.x - ray.origin.x) * invDirX;
			tmax = (aabb.max.x - ray.origin.x) * invDirX;
		}
	}
	else if (ray.origin.x < aabb.min.x || ray.origin.x > aabb.max.x)
		return false;

	if (ray.direction.y != 0.0f)
	{
		float tymin, tymax;
		float invDirY = 1.0f / ray.direction.y;
		if (invDirY < 0.0f)
		{
			tymin = (aabb.max.y - ray.origin.y) * invDirY;
			tymax = (aabb.min.y - ray.origin.y) * invDirY;
		}
		else
		{
			tymin = (aabb.min.y - ray.origin.y) * invDirY;
			tymax = (aabb.max.y - ray.origin.y) * invDirY;
		}
		if ( (tmin > tymax) || (tymin > tmax) )
			return false;
		if (tymin > tmin)
			tmin = tymin;
		if (tymax < tmax)
			tmax = tymax;
	}
	else if (ray.origin.y < aabb.min.y || ray.origin.y > aabb.max.y)
		return false;

	if (ray.direction.z != 0.0f)
	{
		float tzmin, tzmax;
		float invDirZ = 1.0f / ray.direction.z;
		if (invDirZ < 0.0f)
		{
			tzmin = (aabb.max.z - ray.origin.z) * invDirZ;
			tzmax = (aabb.min.z - ray.origin.z) * invDirZ;
		}
		else
		{
			tzmin = (aabb.min.z - ray.origin.z) * invDirZ;
			tzmax = (aabb.max.z - ray.origin.z) * invDirZ;
		}
		if ( (tmin > tzmax) || (tzmin > tmax) )
			return false;
		if (tzmin > tmin)
			tmin = tzmin;
		if (tzmax < tmax)
			tmax = tzmax;
	}
	else if (ray.origin.z < aabb.min.z || ray.origin.z > aabb.max.z)
		return false;

	//compute coordinates
	//float len = ray->direction.Length();//the direction must be normalized so len=1


	if (PointIsIn(ray.origin, aabb))
		*A = ray.origin;
	else
		*A = ray.origin + (tmin * ray.direction);

	*B = ray.origin + (tmax * ray.direction);

	return true;
}

int findIntersect(Ray ray, float3 *info, int maxDepth, int idRoot, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, __global cl_vertex *verts, __global cl_triangle *triIndex, int8 typeR, cl_AABB aabb)
{
	typedef struct
	{
		int idNode;
		float tmin;
		float tmax;
	} ToDo;

	float tmin=0.0f, tmax=FLT_MAX;//info.dist;
	if (!IntersectionAABB3(ray, &aabb, &tmin, &tmax))
		return -1;

	int bestIndex = -1;
	float bestDist = FLT_MAX;  //best intersection distance, TODO: it must be in info
	float3 invrd;
	invrd.x=1.0f/ray.direction.x;
	invrd.y=1.0f/ray.direction.y;
	invrd.z=1.0f/ray.direction.z;

	ToDo todo[256];
	int todoPos = 0;

	int idNode = idRoot;
	while (true)
	{
		if (bestDist < tmin) 
			break;

		if (kdtree[idNode].SCA.z != 3 ) //(!node->IsLeaf()) 
		{
			float split = kdtree[idNode].splitPlanePos.x;
			int axis = kdtree[idNode].SCA.z;  //division axis
			float tplane;
			int firstChild;
			int secondChild;
			int belowFirst;
			if (axis==0)
			{
				tplane = (split - ray.origin.x) * invrd.x;
				belowFirst = (ray.origin.x < split) || ((ray.origin.x == split) && (ray.direction.x <= 0));
			}
			else if (axis==1)
			{
				tplane = (split - ray.origin.y) * invrd.y;
				belowFirst = (ray.origin.y < split) || ((ray.origin.y == split) && (ray.direction.y <= 0));
			}
			else
			{
				tplane = (split - ray.origin.z) * invrd.z;
				belowFirst = (ray.origin.z < split) || ((ray.origin.z == split) && (ray.direction.z <= 0));
			}


			if (belowFirst) 
			{
				firstChild = kdtree[idNode].IPLR.z;
				secondChild = kdtree[idNode].IPLR.w;
			}
			else 
			{
				firstChild = kdtree[idNode].IPLR.w;
				secondChild = kdtree[idNode].IPLR.z;
			}

			if (tplane > tmax || tplane <= 0)
				idNode = firstChild;
			else if (tplane < tmin)
				idNode = secondChild;
			else 
			{
				todo[todoPos].idNode = secondChild;
				todo[todoPos].tmin = tplane;
				todo[todoPos].tmax = tmax;
				++todoPos;
				idNode = firstChild;
				tmax = tplane;
			}
		}
		else
		{
			for (int i = 0; i < kdtree[idNode].SCA.y; i++)
			{			
				int leafTriiAIT =  kdtree[idNode].SCA.x + i;

				int leafTrii = arrayIndexTriangle[leafTriiAIT];
				float3 testI;

				testI = IntersectionTriangle(ray, verts[triIndex[leafTrii].A].position,  verts[triIndex[leafTrii].B].position, verts[triIndex[leafTrii].C].position, typeR);
				if (testI.x > 0.0f && testI.x < bestDist)
				{
					bestDist = testI.x;
					*info = testI;
					bestIndex = leafTrii;
				}
			}			
			if (todoPos > 0) 
			{
				--todoPos;
				idNode = todo[todoPos].idNode;
				tmin = todo[todoPos].tmin;
				tmax = todo[todoPos].tmax;
			}
			else
				break;
		}
	}
	return bestIndex;
}

//bubble sort
void bSortBVH(cl_bvhCache* array, size_t count)
{
	size_t last = count ;
	size_t first = 0;
	size_t prelast;
	size_t new_last;
	size_t new_first;
	size_t postiter;
	size_t preiter;
	cl_bvhCache tmp;
	bool is_swap;
	do
	{
		is_swap = false;

		prelast = last;
		--prelast;

		new_last = first;
		for (size_t iter = first; iter != prelast; )
		{
			postiter = iter;
			++postiter;

			if (array[postiter].dist < array[iter].dist)
			{
				//std::swap(*iter, *postiter);
				tmp = array[iter];
				array[iter] = array[postiter];
				array[postiter] = tmp;

				new_last = postiter;
				is_swap = true;
			}

			iter = postiter;
		}

		if (!is_swap)	break;
		is_swap = false;

		last = new_last;
		prelast = last;
		--prelast;

		new_first = prelast;
		for (size_t iter = prelast; iter != first; )
		{
			preiter = iter;
			--preiter;

			if (array[iter].dist < array[preiter].dist)
			{
				//std::swap(*preiter, *iter);
				tmp = array[iter];
				array[iter] = array[preiter];
				array[preiter] = tmp;
				new_first = iter;
				is_swap = true;
			}

			iter = preiter;
		}

		first = new_first;
	} while(is_swap);
}

void bSort(cl_colorDistance* array, size_t count)
{
	size_t last = count ;
	size_t first = 0;
	size_t prelast;
	size_t new_last;
	size_t new_first;
	size_t postiter;
	size_t preiter;
	cl_colorDistance tmp;
	bool is_swap;
	do
	{
		is_swap = false;

		prelast = last;
		--prelast;

		new_last = first;
		for (size_t iter = first; iter != prelast; )
		{
			postiter = iter;
			++postiter;

			if (array[postiter].dist < array[iter].dist)
			{
				//std::swap(*iter, *postiter);
				tmp = array[iter];
				array[iter] = array[postiter];
				array[postiter] = tmp;

				new_last = postiter;
				is_swap = true;
			}

			iter = postiter;
		}

		if (!is_swap)	break;
		is_swap = false;

		last = new_last;
		prelast = last;
		--prelast;

		new_first = prelast;
		for (size_t iter = prelast; iter != first; )
		{
			preiter = iter;
			--preiter;

			if (array[iter].dist < array[preiter].dist)
			{
				//std::swap(*preiter, *iter);
				tmp = array[iter];
				array[iter] = array[preiter];
				array[preiter] = tmp;
				new_first = iter;
				is_swap = true;
			}

			iter = preiter;
		}

		first = new_first;
	} while(is_swap);
}

bool checkShadowKDTrans(cl_colorDistance * array, size_t *count, cl_material material, float distanceRay, Ray ray, float3 *info, int maxDepth, int idRoot, __global cl_kdtree *kdtree,
	__global uint * arrayIndexTriangle, __global cl_vertex *verts, __global cl_triangle *triIndex, __global float4 *dbtextures, int8 typeR)
{
	return true;
/*	float3 A, B;
	if (!findEntryAndExitPoints(&A, &B, kdtree[idRoot].aabb, ray))
		return true;

	cl_traversalInfo traversaldata[MAXSTACKKD];


	float bestDist = FLT_MAX;
	int bestIndex = -1;
	int index = 0;
	cl_traversalInfo nodeInfo;//(A, B, idRoot);

	nodeInfo.A[0] = A.x;
	nodeInfo.A[1] = A.y;
	nodeInfo.A[2] = A.z;
	nodeInfo.B[0] = B.x;
	nodeInfo.B[1] = B.y;
	nodeInfo.B[2] = B.z;
	nodeInfo.node.x = idRoot;

	while (true)
	{
		if (kdtree[nodeInfo.node.x].SCA.z != 3 )
		{			
			uint interior = nodeInfo.node.x;
			int coord = kdtree[interior].SCA.z;  //division axis
			if (nodeInfo.A[coord] <=  kdtree[interior].splitPlanePos.x)
			{

				if (nodeInfo.B[coord] <  kdtree[interior].splitPlanePos.x)
				{
					//visit left node
					nodeInfo.node.x = kdtree[interior].IPLR.z;
				}
				else
				{
					//compute and store location of s
					//visit left node and right node
					findEntryAndExitPoints(&A, &B, kdtree[kdtree[interior].IPLR.z].aabb, ray);					

					traversaldata[index].A[0] = B.x;
					traversaldata[index].A[1] = B.y;
					traversaldata[index].A[2] = B.z;
					traversaldata[index].B[0] = nodeInfo.B[0];
					traversaldata[index].B[1] = nodeInfo.B[1];
					traversaldata[index].B[2] = nodeInfo.B[2];
					traversaldata[index].node.x = kdtree[interior].IPLR.w;
					index++;

					nodeInfo.B[0] = B.x;
					nodeInfo.B[1] = B.y;
					nodeInfo.B[2] = B.z;
					nodeInfo.node.x = kdtree[interior].IPLR.z; //left
				}

			}
			else
			{

				if (nodeInfo.B[coord] > kdtree[interior].splitPlanePos.x)
				{
					//visit right node
					nodeInfo.node.x = kdtree[interior].IPLR.w;
				}
				else
				{
					//compute and store location of s
					//visit right node and left node
					findEntryAndExitPoints(&A, &B, kdtree[kdtree[interior].IPLR.w].aabb, ray);					

					traversaldata[index].A[0] = B.x;
					traversaldata[index].A[1] = B.y;
					traversaldata[index].A[2] = B.z;
					traversaldata[index].B[0] = nodeInfo.B[0];
					traversaldata[index].B[1] = nodeInfo.B[1];
					traversaldata[index].B[2] = nodeInfo.B[2];
					traversaldata[index].node.x = kdtree[interior].IPLR.z;
					index++;
					nodeInfo.B[0] = B.x;
					nodeInfo.B[1] = B.y;
					nodeInfo.B[2] = B.z;
					nodeInfo.node.x = kdtree[interior].IPLR.w;//right
				}

			}
		}
		else
		{			
			uint leaf = nodeInfo.node.x;

			for (int i = 0; i < kdtree[leaf].SCA.y; i++)
			{
				int leafTriiAIT =  kdtree[leaf].SCA.x + i;

				int leafTrii = arrayIndexTriangle[leafTriiAIT];

				float3 testI;
				testI = IntersectionTriangle( ray, verts[triIndex[leafTrii].A].position,  verts[triIndex[leafTrii].B].position, verts[triIndex[leafTrii].C].position, typeR);
				bool dupli = true;
				if (testI.x > 0.0f && testI.x < distanceRay)
				{
					for (size_t k = 0;k < *count; k++)
					{
						if(array[k].dist == testI.x)
						{
							dupli = false;
							break;
						}
					}

					if(dupli)
					{
						//add
						float4 color = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
						if(material.TexWHOT.x != 0)
						{
							UVface tu;
							tu.A = verts[triIndex[leafTrii].A].uv.xy;
							tu.B = verts[triIndex[leafTrii].B].uv.xy;
							tu.C = verts[triIndex[leafTrii].C].uv.xy;
							float2 pos = (float2)GetUV( tu,  testI);
							color = readImage(dbtextures, material.TexWHOT.xy , material.TexWHOT.z, material.TexWHOT.w, pos);
						}
						color *= material.diffuse;
						if( color.w == 1)
							return false;

						array[*count].dist = testI.x;
						array[*count].color = color;
						*count += 1;
					}				
				}				
			}			
			if (index > 0)
				nodeInfo = traversaldata[--index];
			else
				break;
		}
	}

	return true;*/
}

//blood force
void getIndexTriangleTrans(cl_colorDistance * array, size_t *count, float4 color, float distanceRay, Ray ray, uint start, uint countt, __global cl_vertex *verts, __global cl_triangle *triIndex, int8 typeR)
{
	float nearI = FLT_MAX;
	int indexI = -1;
	float3 testI;


	int end = start + countt;
	for(uint i = start; i < end; i++)
	{			
		testI = IntersectionTriangle( ray, verts[triIndex[i].A].position,  verts[triIndex[i].B].position, verts[triIndex[i].C].position, typeR);
		if(testI.x > 0.0f && testI.x < distanceRay)
		{
			array[*count].dist = testI.x;
			array[*count].color = color;
			*count += 1;			
		}		
	}
}

void CastShadowRayMT(Ray ray, float distanceRay, float4 *color, __global cl_vertex *verts, __global cl_triangle *triIndex,
	__global cl_bvh *bvh, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, __global cl_material * materials, __global float4 *dbtextures, int8 typeR)
{

	//cl_stackBVH stack[MAXSTACKBVH];	
	//float4 rColor = (float4)(1.0f, 1.0f, 1.0f, 1.0f);

	cl_colorDistance array[100];
	size_t count = 0;
	float nearTriangle = FLT_MAX;
	float distance;
	int indexTriangle = -1;
	int nearIndexTriangle = -1;
	float3 infoH;
	//float3 nearInfoH;
	int idM;
	//float3 infoKD;
	//int indexKD;

	int stackI[MAXSTACKBVH];
	int cnode = 0;
	int deep = 0;
	//bool returnV = false;
	//return (float3)rColor.xyz;
	for(; ;)
	{
		if(IntersectionAABB(ray, bvh[cnode].aabb, /*&distance*/&distanceRay))
		{
			if (bvh[cnode].IPLR.z == -1 && bvh[cnode].IPLR.w == -1) //find leaf
			{				
				infoH.x = FLT_MAX;				
				if((bvh[cnode].kdRMMA.w==0) || !checkShadowKDTrans(array, &count, materials[bvh[cnode].kdRMMA.z]/*.diffuse*/, distanceRay, ray, &infoH, bvh[cnode].kdRMMA.y, bvh[cnode].kdRMMA.x, kdtree, arrayIndexTriangle, verts, triIndex, dbtextures, typeR))
				{
					*color = (float4)(0.0f, 0.0f, 0.0f, 1.0f);
					return;
				}			

				if(deep == 0)
					break;
				cnode = stackI[deep];
				deep --;
			}
			else
			{
				deep ++;
				stackI[deep] = bvh[cnode].IPLR.w;
				cnode = bvh[cnode].IPLR.z;				

			}

		}
		else
		{
			if(deep == 0)
				break;

			cnode = stackI[deep];
			deep --;
		}

	}
	//return (float3)rColor.xyz / count;
	//return false;
	if(count == 0 )
		return;//(float3)rColor.xyz ;
	bSort(array, count);
	//return (float3)rColor.xyz / count;



	for(size_t i = 0; i < count; i++)
	{
		float maxC= fmax(array[i].color.x, fmax(array[i].color.y, array[i].color.z));
		if (maxC==0)
			array[i].color=(float4)(1.0f, 1.0f , 1.0f, 1.0f) * (1.0f - array[i].color.w);
		else
		{
			float transp = array[i].color.w;
			array[i].color /= maxC;
			array[i].color=(array[i].color * array[i].color.w + (float4)(1.0f, 1.0f , 1.0f, 1.0f) * (1.0f - array[i].color.w)) * (1.0f - array[i].color.w);
			array[i].color.w= 1.0f - transp;
		}
		*color *= array[i].color;
	}

	//return (float3)rColor.xyz;
}

int2 CastRayMT(Ray ray, float3 * nearInfoH, __global cl_vertex *verts, __global cl_triangle *triIndex,
	__global cl_bvh *bvh, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, int8 typeR)
{
	//cl_stackBVH stack[MAXSTACKBVH];

	float nearTriangle = FLT_MAX;
	float distance;
	int indexTriangle = -1;
	int nearIndexTriangle = -1;
	float3 infoH;
	//float3 nearInfoH;
	int idM;
	//float3 infoKD;
	//int indexKD;

	int stackI[MAXSTACKBVH];
	int cnode = 0;
	int deep = 0;
	//float3 debug;
	for(; ;)
	{
		if(IntersectionAABB(ray, bvh[cnode].aabb, /*&distance*/ &nearTriangle))
		{			
			if (bvh[cnode].IPLR.z == -1 && bvh[cnode].IPLR.w == -1 && (bvh[cnode].kdRMMA.w==1)) //find leaf
			{			
				infoH.x = FLT_MAX;
				indexTriangle = findIntersect( ray, &infoH, bvh[cnode].kdRMMA.y, bvh[cnode].kdRMMA.x, kdtree, arrayIndexTriangle, verts, triIndex, typeR, bvh[cnode].aabb);

				if(infoH.x < nearTriangle)
				{
					nearTriangle = infoH.x;
					*nearInfoH = infoH;
					nearIndexTriangle = indexTriangle;
					idM = bvh[cnode].kdRMMA.z;
				}

				if(deep == 0)
					break;
				cnode = stackI[deep];
				deep --;
			}
			else
			{
				deep ++;
				stackI[deep] = bvh[cnode].IPLR.w;
				cnode = bvh[cnode].IPLR.z;				
			}


		}
		else
		{
			if(deep == 0)
				break;

			cnode = stackI[deep];
			deep --;
		}

	}

	return (int2)(nearIndexTriangle, idM);
}

/////////hahahah
float Radius(cl_AABB aabb)
{
	return length(aabb.max - aabb.min) * 0.5f;
}

float GetIntensity(float ha, float va,/*__global cl_IESProfile *IESProfile*//*cl_IESProfile IESProfile*/uint16 IESP, __global float *IESProfileData)
{
	float tmpReturn = 0.0f;
	//uint16 IESP = IESProfile.indexes;
	int hf = 0;//horizontal first
	int hl = 0; // horizontal last
	if (/*horizontalAngles.empty() || horizontalAngles.empty()*/ IESP.s1 == 0 || IESP.s3 == 0)
		return 0.0f;
	if (ha < /*horizontalAngles[0]*/IESProfileData[IESP.s2]) //miss
		return 0.0f;
	else if (ha > /*horizontalAngles[numberOfHorizontal -1]*/IESProfileData[IESP.s2 + IESP.s3 - 1]) //miss
		return 0.0f;
	else
	{
		float halfI = 0.0f;
		uint start = IESP.s2;
		halfI = IESP.s3 - 1;
		halfI /= 2;
		int i = halfI;
		bool test = false;
		while(true)
		{

			if (ha == IESProfileData[start + i])
			{
				test = true;
				break;
			}
			else if ((ha <= IESProfileData[start + i]) && (ha >= IESProfileData[start + i - 1]))
			{
				break;
			}
			else if (ha < IESProfileData[start + i])
			{
				if (halfI > 2)
					halfI /= 2;
				i -=halfI;
			}
			else if (ha > IESProfileData[start + i])
			{
				if (halfI > 2)
					halfI /= 2;
				i += halfI;
			}
		}

		hl = i;
		if (i == 0)
		{
			hl = 1;
			hf = 0;
		}
		else
			hf = i - 1;
	}

	int vf = 0; // vertical first
	int vl = 0; // vertical last
	if (va < IESProfileData[IESP.s0])	//miss
		return 0.0f;
	else if (va > IESProfileData[IESP.s0 + IESP.s1 - 1]) //miss
		return 0.0f;
	else
	{
		float halfI = 0.0f;
		uint start = IESP.s0;
		halfI = IESP.s1 - 1;
		halfI /= 2;
		int i = halfI;
		bool test = false;
		while (true)
		{
			if (va == IESProfileData[start + i])
			{
				test = true;
				break;
			}
			else if ((va <= IESProfileData[start + i]) && (va >= IESProfileData[start + i - 1]))
			{
				break;
			}
			else if (va < IESProfileData[start + i])
			{
				if (halfI > 2)
					halfI /= 2;
				i -=halfI;
			}
			else if (va > IESProfileData[start + i])
			{
				if (halfI > 2)
					halfI /= 2;
				i +=halfI;
			}
		}

		vl = i;
		if (vl == 0)
		{
			vl = 1;
			vf = 0;
		}
		else
			vf = i - 1;
	}

	//merge point
	float Q11 = 0.0f;
	float Q12 = 0.0f;
	float Q21 = 0.0f;
	float Q22 = 0.0f;
	uint startI = IESP.s4;
	uint startV = IESP.s0;
	uint startH = IESP.s2;
	Q11 = IESProfileData[startI + ((hf * IESP.s1) + vf)];
	Q22 = IESProfileData[startI + ((hl * IESP.s1) + vl)];
	Q21 = IESProfileData[startI + ((hf * IESP.s1) + vl)];
	Q12 = IESProfileData[startI + ((hl * IESP.s1) + vf)];

	float a = (va - IESProfileData[startV + vf]) / (IESProfileData[startV + vl] - IESProfileData[startV + vf]);
	float b = (ha - IESProfileData[startH + hf]) / (IESProfileData[startH + hl] - IESProfileData[startH + hf]);
	tmpReturn = a * ( b * Q22 + (1 - b) * Q21) + (1 - a) * ( b * Q12 + ( 1 - b) * Q11);

	return tmpReturn;
}

/*
void ClampColor4(float4 *color)
		{
			*color.x=min(*color.x, 1.0f);
			*color.y=min(*color.y, 1.0f);
			*color.z=min(*color.z, 1.0f);
			*color.w=min(*color.w, 1.0f);
		}*/

void ClampColor3(float3 *color)
{
	*color=min(*color, (float3)(1.0f, 1.0f, 1.0f));
	/*
	*color.x=min(*color.x, 1.0f);
	*color.y=min(*color.y, 1.0f);
	*color.z=min(*color.z, 1.0f);
	*/
	//*color.w=min(color.w, 1);
}

float3 CastAndShadeMT(uint indexI,float3 info, Ray ray, __global cl_vertex *verts, __global cl_triangle *triIndex,
	__global cl_bvh *bvh, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, __global cl_material *materials, size_t idM,
	__global cl_light * lights, uint countL, __global cl_IESProfile *IESProfile, __global float *IESProfileData,  float3 *normal, float3 *pointI, size_t idRay, int8 typeR, __global float4 * dbtextures
#ifdef IMAGE_SUPPORT
	,__read_only image2d_t textureCL,
	sampler_t textureSamplerCL
#endif						
	)
{
	float3 rrColor = (float3)(0.0f, 0.0f, 0.0f);
	float4 texColor = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
	TriangleFace tf;
	tf.A = verts[triIndex[indexI].A].normal;
	tf.B = verts[triIndex[indexI].B].normal;
	tf.C = verts[triIndex[indexI].C].normal;
	float3 N = GetPhongNormal(tf, info);
	cl_material material = materials[idM];

#ifdef IMAGE_SUPPORT
	UVface tu;
	tu.A = verts[triIndex[indexI].A].uv.xy;
	tu.B = verts[triIndex[indexI].B].uv.xy;
	tu.C = verts[triIndex[indexI].C].uv.xy;
	float2 pos = (float2)GetUV( tu,  info);
	texColor = read_imagef(textureCL, textureSamplerCL, pos);
#else
	texColor = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
#endif

	float3 point = ( ray.origin + info.x * ray.direction );
	//float3 V = normalize( ray.origin - point );
	float3 V = normalize(point - ray.origin);//je to skoro eyeToPoint
		
	*normal = N;
	*pointI = point;

	float4 matDif = material.diffuse;
	if (material.TexWHOT.x != 0)
	{
		UVface tu;
		tu.A = verts[triIndex[indexI].A].uv.xy;
		tu.B = verts[triIndex[indexI].B].uv.xy;
		tu.C = verts[triIndex[indexI].C].uv.xy;
		float2 pos = (float2)GetUV( tu,  info);
		texColor = readImage(dbtextures, material.TexWHOT.xy, material.TexWHOT.z, material.TexWHOT.w, pos);
		matDif *= texColor;//TEX_ENV_MODULATE
	}
	if (dot(ray.direction, N) > 0.0f)
		N = -N;

	if (typeR.s5 == 1)//headlight
	{
		bool inShadow = false;

		if (dot(N, V) > 0.0f)//info.normal.Dot(V)
			inShadow = true;

		if (dot(ray.direction, N)> 0.0f)//DotProd(ray.direction, info.normal)
			inShadow = !inShadow;

		float4 lightColor = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
		 //cast shadow ray, if shadow is on and is not primary ray
		if (typeR.s1 == 1 && (!inShadow) && (idRay != 0))
		{
			Ray shadowRay;//(camera->GetOrigin(), V);
			shadowRay.origin = ray.origin;
			shadowRay.direction = V;
			CastShadowRayMT(shadowRay, distance(shadowRay.origin, point) - SHADOW_OFFSET, &lightColor, verts, triIndex, bvh, kdtree, arrayIndexTriangle, materials, dbtextures, typeR);
		 	//CastShadowRayMT(shadowRay, 0, Distance(camera->GetOrigin(), info.intersection)-SHADOW_OFFSET, stack, lightColor, 0);
			lightColor*=lightColor.w;
		 	lightColor.w=1;
		}

		if (!inShadow)
		{
			// add lighting caused by this light
			float a = dot( N, V );
			float3 phongVector = normalize(  ( 2.0f * a * N) - V);
			float specularPow = pow(  clamp(dot (  V, phongVector ), 0.0f, 1.0f ), material.RRS.z );
			rrColor += ((material.specular.xyz * specularPow + fabs(a) * matDif.xyz) * lightColor.xyz);//intel
		}
// 		else
// 			rrColor.w = matDif.w;
	}
	for (uint i = 0; i < countL; i++)
	{
		cl_light light = lights[i];
		float4 lColor = light.diffuse;
		float3 L = normalize( point - light.position );

		if (light.type.x == 1)
		{
			L = normalize(light.direction);		
		}
		else if (light.type.x == 3)
		{
			//float3 TransformNormal(float3 v, __global float *mat)
			cl_IESProfile IESP = IESProfile[light.type.y];
			float3 dirL = L;
			float3 shadowRayDirection = TransformNormal16(dirL, IESP.matrix);//lightInfo[lid].inverseWorldMat.TransformNormal( shadowRay.direction);					
			shadowRayDirection = normalize(shadowRayDirection);//.Normalize();				

			float ver=0.0f, hor=0.0f;
			float3 verV = (float3)(0.0f, shadowRayDirection.y, shadowRayDirection.z);
			verV = normalize(verV);

			if (length(verV) != 0)
			{
				//float arg=- verV.Normalize().z;
				float arg=- verV.z;

				if (arg > 1)
					arg = 1.0f;
				if (arg < -1)
					arg = -1.0f;
				ver =  acos(arg) / M_PI * 180.0f;
				if (shadowRayDirection.y < 0)
					ver = -ver;
			}
			float3 horV = (float3)(shadowRayDirection.x, 0, shadowRayDirection.z);
			
			horV = normalize(horV);
			if (length(horV)!=0)
			{
				//float arg=-horV.Normalize().z;
				float arg=-horV.z;
				if (arg > 1)
					arg = 1;
				if (arg < -1)
					arg = -1;
				hor = acos(arg) / M_PI * 180.0f;
				if (shadowRayDirection.x < 0)
					hor = -hor;
			}

			lColor *= GetIntensity(hor, ver, IESP.indexes, IESProfileData);
			lColor.w = 1;
		}

		Ray shadow_ray;
		shadow_ray.direction = L;
		float dist = distance(light.position, point);
		
		if(light.type.x == 1)
		{
			shadow_ray.origin = point - shadow_ray.direction * 2 * Radius(bvh[0].aabb);
			dist = distance(shadow_ray.origin, point);
		}
		else
			shadow_ray.origin = light.position;


		bool inShadow = false;
		if (material.diffuse.w==1)
		{
			if (dot(N, L) > 0.0f)
				inShadow = true;

			if ( dot(ray.direction, N)> 0.0f )
				inShadow = !inShadow;
		}
		if(typeR.s1 == 1) // shadow
		{
			CastShadowRayMT(shadow_ray, dist - SHADOW_OFFSET, &lColor, verts, triIndex, bvh, kdtree, arrayIndexTriangle, materials, dbtextures, typeR);

			lColor *=lColor.w;
			lColor.w = 1.0f;
		}
		
		if(!inShadow)
		{
			//L = -L;
			float a = dot( N, L );
			float3 phongVector = normalize(  ( 2.0f * a * N) - L);
			//Vector3 phongVector = (( 2.0f * a * info.normal) - L);
			//float3 Ri = (normalize( L - 2.0f * dot( L, N ) * N ));

			float specularPow = pow(  clamp(dot (  V, phongVector ), 0.0f, 1.0f ), material.RRS.z );
			//float specularPow = powf(dotEyeNormal, material->shininess );

			//L = -L;
			if (light.type.x == 3)
				specularPow = pow(specularPow, light.AEI.y);
			else if (light.type.x == 2)
			{
				float spota = dot(L, light.direction);//shadowRay.direction.Dot(lightInfo[lid].direction);
				if (acos(spota) > (light.AEI.x / (180.0f * M_PI)))
				{
					spota=0;
					specularPow = 0;
				}
				a*=pow(spota, light.AEI.y);
				specularPow = pow(specularPow, light.AEI.y);
			}
			if(light.type.x != 1)
			{
				a/= (light.attenuation.x + light.attenuation.y * dist + light.attenuation.z * dist * dist);//decrease intensity by distance				
			}
			rrColor += ((material.specular.xyz * specularPow + fabs(a) * matDif.xyz) * lColor.xyz);//intel
		}
	}
	ClampColor3(&rrColor);
	//rrColor *= material.diffuse.w;
	return rrColor;
}

float3 CastAndShadeMT3(uint indexI,float3 info, Ray ray, __global cl_vertex *verts, __global cl_triangle *triIndex,
	__global cl_bvh *bvh, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, __global cl_material *materials, size_t idM,
	__global cl_light * lights, uint countL, __global cl_IESProfile *IESProfile, __global float *IESProfileData,  float3 *normal, float3 *pointI, int8 typeR, __global float4 * dbtextures
#ifdef IMAGE_SUPPORT
	,__read_only image2d_t textureCL,
	sampler_t textureSamplerCL
#endif						
	)
{
	TriangleFace tf;
	//float3 rColor = (float3)(0.0f, 0.0f, 0.0f);
	float4 rColor = (float4)(0.0f, 0.0f, 0.0f, 1.0f);
	float3 rrColor = (float3)(0.0f, 0.0f, 0.0f);
	float4 texColor = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
	float4 matDif = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
	float4 lColor = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
	tf.A = verts[triIndex[indexI].A].normal;
	tf.B = verts[triIndex[indexI].B].normal;
	tf.C = verts[triIndex[indexI].C].normal;
	cl_material material = materials[idM];

#ifdef IMAGE_SUPPORT
	UVface tu;
	tu.A = verts[triIndex[indexI].A].uv.xy;
	tu.B = verts[triIndex[indexI].B].uv.xy;
	tu.C = verts[triIndex[indexI].C].uv.xy;
	float2 pos = (float2)GetUV( tu,  info);
	texColor = read_imagef(textureCL, textureSamplerCL, pos);
#else
	texColor = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
#endif

	float3 point = ( ray.origin + info.x * ray.direction );
	float3 V = normalize( ray.origin - point );
	//float3 V = normalize( point - ray.origin);
	float3 N = GetPhongNormal(tf, info);

	if (dot(ray.direction, N) > 0.0f /*&& material.diffuse.w == 1.0f*/)
	{
		N = -N;
		*normal = N;
		
		if(material.diffuse.w != 1.0f)
		N = -N;
	}
	else
		*normal = N;

	//*normal = N;
	*pointI = point;

	matDif *= material.diffuse;
	if(material.TexWHOT.x != 0)
	{
		UVface tu;
		tu.A = verts[triIndex[indexI].A].uv.xy;
		tu.B = verts[triIndex[indexI].B].uv.xy;
		tu.C = verts[triIndex[indexI].C].uv.xy;
		float2 pos = (float2)GetUV( tu,  info);
		texColor = readImage(dbtextures, material.TexWHOT.xy, material.TexWHOT.z, material.TexWHOT.w, pos);
		matDif *= texColor;
	}

	for(uint i = 0; i < countL; i++)
	{
		cl_light light = lights[i];
		lColor = light.diffuse;
		float3 L = normalize( point - light.position );
		//L = normalize((float3)(0.1f, 1.0f, 1.2f));//light.direction;
		//ambient
		//rColor += 0.2 * material.ambient;
		/*
		switch(light.type.x)
		{
		case 1:
		{
		L = normalize((float3)(0.1f, 1.0f, 1.2f));//light.direction;
		break;
		}
		}
		*/

		if(light.type.x == 1)
		{
			L = normalize(light.direction);
			//L.x = 0.1;
			//L.y = 1.0;
			//L.z = 0.5;
			//N = -N;
		}
		else if(light.type.x == 3)
		{
			//float3 TransformNormal(float3 v, __global float *mat)
			cl_IESProfile IESP = IESProfile[light.type.y];
			float3 dirL = L;
			float3 shadowRayDirection = TransformNormal16(dirL, IESP.matrix);//lightInfo[lid].inverseWorldMat.TransformNormal( shadowRay.direction);					
			shadowRayDirection = normalize(shadowRayDirection);//.Normalize();				

			float ver=0, hor=0;
			float3 verV = (float3)(0.0f, shadowRayDirection.y, shadowRayDirection.z);
			verV = normalize(verV);

			if (length(verV) != 0)
			{
				float arg=- verV/*.Normalize()*/.z;

				if (arg > 1)
					arg = 1;
				if (arg < -1)
					arg = -1;
				ver = acos(arg) / M_PI * 180.0f;
				if (shadowRayDirection.y < 0)
					ver = -ver;
			}
			float3 horV = (float3)(shadowRayDirection.x, 0, shadowRayDirection.z);
			//horV.Normalize();
			horV = normalize(horV);
			if (length(horV)!=0)
			{
				float arg=-horV/*.Normalize()*/.z;
				if (arg > 1)
					arg = 1;
				if (arg < -1)
					arg = -1;
				hor = acos(arg) / M_PI * 180.0f;
				if (shadowRayDirection.x < 0)
					hor = -hor;
			}
			//float GetIntensity(float ha, float va,/*__global cl_IESProfile *IESProfile*/cl_IESProfile IESProfile, __global float *IESProfileData)
			//lightColor *= lightInfo[lid].iesLight->GetIntensity(hor, ver);

			lColor *= GetIntensity(hor, ver, IESP.indexes, IESProfileData);
			lColor.w = 1;
		}

		Ray shadow_ray;
		shadow_ray.direction = L;
		float dist = distance(light.position, point);
		if(light.type.x == 1)
		{
			shadow_ray.origin = point - shadow_ray.direction * 2 * Radius(bvh[0].aabb);
			dist = distance(shadow_ray.origin, point);
		}
		else
			shadow_ray.origin = light.position;

		if(typeR.s1 == 1) // shadow
		{
			/*
			Ray shadow_ray;
			shadow_ray.direction = -L;
			float dist = distance(light.position, point);
			if(light.type.x == 1)
			{
			shadow_ray.origin = point - shadow_ray.direction * 2 * Radius(bvh[0].aabb);
			dist = distance(shadow_ray.origin, point);
			}
			else
			shadow_ray.origin = light.position;
			*/


			CastShadowRayMT(shadow_ray, dist - SHADOW_OFFSET, &lColor, verts, triIndex, bvh, kdtree, arrayIndexTriangle, materials, dbtextures, typeR);


			// 			lColor *= lColor.w;
			//
			//
			// 			rColor += clamp(dot( N, L ), 0.0f, 1.0f) * material.diffuse/*.xyz*/ * /*light.diffuse*/lColor/*.xyz*/;				
			//
			// 			rColor *= texColor;
			//
			// 			float3 Ri = (normalize( L - 2.0f * dot( L, N ) * N ));
			// 			rColor +=  (pow( ( clamp(dot ( - V, Ri ), 0.0f, 1.0f)), material.RRS.z ) * material.specular ) * lColor/*.xyz*/;
			lColor *=lColor.w;
			lColor.w = 1.0f;

		}
		// 		else
		// 		{
		// 			rColor += clamp(dot( L, N ), 0.0f, 1.0f) * material.diffuse/*.xyz*/ * light.diffuse;				
		// 			rColor *= texColor;
		//
		// 			
		// 			float3 Ri = (normalize( L - 2.0f * dot( L, N ) * N ));
		// 			rColor +=  pow( ( clamp(dot ( - V, Ri ), 0.0f, 1.0f)), material.RRS.z ) * material.specular;

		L = -L;
		float dotLN = dot( N, L );
		float a = /*fabs(dotLN);*/clamp(dotLN, 0.0f, 1.0f);
		float3 phongVector = normalize(  ( 2.0f * dotLN * N) - L);
		//Vector3 phongVector = (( 2.0f * a * info.normal) - L);
		//float3 Ri = (normalize( L - 2.0f * dot( L, N ) * N ));

		float specularPow = pow(  clamp(dot (  V, /*Ri*/phongVector ), 0.0f, 1.0f ), material.RRS.z );
		//float specularPow = powf(dotEyeNormal, material->shininess );

		L = -L;
		if (light.type.x == 3)
			specularPow = pow(specularPow, light.AEI.y);
		else if (light.type.x == 2)
		{
			float spota = dot(L, light.direction);//shadowRay.direction.Dot(lightInfo[lid].direction);
			if (acos(spota) > (light.AEI.x/(180.0f * M_PI)))
			{
				spota=0;
				specularPow = 0;
			}
			a*=pow(spota, light.AEI.y);
			specularPow = pow(specularPow, light.AEI.y);
		}
		if(light.type.x != 1)
		{
			a/= (light.attenuation.x + light.attenuation.y * dist + light.attenuation.z * dist * dist);//decrease intensity by distance				
		}
		//rColor = (float4)((material.specular * specularPow + /*fabs(a)*//*clamp(a,0.0f,1.0f)*/ a * /*material.diffuse.xyz * texColor * */matDif) * lColor, matDif.w);//intel

		//		}
		//rColor *= matDif.w;
		rrColor += rColor.xyz;
		//rColor = (float3)(0.0f, 0.0f, 0.0f);
		rColor = (float4)(0.0f, 0.0f, 0.0f, 1.0f);
	}		

	//rrColor *= material.diffuse.w;
	return rrColor;
}

float3 CastAndShadeMT2(uint indexI,float3 info, Ray ray, __global cl_vertex *verts, __global cl_triangle *triIndex,
	__global cl_bvh *bvh, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, __global cl_material *materials, size_t idM,
	__global cl_light * lights, uint countL, __global cl_IESProfile *IESProfile, __global float *IESProfileData,  float3 *normal, float3 *pointI, int8 typeR, __global float4 * dbtextures
#ifdef IMAGE_SUPPORT
	,__read_only image2d_t textureCL,
	sampler_t textureSamplerCL
#endif						
	)
{
	TriangleFace tf;
	//float3 rColor = (float3)(0.0f, 0.0f, 0.0f);
	float4 rColor = (float4)(0.0f, 0.0f, 0.0f, 1.0f);
	float3 rrColor = (float3)(0.0f, 0.0f, 0.0f);
	float4 texColor = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
	//float3 testik;


	tf.A = verts[triIndex[indexI].A].normal;
	tf.B = verts[triIndex[indexI].B].normal;
	tf.C = verts[triIndex[indexI].C].normal;
	cl_material material = materials[idM];


#ifdef IMAGE_SUPPORT
	UVface tu;
	tu.A = verts[triIndex[indexI].A].uv.xy;
	tu.B = verts[triIndex[indexI].B].uv.xy;
	tu.C = verts[triIndex[indexI].C].uv.xy;
	float2 pos = (float2)GetUV( tu,  info);
	texColor = read_imagef(textureCL, textureSamplerCL, pos);
#else
	texColor = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
#endif

	float3 point = ( ray.origin + info.x * ray.direction );
	float3 V = normalize( ray.origin - point );
	float3 N = GetPhongNormal(tf, info);
	//test

	//fuj skarede
	if (dot(ray.direction, N) > 0.0f /*&& material.diffuse.w == 1.0f*/)
	{
		N = -N;
		*normal = N;

		if(material.diffuse.w != 1.0f)
			N = -N;
	}
	else
		*normal = N;

	//*normal = N;
	*pointI = point;
	if(material.TexWHOT.x != 0)
	{
		UVface tu;
		tu.A = verts[triIndex[indexI].A].uv.xy;
		tu.B = verts[triIndex[indexI].B].uv.xy;
		tu.C = verts[triIndex[indexI].C].uv.xy;
		float2 pos = (float2)GetUV( tu,  info);
		texColor = readImage(dbtextures, material.TexWHOT.xy, material.TexWHOT.z, material.TexWHOT.w, pos);
	}

	for(uint i = 0; i < countL; i++)
	{
		cl_light light = lights[i];
		float3 L = normalize( light.position - point );
		//L = normalize((float3)(0.1f, 1.0f, 1.2f));//light.direction;
		//ambient
		//rColor += 0.2 * material.ambient;
		/*
		switch(light.type.x)
		{
		case 1:
		{
		L = normalize((float3)(0.1f, 1.0f, 1.2f));//light.direction;
		break;
		}
		}
		*/

		if(light.type.x == 1)
		{
			L = normalize(light.direction);
			//L.x = 0.1;
			//L.y = 1.0f;
			//L.z = 0.5;
			//N = -N;
		}
		else
		{
		}



		if(typeR.s1 == 1) // shadow
		{
			Ray shadow_ray;
			shadow_ray.direction = -L;
			float dist = distance(light.position, point);
			if(light.type.x == 1)
			{
				shadow_ray.origin = point - shadow_ray.direction * 2 * Radius(bvh[0].aabb);
				dist = distance(shadow_ray.origin, point);
			}
			else
				shadow_ray.origin = light.position;

			/*
			float dist = distance(light.position, point);
			Ray shadow_ray;

			if(light.type.x == 1)
			{
			shadow_ray.direction = (float3)(0.0f, 0.0f, 1.0f); //light.direction;
			//shadow_ray.origin = point - shadow_ray.direction * 2 * Radius(bvh[0].aabb);
			shadow_ray.origin = light.position;//point + L * SHADOW_OFFSET;
			dist = 1000000.0f;

			}
			else
			{
			shadow_ray.origin = light.position;//point + L * SHADOW_OFFSET;
			shadow_ray.direction = -L;
			//}
			*/
			//float4 lColor = (float4)(light.diffuse, 1.0f);
			float4 lColor = (float4)(light.diffuse);
			CastShadowRayMT(shadow_ray, dist - SHADOW_OFFSET, &lColor, verts, triIndex, bvh, kdtree, arrayIndexTriangle, materials, dbtextures, typeR);





			//lColor *= lColor.w;


			//rColor += clamp(dot( N, L ), 0.0f, 1.0f) * material.diffuse/*.xyz*/ * /*light.diffuse*/lColor/*.xyz*/;				

			//rColor *= texColor;

			//float3 Ri = (normalize( L - 2.0f * dot( L, N ) * N ));
			//rColor +=  (pow( ( clamp(dot ( - V, Ri ), 0.0f, 1.0f)), material.RRS.z ) * material.specular ) * lColor/*.xyz*/;


		}
		else
		{
			//rColor += clamp(dot( L, N ), 0.0f, 1.0f) * material.diffuse/*.xyz*/ * light.diffuse;//intel
			//rColor *= texColor;


			//float3 Ri = (normalize( L - 2.0f * dot( L, N ) * N ));
			//rColor +=  pow( ( clamp(dot ( - V, Ri ), 0.0f, 1.0f)), material.RRS.z ) * material.specular;

			//float a = clamp(dot( N, L ), 0.0f, 1.0f);
			//float3 Ri = (normalize( L - 2.0f * dot( L, N ) * N ));
			//float specularPow = pow( ( clamp(dot ( - V, Ri ), 0.0f, 1.0f)), material.RRS.z );
			//rColor = (float4)((material.specular * specularPow + fabs(a) * material.diffuse.xyz * texColor) * light.diffuse, material.diffuse.w);

		}
		//rColor *= rColor.w;
		rrColor += rColor.xyz;
		//rColor = (float3)(0.0f, 0.0f, 0.0f);
		rColor = (float4)(0.0f, 0.0f, 0.0f, 1.0f);
	}		

	//rrColor *= material.diffuse.w;
	return rrColor;
}

float4 GetCubeTexColor(Ray ray, __global float4 *data, uint tw, uint th, uint offset, uint type)
{
	float4 rColor = (float4)(0.0f, 0.0f, 0.0f, 1.0f);

	uint choseSide = 0;
	float s = 0, r = 0;
	if (fabs(ray.direction.x) > fabs(ray.direction.y))
	{
		if (fabs(ray.direction.x) > fabs(ray.direction.z))
		{
			if (ray.direction.x > 0)
			{
				// biggest positive X
				//offset *= 0;
				choseSide = 0;
				r = ray.direction.y / (ray.direction.x);
				s = -ray.direction.z / (ray.direction.x);
				//r = - r;
			}
			else
			{
				// biggest negative X
				//offset *= 1;
				choseSide = 1;
				r = (float) ray.direction.y / fabs(ray.direction.x);
				s = (float) ray.direction.z / fabs(ray.direction.x);
				//r = - r;
			}
		}
		else if (ray.direction.z > 0)
		{
			// biggest positive Z
			//offset *= 4;
			choseSide = 4;
			r = ray.direction.y / (ray.direction.z);
			s = ray.direction.x / (ray.direction.z);
			//r =  - r;
		}
		else
		{
			//biggest negative Z
			//offset *= 5;
			choseSide = 5;
			r = ray.direction.y / fabs(ray.direction.z);
			s = -ray.direction.x / fabs(ray.direction.z);
		}
	}
	else if (fabs(ray.direction.y) > fabs(ray.direction.z))
	{
		if (ray.direction.y < 0)
		{
			// biggest positive Y
			//offset *= 2;
			choseSide = 2;
			// stejne s OpenGL - pro cubeTexturu u Fabia2 nesedi spravne
			r = -ray.direction.z / (ray.direction.y);
			s = -ray.direction.x / (ray.direction.y);
		}
		else
		{
			// biggest negative Y
			//offset *= 3;
			choseSide = 3;
			// stejne s OpenGL - pro cubeTexturu u Fabia2 nesedi spravne
			r = -ray.direction.z / fabs(ray.direction.y);
			s = ray.direction.x / fabs(ray.direction.y);
		}
	}
	else if (ray.direction.z > 0)
	{
		// biggest positive Z
		//offset *= 4;
		choseSide = 4;
		r = ray.direction.y / (ray.direction.z);
		s = ray.direction.x / (ray.direction.z);
		//r = -r;
	}
	else
	{
		//biggest negative Z
		//offset *= 5;
		choseSide = 5;
		r = ray.direction.y / fabs(ray.direction.z);
		s = -ray.direction.x / fabs(ray.direction.z);

	}


	s = (s + 1) / 2.0f;
	r = (r + 1) / 2.0f;

	float fs = (tw-1) * fabs(s);
	float fr = (th-1) * fabs(r);
	uint x = fs;
	uint y = fr;
	float a , b;	
	uint oneSide = tw * th ;

	//point
	if(type == 0)
	{
		int pos = (y * tw + x)  + oneSide * choseSide;		
		return data[pos + offset];
	}



	//interpolate
	float4 Q11, Q12, Q21, Q22; //color
	uint Q11p, Q12p, Q21p, Q22p; //position

	//fs = convert_float(x);
	//fr = convert_float(y);

	a = fs - x;
	b = fr - y;

	switch (choseSide)
	{
	case 0:
		{
			if (x >= (tw - 1) && (y >= 0 && y < ( th - 1)) )
			{
				Q11p = ( y * tw + x ) /*+ oneSide * 0*/;
				Q21p = ( (y + 1) * tw + x ) /*+ oneSide * 0*/;
				Q12p = ( y * tw + (th - x) ) + oneSide * 5;
				Q22p = ( (y + 1) * tw + (th - (x + 1)) ) + oneSide * 5;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];			
			}
			else if ( (x >= 0 && x < (tw - 1)) && y >= (th - 1))
			{
				Q11p = ( y * tw + x ) /** 3*/ /*+ oneSide * 0*/;
				Q21p = ( x * tw + (th - 1) ) /** 3*/  + oneSide * 3;
				Q12p = ( y * tw + (x + 1) ) /** 3*/ /* + oneSide * 0*/;
				Q22p = ( (x + 1) * tw + (th - 1) ) /** 3*/  + oneSide * 3;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];			
			}
			//roh
			else if (x >= (tw - 1) && y >= (th - 1))
			{
				Q11p = ( y * tw + x ) /*+ oneSide * 0*/;
				Q21p = ( x * tw + (th - 1) ) + oneSide * 3;
				Q12p = ( y * tw + (th - x) ) + oneSide * 5;
				//Q22p = ( /*0 * tw +*/ (x + 1) ) /** 3*/  + oneSide * 4;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];

				Q22 = (Q21 + Q12);
				Q22 /= 2.0f;

			}
			else
			{
				Q11p = ( y * tw + x ) /*+ oneSide * 0*/;
				Q21p = ( (y + 1) * tw + x ) /*+ oneSide * 0*/;
				Q12p = ( y * tw + (x + 1) ) /* + oneSide * 0*/;
				Q22p = ( (y + 1) * tw + (x + 1) ) /* + oneSide * 0*/;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}
			break;
		}
	case 1:
		{
			if (x >= (tw - 1) && (y >= 0 && y < ( th - 1)) )
			{
				Q11p = ( y * tw + x ) + oneSide /** 1*/;
				Q21p = ( (y + 1) * tw + x ) + oneSide /** 1*/;
				Q12p = ( y * tw /*+ 0*/ ) + oneSide * 4;
				Q22p = ( (y + 1) * tw /*+ 0*/ ) + oneSide * 4;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}
			else if ( (x >= 0 && x < (tw - 1)) && y >= (th - 1))
			{
				Q11p = ( y * tw + x ) + oneSide /** 1*/;
				Q21p = ( (tw - x ) * tw /*+ (0) */) + oneSide * 3;
				Q12p = ( y * tw + (x + 1) ) + oneSide /** 1*/;
				Q22p = ( (tw - (x + 1)) * tw /*+ (0) */) + oneSide * 3;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}
			//roh
			else if (x >= (tw - 1) && y >= (th - 1))
			{
				Q11p = ( y * tw + x ) + oneSide /** 1*/;
				Q21p = ( (tw - x ) * tw /*+ (0)*/ ) + oneSide * 3;
				Q12p = ( y * tw /*+ 0*/ ) + oneSide * 4;				


				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];

				Q22 = (Q21 + Q12);
				Q22 /= 2.0f;

			}
			else
			{
				Q11p = ( y * tw + x ) + oneSide /** 1*/;
				Q21p = ( (y + 1) * tw + x ) + oneSide /** 1*/;
				Q12p = ( y * tw + (x + 1) ) + oneSide /** 1*/;
				Q22p = ( (y + 1) * tw + (x + 1) ) + oneSide /** 1*/;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}
			break;
		}
	case 2:
		{
			if (x >= (tw - 1) && (y >= 0 && y < ( th - 1)) )
			{
				Q11p = ( y * tw + x ) + oneSide * 2;
				Q21p = ( (y + 1) * tw + x ) + oneSide * 2;
				Q12p = ( /*(0) * tw +*/ (th - y) ) /*+ oneSide * 0*/;
				Q22p = ( /*(0) * tw +*/ (th - (y + 1)) ) /*+ oneSide * 0*/;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}
			else if ( (x >= 0 && x < (tw - 1)) && y >= (th - 1))
			{
				Q11p = ( y * tw + x ) + oneSide * 2;
				Q21p = ( /*(0) * tw +*/ (x ) ) + oneSide * 4;
				Q12p = ( y * tw + (x + 1) ) + oneSide * 2;
				Q22p = ( /*(0) * tw +*/ (x + 1) ) + oneSide * 4;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}
			//roh
			else if (x >= (tw - 1) && y >= (th - 1))
			{
				Q11p = ( y * tw + x ) /** 3*/ + oneSide * 2;
				Q21p = ( /*(0) * tw +*/ (x) ) /** 3*/  + oneSide * 4;
				Q12p = ( /*(0) * tw +*/ (th - y) ) /** 3*/ /*+ oneSide * 0*/;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];

				Q22 = (Q21 + Q12);
				Q22 /= 2.0f;

			}
			else
			{
				Q11p = ( y * tw + x ) + oneSide * 2;
				Q21p = ( (y + 1) * tw + x ) + oneSide * 2;
				Q12p = ( y * tw + (x + 1) ) + oneSide * 2;
				Q22p = ( (y + 1) * tw + (x + 1) ) + oneSide * 2;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}
			break;
		}
	case 3:
		{
			if (x >= (tw - 1) && (y >= 0 && y < ( th - 1)) )
			{
				Q11p = ( y * tw + x ) + oneSide * 3;
				Q21p = ( (y + 1) * tw + x ) + oneSide * 3;
				Q12p = ( (tw - 1) * tw + (y) ) /*+ oneSide * 0*/;
				Q22p = ( (tw - 1) * tw + (y + 1) ) + oneSide * 0;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}
			else if ( (x >= 0 && x < (tw - 1)) && y >= (th - 1))
			{
				Q11p = ( y * tw + x ) + oneSide * 3;
				Q21p = ( (th - 1) * tw + (tw - (x)) ) + oneSide * 5;
				Q12p = ( y * tw + (x + 1) ) + oneSide * 3;
				Q22p = ( (th - 1) * tw  + (tw -(x + 1)) ) + oneSide * 5;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];	
			}
			//roh
			else if (x >= (tw - 1) && y >= (th - 1))
			{
				Q11p = ( y * tw + x ) + oneSide * 3;
				Q21p = ( (th - 1) * tw + (tw - (x)) ) + oneSide * 5;
				Q12p = ( (tw - 1) * tw + (y) ) /*+ oneSide * 0*/;				

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];

				Q22 = (Q21 + Q12);
				Q22 /= 2.0f;

			}
			else
			{
				Q11p = ( y * tw + x ) + oneSide * 3;
				Q21p = ( (y + 1) * tw + x ) + oneSide * 3;
				Q12p = ( y * tw + (x + 1) ) + oneSide * 3;
				Q22p = ( (y + 1) * tw + (x + 1) ) + oneSide * 3;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}
			break;
		}
	case 4:
		{
			if (x >= (tw - 1) && (y >= 0 && y < ( th - 1)) )
			{
				Q11p = ( y * tw + x ) + oneSide * 4;
				Q21p = ( (y + 1) * tw + x ) + oneSide * 4;
				Q12p = ( (y) * tw /*+ 0 */) /*+ oneSide * 0*/;
				Q22p = ( (y + 1) * tw /*+ (0)*/ ) /*+ oneSide * 0*/;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}
			else if ( (x >= 0 && x < (tw - 1)) && y >= (th - 1))
			{
				Q11p = ( y * tw + x ) /** 3*/ + oneSide * 4;
				Q21p = ( /*(0) * tw +*/ ( x) ) /** 3*/  + oneSide * 3;
				Q12p = ( y * tw + (x + 1) ) /** 3*/  + oneSide * 4;
				Q22p = ( /*(0) * tw +*/ (x + 1) ) /** 3*/  + oneSide * 3;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];				
			}
			//roh
			else if (x >= (tw - 1) && y >= (th - 1))
			{
				Q11p = ( y * tw + x ) + oneSide * 4;
				Q21p = ( /*(0) * tw +*/ ( x) ) + oneSide * 3;
				Q12p = ( (y) * tw /*+ 0*/ ) /*+ oneSide * 0*/;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];

				Q22 = (Q21 + Q12);
				Q22 /= 2.0f;

			}
			else
			{
				Q11p = ( y * tw + x ) + oneSide * 4;
				Q21p = ( (y + 1) * tw + x ) + oneSide * 4;
				Q12p = ( y * tw + (x + 1) ) + oneSide * 4;
				Q22p = ( (y + 1) * tw + (x + 1) ) + oneSide * 4;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}					
			break;
		}
	case 5:
		{
			if (x >= (tw - 1) && (y >= 0 && y < ( th - 1)) )
			{
				Q11p = ( y * tw + x ) + oneSide * 5;
				Q21p = ( (y + 1) * tw + x ) + oneSide * 5;
				Q12p = ( (y) * tw /*+ (0)*/ ) + oneSide /** 1*/;
				Q22p = ( (y + 1) * tw /*+ (0)*/ ) + oneSide /** 1*/;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}
			else if ( ((x >= 0) && (x < (tw - 1))) && (y >= (th - 1)))
			{
				Q11p = ( y * tw + x ) /** 3*/ + oneSide * 5;
				Q21p = ( (th - 1) * tw + (tw - x) ) /** 3*/  + oneSide * 3;
				Q12p = ( y * tw + (x + 1) ) /** 3*/  + oneSide * 5;
				Q22p = ( (th - 1) * tw +  (tw - (x + 1)) ) /** 3*/  + oneSide * 3;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];

			}
			else if (x >= (tw - 1) && y >= (th - 1))
			{
				Q11p = ( y * tw + x ) + oneSide * 5;
				Q21p = ( (th - 1) * tw + (tw - x) )  + oneSide * 3;
				Q12p = ( (y) * tw /*+( 0)*/ ) + oneSide /** 1*/;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];

				Q22 = (Q21 + Q12);
				Q22 /= 2.0f;

			}
			else
			{
				Q11p = ( y * tw + x ) + oneSide * 5;
				Q21p = ( (y + 1) * tw + x ) + oneSide * 5;
				Q12p = ( y * tw + (x + 1) ) + oneSide * 5;
				Q22p = ( (y + 1) * tw + (x + 1) ) + oneSide * 5;

				Q11 = data[Q11p + offset];
				Q21 = data[Q21p + offset];
				Q12 = data[Q12p + offset];
				Q22 = data[Q22p + offset];
			}					
			break;
		}
	}
	return a * ( b * Q22 + (1 - b) * Q12 ) + (1 - a) * ( b * Q21 + (1 - b) * Q11 );
}

//bubble sort
void bSortZB(cl_zBuffer* array, size_t *count)
{
	size_t last = *count ;
	size_t first = 0;
	size_t prelast;
	size_t new_last;
	size_t new_first;
	size_t postiter;
	size_t preiter;
	cl_zBuffer tmp;
	bool is_swap;
	do
	{
		is_swap = false;

		prelast = last;
		--prelast;

		new_last = first;
		for (size_t iter = first; iter != prelast; )
		{
			postiter = iter;
			++postiter;

			if (array[postiter].tuv.x < array[iter].tuv.x)
			{
				//std::swap(*iter, *postiter);
				tmp = array[iter];
				array[iter] = array[postiter];
				array[postiter] = tmp;

				new_last = postiter;
				is_swap = true;
			}

			iter = postiter;
		}

		if (!is_swap)	break;
		is_swap = false;

		last = new_last;
		prelast = last;
		--prelast;

		new_first = prelast;
		for (size_t iter = prelast; iter != first; )
		{
			preiter = iter;
			--preiter;

			if (array[iter].tuv.x < array[preiter].tuv.x)
			{
				//std::swap(*preiter, *iter);
				tmp = array[iter];
				array[iter] = array[preiter];
				array[preiter] = tmp;
				new_first = iter;
				is_swap = true;
			}

			iter = preiter;
		}

		first = new_first;
	} while(is_swap);
}

bool checkShadowKDTransZB(cl_zBuffer * array, size_t *count, int idM, float distanceRay, Ray ray, float3 *info, int maxDepth, int idRoot, __global cl_kdtree *kdtree,
	__global uint * arrayIndexTriangle, __global cl_vertex *verts, __global cl_triangle *triIndex, __global float4 *dbtextures, int8 typeR)
{
	return true;
/*	float3 A, B;
	if (!findEntryAndExitPoints(&A, &B, kdtree[idRoot].aabb, ray))
		return true;

	cl_traversalInfo traversaldata[MAXSTACKKD];

	float bestDist = FLT_MAX;
	int bestIndex = -1;
	int index = 0;
	cl_traversalInfo nodeInfo;//(A, B, idRoot);

	nodeInfo.A[0] = A.x;
	nodeInfo.A[1] = A.y;
	nodeInfo.A[2] = A.z;
	nodeInfo.B[0] = B.x;
	nodeInfo.B[1] = B.y;
	nodeInfo.B[2] = B.z;
	nodeInfo.node.x = idRoot;

	while (true)
	{
		if (kdtree[nodeInfo.node.x].SCA.z != 3 )
		{			
			uint interior = nodeInfo.node.x;
			int coord = kdtree[interior].SCA.z;  //division axis
			if (nodeInfo.A[coord] <=  kdtree[interior].splitPlanePos.x)
			{

				if (nodeInfo.B[coord] <  kdtree[interior].splitPlanePos.x)
				{
					//visit left node
					nodeInfo.node.x = kdtree[interior].IPLR.z;
				}
				else
				{
					//compute and store location of s
					//visit left node and right node
					findEntryAndExitPoints(&A, &B, kdtree[kdtree[interior].IPLR.z].aabb, ray);					

					traversaldata[index].A[0] = B.x;
					traversaldata[index].A[1] = B.y;
					traversaldata[index].A[2] = B.z;
					traversaldata[index].B[0] = nodeInfo.B[0];
					traversaldata[index].B[1] = nodeInfo.B[1];
					traversaldata[index].B[2] = nodeInfo.B[2];
					traversaldata[index].node.x = kdtree[interior].IPLR.w;
					index++;

					nodeInfo.B[0] = B.x;
					nodeInfo.B[1] = B.y;
					nodeInfo.B[2] = B.z;
					nodeInfo.node.x = kdtree[interior].IPLR.z; //left
				}
			}
			else
			{
				if (nodeInfo.B[coord] > kdtree[interior].splitPlanePos.x)
				{
					//visit right node
					nodeInfo.node.x = kdtree[interior].IPLR.w;
				}
				else
				{
					//compute and store location of s
					//visit right node and left node
					findEntryAndExitPoints(&A, &B, kdtree[kdtree[interior].IPLR.w].aabb, ray);					

					traversaldata[index].A[0] = B.x;
					traversaldata[index].A[1] = B.y;
					traversaldata[index].A[2] = B.z;
					traversaldata[index].B[0] = nodeInfo.B[0];
					traversaldata[index].B[1] = nodeInfo.B[1];
					traversaldata[index].B[2] = nodeInfo.B[2];
					traversaldata[index].node.x = kdtree[interior].IPLR.z;
					index++;
					nodeInfo.B[0] = B.x;
					nodeInfo.B[1] = B.y;
					nodeInfo.B[2] = B.z;
					nodeInfo.node.x = kdtree[interior].IPLR.w;//right
				}
			}
		}
		else
		{			
			uint leaf = nodeInfo.node.x;
			for (int i = 0; i < kdtree[leaf].SCA.y; i++)
			{			
				int leafTriiAIT =  kdtree[leaf].SCA.x + i;

				int leafTrii = arrayIndexTriangle[leafTriiAIT];

				float3 testI;
				testI = IntersectionTriangle( ray, verts[triIndex[leafTrii].A].position,  verts[triIndex[leafTrii].B].position, verts[triIndex[leafTrii].C].position, typeR);
				bool dupli = true;
				if(testI.x > 0.0f && testI.x < distanceRay)
				{
					for(size_t k = 0; k < *count; k++)
					{
						if(array[k].tuv.x == testI.x)
						{
							dupli = false;
							break;
						}
					}

					if(dupli)
					{
						array[*count].tuv = testI;
						array[*count].idTM.x = leafTrii;
						array[*count].idTM.y = idM;
						*count += 1;
						//add
					}				
				}				
			}			
			if (index > 0)
				nodeInfo = traversaldata[--index];
			else
				break;
		}
	}

	return true;*/
}

//blood force
void getIndexTriangleTransZB(cl_colorDistance * array, size_t *count, float4 color, float distanceRay, Ray ray, uint start, uint countt, __global cl_vertex *verts, __global cl_triangle *triIndex, int8 typeR)
{
	float nearI = FLT_MAX;
	int indexI = -1;
	float3 testI;

	int end = start + countt;
	for(uint i = start; i < end; i++)
	{			
		testI = IntersectionTriangle( ray, verts[triIndex[i].A].position,  verts[triIndex[i].B].position, verts[triIndex[i].C].position, typeR);
		if(testI.x > 0.0f && testI.x < distanceRay)
		{
			array[*count].dist = testI.x;
			array[*count].color = color;
			*count += 1;			
		}		
	}
}

void CastRayMTZB(cl_zBuffer *array, size_t *count, Ray ray, float distanceRay, __global cl_vertex *verts, __global cl_triangle *triIndex,
	__global cl_bvh *bvh, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, __global cl_material * materials, __global float4 *dbtextures, int8 typeR)
{
	//cl_stackBVH stack[MAXSTACKBVH];	
	//float4 rColor = (float4)(1.0f, 1.0f, 1.0f, 1.0f);

	/*
	typedef struct
	{
	int2 idTM; //id triangle , id material
	float3 tuv;
	}cl_zBuffer;
	*/
	//cl_zBuffer array[MAXZBUFFER];
	//size_t count = 0;
	float nearTriangle = FLT_MAX;
	//float distance;
	int indexTriangle = -1;
	//int nearIndexTriangle = -1;
	float3 infoH;
	//float3 nearInfoH;
	//int idM;
	//float3 infoKD;
	//int indexKD;

	int stackI[MAXSTACKBVH];
	int cnode = 0;
	int deep = 0;
	//bool returnV = false;
	//return (float3)rColor.xyz;
	for(; ;)
	{
		if(IntersectionAABB(ray, bvh[cnode].aabb, /*&distance*/&nearTriangle))
		{
			if (bvh[cnode].IPLR.z == -1 && bvh[cnode].IPLR.w == -1 && (bvh[cnode].kdRMMA.w==1)) //find leaf
			{				
				/*
				bool checkShadowKDTransZB(cl_zBuffer * array, size_t *count, int idM, float distanceRay, Ray ray, float3 *info, int maxDepth, int idRoot, __global cl_kdtree *kdtree,
				__global uint * arrayIndexTriangle, __global cl_vertex *verts, __global cl_triangle *triIndex, __global float4 *dbtextures, int8 typeR)

				*/
				if(materials[bvh[cnode].kdRMMA.z].diffuse.w == 1.0f)
				{
					infoH.x = FLT_MAX;					
					indexTriangle = findIntersect( ray, &infoH, bvh[cnode].kdRMMA.y, bvh[cnode].kdRMMA.x, kdtree, arrayIndexTriangle, verts, triIndex, typeR, bvh[cnode].aabb);

					if(infoH.x < nearTriangle)
					{
						*count = 0;
						array[*count].tuv = infoH;
						array[*count].idTM.x = indexTriangle;
						array[*count].idTM.y = bvh[cnode].kdRMMA.z;
						*count += 1;
						nearTriangle = infoH.x;
						//*nearInfoH = infoH;
						//nearIndexTriangle = indexTriangle;
						//idM = bvh[cnode].kdRMMA.z;
					}
				}
				else
				{
					infoH.x = FLT_MAX;				
					if(!checkShadowKDTransZB(array, count, bvh[cnode].kdRMMA.z/*.diffuse*/, distanceRay, ray, &infoH, bvh[cnode].kdRMMA.y, bvh[cnode].kdRMMA.x, kdtree, arrayIndexTriangle, verts, triIndex, dbtextures, typeR))
					{
						//*color = (float4)(0.0f, 0.0f, 0.0f, 1.0f);
						return;
					}			
				}

				if(deep == 0)
					break;
				cnode = stackI[deep];
				deep --;
			}
			else
			{
				deep ++;
				stackI[deep] = bvh[cnode].IPLR.w;
				cnode = bvh[cnode].IPLR.z;				
			}
		}
		else
		{
			if(deep == 0)
				break;

			cnode = stackI[deep];
			deep --;
		}
	}
	//return (float3)rColor.xyz / count;
	//return false;
	//return;
	if(*count < 2 )
		return;//(float3)rColor.xyz ;
	bSortZB(array, count);
	//return (float3)rColor.xyz / count;
}

float3 RayCaster(Ray ray,
	__global cl_vertex *verts, __global cl_triangle *triIndex, uint countT,		//5 -7
	__global cl_bvh *bvh, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, __global cl_material * materials, __global float4 *dbtextures, //8 -12
	__global cl_light * lights, uint countL, __global cl_IESProfile *IESProfile, __global float *IESProfileData,  int8 typeR			//13 - 15			
#ifdef IMAGE_SUPPORT
	,__read_only image2d_t textureCL,
	sampler_t textureSamplerCL
#endif						
	)
{
	float3 rColor = (float3)(0.0f, 0.0f, 0.0f);
	float3 rrColor = (float3)(0.0f, 0.0f, 0.0f);
	cl_zBuffer array[MAXZBUFFER];
	size_t count = 0;
	float3 N;
	float3 point;
	CastRayMTZB(array, &count, ray, 10000000, verts, triIndex, bvh, kdtree, arrayIndexTriangle, materials, dbtextures, typeR);
	//return rColor;
	if(count == 0)
	{
		if(typeR.s2 != -1)			
			rrColor = (GetCubeTexColor(ray, dbtextures, materials[typeR.s2].TexWHOT.x, materials[typeR.s2].TexWHOT.y, materials[typeR.s2].TexWHOT.z, materials[typeR.s2].TexWHOT.w)).xyz;
		else
			rrColor = BCOLOR3;
	}
	else
	{
		float intenz = 1.0f;
		for(size_t i = 0; i < count; i++)
		{
			rColor = (float3)(0.0f, 0.0f, 0.0f);
			cl_material tmpMateril = materials[array[i].idTM.y];

			/*
			CastAndShadeMT(uint indexI,float3 info, Ray ray, __global cl_vertex *verts, __global cl_triangle *triIndex,
			__global cl_bvh *bvh, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, __global cl_material *materials, size_t idM,
			__global cl_light * lights, uint countL, __global cl_IESProfile *IESProfile, __global float *IESProfileData,  float3 *normal, float3 *pointI, int8 typeR, __global float4 * dbtextures
			#ifdef IMAGE_SUPPORT
			,__read_only image2d_t textureCL,
			sampler_t textureSamplerCL
			#endif						
			)*/
			float D = ( tmpMateril.diffuse.w) * (1.0f - tmpMateril.RRS.x);
			rColor += D * CastAndShadeMT(array[i].idTM.x, array[i].tuv, ray, verts, triIndex, bvh, kdtree, arrayIndexTriangle, materials, array[i].idTM.y,
				lights, countL, IESProfile, IESProfileData,  &N, &point, i, typeR, dbtextures
#ifdef IMAGE_SUPPORT
				,image2d_t textureCL,
				textureSamplerCL
#endif						
				);

			if(/*materials[array[i].idTM.y].RRS.x*/tmpMateril.RRS.x > 0.0f)
			{
				float O = ( tmpMateril.diffuse.w) * tmpMateril.RRS.x;
				Ray newRay;
				float cosPhi = dot( N, ray.direction);		
				float cosPhi1 = fabs( cosPhi );	
				float3 new_ray_direction = ray.direction + ( 2 * cosPhi1 ) * N;
				newRay.origin = ( point + new_ray_direction * SHADOW_OFFSET );
				newRay.direction = new_ray_direction;
				if(typeR.s2 != -1)
					rColor += O * materials[array[i].idTM.y].RRS.x * (GetCubeTexColor(newRay, dbtextures, materials[typeR.s2].TexWHOT.x, materials[typeR.s2].TexWHOT.y, materials[typeR.s2].TexWHOT.z, materials[typeR.s2].TexWHOT.w)).xyz;
				else
					rColor += O * materials[array[i].idTM.y].RRS.x * (BCOLOR).xyz;
			}
			if(materials[array[i].idTM.y].diffuse.w == 1.0f)
			{
				rrColor += intenz * rColor;
				break;
			}

			/*
			if(i == 0)
			break;
			*/
			if(i == count - 1)
			{
				if(typeR.s2 != -1)
					rColor += (GetCubeTexColor(ray, dbtextures, materials[typeR.s2].TexWHOT.x, materials[typeR.s2].TexWHOT.y, materials[typeR.s2].TexWHOT.z, materials[typeR.s2].TexWHOT.w)).xyz;
				else
					rColor += BCOLOR3;
			}

			float L = (1.0f - tmpMateril.diffuse.w);			
			rrColor += intenz * rColor;

			intenz *= L;
		}
	}
	return rrColor;
}

float3 RayTracer(Ray ray,
	__global cl_vertex *verts, __global cl_triangle *triIndex, uint countT,		//5 -7
	__global cl_bvh *bvh, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, __global cl_material * materials, __global float4 *dbtextures, //8 -12
	__global cl_light * lights, uint countL, __global cl_IESProfile *IESProfile, __global float *IESProfileData,  int8 typeR			//13 - 15			
#ifdef IMAGE_SUPPORT
	,__read_only image2d_t textureCL,
	sampler_t textureSamplerCL
#endif						
	)
{
	float3 rColor;
	cl_RaytracerStack stack[MAXRAYTRACERSTACK];
	float3  nearInfoH;
	int2 idTM;
	float3 N;
	float3 point;
	uint maxDepth = typeR.s3;
	//size_t depthR = 0;
	size_t i = 0;
	size_t j = 0;
	cl_material tmpMaterial;
	float n1 = 0.0f;
	float n2 = 0.0f;
	Ray newRay;
	stack[i].ray = ray;
	stack[i].depthR = 0;//depthR;
	stack[i].intz = (float2)(1.0f, 1.0f);
	stack[i].intensity = 1.0f;	
	for(; /*i < MAXRAYTRACERSTACK &&*/ stack[j].depthR < maxDepth && j <= i; )
	{		
		ray = stack[j].ray;
		//depthR = stack[j].depthR;
		idTM = CastRayMT(ray, &nearInfoH, verts, triIndex, bvh, kdtree, arrayIndexTriangle, typeR); //idTM.x =id triangle, idTM.y = id matrials
		if(idTM.x >= 0 )
		{
			rColor = CastAndShadeMT(idTM.x, nearInfoH, ray, verts, triIndex, bvh, kdtree, arrayIndexTriangle, materials, idTM.y,
				lights, countL, IESProfile, IESProfileData,  &N, &point, j, typeR, dbtextures
#ifdef IMAGE_SUPPORT
				,image2d_t textureCL,
				textureSamplerCL
#endif						
				) ;

			stack[j].color = rColor;
		}
		else
		{
			float4 tmpc;
			if(typeR.s2 != -1)
			{
				tmpc =  GetCubeTexColor(ray, dbtextures, materials[typeR.s2].TexWHOT.x, materials[typeR.s2].TexWHOT.y, materials[typeR.s2].TexWHOT.z, materials[typeR.s2].TexWHOT.w);
				//reflect background?? (float4)(0.4, 0.3, 0.6, 1.0f);
				stack[j].color = tmpc.xyz;
				stack[j].intz.x = 1.0f;
				j++;
				continue;
			}
			else // temporarily
			{
				tmpc = BCOLOR;//(float4)(0.4, 0.3, 0.6, 1.0f);
				stack[j].color = tmpc.xyz;
				stack[j].intz.x = 1.0f;
				j++;
				continue;
			}
		}

		tmpMaterial = materials[idTM.y];
		stack[j].intz.x = ( tmpMaterial.diffuse.w) * (1.0f - tmpMaterial.RRS.x);		
		float cosPhi = dot( N, ray.direction);		
		float cosPhi1 = fabs( cosPhi );		

		if (tmpMaterial.diffuse.w < 1.0f)
		{
			/*
			if(tmpMaterial.RRS.y == 1.0)
			{
			newRay.origin = ( point + ray.direction * SHADOW_OFFSET );
			newRay.direction = ray.direction;
			}
			else
			{*/
			if(cosPhi > 0)
			{
				n2 = IOR_AIR;
				n1 = tmpMaterial.RRS.y;					
			}
			else
			{
				n1 = IOR_AIR;
				n2 = tmpMaterial.RRS.y;
			}
			float divN1N2 = ( n1 / n2 );
			float sinPhi1 = sqrt(1.0f - (cosPhi1*cosPhi1) );
			//float tmp_n1_n2_sin_phi1 = divN1N2 * sinPhi1;
			float cosPhi2 = sqrt( 1.0f - (divN1N2*divN1N2) * ( 1.0f - (cosPhi1*cosPhi1) ) );
			float3 new_ray_refract = ( divN1N2 ) * ray.direction +  ( ( divN1N2 ) * cosPhi1 - cosPhi2 );

			newRay.origin = ( point + new_ray_refract * SHADOW_OFFSET );
			newRay.direction = new_ray_refract;
			//}
			i++;
			stack[i].ray = newRay;
			stack[i].depthR = stack[j].depthR + 1;
			stack[i].intz.y = (1.0f - tmpMaterial.diffuse.w) * stack[j].intz.y;
		}
		if (tmpMaterial.RRS.x > 1e-5f)
		{
			float3 new_ray_direction = ray.direction + ( 2 * cosPhi1 ) * N;

			newRay.origin = ( point + new_ray_direction * SHADOW_OFFSET );
			newRay.direction = new_ray_direction;						

			i++;
			stack[i].ray = newRay;
			stack[i].depthR = stack[j].depthR + 1;
			stack[i].intz.y = ( tmpMaterial.diffuse.w ) * tmpMaterial.RRS.x * stack[j].intz.y;
		}		
		j++;
	}

	rColor = (float3)(0.0f, 0.0f, 0.0f);
	float3 tmp;
	for (size_t q = 0; q < j; q++)
	{
		tmp = stack[q].color * stack[q].intz.x ;
		if (stack[q].intz.y > 0.0f)
			tmp *= stack[q].intz.y;
		rColor += tmp;
	}
	return rColor;
}

__kernel void RayTracing( __global float *dst, uint width, uint height, __global float *cam, __global float *pcam, //0-4
	__global cl_vertex *verts, __global cl_triangle *triIndex, uint countT,		//5 -7
	__global cl_bvh *bvh, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, __global cl_material * materials, __global float4 *dbtextures, //8 -12
	__global cl_light * lights, uint countL, __global cl_IESProfile *IESProfile, __global float *IESProfileData,  int8 typeR			//13 - 17
#ifdef IMAGE_SUPPORT
	,__read_only image2d_t textureCL, sampler_t textureSamplerCL   //18 - 19
#endif
	)
{		
	int x = get_global_id(0);
	int y = get_global_id(1);	
	uint index = (y * width + x) * 3;

	if(x >= width || y >= height )
	{
		return;
	}	

	Ray ray =  GenerateRay(x, y, width, height, cam, pcam);
	float3 rColor;
	if(typeR.s4 == 0)
	{
		rColor = RayTracer(ray, verts, triIndex, countT, bvh, kdtree, arrayIndexTriangle, materials, dbtextures, lights, countL, IESProfile, IESProfileData,  typeR			//13 - 15			
#ifdef IMAGE_SUPPORT
			,__read_only image2d_t textureCL,
			sampler_t textureSamplerCL
#endif						
			);
	}
	else
	{
		rColor = RayCaster(ray, verts, triIndex, countT, bvh, kdtree, arrayIndexTriangle, materials, dbtextures, lights, countL, IESProfile, IESProfileData,  typeR			//13 - 15			
#ifdef IMAGE_SUPPORT
			,__read_only image2d_t textureCL,
			sampler_t textureSamplerCL
#endif						
			);
	}
	dst[index] =		rColor.x;
	dst[index + 1] =	rColor.y;
	dst[index + 2] =	rColor.z;
}

/*__kernel*/ void RayTracing2( __global float *dst, uint width, uint height, __global float *cam, __global float *pcam, //0-4
	__global cl_vertex *verts, __global cl_triangle *triIndex, uint countT,		//5 -7
	__global cl_bvh *bvh, __global cl_kdtree *kdtree, __global uint * arrayIndexTriangle, __global cl_material * materials, __global float4 *dbtextures, //8 -12
	__global cl_light * lights, uint countL, __global cl_IESProfile *IESProfile, __global float *IESProfileData, uint2 *pixels, size_t sizePixel,  int8 typeR			//13 - 19			
#ifdef IMAGE_SUPPORT
	,__read_only image2d_t textureCL, sampler_t textureSamplerCL    //20 - 21
#endif						
	)
{		
	int x = get_global_id(0);
	//int y = get_global_id(1);
	if(sizePixel < x)
		return;
	uint2 currentPixel= pixels[x];
	uint index = (currentPixel.y * width + currentPixel.x) * 3;

// 	if(x >= width || y >= height )
// 	{
// 		return;
// 	}	

	Ray ray =  GenerateRay(currentPixel.x, currentPixel.y, width, height, cam, pcam);
	float3 rColor;
	if(typeR.s4 == 0)
	{
		rColor = RayTracer(ray, verts, triIndex, countT, bvh, kdtree, arrayIndexTriangle, materials, dbtextures, lights, countL, IESProfile, IESProfileData,  typeR			//13 - 15			
#ifdef IMAGE_SUPPORT
			,__read_only image2d_t textureCL,
			sampler_t textureSamplerCL
#endif						
			);
	}
	else
	{
		rColor = RayCaster(ray, verts, triIndex, countT, bvh, kdtree, arrayIndexTriangle, materials, dbtextures, lights, countL, IESProfile, IESProfileData,  typeR			//13 - 15			
#ifdef IMAGE_SUPPORT
			,__read_only image2d_t textureCL,
			sampler_t textureSamplerCL
#endif						
			);
	}
	dst[index] =		rColor.x;
	dst[index + 1] =	rColor.y;
	dst[index + 2] =	rColor.z;
}
