Replace broken compressor with the latest version from trunk.

Delete images more efficiently?

ChangeLog

@@ -1,3 +1,24 @@
+NVIDIA Texture Tools version 2.0.7
+ * Output correct exit codes. Fixes issue 92.
+ * Fix thread-safety errors. Fixes issue 90.
+ * Add SIMD power method. Fixes issue 94.
+ * Interact better with applications that already use CUDA.
+ * Faster CPU compression.
+
+NVIDIA Texture Tools version 2.0.6
+ * Fix dll version checking.
+ * Detect CUDA 2.1 and future CUDA versions correctly.
+ * Print CUDA detection message in nvcompress.
+ * Select the fastest CUDA device.
+ * Compile squish with -fPIC. Fixes issue 74.
+ * Fix warnings under gcc 4.3.2.
+ * Fix nvzoom option typo by Frank Richter. Fixes issue 81.
+ * Do not use CUDA to compress small mipmaps. Fixes issue 76.
+ * Compute mipmaps of semi-transparent images correctly.
+ * Shutdown CUDA properly. Fixes issue 83.
+ * Fix pixel format converions. Fixes issue 87.
+ * Update single color compression tables. Fixes issue 85.
+
 NVIDIA Texture Tools version 2.0.5
  * Fix error in single color compressor. Fixes issue 66.
  * Detect mismatch between CUDA runtime and driver, and disable CUDA in that case.

VERSION

@@ -1 +1 @@
-2.0.5
+2.0.7

project/vc8/nvtt/nvtt.rc

@@ -53,8 +53,8 @@ END
 //
 
 VS_VERSION_INFO VERSIONINFO
- FILEVERSION 2,0,5,0
- PRODUCTVERSION 2,0,5,0
+ FILEVERSION 2,0,6,0
+ PRODUCTVERSION 2,0,6,0
  FILEFLAGSMASK 0x17L
 #ifdef _DEBUG
  FILEFLAGS 0x1L
@@ -71,12 +71,12 @@ BEGIN
         BEGIN
             VALUE "CompanyName", "NVIDIA Corporation"
             VALUE "FileDescription", "NVIDIA Texture Tools Dynamic Link Library"
-            VALUE "FileVersion", "2, 0, 5, 0"
+            VALUE "FileVersion", "2, 0, 6, 0"
             VALUE "InternalName", "nvtt"
             VALUE "LegalCopyright", "Copyright (C) 2007"
             VALUE "OriginalFilename", "nvtt.dll"
             VALUE "ProductName", "NVIDIA Texture Tools Dynamic Link Library"
-            VALUE "ProductVersion", "2, 0, 5, 0"
+            VALUE "ProductVersion", "2, 0, 6, 0"
         END
     END
     BLOCK "VarFileInfo"

src/CMakeLists.txt

@@ -105,7 +105,8 @@ ENDIF(OPENEXR_FOUND)
 FIND_PACKAGE(Qt4)
 
 # Threads
-FIND_PACKAGE(Threads)
+FIND_PACKAGE(Threads REQUIRED)
+MESSAGE(STATUS "Use thread library: ${CMAKE_THREAD_LIBS_INIT}")
 
 # configuration file
 INCLUDE(CheckIncludeFiles)

src/nvcore/Containers.h

@@ -824,13 +824,13 @@ namespace nv
 		}
 	
 		/// Number of entries in the hash.
-		int	size()
+		int	size() const
 		{
 			return entry_count;
 		}
 	
 		/// Number of entries in the hash.
-		int	count()
+		int	count() const
 		{
 			return size();
 		}

src/nvcore/Debug.cpp

@@ -38,7 +38,7 @@
 #	include <unistd.h>	// getpid
 #	include <sys/types.h>
 #	include <sys/sysctl.h>	// sysctl
-#	include <ucontext.h>
+#	include <sys/ucontext.h>
 #	undef HAVE_EXECINFO_H
 #	if defined(HAVE_EXECINFO_H) // only after OSX 10.5
 #		include <execinfo.h> // backtrace
@@ -136,7 +136,11 @@ namespace
 #if defined(HAVE_EXECINFO_H) // NV_OS_LINUX
 
 	static bool nvHasStackTrace() {
+#if NV_OS_DARWIN
 		return backtrace != NULL;
+#else
+		return true;
+#endif
 	}
 
 	static void nvPrintStackTrace(void * trace[], int size, int start=0) {
@@ -401,7 +405,7 @@ namespace
 			{
 				void * trace[64];
 				int size = backtrace(trace, 64);
-				nvPrintStackTrace(trace, size, 3);
+				nvPrintStackTrace(trace, size, 2);
 			}
 #		endif
 

src/nvcore/Debug.h

@@ -115,6 +115,7 @@ namespace nv
 	{
 		NVCORE_API void dumpInfo();
 	
+		// These functions are not thread safe.
 		NVCORE_API void setMessageHandler( MessageHandler * messageHandler );
 		NVCORE_API void resetMessageHandler();
 	

src/nvcore/DefsVcWin32.h

@@ -72,8 +72,6 @@ typedef uint32              uint;
 #pragma warning(disable : 4711)		// function selected for automatic inlining
 #pragma warning(disable : 4725)		// Pentium fdiv bug
 
-#pragma warning(disable : 4345)		// behavior change: an object of POD type constructed with an initializer of the form () will be default-initialized
-
 #pragma warning(disable : 4786)		// Identifier was truncated and cannot be debugged.
 
 #pragma warning(disable : 4675)		// resolved overload was found by argument-dependent lookup

src/nvcore/StdStream.h

@@ -47,25 +47,25 @@ public:
 
 	/** @name Stream implementation. */
 	//@{
-		virtual void seek( int pos )
+		virtual void seek( uint pos )
 		{
 			nvDebugCheck(m_fp != NULL);
-			nvDebugCheck(pos >= 0 && pos < size());
+			nvDebugCheck(pos < size());
 			fseek(m_fp, pos, SEEK_SET);
 		}
 		
-		virtual int tell() const
+		virtual uint tell() const
 		{
 			nvDebugCheck(m_fp != NULL);
 			return ftell(m_fp);
 		}
 		
-		virtual int size() const
+		virtual uint size() const
 		{
 			nvDebugCheck(m_fp != NULL);
-			int pos = ftell(m_fp);
+			uint pos = ftell(m_fp);
 			fseek(m_fp, 0, SEEK_END);
-			int end = ftell(m_fp);
+			uint end = ftell(m_fp);
 			fseek(m_fp, pos, SEEK_SET);
 			return end;
 		}
@@ -117,11 +117,11 @@ public:
 	/** @name Stream implementation. */
 	//@{
 		/// Write data.
-		virtual void serialize( void * data, int len )
+		virtual uint serialize( void * data, uint len )
 		{
 			nvDebugCheck(data != NULL);
 			nvDebugCheck(m_fp != NULL);
-			fwrite(data, len, 1, m_fp);
+			return (uint)fwrite(data, 1, len, m_fp);
 		}
 		
 		virtual bool isLoading() const
@@ -156,11 +156,11 @@ public:
 	/** @name Stream implementation. */
 	//@{
 		/// Read data.
-		virtual void serialize( void * data, int len )
+		virtual uint serialize( void * data, uint len )
 		{
 			nvDebugCheck(data != NULL);
 			nvDebugCheck(m_fp != NULL);
-			fread(data, len, 1, m_fp);
+			return (uint)fread(data, 1, len, m_fp);
 		}
 		
 		virtual bool isLoading() const
@@ -184,33 +184,40 @@ class NVCORE_CLASS MemoryInputStream : public Stream
 public:
 
 	/// Ctor.
-	MemoryInputStream( const uint8 * mem, int size ) : 
+	MemoryInputStream( const uint8 * mem, uint size ) : 
 		m_mem(mem), m_ptr(mem), m_size(size) { }
 
 	/** @name Stream implementation. */
 	//@{
 		/// Read data.
-		virtual void serialize( void * data, int len )
+		virtual uint serialize( void * data, uint len )
 		{
 			nvDebugCheck(data != NULL);
 			nvDebugCheck(!isError());
+
+			uint left = m_size - tell();
+			if (len > left) len = left;
+			
 			memcpy( data, m_ptr, len );
 			m_ptr += len;
+				
+			return len;
 		}
 		
-		virtual void seek( int pos )
+		virtual void seek( uint pos )
 		{
 			nvDebugCheck(!isError());
 			m_ptr = m_mem + pos;
 			nvDebugCheck(!isError());
 		}
 		
-		virtual int tell() const
+		virtual uint tell() const
 		{
-			return int(m_ptr - m_mem);
+			nvDebugCheck(m_ptr >= m_mem);
+			return uint(m_ptr - m_mem);
 		}
 		
-		virtual int size() const
+		virtual uint size() const
 		{
 			return m_size;
 		}
@@ -252,7 +259,7 @@ private:
 
 	const uint8 * m_mem;
 	const uint8 * m_ptr;
-	int m_size;
+	uint m_size;
 
 };
 
@@ -286,17 +293,19 @@ public:
 	/** @name Stream implementation. */
 	//@{
 		/// Read data.
-		virtual void serialize( void * data, int len )
+		virtual uint serialize( void * data, uint len )
 		{
 			nvDebugCheck(data != NULL);
-			m_s->serialize( data, len );
+			len = m_s->serialize( data, len );
 			
 			if( m_s->isError() ) {
 				throw std::exception();
 			}
+
+			return len;
 		}
 		
-		virtual void seek( int pos )
+		virtual void seek( uint pos )
 		{
 			m_s->seek( pos );
 			
@@ -305,12 +314,12 @@ public:
 			}
 		}
 		
-		virtual int tell() const
+		virtual uint tell() const
 		{
 			return m_s->tell();
 		}
 		
-		virtual int size() const
+		virtual uint size() const
 		{
 			return m_s->size();
 		}

src/nvcore/StrLib.cpp

@@ -209,48 +209,11 @@ StringBuilder::StringBuilder( const StringBuilder & s ) : m_size(0), m_str(NULL)
 }
 
 /** Copy string. */
-StringBuilder::StringBuilder( const char * s )
+StringBuilder::StringBuilder( const char * s ) : m_size(0), m_str(NULL)
 {
 	copy(s);
 }
 
-/** Allocate and copy string. */
-StringBuilder::StringBuilder( int size_hint, const StringBuilder & s) : m_size(size_hint), m_str(NULL)
-{
-	nvDebugCheck(m_size > 0);
-	m_str = strAlloc(m_size);
-	copy(s);
-}
-
-/** Allocate and format string. */
-StringBuilder::StringBuilder( const char * fmt, ... ) : m_size(0), m_str(NULL)
-{
-	nvDebugCheck(fmt != NULL);
-	va_list arg;
-	va_start( arg, fmt );
-
-	format( fmt, arg );
-
-	va_end( arg );
-}
-
-/** Allocate and format string. */
-StringBuilder::StringBuilder( int size_hint, const char * fmt, ... ) : m_size(size_hint), m_str(NULL)
-{
-	nvDebugCheck(m_size > 0);	
-	nvDebugCheck(fmt != NULL);
-	
-	m_str = strAlloc(m_size);
-
-	va_list arg;
-	va_start( arg, fmt );
-
-	format( fmt, arg );
-
-	va_end( arg );
-}
-
-
 /** Delete the string. */
 StringBuilder::~StringBuilder()
 {
@@ -278,8 +241,7 @@ StringBuilder & StringBuilder::format( const char * fmt, ... )
 /** Format a string safely. */
 StringBuilder & StringBuilder::format( const char * fmt, va_list arg )
 {
-	nvCheck(fmt != NULL);
-	nvCheck(m_size >= 0);
+	nvDebugCheck(fmt != NULL);
 
 	if( m_size == 0 ) {
 		m_size = 64;
@@ -327,8 +289,7 @@ StringBuilder & StringBuilder::format( const char * fmt, va_list arg )
 /** Append a string. */
 StringBuilder & StringBuilder::append( const char * s )
 {
-	nvCheck(s != NULL);
-	nvCheck(m_size >= 0);
+	nvDebugCheck(s != NULL);
 
 	const uint slen = uint(strlen( s ));
 
@@ -475,31 +436,6 @@ void StringBuilder::reset()
 }
 
 
-Path::Path(const char * fmt, ...)
-{
-	nvDebugCheck( fmt != NULL );
-
-	va_list arg;
-	va_start( arg, fmt );
-
-	format( fmt, arg );
-
-	va_end( arg );
-}
-
-Path::Path(int size_hint, const char * fmt, ...) : StringBuilder(size_hint)
-{
-	nvDebugCheck( fmt != NULL );
-
-	va_list arg;
-	va_start( arg, fmt );
-
-	format( fmt, arg );
-
-	va_end( arg );
-}
-
-
 /// Get the file name from a path.
 const char * Path::fileName() const
 {
@@ -609,8 +545,6 @@ const char * Path::extension(const char * str)
 }
 
 
-// static
-String String::s_null(String::null);
 
 /// Clone this string
 String String::clone() const
@@ -622,13 +556,13 @@ String String::clone() const
 void String::setString(const char * str)
 {
 	if (str == NULL) {
-		data = s_null.data;
+		data = NULL;
 	}
 	else {
 		allocString( str );
-	}
 		addRef();
 	}
+}
 
 void String::setString(const char * str, int length)
 {
@@ -641,10 +575,10 @@ void String::setString(const char * str, int length)
 void String::setString(const StringBuilder & str)
 {
 	if (str.str() == NULL) {
-		data = s_null.data;
+		data =	NULL;
 	}
 	else {
 		allocString(str);
-	}
 		addRef();
 	}
+}	

src/nvcore/StrLib.h

@@ -47,9 +47,6 @@ namespace nv
 		explicit StringBuilder( int size_hint );
 		StringBuilder( const char * str );
 		StringBuilder( const StringBuilder & );
-		StringBuilder( int size_hint, const StringBuilder & );	
-		StringBuilder( const char * format, ... ) __attribute__((format (printf, 2, 3)));
-		StringBuilder( int size_hint, const char * format, ... ) __attribute__((format (printf, 3, 4)));
 	
 		~StringBuilder();
 	
@@ -122,16 +119,14 @@ namespace nv
 	};
 
 
-	/// Path string.
+	/// Path string. @@ This should be called PathBuilder.
 	class NVCORE_CLASS Path : public StringBuilder
 	{
 	public:
 		Path() : StringBuilder() {}
 		explicit Path(int size_hint) : StringBuilder(size_hint) {}
-		Path(const StringBuilder & str) : StringBuilder(str) {}
-		Path(int size_hint, const StringBuilder & str) : StringBuilder(size_hint, str) {}	
-		Path(const char * format, ...) __attribute__((format (printf, 2, 3)));
-		Path(int size_hint, const char * format, ...) __attribute__((format (printf, 3, 4)));
+		Path(const char * str) : StringBuilder(str) {}
+		Path(const Path & path) : StringBuilder(path) {}
 		
 		const char * fileName() const;
 		const char * extension() const;
@@ -156,15 +151,14 @@ namespace nv
 		/// Constructs a null string. @sa isNull()
 		String()
 		{
-			data = s_null.data;
-			addRef();
+			data = NULL;
 		}
 
 		/// Constructs a shared copy of str.
 		String(const String & str)
 		{
 			data = str.data;
-			addRef();
+			if (data != NULL) addRef();
 		}
 
 		/// Constructs a shared string from a standard string.
@@ -188,7 +182,6 @@ namespace nv
 		/// Dtor.
 		~String()
 		{
-			nvDebugCheck(data != NULL);
 			release();
 		}
 
@@ -225,43 +218,49 @@ namespace nv
 		/// Equal operator.
 		bool operator==( const String & str ) const
 		{
-			nvDebugCheck(data != NULL);
-			nvDebugCheck(str.data != NULL);
 			if( str.data == data ) {
 				return true;
 			}
+			if ((data == NULL) != (str.data == NULL)) {
+				return false;
+			}
 			return strcmp(data, str.data) == 0;
 		}
 
 		/// Equal operator.
 		bool operator==( const char * str ) const
 		{
-			nvDebugCheck(data != NULL);
 			nvCheck(str != NULL);	// Use isNull!
+			if (data == NULL) {
+				return false;
+			}
 			return strcmp(data, str) == 0;
 		}
 
 		/// Not equal operator.
 		bool operator!=( const String & str ) const
 		{
-			nvDebugCheck(data != NULL);
-			nvDebugCheck(str.data != NULL);
 			if( str.data == data ) {
 				return false;
 			}
+			if ((data == NULL) != (str.data == NULL)) {
+				return true;
+			}
 			return strcmp(data, str.data) != 0;
 		}
 	
 		/// Not equal operator.
 		bool operator!=( const char * str ) const
 		{
-			nvDebugCheck(data != NULL);
 			nvCheck(str != NULL);	// Use isNull!
+			if (data == NULL) {
+				return false;
+			}
 			return strcmp(data, str) != 0;
 		}
 	
 		/// Returns true if this string is the null string.
-		bool isNull() const { nvDebugCheck(data != NULL); return data == s_null.data; }
+		bool isNull() const { return data == NULL; }
 	
 		/// Return the exact length.
 		uint length() const { nvDebugCheck(data != NULL); return uint(strlen(data)); }
@@ -270,31 +269,28 @@ namespace nv
 		uint hash() const { nvDebugCheck(data != NULL); return strHash(data); }
 	
 		/// const char * cast operator.
-		operator const char * () const { nvDebugCheck(data != NULL); return data; }
+		operator const char * () const { return data; }
 	
 		/// Get string pointer.
-		const char * str() const { nvDebugCheck(data != NULL); return data; }
+		const char * str() const { return data; }
 	
 
 	private:
 
-		enum null_t { null };
-		
-		// Private constructor for null string.
-		String(null_t) {
-			setString("");
-		}
-
 		// Add reference count.
-		void addRef() {
-			nvDebugCheck(data != NULL);
+		void addRef()
+		{
+			if (data != NULL)
+			{
 				setRefCount(getRefCount() + 1);
 			}
+		}
 		
 		// Decrease reference count.
-		void release() {
-			nvDebugCheck(data != NULL);
-
+		void release()
+		{
+			if (data != NULL)
+			{
 				const uint16 count = getRefCount();
 				setRefCount(count - 1);
 				if (count - 1 == 0) {
@@ -302,12 +298,16 @@ namespace nv
 					data = NULL;
 				}
 			}
+		}
 		
-		uint16 getRefCount() const {
+		uint16 getRefCount() const
+		{
+			nvDebugCheck(data != NULL);
 			return *reinterpret_cast<const uint16 *>(data - 2);
 		}
 		
 		void setRefCount(uint16 count) {
+			nvDebugCheck(data != NULL);
 			nvCheck(count < 0xFFFF);
 			*reinterpret_cast<uint16 *>(const_cast<char *>(data - 2)) = uint16(count);
 		}
@@ -346,8 +346,6 @@ namespace nv
 	
 	private:
 
-		NVCORE_API static String s_null;
-
 		const char * data;
 		
 	};

src/nvcore/Stream.h

@@ -41,17 +41,17 @@ public:
 	ByteOrder byteOrder() const { return m_byteOrder; }
 
 	
-	/// Serialize the given data. @@ Should return bytes serialized?
-	virtual void serialize( void * data, int len ) = 0;
+	/// Serialize the given data.
+	virtual uint serialize( void * data, uint len ) = 0;
 
 	/// Move to the given position in the archive.
-	virtual void seek( int pos ) = 0;
+	virtual void seek( uint pos ) = 0;
 
 	/// Return the current position in the archive.
-	virtual int tell() const = 0;
+	virtual uint tell() const = 0;
 
 	/// Return the current size of the archive.
-	virtual int size() const = 0;
+	virtual uint size() const = 0;
 
 	/// Determine if there has been any error.
 	virtual bool isError() const = 0;
@@ -136,13 +136,13 @@ public:
 protected:
 
 	/// Serialize in the stream byte order.
-	Stream & byteOrderSerialize( void * v, int len ) {
+	Stream & byteOrderSerialize( void * v, uint len ) {
 		if( m_byteOrder == getSystemByteOrder() ) {
 			serialize( v, len );
 		}
 		else {
-			for( int i=len-1; i>=0; i-- ) {
-				serialize( (uint8 *)v + i, 1 );
+			for( uint i = len; i > 0; i-- ) {
+				serialize( (uint8 *)v + i - 1, 1 );
 			}
 		}
 		return *this;

src/nvimage/DirectDrawSurface.cpp

@@ -532,7 +532,7 @@ DDSHeader::DDSHeader()
 
 	// Store version information on the reserved header attributes.
 	this->reserved[9] = MAKEFOURCC('N', 'V', 'T', 'T');
-	this->reserved[10] = (2 << 16) | (0 << 8) | (5);	// major.minor.revision
+	this->reserved[10] = (2 << 16) | (0 << 8) | (7);	// major.minor.revision
 
 	this->pf.size = 32;
 	this->pf.flags = 0;
@@ -989,10 +989,10 @@ void DirectDrawSurface::readLinearImage(Image * img)
 			stream->serialize(&c, byteCount);
 
 			Color32 pixel(0, 0, 0, 0xFF);
-			pixel.r = PixelFormat::convert(c >> rshift, rsize, 8);
-			pixel.g = PixelFormat::convert(c >> gshift, gsize, 8);
-			pixel.b = PixelFormat::convert(c >> bshift, bsize, 8);
-			pixel.a = PixelFormat::convert(c >> ashift, asize, 8);
+			pixel.r = PixelFormat::convert((c & header.pf.rmask) >> rshift, rsize, 8);
+			pixel.g = PixelFormat::convert((c & header.pf.gmask) >> gshift, gsize, 8);
+			pixel.b = PixelFormat::convert((c & header.pf.bmask) >> bshift, bsize, 8);
+			pixel.a = PixelFormat::convert((c & header.pf.amask) >> ashift, asize, 8);
 
 			img->pixel(x, y) = pixel;
 		}

src/nvimage/Filter.cpp

@@ -33,11 +33,10 @@
  * http://www.dspguide.com/ch16.htm
  */
 
+#include "Filter.h"
 
-#include <nvcore/Containers.h>	// swap
-#include <nvmath/nvmath.h>	// fabs
 #include <nvmath/Vector.h>	// Vector4
-#include <nvimage/Filter.h>
+#include <nvcore/Containers.h>	// swap
 
 using namespace nv;
 
@@ -504,7 +503,7 @@ void Kernel2::initBlendedSobel(const Vector4 & scale)
 
 		for (int i = 0; i < 7; i++) {
 			for (int e = 0; e < 7; e++) {
-				m_data[i * 9 + e + 1] += elements[i * 7 + e] * scale.z();
+				m_data[(i + 1) * 9 + e + 1] += elements[i * 7 + e] * scale.z();
 			}
 		}
 	}
@@ -519,7 +518,7 @@ void Kernel2::initBlendedSobel(const Vector4 & scale)
 
 		for (int i = 0; i < 5; i++) {
 			for (int e = 0; e < 5; e++) {
-				m_data[i * 9 + e + 2] += elements[i * 5 + e] * scale.y();
+				m_data[(i + 2) * 9 + e + 2] += elements[i * 5 + e] * scale.y();
 			}
 		}
 	}
@@ -532,7 +531,7 @@ void Kernel2::initBlendedSobel(const Vector4 & scale)
 
 		for (int i = 0; i < 3; i++) {
 			for (int e = 0; e < 3; e++) {
-				m_data[i * 9 + e + 3] += elements[i * 3 + e] * scale.x();
+				m_data[(i + 3) * 9 + e + 3] += elements[i * 3 + e] * scale.x();
 			}
 		}
 	}
@@ -582,7 +581,6 @@ PolyphaseKernel::PolyphaseKernel(const Filter & f, uint srcLength, uint dstLengt
 			m_data[i * m_windowSize + j] /= total;
 		}
 	}
-
 }
 
 PolyphaseKernel::~PolyphaseKernel()

src/nvimage/Filter.h

@@ -11,16 +11,16 @@ namespace nv
 	class Vector4;
 
 	/// Base filter class.
-	class Filter
+	class NVIMAGE_CLASS Filter
 	{
 	public:
-		NVIMAGE_API Filter(float width);
-		NVIMAGE_API virtual ~Filter();
+		Filter(float width);
+		virtual ~Filter();
 
-		NVIMAGE_API float width() const { return m_width; }
-		NVIMAGE_API float sampleDelta(float x, float scale) const;
-		NVIMAGE_API float sampleBox(float x, float scale, int samples) const;
-		NVIMAGE_API float sampleTriangle(float x, float scale, int samples) const;
+		float width() const { return m_width; }
+		float sampleDelta(float x, float scale) const;
+		float sampleBox(float x, float scale, int samples) const;
+		float sampleTriangle(float x, float scale, int samples) const;
 
 		virtual float evaluate(float x) const = 0;
 
@@ -29,56 +29,56 @@ namespace nv
 	};
 
 	// Box filter.
-	class BoxFilter : public Filter
+	class NVIMAGE_CLASS BoxFilter : public Filter
 	{
 	public:
-		NVIMAGE_API BoxFilter();
-		NVIMAGE_API BoxFilter(float width);
-		NVIMAGE_API virtual float evaluate(float x) const;
+		BoxFilter();
+		BoxFilter(float width);
+		virtual float evaluate(float x) const;
 	};
 
 	// Triangle (bilinear/tent) filter.
-	class TriangleFilter : public Filter
+	class NVIMAGE_CLASS TriangleFilter : public Filter
 	{
 	public:
-		NVIMAGE_API TriangleFilter();
-		NVIMAGE_API TriangleFilter(float width);
-		NVIMAGE_API virtual float evaluate(float x) const;
+		TriangleFilter();
+		TriangleFilter(float width);
+		virtual float evaluate(float x) const;
 	};
 
 	// Quadratic (bell) filter.
-	class QuadraticFilter : public Filter
+	class NVIMAGE_CLASS QuadraticFilter : public Filter
 	{
 	public:
-		NVIMAGE_API QuadraticFilter();
-		NVIMAGE_API virtual float evaluate(float x) const;
+		QuadraticFilter();
+		virtual float evaluate(float x) const;
 	};
 
 	// Cubic filter from Thatcher Ulrich.
-	class CubicFilter : public Filter
+	class NVIMAGE_CLASS CubicFilter : public Filter
 	{
 	public:
-		NVIMAGE_API CubicFilter();
-		NVIMAGE_API virtual float evaluate(float x) const;
+		CubicFilter();
+		virtual float evaluate(float x) const;
 	};
 
 	// Cubic b-spline filter from Paul Heckbert.
-	class BSplineFilter : public Filter
+	class NVIMAGE_CLASS BSplineFilter : public Filter
 	{
 	public:
-		NVIMAGE_API BSplineFilter();
-		NVIMAGE_API virtual float evaluate(float x) const;
+		BSplineFilter();
+		virtual float evaluate(float x) const;
 	};
 
 	/// Mitchell & Netravali's two-param cubic
 	/// @see "Reconstruction Filters in Computer Graphics", SIGGRAPH 88
-	class MitchellFilter : public Filter
+	class NVIMAGE_CLASS MitchellFilter : public Filter
 	{
 	public:
-		NVIMAGE_API MitchellFilter();
-		NVIMAGE_API virtual float evaluate(float x) const;
+		MitchellFilter();
+		virtual float evaluate(float x) const;
 
-		NVIMAGE_API void setParameters(float a, float b);
+		void setParameters(float b, float c);
 
 	private:
 		float p0, p2, p3;
@@ -86,29 +86,29 @@ namespace nv
 	};
 
 	// Lanczos3 filter.
-	class LanczosFilter : public Filter
+	class NVIMAGE_CLASS LanczosFilter : public Filter
 	{
 	public:
-		NVIMAGE_API LanczosFilter();
-		NVIMAGE_API virtual float evaluate(float x) const;
+		LanczosFilter();
+		virtual float evaluate(float x) const;
 	};
 
 	// Sinc filter.
-	class SincFilter : public Filter
+	class NVIMAGE_CLASS SincFilter : public Filter
 	{
 	public:
-		NVIMAGE_API SincFilter(float w);
-		NVIMAGE_API virtual float evaluate(float x) const;
+		SincFilter(float w);
+		virtual float evaluate(float x) const;
 	};
 
 	// Kaiser filter.
-	class KaiserFilter : public Filter
+	class NVIMAGE_CLASS KaiserFilter : public Filter
 	{
 	public:
-		NVIMAGE_API KaiserFilter(float w);
-		NVIMAGE_API virtual float evaluate(float x) const;
+		KaiserFilter(float w);
+		virtual float evaluate(float x) const;
 	
-		NVIMAGE_API void setParameters(float a, float stretch);
+		void setParameters(float a, float stretch);
 
 	private:
 		float alpha;
@@ -118,12 +118,12 @@ namespace nv
 
 
 	/// A 1D kernel. Used to precompute filter weights.
-	class Kernel1
+	class NVIMAGE_CLASS Kernel1
 	{
 		NV_FORBID_COPY(Kernel1);
 	public:
-		NVIMAGE_API Kernel1(const Filter & f, int iscale, int samples = 32);
-		NVIMAGE_API ~Kernel1();
+		Kernel1(const Filter & f, int iscale, int samples = 32);
+		~Kernel1();
 		
 		float valueAt(uint x) const {
 			nvDebugCheck(x < (uint)m_windowSize);
@@ -138,7 +138,7 @@ namespace nv
 			return m_width;
 		}
 		
-		NVIMAGE_API void debugPrint();
+		void debugPrint();
 		
 	private:
 		int m_windowSize;
@@ -148,15 +148,15 @@ namespace nv
 
 
 	/// A 2D kernel.
-	class Kernel2 
+	class NVIMAGE_CLASS Kernel2 
 	{
 	public:
-		NVIMAGE_API Kernel2(uint width);
-		NVIMAGE_API Kernel2(const Kernel2 & k);
-		NVIMAGE_API ~Kernel2();
+		Kernel2(uint width);
+		Kernel2(const Kernel2 & k);
+		~Kernel2();
 		
-		NVIMAGE_API void normalize();
-		NVIMAGE_API void transpose();
+		void normalize();
+		void transpose();
 		
 		float valueAt(uint x, uint y) const {
 			return m_data[y * m_windowSize + x];
@@ -166,12 +166,12 @@ namespace nv
 			return m_windowSize;
 		}
 		
-		NVIMAGE_API void initLaplacian();
-		NVIMAGE_API void initEdgeDetection();
-		NVIMAGE_API void initSobel();
-		NVIMAGE_API void initPrewitt();
+		void initLaplacian();
+		void initEdgeDetection();
+		void initSobel();
+		void initPrewitt();
 		
-		NVIMAGE_API void initBlendedSobel(const Vector4 & scale);
+		void initBlendedSobel(const Vector4 & scale);
 		
 	private:
 		const uint m_windowSize;
@@ -180,12 +180,12 @@ namespace nv
 
 
 	/// A 1D polyphase kernel
-	class PolyphaseKernel
+	class NVIMAGE_CLASS PolyphaseKernel
 	{
 		NV_FORBID_COPY(PolyphaseKernel);
 	public:
-		NVIMAGE_API PolyphaseKernel(const Filter & f, uint srcLength, uint dstLength, int samples = 32);
-		NVIMAGE_API ~PolyphaseKernel();
+		PolyphaseKernel(const Filter & f, uint srcLength, uint dstLength, int samples = 32);
+		~PolyphaseKernel();
 		
 		int windowSize() const {
 			return m_windowSize;
@@ -205,7 +205,7 @@ namespace nv
 			return m_data[column * m_windowSize + x];
 		}
 
-		NVIMAGE_API void debugPrint() const;
+		void debugPrint() const;
 		
 	private:
 		int m_windowSize;

src/nvimage/FloatImage.cpp

@@ -1,16 +1,18 @@
 // This code is in the public domain -- castanyo@yahoo.es
 
-#include <nvcore/Containers.h>
-#include <nvcore/Ptr.h>
-
-#include <nvmath/Color.h>
-
 #include "FloatImage.h"
 #include "Filter.h"
 #include "Image.h"
 
+#include <nvmath/Color.h>
+#include <nvmath/Matrix.h>
+
+#include <nvcore/Containers.h>
+#include <nvcore/Ptr.h>
+
 #include <math.h>
 
+
 using namespace nv;
 
 namespace 
@@ -140,7 +142,8 @@ Image * FloatImage::createImageGammaCorrect(float gamma/*= 2.2f*/) const
 /// Allocate a 2d float image of the given format and the given extents.
 void FloatImage::allocate(uint c, uint w, uint h)
 {
-	nvCheck(m_mem == NULL);
+	free();
+
 	m_width = w;
 	m_height = h;
 	m_componentNum = c;
@@ -151,7 +154,6 @@ void FloatImage::allocate(uint c, uint w, uint h)
 /// Free the image, but don't clear the members.
 void FloatImage::free()
 {
-	nvCheck(m_mem != NULL);
 	nv::mem::free( reinterpret_cast<void *>(m_mem) );
 	m_mem = NULL;
 }
@@ -549,6 +551,15 @@ FloatImage * FloatImage::downSample(const Filter & filter, WrapMode wm) const
 	return resize(filter, w, h, wm);
 }
 
+/// Downsample applying a 1D kernel separately in each dimension.
+FloatImage * FloatImage::downSample(const Filter & filter, WrapMode wm, uint alpha) const
+{
+	const uint w = max(1, m_width / 2);
+	const uint h = max(1, m_height / 2);
+
+	return resize(filter, w, h, wm, alpha);
+}
+
 
 /// Downsample applying a 1D kernel separately in each dimension.
 FloatImage * FloatImage::resize(const Filter & filter, uint w, uint h, WrapMode wm) const
@@ -620,10 +631,56 @@ FloatImage * FloatImage::resize(const Filter & filter, uint w, uint h, WrapMode
 	return dst_image.release();
 }
 
+/// Downsample applying a 1D kernel separately in each dimension.
+FloatImage * FloatImage::resize(const Filter & filter, uint w, uint h, WrapMode wm, uint alpha) const
+{
+	nvCheck(alpha < m_componentNum);
+
+	AutoPtr<FloatImage> tmp_image( new FloatImage() );
+	AutoPtr<FloatImage> dst_image( new FloatImage() );	
+	
+	PolyphaseKernel xkernel(filter, m_width, w, 32);
+	PolyphaseKernel ykernel(filter, m_height, h, 32);
+	
+	{
+		tmp_image->allocate(m_componentNum, w, m_height);
+		dst_image->allocate(m_componentNum, w, h);
+		
+		Array<float> tmp_column(h);
+		tmp_column.resize(h);
+		
+		for (uint c = 0; c < m_componentNum; c++)
+		{
+			float * tmp_channel = tmp_image->channel(c);
+			
+			for (uint y = 0; y < m_height; y++) {
+				this->applyKernelHorizontal(xkernel, y, c, alpha, wm, tmp_channel + y * w);
+			}
+		}
+
+		// Process all channels before applying vertical kernel to make sure alpha has been computed.
+
+		for (uint c = 0; c < m_componentNum; c++)
+		{
+			float * dst_channel = dst_image->channel(c);
+			
+			for (uint x = 0; x < w; x++) {
+				tmp_image->applyKernelVertical(ykernel, x, c, alpha, wm, tmp_column.unsecureBuffer());
+				
+				for (uint y = 0; y < h; y++) {
+					dst_channel[y * w + x] = tmp_column[y];
+				}
+			}
+		}
+	}
+	
+	return dst_image.release();
+}
+
 
 
 /// Apply 2D kernel at the given coordinates and return result.
-float FloatImage::applyKernel(const Kernel2 * k, int x, int y, int c, WrapMode wm) const
+float FloatImage::applyKernel(const Kernel2 * k, int x, int y, uint c, WrapMode wm) const
 {
 	nvDebugCheck(k != NULL);
 	
@@ -652,7 +709,7 @@ float FloatImage::applyKernel(const Kernel2 * k, int x, int y, int c, WrapMode w
 
 
 /// Apply 1D vertical kernel at the given coordinates and return result.
-float FloatImage::applyKernelVertical(const Kernel1 * k, int x, int y, int c, WrapMode wm) const
+float FloatImage::applyKernelVertical(const Kernel1 * k, int x, int y, uint c, WrapMode wm) const
 {
 	nvDebugCheck(k != NULL);
 	
@@ -674,7 +731,7 @@ float FloatImage::applyKernelVertical(const Kernel1 * k, int x, int y, int c, Wr
 }
 
 /// Apply 1D horizontal kernel at the given coordinates and return result.
-float FloatImage::applyKernelHorizontal(const Kernel1 * k, int x, int y, int c, WrapMode wm) const
+float FloatImage::applyKernelHorizontal(const Kernel1 * k, int x, int y, uint c, WrapMode wm) const
 {
 	nvDebugCheck(k != NULL);
 	
@@ -697,7 +754,7 @@ float FloatImage::applyKernelHorizontal(const Kernel1 * k, int x, int y, int c,
 
 
 /// Apply 1D vertical kernel at the given coordinates and return result.
-void FloatImage::applyKernelVertical(const PolyphaseKernel & k, int x, int c, WrapMode wm, float * output) const
+void FloatImage::applyKernelVertical(const PolyphaseKernel & k, int x, uint c, WrapMode wm, float * __restrict output) const
 {
 	const uint length = k.length();
 	const float scale = float(length) / float(m_height);
@@ -729,7 +786,7 @@ void FloatImage::applyKernelVertical(const PolyphaseKernel & k, int x, int c, Wr
 }
 
 /// Apply 1D horizontal kernel at the given coordinates and return result.
-void FloatImage::applyKernelHorizontal(const PolyphaseKernel & k, int y, int c, WrapMode wm, float * output) const
+void FloatImage::applyKernelHorizontal(const PolyphaseKernel & k, int y, uint c, WrapMode wm, float * __restrict output) const
 {
 	const uint length = k.length();
 	const float scale = float(length) / float(m_width);
@@ -760,3 +817,93 @@ void FloatImage::applyKernelHorizontal(const PolyphaseKernel & k, int y, int c,
 	}
 }
 
+
+/// Apply 1D vertical kernel at the given coordinates and return result.
+void FloatImage::applyKernelVertical(const PolyphaseKernel & k, int x, uint c, uint a, WrapMode wm, float * __restrict output) const
+{
+	const uint length = k.length();
+	const float scale = float(length) / float(m_height);
+	const float iscale = 1.0f / scale;
+
+	const float width = k.width();
+	const int windowSize = k.windowSize();
+
+	const float * channel = this->channel(c);
+	const float * alpha = this->channel(a);
+
+	for (uint i = 0; i < length; i++)
+	{
+		const float center = (0.5f + i) * iscale;
+		
+		const int left = (int)floorf(center - width);
+		const int right = (int)ceilf(center + width);
+		nvCheck(right - left <= windowSize);
+		
+		float norm = 0;
+		float sum = 0;
+		for (int j = 0; j < windowSize; ++j)
+		{
+			const int idx = this->index(x, j+left, wm);
+			
+			float w = k.valueAt(i, j) * (alpha[idx] + (1.0f / 256.0f));
+			norm += w;
+			sum += w * channel[idx];
+		}
+		
+		output[i] = sum / norm;
+	}
+}
+
+/// Apply 1D horizontal kernel at the given coordinates and return result.
+void FloatImage::applyKernelHorizontal(const PolyphaseKernel & k, int y, uint c, uint a, WrapMode wm, float * __restrict output) const
+{
+	const uint length = k.length();
+	const float scale = float(length) / float(m_width);
+	const float iscale = 1.0f / scale;
+
+	const float width = k.width();
+	const int windowSize = k.windowSize();
+
+	const float * channel = this->channel(c);
+	const float * alpha = this->channel(a);
+
+	for (uint i = 0; i < length; i++)
+	{
+		const float center = (0.5f + i) * iscale;
+		
+		const int left = (int)floorf(center - width);
+		const int right = (int)ceilf(center + width);
+		nvDebugCheck(right - left <= windowSize);
+		
+		float norm = 0.0f;
+		float sum = 0;
+		for (int j = 0; j < windowSize; ++j)
+		{
+			const int idx = this->index(left + j, y, wm);
+
+			float w = k.valueAt(i, j) * (alpha[idx] + (1.0f / 256.0f));
+			norm += w;
+			sum += w * channel[idx];
+		}
+		
+		output[i] = sum / norm;
+	}
+}
+
+FloatImage* FloatImage::clone() const
+{
+	FloatImage* copy = new FloatImage();
+	copy->m_width = m_width;
+	copy->m_height = m_height;
+	copy->m_componentNum = m_componentNum;
+	copy->m_count = m_count;
+	
+	if(m_mem)
+	{
+		copy->allocate(m_componentNum, m_width, m_height);
+		memcpy(copy->m_mem, m_mem, m_count * sizeof(float));
+	}
+		
+	return copy;
+}
+

src/nvimage/FloatImage.h

@@ -3,12 +3,20 @@
 #ifndef NV_IMAGE_FLOATIMAGE_H
 #define NV_IMAGE_FLOATIMAGE_H
 
+#include <nvimage/nvimage.h>
+
+#include <nvmath/Vector.h>
+
 #include <nvcore/Debug.h>
 #include <nvcore/Containers.h> // clamp
-#include <nvimage/nvimage.h>
+
+#include <stdlib.h> // abs
+
 
 namespace nv
 {
+class Vector4;
+class Matrix;
 class Image;
 class Filter;
 class Kernel1;
@@ -63,17 +71,19 @@ public:
 
 	NVIMAGE_API FloatImage * fastDownSample() const;
 	NVIMAGE_API FloatImage * downSample(const Filter & filter, WrapMode wm) const;
+	NVIMAGE_API FloatImage * downSample(const Filter & filter, WrapMode wm, uint alpha) const;
 	NVIMAGE_API FloatImage * resize(const Filter & filter, uint w, uint h, WrapMode wm) const;
 
-	//NVIMAGE_API FloatImage * downSample(const Kernel1 & filter, WrapMode wm) const;
-	//NVIMAGE_API FloatImage * downSample(const Kernel1 & filter, uint w, uint h, WrapMode wm) const;
+	NVIMAGE_API FloatImage * resize(const Filter & filter, uint w, uint h, WrapMode wm, uint alpha) const;
 	//@}
 
-	NVIMAGE_API float applyKernel(const Kernel2 * k, int x, int y, int c, WrapMode wm) const;
-	NVIMAGE_API float applyKernelVertical(const Kernel1 * k, int x, int y, int c, WrapMode wm) const;
-	NVIMAGE_API float applyKernelHorizontal(const Kernel1 * k, int x, int y, int c, WrapMode wm) const;
-	NVIMAGE_API void applyKernelVertical(const PolyphaseKernel & k, int x, int c, WrapMode wm, float * output) const;
-	NVIMAGE_API void applyKernelHorizontal(const PolyphaseKernel & k, int y, int c, WrapMode wm, float * output) const;
+	NVIMAGE_API float applyKernel(const Kernel2 * k, int x, int y, uint c, WrapMode wm) const;
+	NVIMAGE_API float applyKernelVertical(const Kernel1 * k, int x, int y, uint c, WrapMode wm) const;
+	NVIMAGE_API float applyKernelHorizontal(const Kernel1 * k, int x, int y, uint c, WrapMode wm) const;
+	NVIMAGE_API void applyKernelVertical(const PolyphaseKernel & k, int x, uint c, WrapMode wm, float * output) const;
+	NVIMAGE_API void applyKernelHorizontal(const PolyphaseKernel & k, int y, uint c, WrapMode wm, float * output) const;
+	NVIMAGE_API void applyKernelVertical(const PolyphaseKernel & k, int x, uint c, uint a, WrapMode wm, float * output) const;
+	NVIMAGE_API void applyKernelHorizontal(const PolyphaseKernel & k, int y, uint c, uint a, WrapMode wm, float * output) const;
 	
 	
 	uint width() const { return m_width; }
@@ -109,6 +119,9 @@ public:
 	float sampleLinearMirror(float x, float y, int c) const;
 	//@}
 	
+	
+	FloatImage* clone() const;
+	
 public:
 	
 	uint index(uint x, uint y) const;

src/nvimage/Image.cpp

@@ -78,7 +78,7 @@ void Image::unwrap()
 
 void Image::free()
 {
-	::free(m_data);
+	nv::mem::free(m_data);
 	m_data = NULL;
 }
 

src/nvimage/NormalMap.cpp

@@ -21,15 +21,16 @@
 // FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
 // OTHER DEALINGS IN THE SOFTWARE.
 
-#include <nvcore/Ptr.h>
-
-#include <nvmath/Color.h>
-
 #include <nvimage/NormalMap.h>
 #include <nvimage/Filter.h>
 #include <nvimage/FloatImage.h>
 #include <nvimage/Image.h>
 
+#include <nvmath/Color.h>
+
+#include <nvcore/Ptr.h>
+
+
 using namespace nv;
 
 // Create normal map using the given kernels.

src/nvimage/PsdFile.h

@@ -39,7 +39,7 @@ namespace nv
 		bool isSupported() const
 		{
 			if (version != 1) {
-				printf("*** bad version number %u\n", version);
+				nvDebug("*** bad version number %u\n", version);
 				return false;
 			}
 			if (channel_count > 4) {

src/nvimage/Quantize.cpp

@@ -12,11 +12,14 @@ http://www.efg2.com/Lab/Library/ImageProcessing/DHALF.TXT
 @@ This code needs to be reviewed, I'm not sure it's correct.
 */
 
+#include <nvimage/Quantize.h>
+#include <nvimage/Image.h>
+#include <nvimage/PixelFormat.h>
+
 #include <nvmath/Color.h>
 
-#include <nvimage/Image.h>
-#include <nvimage/Quantize.h>
-#include <nvimage/PixelFormat.h>
+#include <nvcore/Containers.h> // swap
+
 
 using namespace nv;
 

src/nvimage/Quantize.h

@@ -3,6 +3,9 @@
 #ifndef NV_IMAGE_QUANTIZE_H
 #define NV_IMAGE_QUANTIZE_H
 
+#include <nvimage/nvimage.h>
+
+
 namespace nv
 {
 	class Image;

src/nvmath/Matrix.h

@@ -332,7 +332,7 @@ inline Matrix transpose(Matrix::Arg m)
 	Matrix r;
 	for (int i = 0; i < 4; i++)
 	{
-		for (int j = 0; j < 4; i++)
+		for (int j = 0; j < 4; j++)
 		{
 			r(i, j) = m(j, i);
 		}

src/nvmath/nvmath.h

@@ -48,19 +48,37 @@
 #define IS_NEGATIVE_FLOAT(x)	(IR(x)&SIGN_BITMASK)
 */
 
-inline float sqrt_assert(const float f)
+inline double sqrt_assert(const double f)
+{
+	nvDebugCheck(f >= 0.0f);
+	return sqrt(f);
+}
+
+inline float sqrtf_assert(const float f)
 {
 	nvDebugCheck(f >= 0.0f);
 	return sqrtf(f);
 }
 
-inline float acos_assert(const float f)
+inline double acos_assert(const double f)
+{
+	nvDebugCheck(f >= -1.0f && f <= 1.0f);
+	return acos(f);
+}
+
+inline float acosf_assert(const float f)
 {
 	nvDebugCheck(f >= -1.0f && f <= 1.0f);
 	return acosf(f);
 }
 
-inline float asin_assert(const float f)
+inline double asin_assert(const double f)
+{
+	nvDebugCheck(f >= -1.0f && f <= 1.0f);
+	return asin(f);
+}
+
+inline float asinf_assert(const float f)
 {
 	nvDebugCheck(f >= -1.0f && f <= 1.0f);
 	return asinf(f);
@@ -68,11 +86,11 @@ inline float asin_assert(const float f)
 
 // Replace default functions with asserting ones.
 #define sqrt sqrt_assert
-#define sqrtf sqrt_assert
+#define sqrtf sqrtf_assert
 #define acos acos_assert
-#define acosf acos_assert
+#define acosf acosf_assert
 #define asin asin_assert
-#define asinf asin_assert
+#define asinf asinf_assert
 
 #if NV_OS_WIN32
 #include <float.h>
@@ -136,6 +154,11 @@ inline float lerp(float f0, float f1, float t)
 	return f0 * s + f1 * t;
 }
 
+inline float square(float f)
+{
+	return f * f;
+}
+
 } // nv
 
 #endif // NV_MATH_H

src/nvtt/CompressDXT.cpp

@@ -205,9 +205,9 @@ void nv::SlowCompressor::compressDXT1(const CompressionOptions::Private & compre
 	ColorBlock rgba;
 	BlockDXT1 block;
 
-	//squish::WeightedClusterFit fit;
+	squish::WeightedClusterFit fit;
 	//squish::ClusterFit fit;
-	squish::FastClusterFit fit;
+	//squish::FastClusterFit fit;
 	fit.SetMetric(compressionOptions.colorWeight.x(), compressionOptions.colorWeight.y(), compressionOptions.colorWeight.z());
 
 	for (uint y = 0; y < h; y += 4) {
@@ -221,7 +221,7 @@ void nv::SlowCompressor::compressDXT1(const CompressionOptions::Private & compre
 			}
 			else
 			{
-				squish::ColourSet colours((uint8 *)rgba.colors(), 0);
+				squish::ColourSet colours((uint8 *)rgba.colors(), 0, true);
 				fit.SetColourSet(&colours, squish::kDxt1);
 				fit.Compress(&block);
 			}

src/nvtt/Compressor.cpp

@@ -207,26 +207,15 @@ Compressor::Compressor() : m(*new Compressor::Private())
 {
 	// CUDA initialization.
 	m.cudaSupported = cuda::isHardwarePresent();
-	m.cudaEnabled = m.cudaSupported;
-
-	if (m.cudaEnabled)
-	{
-		// Select fastest CUDA device.
-		int device = cuda::getFastestDevice();
-		cuda::setDevice(device);
-		
-		m.cuda = new CudaCompressor();
-
-		if (!m.cuda->isValid())
-		{
 	m.cudaEnabled = false;
-			m.cuda = NULL;
-		}
-	}
+	m.cudaDevice = -1;
+
+	enableCudaAcceleration(m.cudaSupported);
 }
 
 Compressor::~Compressor()
 {
+	enableCudaAcceleration(false);
 	delete &m;
 }
 
@@ -236,21 +225,33 @@ void Compressor::enableCudaAcceleration(bool enable)
 {
 	if (m.cudaSupported)
 	{
-		m.cudaEnabled = enable;
-	}
-
-	if (m.cudaEnabled && m.cuda == NULL)
-	{
-		// Select fastest CUDA device.
-		int device = cuda::getFastestDevice();
-		cuda::setDevice(device);
-		
-		m.cuda = new CudaCompressor();
-		
-		if (!m.cuda->isValid())
+		if (m.cudaEnabled && !enable)
 		{
 			m.cudaEnabled = false;
 			m.cuda = NULL;
+
+			if (m.cudaDevice != -1)
+			{
+				// Exit device.
+				cuda::exitDevice();
+			}
+		}
+		else if (!m.cudaEnabled && enable)
+		{
+			// Init the CUDA device. This may return -1 if CUDA was already initialized by the app.
+			m.cudaEnabled = cuda::initDevice(&m.cudaDevice);
+
+			if (m.cudaEnabled)
+			{
+				// Create compressor if initialization succeeds.
+				m.cuda = new CudaCompressor();
+
+				// But cleanup if failed.
+				if (!m.cuda->isValid())
+				{
+					enableCudaAcceleration(false);
+				}
+			}
 		}
 	}
 }
@@ -697,6 +698,7 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const InputOptio
 	SlowCompressor slow;
 	slow.setImage(image, inputOptions.alphaMode);
 
+	const bool useCuda = cudaEnabled && image->width() * image->height() >= 512;
 
 	if (compressionOptions.format == Format_RGBA || compressionOptions.format == Format_RGB)
 	{
@@ -725,7 +727,7 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const InputOptio
 				}
 				else
 				{
-			if (cudaEnabled)
+					if (useCuda)
 					{
 						nvDebugCheck(cudaSupported);
 						cuda->setImage(image, inputOptions.alphaMode);
@@ -745,7 +747,7 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const InputOptio
 		}
 		else
 		{
-			if (cudaEnabled)
+			if (useCuda)
 			{
 				nvDebugCheck(cudaSupported);
 				/*cuda*/slow.compressDXT1a(compressionOptions, outputOptions);
@@ -764,7 +766,7 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const InputOptio
 		}
 		else
 		{
-			if (cudaEnabled)
+			if (useCuda)
 			{
 				nvDebugCheck(cudaSupported);
 				cuda->setImage(image, inputOptions.alphaMode);
@@ -784,7 +786,7 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const InputOptio
 		}
 		else
 		{
-			if (cudaEnabled)
+			if (useCuda)
 			{
 				nvDebugCheck(cudaSupported);
 				cuda->setImage(image, inputOptions.alphaMode);

src/nvtt/Compressor.h

@@ -63,10 +63,12 @@ namespace nvtt
 		bool compressMipmap(const Mipmap & mipmap, const InputOptions::Private & inputOptions, const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions) const;
 
 
+
 	public:
 
 		bool cudaSupported;
 		bool cudaEnabled;
+		int cudaDevice;
 
 		nv::AutoPtr<nv::CudaCompressor> cuda;
 

src/nvtt/InputOptions.cpp

@@ -94,7 +94,7 @@ void InputOptions::reset()
 	m.textureType = TextureType_2D;
 	m.inputFormat = InputFormat_BGRA_8UB;
 
-	m.alphaMode = AlphaMode_Transparency;
+	m.alphaMode = AlphaMode_None;
 
 	m.inputGamma = 2.2f;
 	m.outputGamma = 2.2f;

src/nvtt/cuda/CudaCompressDXT.cpp

@@ -37,7 +37,7 @@
 
 
 #if defined HAVE_CUDA
-#include <cuda_runtime.h>
+#include <cuda_runtime_api.h>
 #endif
 
 #include <time.h>

src/nvtt/cuda/CudaMath.h

@@ -148,7 +148,7 @@ inline __device__ bool singleColor(const float3 * colors)
 	bool sameColor = false;
 	for (int i = 0; i < 16; i++)
 	{
-		sameColor &= (colors[idx] == colors[0]);
+		sameColor &= (colors[i] == colors[0]);
 	}
 	return sameColor;
 #else

src/nvtt/cuda/CudaUtils.cpp

@@ -26,12 +26,14 @@
 #include "CudaUtils.h"
 
 #if defined HAVE_CUDA
-#include <cuda_runtime.h>
+#include <cuda.h>
+#include <cuda_runtime_api.h>
 #endif
 
 using namespace nv;
 using namespace cuda;
 
+/* @@ Move this to win32 utils or somewhere else.
 #if NV_OS_WIN32
 
 #define WINDOWS_LEAN_AND_MEAN
@@ -68,10 +70,12 @@ static bool isWow32()
 }
 
 #endif
+*/
 
 
-static bool isCudaDriverAvailable(uint version)
+static bool isCudaDriverAvailable(int version)
 {
+#if defined HAVE_CUDA
 #if NV_OS_WIN32
 	Library nvcuda("nvcuda.dll");
 #else
@@ -80,20 +84,48 @@ static bool isCudaDriverAvailable(uint version)
 
 	if (!nvcuda.isValid())
 	{
+		nvDebug("*** CUDA driver not found.\n");
 		return false;
 	}
 
-	if (version > 2000)
+	if (version >= 2000)
 	{
 		void * address = nvcuda.bindSymbol("cuStreamCreate");
-		if (address == NULL) return false;
+		if (address == NULL) {
+			nvDebug("*** CUDA driver version < 2.0.\n");
+			return false;
+		}
 	}
 
-	if (version > 2010)
+	if (version >= 2010)
 	{
-		void * address = nvcuda.bindSymbol("cuLoadDataEx");
-		if (address == NULL) return false;
+		void * address = nvcuda.bindSymbol("cuModuleLoadDataEx");
+		if (address == NULL) {
+			nvDebug("*** CUDA driver version < 2.1.\n");
+			return false;
 		}
+	}
+
+	if (version >= 2020)
+	{
+		typedef CUresult (CUDAAPI * PFCU_DRIVERGETVERSION)(int * version);
+
+		PFCU_DRIVERGETVERSION driverGetVersion = (PFCU_DRIVERGETVERSION)nvcuda.bindSymbol("cuDriverGetVersion");
+		if (driverGetVersion == NULL) {
+			nvDebug("*** CUDA driver version < 2.2.\n");
+			return false;
+		}
+
+		int driverVersion;
+		CUresult err = driverGetVersion(&driverVersion);
+		if (err != CUDA_SUCCESS) {
+			nvDebug("*** Error querying driver version: '%s'.\n", cudaGetErrorString((cudaError_t)err));
+			return false;
+		}
+
+		return driverVersion >= version;
+	}
+#endif // HAVE_CUDA
 
 	return true;
 }
@@ -103,10 +135,13 @@ static bool isCudaDriverAvailable(uint version)
 bool nv::cuda::isHardwarePresent()
 {
 #if defined HAVE_CUDA
-#if NV_OS_WIN32
-	//if (isWindowsVista()) return false;
-	//if (isWindowsVista() || !isWow32()) return false;
-#endif
+	// Make sure that CUDA driver matches CUDA runtime.
+	if (!isCudaDriverAvailable(CUDART_VERSION))
+	{
+		nvDebug("CUDA driver not available for CUDA runtime %d\n", CUDART_VERSION);
+		return false;
+	}
+
 	int count = deviceCount();
 	if (count == 1)
 	{
@@ -119,15 +154,11 @@ bool nv::cuda::isHardwarePresent()
 		{
 			return false;
 		}
-
-		// Make sure that CUDA driver matches CUDA runtime.
-		if (!isCudaDriverAvailable(CUDART_VERSION))
-		{
-			return false;
 	}
 
 	// @@ Make sure that warp size == 32
-	}
+
+	// @@ Make sure available GPU is faster than the CPU.
 
 	return count > 0;
 #else
@@ -151,43 +182,119 @@ int nv::cuda::deviceCount()
 	return 0;
 }
 
+
+// Make sure device meets requirements:
+// - Not an emulation device.
+// - Not an integrated device?
+// - Faster than CPU.
+bool nv::cuda::isValidDevice(int i)
+{
+#if defined HAVE_CUDA
+
+	cudaDeviceProp device_properties;
+	cudaGetDeviceProperties(&device_properties, i);
+	int gflops = device_properties.multiProcessorCount * device_properties.clockRate;
+
+	if (device_properties.major == -1 || device_properties.minor == -1) {
+		// Emulation device.
+		return false;
+	}
+
+#if CUDART_VERSION >= 2030 // 2.3
+	/*if (device_properties.integrated)
+	{
+		// Integrated devices.
+		return false;
+	}*/
+#endif
+
+	return true;
+#else
+	return false;
+#endif
+}
+
 int nv::cuda::getFastestDevice()
 {
-	int max_gflops_device = 0;
+	int max_gflops_device = -1;
 #if defined HAVE_CUDA
 	int max_gflops = 0;
 
 	const int device_count = deviceCount();
-	int current_device = 0;
-	while (current_device < device_count)
+	for (int i = 0; i < device_count; i++)
+	{
+		if (isValidDevice(i))
 		{
 			cudaDeviceProp device_properties;
-		cudaGetDeviceProperties(&device_properties, current_device);
+			cudaGetDeviceProperties(&device_properties, i);
 			int gflops = device_properties.multiProcessorCount * device_properties.clockRate;
 
-		if (device_properties.major != -1 && device_properties.minor != -1)
-		{
 			if (gflops > max_gflops)
 			{
 				max_gflops = gflops;
-				max_gflops_device = current_device;
+				max_gflops_device = i;
 			}
 		}
-		
-		current_device++;
 	}
 #endif
 	return max_gflops_device;
 }
 
+
 /// Activate the given devices.
-bool nv::cuda::setDevice(int i)
+bool nv::cuda::initDevice(int * device_ptr)
 {
-	nvCheck(i < deviceCount());
+	nvDebugCheck(device_ptr != NULL);
 #if defined HAVE_CUDA
-	cudaError_t result = cudaSetDevice(i);
-	return result == cudaSuccess;
+
+#if CUDART_VERSION >= 2030 // 2.3
+
+	// Set device flags to yield in order to play nice with other threads and to find out if CUDA was already active.
+	cudaError_t resul = cudaSetDeviceFlags(cudaDeviceScheduleYield);
+
+#endif
+
+	int device = getFastestDevice();
+
+	if (device == -1)
+	{
+		// No device is fast enough.
+		*device_ptr = -1;
+		return false;
+	}
+
+	// Select CUDA device.
+	cudaError_t result = cudaSetDevice(device);
+
+	if (result == cudaErrorSetOnActiveProcess)
+	{
+		int device;
+		result = cudaGetDevice(&device);
+
+		*device_ptr = -1;  // No device to cleanup.
+		return isValidDevice(device); // Return true if device is valid.
+	}
+	else if (result != cudaSuccess)
+	{
+		nvDebug("*** CUDA Error: %s\n", cudaGetErrorString(result));
+		*device_ptr = -1;
+		return false;
+	}
+
+	*device_ptr = device;
+	return true;
 #else
 	return false;
 #endif
 }
+
+void nv::cuda::exitDevice()
+{
+#if defined HAVE_CUDA
+	cudaError_t result = cudaThreadExit();
+
+	if (result != cudaSuccess) {
+		nvDebug("*** CUDA Error: %s\n", cudaGetErrorString(result));
+	}
+#endif
+}

src/nvtt/cuda/CudaUtils.h

@@ -32,7 +32,10 @@ namespace nv
 		bool isHardwarePresent();
 		int deviceCount();
 		int getFastestDevice();
-		bool setDevice(int i);
+		bool isValidDevice(int i);
+
+		bool initDevice(int * device_ptr);
+		void exitDevice();
 	};
 	
 } // nv namespace

src/nvtt/nvtt.h

@@ -73,7 +73,7 @@ namespace nvtt
 		Format_DXT1a,   // DXT1 with binary alpha.
 		Format_DXT3,
 		Format_DXT5,
-		Format_DXT5n,   // Compressed HILO: R=0, G=x, B=0, A=y
+		Format_DXT5n,   // Compressed HILO: R=1, G=y, B=0, A=x
 		
 		// DX10 formats.
 		Format_BC1 = Format_DXT1,
@@ -194,7 +194,7 @@ namespace nvtt
 		// Describe the format of the input.
 		NVTT_API void setFormat(InputFormat format);
 		
-		// Set the way the input alpha channel is interpreted. @@ Not implemented!
+		// Set the way the input alpha channel is interpreted.
 		NVTT_API void setAlphaMode(AlphaMode alphaMode);
 		
 		// Set gamma settings.

src/nvtt/squish/CMakeLists.txt

@@ -1,13 +1,8 @@
 PROJECT(squish)
-ENABLE_TESTING()
 
 INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR})
 
 SET(SQUISH_SRCS
-#	alpha.cpp
-#	alpha.h
-#	clusterfit.cpp
-#	clusterfit.h
 	fastclusterfit.cpp
 	fastclusterfit.h
 	weightedclusterfit.cpp
@@ -21,32 +16,13 @@ SET(SQUISH_SRCS
 	config.h
 	maths.cpp
 	maths.h
-#	rangefit.cpp
-#	rangefit.h
-#	singlecolourfit.cpp
-#	singlecolourfit.h
-#	singlecolourlookup.inl
-#	squish.cpp
-#	squish.h
 	simd.h
 	simd_sse.h
 	simd_ve.h)
 
 ADD_LIBRARY(squish STATIC ${SQUISH_SRCS})
 
-# libpng
-#FIND_PACKAGE(PNG)
-
-#IF(PNG_FOUND)
-#	INCLUDE_DIRECTORIES(${PNG_INCLUDE_DIR})
-#	ADD_EXECUTABLE(squishpng extra/squishpng.cpp)
-#	TARGET_LINK_LIBRARIES(squishpng squish ${PNG_LIBRARY})
-#ENDIF(PNG_FOUND)
-
-##ADD_EXECUTABLE(squishgen extra/squishgen.cpp)
-
-#ADD_EXECUTABLE(squishtest extra/squishtest.cpp)
-#TARGET_LINK_LIBRARIES(squishtest squish)
-
-#ADD_TEST(SQUISHTEST squishtest)
+IF(CMAKE_COMPILER_IS_GNUCXX)
+	SET_TARGET_PROPERTIES(squish PROPERTIES COMPILE_FLAGS -fPIC)
+ENDIF(CMAKE_COMPILER_IS_GNUCXX)
 

src/nvtt/squish/maths.cpp

@@ -24,6 +24,7 @@
    -------------------------------------------------------------------------- */
    
 #include "maths.h"
+#include "simd.h"
 #include <cfloat>
 
 namespace squish {
@@ -60,12 +61,39 @@ Sym3x3 ComputeWeightedCovariance( int n, Vec3 const* points, float const* weight
 }
 
 
+#define POWER_ITERATION_COUNT   8
+
+#if SQUISH_USE_SIMD
+
 Vec3 ComputePrincipleComponent( Sym3x3 const& matrix )
 {
-	const int NUM = 8;
+	Vec4 const row0( matrix[0], matrix[1], matrix[2], 0.0f );
+	Vec4 const row1( matrix[1], matrix[3], matrix[4], 0.0f );
+	Vec4 const row2( matrix[2], matrix[4], matrix[5], 0.0f );
+	Vec4 v = VEC4_CONST( 1.0f );
+	for( int i = 0; i < POWER_ITERATION_COUNT; ++i )
+	{
+		// matrix multiply
+		Vec4 w = row0*v.SplatX();
+		w = MultiplyAdd(row1, v.SplatY(), w);
+		w = MultiplyAdd(row2, v.SplatZ(), w);
 
+		// get max component from xyz in all channels
+		Vec4 a = Max(w.SplatX(), Max(w.SplatY(), w.SplatZ()));
+
+		// divide through and advance
+		v = w*Reciprocal(a);
+	}
+	return v.GetVec3();
+}
+
+#else
+
+Vec3 ComputePrincipleComponent( Sym3x3 const& matrix )
+{
 	Vec3 v(1, 1, 1);
-	for(int i = 0; i < NUM; i++) {
+	for (int i = 0; i < POWER_ITERATION_COUNT; i++)
+	{
 		float x = v.X() * matrix[0] + v.Y() * matrix[1] + v.Z() * matrix[2];
 		float y = v.X() * matrix[1] + v.Y() * matrix[3] + v.Z() * matrix[4];
 		float z = v.X() * matrix[2] + v.Y() * matrix[4] + v.Z() * matrix[5];
@@ -82,5 +110,6 @@ Vec3 ComputePrincipleComponent( Sym3x3 const& matrix )
 	return v;
 }
 
+#endif
 
 } // namespace squish

src/nvtt/squish/weightedclusterfit.cpp

@@ -131,10 +131,13 @@ float WeightedClusterFit::GetBestError() const
 
 	void WeightedClusterFit::Compress3( void* block )
 	{
+		int const count = m_colours->GetCount();
 		Vec4 const one = VEC4_CONST(1.0f);
 		Vec4 const zero = VEC4_CONST(0.0f);
 		Vec4 const half(0.5f, 0.5f, 0.5f, 0.25f);
 		Vec4 const two = VEC4_CONST(2.0);
+		Vec4 const grid( 31.0f, 63.0f, 31.0f, 0.0f );
+		Vec4 const gridrcp( 1.0f/31.0f, 1.0f/63.0f, 1.0f/31.0f, 0.0f );
 
 		// declare variables
 		Vec4 beststart = VEC4_CONST( 0.0f );
@@ -146,11 +149,11 @@ void WeightedClusterFit::Compress3( void* block )
 		int b0 = 0, b1 = 0;
 
 		// check all possible clusters for this total order
-	for( int c0 = 0; c0 <= 16; c0++)
+		for( int c0 = 0; c0 <= count; c0++)
 		{	
 			Vec4 x1 = zero;
 
-		for( int c1 = 0; c1 <= 16-c0; c1++)
+			for( int c1 = 0; c1 <= count-c0; c1++)
 			{
 				Vec4 const x2 = m_xsum - x1 - x0;
 
@@ -173,24 +176,21 @@ void WeightedClusterFit::Compress3( void* block )
 				Vec4 a = NegativeMultiplySubtract(betax_sum, alphabeta_sum, alphax_sum*beta2_sum) * factor;
 				Vec4 b = NegativeMultiplySubtract(alphax_sum, alphabeta_sum, betax_sum*alpha2_sum) * factor;
 
-			// clamp the output to [0, 1]
+				// clamp to the grid
 				a = Min( one, Max( zero, a ) );
 				b = Min( one, Max( zero, b ) );
-			
-			// clamp to the grid
-			Vec4 const grid( 31.0f, 63.0f, 31.0f, 0.0f );
-			Vec4 const gridrcp( 0.03227752766457f, 0.01583151765563f, 0.03227752766457f, 0.0f );
 				a = Truncate( MultiplyAdd( grid, a, half ) ) * gridrcp;
 				b = Truncate( MultiplyAdd( grid, b, half ) ) * gridrcp;
 
-			// compute the error
-			Vec4 e1 = MultiplyAdd( a, alphax_sum, b*betax_sum );
-			Vec4 e2 = MultiplyAdd( a*a, alpha2_sum, b*b*beta2_sum );
-			Vec4 e3 = MultiplyAdd( a*b*alphabeta_sum - e1, two, e2 );
+				// compute the error (we skip the constant xxsum)
+				Vec4 e1 = MultiplyAdd( a*a, alpha2_sum, b*b*beta2_sum );
+				Vec4 e2 = NegativeMultiplySubtract( a, alphax_sum, a*b*alphabeta_sum );
+				Vec4 e3 = NegativeMultiplySubtract( b, betax_sum, e2 );
+				Vec4 e4 = MultiplyAdd( two, e3, e1 );
 
 				// apply the metric to the error term
-			Vec4 e4 = e3 * m_metricSqr;
-			Vec4 error = e4.SplatX() + e4.SplatY() + e4.SplatZ();
+				Vec4 e5 = e4 * m_metricSqr;
+				Vec4 error = e5.SplatX() + e5.SplatY() + e5.SplatZ();
 
 				// keep the solution if it wins
 				if( CompareAnyLessThan( error, besterror ) )
@@ -221,17 +221,17 @@ void WeightedClusterFit::Compress3( void* block )
 				for(; i < b0+b1; i++) {
 					bestindices[i] = 2;
 				}
-			for(; i < 16; i++) {
+				for(; i < count; i++) {
 					bestindices[i] = 1;
 				}
 			}
 
 			// remap the indices
 			u8 ordered[16];
-		for( int i = 0; i < 16; ++i )
+			for( int i = 0; i < count; ++i )
 				ordered[m_order[i]] = bestindices[i];
 
-		m_colours->RemapIndices( ordered, bestindices ); // Set alpha indices.
+			m_colours->RemapIndices( ordered, bestindices );
 
 
 			// save the block
@@ -244,12 +244,16 @@ void WeightedClusterFit::Compress3( void* block )
 
 	void WeightedClusterFit::Compress4( void* block )
 	{
+		int const count = m_colours->GetCount();
 		Vec4 const one = VEC4_CONST(1.0f);
 		Vec4 const zero = VEC4_CONST(0.0f);
 		Vec4 const half = VEC4_CONST(0.5f);
 		Vec4 const two = VEC4_CONST(2.0);
 		Vec4 const onethird( 1.0f/3.0f, 1.0f/3.0f, 1.0f/3.0f, 1.0f/9.0f );
 		Vec4 const twothirds( 2.0f/3.0f, 2.0f/3.0f, 2.0f/3.0f, 4.0f/9.0f );
+		Vec4 const twonineths = VEC4_CONST( 2.0f/9.0f );
+		Vec4 const grid( 31.0f, 63.0f, 31.0f, 0.0f );
+		Vec4 const gridrcp( 1.0f/31.0f, 1.0f/63.0f, 1.0f/31.0f, 0.0f );
 
 		// declare variables
 		Vec4 beststart = VEC4_CONST( 0.0f );
@@ -260,30 +264,30 @@ void WeightedClusterFit::Compress4( void* block )
 		int b0 = 0, b1 = 0, b2 = 0;
 
 		// check all possible clusters for this total order
-	for( int c0 = 0; c0 <= 16; c0++)
+		for( int c0 = 0; c0 <= count; c0++)
 		{	
 			Vec4 x1 = zero;
 
-		for( int c1 = 0; c1 <= 16-c0; c1++)
+			for( int c1 = 0; c1 <= count-c0; c1++)
 			{	
 				Vec4 x2 = zero;
 
-			for( int c2 = 0; c2 <= 16-c0-c1; c2++)
+				for( int c2 = 0; c2 <= count-c0-c1; c2++)
 				{
 					Vec4 const x3 = m_xsum - x2 - x1 - x0;
 
 					//Vec3 const alphax_sum = x0 + x1 * (2.0f / 3.0f) + x2 * (1.0f / 3.0f);
 					//float const alpha2_sum = w0 + w1 * (4.0f/9.0f) + w2 * (1.0f/9.0f);
-				Vec4 const alphax_sum = x0 + MultiplyAdd(x1, twothirds, x2 * onethird); // alphax_sum, alpha2_sum
+					Vec4 const alphax_sum = MultiplyAdd(x2, onethird, MultiplyAdd(x1, twothirds, x0)); // alphax_sum, alpha2_sum
 					Vec4 const alpha2_sum = alphax_sum.SplatW();
 
 					//Vec3 const betax_sum = x3 + x2 * (2.0f / 3.0f) + x1 * (1.0f / 3.0f);
 					//float const beta2_sum = w3 + w2 * (4.0f/9.0f) + w1 * (1.0f/9.0f);
-				Vec4 const betax_sum = x3 + MultiplyAdd(x2, twothirds, x1 * onethird); // betax_sum, beta2_sum
+					Vec4 const betax_sum = MultiplyAdd(x2, twothirds, MultiplyAdd(x1, onethird, x3)); // betax_sum, beta2_sum
 					Vec4 const beta2_sum = betax_sum.SplatW();
 
-				//float const alphabeta_sum = w1 * (2.0f/9.0f) + w2 * (2.0f/9.0f);
-				Vec4 const alphabeta_sum = two * (x1 * onethird + x2 * onethird).SplatW(); // alphabeta_sum
+					//float const alphabeta_sum = (w1 + w2) * (2.0f/9.0f);
+					Vec4 const alphabeta_sum = twonineths*( x1 + x2 ).SplatW(); // alphabeta_sum
 
 					// float const factor = 1.0f / (alpha2_sum * beta2_sum - alphabeta_sum * alphabeta_sum);
 					Vec4 const factor = Reciprocal( NegativeMultiplySubtract(alphabeta_sum, alphabeta_sum, alpha2_sum*beta2_sum) );
@@ -291,24 +295,21 @@ void WeightedClusterFit::Compress4( void* block )
 					Vec4 a = NegativeMultiplySubtract(betax_sum, alphabeta_sum, alphax_sum*beta2_sum) * factor;
 					Vec4 b = NegativeMultiplySubtract(alphax_sum, alphabeta_sum, betax_sum*alpha2_sum) * factor;
 
-				// clamp the output to [0, 1]
+					// clamp to the grid
 					a = Min( one, Max( zero, a ) );
 					b = Min( one, Max( zero, b ) );
-				
-				// clamp to the grid
-				Vec4 const grid( 31.0f, 63.0f, 31.0f, 0.0f );
-				Vec4 const gridrcp( 0.03227752766457f, 0.01583151765563f, 0.03227752766457f, 0.0f );
 					a = Truncate( MultiplyAdd( grid, a, half ) ) * gridrcp;
 					b = Truncate( MultiplyAdd( grid, b, half ) ) * gridrcp;
 
-				// compute the error
-				Vec4 e1 = MultiplyAdd( a, alphax_sum, b*betax_sum );
-				Vec4 e2 = MultiplyAdd( a*a, alpha2_sum, b*b*beta2_sum );
-				Vec4 e3 = MultiplyAdd( a*b*alphabeta_sum - e1, two, e2 );
+					// compute the error (we skip the constant xxsum)
+					Vec4 e1 = MultiplyAdd( a*a, alpha2_sum, b*b*beta2_sum );
+					Vec4 e2 = NegativeMultiplySubtract( a, alphax_sum, a*b*alphabeta_sum );
+					Vec4 e3 = NegativeMultiplySubtract( b, betax_sum, e2 );
+					Vec4 e4 = MultiplyAdd( two, e3, e1 );
 
 					// apply the metric to the error term
-				Vec4 e4 = e3 * m_metricSqr;
-				Vec4 error = e4.SplatX() + e4.SplatY() + e4.SplatZ();
+					Vec4 e5 = e4 * m_metricSqr;
+					Vec4 error = e5.SplatX() + e5.SplatY() + e5.SplatZ();
 
 					// keep the solution if it wins
 					if( CompareAnyLessThan( error, besterror ) )
@@ -346,18 +347,20 @@ void WeightedClusterFit::Compress4( void* block )
 				for(; i < b0+b1+b2; i++) {
 					bestindices[i] = 3;
 				}
-			for(; i < 16; i++) {
+				for(; i < count; i++) {
 					bestindices[i] = 1;
 				}
 			}
 
 			// remap the indices
 			u8 ordered[16];
-		for( int i = 0; i < 16; ++i )
+			for( int i = 0; i < count; ++i )
 				ordered[m_order[i]] = bestindices[i];
 
+			m_colours->RemapIndices( ordered, bestindices );
+
 			// save the block
-		WriteColourBlock4( beststart.GetVec3(), bestend.GetVec3(), ordered, block );
+			WriteColourBlock4( beststart.GetVec3(), bestend.GetVec3(), bestindices, block );
 
 			// save the error
 			m_besterror = besterror;
@@ -368,6 +371,13 @@ void WeightedClusterFit::Compress4( void* block )
 
 	void WeightedClusterFit::Compress3( void* block )
 	{
+		int const count = m_colours->GetCount();
+		Vec3 const one( 1.0f );
+		Vec3 const zero( 0.0f );
+		Vec3 const half( 0.5f );
+		Vec3 const grid( 31.0f, 63.0f, 31.0f );
+		Vec3 const gridrcp( 1.0f/31.0f, 1.0f/63.0f, 1.0f/31.0f );
+
 		// declare variables
 		Vec3 beststart( 0.0f );
 		Vec3 bestend( 0.0f );
@@ -379,12 +389,12 @@ void WeightedClusterFit::Compress3( void* block )
 		int b0 = 0, b1 = 0;
 
 		// check all possible clusters for this total order
-	for( int c0 = 0; c0 <= 16; c0++)
+		for( int c0 = 0; c0 <= count; c0++)
 		{	
 			Vec3 x1(0.0f);
 			float w1 = 0.0f;
 
-		for( int c1 = 0; c1 <= 16-c0; c1++)
+			for( int c1 = 0; c1 <= count-c0; c1++)
 			{	
 				float w2 = m_wsum - w0 - w1;
 
@@ -400,16 +410,9 @@ void WeightedClusterFit::Compress3( void* block )
 				Vec3 a = (alphax_sum*beta2_sum - betax_sum*alphabeta_sum) * factor;
 				Vec3 b = (betax_sum*alpha2_sum - alphax_sum*alphabeta_sum) * factor;
 
-			// clamp the output to [0, 1]
-			Vec3 const one( 1.0f );
-			Vec3 const zero( 0.0f );
+				// clamp to the grid
 				a = Min( one, Max( zero, a ) );
 				b = Min( one, Max( zero, b ) );
-			
-			// clamp to the grid
-			Vec3 const grid( 31.0f, 63.0f, 31.0f );
-			Vec3 const gridrcp( 0.03227752766457f, 0.01583151765563f, 0.03227752766457f );
-			Vec3 const half( 0.5f );
 				a = Floor( grid*a + half )*gridrcp;
 				b = Floor( grid*b + half )*gridrcp;
 
@@ -450,18 +453,20 @@ void WeightedClusterFit::Compress3( void* block )
 				for(; i < b0+b1; i++) {
 					bestindices[i] = 2;
 				}
-			for(; i < 16; i++) {
+				for(; i < count; i++) {
 					bestindices[i] = 1;
 				}
 			}
 
 			// remap the indices
 			u8 ordered[16];
-		for( int i = 0; i < 16; ++i )
+			for( int i = 0; i < count; ++i )
 				ordered[m_order[i]] = bestindices[i];
 
+			m_colours->RemapIndices( ordered, bestindices );
+
 			// save the block
-		WriteColourBlock3( beststart, bestend, ordered, block );
+			WriteColourBlock3( beststart, bestend, bestindices, block );
 
 			// save the error
 			m_besterror = besterror;
@@ -470,6 +475,13 @@ void WeightedClusterFit::Compress3( void* block )
 
 	void WeightedClusterFit::Compress4( void* block )
 	{
+		int const count = m_colours->GetCount();
+		Vec3 const one( 1.0f );
+		Vec3 const zero( 0.0f );
+		Vec3 const half( 0.5f );
+		Vec3 const grid( 31.0f, 63.0f, 31.0f );
+		Vec3 const gridrcp( 1.0f/31.0f, 1.0f/63.0f, 1.0f/31.0f );
+
 		// declare variables
 		Vec3 beststart( 0.0f );
 		Vec3 bestend( 0.0f );
@@ -480,17 +492,17 @@ void WeightedClusterFit::Compress4( void* block )
 		int b0 = 0, b1 = 0, b2 = 0;
 
 		// check all possible clusters for this total order
-	for( int c0 = 0; c0 <= 16; c0++)
+		for( int c0 = 0; c0 <= count; c0++)
 		{	
 			Vec3 x1(0.0f);
 			float w1 = 0.0f;
 
-		for( int c1 = 0; c1 <= 16-c0; c1++)
+			for( int c1 = 0; c1 <= count-c0; c1++)
 			{	
 				Vec3 x2(0.0f);
 				float w2 = 0.0f;
 
-			for( int c2 = 0; c2 <= 16-c0-c1; c2++)
+				for( int c2 = 0; c2 <= count-c0-c1; c2++)
 				{
 					float w3 = m_wsum - w0 - w1 - w2;
 
@@ -505,16 +517,9 @@ void WeightedClusterFit::Compress4( void* block )
 					Vec3 a = ( alphax_sum*beta2_sum - betax_sum*alphabeta_sum )*factor;
 					Vec3 b = ( betax_sum*alpha2_sum - alphax_sum*alphabeta_sum )*factor;
 
-				// clamp the output to [0, 1]
-				Vec3 const one( 1.0f );
-				Vec3 const zero( 0.0f );
+					// clamp to the grid
 					a = Min( one, Max( zero, a ) );
 					b = Min( one, Max( zero, b ) );
-				
-				// clamp to the grid
-				Vec3 const grid( 31.0f, 63.0f, 31.0f );
-				Vec3 const gridrcp( 0.03227752766457f, 0.01583151765563f, 0.03227752766457f );
-				Vec3 const half( 0.5f );
 					a = Floor( grid*a + half )*gridrcp;
 					b = Floor( grid*b + half )*gridrcp;
 
@@ -563,18 +568,20 @@ void WeightedClusterFit::Compress4( void* block )
 				for(; i < b0+b1+b2; i++) {
 					bestindices[i] = 3;
 				}
-			for(; i < 16; i++) {
+				for(; i < count; i++) {
 					bestindices[i] = 1;
 				}
 			}
 
 			// remap the indices
 			u8 ordered[16];
-		for( int i = 0; i < 16; ++i )
+			for( int i = 0; i < count; ++i )
 				ordered[m_order[i]] = bestindices[i];
 
+			m_colours->RemapIndices( ordered, bestindices );
+
 			// save the block
-		WriteColourBlock4( beststart, bestend, ordered, block );
+			WriteColourBlock4( beststart, bestend, bestindices, block );
 
 			// save the error
 			m_besterror = besterror;

src/nvtt/tools/compress.cpp

@@ -87,7 +87,10 @@ struct MyErrorHandler : public nvtt::ErrorHandler
 {
 	virtual void error(nvtt::Error e)
 	{
+#if _DEBUG
 		nvDebugBreak();
+#endif
+		printf("Error: '%s'\n", nvtt::errorString(e));
 	}
 };
 
@@ -254,7 +257,12 @@ int main(int argc, char *argv[])
 		}
 	}
 
-	printf("NVIDIA Texture Tools - Copyright NVIDIA Corporation 2007\n\n");
+	const uint version = nvtt::version();
+	const uint major = version / 100;
+	const uint minor = version % 100;
+	
+
+	printf("NVIDIA Texture Tools %u.%u - Copyright NVIDIA Corporation 2007\n\n", major, minor);
 
 	if (input.isNull())
 	{
@@ -281,7 +289,7 @@ int main(int argc, char *argv[])
 		printf("  -bc4     \tBC4 format (ATI1)\n");
 		printf("  -bc5     \tBC5 format (3Dc/ATI2)\n\n");
 		
-		return 1;
+		return EXIT_FAILURE;
 	}
 
 	// @@ Make sure input file exists.
@@ -296,13 +304,13 @@ int main(int argc, char *argv[])
 		if (!dds.isValid())
 		{
 			fprintf(stderr, "The file '%s' is not a valid DDS file.\n", input.str());
-			return 1;
+			return EXIT_FAILURE;
 		}
 		
 		if (!dds.isSupported() || dds.isTexture3D())
 		{
 			fprintf(stderr, "The file '%s' is not a supported DDS file.\n", input.str());
-			return 1;
+			return EXIT_FAILURE;
 		}
 		
 		uint faceCount;
@@ -339,7 +347,7 @@ int main(int argc, char *argv[])
 		if (!image.load(input))
 		{
 			fprintf(stderr, "The file '%s' is not a supported image type.\n", input.str());
-			return 1;
+			return EXIT_FAILURE;
 		}
 		
 		inputOptions.setTextureLayout(nvtt::TextureType_2D, image.width(), image.height());
@@ -396,33 +404,28 @@ int main(int argc, char *argv[])
 		compressionOptions.setExternalCompressor(externalCompressor);
 	}
 
-	if (format == nvtt::Format_RGB)
-	{
-		compressionOptions.setQuantization(true, false, false);
-		//compressionOptions.setPixelFormat(16, 0xF000, 0x0F00, 0x00F0, 0x000F);
-		compressionOptions.setPixelFormat(16, 
-			0x0F00, 
-			0x00F0, 
-			0x000F, 
-			0xF000);
-		//	0x003F0000, 
-		//	0x00003F00, 
-		//	0x0000003F, 
-		//	0x3F000000);
-	}
-
 	
 	MyErrorHandler errorHandler;
 	MyOutputHandler outputHandler(output);
 	if (outputHandler.stream->isError())
 	{
 		fprintf(stderr, "Error opening '%s' for writting\n", output.str());
-		return 1;
+		return EXIT_FAILURE;
 	}
 
 	nvtt::Compressor compressor;
 	compressor.enableCudaAcceleration(!nocuda);
 
+	printf("CUDA acceleration ");
+	if (compressor.isCudaAccelerationEnabled())
+	{
+		printf("ENABLED\n\n");
+	}
+	else
+	{
+		printf("DISABLED\n\n");
+	}
+	
 	outputHandler.setTotal(compressor.estimateSize(inputOptions, compressionOptions));
 	outputHandler.setDisplayProgress(!silent);
 
@@ -435,27 +438,16 @@ int main(int argc, char *argv[])
 //	fflush(stdout);
 //	getchar();
 
-/*	LARGE_INTEGER temp;
-	QueryPerformanceFrequency((LARGE_INTEGER*) &temp);
-	double freq = ((double) temp.QuadPart) / 1000.0;
-
-    LARGE_INTEGER start_time;
-    QueryPerformanceCounter((LARGE_INTEGER*) &start_time);
-*/
 	clock_t start = clock();
 	
-	compressor.process(inputOptions, compressionOptions, outputOptions);
-/*
-	LARGE_INTEGER end_time;
-	QueryPerformanceCounter((LARGE_INTEGER*) &end_time);
-
-	float diff_time = (float) (((double) end_time.QuadPart - (double) start_time.QuadPart) / freq);
-	printf("\rtime taken: %.3f seconds\n", diff_time/1000);
-*/
+	if (!compressor.process(inputOptions, compressionOptions, outputOptions))
+	{
+		return EXIT_FAILURE;
+	}
 
 	clock_t end = clock();
 	printf("\rtime taken: %.3f seconds\n", float(end-start) / CLOCKS_PER_SEC);
 	
-	return 0;
+	return EXIT_SUCCESS;
 }
 

src/nvtt/tools/imgdiff.cpp

@@ -84,7 +84,7 @@ struct Error
 	{
 		mabse /= samples;
 		mse /= samples;
-		rmse = sqrt(mse);
+		rmse = sqrtf(mse);
 		psnr = (rmse == 0) ? 999.0f : 20.0f * log10(255.0f / rmse);
 	}
 
@@ -134,7 +134,7 @@ struct NormalError
 		{
 			ade /= samples;
 			mse /= samples * 3;
-			rmse = sqrt(mse);
+			rmse = sqrtf(mse);
 			psnr = (rmse == 0) ? 999.0f : 20.0f * log10(255.0f / rmse);
 		}
 	}

src/nvtt/tools/resize.cpp

@@ -113,7 +113,7 @@ int main(int argc, char *argv[])
 				((nv::KaiserFilter *)filter.ptr())->setParameters(4.0f, 1.0f);
 			}
 		}
-		else if (strcmp("-f", argv[i]) == 0)
+		else if (strcmp("-w", argv[i]) == 0)
 		{
 			if (i+1 == argc) break;
 			i++;