Add Qwen3.5-0.8B FP16 INT4 GPTQ experiment

.gitattributes

@@ -41,3 +41,5 @@ qwen35-2b-fp32-int8-g32-home-token-major-v2/graph.json filter=lfs diff=lfs merge
 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-multimodal-fp32-three-grid-state-alias-token-major-v2/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
+qwen35-08b-fp16-int4-g32-gptq-home/graph.json filter=lfs diff=lfs merge=lfs -text
+qwen35-08b-fp16-int4-g32-gptq-home/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text

qwen35-08b-fp16-int4-g32-gptq-home/graph.json

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

qwen35-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/kernels/046b8b39f33faa9e39a553acc1912fbce2c4bc4178588df30cf98717bca6f328.wgsl

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

qwen35-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/kernels/06f0c9ec30de4203c86ec1fa0eaeeea3450d838b78b5c7a02df1c61616405c4d.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 * 1024u;
+if (i >= 1024u) { return; }
+out[i] = f16(b0[i]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/kernels/0bd599ee0e0d3e758477e5b3603c6b8cebf703223656bc3f4b53dc922e2bd314.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 * 1024u;
+if (i >= 1024u) { return; }
+out[i] = f16(f32(b0[i]) + (f32(b1[i]) * 1.0));
+}

qwen35-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/kernels/171389dc1ab21b1521080c899bd9685baa2e05324c4ee95e4b0b042b9ecf4e9b.wgsl

@@ -0,0 +1,29 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@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<f16>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 1024u;
+    let col = index % 1024u;
+    let word = b1[row * 128u + col / 8u];
+    let code = i32((word >> ((col % 8u) * 4u)) & 15u) - 8;
+    return f32(f16(f32(code) * b2[row * 32u + col / 32u]));
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 16u;
+    if (i >= 16u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 16u) * 1024u + p]) * dequant_1((i % 16u) * 1024u + 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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/kernels/197465241bc85d39f754ff558d69abf289859cd7e13f5cf1bdaad66c4c74e94f.wgsl

@@ -0,0 +1,29 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@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<f16>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 1024u;
+    let col = index % 1024u;
+    let word = b1[row * 128u + col / 8u];
+    let code = i32((word >> ((col % 8u) * 4u)) & 15u) - 8;
+    return f32(f16(f32(code) * b2[row * 32u + col / 32u]));
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 512u;
+    if (i >= 512u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 512u) * 1024u + p]) * dequant_1((i % 512u) * 1024u + 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-08b-fp16-int4-g32-gptq-home/kernels/1c3075bda51d817a45cebcbb10fdd9db05093be2843c47c6a0669a058892126b.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 * 51712u;
+    if (i >= 51712u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 51712u) * 1024u + p]) * f32(b1[(i % 51712u) * 1024u + 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-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/kernels/20dd7df6a12031057faa143e5c2c08bfd0d92bbe990d83ea61a6d540e6216649.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 * 1024u;
+if (i >= 1024u) { return; }
+let token = i32(b0[i / 1024u]);
+if (token < 196608 || token >= 248320) { out[i] = f16(0.0); return; }
+out[i] = f16(b1[u32(token - 196608) * 1024u + i % 1024u]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/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-08b-fp16-int4-g32-gptq-home/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]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/26862ccab76186d81a45f06c5a7c9e1cd04e4592dbe151341043725a49a5cd2f.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 * 24576u;
+if (i >= 24576u) { return; }
+out[i] = f16(b0[((i / 24576u) % 1u) * 30720u + ((i / 4u) % 6144u) * 5u + (((i / 1u) % 4u) * 1u + 1u) * 1u]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/2711a4a396071f4e99303d01d4854ccaa37a445f4b718ebd7b61cf1c37e56b48.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 / 16u) % 1u) * 16u + ((i / 1u) % 16u) * 1u + ((i / 1u) % 1u) * 16u]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/27ed9ce421e40975ea99531a57e86dbcde5310faae3826bc81bc60fcdbaa1f4e.wgsl

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

qwen35-08b-fp16-int4-g32-gptq-home/kernels/291830b3cb692bab94344428e329c3dc9417a43f48e1c42e414cf3f956686f1f.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 * 2048u;
+if (i >= 2048u) { return; }
+out[i] = f16(f32(b0[((i / 1u) % 128u) * 1u]) * f32(b1[i]));
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/293b36448f6af1c167040d45840a2ae8f05ca669424bdebd16fc2c7e5bc5bfff.wgsl

@@ -0,0 +1,29 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@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<f16>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 1024u;
+    let col = index % 1024u;
+    let word = b1[row * 128u + col / 8u];
+    let code = i32((word >> ((col % 8u) * 4u)) & 15u) - 8;
+    return f32(f16(f32(code) * b2[row * 32u + col / 32u]));
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 3584u;
+    if (i >= 3584u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 3584u) * 1024u + p]) * dequant_1((i % 3584u) * 1024u + 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-08b-fp16-int4-g32-gptq-home/kernels/2c0cca2a7b0198c317e68992e64f3fe9bba80b67204c2695c20c372471c150d3.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 * 262144u;
+if (i >= 262144u) { return; }
+out[i] = f32(f32(b0[((i / 16384u) % 16u) * 128u + ((i / 128u) % 128u) * 1u]) * f32(b1[((i / 16384u) % 16u) * 128u + ((i / 1u) % 128u) * 1u]));
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/2db8a995aaf1e89bf2da258d8e77c4ccc53a4d61ab9b352757da0e5a7321729d.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(x*x);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/2dd29b61b039fcdd7fbc3c94f629b24b6f81aa5b9923098837597359b295e40d.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) * 128u + ((i / 128u) % 16u) * 128u + ((i / 1u) % 128u) * 1u]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a.wgsl

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

qwen35-08b-fp16-int4-g32-gptq-home/kernels/2ecb3b88f329d2d66b34465ebc62482b3a9d40d37c69a30d55b69d2c85f082cb.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 * 2048u;
+if (i >= 2048u) { return; }
+let x = f32(b0[i]);
+out[i] = f16(1.0 / (1.0 + exp(-x)));
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/32a5027cac5dd969d63fe5d5c88d7c87abb600c43ad1d70e1c56615ea781eff3.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 * 262144u;
+if (i >= 262144u) { return; }
+out[i] = f32(f32(b0[i]) * f32(b1[((i / 16384u) % 16u) * 128u + ((i / 128u) % 128u) * 1u]));
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/344153d8e20e545c4656ce08b7236cd65878f5ea91014447f4a129ea94025970.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 / 32u) % 8u) * 64u + ((i / 32u) % 1u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/370bab74beead2c5e9cc8893585f79ae97ee5eace554cee12555a8ac8293b1c0.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]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/377880fa04b2ba2c6238045a88750ff214243eb739603a04b16b69221ec5a6c2.wgsl

@@ -0,0 +1,18 @@
+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; }
+let batch = i / 4096u;
+let row = (i / 4096u) % 1u;
+let col = i % 4096u;
+var acc = 0.0;
+for (var p = 0u; p < 256u; p++) {
+    acc += f32(b0[(((batch / 1u) % 8u) * 256u) + row * 256u + p]) * f32(b1[(((batch / 1u) % 8u) * 1048576u) + p * 4096u + col]);
+}
+out[i] = f16(acc);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/378330e4bcfd0c2f194f2ca261988c352ba4cdffd71dee16af52c3b7ba25657f.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 * 2048u;
+if (i >= 2048u) { return; }
+out[i] = f16(b0[((i / 2048u) % 1u) * 2048u + ((i / 2048u) % 1u) * 256u + ((i / 256u) % 8u) * 256u + ((i / 1u) % 256u) * 1u]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/3a87037580f20ca38623bdcc88abd77f8cc3a83d8b3d1726756a1ec01656a496.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 * 6144u;
+if (i >= 6144u) { return; }
+out[i] = f16(b0[((i / 6144u) % 1u) * 6144u + ((i / 1u) % 6144u) * 1u + ((i / 1u) % 1u) * 6144u]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/3bb850bda95bf77706a2839616a401b773f9de0428d1b427e61fb04e2c35fe74.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 >= 96u) { return; }
+let coord = (i / 32u) % 3u;
+if (coord == 0u) { out[i] = f32(b0[((i / 32u) % 1u) * 32u + ((i / 32u) % 1u) * 32u + ((i / 1u) % 32u) * 1u]); } else { out[i] = f32(b1[i]); }
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/3f78504435af26dbd6b62c6720a640528ce172b2e9062bd4d09806e3d2553423.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 * 512u;
+    if (i >= 512u) { return; }
+    let token = (i / 256u) % 1u;
+    let outer = i / 256u;
+    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-08b-fp16-int4-g32-gptq-home/kernels/3fa1efae6d90015ba69cb40cef009f49ad7f7a960018fb92a962707deff4f09e.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]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/3ff863808627e030edefbb05f743b5a2cc19412b22779508a9d173a40d2ecc6e.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 >= 2u) { return; }
+var acc = 0.0;
+for (var j = 0u; j < 256u; j++) { acc += f32(b0[(((i / 2u) % 1u) * 512u + ((i / 2u) % 1u) * 512u + ((i / 1u) % 2u) * 256u) + (((j / 1u) % 256u) * 1u)]); }
+out[i] = f32(acc / 256.0);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/408b4edf1721b552fd1406f682e9c1612b2662372860bf14e040f011c1ee0cb9.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 * 2048u;
+if (i >= 2048u) { return; }
+out[i] = f16(f32(b0[i]) * f32(b1[((i / 128u) % 16u) * 1u]));
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/4282aa48d394c7c4b0d6c6f1063526bda4b9da2708f77a1945793a4614107efb.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 * 2048u;
+if (i >= 2048u) { return; }
+out[i] = f16(b0[((i / 2048u) % 1u) * 4096u + ((i / 2048u) % 1u) * 4096u + ((i / 256u) % 8u) * 512u + (((i / 1u) % 256u) * 1u + 256u) * 1u]);
+}

qwen35-08b-fp16-int4-g32-gptq-home/kernels/42c5af4838daf75fff9d4bb9a5df409511c1d6daed617d05eaf033a977503eac.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 * 64u;
+if (i >= 32u) { return; }
+let coord = (i / 1u) % 32u;
+if (coord >= 1u && coord < 34u && (coord - 1u) % 3u == 0u) { out[i] = f32(b0[((i / 32u) % 1u) * 11u + ((i / 32u) % 1u) * 11u + ((((i / 1u) % 32u) - 1u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
+}