#include "../shared.h"
#include "randomnumbers.h"
#include "program.h"

__kernel void clear( write_only image2d_t outimg )
{
	int column = get_global_id( 0 );
	int line = get_global_id( 1 );
	if ((column >= SCRWIDTH) || (line >= SCRHEIGHT)) return;
	write_imagef( outimg, (int2)(column, line), 0 );
}

__kernel void clearGrid( global unsigned int* grid )
{
	unsigned int idx = get_global_id( 0 );
	unsigned int baseIdx = idx * (BALLSPERCELL + 1);
	grid[baseIdx] = 0;
}

__kernel void fillGrid( global float* balls, global unsigned int* grid )
{
	unsigned int idx = get_global_id( 0 );
	unsigned int gx = balls[idx * 4 + 0] * GRIDX;
	unsigned int gy = balls[idx * 4 + 1] * GRIDY;
	if ((gx < 0) || (gy < 0) || (gx >= GRIDX) || (gy >= GRIDY)) return;
	unsigned int baseIdx = (gx + gy * GRIDX) * (BALLSPERCELL + 1);
	if (grid[baseIdx] < BALLSPERCELL)
	{ 
		unsigned int count = atomic_inc( grid + baseIdx );
		if (count < BALLSPERCELL) grid[baseIdx + count + 1] = idx;
	}
	else balls[idx * 4 + 1] = balls[idx * 4 + 3] = 0.1;
}

__kernel void simulate1( global float* balls, float gx, float gy )
{
	int idx = get_global_id( 0 );
	float2 prevPos = { balls[idx * 4 + 0], balls[idx * 4 + 1] };
	float2 delta = { balls[idx * 4 + 0] - balls[idx * 4 + 2] + gx,
					 balls[idx * 4 + 1] - balls[idx * 4 + 3] + gy };
	float speed = length( delta );
	if (speed > 0.01f) delta = 0.01f * normalize( delta );
	balls[idx * 4 + 0] += delta.x * 0.95;
	balls[idx * 4 + 1] += delta.y * 0.95;
	balls[idx * 4 + 2] = prevPos.x;
	balls[idx * 4 + 3] = prevPos.y;
}

__kernel void simulate2( global float* balls, global float* balls2, global unsigned int* grid )
{
	int idx = get_global_id( 0 );
	int baseIdx = idx * (BALLSPERCELL + 1);
	int count = min( (unsigned int)BALLSPERCELL, grid[baseIdx] );
	if (count == 0) return;
	int gx = idx % GRIDX;
	int gy = idx / GRIDX;
	int gx1 = max( 0, gx - 1 ), gx2 = min( GRIDX - 1, gx + 1 );
	int gy1 = max( 0, gy - 1 ), gy2 = min( GRIDY - 1, gy + 1 );
	for( int i = 0; i < count; i++ )
	{
		// get active ball
		int idx1 = grid[baseIdx + i + 1];
		float2 ball1Pos = { balls[idx1 * 4 + 0], balls[idx1 * 4 + 1] };
		// evade other balls
		for( int y = gy1; y <= gy2; y++ ) for( int x = gx1; x <= gx2; x++ )
		{
			int baseIdx = (x + y * GRIDX) * (BALLSPERCELL + 1);
			int count2 = min( (unsigned int)BALLSPERCELL, grid[baseIdx] );
			for( int j = 0; j < count2; j++ )
			{ 
				int idx2 = grid[baseIdx + j + 1];
				if (idx2 != idx1)
				{
					float2 ball2Pos = { balls2[idx2 * 4 + 0], balls2[idx2 * 4 + 1] };
					float2 delta = ball2Pos - ball1Pos;
					float dist = sqrt( delta.x * delta.x + delta.y * delta.y );
					if (dist < 0.01f)
					{
						float2 push = (0.01f - dist) * 0.25f * normalize( delta );
						ball1Pos -= push;
					}
				}
			}
		}
		// respect boundaries
		float jitter = (float)idx1 * 0.000005f;
		ball1Pos.x = min( 0.95f + jitter, max( 0.05f - jitter, ball1Pos.x ) );
		ball1Pos.y = min( 0.95f + jitter, max( 0.05f - jitter, ball1Pos.y ) );
		// update ball position
		balls[idx1 * 4 + 0] = ball1Pos.x;
		balls[idx1 * 4 + 1] = ball1Pos.y;
	}	
}

__kernel void simulateX( global float* balls, global read_only float* balls2, global read_only unsigned int* grid, float mx, float my )
{
	int ballIdx1 = get_global_id( 0 );
	float2 ball1Pos = { balls[ballIdx1 * 4 + 0], balls[ballIdx1 * 4 + 1] };
	int gx = ball1Pos.x * GRIDX, gx1 = max( 0.f, (ball1Pos.x - BALLRADIUS) * GRIDX ), gx2 = min( GRIDX - 1.f, (ball1Pos.x + BALLRADIUS) * GRIDX );
	int gy = ball1Pos.y * GRIDY, gy1 = max( 0.f, (ball1Pos.y - BALLRADIUS) * GRIDY ), gy2 = min( GRIDY - 1.f, (ball1Pos.y + BALLRADIUS) * GRIDY );
	if ((gx < 0) || (gy < 0) || (gx >= GRIDX) || (gy >= GRIDY)) return;
	float2 force = { 0, 0 };
	for( int y = gy1; y <= gy2; y++ ) for( int x = gx1; x <= gx2; x++ )
	{
		int baseIdx = (x + y * GRIDX) * (BALLSPERCELL + 1);
		int count2 = min( (unsigned int)BALLSPERCELL, grid[baseIdx] );
		for( int j = 0; j < count2; j++ )
		{ 
			int ballIdx2 = grid[baseIdx + j + 1];
			if (ballIdx2 != ballIdx1)
			{
				float2 ball2Pos = { balls2[ballIdx2 * 4 + 0], balls2[ballIdx2 * 4 + 1] };
				float2 delta = ball2Pos - ball1Pos;
				float dist = delta.x * delta.x + delta.y * delta.y;
				if (dist < (BALLRADIUS * BALLRADIUS))
				{
					dist = native_sqrt( dist );
					float2 push = (BALLRADIUS - dist) * PUSHFORCE * (delta / dist);
					force -= push;
				}
			}
		}
	}
	// evade mouse
	float2 delta = { ball1Pos.x - mx, ball1Pos.y - my };
	float dist = delta.x * delta.x + delta.y * delta.y;
	if (dist < (0.05f * 0.05f))
	{
		dist = native_sqrt( dist );
		float2 push = (0.05f - dist) * PUSHFORCE * (delta / dist);
		force += push;
	}
	// limit impulse
	float speed = length( force );
	if (speed > MAXFORCE) force = MAXFORCE * normalize( force );
	ball1Pos += force;
	// respect boundaries
	float jitter = (float)ballIdx1 * 0.000001f;
	ball1Pos.x = min( 0.98f + jitter, max( 0.02f - jitter, ball1Pos.x ) );
	ball1Pos.y = min( 0.98f + jitter, max( 0.02f - jitter, ball1Pos.y ) );
	// update ball position
	balls[ballIdx1 * 4 + 0] = ball1Pos.x;
	balls[ballIdx1 * 4 + 1] = ball1Pos.y;
}

__kernel void render( write_only image2d_t outimg, global read_only float* balls, global read_only unsigned int* grid )
{
	int column = get_global_id( 0 );
	int line = get_global_id( 1 );
	if ((column >= SCRWIDTH) || (line >= SCRHEIGHT)) return;
	float2 uv = { (float)column / SCRWIDTH, (float)line / SCRHEIGHT };
	// draw balls using grid
	int gx = uv.x * GRIDX;
	int gy = uv.y * GRIDY;
	int gx1 = max( 0, gx - 1 ), gx2 = min( GRIDX - 1, gx + 1 );
	int gy1 = max( 0, gy - 1 ), gy2 = min( GRIDY - 1, gy + 1 );
	for( int y = gy1; y <= gy2; y++ ) for( int x = gx1; x <= gx2; x++ )
	{
		unsigned int baseIdx = (x + y * GRIDX) * (BALLSPERCELL + 1);
		unsigned int count = min( (unsigned int)BALLSPERCELL, grid[baseIdx] );
		for( int i = 0; i < count; i++ )
		{ 
			unsigned int idx = grid[baseIdx + i + 1];
			float2 delta = { balls[idx * 4] - uv.x, balls[idx * 4 + 1] - uv.y };
			float dist = delta.x * delta.x + delta.y * delta.y;
			if (dist < (BALLRENDERSIZE * BALLRENDERSIZE))
		{
				write_imagef( outimg, (int2)(column, (SCRHEIGHT - 1) - line), (float4)(1,0,0,1) );
			break;
		}
	}
	}
}

__kernel void render2( write_only image2d_t outimg, global float* balls, global unsigned int* grid )
{
	const uint idx = get_global_id( 0 );
	const uint baseIdx = idx * (BALLSPERCELL + 1);
	const uint count = min( (uint)BALLSPERCELL, grid[baseIdx] );
	if (count == 0) return;
	for( int i = 0; i < count; i++ )
	{
		const uint ballIdx = grid[baseIdx + i + 1];
		const float bx = balls[ballIdx * 4], by = balls[ballIdx * 4 + 1];
		const uint x1 = (uint)max( 0.0, (bx - BALLRENDERSIZE) * SCRWIDTH );
		const uint y1 = (uint)max( 0.0, (by - BALLRENDERSIZE) * SCRHEIGHT );
		const uint x2 = (uint)min( SCRWIDTH - 1., (bx + BALLRENDERSIZE) * SCRWIDTH + 1 );
		const uint y2 = (uint)min( SCRHEIGHT - 1., (by + BALLRENDERSIZE) * SCRHEIGHT + 1 );
		for( uint y = y1; y < y2; y++ ) for( uint x = x1; x < x2; x++ )
		{ 
			const float2 uv = { (float)x / SCRWIDTH, (float)y / SCRHEIGHT };
			const float2 delta = { bx - uv.x, by - uv.y };
			if ((delta.x * delta.x + delta.y * delta.y) < (BALLRENDERSIZE * BALLRENDERSIZE))
				write_imagef( outimg, (int2)(x, (SCRHEIGHT - 1) - y), (float4)(1,0,0,1) );
		}
	}
}