diff options
Diffstat (limited to 'tools')
| -rw-r--r-- | tools/cnn_v2_test/index.html | 121 |
1 files changed, 114 insertions, 7 deletions
diff --git a/tools/cnn_v2_test/index.html b/tools/cnn_v2_test/index.html index ba1434c..9a5b2c0 100644 --- a/tools/cnn_v2_test/index.html +++ b/tools/cnn_v2_test/index.html @@ -323,6 +323,17 @@ <div class="drop-zone" id="weightsDrop" onclick="document.getElementById('weightsFile').click()"> Drop .bin Weights or Click to Browse </div> + <div class="panel"> + <div class="panel-header">Features (p0-p3)</div> + <div class="panel-content"> + <label for="mipLevel" style="font-size: 11px;">Mip Level:</label> + <select id="mipLevel" style="width: 100%; background: #1a1a1a; color: #e0e0e0; border: 1px solid #404040; padding: 4px; margin-top: 4px;"> + <option value="0">Mip 0 (original)</option> + <option value="1">Mip 1 (half res)</option> + <option value="2">Mip 2 (quarter res)</option> + </select> + </div> + </div> <div class="panel" id="weightsInfoPanel"> <div class="panel-header">Weights Info</div> <div class="panel-content" id="weightsInfo"> @@ -367,8 +378,10 @@ const DEFAULT_WEIGHTS_B64 = 'Q05OMgEAAAADAAAAEAUAAAMAAAAMAAAABAAAAAAAAACwAQAAAwA const STATIC_SHADER = ` @group(0) @binding(0) var input_tex: texture_2d<f32>; -@group(0) @binding(1) var depth_tex: texture_2d<f32>; -@group(0) @binding(2) var output_tex: texture_storage_2d<rgba32uint, write>; +@group(0) @binding(1) var input_sampler: sampler; +@group(0) @binding(2) var depth_tex: texture_2d<f32>; +@group(0) @binding(3) var output_tex: texture_storage_2d<rgba32uint, write>; +@group(0) @binding(4) var<uniform> mip_level: u32; @compute @workgroup_size(8, 8) fn main(@builtin(global_invocation_id) id: vec3<u32>) { @@ -376,7 +389,8 @@ fn main(@builtin(global_invocation_id) id: vec3<u32>) { let dims = textureDimensions(input_tex); if (coord.x >= i32(dims.x) || coord.y >= i32(dims.y)) { return; } - let rgba = textureLoad(input_tex, coord, 0); + let uv = (vec2<f32>(coord) + 0.5) / vec2<f32>(dims); + let rgba = textureSampleLevel(input_tex, input_sampler, uv, f32(mip_level)); let d = textureLoad(depth_tex, coord, 0).r; let uv_x = f32(coord.x) / f32(dims.x); let uv_y = 1.0 - (f32(coord.y) / f32(dims.y)); @@ -627,6 +641,7 @@ class CNNTester { this.isVideo = false; this.fps = 30; this.isProcessing = false; + this.mipLevel = 0; this.init(); } @@ -908,6 +923,72 @@ class CNNTester { this.visualizeWeights(0); } + generateMipmaps(texture, width, height) { + if (!this.mipmapPipeline) { + const mipmapShader = ` + @group(0) @binding(0) var src: texture_2d<f32>; + @vertex + fn vs_main(@builtin(vertex_index) idx: u32) -> @builtin(position) vec4<f32> { + var pos = array<vec2<f32>, 6>( + vec2<f32>(-1.0, -1.0), vec2<f32>(1.0, -1.0), vec2<f32>(-1.0, 1.0), + vec2<f32>(-1.0, 1.0), vec2<f32>(1.0, -1.0), vec2<f32>(1.0, 1.0) + ); + return vec4<f32>(pos[idx], 0.0, 1.0); + } + @fragment + fn fs_main(@builtin(position) pos: vec4<f32>) -> @location(0) vec4<f32> { + let coord = vec2<i32>(i32(pos.x) * 2, i32(pos.y) * 2); + var sum = vec4<f32>(0.0); + for (var y: i32 = 0; y < 2; y++) { + for (var x: i32 = 0; x < 2; x++) { + sum += textureLoad(src, coord + vec2<i32>(x, y), 0); + } + } + return sum * 0.25; + } + `; + this.mipmapPipeline = this.device.createRenderPipeline({ + layout: 'auto', + vertex: { module: this.device.createShaderModule({ code: mipmapShader }), entryPoint: 'vs_main' }, + fragment: { + module: this.device.createShaderModule({ code: mipmapShader }), + entryPoint: 'fs_main', + targets: [{ format: 'rgba8unorm' }] + } + }); + } + + const encoder = this.device.createCommandEncoder(); + + for (let mip = 1; mip < 3; mip++) { + const mipWidth = Math.max(1, width >> mip); + const mipHeight = Math.max(1, height >> mip); + + const bindGroup = this.device.createBindGroup({ + layout: this.mipmapPipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: texture.createView({ baseMipLevel: mip - 1, mipLevelCount: 1 }) } + ] + }); + + const renderPass = encoder.beginRenderPass({ + colorAttachments: [{ + view: texture.createView({ baseMipLevel: mip, mipLevelCount: 1 }), + loadOp: 'clear', + storeOp: 'store' + }] + }); + + renderPass.setPipeline(this.mipmapPipeline); + renderPass.setBindGroup(0, bindGroup); + renderPass.setViewport(0, 0, mipWidth, mipHeight, 0, 1); + renderPass.draw(6); + renderPass.end(); + } + + this.device.queue.submit([encoder.finish()]); + } + displayOriginal() { const source = this.isVideo ? this.video : this.image; if (!source || !this.device) return; @@ -990,20 +1071,24 @@ class CNNTester { this.context.configure({ device: this.device, format: this.format }); - // Create persistent input texture for original view + // Create persistent input texture for original view with mipmaps if (this.inputTexture) this.inputTexture.destroy(); this.inputTexture = this.device.createTexture({ size: [width, height], format: 'rgba8unorm', + mipLevelCount: 3, usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.COPY_DST | GPUTextureUsage.RENDER_ATTACHMENT }); this.device.queue.copyExternalImageToTexture( { source: source }, - { texture: this.inputTexture }, + { texture: this.inputTexture, mipLevel: 0 }, [width, height] ); + // Generate mipmaps + this.generateMipmaps(this.inputTexture, width, height); + const depthTex = this.device.createTexture({ size: [width, height], format: 'r32float', @@ -1076,12 +1161,28 @@ class CNNTester { const encoder = this.device.createCommandEncoder(); + const mipLevelBuffer = this.device.createBuffer({ + size: 4, + usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST + }); + this.device.queue.writeBuffer(mipLevelBuffer, 0, new Uint32Array([this.mipLevel])); + + if (!this.linearSampler) { + this.linearSampler = this.device.createSampler({ + magFilter: 'linear', + minFilter: 'linear', + mipmapFilter: 'linear' + }); + } + const staticBG = this.device.createBindGroup({ layout: staticPipeline.getBindGroupLayout(0), entries: [ { binding: 0, resource: this.inputTexture.createView() }, - { binding: 1, resource: depthTex.createView() }, - { binding: 2, resource: staticTex.createView() } + { binding: 1, resource: this.linearSampler }, + { binding: 2, resource: depthTex.createView() }, + { binding: 3, resource: staticTex.createView() }, + { binding: 4, resource: { buffer: mipLevelBuffer } } ] }); @@ -1731,6 +1832,12 @@ document.getElementById('depth').addEventListener('input', e => { if ((tester.image || tester.isVideo) && tester.weights) tester.run(); }); +document.getElementById('mipLevel').addEventListener('change', e => { + tester.mipLevel = parseInt(e.target.value); + tester.log(`Mip level changed to ${e.target.value}`); + if ((tester.image || tester.isVideo) && tester.weights) tester.run(); +}); + document.getElementById('playPauseBtn').addEventListener('click', () => tester.togglePlayPause()); document.getElementById('stepBackBtn').addEventListener('click', () => tester.stepFrame(-1)); document.getElementById('stepForwardBtn').addEventListener('click', () => tester.stepFrame(1)); |