                 bestPermutation ^= (~bestPermutation >> 1) & 0x55555555;    // Flip indices.
             }
         }
-
+		
         error = evalPermutation4(colors, permutation, &start, &end);
         
         if (error < bestError)

	// Only write the result of the winner thread.
	if (idx == indices[0])
	{
		if (bestStart == bestEnd)
		{
			bestPermutation = 0;
		}
	
		// Reorder permutation.
		uint perm = 0;
		for(int i = 0; i < 16; i++)
		{
			int ref = xrefs[i];
			perm |= ((bestPermutation >> (2 * ref)) & 3) << (2 * i);
		}
		
		// Write endpoints. (bestStart, bestEnd)
		result[bid].x = (bestEnd << 16) | bestStart;
		
		// Write palette indices (permutation).
		result[bid].y = perm;
	}
}


		} __debugsync(); if (idx < 16) {
#endif
		
		sortColors(dps, xrefs);
		
		float3 tmp = colors[idx];
		colors[xrefs[idx]] = tmp;
	}
}

__device__ void loadColorBlock(const uint * image, float3 colors[16], float weights[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.
		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, xrefs);
		
		float3 tmp = colors[idx];
		colors[xrefs[idx]] = tmp;
		
		float w = weights[idx];
		weights[xrefs[idx]] = tmp;
	}
}


__device__ void saveBlockDXT1(ushort start, ushort end, uint permutation, int xrefs[16])
{
	const int bid = blockIdx.x;
	const int idx = threadIdx.x;

	if (start == end)
	{
		permutation = 0;
	}
	
	// Reorder permutation.
	uint indices = 0;
	for(int i = 0; i < 16; i++)
	{
		int ref = xrefs[i];
		indices |= ((permutation >> (2 * ref)) & 3) << (2 * i);
	}
	
	// Write endpoints.
	result[bid].x = (end << 16) | start;
	
	// Write palette indices.
	result[bid].y = indices;
}


////////////////////////////////////////////////////////////////////////////////
// Compress color block
////////////////////////////////////////////////////////////////////////////////
__global__ void compress(const uint * permutations, const uint * image, uint2 * result)
{
	const int bid = blockIdx.x;
	const int idx = threadIdx.x;
	
	__shared__ float3 colors[16];
	__shared__ int xrefs[16];
	
	loadColorBlock(image, colors, xrefs);
	
	ushort bestStart, bestEnd;
	uint bestPermutation;
