Spaces:

HOLOGRAMTECH
/

q

Running

App Files Files Community

Humuhumu33 commited on Jul 4

Commit

082923c

verified ·

1 Parent(s): d88b0fc

Upload roofline.html with huggingface_hub

Browse files

Files changed (1) hide show

roofline.html +45 -44

roofline.html CHANGED Viewed

@@ -4,7 +4,7 @@
   :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:760px;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}
   .big{font-size:22px;font-weight:700}.k{color:var(--dim)}.v{font-weight:600}
   .ok{color:var(--ok)}.no{color:var(--no)}.warn{color:var(--warn)}
   table{width:100%;border-collapse:collapse;font-family:ui-monospace,monospace;font-size:13px}
@@ -13,78 +13,79 @@
 </style></head><body>
 <h1>BitNet decode <span style="color:var(--ac)">roofline</span> on your GPU</h1>
 <p class="sub">Measures your GPU's achievable VRAM read bandwidth, then computes the batch-1 decode ceiling for BitNet-2B (0.69&nbsp;GB read per token) and how far the current kernel is from it.</p>
-<div id="status" class="card">measuring…</div>
 <div id="out"></div>
 <script type="module">
 const $=s=>document.querySelector(s), st=$("#status"), out=$("#out");
-const MODEL_GB = 0.69;          // BitNet-2B t2 weights read per token
-const CUR_MS = 18.4;            // measured ternary-kernel ms/tok on this app (RDNA-3)
-// A pure VRAM read-bandwidth kernel: grid-stride read the whole buffer, XOR-accumulate (defeats DCE),
-// every thread writes its accumulator so no read is optimized away. Time best-of-N → achievable GB/s.
 const WGSL = `
 @group(0) @binding(0) var<storage, read> data : array<vec4<u32>>;
 @group(0) @binding(1) var<storage, read_write> sink : array<u32>;
-@group(0) @binding(2) var<uniform> P : vec4<u32>;   // x=#vec4 elems, y=total threads
 @compute @workgroup_size(256)
 fn main(@builtin(global_invocation_id) gid:vec3<u32>){
-  let n=P.x; let stride=P.y; var acc=vec4<u32>(0u);
-  var i=gid.x;
-  loop { if(i>=n){break;} acc=acc ^ data[i]; i=i+stride; }
   sink[gid.x]=acc.x ^ acc.y ^ acc.z ^ acc.w;
 }`;
 (async()=>{
-  if(!navigator.gpu){ st.innerHTML='<span class="no">✗ No WebGPU. Open in Chrome/Edge or a recent browser.</span>'; return; }
-  const ad=await navigator.gpu.requestAdapter({powerPreference:"high-performance"});
   const L=ad.limits;
-  const dev=await ad.requestDevice({requiredLimits:{maxStorageBufferBindingSize:L.maxStorageBufferBindingSize,maxBufferSize:L.maxBufferSize,maxComputeWorkgroupsPerDimension:L.maxComputeWorkgroupsPerDimension}});
   const info=ad.info||{};
-  // buffer: as large as the binding allows, capped at 1 GiB, 16-byte aligned
-  const cap=Math.min(L.maxStorageBufferBindingSize, 1024*1024*1024);
-  const bytes=Math.floor(cap/16)*16;
   const nVec=bytes/16;
   const buf=dev.createBuffer({size:bytes, usage:GPUBufferUsage.STORAGE});
-  const TOTAL=Math.min(L.maxComputeWorkgroupsPerDimension,65535)*256;
   const sink=dev.createBuffer({size:TOTAL*4, usage:GPUBufferUsage.STORAGE});
   const P=dev.createBuffer({size:16, usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST});
-  dev.queue.writeBuffer(P,0,new Uint32Array([nVec, TOTAL, 0,0]));
   const mod=dev.createShaderModule({code:WGSL});
   const ci=await mod.getCompilationInfo(); const er=ci.messages.filter(m=>m.type==="error");
-  if(er.length){ st.innerHTML='<span class="no">✗ WGSL: '+er[0].message+'</span>'; return; }
   const pipe=dev.createComputePipeline({layout:"auto",compute:{module:mod,entryPoint:"main"}});
   const bg=dev.createBindGroup({layout:pipe.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:buf}},{binding:1,resource:{buffer:sink}},{binding:2,resource:{buffer:P}}]});
   const rb=dev.createBuffer({size:4, usage:GPUBufferUsage.COPY_DST|GPUBufferUsage.MAP_READ});
-  const wg=Math.min(L.maxComputeWorkgroupsPerDimension,65535);
   async function once(){ const e=dev.createCommandEncoder(); const p=e.beginComputePass(); p.setPipeline(pipe); p.setBindGroup(0,bg); p.dispatchWorkgroups(wg); p.end(); e.copyBufferToBuffer(sink,0,rb,0,4); const t0=performance.now(); dev.queue.submit([e.finish()]); await rb.mapAsync(GPUMapMode.READ); rb.unmap(); return performance.now()-t0; }
-  st.textContent="warming + timing "+(bytes/1073741824).toFixed(2)+" GB reads…";
-  await once(); await once();                        // warm
-  let best=1e9; for(let k=0;k<12;k++){ best=Math.min(best, await once()); }
-  const gbps = (bytes/1073741824)/(best/1000);       // achievable read GB/s
-  const roofTok = gbps/MODEL_GB;                     // batch-1 decode ceiling
-  const curTok = 1000/CUR_MS;                        // current ~54
-  const curGBps = MODEL_GB/(CUR_MS/1000);            // current effective ~37
-  const pct = 100*curGBps/gbps;                      // how much of roofline the kernel reaches
-  const gap = gbps/curGBps;                          // headroom multiplier
-  st.innerHTML='<span class="ok">done</span> · adapter: <b>'+((info.vendor||"?")+" "+(info.architecture||"")+" "+(info.device||"")).trim()+'</b>';
-  const reachable = roofTok>=1000;
-  const specNeeded = Math.max(1, 1000/roofTok);
   out.innerHTML=`
     <div class="card"><span class="k">Achievable VRAM read bandwidth</span><br><span class="big">${gbps.toFixed(0)} GB/s</span></div>
-    <div class="card">
-      <table>
-        <tr><td>Weights read / token (BitNet-2B t2)</td><td class="n">${MODEL_GB} GB</td></tr>
-        <tr><td><b>Bandwidth roofline (batch-1 ceiling)</b></td><td class="n"><b>${roofTok.toFixed(0)} tok/s</b></td></tr>
-        <tr><td>Current kernel</td><td class="n">${curTok.toFixed(0)} tok/s · ${curGBps.toFixed(0)} GB/s</td></tr>
-        <tr><td>Current kernel reaches</td><td class="n ${pct<20?'no':pct<60?'warn':'ok'}">${pct.toFixed(1)}% of roofline</td></tr>
-        <tr><td>Kernel headroom to roofline</td><td class="n">${gap.toFixed(1)}×</td></tr>
-      </table>
-    </div>
     <div class="verdict ${reachable?'ok':'warn'}">${reachable
-      ? '✓ >1000 tok/s is within the bandwidth roofline ('+roofTok.toFixed(0)+') on this GPU — a bandwidth-optimal ternary kernel ('+gap.toFixed(1)+'× headroom) gets there; speculative decode is margin.'
-      : '⚠ This GPU\\'s roofline is '+roofTok.toFixed(0)+' tok/s at batch-1. A perfect kernel reaches ~'+roofTok.toFixed(0)+'; >1000 needs speculative decode accepting ≥'+specNeeded.toFixed(1)+' tokens/pass (or a dedicated card with more bandwidth).'}
     </div>`;
 })();
 </script></body></html>

   :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:760px;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}
   .big{font-size:22px;font-weight:700}.k{color:var(--dim)}.v{font-weight:600}
   .ok{color:var(--ok)}.no{color:var(--no)}.warn{color:var(--warn)}
   table{width:100%;border-collapse:collapse;font-family:ui-monospace,monospace;font-size:13px}
 </style></head><body>
 <h1>BitNet decode <span style="color:var(--ac)">roofline</span> on your GPU</h1>
 <p class="sub">Measures your GPU's achievable VRAM read bandwidth, then computes the batch-1 decode ceiling for BitNet-2B (0.69&nbsp;GB read per token) and how far the current kernel is from it.</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, CUR_MS = 18.4;
+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 WGSL = `
 @group(0) @binding(0) var<storage, read> data : array<vec4<u32>>;
 @group(0) @binding(1) var<storage, read_write> sink : array<u32>;
+@group(0) @binding(2) var<uniform> P : vec4<u32>;   // x=#vec4 elems, y=total threads, z=repeat
 @compute @workgroup_size(256)
 fn main(@builtin(global_invocation_id) gid:vec3<u32>){
+  let n=P.x; let stride=P.y; let R=P.z; var acc=vec4<u32>(0u);
+  for(var rep=0u; rep<R; rep=rep+1u){ var i=gid.x; loop { if(i>=n){break;} acc=acc ^ data[i]; i=i+stride; } }
   sink[gid.x]=acc.x ^ acc.y ^ acc.z ^ acc.w;
 }`;
 (async()=>{
+ try{
+  if(!navigator.gpu){ say("✗ No WebGPU. Open in Chrome/Edge or a recent browser.","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 available.","no"); return; }
   const L=ad.limits;
+  say("requesting device…");
+  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"));
+  dev.pushErrorScope("validation");
   const info=ad.info||{};
+  // buffer capped by BOTH limits and 256 MB (safe on every GPU; we re-read it R times for signal)
+  const bytes=Math.floor(Math.min(L.maxStorageBufferBindingSize, L.maxBufferSize, 256*1024*1024)/16)*16;
   const nVec=bytes/16;
+  say("allocating "+(bytes/1048576).toFixed(0)+" MB…");
   const buf=dev.createBuffer({size:bytes, usage:GPUBufferUsage.STORAGE});
+  const wg=Math.min(L.maxComputeWorkgroupsPerDimension,65535), TOTAL=wg*256, R=16;
   const sink=dev.createBuffer({size:TOTAL*4, usage:GPUBufferUsage.STORAGE});
   const P=dev.createBuffer({size:16, usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST});
+  dev.queue.writeBuffer(P,0,new Uint32Array([nVec, TOTAL, R, 0]));
   const mod=dev.createShaderModule({code:WGSL});
   const ci=await mod.getCompilationInfo(); const er=ci.messages.filter(m=>m.type==="error");
+  if(er.length){ say("✗ WGSL: "+er[0].message,"no"); return; }
   const pipe=dev.createComputePipeline({layout:"auto",compute:{module:mod,entryPoint:"main"}});
   const bg=dev.createBindGroup({layout:pipe.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:buf}},{binding:1,resource:{buffer:sink}},{binding:2,resource:{buffer:P}}]});
   const rb=dev.createBuffer({size:4, usage:GPUBufferUsage.COPY_DST|GPUBufferUsage.MAP_READ});
+  const scopeErr = await dev.popErrorScope();
+  if(scopeErr){ say("✗ GPU validation: "+scopeErr.message,"no"); return; }
   async function once(){ const e=dev.createCommandEncoder(); const p=e.beginComputePass(); p.setPipeline(pipe); p.setBindGroup(0,bg); p.dispatchWorkgroups(wg); p.end(); e.copyBufferToBuffer(sink,0,rb,0,4); const t0=performance.now(); dev.queue.submit([e.finish()]); await rb.mapAsync(GPUMapMode.READ); rb.unmap(); return performance.now()-t0; }
+  say("timing "+(bytes/1048576*R).toFixed(0)+" MB reads…");
+  await once();
+  let best=1e9; for(let k=0;k<10;k++){ best=Math.min(best, await once()); }
+  const gbps = (bytes*R/1073741824)/(best/1000);
+  const roofTok=gbps/MODEL_GB, curTok=1000/CUR_MS, curGBps=MODEL_GB/(CUR_MS/1000);
+  const pct=100*curGBps/gbps, gap=gbps/curGBps, reachable=roofTok>=1000, specNeeded=Math.max(1,1000/roofTok);
+  say("done · adapter: "+((info.vendor||"?")+" "+(info.architecture||"")+" "+(info.device||"")).trim(),"ok");
   out.innerHTML=`
     <div class="card"><span class="k">Achievable VRAM read bandwidth</span><br><span class="big">${gbps.toFixed(0)} GB/s</span></div>
+    <div class="card"><table>
+      <tr><td>Weights read / token (BitNet-2B t2)</td><td class="n">${MODEL_GB} GB</td></tr>
+      <tr><td><b>Bandwidth roofline (batch-1 ceiling)</b></td><td class="n"><b>${roofTok.toFixed(0)} tok/s</b></td></tr>
+      <tr><td>Current kernel</td><td class="n">${curTok.toFixed(0)} tok/s · ${curGBps.toFixed(0)} GB/s</td></tr>
+      <tr><td>Current kernel reaches</td><td class="n ${pct<20?'no':pct<60?'warn':'ok'}">${pct.toFixed(1)}% of roofline</td></tr>
+      <tr><td>Kernel headroom to roofline</td><td class="n">${gap.toFixed(1)}×</td></tr>
+    </table></div>
     <div class="verdict ${reachable?'ok':'warn'}">${reachable
+      ? '✓ >1000 tok/s is within this GPU\\'s roofline ('+roofTok.toFixed(0)+') — a bandwidth-optimal ternary kernel ('+gap.toFixed(1)+'× headroom) gets there; spec-decode is margin.'
+      : '⚠ This GPU\\'s batch-1 roofline is '+roofTok.toFixed(0)+' tok/s. A perfect kernel reaches ~'+roofTok.toFixed(0)+'; >1000 needs speculative decode accepting ≥'+specNeeded.toFixed(1)+' tokens/pass, or a card with more bandwidth.'}
     </div>`;
+ }catch(e){ say("✗ "+(e && (e.message||e)), "no"); }
 })();
 </script></body></html>