drewcassidy/nvidia-texture-tools
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Fix cuda compressor. Start work on new UI.

Browse Source
This commit is contained in:
castano
2010-10-18 08:47:50 +00:00
parent 5d80fb8219
commit 5d498d6824
7 changed files with 1887 additions and 1177 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1977
src/nvtt/cuda/CompressKernel.cu
View File

File diff suppressed because it is too large Load Diff

32
src/nvtt/cuda/CudaCompressorDXT.cpp
View File

@ -34,10 +34,10 @@
#include "nvtt/QuickCompressDXT.h"
#include "nvtt/OptimalCompressDXT.h"
#include <time.h>
#include <stdio.h>
#if defined HAVE_CUDA
#include <cuda_runtime_api.h>
@ -141,7 +141,7 @@ void CudaCompressor::compress(nvtt::InputFormat inputFormat, nvtt::AlphaMode alp
}
else
{
#pragma message(NV_FILE_LINE "FIXME: Floating point textures not really supported by CUDA compressors.")
#pragma message(NV_FILE_LINE "FIXME: Floating point textures not really supported by CUDA compressors.") // @@ What's missing???
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 32, 32, 32, cudaChannelFormatKindFloat);
cudaMallocArray(&d_image, &channelDesc, w, h);
@ -156,7 +156,7 @@ void CudaCompressor::compress(nvtt::InputFormat inputFormat, nvtt::AlphaMode alp
const uint blockNum = bw * bh;
const uint compressedSize = blockNum * bs;
void * h_result = malloc(min(blockNum, MAX_BLOCKS) * bs);
void * h_result = ::malloc(min(blockNum, MAX_BLOCKS) * bs);
setup(d_image, compressionOptions);
@ -164,18 +164,18 @@ void CudaCompressor::compress(nvtt::InputFormat inputFormat, nvtt::AlphaMode alp
// timer.start();
uint bn = 0;
while(bn != blockNum)
while (bn != blockNum)
{
uint count = min(blockNum - bn, MAX_BLOCKS);
compressBlocks(bn, count, w, h, alphaMode, compressionOptions, h_result);
compressBlocks(bn, count, bw, bh, alphaMode, compressionOptions, h_result);
// Check for errors.
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess)
{
//nvDebug("CUDA Error: %s\n", cudaGetErrorString(err));
outputOptions.error(Error_CudaError);
//nvDebug("CUDA Error: %s\n", cudaGetErrorString(err));
outputOptions.error(Error_CudaError);
}
// Output result.
@ -198,10 +198,10 @@ void CudaCompressorDXT1::setup(cudaArray * image, const nvtt::CompressionOptions
bindTextureToArray(image);
}
void CudaCompressorDXT1::compressBlocks(uint first, uint count, uint w, uint h, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output)
void CudaCompressorDXT1::compressBlocks(uint first, uint count, uint bw, uint bh, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output)
{
// Launch kernel.
compressKernelDXT1(first, count, w, m_ctx.result, m_ctx.bitmapTable);
compressKernelDXT1(first, count, bw, m_ctx.result, m_ctx.bitmapTable);
// Copy result to host.
cudaMemcpy(output, m_ctx.result, count * 8, cudaMemcpyDeviceToHost);
@ -214,10 +214,10 @@ void CudaCompressorDXT3::setup(cudaArray * image, const nvtt::CompressionOptions
bindTextureToArray(image);
}
void CudaCompressorDXT3::compressBlocks(uint first, uint count, uint w, uint h, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output)
void CudaCompressorDXT3::compressBlocks(uint first, uint count, uint bw, uint bh, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output)
{
// Launch kernel.
compressKernelDXT3(first, count, w, m_ctx.result, m_ctx.bitmapTable);
compressKernelDXT3(first, count, bw, m_ctx.result, m_ctx.bitmapTable);
// Copy result to host.
cudaMemcpy(output, m_ctx.result, count * 16, cudaMemcpyDeviceToHost);
@ -230,10 +230,10 @@ void CudaCompressorDXT5::setup(cudaArray * image, const nvtt::CompressionOptions
bindTextureToArray(image);
}
void CudaCompressorDXT5::compressBlocks(uint first, uint count, uint w, uint h, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output)
void CudaCompressorDXT5::compressBlocks(uint first, uint count, uint bw, uint bh, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output)
{
/*// Launch kernel.
compressKernelDXT5(first, count, w, m_ctx.result, m_ctx.bitmapTable);
compressKernelDXT5(first, count, bw, m_ctx.result, m_ctx.bitmapTable);
// Copy result to host.
cudaMemcpy(output, m_ctx.result, count * 16, cudaMemcpyDeviceToHost);*/
@ -241,7 +241,7 @@ void CudaCompressorDXT5::compressBlocks(uint first, uint count, uint w, uint h,
// Launch kernel.
if (alphaMode == AlphaMode_Transparency)
{
// compressWeightedKernelDXT1(first, count, w, m_ctx.result, m_ctx.bitmapTable);
// compressWeightedKernelDXT1(first, count, bw, m_ctx.result, m_ctx.bitmapTable);
}
else
{
@ -251,8 +251,8 @@ void CudaCompressorDXT5::compressBlocks(uint first, uint count, uint w, uint h,
// Compress alpha in parallel with the GPU.
for (uint i = 0; i < count; i++)
{
//ColorBlock rgba(blockLinearImage + (first + i) * 16);
//OptimalCompress::compressDXT3A(rgba, alphaBlocks + i);
//ColorBlock rgba(blockLinearImage + (first + i) * 16);
//OptimalCompress::compressDXT3A(rgba, alphaBlocks + i);
}
// Copy result to host.

112
src/nvtt/cuda/CudaCompressorDXT.h
View File

@ -31,78 +31,78 @@ struct cudaArray;
namespace nv
{
class CudaContext
{
public:
CudaContext();
~CudaContext();
class CudaContext
{
public:
CudaContext();
~CudaContext();
bool isValid() const;
bool isValid() const;
public:
// Device pointers.
uint * bitmapTable;
uint * bitmapTableCTX;
uint * data;
uint * result;
};
public:
// Device pointers.
uint * bitmapTable;
uint * bitmapTableCTX;
uint * data;
uint * result;
};
#if defined HAVE_CUDA
struct CudaCompressor : public CompressorInterface
{
CudaCompressor(CudaContext & ctx);
struct CudaCompressor : public CompressorInterface
{
CudaCompressor(CudaContext & ctx);
virtual void compress(nvtt::InputFormat inputFormat, nvtt::AlphaMode alphaMode, uint w, uint h, const void * data, const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
virtual void compress(nvtt::InputFormat inputFormat, nvtt::AlphaMode alphaMode, uint w, uint h, const void * data, const nvtt::CompressionOptions::Private & compressionOptions, const nvtt::OutputOptions::Private & outputOptions);
virtual void setup(cudaArray * image, const nvtt::CompressionOptions::Private & compressionOptions) = 0;
virtual void compressBlocks(uint first, uint count, uint w, uint h, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output) = 0;
virtual uint blockSize() const = 0;
virtual void setup(cudaArray * image, const nvtt::CompressionOptions::Private & compressionOptions) = 0;
virtual void compressBlocks(uint first, uint count, uint w, uint h, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output) = 0;
virtual uint blockSize() const = 0;
protected:
CudaContext & m_ctx;
};
protected:
CudaContext & m_ctx;
};
struct CudaCompressorDXT1 : public CudaCompressor
{
CudaCompressorDXT1(CudaContext & ctx) : CudaCompressor(ctx) {}
struct CudaCompressorDXT1 : public CudaCompressor
{
CudaCompressorDXT1(CudaContext & ctx) : CudaCompressor(ctx) {}
virtual void setup(cudaArray * image, const nvtt::CompressionOptions::Private & compressionOptions);
virtual void compressBlocks(uint first, uint count, uint w, uint h, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output);
virtual uint blockSize() const { return 8; };
};
virtual void setup(cudaArray * image, const nvtt::CompressionOptions::Private & compressionOptions);
virtual void compressBlocks(uint first, uint count, uint w, uint h, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output);
virtual uint blockSize() const { return 8; };
};
/*struct CudaCompressorDXT1n : public CudaCompressor
{
virtual void setup(const CompressionOptions::Private & compressionOptions);
virtual void compressBlocks(uint blockCount, const void * input, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output) = 0;
virtual uint blockSize() const { return 8; };
};*/
/*struct CudaCompressorDXT1n : public CudaCompressor
{
virtual void setup(const CompressionOptions::Private & compressionOptions);
virtual void compressBlocks(uint blockCount, const void * input, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output) = 0;
virtual uint blockSize() const { return 8; };
};*/
struct CudaCompressorDXT3 : public CudaCompressor
{
CudaCompressorDXT3(CudaContext & ctx) : CudaCompressor(ctx) {}
struct CudaCompressorDXT3 : public CudaCompressor
{
CudaCompressorDXT3(CudaContext & ctx) : CudaCompressor(ctx) {}
virtual void setup(cudaArray * image, const nvtt::CompressionOptions::Private & compressionOptions);
virtual void compressBlocks(uint first, uint count, uint w, uint h, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output);
virtual uint blockSize() const { return 16; };
};
virtual void setup(cudaArray * image, const nvtt::CompressionOptions::Private & compressionOptions);
virtual void compressBlocks(uint first, uint count, uint w, uint h, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output);
virtual uint blockSize() const { return 16; };
};
struct CudaCompressorDXT5 : public CudaCompressor
{
CudaCompressorDXT5(CudaContext & ctx) : CudaCompressor(ctx) {}
struct CudaCompressorDXT5 : public CudaCompressor
{
CudaCompressorDXT5(CudaContext & ctx) : CudaCompressor(ctx) {}
virtual void setup(cudaArray * image, const nvtt::CompressionOptions::Private & compressionOptions);
virtual void compressBlocks(uint first, uint count, uint w, uint h, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output);
virtual uint blockSize() const { return 16; };
};
virtual void setup(cudaArray * image, const nvtt::CompressionOptions::Private & compressionOptions);
virtual void compressBlocks(uint first, uint count, uint w, uint h, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output);
virtual uint blockSize() const { return 16; };
};
/*struct CudaCompressorCXT1 : public CudaCompressor
{
virtual void setup(const CompressionOptions::Private & compressionOptions);
virtual void compressBlocks(uint blockCount, const void * input, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output) = 0;
virtual uint blockSize() const { return 8; };
};*/
/*struct CudaCompressorCXT1 : public CudaCompressor
{
virtual void setup(const CompressionOptions::Private & compressionOptions);
virtual void compressBlocks(uint blockCount, const void * input, nvtt::AlphaMode alphaMode, const nvtt::CompressionOptions::Private & compressionOptions, void * output) = 0;
virtual uint blockSize() const { return 8; };
};*/
#endif // defined HAVE_CUDA