kernel void particle(global float4 *in, global float4 *pos)
{
	int i = get_global_id(0) * 2;
	bool change = true;
	const float t = 200000.0f;

	float4 v[8];
	float  len[8];
	for(int j=0; j < 8; j++)
	{
		v[j] = pos[j] - in[i];
		if (j&1) v[j] = -v[j];
		len[j] = dot(v[j],v[j]);
		
		if (len[j] < 1000.0f) { change = false; break; }
	}
 	
	if (change)
	{
		float4 tmp = (float4)(0,0,0,0);
		
		for(int j=0; j < 8; j++)
			tmp += fast_normalize(v[j]) / len[j];

		tmp *= t;
		in[i] = in[i] + tmp/100.0f;
		in[i+1] = in[i] + tmp/50.0f;
	}

	in[i  ].z -= 5.0f;
	in[i+1].z -= 5.0f;
	
	if( in[i].z < -200 )
	{
		in[i].z += 400.0f;
		in[i+1].z += 400.0f;
	}
}