Spaces:

HOLOGRAMTECH
/

q

Running

File size: 11,546 Bytes

<!doctype html><html><head><meta charset=utf8><meta name=viewport content="width=device-width,initial-scale=1">
<title>Ternary GEMV limiter profile — your GPU</title>
<style>
  :root{--bg:#0a0d13;--panel:#111825;--ink:#e8ebf1;--dim:#8b95a7;--ac:#7c5cff;--ok:#48c26c;--no:#f0616d;--warn:#e0a94a;--line:#1e2836}
  *{box-sizing:border-box}body{margin:0;background:var(--bg);color:var(--ink);font:15px/1.6 -apple-system,Segoe UI,Roboto,monospace;padding:24px;max-width:860px;margin:0 auto}
  h1{font-size:20px;margin:0 0 4px}.sub{color:var(--dim);font-size:13px;margin:0 0 18px}
  .card{font-family:ui-monospace,monospace;padding:14px 16px;border:1px solid var(--line);border-radius:10px;background:var(--panel);margin-bottom:12px;word-break:break-word}
  table{width:100%;border-collapse:collapse;font-family:ui-monospace,monospace;font-size:13px}
  td,th{padding:6px 8px;border-bottom:1px solid var(--line);text-align:left}th{color:var(--dim);font-weight:600}
  td.n{text-align:right;font-variant-numeric:tabular-nums}
  .verdict{font-size:16px;font-weight:600;margin-top:14px}.bar{display:inline-block;height:9px;background:var(--ac);border-radius:2px;vertical-align:middle}
</style></head><body>
<h1>Ternary GEMV <span style="color:var(--ac)">limiter profile</span> — your GPU</h1>
<p class="sub">Same GEMV access pattern, one knob at a time. <b>read-only</b> = the reads with zero math (memory ceiling at this pattern). <b>dot ×1</b> = the real kernel. <b>×2/×4</b> = same reads, more ALU. If read-only ≫ dot, the ALU is the wall (int8 is the lever); if read-only ≈ dot, the memory access pattern is the wall (int8 won't help).</p>
<div id="status" class="card">starting…</div>
<div id="out"></div>

<script type="module">
const $=s=>document.querySelector(s), st=$("#status"), out=$("#out");
const MODEL_GB=0.69;
const say=(t,cls)=>{ st.textContent=t; st.className="card"+(cls?" "+cls:""); };
window.addEventListener("unhandledrejection",e=>say("✗ unhandled: "+(e.reason&&(e.reason.message||e.reason)),"no"));

const DOT16=`
fn dot16(word:u32, v:u32) -> f32 {
  var s4=vec4<f32>(0.0);
  s4=s4+x[v]  *(vec4<f32>(f32(word&3u),f32((word>>2u)&3u),f32((word>>4u)&3u),f32((word>>6u)&3u))-vec4<f32>(1.0));
  s4=s4+x[v+1u]*(vec4<f32>(f32((word>>8u)&3u),f32((word>>10u)&3u),f32((word>>12u)&3u),f32((word>>14u)&3u))-vec4<f32>(1.0));
  s4=s4+x[v+2u]*(vec4<f32>(f32((word>>16u)&3u),f32((word>>18u)&3u),f32((word>>20u)&3u),f32((word>>22u)&3u))-vec4<f32>(1.0));
  s4=s4+x[v+3u]*(vec4<f32>(f32((word>>24u)&3u),f32((word>>26u)&3u),f32((word>>28u)&3u),f32((word>>30u)&3u))-vec4<f32>(1.0));
  return s4.x+s4.y+s4.z+s4.w;
}`;
// mode: 'read' (xor raw, no unpack) | 'dot' (R× dot16, same reads). T threads/row, ROWS rows/wg (T*ROWS=256).
const kern=(mode,R,T,ROWS)=>`
@group(0) @binding(0) var<storage,read> x: array<vec4<f32>>;
@group(0) @binding(1) var<storage,read> qw: array<u32>;
@group(0) @binding(2) var<storage,read_write> o: array<f32>;
@group(0) @binding(3) var<uniform> P: vec4<u32>;   // K, N
var<workgroup> red: array<f32, 256>;
${mode==='dot'?DOT16:''}
@compute @workgroup_size(256)
fn main(@builtin(workgroup_id) wg:vec3<u32>, @builtin(local_invocation_id) lid:vec3<u32>){
  let K=P.x; let nw=K>>4u;
  let rr=lid.x/${T}u; let t=lid.x%${T}u;
  let n0=(wg.y*65535u+wg.x)*${ROWS}u+rr; let n=min(n0, P.y-1u);
  let rowW=n*nw;
  ${mode==='read'?'var accu=0u;':'var acc=0.0;'}
  var w=t;
  loop{ if(w>=nw){break;}
    let word=qw[rowW+w]; let v=w<<2u;
    ${mode==='read'?'accu=accu^word;':Array.from({length:R},(_,r)=>`acc=acc+dot16(word^${r}u, v);`).join(' ')}
    w=w+${T}u; }
  ${mode==='read'?'var acc=f32(accu&1u); if(accu==0xffffffffu){ acc=acc+x[0].x; }':''}
  red[lid.x]=acc; workgroupBarrier();
  var s=${T>>1}u; loop{ if(s==0u){break;} if(t<s){ red[rr*${T}u+t]=red[rr*${T}u+t]+red[rr*${T}u+t+s]; } workgroupBarrier(); s=s/2u; }
  if(t==0u && n0<P.y){ o[n0]=red[rr*${T}u]; }
}`;

const FILL=`
@group(0) @binding(0) var<storage,read_write> d: array<u32>;
@group(0) @binding(1) var<uniform> P: vec4<u32>;
@compute @workgroup_size(256)
fn main(@builtin(global_invocation_id) gid:vec3<u32>){ let n=P.x; var i=gid.x; loop{ if(i>=n){break;} d[i]=(i*2654435761u+1u)^((i<<7u)*40503u); i=i+P.y; } }`;

const VARIANTS=[
  {id:"read",  name:"read-only (no ALU) — memory ceiling", mode:"read", R:0, T:64, ROWS:4, grp:"alu"},
  {id:"dot1",  name:"dot ×1 (the real kernel)",            mode:"dot",  R:1, T:64, ROWS:4, grp:"alu"},
  {id:"dot2",  name:"dot ×2 (2× ALU, same reads)",         mode:"dot",  R:2, T:64, ROWS:4, grp:"alu"},
  {id:"dot4",  name:"dot ×4 (4× ALU, same reads)",         mode:"dot",  R:4, T:64, ROWS:4, grp:"alu"},
  {id:"dot8",  name:"dot ×8 (KX=8 batched-verify proxy)",  mode:"dot",  R:8, T:64, ROWS:4, grp:"alu"},
  {id:"occ32", name:"dot ×1 · 32 thr/row · 8 rows/wg",     mode:"dot",  R:1, T:32, ROWS:8, grp:"occ"},
  {id:"occ128",name:"dot ×1 · 128 thr/row · 2 rows/wg",    mode:"dot",  R:1, T:128,ROWS:2, grp:"occ"},
];

(async()=>{
 try{
  if(!navigator.gpu){ say("✗ No WebGPU. Open in Chrome/Edge.","no"); return; }
  say("requesting adapter…");
  let ad=await navigator.gpu.requestAdapter({powerPreference:"high-performance"}); if(!ad) ad=await navigator.gpu.requestAdapter();
  if(!ad){ say("✗ no GPU adapter.","no"); return; }
  const L=ad.limits, info=ad.info||{};
  const dev=await ad.requestDevice({requiredLimits:{maxStorageBufferBindingSize:L.maxStorageBufferBindingSize, maxBufferSize:L.maxBufferSize, maxComputeWorkgroupsPerDimension:L.maxComputeWorkgroupsPerDimension}});
  dev.lost.then(i=>say("✗ device lost: "+(i&&i.message||i.reason||""),"no"));

  const K=4096, nw=K/16;
  const wantBytes=Math.min(L.maxStorageBufferBindingSize, L.maxBufferSize, Math.round(MODEL_GB*1e9));
  const N=Math.floor((wantBytes/4)/nw), qU32=N*nw, wBytes=qU32*4;
  say("allocating "+(wBytes/1048576).toFixed(0)+" MB…");
  const qw=dev.createBuffer({size:wBytes, usage:GPUBufferUsage.STORAGE});
  const xbuf=dev.createBuffer({size:K*4, usage:GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_DST});
  const obuf=dev.createBuffer({size:N*4, usage:GPUBufferUsage.STORAGE});
  const P=dev.createBuffer({size:16, usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST});
  const xa=new Float32Array(K); for(let i=0;i<K;i++) xa[i]=(Math.sin(i*12.9898)*43758.5453)%1;
  dev.queue.writeBuffer(xbuf,0,xa); dev.queue.writeBuffer(P,0,new Uint32Array([K,N,0,0]));

  say("filling weights…");
  { const fmod=dev.createShaderModule({code:FILL}); const fpipe=dev.createComputePipeline({layout:"auto",compute:{module:fmod,entryPoint:"main"}});
    const fp=dev.createBuffer({size:16,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST}); const wgF=Math.min(L.maxComputeWorkgroupsPerDimension,65535);
    dev.queue.writeBuffer(fp,0,new Uint32Array([qU32, wgF*256, 0, 0]));
    const fbg=dev.createBindGroup({layout:fpipe.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:qw}},{binding:1,resource:{buffer:fp}}]});
    const e=dev.createCommandEncoder(); const p=e.beginComputePass(); p.setPipeline(fpipe); p.setBindGroup(0,fbg); p.dispatchWorkgroups(wgF); p.end();
    dev.queue.submit([e.finish()]); await dev.queue.onSubmittedWorkDone(); }

  const PASSES=10, ITERS=5;
  const rows=[];
  for(const V of VARIANTS){
    say("compiling "+V.id+"…");
    dev.pushErrorScope("validation");
    const mod=dev.createShaderModule({code:kern(V.mode,V.R,V.T,V.ROWS)});
    const ci=await mod.getCompilationInfo(); const er=ci.messages.filter(m=>m.type==="error");
    if(er.length){ await dev.popErrorScope(); rows.push({V,err:er[0].message}); continue; }
    const pipe=dev.createComputePipeline({layout:"auto",compute:{module:mod,entryPoint:"main"}});
    const bg=dev.createBindGroup({layout:pipe.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:xbuf}},{binding:1,resource:{buffer:qw}},{binding:2,resource:{buffer:obuf}},{binding:3,resource:{buffer:P}}]});
    const se=await dev.popErrorScope(); if(se){ rows.push({V,err:se.message}); continue; }
    const groups=Math.ceil(N/V.ROWS), wgx=Math.min(groups,65535), wgy=Math.ceil(groups/65535);
    async function run(passes){ const e=dev.createCommandEncoder(); for(let k=0;k<passes;k++){ const p=e.beginComputePass(); p.setPipeline(pipe); p.setBindGroup(0,bg); p.dispatchWorkgroups(wgx,wgy); p.end(); } const t0=performance.now(); dev.queue.submit([e.finish()]); await dev.queue.onSubmittedWorkDone(); return performance.now()-t0; }
    say("timing "+V.id+"…");
    await run(2); let best=1e9; for(let k=0;k<ITERS;k++){ best=Math.min(best, await run(PASSES)); }
    const gbps=(wBytes/1073741824)/((best/PASSES)/1000);
    rows.push({V, gbps, tok:gbps/MODEL_GB});
  }

  say("done · adapter: "+((info.vendor||"?")+" "+(info.architecture||"")+" "+(info.device||"")).trim(),"ok");
  const g=id=>rows.find(r=>r.V.id===id&&r.gbps); const rd=g("read"), d1=g("dot1"), d4=g("dot4"), d8=g("dot8"), o32=g("occ32"), o128=g("occ128");
  const maxG=Math.max(...rows.filter(r=>r.gbps).map(r=>r.gbps));
  const readVsDot = (rd&&d1)? rd.gbps/d1.gbps : 0;     // >~1.5 ⇒ reads themselves fly, ALU is the wall
  const aluSlope = (d1&&d4)? d1.gbps/d4.gbps : 0;      // >~1.5 ⇒ time scales with ALU ⇒ ALU-throughput bound
  const occGain = d1? Math.max(o32?o32.gbps/d1.gbps:0, o128?o128.gbps/d1.gbps:0) : 0;
  const batchFree = (d1&&d8)? d8.gbps/d1.gbps : 0;     // ~1 ⇒ 8-wide batched verify costs ~same as 1 ⇒ spec-decode is ~free
  const readCeil = rd? rd.gbps : 0;
  let verdict, cls;
  const specNote = d8 ? ` Batched verify (dot ×8, the KX=8 spec proxy) runs at ${d8.gbps.toFixed(0)} GB/s = ${batchFree.toFixed(2)}× the single kernel — so verifying 8 tokens costs ≈ 1 token. Spec-decode IS ~free here; its earlier 1.0× was per-window JS/fence OVERHEAD, not GPU cost. Killing that overhead unlocks the ~4× the 93%-acceptance implies.` : "";
  if(aluSlope>=1.6 && readVsDot>=1.6){ verdict = `ALU-BOUND. reads fly (${readCeil.toFixed(0)} GB/s) but dot ×1 only ${d1.gbps.toFixed(0)}, and ×1→×4 scales ${aluSlope.toFixed(1)}×. Fewer instructions per weight (int8 dot4I8Packed) is the lever.` + specNote; cls="warn"; }
  else{ verdict = `MEMORY-LATENCY-BOUND — NOT ALU-bound. Quadrupling the ALU (×1→×4) costs only ${aluSlope.toFixed(1)}× → the f32 unpack is hidden; int8 will NOT help. read-only tops ${readCeil.toFixed(0)} GB/s at this pattern vs 152 pure-stream, and dot ×1 sits at ${d1?d1.gbps.toFixed(0):"?"} (${readCeil?(100*d1.gbps/readCeil).toFixed(0):"?"}% of even the read ceiling). Lever = more memory-level parallelism (register-block several output rows per thread so more loads are in flight before the barrier), NOT int8/occupancy.` + specNote; cls=(occGain>=1.15?"warn":"no"); }

  out.innerHTML=`
    <div class="card"><span class="k" style="color:var(--dim)">Test matrix</span> &nbsp; ${N.toLocaleString()} × ${K} · ${(wBytes/1073741824).toFixed(2)} GB · pure-stream roofline 152 GB/s (220 tok/s)</div>
    <div class="card"><table>
      <tr><th>variant</th><th class="n">GB/s</th><th class="n">tok/s-equiv</th><th></th></tr>
      ${rows.map(r=>r.err
        ? `<tr><td>${r.V.name}</td><td colspan=3 class="no">failed: ${r.err.slice(0,54)}</td></tr>`
        : `<tr><td>${r.V.name}</td><td class="n">${r.gbps.toFixed(0)}</td><td class="n">${r.tok.toFixed(0)}</td><td><span class="bar" style="width:${Math.round(140*r.gbps/maxG)}px"></span></td></tr>`).join("")}
    </table></div>
    <div class="verdict ${cls}">${verdict}</div>`;
 }catch(e){ say("✗ "+(e&&(e.message||e)), "no"); }
})();
</script></body></html>