1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
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));
}
|
