Add exaone4-fp32-int8-g32-release-token-major-v2-v4 (tiled prefill GEMM over the v3 bundle)

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/004e921ad38af5420a00fa5fbba1919f15780198ddfe6df05a31ce667bcd84e2.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 * 4194240u;
+if (i >= 6553600u) { return; }
+let coord = (i / 1u) % 102400u;
+if (coord >= 0u && coord < 65536u) { out[i] = f32(b0[(i / 102400u) * 65536u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 65536u && coord < 102400u) { out[i] = f32(b1[(i / 102400u) * 36864u + (coord - 65536u) * 1u + i % 1u]); }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/0112d93bb1aa347d8887ce1b95610a1a21fc8450ff8f59bcd37747fa6bf6fa52.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 * 2048u;
+if (i >= 2048u) { return; }
+out[i] = f32(f32(b0[i]) + (f32(b1[i]) * 1.0));
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/0337d8a473b6e485bd383533df7d373dbbfe7a5f3a982424004560b7213fbaa3.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 * 4096u;
+if (i >= 4096u) { 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]); }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/036902bfb564d824fa22d0ba0d8ed16a5d58e362fa585035f488ba7b806cefea.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 / 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<array<f32, 16>, 16>;
+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 * 16u;
+    var acc: array<array<f32, 4>, 1>;
+    for (var k0 = 0u; k0 < 2048u; k0 += 16u) {
+        for (var e = 0u; e < 1u; 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 < 4u && col < 2048u) { value = f32(f32(b0[(batch * 4u + row) * 2048u + 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 < 512u && first < 2048u) {
+                word = b1[n_index * 512u + first / 4u];
+                scale = f32(b2[n_index * 64u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 2048u) { 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 < 1u; r++) {
+                let a_value = tile_a[local.y * 1u + r][kk];
+                for (var c = 0u; c < 4u; c++) { acc[r][c] += a_value * b_values[c]; }
+            }
+        }
+        workgroupBarrier();
+    }
+    for (var r = 0u; r < 1u; r++) {
+        let out_row = tile_row + local.y * 1u + r;
+        for (var c = 0u; c < 4u; c++) {
+            let out_col = tile_n + local.x * 4u + c;
+            if (out_row < 4u && out_col < 512u) {
+                let i = (batch * 4u + out_row) * 512u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/07cfabe2e2c2feff22a06c8296600e5ca21efb7444419f0a05e5ff6f48e0292d.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 * 1024u;
+if (i >= 1024u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(-x);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/0b8815b1efe176d55ab08530c0a9e1913cf6c9ef3e773c574843573d237cf196.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(sin(x));
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/0d3892883d1532d173a7dbe85d9ab1daa87f0d9b4923a535b8d2003f24f9500c.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 * 256u;
+if (i >= 256u) { 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]); }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/0e2ba0d1849cd945e9b533bfc599a5664fa7909515c955681084ab8ad82f8755.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 / 512u) % 1u) * 512u + ((i / 64u) % 8u) * 64u + ((i / 64u) % 1u) * 512u + ((i / 1u) % 64u) * 1u]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/1024c1444e2cb17a211ebed96e18db764f1809d6e986d09e75f482ab4e61a07b.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 / 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<array<f32, 16>, 16>;
+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 * 16u;
+    var acc: array<array<f32, 4>, 1>;
+    for (var k0 = 0u; k0 < 2048u; k0 += 16u) {
+        for (var e = 0u; e < 1u; 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 < 16u && col < 2048u) { value = f32(f32(b0[(batch * 16u + row) * 2048u + 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 < 2048u && first < 2048u) {
+                word = b1[n_index * 512u + first / 4u];
+                scale = f32(b2[n_index * 64u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 2048u) { 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 < 1u; r++) {
+                let a_value = tile_a[local.y * 1u + r][kk];
+                for (var c = 0u; c < 4u; c++) { acc[r][c] += a_value * b_values[c]; }
+            }
+        }
+        workgroupBarrier();
+    }
+    for (var r = 0u; r < 1u; r++) {
+        let out_row = tile_row + local.y * 1u + r;
+        for (var c = 0u; c < 4u; c++) {
+            let out_col = tile_n + local.x * 4u + c;
+            if (out_row < 16u && out_col < 2048u) {
+                let i = (batch * 16u + out_row) * 2048u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/133b3d1c9ddb9953a9d3cffa5e3bd5e837cf2474ec6641a311a54d911d958f6b.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) * 131072u + ((i / 2048u) % 32u) * 4096u + ((i / 32u) % 64u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/1533ec909b782828ba90ded15ebe230c6f9dbfbeb74335606f979f2b56c58dbf.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 / 1024u) % 1u) * 2048u + ((i / 32u) % 32u) * 64u + ((i / 32u) % 1u) * 64u + (((i / 1u) % 32u) * 1u + 0u) * 1u]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/18a66ec0875c456c9defd45d15e6ef0c148e64a1618e2825c7f96eadbbe8303a.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 >= 16u) { return; }
+out[i] = i32(b0[i]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/1a73b6fceda28f40bba4ae27a986ffa39842db94fa5903ea5bf689a7e854d4f3.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 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]); }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/1ae471d3a9fd6a17efdb02ac9de0d56089b0df79cd834ba8d34f2417c400a51c.wgsl

@@ -0,0 +1,83 @@
+@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_write> out: array<f32>;
+
+var<workgroup> q_row: array<f32, 64>;
+var<workgroup> scores: array<f32, 1024>;
+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 lane = local.x;
+    let head = group.x / 4u;
+    let row = group.x % 4u;
+    let kv_head = head / 4u;
+    for (var d = lane; d < 64u; d += 64u) { q_row[d] = f32(b0[(head * 4u + row) * 64u + d]); }
+    workgroupBarrier();
+    var running_max = -3.4028234663852886e38;
+    var running_sum = 0.0;
+    var acc: array<f32, 1>;
+    for (var start = 0u; start < 4096u; start += 1024u) {
+        let count = min(1024u, 4096u - start);
+        var tile_max = -3.4028234663852886e38;
+        for (var t = lane; t < count; t += 64u) {
+            let j = start + t;
+            var score = -3.4028234663852886e38;
+            if (f32(b3[row * 4096u + j]) > -1000.0) {
+                var dot = 0.0;
+                for (var d = 0u; d < 64u; d++) { dot += q_row[d] * f32(b1[(((kv_head * 4096u + j) * 64u + d) / 64u) % 4096u * 512u + (((kv_head * 4096u + j) * 64u + d) / 262144u) % 8u * 64u + (((kv_head * 4096u + j) * 64u + d) / 1u) % 64u * 1u]); }
+                score = dot * 0.125 + f32(b3[row * 4096u + j]);
+            }
+            scores[t] = score;
+            tile_max = max(tile_max, score);
+        }
+        partial[lane] = tile_max;
+        workgroupBarrier();
+        for (var stride = 32u; stride > 0u; stride /= 2u) {
+            if (lane < stride) { partial[lane] = max(partial[lane], partial[lane + stride]); }
+            workgroupBarrier();
+        }
+        let next_max = max(running_max, partial[0]);
+        workgroupBarrier();
+        var tile_sum = 0.0;
+        for (var t = lane; t < count; t += 64u) {
+            var p = 0.0;
+            if (scores[t] > -3.0e38) { p = exp(scores[t] - next_max); }
+            scores[t] = p;
+            tile_sum += p;
+        }
+        partial[lane] = tile_sum;
+        workgroupBarrier();
+        for (var stride = 32u; stride > 0u; stride /= 2u) {
+            if (lane < stride) { partial[lane] += partial[lane + stride]; }
+            workgroupBarrier();
+        }
+        var correction = 0.0;
+        if (running_max > -3.0e38) { correction = exp(running_max - next_max); }
+        running_sum = running_sum * correction + partial[0];
+        running_max = next_max;
+        for (var c = 0u; c < 1u; c++) {
+            let d = lane + c * 64u;
+            var total = acc[c] * correction;
+            if (d < 64u) {
+                for (var t = 0u; t < count; t++) {
+                    let p = scores[t];
+                    if (p != 0.0) {
+                        let j = start + t;
+                        total += p * f32(b2[(((kv_head * 4096u + j) * 64u + d) / 64u) % 4096u * 512u + (((kv_head * 4096u + j) * 64u + d) / 262144u) % 8u * 64u + (((kv_head * 4096u + j) * 64u + d) / 1u) % 64u * 1u]);
+                    }
+                }
+            }
+            acc[c] = total;
+        }
+        workgroupBarrier();
+    }
+    for (var c = 0u; c < 1u; c++) {
+        let d = lane + c * 64u;
+        if (d < 64u) {
+            let i = (head * 4u + row) * 64u + d;
+            out[i] = f32(acc[c] / running_sum);
+        }
+    }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/1aec53ecf59335904944de2db37f0d3df64e3445556b40324ce611d4c9fcfc2c.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 / 1024u) % 1u) * 2048u + ((i / 32u) % 32u) * 64u + ((i / 32u) % 1u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/1c5b119286250e322ec8f75a46494057ed2485c4ce8ea081b4858aa5fbeadd0d.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 / 4096u) % 1u) * 8192u + ((i / 128u) % 32u) * 256u + ((i / 32u) % 4u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/1d3137ed94b2dd9113466681ea64a1e9a3086acbadc3c2e2a35f30dc42c67ad5.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 / 4096u) % 1u) * 8192u + ((i / 128u) % 32u) * 256u + ((i / 32u) % 4u) * 64u + (((i / 1u) % 32u) * 1u + 0u) * 1u]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/1f75ad93914b4f5fa3bc44e3995f201125f28abea03eae845b3b829a6ee39b11.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 * 8192u;
+    if (i >= 8192u) { return; }
+    let token = (i / 64u) % 16u;
+    let outer = i / 1024u;
+    let destination = outer * 262144u + u32(b1[token]) * 64u + i % 64u;
+    out[((destination) / 64u) % 4096u * 512u + ((destination) / 262144u) % 8u * 64u + ((destination) / 1u) % 64u * 1u] = f32(b0[i]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/224943172cdf191aa761ddce6dc4bc70265cda77a27b2a5fb69e9824a3dc494a.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 / 4096u;
+    let col = index % 4096u;
+    let word = b1[row * 1024u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 128u + col / 32u];
+}
+
+var<workgroup> tile_a: array<array<f32, 16>, 16>;
+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 * 16u;
+    var acc: array<array<f32, 4>, 1>;
+    for (var k0 = 0u; k0 < 4096u; k0 += 16u) {
+        for (var e = 0u; e < 1u; 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 < 16u && col < 4096u) { value = f32(f32(b0[(batch * 16u + row) * 4096u + 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 < 2048u && first < 4096u) {
+                word = b1[n_index * 1024u + first / 4u];
+                scale = f32(b2[n_index * 128u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 4096u) { 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 < 1u; r++) {
+                let a_value = tile_a[local.y * 1u + r][kk];
+                for (var c = 0u; c < 4u; c++) { acc[r][c] += a_value * b_values[c]; }
+            }
+        }
+        workgroupBarrier();
+    }
+    for (var r = 0u; r < 1u; r++) {
+        let out_row = tile_row + local.y * 1u + r;
+        for (var c = 0u; c < 4u; c++) {
+            let out_col = tile_n + local.x * 4u + c;
+            if (out_row < 16u && out_col < 2048u) {
+                let i = (batch * 16u + out_row) * 2048u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/2602f4b29f873a095f9c0b226b2bcc6f2379956c7469b7d2ee81722d97d06b36.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 / 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];
+}
+
+@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 >= 102400) { out[i] = f32(0.0); return; }
+out[i] = f32(dequant_1(u32(token - 65536) * 2048u + i % 2048u));
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/2790fbaffce918ef00d053c0f83f62375a0a4ac95aa9ec58cb066b2aee41d18a.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 / 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<array<f32, 16>, 16>;
+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 * 16u;
+    var acc: array<array<f32, 4>, 1>;
+    for (var k0 = 0u; k0 < 2048u; k0 += 16u) {
+        for (var e = 0u; e < 1u; 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 < 16u && col < 2048u) { value = f32(f32(b0[(batch * 16u + row) * 2048u + 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 < 512u && first < 2048u) {
+                word = b1[n_index * 512u + first / 4u];
+                scale = f32(b2[n_index * 64u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 2048u) { 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 < 1u; r++) {
+                let a_value = tile_a[local.y * 1u + r][kk];
+                for (var c = 0u; c < 4u; c++) { acc[r][c] += a_value * b_values[c]; }
+            }
+        }
+        workgroupBarrier();
+    }
+    for (var r = 0u; r < 1u; r++) {
+        let out_row = tile_row + local.y * 1u + r;
+        for (var c = 0u; c < 4u; c++) {
+            let out_col = tile_n + local.x * 4u + c;
+            if (out_row < 16u && out_col < 512u) {
+                let i = (batch * 16u + out_row) * 512u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/27a921d3b5edfdca227a99e1c3328959881c09acb02b12cd09d07d598745b2d3.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 >= 1u) { return; }
+out[i] = f32(b0[i]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/28119aa01122c42bd26ceb37129fe0c95b2a9192163f6dabd90fe058b1e14c54.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 * 4096u;
+if (i >= 4096u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(sin(x));
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/284513e54dd254e923b7a86bc5703879e3bfb119b6c71aa234ba744c2dc63615.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 / 32u) % 64u) * 1u + ((i / 1u) % 32u) * 64u]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/28c076a43a25e4705518eefb1c3a2699493c6d0a5734baa14532dfa4826ed566.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 * 102400u;
+if (i >= 102400u) { return; }
+let coord = (i / 1u) % 102400u;
+if (coord >= 0u && coord < 65536u) { out[i] = f32(b0[(i / 102400u) * 65536u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 65536u && coord < 102400u) { out[i] = f32(b1[(i / 102400u) * 36864u + (coord - 65536u) * 1u + i % 1u]); }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/2b9e4b05ef5623d721dbe56b5812e551efcc6fd833d81064320e86fe94d8c670.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(f32(b0[i]) * 1.0);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/2c66dcac8fdd3caa9ff1e2441524070c21c24df9526bee44c49e765a8d4c8222.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 * 16384u;
+if (i >= 16384u) { return; }
+out[i] = f32(b0[((i / 16384u) % 1u) * 32768u + ((i / 2048u) % 8u) * 4096u + ((i / 32u) % 64u) * 64u + (((i / 1u) % 32u) * 1u + 0u) * 1u]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/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);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/30cb10d38275eecebb04ab10dc3246b7a30ff80578096f6efa538f30c602afb8.wgsl

@@ -0,0 +1,41 @@
+@group(0) @binding(0) var<storage, read> b0: array<u32>;
+@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_write> out: array<f32>;
+fn dequant_0(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b0[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b1[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 * 2048u;
+    if (i >= 2048u) { return; }
+    let lane = local.x;
+    let row = i % 2048u;
+    var acc = 0.0;
+    for (var w = lane; w < 512u; w += 64u) {
+        let word = b0[row * 512u + w];
+        let first = w * 4u;
+        var part = 0.0;
+        for (var e = 0u; e < 4u; e++) {
+            let col = first + e;
+            if (col < 2048u) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                part += f32(code) * f32(b2[(i / 2048u) * 2048u + col]);
+            }
+        }
+        acc += part * f32(b1[row * 64u + first / 32u]);
+    }
+    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); }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/3203e08f2c99fa4848f553739b7115a7c677840344f17b1d4391d6130517f5a7.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 >= 4u) { return; }
+out[i] = i32(b0[i]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/3241a118927e109c0dacbb278370abe70d099ee593a07cd023d4530b5fb8f833.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 * 8u;
+    if (row >= 8u) { return; }
+    let lane = local.x;
+    if (lane == 0u) {
+        var total = 0.0;
+        for (var j = 0u; j < 64u; j++) {
+            let v = f32(b0[row * 64u + j]);
+            total += v * v;
+        }
+        factor = inverseSqrt(total / 64.0 + 1e-05);
+    }
+    workgroupBarrier();
+    for (var p = lane; p < 64u; p += 64u) {
+        let i = row * 64u + p;
+        out[i] = f32(f32(b0[row * 64u + p]) * factor * (f32(b1[p]) + 0.0));
+    }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/33e8ec04c66adb4988f562f0f9d3e92d29dfda91a56d0ed62be598c7c737f5ba.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 * 16384u;
+if (i >= 16384u) { return; }
+out[i] = f32(b0[((i / 16384u) % 1u) * 32768u + ((i / 2048u) % 8u) * 4096u + ((i / 32u) % 64u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/37ce78f0b50bc59dc7b6dd62724d6ad3acfbfb73deae12446a4575f5688d8f34.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) * 128u + ((i / 32u) % 4u) * 1u + ((i / 1u) % 32u) * 4u]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/383d0ce2bbe657b6412b8e55e6b0cf06e385cb4efbe45019afff9357f2a77751.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 * 16384u;
+if (i >= 16384u) { return; }
+let x = f32(b0[i]); let s = f32(f32(x / (1.0 + exp(-x))));
+out[i] = f32(s * f32(b1[i]));
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/39646bdfafd1ab0cb874474e3f1e9d79b57f10582a1804bc1d3762a492fec851.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 / 64u) % 64u) * 64u + ((i / 1u) % 64u) * 1u]));
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/3c2cd3f4326cc231268e80b83abfa0520c88c9763ba44fb16df8357610c0a301.wgsl

@@ -0,0 +1,41 @@
+@group(0) @binding(0) var<storage, read> b0: array<u32>;
+@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_write> out: array<f32>;
+fn dequant_0(index: u32) -> f32 {
+    let row = index / 2048u;
+    let col = index % 2048u;
+    let word = b0[row * 512u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b1[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 * 4096u;
+    if (i >= 4096u) { return; }
+    let lane = local.x;
+    let row = i % 4096u;
+    var acc = 0.0;
+    for (var w = lane; w < 512u; w += 64u) {
+        let word = b0[row * 512u + w];
+        let first = w * 4u;
+        var part = 0.0;
+        for (var e = 0u; e < 4u; e++) {
+            let col = first + e;
+            if (col < 2048u) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                part += f32(code) * f32(b2[(i / 4096u) * 2048u + col]);
+            }
+        }
+        acc += part * f32(b1[row * 64u + first / 32u]);
+    }
+    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); }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/3e34409407100d4858c4c109cf62f3eaf2f4c6081166b0ec8679105e6dfbe7fe.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 * 16384u;
+if (i >= 16384u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(-x);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/402b4fbc568183c6fd53e1541bbe6f09d8bb4d70293eb96576db5f0d013d035d.wgsl

@@ -0,0 +1,83 @@
+@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_write> out: array<f32>;
+
+var<workgroup> q_row: array<f32, 64>;
+var<workgroup> scores: array<f32, 1024>;
+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 lane = local.x;
+    let head = group.x / 64u;
+    let row = group.x % 64u;
+    let kv_head = head / 4u;
+    for (var d = lane; d < 64u; d += 64u) { q_row[d] = f32(b0[(head * 64u + row) * 64u + d]); }
+    workgroupBarrier();
+    var running_max = -3.4028234663852886e38;
+    var running_sum = 0.0;
+    var acc: array<f32, 1>;
+    for (var start = 0u; start < 4096u; start += 1024u) {
+        let count = min(1024u, 4096u - start);
+        var tile_max = -3.4028234663852886e38;
+        for (var t = lane; t < count; t += 64u) {
+            let j = start + t;
+            var score = -3.4028234663852886e38;
+            if (f32(b3[row * 4096u + j]) > -1000.0) {
+                var dot = 0.0;
+                for (var d = 0u; d < 64u; d++) { dot += q_row[d] * f32(b1[(((kv_head * 4096u + j) * 64u + d) / 64u) % 4096u * 512u + (((kv_head * 4096u + j) * 64u + d) / 262144u) % 8u * 64u + (((kv_head * 4096u + j) * 64u + d) / 1u) % 64u * 1u]); }
+                score = dot * 0.125 + f32(b3[row * 4096u + j]);
+            }
+            scores[t] = score;
+            tile_max = max(tile_max, score);
+        }
+        partial[lane] = tile_max;
+        workgroupBarrier();
+        for (var stride = 32u; stride > 0u; stride /= 2u) {
+            if (lane < stride) { partial[lane] = max(partial[lane], partial[lane + stride]); }
+            workgroupBarrier();
+        }
+        let next_max = max(running_max, partial[0]);
+        workgroupBarrier();
+        var tile_sum = 0.0;
+        for (var t = lane; t < count; t += 64u) {
+            var p = 0.0;
+            if (scores[t] > -3.0e38) { p = exp(scores[t] - next_max); }
+            scores[t] = p;
+            tile_sum += p;
+        }
+        partial[lane] = tile_sum;
+        workgroupBarrier();
+        for (var stride = 32u; stride > 0u; stride /= 2u) {
+            if (lane < stride) { partial[lane] += partial[lane + stride]; }
+            workgroupBarrier();
+        }
+        var correction = 0.0;
+        if (running_max > -3.0e38) { correction = exp(running_max - next_max); }
+        running_sum = running_sum * correction + partial[0];
+        running_max = next_max;
+        for (var c = 0u; c < 1u; c++) {
+            let d = lane + c * 64u;
+            var total = acc[c] * correction;
+            if (d < 64u) {
+                for (var t = 0u; t < count; t++) {
+                    let p = scores[t];
+                    if (p != 0.0) {
+                        let j = start + t;
+                        total += p * f32(b2[(((kv_head * 4096u + j) * 64u + d) / 64u) % 4096u * 512u + (((kv_head * 4096u + j) * 64u + d) / 262144u) % 8u * 64u + (((kv_head * 4096u + j) * 64u + d) / 1u) % 64u * 1u]);
+                    }
+                }
+            }
+            acc[c] = total;
+        }
+        workgroupBarrier();
+    }
+    for (var c = 0u; c < 1u; c++) {
+        let d = lane + c * 64u;
+        if (d < 64u) {
+            let i = (head * 64u + row) * 64u + d;
+            out[i] = f32(acc[c] / running_sum);
+        }
+    }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/40f115c8647a238bfb3ccb0bb8446bccdc94a0dfadc89111c4920efb2259f41d.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 * 2048u;
+if (i >= 2048u) { return; }
+out[i] = f32(f32(b0[i]) * f32(b1[((i / 64u) % 4u) * 64u + ((i / 1u) % 64u) * 1u]));
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/431a95359c9c41766af53351f3074089c052c44df79b9d1d08205d8e58e90d25.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 * 32768u;
+    if (i >= 32768u) { return; }
+    let token = (i / 64u) % 64u;
+    let outer = i / 4096u;
+    let destination = outer * 262144u + u32(b1[token]) * 64u + i % 64u;
+    out[((destination) / 64u) % 4096u * 512u + ((destination) / 262144u) % 8u * 64u + ((destination) / 1u) % 64u * 1u] = f32(b0[i]);
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/44745cf68171fb657debd4b4fdd15edcd4d5798e18aedc4314d0cb1249948808.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 / 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];
+}
+
+@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 < 0 || token >= 65536) { out[i] = f32(0.0); return; }
+out[i] = f32(dequant_1(u32(token - 0) * 2048u + i % 2048u));
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/453668cbe741eb27e2cc10194eca66f514a3763d97b46c1733c376986defc322.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 * 4u;
+    if (row >= 4u) { 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-05);
+    }
+    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));
+    }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/45b47a1eff94975353fa58c943de182aa0da352efeb23f477ba2c69b8fc52b70.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 * 1638400u;
+if (i >= 1638400u) { return; }
+let coord = (i / 1u) % 102400u;
+if (coord >= 0u && coord < 65536u) { out[i] = f32(b0[(i / 102400u) * 65536u + (coord - 0u) * 1u + i % 1u]); }
+if (coord >= 65536u && coord < 102400u) { out[i] = f32(b1[(i / 102400u) * 36864u + (coord - 65536u) * 1u + i % 1u]); }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/45c3c3f9ae14f5837bb0cfc8cf8ce189010dba8f3b03d640b18883fc1367207b.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 / 4096u;
+    let col = index % 4096u;
+    let word = b1[row * 1024u + col / 4u];
+    let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
+    return f32(code) * b2[row * 128u + 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 < 4096u; 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 < 4096u) { value = f32(f32(b0[(batch * 64u + row) * 4096u + 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 < 2048u && first < 4096u) {
+                word = b1[n_index * 1024u + first / 4u];
+                scale = f32(b2[n_index * 128u + first / 32u]);
+            }
+            for (var e = 0u; e < 4u; e++) {
+                let code = i32((word >> (e * 8u)) & 255u) - 128;
+                var value = 0.0;
+                if (first + e < 4096u) { 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 < 2048u) {
+                let i = (batch * 64u + out_row) * 2048u + out_col;
+                out[i] = f32(acc[r][c] + 0.0);
+            }
+        }
+    }
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/483857d2d9929bc07cd14940dc3bca97abcf9485ce1c805a234d2db4110a7010.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 / 64u) % 16u) * 64u + ((i / 1u) % 64u) * 1u]));
+}

exaone4-fp32-int8-g32-release-token-major-v2-v4/kernels/4ec7950cea0e4257dc30630582df36a6127ce61794f46012db6ad0c0b1ea8380.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 * 64u;
+    if (row >= 64u) { 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-05);
+    }
+    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));
+    }
+}