Add qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2 WebTorch bundle

.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

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/04f26661cc9ac73bb7f39f66b76c4be31ee8774e094eeb41ed7d677e1b1d3fa9.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0554a443ecf6ead32fd2d58178b9ce3c592c89a09b3f63fa509ee947323bb497.wgsl

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

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<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 16u;
+    if (i >= 16u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 16u) * 1024u + p]) * 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); }
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0835fa80eaadb7c0e6d677e1fc622ad5dcedabb431e5d312f3b7aeb2252a7c7c.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0e5468583a1ba929483119c50d4be42a58aed4dde1e15cc231908ed7543d0b71.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0e9b6ee5f92769faaee5c746d1a3b936cdb7731f3f8e70cbf0a3fe96f6c8d921.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0f466ed2664d77386f48afd19d8ab2734312bfd4f5f4d5b063d801a0b0a97cd9.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/0fd78c192784ba4e545091262cf1bac47ef69ca387b850f97b9de39dd0a2a741.wgsl

@@ -0,0 +1,25 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<u32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 16u;
+    if (i >= 16u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 16u) * 1024u + p]) * 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); }
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/10feb208ed0549389f54419ff8985bee4778276913d4c6a613a8b120a809ab5e.wgsl

@@ -0,0 +1,25 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<u32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 2048u;
+    if (i >= 2048u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 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); }
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/12155d86d776e73f23d3bff0268c60dccb54c209ee9ee3a8cfd227ec4aa5b6df.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/1255847146b31732939e573cdb26c483358c1ce07f987bd8c3d3d69ae8c7e141.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/13237b557466d2fef5bfbc74e7ac77219b9821a53bc0843dbb5d14db8b04d414.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/13dc73935b9b3b978232043e13babee3877bd53ca572be440de0590dacce8b17.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/15b4ee3a9b5df59d5213e59951d8874f4eb4b9132e04fc2febddad84b44277f5.wgsl

@@ -0,0 +1,25 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<u32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 512u;
+    if (i >= 512u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 512u) * 1024u + p]) * 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); }
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/17ec6bdaf98edf7657b74a25036c2bec806ecbb4280f33fff6afd2e0de5ed377.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/1ad205080cae29142a3ad87343a0cf3ac4945455398293946110d421ce6d515e.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/20fb9e0d0bd2aab34b7f69eac259898e68cac4a0d6b0bda5e95c9126a660336d.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/222a2645e3e02d0a4063a63ca1b6fa4c89faee6b76905b20de4dd54ab3077ea3.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/234987ec68497956ea7aef8e3f44c8a73dbe4c1519165d123946bedbb4d18e26.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/26cf6eb3b486d1652247ff130b05ed0e8c8560028eeba59acfc075b976bfdfce.wgsl

@@ -0,0 +1,12 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 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]); }
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2711a4a396071f4e99303d01d4854ccaa37a445f4b718ebd7b61cf1c37e56b48.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2bdeb3590b1ba05cf6e9c7e8245f8aa552a12dae84f9578c603c9cee081a82d4.wgsl

@@ -0,0 +1,24 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+var<workgroup> factor: f32;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    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));
+    }
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2c0cca2a7b0198c317e68992e64f3fe9bba80b67204c2695c20c372471c150d3.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2dd29b61b039fcdd7fbc3c94f629b24b6f81aa5b9923098837597359b295e40d.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/32a5027cac5dd969d63fe5d5c88d7c87abb600c43ad1d70e1c56615ea781eff3.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/38bd527686bcade9ed19a52788ccd711ca668e76854cc6274c480860fce88ff2.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/3bb850bda95bf77706a2839616a401b773f9de0428d1b427e61fb04e2c35fe74.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/3fa1efae6d90015ba69cb40cef009f49ad7f7a960018fb92a962707deff4f09e.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/413a0717b830335decd89485fab5fbea4efeb71338961ba7a2d1221eaa45459a.wgsl

@@ -0,0 +1,25 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<u32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 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); }
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/42c5af4838daf75fff9d4bb9a5df409511c1d6daed617d05eaf033a977503eac.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4540d87eaec90c9eb42badfe0abfa70252f26b6e7de4dd31ee6903fe50e96bd1.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4ac112c035a1a4a2f5deae66fac0969000c4666730c2d7b15cb15a157ea0e9b8.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4b1a7529fd5c57d5d4593d4af58a9a3b545c7d106d18251f6452f8a562b97952.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4be963f9a4bb928764fd0efe02769e220b0e69defb2f8b22ab28e41a90d3fc14.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4d3e5f3a99a6c8c9dba5de30be4302f1e442f91d097437e825df97017e389a42.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4f48f08a8aeb51cc47a10db076a46522807d3b5d35952c9718ce1c1cd760681f.wgsl

@@ -0,0 +1,11 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 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]));
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/4fad29e3e337840e25132f10b3e3d55225f727ae13525ff0b0bdf31e03f09d74.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/51fd4d798491c6a80daff4c18b0495926b7f62269dfc3a78a327faa52ba2a759.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/53d6eb0b5e6d49bfd3bb9fd1b781e13317adf528e047da2416e95f5a3a13b1b4.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/552c2f036a125363ca85fd7ea5b1a84465b562dc62ddfffdbaaebf6e3288d321.wgsl

@@ -0,0 +1,17 @@
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<u32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 1024u;
+if (i >= 1024u) { return; }
+let token = i32(b0[i / 1024u]);
+if (token < 131072 || token >= 196608) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 131072) * 1024u + i % 1024u));
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/55b9c48ba429daf80bbae0b4737f895b475abdc9ddfdc88ab2f894f133ec2bf9.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/579685c7f8ca260bd53d1677c43d5fc9a0133c81c7442ffb31daf0c3759267c4.wgsl

@@ -0,0 +1,25 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<u32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 3584u;
+    if (i >= 3584u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 1024u; p += 64u) { acc += f32(b0[(i / 3584u) * 1024u + p]) * 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); }
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/57c5796344c26b430493a3a63dc734d4206380ba89484294f50f67d615eefcf8.wgsl

@@ -0,0 +1,17 @@
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<u32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 1024u;
+if (i >= 1024u) { return; }
+let token = i32(b0[i / 1024u]);
+if (token < 196608 || token >= 248320) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 196608) * 1024u + i % 1024u));
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/57dc406322bf3badd8490493ef924bd7e7e0afd542d2d6e409b4d8a3b4b784f7.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/59ea88d691b919af03bbb26f0253decfbc35d74f9661cdd3a08570ae85c313d2.wgsl

@@ -0,0 +1,11 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 64u;
+if (i >= 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);
+}

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5aa3725287640d289322ab469bd924bd7797f1cef28920ed4526783877c3cf6b.wgsl

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

qwen35-08b-fp32-int2-g32-up-proj-mse-home-token-major-v2/kernels/5adf40822be73f2c85ed1c73c0bd106707e33fcb0e8d1e82a643474b8bfeceeb.wgsl

@@ -0,0 +1,25 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<u32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+fn unpack_bf16_1(index: u32) -> f32 {
+    let pair = b1[index / 2u];
+    let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
+    return bitcast<f32>(bits << 16u);
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 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); }
+}