Spaces:

HOLOGRAMTECH
/

q

Running

App Files Files Community

Humuhumu33 commited on Jul 4

Commit

d88b0fc

verified ·

1 Parent(s): b50807f

Upload roofline.html with huggingface_hub

Browse files

Files changed (1) hide show

roofline.html +90 -0

roofline.html ADDED Viewed

	@@ -0,0 +1,90 @@

+<!doctype html><html><head><meta charset=utf8><meta name=viewport content="width=device-width,initial-scale=1">
+<title>Decode roofline — 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: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}
+  td{padding:5px 8px;border-bottom:1px solid var(--line)}td.n{text-align:right;font-variant-numeric:tabular-nums}
+  .verdict{font-size:17px;font-weight:600;margin-top:14px}
+</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>