9bow commited on
Commit
a2caca3
·
verified ·
1 Parent(s): 05182a7

Add qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4 (tiled prefill GEMM over the v3 bundle) (part 2)

Browse files
This view is limited to 50 files because it contains too many changes.   See raw diff
Files changed (50) hide show
  1. .gitattributes +1 -0
  2. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a1a974f6d8d51fd7c3496ede0966af303ad9f0c0c9ee44cf93ec962d2b0e87e1.wgsl +10 -0
  3. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a1cf6b03baf18d389f0a2b95440c17125a8ecc1e633568de985d42cac6e0d1cb.wgsl +9 -0
  4. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a21c6a35532bc26ccfc46df22f98e1592e3b285e3800544c032edf39425bde1a.wgsl +9 -0
  5. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a24673da4fae6d1c49aa143f8425996b2d036db7b0341987dab330c78ead91d7.wgsl +9 -0
  6. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a27baed9f1bab1d547213060cfb98f64b59479a2b7d7e41486ce8d7e7c1888d8.wgsl +12 -0
  7. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a2bd65fc2709e093644d6cbf0c292b4d71520009a87bedcf4bca8fcb5f5c15b8.wgsl +9 -0
  8. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a2e242b97ac8ec546c9f91b5af599c2231915bdabb57d55cdc21f61a49c21007.wgsl +10 -0
  9. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a37dac611b27158f19aa5bd70c43217286e610e733cbd602b91666b4af42d3ef.wgsl +9 -0
  10. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a3dcb26b08c4e2c0b11e188a362984db75cae2d09ed11e1029d69576b99ed9ac.wgsl +20 -0
  11. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a405c3a6203b40db5812e5299cee829b88a6cee8742a177b8e6bf0395ed52d89.wgsl +83 -0
  12. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a53c1ce51ebaec593ad667efae9603efe3807db8f0ee060a620ac031cc4f2b92.wgsl +10 -0
  13. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a62f6b62277fe048d6bac44a25a2e61399e7e2eab103fccd3a23c29a3c827856.wgsl +44 -0
  14. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a672c3e158716bde7444c8141f3df1da16da739ccdd2eb2ad60f9d42d0775157.wgsl +11 -0
  15. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a77f14c89662d2d3da8abd8f109667793dc662d307e5e1c61123be753e25437b.wgsl +9 -0
  16. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a7b50e8b6e4d56404b6261383f50b3eed4a2a59c3fadaaffcc974575c32a374f.wgsl +9 -0
  17. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a87152b0fd2f78f60d7abb61c2f6c4a9bde8a9dd45d0275b58d760a3da9f71ce.wgsl +17 -0
  18. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a896c93b04922c0f78f1c77fbaabbc79eb71776223748bb32a6d5339e32029e8.wgsl +9 -0
  19. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a9f0eef095f015e2099adf81552a6727c45e70fc11f5e50b777d36b6370ab986.wgsl +9 -0
  20. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a9f6bc361f1ea31f79f0ef6164872eeb04896a45338871a2d779789b47b1920b.wgsl +10 -0
  21. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a9fa5db20fa8de1de429c6850c3e40ff8180bcdb59373919b84d2438b3b17352.wgsl +9 -0
  22. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/aa227ae49fab3f75d86e02e6f390be785eabfeafe70e77a5c7eb4ad51895e382.wgsl +10 -0
  23. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/aac6395f039e396767ee10d0ad95d3d6206ad4a5c1e5913dfd5c3cfeaf227510.wgsl +9 -0
  24. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/abdb17daf29310a7828a55a1efeebaa65b6553e236fc20702e8c712e0a46dd5b.wgsl +10 -0
  25. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/ac173d038b5767711ad9f7cafbceb9ca3ec84aaeb3d07a3037e7cd0f83fead2c.wgsl +9 -0
  26. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/ad72f60d58133272228319f1c066e7cd43ac28573ddeeb7a6f69be01af341f57.wgsl +9 -0
  27. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/adf1e58a124e1301625024ae142dca5fb1dd50799fa94b52eaeffb82ea99852c.wgsl +10 -0
  28. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/adfbba60ade7143599d7de0faa4904afc1400ccc3cf84602f22a2fb004d0398e.wgsl +71 -0
  29. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/ae69e698c92571766e36ad8c078b13e42bde56745d45698daaadaca587b7279c.wgsl +11 -0
  30. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b02b0be460d581c720afd68c8b722e8da47e2ee432b5d61bbdf24b6084cb2f26.wgsl +11 -0
  31. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b06d10513b70659dd70c40974ccc40867e6c17dacc4d4f5c131f6d99c728138d.wgsl +41 -0
  32. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b12eae76f13225c0c5d35fc9309964aba3172712799ff3bd8538bb2c1b34aa33.wgsl +9 -0
  33. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b199d835db70f888c7b99a3ac6e9ec56f6c47d5bba4eaf352e0a8eabed56411e.wgsl +10 -0
  34. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b28ea4a8345a9c9abb2d74bfd620ba5de46616434dbaffdbabf022ea71800839.wgsl +9 -0
  35. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b2ae1b161c8a6315f8863e47511607027dfe583aac1db4c1626371ddfa0c64e2.wgsl +9 -0
  36. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b2b16798cbe94c919025d76b0435afe47966ffdd1ea53ec74a791c03c4709a28.wgsl +9 -0
  37. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b3bb2905555188f535b6b8bc1d050c142b5fea8222d9ef193466f80aa2c54dea.wgsl +11 -0
  38. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b441eef0119fbc7cc4c3757e12f57cffdeeac0bbed3285b5757f7b9a20feb669.wgsl +41 -0
  39. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b57928140994202536cd4f99d2e14f5dba9754f014843673790d5d48c9cbb9ab.wgsl +9 -0
  40. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b5984ff6eb9e56ef0f9930f775c1dc3c203eec15924e5fea4c2525b2ce180857.wgsl +9 -0
  41. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b6c3ce0defd66eef66dc3fd6eb9eab3fee12d179332698c3c764009988628b49.wgsl +10 -0
  42. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bb3e27e020bc22a0c8c21a5f9910a0a26b73ace5f68d9d49d1e4d67d53232060.wgsl +71 -0
  43. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bb7f03e6a2277dc56084e9130086507be58871253867c0ce8ad60ee5ce4521b6.wgsl +9 -0
  44. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bb9029b667eeb13670d1f640cd677c6eabd7a7569686c599e3eeb9482d642b00.wgsl +10 -0
  45. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bcb78b92dbbfeb212c190e064bc2f9d8e79c201398cf38c9aa62df6fdba94469.wgsl +9 -0
  46. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bd70e4fc42ac55f0f52612841bcc3f9fe7df5bbb0fc76f94e56a131a92454658.wgsl +9 -0
  47. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bd8a785d4fbf09918f149c7eb5f34bd25f9358bf3412e27799a7c65436c71759.wgsl +9 -0
  48. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bda7ebd3bcf25c31938fca1a582e82b972ab742f85f826f085417b0e6f651e60.wgsl +12 -0
  49. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/c02ac02d110c0be5d0990d20ab0d8c7a081c74f4e147f083789a203f22fc3258.wgsl +10 -0
  50. qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/c0b2947798baa198598d75ca0e5dcd31a9fa320499665e34f7bb81d27fc15fb0.wgsl +10 -0
.gitattributes CHANGED
@@ -66,3 +66,4 @@ qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2-v3/graph.json fi
66
  qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2-v3/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
67
  qwen35-08b-fp32-8k-token-major-home-v4/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
68
  qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/graph.json filter=lfs diff=lfs merge=lfs -text
 
 
66
  qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2-v3/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
67
  qwen35-08b-fp32-8k-token-major-home-v4/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
68
  qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/graph.json filter=lfs diff=lfs merge=lfs -text
69
+ qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a1a974f6d8d51fd7c3496ede0966af303ad9f0c0c9ee44cf93ec962d2b0e87e1.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 64u;
7
+ if (i >= 64u) { return; }
8
+ let x = f32(b0[i]);
9
+ out[i] = f32(-x);
10
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a1cf6b03baf18d389f0a2b95440c17125a8ecc1e633568de985d42cac6e0d1cb.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 512u;
7
+ if (i >= 512u) { return; }
8
+ out[i] = f32(b0[1u * 512u + ((i / 512u) % 1u) * 512u + ((i / 32u) % 16u) * 32u + ((i / 1u) % 32u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a21c6a35532bc26ccfc46df22f98e1592e3b285e3800544c032edf39425bde1a.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 98304u;
7
+ if (i >= 98304u) { return; }
8
+ out[i] = f32(b0[((i / 98304u) % 1u) * 98304u + ((i / 16u) % 6144u) * 1u + ((i / 1u) % 16u) * 6144u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a24673da4fae6d1c49aa143f8425996b2d036db7b0341987dab330c78ead91d7.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 32768u;
7
+ if (i >= 32768u) { return; }
8
+ out[i] = f32(b0[((i / 32768u) % 1u) * 98304u + ((i / 2048u) % 16u) * 6144u + (((i / 1u) % 2048u) * 1u + 4096u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a27baed9f1bab1d547213060cfb98f64b59479a2b7d7e41486ce8d7e7c1888d8.wgsl ADDED
@@ -0,0 +1,12 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 248320u;
8
+ if (i >= 248320u) { return; }
9
+ let coord = (i / 1u) % 248320u;
10
+ if (coord >= 0u && coord < 131072u) { out[i] = f32(b0[(i / 248320u) * 131072u + (coord - 0u) * 1u + i % 1u]); }
11
+ if (coord >= 131072u && coord < 248320u) { out[i] = f32(b1[(i / 248320u) * 117248u + (coord - 131072u) * 1u + i % 1u]); }
12
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a2bd65fc2709e093644d6cbf0c292b4d71520009a87bedcf4bca8fcb5f5c15b8.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<i32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 64u;
7
+ if (i >= 12u) { return; }
8
+ out[i] = f32(b0[i]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a2e242b97ac8ec546c9f91b5af599c2231915bdabb57d55cdc21f61a49c21007.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 1024u;
7
+ if (i >= 1024u) { return; }
8
+ let x = f32(b0[i]);
9
+ out[i] = f32(1.0 / (1.0 + exp(-x)));
10
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a37dac611b27158f19aa5bd70c43217286e610e733cbd602b91666b4af42d3ef.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<i32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<i32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 64u;
7
+ if (i >= 1u) { return; }
8
+ out[i] = i32(b0[i]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a3dcb26b08c4e2c0b11e188a362984db75cae2d09ed11e1029d69576b99ed9ac.wgsl ADDED
@@ -0,0 +1,20 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<i32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<u32>;
3
+ @group(0) @binding(2) var<storage, read> b2: array<f32>;
4
+ @group(0) @binding(3) var<storage, read_write> out: array<f32>;
5
+ fn dequant_1(index: u32) -> f32 {
6
+ let row = index / 1024u;
7
+ let col = index % 1024u;
8
+ let word = b1[row * 256u + col / 4u];
9
+ let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
10
+ return f32(code) * b2[row * 32u + col / 32u];
11
+ }
12
+
13
+ @compute @workgroup_size(64)
14
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
15
+ let i = gid.x + gid.y * 1024u;
16
+ if (i >= 1024u) { return; }
17
+ let token = i32(b0[i / 1024u]);
18
+ if (token < 131072 || token >= 248320) { out[i] = f32(0.0); return; }
19
+ out[i] = f32(dequant_1(u32(token - 131072) * 1024u + i % 1024u));
20
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a405c3a6203b40db5812e5299cee829b88a6cee8742a177b8e6bf0395ed52d89.wgsl ADDED
@@ -0,0 +1,83 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read> b2: array<f32>;
4
+ @group(0) @binding(3) var<storage, read> b3: array<f32>;
5
+ @group(0) @binding(4) var<storage, read_write> out: array<f32>;
6
+
7
+ var<workgroup> q_row: array<f32, 256>;
8
+ var<workgroup> scores: array<f32, 1024>;
9
+ var<workgroup> partial: array<f32, 64>;
10
+ @compute @workgroup_size(64)
11
+ fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
12
+ let lane = local.x;
13
+ let head = group.x / 16u;
14
+ let row = group.x % 16u;
15
+ let kv_head = head / 4u;
16
+ for (var d = lane; d < 256u; d += 64u) { q_row[d] = f32(b0[(head * 16u + row) * 256u + d]); }
17
+ workgroupBarrier();
18
+ var running_max = -3.4028234663852886e38;
19
+ var running_sum = 0.0;
20
+ var acc: array<f32, 4>;
21
+ for (var start = 0u; start < 4096u; start += 1024u) {
22
+ let count = min(1024u, 4096u - start);
23
+ var tile_max = -3.4028234663852886e38;
24
+ for (var t = lane; t < count; t += 64u) {
25
+ let j = start + t;
26
+ var score = -3.4028234663852886e38;
27
+ if (f32(b3[row * 4096u + j]) > -1000.0) {
28
+ var dot = 0.0;
29
+ for (var d = 0u; d < 256u; d++) { dot += q_row[d] * f32(b1[(((kv_head * 4096u + j) * 256u + d) / 256u) % 4096u * 512u + (((kv_head * 4096u + j) * 256u + d) / 1048576u) % 2u * 256u + (((kv_head * 4096u + j) * 256u + d) / 1u) % 256u * 1u]); }
30
+ score = dot * 0.0625 + f32(b3[row * 4096u + j]);
31
+ }
32
+ scores[t] = score;
33
+ tile_max = max(tile_max, score);
34
+ }
35
+ partial[lane] = tile_max;
36
+ workgroupBarrier();
37
+ for (var stride = 32u; stride > 0u; stride /= 2u) {
38
+ if (lane < stride) { partial[lane] = max(partial[lane], partial[lane + stride]); }
39
+ workgroupBarrier();
40
+ }
41
+ let next_max = max(running_max, partial[0]);
42
+ workgroupBarrier();
43
+ var tile_sum = 0.0;
44
+ for (var t = lane; t < count; t += 64u) {
45
+ var p = 0.0;
46
+ if (scores[t] > -3.0e38) { p = exp(scores[t] - next_max); }
47
+ scores[t] = p;
48
+ tile_sum += p;
49
+ }
50
+ partial[lane] = tile_sum;
51
+ workgroupBarrier();
52
+ for (var stride = 32u; stride > 0u; stride /= 2u) {
53
+ if (lane < stride) { partial[lane] += partial[lane + stride]; }
54
+ workgroupBarrier();
55
+ }
56
+ var correction = 0.0;
57
+ if (running_max > -3.0e38) { correction = exp(running_max - next_max); }
58
+ running_sum = running_sum * correction + partial[0];
59
+ running_max = next_max;
60
+ for (var c = 0u; c < 4u; c++) {
61
+ let d = lane + c * 64u;
62
+ var total = acc[c] * correction;
63
+ if (d < 256u) {
64
+ for (var t = 0u; t < count; t++) {
65
+ let p = scores[t];
66
+ if (p != 0.0) {
67
+ let j = start + t;
68
+ total += p * f32(b2[(((kv_head * 4096u + j) * 256u + d) / 256u) % 4096u * 512u + (((kv_head * 4096u + j) * 256u + d) / 1048576u) % 2u * 256u + (((kv_head * 4096u + j) * 256u + d) / 1u) % 256u * 1u]);
69
+ }
70
+ }
71
+ }
72
+ acc[c] = total;
73
+ }
74
+ workgroupBarrier();
75
+ }
76
+ for (var c = 0u; c < 4u; c++) {
77
+ let d = lane + c * 64u;
78
+ if (d < 256u) {
79
+ let i = (head * 16u + row) * 256u + d;
80
+ out[i] = f32(acc[c] / running_sum);
81
+ }
82
+ }
83
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a53c1ce51ebaec593ad667efae9603efe3807db8f0ee060a620ac031cc4f2b92.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 8192u;
7
+ if (i >= 8192u) { return; }
8
+ let x = f32(b0[i]);
9
+ out[i] = f32(1.0 / (1.0 + exp(-x)));
10
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a62f6b62277fe048d6bac44a25a2e61399e7e2eab103fccd3a23c29a3c827856.wgsl ADDED
@@ -0,0 +1,44 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read> b2: array<f32>;
4
+ @group(0) @binding(3) var<storage, read> b3: array<f32>;
5
+ @group(0) @binding(4) var<storage, read> b4: array<f32>;
6
+ @group(0) @binding(5) var<storage, read> b5: array<f32>;
7
+ @group(0) @binding(6) var<storage, read_write> out: array<f32>;
8
+
9
+ var<workgroup> partial: array<f32, 128>;
10
+ var<workgroup> delta: f32;
11
+ @compute @workgroup_size(128)
12
+ fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
13
+ let column = group.x;
14
+ let head = group.y;
15
+ let batch_index = group.z;
16
+ let lane = local.x;
17
+ var recurrent = 0.0;
18
+ if (lane < 128u) { recurrent = f32(b5[((batch_index * 16u + head) * 128u + lane) * 128u + column]); }
19
+ for (var token = 0u; token < 64u; token++) {
20
+ recurrent *= exp(f32(b3[(batch_index * 64u + token) * 16u + head]));
21
+ partial[lane] = 0.0;
22
+ if (lane < 128u) { partial[lane] = recurrent * f32(b1[((batch_index * 64u + token) * 16u + head) * 128u + lane]); }
23
+ workgroupBarrier();
24
+ for (var stride = 64u; stride > 0u; stride /= 2u) {
25
+ if (lane < stride) { partial[lane] += partial[lane + stride]; }
26
+ workgroupBarrier();
27
+ }
28
+ if (lane == 0u) { delta = (f32(b2[((batch_index * 64u + token) * 16u + head) * 128u + column]) - partial[0]) * f32(b4[(batch_index * 64u + token) * 16u + head]); }
29
+ workgroupBarrier();
30
+ if (lane < 128u) { recurrent += f32(b1[((batch_index * 64u + token) * 16u + head) * 128u + lane]) * delta; }
31
+ partial[lane] = 0.0;
32
+ if (lane < 128u) {
33
+ partial[lane] = recurrent * f32(b0[((batch_index * 64u + token) * 16u + head) * 128u + lane]) * 0.08838834764831845;
34
+ }
35
+ workgroupBarrier();
36
+ for (var stride = 64u; stride > 0u; stride /= 2u) {
37
+ if (lane < stride) { partial[lane] += partial[lane + stride]; }
38
+ workgroupBarrier();
39
+ }
40
+ if (lane == 0u) { out[((batch_index * 16u + head) * 192u + 128u + token) * 128u + column] = f32(partial[0]); }
41
+ workgroupBarrier();
42
+ }
43
+ if (lane < 128u) { out[((batch_index * 16u + head) * 192u + lane) * 128u + column] = f32(recurrent); }
44
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a672c3e158716bde7444c8141f3df1da16da739ccdd2eb2ad60f9d42d0775157.wgsl ADDED
@@ -0,0 +1,11 @@
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 2048u;
8
+ if (i >= 2048u) { return; }
9
+ let coord = (i / 128u) % 1u;
10
+ if (coord == 0u) { out[i] = f32(b0[((i / 2048u) % 1u) * 2048u + ((i / 128u) % 16u) * 128u + ((i / 1u) % 128u) * 1u]); } else { out[i] = f32(b1[i]); }
11
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a77f14c89662d2d3da8abd8f109667793dc662d307e5e1c61123be753e25437b.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 128u;
7
+ if (i >= 96u) { return; }
8
+ out[i] = f32(b0[((i / 1u) % 32u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a7b50e8b6e4d56404b6261383f50b3eed4a2a59c3fadaaffcc974575c32a374f.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 98304u;
7
+ if (i >= 98304u) { return; }
8
+ out[i] = f32(b0[((i / 98304u) % 1u) * 98304u + ((i / 6144u) % 16u) * 1u + ((i / 1u) % 6144u) * 16u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a87152b0fd2f78f60d7abb61c2f6c4a9bde8a9dd45d0275b58d760a3da9f71ce.wgsl ADDED
@@ -0,0 +1,17 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 384u;
8
+ if (i >= 384u) { return; }
9
+ let batch = i / 128u;
10
+ let row = (i / 4u) % 32u;
11
+ let col = i % 4u;
12
+ var acc = 0.0;
13
+ for (var p = 0u; p < 1u; p++) {
14
+ acc += f32(b0[(((batch / 1u) % 3u) * 32u) + row * 1u + p]) * f32(b1[(((batch / 1u) % 3u) * 4u) + p * 4u + col]);
15
+ }
16
+ out[i] = f32(acc);
17
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a896c93b04922c0f78f1c77fbaabbc79eb71776223748bb32a6d5339e32029e8.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<i32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<i32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 256u;
7
+ if (i >= 256u) { return; }
8
+ out[i] = i32(b0[((i / 1u) % 64u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a9f0eef095f015e2099adf81552a6727c45e70fc11f5e50b777d36b6370ab986.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 8192u;
7
+ if (i >= 8192u) { return; }
8
+ out[i] = f32(b0[((i / 8192u) % 1u) * 24576u + ((i / 2048u) % 4u) * 6144u + (((i / 1u) % 2048u) * 1u + 2048u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a9f6bc361f1ea31f79f0ef6164872eeb04896a45338871a2d779789b47b1920b.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 131072u;
8
+ if (i >= 131072u) { return; }
9
+ out[i] = f32(f32(b0[i]) * f32(b1[i]));
10
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/a9fa5db20fa8de1de429c6850c3e40ff8180bcdb59373919b84d2438b3b17352.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<i32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<i32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 64u;
7
+ if (i >= 48u) { return; }
8
+ out[i] = i32(b0[(((i / 16u) % 3u) * 1u + 1u) * 16u + ((i / 16u) % 1u) * 16u + ((i / 1u) % 16u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/aa227ae49fab3f75d86e02e6f390be785eabfeafe70e77a5c7eb4ad51895e382.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 32768u;
8
+ if (i >= 32768u) { return; }
9
+ out[i] = f32(f32(b0[i]) * f32(b1[i]));
10
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/aac6395f039e396767ee10d0ad95d3d6206ad4a5c1e5913dfd5c3cfeaf227510.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 65536u;
7
+ if (i >= 65536u) { return; }
8
+ out[i] = f32(b0[i]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/abdb17daf29310a7828a55a1efeebaa65b6553e236fc20702e8c712e0a46dd5b.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 2048u;
7
+ if (i >= 2048u) { return; }
8
+ let x = f32(b0[i]);
9
+ out[i] = f32(1.0 / (1.0 + exp(-x)));
10
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/ac173d038b5767711ad9f7cafbceb9ca3ec84aaeb3d07a3037e7cd0f83fead2c.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 131072u;
7
+ if (i >= 131072u) { return; }
8
+ out[i] = f32(b0[((i / 131072u) % 1u) * 393216u + ((i / 2048u) % 64u) * 6144u + (((i / 1u) % 2048u) * 1u + 0u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/ad72f60d58133272228319f1c066e7cd43ac28573ddeeb7a6f69be01af341f57.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 512u;
7
+ if (i >= 512u) { return; }
8
+ out[i] = f32(b0[((i / 512u) % 1u) * 2048u + ((i / 256u) % 2u) * 1024u + ((i / 64u) % 4u) * 256u + (((i / 1u) % 64u) * 1u + 0u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/adf1e58a124e1301625024ae142dca5fb1dd50799fa94b52eaeffb82ea99852c.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 4096u;
7
+ if (i >= 4096u) { return; }
8
+ let x = f32(b0[i]);
9
+ out[i] = f32(-x);
10
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/adfbba60ade7143599d7de0faa4904afc1400ccc3cf84602f22a2fb004d0398e.wgsl ADDED
@@ -0,0 +1,71 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<u32>;
3
+ @group(0) @binding(2) var<storage, read> b2: array<f32>;
4
+ @group(0) @binding(3) var<storage, read_write> out: array<f32>;
5
+ fn dequant_1(index: u32) -> f32 {
6
+ let row = index / 1024u;
7
+ let col = index % 1024u;
8
+ let word = b1[row * 256u + col / 4u];
9
+ let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
10
+ return f32(code) * b2[row * 32u + col / 32u];
11
+ }
12
+
13
+ var<workgroup> tile_a: array<array<f32, 16>, 16>;
14
+ var<workgroup> tile_b: array<array<f32, 64>, 16>;
15
+ @compute @workgroup_size(16, 16)
16
+ fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
17
+ let lane = local.y * 16u + local.x;
18
+ let batch = group.z;
19
+ let tile_n = group.x * 64u;
20
+ let tile_row = group.y * 16u;
21
+ var acc: array<array<f32, 4>, 1>;
22
+ for (var k0 = 0u; k0 < 1024u; k0 += 16u) {
23
+ for (var e = 0u; e < 1u; e++) {
24
+ let flat = lane + e * 256u;
25
+ let m_local = flat / 16u;
26
+ let row = tile_row + m_local;
27
+ let col = k0 + flat % 16u;
28
+ var value = 0.0;
29
+ if (row < 4u && col < 1024u) { value = f32(f32(b0[(batch * 4u + row) * 1024u + col])); }
30
+ tile_a[m_local][flat % 16u] = value;
31
+ }
32
+ if (lane < 256u) {
33
+ let n_local = lane / 4u;
34
+ let word_local = lane % 4u;
35
+ let n_index = tile_n + n_local;
36
+ let first = k0 + word_local * 4u;
37
+ var word = 0u;
38
+ var scale = 0.0;
39
+ if (n_index < 131072u && first < 1024u) {
40
+ word = b1[n_index * 256u + first / 4u];
41
+ scale = f32(b2[n_index * 32u + first / 32u]);
42
+ }
43
+ for (var e = 0u; e < 4u; e++) {
44
+ let code = i32((word >> (e * 8u)) & 255u) - 128;
45
+ var value = 0.0;
46
+ if (first + e < 1024u) { value = f32(code) * scale; }
47
+ tile_b[word_local * 4u + e][n_local] = value;
48
+ }
49
+ }
50
+ workgroupBarrier();
51
+ for (var kk = 0u; kk < 16u; kk++) {
52
+ var b_values: array<f32, 4>;
53
+ for (var c = 0u; c < 4u; c++) { b_values[c] = tile_b[kk][local.x * 4u + c]; }
54
+ for (var r = 0u; r < 1u; r++) {
55
+ let a_value = tile_a[local.y * 1u + r][kk];
56
+ for (var c = 0u; c < 4u; c++) { acc[r][c] += a_value * b_values[c]; }
57
+ }
58
+ }
59
+ workgroupBarrier();
60
+ }
61
+ for (var r = 0u; r < 1u; r++) {
62
+ let out_row = tile_row + local.y * 1u + r;
63
+ for (var c = 0u; c < 4u; c++) {
64
+ let out_col = tile_n + local.x * 4u + c;
65
+ if (out_row < 4u && out_col < 131072u) {
66
+ let i = (batch * 4u + out_row) * 131072u + out_col;
67
+ out[i] = f32(acc[r][c] + 0.0);
68
+ }
69
+ }
70
+ }
71
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/ae69e698c92571766e36ad8c078b13e42bde56745d45698daaadaca587b7279c.wgsl ADDED
@@ -0,0 +1,11 @@
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 128u;
8
+ if (i >= 128u) { return; }
9
+ let coord = (i / 1u) % 32u;
10
+ if (coord >= 2u && coord < 32u && (coord - 2u) % 3u == 0u) { out[i] = f32(b0[((i / 128u) % 1u) * 40u + ((i / 32u) % 4u) * 10u + ((((i / 1u) % 32u) - 2u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
11
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b02b0be460d581c720afd68c8b722e8da47e2ee432b5d61bbdf24b6084cb2f26.wgsl ADDED
@@ -0,0 +1,11 @@
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 2048u;
8
+ if (i >= 2048u) { return; }
9
+ let x = f32(b0[i]); let s = f32(f32(x / (1.0 + exp(-x))));
10
+ out[i] = f32(s * f32(b1[i]));
11
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b06d10513b70659dd70c40974ccc40867e6c17dacc4d4f5c131f6d99c728138d.wgsl ADDED
@@ -0,0 +1,41 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<u32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read> b2: array<f32>;
4
+ @group(0) @binding(3) var<storage, read_write> out: array<f32>;
5
+ fn dequant_0(index: u32) -> f32 {
6
+ let row = index / 1024u;
7
+ let col = index % 1024u;
8
+ let word = b0[row * 256u + col / 4u];
9
+ let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
10
+ return f32(code) * b1[row * 32u + col / 32u];
11
+ }
12
+
13
+ var<workgroup> partial: array<f32, 64>;
14
+ @compute @workgroup_size(64)
15
+ fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
16
+ let i = group.x + group.y * 65535u;
17
+ if (i >= 131072u) { return; }
18
+ let lane = local.x;
19
+ let row = i % 131072u;
20
+ var acc = 0.0;
21
+ for (var w = lane; w < 256u; w += 64u) {
22
+ let word = b0[row * 256u + w];
23
+ let first = w * 4u;
24
+ var part = 0.0;
25
+ for (var e = 0u; e < 4u; e++) {
26
+ let col = first + e;
27
+ if (col < 1024u) {
28
+ let code = i32((word >> (e * 8u)) & 255u) - 128;
29
+ part += f32(code) * f32(b2[(i / 131072u) * 1024u + col]);
30
+ }
31
+ }
32
+ acc += part * f32(b1[row * 32u + first / 32u]);
33
+ }
34
+ partial[lane] = acc;
35
+ workgroupBarrier();
36
+ for (var stride = 32u; stride > 0u; stride /= 2u) {
37
+ if (lane < stride) { partial[lane] += partial[lane + stride]; }
38
+ workgroupBarrier();
39
+ }
40
+ if (lane == 0u) { out[i] = f32(partial[0] + 0.0); }
41
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b12eae76f13225c0c5d35fc9309964aba3172712799ff3bd8538bb2c1b34aa33.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 1536u;
7
+ if (i >= 1536u) { return; }
8
+ out[i] = f32(b0[((i / 512u) % 3u) * 512u + ((i / 512u) % 1u) * 512u + ((i / 32u) % 16u) * 1u + ((i / 1u) % 32u) * 16u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b199d835db70f888c7b99a3ac6e9ec56f6c47d5bba4eaf352e0a8eabed56411e.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 2048u;
8
+ if (i >= 2048u) { return; }
9
+ out[i] = f32(f32(b0[i]) * f32(b1[((i / 128u) % 16u) * 1u]));
10
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b28ea4a8345a9c9abb2d74bfd620ba5de46616434dbaffdbabf022ea71800839.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 8192u;
7
+ if (i >= 8192u) { return; }
8
+ out[i] = f32(b0[i]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b2ae1b161c8a6315f8863e47511607027dfe583aac1db4c1626371ddfa0c64e2.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 256u;
7
+ if (i >= 256u) { return; }
8
+ out[i] = f32(b0[((i / 256u) % 1u) * 512u + ((i / 128u) % 2u) * 256u + ((i / 32u) % 4u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b2b16798cbe94c919025d76b0435afe47966ffdd1ea53ec74a791c03c4709a28.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 30720u;
7
+ if (i >= 30720u) { return; }
8
+ out[i] = f32(b0[((i / 30720u) % 1u) * 49152u + ((i / 5u) % 6144u) * 8u + (((i / 1u) % 5u) * 1u + 0u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b3bb2905555188f535b6b8bc1d050c142b5fea8222d9ef193466f80aa2c54dea.wgsl ADDED
@@ -0,0 +1,11 @@
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 2048u;
8
+ if (i >= 2048u) { return; }
9
+ let coord = (i / 1u) % 32u;
10
+ if (coord >= 2u && coord < 32u && (coord - 2u) % 3u == 0u) { out[i] = f32(b0[((i / 2048u) % 1u) * 640u + ((i / 32u) % 64u) * 10u + ((((i / 1u) % 32u) - 2u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
11
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b441eef0119fbc7cc4c3757e12f57cffdeeac0bbed3285b5757f7b9a20feb669.wgsl ADDED
@@ -0,0 +1,41 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<u32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read> b2: array<f32>;
4
+ @group(0) @binding(3) var<storage, read_write> out: array<f32>;
5
+ fn dequant_0(index: u32) -> f32 {
6
+ let row = index / 1024u;
7
+ let col = index % 1024u;
8
+ let word = b0[row * 128u + col / 8u];
9
+ let code = i32((word >> ((col % 8u) * 4u)) & 15u) - 8;
10
+ return f32(code) * b1[row * 32u + col / 32u];
11
+ }
12
+
13
+ var<workgroup> partial: array<f32, 64>;
14
+ @compute @workgroup_size(64)
15
+ fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
16
+ let i = group.x + group.y * 3584u;
17
+ if (i >= 3584u) { return; }
18
+ let lane = local.x;
19
+ let row = i % 3584u;
20
+ var acc = 0.0;
21
+ for (var w = lane; w < 128u; w += 64u) {
22
+ let word = b0[row * 128u + w];
23
+ let first = w * 8u;
24
+ var part = 0.0;
25
+ for (var e = 0u; e < 8u; e++) {
26
+ let col = first + e;
27
+ if (col < 1024u) {
28
+ let code = i32((word >> (e * 4u)) & 15u) - 8;
29
+ part += f32(code) * f32(b2[(i / 3584u) * 1024u + col]);
30
+ }
31
+ }
32
+ acc += part * f32(b1[row * 32u + first / 32u]);
33
+ }
34
+ partial[lane] = acc;
35
+ workgroupBarrier();
36
+ for (var stride = 32u; stride > 0u; stride /= 2u) {
37
+ if (lane < stride) { partial[lane] += partial[lane + stride]; }
38
+ workgroupBarrier();
39
+ }
40
+ if (lane == 0u) { out[i] = f32(partial[0] + 0.0); }
41
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b57928140994202536cd4f99d2e14f5dba9754f014843673790d5d48c9cbb9ab.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 512u;
7
+ if (i >= 512u) { return; }
8
+ out[i] = f32(b0[((i / 512u) % 1u) * 512u + ((i / 256u) % 2u) * 256u + ((i / 256u) % 1u) * 512u + ((i / 1u) % 256u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b5984ff6eb9e56ef0f9930f775c1dc3c203eec15924e5fea4c2525b2ce180857.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 4096u;
7
+ if (i >= 4096u) { return; }
8
+ out[i] = f32(b0[((i / 4096u) % 1u) * 8192u + ((i / 512u) % 8u) * 1024u + ((i / 32u) % 16u) * 64u + (((i / 1u) % 32u) * 1u + 0u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/b6c3ce0defd66eef66dc3fd6eb9eab3fee12d179332698c3c764009988628b49.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 49152u;
7
+ if (i >= 49152u) { return; }
8
+ let x = f32(b0[i]);
9
+ out[i] = f32(x / (1.0 + exp(-x)));
10
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bb3e27e020bc22a0c8c21a5f9910a0a26b73ace5f68d9d49d1e4d67d53232060.wgsl ADDED
@@ -0,0 +1,71 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<u32>;
3
+ @group(0) @binding(2) var<storage, read> b2: array<f32>;
4
+ @group(0) @binding(3) var<storage, read_write> out: array<f32>;
5
+ fn dequant_1(index: u32) -> f32 {
6
+ let row = index / 2048u;
7
+ let col = index % 2048u;
8
+ let word = b1[row * 512u + col / 4u];
9
+ let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
10
+ return f32(code) * b2[row * 64u + col / 32u];
11
+ }
12
+
13
+ var<workgroup> tile_a: array<array<f32, 16>, 16>;
14
+ var<workgroup> tile_b: array<array<f32, 64>, 16>;
15
+ @compute @workgroup_size(16, 16)
16
+ fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
17
+ let lane = local.y * 16u + local.x;
18
+ let batch = group.z;
19
+ let tile_n = group.x * 64u;
20
+ let tile_row = group.y * 16u;
21
+ var acc: array<array<f32, 4>, 1>;
22
+ for (var k0 = 0u; k0 < 2048u; k0 += 16u) {
23
+ for (var e = 0u; e < 1u; e++) {
24
+ let flat = lane + e * 256u;
25
+ let m_local = flat / 16u;
26
+ let row = tile_row + m_local;
27
+ let col = k0 + flat % 16u;
28
+ var value = 0.0;
29
+ if (row < 4u && col < 2048u) { value = f32(f32(b0[(batch * 4u + row) * 2048u + col])); }
30
+ tile_a[m_local][flat % 16u] = value;
31
+ }
32
+ if (lane < 256u) {
33
+ let n_local = lane / 4u;
34
+ let word_local = lane % 4u;
35
+ let n_index = tile_n + n_local;
36
+ let first = k0 + word_local * 4u;
37
+ var word = 0u;
38
+ var scale = 0.0;
39
+ if (n_index < 1024u && first < 2048u) {
40
+ word = b1[n_index * 512u + first / 4u];
41
+ scale = f32(b2[n_index * 64u + first / 32u]);
42
+ }
43
+ for (var e = 0u; e < 4u; e++) {
44
+ let code = i32((word >> (e * 8u)) & 255u) - 128;
45
+ var value = 0.0;
46
+ if (first + e < 2048u) { value = f32(code) * scale; }
47
+ tile_b[word_local * 4u + e][n_local] = value;
48
+ }
49
+ }
50
+ workgroupBarrier();
51
+ for (var kk = 0u; kk < 16u; kk++) {
52
+ var b_values: array<f32, 4>;
53
+ for (var c = 0u; c < 4u; c++) { b_values[c] = tile_b[kk][local.x * 4u + c]; }
54
+ for (var r = 0u; r < 1u; r++) {
55
+ let a_value = tile_a[local.y * 1u + r][kk];
56
+ for (var c = 0u; c < 4u; c++) { acc[r][c] += a_value * b_values[c]; }
57
+ }
58
+ }
59
+ workgroupBarrier();
60
+ }
61
+ for (var r = 0u; r < 1u; r++) {
62
+ let out_row = tile_row + local.y * 1u + r;
63
+ for (var c = 0u; c < 4u; c++) {
64
+ let out_col = tile_n + local.x * 4u + c;
65
+ if (out_row < 4u && out_col < 1024u) {
66
+ let i = (batch * 4u + out_row) * 1024u + out_col;
67
+ out[i] = f32(acc[r][c] + 0.0);
68
+ }
69
+ }
70
+ }
71
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bb7f03e6a2277dc56084e9130086507be58871253867c0ce8ad60ee5ce4521b6.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 262144u;
7
+ if (i >= 262144u) { return; }
8
+ out[i] = f32(b0[((i / 262144u) % 1u) * 393216u + ((i / 16384u) % 16u) * 24576u + (((i / 128u) % 128u) * 1u + 0u) * 128u + ((i / 1u) % 128u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bb9029b667eeb13670d1f640cd677c6eabd7a7569686c599e3eeb9482d642b00.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 2048u;
8
+ if (i >= 2048u) { return; }
9
+ out[i] = f32(f32(b0[i]) - (f32(b1[i]) * 1.0));
10
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bcb78b92dbbfeb212c190e064bc2f9d8e79c201398cf38c9aa62df6fdba94469.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 32768u;
7
+ if (i >= 32768u) { return; }
8
+ out[i] = f32(b0[((i / 32768u) % 1u) * 32768u + ((i / 16384u) % 2u) * 256u + ((i / 256u) % 64u) * 512u + ((i / 1u) % 256u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bd70e4fc42ac55f0f52612841bcc3f9fe7df5bbb0fc76f94e56a131a92454658.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 24576u;
7
+ if (i >= 24576u) { return; }
8
+ out[i] = f32(b0[((i / 24576u) % 1u) * 32768u + ((i / 12288u) % 2u) * 16384u + ((i / 192u) % 64u) * 256u + (((i / 1u) % 192u) * 1u + 64u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bd8a785d4fbf09918f149c7eb5f34bd25f9358bf3412e27799a7c65436c71759.wgsl ADDED
@@ -0,0 +1,9 @@
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 24576u;
7
+ if (i >= 24576u) { return; }
8
+ out[i] = f32(b0[((i / 24576u) % 1u) * 417792u + ((i / 4u) % 6144u) * 68u + (((i / 1u) % 4u) * 1u + 64u) * 1u]);
9
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/bda7ebd3bcf25c31938fca1a582e82b972ab742f85f826f085417b0e6f651e60.wgsl ADDED
@@ -0,0 +1,12 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 2048u;
8
+ if (i >= 2048u) { return; }
9
+ let coord = (i / 1u) % 64u;
10
+ if (coord >= 0u && coord < 32u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 0u) * 1u + i % 1u]); }
11
+ if (coord >= 32u && coord < 64u) { out[i] = f32(b1[(i / 64u) * 32u + (coord - 32u) * 1u + i % 1u]); }
12
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/c02ac02d110c0be5d0990d20ab0d8c7a081c74f4e147f083789a203f22fc3258.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read> b1: array<f32>;
3
+ @group(0) @binding(2) var<storage, read_write> out: array<f32>;
4
+
5
+ @compute @workgroup_size(64)
6
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
7
+ let i = gid.x + gid.y * 256u;
8
+ if (i >= 256u) { return; }
9
+ out[i] = f32(f32(b0[((i / 1u) % 16u) * 1u]) * f32(b1[i]));
10
+ }
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v4/kernels/c0b2947798baa198598d75ca0e5dcd31a9fa320499665e34f7bb81d27fc15fb0.wgsl ADDED
@@ -0,0 +1,10 @@
 
 
 
 
 
 
 
 
 
 
 
1
+ @group(0) @binding(0) var<storage, read> b0: array<f32>;
2
+ @group(0) @binding(1) var<storage, read_write> out: array<f32>;
3
+
4
+ @compute @workgroup_size(64)
5
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
6
+ let i = gid.x + gid.y * 64u;
7
+ if (i >= 64u) { return; }
8
+ let x = f32(b0[i]);
9
+ out[i] = f32(cos(x));
10
+ }