| <!doctype html> |
| <html lang="en"> |
| <head> |
| <meta charset="utf-8"> |
| <meta name="viewport" content="width=device-width, initial-scale=1"> |
| <title>İvme-Conversate-v2 · raw WebGPU inference</title> |
| <style> |
| @import url('https://fonts.googleapis.com/css2?family=Newsreader:ital,opsz,wght@0,6..72,380;0,6..72,500;0,6..72,600;0,6..72,700;1,6..72,400&family=JetBrains+Mono:wght@400;500;600&display=swap'); |
| |
| html{ background:#f7f4ec; } |
| *,*::before,*::after{ box-sizing:border-box; } |
| body{ |
| margin:0; |
| background: |
| repeating-linear-gradient(180deg, rgba(0,0,0,0.014) 0px, rgba(0,0,0,0.014) 1px, transparent 1px, transparent 3px), |
| #f7f4ec; |
| color:#1c1a16; |
| font-family:'Newsreader', Georgia, serif; |
| min-height:100vh; |
| -webkit-font-smoothing:antialiased; |
| } |
| ::selection{ background:#ff0000; color:#f7f4ec; } |
| a{ color:#c40000; text-decoration:none; border-bottom:1px solid rgba(255,0,0,.35); } |
| a:hover{ border-bottom-color:#ff0000; } |
| |
| code{ |
| font-family:'JetBrains Mono', monospace !important; |
| font-size:.86em; |
| background:#ece5d3; |
| color:#a02020; |
| padding:1px 6px; |
| border-radius:2px; |
| border:1px solid #d9d2bf; |
| } |
| |
| .page{ max-width:880px; margin:0 auto; padding:0 32px 70px; } |
| |
| |
| .masthead{ padding:52px 0 0; text-align:center; } |
| .eyebrow{ |
| font-family:'JetBrains Mono', monospace; |
| font-size:11px; letter-spacing:.28em; text-transform:uppercase; |
| color:#ff0000; margin-bottom:22px; |
| display:flex; align-items:center; justify-content:center; gap:10px; |
| } |
| .eyebrow .dot-live{ |
| width:6px; height:6px; border-radius:50%; background:#c7c1ae; display:inline-block; |
| } |
| .eyebrow .dot-live.on{ background:#ff0000; animation:pulse 2.4s ease-in-out infinite; } |
| @keyframes pulse{ 0%,100%{opacity:1;} 50%{opacity:.35;} } |
| @media (prefers-reduced-motion: reduce){ .eyebrow .dot-live{ animation:none; } } |
| |
| .wordmark{ |
| font-size:clamp(40px, 7vw, 68px); |
| font-weight:700; line-height:.98; letter-spacing:-.015em; margin:0; |
| } |
| .wordmark .dot{ color:#ff0000; } |
| .tagline{ |
| font-size:16.5px; font-style:italic; color:#4a463d; |
| max-width:520px; margin:18px auto 0; line-height:1.6; |
| } |
| .tagline code{ font-style:normal; } |
| |
| .tickline{ |
| height:14px; margin:40px 0 0; |
| background-image: repeating-linear-gradient(90deg, #ff0000 0, #ff0000 1px, transparent 1px, transparent 10px); |
| background-position: center; background-size:10px 1px; background-repeat:repeat-x; |
| background-position-y:6px; opacity:.9; |
| } |
| |
| |
| .meta-strip{ |
| display:flex; justify-content:center; gap:0; padding:26px 0 6px; flex-wrap:wrap; |
| } |
| .meta-item{ |
| font-family:'JetBrains Mono', monospace; font-size:11px; letter-spacing:.05em; |
| color:#4a463d; padding:0 22px; border-right:1px solid #d9d2bf; text-align:center; |
| } |
| .meta-item:last-child{ border-right:none; } |
| .meta-item b{ |
| display:block; color:#1c1a16; font-weight:600; font-size:12.5px; |
| margin-top:2px; letter-spacing:.02em; |
| } |
| |
| |
| section{ padding:48px 0; border-top:1px solid #d9d2bf; } |
| .section-head{ display:flex; align-items:baseline; gap:16px; margin-bottom:8px; } |
| .section-index{ font-family:'JetBrains Mono', monospace; font-size:13px; color:#ff0000; letter-spacing:.05em; } |
| h2{ font-size:24px; font-weight:600; margin:0; letter-spacing:-.01em; } |
| .section-desc{ |
| font-size:14.5px; color:#4a463d; max-width:580px; margin:12px 0 30px; line-height:1.65; |
| } |
| p.lede{ font-size:16.5px; line-height:1.72; color:#2b2822; max-width:640px; margin:0 0 30px; } |
| p.lede b{ font-weight:600; box-shadow: inset 0 -.35em 0 rgba(255,0,0,.16); } |
| |
| |
| .stage-list{ border-top:1px solid #1c1a16; } |
| .stage{ |
| display:grid; grid-template-columns:100px 1fr auto; gap:20px; align-items:baseline; |
| padding:13px 0; border-bottom:1px solid #d9d2bf; |
| } |
| .stage .mark{ |
| font-family:'JetBrains Mono', monospace; font-size:10px; letter-spacing:.06em; text-transform:uppercase; |
| color:#8a8474; position:relative; padding-left:14px; |
| } |
| .stage .mark::before{ |
| content:""; position:absolute; left:0; top:4px; width:6px; height:6px; border-radius:50%; background:#c7c1ae; |
| } |
| .stage.active .mark{ color:#1c1a16; } |
| .stage.active .mark::before{ background:#ff0000; animation:pulse 1.1s ease-in-out infinite; } |
| .stage.done .mark{ color:#1c1a16; } |
| .stage.done .mark::before{ background:#ff0000; } |
| .stage .label{ font-size:14.5px; color:#4a463d; } |
| .stage.done .label, .stage.active .label{ color:#1c1a16; } |
| .stage .meta{ |
| font-family:'JetBrains Mono', monospace; font-size:11px; color:#8a8474; white-space:nowrap; text-align:right; |
| } |
| |
| .progress-track{ height:2px; background:#ece7d8; margin-top:22px; position:relative; overflow:hidden; } |
| .progress-fill{ position:absolute; left:0; top:0; bottom:0; width:0%; background:#ff0000; transition:width .15s ease; } |
| .progress-num{ |
| font-family:'JetBrains Mono', monospace; font-size:11px; color:#8a8474; margin-top:9px; |
| display:flex; justify-content:space-between; |
| } |
| |
| .load-row{ margin-top:26px; display:flex; align-items:center; gap:16px; flex-wrap:wrap; } |
| |
| button{ |
| font-family:'JetBrains Mono', monospace; font-size:11.5px; letter-spacing:.05em; text-transform:uppercase; |
| background:#1c1a16; color:#f7f4ec; border:1px solid #1c1a16; |
| padding:13px 24px; cursor:pointer; transition:background .15s ease, color .15s ease, border-color .15s ease; |
| } |
| button:hover:not(:disabled){ background:#ff0000; border-color:#ff0000; } |
| button:disabled{ opacity:.35; cursor:not-allowed; } |
| button.ghost{ background:transparent; color:#1c1a16; border:1px solid #d9d2bf; } |
| button.ghost:hover:not(:disabled){ background:#1c1a16; color:#f7f4ec; border-color:#1c1a16; } |
| .load-hint{ font-family:'JetBrains Mono', monospace; font-size:11px; color:#8a8474; } |
| |
| .note-box{ |
| font-size:13.5px; color:#4a463d; line-height:1.6; margin:18px 0 0; |
| padding-left:16px; border-left:2px solid #ff0000; |
| } |
| .err-box{ |
| font-family:'JetBrains Mono', monospace; font-size:12px; line-height:1.6; color:#7a1010; |
| margin:18px 0 0; padding-left:16px; border-left:2px solid #c40000; white-space:pre-wrap; word-break:break-word; |
| } |
| |
| |
| table.kv{ width:100%; border-collapse:collapse; font-size:13.5px; } |
| table.kv td{ padding:11px 12px 11px 0; border-bottom:1px solid #ece7d8; vertical-align:top; } |
| table.kv tr:first-child td{ padding-top:0; } |
| table.kv tr:last-child td{ border-bottom:none; } |
| table.kv td:first-child{ |
| font-family:'JetBrains Mono', monospace; font-size:11px; color:#8a8474; width:38%; white-space:nowrap; padding-top:12px; |
| } |
| table.kv td:last-child{ color:#4a463d; line-height:1.55; } |
| |
| |
| .chat-shell{ display:none; } |
| .chat-shell.active{ display:block; } |
| |
| .stat-strip{ display:flex; justify-content:flex-start; gap:0; margin-bottom:32px; } |
| .stat{ padding:0 26px 0 0; margin-right:26px; border-right:1px solid #d9d2bf; } |
| .stat:last-child{ border-right:none; margin-right:0; padding-right:0; } |
| .stat-num{ font-family:'JetBrains Mono', monospace; font-size:16px; font-weight:600; color:#1c1a16; display:block; } |
| .stat-label{ font-family:'JetBrains Mono', monospace; font-size:9px; letter-spacing:.06em; text-transform:uppercase; color:#8a8474; margin-top:3px; display:block; } |
| |
| .transcript{ |
| border-top:1px solid #1c1a16; border-bottom:1px solid #1c1a16; |
| min-height:220px; max-height:52vh; overflow-y:auto; |
| padding:26px 2px; font-size:16px; line-height:1.8; |
| } |
| .msg{ margin-bottom:24px; } |
| .msg:last-child{ margin-bottom:0; } |
| .msg .who{ |
| font-family:'JetBrains Mono', monospace; font-size:9.5px; letter-spacing:.12em; text-transform:uppercase; |
| color:#8a8474; margin-bottom:5px; |
| } |
| .msg.user .who{ color:#ff0000; } |
| .msg .txt{ color:#1c1a16; white-space:pre-wrap; } |
| .msg .txt .cursor{ |
| display:inline-block; width:7px; height:1em; background:#ff0000; vertical-align:text-bottom; |
| animation:blink .85s steps(1) infinite; margin-left:1px; |
| } |
| @keyframes blink{ 50%{ opacity:.2; } } |
| .placeholder-hint{ color:#8a8474; font-size:14px; font-style:italic; } |
| |
| .composer{ display:flex; gap:12px; margin-top:20px; } |
| .composer textarea{ |
| flex:1; resize:none; font-family:'Newsreader', Georgia, serif; font-size:16px; |
| border:none; border-bottom:1px solid #1c1a16; background:transparent; color:#1c1a16; |
| padding:8px 2px; line-height:1.5; height:44px; |
| } |
| .composer textarea:focus{ outline:none; border-bottom-color:#ff0000; } |
| |
| .control-row{ display:flex; align-items:center; justify-content:space-between; gap:16px; margin-top:22px; flex-wrap:wrap; } |
| .sliders{ display:flex; gap:30px; flex-wrap:wrap; } |
| .slider-group{ font-family:'JetBrains Mono', monospace; font-size:10.5px; color:#8a8474; } |
| .slider-group .val{ color:#1c1a16; font-weight:600; } |
| .slider-group input[type=range]{ display:block; width:130px; margin-top:7px; accent-color:#ff0000; } |
| |
| .perf-line{ |
| font-family:'JetBrains Mono', monospace; font-size:11px; color:#8a8474; |
| margin-top:18px; padding-top:14px; border-top:1px solid #ece7d8; |
| display:flex; gap:20px; flex-wrap:wrap; |
| } |
| .perf-line b{ color:#4a463d; font-weight:600; } |
| |
| |
| details.expivme{ margin-top:22px; border-top:1px solid #d9d2bf; padding-top:18px; } |
| details.expivme summary{ |
| cursor:pointer; font-family:'JetBrains Mono', monospace; font-size:12px; letter-spacing:.03em; |
| color:#1c1a16; list-style:none; display:flex; align-items:center; gap:10px; |
| } |
| details.expivme summary::-webkit-details-marker{ display:none; } |
| details.expivme summary::before{ |
| content:"+"; color:#ff0000; font-family:'JetBrains Mono', monospace; font-weight:600; width:14px; |
| } |
| details.expivme[open] summary::before{ content:"–"; } |
| details.expivme pre{ |
| margin:18px 0 0; padding:18px 20px; background:#1c1a16; overflow-x:auto; max-height:340px; overflow-y:auto; |
| } |
| details.expivme pre code{ |
| background:transparent; border:none; color:#e9e4d6; font-size:11.5px; line-height:1.6; padding:0; white-space:pre; |
| } |
| |
| |
| ul.limits{ margin:0; padding:0; list-style:none; } |
| ul.limits li{ |
| font-size:14px; line-height:1.6; color:#4a463d; padding:11px 0 11px 20px; |
| border-bottom:1px solid #ece7d8; position:relative; |
| } |
| ul.limits li:last-child{ border-bottom:none; } |
| ul.limits li::before{ content:"–"; position:absolute; left:0; color:#ff0000; font-weight:600; } |
| |
| |
| footer{ |
| padding:44px 0 10px; display:flex; justify-content:space-between; align-items:center; |
| font-family:'JetBrains Mono', monospace; font-size:11px; color:#8a8474; letter-spacing:.02em; |
| border-top:1px solid #d9d2bf; flex-wrap:wrap; gap:8px; |
| } |
| footer .fmark{ color:#ff0000; } |
| |
| @media (max-width:620px){ |
| .page{ padding:0 18px 60px; } |
| .masthead{ padding:36px 0 0; } |
| .stage{ grid-template-columns:1fr; gap:4px; } |
| .stage .meta{ text-align:left; } |
| .stat-strip{ flex-wrap:wrap; row-gap:16px; } |
| table.kv td:first-child{ width:auto; white-space:normal; } |
| footer{ flex-direction:column; text-align:center; } |
| } |
| </style> |
| </head> |
| <body> |
| <div class="page"> |
|
|
| <div class="masthead"> |
| <div class="eyebrow"><span class="dot-live" id="live-dot"></span> Raw WebGPU · No ONNX · No transformers.js</div> |
| <h1 class="wordmark">İvme-Conversate<span class="dot">.</span></h1> |
| <p class="tagline">Hand-written WGSL kernels reading <code>model.safetensors</code> straight off the wire — the tensor bytes go into GPU buffers directly, no export step, no runtime graph compiler.</p> |
| </div> |
|
|
| <div class="tickline"></div> |
|
|
| <div class="meta-strip"> |
| <div class="meta-item">MODEL<br><b>Conversate-v2-Base</b></div> |
| <div class="meta-item">PARAMS<br><b>23.8M</b></div> |
| <div class="meta-item">RUNTIME<br><b>Browser / WebGPU</b></div> |
| <div class="meta-item">WEIGHTS<br><b><a href="https://huggingface.co/IvmeLabs/Ivme-Conversate-v2-Base" target="_blank" rel="noopener">İvmeLabs</a></b></div> |
| </div> |
|
|
| |
| <section id="boot-section"> |
| <div class="section-head"><span class="section-index">01</span><h2>Load model</h2></div> |
| <p class="section-desc">Fetches <code>config.json</code>, <code>tokenizer.json</code>, and <code>model.safetensors</code> from the Hugging Face CDN, parses the safetensors header by hand, and uploads all 92 tensors straight into <code>GPUBuffer</code>s.</p> |
|
|
| <div class="stage-list" id="stage-list"> |
| <div class="stage" data-stage="webgpu"><span class="mark">idle</span><span class="label">Request WebGPU adapter + device</span><span class="meta"></span></div> |
| <div class="stage" data-stage="config"><span class="mark">idle</span><span class="label">Fetch config.json</span><span class="meta"></span></div> |
| <div class="stage" data-stage="tokenizer"><span class="mark">idle</span><span class="label">Fetch + build tokenizer (byte-level BPE)</span><span class="meta"></span></div> |
| <div class="stage" data-stage="weights"><span class="mark">idle</span><span class="label">Stream model.safetensors</span><span class="meta"></span></div> |
| <div class="stage" data-stage="upload"><span class="mark">idle</span><span class="label">Parse header, upload tensors to GPU</span><span class="meta"></span></div> |
| <div class="stage" data-stage="pipelines"><span class="mark">idle</span><span class="label">Compile WGSL, build compute pipelines</span><span class="meta"></span></div> |
| <div class="stage" data-stage="warmup"><span class="mark">idle</span><span class="label">Warm-up forward pass</span><span class="meta"></span></div> |
| </div> |
|
|
| <div class="progress-track"><div class="progress-fill" id="progress-fill"></div></div> |
| <div class="progress-num"><span id="progress-label">idle</span><span id="progress-bytes"></span></div> |
|
|
| <div class="load-row"> |
| <button id="load-btn">Load model & run in-browser</button> |
| <span class="load-hint">~96MB, cached by the browser after first load</span> |
| </div> |
|
|
| <div class="err-box" id="err-box" style="display:none;"></div> |
| <div class="note-box" id="webgpu-warning" style="display:none;"> |
| This browser doesn't expose <code>navigator.gpu</code>. WebGPU currently ships in Chrome/Edge 113+, and behind a flag in Firefox Nightly and Safari Technology Preview. |
| </div> |
| </section> |
|
|
| |
| <section class="chat-shell" id="chat-shell"> |
| <div class="section-head"><span class="section-index">02</span><h2>Generate</h2></div> |
| <p class="section-desc">Base model — it continues text, it doesn't follow instructions. Try a story opener or a sentence fragment and let it run.</p> |
|
|
| <div class="stat-strip"> |
| <div class="stat"><span class="stat-num" id="stat-params">23.8M</span><span class="stat-label">Parameters</span></div> |
| <div class="stat"><span class="stat-num" id="stat-backend">–</span><span class="stat-label">GPU backend</span></div> |
| <div class="stat"><span class="stat-num" id="stat-loadtime">–</span><span class="stat-label">Load time</span></div> |
| <div class="stat"><span class="stat-num" id="stat-toksec">–</span><span class="stat-label">tok/s (last run)</span></div> |
| </div> |
|
|
| <div class="transcript" id="transcript"> |
| <div class="placeholder-hint">This is a base model — it continues text, it doesn't follow instructions. Try a story opener, a sentence fragment, anything it can run with.</div> |
| </div> |
|
|
| <div class="composer"> |
| <textarea id="prompt-input" placeholder="Once upon a time, there was a" spellcheck="false"></textarea> |
| <button id="gen-btn">Run</button> |
| </div> |
|
|
| <div class="control-row"> |
| <div class="sliders"> |
| <div class="slider-group">temperature <span class="val" id="temp-val">0.80</span> |
| <input type="range" id="temp-slider" min="0.05" max="1.5" step="0.05" value="0.8"> |
| </div> |
| <div class="slider-group">top-k <span class="val" id="topk-val">50</span> |
| <input type="range" id="topk-slider" min="1" max="200" step="1" value="50"> |
| </div> |
| <div class="slider-group">max new tokens <span class="val" id="maxtok-val">200</span> |
| <input type="range" id="maxtok-slider" min="10" max="1024" step="10" value="200"> |
| </div> |
| </div> |
| <button class="ghost" id="stop-btn" disabled>Stop</button> |
| </div> |
|
|
| <div class="perf-line" id="perf-line"></div> |
| </section> |
|
|
| |
| <section> |
| <div class="section-head"><span class="section-index">03</span><h2>What's actually running</h2></div> |
| <table class="kv"> |
| <tr><td>Architecture</td><td>Dense decoder-only, 10 layers, hidden 384, 6 heads (head_dim 64), FFN inner 1024, tied embeddings</td></tr> |
| <tr><td>Positional encoding</td><td>RoPE, θ=10000, applied as complex-plane rotation on interleaved (x0,x1),(x2,x3)… pairs — matched exactly, not the "rotate-half" GPT-NeoX layout</td></tr> |
| <tr><td>Normalization</td><td>RMSNorm, pre-norm, eps=1e-5, computed in f32</td></tr> |
| <tr><td>Attention</td><td>Full scaled dot-product attention, single flash-style kernel per head with an online softmax and a causal early-exit on the key loop</td></tr> |
| <tr><td>Feedforward</td><td>SwiGLU: down_proj(silu(gate_proj(x)) ⊙ up_proj(x)), inner dim 1024</td></tr> |
| <tr><td>Sampling</td><td>Temperature scaling + top-k, categorical draw via cumulative sum on CPU over the top-k slice returned from GPU</td></tr> |
| <tr><td>Weight loading</td><td>Raw <code>fetch()</code> of <code>model.safetensors</code>, JSON header parsed by hand, each tensor's byte range sliced from the <code>ArrayBuffer</code> and copied into a <code>GPUBuffer</code> via <code>writeBuffer</code> — no ONNX export, no ORT runtime, no transformers.js</td></tr> |
| </table> |
|
|
| <details class="expivme"> |
| <summary>View the WGSL kernel source (matmul, RMSNorm, RoPE, attention, SwiGLU, sampling)</summary> |
| <pre><code id="kernel-source-view"></code></pre> |
| </details> |
| </section> |
|
|
| <section> |
| <div class="section-head"><span class="section-index">04</span><h2>Limitations of this demo</h2></div> |
| <ul class="limits"> |
| <li>Base model only — it completes text, it does not follow instructions or answer questions</li> |
| <li>1024-token context window, enforced by refusing to generate past it</li> |
| <li>All matmuls run in f32 with a naive tiled kernel — no flash-attention-grade fusion, no int8/int4 quantization, so this favors correctness and readability over raw throughput</li> |
| <li>KV cache lives in a fixed-size GPU buffer sized to <code>context_len</code>; no paging, no batching, single sequence only</li> |
| <li>Tokenizer is a from-scratch byte-level BPE implementation of the GPT-2 scheme — matches <code>tokenizer.json</code>'s vocab and merges exactly, tested against the reference tokenizer output</li> |
| </ul> |
| </section> |
|
|
| <footer> |
| <span>İvmeLabs <span class="fmark">/</span> Conversate <span class="fmark">/</span> WebGPU</span> |
| <span>Weights: <a href="https://huggingface.co/IvmeLabs/Ivme-Conversate-v2-Base" target="_blank" rel="noopener">IvmeLabs/Ivme-Conversate-v2-Base</a></span> |
| </footer> |
|
|
| </div> |
| <script type="module"> |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| |
| const REPO = "IvmeLabs/Ivme-Conversate-v2-Base"; |
| const CDN = `https://huggingface.co/${REPO}/resolve/main`; |
| |
| |
| |
| |
| |
| |
| const WGSL = {}; |
| |
| |
| |
| WGSL.rmsnorm = ` |
| struct Dims { rows: u32, dim: u32, eps: f32, _pad: u32 }; |
| @group(0) @binding(0) var<uniform> dims: Dims; |
| @group(0) @binding(1) var<storage, read> x: array<f32>; |
| @group(0) @binding(2) var<storage, read> weight: array<f32>; |
| @group(0) @binding(3) var<storage, read_write> y: array<f32>; |
| |
| var<workgroup> partial: array<f32, 256>; |
| |
| @compute @workgroup_size(256) |
| fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) { |
| let row = wg.x; |
| if (row >= dims.rows) { return; } |
| let base = row * dims.dim; |
| |
| var acc: f32 = 0.0; |
| var i: u32 = lid.x; |
| loop { |
| if (i >= dims.dim) { break; } |
| let v = x[base + i]; |
| acc = acc + v * v; |
| i = i + 256u; |
| } |
| partial[lid.x] = acc; |
| workgroupBarrier(); |
| |
| var stride: u32 = 128u; |
| loop { |
| if (stride == 0u) { break; } |
| if (lid.x < stride) { partial[lid.x] = partial[lid.x] + partial[lid.x + stride]; } |
| workgroupBarrier(); |
| stride = stride / 2u; |
| } |
| |
| let meanSq = partial[0] / f32(dims.dim); |
| let rms = inverseSqrt(meanSq + dims.eps); |
| |
| i = lid.x; |
| loop { |
| if (i >= dims.dim) { break; } |
| y[base + i] = x[base + i] * rms * weight[i]; |
| i = i + 256u; |
| } |
| } |
| `; |
| |
| |
| WGSL.matmul = ` |
| struct Dims { rows: u32, K: u32, N: u32, addBias: u32 }; |
| @group(0) @binding(0) var<uniform> dims: Dims; |
| @group(0) @binding(1) var<storage, read> X: array<f32>; // [rows, K] |
| @group(0) @binding(2) var<storage, read> W: array<f32>; // [N, K] (nn.Linear weight, row-major) |
| @group(0) @binding(3) var<storage, read_write> Y: array<f32>; // [rows, N] |
| |
| const TILE = 16u; |
| var<workgroup> Xs: array<array<f32, TILE>, TILE>; |
| var<workgroup> Ws: array<array<f32, TILE>, TILE>; |
| |
| @compute @workgroup_size(TILE, TILE) |
| fn main(@builtin(global_invocation_id) gid: vec3<u32>, |
| @builtin(local_invocation_id) lid: vec3<u32>) { |
| let row = gid.y; // output row (token) |
| let col = gid.x; // output col (feature n) |
| |
| var acc: f32 = 0.0; |
| let numTiles = (dims.K + TILE - 1u) / TILE; |
| |
| for (var t: u32 = 0u; t < numTiles; t = t + 1u) { |
| let kX = t * TILE + lid.x; |
| if (row < dims.rows && kX < dims.K) { |
| Xs[lid.y][lid.x] = X[row * dims.K + kX]; |
| } else { |
| Xs[lid.y][lid.x] = 0.0; |
| } |
| |
| let kW = t * TILE + lid.y; |
| if (col < dims.N && kW < dims.K) { |
| // W is [N, K] row-major -> W[col, kW] |
| Ws[lid.y][lid.x] = W[col * dims.K + kW]; |
| } else { |
| Ws[lid.y][lid.x] = 0.0; |
| } |
| |
| workgroupBarrier(); |
| |
| for (var k: u32 = 0u; k < TILE; k = k + 1u) { |
| acc = acc + Xs[lid.y][k] * Ws[k][lid.x]; |
| } |
| workgroupBarrier(); |
| } |
| |
| if (row < dims.rows && col < dims.N) { |
| Y[row * dims.N + col] = acc; |
| } |
| } |
| `; |
| |
| |
| |
| |
| WGSL.rope = ` |
| struct Dims { T: u32, H: u32, headDim: u32, posOffset: u32 }; |
| @group(0) @binding(0) var<uniform> dims: Dims; |
| @group(0) @binding(1) var<storage, read_write> qk: array<f32>; // [T, H, headDim] |
| @group(0) @binding(2) var<storage, read> invFreq: array<f32>; // [headDim/2] |
| |
| @compute @workgroup_size(8, 8) |
| fn main(@builtin(global_invocation_id) gid: vec3<u32>) { |
| let t = gid.x; // token index within this dispatch chunk |
| let h = gid.y; // head index |
| if (t >= dims.T || h >= dims.H) { return; } |
| |
| let pos = f32(t + dims.posOffset); |
| let half = dims.headDim / 2u; |
| let base = (t * dims.H + h) * dims.headDim; |
| |
| for (var i: u32 = 0u; i < half; i = i + 1u) { |
| let angle = pos * invFreq[i]; |
| let c = cos(angle); |
| let s = sin(angle); |
| let x0 = qk[base + 2u * i]; |
| let x1 = qk[base + 2u * i + 1u]; |
| qk[base + 2u * i] = x0 * c - x1 * s; |
| qk[base + 2u * i + 1u] = x0 * s + x1 * c; |
| } |
| } |
| `; |
| |
| |
| |
| |
| WGSL.attention = ` |
| struct Dims { qLen: u32, kvLen: u32, H: u32, headDim: u32, scale: f32, queryOffset: u32 }; |
| @group(0) @binding(0) var<uniform> dims: Dims; |
| @group(0) @binding(1) var<storage, read> Q: array<f32>; // [qLen, H, headDim] |
| @group(0) @binding(2) var<storage, read> Kc: array<f32>; // [kvLen, H, headDim] cache |
| @group(0) @binding(3) var<storage, read> Vc: array<f32>; // [kvLen, H, headDim] cache |
| @group(0) @binding(4) var<storage, read_write> O: array<f32>; // [qLen, H, headDim] |
| |
| const MAX_HD = 128u; |
| |
| @compute @workgroup_size(64) |
| fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) { |
| let qi = wg.x; // query token index (0..qLen) |
| let h = wg.y; // head index |
| if (qi >= dims.qLen || h >= dims.H) { return; } |
| |
| let qBase = (qi * dims.H + h) * dims.headDim; |
| var qvec: array<f32, MAX_HD>; |
| for (var d: u32 = 0u; d < dims.headDim; d = d + 1u) { |
| qvec[d] = Q[qBase + d]; |
| } |
| |
| // causal boundary: this query (absolute position queryOffset+qi) may attend |
| // to all cached keys up through its own absolute position. |
| let selfPos = dims.queryOffset + qi; |
| var validLen = dims.kvLen; |
| if (selfPos + 1u < validLen) { validLen = selfPos + 1u; } |
| |
| var m: f32 = -3.4e38; // running max |
| var l: f32 = 0.0; // running normalizer |
| var acc: array<f32, MAX_HD>; |
| for (var d: u32 = 0u; d < dims.headDim; d = d + 1u) { acc[d] = 0.0; } |
| |
| var ki: u32 = 0u; |
| loop { |
| if (ki >= validLen) { break; } |
| let kBase = (ki * dims.H + h) * dims.headDim; |
| var score: f32 = 0.0; |
| for (var d: u32 = 0u; d < dims.headDim; d = d + 1u) { |
| score = score + qvec[d] * Kc[kBase + d]; |
| } |
| score = score * dims.scale; |
| |
| let newM = max(m, score); |
| let corr = exp(m - newM); |
| let p = exp(score - newM); |
| |
| l = l * corr + p; |
| for (var d: u32 = 0u; d < dims.headDim; d = d + 1u) { |
| acc[d] = acc[d] * corr + p * Vc[kBase + d]; |
| } |
| m = newM; |
| ki = ki + 1u; |
| } |
| |
| let oBase = (qi * dims.H + h) * dims.headDim; |
| let invL = 1.0 / max(l, 1e-20); |
| for (var d: u32 = 0u; d < dims.headDim; d = d + 1u) { |
| O[oBase + d] = acc[d] * invL; |
| } |
| } |
| `; |
| |
| |
| WGSL.swiglu_gate = ` |
| @group(0) @binding(0) var<storage, read> gate: array<f32>; |
| @group(0) @binding(1) var<storage, read> up: array<f32>; |
| @group(0) @binding(2) var<storage, read_write> out: array<f32>; |
| |
| @compute @workgroup_size(256) |
| fn main(@builtin(global_invocation_id) gid: vec3<u32>) { |
| let i = gid.x; |
| if (i >= arrayLength(&gate)) { return; } |
| let g = gate[i]; |
| let silu = g / (1.0 + exp(-g)); |
| out[i] = silu * up[i]; |
| } |
| `; |
| |
| |
| WGSL.add = ` |
| @group(0) @binding(0) var<storage, read> a: array<f32>; |
| @group(0) @binding(1) var<storage, read> b: array<f32>; |
| @group(0) @binding(2) var<storage, read_write> out: array<f32>; |
| |
| @compute @workgroup_size(256) |
| fn main(@builtin(global_invocation_id) gid: vec3<u32>) { |
| let i = gid.x; |
| if (i >= arrayLength(&a)) { return; } |
| out[i] = a[i] + b[i]; |
| } |
| `; |
| |
| |
| |
| |
| WGSL.kvwrite = ` |
| struct Dims { T: u32, H: u32, headDim: u32, destOffset: u32 }; |
| @group(0) @binding(0) var<uniform> dims: Dims; |
| @group(0) @binding(1) var<storage, read> src: array<f32>; // [T, H, headDim] |
| @group(0) @binding(2) var<storage, read_write> cache: array<f32>; // [context_len, H, headDim] |
| |
| @compute @workgroup_size(256) |
| fn main(@builtin(global_invocation_id) gid: vec3<u32>) { |
| let i = gid.x; |
| let total = dims.T * dims.H * dims.headDim; |
| if (i >= total) { return; } |
| |
| let perTok = dims.H * dims.headDim; |
| let t = i / perTok; |
| let rem = i % perTok; |
| let destBase = (dims.destOffset + t) * perTok; |
| cache[destBase + rem] = src[i]; |
| } |
| `; |
| |
| |
| |
| |
| |
| |
| const KERNEL_ORDER = ["rmsnorm","matmul","rope","attention","swiglu_gate","add","kvwrite"]; |
| |
| |
| |
| |
| |
| class ByteLevelBPE { |
| constructor(tokJson) { |
| this.vocab = tokJson.model.vocab; |
| this.idToToken = new Array(Object.keys(this.vocab).length); |
| for (const [tok, id] of Object.entries(this.vocab)) this.idToToken[id] = tok; |
| |
| this.merges = tokJson.model.merges.map(pair => |
| Array.isArray(pair) ? pair : pair.split(" ") |
| ); |
| this.mergeRank = new Map(); |
| this.merges.forEach(([a, b], i) => this.mergeRank.set(a + "\u0001" + b, i)); |
| |
| this.specials = new Map(); |
| for (const t of (tokJson.added_tokens || [])) this.specials.set(t.content, t.id); |
| |
| this.byteEncoder = ByteLevelBPE._buildByteEncoder(); |
| this.byteDecoder = new Map(); |
| for (const [b, ch] of this.byteEncoder.entries()) this.byteDecoder.set(ch, b); |
| |
| |
| this.splitRe = /'s|'t|'re|'ve|'m|'ll|'d| ?\p{L}+| ?\p{N}+| ?[^\s\p{L}\p{N}]+|\s+(?!\S)|\s+/gu; |
| |
| this._cache = new Map(); |
| } |
| |
| static _buildByteEncoder() { |
| const bs = []; |
| for (let i = "!".charCodeAt(0); i <= "~".charCodeAt(0); i++) bs.push(i); |
| for (let i = "\u00A1".charCodeAt(0); i <= "\u00AC".charCodeAt(0); i++) bs.push(i); |
| for (let i = "\u00AE".charCodeAt(0); i <= "\u00FF".charCodeAt(0); i++) bs.push(i); |
| const cs = bs.slice(); |
| let n = 0; |
| for (let b = 0; b < 256; b++) { |
| if (!bs.includes(b)) { |
| bs.push(b); |
| cs.push(256 + n); |
| n++; |
| } |
| } |
| const map = new Map(); |
| for (let i = 0; i < bs.length; i++) map.set(bs[i], String.fromCodePoint(cs[i])); |
| return map; |
| } |
| |
| _bpe(token) { |
| if (this._cache.has(token)) return this._cache.get(token); |
| let word = Array.from(token); |
| if (word.length === 1) { this._cache.set(token, word); return word; } |
| |
| while (true) { |
| let bestRank = Infinity, bestIdx = -1; |
| for (let i = 0; i < word.length - 1; i++) { |
| const key = word[i] + "\u0001" + word[i + 1]; |
| const r = this.mergeRank.get(key); |
| if (r !== undefined && r < bestRank) { bestRank = r; bestIdx = i; } |
| } |
| if (bestIdx === -1) break; |
| const merged = word[bestIdx] + word[bestIdx + 1]; |
| word = [...word.slice(0, bestIdx), merged, ...word.slice(bestIdx + 2)]; |
| } |
| this._cache.set(token, word); |
| return word; |
| } |
| |
| encode(text) { |
| const ids = []; |
| const chunks = text.match(this.splitRe) || []; |
| for (const chunk of chunks) { |
| const bytes = new TextEncoder().encode(chunk); |
| let mapped = ""; |
| for (const b of bytes) mapped += this.byteEncoder.get(b); |
| const pieces = this._bpe(mapped); |
| for (const p of pieces) { |
| const id = this.vocab[p]; |
| if (id !== undefined) ids.push(id); |
| else { |
| |
| const unk = this.specials.get("<|unk|>"); |
| if (unk !== undefined) ids.push(unk); |
| } |
| } |
| } |
| return ids; |
| } |
| |
| decode(ids) { |
| let mapped = ""; |
| for (const id of ids) { |
| const tok = this.idToToken[id]; |
| if (tok === undefined) continue; |
| if (this.specials.has(tok)) continue; |
| mapped += tok; |
| } |
| const bytes = []; |
| for (const ch of mapped) { |
| const b = this.byteDecoder.get(ch); |
| if (b !== undefined) bytes.push(b); |
| } |
| return new TextDecoder("utf-8", { fatal: false }).decode(new Uint8Array(bytes)); |
| } |
| } |
| |
| |
| |
| |
| |
| |
| function parseSafetensors(buffer) { |
| const view = new DataView(buffer); |
| const headerLen = Number(view.getBigUint64(0, true)); |
| const headerBytes = new Uint8Array(buffer, 8, headerLen); |
| const header = JSON.parse(new TextDecoder("utf-8").decode(headerBytes)); |
| const dataStart = 8 + headerLen; |
| |
| const tensors = {}; |
| for (const [name, info] of Object.entries(header)) { |
| if (name === "__metadata__") continue; |
| const [begin, end] = info.data_offsets; |
| tensors[name] = { |
| dtype: info.dtype, |
| shape: info.shape, |
| byteOffset: dataStart + begin, |
| byteLength: end - begin, |
| }; |
| } |
| return { tensors, buffer }; |
| } |
| |
| |
| |
| |
| |
| class IvmeEngine { |
| constructor(device) { |
| this.device = device; |
| this.pipelines = {}; |
| this.weights = {}; |
| this.cfg = null; |
| this.tokenizer = null; |
| this.kvCacheK = []; |
| this.kvCacheV = []; |
| this.curLen = 0; |
| this.stopFlag = false; |
| } |
| |
| async buildPipelines() { |
| const mk = (code) => { |
| const module = this.device.createShaderModule({ code }); |
| return this.device.createComputePipeline({ layout: "auto", compute: { module, entryPoint: "main" } }); |
| }; |
| for (const name of KERNEL_ORDER) { |
| this.pipelines[name] = mk(WGSL[name]); |
| } |
| } |
| |
| loadWeightsFromSafetensors(parsed) { |
| const { tensors, buffer } = parsed; |
| for (const [name, t] of Object.entries(tensors)) { |
| if (t.dtype !== "F32") { |
| throw new Error(`unexpected dtype ${t.dtype} for ${name}; this loader only handles F32 safetensors`); |
| } |
| const bytes = buffer.slice(t.byteOffset, t.byteOffset + t.byteLength); |
| const gbuf = this.device.createBuffer({ |
| size: Math.ceil(bytes.byteLength / 4) * 4, |
| usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, |
| mappedAtCreation: true, |
| }); |
| new Uint8Array(gbuf.getMappedRange()).set(new Uint8Array(bytes)); |
| gbuf.unmap(); |
| this.weights[name] = { buffer: gbuf, shape: t.shape }; |
| } |
| } |
| |
| _uniform(dataArray) { |
| const buf = this.device.createBuffer({ |
| size: Math.max(16, dataArray.byteLength), |
| usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, |
| }); |
| this.device.queue.writeBuffer(buf, 0, dataArray); |
| return buf; |
| } |
| |
| _storage(size, extraUsage = 0) { |
| return this.device.createBuffer({ |
| size: Math.ceil(size / 4) * 4, |
| usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC | extraUsage, |
| }); |
| } |
| |
| initKVCache() { |
| const { n_layers, context_len, hidden_dim, n_heads } = this.cfg; |
| const headDim = hidden_dim / n_heads; |
| const bytesPerLayer = context_len * n_heads * headDim * 4; |
| this.kvCacheK = []; |
| this.kvCacheV = []; |
| for (let l = 0; l < n_layers; l++) { |
| this.kvCacheK.push(this._storage(bytesPerLayer)); |
| this.kvCacheV.push(this._storage(bytesPerLayer)); |
| } |
| this.curLen = 0; |
| } |
| |
| _dispatchMatmul(enc, X, rows, K, W, N, Y) { |
| const dims = this._uniform(new Uint32Array([rows, K, N, 0])); |
| const bg = this.device.createBindGroup({ |
| layout: this.pipelines.matmul.getBindGroupLayout(0), |
| entries: [ |
| { binding: 0, resource: { buffer: dims } }, |
| { binding: 1, resource: { buffer: X } }, |
| { binding: 2, resource: { buffer: W } }, |
| { binding: 3, resource: { buffer: Y } }, |
| ], |
| }); |
| const pass = enc.beginComputePass(); |
| pass.setPipeline(this.pipelines.matmul); |
| pass.setBindGroup(0, bg); |
| pass.dispatchWorkgroups(Math.ceil(N / 16), Math.ceil(rows / 16)); |
| pass.end(); |
| } |
| |
| _dispatchRMSNorm(enc, X, rows, dim, weightBuf, Y) { |
| const dims = this._uniform(new Uint32Array([rows, dim, 0, 0])); |
| |
| const dimsF = new ArrayBuffer(16); |
| const dv = new DataView(dimsF); |
| dv.setUint32(0, rows, true); |
| dv.setUint32(4, dim, true); |
| dv.setFloat32(8, this.cfg.norm_eps, true); |
| dv.setUint32(12, 0, true); |
| const dimsBuf = this._uniform(new Uint8Array(dimsF)); |
| const bg = this.device.createBindGroup({ |
| layout: this.pipelines.rmsnorm.getBindGroupLayout(0), |
| entries: [ |
| { binding: 0, resource: { buffer: dimsBuf } }, |
| { binding: 1, resource: { buffer: X } }, |
| { binding: 2, resource: { buffer: weightBuf } }, |
| { binding: 3, resource: { buffer: Y } }, |
| ], |
| }); |
| const pass = enc.beginComputePass(); |
| pass.setPipeline(this.pipelines.rmsnorm); |
| pass.setBindGroup(0, bg); |
| pass.dispatchWorkgroups(rows); |
| pass.end(); |
| } |
| |
| _dispatchRoPE(enc, qkBuf, T, H, headDim, posOffset, invFreqBuf) { |
| const dims = this._uniform(new Uint32Array([T, H, headDim, posOffset])); |
| const bg = this.device.createBindGroup({ |
| layout: this.pipelines.rope.getBindGroupLayout(0), |
| entries: [ |
| { binding: 0, resource: { buffer: dims } }, |
| { binding: 1, resource: { buffer: qkBuf } }, |
| { binding: 2, resource: { buffer: invFreqBuf } }, |
| ], |
| }); |
| const pass = enc.beginComputePass(); |
| pass.setPipeline(this.pipelines.rope); |
| pass.setBindGroup(0, bg); |
| pass.dispatchWorkgroups(Math.ceil(T / 8), Math.ceil(H / 8)); |
| pass.end(); |
| } |
| |
| _dispatchKVWrite(enc, srcBuf, T, H, headDim, destOffset, cacheBuf) { |
| const dims = this._uniform(new Uint32Array([T, H, headDim, destOffset])); |
| const bg = this.device.createBindGroup({ |
| layout: this.pipelines.kvwrite.getBindGroupLayout(0), |
| entries: [ |
| { binding: 0, resource: { buffer: dims } }, |
| { binding: 1, resource: { buffer: srcBuf } }, |
| { binding: 2, resource: { buffer: cacheBuf } }, |
| ], |
| }); |
| const pass = enc.beginComputePass(); |
| pass.setPipeline(this.pipelines.kvwrite); |
| pass.setBindGroup(0, bg); |
| pass.dispatchWorkgroups(Math.ceil((T * H * headDim) / 256)); |
| pass.end(); |
| } |
| |
| _dispatchAttention(enc, Q, kvLen, H, headDim, queryOffset, Kc, Vc, O, qLen) { |
| const dimsF = new ArrayBuffer(24); |
| const dv = new DataView(dimsF); |
| dv.setUint32(0, qLen, true); |
| dv.setUint32(4, kvLen, true); |
| dv.setUint32(8, H, true); |
| dv.setUint32(12, headDim, true); |
| dv.setFloat32(16, 1.0 / Math.sqrt(headDim), true); |
| dv.setUint32(20, queryOffset, true); |
| const dimsBuf = this._uniform(new Uint8Array(dimsF)); |
| const bg = this.device.createBindGroup({ |
| layout: this.pipelines.attention.getBindGroupLayout(0), |
| entries: [ |
| { binding: 0, resource: { buffer: dimsBuf } }, |
| { binding: 1, resource: { buffer: Q } }, |
| { binding: 2, resource: { buffer: Kc } }, |
| { binding: 3, resource: { buffer: Vc } }, |
| { binding: 4, resource: { buffer: O } }, |
| ], |
| }); |
| const pass = enc.beginComputePass(); |
| pass.setPipeline(this.pipelines.attention); |
| pass.setBindGroup(0, bg); |
| pass.dispatchWorkgroups(qLen, H); |
| pass.end(); |
| } |
| |
| _dispatchElemwise(enc, pipelineName, a, b, out, count) { |
| const bg = this.device.createBindGroup({ |
| layout: this.pipelines[pipelineName].getBindGroupLayout(0), |
| entries: [ |
| { binding: 0, resource: { buffer: a } }, |
| { binding: 1, resource: { buffer: b } }, |
| { binding: 2, resource: { buffer: out } }, |
| ], |
| }); |
| const pass = enc.beginComputePass(); |
| pass.setPipeline(this.pipelines[pipelineName]); |
| pass.setBindGroup(0, bg); |
| pass.dispatchWorkgroups(Math.ceil(count / 256)); |
| pass.end(); |
| } |
| |
| |
| _buildInvFreq() { |
| const headDim = this.cfg.hidden_dim / this.cfg.n_heads; |
| const half = headDim / 2; |
| const arr = new Float32Array(half); |
| for (let i = 0; i < half; i++) { |
| arr[i] = 1.0 / Math.pow(this.cfg.rope_theta, (2 * i) / headDim); |
| } |
| const buf = this.device.createBuffer({ |
| size: arr.byteLength, |
| usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, |
| mappedAtCreation: true, |
| }); |
| new Float32Array(buf.getMappedRange()).set(arr); |
| buf.unmap(); |
| this.invFreqBuf = buf; |
| } |
| |
| |
| |
| |
| |
| |
| async forward(tokenIds) { |
| const { hidden_dim, n_heads, n_layers, vocab_size, context_len } = this.cfg; |
| const headDim = hidden_dim / n_heads; |
| const T = tokenIds.length; |
| const posOffset = this.curLen; |
| |
| if (posOffset + T > context_len) { |
| throw new Error(`sequence length ${posOffset + T} exceeds context_len ${context_len}`); |
| } |
| |
| const enc = this.device.createCommandEncoder(); |
| |
| |
| const embedWeight = this.weights["model.tok_embed.weight"]; |
| |
| |
| const hidden = this._storage(T * hidden_dim * 4); |
| { |
| const embedBytes = this._embedRaw; |
| const rowBuf = new Float32Array(T * hidden_dim); |
| for (let t = 0; t < T; t++) { |
| const id = tokenIds[t]; |
| rowBuf.set(embedBytes.subarray(id * hidden_dim, (id + 1) * hidden_dim), t * hidden_dim); |
| } |
| this.device.queue.writeBuffer(hidden, 0, rowBuf); |
| } |
| |
| let x = hidden; |
| |
| for (let l = 0; l < n_layers; l++) { |
| const p = `model.blocks.${l}.`; |
| |
| |
| const normed = this._storage(T * hidden_dim * 4); |
| this._dispatchRMSNorm(enc, x, T, hidden_dim, this.weights[p + "attn_norm.weight"].buffer, normed); |
| |
| |
| const q = this._storage(T * hidden_dim * 4); |
| const k = this._storage(T * hidden_dim * 4); |
| const v = this._storage(T * hidden_dim * 4); |
| this._dispatchMatmul(enc, normed, T, hidden_dim, this.weights[p + "attn.q_proj.weight"].buffer, hidden_dim, q); |
| this._dispatchMatmul(enc, normed, T, hidden_dim, this.weights[p + "attn.k_proj.weight"].buffer, hidden_dim, k); |
| this._dispatchMatmul(enc, normed, T, hidden_dim, this.weights[p + "attn.v_proj.weight"].buffer, hidden_dim, v); |
| |
| |
| |
| this._dispatchRoPE(enc, q, T, n_heads, headDim, posOffset, this.invFreqBuf); |
| this._dispatchRoPE(enc, k, T, n_heads, headDim, posOffset, this.invFreqBuf); |
| |
| |
| this._dispatchKVWrite(enc, k, T, n_heads, headDim, posOffset, this.kvCacheK[l]); |
| this._dispatchKVWrite(enc, v, T, n_heads, headDim, posOffset, this.kvCacheV[l]); |
| |
| |
| const attnOut = this._storage(T * hidden_dim * 4); |
| this._dispatchAttention(enc, q, posOffset + T, n_heads, headDim, posOffset, this.kvCacheK[l], this.kvCacheV[l], attnOut, T); |
| |
| |
| const attnProj = this._storage(T * hidden_dim * 4); |
| this._dispatchMatmul(enc, attnOut, T, hidden_dim, this.weights[p + "attn.out_proj.weight"].buffer, hidden_dim, attnProj); |
| |
| |
| const afterAttn = this._storage(T * hidden_dim * 4); |
| this._dispatchElemwise(enc, "add", x, attnProj, afterAttn, T * hidden_dim); |
| |
| |
| const ffnNormed = this._storage(T * hidden_dim * 4); |
| this._dispatchRMSNorm(enc, afterAttn, T, hidden_dim, this.weights[p + "ffn_norm.weight"].buffer, ffnNormed); |
| |
| |
| const ffnInner = this.weights[p + "ffn.gate_proj.weight"].shape[0]; |
| const gate = this._storage(T * ffnInner * 4); |
| const up = this._storage(T * ffnInner * 4); |
| this._dispatchMatmul(enc, ffnNormed, T, hidden_dim, this.weights[p + "ffn.gate_proj.weight"].buffer, ffnInner, gate); |
| this._dispatchMatmul(enc, ffnNormed, T, hidden_dim, this.weights[p + "ffn.up_proj.weight"].buffer, ffnInner, up); |
| const gated = this._storage(T * ffnInner * 4); |
| this._dispatchElemwise(enc, "swiglu_gate", gate, up, gated, T * ffnInner); |
| const ffnOut = this._storage(T * hidden_dim * 4); |
| this._dispatchMatmul(enc, gated, T, ffnInner, this.weights[p + "ffn.down_proj.weight"].buffer, hidden_dim, ffnOut); |
| |
| |
| const afterFFN = this._storage(T * hidden_dim * 4); |
| this._dispatchElemwise(enc, "add", afterAttn, ffnOut, afterFFN, T * hidden_dim); |
| |
| x = afterFFN; |
| } |
| |
| |
| const finalNormed = this._storage(T * hidden_dim * 4); |
| this._dispatchRMSNorm(enc, x, T, hidden_dim, this.weights["model.final_norm.weight"].buffer, finalNormed); |
| |
| |
| const lastHidden = this._storage(hidden_dim * 4); |
| enc.copyBufferToBuffer(finalNormed, (T - 1) * hidden_dim * 4, lastHidden, 0, hidden_dim * 4); |
| |
| const logits = this._storage(vocab_size * 4, GPUBufferUsage.COPY_SRC); |
| this._dispatchMatmul(enc, lastHidden, 1, hidden_dim, embedWeight.buffer, vocab_size, logits); |
| |
| |
| const readBuf = this.device.createBuffer({ |
| size: vocab_size * 4, |
| usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, |
| }); |
| enc.copyBufferToBuffer(logits, 0, readBuf, 0, vocab_size * 4); |
| |
| this.device.queue.submit([enc.finish()]); |
| |
| await readBuf.mapAsync(GPUMapMode.READ); |
| const out = new Float32Array(readBuf.getMappedRange().slice(0)); |
| readBuf.unmap(); |
| |
| this.curLen = posOffset + T; |
| return out; |
| } |
| |
| sample(logits, temperature, topK) { |
| const n = logits.length; |
| const scaled = new Float32Array(n); |
| const invT = 1.0 / Math.max(temperature, 1e-6); |
| for (let i = 0; i < n; i++) scaled[i] = logits[i] * invT; |
| |
| |
| const idx = Array.from({ length: n }, (_, i) => i); |
| idx.sort((a, b) => scaled[b] - scaled[a]); |
| const top = idx.slice(0, Math.min(topK, n)); |
| |
| const maxLogit = scaled[top[0]]; |
| let sum = 0; |
| const probs = top.map(i => { |
| const p = Math.exp(scaled[i] - maxLogit); |
| sum += p; |
| return p; |
| }); |
| let r = Math.random() * sum; |
| for (let j = 0; j < top.length; j++) { |
| r -= probs[j]; |
| if (r <= 0) return top[j]; |
| } |
| return top[top.length - 1]; |
| } |
| } |
| |
| |
| |
| |
| |
| const stageListEl = document.getElementById("stage-list"); |
| const progressFill = document.getElementById("progress-fill"); |
| const progressLabel = document.getElementById("progress-label"); |
| const progressBytes = document.getElementById("progress-bytes"); |
| const loadBtn = document.getElementById("load-btn"); |
| const errBox = document.getElementById("err-box"); |
| const webgpuWarning = document.getElementById("webgpu-warning"); |
| const chatShell = document.getElementById("chat-shell"); |
| const bootSection = document.getElementById("boot-section"); |
| const liveDot = document.getElementById("live-dot"); |
| |
| function setStage(name, state, meta = "") { |
| const el = stageListEl.querySelector(`[data-stage="${name}"]`); |
| if (!el) return; |
| el.classList.remove("active", "done"); |
| if (state === "active") { el.classList.add("active"); el.querySelector(".mark").textContent = "running"; } |
| if (state === "done") { el.classList.add("done"); el.querySelector(".mark").textContent = "done"; } |
| if (state === "idle") { el.querySelector(".mark").textContent = "idle"; } |
| el.querySelector(".meta").textContent = meta; |
| } |
| |
| function setProgress(frac, label, bytesLabel = "") { |
| progressFill.style.width = `${Math.round(frac * 100)}%`; |
| progressLabel.textContent = label; |
| progressBytes.textContent = bytesLabel; |
| } |
| |
| function showError(err) { |
| console.error(err); |
| errBox.style.display = "block"; |
| errBox.textContent = (err && err.message) ? err.message : String(err); |
| } |
| |
| function fmtBytes(n) { |
| if (n > 1e6) return (n / 1e6).toFixed(1) + " MB"; |
| if (n > 1e3) return (n / 1e3).toFixed(1) + " KB"; |
| return n + " B"; |
| } |
| |
| async function fetchWithProgress(url, onProgress) { |
| const resp = await fetch(url); |
| if (!resp.ok) throw new Error(`fetch failed: ${resp.status} ${resp.statusText} — ${url}`); |
| const total = Number(resp.headers.get("content-length")) || 0; |
| const reader = resp.body.getReader(); |
| const chunks = []; |
| let received = 0; |
| while (true) { |
| const { done, value } = await reader.read(); |
| if (done) break; |
| chunks.push(value); |
| received += value.length; |
| if (onProgress) onProgress(received, total); |
| } |
| const buf = new Uint8Array(received); |
| let offset = 0; |
| for (const c of chunks) { buf.set(c, offset); offset += c.length; } |
| return buf.buffer; |
| } |
| |
| let engine = null; |
| let genController = { stop: false }; |
| |
| async function boot() { |
| errBox.style.display = "none"; |
| loadBtn.disabled = true; |
| const t0 = performance.now(); |
| |
| try { |
| if (!("gpu" in navigator)) { |
| webgpuWarning.style.display = "block"; |
| throw new Error("navigator.gpu is undefined — WebGPU is not available in this browser."); |
| } |
| |
| setStage("webgpu", "active"); |
| const adapter = await navigator.gpu.requestAdapter(); |
| if (!adapter) throw new Error("navigator.gpu.requestAdapter() returned null — no compatible GPU found."); |
| const device = await adapter.requestDevice(); |
| let backendLabel = "WebGPU"; |
| try { |
| const info = adapter.info || (adapter.requestAdapterInfo ? await adapter.requestAdapterInfo() : null); |
| if (info && (info.vendor || info.architecture)) { |
| backendLabel = [info.vendor, info.architecture].filter(Boolean).join(" / "); |
| } |
| } catch (_) {} |
| document.getElementById("stat-backend").textContent = backendLabel || "WebGPU"; |
| setStage("webgpu", "done", "device acquired"); |
| |
| engine = new IvmeEngine(device); |
| |
| setStage("config", "active"); |
| const cfgResp = await fetch(`${CDN}/config.json`); |
| if (!cfgResp.ok) throw new Error(`config.json fetch failed: ${cfgResp.status}`); |
| engine.cfg = await cfgResp.json(); |
| setStage("config", "done", `${engine.cfg.n_layers}L / d${engine.cfg.hidden_dim}`); |
| |
| setStage("tokenizer", "active"); |
| const tokBuf = await fetchWithProgress(`${CDN}/tokenizer.json`, (r, t) => { |
| setProgress(t ? r / t : 0, "fetching tokenizer.json", `${fmtBytes(r)}${t ? " / " + fmtBytes(t) : ""}`); |
| }); |
| const tokJson = JSON.parse(new TextDecoder().decode(tokBuf)); |
| engine.tokenizer = new ByteLevelBPE(tokJson); |
| setStage("tokenizer", "done", `${Object.keys(tokJson.model.vocab).length} vocab, ${tokJson.model.merges.length} merges`); |
| |
| setStage("weights", "active"); |
| const weightsBuf = await fetchWithProgress(`${CDN}/model.safetensors`, (r, t) => { |
| setProgress(t ? r / t : 0, "streaming model.safetensors", `${fmtBytes(r)}${t ? " / " + fmtBytes(t) : ""}`); |
| }); |
| setStage("weights", "done", fmtBytes(weightsBuf.byteLength)); |
| |
| setStage("upload", "active"); |
| const parsed = parseSafetensors(weightsBuf); |
| const nTensors = Object.keys(parsed.tensors).length; |
| engine.loadWeightsFromSafetensors(parsed); |
| |
| |
| |
| { |
| const t = parsed.tensors["model.tok_embed.weight"]; |
| engine._embedRaw = new Float32Array(parsed.buffer.slice(t.byteOffset, t.byteOffset + t.byteLength)); |
| } |
| setStage("upload", "done", `${nTensors} tensors`); |
| |
| setStage("pipelines", "active"); |
| await engine.buildPipelines(); |
| engine._buildInvFreq(); |
| engine.initKVCache(); |
| setStage("pipelines", "done", `${KERNEL_ORDER.length} kernels compiled`); |
| |
| setStage("warmup", "active"); |
| const bosId = tokJson.added_tokens?.[0]?.id ?? 0; |
| await engine.forward([bosId]); |
| engine.initKVCache(); |
| setStage("warmup", "done", "ok"); |
| |
| setProgress(1, "ready", ""); |
| liveDot.classList.add("live"); |
| |
| const loadMs = performance.now() - t0; |
| document.getElementById("stat-loadtime").textContent = loadMs > 1000 ? (loadMs / 1000).toFixed(1) + "s" : Math.round(loadMs) + "ms"; |
| |
| document.getElementById("kernel-source-view").textContent = KERNEL_ORDER.map(k => `// ==== ${k} ====\n${WGSL[k]}`).join("\n\n"); |
| |
| bootSection.style.opacity = "0.55"; |
| loadBtn.textContent = "Loaded ✓"; |
| chatShell.classList.add("active"); |
| chatShell.scrollIntoView({ behavior: "smooth", block: "start" }); |
| |
| } catch (err) { |
| showError(err); |
| loadBtn.disabled = false; |
| loadBtn.textContent = "Retry"; |
| } |
| } |
| |
| loadBtn.addEventListener("click", boot); |
| |
| |
| |
| const transcriptEl = document.getElementById("transcript"); |
| const promptInput = document.getElementById("prompt-input"); |
| const genBtn = document.getElementById("gen-btn"); |
| const stopBtn = document.getElementById("stop-btn"); |
| const tempSlider = document.getElementById("temp-slider"); |
| const topkSlider = document.getElementById("topk-slider"); |
| const maxtokSlider = document.getElementById("maxtok-slider"); |
| const tempVal = document.getElementById("temp-val"); |
| const topkVal = document.getElementById("topk-val"); |
| const maxtokVal = document.getElementById("maxtok-val"); |
| const perfLine = document.getElementById("perf-line"); |
| |
| tempSlider.addEventListener("input", () => tempVal.textContent = Number(tempSlider.value).toFixed(2)); |
| topkSlider.addEventListener("input", () => topkVal.textContent = topkSlider.value); |
| maxtokSlider.addEventListener("input", () => maxtokVal.textContent = maxtokSlider.value); |
| |
| function clearPlaceholder() { |
| const hint = transcriptEl.querySelector(".placeholder-hint"); |
| if (hint) hint.remove(); |
| } |
| |
| async function runGeneration() { |
| const prompt = promptInput.value; |
| if (!prompt.trim() || !engine) return; |
| |
| clearPlaceholder(); |
| genBtn.disabled = true; |
| stopBtn.disabled = false; |
| promptInput.disabled = true; |
| genController.stop = false; |
| |
| const userMsg = document.createElement("div"); |
| userMsg.className = "msg user"; |
| userMsg.innerHTML = `<div class="who">prompt</div><div class="txt"></div>`; |
| userMsg.querySelector(".txt").textContent = prompt; |
| transcriptEl.appendChild(userMsg); |
| |
| const genMsg = document.createElement("div"); |
| genMsg.className = "msg gen"; |
| genMsg.innerHTML = `<div class="who">İvme-Conversate-v2</div><div class="txt"><span class="cursor"></span></div>`; |
| transcriptEl.appendChild(genMsg); |
| const txtEl = genMsg.querySelector(".txt"); |
| transcriptEl.scrollTop = transcriptEl.scrollHeight; |
| |
| try { |
| |
| |
| engine.initKVCache(); |
| |
| const ids = engine.tokenizer.encode(prompt); |
| const maxNew = Number(maxtokSlider.value); |
| const temperature = Number(tempSlider.value); |
| const topK = Number(topkSlider.value); |
| const contextLen = engine.cfg.context_len; |
| |
| const eosId = 0; |
| |
| let generated = ""; |
| let tokenCount = 0; |
| const t0 = performance.now(); |
| |
| |
| let logits = await engine.forward(ids); |
| |
| while (tokenCount < maxNew && !genController.stop) { |
| const nextId = engine.sample(logits, temperature, topK); |
| if (nextId === eosId) break; |
| |
| const piece = engine.tokenizer.decode([nextId]); |
| generated += piece; |
| tokenCount++; |
| txtEl.innerHTML = ""; |
| txtEl.appendChild(document.createTextNode(generated)); |
| const cursor = document.createElement("span"); |
| cursor.className = "cursor"; |
| txtEl.appendChild(cursor); |
| transcriptEl.scrollTop = transcriptEl.scrollHeight; |
| |
| if (engine.curLen >= contextLen) break; |
| |
| logits = await engine.forward([nextId]); |
| } |
| |
| txtEl.querySelector(".cursor")?.remove(); |
| |
| const elapsed = (performance.now() - t0) / 1000; |
| const toksPerSec = tokenCount / Math.max(elapsed, 1e-6); |
| document.getElementById("stat-toksec").textContent = toksPerSec.toFixed(1); |
| perfLine.innerHTML = `<span><b>${tokenCount}</b> tokens generated</span><span><b>${elapsed.toFixed(2)}s</b> wall time</span><span><b>${toksPerSec.toFixed(2)}</b> tok/s</span><span><b>${ids.length}</b> prompt tokens</span>`; |
| |
| } catch (err) { |
| showError(err); |
| txtEl.querySelector(".cursor")?.remove(); |
| } finally { |
| genBtn.disabled = false; |
| stopBtn.disabled = true; |
| promptInput.disabled = false; |
| } |
| } |
| |
| genBtn.addEventListener("click", runGeneration); |
| promptInput.addEventListener("keydown", (e) => { |
| if (e.key === "Enter" && !e.shiftKey) { |
| e.preventDefault(); |
| runGeneration(); |
| } |
| }); |
| stopBtn.addEventListener("click", () => { genController.stop = true; }); |
| </script> |
| </body> |
| </html> |
