Publish WebTorch model catalog bundles (part 3)

exaone4-fp32-8k-state-alias-token-major-v2/kernels/033ad0e2dfc972b75e5a7861b4f0f237c46b27fd9b1dcd22f01da3a1c2ad3035.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@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 = i32(b0[i / 2048u]);
+if (token < 98304 || token >= 102400) { out[i] = f32(0.0); return; }
+out[i] = f32(b1[u32(token - 98304) * 2048u + i % 2048u]);
+}

exaone4-fp32-8k-state-alias-token-major-v2/kernels/062917efbd690ad27ed43b70ec6f4fc9f5e0c592d87b22d2063d35eca11bd551.wgsl

@@ -0,0 +1,24 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+var<workgroup> 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) * 32768u + 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 < 32768u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = f32(b1[(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 < 32768u) { out[i] = f32(acc + 0.0); }
+}

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/097177b51ca66686a497098825ab4be5f52aa06560bbfd449e0cedc541a40c57.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 / 4096u;
+let row = (i / 64u) % 64u;
+let col = i % 64u;
+var acc = 0.0;
+for (var p = 0u; p < 8192u; p++) {
+    acc += f32(b0[(((batch / 1u) % 32u) * 524288u) + row * 8192u + p]) * f32(b1[(((batch / 1u) % 32u) * 524288u) + p * 64u + col]);
+}
+out[i] = f32(acc);
+}

exaone4-fp32-8k-state-alias-token-major-v2/kernels/09d52a47128189748341738ed18c539c6adb15107224cdf66db5c3f272ede305.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 * 2048u;
+    if (i >= 2048u) { return; }
+    let token = (i / 64u) % 4u;
+    let outer = i / 256u;
+    let destination = outer * 524288u + u32(b1[token]) * 64u + i % 64u;
+    out[((destination) / 64u) % 8192u * 512u + ((destination) / 524288u) % 8u * 64u + ((destination) / 1u) % 64u * 1u] = f32(b0[i]);
+}

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/0bf4667c6913011c300917b4b727dba8f21e2d6f0b36288e5140057654cf28a7.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 * 4194240u;
+if (i >= 4194304u) { return; }
+out[i] = f32(b0[((i) / 64u) % 8192u * 512u + ((i) / 524288u) % 8u * 64u + ((i) / 1u) % 64u * 1u]);
+}

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/0e5468583a1ba929483119c50d4be42a58aed4dde1e15cc231908ed7543d0b71.wgsl

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

exaone4-fp32-8k-state-alias-token-major-v2/kernels/0f0a1d8c9af01849a6831f1cdb0c409f56fa61ba180e05e1f61eb2cbf3db488c.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 * 262144u;
+if (i >= 262144u) { return; }
+let batch = i / 8192u;
+let row = (i / 8192u) % 1u;
+let col = i % 8192u;
+var acc = 0.0;
+for (var p = 0u; p < 64u; p++) {
+    acc += f32(b0[(((batch / 1u) % 32u) * 64u) + row * 64u + p]) * f32(b1[(((batch / 1u) % 32u) * 524288u) + p * 8192u + col]);
+}
+out[i] = f32(acc);
+}

exaone4-fp32-8k-state-alias-token-major-v2/kernels/104f307af02cae1435117d2cfafd3238b0a325bde8eb56746e4d33df51820c90.wgsl

@@ -0,0 +1,24 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+var<workgroup> 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 < 4096u; 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 < 4096u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 4u + row) * 4096u + tile + local.x]); }
+        if (group.x * 8u + local.y < 2048u && tile + local.x < 4096u) { tile_b[local.x * 8u + local.y] = f32(b1[(group.x * 8u + local.y) * 4096u + 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); }
+}

exaone4-fp32-8k-state-alias-token-major-v2/kernels/12155d86d776e73f23d3bff0268c60dccb54c209ee9ee3a8cfd227ec4aa5b6df.wgsl

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

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/14d88fb2ba88841a93259cafdf9a2cbf86411c82ec5da801fcd8a887de2ec80d.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@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 = i32(b0[i / 2048u]);
+if (token < 32768 || token >= 65536) { out[i] = f32(0.0); return; }
+out[i] = f32(b1[u32(token - 32768) * 2048u + i % 2048u]);
+}

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/1615236e57601bc3bf4f9f931c34ef9c3122e833d24f228064c0153ebe739005.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 * 1048576u;
+if (i >= 1048576u) { return; }
+out[i] = f32(f32(b0[i]) + (f32(b1[((i / 8192u) % 4u) * 8192u + ((i / 1u) % 8192u) * 1u]) * 1.0));
+}

exaone4-fp32-8k-state-alias-token-major-v2/kernels/1701e3705e0212caf2880d0a27ccb420ef83e5c95aeb2e9a0679eacf9c2e0a87.wgsl

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

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/193f33ea1b99ff98b1c55233fdfe8d81821056ba970909ee4bded551c2ff4811.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 / 131072u;
+let row = (i / 8192u) % 16u;
+let col = i % 8192u;
+var acc = 0.0;
+for (var p = 0u; p < 64u; p++) {
+    acc += f32(b0[(((batch / 1u) % 32u) * 1024u) + row * 64u + p]) * f32(b1[(((batch / 1u) % 32u) * 524288u) + p * 8192u + col]);
+}
+out[i] = f32(acc);
+}

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/214ac9825bc5014847277240f7d17b1f0aa4f2272175a7d77e94116dc75f9fff.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@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 token = i32(b0[i / 2048u]);
+if (token < 0 || token >= 32768) { out[i] = f32(0.0); return; }
+out[i] = f32(b1[u32(token - 0) * 2048u + i % 2048u]);
+}

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/2d74b64b4a60c8c6ab0338d5438b07fcf8fd140687826128ee7de0d3f8cd786c.wgsl

@@ -0,0 +1,24 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+
+var<workgroup> 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) * 512u + 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 < 512u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = f32(b1[(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 < 512u) { out[i] = f32(acc + 0.0); }
+}

exaone4-fp32-8k-state-alias-token-major-v2/kernels/2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a.wgsl

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

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/349a8efca8a73e71ccf2a924271666d936401c12b489afcf7581bdd3da2cbc27.wgsl

@@ -0,0 +1,34 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+var<workgroup> partial: array<f32, 64>;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
+    let row = group.x + group.y * 2048u;
+    if (row >= 2048u) { return; }
+    let lane = local.x;
+    var maximum = -3.4028234663852886e38;
+    for (var p = lane; p < 8192u; p += 64u) { maximum = max(maximum, f32(b0[row * 8192u + p])); }
+    partial[lane] = maximum;
+    workgroupBarrier();
+    for (var stride = 32u; stride > 0u; stride /= 2u) {
+        if (lane < stride) { partial[lane] = max(partial[lane], partial[lane + stride]); }
+        workgroupBarrier();
+    }
+    maximum = partial[0];
+    // All lanes must finish reading the maximum before the scratch array is reused.
+    workgroupBarrier();
+    var total = 0.0;
+    for (var p = lane; p < 8192u; p += 64u) { total += exp(f32(b0[row * 8192u + p]) - maximum); }
+    partial[lane] = total;
+    workgroupBarrier();
+    for (var stride = 32u; stride > 0u; stride /= 2u) {
+        if (lane < stride) { partial[lane] += partial[lane + stride]; }
+        workgroupBarrier();
+    }
+    total = partial[0];
+    for (var p = lane; p < 8192u; p += 64u) {
+        let i = row * 8192u + p;
+        out[i] = f32(exp(f32(b0[row * 8192u + p]) - maximum) / total);
+    }
+}

exaone4-fp32-8k-state-alias-token-major-v2/kernels/358e73d0a8169937dce5d2c7fd5f4dea41f7c8307696f1698967bc6c4132efc8.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@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 = i32(b0[i / 2048u]);
+if (token < 32768 || token >= 65536) { out[i] = f32(0.0); return; }
+out[i] = f32(b1[u32(token - 32768) * 2048u + i % 2048u]);
+}

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/3866ae40efed23181bf612efca4329d5348debc12468bf7565a2f13889b73bea.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@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 = i32(b0[i / 2048u]);
+if (token < 65536 || token >= 98304) { out[i] = f32(0.0); return; }
+out[i] = f32(b1[u32(token - 65536) * 2048u + i % 2048u]);
+}

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/3b495259612ab4c2e36bc153ec19119e95ab0d425693674839363ee4339a2770.wgsl

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

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/411153b2ab28bcb2ad154d237f976c034ae79a25f131bd899331278b53a0e15f.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 * 4194240u;
+if (i >= 16777216u) { return; }
+out[i] = f32(f32(b0[i]) * 0.125);
+}

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/463d2b26b08953bffa43b37837cf818d5e3af8498bf46b089e0b8e05d068c799.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 * 4194240u;
+if (i >= 16777216u) { return; }
+out[i] = f32(b0[i]);
+}

exaone4-fp32-8k-state-alias-token-major-v2/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-8k-state-alias-token-major-v2/kernels/4c3a20bb702b43bacecf9c7df989b55517469eba683bb2cb596e1f7a227347c7.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@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 token = i32(b0[i / 2048u]);
+if (token < 32768 || token >= 65536) { out[i] = f32(0.0); return; }
+out[i] = f32(b1[u32(token - 32768) * 2048u + i % 2048u]);
+}

exaone4-fp32-8k-state-alias-token-major-v2/kernels/4c526eba157bfd1aee20a2637586d6f4cdcdc68dcf238dcef84647649ffc6f9e.wgsl

@@ -0,0 +1,13 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<i32>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@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 token = i32(b0[i / 2048u]);
+if (token < 98304 || token >= 102400) { out[i] = f32(0.0); return; }
+out[i] = f32(b1[u32(token - 98304) * 2048u + i % 2048u]);
+}