Publish WebTorch model catalog bundles (part 8)
Browse filesThis view is limited to 50 files because it contains too many changes. See raw diff
- .gitattributes +1 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/228e9c354122040cd7c70f0bcd9445c92011bf85343add619aeb51846c885636.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/234987ec68497956ea7aef8e3f44c8a73dbe4c1519165d123946bedbb4d18e26.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/237458dff084b2aa9ccba5578da3b2a0b062951575364d1464250baca3fdac5a.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/242b7d51c19c9857a096c2ae1f7b72ebd410100a1542c51390950adbb8aed35c.wgsl +17 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/246ea07912fcd8a9073926e863141111eb4091b52cc980a6eef148beed122fc4.wgsl +24 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2505a613f10a5ed1c70a3f80a16dd566da65caa2d7d85083e23f85a30844d63e.wgsl +11 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2622c3df557e0d2b8c61c86cb748cfcf7970ebce8750a37ef611b8cabdd152d6.wgsl +10 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/262399e4b2661342c29fddf17556e93608f7aead68939859c29fcc594f455f69.wgsl +28 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/26cf6eb3b486d1652247ff130b05ed0e8c8560028eeba59acfc075b976bfdfce.wgsl +12 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2711a4a396071f4e99303d01d4854ccaa37a445f4b718ebd7b61cf1c37e56b48.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2792ae5841828bd5af938c10096abe385e63bb3ec2a22e850f146739433c12bf.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/28119aa01122c42bd26ceb37129fe0c95b2a9192163f6dabd90fe058b1e14c54.wgsl +10 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2927b33e105f6ccb9c4352ebdf028ea856d0f47c23a7f6cb140d5acb45356c35.wgsl +17 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2af7517eb26022ada8f8870fd9e1e1ef3da66b9c466336ee18d5295957820b27.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2b02ece9784fc350cef18e08c962031f74dbc27c7d1cad3a13dda1cebb84b403.wgsl +11 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2b9e4b05ef5623d721dbe56b5812e551efcc6fd833d81064320e86fe94d8c670.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2bdeb3590b1ba05cf6e9c7e8245f8aa552a12dae84f9578c603c9cee081a82d4.wgsl +24 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2c0cca2a7b0198c317e68992e64f3fe9bba80b67204c2695c20c372471c150d3.wgsl +10 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2c66dcac8fdd3caa9ff1e2441524070c21c24df9526bee44c49e765a8d4c8222.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2c81f71954e9a2503ae8f3d84f319ffe366cdbb428ab45c22bddee91751c035c.wgsl +12 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2cd4a533d32097a17641196f0cd315d170c587bc2b1c62b42b79d566d6fbf140.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2dd29b61b039fcdd7fbc3c94f629b24b6f81aa5b9923098837597359b295e40d.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2e1ce6ef01198780289fd4a35d8f86b5172595ed189cf931e8444091d49e8d3a.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2e6f807cdf7ba91d7f531d9fb37768171ca71dc6278a9863bdba4a3dbac34fb3.wgsl +10 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2eeb54034621ff097879064a60fd4e0dd5d68201e48296ad234413a39579acc8.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3087f59456d6581e7cdad1a14086af53281e39d2041a6a93387ba6e1daf5f9cb.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/309f1b67894a30dd3e88cd18af63ffc51e384b6891b5f34dd1a0588344188f6e.wgsl +17 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/32a5027cac5dd969d63fe5d5c88d7c87abb600c43ad1d70e1c56615ea781eff3.wgsl +10 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/32cb7d055e8ecf035e10d05724e64544393bd3916079992a3d8dce427bb14181.wgsl +11 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/33e8ec04c66adb4988f562f0f9d3e92d29dfda91a56d0ed62be598c7c737f5ba.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3560fe2bb07bf6fb92598b9f9c006029892f16dcb55b418973ed89b2b5086b1b.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/359d61f2ba174e54e2a0920fcf61e84e436fd3eaf7e2832e21fdb24094493667.wgsl +11 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/35b866864f2f36fad2f64440227cd6088083c96b0ef50192ce71ef159e1d1b11.wgsl +24 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/360a02a53c176d885331b307a6972d974ce5c8dad0d961b8521c6fd8b37d7ebe.wgsl +17 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3622ffa2072335a35d6c9db5d284ed92b3fb5ed304acc2c606408063e7e21dcd.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/36c55d910d6fdaf8d3a2bba5a559bf7dcf54364986a7f2c0610b523600667d91.wgsl +22 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3830061b0fc54c85d3d8a783b927c44484b1a3b68b52b8d8a1b980f7553deebf.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/38bd527686bcade9ed19a52788ccd711ca668e76854cc6274c480860fce88ff2.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/39646bdfafd1ab0cb874474e3f1e9d79b57f10582a1804bc1d3762a492fec851.wgsl +10 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/39e41393f5850cc4f0b2e8172a3c53c8c7fb874ceb4df2c6f13f9af6128ad7fc.wgsl +10 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3a76085983a5330bf07a4166ea4975700403f9d8035a8ef450f7f18b9cf47c66.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3b2066052a48ffaf36270fc0f07d7725e7146f7d0cabef87ecc8507eda259a1d.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3b2b13d75a6ce00e35c674883bbf1969555c8fc9e088e7d2003a5ea57b4552ff.wgsl +22 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3b8210d9acf57a3c1cca9e926547520d5994967f27ce92897a15b14137f7da66.wgsl +9 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3bb850bda95bf77706a2839616a401b773f9de0428d1b427e61fb04e2c35fe74.wgsl +11 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3d4359c4f1d0d65e6b6049b73a4e3f5fb67573fb80e29c9321f3671f40786f65.wgsl +28 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3d5b97cfaf8dd0e680b6a9c84e7eca497f9f6e49817d20f1a72dddd7a3303f73.wgsl +11 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3daeb9cba5c07ed47ac55e1c2565aa82f7d8e8409037707e7339d1d96d122144.wgsl +11 -0
- qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3e34409407100d4858c4c109cf62f3eaf2f4c6081166b0ec8679105e6dfbe7fe.wgsl +10 -0
.gitattributes
CHANGED
|
@@ -40,3 +40,4 @@ qwen35-08b-fp32-int8-g32-home-token-major-v2/tokenizer/tokenizer.json filter=lfs
|
|
| 40 |
qwen35-2b-fp32-int8-g32-home-token-major-v2/graph.json filter=lfs diff=lfs merge=lfs -text
|
| 41 |
qwen35-2b-fp32-int8-g32-home-token-major-v2/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
|
| 42 |
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/graph.json filter=lfs diff=lfs merge=lfs -text
|
|
|
|
|
|
| 40 |
qwen35-2b-fp32-int8-g32-home-token-major-v2/graph.json filter=lfs diff=lfs merge=lfs -text
|
| 41 |
qwen35-2b-fp32-int8-g32-home-token-major-v2/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
|
| 42 |
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/graph.json filter=lfs diff=lfs merge=lfs -text
|
| 43 |
+
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/tokenizer/tokenizer.json filter=lfs diff=lfs merge=lfs -text
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/228e9c354122040cd7c70f0bcd9445c92011bf85343add619aeb51846c885636.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 >= 48u) { return; }
|
| 8 |
+
out[i] = i32(b0[i]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/234987ec68497956ea7aef8e3f44c8a73dbe4c1519165d123946bedbb4d18e26.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 / 128u) % 16u) * 128u + 0u * 128u + ((i / 1u) % 128u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/237458dff084b2aa9ccba5578da3b2a0b062951575364d1464250baca3fdac5a.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 * 32768u;
|
| 7 |
+
if (i >= 32768u) { return; }
|
| 8 |
+
out[i] = f32(f32(b0[i]) * f32(b0[i]));
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/242b7d51c19c9857a096c2ae1f7b72ebd410100a1542c51390950adbb8aed35c.wgsl
ADDED
|
@@ -0,0 +1,17 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<i32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<u32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
fn unpack_bf16_1(index: u32) -> f32 {
|
| 5 |
+
let pair = b1[index / 2u];
|
| 6 |
+
let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
|
| 7 |
+
return bitcast<f32>(bits << 16u);
|
| 8 |
+
}
|
| 9 |
+
|
| 10 |
+
@compute @workgroup_size(64)
|
| 11 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 12 |
+
let i = gid.x + gid.y * 2048u;
|
| 13 |
+
if (i >= 2048u) { return; }
|
| 14 |
+
let token = i32(b0[i / 2048u]);
|
| 15 |
+
if (token < 32768 || token >= 65536) { out[i] = f32(0.0); return; }
|
| 16 |
+
out[i] = f32(unpack_bf16_1(u32(token - 32768) * 2048u + i % 2048u));
|
| 17 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/246ea07912fcd8a9073926e863141111eb4091b52cc980a6eef148beed122fc4.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 * 32u;
|
| 9 |
+
if (row >= 32u) { return; }
|
| 10 |
+
let lane = local.x;
|
| 11 |
+
if (lane == 0u) {
|
| 12 |
+
var total = 0.0;
|
| 13 |
+
for (var j = 0u; j < 256u; j++) {
|
| 14 |
+
let v = f32(b0[row * 256u + j]);
|
| 15 |
+
total += v * v;
|
| 16 |
+
}
|
| 17 |
+
factor = inverseSqrt(total / 256.0 + 1e-06);
|
| 18 |
+
}
|
| 19 |
+
workgroupBarrier();
|
| 20 |
+
for (var p = lane; p < 256u; p += 64u) {
|
| 21 |
+
let i = row * 256u + p;
|
| 22 |
+
out[i] = f32(f32(b0[row * 256u + p]) * factor * (f32(b1[p]) + 1.0));
|
| 23 |
+
}
|
| 24 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2505a613f10a5ed1c70a3f80a16dd566da65caa2d7d85083e23f85a30844d63e.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 * 98304u;
|
| 8 |
+
if (i >= 98304u) { 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 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2622c3df557e0d2b8c61c86cb748cfcf7970ebce8750a37ef611b8cabdd152d6.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 / 4096u) % 4u) * 4096u + ((i / 1u) % 4096u) * 1u]) * 1.0));
|
| 10 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/262399e4b2661342c29fddf17556e93608f7aead68939859c29fcc594f455f69.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_write> out: array<f32>;
|
| 4 |
+
fn unpack_bf16_1(index: u32) -> f32 {
|
| 5 |
+
let pair = b1[index / 2u];
|
| 6 |
+
let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
|
| 7 |
+
return bitcast<f32>(bits << 16u);
|
| 8 |
+
}
|
| 9 |
+
|
| 10 |
+
var<workgroup> tile_a: array<f32, 64>;
|
| 11 |
+
var<workgroup> tile_b: array<f32, 64>;
|
| 12 |
+
@compute @workgroup_size(8, 8)
|
| 13 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 14 |
+
let row = group.y * 8u + local.y;
|
| 15 |
+
let col = group.x * 8u + local.x;
|
| 16 |
+
let i = (group.z * 4u + row) * 512u + col;
|
| 17 |
+
var acc = 0.0;
|
| 18 |
+
for (var tile = 0u; tile < 2048u; tile += 8u) {
|
| 19 |
+
tile_a[local.y * 8u + local.x] = 0.0;
|
| 20 |
+
tile_b[local.x * 8u + local.y] = 0.0;
|
| 21 |
+
if (row < 4u && tile + local.x < 2048u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 4u + row) * 2048u + tile + local.x]); }
|
| 22 |
+
if (group.x * 8u + local.y < 512u && tile + local.x < 2048u) { tile_b[local.x * 8u + local.y] = unpack_bf16_1((group.x * 8u + local.y) * 2048u + tile + local.x); }
|
| 23 |
+
workgroupBarrier();
|
| 24 |
+
for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
|
| 25 |
+
workgroupBarrier();
|
| 26 |
+
}
|
| 27 |
+
if (row < 4u && col < 512u) { out[i] = f32(acc + 0.0); }
|
| 28 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/26cf6eb3b486d1652247ff130b05ed0e8c8560028eeba59acfc075b976bfdfce.wgsl
ADDED
|
@@ -0,0 +1,12 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 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 * 30720u;
|
| 8 |
+
if (i >= 30720u) { return; }
|
| 9 |
+
let coord = (i / 1u) % 5u;
|
| 10 |
+
if (coord >= 0u && coord < 4u) { out[i] = f32(b0[(i / 5u) * 4u + (coord - 0u) * 1u + i % 1u]); }
|
| 11 |
+
if (coord >= 4u && coord < 5u) { out[i] = f32(b1[(i / 5u) * 1u + (coord - 4u) * 1u + i % 1u]); }
|
| 12 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2711a4a396071f4e99303d01d4854ccaa37a445f4b718ebd7b61cf1c37e56b48.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 >= 16u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 16u) % 1u) * 16u + ((i / 1u) % 16u) * 1u + ((i / 1u) % 1u) * 16u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2792ae5841828bd5af938c10096abe385e63bb3ec2a22e850f146739433c12bf.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 / 2048u) % 4u) * 256u + ((i / 256u) % 8u) * 1024u + ((i / 1u) % 256u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/28119aa01122c42bd26ceb37129fe0c95b2a9192163f6dabd90fe058b1e14c54.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 * 4096u;
|
| 7 |
+
if (i >= 4096u) { return; }
|
| 8 |
+
let x = f32(b0[i]);
|
| 9 |
+
out[i] = f32(sin(x));
|
| 10 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2927b33e105f6ccb9c4352ebdf028ea856d0f47c23a7f6cb140d5acb45356c35.wgsl
ADDED
|
@@ -0,0 +1,17 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<i32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<u32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
fn unpack_bf16_1(index: u32) -> f32 {
|
| 5 |
+
let pair = b1[index / 2u];
|
| 6 |
+
let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
|
| 7 |
+
return bitcast<f32>(bits << 16u);
|
| 8 |
+
}
|
| 9 |
+
|
| 10 |
+
@compute @workgroup_size(64)
|
| 11 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 12 |
+
let i = gid.x + gid.y * 131072u;
|
| 13 |
+
if (i >= 131072u) { return; }
|
| 14 |
+
let token = i32(b0[i / 2048u]);
|
| 15 |
+
if (token < 229376 || token >= 248320) { out[i] = f32(0.0); return; }
|
| 16 |
+
out[i] = f32(unpack_bf16_1(u32(token - 229376) * 2048u + i % 2048u));
|
| 17 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2af7517eb26022ada8f8870fd9e1e1ef3da66b9c466336ee18d5295957820b27.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 * 417792u;
|
| 7 |
+
if (i >= 417792u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 417792u) % 1u) * 436224u + ((i / 68u) % 6144u) * 71u + (((i / 1u) % 68u) * 1u + 0u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2b02ece9784fc350cef18e08c962031f74dbc27c7d1cad3a13dda1cebb84b403.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 * 393216u;
|
| 8 |
+
if (i >= 393216u) { 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 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2b9e4b05ef5623d721dbe56b5812e551efcc6fd833d81064320e86fe94d8c670.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(f32(b0[i]) * 1.0);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2bdeb3590b1ba05cf6e9c7e8245f8aa552a12dae84f9578c603c9cee081a82d4.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 * 16u;
|
| 9 |
+
if (row >= 16u) { return; }
|
| 10 |
+
let lane = local.x;
|
| 11 |
+
if (lane == 0u) {
|
| 12 |
+
var total = 0.0;
|
| 13 |
+
for (var j = 0u; j < 128u; j++) {
|
| 14 |
+
let v = f32(b0[row * 128u + j]);
|
| 15 |
+
total += v * v;
|
| 16 |
+
}
|
| 17 |
+
factor = inverseSqrt(total / 128.0 + 1e-06);
|
| 18 |
+
}
|
| 19 |
+
workgroupBarrier();
|
| 20 |
+
for (var p = lane; p < 128u; p += 64u) {
|
| 21 |
+
let i = row * 128u + p;
|
| 22 |
+
out[i] = f32(f32(b0[row * 128u + p]) * factor * (f32(b1[p]) + 0.0));
|
| 23 |
+
}
|
| 24 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2c0cca2a7b0198c317e68992e64f3fe9bba80b67204c2695c20c372471c150d3.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 * 262144u;
|
| 8 |
+
if (i >= 262144u) { return; }
|
| 9 |
+
out[i] = f32(f32(b0[((i / 16384u) % 16u) * 128u + ((i / 128u) % 128u) * 1u]) * f32(b1[((i / 16384u) % 16u) * 128u + ((i / 1u) % 128u) * 1u]));
|
| 10 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2c66dcac8fdd3caa9ff1e2441524070c21c24df9526bee44c49e765a8d4c8222.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 + 0u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2c81f71954e9a2503ae8f3d84f319ffe366cdbb428ab45c22bddee91751c035c.wgsl
ADDED
|
@@ -0,0 +1,12 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 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 * 3973120u;
|
| 8 |
+
if (i >= 3973120u) { return; }
|
| 9 |
+
let coord = (i / 1u) % 248320u;
|
| 10 |
+
if (coord >= 0u && coord < 229376u) { out[i] = f32(b0[(i / 248320u) * 229376u + (coord - 0u) * 1u + i % 1u]); }
|
| 11 |
+
if (coord >= 229376u && coord < 248320u) { out[i] = f32(b1[(i / 248320u) * 18944u + (coord - 229376u) * 1u + i % 1u]); }
|
| 12 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2cd4a533d32097a17641196f0cd315d170c587bc2b1c62b42b79d566d6fbf140.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 * 384u;
|
| 7 |
+
if (i >= 384u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 128u) % 3u) * 128u + ((i / 128u) % 1u) * 128u + ((i / 4u) % 32u) * 1u + ((i / 1u) % 4u) * 32u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2dd29b61b039fcdd7fbc3c94f629b24b6f81aa5b9923098837597359b295e40d.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 / 2048u) % 1u) * 128u + ((i / 128u) % 16u) * 128u + ((i / 1u) % 128u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-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 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2e6f807cdf7ba91d7f531d9fb37768171ca71dc6278a9863bdba4a3dbac34fb3.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) % 16u) * 4096u + ((i / 1u) % 4096u) * 1u]) * 1.0));
|
| 10 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/2eeb54034621ff097879064a60fd4e0dd5d68201e48296ad234413a39579acc8.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) * 16384u + ((i / 2048u) % 4u) * 4096u + ((i / 256u) % 8u) * 512u + (((i / 1u) % 256u) * 1u + 0u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3087f59456d6581e7cdad1a14086af53281e39d2041a6a93387ba6e1daf5f9cb.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 * 262144u;
|
| 7 |
+
if (i >= 262144u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 262144u) % 1u) * 270336u + ((i / 16384u) % 16u) * 16896u + (((i / 128u) % 128u) * 1u + 0u) * 128u + ((i / 1u) % 128u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/309f1b67894a30dd3e88cd18af63ffc51e384b6891b5f34dd1a0588344188f6e.wgsl
ADDED
|
@@ -0,0 +1,17 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<i32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<u32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
fn unpack_bf16_1(index: u32) -> f32 {
|
| 5 |
+
let pair = b1[index / 2u];
|
| 6 |
+
let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
|
| 7 |
+
return bitcast<f32>(bits << 16u);
|
| 8 |
+
}
|
| 9 |
+
|
| 10 |
+
@compute @workgroup_size(64)
|
| 11 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 12 |
+
let i = gid.x + gid.y * 32768u;
|
| 13 |
+
if (i >= 32768u) { return; }
|
| 14 |
+
let token = i32(b0[i / 2048u]);
|
| 15 |
+
if (token < 65536 || token >= 98304) { out[i] = f32(0.0); return; }
|
| 16 |
+
out[i] = f32(unpack_bf16_1(u32(token - 65536) * 2048u + i % 2048u));
|
| 17 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/32a5027cac5dd969d63fe5d5c88d7c87abb600c43ad1d70e1c56615ea781eff3.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 * 262144u;
|
| 8 |
+
if (i >= 262144u) { return; }
|
| 9 |
+
out[i] = f32(f32(b0[i]) * f32(b1[((i / 16384u) % 16u) * 128u + ((i / 128u) % 128u) * 1u]));
|
| 10 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/32cb7d055e8ecf035e10d05724e64544393bd3916079992a3d8dce427bb14181.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 * 1024u;
|
| 7 |
+
if (i >= 1024u) { return; }
|
| 8 |
+
var acc = 0.0;
|
| 9 |
+
for (var j = 0u; j < 128u; j++) { acc += f32(b0[(((i / 1024u) % 1u) * 131072u + ((i / 16u) % 64u) * 2048u + ((i / 1u) % 16u) * 128u) + (((j / 1u) % 128u) * 1u)]); }
|
| 10 |
+
out[i] = f32(acc);
|
| 11 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-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 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3560fe2bb07bf6fb92598b9f9c006029892f16dcb55b418973ed89b2b5086b1b.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) * 393216u + ((i / 2048u) % 64u) * 6144u + (((i / 1u) % 2048u) * 1u + 4096u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/359d61f2ba174e54e2a0920fcf61e84e436fd3eaf7e2832e21fdb24094493667.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 * 512u;
|
| 8 |
+
if (i >= 512u) { return; }
|
| 9 |
+
let coord = (i / 1u) % 32u;
|
| 10 |
+
if (coord >= 2u && coord < 32u && (coord - 2u) % 3u == 0u) { out[i] = f32(b0[((i / 512u) % 1u) * 160u + ((i / 32u) % 16u) * 10u + ((((i / 1u) % 32u) - 2u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
|
| 11 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/35b866864f2f36fad2f64440227cd6088083c96b0ef50192ce71ef159e1d1b11.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-06);
|
| 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]) + 1.0));
|
| 23 |
+
}
|
| 24 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/360a02a53c176d885331b307a6972d974ce5c8dad0d961b8521c6fd8b37d7ebe.wgsl
ADDED
|
@@ -0,0 +1,17 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
@group(0) @binding(0) var<storage, read> b0: array<i32>;
|
| 2 |
+
@group(0) @binding(1) var<storage, read> b1: array<u32>;
|
| 3 |
+
@group(0) @binding(2) var<storage, read_write> out: array<f32>;
|
| 4 |
+
fn unpack_bf16_1(index: u32) -> f32 {
|
| 5 |
+
let pair = b1[index / 2u];
|
| 6 |
+
let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
|
| 7 |
+
return bitcast<f32>(bits << 16u);
|
| 8 |
+
}
|
| 9 |
+
|
| 10 |
+
@compute @workgroup_size(64)
|
| 11 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 12 |
+
let i = gid.x + gid.y * 2048u;
|
| 13 |
+
if (i >= 2048u) { return; }
|
| 14 |
+
let token = i32(b0[i / 2048u]);
|
| 15 |
+
if (token < 65536 || token >= 98304) { out[i] = f32(0.0); return; }
|
| 16 |
+
out[i] = f32(unpack_bf16_1(u32(token - 65536) * 2048u + i % 2048u));
|
| 17 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3622ffa2072335a35d6c9db5d284ed92b3fb5ed304acc2c606408063e7e21dcd.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 * 32768u;
|
| 7 |
+
if (i >= 32768u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 32768u) % 1u) * 65536u + ((i / 2048u) % 16u) * 4096u + ((i / 256u) % 8u) * 512u + (((i / 1u) % 256u) * 1u + 256u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/36c55d910d6fdaf8d3a2bba5a559bf7dcf54364986a7f2c0610b523600667d91.wgsl
ADDED
|
@@ -0,0 +1,22 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 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> b4: array<f32>;
|
| 6 |
+
@group(0) @binding(5) var<storage, read> b5: array<f32>;
|
| 7 |
+
@group(0) @binding(6) var<storage, read> b6: array<f32>;
|
| 8 |
+
@group(0) @binding(7) var<storage, read_write> out: array<f32>;
|
| 9 |
+
|
| 10 |
+
@compute @workgroup_size(64)
|
| 11 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 12 |
+
let i = gid.x + gid.y * 3670016u;
|
| 13 |
+
if (i >= 3670016u) { return; }
|
| 14 |
+
let coord = (i / 1u) % 229376u;
|
| 15 |
+
if (coord >= 0u && coord < 32768u) { out[i] = f32(b0[(i / 229376u) * 32768u + (coord - 0u) * 1u + i % 1u]); }
|
| 16 |
+
if (coord >= 32768u && coord < 65536u) { out[i] = f32(b1[(i / 229376u) * 32768u + (coord - 32768u) * 1u + i % 1u]); }
|
| 17 |
+
if (coord >= 65536u && coord < 98304u) { out[i] = f32(b2[(i / 229376u) * 32768u + (coord - 65536u) * 1u + i % 1u]); }
|
| 18 |
+
if (coord >= 98304u && coord < 131072u) { out[i] = f32(b3[(i / 229376u) * 32768u + (coord - 98304u) * 1u + i % 1u]); }
|
| 19 |
+
if (coord >= 131072u && coord < 163840u) { out[i] = f32(b4[(i / 229376u) * 32768u + (coord - 131072u) * 1u + i % 1u]); }
|
| 20 |
+
if (coord >= 163840u && coord < 196608u) { out[i] = f32(b5[(i / 229376u) * 32768u + (coord - 163840u) * 1u + i % 1u]); }
|
| 21 |
+
if (coord >= 196608u && coord < 229376u) { out[i] = f32(b6[(i / 229376u) * 32768u + (coord - 196608u) * 1u + i % 1u]); }
|
| 22 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3830061b0fc54c85d3d8a783b927c44484b1a3b68b52b8d8a1b980f7553deebf.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 * 1024u;
|
| 7 |
+
if (i >= 1024u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 1024u) % 1u) * 2048u + ((i / 512u) % 2u) * 1024u + ((i / 32u) % 16u) * 64u + (((i / 1u) % 32u) * 1u + 32u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/38bd527686bcade9ed19a52788ccd711ca668e76854cc6274c480860fce88ff2.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) * 6144u + ((i / 2048u) % 1u) * 6144u + (((i / 1u) % 2048u) * 1u + 0u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-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 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/39e41393f5850cc4f0b2e8172a3c53c8c7fb874ceb4df2c6f13f9af6128ad7fc.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 / 64u) % 64u) * 64u + ((i / 1u) % 64u) * 1u]));
|
| 10 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3a76085983a5330bf07a4166ea4975700403f9d8035a8ef450f7f18b9cf47c66.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) * 262144u + ((i / 2048u) % 64u) * 4096u + ((i / 256u) % 8u) * 512u + (((i / 1u) % 256u) * 1u + 0u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3b2066052a48ffaf36270fc0f07d7725e7146f7d0cabef87ecc8507eda259a1d.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 * 393216u;
|
| 7 |
+
if (i >= 393216u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 393216u) % 1u) * 393216u + ((i / 6144u) % 64u) * 1u + ((i / 1u) % 6144u) * 64u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3b2b13d75a6ce00e35c674883bbf1969555c8fc9e088e7d2003a5ea57b4552ff.wgsl
ADDED
|
@@ -0,0 +1,22 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 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> b4: array<f32>;
|
| 6 |
+
@group(0) @binding(5) var<storage, read> b5: array<f32>;
|
| 7 |
+
@group(0) @binding(6) var<storage, read> b6: array<f32>;
|
| 8 |
+
@group(0) @binding(7) var<storage, read_write> out: array<f32>;
|
| 9 |
+
|
| 10 |
+
@compute @workgroup_size(64)
|
| 11 |
+
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
| 12 |
+
let i = gid.x + gid.y * 4194240u;
|
| 13 |
+
if (i >= 14680064u) { return; }
|
| 14 |
+
let coord = (i / 1u) % 229376u;
|
| 15 |
+
if (coord >= 0u && coord < 32768u) { out[i] = f32(b0[(i / 229376u) * 32768u + (coord - 0u) * 1u + i % 1u]); }
|
| 16 |
+
if (coord >= 32768u && coord < 65536u) { out[i] = f32(b1[(i / 229376u) * 32768u + (coord - 32768u) * 1u + i % 1u]); }
|
| 17 |
+
if (coord >= 65536u && coord < 98304u) { out[i] = f32(b2[(i / 229376u) * 32768u + (coord - 65536u) * 1u + i % 1u]); }
|
| 18 |
+
if (coord >= 98304u && coord < 131072u) { out[i] = f32(b3[(i / 229376u) * 32768u + (coord - 98304u) * 1u + i % 1u]); }
|
| 19 |
+
if (coord >= 131072u && coord < 163840u) { out[i] = f32(b4[(i / 229376u) * 32768u + (coord - 131072u) * 1u + i % 1u]); }
|
| 20 |
+
if (coord >= 163840u && coord < 196608u) { out[i] = f32(b5[(i / 229376u) * 32768u + (coord - 163840u) * 1u + i % 1u]); }
|
| 21 |
+
if (coord >= 196608u && coord < 229376u) { out[i] = f32(b6[(i / 229376u) * 32768u + (coord - 196608u) * 1u + i % 1u]); }
|
| 22 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3b8210d9acf57a3c1cca9e926547520d5994967f27ce92897a15b14137f7da66.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 * 192u;
|
| 7 |
+
if (i >= 176u) { return; }
|
| 8 |
+
out[i] = f32(b0[((i / 176u) % 1u) * 512u + ((i / 11u) % 16u) * 32u + (((i / 1u) % 11u) * 3u + 1u) * 1u]);
|
| 9 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3bb850bda95bf77706a2839616a401b773f9de0428d1b427e61fb04e2c35fe74.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 * 128u;
|
| 8 |
+
if (i >= 96u) { return; }
|
| 9 |
+
let coord = (i / 32u) % 3u;
|
| 10 |
+
if (coord == 0u) { out[i] = f32(b0[((i / 32u) % 1u) * 32u + ((i / 32u) % 1u) * 32u + ((i / 1u) % 32u) * 1u]); } else { out[i] = f32(b1[i]); }
|
| 11 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3d4359c4f1d0d65e6b6049b73a4e3f5fb67573fb80e29c9321f3671f40786f65.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_write> out: array<f32>;
|
| 4 |
+
fn unpack_bf16_1(index: u32) -> f32 {
|
| 5 |
+
let pair = b1[index / 2u];
|
| 6 |
+
let bits = (pair >> ((index % 2u) * 16u)) & 65535u;
|
| 7 |
+
return bitcast<f32>(bits << 16u);
|
| 8 |
+
}
|
| 9 |
+
|
| 10 |
+
var<workgroup> tile_a: array<f32, 64>;
|
| 11 |
+
var<workgroup> tile_b: array<f32, 64>;
|
| 12 |
+
@compute @workgroup_size(8, 8)
|
| 13 |
+
fn main(@builtin(workgroup_id) group: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
|
| 14 |
+
let row = group.y * 8u + local.y;
|
| 15 |
+
let col = group.x * 8u + local.x;
|
| 16 |
+
let i = (group.z * 16u + row) * 2048u + col;
|
| 17 |
+
var acc = 0.0;
|
| 18 |
+
for (var tile = 0u; tile < 6144u; tile += 8u) {
|
| 19 |
+
tile_a[local.y * 8u + local.x] = 0.0;
|
| 20 |
+
tile_b[local.x * 8u + local.y] = 0.0;
|
| 21 |
+
if (row < 16u && tile + local.x < 6144u) { tile_a[local.y * 8u + local.x] = f32(b0[(group.z * 16u + row) * 6144u + tile + local.x]); }
|
| 22 |
+
if (group.x * 8u + local.y < 2048u && tile + local.x < 6144u) { tile_b[local.x * 8u + local.y] = unpack_bf16_1((group.x * 8u + local.y) * 6144u + tile + local.x); }
|
| 23 |
+
workgroupBarrier();
|
| 24 |
+
for (var p = 0u; p < 8u; p++) { acc += tile_a[local.y * 8u + p] * tile_b[p * 8u + local.x]; }
|
| 25 |
+
workgroupBarrier();
|
| 26 |
+
}
|
| 27 |
+
if (row < 16u && col < 2048u) { out[i] = f32(acc + 0.0); }
|
| 28 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3d5b97cfaf8dd0e680b6a9c84e7eca497f9f6e49817d20f1a72dddd7a3303f73.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 * 24576u;
|
| 8 |
+
if (i >= 24576u) { 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 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-token-major-v2/kernels/3daeb9cba5c07ed47ac55e1c2565aa82f7d8e8409037707e7339d1d96d122144.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 * 512u;
|
| 8 |
+
if (i >= 512u) { return; }
|
| 9 |
+
let coord = (i / 1u) % 32u;
|
| 10 |
+
if (coord >= 1u && coord < 34u && (coord - 1u) % 3u == 0u) { out[i] = f32(b0[((i / 512u) % 1u) * 176u + ((i / 32u) % 16u) * 11u + ((((i / 1u) % 32u) - 1u) / 3u) * 1u]); } else { out[i] = f32(b1[i]); }
|
| 11 |
+
}
|
qwen35-2b-multimodal-fp32-three-grid-state-alias-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 |
+
}
|