Add qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5 (attention bounded by the used context) (part 2)

.gitattributes

@@ -75,3 +75,4 @@ qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2-v4/graph.json fi
 qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2-v4/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
 qwen35-08b-fp32-8k-token-major-home-v5/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
 qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/graph.json filter=lfs diff=lfs merge=lfs -text
+qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/a77f14c89662d2d3da8abd8f109667793dc662d307e5e1c61123be753e25437b.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 * 128u;
+if (i >= 96u) { return; }
+out[i] = f32(b0[((i / 1u) % 32u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/a7b50e8b6e4d56404b6261383f50b3eed4a2a59c3fadaaffcc974575c32a374f.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 * 98304u;
+if (i >= 98304u) { return; }
+out[i] = f32(b0[((i / 98304u) % 1u) * 98304u + ((i / 6144u) % 16u) * 1u + ((i / 1u) % 6144u) * 16u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/a87152b0fd2f78f60d7abb61c2f6c4a9bde8a9dd45d0275b58d760a3da9f71ce.wgsl

@@ -0,0 +1,17 @@
+@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 * 384u;
+if (i >= 384u) { return; }
+let batch = i / 128u;
+let row = (i / 4u) % 32u;
+let col = i % 4u;
+var acc = 0.0;
+for (var p = 0u; p < 1u; p++) {
+    acc += f32(b0[(((batch / 1u) % 3u) * 32u) + row * 1u + p]) * f32(b1[(((batch / 1u) % 3u) * 4u) + p * 4u + col]);
+}
+out[i] = f32(acc);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/a896c93b04922c0f78f1c77fbaabbc79eb71776223748bb32a6d5339e32029e8.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 * 256u;
+if (i >= 256u) { return; }
+out[i] = i32(b0[((i / 1u) % 64u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/a9f0eef095f015e2099adf81552a6727c45e70fc11f5e50b777d36b6370ab986.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) * 24576u + ((i / 2048u) % 4u) * 6144u + (((i / 1u) % 2048u) * 1u + 2048u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/a9f6bc361f1ea31f79f0ef6164872eeb04896a45338871a2d779789b47b1920b.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 * 131072u;
+if (i >= 131072u) { return; }
+out[i] = f32(f32(b0[i]) * f32(b1[i]));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/a9fa5db20fa8de1de429c6850c3e40ff8180bcdb59373919b84d2438b3b17352.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 >= 48u) { return; }
+out[i] = i32(b0[(((i / 16u) % 3u) * 1u + 1u) * 16u + ((i / 16u) % 1u) * 16u + ((i / 1u) % 16u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/aa227ae49fab3f75d86e02e6f390be785eabfeafe70e77a5c7eb4ad51895e382.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]));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/abdb17daf29310a7828a55a1efeebaa65b6553e236fc20702e8c712e0a46dd5b.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 * 2048u;
+if (i >= 2048u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(1.0 / (1.0 + exp(-x)));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/ac0f1760c26da88505cd869eb44fa883b9fc8100196da2be6f68196db30794f5.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 / 1024u;
+    let col = index % 1024u;
+    let word = b0[row * 256u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b1[row * 32u + 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 * 4u) * 4u + slot;
+    var acc = 0.0;
+    if (i < 16u) {
+        let row = i % 16u;
+        for (var w = lane; w < 256u; w += 16u) {
+            let word = b0[row * 256u + w];
+            let first = w * 4u;
+            var part = 0.0;
+            for (var e = 0u; e < 4u; e++) {
+                let col = first + e;
+                if (col < 1024u) {
+                    let code = i32((word >> (e * 8u)) & 255u) - 128;
+                    part += f32(code) * f32(b2[(i / 16u) * 1024u + col]);
+                }
+            }
+            acc += part * f32(b1[row * 32u + 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 < 16u) { out[i] = f32(partial[local.x] + 0.0); }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/ac173d038b5767711ad9f7cafbceb9ca3ec84aaeb3d07a3037e7cd0f83fead2c.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) * 393216u + ((i / 2048u) % 64u) * 6144u + (((i / 1u) % 2048u) * 1u + 0u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/ad72f60d58133272228319f1c066e7cd43ac28573ddeeb7a6f69be01af341f57.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) * 2048u + ((i / 256u) % 2u) * 1024u + ((i / 64u) % 4u) * 256u + (((i / 1u) % 64u) * 1u + 0u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/adf1e58a124e1301625024ae142dca5fb1dd50799fa94b52eaeffb82ea99852c.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 * 4096u;
+if (i >= 4096u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(-x);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/adfbba60ade7143599d7de0faa4904afc1400ccc3cf84602f22a2fb004d0398e.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 / 1024u;
+    let col = index % 1024u;
+    let word = b1[row * 256u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 32u + 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 < 1024u; 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 < 4u && col < 1024u) { value = f32(f32(b0[(batch * 4u + row) * 1024u + 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 < 131072u && first < 1024u) {
+                word = b1[n_index * 256u + first / 4u];
+                scale = f32(b2[n_index * 32u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 1024u) { 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 < 4u && out_col < 131072u) {
+                let i = (batch * 4u + out_row) * 131072u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/ae69e698c92571766e36ad8c078b13e42bde56745d45698daaadaca587b7279c.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 * 128u;
+if (i >= 128u) { return; }
+let coord = (i / 1u) % 32u;
+if (coord >= 2u && coord < 32u && (coord - 2u) % 3u == 0u) { out[i] = f32(b0[((i / 128u) % 1u) * 40u + ((i / 32u) % 4u) * 10u + ((((i / 1u) % 32u) - 2u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/b02b0be460d581c720afd68c8b722e8da47e2ee432b5d61bbdf24b6084cb2f26.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 * 2048u;
+if (i >= 2048u) { return; }
+let x = f32(b0[i]); let s = f32(f32(x / (1.0 + exp(-x))));
+out[i] = f32(s * f32(b1[i]));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/b12eae76f13225c0c5d35fc9309964aba3172712799ff3bd8538bb2c1b34aa33.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 * 1536u;
+if (i >= 1536u) { return; }
+out[i] = f32(b0[((i / 512u) % 3u) * 512u + ((i / 512u) % 1u) * 512u + ((i / 32u) % 16u) * 1u + ((i / 1u) % 32u) * 16u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/b199d835db70f888c7b99a3ac6e9ec56f6c47d5bba4eaf352e0a8eabed56411e.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 / 128u) % 16u) * 1u]));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/b2ae1b161c8a6315f8863e47511607027dfe583aac1db4c1626371ddfa0c64e2.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 * 256u;
+if (i >= 256u) { return; }
+out[i] = f32(b0[((i / 256u) % 1u) * 512u + ((i / 128u) % 2u) * 256u + ((i / 32u) % 4u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/b2b16798cbe94c919025d76b0435afe47966ffdd1ea53ec74a791c03c4709a28.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 * 30720u;
+if (i >= 30720u) { return; }
+out[i] = f32(b0[((i / 30720u) % 1u) * 49152u + ((i / 5u) % 6144u) * 8u + (((i / 1u) % 5u) * 1u + 0u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/b3bb2905555188f535b6b8bc1d050c142b5fea8222d9ef193466f80aa2c54dea.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 * 2048u;
+if (i >= 2048u) { return; }
+let coord = (i / 1u) % 32u;
+if (coord >= 2u && coord < 32u && (coord - 2u) % 3u == 0u) { out[i] = f32(b0[((i / 2048u) % 1u) * 640u + ((i / 32u) % 64u) * 10u + ((((i / 1u) % 32u) - 2u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/b57928140994202536cd4f99d2e14f5dba9754f014843673790d5d48c9cbb9ab.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) * 512u + ((i / 256u) % 2u) * 256u + ((i / 256u) % 1u) * 512u + ((i / 1u) % 256u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/b5984ff6eb9e56ef0f9930f775c1dc3c203eec15924e5fea4c2525b2ce180857.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) * 8192u + ((i / 512u) % 8u) * 1024u + ((i / 32u) % 16u) * 64u + (((i / 1u) % 32u) * 1u + 0u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/b6aefeed0ff00bd9d8e09ef5cee0a3e3883f7e2ad15a16918acc82d773b89fba.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 / 3584u;
+    let col = index % 3584u;
+    let word = b0[row * 448u + col / 8u];
+    let code = i32((word >> ((col % 8u) * 4u)) & 15u) - 8;
+    return f32(code) * b1[row * 112u + 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 * 256u) * 4u + slot;
+    var acc = 0.0;
+    if (i < 1024u) {
+        let row = i % 1024u;
+        for (var w = lane; w < 448u; w += 16u) {
+            let word = b0[row * 448u + w];
+            let first = w * 8u;
+            var part = 0.0;
+            for (var e = 0u; e < 8u; e++) {
+                let col = first + e;
+                if (col < 3584u) {
+                    let code = i32((word >> (e * 4u)) & 15u) - 8;
+                    part += f32(code) * f32(b2[(i / 1024u) * 3584u + col]);
+                }
+            }
+            acc += part * f32(b1[row * 112u + 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 < 1024u) { out[i] = f32(partial[local.x] + 0.0); }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/b6c3ce0defd66eef66dc3fd6eb9eab3fee12d179332698c3c764009988628b49.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 * 49152u;
+if (i >= 49152u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(x / (1.0 + exp(-x)));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/bb3e27e020bc22a0c8c21a5f9910a0a26b73ace5f68d9d49d1e4d67d53232060.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 < 4u && col < 2048u) { value = f32(f32(b0[(batch * 4u + 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 < 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 < 4u && out_col < 1024u) {
+                let i = (batch * 4u + out_row) * 1024u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/bb7f03e6a2277dc56084e9130086507be58871253867c0ce8ad60ee5ce4521b6.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 * 262144u;
+if (i >= 262144u) { return; }
+out[i] = f32(b0[((i / 262144u) % 1u) * 393216u + ((i / 16384u) % 16u) * 24576u + (((i / 128u) % 128u) * 1u + 0u) * 128u + ((i / 1u) % 128u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/bb9029b667eeb13670d1f640cd677c6eabd7a7569686c599e3eeb9482d642b00.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));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/bcb78b92dbbfeb212c190e064bc2f9d8e79c201398cf38c9aa62df6fdba94469.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 * 32768u;
+if (i >= 32768u) { return; }
+out[i] = f32(b0[((i / 32768u) % 1u) * 32768u + ((i / 16384u) % 2u) * 256u + ((i / 256u) % 64u) * 512u + ((i / 1u) % 256u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/bd70e4fc42ac55f0f52612841bcc3f9fe7df5bbb0fc76f94e56a131a92454658.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 * 24576u;
+if (i >= 24576u) { return; }
+out[i] = f32(b0[((i / 24576u) % 1u) * 32768u + ((i / 12288u) % 2u) * 16384u + ((i / 192u) % 64u) * 256u + (((i / 1u) % 192u) * 1u + 64u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/bd8a785d4fbf09918f149c7eb5f34bd25f9358bf3412e27799a7c65436c71759.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 * 24576u;
+if (i >= 24576u) { return; }
+out[i] = f32(b0[((i / 24576u) % 1u) * 417792u + ((i / 4u) % 6144u) * 68u + (((i / 1u) % 4u) * 1u + 64u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/bda7ebd3bcf25c31938fca1a582e82b972ab742f85f826f085417b0e6f651e60.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) % 64u;
+if (coord >= 0u && coord < 32u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 32u && coord < 64u) { out[i] = f32(b1[(i / 64u) * 32u + (coord - 32u) * 1u + i % 1u]); }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c02ac02d110c0be5d0990d20ab0d8c7a081c74f4e147f083789a203f22fc3258.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 * 256u;
+if (i >= 256u) { return; }
+out[i] = f32(f32(b0[((i / 1u) % 16u) * 1u]) * f32(b1[i]));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c0b2947798baa198598d75ca0e5dcd31a9fa320499665e34f7bb81d27fc15fb0.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 * 64u;
+if (i >= 64u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(cos(x));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c1096000361a304457865401d6266a640469ca04c0bbfe2eaa657e656b2fca33.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 * 64u;
+if (i >= 64u) { return; }
+var acc = 0.0;
+for (var j = 0u; j < 128u; j++) { acc += f32(b0[(((i / 64u) % 1u) * 8192u + ((i / 16u) % 4u) * 2048u + ((i / 1u) % 16u) * 128u) + (((j / 1u) % 128u) * 1u)]); }
+out[i] = f32(acc);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c23bee3c3664172bf0a078a96acd2b3ad3136732234975f757d313e696c2c37a.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 * 64u;
+if (i >= 16u) { return; }
+out[i] = f32(f32(b0[((i / 1u) % 16u) * 1u]) * f32(b1[i]));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c24c7e6057a3ab4778def9751ec5123025ec0e36139db626eccc4a9151bc6cc3.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 * 32768u;
+if (i >= 32768u) { return; }
+out[i] = f32(b0[((i / 32768u) % 1u) * 98304u + ((i / 2048u) % 16u) * 6144u + (((i / 1u) % 2048u) * 1u + 0u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c2c806bbe93f09247867b4ace45489a5f2fe3ee67f265418eb7aa71a29880276.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 * 64u;
+if (i >= 16u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(inverseSqrt(x));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c2e7cdc03edb25ef7a160965c17c559b6ec7e139c145dc7390b73523bacba5bb.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 * 192u;
+if (i >= 160u) { return; }
+out[i] = f32(b0[i]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c342e2d29647509fe0f6ac1e2dbbe74032dbe9ecd6229c03557b86f160381d2e.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 * 1024u;
+if (i >= 1024u) { return; }
+out[i] = f32(f32(b0[((i / 1u) % 16u) * 1u]) * f32(b1[i]));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c3735f34cafbac42efa5928641cefdaf7277659cc2675a36756bba97240ba5ae.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 / 2048u) % 4u) * 16u + ((i / 128u) % 16u) * 1u]));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c3b8a95df30cf49f5ba2a9aa459e227338c340989ea076b9cf7fc3d09ba103a4.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 * 1024u;
+    if (row >= 1024u) { 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));
+    }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c4862be63d2532b6891da039edbef2f46737a34518f8a89602821e00447b5933.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) * 6144u + ((i / 2048u) % 1u) * 6144u + (((i / 1u) % 2048u) * 1u + 2048u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c5a289d1fb6564db68cc05de03941dff72def270e328ba8fb520e5dd0efb075f.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 / 1024u;
+    let col = index % 1024u;
+    let word = b1[row * 128u + col / 8u];
+    let code = i32((word >> ((col % 8u) * 4u)) & 15u) - 8;
+    return f32(code) * b2[row * 32u + 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 < 1024u; 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 < 1024u) { value = f32(f32(b0[(batch * 64u + row) * 1024u + col])); }
+            tile_a[m_local][flat % 16u] = value;
+        }
+        if (lane < 128u) {
+            let n_local = lane / 2u;
+            let word_local = lane % 2u;
+            let n_index = tile_n + n_local;
+            let first = k0 + word_local * 8u;
+            var word = 0u;
+            var scale = 0.0;
+            if (n_index < 3584u && first < 1024u) {
+                word = b1[n_index * 128u + first / 8u];
+                scale = f32(b2[n_index * 32u + first / 32u]);
+            }
+            for (var e = 0u; e < 8u; e++) {
+                let code = i32((word >> (e * 4u)) & 15u) - 8;
+                var value = 0.0;
+                if (first + e < 1024u) { value = f32(code) * scale; }
+                tile_b[word_local * 8u + 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 < 3584u) {
+                let i = (batch * 64u + out_row) * 3584u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c5b7564d87c89606da1cba30e05e401e4fb8c83957e779a694d2ec2c472d88ad.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 / 1024u;
+    let col = index % 1024u;
+    let word = b1[row * 256u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 32u + 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 < 1024u; 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 < 1024u) { value = f32(f32(b0[(batch * 64u + row) * 1024u + 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 < 512u && first < 1024u) {
+                word = b1[n_index * 256u + first / 4u];
+                scale = f32(b2[n_index * 32u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 1024u) { 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 < 512u) {
+                let i = (batch * 64u + out_row) * 512u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c7c966b2f9c54a16d0078728edc1f2a7f79265fed388a815a0d52bcd73d1e0f6.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 / 1024u;
+    let col = index % 1024u;
+    let word = b1[row * 256u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 32u + 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 < 1024u; 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 < 1024u) { value = f32(f32(b0[(batch * 16u + row) * 1024u + 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 < 131072u && first < 1024u) {
+                word = b1[n_index * 256u + first / 4u];
+                scale = f32(b2[n_index * 32u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 1024u) { 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 < 131072u) {
+                let i = (batch * 16u + out_row) * 131072u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c83a2a6faf563ec3bbad43a3e108d58c09cd8aeed01b38e4e82d1c3c9b650c21.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[2u * 2048u + ((i / 2048u) % 1u) * 2048u + ((i / 32u) % 64u) * 32u + ((i / 1u) % 32u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c90428425f6d220b441e97e3b894adfb090869c98b17f53afbd2f9fb202b4525.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 * 64u;
+if (i >= 64u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(sin(x));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/c93686edb422c5385c7a6f484e76be0aa5cfa478cc80bc900a707e19fa2de99c.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 * 4096u;
+if (i >= 4096u) { return; }
+out[i] = f32(f32(b0[i]) + (f32(b1[i]) * 1.0));
+}