Add qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5 (attention bounded by the used context)
Browse filesThis view is limited to 50 files because it contains too many changes. See raw diff
- .gitattributes +1 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/entrypoints/prefill-16.json +0 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/entrypoints/prefill-4.json +0 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/entrypoints/prefill-64.json +0 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/graph.json +3 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/004c06d76ce77979b63fe4450df54986355fd10cd2ef62adae86780e49e4a4a7.wgsl +10 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0112d93bb1aa347d8887ce1b95610a1a21fc8450ff8f59bcd37747fa6bf6fa52.wgsl +10 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0151a813c4a92adeb44f0e3b6108c77b29804118b588c9025805f724fce839b4.wgsl +24 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0153b6b641515b31ea2980613978c809e225435001b48fbec69afa520251a7af.wgsl +71 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/026b53c4ecb0527a45fbd6020a725ad4d0ea2ea28c8cfa0a56e19cb0c1ad08ea.wgsl +12 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0337d8a473b6e485bd383533df7d373dbbfe7a5f3a982424004560b7213fbaa3.wgsl +11 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/04f26661cc9ac73bb7f39f66b76c4be31ee8774e094eeb41ed7d677e1b1d3fa9.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/04f5cb70141e507d48bd44a76bfdf264a404a706e8a012502dce5a22ee75bab0.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/053386c9a33cc1c62378ebe3b3614c6b9ee662d6004a1f27bc616653e62cc24b.wgsl +71 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0554a443ecf6ead32fd2d58178b9ce3c592c89a09b3f63fa509ee947323bb497.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/06eecde176beab9ff488ad533f2f4513e09bc49a30c4358f9648076f619f9992.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0725908abf78c022f86deb338a06779f6f1346ebfe54ff5a43ee6b9eeb93c972.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/07cfabe2e2c2feff22a06c8296600e5ca21efb7444419f0a05e5ff6f48e0292d.wgsl +10 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0835fa80eaadb7c0e6d677e1fc622ad5dcedabb431e5d312f3b7aeb2252a7c7c.wgsl +8 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/08390b1553d11820d1d96b10add657b92acb8d26cd9f11acab8f16777ecfc10a.wgsl +21 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/09ea34c2533b929bd309e2bcbcf6d27fa02c731d82276ba6425db85961ea121b.wgsl +10 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0b8815b1efe176d55ab08530c0a9e1913cf6c9ef3e773c574843573d237cf196.wgsl +10 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0bbd6a5e356c1205a61dc3580a3ff68d88b98a952a5692d66b56ded171d9fc56.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0d1edb63bf4dc212d55f8a4c79a50e04c0d822bd9f6e193003e2071ba0d57d69.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0d3892883d1532d173a7dbe85d9ab1daa87f0d9b4923a535b8d2003f24f9500c.wgsl +11 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0ef8fd3c2fe267549e7de4eb9a074d2ada8e80c91ce2f9c9c4d6b24d17742821.wgsl +71 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0f943f386d2f37a8e56b4285b26e1c093c2f2fa2a5ce81cc43c44cc9bb0137f5.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0fbc9484ef8ca703ced4ffd123f1ae733c44556b6cb160ee803b820ff5ec867d.wgsl +10 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/11a92e9cba6bfc381c98c5127db137639d6c87f1cb220459880ead039dc81256.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/121285fc94b4a9f873c60a1d6e0de9d764286607718405fd18288e7ff419c483.wgsl +11 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1255847146b31732939e573cdb26c483358c1ce07f987bd8c3d3d69ae8c7e141.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/13237b557466d2fef5bfbc74e7ac77219b9821a53bc0843dbb5d14db8b04d414.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1326034f3abab37629cd88f869d1b0bafcde724f85c07904b54feaaf71d89c62.wgsl +44 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/13dc73935b9b3b978232043e13babee3877bd53ca572be440de0590dacce8b17.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/140ebe12d31f1420a8a159b91ae1773390190f496e06e0e5d7b5351f38245068.wgsl +96 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/149c3fdb542967fc038109de3f2b7b2fbee1a6f594baffe28f6f16bafb84d355.wgsl +24 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1545025e0a625ccc73b42ae112c6f9d78a1504c10b759979ae2247984f25f056.wgsl +11 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/179d5bef469ee5199c2a198580bc5876a1e3d17a7a545c045152d59b246b869f.wgsl +43 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/17ec6bdaf98edf7657b74a25036c2bec806ecbb4280f33fff6afd2e0de5ed377.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/18a0b02a12ebf73203283d267342582d54ba411aeee9cbe21ad89ae890e4cfea.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/18a66ec0875c456c9defd45d15e6ef0c148e64a1618e2825c7f96eadbbe8303a.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/18b1cef2bae996f56664d9215180e3bfc77b4e514db52fdf9b6c58be51b53edb.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/19564c145248f065669be23f5daeeda96e025ee4389bef80983001e83fddb31d.wgsl +11 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1a73b6fceda28f40bba4ae27a986ffa39842db94fa5903ea5bf689a7e854d4f3.wgsl +11 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1ad205080cae29142a3ad87343a0cf3ac4945455398293946110d421ce6d515e.wgsl +10 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1ada8fd10cb41796bc524ec61c2001d2a3ee08e9c062d35d14066088ace72436.wgsl +43 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1b9f4e330debc7235cfaf8f51d00fec575e09aa6116a45ea001e10ff5f09eb8c.wgsl +17 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1c441e95ec3ce8fee9348560af48191844baf402715f575940a0d4eb394c82d2.wgsl +71 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1d0cbb068fe17123e646934bfa1259f19e5db04e45d72a157d3ea4fa48c81aa0.wgsl +9 -0
- qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1d7df24388934e162899442d3cc6c6576791aac9f1a5dff86766d840f08fd46e.wgsl +11 -0
.gitattributes
CHANGED
|
@@ -74,3 +74,4 @@ qwen35-2b-fp32-int8-g32-exclude-l6-l9-l10up-l15o-home-token-major-v2-v4/tokenize
|
|
| 74 |
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2-v4/graph.json filter=lfs diff=lfs merge=lfs -text
|
| 75 |
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2-v4/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
|
| 76 |
qwen35-08b-fp32-8k-token-major-home-v5/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
|
|
|
|
|
|
| 74 |
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2-v4/graph.json filter=lfs diff=lfs merge=lfs -text
|
| 75 |
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2-v4/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
|
| 76 |
qwen35-08b-fp32-8k-token-major-home-v5/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
|
| 77 |
+
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/graph.json filter=lfs diff=lfs merge=lfs -text
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/entrypoints/prefill-16.json
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/entrypoints/prefill-4.json
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/entrypoints/prefill-64.json
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/graph.json
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:32d7ed995ec990c9ef6751dd8c2dd5824068d5816a4c070f43b8f8f74bdb59fe
|
| 3 |
+
size 10602779
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/004c06d76ce77979b63fe4450df54986355fd10cd2ef62adae86780e49e4a4a7.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 * 16384u;
|
| 8 |
+
if (i >= 16384u) { 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-v5/kernels/0112d93bb1aa347d8887ce1b95610a1a21fc8450ff8f59bcd37747fa6bf6fa52.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-v5/kernels/0151a813c4a92adeb44f0e3b6108c77b29804118b588c9025805f724fce839b4.wgsl
ADDED
|
@@ -0,0 +1,24 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 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 |
+
var<workgroup> factor: f32;
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 8 |
+
let row = group.x + group.y * 128u;
|
| 9 |
+
if (row >= 128u) { return; }
|
| 10 |
+
let lane = local.x;
|
| 11 |
+
if (lane == 0u) {
|
| 12 |
+
var total = 0.0;
|
| 13 |
+
for (var j = 0u; j < 256u; j++) {
|
| 14 |
+
let v = f32(b0[row * 256u + j]);
|
| 15 |
+
total += v * v;
|
| 16 |
+
}
|
| 17 |
+
factor = inverseSqrt(total / 256.0 + 1e-06);
|
| 18 |
+
}
|
| 19 |
+
workgroupBarrier();
|
| 20 |
+
for (var p = lane; p < 256u; p += 64u) {
|
| 21 |
+
let i = row * 256u + p;
|
| 22 |
+
out[i] = f32(f32(b0[row * 256u + p]) * factor * (f32(b1[p]) + 1.0));
|
| 23 |
+
}
|
| 24 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0153b6b641515b31ea2980613978c809e225435001b48fbec69afa520251a7af.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 < 117248u && 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 < 117248u) {
|
| 66 |
+
let i = (batch * 4u + out_row) * 117248u + 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-v5/kernels/026b53c4ecb0527a45fbd6020a725ad4d0ea2ea28c8cfa0a56e19cb0c1ad08ea.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 * 3973120u;
|
| 8 |
+
if (i >= 3973120u) { 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-v5/kernels/0337d8a473b6e485bd383533df7d373dbbfe7a5f3a982424004560b7213fbaa3.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_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 coord = (i / 1u) % 64u;
|
| 9 |
+
if (coord >= 0u && coord < 32u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 0u) * 1u + i % 1u]); }
|
| 10 |
+
if (coord >= 32u && coord < 64u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 32u) * 1u + i % 1u]); }
|
| 11 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/04f26661cc9ac73bb7f39f66b76c4be31ee8774e094eeb41ed7d677e1b1d3fa9.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 * 64u;
|
| 7 |
+
if (i >= 10u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 10u) % 1u) * 32u + ((i / 10u) % 1u) * 32u + (((i / 1u) % 10u) * 3u + 2u) * 1u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/04f5cb70141e507d48bd44a76bfdf264a404a706e8a012502dce5a22ee75bab0.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) * 24576u + ((i / 4u) % 6144u) * 1u + ((i / 1u) % 4u) * 6144u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/053386c9a33cc1c62378ebe3b3614c6b9ee662d6004a1f27bc616653e62cc24b.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>, 64>;
|
| 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 * 64u;
|
| 21 |
+
var acc: array<array<f32, 4>, 4>;
|
| 22 |
+
for (var k0 = 0u; k0 < 1024u; k0 += 16u) {
|
| 23 |
+
for (var e = 0u; e < 4u; 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 < 64u && col < 1024u) { value = f32(f32(b0[(batch * 64u + 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 < 2048u && 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 < 4u; r++) {
|
| 55 |
+
let a_value = tile_a[local.y * 4u + 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 < 4u; r++) {
|
| 62 |
+
let out_row = tile_row + local.y * 4u + r;
|
| 63 |
+
for (var c = 0u; c < 4u; c++) {
|
| 64 |
+
let out_col = tile_n + local.x * 4u + c;
|
| 65 |
+
if (out_row < 64u && out_col < 2048u) {
|
| 66 |
+
let i = (batch * 64u + out_row) * 2048u + 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-v5/kernels/0554a443ecf6ead32fd2d58178b9ce3c592c89a09b3f63fa509ee947323bb497.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 >= 128u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 128u) % 1u) * 512u + ((i / 64u) % 2u) * 256u + ((i / 64u) % 1u) * 256u + (((i / 1u) % 64u) * 1u + 0u) * 1u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/06eecde176beab9ff488ad533f2f4513e09bc49a30c4358f9648076f619f9992.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 * 6144u;
|
| 7 |
+
if (i >= 6144u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 2048u) % 3u) * 2048u + ((i / 2048u) % 1u) * 2048u + ((i / 64u) % 32u) * 1u + ((i / 1u) % 64u) * 32u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0725908abf78c022f86deb338a06779f6f1346ebfe54ff5a43ee6b9eeb93c972.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 * 64u;
|
| 7 |
+
if (i >= 44u) { return; }
|
| 8 |
+
out[i] = f32(b0[i]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/07cfabe2e2c2feff22a06c8296600e5ca21efb7444419f0a05e5ff6f48e0292d.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(-x);
|
| 10 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0835fa80eaadb7c0e6d677e1fc622ad5dcedabb431e5d312f3b7aeb2252a7c7c.wgsl
ADDED
|
@@ -0,0 +1,8 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read_write> out: array<f32>;
|
| 2 |
+
|
| 3 |
+
@compute @workgroup_size(64)
|
| 4 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 5 |
+
let i = gid.x + gid.y * 2048u;
|
| 6 |
+
if (i >= 2048u) { return; }
|
| 7 |
+
out[i] = f32(0.0);
|
| 8 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/08390b1553d11820d1d96b10add657b92acb8d26cd9f11acab8f16777ecfc10a.wgsl
ADDED
|
@@ -0,0 +1,21 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 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 * 436224u;
|
| 8 |
+
if (i >= 436224u) { return; }
|
| 9 |
+
let batch = i / 436224u;
|
| 10 |
+
let channel = (i / 71u) % 6144u;
|
| 11 |
+
let group = channel / 1u;
|
| 12 |
+
var acc = 0.0;
|
| 13 |
+
for (var c = 0u; c < 1u; c++) {
|
| 14 |
+
for (var p = 0u; p < 4u; p++) {
|
| 15 |
+
let pos0 = i32((i / 1u) % 71u) * 1 + i32((p / 1u) % 4u) * 1 - 3;
|
| 16 |
+
if (pos0 >= 0 && pos0 < 68) { acc += f32(b0[batch * 417792u + (group * 1u + c) * 68u + u32(pos0) * 1u]) * f32(b1[(channel * 1u + c) * 4u + p]); }
|
| 17 |
+
}
|
| 18 |
+
}
|
| 19 |
+
out[i] = f32(acc + 0.0);
|
| 20 |
+
|
| 21 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/09ea34c2533b929bd309e2bcbcf6d27fa02c731d82276ba6425db85961ea121b.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 * 256u;
|
| 7 |
+
if (i >= 256u) { return; }
|
| 8 |
+
let x = f32(b0[i]);
|
| 9 |
+
out[i] = f32(inverseSqrt(x));
|
| 10 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0b8815b1efe176d55ab08530c0a9e1913cf6c9ef3e773c574843573d237cf196.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 * 256u;
|
| 7 |
+
if (i >= 256u) { return; }
|
| 8 |
+
let x = f32(b0[i]);
|
| 9 |
+
out[i] = f32(sin(x));
|
| 10 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0bbd6a5e356c1205a61dc3580a3ff68d88b98a952a5692d66b56ded171d9fc56.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 * 1024u;
|
| 7 |
+
if (i >= 1024u) { return; }
|
| 8 |
+
out[i] = f32(f32(b0[i]) + (1e-06 * 1.0));
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0d1edb63bf4dc212d55f8a4c79a50e04c0d822bd9f6e193003e2071ba0d57d69.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) * 16384u + ((i / 2048u) % 4u) * 4096u + ((i / 256u) % 8u) * 512u + (((i / 1u) % 256u) * 1u + 256u) * 1u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0d3892883d1532d173a7dbe85d9ab1daa87f0d9b4923a535b8d2003f24f9500c.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_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 |
+
let coord = (i / 1u) % 64u;
|
| 9 |
+
if (coord >= 0u && coord < 32u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 0u) * 1u + i % 1u]); }
|
| 10 |
+
if (coord >= 32u && coord < 64u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 32u) * 1u + i % 1u]); }
|
| 11 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0ef8fd3c2fe267549e7de4eb9a074d2ada8e80c91ce2f9c9c4d6b24d17742821.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>, 64>;
|
| 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 * 64u;
|
| 21 |
+
var acc: array<array<f32, 4>, 4>;
|
| 22 |
+
for (var k0 = 0u; k0 < 2048u; k0 += 16u) {
|
| 23 |
+
for (var e = 0u; e < 4u; 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 < 64u && col < 2048u) { value = f32(f32(b0[(batch * 64u + 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 < 4u; r++) {
|
| 55 |
+
let a_value = tile_a[local.y * 4u + 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 < 4u; r++) {
|
| 62 |
+
let out_row = tile_row + local.y * 4u + r;
|
| 63 |
+
for (var c = 0u; c < 4u; c++) {
|
| 64 |
+
let out_col = tile_n + local.x * 4u + c;
|
| 65 |
+
if (out_row < 64u && out_col < 1024u) {
|
| 66 |
+
let i = (batch * 64u + 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-v5/kernels/0f943f386d2f37a8e56b4285b26e1c093c2f2fa2a5ce81cc43c44cc9bb0137f5.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 * 384u;
|
| 7 |
+
if (i >= 384u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 128u) % 3u) * 128u + ((i / 128u) % 1u) * 128u + ((i / 32u) % 4u) * 1u + ((i / 1u) % 32u) * 4u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/0fbc9484ef8ca703ced4ffd123f1ae733c44556b6cb160ee803b820ff5ec867d.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 / 64u) % 16u) * 64u + ((i / 1u) % 64u) * 1u]));
|
| 10 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/11a92e9cba6bfc381c98c5127db137639d6c87f1cb220459880ead039dc81256.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 >= 12u) { return; }
|
| 8 |
+
out[i] = i32(b0[(((i / 4u) % 3u) * 1u + 1u) * 4u + ((i / 4u) % 1u) * 4u + ((i / 1u) % 4u) * 1u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/121285fc94b4a9f873c60a1d6e0de9d764286607718405fd18288e7ff419c483.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_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 |
+
let x = f32(b0[i]);
|
| 9 |
+
if (x * 1.0 > 20.0) { out[i] = f32(x); return; }
|
| 10 |
+
out[i] = f32(log(1.0 + exp(x * 1.0)) / 1.0);
|
| 11 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1255847146b31732939e573cdb26c483358c1ce07f987bd8c3d3d69ae8c7e141.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) * 30720u + ((i / 4u) % 6144u) * 5u + (((i / 1u) % 4u) * 1u + 1u) * 1u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/13237b557466d2fef5bfbc74e7ac77219b9821a53bc0843dbb5d14db8b04d414.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 * 2048u;
|
| 7 |
+
if (i >= 2048u) { return; }
|
| 8 |
+
out[i] = f32(f32(b0[i]) * f32(b0[i]));
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1326034f3abab37629cd88f869d1b0bafcde724f85c07904b54feaaf71d89c62.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 < 4u; token++) {
|
| 20 |
+
recurrent *= exp(f32(b3[(batch_index * 4u + token) * 16u + head]));
|
| 21 |
+
partial[lane] = 0.0;
|
| 22 |
+
if (lane < 128u) { partial[lane] = recurrent * f32(b1[((batch_index * 4u + 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 * 4u + token) * 16u + head) * 128u + column]) - partial[0]) * f32(b4[(batch_index * 4u + token) * 16u + head]); }
|
| 29 |
+
workgroupBarrier();
|
| 30 |
+
if (lane < 128u) { recurrent += f32(b1[((batch_index * 4u + token) * 16u + head) * 128u + lane]) * delta; }
|
| 31 |
+
partial[lane] = 0.0;
|
| 32 |
+
if (lane < 128u) {
|
| 33 |
+
partial[lane] = recurrent * f32(b0[((batch_index * 4u + 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) * 132u + 128u + token) * 128u + column] = f32(partial[0]); }
|
| 41 |
+
workgroupBarrier();
|
| 42 |
+
}
|
| 43 |
+
if (lane < 128u) { out[((batch_index * 16u + head) * 132u + lane) * 128u + column] = f32(recurrent); }
|
| 44 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/13dc73935b9b3b978232043e13babee3877bd53ca572be440de0590dacce8b17.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 * 64u;
|
| 7 |
+
if (i >= 32u) { return; }
|
| 8 |
+
out[i] = f32(b0[0u * 32u + ((i / 32u) % 1u) * 32u + ((i / 32u) % 1u) * 32u + ((i / 1u) % 32u) * 1u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/140ebe12d31f1420a8a159b91ae1773390190f496e06e0e5d7b5351f38245068.wgsl
ADDED
|
@@ -0,0 +1,96 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 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 |
+
var<workgroup> limit_shared: u32;
|
| 11 |
+
@compute @workgroup_size(64)
|
| 12 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 13 |
+
let lane = local.x;
|
| 14 |
+
let head = group.x / 64u;
|
| 15 |
+
let row = group.x % 64u;
|
| 16 |
+
let kv_head = head / 4u;
|
| 17 |
+
for (var d = lane; d < 256u; d += 64u) { q_row[d] = f32(b0[(head * 64u + row) * 256u + d]); }
|
| 18 |
+
// Bound the key loop by the last unmasked position so work follows the used context.
|
| 19 |
+
var last = 0.0;
|
| 20 |
+
for (var j = lane; j < 4096u; j += 64u) {
|
| 21 |
+
if (f32(b3[row * 4096u + j]) > -1000.0) { last = max(last, f32(j + 1u)); }
|
| 22 |
+
}
|
| 23 |
+
partial[lane] = last;
|
| 24 |
+
workgroupBarrier();
|
| 25 |
+
for (var stride = 32u; stride > 0u; stride /= 2u) {
|
| 26 |
+
if (lane < stride) { partial[lane] = max(partial[lane], partial[lane + stride]); }
|
| 27 |
+
workgroupBarrier();
|
| 28 |
+
}
|
| 29 |
+
if (lane == 0u) { limit_shared = u32(partial[0]); }
|
| 30 |
+
let limit = workgroupUniformLoad(&limit_shared);
|
| 31 |
+
var running_max = -3.4028234663852886e38;
|
| 32 |
+
var running_sum = 0.0;
|
| 33 |
+
var acc: array<f32, 4>;
|
| 34 |
+
for (var start = 0u; start < limit; start += 1024u) {
|
| 35 |
+
let count = min(1024u, limit - start);
|
| 36 |
+
var tile_max = -3.4028234663852886e38;
|
| 37 |
+
for (var t = lane; t < count; t += 64u) {
|
| 38 |
+
let j = start + t;
|
| 39 |
+
var score = -3.4028234663852886e38;
|
| 40 |
+
if (f32(b3[row * 4096u + j]) > -1000.0) {
|
| 41 |
+
var dot = 0.0;
|
| 42 |
+
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]); }
|
| 43 |
+
score = dot * 0.0625 + f32(b3[row * 4096u + j]);
|
| 44 |
+
}
|
| 45 |
+
scores[t] = score;
|
| 46 |
+
tile_max = max(tile_max, score);
|
| 47 |
+
}
|
| 48 |
+
partial[lane] = tile_max;
|
| 49 |
+
workgroupBarrier();
|
| 50 |
+
for (var stride = 32u; stride > 0u; stride /= 2u) {
|
| 51 |
+
if (lane < stride) { partial[lane] = max(partial[lane], partial[lane + stride]); }
|
| 52 |
+
workgroupBarrier();
|
| 53 |
+
}
|
| 54 |
+
let next_max = max(running_max, partial[0]);
|
| 55 |
+
workgroupBarrier();
|
| 56 |
+
var tile_sum = 0.0;
|
| 57 |
+
for (var t = lane; t < count; t += 64u) {
|
| 58 |
+
var p = 0.0;
|
| 59 |
+
if (scores[t] > -3.0e38) { p = exp(scores[t] - next_max); }
|
| 60 |
+
scores[t] = p;
|
| 61 |
+
tile_sum += p;
|
| 62 |
+
}
|
| 63 |
+
partial[lane] = tile_sum;
|
| 64 |
+
workgroupBarrier();
|
| 65 |
+
for (var stride = 32u; stride > 0u; stride /= 2u) {
|
| 66 |
+
if (lane < stride) { partial[lane] += partial[lane + stride]; }
|
| 67 |
+
workgroupBarrier();
|
| 68 |
+
}
|
| 69 |
+
var correction = 0.0;
|
| 70 |
+
if (running_max > -3.0e38) { correction = exp(running_max - next_max); }
|
| 71 |
+
running_sum = running_sum * correction + partial[0];
|
| 72 |
+
running_max = next_max;
|
| 73 |
+
for (var c = 0u; c < 4u; c++) {
|
| 74 |
+
let d = lane + c * 64u;
|
| 75 |
+
var total = acc[c] * correction;
|
| 76 |
+
if (d < 256u) {
|
| 77 |
+
for (var t = 0u; t < count; t++) {
|
| 78 |
+
let p = scores[t];
|
| 79 |
+
if (p != 0.0) {
|
| 80 |
+
let j = start + t;
|
| 81 |
+
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]);
|
| 82 |
+
}
|
| 83 |
+
}
|
| 84 |
+
}
|
| 85 |
+
acc[c] = total;
|
| 86 |
+
}
|
| 87 |
+
workgroupBarrier();
|
| 88 |
+
}
|
| 89 |
+
for (var c = 0u; c < 4u; c++) {
|
| 90 |
+
let d = lane + c * 64u;
|
| 91 |
+
if (d < 256u) {
|
| 92 |
+
let i = (head * 64u + row) * 256u + d;
|
| 93 |
+
out[i] = f32(acc[c] / running_sum);
|
| 94 |
+
}
|
| 95 |
+
}
|
| 96 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/149c3fdb542967fc038109de3f2b7b2fbee1a6f594baffe28f6f16bafb84d355.wgsl
ADDED
|
@@ -0,0 +1,24 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 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 |
+
var<workgroup> factor: f32;
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 8 |
+
let row = group.x + group.y * 4u;
|
| 9 |
+
if (row >= 4u) { return; }
|
| 10 |
+
let lane = local.x;
|
| 11 |
+
if (lane == 0u) {
|
| 12 |
+
var total = 0.0;
|
| 13 |
+
for (var j = 0u; j < 1024u; j++) {
|
| 14 |
+
let v = f32(b0[row * 1024u + j]);
|
| 15 |
+
total += v * v;
|
| 16 |
+
}
|
| 17 |
+
factor = inverseSqrt(total / 1024.0 + 1e-06);
|
| 18 |
+
}
|
| 19 |
+
workgroupBarrier();
|
| 20 |
+
for (var p = lane; p < 1024u; p += 64u) {
|
| 21 |
+
let i = row * 1024u + p;
|
| 22 |
+
out[i] = f32(f32(b0[row * 1024u + p]) * factor * (f32(b1[p]) + 1.0));
|
| 23 |
+
}
|
| 24 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1545025e0a625ccc73b42ae112c6f9d78a1504c10b759979ae2247984f25f056.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 >= 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]); }
|
| 11 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/179d5bef469ee5199c2a198580bc5876a1e3d17a7a545c045152d59b246b869f.wgsl
ADDED
|
@@ -0,0 +1,43 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 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 slot = local.x / 16u;
|
| 17 |
+
let lane = local.x % 16u;
|
| 18 |
+
let i = (group.x + group.y * 1536u) * 4u + slot;
|
| 19 |
+
var acc = 0.0;
|
| 20 |
+
if (i < 6144u) {
|
| 21 |
+
let row = i % 6144u;
|
| 22 |
+
for (var w = lane; w < 256u; w += 16u) {
|
| 23 |
+
let word = b0[row * 256u + w];
|
| 24 |
+
let first = w * 4u;
|
| 25 |
+
var part = 0.0;
|
| 26 |
+
for (var e = 0u; e < 4u; e++) {
|
| 27 |
+
let col = first + e;
|
| 28 |
+
if (col < 1024u) {
|
| 29 |
+
let code = i32((word >> (e * 8u)) & 255u) - 128;
|
| 30 |
+
part += f32(code) * f32(b2[(i / 6144u) * 1024u + col]);
|
| 31 |
+
}
|
| 32 |
+
}
|
| 33 |
+
acc += part * f32(b1[row * 32u + first / 32u]);
|
| 34 |
+
}
|
| 35 |
+
}
|
| 36 |
+
partial[local.x] = acc;
|
| 37 |
+
workgroupBarrier();
|
| 38 |
+
for (var stride = 8u; stride > 0u; stride /= 2u) {
|
| 39 |
+
if (lane < stride) { partial[local.x] += partial[local.x + stride]; }
|
| 40 |
+
workgroupBarrier();
|
| 41 |
+
}
|
| 42 |
+
if (lane == 0u && i < 6144u) { out[i] = f32(partial[local.x] + 0.0); }
|
| 43 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/17ec6bdaf98edf7657b74a25036c2bec806ecbb4280f33fff6afd2e0de5ed377.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 * 2048u;
|
| 7 |
+
if (i >= 2048u) { return; }
|
| 8 |
+
out[i] = f32(f32(b0[i]) * 0.08838834764831843);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/18a0b02a12ebf73203283d267342582d54ba411aeee9cbe21ad89ae890e4cfea.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) * 8192u + ((i / 1024u) % 8u) * 256u + ((i / 256u) % 4u) * 2048u + ((i / 1u) % 256u) * 1u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/18a66ec0875c456c9defd45d15e6ef0c148e64a1618e2825c7f96eadbbe8303a.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 >= 16u) { return; }
|
| 8 |
+
out[i] = i32(b0[i]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/18b1cef2bae996f56664d9215180e3bfc77b4e514db52fdf9b6c58be51b53edb.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 * 64u;
|
| 7 |
+
if (i >= 44u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 44u) % 1u) * 128u + ((i / 11u) % 4u) * 32u + (((i / 1u) % 11u) * 3u + 1u) * 1u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/19564c145248f065669be23f5daeeda96e025ee4389bef80983001e83fddb31d.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_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 |
+
if (x * 1.0 > 20.0) { out[i] = f32(x); return; }
|
| 10 |
+
out[i] = f32(log(1.0 + exp(x * 1.0)) / 1.0);
|
| 11 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1a73b6fceda28f40bba4ae27a986ffa39842db94fa5903ea5bf689a7e854d4f3.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_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 coord = (i / 1u) % 64u;
|
| 9 |
+
if (coord >= 0u && coord < 32u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 0u) * 1u + i % 1u]); }
|
| 10 |
+
if (coord >= 32u && coord < 64u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 32u) * 1u + i % 1u]); }
|
| 11 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1ad205080cae29142a3ad87343a0cf3ac4945455398293946110d421ce6d515e.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 * 128u;
|
| 8 |
+
if (i >= 128u) { 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-v5/kernels/1ada8fd10cb41796bc524ec61c2001d2a3ee08e9c062d35d14066088ace72436.wgsl
ADDED
|
@@ -0,0 +1,43 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 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 slot = local.x / 16u;
|
| 17 |
+
let lane = local.x % 16u;
|
| 18 |
+
let i = (group.x + group.y * 128u) * 4u + slot;
|
| 19 |
+
var acc = 0.0;
|
| 20 |
+
if (i < 512u) {
|
| 21 |
+
let row = i % 512u;
|
| 22 |
+
for (var w = lane; w < 256u; w += 16u) {
|
| 23 |
+
let word = b0[row * 256u + w];
|
| 24 |
+
let first = w * 4u;
|
| 25 |
+
var part = 0.0;
|
| 26 |
+
for (var e = 0u; e < 4u; e++) {
|
| 27 |
+
let col = first + e;
|
| 28 |
+
if (col < 1024u) {
|
| 29 |
+
let code = i32((word >> (e * 8u)) & 255u) - 128;
|
| 30 |
+
part += f32(code) * f32(b2[(i / 512u) * 1024u + col]);
|
| 31 |
+
}
|
| 32 |
+
}
|
| 33 |
+
acc += part * f32(b1[row * 32u + first / 32u]);
|
| 34 |
+
}
|
| 35 |
+
}
|
| 36 |
+
partial[local.x] = acc;
|
| 37 |
+
workgroupBarrier();
|
| 38 |
+
for (var stride = 8u; stride > 0u; stride /= 2u) {
|
| 39 |
+
if (lane < stride) { partial[local.x] += partial[local.x + stride]; }
|
| 40 |
+
workgroupBarrier();
|
| 41 |
+
}
|
| 42 |
+
if (lane == 0u && i < 512u) { out[i] = f32(partial[local.x] + 0.0); }
|
| 43 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1b9f4e330debc7235cfaf8f51d00fec575e09aa6116a45ea001e10ff5f09eb8c.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 * 1536u;
|
| 8 |
+
if (i >= 1536u) { return; }
|
| 9 |
+
let batch = i / 512u;
|
| 10 |
+
let row = (i / 16u) % 32u;
|
| 11 |
+
let col = i % 16u;
|
| 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) * 16u) + p * 16u + col]);
|
| 15 |
+
}
|
| 16 |
+
out[i] = f32(acc);
|
| 17 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1c441e95ec3ce8fee9348560af48191844baf402715f575940a0d4eb394c82d2.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 < 2048u && 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 < 2048u) {
|
| 66 |
+
let i = (batch * 4u + out_row) * 2048u + 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-v5/kernels/1d0cbb068fe17123e646934bfa1259f19e5db04e45d72a157d3ea4fa48c81aa0.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) * 131072u + ((i / 16384u) % 8u) * 256u + ((i / 256u) % 64u) * 2048u + ((i / 1u) % 256u) * 1u]);
|
| 9 |
+
}
|
qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v5/kernels/1d7df24388934e162899442d3cc6c6576791aac9f1a5dff86766d840f08fd46e.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 * 57344u;
|
| 8 |
+
if (i >= 57344u) { 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 |
+
}
|