Add qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6 (column-per-lane gated delta rule) (part 3)

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/96924710cde2ae782e03138b4399e2ed2406f8b37c31401f5a79af528c2193fd.wgsl

@@ -0,0 +1,20 @@
+@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> b2: array<f32>;
+@group(0) @binding(3) var<storage, read_write> out: array<f32>;
+fn dequant_1(index: u32) -> f32 {
+    let row = index / 1024u;
+    let col = index % 1024u;
+    let word = b1[row * 256u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 32u + col / 32u];
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 4096u;
+if (i >= 4096u) { return; }
+let token = i32(b0[i / 1024u]);
+if (token < 131072 || token >= 248320) { out[i] = f32(0.0); return; }
+out[i] = f32(dequant_1(u32(token - 131072) * 1024u + i % 1024u));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/9701df2bcd72ed6a0b72dca81006d9fcc2336e0250fbd66cf37d4fd80667358f.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) * 393216u + ((i / 8192u) % 16u) * 24576u + (((i / 128u) % 64u) * 1u + 128u) * 128u + ((i / 1u) % 128u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/970aa8fd65c7e96611a4b6f30cb0b71e24bdaefff0ca117243cab1a5a7e645a9.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 >= 48u) { return; }
+out[i] = f32(b0[i]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/97aa9b5cb2a41dbf920a6c15c61964d33dd0d87990a00183ac4b87c614bc376d.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 * 8192u;
+if (i >= 8192u) { 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]); }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/985bb9ad33654b9adbd4a256d626e346556a10dad514692800da8ebd7910f0f0.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 / 1u) % 3u) * 1u + 1u) * 1u + ((i / 1u) % 1u) * 1u + ((i / 1u) % 1u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/987d27946334b53b3006cdca1a7158b13a27e01e688047aeafc4b029761b80af.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) * 49152u + ((i / 4u) % 6144u) * 8u + (((i / 1u) % 4u) * 1u + 4u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/99e423f03f9fdfee41f3dcaab73257064ffb2a64110066752208324724505a8b.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) * 122880u + ((i / 4u) % 6144u) * 20u + (((i / 1u) % 4u) * 1u + 16u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/9af812cf7d1d31828c41aac6e5c1b9a55ca5cd015f8658d06bb078a5872becfb.wgsl

@@ -0,0 +1,71 @@
+@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 / 1024u;
+    let col = index % 1024u;
+    let word = b1[row * 256u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 32u + col / 32u];
+}
+
+var<workgroup> tile_a: array<array<f32, 16>, 64>;
+var<workgroup> tile_b: array<array<f32, 64>, 16>;
+@compute @workgroup_size(16, 16)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let lane = local.y * 16u + local.x;
+    let batch = group.z;
+    let tile_n = group.x * 64u;
+    let tile_row = group.y * 64u;
+    var acc: array<array<f32, 4>, 4>;
+    for (var k0 = 0u; k0 < 1024u; k0 += 16u) {
+        for (var e = 0u; e < 4u; e++) {
+            let flat = lane + e * 256u;
+            let m_local = flat / 16u;
+            let row = tile_row + m_local;
+            let col = k0 + flat % 16u;
+            var value = 0.0;
+            if (row < 64u && col < 1024u) { value = f32(f32(b0[(batch * 64u + row) * 1024u + col])); }
+            tile_a[m_local][flat % 16u] = value;
+        }
+        if (lane < 256u) {
+            let n_local = lane / 4u;
+            let word_local = lane % 4u;
+            let n_index = tile_n + n_local;
+            let first = k0 + word_local * 4u;
+            var word = 0u;
+            var scale = 0.0;
+            if (n_index < 6144u && first < 1024u) {
+                word = b1[n_index * 256u + first / 4u];
+                scale = f32(b2[n_index * 32u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 1024u) { value = f32(code) * scale; }
+                tile_b[word_local * 4u + e][n_local] = value;
+            }
+        }
+        workgroupBarrier();
+        for (var kk = 0u; kk < 16u; kk++) {
+            var b_values: array<f32, 4>;
+            for (var c = 0u; c < 4u; c++) { b_values[c] = tile_b[kk][local.x * 4u + c]; }
+            for (var r = 0u; r < 4u; r++) {
+                let a_value = tile_a[local.y * 4u + r][kk];
+                for (var c = 0u; c < 4u; c++) { acc[r][c] += a_value * b_values[c]; }
+            }
+        }
+        workgroupBarrier();
+    }
+    for (var r = 0u; r < 4u; r++) {
+        let out_row = tile_row + local.y * 4u + r;
+        for (var c = 0u; c < 4u; c++) {
+            let out_col = tile_n + local.x * 4u + c;
+            if (out_row < 64u && out_col < 6144u) {
+                let i = (batch * 64u + out_row) * 6144u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/9b7697d3e1de37277befbd725307f40a5143faae3dcf43182d4ad89d8572ed42.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(f32(b0[i]) * 1.0);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/9cd8236a30e65989f2ed9c7afcfcf2369b6036ade2a7d81c8e203c2789178074.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 * 417792u;
+if (i >= 417792u) { return; }
+let coord = (i / 1u) % 68u;
+if (coord >= 0u && coord < 4u) { out[i] = f32(b0[(i / 68u) * 4u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 4u && coord < 68u) { out[i] = f32(b1[(i / 68u) * 64u + (coord - 4u) * 1u + i % 1u]); }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/9d12ba7fd89692e99bb291808c5a80ffd5b95478671dde2c936fc3deb8383fa0.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 * 1024u;
+if (i >= 1024u) { 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-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/9d3d7b777701cd320ca11358e554e1019fb041ee4d765e66071c215242906773.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 / 32768u) % 1u) * 65536u + ((i / 2048u) % 16u) * 4096u + ((i / 256u) % 8u) * 512u + (((i / 1u) % 256u) * 1u + 0u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/9d504391a987c7be7ccfbb7972c27cf54cd62730c19cbe7949fee2ea059f0c43.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 * 122880u;
+if (i >= 122880u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(x / (1.0 + exp(-x)));
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/9d9666074828f92b63a676348da9d7bb80e5a03213f1bc88f4dd8607e8a26711.wgsl

@@ -0,0 +1,71 @@
+@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 / 1024u;
+    let col = index % 1024u;
+    let word = b1[row * 256u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 32u + col / 32u];
+}
+
+var<workgroup> tile_a: array<array<f32, 16>, 64>;
+var<workgroup> tile_b: array<array<f32, 64>, 16>;
+@compute @workgroup_size(16, 16)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let lane = local.y * 16u + local.x;
+    let batch = group.z;
+    let tile_n = group.x * 64u;
+    let tile_row = group.y * 64u;
+    var acc: array<array<f32, 4>, 4>;
+    for (var k0 = 0u; k0 < 1024u; k0 += 16u) {
+        for (var e = 0u; e < 4u; e++) {
+            let flat = lane + e * 256u;
+            let m_local = flat / 16u;
+            let row = tile_row + m_local;
+            let col = k0 + flat % 16u;
+            var value = 0.0;
+            if (row < 64u && col < 1024u) { value = f32(f32(b0[(batch * 64u + row) * 1024u + col])); }
+            tile_a[m_local][flat % 16u] = value;
+        }
+        if (lane < 256u) {
+            let n_local = lane / 4u;
+            let word_local = lane % 4u;
+            let n_index = tile_n + n_local;
+            let first = k0 + word_local * 4u;
+            var word = 0u;
+            var scale = 0.0;
+            if (n_index < 131072u && first < 1024u) {
+                word = b1[n_index * 256u + first / 4u];
+                scale = f32(b2[n_index * 32u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 1024u) { value = f32(code) * scale; }
+                tile_b[word_local * 4u + e][n_local] = value;
+            }
+        }
+        workgroupBarrier();
+        for (var kk = 0u; kk < 16u; kk++) {
+            var b_values: array<f32, 4>;
+            for (var c = 0u; c < 4u; c++) { b_values[c] = tile_b[kk][local.x * 4u + c]; }
+            for (var r = 0u; r < 4u; r++) {
+                let a_value = tile_a[local.y * 4u + r][kk];
+                for (var c = 0u; c < 4u; c++) { acc[r][c] += a_value * b_values[c]; }
+            }
+        }
+        workgroupBarrier();
+    }
+    for (var r = 0u; r < 4u; r++) {
+        let out_row = tile_row + local.y * 4u + r;
+        for (var c = 0u; c < 4u; c++) {
+            let out_col = tile_n + local.x * 4u + c;
+            if (out_row < 64u && out_col < 131072u) {
+                let i = (batch * 64u + out_row) * 131072u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/9f040facf0594804dba5b0f10f53b11b7223a1e1705896f95a8f95aa0fff5eec.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 * 2048u;
+if (i >= 2048u) { return; }
+let coord = (i / 1u) % 32u;
+if (coord >= 1u && coord < 34u && (coord - 1u) % 3u == 0u) { out[i] = f32(b0[((i / 2048u) % 1u) * 704u + ((i / 32u) % 64u) * 11u + ((((i / 1u) % 32u) - 1u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/a16a539c1a662c846d6c661e572b0b8e19335f9ef41742f784d738bcfd3cd7a6.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 * 393216u;
+if (i >= 393216u) { return; }
+out[i] = f32(b0[((i / 393216u) % 1u) * 417792u + ((i / 64u) % 6144u) * 68u + (((i / 1u) % 64u) * 1u + 4u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/a1a974f6d8d51fd7c3496ede0966af303ad9f0c0c9ee44cf93ec962d2b0e87e1.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 >= 64u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(-x);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/a1cf6b03baf18d389f0a2b95440c17125a8ecc1e633568de985d42cac6e0d1cb.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[1u * 512u + ((i / 512u) % 1u) * 512u + ((i / 32u) % 16u) * 32u + ((i / 1u) % 32u) * 1u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/a21c6a35532bc26ccfc46df22f98e1592e3b285e3800544c032edf39425bde1a.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 * 98304u;
+if (i >= 98304u) { return; }
+out[i] = f32(b0[((i / 98304u) % 1u) * 98304u + ((i / 16u) % 6144u) * 1u + ((i / 1u) % 16u) * 6144u]);
+}

qwen35-08b-fp32-int8-attn-int4-mlp-gptq-home-token-major-v2-v6/kernels/a24673da4fae6d1c49aa143f8425996b2d036db7b0341987dab330c78ead91d7.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 / 32768u) % 1u) * 98304u + ((i / 2048u) % 16u) * 6144u + (((i / 1u) % 2048u) * 1u + 4096u) * 1u]);
+}