Tag 2.0.6 for release.

Shutdown CUDA properly when nvtt context is destroyed.
Fixes issue 83.

CMakeLists.txt

@@ -1,4 +1,4 @@
-CMAKE_MINIMUM_REQUIRED(VERSION 2.4.0)
+CMAKE_MINIMUM_REQUIRED(VERSION 2.6.0)
 PROJECT(NV)
 ENABLE_TESTING()
 
@@ -16,6 +16,13 @@ MESSAGE(STATUS "Setting optimal options")
 MESSAGE(STATUS "  Processor: ${NV_SYSTEM_PROCESSOR}")
 MESSAGE(STATUS "  Compiler Flags: ${CMAKE_CXX_FLAGS}")
 
+IF(NVTT_SHARED)
+	SET(NVCORE_SHARED TRUE)
+	SET(NVMATH_SHARED TRUE)
+	SET(NVIMAGE_SHARED TRUE)
+ENDIF(NVTT_SHARED)
+
+
 ADD_SUBDIRECTORY(src)
 
 IF(WIN32)

ChangeLog

@@ -1,3 +1,36 @@
+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.
+ * Fix cmake files when compiling NVTT as a shared library.
+ * When linking nvtt dynamically on unix, link all libraries dynamically.
+ * Select fastest CUDA device.
+
+NVIDIA Texture Tools version 2.0.4
+ * Fix error in RGB format output; reported by jonsoh. See issue 49.
+ * Added support RGB format dithering by jonsoh. Fixes issue 50 and 51.
+ * Prevent infinite loop in indexMirror when width equal 1. Fixes issue 65.
+ * Implement general scale filter, including upsampling.
+
+NVIDIA Texture Tools version 2.0.3
+ * More accurate DXT3 compressor. Fixes issue 38.
+ * Remove legacy compressors. Fix issue 34.
+ * Check for single color in all compressors. Fixes issue 43.
+ * Fix error in fast downsample filter, reported by Noel Llopis.
+
 NVIDIA Texture Tools version 2.0.2
  * Fix copy ctor error reported by Richard Sim.
  * Fix indexMirror error reported by Chris Lambert.

VERSION

@@ -1 +1 @@
-2.0.2
+2.0.6

cmake/FindCUDA.cmake

@@ -57,7 +57,7 @@ MARK_AS_ADVANCED (CUDA_FOUND CUDA_COMPILER CUDA_RUNTIME_LIBRARY)
 
 
 #SET(CUDA_OPTIONS "-ncfe")
-SET(CUDA_OPTIONS "")
+SET(CUDA_OPTIONS "--host-compilation=C")
 
 IF (CUDA_EMULATION)
 	SET (CUDA_OPTIONS "${CUDA_OPTIONS} -deviceemu")

configure

@@ -53,7 +53,7 @@ echo "-- Configuring nvidia-texture-tools "`cat VERSION`
 
 mkdir -p ./build
 cd ./build
-$CMAKE .. -DCMAKE_BUILD_TYPE=$build -DCMAKE_INSTALL_PREFIX=$prefix -G "Unix Makefiles" || exit 1
+$CMAKE .. -DNVTT_SHARED=1 -DCMAKE_BUILD_TYPE=$build -DCMAKE_INSTALL_PREFIX=$prefix -G "Unix Makefiles" || exit 1
 cd ..
 
 echo ""

project/vc8/nvcore/nvcore.vcproj

@@ -281,6 +281,10 @@
 				RelativePath="..\..\..\src\nvcore\Debug.cpp"
 				>
 			</File>
+			<File
+				RelativePath="..\..\..\src\nvcore\Library.cpp"
+				>
+			</File>
 			<File
 				RelativePath="..\..\..\src\nvcore\Memory.cpp"
 				>
@@ -315,6 +319,10 @@
 				RelativePath="..\..\..\src\nvcore\DefsVcWin32.h"
 				>
 			</File>
+			<File
+				RelativePath="..\..\..\src\nvcore\Library.h"
+				>
+			</File>
 			<File
 				RelativePath="..\..\..\src\nvcore\Memory.h"
 				>

project/vc8/nvimage/nvimage.vcproj

@@ -355,6 +355,10 @@
 				RelativePath="..\..\..\src\nvimage\nvimage.h"
 				>
 			</File>
+			<File
+				RelativePath="..\..\..\src\nvimage\PixelFormat.h"
+				>
+			</File>
 			<File
 				RelativePath="..\..\..\src\nvimage\PsdFile.h"
 				>

project/vc8/nvtt/nvtt.rc

@@ -53,8 +53,8 @@ END
 //
 
 VS_VERSION_INFO VERSIONINFO
- FILEVERSION 2,0,2,0
- PRODUCTVERSION 2,0,2,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, 2, 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, 2, 0"
+            VALUE "ProductVersion", "2, 0, 6, 0"
         END
     END
     BLOCK "VarFileInfo"

project/vc8/nvtt/nvtt.vcproj

@@ -711,7 +711,7 @@
 					>
 					<Tool
 						Name="VCCustomBuildTool"
-						CommandLine="&quot;$(CUDA_BIN_PATH)\nvcc.exe&quot; -m32 -ccbin &quot;$(VCInstallDir)bin&quot; -c -DNDEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MD -I&quot;$(CUDA_INC_PATH)&quot; -I./ -o $(IntDir)\$(InputName).obj ..\\..\\..\\src\\nvtt\\cuda\\CompressKernel.cu"
+						CommandLine="&quot;$(CUDA_BIN_PATH)\nvcc.exe&quot; -m32 -ccbin &quot;$(VCInstallDir)bin&quot; -c -DNDEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MD -I&quot;$(CUDA_INC_PATH)&quot; -I./ -o $(IntDir)\$(InputName).obj ..\\..\\..\\src\\nvtt\\cuda\\CompressKernel.cu&#x0D;&#x0A;"
 						AdditionalDependencies="CudaMath.h"
 						Outputs="$(IntDir)\$(InputName).obj"
 					/>
@@ -849,10 +849,6 @@
 				RelativePath="..\..\..\src\nvtt\cuda\CudaUtils.cpp"
 				>
 			</File>
-			<File
-				RelativePath="..\..\..\src\nvtt\FastCompressDXT.cpp"
-				>
-			</File>
 			<File
 				RelativePath="..\..\..\src\nvtt\InputOptions.cpp"
 				>
@@ -865,6 +861,10 @@
 				RelativePath="..\..\..\src\nvtt\nvtt_wrapper.cpp"
 				>
 			</File>
+			<File
+				RelativePath="..\..\..\src\nvtt\OptimalCompressDXT.cpp"
+				>
+			</File>
 			<File
 				RelativePath="..\..\..\src\nvtt\OutputOptions.cpp"
 				>
@@ -911,10 +911,6 @@
 				RelativePath="..\..\..\src\nvtt\cuda\CudaUtils.h"
 				>
 			</File>
-			<File
-				RelativePath="..\..\..\src\nvtt\FastCompressDXT.h"
-				>
-			</File>
 			<File
 				RelativePath="..\..\..\src\nvtt\InputOptions.h"
 				>
@@ -927,6 +923,10 @@
 				RelativePath="..\..\..\src\nvtt\nvtt_wrapper.h"
 				>
 			</File>
+			<File
+				RelativePath="..\..\..\src\nvtt\OptimalCompressDXT.h"
+				>
+			</File>
 			<File
 				RelativePath="..\..\..\src\nvtt\OutputOptions.h"
 				>

src/nvcore/CMakeLists.txt

@@ -19,13 +19,19 @@ SET(CORE_SRCS
 	TextWriter.h
 	TextWriter.cpp
 	Radix.h
-	Radix.cpp)
+	Radix.cpp
+	Library.h
+	Library.cpp)
 
 INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR})
 
 # targets
 ADD_DEFINITIONS(-DNVCORE_EXPORTS)
 
+IF(UNIX)
+	SET(LIBS ${LIBS} ${CMAKE_DL_LIBS})
+ENDIF(UNIX)
+
 IF(NVCORE_SHARED)
 	ADD_LIBRARY(nvcore SHARED ${CORE_SRCS})
 ELSE(NVCORE_SHARED)

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

@@ -28,7 +28,7 @@
 #endif
 
 #if NV_OS_LINUX && defined(HAVE_EXECINFO_H)
-#	include <execinfo.h>
+#	include <execinfo.h> // backtrace
 #	if NV_CC_GNUC // defined(HAVE_CXXABI_H)
 #		include <cxxabi.h>
 #	endif
@@ -39,6 +39,13 @@
 #	include <sys/types.h>
 #	include <sys/sysctl.h>	// sysctl
 #	include <ucontext.h>
+#	undef HAVE_EXECINFO_H
+#	if defined(HAVE_EXECINFO_H) // only after OSX 10.5
+#		include <execinfo.h> // backtrace
+#		if NV_CC_GNUC // defined(HAVE_CXXABI_H)
+#			include <cxxabi.h>
+#		endif
+#	endif
 #endif
 
 #include <stdexcept> // std::runtime_error
@@ -128,6 +135,14 @@ 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) {
 		char ** string_array = backtrace_symbols(trace, size);
 	
@@ -166,24 +181,36 @@ namespace
 
 	static void * callerAddress(void * secret)
 	{
-#	if NV_OS_DARWIN && NV_CPU_PPC
-		ucontext_t * ucp = (ucontext_t *)secret;
-        return (void *) ucp->uc_mcontext->ss.srr0;
-#	elif NV_OS_DARWIN && NV_CPU_X86
-		ucontext_t * ucp = (ucontext_t *)secret;
-        return (void *) ucp->uc_mcontext->ss.eip;
-#	elif NV_CPU_X86_64
-		// #define REG_RIP REG_INDEX(rip) // seems to be 16
-		ucontext_t * ucp = (ucontext_t *)secret;
-		return (void *)ucp->uc_mcontext.gregs[REG_RIP];
-#	elif NV_CPU_X86
-		ucontext_t * ucp = (ucontext_t *)secret;
-		return (void *)ucp->uc_mcontext.gregs[14/*REG_EIP*/];
-#	elif NV_CPU_PPC
-		ucontext_t * ucp = (ucontext_t *)secret;
-		return (void *) ucp->uc_mcontext.regs->nip;
+#	if NV_OS_DARWIN
+#		if defined(_STRUCT_MCONTEXT)
+#			if NV_CPU_PPC
+				ucontext_t * ucp = (ucontext_t *)secret;
+				return (void *) ucp->uc_mcontext->__ss.__srr0;
+#			elif NV_CPU_X86
+				ucontext_t * ucp = (ucontext_t *)secret;
+				return (void *) ucp->uc_mcontext->__ss.__eip;
+#			endif
+#		else
+#			if NV_CPU_PPC
+				ucontext_t * ucp = (ucontext_t *)secret;
+				return (void *) ucp->uc_mcontext->ss.srr0;
+#			elif NV_CPU_X86
+				ucontext_t * ucp = (ucontext_t *)secret;
+				return (void *) ucp->uc_mcontext->ss.eip;
+#			endif
+#		endif
 #	else
-		return NULL;
+#		if NV_CPU_X86_64
+			// #define REG_RIP REG_INDEX(rip) // seems to be 16
+			ucontext_t * ucp = (ucontext_t *)secret;
+			return (void *)ucp->uc_mcontext.gregs[REG_RIP];
+#		elif NV_CPU_X86
+			ucontext_t * ucp = (ucontext_t *)secret;
+			return (void *)ucp->uc_mcontext.gregs[14/*REG_EIP*/];
+#		elif NV_CPU_PPC
+			ucontext_t * ucp = (ucontext_t *)secret;
+			return (void *) ucp->uc_mcontext.regs->nip;
+#		endif
 #	endif
 		
 		// How to obtain the instruction pointers in different platforms, from mlton's source code.
@@ -228,17 +255,18 @@ namespace
 		}
 		
 #	if defined(HAVE_EXECINFO_H)
+		if (nvHasStackTrace()) // in case of weak linking
+		{
+			void * trace[64];
+			int size = backtrace(trace, 64);
 		
-		void * trace[64];
-		int size = backtrace(trace, 64);
-		
-		if (pnt != NULL) {
-			// Overwrite sigaction with caller's address.
-			trace[1] = pnt;
+			if (pnt != NULL) {
+				// Overwrite sigaction with caller's address.
+				trace[1] = pnt;
+			}
+			
+			nvPrintStackTrace(trace, size, 1);
 		}
-		
-		nvPrintStackTrace(trace, size, 1);
-		
 #	endif // defined(HAVE_EXECINFO_H)
 		
 		exit(0);
@@ -373,9 +401,12 @@ namespace
 #		endif
 
 #		if defined(HAVE_EXECINFO_H)
-			void * trace[64];
-			int size = backtrace(trace, 64);
-			nvPrintStackTrace(trace, size, 3);
+			if (nvHasStackTrace())
+			{
+				void * trace[64];
+				int size = backtrace(trace, 64);
+				nvPrintStackTrace(trace, size, 2);
+			}
 #		endif
 
 			// Exit cleanly.
@@ -422,9 +453,12 @@ void NV_CDECL nvDebug(const char *msg, ...)
 void debug::dumpInfo()
 {
 #if !NV_OS_WIN32 && defined(HAVE_SIGNAL_H) && defined(HAVE_EXECINFO_H)
-	void * trace[64];
-	int size = backtrace(trace, 64);
-	nvPrintStackTrace(trace, size, 1);
+	if (nvHasStackTrace())
+	{
+		void * trace[64];
+		int size = backtrace(trace, 64);
+		nvPrintStackTrace(trace, size, 1);
+	}
 #endif
 }
 

src/nvcore/DefsGnucDarwin.h

@@ -2,8 +2,7 @@
 #error "Do not include this file directly."
 #endif
 
-#include <stdlib.h>	// uint8_t, int8_t, ...
-
+#include <stdint.h> // uint8_t, int8_t, ...
 
 // Function linkage
 #define DLL_IMPORT

src/nvcore/DefsVcWin32.h

@@ -19,7 +19,9 @@
 
 // Set standard function names.
 #define snprintf _snprintf
-#define vsnprintf _vsnprintf
+#if _MSC_VER < 1500
+#	define vsnprintf _vsnprintf
+#endif
 #define vsscanf _vsscanf
 #define chdir _chdir
 #define getcwd _getcwd 
@@ -70,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/Library.cpp

@@ -0,0 +1,41 @@
+
+#include "Library.h"
+#include "Debug.h"
+
+#if NV_OS_WIN32
+#define WIN32_LEAN_AND_MEAN
+#define VC_EXTRALEAN
+#include <windows.h>
+#else
+#include <dlfcn.h>
+#endif
+
+
+
+void * nvLoadLibrary(const char * name)
+{
+#if NV_OS_WIN32
+	return (void *)LoadLibraryExA( name, NULL, 0 );
+#else
+	return dlopen(name, RTLD_LAZY);
+#endif
+}
+
+void nvUnloadLibrary(void * handle)
+{
+	nvDebugCheck(handle != NULL);
+#if NV_OS_WIN32
+	FreeLibrary((HMODULE)handle);
+#else
+	dlclose(handle);
+#endif
+}
+
+void * nvBindSymbol(void * handle, const char * symbol)
+{
+#if NV_OS_WIN32
+	return (void *)GetProcAddress((HMODULE)handle, symbol);
+#else
+	return (void *)dlsym(handle, symbol);
+#endif	
+}

src/nvcore/Library.h

@@ -0,0 +1,50 @@
+// This code is in the public domain -- castano@gmail.com
+
+#ifndef NV_CORE_LIBRARY_H
+#define NV_CORE_LIBRARY_H
+
+#include <nvcore/nvcore.h>
+
+#if NV_OS_WIN32
+#define LIBRARY_NAME(name)	#name ".dll"
+#elif NV_OS_DARWIN
+#define NV_LIBRARY_NAME(name)	"lib" #name ".dylib"
+#else
+#define NV_LIBRARY_NAME(name)	"lib" #name ".so"
+#endif
+
+NVCORE_API void * nvLoadLibrary(const char * name);
+NVCORE_API void nvUnloadLibrary(void * lib);
+NVCORE_API void * nvBindSymbol(void * lib, const char * symbol);
+
+class NVCORE_CLASS Library
+{
+public:
+	Library(const char * name)
+	{
+		handle = nvLoadLibrary(name);
+	}
+	~Library()
+	{
+		if (isValid())
+		{
+			nvUnloadLibrary(handle);
+		}
+	}
+	
+	bool isValid() const
+	{
+		return handle != NULL;
+	}
+	
+	void * bindSymbol(const char * symbol)
+	{
+		return nvBindSymbol(handle, symbol);
+	}
+	
+private:
+	void * handle;
+};
+
+
+#endif // NV_CORE_LIBRARY_H

src/nvcore/Prefetch.h

@@ -24,7 +24,7 @@ __forceinline void nvPrefetch(const void * mem)
 #else // NV_CC_MSVC
 
 // do nothing in other case.
-#define piPrefetch(ptr)
+#define nvPrefetch(ptr)
 
 #endif // NV_CC_MSVC
 

src/nvcore/Ptr.h

@@ -43,8 +43,11 @@ public:
 
 	/** Delete owned pointer and assign new one. */
 	void operator=( T * p ) {
-		delete m_ptr;
-		m_ptr = p;
+		if (p != m_ptr)
+		{
+			delete m_ptr;
+			m_ptr = p;
+		}
 	}
 
 	/** Member access. */
@@ -249,14 +252,14 @@ public:
 		/** -> operator. */
 		BaseClass * operator -> () const
 		{
-			piCheck( m_ptr != NULL );
+			nvCheck( m_ptr != NULL );
 			return m_ptr;
 		}
 
 		/** * operator. */
 		BaseClass & operator*() const
 		{
-			piCheck( m_ptr != NULL );
+			nvCheck( m_ptr != NULL );
 			return *m_ptr;
 		}
 

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
 {

src/nvcore/StrLib.h

@@ -14,7 +14,7 @@ namespace nv
 
 	uint strHash(const char * str, uint h) NV_PURE;
 
-	/// String hash vased on Bernstein's hash.
+	/// String hash based on Bernstein's hash.
 	inline uint strHash(const char * data, uint h = 5381)
 	{
 		uint i;
@@ -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();
 	
@@ -120,18 +117,16 @@ namespace nv
 		char * m_str;
 		
 	};
-	
 
-	/// 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;
@@ -140,7 +135,7 @@ namespace nv
 		
 		void stripFileName();
 		void stripExtension();
-		
+
 		// statics
 		NVCORE_API static char separator();
 		NVCORE_API static const char * fileName(const char *);
@@ -213,9 +208,12 @@ namespace nv
 		/// Implement value semantics.
 		String & operator=( const String & str )
 		{
-			release();
-			data = str.data;
-			addRef();
+			if (str.data != data)
+			{
+				release();
+				data = str.data;
+				addRef();
+			}
 			return *this;
 		}
 

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) | (2);	// major.minor.revision
+	this->reserved[10] = (2 << 16) | (0 << 8) | (6);	// 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

@@ -26,18 +26,17 @@
  * http://www.xmission.com/~legalize/zoom.html
  * 
  * Reconstruction Filters in Computer Graphics
- * http://www.mentallandscape.com/Papers_siggraph88.pdf 
+ * http://www.mentallandscape.com/Papers_siggraph88.pdf
  *
  * More references:
- * http://www.worldserver.com/turk/computergraphics/ResamplingFilters.pdf 
+ * http://www.worldserver.com/turk/computergraphics/ResamplingFilters.pdf
  * 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;
 
@@ -244,7 +243,7 @@ SincFilter::SincFilter(float w) : Filter(w) {}
 
 float SincFilter::evaluate(float x) const
 {
-	return 0.0f;
+	return sincf(PI * x);
 }
 
 
@@ -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();
 			}
 		}
 	}
@@ -541,12 +540,17 @@ void Kernel2::initBlendedSobel(const Vector4 & scale)
 
 PolyphaseKernel::PolyphaseKernel(const Filter & f, uint srcLength, uint dstLength, int samples/*= 32*/)
 {
-	nvCheck(srcLength >= dstLength);	// @@ Upsampling not implemented!
 	nvDebugCheck(samples > 0);
-	
-	const float scale = float(dstLength) / float(srcLength);
+
+	float scale = float(dstLength) / float(srcLength);
 	const float iscale = 1.0f / scale;
 
+	if (scale > 1) {
+		// Upsampling.
+		samples = 1;
+		scale = 1;
+	}
+
 	m_length = dstLength;
 	m_width = f.width() * iscale;
 	m_windowSize = (int)ceilf(m_width * 2) + 1;

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;
 }
@@ -376,7 +378,7 @@ FloatImage * FloatImage::fastDownSample() const
 	{
 		const uint n = w * h;
 		
-		if (n & 1)
+		if ((m_width * m_height) & 1)
 		{
 			const float scale = 1.0f / (2 * n + 1);
 			
@@ -540,73 +542,27 @@ FloatImage * FloatImage::fastDownSample() const
 	return dst_image.release();
 }
 
-/*
-/// Downsample applying a 1D kernel separately in each dimension.
-FloatImage * FloatImage::downSample(const Kernel1 & kernel, WrapMode wm) const
-{
-	const uint w = max(1, m_width / 2);
-	const uint h = max(1, m_height / 2);
-	
-	return downSample(kernel, w, h, wm);
-}
-
-
-/// Downsample applying a 1D kernel separately in each dimension.
-FloatImage * FloatImage::downSample(const Kernel1 & kernel, uint w, uint h, WrapMode wm) const
-{
-	nvCheck(!(kernel.windowSize() & 1));	// Make sure that kernel m_width is even.
-
-	AutoPtr<FloatImage> tmp_image( new FloatImage() );
-	tmp_image->allocate(m_componentNum, w, m_height);
-	
-	AutoPtr<FloatImage> dst_image( new FloatImage() );	
-	dst_image->allocate(m_componentNum, w, h);
-	
-	const float xscale = float(m_width) / float(w);
-	const float yscale = float(m_height) / float(h);
-	
-	for(uint c = 0; c < m_componentNum; c++) {
-		float * tmp_channel = tmp_image->channel(c);
-		
-		for(uint y = 0; y < m_height; y++) {
-			for(uint x = 0; x < w; x++) {
-				
-				float sum = this->applyKernelHorizontal(&kernel, uint(x*xscale), y, c, wm);
-				
-				const uint tmp_index = tmp_image->index(x, y);
-				tmp_channel[tmp_index] = sum;
-			}
-		}
-		
-		float * dst_channel = dst_image->channel(c);
-		
-		for(uint y = 0; y < h; y++) {
-			for(uint x = 0; x < w; x++) {
-				
-				float sum = tmp_image->applyKernelVertical(&kernel, uint(x*xscale), uint(y*yscale), c, wm);
-				
-				const uint dst_index = dst_image->index(x, y);
-				dst_channel[dst_index] = sum;
-			}
-		}
-	}
-	
-	return dst_image.release();
-}
-*/
-
 /// Downsample applying a 1D kernel separately in each dimension.
 FloatImage * FloatImage::downSample(const Filter & filter, WrapMode wm) const
 {
 	const uint w = max(1, m_width / 2);
 	const uint h = max(1, m_height / 2);
 
-	return downSample(filter, w, h, wm);
+	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::downSample(const Filter & filter, uint w, uint h, WrapMode wm) const
+FloatImage * FloatImage::resize(const Filter & filter, uint w, uint h, WrapMode wm) const
 {
 	// @@ Use monophase filters when frac(m_width / w) == 0
 
@@ -675,10 +631,56 @@ FloatImage * FloatImage::downSample(const Filter & filter, uint w, uint h, WrapM
 	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);
 	
@@ -707,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);
 	
@@ -729,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);
 	
@@ -752,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);
@@ -784,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);
@@ -815,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;
@@ -60,20 +68,22 @@ public:
 	NVIMAGE_API void toGamma(uint base_component, uint num, float gamma = 2.2f);
 	NVIMAGE_API void exponentiate(uint base_component, uint num, float power);
 	
-	
+
 	NVIMAGE_API FloatImage * fastDownSample() const;
 	NVIMAGE_API FloatImage * downSample(const Filter & filter, WrapMode wm) const;
-	NVIMAGE_API FloatImage * downSample(const Filter & filter, uint w, uint h, 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;
@@ -226,11 +239,15 @@ inline uint FloatImage::indexRepeat(int x, int y) const
 
 inline uint FloatImage::indexMirror(int x, int y) const
 {
+	if (m_width == 1) x = 0;
+
 	x = abs(x);
 	while (x >= m_width) {
 		x = abs(m_width + m_width - x - 2);
 	}
 
+	if (m_height == 1) y = 0;
+
 	y = abs(y);
 	while (y >= m_height) {
 		y = abs(m_height + m_height - y - 2);

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,10 +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 <nvcore/Containers.h> // swap
+
 
 using namespace nv;
 
@@ -47,94 +51,20 @@ void nv::Quantize::BinaryAlpha( Image * image, int alpha_threshold /*= 127*/ )
 // Simple quantization.
 void nv::Quantize::RGB16( Image * image )
 {
-	nvCheck(image != NULL);
-	
-	const uint w = image->width();
-	const uint h = image->height();
-	
-	for(uint y = 0; y < h; y++) {
-		for(uint x = 0; x < w; x++) {
-			
-			Color32 pixel32 = image->pixel(x, y);
-			
-			// Convert to 16 bit and back to 32 using regular bit expansion.
-			Color32 pixel16 = toColor32( toColor16(pixel32) );
-			
-			// Store color.
-			image->pixel(x, y) = pixel16;
-		}
-	}
+	Truncate(image, 5, 6, 5, 8);
 }
 
 // Alpha quantization.
 void nv::Quantize::Alpha4( Image * image )
 {
-	nvCheck(image != NULL);
-	
-	const uint w = image->width();
-	const uint h = image->height();
-	
-	for(uint y = 0; y < h; y++) {
-		for(uint x = 0; x < w; x++) {
-			
-			Color32 pixel = image->pixel(x, y);
-			
-			// Convert to 4 bit using regular bit expansion.
-			pixel.a = (pixel.a & 0xF0) | ((pixel.a & 0xF0) >> 4);
-			
-			// Store color.
-			image->pixel(x, y) = pixel;
-		}
-	}
+	Truncate(image, 8, 8, 8, 4);
 }
 
 
 // Error diffusion. Floyd Steinberg.
 void nv::Quantize::FloydSteinberg_RGB16( Image * image )
 {
-	nvCheck(image != NULL);
-	
-	const uint w = image->width();
-	const uint h = image->height();
-	
-	// @@ Use fixed point?
-	Vector3 * row0 = new Vector3[w+2];
-	Vector3 * row1 = new Vector3[w+2];
-	memset(row0, 0, sizeof(Vector3)*(w+2));
-	memset(row1, 0, sizeof(Vector3)*(w+2));
-	
-	for(uint y = 0; y < h; y++) {
-		for(uint x = 0; x < w; x++) {
-			
-			Color32 pixel32 = image->pixel(x, y);
-			
-			// Add error.	// @@ We shouldn't clamp here!
-			pixel32.r = clamp(int(pixel32.r) + int(row0[1+x].x()), 0, 255);
-			pixel32.g = clamp(int(pixel32.g) + int(row0[1+x].y()), 0, 255);
-			pixel32.b = clamp(int(pixel32.b) + int(row0[1+x].z()), 0, 255);
-			
-			// Convert to 16 bit. @@ Use regular clamp?
-			Color32 pixel16 = toColor32( toColor16(pixel32) );
-			
-			// Store color.
-			image->pixel(x, y) = pixel16;
-			
-			// Compute new error.
-			Vector3 diff(float(pixel32.r - pixel16.r), float(pixel32.g - pixel16.g), float(pixel32.b - pixel16.b));
-			
-			// Propagate new error.
-			row0[1+x+1] += 7.0f / 16.0f * diff;
-			row1[1+x-1] += 3.0f / 16.0f * diff;
-			row1[1+x+0] += 5.0f / 16.0f * diff;
-			row1[1+x+1] += 1.0f / 16.0f * diff;
-		}
-		
-		swap(row0, row1);
-		memset(row1, 0, sizeof(Vector3)*(w+2));
-	}
-	
-	delete [] row0;
-	delete [] row1;
+	FloydSteinberg(image, 5, 6, 5, 8);
 }
 
 
@@ -188,34 +118,90 @@ void nv::Quantize::FloydSteinberg_BinaryAlpha( Image * image, int alpha_threshol
 
 // Error diffusion. Floyd Steinberg.
 void nv::Quantize::FloydSteinberg_Alpha4( Image * image )
+{
+	FloydSteinberg(image, 8, 8, 8, 4);
+}
+
+
+void nv::Quantize::Truncate(Image * image, uint rsize, uint gsize, uint bsize, uint asize)
 {
 	nvCheck(image != NULL);
 	
 	const uint w = image->width();
 	const uint h = image->height();
 	
-	// @@ Use fixed point?
-	float * row0 = new float[(w+2)];
-	float * row1 = new float[(w+2)];
-	memset(row0, 0, sizeof(float)*(w+2));
-	memset(row1, 0, sizeof(float)*(w+2));
-	
 	for(uint y = 0; y < h; y++) {
 		for(uint x = 0; x < w; x++) {
 			
 			Color32 pixel = image->pixel(x, y);
-			
-			// Add error.
-			int alpha = int(pixel.a) + int(row0[1+x]);
-			
-			// Convert to 4 bit using regular bit expansion.
-			pixel.a = (pixel.a & 0xF0) | ((pixel.a & 0xF0) >> 4);
-			
+
+			// Convert to our desired size, and reconstruct.
+			pixel.r = PixelFormat::convert(pixel.r, 8, rsize);
+			pixel.r = PixelFormat::convert(pixel.r, rsize, 8);
+
+			pixel.g = PixelFormat::convert(pixel.g, 8, gsize);
+			pixel.g = PixelFormat::convert(pixel.g, gsize, 8);
+
+			pixel.b = PixelFormat::convert(pixel.b, 8, bsize);
+			pixel.b = PixelFormat::convert(pixel.b, bsize, 8);
+
+			pixel.a = PixelFormat::convert(pixel.a, 8, asize);
+			pixel.a = PixelFormat::convert(pixel.a, asize, 8);
+
 			// Store color.
 			image->pixel(x, y) = pixel;
+		}
+	}
+}
+
+
+// Error diffusion. Floyd Steinberg.
+void nv::Quantize::FloydSteinberg(Image * image, uint rsize, uint gsize, uint bsize, uint asize)
+{
+	nvCheck(image != NULL);
+	
+	const uint w = image->width();
+	const uint h = image->height();
+	
+	Vector4 * row0 = new Vector4[w+2];
+	Vector4 * row1 = new Vector4[w+2];
+	memset(row0, 0, sizeof(Vector4)*(w+2));
+	memset(row1, 0, sizeof(Vector4)*(w+2));
+	
+	for (uint y = 0; y < h; y++) {
+		for (uint x = 0; x < w; x++) {
+			
+			Color32 pixel = image->pixel(x, y);
+
+			// Add error.
+			pixel.r = clamp(int(pixel.r) + int(row0[1+x].x()), 0, 255);
+			pixel.g = clamp(int(pixel.g) + int(row0[1+x].y()), 0, 255);
+			pixel.b = clamp(int(pixel.b) + int(row0[1+x].z()), 0, 255);
+			pixel.a = clamp(int(pixel.a) + int(row0[1+x].w()), 0, 255);
+			
+			int r = pixel.r;
+			int g = pixel.g;
+			int b = pixel.b;
+			int a = pixel.a;
+
+			// Convert to our desired size, and reconstruct.
+			r = PixelFormat::convert(r, 8, rsize);
+			r = PixelFormat::convert(r, rsize, 8);
+
+			g = PixelFormat::convert(g, 8, gsize);
+			g = PixelFormat::convert(g, gsize, 8);
+
+			b = PixelFormat::convert(b, 8, bsize);
+			b = PixelFormat::convert(b, bsize, 8);
+
+			a = PixelFormat::convert(a, 8, asize);
+			a = PixelFormat::convert(a, asize, 8);
+
+			// Store color.
+			image->pixel(x, y) = Color32(r, g, b, a);
 			
 			// Compute new error.
-			float diff = float(alpha - pixel.a);
+			Vector4 diff(float(int(pixel.r) - r), float(int(pixel.g) - g), float(int(pixel.b) - b), float(int(pixel.a) - a));
 			
 			// Propagate new error.
 			row0[1+x+1] += 7.0f / 16.0f * diff;
@@ -225,10 +211,9 @@ void nv::Quantize::FloydSteinberg_Alpha4( Image * image )
 		}
 		
 		swap(row0, row1);
-		memset(row1, 0, sizeof(float)*(w+2));
+		memset(row1, 0, sizeof(Vector4)*(w+2));
 	}
 	
 	delete [] row0;
 	delete [] row1;
 }
-

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;
@@ -17,6 +20,9 @@ namespace nv
 		void FloydSteinberg_BinaryAlpha(Image * img, int alpha_threshold = 127);
 		void FloydSteinberg_Alpha4(Image * img);
 
+		void Truncate(Image * image, uint rsize, uint gsize, uint bsize, uint asize);
+		void FloydSteinberg(Image * image, uint rsize, uint gsize, uint bsize, uint asize);
+
 		// @@ Add palette quantization algorithms!
 	}
 }

src/nvmath/Box.h

@@ -108,7 +108,7 @@ public:
 	float area() const
 	{
 		const Vector3 d = extents();
-		return 4.0f * (d.x()*d.y() + d.x()*d.z() + d.y()*d.z());
+		return 8.0f * (d.x()*d.y() + d.x()*d.z() + d.y()*d.z());
 	}	
 
 	/// Get the volume of the box.
@@ -118,6 +118,14 @@ public:
 		return 8.0f * (d.x() * d.y() * d.z());
 	}
 	
+	/// Return true if the box contains the given point.
+	bool contains(Vector3::Arg p) const
+	{
+		return 
+			m_mins.x() < p.x() && m_mins.y() < p.y() && m_mins.z() < p.z() &&
+			m_maxs.x() > p.x() && m_maxs.y() > p.y() && m_maxs.z() > p.z();
+	}
+
 private:
 
 	Vector3 m_mins;
@@ -125,15 +133,6 @@ private:
 };
 
 
-/*
-/// Point inside box test.
-inline bool pointInsideBox(const Box & b, Vector3::Arg p) const
-{
-	return (m_mins.x() < p.x() && m_mins.y() < p.y() && m_mins.z() < p.z() &&
-		m_maxs.x() > p.x() && m_maxs.y() > p.y() && m_maxs.z() > p.z());
-}
-*/
-
 
 } // nv namespace
 

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/CMakeLists.txt

@@ -13,10 +13,10 @@ SET(NVTT_SRCS
 	CompressDXT.cpp
 	CompressRGB.h
 	CompressRGB.cpp
-	FastCompressDXT.h
-	FastCompressDXT.cpp
 	QuickCompressDXT.h
 	QuickCompressDXT.cpp
+	OptimalCompressDXT.h
+	OptimalCompressDXT.cpp
 	SingleColorLookup.h
 	CompressionOptions.h
 	CompressionOptions.cpp
@@ -43,8 +43,9 @@ INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR})
 
 ADD_DEFINITIONS(-DNVTT_EXPORTS)
 
-IF(NVTT_SHARED)	
-	ADD_LIBRARY(nvtt SHARED ${DXT_SRCS})
+IF(NVTT_SHARED)
+	ADD_DEFINITIONS(-DNVTT_SHARED=1)
+	ADD_LIBRARY(nvtt SHARED ${NVTT_SRCS})
 ELSE(NVTT_SHARED)
 	ADD_LIBRARY(nvtt ${NVTT_SRCS})
 ENDIF(NVTT_SHARED)
@@ -84,7 +85,7 @@ TARGET_LINK_LIBRARIES(nvzoom nvcore nvmath nvimage)
 
 INSTALL(TARGETS nvcompress nvdecompress nvddsinfo nvimgdiff nvassemble nvzoom DESTINATION bin)
 
-# UI tools
+# UI tools
 IF(QT4_FOUND AND NOT MSVC)
 	SET(QT_USE_QTOPENGL TRUE)
 	INCLUDE_DIRECTORIES(${QT_INCLUDE_DIR} ${CMAKE_CURRENT_BINARY_DIR})
@@ -106,7 +107,7 @@ IF(QT4_FOUND AND NOT MSVC)
 
 	ADD_EXECUTABLE(nvcompressui MACOSX_BUNDLE ${SRCS} ${UICS} ${MOCS})
 	TARGET_LINK_LIBRARIES(nvcompressui ${LIBS})
-
+
 ENDIF(QT4_FOUND AND NOT MSVC)
 
 

src/nvtt/CompressDXT.cpp

@@ -29,8 +29,8 @@
 
 #include "nvtt.h"
 #include "CompressDXT.h"
-#include "FastCompressDXT.h"
 #include "QuickCompressDXT.h"
+#include "OptimalCompressDXT.h"
 #include "CompressionOptions.h"
 #include "OutputOptions.h"
 
@@ -57,26 +57,33 @@ using namespace nv;
 using namespace nvtt;
 
 
-void nv::fastCompressDXT1(const Image * image, const OutputOptions::Private & outputOptions)
+nv::FastCompressor::FastCompressor() : m_image(NULL), m_alphaMode(AlphaMode_None)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+}
+
+nv::FastCompressor::~FastCompressor()
+{
+}
+
+void nv::FastCompressor::setImage(const Image * image, nvtt::AlphaMode alphaMode)
+{
+	m_image = image;
+	m_alphaMode = alphaMode;
+}
+
+void nv::FastCompressor::compressDXT1(const OutputOptions::Private & outputOptions)
+{
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 	
 	ColorBlock rgba;
 	BlockDXT1 block;
 
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
-			rgba.init(image, x, y);
+			rgba.init(m_image, x, y);
 			
-			if (rgba.isSingleColor())
-			{
-				QuickCompress::compressDXT1(rgba.color(0), &block);
-			}
-			else
-			{
-				QuickCompress::compressDXT1(rgba, &block);
-			}
+			QuickCompress::compressDXT1(rgba, &block);
 			
 			if (outputOptions.outputHandler != NULL) {
 				outputOptions.outputHandler->writeData(&block, sizeof(block));
@@ -86,27 +93,19 @@ void nv::fastCompressDXT1(const Image * image, const OutputOptions::Private & ou
 }
 
 
-void nv::fastCompressDXT1a(const Image * image, const OutputOptions::Private & outputOptions)
+void nv::FastCompressor::compressDXT1a(const OutputOptions::Private & outputOptions)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 	
 	ColorBlock rgba;
 	BlockDXT1 block;
 
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
-			rgba.init(image, x, y);
+			rgba.init(m_image, x, y);
 			
-			// @@ We could do better here: check for single RGB, but varying alpha.
-			if (rgba.isSingleColor())
-			{
-				QuickCompress::compressDXT1a(rgba.color(0), &block);
-			}
-			else
-			{
-				QuickCompress::compressDXT1a(rgba, &block);
-			}
+			QuickCompress::compressDXT1a(rgba, &block);
 			
 			if (outputOptions.outputHandler != NULL) {
 				outputOptions.outputHandler->writeData(&block, sizeof(block));
@@ -116,17 +115,18 @@ void nv::fastCompressDXT1a(const Image * image, const OutputOptions::Private & o
 }
 
 
-void nv::fastCompressDXT3(const Image * image, const nvtt::OutputOptions::Private & outputOptions)
+void nv::FastCompressor::compressDXT3(const nvtt::OutputOptions::Private & outputOptions)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 	
 	ColorBlock rgba;
 	BlockDXT3 block;
 
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
-			rgba.init(image, x, y);
+			rgba.init(m_image, x, y);
+
 			QuickCompress::compressDXT3(rgba, &block);
 			
 			if (outputOptions.outputHandler != NULL) {
@@ -137,19 +137,19 @@ void nv::fastCompressDXT3(const Image * image, const nvtt::OutputOptions::Privat
 }
 
 
-void nv::fastCompressDXT5(const Image * image, const nvtt::OutputOptions::Private & outputOptions)
+void nv::FastCompressor::compressDXT5(const nvtt::OutputOptions::Private & outputOptions)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 	
 	ColorBlock rgba;
 	BlockDXT5 block;
 
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
-			rgba.init(image, x, y);
-			//QuickCompress::compressDXT5(rgba, &block);	// @@ Use fast version!!
-			nv::compressBlock_BoundsRange(rgba, &block);
+			rgba.init(m_image, x, y);
+			
+			QuickCompress::compressDXT5(rgba, &block, 0);
 			
 			if (outputOptions.outputHandler != NULL) {
 				outputOptions.outputHandler->writeData(&block, sizeof(block));
@@ -159,23 +159,21 @@ void nv::fastCompressDXT5(const Image * image, const nvtt::OutputOptions::Privat
 }
 
 
-void nv::fastCompressDXT5n(const Image * image, const nvtt::OutputOptions::Private & outputOptions)
+void nv::FastCompressor::compressDXT5n(const nvtt::OutputOptions::Private & outputOptions)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 	
 	ColorBlock rgba;
 	BlockDXT5 block;
 
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
-			rgba.init(image, x, y);
+			rgba.init(m_image, x, y);
 			
-			// copy X coordinate to alpha channel and Y coordinate to green channel.
 			rgba.swizzleDXT5n();
 
-			//QuickCompress::compressDXT5(rgba, &block);	// @@ Use fast version!!
-			nv::compressBlock_BoundsRange(rgba, &block);
+			QuickCompress::compressDXT5(rgba, &block, 0);
 			
 			if (outputOptions.outputHandler != NULL) {
 				outputOptions.outputHandler->writeData(&block, sizeof(block));
@@ -185,42 +183,28 @@ void nv::fastCompressDXT5n(const Image * image, const nvtt::OutputOptions::Priva
 }
 
 
-void nv::fastCompressBC4(const Image * image, const nvtt::OutputOptions::Private & outputOptions)
+nv::SlowCompressor::SlowCompressor() : m_image(NULL), m_alphaMode(AlphaMode_None)
 {
-	// @@ TODO
-	// compress red channel (X)
 }
 
-
-void nv::fastCompressBC5(const Image * image, const nvtt::OutputOptions::Private & outputOptions)
+nv::SlowCompressor::~SlowCompressor()
 {
-	// @@ TODO
-	// compress red, green channels (X,Y)
 }
 
-
-void nv::doPrecomputation()
+void nv::SlowCompressor::setImage(const Image * image, nvtt::AlphaMode alphaMode)
 {
-	static bool done = false;	// @@ Stop using statics for reentrancy. Although the worst that could happen is that this stuff is precomputed multiple times.
-	
-	if (!done)
-	{
-		done = true;
-		squish::FastClusterFit::DoPrecomputation();
-	}
+	m_image = image;
+	m_alphaMode = alphaMode;
 }
 
-
-void nv::compressDXT1(const Image * image, const OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions)
+void nv::SlowCompressor::compressDXT1(const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 	
 	ColorBlock rgba;
 	BlockDXT1 block;
 
-	doPrecomputation();
-
 	//squish::WeightedClusterFit fit;
 	//squish::ClusterFit fit;
 	squish::FastClusterFit fit;
@@ -229,11 +213,11 @@ void nv::compressDXT1(const Image * image, const OutputOptions::Private & output
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
 			
-			rgba.init(image, x, y);
+			rgba.init(m_image, x, y);
 			
 			if (rgba.isSingleColor())
 			{
-				QuickCompress::compressDXT1(rgba.color(0), &block);
+				OptimalCompress::compressDXT1(rgba.color(0), &block);
 			}
 			else
 			{
@@ -250,10 +234,10 @@ void nv::compressDXT1(const Image * image, const OutputOptions::Private & output
 }
 
 
-void nv::compressDXT1a(const Image * image, const OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions)
+void nv::SlowCompressor::compressDXT1a(const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 	
 	ColorBlock rgba;
 	BlockDXT1 block;
@@ -264,11 +248,20 @@ void nv::compressDXT1a(const Image * image, const OutputOptions::Private & outpu
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
 			
-			rgba.init(image, x, y);
+			rgba.init(m_image, x, y);
 			
-			if (rgba.isSingleColor())
+			bool anyAlpha = false;
+			bool allAlpha = true;
+			
+			for (uint i = 0; i < 16; i++)
 			{
-				QuickCompress::compressDXT1a(rgba.color(0), &block);
+				if (rgba.color(i).a < 128) anyAlpha = true;
+				else allAlpha = false;
+			}
+			
+			if ((!anyAlpha && rgba.isSingleColor() || allAlpha))
+			{
+				OptimalCompress::compressDXT1a(rgba.color(0), &block);
 			}
 			else
 			{
@@ -285,29 +278,37 @@ void nv::compressDXT1a(const Image * image, const OutputOptions::Private & outpu
 }
 
 
-void nv::compressDXT3(const Image * image, const OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions)
+void nv::SlowCompressor::compressDXT3(const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 	
 	ColorBlock rgba;
 	BlockDXT3 block;
 	
 	squish::WeightedClusterFit fit;
+	//squish::FastClusterFit fit;
 	fit.SetMetric(compressionOptions.colorWeight.x(), compressionOptions.colorWeight.y(), compressionOptions.colorWeight.z());
 
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
 			
-			rgba.init(image, x, y);
+			rgba.init(m_image, x, y);
 			
 			// Compress explicit alpha.
-			QuickCompress::compressDXT3A(rgba, &block.alpha);
-			
+			OptimalCompress::compressDXT3A(rgba, &block.alpha);
+
 			// Compress color.
-			squish::ColourSet colours((uint8 *)rgba.colors(), squish::kWeightColourByAlpha);
-			fit.SetColourSet(&colours, 0);
-			fit.Compress(&block.color);
+			if (rgba.isSingleColor())
+			{
+				OptimalCompress::compressDXT1(rgba.color(0), &block.color);
+			}
+			else
+			{
+				squish::ColourSet colours((uint8 *)rgba.colors(), squish::kWeightColourByAlpha);
+				fit.SetColourSet(&colours, 0);
+				fit.Compress(&block.color);
+			}
 			
 			if (outputOptions.outputHandler != NULL) {
 				outputOptions.outputHandler->writeData(&block, sizeof(block));
@@ -316,10 +317,10 @@ void nv::compressDXT3(const Image * image, const OutputOptions::Private & output
 	}
 }
 
-void nv::compressDXT5(const Image * image, const OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions)
+void nv::SlowCompressor::compressDXT5(const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 	
 	ColorBlock rgba;
 	BlockDXT5 block;
@@ -330,22 +331,29 @@ void nv::compressDXT5(const Image * image, const OutputOptions::Private & output
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
 			
-			rgba.init(image, x, y);
+			rgba.init(m_image, x, y);
 
 			// Compress alpha.
 			if (compressionOptions.quality == Quality_Highest)
 			{
-				compressBlock_BruteForce(rgba, &block.alpha);
+				OptimalCompress::compressDXT5A(rgba, &block.alpha);
 			}
 			else
 			{
 				QuickCompress::compressDXT5A(rgba, &block.alpha);
 			}
-
+		
 			// Compress color.
-			squish::ColourSet colours((uint8 *)rgba.colors(), squish::kWeightColourByAlpha);
-			fit.SetColourSet(&colours, 0);
-			fit.Compress(&block.color);
+			if (rgba.isSingleColor())
+			{
+				OptimalCompress::compressDXT1(rgba.color(0), &block.color);
+			}
+			else
+			{
+				squish::ColourSet colours((uint8 *)rgba.colors(), squish::kWeightColourByAlpha);
+				fit.SetColourSet(&colours, 0);
+				fit.Compress(&block.color);
+			}
 			
 			if (outputOptions.outputHandler != NULL) {
 				outputOptions.outputHandler->writeData(&block, sizeof(block));
@@ -355,28 +363,25 @@ void nv::compressDXT5(const Image * image, const OutputOptions::Private & output
 }
 
 
-void nv::compressDXT5n(const Image * image, const OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions)
+void nv::SlowCompressor::compressDXT5n(const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 	
 	ColorBlock rgba;
 	BlockDXT5 block;
 	
-	doPrecomputation();
-
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
 			
-			rgba.init(image, x, y);
+			rgba.init(m_image, x, y);
 			
-			// copy X coordinate to green channel and Y coordinate to alpha channel.
 			rgba.swizzleDXT5n();			
 			
 			// Compress X.
 			if (compressionOptions.quality == Quality_Highest)
 			{
-				compressBlock_BruteForce(rgba, &block.alpha);
+				OptimalCompress::compressDXT5A(rgba, &block.alpha);
 			}
 			else
 			{
@@ -384,7 +389,7 @@ void nv::compressDXT5n(const Image * image, const OutputOptions::Private & outpu
 			}
 			
 			// Compress Y.
-			QuickCompress::compressDXT1G(rgba, &block.color);
+			OptimalCompress::compressDXT1G(rgba, &block.color);
 			
 			if (outputOptions.outputHandler != NULL) {
 				outputOptions.outputHandler->writeData(&block, sizeof(block));
@@ -394,10 +399,10 @@ void nv::compressDXT5n(const Image * image, const OutputOptions::Private & outpu
 }
 
 
-void nv::compressBC4(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions)
+void nv::SlowCompressor::compressBC4(const CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 	
 	ColorBlock rgba;
 	AlphaBlockDXT5 block;
@@ -405,11 +410,11 @@ void nv::compressBC4(const Image * image, const nvtt::OutputOptions::Private & o
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
 			
-			rgba.init(image, x, y);
+			rgba.init(m_image, x, y);
 
 			if (compressionOptions.quality == Quality_Highest)
 			{
-				compressBlock_BruteForce(rgba, &block);
+				OptimalCompress::compressDXT5A(rgba, &block);
 			}
 			else
 			{
@@ -424,10 +429,10 @@ void nv::compressBC4(const Image * image, const nvtt::OutputOptions::Private & o
 }
 
 
-void nv::compressBC5(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions)
+void nv::SlowCompressor::compressBC5(const CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions)
 {
-	const uint w = image->width();
-	const uint h = image->height();
+	const uint w = m_image->width();
+	const uint h = m_image->height();
 
 	ColorBlock xcolor;
 	ColorBlock ycolor;
@@ -437,16 +442,16 @@ void nv::compressBC5(const Image * image, const nvtt::OutputOptions::Private & o
 	for (uint y = 0; y < h; y += 4) {
 		for (uint x = 0; x < w; x += 4) {
 			
-			xcolor.init(image, x, y);
+			xcolor.init(m_image, x, y);
 			xcolor.splatX();
 			
-			ycolor.init(image, x, y);
+			ycolor.init(m_image, x, y);
 			ycolor.splatY();
 
 			if (compressionOptions.quality == Quality_Highest)
 			{
-				compressBlock_BruteForce(xcolor, &block.x);
-				compressBlock_BruteForce(ycolor, &block.y);
+				OptimalCompress::compressDXT5A(xcolor, &block.x);
+				OptimalCompress::compressDXT5A(ycolor, &block.y);
 			}
 			else
 			{

src/nvtt/CompressDXT.h

@@ -32,26 +32,46 @@ namespace nv
 	class Image;
 	class FloatImage;
 
-	void doPrecomputation();
-	
-	// Fast compressors.
-	void fastCompressDXT1(const Image * image, const nvtt::OutputOptions::Private & outputOptions);
-	void fastCompressDXT1a(const Image * image, const nvtt::OutputOptions::Private & outputOptions);
-	void fastCompressDXT3(const Image * image, const nvtt::OutputOptions::Private & outputOptions);
-	void fastCompressDXT5(const Image * image, const nvtt::OutputOptions::Private & outputOptions);
-	void fastCompressDXT5n(const Image * image, const nvtt::OutputOptions::Private & outputOptions);
-	void fastCompressBC4(const Image * image, const nvtt::OutputOptions::Private & outputOptions);
-	void fastCompressBC5(const Image * image, const nvtt::OutputOptions::Private & outputOptions);
+	class FastCompressor
+	{
+	public:
+		FastCompressor();
+		~FastCompressor();
+
+		void setImage(const Image * image, nvtt::AlphaMode alphaMode);
+
+		void compressDXT1(const nvtt::OutputOptions::Private & outputOptions);
+		void compressDXT1a(const nvtt::OutputOptions::Private & outputOptions);
+		void compressDXT3(const nvtt::OutputOptions::Private & outputOptions);
+		void compressDXT5(const nvtt::OutputOptions::Private & outputOptions);
+		void compressDXT5n(const nvtt::OutputOptions::Private & outputOptions);
+
+	private:
+		const Image * m_image;
+		nvtt::AlphaMode m_alphaMode;
+	};
+
+	class SlowCompressor
+	{
+	public:
+		SlowCompressor();
+		~SlowCompressor();
+
+		void setImage(const Image * image, nvtt::AlphaMode alphaMode);
+
+		void compressDXT1(const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
+		void compressDXT1a(const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
+		void compressDXT3(const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
+		void compressDXT5(const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
+		void compressDXT5n(const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
+		void compressBC4(const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
+		void compressBC5(const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
+
+	private:
+		const Image * m_image;
+		nvtt::AlphaMode m_alphaMode;
+	};
 
-	// Normal compressors.
-	void compressDXT1(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions);
-	void compressDXT1a(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 compressDXT5n(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions);
-	void compressBC4(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions);
-	void compressBC5(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions);
-	
 	// External compressors.
 #if defined(HAVE_S3QUANT)
 	void s3CompressDXT1(const Image * image, const nvtt::OutputOptions::Private & outputOptions);

src/nvtt/CompressRGB.cpp

@@ -123,7 +123,7 @@ void nv::compressRGB(const Image * image, const OutputOptions::Private & outputO
 			}
 			
 			// Zero padding.
-			for (uint x = w; x < pitch; x++)
+			for (uint x = w * byteCount; x < pitch; x++)
 			{
 				*(dst + x) = 0;
 			}

src/nvtt/Compressor.cpp

@@ -34,6 +34,7 @@
 #include <nvimage/Filter.h>
 #include <nvimage/Quantize.h>
 #include <nvimage/NormalMap.h>
+#include <nvimage/PixelFormat.h>
 
 #include "Compressor.h"
 #include "InputOptions.h"
@@ -41,7 +42,6 @@
 #include "OutputOptions.h"
 
 #include "CompressDXT.h"
-#include "FastCompressDXT.h"
 #include "CompressRGB.h"
 #include "cuda/CudaUtils.h"
 #include "cuda/CudaCompressDXT.h"
@@ -200,7 +200,7 @@ namespace nvtt
 		AutoPtr<FloatImage> m_floatImage;
 	};
 
-}
+} // nvtt namespace
 
 
 Compressor::Compressor() : m(*new Compressor::Private())
@@ -211,6 +211,10 @@ Compressor::Compressor() : m(*new Compressor::Private())
 
 	if (m.cudaEnabled)
 	{
+		// Select fastest CUDA device.
+		int device = cuda::getFastestDevice();
+		cuda::setDevice(device);
+		
 		m.cuda = new CudaCompressor();
 
 		if (!m.cuda->isValid())
@@ -224,6 +228,7 @@ Compressor::Compressor() : m(*new Compressor::Private())
 Compressor::~Compressor()
 {
 	delete &m;
+	cuda::exit();
 }
 
 
@@ -237,6 +242,10 @@ void Compressor::enableCudaAcceleration(bool 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())
@@ -422,7 +431,7 @@ bool Compressor::Private::compressMipmaps(uint f, const InputOptions::Private &
 		
 		quantizeMipmap(mipmap, compressionOptions);
 
-		compressMipmap(mipmap, compressionOptions, outputOptions);
+		compressMipmap(mipmap, inputOptions, compressionOptions, outputOptions);
 
 		// Compute extents of next mipmap:
 		w = max(1U, w / 2);
@@ -571,7 +580,7 @@ void Compressor::Private::scaleMipmap(Mipmap & mipmap, const InputOptions::Priva
 
 	// Resize image. 
 	BoxFilter boxFilter;
-	mipmap.setImage(mipmap.asFloatImage()->downSample(boxFilter, w, h, (FloatImage::WrapMode)inputOptions.wrapMode));
+	mipmap.setImage(mipmap.asFloatImage()->resize(boxFilter, w, h, (FloatImage::WrapMode)inputOptions.wrapMode));
 }
 
 
@@ -618,13 +627,6 @@ void Compressor::Private::quantizeMipmap(Mipmap & mipmap, const CompressionOptio
 {
 	nvDebugCheck(mipmap.asFixedImage() != NULL);
 
-	if (compressionOptions.enableColorDithering)
-	{
-		if (compressionOptions.format >= Format_DXT1 && compressionOptions.format <= Format_DXT5)
-		{
-			Quantize::FloydSteinberg_RGB16(mipmap.asMutableFixedImage());
-		}
-	}
 	if (compressionOptions.binaryAlpha)
 	{
 		if (compressionOptions.enableAlphaDithering)
@@ -636,30 +638,68 @@ void Compressor::Private::quantizeMipmap(Mipmap & mipmap, const CompressionOptio
 			Quantize::BinaryAlpha(mipmap.asMutableFixedImage(), compressionOptions.alphaThreshold);
 		}
 	}
-	else
+
+	if (compressionOptions.enableColorDithering || compressionOptions.enableAlphaDithering)
 	{
+		uint rsize = 8;
+		uint gsize = 8;
+		uint bsize = 8;
+		uint asize = 8;
+
+		if (compressionOptions.enableColorDithering)
+		{
+			if (compressionOptions.format >= Format_DXT1 && compressionOptions.format <= Format_DXT5)
+			{
+				rsize = 5;
+				gsize = 6;
+				bsize = 5;
+			}
+			else if (compressionOptions.format == Format_RGB)
+			{
+				uint rshift, gshift, bshift;
+				PixelFormat::maskShiftAndSize(compressionOptions.rmask, &rshift, &rsize);
+				PixelFormat::maskShiftAndSize(compressionOptions.gmask, &gshift, &gsize);
+				PixelFormat::maskShiftAndSize(compressionOptions.bmask, &bshift, &bsize);
+			}
+		}
+
 		if (compressionOptions.enableAlphaDithering)
 		{
 			if (compressionOptions.format == Format_DXT3)
 			{
-				Quantize::Alpha4(mipmap.asMutableFixedImage());
+				asize = 4;
 			}
-			else if (compressionOptions.format == Format_DXT1a)
+			else if (compressionOptions.format == Format_RGB)
 			{
-				Quantize::BinaryAlpha(mipmap.asMutableFixedImage(), compressionOptions.alphaThreshold);
+				uint ashift;
+				PixelFormat::maskShiftAndSize(compressionOptions.amask, &ashift, &asize);
 			}
 		}
+
+		if (compressionOptions.binaryAlpha)
+		{
+			asize = 8; // Already quantized.
+		}
+
+		Quantize::FloydSteinberg(mipmap.asMutableFixedImage(), rsize, gsize, bsize, asize);
 	}
 }
 
 
 // Compress the given mipmap.
-bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions) const
+bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const InputOptions::Private & inputOptions, const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions) const
 {
 	const Image * image = mipmap.asFixedImage();
-
 	nvDebugCheck(image != NULL);
 
+	FastCompressor fast;
+	fast.setImage(image, inputOptions.alphaMode);
+
+	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)
 	{
 		compressRGB(image, outputOptions, compressionOptions);
@@ -683,18 +723,19 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio
 #endif
 		if (compressionOptions.quality == Quality_Fastest)
 		{
-			fastCompressDXT1(image, outputOptions);
+			fast.compressDXT1(outputOptions);
 		}
 		else
 		{
-			if (cudaEnabled)
+			if (useCuda)
 			{
 				nvDebugCheck(cudaSupported);
-				cuda->compressDXT1(image, outputOptions, compressionOptions);
+				cuda->setImage(image, inputOptions.alphaMode);
+				cuda->compressDXT1(compressionOptions, outputOptions);
 			}
 			else
 			{
-				compressDXT1(image, outputOptions, compressionOptions);
+				slow.compressDXT1(compressionOptions, outputOptions);
 			}
 		}
 	}
@@ -702,18 +743,18 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio
 	{
 		if (compressionOptions.quality == Quality_Fastest)
 		{
-			fastCompressDXT1a(image, outputOptions);
+			fast.compressDXT1a(outputOptions);
 		}
 		else
 		{
-			if (cudaEnabled)
+			if (useCuda)
 			{
 				nvDebugCheck(cudaSupported);
-				/*cuda*/compressDXT1a(image, outputOptions, compressionOptions);
+				/*cuda*/slow.compressDXT1a(compressionOptions, outputOptions);
 			}
 			else
 			{
-				compressDXT1a(image, outputOptions, compressionOptions);
+				slow.compressDXT1a(compressionOptions, outputOptions);
 			}
 		}
 	}
@@ -721,18 +762,19 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio
 	{
 		if (compressionOptions.quality == Quality_Fastest)
 		{
-			fastCompressDXT3(image, outputOptions);
+			fast.compressDXT3(outputOptions);
 		}
 		else
 		{
-			if (cudaEnabled)
+			if (useCuda)
 			{
 				nvDebugCheck(cudaSupported);
-				cuda->compressDXT3(image, outputOptions, compressionOptions);
+				cuda->setImage(image, inputOptions.alphaMode);
+				cuda->compressDXT3(compressionOptions, outputOptions);
 			}
 			else
 			{
-				compressDXT3(image, outputOptions, compressionOptions);
+				slow.compressDXT3(compressionOptions, outputOptions);
 			}
 		}
 	}
@@ -740,18 +782,19 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio
 	{
 		if (compressionOptions.quality == Quality_Fastest)
 		{
-			fastCompressDXT5(image, outputOptions);
+			fast.compressDXT5(outputOptions);
 		}
 		else
 		{
-			if (cudaEnabled)
+			if (useCuda)
 			{
 				nvDebugCheck(cudaSupported);
-				cuda->compressDXT5(image, outputOptions, compressionOptions);
+				cuda->setImage(image, inputOptions.alphaMode);
+				cuda->compressDXT5(compressionOptions, outputOptions);
 			}
 			else
 			{
-				compressDXT5(image, outputOptions, compressionOptions);
+				slow.compressDXT5(compressionOptions, outputOptions);
 			}
 		}
 	}
@@ -759,20 +802,20 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio
 	{
 		if (compressionOptions.quality == Quality_Fastest)
 		{
-			fastCompressDXT5n(image, outputOptions);
+			fast.compressDXT5n(outputOptions);
 		}
 		else
 		{
-			compressDXT5n(image, outputOptions, compressionOptions);
+			slow.compressDXT5n(compressionOptions, outputOptions);
 		}
 	}
 	else if (compressionOptions.format == Format_BC4)
 	{
-		compressBC4(image, outputOptions, compressionOptions);
+		slow.compressBC4(compressionOptions, outputOptions);
 	}
 	else if (compressionOptions.format == Format_BC5)
 	{
-		compressBC5(image, outputOptions, compressionOptions);
+		slow.compressBC5(compressionOptions, outputOptions);
 	}
 
 	return true;

src/nvtt/Compressor.h

@@ -60,7 +60,7 @@ namespace nvtt
 		void scaleMipmap(Mipmap & mipmap, const InputOptions::Private & inputOptions, uint w, uint h, uint d) const;
 		void processInputImage(Mipmap & mipmap, const InputOptions::Private & inputOptions) const;
 		void quantizeMipmap(Mipmap & mipmap, const CompressionOptions::Private & compressionOptions) const;
-		bool compressMipmap(const Mipmap & mipmap, const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions) const;
+		bool compressMipmap(const Mipmap & mipmap, const InputOptions::Private & inputOptions, const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions) const;
 
 
 	public:

src/nvtt/FastCompressDXT.cpp

@@ -1,456 +0,0 @@
-// Copyright NVIDIA Corporation 2007 -- Ignacio Castano <icastano@nvidia.com>
-// 
-// Permission is hereby granted, free of charge, to any person
-// obtaining a copy of this software and associated documentation
-// files (the "Software"), to deal in the Software without
-// restriction, including without limitation the rights to use,
-// copy, modify, merge, publish, distribute, sublicense, and/or sell
-// copies of the Software, and to permit persons to whom the
-// Software is furnished to do so, subject to the following
-// conditions:
-// 
-// The above copyright notice and this permission notice shall be
-// included in all copies or substantial portions of the Software.
-// 
-// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
-// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
-// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
-// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
-// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
-// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
-// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
-// OTHER DEALINGS IN THE SOFTWARE.
-
-#include <nvmath/Color.h>
-#include <nvimage/ColorBlock.h>
-#include <nvimage/BlockDXT.h>
-
-#include "FastCompressDXT.h"
-
-#if defined(__SSE2__)
-#include <emmintrin.h>
-#endif
-
-#if defined(__SSE__)
-#include <xmmintrin.h>
-#endif
-
-#if defined(__MMX__)
-#include <mmintrin.h>
-#endif
-
-#undef __VEC__
-#if defined(__VEC__)
-#include <altivec.h>
-#undef bool
-#endif
-// Online Resources:
-// - http://www.jasondorie.com/ImageLib.zip
-// - http://homepage.hispeed.ch/rscheidegger/dri_experimental/s3tc_index.html
-// - http://www.sjbrown.co.uk/?article=dxt
-
-using namespace nv;
-
-
-#if defined(__SSE2__) && 0
-
-// @@ TODO
-
-typedef __m128i VectorColor;
-
-inline static __m128i loadColor(Color32 c)
-{
-	return ...;
-}
-
-inline static __m128i absoluteDifference(__m128i a, __m128i b)
-{
-	return ...;
-}
-	
-inline uint colorDistance(__m128i a, __m128i b)
-{
-	return 0;
-}
-
-#elif defined(__MMX__) && 0
-
-typedef __m64 VectorColor;
-
-inline static __m64 loadColor(Color32 c)
-{
-	return _mm_unpacklo_pi8(_mm_cvtsi32_si64(c), _mm_setzero_si64());
-}
-
-inline static __m64 absoluteDifference(__m64 a, __m64 b)
-{
-	// = |a-b| or |b-a|
-	return _mm_or_si64(_mm_subs_pu16(a, b), _mm_subs_pu16(b, a));
-}
-	
-inline uint colorDistance(__m64 a, __m64 b)
-{
-	union {
-		__m64 v;
-		uint16 part[4];
-	} s;
-	
-	s.v = absoluteDifference(a, b);
-		
-	// @@ This is very slow!
-	return s.part[0] + s.part[1] + s.part[2] + s.part[3];
-}
-
-#define vectorEnd	_mm_empty
-
-#elif defined(__VEC__)
-
-typedef vector signed int VectorColor;
-
-inline static vector signed int loadColor(Color32 c)
-{
-	return (vector signed int) (c.r, c.g, c.b, c.a);
-}
-
-// Get the absolute distance between the given colors.
-inline static uint colorDistance(vector signed int c0, vector signed int c1)
-{
-	int result;
-	vector signed int v = vec_sums(vec_abs(vec_sub(c0, c1)), (vector signed int)0);
-	vec_ste(vec_splat(v, 3), 0, &result);
-	return result;
-}
-
-inline void vectorEnd()
-{
-}
-
-#else
-
-typedef Color32 VectorColor;
-
-inline static Color32 loadColor(Color32 c)
-{
-	return c;
-}
-
-inline static Color32 premultiplyAlpha(Color32 c)
-{
-	Color32 pm;
-	pm.r = (c.r * c.a) >> 8;
-	pm.g = (c.g * c.a) >> 8;
-	pm.b = (c.b * c.a) >> 8;
-	pm.a = c.a;
-	return pm;
-}
-
-inline static uint sqr(uint s)
-{
-	return s*s;
-}
-
-// Get the absolute distance between the given colors.
-inline static uint colorDistance(Color32 c0, Color32 c1)
-{
-	return sqr(c0.r - c1.r) + sqr(c0.g - c1.g) + sqr(c0.b - c1.b);
-	//return abs(c0.r - c1.r) + abs(c0.g - c1.g) + abs(c0.b - c1.b);
-}
-
-inline void vectorEnd()
-{
-}
-
-#endif
-
-
-inline static uint computeIndices(const ColorBlock & rgba, const Color32 palette[4])
-{
-	const VectorColor vcolor0 = loadColor(palette[0]);
-	const VectorColor vcolor1 = loadColor(palette[1]);
-	const VectorColor vcolor2 = loadColor(palette[2]);
-	const VectorColor vcolor3 = loadColor(palette[3]);
-	
-	uint indices = 0;
-	for(int i = 0; i < 16; i++) {
-		const VectorColor vcolor = loadColor(rgba.color(i));
-		
-		uint d0 = colorDistance(vcolor0, vcolor);
-		uint d1 = colorDistance(vcolor1, vcolor);
-		uint d2 = colorDistance(vcolor2, vcolor);
-		uint d3 = colorDistance(vcolor3, vcolor);
-		
-		uint b0 = d0 > d3;
-		uint b1 = d1 > d2;
-		uint b2 = d0 > d2;
-		uint b3 = d1 > d3;
-		uint b4 = d2 > d3;
-		
-		uint x0 = b1 & b2;
-		uint x1 = b0 & b3;
-		uint x2 = b0 & b4;
-		
-		indices |= (x2 | ((x0 | x1) << 1)) << (2 * i);
-	}
-
-	vectorEnd();
-	return indices;
-}
-
-
-// Compressor that uses bounding box.
-void nv::compressBlock_BoundsRange(const ColorBlock & rgba, BlockDXT1 * block)
-{
-	Color32 c0, c1;
-	rgba.boundsRange(&c1, &c0);
-	
-	block->col0 = toColor16(c0);
-	block->col1 = toColor16(c1);
-	
-	nvDebugCheck(block->col0.u > block->col1.u);
-	
-	// Use 4 color mode only.
-	//if (block->col0.u < block->col1.u) {
-	//	swap(block->col0.u, block->col1.u);
-	//}
-	
-	Color32 palette[4];
-	block->evaluatePalette4(palette);
-	
-	block->indices = computeIndices(rgba, palette);
-}
-
-
-
-
-
-
-
-
-// Encode DXT3 block.
-void nv::compressBlock_BoundsRange(const ColorBlock & rgba, BlockDXT3 * block)
-{
-	compressBlock_BoundsRange(rgba, &block->color);
-	compressBlock(rgba, &block->alpha);
-}
-
-// Encode DXT3 alpha block.
-void nv::compressBlock(const ColorBlock & rgba, AlphaBlockDXT3 * block)
-{
-	block->alpha0 = rgba.color(0).a >> 4;
-	block->alpha1 = rgba.color(1).a >> 4;
-	block->alpha2 = rgba.color(2).a >> 4;
-	block->alpha3 = rgba.color(3).a >> 4;
-	block->alpha4 = rgba.color(4).a >> 4;
-	block->alpha5 = rgba.color(5).a >> 4;
-	block->alpha6 = rgba.color(6).a >> 4;
-	block->alpha7 = rgba.color(7).a >> 4;
-	block->alpha8 = rgba.color(8).a >> 4;
-	block->alpha9 = rgba.color(9).a >> 4;
-	block->alphaA = rgba.color(10).a >> 4;
-	block->alphaB = rgba.color(11).a >> 4;
-	block->alphaC = rgba.color(12).a >> 4;
-	block->alphaD = rgba.color(13).a >> 4;
-	block->alphaE = rgba.color(14).a >> 4;
-	block->alphaF = rgba.color(15).a >> 4;
-}
-
-
-
-static uint computeAlphaIndices(const ColorBlock & rgba, AlphaBlockDXT5 * block)
-{
-	uint8 alphas[8];
-	block->evaluatePalette(alphas);
-
-	uint totalError = 0;
-
-	for (uint i = 0; i < 16; i++)
-	{
-		uint8 alpha = rgba.color(i).a;
-
-		uint besterror = 256*256;
-		uint best = 8;
-		for(uint p = 0; p < 8; p++)
-		{
-			int d = alphas[p] - alpha;
-			uint error = d * d;
-
-			if (error < besterror)
-			{
-				besterror = error;
-				best = p;
-			}
-		}
-		nvDebugCheck(best < 8);
-
-		totalError += besterror;
-		block->setIndex(i, best);
-	}
-
-	return totalError;
-}
-
-static uint computeAlphaError(const ColorBlock & rgba, const AlphaBlockDXT5 * block)
-{
-	uint8 alphas[8];
-	block->evaluatePalette(alphas);
-
-	uint totalError = 0;
-
-	for (uint i = 0; i < 16; i++)
-	{
-		uint8 alpha = rgba.color(i).a;
-
-		uint besterror = 256*256;
-		uint best;
-		for(uint p = 0; p < 8; p++)
-		{
-			int d = alphas[p] - alpha;
-			uint error = d * d;
-
-			if (error < besterror)
-			{
-				besterror = error;
-				best = p;
-			}
-		}
-
-		totalError += besterror;
-	}
-
-	return totalError;
-}
-
-
-void nv::compressBlock_BoundsRange(const ColorBlock & rgba, BlockDXT5 * block)
-{
-	Color32 c0, c1;
-	rgba.boundsRangeAlpha(&c1, &c0);
-	
-	block->color.col0 = toColor16(c0);
-	block->color.col1 = toColor16(c1);
-	
-	nvDebugCheck(block->color.col0.u > block->color.col1.u);
-	
-	Color32 palette[4];
-	block->color.evaluatePalette4(palette);
-	
-	block->color.indices = computeIndices(rgba, palette);
-	
-	nvDebugCheck(c0.a <= c1.a);
-	
-	block->alpha.alpha0 = c0.a;
-	block->alpha.alpha1 = c1.a;
-	
-	computeAlphaIndices(rgba, &block->alpha);
-}
-
-
-uint nv::compressBlock_BoundsRange(const ColorBlock & rgba, AlphaBlockDXT5 * block)
-{
-	uint8 alpha0 = 0;
-	uint8 alpha1 = 255;
-
-	// Get min/max alpha.
-	for (uint i = 0; i < 16; i++)
-	{
-		uint8 alpha = rgba.color(i).a;
-		alpha0 = max(alpha0, alpha);
-		alpha1 = min(alpha1, alpha);
-	}
-
-	alpha0 = alpha0 - (alpha0 - alpha1) / 32;
-	alpha1 = alpha1 + (alpha0 - alpha1) / 32;
-
-	AlphaBlockDXT5 block0;
-	block0.alpha0 = alpha0;
-	block0.alpha1 = alpha1;
-	uint error0 = computeAlphaIndices(rgba, &block0);
-
-	AlphaBlockDXT5 block1;
-	block1.alpha0 = alpha1;
-	block1.alpha1 = alpha0;
-	uint error1 = computeAlphaIndices(rgba, &block1);
-
-	if (error0 < error1)
-	{
-		*block = block0;
-		return error0;
-	}
-	else
-	{
-		*block = block1;
-		return error1;
-	}
-}
-
-uint nv::compressBlock_BruteForce(const ColorBlock & rgba, AlphaBlockDXT5 * block)
-{
-	uint8 mina = 255;
-	uint8 maxa = 0;
-
-	// Get min/max alpha.
-	for (uint i = 0; i < 16; i++)
-	{
-		uint8 alpha = rgba.color(i).a;
-		mina = min(mina, alpha);
-		maxa = max(maxa, alpha);
-	}
-
-	block->alpha0 = maxa;
-	block->alpha1 = mina;
-
-	/*int centroidDist = 256;
-	int centroid;
-
-	// Get the closest to the centroid.
-	for (uint i = 0; i < 16; i++)
-	{
-		uint8 alpha = rgba.color(i).a;
-		int dist = abs(alpha - (maxa + mina) / 2);
-		if (dist < centroidDist)
-		{
-			centroidDist = dist;
-			centroid = alpha;
-		}
-	}*/
-
-	if (maxa - mina > 8)
-	{
-		int besterror = computeAlphaError(rgba, block);
-		int besta0 = maxa;
-		int besta1 = mina;
-
-		for (int a0 = mina+9; a0 < maxa; a0++)
-		{
-			for (int a1 = mina; a1 < a0-8; a1++)
-			//for (int a1 = mina; a1 < maxa; a1++)
-			{
-				//nvCheck(abs(a1-a0) > 8);
-
-				//if (abs(a0 - a1) < 8) continue;
-				//if ((maxa-a0) + (a1-mina) + min(abs(centroid-a0), abs(centroid-a1)) > besterror)
-				if ((maxa-a0) + (a1-mina) > besterror)
-					continue;
-
-				block->alpha0 = a0;
-				block->alpha1 = a1;
-				int error = computeAlphaError(rgba, block);
-
-				if (error < besterror)
-				{
-					besterror = error;
-					besta0 = a0;
-					besta1 = a1;
-				}
-			}
-		}
-
-		block->alpha0 = besta0;
-		block->alpha1 = besta1;
-	}
-
-	return computeAlphaIndices(rgba, block);
-}
-
-
-

src/nvtt/FastCompressDXT.h

@@ -1,84 +0,0 @@
-// Copyright NVIDIA Corporation 2007 -- Ignacio Castano <icastano@nvidia.com>
-// 
-// Permission is hereby granted, free of charge, to any person
-// obtaining a copy of this software and associated documentation
-// files (the "Software"), to deal in the Software without
-// restriction, including without limitation the rights to use,
-// copy, modify, merge, publish, distribute, sublicense, and/or sell
-// copies of the Software, and to permit persons to whom the
-// Software is furnished to do so, subject to the following
-// conditions:
-// 
-// The above copyright notice and this permission notice shall be
-// included in all copies or substantial portions of the Software.
-// 
-// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
-// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
-// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
-// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
-// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
-// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
-// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
-// OTHER DEALINGS IN THE SOFTWARE.
-
-#ifndef NV_TT_FASTCOMPRESSDXT_H
-#define NV_TT_FASTCOMPRESSDXT_H
-
-#include <nvimage/nvimage.h>
-
-namespace nv
-{
-	struct ColorBlock;
-	struct BlockDXT1;
-	struct BlockDXT3;
-	struct BlockDXT5;
-	struct AlphaBlockDXT3;
-	struct AlphaBlockDXT5;
-
-	// Color compression:
-
-	// Compressor that uses the extremes of the luminance axis.
-//	void compressBlock_DiameterAxis(const ColorBlock & rgba, BlockDXT1 * block);
-
-	// Compressor that uses the extremes of the luminance axis.
-//	void compressBlock_LuminanceAxis(const ColorBlock & rgba, BlockDXT1 * block);
-
-	// Compressor that uses bounding box.
-	void compressBlock_BoundsRange(const ColorBlock & rgba, BlockDXT1 * block);
-
-	// Compressor that uses bounding box and takes alpha into account.
-//	void compressBlock_BoundsRangeAlpha(const ColorBlock & rgba, BlockDXT1 * block);
-
-
-	// Simple, but slow compressor that tests all color pairs.
-//	void compressBlock_TestAllPairs(const ColorBlock & rgba, BlockDXT1 * block);
-	
-	// Brute force 6d search along the best fit axis.
-//	void compressBlock_AnalyzeBestFitAxis(const ColorBlock & rgba, BlockDXT1 * block);
-
-	// Spatial greedy search.
-//	void refineSolution_1dSearch(const ColorBlock & rgba, BlockDXT1 * block);
-//	void refineSolution_3dSearch(const ColorBlock & rgba, BlockDXT1 * block);
-//	void refineSolution_6dSearch(const ColorBlock & rgba, BlockDXT1 * block);
-	
-	// Brute force compressor for DXT5n
-//	void compressGreenBlock_BruteForce(const ColorBlock & rgba, BlockDXT1 * block);
-
-	// Minimize error of the endpoints.
-//	void optimizeEndPoints(const ColorBlock & rgba, BlockDXT1 * block);
-	
-//	uint blockError(const ColorBlock & rgba, const BlockDXT1 & block);
-//	uint blockError(const ColorBlock & rgba, const AlphaBlockDXT5 & block);
-
-	// Alpha compression:
-	void compressBlock(const ColorBlock & rgba, AlphaBlockDXT3 * block);
-	void compressBlock_BoundsRange(const ColorBlock & rgba, BlockDXT3 * block);
-	void compressBlock_BoundsRange(const ColorBlock & rgba, BlockDXT5 * block);
-
-	uint compressBlock_BoundsRange(const ColorBlock & rgba, AlphaBlockDXT5 * block);
-	uint compressBlock_BruteForce(const ColorBlock & rgba, AlphaBlockDXT5 * block);
-//	uint compressBlock_Iterative(const ColorBlock & rgba, AlphaBlockDXT5 * block);
-
-} // nv namespace
-
-#endif // NV_TT_FASTCOMPRESSDXT_H

src/nvtt/OptimalCompressDXT.cpp

@@ -0,0 +1,368 @@
+// Copyright NVIDIA Corporation 2007 -- Ignacio Castano <icastano@nvidia.com>
+// 
+// Permission is hereby granted, free of charge, to any person
+// obtaining a copy of this software and associated documentation
+// files (the "Software"), to deal in the Software without
+// restriction, including without limitation the rights to use,
+// copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the
+// Software is furnished to do so, subject to the following
+// conditions:
+// 
+// The above copyright notice and this permission notice shall be
+// included in all copies or substantial portions of the Software.
+// 
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
+// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
+// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+// OTHER DEALINGS IN THE SOFTWARE.
+
+#include <nvcore/Containers.h> // swap
+
+#include <nvmath/Color.h>
+
+#include <nvimage/ColorBlock.h>
+#include <nvimage/BlockDXT.h>
+
+#include "OptimalCompressDXT.h"
+#include "SingleColorLookup.h"
+
+
+using namespace nv;
+using namespace OptimalCompress;
+
+
+
+namespace
+{
+	static int computeGreenError(const ColorBlock & rgba, const BlockDXT1 * block)
+	{
+		nvDebugCheck(block != NULL);
+
+		int palette[4];
+		palette[0] = (block->col0.g << 2) | (block->col0.g >> 4);
+		palette[1] = (block->col1.g << 2) | (block->col1.g >> 4);
+		palette[2] = (2 * palette[0] + palette[1]) / 3;
+		palette[3] = (2 * palette[1] + palette[0]) / 3;
+
+		int totalError = 0;
+
+		for (int i = 0; i < 16; i++)
+		{
+			const int green = rgba.color(i).g;
+			
+			int error = abs(green - palette[0]);
+			error = min(error, abs(green - palette[1]));
+			error = min(error, abs(green - palette[2]));
+			error = min(error, abs(green - palette[3]));
+			
+			totalError += error;
+		}
+
+		return totalError;
+	}
+
+	static uint computeGreenIndices(const ColorBlock & rgba, const Color32 palette[4])
+	{
+		const int color0 = palette[0].g;
+		const int color1 = palette[1].g;
+		const int color2 = palette[2].g;
+		const int color3 = palette[3].g;
+		
+		uint indices = 0;
+		for (int i = 0; i < 16; i++)
+		{
+			const int color = rgba.color(i).g;
+			
+			uint d0 = abs(color0 - color);
+			uint d1 = abs(color1 - color);
+			uint d2 = abs(color2 - color);
+			uint d3 = abs(color3 - color);
+			
+			uint b0 = d0 > d3;
+			uint b1 = d1 > d2;
+			uint b2 = d0 > d2;
+			uint b3 = d1 > d3;
+			uint b4 = d2 > d3;
+			
+			uint x0 = b1 & b2;
+			uint x1 = b0 & b3;
+			uint x2 = b0 & b4;
+			
+			indices |= (x2 | ((x0 | x1) << 1)) << (2 * i);
+		}
+
+		return indices;
+	}
+
+	// Choose quantized color that produces less error. Used by DXT3 compressor.
+	inline static uint quantize4(uint8 a)
+	{
+		int q0 = (a >> 4) - 1;
+		int q1 = (a >> 4);
+		int q2 = (a >> 4) + 1;
+		
+		q0 = (q0 << 4) | q0;
+		q1 = (q1 << 4) | q1;
+		q2 = (q2 << 4) | q2;
+		
+		int d0 = abs(q0 - a);
+		int d1 = abs(q1 - a);
+		int d2 = abs(q2 - a);
+
+		if (d0 < d1 && d0 < d2) return q0 >> 4;
+		if (d1 < d2) return q1 >> 4;
+		return q2 >> 4;
+	}
+	
+	static uint computeAlphaError(const ColorBlock & rgba, const AlphaBlockDXT5 * block)
+	{
+		uint8 alphas[8];
+		block->evaluatePalette(alphas);
+
+		uint totalError = 0;
+
+		for (uint i = 0; i < 16; i++)
+		{
+			uint8 alpha = rgba.color(i).a;
+
+			uint besterror = 256*256;
+			uint best;
+			for (uint p = 0; p < 8; p++)
+			{
+				int d = alphas[p] - alpha;
+				uint error = d * d;
+
+				if (error < besterror)
+				{
+					besterror = error;
+					best = p;
+				}
+			}
+
+			totalError += besterror;
+		}
+
+		return totalError;
+	}
+	
+	static void computeAlphaIndices(const ColorBlock & rgba, AlphaBlockDXT5 * block)
+	{
+		uint8 alphas[8];
+		block->evaluatePalette(alphas);
+
+		for (uint i = 0; i < 16; i++)
+		{
+			uint8 alpha = rgba.color(i).a;
+
+			uint besterror = 256*256;
+			uint best = 8;
+			for(uint p = 0; p < 8; p++)
+			{
+				int d = alphas[p] - alpha;
+				uint error = d * d;
+
+				if (error < besterror)
+				{
+					besterror = error;
+					best = p;
+				}
+			}
+			nvDebugCheck(best < 8);
+
+			block->setIndex(i, best);
+		}
+	}
+
+} // namespace
+
+
+
+
+
+// Single color compressor, based on:
+// https://mollyrocket.com/forums/viewtopic.php?t=392
+void OptimalCompress::compressDXT1(Color32 c, BlockDXT1 * dxtBlock)
+{
+	dxtBlock->col0.r = OMatch5[c.r][0];
+	dxtBlock->col0.g = OMatch6[c.g][0];
+	dxtBlock->col0.b = OMatch5[c.b][0];
+	dxtBlock->col1.r = OMatch5[c.r][1];
+	dxtBlock->col1.g = OMatch6[c.g][1];
+	dxtBlock->col1.b = OMatch5[c.b][1];
+	dxtBlock->indices = 0xaaaaaaaa;
+
+	if (dxtBlock->col0.u < dxtBlock->col1.u)
+	{
+		swap(dxtBlock->col0.u, dxtBlock->col1.u);
+		dxtBlock->indices ^= 0x55555555;
+	}
+}
+
+void OptimalCompress::compressDXT1a(Color32 rgba, BlockDXT1 * dxtBlock)
+{
+	if (rgba.a < 128)
+	{
+		dxtBlock->col0.u = 0;
+		dxtBlock->col1.u = 0;
+		dxtBlock->indices = 0xFFFFFFFF;
+	}
+	else
+	{
+		compressDXT1(rgba, dxtBlock);
+	}
+}
+
+
+// Brute force green channel compressor
+void OptimalCompress::compressDXT1G(const ColorBlock & rgba, BlockDXT1 * block)
+{
+	nvDebugCheck(block != NULL);
+	
+	uint8 ming = 63;
+	uint8 maxg = 0;
+	
+	// Get min/max green.
+	for (uint i = 0; i < 16; i++)
+	{
+		uint8 green = rgba.color(i).g >> 2;
+		ming = min(ming, green);
+		maxg = max(maxg, green);
+	}
+
+	block->col0.r = 31;
+	block->col1.r = 31;
+	block->col0.g = maxg;
+	block->col1.g = ming;
+	block->col0.b = 0;
+	block->col1.b = 0;
+
+	if (maxg - ming > 4)
+	{
+		int besterror = computeGreenError(rgba, block);
+		int bestg0 = maxg;
+		int bestg1 = ming;
+		
+		for (int g0 = ming+5; g0 < maxg; g0++)
+		{
+			for (int g1 = ming; g1 < g0-4; g1++)
+			{
+				if ((maxg-g0) + (g1-ming) > besterror)
+					continue;
+				
+				block->col0.g = g0;
+				block->col1.g = g1;
+				int error = computeGreenError(rgba, block);
+				
+				if (error < besterror)
+				{
+					besterror = error;
+					bestg0 = g0;
+					bestg1 = g1;
+				}
+			}
+		}
+		
+		block->col0.g = bestg0;
+		block->col1.g = bestg1;
+	}
+	
+	Color32 palette[4];
+	block->evaluatePalette(palette);
+	block->indices = computeGreenIndices(rgba, palette);
+}
+
+void OptimalCompress::compressDXT3A(const ColorBlock & rgba, AlphaBlockDXT3 * dxtBlock)
+{
+	dxtBlock->alpha0 = quantize4(rgba.color(0).a);
+	dxtBlock->alpha1 = quantize4(rgba.color(1).a);
+	dxtBlock->alpha2 = quantize4(rgba.color(2).a);
+	dxtBlock->alpha3 = quantize4(rgba.color(3).a);
+	dxtBlock->alpha4 = quantize4(rgba.color(4).a);
+	dxtBlock->alpha5 = quantize4(rgba.color(5).a);
+	dxtBlock->alpha6 = quantize4(rgba.color(6).a);
+	dxtBlock->alpha7 = quantize4(rgba.color(7).a);
+	dxtBlock->alpha8 = quantize4(rgba.color(8).a);
+	dxtBlock->alpha9 = quantize4(rgba.color(9).a);
+	dxtBlock->alphaA = quantize4(rgba.color(10).a);
+	dxtBlock->alphaB = quantize4(rgba.color(11).a);
+	dxtBlock->alphaC = quantize4(rgba.color(12).a);
+	dxtBlock->alphaD = quantize4(rgba.color(13).a);
+	dxtBlock->alphaE = quantize4(rgba.color(14).a);
+	dxtBlock->alphaF = quantize4(rgba.color(15).a);
+}
+
+
+void OptimalCompress::compressDXT5A(const ColorBlock & rgba, AlphaBlockDXT5 * dxtBlock)
+{
+	uint8 mina = 255;
+	uint8 maxa = 0;
+
+	// Get min/max alpha.
+	for (uint i = 0; i < 16; i++)
+	{
+		uint8 alpha = rgba.color(i).a;
+		mina = min(mina, alpha);
+		maxa = max(maxa, alpha);
+	}
+
+	dxtBlock->alpha0 = maxa;
+	dxtBlock->alpha1 = mina;
+
+	/*int centroidDist = 256;
+	int centroid;
+
+	// Get the closest to the centroid.
+	for (uint i = 0; i < 16; i++)
+	{
+		uint8 alpha = rgba.color(i).a;
+		int dist = abs(alpha - (maxa + mina) / 2);
+		if (dist < centroidDist)
+		{
+			centroidDist = dist;
+			centroid = alpha;
+		}
+	}*/
+
+	if (maxa - mina > 8)
+	{
+		int besterror = computeAlphaError(rgba, dxtBlock);
+		int besta0 = maxa;
+		int besta1 = mina;
+
+		for (int a0 = mina+9; a0 < maxa; a0++)
+		{
+			for (int a1 = mina; a1 < a0-8; a1++)
+			//for (int a1 = mina; a1 < maxa; a1++)
+			{
+				//nvCheck(abs(a1-a0) > 8);
+
+				//if (abs(a0 - a1) < 8) continue;
+				//if ((maxa-a0) + (a1-mina) + min(abs(centroid-a0), abs(centroid-a1)) > besterror)
+				if ((maxa-a0) + (a1-mina) > besterror)
+					continue;
+
+				dxtBlock->alpha0 = a0;
+				dxtBlock->alpha1 = a1;
+				int error = computeAlphaError(rgba, dxtBlock);
+
+				if (error < besterror)
+				{
+					besterror = error;
+					besta0 = a0;
+					besta1 = a1;
+				}
+			}
+		}
+
+		dxtBlock->alpha0 = besta0;
+		dxtBlock->alpha1 = besta1;
+	}
+
+	computeAlphaIndices(rgba, dxtBlock);
+}
+

src/nvtt/OptimalCompressDXT.h

@@ -0,0 +1,49 @@
+// Copyright NVIDIA Corporation 2008 -- Ignacio Castano <icastano@nvidia.com>
+// 
+// Permission is hereby granted, free of charge, to any person
+// obtaining a copy of this software and associated documentation
+// files (the "Software"), to deal in the Software without
+// restriction, including without limitation the rights to use,
+// copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the
+// Software is furnished to do so, subject to the following
+// conditions:
+// 
+// The above copyright notice and this permission notice shall be
+// included in all copies or substantial portions of the Software.
+// 
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
+// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
+// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
+// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+// OTHER DEALINGS IN THE SOFTWARE.
+
+#ifndef NV_TT_OPTIMALCOMPRESSDXT_H
+#define NV_TT_OPTIMALCOMPRESSDXT_H
+
+#include <nvimage/nvimage.h>
+
+namespace nv
+{
+	struct ColorBlock;
+	struct BlockDXT1;
+	struct BlockDXT3;
+	struct BlockDXT5;
+	struct AlphaBlockDXT3;
+	struct AlphaBlockDXT5;
+
+	namespace OptimalCompress
+	{
+		void compressDXT1(Color32 rgba, BlockDXT1 * dxtBlock);
+		void compressDXT1a(Color32 rgba, BlockDXT1 * dxtBlock);
+		
+		void compressDXT1G(const ColorBlock & rgba, BlockDXT1 * block);
+		void compressDXT3A(const ColorBlock & rgba, AlphaBlockDXT3 * dxtBlock);
+		void compressDXT5A(const ColorBlock & rgba, AlphaBlockDXT5 * dxtBlock);
+	}
+} // nv namespace
+
+#endif // NV_TT_OPTIMALCOMPRESSDXT_H

src/nvtt/QuickCompressDXT.cpp

@@ -27,7 +27,7 @@
 #include <nvimage/BlockDXT.h>
 
 #include "QuickCompressDXT.h"
-#include "SingleColorLookup.h"
+#include "OptimalCompressDXT.h"
 
 
 using namespace nv;
@@ -288,70 +288,6 @@ static void optimizeEndPoints4(Vector3 block[16], BlockDXT1 * dxtBlock)
 	dxtBlock->indices = computeIndices3(block, a, b);
 }*/
 
-namespace
-{
-	static int computeGreenError(const ColorBlock & rgba, const BlockDXT1 * block)
-	{
-		nvDebugCheck(block != NULL);
-
-		int palette[4];
-		palette[0] = (block->col0.g << 2) | (block->col0.g >> 4);
-		palette[1] = (block->col1.g << 2) | (block->col1.g >> 4);
-		palette[2] = (2 * palette[0] + palette[1]) / 3;
-		palette[3] = (2 * palette[1] + palette[0]) / 3;
-
-		int totalError = 0;
-
-		for (int i = 0; i < 16; i++)
-		{
-			const int green = rgba.color(i).g;
-			
-			int error = abs(green - palette[0]);
-			error = min(error, abs(green - palette[1]));
-			error = min(error, abs(green - palette[2]));
-			error = min(error, abs(green - palette[3]));
-			
-			totalError += error;
-		}
-
-		return totalError;
-	}
-
-	static uint computeGreenIndices(const ColorBlock & rgba, const Color32 palette[4])
-	{
-		const int color0 = palette[0].g;
-		const int color1 = palette[1].g;
-		const int color2 = palette[2].g;
-		const int color3 = palette[3].g;
-		
-		uint indices = 0;
-		for (int i = 0; i < 16; i++)
-		{
-			const int color = rgba.color(i).g;
-			
-			uint d0 = abs(color0 - color);
-			uint d1 = abs(color1 - color);
-			uint d2 = abs(color2 - color);
-			uint d3 = abs(color3 - color);
-			
-			uint b0 = d0 > d3;
-			uint b1 = d1 > d2;
-			uint b2 = d0 > d2;
-			uint b3 = d1 > d3;
-			uint b4 = d2 > d3;
-			
-			uint x0 = b1 & b2;
-			uint x1 = b0 & b3;
-			uint x2 = b0 & b4;
-			
-			indices |= (x2 | ((x0 | x1) << 1)) << (2 * i);
-		}
-
-		return indices;
-	}
-
-} // namespace
-
 namespace
 {
 
@@ -503,79 +439,63 @@ namespace
 
 
 
-
-
-// Single color compressor, based on:
-// https://mollyrocket.com/forums/viewtopic.php?t=392
-void QuickCompress::compressDXT1(Color32 c, BlockDXT1 * dxtBlock)
-{
-	dxtBlock->col0.r = OMatch5[c.r][0];
-	dxtBlock->col0.g = OMatch6[c.g][0];
-	dxtBlock->col0.b = OMatch5[c.b][0];
-	dxtBlock->col1.r = OMatch5[c.r][1];
-	dxtBlock->col1.g = OMatch6[c.g][1];
-	dxtBlock->col1.b = OMatch5[c.b][1];
-	dxtBlock->indices = 0xaaaaaaaa;
-
-	if (dxtBlock->col0.u < dxtBlock->col1.u)
-	{
-		swap(dxtBlock->col0.u, dxtBlock->col1.u);
-		dxtBlock->indices ^= 0x55555555;
-	}
-}
-
 void QuickCompress::compressDXT1(const ColorBlock & rgba, BlockDXT1 * dxtBlock)
 {
-	// read block
-	Vector3 block[16];
-	extractColorBlockRGB(rgba, block);
-	
-	// find min and max colors
-	Vector3 maxColor, minColor;
-	findMinMaxColorsBox(block, 16, &maxColor, &minColor);
-	
-	selectDiagonal(block, 16, &maxColor, &minColor);
-	
-	insetBBox(&maxColor, &minColor);
-	
-	uint16 color0 = roundAndExpand(&maxColor);
-	uint16 color1 = roundAndExpand(&minColor);
-
-	if (color0 < color1)
+	if (rgba.isSingleColor())
 	{
-		swap(maxColor, minColor);
-		swap(color0, color1);
-	}
-
-	dxtBlock->col0 = Color16(color0);
-	dxtBlock->col1 = Color16(color1);
-	dxtBlock->indices = computeIndices4(block, maxColor, minColor);
-
-	optimizeEndPoints4(block, dxtBlock);
-}
-
-
-void QuickCompress::compressDXT1a(Color32 rgba, BlockDXT1 * dxtBlock)
-{
-	if (rgba.a == 0)
-	{
-		dxtBlock->col0.u = 0;
-		dxtBlock->col1.u = 0;
-		dxtBlock->indices = 0xFFFFFFFF;
+		OptimalCompress::compressDXT1(rgba.color(0), dxtBlock);
 	}
 	else
 	{
-		compressDXT1(rgba, dxtBlock);
+		// read block
+		Vector3 block[16];
+		extractColorBlockRGB(rgba, block);
+		
+		// find min and max colors
+		Vector3 maxColor, minColor;
+		findMinMaxColorsBox(block, 16, &maxColor, &minColor);
+		
+		selectDiagonal(block, 16, &maxColor, &minColor);
+		
+		insetBBox(&maxColor, &minColor);
+		
+		uint16 color0 = roundAndExpand(&maxColor);
+		uint16 color1 = roundAndExpand(&minColor);
+
+		if (color0 < color1)
+		{
+			swap(maxColor, minColor);
+			swap(color0, color1);
+		}
+
+		dxtBlock->col0 = Color16(color0);
+		dxtBlock->col1 = Color16(color1);
+		dxtBlock->indices = computeIndices4(block, maxColor, minColor);
+
+		optimizeEndPoints4(block, dxtBlock);
 	}
 }
 
+
 void QuickCompress::compressDXT1a(const ColorBlock & rgba, BlockDXT1 * dxtBlock)
 {
-	if (!rgba.hasAlpha())
+	bool hasAlpha = false;
+	
+	for (uint i = 0; i < 16; i++)
+	{
+		if (rgba.color(i).a < 128) {
+			hasAlpha = true;
+			break;
+		}
+	}
+	
+	if (!hasAlpha)
 	{
 		compressDXT1(rgba, dxtBlock);
 	}
-	else
+	// @@ Handle single RGB, with varying alpha? We need tables for single color compressor in 3 color mode.
+	//else if (rgba.isSingleColorNoAlpha()) { ... }
+	else 
 	{
 		// read block
 		Vector3 block[16];
@@ -607,95 +527,14 @@ void QuickCompress::compressDXT1a(const ColorBlock & rgba, BlockDXT1 * dxtBlock)
 }
 
 
-// Brute force green channel compressor
-void QuickCompress::compressDXT1G(const ColorBlock & rgba, BlockDXT1 * block)
-{
-	nvDebugCheck(block != NULL);
-	
-	uint8 ming = 63;
-	uint8 maxg = 0;
-	
-	// Get min/max green.
-	for (uint i = 0; i < 16; i++)
-	{
-		uint8 green = rgba.color(i).g >> 2;
-		ming = min(ming, green);
-		maxg = max(maxg, green);
-	}
-
-	block->col0.r = 31;
-	block->col1.r = 31;
-	block->col0.g = maxg;
-	block->col1.g = ming;
-	block->col0.b = 0;
-	block->col1.b = 0;
-
-	if (maxg - ming > 4)
-	{
-		int besterror = computeGreenError(rgba, block);
-		int bestg0 = maxg;
-		int bestg1 = ming;
-		
-		for (int g0 = ming+5; g0 < maxg; g0++)
-		{
-			for (int g1 = ming; g1 < g0-4; g1++)
-			{
-				if ((maxg-g0) + (g1-ming) > besterror)
-					continue;
-				
-				block->col0.g = g0;
-				block->col1.g = g1;
-				int error = computeGreenError(rgba, block);
-				
-				if (error < besterror)
-				{
-					besterror = error;
-					bestg0 = g0;
-					bestg1 = g1;
-				}
-			}
-		}
-		
-		block->col0.g = bestg0;
-		block->col1.g = bestg1;
-	}
-	
-	Color32 palette[4];
-	block->evaluatePalette(palette);
-	block->indices = computeGreenIndices(rgba, palette);
-}
-
-
-
-void QuickCompress::compressDXT3A(const ColorBlock & rgba, AlphaBlockDXT3 * dxtBlock)
-{
-	// @@ Round instead of truncate. When rounding take into account bit expansion.
-	dxtBlock->alpha0 = rgba.color(0).a >> 4;
-	dxtBlock->alpha1 = rgba.color(1).a >> 4;
-	dxtBlock->alpha2 = rgba.color(2).a >> 4;
-	dxtBlock->alpha3 = rgba.color(3).a >> 4;
-	dxtBlock->alpha4 = rgba.color(4).a >> 4;
-	dxtBlock->alpha5 = rgba.color(5).a >> 4;
-	dxtBlock->alpha6 = rgba.color(6).a >> 4;
-	dxtBlock->alpha7 = rgba.color(7).a >> 4;
-	dxtBlock->alpha8 = rgba.color(8).a >> 4;
-	dxtBlock->alpha9 = rgba.color(9).a >> 4;
-	dxtBlock->alphaA = rgba.color(10).a >> 4;
-	dxtBlock->alphaB = rgba.color(11).a >> 4;
-	dxtBlock->alphaC = rgba.color(12).a >> 4;
-	dxtBlock->alphaD = rgba.color(13).a >> 4;
-	dxtBlock->alphaE = rgba.color(14).a >> 4;
-	dxtBlock->alphaF = rgba.color(15).a >> 4;
-}
-
 void QuickCompress::compressDXT3(const ColorBlock & rgba, BlockDXT3 * dxtBlock)
 {
 	compressDXT1(rgba, &dxtBlock->color);
-	compressDXT3A(rgba, &dxtBlock->alpha);
+	OptimalCompress::compressDXT3A(rgba, &dxtBlock->alpha);
 }
 
 
-void QuickCompress::compressDXT5A(const ColorBlock & rgba, AlphaBlockDXT5 * dxtBlock)
+void QuickCompress::compressDXT5A(const ColorBlock & rgba, AlphaBlockDXT5 * dxtBlock, int iterationCount/*=8*/)
 {
 	uint8 alpha0 = 0;
 	uint8 alpha1 = 255;
@@ -714,8 +553,8 @@ void QuickCompress::compressDXT5A(const ColorBlock & rgba, AlphaBlockDXT5 * dxtB
 	uint besterror = computeAlphaIndices(rgba, &block);
 	
 	AlphaBlockDXT5 bestblock = block;
-	
-	while(true)
+
+	for (int i = 0; i < iterationCount; i++)
 	{
 		optimizeAlpha8(rgba, &block);
 		uint error = computeAlphaIndices(rgba, &block);
@@ -739,9 +578,8 @@ void QuickCompress::compressDXT5A(const ColorBlock & rgba, AlphaBlockDXT5 * dxtB
 	*dxtBlock = bestblock;
 }
 
-void QuickCompress::compressDXT5(const ColorBlock & rgba, BlockDXT5 * dxtBlock)
+void QuickCompress::compressDXT5(const ColorBlock & rgba, BlockDXT5 * dxtBlock, int iterationCount/*=8*/)
 {
 	compressDXT1(rgba, &dxtBlock->color);
-	compressDXT5A(rgba, &dxtBlock->alpha);
+	compressDXT5A(rgba, &dxtBlock->alpha, iterationCount);
 }
-

src/nvtt/QuickCompressDXT.h

@@ -37,17 +37,13 @@ namespace nv
 
 	namespace QuickCompress
 	{
-		void compressDXT1(Color32 rgba, BlockDXT1 * dxtBlock);
 		void compressDXT1(const ColorBlock & rgba, BlockDXT1 * dxtBlock);
-		void compressDXT1a(Color32 rgba, BlockDXT1 * dxtBlock);
 		void compressDXT1a(const ColorBlock & rgba, BlockDXT1 * dxtBlock);
-		void compressDXT1G(const ColorBlock & rgba, BlockDXT1 * block);
 		
-		void compressDXT3A(const ColorBlock & rgba, AlphaBlockDXT3 * dxtBlock);
 		void compressDXT3(const ColorBlock & rgba, BlockDXT3 * dxtBlock);
 		
-		void compressDXT5A(const ColorBlock & rgba, AlphaBlockDXT5 * dxtBlock);
-		void compressDXT5(const ColorBlock & rgba, BlockDXT5 * dxtBlock);
+		void compressDXT5A(const ColorBlock & rgba, AlphaBlockDXT5 * dxtBlock, int iterationCount=8);
+		void compressDXT5(const ColorBlock & rgba, BlockDXT5 * dxtBlock, int iterationCount=8);
 	}
 } // nv namespace
 

src/nvtt/cuda/CompressKernel.cu

@@ -159,7 +159,7 @@ __device__ void loadColorBlock(const uint * image, float3 colors[16], float3 sum
 	}
 }
 
-__device__ void loadColorBlock(const uint * image, float3 colors[16], float3 sums[16], float weights[16], int xrefs[16])
+__device__ void loadColorBlock(const uint * image, float3 colors[16], float3 sums[16], float weights[16], int xrefs[16], int * sameColor)
 {
 	const int bid = blockIdx.x;
 	const int idx = threadIdx.x;
@@ -189,6 +189,11 @@ __device__ void loadColorBlock(const uint * image, float3 colors[16], float3 sum
 		colorSums(colors, sums);
 		float3 axis = bestFitLine(colors, sums[0], kColorMetric);
 
+		*sameColor = (axis == make_float3(0, 0, 0));
+
+		// Single color compressor needs unweighted colors.
+		if (*sameColor) colors[idx] = rawColors[idx];
+
 		dps[idx] = dot(rawColors[idx], axis);
 		
 #if __DEVICE_EMULATION__
@@ -592,6 +597,40 @@ __device__ void evalAllPermutations(const float3 * colors, const float * weights
 }
 */
 
+__device__ void evalLevel4Permutations(const float3 * colors, float3 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 = 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.
+	}
+
+	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;
@@ -627,7 +666,6 @@ __device__ void evalLevel4Permutations(const float3 * colors, const float * weig
 }
 
 
-
 ////////////////////////////////////////////////////////////////////////////////
 // Find index with minimum error
 ////////////////////////////////////////////////////////////////////////////////
@@ -798,6 +836,39 @@ __global__ void compressDXT1(const uint * permutations, const uint * image, uint
 	}
 }
 
+__global__ void compressLevel4DXT1(const uint * permutations, const uint * image, uint2 * result)
+{
+	__shared__ float3 colors[16];
+	__shared__ float3 sums[16];
+	__shared__ int xrefs[16];
+	__shared__ int sameColor;
+	
+	loadColorBlock(image, colors, sums, xrefs, &sameColor);
+
+	__syncthreads();
+
+	if (sameColor)
+	{
+		if (threadIdx.x == 0) saveSingleColorBlockDXT1(colors[0], result);
+		return;
+	}
+
+	ushort bestStart, bestEnd;
+	uint bestPermutation;
+
+	__shared__ float errors[NUM_THREADS];
+
+	evalLevel4Permutations(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 compressWeightedDXT1(const uint * permutations, const uint * image, uint2 * result)
 {
@@ -805,11 +876,18 @@ __global__ void compressWeightedDXT1(const uint * permutations, const uint * ima
 	__shared__ float3 sums[16];
 	__shared__ float weights[16];
 	__shared__ int xrefs[16];
+	__shared__ int sameColor;
 	
-	loadColorBlock(image, colors, sums, weights, xrefs);
+	loadColorBlock(image, colors, sums, weights, xrefs, &sameColor);
 	
 	__syncthreads();
 
+	if (sameColor)
+	{
+		if (threadIdx.x == 0) saveSingleColorBlockDXT1(colors[0], result);
+		return;
+	}
+
 	ushort bestStart, bestEnd;
 	uint bestPermutation;
 
@@ -1033,6 +1111,11 @@ extern "C" void compressKernelDXT1(uint blockNum, uint * d_data, uint * d_result
 	compressDXT1<<<blockNum, NUM_THREADS>>>(d_bitmaps, d_data, (uint2 *)d_result);
 }
 
+extern "C" void compressKernelDXT1_Level4(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps)
+{
+	compressLevel4DXT1<<<blockNum, NUM_THREADS>>>(d_bitmaps, d_data, (uint2 *)d_result);
+}
+
 extern "C" void compressWeightedKernelDXT1(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps)
 {
 	compressWeightedDXT1<<<blockNum, NUM_THREADS>>>(d_bitmaps, d_data, (uint2 *)d_result);

src/nvtt/cuda/CudaCompressDXT.cpp

@@ -30,13 +30,14 @@
 #include <nvtt/CompressionOptions.h>
 #include <nvtt/OutputOptions.h>
 #include <nvtt/QuickCompressDXT.h>
+#include <nvtt/OptimalCompressDXT.h>
 
 #include "CudaCompressDXT.h"
 #include "CudaUtils.h"
 
 
 #if defined HAVE_CUDA
-#include <cuda_runtime.h>
+#include <cuda_runtime_api.h>
 #endif
 
 #include <time.h>
@@ -52,6 +53,7 @@ using namespace nvtt;
 
 extern "C" void setupCompressKernel(const float weights[3]);
 extern "C" void compressKernelDXT1(uint blockNum, uint * d_data, uint * d_result, uint * d_bitmaps);
+extern "C" void compressKernelDXT1_Level4(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);
 
 #include "Bitmaps.h"	// @@ Rename to BitmapTable.h
@@ -118,20 +120,25 @@ bool CudaCompressor::isValid() const
 
 // @@ This code is very repetitive and needs to be cleaned up.
 
+void CudaCompressor::setImage(const Image * image, nvtt::AlphaMode alphaMode)
+{
+	m_image = image;
+	m_alphaMode = alphaMode;
+}
 
 /// Compress image using CUDA.
-void CudaCompressor::compressDXT1(const Image * image, const OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions)
+void CudaCompressor::compressDXT1(const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions)
 {
 	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;
+	const uint w = (m_image->width() + 3) / 4;
+	const uint h = (m_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!
+	convertToBlockLinear(m_image, blockLinearImage);	// @@ Do this in parallel with the GPU, or in the GPU!
 
 	const uint blockNum = w * h;
 	const uint compressedSize = blockNum * 8;
@@ -176,7 +183,7 @@ void CudaCompressor::compressDXT1(const Image * image, const OutputOptions::Priv
 	}
 
 	clock_t end = clock();
-	printf("\rCUDA time taken: %.3f seconds\n", float(end-start) / CLOCKS_PER_SEC);
+	//printf("\rCUDA time taken: %.3f seconds\n", float(end-start) / CLOCKS_PER_SEC);
 
 	free(blockLinearImage);
 
@@ -190,18 +197,18 @@ void CudaCompressor::compressDXT1(const Image * image, const OutputOptions::Priv
 
 
 /// Compress image using CUDA.
-void CudaCompressor::compressDXT3(const Image * image, const OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions)
+void CudaCompressor::compressDXT3(const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions)
 {
 	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;
+	const uint w = (m_image->width() + 3) / 4;
+	const uint h = (m_image->height() + 3) / 4;
 
 	uint imageSize = w * h * 16 * sizeof(Color32);
     uint * blockLinearImage = (uint *) malloc(imageSize);
-	convertToBlockLinear(image, blockLinearImage);
+	convertToBlockLinear(m_image, blockLinearImage);
 
 	const uint blockNum = w * h;
 	const uint compressedSize = blockNum * 8;
@@ -221,13 +228,20 @@ void CudaCompressor::compressDXT3(const Image * image, const OutputOptions::Priv
 	    cudaMemcpy(m_data, blockLinearImage + bn * 16, count * 64, cudaMemcpyHostToDevice);
 
 		// Launch kernel.
-		compressWeightedKernelDXT1(count, m_data, m_result, m_bitmapTable);
+		if (m_alphaMode == AlphaMode_Transparency)
+		{
+			compressWeightedKernelDXT1(count, m_data, m_result, m_bitmapTable);
+		}
+		else
+		{
+			compressKernelDXT1_Level4(count, m_data, m_result, m_bitmapTable);
+		}
 
 		// Compress alpha in parallel with the GPU.
 		for (uint i = 0; i < count; i++)
 		{
 			ColorBlock rgba(blockLinearImage + (bn + i) * 16);
-			QuickCompress::compressDXT3A(rgba, alphaBlocks + i);
+			OptimalCompress::compressDXT3A(rgba, alphaBlocks + i);
 		}
 
 		// Check for errors.
@@ -259,7 +273,7 @@ void CudaCompressor::compressDXT3(const Image * image, const OutputOptions::Priv
 	}
 
 	clock_t end = clock();
-	printf("\rCUDA time taken: %.3f seconds\n", float(end-start) / CLOCKS_PER_SEC);
+	//printf("\rCUDA time taken: %.3f seconds\n", float(end-start) / CLOCKS_PER_SEC);
 
 	free(alphaBlocks);
 	free(blockLinearImage);
@@ -274,18 +288,18 @@ void CudaCompressor::compressDXT3(const Image * image, const OutputOptions::Priv
 
 
 /// Compress image using CUDA.
-void CudaCompressor::compressDXT5(const Image * image, const OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions)
+void CudaCompressor::compressDXT5(const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions)
 {
 	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;
+	const uint w = (m_image->width() + 3) / 4;
+	const uint h = (m_image->height() + 3) / 4;
 
 	uint imageSize = w * h * 16 * sizeof(Color32);
     uint * blockLinearImage = (uint *) malloc(imageSize);
-	convertToBlockLinear(image, blockLinearImage);
+	convertToBlockLinear(m_image, blockLinearImage);
 
 	const uint blockNum = w * h;
 	const uint compressedSize = blockNum * 8;
@@ -305,7 +319,14 @@ void CudaCompressor::compressDXT5(const Image * image, const OutputOptions::Priv
 	    cudaMemcpy(m_data, blockLinearImage + bn * 16, count * 64, cudaMemcpyHostToDevice);
 
 		// Launch kernel.
-		compressWeightedKernelDXT1(count, m_data, m_result, m_bitmapTable);
+		if (m_alphaMode == AlphaMode_Transparency)
+		{
+			compressWeightedKernelDXT1(count, m_data, m_result, m_bitmapTable);
+		}
+		else
+		{
+			compressKernelDXT1_Level4(count, m_data, m_result, m_bitmapTable);
+		}
 
 		// Compress alpha in parallel with the GPU.
 		for (uint i = 0; i < count; i++)
@@ -343,7 +364,7 @@ void CudaCompressor::compressDXT5(const Image * image, const OutputOptions::Priv
 	}
 
 	clock_t end = clock();
-	printf("\rCUDA time taken: %.3f seconds\n", float(end-start) / CLOCKS_PER_SEC);
+	//printf("\rCUDA time taken: %.3f seconds\n", float(end-start) / CLOCKS_PER_SEC);
 
 	free(alphaBlocks);
 	free(blockLinearImage);
@@ -357,185 +378,3 @@ void CudaCompressor::compressDXT5(const Image * image, const OutputOptions::Priv
 }
 
 
-
-#if 0
-
-class Task
-{
-public:
-	explicit Task(uint numBlocks) : blockMaxCount(numBlocks), blockCount(0)
-	{
-		// System memory allocations.
-		blockLinearImage = new uint[blockMaxCount * 16];
-		xrefs = new uint[blockMaxCount * 16];
-		
-		// Device memory allocations.
-		cudaMalloc((void**) &d_blockLinearImage, blockMaxCount * 16 * sizeof(uint));
-		cudaMalloc((void**) &d_compressedImage, blockMaxCount * 8U);
-		
-		// @@ Check for allocation errors.
-	}
-	
-	~Task()
-	{
-		delete [] blockLinearImage;
-		delete [] xrefs;
-		
-		cudaFree(d_blockLinearImage);
-		cudaFree(d_compressedImage);
-	}
-	
-	
-	
-	void addColorBlock(const ColorBlock & rgba)
-	{
-		nvDebugCheck(!isFull());
-		
-		// @@ Count unique colors?
-		/*
-		// Convert colors to vectors.
-		Array<Vector3> pointArray(16);
-		
-		for(int i = 0; i < 16; i++) {
-			const Color32 color = rgba.color(i);
-			pointArray.append(Vector3(color.r, color.g, color.b));
-		}
-		
-		// Find best fit line.
-		const Vector3 axis = Fit::bestLine(pointArray).direction();
-		
-		// Project points to axis.
-		float dps[16];
-		uint * order = &xrefs[blockCount * 16];
-		
-		for (uint i = 0; i < 16; ++i)
-		{
-			dps[i] = dot(pointArray[i], axis);
-			order[i] = i;
-		}
-		
-		// Sort them.
-		for (uint i = 0; i < 16; ++i)
-		{
-			for (uint j = i; j > 0 && dps[j] < dps[j - 1]; --j)
-			{
-				swap(dps[j], dps[j - 1]);
-				swap(order[j], order[j - 1]);
-			}
-		}
-		*/
-		// Write sorted colors to blockLinearImage.
-		for(uint i = 0; i < 16; ++i)
-		{
-		//	blockLinearImage[blockCount * 16 + i] = rgba.color(order[i]);
-			blockLinearImage[blockCount * 16 + i] = rgba.color(i);
-		}
-		
-		++blockCount;
-	}
-	
-	bool isFull()
-	{
-		nvDebugCheck(blockCount <= blockMaxCount);
-		return blockCount == blockMaxCount;
-	}
-	
-	void flush(const OutputOptions::Private & outputOptions)
-	{
-		if (blockCount == 0)
-		{
-			// Nothing to do.
-			return;
-		}
-		
-		// Copy input color blocks.
-		cudaMemcpy(d_blockLinearImage, blockLinearImage, blockCount * 64, cudaMemcpyHostToDevice);
-		
-		// Launch kernel.
-		compressKernelDXT1(blockCount, d_blockLinearImage, d_compressedImage, d_bitmaps);
-		
-		// 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.
-		uint * compressedImage = blockLinearImage;
-		cudaMemcpy(compressedImage, d_compressedImage, blockCount * 8, cudaMemcpyDeviceToHost);
-		
-		// @@ Sort block indices.
-		
-		// Output result.
-		if (outputOptions.outputHandler != NULL)
-		{
-		//	outputOptions.outputHandler->writeData(compressedImage, blockCount * 8);
-		}
-
-		blockCount = 0;
-	}
-	
-private:
-	
-	const uint blockMaxCount;
-	uint blockCount;
-	
-	uint * blockLinearImage;
-	uint * xrefs;
-	
-	uint * d_blockLinearImage;
-	uint * d_compressedImage;
-	
-};
-
-
-void nv::cudaCompressDXT1_2(const Image * image, const OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions)
-{
-#if defined HAVE_CUDA	
-	const uint w = image->width();
-	const uint h = image->height();
-	
-	const uint blockNum = ((w + 3) / 4) * ((h + 3) / 4);
-	const uint blockMax = 32768; // 49152, 65535
-		
-	setupCompressKernelDXT1(compressionOptions.colorWeight.ptr());
-
-	ColorBlock rgba;
-	Task task(min(blockNum, blockMax));
-
-	clock_t start = clock();
-
-	for (uint y = 0; y < h; y += 4) {
-		for (uint x = 0; x < w; x += 4) {
-			
-			rgba.init(image, x, y);
-			
-			task.addColorBlock(rgba);
-			
-			if (task.isFull())
-			{
-				task.flush(outputOptions);
-			}
-		}
-	}
-	
-	task.flush(outputOptions);
-
-	clock_t end = clock();
-	printf("\rCUDA time taken: %.3f seconds\n", float(end-start) / CLOCKS_PER_SEC);
-
-#else
-	if (outputOptions.errorHandler != NULL)
-	{
-		outputOptions.errorHandler->error(Error_CudaError);
-	}
-#endif
-}
-
-#endif // 0

src/nvtt/cuda/CudaCompressDXT.h

@@ -39,15 +39,20 @@ namespace nv
 
 		bool isValid() const;
 
-		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 setImage(const Image * image, nvtt::AlphaMode alphaMode);
+
+		void compressDXT1(const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
+		void compressDXT3(const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
+		void compressDXT5(const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
 
 	private:
 
 		uint * m_bitmapTable;
 		uint * m_data;
 		uint * m_result;
+		
+		const Image * m_image;
+		nvtt::AlphaMode m_alphaMode;
 	};
 
 } // nv namespace

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

@@ -22,15 +22,18 @@
 // OTHER DEALINGS IN THE SOFTWARE.
 
 #include <nvcore/Debug.h>
+#include <nvcore/Library.h>
 #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
@@ -52,31 +55,93 @@ static bool isWow32()
 {
 	LPFN_ISWOW64PROCESS fnIsWow64Process = (LPFN_ISWOW64PROCESS)GetProcAddress(GetModuleHandle("kernel32"), "IsWow64Process");
 
-    BOOL bIsWow64 = FALSE;
+	BOOL bIsWow64 = FALSE;
  
-    if (NULL != fnIsWow64Process)
-    {
-        if (!fnIsWow64Process(GetCurrentProcess(), &bIsWow64))
-        {
+	if (NULL != fnIsWow64Process)
+	{
+		if (!fnIsWow64Process(GetCurrentProcess(), &bIsWow64))
+		{
 			// Assume 32 bits.
-            return true;
-        }
-    }
+			return true;
+		}
+	}
 
-    return !bIsWow64;
+	return !bIsWow64;
 }
 
 #endif
+*/
+
+
+static bool isCudaDriverAvailable(int version)
+{
+#if defined HAVE_CUDA
+#if NV_OS_WIN32
+	Library nvcuda("nvcuda.dll");
+#else
+	Library nvcuda(NV_LIBRARY_NAME(cuda));
+#endif
+	
+	if (!nvcuda.isValid())
+	{
+		nvDebug("*** CUDA driver not found.\n");
+		return false;
+	}
+	
+	if (version >= 2000)
+	{
+		void * address = nvcuda.bindSymbol("cuStreamCreate");
+		if (address == NULL) {
+			nvDebug("*** CUDA driver version < 2.0.\n");
+			return false;
+		}
+	}
+
+	if (version >= 2010)
+	{
+		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;
+}
 
 
 /// Determine if CUDA is available.
 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)
 	{
@@ -89,10 +154,10 @@ bool nv::cuda::isHardwarePresent()
 		{
 			return false;
 		}
-
-		// @@ Make sure that warp size == 32
 	}
 
+	// @@ Make sure that warp size == 32
+
 	return count > 0;
 #else
 	return false;
@@ -115,14 +180,60 @@ int nv::cuda::deviceCount()
 	return 0;
 }
 
+int nv::cuda::getFastestDevice()
+{
+	int max_gflops_device = 0;
+#if defined HAVE_CUDA
+	int max_gflops = 0;
+
+	const int device_count = deviceCount();
+	int current_device = 0;
+	while (current_device < device_count)
+	{
+		cudaDeviceProp device_properties;
+		cudaGetDeviceProperties(&device_properties, current_device);
+		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;
+			}
+		}
+		
+		current_device++;
+	}
+#endif
+	return max_gflops_device;
+}
+
+
 /// Activate the given devices.
 bool nv::cuda::setDevice(int i)
 {
 	nvCheck(i < deviceCount());
 #if defined HAVE_CUDA
 	cudaError_t result = cudaSetDevice(i);
+
+	if (result != cudaSuccess) {
+		nvDebug("*** CUDA Error: %s\n", cudaGetErrorString(result));
+	}
+
 	return result == cudaSuccess;
 #else
 	return false;
 #endif
 }
+
+void nv::cuda::exit()
+{
+#if defined HAVE_CUDA
+	cudaError_t result = cudaThreadExit();
+
+	if (result != cudaSuccess) {
+		nvDebug("*** CUDA Error: %s\n", cudaGetErrorString(result));
+	}
+#endif
+}

src/nvtt/cuda/CudaUtils.h

@@ -31,7 +31,9 @@ namespace nv
 	{
 		bool isHardwarePresent();
 		int deviceCount();
+		int getFastestDevice();
 		bool setDevice(int i);
+		void exit();
 	};
 	
 } // nv namespace

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/fastclusterfit.cpp

@@ -29,6 +29,8 @@
 #include "colourblock.h"
 #include <cfloat>
 
+#include "fastclusterlookup.inl"
+
 namespace squish {
 
 FastClusterFit::FastClusterFit()
@@ -97,91 +99,6 @@ void FastClusterFit::SetColourSet( ColourSet const* colours, int flags )
 }
 
 
-struct Precomp {
-	float alpha2_sum;
-	float beta2_sum;
-	float alphabeta_sum;
-	float factor;
-};
-
-static SQUISH_ALIGN_16 Precomp s_threeElement[153];
-static SQUISH_ALIGN_16 Precomp s_fourElement[969];
-
-void FastClusterFit::DoPrecomputation()
-{
-	int i = 0;
-	
-	// Three element clusters:
-	for( int c0 = 0; c0 <= 16; c0++)	// At least two clusters.
-	{
-		for( int c1 = 0; c1 <=  16-c0; c1++)
-		{
-			int c2 = 16 - c0 - c1;
-
-			/*if (c2 == 16) {
-				// a = b = x2 / 16
-				s_threeElement[i].alpha2_sum = 0;
-				s_threeElement[i].beta2_sum = 16;
-				s_threeElement[i].alphabeta_sum = -16;
-				s_threeElement[i].factor = 1.0f / 256.0f;
-			}
-			else if (c0 == 16) {
-				// a = b = x0 / 16
-				s_threeElement[i].alpha2_sum = 16;
-				s_threeElement[i].beta2_sum = 0;
-				s_threeElement[i].alphabeta_sum = -16;
-				s_threeElement[i].factor = 1.0f / 256.0f;
-			}
-			else*/ {
-				s_threeElement[i].alpha2_sum = c0 + c1 * 0.25f;
-				s_threeElement[i].beta2_sum = c2 + c1 * 0.25f;
-				s_threeElement[i].alphabeta_sum = c1 * 0.25f;
-				s_threeElement[i].factor = 1.0f / (s_threeElement[i].alpha2_sum * s_threeElement[i].beta2_sum - s_threeElement[i].alphabeta_sum * s_threeElement[i].alphabeta_sum);
-			}
-			
-			i++;
-		}
-	}
-	//printf("%d three cluster elements\n", i);
-	
-	// Four element clusters:
-	i = 0;
-	for( int c0 = 0; c0 <= 16; c0++)
-	{
-		for( int c1 = 0; c1 <=  16-c0; c1++)
-		{
-			for( int c2 = 0; c2 <=  16-c0-c1; c2++)
-			{
-				int c3 = 16 - c0 - c1 - c2;
-				
-				/*if (c3 == 16) {
-					// a = b = x3 / 16
-					s_fourElement[i].alpha2_sum = 16.0f;
-					s_fourElement[i].beta2_sum = 0.0f;
-					s_fourElement[i].alphabeta_sum = -16.0f;
-					s_fourElement[i].factor = 1.0f / 256.0f;					
-				}
-				else if (c0 == 16) {
-					// a = b = x0 / 16
-					s_fourElement[i].alpha2_sum = 0.0f;
-					s_fourElement[i].beta2_sum = 16.0f;
-					s_fourElement[i].alphabeta_sum = -16.0f;
-					s_fourElement[i].factor = 1.0f / 256.0f;					
-				}
-				else*/ {
-					s_fourElement[i].alpha2_sum = c0 + c1 * (4.0f/9.0f) + c2 * (1.0f/9.0f);
-					s_fourElement[i].beta2_sum = c3 + c2 * (4.0f/9.0f) + c1 * (1.0f/9.0f);
-					s_fourElement[i].alphabeta_sum = (c1 + c2) * (2.0f/9.0f);
-					s_fourElement[i].factor = 1.0f / (s_fourElement[i].alpha2_sum * s_fourElement[i].beta2_sum - s_fourElement[i].alphabeta_sum * s_fourElement[i].alphabeta_sum);
-				}
-
-				i++;
-			}
-		}
-	}
-	//printf("%d four cluster elements\n", i);
-}
-
 void FastClusterFit::SetMetric(float r, float g, float b)
 {
 #if SQUISH_USE_SIMD

src/nvtt/squish/fastclusterfit.h

@@ -44,8 +44,6 @@ public:
 	void SetMetric(float r, float g, float b);
 	float GetBestError() const;
 
-	static void DoPrecomputation();
-
 	// Make them public
 	virtual void Compress3( void* block );
 	virtual void Compress4( void* block );

src/nvtt/tools/compress.cpp

@@ -42,11 +42,11 @@ struct MyOutputHandler : public nvtt::OutputHandler
 	MyOutputHandler(const char * name) : total(0), progress(0), percentage(0), stream(new nv::StdOutputStream(name)) {}
 	virtual ~MyOutputHandler() { delete stream; }
 	
-	virtual void setTotal(int64 t)
+	void setTotal(int64 t)
 	{
 		total = t + 128;
 	}
-	virtual void setDisplayProgress(bool b)
+	void setDisplayProgress(bool b)
 	{
 		verbose = b;
 	}
@@ -373,7 +373,6 @@ int main(int argc, char *argv[])
 		inputOptions.setMipmapGeneration(false);
 	}
 
-
 	nvtt::CompressionOptions compressionOptions;
 	compressionOptions.setFormat(format);
 	if (fast)
@@ -409,6 +408,16 @@ int main(int argc, char *argv[])
 	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);
 

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

@@ -73,10 +73,12 @@ int main(int argc, char *argv[])
 
 	float scale = 0.5f;
 	float gamma = 2.2f;
-	nv::Filter * filter = NULL;
+	nv::AutoPtr<nv::Filter> filter;
 	nv::Path input;
 	nv::Path output;
 	
+	nv::FloatImage::WrapMode wrapMode = nv::FloatImage::WrapMode_Mirror;
+
 	// Parse arguments.
 	for (int i = 1; i < argc; i++)
 	{
@@ -108,9 +110,18 @@ int main(int argc, char *argv[])
 			else if (strcmp("lanczos", argv[i]) == 0) filter = new nv::LanczosFilter();
 			else if (strcmp("kaiser", argv[i]) == 0) {
 				filter = new nv::KaiserFilter(3);
-				((nv::KaiserFilter *)filter)->setParameters(4.0f, 1.0f);
+				((nv::KaiserFilter *)filter.ptr())->setParameters(4.0f, 1.0f);
 			}
 		}
+		else if (strcmp("-w", argv[i]) == 0)
+		{
+			if (i+1 == argc) break;
+			i++;
+
+			if (strcmp("mirror", argv[i]) == 0) wrapMode = nv::FloatImage::WrapMode_Mirror;
+			else if (strcmp("repeat", argv[i]) == 0) wrapMode = nv::FloatImage::WrapMode_Repeat;
+			else if (strcmp("clamp", argv[i]) == 0) wrapMode = nv::FloatImage::WrapMode_Clamp;
+		}
 		else if (argv[i][0] != '-')
 		{
 			input = argv[i];
@@ -140,6 +151,10 @@ int main(int argc, char *argv[])
 		printf("                * mitchell\n");
 		printf("                * lanczos\n");
 		printf("                * kaiser\n");
+		printf("  -w mode      One of the following: (default = 'mirror')\n");
+		printf("                * mirror\n");
+		printf("                * repeat\n");
+		printf("                * clamp\n");
 
 		return 1;
 	}
@@ -155,15 +170,14 @@ int main(int argc, char *argv[])
 	nv::FloatImage fimage(&image);
 	fimage.toLinear(0, 3, gamma);
 	
-	nv::AutoPtr<nv::FloatImage> fresult(fimage.downSample(*filter, uint(image.width() * scale), uint(image.height() * scale), nv::FloatImage::WrapMode_Mirror));
+	nv::AutoPtr<nv::FloatImage> fresult(fimage.resize(*filter, uint(image.width() * scale), uint(image.height() * scale), wrapMode));
 	
 	nv::AutoPtr<nv::Image> result(fresult->createImageGammaCorrect(gamma));
+	result->setFormat(nv::Image::Format_ARGB);
 
 	nv::StdOutputStream stream(output);
 	nv::ImageIO::saveTGA(stream, result.ptr());	// @@ Add generic save function. Add support for png too.
 	
-	delete filter;
-	
 	return 0;
 }