Example #1
0
// Ugly quick solution, since atomic max does not work for floats
__kernel void CalculateMaxAtomic(volatile __global int* max_value,
	                             __global const float* Volume,
								 __global const float* Mask,
								 __private int DATA_W,
								 __private int DATA_H,
								 __private int DATA_D)
{
	int x = get_global_id(0);	
	int y = get_global_id(1);
	int z = get_global_id(2);

	if ( x >= DATA_W || y >= DATA_H || z >= DATA_D )
		return;

	if ( Mask[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] != 1.0f )
		return;

	int value = (int)(Volume[Calculate3DIndex(x,y,z,DATA_W,DATA_H)] * 10000.0f);
	atomic_max(max_value, value);
}
__kernel void CalculateLargestCluster(__global unsigned int* Cluster_Sizes,
								      volatile global unsigned int* Largest_Cluster,
   						  	  	      __private int DATA_W,
									  __private int DATA_H,
									  __private int DATA_D)
{
	int x = get_global_id(0);
	int y = get_global_id(1);
	int z = get_global_id(2);

	if (x >= DATA_W || y >= DATA_H || z >= DATA_D)
		return;

	unsigned int cluster_size = Cluster_Sizes[Calculate3DIndex(x,y,z,DATA_W,DATA_H)];

	// Most cluster size counters are zero, so avoid running atomic max for those
	if (cluster_size == 0)
		return;

	atomic_max(Largest_Cluster,cluster_size);
}