cuda_dxt.cu File Reference

CUDA implementation of DXT compression. More...

#include <cuda_runtime_api.h>
#include <stdio.h>
#include "cuda_dxt.h"

Include dependency graph for cuda_dxt.cu:

Go to the source code of this file.

Classes
struct	uvec2
struct	uvec3
struct	vec2
struct	vec3

Typedefs
typedef unsigned int	u32
typedef u32	UINT

Functions
static __device__ vec2	abs (const vec2 &v)
static __device__ float	clamp (const float &v, float min_val, float max_val)
static __device__ vec2	clamp (const vec2 &v, float min_val, float max_val)
static __device__ float	colorDistance (vec2 c0, vec2 c1)
static __device__ vec3	ConvertRGBToYCoCg (vec3 color)
int	cuda_rgb_to_dxt1 (const void *src, void *out, int size_x, int size_y, cuda_wrapper_stream_t stream)
int	cuda_rgb_to_dxt6 (const void *src, void *out, int size_x, int size_y, cuda_wrapper_stream_t stream)
int	cuda_yuv422_to_yuv444 (const void *src, void *out, int pix_count, cuda_wrapper_stream_t str)
int	cuda_yuv_to_dxt1 (const void *src, void *out, int size_x, int size_y, cuda_wrapper_stream_t stream)
int	cuda_yuv_to_dxt6 (const void *src, void *out, int size_x, int size_y, cuda_wrapper_stream_t stream)
static __device__ float	dot (const vec3 &a, const vec3 &b)
template<int DXT_TYPE>
__device__ void	dxt_encode (void *out, const int block_idx, float r[16], float g[16], float b[16])
	Encodes and saves the block.
template<>
__device__ void	dxt_encode< 1 > (void *out, const int block_idx, float r[16], float g[16], float b[16])
	Encodes the block into DXT1 format and saves it into output buffer.
template<>
__device__ void	dxt_encode< 6 > (void *out, const int block_idx, float r[16], float g[16], float b[16])
template<bool YUV_TO_RGB, bool VERTICAL_MIRRORING, int DXT_TYPE>
static __global__ void	dxt_kernel (const void *src, void *out, int size_x, int size_y)
template<bool YUV_TO_RGB, int DXT_TYPE>
static int	dxt_launch (const void *src, void *out, int sx, int sy, cudaStream_t str)
	Compute grid size and launch DXT kernel.
static __device__ UINT	EmitAlphaEndPointsYCoCgDXT5 (float mincol, float maxcol)
static __device__ uvec2	EmitAlphaIndicesYCoCgDXT5 (vec3 block[16], float minAlpha, float maxAlpha)
static __device__ UINT	EmitEndPointsYCoCgDXT5 (float &mincol_r, float &mincol_g, float &maxcol_r, float &maxcol_g, UINT scale)
static __device__ UINT	EmitIndicesYCoCgDXT5 (vec3 block[16], vec2 mincol, vec2 maxcol)
static __device__ u32	encode_endpoint (float &r, float &g, float &b)
	Encodes color palette endpoint into 565 code and adjusts input values.
static __device__ void	FindMinMaxColorsBox (vec3 block[16], vec3 &mincol, vec3 &maxcol)
static __device__ void	InsetCoCgBBox (vec2 &mincol, vec2 &maxcol)
static __device__ void	InsetYBBox (float &mincol, float &maxcol)
static __device__ vec3	lerp (const vec3 &a, const vec3 &b, const float q)
static __device__ vec2	max (const vec2 &a, const vec2 &b)
static __device__ vec3	max (const vec3 &a, const vec3 &b)
static __device__ vec2	min (const vec2 &a, const vec2 &b)
static __device__ vec3	min (const vec3 &a, const vec3 &b)
static __device__ vec3	round (const vec3 &v)
static __device__ UINT	ScaleYCoCg (vec2 minColor, vec2 maxColor)
static __device__ bool	SelectYCoCgDiagonal (const vec3 block[16], vec2 minColor, vec2 maxColor)
template<typename T >
static __device__ void	swap (T &a, T &b)
	Swaps two referenced values.
static __global__ void	yuv422_to_yuv444_kernel (const void *src, void *out, int pix_count)
static __device__ void	yuv_to_rgb (float &r, float &g, float &b)
	Transform YUV to RGB.

Variables
static __device__ const float	offset = 128.0 / 255.0

Detailed Description

CUDA implementation of DXT compression.

Author

Martin Jirman jirma.nosp@m.n@ce.nosp@m.snet..nosp@m.cz

Definition in file cuda_dxt.cu.

Typedef Documentation

u32

typedef unsigned int u32

Definition at line 12 of file cuda_dxt.cu.

UINT

typedef u32 UINT

Definition at line 13 of file cuda_dxt.cu.

Function Documentation

abs()

static __device__ vec2 abs ( const vec2 & v )

static

Definition at line 118 of file cuda_dxt.cu.

                                           {
    return vec2(fabsf(v.r), fabsf(v.g));
}

References vec2::g, and vec2::r.

Referenced by ldgm::ldgm(), downshift_with_dither(), main(), misc_test_color_coeff_range(), pbuf_process_stats(), AudioDriftFixer::scale_buffer_delta(), ScaleYCoCg(), set_audio_delay(), tx_update(), and AudioDriftFixer::update().

clamp() [1/2]

static __device__ float clamp ( const float & v, float min_val, float max_val )

static

Definition at line 114 of file cuda_dxt.cu.

                                                                             {
    return min(max_val, max(min_val, v));
}

References max, and min.

clamp() [2/2]

static __device__ vec2 clamp ( const vec2 & v, float min_val, float max_val )

static

Definition at line 110 of file cuda_dxt.cu.

                                                                           {
    return min(vec2(max_val), max(vec2(min_val), v));
}

References max, and min.

Referenced by fill_full_to_limited_lut(), fill_limited_to_full_lut(), InsetCoCgBBox(), and InsetYBBox().

colorDistance()

static __device__ float colorDistance ( vec2 c0, vec2 c1 )

static

Definition at line 154 of file cuda_dxt.cu.

{
    return dot(c0-c1, c0-c1);
}

References dot().

Referenced by EmitIndicesYCoCgDXT5().

Here is the call graph for this function:

ConvertRGBToYCoCg()

static __device__ vec3 ConvertRGBToYCoCg ( vec3 color )

static

Definition at line 141 of file cuda_dxt.cu.

{
    float Y = (color.r + 2.0 * color.g + color.b) * 0.25;
    float Co = ( ( 2.0 * color.r - 2.0 * color.b      ) * 0.25 + offset );
    float Cg = ( (    -color.r + 2.0 * color.g - color.b) * 0.25 + offset );
 
    return vec3(Y, Co, Cg);
}

References vec3::b, vec2::g, offset, and vec2::r.

Referenced by dxt_encode< 6 >().

cuda_rgb_to_dxt1()

int cuda_rgb_to_dxt1	(	const void *	src,
		void *	out,
		int	size_x,
		int	size_y,
		cuda_wrapper_stream_t	stream )

CUDA DXT1 compression (only RGB without alpha).

Parameters

src	Pointer to top-left source pixel in device-memory buffer. 8bit RGB samples are expected (no alpha and no padding). (Pointer must be aligned to multiples of 16 bytes.)
out	Pointer to output buffer in device memory. (Must be aligned to multiples of 8 bytes.)
size_x	Width of the input image (must be divisible by 4).
size_y	Height of the input image (must be divisible by 4).
stream	CUDA stream to run in, or 0 for default stream.

Returns

0 if OK, nonzero if failed.

Definition at line 782 of file cuda_dxt.cu.

                                                                      {
    return dxt_launch<false, 1>(src, out, size_x, size_y, (cudaStream_t) stream);
}

References dxt_launch().

Referenced by main(), and worker_thread().

Here is the call graph for this function:

cuda_rgb_to_dxt6()

int cuda_rgb_to_dxt6	(	const void *	src,
		void *	out,
		int	size_x,
		int	size_y,
		cuda_wrapper_stream_t	stream )

CUDA DXT6 (DXT5-YcOcG) compression (only RGB without alpha).

Parameters

src	Pointer to top-left source pixel in device-memory buffer. 8bit RGB samples are expected (no alpha and no padding). (Pointer must be aligned to multiples of 16 bytes.)
out	Pointer to output buffer in device memory. (Must be aligned to multiples of 8 bytes.)
size_x	Width of the input image (must be divisible by 4).
size_y	Height of the input image (must be divisible by 4). (Input is read bottom up if negative)
stream	CUDA stream to run in, or 0 for default stream.

Returns

0 if OK, nonzero if failed.

Definition at line 816 of file cuda_dxt.cu.

                                                                      {
    return dxt_launch<false, 6>(src, out, size_x, size_y, (cudaStream_t) stream);
}

References dxt_launch().

Referenced by main(), and worker_thread().

Here is the call graph for this function:

cuda_yuv422_to_yuv444()

int cuda_yuv422_to_yuv444	(	const void *	src,
		void *	out,
		int	pix_count,
		cuda_wrapper_stream_t	str )

Definition at line 762 of file cuda_dxt.cu.

                                                          {
    // grid and threadblock sizes
    const dim3 tsiz(64, 1);
    int thread_count = pix_count / 4; // we process block of 4 pixels
    const dim3 gsiz((thread_count + tsiz.x - 1) / tsiz.x, 1);
    yuv422_to_yuv444_kernel<<<gsiz, tsiz, 0, (cudaStream_t) str>>>(src, out, pix_count);
    return cudaSuccess != cudaStreamSynchronize((cudaStream_t) str) ? -3 : 0;
}

References cudaSuccess, and yuv422_to_yuv444_kernel().

Here is the call graph for this function:

cuda_yuv_to_dxt1()

int cuda_yuv_to_dxt1	(	const void *	src,
		void *	out,
		int	size_x,
		int	size_y,
		cuda_wrapper_stream_t	stream )

CUDA DXT1 compression (only RGB without alpha). Converts input from YUV to RGB color space.

Parameters

src	Pointer to top-left source pixel in device-memory buffer. 8bit RGB samples are expected (no alpha and no padding). (Pointer must be aligned to multiples of 16 bytes.)
out	Pointer to output buffer in device memory. (Must be aligned to multiples of 8 bytes.)
size_x	Width of the input image (must be divisible by 4).
size_y	Height of the input image (must be divisible by 4).
stream	CUDA stream to run in, or 0 for default stream.

Returns

0 if OK, nonzero if failed.

Definition at line 799 of file cuda_dxt.cu.

                                                                      {
    return dxt_launch<true, 1>(src, out, size_x, size_y, (cudaStream_t) stream);
}

References dxt_launch().

Here is the call graph for this function:

cuda_yuv_to_dxt6()

int cuda_yuv_to_dxt6	(	const void *	src,
		void *	out,
		int	size_x,
		int	size_y,
		cuda_wrapper_stream_t	stream )

Definition at line 821 of file cuda_dxt.cu.

                                                                      {
    return dxt_launch<true, 6>(src, out, size_x, size_y, (cudaStream_t) stream);
}

References dxt_launch().

Here is the call graph for this function:

dot()

static __device__ float dot ( const vec3 & a, const vec3 & b )

static

Definition at line 106 of file cuda_dxt.cu.

                                                            {
    return a.r * b.r + a.g * b.g + a.b * b.b;
}

References vec3::b, vec2::g, and vec2::r.

Referenced by colorDistance().

dxt_encode()

template<int DXT_TYPE>

__device__ void dxt_encode	(	void *	out,
		const int	block_idx,
		float	r[16],
		float	g[16],
		float	b[16] )

Encodes and saves the block.

Referenced by dxt_kernel().

dxt_encode< 1 >()

template<>

__device__ void dxt_encode< 1 >	(	void *	out,
		const int	block_idx,
		float	r[16],
		float	g[16],
		float	b[16] )

Encodes the block into DXT1 format and saves it into output buffer.

Definition at line 513 of file cuda_dxt.cu.

                                                                            {
    // find min and max sample values for each component
    float mincol_r = r[0];
    float mincol_g = g[0];
    float mincol_b = b[0];
    float maxcol_r = r[0];
    float maxcol_g = g[0];
    float maxcol_b = b[0];
    for(int i = 1; i < 16; i++) {
        mincol_r = min(mincol_r, r[i]);
        mincol_g = min(mincol_g, g[i]);
        mincol_b = min(mincol_b, b[i]);
        maxcol_r = max(maxcol_r, r[i]);
        maxcol_g = max(maxcol_g, g[i]);
        maxcol_b = max(maxcol_b, b[i]);
    }
 
    // inset the bounding box
    const float inset_r = (maxcol_r - mincol_r) * 0.0625f;
    const float inset_g = (maxcol_g - mincol_g) * 0.0625f;
    const float inset_b = (maxcol_b - mincol_b) * 0.0625f;
    mincol_r += inset_r;
    mincol_g += inset_g;
    mincol_b += inset_b;
    maxcol_r -= inset_r;
    maxcol_g -= inset_g;
    maxcol_b -= inset_b;
 
    // select diagonal
    const float center_r = (mincol_r + maxcol_r) * 0.5f;
    const float center_g = (mincol_g + maxcol_g) * 0.5f;
    const float center_b = (mincol_b + maxcol_b) * 0.5f;
    float cov_x = 0.0f;
    float cov_y = 0.0f;
    for(int i = 0; i < 16; i++) {
        const float dir_r = r[i] - center_r;
        const float dir_g = g[i] - center_g;
        const float dir_b = b[i] - center_b;
        cov_x += dir_r * dir_b;
        cov_y += dir_g * dir_b;
    }
    if(cov_x < 0.0f) {
        swap(maxcol_r, mincol_r);
    }
    if(cov_y < 0.0f) {
        swap(maxcol_g, mincol_g);
    }
 
    // encode both endpoints into 565 color format
    const u32 max_code = encode_endpoint(maxcol_r, maxcol_g, maxcol_b);
    const u32 min_code = encode_endpoint(mincol_r, mincol_g, mincol_b);
 
    // swap palette end colors if 'max' code is less than 'min' color code
    // (Palette color #3 would otherwise be interpreted as 'transparent'.)
    const bool swap_end_colors = max_code < min_code;
    
    // encode the palette into 32 bits (Only 2 end colors are stored.)
    const u32 palette_code = swap_end_colors ?
            min_code + (max_code << 16): max_code + (min_code << 16);
    
    // pack palette color indices (if both endpoint colors are not equal)
    u32 indices = 0;
    if(max_code != min_code) {
        // project each color to line maxcol-mincol, represent it as
        // "mincol + t * (maxcol - mincol)" and then use 't' to find closest 
        // palette color index.
        const float dir_r = mincol_r - maxcol_r;
        const float dir_g = mincol_g - maxcol_g;
        const float dir_b = mincol_b - maxcol_b;
        const float dir_sqr_len = dir_r * dir_r + dir_g * dir_g + dir_b * dir_b;
        const float dir_inv_sqr_len = __fdividef(1.0f, dir_sqr_len);
        const float t_r = dir_r * dir_inv_sqr_len;
        const float t_g = dir_g * dir_inv_sqr_len;
        const float t_b = dir_b * dir_inv_sqr_len;
        const float t_bias = t_r * maxcol_r + t_g * maxcol_g + t_b * maxcol_b;
        
        // for each pixel color:
        for(int i = 0; i < 16; i++) {
            // get 't' for the color
            const float col_t = r[i] * t_r + g[i] * t_g + b[i] * t_b - t_bias;
            
            // scale the range of the 't' to [0..3] and convert to integer
            // to get the index of palette color
            const u32 col_idx = (u32)(3.0f * __saturatef(col_t) + 0.5f);
            
            // pack the color palette index with others
            indices += col_idx << (i * 2);
        }
    }
    
    // possibly invert indices if end colors must be swapped
    if(swap_end_colors) {
        indices = ~indices;
    }
    
    // substitute all packed indices (each index is packed into two bits)
    // 00 -> 00, 01 -> 10, 10 -> 11 and 11 -> 01
    const u32 lsbs = indices & 0x55555555;
    const u32 msbs = indices & 0xaaaaaaaa;
    indices = msbs ^ (2 * lsbs + (msbs >> 1));
 
    // compose and save output
    ((uint2*)out)[block_idx] = make_uint2(palette_code, indices);
}

References encode_endpoint(), max, min, and swap().

Here is the call graph for this function:

dxt_encode< 6 >()

template<>

__device__ void dxt_encode< 6 >	(	void *	out,
		const int	block_idx,
		float	r[16],
		float	g[16],
		float	b[16] )

Encodes the block into DXT6 format (DXT5-YcOcG) and saves it into output buffer

Definition at line 472 of file cuda_dxt.cu.

                                                                            {
     // Read block of data
    vec3 block[16];
    
    for(int i = 0; i < 16; i++) {
        block[i] = ConvertRGBToYCoCg(vec3(r[i], g[i], b[i]));
    }
 
    // Find min and max colors
    vec3 mincol, maxcol;
    FindMinMaxColorsBox(block, mincol, maxcol);
 
    if(SelectYCoCgDiagonal(block, mincol.gb(), maxcol.gb()))  {
        float tmp = maxcol.b;
        maxcol.b = mincol.b;
        mincol.b = tmp;
    }
 
    u32 scale = ScaleYCoCg(mincol.gb(), maxcol.gb());
//     printf("Scale: %u.\n", scale);
    // Output CoCg in DXT1 block.
    uint4 outp;
    outp.z = EmitEndPointsYCoCgDXT5(mincol.g, mincol.b, maxcol.g, maxcol.b, scale);
    outp.w = EmitIndicesYCoCgDXT5(block, mincol.gb(), maxcol.gb());
 
    InsetYBBox(mincol.r, maxcol.r);
 
    // Output Y in DXT5 alpha block.
    outp.x = EmitAlphaEndPointsYCoCgDXT5(mincol.r, maxcol.r);
 
    uvec2 indices = EmitAlphaIndicesYCoCgDXT5(block, mincol.r, maxcol.r);
    outp.x |= indices.r;
    outp.y = indices.g;
    
    ((uint4*)out)[block_idx] = outp;
    
}

References vec3::b, ConvertRGBToYCoCg(), EmitAlphaEndPointsYCoCgDXT5(), EmitAlphaIndicesYCoCgDXT5(), EmitEndPointsYCoCgDXT5(), EmitIndicesYCoCgDXT5(), FindMinMaxColorsBox(), uvec2::g, vec2::g, vec3::gb(), InsetYBBox(), uvec2::r, vec2::r, ScaleYCoCg(), and SelectYCoCgDiagonal().

Here is the call graph for this function:

dxt_kernel()

template<bool YUV_TO_RGB, bool VERTICAL_MIRRORING, int DXT_TYPE>

static __global__ void dxt_kernel ( const void * src, void * out, int size_x, int size_y )

static

DXT compression - each thread compresses one 4x4 DXT block. Alpha-color palette mode is not used (always emmits 4color palette code).

Definition at line 623 of file cuda_dxt.cu.

                                                                                        {
    // coordinates of this thread's 4x4 block
    const int block_idx_x = threadIdx.x + blockIdx.x * blockDim.x;
    const int block_idx_y = threadIdx.y + blockIdx.y * blockDim.y;
    
    // coordinates of block's top-left pixel
    const int block_x = block_idx_x * 4;
    const int block_y = block_idx_y * 4;
    
    // raster order index of the block
    const int block_idx = block_idx_x + (size_x >> 2) * block_idx_y;
    
    // skip if out of bounds
    if(block_y >= size_y || block_x >= size_x) {
        return;
    }
 
    // samples of 16 pixels
    float r[16];
    float g[16];
    float b[16];
 
    // load RGB samples for all 16 input pixels
    const int src_stride = (size_x >> 2) * 3;
    for(int y = 0; y < 4; y++) {
        // offset of loaded pixels in the buffer
        const int load_offset = y * 4;
        
        // pointer to source of this input row
        int row_idx = block_y + y;
        if(VERTICAL_MIRRORING) {
            row_idx = size_y - 1 - row_idx;
        }
        const uchar4 * const row_src = (uchar4*)src
                                     + src_stride * row_idx
                                     + block_idx_x * 3;
        
        // load all 4 3component pixels of the row
        const uchar4 p0 = row_src[0];
        const uchar4 p1 = row_src[1];
        const uchar4 p2 = row_src[2];
        
        // pixel #0
        r[load_offset + 0] = p0.x * 0.00392156862745f;
        g[load_offset + 0] = p0.y * 0.00392156862745f;
        b[load_offset + 0] = p0.z * 0.00392156862745f;
        
        // pixel #1
        r[load_offset + 1] = p0.w * 0.00392156862745f;
        g[load_offset + 1] = p1.x * 0.00392156862745f;
        b[load_offset + 1] = p1.y * 0.00392156862745f;
        
        // pixel #2
        r[load_offset + 2] = p1.z * 0.00392156862745f;
        g[load_offset + 2] = p1.w * 0.00392156862745f;
        b[load_offset + 2] = p2.x * 0.00392156862745f;
        
        // pixel #3
        r[load_offset + 3] = p2.y * 0.00392156862745f;
        g[load_offset + 3] = p2.z * 0.00392156862745f;
        b[load_offset + 3] = p2.w * 0.00392156862745f;
    }
    
    // transform colors from YUV to RGB if required
    if(YUV_TO_RGB) {
        for(int i = 0; i < 16; i++) {
            yuv_to_rgb(r[i], g[i], b[i]);
        }
    }
    
    // Select the right DXT type transform
    dxt_encode<DXT_TYPE>(out, block_idx, r, g, b);
}

References dxt_encode(), and yuv_to_rgb().

Referenced by dxt_launch().

Here is the call graph for this function:

dxt_launch()

template<bool YUV_TO_RGB, int DXT_TYPE>

static int dxt_launch ( const void * src, void * out, int sx, int sy, cudaStream_t str )

static

Compute grid size and launch DXT kernel.

Definition at line 736 of file cuda_dxt.cu.

                                                                                      {
    // vertical mirroring?
    bool mirrored = false;
    if(sy < 0) {
        mirrored = true;
        sy = -sy;
    }
    
    // check image size and alignment
    if((sx & 3) || (sy & 3) || (15 & (size_t)src) || (7 & (size_t)out)) {
        return -1;
    }
    
    // grid and threadblock sizes
    const dim3 tsiz(16, 16);
    const dim3 gsiz((sx + tsiz.x - 1) / tsiz.x, (sy + tsiz.y - 1) / tsiz.y);
    
    // launch kernel, sync and check the result
    if(mirrored) {
        dxt_kernel<YUV_TO_RGB, true, DXT_TYPE><<<gsiz, tsiz, 0, str>>>(src, out, sx, sy);
    } else {
        dxt_kernel<YUV_TO_RGB, false, DXT_TYPE><<<gsiz, tsiz, 0, str>>>(src, out, sx, sy);
    }
    return cudaSuccess != cudaStreamSynchronize(str) ? -3 : 0;
}

References cudaSuccess, and dxt_kernel().

Referenced by cuda_rgb_to_dxt1(), cuda_rgb_to_dxt6(), cuda_yuv_to_dxt1(), and cuda_yuv_to_dxt6().

Here is the call graph for this function:

EmitAlphaEndPointsYCoCgDXT5()

static __device__ UINT EmitAlphaEndPointsYCoCgDXT5 ( float mincol, float maxcol )

static

Definition at line 350 of file cuda_dxt.cu.

{
    uvec2 tmp = uvec2(roundf(mincol * 255.0), roundf(maxcol * 255.0));
    UINT c0 = tmp.r;
    UINT c1 = tmp.g;
 
    return (c0 << 8u) | c1;
}

References uvec2::g, and uvec2::r.

Referenced by dxt_encode< 6 >().

EmitAlphaIndicesYCoCgDXT5()

static __device__ uvec2 EmitAlphaIndicesYCoCgDXT5 ( vec3 block[16], float minAlpha, float maxAlpha )

static

Definition at line 360 of file cuda_dxt.cu.

{
    const float ALPHA_RANGE = 7.0;
 
    float mid = (maxAlpha - minAlpha) / (2.0 * ALPHA_RANGE);
 
    float ab1 = minAlpha + mid;
    float ab2 = (6.0 * maxAlpha + 1.0 * minAlpha) * (1.0 / ALPHA_RANGE) + mid;
    float ab3 = (5.0 * maxAlpha + 2.0 * minAlpha) * (1.0 / ALPHA_RANGE) + mid;
    float ab4 = (4.0 * maxAlpha + 3.0 * minAlpha) * (1.0 / ALPHA_RANGE) + mid;
    float ab5 = (3.0 * maxAlpha + 4.0 * minAlpha) * (1.0 / ALPHA_RANGE) + mid;
    float ab6 = (2.0 * maxAlpha + 5.0 * minAlpha) * (1.0 / ALPHA_RANGE) + mid;
    float ab7 = (1.0 * maxAlpha + 6.0 * minAlpha) * (1.0 / ALPHA_RANGE) + mid;
 
    uvec2 indices = uvec2(0, 0);
 
    UINT index = 1u;
    for ( int i = 0; i < 6; i++ ) {
        float a = block[i].r;
        index = 1u;
        index += (a <= ab1) ? 1u : 0u;
        index += (a <= ab2) ? 1u : 0u;
        index += (a <= ab3) ? 1u : 0u;
        index += (a <= ab4) ? 1u : 0u;
        index += (a <= ab5) ? 1u : 0u;
        index += (a <= ab6) ? 1u : 0u;
        index += (a <= ab7) ? 1u : 0u;
        index &= 7u;
        index ^= (2u > index) ? 1u : 0u;
        indices.r |= index << (3u * UINT(i) + 16u);
    }
 
    indices.g = index >> 1u;
 
    for ( int i = 6; i < 16; i++ ) {
        float a = block[i].r;
        index = 1u;
        index += (a <= ab1) ? 1u : 0u;
        index += (a <= ab2) ? 1u : 0u;
        index += (a <= ab3) ? 1u : 0u;
        index += (a <= ab4) ? 1u : 0u;
        index += (a <= ab5) ? 1u : 0u;
        index += (a <= ab6) ? 1u : 0u;
        index += (a <= ab7) ? 1u : 0u;
        index &= 7u;
        index ^= (2u > index) ? 1u : 0u;
        indices.g |= index << (3u * UINT(i) - 16u);
    }
 
    return indices;
}

References uvec2::g, uvec2::r, and vec2::r.

Referenced by dxt_encode< 6 >().

EmitEndPointsYCoCgDXT5()

static __device__ UINT EmitEndPointsYCoCgDXT5 ( float & mincol_r, float & mincol_g, float & maxcol_r, float & maxcol_g, UINT scale )

static

Definition at line 272 of file cuda_dxt.cu.

{
    vec2 mincol = vec2(mincol_r, mincol_g);
    vec2 maxcol = vec2(maxcol_r, maxcol_g);
    
    maxcol = (maxcol - offset) * float(scale) + float(offset);
    mincol = (mincol - offset) * float(scale) + float(offset);
 
    InsetCoCgBBox(mincol, maxcol);
 
    maxcol = round(maxcol * vec2(31, 63));
    mincol = round(mincol * vec2(31, 63));
 
    uvec2 imaxcol = uvec2(maxcol);
    uvec2 imincol = uvec2(mincol);
 
    uvec2 outp;
    outp.r = (imaxcol.r << 11u) | (imaxcol.g << 5u) | (scale - UINT(1));
    outp.g = (imincol.r << 11u) | (imincol.g << 5u) | (scale - UINT(1));
 
    imaxcol.r = (imaxcol.r << 3u) | (imaxcol.r >> 2u);
    imaxcol.g = (imaxcol.g << 2u) | (imaxcol.g >> 4u);
    imincol.r = (imincol.r << 3u) | (imincol.r >> 2u);
    imincol.g = (imincol.g << 2u) | (imincol.g >> 4u);
 
    maxcol = vec2(imaxcol) * (1.0 / 255.0);
    mincol = vec2(imincol) * (1.0 / 255.0);
 
    // Undo rescale.
    maxcol = (maxcol - offset) / float(scale) + float(offset);
    mincol = (mincol - offset) / float(scale) + float(offset);
    
    // distribute back
    mincol_r = mincol.r;
    mincol_g = mincol.g;
    maxcol_r = maxcol.r;
    maxcol_g = maxcol.g;
 
    return outp.r | (outp.g << 16u);
}

References uvec2::g, vec2::g, InsetCoCgBBox(), offset, uvec2::r, vec2::r, and round().

Referenced by dxt_encode< 6 >().

Here is the call graph for this function:

EmitIndicesYCoCgDXT5()

static __device__ UINT EmitIndicesYCoCgDXT5 ( vec3 block[16], vec2 mincol, vec2 maxcol )

static

Definition at line 315 of file cuda_dxt.cu.

{
    // Compute palette
    vec2 c[4];
    c[0] = maxcol;
    c[1] = mincol;
    c[2] = lerp(c[0], c[1], 1.0/3.0);
    c[3] = lerp(c[0], c[1], 2.0/3.0);
 
    // Compute indices
    UINT indices = 0u;
    for ( int i = 0; i < 16; i++ )
    {
        // find index of closest color
        float4 dist;
        dist.x = colorDistance(block[i].gb(), c[0]);
        dist.y = colorDistance(block[i].gb(), c[1]);
        dist.z = colorDistance(block[i].gb(), c[2]);
        dist.w = colorDistance(block[i].gb(), c[3]);
 
        uint4 b;
        b.x = dist.x > dist.w ? 1u : 0u;
        b.y = dist.y > dist.z ? 1u : 0u;
        b.z = dist.x > dist.z ? 1u : 0u;
        b.w = dist.y > dist.w ? 1u : 0u;
        UINT b4 = dist.z > dist.w ? 1u : 0u;
 
        UINT index = (b.x & b4) | (((b.y & b.z) | (b.x & b.w)) << 1u);
        indices |= index << (UINT(i) * 2u);
    }
 
    // Output indices
    return indices;
}

References colorDistance(), and lerp().

Referenced by dxt_encode< 6 >().

Here is the call graph for this function:

encode_endpoint()

static __device__ u32 encode_endpoint ( float & r, float & g, float & b )

static

Encodes color palette endpoint into 565 code and adjusts input values.

Definition at line 424 of file cuda_dxt.cu.

                                                                       {
    // clamp to range [0,1] and use full output range for each component
    r = rintf(__saturatef(r) * 31.0f);
    g = rintf(__saturatef(g) * 63.0f);  // 6 bits for green sample
    b = rintf(__saturatef(b) * 31.0f);
    
    // compose output 16bit code representing the endpoint color
    const u32 code = ((u32)r << 11) + ((u32)g << 5) + (u32)b;
 
    // convert all 3 endpoint component samples back to unit range
    r *= 0.0322580645161f;  // divide by 31
    g *= 0.015873015873f;   // divide by 63
    b *= 0.0322580645161f;  // divide by 31
    
    // return output 16bit code for the endpoint
    return code;
}

Referenced by dxt_encode< 1 >().

FindMinMaxColorsBox()

static __device__ void FindMinMaxColorsBox ( vec3 block[16], vec3 & mincol, vec3 & maxcol )

static

Definition at line 159 of file cuda_dxt.cu.

{
    mincol = block[0];
    maxcol = block[0];
    
    for ( int i = 1; i < 16; i++ ) {
        mincol = min(mincol, block[i]);
        maxcol = max(maxcol, block[i]);
    }
}

References max, and min.

Referenced by dxt_encode< 6 >().

InsetCoCgBBox()

static __device__ void InsetCoCgBBox ( vec2 & mincol, vec2 & maxcol )

static

Definition at line 182 of file cuda_dxt.cu.

{
    vec2 inset = (maxcol - mincol) / 16.0 - (8.0 / 255.0) / 16.0;
    mincol = clamp(mincol + inset, 0.0, 1.0);
    maxcol = clamp(maxcol - inset, 0.0, 1.0);
}

References clamp().

Referenced by EmitEndPointsYCoCgDXT5().

Here is the call graph for this function:

InsetYBBox()

static __device__ void InsetYBBox ( float & mincol, float & maxcol )

static

Definition at line 176 of file cuda_dxt.cu.

{
    float inset = (maxcol - mincol) / 32.0 - (16.0 / 255.0) / 32.0;
    mincol = clamp(mincol + inset, 0.0, 1.0);
    maxcol = clamp(maxcol - inset, 0.0, 1.0);
}

References clamp().

Referenced by dxt_encode< 6 >().

Here is the call graph for this function:

lerp()

static __device__ vec3 lerp ( const vec3 & a, const vec3 & b, const float q )

static

Definition at line 126 of file cuda_dxt.cu.

                                                                           {
    return a * (1.0f - q) + b * q;
}

Referenced by EmitIndicesYCoCgDXT5().

max() [1/2]

static __device__ vec2 max ( const vec2 & a, const vec2 & b )

static

Definition at line 102 of file cuda_dxt.cu.

                                                           {
    return vec2(max(a.r, b.r), max(a.g, b.g));
}

References vec2::g, max, and vec2::r.

max() [2/2]

static __device__ vec3 max ( const vec3 & a, const vec3 & b )

static

Definition at line 94 of file cuda_dxt.cu.

                                                           {
    return vec3(max(a.r, b.r), max(a.g, b.g), max(a.b, b.b));
}

References vec3::b, vec2::g, max, and vec2::r.

min() [1/2]

static __device__ vec2 min ( const vec2 & a, const vec2 & b )

static

Definition at line 98 of file cuda_dxt.cu.

                                                           {
    return vec2(min(a.r, b.r), min(a.g, b.g));
}

References vec2::g, min, and vec2::r.

min() [2/2]

static __device__ vec3 min ( const vec3 & a, const vec3 & b )

static

Definition at line 90 of file cuda_dxt.cu.

                                                           {
    return vec3(min(a.r, b.r), min(a.g, b.g), min(a.b, b.b));
}

References vec3::b, vec2::g, min, and vec2::r.

round()

static __device__ vec3 round ( const vec3 & v )

static

Definition at line 122 of file cuda_dxt.cu.

                                             {
    return vec3(roundf(v.r), roundf(v.g), roundf(v.b));
}

References vec3::b, vec2::g, and vec2::r.

Referenced by delay_init(), EmitEndPointsYCoCgDXT5(), format_interl_fps_hdr_row(), generate_audio_sine(), get_ebu_signal(), get_framerate_d(), get_framerate_n(), get_framerate_test_free(), get_refresh_rate(), gl_reconfigure_screen(), glfw_resize_window(), and DecklinkAudioSummary::set_buffer_average().

ScaleYCoCg()

static __device__ UINT ScaleYCoCg ( vec2 minColor, vec2 maxColor )

static

Definition at line 241 of file cuda_dxt.cu.

{
    vec2 m0 = abs(minColor - offset);
    vec2 m1 = abs(maxColor - offset);
 
    float m = max(max(m0.r, m0.g), max(m1.r, m1.g));
 
    const float s0 = 64.0 / 255.0;
    const float s1 = 32.0 / 255.0;
 
    UINT scale = 1u;
    if ( m < s0 )
        scale = 2u;
    if ( m < s1 )
        scale = 4u;
 
    return scale;
}

References abs(), vec2::g, max, offset, and vec2::r.

Referenced by dxt_encode< 6 >().

Here is the call graph for this function:

SelectYCoCgDiagonal()

static __device__ bool SelectYCoCgDiagonal ( const vec3 block[16], vec2 minColor, vec2 maxColor )

static

Definition at line 260 of file cuda_dxt.cu.

{
    vec2 mid = (maxColor + minColor) * 0.5;
 
    float cov = 0.0;
    for ( int i = 0; i < 16; i++ ) {
        vec2 t = block[i].gb() - mid;
        cov += t.r * t.g;
    }
    return cov < 0.0;
}

References vec2::g, vec3::gb(), and vec2::r.

Referenced by dxt_encode< 6 >().

Here is the call graph for this function:

swap()

template<typename T >

static __device__ void swap ( T & a, T & b )

static

Swaps two referenced values.

Definition at line 456 of file cuda_dxt.cu.

                                          {
    const T temp = a;
    a = b;
    b = temp;
}

Referenced by dxt_encode< 1 >(), gl_process_frames(), and vidcap_deltacast_dvi_grab().

yuv422_to_yuv444_kernel()

static __global__ void yuv422_to_yuv444_kernel ( const void * src, void * out, int pix_count )

static

Definition at line 697 of file cuda_dxt.cu.

                                                                                            {
    // coordinates of this thread
    const int block_idx_x = threadIdx.x + blockIdx.x * blockDim.x;
 
    // skip if out of bounds
    if(block_idx_x >= pix_count / 2) {
        return;
    }
 
    uchar4 *this_src = ((uchar4 *) src) + block_idx_x * 2;
    uchar4 *this_out = ((uchar4 *) out) + block_idx_x * 3;
 
    uchar4 pix12 = this_src[0];
    uchar4 pix34 = this_src[1];
 
    uchar4 out_pix[3];
    out_pix[0].x = pix12.y;
    out_pix[0].y = pix12.x;
    out_pix[0].z = pix12.z;
 
    out_pix[0].w = pix12.w;
    out_pix[1].x = pix12.x;
    out_pix[1].y = pix12.z;
 
    out_pix[1].z = pix34.y;
    out_pix[1].w = pix34.x;
    out_pix[2].x = pix34.z;
 
    out_pix[2].y = pix34.w;
    out_pix[2].z = pix34.x;
    out_pix[2].w = pix34.z;
 
    this_out[0] = out_pix[0];
    this_out[1] = out_pix[1];
    this_out[2] = out_pix[2];
}

Referenced by cuda_yuv422_to_yuv444().

yuv_to_rgb()

static __device__ void yuv_to_rgb ( float & r, float & g, float & b )

static

Transform YUV to RGB.

Definition at line 444 of file cuda_dxt.cu.

                                                                   {
    const float y = 1.1643f * (r - 0.0625f);  // TODO: convert to FFMA
    const float u = g - 0.5f;
    const float v = b - 0.5f;
    r = y + 1.7926f * v;
    g = y - 0.2132f * u - 0.5328f * v;
    b = y + 2.1124f * u;
}

Referenced by dxt_kernel().

Variable Documentation

offset

__device__ const float offset = 128.0 / 255.0

static

Definition at line 139 of file cuda_dxt.cu.

Referenced by audio_tx_send_standard(), ConvertRGBToYCoCg(), rs::decode(), decode_audio_frame(), decode_video_frame(), draw_line(), EmitEndPointsYCoCgDXT5(), frame_decode_int(), frame_encode_int(), frame_encode_int_big(), frame_encode_staircase(), libavcodec_compress(), packet_counter_register_packet(), pw_frame_to_uv_frame_memcpy(), packet_counter::register_packet(), audio_frame2::replace(), ScaleYCoCg(), vidcap_state_aja::SetupHDMI(), vidcap_state_aja::SetupVideo(), ts_add_nsec(), tv_add_usec(), video_pattern_generator_create(), and wav_seek().