
#define NUM_BANKS 16
#define NUM_ELEMENTS (NUM_BANKS+1)

void generateSkipList(int idx, __local int* list, __local int* skipList, int invalidValue){
	int dim = get_local_size(0);
	// default value for debugging reasons
	skipList[idx] = 0;
	// sync threads to ensure that
	// the input list is complete
	barrier(CLK_LOCAL_MEM_FENCE);
	// the start element in the list
	// {0, NUM_ELEMENTS, NUM_ELEMENTS*2, ...}
	int cur = NUM_ELEMENTS*idx, last = cur;
	// find the next valid element
	for(int i=0; cur<dim; cur++, i++){
		// element is valid
		if(list[cur] != invalidValue){
			// store the relative position
			// to this element in the skip list
			skipList[last] = cur-last;
			// if the element is invalid sign it
			// so it can be skipt (this happens when
			// the start element of the loop is invalid)
			if(list[last] == invalidValue)
				skipList[last] = last-cur;
			// stop when an element is found in
			// the searching area of another thread
			if(NUM_ELEMENTS < i)
				break;
			// if the current element still lies
			// within the searchin area we continue
			last = cur;
		}
	}
	// no valid element found in the whole list
	if(cur == dim)
		// skip the whole list
		skipList[last] = cur-last;
	// make sure every thread has finished
	// writing to the skip list so it can be stored
	barrier(CLK_LOCAL_MEM_FENCE);
}

__kernel void cullSphereSpheres(float4 sphere, __global float4* spheres, __global int* result, __local int* list, __local int* skipList) {
	// position in the global memory
	int globalIdx = get_global_id(0);
	int localIdx = get_local_id(0);
	// cache global memory value
	float4 instance = spheres[globalIdx];
	float R = sphere.w + instance.w;
	instance -= sphere;
	// set result to prevent 'else'
	list[localIdx] = 0;
	// is there an intersection
	if(    (instance.w > 0.f)
		&& (instance.x*instance.x + instance.y*instance.y + instance.z*instance.z < R*R))
		list[localIdx] = 1;
	// generate skip list
	generateSkipList(localIdx, list, skipList, 0);
	// write result
	result[globalIdx] = skipList[localIdx];
}

__kernel void cullFrustumSpheres(float4 left, float4 right, float4 top, float4 bottom, float4 far, float4 near, __global float4* spheres, __global int* result, __local int* list, __local int* skipList) {
	// position in the global memory
	int globalIdx = get_global_id(0);
	int localIdx = get_local_id(0);
	// cache global memory value
	float4 instance = spheres[globalIdx];
	// set result to prevent 'else'
	list[localIdx] = 0;
	// is there an intersection
	if(    (instance.w > 0.f)
		&& (left.x  *instance.x + left.y  *instance.y + left.z  *instance.z + left.w   + instance.w > 0.f)
		&& (right.x *instance.x + right.y *instance.y + right.z *instance.z + right.w  + instance.w > 0.f)
		&& (top.x   *instance.x + top.y   *instance.y + top.z   *instance.z + top.w    + instance.w > 0.f)
		&& (bottom.x*instance.x + bottom.y*instance.y + bottom.z*instance.z + bottom.w + instance.w > 0.f)
		&& (far.x   *instance.x + far.y   *instance.y + far.z   *instance.z + far.w    + instance.w > 0.f)
		&& (near.x  *instance.x + near.y  *instance.y + near.z  *instance.z + near.w   + instance.w > 0.f))
		list[localIdx] = 1;
	// generate skip list
	generateSkipList(localIdx, list, skipList, 0);
	// write result
	result[globalIdx] = skipList[localIdx];
}
