 
 __constant__ float3 kColorMetric = { 1.0f, 1.0f, 1.0f };
 
-
-////////////////////////////////////////////////////////////////////////////////
-// Round color to RGB565 and expand
-////////////////////////////////////////////////////////////////////////////////
-inline __device__ float3 roundAndExpand(float3 v, ushort * w)
-{
-    v.x = rintf(__saturatef(v.x) * 31.0f);
-    v.y = rintf(__saturatef(v.y) * 63.0f);
-    v.z = rintf(__saturatef(v.z) * 31.0f);
-    *w = ((ushort)v.x << 11) | ((ushort)v.y << 5) | (ushort)v.z;
-    v.x *= 0.03227752766457f; // approximate integer bit expansion.
-    v.y *= 0.01583151765563f;
-    v.z *= 0.03227752766457f;
-    return v;
-}
-
-
-////////////////////////////////////////////////////////////////////////////////
-// Evaluate permutations
-////////////////////////////////////////////////////////////////////////////////
-static __device__ float evalPermutation4(const float3 * colors, uint permutation, ushort * start, ushort * end)
-{
-    // Compute endpoints using least squares.
-    float alpha2_sum = 0.0f;
-    float beta2_sum = 0.0f;
-    float alphabeta_sum = 0.0f;
-    float3 alphax_sum = make_float3(0.0f, 0.0f, 0.0f);
-    float3 betax_sum = make_float3(0.0f, 0.0f, 0.0f);
-
-    // Compute alpha & beta for this permutation.
-    for (int i = 0; i < 16; i++)
-    {
-        const uint bits = permutation >> (2*i);
-
-        float beta = (bits & 1);
-        if (bits & 2) beta = (1 + beta) / 3.0f;
-        float alpha = 1.0f - beta;
-    
-        alpha2_sum += alpha * alpha;
-        beta2_sum += beta * beta;
-        alphabeta_sum += alpha * beta;
-        alphax_sum += alpha * colors[i];
-        betax_sum += beta * colors[i];
-    }
-
-    const float factor = 1.0f / (alpha2_sum * beta2_sum - alphabeta_sum * alphabeta_sum);
-
-    float3 a = (alphax_sum * beta2_sum - betax_sum * alphabeta_sum) * factor;
-    float3 b = (betax_sum * alpha2_sum - alphax_sum * alphabeta_sum) * factor;
-    
-    // Round a, b to the closest 5-6-5 color and expand...
-    a = roundAndExpand(a, start);
-    b = roundAndExpand(b, end);
-
-    // compute the error
-    float3 e = a * a * alpha2_sum + b * b * beta2_sum + 2.0f * (a * b * alphabeta_sum - a * alphax_sum - b * betax_sum);
-
-    return dot(e, kColorMetric);
-}
-
-
-static __device__ float evalPermutation3(const float3 * colors, uint permutation, ushort * start, ushort * end)
-{
-    // Compute endpoints using least squares.
-    float alpha2_sum = 0.0f;
-    float beta2_sum = 0.0f;
-    float alphabeta_sum = 0.0f;
-    float3 alphax_sum = make_float3(0.0f, 0.0f, 0.0f);
-    float3 betax_sum = make_float3(0.0f, 0.0f, 0.0f);
-
-    // Compute alpha & beta for this permutation.
-    for (int i = 0; i < 16; i++)
-    {
-        const uint bits = permutation >> (2*i);
-
-        float beta = (bits & 1);
-        if (bits & 2) beta = 0.5f;
-        float alpha = 1.0f - beta;
-    
-        alpha2_sum += alpha * alpha;
-        beta2_sum += beta * beta;
-        alphabeta_sum += alpha * beta;
-        alphax_sum += alpha * colors[i];
-        betax_sum += beta * colors[i];
-    }
-
-    const float factor = 1.0f / (alpha2_sum * beta2_sum - alphabeta_sum * alphabeta_sum);
-
-    float3 a = (alphax_sum * beta2_sum - betax_sum * alphabeta_sum) * factor;
-    float3 b = (betax_sum * alpha2_sum - alphax_sum * alphabeta_sum) * factor;
-    
-    // Round a, b to the closest 5-6-5 color and expand...
-    a = roundAndExpand(a, start);
-    b = roundAndExpand(b, end);
-
-    // compute the error
-    float3 e = a * a * alpha2_sum + b * b * beta2_sum + 2.0f * (a * b * alphabeta_sum - a * alphax_sum - b * betax_sum);
-
-    return dot(e, kColorMetric);
-}
-
-static __device__ float evalPermutation4(const float3 * colors, const float * weights, uint permutation, ushort * start, ushort * end)
-{
-    // Compute endpoints using least squares.
-    float alpha2_sum = 0.0f;
-    float beta2_sum = 0.0f;
-    float alphabeta_sum = 0.0f;
-    float3 alphax_sum = make_float3(0.0f, 0.0f, 0.0f);
-    float3 betax_sum = make_float3(0.0f, 0.0f, 0.0f);
-
-    // Compute alpha & beta for this permutation.
-    for (int i = 0; i < 16; i++)
-    {
-        const uint bits = permutation >> (2*i);
-		
-        float beta = (bits & 1);
-        if (bits & 2) beta = (1 + beta) / 3.0f;
-        float alpha = 1.0f - beta;
-		
-		alpha2_sum += alpha * alpha * weights[i];
-        beta2_sum += beta * beta * weights[i];
-        alphabeta_sum += alpha * beta * weights[i];
-        alphax_sum += alpha * colors[i] * weights[i];
-        betax_sum += beta * colors[i] * weights[i];
-    }
-
-    const float factor = 1.0f / (alpha2_sum * beta2_sum - alphabeta_sum * alphabeta_sum);
-
-    float3 a = (alphax_sum * beta2_sum - betax_sum * alphabeta_sum) * factor;
-    float3 b = (betax_sum * alpha2_sum - alphax_sum * alphabeta_sum) * factor;
-    
-    // Round a, b to the closest 5-6-5 color and expand...
-    a = roundAndExpand(a, start);
-    b = roundAndExpand(b, end);
-
-    // compute the error
-    float3 e = a * a * alpha2_sum + b * b * beta2_sum + 2.0f * (a * b * alphabeta_sum - a * alphax_sum - b * betax_sum);
-
-    return dot(e, kColorMetric);
-}
-
-
-static __device__ float evalPermutation3(const float3 * colors, const float * weights, uint permutation, ushort * start, ushort * end)
-{
-    // Compute endpoints using least squares.
-    float alpha2_sum = 0.0f;
-    float beta2_sum = 0.0f;
-    float alphabeta_sum = 0.0f;
-    float3 alphax_sum = make_float3(0.0f, 0.0f, 0.0f);
-    float3 betax_sum = make_float3(0.0f, 0.0f, 0.0f);
-
-    // Compute alpha & beta for this permutation.
-    for (int i = 0; i < 16; i++)
-    {
-        const uint bits = permutation >> (2*i);
-
-        float beta = (bits & 1);
-        if (bits & 2) beta = 0.5f;
-        float alpha = 1.0f - beta;
-    
-        alpha2_sum += alpha * alpha * weights[i];
-        beta2_sum += beta * beta * weights[i];
-        alphabeta_sum += alpha * beta * weights[i];
-        alphax_sum += alpha * colors[i] * weights[i];
-        betax_sum += beta * colors[i] * weights[i];
-    }
-
-    const float factor = 1.0f / (alpha2_sum * beta2_sum - alphabeta_sum * alphabeta_sum);
-
-    float3 a = (alphax_sum * beta2_sum - betax_sum * alphabeta_sum) * factor;
-    float3 b = (betax_sum * alpha2_sum - alphax_sum * alphabeta_sum) * factor;
-    
-    // Round a, b to the closest 5-6-5 color and expand...
-    a = roundAndExpand(a, start);
-    b = roundAndExpand(b, end);
-
-    // compute the error
-    float3 e = a * a * alpha2_sum + b * b * beta2_sum + 2.0f * (a * b * alphabeta_sum - a * alphax_sum - b * betax_sum);
-
-    return dot(e, kColorMetric);
-}
 
 
 ////////////////////////////////////////////////////////////////////////////////



////////////////////////////////////////////////////////////////////////////////
// Find index with minimum error
////////////////////////////////////////////////////////////////////////////////
__device__ void minimizeError(float * errors, int * indices)
{
	const int idx = threadIdx.x;

#if __DEVICE_EMULATION__
	for(int d = NUM_THREADS/2; d > 0; d >>= 1)

	}
	
	// Reorder permutation.
	uint indices = permutation;
	for(int i = 0; i < 16; i++)
	{
		int ref = xrefs[i];
