Add DXT1n and CTX1 CUDA compressors.

src/nvtt/Compressor.cpp

@@ -56,7 +56,7 @@ namespace
 	
 	static int blockSize(Format format)
 	{
-		if (format == Format_DXT1 || format == Format_DXT1a) {
+		if (format == Format_DXT1 || format == Format_DXT1a || format == Format_DXT1n) {
 			return 8;
 		}
 		else if (format == Format_DXT3) {
@@ -71,6 +71,9 @@ namespace
 		else if (format == Format_BC5) {
 			return 16;
 		}
+		else if (format == Format_CTX1) {
+			return 8;
+		}
 		return 0;
 	}
 
@@ -333,7 +336,7 @@ bool Compressor::Private::outputHeader(const InputOptions::Private & inputOption
 	{
 		header.setLinearSize(computeImageSize(inputOptions.targetWidth, inputOptions.targetHeight, inputOptions.targetDepth, compressionOptions.bitcount, compressionOptions.format));
 		
-		if (compressionOptions.format == Format_DXT1 || compressionOptions.format == Format_DXT1a) {
+		if (compressionOptions.format == Format_DXT1 || compressionOptions.format == Format_DXT1a || compressionOptions.format == Format_DXT1n) {
 			header.setFourCC('D', 'X', 'T', '1');
 			if (inputOptions.isNormalMap) header.setNormalFlag(true);
 		}
@@ -354,6 +357,10 @@ bool Compressor::Private::outputHeader(const InputOptions::Private & inputOption
 			header.setFourCC('A', 'T', 'I', '2');
 			if (inputOptions.isNormalMap) header.setNormalFlag(true);
 		}
+		else if (compressionOptions.format == Format_CTX1) {
+			header.setFourCC('C', 'T', 'X', '1');
+			if (inputOptions.isNormalMap) header.setNormalFlag(true);
+		}
 	}
 	
 	// Swap bytes if necessary.
@@ -705,6 +712,18 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio
 			}
 		}
 	}
+	else if (compressionOptions.format == Format_DXT1n)
+	{
+		if (cudaEnabled)
+		{
+			nvDebugCheck(cudaSupported);
+			cuda->compressDXT1n(image, outputOptions, compressionOptions);
+		}
+		else
+		{
+			if (outputOptions.errorHandler) outputOptions.errorHandler->error(Error_UnsupportedFeature);
+		}
+	}
 	else if (compressionOptions.format == Format_DXT3)
 	{
 		if (compressionOptions.quality == Quality_Fastest)
@@ -762,6 +781,18 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio
 	{
 		compressBC5(image, outputOptions, compressionOptions);
 	}
+	else if (compressionOptions.format == Format_CTX1)
+	{
+		if (cudaEnabled)
+		{
+			nvDebugCheck(cudaSupported);
+			cuda->compressCTX1(image, outputOptions, compressionOptions);
+		}
+		else
+		{
+			if (outputOptions.errorHandler) outputOptions.errorHandler->error(Error_UnsupportedFeature);
+		}
+	}
 
 	return true;
 }

src/nvtt/cuda/Bitmaps.h

@@ -122,7 +122,7 @@ static void doPrecomputation()
 */
 
 
-const static uint bitmaps[992] =
+const static uint s_bitmapTable[992] =
 {
 	0x80000000,
 	0x40000000,

src/nvtt/cuda/CompressKernel.cu

@@ -60,6 +60,7 @@ __device__ void sortColors(const float * values, int * cmp)
 {
 	int tid = threadIdx.x;
 
+#if 1
 	cmp[tid] = (values[0] < values[tid]);
 	cmp[tid] += (values[1] < values[tid]);
 	cmp[tid] += (values[2] < values[tid]);
@@ -93,6 +94,23 @@ __device__ void sortColors(const float * values, int * cmp)
 	if (tid > 12 && cmp[tid] == cmp[12]) ++cmp[tid];
 	if (tid > 13 && cmp[tid] == cmp[13]) ++cmp[tid];
 	if (tid > 14 && cmp[tid] == cmp[14]) ++cmp[tid];
+#else
+
+	cmp[tid] = 0;
+
+	#pragma unroll
+	for (int i = 0; i < 16; i++)
+	{
+		cmp[tid] += (values[i] < values[tid]);
+	}
+
+	// Resolve elements with the same index.
+	#pragma unroll
+	for (int i = 0; i < 15; i++)
+	{
+		if (tid > 0 && cmp[tid] == cmp[i]) ++cmp[tid];
+	}
+#endif
 }
 
 
@@ -183,11 +201,48 @@ __device__ void loadColorBlock(const uint * image, float3 colors[16], float3 sum
 	}
 }
 
+__device__ void loadColorBlock(const uint * image, float2 colors[16], float2 sums[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].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.
+		colorSums(colors, sums);
+		float2 axis = bestFitLine(colors, sums[0]);
+		
+		dps[idx] = dot(colors[idx], axis);
+		
+#if __DEVICE_EMULATION__
+		} __debugsync(); if (idx < 16) {
+#endif
+		
+		sortColors(dps, xrefs);
+		
+		float2 tmp = colors[idx];
+		colors[xrefs[idx]] = tmp;
+	}
+}
+
 
 ////////////////////////////////////////////////////////////////////////////////
 // Round color to RGB565 and expand
 ////////////////////////////////////////////////////////////////////////////////
-inline __device__ float3 roundAndExpand(float3 v, ushort * w)
+inline __device__ float3 roundAndExpand565(float3 v, ushort * w)
 {
 	v.x = rintf(__saturatef(v.x) * 31.0f);
 	v.y = rintf(__saturatef(v.y) * 63.0f);
@@ -199,6 +254,26 @@ inline __device__ float3 roundAndExpand(float3 v, ushort * w)
 	return v;
 }
 
+inline __device__ float2 roundAndExpand56(float2 v, ushort * w)
+{
+	v.x = rintf(__saturatef(v.x) * 31.0f);
+	v.y = rintf(__saturatef(v.y) * 63.0f);
+	*w = ((ushort)v.x << 11) | ((ushort)v.y << 5);
+	v.x *= 0.03227752766457f; // approximate integer bit expansion.
+	v.y *= 0.01583151765563f;
+	return v;
+}
+
+inline __device__ float2 roundAndExpand88(float2 v, ushort * w)
+{
+	v.x = rintf(__saturatef(v.x) * 255.0f);
+	v.y = rintf(__saturatef(v.y) * 255.0f);
+	*w = ((ushort)v.x << 8) | ((ushort)v.y);
+	v.x *= 1.0f / 255.0f;
+	v.y *= 1.0f / 255.0f;
+	return v;
+}
+
 
 ////////////////////////////////////////////////////////////////////////////////
 // Evaluate permutations
@@ -234,8 +309,8 @@ __device__ float evalPermutation4(const float3 * colors, uint permutation, ushor
 	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);
+	a = roundAndExpand565(a, start);
-	b = roundAndExpand(b, end);
+	b = roundAndExpand565(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);
@@ -274,8 +349,8 @@ __device__ float evalPermutation3(const float3 * colors, uint permutation, ushor
 	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);
+	a = roundAndExpand565(a, start);
-	b = roundAndExpand(b, end);
+	b = roundAndExpand565(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);
@@ -315,8 +390,8 @@ __device__ float evalPermutation4(const float3 * colors, float3 color_sum, uint
 	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);
+	a = roundAndExpand565(a, start);
-	b = roundAndExpand(b, end);
+	b = roundAndExpand565(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);
@@ -351,8 +426,8 @@ __device__ float evalPermutation3(const float3 * colors, float3 color_sum, uint
 	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);
+	a = roundAndExpand565(a, start);
-	b = roundAndExpand(b, end);
+	b = roundAndExpand565(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);
@@ -391,8 +466,8 @@ __device__ float evalPermutation4(const float3 * colors, const float * weights,
 	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);
+	a = roundAndExpand565(a, start);
-	b = roundAndExpand(b, end);
+	b = roundAndExpand565(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);
@@ -432,8 +507,8 @@ __device__ float evalPermutation3(const float3 * colors, const float * weights,
 	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);
+	a = roundAndExpand565(a, start);
-	b = roundAndExpand(b, end);
+	b = roundAndExpand565(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);
@@ -442,6 +517,114 @@ __device__ float evalPermutation3(const float3 * colors, const float * weights,
 }
 */
 
+__device__ float evalPermutation4(const float2 * colors, float2 color_sum, uint permutation, ushort * start, ushort * end)
+{
+	// Compute endpoints using least squares.
+	float2 alphax_sum = make_float2(0.0f, 0.0f);
+	uint akku = 0;
+
+	// Compute alpha & beta for this permutation.
+	#pragma unroll
+	for (int i = 0; i < 16; i++)
+	{
+		const uint bits = permutation >> (2*i);
+		
+		alphax_sum += alphaTable4[bits & 3] * colors[i];
+		akku += prods4[bits & 3];
+	}
+
+	float alpha2_sum = float(akku >> 16);
+	float beta2_sum = float((akku >> 8) & 0xff);
+	float alphabeta_sum = float(akku & 0xff);
+	float2 betax_sum = 9.0f * color_sum - alphax_sum;
+
+	const float factor = 1.0f / (alpha2_sum * beta2_sum - alphabeta_sum * alphabeta_sum);
+
+	float2 a = (alphax_sum * beta2_sum - betax_sum * alphabeta_sum) * factor;
+	float2 b = (betax_sum * alpha2_sum - alphax_sum * alphabeta_sum) * factor;
+	
+	// Round a, b to the closest 5-6 color and expand...
+	a = roundAndExpand56(a, start);
+	b = roundAndExpand56(b, end);
+
+	// compute the error
+	float2 e = a * a * alpha2_sum + b * b * beta2_sum + 2.0f * (a * b * alphabeta_sum - a * alphax_sum - b * betax_sum);
+
+	return (1.0f / 9.0f) * (e.x + e.y);
+}
+
+__device__ float evalPermutation3(const float2 * colors, float2 color_sum, uint permutation, ushort * start, ushort * end)
+{
+	// Compute endpoints using least squares.
+	float2 alphax_sum = make_float2(0.0f, 0.0f);
+	uint akku = 0;
+
+	// Compute alpha & beta for this permutation.
+	#pragma unroll
+	for (int i = 0; i < 16; i++)
+	{
+		const uint bits = permutation >> (2*i);
+
+		alphax_sum += alphaTable3[bits & 3] * colors[i];
+		akku += prods3[bits & 3];
+	}
+
+	float alpha2_sum = float(akku >> 16);
+	float beta2_sum = float((akku >> 8) & 0xff);
+	float alphabeta_sum = float(akku & 0xff);
+	float2 betax_sum = 4.0f * color_sum - alphax_sum;
+
+	const float factor = 1.0f / (alpha2_sum * beta2_sum - alphabeta_sum * alphabeta_sum);
+
+	float2 a = (alphax_sum * beta2_sum - betax_sum * alphabeta_sum) * factor;
+	float2 b = (betax_sum * alpha2_sum - alphax_sum * alphabeta_sum) * factor;
+	
+	// Round a, b to the closest 5-6 color and expand...
+	a = roundAndExpand56(a, start);
+	b = roundAndExpand56(b, end);
+
+	// compute the error
+	float2 e = a * a * alpha2_sum + b * b * beta2_sum + 2.0f * (a * b * alphabeta_sum - a * alphax_sum - b * betax_sum);
+
+	return (1.0f / 4.0f) * (e.x + e.y);
+}
+
+__device__ float evalPermutationCTX(const float2 * colors, float2 color_sum, uint permutation, ushort * start, ushort * end)
+{
+	// Compute endpoints using least squares.
+	float2 alphax_sum = make_float2(0.0f, 0.0f);
+	uint akku = 0;
+
+	// Compute alpha & beta for this permutation.
+	#pragma unroll
+	for (int i = 0; i < 16; i++)
+	{
+		const uint bits = permutation >> (2*i);
+		
+		alphax_sum += alphaTable4[bits & 3] * colors[i];
+		akku += prods4[bits & 3];
+	}
+
+	float alpha2_sum = float(akku >> 16);
+	float beta2_sum = float((akku >> 8) & 0xff);
+	float alphabeta_sum = float(akku & 0xff);
+	float2 betax_sum = 9.0f * color_sum - alphax_sum;
+
+	const float factor = 1.0f / (alpha2_sum * beta2_sum - alphabeta_sum * alphabeta_sum);
+
+	float2 a = (alphax_sum * beta2_sum - betax_sum * alphabeta_sum) * factor;
+	float2 b = (betax_sum * alpha2_sum - alphax_sum * alphabeta_sum) * factor;
+	
+	// Round a, b to the closest 8-8 color and expand...
+	a = roundAndExpand88(a, start);
+	b = roundAndExpand88(b, end);
+
+	// compute the error
+	float2 e = a * a * alpha2_sum + b * b * beta2_sum + 2.0f * (a * b * alphabeta_sum - a * alphax_sum - b * betax_sum);
+
+	return (1.0f / 9.0f) * (e.x + e.y);
+}
+
 
 ////////////////////////////////////////////////////////////////////////////////
 // Evaluate all permutations
@@ -570,6 +753,67 @@ __device__ void evalAllPermutations(const float3 * colors, const float * weights
 }
 */
 
+__device__ void evalAllPermutations(const float2 * colors, float2 colorSum, const uint * permutations, ushort & bestStart, ushort & bestEnd, uint & bestPermutation, float * errors)
+{
+	const int idx = threadIdx.x;
+	
+	float bestError = FLT_MAX;
+	
+	__shared__ uint s_permutations[160];
+
+	for(int i = 0; i < 16; i++)
+	{
+		int pidx = idx + NUM_THREADS * i;
+		if (pidx >= 992) break;
+		
+		ushort start, end;
+		uint permutation = permutations[pidx];
+		if (pidx < 160) s_permutations[pidx] = permutation;
+				
+		float error = evalPermutation4(colors, 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.
+	}
+
+	for(int i = 0; i < 3; i++)
+	{
+		int pidx = idx + NUM_THREADS * i;
+		if (pidx >= 160) break;
+		
+		ushort start, end;
+		uint permutation = s_permutations[pidx];
+		float error = evalPermutation3(colors, colorSum, permutation, &start, &end);
+		
+		if (error < bestError)
+		{
+			bestError = error;
+			bestPermutation = permutation;
+			bestStart = start;
+			bestEnd = end;
+			
+			if (bestStart > bestEnd)
+			{
+				swap(bestEnd, bestStart);
+				bestPermutation ^= (~bestPermutation >> 1) & 0x55555555;	// Flip indices.
+			}
+		}
+	}
+
+	errors[idx] = bestError;
+}
+
 __device__ void evalLevel4Permutations(const float3 * colors, const float * weights, float3 colorSum, const uint * permutations, ushort & bestStart, ushort & bestEnd, uint & bestPermutation, float * errors)
 {
 	const int idx = threadIdx.x;
@@ -604,6 +848,39 @@ __device__ void evalLevel4Permutations(const float3 * colors, const float * weig
 	errors[idx] = bestError;
 }
 
+__device__ void evalAllPermutationsCTX(const float2 * colors, float2 colorSum, const uint * permutations, ushort & bestStart, ushort & bestEnd, uint & bestPermutation, float * errors)
+{
+	const int idx = threadIdx.x;
+	
+	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 = evalPermutationCTX(colors, 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;
+}
 
 
 ////////////////////////////////////////////////////////////////////////////////
@@ -715,13 +992,17 @@ __device__ void saveBlockDXT1(ushort start, ushort end, uint permutation, int xr
 	result[bid].y = indices;
 }
 
+__device__ void saveBlockCTX1(ushort start, ushort end, uint permutation, int xrefs[16], uint2 * result)
+{
+	saveBlockDXT1(start, end, permutation, xrefs, result);
+}
 
 
 
 ////////////////////////////////////////////////////////////////////////////////
 // Compress color block
 ////////////////////////////////////////////////////////////////////////////////
-__global__ void compress(const uint * permutations, const uint * image, uint2 * result)
+__global__ void compressDXT1(const uint * permutations, const uint * image, uint2 * result)
 {
 	__shared__ float3 colors[16];
 	__shared__ float3 sums[16];
@@ -749,7 +1030,7 @@ __global__ void compress(const uint * permutations, const uint * image, uint2 *
 }
 
 
-__global__ void compressWeighted(const uint * permutations, const uint * image, uint2 * result)
+__global__ void compressWeightedDXT1(const uint * permutations, const uint * image, uint2 * result)
 {
 	__shared__ float3 colors[16];
 	__shared__ float3 sums[16];
@@ -778,6 +1059,61 @@ __global__ void compressWeighted(const uint * permutations, const uint * image,
 }
 
 
+__global__ void compressNormalDXT1(const uint * permutations, const uint * image, uint2 * result)
+{
+	__shared__ float2 colors[16];
+	__shared__ float2 sums[16];
+	__shared__ int xrefs[16];
+	
+	loadColorBlock(image, colors, sums, xrefs);
+	
+	__syncthreads();
+
+	ushort bestStart, bestEnd;
+	uint bestPermutation;
+
+	__shared__ float errors[NUM_THREADS];
+
+	evalAllPermutations(colors, sums[0], 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 compressCTX1(const uint * permutations, const uint * image, uint2 * result)
+{
+	__shared__ float2 colors[16];
+	__shared__ float2 sums[16];
+	__shared__ int xrefs[16];
+	
+	loadColorBlock(image, colors, sums, xrefs);
+	
+	__syncthreads();
+
+	ushort bestStart, bestEnd;
+	uint bestPermutation;
+
+	__shared__ float errors[NUM_THREADS];
+
+	evalAllPermutationsCTX(colors, sums[0], 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)
+	{
+		saveBlockCTX1(bestStart, bestEnd, bestPermutation, xrefs, result);
+	}
+}
+
+
 /*
 __device__ float computeError(const float weights[16], uchar a0, uchar a1)
 {
@@ -845,8 +1181,8 @@ __device__ void optimizeAlpha8(const float alphas[16], uchar & a0, uchar & a1)
 	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);
+	a0 = roundAndExpand8(a);
-	a1 = roundAndExpand(b);
+	a1 = roundAndExpand8(b);
 }
 */
 /*
@@ -978,12 +1314,22 @@ extern "C" void setupCompressKernel(const float weights[3])
 // Launch kernel
 ////////////////////////////////////////////////////////////////////////////////
 
-extern "C" void compressKernel(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps)
+extern "C" void compressKernelDXT1(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps)
 {
-	compress<<<blockNum, NUM_THREADS>>>(d_bitmaps, d_data, (uint2 *)d_result);
+	compressDXT1<<<blockNum, NUM_THREADS>>>(d_bitmaps, d_data, (uint2 *)d_result);
 }
 
-extern "C" void compressWeightedKernel(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps)
+extern "C" void compressWeightedKernelDXT1(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps)
 {
-	compressWeighted<<<blockNum, NUM_THREADS>>>(d_bitmaps, d_data, (uint2 *)d_result);
+	compressWeightedDXT1<<<blockNum, NUM_THREADS>>>(d_bitmaps, d_data, (uint2 *)d_result);
+}
+
+extern "C" void compressNormalKernelDXT1(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps)
+{
+	compressNormalDXT1<<<blockNum, NUM_THREADS>>>(d_bitmaps, d_data, (uint2 *)d_result);
+}
+
+extern "C" void compressKernelCTX1(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps)
+{
+	compressCTX1<<<blockNum, NUM_THREADS>>>(d_bitmaps, d_data, (uint2 *)d_result);
 }

src/nvtt/cuda/CudaCompressDXT.cpp

@@ -48,11 +48,16 @@ using namespace nvtt;
 
 #if defined HAVE_CUDA
 
-#define MAX_BLOCKS 32768 // 49152, 65535
+//#define MAX_BLOCKS 32768U // 49152, 65535
+#define MAX_BLOCKS 8192U // 49152, 65535
+
 
 extern "C" void setupCompressKernel(const float weights[3]);
-extern "C" void compressKernel(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps);
+extern "C" void compressKernelDXT1(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps);
-extern "C" void compressWeightedKernel(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps);
+extern "C" void compressWeightedKernelDXT1(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps);
+extern "C" void compressNormalKernelDXT1(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps);
+extern "C" void compressKernelCTX1(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps);
+
 
 #include "Bitmaps.h"	// @@ Rename to BitmapTable.h
 
@@ -84,7 +89,7 @@ CudaCompressor::CudaCompressor()
 #if defined HAVE_CUDA
     // Allocate and upload bitmaps.
     cudaMalloc((void**) &m_bitmapTable, 992 * sizeof(uint));
-    cudaMemcpy(m_bitmapTable, bitmaps, 992 * sizeof(uint), cudaMemcpyHostToDevice);
+    cudaMemcpy(m_bitmapTable, s_bitmapTable, 992 * sizeof(uint), cudaMemcpyHostToDevice);
 
 	// Allocate scratch buffers.
     cudaMalloc((void**) &m_data, MAX_BLOCKS * 64U);
@@ -119,11 +124,10 @@ void CudaCompressor::compressDXT1(const Image * image, const OutputOptions::Priv
 
 	uint imageSize = w * h * 16 * sizeof(Color32);
     uint * blockLinearImage = (uint *) malloc(imageSize);
-	convertToBlockLinear(image, blockLinearImage);	// @@ Do this on the GPU!
+	convertToBlockLinear(image, blockLinearImage);	// @@ Do this in parallel with the GPU, or in the GPU!
 
 	const uint blockNum = w * h;
 	const uint compressedSize = blockNum * 8;
-	const uint blockMax = 32768; // 49152, 65535
 
 	clock_t start = clock();
 
@@ -133,12 +137,12 @@ void CudaCompressor::compressDXT1(const Image * image, const OutputOptions::Priv
 	uint bn = 0;
 	while(bn != blockNum)
 	{
-		uint count = min(blockNum - bn, blockMax);
+		uint count = min(blockNum - bn, MAX_BLOCKS);
 
 	    cudaMemcpy(m_data, blockLinearImage + bn * 16, count * 64, cudaMemcpyHostToDevice);
 
 		// Launch kernel.
-		compressKernel(count, m_data, m_result, m_bitmapTable);
+		compressKernelDXT1(count, m_data, m_result, m_bitmapTable);
 
 		// Check for errors.
 		cudaError_t err = cudaGetLastError();
@@ -194,10 +198,9 @@ void CudaCompressor::compressDXT3(const Image * image, const OutputOptions::Priv
 
 	const uint blockNum = w * h;
 	const uint compressedSize = blockNum * 8;
-	const uint blockMax = 32768; // 49152, 65535
 
 	AlphaBlockDXT3 * alphaBlocks = NULL;
-	alphaBlocks = (AlphaBlockDXT3 *)malloc(min(compressedSize, blockMax * 8U));
+	alphaBlocks = (AlphaBlockDXT3 *)malloc(min(compressedSize, MAX_BLOCKS * 8U));
 
 	setupCompressKernel(compressionOptions.colorWeight.ptr());
 	
@@ -206,12 +209,12 @@ void CudaCompressor::compressDXT3(const Image * image, const OutputOptions::Priv
 	uint bn = 0;
 	while(bn != blockNum)
 	{
-		uint count = min(blockNum - bn, blockMax);
+		uint count = min(blockNum - bn, MAX_BLOCKS);
 
 	    cudaMemcpy(m_data, blockLinearImage + bn * 16, count * 64, cudaMemcpyHostToDevice);
 
 		// Launch kernel.
-		compressWeightedKernel(count, m_data, m_result, m_bitmapTable);
+		compressWeightedKernelDXT1(count, m_data, m_result, m_bitmapTable);
 
 		// Compress alpha in parallel with the GPU.
 		for (uint i = 0; i < count; i++)
@@ -279,10 +282,9 @@ void CudaCompressor::compressDXT5(const Image * image, const OutputOptions::Priv
 
 	const uint blockNum = w * h;
 	const uint compressedSize = blockNum * 8;
-	const uint blockMax = 32768; // 49152, 65535
 
 	AlphaBlockDXT5 * alphaBlocks = NULL;
-	alphaBlocks = (AlphaBlockDXT5 *)malloc(min(compressedSize, blockMax * 8U));
+	alphaBlocks = (AlphaBlockDXT5 *)malloc(min(compressedSize, MAX_BLOCKS * 8U));
 
 	setupCompressKernel(compressionOptions.colorWeight.ptr());
 	
@@ -291,12 +293,12 @@ void CudaCompressor::compressDXT5(const Image * image, const OutputOptions::Priv
 	uint bn = 0;
 	while(bn != blockNum)
 	{
-		uint count = min(blockNum - bn, blockMax);
+		uint count = min(blockNum - bn, MAX_BLOCKS);
 
 	    cudaMemcpy(m_data, blockLinearImage + bn * 16, count * 64, cudaMemcpyHostToDevice);
 
 		// Launch kernel.
-		compressWeightedKernel(count, m_data, m_result, m_bitmapTable);
+		compressWeightedKernelDXT1(count, m_data, m_result, m_bitmapTable);
 
 		// Compress alpha in parallel with the GPU.
 		for (uint i = 0; i < count; i++)
@@ -348,6 +350,144 @@ void CudaCompressor::compressDXT5(const Image * image, const OutputOptions::Priv
 }
 
 
+void CudaCompressor::compressDXT1n(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions)
+{
+	nvDebugCheck(cuda::isHardwarePresent());
+#if defined HAVE_CUDA
+
+	// Image size in blocks.
+	const uint w = (image->width() + 3) / 4;
+	const uint h = (image->height() + 3) / 4;
+
+	uint imageSize = w * h * 16 * sizeof(Color32);
+    uint * blockLinearImage = (uint *) malloc(imageSize);
+	convertToBlockLinear(image, blockLinearImage);	// @@ Do this in parallel with the GPU, or in the GPU!
+
+	const uint blockNum = w * h;
+	const uint compressedSize = blockNum * 8;
+
+	clock_t start = clock();
+
+	setupCompressKernel(compressionOptions.colorWeight.ptr());
+	
+	// TODO: Add support for multiple GPUs.
+	uint bn = 0;
+	while(bn != blockNum)
+	{
+		uint count = min(blockNum - bn, MAX_BLOCKS);
+
+	    cudaMemcpy(m_data, blockLinearImage + bn * 16, count * 64, cudaMemcpyHostToDevice);
+
+		// Launch kernel.
+		compressNormalKernelDXT1(count, m_data, m_result, m_bitmapTable);
+
+		// Check for errors.
+		cudaError_t err = cudaGetLastError();
+		if (err != cudaSuccess)
+		{
+			nvDebug("CUDA Error: %s\n", cudaGetErrorString(err));
+
+			if (outputOptions.errorHandler != NULL)
+			{
+				outputOptions.errorHandler->error(Error_CudaError);
+			}
+		}
+
+		// Copy result to host, overwrite swizzled image.
+		cudaMemcpy(blockLinearImage, m_result, count * 8, cudaMemcpyDeviceToHost);
+
+		// Output result.
+		if (outputOptions.outputHandler != NULL)
+		{
+			outputOptions.outputHandler->writeData(blockLinearImage, count * 8);
+		}
+
+		bn += count;
+	}
+
+	clock_t end = clock();
+	printf("\rCUDA time taken: %.3f seconds\n", float(end-start) / CLOCKS_PER_SEC);
+
+	free(blockLinearImage);
+
+#else
+	if (outputOptions.errorHandler != NULL)
+	{
+		outputOptions.errorHandler->error(Error_CudaError);
+	}
+#endif
+}
+
+
+void CudaCompressor::compressCTX1(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions)
+{
+	nvDebugCheck(cuda::isHardwarePresent());
+#if defined HAVE_CUDA
+
+	// Image size in blocks.
+	const uint w = (image->width() + 3) / 4;
+	const uint h = (image->height() + 3) / 4;
+
+	uint imageSize = w * h * 16 * sizeof(Color32);
+    uint * blockLinearImage = (uint *) malloc(imageSize);
+	convertToBlockLinear(image, blockLinearImage);	// @@ Do this in parallel with the GPU, or in the GPU!
+
+	const uint blockNum = w * h;
+	const uint compressedSize = blockNum * 8;
+
+	clock_t start = clock();
+
+	setupCompressKernel(compressionOptions.colorWeight.ptr());
+	
+	// TODO: Add support for multiple GPUs.
+	uint bn = 0;
+	while(bn != blockNum)
+	{
+		uint count = min(blockNum - bn, MAX_BLOCKS);
+
+	    cudaMemcpy(m_data, blockLinearImage + bn * 16, count * 64, cudaMemcpyHostToDevice);
+
+		// Launch kernel.
+		compressKernelCTX1(count, m_data, m_result, m_bitmapTable);
+
+		// Check for errors.
+		cudaError_t err = cudaGetLastError();
+		if (err != cudaSuccess)
+		{
+			nvDebug("CUDA Error: %s\n", cudaGetErrorString(err));
+
+			if (outputOptions.errorHandler != NULL)
+			{
+				outputOptions.errorHandler->error(Error_CudaError);
+			}
+		}
+
+		// Copy result to host, overwrite swizzled image.
+		cudaMemcpy(blockLinearImage, m_result, count * 8, cudaMemcpyDeviceToHost);
+
+		// Output result.
+		if (outputOptions.outputHandler != NULL)
+		{
+			outputOptions.outputHandler->writeData(blockLinearImage, count * 8);
+		}
+
+		bn += count;
+	}
+
+	clock_t end = clock();
+	printf("\rCUDA time taken: %.3f seconds\n", float(end-start) / CLOCKS_PER_SEC);
+
+	free(blockLinearImage);
+
+#else
+	if (outputOptions.errorHandler != NULL)
+	{
+		outputOptions.errorHandler->error(Error_CudaError);
+	}
+#endif
+}
+
+
 
 #if 0
 
@@ -443,7 +583,7 @@ public:
 		cudaMemcpy(d_blockLinearImage, blockLinearImage, blockCount * 64, cudaMemcpyHostToDevice);
 		
 		// Launch kernel.
-		compressKernel(blockCount, d_blockLinearImage, d_compressedImage, d_bitmaps);
+		compressKernelDXT1(blockCount, d_blockLinearImage, d_compressedImage, d_bitmaps);
 		
 		// Check for errors.
 		cudaError_t err = cudaGetLastError();
@@ -495,7 +635,7 @@ void nv::cudaCompressDXT1_2(const Image * image, const OutputOptions::Private &
 	const uint blockNum = ((w + 3) / 4) * ((h + 3) / 4);
 	const uint blockMax = 32768; // 49152, 65535
 		
-	setupCompressKernel(compressionOptions.colorWeight.ptr());
+	setupCompressKernelDXT1(compressionOptions.colorWeight.ptr());
 
 	ColorBlock rgba;
 	Task task(min(blockNum, blockMax));

src/nvtt/cuda/CudaCompressDXT.h

@@ -40,6 +40,8 @@ namespace nv
 		void compressDXT1(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions);
 		void compressDXT3(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions);
 		void compressDXT5(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions);
+		void compressDXT1n(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions);
+		void compressCTX1(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions);
 
 	private:
 

src/nvtt/cuda/CudaMath.h

@@ -82,6 +82,62 @@ inline __device__ __host__ void operator /=(float3 & b, float f)
     b.z *= inv;
 }
 
+// float2 operators
+inline __device__ __host__ float2 operator *(float2 a, float2 b)
+{
+    return make_float2(a.x*b.x, a.y*b.y);
+}
+
+inline __device__ __host__ float2 operator *(float f, float2 v)
+{
+    return make_float2(v.x*f, v.y*f);
+}
+
+inline __device__ __host__ float2 operator *(float2 v, float f)
+{
+    return make_float2(v.x*f, v.y*f);
+}
+
+inline __device__ __host__ float2 operator +(float2 a, float2 b)
+{
+    return make_float2(a.x+b.x, a.y+b.y);
+}
+
+inline __device__ __host__ void operator +=(float2 & b, float2 a)
+{
+    b.x += a.x;
+    b.y += a.y;
+}
+
+inline __device__ __host__ float2 operator -(float2 a, float2 b)
+{
+    return make_float2(a.x-b.x, a.y-b.y);
+}
+
+inline __device__ __host__ void operator -=(float2 & b, float2 a)
+{
+    b.x -= a.x;
+    b.y -= a.y;
+}
+
+inline __device__ __host__ float2 operator /(float2 v, float f)
+{
+    float inv = 1.0f / f;
+    return v * inv;
+}
+
+inline __device__ __host__ void operator /=(float2 & b, float f)
+{
+    float inv = 1.0f / f;
+    b.x *= inv;
+	b.y *= inv;
+}
+
+
+inline __device__ __host__ float dot(float2 a, float2 b)
+{
+    return a.x * b.x + a.y * b.y;
+}
 
 inline __device__ __host__ float dot(float3 a, float3 b)
 {
@@ -217,5 +273,91 @@ inline __device__ float3 bestFitLine(const float3 * colors, float3 color_sum, fl
 	return firstEigenVector(covariance);
 }
 
+// @@ For 2D this may not be the most efficient method. It's a quadratic equation, right?
+inline __device__ __host__ float2 firstEigenVector2D( float matrix[3] )
+{
+	// @@ 8 iterations is probably more than enough.
+
+	float2 v = make_float2(1.0f, 1.0f);
+	for(int i = 0; i < 8; i++) {
+		float x = v.x * matrix[0] + v.y * matrix[1];
+		float y = v.x * matrix[1] + v.y * matrix[2];
+		float m = max(x, y);        
+		float iv = 1.0f / m;
+		#if __DEVICE_EMULATION__
+		if (m == 0.0f) iv = 0.0f;
+		#endif
+		v = make_float2(x*iv, y*iv);
+	}
+
+	return v;
+}
+
+inline __device__ void colorSums(const float2 * colors, float2 * sums)
+{
+#if __DEVICE_EMULATION__
+	float2 color_sum = make_float2(0.0f, 0.0f, 0.0f);
+	for (int i = 0; i < 16; i++)
+	{
+		color_sum += colors[i];
+	}
+
+	for (int i = 0; i < 16; i++)
+	{
+		sums[i] = color_sum;
+	}
+#else
+
+	const int idx = threadIdx.x;
+
+	sums[idx] = colors[idx];
+	sums[idx] += sums[idx^8];
+	sums[idx] += sums[idx^4];
+	sums[idx] += sums[idx^2];
+	sums[idx] += sums[idx^1];
+
+#endif
+}
+
+inline __device__ float2 bestFitLine(const float2 * colors, float2 color_sum)
+{
+	// Compute covariance matrix of the given colors.
+#if __DEVICE_EMULATION__
+	float covariance[3] = {0, 0, 0};
+	for (int i = 0; i < 16; i++)
+	{
+		float2 a = (colors[i] - color_sum * (1.0f / 16.0f));
+		covariance[0] += a.x * a.x;
+		covariance[1] += a.x * a.y;
+		covariance[3] += a.y * a.y;
+	}
+#else
+
+	const int idx = threadIdx.x;
+
+	float2 diff = (colors[idx] - color_sum * (1.0f / 16.0f));
+
+	__shared__ float covariance[16*3];
+
+	covariance[3 * idx + 0] = diff.x * diff.x;
+	covariance[3 * idx + 1] = diff.x * diff.y;
+	covariance[3 * idx + 2] = diff.y * diff.y;
+
+	for(int d = 8; d > 0; d >>= 1)
+	{
+		if (idx < d)
+		{
+			covariance[3 * idx + 0] += covariance[3 * (idx+d) + 0];
+			covariance[3 * idx + 1] += covariance[3 * (idx+d) + 1];
+			covariance[3 * idx + 2] += covariance[3 * (idx+d) + 2];
+		}
+	}
+
+#endif
+
+	// Compute first eigen vector.
+	return firstEigenVector2D(covariance);
+}
+
 
 #endif // CUDAMATH_H

src/nvtt/nvtt.h

@@ -75,6 +75,9 @@ namespace nvtt
 		Format_BC3n = Format_DXT5n,
 		Format_BC4,     // ATI1
 		Format_BC5,     // 3DC, ATI2
+
+		Format_DXT1n,
+		Format_CTX1,
 	};
 	
 	/// Quality modes.

src/nvtt/tests/stress.cpp

@@ -83,9 +83,102 @@ struct MyOutputHandler : public nvtt::OutputHandler
 
 };
 
+void precomp()
+{
+	unsigned int bitmaps[1024];
+
+	int num = 0;
+
+	printf("{\n");
+	printf("\t%8X,\n", 0);
+
+	bitmaps[0] = 0;
+
+	num = 1;
+	for (int a = 1; a <= 15; a++)
+	{
+		  for (int b = a; b <= 15; b++)
+		  {
+				for (int c = b; c <= 15; c++)
+				{
+					int indices[16];
+
+					int i = 0;
+					for(; i < a; i++) {
+						indices[i] = 0;
+					}
+					for(; i < a+b; i++) {
+						indices[i] = 2;
+					}
+					for(; i < a+b+c; i++) {
+						indices[i] = 3;
+					}
+					for(; i < 16; i++) {
+						indices[i] = 1;
+					}
+
+					unsigned int bm = 0;
+					for(i = 0; i < 16; i++) {
+						bm |= indices[i] << (i * 2);
+					}
+
+					printf("\t0x%8X, // %d %d %d %d\n", bm, a-0, b-a, c-b, 16-c);
+
+					bitmaps[num] = bm;
+					num++;
+				}
+		  }
+	}
+
+	printf("}\n");
+
+	printf("// num = %d\n", num);
+
+/*
+	for( int i = imax; i >= 0; --i )
+	{
+		// second cluster [i,j) is one third along
+		for( int m = i; m < 16; ++m )
+		{
+			indices[m] = 2;
+		}
+		const int jmax = ( i == 0 ) ? 15 : 16;
+		for( int j = jmax; j >= i; --j )
+		{
+			// third cluster [j,k) is two thirds along
+			for( int m = j; m < 16; ++m )
+			{
+				indices[m] = 3;
+			}
+			
+			int kmax = ( j == 0 ) ? 15 : 16;
+			for( int k = kmax; k >= j; --k )
+			{
+				// last cluster [k,n) is at the end
+				if( k < 16 )
+				{
+					indices[k] = 1;
+				}
+				
+				uint bitmap = 0;
+				
+				bool hasThree = false;
+				for(int p = 0; p < 16; p++) {
+					bitmap |= indices[p] << (p * 2);
+				}
+				
+				bitmaps[num] = bitmap;
+				num++;
+			}
+		}
+	}
+*/
+}
 
 int main(int argc, char *argv[])
 {
+	precomp();
+
 	nvtt::InputOptions inputOptions;
 	inputOptions.setTextureLayout(nvtt::TextureType_2D, 1024, 1024);
 
@@ -98,6 +191,9 @@ int main(int argc, char *argv[])
 	inputOptions.setMipmapGeneration(false);
 
 	nvtt::CompressionOptions compressionOptions;
+//	compressionOptions.setFormat(nvtt::Format_DXT1);
+//	compressionOptions.setFormat(nvtt::Format_DXT1n);
+	compressionOptions.setFormat(nvtt::Format_CTX1);
 	
 	nvtt::OutputOptions outputOptions;
 	outputOptions.setOutputHeader(false);

src/nvtt/tools/imgdiff.cpp

@@ -129,12 +129,15 @@ struct NormalError
 	}
 
 	void done()
+	{
+		if (samples)
 		{
 			ade /= samples;
 			mse /= samples * 3;
 			rmse = sqrt(mse);
 			psnr = (rmse == 0) ? 999.0f : 20.0f * log10(255.0f / rmse);
 		}
+	}
 
 	void print()
 	{