typedef struct LatticeData
{
	float4 pos;
	float4 dir;
} LatticeData;


LatticeData GetLatticeData(__global struct LatticeData *inLattice, uint num_x_lattice, uint vertStep, uint x, uint y )
{
	return inLattice[(x*vertStep) + (y*vertStep)*num_x_lattice];
}

float3 Pos( LatticeData datum )
{
	return datum.pos.xyz + datum.dir.xyz;
}

__kernel void krnl( __global struct LatticeData *inLattice, 
					__global float8 *outVertData,
					uint vertStep
					  )
{
	uint x = get_global_id(0);
	uint y = get_global_id(1);
    
	uint num_x = get_global_size(0);
	uint num_y = get_global_size(1);
	uint num_x_lattice = num_x*vertStep+1;
	uint num_y_lattice = num_y*vertStep+1;

	LatticeData latData[6];
	float3 normals[2];

	if ( (x+y)&0x1 )
	{
		// bottom-left
		latData[0] = GetLatticeData(inLattice, num_x_lattice, vertStep, x, y );
		latData[1] = GetLatticeData(inLattice, num_x_lattice, vertStep, x+1, y );
		latData[2] = GetLatticeData(inLattice, num_x_lattice, vertStep, x, y+1 );

		// top-right
		latData[3] = GetLatticeData(inLattice, num_x_lattice, vertStep, x+1, y );
		latData[4] = GetLatticeData(inLattice, num_x_lattice, vertStep, x+1, y+1 );
		latData[5] = GetLatticeData(inLattice, num_x_lattice, vertStep, x, y+1 );
	}
	else
	{
		// bottom-right
		latData[0] = GetLatticeData(inLattice, num_x_lattice, vertStep, x, y );
		latData[1] = GetLatticeData(inLattice, num_x_lattice, vertStep, x+1, y );
		latData[2] = GetLatticeData(inLattice, num_x_lattice, vertStep, x+1, y+1 );

		// top-left
		latData[3] = GetLatticeData(inLattice, num_x_lattice, vertStep, x, y );
		latData[4] = GetLatticeData(inLattice, num_x_lattice, vertStep, x+1, y+1 );
		latData[5] = GetLatticeData(inLattice, num_x_lattice, vertStep, x, y+1 );
	}

	// calc normals
	normals[0] = normalize( cross( Pos(latData[1])-Pos(latData[0]),Pos(latData[2])-Pos(latData[0]) ) );
	normals[1] = normalize( cross( Pos(latData[4])-Pos(latData[3]),Pos(latData[5])-Pos(latData[3]) ) );

	uint vboOffset = (x + y*num_x)*6;


	float2 splitDir[2] =
	{
		(float2)( (float)(x%2) * 2.0f - 1.0f, (float)(y%2) * 2.0f - 1.0f ),
		(float2)( (float)((x+1)%2) * 2.0f - 1.0f, (float)((y+1)%2) * 2.0f - 1.0f )
	};

	float8 vboData[6];
	for( uint i = 0; i < 6; ++i )
	{
		float8 vert;

		// position
		vert.s012 = Pos(latData[i]);

		// normal
		vert.s345 = normals[i/3];

		// split dir
		vert.s67 = splitDir[0];

		outVertData[vboOffset+i] = vert;
	}
}