diff options
| author | skal <pascal.massimino@gmail.com> | 2026-02-15 18:44:17 +0100 |
|---|---|---|
| committer | skal <pascal.massimino@gmail.com> | 2026-02-15 18:44:17 +0100 |
| commit | 161a59fa50bb92e3664c389fa03b95aefe349b3f (patch) | |
| tree | 71548f64b2bdea958388f9063b74137659d70306 /cnn_v2/shaders/cnn_v2_layer_template.wgsl | |
| parent | 9c3b72c710bf1ffa7e18f7c7390a425d57487eba (diff) | |
refactor(cnn): isolate CNN v2 to cnn_v2/ subdirectory
Move all CNN v2 files to dedicated cnn_v2/ directory to prepare for CNN v3 development. Zero functional changes.
Structure:
- cnn_v2/src/ - C++ effect implementation
- cnn_v2/shaders/ - WGSL shaders (6 files)
- cnn_v2/weights/ - Binary weights (3 files)
- cnn_v2/training/ - Python training scripts (4 files)
- cnn_v2/scripts/ - Shell scripts (train_cnn_v2_full.sh)
- cnn_v2/tools/ - Validation tools (HTML)
- cnn_v2/docs/ - Documentation (4 markdown files)
Changes:
- Update CMake source list to cnn_v2/src/cnn_v2_effect.cc
- Update assets.txt with relative paths to cnn_v2/
- Update includes to ../../cnn_v2/src/cnn_v2_effect.h
- Add PROJECT_ROOT resolution to Python/shell scripts
- Update doc references in HOWTO.md, TODO.md
- Add cnn_v2/README.md
Verification: 34/34 tests passing, demo runs correctly.
Co-Authored-By: Claude Sonnet 4.5 <noreply@anthropic.com>
Diffstat (limited to 'cnn_v2/shaders/cnn_v2_layer_template.wgsl')
| -rw-r--r-- | cnn_v2/shaders/cnn_v2_layer_template.wgsl | 68 |
1 files changed, 68 insertions, 0 deletions
diff --git a/cnn_v2/shaders/cnn_v2_layer_template.wgsl b/cnn_v2/shaders/cnn_v2_layer_template.wgsl new file mode 100644 index 0000000..1bf6819 --- /dev/null +++ b/cnn_v2/shaders/cnn_v2_layer_template.wgsl @@ -0,0 +1,68 @@ +// CNN v2 Layer Template (placeholder for generated shaders) +// This file documents the structure - actual layers generated by export script + +// Example: Layer 0 (1×1 kernel, 8→16 channels) +// const KERNEL_SIZE: u32 = 1u; +// const IN_CHANNELS: u32 = 8u; // 7 features + bias +// const OUT_CHANNELS: u32 = 16u; +// const weights: array<f32, 128> = array(...); + +@group(0) @binding(0) var static_features: texture_2d<u32>; +@group(0) @binding(1) var layer_input: texture_2d<u32>; // Previous layer output +@group(0) @binding(2) var output_tex: texture_storage_2d<rgba32uint, write>; + +fn unpack_static_features(coord: vec2<i32>) -> array<f32, 8> { + let packed = textureLoad(static_features, coord, 0); + let v0 = unpack2x16float(packed.x); + let v1 = unpack2x16float(packed.y); + let v2 = unpack2x16float(packed.z); + let v3 = unpack2x16float(packed.w); + return array<f32, 8>(v0.x, v0.y, v1.x, v1.y, v2.x, v2.y, v3.x, v3.y); +} + +fn unpack_layer_channels(coord: vec2<i32>) -> array<f32, 8> { + let packed = textureLoad(layer_input, coord, 0); + let v0 = unpack2x16float(packed.x); + let v1 = unpack2x16float(packed.y); + let v2 = unpack2x16float(packed.z); + let v3 = unpack2x16float(packed.w); + return array<f32, 8>(v0.x, v0.y, v1.x, v1.y, v2.x, v2.y, v3.x, v3.y); +} + +fn pack_channels(values: array<f32, 8>) -> vec4<u32> { + return vec4<u32>( + pack2x16float(vec2<f32>(values[0], values[1])), + pack2x16float(vec2<f32>(values[2], values[3])), + pack2x16float(vec2<f32>(values[4], values[5])), + pack2x16float(vec2<f32>(values[6], values[7])) + ); +} + +@compute @workgroup_size(8, 8) +fn main(@builtin(global_invocation_id) id: vec3<u32>) { + let coord = vec2<i32>(id.xy); + let dims = textureDimensions(static_features); + + if (coord.x >= i32(dims.x) || coord.y >= i32(dims.y)) { + return; + } + + // Load static features (always available) + let static_feat = unpack_static_features(coord); + + // Convolution loop (example for generated code) + // var output: array<f32, OUT_CHANNELS>; + // for (var c: u32 = 0u; c < OUT_CHANNELS; c++) { + // var sum: f32 = 0.0; + // for (var ky: i32 = -radius; ky <= radius; ky++) { + // for (var kx: i32 = -radius; kx <= radius; kx++) { + // let sample_coord = coord + vec2<i32>(kx, ky); + // // Load static + prev layer, multiply weights, accumulate + // } + // } + // output[c] = max(0.0, sum); // ReLU + // } + + // Placeholder output + textureStore(output_tex, coord, vec4<u32>(0u)); +} |
