From 14fbcf9c37fe1b4a37317f27da7b291487c5e97a Mon Sep 17 00:00:00 2001 From: seedds Date: Sun, 18 Oct 2026 10:29:30 +0800 Subject: [PATCH] Only GPUs with BC textures are supported: no-BC decompression paths removed Only the latest iPhones are supported now: A17 Pro and later, i.e. iPhone 15 Pro and newer, and M1 or later iPads. All of them read BC textures natively. On older GPUs the decompressed textures didn't fit in memory anyway: Breath of the Wild needs 3.9 GB instead of 2.8 GB. - nvn_mtl.m: BC textures are always created and uploaded as BC. Removed: no_bc(), the GPU bcd kernel for BC1-5, CPU decompression of BC6H/BC7 at upload and in buffer-to-texture copies, and the ASTC to decompressed-format copy. The rule that a BC texture is never a render target or viewed raw stays (GitHub issue #4). - dm/gfx/tex/nvn_bcdec.c/h, dm/third_party/bcdec and DM_TEST_NO_BC are gone. The docs are updated. --- dm/gfx/README.md | 1 - dm/gfx/nvn/nvn_mtl.m | 158 +--- dm/gfx/tex/README.md | 10 +- dm/gfx/tex/nvn_bcdec.c | 56 -- dm/gfx/tex/nvn_bcdec.h | 10 - dm/third_party/README.md | 3 +- dm/third_party/bcdec/LICENSE | 52 -- dm/third_party/bcdec/README.md | 2 - dm/third_party/bcdec/bcdec.h | 1512 -------------------------------- docs/agents/04-graphics.md | 2 +- 10 files changed, 15 insertions(+), 1791 deletions(-) delete mode 100644 dm/gfx/tex/nvn_bcdec.c delete mode 100644 dm/gfx/tex/nvn_bcdec.h delete mode 100644 dm/third_party/bcdec/LICENSE delete mode 100644 dm/third_party/bcdec/README.md delete mode 100644 dm/third_party/bcdec/bcdec.h diff --git a/dm/gfx/README.md b/dm/gfx/README.md index 8d7ac1f..3ce07ed 100644 --- a/dm/gfx/README.md +++ b/dm/gfx/README.md @@ -169,7 +169,6 @@ Comparison and testing: | `DM_NVN_TBDR=0` | Turns off all pass optimizations in the replay layer and takes the plainest path of "all attachments attached, all loaded, all stored", for equivalence comparison | | `DM_NVN_UNWRITTEN_DC=0` | Turns off the single item "attachments not written in this pass are not stored" | | `DM_NVN_FUSE_CHECK=1` | Same-frame check of pass merging; logs the comparison result every time | -| `DM_TEST_NO_BC=1` | Forces the decompression path for "the GPU does not support BC compressed textures" on any machine | | `DM_FPS=30` or `60` | Pins the target frame rate | | `DM_PRESENT_HZ=N` | The frequency of presentation pacing (default 60); 0 means no frame limit and is for performance profiling only | diff --git a/dm/gfx/nvn/nvn_mtl.m b/dm/gfx/nvn/nvn_mtl.m index 613d64e..9df6747 100644 --- a/dm/gfx/nvn/nvn_mtl.m +++ b/dm/gfx/nvn/nvn_mtl.m @@ -17,7 +17,6 @@ #include "../tex/nvn_fmt_mtl.h" #include "../tex/nvn_deswizzle.h" #include "../tex/nvn_deswizzle_gpu.h" -#include "../tex/nvn_bcdec.h" #include "nvn_builtin.h" #include "nvn_dirty.h" #include "nvn_refl.h" @@ -884,19 +883,9 @@ static uint64_t g_pfv_upgrades; // DM_NVN_DIRTYCHECK=1 (diagnostic, compared against ground truth): fully hash the pool content on every use; changed but not detected by dirty-page tracking = a missed detection static int dirtycheck(void) { static int v = -1; if (v < 0) v = getenv("DM_NVN_DIRTYCHECK") != NULL; return v; } static uint64_t full_hash(const uint8_t *m, uint64_t n) { uint64_t h = 1469598103934665603ull; for (uint64_t i = 0; i < n; i++) h = (h ^ m[i]) * 1099511628211ull; return h; } -// The GPU does not support BC compressed formats (A15 / A16 etc.): BC textures are created in the decompressed format and decompressed on the CPU at upload / buffer→texture copy time (nvn_bcdec.c). -// (d2r-mobile converted most of its textures to ASTC offline with a supplementary pack; that tool is not part of this repository, so here every BC texture takes this path). The test switch DM_TEST_NO_BC=1 forces this path on any machine -static int no_bc(void) { - static int v = -1; - if (v < 0) { const char *e = getenv("DM_TEST_NO_BC"); v = (e && *e && strcmp(e, "0")) || ![g_dev supportsBCTextureCompression]; - if (v) MLOG("显卡不支持 BC 压缩格式%s：BC 贴图建成解压格式（BC1–5 在 GPU 上解压，BC6H / BC7 / 3D 在 CPU 上）", e && *e && strcmp(e, "0") ? "（测试开关 DM_TEST_NO_BC）" : ""); } - return v; -} -static int bc_unpack(uint32_t nvn_fmt, uint32_t *pf, int *bpp) { return no_bc() && nvn_bc_unpacked(nvn_fmt, pf, bpp); } static MTLTextureDescriptor *tex_desc(const tex_t *t, int packaged) { uint32_t pf = 0; uint8_t swz[4]; if (!nvn_fmt_to_mtl(t->d.format, &pf, swz) || !pf) { RED_ONCE(0, "贴图格式 %#x Metal 无对应（首例）", t->d.format); return nil; } - int unp = 0; { uint32_t upf; int ubpp; if (bc_unpack(t->d.format, &upf, &ubpp)) { pf = upf; unp = 1; } } // Metal's limits: 16384 in 2D, 2048 in 3D, 2048 array layers; beyond them Metal aborts the process on an assertion, so they are caught here with a red line if (t->d.width > 16384 || t->d.height > 16384 || t->d.depth > 2048 || t->d.levels > 15) { RED_ONCE(71, "贴图尺寸超出 Metal 上限：%u×%u×%u 级 %u 格式 %#x 目标 %d（首例 ser=%llu）", t->d.width, t->d.height, t->d.depth, t->d.levels, t->d.format, t->d.target, (unsigned long long)t->serial); return nil; } MTLTextureDescriptor *d = [MTLTextureDescriptor new]; @@ -921,8 +910,7 @@ static MTLTextureDescriptor *tex_desc(const tex_t *t, int packaged) { } d.storageMode = MTLStorageModePrivate; // Render targets and assets both go in GPU private storage (lossless compression available) d.usage = MTLTextureUsageShaderRead | (fmt_compressed(t->d.format) || packaged ? 0 : MTLTextureUsageRenderTarget) - | (t->d.flags & NVN_TEXTURE_FLAGS_IMAGE ? MTLTextureUsageShaderWrite : 0) // Image usage: shader-writable (compute kernels generating the 3D color lookup table, etc.) - | (unp ? MTLTextureUsageShaderWrite | MTLTextureUsagePixelFormatView : 0); // No BC: textures in the decompressed format are written by the GPU decompression kernel (up_dec) / receive whole-block copies from ASTC sources (copy_t2t_astc) + | (t->d.flags & NVN_TEXTURE_FLAGS_IMAGE ? MTLTextureUsageShaderWrite : 0); // Image usage: shader-writable (compute kernels generating the 3D color lookup table, etc. return d; } // In-place upgrade: when the original texture was not created with PixelFormatView (off by default, to keep Apple GPU lossless compression) but a different-format view is needed, create a new one with the original spec + PixelFormatView, copy over every slice × level, and replace the registration. @@ -973,45 +961,6 @@ typedef struct { NSUInteger slice, level; } up_blit; static up_blit *g_up_blits; static int g_nup_blits, g_up_blitcap; -// Without BC, BC1–5 are decompressed on the GPU (A16 logs: CPU deswizzle + CPU decompression took over 300 ms in a single frame): the GPU deswizzle writes the compressed blocks into the staging area, and at flush time, after the barrier, -// the compute kernel bcd decodes them into the texture one thread per block (the target is a 2D view in the decompressed format; an sRGB base texture is written with raw values through an RGBA8Unorm view). BC6H / BC7 / 3D textures are very few and still go through the CPU (nvn_bcdec.c) -typedef struct { __unsafe_unretained id src, params; size_t off, poff; uint32_t bw, bh; __unsafe_unretained id view; } up_dec; // view: held by g_up_texs within the batch, and afterwards by defer_obj_unref until the GPU finishes -static up_dec *g_up_decs; static int g_nup_decs, g_up_deccap; -static const char k_bcd_msl[] = -"#include \nusing namespace metal;\n" -"struct P { uint off, pitch, w, h, fmt; };\n" // fmt: 1 BC1, 2 BC2, 3 BC3, 4 BC4, 5 BC4 signed, 6 BC5, 7 BC5 signed -"static float3 c565(uint c) { return float3((c >> 11) & 31, (c >> 5) & 63, c & 31) / float3(31, 63, 31); }\n" -"static void col(device const uchar *b, thread float4 *o, bool four) {\n" -" uint c0 = b[0] | (b[1] << 8), c1 = b[2] | (b[3] << 8), ix = b[4] | (b[5] << 8) | (b[6] << 16) | (b[7] << 24); float3 e0 = c565(c0), e1 = c565(c1); float4 p[4];\n" -" p[0] = float4(e0, 1); p[1] = float4(e1, 1);\n" -" if (four || c0 > c1) { p[2] = float4((2 * e0 + e1) / 3, 1); p[3] = float4((e0 + 2 * e1) / 3, 1); } else { p[2] = float4((e0 + e1) / 2, 1); p[3] = float4(0); }\n" -" for (int i = 0; i < 16; i++) o[i] = p[(ix >> (2 * i)) & 3]; }\n" -"static void a8(device const uchar *b, thread float *o, bool sn) {\n" // Single-channel block of BC4 / BC3 alpha / BC5 -" float a0 = sn ? max(float(as_type(b[0])), -127.0) / 127 : b[0] / 255.0, a1 = sn ? max(float(as_type(b[1])), -127.0) / 127 : b[1] / 255.0;\n" -" bool eight = sn ? as_type(b[0]) > as_type(b[1]) : b[0] > b[1]; float lo = sn ? -1 : 0; float p[8]; p[0] = a0; p[1] = a1;\n" -" for (int k = 2; k < 8; k++) p[k] = eight ? ((8 - k) * a0 + (k - 1) * a1) / 7 : (k < 6 ? ((6 - k) * a0 + (k - 1) * a1) / 5 : (k == 6 ? lo : 1));\n" -" uint lo24 = b[2] | (b[3] << 8) | (b[4] << 16), hi24 = b[5] | (b[6] << 8) | (b[7] << 16);\n" -" for (int i = 0; i < 16; i++) o[i] = p[i < 8 ? (lo24 >> (3 * i)) & 7 : (hi24 >> (3 * (i - 8))) & 7]; }\n" -"kernel void bcd(device const uchar *src [[buffer(0)]], device const P &p [[buffer(1)]], texture2d dst [[texture(0)]], uint2 g [[thread_position_in_grid]]) {\n" -" uint nbx = (p.w + 3) / 4, nby = (p.h + 3) / 4; if (g.x >= nbx || g.y >= nby) return;\n" -" uint bb = (p.fmt == 1 || p.fmt == 4 || p.fmt == 5) ? 8 : 16; device const uchar *b = src + p.off + g.y * p.pitch + g.x * bb; float4 o[16]; float a[16], c[16];\n" -" if (p.fmt == 1) col(b, o, false);\n" -" else if (p.fmt == 2) { col(b + 8, o, true); for (int i = 0; i < 16; i++) o[i].a = ((b[i / 2] >> (4 * (i & 1))) & 15) / 15.0; }\n" -" else if (p.fmt == 3) { col(b + 8, o, true); a8(b, a, false); for (int i = 0; i < 16; i++) o[i].a = a[i]; }\n" -" else if (p.fmt == 4 || p.fmt == 5) { a8(b, a, p.fmt == 5); for (int i = 0; i < 16; i++) o[i] = float4(a[i], 0, 0, 1); }\n" -" else { a8(b, a, p.fmt == 7); a8(b + 8, c, p.fmt == 7); for (int i = 0; i < 16; i++) o[i] = float4(a[i], c[i], 0, 1); }\n" -" for (int i = 0; i < 16; i++) { uint2 q = g * 4 + uint2(i & 3, i >> 2); if (q.x < p.w && q.y < p.h) dst.write(o[i], q); } }\n"; -static int bcd_fmt(uint32_t f) { // NVN format → bcd fmt; 0 = not decoded on the GPU (BC6H / BC7 / non-BC) - switch (f) { case 0x41: case 0x42: case 0x45: case 0x46: return 1; case 0x43: case 0x47: return 2; case 0x44: case 0x48: return 3; - case 0x49: return 4; case 0x4a: return 5; case 0x4b: return 6; case 0x4c: return 7; default: return 0; } -} -static id bcd_pso(void) { - static id ps; static int failed; - if (!ps && !failed) { NSError *e = nil; id lib = [g_dev newLibraryWithSource:@(k_bcd_msl) options:nil error:&e]; id fn = lib ? [lib newFunctionWithName:@"bcd"] : nil; - ps = fn ? [g_dev newComputePipelineStateWithFunction:fn error:&e] : nil; - if (!ps) { failed = 1; MLOG("【红】无 BC：GPU 解压管线建不出 %s（退回 CPU 解压）", e.localizedDescription.UTF8String); } else MLOG("无 BC：BC1–5 改在 GPU 上解压"); } - return ps; -} static NSMutableArray *g_up_texs; // Target textures in the batch: before the flush, textures may be replaced / released within the same submission static id g_up_block; static size_t g_up_block_off, g_up_block_size; @@ -1032,18 +981,12 @@ static void up_flush(void) { if (g_up_ce) { m4_end(g_up_ce); g_up_ce = nil; } - id b = g_nup_blits || g_nup_decs ? g_up_ce = [g_up_cb computeCommandEncoder] : nil; + id b = g_nup_blits ? g_up_ce = [g_up_cb computeCommandEncoder] : nil; if (b) [b barrierAfterQueueStages:MTLStageAll beforeStages:MTLStageAll visibilityOptions:MTL4VisibilityOptionDevice]; for (int i = 0; b && i < g_nup_blits; i++) { up_blit *x = &g_up_blits[i]; enc_b_copy_buffer_texture(b, x->src, x->off, x->row, x->image, x->size, x->dst, x->slice, x->level, MTLOriginMake(0, 0, 0)); } - if (b && g_nup_decs) { // GPU decompression without BC (see up_dec): offsets in the parameter block are relative to the start of the staging block - [b setArgumentTable:g_up_ct]; [b setComputePipelineState:bcd_pso()]; - for (int i = 0; i < g_nup_decs; i++) { up_dec *x = &g_up_decs[i]; - [g_up_ct setAddress:x->src.gpuAddress atIndex:0]; [g_up_ct setAddress:x->params.gpuAddress + x->poff atIndex:1]; [g_up_ct setTexture:x->view.gpuResourceID atIndex:0]; - [b dispatchThreadgroups:MTLSizeMake((x->bw + 7) / 8, (x->bh + 7) / 8, 1) threadsPerThreadgroup:MTLSizeMake(8, 8, 1)]; } - } if (b) { m4_end(b); g_up_ce = nil; } id cb = g_up_cb; unsigned slot = (g_up_seq - 1) % 3; @@ -1071,7 +1014,7 @@ static void up_flush(void) { for (id obj in g_up_keep) defer_obj_release(obj); [g_up_keep removeAllObjects]; up_diag(); - g_nup_blits = 0; g_nup_decs = 0; + g_nup_blits = 0; [g_up_texs removeAllObjects]; } @@ -1170,40 +1113,27 @@ static id pkg_upload_(const tex_t *t, uint64_t *stamp) { int bw, bh, bpb = nvn_fmt_bpb(t->d.format, &bw, &bh); unsigned layers = (t->d.target == 4 || t->d.target == 9) ? t->d.depth : t->d.target == 8 ? 6 : t->d.target == 3 ? t->d.height : 1; if (t->d.target == 9) layers *= 6; - uint32_t upf; int ubpp = 0, unp = bc_unpack(t->d.format, &upf, &ubpp); // No BC: decompress on the CPU after deswizzling, with the staging area sized for the decompressed data size_t total = 0; for (unsigned l = 0; l < d.mipmapLevelCount; l++) { uint32_t w = MAX(t->d.width >> l, 1u), h = MAX(t->d.height >> l, 1u), z = t->d.target == 2 ? MAX(t->d.depth >> l, 1u) : 1; - total += unp ? (size_t)w * ubpp * h * z * layers : (size_t)((w + bw - 1) / bw) * bpb * ((h + bh - 1) / bh) * z * layers; + total += (size_t)((w + bw - 1) / bw) * bpb * ((h + bh - 1) / bh) * z * layers; } g_bp_up_bytes += total; void *mb; uint64_t mo; - int gdec = unp && t->d.target != 2 && bcd_fmt(t->d.format) && bcd_pso(); // BC1–5 without BC: GPU deswizzle + GPU decompression - size_t ctotal = 0; // Staging size of the compressed blocks (with GPU decompression the staging area holds compressed blocks) - if (gdec) for (unsigned l = 0; l < d.mipmapLevelCount; l++) { uint32_t w = MAX(t->d.width >> l, 1u), h = MAX(t->d.height >> l, 1u); ctotal += (size_t)((w + bw - 1) / bw) * bpb * ((h + bh - 1) / bh) * layers; } - int use_gpu = (!unp || gdec) && g_nvn_deswz_gpu && mtl_addr((uint64_t)t->mem, t->L.size, &mb, &mo); - if (g_nvn_deswz_gpu && !use_gpu && !unp) RED_ONCE(40, "预排布贴图存储不在内存池里，回退 CPU"); + int use_gpu = g_nvn_deswz_gpu && mtl_addr((uint64_t)t->mem, t->L.size, &mb, &mo); + if (g_nvn_deswz_gpu && !use_gpu) RED_ONCE(40, "预排布贴图存储不在内存池里，回退 CPU"); if (use_gpu) { // The GPU reads the memory pool at execution time, which is closer to a real Switch than a snapshot at CPU upload time; the dirty-page stamp is still taken before encoding, and later writes still trigger a re-upload. up_begin(); size_t stg_off; - id stg = up_alloc(gdec ? ctotal : total ? total : 16, &stg_off); - size_t off = 0; int enc_ok = g_up_ce != nil, nb0 = g_nup_blits, nd0 = g_nup_decs; + id stg = up_alloc(total ? total : 16, &stg_off); + size_t off = 0; int enc_ok = g_up_ce != nil, nb0 = g_nup_blits; for (unsigned l = 0; enc_ok && l < d.mipmapLevelCount; l++) { uint32_t w = MAX(t->d.width >> l, 1u), h = MAX(t->d.height >> l, 1u), z = t->d.target == 2 ? MAX(t->d.depth >> l, 1u) : 1; size_t pitch = (size_t)((w + bw - 1) / bw) * bpb, rows = (h + bh - 1) / bh; for (unsigned y = 0; y < layers; y++) { size_t poff; id params = up_param_alloc(&poff); // Metal 4 has no setBytes: the parameter block of each deswizzle is cut from a shared block (no new GPU buffer each time) if (nvn_deswz_gpu_encode_m4(g_up_ce, g_up_ct, params, poff, (__bridge id)mb, mo, &t->d, l, y, stg, stg_off + off, pitch)) { enc_ok = 0; break; } - if (gdec) { // Decompression job: dispatched at flush time after the barrier (see up_flush) - MTLPixelFormat vf = mt.pixelFormat == MTLPixelFormatRGBA8Unorm_sRGB ? MTLPixelFormatRGBA8Unorm : mt.pixelFormat; - id v = [mt newTextureViewWithPixelFormat:vf textureType:MTLTextureType2D levels:NSMakeRange(l, 1) slices:NSMakeRange(y, 1)]; - size_t qoff; id q = up_param_alloc(&qoff); - if (!v || !q) { enc_ok = 0; break; } - uint32_t *pp = (uint32_t *)((uint8_t *)q.contents + qoff); pp[0] = (uint32_t)(stg_off + off); pp[1] = (uint32_t)pitch; pp[2] = w; pp[3] = h; pp[4] = (uint32_t)bcd_fmt(t->d.format); - if (g_nup_decs == g_up_deccap) { g_up_deccap = g_up_deccap ? g_up_deccap * 2 : 256; g_up_decs = realloc(g_up_decs, g_up_deccap * sizeof *g_up_decs); } - g_up_decs[g_nup_decs++] = (up_dec){stg, q, 0, qoff, (w + 3) / 4, (h + 3) / 4, v}; [g_up_texs addObject:v]; defer_obj_unref(v); // View: released after this submission finishes running, residency not revoked (what is resident is the base texture) - } else up_blit_add(stg, stg_off + off, pitch, pitch * rows, MTLSizeMake(w, h, z), mt, y, l); off += pitch * rows * z; } @@ -1213,7 +1143,7 @@ static id pkg_upload_(const tex_t *t, uint64_t *stamp) { if (!g_cb) up_flush(); return mt; } - g_nup_blits = nb0; g_nup_decs = nd0; // Roll back the copies / decompressions already recorded for this texture: the already-encoded deswizzle only writes the staging area and is harmless; this texture takes the CPU path below instead + g_nup_blits = nb0; // Roll back the copies already recorded for this texture: the already-encoded deswizzle only writes the staging area and is harmless; this texture takes the CPU path below instead if (!g_cb) up_flush(); RED_ONCE(41, "预排布贴图 GPU 解排布不支持格式 %#x 目标 %d，回退 CPU", t->d.format, t->d.target); } @@ -1226,15 +1156,6 @@ static id pkg_upload_(const tex_t *t, uint64_t *stamp) { uint32_t w = MAX(t->d.width >> l, 1u), h = MAX(t->d.height >> l, 1u), z = t->d.target == 2 ? MAX(t->d.depth >> l, 1u) : 1; size_t pitch = (size_t)((w + bw - 1) / bw) * bpb, rows = (h + bh - 1) / bh; for (unsigned y = 0; y < layers; y++) { - if (unp) { // No BC: first deswizzle into a temporary area (compressed blocks), then decompress layer by layer into the staging area - uint8_t *cb = malloc(pitch * rows * z); size_t up = (size_t)w * ubpp; - if (!cb || nvn_deswizzle(&t->d, l, y, t->mem, cb, pitch)) { free(cb); RED_ONCE(9, "解排布失败 格式 %#x 目标 %d", t->d.format, t->d.target); break; } - for (uint32_t zi = 0; zi < z; zi++) nvn_bc_decode(t->d.format, cb + zi * pitch * rows, pitch, w, h, (uint8_t *)stg.contents + off + zi * up * h, up); - free(cb); - up_blit_add(stg, off, up, up * h, MTLSizeMake(w, h, z), mt, y, l); - off += up * h * z; - continue; - } if (nvn_deswizzle(&t->d, l, y, t->mem, (uint8_t *)stg.contents + off, pitch)) { RED_ONCE(9, "解排布失败 格式 %#x 目标 %d", t->d.format, t->d.target); break; } up_blit_add(stg, off, pitch, pitch * rows, MTLSizeMake(w, h, z), mt, y, l); off += pitch * rows * z; @@ -1818,10 +1739,8 @@ void pg_commit_ext(id cb) { pg_commit(cb); } // A compressed texture as a render target (The Legend of Zelda: Breath of the Wild: a 32x32 RGTC1 (BC4) texture written through an RGBA16UI view, one texel = one 8-byte block, i.e. BC4 encoded on the GPU): // Metal can neither attach a compressed format as a colour attachment nor create compressed <-> uncompressed views (the pipeline descriptor validation aborts the process). For now the attachment is left out (a pass with no attachment at all skips its draws), with a red line; // drawing it correctly needs an uncompressed proxy texture with the same block size, copied into the compressed texture through a buffer after the pass -// On a GPU without BC the same texture is stored decompressed (R8 for BC4), so the attachment check must go by the NVN format, not the Metal one: attaching it through the RGBA16UI view -// asked Metal for an R8Unorm -> RGBA16Uint view, which aborts the process (GitHub issue #4, iPhone 14 Pro; DM_TEST_NO_BC=1 reproduces it on the Mac) static int mtl_fmt_block(MTLPixelFormat f) { return (f >= MTLPixelFormatBC1_RGBA && f <= MTLPixelFormatBC7_RGBAUnorm_sRGB) || (f >= MTLPixelFormatPVRTC_RGB_2BPP && f < MTLPixelFormatGBGR422); } -static int nvn_fmt_bc(uint32_t f) { uint32_t pf; int bpp; return nvn_bc_unpacked(f, &pf, &bpp); } // NVN BC1-7 / BC6H formats, whether or not this GPU stores them compressed +static int nvn_fmt_bc(uint32_t f) { return f >= 0x41 && f <= 0x50; } // NVN BC1–7 / BC6H formats (0x41 RGB_DXT1 … 0x50 BC6H unsigned) static MTLPixelFormat rt_view_fmt(id base, const view_t *v, int has_v) { if (base && mtl_fmt_block(base.pixelFormat)) { RED_ONCE(78, "渲染目标是压缩格式贴图（Metal 格式 %lu %lux%lu 层 %lu，视图换格式 %d → NVN %#x）：不挂，该附件的绘制落空", (unsigned long)base.pixelFormat, (unsigned long)base.width, (unsigned long)base.height, @@ -1851,11 +1770,6 @@ static int rt_dup_slot(const p_set_rt *r, int i) { static void attach(MTLRenderPassAttachmentDescriptor *a, const tex_t *t, const view_t *v, int has_v) { id base = tex_mtl(t); a.texture = base; if (base && mtl_fmt_block(base.pixelFormat)) { rt_view_fmt(base, v, has_v); a.texture = nil; return; } // compressed formats cannot be attached (see rt_view_fmt) - if (base && nvn_fmt_bc(t->d.format)) { // the same texture stored decompressed on a GPU without BC: not attached either (see above rt_view_fmt) - RED_ONCE(78, "渲染目标是压缩格式贴图（无 BC，按解压格式存：Metal 格式 %lu %lux%lu，视图换格式 %d → NVN %#x）：不挂，该附件的绘制落空", (unsigned long)base.pixelFormat, (unsigned long)base.width, (unsigned long)base.height, - has_v ? v->set_format : 0, has_v ? v->format : 0); - a.texture = nil; return; - } if (has_v && v->set_levels) a.level = v->base_level; if (has_v && v->set_layers) { if (t->d.target == 2) a.depthPlane = v->min_layer; else a.slice = v->min_layer; } // A view that changes the format is attached as a Metal view of that format (pipeline formats follow the attachment). SMBW's last pass writes the RGBA8 swapchain texture through an RGBA8_SRGB view, so the hardware encodes sRGB on write; @@ -2458,8 +2372,7 @@ static id tex_by_id(int tid) { // Caller holds g_rmu if (!r.has_v) memset(&r.v, 0, sizeof r.v); uint32_t pf = (uint32_t)base.pixelFormat; uint8_t swz[4]; if (r.v.set_format && !(nvn_fmt_to_mtl(r.v.format, &pf, swz) && pf)) pf = (uint32_t)base.pixelFormat; - { uint32_t upf; int ubpp; if (r.v.set_format && bc_unpack(r.v.format, &upf, &ubpp)) pf = upf; } // No BC: a BC view is switched to the decompressed format just like the base texture - if (r.v.set_format && nvn_fmt_bc(rp->t.d.format) && !nvn_fmt_bc(r.v.format)) { // the raw blocks of a BC texture through an uncompressed view (see rt_view_fmt): Metal has no such view, with or without BC + if (r.v.set_format && nvn_fmt_bc(rp->t.d.format) && !nvn_fmt_bc(r.v.format)) { // the raw blocks of a BC texture through an uncompressed view (see rt_view_fmt): Metal has no such view RED_ONCE(86, "压缩格式贴图的原始块视图（NVN 格式 %#x → 视图 %#x）Metal 建不出：按空槽", rp->t.d.format, r.v.format); return nil; } @@ -5255,20 +5168,6 @@ static void copy_b2t(const p_copy_b2t *c) { int is3d = c->tex->d.target == 2, nz = c->r.depth ? c->r.depth : 1; void *mb; uint64_t mo; if (!mtl_addr(c->src, img * nz, &mb, &mo)) { RED_ONCE(23, "缓冲→贴图拷贝的源不在内存池里"); return; } - uint32_t upf; int ubpp; - if (bc_unpack(c->tex->d.format, &upf, &ubpp)) { // No BC: the compressed blocks in the source are decompressed on the CPU into a temporary buffer, then copied pixel-wise into the decompressed-format texture - const uint8_t *src = [(__bridge id)mb contents]; - if (!src) { RED_ONCE(46, "无 BC 解压：缓冲→贴图拷贝的源缓冲 CPU 读不到"); return; } - size_t up = (size_t)c->r.width * ubpp, uimg = up * c->r.height; - id tb = [g_dev newBufferWithLength:uimg * nz ?: 16 options:MTLResourceStorageModeShared]; - for (int z = 0; z < nz; z++) nvn_bc_decode(c->tex->d.format, src + mo + z * img, row, c->r.width, c->r.height, (uint8_t *)tb.contents + z * uimg, up); - res_add(tb); - id b = blit(); - if (is3d) enc_b_copy_buffer_texture(b, tb, 0, up, uimg, MTLSizeMake(c->r.width, c->r.height, nz), t, 0, level, MTLOriginMake(c->r.xoffset, c->r.yoffset, c->r.zoffset)); - else for (int z = 0; z < nz; z++) enc_b_copy_buffer_texture(b, tb, z * uimg, up, uimg, MTLSizeMake(c->r.width, c->r.height, 1), t, layer0 + c->r.zoffset + z, level, MTLOriginMake(c->r.xoffset, c->r.yoffset, 0)); - defer_obj_release(tb); // Release after this submit has finished running (and remove from residency) - return; - } id b = blit(); if (is3d) enc_b_copy_buffer_texture(b, (__bridge id)mb, mo, row, img, MTLSizeMake(c->r.width, c->r.height, nz), t, 0, level, MTLOriginMake(c->r.xoffset, c->r.yoffset, c->r.zoffset)); @@ -5277,35 +5176,6 @@ static void copy_b2t(const p_copy_b2t *c) { enc_b_copy_buffer_texture(b, (__bridge id)mb, mo + z * img, row, img, MTLSizeMake(c->r.width, c->r.height, 1), t, layer0 + c->r.zoffset + z, level, MTLOriginMake(c->r.xoffset, c->r.yoffset, 0)); } -// No BC + Legacy Device Pack: the game copies textures from files (already ASTC 4x4) whole into a BC3 texture array it creates itself (created as RGBA8 when there is no BC). On Switch the two formats have the same block size and -// the copy moves raw blocks; here it is changed to a compute kernel that decodes pixel by pixel and writes them in, with sRGB sources first read out as linear values and then encoded back to sRGB values, so the result = the values read in the destination format after moving the blocks as-is -static const char k_a2u_msl[] = -"#include \nusing namespace metal;\n" -"static float3 enc(float3 c) { return select(1.055 * pow(c, 1.0 / 2.4) - 0.055, c * 12.92, c <= 0.0031308); }\n" -"kernel void a2u(texture2d src [[texture(0)]], texture2d dst [[texture(1)]], constant uint4 *p [[buffer(0)]], uint2 g [[thread_position_in_grid]]) {\n" -" if (g.x >= p[1].x || g.y >= p[1].y) return; float4 v = src.read(g + p[0].xy); if (p[1].z) v.rgb = enc(v.rgb); dst.write(v, g + p[0].zw); }\n"; -static int copy_t2t_astc(id s, id d, unsigned sl, unsigned ss, unsigned dl, unsigned ds, int sx, int sy, int dx, int dy, int w, int h) { - static int failed; static id ps; static id at; - if (failed) return 0; - if (!ps) { NSError *e = nil; id lib = [g_dev newLibraryWithSource:@(k_a2u_msl) options:nil error:&e]; id f = lib ? [lib newFunctionWithName:@"a2u"] : nil; - ps = f ? [g_dev newComputePipelineStateWithFunction:f error:&e] : nil; - MTL4ArgumentTableDescriptor *ad = [MTL4ArgumentTableDescriptor new]; ad.maxTextureBindCount = 2; ad.maxBufferBindCount = 1; at = [g_dev newArgumentTableWithDescriptor:ad error:&e]; - if (!ps || !at) { failed = 1; MLOG("【红】无 BC：ASTC → 解压格式拷贝的计算管线建不出 %s", e.localizedDescription.UTF8String); return 0; } } - MTLTextureType tt = MTLTextureType2D; - id sv = [s newTextureViewWithPixelFormat:s.pixelFormat textureType:tt levels:NSMakeRange(sl, 1) slices:NSMakeRange(ss, 1)]; - id dv = [d newTextureViewWithPixelFormat:MTLPixelFormatRGBA8Unorm textureType:tt levels:NSMakeRange(dl, 1) slices:NSMakeRange(ds, 1)]; - id pb = [g_dev newBufferWithLength:32 options:MTLResourceStorageModeShared]; - if (!sv || !dv || !pb) { RED_ONCE(47, "无 BC：ASTC → 解压格式拷贝建不出视图 / 参数块"); return 0; } - uint32_t *p = pb.contents; p[0] = sx; p[1] = sy; p[2] = dx; p[3] = dy; p[4] = w; p[5] = h; p[6] = s.pixelFormat == MTLPixelFormatASTC_4x4_sRGB; p[7] = 0; - res_add(pb); - rp_end(); if (g_blit) { enc_b_end(g_blit); g_blit = nil; } - id ce = (id)enc_c_begin(NULL); [ce setArgumentTable:at]; - [at setTexture:sv.gpuResourceID atIndex:0]; [at setTexture:dv.gpuResourceID atIndex:1]; [at setAddress:pb.gpuAddress atIndex:0]; - enc_c_pso(ce, ps); enc_c_dispatch(ce, MTLSizeMake((w + 7) / 8, (h + 7) / 8, 1), MTLSizeMake(8, 8, 1)); enc_c_end(ce); - args_dirty_all(); - defer_obj_unref(sv); defer_obj_unref(dv); defer_obj_release(pb); // Release after this submit has finished running (views are not made resident separately; the parameter block is removed from residency) - return 1; -} // Texture copies a blit cannot do -- a format change, a scaled region or a mirrored one (MIRROR_X / MIRROR_Y) (Link's Awakening converts its 1408x792 RG11B10F scene to RGBA16F every frame; Super Mario 3D World scales 1920x1080 to 1280x720 // when leaving the title): one render pass per layer that draws the source region into the destination region, nearest or (LINEAR_FILTER) bilinear. Float, unorm and sRGB formats only, 2D / 2D array, // destinations that can be render targets; anything else returns 0 and keeps the red line. Depth (Depth32Float / Depth32Float_Stencil8 into the same format) is drawn by writing depth from the fragment shader, @@ -5437,12 +5307,6 @@ static void copy_t2t(const p_copy_t2t *c) { PD("t src=%llu lv%u sl%u xyz=%d,%d,%d dst=%llu lv%u sl%u xyz=%d,%d,%d size=%d,%d,%d", c->src->serial, sl, sz0, c->sr.xoffset, c->sr.yoffset, c->sr.zoffset, c->dst->serial, dl, dz0, c->dr.xoffset, c->dr.yoffset, c->dr.zoffset, c->sr.width, c->sr.height, c->sr.depth); int s3 = c->src->d.target == 2, d3 = c->dst->d.target == 2, nz = c->sr.depth ? c->sr.depth : 1; - if (no_bc() && (s.pixelFormat == MTLPixelFormatASTC_4x4_LDR || s.pixelFormat == MTLPixelFormatASTC_4x4_sRGB) && (d.pixelFormat == MTLPixelFormatRGBA8Unorm || d.pixelFormat == MTLPixelFormatRGBA8Unorm_sRGB) - && !s3 && !d3 && c->sr.width == c->dr.width && c->sr.height == c->dr.height && (d.usage & MTLTextureUsageShaderWrite)) { // No BC: ASTC source -> decompressed-format destination (see copy_t2t_astc) - int ok = 1; - for (int z = 0; z < nz && ok; z++) ok = copy_t2t_astc(s, d, sl, sz0 + c->sr.zoffset + z, dl, dz0 + c->dr.zoffset + z, c->sr.xoffset, c->sr.yoffset, c->dr.xoffset, c->dr.yoffset, c->sr.width, c->sr.height); - if (ok) return; - } if ((c->sr.width != c->dr.width || c->sr.height != c->dr.height || s.pixelFormat != d.pixelFormat || (c->flags & (NVN_COPY_FLAGS_MIRROR_X_BIT | NVN_COPY_FLAGS_MIRROR_Y_BIT))) && !s3 && !d3 && (c->sr.depth ? c->sr.depth : 1) == (c->dr.depth ? c->dr.depth : 1) && !(c->flags & ~(NVN_COPY_FLAGS_ENGINE_2D_BIT | NVN_COPY_FLAGS_LINEAR_FILTER_BIT | NVN_COPY_FLAGS_MIRROR_X_BIT | NVN_COPY_FLAGS_MIRROR_Y_BIT)) diff --git a/dm/gfx/tex/README.md b/dm/gfx/tex/README.md index 3629c8c..641a071 100644 --- a/dm/gfx/tex/README.md +++ b/dm/gfx/tex/README.md @@ -15,19 +15,13 @@ Block linear is how the Tegra X1 GPU stores textures: the data is not one row af | `nvn_fmt_mtl.h`, `nvn_fmt_mtl.c` | `nvn_fmt_to_mtl()`: NVN format (0x01–0x9C) → Metal pixel format + component remapping. Returns 1 for an exact match, 2 for a substitute (precision, bit order or interpretation differs and a conversion is needed), 0 when Metal has no corresponding format. How each format is handled is written in the comments inside the table. Pure C, no dependency on Metal headers | | `nvn_deswizzle.h`, `nvn_deswizzle.c` | `nvn_deswizzle()`: deswizzles one level and one layer from block linear to linear; `nvn_swizzle()`: the reverse, used for round trips in unit tests; `nvn_bl_offset()`: the block-linear address of a single point. Pure C | | `nvn_deswizzle_gpu.h`, `nvn_deswizzle_gpu.m` | A Metal compute shader version of the same deswizzling algorithm, used by the replay layer when uploading packaged textures; when initialization fails or an unsupported format is encountered, the replay layer falls back to the CPU version | -| `nvn_bcdec.h`, `nvn_bcdec.c` | `nvn_bc_unpacked()` / `nvn_bc_decode()`: decompress BC textures on the CPU into uncompressed formats; decoding uses `dm/third_party/bcdec` | | `tests/` | Unit tests | The offsets of each level and layer within the storage, and the block height and block depth actually used by each level, are all taken from `../nvn/nvn_tex.c` (`nvn_tex_layout`, `nvn_tex_level_offset`, `nvn_tex_level_tiles`); this directory only does addressing and moving within a level. The game's packaged textures carry their own layout parameters, whose byte 0 is the log2 of the level 0 block height; when it is present it takes precedence (the `pkg_block_h` field of `nvn_texdesc`). -## Devices without BC support +## BC textures -Some older Apple GPUs do not support BC compressed textures (determined by `supportsBCTextureCompression` of `MTLDevice`). On these devices: - -- BC textures are created by the replay layer in their decompressed format and decompressed at upload: BC1–5 on the GPU (the `bcd` kernel in `../nvn/nvn_mtl.m`), BC6H, BC7 and 3D textures on the CPU through `nvn_bcdec.c` in this directory. Decompressed textures take 2–8 times the memory (Breath of the Wild: about 3.9 GB instead of 2.8 GB, which iOS does not allow on a 6 GB iPhone without the increased-memory-limit entitlement); d2r-mobile's offline ASTC pack (`astc_pack.py`) is not part of this repository; -- a BC texture is never a render target and its raw blocks are not sampled through an uncompressed view, with or without BC (Metal has no such views; Breath of the Wild encodes one 32×32 BC4 texture on the GPU, GitHub issue #4). - -The environment variable `DM_TEST_NO_BC=1` forces this path on any machine. +Only GPUs with BC compressed textures are supported (`supportsBCTextureCompression` of `MTLDevice`: Apple silicon Macs and the latest iPhones), so BC textures are created and uploaded as BC; there is no decompression path. A BC texture is never a render target and its raw blocks are not sampled through an uncompressed view (Metal has no such views; Breath of the Wild encodes one 32×32 BC4 texture on the GPU, GitHub issue #4). ## Tests diff --git a/dm/gfx/tex/nvn_bcdec.c b/dm/gfx/tex/nvn_bcdec.c deleted file mode 100644 index 6f1cb09..0000000 --- a/dm/gfx/tex/nvn_bcdec.c +++ /dev/null @@ -1,56 +0,0 @@ -// See nvn_bcdec.h. Decoding uses bcdec (dm/third_party/bcdec, single header, MIT / public domain). -#define BCDEC_IMPLEMENTATION -#define BCDEC_BC4BC5_PRECISE // interface with signed variants of BC4 / BC5 -#include "../../third_party/bcdec/bcdec.h" -#include "nvn_bcdec.h" -#include - -enum { R8Unorm = 10, R8Snorm = 12, RG8Unorm = 30, RG8Snorm = 32, RGBA8Unorm = 70, RGBA8Unorm_sRGB = 71, RGBA16Float = 115 }; -// NVN: 0x41 RGB_DXT1, 0x42 RGBA_DXT1, 0x43 DXT3, 0x44 DXT5, 0x45–0x48 the same in sRGB, 0x49 / 0x4a BC4 unsigned / signed, 0x4b / 0x4c BC5, 0x4d / 0x4e BC7 (sRGB), 0x4f / 0x50 BC6H signed / unsigned -int nvn_bc_unpacked(uint32_t f, uint32_t *pf, int *bpp) { - switch (f) { - case 0x41: case 0x42: case 0x43: case 0x44: case 0x4d: *pf = RGBA8Unorm; *bpp = 4; return 1; - case 0x45: case 0x46: case 0x47: case 0x48: case 0x4e: *pf = RGBA8Unorm_sRGB; *bpp = 4; return 1; - case 0x49: *pf = R8Unorm; *bpp = 1; return 1; - case 0x4a: *pf = R8Snorm; *bpp = 1; return 1; - case 0x4b: *pf = RG8Unorm; *bpp = 2; return 1; - case 0x4c: *pf = RG8Snorm; *bpp = 2; return 1; - case 0x4f: case 0x50: *pf = RGBA16Float; *bpp = 8; return 1; - default: return 0; - } -} - -static void block(uint32_t f, const uint8_t *b, uint8_t *o, int op) { // one block → 4×4 pixels (o has op bytes per row) - switch (f) { - case 0x41: case 0x42: case 0x45: case 0x46: bcdec_bc1(b, o, op); break; - case 0x43: case 0x47: bcdec_bc2(b, o, op); break; - case 0x44: case 0x48: bcdec_bc3(b, o, op); break; - case 0x49: bcdec_bc4(b, o, op, 0); break; - case 0x4a: bcdec_bc4(b, o, op, 1); break; - case 0x4b: bcdec_bc5(b, o, op, 0); break; - case 0x4c: bcdec_bc5(b, o, op, 1); break; - case 0x4d: case 0x4e: bcdec_bc7(b, o, op); break; - case 0x4f: case 0x50: { // bcdec outputs RGB half-float; add alpha = 1.0 - uint16_t rgb[16 * 3]; bcdec_bc6h_half(b, rgb, 4 * 3, f == 0x4f); - for (int y = 0; y < 4; y++) for (int x = 0; x < 4; x++) { - uint16_t *d = (uint16_t *)(o + y * op) + x * 4; const uint16_t *s = rgb + (y * 4 + x) * 3; - d[0] = s[0]; d[1] = s[1]; d[2] = s[2]; d[3] = 0x3c00; - } - break; } - } -} - -int nvn_bc_decode(uint32_t f, const uint8_t *src, size_t pitch, uint32_t w, uint32_t h, uint8_t *dst, size_t dpitch) { - uint32_t pf; int bpp; if (!nvn_bc_unpacked(f, &pf, &bpp)) return -1; - int bb = (f == 0x41 || f == 0x42 || f == 0x45 || f == 0x46 || f == 0x49 || f == 0x4a) ? 8 : 16; // bytes per block - uint8_t tmp[4 * 4 * 8]; - for (uint32_t by = 0; by * 4 < h; by++) for (uint32_t bx = 0; bx * 4 < w; bx++) { - const uint8_t *b = src + by * pitch + bx * bb; - uint32_t cw = w - bx * 4 < 4 ? w - bx * 4 : 4, ch = h - by * 4 < 4 ? h - by * 4 : 4; - uint8_t *o = dst + (size_t)by * 4 * dpitch + (size_t)bx * 4 * bpp; - if (cw == 4 && ch == 4) { block(f, b, o, (int)dpitch); continue; } - block(f, b, tmp, 4 * bpp); // blocks at the edge that are not a full 4×4: decode into a temporary block first, then copy the needed part - for (uint32_t y = 0; y < ch; y++) memcpy(o + y * dpitch, tmp + y * 4 * bpp, cw * bpp); - } - return 0; -} diff --git a/dm/gfx/tex/nvn_bcdec.h b/dm/gfx/tex/nvn_bcdec.h deleted file mode 100644 index 2b44888..0000000 --- a/dm/gfx/tex/nvn_bcdec.h +++ /dev/null @@ -1,10 +0,0 @@ -// GPUs without BC compressed format support (A15 / A16 etc., Apple GPU family Apple8 and below): BC textures are decompressed on the CPU into uncompressed formats. -// On a GPU without BC every BC texture is stored decompressed; BC1-5 are decoded on the GPU at upload (nvn_mtl.m, bcd), BC6H / BC7 / 3D textures go through here (d2r-mobile's offline ASTC pack is not part of this repository). -#pragma once -#include -#include - -// NVN format is BC → returns 1, *mtl_pf = Metal pixel format after decompression, *bpp = bytes per pixel after decompression; returns 0 if not BC -int nvn_bc_unpacked(uint32_t nvn_fmt, uint32_t *mtl_pf, int *bpp); -// Decode a w×h pixel region: src is laid out by block rows (pitch bytes per block row), dst has dpitch bytes per pixel row. Returns 0 on success -int nvn_bc_decode(uint32_t nvn_fmt, const uint8_t *src, size_t pitch, uint32_t w, uint32_t h, uint8_t *dst, size_t dpitch); diff --git a/dm/third_party/README.md b/dm/third_party/README.md index 811433b..1ba5cf6 100644 --- a/dm/third_party/README.md +++ b/dm/third_party/README.md @@ -4,7 +4,6 @@ This is code written by others, shipped with the repository so there is nothing | Directory | What it is | Version and changes | Where it is used | License file | |---|---|---|---|---| -| `bcdec/` | Decompression of BC-compressed textures, a single header `bcdec.h` | 0.985 (upstream commit `80859ed`), unmodified | Graphics layer: decompresses textures on the CPU when the GPU does not support BC formats (`dm/gfx/tex/nvn_bcdec.c`) | `bcdec/LICENSE` (MIT or Unlicense, your choice) | | `flip/` | NVIDIA FLIP, an image difference metric that models human viewing. `FLIP.h` plus the source of the command-line tool in `tool/` | Not recorded in the directory | Verification: the image quality measurements in `dm/verify/gfxq/`; build `flip` first with `make -C dm/third_party/flip` | `flip/LICENSE` (BSD 3-Clause) | | `libdeflate/` | Decompression library | Upstream commit `92e6a0d`, unmodified | Host: zlib decoding of the game's resource blocks (`dm/host/hle_tact.c`). Only the five decompression-related source files are built; they are listed in `dm/host/Makefile` | `libdeflate/COPYING` (MIT) | | `libvpx/` | The VP9 decoder from libvpx: the 71 source files a VP9-decoder-only build compiles, their headers, and the generated configuration in `config/` | v1.17.0 (upstream commit `6df3ec3`), unmodified | Host: decodes the games' VP9 movies (`dm/host/hle_movie.c`). `dm/host/Makefile` builds it into `gen/vpx/libvpx.a`; the iOS build compiles the same files | `libvpx/LICENSE` (BSD 3-Clause), with the patent grant in `libvpx/PATENTS` | @@ -15,5 +14,5 @@ This is code written by others, shipped with the repository so there is nothing A few more notes: - `flip/tool/` also carries three single-header libraries used by the FLIP command-line tool: `stb_image.h` (2.23), `stb_image_write.h` (1.13) and `tinyexr.h`. Their license notices are in the files themselves. -- `bcdec/README.md`, `libdeflate/VERSION.md`, `libvpx/README.md`, `stb_vorbis/README.md`, `ryujinx/README.md` and `vulkan-headers/README.md` record where each came from and what it is used for. +- `libdeflate/VERSION.md`, `libvpx/README.md`, `stb_vorbis/README.md`, `ryujinx/README.md` and `vulkan-headers/README.md` record where each came from and what it is used for. - Two pieces of third-party code live outside this directory: `dm/gfx/nvn/include/nvn.h` (MIT, author named in the file header) and the layout size algorithm in `dm/gfx/nvn/nvn_tex.c`, ported from deko3d (zlib license). diff --git a/dm/third_party/bcdec/LICENSE b/dm/third_party/bcdec/LICENSE deleted file mode 100644 index 78e53f1..0000000 --- a/dm/third_party/bcdec/LICENSE +++ /dev/null @@ -1,52 +0,0 @@ -This software is available under 2 licenses -- choose whichever you prefer. - ------------------------------------------------------------------------------- -ALTERNATIVE A - MIT License - -Copyright (c) 2022 Sergii Kudlai - -Permission is hereby granted, free of charge, to any person obtaining a copy of -this software and associated documentation files (the "Software"), to deal in -the Software without restriction, including without limitation the rights to -use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies -of the Software, and to permit persons to whom the Software is furnished to do -so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in all -copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -SOFTWARE. - ------------------------------------------------------------------------------- -ALTERNATIVE B - The Unlicense - -This is free and unencumbered software released into the public domain. - -Anyone is free to copy, modify, publish, use, compile, sell, or -distribute this software, either in source code form or as a compiled -binary, for any purpose, commercial or non-commercial, and by any -means. - -In jurisdictions that recognize copyright laws, the author or authors -of this software dedicate any and all copyright interest in the -software to the public domain. We make this dedication for the benefit -of the public at large and to the detriment of our heirs and -successors. We intend this dedication to be an overt act of -relinquishment in perpetuity of all present and future rights to this -software under copyright law. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, -EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF -MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. -IN NO EVENT SHALL THE AUTHORS BE LIABLE FOR ANY CLAIM, DAMAGES OR -OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, -ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR -OTHER DEALINGS IN THE SOFTWARE. - -For more information, please refer to diff --git a/dm/third_party/bcdec/README.md b/dm/third_party/bcdec/README.md deleted file mode 100644 index 85fb9c7..0000000 --- a/dm/third_party/bcdec/README.md +++ /dev/null @@ -1,2 +0,0 @@ -Source: https://github.com/iOrange/bcdec commit 80859ed (version 0.985), unmodified. License in LICENSE (MIT or public domain, at your choice). -Use: on GPUs without BC compressed formats, the translation layer decompresses BC textures on the CPU (dm/gfx/tex/nvn_bcdec.c). diff --git a/dm/third_party/bcdec/bcdec.h b/dm/third_party/bcdec/bcdec.h deleted file mode 100644 index 553c3ec..0000000 --- a/dm/third_party/bcdec/bcdec.h +++ /dev/null @@ -1,1512 +0,0 @@ -/* bcdec.h - v0.985 - provides functions to decompress blocks of BC compressed images - written by Sergii "iOrange" Kudlai in 2022 - - This library does not allocate memory and is trying to use as less stack as possible - - This library was never optimized specifically for speed but for the overall size - it has zero external dependencies and is not using any runtime functions - - This library was writtem by humans and will never contain any LLM-generated code - - Supported BC formats: - BC1 (also known as DXT1) + it's "binary alpha" variant BC1A (DXT1A) - BC2 (also known as DXT3) - BC3 (also known as DXT5) - BC4 (also known as ATI1N) - BC5 (also known as ATI2N) - BC6H (HDR format) - BC7 - - BC1/BC2/BC3/BC7 are expected to decompress into 4*4 RGBA blocks 8bit per component (32bit pixel) - BC4/BC5 are expected to decompress into 4*4 R/RG blocks 8bit per component (8bit and 16bit pixel) - BC6H is expected to decompress into 4*4 RGB blocks of either 32bit float or 16bit "half" per - component (96bit or 48bit pixel) - - For more info, issues and suggestions please visit https://github.com/iOrange/bcdec - - Configuration: - #define BCDEC_BC4BC5_PRECISE: - enables more precise but slower BC4/BC5 decoding + signed/unsigned mode - - CREDITS: - Aras Pranckevicius (@aras-p) - BC1/BC3 decoders optimizations (up to 3x the speed) - - BC6H/BC7 bits pulling routines optimizations - - optimized BC6H by moving unquantize out of the loop - - Split BC6H decompression function into 'half' and - 'float' variants - - Michael Schmidt (@RunDevelopment) - Found better "magic" coefficients for integer interpolation - of reference colors in BC1 color block, that match with - the floating point interpolation. This also made it faster - than integer division by 3! - - bugfixes: - @linkmauve - @wkjarosz - - LICENSE: See end of file for license information. -*/ - -#ifndef BCDEC_HEADER_INCLUDED -#define BCDEC_HEADER_INCLUDED - -#define BCDEC_VERSION_MAJOR 0 -#define BCDEC_VERSION_MINOR 985 - -/* if BCDEC_STATIC causes problems, try defining BCDECDEF to 'inline' or 'static inline' */ -#ifndef BCDECDEF -#ifdef BCDEC_STATIC -#define BCDECDEF static -#else -#ifdef __cplusplus -#define BCDECDEF extern "C" -#else -#define BCDECDEF extern -#endif -#endif -#endif - -/* Used information sources: - https://docs.microsoft.com/en-us/windows/win32/direct3d10/d3d10-graphics-programming-guide-resources-block-compression - https://docs.microsoft.com/en-us/windows/win32/direct3d11/bc6h-format - https://docs.microsoft.com/en-us/windows/win32/direct3d11/bc7-format - https://docs.microsoft.com/en-us/windows/win32/direct3d11/bc7-format-mode-reference - - ! WARNING ! Khronos's BPTC partitions tables contain mistakes, do not use them! - https://www.khronos.org/registry/DataFormat/specs/1.1/dataformat.1.1.html#BPTC - - ! Use tables from here instead ! - https://www.khronos.org/registry/OpenGL/extensions/ARB/ARB_texture_compression_bptc.txt - - Leaving it here as it's a nice read - https://fgiesen.wordpress.com/2021/10/04/gpu-bcn-decoding/ - - Fast half to float function from here - https://gist.github.com/rygorous/2144712 -*/ - -#define BCDEC_BC1_BLOCK_SIZE 8 -#define BCDEC_BC2_BLOCK_SIZE 16 -#define BCDEC_BC3_BLOCK_SIZE 16 -#define BCDEC_BC4_BLOCK_SIZE 8 -#define BCDEC_BC5_BLOCK_SIZE 16 -#define BCDEC_BC6H_BLOCK_SIZE 16 -#define BCDEC_BC7_BLOCK_SIZE 16 - -#define BCDEC_BC1_COMPRESSED_SIZE(w, h) ((((w)>>2)*((h)>>2))*BCDEC_BC1_BLOCK_SIZE) -#define BCDEC_BC2_COMPRESSED_SIZE(w, h) ((((w)>>2)*((h)>>2))*BCDEC_BC2_BLOCK_SIZE) -#define BCDEC_BC3_COMPRESSED_SIZE(w, h) ((((w)>>2)*((h)>>2))*BCDEC_BC3_BLOCK_SIZE) -#define BCDEC_BC4_COMPRESSED_SIZE(w, h) ((((w)>>2)*((h)>>2))*BCDEC_BC4_BLOCK_SIZE) -#define BCDEC_BC5_COMPRESSED_SIZE(w, h) ((((w)>>2)*((h)>>2))*BCDEC_BC5_BLOCK_SIZE) -#define BCDEC_BC6H_COMPRESSED_SIZE(w, h) ((((w)>>2)*((h)>>2))*BCDEC_BC6H_BLOCK_SIZE) -#define BCDEC_BC7_COMPRESSED_SIZE(w, h) ((((w)>>2)*((h)>>2))*BCDEC_BC7_BLOCK_SIZE) - -BCDECDEF void bcdec_bc1(const void* compressedBlock, void* decompressedBlock, int destinationPitch); -BCDECDEF void bcdec_bc2(const void* compressedBlock, void* decompressedBlock, int destinationPitch); -BCDECDEF void bcdec_bc3(const void* compressedBlock, void* decompressedBlock, int destinationPitch); -#ifndef BCDEC_BC4BC5_PRECISE -BCDECDEF void bcdec_bc4(const void* compressedBlock, void* decompressedBlock, int destinationPitch); -BCDECDEF void bcdec_bc5(const void* compressedBlock, void* decompressedBlock, int destinationPitch); -#else -BCDECDEF void bcdec_bc4(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned); -BCDECDEF void bcdec_bc5(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned); -BCDECDEF void bcdec_bc4_float(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned); -BCDECDEF void bcdec_bc5_float(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned); -#endif -BCDECDEF void bcdec_bc6h_float(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned); -BCDECDEF void bcdec_bc6h_half(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned); -BCDECDEF void bcdec_bc7(const void* compressedBlock, void* decompressedBlock, int destinationPitch); - -#endif /* BCDEC_HEADER_INCLUDED */ - -#ifdef BCDEC_IMPLEMENTATION - -static unsigned short bcdec__getu16le(const void* s) { - unsigned short l = (unsigned short)((const unsigned char*)s)[0]; - unsigned short h = (unsigned short)((const unsigned char*)s)[1]; - return l | (h << 8); -} -static unsigned int bcdec__getu32le(const void* s) { - unsigned int l = bcdec__getu16le(s); - unsigned int h = bcdec__getu16le((const unsigned char*)s + 2); - return l | (h << 16); -} -static unsigned long long bcdec__getu64le(const void* s) { - unsigned long long l = bcdec__getu32le(s); - unsigned long long h = bcdec__getu32le((const unsigned char*)s + 4); - return l | (h << 32); -} - -static void bcdec__color_block(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int onlyOpaqueMode) { - unsigned short c0, c1; - unsigned int refColors[4]; /* 0xAABBGGRR */ - unsigned char* dstColors; - unsigned int colorIndices; - int i, j, idx; - unsigned int r0, g0, b0, r1, g1, b1, r, g, b; - - c0 = bcdec__getu16le(compressedBlock); - c1 = bcdec__getu16le((const unsigned char*)compressedBlock + 2); - - /* Unpack 565 ref colors */ - r0 = (c0 >> 11) & 0x1F; - g0 = (c0 >> 5) & 0x3F; - b0 = c0 & 0x1F; - - r1 = (c1 >> 11) & 0x1F; - g1 = (c1 >> 5) & 0x3F; - b1 = c1 & 0x1F; - - /* Expand 565 ref colors to 888 */ - r = (r0 * 527 + 23) >> 6; - g = (g0 * 259 + 33) >> 6; - b = (b0 * 527 + 23) >> 6; - refColors[0] = 0xFF000000 | (b << 16) | (g << 8) | r; - - r = (r1 * 527 + 23) >> 6; - g = (g1 * 259 + 33) >> 6; - b = (b1 * 527 + 23) >> 6; - refColors[1] = 0xFF000000 | (b << 16) | (g << 8) | r; - - if (c0 > c1 || onlyOpaqueMode) { /* Standard BC1 mode (also BC3 color block uses ONLY this mode) */ - /* color_2 = 2/3*color_0 + 1/3*color_1 - color_3 = 1/3*color_0 + 2/3*color_1 */ - r = ((2 * r0 + r1) * 351 + 61) >> 7; - g = ((2 * g0 + g1) * 2763 + 1039) >> 11; - b = ((2 * b0 + b1) * 351 + 61) >> 7; - refColors[2] = 0xFF000000 | (b << 16) | (g << 8) | r; - - r = ((r0 + r1 * 2) * 351 + 61) >> 7; - g = ((g0 + g1 * 2) * 2763 + 1039) >> 11; - b = ((b0 + b1 * 2) * 351 + 61) >> 7; - refColors[3] = 0xFF000000 | (b << 16) | (g << 8) | r; - } else { /* Quite rare BC1A mode */ - /* color_2 = 1/2*color_0 + 1/2*color_1; - color_3 = 0; */ - r = ((r0 + r1) * 1053 + 125) >> 8; - g = ((g0 + g1) * 4145 + 1019) >> 11; - b = ((b0 + b1) * 1053 + 125) >> 8; - refColors[2] = 0xFF000000 | (b << 16) | (g << 8) | r; - - refColors[3] = 0x00000000; - } - - colorIndices = bcdec__getu32le((const unsigned char*)compressedBlock + 4); - - /* Fill out the decompressed color block */ - dstColors = (unsigned char*)decompressedBlock; - for (i = 0; i < 4; ++i) { - for (j = 0; j < 4; ++j) { - idx = colorIndices & 0x03; - ((unsigned int*)dstColors)[j] = refColors[idx]; - colorIndices >>= 2; - } - - dstColors += destinationPitch; - } -} - -static void bcdec__sharp_alpha_block(const void* compressedBlock, void* decompressedBlock, int destinationPitch) { - unsigned short alpha; - const unsigned char* alphaPtr; - unsigned char* decompressed; - int i, j; - - alphaPtr = (const unsigned char*)compressedBlock; - decompressed = (unsigned char*)decompressedBlock; - - for (i = 0; i < 4; ++i) { - alpha = bcdec__getu16le(alphaPtr); - - for (j = 0; j < 4; ++j) { - decompressed[j * 4] = ((alpha >> (4 * j)) & 0x0F) * 17; - } - - alphaPtr += 2; - decompressed += destinationPitch; - } -} - -static void bcdec__smooth_alpha_block(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int pixelSize) { - unsigned char* decompressed; - unsigned char alpha[8]; - int i, j; - unsigned long long block, indices; - - block = bcdec__getu64le(compressedBlock); - decompressed = (unsigned char*)decompressedBlock; - - alpha[0] = block & 0xFF; - alpha[1] = (block >> 8) & 0xFF; - - if (alpha[0] > alpha[1]) { - /* 6 interpolated alpha values. */ - alpha[2] = (6 * alpha[0] + alpha[1]) / 7; /* 6/7*alpha_0 + 1/7*alpha_1 */ - alpha[3] = (5 * alpha[0] + 2 * alpha[1]) / 7; /* 5/7*alpha_0 + 2/7*alpha_1 */ - alpha[4] = (4 * alpha[0] + 3 * alpha[1]) / 7; /* 4/7*alpha_0 + 3/7*alpha_1 */ - alpha[5] = (3 * alpha[0] + 4 * alpha[1]) / 7; /* 3/7*alpha_0 + 4/7*alpha_1 */ - alpha[6] = (2 * alpha[0] + 5 * alpha[1]) / 7; /* 2/7*alpha_0 + 5/7*alpha_1 */ - alpha[7] = ( alpha[0] + 6 * alpha[1]) / 7; /* 1/7*alpha_0 + 6/7*alpha_1 */ - } - else { - /* 4 interpolated alpha values. */ - alpha[2] = (4 * alpha[0] + alpha[1]) / 5; /* 4/5*alpha_0 + 1/5*alpha_1 */ - alpha[3] = (3 * alpha[0] + 2 * alpha[1]) / 5; /* 3/5*alpha_0 + 2/5*alpha_1 */ - alpha[4] = (2 * alpha[0] + 3 * alpha[1]) / 5; /* 2/5*alpha_0 + 3/5*alpha_1 */ - alpha[5] = ( alpha[0] + 4 * alpha[1]) / 5; /* 1/5*alpha_0 + 4/5*alpha_1 */ - alpha[6] = 0x00; - alpha[7] = 0xFF; - } - - indices = block >> 16; - for (i = 0; i < 4; ++i) { - for (j = 0; j < 4; ++j) { - decompressed[j * pixelSize] = alpha[indices & 0x07]; - indices >>= 3; - } - - decompressed += destinationPitch; - } -} - -#ifdef BCDEC_BC4BC5_PRECISE -static void bcdec__bc4_block(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int pixelSize, int isSigned) { - signed char* sblock; - unsigned char* ublock; - int alpha[8]; - int i, j; - unsigned long long block, indices; - - static int aWeights4[4] = { 13107, 26215, 39321, 52429 }; - static int aWeights6[6] = { 9363, 18724, 28086, 37450, 46812, 56173 }; - - block = bcdec__getu64le(compressedBlock); - - if (isSigned) { - alpha[0] = (char)(block & 0xFF); - alpha[1] = (char)((block >> 8) & 0xFF); - if (alpha[0] < -127) alpha[0] = -127; /* -128 clamps to -127 */ - if (alpha[1] < -127) alpha[1] = -127; /* -128 clamps to -127 */ - } else { - alpha[0] = block & 0xFF; - alpha[1] = (block >> 8) & 0xFF; - } - - if (alpha[0] > alpha[1]) { - /* 6 interpolated alpha values. */ - alpha[2] = (aWeights6[5] * alpha[0] + aWeights6[0] * alpha[1] + 32768) >> 16; /* 6/7*alpha_0 + 1/7*alpha_1 */ - alpha[3] = (aWeights6[4] * alpha[0] + aWeights6[1] * alpha[1] + 32768) >> 16; /* 5/7*alpha_0 + 2/7*alpha_1 */ - alpha[4] = (aWeights6[3] * alpha[0] + aWeights6[2] * alpha[1] + 32768) >> 16; /* 4/7*alpha_0 + 3/7*alpha_1 */ - alpha[5] = (aWeights6[2] * alpha[0] + aWeights6[3] * alpha[1] + 32768) >> 16; /* 3/7*alpha_0 + 4/7*alpha_1 */ - alpha[6] = (aWeights6[1] * alpha[0] + aWeights6[4] * alpha[1] + 32768) >> 16; /* 2/7*alpha_0 + 5/7*alpha_1 */ - alpha[7] = (aWeights6[0] * alpha[0] + aWeights6[5] * alpha[1] + 32768) >> 16; /* 1/7*alpha_0 + 6/7*alpha_1 */ - } else { - /* 4 interpolated alpha values. */ - alpha[2] = (aWeights4[3] * alpha[0] + aWeights4[0] * alpha[1] + 32768) >> 16; /* 4/5*alpha_0 + 1/5*alpha_1 */ - alpha[3] = (aWeights4[2] * alpha[0] + aWeights4[1] * alpha[1] + 32768) >> 16; /* 3/5*alpha_0 + 2/5*alpha_1 */ - alpha[4] = (aWeights4[1] * alpha[0] + aWeights4[2] * alpha[1] + 32768) >> 16; /* 2/5*alpha_0 + 3/5*alpha_1 */ - alpha[5] = (aWeights4[0] * alpha[0] + aWeights4[3] * alpha[1] + 32768) >> 16; /* 1/5*alpha_0 + 4/5*alpha_1 */ - alpha[6] = isSigned ? -127 : 0; - alpha[7] = isSigned ? 127 : 255; - } - - indices = block >> 16; - if (isSigned) { - sblock = (signed char*)decompressedBlock; - for (i = 0; i < 4; ++i) { - for (j = 0; j < 4; ++j) { - sblock[j * pixelSize] = (signed char)alpha[indices & 0x07]; - indices >>= 3; - } - sblock += destinationPitch; - } - } else { - ublock = (unsigned char*)decompressedBlock; - for (i = 0; i < 4; ++i) { - for (j = 0; j < 4; ++j) { - ublock[j * pixelSize] = (unsigned char)alpha[indices & 0x07]; - indices >>= 3; - } - ublock += destinationPitch; - } - } -} - -static void bcdec__bc4_block_float(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int pixelSize, int isSigned) { - float* decompressed; - float alpha[8]; - int i, j; - unsigned long long block, indices; - - block = bcdec__getu64le(compressedBlock); - decompressed = (float*)decompressedBlock; - - if (isSigned) { - alpha[0] = (float)((char)(block & 0xFF)) / 127.0f; - alpha[1] = (float)((char)((block >> 8) & 0xFF)) / 127.0f; - if (alpha[0] < -1.0f) alpha[0] = -1.0f; /* -128 clamps to -127 */ - if (alpha[1] < -1.0f) alpha[1] = -1.0f; /* -128 clamps to -127 */ - } else { - alpha[0] = (float)(block & 0xFF) / 255.0f; - alpha[1] = (float)((block >> 8) & 0xFF) / 255.0f; - } - - if (alpha[0] > alpha[1]) { - /* 6 interpolated alpha values. */ - alpha[2] = (6.0f * alpha[0] + alpha[1]) / 7.0f; /* 6/7*alpha_0 + 1/7*alpha_1 */ - alpha[3] = (5.0f * alpha[0] + 2.0f * alpha[1]) / 7.0f; /* 5/7*alpha_0 + 2/7*alpha_1 */ - alpha[4] = (4.0f * alpha[0] + 3.0f * alpha[1]) / 7.0f; /* 4/7*alpha_0 + 3/7*alpha_1 */ - alpha[5] = (3.0f * alpha[0] + 4.0f * alpha[1]) / 7.0f; /* 3/7*alpha_0 + 4/7*alpha_1 */ - alpha[6] = (2.0f * alpha[0] + 5.0f * alpha[1]) / 7.0f; /* 2/7*alpha_0 + 5/7*alpha_1 */ - alpha[7] = ( alpha[0] + 6.0f * alpha[1]) / 7.0f; /* 1/7*alpha_0 + 6/7*alpha_1 */ - } else { - /* 4 interpolated alpha values. */ - alpha[2] = (4.0f * alpha[0] + alpha[1]) / 5.0f; /* 4/5*alpha_0 + 1/5*alpha_1 */ - alpha[3] = (3.0f * alpha[0] + 2.0f * alpha[1]) / 5.0f; /* 3/5*alpha_0 + 2/5*alpha_1 */ - alpha[4] = (2.0f * alpha[0] + 3.0f * alpha[1]) / 5.0f; /* 2/5*alpha_0 + 3/5*alpha_1 */ - alpha[5] = ( alpha[0] + 4.0f * alpha[1]) / 5.0f; /* 1/5*alpha_0 + 4/5*alpha_1 */ - alpha[6] = isSigned ? -1.0f : 0.0f; - alpha[7] = 1.0f; - } - - indices = block >> 16; - for (i = 0; i < 4; ++i) { - for (j = 0; j < 4; ++j) { - decompressed[j * pixelSize] = alpha[indices & 0x07]; - indices >>= 3; - } - decompressed += destinationPitch; - } -} -#endif /* BCDEC_BC4BC5_PRECISE */ - -typedef struct bcdec__bitstream { - unsigned long long low; - unsigned long long high; -} bcdec__bitstream_t; - -static int bcdec__bitstream_read_bits(bcdec__bitstream_t* bstream, int numBits) { - unsigned int mask = (1 << numBits) - 1; - /* Read the low N bits */ - unsigned int bits = (bstream->low & mask); - - bstream->low >>= numBits; - /* Put the low N bits of "high" into the high 64-N bits of "low". */ - bstream->low |= (bstream->high & mask) << (sizeof(bstream->high) * 8 - numBits); - bstream->high >>= numBits; - - return bits; -} - -static int bcdec__bitstream_read_bit(bcdec__bitstream_t* bstream) { - return bcdec__bitstream_read_bits(bstream, 1); -} - -/* reversed bits pulling, used in BC6H decoding - why ?? just why ??? */ -static int bcdec__bitstream_read_bits_r(bcdec__bitstream_t* bstream, int numBits) { - int bits = bcdec__bitstream_read_bits(bstream, numBits); - /* Reverse the bits. */ - int result = 0; - while (numBits--) { - result <<= 1; - result |= (bits & 1); - bits >>= 1; - } - return result; -} - - - -BCDECDEF void bcdec_bc1(const void* compressedBlock, void* decompressedBlock, int destinationPitch) { - bcdec__color_block(compressedBlock, decompressedBlock, destinationPitch, 0); -} - -BCDECDEF void bcdec_bc2(const void* compressedBlock, void* decompressedBlock, int destinationPitch) { - bcdec__color_block(((char*)compressedBlock) + 8, decompressedBlock, destinationPitch, 1); - bcdec__sharp_alpha_block(compressedBlock, ((char*)decompressedBlock) + 3, destinationPitch); -} - -BCDECDEF void bcdec_bc3(const void* compressedBlock, void* decompressedBlock, int destinationPitch) { - bcdec__color_block(((char*)compressedBlock) + 8, decompressedBlock, destinationPitch, 1); - bcdec__smooth_alpha_block(compressedBlock, ((char*)decompressedBlock) + 3, destinationPitch, 4); -} - -#ifndef BCDEC_BC4BC5_PRECISE -BCDECDEF void bcdec_bc4(const void* compressedBlock, void* decompressedBlock, int destinationPitch) { - bcdec__smooth_alpha_block(compressedBlock, decompressedBlock, destinationPitch, 1); -#else -BCDECDEF void bcdec_bc4(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned) { - bcdec__bc4_block(compressedBlock, decompressedBlock, destinationPitch, 1, isSigned); -#endif -} - -#ifndef BCDEC_BC4BC5_PRECISE -BCDECDEF void bcdec_bc5(const void* compressedBlock, void* decompressedBlock, int destinationPitch) { - bcdec__smooth_alpha_block(compressedBlock, decompressedBlock, destinationPitch, 2); - bcdec__smooth_alpha_block(((char*)compressedBlock) + 8, ((char*)decompressedBlock) + 1, destinationPitch, 2); -#else -BCDECDEF void bcdec_bc5(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned) { - bcdec__bc4_block(compressedBlock, decompressedBlock, destinationPitch, 2, isSigned); - bcdec__bc4_block(((char*)compressedBlock) + 8, ((char*)decompressedBlock) + 1, destinationPitch, 2, isSigned); -#endif -} - -#ifdef BCDEC_BC4BC5_PRECISE -BCDECDEF void bcdec_bc4_float(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned) { - bcdec__bc4_block_float(compressedBlock, decompressedBlock, destinationPitch, 1, isSigned); -} - -BCDECDEF void bcdec_bc5_float(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned) { - bcdec__bc4_block_float(compressedBlock, decompressedBlock, destinationPitch, 2, isSigned); - bcdec__bc4_block_float(((char*)compressedBlock) + 8, ((float*)decompressedBlock) + 1, destinationPitch, 2, isSigned); -} -#endif /* BCDEC_BC4BC5_PRECISE */ - -/* http://graphics.stanford.edu/~seander/bithacks.html#VariableSignExtend */ -static int bcdec__extend_sign(int val, int bits) { - return (val << (32 - bits)) >> (32 - bits); -} - -static int bcdec__transform_inverse(int val, int a0, int bits, int isSigned) { - /* If the precision of A0 is "p" bits, then the transform algorithm is: - B0 = (B0 + A0) & ((1 << p) - 1) */ - val = (val + a0) & ((1 << bits) - 1); - if (isSigned) { - val = bcdec__extend_sign(val, bits); - } - return val; -} - -/* pretty much copy-paste from documentation */ -static int bcdec__unquantize(int val, int bits, int isSigned) { - int unq, s = 0; - - if (!isSigned) { - if (bits >= 15) { - unq = val; - } else if (!val) { - unq = 0; - } else if (val == ((1 << bits) - 1)) { - unq = 0xFFFF; - } else { - unq = ((val << 16) + 0x8000) >> bits; - } - } else { - if (bits >= 16) { - unq = val; - } else { - if (val < 0) { - s = 1; - val = -val; - } - - if (val == 0) { - unq = 0; - } else if (val >= ((1 << (bits - 1)) - 1)) { - unq = 0x7FFF; - } else { - unq = ((val << 15) + 0x4000) >> (bits - 1); - } - - if (s) { - unq = -unq; - } - } - } - return unq; -} - -static int bcdec__interpolate(int a, int b, int* weights, int index) { - return (a * (64 - weights[index]) + b * weights[index] + 32) >> 6; -} - -static unsigned short bcdec__finish_unquantize(int val, int isSigned) { - int s; - - if (!isSigned) { - return (unsigned short)((val * 31) >> 6); /* scale the magnitude by 31 / 64 */ - } else { - val = (val < 0) ? -(((-val) * 31) >> 5) : (val * 31) >> 5; /* scale the magnitude by 31 / 32 */ - s = 0; - if (val < 0) { - s = 0x8000; - val = -val; - } - return (unsigned short)(s | val); - } -} - -/* modified half_to_float_fast4 from https://gist.github.com/rygorous/2144712 */ -static float bcdec__half_to_float_quick(unsigned short half) { - typedef union { - unsigned int u; - float f; - } FP32; - - static const FP32 magic = { 113 << 23 }; - static const unsigned int shifted_exp = 0x7c00 << 13; /* exponent mask after shift */ - FP32 o; - unsigned int exp; - - o.u = (half & 0x7fff) << 13; /* exponent/mantissa bits */ - exp = shifted_exp & o.u; /* just the exponent */ - o.u += (127 - 15) << 23; /* exponent adjust */ - - /* handle exponent special cases */ - if (exp == shifted_exp) { /* Inf/NaN? */ - o.u += (128 - 16) << 23; /* extra exp adjust */ - } else if (exp == 0) { /* Zero/Denormal? */ - o.u += 1 << 23; /* extra exp adjust */ - o.f -= magic.f; /* renormalize */ - } - - o.u |= (half & 0x8000) << 16; /* sign bit */ - return o.f; -} - -BCDECDEF void bcdec_bc6h_half(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned) { - static char actual_bits_count[4][14] = { - { 10, 7, 11, 11, 11, 9, 8, 8, 8, 6, 10, 11, 12, 16 }, /* W */ - { 5, 6, 5, 4, 4, 5, 6, 5, 5, 6, 10, 9, 8, 4 }, /* dR */ - { 5, 6, 4, 5, 4, 5, 5, 6, 5, 6, 10, 9, 8, 4 }, /* dG */ - { 5, 6, 4, 4, 5, 5, 5, 5, 6, 6, 10, 9, 8, 4 } /* dB */ - }; - - /* There are 32 possible partition sets for a two-region tile. - Each 4x4 block represents a single shape. - Here also every fix-up index has MSB bit set. */ - static unsigned char partition_sets[32][4][4] = { - { {128, 0, 1, 1}, {0, 0, 1, 1}, { 0, 0, 1, 1}, {0, 0, 1, 129} }, /* 0 */ - { {128, 0, 0, 1}, {0, 0, 0, 1}, { 0, 0, 0, 1}, {0, 0, 0, 129} }, /* 1 */ - { {128, 1, 1, 1}, {0, 1, 1, 1}, { 0, 1, 1, 1}, {0, 1, 1, 129} }, /* 2 */ - { {128, 0, 0, 1}, {0, 0, 1, 1}, { 0, 0, 1, 1}, {0, 1, 1, 129} }, /* 3 */ - { {128, 0, 0, 0}, {0, 0, 0, 1}, { 0, 0, 0, 1}, {0, 0, 1, 129} }, /* 4 */ - { {128, 0, 1, 1}, {0, 1, 1, 1}, { 0, 1, 1, 1}, {1, 1, 1, 129} }, /* 5 */ - { {128, 0, 0, 1}, {0, 0, 1, 1}, { 0, 1, 1, 1}, {1, 1, 1, 129} }, /* 6 */ - { {128, 0, 0, 0}, {0, 0, 0, 1}, { 0, 0, 1, 1}, {0, 1, 1, 129} }, /* 7 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, { 0, 0, 0, 1}, {0, 0, 1, 129} }, /* 8 */ - { {128, 0, 1, 1}, {0, 1, 1, 1}, { 1, 1, 1, 1}, {1, 1, 1, 129} }, /* 9 */ - { {128, 0, 0, 0}, {0, 0, 0, 1}, { 0, 1, 1, 1}, {1, 1, 1, 129} }, /* 10 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, { 0, 0, 0, 1}, {0, 1, 1, 129} }, /* 11 */ - { {128, 0, 0, 1}, {0, 1, 1, 1}, { 1, 1, 1, 1}, {1, 1, 1, 129} }, /* 12 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, { 1, 1, 1, 1}, {1, 1, 1, 129} }, /* 13 */ - { {128, 0, 0, 0}, {1, 1, 1, 1}, { 1, 1, 1, 1}, {1, 1, 1, 129} }, /* 14 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, { 0, 0, 0, 0}, {1, 1, 1, 129} }, /* 15 */ - { {128, 0, 0, 0}, {1, 0, 0, 0}, { 1, 1, 1, 0}, {1, 1, 1, 129} }, /* 16 */ - { {128, 1, 129, 1}, {0, 0, 0, 1}, { 0, 0, 0, 0}, {0, 0, 0, 0} }, /* 17 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, {129, 0, 0, 0}, {1, 1, 1, 0} }, /* 18 */ - { {128, 1, 129, 1}, {0, 0, 1, 1}, { 0, 0, 0, 1}, {0, 0, 0, 0} }, /* 19 */ - { {128, 0, 129, 1}, {0, 0, 0, 1}, { 0, 0, 0, 0}, {0, 0, 0, 0} }, /* 20 */ - { {128, 0, 0, 0}, {1, 0, 0, 0}, {129, 1, 0, 0}, {1, 1, 1, 0} }, /* 21 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, {129, 0, 0, 0}, {1, 1, 0, 0} }, /* 22 */ - { {128, 1, 1, 1}, {0, 0, 1, 1}, { 0, 0, 1, 1}, {0, 0, 0, 129} }, /* 23 */ - { {128, 0, 129, 1}, {0, 0, 0, 1}, { 0, 0, 0, 1}, {0, 0, 0, 0} }, /* 24 */ - { {128, 0, 0, 0}, {1, 0, 0, 0}, {129, 0, 0, 0}, {1, 1, 0, 0} }, /* 25 */ - { {128, 1, 129, 0}, {0, 1, 1, 0}, { 0, 1, 1, 0}, {0, 1, 1, 0} }, /* 26 */ - { {128, 0, 129, 1}, {0, 1, 1, 0}, { 0, 1, 1, 0}, {1, 1, 0, 0} }, /* 27 */ - { {128, 0, 0, 1}, {0, 1, 1, 1}, {129, 1, 1, 0}, {1, 0, 0, 0} }, /* 28 */ - { {128, 0, 0, 0}, {1, 1, 1, 1}, {129, 1, 1, 1}, {0, 0, 0, 0} }, /* 29 */ - { {128, 1, 129, 1}, {0, 0, 0, 1}, { 1, 0, 0, 0}, {1, 1, 1, 0} }, /* 30 */ - { {128, 0, 129, 1}, {1, 0, 0, 1}, { 1, 0, 0, 1}, {1, 1, 0, 0} } /* 31 */ - }; - - static int aWeight3[8] = { 0, 9, 18, 27, 37, 46, 55, 64 }; - static int aWeight4[16] = { 0, 4, 9, 13, 17, 21, 26, 30, 34, 38, 43, 47, 51, 55, 60, 64 }; - - bcdec__bitstream_t bstream; - int mode, partition, numPartitions, i, j, partitionSet, indexBits, index, ep_i, actualBits0Mode; - int r[4], g[4], b[4]; /* wxyz */ - unsigned short* decompressed; - int* weights; - - decompressed = (unsigned short*)decompressedBlock; - - bstream.low = bcdec__getu64le(compressedBlock); - bstream.high = bcdec__getu64le((const unsigned char*)compressedBlock + 8); - - r[0] = r[1] = r[2] = r[3] = 0; - g[0] = g[1] = g[2] = g[3] = 0; - b[0] = b[1] = b[2] = b[3] = 0; - - mode = bcdec__bitstream_read_bits(&bstream, 2); - if (mode > 1) { - mode |= (bcdec__bitstream_read_bits(&bstream, 3) << 2); - } - - /* modes >= 11 (10 in my code) are using 0 one, others will read it from the bitstream */ - partition = 0; - - switch (mode) { - /* mode 1 */ - case 0b00: { - /* Partitition indices: 46 bits - Partition: 5 bits - Color Endpoints: 75 bits (10.555, 10.555, 10.555) */ - g[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gy[4] */ - b[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* by[4] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* bz[4] */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* rw[9:0] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* gw[9:0] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* bw[9:0] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* rx[4:0] */ - g[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gz[4] */ - g[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* gy[3:0] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* gx[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream); /* bz[0] */ - g[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* gz[3:0] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* bx[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 1; /* bz[1] */ - b[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* by[3:0] */ - r[2] |= bcdec__bitstream_read_bits(&bstream, 5); /* ry[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 2; /* bz[2] */ - r[3] |= bcdec__bitstream_read_bits(&bstream, 5); /* rz[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 3; /* bz[3] */ - partition = bcdec__bitstream_read_bits(&bstream, 5); /* d[4:0] */ - mode = 0; - } break; - - /* mode 2 */ - case 0b01: { - /* Partitition indices: 46 bits - Partition: 5 bits - Color Endpoints: 75 bits (7666, 7666, 7666) */ - g[2] |= bcdec__bitstream_read_bit(&bstream) << 5; /* gy[5] */ - g[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gz[4] */ - g[3] |= bcdec__bitstream_read_bit(&bstream) << 5; /* gz[5] */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 7); /* rw[6:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream); /* bz[0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 1; /* bz[1] */ - b[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* by[4] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 7); /* gw[6:0] */ - b[2] |= bcdec__bitstream_read_bit(&bstream) << 5; /* by[5] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 2; /* bz[2] */ - g[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gy[4] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 7); /* bw[6:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 3; /* bz[3] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 5; /* bz[5] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* bz[4] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 6); /* rx[5:0] */ - g[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* gy[3:0] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 6); /* gx[5:0] */ - g[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* gz[3:0] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 6); /* bx[5:0] */ - b[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* by[3:0] */ - r[2] |= bcdec__bitstream_read_bits(&bstream, 6); /* ry[5:0] */ - r[3] |= bcdec__bitstream_read_bits(&bstream, 6); /* rz[5:0] */ - partition = bcdec__bitstream_read_bits(&bstream, 5); /* d[4:0] */ - mode = 1; - } break; - - /* mode 3 */ - case 0b00010: { - /* Partitition indices: 46 bits - Partition: 5 bits - Color Endpoints: 72 bits (11.555, 11.444, 11.444) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* rw[9:0] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* gw[9:0] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* bw[9:0] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* rx[4:0] */ - r[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* rw[10] */ - g[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* gy[3:0] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 4); /* gx[3:0] */ - g[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* gw[10] */ - b[3] |= bcdec__bitstream_read_bit(&bstream); /* bz[0] */ - g[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* gz[3:0] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 4); /* bx[3:0] */ - b[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* bw[10] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 1; /* bz[1] */ - b[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* by[3:0] */ - r[2] |= bcdec__bitstream_read_bits(&bstream, 5); /* ry[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 2; /* bz[2] */ - r[3] |= bcdec__bitstream_read_bits(&bstream, 5); /* rz[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 3; /* bz[3] */ - partition = bcdec__bitstream_read_bits(&bstream, 5); /* d[4:0] */ - mode = 2; - } break; - - /* mode 4 */ - case 0b00110: { - /* Partitition indices: 46 bits - Partition: 5 bits - Color Endpoints: 72 bits (11.444, 11.555, 11.444) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* rw[9:0] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* gw[9:0] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* bw[9:0] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 4); /* rx[3:0] */ - r[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* rw[10] */ - g[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gz[4] */ - g[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* gy[3:0] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* gx[4:0] */ - g[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* gw[10] */ - g[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* gz[3:0] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 4); /* bx[3:0] */ - b[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* bw[10] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 1; /* bz[1] */ - b[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* by[3:0] */ - r[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* ry[3:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream); /* bz[0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 2; /* bz[2] */ - r[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* rz[3:0] */ - g[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gy[4] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 3; /* bz[3] */ - partition = bcdec__bitstream_read_bits(&bstream, 5); /* d[4:0] */ - mode = 3; - } break; - - /* mode 5 */ - case 0b01010: { - /* Partitition indices: 46 bits - Partition: 5 bits - Color Endpoints: 72 bits (11.444, 11.444, 11.555) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* rw[9:0] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* gw[9:0] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* bw[9:0] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 4); /* rx[3:0] */ - r[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* rw[10] */ - b[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* by[4] */ - g[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* gy[3:0] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 4); /* gx[3:0] */ - g[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* gw[10] */ - b[3] |= bcdec__bitstream_read_bit(&bstream); /* bz[0] */ - g[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* gz[3:0] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* bx[4:0] */ - b[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* bw[10] */ - b[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* by[3:0] */ - r[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* ry[3:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 1; /* bz[1] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 2; /* bz[2] */ - r[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* rz[3:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* bz[4] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 3; /* bz[3] */ - partition = bcdec__bitstream_read_bits(&bstream, 5); /* d[4:0] */ - mode = 4; - } break; - - /* mode 6 */ - case 0b01110: { - /* Partitition indices: 46 bits - Partition: 5 bits - Color Endpoints: 72 bits (9555, 9555, 9555) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 9); /* rw[8:0] */ - b[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* by[4] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 9); /* gw[8:0] */ - g[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gy[4] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 9); /* bw[8:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* bz[4] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* rx[4:0] */ - g[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gz[4] */ - g[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* gy[3:0] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* gx[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream); /* bz[0] */ - g[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* gx[3:0] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* bx[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 1; /* bz[1] */ - b[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* by[3:0] */ - r[2] |= bcdec__bitstream_read_bits(&bstream, 5); /* ry[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 2; /* bz[2] */ - r[3] |= bcdec__bitstream_read_bits(&bstream, 5); /* rz[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 3; /* bz[3] */ - partition = bcdec__bitstream_read_bits(&bstream, 5); /* d[4:0] */ - mode = 5; - } break; - - /* mode 7 */ - case 0b10010: { - /* Partitition indices: 46 bits - Partition: 5 bits - Color Endpoints: 72 bits (8666, 8555, 8555) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 8); /* rw[7:0] */ - g[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gz[4] */ - b[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* by[4] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 8); /* gw[7:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 2; /* bz[2] */ - g[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gy[4] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 8); /* bw[7:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 3; /* bz[3] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* bz[4] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 6); /* rx[5:0] */ - g[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* gy[3:0] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* gx[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream); /* bz[0] */ - g[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* gz[3:0] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* bx[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 1; /* bz[1] */ - b[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* by[3:0] */ - r[2] |= bcdec__bitstream_read_bits(&bstream, 6); /* ry[5:0] */ - r[3] |= bcdec__bitstream_read_bits(&bstream, 6); /* rz[5:0] */ - partition = bcdec__bitstream_read_bits(&bstream, 5); /* d[4:0] */ - mode = 6; - } break; - - /* mode 8 */ - case 0b10110: { - /* Partitition indices: 46 bits - Partition: 5 bits - Color Endpoints: 72 bits (8555, 8666, 8555) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 8); /* rw[7:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream); /* bz[0] */ - b[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* by[4] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 8); /* gw[7:0] */ - g[2] |= bcdec__bitstream_read_bit(&bstream) << 5; /* gy[5] */ - g[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gy[4] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 8); /* bw[7:0] */ - g[3] |= bcdec__bitstream_read_bit(&bstream) << 5; /* gz[5] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* bz[4] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* rx[4:0] */ - g[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gz[4] */ - g[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* gy[3:0] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 6); /* gx[5:0] */ - g[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* zx[3:0] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* bx[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 1; /* bz[1] */ - b[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* by[3:0] */ - r[2] |= bcdec__bitstream_read_bits(&bstream, 5); /* ry[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 2; /* bz[2] */ - r[3] |= bcdec__bitstream_read_bits(&bstream, 5); /* rz[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 3; /* bz[3] */ - partition = bcdec__bitstream_read_bits(&bstream, 5); /* d[4:0] */ - mode = 7; - } break; - - /* mode 9 */ - case 0b11010: { - /* Partitition indices: 46 bits - Partition: 5 bits - Color Endpoints: 72 bits (8555, 8555, 8666) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 8); /* rw[7:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 1; /* bz[1] */ - b[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* by[4] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 8); /* gw[7:0] */ - b[2] |= bcdec__bitstream_read_bit(&bstream) << 5; /* by[5] */ - g[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gy[4] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 8); /* bw[7:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 5; /* bz[5] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* bz[4] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* bw[4:0] */ - g[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gz[4] */ - g[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* gy[3:0] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 5); /* gx[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream); /* bz[0] */ - g[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* gz[3:0] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 6); /* bx[5:0] */ - b[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* by[3:0] */ - r[2] |= bcdec__bitstream_read_bits(&bstream, 5); /* ry[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 2; /* bz[2] */ - r[3] |= bcdec__bitstream_read_bits(&bstream, 5); /* rz[4:0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 3; /* bz[3] */ - partition = bcdec__bitstream_read_bits(&bstream, 5); /* d[4:0] */ - mode = 8; - } break; - - /* mode 10 */ - case 0b11110: { - /* Partitition indices: 46 bits - Partition: 5 bits - Color Endpoints: 72 bits (6666, 6666, 6666) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 6); /* rw[5:0] */ - g[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gz[4] */ - b[3] |= bcdec__bitstream_read_bit(&bstream); /* bz[0] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 1; /* bz[1] */ - b[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* by[4] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 6); /* gw[5:0] */ - g[2] |= bcdec__bitstream_read_bit(&bstream) << 5; /* gy[5] */ - b[2] |= bcdec__bitstream_read_bit(&bstream) << 5; /* by[5] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 2; /* bz[2] */ - g[2] |= bcdec__bitstream_read_bit(&bstream) << 4; /* gy[4] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 6); /* bw[5:0] */ - g[3] |= bcdec__bitstream_read_bit(&bstream) << 5; /* gz[5] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 3; /* bz[3] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 5; /* bz[5] */ - b[3] |= bcdec__bitstream_read_bit(&bstream) << 4; /* bz[4] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 6); /* rx[5:0] */ - g[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* gy[3:0] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 6); /* gx[5:0] */ - g[3] |= bcdec__bitstream_read_bits(&bstream, 4); /* gz[3:0] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 6); /* bx[5:0] */ - b[2] |= bcdec__bitstream_read_bits(&bstream, 4); /* by[3:0] */ - r[2] |= bcdec__bitstream_read_bits(&bstream, 6); /* ry[5:0] */ - r[3] |= bcdec__bitstream_read_bits(&bstream, 6); /* rz[5:0] */ - partition = bcdec__bitstream_read_bits(&bstream, 5); /* d[4:0] */ - mode = 9; - } break; - - /* mode 11 */ - case 0b00011: { - /* Partitition indices: 63 bits - Partition: 0 bits - Color Endpoints: 60 bits (10.10, 10.10, 10.10) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* rw[9:0] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* gw[9:0] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* bw[9:0] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 10); /* rx[9:0] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 10); /* gx[9:0] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 10); /* bx[9:0] */ - mode = 10; - } break; - - /* mode 12 */ - case 0b00111: { - /* Partitition indices: 63 bits - Partition: 0 bits - Color Endpoints: 60 bits (11.9, 11.9, 11.9) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* rw[9:0] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* gw[9:0] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* bw[9:0] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 9); /* rx[8:0] */ - r[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* rw[10] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 9); /* gx[8:0] */ - g[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* gw[10] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 9); /* bx[8:0] */ - b[0] |= bcdec__bitstream_read_bit(&bstream) << 10; /* bw[10] */ - mode = 11; - } break; - - /* mode 13 */ - case 0b01011: { - /* Partitition indices: 63 bits - Partition: 0 bits - Color Endpoints: 60 bits (12.8, 12.8, 12.8) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* rw[9:0] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* gw[9:0] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* bw[9:0] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 8); /* rx[7:0] */ - r[0] |= bcdec__bitstream_read_bits_r(&bstream, 2) << 10;/* rx[10:11] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 8); /* gx[7:0] */ - g[0] |= bcdec__bitstream_read_bits_r(&bstream, 2) << 10;/* gx[10:11] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 8); /* bx[7:0] */ - b[0] |= bcdec__bitstream_read_bits_r(&bstream, 2) << 10;/* bx[10:11] */ - mode = 12; - } break; - - /* mode 14 */ - case 0b01111: { - /* Partitition indices: 63 bits - Partition: 0 bits - Color Endpoints: 60 bits (16.4, 16.4, 16.4) */ - r[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* rw[9:0] */ - g[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* gw[9:0] */ - b[0] |= bcdec__bitstream_read_bits(&bstream, 10); /* bw[9:0] */ - r[1] |= bcdec__bitstream_read_bits(&bstream, 4); /* rx[3:0] */ - r[0] |= bcdec__bitstream_read_bits_r(&bstream, 6) << 10;/* rw[10:15] */ - g[1] |= bcdec__bitstream_read_bits(&bstream, 4); /* gx[3:0] */ - g[0] |= bcdec__bitstream_read_bits_r(&bstream, 6) << 10;/* gw[10:15] */ - b[1] |= bcdec__bitstream_read_bits(&bstream, 4); /* bx[3:0] */ - b[0] |= bcdec__bitstream_read_bits_r(&bstream, 6) << 10;/* bw[10:15] */ - mode = 13; - } break; - - default: { - /* Modes 10011, 10111, 11011, and 11111 (not shown) are reserved. - Do not use these in your encoder. If the hardware is passed blocks - with one of these modes specified, the resulting decompressed block - must contain all zeroes in all channels except for the alpha channel. */ - for (i = 0; i < 4; ++i) { - for (j = 0; j < 4; ++j) { - decompressed[j * 3 + 0] = 0; - decompressed[j * 3 + 1] = 0; - decompressed[j * 3 + 2] = 0; - } - decompressed += destinationPitch; - } - - return; - } - } - - numPartitions = (mode >= 10) ? 0 : 1; - - actualBits0Mode = actual_bits_count[0][mode]; - if (isSigned) { - r[0] = bcdec__extend_sign(r[0], actualBits0Mode); - g[0] = bcdec__extend_sign(g[0], actualBits0Mode); - b[0] = bcdec__extend_sign(b[0], actualBits0Mode); - } - - /* Mode 11 (like Mode 10) does not use delta compression, - and instead stores both color endpoints explicitly. */ - if ((mode != 9 && mode != 10) || isSigned) { - for (i = 1; i < (numPartitions + 1) * 2; ++i) { - r[i] = bcdec__extend_sign(r[i], actual_bits_count[1][mode]); - g[i] = bcdec__extend_sign(g[i], actual_bits_count[2][mode]); - b[i] = bcdec__extend_sign(b[i], actual_bits_count[3][mode]); - } - } - - if (mode != 9 && mode != 10) { - for (i = 1; i < (numPartitions + 1) * 2; ++i) { - r[i] = bcdec__transform_inverse(r[i], r[0], actualBits0Mode, isSigned); - g[i] = bcdec__transform_inverse(g[i], g[0], actualBits0Mode, isSigned); - b[i] = bcdec__transform_inverse(b[i], b[0], actualBits0Mode, isSigned); - } - } - - for (i = 0; i < (numPartitions + 1) * 2; ++i) { - r[i] = bcdec__unquantize(r[i], actualBits0Mode, isSigned); - g[i] = bcdec__unquantize(g[i], actualBits0Mode, isSigned); - b[i] = bcdec__unquantize(b[i], actualBits0Mode, isSigned); - } - - weights = (mode >= 10) ? aWeight4 : aWeight3; - for (i = 0; i < 4; ++i) { - for (j = 0; j < 4; ++j) { - partitionSet = (mode >= 10) ? ((i|j) ? 0 : 128) : partition_sets[partition][i][j]; - - indexBits = (mode >= 10) ? 4 : 3; - /* fix-up index is specified with one less bit */ - /* The fix-up index for subset 0 is always index 0 */ - if (partitionSet & 0x80) { - indexBits--; - } - partitionSet &= 0x01; - - index = bcdec__bitstream_read_bits(&bstream, indexBits); - - ep_i = partitionSet * 2; - decompressed[j * 3 + 0] = bcdec__finish_unquantize( - bcdec__interpolate(r[ep_i], r[ep_i+1], weights, index), isSigned); - decompressed[j * 3 + 1] = bcdec__finish_unquantize( - bcdec__interpolate(g[ep_i], g[ep_i+1], weights, index), isSigned); - decompressed[j * 3 + 2] = bcdec__finish_unquantize( - bcdec__interpolate(b[ep_i], b[ep_i+1], weights, index), isSigned); - } - - decompressed += destinationPitch; - } -} - -BCDECDEF void bcdec_bc6h_float(const void* compressedBlock, void* decompressedBlock, int destinationPitch, int isSigned) { - unsigned short block[16*3]; - float* decompressed; - const unsigned short* b; - int i, j; - - bcdec_bc6h_half(compressedBlock, block, 4*3, isSigned); - b = block; - decompressed = (float*)decompressedBlock; - for (i = 0; i < 4; ++i) { - for (j = 0; j < 4; ++j) { - decompressed[j * 3 + 0] = bcdec__half_to_float_quick(*b++); - decompressed[j * 3 + 1] = bcdec__half_to_float_quick(*b++); - decompressed[j * 3 + 2] = bcdec__half_to_float_quick(*b++); - } - decompressed += destinationPitch; - } -} - -static void bcdec__swap_values(int* a, int* b) { - a[0] ^= b[0], b[0] ^= a[0], a[0] ^= b[0]; -} - -BCDECDEF void bcdec_bc7(const void* compressedBlock, void* decompressedBlock, int destinationPitch) { - static char actual_bits_count[2][8] = { - { 4, 6, 5, 7, 5, 7, 7, 5 }, /* RGBA */ - { 0, 0, 0, 0, 6, 8, 7, 5 }, /* Alpha */ - }; - - /* There are 64 possible partition sets for a two-region tile. - Each 4x4 block represents a single shape. - Here also every fix-up index has MSB bit set. */ - static unsigned char partition_sets[2][64][4][4] = { - { /* Partition table for 2-subset BPTC */ - { {128, 0, 1, 1}, {0, 0, 1, 1}, { 0, 0, 1, 1}, {0, 0, 1, 129} }, /* 0 */ - { {128, 0, 0, 1}, {0, 0, 0, 1}, { 0, 0, 0, 1}, {0, 0, 0, 129} }, /* 1 */ - { {128, 1, 1, 1}, {0, 1, 1, 1}, { 0, 1, 1, 1}, {0, 1, 1, 129} }, /* 2 */ - { {128, 0, 0, 1}, {0, 0, 1, 1}, { 0, 0, 1, 1}, {0, 1, 1, 129} }, /* 3 */ - { {128, 0, 0, 0}, {0, 0, 0, 1}, { 0, 0, 0, 1}, {0, 0, 1, 129} }, /* 4 */ - { {128, 0, 1, 1}, {0, 1, 1, 1}, { 0, 1, 1, 1}, {1, 1, 1, 129} }, /* 5 */ - { {128, 0, 0, 1}, {0, 0, 1, 1}, { 0, 1, 1, 1}, {1, 1, 1, 129} }, /* 6 */ - { {128, 0, 0, 0}, {0, 0, 0, 1}, { 0, 0, 1, 1}, {0, 1, 1, 129} }, /* 7 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, { 0, 0, 0, 1}, {0, 0, 1, 129} }, /* 8 */ - { {128, 0, 1, 1}, {0, 1, 1, 1}, { 1, 1, 1, 1}, {1, 1, 1, 129} }, /* 9 */ - { {128, 0, 0, 0}, {0, 0, 0, 1}, { 0, 1, 1, 1}, {1, 1, 1, 129} }, /* 10 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, { 0, 0, 0, 1}, {0, 1, 1, 129} }, /* 11 */ - { {128, 0, 0, 1}, {0, 1, 1, 1}, { 1, 1, 1, 1}, {1, 1, 1, 129} }, /* 12 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, { 1, 1, 1, 1}, {1, 1, 1, 129} }, /* 13 */ - { {128, 0, 0, 0}, {1, 1, 1, 1}, { 1, 1, 1, 1}, {1, 1, 1, 129} }, /* 14 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, { 0, 0, 0, 0}, {1, 1, 1, 129} }, /* 15 */ - { {128, 0, 0, 0}, {1, 0, 0, 0}, { 1, 1, 1, 0}, {1, 1, 1, 129} }, /* 16 */ - { {128, 1, 129, 1}, {0, 0, 0, 1}, { 0, 0, 0, 0}, {0, 0, 0, 0} }, /* 17 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, {129, 0, 0, 0}, {1, 1, 1, 0} }, /* 18 */ - { {128, 1, 129, 1}, {0, 0, 1, 1}, { 0, 0, 0, 1}, {0, 0, 0, 0} }, /* 19 */ - { {128, 0, 129, 1}, {0, 0, 0, 1}, { 0, 0, 0, 0}, {0, 0, 0, 0} }, /* 20 */ - { {128, 0, 0, 0}, {1, 0, 0, 0}, {129, 1, 0, 0}, {1, 1, 1, 0} }, /* 21 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, {129, 0, 0, 0}, {1, 1, 0, 0} }, /* 22 */ - { {128, 1, 1, 1}, {0, 0, 1, 1}, { 0, 0, 1, 1}, {0, 0, 0, 129} }, /* 23 */ - { {128, 0, 129, 1}, {0, 0, 0, 1}, { 0, 0, 0, 1}, {0, 0, 0, 0} }, /* 24 */ - { {128, 0, 0, 0}, {1, 0, 0, 0}, {129, 0, 0, 0}, {1, 1, 0, 0} }, /* 25 */ - { {128, 1, 129, 0}, {0, 1, 1, 0}, { 0, 1, 1, 0}, {0, 1, 1, 0} }, /* 26 */ - { {128, 0, 129, 1}, {0, 1, 1, 0}, { 0, 1, 1, 0}, {1, 1, 0, 0} }, /* 27 */ - { {128, 0, 0, 1}, {0, 1, 1, 1}, {129, 1, 1, 0}, {1, 0, 0, 0} }, /* 28 */ - { {128, 0, 0, 0}, {1, 1, 1, 1}, {129, 1, 1, 1}, {0, 0, 0, 0} }, /* 29 */ - { {128, 1, 129, 1}, {0, 0, 0, 1}, { 1, 0, 0, 0}, {1, 1, 1, 0} }, /* 30 */ - { {128, 0, 129, 1}, {1, 0, 0, 1}, { 1, 0, 0, 1}, {1, 1, 0, 0} }, /* 31 */ - { {128, 1, 0, 1}, {0, 1, 0, 1}, { 0, 1, 0, 1}, {0, 1, 0, 129} }, /* 32 */ - { {128, 0, 0, 0}, {1, 1, 1, 1}, { 0, 0, 0, 0}, {1, 1, 1, 129} }, /* 33 */ - { {128, 1, 0, 1}, {1, 0, 129, 0}, { 0, 1, 0, 1}, {1, 0, 1, 0} }, /* 34 */ - { {128, 0, 1, 1}, {0, 0, 1, 1}, {129, 1, 0, 0}, {1, 1, 0, 0} }, /* 35 */ - { {128, 0, 129, 1}, {1, 1, 0, 0}, { 0, 0, 1, 1}, {1, 1, 0, 0} }, /* 36 */ - { {128, 1, 0, 1}, {0, 1, 0, 1}, {129, 0, 1, 0}, {1, 0, 1, 0} }, /* 37 */ - { {128, 1, 1, 0}, {1, 0, 0, 1}, { 0, 1, 1, 0}, {1, 0, 0, 129} }, /* 38 */ - { {128, 1, 0, 1}, {1, 0, 1, 0}, { 1, 0, 1, 0}, {0, 1, 0, 129} }, /* 39 */ - { {128, 1, 129, 1}, {0, 0, 1, 1}, { 1, 1, 0, 0}, {1, 1, 1, 0} }, /* 40 */ - { {128, 0, 0, 1}, {0, 0, 1, 1}, {129, 1, 0, 0}, {1, 0, 0, 0} }, /* 41 */ - { {128, 0, 129, 1}, {0, 0, 1, 0}, { 0, 1, 0, 0}, {1, 1, 0, 0} }, /* 42 */ - { {128, 0, 129, 1}, {1, 0, 1, 1}, { 1, 1, 0, 1}, {1, 1, 0, 0} }, /* 43 */ - { {128, 1, 129, 0}, {1, 0, 0, 1}, { 1, 0, 0, 1}, {0, 1, 1, 0} }, /* 44 */ - { {128, 0, 1, 1}, {1, 1, 0, 0}, { 1, 1, 0, 0}, {0, 0, 1, 129} }, /* 45 */ - { {128, 1, 1, 0}, {0, 1, 1, 0}, { 1, 0, 0, 1}, {1, 0, 0, 129} }, /* 46 */ - { {128, 0, 0, 0}, {0, 1, 129, 0}, { 0, 1, 1, 0}, {0, 0, 0, 0} }, /* 47 */ - { {128, 1, 0, 0}, {1, 1, 129, 0}, { 0, 1, 0, 0}, {0, 0, 0, 0} }, /* 48 */ - { {128, 0, 129, 0}, {0, 1, 1, 1}, { 0, 0, 1, 0}, {0, 0, 0, 0} }, /* 49 */ - { {128, 0, 0, 0}, {0, 0, 129, 0}, { 0, 1, 1, 1}, {0, 0, 1, 0} }, /* 50 */ - { {128, 0, 0, 0}, {0, 1, 0, 0}, {129, 1, 1, 0}, {0, 1, 0, 0} }, /* 51 */ - { {128, 1, 1, 0}, {1, 1, 0, 0}, { 1, 0, 0, 1}, {0, 0, 1, 129} }, /* 52 */ - { {128, 0, 1, 1}, {0, 1, 1, 0}, { 1, 1, 0, 0}, {1, 0, 0, 129} }, /* 53 */ - { {128, 1, 129, 0}, {0, 0, 1, 1}, { 1, 0, 0, 1}, {1, 1, 0, 0} }, /* 54 */ - { {128, 0, 129, 1}, {1, 0, 0, 1}, { 1, 1, 0, 0}, {0, 1, 1, 0} }, /* 55 */ - { {128, 1, 1, 0}, {1, 1, 0, 0}, { 1, 1, 0, 0}, {1, 0, 0, 129} }, /* 56 */ - { {128, 1, 1, 0}, {0, 0, 1, 1}, { 0, 0, 1, 1}, {1, 0, 0, 129} }, /* 57 */ - { {128, 1, 1, 1}, {1, 1, 1, 0}, { 1, 0, 0, 0}, {0, 0, 0, 129} }, /* 58 */ - { {128, 0, 0, 1}, {1, 0, 0, 0}, { 1, 1, 1, 0}, {0, 1, 1, 129} }, /* 59 */ - { {128, 0, 0, 0}, {1, 1, 1, 1}, { 0, 0, 1, 1}, {0, 0, 1, 129} }, /* 60 */ - { {128, 0, 129, 1}, {0, 0, 1, 1}, { 1, 1, 1, 1}, {0, 0, 0, 0} }, /* 61 */ - { {128, 0, 129, 0}, {0, 0, 1, 0}, { 1, 1, 1, 0}, {1, 1, 1, 0} }, /* 62 */ - { {128, 1, 0, 0}, {0, 1, 0, 0}, { 0, 1, 1, 1}, {0, 1, 1, 129} } /* 63 */ - }, - { /* Partition table for 3-subset BPTC */ - { {128, 0, 1, 129}, {0, 0, 1, 1}, { 0, 2, 2, 1}, { 2, 2, 2, 130} }, /* 0 */ - { {128, 0, 0, 129}, {0, 0, 1, 1}, {130, 2, 1, 1}, { 2, 2, 2, 1} }, /* 1 */ - { {128, 0, 0, 0}, {2, 0, 0, 1}, {130, 2, 1, 1}, { 2, 2, 1, 129} }, /* 2 */ - { {128, 2, 2, 130}, {0, 0, 2, 2}, { 0, 0, 1, 1}, { 0, 1, 1, 129} }, /* 3 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, {129, 1, 2, 2}, { 1, 1, 2, 130} }, /* 4 */ - { {128, 0, 1, 129}, {0, 0, 1, 1}, { 0, 0, 2, 2}, { 0, 0, 2, 130} }, /* 5 */ - { {128, 0, 2, 130}, {0, 0, 2, 2}, { 1, 1, 1, 1}, { 1, 1, 1, 129} }, /* 6 */ - { {128, 0, 1, 1}, {0, 0, 1, 1}, {130, 2, 1, 1}, { 2, 2, 1, 129} }, /* 7 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, {129, 1, 1, 1}, { 2, 2, 2, 130} }, /* 8 */ - { {128, 0, 0, 0}, {1, 1, 1, 1}, {129, 1, 1, 1}, { 2, 2, 2, 130} }, /* 9 */ - { {128, 0, 0, 0}, {1, 1, 129, 1}, { 2, 2, 2, 2}, { 2, 2, 2, 130} }, /* 10 */ - { {128, 0, 1, 2}, {0, 0, 129, 2}, { 0, 0, 1, 2}, { 0, 0, 1, 130} }, /* 11 */ - { {128, 1, 1, 2}, {0, 1, 129, 2}, { 0, 1, 1, 2}, { 0, 1, 1, 130} }, /* 12 */ - { {128, 1, 2, 2}, {0, 129, 2, 2}, { 0, 1, 2, 2}, { 0, 1, 2, 130} }, /* 13 */ - { {128, 0, 1, 129}, {0, 1, 1, 2}, { 1, 1, 2, 2}, { 1, 2, 2, 130} }, /* 14 */ - { {128, 0, 1, 129}, {2, 0, 0, 1}, {130, 2, 0, 0}, { 2, 2, 2, 0} }, /* 15 */ - { {128, 0, 0, 129}, {0, 0, 1, 1}, { 0, 1, 1, 2}, { 1, 1, 2, 130} }, /* 16 */ - { {128, 1, 1, 129}, {0, 0, 1, 1}, {130, 0, 0, 1}, { 2, 2, 0, 0} }, /* 17 */ - { {128, 0, 0, 0}, {1, 1, 2, 2}, {129, 1, 2, 2}, { 1, 1, 2, 130} }, /* 18 */ - { {128, 0, 2, 130}, {0, 0, 2, 2}, { 0, 0, 2, 2}, { 1, 1, 1, 129} }, /* 19 */ - { {128, 1, 1, 129}, {0, 1, 1, 1}, { 0, 2, 2, 2}, { 0, 2, 2, 130} }, /* 20 */ - { {128, 0, 0, 129}, {0, 0, 0, 1}, {130, 2, 2, 1}, { 2, 2, 2, 1} }, /* 21 */ - { {128, 0, 0, 0}, {0, 0, 129, 1}, { 0, 1, 2, 2}, { 0, 1, 2, 130} }, /* 22 */ - { {128, 0, 0, 0}, {1, 1, 0, 0}, {130, 2, 129, 0}, { 2, 2, 1, 0} }, /* 23 */ - { {128, 1, 2, 130}, {0, 129, 2, 2}, { 0, 0, 1, 1}, { 0, 0, 0, 0} }, /* 24 */ - { {128, 0, 1, 2}, {0, 0, 1, 2}, {129, 1, 2, 2}, { 2, 2, 2, 130} }, /* 25 */ - { {128, 1, 1, 0}, {1, 2, 130, 1}, {129, 2, 2, 1}, { 0, 1, 1, 0} }, /* 26 */ - { {128, 0, 0, 0}, {0, 1, 129, 0}, { 1, 2, 130, 1}, { 1, 2, 2, 1} }, /* 27 */ - { {128, 0, 2, 2}, {1, 1, 0, 2}, {129, 1, 0, 2}, { 0, 0, 2, 130} }, /* 28 */ - { {128, 1, 1, 0}, {0, 129, 1, 0}, { 2, 0, 0, 2}, { 2, 2, 2, 130} }, /* 29 */ - { {128, 0, 1, 1}, {0, 1, 2, 2}, { 0, 1, 130, 2}, { 0, 0, 1, 129} }, /* 30 */ - { {128, 0, 0, 0}, {2, 0, 0, 0}, {130, 2, 1, 1}, { 2, 2, 2, 129} }, /* 31 */ - { {128, 0, 0, 0}, {0, 0, 0, 2}, {129, 1, 2, 2}, { 1, 2, 2, 130} }, /* 32 */ - { {128, 2, 2, 130}, {0, 0, 2, 2}, { 0, 0, 1, 2}, { 0, 0, 1, 129} }, /* 33 */ - { {128, 0, 1, 129}, {0, 0, 1, 2}, { 0, 0, 2, 2}, { 0, 2, 2, 130} }, /* 34 */ - { {128, 1, 2, 0}, {0, 129, 2, 0}, { 0, 1, 130, 0}, { 0, 1, 2, 0} }, /* 35 */ - { {128, 0, 0, 0}, {1, 1, 129, 1}, { 2, 2, 130, 2}, { 0, 0, 0, 0} }, /* 36 */ - { {128, 1, 2, 0}, {1, 2, 0, 1}, {130, 0, 129, 2}, { 0, 1, 2, 0} }, /* 37 */ - { {128, 1, 2, 0}, {2, 0, 1, 2}, {129, 130, 0, 1}, { 0, 1, 2, 0} }, /* 38 */ - { {128, 0, 1, 1}, {2, 2, 0, 0}, { 1, 1, 130, 2}, { 0, 0, 1, 129} }, /* 39 */ - { {128, 0, 1, 1}, {1, 1, 130, 2}, { 2, 2, 0, 0}, { 0, 0, 1, 129} }, /* 40 */ - { {128, 1, 0, 129}, {0, 1, 0, 1}, { 2, 2, 2, 2}, { 2, 2, 2, 130} }, /* 41 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, {130, 1, 2, 1}, { 2, 1, 2, 129} }, /* 42 */ - { {128, 0, 2, 2}, {1, 129, 2, 2}, { 0, 0, 2, 2}, { 1, 1, 2, 130} }, /* 43 */ - { {128, 0, 2, 130}, {0, 0, 1, 1}, { 0, 0, 2, 2}, { 0, 0, 1, 129} }, /* 44 */ - { {128, 2, 2, 0}, {1, 2, 130, 1}, { 0, 2, 2, 0}, { 1, 2, 2, 129} }, /* 45 */ - { {128, 1, 0, 1}, {2, 2, 130, 2}, { 2, 2, 2, 2}, { 0, 1, 0, 129} }, /* 46 */ - { {128, 0, 0, 0}, {2, 1, 2, 1}, {130, 1, 2, 1}, { 2, 1, 2, 129} }, /* 47 */ - { {128, 1, 0, 129}, {0, 1, 0, 1}, { 0, 1, 0, 1}, { 2, 2, 2, 130} }, /* 48 */ - { {128, 2, 2, 130}, {0, 1, 1, 1}, { 0, 2, 2, 2}, { 0, 1, 1, 129} }, /* 49 */ - { {128, 0, 0, 2}, {1, 129, 1, 2}, { 0, 0, 0, 2}, { 1, 1, 1, 130} }, /* 50 */ - { {128, 0, 0, 0}, {2, 129, 1, 2}, { 2, 1, 1, 2}, { 2, 1, 1, 130} }, /* 51 */ - { {128, 2, 2, 2}, {0, 129, 1, 1}, { 0, 1, 1, 1}, { 0, 2, 2, 130} }, /* 52 */ - { {128, 0, 0, 2}, {1, 1, 1, 2}, {129, 1, 1, 2}, { 0, 0, 0, 130} }, /* 53 */ - { {128, 1, 1, 0}, {0, 129, 1, 0}, { 0, 1, 1, 0}, { 2, 2, 2, 130} }, /* 54 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, { 2, 1, 129, 2}, { 2, 1, 1, 130} }, /* 55 */ - { {128, 1, 1, 0}, {0, 129, 1, 0}, { 2, 2, 2, 2}, { 2, 2, 2, 130} }, /* 56 */ - { {128, 0, 2, 2}, {0, 0, 1, 1}, { 0, 0, 129, 1}, { 0, 0, 2, 130} }, /* 57 */ - { {128, 0, 2, 2}, {1, 1, 2, 2}, {129, 1, 2, 2}, { 0, 0, 2, 130} }, /* 58 */ - { {128, 0, 0, 0}, {0, 0, 0, 0}, { 0, 0, 0, 0}, { 2, 129, 1, 130} }, /* 59 */ - { {128, 0, 0, 130}, {0, 0, 0, 1}, { 0, 0, 0, 2}, { 0, 0, 0, 129} }, /* 60 */ - { {128, 2, 2, 2}, {1, 2, 2, 2}, { 0, 2, 2, 2}, {129, 2, 2, 130} }, /* 61 */ - { {128, 1, 0, 129}, {2, 2, 2, 2}, { 2, 2, 2, 2}, { 2, 2, 2, 130} }, /* 62 */ - { {128, 1, 1, 129}, {2, 0, 1, 1}, {130, 2, 0, 1}, { 2, 2, 2, 0} } /* 63 */ - } - }; - - static int aWeight2[] = { 0, 21, 43, 64 }; - static int aWeight3[] = { 0, 9, 18, 27, 37, 46, 55, 64 }; - static int aWeight4[] = { 0, 4, 9, 13, 17, 21, 26, 30, 34, 38, 43, 47, 51, 55, 60, 64 }; - - static unsigned char sModeHasPBits = 0b11001011; - - bcdec__bitstream_t bstream; - int mode, partition, numPartitions, numEndpoints, i, j, k, rotation, partitionSet; - int indexSelectionBit, indexBits, indexBits2, index, index2; - int endpoints[6][4]; - char indices[4][4]; - int r, g, b, a; - int* weights, * weights2; - unsigned char* decompressed; - - decompressed = (unsigned char*)decompressedBlock; - - bstream.low = bcdec__getu64le(compressedBlock); - bstream.high = bcdec__getu64le((const unsigned char*)compressedBlock + 8); - - for (mode = 0; mode < 8 && (0 == bcdec__bitstream_read_bit(&bstream)); ++mode); - - /* unexpected mode, clear the block (transparent black) */ - if (mode >= 8) { - for (i = 0; i < 4; ++i) { - for (j = 0; j < 4; ++j) { - decompressed[j * 4 + 0] = 0; - decompressed[j * 4 + 1] = 0; - decompressed[j * 4 + 2] = 0; - decompressed[j * 4 + 3] = 0; - } - decompressed += destinationPitch; - } - - return; - } - - partition = 0; - numPartitions = 1; - rotation = 0; - indexSelectionBit = 0; - - if (mode == 0 || mode == 1 || mode == 2 || mode == 3 || mode == 7) { - numPartitions = (mode == 0 || mode == 2) ? 3 : 2; - partition = bcdec__bitstream_read_bits(&bstream, (mode == 0) ? 4 : 6); - } - - numEndpoints = numPartitions * 2; - - if (mode == 4 || mode == 5) { - rotation = bcdec__bitstream_read_bits(&bstream, 2); - - if (mode == 4) { - indexSelectionBit = bcdec__bitstream_read_bit(&bstream); - } - } - - /* Extract endpoints */ - /* RGB */ - for (i = 0; i < 3; ++i) { - for (j = 0; j < numEndpoints; ++j) { - endpoints[j][i] = bcdec__bitstream_read_bits(&bstream, actual_bits_count[0][mode]); - } - } - /* Alpha (if any) */ - if (actual_bits_count[1][mode] > 0) { - for (j = 0; j < numEndpoints; ++j) { - endpoints[j][3] = bcdec__bitstream_read_bits(&bstream, actual_bits_count[1][mode]); - } - } - - /* Fully decode endpoints */ - /* First handle modes that have P-bits */ - if (mode == 0 || mode == 1 || mode == 3 || mode == 6 || mode == 7) { - for (i = 0; i < numEndpoints; ++i) { - /* component-wise left-shift */ - for (j = 0; j < 4; ++j) { - endpoints[i][j] <<= 1; - } - } - - /* if P-bit is shared */ - if (mode == 1) { - i = bcdec__bitstream_read_bit(&bstream); - j = bcdec__bitstream_read_bit(&bstream); - - /* rgb component-wise insert pbits */ - for (k = 0; k < 3; ++k) { - endpoints[0][k] |= i; - endpoints[1][k] |= i; - endpoints[2][k] |= j; - endpoints[3][k] |= j; - } - } else if (sModeHasPBits & (1 << mode)) { - /* unique P-bit per endpoint */ - for (i = 0; i < numEndpoints; ++i) { - j = bcdec__bitstream_read_bit(&bstream); - for (k = 0; k < 4; ++k) { - endpoints[i][k] |= j; - } - } - } - } - - for (i = 0; i < numEndpoints; ++i) { - /* get color components precision including pbit */ - j = actual_bits_count[0][mode] + ((sModeHasPBits >> mode) & 1); - - for (k = 0; k < 3; ++k) { - /* left shift endpoint components so that their MSB lies in bit 7 */ - endpoints[i][k] = endpoints[i][k] << (8 - j); - /* Replicate each component's MSB into the LSBs revealed by the left-shift operation above */ - endpoints[i][k] = endpoints[i][k] | (endpoints[i][k] >> j); - } - - /* get alpha component precision including pbit */ - j = actual_bits_count[1][mode] + ((sModeHasPBits >> mode) & 1); - - /* left shift endpoint components so that their MSB lies in bit 7 */ - endpoints[i][3] = endpoints[i][3] << (8 - j); - /* Replicate each component's MSB into the LSBs revealed by the left-shift operation above */ - endpoints[i][3] = endpoints[i][3] | (endpoints[i][3] >> j); - } - - /* If this mode does not explicitly define the alpha component */ - /* set alpha equal to 1.0 */ - if (!actual_bits_count[1][mode]) { - for (j = 0; j < numEndpoints; ++j) { - endpoints[j][3] = 0xFF; - } - } - - /* Determine weights tables */ - indexBits = (mode == 0 || mode == 1) ? 3 : ((mode == 6) ? 4 : 2); - indexBits2 = (mode == 4) ? 3 : ((mode == 5) ? 2 : 0); - weights = (indexBits == 2) ? aWeight2 : ((indexBits == 3) ? aWeight3 : aWeight4); - weights2 = (indexBits2 == 2) ? aWeight2 : aWeight3; - - /* Quite inconvenient that indices aren't interleaved so we have to make 2 passes here */ - /* Pass #1: collecting color indices */ - for (i = 0; i < 4; ++i) { - for (j = 0; j < 4; ++j) { - partitionSet = (numPartitions == 1) ? ((i | j) ? 0 : 128) : partition_sets[numPartitions - 2][partition][i][j]; - - indexBits = (mode == 0 || mode == 1) ? 3 : ((mode == 6) ? 4 : 2); - /* fix-up index is specified with one less bit */ - /* The fix-up index for subset 0 is always index 0 */ - if (partitionSet & 0x80) { - indexBits--; - } - - indices[i][j] = bcdec__bitstream_read_bits(&bstream, indexBits); - } - } - - /* Pass #2: reading alpha indices (if any) and interpolating & rotating */ - for (i = 0; i < 4; ++i) { - for (j = 0; j < 4; ++j) { - partitionSet = (numPartitions == 1) ? ((i|j) ? 0 : 128) : partition_sets[numPartitions - 2][partition][i][j]; - partitionSet &= 0x03; - - index = indices[i][j]; - - if (!indexBits2) { - r = bcdec__interpolate(endpoints[partitionSet * 2][0], endpoints[partitionSet * 2 + 1][0], weights, index); - g = bcdec__interpolate(endpoints[partitionSet * 2][1], endpoints[partitionSet * 2 + 1][1], weights, index); - b = bcdec__interpolate(endpoints[partitionSet * 2][2], endpoints[partitionSet * 2 + 1][2], weights, index); - a = bcdec__interpolate(endpoints[partitionSet * 2][3], endpoints[partitionSet * 2 + 1][3], weights, index); - } else { - index2 = bcdec__bitstream_read_bits(&bstream, (i|j) ? indexBits2 : (indexBits2 - 1)); - /* The index value for interpolating color comes from the secondary index bits for the texel - if the mode has an index selection bit and its value is one, and from the primary index bits otherwise. - The alpha index comes from the secondary index bits if the block has a secondary index and - the block either doesn’t have an index selection bit or that bit is zero, and from the primary index bits otherwise. */ - if (!indexSelectionBit) { - r = bcdec__interpolate(endpoints[partitionSet * 2][0], endpoints[partitionSet * 2 + 1][0], weights, index); - g = bcdec__interpolate(endpoints[partitionSet * 2][1], endpoints[partitionSet * 2 + 1][1], weights, index); - b = bcdec__interpolate(endpoints[partitionSet * 2][2], endpoints[partitionSet * 2 + 1][2], weights, index); - a = bcdec__interpolate(endpoints[partitionSet * 2][3], endpoints[partitionSet * 2 + 1][3], weights2, index2); - } else { - r = bcdec__interpolate(endpoints[partitionSet * 2][0], endpoints[partitionSet * 2 + 1][0], weights2, index2); - g = bcdec__interpolate(endpoints[partitionSet * 2][1], endpoints[partitionSet * 2 + 1][1], weights2, index2); - b = bcdec__interpolate(endpoints[partitionSet * 2][2], endpoints[partitionSet * 2 + 1][2], weights2, index2); - a = bcdec__interpolate(endpoints[partitionSet * 2][3], endpoints[partitionSet * 2 + 1][3], weights, index); - } - } - - switch (rotation) { - case 1: { /* 01 – Block format is Scalar(R) Vector(AGB) - swap A and R */ - bcdec__swap_values(&a, &r); - } break; - case 2: { /* 10 – Block format is Scalar(G) Vector(RAB) - swap A and G */ - bcdec__swap_values(&a, &g); - } break; - case 3: { /* 11 - Block format is Scalar(B) Vector(RGA) - swap A and B */ - bcdec__swap_values(&a, &b); - } break; - } - - decompressed[j * 4 + 0] = r; - decompressed[j * 4 + 1] = g; - decompressed[j * 4 + 2] = b; - decompressed[j * 4 + 3] = a; - } - - decompressed += destinationPitch; - } -} - -#endif /* BCDEC_IMPLEMENTATION */ - -/* LICENSE: - -This software is available under 2 licenses -- choose whichever you prefer. - ------------------------------------------------------------------------------- -ALTERNATIVE A - MIT License - -Copyright (c) 2022 Sergii Kudlai - -Permission is hereby granted, free of charge, to any person obtaining a copy of -this software and associated documentation files (the "Software"), to deal in -the Software without restriction, including without limitation the rights to -use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies -of the Software, and to permit persons to whom the Software is furnished to do -so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in all -copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -SOFTWARE. - ------------------------------------------------------------------------------- -ALTERNATIVE B - The Unlicense - -This is free and unencumbered software released into the public domain. - -Anyone is free to copy, modify, publish, use, compile, sell, or -distribute this software, either in source code form or as a compiled -binary, for any purpose, commercial or non-commercial, and by any -means. - -In jurisdictions that recognize copyright laws, the author or authors -of this software dedicate any and all copyright interest in the -software to the public domain. We make this dedication for the benefit -of the public at large and to the detriment of our heirs and -successors. We intend this dedication to be an overt act of -relinquishment in perpetuity of all present and future rights to this -software under copyright law. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, -EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF -MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. -IN NO EVENT SHALL THE AUTHORS BE LIABLE FOR ANY CLAIM, DAMAGES OR -OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, -ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR -OTHER DEALINGS IN THE SOFTWARE. - -For more information, please refer to - -*/ diff --git a/docs/agents/04-graphics.md b/docs/agents/04-graphics.md index a129097..7ff789a 100644 --- a/docs/agents/04-graphics.md +++ b/docs/agents/04-graphics.md @@ -131,7 +131,7 @@ So the translator leaves z alone (`QueryGraphicsState` depth mode false in `Prog - **Frame rate in the Shrine of Resurrection on the Mac (M4 Pro)**, after the above: docked (1600×900) GPU about 13 ms, frame p50 16.7 ms (60 fps; p90 17.6 ms), was 20.8 ms before segmenting; handheld (`DM_OPMODE=0`, 1280×720) GPU p50 9.3 ms and a steady 60. Host translation is 8 ms per frame (22,000 instructions per draw), spread over Metal's encoder creation and draw calls, argument encoding, Objective-C retain / release and pool lookups. The per-pass GPU timer (`DM_NVN_PASSGPU`) is switched off under Metal 4 (`pg_render` returns at once), so GPU time cannot yet be attributed to passes. The frame loads 376 MB and stores 432 MB of render-target data (`@rt_load_mb` / `@rt_store_mb`; 325 of 347 passes load and store everything), which is the next thing to look at. - **Fixed: hard-edged dark boxes behind every HUD prompt (Mac and iPhone).** "Move / Controls / Camera" and the "A" action prompt sat on translucent black rectangles that overlapped each other. Each prompt is drawn in two layers: the text and icons, and under them a shadow pane, a 198×40 layout-unit picture pane drawn with `e6a88e12dbae9743` (a NintendoWare layout shader: `out = (tex × c0 + c1) × vertex colour`, here c0 = (0, 0, 0, 1), c1 = 0, vertex alpha 180/255, so the output is black with alpha = texture alpha × 0.71). The texture is a mask the game builds once when the prompt appears (`DM_NVN_TEXDUMP=size=198x40:…` catches that frame): the text is drawn into a 198×40 RGBA8 target, shrunk to 99×20, blurred in six ping-pong passes, enlarged again, and its alpha copied (through a view mapped (A, 1, 1, 1)) into the red channel of a 99×20 R8 texture. The game creates that R8 texture with the builder channel mapping (R, 0, 0, 1), the same per-format mapping it gives every R8 texture (one call site, `main+0xecb0bc`, sets R8 → (R, 0, 0, 1), RG8 → (R, G, 0, 1), and so on), and registers it without a view. Read that way the pane's alpha is always 1: a box the size of the pane. The original Maxwell code of the pane shader takes alpha from the texture's fourth component (TEXS writes all four components to R0–R3 and R3 becomes the output alpha), so the translation is faithful. Switch footage (an "Ⓐ Take" prompt) shows the prompt's own dark bar and the ghosted X / Y / B, but no box around them, so the hardware evidently ends up with the mask as alpha; how NVN gets there is not understood. The profile's new `alpha_from_red` key lists the pane shader, and for it a single-channel texture mapped exactly (R, 0, 0, 1) is bound through a view mapped (R, 0, 0, R) (`tex_alpha_red` in `nvn_mtl.m`). The boxes become soft dark glows around the letters and icons (Mac, frame 7000 of the shrine route, compared with a run without the pane: the difference is a glow hugging the text). `DM_NVN_ALPHA_RED=0` restores the literal reading. - **Fixed: cream cloud rectangles behind HUD prompts (iPhone, after the fix above; Mac during a bar's fade-in).** The mask is a layout *capture*: `PaGuideKeySmoke_00` and about 140 other layouts have a pane with user data `CaptureOn`, `A8`, `CaptureScale` 0.5 and `MultiFilter`, and their materials sample the capture through the placeholder texture `FBLayout_00`. The prompt bar draws three layers per key: a glow `P_Glow_00` (`aa209f57b133e862`, material colour (255, 252, 198), texture A = the 64×64 cloud `Kumo64_00^r` (BC4, mapped R R R 1), texture B = the capture; colour = mix(B, A, A.alpha) = the cloud, alpha = min(A.alpha, B.alpha), additive), then the shadow `P_Sh_00` (`e6a88e12dbae9743`), then the text. With only the shadow shader listed, the glow kept alpha 1 and filled each key's rectangle with cream cloud; before that, the opaque shadow had hidden it. Every layout shader that sampled a capture in the shrine (`e6a88e12`, `aa209f57`, `11bb9006c60d49d6` = the glow with vertex colour, `87d900168952ea49` = a projected capture pane) uses its alpha as a mask, which matches the `A8` tag, so `alpha_from_red` now lists all 26 pixel-shader variants of the layout library (`main`'s BNSH at 0x1d83000) and `tex_alpha_red` only converts render targets (captures always are; asset and CPU-written textures are not, and none of the 649 single-channel layout textures in `Bootup.pack` is mapped R 0 0 1). Checked on the Mac: at the bar's fade-in the old list shows the cloud rectangles and the full list a glow around the words; the steady bar, the location title and the pause / controls screen are unchanged (the controls screen pixel-identical). -- **Fixed: crash at start on iPhones without BC (GitHub issue #4, iPhone 14 Pro / A16).** The game encodes one 32×32 BC4 texture on the GPU, drawing its blocks through an RGBA16UI view (one texel = one 8-byte block). Metal can neither attach a compressed texture nor make an uncompressed view of one, so on BC GPUs that attachment was already left out (red line). Without BC the texture is stored decompressed (R8), the check on the Metal format did not see it as compressed, and the host asked Metal for an R8Unorm → RGBA16Uint view, which aborts the process. `attach`, the pipeline prediction and the sampling views now go by the NVN format: a BC texture is never attached and its raw blocks are not sampled, on any GPU. `DM_TEST_NO_BC=1` reproduces it on the Mac; afterwards the game runs with the same two red lines as with BC. Drawing that texture for real needs an uncompressed proxy texture per block, copied (or, without BC, decoded with the `bcd` kernel) into the texture after the pass. The same report measured the memory cost of no BC: about 3.9 GB instead of 2.8 GB, over iOS's limit on a 6 GB iPhone unless the app has `com.apple.developer.kernel.increased-memory-limit` (about 4.1 GB, and the game then ran for about 200 s); converting BC to ASTC, which every Apple GPU decodes, would remove most of it. +- **BC textures are never attached or viewed raw (GitHub issue #4, reported from an iPhone 14 Pro / A16).** The game encodes one 32×32 BC4 texture on the GPU, drawing its blocks through an RGBA16UI view (one texel = one 8-byte block). Metal can neither attach a compressed texture nor make an uncompressed view of one, so that attachment is left out (red line), and `attach`, the pipeline prediction and the sampling views go by the NVN format. Drawing that texture for real needs an uncompressed proxy texture per block, copied into the texture after the pass. The A16 has no BC formats; the host used to store BC textures decompressed on such GPUs (about 3.9 GB instead of 2.8 GB for this game, over iOS's limit on a 6 GB iPhone), and that path has since been removed: only GPUs with BC are supported. ### Seen in The Legend of Zelda: Link's Awakening