Publish WebTorch model catalog bundles (part 7)

.gitattributes

@@ -38,3 +38,5 @@ qwen35-08b-fp32-8k-token-major-home/tokenizer/tokenizer.json filter=lfs diff=lfs
 qwen35-08b-fp32-int8-g32-home-token-major-v2/graph.json filter=lfs diff=lfs merge=lfs -text
 qwen35-08b-fp32-int8-g32-home-token-major-v2/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
 qwen35-2b-fp32-int8-g32-home-token-major-v2/graph.json filter=lfs diff=lfs merge=lfs -text
+qwen35-2b-fp32-int8-g32-home-token-major-v2/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
+qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/graph.json filter=lfs diff=lfs merge=lfs -text

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/4d46494a55ac0b66fdddd47b9422106ac39e0c7d206f01004e13f93a010e1bf5.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 * 49152u;
+if (i >= 49152u) { return; }
+let coord = (i / 1u) % 8u;
+if (coord >= 0u && coord < 4u) { out[i] = f32(b0[(i / 8u) * 4u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 4u && coord < 8u) { out[i] = f32(b1[(i / 8u) * 4u + (coord - 4u) * 1u + i % 1u]); }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/4d8ebffbbf193b18c13a5cc17d286bf5c92c5ba53ca022af78a9151ec1e5b9c0.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) * 131072u + ((i / 4096u) % 8u) * 16384u + ((i / 64u) % 64u) * 256u + (((i / 1u) % 64u) * 1u + 0u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/4f27100b1e707c866f7cdd701f722bed2969be6e71245d4e4a14e1ef7ebe3fcf.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 * 131072u;
+if (i >= 131072u) { return; }
+let token = i32(b0[i / 2048u]);
+if (token < 131072 || token >= 163840) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 131072) * 2048u + i % 2048u));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/4f2c0d08150d4b97d90c244db2d41baca81484c2630dbbd6d586f2d8a4b07b8b.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 / 2048u) % 64u) * 256u + ((i / 256u) % 8u) * 16384u + ((i / 1u) % 256u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/4fa24e8f09b07271534341707a7a688af67c7deca1f114bfaeba20a2e066c9f1.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 * 1536u;
+if (i >= 1536u) { return; }
+let coord = (i / 512u) % 3u;
+if (coord == 0u) { out[i] = f32(b0[((i / 512u) % 1u) * 512u + ((i / 32u) % 16u) * 32u + ((i / 1u) % 32u) * 1u]); } else { out[i] = f32(b1[i]); }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/4fad29e3e337840e25132f10b3e3d55225f727ae13525ff0b0bdf31e03f09d74.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 * 256u;
+if (i >= 256u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(-x);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/50535e4c99d9fb4a272d8260a1415f6b38a7ad1025da64e145f230df697d9b73.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 * 131072u;
+if (i >= 131072u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(1.0 / (1.0 + exp(-x)));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/51fd4d798491c6a80daff4c18b0495926b7f62269dfc3a78a327faa52ba2a759.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(exp(x));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/52210a661e4d5241f53b3b1caf65456e1988ddd663ded8ead63ee7f7c7e3ff57.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 / 2048u) % 16u) * 128u + ((i / 128u) % 16u) * 2048u + ((i / 1u) % 128u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5244c41c17f08956963dde4c3784e8e122ba0e3c0d76794c8b3673252f672f27.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) * 18944u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 2048u; 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 < 2048u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 64u + row) * 2048u + tile + local.x]); }
+        if (group.x * 8u + local.y < 18944u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = unpack_bf16_1((group.x * 8u + local.y) * 2048u + 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 < 18944u) { out[i] = f32(acc + 0.0); }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/52b439507af68e969a6e116d3c8ea5338a4cde5a3fc1e97b729512d17ab2f4f9.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 / 1u) % 4u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/52d5f0189ff49849070362fea5585100e0c7a9eb6ff6722e3b545c113b965292.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 * 4u + row) * 32768u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 2048u; tile += 8u) {
+        tile_a[local.y * 8u + local.x] = 0.0;
+        tile_b[local.x * 8u + local.y] = 0.0;
+        if (row < 4u && tile + local.x < 2048u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 4u + row) * 2048u + tile + local.x]); }
+        if (group.x * 8u + local.y < 32768u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = unpack_bf16_1((group.x * 8u + local.y) * 2048u + 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 < 4u && col < 32768u) { out[i] = f32(acc + 0.0); }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/53d6eb0b5e6d49bfd3bb9fd1b781e13317adf528e047da2416e95f5a3a13b1b4.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 / 32u) % 1u) * 1u + ((i / 1u) % 32u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/55b9c48ba429daf80bbae0b4737f895b475abdc9ddfdc88ab2f894f133ec2bf9.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 * 30720u;
+if (i >= 30720u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(x / (1.0 + exp(-x)));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/56167a7e8fb1c67fa72948a8baa71866d54cc775be83f8cfa75a947ac0f58b12.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 * 640u;
+if (i >= 640u) { return; }
+out[i] = f32(b0[i]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/564b62d37d6845f8049f903d4ba35762c607937d8a9644bfe511fbc7bf90dacd.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 * 64u;
+    if (row >= 64u) { 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-2b-fp32-int8-g32-home-token-major-v2/kernels/56ea85fd6d5f7cc9394d6aa0ee2fdeee921dc84957548088041c975ba925c6a0.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 * 524288u;
+if (i >= 524288u) { return; }
+let batch = i / 65536u;
+let row = (i / 4096u) % 16u;
+let col = i % 4096u;
+var acc = 0.0;
+for (var p = 0u; p < 256u; p++) {
+    acc += f32(b0[(((batch / 1u) % 8u) * 4096u) + row * 256u + p]) * f32(b1[(((batch / 1u) % 8u) * 1048576u) + p * 4096u + col]);
+}
+out[i] = f32(acc);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/577120bb3cdd5d1fb64df75a3d9d9bfb7df9a47cbe10c3df4702cdedda8742f2.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 * 32768u;
+if (i >= 32768u) { return; }
+let token = i32(b0[i / 2048u]);
+if (token < 131072 || token >= 163840) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 131072) * 2048u + i % 2048u));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/57dc406322bf3badd8490493ef924bd7e7e0afd542d2d6e409b4d8a3b4b784f7.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 >= 16u) { return; }
+out[i] = f32(b0[i]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/58077eb80a5fa253f535c97c82dc07edd48455a17cc52583aeee41c7d1bbd0b3.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 / 512u) % 2u) * 1024u + ((i / 32u) % 16u) * 64u + (((i / 1u) % 32u) * 1u + 0u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/58a04f81024c554fa7384cf7989536e51ab65a353203039f95c9787453ee627c.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) * 18944u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 2048u; 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 < 2048u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 16u + row) * 2048u + tile + local.x]); }
+        if (group.x * 8u + local.y < 18944u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = unpack_bf16_1((group.x * 8u + local.y) * 2048u + 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 < 18944u) { out[i] = f32(acc + 0.0); }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/59ea88d691b919af03bbb26f0253decfbc35d74f9661cdd3a08570ae85c313d2.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 >= 16u) { 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-2b-fp32-int8-g32-home-token-major-v2/kernels/5a3a2ba3e1cb848fbbeeafe66425eae65d9fd196646d15eef291188a8179e044.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 * 256u;
+    if (row >= 256u) { 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-2b-fp32-int8-g32-home-token-major-v2/kernels/5b490e5074446b3ca05ffdd5f79c18e32aa5c23339dd93c75956125f4caba51f.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 * 1u;
+    if (row >= 1u) { 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]) + 1.0));
+    }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5b657fa4e308a89ca0d0edb8cb2a485d6f048dcdbcffb9e2c3a50073be7ad3ff.wgsl

@@ -0,0 +1,22 @@
+@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> b6: array<f32>;
+@group(0) @binding(7) 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 * 917504u;
+if (i >= 917504u) { return; }
+let coord = (i / 1u) % 229376u;
+if (coord >= 0u && coord < 32768u) { out[i] = f32(b0[(i / 229376u) * 32768u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 32768u && coord < 65536u) { out[i] = f32(b1[(i / 229376u) * 32768u + (coord - 32768u) * 1u + i % 1u]); }
+if (coord >= 65536u && coord < 98304u) { out[i] = f32(b2[(i / 229376u) * 32768u + (coord - 65536u) * 1u + i % 1u]); }
+if (coord >= 98304u && coord < 131072u) { out[i] = f32(b3[(i / 229376u) * 32768u + (coord - 98304u) * 1u + i % 1u]); }
+if (coord >= 131072u && coord < 163840u) { out[i] = f32(b4[(i / 229376u) * 32768u + (coord - 131072u) * 1u + i % 1u]); }
+if (coord >= 163840u && coord < 196608u) { out[i] = f32(b5[(i / 229376u) * 32768u + (coord - 163840u) * 1u + i % 1u]); }
+if (coord >= 196608u && coord < 229376u) { out[i] = f32(b6[(i / 229376u) * 32768u + (coord - 196608u) * 1u + i % 1u]); }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5c22e6f60217881c0351e38d19057626a6c2292771c1a83816995245777ac2fd.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 / 32u) % 8u) * 64u + ((i / 32u) % 1u) * 64u + (((i / 1u) % 32u) * 1u + 0u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5c9a5ae6121b3b10611e09e5f1bc563b70191f2506a2afc0f32de823a9d2dac4.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 * 128u;
+if (i >= 128u) { 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-2b-fp32-int8-g32-home-token-major-v2/kernels/5cf28cf1e0aa6d5f0ce563a179880681416b707a70113119e68f1d62dc368c51.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 * 122880u;
+if (i >= 122880u) { return; }
+let coord = (i / 1u) % 20u;
+if (coord >= 0u && coord < 4u) { out[i] = f32(b0[(i / 20u) * 4u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 4u && coord < 20u) { out[i] = f32(b1[(i / 20u) * 16u + (coord - 4u) * 1u + i % 1u]); }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5eac5f92aa4d18895369b11adb99ac0d9639d8b615f7240cffe6f6be7046e54a.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 * 2u;
+    if (row >= 2u) { return; }
+    let lane = local.x;
+    if (lane == 0u) {
+        var total = 0.0;
+        for (var j = 0u; j < 256u; j++) {
+            let v = f32(b0[row * 256u + j]);
+            total += v * v;
+        }
+        factor = inverseSqrt(total / 256.0 + 1e-06);
+    }
+    workgroupBarrier();
+    for (var p = lane; p < 256u; p += 64u) {
+        let i = row * 256u + p;
+        out[i] = f32(f32(b0[row * 256u + p]) * factor * (f32(b1[p]) + 1.0));
+    }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5f4403d78557e88e66d71eb5e53746cc07449a7c421b57eda4f093a09a68a99a.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]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5f57b46016eccc3b50e3685e142b7deadce8f074f124e643e6e7015e2cc27357.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[0u * 512u + ((i / 512u) % 1u) * 512u + ((i / 32u) % 16u) * 32u + ((i / 1u) % 32u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5f657afde3ab80a1fc4b6739c0b702bc5bd60e52630762e66b210fbb8b42b52a.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 * 384u;
+if (i >= 384u) { return; }
+out[i] = f32(b0[((i / 384u) % 1u) * 512u + ((i / 192u) % 2u) * 256u + ((i / 192u) % 1u) * 256u + (((i / 1u) % 192u) * 1u + 64u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5f9d7a23bba38f692aefbf2617133c7884ddbe4ea6d985a89a59c1683dcfd6bf.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]) + (f32(b1[((i / 1u) % 16u) * 1u]) * 1.0));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5fbfe4534ce5bfff645563250dd536998f537f4d690c4c2fc422723d5af8e3ec.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 / 160u) % 1u) * 512u + ((i / 10u) % 16u) * 32u + (((i / 1u) % 10u) * 3u + 2u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5fd6e068841f5c054392bef564909dbf086bae9e4030f3625d1143244a08cbbf.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; }
+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(b0[(i / 64u) * 32u + (coord - 32u) * 1u + i % 1u]); }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/5ffab5e011b6e44aae54b4287fcbe68c4a20cee0779818a3178470f3b6d5ffd6.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 * 8192u;
+if (i >= 8192u) { return; }
+let token = i32(b0[i / 2048u]);
+if (token < 131072 || token >= 163840) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 131072) * 2048u + i % 2048u));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/605f51996eef0769d3b456a83f7540262d1761589b55d1fd41fb2b4dd8fa249f.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]) * 1.0));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/60a9a411c756055d21d8ec478436ba82e0fd522ccf1eb4488aef6b25769d96c5.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 * 8192u;
+if (i >= 8192u) { return; }
+let batch = i / 1024u;
+let row = (i / 256u) % 4u;
+let col = i % 256u;
+var acc = 0.0;
+for (var p = 0u; p < 4096u; p++) {
+    acc += f32(b0[(((batch / 1u) % 8u) * 16384u) + row * 4096u + p]) * f32(b1[(((batch / 1u) % 8u) * 1048576u) + p * 256u + col]);
+}
+out[i] = f32(acc);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/62708aacbfcd0063a897e3348ae69a00d4dd155fce6665b18c5176224b039e12.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 * 131072u;
+if (i >= 131072u) { return; }
+let token = i32(b0[i / 2048u]);
+if (token < 32768 || token >= 65536) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 32768) * 2048u + i % 2048u));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/62f30a62c18ed78c9bf37bdc56529accbee2264b9eac931ba0b35a80bb917e7c.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[2u * 128u + ((i / 128u) % 1u) * 128u + ((i / 32u) % 4u) * 32u + ((i / 1u) % 32u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/63396da046b345383e1c7b09aee551d08e4c3c45030ff5c8bb76bbc34bb2abc6.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 / 4096u) % 8u) * 256u + ((i / 256u) % 16u) * 2048u + ((i / 1u) % 256u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/64f29e6a6e8772f8cfc7449941fd85f70caff7c7e20f316c39b394500e9a8ce1.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(inverseSqrt(x));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/6633b8c4bd21f4e281eb976cdf9e94037ea3de39224dc79ad89e46b6569f5a60.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 / 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<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 * 4u + row) * 6144u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 2048u; tile += 8u) {
+        tile_a[local.y * 8u + local.x] = 0.0;
+        tile_b[local.x * 8u + local.y] = 0.0;
+        if (row < 4u && tile + local.x < 2048u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 4u + row) * 2048u + tile + local.x]); }
+        if (group.x * 8u + local.y < 6144u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = dequant_1((group.x * 8u + local.y) * 2048u + 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 < 4u && col < 6144u) { out[i] = f32(acc + 0.0); }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/66978c862f8d1e391ded2424a4132cfea5f0cc7686c76610502bb596224db368.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> 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> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 4096u;
+    if (i >= 4096u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 2048u; p += 64u) { acc += f32(b0[(i / 4096u) * 2048u + p]) * dequant_1((i % 4096u) * 2048u + p); }
+    partial[lane] = acc;
+    workgroupBarrier();
+    for (var stride = 32u; stride > 0u; stride /= 2u) {
+        if (lane < stride) { partial[lane] += partial[lane + stride]; }
+        workgroupBarrier();
+    }
+    if (lane == 0u) { out[i] = f32(partial[0] + 0.0); }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/66d2c54f014fcbb9053b0710f63280a933d6ad881f5af3c6df4b553c56b53a53.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 / 6144u;
+    let col = index % 6144u;
+    let word = b1[row * 1536u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 192u + 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 * 4u + row) * 2048u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 6144u; tile += 8u) {
+        tile_a[local.y * 8u + local.x] = 0.0;
+        tile_b[local.x * 8u + local.y] = 0.0;
+        if (row < 4u && tile + local.x < 6144u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 4u + row) * 6144u + tile + local.x]); }
+        if (group.x * 8u + local.y < 2048u && tile + local.x < 6144u) { tile_b[local.x * 8u + local.y] = dequant_1((group.x * 8u + local.y) * 6144u + 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 < 4u && col < 2048u) { out[i] = f32(acc + 0.0); }
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/67816e11d408a6230572433502ace2f9a337285eb623360a1cec493e95add139.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 / 1024u) % 2u) * 256u + ((i / 256u) % 4u) * 512u + ((i / 1u) % 256u) * 1u]);
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/67978e951813ede8016ebf553fb8d3b4ce9e15dd184810544879559208d6b7e1.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(1.0 / (1.0 + exp(-x)));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/687a9faf59b35be173c2861749834333b53400d1dd22bf430cef04722237a309.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 * 8192u;
+if (i >= 8192u) { return; }
+let token = i32(b0[i / 2048u]);
+if (token < 98304 || token >= 131072) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 98304) * 2048u + i % 2048u));
+}

qwen35-2b-fp32-int8-g32-home-token-major-v2/kernels/690d8dfb579e55be9c90fa8375f6e2ddde55256a6552a59d83e054f34b485ea0.wgsl

@@ -0,0 +1,22 @@
+@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> b6: array<f32>;
+@group(0) @binding(7) 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 * 229376u;
+if (i >= 229376u) { return; }
+let coord = (i / 1u) % 229376u;
+if (coord >= 0u && coord < 32768u) { out[i] = f32(b0[(i / 229376u) * 32768u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 32768u && coord < 65536u) { out[i] = f32(b1[(i / 229376u) * 32768u + (coord - 32768u) * 1u + i % 1u]); }
+if (coord >= 65536u && coord < 98304u) { out[i] = f32(b2[(i / 229376u) * 32768u + (coord - 65536u) * 1u + i % 1u]); }
+if (coord >= 98304u && coord < 131072u) { out[i] = f32(b3[(i / 229376u) * 32768u + (coord - 98304u) * 1u + i % 1u]); }
+if (coord >= 131072u && coord < 163840u) { out[i] = f32(b4[(i / 229376u) * 32768u + (coord - 131072u) * 1u + i % 1u]); }
+if (coord >= 163840u && coord < 196608u) { out[i] = f32(b5[(i / 229376u) * 32768u + (coord - 163840u) * 1u + i % 1u]); }
+if (coord >= 196608u && coord < 229376u) { out[i] = f32(b6[(i / 229376u) * 32768u + (coord - 196608u) * 1u + i % 1u]); }
+}