Publish WebTorch model catalog bundles (part 4)
Browse filesThis view is limited to 50 files because it contains too many changes. See raw diff
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a.wgsl +9 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/3203e08f2c99fa4848f553739b7115a7c677840344f17b1d4391d6130517f5a7.wgsl +9 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/3241a118927e109c0dacbb278370abe70d099ee593a07cd023d4530b5fb8f833.wgsl +24 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/33309cf63b78ee01282be4443e097b3ecfc78e5397b07229c915578f1ab5ebfb.wgsl +17 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/33e8ec04c66adb4988f562f0f9d3e92d29dfda91a56d0ed62be598c7c737f5ba.wgsl +9 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/358e73d0a8169937dce5d2c7fd5f4dea41f7c8307696f1698967bc6c4132efc8.wgsl +13 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/3612e5e8583c88858acf8b9868817b59e590adc115964205d24b5ce23659916d.wgsl +31 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/37ce78f0b50bc59dc7b6dd62724d6ad3acfbfb73deae12446a4575f5688d8f34.wgsl +9 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/383d0ce2bbe657b6412b8e55e6b0cf06e385cb4efbe45019afff9357f2a77751.wgsl +11 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/3866ae40efed23181bf612efca4329d5348debc12468bf7565a2f13889b73bea.wgsl +13 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/387ae1bba79da2cd17c9301d5f5afbba6597939d80cfc4fc4cbca9bc09cba4aa.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/39646bdfafd1ab0cb874474e3f1e9d79b57f10582a1804bc1d3762a492fec851.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/3b495259612ab4c2e36bc153ec19119e95ab0d425693674839363ee4339a2770.wgsl +21 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/3e34409407100d4858c4c109cf62f3eaf2f4c6081166b0ec8679105e6dfbe7fe.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/3fa622a834e456fdb19b07df0a199659495aef231294b3bf3123a067c611f13a.wgsl +31 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/40f115c8647a238bfb3ccb0bb8446bccdc94a0dfadc89111c4920efb2259f41d.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/41762d906bc04bda322951d466c6ff5f628c5a37db58871ee2f60e6e7310b368.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/431a95359c9c41766af53351f3074089c052c44df79b9d1d08205d8e58e90d25.wgsl +13 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/453668cbe741eb27e2cc10194eca66f514a3763d97b46c1733c376986defc322.wgsl +24 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/46da6fd96b23846a4f194f7ee18357aabf58002ce3cc430fe2508b485dcd5a82.wgsl +28 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/483857d2d9929bc07cd14940dc3bca97abcf9485ce1c805a234d2db4110a7010.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/4c3a20bb702b43bacecf9c7df989b55517469eba683bb2cb596e1f7a227347c7.wgsl +13 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/4c526eba157bfd1aee20a2637586d6f4cdcdc68dcf238dcef84647649ffc6f9e.wgsl +13 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/4ec7950cea0e4257dc30630582df36a6127ce61794f46012db6ad0c0b1ea8380.wgsl +24 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/4fad29e3e337840e25132f10b3e3d55225f727ae13525ff0b0bdf31e03f09d74.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/51af37af0ad3d33571d8dd0d261e73f584605ebcd12174b25551438d190eb7d7.wgsl +13 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/53202349e6a251956a220a83f711afd301d6409ac3f8280ea615e7c2924073ba.wgsl +9 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/5453b6b404f31128f8f078d0fd1d167369d60314eeef96913cbf43c1497c9ba1.wgsl +11 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/5840702a9664fc244e8ea56899c5ebc5410b69fba0cab2184aa575635398be18.wgsl +34 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/59c2a7b048981784fb950a5ca345ace028307a715923c76eec6886b4e7d4ef77.wgsl +16 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/5a213200cd8d0aed4fa53959ff55af996fc5e841ab19e8966edf791aa8589ece.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/5b38fbf5d669e747694382155d845113c7f75356bea9c989c685769dd8514a7e.wgsl +17 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/5c22e6f60217881c0351e38d19057626a6c2292771c1a83816995245777ac2fd.wgsl +9 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/5c7660829aa2cc487192227f88305fc4fca5736388eb82bffbc7b4bbde33624f.wgsl +17 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/5dbfeacd0dff900672cf198eb01bda26237de50a661639c3aa24bfde0bff09b5.wgsl +9 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/5f4403d78557e88e66d71eb5e53746cc07449a7c421b57eda4f093a09a68a99a.wgsl +9 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/5fd6e068841f5c054392bef564909dbf086bae9e4030f3625d1143244a08cbbf.wgsl +11 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/605f51996eef0769d3b456a83f7540262d1761589b55d1fd41fb2b4dd8fa249f.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/61f79c3b6844829354ffb9a676195191ed203edd062b78faa8424f758bc17a30.wgsl +9 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/66978c862f8d1e391ded2424a4132cfea5f0cc7686c76610502bb596224db368.wgsl +28 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/6726b52fb13480f5670b2dd4c37aaa8b3d39b0705cc47654f117c5781793aad5.wgsl +9 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/6968677e60103e0726f3c8a45bd76472ec8790313e1eb43b142500860f098d72.wgsl +13 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/696b0d070511e345cddeb934dc99fc629d0ec7e90ee55799a41d6726d602eb5b.wgsl +24 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/6e99c56e3571c8da5671f66929f70a2fd55b3f2fb17f405008ff7195e0900728.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/6f49fdd6312fb3bfeef8d2bbb2ee3bd9327b5def9064b531b7cb5413524479b7.wgsl +31 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/7002f8fdf09f5124ce2875c3540bce3e77750260556fc1a00a8b13b6c1871c30.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/70dd45c48a989e6617431d39bc63f0187859f791c0915bcfc95fd6ebf4153403.wgsl +31 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/73eead1530c6564c2c34dd4b18bfa56e96844abc065ac25ad4bed1eeb1f4402c.wgsl +9 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/764403c33cd8994568e068810380cb57e9f2e326dfc09bed63eed1fbed044da1.wgsl +10 -0
- exaone4-fp32-int8-g32-release-token-major-v2/kernels/7ab61f48fc60ed8197fe395932ca6d1b90869c9d307fb397eb8f6a26811049e7.wgsl +13 -0
exaone4-fp32-int8-g32-release-token-major-v2/kernels/2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a.wgsl
ADDED
|
@@ -0,0 +1,9 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 64u;
|
| 7 |
+
if (i >= 64u) { return; }
|
| 8 |
+
out[i] = f32(f32(b0[i]) * 1.0);
|
| 9 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/3203e08f2c99fa4848f553739b7115a7c677840344f17b1d4391d6130517f5a7.wgsl
ADDED
|
@@ -0,0 +1,9 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<i32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<i32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 64u;
|
| 7 |
+
if (i >= 4u) { return; }
|
| 8 |
+
out[i] = i32(b0[i]);
|
| 9 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/3241a118927e109c0dacbb278370abe70d099ee593a07cd023d4530b5fb8f833.wgsl
ADDED
|
@@ -0,0 +1,24 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
var<workgroup> factor: f32;
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 8 |
+
let row = group.x + group.y * 8u;
|
| 9 |
+
if (row >= 8u) { return; }
|
| 10 |
+
let lane = local.x;
|
| 11 |
+
if (lane == 0u) {
|
| 12 |
+
var total = 0.0;
|
| 13 |
+
for (var j = 0u; j < 64u; j++) {
|
| 14 |
+
let v = f32(b0[row * 64u + j]);
|
| 15 |
+
total += v * v;
|
| 16 |
+
}
|
| 17 |
+
factor = inverseSqrt(total / 64.0 + 1e-05);
|
| 18 |
+
}
|
| 19 |
+
workgroupBarrier();
|
| 20 |
+
for (var p = lane; p < 64u; p += 64u) {
|
| 21 |
+
let i = row * 64u + p;
|
| 22 |
+
out[i] = f32(f32(b0[row * 64u + p]) * factor * (f32(b1[p]) + 0.0));
|
| 23 |
+
}
|
| 24 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/33309cf63b78ee01282be4443e097b3ecfc78e5397b07229c915578f1ab5ebfb.wgsl
ADDED
|
@@ -0,0 +1,17 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 8192u;
|
| 8 |
+
if (i >= 8192u) { return; }
|
| 9 |
+
let batch = i / 256u;
|
| 10 |
+
let row = (i / 64u) % 4u;
|
| 11 |
+
let col = i % 64u;
|
| 12 |
+
var acc = 0.0;
|
| 13 |
+
for (var p = 0u; p < 4096u; p++) {
|
| 14 |
+
acc += f32(b0[(((batch / 1u) % 32u) * 16384u) + row * 4096u + p]) * f32(b1[(((batch / 1u) % 32u) * 262144u) + p * 64u + col]);
|
| 15 |
+
}
|
| 16 |
+
out[i] = f32(acc);
|
| 17 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/33e8ec04c66adb4988f562f0f9d3e92d29dfda91a56d0ed62be598c7c737f5ba.wgsl
ADDED
|
@@ -0,0 +1,9 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 16384u;
|
| 7 |
+
if (i >= 16384u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 16384u) % 1u) * 32768u + ((i / 2048u) % 8u) * 4096u + ((i / 32u) % 64u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]);
|
| 9 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/358e73d0a8169937dce5d2c7fd5f4dea41f7c8307696f1698967bc6c4132efc8.wgsl
ADDED
|
@@ -0,0 +1,13 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
enable f16;
|
| 2 |
+
@group(0) @binding(0) var<storage, read> b0: array<i32>;
|
| 3 |
+
@group(0) @binding(1) var<storage, read> b1: array<f16>;
|
| 4 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 5 |
+
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 8 |
+
let i = gid.x + gid.y * 8192u;
|
| 9 |
+
if (i >= 8192u) { return; }
|
| 10 |
+
let token = i32(b0[i / 2048u]);
|
| 11 |
+
if (token < 32768 || token >= 65536) { out[i] = f32(0.0); return; }
|
| 12 |
+
out[i] = f32(b1[u32(token - 32768) * 2048u + i % 2048u]);
|
| 13 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/3612e5e8583c88858acf8b9868817b59e590adc115964205d24b5ce23659916d.wgsl
ADDED
|
@@ -0,0 +1,31 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<u32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read> b2: array<f32>;
|
| 4 |
+
@group(0) @binding(3) var<storage, read_write> out: array<f32>;
|
| 5 |
+
fn dequant_1(index: u32) -> f32 {
|
| 6 |
+
let row = index / 2048u;
|
| 7 |
+
let col = index % 2048u;
|
| 8 |
+
let word = b1[row * 512u + col / 4u];
|
| 9 |
+
let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
|
| 10 |
+
return f32(code) * b2[row * 64u + col / 32u];
|
| 11 |
+
}
|
| 12 |
+
|
| 13 |
+
var<workgroup> tile_a: array<f32, 64>;
|
| 14 |
+
var<workgroup> tile_b: array<f32, 64>;
|
| 15 |
+
@compute @workgroup_size(8, 8)
|
| 16 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 17 |
+
let row = group.y * 8u + local.y;
|
| 18 |
+
let col = group.x * 8u + local.x;
|
| 19 |
+
let i = (group.z * 16u + row) * 4096u + col;
|
| 20 |
+
var acc = 0.0;
|
| 21 |
+
for (var tile = 0u; tile < 2048u; tile += 8u) {
|
| 22 |
+
tile_a[local.y * 8u + local.x] = 0.0;
|
| 23 |
+
tile_b[local.x * 8u + local.y] = 0.0;
|
| 24 |
+
if (row < 16u && tile + local.x < 2048u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 16u + row) * 2048u + tile + local.x]); }
|
| 25 |
+
if (group.x * 8u + local.y < 4096u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = dequant_1((group.x * 8u + local.y) * 2048u + tile + local.x); }
|
| 26 |
+
workgroupBarrier();
|
| 27 |
+
for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
|
| 28 |
+
workgroupBarrier();
|
| 29 |
+
}
|
| 30 |
+
if (row < 16u && col < 4096u) { out[i] = f32(acc + 0.0); }
|
| 31 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/37ce78f0b50bc59dc7b6dd62724d6ad3acfbfb73deae12446a4575f5688d8f34.wgsl
ADDED
|
@@ -0,0 +1,9 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 128u;
|
| 7 |
+
if (i >= 128u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 128u) % 1u) * 128u + ((i / 32u) % 4u) * 1u + ((i / 1u) % 32u) * 4u]);
|
| 9 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/383d0ce2bbe657b6412b8e55e6b0cf06e385cb4efbe45019afff9357f2a77751.wgsl
ADDED
|
@@ -0,0 +1,11 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 16384u;
|
| 8 |
+
if (i >= 16384u) { return; }
|
| 9 |
+
let x = f32(b0[i]); let s = f32(f32(x / (1.0 + exp(-x))));
|
| 10 |
+
out[i] = f32(s * f32(b1[i]));
|
| 11 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/3866ae40efed23181bf612efca4329d5348debc12468bf7565a2f13889b73bea.wgsl
ADDED
|
@@ -0,0 +1,13 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
enable f16;
|
| 2 |
+
@group(0) @binding(0) var<storage, read> b0: array<i32>;
|
| 3 |
+
@group(0) @binding(1) var<storage, read> b1: array<f16>;
|
| 4 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 5 |
+
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 8 |
+
let i = gid.x + gid.y * 8192u;
|
| 9 |
+
if (i >= 8192u) { return; }
|
| 10 |
+
let token = i32(b0[i / 2048u]);
|
| 11 |
+
if (token < 65536 || token >= 98304) { out[i] = f32(0.0); return; }
|
| 12 |
+
out[i] = f32(b1[u32(token - 65536) * 2048u + i % 2048u]);
|
| 13 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/387ae1bba79da2cd17c9301d5f5afbba6597939d80cfc4fc4cbca9bc09cba4aa.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 4194240u;
|
| 8 |
+
if (i >= 8388608u) { return; }
|
| 9 |
+
out[i] = f32(f32(b0[i]) + (f32(b1[((i / 4096u) % 64u) * 4096u + ((i / 1u) % 4096u) * 1u]) * 1.0));
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/39646bdfafd1ab0cb874474e3f1e9d79b57f10582a1804bc1d3762a492fec851.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 32768u;
|
| 8 |
+
if (i >= 32768u) { return; }
|
| 9 |
+
out[i] = f32(f32(b0[i]) * f32(b1[((i / 64u) % 64u) * 64u + ((i / 1u) % 64u) * 1u]));
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/3b495259612ab4c2e36bc153ec19119e95ab0d425693674839363ee4339a2770.wgsl
ADDED
|
@@ -0,0 +1,21 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
enable f16;
|
| 2 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 3 |
+
@group(0) @binding(1) var<storage, read> b1: array<f16>;
|
| 4 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 5 |
+
|
| 6 |
+
var<workgroup> partial: array<f32, 64>;
|
| 7 |
+
@compute @workgroup_size(64)
|
| 8 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 9 |
+
let i = group.x + group.y * 4096u;
|
| 10 |
+
if (i >= 4096u) { return; }
|
| 11 |
+
let lane = local.x;
|
| 12 |
+
var acc = 0.0;
|
| 13 |
+
for (var p = lane; p < 2048u; p += 64u) { acc += f32(b0[(i / 4096u) * 2048u + p]) * f32(b1[(i % 4096u) * 2048u + p]); }
|
| 14 |
+
partial[lane] = acc;
|
| 15 |
+
workgroupBarrier();
|
| 16 |
+
for (var stride = 32u; stride > 0u; stride /= 2u) {
|
| 17 |
+
if (lane < stride) { partial[lane] += partial[lane + stride]; }
|
| 18 |
+
workgroupBarrier();
|
| 19 |
+
}
|
| 20 |
+
if (lane == 0u) { out[i] = f32(partial[0] + 0.0); }
|
| 21 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/3e34409407100d4858c4c109cf62f3eaf2f4c6081166b0ec8679105e6dfbe7fe.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 16384u;
|
| 7 |
+
if (i >= 16384u) { return; }
|
| 8 |
+
let x = f32(b0[i]);
|
| 9 |
+
out[i] = f32(-x);
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/3fa622a834e456fdb19b07df0a199659495aef231294b3bf3123a067c611f13a.wgsl
ADDED
|
@@ -0,0 +1,31 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<u32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read> b2: array<f32>;
|
| 4 |
+
@group(0) @binding(3) var<storage, read_write> out: array<f32>;
|
| 5 |
+
fn dequant_1(index: u32) -> f32 {
|
| 6 |
+
let row = index / 2048u;
|
| 7 |
+
let col = index % 2048u;
|
| 8 |
+
let word = b1[row * 512u + col / 4u];
|
| 9 |
+
let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
|
| 10 |
+
return f32(code) * b2[row * 64u + col / 32u];
|
| 11 |
+
}
|
| 12 |
+
|
| 13 |
+
var<workgroup> tile_a: array<f32, 64>;
|
| 14 |
+
var<workgroup> tile_b: array<f32, 64>;
|
| 15 |
+
@compute @workgroup_size(8, 8)
|
| 16 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 17 |
+
let row = group.y * 8u + local.y;
|
| 18 |
+
let col = group.x * 8u + local.x;
|
| 19 |
+
let i = (group.z * 64u + row) * 4096u + col;
|
| 20 |
+
var acc = 0.0;
|
| 21 |
+
for (var tile = 0u; tile < 2048u; tile += 8u) {
|
| 22 |
+
tile_a[local.y * 8u + local.x] = 0.0;
|
| 23 |
+
tile_b[local.x * 8u + local.y] = 0.0;
|
| 24 |
+
if (row < 64u && tile + local.x < 2048u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 64u + row) * 2048u + tile + local.x]); }
|
| 25 |
+
if (group.x * 8u + local.y < 4096u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = dequant_1((group.x * 8u + local.y) * 2048u + tile + local.x); }
|
| 26 |
+
workgroupBarrier();
|
| 27 |
+
for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
|
| 28 |
+
workgroupBarrier();
|
| 29 |
+
}
|
| 30 |
+
if (row < 64u && col < 4096u) { out[i] = f32(acc + 0.0); }
|
| 31 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/40f115c8647a238bfb3ccb0bb8446bccdc94a0dfadc89111c4920efb2259f41d.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 2048u;
|
| 8 |
+
if (i >= 2048u) { return; }
|
| 9 |
+
out[i] = f32(f32(b0[i]) * f32(b1[((i / 64u) % 4u) * 64u + ((i / 1u) % 64u) * 1u]));
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/41762d906bc04bda322951d466c6ff5f628c5a37db58871ee2f60e6e7310b368.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 524288u;
|
| 8 |
+
if (i >= 524288u) { return; }
|
| 9 |
+
out[i] = f32(f32(b0[i]) + (f32(b1[((i / 4096u) % 4u) * 4096u + ((i / 1u) % 4096u) * 1u]) * 1.0));
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/431a95359c9c41766af53351f3074089c052c44df79b9d1d08205d8e58e90d25.wgsl
ADDED
|
@@ -0,0 +1,13 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<i32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 32768u;
|
| 8 |
+
if (i >= 32768u) { return; }
|
| 9 |
+
let token = (i / 64u) % 64u;
|
| 10 |
+
let outer = i / 4096u;
|
| 11 |
+
let destination = outer * 262144u + u32(b1[token]) * 64u + i % 64u;
|
| 12 |
+
out[((destination) / 64u) % 4096u * 512u + ((destination) / 262144u) % 8u * 64u + ((destination) / 1u) % 64u * 1u] = f32(b0[i]);
|
| 13 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/453668cbe741eb27e2cc10194eca66f514a3763d97b46c1733c376986defc322.wgsl
ADDED
|
@@ -0,0 +1,24 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
var<workgroup> factor: f32;
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 8 |
+
let row = group.x + group.y * 4u;
|
| 9 |
+
if (row >= 4u) { return; }
|
| 10 |
+
let lane = local.x;
|
| 11 |
+
if (lane == 0u) {
|
| 12 |
+
var total = 0.0;
|
| 13 |
+
for (var j = 0u; j < 2048u; j++) {
|
| 14 |
+
let v = f32(b0[row * 2048u + j]);
|
| 15 |
+
total += v * v;
|
| 16 |
+
}
|
| 17 |
+
factor = inverseSqrt(total / 2048.0 + 1e-05);
|
| 18 |
+
}
|
| 19 |
+
workgroupBarrier();
|
| 20 |
+
for (var p = lane; p < 2048u; p += 64u) {
|
| 21 |
+
let i = row * 2048u + p;
|
| 22 |
+
out[i] = f32(f32(b0[row * 2048u + p]) * factor * (f32(b1[p]) + 0.0));
|
| 23 |
+
}
|
| 24 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/46da6fd96b23846a4f194f7ee18357aabf58002ce3cc430fe2508b485dcd5a82.wgsl
ADDED
|
@@ -0,0 +1,28 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<u32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read> b2: array<f32>;
|
| 4 |
+
@group(0) @binding(3) var<storage, read_write> out: array<f32>;
|
| 5 |
+
fn dequant_1(index: u32) -> f32 {
|
| 6 |
+
let row = index / 4096u;
|
| 7 |
+
let col = index % 4096u;
|
| 8 |
+
let word = b1[row * 1024u + col / 4u];
|
| 9 |
+
let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
|
| 10 |
+
return f32(code) * b2[row * 128u + col / 32u];
|
| 11 |
+
}
|
| 12 |
+
|
| 13 |
+
var<workgroup> partial: array<f32, 64>;
|
| 14 |
+
@compute @workgroup_size(64)
|
| 15 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 16 |
+
let i = group.x + group.y * 2048u;
|
| 17 |
+
if (i >= 2048u) { return; }
|
| 18 |
+
let lane = local.x;
|
| 19 |
+
var acc = 0.0;
|
| 20 |
+
for (var p = lane; p < 4096u; p += 64u) { acc += f32(b0[(i / 2048u) * 4096u + p]) * dequant_1((i % 2048u) * 4096u + p); }
|
| 21 |
+
partial[lane] = acc;
|
| 22 |
+
workgroupBarrier();
|
| 23 |
+
for (var stride = 32u; stride > 0u; stride /= 2u) {
|
| 24 |
+
if (lane < stride) { partial[lane] += partial[lane + stride]; }
|
| 25 |
+
workgroupBarrier();
|
| 26 |
+
}
|
| 27 |
+
if (lane == 0u) { out[i] = f32(partial[0] + 0.0); }
|
| 28 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/483857d2d9929bc07cd14940dc3bca97abcf9485ce1c805a234d2db4110a7010.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 32768u;
|
| 8 |
+
if (i >= 32768u) { return; }
|
| 9 |
+
out[i] = f32(f32(b0[i]) * f32(b1[((i / 64u) % 16u) * 64u + ((i / 1u) % 64u) * 1u]));
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/4c3a20bb702b43bacecf9c7df989b55517469eba683bb2cb596e1f7a227347c7.wgsl
ADDED
|
@@ -0,0 +1,13 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
enable f16;
|
| 2 |
+
@group(0) @binding(0) var<storage, read> b0: array<i32>;
|
| 3 |
+
@group(0) @binding(1) var<storage, read> b1: array<f16>;
|
| 4 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 5 |
+
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 8 |
+
let i = gid.x + gid.y * 131072u;
|
| 9 |
+
if (i >= 131072u) { return; }
|
| 10 |
+
let token = i32(b0[i / 2048u]);
|
| 11 |
+
if (token < 32768 || token >= 65536) { out[i] = f32(0.0); return; }
|
| 12 |
+
out[i] = f32(b1[u32(token - 32768) * 2048u + i % 2048u]);
|
| 13 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/4c526eba157bfd1aee20a2637586d6f4cdcdc68dcf238dcef84647649ffc6f9e.wgsl
ADDED
|
@@ -0,0 +1,13 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
enable f16;
|
| 2 |
+
@group(0) @binding(0) var<storage, read> b0: array<i32>;
|
| 3 |
+
@group(0) @binding(1) var<storage, read> b1: array<f16>;
|
| 4 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 5 |
+
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 8 |
+
let i = gid.x + gid.y * 131072u;
|
| 9 |
+
if (i >= 131072u) { return; }
|
| 10 |
+
let token = i32(b0[i / 2048u]);
|
| 11 |
+
if (token < 98304 || token >= 102400) { out[i] = f32(0.0); return; }
|
| 12 |
+
out[i] = f32(b1[u32(token - 98304) * 2048u + i % 2048u]);
|
| 13 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/4ec7950cea0e4257dc30630582df36a6127ce61794f46012db6ad0c0b1ea8380.wgsl
ADDED
|
@@ -0,0 +1,24 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
var<workgroup> factor: f32;
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 8 |
+
let row = group.x + group.y * 64u;
|
| 9 |
+
if (row >= 64u) { return; }
|
| 10 |
+
let lane = local.x;
|
| 11 |
+
if (lane == 0u) {
|
| 12 |
+
var total = 0.0;
|
| 13 |
+
for (var j = 0u; j < 2048u; j++) {
|
| 14 |
+
let v = f32(b0[row * 2048u + j]);
|
| 15 |
+
total += v * v;
|
| 16 |
+
}
|
| 17 |
+
factor = inverseSqrt(total / 2048.0 + 1e-05);
|
| 18 |
+
}
|
| 19 |
+
workgroupBarrier();
|
| 20 |
+
for (var p = lane; p < 2048u; p += 64u) {
|
| 21 |
+
let i = row * 2048u + p;
|
| 22 |
+
out[i] = f32(f32(b0[row * 2048u + p]) * factor * (f32(b1[p]) + 0.0));
|
| 23 |
+
}
|
| 24 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/4fad29e3e337840e25132f10b3e3d55225f727ae13525ff0b0bdf31e03f09d74.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 256u;
|
| 7 |
+
if (i >= 256u) { return; }
|
| 8 |
+
let x = f32(b0[i]);
|
| 9 |
+
out[i] = f32(-x);
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/51af37af0ad3d33571d8dd0d261e73f584605ebcd12174b25551438d190eb7d7.wgsl
ADDED
|
@@ -0,0 +1,13 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
enable f16;
|
| 2 |
+
@group(0) @binding(0) var<storage, read> b0: array<i32>;
|
| 3 |
+
@group(0) @binding(1) var<storage, read> b1: array<f16>;
|
| 4 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 5 |
+
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 8 |
+
let i = gid.x + gid.y * 2048u;
|
| 9 |
+
if (i >= 2048u) { return; }
|
| 10 |
+
let token = i32(b0[i / 2048u]);
|
| 11 |
+
if (token < 0 || token >= 32768) { out[i] = f32(0.0); return; }
|
| 12 |
+
out[i] = f32(b1[u32(token - 0) * 2048u + i % 2048u]);
|
| 13 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/53202349e6a251956a220a83f711afd301d6409ac3f8280ea615e7c2924073ba.wgsl
ADDED
|
@@ -0,0 +1,9 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 2097152u;
|
| 7 |
+
if (i >= 2097152u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i) / 64u) % 4096u * 512u + ((i) / 262144u) % 8u * 64u + ((i) / 1u) % 64u * 1u]);
|
| 9 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/5453b6b404f31128f8f078d0fd1d167369d60314eeef96913cbf43c1497c9ba1.wgsl
ADDED
|
@@ -0,0 +1,11 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 262144u;
|
| 8 |
+
if (i >= 262144u) { return; }
|
| 9 |
+
let x = f32(b0[i]); let s = f32(f32(x / (1.0 + exp(-x))));
|
| 10 |
+
out[i] = f32(s * f32(b1[i]));
|
| 11 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/5840702a9664fc244e8ea56899c5ebc5410b69fba0cab2184aa575635398be18.wgsl
ADDED
|
@@ -0,0 +1,34 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
var<workgroup> partial: array<f32, 64>;
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 7 |
+
let row = group.x + group.y * 2048u;
|
| 8 |
+
if (row >= 2048u) { return; }
|
| 9 |
+
let lane = local.x;
|
| 10 |
+
var maximum = -3.4028234663852886e38;
|
| 11 |
+
for (var p = lane; p < 4096u; p += 64u) { maximum = max(maximum, f32(b0[row * 4096u + p])); }
|
| 12 |
+
partial[lane] = maximum;
|
| 13 |
+
workgroupBarrier();
|
| 14 |
+
for (var stride = 32u; stride > 0u; stride /= 2u) {
|
| 15 |
+
if (lane < stride) { partial[lane] = max(partial[lane], partial[lane + stride]); }
|
| 16 |
+
workgroupBarrier();
|
| 17 |
+
}
|
| 18 |
+
maximum = partial[0];
|
| 19 |
+
// All lanes must finish reading the maximum before the scratch array is reused.
|
| 20 |
+
workgroupBarrier();
|
| 21 |
+
var total = 0.0;
|
| 22 |
+
for (var p = lane; p < 4096u; p += 64u) { total += exp(f32(b0[row * 4096u + p]) - maximum); }
|
| 23 |
+
partial[lane] = total;
|
| 24 |
+
workgroupBarrier();
|
| 25 |
+
for (var stride = 32u; stride > 0u; stride /= 2u) {
|
| 26 |
+
if (lane < stride) { partial[lane] += partial[lane + stride]; }
|
| 27 |
+
workgroupBarrier();
|
| 28 |
+
}
|
| 29 |
+
total = partial[0];
|
| 30 |
+
for (var p = lane; p < 4096u; p += 64u) {
|
| 31 |
+
let i = row * 4096u + p;
|
| 32 |
+
out[i] = f32(exp(f32(b0[row * 4096u + p]) - maximum) / total);
|
| 33 |
+
}
|
| 34 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/59c2a7b048981784fb950a5ca345ace028307a715923c76eec6886b4e7d4ef77.wgsl
ADDED
|
@@ -0,0 +1,16 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read> b2: array<f32>;
|
| 4 |
+
@group(0) @binding(3) var<storage, read> b3: array<f32>;
|
| 5 |
+
@group(0) @binding(4) var<storage, read_write> out: array<f32>;
|
| 6 |
+
|
| 7 |
+
@compute @workgroup_size(64)
|
| 8 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 9 |
+
let i = gid.x + gid.y * 4194240u;
|
| 10 |
+
if (i >= 6553600u) { return; }
|
| 11 |
+
let coord = (i / 1u) % 102400u;
|
| 12 |
+
if (coord >= 0u && coord < 32768u) { out[i] = f32(b0[(i / 102400u) * 32768u + (coord - 0u) * 1u + i % 1u]); }
|
| 13 |
+
if (coord >= 32768u && coord < 65536u) { out[i] = f32(b1[(i / 102400u) * 32768u + (coord - 32768u) * 1u + i % 1u]); }
|
| 14 |
+
if (coord >= 65536u && coord < 98304u) { out[i] = f32(b2[(i / 102400u) * 32768u + (coord - 65536u) * 1u + i % 1u]); }
|
| 15 |
+
if (coord >= 98304u && coord < 102400u) { out[i] = f32(b3[(i / 102400u) * 4096u + (coord - 98304u) * 1u + i % 1u]); }
|
| 16 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/5a213200cd8d0aed4fa53959ff55af996fc5e841ab19e8966edf791aa8589ece.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 131072u;
|
| 8 |
+
if (i >= 131072u) { return; }
|
| 9 |
+
out[i] = f32(f32(b0[i]) + (f32(b1[((i / 1u) % 4096u) * 1u]) * 1.0));
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/5b38fbf5d669e747694382155d845113c7f75356bea9c989c685769dd8514a7e.wgsl
ADDED
|
@@ -0,0 +1,17 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 524288u;
|
| 8 |
+
if (i >= 524288u) { return; }
|
| 9 |
+
let batch = i / 16384u;
|
| 10 |
+
let row = (i / 4096u) % 4u;
|
| 11 |
+
let col = i % 4096u;
|
| 12 |
+
var acc = 0.0;
|
| 13 |
+
for (var p = 0u; p < 64u; p++) {
|
| 14 |
+
acc += f32(b0[(((batch / 1u) % 32u) * 256u) + row * 64u + p]) * f32(b1[(((batch / 1u) % 32u) * 262144u) + p * 4096u + col]);
|
| 15 |
+
}
|
| 16 |
+
out[i] = f32(acc);
|
| 17 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/5c22e6f60217881c0351e38d19057626a6c2292771c1a83816995245777ac2fd.wgsl
ADDED
|
@@ -0,0 +1,9 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 256u;
|
| 7 |
+
if (i >= 256u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 256u) % 1u) * 512u + ((i / 32u) % 8u) * 64u + ((i / 32u) % 1u) * 64u + (((i / 1u) % 32u) * 1u + 0u) * 1u]);
|
| 9 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/5c7660829aa2cc487192227f88305fc4fca5736388eb82bffbc7b4bbde33624f.wgsl
ADDED
|
@@ -0,0 +1,17 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 4194240u;
|
| 8 |
+
if (i >= 8388608u) { return; }
|
| 9 |
+
let batch = i / 262144u;
|
| 10 |
+
let row = (i / 4096u) % 64u;
|
| 11 |
+
let col = i % 4096u;
|
| 12 |
+
var acc = 0.0;
|
| 13 |
+
for (var p = 0u; p < 64u; p++) {
|
| 14 |
+
acc += f32(b0[(((batch / 1u) % 32u) * 4096u) + row * 64u + p]) * f32(b1[(((batch / 1u) % 32u) * 262144u) + p * 4096u + col]);
|
| 15 |
+
}
|
| 16 |
+
out[i] = f32(acc);
|
| 17 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/5dbfeacd0dff900672cf198eb01bda26237de50a661639c3aa24bfde0bff09b5.wgsl
ADDED
|
@@ -0,0 +1,9 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 131072u;
|
| 7 |
+
if (i >= 131072u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 131072u) % 1u) * 131072u + ((i / 2048u) % 64u) * 64u + ((i / 64u) % 32u) * 4096u + ((i / 1u) % 64u) * 1u]);
|
| 9 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/5f4403d78557e88e66d71eb5e53746cc07449a7c421b57eda4f093a09a68a99a.wgsl
ADDED
|
@@ -0,0 +1,9 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 4096u;
|
| 7 |
+
if (i >= 4096u) { return; }
|
| 8 |
+
out[i] = f32(b0[i]);
|
| 9 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/5fd6e068841f5c054392bef564909dbf086bae9e4030f3625d1143244a08cbbf.wgsl
ADDED
|
@@ -0,0 +1,11 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 64u;
|
| 7 |
+
if (i >= 64u) { return; }
|
| 8 |
+
let coord = (i / 1u) % 64u;
|
| 9 |
+
if (coord >= 0u && coord < 32u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 0u) * 1u + i % 1u]); }
|
| 10 |
+
if (coord >= 32u && coord < 64u) { out[i] = f32(b0[(i / 64u) * 32u + (coord - 32u) * 1u + i % 1u]); }
|
| 11 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/605f51996eef0769d3b456a83f7540262d1761589b55d1fd41fb2b4dd8fa249f.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 8192u;
|
| 8 |
+
if (i >= 8192u) { return; }
|
| 9 |
+
out[i] = f32(f32(b0[i]) + (f32(b1[i]) * 1.0));
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/61f79c3b6844829354ffb9a676195191ed203edd062b78faa8424f758bc17a30.wgsl
ADDED
|
@@ -0,0 +1,9 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 8192u;
|
| 7 |
+
if (i >= 8192u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 8192u) % 1u) * 8192u + ((i / 256u) % 32u) * 64u + ((i / 64u) % 4u) * 2048u + ((i / 1u) % 64u) * 1u]);
|
| 9 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/66978c862f8d1e391ded2424a4132cfea5f0cc7686c76610502bb596224db368.wgsl
ADDED
|
@@ -0,0 +1,28 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<u32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read> b2: array<f32>;
|
| 4 |
+
@group(0) @binding(3) var<storage, read_write> out: array<f32>;
|
| 5 |
+
fn dequant_1(index: u32) -> f32 {
|
| 6 |
+
let row = index / 2048u;
|
| 7 |
+
let col = index % 2048u;
|
| 8 |
+
let word = b1[row * 512u + col / 4u];
|
| 9 |
+
let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
|
| 10 |
+
return f32(code) * b2[row * 64u + col / 32u];
|
| 11 |
+
}
|
| 12 |
+
|
| 13 |
+
var<workgroup> partial: array<f32, 64>;
|
| 14 |
+
@compute @workgroup_size(64)
|
| 15 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 16 |
+
let i = group.x + group.y * 4096u;
|
| 17 |
+
if (i >= 4096u) { return; }
|
| 18 |
+
let lane = local.x;
|
| 19 |
+
var acc = 0.0;
|
| 20 |
+
for (var p = lane; p < 2048u; p += 64u) { acc += f32(b0[(i / 4096u) * 2048u + p]) * dequant_1((i % 4096u) * 2048u + p); }
|
| 21 |
+
partial[lane] = acc;
|
| 22 |
+
workgroupBarrier();
|
| 23 |
+
for (var stride = 32u; stride > 0u; stride /= 2u) {
|
| 24 |
+
if (lane < stride) { partial[lane] += partial[lane + stride]; }
|
| 25 |
+
workgroupBarrier();
|
| 26 |
+
}
|
| 27 |
+
if (lane == 0u) { out[i] = f32(partial[0] + 0.0); }
|
| 28 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/6726b52fb13480f5670b2dd4c37aaa8b3d39b0705cc47654f117c5781793aad5.wgsl
ADDED
|
@@ -0,0 +1,9 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 4194240u;
|
| 7 |
+
if (i >= 8388608u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 8388608u) % 1u) * 8388608u + ((i / 262144u) % 32u) * 262144u + ((i / 4096u) % 64u) * 1u + ((i / 1u) % 4096u) * 64u]);
|
| 9 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/6968677e60103e0726f3c8a45bd76472ec8790313e1eb43b142500860f098d72.wgsl
ADDED
|
@@ -0,0 +1,13 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<i32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 512u;
|
| 8 |
+
if (i >= 512u) { return; }
|
| 9 |
+
let token = (i / 64u) % 1u;
|
| 10 |
+
let outer = i / 64u;
|
| 11 |
+
let destination = outer * 262144u + u32(b1[token]) * 64u + i % 64u;
|
| 12 |
+
out[((destination) / 64u) % 4096u * 512u + ((destination) / 262144u) % 8u * 64u + ((destination) / 1u) % 64u * 1u] = f32(b0[i]);
|
| 13 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/696b0d070511e345cddeb934dc99fc629d0ec7e90ee55799a41d6726d602eb5b.wgsl
ADDED
|
@@ -0,0 +1,24 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
var<workgroup> factor: f32;
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 8 |
+
let row = group.x + group.y * 2048u;
|
| 9 |
+
if (row >= 2048u) { return; }
|
| 10 |
+
let lane = local.x;
|
| 11 |
+
if (lane == 0u) {
|
| 12 |
+
var total = 0.0;
|
| 13 |
+
for (var j = 0u; j < 64u; j++) {
|
| 14 |
+
let v = f32(b0[row * 64u + j]);
|
| 15 |
+
total += v * v;
|
| 16 |
+
}
|
| 17 |
+
factor = inverseSqrt(total / 64.0 + 1e-05);
|
| 18 |
+
}
|
| 19 |
+
workgroupBarrier();
|
| 20 |
+
for (var p = lane; p < 64u; p += 64u) {
|
| 21 |
+
let i = row * 64u + p;
|
| 22 |
+
out[i] = f32(f32(b0[row * 64u + p]) * factor * (f32(b1[p]) + 0.0));
|
| 23 |
+
}
|
| 24 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/6e99c56e3571c8da5671f66929f70a2fd55b3f2fb17f405008ff7195e0900728.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 131072u;
|
| 8 |
+
if (i >= 131072u) { return; }
|
| 9 |
+
out[i] = f32(f32(b0[i]) * f32(b1[((i / 64u) % 64u) * 64u + ((i / 1u) % 64u) * 1u]));
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/6f49fdd6312fb3bfeef8d2bbb2ee3bd9327b5def9064b531b7cb5413524479b7.wgsl
ADDED
|
@@ -0,0 +1,31 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<u32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read> b2: array<f32>;
|
| 4 |
+
@group(0) @binding(3) var<storage, read_write> out: array<f32>;
|
| 5 |
+
fn dequant_1(index: u32) -> f32 {
|
| 6 |
+
let row = index / 4096u;
|
| 7 |
+
let col = index % 4096u;
|
| 8 |
+
let word = b1[row * 1024u + col / 4u];
|
| 9 |
+
let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
|
| 10 |
+
return f32(code) * b2[row * 128u + col / 32u];
|
| 11 |
+
}
|
| 12 |
+
|
| 13 |
+
var<workgroup> tile_a: array<f32, 64>;
|
| 14 |
+
var<workgroup> tile_b: array<f32, 64>;
|
| 15 |
+
@compute @workgroup_size(8, 8)
|
| 16 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 17 |
+
let row = group.y * 8u + local.y;
|
| 18 |
+
let col = group.x * 8u + local.x;
|
| 19 |
+
let i = (group.z * 4u + row) * 2048u + col;
|
| 20 |
+
var acc = 0.0;
|
| 21 |
+
for (var tile = 0u; tile < 4096u; tile += 8u) {
|
| 22 |
+
tile_a[local.y * 8u + local.x] = 0.0;
|
| 23 |
+
tile_b[local.x * 8u + local.y] = 0.0;
|
| 24 |
+
if (row < 4u && tile + local.x < 4096u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 4u + row) * 4096u + tile + local.x]); }
|
| 25 |
+
if (group.x * 8u + local.y < 2048u && tile + local.x < 4096u) { tile_b[local.x * 8u + local.y] = dequant_1((group.x * 8u + local.y) * 4096u + tile + local.x); }
|
| 26 |
+
workgroupBarrier();
|
| 27 |
+
for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
|
| 28 |
+
workgroupBarrier();
|
| 29 |
+
}
|
| 30 |
+
if (row < 4u && col < 2048u) { out[i] = f32(acc + 0.0); }
|
| 31 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/7002f8fdf09f5124ce2875c3540bce3e77750260556fc1a00a8b13b6c1871c30.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<f32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
|
| 5 |
+
@compute @workgroup_size(64)
|
| 6 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 7 |
+
let i = gid.x + gid.y * 512u;
|
| 8 |
+
if (i >= 512u) { return; }
|
| 9 |
+
out[i] = f32(f32(b0[i]) * f32(b1[((i / 1u) % 64u) * 1u]));
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/70dd45c48a989e6617431d39bc63f0187859f791c0915bcfc95fd6ebf4153403.wgsl
ADDED
|
@@ -0,0 +1,31 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<u32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read> b2: array<f32>;
|
| 4 |
+
@group(0) @binding(3) var<storage, read_write> out: array<f32>;
|
| 5 |
+
fn dequant_1(index: u32) -> f32 {
|
| 6 |
+
let row = index / 2048u;
|
| 7 |
+
let col = index % 2048u;
|
| 8 |
+
let word = b1[row * 512u + col / 4u];
|
| 9 |
+
let code = i32((word >> ((col % 4u) * 8u)) & 255u) - 128;
|
| 10 |
+
return f32(code) * b2[row * 64u + col / 32u];
|
| 11 |
+
}
|
| 12 |
+
|
| 13 |
+
var<workgroup> tile_a: array<f32, 64>;
|
| 14 |
+
var<workgroup> tile_b: array<f32, 64>;
|
| 15 |
+
@compute @workgroup_size(8, 8)
|
| 16 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 17 |
+
let row = group.y * 8u + local.y;
|
| 18 |
+
let col = group.x * 8u + local.x;
|
| 19 |
+
let i = (group.z * 4u + row) * 512u + col;
|
| 20 |
+
var acc = 0.0;
|
| 21 |
+
for (var tile = 0u; tile < 2048u; tile += 8u) {
|
| 22 |
+
tile_a[local.y * 8u + local.x] = 0.0;
|
| 23 |
+
tile_b[local.x * 8u + local.y] = 0.0;
|
| 24 |
+
if (row < 4u && tile + local.x < 2048u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 4u + row) * 2048u + tile + local.x]); }
|
| 25 |
+
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); }
|
| 26 |
+
workgroupBarrier();
|
| 27 |
+
for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
|
| 28 |
+
workgroupBarrier();
|
| 29 |
+
}
|
| 30 |
+
if (row < 4u && col < 512u) { out[i] = f32(acc + 0.0); }
|
| 31 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/73eead1530c6564c2c34dd4b18bfa56e96844abc065ac25ad4bed1eeb1f4402c.wgsl
ADDED
|
@@ -0,0 +1,9 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 2048u;
|
| 7 |
+
if (i >= 2048u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 2048u) % 1u) * 2048u + ((i / 256u) % 8u) * 64u + ((i / 64u) % 4u) * 512u + ((i / 1u) % 64u) * 1u]);
|
| 9 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/764403c33cd8994568e068810380cb57e9f2e326dfc09bed63eed1fbed044da1.wgsl
ADDED
|
@@ -0,0 +1,10 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<f32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read_write> out: array<f32>;
|
| 3 |
+
|
| 4 |
+
@compute @workgroup_size(64)
|
| 5 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 6 |
+
let i = gid.x + gid.y * 1024u;
|
| 7 |
+
if (i >= 1024u) { return; }
|
| 8 |
+
let x = f32(b0[i]);
|
| 9 |
+
out[i] = f32(cos(x));
|
| 10 |
+
}
|
exaone4-fp32-int8-g32-release-token-major-v2/kernels/7ab61f48fc60ed8197fe395932ca6d1b90869c9d307fb397eb8f6a26811049e7.wgsl
ADDED
|
@@ -0,0 +1,13 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
enable f16;
|
| 2 |
+
@group(0) @binding(0) var<storage, read> b0: array<i32>;
|
| 3 |
+
@group(0) @binding(1) var<storage, read> b1: array<f16>;
|
| 4 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 5 |
+
|
| 6 |
+
@compute @workgroup_size(64)
|
| 7 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 8 |
+
let i = gid.x + gid.y * 2048u;
|
| 9 |
+
if (i >= 2048u) { return; }
|
| 10 |
+
let token = i32(b0[i / 2048u]);
|
| 11 |
+
if (token < 98304 || token >= 102400) { out[i] = f32(0.0); return; }
|
| 12 |
+
out[i] = f32(b1[u32(token - 98304) * 2048u + i % 2048u]);
|
| 13 |
+
}
|