Add Qwen3.5-2B mixed FP16 INT4 candidate

.gitattributes

@@ -45,3 +45,4 @@ qwen35-08b-fp16-int4-g32-gptq-home/graph.json filter=lfs diff=lfs merge=lfs -tex
 qwen35-08b-fp16-int4-g32-gptq-home/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
 qwen35-08b-fp16-int4-g8-gptq-up-proj-exclude-l15-candidate/graph.json filter=lfs diff=lfs merge=lfs -text
 qwen35-08b-fp16-int4-g8-gptq-up-proj-exclude-l15-candidate/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
+qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/graph.json filter=lfs diff=lfs merge=lfs -text

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/graph.json

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:ee2e4f0cd06a4ac73213766c544fb6cfca8e2cf39104961bc7469ea3c5c0c6a2
+size 11991700

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/002f3dff1b027a15b1b92e3f04c88efa1ecd7199d6852603412b700547bdc70d.wgsl

@@ -0,0 +1,11 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@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 >= 64u) { return; }
+out[i] = f32(f32(b0[i]) + (f32(b1[((i / 1u) % 16u) * 1u]) * 1.0));
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/023d7b2d9e71a92f74db0e4583d453118915f46ba0c8c96c796632d958654ea1.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 >= 8u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(inverseSqrt(x));
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/02915a0f22b3245bc6490c9bfbedc331f50293e3e7f87525fe4482f0c94e8aff.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[((i / 1536u) % 1u) * 2048u + ((i / 192u) % 8u) * 256u + ((i / 192u) % 1u) * 256u + (((i / 1u) % 192u) * 1u + 64u) * 1u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/02b89cf7780686b8cf46c4db784144cef988e047230cf92ecd1e25e45cabd224.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(f32(b0[i]) * 0.0625);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/035979e9647fb0752d2135ad95461bae23933ee3fc591c65963104942287d12f.wgsl

@@ -0,0 +1,24 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+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] = f32(b1[(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] = f16(acc + 0.0); }
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/037256742da99812bfb6619450f2af103bf17e3d38333c9d309f07aba3a3ae20.wgsl

@@ -0,0 +1,11 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(x / (1.0 + exp(-x)));
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/03d3272d16bc278fcc977e95fcb512bac183f503c7d6ced229f9f01f8e5011bf.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+@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 < 0 || token >= 32768) { out[i] = f16(0.0); return; }
+out[i] = f16(b1[u32(token - 0) * 2048u + i % 2048u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/044fcdf70bcfbb2e3fa472f622959964b1e5929797fcd65fa2eab7016910ee34.wgsl

@@ -0,0 +1,21 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+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 * 18944u;
+    if (i >= 18944u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 2048u; p += 64u) { acc += f32(b0[(i / 18944u) * 2048u + p]) * f32(b1[(i % 18944u) * 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] = f16(partial[0] + 0.0); }
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/04f26661cc9ac73bb7f39f66b76c4be31ee8774e094eeb41ed7d677e1b1d3fa9.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 >= 10u) { return; }
+out[i] = f32(b0[((i / 10u) % 1u) * 32u + ((i / 10u) % 1u) * 32u + (((i / 1u) % 10u) * 3u + 2u) * 1u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/05a8da37281cc1281b5928ab86a013be4a4648b3a422652a3f525829a8981036.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 >= 32u) { return; }
+var acc = 0.0;
+for (var j = 0u; j < 256u; j++) { acc += f32(b0[(((i / 32u) % 1u) * 8192u + ((i / 8u) % 4u) * 2048u + ((i / 1u) % 8u) * 256u) + (((j / 1u) % 256u) * 1u)]); }
+out[i] = f32(acc / 256.0);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0725908abf78c022f86deb338a06779f6f1346ebfe54ff5a43ee6b9eeb93c972.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 >= 44u) { return; }
+out[i] = f32(b0[i]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/078c19606137dad9d459e8041919bac37c31e51702a6d466d4d3df095a6d1f5a.wgsl

@@ -0,0 +1,21 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+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 * 32768u;
+    if (i >= 32768u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 2048u; p += 64u) { acc += f32(b0[(i / 32768u) * 2048u + p]) * f32(b1[(i % 32768u) * 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] = f16(partial[0] + 0.0); }
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0818b2d5166084f4c42b7d514426619af003f77e54291df8f45b72f91dfe5736.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[((i / 256u) % 1u) * 512u + ((i / 128u) % 2u) * 256u + ((i / 32u) % 4u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0835fa80eaadb7c0e6d677e1fc622ad5dcedabb431e5d312f3b7aeb2252a7c7c.wgsl

@@ -0,0 +1,8 @@
+@group(0) @binding(0) 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(0.0);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/08476f486255074c433aee58f82b8ec4150571dc774680f5d8b183a0a7e63e18.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[((i / 8192u) % 1u) * 24576u + ((i / 2048u) % 4u) * 6144u + (((i / 1u) % 2048u) * 1u + 2048u) * 1u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0888235d9d0cb75cbdb598cbb2d235a2db071bed0dd8193164f8f21eab5a8ba9.wgsl

@@ -0,0 +1,11 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+@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] = f16(f32(b0[i]) + (f32(b1[((i / 1u) % 4096u) * 1u]) * 1.0));
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/09e236dd54b3abd945c2d4fef393e9e9810f0d90cdec0d5deafd68ed3b95083d.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 * 512u;
+if (i >= 512u) { return; }
+out[i] = f32(f32(b0[i]) * f32(b1[((i / 256u) % 2u) * 1u]));
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/09e48926914d3759efd11d83aa3ba539f06f33c6f2c28fcc0c12811f9cd1bc44.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[i]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0a9068c3d7c002ffe2631245f0873bf56d726b5d933b040851e20cd8fad9fe1a.wgsl

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

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0b8815b1efe176d55ab08530c0a9e1913cf6c9ef3e773c574843573d237cf196.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(sin(x));
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0d3892883d1532d173a7dbe85d9ab1daa87f0d9b4923a535b8d2003f24f9500c.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 * 256u;
+if (i >= 256u) { 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-fp16-int4-g8-gptq-up-proj-candidate/kernels/0da9c6227714e2a6c031928cb6a6db229c1cb0a53a9309aea4e833a9b2b1f2f1.wgsl

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

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0e6c14d2118fe801932c1b41995a65c8e165bf8df3989e40793f509448988218.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+@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) % 256u;
+if (coord >= 0u && coord < 64u) { out[i] = f16(b0[(i / 256u) * 64u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 64u && coord < 256u) { out[i] = f16(b1[(i / 256u) * 192u + (coord - 64u) * 1u + i % 1u]); }
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0ecd1d59b46e658df21ec09d4256b75bf9a57138f7bfe5e2bcc309a352a37468.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+@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] = f16(0.0); return; }
+out[i] = f16(b1[u32(token - 131072) * 2048u + i % 2048u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0f943f386d2f37a8e56b4285b26e1c093c2f2fa2a5ce81cc43c44cc9bb0137f5.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 / 128u) % 3u) * 128u + ((i / 128u) % 1u) * 128u + ((i / 32u) % 4u) * 1u + ((i / 1u) % 32u) * 4u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0fb8222de35c84fb88807d4a3880f21188cf48d64c89a21aa208d7d72a48c273.wgsl

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

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/0ff3d9cacb67aa460769fcf57dcfe10c5354e7e98681a4647a9acf1045d60892.wgsl

@@ -0,0 +1,21 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+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 * 2048u;
+    if (i >= 2048u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 6144u; p += 64u) { acc += f32(b0[(i / 2048u) * 6144u + p]) * f32(b1[(i % 2048u) * 6144u + 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] = f16(partial[0] + 0.0); }
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/11a92e9cba6bfc381c98c5127db137639d6c87f1cb220459880ead039dc81256.wgsl

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

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/1326034f3abab37629cd88f869d1b0bafcde724f85c07904b54feaaf71d89c62.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 < 4u; token++) {
+        recurrent *= exp(f32(b3[(batch_index * 4u + token) * 16u + head]));
+        partial[lane] = 0.0;
+        if (lane < 128u) { partial[lane] = recurrent * f32(b1[((batch_index * 4u + 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 * 4u + token) * 16u + head) * 128u + column]) - partial[0]) * f32(b4[(batch_index * 4u + token) * 16u + head]); }
+        workgroupBarrier();
+        if (lane < 128u) { recurrent += f32(b1[((batch_index * 4u + token) * 16u + head) * 128u + lane]) * delta; }
+        partial[lane] = 0.0;
+        if (lane < 128u) {
+            partial[lane] = recurrent * f32(b0[((batch_index * 4u + 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) * 132u + 128u + token) * 128u + column] = f32(partial[0]); }
+        workgroupBarrier();
+    }
+    if (lane < 128u) { out[((batch_index * 16u + head) * 132u + lane) * 128u + column] = f32(recurrent); }
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/13dc73935b9b3b978232043e13babee3877bd53ca572be440de0590dacce8b17.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[0u * 32u + ((i / 32u) % 1u) * 32u + ((i / 32u) % 1u) * 32u + ((i / 1u) % 32u) * 1u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/1545025e0a625ccc73b42ae112c6f9d78a1504c10b759979ae2247984f25f056.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 >= 1u && coord < 34u && (coord - 1u) % 3u == 0u) { out[i] = f32(b0[((i / 128u) % 1u) * 44u + ((i / 32u) % 4u) * 11u + ((((i / 1u) % 32u) - 1u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/1608c6cf00b5863ad002f089be58fce76cff2c61fd43e3bb54ba90add6f41b02.wgsl

@@ -0,0 +1,11 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(-x);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/16757bdfacc94a92cfc58c9a3ba83ca377b8ea7ddf82fa8f6611f797a6934f14.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+@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 < 32768 || token >= 65536) { out[i] = f16(0.0); return; }
+out[i] = f16(b1[u32(token - 32768) * 2048u + i % 2048u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/16de135b328f53749dc017ed0a624ea63f39df86947f72a0d5d4aa06bedeadd1.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[i]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/17d9b9ec53d19b9ce8e3e750fd7beeb668dbc872a19e0ab5d29e566d84490e74.wgsl

@@ -0,0 +1,21 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+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 * 2048u;
+    if (i >= 2048u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 2048u; p += 64u) { acc += f32(b0[(i / 2048u) * 2048u + p]) * f32(b1[(i % 2048u) * 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] = f16(partial[0] + 0.0); }
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/17ec6bdaf98edf7657b74a25036c2bec806ecbb4280f33fff6afd2e0de5ed377.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(f32(b0[i]) * 0.08838834764831843);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/18b1cef2bae996f56664d9215180e3bfc77b4e514db52fdf9b6c58be51b53edb.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 >= 44u) { return; }
+out[i] = f32(b0[((i / 44u) % 1u) * 128u + ((i / 11u) % 4u) * 32u + (((i / 1u) % 11u) * 3u + 1u) * 1u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/19564c145248f065669be23f5daeeda96e025ee4389bef80983001e83fddb31d.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 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-fp16-int4-g8-gptq-up-proj-candidate/kernels/1a3351b3225f169402dd4c0b1bfbcc407f8b5fdccfb10c74ccb80b7c3bd85a37.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+@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 = i32(b0[i / 2048u]);
+if (token < 196608 || token >= 229376) { out[i] = f16(0.0); return; }
+out[i] = f16(b1[u32(token - 196608) * 2048u + i % 2048u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/1c7201caab74d50198d8bc7168ba7da0c97ece74a5ebdd791f88d6cfb5086975.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 >= 8u) { return; }
+out[i] = f32(f32(b0[i]) + (1e-06 * 1.0));
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/1cb76debd822c58d652ea0e82d2f5bdebc3dd67a8cf540154fd85f504fc72cd2.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[((i / 8192u) % 1u) * 8192u + ((i / 2048u) % 4u) * 256u + ((i / 256u) % 8u) * 1024u + ((i / 1u) % 256u) * 1u]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/1ce543c34cc4033e22afd4fc00a62ed19f44307cc4aeb52e7a503538e5831e9d.wgsl

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

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/1df96eddb194a002a44f5a610245f84d192a38d2fff876c30fd55a4f518e8b25.wgsl

@@ -0,0 +1,11 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(-x);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/1f252e46a7c6dc933462fc88d57b832b6e87ffe79232ad52b4dad28e5083cc69.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(f32(b0[i]) * 0.0625);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/1f821653e02930a221ffe24970bbbd13b88c5e7bd119af97f624b075b1ab6534.wgsl

@@ -0,0 +1,14 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<i32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[i]);
+}

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/222a2645e3e02d0a4063a63ca1b6fa4c89faee6b76905b20de4dd54ab3077ea3.wgsl

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

qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/kernels/234987ec68497956ea7aef8e3f44c8a73dbe4c1519165d123946bedbb4d18e26.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 / 128u) % 16u) * 128u + 0u * 128u + ((i / 1u) % 128u) * 1u]);
+}