Add exaone4-fp16-int4-g32-home-token-major-v2 WebTorch bundle

exaone4-fp16-int4-g32-home-token-major-v2/kernels/00dc173132a14e7104761434f61a56f267c2acb395ce97ef9d716a0a60c250f7.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/011b6264d946045ea160373e0d10175929e27ca18caa25582bc6a7ed5ecad4bf.wgsl

@@ -0,0 +1,18 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+@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 batch = i / 256u;
+let row = (i / 64u) % 4u;
+let col = i % 64u;
+var acc = 0.0;
+for (var p = 0u; p < 4096u; p++) {
+    acc += f32(b0[(((batch / 1u) % 32u) * 16384u) + row * 4096u + p]) * f32(b1[(((batch / 1u) % 32u) * 262144u) + p * 64u + col]);
+}
+out[i] = f16(acc);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/0166b41608be1a215345f1d90cb44a551dc8a191843a2804d1183a1d9719d84f.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[((i / 1024u) % 1u) * 2048u + ((i / 32u) % 32u) * 64u + ((i / 32u) % 1u) * 64u + (((i / 1u) % 32u) * 1u + 0u) * 1u]);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/019e4fd7e957ed2b0b1824813f9256bb8a1d6800c8dd705ee8a6458066500e1b.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/023d7b2d9e71a92f74db0e4583d453118915f46ba0c8c96c796632d958654ea1.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 >= 8u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(inverseSqrt(x));
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/03d3272d16bc278fcc977e95fcb512bac183f503c7d6ced229f9f01f8e5011bf.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<f16>;
+
+@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 >= 32768) { out[i] = f16(0.0); return; }
+out[i] = f16(b1[u32(token - 0) * 2048u + i % 2048u]);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/060d133d7c3df3be563af3249da9f84a6a2b31bfa4fa4a555f17dd524dd91e80.wgsl

@@ -0,0 +1,12 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+@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(f16(x / (1.0 + exp(-x))));
+out[i] = f16(s * f32(b1[i]));
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/065cbd92c6d555684b08098dc44c583dac58a722ef4c45f6787b6876994bc2b2.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[((i / 2048u) % 1u) * 2048u + ((i / 256u) % 8u) * 64u + ((i / 64u) % 4u) * 512u + ((i / 1u) % 64u) * 1u]);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/06e77154e0ee6052d86084b89d9d52dc46b08eba9a1c9eb0ad46f0bbb21b5e38.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/078c19606137dad9d459e8041919bac37c31e51702a6d466d4d3df095a6d1f5a.wgsl

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

exaone4-fp16-int4-g32-home-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-fp16-int4-g32-home-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-fp16-int4-g32-home-token-major-v2/kernels/0fb8222de35c84fb88807d4a3880f21188cf48d64c89a21aa208d7d72a48c273.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[i]);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/1589b6698aa9d57278b1f19482a7091fe7ecd424479d1f3adf7e7296a1c57957.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/16757bdfacc94a92cfc58c9a3ba83ca377b8ea7ddf82fa8f6611f797a6934f14.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<f16>;
+
+@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] = f16(0.0); return; }
+out[i] = f16(b1[u32(token - 32768) * 2048u + i % 2048u]);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/16de135b328f53749dc017ed0a624ea63f39df86947f72a0d5d4aa06bedeadd1.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/177e2dbd02302cd1c6265992d4a1b1c9276880d7556fa0d95865ca1bf7b2bf6f.wgsl

@@ -0,0 +1,11 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+@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] = f16(f32(b0[i]) + (f32(b1[((i / 1u) % 4096u) * 1u]) * 1.0));
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/1a1f8cbcc944c27cf3535d5efee191677cbd1b4e620cc906f71a4d6521dfc788.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) % 32u) * 1u]));
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/1df96eddb194a002a44f5a610245f84d192a38d2fff876c30fd55a4f518e8b25.wgsl

@@ -0,0 +1,11 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(-x);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/26655d5856c1c506477f9af344dc55197aa9220fe6e447eb38b11add83e278f3.wgsl

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

exaone4-fp16-int4-g32-home-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-fp16-int4-g32-home-token-major-v2/kernels/283d30fe187556e3ef23103af43381241093e52a62462461bcd16625a1c802bd.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/2a75f53dd6879f1759de2bcb4a2e141ca034c960393bc1107c56554ff3c60654.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[((i / 2048u) % 1u) * 2048u + ((i / 64u) % 32u) * 64u + ((i / 64u) % 1u) * 2048u + ((i / 1u) % 64u) * 1u]);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/2b707696a2df65d49bff973579e864e32bec86f16b559093c983e36ccae31a4a.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/2db8a995aaf1e89bf2da258d8e77c4ccc53a4d61ab9b352757da0e5a7321729d.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 * 2048u;
+if (i >= 2048u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(x*x);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/2fc9bfdf3b744d6a236488c5990f6534e10e652dfc6bc149d4587af5debb72d3.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/30c75242190dc245eaae002b6b370f5071b9091a315a677b837113f7460bfa77.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<f16>;
+
+@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 = i32(b0[i / 2048u]);
+if (token < 98304 || token >= 102400) { out[i] = f16(0.0); return; }
+out[i] = f16(b1[u32(token - 98304) * 2048u + i % 2048u]);
+}

exaone4-fp16-int4-g32-home-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-fp16-int4-g32-home-token-major-v2/kernels/344153d8e20e545c4656ce08b7236cd65878f5ea91014447f4a129ea94025970.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/370bab74beead2c5e9cc8893585f79ae97ee5eace554cee12555a8ac8293b1c0.wgsl

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

exaone4-fp16-int4-g32-home-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-fp16-int4-g32-home-token-major-v2/kernels/3976659e4383cb6f1c4f9c58ccef28215f246d506dae2455e18f51a78217549d.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/3be2a78373953009b427ed157b12f73e61fd2d2586572ea03746c0de9d31d583.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/3fe8cec0f9e8b563387f62876982a8f87cb5f03a1030cedde9cdb8e88c52e100.wgsl

@@ -0,0 +1,21 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+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] = f16(partial[0] + 0.0); }
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/40a1662d6cc89b15b300b62d9ca952abd0253b13a114fd1a5f942e5161a6d76c.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/447374683b0dcaacc17edb104c3ec2c932850855bb4f2750f67720678dd26326.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/44a50c95f6df65dd90f18b3425789a3b213a25a0e3c80080df677f2972aa63e9.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<f16>;
+
+@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 = i32(b0[i / 2048u]);
+if (token < 0 || token >= 32768) { out[i] = f16(0.0); return; }
+out[i] = f16(b1[u32(token - 0) * 2048u + i % 2048u]);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/48f34a333ff70495892101c428694fcddedec29f954b659093d206abbe1f7c15.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 >= 1u) { return; }
+let x = f32(b0[i]);
+out[i] = f32(inverseSqrt(x));
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/4cd438e7ca1c026269bc19fdf9fb0fdf0041d106ab97f168b3ee259159689b08.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[0u]));
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/4f5d44c78c9a1b2c2869666fb8ff4f626d6cbd9eb108e0b137980dbc2dc0c37b.wgsl

@@ -0,0 +1,11 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(-x);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/52ba9bc1eac03145df4ec09893e562a9c2a42758c6f6b0ec500a4acc9a4cb175.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/55ade971d1552c05370625247f20c24bab2e8ea7efde0b46cd9763527ebb6458.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/56ff2181d5d1a5ffa3b807eabdc3e5dac40edeb0d289824602e7d8ae63a58757.wgsl

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

exaone4-fp16-int4-g32-home-token-major-v2/kernels/578ba670650363664652f61f499af0f982bea5e68d2448225f4311b76bc6754b.wgsl

@@ -0,0 +1,24 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read> b1: array<f16>;
+@group(0) @binding(2) var<storage, read_write> out: array<f16>;
+
+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) * 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 < 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 < 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 < 4u && col < 32768u) { out[i] = f16(acc + 0.0); }
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/5fcece03bcdac8f6126ca4a1d087afd8b57ed402d87906641e15496d8cfcad21.wgsl

@@ -0,0 +1,10 @@
+enable f16;
+@group(0) @binding(0) var<storage, read> b0: array<f16>;
+@group(0) @binding(1) var<storage, read_write> out: array<f16>;
+
+@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] = f16(b0[((i / 1024u) % 1u) * 2048u + ((i / 32u) % 32u) * 64u + ((i / 32u) % 1u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]);
+}

exaone4-fp16-int4-g32-home-token-major-v2/kernels/5fd6e068841f5c054392bef564909dbf086bae9e4030f3625d1143244a08cbbf.wgsl

@@ -0,0 +1,11 @@
+@group(0) @binding(0) var<storage, read> b0: array<f32>;
+@group(0) @binding(1) var<storage, read_write> out: array<f32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+let i = gid.x + gid.y * 64u;
+if (i >= 64u) { 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-fp16-int4-g32-home-token-major-v2/kernels/6076d0bf36bd399e99fc409a95634cc3a508bec06ca558f0177434547983cd8b.wgsl

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