Publish WebTorch model catalog bundles (part 10)

qwen3vl-instruct-fp32-int8-g32-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));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/5bf483bfe0c68c1f7a9db4dac4ab0b9bd17c5f89eac369b179863d429e5eb6bc.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 * 4194240u;
+if (i >= 4194304u) { return; }
+out[i] = f32(f32(b0[i]) + (f32(b1[((i / 4096u) % 64u) * 4096u + ((i / 1u) % 4096u) * 1u]) * 1.0));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/5cd164e28c23e9b597089518e2f3abe538b18e0411ebf363ad2238b9fc99c002.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 * 131072u;
+if (i >= 131072u) { return; }
+let token = i32(b0[i / 2048u]);
+if (token < 131072 || token >= 151936) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 131072) * 2048u + i % 2048u));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/5de87cb506e20f87b8e8f6905d8631ed3c3f4e3697a780d9c084b4435e7077a9.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 * 768u;
+if (i >= 768u) { return; }
+out[i] = f32(b0[((i / 256u) % 3u) * 256u + ((i / 256u) % 1u) * 256u + ((i / 64u) % 4u) * 1u + ((i / 1u) % 64u) * 4u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/5f4403d78557e88e66d71eb5e53746cc07449a7c421b57eda4f093a09a68a99a.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 * 4096u;
+if (i >= 4096u) { return; }
+out[i] = f32(b0[i]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/605f51996eef0769d3b456a83f7540262d1761589b55d1fd41fb2b4dd8fa249f.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 * 8192u;
+if (i >= 8192u) { return; }
+out[i] = f32(f32(b0[i]) + (f32(b1[i]) * 1.0));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/6075dfc092685290714c66500c8308698503814bcc9d40bc8d4328b1cb820a2e.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 * 256u;
+if (i >= 256u) { return; }
+let coord = (i / 1u) % 64u;
+if (coord >= 1u && coord < 61u && (coord - 1u) % 3u == 0u) { out[i] = f32(b0[((i / 256u) % 1u) * 80u + ((i / 64u) % 4u) * 20u + ((((i / 1u) % 64u) - 1u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/61238ff310ce900b6d6a9bec5bc1d707953d54845d8992948d63a445232d57f8.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(f32(b0[i]) * 1.0);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/62708aacbfcd0063a897e3348ae69a00d4dd155fce6665b18c5176224b039e12.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 * 131072u;
+if (i >= 131072u) { return; }
+let token = i32(b0[i / 2048u]);
+if (token < 32768 || token >= 65536) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 32768) * 2048u + i % 2048u));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/62e3f4ed9160fc2d07405ba0037858fbd4ea35ea56e565b533ebd1c42bf52d54.wgsl

@@ -0,0 +1,31 @@
+@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> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b1[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 64u + col / 32u];
+}
+
+var<workgroup> tile_a: array<f32, 64>;
+var<workgroup> tile_b: array<f32, 64>;
+@compute @workgroup_size(8, 8)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.y * 8u + local.y;
+    let col = group.x * 8u + local.x;
+    let i = (group.z * 64u + row) * 1024u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 2048u; tile += 8u) {
+        tile_a[local.y * 8u + local.x] = 0.0;
+        tile_b[local.x * 8u + local.y] = 0.0;
+        if (row < 64u && tile + local.x < 2048u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 64u + row) * 2048u + tile + local.x]); }
+        if (group.x * 8u + local.y < 1024u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = dequant_1((group.x * 8u + local.y) * 2048u + tile + local.x); }
+        workgroupBarrier();
+        for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
+        workgroupBarrier();
+    }
+    if (row < 64u && col < 1024u) { out[i] = f32(acc + 0.0); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/63324a9b275240b6f22bd9cdfef03aa98c69c89fbce0dc38f96e7b9eeabd09d8.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 * 128u;
+if (i >= 128u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(cos(x));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/6460788c21ff27c6ed7a52ce76e1ebc79b04274ccc1dedc81729c7df37e0325f.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) * 4096u + ((i / 256u) % 8u) * 512u + ((i / 64u) % 4u) * 128u + (((i / 1u) % 64u) * 1u + 64u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/646a62a1ea8c9249e411bda1c429f7412c618fc87d03bc0fd0d8fcc3c57027f3.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 * 1u;
+    if (row >= 1u) { return; }
+    let lane = local.x;
+    if (lane == 0u) {
+        var total = 0.0;
+        for (var j = 0u; j < 2048u; j++) {
+            let v = f32(b0[row * 2048u + j]);
+            total += v * v;
+        }
+        factor = inverseSqrt(total / 2048.0 + 1e-06);
+    }
+    workgroupBarrier();
+    for (var p = lane; p < 2048u; p += 64u) {
+        let i = row * 2048u + p;
+        out[i] = f32(f32(b0[row * 2048u + p]) * factor * (f32(b1[p]) + 0.0));
+    }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/655b09ac131b42d138be905d824f315d4d55ceafbd2bf4315aa07be144b1813b.wgsl

@@ -0,0 +1,17 @@
+@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 * 4194240u;
+if (i >= 4194304u) { return; }
+let batch = i / 262144u;
+let row = (i / 4096u) % 64u;
+let col = i % 4096u;
+var acc = 0.0;
+for (var p = 0u; p < 128u; p++) {
+    acc += f32(b0[(((batch / 1u) % 16u) * 8192u) + row * 128u + p]) * f32(b1[(((batch / 1u) % 16u) * 524288u) + p * 4096u + col]);
+}
+out[i] = f32(acc);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/65925fb31e923a32c9aab8a7ed72d255f6439dcbb3e604523d87995833be9c05.wgsl

@@ -0,0 +1,31 @@
+@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> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b1[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 64u + col / 32u];
+}
+
+var<workgroup> tile_a: array<f32, 64>;
+var<workgroup> tile_b: array<f32, 64>;
+@compute @workgroup_size(8, 8)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.y * 8u + local.y;
+    let col = group.x * 8u + local.x;
+    let i = (group.z * 16u + row) * 1024u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 2048u; tile += 8u) {
+        tile_a[local.y * 8u + local.x] = 0.0;
+        tile_b[local.x * 8u + local.y] = 0.0;
+        if (row < 16u && tile + local.x < 2048u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 16u + row) * 2048u + tile + local.x]); }
+        if (group.x * 8u + local.y < 1024u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = dequant_1((group.x * 8u + local.y) * 2048u + tile + local.x); }
+        workgroupBarrier();
+        for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
+        workgroupBarrier();
+    }
+    if (row < 16u && col < 1024u) { out[i] = f32(acc + 0.0); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/6633b8c4bd21f4e281eb976cdf9e94037ea3de39224dc79ad89e46b6569f5a60.wgsl

@@ -0,0 +1,31 @@
+@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> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b1[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 64u + col / 32u];
+}
+
+var<workgroup> tile_a: array<f32, 64>;
+var<workgroup> tile_b: array<f32, 64>;
+@compute @workgroup_size(8, 8)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.y * 8u + local.y;
+    let col = group.x * 8u + local.x;
+    let i = (group.z * 4u + row) * 6144u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 2048u; tile += 8u) {
+        tile_a[local.y * 8u + local.x] = 0.0;
+        tile_b[local.x * 8u + local.y] = 0.0;
+        if (row < 4u && tile + local.x < 2048u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 4u + row) * 2048u + tile + local.x]); }
+        if (group.x * 8u + local.y < 6144u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = dequant_1((group.x * 8u + local.y) * 2048u + tile + local.x); }
+        workgroupBarrier();
+        for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
+        workgroupBarrier();
+    }
+    if (row < 4u && col < 6144u) { out[i] = f32(acc + 0.0); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/6691194f7f41d1b69e72001549fe5d15614e48b5e1b629c18434fcbdb76f5818.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 * 320u;
+if (i >= 320u) { return; }
+out[i] = f32(b0[((i / 320u) % 1u) * 1024u + ((i / 20u) % 16u) * 64u + (((i / 1u) % 20u) * 3u + 2u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/66d2c54f014fcbb9053b0710f63280a933d6ad881f5af3c6df4b553c56b53a53.wgsl

@@ -0,0 +1,31 @@
+@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> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 6144u;
+    let col = index % 6144u;
+    let word = b1[row * 1536u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 192u + col / 32u];
+}
+
+var<workgroup> tile_a: array<f32, 64>;
+var<workgroup> tile_b: array<f32, 64>;
+@compute @workgroup_size(8, 8)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.y * 8u + local.y;
+    let col = group.x * 8u + local.x;
+    let i = (group.z * 4u + row) * 2048u + col;
+    var acc = 0.0;
+    for (var tile = 0u; tile < 6144u; tile += 8u) {
+        tile_a[local.y * 8u + local.x] = 0.0;
+        tile_b[local.x * 8u + local.y] = 0.0;
+        if (row < 4u && tile + local.x < 6144u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 4u + row) * 6144u + tile + local.x]); }
+        if (group.x * 8u + local.y < 2048u && tile + local.x < 6144u) { tile_b[local.x * 8u + local.y] = dequant_1((group.x * 8u + local.y) * 6144u + tile + local.x); }
+        workgroupBarrier();
+        for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
+        workgroupBarrier();
+    }
+    if (row < 4u && col < 2048u) { out[i] = f32(acc + 0.0); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/68478053918a18ace8851c6bd134cce3d9a19cdc8759e21d4ce6929fc12826a9.wgsl

@@ -0,0 +1,17 @@
+@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 * 192u;
+if (i >= 192u) { return; }
+let batch = i / 64u;
+let row = (i / 1u) % 64u;
+let col = i % 1u;
+var acc = 0.0;
+for (var p = 0u; p < 1u; p++) {
+    acc += f32(b0[(((batch / 1u) % 3u) * 64u) + row * 1u + p]) * f32(b1[(((batch / 1u) % 3u) * 1u) + p * 1u + col]);
+}
+out[i] = f32(acc);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/687a9faf59b35be173c2861749834333b53400d1dd22bf430cef04722237a309.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 * 8192u;
+if (i >= 8192u) { return; }
+let token = i32(b0[i / 2048u]);
+if (token < 98304 || token >= 131072) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 98304) * 2048u + i % 2048u));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/697e591c5cad8129c8363ebdde06e6dce400c2f33c323b4ac13f2a770506674e.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 * 512u;
+    if (row >= 512u) { 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));
+    }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/6ad3e0f21f70c7cb284d8f0115a3e89349d86bb63af4eae8aa70d6c3b1ae1131.wgsl

@@ -0,0 +1,28 @@
+@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> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 6144u;
+    let col = index % 6144u;
+    let word = b1[row * 1536u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 192u + col / 32u];
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 2048u;
+    if (i >= 2048u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 6144u; p += 64u) { acc += f32(b0[(i / 2048u) * 6144u + p]) * dequant_1((i % 2048u) * 6144u + 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); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/6b10aa16f1d7480a310e4afb48f199a3cfe80b086a7a1a927d85c551feaa3c14.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 * 192u;
+if (i >= 192u) { return; }
+out[i] = f32(b0[((i / 64u) % 3u) * 64u + ((i / 64u) % 1u) * 64u + ((i / 64u) % 1u) * 1u + ((i / 1u) % 64u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/6d4b992e3d432d1a6fcae8a5477aaef20bdfbec47bf39a06cd5be4f9661cc034.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 * 320u;
+if (i >= 320u) { return; }
+out[i] = f32(b0[((i / 320u) % 1u) * 1024u + ((i / 20u) % 16u) * 64u + (((i / 1u) % 20u) * 3u + 1u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/6d7c323a78bcae8ff287097ca35860e4a282d37acd4aa6e94680904e7e5afd68.wgsl

@@ -0,0 +1,28 @@
+@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> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b1[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 64u + col / 32u];
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 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]) * dequant_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); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/6e158cae64ceafa60c5498129bbc4482684844432772c61a9f2cfb62482bdeb4.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 * 8192u;
+if (i >= 8192u) { return; }
+let token = i32(b0[i / 2048u]);
+if (token < 32768 || token >= 65536) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 32768) * 2048u + i % 2048u));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/7183997261247fa4df7e081c0d6d01036fee95b9b57ef944d5f44f72fc80d5fe.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<i32>;
+
+@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] = i32(b0[i]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/71d38690bcc7fc30f01f1c8fa2f13a38beff889fcb3ff387861bb70038e92129.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<i32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 192u;
+if (i >= 192u) { return; }
+out[i] = i32(b0[(((i / 64u) % 3u) * 1u + 1u) * 64u + ((i / 64u) % 1u) * 64u + ((i / 1u) % 64u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/72474f821f837da4ebaab4cc925fdac78b707346be174a7c033c4c6518935abc.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 * 131072u;
+if (i >= 131072u) { return; }
+out[i] = f32(b0[((i / 131072u) % 1u) * 131072u + ((i / 2048u) % 64u) * 128u + ((i / 128u) % 16u) * 8192u + ((i / 1u) % 128u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/72ec1c10edeab7406245c736064375f9f5903562b2e27c1a4978d20a2bf313e7.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 >= 64u) { return; }
+let coord = (i / 1u) % 64u;
+if (coord >= 1u && coord < 61u && (coord - 1u) % 3u == 0u) { out[i] = f32(b0[((i / 64u) % 1u) * 20u + ((i / 64u) % 1u) * 20u + ((((i / 1u) % 64u) - 1u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/770ce07c0ddf504d2c14fd8e896c5b536200e3de22645024b090d26fa4f06f1a.wgsl

@@ -0,0 +1,13 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<i32>;
+@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 * 16384u;
+    if (i >= 16384u) { return; }
+    let token = (i / 128u) % 16u;
+    let outer = i / 2048u;
+    let destination = outer * 524288u + u32(b1[token]) * 128u + i % 128u;
+    out[((destination) / 128u) % 4096u * 1024u + ((destination) / 524288u) % 8u * 128u + ((destination) / 1u) % 128u * 1u] = f32(b0[i]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/77ae1020e9b490014bbfc055f2f83ebe28e564ff43549d584c1286dc318468ca.wgsl

@@ -0,0 +1,18 @@
+@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> b2: array<f32>;
+@group(0) @binding(3) var<storage, read> b3: array<f32>;
+@group(0) @binding(4) var<storage, read> b4: array<f32>;
+@group(0) @binding(5) 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 * 151936u;
+if (i >= 151936u) { return; }
+let coord = (i / 1u) % 151936u;
+if (coord >= 0u && coord < 32768u) { out[i] = f32(b0[(i / 151936u) * 32768u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 32768u && coord < 65536u) { out[i] = f32(b1[(i / 151936u) * 32768u + (coord - 32768u) * 1u + i % 1u]); }
+if (coord >= 65536u && coord < 98304u) { out[i] = f32(b2[(i / 151936u) * 32768u + (coord - 65536u) * 1u + i % 1u]); }
+if (coord >= 98304u && coord < 131072u) { out[i] = f32(b3[(i / 151936u) * 32768u + (coord - 98304u) * 1u + i % 1u]); }
+if (coord >= 131072u && coord < 151936u) { out[i] = f32(b4[(i / 151936u) * 20864u + (coord - 131072u) * 1u + i % 1u]); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/79cf7e6470bd1d26eab0b3cba06973b0e5247e5b01f9d7495d5fbf188b77456e.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 * 65536u;
+if (i >= 65536u) { return; }
+out[i] = f32(f32(b0[i]) * f32(b1[((i / 128u) % 64u) * 128u + ((i / 1u) % 128u) * 1u]));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/7a1f3615976306dc135893dfe7b56152d40162e002905883253b70fb3b450462.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 >= 20u) { return; }
+out[i] = f32(b0[((i / 20u) % 1u) * 64u + ((i / 20u) % 1u) * 64u + (((i / 1u) % 20u) * 3u + 1u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/7b2377ed11d1233f699351e3b23e150bdc0433606c21f49d43459cb31287d655.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 * 65536u;
+if (i >= 65536u) { return; }
+out[i] = f32(b0[((i / 65536u) % 1u) * 65536u + ((i / 8192u) % 8u) * 128u + ((i / 128u) % 64u) * 1024u + ((i / 1u) % 128u) * 1u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/7df80ebb55df6d6f9f5734d43f2529398210d7fa7b4a772c68b7f48263d3c10b.wgsl

@@ -0,0 +1,17 @@
+@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 * 131072u;
+if (i >= 131072u) { return; }
+let batch = i / 8192u;
+let row = (i / 128u) % 64u;
+let col = i % 128u;
+var acc = 0.0;
+for (var p = 0u; p < 4096u; p++) {
+    acc += f32(b0[(((batch / 1u) % 16u) * 262144u) + row * 4096u + p]) * f32(b1[(((batch / 1u) % 16u) * 524288u) + p * 128u + col]);
+}
+out[i] = f32(acc);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/7fe183583b43b627721ba6853b306288d7ceea4b42b630410f47ea9637220a4b.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]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/800fcd669d8d837a4c8767430651b5988e031388e3bf929f931754cabfade9fe.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 * 12288u;
+if (i >= 12288u) { return; }
+out[i] = f32(b0[((i / 4096u) % 3u) * 4096u + ((i / 4096u) % 1u) * 4096u + ((i / 64u) % 64u) * 1u + ((i / 1u) % 64u) * 64u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/8149829caefecbe2a94b68cec2ab6913de07685230d9aba67e4f834cad2824e8.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 * 8192u;
+if (i >= 8192u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(sin(x));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/82bc226410ee45662442e25040b6a6456c2532b87e96860e10f8f0135dde892b.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<i32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 64u;
+if (i >= 12u) { return; }
+out[i] = i32(b0[i]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/83241bf2399b1a681ba9b03d0bb4f2c380a62fb1e15afa5dc29628af17696371.wgsl

@@ -0,0 +1,17 @@
+@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 * 65536u;
+if (i >= 65536u) { return; }
+let batch = i / 4096u;
+let row = (i / 4096u) % 1u;
+let col = i % 4096u;
+var acc = 0.0;
+for (var p = 0u; p < 128u; p++) {
+    acc += f32(b0[(((batch / 1u) % 16u) * 128u) + row * 128u + p]) * f32(b1[(((batch / 1u) % 16u) * 524288u) + p * 4096u + col]);
+}
+out[i] = f32(acc);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/8485c3edc4d8013903fedc22091c7530a69a97da8b3dbd0cc5c55f54220709fb.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 * 131072u;
+if (i >= 131072u) { return; }
+let token = i32(b0[i / 2048u]);
+if (token < 65536 || token >= 98304) { out[i] = f32(0.0); return; }
+out[i] = f32(unpack_bf16_1(u32(token - 65536) * 2048u + i % 2048u));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/872b58c4d6944bc801a906cd980859e91ed8d44abccd0498bafa42d19f37eb46.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 * 65536u;
+if (i >= 65536u) { return; }
+out[i] = f32(f32(b0[i]) + (f32(b1[((i / 1u) % 4096u) * 1u]) * 1.0));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/897a6219da4cd27253c3b37bb46f07c7327701b764b5207eefa2cdef4c3ace99.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 * 131072u;
+if (i >= 131072u) { return; }
+out[i] = f32(f32(b0[i]) * f32(b1[((i / 128u) % 64u) * 128u + ((i / 1u) % 128u) * 1u]));
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/8acff95e1c324ceae0104784991a5262b08fc893c5b8bd3c7785579cc9958b0f.wgsl

@@ -0,0 +1,28 @@
+@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> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b1[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 64u + col / 32u];
+}
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let i = group.x + group.y * 6144u;
+    if (i >= 6144u) { return; }
+    let lane = local.x;
+    var acc = 0.0;
+    for (var p = lane; p < 2048u; p += 64u) { acc += f32(b0[(i / 6144u) * 2048u + p]) * dequant_1((i % 6144u) * 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); }
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/8cd9502df94e141090c88a4f43475f9b148c2ddfee89de96f96fe508fddb0675.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 * 1280u;
+if (i >= 1280u) { return; }
+out[i] = f32(b0[i]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/8dee3a06615c8adc6a1d111710c2f50ce1520d2ef5884b4da3d07c57d3816285.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 * 3072u;
+if (i >= 3072u) { return; }
+out[i] = f32(b0[((i / 1024u) % 3u) * 1024u + ((i / 1024u) % 1u) * 1024u + ((i / 64u) % 16u) * 1u + ((i / 1u) % 64u) * 16u]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/8df059dc71073b4ab38f6ff5685a6f52006e0b2264241c26bbc4ee288b9470a9.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(f32(b0[i]) * 1.0);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/8edd8a2bae8f26ae570c97da27c254da93d6b2df00784f505002f97830f806b3.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 * 32768u;
+if (i >= 32768u) { return; }
+out[i] = f32(b0[i]);
+}

qwen3vl-instruct-fp32-int8-g32-home-token-major-v2/kernels/90b5e8163a29d55ef2607822c57dd3ff7a824f7b0dbf46ac9aaf264d51d89bbd.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 * 131072u;
+if (i >= 131072u) { return; }
+let coord = (i / 1u) % 128u;
+if (coord >= 0u && coord < 64u) { out[i] = f32(b0[(i / 128u) * 64u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 64u && coord < 128u) { out[i] = f32(b1[(i / 128u) * 64u + (coord - 64u) * 1u + i % 1u]); }
+}