OPENCL : CALCUL ET AFFICHAGE DU GRADIENT COULEUR D'UNE WEBCAM (RGB24)

//512
//OpenCL

const sampler_t samp = CLK_ADDRESS_CLAMP_TO_EDGE ;
const sampler_t sampi = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;


__kernel void load (__read_only image2d_t in0, __write_only image2d_t out, __local uchar *ld, __constant char *gld, int spc, __read_only image2d_t cfc)
{
	int2 po = (int2)(get_global_id(0), get_global_id(1));
	int dcy = 4+get_local_size (0);
	int dcl = dcy*3;

	const int pc = (dcy*(2+get_local_id (1))+2+get_local_id (0));
	char sx=1,sy=1;
	if (get_local_id(0)<4) sx=-1;
	if (get_local_id(1)<6) sy=-1;

	float4 lv;
	int pp;
	char cas=gld[(get_local_id(1)*get_local_size(0)+get_local_id(0))];

	switch (cas)
	{
		case 11: 
			{
				lv = read_imagef (in0,samp,po+(int2)(5*sx,5)); pp = 3*(pc+5*sx+5*dcy);
				ld[pp] = (lv.x*255);
				ld[pp+1]=(lv.y*255);
				ld[pp+2]=(lv.z*255);
			}
		break;
		case 10: 
			{
				lv = read_imagef (in0,samp,po+(int2)(4*sx,-5)); pp = 3*(pc+4*sx-5*dcy);
				ld[pp] = (lv.x*255);
				ld[pp+1]=(lv.y*255);
				ld[pp+2]=(lv.z*255);
			}
		break;
		case  9: 
			{
				lv = read_imagef (in0,samp,po+(int2)(3,5*sy)); pp=3*(pc+5*sy*dcy+3);
				ld[pp] = (lv.x*255);
				ld[pp+1]=(lv.y*255);
				ld[pp+2]=(lv.z*255);
			}
		break;
		case  8: 
			{
				lv = read_imagef (in0,samp,po+(int2)(-3,4*sy)); pp = 3*(pc+4*sy*dcy-3);		 
				ld[pp] = (lv.x*255);
				ld[pp+1]=(lv.y*255);
				ld[pp+2]=(lv.z*255);
			}
			break;
		case  7: 
			{
				lv = read_imagef (in0,samp,po+(int2)(3*sx,0));  pp = 3*(pc+3*sx);			
				ld[pp] = (lv.x*255);
				ld[pp+1]=(lv.y*255);
				ld[pp+2]=(lv.z*255);
			}
		break;
		case  6: 
			{
				lv = read_imagef (in0,samp,po+(int2)(0,3*sy));  pp = 3*(pc+3*sy*dcy);		 
				ld[pp] = (lv.x*255);
				ld[pp+1]=(lv.y*255);
				ld[pp+2]=(lv.z*255);
			}
		break;
		case  5:
			{
				lv = read_imagef (in0,samp,po+(int2)(3*sx,3*sy)); pp=3*(pc+3*(sx+sy*dcy)); 
				ld[pp] = (lv.x*255);
				ld[pp+1]=(lv.y*255);
				ld[pp+2]=(lv.z*255);
			}
		break;
		case  4: 
			{
				lv = read_imagef (in0,samp,po+(int2)(2*sx,2*sy)); pp=3*(pc+2*(sx+sy*dcy));
				ld[pp] = (lv.x*255);
				ld[pp+1]=(lv.y*255);
				ld[pp+2]=(lv.z*255);
			}
		break;
		case  3: 
			{
				lv = read_imagef (in0,samp,po+(int2)(sx,sy)); pp=3*(pc+sx+sy*dcy);
				ld[pp] = (lv.x*255);
				ld[pp+1]=(lv.y*255);
				ld[pp+2]=(lv.z*255);
			}
		break;
		case  2: 
			{
				lv = read_imagef (in0,samp,po+(int2)(0,sy));  pp=3*(pc+sy*dcy);
				ld[pp] = (lv.x*255);
				ld[pp+1]=(lv.y*255);
				ld[pp+2]=(lv.z*255);
			}
		break;
		case 1: 
			{
				lv = read_imagef (in0,samp,po+(int2)(sx,0)); pp = 3*(pc+sx);
				ld[pp] = (lv.x*255);
				ld[pp+1]=(lv.y*255);
				ld[pp+2]=(lv.z*255);
			}
		break;
		default:
		break;
	}
	lv = read_imagef (in0,samp,po)*255;
	int r = lv.x,v = lv.y,b = lv.z;
	ld[3*pc+0] = r;
	ld[3*pc+1] = v;
	ld[3*pc+2] = b;

	barrier (CLK_LOCAL_MEM_FENCE);

	float mr=0,mv=0,mb=0;
	float som=0,cf=1;
	int rr,vv,bb;

	for (int j=-2; j<3; j++)
	{
		for (int i=-2; i<3; i++)
		{
			pp = j*dcl+3*(pc+i);
			rr=ld[pp];
			vv=ld[pp+1];
			bb=ld[pp+2];

			cf=read_imagef (cfc,samp,(int2)((abs(rr-r)+abs(vv-v)+abs(bb-b)),0)).x;//lmv[vp+i+5*j]
			
			mr+=rr*cf;
			mv+=vv*cf;
			mb+=bb*cf;
			som+=cf;

		}
	}

	if (spc) {mr = lv.x*som;mv = lv.y*som;mb = lv.z*som;}

	write_imagef (out, (int2)(get_global_id(0), 480-get_global_id(1)), (float4)(mb,mv,mr,0)/som/255);

	barrier (CLK_LOCAL_MEM_FENCE);
}

__kernel void gradient (__read_only image2d_t in, __write_only image2d_t grd)
{
	int pw = get_global_id(0);
	int ph = get_global_id(1);

	float4 th = read_imagef (in,samp,(int2)(pw,ph));

	float4 rsx = th-read_imagef (in,samp,(int2)(pw-1,ph));
	float4 rsy = th-read_imagef (in,samp,(int2)(pw,ph-1));
	float4 rtx = th-read_imagef (in,samp,(int2)(pw+1,ph));
	float4 rty = th-read_imagef (in,samp,(int2)(pw,ph+1));
	
	float k1 = fast_length (rsx);
	float l1 = fast_length (rsy);
	float k2 = fast_length (rtx);
	float l2 = fast_length (rty);
	
	float4 gr = (float4)0;
	
	gr.z = gr.x += k1>k2?
	fast_length(rsx+read_imagef (in,samp,(int2)(pw,ph+1))-read_imagef (in,samp,(int2)(pw-1,ph+1))+read_imagef (in,samp,(int2)(pw,ph-1))-read_imagef (in,samp,(int2)(pw-1,ph-1))):
	fast_length(rtx+read_imagef (in,samp,(int2)(pw,ph+1))-read_imagef (in,samp,(int2)(pw+1,ph+1))+read_imagef (in,samp,(int2)(pw,ph-1))-read_imagef (in,samp,(int2)(pw+1,ph-1)));
	
	gr.z += gr.y += l1>l2?
	fast_length(rsy+read_imagef (in,samp,(int2)(pw+1,ph))-read_imagef (in,samp,(int2)(pw+1,ph-1))+read_imagef (in,samp,(int2)(pw-1,ph))-read_imagef (in,samp,(int2)(pw-1,ph-1))):
	fast_length(rty+read_imagef (in,samp,(int2)(pw+1,ph))-read_imagef (in,samp,(int2)(pw+1,ph+1))+read_imagef (in,samp,(int2)(pw-1,ph))-read_imagef (in,samp,(int2)(pw-1,ph+1)));

	write_imagef (grd, (int2)(pw,ph), gr);
}

__kernel void visu (__read_only image2d_t grd, __write_only image2d_t out, int nrm, __read_only image2d_t pcl)
{
	int pw = get_global_id(0);
	int ph = get_global_id(1);
	float amp,lbd=1;
	
	float4 dx = read_imagef (grd,samp,(int2)(pw,ph));

	if (nrm==1)
	{
		dx.z *= 100;
		float dz = fabs(dx.z-(int)dx.z);
		amp = ((int)dx.z)%896+dz;
		amp = log2(2+amp/4)*192;
		amp -= (int)(amp/896)*896;
	}
	else
	{
		if (nrm==2)
		{
			//dx.z *= 100;
			lbd = sqrt(dx.z);
			amp = atan (dx.y/dx.x)/3.1415926535*896.0*2;
		}
		else
		{
			amp = atan (dx.y/dx.x)/3.1415926535*896.0*2;
		}
	}

	if (lbd>1) lbd=1;

	write_imagef (out, (int2)(pw,ph), read_imagef (pcl,samp,(int2)(amp,0))*lbd);
}