Add qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5 (attention bounded by the used context)

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/004c06d76ce77979b63fe4450df54986355fd10cd2ef62adae86780e49e4a4a7.wgsl

@@ -0,0 +1,10 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 16384u;
+if (i >= 16384u) { return; }
+out[i] = f32(f32(b0[i]) + (f32(b1[i]) * 1.0));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/0112d93bb1aa347d8887ce1b95610a1a21fc8450ff8f59bcd37747fa6bf6fa52.wgsl

@@ -0,0 +1,10 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 2048u;
+if (i >= 2048u) { return; }
+out[i] = f32(f32(b0[i]) + (f32(b1[i]) * 1.0));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/07cfabe2e2c2feff22a06c8296600e5ca21efb7444419f0a05e5ff6f48e0292d.wgsl

@@ -0,0 +1,10 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 1024u;
+if (i >= 1024u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(-x);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/0822a85cfcb62e2c4cecbf9eaa8225c089c1bb50eb728967d0344d881d0b5d89.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 64u;
+if (i >= 64u) { return; }
+out[i] = f32(b0[2u * 64u + ((i / 64u) % 1u) * 64u + ((i / 64u) % 1u) * 64u + ((i / 1u) % 64u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/091df2c4538eedd1af1dc3f6f6145095ed29510a93a911c0f6bb5183e9de29a3.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 4096u;
+if (i >= 4096u) { return; }
+out[i] = f32(b0[((i / 4096u) % 1u) * 4096u + ((i / 512u) % 8u) * 128u + ((i / 128u) % 4u) * 1024u + ((i / 1u) % 128u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/0baaeed75a52bdecaf32e25a919cc33aad5ddc8da4296d3373f9a6b4395ba696.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 4096u;
+if (i >= 4096u) { return; }
+out[i] = f32(b0[0u * 4096u + ((i / 4096u) % 1u) * 4096u + ((i / 64u) % 64u) * 64u + ((i / 1u) % 64u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/0ef8fd3c2fe267549e7de4eb9a074d2ada8e80c91ce2f9c9c4d6b24d17742821.wgsl

@@ -0,0 +1,71 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<u32>;
+@group(0) @binding(2) var<storage, read> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b1[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 64u + col / 32u];
+}
+
+var<workgroup> tile_a: array<array<f32, 16>, 64>;
+var<workgroup> tile_b: array<array<f32, 64>, 16>;
+@compute @workgroup_size(16, 16)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let lane = local.y * 16u + local.x;
+    let batch = group.z;
+    let tile_n = group.x * 64u;
+    let tile_row = group.y * 64u;
+    var acc: array<array<f32, 4>, 4>;
+    for (var k0 = 0u; k0 < 2048u; k0 += 16u) {
+        for (var e = 0u; e < 4u; e++) {
+            let flat = lane + e * 256u;
+            let m_local = flat / 16u;
+            let row = tile_row + m_local;
+            let col = k0 + flat % 16u;
+            var value = 0.0;
+            if (row < 64u && col < 2048u) { value = f32(f32(b0[(batch * 64u + row) * 2048u + col])); }
+            tile_a[m_local][flat % 16u] = value;
+        }
+        if (lane < 256u) {
+            let n_local = lane / 4u;
+            let word_local = lane % 4u;
+            let n_index = tile_n + n_local;
+            let first = k0 + word_local * 4u;
+            var word = 0u;
+            var scale = 0.0;
+            if (n_index < 1024u && first < 2048u) {
+                word = b1[n_index * 512u + first / 4u];
+                scale = f32(b2[n_index * 64u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 2048u) { value = f32(code) * scale; }
+                tile_b[word_local * 4u + e][n_local] = value;
+            }
+        }
+        workgroupBarrier();
+        for (var kk = 0u; kk < 16u; kk++) {
+            var b_values: array<f32, 4>;
+            for (var c = 0u; c < 4u; c++) { b_values[c] = tile_b[kk][local.x * 4u + c]; }
+            for (var r = 0u; r < 4u; r++) {
+                let a_value = tile_a[local.y * 4u + r][kk];
+                for (var c = 0u; c < 4u; c++) { acc[r][c] += a_value * b_values[c]; }
+            }
+        }
+        workgroupBarrier();
+    }
+    for (var r = 0u; r < 4u; r++) {
+        let out_row = tile_row + local.y * 4u + r;
+        for (var c = 0u; c < 4u; c++) {
+            let out_col = tile_n + local.x * 4u + c;
+            if (out_row < 64u && out_col < 1024u) {
+                let i = (batch * 64u + out_row) * 1024u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/0f1b24ac8390100761206441b3050b3efa07a167129149f4900aff8d118ee8a4.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 8192u;
+if (i >= 8192u) { return; }
+out[i] = f32(b0[((i / 8192u) % 1u) * 16384u + ((i / 1024u) % 8u) * 2048u + ((i / 64u) % 16u) * 128u + (((i / 1u) % 64u) * 1u + 64u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/1024c1444e2cb17a211ebed96e18db764f1809d6e986d09e75f482ab4e61a07b.wgsl

@@ -0,0 +1,71 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<u32>;
+@group(0) @binding(2) var<storage, read> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b1[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 64u + col / 32u];
+}
+
+var<workgroup> tile_a: array<array<f32, 16>, 16>;
+var<workgroup> tile_b: array<array<f32, 64>, 16>;
+@compute @workgroup_size(16, 16)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let lane = local.y * 16u + local.x;
+    let batch = group.z;
+    let tile_n = group.x * 64u;
+    let tile_row = group.y * 16u;
+    var acc: array<array<f32, 4>, 1>;
+    for (var k0 = 0u; k0 < 2048u; k0 += 16u) {
+        for (var e = 0u; e < 1u; e++) {
+            let flat = lane + e * 256u;
+            let m_local = flat / 16u;
+            let row = tile_row + m_local;
+            let col = k0 + flat % 16u;
+            var value = 0.0;
+            if (row < 16u && col < 2048u) { value = f32(f32(b0[(batch * 16u + row) * 2048u + col])); }
+            tile_a[m_local][flat % 16u] = value;
+        }
+        if (lane < 256u) {
+            let n_local = lane / 4u;
+            let word_local = lane % 4u;
+            let n_index = tile_n + n_local;
+            let first = k0 + word_local * 4u;
+            var word = 0u;
+            var scale = 0.0;
+            if (n_index < 2048u && first < 2048u) {
+                word = b1[n_index * 512u + first / 4u];
+                scale = f32(b2[n_index * 64u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 2048u) { value = f32(code) * scale; }
+                tile_b[word_local * 4u + e][n_local] = value;
+            }
+        }
+        workgroupBarrier();
+        for (var kk = 0u; kk < 16u; kk++) {
+            var b_values: array<f32, 4>;
+            for (var c = 0u; c < 4u; c++) { b_values[c] = tile_b[kk][local.x * 4u + c]; }
+            for (var r = 0u; r < 1u; r++) {
+                let a_value = tile_a[local.y * 1u + r][kk];
+                for (var c = 0u; c < 4u; c++) { acc[r][c] += a_value * b_values[c]; }
+            }
+        }
+        workgroupBarrier();
+    }
+    for (var r = 0u; r < 1u; r++) {
+        let out_row = tile_row + local.y * 1u + r;
+        for (var c = 0u; c < 4u; c++) {
+            let out_col = tile_n + local.x * 4u + c;
+            if (out_row < 16u && out_col < 2048u) {
+                let i = (batch * 16u + out_row) * 2048u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/104f3ed9412b8189581ee7fe358956b80adbf7cdc45a8f7770cca2737bbed4bd.wgsl

@@ -0,0 +1,12 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 2048u;
+if (i >= 2048u) { return; }
+let coord = (i / 1u) % 128u;
+if (coord >= 0u && coord < 64u) { out[i] = f32(b0[(i / 128u) * 64u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 64u && coord < 128u) { out[i] = f32(b1[(i / 128u) * 64u + (coord - 64u) * 1u + i % 1u]); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/11a92e9cba6bfc381c98c5127db137639d6c87f1cb220459880ead039dc81256.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read_write> out: array<i32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 64u;
+if (i >= 12u) { return; }
+out[i] = i32(b0[(((i / 4u) % 3u) * 1u + 1u) * 4u + ((i / 4u) % 1u) * 4u + ((i / 1u) % 4u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/125e41eee9e85ba0717b281d3b44d7ab79158d0c7038662a281fb760ffdc594f.wgsl

@@ -0,0 +1,12 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 32768u;
+if (i >= 32768u) { return; }
+let coord = (i / 1u) % 128u;
+if (coord >= 0u && coord < 64u) { out[i] = f32(b0[(i / 128u) * 64u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 64u && coord < 128u) { out[i] = f32(b1[(i / 128u) * 64u + (coord - 64u) * 1u + i % 1u]); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/13d42225cdbd1cbba4db753aa6159cce0bfeaaee44f82f9504ea29bf7dfacbbc.wgsl

@@ -0,0 +1,11 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 4096u;
+if (i >= 4096u) { return; }
+let coord = (i / 1u) % 64u;
+if (coord >= 1u && coord < 61u && (coord - 1u) % 3u == 0u) { out[i] = f32(b0[((i / 4096u) % 1u) * 1280u + ((i / 64u) % 64u) * 20u + ((((i / 1u) % 64u) - 1u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/14daf0c4c30f421761cbfd4118c412fc2946dfe95e41e8a6cf1bae767686194f.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 512u;
+if (i >= 512u) { return; }
+out[i] = f32(b0[((i / 512u) % 1u) * 1024u + ((i / 64u) % 8u) * 128u + ((i / 64u) % 1u) * 128u + (((i / 1u) % 64u) * 1u + 0u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/168567aed41eba545c80c8fcd083c67f3c0f14784912d9e76a5f7bd74021fd0b.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 16384u;
+if (i >= 16384u) { return; }
+out[i] = f32(b0[((i / 16384u) % 1u) * 32768u + ((i / 1024u) % 16u) * 2048u + ((i / 64u) % 16u) * 128u + (((i / 1u) % 64u) * 1u + 64u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/168e0daa530a7310c0ce3202ce4ddb0e26589426cc61d0a32c5958dadbeeb318.wgsl

@@ -0,0 +1,24 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+var<workgroup> factor: f32;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.x + group.y * 32u;
+    if (row >= 32u) { return; }
+    let lane = local.x;
+    if (lane == 0u) {
+        var total = 0.0;
+        for (var j = 0u; j < 128u; j++) {
+            let v = f32(b0[row * 128u + j]);
+            total += v * v;
+        }
+        factor = inverseSqrt(total / 128.0 + 1e-06);
+    }
+    workgroupBarrier();
+    for (var p = lane; p < 128u; p += 64u) {
+        let i = row * 128u + p;
+        out[i] = f32(f32(b0[row * 128u + p]) * factor * (f32(b1[p]) + 0.0));
+    }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/18a66ec0875c456c9defd45d15e6ef0c148e64a1618e2825c7f96eadbbe8303a.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read_write> out: array<i32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 64u;
+if (i >= 16u) { return; }
+out[i] = i32(b0[i]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/1aebe6301a5cb049345d5cf687b96ffb6a97e16baa18efab9ea940c4b5aaec0a.wgsl

@@ -0,0 +1,24 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+var<workgroup> factor: f32;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.x + group.y * 16u;
+    if (row >= 16u) { return; }
+    let lane = local.x;
+    if (lane == 0u) {
+        var total = 0.0;
+        for (var j = 0u; j < 2048u; j++) {
+            let v = f32(b0[row * 2048u + j]);
+            total += v * v;
+        }
+        factor = inverseSqrt(total / 2048.0 + 1e-06);
+    }
+    workgroupBarrier();
+    for (var p = lane; p < 2048u; p += 64u) {
+        let i = row * 2048u + p;
+        out[i] = f32(f32(b0[row * 2048u + p]) * factor * (f32(b1[p]) + 0.0));
+    }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/1b6b51d3ab0277a5fb7418794002fb1ab1c8791ebee7c7adb312d99b002b2897.wgsl

@@ -0,0 +1,14 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 607744u;
+if (i >= 607744u) { return; }
+let coord = (i / 1u) % 151936u;
+if (coord >= 0u && coord < 65536u) { out[i] = f32(b0[(i / 151936u) * 65536u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 65536u && coord < 131072u) { out[i] = f32(b1[(i / 151936u) * 65536u + (coord - 65536u) * 1u + i % 1u]); }
+if (coord >= 131072u && coord < 151936u) { out[i] = f32(b2[(i / 151936u) * 20864u + (coord - 131072u) * 1u + i % 1u]); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/1d104e052ac8ec4f33800dec261af1152387fb76cc709e668840f36680e642b8.wgsl

@@ -0,0 +1,24 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+var<workgroup> factor: f32;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.x + group.y * 8u;
+    if (row >= 8u) { return; }
+    let lane = local.x;
+    if (lane == 0u) {
+        var total = 0.0;
+        for (var j = 0u; j < 128u; j++) {
+            let v = f32(b0[row * 128u + j]);
+            total += v * v;
+        }
+        factor = inverseSqrt(total / 128.0 + 1e-06);
+    }
+    workgroupBarrier();
+    for (var p = lane; p < 128u; p += 64u) {
+        let i = row * 128u + p;
+        out[i] = f32(f32(b0[row * 128u + p]) * factor * (f32(b1[p]) + 0.0));
+    }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/222a2645e3e02d0a4063a63ca1b6fa4c89faee6b76905b20de4dd54ab3077ea3.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 64u;
+if (i >= 3u) { return; }
+out[i] = f32(b0[i]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/2332da1a668445655c3381ea4957b5e7b386a0f0d3340dc5ada581dcbbe85423.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 1024u;
+if (i >= 1024u) { return; }
+out[i] = f32(b0[((i / 1024u) % 1u) * 2048u + ((i / 64u) % 16u) * 128u + ((i / 64u) % 1u) * 128u + (((i / 1u) % 64u) * 1u + 0u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/23bec6a3755dde9c4ce930c33434f86d5658d7602451306b35596ab18a5c6f7f.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 8192u;
+if (i >= 8192u) { return; }
+out[i] = f32(b0[((i / 8192u) % 1u) * 8192u + ((i / 512u) % 16u) * 128u + ((i / 128u) % 4u) * 2048u + ((i / 1u) % 128u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/2505a613f10a5ed1c70a3f80a16dd566da65caa2d7d85083e23f85a30844d63e.wgsl

@@ -0,0 +1,11 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 98304u;
+if (i >= 98304u) { return; }
+let x = f32(b0[i]); let s = f32(f32(x / (1.0 + exp(-x))));
+out[i] = f32(s * f32(b1[i]));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/25532abe433506b9b9ba939343500c2c3a46b6acb963e4aefebb913ebea672e0.wgsl

@@ -0,0 +1,20 @@
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<u32>;
+@group(0) @binding(2) var<storage, read> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b1[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 64u + col / 32u];
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 131072u;
+if (i >= 131072u) { return; }
+let token = i32(b0[i / 2048u]);
+if (token < 65536 || token >= 131072) { out[i] = f32(0.0); return; }
+out[i] = f32(dequant_1(u32(token - 65536) * 2048u + i % 2048u));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/26e7750c2eb32d9a20e2e2e4094f98aab2a61d2ce78ccf72d48aeaa216c1c805.wgsl

@@ -0,0 +1,43 @@
+@group(0) @binding(0) var<storage, read> b0: array<u32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_0(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b0[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b1[row * 64u + col / 32u];
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let slot = local.x / 16u;
+    let lane = local.x % 16u;
+    let i = (group.x + group.y * 1536u) * 4u + slot;
+    var acc = 0.0;
+    if (i < 6144u) {
+        let row = i % 6144u;
+        for (var w = lane; w < 512u; w += 16u) {
+            let word = b0[row * 512u + w];
+            let first = w * 4u;
+            var part = 0.0;
+            for (var e = 0u; e < 4u; e++) {
+                let col = first + e;
+                if (col < 2048u) {
+                    let code = i32((word >> (e * 8u)) & 255u) - 128;
+                    part += f32(code) * f32(b2[(i / 6144u) * 2048u + col]);
+                }
+            }
+            acc += part * f32(b1[row * 64u + first / 32u]);
+        }
+    }
+    partial[local.x] = acc;
+    workgroupBarrier();
+    for (var stride = 8u; stride > 0u; stride /= 2u) {
+        if (lane < stride) { partial[local.x] += partial[local.x + stride]; }
+        workgroupBarrier();
+    }
+    if (lane == 0u && i < 6144u) { out[i] = f32(partial[local.x] + 0.0); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/276e9e1e05c3835d7a10dd0e79be4657ff9e812111751e30425c6dc609f1b79e.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 1280u;
+if (i >= 1280u) { return; }
+out[i] = f32(b0[((i / 1280u) % 1u) * 4096u + ((i / 20u) % 64u) * 64u + (((i / 1u) % 20u) * 3u + 1u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/28b69aba1aba878089e1293d59fce422e76487f62cda6ef45c4a1f25cb733f3b.wgsl

@@ -0,0 +1,10 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 65536u;
+if (i >= 65536u) { return; }
+out[i] = f32(f32(b0[i]) + (f32(b1[i]) * 1.0));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/28f11beb98a242e2cb799b2756b47748a81738ff2d0a62a46c2a24427d6af843.wgsl

@@ -0,0 +1,10 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 8192u;
+if (i >= 8192u) { return; }
+out[i] = f32(f32(b0[i]) * f32(b1[((i / 128u) % 4u) * 128u + ((i / 1u) % 128u) * 1u]));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/2b02ece9784fc350cef18e08c962031f74dbc27c7d1cad3a13dda1cebb84b403.wgsl

@@ -0,0 +1,11 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 393216u;
+if (i >= 393216u) { return; }
+let x = f32(b0[i]); let s = f32(f32(x / (1.0 + exp(-x))));
+out[i] = f32(s * f32(b1[i]));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/2bdeb3590b1ba05cf6e9c7e8245f8aa552a12dae84f9578c603c9cee081a82d4.wgsl

@@ -0,0 +1,24 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+var<workgroup> factor: f32;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.x + group.y * 16u;
+    if (row >= 16u) { return; }
+    let lane = local.x;
+    if (lane == 0u) {
+        var total = 0.0;
+        for (var j = 0u; j < 128u; j++) {
+            let v = f32(b0[row * 128u + j]);
+            total += v * v;
+        }
+        factor = inverseSqrt(total / 128.0 + 1e-06);
+    }
+    workgroupBarrier();
+    for (var p = lane; p < 128u; p += 64u) {
+        let i = row * 128u + p;
+        out[i] = f32(f32(b0[row * 128u + p]) * factor * (f32(b1[p]) + 0.0));
+    }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/2ce5bb171c1116e8bbab4652f5bbbbec1423238169a611533f0cda5f4e6e23bb.wgsl

@@ -0,0 +1,10 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 32768u;
+if (i >= 32768u) { return; }
+out[i] = f32(f32(b0[i]) * f32(b1[((i / 128u) % 16u) * 128u + ((i / 1u) % 128u) * 1u]));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/3203e08f2c99fa4848f553739b7115a7c677840344f17b1d4391d6130517f5a7.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read_write> out: array<i32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 64u;
+if (i >= 4u) { return; }
+out[i] = i32(b0[i]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/33245a8fb16de10d97ccedb17a9b479b3b7024b2b88a26c738a95f94f248fa8d.wgsl

@@ -0,0 +1,43 @@
+@group(0) @binding(0) var<storage, read> b0: array<u32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_0(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b0[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b1[row * 64u + col / 32u];
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let slot = local.x / 16u;
+    let lane = local.x % 16u;
+    let i = (group.x + group.y * 512u) * 4u + slot;
+    var acc = 0.0;
+    if (i < 2048u) {
+        let row = i % 2048u;
+        for (var w = lane; w < 512u; w += 16u) {
+            let word = b0[row * 512u + w];
+            let first = w * 4u;
+            var part = 0.0;
+            for (var e = 0u; e < 4u; e++) {
+                let col = first + e;
+                if (col < 2048u) {
+                    let code = i32((word >> (e * 8u)) & 255u) - 128;
+                    part += f32(code) * f32(b2[(i / 2048u) * 2048u + col]);
+                }
+            }
+            acc += part * f32(b1[row * 64u + first / 32u]);
+        }
+    }
+    partial[local.x] = acc;
+    workgroupBarrier();
+    for (var stride = 8u; stride > 0u; stride /= 2u) {
+        if (lane < stride) { partial[local.x] += partial[local.x + stride]; }
+        workgroupBarrier();
+    }
+    if (lane == 0u && i < 2048u) { out[i] = f32(partial[local.x] + 0.0); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/34f69d1969e86c8f4c8d1781bbc4151e6636a829db146395f83ed0d3b08032be.wgsl

@@ -0,0 +1,11 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 1024u;
+if (i >= 1024u) { return; }
+let coord = (i / 1u) % 64u;
+if (coord >= 1u && coord < 61u && (coord - 1u) % 3u == 0u) { out[i] = f32(b0[((i / 1024u) % 1u) * 320u + ((i / 64u) % 16u) * 20u + ((((i / 1u) % 64u) - 1u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/362a01a15696bbdeba794896dfd844b25f0400b1d51b510ced756acd0845aa23.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 2048u;
+if (i >= 2048u) { return; }
+out[i] = f32(b0[((i / 2048u) % 1u) * 4096u + ((i / 256u) % 8u) * 512u + ((i / 64u) % 4u) * 128u + (((i / 1u) % 64u) * 1u + 0u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/3669082b9a10903156fc5a35ec21f6dd470e11be7cc903372a6c2d410ed3287a.wgsl

@@ -0,0 +1,43 @@
+@group(0) @binding(0) var<storage, read> b0: array<u32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_0(index: u32) -> f32 {
+    let row = index / 6144u;
+    let col = index % 6144u;
+    let word = b0[row * 1536u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b1[row * 192u + col / 32u];
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let slot = local.x / 16u;
+    let lane = local.x % 16u;
+    let i = (group.x + group.y * 512u) * 4u + slot;
+    var acc = 0.0;
+    if (i < 2048u) {
+        let row = i % 2048u;
+        for (var w = lane; w < 1536u; w += 16u) {
+            let word = b0[row * 1536u + w];
+            let first = w * 4u;
+            var part = 0.0;
+            for (var e = 0u; e < 4u; e++) {
+                let col = first + e;
+                if (col < 6144u) {
+                    let code = i32((word >> (e * 8u)) & 255u) - 128;
+                    part += f32(code) * f32(b2[(i / 2048u) * 6144u + col]);
+                }
+            }
+            acc += part * f32(b1[row * 192u + first / 32u]);
+        }
+    }
+    partial[local.x] = acc;
+    workgroupBarrier();
+    for (var stride = 8u; stride > 0u; stride /= 2u) {
+        if (lane < stride) { partial[local.x] += partial[local.x + stride]; }
+        workgroupBarrier();
+    }
+    if (lane == 0u && i < 2048u) { out[i] = f32(partial[local.x] + 0.0); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/3852e142e2cc2bdb19ade848500718b06f307b4abeeeb7cdb9057c1935c9bde7.wgsl

@@ -0,0 +1,12 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 1024u;
+if (i >= 1024u) { return; }
+let coord = (i / 1u) % 128u;
+if (coord >= 0u && coord < 64u) { out[i] = f32(b0[(i / 128u) * 64u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 64u && coord < 128u) { out[i] = f32(b1[(i / 128u) * 64u + (coord - 64u) * 1u + i % 1u]); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/38e27417786aaea5ad0e1e556c832f4e27fb341aac1f673e05ab3fed7d3d45ff.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 64u;
+if (i >= 20u) { return; }
+out[i] = f32(b0[i]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/3b797f20f83268b3aed99f37180b0ed472ece9567c9a2b3802b34c84dad22b4d.wgsl

@@ -0,0 +1,14 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 2430976u;
+if (i >= 2430976u) { return; }
+let coord = (i / 1u) % 151936u;
+if (coord >= 0u && coord < 65536u) { out[i] = f32(b0[(i / 151936u) * 65536u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 65536u && coord < 131072u) { out[i] = f32(b1[(i / 151936u) * 65536u + (coord - 65536u) * 1u + i % 1u]); }
+if (coord >= 131072u && coord < 151936u) { out[i] = f32(b2[(i / 151936u) * 20864u + (coord - 131072u) * 1u + i % 1u]); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/3caf74924884e14997a7fcbb2625a36e3b8a2bcf92da835bf6f787ddae437cdb.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read_write> out: array<i32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 64u;
+if (i >= 64u) { return; }
+out[i] = i32(b0[((i / 1u) % 16u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/3cdd7e7cfd35a1c5035d9d6af6d244348fedcfd7fdd4863953b1a29c131b914f.wgsl

@@ -0,0 +1,12 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 4096u;
+if (i >= 4096u) { return; }
+let coord = (i / 1u) % 128u;
+if (coord >= 0u && coord < 64u) { out[i] = f32(b0[(i / 128u) * 64u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 64u && coord < 128u) { out[i] = f32(b1[(i / 128u) * 64u + (coord - 64u) * 1u + i % 1u]); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/3d5b97cfaf8dd0e680b6a9c84e7eca497f9f6e49817d20f1a72dddd7a3303f73.wgsl

@@ -0,0 +1,11 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 24576u;
+if (i >= 24576u) { return; }
+let x = f32(b0[i]); let s = f32(f32(x / (1.0 + exp(-x))));
+out[i] = f32(s * f32(b1[i]));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/3e34409407100d4858c4c109cf62f3eaf2f4c6081166b0ec8679105e6dfbe7fe.wgsl

@@ -0,0 +1,10 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 16384u;
+if (i >= 16384u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(-x);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/41395d069ce32ecf4ddc74fddd167040742fc8b1a6bcd34e735b58c8feac1030.wgsl

@@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 131072u;
+if (i >= 131072u) { return; }
+out[i] = f32(b0[((i / 131072u) % 1u) * 131072u + ((i / 8192u) % 16u) * 128u + ((i / 128u) % 64u) * 2048u + ((i / 1u) % 128u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2-v5/kernels/41e68053bc9465358f12a90e0efa6f471e5bb8ac763f9e9ec12e3141fb99b68e.wgsl

@@ -0,0 +1,11 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 128u;
+if (i >= 128u) { return; }
+let coord = (i / 1u) % 128u;
+if (coord >= 0u && coord < 64u) { out[i] = f32(b0[(i / 128u) * 64u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 64u && coord < 128u) { out[i] = f32(b0[(i / 128u) * 64u + (coord - 64u) * 1u + i % 1u]); }
+}