 		}
 	}
 #endif
-}
-
-
-////////////////////////////////////////////////////////////////////////////////
-// Load color block to shared mem
-////////////////////////////////////////////////////////////////////////////////
-__device__ void loadColorBlock(const uint * image, float3 colors[16], int xrefs[16])
+
+	__syncthreads();
+
+	return indices[0];
+}
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Save DXT block
+////////////////////////////////////////////////////////////////////////////////
+__device__ void saveBlockDXT1(ushort start, ushort end, uint permutation, int xrefs[16], uint2 * result)
 {
 	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);
-		
-		// 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;
-	}
-}
-
-__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)
 	{

#define __debugsync()
#endif

typedef unsigned short ushort;
typedef unsigned int uint;


}




////////////////////////////////////////////////////////////////////////////////
// Compress color block
////////////////////////////////////////////////////////////////////////////////
__global__ void compress(const uint * permutations, const uint * image, uint2 * result)
{
	__shared__ float3 colors[16];
	__shared__ int xrefs[16];
	
	loadColorBlock(image, colors, xrefs);
	
	__syncthreads();

	ushort bestStart, bestEnd;
	uint bestPermutation;

	__shared__ float errors[NUM_THREADS];

	evalAllPermutations(colors, permutations, bestStart, bestEnd, bestPermutation, errors);
	
	// Use a parallel reduction to find minimum error.
	const int minIdx = findMinError(errors);
	
	// Only write the result of the winner thread.
	if (threadIdx.x == minIdx)
	{
		saveBlockDXT1(bestStart, bestEnd, bestPermutation, xrefs, result);
	}
}


__global__ void compressWeighted(const uint * permutations, const uint * image, uint2 * result)
{
	__shared__ float3 colors[16];
	__shared__ float weights[16];
	__shared__ int xrefs[16];
	
	loadColorBlock(image, colors, weights, xrefs);
	
	__syncthreads();

	ushort bestStart, bestEnd;
	uint bestPermutation;

	__shared__ float errors[NUM_THREADS];
	
	evalLevel4Permutations(colors, weights, permutations, bestStart, bestEnd, bestPermutation, errors);
	
	// Use a parallel reduction to find minimum error.
	int minIdx = findMinError(errors);
	
	// Only write the result of the winner thread.
	if (threadIdx.x == minIdx)
	{
		saveBlockDXT1(bestStart, bestEnd, bestPermutation, xrefs, result);
	}
}


__device__ float computeError(const float weights[16], uchar a0, uchar a1)
{
	float palette[6];
	palette[0] = (6.0f/7.0f * a0 + 1.0f/7.0f * a1);
	palette[1] = (5.0f/7.0f * a0 + 2.0f/7.0f * a1);
	palette[2] = (4.0f/7.0f * a0 + 3.0f/7.0f * a1);
	palette[3] = (3.0f/7.0f * a0 + 4.0f/7.0f * a1);
	palette[4] = (2.0f/7.0f * a0 + 5.0f/7.0f * a1);
	palette[5] = (1.0f/7.0f * a0 + 6.0f/7.0f * a1);

	float total = 0.0f;

	for (uint i = 0; i < 16; i++)
	{
		float alpha = weights[i];

		float error = a0 - alpha;
		error = min(error, palette[0] - alpha);
		error = min(error, palette[1] - alpha);
		error = min(error, palette[2] - alpha);
		error = min(error, palette[3] - alpha);
		error = min(error, palette[4] - alpha);
		error = min(error, palette[5] - alpha);
		error = min(error, a1 - alpha);
		
		total += error;
	}
	
	return total;
}

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

/*
__device__ void optimizeAlpha8(const float alphas[16], uchar & a0, uchar & a1)
{
	float alpha2_sum = 0;
	float beta2_sum = 0;
	float alphabeta_sum = 0;
	float alphax_sum = 0;
	float betax_sum = 0;

	for (int i = 0; i < 16; i++)
	{
		uint idx = index[i];
		float alpha;
		if (idx < 2) alpha = 1.0f - idx;
		else alpha = (8.0f - idx) / 7.0f;
		
		float beta = 1 - alpha;

        alpha2_sum += alpha * alpha;
        beta2_sum += beta * beta;
        alphabeta_sum += alpha * beta;
        alphax_sum += alpha * alphas[i];
        betax_sum += beta * alphas[i];
	}

    const float factor = 1.0f / (alpha2_sum * beta2_sum - alphabeta_sum * alphabeta_sum);

    float a = (alphax_sum * beta2_sum - betax_sum * alphabeta_sum) * factor;
	float b = (betax_sum * alpha2_sum - alphax_sum * alphabeta_sum) * factor;

	a0 = roundAndExpand(a);
	a1 = roundAndExpand(b);
}
*/

__device__ void compressAlpha(const float alphas[16], uint4 * result)
{
	const int tid = threadIdx.x;
	
	// Compress alpha block!
	// Brute force approach:
	// Try all color pairs: 256*256/2 = 32768, 32768/64 = 512 iterations?

	// Determine min & max alphas

	float A0, A1;

	if (tid < 16)
	{
		__shared__ uint s_alphas[16];
		
		s_alphas[tid] = alphas[tid];
		s_alphas[tid] = min(s_alphas[tid], s_alphas[tid^8]);
		s_alphas[tid] = min(s_alphas[tid], s_alphas[tid^4]);
		s_alphas[tid] = min(s_alphas[tid], s_alphas[tid^2]);
		s_alphas[tid] = min(s_alphas[tid], s_alphas[tid^1]);
		A0 = s_alphas[tid];
		
		s_alphas[tid] = alphas[tid];
		s_alphas[tid] = max(s_alphas[tid], s_alphas[tid^8]);
		s_alphas[tid] = max(s_alphas[tid], s_alphas[tid^4]);
		s_alphas[tid] = max(s_alphas[tid], s_alphas[tid^2]);
		s_alphas[tid] = max(s_alphas[tid], s_alphas[tid^1]);
		A1 = s_alphas[tid];
	}

	__syncthreads();

	int minIdx = 0;

	if (A1 - A0 > 8)
	{
		float bestError = FLT_MAX;

		// 64 threads -> 8x8
		// divide [A1-A0] in partitions.
		// test endpoints 
		
		for (int i = 0; i < 128; i++)
		{
			uint idx = (i * NUM_THREADS + tid) * 4;
			uchar a0 = idx & 255;
			uchar a1 = idx >> 8;
			
			float error = computeError(alphas, a0, a1);
			
			if (error < bestError)
			{
				bestError = error;
				A0 = a0;
				A1 = a1;
			}
		}
		
		__shared__ float errors[NUM_THREADS];
		errors[tid] = bestError;
		
		// Minimize error.
		minIdx = findMinError(errors);

	}

	if (minIdx == tid)
	{
		// @@ Compute indices.
	
		// @@ Write alpha block.
	}
}

__global__ void compressDXT5(const uint * permutations, const uint * image, uint4 * result)
{
	__shared__ float3 colors[16];
	__shared__ float weights[16];
	__shared__ int xrefs[16];
	
	loadColorBlock(image, colors, weights, xrefs);
		
	__syncthreads();

	compressAlpha(weights, result);	

	ushort bestStart, bestEnd;
	uint bestPermutation;

	__shared__ float errors[NUM_THREADS];
	
	evalLevel4Permutations(colors, weights, permutations, bestStart, bestEnd, bestPermutation, errors);
	
	// Use a parallel reduction to find minimum error.
	int minIdx = findMinError(errors);
	
	// Only write the result of the winner thread.
	if (threadIdx.x == minIdx)
	{
		saveBlockDXT1(bestStart, bestEnd, bestPermutation, xrefs, (uint2 *)result);
	}
}


////////////////////////////////////////////////////////////////////////////////
