Create 2.0.5 release.

Rename our squish version to nvsquish.

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,16 @@
+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.

VERSION

@@ -1 +1 @@
-2.0.3
+2.0.5

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,3,0
- PRODUCTVERSION 2,0,3,0
+ FILEVERSION 2,0,5,0
+ PRODUCTVERSION 2,0,5,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, 3, 0"
+            VALUE "FileVersion", "2, 0, 5, 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, 3, 0"
+            VALUE "ProductVersion", "2, 0, 5, 0"
         END
     END
     BLOCK "VarFileInfo"

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/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,10 @@ namespace
 
 #if defined(HAVE_EXECINFO_H) // NV_OS_LINUX
 
+	static bool nvHasStackTrace() {
+		return backtrace != NULL;
+	}
+
 	static void nvPrintStackTrace(void * trace[], int size, int start=0) {
 		char ** string_array = backtrace_symbols(trace, size);
 	
@@ -166,24 +177,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 +251,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 +397,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, 3);
+			}
 #		endif
 
 			// Exit cleanly.
@@ -422,9 +449,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 

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/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;
@@ -213,9 +213,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/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) | (3);	// major.minor.revision
+	this->reserved[10] = (2 << 16) | (0 << 8) | (5);	// major.minor.revision
 
 	this->pf.size = 32;
 	this->pf.flags = 0;

src/nvimage/Filter.cpp

@@ -244,7 +244,7 @@ SincFilter::SincFilter(float w) : Filter(w) {}
 
 float SincFilter::evaluate(float x) const
 {
-	return 0.0f;
+	return sincf(PI * x);
 }
 
 
@@ -541,12 +541,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;
@@ -577,6 +582,7 @@ PolyphaseKernel::PolyphaseKernel(const Filter & f, uint srcLength, uint dstLengt
 			m_data[i * m_windowSize + j] /= total;
 		}
 	}
+
 }
 
 PolyphaseKernel::~PolyphaseKernel()

src/nvimage/FloatImage.cpp

@@ -540,73 +540,18 @@ 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, 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
 

src/nvimage/FloatImage.h

@@ -63,7 +63,7 @@ public:
 	
 	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 * 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;
@@ -226,11 +226,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/Quantize.cpp

@@ -16,6 +16,7 @@ http://www.efg2.com/Lab/Library/ImageProcessing/DHALF.TXT
 
 #include <nvimage/Image.h>
 #include <nvimage/Quantize.h>
+#include <nvimage/PixelFormat.h>
 
 using namespace nv;
 
@@ -47,94 +48,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 +115,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 our desired size, and reconstruct.
+			pixel.r = PixelFormat::convert(pixel.r, 8, rsize);
+			pixel.r = PixelFormat::convert(pixel.r, rsize, 8);
 
-			// Convert to 4 bit using regular bit expansion.
-			pixel.a = (pixel.a & 0xF0) | ((pixel.a & 0xF0) >> 4);
+			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 +208,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

@@ -17,6 +17,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/nvtt/CMakeLists.txt

@@ -44,7 +44,8 @@ INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR})
 ADD_DEFINITIONS(-DNVTT_EXPORTS)
 
 IF(NVTT_SHARED)
-	ADD_LIBRARY(nvtt SHARED ${DXT_SRCS})
+	ADD_DEFINITIONS(-DNVTT_SHARED=1)
+	ADD_LIBRARY(nvtt SHARED ${NVTT_SRCS})
 ELSE(NVTT_SHARED)
 	ADD_LIBRARY(nvtt ${NVTT_SRCS})
 ENDIF(NVTT_SHARED)

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"
@@ -210,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())
@@ -236,6 +241,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())
@@ -570,7 +579,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));
 }
 
 
@@ -617,13 +626,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)
@@ -635,19 +637,50 @@ 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);
 	}
 }
 

src/nvtt/cuda/CompressKernel.cu

@@ -191,6 +191,9 @@ __device__ void loadColorBlock(const uint * image, float3 colors[16], float3 sum
 
 		*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__

src/nvtt/cuda/CudaUtils.cpp

@@ -22,6 +22,7 @@
 // OTHER DEALINGS IN THE SOFTWARE.
 
 #include <nvcore/Debug.h>
+#include <nvcore/Library.h>
 #include "CudaUtils.h"
 
 #if defined HAVE_CUDA
@@ -52,23 +53,52 @@ 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(uint version)
+{
+#if NV_OS_WIN32
+	Library nvcuda("nvcuda.dll");
+#else
+	Library nvcuda(NV_LIBRARY_NAME(cuda));
+#endif
+	
+	if (!nvcuda.isValid())
+	{
+		return false;
+	}
+	
+	if (version > 2000)
+	{
+		void * address = nvcuda.bindSymbol("cuStreamCreate");
+		if (address == NULL) return false;
+	}
+
+	if (version > 2010)
+	{
+		void * address = nvcuda.bindSymbol("cuLoadDataEx");
+		if (address == NULL) return false;
+	}
+	
+	return true;
+}
+
+
 /// Determine if CUDA is available.
 bool nv::cuda::isHardwarePresent()
 {
@@ -90,6 +120,12 @@ bool nv::cuda::isHardwarePresent()
 			return false;
 		}
 
+		// Make sure that CUDA driver matches CUDA runtime.
+		if (!isCudaDriverAvailable(CUDART_VERSION))
+		{
+			return false;
+		}
+
 		// @@ Make sure that warp size == 32
 	}
 
@@ -115,6 +151,35 @@ 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)
 {

src/nvtt/cuda/CudaUtils.h

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

src/nvtt/tools/compress.cpp

@@ -373,7 +373,6 @@ int main(int argc, char *argv[])
 		inputOptions.setMipmapGeneration(false);
 	}
 
-
 	nvtt::CompressionOptions compressionOptions;
 	compressionOptions.setFormat(format);
 	if (fast)
@@ -397,6 +396,21 @@ int main(int argc, char *argv[])
 		compressionOptions.setExternalCompressor(externalCompressor);
 	}
 
+	if (format == nvtt::Format_RGB)
+	{
+		compressionOptions.setQuantization(true, false, false);
+		//compressionOptions.setPixelFormat(16, 0xF000, 0x0F00, 0x00F0, 0x000F);
+		compressionOptions.setPixelFormat(16, 
+			0x0F00, 
+			0x00F0, 
+			0x000F, 
+			0xF000);
+		//	0x003F0000, 
+		//	0x00003F00, 
+		//	0x0000003F, 
+		//	0x3F000000);
+	}
+
 	
 	MyErrorHandler errorHandler;
 	MyOutputHandler outputHandler(output);

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("-f", 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;
 }