Equal endpoints

README.md

@@ -198,6 +198,10 @@ For each set bit in the mask, the corresponding difference byte is encoded using
 
 This avoids storing zero bytes in the residual and makes the representation particularly compact when the block is close to one of the top-table patterns.
 
+#### Flat blocks
+
+A block whose two endpoint colors are identical carries a uniform index word: all 16 indices hold the same 2-bit value. A global flat index mode, decided once in the first pass, collapses this: mode 1 is selected when at least one block has identical endpoints and every such block holds the same 2-bit value. In mode 1 a flat block stores no index payload at all: the value, stored once for the whole image in the stream header, is expanded to all 16 indices on decode, instead of the top-table reference + mask + differences; in mode 0, or for a non-flat block, the indices are stored as usual. The mode and the value share one header byte (mode | value << 1) right after the top-table size. A decoder never needs the mode for a non-flat block, so the encoder and decoder always pick the same path; an encoder that leaves a non-uniform or a differing index word on a flat block forces mode 0 and round-trips byte-exact.
+
 ## Compressed stream
 
 Pseudo-description of the stream.
@@ -211,7 +215,8 @@ The output buffer must be at least `packed_bc1_maxsize(width, height)` bytes. Th
 ```text
 * 19 static models, in order: for each endpoint (0, 1): red (3 context buckets), green (3 context buckets), blue (3 context buckets) — 16 color models — then top-table reference, difference mask, table difference.
   Each is 3 + num_symbols bytes, no padding (max 259 bytes).
-* Top table size: one uint8_t (number of entries − 1), padded to a 4-byte boundary.
+* Top table size: one uint8_t (number of entries − 1).
+* Flat index mode + value: one uint8_t (mode | value << 1), then padded to a 4-byte boundary. See Block indices → Flat blocks.
 * Tile seed: one uint32_t (little-endian) — the two endpoint colors of the "previous block" used to predict the first block of each tile, averaged over the 64 tiles at compression time. Low 16 bits = color[0], high 16 bits = color[1].
 * Tile predictor flags: 8 bytes (little-endian uint64_t) — one bit per tile: 1 → endpoint colors of that tile are predicted from the row above (same zigzag position) after the first scanline, 0 → zigzag-previous block. See Per-tile predictor.
 * Tile dictionary enable flags: 8 bytes (little-endian uint64_t) — one bit per tile: 1 → the tile's blocks write the color-pair dictionary flag + reference per block, 0 → the six color deltas only (byte-identical to the dictionary-less format). See Per-tile enable flag.
@@ -233,9 +238,9 @@ Within a tile, the blocks are traversed in the zigzag row order. For each block:
   * If 0: the six color deltas, as below.
 * For a tile with the dictionary disabled: no flag.
 * Two endpoints, each: green, red, blue deltas (7-bit values offset by 64).
-* Top-table reference (8-bit index).
-* Difference mask (4 bits).
-* One difference byte per set mask bit.
+* If the block is flat (its two endpoint colors are identical) and the flat index mode is 1:
+  no index payload, the header value is expanded to all 16 indices on decode.
+* Otherwise: top-table reference (8-bit index), difference mask (4 bits), and one difference byte per set mask bit.
 ```
 
 See Color-pair dictionary. The six color deltas, the top-table reference, the mask, and the difference bytes are all encoded with Rice-Golomb using one of the 19 static models; each red/blue uses the model of its (channel, endpoint, context) — the three-bucket green delta described in Endpoints — and each green uses the model of its (endpoint, context), the first endpoint uncontexted and the second contextual on the sign of the first endpoint's green delta. The dictionary flag and reference are raw bits, not Rice-Golomb symbols.
@@ -291,6 +296,10 @@ That is an 8 x 8 matrix = 64 index models per channel. The context is derived fr
 
 The 8-symbol models use a compact rank table (8 entries instead of 256): `3 + num_symbols ≤ 11` bytes each in the stream.
 
+#### Flat blocks
+
+A block whose two endpoints are identical carries a uniform index word: all 16 indices hold the same 3-bit value. A per-channel flat index mode, decided once in the first pass, collapses this: mode 1 is selected for a channel when at least one block of that channel has identical endpoints and every such block holds the same 3-bit value, else mode 0. In mode 1 a flat block stores no index payload at all: the value, stored once per channel for the whole image in the stream header, is expanded to all 16 indices on decode, instead of the 15 contextual symbols; in mode 0, or for a non-flat block, the indices are stored as usual. The mode and the value share one header byte per channel (mode | value << 1) after the tile seeds. A decoder never needs the mode for a non-flat block, so the encoder and decoder always pick the same path; an encoder that leaves a non-uniform or a differing index word on a flat block forces mode 0 and round-trips byte-exact.
+
 ### Compressed stream
 
 As BC1, the stream does not include width, height or format: the BC4 / BC5 decompressors are separate entry points and know their own header.
@@ -299,13 +308,14 @@ As BC1, the stream does not include width, height or format: the BC4 / BC5 decom
 * 2 endpoint models per channel — 256-symbol static models, 3 + num_symbols bytes each, no padding (max 259).
 * 64 index models per channel — 8-symbol static models, 3 + num_symbols bytes each, no padding (max 11).
 * Tile seed: 2 bytes per channel (endpoint 0, endpoint 1).
+* Flat index mode + value: one uint8_t (mode | value << 1) per channel. See Indices → Flat blocks.
 * Padded to a 2-byte boundary.
 * 63 tile offsets (uint16_t each), relative deltas in dword units. Tile 0 always starts the tile data; its offset is zero and not stored.
 * Padded to a 4-byte boundary.
 * 64 tile bitstreams, each 4-byte aligned, each flushed independently.
 ```
 
-For BC5 the red channel's models and seed come first, then the green channel's, then the shared offsets and tiles. Within a tile the blocks follow the zigzag row order; per block the payload is the red channel complete, then the green one, and per channel: endpoint 0, endpoint 1, index 0 (3 raw bits), then indices 1..15.
+For BC5 the red channel's models and seed come first, then the green channel's, then the shared offsets and tiles. Within a tile the blocks follow the zigzag row order; per block the payload is the red channel complete, then the green one, and per channel: endpoint 0, endpoint 1, then either no index payload (a flat block in mode 1, the header value is expanded to all 16 indices on decode) or index 0 (3 raw bits) followed by indices 1..15 (each XORed with its intra-block prediction, as above).
 
 ### Status
 

src/bc_decompress.metal

@@ -25,10 +25,18 @@
 // buffer(2): constant decode_args_t
 //
 // The stream layout parsed here mirrors the CPU compressors / decompressors:
-//   BC1       [19 static models][top table size u8][4-align][strip seed u32][strip predictor u64][strip dict enable u64]
-//             [top table u32s][2-align][63 strip offsets u16, deltas in dword units][4-align][strip bitstreams]
+//   BC1       [19 static models][top table size u8][flat index mode + value u8][4-align][strip seed u32]
+//             [strip predictor u64][strip dict enable u64][top table u32s][2-align][63 strip offsets u16,
+//             deltas in dword units][4-align][strip bitstreams]
 //   BC4 / BC5 [2 endpoint models per channel][64 index models per channel][strip seed, 2 bytes per channel]
-//             [2-align][63 strip offsets u16, deltas in dword units][4-align][strip bitstreams]
+//             [flat index mode + value, 1 byte per channel][2-align][63 strip offsets u16, deltas in dword units]
+//             [4-align][strip bitstreams]
+//
+// The flat index mode (bit 0 of the mode + value byte, the value in the upper bits) is decided once per
+// stream (per channel for BC4 / BC5) in the first pass: 1 when at least one block has identical
+// endpoints and every such block holds the same uniform index value (2-bit for BC1, 3-bit per channel
+// for BC4 / BC5). A block with identical endpoints then stores no index payload: the header value is
+// expanded to all 16 indices, instead of the full index payload.
 //
 // On a corrupted stream the decode runs to completion and the output is garbage, mirroring the
 // CPU version which only reports the failure after the whole strip was decoded. A valid stream
@@ -391,6 +399,11 @@ kernel void bcp_decompress_bc1(const device uint8_t* input [[buffer(0)]],
         return;
 
     uint top_table_size = input[pos++] + 1;
+    if (pos + 1 > length)
+        return;
+    const uint8_t flat_index_byte = input[pos++]; // flat-index mode + value, right after the top-table size
+    const bool flat_index_mode = (flat_index_byte & 1) != 0;
+    const uint8_t flat_index_value = (uint8_t)((flat_index_byte >> 1) & 3);
     pos = (pos + 3) & ~(uint64_t)3; // 4-byte align
 
     if (pos + 20 > length)
@@ -541,16 +554,24 @@ kernel void bcp_decompress_bc1(const device uint8_t* input [[buffer(0)]],
                 }
             }
 
-            // indices difference with top table
-            uint8_t reference = decode_symbol(&compressed_stream, &reference_model);
-            uint8_t mask = decode_symbol(&compressed_stream, &mask_model);
+            // a flat block (identical endpoints, reconstructed exactly like the encoder sees the source) in
+            // mode 1 stores no index payload: the header value is expanded to all 16 indices; otherwise
+            // the full reference + mask + differences
+            uint32_t indices;
+            if ((current_color0 == current_color1) && flat_index_mode)
+                indices = (uint32_t)flat_index_value * 0x11111111u;
+            else
+            {
+                uint8_t reference = decode_symbol(&compressed_stream, &reference_model);
+                uint8_t mask = decode_symbol(&compressed_stream, &mask_model);
 
-            uint32_t difference = 0;
-            for (uint j = 0; j < 4; ++j)
-                if (mask & (1 << j))
-                    difference |= (uint32_t)decode_symbol(&compressed_stream, &difference_model) << (j * 8);
+                uint32_t difference = 0;
+                for (uint j = 0; j < 4; ++j)
+                    if (mask & (1 << j))
+                        difference |= (uint32_t)decode_symbol(&compressed_stream, &difference_model) << (j * 8);
 
-            uint32_t indices = difference ^ top_table[reference];
+                indices = difference ^ top_table[reference];
+            }
 
             write_block(output, block_index, current_color0, current_color1, indices);
 
@@ -640,6 +661,19 @@ void bc45_decode_strip(const device uint8_t* input,
             seed[c][e] = input[pos++];
     }
 
+    // flat-index mode + value: one byte per channel (mode | value << 1), after the seeds, before the 2-align
+    if (pos + NumChannels > length)
+        return;
+
+    bool flat_mode[NumChannels];
+    uint8_t flat_value[NumChannels];
+    for (uint c = 0; c < NumChannels; ++c)
+    {
+        uint8_t flat_byte = input[pos++];
+        flat_mode[c] = (flat_byte & 1) != 0;
+        flat_value[c] = (flat_byte >> 1) & 7;
+    }
+
     //-------------------------------------------------------------------------------------------------------------------------
     // strip offsets (deltas in dword units), then the strip bitstream of this thread
     pos = (pos + 1) & ~(uint64_t)1; // 2-byte align
@@ -721,16 +755,26 @@ void bc45_decode_strip(const device uint8_t* input,
 
                 // index payload: the first index in raw 3 bits, then each index xored with its
                 // intra-block prediction, contextual on the 8 x 8 [left][top] matrix
-                idx[0] = bs_read_bits(&compressed_stream, 3);
-
-                for (uint i = 1; i < 16; ++i)
+                // a flat block (identical endpoints, reconstructed exactly like the encoder sees the
+                // source) in mode 1 stores no index payload: the header value is expanded to all 16
+                if ((e[0] == e[1]) && flat_mode[c])
                 {
-                    // the missing neighbor of the pixel (first row, row start) defaults to 0
-                    uint8_t left = ((i & 3) == 0) ? 0 : idx[i - 1];
-                    uint8_t top = (i < 4) ? 0 : idx[i - 4];
-                    uint8_t p = ((i & 3) == 0) ? top : left;
+                    for (uint i = 0; i < 16; ++i)
+                        idx[i] = flat_value[c];
+                }
+                else
+                {
+                    idx[0] = bs_read_bits(&compressed_stream, 3);
+
+                    for (uint i = 1; i < 16; ++i)
+                    {
+                        // the missing neighbor of the pixel (first row, row start) defaults to 0
+                        uint8_t left = ((i & 3) == 0) ? 0 : idx[i - 1];
+                        uint8_t top = (i < 4) ? 0 : idx[i - 4];
+                        uint8_t p = ((i & 3) == 0) ? top : left;
 
-                    idx[i] = (uint8_t)(p ^ decode_symbol8(&compressed_stream, &index[c][bc4_index_context(left, top)]));
+                        idx[i] = (uint8_t)(p ^ decode_symbol8(&compressed_stream, &index[c][bc4_index_context(left, top)]));
+                    }
                 }
 
                 // the standard block: two endpoints + the 48-bit little-endian index field, as two aligned u32 words