unsigned char fx_plasma_run(__global unsigned int* fsin1, __global unsigned int* fsin2, __global unsigned int* fsin3, int x, int y, int z, int r, int t)
{
	unsigned char c = fsin1[x] + fsin2[y] + fsin3[z] + fsin1[x+y+t] + fsin2[x+z+t+t] + fsin3[y+z+t+t+t];

	return c;
}

__kernel void fx_plasma(__global unsigned int* volume, __global unsigned int* fsin1, __global unsigned int* fsin2, __global unsigned int* fsin3, int vol_width, int ptime)
{
	int x = get_global_id(0) << 2;
	int y = get_global_id(1);
	int z = get_global_id(2);

	int t = ptime >> 6;
	int r = vol_width >> 1;

	unsigned char c3 = fx_plasma_run(fsin1, fsin2, fsin3, x, y, z, r, t);
	unsigned char c2 = fx_plasma_run(fsin1, fsin2, fsin3, x+1, y, z, r, t);
	unsigned char c1 = fx_plasma_run(fsin1, fsin2, fsin3, x+2, y, z, r, t);
	unsigned char c0 = fx_plasma_run(fsin1, fsin2, fsin3, x+3, y, z, r, t);

	int addr = (z * vol_width * vol_width + y * vol_width + x) >> 2;
	volume[addr] = (c0 << 24) | (c1 << 16) | (c2 << 8) | c3;
}
