 
 
 ////////////////////////////////////////////////////////////////////////////////
-// Compress color block
-////////////////////////////////////////////////////////////////////////////////
-__global__ void compress(const uint * permutations, const uint * image, uint2 * result)
+// 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__ float3 colors[16];
+
 	__shared__ float dps[16];
-	__shared__ int xrefs[16];
-	
+
 	if (idx < 16)
 	{
 		// Read color and copy to shared mem.

            bestPermutation = permutation;
            bestStart = start;
            bestEnd = end;

			if (bestStart < bestEnd)
			{
				swap(bestEnd, bestStart);

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

