// 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); }