 
 inline __device__ uchar roundAndExpand(float a)
 {
-    return rintf(__saturatef(a) * 255.0f);
+	return rintf(__saturatef(a) * 255.0f);
 }
 */
 /*

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



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