From a2c4893644f733992d1d1794ec1529c2363844ec Mon Sep 17 00:00:00 2001 From: SamG-Coder <13449963+SamG-Coder@users.noreply.github.com> Date: Wed, 23 Sep 2026 08:48:41 +1000 Subject: [PATCH 1/2] Add CUDA WebShader browser encoder and integrated NNTC viewer --- .gitignore | 1 + .gitmodules | 3 + README.md | 9 +++ docs/BROWSER_ENCODER.md | 141 +++++++++++++++++++++++++++++++++++ package-lock.json | 63 ++++++++++++++++ package.json | 13 ++++ scripts/serve-web.mjs | 23 ++++++ tests/web-encoder-gpu.html | 59 +++++++++++++++ tests/web-encoder-gpu.mjs | 24 ++++++ tests/web-encoder-ui.mjs | 41 ++++++++++ tests/web-encoder.test.mjs | 66 +++++++++++++++++ vendor/cuda-webshader | 1 + webgpu/README.md | 6 ++ webgpu/encode.html | 41 ++++++++++ webgpu/encoder/app.js | 76 +++++++++++++++++++ webgpu/encoder/encoder.js | 119 +++++++++++++++++++++++++++++ webgpu/encoder/format.js | 111 ++++++++++++++++++++++++++++ webgpu/encoder/kernels.cu | 148 +++++++++++++++++++++++++++++++++++++ webgpu/encoder/prepare.js | 87 ++++++++++++++++++++++ webgpu/encoder/style.css | 1 + webgpu/encoder/zip.js | 22 ++++++ webgpu/index.html | 1 + webgpu/main.js | 12 +++ 23 files changed, 1068 insertions(+) create mode 100644 .gitmodules create mode 100644 docs/BROWSER_ENCODER.md create mode 100644 package-lock.json create mode 100644 package.json create mode 100644 scripts/serve-web.mjs create mode 100644 tests/web-encoder-gpu.html create mode 100644 tests/web-encoder-gpu.mjs create mode 100644 tests/web-encoder-ui.mjs create mode 100644 tests/web-encoder.test.mjs create mode 160000 vendor/cuda-webshader create mode 100644 webgpu/encode.html create mode 100644 webgpu/encoder/app.js create mode 100644 webgpu/encoder/encoder.js create mode 100644 webgpu/encoder/format.js create mode 100644 webgpu/encoder/kernels.cu create mode 100644 webgpu/encoder/prepare.js create mode 100644 webgpu/encoder/style.css create mode 100644 webgpu/encoder/zip.js diff --git a/.gitignore b/.gitignore index f75c8cb..b0202cc 100644 --- a/.gitignore +++ b/.gitignore @@ -3,3 +3,4 @@ build_*/ out/ out_*/ __pycache__/ +node_modules/ diff --git a/.gitmodules b/.gitmodules new file mode 100644 index 0000000..0188b49 --- /dev/null +++ b/.gitmodules @@ -0,0 +1,3 @@ +[submodule "vendor/cuda-webshader"] + path = vendor/cuda-webshader + url = https://github.com/SamG-Coder/cuda-webshader.git diff --git a/README.md b/README.md index ec7803a..5a4ff11 100644 --- a/README.md +++ b/README.md @@ -3,6 +3,15 @@ # NNTC: Non-Neural Texture Compression +> **CUDA WebShader browser encoder fork.** This fork adds browser encoding and +> direct handoff to the existing WebGPU viewer. Clone with `--recurse-submodules`, +> run `npm ci` then `npm start`, and open `http://127.0.0.1:5174`. +> See [the browser encoder guide](docs/BROWSER_ENCODER.md) for supported features, +> tests and differences from the native encoder. The browser implementation uses +> the same NNTC asset format with a different optimizer; it does not reproduce +> the native encoder's full option set or exact output. The original project +> documentation follows below. + NNTC is a generalized latent/modulation codec for PBR material texture sets, in the same broad family of ideas as PVRTC1, but with more latent dimensions and a fixed bilinear feature expansion followed by a fitted affine decoder. It builds directly off the author's [public texture work from 2012](https://web.archive.org/web/20201024153426/https://sites.google.com/site/richgel99/luma_chroma_texture_compression). It uses a bilinear/degree-2 polynomial reconstruction: a degree-2 polynomial whose only quadratic terms are products between the two latents, which are stored as standard BC4/BC5 or uncompressed mipmapped textures. The full-resolution latent texture(s) carry selector/detail information, while the quarter-resolution latent texture carries more slowly varying material/color state. Decoding consists of a fixed, PVRTC1-like bilinear modulation/feature expansion followed by a single fitted affine transform, implemented as a GEMV plus bias. The fitted decoder itself contains no nonlinearity at all; there are no fitted nonlinear functions, activation functions, hidden layers, or neural-network layers. The `nntc_encode` tool (for Windows or Linux) compresses a PBR material texture set of up to six 24-bpp textures into **two mipmapped latent planes and (typically) a few dozen diff --git a/docs/BROWSER_ENCODER.md b/docs/BROWSER_ENCODER.md new file mode 100644 index 0000000..f12ecbe --- /dev/null +++ b/docs/BROWSER_ENCODER.md @@ -0,0 +1,141 @@ +# Browser encoder using CUDA WebShader + +This fork adds a working browser encoder at `webgpu/encode.html`. It compiles +the CUDA device kernels in `webgpu/encoder/kernels.cu` with CUDA WebShader and +executes them on WebGPU. The existing WebGPU viewer can open the exported files +or receive them directly from the encoder's **Open interactive viewer** button. +No CUDA installation or NVIDIA GPU is required by this path. + +## Run + +```sh +git clone --recurse-submodules --branch feat/cuda-webshader-encoder https://github.com/SamG-Coder/nntc.git +cd nntc +npm ci +npm start +``` + +Open http://127.0.0.1:5174. For an existing checkout, first run +`git submodule update --init --recursive`. Node 20+ is needed for the local +server and tests. The actual application is static JavaScript and needs only +HTTPS or localhost and a working WebGPU browser. A static deployment must serve +both `webgpu/` and `vendor/cuda-webshader/src/` with JavaScript MIME types; do not +serve only the `webgpu/` directory with the old viewer-only server. + +Choose one to six same-sized images, or try the bundled material. The browser +reads RGB bytes, ignores alpha with a notice, and pads the right/bottom edges to +multiples of four with clamped source pixels. Choose the latent channel counts, +number of fitting passes, mipmaps, and BC or uncompressed storage. Download the +ZIP and extract the descriptor and DDS files together. The output descriptor is +the existing `nntc-dds-1` format, with an informational `encoder` field. + +The app does not upload source images. Cancellation takes effect between GPU +dispatches and host stages. Images beyond the adapter's buffer/dispatch limits +or the encoder's estimated 1 GiB working-memory ceiling are refused before GPU +allocation. The UI also bounds source image memory before reading pixels. + +## Implemented path + +* One to six RGB maps; C0 and C1 each support one to four latent channels. +* Shared PCA initialization and an area-averaged source mip chain. +* The same ordered bilinear feature vector `[c; s; s*c]` and affine decoder. +* Center-only fitting or the center plus NNTC's four rotated-grid fractional + positions. Both latents and the source use bilinear clamp sampling. +* Streaming GPU normal equations; the at-most-25-variable dense solve runs in + JavaScript's double precision. The GPU never holds all observations at once. +* Four-color coordinate descent on each latent plane. Each texel gathers its + affected filter sites and solves its channel vector's small quadratic. +* BC4/BC5 endpoint/selector fitting (split into two DDS files for C0=3 or 4), or + R8/RG8/RGBA8 storage, plus quantized R8/RG8/RGBA8 for the quarter-size plane. +* A final decoder fit using the actual quantized/BC-decoded values that are + exported. Ranges and decoder coefficients are shared by all stored mips. +* Existing viewer validation of every exported descriptor and DDS file, source + comparison, base-level PSNR, ZIP download and interactive viewer handoff. + +The GPU performs observation generation, normal-equation accumulation, latent +updates and reconstruction. JavaScript performs image intake, source mip/PCA +initialization, small dense decoder solves, quantization/BC packing and file I/O. +The `cuda-webshader` submodule is pinned to the commit that supplies the new +`NormalEquations` streaming operation. + +## Differences from the native encoder + +This is a **format-compatible browser implementation, not native encoder parity**. +The native C++/CUDA encoder remains available and unchanged. The browser path +does not implement its full command-line/material-JSON option set or reproduce +its exact initialization, global sparse latent solves, ridge acceptance ladder, +palette modes, per-map weighting, sRGB/normal-specific mip filters, or joint +decoder-objective BC block refinement. BC packing here minimizes latent-channel +error, then refits the decoder. Raising fitting passes does not add those native +algorithms. + +GPU normal-equation products and sums use f32, with a fixed reduction order; +only the small host solve uses f64. Results are not promised to match native +CUDA bytes, PSNR or performance, or to be bit-identical across GPU vendors. +The four-color solver uses a bounded, ridge-regularized update, rather than the +native global sparse solve. Quality must be measured for the intended materials. + +The source mip filter is a box/area filter on the supplied RGB bytes; it does not +renormalize normals or linearize sRGB. The reported PSNR compares unclamped +reconstruction with original RGB at base-level pixel centers, excluding padding. +It uses the BC specification's palette values; actual hardware BC interpolation +and the viewer's byte-unpack fallback can vary slightly. It is not a fractional +site or whole-mip-chain error measurement. + +## Verification + +```sh +npm test +npm run test:gpu +npm run test:ui +``` + +The Node tests compile every CUDA entry and check all 32 combinations of latent +layout and storage type through the existing viewer parser, partial BC blocks, +DDS mip sizes, malformed input rejection and ZIP CRCs. GPU tests exercise +one/six outputs, split BC planes, uncompressed RGB, odd sizes, mip chains, +constant input, repeatability, cancellation and buffer cleanup. An independent +JavaScript decoder reads the generated DDS bytes and descriptor, and checks its +reconstruction against the GPU. The UI test encodes the bundled single-map and +four-map materials, downloads ZIPs, loads the interactive viewer, and checks a +mobile viewport and browser errors. + +On Windows the browser tests use installed Edge. Set `NNTC_BROWSER` to another +Playwright browser channel. On other platforms they use Playwright Chromium; +install it with `npx playwright install chromium`. `NNTC_SOFTWARE_GPU=1` requests +SwiftShader for the small numerical GPU suite. Reports, downloaded assets and +screenshots go to ignored `out/web-tests/`. + +Measured on the development machine with Edge/WebGPU: the bundled 512×512 +`m1.png`, default C0=2/C1=4, six fitting passes, fractional sites and mipmaps, +encoded in approximately 2.8 seconds at 35.2 dB, producing approximately +431.2 KiB including its JSON. This is a local smoke measurement, not a native +CUDA comparison or a claim for other hardware. + +The four-map `m1.png` through `m4.png` material took approximately 3.5 seconds +with the same options, producing 436.0 KiB and per-map PSNR of 25.9, 33.9, 30.7 +and 34.9 dB. The optional SwiftShader test could not obtain a WebGPU adapter in +the installed Edge browser; it is not counted as a passing software-adapter test. + +## JavaScript API + +```js +import {BrowserEncoder} from './webgpu/encoder/encoder.js'; + +const encoder = await BrowserEncoder.create(); +try { + const result = await encoder.encode([ + { name: 'albedo.png', width, height, data: rgbaUint8Array } + ], { name: 'material', c0: 2, c1: 4, iterations: 6, + mips: true, bc: true, sites: 5, + signal: abortController.signal, onProgress: console.log }); + // result.files: Map, containing DDS and JSON files. + // result.psnr: per-source-map PSNR; result.reconstructed: interleaved floats. +} finally { + await encoder.dispose(); +} +``` + +An existing CUDA WebShader runtime can be passed to `BrowserEncoder.create({runtime})`; +in that case the encoder does not own or dispose that runtime. One encoder accepts +one encode at a time; cancellation or failure releases all per-encode GPU buffers. diff --git a/package-lock.json b/package-lock.json new file mode 100644 index 0000000..33af057 --- /dev/null +++ b/package-lock.json @@ -0,0 +1,63 @@ +{ + "name": "nntc-cuda-webshader", + "lockfileVersion": 3, + "requires": true, + "packages": { + "": { + "name": "nntc-cuda-webshader", + "devDependencies": { + "playwright": "1.56.1" + }, + "engines": { + "node": ">=20" + } + }, + "node_modules/fsevents": { + "version": "2.3.2", + "resolved": "https://registry.npmjs.org/fsevents/-/fsevents-2.3.2.tgz", + "integrity": "sha512-xiqMQR4xAeHTuB9uWm+fFRcIOgKBMiOBP+eXiyT7jsgVCq1bkVygt00oASowB7EdtpOHaaPgKt812P9ab+DDKA==", + "dev": true, + "hasInstallScript": true, + "license": "MIT", + "optional": true, + "os": [ + "darwin" + ], + "engines": { + "node": "^8.16.0 || ^10.6.0 || >=11.0.0" + } + }, + "node_modules/playwright": { + "version": "1.56.1", + "resolved": "https://registry.npmjs.org/playwright/-/playwright-1.56.1.tgz", + "integrity": "sha512-aFi5B0WovBHTEvpM3DzXTUaeN6eN0qWnTkKx4NQaH4Wvcmc153PdaY2UBdSYKaGYw+UyWXSVyxDUg5DoPEttjw==", + "dev": true, + "license": "Apache-2.0", + "dependencies": { + "playwright-core": "1.56.1" + }, + "bin": { + "playwright": "cli.js" + }, + "engines": { + "node": ">=18" + }, + "optionalDependencies": { + "fsevents": "2.3.2" + } + }, + "node_modules/playwright-core": { + "version": "1.56.1", + "resolved": "https://registry.npmjs.org/playwright-core/-/playwright-core-1.56.1.tgz", + "integrity": "sha512-hutraynyn31F+Bifme+Ps9Vq59hKuUCz7H1kDOcBs+2oGguKkWTU50bBWrtz34OUWmIwpBTWDxaRPXrIXkgvmQ==", + "dev": true, + "license": "Apache-2.0", + "bin": { + "playwright-core": "cli.js" + }, + "engines": { + "node": ">=18" + } + } + } +} diff --git a/package.json b/package.json new file mode 100644 index 0000000..b59552d --- /dev/null +++ b/package.json @@ -0,0 +1,13 @@ +{ + "name": "nntc-cuda-webshader", + "private": true, + "type": "module", + "scripts": { + "start": "node scripts/serve-web.mjs", + "test": "node --test tests/web-encoder.test.mjs", + "test:gpu": "node tests/web-encoder-gpu.mjs", + "test:ui": "node tests/web-encoder-ui.mjs" + }, + "devDependencies": {"playwright": "1.56.1"}, + "engines": {"node": ">=20"} +} diff --git a/scripts/serve-web.mjs b/scripts/serve-web.mjs new file mode 100644 index 0000000..a60a45e --- /dev/null +++ b/scripts/serve-web.mjs @@ -0,0 +1,23 @@ +import http from 'node:http'; +import {readFile,stat} from 'node:fs/promises'; +import path from 'node:path'; +import {fileURLToPath} from 'node:url'; +const root=fileURLToPath(new URL('../',import.meta.url)); +const mime={'.html':'text/html','.js':'text/javascript','.mjs':'text/javascript','.json':'application/json','.css':'text/css','.cu':'text/plain','.wgsl':'text/plain','.png':'image/png','.jpg':'image/jpeg','.dds':'application/octet-stream'}; +export function createServer() { + return http.createServer(async(req,res)=>{ + try { + const pathname=decodeURIComponent(new URL(req.url,'http://localhost').pathname); + if(pathname==='/') {res.writeHead(302,{Location:'/webgpu/encode.html'}).end();return;} + const relative=pathname==='/'?'webgpu/encode.html':pathname.replace(/^\/+/,''),file=path.resolve(root,relative); + if(!file.startsWith(root)||relative.split(/[\\/]/).some(p=>p.startsWith('.'))) {res.writeHead(403).end();return;} + const target=(await stat(file)).isDirectory()?path.join(file,'index.html'):file; + const bytes=await readFile(target); + res.writeHead(200,{'Content-Type':mime[path.extname(target)]||'application/octet-stream','Cache-Control':'no-store'}).end(bytes); + } catch {res.writeHead(404).end('Not found');} + }); +} +if(process.argv[1]&&path.resolve(process.argv[1])===fileURLToPath(import.meta.url)) { + const port=Number(process.env.PORT||5174); + createServer().listen(port,'127.0.0.1',()=>console.log(`NNTC encoder: http://127.0.0.1:${port}/`)); +} diff --git a/tests/web-encoder-gpu.html b/tests/web-encoder-gpu.html new file mode 100644 index 0000000..b9d8466 --- /dev/null +++ b/tests/web-encoder-gpu.html @@ -0,0 +1,59 @@ +NNTC GPU tests + diff --git a/tests/web-encoder-gpu.mjs b/tests/web-encoder-gpu.mjs new file mode 100644 index 0000000..fecc318 --- /dev/null +++ b/tests/web-encoder-gpu.mjs @@ -0,0 +1,24 @@ +import {chromium} from 'playwright'; +import {createServer} from '../scripts/serve-web.mjs'; +import {mkdir,writeFile} from 'node:fs/promises'; +const server=createServer();await new Promise(resolve=>server.listen(0,'127.0.0.1',resolve)); +let browser; +try { + browser=await chromium.launch({...(process.env.NNTC_BROWSER?{channel:process.env.NNTC_BROWSER}:process.platform==='win32'?{channel:'msedge'}:{}),headless:true,args:['--enable-unsafe-webgpu',...(process.env.NNTC_SOFTWARE_GPU==='1'?['--use-angle=swiftshader','--enable-unsafe-swiftshader']:[])]}); + const page=await browser.newPage(); + page.on('console',message=>{if(message.type()==='error') console.error(message.text());}); + page.on('pageerror',error=>console.error(error)); + await page.goto(`http://127.0.0.1:${server.address().port}/tests/web-encoder-gpu.html`); + await page.waitForFunction(()=>window.runTest); + const results=[]; + for(const options of [{repeat:true},{c0:4,c1:3,width:16,height:16},{c0:1,c1:1,width:9,height:7,bc:false}, + {textures:6,c0:4,c1:4,width:20,height:12},{width:1,height:1,constant:true},{width:36,height:20,c0:3,c1:2}, + {sites:1,mips:false,bc:false,c0:3,c1:3}]) { + const result=await page.evaluate(options=>window.runTest(options),options);results.push({options,...result}); + console.log(JSON.stringify({options,psnr:result.psnr,maxDecodeDifference:result.difference,files:result.files.length})); + if(!result.psnr.every(v=>v>20)) throw new Error('Ramp quality below 20 dB.'); + } + await mkdir('out/web-tests',{recursive:true}); + await writeFile('out/web-tests/gpu.json',JSON.stringify(results,null,2)); +} finally {await browser?.close();await new Promise(resolve=>server.close(resolve));} + diff --git a/tests/web-encoder-ui.mjs b/tests/web-encoder-ui.mjs new file mode 100644 index 0000000..478d662 --- /dev/null +++ b/tests/web-encoder-ui.mjs @@ -0,0 +1,41 @@ +import {chromium} from 'playwright'; +import {createServer} from '../scripts/serve-web.mjs'; +import {mkdir,writeFile} from 'node:fs/promises'; +import assert from 'node:assert/strict'; +const server=createServer();await new Promise(resolve=>server.listen(0,'127.0.0.1',resolve)); +let browser; +try { + browser=await chromium.launch({...(process.env.NNTC_BROWSER?{channel:process.env.NNTC_BROWSER}:process.platform==='win32'?{channel:'msedge'}:{}),headless:true,args:['--enable-unsafe-webgpu',...(process.env.NNTC_SOFTWARE_GPU==='1'?['--use-angle=swiftshader','--enable-unsafe-swiftshader']:[])]}); + const page=await browser.newPage({viewport:{width:1440,height:1100}}),errors=[]; + page.on('pageerror',error=>errors.push(error.message)); + page.on('console',message=>{if(message.type()==='error')errors.push(message.text());}); + await page.goto(`http://127.0.0.1:${server.address().port}/`); + await page.locator('#example').click();await page.waitForFunction(()=>!document.querySelector('#encode').disabled); + await page.locator('#encode').click(); + await page.waitForFunction(()=>document.querySelector('#status').textContent.startsWith('Encoding complete')||document.body.classList.contains('error'),null,{timeout:300000}); + const status=await page.locator('#status').textContent();assert.match(status,/Encoding complete/); + await mkdir('out/web-tests',{recursive:true}); + await page.screenshot({path:'out/web-tests/encoder.png',fullPage:true}); + const downloadPromise=page.waitForEvent('download');await page.locator('#download').click();const download=await downloadPromise;await download.saveAs('out/web-tests/m1.zip'); + const metrics=await page.locator('#metrics').innerText();console.log(metrics); + await page.locator('#view').click(); + const frame=page.frameLocator('#viewer');await frame.locator('#status-line').filter({hasText:'m1_nntc.json'}).waitFor({timeout:60000}); + const viewerStatus=await frame.locator('#status-line').innerText();console.log(viewerStatus); + assert.match(viewerStatus,/Loaded/); + await page.screenshot({path:'out/web-tests/viewer.png',fullPage:true}); + await page.setViewportSize({width:390,height:844});await page.locator('#close-viewer').click();await page.screenshot({path:'out/web-tests/mobile.png',fullPage:true}); + assert.equal(await page.evaluate(()=>document.documentElement.scrollWidth>innerWidth),false); + await page.setViewportSize({width:1440,height:1100}); + await page.locator('#images').setInputFiles(['examples/m1.png','examples/m2.png','examples/m3.png','examples/m4.png']); + await page.waitForFunction(()=>document.querySelector('#status').textContent.startsWith('4 maps')); + await page.locator('#encode').click(); + await page.waitForFunction(()=>document.querySelector('#status').textContent.startsWith('Encoding complete')||document.body.classList.contains('error'),null,{timeout:300000}); + assert.match(await page.locator('#status').innerText(),/Encoding complete/); + const material=[]; + for(let i=0;i<4;i++) {await page.locator('#texture').selectOption(String(i));material.push(await page.locator('#metrics').innerText());} + console.log(JSON.stringify({fourMapMaterial:material})); + const multiDownloadPromise=page.waitForEvent('download');await page.locator('#download').click();await (await multiDownloadPromise).saveAs('out/web-tests/m1-m4.zip'); + assert.deepEqual(errors,[]); + await writeFile('out/web-tests/ui.json',JSON.stringify({status,metrics,viewerStatus,fourMapMaterial:material,errors},null,2)); +}finally{await browser?.close();await new Promise(resolve=>server.close(resolve));} + diff --git a/tests/web-encoder.test.mjs b/tests/web-encoder.test.mjs new file mode 100644 index 0000000..250c8b9 --- /dev/null +++ b/tests/web-encoder.test.mjs @@ -0,0 +1,66 @@ +import test from 'node:test'; +import assert from 'node:assert/strict'; +import {readFile} from 'node:fs/promises'; +import {compile} from '../vendor/cuda-webshader/src/compiler/compiler.js'; +import {packBC,quantize,writeDDS,makeDescriptor,ranges} from '../webgpu/encoder/format.js'; +import {prepare,validateInputs} from '../webgpu/encoder/prepare.js'; +import {ddsRead} from '../webgpu/dds.js'; +import {unpackImage} from '../webgpu/bcdec.js'; +import {buildAsset} from '../webgpu/descriptor.js'; +import {zipFiles} from '../webgpu/encoder/zip.js'; + +test('all browser CUDA entries compile through cuda-webshader',async()=>{ + const source=await readFile(new URL('../webgpu/encoder/kernels.cu',import.meta.url),'utf8'); + for(const entry of ['make_rows','update_latent','decode']) assert.match(compile(source,{entry,workgroupSize:[64,1,1]}).wgsl,/@compute/); +}); +test('BC packing round-trips partial blocks, constant channels and split BC5/BC4 planes',()=>{ + for(let channels=1;channels<=4;channels++) { + const w=7,h=5,values=Float32Array.from({length:w*h*channels},(_,i)=>i%channels===0?0.25:((i*17)%256)/255); + const {payloads,decoded}=packBC(values,w,h,channels,Array(channels).fill(0),Array(channels).fill(1)); + let firstChannel=0; + for(const payload of payloads) { + const count=Math.min(2,channels-firstChannel),bytes=writeDDS(w,h,count===1?80:83,count,[payload]),dds=ddsRead('test.dds',bytes),unpacked=unpackImage(dds)[0]; + assert.equal(new DataView(bytes.buffer).getUint32(24,true),1,'upstream Python decoder requires DDS depth=1'); + for(let i=0;i{ + for(let c0=1;c0<=4;c0++)for(let c1=1;c1<=4;c1++)for(const bc of [false,true]) { + const planes=[{w:16,h:16,w1:4,h1:4},{w:8,h:8,w1:2,h1:2}],r0={lo:Array(c0).fill(-1),hi:Array(c0).fill(1)},r1={lo:Array(c1).fill(-1),hi:Array(c1).fill(1)}; + const d=makeDescriptor({name:'test',width:16,height:16,inputs:['a.png','b.png'],c0,c1,planes,range0:r0,range1:r1,weights:new Float32Array((c0+c1+c0*c1+1)*6),bc}),files=new Map(); + for(let level=0;level<2;level++)for(const e of d.textures[level].files||[d.textures[level]]) { + const compressed=e.dxgi_format_id===80||e.dxgi_format_id===83,texture=d.textures[level]; + files.set(e.file,writeDDS(texture.width,texture.height,e.dxgi_format_id,e.channels_stored,texture.mip_sizes.map(([w,h])=>new Uint8Array(compressed?Math.ceil(w/4)*Math.ceil(h/4)*e.channels_stored*8:w*h*e.channels_stored)))); + } + assert.equal(buildAsset('test_nntc.json',d,files,true).texturesOut,2); + } +}); +test('preparation clamps padding and generates matching DDS mip dimensions',()=>{ + const width=35,height=19,data=new Uint8Array(width*height*4).fill(127); + const result=prepare([{width,height,data}],{c0:3,c1:4}); + assert.equal(result.planes[0].w,36);assert.equal(result.planes[0].h,20); + for(const [i,p] of result.planes.entries()) { + assert.ok(p.values0.every(Number.isFinite));assert.ok(p.values1.every(Number.isFinite)); + assert.equal(p.w1,Math.max(1,result.planes[0].w1>>i));assert.equal(p.h1,Math.max(1,result.planes[0].h1>>i)); + assert.ok(p.source.every(v=>Math.abs(v-127/255)<1e-6)); + } +}); +test('reject malformed inputs, nonfinite latents and incorrect DDS payloads',()=>{ + assert.throws(()=>validateInputs([]),/one to six/); + const image={width:1,height:1,data:new Uint8Array(4)}; + assert.throws(()=>validateInputs([image],{c0:5}),/1..4/); + assert.throws(()=>validateInputs([image],{iterations:-1}),/Iterations/); + assert.throws(()=>validateInputs([image,{...image,width:2}]),/same dimensions/); + assert.throws(()=>ranges([new Float32Array([NaN])],1),/non-finite/); + assert.throws(()=>writeDDS(4,4,80,1,[new Uint8Array(7)]),/payload/); + const r=ranges([new Float32Array([0,0])],1);assert.ok(r.hi[0]>r.lo[0]); + assert.ok(quantize(new Float32Array([0,0]),1,r.lo,r.hi).decoded.every(Number.isFinite)); +}); +test('ZIP contains standard headers, correct CRC32 and intact bytes',async()=>{ + const data=new TextEncoder().encode('123456789'),zip=new Uint8Array(await zipFiles(new Map([['test.txt',data]])).arrayBuffer()),v=new DataView(zip.buffer); + assert.equal(v.getUint32(0,true),0x04034b50);assert.equal(v.getUint32(14,true),0xcbf43926); + assert.deepEqual(zip.slice(38,47),data);assert.equal(v.getUint32(47,true),0x02014b50); + assert.equal(v.getUint32(zip.length-22,true),0x06054b50);assert.equal(v.getUint16(zip.length-12,true),1); +}); diff --git a/vendor/cuda-webshader b/vendor/cuda-webshader new file mode 160000 index 0000000..896432c --- /dev/null +++ b/vendor/cuda-webshader @@ -0,0 +1 @@ +Subproject commit 896432c3d7a78faaf72b43ae6bcfd442b31e8428 diff --git a/webgpu/README.md b/webgpu/README.md index 9c51e66..1850fd8 100644 --- a/webgpu/README.md +++ b/webgpu/README.md @@ -1,5 +1,11 @@ # The NNTC viewer in the browser (WebGPU) +> This fork also includes a browser encoder at `encode.html`, powered by CUDA +> WebShader. Start it from the repository root with `npm ci` and `npm start`; +> initialize the pinned `vendor/cuda-webshader` submodule first. See +> [the encoder guide](../docs/BROWSER_ENCODER.md). The instructions below apply +> to the viewer by itself. + This page draws an NNTC asset the way a game would: two latent textures on a real hardware sampler, and one small affine layer over two samples. It is the same decode the Direct3D 11, Direct3D 12 and Vulkan viewers run, ported to WGSL. The controls are ordinary HTML on the left, the picture is WebGPU on the right, and a two-line strip is drawn diff --git a/webgpu/encode.html b/webgpu/encode.html new file mode 100644 index 0000000..7ea48cf --- /dev/null +++ b/webgpu/encode.html @@ -0,0 +1,41 @@ + + + +NNTC · Browser encoder + +
NNTC TEXTURE TOOLSOpen asset viewer ↗
+
+
CUDA WEBSHADER / WEBGPU

One material.
Shared texture space.

+

Compress up to six material maps into two latent planes and a small decoder. Encode, inspect, and export directly in your browser.

+
+
+
+
01

Source material

+ + +
+

Images stay on your device. RGB channels are encoded; alpha is ignored.

+
02

Compression

+ +
+
+
+
+ + +

This browser optimizer uses the NNTC format, with a different fitting strategy from the native encoder. Quality and speed can differ.

+
+
+

Choose images to begin.

+
+
+
03

Inspect the result

+

Your material, reconstructed.

Compare the source with the exported representation here.

+ +
+
+ +
Non-Neural Texture Compression · Based on Richard Geldreich Jr.’s NNTC · Source
+
+ diff --git a/webgpu/encoder/app.js b/webgpu/encoder/app.js new file mode 100644 index 0000000..baf2d37 --- /dev/null +++ b/webgpu/encoder/app.js @@ -0,0 +1,76 @@ +// SPDX-License-Identifier: Apache-2.0 +import {BrowserEncoder} from './encoder.js'; +import {zipFiles} from './zip.js'; +const $=id=>document.getElementById(id); +let images=[],result=null,controller=null,encoder=null,elapsed=0,previewURLs=[]; +const status=(message,error=false)=>{$('status').textContent=message;document.body.classList.toggle('error',error);}; +async function loadImages(files) { + if(controller) return; + if(files.length<1||files.length>6) throw new Error('Choose one to six images.'); + const next=[]; + for(const file of files) { + const bitmap=await createImageBitmap(file,{colorSpaceConversion:'none',premultiplyAlpha:'none'}); + try { + if(bitmap.width>16384||bitmap.height>16384||bitmap.width*bitmap.height*files.length>32*1024*1024) throw new Error('This material is too large for the browser encoder; reduce the image dimensions.'); + const canvas=new OffscreenCanvas(bitmap.width,bitmap.height),ctx=canvas.getContext('2d',{willReadFrequently:true}); + ctx.drawImage(bitmap,0,0);const data=ctx.getImageData(0,0,canvas.width,canvas.height).data; + next.push({name:file.name,width:canvas.width,height:canvas.height,data,file}); + } finally {bitmap.close();} + } + if(next.some(image=>image.width!==next[0].width||image.height!==next[0].height)) throw new Error('All material images must have the same dimensions.'); + images=next;for(const url of previewURLs) URL.revokeObjectURL(url);previewURLs=[];$('inputs').replaceChildren(); + for(const image of images) {const thumbnail=document.createElement('img');thumbnail.src=URL.createObjectURL(image.file);previewURLs.push(thumbnail.src);thumbnail.alt=image.name;thumbnail.title=image.name;$('inputs').append(thumbnail);} + $('encode').disabled=false; + $('name').value=images[0].name.replace(/\.[^.]+$/,'').replace(/[^\w.-]/g,'_')||'material'; + result=null;$('result').hidden=true;$('empty').hidden=false;$('texture').hidden=true;$('viewer-section').hidden=true;$('viewer').removeAttribute('src'); + const alpha=images.some(image=>image.data.some((value,i)=>i%4===3&&value!==255)); + status(`${images.length} map${images.length===1?'':'s'} · ${images[0].width} × ${images[0].height}. ${alpha?'Alpha will be ignored.':'Ready to encode.'}`); +} +$('images').addEventListener('change',()=>loadImages([...$('images').files]).catch(error=>status(error.message,true))); +for(const kind of ['dragover','dragleave','drop']) $('dropzone').addEventListener(kind,event=>{ + event.preventDefault();$('dropzone').classList.toggle('drag',kind==='dragover'); + if(kind==='drop') loadImages([...event.dataTransfer.files]).catch(error=>status(error.message,true)); +}); +$('example').addEventListener('click',async()=>{ + try {status('Loading the bundled 512 × 512 material…');const response=await fetch('../examples/m1.png');if(!response.ok) throw new Error('Could not load the example image.');await loadImages([new File([await response.blob()],'m1.png')]);} + catch(error){status(error.message,true);} +}); +function showComparison() { + const texture=Number($('texture').value),image=images[texture]; + const source=$('source');source.width=image.width;source.height=image.height;source.getContext('2d').putImageData(new ImageData(image.data,image.width,image.height),0,0); + const target=$('reconstruction');target.width=result.width;target.height=result.height; + const pixels=new Uint8ClampedArray(result.width*result.height*4); + for(let y=0;yn+data.length,0),metric=(value,label)=>{const block=document.createElement('div'),strong=document.createElement('strong'),span=document.createElement('span');strong.textContent=value;span.textContent=label;block.append(strong,span);return block;}; + $('metrics').replaceChildren(metric(`${result.psnr[texture].toFixed(1)} dB`,'Reconstruction PSNR'),metric(`${(bytes/1024).toFixed(1)} KB`,'DDS + descriptor'),metric(`${elapsed.toFixed(1)} s`,'Encoding time')); +} +$('texture').addEventListener('change',showComparison); +$('form').addEventListener('submit',async event=>{ + event.preventDefault();if(controller||!images.length) return; + controller=new AbortController();$('settings').disabled=true;$('encode').disabled=true;$('cancel').hidden=false;document.body.classList.add('busy'); + const start=performance.now(); + try { + status('Initializing CUDA WebShader and validating GPU kernels…'); + encoder ||= await BrowserEncoder.create(); + const next=await encoder.encode(images,{name:$('name').value,c0:Number($('c0').value),c1:Number($('c1').value),iterations:Number($('iterations').value), + mips:$('mips').checked,sites:$('filtered').checked?5:1,bc:$('storage').value==='bc',signal:controller.signal,onProgress:status}); + result=next;elapsed=(performance.now()-start)/1000; + $('texture').replaceChildren(...images.map((image,i)=>new Option(image.name,String(i))));$('texture').hidden=false; + $('empty').hidden=true;$('result').hidden=false;showComparison();status('Encoding complete. The exported descriptor and DDS files passed the viewer’s format checks.'); + $('viewer-section').hidden=true;$('viewer').removeAttribute('src'); + } catch(error) {status(error.message,error.name!=='AbortError');} + finally {controller=null;$('settings').disabled=false;$('encode').disabled=false;$('cancel').hidden=true;document.body.classList.remove('busy');} +}); +$('cancel').addEventListener('click',()=>{controller?.abort();status('Cancelling after the current GPU dispatch…');}); +$('download').addEventListener('click',()=>{ + if(!result)return; + const url=URL.createObjectURL(zipFiles(result.files)),link=document.createElement('a');link.href=url;link.download=result.descriptorName.replace('_nntc.json','.zip');link.click();setTimeout(()=>URL.revokeObjectURL(url),30000); +}); +$('view').addEventListener('click',()=>{$('viewer-section').hidden=false;$('viewer').src='index.html?encoder=1';$('viewer-section').scrollIntoView({behavior:'smooth'});}); +$('close-viewer').addEventListener('click',()=>{$('viewer-section').hidden=true;$('viewer').removeAttribute('src');}); +window.addEventListener('message',event=>{ + if(event.origin!==location.origin||event.source!==$('viewer').contentWindow||event.data?.type!=='nntc-viewer-ready'||!result) return; + $('viewer').contentWindow.postMessage({type:'nntc-encoded-files',files:[...result.files].map(([name,bytes])=>new File([bytes],name))},location.origin); +}); +if(!navigator.gpu) {status('WebGPU is unavailable. Open localhost or HTTPS in a WebGPU-enabled browser.',true);$('settings').disabled=true;} diff --git a/webgpu/encoder/encoder.js b/webgpu/encoder/encoder.js new file mode 100644 index 0000000..fdfc13e --- /dev/null +++ b/webgpu/encoder/encoder.js @@ -0,0 +1,119 @@ +// SPDX-License-Identifier: Apache-2.0 +import {GpuRuntime} from '../../vendor/cuda-webshader/src/runtime/runtime.js'; +import {NormalEquations} from '../../vendor/cuda-webshader/src/runtime/least-squares.js'; +import {prepare,validateInputs} from './prepare.js'; +import {ranges,quantize,packBC,makeDescriptor,writeDDS} from './format.js'; +import {buildAsset} from '../descriptor.js'; + +const TILE=4096; +const tick=()=>new Promise(resolve=>setTimeout(resolve,0)); +const abort=signal=>{if(signal?.aborted) throw new DOMException('Encoding cancelled.','AbortError');}; + +export class BrowserEncoder { + static async create({runtime}={}) { + const own=!runtime; + runtime ||= await GpuRuntime.create({useAdapterBufferLimits:true}); + try { + const response=await fetch(new URL('./kernels.cu',import.meta.url)); + if(!response.ok) throw new Error(`Cannot load encoder kernels: HTTP ${response.status}`); + const source=await response.text(),kernels={}; + for(const entry of ['make_rows','update_latent','decode']) kernels[entry]=await runtime.kernel(source,{entry,workgroupSize:[64,1,1]}); + return new BrowserEncoder(runtime,kernels,own); + } catch(error) {if(own) runtime.dispose();throw error;} + } + constructor(runtime,kernels,own) {this.runtime=runtime;this.kernels=kernels;this.own=own;this.busy=false;} + async dispose() { + if(this.busy) throw new Error('Wait for the encode or cancellation before disposing.'); + await this.runtime.idle();if(this.own) this.runtime.dispose(); + } + async encode(images,{name='material',c0=2,c1=4,iterations=6,sites=5,mips=true,bc=true,signal,onProgress=()=>{}}={}) { + if(this.busy) throw new Error('An encode is already running.'); + const {width,height}=validateInputs(images,{c0,c1,iterations,sites}); + if(typeof name!=='string'||!name.length||!/^[\w.-]+$/.test(name)||name==='.'||name==='..') throw new Error('Output name must contain letters, digits, dots, underscores or hyphens.'); + const rt=this.runtime,limits=rt.device.limits,w=Math.ceil(width/4)*4,h=Math.ceil(height/4)*4,nout=images.length*3; + const largest=w*h*Math.max(nout,c0)*4; + if(largest>Math.min(limits.maxStorageBufferBindingSize,limits.maxBufferSize)||Math.ceil(w*h/64)>limits.maxComputeWorkgroupsPerDimension) throw new Error('Images exceed this adapter’s buffer or dispatch limits; reduce their dimensions.'); + // Includes source chain, both latent ping-pong buffers, host preparation, + // readback and reconstruction; refuse before allocating those arrays. + if(largest*5+w*h*(c0+c1)*16>1024**3) throw new Error('This browser encoder limits working memory to 1 GiB; reduce image dimensions.'); + this.busy=true; + const resources=new Set(),alloc=data=>{const b=rt.createBuffer(data);resources.add(b);return b;}; + const report=async message=>{abort(signal);onProgress(message);await tick();abort(signal);}; + const launch=(entry,buffers,scalars,count)=>rt.batch().dispatch(this.kernels[entry].bind(buffers,scalars),[Math.ceil(count/64),1,1]).submit(); + try { + await report('Preparing source mip chain and shared PCA basis…'); + const model=prepare(images,{c0,c1,iterations,sites,mips}),planes=model.planes,mm=c0+c1+c0*c1+1; + const decoder=alloc(mm*nout*4),rows=alloc(TILE*(mm+nout)*4); + for(const p of planes) { + p.src=alloc(p.source);p.l0=alloc(p.values0);p.l1=alloc(p.values1); + p.scratch0=alloc(p.values0.byteLength);p.scratch1=alloc(p.values1.byteLength); + p.scalars={w:p.w,h:p.h,w1:p.w1,h1:p.h1,c0,c1,nout}; + } + let weights; + const fit=async()=>{ + const equations=await NormalEquations.create(rt,mm,nout); + try { + for(const p of planes) for(let start=0;startp.values0),c0),range1=ranges(planes.map(p=>p.values1),c1); + for(const p of planes) { + abort(signal); + if(bc) {const packed=packBC(p.values0,p.w,p.h,c0,range0.lo,range0.hi);p.payload0=packed.payloads;p.values0=packed.decoded;} + else {const packed=quantize(p.values0,c0,range0.lo,range0.hi);p.payload0=[packed.bytes];p.values0=packed.decoded;} + const packed=quantize(p.values1,c1,range1.lo,range1.hi);p.payload1=packed.bytes;p.values1=packed.decoded; + rt.write(p.l0,p.values0);rt.write(p.l1,p.values1);await tick(); + } + await report('Refitting the decoder against the stored, quantized representation…');await fit(); + const descriptor=makeDescriptor({name,width,height,inputs:images.map((image,i)=>image.name||`input${i}.png`),c0,c1,planes,range0,range1,weights,bc}); + const files=new Map(); + for(let level=0;level<2;level++) { + const texture=descriptor.textures[level]; + (texture.files||[texture]).forEach((entry,index)=>files.set(entry.file,writeDDS(texture.width,texture.height,entry.dxgi_format_id,entry.channels_stored, + planes.map(p=>level?p.payload1:p.payload0[index])))); + } + const descriptorName=`${name}_nntc.json`; + // Independent, existing viewer parser checks each DDS header and model. + buildAsset(descriptorName,descriptor,files,true); + files.set(descriptorName,new TextEncoder().encode(JSON.stringify(descriptor,null,2)+'\n')); + await report('Measuring the exported representation…'); + const base=planes[0],output=alloc(base.source.byteLength); + launch('decode',{latent0:base.l0,latent1:base.l1,decoder,output},base.scalars,base.w*base.h); + const reconstructed=await rt.read(output),psnr=[]; + if(reconstructed.some(v=>!Number.isFinite(v))) throw new Error('Non-finite reconstruction; the encode was not exported.'); + for(let t=0;t>1); h = Math.max(1,h>>1); + } + const bytes = new Uint8Array(148 + levels.reduce((n,data)=>n+data.byteLength,0)); + const header = new DataView(bytes.buffer); + const set = (i,v)=>header.setUint32(i*4,v,true); + set(0,0x20534444); set(1,124); set(2,0x1007 | (bc?0x80000:8) | (levels.length>1?0x20000:0)); + set(3,height); set(4,width); set(5,bc?levels[0].length:width*channels); set(6,1); set(7,levels.length); + set(19,32); set(20,4); set(21,0x30315844); + set(27,0x1000 | (levels.length>1?0x400008:0)); + set(32,dxgi); set(33,3); set(35,1); + let offset = 148; + for (const data of levels) { bytes.set(data,offset); offset += data.byteLength; } + return bytes; +} + +// Endpoint / selector least squares in each BC4 channel. The final decoder is +// refitted against these exact standard palette values, not the prepack plane. +export function packBC(values, w, h, channels, lo, hi) { + const bx = Math.ceil(w/4), by = Math.ceil(h/4); + const groups = channels > 2 ? [2,channels-2] : [channels]; + const payloads = groups.map(n=>new Uint8Array(bx*by*n*8)); + const decoded = new Float32Array(values.length); + for (let y = 0; y < by; y++) for (let x = 0; x < bx; x++) for (let ch = 0; ch < channels; ch++) { + const samples = new Float64Array(16), range = hi[ch]-lo[ch]; + for (let t = 0; t < 16; t++) { + const px = Math.min(w-1,x*4+(t&3)), py = Math.min(h-1,y*4+(t>>2)); + samples[t] = Math.max(0,Math.min(255,(values[(py*w+px)*channels+ch]-lo[ch])/range*255)); + } + let a = Math.round(Math.max(...samples)), b = Math.round(Math.min(...samples)); + let best = null; + for (let iteration = 0; iteration < 5; iteration++) { + if (a <= b) { if (b < 255) a = b+1; else b = a-1; } + const palette = [a,b,...Array.from({length:6},(_,i)=>((6-i)*a+(i+1)*b)/7)]; + const indices = new Uint8Array(16); + let error=0, aa=0, ab=0, bb=0, ay=0, byValue=0; + for (let t = 0; t < 16; t++) { + let pick=0; + for (let k=1;k<8;k++) if (Math.abs(samples[t]-palette[k]) < Math.abs(samples[t]-palette[pick])) pick=k; + indices[t]=pick; error+=(samples[t]-palette[pick])**2; + const u=pick===0?1:pick===1?0:(8-pick)/7, v=1-u; + aa+=u*u; ab+=u*v; bb+=v*v; ay+=u*samples[t]; byValue+=v*samples[t]; + } + if (!best || error < best.error) best={a,b,indices,palette,error}; + const determinant=aa*bb-ab*ab; + if (determinant<1e-9) break; + a=Math.max(0,Math.min(255,Math.round((ay*bb-byValue*ab)/determinant))); + b=Math.max(0,Math.min(255,Math.round((byValue*aa-ay*ab)/determinant))); + } + const group=ch<2?0:1, gc=groups[group], local=ch%2, bytes=payloads[group]; + const offset=((y*bx+x)*gc+local)*8; + bytes[offset]=best.a; bytes[offset+1]=best.b; + for(let t=0;t<16;t++) { + const bit=t*3, byte=offset+2+(bit>>3), shift=bit&7; + bytes[byte] |= best.indices[t]<5) bytes[byte+1] |= best.indices[t]>>(8-shift); + const px=x*4+(t&3), py=y*4+(t>>2); + if(px({level,width:planes[0][level?'w1':'w'],height:planes[0][level?'h1':'h'], + channels_used:channels,bits_per_channel:Array(channels).fill(8),mip_count:planes.length, + mip_sizes:planes.map(p=>[p[level?'w1':'w'],p[level?'h1':'h']]), + dequantise:{kind:'range',lo:range.lo,hi:range.hi,levels:255}}); + const first=texture(0,c0,range0),second=texture(1,c1,range1); + const entry=(file,channels,compressed)=>({file,channels_stored:channels,dxgi_format_id:compressed?(channels===1?80:83):channels===1?61:channels===2?49:28}); + if(bc && c0>2) first.files=[entry(`${name}_lat0a.dds`,2,true),entry(`${name}_lat0b.dds`,c0-2,true)]; + else Object.assign(first,entry(`${name}_lat0.dds`,bc?c0:c0===3?4:c0,bc)); + Object.assign(second,entry(`${name}_lat1.dds`,c1===3?4:c1,false)); + for(const t of [first,second]) for(const e of t.files||[t]) e.dxgi_format=dxgiName(e.dxgi_format_id); + if(bc) first.bc_palette='standard'; + const matrix=[],bias=[]; + for(let o=0;o({file,filter:'box',srgb:false,edge:'clamp',normal_map:false,weight:1,rgb_weights:[1,1,1]}))}, + decode:{width:planes[0].w,height:planes[0].h,textures_out:inputs.length},block:4,lod_bias_level1:2, + textures:[first,second],decoder:{type:'bilinear',terms:'a b sc',C0:c0,C1:c1,nin,nout,hidden:[],output:'identity', + layers:[{rows:nout,cols:nin,weights:matrix,bias}]}}; +} diff --git a/webgpu/encoder/kernels.cu b/webgpu/encoder/kernels.cu new file mode 100644 index 0000000..55ba9b9 --- /dev/null +++ b/webgpu/encoder/kernels.cu @@ -0,0 +1,148 @@ +// SPDX-License-Identifier: Apache-2.0 +// Browser port of NNTC's bilinear features and fixed filter-aware site objective. +// Derived from Richard Geldreich Jr.'s src/sample.cuh (Copyright 2026). +// Host orchestration lives in encoder.js; cuda-webshader compiles these kernels. +__device__ float sample(const float* image, int w, int h, int channels, int ch, float x, float y) { + int ix = (int)floorf(x), iy = (int)floorf(y); + float fx = x - (float)ix, fy = y - (float)iy; + int x0 = max(0, min(w - 1, ix)), x1 = max(0, min(w - 1, ix + 1)); + int y0 = max(0, min(h - 1, iy)), y1 = max(0, min(h - 1, iy + 1)); + float a = image[(y0 * w + x0) * channels + ch]; + float b = image[(y0 * w + x1) * channels + ch]; + float c = image[(y1 * w + x0) * channels + ch]; + float d = image[(y1 * w + x1) * channels + ch]; + return (a + (b - a) * fx) * (1.0f - fy) + (c + (d - c) * fx) * fy; +} +__device__ float site_x(int j) { + if (j == 1) return -0.375f; + if (j == 2) return -0.125f; + if (j == 3) return 0.125f; + if (j == 4) return 0.375f; + return 0.0f; +} +__device__ float site_y(int j) { + if (j == 1) return -0.125f; + if (j == 2) return 0.375f; + if (j == 3) return -0.375f; + if (j == 4) return 0.125f; + return 0.0f; +} +__device__ float influence(float coordinate, int texel, int extent) { + int low = (int)floorf(coordinate); + float f = coordinate - (float)low; + float weight = 0.0f; + if (max(0, min(extent - 1, low)) == texel) weight += 1.0f - f; + if (max(0, min(extent - 1, low + 1)) == texel) weight += f; + return weight; +} +__global__ void make_rows(const float* source, const float* latent0, const float* latent1, float* rows, + int w, int h, int w1, int h1, int c0, int c1, int nout, int sites, int start, int count, float weight) { + int id = blockIdx.x * blockDim.x + threadIdx.x; + if (id >= count) return; + int si = start + id, pixel = si / sites, j = si % sites; + float x = (float)(pixel % w) + site_x(j), y = (float)(pixel / w) + site_y(j); + float s[4], c[4]; + for (int i = 0; i < c0; i++) s[i] = sample(latent0,w,h,c0,i,x,y); + for (int i = 0; i < c1; i++) c[i] = sample(latent1,w1,h1,c1,i,(x+0.5f)*(float)w1/(float)w-0.5f,(y+0.5f)*(float)h1/(float)h-0.5f); + int mm = c0 + c1 + c0*c1 + 1, at = id*(mm+nout), k = 0; + for (int i = 0; i < c1; i++) rows[at+k++] = c[i]*weight; + for (int i = 0; i < c0; i++) rows[at+k++] = s[i]*weight; + for (int i = 0; i < c0; i++) for (int t = 0; t < c1; t++) rows[at+k++] = s[i]*c[t]*weight; + rows[at+k++] = weight; + for (int o = 0; o < nout; o++) rows[at+k++] = sample(source,w,h,nout,o,x,y)*weight; +} + +// Four-color block coordinate descent. With the opposite latent fixed, each +// texel's channel vector has a small convex quadratic. Gather all affected sites +// so neighboring work items never scatter/atomically accumulate floating values. +// Read and write distinct buffers: a dispatch sees a coherent previous plane. +__global__ void update_latent(const float* source, const float* latent0, const float* latent1, + const float* decoder, float* result, int w, int h, int w1, int h1, + int c0, int c1, int nout, int sites, int level, int color) { + int id = blockIdx.x * blockDim.x + threadIdx.x; + int tw = level == 0 ? w : w1, th = level == 0 ? h : h1, channels = level == 0 ? c0 : c1; + if (id >= tw*th) return; + int tx = id % tw, ty = id / tw; + for (int a = 0; a < channels; a++) result[id*channels+a] = level == 0 ? latent0[id*c0+a] : latent1[id*c1+a]; + if ((tx % 2) + 2*(ty % 2) != color) return; + int mm = c0 + c1 + c0*c1 + 1; + float matrix[16], rhs[4]; + for (int a = 0; a < 16; a++) matrix[a] = 0.0f; + for (int a = 0; a < 4; a++) rhs[a] = 0.0f; + float scaleX = (float)w/(float)tw, scaleY = (float)h/(float)th; + int left = max(0,(int)floorf(((float)tx-0.5f)*scaleX-1.0f)); + int right = min(w-1,(int)ceilf(((float)tx+1.5f)*scaleX+1.0f)); + int top = max(0,(int)floorf(((float)ty-0.5f)*scaleY-1.0f)); + int bottom = min(h-1,(int)ceilf(((float)ty+1.5f)*scaleY+1.0f)); + for (int py = top; py <= bottom; py++) for (int px = left; px <= right; px++) for (int j = 0; j < sites; j++) { + float x = (float)px+site_x(j), y = (float)py+site_y(j); + float gx = (x+0.5f)/scaleX-0.5f, gy = (y+0.5f)/scaleY-0.5f; + float q = influence(gx,tx,tw)*influence(gy,ty,th); + if (q <= 0.0f) continue; + float s[4], c[4], phi[25]; + for (int a = 0; a < c0; a++) s[a] = sample(latent0,w,h,c0,a,x,y); + for (int a = 0; a < c1; a++) c[a] = sample(latent1,w1,h1,c1,a,(x+0.5f)*(float)w1/(float)w-0.5f,(y+0.5f)*(float)h1/(float)h-0.5f); + int k = 0; + for (int a = 0; a < c1; a++) phi[k++] = c[a]; + for (int a = 0; a < c0; a++) phi[k++] = s[a]; + for (int a = 0; a < c0; a++) for (int b = 0; b < c1; b++) phi[k++] = s[a]*c[b]; + phi[k] = 1.0f; + for (int o = 0; o < nout; o++) { + float predicted = 0.0f; + for (int a = 0; a < mm; a++) predicted += decoder[o*mm+a]*phi[a]; + float residual = sample(source,w,h,nout,o,x,y)-predicted; + float slope[4]; + for (int a = 0; a < channels; a++) { + float derivative = decoder[o*mm+(level == 0 ? c1+a : a)]; + if (level == 0) { + for (int b = 0; b < c1; b++) derivative += decoder[o*mm+c1+c0+a*c1+b]*c[b]; + } else { + for (int b = 0; b < c0; b++) derivative += decoder[o*mm+c1+c0+b*c1+a]*s[b]; + } + slope[a] = q*derivative; + } + for (int a = 0; a < channels; a++) { + rhs[a] += slope[a]*residual; + for (int b = 0; b < channels; b++) matrix[a*4+b] += slope[a]*slope[b]; + } + } + } + float trace = 0.0f; + for (int a = 0; a < channels; a++) trace += matrix[a*4+a]; + // Positive ridge makes the symmetric system definite even for constant input. + for (int a = 0; a < channels; a++) matrix[a*4+a] += max(1.0e-8f,trace*1.0e-5f); + for (int a = 0; a < channels; a++) { + float diagonal = matrix[a*4+a]; + for (int b = a; b < channels; b++) matrix[a*4+b] /= diagonal; + rhs[a] /= diagonal; + for (int row = 0; row < channels; row++) if (row != a) { + float factor = matrix[row*4+a]; + for (int b = a; b < channels; b++) matrix[row*4+b] -= factor*matrix[a*4+b]; + rhs[row] -= factor*rhs[a]; + } + } + // Bounded trust step retains numerical stability in nearly rank-deficient fits. + float largest = 0.5f; + for (int a = 0; a < channels; a++) largest = fmaxf(largest,fabsf(rhs[a])); + for (int a = 0; a < channels; a++) result[id*channels+a] += rhs[a]*(0.5f/largest); +} + +__global__ void decode(const float* latent0, const float* latent1, const float* decoder, float* output, + int w, int h, int w1, int h1, int c0, int c1, int nout) { + int id = blockIdx.x*blockDim.x+threadIdx.x; + if (id >= w*h) return; + float x = (float)(id%w), y = (float)(id/w), s[4], c[4], phi[25]; + for (int i = 0; i < c0; i++) s[i] = latent0[id*c0+i]; + for (int i = 0; i < c1; i++) c[i] = sample(latent1,w1,h1,c1,i,(x+0.5f)*(float)w1/(float)w-0.5f,(y+0.5f)*(float)h1/(float)h-0.5f); + int k = 0; + for (int i = 0; i < c1; i++) phi[k++] = c[i]; + for (int i = 0; i < c0; i++) phi[k++] = s[i]; + for (int i = 0; i < c0; i++) for (int j = 0; j < c1; j++) phi[k++] = s[i]*c[j]; + phi[k++] = 1.0f; + for (int o = 0; o < nout; o++) { + float value = 0.0f; + for (int i = 0; i < k; i++) value += decoder[o*k+i]*phi[i]; + output[id*nout+o] = value; + } +} + diff --git a/webgpu/encoder/prepare.js b/webgpu/encoder/prepare.js new file mode 100644 index 0000000..ff773dd --- /dev/null +++ b/webgpu/encoder/prepare.js @@ -0,0 +1,87 @@ +// SPDX-License-Identifier: Apache-2.0 +export function validateInputs(images,{c0=2,c1=4,iterations=6,sites=5}={}) { + if(!Array.isArray(images)||images.length<1||images.length>6) throw new Error('Select one to six RGB images.'); + if(![c0,c1].every(c=>Number.isInteger(c)&&c>=1&&c<=4)) throw new Error('Latent channels must be 1..4.'); + if(!Number.isInteger(iterations)||iterations<0||iterations>100) throw new Error('Iterations must be 0..100.'); + if(![1,5].includes(sites)) throw new Error('Use one center site or five filtered sites.'); + const {width,height}=images[0]; + if(![width,height].every(v=>Number.isInteger(v)&&v>0&&v<=16384)) throw new Error('Image dimensions must be 1..16384.'); + for(const image of images) { + if(image.width!==width||image.height!==height) throw new Error('All material images must have the same dimensions.'); + if(!(image.data instanceof Uint8Array || image.data instanceof Uint8ClampedArray)||image.data.length!==width*height*4) throw new Error('Expected RGBA8 image data.'); + } + return {width,height}; +} + +// A shared PCA basis across the mip chain, with deterministic Jacobi rotations. +function basis(data,channels,count) { + const mean=new Float64Array(channels),cov=new Float64Array(channels*channels),vectors=new Float64Array(channels*channels),n=data.length/channels; + for(let i=0;ilargest) {p=a;q=b;largest=Math.abs(cov[a*channels+b]);} + if(largest<1e-12) break; + const angle=0.5*Math.atan2(2*cov[p*channels+q],cov[q*channels+q]-cov[p*channels+p]),c=Math.cos(angle),s=Math.sin(angle); + for(let a=0;ai).sort((a,b)=>cov[b*channels+b]-cov[a*channels+a]); + return {mean,axes:Array.from({length:count},(_,i)=>Array.from({length:channels},(_,j)=>vectors[j*channels+order[i%channels]]))}; +} + +function resize(data,w,h,nw,nh,channels) { + const result=new Float32Array(nw*nh*channels); + // Area averaging including partial source cells for non-power-of-two mips. + for(let y=0;y>1,nh=ph>>1; + // DDS halves each latent independently. Stop before rounding would make + // floor(full/4) disagree with the DDS half-size of the previous latent. + if(Math.floor(nw/4)!==(w1>>1)||Math.floor(nh/4)!==(h1>>1)) break; + src=resize(src,pw,ph,nw,nh,nout);pw=nw;ph=nh; + } + const pca1=basis(planes[0].coarse,nout,c1),baseResidual=detailResidual(planes[0],nout),pca0=basis(baseResidual,nout,c0); + for(const p of planes) { + p.values0=new Float32Array(p.w*p.h*c0);p.values1=new Float32Array(p.w1*p.h1*c1); + for(const [data,target,channels,pca] of [[p===planes[0]?baseResidual:detailResidual(p,nout),p.values0,c0,pca0],[p.coarse,p.values1,c1,pca1]]) for(let i=0;ia:last-child{font-size:13px;color:#bcc5c0}main{max-width:1500px;margin:auto;padding:52px 5vw 20px}.intro{position:relative;max-width:820px;margin-bottom:42px}.eyebrow{color:#b9ee96;font-size:11px;letter-spacing:2px;font-weight:650}h1{font-size:clamp(36px,4.5vw,60px);font-weight:500;letter-spacing:-2px;line-height:1.06;margin:20px 0}.intro p{color:#a1ada6;max-width:570px;line-height:1.7;font-size:15px}.workspace{display:grid;grid-template-columns:350px minmax(0,1fr);border:1px solid #343c38;border-radius:12px;overflow:hidden;background:#151b18}form{padding:26px;border-right:1px solid #343c38}fieldset{padding:0;border:0;margin:0;min-width:0}.section-heading{display:flex;align-items:center;gap:10px;margin-bottom:18px}.section-heading>span{font-size:10px;color:#b9ee96;border:1px solid #4a5e42;border-radius:20px;padding:4px 7px}h2{font-size:14px;font-weight:550;margin:0}.dropzone{border:1px dashed #5b6d56;border-radius:8px;display:flex;align-items:center;text-align:center;padding:26px 10px;cursor:pointer;background:#1b251b;gap:8px}.dropzone:hover,.dropzone.drag{border-color:#b9ee96;background:#253522}.dropzone strong{font-size:13px;font-weight:500}.dropzone span{font-size:10px;color:#9ca895}.dropzone input{font-size:10px;max-width:100%;border:0;background:transparent;margin-top:8px}label{display:flex;flex-direction:column;gap:8px;font-size:11px;color:#c5cec8;margin-bottom:16px}input,select{border:1px solid #3a443e;background:#111814;color:#edf2ef;border-radius:5px;padding:10px;font:inherit;min-width:0}input:focus,select:focus,button:focus-visible{outline:2px solid #b9ee96;outline-offset:2px}.fields{display:grid;grid-template-columns:1fr 1fr;gap:12px}.check{flex-direction:row;align-items:center;margin:12px 0}.check input{accent-color:#b9ee96}.note{font-size:11px;line-height:1.65;color:#89978e;margin:16px 0 24px}.inputs{display:flex;gap:6px;flex-wrap:wrap}.inputs img{width:42px;height:42px;object-fit:cover;border-radius:4px}.text-button{background:none!important;color:#b9ee96!important;padding:0!important;font-size:11px!important;border:0!important;margin-bottom:14px}.actions{display:flex;gap:10px;flex-wrap:wrap}button{cursor:pointer;border:1px solid #b9ee96;border-radius:5px;background:#b9ee96;color:#152013;font:600 12px inherit;padding:12px 16px;font-family:inherit;font-size:12px}button:hover{filter:brightness(1.08)}button:disabled{opacity:.4;cursor:default}.secondary,#cancel{background:transparent;color:#b9ee96}#encode{flex:1;text-align:left;display:flex;justify-content:space-between}.status{display:flex;gap:8px;align-items:center;min-height:52px}.status p{font-size:11px;line-height:1.5;color:#a9b7ad}#indicator{width:6px;height:6px;flex-shrink:0;background:#68736c;border-radius:50%}.busy #indicator{background:#b9ee96;animation:pulse 1s infinite}.error #indicator{background:#ff9b7c}.error #status{color:#ffb49c}.inspection{padding:26px;min-height:620px}#empty{min-height:460px;display:flex;flex-direction:column;align-items:center;justify-content:center;color:#85978b;text-align:center}#empty h3{color:#becac2;font-size:17px;font-weight:400;margin-bottom:0}#empty p{font-size:12px;line-height:1.7;max-width:250px}.grid-art{width:120px;height:120px;border:1px solid #51664b;transform:rotate(-12deg);background:repeating-linear-gradient(0deg,transparent 0 23px,#35422f 24px 25px),repeating-linear-gradient(90deg,transparent 0 23px,#35422f 24px 25px);box-shadow:13px 13px 0 #1d291c}.comparison{display:grid;grid-template-columns:1fr 1fr;gap:16px;margin-top:30px}figure{margin:0;min-width:0}canvas{display:block;width:100%;height:auto;max-height:500px;object-fit:contain;background:#101612;border:1px solid #344136;border-radius:4px}figcaption{font-size:9px;letter-spacing:1.5px;color:#8d9e92;margin-top:12px}.metrics{display:grid;grid-template-columns:repeat(3,1fr);padding:26px 0;margin:28px 0;border-top:1px solid #344136;border-bottom:1px solid #344136;gap:14px}.metrics strong{display:block;font-size:22px;font-weight:400;color:#d3edc6}.metrics span{font-size:10px;color:#91a28f;display:block;margin-top:6px}#texture{margin-left:auto;font-size:11px;padding:6px;max-width:160px}#viewer-section{margin-top:30px}#viewer{width:100%;height:740px;border:1px solid #394139;border-radius:8px}#close-viewer{margin:0 0 0 auto}footer{font-size:10px;color:#748479;padding:30px 0}footer a{text-decoration:underline}[hidden]{display:none!important}@keyframes pulse{50%{opacity:.25}}@media(max-width:900px){.workspace{grid-template-columns:300px minmax(0,1fr)}.comparison{grid-template-columns:1fr}.metrics{grid-template-columns:1fr}.inspection{padding:20px}}@media(max-width:650px){main{padding:30px 20px}.workspace{grid-template-columns:1fr}form{border-right:0;border-bottom:1px solid #343c38}.comparison{grid-template-columns:1fr 1fr}.metrics{grid-template-columns:repeat(3,1fr)}header{padding:0 20px}.brand span{display:none}#viewer{height:700px}}@media(prefers-reduced-motion:reduce){.busy #indicator{animation:none}} diff --git a/webgpu/encoder/zip.js b/webgpu/encoder/zip.js new file mode 100644 index 0000000..bf4d6c2 --- /dev/null +++ b/webgpu/encoder/zip.js @@ -0,0 +1,22 @@ +// SPDX-License-Identifier: Apache-2.0 +// Uncompressed ZIP (DDS is already block compressed); one explicit download. +export function zipFiles(files) { + const table=Uint32Array.from({length:256},(_,n)=>{for(let k=0;k<8;k++) n=(n&1)?0xedb88320^(n>>>1):n>>>1;return n>>>0;}); + const local=[],central=[];let offset=0,centralSize=0; + for(const [name,data] of files) { + const encoded=new TextEncoder().encode(name);let crc=0xffffffff; + for(const byte of data) crc=table[(crc^byte)&255]^(crc>>>8); + crc=(crc^0xffffffff)>>>0; + const header=new Uint8Array(30+encoded.length),v=new DataView(header.buffer); + v.setUint32(0,0x04034b50,true);v.setUint16(4,20,true);v.setUint16(6,0x800,true);v.setUint16(12,33,true); + v.setUint32(14,crc,true);v.setUint32(18,data.length,true);v.setUint32(22,data.length,true);v.setUint16(26,encoded.length,true);header.set(encoded,30); + local.push(header,data); + const entry=new Uint8Array(46+encoded.length),e=new DataView(entry.buffer); + e.setUint32(0,0x02014b50,true);e.setUint16(4,20,true);e.setUint16(6,20,true);e.setUint16(8,0x800,true);e.setUint16(14,33,true); + e.setUint32(16,crc,true);e.setUint32(20,data.length,true);e.setUint32(24,data.length,true);e.setUint16(28,encoded.length,true);e.setUint32(42,offset,true);entry.set(encoded,46); + central.push(entry);centralSize+=entry.length;offset+=header.length+data.length; + } + const end=new Uint8Array(22),v=new DataView(end.buffer); + v.setUint32(0,0x06054b50,true);v.setUint16(8,files.size,true);v.setUint16(10,files.size,true);v.setUint32(12,centralSize,true);v.setUint32(16,offset,true); + return new Blob([...local,...central,end],{type:'application/zip'}); +} diff --git a/webgpu/index.html b/webgpu/index.html index 2f5cab4..1a9b587 100644 --- a/webgpu/index.html +++ b/webgpu/index.html @@ -11,6 +11,7 @@

NNTC latent viewer WebGPU v1.00

+ Encode
diff --git a/webgpu/main.js b/webgpu/main.js index ed4c1a0..7f095ac 100644 --- a/webgpu/main.js +++ b/webgpu/main.js @@ -1646,6 +1646,18 @@ async function main() { requestAnimationFrame(frame); canvas.focus(); + // The same-origin encoder hands off in-memory Files through the normal intake. + // No uploaded data or persistent browser storage is needed for this preview. + if (flag('encoder') && window.parent !== window) { + window.addEventListener('message', (event) => { + if (event.origin !== location.origin || event.source !== window.parent || + event.data?.type !== 'nntc-encoded-files' || !Array.isArray(event.data.files) || + !event.data.files.every(file => file instanceof File)) return; + openFiles(event.data.files); + }); + window.parent.postMessage({type:'nntc-viewer-ready'}, location.origin); + } + // LAST, and after the first frame is scheduled: the shot renders offscreen and does not depend on the canvas // having drawn, but everything it does depend on - the adapter, the shader, the asset - has happened by here, // and a failure before this point has already written its own message and left the title alone. From bcf15583e44d3526e3639485c462328356fdcca9 Mon Sep 17 00:00:00 2001 From: SamG-Coder <13449963+SamG-Coder@users.noreply.github.com> Date: Wed, 23 Sep 2026 08:49:28 +1000 Subject: [PATCH 2/2] Add upstream decoder verification for browser ZIP exports --- docs/BROWSER_ENCODER.md | 6 ++++++ tests/check-web-exports.py | 26 ++++++++++++++++++++++++++ tests/web-encoder-gpu.mjs | 1 - tests/web-encoder-ui.mjs | 1 - webgpu/encoder/kernels.cu | 1 - 5 files changed, 32 insertions(+), 3 deletions(-) create mode 100644 tests/check-web-exports.py diff --git a/docs/BROWSER_ENCODER.md b/docs/BROWSER_ENCODER.md index f12ecbe..197f3ce 100644 --- a/docs/BROWSER_ENCODER.md +++ b/docs/BROWSER_ENCODER.md @@ -88,6 +88,8 @@ site or whole-mip-chain error measurement. npm test npm run test:gpu npm run test:ui +# Optional independent upstream decoder (requires numpy and Pillow): +python tests/check-web-exports.py ``` The Node tests compile every CUDA entry and check all 32 combinations of latent @@ -100,6 +102,10 @@ reconstruction against the GPU. The UI test encodes the bundled single-map and four-map materials, downloads ZIPs, loads the interactive viewer, and checks a mobile viewport and browser errors. +The Python cross-check also validates ZIP CRCs and runs the original repository's +strict DDS reader and decoder on both downloaded assets. It reports saturated, +rounded RGB8 PSNR, which can differ from the browser's unclamped float PSNR. + On Windows the browser tests use installed Edge. Set `NNTC_BROWSER` to another Playwright browser channel. On other platforms they use Playwright Chromium; install it with `npx playwright install chromium`. `NNTC_SOFTWARE_GPU=1` requests diff --git a/tests/check-web-exports.py b/tests/check-web-exports.py new file mode 100644 index 0000000..3ee9d76 --- /dev/null +++ b/tests/check-web-exports.py @@ -0,0 +1,26 @@ +"""Verify browser ZIP exports with upstream's independent Python decoder. + +Run npm run test:ui first, then python tests/check-web-exports.py. +Requires numpy and Pillow (the existing decoder's dependencies). +""" +import pathlib +import sys +import tempfile +import zipfile + +ROOT = pathlib.Path(__file__).resolve().parents[1] +sys.path.insert(0, str(ROOT / 'tools')) +import dds_decode + +for archive, sources in [('m1.zip', ['m1.png']), ('m1-m4.zip', ['m1.png', 'm2.png', 'm3.png', 'm4.png'])]: + with tempfile.TemporaryDirectory() as temporary: + directory = pathlib.Path(temporary) + with zipfile.ZipFile(ROOT / 'out' / 'web-tests' / archive) as bundle: + assert bundle.testzip() is None, 'ZIP CRC failure' + for name in bundle.namelist(): + assert pathlib.PurePosixPath(name).name == name and '\\' not in name, 'Expected flat ZIP entries' + (directory / name).write_bytes(bundle.read(name)) + descriptor = directory / 'm1_nntc.json' + # Exercise upstream's header, format and every mip decoder, not just our JS reader. + sys.argv = ['dds_decode.py', str(descriptor), '--psnr'] + [str(ROOT / 'examples' / name) for name in sources] + assert dds_decode.main() == 0 diff --git a/tests/web-encoder-gpu.mjs b/tests/web-encoder-gpu.mjs index fecc318..4eb8f84 100644 --- a/tests/web-encoder-gpu.mjs +++ b/tests/web-encoder-gpu.mjs @@ -21,4 +21,3 @@ try { await mkdir('out/web-tests',{recursive:true}); await writeFile('out/web-tests/gpu.json',JSON.stringify(results,null,2)); } finally {await browser?.close();await new Promise(resolve=>server.close(resolve));} - diff --git a/tests/web-encoder-ui.mjs b/tests/web-encoder-ui.mjs index 478d662..5a2fefa 100644 --- a/tests/web-encoder-ui.mjs +++ b/tests/web-encoder-ui.mjs @@ -38,4 +38,3 @@ try { assert.deepEqual(errors,[]); await writeFile('out/web-tests/ui.json',JSON.stringify({status,metrics,viewerStatus,fourMapMaterial:material,errors},null,2)); }finally{await browser?.close();await new Promise(resolve=>server.close(resolve));} - diff --git a/webgpu/encoder/kernels.cu b/webgpu/encoder/kernels.cu index 55ba9b9..bf664e4 100644 --- a/webgpu/encoder/kernels.cu +++ b/webgpu/encoder/kernels.cu @@ -145,4 +145,3 @@ __global__ void decode(const float* latent0, const float* latent1, const float* output[id*nout+o] = value; } } -