roofline: fill buffer non-zero + exceed cache for real VRAM bw

roofline.html

@@ -22,6 +22,19 @@ 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"));
 
+// FILL: write a non-zero, non-compressible pattern so the read pass hits real VRAM
+// (an uninitialised STORAGE buffer is zero/DCC-compressed → the memory controller
+//  never fetches it → measured bandwidth is fictional). Must run before timing.
+const FILL = `
+@group(0) @binding(0) var<storage, read_write> d : array<vec4<u32>>;
+@group(0) @binding(1) 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 i=gid.x;
+  loop { if(i>=n){break;}
+    d[i]=vec4<u32>(i*2654435761u+1u, i*40503u+7u, i^0x9e3779b9u, i*2246822519u+3u);
+    i=i+stride; }
+}`;
 const WGSL = `
 @group(0) @binding(0) var<storage, read> data : array<vec4<u32>>;
 @group(0) @binding(1) var<storage, read_write> sink : array<u32>;
@@ -48,15 +61,23 @@ fn main(@builtin(global_invocation_id) gid:vec3<u32>){
   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;
+  // Buffer must exceed the GPU's last-level cache (RDNA-3 Infinity Cache is up to 96 MB)
+  // so re-reads actually miss to VRAM. Aim as large as the device allows, floor 512 MB.
+  const wantBytes=Math.min(L.maxStorageBufferBindingSize, L.maxBufferSize, 1024*1024*1024);
+  const bytes=Math.floor(Math.max(512*1024*1024, wantBytes)/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 wg=Math.min(L.maxComputeWorkgroupsPerDimension,65535), TOTAL=wg*256, R=4;
   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]));
+  // fill pipeline (non-zero pattern) — its own bind group, same buffer
+  const fmod=dev.createShaderModule({code:FILL});
+  const fci=await fmod.getCompilationInfo(); const fer=fci.messages.filter(m=>m.type==="error");
+  if(fer.length){ say("✗ FILL WGSL: "+fer[0].message,"no"); return; }
+  const fpipe=dev.createComputePipeline({layout:"auto",compute:{module:fmod,entryPoint:"main"}});
+  const fbg=dev.createBindGroup({layout:fpipe.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:buf}},{binding:1,resource:{buffer:P}}]});
   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; }
@@ -65,6 +86,9 @@ fn main(@builtin(global_invocation_id) gid:vec3<u32>){
   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; }
+  // write the non-zero pattern into VRAM before we time reads
+  say("filling "+(bytes/1048576).toFixed(0)+" MB with non-zero pattern…");
+  { const e=dev.createCommandEncoder(); const p=e.beginComputePass(); p.setPipeline(fpipe); p.setBindGroup(0,fbg); p.dispatchWorkgroups(wg); p.end(); dev.queue.submit([e.finish()]); await dev.queue.onSubmittedWorkDone(); }
   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();