 		} __debugsync(); if (idx < 16) {
 #endif
 		
-		sortColors(dps, colors, xrefs);
-	}
+		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;

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

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

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