Add qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2 WebTorch bundle (part 2)

.gitattributes

@@ -54,3 +54,4 @@ qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/tokenizer/tokenizer.jso
 qwen35-2b-fp32-int8-g32-exclude-l6-l9-l10up-l15o-home-token-major-v2/graph.json filter=lfs diff=lfs merge=lfs -text
 qwen35-2b-fp32-int8-g32-exclude-l6-l9-l10up-l15o-home-token-major-v2/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
 qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2/graph.json filter=lfs diff=lfs merge=lfs -text
+qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text

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

@@ -0,0 +1,13 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<i32>;
+@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 token = (i / 256u) % 4u;
+    let outer = i / 1024u;
+    let destination = outer * 1048576u + u32(b1[token]) * 256u + i % 256u;
+    out[((destination) / 256u) % 4096u * 512u + ((destination) / 1048576u) % 2u * 256u + ((destination) / 1u) % 256u * 1u] = f32(b0[i]);
+}

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

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

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

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

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

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

@@ -0,0 +1,31 @@
+@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<f32, 64>;
+var<workgroup> tile_b: array<f32, 64>;
+@compute @workgroup_size(8, 8)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.y * 8u + local.y;
+    let col = group.x * 8u + local.x;
+    let i = (group.z * 64u + row) * 16u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 1024u; tile += 8u) {
+        tile_a[local.y * 8u + local.x] = 0.0;
+        tile_b[local.x * 8u + local.y] = 0.0;
+        if (row < 64u && tile + local.x < 1024u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 64u + row) * 1024u + tile + local.x]); }
+        if (group.x * 8u + local.y < 16u && tile + local.x < 1024u) { tile_b[local.x * 8u + local.y] = dequant_1((group.x * 8u + local.y) * 1024u + tile + local.x); }
+        workgroupBarrier();
+        for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
+        workgroupBarrier();
+    }
+    if (row < 64u && col < 16u) { out[i] = f32(acc + 0.0); }
+}

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

@@ -0,0 +1,28 @@
+@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_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+var<workgroup> tile_a: array<f32, 64>;
+var<workgroup> tile_b: array<f32, 64>;
+@compute @workgroup_size(8, 8)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.y * 8u + local.y;
+    let col = group.x * 8u + local.x;
+    let i = (group.z * 64u + row) * 65536u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 1024u; tile += 8u) {
+        tile_a[local.y * 8u + local.x] = 0.0;
+        tile_b[local.x * 8u + local.y] = 0.0;
+        if (row < 64u && tile + local.x < 1024u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 64u + row) * 1024u + tile + local.x]); }
+        if (group.x * 8u + local.y < 65536u && tile + local.x < 1024u) { tile_b[local.x * 8u + local.y] = unpack_bf16_1((group.x * 8u + local.y) * 1024u + tile + local.x); }
+        workgroupBarrier();
+        for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
+        workgroupBarrier();
+    }
+    if (row < 64u && col < 65536u) { out[i] = f32(acc + 0.0); }
+}

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

@@ -0,0 +1,28 @@
+@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_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+var<workgroup> tile_a: array<f32, 64>;
+var<workgroup> tile_b: array<f32, 64>;
+@compute @workgroup_size(8, 8)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.y * 8u + local.y;
+    let col = group.x * 8u + local.x;
+    let i = (group.z * 16u + row) * 51712u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 1024u; tile += 8u) {
+        tile_a[local.y * 8u + local.x] = 0.0;
+        tile_b[local.x * 8u + local.y] = 0.0;
+        if (row < 16u && tile + local.x < 1024u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 16u + row) * 1024u + tile + local.x]); }
+        if (group.x * 8u + local.y < 51712u && tile + local.x < 1024u) { tile_b[local.x * 8u + local.y] = unpack_bf16_1((group.x * 8u + local.y) * 1024u + tile + local.x); }
+        workgroupBarrier();
+        for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
+        workgroupBarrier();
+    }
+    if (row < 16u && col < 51712u) { out[i] = f32(acc + 0.0); }
+}

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

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

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

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

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

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

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2/kernels/9b7697d3e1de37277befbd725307f40a5143faae3dcf43182d4ad89d8572ed42.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(f32(b0[i]) * 1.0);
+}

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

@@ -0,0 +1,17 @@
+@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_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+@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 token = i32(b0[i / 1024u]);
+if (token < 131072 || token >= 196608) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 131072) * 1024u + i % 1024u));
+}

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

@@ -0,0 +1,17 @@
+@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_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+@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 token = i32(b0[i / 1024u]);
+if (token < 65536 || token >= 131072) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 65536) * 1024u + i % 1024u));
+}

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

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

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2/kernels/9d3d7b777701cd320ca11358e554e1019fb041ee4d765e66071c215242906773.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) * 65536u + ((i / 2048u) % 16u) * 4096u + ((i / 256u) % 8u) * 512u + (((i / 1u) % 256u) * 1u + 0u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2/kernels/9d504391a987c7be7ccfbb7972c27cf54cd62730c19cbe7949fee2ea059f0c43.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 * 122880u;
+if (i >= 122880u) { 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/kernels/9f02cd7a8eb5e6d20c2b01559fc58db5b0447ae41a51e2cd479ed519fc284791.wgsl

@@ -0,0 +1,16 @@
+@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> b3: array<f32>;
+@group(0) @binding(4) 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 * 248320u;
+if (i >= 248320u) { return; }
+let coord = (i / 1u) % 248320u;
+if (coord >= 0u && coord < 65536u) { out[i] = f32(b0[(i / 248320u) * 65536u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 65536u && coord < 131072u) { out[i] = f32(b1[(i / 248320u) * 65536u + (coord - 65536u) * 1u + i % 1u]); }
+if (coord >= 131072u && coord < 196608u) { out[i] = f32(b2[(i / 248320u) * 65536u + (coord - 131072u) * 1u + i % 1u]); }
+if (coord >= 196608u && coord < 248320u) { out[i] = f32(b3[(i / 248320u) * 51712u + (coord - 196608u) * 1u + i % 1u]); }
+}

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

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

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

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

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

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

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

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

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

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2/kernels/a53c1ce51ebaec593ad667efae9603efe3807db8f0ee060a620ac031cc4f2b92.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 * 8192u;
+if (i >= 8192u) { 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/kernels/a62f6b62277fe048d6bac44a25a2e61399e7e2eab103fccd3a23c29a3c827856.wgsl

@@ -0,0 +1,44 @@
+@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> b3: array<f32>;
+@group(0) @binding(4) var<storage, read> b4: array<f32>;
+@group(0) @binding(5) var<storage, read> b5: array<f32>;
+@group(0) @binding(6) var<storage, read_write> out: array<f32>;
+
+var<workgroup> partial: array<f32, 128>;
+var<workgroup> delta: f32;
+@compute @workgroup_size(128)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let column = group.x;
+    let head = group.y;
+    let batch_index = group.z;
+    let lane = local.x;
+    var recurrent = 0.0;
+    if (lane < 128u) { recurrent = f32(b5[((batch_index * 16u + head) * 128u + lane) * 128u + column]); }
+    for (var token = 0u; token < 64u; token++) {
+        recurrent *= exp(f32(b3[(batch_index * 64u + token) * 16u + head]));
+        partial[lane] = 0.0;
+        if (lane < 128u) { partial[lane] = recurrent * f32(b1[((batch_index * 64u + token) * 16u + head) * 128u + lane]); }
+        workgroupBarrier();
+        for (var stride = 64u; stride > 0u; stride /= 2u) {
+            if (lane < stride) { partial[lane] += partial[lane + stride]; }
+            workgroupBarrier();
+        }
+        if (lane == 0u) { delta = (f32(b2[((batch_index * 64u + token) * 16u + head) * 128u + column]) - partial[0]) * f32(b4[(batch_index * 64u + token) * 16u + head]); }
+        workgroupBarrier();
+        if (lane < 128u) { recurrent += f32(b1[((batch_index * 64u + token) * 16u + head) * 128u + lane]) * delta; }
+        partial[lane] = 0.0;
+        if (lane < 128u) {
+            partial[lane] = recurrent * f32(b0[((batch_index * 64u + token) * 16u + head) * 128u + lane]) * 0.08838834764831845;
+        }
+        workgroupBarrier();
+        for (var stride = 64u; stride > 0u; stride /= 2u) {
+            if (lane < stride) { partial[lane] += partial[lane + stride]; }
+            workgroupBarrier();
+        }
+        if (lane == 0u) { out[((batch_index * 16u + head) * 192u + 128u + token) * 128u + column] = f32(partial[0]); }
+        workgroupBarrier();
+    }
+    if (lane < 128u) { out[((batch_index * 16u + head) * 192u + lane) * 128u + column] = f32(recurrent); }
+}

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

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

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

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2/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/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/kernels/ac61ff8b1e3d15ef1f43e9d610e97d84fcae08527ed0a992d3cebbce5fdd286a.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 * 32768u;
+if (i >= 32768u) { return; }
+let batch = i / 4096u;
+let row = (i / 256u) % 16u;
+let col = i % 256u;
+var acc = 0.0;
+for (var p = 0u; p < 4096u; p++) {
+    acc += f32(b0[(((batch / 1u) % 8u) * 65536u) + row * 4096u + p]) * f32(b1[(((batch / 1u) % 8u) * 1048576u) + p * 256u + col]);
+}
+out[i] = f32(acc);
+}

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