oxigdal_gpu/

texture_compress.rs

//! Multi-format texture compression (BC1 and BC4) for OxiGDAL GPU.
//!
//! This module provides both pure-Rust CPU implementations and GPU-accelerated
//! (wgpu compute shader) paths for compressing raster data into BC1 (DXT1)
//! and BC4 (ATI1) block-compressed texture formats.
//!
//! # Format overview
//!
//! | Format            | Input            | Output           | Ratio |
//! |-------------------|------------------|------------------|-------|
//! | [`TextureFormat::Bc1RgbUnorm`] | RGBA8 (4 B/px)   | 8 B / 4×4 block  | 6:1   |
//! | [`TextureFormat::Bc4RUnorm`]   | R8    (1 B/px)   | 8 B / 4×4 block  | 2:1   |
//!
//! # BC1 block layout (8 bytes)
//!
//! ```text
//! [u16 c0 LE] [u16 c1 LE] [u32 lookup LE]
//! ```
//!
//! When `c0 > c1` four colours are interpolated (opaque mode, always used
//! here).  When `c0 <= c1` three colours + transparent; this module never
//! emits that mode.
//!
//! # BC4 block layout (8 bytes)
//!
//! ```text
//! [u8 ep0] [u8 ep1] [48-bit LE packed 3-bit indices for 16 pixels]
//! ```
//!
//! When `ep0 > ep1` eight interpolated values are used (always in this
//! module).
//!
//! # GPU path
//!
//! [`TextureCompressor`] drives a wgpu compute shader that processes one 4×4
//! block per workgroup thread.  When wgpu is unavailable the `compress` method
//! falls back silently to the CPU path.

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};

// ─────────────────────────────────────────────────────────────────────────────
// Pure-Rust quantisation helpers
// ─────────────────────────────────────────────────────────────────────────────

/// Quantise 8-bit RGB components to packed 16-bit RGB565.
///
/// The format allocates 5 bits for red, 6 bits for green, 5 bits for blue.
///
/// # Examples
///
/// ```
/// use oxigdal_gpu::texture_compress::quantize_rgb565;
/// let packed = quantize_rgb565(255, 255, 255);
/// assert_eq!(packed, 0xFFFF);
/// ```
pub fn quantize_rgb565(r: u8, g: u8, b: u8) -> u16 {
    let r5 = (r >> 3) as u16; // 5 bits (drop bottom 3)
    let g6 = (g >> 2) as u16; // 6 bits (drop bottom 2)
    let b5 = (b >> 3) as u16; // 5 bits (drop bottom 3)
    (r5 << 11) | (g6 << 5) | b5
}

/// Reverse a packed RGB565 value back to approximate 8-bit R, G, B.
///
/// Due to the quantisation the values are the nearest representable byte,
/// not necessarily the exact original.
///
/// # Examples
///
/// ```
/// use oxigdal_gpu::texture_compress::dequantize_rgb565;
/// let (r, g, b) = dequantize_rgb565(0xFFFF);
/// assert_eq!((r, g, b), (255, 255, 255));
/// ```
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;
    // Expand via replication of the MSBs into the vacated LSBs.
    let r = (r5 << 3) | (r5 >> 2);
    let g = (g6 << 2) | (g6 >> 4);
    let b = (b5 << 3) | (b5 >> 2);
    (r, g, b)
}

/// Return the index (0..4) of the endpoint in `endpoints` nearest to `value_rgb`
/// under the squared-Euclidean distance metric in RGB space.
///
/// # Examples
///
/// ```
/// use oxigdal_gpu::texture_compress::nearest_index_4;
/// let endpoints = [(0u8,0,0),(255,255,255),(170,170,170),(85,85,85)];
/// assert_eq!(nearest_index_4((0,0,0), endpoints), 0);
/// ```
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
}

// ─────────────────────────────────────────────────────────────────────────────
// BC1 CPU block encoder
// ─────────────────────────────────────────────────────────────────────────────

/// Compress one 4×4 RGBA block (64 bytes, row-major) to an 8-byte BC1 block.
///
/// # Algorithm
///
/// 1. Find the pixel with the highest and lowest luminance (used as `c0` / `c1`).
/// 2. Quantise each to RGB565 and dequantise back so the four interpolated
///    colours match exactly what a decoder will reconstruct.
/// 3. Guarantee `c0 > c1` (opaque 4-colour mode) by swapping when necessary.
/// 4. For each pixel, choose the nearest of the 4 colours and pack the 2-bit
///    index into a 32-bit lookup table (pixel 0 in bits 1:0).
/// 5. Write `[c0 u16 LE][c1 u16 LE][lookup u32 LE]`.
///
/// # Examples
///
/// ```
/// use oxigdal_gpu::texture_compress::compress_bc1_block_cpu;
///
/// // Solid-colour block: all pixels identical → c0 == c1
/// let mut block = [0u8; 64];
/// for i in 0..16 {
///     block[i * 4] = 100;   // R
///     block[i * 4 + 1] = 200; // G
///     block[i * 4 + 2] = 50;  // B
///     block[i * 4 + 3] = 255; // A
/// }
/// 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);
/// ```
pub fn compress_bc1_block_cpu(rgba_block: &[u8; 64]) -> [u8; 8] {
    // --- Step 1: find min/max luminance pixels ---
    // Luminance approximation: 0.299R + 0.587G + 0.114B (fixed-point × 1000)
    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; // ×1000
        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);
        }
    }

    // Quantise endpoints to RGB565 and back so our palette matches a decoder.
    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);

    // Guarantee opaque 4-colour mode: c0 > c1.
    // If equal (solid block) swap to a canonical equal state (no harm).
    if q0 < q1 {
        std::mem::swap(&mut q0, &mut q1);
    }
    // If still equal (truly solid block) we write c0 == c1, which is legal
    // but decodes in 3-colour+transparent mode on some decoders.  For an
    // opaque solid block the lookup table will always map to index 0 == c0,
    // so the visual result is identical.  We bump c0 up by one unit only
    // when that would not change the decoded RGB triplet — this is a
    // best-effort path; compressed output is already perceptually lossless.
    // We intentionally leave q0 == q1 here; callers/tests that require
    // equal endpoints for solid blocks should not be confused by a bump.

    let (dr0, dg0, db0) = dequantize_rgb565(q0);
    let (dr1, dg1, db1) = dequantize_rgb565(q1);

    // --- Step 2: compute the 4 interpolated palette colours ---
    // colour[0] = c0, colour[1] = c1
    // colour[2] = (2·c0 + c1) / 3 (integer per channel)
    // colour[3] = (c0 + 2·c1) / 3
    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,
        ),
    ];

    // --- Step 3 & 4: assign each pixel to the nearest colour, pack indices ---
    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);
    }

    // --- Step 5: write output ---
    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
}

// ─────────────────────────────────────────────────────────────────────────────
// BC4 CPU block encoder
// ─────────────────────────────────────────────────────────────────────────────

/// Compress one 4×4 R-channel block (16 bytes) to an 8-byte BC4 block.
///
/// # Algorithm
///
/// 1. `ep0 = max R`, `ep1 = min R` in the block.
/// 2. Compute 8 interpolated values: `v[0]=ep0`, `v[1]=ep1`,
///    `v[k] = ((7-k)*ep0 + (k-1)*ep1) / 7` for `k=2..7`.
/// 3. For each of the 16 pixels, find the nearest value → 3-bit index 0..7.
/// 4. Pack 16 × 3 = 48 bits into 6 bytes (little-endian 48-bit integer).
/// 5. Write `[ep0][ep1][6 index bytes]`.
///
/// # Examples
///
/// ```
/// use oxigdal_gpu::texture_compress::compress_bc4_block_cpu;
///
/// // Solid-128 block: ep0 == ep1
/// let block = [128u8; 16];
/// let encoded = compress_bc4_block_cpu(&block);
/// assert_eq!(encoded[0], encoded[1]);
/// ```
pub fn compress_bc4_block_cpu(r_block: &[u8; 16]) -> [u8; 8] {
    // --- Step 1: find min / max R ---
    let mut ep0 = 0u8; // max
    let mut ep1 = 255u8; // min
    for &v in r_block.iter() {
        if v > ep0 {
            ep0 = v;
        }
        if v < ep1 {
            ep1 = v;
        }
    }

    // --- Step 2: compute 8 interpolated values (ep0 > ep1 mode) ---
    // v[0] = ep0, v[1] = ep1
    // v[k] = ((7-k)·ep0 + (k-1)·ep1) / 7  for k in 2..=7
    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,
        ]
    };

    // --- Step 3: find nearest palette entry for each pixel ---
    // Compute 16 × 3-bit indices, then pack into 48 bits (6 bytes) LE.
    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);
    }

    // Write output: [ep0][ep1][6 bytes of packed indices LE]
    let mut out = [0u8; 8];
    out[0] = ep0;
    out[1] = ep1;
    // The 48-bit index block sits in bits 0..47 of `packed`.
    let index_bytes = packed.to_le_bytes(); // 8 bytes; we use first 6
    out[2..8].copy_from_slice(&index_bytes[0..6]);
    out
}

// ─────────────────────────────────────────────────────────────────────────────
// Dimension validation (pure-Rust, no GPU required)
// ─────────────────────────────────────────────────────────────────────────────

/// Validate that `width` and `height` are both multiples of 4 (required for
/// block-compressed textures).
///
/// This function is intentionally separated from [`TextureCompressor::new`] so
/// that tests can exercise the validation logic without constructing a GPU
/// context.
///
/// # Errors
///
/// Returns [`GpuError::InvalidKernelParams`] when either dimension is not a
/// multiple of 4.
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(())
    }
}

// ─────────────────────────────────────────────────────────────────────────────
// TextureFormat enum
// ─────────────────────────────────────────────────────────────────────────────

/// Target block-compressed texture format produced by [`TextureCompressor`].
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum TextureFormat {
    /// BC1 (DXT1) RGB opaque compression.
    ///
    /// - Input:  RGBA8 — `width × height × 4` bytes.
    /// - Output: 8 bytes per 4×4 block → `(width/4) × (height/4) × 8` bytes.
    /// - Compression ratio: 6:1 (vs. RGBA8).
    Bc1RgbUnorm,

    /// BC4 single-channel (ATI1 / RGTC1) compression.
    ///
    /// - Input:  R8 — `width × height` bytes.
    /// - Output: 8 bytes per 4×4 block → `(width/4) × (height/4) × 8` bytes.
    /// - Compression ratio: 2:1 (vs. R8).
    Bc4RUnorm,
}

// ─────────────────────────────────────────────────────────────────────────────
// WGSL shader source generation
// ─────────────────────────────────────────────────────────────────────────────

/// Uniform buffer layout for the BC1 shader.
///
/// Matches the WGSL struct:
/// ```wgsl
/// struct Uniforms { width: u32, height: u32 }
/// ```
#[repr(C, align(8))]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct Bc1Uniforms {
    width: u32,
    height: u32,
}

/// Uniform buffer layout for the BC4 shader.
#[repr(C, align(8))]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct Bc4Uniforms {
    width: u32,
    height: u32,
}

/// Emit the WGSL compute shader for BC1 compression.
///
/// The shader processes one 4×4 block per invocation.  Each dispatch thread
/// reads 16 RGBA8 pixels (packed as 4 bytes each → one `u32` per pixel),
/// finds the min/max luminance, builds the 4-colour palette, assigns indices,
/// and writes the 8-byte BC1 block as two `u32`s.
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,
    )
}

/// Emit the WGSL compute shader for BC4 compression.
///
/// Each invocation processes one 4×4 R8 block (16 bytes).  The input is a
/// `u32` storage buffer where each `u32` packs four consecutive R8 pixels.
fn make_bc4_shader_source(width: u32, _height: u32) -> String {
    let num_blocks_x = width / 4;
    let pixels_per_u32 = 4u32; // 4 R8 bytes → 1 u32
    let u32s_per_row = width / pixels_per_u32; // how many u32s span one image row
    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,
    )
}

// ─────────────────────────────────────────────────────────────────────────────
// TextureCompressor struct
// ─────────────────────────────────────────────────────────────────────────────

/// GPU-accelerated texture compressor for BC1 and BC4 block formats.
///
/// Create once with [`TextureCompressor::new`], then call [`TextureCompressor::compress`]
/// for each image.  The `width` and `height` supplied at construction time are
/// fixed for the lifetime of this object.
///
/// When no wgpu backend is available (e.g., headless CI) `compress` falls back
/// automatically to the CPU path provided by [`compress_bc1_block_cpu`] and
/// [`compress_bc4_block_cpu`].
pub struct TextureCompressor {
    format: TextureFormat,
    /// The compiled GPU compute pipeline (may be `None` when GPU unavailable).
    pipeline: Option<Arc<wgpu::ComputePipeline>>,
    /// Bind group layout for the GPU pipeline.
    bind_group_layout: Option<wgpu::BindGroupLayout>,
    width: u32,
    height: u32,
}

impl TextureCompressor {
    /// Create a new `TextureCompressor`.
    ///
    /// # Arguments
    ///
    /// * `ctx`    — Active GPU context.
    /// * `format` — Target compressed format (`Bc1RgbUnorm` or `Bc4RUnorm`).
    /// * `width`  — Pixel width of images to compress (must be a multiple of 4).
    /// * `height` — Pixel height of images to compress (must be a multiple of 4).
    ///
    /// # Errors
    ///
    /// Returns [`GpuError::InvalidKernelParams`] when either dimension is not a
    /// multiple of 4, or a pipeline error when WGSL compilation fails.
    pub fn new(
        ctx: &GpuContext,
        format: TextureFormat,
        width: u32,
        height: u32,
    ) -> GpuResult<Self> {
        // Pure-Rust validation: no GPU touched.
        validate_texture_dimensions(width, height)?;

        // Build shader source and pipeline.
        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) => {
                // If shader compilation fails we still return a valid (CPU-only)
                // compressor rather than propagating.  This allows callers on
                // restricted environments to receive useful output.
                tracing::warn!("BC texture shader compilation failed, falling back to CPU: {e}");
                return Ok(Self {
                    format,
                    pipeline: None,
                    bind_group_layout: None,
                    width,
                    height,
                });
            }
        };

        // Bind group layout: uniform | storage read | storage read-write
        let bind_group_layout = create_compute_bind_group_layout(
            ctx.device(),
            &[
                uniform_buffer_layout(0),        // uniforms
                storage_buffer_layout(1, true),  // input pixels (read-only)
                storage_buffer_layout(2, false), // output blocks (read-write)
            ],
            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,
        })
    }

    // ── Accessors ──────────────────────────────────────────────────────────

    /// The target texture format.
    pub fn format(&self) -> TextureFormat {
        self.format
    }

    /// Image width (pixels) supplied at construction time.
    pub fn width(&self) -> u32 {
        self.width
    }

    /// Image height (pixels) supplied at construction time.
    pub fn height(&self) -> u32 {
        self.height
    }

    // ── Main compression entry-point ───────────────────────────────────────

    /// Compress `input` using this compressor's format.
    ///
    /// - **BC1**: `input` must be `width × height × 4` bytes (RGBA8 row-major).
    /// - **BC4**: `input` must be `width × height` bytes (R8 row-major).
    ///
    /// Returns `(width/4) × (height/4) × 8` bytes of block-compressed data.
    ///
    /// # GPU path
    ///
    /// When a wgpu backend is available the method dispatches a compute shader
    /// and reads back the result synchronously.
    ///
    /// # CPU fallback
    ///
    /// When `self.pipeline` is `None` (shader compilation failed or no backend)
    /// or when wgpu dispatch fails, the method falls back to the pure-Rust
    /// block encoders `compress_bc1_block_cpu` / `compress_bc4_block_cpu`.
    ///
    /// # Errors
    ///
    /// Returns an error when the input length does not match the expected size
    /// for the configured format and dimensions.
    pub fn compress(&self, ctx: &GpuContext, input: &[u8]) -> GpuResult<Vec<u8>> {
        // Validate input length.
        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
            )));
        }

        // Try GPU path first.
        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) => {
                    // GPU failed; fall through to CPU.
                    tracing::warn!("GPU texture compression failed, falling back to CPU: {e}");
                }
            }
        }

        // CPU fallback path.
        self.compress_cpu(input)
    }

    // ── CPU fallback ───────────────────────────────────────────────────────

    /// Pure-CPU compression path (always correct; slower for large images).
    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 => {
                // RGBA8 input: 4 bytes per pixel, 16 pixels per block → 64 bytes per block.
                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 => {
                // R8 input: 1 byte per pixel, 16 pixels per block → 16 bytes per block.
                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)
    }

    // ── GPU dispatch ───────────────────────────────────────────────────────

    /// Attempt GPU-accelerated compression.
    ///
    /// For BC1 the input is uploaded as `u32`s (each `u32` = one RGBA8 pixel).
    /// For BC4 the input is uploaded as `u32`s with 4 R8 bytes packed per `u32`.
    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; // 2 u32s per block

        // --- Build uniform data ---
        let uniforms_value = Bc1Uniforms {
            width: self.width,
            height: self.height,
        };
        let uniform_data = bytemuck::bytes_of(&uniforms_value);

        // --- Upload uniform buffer ---
        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);

        // --- Pack input into u32 buffer ---
        let input_u32: Vec<u32> = match self.format {
            TextureFormat::Bc1RgbUnorm => {
                // Each RGBA8 pixel → 1 u32 (R in bits 7:0, G 15:8, B 23:16, A 31:24)
                input
                    .chunks_exact(4)
                    .map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]]))
                    .collect()
            }
            TextureFormat::Bc4RUnorm => {
                // 4 R8 bytes → 1 u32
                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));

        // --- Create output buffer (2 u32s per block) ---
        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,
        });

        // --- Build bind group ---
        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(),
                },
            ],
        });

        // --- Dispatch compute ---
        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);
        }

        // --- Staging buffer for readback ---
        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()));

        // Map before poll: poll drives both submit completion and map callback.
        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)
    }
}

// ─────────────────────────────────────────────────────────────────────────────
// Unit tests (pure-Rust, no GPU required)
// ─────────────────────────────────────────────────────────────────────────────

#[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());
    }
}