From bd3314f4af9d7bc43b0a68bf71323673772be53a Mon Sep 17 00:00:00 2001 From: castano Date: Thu, 27 Mar 2008 04:28:17 +0000 Subject: [PATCH] Add inputOptions argument to compressors, so that they can access alpha mode. --- src/nvtt/Compressor.cpp | 14 ++-- src/nvtt/Compressor.h | 2 +- src/nvtt/cuda/CudaCompressDXT.cpp | 103 ++++++++++++++++++++++++++++-- src/nvtt/cuda/CudaCompressDXT.h | 10 +-- 4 files changed, 111 insertions(+), 18 deletions(-) diff --git a/src/nvtt/Compressor.cpp b/src/nvtt/Compressor.cpp index 357630f..beab2c3 100644 --- a/src/nvtt/Compressor.cpp +++ b/src/nvtt/Compressor.cpp @@ -429,7 +429,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); @@ -689,7 +689,7 @@ void Compressor::Private::quantizeMipmap(Mipmap & mipmap, const CompressionOptio // 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(); @@ -725,7 +725,7 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio if (cudaEnabled) { nvDebugCheck(cudaSupported); - cuda->compressDXT1(image, outputOptions, compressionOptions); + cuda->compressDXT1(image, compressionOptions, outputOptions); } else { @@ -757,7 +757,7 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio if (cudaEnabled) { nvDebugCheck(cudaSupported); - cuda->compressDXT1n(image, outputOptions, compressionOptions); + cuda->compressDXT1n(image, compressionOptions, outputOptions); } else { @@ -775,7 +775,7 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio if (cudaEnabled) { nvDebugCheck(cudaSupported); - cuda->compressDXT3(image, outputOptions, compressionOptions); + cuda->compressDXT3(image, inputOptions, compressionOptions, outputOptions); } else { @@ -794,7 +794,7 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio if (cudaEnabled) { nvDebugCheck(cudaSupported); - cuda->compressDXT5(image, outputOptions, compressionOptions); + cuda->compressDXT5(image, inputOptions, compressionOptions, outputOptions); } else { @@ -826,7 +826,7 @@ bool Compressor::Private::compressMipmap(const Mipmap & mipmap, const Compressio if (cudaEnabled) { nvDebugCheck(cudaSupported); - cuda->compressCTX1(image, outputOptions, compressionOptions); + cuda->compressCTX1(image, compressionOptions, outputOptions); } else { diff --git a/src/nvtt/Compressor.h b/src/nvtt/Compressor.h index 6da7dd1..9b98bdb 100644 --- a/src/nvtt/Compressor.h +++ b/src/nvtt/Compressor.h @@ -61,7 +61,7 @@ namespace nvtt void premultiplyAlphaMipmap(Mipmap & mipmap, const InputOptions::Private & inputOptions) 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: diff --git a/src/nvtt/cuda/CudaCompressDXT.cpp b/src/nvtt/cuda/CudaCompressDXT.cpp index 65af159..6119756 100644 --- a/src/nvtt/cuda/CudaCompressDXT.cpp +++ b/src/nvtt/cuda/CudaCompressDXT.cpp @@ -121,9 +121,102 @@ bool CudaCompressor::isValid() const // @@ This code is very repetitive and needs to be cleaned up. +#if 0 + +struct CudaCompressionKernel +{ + virtual void setup(const CompressionOptions::Private & compressionOptions) + { + setupCompressKernel(compressionOptions.colorWeight.ptr()); + } + + virtual void setBitmapTable(); + + virtual void runDeviceCode(int count); + + virtual void runHostCode(int count); + +}; + +void CudaCompressor::compressKernel(CudaCompressionKernel * kernel) +{ + nvDebugCheck(cuda::isHardwarePresent()); +#if defined HAVE_CUDA + + // Image size in blocks. + const uint w = (image->width() + 3) / 4; + const uint h = (image->height() + 3) / 4; + + uint imageSize = w * h * 16 * sizeof(Color32); + uint * blockLinearImage = (uint *) malloc(imageSize); + convertToBlockLinear(image, blockLinearImage); // @@ Do this in parallel with the GPU, or in the GPU! + + const uint blockNum = w * h; + const uint compressedSize = blockNum * 8; + + clock_t start = clock(); + + kernel->setup(compressionOptions); + kernel->setBitmapTable(m_bitmapTable); + + // TODO: Add support for multiple GPUs. + uint bn = 0; + while(bn != blockNum) + { + uint count = min(blockNum - bn, MAX_BLOCKS); + + cudaMemcpy(m_data, blockLinearImage + bn * 16, count * 64, cudaMemcpyHostToDevice); + + kernel->runDeviceCode(count, m_data, m_result); + + kernel->runHostCode(count); + + // Check for errors. + cudaError_t err = cudaGetLastError(); + if (err != cudaSuccess) + { + nvDebug("CUDA Error: %s\n", cudaGetErrorString(err)); + + if (outputOptions.errorHandler != NULL) + { + outputOptions.errorHandler->error(Error_CudaError); + } + } + + // Copy result to host, overwrite swizzled image. + cudaMemcpy(blockLinearImage, m_result, count * 8, cudaMemcpyDeviceToHost); + + // Output result. + kernel->outputResult(outputOptions.outputHandler); + + if (outputOptions.outputHandler != NULL) + { + outputOptions.outputHandler->writeData(blockLinearImage, count * 8); + } + + bn += count; + } + + clock_t end = clock(); + printf("\rCUDA time taken: %.3f seconds\n", float(end-start) / CLOCKS_PER_SEC); + + free(blockLinearImage); + +#else + if (outputOptions.errorHandler != NULL) + { + outputOptions.errorHandler->error(Error_CudaError); + } +#endif +} + +#endif // 0 + + + /// Compress image using CUDA. -void CudaCompressor::compressDXT1(const Image * image, const OutputOptions::Private & outputOptions, const CompressionOptions::Private & compressionOptions) +void CudaCompressor::compressDXT1(const Image * image, const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions) { nvDebugCheck(cuda::isHardwarePresent()); #if defined HAVE_CUDA @@ -193,7 +286,7 @@ 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 Image * image, const InputOptions::Private & inputOptions, const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions) { nvDebugCheck(cuda::isHardwarePresent()); #if defined HAVE_CUDA @@ -277,7 +370,7 @@ 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 Image * image, const InputOptions::Private & inputOptions, const CompressionOptions::Private & compressionOptions, const OutputOptions::Private & outputOptions) { nvDebugCheck(cuda::isHardwarePresent()); #if defined HAVE_CUDA @@ -360,7 +453,7 @@ void CudaCompressor::compressDXT5(const Image * image, const OutputOptions::Priv } -void CudaCompressor::compressDXT1n(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions) +void CudaCompressor::compressDXT1n(const Image * image, const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions) { nvDebugCheck(cuda::isHardwarePresent()); #if defined HAVE_CUDA @@ -429,7 +522,7 @@ void CudaCompressor::compressDXT1n(const Image * image, const nvtt::OutputOption } -void CudaCompressor::compressCTX1(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions) +void CudaCompressor::compressCTX1(const Image * image, const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions) { nvDebugCheck(cuda::isHardwarePresent()); #if defined HAVE_CUDA diff --git a/src/nvtt/cuda/CudaCompressDXT.h b/src/nvtt/cuda/CudaCompressDXT.h index 2d1691d..fee5f10 100644 --- a/src/nvtt/cuda/CudaCompressDXT.h +++ b/src/nvtt/cuda/CudaCompressDXT.h @@ -39,11 +39,11 @@ 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 compressDXT1n(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions); - void compressCTX1(const Image * image, const nvtt::OutputOptions::Private & outputOptions, const nvtt::CompressionOptions::Private & compressionOptions); + void compressDXT1(const Image * image, const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions); + void compressDXT3(const Image * image, const nvtt::InputOptions::Private & inputOptions, const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions); + void compressDXT5(const Image * image, const nvtt::InputOptions::Private & inputOptions, const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions); + void compressDXT1n(const Image * image, const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions); + void compressCTX1(const Image * image, const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions); private: