 ////////////////////////////////////////////////////////////////////////////////
 // Sort colors
 ////////////////////////////////////////////////////////////////////////////////
-__device__ void sortColors(float * values, float3 * colors, int * cmp)
+__device__ void sortColors(float * values, int * cmp)
 {
 	int tid = threadIdx.x;
 

	
	__shared__ float3 colors[16];
	__shared__ float weights[16];
	__shared__ float dps[16];
	__shared__ int xrefs[16];
	
	if (idx < 16)
	{
		// Read color and copy to shared mem.
		uint c = image[(bid) * 16 + idx];
		
		colors[idx].z = ((c >> 0) & 0xFF) * (1.0f / 255.0f);
		colors[idx].y = ((c >> 8) & 0xFF) * (1.0f / 255.0f);
		colors[idx].x = ((c >> 16) & 0xFF) * (1.0f / 255.0f);
		weights[idx] = ((c >> 24) & 0xFF) * (1.0f / 255.0f);
		
		// No need to synchronize, 16 < warp size.
#if __DEVICE_EMULATION__
		} __debugsync(); if (idx < 16) {
#endif
		
		// Sort colors along the best fit line.
		float3 axis = bestFitLine(colors);
		
		dps[idx] = dot(colors[idx], axis);
		
#if __DEVICE_EMULATION__
		} __debugsync(); if (idx < 16) {
#endif
		
		sortColors(dps, colors, weights, xrefs);
	}
	
	ushort bestStart, bestEnd;
	uint bestPermutation;
	float bestError = FLT_MAX;



////////////////////////////////////////////////////////////////////////////////
// Load color block to shared mem
////////////////////////////////////////////////////////////////////////////////
__device__ void loadColorBlock(const uint * image, float3 colors[16], int xrefs[16])
{
	const int bid = blockIdx.x;
	const int idx = threadIdx.x;

	__shared__ float dps[16];

	if (idx < 16)
	{
		// Read color and copy to shared mem.
