diff --git a/.gitattributes b/.gitattributes index 4f662b336efdf4d56cdb180c5be74973c56b3b36..b9eccaa0a37aa41bcfadf44ad74b34e64d34a20e 100644 --- a/.gitattributes +++ b/.gitattributes @@ -49,3 +49,5 @@ qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/graph.json filter=lfs diff=lfs mer qwen35-2b-fp16-int4-g8-gptq-up-proj-candidate/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text qwen35-08b-fp32-int2-g32-mse-home-token-major-v2/graph.json filter=lfs diff=lfs merge=lfs -text qwen35-08b-fp32-int2-g32-mse-home-token-major-v2/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text +qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/graph.json filter=lfs diff=lfs merge=lfs -text +qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/graph.json b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/graph.json new file mode 100644 index 0000000000000000000000000000000000000000..81321b712a648f1ae0523cd94c73395d6a5693aa --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/graph.json @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:78123794a111311b0bdafead59f8fa1b93b71c31cc1b293c1c11221e3ccc33de +size 10577701 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/04f26661cc9ac73bb7f39f66b76c4be31ee8774e094eeb41ed7d677e1b1d3fa9.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/04f26661cc9ac73bb7f39f66b76c4be31ee8774e094eeb41ed7d677e1b1d3fa9.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..67133bd5fb075efe996a921acde0d02717ffa6f3 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/04f26661cc9ac73bb7f39f66b76c4be31ee8774e094eeb41ed7d677e1b1d3fa9.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 10u) { return; } +out[i] = f32(b0[((i / 10u) % 1u) * 32u + ((i / 10u) % 1u) * 32u + (((i / 1u) % 10u) * 3u + 2u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0554a443ecf6ead32fd2d58178b9ce3c592c89a09b3f63fa509ee947323bb497.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0554a443ecf6ead32fd2d58178b9ce3c592c89a09b3f63fa509ee947323bb497.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..586ebe23e7ab6f7d0ea3d4db2ae21148fc2ede1a --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0554a443ecf6ead32fd2d58178b9ce3c592c89a09b3f63fa509ee947323bb497.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 128u; +if (i >= 128u) { return; } +out[i] = f32(b0[((i / 128u) % 1u) * 512u + ((i / 64u) % 2u) * 256u + ((i / 64u) % 1u) * 256u + (((i / 1u) % 64u) * 1u + 0u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0680219da63feb9bb85653407739c5e5cc93ed6ca0971c1b03c512068d97f69f.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0680219da63feb9bb85653407739c5e5cc93ed6ca0971c1b03c512068d97f69f.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..b9fdb3d0db7aac20f07e3ed7d65c62a8689337f1 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0680219da63feb9bb85653407739c5e5cc93ed6ca0971c1b03c512068d97f69f.wgsl @@ -0,0 +1,21 @@ +enable f16; +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 16u; + if (i >= 16u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 16u) * 1024u + p]) * f32(b1[(i % 16u) * 1024u + p]); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0835fa80eaadb7c0e6d677e1fc622ad5dcedabb431e5d312f3b7aeb2252a7c7c.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0835fa80eaadb7c0e6d677e1fc622ad5dcedabb431e5d312f3b7aeb2252a7c7c.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..91f740aa7006d14dca451533719861c5823ff016 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0835fa80eaadb7c0e6d677e1fc622ad5dcedabb431e5d312f3b7aeb2252a7c7c.wgsl @@ -0,0 +1,8 @@ +@group(0) @binding(0) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(0.0); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0e5468583a1ba929483119c50d4be42a58aed4dde1e15cc231908ed7543d0b71.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0e5468583a1ba929483119c50d4be42a58aed4dde1e15cc231908ed7543d0b71.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..ade08a974329726324b1351190b859421467cb83 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0e5468583a1ba929483119c50d4be42a58aed4dde1e15cc231908ed7543d0b71.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 512u; +if (i >= 512u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0e9b6ee5f92769faaee5c746d1a3b936cdb7731f3f8e70cbf0a3fe96f6c8d921.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0e9b6ee5f92769faaee5c746d1a3b936cdb7731f3f8e70cbf0a3fe96f6c8d921.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..4f53d9e9c23631cf5a05331a915d782100def317 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0e9b6ee5f92769faaee5c746d1a3b936cdb7731f3f8e70cbf0a3fe96f6c8d921.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 32768u; +if (i >= 32768u) { return; } +out[i] = f32(f32(b0[i]) + (f32(b1[((i / 1u) % 4096u) * 1u]) * 1.0)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0f466ed2664d77386f48afd19d8ab2734312bfd4f5f4d5b063d801a0b0a97cd9.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0f466ed2664d77386f48afd19d8ab2734312bfd4f5f4d5b063d801a0b0a97cd9.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..a4fdbfd039369e36f1cdc9eba26852986c005d38 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0f466ed2664d77386f48afd19d8ab2734312bfd4f5f4d5b063d801a0b0a97cd9.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(b0[((i / 2048u) % 1u) * 2048u + ((i / 128u) % 16u) * 128u + ((i / 128u) % 1u) * 2048u + ((i / 1u) % 128u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0fd78c192784ba4e545091262cf1bac47ef69ca387b850f97b9de39dd0a2a741.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0fd78c192784ba4e545091262cf1bac47ef69ca387b850f97b9de39dd0a2a741.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..ff7be2a7dfd4febc209d987ef4330012d7dc722a --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0fd78c192784ba4e545091262cf1bac47ef69ca387b850f97b9de39dd0a2a741.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 16u; + if (i >= 16u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 16u) * 1024u + p]) * unpack_bf16_1((i % 16u) * 1024u + p); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/10feb208ed0549389f54419ff8985bee4778276913d4c6a613a8b120a809ab5e.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/10feb208ed0549389f54419ff8985bee4778276913d4c6a613a8b120a809ab5e.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..1d5bc2bc511d322d68a45fe9eb15628391bf5c85 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/10feb208ed0549389f54419ff8985bee4778276913d4c6a613a8b120a809ab5e.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 2048u; + if (i >= 2048u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 2048u) * 1024u + p]) * unpack_bf16_1((i % 2048u) * 1024u + p); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/12155d86d776e73f23d3bff0268c60dccb54c209ee9ee3a8cfd227ec4aa5b6df.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/12155d86d776e73f23d3bff0268c60dccb54c209ee9ee3a8cfd227ec4aa5b6df.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..6452ce1456ec5251372f66c853bffb54bd0465e4 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/12155d86d776e73f23d3bff0268c60dccb54c209ee9ee3a8cfd227ec4aa5b6df.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 1024u; +if (i >= 1024u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/1255847146b31732939e573cdb26c483358c1ce07f987bd8c3d3d69ae8c7e141.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/1255847146b31732939e573cdb26c483358c1ce07f987bd8c3d3d69ae8c7e141.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..19ce143a6079d1d8518ef9ca512fcc3cdfe1205c --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/1255847146b31732939e573cdb26c483358c1ce07f987bd8c3d3d69ae8c7e141.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 24576u; +if (i >= 24576u) { return; } +out[i] = f32(b0[((i / 24576u) % 1u) * 30720u + ((i / 4u) % 6144u) * 5u + (((i / 1u) % 4u) * 1u + 1u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/13237b557466d2fef5bfbc74e7ac77219b9821a53bc0843dbb5d14db8b04d414.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/13237b557466d2fef5bfbc74e7ac77219b9821a53bc0843dbb5d14db8b04d414.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..b455658f02601c7cfd2493171e0ae0ce54e635a5 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/13237b557466d2fef5bfbc74e7ac77219b9821a53bc0843dbb5d14db8b04d414.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(f32(b0[i]) * f32(b0[i])); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/13dc73935b9b3b978232043e13babee3877bd53ca572be440de0590dacce8b17.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/13dc73935b9b3b978232043e13babee3877bd53ca572be440de0590dacce8b17.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..de89efb15b0a288331488e51bd02760c3dded87f --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/13dc73935b9b3b978232043e13babee3877bd53ca572be440de0590dacce8b17.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 32u) { return; } +out[i] = f32(b0[0u * 32u + ((i / 32u) % 1u) * 32u + ((i / 32u) % 1u) * 32u + ((i / 1u) % 32u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/15b4ee3a9b5df59d5213e59951d8874f4eb4b9132e04fc2febddad84b44277f5.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/15b4ee3a9b5df59d5213e59951d8874f4eb4b9132e04fc2febddad84b44277f5.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..529a46d143f1017feeace76f912434dc166607bb --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/15b4ee3a9b5df59d5213e59951d8874f4eb4b9132e04fc2febddad84b44277f5.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 512u; + if (i >= 512u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 512u) * 1024u + p]) * unpack_bf16_1((i % 512u) * 1024u + p); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/17ec6bdaf98edf7657b74a25036c2bec806ecbb4280f33fff6afd2e0de5ed377.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/17ec6bdaf98edf7657b74a25036c2bec806ecbb4280f33fff6afd2e0de5ed377.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..f8b87c0bd2b9575bbbc8395e5d566f057aeb69dd --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/17ec6bdaf98edf7657b74a25036c2bec806ecbb4280f33fff6afd2e0de5ed377.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(f32(b0[i]) * 0.08838834764831843); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/1ad205080cae29142a3ad87343a0cf3ac4945455398293946110d421ce6d515e.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/1ad205080cae29142a3ad87343a0cf3ac4945455398293946110d421ce6d515e.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..04dc360880644481a17005996f0dece2fe491df3 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/1ad205080cae29142a3ad87343a0cf3ac4945455398293946110d421ce6d515e.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 128u; +if (i >= 128u) { return; } +out[i] = f32(f32(b0[i]) + (f32(b1[i]) * 1.0)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/20fb9e0d0bd2aab34b7f69eac259898e68cac4a0d6b0bda5e95c9126a660336d.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/20fb9e0d0bd2aab34b7f69eac259898e68cac4a0d6b0bda5e95c9126a660336d.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..5dc4ef605d0aeb8c1e6544aebe9f0ade791f4095 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/20fb9e0d0bd2aab34b7f69eac259898e68cac4a0d6b0bda5e95c9126a660336d.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 64u) { return; } +out[i] = f32(b0[((i / 64u) % 1u) * 128u + ((i / 32u) % 2u) * 64u + ((i / 32u) % 1u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/222a2645e3e02d0a4063a63ca1b6fa4c89faee6b76905b20de4dd54ab3077ea3.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/222a2645e3e02d0a4063a63ca1b6fa4c89faee6b76905b20de4dd54ab3077ea3.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..1447aa7dfa8f3765ba1ecc1e492c6fce42213779 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/222a2645e3e02d0a4063a63ca1b6fa4c89faee6b76905b20de4dd54ab3077ea3.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 3u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/234987ec68497956ea7aef8e3f44c8a73dbe4c1519165d123946bedbb4d18e26.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/234987ec68497956ea7aef8e3f44c8a73dbe4c1519165d123946bedbb4d18e26.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..f6c451610050d5332f2e96fe27fb0fc3d8359841 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/234987ec68497956ea7aef8e3f44c8a73dbe4c1519165d123946bedbb4d18e26.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(b0[((i / 2048u) % 1u) * 2048u + ((i / 128u) % 16u) * 128u + 0u * 128u + ((i / 1u) % 128u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/26cf6eb3b486d1652247ff130b05ed0e8c8560028eeba59acfc075b976bfdfce.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/26cf6eb3b486d1652247ff130b05ed0e8c8560028eeba59acfc075b976bfdfce.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..7018b9532a9e500c3ab541a52cb08ceeedd4810f --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/26cf6eb3b486d1652247ff130b05ed0e8c8560028eeba59acfc075b976bfdfce.wgsl @@ -0,0 +1,12 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 30720u; +if (i >= 30720u) { return; } +let coord = (i / 1u) % 5u; +if (coord >= 0u && coord < 4u) { out[i] = f32(b0[(i / 5u) * 4u + (coord - 0u) * 1u + i % 1u]); } +if (coord >= 4u && coord < 5u) { out[i] = f32(b1[(i / 5u) * 1u + (coord - 4u) * 1u + i % 1u]); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2711a4a396071f4e99303d01d4854ccaa37a445f4b718ebd7b61cf1c37e56b48.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2711a4a396071f4e99303d01d4854ccaa37a445f4b718ebd7b61cf1c37e56b48.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..8a950952abaa36f3937eab9390ad3807c223cb72 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2711a4a396071f4e99303d01d4854ccaa37a445f4b718ebd7b61cf1c37e56b48.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +out[i] = f32(b0[((i / 16u) % 1u) * 16u + ((i / 1u) % 16u) * 1u + ((i / 1u) % 1u) * 16u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2bdeb3590b1ba05cf6e9c7e8245f8aa552a12dae84f9578c603c9cee081a82d4.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2bdeb3590b1ba05cf6e9c7e8245f8aa552a12dae84f9578c603c9cee081a82d4.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..2b4e4e01f4485a846699c215551f49c2f1452a40 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2bdeb3590b1ba05cf6e9c7e8245f8aa552a12dae84f9578c603c9cee081a82d4.wgsl @@ -0,0 +1,24 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +var factor: f32; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let row = group.x + group.y * 16u; + if (row >= 16u) { return; } + let lane = local.x; + if (lane == 0u) { + var total = 0.0; + for (var j = 0u; j < 128u; j++) { + let v = f32(b0[row * 128u + j]); + total += v * v; + } + factor = inverseSqrt(total / 128.0 + 1e-06); + } + workgroupBarrier(); + for (var p = lane; p < 128u; p += 64u) { + let i = row * 128u + p; + out[i] = f32(f32(b0[row * 128u + p]) * factor * (f32(b1[p]) + 0.0)); + } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2c0cca2a7b0198c317e68992e64f3fe9bba80b67204c2695c20c372471c150d3.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2c0cca2a7b0198c317e68992e64f3fe9bba80b67204c2695c20c372471c150d3.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..05c73e3e7e256b462ade35867855db48ad811b19 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2c0cca2a7b0198c317e68992e64f3fe9bba80b67204c2695c20c372471c150d3.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 262144u; +if (i >= 262144u) { return; } +out[i] = f32(f32(b0[((i / 16384u) % 16u) * 128u + ((i / 128u) % 128u) * 1u]) * f32(b1[((i / 16384u) % 16u) * 128u + ((i / 1u) % 128u) * 1u])); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2dd29b61b039fcdd7fbc3c94f629b24b6f81aa5b9923098837597359b295e40d.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2dd29b61b039fcdd7fbc3c94f629b24b6f81aa5b9923098837597359b295e40d.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..9548ed074fc5f4321188aa402aa399a4525a86b3 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2dd29b61b039fcdd7fbc3c94f629b24b6f81aa5b9923098837597359b295e40d.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(b0[((i / 2048u) % 1u) * 2048u + ((i / 2048u) % 1u) * 128u + ((i / 128u) % 16u) * 128u + ((i / 1u) % 128u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..2269a620c58e84a4903bb1135733ef2385d6af46 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 64u) { return; } +out[i] = f32(f32(b0[i]) * 1.0); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/32a5027cac5dd969d63fe5d5c88d7c87abb600c43ad1d70e1c56615ea781eff3.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/32a5027cac5dd969d63fe5d5c88d7c87abb600c43ad1d70e1c56615ea781eff3.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..acbc0b91943f7202976d3a0119d1ad2f45215e18 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/32a5027cac5dd969d63fe5d5c88d7c87abb600c43ad1d70e1c56615ea781eff3.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 262144u; +if (i >= 262144u) { return; } +out[i] = f32(f32(b0[i]) * f32(b1[((i / 16384u) % 16u) * 128u + ((i / 128u) % 128u) * 1u])); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/38bd527686bcade9ed19a52788ccd711ca668e76854cc6274c480860fce88ff2.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/38bd527686bcade9ed19a52788ccd711ca668e76854cc6274c480860fce88ff2.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..934e1895005499ddd0c97450335aaf4cb8e9454e --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/38bd527686bcade9ed19a52788ccd711ca668e76854cc6274c480860fce88ff2.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(b0[((i / 2048u) % 1u) * 6144u + ((i / 2048u) % 1u) * 6144u + (((i / 1u) % 2048u) * 1u + 0u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/3bb850bda95bf77706a2839616a401b773f9de0428d1b427e61fb04e2c35fe74.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/3bb850bda95bf77706a2839616a401b773f9de0428d1b427e61fb04e2c35fe74.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..a3e680fcc3fb413659faf5613327f7d216420af9 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/3bb850bda95bf77706a2839616a401b773f9de0428d1b427e61fb04e2c35fe74.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 128u; +if (i >= 96u) { return; } +let coord = (i / 32u) % 3u; +if (coord == 0u) { out[i] = f32(b0[((i / 32u) % 1u) * 32u + ((i / 32u) % 1u) * 32u + ((i / 1u) % 32u) * 1u]); } else { out[i] = f32(b1[i]); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/3fa1efae6d90015ba69cb40cef009f49ad7f7a960018fb92a962707deff4f09e.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/3fa1efae6d90015ba69cb40cef009f49ad7f7a960018fb92a962707deff4f09e.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..dd0f6856d5316f3d610290ed98cf1a9c465f554d --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/3fa1efae6d90015ba69cb40cef009f49ad7f7a960018fb92a962707deff4f09e.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 10u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/413a0717b830335decd89485fab5fbea4efeb71338961ba7a2d1221eaa45459a.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/413a0717b830335decd89485fab5fbea4efeb71338961ba7a2d1221eaa45459a.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..220fe7f611e59561eb8b2b532830e97456170e0f --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/413a0717b830335decd89485fab5fbea4efeb71338961ba7a2d1221eaa45459a.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 1024u; + if (i >= 1024u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 3584u; p += 64u) { acc += f32(b0[(i / 1024u) * 3584u + p]) * unpack_bf16_1((i % 1024u) * 3584u + p); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/42c5af4838daf75fff9d4bb9a5df409511c1d6daed617d05eaf033a977503eac.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/42c5af4838daf75fff9d4bb9a5df409511c1d6daed617d05eaf033a977503eac.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..ab4b18691ec8371839272e82e1191ffc00772c73 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/42c5af4838daf75fff9d4bb9a5df409511c1d6daed617d05eaf033a977503eac.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 32u) { return; } +let coord = (i / 1u) % 32u; +if (coord >= 1u && coord < 34u && (coord - 1u) % 3u == 0u) { out[i] = f32(b0[((i / 32u) % 1u) * 11u + ((i / 32u) % 1u) * 11u + ((((i / 1u) % 32u) - 1u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4540d87eaec90c9eb42badfe0abfa70252f26b6e7de4dd31ee6903fe50e96bd1.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4540d87eaec90c9eb42badfe0abfa70252f26b6e7de4dd31ee6903fe50e96bd1.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..a0caa5c08df286ab5dbba084e1b6b90da802badc --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4540d87eaec90c9eb42badfe0abfa70252f26b6e7de4dd31ee6903fe50e96bd1.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 262144u; +if (i >= 262144u) { return; } +out[i] = f32(f32(b0[i]) + (f32(b1[i]) * 1.0)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4ac112c035a1a4a2f5deae66fac0969000c4666730c2d7b15cb15a157ea0e9b8.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4ac112c035a1a4a2f5deae66fac0969000c4666730c2d7b15cb15a157ea0e9b8.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..bc06995ea01af93b264998715f5789c5f01e0b80 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4ac112c035a1a4a2f5deae66fac0969000c4666730c2d7b15cb15a157ea0e9b8.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(b0[((i / 2048u) % 1u) * 2048u + ((i / 256u) % 8u) * 256u + ((i / 256u) % 1u) * 2048u + ((i / 1u) % 256u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4b1a7529fd5c57d5d4593d4af58a9a3b545c7d106d18251f6452f8a562b97952.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4b1a7529fd5c57d5d4593d4af58a9a3b545c7d106d18251f6452f8a562b97952.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..a2c7dcea92e4a5fa9a904796505efde56c5a3ae5 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4b1a7529fd5c57d5d4593d4af58a9a3b545c7d106d18251f6452f8a562b97952.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +out[i] = f32(f32(b0[i]) + (f32(b1[((i / 1u) % 16u) * 1u]) * 1.0)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4be963f9a4bb928764fd0efe02769e220b0e69defb2f8b22ab28e41a90d3fc14.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4be963f9a4bb928764fd0efe02769e220b0e69defb2f8b22ab28e41a90d3fc14.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..24b926b3b907f9052d46718955950614e3f4d5b2 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4be963f9a4bb928764fd0efe02769e220b0e69defb2f8b22ab28e41a90d3fc14.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 11u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4d3e5f3a99a6c8c9dba5de30be4302f1e442f91d097437e825df97017e389a42.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4d3e5f3a99a6c8c9dba5de30be4302f1e442f91d097437e825df97017e389a42.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..ad9eb2deb4d8a151884a987a9bff08c50ffd660f --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4d3e5f3a99a6c8c9dba5de30be4302f1e442f91d097437e825df97017e389a42.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 6144u; +if (i >= 6144u) { return; } +out[i] = f32(b0[((i / 6144u) % 1u) * 6144u + ((i / 6144u) % 1u) * 1u + ((i / 1u) % 6144u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4f48f08a8aeb51cc47a10db076a46522807d3b5d35952c9718ce1c1cd760681f.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4f48f08a8aeb51cc47a10db076a46522807d3b5d35952c9718ce1c1cd760681f.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..1fedbe2d560029fc3bb96a1cacb1cb911d5f0853 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4f48f08a8aeb51cc47a10db076a46522807d3b5d35952c9718ce1c1cd760681f.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 3584u; +if (i >= 3584u) { return; } +let x = f32(b0[i]); let s = f32(f32(x / (1.0 + exp(-x)))); +out[i] = f32(s * f32(b1[i])); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4fad29e3e337840e25132f10b3e3d55225f727ae13525ff0b0bdf31e03f09d74.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4fad29e3e337840e25132f10b3e3d55225f727ae13525ff0b0bdf31e03f09d74.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..77c9771680a7b21ef657c6b7303a2db2a2eeb5c4 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4fad29e3e337840e25132f10b3e3d55225f727ae13525ff0b0bdf31e03f09d74.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 256u; +if (i >= 256u) { return; } +let x = f32(b0[i]); +out[i] = f32(-x); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/51fd4d798491c6a80daff4c18b0495926b7f62269dfc3a78a327faa52ba2a759.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/51fd4d798491c6a80daff4c18b0495926b7f62269dfc3a78a327faa52ba2a759.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..fc751a3c5de6830de0104027f986f119ec7276f2 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/51fd4d798491c6a80daff4c18b0495926b7f62269dfc3a78a327faa52ba2a759.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +let x = f32(b0[i]); +out[i] = f32(exp(x)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/53d6eb0b5e6d49bfd3bb9fd1b781e13317adf528e047da2416e95f5a3a13b1b4.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/53d6eb0b5e6d49bfd3bb9fd1b781e13317adf528e047da2416e95f5a3a13b1b4.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..b42167010260a49ed2060dcc92a7440b82333b7e --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/53d6eb0b5e6d49bfd3bb9fd1b781e13317adf528e047da2416e95f5a3a13b1b4.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 128u; +if (i >= 96u) { return; } +out[i] = f32(b0[((i / 32u) % 3u) * 32u + ((i / 32u) % 1u) * 32u + ((i / 32u) % 1u) * 1u + ((i / 1u) % 32u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/552c2f036a125363ca85fd7ea5b1a84465b562dc62ddfffdbaaebf6e3288d321.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/552c2f036a125363ca85fd7ea5b1a84465b562dc62ddfffdbaaebf6e3288d321.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..58488000520b2483d92b7c966f4a23f6a46c2143 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/552c2f036a125363ca85fd7ea5b1a84465b562dc62ddfffdbaaebf6e3288d321.wgsl @@ -0,0 +1,17 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 1024u; +if (i >= 1024u) { return; } +let token = i32(b0[i / 1024u]); +if (token < 131072 || token >= 196608) { out[i] = f32(0.0); return; } +out[i] = f32(unpack_bf16_1(u32(token - 131072) * 1024u + i % 1024u)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/55b9c48ba429daf80bbae0b4737f895b475abdc9ddfdc88ab2f894f133ec2bf9.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/55b9c48ba429daf80bbae0b4737f895b475abdc9ddfdc88ab2f894f133ec2bf9.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..2af2397f376be638d8bdd9858e37047571ac05c7 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/55b9c48ba429daf80bbae0b4737f895b475abdc9ddfdc88ab2f894f133ec2bf9.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 30720u; +if (i >= 30720u) { return; } +let x = f32(b0[i]); +out[i] = f32(x / (1.0 + exp(-x))); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/579685c7f8ca260bd53d1677c43d5fc9a0133c81c7442ffb31daf0c3759267c4.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/579685c7f8ca260bd53d1677c43d5fc9a0133c81c7442ffb31daf0c3759267c4.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..191f3c622c2d82a894c6204a475d1b76af6bd95e --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/579685c7f8ca260bd53d1677c43d5fc9a0133c81c7442ffb31daf0c3759267c4.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 3584u; + if (i >= 3584u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 3584u) * 1024u + p]) * unpack_bf16_1((i % 3584u) * 1024u + p); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/57c5796344c26b430493a3a63dc734d4206380ba89484294f50f67d615eefcf8.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/57c5796344c26b430493a3a63dc734d4206380ba89484294f50f67d615eefcf8.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..8dc2442375bc0e019777ac15b4efd409d8a18dd7 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/57c5796344c26b430493a3a63dc734d4206380ba89484294f50f67d615eefcf8.wgsl @@ -0,0 +1,17 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 1024u; +if (i >= 1024u) { return; } +let token = i32(b0[i / 1024u]); +if (token < 196608 || token >= 248320) { out[i] = f32(0.0); return; } +out[i] = f32(unpack_bf16_1(u32(token - 196608) * 1024u + i % 1024u)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/57dc406322bf3badd8490493ef924bd7e7e0afd542d2d6e409b4d8a3b4b784f7.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/57dc406322bf3badd8490493ef924bd7e7e0afd542d2d6e409b4d8a3b4b784f7.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..5442fb5e440bbbbc406fb5988a08504f70fb6527 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/57dc406322bf3badd8490493ef924bd7e7e0afd542d2d6e409b4d8a3b4b784f7.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/59ea88d691b919af03bbb26f0253decfbc35d74f9661cdd3a08570ae85c313d2.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/59ea88d691b919af03bbb26f0253decfbc35d74f9661cdd3a08570ae85c313d2.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..5330d07c480a06ace52bbbeb6bf873fea1416d6c --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/59ea88d691b919af03bbb26f0253decfbc35d74f9661cdd3a08570ae85c313d2.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +let x = f32(b0[i]); +if (x * 1.0 > 20.0) { out[i] = f32(x); return; } +out[i] = f32(log(1.0 + exp(x * 1.0)) / 1.0); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5aa3725287640d289322ab469bd924bd7797f1cef28920ed4526783877c3cf6b.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5aa3725287640d289322ab469bd924bd7797f1cef28920ed4526783877c3cf6b.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..6def4ac27ef904a2b09da3fd89c8ca8c0eb2a2db --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5aa3725287640d289322ab469bd924bd7797f1cef28920ed4526783877c3cf6b.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 1024u; +if (i >= 1024u) { return; } +out[i] = f32(f32(b0[i]) + (f32(b1[i]) * 1.0)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5adf40822be73f2c85ed1c73c0bd106707e33fcb0e8d1e82a643474b8bfeceeb.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5adf40822be73f2c85ed1c73c0bd106707e33fcb0e8d1e82a643474b8bfeceeb.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..aa22ab6283b175cf72faef9edbcc001f23aacbb2 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5adf40822be73f2c85ed1c73c0bd106707e33fcb0e8d1e82a643474b8bfeceeb.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 65535u; + if (i >= 65536u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 65536u) * 1024u + p]) * unpack_bf16_1((i % 65536u) * 1024u + p); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5c22e6f60217881c0351e38d19057626a6c2292771c1a83816995245777ac2fd.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5c22e6f60217881c0351e38d19057626a6c2292771c1a83816995245777ac2fd.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..a0a200bcf1771e18ac73794fa78eac0151343d48 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5c22e6f60217881c0351e38d19057626a6c2292771c1a83816995245777ac2fd.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 256u; +if (i >= 256u) { return; } +out[i] = f32(b0[((i / 256u) % 1u) * 512u + ((i / 32u) % 8u) * 64u + ((i / 32u) % 1u) * 64u + (((i / 1u) % 32u) * 1u + 0u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5c9a5ae6121b3b10611e09e5f1bc563b70191f2506a2afc0f32de823a9d2dac4.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5c9a5ae6121b3b10611e09e5f1bc563b70191f2506a2afc0f32de823a9d2dac4.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..fd1c96b16a3a569b268cd100091060be8f6f800c --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5c9a5ae6121b3b10611e09e5f1bc563b70191f2506a2afc0f32de823a9d2dac4.wgsl @@ -0,0 +1,12 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 128u; +if (i >= 128u) { return; } +let coord = (i / 1u) % 64u; +if (coord >= 0u && coord < 32u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 0u) * 1u + i % 1u]); } +if (coord >= 32u && coord < 64u) { out[i] = f32(b1[(i / 64u) * 32u + (coord - 32u) * 1u + i % 1u]); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5eac5f92aa4d18895369b11adb99ac0d9639d8b615f7240cffe6f6be7046e54a.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5eac5f92aa4d18895369b11adb99ac0d9639d8b615f7240cffe6f6be7046e54a.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..44b1efb36c66e24c004b699b692f008a68d352e6 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5eac5f92aa4d18895369b11adb99ac0d9639d8b615f7240cffe6f6be7046e54a.wgsl @@ -0,0 +1,24 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +var factor: f32; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let row = group.x + group.y * 2u; + if (row >= 2u) { return; } + let lane = local.x; + if (lane == 0u) { + var total = 0.0; + for (var j = 0u; j < 256u; j++) { + let v = f32(b0[row * 256u + j]); + total += v * v; + } + factor = inverseSqrt(total / 256.0 + 1e-06); + } + workgroupBarrier(); + for (var p = lane; p < 256u; p += 64u) { + let i = row * 256u + p; + out[i] = f32(f32(b0[row * 256u + p]) * factor * (f32(b1[p]) + 1.0)); + } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5f4403d78557e88e66d71eb5e53746cc07449a7c421b57eda4f093a09a68a99a.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5f4403d78557e88e66d71eb5e53746cc07449a7c421b57eda4f093a09a68a99a.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..551de5c0fa29e8a924b81b6fe11c9a2cc3673d1d --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5f4403d78557e88e66d71eb5e53746cc07449a7c421b57eda4f093a09a68a99a.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 4096u; +if (i >= 4096u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5f657afde3ab80a1fc4b6739c0b702bc5bd60e52630762e66b210fbb8b42b52a.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5f657afde3ab80a1fc4b6739c0b702bc5bd60e52630762e66b210fbb8b42b52a.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..c35d470700b5e1330d30c2ea3f7e0c5d0f77d4a6 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5f657afde3ab80a1fc4b6739c0b702bc5bd60e52630762e66b210fbb8b42b52a.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 384u; +if (i >= 384u) { return; } +out[i] = f32(b0[((i / 384u) % 1u) * 512u + ((i / 192u) % 2u) * 256u + ((i / 192u) % 1u) * 256u + (((i / 1u) % 192u) * 1u + 64u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5fd6e068841f5c054392bef564909dbf086bae9e4030f3625d1143244a08cbbf.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5fd6e068841f5c054392bef564909dbf086bae9e4030f3625d1143244a08cbbf.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..d60668fd46afd65337a293687e115ac26feb2dab --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5fd6e068841f5c054392bef564909dbf086bae9e4030f3625d1143244a08cbbf.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 64u) { return; } +let coord = (i / 1u) % 64u; +if (coord >= 0u && coord < 32u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 0u) * 1u + i % 1u]); } +if (coord >= 32u && coord < 64u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 32u) * 1u + i % 1u]); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/67978e951813ede8016ebf553fb8d3b4ce9e15dd184810544879559208d6b7e1.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/67978e951813ede8016ebf553fb8d3b4ce9e15dd184810544879559208d6b7e1.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..3decac1417ec3c8a4798040f6305fb178904c916 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/67978e951813ede8016ebf553fb8d3b4ce9e15dd184810544879559208d6b7e1.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +let x = f32(b0[i]); +out[i] = f32(1.0 / (1.0 + exp(-x))); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/6a045508377e47e95d411257071b7ed8f9097fd220c00de6a5167abb548ebe37.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/6a045508377e47e95d411257071b7ed8f9097fd220c00de6a5167abb548ebe37.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..6bd34a6b3bb77a79fdcd3a519ca2897dc0b8084d --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/6a045508377e47e95d411257071b7ed8f9097fd220c00de6a5167abb548ebe37.wgsl @@ -0,0 +1,21 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 49152u; +if (i >= 49152u) { return; } +let batch = i / 49152u; +let channel = (i / 8u) % 6144u; +let group = channel / 1u; +var acc = 0.0; +for (var c = 0u; c < 1u; c++) { + for (var p = 0u; p < 4u; p++) { + let pos0 = i32((i / 1u) % 8u) * 1 + i32((p / 1u) % 4u) * 1 - 3; + if (pos0 >= 0 && pos0 < 5) { acc += f32(b0[batch * 30720u + (group * 1u + c) * 5u + u32(pos0) * 1u]) * f32(b1[(channel * 1u + c) * 4u + p]); } + } +} +out[i] = f32(acc + 0.0); + +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/6a7b2e3778a02c76155d46fa84aeb33f594954e82e09b6b7f2915c0a606eb3f5.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/6a7b2e3778a02c76155d46fa84aeb33f594954e82e09b6b7f2915c0a606eb3f5.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..dbb8e11929641cd77e1054eb002c379556e71409 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/6a7b2e3778a02c76155d46fa84aeb33f594954e82e09b6b7f2915c0a606eb3f5.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +var acc = 0.0; +for (var j = 0u; j < 128u; j++) { acc += f32(b0[(((i / 2048u) % 1u) * 262144u + ((i / 128u) % 16u) * 16384u + ((i / 1u) % 128u) * 1u) + (((j / 1u) % 128u) * 128u)]); } +out[i] = f32(acc); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/6ae4c7f13978e60e3f4ccd5b8172af88d7dc24cfdb236e51b826f13a03d8cd02.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/6ae4c7f13978e60e3f4ccd5b8172af88d7dc24cfdb236e51b826f13a03d8cd02.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..1d620a8021e4cd0706b810a2dbc2ee727980b6c2 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/6ae4c7f13978e60e3f4ccd5b8172af88d7dc24cfdb236e51b826f13a03d8cd02.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 512u; +if (i >= 512u) { return; } +out[i] = f32(b0[((i / 512u) % 1u) * 2048u + ((i / 64u) % 8u) * 256u + ((i / 64u) % 1u) * 256u + (((i / 1u) % 64u) * 1u + 0u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7002f8fdf09f5124ce2875c3540bce3e77750260556fc1a00a8b13b6c1871c30.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7002f8fdf09f5124ce2875c3540bce3e77750260556fc1a00a8b13b6c1871c30.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..5296bf14fb8b70f8480bfd7ccbbb1548cc142bca --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7002f8fdf09f5124ce2875c3540bce3e77750260556fc1a00a8b13b6c1871c30.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 512u; +if (i >= 512u) { return; } +out[i] = f32(f32(b0[i]) * f32(b1[((i / 1u) % 64u) * 1u])); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7121b5c223da0b0f28991f55cfed3c3f3390caa0bf65a8e58ae380db281ceb6b.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7121b5c223da0b0f28991f55cfed3c3f3390caa0bf65a8e58ae380db281ceb6b.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..69d777ae7caf63f2db83f47bb596dbe1e85a7c4b --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7121b5c223da0b0f28991f55cfed3c3f3390caa0bf65a8e58ae380db281ceb6b.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 11u) { return; } +out[i] = f32(b0[((i / 11u) % 1u) * 32u + ((i / 11u) % 1u) * 32u + (((i / 1u) % 11u) * 3u + 1u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7183997261247fa4df7e081c0d6d01036fee95b9b57ef944d5f44f72fc80d5fe.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7183997261247fa4df7e081c0d6d01036fee95b9b57ef944d5f44f72fc80d5fe.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..99d3e6e77543d2cc58b39e9dcb97109efd21a2fe --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7183997261247fa4df7e081c0d6d01036fee95b9b57ef944d5f44f72fc80d5fe.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 3u) { return; } +out[i] = i32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7194f0a413c1d8c6013b45b42b129c0a7557141a31bf3f08946fa1dfd1071393.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7194f0a413c1d8c6013b45b42b129c0a7557141a31bf3f08946fa1dfd1071393.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..3fb94c64c32b779b21c8bf8a58afb0556c82b14a --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7194f0a413c1d8c6013b45b42b129c0a7557141a31bf3f08946fa1dfd1071393.wgsl @@ -0,0 +1,24 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +var factor: f32; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let row = group.x + group.y * 1u; + if (row >= 1u) { return; } + let lane = local.x; + if (lane == 0u) { + var total = 0.0; + for (var j = 0u; j < 1024u; j++) { + let v = f32(b0[row * 1024u + j]); + total += v * v; + } + factor = inverseSqrt(total / 1024.0 + 1e-06); + } + workgroupBarrier(); + for (var p = lane; p < 1024u; p += 64u) { + let i = row * 1024u + p; + out[i] = f32(f32(b0[row * 1024u + p]) * factor * (f32(b1[p]) + 1.0)); + } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/720dc88131b6bf06e6b09185aa7d5986eb755dcfa60d0e040cbbf621c8f3f761.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/720dc88131b6bf06e6b09185aa7d5986eb755dcfa60d0e040cbbf621c8f3f761.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..77d2c6148b0bb0637e96ebb50a802068507c98e0 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/720dc88131b6bf06e6b09185aa7d5986eb755dcfa60d0e040cbbf621c8f3f761.wgsl @@ -0,0 +1,12 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 512u; +if (i >= 512u) { return; } +let coord = (i / 1u) % 256u; +if (coord >= 0u && coord < 64u) { out[i] = f32(b0[(i / 256u) * 64u + (coord - 0u) * 1u + i % 1u]); } +if (coord >= 64u && coord < 256u) { out[i] = f32(b1[(i / 256u) * 192u + (coord - 64u) * 1u + i % 1u]); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/742db031d819fdf5a756f37e8e824e2dcb2ca3f30efac3ffbd5918181059ae3a.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/742db031d819fdf5a756f37e8e824e2dcb2ca3f30efac3ffbd5918181059ae3a.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..0879520c73bf4f1395e8a15cc5874fa97dbecde1 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/742db031d819fdf5a756f37e8e824e2dcb2ca3f30efac3ffbd5918181059ae3a.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 262144u; +if (i >= 262144u) { return; } +out[i] = f32(f32(b0[i]) * f32(b1[((i / 16384u) % 16u) * 1u])); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7ea4a12734f1206e9c94a060bf6b24732f46685f0c46a78c3c5fbcd64489d8c3.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7ea4a12734f1206e9c94a060bf6b24732f46685f0c46a78c3c5fbcd64489d8c3.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..8e8a838dfb6ac9c617602bd0b2a35a31a9482fbd --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7ea4a12734f1206e9c94a060bf6b24732f46685f0c46a78c3c5fbcd64489d8c3.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 1536u; +if (i >= 1536u) { return; } +out[i] = f32(b0[((i / 1536u) % 1u) * 2048u + ((i / 192u) % 8u) * 256u + ((i / 192u) % 1u) * 256u + (((i / 1u) % 192u) * 1u + 64u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7fab02b2d46f9419002e2fc6f6e0e4e745ddb737a2c9e7060e45e7125c1be187.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7fab02b2d46f9419002e2fc6f6e0e4e745ddb737a2c9e7060e45e7125c1be187.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..3c05340808c48d12676c5acad9a623f97573efc1 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7fab02b2d46f9419002e2fc6f6e0e4e745ddb737a2c9e7060e45e7125c1be187.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2097152u; +if (i >= 2097152u) { return; } +out[i] = f32(b0[((i) / 256u) % 4096u * 512u + ((i) / 1048576u) % 2u * 256u + ((i) / 1u) % 256u * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7fe183583b43b627721ba6853b306288d7ceea4b42b630410f47ea9637220a4b.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7fe183583b43b627721ba6853b306288d7ceea4b42b630410f47ea9637220a4b.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..afbf92723660dbb4d364ef0cbeaf40d4cdc4891a --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/7fe183583b43b627721ba6853b306288d7ceea4b42b630410f47ea9637220a4b.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/8898353d0c367a3982a48d416f6777c3ce7a854572eeaf55843f3e7ee458c65d.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/8898353d0c367a3982a48d416f6777c3ce7a854572eeaf55843f3e7ee458c65d.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..ef9b5984a050b004ec691b6fd90a5704e441c981 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/8898353d0c367a3982a48d416f6777c3ce7a854572eeaf55843f3e7ee458c65d.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 24576u; +if (i >= 24576u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/88b50ebddd13f3ff5e2819b8bc956e8fc1346e8f75cacbc3d0edf031b7f1f248.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/88b50ebddd13f3ff5e2819b8bc956e8fc1346e8f75cacbc3d0edf031b7f1f248.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..af3f555638ab88bb67a3930d8f8766ec3b353d87 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/88b50ebddd13f3ff5e2819b8bc956e8fc1346e8f75cacbc3d0edf031b7f1f248.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 32768u; +if (i >= 32768u) { return; } +out[i] = f32(f32(b0[i]) * 0.0625); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/896006749045a93d89dca7ed7326aa05c5ff29423400569d34ab2256dea6b0e4.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/896006749045a93d89dca7ed7326aa05c5ff29423400569d34ab2256dea6b0e4.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..153b27098a91332adc33e93c50826711f884ea94 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/896006749045a93d89dca7ed7326aa05c5ff29423400569d34ab2256dea6b0e4.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 128u; +if (i >= 128u) { return; } +out[i] = f32(f32(b0[i]) * f32(b1[((i / 1u) % 64u) * 1u])); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/898086c99e03b2cdb78e3e5658c3fff202f1c87d09d8da08659bd1ba85e01b09.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/898086c99e03b2cdb78e3e5658c3fff202f1c87d09d8da08659bd1ba85e01b09.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..b5f7727d651290008885c57d7eb3d4d47da40e0d --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/898086c99e03b2cdb78e3e5658c3fff202f1c87d09d8da08659bd1ba85e01b09.wgsl @@ -0,0 +1,12 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 512u; +if (i >= 512u) { return; } +let coord = (i / 1u) % 64u; +if (coord >= 0u && coord < 32u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 0u) * 1u + i % 1u]); } +if (coord >= 32u && coord < 64u) { out[i] = f32(b1[(i / 64u) * 32u + (coord - 32u) * 1u + i % 1u]); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/8a46c4e24d014fd6af88a3f370c4b6161be7c63b6c30211bd62af3dbd87293a8.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/8a46c4e24d014fd6af88a3f370c4b6161be7c63b6c30211bd62af3dbd87293a8.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..879d6b86b22a39f635c45604888a324389324c38 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/8a46c4e24d014fd6af88a3f370c4b6161be7c63b6c30211bd62af3dbd87293a8.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 256u; +if (i >= 256u) { return; } +out[i] = f32(b0[((i / 256u) % 1u) * 512u + ((i / 32u) % 8u) * 64u + ((i / 32u) % 1u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/8d0a37e497bcb37e2d5a1746fba59fd69680ed8eeada4835c5d0eeef840d57ab.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/8d0a37e497bcb37e2d5a1746fba59fd69680ed8eeada4835c5d0eeef840d57ab.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..8a01c2daa51364f3125cbebcab9dbf7871446a9a --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/8d0a37e497bcb37e2d5a1746fba59fd69680ed8eeada4835c5d0eeef840d57ab.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +let x = f32(b0[i]); +out[i] = f32(-x); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/9144d827dae4c52b80abb829178a9ed901ce5dbc6702ad94b0b355e8eaad5ea5.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/9144d827dae4c52b80abb829178a9ed901ce5dbc6702ad94b0b355e8eaad5ea5.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..0e196cbf9c5f57d582e8784874efe3d9797e27fe --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/9144d827dae4c52b80abb829178a9ed901ce5dbc6702ad94b0b355e8eaad5ea5.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 32u) { return; } +let coord = (i / 1u) % 32u; +if (coord >= 2u && coord < 32u && (coord - 2u) % 3u == 0u) { out[i] = f32(b0[((i / 32u) % 1u) * 10u + ((i / 32u) % 1u) * 10u + ((((i / 1u) % 32u) - 2u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/93d02cadeb259e826a8d196998a25cbb09a7053b0fa50174dc6c0d9ec05f8d6d.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/93d02cadeb259e826a8d196998a25cbb09a7053b0fa50174dc6c0d9ec05f8d6d.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..bca0d8a07f14145b06cebdb254f80a0a73f05f96 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/93d02cadeb259e826a8d196998a25cbb09a7053b0fa50174dc6c0d9ec05f8d6d.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 4096u; + if (i >= 4096u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 4096u) * 1024u + p]) * unpack_bf16_1((i % 4096u) * 1024u + p); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/947f5cc8b3f87e1a0d4d322d9d9e71e2ed16ce6e1f7c91118678c164dc1c8752.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/947f5cc8b3f87e1a0d4d322d9d9e71e2ed16ce6e1f7c91118678c164dc1c8752.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..2bb5bd4f5d110df131987419e6018288811ea172 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/947f5cc8b3f87e1a0d4d322d9d9e71e2ed16ce6e1f7c91118678c164dc1c8752.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 128u; +if (i >= 96u) { return; } +out[i] = f32(b0[((i / 32u) % 3u) * 32u + ((i / 32u) % 1u) * 32u + ((i / 1u) % 32u) * 1u + ((i / 1u) % 1u) * 32u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/95389a9d24d01ffc452bc9cea28e01f8bbca2202271cbe96b1399afbbe725534.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/95389a9d24d01ffc452bc9cea28e01f8bbca2202271cbe96b1399afbbe725534.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..d597b239a895fd2fd83f160cfff0cfdee12563e4 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/95389a9d24d01ffc452bc9cea28e01f8bbca2202271cbe96b1399afbbe725534.wgsl @@ -0,0 +1,28 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var b2: array; +@group(0) @binding(3) var out: array; +fn dequant_1(index: u32) -> f32 { + let row = index / 1024u; + let col = index % 1024u; + let word = b1[row * 64u + col / 16u]; + let code = i32((word >> ((col % 16u) * 2u)) & 3u) - 2; + return f32(code) * b2[row * 32u + col / 32u]; +} + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 3584u; + if (i >= 3584u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 3584u) * 1024u + p]) * dequant_1((i % 3584u) * 1024u + p); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/95addfdb139df33bb3d5c05859412434d318cee0340560537dce4d56a05e159c.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/95addfdb139df33bb3d5c05859412434d318cee0340560537dce4d56a05e159c.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..a5e5b4d6f13125cc2cdaea02ee4549dd7bf38f8b --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/95addfdb139df33bb3d5c05859412434d318cee0340560537dce4d56a05e159c.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(b0[((i / 2048u) % 1u) * 2048u + ((i / 2048u) % 1u) * 256u + ((i / 256u) % 8u) * 256u + ((i / 1u) % 256u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/9625ca0b071511e5261877fd21627ca354ea0219da86a10887e8054afa15dcfd.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/9625ca0b071511e5261877fd21627ca354ea0219da86a10887e8054afa15dcfd.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..e1f0cf834667ef3ff29080e64d8f9d08e6702c3e --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/9625ca0b071511e5261877fd21627ca354ea0219da86a10887e8054afa15dcfd.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 32u) { return; } +out[i] = f32(b0[1u * 32u + ((i / 32u) % 1u) * 32u + ((i / 32u) % 1u) * 32u + ((i / 1u) % 32u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/985bb9ad33654b9adbd4a256d626e346556a10dad514692800da8ebd7910f0f0.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/985bb9ad33654b9adbd4a256d626e346556a10dad514692800da8ebd7910f0f0.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..9018c3fdbf086b98282a94876e8a6631eb5442af --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/985bb9ad33654b9adbd4a256d626e346556a10dad514692800da8ebd7910f0f0.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 3u) { return; } +out[i] = i32(b0[(((i / 1u) % 3u) * 1u + 1u) * 1u + ((i / 1u) % 1u) * 1u + ((i / 1u) % 1u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/9f02cd7a8eb5e6d20c2b01559fc58db5b0447ae41a51e2cd479ed519fc284791.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/9f02cd7a8eb5e6d20c2b01559fc58db5b0447ae41a51e2cd479ed519fc284791.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..928d1ad581850d91a2b178bed4f5d3f5e81e28b1 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/9f02cd7a8eb5e6d20c2b01559fc58db5b0447ae41a51e2cd479ed519fc284791.wgsl @@ -0,0 +1,16 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var b2: array; +@group(0) @binding(3) var b3: array; +@group(0) @binding(4) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 248320u; +if (i >= 248320u) { return; } +let coord = (i / 1u) % 248320u; +if (coord >= 0u && coord < 65536u) { out[i] = f32(b0[(i / 248320u) * 65536u + (coord - 0u) * 1u + i % 1u]); } +if (coord >= 65536u && coord < 131072u) { out[i] = f32(b1[(i / 248320u) * 65536u + (coord - 65536u) * 1u + i % 1u]); } +if (coord >= 131072u && coord < 196608u) { out[i] = f32(b2[(i / 248320u) * 65536u + (coord - 131072u) * 1u + i % 1u]); } +if (coord >= 196608u && coord < 248320u) { out[i] = f32(b3[(i / 248320u) * 51712u + (coord - 196608u) * 1u + i % 1u]); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a1a974f6d8d51fd7c3496ede0966af303ad9f0c0c9ee44cf93ec962d2b0e87e1.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a1a974f6d8d51fd7c3496ede0966af303ad9f0c0c9ee44cf93ec962d2b0e87e1.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..12d890f4a402b0675a5fd2b6d98f661c3f315dd5 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a1a974f6d8d51fd7c3496ede0966af303ad9f0c0c9ee44cf93ec962d2b0e87e1.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 64u) { return; } +let x = f32(b0[i]); +out[i] = f32(-x); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a37dac611b27158f19aa5bd70c43217286e610e733cbd602b91666b4af42d3ef.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a37dac611b27158f19aa5bd70c43217286e610e733cbd602b91666b4af42d3ef.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..c0eaa7ea6d730c603c8ce3193ea9e00a3ec74433 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a37dac611b27158f19aa5bd70c43217286e610e733cbd602b91666b4af42d3ef.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 1u) { return; } +out[i] = i32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a672c3e158716bde7444c8141f3df1da16da739ccdd2eb2ad60f9d42d0775157.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a672c3e158716bde7444c8141f3df1da16da739ccdd2eb2ad60f9d42d0775157.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..aa456f658303b9267eb9faab7425cd740bbdc6f5 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a672c3e158716bde7444c8141f3df1da16da739ccdd2eb2ad60f9d42d0775157.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +let coord = (i / 128u) % 1u; +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]); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a77f14c89662d2d3da8abd8f109667793dc662d307e5e1c61123be753e25437b.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a77f14c89662d2d3da8abd8f109667793dc662d307e5e1c61123be753e25437b.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..8f7e080e8cee8c2acabacf51697ca01fe402a237 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/a77f14c89662d2d3da8abd8f109667793dc662d307e5e1c61123be753e25437b.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 128u; +if (i >= 96u) { return; } +out[i] = f32(b0[((i / 1u) % 32u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/abdb17daf29310a7828a55a1efeebaa65b6553e236fc20702e8c712e0a46dd5b.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/abdb17daf29310a7828a55a1efeebaa65b6553e236fc20702e8c712e0a46dd5b.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..0e63cebe9c685e72f2ef3807fe8a83b8853e6dd9 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/abdb17daf29310a7828a55a1efeebaa65b6553e236fc20702e8c712e0a46dd5b.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +let x = f32(b0[i]); +out[i] = f32(1.0 / (1.0 + exp(-x))); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b02b0be460d581c720afd68c8b722e8da47e2ee432b5d61bbdf24b6084cb2f26.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b02b0be460d581c720afd68c8b722e8da47e2ee432b5d61bbdf24b6084cb2f26.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..149a4ba7b993cd3c7a5ae3d437265ae0e1001e0e --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b02b0be460d581c720afd68c8b722e8da47e2ee432b5d61bbdf24b6084cb2f26.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +let x = f32(b0[i]); let s = f32(f32(x / (1.0 + exp(-x)))); +out[i] = f32(s * f32(b1[i])); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b199d835db70f888c7b99a3ac6e9ec56f6c47d5bba4eaf352e0a8eabed56411e.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b199d835db70f888c7b99a3ac6e9ec56f6c47d5bba4eaf352e0a8eabed56411e.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..54526e84e0f57f0b27c2e0d57273b25cde83790f --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b199d835db70f888c7b99a3ac6e9ec56f6c47d5bba4eaf352e0a8eabed56411e.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(f32(b0[i]) * f32(b1[((i / 128u) % 16u) * 1u])); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b2b16798cbe94c919025d76b0435afe47966ffdd1ea53ec74a791c03c4709a28.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b2b16798cbe94c919025d76b0435afe47966ffdd1ea53ec74a791c03c4709a28.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..4d8a5bb75deb878ba50367abeb2220352b270284 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b2b16798cbe94c919025d76b0435afe47966ffdd1ea53ec74a791c03c4709a28.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 30720u; +if (i >= 30720u) { return; } +out[i] = f32(b0[((i / 30720u) % 1u) * 49152u + ((i / 5u) % 6144u) * 8u + (((i / 1u) % 5u) * 1u + 0u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b3128f8b5ac3b671947588c52e53f26426877fd37a2e7bdfcd9819a6cbb7e3a7.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b3128f8b5ac3b671947588c52e53f26426877fd37a2e7bdfcd9819a6cbb7e3a7.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..71e65d6a93cf6c0306d093bbd7707c022315afe0 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b3128f8b5ac3b671947588c52e53f26426877fd37a2e7bdfcd9819a6cbb7e3a7.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 4194240u; +if (i >= 8388608u) { return; } +out[i] = f32(b0[((i / 8388608u) % 1u) * 8388608u + ((i / 1048576u) % 8u) * 1048576u + ((i / 4096u) % 256u) * 1u + ((i / 1u) % 4096u) * 256u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b57928140994202536cd4f99d2e14f5dba9754f014843673790d5d48c9cbb9ab.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b57928140994202536cd4f99d2e14f5dba9754f014843673790d5d48c9cbb9ab.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..a5d1a988399067b52e6658a8d39dcbfac3d81a20 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b57928140994202536cd4f99d2e14f5dba9754f014843673790d5d48c9cbb9ab.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 512u; +if (i >= 512u) { return; } +out[i] = f32(b0[((i / 512u) % 1u) * 512u + ((i / 256u) % 2u) * 256u + ((i / 256u) % 1u) * 512u + ((i / 1u) % 256u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b71488c9b47fa85c9dcdbbf9185c31ab819169ab631e1aa901f65e2999c5d22f.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b71488c9b47fa85c9dcdbbf9185c31ab819169ab631e1aa901f65e2999c5d22f.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..68cd0341550bad4178f3ef83426d09d045b7b659 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b71488c9b47fa85c9dcdbbf9185c31ab819169ab631e1aa901f65e2999c5d22f.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 1024u; + if (i >= 1024u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 2048u; p += 64u) { acc += f32(b0[(i / 1024u) * 2048u + p]) * unpack_bf16_1((i % 1024u) * 2048u + p); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b75a05038c8d340aa61c4cdf617403ed9df75a05ecb9284c659b456ee9e120cd.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b75a05038c8d340aa61c4cdf617403ed9df75a05ecb9284c659b456ee9e120cd.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..ad300da6f1af3bf2cdfbf4fc7fd7d4b5163052ee --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/b75a05038c8d340aa61c4cdf617403ed9df75a05ecb9284c659b456ee9e120cd.wgsl @@ -0,0 +1,17 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 1024u; +if (i >= 1024u) { return; } +let token = i32(b0[i / 1024u]); +if (token < 0 || token >= 65536) { out[i] = f32(0.0); return; } +out[i] = f32(unpack_bf16_1(u32(token - 0) * 1024u + i % 1024u)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/bb9029b667eeb13670d1f640cd677c6eabd7a7569686c599e3eeb9482d642b00.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/bb9029b667eeb13670d1f640cd677c6eabd7a7569686c599e3eeb9482d642b00.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..70d42d4c763bcca299e524c1695c1ae352353cd9 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/bb9029b667eeb13670d1f640cd677c6eabd7a7569686c599e3eeb9482d642b00.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(f32(b0[i]) - (f32(b1[i]) * 1.0)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/bbe7cf742bbc4d053e33018f81ca9e71a5f66cf38db88df0fe1b4584b27641f0.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/bbe7cf742bbc4d053e33018f81ca9e71a5f66cf38db88df0fe1b4584b27641f0.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..b65c3c5e7920f28a3c0ad4dafb598bd0112252df --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/bbe7cf742bbc4d053e33018f81ca9e71a5f66cf38db88df0fe1b4584b27641f0.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 51712u; + if (i >= 51712u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 51712u) * 1024u + p]) * unpack_bf16_1((i % 51712u) * 1024u + p); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c0a7f99d0ce8f9a199b9baa52effe47afa9605f77405c32b32d132be1f7404f1.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c0a7f99d0ce8f9a199b9baa52effe47afa9605f77405c32b32d132be1f7404f1.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..bf490a01533f687f754f69957cc65aa09be66c54 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c0a7f99d0ce8f9a199b9baa52effe47afa9605f77405c32b32d132be1f7404f1.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 64u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c0b2947798baa198598d75ca0e5dcd31a9fa320499665e34f7bb81d27fc15fb0.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c0b2947798baa198598d75ca0e5dcd31a9fa320499665e34f7bb81d27fc15fb0.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..bb8b3676ab0dfbc0c5c78d28f96387bc3a721046 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c0b2947798baa198598d75ca0e5dcd31a9fa320499665e34f7bb81d27fc15fb0.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 64u) { return; } +let x = f32(b0[i]); +out[i] = f32(cos(x)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c23bee3c3664172bf0a078a96acd2b3ad3136732234975f757d313e696c2c37a.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c23bee3c3664172bf0a078a96acd2b3ad3136732234975f757d313e696c2c37a.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..2d80755364a4fee024c589b10f9f00932c251576 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c23bee3c3664172bf0a078a96acd2b3ad3136732234975f757d313e696c2c37a.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +out[i] = f32(f32(b0[((i / 1u) % 16u) * 1u]) * f32(b1[i])); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c295fd955666232c92869895d186ebd33a25e4a5bf14df41d5b455fbe304d007.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c295fd955666232c92869895d186ebd33a25e4a5bf14df41d5b455fbe304d007.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..22b9338f6e2f9c6948f9be91fb996e767aa1f5dc --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c295fd955666232c92869895d186ebd33a25e4a5bf14df41d5b455fbe304d007.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 4194240u; +if (i >= 8388608u) { return; } +out[i] = f32(b0[((i / 4194304u) % 2u) * 1048576u + ((i / 256u) % 4096u) * 256u + ((i / 1u) % 256u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c2c806bbe93f09247867b4ace45489a5f2fe3ee67f265418eb7aa71a29880276.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c2c806bbe93f09247867b4ace45489a5f2fe3ee67f265418eb7aa71a29880276.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..fa4df64a06b41dc3305ec748ea504d435a19aa23 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c2c806bbe93f09247867b4ace45489a5f2fe3ee67f265418eb7aa71a29880276.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +let x = f32(b0[i]); +out[i] = f32(inverseSqrt(x)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c4862be63d2532b6891da039edbef2f46737a34518f8a89602821e00447b5933.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c4862be63d2532b6891da039edbef2f46737a34518f8a89602821e00447b5933.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..7e12c353c29523774d93254e8dc9f119ce04b4fb --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c4862be63d2532b6891da039edbef2f46737a34518f8a89602821e00447b5933.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(b0[((i / 2048u) % 1u) * 6144u + ((i / 2048u) % 1u) * 6144u + (((i / 1u) % 2048u) * 1u + 2048u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c72041e82befee6919af48775291256e22a43dc6445ef0aa288db6b96585a2a9.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c72041e82befee6919af48775291256e22a43dc6445ef0aa288db6b96585a2a9.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..ceffcce3113d363ae38d0760c9af3a93e2a924d3 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c72041e82befee6919af48775291256e22a43dc6445ef0aa288db6b96585a2a9.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 32u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c90428425f6d220b441e97e3b894adfb090869c98b17f53afbd2f9fb202b4525.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c90428425f6d220b441e97e3b894adfb090869c98b17f53afbd2f9fb202b4525.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..6bbc472bea9c8d96ae70e94b1ca3e146cd081981 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/c90428425f6d220b441e97e3b894adfb090869c98b17f53afbd2f9fb202b4525.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 64u) { return; } +let x = f32(b0[i]); +out[i] = f32(sin(x)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/cc250879ffbf43ba05860b5a86f6a824e8ca7b4c9f820137a68e1ba633e09738.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/cc250879ffbf43ba05860b5a86f6a824e8ca7b4c9f820137a68e1ba633e09738.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..f99a31a5530ee04c3f3573c324dd7c6471dc1049 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/cc250879ffbf43ba05860b5a86f6a824e8ca7b4c9f820137a68e1ba633e09738.wgsl @@ -0,0 +1,17 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 128u; +if (i >= 96u) { return; } +let batch = i / 32u; +let row = (i / 1u) % 32u; +let col = i % 1u; +var acc = 0.0; +for (var p = 0u; p < 1u; p++) { + acc += f32(b0[(((batch / 1u) % 3u) * 32u) + row * 1u + p]) * f32(b1[(((batch / 1u) % 3u) * 1u) + p * 1u + col]); +} +out[i] = f32(acc); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/ccb303d2d38ef10c1cf514f01ec2fa3b2541734aca8a34b3e8bb8d71d03a6688.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/ccb303d2d38ef10c1cf514f01ec2fa3b2541734aca8a34b3e8bb8d71d03a6688.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..228b4bbd0824c9c80c477855a3d00e3274aa4588 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/ccb303d2d38ef10c1cf514f01ec2fa3b2541734aca8a34b3e8bb8d71d03a6688.wgsl @@ -0,0 +1,13 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x + gid.y * 512u; + if (i >= 512u) { return; } + let token = (i / 256u) % 1u; + let outer = i / 256u; + let destination = outer * 1048576u + u32(b1[token]) * 256u + i % 256u; + out[((destination) / 256u) % 4096u * 512u + ((destination) / 1048576u) % 2u * 256u + ((destination) / 1u) % 256u * 1u] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/cdc784b39918d10424edc08d2e44da0075863e8fd41ae362f5695a248363531d.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/cdc784b39918d10424edc08d2e44da0075863e8fd41ae362f5695a248363531d.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..ee515425cf5358aa221f3366d77a028e1327f04d --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/cdc784b39918d10424edc08d2e44da0075863e8fd41ae362f5695a248363531d.wgsl @@ -0,0 +1,34 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let row = group.x + group.y * 8u; + if (row >= 8u) { return; } + let lane = local.x; + var maximum = -3.4028234663852886e38; + for (var p = lane; p < 4096u; p += 64u) { maximum = max(maximum, f32(b0[row * 4096u + p])); } + partial[lane] = maximum; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] = max(partial[lane], partial[lane + stride]); } + workgroupBarrier(); + } + maximum = partial[0]; + // All lanes must finish reading the maximum before the scratch array is reused. + workgroupBarrier(); + var total = 0.0; + for (var p = lane; p < 4096u; p += 64u) { total += exp(f32(b0[row * 4096u + p]) - maximum); } + partial[lane] = total; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + total = partial[0]; + for (var p = lane; p < 4096u; p += 64u) { + let i = row * 4096u + p; + out[i] = f32(exp(f32(b0[row * 4096u + p]) - maximum) / total); + } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/d4c80950a743190f3c54e4675a4fb139edf142a4737c55d803326c1a1cedcc2e.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/d4c80950a743190f3c54e4675a4fb139edf142a4737c55d803326c1a1cedcc2e.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..22592bb6be29c49767f5012f185f69bd03d17d7c --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/d4c80950a743190f3c54e4675a4fb139edf142a4737c55d803326c1a1cedcc2e.wgsl @@ -0,0 +1,17 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 1024u; +if (i >= 1024u) { return; } +let token = i32(b0[i / 1024u]); +if (token < 65536 || token >= 131072) { out[i] = f32(0.0); return; } +out[i] = f32(unpack_bf16_1(u32(token - 65536) * 1024u + i % 1024u)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/dc4ee648b97b273931e31c95b910d66f11e34a1cb036589e9ae9cc250b6e6291.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/dc4ee648b97b273931e31c95b910d66f11e34a1cb036589e9ae9cc250b6e6291.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..0e13dfcb4085c03a943bab386ba2afe5e2a6133c --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/dc4ee648b97b273931e31c95b910d66f11e34a1cb036589e9ae9cc250b6e6291.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 4u) { return; } +out[i] = i32(b0[0u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/de276eae1430ed164ae977ce4b005301ea0f2652bed1f879acb3a355e600ff5e.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/de276eae1430ed164ae977ce4b005301ea0f2652bed1f879acb3a355e600ff5e.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..684e67f205ed8e5d2b379a1b312f0339ea744c07 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/de276eae1430ed164ae977ce4b005301ea0f2652bed1f879acb3a355e600ff5e.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(f32(b0[i]) * f32(b1[i])); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/deee894ecb86334ff82870260fc5edbc36439c1d3bce09fab5708c03d699b390.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/deee894ecb86334ff82870260fc5edbc36439c1d3bce09fab5708c03d699b390.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..edaf6eca29e64fa2e4a563af4eaf9fe76bfb99b7 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/deee894ecb86334ff82870260fc5edbc36439c1d3bce09fab5708c03d699b390.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 32u) { return; } +out[i] = f32(b0[2u * 32u + ((i / 32u) % 1u) * 32u + ((i / 32u) % 1u) * 32u + ((i / 1u) % 32u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e01dc46d731ca9d3358814482da28b4edb8496243937be85ddde86a02a215169.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e01dc46d731ca9d3358814482da28b4edb8496243937be85ddde86a02a215169.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..5bc58ca8a83e46e3403c1f4ea7a6fd95358ac3d0 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e01dc46d731ca9d3358814482da28b4edb8496243937be85ddde86a02a215169.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 6144u; +if (i >= 6144u) { return; } +out[i] = f32(b0[((i / 6144u) % 1u) * 30720u + ((i / 1u) % 6144u) * 5u + (((i / 1u) % 1u) * 1u + 4u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e23f8d0e93dd1cec22eea5c04988cc373462885008c37beb9e5f012773c6af1f.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e23f8d0e93dd1cec22eea5c04988cc373462885008c37beb9e5f012773c6af1f.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..5888b024d6312049693e552cd2cdaa1d267bf4ac --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e23f8d0e93dd1cec22eea5c04988cc373462885008c37beb9e5f012773c6af1f.wgsl @@ -0,0 +1,25 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; +fn unpack_bf16_1(index: u32) -> f32 { + let pair = b1[index / 2u]; + let bits = (pair >> ((index % 2u) * 16u)) & 65535u; + return bitcast(bits << 16u); +} + +var partial: array; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let i = group.x + group.y * 6144u; + if (i >= 6144u) { return; } + let lane = local.x; + var acc = 0.0; + for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 6144u) * 1024u + p]) * unpack_bf16_1((i % 6144u) * 1024u + p); } + partial[lane] = acc; + workgroupBarrier(); + for (var stride = 32u; stride > 0u; stride /= 2u) { + if (lane < stride) { partial[lane] += partial[lane + stride]; } + workgroupBarrier(); + } + if (lane == 0u) { out[i] = f32(partial[0] + 0.0); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e38b21c1b50403e478d2d4956df390560aa740de2dcba8d69322d996c998adf4.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e38b21c1b50403e478d2d4956df390560aa740de2dcba8d69322d996c998adf4.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..2eb119d2449221b9b8bace6cfb92a13e33a3c95c --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e38b21c1b50403e478d2d4956df390560aa740de2dcba8d69322d996c998adf4.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(b0[((i / 2048u) % 1u) * 4096u + ((i / 2048u) % 1u) * 4096u + ((i / 256u) % 8u) * 512u + (((i / 1u) % 256u) * 1u + 0u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e6571e9ffa79d4f7888d606a5845a6fe7d39e3b7944eafb979c162a384245c66.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e6571e9ffa79d4f7888d606a5845a6fe7d39e3b7944eafb979c162a384245c66.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..1cca9e30814dd2231912899d446310b2c59b8009 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e6571e9ffa79d4f7888d606a5845a6fe7d39e3b7944eafb979c162a384245c66.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(b0[((i / 2048u) % 1u) * 6144u + ((i / 2048u) % 1u) * 6144u + (((i / 1u) % 2048u) * 1u + 4096u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e829b3bd2efe5b094dc36238e2a933a1e30c942efddfc80e79862a55c57e6e20.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e829b3bd2efe5b094dc36238e2a933a1e30c942efddfc80e79862a55c57e6e20.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..c061788bda0478004560218a522ad6db46d1c40e --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/e829b3bd2efe5b094dc36238e2a933a1e30c942efddfc80e79862a55c57e6e20.wgsl @@ -0,0 +1,11 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +var acc = 0.0; +for (var j = 0u; j < 128u; j++) { acc += f32(b0[(((i / 16u) % 1u) * 2048u + ((i / 16u) % 1u) * 2048u + ((i / 1u) % 16u) * 128u) + (((j / 1u) % 128u) * 1u)]); } +out[i] = f32(acc); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/ec606c53f8610fff07d8464d8f1cd2a909203cd4064d1be19ad9bdb42263794d.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/ec606c53f8610fff07d8464d8f1cd2a909203cd4064d1be19ad9bdb42263794d.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..5d9693423508c7bd74ec4f23c60f4d86602c81e1 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/ec606c53f8610fff07d8464d8f1cd2a909203cd4064d1be19ad9bdb42263794d.wgsl @@ -0,0 +1,17 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +let batch = i / 256u; +let row = (i / 256u) % 1u; +let col = i % 256u; +var acc = 0.0; +for (var p = 0u; p < 4096u; p++) { + acc += f32(b0[(((batch / 1u) % 8u) * 4096u) + row * 4096u + p]) * f32(b1[(((batch / 1u) % 8u) * 1048576u) + p * 256u + col]); +} +out[i] = f32(acc); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/edfa4ff396bf5cc3be28d9051639c4dda6fec1302d7a2bd34d89ec903af94af1.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/edfa4ff396bf5cc3be28d9051639c4dda6fec1302d7a2bd34d89ec903af94af1.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..fa04fd0adecdb9f18898bad16cb516ccdf4bf468 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/edfa4ff396bf5cc3be28d9051639c4dda6fec1302d7a2bd34d89ec903af94af1.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +out[i] = f32(f32(b0[i]) + (1e-06 * 1.0)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/eea033222d278c8aa8f4b0de36243adba4e5b6ce274d24a3dcd365ad42c19a38.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/eea033222d278c8aa8f4b0de36243adba4e5b6ce274d24a3dcd365ad42c19a38.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..7dd035a2b7b4ce728303f33a2f815174cbedc865 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/eea033222d278c8aa8f4b0de36243adba4e5b6ce274d24a3dcd365ad42c19a38.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +out[i] = f32(b0[((i / 2048u) % 1u) * 4096u + ((i / 2048u) % 1u) * 4096u + ((i / 256u) % 8u) * 512u + (((i / 1u) % 256u) * 1u + 256u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f114c5cb9f1a7e377d9a3bcb28ab15b72807ece1796cf9c02fc74e7ce7318cc3.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f114c5cb9f1a7e377d9a3bcb28ab15b72807ece1796cf9c02fc74e7ce7318cc3.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..71fee4a974a4ef03d91dd3ce0680762c347a9577 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f114c5cb9f1a7e377d9a3bcb28ab15b72807ece1796cf9c02fc74e7ce7318cc3.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 16u) { return; } +out[i] = f32(b0[((i / 16u) % 1u) * 16u + ((i / 1u) % 16u) * 1u + 0u * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f276cf6656b999339df464b61934487f5c8eb7a0957cf6d902254ab1aa2ffcd0.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f276cf6656b999339df464b61934487f5c8eb7a0957cf6d902254ab1aa2ffcd0.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..15382fc197ef648b58ce80bf82efcabed936514c --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f276cf6656b999339df464b61934487f5c8eb7a0957cf6d902254ab1aa2ffcd0.wgsl @@ -0,0 +1,12 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 2048u; +if (i >= 2048u) { return; } +let coord = (i / 1u) % 256u; +if (coord >= 0u && coord < 64u) { out[i] = f32(b0[(i / 256u) * 64u + (coord - 0u) * 1u + i % 1u]); } +if (coord >= 64u && coord < 256u) { out[i] = f32(b1[(i / 256u) * 192u + (coord - 64u) * 1u + i % 1u]); } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f43cdd9c75fc256957b1b144a0587e8c8ddea861f991e1cdb6c9098e8b5d6ed7.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f43cdd9c75fc256957b1b144a0587e8c8ddea861f991e1cdb6c9098e8b5d6ed7.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..b2f14ad5c01e9efb2e5f7424d8124543f564a86d --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f43cdd9c75fc256957b1b144a0587e8c8ddea861f991e1cdb6c9098e8b5d6ed7.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 64u; +if (i >= 64u) { return; } +out[i] = f32(b0[((i / 64u) % 1u) * 128u + ((i / 32u) % 2u) * 64u + ((i / 32u) % 1u) * 64u + (((i / 1u) % 32u) * 1u + 0u) * 1u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f5b51b1e0e52bd87a70a0ca169186ebb5734cbdc2b8f375ba0241485b1df4e66.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f5b51b1e0e52bd87a70a0ca169186ebb5734cbdc2b8f375ba0241485b1df4e66.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..92790934326cbc014348566d54624e3d2d8a85c1 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f5b51b1e0e52bd87a70a0ca169186ebb5734cbdc2b8f375ba0241485b1df4e66.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 512u; +if (i >= 512u) { return; } +out[i] = f32(f32(b0[i]) + (f32(b1[i]) * 1.0)); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f6ccc1a43bc7cbaab21e08ce5de522f383236035c835289025961d70c49c60e8.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f6ccc1a43bc7cbaab21e08ce5de522f383236035c835289025961d70c49c60e8.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..6af5f79d6fa16007e949c9cac3ac0f2bc1a6910e --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f6ccc1a43bc7cbaab21e08ce5de522f383236035c835289025961d70c49c60e8.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 4194240u; +if (i >= 8388608u) { return; } +out[i] = f32(b0[i]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f73912607dffcf0ee27e9e1e1d342a3bef8f998e02cd06a8352524ef35a96d17.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f73912607dffcf0ee27e9e1e1d342a3bef8f998e02cd06a8352524ef35a96d17.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..e5eca2c179c0804dcd0c49c39d7427ce2673747c --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f73912607dffcf0ee27e9e1e1d342a3bef8f998e02cd06a8352524ef35a96d17.wgsl @@ -0,0 +1,17 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 32768u; +if (i >= 32768u) { return; } +let batch = i / 4096u; +let row = (i / 4096u) % 1u; +let col = i % 4096u; +var acc = 0.0; +for (var p = 0u; p < 256u; p++) { + acc += f32(b0[(((batch / 1u) % 8u) * 256u) + row * 256u + p]) * f32(b1[(((batch / 1u) % 8u) * 1048576u) + p * 4096u + col]); +} +out[i] = f32(acc); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f84306aa81033babc4c37ecfb91a411dba2ccee17202615deefb2206add106e2.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f84306aa81033babc4c37ecfb91a411dba2ccee17202615deefb2206add106e2.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..58968344810b5ec1eb718ef062079a8fd93b02a0 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f84306aa81033babc4c37ecfb91a411dba2ccee17202615deefb2206add106e2.wgsl @@ -0,0 +1,9 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var out: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { +let i = gid.x + gid.y * 6144u; +if (i >= 6144u) { return; } +out[i] = f32(b0[((i / 6144u) % 1u) * 6144u + ((i / 1u) % 6144u) * 1u + ((i / 1u) % 1u) * 6144u]); +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f994142054bcb47d6a554524dd8de9c3054467a1d07011027c47f475d0909c19.wgsl b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f994142054bcb47d6a554524dd8de9c3054467a1d07011027c47f475d0909c19.wgsl new file mode 100644 index 0000000000000000000000000000000000000000..cf15e04cdeba75c873e80fc764f945a9f56b3a48 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/f994142054bcb47d6a554524dd8de9c3054467a1d07011027c47f475d0909c19.wgsl @@ -0,0 +1,24 @@ +@group(0) @binding(0) var b0: array; +@group(0) @binding(1) var b1: array; +@group(0) @binding(2) var out: array; + +var factor: f32; +@compute @workgroup_size(64) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_id) local: vec3) { + let row = group.x + group.y * 8u; + if (row >= 8u) { return; } + let lane = local.x; + if (lane == 0u) { + var total = 0.0; + for (var j = 0u; j < 256u; j++) { + let v = f32(b0[row * 256u + j]); + total += v * v; + } + factor = inverseSqrt(total / 256.0 + 1e-06); + } + workgroupBarrier(); + for (var p = lane; p < 256u; p += 64u) { + let i = row * 256u + p; + out[i] = f32(f32(b0[row * 256u + p]) * factor * (f32(b1[p]) + 1.0)); + } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/manifest.json b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/manifest.json new file mode 100644 index 0000000000000000000000000000000000000000..367becf4fb9935ce6c7cdcf83ec511f1cc6b4a07 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/manifest.json @@ -0,0 +1,829 @@ +{ + "format": "webtorch", + "version": 1, + "producer": { + "name": "webtorch", + "version": "0.1.0a1", + "torch": "2.14.0" + }, + "entrypoints": { + "forward": "graph.json" + }, + "requiredFeatures": [ + "shader-f16" + ], + "assets": [ + { + "path": "weights/00000.bin", + "byteLength": 67108864, + "sha256": "998effe0875f489e58eb9592d73a2a2ef51b627520bd119524651c8ffc886282" + }, + { + "path": "weights/00001.bin", + "byteLength": 67108864, + "sha256": "7b669e08c1c1103797bd533fad1d52fe36b3f1293db9d09065ec24c7c5f2db8d" + }, + { + "path": "weights/00002.bin", + "byteLength": 67108864, + "sha256": "a281c127a0a9bf7f5498534c1a350d922f251e19343c303e70663486d609260e" + }, + { + "path": "weights/00003.bin", + "byteLength": 67108864, + "sha256": "1b4b102afabd0edfc06a42c1c14f6cc459e9e1e2cf0b8193fc5aa3046c04cc76" + }, + { + "path": "weights/00004.bin", + "byteLength": 67108864, + "sha256": "6480b9ca38668264d75fbc1ecdd78870fe1a8b95835f485c5543d11e2abbbe1f" + }, + { + "path": "weights/00005.bin", + "byteLength": 67108864, + "sha256": "c8f2a7532918fb3acc76952b3e685030c001abdae0b74ab41fc84c64a4ab8c2e" + }, + { + "path": "weights/00006.bin", + "byteLength": 67108864, + "sha256": "54fa18728d34cb3f67db0c70064574e8f34b786d9bf5d92e61241472e1249aa4" + }, + { + "path": "weights/00007.bin", + "byteLength": 67108864, + "sha256": "89edafc7a5ed368e67498221df821521c36df98216fcf4e55efbd3d8e2925725" + }, + { + "path": "weights/00008.bin", + "byteLength": 67108864, + "sha256": "273556893c277e1d8273045bcdc0943a4fb93d995655a870309dd8ea7f6b56d0" + }, + { + "path": "weights/00009.bin", + "byteLength": 67108864, + "sha256": "9207ae6e82ef7ce44d4057eb208a0c2f35ab06d2dd571aee8b16191116133536" + }, + { + "path": "weights/00010.bin", + "byteLength": 67108864, + "sha256": "7a64f841b6a5d951d488428e0e6eb5062717935cbaf8194931797fd8fb3eec53" + }, + { + "path": "weights/00011.bin", + "byteLength": 67108864, + "sha256": "718459e02d1bd28ee721759d716fc6e124523709e4bb3f7b60b36b7d206c63ec" + }, + { + "path": "weights/00012.bin", + "byteLength": 67108864, + "sha256": "2ee07dac007a616667f2c3585641332e7df05b34b6ef0e6f0904eb7a029009ad" + }, + { + "path": "weights/00013.bin", + "byteLength": 67108864, + "sha256": "e054e1b2605cd22d63e477299218d5426b042968acd0d53942d01ec5a29db3cc" + }, + { + "path": "weights/00014.bin", + "byteLength": 67108864, + "sha256": "570d6d4ff3455dcb089501e0a30efb71a00e8ed835478e6fa36348d4f3a6dc26" + }, + { + "path": "weights/00015.bin", + "byteLength": 67108864, + "sha256": "de850adc45c0f4b78b99be441605d5c843e1f2a167b017b030cbc897654ebee9" + }, + { + "path": "weights/00016.bin", + "byteLength": 67108864, + "sha256": "01dfe7fb384c6016183de73de8cfb43ef410b697391103f3fd7c269983706651" + }, + { + "path": "weights/00017.bin", + "byteLength": 67108864, + "sha256": "aedc4a87a0492afaab3aeabe395ae647f9d260b3324f5c56fb6079daae9283e3" + }, + { + "path": "weights/00018.bin", + "byteLength": 67108864, + "sha256": "a2b97aff1515a8441f067b40b71eb10ffeb1e60ff0e94579cd1d2b9b10514683" + }, + { + "path": "weights/00019.bin", + "byteLength": 67108864, + "sha256": "58b262bf15ce55ec421276df579123f6cfcb4025b27dd86c26f0a02a18ad766b" + }, + { + "path": "weights/00020.bin", + "byteLength": 20475264, + "sha256": "d99fbd379a8c7234b904b0375c9da6b2d212e75466d0bf6eaee7d65cba91ff12" + }, + { + "path": "tokenizer/tokenizer.json", + "byteLength": 12807982, + "sha256": "5f9e4d4901a92b997e463c1f46055088b6cca5ca61a6522d1b9f64c4bb81cb42" + }, + { + "path": "tokenizer/tokenizer_config.json", + "byteLength": 16709, + "sha256": "49e2b6e395f959f077f1e992b338919c0d4a9732fc6e613995e06557f843500c" + }, + { + "path": "tokenizer/chat_template.jinja", + "byteLength": 7755, + "sha256": "273d8e0e683b885071fb17e08d71e5f2a5ddfb5309756181681de4f5a1822d80" + }, + { + "path": "source/config.json", + "byteLength": 2907, + "sha256": "b90b86f35c8e6925ef74ee04d0e758f0a845c83a42089ad82bbaa948de9b4204" + }, + { + "path": "source/LICENSE", + "byteLength": 11544, + "sha256": "bbedc3fda3305820b977265f01b8619d87570a6739de3a5582c3464840f1e57a" + }, + { + "path": "source/provenance.json", + "byteLength": 201, + "sha256": "cc031c3b7563aac8d5663b31e2f4fc36eff716921d0770718de26f43d2d883c5" + }, + { + "path": "graph.json", + "byteLength": 10577701, + "sha256": "78123794a111311b0bdafead59f8fa1b93b71c31cc1b293c1c11221e3ccc33de" + }, + { + "path": "kernels/04f26661cc9ac73bb7f39f66b76c4be31ee8774e094eeb41ed7d677e1b1d3fa9.wgsl", + "byteLength": 368, + "sha256": "04f26661cc9ac73bb7f39f66b76c4be31ee8774e094eeb41ed7d677e1b1d3fa9" + }, + { + "path": "kernels/0554a443ecf6ead32fd2d58178b9ce3c592c89a09b3f63fa509ee947323bb497.wgsl", + "byteLength": 399, + "sha256": "0554a443ecf6ead32fd2d58178b9ce3c592c89a09b3f63fa509ee947323bb497" + }, + { + "path": "kernels/0680219da63feb9bb85653407739c5e5cc93ed6ca0971c1b03c512068d97f69f.wgsl", + "byteLength": 858, + "sha256": "0680219da63feb9bb85653407739c5e5cc93ed6ca0971c1b03c512068d97f69f" + }, + { + "path": "kernels/0835fa80eaadb7c0e6d677e1fc622ad5dcedabb431e5d312f3b7aeb2252a7c7c.wgsl", + "byteLength": 231, + "sha256": "0835fa80eaadb7c0e6d677e1fc622ad5dcedabb431e5d312f3b7aeb2252a7c7c" + }, + { + "path": "kernels/0e5468583a1ba929483119c50d4be42a58aed4dde1e15cc231908ed7543d0b71.wgsl", + "byteLength": 288, + "sha256": "0e5468583a1ba929483119c50d4be42a58aed4dde1e15cc231908ed7543d0b71" + }, + { + "path": "kernels/0e9b6ee5f92769faaee5c746d1a3b936cdb7731f3f8e70cbf0a3fe96f6c8d921.wgsl", + "byteLength": 397, + "sha256": "0e9b6ee5f92769faaee5c746d1a3b936cdb7731f3f8e70cbf0a3fe96f6c8d921" + }, + { + "path": "kernels/0f466ed2664d77386f48afd19d8ab2734312bfd4f5f4d5b063d801a0b0a97cd9.wgsl", + "byteLength": 396, + "sha256": "0f466ed2664d77386f48afd19d8ab2734312bfd4f5f4d5b063d801a0b0a97cd9" + }, + { + "path": "kernels/0fd78c192784ba4e545091262cf1bac47ef69ca387b850f97b9de39dd0a2a741.wgsl", + "byteLength": 1017, + "sha256": "0fd78c192784ba4e545091262cf1bac47ef69ca387b850f97b9de39dd0a2a741" + }, + { + "path": "kernels/10feb208ed0549389f54419ff8985bee4778276913d4c6a613a8b120a809ab5e.wgsl", + "byteLength": 1025, + "sha256": "10feb208ed0549389f54419ff8985bee4778276913d4c6a613a8b120a809ab5e" + }, + { + "path": "kernels/12155d86d776e73f23d3bff0268c60dccb54c209ee9ee3a8cfd227ec4aa5b6df.wgsl", + "byteLength": 290, + "sha256": "12155d86d776e73f23d3bff0268c60dccb54c209ee9ee3a8cfd227ec4aa5b6df" + }, + { + "path": "kernels/1255847146b31732939e573cdb26c483358c1ce07f987bd8c3d3d69ae8c7e141.wgsl", + "byteLength": 380, + "sha256": "1255847146b31732939e573cdb26c483358c1ce07f987bd8c3d3d69ae8c7e141" + }, + { + "path": "kernels/13237b557466d2fef5bfbc74e7ac77219b9821a53bc0843dbb5d14db8b04d414.wgsl", + "byteLength": 308, + "sha256": "13237b557466d2fef5bfbc74e7ac77219b9821a53bc0843dbb5d14db8b04d414" + }, + { + "path": "kernels/13dc73935b9b3b978232043e13babee3877bd53ca572be440de0590dacce8b17.wgsl", + "byteLength": 367, + "sha256": "13dc73935b9b3b978232043e13babee3877bd53ca572be440de0590dacce8b17" + }, + { + "path": "kernels/15b4ee3a9b5df59d5213e59951d8874f4eb4b9132e04fc2febddad84b44277f5.wgsl", + "byteLength": 1021, + "sha256": "15b4ee3a9b5df59d5213e59951d8874f4eb4b9132e04fc2febddad84b44277f5" + }, + { + "path": "kernels/17ec6bdaf98edf7657b74a25036c2bec806ecbb4280f33fff6afd2e0de5ed377.wgsl", + "byteLength": 317, + "sha256": "17ec6bdaf98edf7657b74a25036c2bec806ecbb4280f33fff6afd2e0de5ed377" + }, + { + "path": "kernels/1ad205080cae29142a3ad87343a0cf3ac4945455398293946110d421ce6d515e.wgsl", + "byteLength": 371, + "sha256": "1ad205080cae29142a3ad87343a0cf3ac4945455398293946110d421ce6d515e" + }, + { + "path": "kernels/20fb9e0d0bd2aab34b7f69eac259898e68cac4a0d6b0bda5e95c9126a660336d.wgsl", + "byteLength": 395, + "sha256": "20fb9e0d0bd2aab34b7f69eac259898e68cac4a0d6b0bda5e95c9126a660336d" + }, + { + "path": "kernels/222a2645e3e02d0a4063a63ca1b6fa4c89faee6b76905b20de4dd54ab3077ea3.wgsl", + "byteLength": 285, + "sha256": "222a2645e3e02d0a4063a63ca1b6fa4c89faee6b76905b20de4dd54ab3077ea3" + }, + { + "path": "kernels/234987ec68497956ea7aef8e3f44c8a73dbe4c1519165d123946bedbb4d18e26.wgsl", + "byteLength": 380, + "sha256": "234987ec68497956ea7aef8e3f44c8a73dbe4c1519165d123946bedbb4d18e26" + }, + { + "path": "kernels/26cf6eb3b486d1652247ff130b05ed0e8c8560028eeba59acfc075b976bfdfce.wgsl", + "byteLength": 549, + "sha256": "26cf6eb3b486d1652247ff130b05ed0e8c8560028eeba59acfc075b976bfdfce" + }, + { + "path": "kernels/2711a4a396071f4e99303d01d4854ccaa37a445f4b718ebd7b61cf1c37e56b48.wgsl", + "byteLength": 355, + "sha256": "2711a4a396071f4e99303d01d4854ccaa37a445f4b718ebd7b61cf1c37e56b48" + }, + { + "path": "kernels/2bdeb3590b1ba05cf6e9c7e8245f8aa552a12dae84f9578c603c9cee081a82d4.wgsl", + "byteLength": 847, + "sha256": "2bdeb3590b1ba05cf6e9c7e8245f8aa552a12dae84f9578c603c9cee081a82d4" + }, + { + "path": "kernels/2c0cca2a7b0198c317e68992e64f3fe9bba80b67204c2695c20c372471c150d3.wgsl", + "byteLength": 473, + "sha256": "2c0cca2a7b0198c317e68992e64f3fe9bba80b67204c2695c20c372471c150d3" + }, + { + "path": "kernels/2dd29b61b039fcdd7fbc3c94f629b24b6f81aa5b9923098837597359b295e40d.wgsl", + "byteLength": 396, + "sha256": "2dd29b61b039fcdd7fbc3c94f629b24b6f81aa5b9923098837597359b295e40d" + }, + { + "path": "kernels/2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a.wgsl", + "byteLength": 297, + "sha256": "2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a" + }, + { + "path": "kernels/32a5027cac5dd969d63fe5d5c88d7c87abb600c43ad1d70e1c56615ea781eff3.wgsl", + "byteLength": 422, + "sha256": "32a5027cac5dd969d63fe5d5c88d7c87abb600c43ad1d70e1c56615ea781eff3" + }, + { + "path": "kernels/38bd527686bcade9ed19a52788ccd711ca668e76854cc6274c480860fce88ff2.wgsl", + "byteLength": 382, + "sha256": "38bd527686bcade9ed19a52788ccd711ca668e76854cc6274c480860fce88ff2" + }, + { + "path": "kernels/3bb850bda95bf77706a2839616a401b773f9de0428d1b427e61fb04e2c35fe74.wgsl", + "byteLength": 493, + "sha256": "3bb850bda95bf77706a2839616a401b773f9de0428d1b427e61fb04e2c35fe74" + }, + { + "path": "kernels/3fa1efae6d90015ba69cb40cef009f49ad7f7a960018fb92a962707deff4f09e.wgsl", + "byteLength": 286, + "sha256": "3fa1efae6d90015ba69cb40cef009f49ad7f7a960018fb92a962707deff4f09e" + }, + { + "path": "kernels/413a0717b830335decd89485fab5fbea4efeb71338961ba7a2d1221eaa45459a.wgsl", + "byteLength": 1025, + "sha256": "413a0717b830335decd89485fab5fbea4efeb71338961ba7a2d1221eaa45459a" + }, + { + "path": "kernels/42c5af4838daf75fff9d4bb9a5df409511c1d6daed617d05eaf033a977503eac.wgsl", + "byteLength": 548, + "sha256": "42c5af4838daf75fff9d4bb9a5df409511c1d6daed617d05eaf033a977503eac" + }, + { + "path": "kernels/4540d87eaec90c9eb42badfe0abfa70252f26b6e7de4dd31ee6903fe50e96bd1.wgsl", + "byteLength": 377, + "sha256": "4540d87eaec90c9eb42badfe0abfa70252f26b6e7de4dd31ee6903fe50e96bd1" + }, + { + "path": "kernels/4ac112c035a1a4a2f5deae66fac0969000c4666730c2d7b15cb15a157ea0e9b8.wgsl", + "byteLength": 395, + "sha256": "4ac112c035a1a4a2f5deae66fac0969000c4666730c2d7b15cb15a157ea0e9b8" + }, + { + "path": "kernels/4b1a7529fd5c57d5d4593d4af58a9a3b545c7d106d18251f6452f8a562b97952.wgsl", + "byteLength": 389, + "sha256": "4b1a7529fd5c57d5d4593d4af58a9a3b545c7d106d18251f6452f8a562b97952" + }, + { + "path": "kernels/4be963f9a4bb928764fd0efe02769e220b0e69defb2f8b22ab28e41a90d3fc14.wgsl", + "byteLength": 286, + "sha256": "4be963f9a4bb928764fd0efe02769e220b0e69defb2f8b22ab28e41a90d3fc14" + }, + { + "path": "kernels/4d3e5f3a99a6c8c9dba5de30be4302f1e442f91d097437e825df97017e389a42.wgsl", + "byteLength": 367, + "sha256": "4d3e5f3a99a6c8c9dba5de30be4302f1e442f91d097437e825df97017e389a42" + }, + { + "path": "kernels/4f48f08a8aeb51cc47a10db076a46522807d3b5d35952c9718ce1c1cd760681f.wgsl", + "byteLength": 415, + "sha256": "4f48f08a8aeb51cc47a10db076a46522807d3b5d35952c9718ce1c1cd760681f" + }, + { + "path": "kernels/4fad29e3e337840e25132f10b3e3d55225f727ae13525ff0b0bdf31e03f09d74.wgsl", + "byteLength": 305, + "sha256": "4fad29e3e337840e25132f10b3e3d55225f727ae13525ff0b0bdf31e03f09d74" + }, + { + "path": "kernels/51fd4d798491c6a80daff4c18b0495926b7f62269dfc3a78a327faa52ba2a759.wgsl", + "byteLength": 307, + "sha256": "51fd4d798491c6a80daff4c18b0495926b7f62269dfc3a78a327faa52ba2a759" + }, + { + "path": "kernels/53d6eb0b5e6d49bfd3bb9fd1b781e13317adf528e047da2416e95f5a3a13b1b4.wgsl", + "byteLength": 381, + "sha256": "53d6eb0b5e6d49bfd3bb9fd1b781e13317adf528e047da2416e95f5a3a13b1b4" + }, + { + "path": "kernels/552c2f036a125363ca85fd7ea5b1a84465b562dc62ddfffdbaaebf6e3288d321.wgsl", + "byteLength": 663, + "sha256": "552c2f036a125363ca85fd7ea5b1a84465b562dc62ddfffdbaaebf6e3288d321" + }, + { + "path": "kernels/55b9c48ba429daf80bbae0b4737f895b475abdc9ddfdc88ab2f894f133ec2bf9.wgsl", + "byteLength": 326, + "sha256": "55b9c48ba429daf80bbae0b4737f895b475abdc9ddfdc88ab2f894f133ec2bf9" + }, + { + "path": "kernels/579685c7f8ca260bd53d1677c43d5fc9a0133c81c7442ffb31daf0c3759267c4.wgsl", + "byteLength": 1025, + "sha256": "579685c7f8ca260bd53d1677c43d5fc9a0133c81c7442ffb31daf0c3759267c4" + }, + { + "path": "kernels/57c5796344c26b430493a3a63dc734d4206380ba89484294f50f67d615eefcf8.wgsl", + "byteLength": 663, + "sha256": "57c5796344c26b430493a3a63dc734d4206380ba89484294f50f67d615eefcf8" + }, + { + "path": "kernels/57dc406322bf3badd8490493ef924bd7e7e0afd542d2d6e409b4d8a3b4b784f7.wgsl", + "byteLength": 286, + "sha256": "57dc406322bf3badd8490493ef924bd7e7e0afd542d2d6e409b4d8a3b4b784f7" + }, + { + "path": "kernels/59ea88d691b919af03bbb26f0253decfbc35d74f9661cdd3a08570ae85c313d2.wgsl", + "byteLength": 379, + "sha256": "59ea88d691b919af03bbb26f0253decfbc35d74f9661cdd3a08570ae85c313d2" + }, + { + "path": "kernels/5aa3725287640d289322ab469bd924bd7797f1cef28920ed4526783877c3cf6b.wgsl", + "byteLength": 373, + "sha256": "5aa3725287640d289322ab469bd924bd7797f1cef28920ed4526783877c3cf6b" + }, + { + "path": "kernels/5adf40822be73f2c85ed1c73c0bd106707e33fcb0e8d1e82a643474b8bfeceeb.wgsl", + "byteLength": 1029, + "sha256": "5adf40822be73f2c85ed1c73c0bd106707e33fcb0e8d1e82a643474b8bfeceeb" + }, + { + "path": "kernels/5c22e6f60217881c0351e38d19057626a6c2292771c1a83816995245777ac2fd.wgsl", + "byteLength": 397, + "sha256": "5c22e6f60217881c0351e38d19057626a6c2292771c1a83816995245777ac2fd" + }, + { + "path": "kernels/5c9a5ae6121b3b10611e09e5f1bc563b70191f2506a2afc0f32de823a9d2dac4.wgsl", + "byteLength": 554, + "sha256": "5c9a5ae6121b3b10611e09e5f1bc563b70191f2506a2afc0f32de823a9d2dac4" + }, + { + "path": "kernels/5eac5f92aa4d18895369b11adb99ac0d9639d8b615f7240cffe6f6be7046e54a.wgsl", + "byteLength": 845, + "sha256": "5eac5f92aa4d18895369b11adb99ac0d9639d8b615f7240cffe6f6be7046e54a" + }, + { + "path": "kernels/5f4403d78557e88e66d71eb5e53746cc07449a7c421b57eda4f093a09a68a99a.wgsl", + "byteLength": 290, + "sha256": "5f4403d78557e88e66d71eb5e53746cc07449a7c421b57eda4f093a09a68a99a" + }, + { + "path": "kernels/5f657afde3ab80a1fc4b6739c0b702bc5bd60e52630762e66b210fbb8b42b52a.wgsl", + "byteLength": 403, + "sha256": "5f657afde3ab80a1fc4b6739c0b702bc5bd60e52630762e66b210fbb8b42b52a" + }, + { + "path": "kernels/5fd6e068841f5c054392bef564909dbf086bae9e4030f3625d1143244a08cbbf.wgsl", + "byteLength": 495, + "sha256": "5fd6e068841f5c054392bef564909dbf086bae9e4030f3625d1143244a08cbbf" + }, + { + "path": "kernels/67978e951813ede8016ebf553fb8d3b4ce9e15dd184810544879559208d6b7e1.wgsl", + "byteLength": 322, + "sha256": "67978e951813ede8016ebf553fb8d3b4ce9e15dd184810544879559208d6b7e1" + }, + { + "path": "kernels/6a045508377e47e95d411257071b7ed8f9097fd220c00de6a5167abb548ebe37.wgsl", + "byteLength": 750, + "sha256": "6a045508377e47e95d411257071b7ed8f9097fd220c00de6a5167abb548ebe37" + }, + { + "path": "kernels/6a7b2e3778a02c76155d46fa84aeb33f594954e82e09b6b7f2915c0a606eb3f5.wgsl", + "byteLength": 471, + "sha256": "6a7b2e3778a02c76155d46fa84aeb33f594954e82e09b6b7f2915c0a606eb3f5" + }, + { + "path": "kernels/6ae4c7f13978e60e3f4ccd5b8172af88d7dc24cfdb236e51b826f13a03d8cd02.wgsl", + "byteLength": 400, + "sha256": "6ae4c7f13978e60e3f4ccd5b8172af88d7dc24cfdb236e51b826f13a03d8cd02" + }, + { + "path": "kernels/7002f8fdf09f5124ce2875c3540bce3e77750260556fc1a00a8b13b6c1871c30.wgsl", + "byteLength": 383, + "sha256": "7002f8fdf09f5124ce2875c3540bce3e77750260556fc1a00a8b13b6c1871c30" + }, + { + "path": "kernels/7121b5c223da0b0f28991f55cfed3c3f3390caa0bf65a8e58ae380db281ceb6b.wgsl", + "byteLength": 368, + "sha256": "7121b5c223da0b0f28991f55cfed3c3f3390caa0bf65a8e58ae380db281ceb6b" + }, + { + "path": "kernels/7183997261247fa4df7e081c0d6d01036fee95b9b57ef944d5f44f72fc80d5fe.wgsl", + "byteLength": 285, + "sha256": "7183997261247fa4df7e081c0d6d01036fee95b9b57ef944d5f44f72fc80d5fe" + }, + { + "path": "kernels/7194f0a413c1d8c6013b45b42b129c0a7557141a31bf3f08946fa1dfd1071393.wgsl", + "byteLength": 851, + "sha256": "7194f0a413c1d8c6013b45b42b129c0a7557141a31bf3f08946fa1dfd1071393" + }, + { + "path": "kernels/720dc88131b6bf06e6b09185aa7d5986eb755dcfa60d0e040cbbf621c8f3f761.wgsl", + "byteLength": 559, + "sha256": "720dc88131b6bf06e6b09185aa7d5986eb755dcfa60d0e040cbbf621c8f3f761" + }, + { + "path": "kernels/742db031d819fdf5a756f37e8e824e2dcb2ca3f30efac3ffbd5918181059ae3a.wgsl", + "byteLength": 393, + "sha256": "742db031d819fdf5a756f37e8e824e2dcb2ca3f30efac3ffbd5918181059ae3a" + }, + { + "path": "kernels/7ea4a12734f1206e9c94a060bf6b24732f46685f0c46a78c3c5fbcd64489d8c3.wgsl", + "byteLength": 407, + "sha256": "7ea4a12734f1206e9c94a060bf6b24732f46685f0c46a78c3c5fbcd64489d8c3" + }, + { + "path": "kernels/7fab02b2d46f9419002e2fc6f6e0e4e745ddb737a2c9e7060e45e7125c1be187.wgsl", + "byteLength": 378, + "sha256": "7fab02b2d46f9419002e2fc6f6e0e4e745ddb737a2c9e7060e45e7125c1be187" + }, + { + "path": "kernels/7fe183583b43b627721ba6853b306288d7ceea4b42b630410f47ea9637220a4b.wgsl", + "byteLength": 290, + "sha256": "7fe183583b43b627721ba6853b306288d7ceea4b42b630410f47ea9637220a4b" + }, + { + "path": "kernels/8898353d0c367a3982a48d416f6777c3ce7a854572eeaf55843f3e7ee458c65d.wgsl", + "byteLength": 292, + "sha256": "8898353d0c367a3982a48d416f6777c3ce7a854572eeaf55843f3e7ee458c65d" + }, + { + "path": "kernels/88b50ebddd13f3ff5e2819b8bc956e8fc1346e8f75cacbc3d0edf031b7f1f248.wgsl", + "byteLength": 306, + "sha256": "88b50ebddd13f3ff5e2819b8bc956e8fc1346e8f75cacbc3d0edf031b7f1f248" + }, + { + "path": "kernels/896006749045a93d89dca7ed7326aa05c5ff29423400569d34ab2256dea6b0e4.wgsl", + "byteLength": 383, + "sha256": "896006749045a93d89dca7ed7326aa05c5ff29423400569d34ab2256dea6b0e4" + }, + { + "path": "kernels/898086c99e03b2cdb78e3e5658c3fff202f1c87d09d8da08659bd1ba85e01b09.wgsl", + "byteLength": 554, + "sha256": "898086c99e03b2cdb78e3e5658c3fff202f1c87d09d8da08659bd1ba85e01b09" + }, + { + "path": "kernels/8a46c4e24d014fd6af88a3f370c4b6161be7c63b6c30211bd62af3dbd87293a8.wgsl", + "byteLength": 398, + "sha256": "8a46c4e24d014fd6af88a3f370c4b6161be7c63b6c30211bd62af3dbd87293a8" + }, + { + "path": "kernels/8d0a37e497bcb37e2d5a1746fba59fd69680ed8eeada4835c5d0eeef840d57ab.wgsl", + "byteLength": 303, + "sha256": "8d0a37e497bcb37e2d5a1746fba59fd69680ed8eeada4835c5d0eeef840d57ab" + }, + { + "path": "kernels/9144d827dae4c52b80abb829178a9ed901ce5dbc6702ad94b0b355e8eaad5ea5.wgsl", + "byteLength": 548, + "sha256": "9144d827dae4c52b80abb829178a9ed901ce5dbc6702ad94b0b355e8eaad5ea5" + }, + { + "path": "kernels/93d02cadeb259e826a8d196998a25cbb09a7053b0fa50174dc6c0d9ec05f8d6d.wgsl", + "byteLength": 1025, + "sha256": "93d02cadeb259e826a8d196998a25cbb09a7053b0fa50174dc6c0d9ec05f8d6d" + }, + { + "path": "kernels/947f5cc8b3f87e1a0d4d322d9d9e71e2ed16ce6e1f7c91118678c164dc1c8752.wgsl", + "byteLength": 381, + "sha256": "947f5cc8b3f87e1a0d4d322d9d9e71e2ed16ce6e1f7c91118678c164dc1c8752" + }, + { + "path": "kernels/95389a9d24d01ffc452bc9cea28e01f8bbca2202271cbe96b1399afbbe725534.wgsl", + "byteLength": 1158, + "sha256": "95389a9d24d01ffc452bc9cea28e01f8bbca2202271cbe96b1399afbbe725534" + }, + { + "path": "kernels/95addfdb139df33bb3d5c05859412434d318cee0340560537dce4d56a05e159c.wgsl", + "byteLength": 395, + "sha256": "95addfdb139df33bb3d5c05859412434d318cee0340560537dce4d56a05e159c" + }, + { + "path": "kernels/9625ca0b071511e5261877fd21627ca354ea0219da86a10887e8054afa15dcfd.wgsl", + "byteLength": 367, + "sha256": "9625ca0b071511e5261877fd21627ca354ea0219da86a10887e8054afa15dcfd" + }, + { + "path": "kernels/985bb9ad33654b9adbd4a256d626e346556a10dad514692800da8ebd7910f0f0.wgsl", + "byteLength": 362, + "sha256": "985bb9ad33654b9adbd4a256d626e346556a10dad514692800da8ebd7910f0f0" + }, + { + "path": "kernels/9f02cd7a8eb5e6d20c2b01559fc58db5b0447ae41a51e2cd479ed519fc284791.wgsl", + "byteLength": 947, + "sha256": "9f02cd7a8eb5e6d20c2b01559fc58db5b0447ae41a51e2cd479ed519fc284791" + }, + { + "path": "kernels/a1a974f6d8d51fd7c3496ede0966af303ad9f0c0c9ee44cf93ec962d2b0e87e1.wgsl", + "byteLength": 303, + "sha256": "a1a974f6d8d51fd7c3496ede0966af303ad9f0c0c9ee44cf93ec962d2b0e87e1" + }, + { + "path": "kernels/a37dac611b27158f19aa5bd70c43217286e610e733cbd602b91666b4af42d3ef.wgsl", + "byteLength": 285, + "sha256": "a37dac611b27158f19aa5bd70c43217286e610e733cbd602b91666b4af42d3ef" + }, + { + "path": "kernels/a672c3e158716bde7444c8141f3df1da16da739ccdd2eb2ad60f9d42d0775157.wgsl", + "byteLength": 505, + "sha256": "a672c3e158716bde7444c8141f3df1da16da739ccdd2eb2ad60f9d42d0775157" + }, + { + "path": "kernels/a77f14c89662d2d3da8abd8f109667793dc662d307e5e1c61123be753e25437b.wgsl", + "byteLength": 307, + "sha256": "a77f14c89662d2d3da8abd8f109667793dc662d307e5e1c61123be753e25437b" + }, + { + "path": "kernels/abdb17daf29310a7828a55a1efeebaa65b6553e236fc20702e8c712e0a46dd5b.wgsl", + "byteLength": 326, + "sha256": "abdb17daf29310a7828a55a1efeebaa65b6553e236fc20702e8c712e0a46dd5b" + }, + { + "path": "kernels/b02b0be460d581c720afd68c8b722e8da47e2ee432b5d61bbdf24b6084cb2f26.wgsl", + "byteLength": 415, + "sha256": "b02b0be460d581c720afd68c8b722e8da47e2ee432b5d61bbdf24b6084cb2f26" + }, + { + "path": "kernels/b199d835db70f888c7b99a3ac6e9ec56f6c47d5bba4eaf352e0a8eabed56411e.wgsl", + "byteLength": 387, + "sha256": "b199d835db70f888c7b99a3ac6e9ec56f6c47d5bba4eaf352e0a8eabed56411e" + }, + { + "path": "kernels/b2b16798cbe94c919025d76b0435afe47966ffdd1ea53ec74a791c03c4709a28.wgsl", + "byteLength": 380, + "sha256": "b2b16798cbe94c919025d76b0435afe47966ffdd1ea53ec74a791c03c4709a28" + }, + { + "path": "kernels/b3128f8b5ac3b671947588c52e53f26426877fd37a2e7bdfcd9819a6cbb7e3a7.wgsl", + "byteLength": 418, + "sha256": "b3128f8b5ac3b671947588c52e53f26426877fd37a2e7bdfcd9819a6cbb7e3a7" + }, + { + "path": "kernels/b57928140994202536cd4f99d2e14f5dba9754f014843673790d5d48c9cbb9ab.wgsl", + "byteLength": 390, + "sha256": "b57928140994202536cd4f99d2e14f5dba9754f014843673790d5d48c9cbb9ab" + }, + { + "path": "kernels/b71488c9b47fa85c9dcdbbf9185c31ab819169ab631e1aa901f65e2999c5d22f.wgsl", + "byteLength": 1025, + "sha256": "b71488c9b47fa85c9dcdbbf9185c31ab819169ab631e1aa901f65e2999c5d22f" + }, + { + "path": "kernels/b75a05038c8d340aa61c4cdf617403ed9df75a05ecb9284c659b456ee9e120cd.wgsl", + "byteLength": 652, + "sha256": "b75a05038c8d340aa61c4cdf617403ed9df75a05ecb9284c659b456ee9e120cd" + }, + { + "path": "kernels/bb9029b667eeb13670d1f640cd677c6eabd7a7569686c599e3eeb9482d642b00.wgsl", + "byteLength": 373, + "sha256": "bb9029b667eeb13670d1f640cd677c6eabd7a7569686c599e3eeb9482d642b00" + }, + { + "path": "kernels/bbe7cf742bbc4d053e33018f81ca9e71a5f66cf38db88df0fe1b4584b27641f0.wgsl", + "byteLength": 1029, + "sha256": "bbe7cf742bbc4d053e33018f81ca9e71a5f66cf38db88df0fe1b4584b27641f0" + }, + { + "path": "kernels/c0a7f99d0ce8f9a199b9baa52effe47afa9605f77405c32b32d132be1f7404f1.wgsl", + "byteLength": 286, + "sha256": "c0a7f99d0ce8f9a199b9baa52effe47afa9605f77405c32b32d132be1f7404f1" + }, + { + "path": "kernels/c0b2947798baa198598d75ca0e5dcd31a9fa320499665e34f7bb81d27fc15fb0.wgsl", + "byteLength": 307, + "sha256": "c0b2947798baa198598d75ca0e5dcd31a9fa320499665e34f7bb81d27fc15fb0" + }, + { + "path": "kernels/c23bee3c3664172bf0a078a96acd2b3ad3136732234975f757d313e696c2c37a.wgsl", + "byteLength": 381, + "sha256": "c23bee3c3664172bf0a078a96acd2b3ad3136732234975f757d313e696c2c37a" + }, + { + "path": "kernels/c295fd955666232c92869895d186ebd33a25e4a5bf14df41d5b455fbe304d007.wgsl", + "byteLength": 382, + "sha256": "c295fd955666232c92869895d186ebd33a25e4a5bf14df41d5b455fbe304d007" + }, + { + "path": "kernels/c2c806bbe93f09247867b4ace45489a5f2fe3ee67f265418eb7aa71a29880276.wgsl", + "byteLength": 315, + "sha256": "c2c806bbe93f09247867b4ace45489a5f2fe3ee67f265418eb7aa71a29880276" + }, + { + "path": "kernels/c4862be63d2532b6891da039edbef2f46737a34518f8a89602821e00447b5933.wgsl", + "byteLength": 385, + "sha256": "c4862be63d2532b6891da039edbef2f46737a34518f8a89602821e00447b5933" + }, + { + "path": "kernels/c72041e82befee6919af48775291256e22a43dc6445ef0aa288db6b96585a2a9.wgsl", + "byteLength": 286, + "sha256": "c72041e82befee6919af48775291256e22a43dc6445ef0aa288db6b96585a2a9" + }, + { + "path": "kernels/c90428425f6d220b441e97e3b894adfb090869c98b17f53afbd2f9fb202b4525.wgsl", + "byteLength": 307, + "sha256": "c90428425f6d220b441e97e3b894adfb090869c98b17f53afbd2f9fb202b4525" + }, + { + "path": "kernels/cc250879ffbf43ba05860b5a86f6a824e8ca7b4c9f820137a68e1ba633e09738.wgsl", + "byteLength": 573, + "sha256": "cc250879ffbf43ba05860b5a86f6a824e8ca7b4c9f820137a68e1ba633e09738" + }, + { + "path": "kernels/ccb303d2d38ef10c1cf514f01ec2fa3b2541734aca8a34b3e8bb8d71d03a6688.wgsl", + "byteLength": 603, + "sha256": "ccb303d2d38ef10c1cf514f01ec2fa3b2541734aca8a34b3e8bb8d71d03a6688" + }, + { + "path": "kernels/cdc784b39918d10424edc08d2e44da0075863e8fd41ae362f5695a248363531d.wgsl", + "byteLength": 1406, + "sha256": "cdc784b39918d10424edc08d2e44da0075863e8fd41ae362f5695a248363531d" + }, + { + "path": "kernels/d4c80950a743190f3c54e4675a4fb139edf142a4737c55d803326c1a1cedcc2e.wgsl", + "byteLength": 661, + "sha256": "d4c80950a743190f3c54e4675a4fb139edf142a4737c55d803326c1a1cedcc2e" + }, + { + "path": "kernels/dc4ee648b97b273931e31c95b910d66f11e34a1cb036589e9ae9cc250b6e6291.wgsl", + "byteLength": 286, + "sha256": "dc4ee648b97b273931e31c95b910d66f11e34a1cb036589e9ae9cc250b6e6291" + }, + { + "path": "kernels/de276eae1430ed164ae977ce4b005301ea0f2652bed1f879acb3a355e600ff5e.wgsl", + "byteLength": 365, + "sha256": "de276eae1430ed164ae977ce4b005301ea0f2652bed1f879acb3a355e600ff5e" + }, + { + "path": "kernels/deee894ecb86334ff82870260fc5edbc36439c1d3bce09fab5708c03d699b390.wgsl", + "byteLength": 367, + "sha256": "deee894ecb86334ff82870260fc5edbc36439c1d3bce09fab5708c03d699b390" + }, + { + "path": "kernels/e01dc46d731ca9d3358814482da28b4edb8496243937be85ddde86a02a215169.wgsl", + "byteLength": 377, + "sha256": "e01dc46d731ca9d3358814482da28b4edb8496243937be85ddde86a02a215169" + }, + { + "path": "kernels/e23f8d0e93dd1cec22eea5c04988cc373462885008c37beb9e5f012773c6af1f.wgsl", + "byteLength": 1025, + "sha256": "e23f8d0e93dd1cec22eea5c04988cc373462885008c37beb9e5f012773c6af1f" + }, + { + "path": "kernels/e38b21c1b50403e478d2d4956df390560aa740de2dcba8d69322d996c998adf4.wgsl", + "byteLength": 408, + "sha256": "e38b21c1b50403e478d2d4956df390560aa740de2dcba8d69322d996c998adf4" + }, + { + "path": "kernels/e6571e9ffa79d4f7888d606a5845a6fe7d39e3b7944eafb979c162a384245c66.wgsl", + "byteLength": 385, + "sha256": "e6571e9ffa79d4f7888d606a5845a6fe7d39e3b7944eafb979c162a384245c66" + }, + { + "path": "kernels/e829b3bd2efe5b094dc36238e2a933a1e30c942efddfc80e79862a55c57e6e20.wgsl", + "byteLength": 459, + "sha256": "e829b3bd2efe5b094dc36238e2a933a1e30c942efddfc80e79862a55c57e6e20" + }, + { + "path": "kernels/ec606c53f8610fff07d8464d8f1cd2a909203cd4064d1be19ad9bdb42263794d.wgsl", + "byteLength": 596, + "sha256": "ec606c53f8610fff07d8464d8f1cd2a909203cd4064d1be19ad9bdb42263794d" + }, + { + "path": "kernels/edfa4ff396bf5cc3be28d9051639c4dda6fec1302d7a2bd34d89ec903af94af1.wgsl", + "byteLength": 307, + "sha256": "edfa4ff396bf5cc3be28d9051639c4dda6fec1302d7a2bd34d89ec903af94af1" + }, + { + "path": "kernels/eea033222d278c8aa8f4b0de36243adba4e5b6ce274d24a3dcd365ad42c19a38.wgsl", + "byteLength": 410, + "sha256": "eea033222d278c8aa8f4b0de36243adba4e5b6ce274d24a3dcd365ad42c19a38" + }, + { + "path": "kernels/f114c5cb9f1a7e377d9a3bcb28ab15b72807ece1796cf9c02fc74e7ce7318cc3.wgsl", + "byteLength": 341, + "sha256": "f114c5cb9f1a7e377d9a3bcb28ab15b72807ece1796cf9c02fc74e7ce7318cc3" + }, + { + "path": "kernels/f276cf6656b999339df464b61934487f5c8eb7a0957cf6d902254ab1aa2ffcd0.wgsl", + "byteLength": 561, + "sha256": "f276cf6656b999339df464b61934487f5c8eb7a0957cf6d902254ab1aa2ffcd0" + }, + { + "path": "kernels/f43cdd9c75fc256957b1b144a0587e8c8ddea861f991e1cdb6c9098e8b5d6ed7.wgsl", + "byteLength": 394, + "sha256": "f43cdd9c75fc256957b1b144a0587e8c8ddea861f991e1cdb6c9098e8b5d6ed7" + }, + { + "path": "kernels/f5b51b1e0e52bd87a70a0ca169186ebb5734cbdc2b8f375ba0241485b1df4e66.wgsl", + "byteLength": 371, + "sha256": "f5b51b1e0e52bd87a70a0ca169186ebb5734cbdc2b8f375ba0241485b1df4e66" + }, + { + "path": "kernels/f6ccc1a43bc7cbaab21e08ce5de522f383236035c835289025961d70c49c60e8.wgsl", + "byteLength": 296, + "sha256": "f6ccc1a43bc7cbaab21e08ce5de522f383236035c835289025961d70c49c60e8" + }, + { + "path": "kernels/f73912607dffcf0ee27e9e1e1d342a3bef8f998e02cd06a8352524ef35a96d17.wgsl", + "byteLength": 599, + "sha256": "f73912607dffcf0ee27e9e1e1d342a3bef8f998e02cd06a8352524ef35a96d17" + }, + { + "path": "kernels/f84306aa81033babc4c37ecfb91a411dba2ccee17202615deefb2206add106e2.wgsl", + "byteLength": 367, + "sha256": "f84306aa81033babc4c37ecfb91a411dba2ccee17202615deefb2206add106e2" + }, + { + "path": "kernels/f994142054bcb47d6a554524dd8de9c3054467a1d07011027c47f475d0909c19.wgsl", + "byteLength": 845, + "sha256": "f994142054bcb47d6a554524dd8de9c3054467a1d07011027c47f475d0909c19" + } + ], + "metadata": { + "repository": "Qwen/Qwen3.5-0.8B", + "revision": "2fc06364715b967f1860aea9cf38778875588b17", + "architecture": "qwen3_5", + "modalities": [ + "text" + ], + "license": "apache-2.0", + "sourceWeights": [ + { + "path": "model.safetensors-00001-of-00001.safetensors", + "sha256": "04b1c301231dd422b8860db31311ab2721511346a32cb1e079c4c4e5f1fe4696" + } + ], + "profile": "fp32-int2-g32-mixed-8f56efd4", + "prefill": "queued-incremental", + "checkpointCertified": false, + "fusion": true, + "quantizationSelection": { + "include": [ + "*_mlp_up_proj_weight" + ], + "exclude": [] + } + }, + "validation": { + "checkpointCertified": false, + "hardware": [] + }, + "generation": { + "contextLength": 4096, + "inputIds": "input_ids", + "positionIds": "position_ids", + "attentionMask": "attention_mask", + "logits": "linear_186", + "eosTokenIds": [ + 248044 + ], + "tokenizer": "tokenizer/tokenizer.json", + "tokenizerConfig": "tokenizer/tokenizer_config.json", + "chatTemplate": "tokenizer/chat_template.jinja" + } +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/source/LICENSE b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/source/LICENSE new file mode 100644 index 0000000000000000000000000000000000000000..f938136e3adacfd92be087f6e113b5d6d97f678f --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/source/LICENSE @@ -0,0 +1,202 @@ + + Apache License + Version 2.0, January 2004 + http://www.apache.org/licenses/ + + TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION + + 1. Definitions. + + "License" shall mean the terms and conditions for use, reproduction, + and distribution as defined by Sections 1 through 9 of this document. + + "Licensor" shall mean the copyright owner or entity authorized by + the copyright owner that is granting the License. + + "Legal Entity" shall mean the union of the acting entity and all + other entities that control, are controlled by, or are under common + control with that entity. For the purposes of this definition, + "control" means (i) the power, direct or indirect, to cause the + direction or management of such entity, whether by contract or + otherwise, or (ii) ownership of fifty percent (50%) or more of the + outstanding shares, or (iii) beneficial ownership of such entity. + + "You" (or "Your") shall mean an individual or Legal Entity + exercising permissions granted by this License. + + "Source" form shall mean the preferred form for making modifications, + including but not limited to software source code, documentation + source, and configuration files. + + "Object" form shall mean any form resulting from mechanical + transformation or translation of a Source form, including but + not limited to compiled object code, generated documentation, + and conversions to other media types. + + "Work" shall mean the work of authorship, whether in Source or + Object form, made available under the License, as indicated by a + copyright notice that is included in or attached to the work + (an example is provided in the Appendix below). + + "Derivative Works" shall mean any work, whether in Source or Object + form, that is based on (or derived from) the Work and for which the + editorial revisions, annotations, elaborations, or other modifications + represent, as a whole, an original work of authorship. For the purposes + of this License, Derivative Works shall not include works that remain + separable from, or merely link (or bind by name) to the interfaces of, + the Work and Derivative Works thereof. + + "Contribution" shall mean any work of authorship, including + the original version of the Work and any modifications or additions + to that Work or Derivative Works thereof, that is intentionally + submitted to Licensor for inclusion in the Work by the copyright owner + or by an individual or Legal Entity authorized to submit on behalf of + the copyright owner. For the purposes of this definition, "submitted" + means any form of electronic, verbal, or written communication sent + to the Licensor or its representatives, including but not limited to + communication on electronic mailing lists, source code control systems, + and issue tracking systems that are managed by, or on behalf of, the + Licensor for the purpose of discussing and improving the Work, but + excluding communication that is conspicuously marked or otherwise + designated in writing by the copyright owner as "Not a Contribution." + + "Contributor" shall mean Licensor and any individual or Legal Entity + on behalf of whom a Contribution has been received by Licensor and + subsequently incorporated within the Work. + + 2. Grant of Copyright License. Subject to the terms and conditions of + this License, each Contributor hereby grants to You a perpetual, + worldwide, non-exclusive, no-charge, royalty-free, irrevocable + copyright license to reproduce, prepare Derivative Works of, + publicly display, publicly perform, sublicense, and distribute the + Work and such Derivative Works in Source or Object form. + + 3. Grant of Patent License. Subject to the terms and conditions of + this License, each Contributor hereby grants to You a perpetual, + worldwide, non-exclusive, no-charge, royalty-free, irrevocable + (except as stated in this section) patent license to make, have made, + use, offer to sell, sell, import, and otherwise transfer the Work, + where such license applies only to those patent claims licensable + by such Contributor that are necessarily infringed by their + Contribution(s) alone or by combination of their Contribution(s) + with the Work to which such Contribution(s) was submitted. If You + institute patent litigation against any entity (including a + cross-claim or counterclaim in a lawsuit) alleging that the Work + or a Contribution incorporated within the Work constitutes direct + or contributory patent infringement, then any patent licenses + granted to You under this License for that Work shall terminate + as of the date such litigation is filed. + + 4. Redistribution. You may reproduce and distribute copies of the + Work or Derivative Works thereof in any medium, with or without + modifications, and in Source or Object form, provided that You + meet the following conditions: + + (a) You must give any other recipients of the Work or + Derivative Works a copy of this License; and + + (b) You must cause any modified files to carry prominent notices + stating that You changed the files; and + + (c) You must retain, in the Source form of any Derivative Works + that You distribute, all copyright, patent, trademark, and + attribution notices from the Source form of the Work, + excluding those notices that do not pertain to any part of + the Derivative Works; and + + (d) If the Work includes a "NOTICE" text file as part of its + distribution, then any Derivative Works that You distribute must + include a readable copy of the attribution notices contained + within such NOTICE file, excluding those notices that do not + pertain to any part of the Derivative Works, in at least one + of the following places: within a NOTICE text file distributed + as part of the Derivative Works; within the Source form or + documentation, if provided along with the Derivative Works; or, + within a display generated by the Derivative Works, if and + wherever such third-party notices normally appear. The contents + of the NOTICE file are for informational purposes only and + do not modify the License. You may add Your own attribution + notices within Derivative Works that You distribute, alongside + or as an addendum to the NOTICE text from the Work, provided + that such additional attribution notices cannot be construed + as modifying the License. + + You may add Your own copyright statement to Your modifications and + may provide additional or different license terms and conditions + for use, reproduction, or distribution of Your modifications, or + for any such Derivative Works as a whole, provided Your use, + reproduction, and distribution of the Work otherwise complies with + the conditions stated in this License. + + 5. Submission of Contributions. Unless You explicitly state otherwise, + any Contribution intentionally submitted for inclusion in the Work + by You to the Licensor shall be under the terms and conditions of + this License, without any additional terms or conditions. + Notwithstanding the above, nothing herein shall supersede or modify + the terms of any separate license agreement you may have executed + with Licensor regarding such Contributions. + + 6. Trademarks. This License does not grant permission to use the trade + names, trademarks, service marks, or product names of the Licensor, + except as required for reasonable and customary use in describing the + origin of the Work and reproducing the content of the NOTICE file. + + 7. Disclaimer of Warranty. Unless required by applicable law or + agreed to in writing, Licensor provides the Work (and each + Contributor provides its Contributions) on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or + implied, including, without limitation, any warranties or conditions + of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A + PARTICULAR PURPOSE. You are solely responsible for determining the + appropriateness of using or redistributing the Work and assume any + risks associated with Your exercise of permissions under this License. + + 8. Limitation of Liability. In no event and under no legal theory, + whether in tort (including negligence), contract, or otherwise, + unless required by applicable law (such as deliberate and grossly + negligent acts) or agreed to in writing, shall any Contributor be + liable to You for damages, including any direct, indirect, special, + incidental, or consequential damages of any character arising as a + result of this License or out of the use or inability to use the + Work (including but not limited to damages for loss of goodwill, + work stoppage, computer failure or malfunction, or any and all + other commercial damages or losses), even if such Contributor + has been advised of the possibility of such damages. + + 9. Accepting Warranty or Additional Liability. While redistributing + the Work or Derivative Works thereof, You may choose to offer, + and charge a fee for, acceptance of support, warranty, indemnity, + or other liability obligations and/or rights consistent with this + License. However, in accepting such obligations, You may act only + on Your own behalf and on Your sole responsibility, not on behalf + of any other Contributor, and only if You agree to indemnify, + defend, and hold each Contributor harmless for any liability + incurred by, or claims asserted against, such Contributor by reason + of your accepting any such warranty or additional liability. + + END OF TERMS AND CONDITIONS + + APPENDIX: How to apply the Apache License to your work. + + To apply the Apache License to your work, attach the following + boilerplate notice, with the fields enclosed by brackets "[]" + replaced with your own identifying information. (Don't include + the brackets!) The text should be enclosed in the appropriate + comment syntax for the file format. We also recommend that a + file or class name and description of purpose be included on the + same "printed page" as the copyright notice for easier + identification within third-party archives. + + Copyright 2026 Alibaba Cloud + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. \ No newline at end of file diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/source/config.json b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/source/config.json new file mode 100644 index 0000000000000000000000000000000000000000..715f0448b9d38103211f0ad88bbb4d6e4f4be8c9 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/source/config.json @@ -0,0 +1,96 @@ +{ + "architectures": [ + "Qwen3_5ForConditionalGeneration" + ], + "image_token_id": 248056, + "model_type": "qwen3_5", + "text_config": { + "attention_bias": false, + "attention_dropout": 0.0, + "attn_output_gate": true, + "dtype": "bfloat16", + "eos_token_id": 248044, + "full_attention_interval": 4, + "head_dim": 256, + "hidden_act": "silu", + "hidden_size": 1024, + "initializer_range": 0.02, + "intermediate_size": 3584, + "layer_types": [ + "linear_attention", + "linear_attention", + "linear_attention", + "full_attention", + "linear_attention", + "linear_attention", + "linear_attention", + "full_attention", + "linear_attention", + "linear_attention", + "linear_attention", + "full_attention", + "linear_attention", + "linear_attention", + "linear_attention", + "full_attention", + "linear_attention", + "linear_attention", + "linear_attention", + "full_attention", + "linear_attention", + "linear_attention", + "linear_attention", + "full_attention" + ], + "linear_conv_kernel_dim": 4, + "linear_key_head_dim": 128, + "linear_num_key_heads": 16, + "linear_num_value_heads": 16, + "linear_value_head_dim": 128, + "max_position_embeddings": 262144, + "mlp_only_layers": [], + "model_type": "qwen3_5_text", + "mtp_num_hidden_layers": 1, + "mtp_use_dedicated_embeddings": false, + "num_attention_heads": 8, + "num_hidden_layers": 24, + "num_key_value_heads": 2, + "rms_norm_eps": 1e-06, + "tie_word_embeddings": true, + "use_cache": true, + "vocab_size": 248320, + "mamba_ssm_dtype": "float32", + "rope_parameters": { + "mrope_interleaved": true, + "mrope_section": [ + 11, + 11, + 10 + ], + "rope_type": "default", + "rope_theta": 10000000, + "partial_rotary_factor": 0.25 + } + }, + "tie_word_embeddings": true, + "transformers_version": "4.57.0.dev0", + "video_token_id": 248057, + "vision_config": { + "deepstack_visual_indexes": [], + "depth": 12, + "hidden_act": "gelu_pytorch_tanh", + "hidden_size": 768, + "in_channels": 3, + "initializer_range": 0.02, + "intermediate_size": 3072, + "model_type": "qwen3_5", + "num_heads": 12, + "num_position_embeddings": 2304, + "out_hidden_size": 1024, + "patch_size": 16, + "spatial_merge_size": 2, + "temporal_patch_size": 2 + }, + "vision_end_token_id": 248054, + "vision_start_token_id": 248053 +} \ No newline at end of file diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/source/provenance.json b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/source/provenance.json new file mode 100644 index 0000000000000000000000000000000000000000..6b5a8add65c7ba0efb165c5fe3f96c6ceaf7628d --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/source/provenance.json @@ -0,0 +1,10 @@ +{ + "repository": "Qwen/Qwen3.5-0.8B", + "revision": "2fc06364715b967f1860aea9cf38778875588b17", + "architecture": "qwen3_5", + "modalities": [ + "text", + "image" + ], + "license": "apache-2.0" +} diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/tokenizer/chat_template.jinja b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/tokenizer/chat_template.jinja new file mode 100644 index 0000000000000000000000000000000000000000..0ef09f214eaa6d9bca297988afc1454b5827b2c7 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/tokenizer/chat_template.jinja @@ -0,0 +1,154 @@ +{%- set image_count = namespace(value=0) %} +{%- set video_count = namespace(value=0) %} +{%- macro render_content(content, do_vision_count, is_system_content=false) %} + {%- if content is string %} + {{- content }} + {%- elif content is iterable and content is not mapping %} + {%- for item in content %} + {%- if 'image' in item or 'image_url' in item or item.type == 'image' %} + {%- if is_system_content %} + {{- raise_exception('System message cannot contain images.') }} + {%- endif %} + {%- if do_vision_count %} + {%- set image_count.value = image_count.value + 1 %} + {%- endif %} + {%- if add_vision_id %} + {{- 'Picture ' ~ image_count.value ~ ': ' }} + {%- endif %} + {{- '<|vision_start|><|image_pad|><|vision_end|>' }} + {%- elif 'video' in item or item.type == 'video' %} + {%- if is_system_content %} + {{- raise_exception('System message cannot contain videos.') }} + {%- endif %} + {%- if do_vision_count %} + {%- set video_count.value = video_count.value + 1 %} + {%- endif %} + {%- if add_vision_id %} + {{- 'Video ' ~ video_count.value ~ ': ' }} + {%- endif %} + {{- '<|vision_start|><|video_pad|><|vision_end|>' }} + {%- elif 'text' in item %} + {{- item.text }} + {%- else %} + {{- raise_exception('Unexpected item type in content.') }} + {%- endif %} + {%- endfor %} + {%- elif content is none or content is undefined %} + {{- '' }} + {%- else %} + {{- raise_exception('Unexpected content type.') }} + {%- endif %} +{%- endmacro %} +{%- if not messages %} + {{- raise_exception('No messages provided.') }} +{%- endif %} +{%- if tools and tools is iterable and tools is not mapping %} + {{- '<|im_start|>system\n' }} + {{- "# Tools\n\nYou have access to the following functions:\n\n" }} + {%- for tool in tools %} + {{- "\n" }} + {{- tool | tojson }} + {%- endfor %} + {{- "\n" }} + {{- '\n\nIf you choose to call a function ONLY reply in the following format with NO suffix:\n\n\n\n\nvalue_1\n\n\nThis is the value for the second parameter\nthat can span\nmultiple lines\n\n\n\n\n\nReminder:\n- Function calls MUST follow the specified format: an inner block must be nested within XML tags\n- Required parameters MUST be specified\n- You may provide optional reasoning for your function call in natural language BEFORE the function call, but NOT after\n- If there is no function call available, answer the question like normal with your current knowledge and do not tell the user about function calls\n' }} + {%- if messages[0].role == 'system' %} + {%- set content = render_content(messages[0].content, false, true)|trim %} + {%- if content %} + {{- '\n\n' + content }} + {%- endif %} + {%- endif %} + {{- '<|im_end|>\n' }} +{%- else %} + {%- if messages[0].role == 'system' %} + {%- set content = render_content(messages[0].content, false, true)|trim %} + {{- '<|im_start|>system\n' + content + '<|im_end|>\n' }} + {%- endif %} +{%- endif %} +{%- set ns = namespace(multi_step_tool=true, last_query_index=messages|length - 1) %} +{%- for message in messages[::-1] %} + {%- set index = (messages|length - 1) - loop.index0 %} + {%- if ns.multi_step_tool and message.role == "user" %} + {%- set content = render_content(message.content, false)|trim %} + {%- if not(content.startswith('') and content.endswith('')) %} + {%- set ns.multi_step_tool = false %} + {%- set ns.last_query_index = index %} + {%- endif %} + {%- endif %} +{%- endfor %} +{%- if ns.multi_step_tool %} + {{- raise_exception('No user query found in messages.') }} +{%- endif %} +{%- for message in messages %} + {%- set content = render_content(message.content, true)|trim %} + {%- if message.role == "system" %} + {%- if not loop.first %} + {{- raise_exception('System message must be at the beginning.') }} + {%- endif %} + {%- elif message.role == "user" %} + {{- '<|im_start|>' + message.role + '\n' + content + '<|im_end|>' + '\n' }} + {%- elif message.role == "assistant" %} + {%- set reasoning_content = '' %} + {%- if message.reasoning_content is string %} + {%- set reasoning_content = message.reasoning_content %} + {%- else %} + {%- if '' in content %} + {%- set reasoning_content = content.split('')[0].rstrip('\n').split('')[-1].lstrip('\n') %} + {%- set content = content.split('')[-1].lstrip('\n') %} + {%- endif %} + {%- endif %} + {%- set reasoning_content = reasoning_content|trim %} + {%- if loop.index0 > ns.last_query_index %} + {{- '<|im_start|>' + message.role + '\n\n' + reasoning_content + '\n\n\n' + content }} + {%- else %} + {{- '<|im_start|>' + message.role + '\n' + content }} + {%- endif %} + {%- if message.tool_calls and message.tool_calls is iterable and message.tool_calls is not mapping %} + {%- for tool_call in message.tool_calls %} + {%- if tool_call.function is defined %} + {%- set tool_call = tool_call.function %} + {%- endif %} + {%- if loop.first %} + {%- if content|trim %} + {{- '\n\n\n\n' }} + {%- else %} + {{- '\n\n' }} + {%- endif %} + {%- else %} + {{- '\n\n\n' }} + {%- endif %} + {%- if tool_call.arguments is defined %} + {%- for args_name, args_value in tool_call.arguments|items %} + {{- '\n' }} + {%- set args_value = args_value | tojson | safe if args_value is mapping or (args_value is sequence and args_value is not string) else args_value | string %} + {{- args_value }} + {{- '\n\n' }} + {%- endfor %} + {%- endif %} + {{- '\n' }} + {%- endfor %} + {%- endif %} + {{- '<|im_end|>\n' }} + {%- elif message.role == "tool" %} + {%- if loop.previtem and loop.previtem.role != "tool" %} + {{- '<|im_start|>user' }} + {%- endif %} + {{- '\n\n' }} + {{- content }} + {{- '\n' }} + {%- if not loop.last and loop.nextitem.role != "tool" %} + {{- '<|im_end|>\n' }} + {%- elif loop.last %} + {{- '<|im_end|>\n' }} + {%- endif %} + {%- else %} + {{- raise_exception('Unexpected message role.') }} + {%- endif %} +{%- endfor %} +{%- if add_generation_prompt %} + {{- '<|im_start|>assistant\n' }} + {%- if enable_thinking is defined and enable_thinking is true %} + {{- '\n' }} + {%- else %} + {{- '\n\n\n\n' }} + {%- endif %} +{%- endif %} \ No newline at end of file diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/tokenizer/tokenizer.json b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/tokenizer/tokenizer.json new file mode 100644 index 0000000000000000000000000000000000000000..a73a846725794819aa6e1c9e97d8dc9671c2006d --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/tokenizer/tokenizer.json @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:5f9e4d4901a92b997e463c1f46055088b6cca5ca61a6522d1b9f64c4bb81cb42 +size 12807982 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/tokenizer/tokenizer_config.json b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/tokenizer/tokenizer_config.json new file mode 100644 index 0000000000000000000000000000000000000000..fae3ce993e07c092ad024dde45e592379fde91bb --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/tokenizer/tokenizer_config.json @@ -0,0 +1,305 @@ +{ + "add_prefix_space": false, + "added_tokens_decoder": { + "248044": { + "content": "<|endoftext|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248045": { + "content": "<|im_start|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248046": { + "content": "<|im_end|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248047": { + "content": "<|object_ref_start|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248048": { + "content": "<|object_ref_end|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248049": { + "content": "<|box_start|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248050": { + "content": "<|box_end|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248051": { + "content": "<|quad_start|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248052": { + "content": "<|quad_end|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248053": { + "content": "<|vision_start|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248054": { + "content": "<|vision_end|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248055": { + "content": "<|vision_pad|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248056": { + "content": "<|image_pad|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248057": { + "content": "<|video_pad|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248058": { + "content": "", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248059": { + "content": "", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248060": { + "content": "<|fim_prefix|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248061": { + "content": "<|fim_middle|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248062": { + "content": "<|fim_suffix|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248063": { + "content": "<|fim_pad|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248064": { + "content": "<|repo_name|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248065": { + "content": "<|file_sep|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248066": { + "content": "", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248067": { + "content": "", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248068": { + "content": "", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248069": { + "content": "", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": false + }, + "248070": { + "content": "<|audio_start|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248071": { + "content": "<|audio_end|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248072": { + "content": "", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248073": { + "content": "", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248074": { + "content": "", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248075": { + "content": "", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + }, + "248076": { + "content": "<|audio_pad|>", + "lstrip": false, + "normalized": false, + "rstrip": false, + "single_word": false, + "special": true + } + }, + "additional_special_tokens": [ + "<|im_start|>", + "<|im_end|>", + "<|object_ref_start|>", + "<|object_ref_end|>", + "<|box_start|>", + "<|box_end|>", + "<|quad_start|>", + "<|quad_end|>", + "<|vision_start|>", + "<|vision_end|>", + "<|vision_pad|>", + "<|image_pad|>", + "<|video_pad|>" + ], + "bos_token": null, + "chat_template": "{%- set image_count = namespace(value=0) %}\n{%- set video_count = namespace(value=0) %}\n{%- macro render_content(content, do_vision_count, is_system_content=false) %}\n {%- if content is string %}\n {{- content }}\n {%- elif content is iterable and content is not mapping %}\n {%- for item in content %}\n {%- if 'image' in item or 'image_url' in item or item.type == 'image' %}\n {%- if is_system_content %}\n {{- raise_exception('System message cannot contain images.') }}\n {%- endif %}\n {%- if do_vision_count %}\n {%- set image_count.value = image_count.value + 1 %}\n {%- endif %}\n {%- if add_vision_id %}\n {{- 'Picture ' ~ image_count.value ~ ': ' }}\n {%- endif %}\n {{- '<|vision_start|><|image_pad|><|vision_end|>' }}\n {%- elif 'video' in item or item.type == 'video' %}\n {%- if is_system_content %}\n {{- raise_exception('System message cannot contain videos.') }}\n {%- endif %}\n {%- if do_vision_count %}\n {%- set video_count.value = video_count.value + 1 %}\n {%- endif %}\n {%- if add_vision_id %}\n {{- 'Video ' ~ video_count.value ~ ': ' }}\n {%- endif %}\n {{- '<|vision_start|><|video_pad|><|vision_end|>' }}\n {%- elif 'text' in item %}\n {{- item.text }}\n {%- else %}\n {{- raise_exception('Unexpected item type in content.') }}\n {%- endif %}\n {%- endfor %}\n {%- elif content is none or content is undefined %}\n {{- '' }}\n {%- else %}\n {{- raise_exception('Unexpected content type.') }}\n {%- endif %}\n{%- endmacro %}\n{%- if not messages %}\n {{- raise_exception('No messages provided.') }}\n{%- endif %}\n{%- if tools and tools is iterable and tools is not mapping %}\n {{- '<|im_start|>system\\n' }}\n {{- \"# Tools\\n\\nYou have access to the following functions:\\n\\n\" }}\n {%- for tool in tools %}\n {{- \"\\n\" }}\n {{- tool | tojson }}\n {%- endfor %}\n {{- \"\\n\" }}\n {{- '\\n\\nIf you choose to call a function ONLY reply in the following format with NO suffix:\\n\\n\\n\\n\\nvalue_1\\n\\n\\nThis is the value for the second parameter\\nthat can span\\nmultiple lines\\n\\n\\n\\n\\n\\nReminder:\\n- Function calls MUST follow the specified format: an inner block must be nested within XML tags\\n- Required parameters MUST be specified\\n- You may provide optional reasoning for your function call in natural language BEFORE the function call, but NOT after\\n- If there is no function call available, answer the question like normal with your current knowledge and do not tell the user about function calls\\n' }}\n {%- if messages[0].role == 'system' %}\n {%- set content = render_content(messages[0].content, false, true)|trim %}\n {%- if content %}\n {{- '\\n\\n' + content }}\n {%- endif %}\n {%- endif %}\n {{- '<|im_end|>\\n' }}\n{%- else %}\n {%- if messages[0].role == 'system' %}\n {%- set content = render_content(messages[0].content, false, true)|trim %}\n {{- '<|im_start|>system\\n' + content + '<|im_end|>\\n' }}\n {%- endif %}\n{%- endif %}\n{%- set ns = namespace(multi_step_tool=true, last_query_index=messages|length - 1) %}\n{%- for message in messages[::-1] %}\n {%- set index = (messages|length - 1) - loop.index0 %}\n {%- if ns.multi_step_tool and message.role == \"user\" %}\n {%- set content = render_content(message.content, false)|trim %}\n {%- if not(content.startswith('') and content.endswith('')) %}\n {%- set ns.multi_step_tool = false %}\n {%- set ns.last_query_index = index %}\n {%- endif %}\n {%- endif %}\n{%- endfor %}\n{%- if ns.multi_step_tool %}\n {{- raise_exception('No user query found in messages.') }}\n{%- endif %}\n{%- for message in messages %}\n {%- set content = render_content(message.content, true)|trim %}\n {%- if message.role == \"system\" %}\n {%- if not loop.first %}\n {{- raise_exception('System message must be at the beginning.') }}\n {%- endif %}\n {%- elif message.role == \"user\" %}\n {{- '<|im_start|>' + message.role + '\\n' + content + '<|im_end|>' + '\\n' }}\n {%- elif message.role == \"assistant\" %}\n {%- set reasoning_content = '' %}\n {%- if message.reasoning_content is string %}\n {%- set reasoning_content = message.reasoning_content %}\n {%- else %}\n {%- if '' in content %}\n {%- set reasoning_content = content.split('')[0].rstrip('\\n').split('')[-1].lstrip('\\n') %}\n {%- set content = content.split('')[-1].lstrip('\\n') %}\n {%- endif %}\n {%- endif %}\n {%- set reasoning_content = reasoning_content|trim %}\n {%- if loop.index0 > ns.last_query_index %}\n {{- '<|im_start|>' + message.role + '\\n\\n' + reasoning_content + '\\n\\n\\n' + content }}\n {%- else %}\n {{- '<|im_start|>' + message.role + '\\n' + content }}\n {%- endif %}\n {%- if message.tool_calls and message.tool_calls is iterable and message.tool_calls is not mapping %}\n {%- for tool_call in message.tool_calls %}\n {%- if tool_call.function is defined %}\n {%- set tool_call = tool_call.function %}\n {%- endif %}\n {%- if loop.first %}\n {%- if content|trim %}\n {{- '\\n\\n\\n\\n' }}\n {%- else %}\n {{- '\\n\\n' }}\n {%- endif %}\n {%- else %}\n {{- '\\n\\n\\n' }}\n {%- endif %}\n {%- if tool_call.arguments is defined %}\n {%- for args_name, args_value in tool_call.arguments|items %}\n {{- '\\n' }}\n {%- set args_value = args_value | tojson | safe if args_value is mapping or (args_value is sequence and args_value is not string) else args_value | string %}\n {{- args_value }}\n {{- '\\n\\n' }}\n {%- endfor %}\n {%- endif %}\n {{- '\\n' }}\n {%- endfor %}\n {%- endif %}\n {{- '<|im_end|>\\n' }}\n {%- elif message.role == \"tool\" %}\n {%- if loop.previtem and loop.previtem.role != \"tool\" %}\n {{- '<|im_start|>user' }}\n {%- endif %}\n {{- '\\n\\n' }}\n {{- content }}\n {{- '\\n' }}\n {%- if not loop.last and loop.nextitem.role != \"tool\" %}\n {{- '<|im_end|>\\n' }}\n {%- elif loop.last %}\n {{- '<|im_end|>\\n' }}\n {%- endif %}\n {%- else %}\n {{- raise_exception('Unexpected message role.') }}\n {%- endif %}\n{%- endfor %}\n{%- if add_generation_prompt %}\n {{- '<|im_start|>assistant\\n' }}\n {%- if enable_thinking is defined and enable_thinking is true %}\n {{- '\\n' }}\n {%- else %}\n {{- '\\n\\n\\n\\n' }}\n {%- endif %}\n{%- endif %}", + "clean_up_tokenization_spaces": false, + "eos_token": "<|im_end|>", + "errors": "replace", + "model_max_length": 262144, + "pad_token": "<|endoftext|>", + "split_special_tokens": false, + "tokenizer_class": "Qwen2Tokenizer", + "unk_token": null, + "add_bos_token": false, + "pretokenize_regex": "(?i:'s|'t|'re|'ve|'m|'ll|'d)|[^\\r\\n\\p{L}\\p{N}]?[\\p{L}\\p{M}]+|\\p{N}| ?[^\\s\\p{L}\\p{M}\\p{N}]+[\\r\\n]*|\\s*[\\r\\n]+|\\s+(?!\\S)|\\s+", + "extra_special_tokens": { + "audio_bos_token": "<|audio_start|>", + "audio_eos_token": "<|audio_end|>", + "audio_token": "<|audio_pad|>", + "image_token": "<|image_pad|>", + "video_token": "<|video_pad|>", + "vision_bos_token": "<|vision_start|>", + "vision_eos_token": "<|vision_end|>" + } +} \ No newline at end of file diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00000.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00000.bin new file mode 100644 index 0000000000000000000000000000000000000000..5db2ecb6f26a1c25d7bb3a3cd3583301fe369a16 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00000.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:998effe0875f489e58eb9592d73a2a2ef51b627520bd119524651c8ffc886282 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00001.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00001.bin new file mode 100644 index 0000000000000000000000000000000000000000..c66d6bf828e3954433fc9a1cbc977897111dcce8 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00001.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:7b669e08c1c1103797bd533fad1d52fe36b3f1293db9d09065ec24c7c5f2db8d +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00002.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00002.bin new file mode 100644 index 0000000000000000000000000000000000000000..31d767aba2f1556e144aca0f3836482a736e8d3b --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00002.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:a281c127a0a9bf7f5498534c1a350d922f251e19343c303e70663486d609260e +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00003.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00003.bin new file mode 100644 index 0000000000000000000000000000000000000000..852884880097dc1e1012ba199d1c8f1ae1bd7514 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00003.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:1b4b102afabd0edfc06a42c1c14f6cc459e9e1e2cf0b8193fc5aa3046c04cc76 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00004.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00004.bin new file mode 100644 index 0000000000000000000000000000000000000000..7dd39ed1515eb63706e647cef15a3ce06cffe4fa --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00004.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:6480b9ca38668264d75fbc1ecdd78870fe1a8b95835f485c5543d11e2abbbe1f +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00005.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00005.bin new file mode 100644 index 0000000000000000000000000000000000000000..97f5a7d9e18d5485309c2582896c4e57de2cf0b7 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00005.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:c8f2a7532918fb3acc76952b3e685030c001abdae0b74ab41fc84c64a4ab8c2e +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00006.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00006.bin new file mode 100644 index 0000000000000000000000000000000000000000..d3c1d66bcc12f63cb4a37a3e212a3f508f46a415 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00006.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:54fa18728d34cb3f67db0c70064574e8f34b786d9bf5d92e61241472e1249aa4 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00007.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00007.bin new file mode 100644 index 0000000000000000000000000000000000000000..40f67786390d40cf56257daf59a152249410eda6 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00007.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:89edafc7a5ed368e67498221df821521c36df98216fcf4e55efbd3d8e2925725 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00008.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00008.bin new file mode 100644 index 0000000000000000000000000000000000000000..9b067723851bde97e3a855788c6be4c7057538fc --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00008.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:273556893c277e1d8273045bcdc0943a4fb93d995655a870309dd8ea7f6b56d0 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00009.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00009.bin new file mode 100644 index 0000000000000000000000000000000000000000..73ed735ff641345da7f6f18a59063b36c5d48895 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00009.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:9207ae6e82ef7ce44d4057eb208a0c2f35ab06d2dd571aee8b16191116133536 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00010.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00010.bin new file mode 100644 index 0000000000000000000000000000000000000000..c34ea852439b35e7ff3ee6cf3f3c427c0f554f4b --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00010.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:7a64f841b6a5d951d488428e0e6eb5062717935cbaf8194931797fd8fb3eec53 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00011.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00011.bin new file mode 100644 index 0000000000000000000000000000000000000000..b9d480dc30eff157a1e60193de555f5410b64133 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00011.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:718459e02d1bd28ee721759d716fc6e124523709e4bb3f7b60b36b7d206c63ec +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00012.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00012.bin new file mode 100644 index 0000000000000000000000000000000000000000..56271e898f70e84aba086d9379efc2caeac1a320 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00012.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:2ee07dac007a616667f2c3585641332e7df05b34b6ef0e6f0904eb7a029009ad +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00013.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00013.bin new file mode 100644 index 0000000000000000000000000000000000000000..177c023a1c6526c3ea693b60b3eb8e52c5fa1a2d --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00013.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:e054e1b2605cd22d63e477299218d5426b042968acd0d53942d01ec5a29db3cc +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00014.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00014.bin new file mode 100644 index 0000000000000000000000000000000000000000..e428f2e7f1c2397854b11573b346d4b35d91832d --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00014.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:570d6d4ff3455dcb089501e0a30efb71a00e8ed835478e6fa36348d4f3a6dc26 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00015.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00015.bin new file mode 100644 index 0000000000000000000000000000000000000000..192e31b84429eb780d04ad65e198df93618f12cb --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00015.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:de850adc45c0f4b78b99be441605d5c843e1f2a167b017b030cbc897654ebee9 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00016.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00016.bin new file mode 100644 index 0000000000000000000000000000000000000000..6c96bfa77e2ca854b0e1ba8a9d3ea72ed5bc750d --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00016.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:01dfe7fb384c6016183de73de8cfb43ef410b697391103f3fd7c269983706651 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00017.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00017.bin new file mode 100644 index 0000000000000000000000000000000000000000..3ac4e2a2a2aba39946cb47db4d5519cab3434784 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00017.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:aedc4a87a0492afaab3aeabe395ae647f9d260b3324f5c56fb6079daae9283e3 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00018.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00018.bin new file mode 100644 index 0000000000000000000000000000000000000000..1bfddebe0478ac6a926afae244adf3731373daa6 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00018.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:a2b97aff1515a8441f067b40b71eb10ffeb1e60ff0e94579cd1d2b9b10514683 +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00019.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00019.bin new file mode 100644 index 0000000000000000000000000000000000000000..dfbd9f229f319ac784c2d4bce36f6d8072e01919 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00019.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:58b262bf15ce55ec421276df579123f6cfcb4025b27dd86c26f0a02a18ad766b +size 67108864 diff --git a/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00020.bin b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00020.bin new file mode 100644 index 0000000000000000000000000000000000000000..d9bf9975a1d7e4dea7b64ea1b250eb5cacd53580 --- /dev/null +++ b/qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/weights/00020.bin @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:d99fbd379a8c7234b904b0375c9da6b2d212e75466d0bf6eaee7d65cba91ff12 +size 20475264