 	// 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;
+		saveBlockDXT1(bestStart, bestEnd, bestPermutation, xrefs);
 	}
 }
 

	
	__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;

		}
	}

	// unroll last 6 iterations
	if (idx <= 32)
	{
		if (errors[idx + 32] < errors[idx]) {
