 ////////////////////////////////////////////////////////////////////////////////
 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;
+	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;
 }
 
 

	
	float bestError = FLT_MAX;
	
    for(int i = 0; i < 16; i++)
    {
		int pidx = idx + NUM_THREADS * i;
		if (pidx >= 992) break;
        
        ushort start, end;
        uint permutation = permutations[pidx];

        float error = evalPermutation4(colors, weights, colorSum, permutation, &start, &end);
        
        if (error < bestError)
        {
            bestError = error;
            bestPermutation = permutation;
            bestStart = start;
            bestEnd = end;
        }
    }

    if (bestStart < bestEnd)
    {
        swap(bestEnd, bestStart);
        bestPermutation ^= 0x55555555;    // Flip indices.
    }

	errors[idx] = bestError;
}

////////////////////////////////////////////////////////////////////////////////
__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, kColorMetricSqr);
}

__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, kColorMetricSqr);
}

__constant__ const float alphaTable4[4] = { 9.0f, 0.0f, 6.0f, 3.0f };
