use crate::context::GpuContext;
use crate::error::{GpuError, GpuResult};
use crate::shaders::{
ComputePipelineBuilder, WgslShader, create_compute_bind_group_layout, storage_buffer_layout,
uniform_buffer_layout,
};
use std::sync::Arc;
use wgpu::{BindGroupDescriptor, BindGroupEntry, BufferDescriptor, BufferUsages};
pub fn quantize_rgb565(r: u8, g: u8, b: u8) -> u16 {
let r5 = (r >> 3) as u16; let g6 = (g >> 2) as u16; let b5 = (b >> 3) as u16; (r5 << 11) | (g6 << 5) | b5
}
pub fn dequantize_rgb565(rgb565: u16) -> (u8, u8, u8) {
let r5 = ((rgb565 >> 11) & 0x1F) as u8;
let g6 = ((rgb565 >> 5) & 0x3F) as u8;
let b5 = (rgb565 & 0x1F) as u8;
let r = (r5 << 3) | (r5 >> 2);
let g = (g6 << 2) | (g6 >> 4);
let b = (b5 << 3) | (b5 >> 2);
(r, g, b)
}
pub fn nearest_index_4(value_rgb: (u8, u8, u8), endpoints: [(u8, u8, u8); 4]) -> u8 {
let (vr, vg, vb) = value_rgb;
let mut best_idx = 0u8;
let mut best_dist = u32::MAX;
for (i, (er, eg, eb)) in endpoints.iter().enumerate() {
let dr = (vr as i32) - (*er as i32);
let dg = (vg as i32) - (*eg as i32);
let db = (vb as i32) - (*eb as i32);
let dist = (dr * dr + dg * dg + db * db) as u32;
if dist < best_dist {
best_dist = dist;
best_idx = i as u8;
}
}
best_idx
}
pub fn compress_bc1_block_cpu(rgba_block: &[u8; 64]) -> [u8; 8] {
let mut max_lum: u32 = 0;
let mut min_lum: u32 = u32::MAX;
let mut max_px = (0u8, 0u8, 0u8);
let mut min_px = (0u8, 0u8, 0u8);
for i in 0..16usize {
let r = rgba_block[i * 4] as u32;
let g = rgba_block[i * 4 + 1] as u32;
let b = rgba_block[i * 4 + 2] as u32;
let lum = r * 299 + g * 587 + b * 114; if lum > max_lum {
max_lum = lum;
max_px = (r as u8, g as u8, b as u8);
}
if lum < min_lum {
min_lum = lum;
min_px = (r as u8, g as u8, b as u8);
}
}
let mut q0 = quantize_rgb565(max_px.0, max_px.1, max_px.2);
let mut q1 = quantize_rgb565(min_px.0, min_px.1, min_px.2);
if q0 < q1 {
std::mem::swap(&mut q0, &mut q1);
}
let (dr0, dg0, db0) = dequantize_rgb565(q0);
let (dr1, dg1, db1) = dequantize_rgb565(q1);
let colours: [(u8, u8, u8); 4] = [
(dr0, dg0, db0),
(dr1, dg1, db1),
(
((2 * dr0 as u32 + dr1 as u32) / 3) as u8,
((2 * dg0 as u32 + dg1 as u32) / 3) as u8,
((2 * db0 as u32 + db1 as u32) / 3) as u8,
),
(
((dr0 as u32 + 2 * dr1 as u32) / 3) as u8,
((dg0 as u32 + 2 * dg1 as u32) / 3) as u8,
((db0 as u32 + 2 * db1 as u32) / 3) as u8,
),
];
let mut lookup: u32 = 0;
for i in 0..16usize {
let r = rgba_block[i * 4];
let g = rgba_block[i * 4 + 1];
let b = rgba_block[i * 4 + 2];
let idx = nearest_index_4((r, g, b), colours) as u32;
lookup |= idx << (i * 2);
}
let mut out = [0u8; 8];
out[0..2].copy_from_slice(&q0.to_le_bytes());
out[2..4].copy_from_slice(&q1.to_le_bytes());
out[4..8].copy_from_slice(&lookup.to_le_bytes());
out
}
pub fn compress_bc4_block_cpu(r_block: &[u8; 16]) -> [u8; 8] {
let mut ep0 = 0u8; let mut ep1 = 255u8; for &v in r_block.iter() {
if v > ep0 {
ep0 = v;
}
if v < ep1 {
ep1 = v;
}
}
let palette: [u8; 8] = {
let e0 = ep0 as u32;
let e1 = ep1 as u32;
[
ep0,
ep1,
((6 * e0 + e1) / 7) as u8,
((5 * e0 + 2 * e1) / 7) as u8,
((4 * e0 + 3 * e1) / 7) as u8,
((3 * e0 + 4 * e1) / 7) as u8,
((2 * e0 + 5 * e1) / 7) as u8,
((e0 + 6 * e1) / 7) as u8,
]
};
let mut packed: u64 = 0u64;
for (i, &pixel_val) in r_block.iter().enumerate() {
let val = pixel_val as i32;
let mut best_idx = 0u8;
let mut best_dist = i32::MAX;
for (j, &pv) in palette.iter().enumerate() {
let d = (val - pv as i32).abs();
if d < best_dist {
best_dist = d;
best_idx = j as u8;
}
}
packed |= (best_idx as u64) << (i * 3);
}
let mut out = [0u8; 8];
out[0] = ep0;
out[1] = ep1;
let index_bytes = packed.to_le_bytes(); out[2..8].copy_from_slice(&index_bytes[0..6]);
out
}
pub fn validate_texture_dimensions(width: u32, height: u32) -> GpuResult<()> {
if width % 4 != 0 || height % 4 != 0 {
Err(GpuError::invalid_kernel_params(
"width and height must be multiples of 4",
))
} else {
Ok(())
}
}
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum TextureFormat {
Bc1RgbUnorm,
Bc4RUnorm,
}
#[repr(C, align(8))]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct Bc1Uniforms {
width: u32,
height: u32,
}
#[repr(C, align(8))]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct Bc4Uniforms {
width: u32,
height: u32,
}
fn make_bc1_shader_source(width: u32, _height: u32) -> String {
let num_blocks_x = width / 4;
format!(
r#"// BC1 (DXT1) block-compression compute shader.
// One thread per 4×4 block; input is RGBA8 packed as u32 (1 u32 per pixel).
struct Uniforms {{
width: u32,
height: u32,
}};
@group(0) @binding(0) var<uniform> uniforms: Uniforms;
@group(0) @binding(1) var<storage, read> input_pixels: array<u32>;
@group(0) @binding(2) var<storage, read_write> output_blocks: array<u32>;
// Squared Euclidean distance in RGB space between two packed RGBA pixels.
fn rgb_dist_sq(a: u32, b: u32) -> u32 {{
let ar = (a ) & 0xFFu;
let ag = (a >> 8u) & 0xFFu;
let ab = (a >> 16u) & 0xFFu;
let br = (b ) & 0xFFu;
let bg = (b >> 8u) & 0xFFu;
let bb = (b >> 16u) & 0xFFu;
let dr = select(ar - br, br - ar, br > ar);
let dg = select(ag - bg, bg - ag, bg > ag);
let db = select(ab - bb, bb - ab, bb > ab);
return dr*dr + dg*dg + db*db;
}}
// Compute luminance proxy (×1000) from packed RGBA.
fn luminance(px: u32) -> u32 {{
let r = (px ) & 0xFFu;
let g = (px >> 8u) & 0xFFu;
let b = (px >> 16u) & 0xFFu;
return r * 299u + g * 587u + b * 114u;
}}
// Quantise RGB8 (packed in u32 low 24 bits) to RGB565.
fn quantize_565(px: u32) -> u32 {{
let r = (px ) & 0xFFu;
let g = (px >> 8u) & 0xFFu;
let b = (px >> 16u) & 0xFFu;
let r5 = r >> 3u;
let g6 = g >> 2u;
let b5 = b >> 3u;
return (r5 << 11u) | (g6 << 5u) | b5;
}}
// Dequantise RGB565 → u32 with R in bits 7:0, G in bits 15:8, B in bits 23:16.
fn dequantize_565(c: u32) -> u32 {{
let r5 = (c >> 11u) & 0x1Fu;
let g6 = (c >> 5u) & 0x3Fu;
let b5 = c & 0x1Fu;
let r = (r5 << 3u) | (r5 >> 2u);
let g = (g6 << 2u) | (g6 >> 4u);
let b = (b5 << 3u) | (b5 >> 2u);
return r | (g << 8u) | (b << 16u);
}}
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {{
let block_idx = gid.x;
let num_blocks_x = {num_blocks_x}u;
let num_blocks_y = uniforms.height / 4u;
let total_blocks = num_blocks_x * num_blocks_y;
if block_idx >= total_blocks {{
return;
}}
let bx = block_idx % num_blocks_x;
let by = block_idx / num_blocks_x;
// --- Gather 16 pixels ---
var pixels: array<u32, 16>;
for (var iy = 0u; iy < 4u; iy++) {{
for (var ix = 0u; ix < 4u; ix++) {{
let px_x = bx * 4u + ix;
let px_y = by * 4u + iy;
let px_idx = px_y * uniforms.width + px_x;
pixels[iy * 4u + ix] = input_pixels[px_idx];
}}
}}
// --- Find min/max luminance pixel ---
var max_lum = luminance(pixels[0]);
var min_lum = max_lum;
var max_px = pixels[0];
var min_px = pixels[0];
for (var i = 1u; i < 16u; i++) {{
let lum = luminance(pixels[i]);
if lum > max_lum {{ max_lum = lum; max_px = pixels[i]; }}
if lum < min_lum {{ min_lum = lum; min_px = pixels[i]; }}
}}
// --- Quantise to RGB565 (c0 = max luminance = bright, c1 = dark) ---
var q0 = quantize_565(max_px);
var q1 = quantize_565(min_px);
// Guarantee c0 > c1 (opaque 4-colour mode).
if q0 < q1 {{ let tmp = q0; q0 = q1; q1 = tmp; }}
// Dequantise to reconstruct the palette the decoder will use.
let d0 = dequantize_565(q0);
let d1 = dequantize_565(q1);
// --- Build 4-colour palette (colour[2] and colour[3] interpolated) ---
let r0 = (d0 ) & 0xFFu;
let g0 = (d0 >> 8u) & 0xFFu;
let b0 = (d0 >> 16u) & 0xFFu;
let r1 = (d1 ) & 0xFFu;
let g1 = (d1 >> 8u) & 0xFFu;
let b1 = (d1 >> 16u) & 0xFFu;
var palette: array<u32, 4>;
palette[0] = d0;
palette[1] = d1;
palette[2] = ((2u*r0+r1)/3u) | (((2u*g0+g1)/3u)<<8u) | (((2u*b0+b1)/3u)<<16u);
palette[3] = ((r0+2u*r1)/3u) | (((g0+2u*g1)/3u)<<8u) | (((b0+2u*b1)/3u)<<16u);
// --- Assign each pixel to the nearest palette entry ---
var lookup = 0u;
for (var i = 0u; i < 16u; i++) {{
let px = pixels[i];
var best_idx = 0u;
var best_dist = rgb_dist_sq(px, palette[0]);
for (var j = 1u; j < 4u; j++) {{
let d = rgb_dist_sq(px, palette[j]);
if d < best_dist {{ best_dist = d; best_idx = j; }}
}}
lookup |= best_idx << (i * 2u);
}}
// --- Write BC1 block: [c0 u16 LE | c1 u16 LE] as u32, then lookup u32 ---
let word0 = q0 | (q1 << 16u);
output_blocks[block_idx * 2u ] = word0;
output_blocks[block_idx * 2u + 1u] = lookup;
}}
"#,
num_blocks_x = num_blocks_x,
)
}
fn make_bc4_shader_source(width: u32, _height: u32) -> String {
let num_blocks_x = width / 4;
let pixels_per_u32 = 4u32; let u32s_per_row = width / pixels_per_u32; format!(
r#"// BC4 (ATI1 / RGTC1) single-channel block-compression compute shader.
// One thread per 4×4 block; input R8 pixels are packed 4-per-u32.
struct Uniforms {{
width: u32,
height: u32,
}};
@group(0) @binding(0) var<uniform> uniforms: Uniforms;
@group(0) @binding(1) var<storage, read> input_r8: array<u32>;
@group(0) @binding(2) var<storage, read_write> output_blocks: array<u32>;
// Extract byte `lane` (0..3) from a packed u32 (little-endian byte order).
fn extract_byte(word: u32, lane: u32) -> u32 {{
return (word >> (lane * 8u)) & 0xFFu;
}}
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {{
let block_idx = gid.x;
let num_blocks_x = {num_blocks_x}u;
let num_blocks_y = uniforms.height / 4u;
let total_blocks = num_blocks_x * num_blocks_y;
if block_idx >= total_blocks {{
return;
}}
let bx = block_idx % num_blocks_x;
let by = block_idx / num_blocks_x;
// --- Gather 16 R8 pixels from packed u32 storage ---
// Each u32 holds 4 consecutive R8 bytes in little-endian order.
var pixels: array<u32, 16>;
let u32s_per_row = {u32s_per_row}u;
for (var iy = 0u; iy < 4u; iy++) {{
let row_base = (by * 4u + iy) * u32s_per_row;
let col_u32 = (bx * 4u) / 4u; // which u32 word holds column bx*4
let col_lane = (bx * 4u) % 4u; // which byte lane within that word
// We need 4 consecutive bytes starting at (bx*4, iy).
// Since bx*4 is always a multiple of 4 (width is a multiple of 4),
// all 4 bytes live in the same u32 word.
let word = input_r8[row_base + col_u32];
for (var ix = 0u; ix < 4u; ix++) {{
pixels[iy * 4u + ix] = extract_byte(word, col_lane + ix);
}}
}}
// --- Find ep0 (max) and ep1 (min) ---
var ep0 = pixels[0];
var ep1 = pixels[0];
for (var i = 1u; i < 16u; i++) {{
if pixels[i] > ep0 {{ ep0 = pixels[i]; }}
if pixels[i] < ep1 {{ ep1 = pixels[i]; }}
}}
// --- Build 8-value palette (ep0 > ep1 mode) ---
// v[0]=ep0, v[1]=ep1, v[2..7] interpolated
var palette: array<u32, 8>;
palette[0] = ep0;
palette[1] = ep1;
palette[2] = (6u*ep0 + 1u*ep1) / 7u;
palette[3] = (5u*ep0 + 2u*ep1) / 7u;
palette[4] = (4u*ep0 + 3u*ep1) / 7u;
palette[5] = (3u*ep0 + 4u*ep1) / 7u;
palette[6] = (2u*ep0 + 5u*ep1) / 7u;
palette[7] = (1u*ep0 + 6u*ep1) / 7u;
// --- Assign 3-bit indices ---
// Pack 16 × 3-bit indices into 48 bits → stored as two u32s.
// We fill a 64-bit logical value then split it: lo = bits 31:0, hi = bits 47:32.
var indices_lo = 0u; // bits 0..31 (10.67 indices)
var indices_hi = 0u; // bits 32..47 ( 5.33 indices) — only low 16 bits used
for (var i = 0u; i < 16u; i++) {{
let val = pixels[i];
// Find nearest palette entry.
var best_idx = 0u;
var best_dist = select(ep0 - val, val - ep0, val > ep0); // abs
for (var j = 1u; j < 8u; j++) {{
let pv = palette[j];
let d = select(pv - val, val - pv, val > pv);
if d < best_dist {{ best_dist = d; best_idx = j; }}
}}
// Bit position of this index in the 48-bit packed field.
let bit_pos = i * 3u;
if bit_pos < 32u {{
indices_lo |= best_idx << bit_pos;
// Handle straddle: if best_idx's 3 bits cross the 32-bit boundary.
if bit_pos > 29u {{
let overflow_bits = bit_pos + 3u - 32u;
indices_hi |= best_idx >> (3u - overflow_bits);
}}
}} else {{
indices_hi |= best_idx << (bit_pos - 32u);
}}
}}
// --- Write BC4 block ---
// Layout: [ep0 u8][ep1 u8][6 index bytes LE]
// We emit as two u32s:
// word0 = ep0 | (ep1 << 8) | (index_byte0 << 16) | (index_byte1 << 24)
// word1 = index_byte2..5
// indices_lo holds index bytes 0,1,2,3; indices_hi holds bytes 4,5 in low 16 bits.
let word0 = ep0 | (ep1 << 8u) | ((indices_lo & 0xFFFFu) << 16u);
let word1 = (indices_lo >> 16u) | (indices_hi << 16u);
// Note: word1 bits 16..31 hold index bytes 4&5 (indices_hi low 16 bits shifted up).
// Actual layout is: word0[16:31] = idx_bytes[0:1], word1[0:15] = idx_bytes[2:3],
// word1[16:31] = idx_bytes[4:5].
output_blocks[block_idx * 2u ] = word0;
output_blocks[block_idx * 2u + 1u] = word1;
}}
"#,
num_blocks_x = num_blocks_x,
u32s_per_row = u32s_per_row,
)
}
pub struct TextureCompressor {
format: TextureFormat,
pipeline: Option<Arc<wgpu::ComputePipeline>>,
bind_group_layout: Option<wgpu::BindGroupLayout>,
width: u32,
height: u32,
}
impl TextureCompressor {
pub fn new(
ctx: &GpuContext,
format: TextureFormat,
width: u32,
height: u32,
) -> GpuResult<Self> {
validate_texture_dimensions(width, height)?;
let shader_src = match format {
TextureFormat::Bc1RgbUnorm => make_bc1_shader_source(width, height),
TextureFormat::Bc4RUnorm => make_bc4_shader_source(width, height),
};
let mut shader = WgslShader::new(shader_src, "main");
let shader_module = match shader.compile(ctx.device()) {
Ok(m) => m,
Err(e) => {
tracing::warn!("BC texture shader compilation failed, falling back to CPU: {e}");
return Ok(Self {
format,
pipeline: None,
bind_group_layout: None,
width,
height,
});
}
};
let bind_group_layout = create_compute_bind_group_layout(
ctx.device(),
&[
uniform_buffer_layout(0), storage_buffer_layout(1, true), storage_buffer_layout(2, false), ],
Some("TextureCompressor BGL"),
)?;
let pipeline = ComputePipelineBuilder::new(ctx.device(), shader_module, "main")
.bind_group_layout(&bind_group_layout)
.label("TextureCompressor Pipeline")
.build()?;
Ok(Self {
format,
pipeline: Some(Arc::new(pipeline)),
bind_group_layout: Some(bind_group_layout),
width,
height,
})
}
pub fn format(&self) -> TextureFormat {
self.format
}
pub fn width(&self) -> u32 {
self.width
}
pub fn height(&self) -> u32 {
self.height
}
pub fn compress(&self, ctx: &GpuContext, input: &[u8]) -> GpuResult<Vec<u8>> {
let expected = match self.format {
TextureFormat::Bc1RgbUnorm => (self.width * self.height * 4) as usize,
TextureFormat::Bc4RUnorm => (self.width * self.height) as usize,
};
if input.len() != expected {
return Err(GpuError::invalid_kernel_params(format!(
"input length {} != expected {} for {:?} {}×{}",
input.len(),
expected,
self.format,
self.width,
self.height
)));
}
if let (Some(pipeline), Some(bgl)) = (&self.pipeline, &self.bind_group_layout) {
match self.compress_gpu(ctx, input, pipeline, bgl) {
Ok(output) => return Ok(output),
Err(e) => {
tracing::warn!("GPU texture compression failed, falling back to CPU: {e}");
}
}
}
self.compress_cpu(input)
}
fn compress_cpu(&self, input: &[u8]) -> GpuResult<Vec<u8>> {
let blocks_x = (self.width / 4) as usize;
let blocks_y = (self.height / 4) as usize;
let total_blocks = blocks_x * blocks_y;
let mut output = vec![0u8; total_blocks * 8];
match self.format {
TextureFormat::Bc1RgbUnorm => {
for by in 0..blocks_y {
for bx in 0..blocks_x {
let mut rgba_block = [0u8; 64];
for iy in 0..4 {
for ix in 0..4 {
let px_x = bx * 4 + ix;
let px_y = by * 4 + iy;
let src = (px_y * self.width as usize + px_x) * 4;
let dst = (iy * 4 + ix) * 4;
rgba_block[dst..dst + 4].copy_from_slice(&input[src..src + 4]);
}
}
let encoded = compress_bc1_block_cpu(&rgba_block);
let block_idx = by * blocks_x + bx;
output[block_idx * 8..block_idx * 8 + 8].copy_from_slice(&encoded);
}
}
}
TextureFormat::Bc4RUnorm => {
for by in 0..blocks_y {
for bx in 0..blocks_x {
let mut r_block = [0u8; 16];
for iy in 0..4 {
for ix in 0..4 {
let px_x = bx * 4 + ix;
let px_y = by * 4 + iy;
let src = px_y * self.width as usize + px_x;
r_block[iy * 4 + ix] = input[src];
}
}
let encoded = compress_bc4_block_cpu(&r_block);
let block_idx = by * blocks_x + bx;
output[block_idx * 8..block_idx * 8 + 8].copy_from_slice(&encoded);
}
}
}
}
Ok(output)
}
fn compress_gpu(
&self,
ctx: &GpuContext,
input: &[u8],
pipeline: &wgpu::ComputePipeline,
bgl: &wgpu::BindGroupLayout,
) -> GpuResult<Vec<u8>> {
let blocks_x = self.width / 4;
let blocks_y = self.height / 4;
let total_blocks = (blocks_x * blocks_y) as usize;
let output_bytes = total_blocks * 8;
let uniforms_value = Bc1Uniforms {
width: self.width,
height: self.height,
};
let uniform_data = bytemuck::bytes_of(&uniforms_value);
let uniform_buf_size = (std::mem::size_of::<Bc1Uniforms>() as u64 + 15) & !15;
let uniform_buf = ctx.device().create_buffer(&BufferDescriptor {
label: Some("TC uniform"),
size: uniform_buf_size,
usage: BufferUsages::UNIFORM | BufferUsages::COPY_DST,
mapped_at_creation: false,
});
ctx.queue().write_buffer(&uniform_buf, 0, uniform_data);
let input_u32: Vec<u32> = match self.format {
TextureFormat::Bc1RgbUnorm => {
input
.chunks_exact(4)
.map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]]))
.collect()
}
TextureFormat::Bc4RUnorm => {
let padded_len = (input.len() + 3) & !3;
let mut padded = input.to_vec();
padded.resize(padded_len, 0);
padded
.chunks_exact(4)
.map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]]))
.collect()
}
};
let input_buf_size = ((input_u32.len() * std::mem::size_of::<u32>()) as u64 + 255) & !255;
let input_buf = ctx.device().create_buffer(&BufferDescriptor {
label: Some("TC input"),
size: input_buf_size,
usage: BufferUsages::STORAGE | BufferUsages::COPY_DST,
mapped_at_creation: false,
});
ctx.queue()
.write_buffer(&input_buf, 0, bytemuck::cast_slice(&input_u32));
let output_buf_size = ((total_blocks * 2 * std::mem::size_of::<u32>()) as u64 + 255) & !255;
let output_buf = ctx.device().create_buffer(&BufferDescriptor {
label: Some("TC output"),
size: output_buf_size,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
let bind_group = ctx.device().create_bind_group(&BindGroupDescriptor {
label: Some("TC BindGroup"),
layout: bgl,
entries: &[
BindGroupEntry {
binding: 0,
resource: uniform_buf.as_entire_binding(),
},
BindGroupEntry {
binding: 1,
resource: input_buf.as_entire_binding(),
},
BindGroupEntry {
binding: 2,
resource: output_buf.as_entire_binding(),
},
],
});
let workgroup_size = 64u32;
let dispatch_count = (total_blocks as u32 + workgroup_size - 1) / workgroup_size;
let mut encoder = ctx
.device()
.create_command_encoder(&wgpu::CommandEncoderDescriptor {
label: Some("TC encoder"),
});
{
let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: Some("TC compute pass"),
timestamp_writes: None,
});
cpass.set_pipeline(pipeline);
cpass.set_bind_group(0, &bind_group, &[]);
cpass.dispatch_workgroups(dispatch_count, 1, 1);
}
let staging = ctx.device().create_buffer(&BufferDescriptor {
label: Some("TC staging"),
size: output_buf_size,
usage: BufferUsages::MAP_READ | BufferUsages::COPY_DST,
mapped_at_creation: false,
});
encoder.copy_buffer_to_buffer(&output_buf, 0, &staging, 0, output_buf_size);
ctx.queue().submit(Some(encoder.finish()));
let slice = staging.slice(..);
let (tx, rx) = std::sync::mpsc::sync_channel(1);
slice.map_async(wgpu::MapMode::Read, move |result| {
let _ = tx.send(result);
});
ctx.device()
.poll(wgpu::PollType::wait_indefinitely())
.map_err(|e| GpuError::execution_failed(format!("device poll failed: {e}")))?;
rx.recv()
.map_err(|_| GpuError::buffer_mapping("TC staging channel closed"))?
.map_err(|e| GpuError::buffer_mapping(e.to_string()))?;
let mapped = slice.get_mapped_range();
let output_u32: &[u32] = bytemuck::cast_slice(&mapped[..total_blocks * 8]);
let out_bytes: Vec<u8> = bytemuck::cast_slice(output_u32)[..output_bytes].to_vec();
drop(mapped);
staging.unmap();
Ok(out_bytes)
}
}
#[cfg(test)]
mod tests {
use super::*;
#[test]
fn test_quantize_dequantize_round_trip_basics() {
let packed = quantize_rgb565(255, 255, 255);
let (r, g, b) = dequantize_rgb565(packed);
assert_eq!((r, g, b), (255, 255, 255));
let packed0 = quantize_rgb565(0, 0, 0);
let (r0, g0, b0) = dequantize_rgb565(packed0);
assert_eq!((r0, g0, b0), (0, 0, 0));
}
#[test]
fn test_nearest_index_4_black_is_index_0() {
let endpoints = [
(0u8, 0u8, 0u8),
(255, 255, 255),
(170, 170, 170),
(85, 85, 85),
];
assert_eq!(nearest_index_4((0, 0, 0), endpoints), 0);
}
#[test]
fn test_bc1_solid_block_equal_endpoints() {
let mut block = [0u8; 64];
for i in 0..16 {
block[i * 4] = 100;
block[i * 4 + 1] = 200;
block[i * 4 + 2] = 50;
block[i * 4 + 3] = 255;
}
let encoded = compress_bc1_block_cpu(&block);
let c0 = u16::from_le_bytes([encoded[0], encoded[1]]);
let c1 = u16::from_le_bytes([encoded[2], encoded[3]]);
assert_eq!(c0, c1, "solid block must produce equal endpoints");
}
#[test]
fn test_bc4_solid_block_equal_endpoints() {
let block = [128u8; 16];
let encoded = compress_bc4_block_cpu(&block);
assert_eq!(encoded[0], encoded[1], "solid BC4 block must have ep0==ep1");
}
#[test]
fn test_validate_texture_dimensions_multiples_ok() {
assert!(validate_texture_dimensions(4, 4).is_ok());
assert!(validate_texture_dimensions(256, 512).is_ok());
}
#[test]
fn test_validate_texture_dimensions_rejects_non_multiple() {
assert!(validate_texture_dimensions(7, 8).is_err());
assert!(validate_texture_dimensions(8, 5).is_err());
assert!(validate_texture_dimensions(1, 1).is_err());
}
}