inductiveML/monarch-webgpu
1
1var an=["timestamp-query","shader-f16","subgroups"];function un(e){return an.includes(e)}function Je(e){if(!un(e))throw new Error(`${e} is not a registered requestable feature`);return e}function nt(){return{dispatches:0,dispatchesIndirect:0,submits:0,commandBuffers:0,passes:0,bindGroups:0,writeBuffers:0,writeBufferBytes:0,createBuffers:0,createBufferBytes:0,destroyBuffers:0,pipelines:0,pipelinesAsync:0,pipelineCompileMs:0,shaderModules:0,mapAsyncs:0,copyBufferToBuffer:0,submitOffsets:[],injected:0}}function ln(e){let r={},t=Object.getPrototypeOf(e);for(let o of Object.getOwnPropertyNames(t)){if(o==="constructor")continue;let n=Object.getOwnPropertyDescriptor(t,o);if(n===void 0||typeof n.get!="function")continue;let s=n.get.call(e);typeof s=="number"&&(r[o]=s)}return r}function cn(e){let r={};if(e===void 0)return r;for(let[t,o]of Object.entries(e))typeof o=="number"&&(r[t]=o);return r}(()=>{if(typeof navigator>"u"||navigator.gpu===void 0){globalThis.__probeHooks={installed:!1,configure:()=>{},beginStep:()=>{},endStep:()=>{throw new Error("WebGPU unavailable")},snapshot:()=>{throw new Error("WebGPU unavailable")},setInjection:()=>{},enableTimestamps:()=>{},disableTimestamps:()=>{},collectTimestamps:()=>Promise.resolve([]),devices:()=>[],descriptors:()=>[],errors:()=>[],lost:()=>[],liveBytes:()=>0,suspend:p=>p(),suspendAsync:p=>p()};return}let e=!1,r=!1,t=0,o=0,n=nt(),s=null,a=null,i=0,u=[],c=[],f=[],k=[],y=[],P=new WeakMap,B=new WeakMap,S=()=>e&&!r,v=GPUAdapter.prototype,$=v.requestDevice;v.requestDevice=async function(p){let g=p===void 0?[]:[...p.requiredFeatures??[]],w=[...g],R=[];for(let M of u){let te=Je(M);!w.includes(te)&&this.features.has(te)&&(w.push(te),R.push(te))}let L={...p??{},requiredFeatures:w},h=this.info,A={requestedFeatures:g.map(String),requiredFeatures:w.map(String),addedFeatures:R.map(String),requiredLimits:cn(p?.requiredLimits),label:p?.label??"",adapter:{vendor:h.vendor,architecture:h.architecture,device:h.device,description:h.description,isFallbackAdapter:h.isFallbackAdapter,subgroupMinSize:h.subgroupMinSize,subgroupMaxSize:h.subgroupMaxSize,features:[...this.features].map(String).sort(),limits:ln(this.limits)}},N=await $.call(this,L);return c.push(N),f.push(A),N.addEventListener("uncapturederror",M=>{k.push(`${M.error.constructor.name}: ${M.error.message}`)}),N.lost.then(M=>{y.push(`${M.reason}: ${M.message}`)}),N};let x=GPUDevice.prototype,X=x.createBuffer;x.createBuffer=function(p){let g=X.call(this,p);return P.set(g,p.size),i+=p.size,S()&&(n.createBuffers+=1,n.createBufferBytes+=p.size),g};let Z=x.createBindGroup;x.createBindGroup=function(p){return S()&&(n.bindGroups+=1),Z.call(this,p)};let I=x.createShaderModule;x.createShaderModule=function(p){return S()&&(n.shaderModules+=1),I.call(this,p)};let q=x.createComputePipeline;x.createComputePipeline=function(p){let g=performance.now(),w=q.call(this,p);return S()&&(n.pipelines+=1,n.pipelineCompileMs+=performance.now()-g),w};let ee=x.createComputePipelineAsync;x.createComputePipelineAsync=async function(p){let g=performance.now(),w=await ee.call(this,p);return S()&&(n.pipelinesAsync+=1,n.pipelineCompileMs+=performance.now()-g),w};let pe=GPUBuffer.prototype,Ae=pe.destroy;pe.destroy=function(){let p=P.get(this);return p!==void 0&&(i-=p,P.delete(this)),S()&&(n.destroyBuffers+=1),Ae.call(this)};let Pe=pe.mapAsync;pe.mapAsync=function(p,g,w){return S()&&(n.mapAsyncs+=1),Pe.call(this,p,g,w)};let E=GPUQueue.prototype,ke=E.submit;E.submit=function(p){let g=[...p];return S()&&(s!==null&&(g=[...g,s],s=null,n.injected+=1),n.submits+=1,n.commandBuffers+=g.length,n.submitOffsets.push(performance.now()-t)),ke.call(this,g)};let Q=E.writeBuffer;E.writeBuffer=function(p,g,w,R,L){if(S()){n.writeBuffers+=1;let h=L!==void 0?L:w.byteLength-(R??0);n.writeBufferBytes+=h}return Q.call(this,p,g,w,R,L)};let K=GPUCommandEncoder.prototype,re=K.beginComputePass;K.beginComputePass=function(p){let g=p,w=null;if(S()&&(n.passes+=1,a!==null&&a.next<a.capacity&&(p===void 0||p.timestampWrites===void 0))){let L=a.next;a.next+=1,g={...p??{},timestampWrites:{querySet:a.querySet,beginningOfPassWriteIndex:2*L,endOfPassWriteIndex:2*L+1}},w={index:L,dispatches:0},a.passes.push(w)}let R=re.call(this,g);return w!==null&&B.set(R,w),R};let ue=K.copyBufferToBuffer,ze=ue,ie=function(p,g,w,R,L){if(S()&&(n.copyBufferToBuffer+=1),typeof g=="number"){if(!(w instanceof GPUBuffer)||R===void 0)throw new TypeError("copyBufferToBuffer: bad five-argument form");return ue.call(this,p,g,w,R,L)}if(w!==void 0&&typeof w!="number")throw new TypeError("copyBufferToBuffer: bad three-argument form");return ze.call(this,p,g,w)};K.copyBufferToBuffer=ie;let le=GPUComputePassEncoder.prototype,de=le.dispatchWorkgroups;le.dispatchWorkgroups=function(p,g,w){if(S()){n.dispatches+=1;let R=B.get(this);R!==void 0&&(R.dispatches+=1)}return de.call(this,p,g,w)};let Ee=le.dispatchWorkgroupsIndirect;le.dispatchWorkgroupsIndirect=function(p,g){if(S()){n.dispatches+=1,n.dispatchesIndirect+=1;let w=B.get(this);w!==void 0&&(w.dispatches+=1)}return Ee.call(this,p,g)};let fe=()=>{let p=e?performance.now():o,g=n.submitOffsets[0],w=n.submitOffsets[n.submitOffsets.length-1];return{dispatches:n.dispatches,dispatches_indirect:n.dispatchesIndirect,submits:n.submits,command_buffers:n.commandBuffers,passes:n.passes,bind_groups:n.bindGroups,write_buffers:n.writeBuffers,write_buffer_bytes:n.writeBufferBytes,create_buffers:n.createBuffers,create_buffer_bytes:n.createBufferBytes,destroy_buffers:n.destroyBuffers,pipelines:n.pipelines,pipelines_async:n.pipelinesAsync,pipeline_compile_ms:n.pipelineCompileMs,shader_modules:n.shaderModules,map_asyncs:n.mapAsyncs,copy_buffer_to_buffer:n.copyBufferToBuffer,first_submit_ms:g===void 0?-1:g,last_submit_ms:w===void 0?-1:w,submit_offsets_ms:[...n.submitOffsets],injected:n.injected,live_gpu_bytes:i,step_ms:p-t}},_={installed:!0,configure(p){u=[...p]},beginStep(){n=nt(),a!==null&&(a.next=0,a.passes=[]),t=performance.now(),o=t,e=!0},endStep(){return o=performance.now(),e=!1,fe()},snapshot:fe,setInjection(p){s=p},enableTimestamps(p,g){if(a!==null&&_.disableTimestamps(),g<1||g>2048)throw new Error(`timestamp capacity ${g} pairs outside 1..2048`);let w=p.createQuerySet({type:"timestamp",count:2*g}),R=p.createBuffer({size:16*g,usage:GPUBufferUsage.QUERY_RESOLVE|GPUBufferUsage.COPY_SRC}),L=p.createBuffer({size:16*g,usage:GPUBufferUsage.MAP_READ|GPUBufferUsage.COPY_DST});a={device:p,querySet:w,capacity:g,next:0,resolveBuffer:R,stagingBuffer:L,passes:[]}},disableTimestamps(){a!==null&&(a.querySet.destroy(),a.resolveBuffer.destroy(),a.stagingBuffer.destroy(),a=null)},async collectTimestamps(){let p=a;if(p===null||p.next===0)return[];let g=r;r=!0;try{let w=p.device.createCommandEncoder();w.resolveQuerySet(p.querySet,0,2*p.next,p.resolveBuffer,0),w.copyBufferToBuffer(p.resolveBuffer,0,p.stagingBuffer,0,16*p.next),p.device.queue.submit([w.finish()]),await p.stagingBuffer.mapAsync(GPUMapMode.READ,0,16*p.next);let R=new BigUint64Array(p.stagingBuffer.getMappedRange(0,16*p.next).slice(0));p.stagingBuffer.unmap();let L=[];for(let h of p.passes){let A=R[2*h.index],N=R[2*h.index+1];if(A===void 0||N===void 0)throw new Error("timestamp index out of range");L.push({index:h.index,dispatches:h.dispatches,begin_ns:Number(A),end_ns:Number(N)})}return L}finally{r=g}},devices:()=>c,descriptors:()=>f,errors:()=>k,lost:()=>y,liveBytes:()=>i,suspend(p){let g=r;r=!0;try{return p()}finally{r=g}},async suspendAsync(p){let g=r;r=!0;try{return await p()}finally{r=g}}};globalThis.__probeHooks=_})();var st={host:"127.0.0.1",port:47315,transformersBundle:"dist/transformers.js",transformersRoute:"/vendor/transformers/transformers.js",ortRoute:"/vendor/ort/",ortWasmFiles:["ort-wasm-simd-threaded.asyncify.mjs","ort-wasm-simd-threaded.asyncify.wasm"],modelsRoute:"/models/",weightsRoute:"/weights/",pageRoute:"/page.js",hooksRoute:"/hooks.js"};var Ze={layers:14,layerTypes:["conv","conv","attention","conv","attention","conv","attention","conv","attention","conv","attention","conv","attention","conv"],hidden:1024,ffn:2560,vocab:65536,heads:16,kvHeads:8,headDim:64,convCache:3,convBias:!1,normEps:1e-5,ropeTheta:1e6,ropeInterleaved:!1,tieEmbedding:!0,quantBits:4,quantBlock:32,maxPositionEmbeddings:128e3,bosTokenId:1,eosTokenId:7},Y={weightBits:4,residualBytes:4,hiddenBytes:2,kvCacheBytes:2,convCacheBytes:2,logitsBytes:4,accumulate:"f32"},gr={commandBuffersPerToken:1,passesPerToken:1,readbackEveryTokens:32,readbackRingTokens:64,concatQkv:!0,concatGateUp:!0},C={mlpFused:{workgroupSize:256,slices:80},attnSplit:{coreWorkgroupSize:256},flash:{coreWorkgroupSize:64,tilePositions:64,maxBlocks:32,positionsPerBlockRule:64,mergeWorkgroupSize:64,partialWords:68},fold:{workgroupSize:256,threadsPerRow:8},geometry:{base:{workgroupSize:128,subgroupSize:32,rowsPerSubgroup:4,subgroupsPerRow:1},wide:{workgroupSize:128,subgroupSize:32,rowsPerSubgroup:1,subgroupsPerRow:1}}};var at=[{key:"A2",splitN:!1,flash:!1,dispatches:74},{key:"A2N1",splitN:!0,flash:!1,dispatches:74},{key:"A2N2",splitN:!1,flash:!0,dispatches:80},{key:"A2N1N2",splitN:!0,flash:!0,dispatches:80}];var Xe={ropeRows:1024,ropeRowsExtended:2560,headPreScale:2,sectionAlignment:256,weightsFile:"weights.bin",manifestFile:"weights_manifest.json",maxQueries:1},T={matvec:{workgroupSize:128,subgroupSize:32,rowsPerSubgroup:4,subgroupsPerRow:1,maxCols:2560,quantBlock:32},matvecWide:{workgroupSize:128,subgroupSize:32,rowsPerSubgroup:1,subgroupsPerRow:1,maxCols:2560,quantBlock:32},conv:{workgroupSize:256},embed:{workgroupSize:256},sample:{workgroupSize:256,partials:256},maxPositions:2560};var it={maxStorageBuffersPerShaderStage:10,maxBindingsPerBindGroup:18,maxBindGroups:1,maxComputeWorkgroupStorageSize:23552,maxComputeWorkgroupStorageSizeFlash:18176,maxComputeInvocationsPerWorkgroup:256,maxComputeWorkgroupSizeX:256,maxComputeWorkgroupsPerDimension:4096};var pn={warmupSteps:20,m1Tokens:256,interleaveOrder:["A_arm","B_baseline","A_arm","B_baseline"],m1Blocks:16,m1BlocksAfterL4:8,m1PromptIndex:0,thermalRepeatBlocks:1,thermalDriftThreshold:.1,thermalCooldownSeconds:120,thermalMaxReruns:1,m3PerDispatchReps:64,mSyncRuns:4,m5RunsPerVariant:4,readbackVariants:[1,32],determinismRuns:3,determinismTokens:256,correctnessTokens:256,correctnessPrompts:8,correctnessTop1Min:.99,correctnessLogitStride:16,dequantLayerIndex:0,dequantMatrix:"conv.in_proj",dequantRowSamples:64,dequantSeed:111006,dequantTolerance:.005,deferredTolerance:.005,flashCoreTolerance:.005,flashCoreLayerIndex:2,flashCorePositions:[192,2048],splitNTolerance:.005,splitNLayerIndex:2,correctnessPositions:[192,1024],correctnessCompareTokens:256,bootstrapReplicates:1e4,seedTiming:111001,seedBootstrap:111002,seedSampling:111004,timestampCapacityPairs:2048,liveGpuBytesGrowthMaxPerBlock:16*1024*1024};var dn={chainCounts:[8,16,24,32,48,64,96,128],rounds:8,repsPerRound:25,warmupReps:20,workgroupSize:64,shapes:[{key:"attn_o_proj",mode:"matvec_residual",rows:1024,cols:1024,emits:1024,zeroPoints:!1,preScale:1,weightBytes:655360,dispatchBytes:663552},{key:"mlp_down",mode:"matvec_residual",rows:1024,cols:2560,emits:1024,zeroPoints:!1,preScale:1,weightBytes:1638400,dispatchBytes:1651712},{key:"attn_qkv",mode:"norm_projection",rows:2048,cols:1024,emits:2048,zeroPoints:!1,preScale:1,weightBytes:1310720,dispatchBytes:1327104},{key:"conv_in_proj",mode:"norm_matvec",rows:3072,cols:1024,emits:3072,zeroPoints:!1,preScale:1,weightBytes:1966080,dispatchBytes:1980416},{key:"mlp_gate_up",mode:"norm_swiglu",rows:5120,cols:1024,emits:2560,zeroPoints:!1,preScale:1,weightBytes:3276800,dispatchBytes:3290112},{key:"head",mode:"norm_head",rows:65536,cols:1024,emits:65536,zeroPoints:!0,preScale:Xe.headPreScale,weightBytes:42991616,dispatchBytes:43257856}],rowsPerWorkgroup:[1,2,4,8,16,32],rowsPerWorkgroupHead:[16,32,64,128,256,512,1024],workgroupSizes:[64,128,256],reachabilityRule:"subgroupsPerRow * subgroupSize <= cols / quantBlock",dispatchCounts:[4,16],footprintBytes:754974720,fillChunkBytes:67108864,curveRounds:6,curveTargetBytesPerRound:300*1024*1024,curveMinReps:3,curveMaxReps:25,curveWarmupReps:8,kneeFraction:.9,seed:111007};var Ri={totalSeconds:12*3600,l1MinRemainingBeforeP0b:10*3600,l2MinRemainingBeforeContext:4200,l3MinRemainingBeforeThirdArm:3*3600,l4MinRemainingBeforeM:4*3600,stageBudgetSeconds:{e1:300,p0:5700,weights:1500,preflight:1800,correctness:5400,m:13800,context:3600,finalize:600,closure:300}};function Ye(e,r,t,o,n){if(!Number.isInteger(e)||e<=0||!Number.isInteger(r)||r<=0)throw new Error(`quantMatrixBytes: bad shape ${e}x${r}`);if(r%n!==0)throw new Error(`quantMatrixBytes: cols ${r} not a multiple of block ${n}`);if(o!==4)throw new Error(`quantMatrixBytes: only 4-bit packing is registered, got ${o}`);let s=r/n,a=e*r*o/8,i=e*s*4,u=t?e*Math.ceil(s/2):0;return{quant:a,scales:i,zeroPoints:u,total:a+i+u}}function fn(e){if(!Number.isInteger(e)||e<3)throw new Error(`layerTypesFor: layers must be an integer >= 3, got ${e}`);let r=[];for(let t=0;t<e;t+=1)r.push(t>=2&&t%2===0?"attention":"conv");return r}function ut(e,r){let t=e.layerTypes[r];if(t===void 0)throw new Error(`layerMatrices: layer ${r} out of range`);let{hidden:o,ffn:n,kvHeads:s,headDim:a,quantBits:i,quantBlock:u}=e,c=s*a,f=(y,P,B,S)=>({name:P,layer:r,role:y,rows:B,cols:S,zeroPoints:!1,bytes:Ye(B,S,!1,i,u)}),k=[f("mlp.gate_proj",`/model/layers.${r}/mlp/gate_proj/MatMul_Q4`,n,o),f("mlp.up_proj",`/model/layers.${r}/mlp/up_proj/MatMul_Q4`,n,o),f("mlp.down_proj",`/model/layers.${r}/mlp/down_proj/MatMul_Q4`,o,n)];return t==="conv"?[f("conv.in_proj",`/model/layers.${r}/conv/in_proj/MatMul_Q4`,3*o,o),f("conv.out_proj",`/model/layers.${r}/conv/out_proj/MatMul_Q4`,o,o),...k]:[f("attn.q_proj",`/model/layers.${r}/attn/q_proj/MatMul_Q4`,o,o),f("attn.k_proj",`/model/layers.${r}/attn/k_proj/MatMul_Q4`,c,o),f("attn.v_proj",`/model/layers.${r}/attn/v_proj/MatMul_Q4`,c,o),f("attn.o_proj",`/model/layers.${r}/attn/o_proj/MatMul_Q4`,o,o),...k]}function lt(e){return{name:"logits",layer:-1,role:"lm_head",rows:e.vocab,cols:e.hidden,zeroPoints:!0,bytes:Ye(e.vocab,e.hidden,!0,e.quantBits,e.quantBlock)}}var Re={layerTypes:fn(14),hidden:1024,ffn:2560,vocab:65536,heads:16,kvHeads:8,headDim:64,convCache:3,quantBits:4,quantBlock:32};var ct=["attn_o_proj","mlp_down","attn_qkv","conv_in_proj","mlp_gate_up","head"];var mn={conv_in_proj:"conv_in_proj",conv_out_proj:"attn_o_proj",attn_qkv:"attn_qkv",attn_o_proj:"attn_o_proj",mlp_gate_up:"mlp_gate_up",mlp_down:"mlp_down",head:"head"},gn=["conv_in_proj","conv_out_proj","attn_o_proj","mlp_gate_up","mlp_down"];var zr=8;function hr(e){let{workgroupSize:r,subgroupSize:t,rowsPerSubgroup:o,subgroupsPerRow:n}=e;if(r%t!==0)throw new Error(`rowsPerWorkgroup: workgroup ${r} is not a multiple of subgroup ${t}`);let s=r/t;if(!Number.isInteger(n)||n<1)throw new Error(`rowsPerWorkgroup: subgroupsPerRow must be a positive integer, got ${n}`);if(s%n!==0)throw new Error(`rowsPerWorkgroup: ${s} subgroups do not partition into groups of ${n}`);if(!Number.isInteger(o)||o<1)throw new Error(`rowsPerWorkgroup: rowsPerSubgroup must be a positive integer, got ${o}`);return s/n*o}function pt(e,r,t){if(r%t!==0)throw new Error(`cellReachable: ${r} columns are not a multiple of the ${t}-wide block`);return e.subgroupsPerRow*e.subgroupSize<=r/t}function hn(e,r,t){return r.splitN&&gn.includes(e)?t.knee[mn[e]]:t.carried[e]}function dt(e,r,t){if(!Number.isInteger(e)||e<1)throw new Error(`blocksForPositions: ${e} is not a positive position count`);if(!Number.isInteger(r)||r<1)throw new Error(`blocksForPositions: ${r} is not a positive block length`);return Math.min(t,Math.max(1,Math.ceil(e/r)))}function Fe(e,r,t){let o=e[0];if(o===void 0)throw new Error("fuse: empty source list");for(let i of e){if(i.cols!==o.cols)throw new Error(`fuse: column mismatch ${i.name} ${i.cols} != ${o.cols}`);if(i.zeroPoints!==o.zeroPoints)throw new Error(`fuse: zero-point mismatch in ${i.name}`)}let n=e.reduce((i,u)=>i+u.rows,0),s=Ye(n,o.cols,o.zeroPoints,t,r),a=e.reduce((i,u)=>i+u.bytes.total,0);if(s.total!==a)throw new Error(`fuse: concatenated bytes ${s.total} != sum of sources ${a}`);return{sources:e,rows:n,cols:o.cols,zeroPoints:o.zeroPoints,bytes:s.total}}function Me(e,r){let t=e.find(o=>o.role===r);if(t===void 0)throw new Error(`plan: no matrix with role ${r}`);return t}function Ge(e,r,t,o,n,s){let a=Ye(r,t,o,n.quantBits,n.quantBlock).total;return{section:e,rows:r,cols:t,zeroPoints:o,bytes:a,kSlices:s}}function ft(e,r){return r.splitN?0:e.layerTypes.length}function bn(e,r){return ft(e,r)%2===0}function br(e,r){let{queries:t,position:o,precision:n,levers:s,slices:a,flash:i,geometry:u}=r;if(!Number.isInteger(t)||t<1)throw new Error(`planStep: queries must be a positive integer, got ${t}`);let c=!s.splitN;if(!bn(e,s))throw new Error(`planStep: ${ft(e,s)} fused blocks at ${e.layerTypes.length} layers is odd, so the residual ping-pong does not close and the token cannot be encoded once and replayed`);if(c&&t!==1)throw new Error(`planStep: the fused MLP is a single-query decode specialisation, got q = ${t}`);if(!Number.isInteger(i.blocks)||i.blocks<1||i.blocks>i.maxBlocks)throw new Error(`planStep: ${i.blocks} position blocks is outside 1..${i.maxBlocks}`);let{hidden:f,ffn:k,vocab:y,heads:P,headDim:B,kvHeads:S,convCache:v,quantBlock:$,quantBits:x}=e;if(P%S!==0)throw new Error(`planStep: ${P} q heads do not divide into ${S} kv heads`);let X=P/S,Z=S*B,I=f*n.residualBytes*t,q=f*n.hiddenBytes*t,ee=f*4,Ae=(P*B+2*S*B)*4*t,Pe=[],E={residual:0,partials:0,pending:0},ke=(_,p)=>Math.ceil(_/hr(p))*t,Q=(_,p,g,w,R,L,h,A,N,M,te,me,Oe,pr)=>{if(!Number.isInteger(M)||M<1)throw new Error(`planStep: ${_} has a non-positive workgroup count ${M}`);if(me===null!=(Oe===null))throw new Error(`planStep: ${_} carries a site without a geometry, or the reverse`);Pe.push({index:Pe.length,name:_,kind:p,layer:g,weight:w,reads:R,constantBytes:L,readBytes:h,writeBytes:A,fusion:N,workgroups:M,workgroupSize:te,site:me,geometry:Oe,attnPart:pr})},K=(_,p,g,w,R,L,h,A,N,M,te)=>{let me=hn(M,s,u);Q(_,p,g,w,R,L,h,A,re(),ke(N,me),me.workgroupSize,M,me,te)},re=()=>({foldSlices:0,slices:0,residualIn:E.residual,residualOut:E.residual,partialsIn:E.partials,partialsOut:E.partials}),ue=_=>{if(E.pending===0)return;let p=E.pending;Q(_,"fold",-1,null,[],0,I+p*ee,I,{foldSlices:p,slices:0,residualIn:E.residual,residualOut:E.residual,partialsIn:E.partials,partialsOut:E.partials},Math.ceil(f/(u.foldGeometry.workgroupSize/u.foldGeometry.threadsPerRow)),u.foldGeometry.workgroupSize,null,null,"none"),E.pending=0},ze=_=>{let p={foldSlices:E.pending,slices:_,residualIn:E.residual,residualOut:1-E.residual,partialsIn:E.partials,partialsOut:1-E.partials};return E.residual=1-E.residual,E.partials=1-E.partials,E.pending=_,p},ie=f/$,le=f/2+ie*4+Math.ceil(ie/2);Q("embed","embed",-1,null,[],le*t,4*t,I,re(),t,u.embedWorkgroupSize,null,null,"none");for(let _=0;_<e.layerTypes.length;_+=1){let p=e.layerTypes[_],g=ut(e,_),w=f*4,R=2*(B/2)*4*t;if(p==="conv"){let A=Fe([Me(g,"conv.in_proj")],$,x),N=Fe([Me(g,"conv.out_proj")],$,x),M=f*v*n.convCacheBytes;ue(`L${_}.conv.fold`),K(`L${_}.conv.in_proj`,"norm_matvec",_,A,[Ge(`L${_}.conv.in_proj`,A.rows,A.cols,!1,e,0)],w,I,3*q,A.rows,"conv_in_proj","none"),Q(`L${_}.conv.core`,"conv_core",_,null,[],f*v*4,3*q+M,q+M,re(),Math.ceil(f/u.convWorkgroupSize)*t,u.convWorkgroupSize,null,null,"none"),K(`L${_}.conv.out_proj`,"matvec_residual",_,N,[Ge(`L${_}.conv.out_proj`,N.rows,N.cols,!1,e,0)],0,q+I,I,N.rows,"conv_out_proj","none")}else{let A=Fe([Me(g,"attn.q_proj"),Me(g,"attn.k_proj"),Me(g,"attn.v_proj")],$,x),N=Fe([Me(g,"attn.o_proj")],$,x),M=o+t,te=2*Z*n.kvCacheBytes*t;if(ue(`L${_}.attn.fold`),K(`L${_}.attn.qkv`,"attn_proj",_,A,[Ge(`L${_}.attn.qkv`,A.rows,A.cols,!1,e,0)],w,I,Ae,A.rows,"attn_qkv","projection"),s.flash){let me=X*B*4*t,Oe=i.blocks*P*i.partialWords*4*t;Q(`L${_}.attn.flash`,"attn_core_flash",_,null,[],2*B*4+R,S*i.blocks*me+S*2*B*4*t+2*S*o*B*n.kvCacheBytes,Oe+te,re(),S*i.blocks*t,u.flashCoreWorkgroupSize,null,null,"core"),Q(`L${_}.attn.merge`,"attn_merge",_,null,[],0,Oe,q,re(),P*t,u.mergeWorkgroupSize,null,null,"merge")}else Q(`L${_}.attn.core`,"attn_core_qkv",_,null,[],2*B*4+R,Ae+2*S*M*B*n.kvCacheBytes,q+te,re(),S*t,u.attnCoreWorkgroupSize,null,null,"core");K(`L${_}.attn.o_proj`,"matvec_residual",_,N,[Ge(`L${_}.attn.o_proj`,N.rows,N.cols,!1,e,0)],0,q+I,I,N.rows,"attn_o_proj","out_projection")}let L=Fe([Me(g,"mlp.gate_proj"),Me(g,"mlp.up_proj")],$,x),h=Fe([Me(g,"mlp.down_proj")],$,x);if(s.splitN)ue(`L${_}.mlp.fold`),K(`L${_}.mlp.gate_up`,"norm_matvec_swiglu",_,L,[Ge(`L${_}.mlp.gate_up`,L.rows,L.cols,!1,e,0)],w,I,k*n.hiddenBytes*t,k,"mlp_gate_up","none"),K(`L${_}.mlp.down`,"matvec_residual",_,h,[Ge(`L${_}.mlp.down`,h.rows,h.cols,!1,e,0)],0,k*n.hiddenBytes*t+I,I,h.rows,"mlp_down","none");else{let A=ze(a.mlp);Q(`L${_}.mlp.fused`,"mlp_fused",_,null,[Ge(`L${_}.mlp.gate_up`,L.rows,L.cols,!1,e,0),Ge(`L${_}.mlp.down`,h.rows,h.cols,!1,e,a.mlp)],w,I+A.foldSlices*ee,I+a.mlp*ee,A,a.mlp*t,u.blockWorkgroupSize,null,null,"none")}}ue("head.fold");let de=Fe([lt(e)],$,x),Ee=y*n.logitsBytes*t,fe=r.samplePartials.plain;if(K("head","norm_head",-1,de,[Ge("head",de.rows,de.cols,!0,e,0)],f*4,I,Ee,de.rows,"head","none"),Q("sample.partial","sample_partial",-1,null,[],0,Ee,zr*fe,re(),fe,u.sampleWorkgroupSize,null,null,"none"),Q("sample.final","sample_final",-1,null,[],0,zr*fe,zr,re(),1,u.sampleWorkgroupSize,null,null,"none"),E.residual!==0||E.pending!==0)throw new Error(`planStep: token does not close on its starting parity (residual ${E.residual}, pending ${E.pending})`);return Pe}function wn(e){let r=at.find(t=>t.key===e);if(r===void 0)throw new Error(`armSpec: unknown arm ${e}`);return r}function yn(e){let r=wn(e);return{splitN:r.splitN,flash:r.flash}}function rr(e){let r=e.ffn/e.quantBlock;if(!Number.isInteger(r))throw new Error(`planSlicesFor: ffn ${e.ffn} must be a multiple of the quantisation block ${e.quantBlock}`);if(e.hidden%e.heads!==0)throw new Error(`planSlicesFor: hidden ${e.hidden} does not divide into ${e.heads} heads`);if(e.heads%e.kvHeads!==0)throw new Error(`planSlicesFor: ${e.heads} q heads do not divide into ${e.kvHeads} kv heads`);return{mlp:r}}function Or(e){return rr(e).mlp}var mt=rr(Re),Ui=Or(Re);if(C.mlpFused.slices!==mt.mlp)throw new Error(`plan_options: mlpFused registers ${C.mlpFused.slices} slices, the shape rule gives ${mt.mlp}`);if(C.flash.coreWorkgroupSize!==Re.headDim)throw new Error(`plan_options: the flash core registers workgroup ${C.flash.coreWorkgroupSize} against head dimension ${Re.headDim}`);if(C.flash.mergeWorkgroupSize!==Re.headDim)throw new Error(`plan_options: the merge registers workgroup ${C.flash.mergeWorkgroupSize} against head dimension ${Re.headDim}`);if(C.flash.partialWords<Re.headDim+2)throw new Error(`plan_options: ${C.flash.partialWords} partial words cannot hold ${Re.headDim} lanes plus a maximum and a denominator`);if(!gr.concatQkv||!gr.concatGateUp)throw new Error("plan_options: planStep has no unconcatenated form; FUSION registers both concatenations as true");var kn={plain:T.sample.partials},er={workgroupSize:T.matvec.workgroupSize,subgroupSize:T.matvec.subgroupSize,rowsPerSubgroup:T.matvec.rowsPerSubgroup,subgroupsPerRow:T.matvec.subgroupsPerRow},Ie={workgroupSize:T.matvecWide.workgroupSize,subgroupSize:T.matvecWide.subgroupSize,rowsPerSubgroup:T.matvecWide.rowsPerSubgroup,subgroupsPerRow:T.matvecWide.subgroupsPerRow},vn={conv_in_proj:er,conv_out_proj:er,attn_qkv:Ie,attn_o_proj:Ie,mlp_gate_up:Ie,mlp_down:Ie,head:er},Sn={attn_o_proj:Ie,mlp_down:Ie,attn_qkv:Ie,conv_in_proj:er,mlp_gate_up:Ie,head:er};function gt(e){let r=new Map;for(let o of e){if(r.has(o.shape))throw new Error(`knee table: ${o.shape} appears twice`);r.set(o.shape,o.geometry)}let t=o=>{let n=r.get(o);if(n===void 0)throw new Error(`knee table: no measured cell for ${o}`);return n};return{attn_o_proj:t("attn_o_proj"),mlp_down:t("mlp_down"),attn_qkv:t("attn_qkv"),conv_in_proj:t("conv_in_proj"),mlp_gate_up:t("mlp_gate_up"),head:t("head")}}function Pn(e,r){let t=e.flashOptima.find(o=>o.positions===r);if(t!==void 0){if(!Number.isInteger(t.blocks)||t.blocks<1||t.blocks>C.flash.maxBlocks)throw new Error(`flashBlocksFor: P0c optimum ${t.blocks} at ${r} positions is outside 1..${C.flash.maxBlocks}`);return t.blocks}return dt(r,C.flash.positionsPerBlockRule,C.flash.maxBlocks)}function ht(e){for(let r of ct){let t=e[r];if(t.workgroupSize%t.subgroupSize!==0)throw new Error(`planGeometry: knee for ${r} has workgroup ${t.workgroupSize} against subgroup ${t.subgroupSize}`)}return{carried:vn,knee:e,blockWorkgroupSize:C.mlpFused.workgroupSize,attnCoreWorkgroupSize:C.attnSplit.coreWorkgroupSize,flashCoreWorkgroupSize:C.flash.coreWorkgroupSize,mergeWorkgroupSize:C.flash.mergeWorkgroupSize,convWorkgroupSize:T.conv.workgroupSize,embedWorkgroupSize:T.embed.workgroupSize,sampleWorkgroupSize:T.sample.workgroupSize,foldGeometry:{workgroupSize:C.fold.workgroupSize,threadsPerRow:C.fold.threadsPerRow}}}var Fi=ht(Sn);function En(e,r){return{blocks:Pn(e,r),maxBlocks:C.flash.maxBlocks,partialWords:C.flash.partialWords}}function bt(e,r,t,o,n,s){return{queries:r,position:t,precision:{residualBytes:Y.residualBytes,hiddenBytes:Y.hiddenBytes,kvCacheBytes:Y.kvCacheBytes,convCacheBytes:Y.convCacheBytes,logitsBytes:Y.logitsBytes},levers:yn(n),slices:rr(e),flash:En(s,o),geometry:ht(s.knee),samplePartials:kn}}var _r=[];function _t(e){if(!e.length)throw new Error("No registered candidates");_r=e}function tr(e){let r=_r.find(t=>t.key===e);if(!r)throw new Error("Unknown candidate "+e);return r}function wt(e,r,t,o){return _r.map(n=>{let s=bt(e,1,r,t,n.unfused?"A2N1":n.flash?"A2N2":"A2",o),a={...s.geometry.carried};for(let[i,u]of Object.entries(n.matvec))a[i]=u.geometry;return{key:n.key,options:{...s,flash:{...s.flash,blocks:n.flashBlocks??s.flash.blocks},slices:{...s.slices,mlp:n.mlpSlices??s.slices.mlp},geometry:{...s.geometry,foldGeometry:{...s.geometry.foldGeometry,threadsPerRow:n.foldLanes??s.geometry.foldGeometry.threadsPerRow},carried:a,blockWorkgroupSize:n.mlpWorkgroup??s.geometry.blockWorkgroupSize}}}})}function yt(){return Math.max(80,..._r.map(e=>e.mlpSlices??80))}var tu={...st,port:47320};var{maxComputeWorkgroupStorageSizeFlash:Bn,...kt}=it,vt={...kt,maxComputeWorkgroupStorageSize:Math.max(kt.maxComputeWorkgroupStorageSize,Bn)};function Dr(e,r){return r.kind==="mlp_fused"?e.normOnceMlp??!1:["norm_matvec","attn_proj","norm_head"].includes(r.kind)?(e.normOnceMatvec||r.kind==="norm_head"&&e.normOnceHead)??!1:!1}function Cr(e,r,t,o,n){let s=[];for(let a of e){if(r.flashOnline&&a.kind==="attn_core_flash"){if(n!==1)throw new Error("Online flash accounting requires one decode query");let i=r.flashBlocks,u=t.heads*i*t.headDim*4*3,c=t.heads*2*o*t.headDim*2,f=t.heads*i*(t.headDim+2)*4+t.kvHeads*2*t.headDim*2;s.push({...a,index:s.length,readBytes:u+c,writeBytes:f});continue}if(!(r.convFuse&&a.kind==="conv_core")){if(r.convFuse&&a.kind==="norm_matvec"){s.push({...a,index:s.length,kind:"conv_proj_core",constantBytes:a.constantBytes+12288,readBytes:a.readBytes+6144,writeBytes:8192});continue}Dr(r,a)&&s.push({...a,index:s.length,name:a.name+".input_norm",kind:"input_norm",weight:null,reads:[],constantBytes:4096,readBytes:4096,writeBytes:4096,workgroups:1,site:null,geometry:null,fusion:{...a.fusion,slices:0,foldSlices:0}}),s.push({...a,index:s.length})}}return s}var xn=["maxBufferSize","maxStorageBufferBindingSize","maxStorageBuffersPerShaderStage","maxComputeWorkgroupStorageSize","maxComputeInvocationsPerWorkgroup","maxComputeWorkgroupSizeX","maxComputeWorkgroupsPerDimension","maxBindGroups","maxBindingsPerBindGroup"];async function Rn(e){if(navigator.gpu===void 0)throw new Error("navigator.gpu unavailable");let r=await navigator.gpu.requestAdapter({powerPreference:e});if(r===null)throw new Error("requestAdapter returned null");return r}function Mn(e,r,t){let o=e.descriptor.adapter;if(o.isFallbackAdapter)throw new Error("SOFTWARE_FALLBACK: adapter.info.isFallbackAdapter");if(o.subgroupMinSize!==t||o.subgroupMaxSize!==t)throw new Error(`subgroup size ${o.subgroupMinSize}..${o.subgroupMaxSize} != compiled ${t}`);for(let n of r.features){let s=Je(n);if(!e.device.features.has(s))throw new Error(`device lacks required feature ${s}`)}for(let[n,s]of Object.entries(r.limits)){let a=Reflect.get(e.device.limits,n);if(typeof a!="number")throw new Error(`device limit ${n} missing`);if(a<s)throw new Error(`device limit ${n} = ${a} < required ${s}`)}}async function Gn(e,r){let t=await Rn(r),o=[];for(let c of e.features){let f=Je(c);if(!t.features.has(f))throw new Error(`adapter lacks required feature ${f}`);o.push(f)}let n={};for(let c of xn){let f=t.limits[c];if(typeof f!="number")throw new Error(`adapter limit ${c} missing`);n[c]=f}let s=globalThis.__probeHooks,a=s.devices().length,i=await t.requestDevice({requiredFeatures:o,requiredLimits:n,label:"probe"}),u=s.descriptors()[a];if(u===void 0||s.devices()[a]!==i)throw new Error("hooks did not record the probe device");return{device:i,descriptor:u,owned:!0}}async function wr(e,r,t){let o=globalThis.__probeHooks,n=o.devices();if(n.length>1)throw new Error(`SHARED_DEVICE: ${n.length} GPUDevices exist; the session must hold one`);let s=n[0],a=o.descriptors()[0],i=s===void 0||a===void 0?await Gn(e,r):{device:s,descriptor:a,owned:!1};return Mn(i,e,t),i}async function je(e){await e.queue.onSubmittedWorkDone()}async function St(e,r,t,o){e.pushErrorScope("validation"),e.pushErrorScope("out-of-memory"),e.pushErrorScope("internal");try{return await o()}finally{let n=await e.popErrorScope(),s=await e.popErrorScope(),a=await e.popErrorScope();for(let i of[a,s,n])i!==null&&r.push(`${t}: ${i.constructor.name}: ${i.message}`)}}function $n(){let e=Number.POSITIVE_INFINITY,r=performance.now();for(let t=0;t<2e4;t+=1){let o=performance.now(),n=o-r;n>0&&n<e&&(e=n),r=o}return e}function Pt(){return{user_agent:navigator.userAgent,cross_origin_isolated:globalThis.crossOriginIsolated,timer_resolution_ms:$n(),hardware_concurrency:navigator.hardwareConcurrency,webgpu:navigator.gpu!==void 0}}async function or(e,r,t,o){let n=e.createBuffer({size:o,usage:GPUBufferUsage.MAP_READ|GPUBufferUsage.COPY_DST,label:"readback"}),s=e.createCommandEncoder();s.copyBufferToBuffer(r,t,n,0,o),e.queue.submit([s.finish()]),await n.mapAsync(GPUMapMode.READ,0,o);let a=n.getMappedRange(0,o).slice(0);return n.unmap(),n.destroy(),a}async function Ir(e,r,t){let o=t*2*8,n=e.createBuffer({size:o,usage:GPUBufferUsage.QUERY_RESOLVE|GPUBufferUsage.COPY_SRC,label:"timestamps"}),s=e.createCommandEncoder();s.resolveQuerySet(r,0,t*2,n,0),e.queue.submit([s.finish()]);let a=await or(e,n,0,o);n.destroy();let i=new BigUint64Array(a),u=new Float64Array(t*2);for(let c=0;c<u.length;c+=1){let f=i[c];u[c]=f===void 0?0:Number(f)}return u}function Et(e,r,t){let o=t==="gate"?2*r:e,n=t==="gate"?e/32:r/32;return{workgroupSize:256,bindings:[{binding:0,name:"weights",access:"read"},{binding:1,name:"destination",access:"read_write"}],code:`2 const ROWS:u32=${o}u;const BLOCKS:u32=${n}u;3 @group(0) @binding(0) var<storage,read> input_words:array<u32>;4 @group(0) @binding(1) var<storage,read_write> output_words:array<u32>;5 @compute @workgroup_size(256) fn main(@builtin(global_invocation_id) gid:vec3<u32>){6 let i=gid.x;if(i>=ROWS*BLOCKS*4u){return;}${t==="gate"?"let row=i/(BLOCKS*4u);let word=(i/BLOCKS)%4u;let block=i%BLOCKS;let at=(row*BLOCKS+block)*4u+word;":"let block=i/(ROWS*4u);let word=(i/ROWS)%4u;let row=i%ROWS;let at=(block*ROWS+row)*4u+word;"}output_words[i]=input_words[at];7 }`}}function He(e){return e==="v1_vec4_f16"||e==="v3_both"}function Ln(e){return e==="v2_vec4_accumulator"||e==="v3_both"}function Rt(e,r){return{workgroupSize:e.workgroupSize,subgroupSize:e.subgroupSize,rowsPerSubgroup:e.rowsPerSubgroup,subgroupsPerRow:e.subgroupsPerRow,maxCols:r}}var he=`enable f16;8enable subgroups;9`,be=`10struct State {11 position: u32,12 step: u32,13 queries: u32,14 sample_mode: u32,15 rng: u32,16 freerun: u32,17 advance: u32,18 reserved: u32,19}20`,An={weights:"array<u32>",scales:"array<f32>",zero_points:"array<u32>",gamma:"array<f32>",residual:"array<f32>",source:"array<f16>",destination:"array<f16>",state:"State",tokens:"array<u32>",cos:"array<f32>",sin:"array<f32>",key_cache:"array<f16>",value_cache:"array<f16>",cache:"array<f16>",logits:"array<f32>",partials:"array<u32>",sampled:"array<u32>",blob:"array<u32>",residual_in:"array<f32>",residual_out:"array<f32>",partials_in:"array<f32>",partials_out:"array<f32>",projection:"array<f32>",flash:"array<f32>"};function _e(e,r){return e.map(t=>{let o=r[t.name],n=o===void 0?An[t.name]:o;return`@group(0) @binding(${t.binding}) var<storage, ${t.access}> ${t.name}: ${n};`}).join(`21`)}var Tr=`22fn dot8(packed: u32, zp: f32, a0: vec4<f32>, a1: vec4<f32>) -> f32 {23 let lo = vec4<f32>(f32(packed & 15u), f32((packed >> 4u) & 15u), f32((packed >> 8u) & 15u), f32((packed >> 12u) & 15u));24 let hi = vec4<f32>(f32((packed >> 16u) & 15u), f32((packed >> 20u) & 15u), f32((packed >> 24u) & 15u), f32((packed >> 28u) & 15u));25 let z = vec4<f32>(zp, zp, zp, zp);26 return dot(lo - z, a0) + dot(hi - z, a1);27}28`,nr=`29fn dot8h(packed: u32, zp: f16, a0: vec4<f16>, a1: vec4<f16>) -> f16 {30 let lo = vec4<f16>(f16(packed & 15u), f16((packed >> 4u) & 15u), f16((packed >> 8u) & 15u), f16((packed >> 12u) & 15u));31 let hi = vec4<f16>(f16((packed >> 16u) & 15u), f16((packed >> 20u) & 15u), f16((packed >> 24u) & 15u), f16((packed >> 28u) & 15u));32 let z = vec4<f16>(zp, zp, zp, zp);33 return dot(lo - z, a0) + dot(hi - z, a1);34}35`,zn=`36fn hash(seed: u32, index: u32) -> u32 {37 var x = seed ^ (index * 2654435761u);38 x = x ^ (x >> 16u);39 x = x * 2246822519u;40 x = x ^ (x >> 13u);41 x = x * 3266489917u;42 return x ^ (x >> 16u);43}44 45fn gumbel(seed: u32, index: u32) -> f32 {46 let u = (f32(hash(seed, index) >> 8u) + 0.5) * ${1/16777216};47 return -log(-log(u));48}49`;function On(e,r){return e?`50const ZP_BYTES_PER_ROW: u32 = ${Math.ceil(r/2)}u;51fn zeroPoint(row: u32, block: u32) -> f32 {52 let index = row * ZP_BYTES_PER_ROW + block / 2u;53 let byte = (zero_points[index / 4u] >> ((index % 4u) * 8u)) & 255u;54 if ((block & 1u) == 0u) { return f32(byte & 15u); }55 return f32(byte >> 4u);56}57`:`fn zeroPoint(row: u32, block: u32) -> f32 { return 8.0; }58`}function Dn(e){return`59fn stageInput(query: u32, local: u32) {60 var sum: f32 = 0.0;61 for (var i = local; i < COLS; i = i + WG) {62 let v = residual[query * COLS + i] * ${e.toFixed(1)};63 stage[i] = v;64 sum = sum + v * v;65 }66 reduce[local] = sum;67 workgroupBarrier();68 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {69 if (local < stride) { reduce[local] = reduce[local] + reduce[local + stride]; }70 workgroupBarrier();71 }72 let inv = inverseSqrt(reduce[0] / f32(COLS) + EPSILON);73 workgroupBarrier();74 for (var i = local; i < COLS; i = i + WG) {75 stage[i] = stage[i] * inv * gamma[i];76 }77 workgroupBarrier();78}79`}var Cn=`80fn stageInput(query: u32, local: u32) {81 for (var i = local; i < COLS; i = i + WG) {82 stage[i] = f32(source[query * COLS + i]);83 }84 workgroupBarrier();85}86`;function In(e){return e===1?"lane + i * SG":"slice * SG + lane + i * (SG * SUBGROUPS_PER_ROW)"}function Mt(e){return e===0?"stage[o]":`stage[o + ${e}u]`}function Bt(e){return`vec4<f32>(${[0,1,2,3].map(r=>Mt(e+r)).join(", ")})`}function xt(e){return`vec4<f16>(${[0,1,2,3].map(r=>`f16(${Mt(e+r)})`).join(", ")})`}function Tn(e,r){let t=r*8,o=r*8+4;return He(e)?`dot8h(packed.${["x","y","z","w"][r]}, zpv,87 ${xt(t)},88 ${xt(o)})`:`dot8(${r===0?"weights[base]":`weights[base + ${r}u]`}, zpv,89 ${Bt(t)},90 ${Bt(o)})`}function qn(e,r){let t=r===1?"row: u32, lane: u32":"row: u32, slice: u32, lane: u32",o=He(e)?` let packed = weights[row * BLOCKS + block];91 let zpv = f16(zeroPoint(row, block));92 let o = block * 32u;`:` let base = row * BLOCKS * 4u + block * 4u;93 let zpv = zeroPoint(row, block);94 let o = block * 32u;`,n=[0,1,2,3].map(c=>Tn(e,c)),s,a,i;Ln(e)?(s=" var acc: vec4<f32> = vec4<f32>(0.0, 0.0, 0.0, 0.0);",a=` let d = vec4<f32>(95 ${n.map(c=>He(e)?`f32(${c})`:c).join(`,96 `)});97 acc = acc + d * scales[row * BLOCKS + block];`,i=` let total = subgroupAdd(acc);98 return total.x + total.y + total.z + total.w;`):(s=" var acc: f32 = 0.0;",a=` var part: ${He(e)?"f16 = f16(0.0)":"f32 = 0.0"};99${n.map(c=>` part = part + ${c};`).join(`100`)}101 acc = acc + ${He(e)?"f32(part)":"part"} * scales[row * BLOCKS + block];`,i=" return subgroupAdd(acc);");let u=`102fn rowSum(${t}) -> f32 {103${s}104 for (var i = 0u; i < BLOCK_ITERS; i = i + 1u) {105 let block = ${In(r)};106 if (block < BLOCKS) {107${o}108${a}109 }110 }111${i}112}113`;return r===1?u:`${u}114fn rowTotal(row: u32, subgroup: u32, slice: u32, lane: u32) -> f32 {115 let mine = rowSum(row, slice, lane);116 if (lane == 0u) { cross[subgroup] = mine; }117 workgroupBarrier();118 let base = subgroup - slice;119 var total: f32 = 0.0;120 for (var p = 0u; p < SUBGROUPS_PER_ROW; p = p + 1u) {121 total = total + cross[base + p];122 }123 workgroupBarrier();124 return total;125}126`}function Nn(e,r,t,o){let n=e==="norm_swiglu"?t.ffn:r;return Math.ceil(n/hr(o))}function Gt(e,r,t,o,n){return[Nn(e,r,t,o),n,1]}function $t(e,r,t,o,n,s,a,i){let{workgroupSize:u,subgroupSize:c,rowsPerSubgroup:f,subgroupsPerRow:k}=n;if(u%c!==0)throw new Error(`matvecKernel: workgroup ${u} is not a multiple of subgroup ${c}`);if(t%s.quantBlock!==0)throw new Error(`matvecKernel: cols ${t} is not a multiple of block ${s.quantBlock}`);if(t>n.maxCols)throw new Error(`matvecKernel: cols ${t} exceeds the registered stage width ${n.maxCols}`);if(e==="norm_swiglu"&&r!==2*s.ffn)throw new Error(`matvecKernel: swiglu expects ${2*s.ffn} rows, got ${r}`);if(e==="norm_projection"&&r%s.headDim!==0)throw new Error(`matvecKernel: a q/k/v projection emits whole heads, ${r} rows is not a multiple of ${s.headDim}`);if(!pt(n,t,s.quantBlock))throw new Error(`matvecKernel: ${k} subgroups of ${c} lanes cannot each take a block of a ${t}-wide row`);let y=hr(n),P=u/c,B=t/s.quantBlock,S=e==="norm_swiglu"?s.ffn:r,v=k>1;if(v&&f!==1)throw new Error(`matvecKernel: cooperating subgroups own one row each, got ${f} per subgroup`);let $=[{binding:0,name:"weights",access:"read"},{binding:1,name:"scales",access:"read"}];o&&$.push({binding:2,name:"zero_points",access:"read"}),e==="matvec_residual"?($.push({binding:4,name:"residual",access:"read_write"}),$.push({binding:5,name:"source",access:"read"})):($.push({binding:3,name:"gamma",access:"read"}),$.push({binding:4,name:"residual",access:"read"})),e==="norm_head"?$.push({binding:16,name:"logits",access:"read_write"}):e==="norm_projection"?$.push({binding:17,name:"projection",access:"read_write"}):e!=="matvec_residual"&&$.push({binding:6,name:"destination",access:"read_write"}),$.push({binding:7,name:"state",access:"read"});let x=He(i)?{weights:"array<vec4<u32>>"}:{},X=e==="norm_head"?"logits[query * EMITS + row] = value;":e==="norm_projection"?"projection[query * EMITS + row] = value;":e==="matvec_residual"?"residual[query * EMITS + row] = residual[query * EMITS + row] + value;":"destination[query * EMITS + row] = f16(value);",Z=v?"rowTotal(ROWEXPR, subgroup, slice, lane)":"rowSum(ROWEXPR, lane)",I=K=>Z.replace("ROWEXPR",K),q=v?"lane == 0u && slice == 0u":"lane == 0u",ee=e==="norm_swiglu"?`127 for (var r = 0u; r < ROWS_PER_SG; r = r + 1u) {128 let want = rowBase + r;129 let row = min(want, EMITS - 1u);130 let gate = ${I("row")};131 let up = ${I("row + EMITS")};132 if (${q} && want < EMITS) {133 destination[query * EMITS + row] = f16((gate / (1.0 + exp(-gate))) * up);134 }135 }136`:` for (var r = 0u; r < ROWS_PER_SG; r = r + 1u) {137 let want = rowBase + r;138 let row = min(want, EMITS - 1u);139 let value = ${I("row")};140 if (${q} && want < EMITS) {141 ${X}142 }143 }144`,pe=He(i)?nr:Tr,Ae=v?`var<workgroup> cross: array<f32, ${P}>;145`:"",Pe=v?"let rowBase = wid.x * ROWS_PER_WG + (subgroup / SUBGROUPS_PER_ROW) * ROWS_PER_SG;":"let rowBase = wid.x * ROWS_PER_WG + (local / SG) * ROWS_PER_SG;",E=v?` let subgroup = local / SG;146 let slice = subgroup % SUBGROUPS_PER_ROW;147`:"",ke=v?`const SUBGROUPS_PER_ROW: u32 = ${k}u;148`:"";return{code:`${he}${be}${pe}149const COLS: u32 = ${t}u;150const BLOCKS: u32 = ${B}u;151const BLOCK_ITERS: u32 = ${Math.ceil(B/(c*k))}u;152const WG: u32 = ${u}u;153const SG: u32 = ${c}u;154${ke}const ROWS_PER_SG: u32 = ${f}u;155const ROWS_PER_WG: u32 = ${y}u;156const EMITS: u32 = ${S}u;157const EPSILON: f32 = ${s.epsilon};158 159${_e($,x)}160 161var<workgroup> stage: array<f32, ${t}>;162var<workgroup> reduce: array<f32, ${u}>;163${Ae}${On(o,B)}${e==="matvec_residual"?Cn:Dn(a)}${qn(i,k)}164@compute @workgroup_size(WG)165fn main(166 @builtin(workgroup_id) wid: vec3<u32>,167 @builtin(local_invocation_index) local: u32,168 @builtin(subgroup_invocation_id) lane: u32,169) {170 let query = wid.y;171 if (query >= state.queries) { return; }172 stageInput(query, local);173${E} ${Pe}174${ee}}175`,bindings:$,workgroupSize:u}}function Lt(e,r){let t=e.hidden/e.quantBlock,o=[{binding:0,name:"weights",access:"read"},{binding:1,name:"scales",access:"read"},{binding:2,name:"zero_points",access:"read"},{binding:4,name:"residual",access:"read_write"},{binding:7,name:"state",access:"read"},{binding:8,name:"tokens",access:"read"}];return{code:`${he}${be}176const HIDDEN: u32 = ${e.hidden}u;177const BLOCK: u32 = ${e.quantBlock}u;178const BLOCKS: u32 = ${t}u;179const U32_PER_ROW: u32 = ${e.hidden/8}u;180const ZP_BYTES_PER_ROW: u32 = ${Math.ceil(t/2)}u;181const WG: u32 = ${r}u;182 183${_e(o,{})}184 185@compute @workgroup_size(WG)186fn main(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) local: u32) {187 let query = wid.x;188 if (query >= state.queries) { return; }189 let token = tokens[state.step + query];190 for (var i = local; i < HIDDEN; i = i + WG) {191 let block = i / BLOCK;192 let packed = weights[token * U32_PER_ROW + i / 8u];193 let nibble = (packed >> ((i % 8u) * 4u)) & 15u;194 let zpIndex = token * ZP_BYTES_PER_ROW + block / 2u;195 let zpByte = (zero_points[zpIndex / 4u] >> ((zpIndex % 4u) * 8u)) & 255u;196 var zpv: u32 = zpByte & 15u;197 if ((block & 1u) == 1u) { zpv = zpByte >> 4u; }198 residual[query * HIDDEN + i] = (f32(nibble) - f32(zpv)) * scales[token * BLOCKS + block];199 }200}201`,bindings:o,workgroupSize:r}}function At(e){return[e,1,1]}function zt(e,r){if(e.convTaps!==3)throw new Error(`convKernel: registered for 3 taps, got ${e.convTaps}`);let t=[{binding:3,name:"gamma",access:"read"},{binding:5,name:"source",access:"read"},{binding:6,name:"destination",access:"read_write"},{binding:7,name:"state",access:"read"},{binding:9,name:"cache",access:"read_write"}];return{code:`${he}${be}202const HIDDEN: u32 = ${e.hidden}u;203const TAPS: u32 = ${e.convTaps}u;204const WG: u32 = ${r}u;205 206${_e(t,{})}207 208@compute @workgroup_size(WG)209fn main(@builtin(global_invocation_id) gid: vec3<u32>) {210 let channel = gid.x;211 if (channel >= HIDDEN) { return; }212 let w0 = gamma[channel * TAPS];213 let w1 = gamma[channel * TAPS + 1u];214 let w2 = gamma[channel * TAPS + 2u];215 var s1 = f32(cache[channel * TAPS + 1u]);216 var s2 = f32(cache[channel * TAPS + 2u]);217 for (var q = 0u; q < state.queries; q = q + 1u) {218 let base = q * 3u * HIDDEN;219 let bx = f32(source[base + channel]) * f32(source[base + 2u * HIDDEN + channel]);220 let y = w0 * s1 + w1 * s2 + w2 * bx;221 destination[q * HIDDEN + channel] = f16(f32(source[base + HIDDEN + channel]) * y);222 s1 = s2;223 s2 = bx;224 }225 cache[channel * TAPS] = f16(0.0);226 cache[channel * TAPS + 1u] = f16(s1);227 cache[channel * TAPS + 2u] = f16(s2);228}229`,bindings:t,workgroupSize:r}}function Ot(e,r){return[Math.ceil(e.hidden/r),1,1]}function Dt(e,r,t){if(e.vocab%t!==0)throw new Error(`samplePartialKernel: vocab ${e.vocab} is not a multiple of ${t}`);let o=[{binding:7,name:"state",access:"read"},{binding:15,name:"partials",access:"read_write"},{binding:16,name:"logits",access:"read"}];return{code:`${he}${be}230const VOCAB: u32 = ${e.vocab}u;231const WG: u32 = ${r}u;232const PER_GROUP: u32 = ${e.vocab/t}u;233 234${_e(o,{})}235 236var<workgroup> bestValue: array<f32, ${r}>;237var<workgroup> bestIndex: array<u32, ${r}>;238${zn}239@compute @workgroup_size(WG)240fn main(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) local: u32) {241 let offset = (state.queries - 1u) * VOCAB;242 var value: f32 = -3.0e38;243 var index: u32 = 0u;244 for (var i = wid.x * PER_GROUP + local; i < (wid.x + 1u) * PER_GROUP; i = i + WG) {245 var s = logits[offset + i];246 if (state.sample_mode == 1u) { s = s + gumbel(state.rng, i); }247 if (s > value) { value = s; index = i; }248 }249 bestValue[local] = value;250 bestIndex[local] = index;251 workgroupBarrier();252 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {253 if (local < stride) {254 if (bestValue[local + stride] > bestValue[local]) {255 bestValue[local] = bestValue[local + stride];256 bestIndex[local] = bestIndex[local + stride];257 }258 }259 workgroupBarrier();260 }261 if (local == 0u) {262 partials[wid.x * 2u] = bitcast<u32>(bestValue[0]);263 partials[wid.x * 2u + 1u] = bestIndex[0];264 }265}266`,bindings:o,workgroupSize:r}}function Ct(e,r,t){let o=[{binding:7,name:"state",access:"read_write"},{binding:8,name:"tokens",access:"read_write"},{binding:15,name:"partials",access:"read"},{binding:17,name:"sampled",access:"read_write"}];return{code:`${he}${be}267const WG: u32 = ${r}u;268const PARTIALS: u32 = ${t}u;269const MAX_POSITIONS: u32 = ${e.maxPositions}u;270 271${_e(o,{})}272 273var<workgroup> bestValue: array<f32, ${r}>;274var<workgroup> bestIndex: array<u32, ${r}>;275 276@compute @workgroup_size(WG)277fn main(@builtin(local_invocation_index) local: u32) {278 var value: f32 = -3.0e38;279 var index: u32 = 0u;280 for (var i = local; i < PARTIALS; i = i + WG) {281 let candidate = bitcast<f32>(partials[i * 2u]);282 if (candidate > value) { value = candidate; index = partials[i * 2u + 1u]; }283 }284 bestValue[local] = value;285 bestIndex[local] = index;286 workgroupBarrier();287 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {288 if (local < stride) {289 if (bestValue[local + stride] > bestValue[local]) {290 bestValue[local] = bestValue[local + stride];291 bestIndex[local] = bestIndex[local + stride];292 }293 }294 workgroupBarrier();295 }296 if (local == 0u) {297 let token = bestIndex[0];298 let last = state.step + state.queries - 1u;299 sampled[last] = token;300 if (state.freerun == 1u && last + 1u < MAX_POSITIONS) {301 tokens[last + 1u] = token;302 }303 if (state.advance == 1u) {304 state.position = state.position + state.queries;305 state.step = state.step + state.queries;306 state.rng = state.rng * 1664525u + 1013904223u;307 }308 }309}310`,bindings:o,workgroupSize:r}}function It(e){return[e,1,1]}function Tt(){return[1,1,1]}function qt(e,r,t){if(!t)return e;if(r>32)throw new Error("Subgroup merge supports at most 32 splits");let o=e.code.indexOf("@compute");if(o<0)throw new Error("Merge source seam");let n=e.code.slice(0,o)+`@compute @workgroup_size(WG)311 fn main(@builtin(workgroup_id) wid:vec3<u32>,@builtin(local_invocation_index) local:u32,@builtin(subgroup_invocation_id) lane:u32){312 let first=(wid.y*SPLITS*HEADS+wid.x)*RECORD;313 var mine=NEG_INF;if(lane<SPLITS){mine=flash[first+lane*HEADS*RECORD+HEAD_DIM];}314 let peak=subgroupMax(mine);var weight=0.0;var den=0.0;315 if(lane<SPLITS){weight=exp(mine-peak);den=weight*flash[first+lane*HEADS*RECORD+HEAD_DIM+1u];}316 let denominator=subgroupAdd(den);var numerator=0.0;317 for(var s=0u;s<SPLITS;s++){numerator+=subgroupShuffle(weight,s)*flash[first+s*HEADS*RECORD+local];}318 destination[(wid.y*HEADS+wid.x)*HEAD_DIM+local]=f16(numerator/denominator);319 }`;return{...e,code:n}}function Nt(e,r,t){if(e.headDim!==64||e.heads/e.kvHeads!==2)throw new Error("Online flash requires 64-d GQA2");return{workgroupSize:64,bindings:[{binding:0,name:"blob",access:"read"},{binding:5,name:"state",access:"read"},{binding:6,name:"key_cache",access:"read_write"},{binding:7,name:"value_cache",access:"read_write"},{binding:8,name:"projection",access:"read"},{binding:10,name:"flash",access:"read_write"}],code:`enable f16;320 enable subgroups;321 ${be}322 const SPLITS:u32=${t}u;const MAX_POSITIONS:u32=${e.maxPositions}u;323 @group(0) @binding(0) var<storage,read> blob:array<u32>;324 @group(0) @binding(5) var<storage,read> state:State;325 @group(0) @binding(6) var<storage,read_write> key_cache:array<f16>;326 @group(0) @binding(7) var<storage,read_write> value_cache:array<f16>;327 @group(0) @binding(8) var<storage,read> projection:array<f32>;328 @group(0) @binding(10) var<storage,read_write> flash:array<f32>;329 @compute @workgroup_size(64)330 fn main(@builtin(workgroup_id) wid:vec3<u32>,@builtin(local_invocation_index) local:u32,@builtin(subgroup_invocation_id) lane:u32){331 let kvHead=wid.x/SPLITS;let split=wid.x%SPLITS;let group=local/32u;let head=kvHead*2u+group;332 let position=subgroupBroadcastFirst(state.position)+wid.y;let total=position+1u;333 let blockLen=(total+SPLITS-1u)/SPLITS;let start=split*blockLen;let stop=min(start+blockLen,total);let owner=(total-1u)/blockLen;334 let base=wid.y*2048u;335 var qlo=projection[base+head*64u+lane];var qhi=projection[base+head*64u+lane+32u];336 let qinv=inverseSqrt(subgroupAdd(qlo*qlo+qhi*qhi)/64.0+${e.epsilon});337 qlo=qlo*qinv*bitcast<f32>(blob[${r.qNorm}u+lane]);qhi=qhi*qinv*bitcast<f32>(blob[${r.qNorm}u+lane+32u]);338 let co=bitcast<f32>(blob[${r.cos}u+position*32u+lane]);let si=bitcast<f32>(blob[${r.sin}u+position*32u+lane]);339 let qrot=qlo*co-qhi*si;qhi=f32(f16(qhi*co+qlo*si));qlo=f32(f16(qrot));340 var klo=projection[base+1024u+kvHead*64u+lane];var khi=projection[base+1024u+kvHead*64u+lane+32u];341 let kinv=inverseSqrt(subgroupAdd(klo*klo+khi*khi)/64.0+${e.epsilon});342 klo=klo*kinv*bitcast<f32>(blob[${r.kNorm}u+lane]);khi=khi*kinv*bitcast<f32>(blob[${r.kNorm}u+lane+32u]);343 let krot=klo*co-khi*si;khi=f32(f16(khi*co+klo*si));klo=f32(f16(krot));344 let vlo=f32(f16(projection[base+1536u+kvHead*64u+lane]));let vhi=f32(f16(projection[base+1536u+kvHead*64u+lane+32u]));345 if(split==owner&&group==0u){let at=(kvHead*MAX_POSITIONS+position)*64u+lane;346 key_cache[at]=f16(klo);key_cache[at+32u]=f16(khi);value_cache[at]=f16(vlo);value_cache[at+32u]=f16(vhi);347 }348 var peak=-3.0e38;var denom=0.0;var nlo=0.0;var nhi=0.0;349 for(var p=start;p<stop;p++){350 var kl=klo;var kh=khi;var vl=vlo;var vh=vhi;351 if(p!=position){let at=(kvHead*MAX_POSITIONS+p)*64u+lane;kl=f32(key_cache[at]);kh=f32(key_cache[at+32u]);vl=f32(value_cache[at]);vh=f32(value_cache[at+32u]);}352 let score=subgroupAdd(qlo*kl+qhi*kh)*0.125;353 let next=max(peak,score);let alpha=exp(peak-next);let beta=exp(score-next);354 denom=denom*alpha+beta;nlo=nlo*alpha+vl*beta;nhi=nhi*alpha+vh*beta;peak=next;355 }356 let rec=((wid.y*SPLITS+split)*16u+head)*68u;357 flash[rec+lane]=nlo;flash[rec+lane+32u]=nhi;358 if(lane==0u){flash[rec+64u]=peak;flash[rec+65u]=denom;}359 }`}}function qr(e,r){if(!r)return e;let t=e.code;for(let o of["dot8","dot8h"]){if(!t.includes("fn "+o+"("))continue;let n=o==="dot8h",s=n?"f16":"f32",a=Array.from({length:4},(i,u)=>{let c=`(((packed >> ${u*4}u)&0x000f000fu)|0x64006400u)`,f=n?`bitcast<vec2<f16>>(${c})`:`unpack2x16float(${c})`;return`let p${u}=${f}-vec2<${s}>(${s}(1024.0)+zp);`}).join(`360`);t=Ke(t,o,`fn ${o}(packed:u32,zp:${s},a0:vec4<${s}>,a1:vec4<${s}>)->${s}{${a}return dot(vec4<${s}>(p0.x,p1.x,p2.x,p3.x),a0)+dot(vec4<${s}>(p0.y,p1.y,p2.y,p3.y),a1);}`)}return{...e,code:t}}function Nr(e,r){if(!r||!e.code.includes("fn dot8("))return e;let t=Ke(e.code,"dot8",`fn dot8(packed:u32,zp:f32,a0:vec4<f32>,a1:vec4<f32>)->f32{361 let mask=vec4<u32>(15u,240u,3840u,61440u);362 let factor=vec4<f32>(1.0,16.0,256.0,4096.0);let inv=vec4<f32>(1.0,0.0625,0.00390625,0.000244140625);363 let lo=vec4<f32>(vec4<u32>(packed)&mask)-vec4<f32>(zp)*factor;364 let hi=vec4<f32>(vec4<u32>(packed>>16u)&mask)-vec4<f32>(zp)*factor;365 return dot(lo,a0*inv)+dot(hi,a1*inv);366 }`);return{...e,code:t}}function Ke(e,r,t){let o=e.indexOf("fn "+r+"(");if(o<0)throw new Error("Missing function "+r);let n=e.indexOf("{",o)+1,s=1;for(;s&&n<e.length;)e[n]==="{"&&s++,e[n]==="}"&&s--,n++;if(s)throw new Error("Unclosed function "+r);return e.slice(0,o)+t+e.slice(n)}function Wt(e){let r=0;if(e=e.replace(/\bstage\[o(?:\s*\+\s*(\d+)u)?\]/g,(t,o)=>(r++,`laneBlock[${o??"0"}u]`)),r!==32)throw new Error("Private input source seam "+r);return e+`367var<private> laneBlock:array<f32,32>;368`}function Ut(e,r,t,o){if(!o)return e;if(!e.code.includes("const COLS: u32 = 1024u;"))throw new Error("Private matvec requires 1024 columns");let n=r==="matvec_residual"?"":`var square=0.0;for(var i=lane;i<COLS;i+=SG){let v=residual[query*COLS+i]*${t.toFixed(1)};square+=v*v;}let inv=inverseSqrt(subgroupAdd(square)/f32(COLS)+EPSILON);`,s=r==="matvec_residual"?"f32(source[query*COLS+at])":`(residual[query*COLS+at]*${t.toFixed(1)})*inv*gamma[at]`,a=Ke(e.code,"stageInput",`fn stageInput(query:u32,local:u32){let lane=local%SG;${n}for(var i=0u;i<32u;i++){let at=lane*32u+i;laneBlock[i]=${s};}}`);return{...e,code:Wt(a)}}function Ft(e,r){return r?{...e,code:Ke(e.code,"stageInput","fn stageInput(query:u32,local:u32){for(var i=local;i<COLS;i+=WG){stage[i]=residual[query*COLS+i];}workgroupBarrier();}")}:e}function jt(e,r){if(!r)return e;let t=e.code,o=`@compute @workgroup_size(WG)369 fn main(@builtin(workgroup_id) wid:vec3<u32>,@builtin(local_invocation_index) local:u32,@builtin(subgroup_invocation_id) lane:u32){370 let query=wid.y;if(query>=state.queries){return;}stageInput(query,local);371 let rowBase=wid.x*(ROWS_PER_WG/3u)+(local/SG)*(ROWS_PER_SG/3u);372 for(var r=0u;r<ROWS_PER_SG/3u;r++){373 let want=rowBase+r;let channel=min(want,1023u);374 let b=f32(f16(rowSum(channel,lane)));let c=f32(f16(rowSum(channel+1024u,lane)));let x=f32(f16(rowSum(channel+2048u,lane)));375 if(lane==0u&&want<1024u){376 let s1=f32(cache[channel*3u+1u]);let s2=f32(cache[channel*3u+2u]);let bx=b*x;377 let w0=bitcast<f32>(blob[channel*3u]);let w1=bitcast<f32>(blob[channel*3u+1u]);let w2=bitcast<f32>(blob[channel*3u+2u]);378 let y=w0*s1+w1*s2+w2*bx;destination[query*1024u+channel]=f16(c*y);379 cache[channel*3u]=f16(0.0);cache[channel*3u+1u]=f16(s2);cache[channel*3u+2u]=f16(bx);380 }381 }382 }`,n=t.indexOf("@compute");if(n<0)throw new Error("Conv projection source seam");return t=t.slice(0,n)+o+`383@group(0) @binding(18) var<storage,read> blob:array<u32>;384@group(0) @binding(9) var<storage,read_write> cache:array<f16>;385`,{...e,code:t,bindings:[...e.bindings,{binding:18,name:"blob",access:"read"},{binding:9,name:"cache",access:"read_write"}]}}function Ht(e,r){return{workgroupSize:e,bindings:[{binding:0,name:"residual",access:"read"},{binding:1,name:"gamma",access:"read"},{binding:2,name:"projection",access:"read_write"}],code:`enable subgroups;386 const WG:u32=${e}u;387 @group(0) @binding(0) var<storage,read> residual:array<f32>;388 @group(0) @binding(1) var<storage,read> gamma:array<f32>;389 @group(0) @binding(2) var<storage,read_write> projection:array<f32>;390 var<workgroup> squares:array<f32,${e/32}>;391 @compute @workgroup_size(WG) fn main(@builtin(local_invocation_index) local:u32){392 var sum=0.0;for(var i=local;i<1024u;i+=WG){let v=residual[i]*${r.toFixed(1)};sum+=v*v;}393 let sg=subgroupAdd(sum);if(local%32u==0u){squares[local/32u]=sg;}workgroupBarrier();394 var total=0.0;for(var g=0u;g<WG/32u;g++){total+=squares[g];}395 let inv=inverseSqrt(total/1024.0+0.00001);396 for(var i=local;i<1024u;i+=WG){projection[i]=(residual[i]*${r.toFixed(1)})*inv*gamma[i];}397 }`}}function Wr(e,r){if(!r)return e;let t=0,o=e.code.replace(/\bstage\[([^\]]+)\]/g,(n,s)=>(t++,`stage[stageIndex(${s})]`));if(t<8)throw new Error("Stage swizzle source seam");return{...e,code:o+`398fn stageIndex(i:u32)->u32{return i ^ ((i >> 5u) & 31u);}399`}}function Kt(e,r){let t=e.code;if(r.mlpVecLoad){if(r.mlpPair||r.projectPair)throw new Error("Vector-load arm uses unpaired baseline arithmetic");let o=0;if(t=t.replace(/var part: f32 = 0\.0;/g,n=>`let packed = weights[base / 4u];400 `+n),t=t.replace(/blob\[base(?: \+ ([123])u)?\]/g,(n,s)=>(o++,"packed."+["x","y","z","w"][Number(s??0)])),o!==8)throw new Error("MLP vector load seam "+o);t+=`401@group(0) @binding(19) var<storage,read> weights:array<vec4<u32>>;402`,e={...e,bindings:[...e.bindings,{binding:19,name:"weights",access:"read"}]}}if(r.normOnceMlp&&(t=Ke(t,"stageBlock",`fn stageBlock(local:u32,group:u32){403 let rowBegin=group*HIDDEN/SLICES;let rowEnd=(group+1u)*HIDDEN/SLICES;404 for(var i=rowBegin+local;i<rowEnd;i+=WG){residual_out[i]=residual_in[i];}405 for(var i=local;i<HIDDEN;i+=WG){stage[i]=projection[i];}workgroupBarrier();406 }`),t+=`407@group(0) @binding(8) var<storage,read> projection:array<f32>;408`,e={...e,bindings:[...e.bindings,{binding:8,name:"projection",access:"read"}]}),r.mlpPrivate){if(r.mlpPair||!t.includes("const FOLD_SLICES: u32 = 0u;"))throw new Error("Private MLP requires a plain no-fold input");t=Ke(t,"stageBlock",`fn stageBlock(local:u32,group:u32){409 let lane=local%SG;var square=0.0;410 for(var i=lane;i<HIDDEN;i+=SG){let v=residual_in[i];square+=v*v;}411 let inv=inverseSqrt(subgroupAdd(square)/f32(HIDDEN)+EPSILON);412 let rowBegin=group*HIDDEN/SLICES;let rowEnd=(group+1u)*HIDDEN/SLICES;413 for(var i=rowBegin+local;i<rowEnd;i+=WG){residual_out[i]=residual_in[i];}414 for(var i=0u;i<32u;i++){let at=lane*32u+i;laneBlock[i]=residual_in[at]*inv*bitcast<f32>(blob[NORM+at]);}415 }`),t=Wt(t)}if(r.mlpPair){let o=t.indexOf("fn rowSum("),n=t.indexOf("const KQUANT:",o);if(o<0||n<0)throw new Error("PAIR source seam");let s=Array.from({length:4},(a,i)=>{let u=i*8,c=Array.from({length:4},(k,y)=>`stage[o+${u+y}u]`).join(","),f=Array.from({length:4},(k,y)=>`stage[o+${u+y+4}u]`).join(",");return`let x${i}=vec4<f32>(${c}); let y${i}=vec4<f32>(${f});416part.x += dot8(blob[base+${i}u],8.0,x${i},y${i});417part.y += dot8(blob[upbase+${i}u],8.0,x${i},y${i});`}).join(`418`);t=t.slice(0,o)+`fn rowPair(row:u32,lane:u32)->vec2<f32>{419 var acc=vec2<f32>(0.0);420 for(var i=0u;i<BLOCK_ITERS;i++){421 let block=lane+i*SG;422 if(block<BLOCKS){let base=QUANT+row*BLOCKS*4u+block*4u;let upbase=base+FFN*BLOCKS*4u;let o=block*32u;var part=vec2<f32>(0.0);423 ${s}424 acc += part*vec2<f32>(bitcast<f32>(blob[SCALES+row*BLOCKS+block]),bitcast<f32>(blob[SCALES+(FFN+row)*BLOCKS+block]));}425 }426 return subgroupAdd(acc);427}428`+t.slice(n),t=t.replace(`let gate = rowSum(rowBase + r, lane);429 let up = rowSum(FFN + rowBase + r, lane);`,`let pair = rowPair(rowBase + r, lane);430 let gate = pair.x; let up = pair.y;`)}if(r.projectPair){let o=t.indexOf("fn project("),n=t.indexOf("@compute",o);if(o<0||n<0)throw new Error("PROJECT source seam");let s=Array.from({length:4},(a,i)=>{let u=Array.from({length:4},(f,k)=>`hidden_slice[o+${i*8+k}u]`).join(","),c=Array.from({length:4},(f,k)=>`hidden_slice[o+${i*8+k+4}u]`).join(",");return`let x${i}=vec4<f32>(${u});let y${i}=vec4<f32>(${c});part.x+=dot8(blob[base+${i}u],8.0,x${i},y${i});part.y+=dot8(blob[base2+${i}u],8.0,x${i},y${i});`}).join(`431`);t=t.slice(0,o)+`fn project(local:u32,group:u32){432 let sliceBase=group*HIDDEN;433 for(var row=local;row<HIDDEN;row+=2u*WG){var acc=vec2<f32>(0.0);434 for(var b=0u;b<BLOCKS_PER_SLICE;b++){let base=KQUANT+(sliceBase+row)*WORDS_PER_ROW_SLICE+b*4u;let base2=base+WG*WORDS_PER_ROW_SLICE;let o=b*32u;var part=vec2<f32>(0.0);${s}435 acc+=part*vec2<f32>(bitcast<f32>(blob[KSCALES+(sliceBase+row)*BLOCKS_PER_SLICE+b]),bitcast<f32>(blob[KSCALES+(sliceBase+row+WG)*BLOCKS_PER_SLICE+b]));}436 partials_out[sliceBase+row]=acc.x;partials_out[sliceBase+row+WG]=acc.y;}437}438`+t.slice(n)}if(r.mlpDotIlp){if(!r.mlpPair||r.projectPair||r.mlpVecLoad)throw new Error("MLP ILP requires paired gate and plain down");if(r.mlpDotIlp==="gate"||r.mlpDotIlp==="both"){let o=0;if(t=t.replace("var part=vec2<f32>(0.0);","var gatePart=vec4<f32>(0.0);var upPart=vec4<f32>(0.0);"),t=t.replace(/part\.([xy]) \+= dot8\(blob\[(base|upbase)\+(\d)u\],8\.0,x(\d),y(\d)\);/g,(n,s,a,i,u,c)=>{if(i!==u||i!==c)throw new Error("Gate ILP lane mismatch");return o++,`${s==="x"?"gatePart":"upPart"}.${["x","y","z","w"][Number(i)]}=dot8(blob[${s==="x"?"base":"upbase"}+${i}u],8.0,x${i},y${i});`}),o!==8)throw new Error("Gate ILP source seam "+o);t=t.replace("acc += part*vec2<f32>","let part=vec2<f32>((gatePart.x+gatePart.y)+(gatePart.z+gatePart.w),(upPart.x+upPart.y)+(upPart.z+upPart.w));acc += part*vec2<f32>")}if(r.mlpDotIlp==="down"||r.mlpDotIlp==="both"){let o=Array.from({length:4},(n,s)=>{let a=Array.from({length:4},(u,c)=>`hidden_slice[o+${s*8+c}u]`).join(","),i=Array.from({length:4},(u,c)=>`hidden_slice[o+${s*8+c+4}u]`).join(",");return`dot8(blob[base+${s}u],8.0,vec4<f32>(${a}),vec4<f32>(${i}))`}).join(",");t=Ke(t,"project",`fn project(local:u32,group:u32){439 let sliceBase=group*HIDDEN;440 for(var row=local;row<HIDDEN;row+=WG){var acc=0.0;441 for(var b=0u;b<BLOCKS_PER_SLICE;b++){let base=KQUANT+(sliceBase+row)*WORDS_PER_ROW_SLICE+b*4u;let o=b*32u;442 let parts=vec4<f32>(${o});let part=(parts.x+parts.y)+(parts.z+parts.w);443 acc+=part*bitcast<f32>(blob[KSCALES+(sliceBase+row)*BLOCKS_PER_SLICE+b]);444 }partials_out[sliceBase+row]=acc;445 }446 }`)}}if(r.mlpHalfProduct){t+=`447fn dot8halfProduct(packed:u32,zp:f32,a0:vec4<f32>,a1:vec4<f32>)->f32{448 let lo=vec4<f16>(f16(packed&15u),f16((packed>>4u)&15u),f16((packed>>8u)&15u),f16((packed>>12u)&15u))-vec4<f16>(f16(zp));449 let hi=vec4<f16>(f16((packed>>16u)&15u),f16((packed>>20u)&15u),f16((packed>>24u)&15u),f16((packed>>28u)&15u))-vec4<f16>(f16(zp));450 let p0=vec4<f32>(lo*vec4<f16>(a0));let p1=vec4<f32>(hi*vec4<f16>(a1));451 return dot(p0,vec4<f32>(1.0))+dot(p1,vec4<f32>(1.0));452 }`;let o=t.indexOf(r.mlpPair?"fn rowPair(":"fn rowSum("),n=t.indexOf("fn project("),s=t.indexOf("@compute",n);if(o<0||n<0||s<0)throw new Error("Half product source seam");(r.mlpHalfProduct==="down"||r.mlpHalfProduct==="both")&&(t=t.slice(0,n)+t.slice(n,s).replaceAll("dot8(","dot8halfProduct(")+t.slice(s)),(r.mlpHalfProduct==="gate"||r.mlpHalfProduct==="both")&&(t=t.slice(0,o)+t.slice(o,n).replaceAll("dot8(","dot8halfProduct(")+t.slice(n))}if(r.mlpHalfDot){t+=nr+`453fn dot8half(packed:u32,zp:f32,a0:vec4<f32>,a1:vec4<f32>)->f32{454 return f32(dot8h(packed,f16(zp),vec4<f16>(a0),vec4<f16>(a1)));455 }`;let o=t.indexOf(r.mlpPair?"fn rowPair(":"fn rowSum("),n=t.indexOf("fn project("),s=t.indexOf("@compute",n);if(o<0||n<0||s<0)throw new Error("Half dot source seam");(r.mlpHalfDot==="down"||r.mlpHalfDot==="both")&&(t=t.slice(0,n)+t.slice(n,s).replaceAll("dot8(","dot8half(")+t.slice(s)),(r.mlpHalfDot==="gate"||r.mlpHalfDot==="both")&&(t=t.slice(0,o)+t.slice(o,n).replaceAll("dot8(","dot8half(")+t.slice(n))}if(r.mlpSoA){if(r.mlpVecLoad||r.mlpSlices===160)throw new Error("SoA requires ordinary whole quant blocks and no competing weight binding");let o=t.indexOf(r.mlpPair?"fn rowPair(":"fn rowSum("),n=t.indexOf("const KQUANT:",o),s=t.indexOf("fn project("),a=t.indexOf("@compute",s);if(o<0||n<0||s<0||a<0)throw new Error("SoA source seam");if(r.mlpSoA==="down"||r.mlpSoA==="both"){let i=0,u=t.slice(s,a).replace(/blob\[(base|base2)(?:\s*\+\s*(\d)u)?\]/g,(c,f,k)=>(i++,`down_soa[((group*BLOCKS_PER_SLICE+b)*4u+${k??0}u)*HIDDEN+row${f==="base2"?"+WG":""}]`));if(i!==(r.projectPair?8:4))throw new Error("SoA down reads "+i);t=t.slice(0,s)+u+t.slice(a),t+=`456@group(0) @binding(20) var<storage,read> down_soa:array<u32>;457`,e={...e,bindings:[...e.bindings,{binding:20,name:"source",access:"read"}]}}if(r.mlpSoA==="gate"||r.mlpSoA==="both"){let i=0,u=t.slice(o,n).replace(/blob\[(base|upbase)(?:\s*\+\s*(\d)u)?\]/g,(c,f,k)=>(i++,`gate_soa[(${f==="upbase"?"row+FFN":"row"})*BLOCKS*4u+${k??0}u*BLOCKS+block]`));if(i!==(r.mlpPair?8:4))throw new Error("SoA gate reads "+i);t=t.slice(0,o)+u+t.slice(n),t+=`458@group(0) @binding(19) var<storage,read> gate_soa:array<u32>;459`,e={...e,bindings:[...e.bindings,{binding:19,name:"weights",access:"read"}]}}}if(r.mlpTranspose){if(r.mlpStageSwizzle||r.mlpPrivate)throw new Error("Transposed stage cannot combine with alternate stage addressing");let o=0;if(t=t.replace(/\bstage\[o(?:\s*\+\s*(\d+)u)?\]/g,(n,s)=>(o++,`stage[block+${Number(s??0)*32}u]`)),o!==32)throw new Error("Transposed MLP input seam "+o);t=t.replace(/\bstage\[i\]/g,"stage[((i&31u)<<5u)+(i>>5u)]")}return qr(Nr(Wr({...e,code:t},r.mlpStageSwizzle??!1),r.maskMlp??!1),r.magicMlp??!1)}function Vt(e,r){if(r.foldSubgroup){let n=e.code.indexOf("var<workgroup> part:");if(n<0)throw new Error("Fold subgroup seam");let s=e.code.slice(0,n)+`var<workgroup> part:array<f32,${e.workgroupSize}>;460 @compute @workgroup_size(WG)461 fn main(@builtin(workgroup_id) wid:vec3<u32>,@builtin(local_invocation_index) local:u32){462 let slot=local%ROWS_PER_WG;let lane=local/ROWS_PER_WG;let row=wid.x*ROWS_PER_WG+slot;463 var acc=0.0;if(row<HIDDEN){for(var s=lane;s<FOLD_SLICES;s+=LANES){acc+=partials_in[s*HIDDEN+row];}}464 for(var delta=ROWS_PER_WG;delta<32u;delta*=2u){acc+=subgroupShuffleXor(acc,delta);}465 if(local%32u<ROWS_PER_WG){part[(local/32u)*ROWS_PER_WG+slot]=acc;}workgroupBarrier();466 if(local<ROWS_PER_WG&&row<HIDDEN){var sum=residual_out[row];for(var sg=0u;sg<WG/32u;sg++){sum+=part[sg*ROWS_PER_WG+local];}residual_out[row]=sum;}467 }`;return{...e,code:s}}if(!r.foldFast)return e;let t=e.code.indexOf("var<workgroup> part:");if(t<0)throw new Error("FOLD source seam");let o=e.code.slice(0,t)+`@compute @workgroup_size(WG)468 fn main(@builtin(workgroup_id) wid:vec3<u32>,@builtin(local_invocation_index) local:u32){469 if(local>=ROWS_PER_WG){return;}let row=wid.x*ROWS_PER_WG+local;if(row>=HIDDEN){return;}470 var sum=residual_out[row];for(var lane=0u;lane<LANES;lane++){var acc=0.0;for(var s=lane;s<FOLD_SLICES;s+=LANES){acc+=partials_in[s*HIDDEN+row];}sum+=acc;}residual_out[row]=sum;471 }`;return{...e,code:o}}var j={blob:0,residualIn:1,residualOut:2,partialsIn:3,partialsOut:4,state:5,keyCache:6,valueCache:7,projection:8,destination:9,flash:10};function Wn(e,r,t,o){let{workgroupSize:n,subgroupSize:s,slices:a,foldSlices:i}=t;if(n%s!==0)throw new Error(`${e}: workgroup ${n} is not a multiple of subgroup ${s}`);if(o%a!==0)throw new Error(`${e}: ${o} is not divisible by ${a} slices`);let u=o/a,c=n/s;if(u%c!==0)throw new Error(`${e}: ${u} rows per slice is not a multiple of the ${c} rows a workgroup sums at once`);if(r.hidden%r.quantBlock!==0)throw new Error(`${e}: hidden ${r.hidden} is not a multiple of block ${r.quantBlock}`);if(!Number.isInteger(i)||i<0)throw new Error(`${e}: fold slices must be a non-negative integer, got ${i}`)}function Qt(e,r){let t=s=>s===0?`${e}[o]`:`${e}[o + ${s}u]`,o=s=>`vec4<f32>(${t(s)}, ${t(s+1)}, ${t(s+2)}, ${t(s+3)})`,n=[];for(let s=0;s<4;s+=1){let a=s===0?"base":`base + ${s}u`;n.push(`${r}part = part + dot8(blob[${a}], 8.0,472${r} ${o(s*8)},473${r} ${o(s*8+4)});`)}return n.join(`474`)}function Un(e,r,t,o){return`475const QUANT: u32 = ${e}u;476const SCALES: u32 = ${r}u;477const BLOCKS: u32 = ${t}u;478const BLOCK_ITERS: u32 = ${Math.ceil(t/o)}u;479 480fn rowSum(row: u32, lane: u32) -> f32 {481 var acc: f32 = 0.0;482 for (var i = 0u; i < BLOCK_ITERS; i = i + 1u) {483 let block = lane + i * SG;484 if (block < BLOCKS) {485 let base = QUANT + row * BLOCKS * 4u + block * 4u;486 let o = block * 32u;487 var part: f32 = 0.0;488${Qt("stage"," ")}489 acc = acc + part * bitcast<f32>(blob[SCALES + row * BLOCKS + block]);490 }491 }492 return subgroupAdd(acc);493}494`}function Fn(e,r,t,o){return`495const KQUANT: u32 = ${e}u;496const KSCALES: u32 = ${r}u;497const BLOCKS_PER_SLICE: u32 = ${t}u;498const WORDS_PER_ROW_SLICE: u32 = ${t*4}u;499 500fn project(local: u32, group: u32) {501 let sliceBase = group * HIDDEN;502 for (var row = local; row < HIDDEN; row = row + WG) {503 var acc: f32 = 0.0;504 for (var b = 0u; b < BLOCKS_PER_SLICE; b = b + 1u) {505 let base = KQUANT + (sliceBase + row) * WORDS_PER_ROW_SLICE + b * 4u;506 let o = b * 32u;507 var part: f32 = 0.0;508${Qt(o," ")}509 acc = acc + part * bitcast<f32>(blob[KSCALES + (sliceBase + row) * BLOCKS_PER_SLICE + b]);510 }511 partials_out[sliceBase + row] = acc;512 }513}514`}function jn(e,r,t){let o=t===0?" let v = residual_in[i];":` var v = residual_in[i];515 for (var s = 0u; s < FOLD_SLICES; s = s + 1u) { v = v + partials_in[s * HIDDEN + i]; }`;return`516const NORM: u32 = ${r}u;517const FOLD_SLICES: u32 = ${t}u;518const EPSILON: f32 = ${e.epsilon};519 520fn stageBlock(local: u32, group: u32) {521 var sum: f32 = 0.0;522 for (var i = local; i < HIDDEN; i = i + WG) {523${o}524 stage[i] = v;525 sum = sum + v * v;526 }527 reduce[local] = sum;528 workgroupBarrier();529 let rowBegin = (group * HIDDEN) / SLICES;530 let rowEnd = ((group + 1u) * HIDDEN) / SLICES;531 for (var i = rowBegin + local; i < rowEnd; i = i + WG) {532 residual_out[i] = stage[i];533 }534 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {535 if (local < stride) { reduce[local] = reduce[local] + reduce[local + stride]; }536 workgroupBarrier();537 }538 let inv = inverseSqrt(reduce[0] / f32(HIDDEN) + EPSILON);539 workgroupBarrier();540 for (var i = local; i < HIDDEN; i = i + WG) {541 stage[i] = stage[i] * inv * bitcast<f32>(blob[NORM + i]);542 }543 workgroupBarrier();544}545`}function Hn(e,r){return`const HIDDEN: u32 = ${e.hidden}u;546const WG: u32 = ${r.workgroupSize}u;547const SG: u32 = ${r.subgroupSize}u;548const SLICES: u32 = ${r.slices}u;549const ROWS_AT_ONCE: u32 = ${r.workgroupSize/r.subgroupSize}u;550`}function Kn(e){let r=[{binding:j.blob,name:"blob",access:"read"},{binding:j.residualIn,name:"residual_in",access:"read"},{binding:j.residualOut,name:"residual_out",access:"read_write"}];return e.foldSlices>0&&r.push({binding:j.partialsIn,name:"partials_in",access:"read"}),r.push({binding:j.partialsOut,name:"partials_out",access:"read_write"}),r}function Xt(e,r,t){let o=e.heads/e.kvHeads;return`const HEAD_DIM: u32 = ${e.headDim}u;551const HALF_DIM: u32 = ${e.headDim/2}u;552const GROUP: u32 = ${o}u;553const GROUP_ROWS: u32 = ${o+1}u;554const K_ROW: u32 = ${o*e.headDim}u;555const KV_WIDTH: u32 = ${e.kvHeads*e.headDim}u;556const MAX_POSITIONS: u32 = ${e.maxPositions}u;557const POS_GROUPS: u32 = ${t/e.headDim}u;558const SCALE: f32 = ${1/Math.sqrt(e.headDim)};559const Q_NORM: u32 = ${r.qNorm}u;560const K_NORM: u32 = ${r.kNorm}u;561const COS: u32 = ${r.cos}u;562const SIN: u32 = ${r.sin}u;563`}function Yt(e,r,t){if(r.heads%r.kvHeads!==0)throw new Error(`${e}: ${r.heads} q heads is not a multiple of ${r.kvHeads} kv heads`);if(t%r.headDim!==0)throw new Error(`${e}: workgroup ${t} is not a multiple of head dim ${r.headDim}`);if(r.headDim%2!==0)throw new Error(`${e}: head dim ${r.headDim} does not split into the two halves the rotation pairs`);return r.heads/r.kvHeads}function Jt(){return`564fn headRows(local: u32, position: u32) {565 for (var which = 0u; which < GROUP_ROWS; which = which + 1u) {566 let normBase = select(K_NORM, Q_NORM, which < GROUP);567 let off = which * HEAD_DIM;568 reduce[local] = 0.0;569 workgroupBarrier();570 if (local < HEAD_DIM) { reduce[local] = qk[off + local] * qk[off + local]; }571 workgroupBarrier();572 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {573 if (local < stride) { reduce[local] = reduce[local] + reduce[local + stride]; }574 workgroupBarrier();575 }576 let inv = inverseSqrt(reduce[0] / f32(HEAD_DIM) + EPSILON);577 workgroupBarrier();578 if (local < HEAD_DIM) { qk[off + local] = qk[off + local] * inv * bitcast<f32>(blob[normBase + local]); }579 workgroupBarrier();580 if (local < HALF_DIM) {581 let c = bitcast<f32>(blob[COS + position * HALF_DIM + local]);582 let s = bitcast<f32>(blob[SIN + position * HALF_DIM + local]);583 let lo = qk[off + local];584 let hi = qk[off + local + HALF_DIM];585 reduce[local] = lo * c - hi * s;586 reduce[local + HALF_DIM] = hi * c + lo * s;587 }588 workgroupBarrier();589 if (local < HEAD_DIM) { qk[off + local] = f32(f16(reduce[local])); }590 workgroupBarrier();591 }592}593`}function Vn(e,r){return`${Jt()}594fn attnCore(local: u32, kvHead: u32, position: u32, queryIndex: u32) {595 let total = position + 1u;596 headRows(local, position);597 598 // One workgroup per KV head, so the current slot is written exactly once and read by nobody: the599 // walks below score and weight the current position against the k and v still in registers.600 if (local < HEAD_DIM) {601 key_cache[(kvHead * MAX_POSITIONS + position) * HEAD_DIM + local] = f16(qk[K_ROW + local]);602 value_cache[(kvHead * MAX_POSITIONS + position) * HEAD_DIM + local] = f16(vcur[local]);603 }604 605 var best: array<f32, ${e}>;606 for (var g = 0u; g < GROUP; g = g + 1u) { best[g] = -3.0e38; }607 for (var p = local; p < total; p = p + WG) {608 var acc: array<f32, ${e}>;609 for (var g = 0u; g < GROUP; g = g + 1u) { acc[g] = 0.0; }610 let base = (kvHead * MAX_POSITIONS + p) * HEAD_DIM;611 for (var d = 0u; d < HEAD_DIM; d = d + 1u) {612 var kd: f32 = qk[K_ROW + d];613 if (p != position) { kd = f32(key_cache[base + d]); }614 for (var g = 0u; g < GROUP; g = g + 1u) { acc[g] = acc[g] + qk[g * HEAD_DIM + d] * kd; }615 }616 for (var g = 0u; g < GROUP; g = g + 1u) {617 let s = acc[g] * SCALE;618 scores[g * MAX_POSITIONS + p] = s;619 best[g] = max(best[g], s);620 }621 }622 623 var denom: array<f32, ${e}>;624 for (var g = 0u; g < GROUP; g = g + 1u) {625 reduce[local] = best[g];626 workgroupBarrier();627 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {628 if (local < stride) { reduce[local] = max(reduce[local], reduce[local + stride]); }629 workgroupBarrier();630 }631 let maximum = reduce[0];632 workgroupBarrier();633 var sum: f32 = 0.0;634 for (var p = local; p < total; p = p + WG) {635 let e = exp(scores[g * MAX_POSITIONS + p] - maximum);636 scores[g * MAX_POSITIONS + p] = e;637 sum = sum + e;638 }639 reduce[local] = sum;640 workgroupBarrier();641 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {642 if (local < stride) { reduce[local] = reduce[local] + reduce[local + stride]; }643 workgroupBarrier();644 }645 denom[g] = reduce[0];646 workgroupBarrier();647 }648 649 // One pass over the value cache for all GROUP heads, weights held in registers, then one fold per650 // head through the shared array in position-group order so the sum is fixed.651 let dim = local % HEAD_DIM;652 let posGroup = local / HEAD_DIM;653 var weighted: array<f32, ${e}>;654 for (var g = 0u; g < GROUP; g = g + 1u) { weighted[g] = 0.0; }655 for (var p = posGroup; p < total; p = p + POS_GROUPS) {656 var vd: f32 = vcur[dim];657 if (p != position) { vd = f32(value_cache[(kvHead * MAX_POSITIONS + p) * HEAD_DIM + dim]); }658 for (var g = 0u; g < GROUP; g = g + 1u) { weighted[g] = weighted[g] + scores[g * MAX_POSITIONS + p] * vd; }659 }660 for (var g = 0u; g < GROUP; g = g + 1u) {661 workgroupBarrier();662 vacc[posGroup * HEAD_DIM + dim] = weighted[g];663 workgroupBarrier();664 if (local < HEAD_DIM) {665 var acc: f32 = 0.0;666 for (var pg = 0u; pg < POS_GROUPS; pg = pg + 1u) { acc = acc + vacc[pg * HEAD_DIM + local]; }667 ${r}668 }669 }670 workgroupBarrier();671}672`}function Zt(e,r,t){let o=Yt("attnCoreQkvKernel",e,t),n=[{binding:j.blob,name:"blob",access:"read"},{binding:j.state,name:"state",access:"read"},{binding:j.keyCache,name:"key_cache",access:"read_write"},{binding:j.valueCache,name:"value_cache",access:"read_write"},{binding:j.projection,name:"projection",access:"read"},{binding:j.destination,name:"destination",access:"read_write"}];return{code:`${he}${be}673const HIDDEN: u32 = ${e.hidden}u;674const HEADS: u32 = ${e.heads}u;675const WG: u32 = ${t}u;676const EPSILON: f32 = ${e.epsilon};677const PROJ_ROWS: u32 = ${e.hidden+2*e.kvHeads*e.headDim}u;678${Xt(e,r,t)}679${_e(n,{})}680 681var<workgroup> reduce: array<f32, ${t}>;682var<workgroup> scores: array<f32, ${o*e.maxPositions}>;683var<workgroup> qk: array<f32, ${(o+1)*e.headDim}>;684var<workgroup> vcur: array<f32, ${e.headDim}>;685var<workgroup> vacc: array<f32, ${t}>;686${Vn(o,"destination[(queryIndex * HEADS + kvHead * GROUP + g) * HEAD_DIM + local] = f16(acc / denom[g]);")}687@compute @workgroup_size(WG)688fn main(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) local: u32) {689 let kvHead = wid.x;690 let queryIndex = wid.y;691 let position = state.position + queryIndex;692 let base = queryIndex * PROJ_ROWS;693 if (local < HEAD_DIM) {694 for (var g = 0u; g < GROUP; g = g + 1u) {695 qk[g * HEAD_DIM + local] = projection[base + (kvHead * GROUP + g) * HEAD_DIM + local];696 }697 qk[K_ROW + local] = projection[base + HIDDEN + kvHead * HEAD_DIM + local];698 vcur[local] = f32(f16(projection[base + HIDDEN + KV_WIDTH + kvHead * HEAD_DIM + local]));699 }700 workgroupBarrier();701 attnCore(local, kvHead, position, queryIndex);702}703`,bindings:n,workgroupSize:t}}function eo(e,r){return[e.kvHeads,r,1]}var Te=64,Ur=68;function ro(e,r,t){return t*r*e.heads*Ur*4}function Qn(e,r,t,o){let n=Yt("attnCoreFlashKernel",e,r);if(r!==e.headDim)throw new Error(`attnCoreFlashKernel: the split-position core is one thread per head dim, so workgroup ${r} must be ${e.headDim}`);if(Te!==r)throw new Error(`attnCoreFlashKernel: a ${Te}-position tile needs ${Te} threads, the workgroup has ${r}`);if(!Number.isInteger(t)||t<1)throw new Error(`attnCoreFlashKernel: splits must be a positive integer, got ${t}`);if(!Number.isInteger(o)||o<1)throw new Error(`attnCoreFlashKernel: maxTotal must be a positive integer, got ${o}`);if(o>e.maxPositions)throw new Error(`attnCoreFlashKernel: maxTotal ${o} is past the ${e.maxPositions}-position cache`);return n}function Xn(e,r){return Math.ceil(Math.ceil(r/e)/Te)}function to(e,r,t,o,n){let s=Qn(e,t,o,n),a=[{binding:j.blob,name:"blob",access:"read"},{binding:j.state,name:"state",access:"read"},{binding:j.keyCache,name:"key_cache",access:"read_write"},{binding:j.valueCache,name:"value_cache",access:"read_write"},{binding:j.projection,name:"projection",access:"read"},{binding:j.flash,name:"flash",access:"read_write"}];return{code:`${he}${be}704const HIDDEN: u32 = ${e.hidden}u;705const HEADS: u32 = ${e.heads}u;706const WG: u32 = ${t}u;707const EPSILON: f32 = ${e.epsilon};708const PROJ_ROWS: u32 = ${e.hidden+2*e.kvHeads*e.headDim}u;709const SPLITS: u32 = ${o}u;710const TILE: u32 = ${Te}u;711const TILE_ITERS: u32 = ${Xn(o,n)}u;712const RECORD: u32 = ${Ur}u;713const NEG_INF: f32 = -3.0e38;714${Xt(e,r,t)}715${_e(a,{})}716 717var<workgroup> ktile: array<f16, ${Te*e.headDim}>;718var<workgroup> vtile: array<f16, ${Te*e.headDim}>;719var<workgroup> qk: array<f32, ${(s+1)*e.headDim}>;720var<workgroup> vcur: array<f32, ${e.headDim}>;721var<workgroup> reduce: array<f32, ${t}>;722var<workgroup> sc: array<f32, ${s*Te}>;723${Jt()}724@compute @workgroup_size(WG)725fn main(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) local: u32) {726 let kvHead = wid.x / SPLITS;727 let s = wid.x % SPLITS;728 let queryIndex = wid.y;729 let position = state.position + queryIndex;730 let total = position + 1u;731 let blockLen = (total + SPLITS - 1u) / SPLITS;732 let start = s * blockLen;733 let stop = min(start + blockLen, total);734 // The block that holds this token's own position. Under the registered S rule it is the last735 // one, but a short token under a large S leaves trailing blocks empty and it is not, so it is736 // computed rather than assumed.737 let owner = (total - 1u) / blockLen;738 739 // Every block needs q. Only the block that owns the current position needs this token's k and v,740 // and reading them in every block would multiply the projection scratch by S. The rows the other741 // blocks skip are written to zero rather than left to the implementation: they still take the742 // norm and the rotation, because dropping them would put a barrier under a non-uniform branch,743 // and nothing reads the result.744 let base = queryIndex * PROJ_ROWS;745 for (var g = 0u; g < GROUP; g = g + 1u) {746 qk[g * HEAD_DIM + local] = projection[base + (kvHead * GROUP + g) * HEAD_DIM + local];747 }748 if (s == owner) {749 qk[K_ROW + local] = projection[base + HIDDEN + kvHead * HEAD_DIM + local];750 vcur[local] = f32(f16(projection[base + HIDDEN + KV_WIDTH + kvHead * HEAD_DIM + local]));751 } else {752 qk[K_ROW + local] = 0.0;753 vcur[local] = 0.0;754 }755 workgroupBarrier();756 headRows(local, position);757 if (s == owner) {758 key_cache[(kvHead * MAX_POSITIONS + position) * HEAD_DIM + local] = f16(qk[K_ROW + local]);759 value_cache[(kvHead * MAX_POSITIONS + position) * HEAD_DIM + local] = f16(vcur[local]);760 }761 762 // The running state of the online softmax. Every one of these is computed from workgroup-uniform763 // data and is identical in every lane, so it lives in registers and costs no workgroup memory --764 // except vacc, which is per dim, and thread local owns dim local in every tile.765 var m: array<f32, ${s}>;766 var l: array<f32, ${s}>;767 var vacc: array<f32, ${s}>;768 var corr: array<f32, ${s}>;769 var peak: array<f32, ${s}>;770 for (var g = 0u; g < GROUP; g = g + 1u) { m[g] = NEG_INF; l[g] = 0.0; vacc[g] = 0.0; }771 772 for (var t = 0u; t < TILE_ITERS; t = t + 1u) {773 let tileStart = start + t * TILE;774 // TILE_ITERS is baked from the deepest total the run reaches, because one encoded command775 // buffer serves every token of a run that sweeps depth. At shallower positions the trailing776 // tiles are empty, and an empty tile must cost a branch rather than two full workgroup777 // reductions per q head -- otherwise the core's price would be set by the end of the run778 // instead of by the position it is at. The test is workgroup-uniform: start and stop come from779 // workgroup_id and a read-only storage load, and t is a uniform loop counter, so every barrier780 // below stays in uniform control flow.781 if (tileStart >= stop) { continue; }782 let len = min(TILE, stop - tileStart);783 784 // No barrier in this loop, so a runtime bound is legal.785 for (var p = 0u; p < len; p = p + 1u) {786 let row = (kvHead * MAX_POSITIONS + tileStart + p) * HEAD_DIM + local;787 ktile[p * HEAD_DIM + local] = key_cache[row];788 vtile[p * HEAD_DIM + local] = value_cache[row];789 }790 workgroupBarrier();791 792 // Thread local scores position local of the tile against all GROUP q heads in one pass793 // over its staged k row, in ascending d -- the same term order as the core this replaces.794 var sv: array<f32, ${s}>;795 for (var g = 0u; g < GROUP; g = g + 1u) { sv[g] = NEG_INF; }796 if (local < len) {797 var a: array<f32, ${s}>;798 for (var g = 0u; g < GROUP; g = g + 1u) { a[g] = 0.0; }799 for (var d = 0u; d < HEAD_DIM; d = d + 1u) {800 let kd = f32(ktile[local * HEAD_DIM + d]);801 for (var g = 0u; g < GROUP; g = g + 1u) { a[g] = a[g] + qk[g * HEAD_DIM + d] * kd; }802 }803 for (var g = 0u; g < GROUP; g = g + 1u) { sv[g] = a[g] * SCALE; }804 }805 806 for (var g = 0u; g < GROUP; g = g + 1u) {807 reduce[local] = sv[g];808 workgroupBarrier();809 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {810 if (local < stride) { reduce[local] = max(reduce[local], reduce[local + stride]); }811 workgroupBarrier();812 }813 peak[g] = reduce[0];814 workgroupBarrier();815 }816 817 // The online update. An empty tile leaves peak at NEG_INF, so corr is exp(0) = 1 and the818 // denominator gains nothing; an empty BLOCK never leaves NEG_INF, and the merge prices its819 // weight at exp(NEG_INF - M) = 0 against any real maximum.820 for (var g = 0u; g < GROUP; g = g + 1u) {821 let next = max(m[g], peak[g]);822 corr[g] = exp(m[g] - next);823 var e: f32 = 0.0;824 if (local < len) { e = exp(sv[g] - next); }825 sc[g * TILE + local] = e;826 reduce[local] = e;827 workgroupBarrier();828 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {829 if (local < stride) { reduce[local] = reduce[local] + reduce[local + stride]; }830 workgroupBarrier();831 }832 l[g] = l[g] * corr[g] + reduce[0];833 m[g] = next;834 workgroupBarrier();835 }836 837 // Thread local owns dim local and walks the tile's positions in index order, so the sum is838 // fixed and the staged V row it reads is contiguous across the workgroup.839 var w: array<f32, ${s}>;840 for (var g = 0u; g < GROUP; g = g + 1u) { w[g] = 0.0; }841 for (var p = 0u; p < len; p = p + 1u) {842 let vd = f32(vtile[p * HEAD_DIM + local]);843 for (var g = 0u; g < GROUP; g = g + 1u) { w[g] = w[g] + sc[g * TILE + p] * vd; }844 }845 for (var g = 0u; g < GROUP; g = g + 1u) { vacc[g] = vacc[g] * corr[g] + w[g]; }846 workgroupBarrier();847 }848 849 for (var g = 0u; g < GROUP; g = g + 1u) {850 let rec = ((queryIndex * SPLITS + s) * HEADS + kvHead * GROUP + g) * RECORD;851 flash[rec + local] = vacc[g];852 if (local == 0u) {853 flash[rec + HEAD_DIM] = m[g];854 flash[rec + HEAD_DIM + 1u] = l[g];855 }856 }857}858`,bindings:a,workgroupSize:t}}function oo(e,r,t){return[e.kvHeads*r,t,1]}function no(e,r,t){if(r!==e.headDim)throw new Error(`attnMergeKernel: one thread per head dim, so workgroup ${r} must be ${e.headDim}`);if(!Number.isInteger(t)||t<1)throw new Error(`attnMergeKernel: splits must be a positive integer, got ${t}`);let o=[{binding:j.destination,name:"destination",access:"read_write"},{binding:j.flash,name:"flash",access:"read"}];return{code:`${he}859const HEADS: u32 = ${e.heads}u;860const HEAD_DIM: u32 = ${e.headDim}u;861const WG: u32 = ${r}u;862const SPLITS: u32 = ${t}u;863const RECORD: u32 = ${Ur}u;864const NEG_INF: f32 = -3.0e38;865 866${_e(o,{})}867 868@compute @workgroup_size(WG)869fn main(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) local: u32) {870 let head = wid.x;871 let queryIndex = wid.y;872 let first = (queryIndex * SPLITS * HEADS + head) * RECORD;873 var peak: f32 = NEG_INF;874 for (var s = 0u; s < SPLITS; s = s + 1u) {875 peak = max(peak, flash[first + s * HEADS * RECORD + HEAD_DIM]);876 }877 var num: f32 = 0.0;878 var den: f32 = 0.0;879 for (var s = 0u; s < SPLITS; s = s + 1u) {880 let rec = first + s * HEADS * RECORD;881 let weight = exp(flash[rec + HEAD_DIM] - peak);882 num = num + weight * flash[rec + local];883 den = den + weight * flash[rec + HEAD_DIM + 1u];884 }885 destination[(queryIndex * HEADS + head) * HEAD_DIM + local] = f16(num / den);886}887`,bindings:o,workgroupSize:r}}function so(e,r){return[e.heads,r,1]}function ao(e,r,t){Wn("mlpFusedKernel",e,t,e.ffn);let{workgroupSize:o,subgroupSize:n,slices:s}=t,a=e.ffn/s;if(a%e.quantBlock!==0)throw new Error(`mlpFusedKernel: ${a} rows per slice is not a whole number of ${e.quantBlock}-wide blocks`);let i=Kn(t);return{code:`${he}${be}${Tr}888${Hn(e,t)}const FFN: u32 = ${e.ffn}u;889const ROWS_PER_SLICE: u32 = ${a}u;890const SLICE_ITERS: u32 = ${a/(o/n)}u;891 892${_e(i,{})}893 894var<workgroup> stage: array<f32, ${e.hidden}>;895var<workgroup> reduce: array<f32, ${o}>;896var<workgroup> hidden_slice: array<f32, ${a}>;897${jn(e,r.norm,t.foldSlices)}${Un(r.gateUpQuant,r.gateUpScales,e.hidden/e.quantBlock,n)}${Fn(r.downKQuant,r.downKScales,a/e.quantBlock,"hidden_slice")}898@compute @workgroup_size(WG)899fn main(900 @builtin(workgroup_id) wid: vec3<u32>,901 @builtin(local_invocation_index) local: u32,902 @builtin(subgroup_invocation_id) lane: u32,903) {904 let group = wid.x;905 stageBlock(local, group);906 let sub = local / SG;907 let rowBase = group * ROWS_PER_SLICE;908 for (var i = 0u; i < SLICE_ITERS; i = i + 1u) {909 let r = sub + i * ROWS_AT_ONCE;910 let gate = rowSum(rowBase + r, lane);911 let up = rowSum(FFN + rowBase + r, lane);912 if (lane == 0u) { hidden_slice[r] = f32(f16((gate / (1.0 + exp(-gate))) * up)); }913 }914 workgroupBarrier();915 project(local, group);916}917`,bindings:i,workgroupSize:o}}function io(e,r,t,o){if(r%t!==0)throw new Error(`foldKernel: workgroup ${r} is not a multiple of ${t} threads per row`);if(o<1)throw new Error(`foldKernel: nothing to fold, got ${o} slices`);let n=r/t,s=[{binding:j.residualOut,name:"residual_out",access:"read_write"},{binding:j.partialsIn,name:"partials_in",access:"read"}];return{code:`${he}918const HIDDEN: u32 = ${e.hidden}u;919const WG: u32 = ${r}u;920const LANES: u32 = ${t}u;921const ROWS_PER_WG: u32 = ${n}u;922const FOLD_SLICES: u32 = ${o}u;923 924${_e(s,{})}925 926var<workgroup> part: array<f32, ${r}>;927 928@compute @workgroup_size(WG)929fn main(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_index) local: u32) {930 let slot = local % ROWS_PER_WG;931 let lane = local / ROWS_PER_WG;932 let row = wid.x * ROWS_PER_WG + slot;933 var acc: f32 = 0.0;934 if (row < HIDDEN) {935 for (var s = lane; s < FOLD_SLICES; s = s + LANES) { acc = acc + partials_in[s * HIDDEN + row]; }936 }937 part[local] = acc;938 workgroupBarrier();939 if (lane == 0u && row < HIDDEN) {940 var sum: f32 = residual_out[row];941 for (var l = 0u; l < LANES; l = l + 1u) { sum = sum + part[l * ROWS_PER_WG + slot]; }942 residual_out[row] = sum;943 }944}945`,bindings:s,workgroupSize:r}}function uo(e){return[e,1,1]}function lo(e,r,t){return[Math.ceil(e.hidden/(r/t)),1,1]}var yr=0,sr="carried";function co(e){if(![0,2,4,8].includes(e))throw new Error("Unregistered padding");yr=e}function po(e){if(!["carried","subgroup","vec4","both"].includes(e))throw new Error("Unregistered flash variant");sr=e}function J(e,r,t){if(e.split(r).length!==2)throw new Error("Flash source seam drift: "+r.slice(0,70));return e.replace(r,t)}function fo(...e){let r=to(...e),t=r.code,o=e[0],n=o.headDim+yr;if(yr&&(t=J(t,`var<workgroup> ktile: array<f16, ${64*o.headDim}>;`,`var<workgroup> ktile: array<f16, ${64*n}>;`),t=J(t,"ktile[p * HEAD_DIM + local]",`ktile[p * ${n}u + local]`),t=J(t,"ktile[local * HEAD_DIM + d]",`ktile[local * ${n}u + d]`)),(sr==="subgroup"||sr==="both")&&(t=J(t,` reduce[local] = 0.0;946 workgroupBarrier();947 if (local < HEAD_DIM) { reduce[local] = qk[off + local] * qk[off + local]; }948 workgroupBarrier();949 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {950 if (local < stride) { reduce[local] = reduce[local] + reduce[local + stride]; }951 workgroupBarrier();952 }953 let inv = inverseSqrt(reduce[0] / f32(HEAD_DIM) + EPSILON);`,` let square = qk[off + local] * qk[off + local];954 let subgroupSquare = subgroupAdd(square);955 if (local % 32u == 0u) { reduce[local / 32u] = subgroupSquare; }956 workgroupBarrier();957 let inv = inverseSqrt((reduce[0] + reduce[1]) / f32(HEAD_DIM) + EPSILON);`),t=J(t,` reduce[local] = sv[g];958 workgroupBarrier();959 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {960 if (local < stride) { reduce[local] = max(reduce[local], reduce[local + stride]); }961 workgroupBarrier();962 }963 peak[g] = reduce[0];`,` let subgroupPeak = subgroupMax(sv[g]);964 if (local % 32u == 0u) { reduce[local / 32u] = subgroupPeak; }965 workgroupBarrier();966 peak[g] = max(reduce[0], reduce[1]);`),t=J(t,` reduce[local] = e;967 workgroupBarrier();968 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {969 if (local < stride) { reduce[local] = reduce[local] + reduce[local + stride]; }970 workgroupBarrier();971 }972 l[g] = l[g] * corr[g] + reduce[0];`,` let subgroupDenom = subgroupAdd(e);973 if (local % 32u == 0u) { reduce[local / 32u] = subgroupDenom; }974 workgroupBarrier();975 l[g] = l[g] * corr[g] + (reduce[0] + reduce[1]);`)),sr==="vec4"||sr==="both"){let s=o.heads/o.kvHeads,a=yr?`${n}u`:"HEAD_DIM";t=J(t,` var a: array<f32, ${s}>;976 for (var g = 0u; g < GROUP; g = g + 1u) { a[g] = 0.0; }977 for (var d = 0u; d < HEAD_DIM; d = d + 1u) {978 let kd = f32(ktile[local * ${a} + d]);979 for (var g = 0u; g < GROUP; g = g + 1u) { a[g] = a[g] + qk[g * HEAD_DIM + d] * kd; }980 }981 for (var g = 0u; g < GROUP; g = g + 1u) { sv[g] = a[g] * SCALE; }`,` var a: array<vec4<f32>, ${s}>;982 for (var g = 0u; g < GROUP; g = g + 1u) { a[g] = vec4<f32>(0.0); }983 for (var d = 0u; d < HEAD_DIM; d = d + 4u) {984 let b = local * ${a} + d;985 let kd = vec4<f32>(f32(ktile[b]), f32(ktile[b+1u]), f32(ktile[b+2u]), f32(ktile[b+3u]));986 for (var g = 0u; g < GROUP; g = g + 1u) {987 let q = g * HEAD_DIM + d;988 a[g] = a[g] + vec4<f32>(qk[q],qk[q+1u],qk[q+2u],qk[q+3u]) * kd;989 }990 }991 for (var g = 0u; g < GROUP; g = g + 1u) { sv[g] = ((a[g].x+a[g].y)+(a[g].z+a[g].w)) * SCALE; }`)}return{...r,code:t}}var ar="carried";function mo(e){if(!["carried","subgroup","packed","both"].includes(e))throw new Error("Unknown MLP mode");ar=e}function go(...e){let r=e[2].slices===160,t=ao(e[0],e[1],r?{...e[2],slices:80}:e[2]),o=t.code;if(r){let n=e[2].workgroupSize;o=J(o,"const SLICES: u32 = 80u;","const SLICES: u32 = 160u;"),o=J(o,"const ROWS_PER_SLICE: u32 = 32u;","const ROWS_PER_SLICE: u32 = 16u;"),o=J(o,`const SLICE_ITERS: u32 = ${32/(n/32)}u;`,`const SLICE_ITERS: u32 = ${16/(n/32)}u;`),o=J(o,"var<workgroup> hidden_slice: array<f32, 32>;","var<workgroup> hidden_slice: array<f32, 16>;"),o=J(o,"KQUANT + (sliceBase + row) * WORDS_PER_ROW_SLICE + b * 4u","KQUANT + ((group / 2u) * HIDDEN + row) * 4u + (group % 2u) * 2u"),o=J(o,"KSCALES + (sliceBase + row) * BLOCKS_PER_SLICE + b","KSCALES + (group / 2u) * HIDDEN + row");let s=o.indexOf("fn project("),a=o.indexOf("@compute",s),i=o.slice(s,a),u=0;if(i=i.replace(/ part = part \+ dot8\(blob\[base \+ [23]u\], 8\.0,[\s\S]*?\);/g,()=>(u++,"")),u!==2)throw new Error("Half-block source seam drift");o=o.slice(0,s)+i+o.slice(a)}if(!r&&e[2].slices!==80){if(![20,40].includes(e[2].slices))throw new Error("Unregistered MLP slices");o=J(o,"KQUANT + (sliceBase + row) * WORDS_PER_ROW_SLICE + b * 4u","KQUANT + ((group * BLOCKS_PER_SLICE + b) * HIDDEN + row) * 4u"),o=J(o,"KSCALES + (sliceBase + row) * BLOCKS_PER_SLICE + b","KSCALES + (group * BLOCKS_PER_SLICE + b) * HIDDEN + row")}if((ar==="subgroup"||ar==="both")&&(o=J(o," reduce[local] = sum;",` let sgSum = subgroupAdd(sum);992 if (local % SG == 0u) { reduce[local / SG] = sgSum; }`),o=J(o,` for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {993 if (local < stride) { reduce[local] = reduce[local] + reduce[local + stride]; }994 workgroupBarrier();995 }996 let inv = inverseSqrt(reduce[0] / f32(HIDDEN) + EPSILON);`,` var squareSum = 0.0;997 for (var subgroup = 0u; subgroup < WG / SG; subgroup++) { squareSum += reduce[subgroup]; }998 let inv = inverseSqrt(squareSum / f32(HIDDEN) + EPSILON);`)),ar==="packed"||ar==="both"){let n=0;if(o=o.replace(/part = part \+ dot8\((blob\[[^\]]+\]), 8\.0,\s*(vec4<f32>\([^;]+?)\);/g,(s,a,i)=>{n++;let u=i.replaceAll("vec4<f32>","vec4<f16>").replace(/(stage|hidden_slice)\[[^\]]+\]/g,c=>`f16(${c})`);return`part = part + f32(dot8h(${a}, f16(8.0), ${u}));`}),n!==8)throw new Error("MLP dot source seam drift "+n);o+=`999`+nr}return{...t,code:o}}function ho(e,r){let t=$t(...e);if(!r||e[0]==="matvec_residual")return t;let o=t.code,n=` reduce[local] = sum;1000 workgroupBarrier();1001 for (var stride = WG / 2u; stride > 0u; stride = stride / 2u) {1002 if (local < stride) { reduce[local] = reduce[local] + reduce[local + stride]; }1003 workgroupBarrier();1004 }1005 let inv = inverseSqrt(reduce[0] / f32(COLS) + EPSILON);`;if(o.split(n).length!==2)throw new Error("Matvec RMS source seam drift");return o=o.replace(n,` let sgSum = subgroupAdd(sum);1006 if (local % SG == 0u) { reduce[local / SG] = sgSum; }1007 workgroupBarrier();1008 var squareSum = 0.0;1009 for (var subgroup = 0u; subgroup < WG / SG; subgroup++) { squareSum += reduce[subgroup]; }1010 let inv = inverseSqrt(squareSum / f32(COLS) + EPSILON);`),{...t,code:o}}var bo=8;function ne(e,r){if(r.offset%256!==0)throw new Error(`runtime: section offset ${r.offset} is not 256-aligned`);return{buffer:e,offset:r.offset,size:r.bytes}}function Yn(e,r,t){return e.createBindGroupLayout({label:t,entries:r.bindings.map(o=>({binding:o.binding,visibility:GPUShaderStage.COMPUTE,buffer:{type:o.access==="read"?"read-only-storage":"storage"}}))})}function _o(e,r,t){let o=Yn(e,r,`${t}.layout`),n=e.createShaderModule({code:r.code,label:t});return{pipeline:e.createComputePipeline({label:t,layout:e.createPipelineLayout({bindGroupLayouts:[o]}),compute:{module:n,entryPoint:"main"}}),layout:o}}function Jn(e,r){return{hidden:e.hidden,ffn:e.ffn,vocab:e.vocab,heads:e.heads,kvHeads:e.kvHeads,headDim:e.headDim,convTaps:e.convCache,quantBlock:e.quantBlock,epsilon:Ze.normEps,maxPositions:r.maxPositions,maxQueries:r.maxQueries}}var Fr=class{constructor(r){this.device=r}entries=new Map;get(r,t){let o=this.entries.get(r);if(o!==void 0)return o;let n=t(),{pipeline:s,layout:a}=_o(this.device,n,r),i={pipeline:s,layout:a,kernel:n};return this.entries.set(r,i),i}get size(){return this.entries.size}};function wo(e,r,t,o,n,s,a){if(s.length===0)throw new Error("runtime: no arms to build");let i=Jn(r,a),{maxQueries:u,maxPositions:c}=a,f=r.layerTypes.filter(l=>l==="attention").length,k=r.layerTypes.length-f,y=l=>{let d=n.get(l);if(d===void 0)throw new Error(`runtime: weight section ${l} missing`);return d},P=l=>{let d=y(l);if(d.offset%4!==0)throw new Error(`runtime: section ${l} at ${d.offset} is not word aligned`);return d.offset/4},B=l=>{let d=y(`${l}.norm`),G=y(`${l}.q_norm`),z=y(`${l}.k_norm`);if(G.offset!==d.offset+d.bytes||z.offset!==G.offset+G.bytes)throw new Error(`runtime: ${l} norm/q_norm/k_norm are not contiguous`);return{offset:d.offset,bytes:d.bytes+G.bytes+z.bytes}},S=GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_DST|GPUBufferUsage.COPY_SRC,v=(l,d)=>e.createBuffer({size:l,usage:S,label:d}),$=u*r.hidden*Y.residualBytes,x=a.maxSlices*r.hidden*Y.residualBytes,X=u*3*r.hidden*Y.hiddenBytes,Z=u*Math.max(r.hidden,r.ffn)*Y.hiddenBytes,I=u*(r.hidden+2*r.kvHeads*r.headDim)*4,q=r.kvHeads*c*r.headDim*Y.kvCacheBytes,ee=r.hidden*r.convCache*Y.convCacheBytes,pe=u*r.vocab*Y.logitsBytes,Ae=s.reduce((l,d)=>Math.max(l,d.options.flash.blocks),1),Pe=ro(i,Ae,u),E=a.samplePartials;if(q%256!==0||ee%256!==0)throw new Error("runtime: per-layer cache stride is not 256-aligned");let ke=v($,"residual_a"),Q=v($,"residual_b"),K=v($,"normalized"),re=v(x,"deferred_a"),ue=v(x,"deferred_b"),ze=v(X,"proj"),ie=v(Z,"mid"),le=v(I,"qkv_proj"),de=v(f*q,"key_cache"),Ee=v(f*q,"value_cache"),fe=v(k*ee,"conv_cache"),_=v(pe,"logits"),p=v(bo*4,"state"),g=v(c*4,"tokens"),w=v(c*4,"sampled"),R=v(E*2*4,"partials"),L=v(Pe,"flash"),h=l=>({buffer:l,offset:0,size:l.size}),A=(l,d,G)=>({buffer:l,offset:d*G,size:G}),N=(l,d,G,z)=>{if(l===0)return d;if(l===1)return G;throw new Error(`runtime: ${z} parity ${l} is not 0 or 1`)},M=l=>h(N(l,ke,Q,"residual")),te=l=>h(N(l,re,ue,"partials")),me=2*r.ffn*r.hidden/2,Oe=r.ffn*r.hidden/2,pr=l=>{if(!s.some(F=>{let se=tr(F.key).mlpSoA;return se===l||se==="both"}))return null;let d=l==="gate"?me:Oe,G=v(d*r.layerTypes.length,"mlp_"+l+"_soa"),z=Et(r.hidden,r.ffn,l),D=_o(e,z,"repack_"+l),U=e.createCommandEncoder(),W=U.beginComputePass();W.setPipeline(D.pipeline);for(let F=0;F<r.layerTypes.length;F++){let se=y(`L${F}.mlp.${l==="gate"?"gate_up.quant":"down.kquant"}`);if(se.bytes!==d)throw new Error("SoA source size mismatch");let Ue=e.createBindGroup({layout:D.layout,entries:[{binding:0,resource:ne(t,se)},{binding:1,resource:A(G,F,d)}]});W.setBindGroup(0,Ue),W.dispatchWorkgroups(Math.ceil(d/4/256))}return W.end(),e.queue.submit([U.finish()]),G},dr=pr("gate"),fr=pr("down"),ce=new Fr(e),Ho=(l,d,G,z)=>{let D=d.map(U=>{let W=G[U.name];if(W===void 0)throw new Error(`runtime: ${z} has no resource for ${U.name}`);return{binding:U.binding,resource:W}});return e.createBindGroup({layout:l,entries:D,label:z})},Ko=l=>({norm:P(`L${l}.mlp.gate_up.norm`),gateUpQuant:P(`L${l}.mlp.gate_up.quant`),gateUpScales:P(`L${l}.mlp.gate_up.scales`),downKQuant:P(`L${l}.mlp.down.kquant`),downKScales:P(`L${l}.mlp.down.kscales`)}),Mr=l=>({qNorm:P(`L${l}.attn.qkv.q_norm`),kNorm:P(`L${l}.attn.qkv.k_norm`),cos:P("rope.cos_ext"),sin:P("rope.sin_ext")}),Vo=(l,d)=>{if(d.fusion.slices!==l.slices)throw new Error(`runtime: ${d.name} publishes ${d.fusion.slices} slices, the kernel is built for ${l.slices}`);if(l.slices>a.maxSlices)throw new Error(`runtime: ${d.name} publishes ${l.slices} slices into a ${a.maxSlices}-row partial buffer`);return{workgroupSize:l.workgroupSize,subgroupSize:l.subgroupSize,slices:l.slices,foldSlices:d.fusion.foldSlices}},Ve=l=>{if(l.weight===null)throw new Error(`runtime: ${l.name} is a matmul with no weight`);return l.weight},Qo=l=>{let d=tr(l.key);co(d.padding),po(d.flashMode),mo(d.mlpMode??"carried");let G=Cr(br(r,l.options),d,r,l.options.position,l.options.queries),z=l.options.queries,D=l.options.flash.blocks,U=a.flashMaxTotalPositions,W=[],F=0,se=0,Ue=!1,ge=(m,O,b,V)=>{if(O.kernel.workgroupSize!==m.workgroupSize)throw new Error(`runtime: ${m.name} is planned at workgroup size ${m.workgroupSize}, its kernel is ${O.kernel.workgroupSize}`);let[oe,xe,Ce]=V(z);if(oe*xe*Ce!==m.workgroups)throw new Error(`runtime: ${m.name} is planned at ${m.workgroups} workgroups, the runtime encodes ${oe*xe*Ce}`);W.push({plan:m,pipeline:O.pipeline,bindGroup:Ho(O.layout,O.kernel.bindings,b,`${l.key}.${m.name}`),workgroups:V})},Qe=(m,O,b,V,oe,xe,Ce,Gr)=>{if(m.site===null||m.geometry===null)throw new Error(`runtime: ${m.name} is a matvec, so the plan must name the site and the geometry it runs at`);let $r=Rt(m.geometry,a.matvecMaxCols),{workgroupSize:Jo,subgroupSize:Zo,rowsPerSubgroup:en,subgroupsPerRow:rn}=$r,tn=`w${Jo}s${Zo}r${en}c${rn}`,tt=d.matvec[m.site]?.variant??"carried",ot=m.kind==="conv_proj_core",Lr=Dr(d,m),on=`${O}_${b}x${V}${oe?"_zp":""}_${tn}_${tt}_n${d.normSubgroup??!1}_sw${d.matvecStageSwizzle??!1}_pv${d.matvecPrivate??!1}_once${Lr}_conv${ot}_mask${d.maskMatvec??!1}_magic${d.magicMatvec??!1}`,nn=ce.get(on,()=>qr(Nr(jt(Ft(Ut(Wr(ho([O,b,V,oe,$r,i,xe,tt],d.normSubgroup??!1),d.matvecStageSwizzle??!1),O,xe,d.matvecPrivate??!1),Lr),ot),d.maskMatvec??!1),d.magicMatvec??!1)),Ar={weights:ne(t,y(`${Ce}.quant`)),scales:ne(t,y(`${Ce}.scales`)),state:h(p),partials:h(R),...Gr};Lr&&(Ar.residual=h(K)),oe&&(Ar.zero_points=ne(t,y(`${Ce}.zero_points`))),ge(m,nn,Ar,sn=>Gt(O,b,i,$r,sn))},Yo=(m,O,b,V,oe)=>{let xe=Vo(b,m),Ce=ce.get(O,()=>V(xe)),Gr={blob:h(t),weights:h(t),projection:h(K),residual_in:M(m.fusion.residualIn),residual_out:M(m.fusion.residualOut),partials_in:te(m.fusion.partialsIn),partials_out:te(m.fusion.partialsOut),state:h(p),...oe};ge(m,Ce,Gr,()=>uo(b.slices)),Ue=!0};for(let m of G){let O=m.fusion;switch(m.kind){case"input_norm":{let b=m.name.replace(/\.input_norm$/,""),V=b.endsWith(".mlp.fused")?b.replace(/\.fused$/,".gate_up.norm"):b+".norm",oe=b==="head"?a.headPreScale:1,xe=ce.get("input_norm_"+m.workgroupSize+"_"+oe,()=>Ht(m.workgroupSize,oe));ge(m,xe,{residual:M(O.residualIn),gamma:ne(t,y(V)),projection:h(K)},()=>[1,1,1]);break}case"embed":{let b=ce.get("embed",()=>Lt(i,a.embedWorkgroup));ge(m,b,{weights:ne(t,y("embed.quant")),scales:ne(t,y("embed.scales")),zero_points:ne(t,y("embed.zero_points")),residual:M(O.residualOut),state:h(p),tokens:h(g)},V=>At(V));break}case"conv_proj_core":{let b=Ve(m);Qe(m,"norm_matvec",b.rows,b.cols,b.zeroPoints,1,m.name,{gamma:ne(t,y(m.name+".norm")),residual:M(O.residualIn),destination:h(ie),blob:ne(t,y(m.name.replace(/\.in_proj$/,".core.taps"))),cache:A(fe,se,ee)}),se++;break}case"norm_matvec":{let b=Ve(m);Qe(m,"norm_matvec",b.rows,b.cols,b.zeroPoints,1,m.name,{gamma:ne(t,y(`${m.name}.norm`)),residual:M(O.residualIn),destination:h(ze)});break}case"conv_core":{let b=ce.get("conv_core",()=>zt(i,a.convWorkgroup));ge(m,b,{gamma:ne(t,y(`${m.name}.taps`)),source:h(ze),destination:h(ie),state:h(p),cache:A(fe,se,ee)},()=>Ot(i,a.convWorkgroup)),se+=1;break}case"matvec_residual":{let b=Ve(m);Qe(m,"matvec_residual",b.rows,b.cols,b.zeroPoints,1,m.name,{residual:M(O.residualOut),source:h(ie)});break}case"norm_matvec_swiglu":{let b=Ve(m);Qe(m,"norm_swiglu",b.rows,b.cols,b.zeroPoints,1,m.name,{gamma:ne(t,y(`${m.name}.norm`)),residual:M(O.residualIn),destination:h(ie)});break}case"attn_proj":{let b=Ve(m);Qe(m,"norm_projection",b.rows,b.cols,b.zeroPoints,1,m.name,{gamma:ne(t,y(`${m.name}.norm`)),residual:M(O.residualIn),projection:h(le)});break}case"attn_core_qkv":{let b=m.layer,V=ce.get(`attn_core_qkv_L${b}`,()=>Zt(i,Mr(b),a.attnCoreWorkgroup));ge(m,V,{blob:h(t),state:h(p),key_cache:A(de,F,q),value_cache:A(Ee,F,q),projection:h(le),destination:h(ie)},oe=>eo(i,oe)),F+=1;break}case"attn_core_flash":{let b=m.layer,V=ce.get(`attn_core_flash_L${b}_s${D}_t${U}_p${d.padding}_${d.flashMode}_online${d.flashOnline??!1}`,()=>d.flashOnline?Nt(i,Mr(b),D):fo(i,Mr(b),a.flashCoreWorkgroup,D,U));ge(m,V,{blob:h(t),state:h(p),key_cache:A(de,F,q),value_cache:A(Ee,F,q),projection:h(le),flash:h(L)},oe=>oo(i,D,oe)),F+=1;break}case"attn_merge":{let b=ce.get(`attn_merge_s${D}_sg${d.mergeSubgroup??!1}`,()=>qt(no(i,a.mergeWorkgroup,D),D,d.mergeSubgroup??!1));ge(m,b,{flash:h(L),destination:h(ie)},V=>so(i,V));break}case"mlp_fused":{let b=m.layer;Yo(m,`mlp_fused_${d.key}_L${b}_f${O.foldSlices}_${d.mlpMode??"carried"}_s${d.mlpSlices??80}_w${d.mlpWorkgroup??256}`,{...a.mlpFused,slices:d.mlpSlices??a.mlpFused.slices,workgroupSize:d.mlpWorkgroup??a.mlpFused.workgroupSize},V=>Kt(go(i,Ko(b),V),d),{...(d.mlpSoA==="gate"||d.mlpSoA==="both")&&dr?{weights:A(dr,b,me)}:{},...(d.mlpSoA==="down"||d.mlpSoA==="both")&&fr?{source:A(fr,b,Oe)}:{}});break}case"fold":{if(O.residualIn!==O.residualOut||O.partialsIn!==O.partialsOut)throw new Error(`runtime: ${m.name} is in place, so its parities must match`);let b=ce.get(`fold_${d.key}_${O.foldSlices}`,()=>Vt(io(i,a.fold.workgroupSize,d.foldLanes??a.fold.threadsPerRow,O.foldSlices),d));ge(m,b,{residual_out:M(O.residualOut),partials_in:te(O.partialsIn)},()=>lo(i,a.fold.workgroupSize,d.foldLanes??a.fold.threadsPerRow));break}case"norm_head":{let b=Ve(m);Qe(m,"norm_head",b.rows,b.cols,b.zeroPoints,a.headPreScale,"head",{gamma:ne(t,y("head.norm")),residual:M(O.residualIn),logits:h(_)});break}case"sample_partial":{let b=ce.get(`sample_partial_${E}`,()=>Dt(i,a.sampleWorkgroup,E));ge(m,b,{state:h(p),partials:h(R),logits:h(_)},()=>It(E));break}case"sample_final":{let b=ce.get(`sample_final_${E}`,()=>Ct(i,a.sampleWorkgroup,E));ge(m,b,{state:h(p),tokens:h(g),partials:h(R),sampled:h(w)},()=>Tt());break}}}if(W.length!==G.length)throw new Error(`runtime: arm ${l.key} encoded ${W.length} of ${G.length} planned dispatches`);if(Ue&&z!==1)throw new Error(`runtime: arm ${l.key} is fused, which is a single-query decode specialisation, and plans ${z} queries`);return{key:l.key,plan:G,steps:W,fused:Ue}},De=new Map;for(let l of s){if(De.has(l.key))throw new Error(`runtime: arm ${l.key} built twice`);De.set(l.key,Qo(l))}let Zr=De.get(s[0]?.key??"");if(Zr===void 0)throw new Error("runtime: the first arm did not build");let ve=Zr,et=[...dr?[dr]:[],...fr?[fr]:[],K,ke,Q,re,ue,ze,ie,le,de,Ee,fe,_,p,g,w,R,L],Xo=et.reduce((l,d)=>l+d.size,o),rt=0;for(let l of De.values())rt+=l.steps.length;let Be=new Uint32Array(bo),mr=l=>{if(l!==1&&ve.fused)throw new Error(`runtime: arm ${ve.key} is fused and encodes one query, not ${l}`);return ve.steps};return{shape:r,armKeys:[...De.keys()],get arm(){return ve.key},get steps(){return ve.steps},get dispatchCount(){return ve.steps.length},get tokensPerPass(){return tr(ve.key).tokensPerPass??1},pipelineCount:ce.size,bindGroupCount:rt,weightBytes:o,liveGpuBytes:Xo,buffers:{state:p,tokens:g,sampled:w,logits:_,residual:ke,residualB:Q,deferredA:re,deferredB:ue,partials:R,proj:ze,qkvProj:le,flash:L,mid:ie},setArm(l){let d=De.get(l);if(d===void 0)throw new Error(`runtime: no arm ${l}`);ve=d},dispatchCountOf(l){let d=De.get(l);if(d===void 0)throw new Error(`runtime: no arm ${l}`);return d.steps.length},planOf(l){let d=De.get(l);if(d===void 0)throw new Error(`runtime: no arm ${l}`);return d.plan},writeState(l){Be[0]=l.position,Be[1]=l.step,Be[2]=l.queries,Be[3]=l.sampleMode,Be[4]=l.rng,Be[5]=l.freerun,Be[6]=l.advance,Be[7]=0,e.queue.writeBuffer(p,0,Be)},writeTokens(l,d){e.queue.writeBuffer(g,d*4,l)},encodeChunk(l,d,G){let z=l.beginComputePass(G?{timestampWrites:G}:{});for(let D=0;D<d;D++)for(let U of mr(1)){z.setPipeline(U.pipeline),z.setBindGroup(0,U.bindGroup);let[W,F,se]=U.workgroups(1);z.dispatchWorkgroups(W,F,se)}z.end()},encodeToken(l,d,G){let z=G===null?l.beginComputePass():l.beginComputePass({timestampWrites:G});for(let D of mr(d)){z.setPipeline(D.pipeline),z.setBindGroup(0,D.bindGroup);let[U,W,F]=D.workgroups(d);z.dispatchWorkgroups(U,W,F)}z.end()},assertReach(l){if(!Number.isInteger(l)||l<1)throw new Error(`runtime: ${l} is not a positive total position count`);if(l>a.flashMaxTotalPositions)throw new Error(`runtime: a run reaching ${l} total positions was built for ${a.flashMaxTotalPositions}, so the split-position core would drop every position past its last tile`);if(l>c)throw new Error(`runtime: a run reaching ${l} total positions is past the ${c}-position cache`)},stepIndex(l){let d=ve.steps.findIndex(G=>G.plan.name===l);if(d<0)throw new Error(`runtime: arm ${ve.key} has no planned dispatch named ${l}`);return d},encodeOne(l,d,G){let z=mr(G)[d];if(z===void 0)throw new Error(`runtime: dispatch index ${d} out of range`);let D=l.beginComputePass();D.setPipeline(z.pipeline),D.setBindGroup(0,z.bindGroup);let[U,W,F]=z.workgroups(G);D.dispatchWorkgroups(U,W,F),D.end()},encodePerDispatchPasses(l,d,G,z){let D=z;for(let U of mr(d)){let W=l.beginComputePass({timestampWrites:{querySet:G,beginningOfPassWriteIndex:D,endOfPassWriteIndex:D+1}});W.setPipeline(U.pipeline),W.setBindGroup(0,U.bindGroup);let[F,se,Ue]=U.workgroups(d);W.dispatchWorkgroups(F,se,Ue),W.end(),D+=2}return D-z},clearCaches(){let l=e.createCommandEncoder();l.clearBuffer(de),l.clearBuffer(Ee),l.clearBuffer(fe),l.clearBuffer(w),e.queue.submit([l.finish()])},clearResidual(){let l=e.createCommandEncoder();l.clearBuffer(ke),l.clearBuffer(Q),l.clearBuffer(re),l.clearBuffer(ue),e.queue.submit([l.finish()])},destroy(){for(let l of et)l.destroy()}}}function yo(e,r,t){let o=rr(e);if(T.matvec.maxCols!==T.matvecWide.maxCols)throw new Error(`registeredRuntimeOptions: the two registered geometries stage ${T.matvec.maxCols} and ${T.matvecWide.maxCols} columns, but the stage width is a kernel bound and must be one number`);return{maxQueries:Xe.maxQueries,maxPositions:T.maxPositions,matvecMaxCols:T.matvec.maxCols,dequantVariant:r,samplePartials:T.sample.partials,sampleWorkgroup:T.sample.workgroupSize,embedWorkgroup:T.embed.workgroupSize,convWorkgroup:T.conv.workgroupSize,attnCoreWorkgroup:C.attnSplit.coreWorkgroupSize,flashCoreWorkgroup:C.flash.coreWorkgroupSize,mergeWorkgroup:C.flash.mergeWorkgroupSize,flashMaxTotalPositions:t,headPreScale:Xe.headPreScale,mlpFused:{workgroupSize:C.mlpFused.workgroupSize,subgroupSize:T.matvec.subgroupSize,slices:o.mlp},fold:C.fold,maxSlices:Or(e)}}function ko(e){let r=new Map;for(let t of e){if(r.has(t.name))throw new Error(`duplicate weight section ${t.name}`);r.set(t.name,{offset:t.offset,bytes:t.bytes})}return r}function Zn(e){return Array.from(new Uint8Array(e)).map(r=>r.toString(16).padStart(2,"0")).join("")}var es=GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_DST|GPUBufferUsage.COPY_SRC;async function vo(e,r,t){let o=performance.now(),n=await fetch(r);if(!n.ok)throw new Error(`weights fetch ${r}: ${n.status}`);let s=await n.arrayBuffer(),a=performance.now()-o;if(s.byteLength!==t)throw new Error(`weights are ${s.byteLength} B, manifest says ${t} B`);let i=Zn(await crypto.subtle.digest("SHA-256",s)),u=e.createBuffer({size:s.byteLength,usage:es,mappedAtCreation:!0,label:"weights"});return new Uint8Array(u.getMappedRange()).set(new Uint8Array(s)),u.unmap(),{buffer:u,bytes:s.byteLength,sha256:i,fetchMs:a}}function rs(e){return{knee:gt(e.knee),flashOptima:e.flashOptima.map(r=>({...r}))}}function ts(e){return{layerTypes:e.shape.layerTypes,hidden:e.shape.hidden,ffn:e.shape.ffn,vocab:e.shape.vocab,heads:e.shape.heads,kvHeads:e.shape.kvHeads,headDim:e.shape.headDim,convCache:e.shape.convCache,quantBits:e.shape.quantBits,quantBlock:e.shape.quantBlock}}function os(e,r,t,o){let n=ts(r),s=ko(r.sections),a=rs(r),i={...yo(n,r.dequantVariant,r.maxTotalPositions),maxSlices:yt()},u=wo(e.device,n,t.buffer,t.bytes,s,wt(n,r.planPosition-r.maxQueries,r.planPosition,a),{...i,maxQueries:r.maxQueries,maxPositions:r.maxPositions,samplePartials:r.samplePartials});return{gpu:e,runtime:u,blob:t,shape:n,sections:s,tuning:a,planPosition:r.planPosition,errors:o,destroy(){u.destroy(),t.buffer.destroy()}}}function ns(e,r,t){return{features:e,limits:{...vt,maxBufferSize:r,maxStorageBufferBindingSize:t}}}function ss(e,r,t,o){let n=Math.max(e*r*Y.logitsBytes,o);for(let s of t)n=Math.max(n,s.bytes);return n}function Hr(e){return ns(e.extraFeatures,e.weightBytes,ss(e.shape.vocab,e.maxQueries,e.sections,e.weightBytes))}async function So(e){let r=[],t=await wr(Hr(e),e.powerPreference,e.expectedSubgroupSize),o=await vo(t.device,e.weightsUrl,e.weightBytes),n=performance.now(),s=null;if(await St(t.device,r,"runtime_build",async()=>{s=os(t,e,o,r),await je(t.device)}),s===null)throw new Error(`runtime build failed: ${r.join("; ")}`);let a=s,i=performance.now()-n,u={env:Pt(),device:t.descriptor,dispatches_per_token:a.runtime.dispatchCount,arm_dispatches:a.runtime.armKeys.map(c=>a.runtime.dispatchCountOf(c)),pipelines:a.runtime.pipelineCount,bind_groups:a.runtime.bindGroupCount,weight_bytes:o.bytes,live_gpu_bytes:a.runtime.liveGpuBytes,build_ms:i,weights_fetch_ms:o.fetchMs,weights_sha256:o.sha256,subgroup_min:t.descriptor.adapter.subgroupMinSize,subgroup_max:t.descriptor.adapter.subgroupMaxSize,validation_errors:[...r]};return{session:a,ready:u}}var jr=class{constructor(r,t,o){this.device=r;this.slotCount=t;this.tokensPerSlot=o;for(let n=0;n<t;n+=1)this.slots.push(r.createBuffer({size:o*4,usage:GPUBufferUsage.MAP_READ|GPUBufferUsage.COPY_DST,label:`readback${n}`})),this.pending.push(null),this.chunkAt.push(-1)}slots=[];pending=[];chunkAt=[];ids=[];slot(r){let t=this.slots[r%this.slotCount];if(t===void 0)throw new Error("readback ring index");return t}async recycle(r){let t=r%this.slotCount,o=this.pending[t];return o!=null&&(await o,this.drainSlot(t)),this.slot(r)}drainSlot(r){let t=this.slots[r],o=this.chunkAt[r];if(t===void 0||o===void 0||o<0)return;let n=new Uint32Array(t.getMappedRange().slice(0));t.unmap();for(let s=0;s<n.length;s+=1){let a=n[s];if(a===void 0)throw new Error("readback view index");this.ids[o*this.tokensPerSlot+s]=a}this.pending[r]=null,this.chunkAt[r]=-1}start(r){let t=r%this.slotCount;this.chunkAt[t]=r,this.pending[t]=this.slot(r).mapAsync(GPUMapMode.READ)}async drain(){for(let r=0;r<this.slotCount;r+=1){let t=this.pending[r];t!=null&&(await t,this.drainSlot(r))}return this.ids}destroy(){for(let r of this.slots)r.destroy()}};async function kr(e,r){let{runtime:t,gpu:o}=e;t.assertReach(r.promptLength),t.clearCaches(),t.clearResidual(),t.writeTokens(r.tokenIds,0),t.writeState({position:0,step:0,queries:1,sampleMode:0,rng:r.seed,freerun:0,advance:1});let n=o.device.createCommandEncoder();for(let s=0;s<r.promptLength;s+=1)t.encodeToken(n,1,null);if(r.freerun){if(r.promptLength<1)throw new Error("Free generation needs a nonempty prompt");n.copyBufferToBuffer(t.buffers.sampled,(r.promptLength-1)*4,t.buffers.tokens,r.promptLength*4,4)}o.device.queue.submit([n.finish()]),await je(o.device)}function Po(e,r){e.runtime.writeState({position:r.startPosition,step:r.startStep,queries:1,sampleMode:0,rng:r.seed,freerun:r.freerun?1:0,advance:1})}function Eo(e,r){if(r.ringTokens%r.readbackEvery!==0)throw new Error(`ring of ${r.ringTokens} tokens does not divide by readback every ${r.readbackEvery}`);return new jr(e,r.ringTokens/r.readbackEvery,r.readbackEvery)}async function as(e,r,t,o){if(r===null)return{perTokenMs:[],gpuBusyMs:-1,gpuSpanMs:-1};let n=await Ir(e,r,t),s=[],a=0;for(let c=0;c<t;c+=1){let f=n[c*2],k=n[c*2+1],y=c+1<t?n[(c+1)*2]:void 0;if(f===void 0||k===void 0)throw new Error("timestamp index");a+=k-f,s.push(((y===void 0?k:y)-f)/1e6)}let i=n[0],u=n[t*2-1];if(i===void 0||u===void 0)throw new Error("timestamp span index");return{perTokenMs:s,gpuBusyMs:a/1e6,gpuSpanMs:(u-i)/1e6}}async function Bo(e,r){if(e.runtime.tokensPerPass>1)return is(e,r);let{runtime:t,gpu:o}=e;t.assertReach(r.startPosition+r.tokens);let n=o.device,s=globalThis.__probeHooks,a=r.timestamps?n.createQuerySet({type:"timestamp",count:r.tokens*2,label:"token_pass"}):null,i=Eo(n,r);Po(e,r),await je(n),s.beginStep();let u=performance.now();for(let P=0;P<r.tokens;P+=1){let B=n.createCommandEncoder();t.encodeToken(B,1,a===null?null:{querySet:a,beginningOfPassWriteIndex:P*2,endOfPassWriteIndex:P*2+1});let S=(P+1)%r.readbackEvery===0,v=-1;if(S){v=(P+1)/r.readbackEvery-1;let $=await i.recycle(v);B.copyBufferToBuffer(t.buffers.sampled,(r.startStep+P+1-r.readbackEvery)*4,$,0,r.readbackEvery*4)}n.queue.submit([B.finish()]),S&&i.start(v)}await je(n);let c=performance.now()-u,f=s.endStep(),k=await i.drain();i.destroy();let y=await as(n,a,r.tokens,c);return a!==null&&a.destroy(),{wallMs:c,sampled:[...k],snapshot:f,...y}}async function is(e,r){let{runtime:t,gpu:o}=e,n=o.device,s=globalThis.__probeHooks;t.assertReach(r.startPosition+r.tokens);let a=Math.min(t.tokensPerPass,r.readbackEvery);if(r.readbackEvery%a)throw new Error("Chunk does not divide readback cadence");let i=Math.ceil(r.tokens/a),u=r.timestamps?n.createQuerySet({type:"timestamp",count:i*2,label:"decode_chunk"}):null,c=Eo(n,r);Po(e,r),await je(n),s.beginStep();let f=performance.now(),k=[];for(let x=0,X=0;x<r.tokens;x+=a,X++){let Z=Math.min(a,r.tokens-x),I=(x+Z)%r.readbackEvery===0,q=n.createCommandEncoder();t.encodeChunk(q,Z,u?{querySet:u,beginningOfPassWriteIndex:X*2,endOfPassWriteIndex:X*2+1}:null),k.push(Z);let ee=-1;if(I){ee=(x+Z)/r.readbackEvery-1;let pe=await c.recycle(ee);q.copyBufferToBuffer(t.buffers.sampled,(r.startStep+x+Z-r.readbackEvery)*4,pe,0,r.readbackEvery*4)}n.queue.submit([q.finish()]),I&&c.start(ee)}await je(n);let y=performance.now()-f,P=s.endStep(),B=await c.drain();c.destroy();let S=[],v=-1,$=-1;if(u){let x=await Ir(n,u,i);v=0;for(let X=0;X<i;X++){let Z=(x[X*2+1]-x[X*2])/1e6;S.push({tokens:k[X],ms:Z}),v+=Z}$=(x[i*2-1]-x[0])/1e6,u.destroy()}return{wallMs:y,snapshot:P,sampled:[...B],perTokenMs:[],gpuBusyMs:v,gpuSpanMs:$,gpuChunks:S}}var al=Ze.normEps;var us=class{constructor(e){this.trie=this._build_trie(e)}_build_trie(e){let r=Object.create(null);for(let t of e){let o=r;for(let n=0;n<t.length;++n){let s=t[n];o=o[s]??=Object.create(null)}o.end=t}return r}split(e){let r=[],t=e.length,o=0,n=0;for(;n<t;){let s=this.trie,a=null,i=n;for(;i<t&&(s=s[e[i]]);)s.end&&(a=s.end),++i;a?(n>o&&r.push(e.slice(o,n)),r.push(a),n+=a.length,o=n):++n}return o<t&&r.push(e.slice(o)),r}},xo=us,ls=class{constructor(e){this.content=e.content,this.id=e.id,this.single_word=e.single_word??!1,this.lstrip=e.lstrip??!1,this.rstrip=e.rstrip??!1,this.special=e.special??!1,this.normalized=e.normalized??!this.special}},cs=ls,zo=(()=>{let e=[...Array.from({length:94},(n,s)=>s+33),...Array.from({length:12},(n,s)=>s+161),...Array.from({length:82},(n,s)=>s+174)],r=e.slice(),t=0;for(let n=0;n<256;++n)e.includes(n)||(e.push(n),r.push(256+t),t+=1);let o=r.map(n=>String.fromCharCode(n));return Object.fromEntries(e.map((n,s)=>[n,o[s]]))})(),ps=e=>Object.fromEntries(Object.entries(e).map(([r,t])=>[t,r])),ds=ps(zo),Ro=".,!?\u2026\u3002\uFF0C\u3001\u0964\u06D4\u060C",fs=new Map([["(?i:'s|'t|'re|'ve|'m|'ll|'d)","(?:'([sS]|[tT]|[rR][eE]|[vV][eE]|[mM]|[lL][lL]|[dD]))"],["(?i:[sdmt]|ll|ve|re)","(?:[sS]|[dD]|[mM]|[tT]|[lL][lL]|[vV][eE]|[rR][eE])"],["[^\\r\\n\\p{L}\\p{N}]?+","[^\\r\\n\\p{L}\\p{N}]?"],["[^\\s\\p{L}\\p{N}]++","[^\\s\\p{L}\\p{N}]+"],["(?>\\p{Nd}{510})","(?:\\p{Nd}{510})"],["\\p{Nd}{3}+","(?:\\p{Nd}{3})+"],["\\G",""],[` ?[^(\\s|[${Ro}])]+`,` ?[^\\s${Ro}]+`]]),vr="\\p{P}\\u0021-\\u002F\\u003A-\\u0040\\u005B-\\u0060\\u007B-\\u007E",Vr=e=>e.replace(/ \./g,".").replace(/ \?/g,"?").replace(/ \!/g,"!").replace(/ ,/g,",").replace(/ \' /g,"'").replace(/ n't/g,"n't").replace(/ 'm/g,"'m").replace(/ 's/g,"'s").replace(/ 've/g,"'ve").replace(/ 're/g,"'re"),Sr=(e,r=!0)=>{if(e.Regex!==void 0){let t=e.Regex.replace(/\\([#&~])/g,"$1");t=t.replace(/\\A/g,"^").replace(/\\z/g,"$").replace(/\\Z/g,"(?=\\r?\\n?$)");for(let[o,n]of fs)t=t.replaceAll(o,n);try{return new RegExp(t,"gu")}catch(o){if(!(o instanceof SyntaxError)||!o.message.toLowerCase().includes("invalid property name"))throw o;let n=!1,s=t.replace(/(\\[pP])\{([^}=]+)\}/g,(a,i,u)=>{try{return new RegExp(`\\p{${u}}`,"u"),`${i}{${u}}`}catch{return n=!0,`${i}{Script=${u}}`}});if(!n)throw o;try{return new RegExp(s,"gu")}catch{throw o}}}else if(e.String!==void 0){let t=ms(e.String);return new RegExp(r?t:`(${t})`,"gu")}else return console.warn("Unknown pattern type:",e),null},ms=e=>e.replace(/[.*+?^${}()|[\]\\]/g,"\\$&"),gs=(e,r,t)=>{let o=[],n=0;for(;n<e.length;){if(o.push(e[n]),(r.get(e[n])??t)!==t){++n;continue}for(;++n<e.length&&(r.get(e[n])??t)===t;)r.get(o.at(-1))!==t&&(o[o.length-1]+=e[n])}return o},hs=e=>e>=19968&&e<=40959||e>=13312&&e<=19903||e>=131072&&e<=173791||e>=173824&&e<=177983||e>=177984&&e<=178207||e>=178208&&e<=183983||e>=63744&&e<=64255||e>=194560&&e<=195103,bs=e=>Number.isInteger(e)||typeof e=="bigint",_s=e=>{let r=0;for(let t of e)++r;return r},ws=e=>Oo(e.toLowerCase()),we=(...e)=>Array.prototype.concat.apply([],e),Qr=e=>new Map(Object.entries(e)),ys=(e,r)=>{let t=[],o=0;for(let n of e.matchAll(r)){let s=n[0];o<n.index&&t.push(e.slice(o,n.index)),s.length>0&&t.push(s),o=n.index+s.length}return o<e.length&&t.push(e.slice(o)),t},Oo=e=>e.replace(/\p{M}/gu,""),Mo=(e,r,t=[])=>{if(!e||Array.isArray(e)||typeof e!="object")return`${r} must be a valid object`;for(let o of t)if(!(o in e))return`${r} must contain a "${o}" property`;return null},ks=e=>e.match(/\S+/g)||[],vs=class{constructor(){let e=function(...r){return e._call(...r)};return Object.setPrototypeOf(e,new.target.prototype)}},ir=vs,Ss=class extends ir{constructor(e){super(),this.config=e}_call(e){return this.normalize(e)}},$e=Ss,Ps=class extends $e{tokenize_chinese_chars(e){let r=[];for(let t=0;t<e.length;++t){let o=e[t],n=o.charCodeAt(0);hs(n)?(r.push(" "),r.push(o),r.push(" ")):r.push(o)}return r.join("")}strip_accents(e){return e.normalize("NFD").replace(/\p{Mn}/gu,"")}is_control(e){switch(e){case" ":case`1011`:case"\r":return!1;default:return/^\p{Cc}|\p{Cf}|\p{Co}|\p{Cs}$/u.test(e)}}clean_text(e){let r=[];for(let t of e){let o=t.charCodeAt(0);o===0||o===65533||this.is_control(t)||(/^\s$/.test(t)?r.push(" "):r.push(t))}return r.join("")}normalize(e){return this.config.clean_text&&(e=this.clean_text(e)),this.config.handle_chinese_chars&&(e=this.tokenize_chinese_chars(e)),this.config.lowercase?(e=e.toLowerCase(),this.config.strip_accents!==!1&&(e=this.strip_accents(e))):this.config.strip_accents&&(e=this.strip_accents(e)),e}},Es=Ps,Bs=class extends $e{constructor(e){super(e),this.charsmap=e.precompiled_charsmap??null}normalize(e){return e=e.replace(/[\u0001-\u0008\u000B\u000E-\u001F\u007F\u008F\u009F]/gm,""),e=e.replace(/[\u0009\u000A\u000C\u000D\u00A0\u1680\u2000-\u200F\u2028\u2029\u202F\u205F\u2581\u3000\uFEFF\uFFFD]/gm," "),e.includes("\uFF5E")?e=e.split("\uFF5E").map(t=>t.normalize("NFKC")).join("\uFF5E"):e=e.normalize("NFKC"),e}},xs=Bs,Rs=class extends $e{constructor(e){super(e),this.normalizers=(e.normalizers??[]).map(r=>Do(r))}normalize(e){return this.normalizers.reduce((r,t)=>t?t.normalize(r):r,e)}},Ms=Rs,Gs=class extends $e{normalize(e){let r=Sr(this.config.pattern??{});return r===null?e:e.replaceAll(r,this.config.content??"")}},$s=Gs,Ls=class extends $e{constructor(){super(...arguments),this.form="NFC"}normalize(e){return e=e.normalize(this.form),e}},Pr=Ls,As=class extends Pr{constructor(){super(...arguments),this.form="NFC"}},zs=As,Os=class extends Pr{constructor(){super(...arguments),this.form="NFD"}},Ds=Os,Cs=class extends Pr{constructor(){super(...arguments),this.form="NFKC"}},Is=Cs,Ts=class extends Pr{constructor(){super(...arguments),this.form="NFKD"}},qs=Ts,Ns=class extends $e{normalize(e){return this.config.strip_left&&this.config.strip_right?e=e.trim():(this.config.strip_left&&(e=e.trimStart()),this.config.strip_right&&(e=e.trimEnd())),e}},Ws=Ns,Us=class extends $e{normalize(e){return Oo(e)}},Fs=Us,js=class extends $e{normalize(e){return e.toLowerCase()}},Hs=js,Ks=class extends $e{normalize(e){return e=this.config.prepend+e,e}},Vs=Ks;function Qs(e){if(e===null)return null;switch(e.type){case"BertNormalizer":return new Es(e);case"Precompiled":return new xs(e);case"Sequence":return new Ms(e);case"Replace":return new $s(e);case"NFC":return new zs(e);case"NFD":return new Ds(e);case"NFKC":return new Is(e);case"NFKD":return new qs(e);case"Strip":return new Ws(e);case"StripAccents":return new Fs(e);case"Lowercase":return new Hs(e);case"Prepend":return new Vs(e);default:throw new Error(`Unknown Normalizer type: ${e.type}`)}}var Do=Qs,Xs=class extends ir{pre_tokenize(e,r){return(Array.isArray(e)?e.map(t=>this.pre_tokenize_text(t,r)):this.pre_tokenize_text(e,r)).flat()}_call(e,r){return this.pre_tokenize(e,r)}},ye=Xs,Ys=class extends ye{constructor(e){super(),this.config=e,this.add_prefix_space=this.config.add_prefix_space??!1,this.trim_offsets=this.config.trim_offsets??!1,this.use_regex=this.config.use_regex??!0,this.pattern=/'s|'t|'re|'ve|'m|'ll|'d| ?\p{L}+| ?\p{N}+| ?[^\s\p{L}\p{N}]+|\s+(?!\S)|\s+/gu,this.byte_encoder=zo,this.text_encoder=new TextEncoder}pre_tokenize_text(e,r){return this.add_prefix_space&&!e.startsWith(" ")&&(e=" "+e),(this.use_regex?e.match(this.pattern)||[]:[e]).map(o=>Array.from(this.text_encoder.encode(o),n=>this.byte_encoder[n]).join(""))}},Js=Ys,Zs=class extends ye{pre_tokenize_text(e,r){return e.match(/\w+|[^\w\s]+/g)||[]}},ea=Zs,ra=class extends ye{constructor(e){super(),this.replacement=e.replacement??"\u2581",this.str_rep=e.str_rep||this.replacement,this.prepend_scheme=e.prepend_scheme??"always"}pre_tokenize_text(e,r){let{section_index:t=void 0}=r??{},o=e.replaceAll(" ",this.str_rep);return!o.startsWith(this.replacement)&&(this.prepend_scheme==="always"||this.prepend_scheme==="first"&&t===0)&&(o=this.str_rep+o),[o]}},ta=ra,oa=class extends ye{constructor(e){super(),this.config=e,this.pattern=Sr(this.config.pattern??{},this.config.invert??!0)}pre_tokenize_text(e){return this.pattern===null?[]:this.config.invert?e.match(this.pattern)||[]:this.config.behavior?.toLowerCase()==="removed"?e.split(this.pattern).filter(r=>r):ys(e,this.pattern)}},na=oa,sa=class extends ye{constructor(e){super(),this.config=e,this.pattern=new RegExp(`[^${vr}]+|[${vr}]+`,"gu")}pre_tokenize_text(e){return e.match(this.pattern)||[]}},aa=sa,ia=class extends ye{constructor(e){super(),this.config=e;let r=`[^\\d]+|\\d${this.config.individual_digits?"":"+"}`;this.pattern=new RegExp(r,"gu")}pre_tokenize_text(e){return e.match(this.pattern)||[]}},ua=ia,la=class extends ye{constructor(){super(),this.pattern=new RegExp(`[^\\s${vr}]+|[${vr}]`,"gu")}pre_tokenize_text(e,r){return e.trim().match(this.pattern)||[]}},ca=la,pa=class extends ye{constructor(e){super(),this.config=e,this.pattern=Sr(this.config.pattern??{}),this.content=this.config.content??""}pre_tokenize_text(e){return this.pattern===null?[e]:[e.replaceAll(this.pattern,this.config.content??"")]}},da=pa,fa=class extends ye{constructor(e){super(),this.tokenizers=(e.pretokenizers??[]).map(r=>Co(r))}pre_tokenize_text(e,r){return this.tokenizers.reduce((t,o)=>o?o.pre_tokenize(t,r):t,[e])}},ma=fa,ga=class extends ye{pre_tokenize_text(e){return ks(e)}},ha=ga,ba=class extends ye{constructor(e){super(),this.config=e,this._length=e.length}pre_tokenize_text(e){let r=[];for(let t=0;t<e.length;t+=this._length)r.push(e.slice(t,t+this._length));return r}},_a=ba;function wa(e){if(e===null)return null;switch(e.type){case"BertPreTokenizer":return new ca;case"Sequence":return new ma(e);case"Whitespace":return new ea;case"WhitespaceSplit":return new ha;case"Metaspace":return new ta(e);case"ByteLevel":return new Js(e);case"Split":return new na(e);case"Punctuation":return new aa(e);case"Digits":return new ua(e);case"Replace":return new da(e);case"FixedLength":return new _a(e);default:throw new Error(`Unknown PreTokenizer type: ${e.type}`)}}var Co=wa,ya=class extends ir{constructor(e){super(),this.config=e,this.vocab=[],this.tokens_to_ids=new Map,this.unk_token_id=void 0,this.unk_token=void 0,this.end_of_word_suffix=void 0,this.fuse_unk=this.config.fuse_unk??!1}_call(e){let r=this.encode(e);return this.fuse_unk&&(r=gs(r,this.tokens_to_ids,this.unk_token_id)),r}},Er=ya,ka=class extends Er{constructor(e){super(e),this.max_input_chars_per_word=100,this.tokens_to_ids=Qr(e.vocab),this.unk_token_id=this.tokens_to_ids.get(e.unk_token),this.unk_token=e.unk_token,this.max_input_chars_per_word=e.max_input_chars_per_word??100,this.vocab=new Array(this.tokens_to_ids.size);for(let[r,t]of this.tokens_to_ids)this.vocab[t]=r}encode(e){let r=[];for(let t of e){let o=[...t];if(o.length>this.max_input_chars_per_word){r.push(this.unk_token);continue}let n=!1,s=0,a=[];for(;s<o.length;){let i=o.length,u=null;for(;s<i;){let c=o.slice(s,i).join("");if(s>0&&(c=this.config.continuing_subword_prefix+c),this.tokens_to_ids.has(c)){u=c;break}--i}if(u===null){n=!0;break}a.push(u),s=i}n?r.push(this.unk_token):r.push(...a)}return r}},Go=ka,$o=class Io{constructor(r,t){this.is_leaf=r,this.children=t}static default(){return new Io(!1,new Map)}},va=class{constructor(){this.root=$o.default()}extend(e){for(let r of e)this.push(r)}push(e){let r=this.root;for(let t of e){let o=r.children.get(t);o===void 0&&(o=$o.default(),r.children.set(t,o)),r=o}r.is_leaf=!0}*common_prefix_search(e){let r=this.root;if(r===void 0)return;let t="";for(let o of e){if(t+=o,r=r.children.get(o),r===void 0)return;r.is_leaf&&(yield t)}}},Sa=va,Kr=class To{constructor(r,t,o,n,s){this.token_id=r,this.node_id=t,this.pos=o,this.length=n,this.score=s,this.prev=null,this.backtrace_score=0}clone(){let r=new To(this.token_id,this.node_id,this.pos,this.length,this.score);return r.prev=this.prev,r.backtrace_score=this.backtrace_score,r}},Pa=class{constructor(e,r,t){this.chars=Array.from(e),this.len=this.chars.length,this.bos_token_id=r,this.eos_token_id=t,this.nodes=[],this.begin_nodes=Array.from({length:this.len+1},()=>[]),this.end_nodes=Array.from({length:this.len+1},()=>[]);let o=new Kr(this.bos_token_id??0,0,0,0,0),n=new Kr(this.eos_token_id??0,1,this.len,0,0);this.nodes.push(o.clone()),this.nodes.push(n.clone()),this.begin_nodes[this.len].push(n),this.end_nodes[0].push(o)}insert(e,r,t,o){let n=this.nodes.length,s=new Kr(o,n,e,r,t);this.begin_nodes[e].push(s),this.end_nodes[e+r].push(s),this.nodes.push(s)}viterbi(){let e=this.len,r=0;for(;r<=e;){if(this.begin_nodes[r].length==0)return[];for(let a of this.begin_nodes[r]){a.prev=null;let i=0,u=null;for(let c of this.end_nodes[r]){let f=c.backtrace_score+a.score;(u===null||f>i)&&(u=c.clone(),i=f)}if(u!==null)a.prev=u,a.backtrace_score=i;else return[]}++r}let t=[],n=this.begin_nodes[e][0].prev;if(n===null)return[];let s=n.clone();for(;s.prev!==null;)t.push(s.clone()),s=s.clone().prev.clone();return t.reverse(),t}piece(e){return this.chars.slice(e.pos,e.pos+e.length).join("")}tokens(){return this.viterbi().map(r=>this.piece(r))}token_ids(){return this.viterbi().map(r=>r.token_id)}},Ea=Pa;function Ba(e){if(e.length===0)throw new Error("Array must not be empty");let r=e[0],t=0;for(let o=1;o<e.length;++o)e[o]<r&&(r=e[o],t=o);return[r,t]}var xa=class extends Er{constructor(e,r){super(e);let t=e.vocab.length;this.vocab=new Array(t),this.scores=new Array(t);for(let o=0;o<t;++o)[this.vocab[o],this.scores[o]]=e.vocab[o];this.unk_token_id=e.unk_id,this.unk_token=this.vocab[e.unk_id],this.tokens_to_ids=new Map(this.vocab.map((o,n)=>[o,n])),this.bos_token=" ",this.bos_token_id=this.tokens_to_ids.get(this.bos_token),this.eos_token=r,this.eos_token_id=this.tokens_to_ids.get(this.eos_token),this.unk_token=this.vocab[this.unk_token_id],this.min_score=Ba(this.scores)[0],this.unk_score=this.min_score-10,this.scores[this.unk_token_id]=this.unk_score,this.trie=new Sa,this.trie.extend(this.vocab),this.fuse_unk=!0}populate_nodes(e){let r=e.chars,t=1,o=0;for(;o<r.length;){let n=!1,s=[],a=r.slice(o).join(""),i=this.trie.common_prefix_search(a);for(let u of i){s.push(u);let c=this.tokens_to_ids.get(u),f=this.scores[c],k=_s(u);e.insert(o,k,f,c),!n&&k===t&&(n=!0)}n||e.insert(o,t,this.unk_score,this.unk_token_id),o+=t}}tokenize(e){let r=new Ea(e,this.bos_token_id,this.eos_token_id);return this.populate_nodes(r),r.tokens()}encode(e){let r=[];for(let t of e){let o=this.tokenize(t);r.push(...o)}return r}},Lo=xa,Ra=class{constructor(e=(t,o)=>t>o,r=1/0){this._heap=[],this._comparator=e,this._max_size=r}get size(){return this._heap.length}is_empty(){return this.size===0}peek(){return this._heap[0]}push(...e){return this.extend(e)}extend(e){for(let r of e)if(this.size<this._max_size)this._heap.push(r),this._sift_up();else{let t=this._smallest();this._comparator(r,this._heap[t])&&(this._heap[t]=r,this._sift_up_from(t))}return this.size}pop(){let e=this.peek(),r=this.size-1;return r>0&&this._swap(0,r),this._heap.pop(),this._sift_down(),e}replace(e){let r=this.peek();return this._heap[0]=e,this._sift_down(),r}_parent(e){return(e+1>>>1)-1}_left(e){return(e<<1)+1}_right(e){return e+1<<1}_greater(e,r){return this._comparator(this._heap[e],this._heap[r])}_swap(e,r){let t=this._heap[e];this._heap[e]=this._heap[r],this._heap[r]=t}_sift_up(){this._sift_up_from(this.size-1)}_sift_up_from(e){for(;e>0&&this._greater(e,this._parent(e));)this._swap(e,this._parent(e)),e=this._parent(e)}_sift_down(){let e=0;for(;this._left(e)<this.size&&this._greater(this._left(e),e)||this._right(e)<this.size&&this._greater(this._right(e),e);){let r=this._right(e)<this.size&&this._greater(this._right(e),this._left(e))?this._right(e):this._left(e);this._swap(e,r),e=r}}_smallest(){return 2**Math.floor(Math.log2(this.size))-1}},Ma=Ra,Ga=class{constructor(e){this.capacity=e,this.cache=new Map}get(e){if(!this.cache.has(e))return;let r=this.cache.get(e);return this.cache.delete(e),this.cache.set(e,r),r}put(e,r){this.cache.has(e)&&this.cache.delete(e),this.cache.set(e,r),this.cache.size>this.capacity&&this.cache.delete(this.cache.keys().next().value)}clear(){this.cache.clear()}},$a=Ga,La=class extends Er{constructor(e){super(e),this.tokens_to_ids=Qr(e.vocab),this.unk_token_id=this.tokens_to_ids.get(e.unk_token),this.unk_token=e.unk_token,this.vocab=new Array(this.tokens_to_ids.size);for(let[t,o]of this.tokens_to_ids)this.vocab[o]=t;let r=Array.isArray(e.merges[0]);this.merges=r?e.merges:e.merges.map(t=>t.split(" ",2)),this.bpe_ranks=new Map(this.merges.map((t,o)=>[JSON.stringify(t),o])),this.end_of_word_suffix=e.end_of_word_suffix,this.continuing_subword_suffix=e.continuing_subword_suffix??null,this.byte_fallback=this.config.byte_fallback??!1,this.byte_fallback&&(this.text_encoder=new TextEncoder),this.ignore_merges=this.config.ignore_merges??!1,this.max_length_to_cache=256,this.cache_capacity=1e4,this.cache=new $a(this.cache_capacity)}clear_cache(){this.cache.clear()}bpe(e){if(e.length===0)return[];let r=this.cache.get(e);if(r!==void 0)return r;let t=Array.from(e);this.end_of_word_suffix&&(t[t.length-1]+=this.end_of_word_suffix);let o=[];if(t.length>1){let n=new Ma((i,u)=>i.score<u.score),s={token:t[0],bias:0,prev:null,next:null},a=s;for(let i=1;i<t.length;++i){let u={bias:i/t.length,token:t[i],prev:a,next:null};a.next=u,this.add_node(n,a),a=u}for(;!n.is_empty();){let i=n.pop();if(i.deleted||!i.next||i.next.deleted)continue;if(i.deleted=!0,i.next.deleted=!0,i.prev){let c={...i.prev};i.prev.deleted=!0,i.prev=c,c.prev?c.prev.next=c:s=c}let u={token:i.token+i.next.token,bias:i.bias,prev:i.prev,next:i.next.next};u.prev?(u.prev.next=u,this.add_node(n,u.prev)):s=u,u.next&&(u.next.prev=u,this.add_node(n,u))}for(let i=s;i!==null;i=i.next)o.push(i.token)}else o=t;if(this.continuing_subword_suffix)for(let n=0;n<o.length-1;++n)o[n]+=this.continuing_subword_suffix;return e.length<this.max_length_to_cache&&this.cache.put(e,o),o}add_node(e,r){let t=this.bpe_ranks.get(JSON.stringify([r.token,r.next.token]));t!==void 0&&(r.score=t+r.bias,e.push(r))}encode(e){let r=[];for(let t of e){if(this.ignore_merges&&this.tokens_to_ids.has(t)){r.push(t);continue}let o=this.bpe(t);for(let n of o)if(this.tokens_to_ids.has(n))r.push(n);else if(this.byte_fallback){let s=Array.from(this.text_encoder.encode(n)).map(a=>`<0x${a.toString(16).toUpperCase().padStart(2,"0")}>`);s.every(a=>this.tokens_to_ids.has(a))?r.push(...s):this.unk_token!=null&&r.push(this.unk_token)}else this.unk_token!=null&&r.push(this.unk_token)}return r}},Ao=La,Aa=class extends Er{constructor(e,r){super(e);let t=e.vocab;this.tokens_to_ids=Qr(r.target_lang?t[r.target_lang]:t),this.bos_token=r.bos_token,this.bos_token_id=this.tokens_to_ids.get(this.bos_token),this.eos_token=r.eos_token,this.eos_token_id=this.tokens_to_ids.get(this.eos_token),this.pad_token=r.pad_token,this.pad_token_id=this.tokens_to_ids.get(this.pad_token),this.unk_token=r.unk_token,this.unk_token_id=this.tokens_to_ids.get(this.unk_token),this.vocab=new Array(this.tokens_to_ids.size);for(let[o,n]of this.tokens_to_ids)this.vocab[n]=o}encode(e){return e}},za=Aa;function Oa(e,r){switch(e.type){case"WordPiece":return new Go(e);case"Unigram":return new Lo(e,r.eos_token);case"BPE":return new Ao(e);default:if(e.vocab)return Array.isArray(e.vocab)?new Lo(e,r.eos_token):Object.hasOwn(e,"continuing_subword_prefix")&&Object.hasOwn(e,"unk_token")?Object.hasOwn(e,"merges")?new Ao(e):new Go(e):new za(e,{target_lang:r.target_lang,bos_token:r.bos_token,eos_token:r.eos_token,pad_token:r.pad_token,unk_token:r.unk_token});throw new Error(`Unknown TokenizerModel type: ${e?.type}`)}}var Da=Oa,Ca=class extends ir{constructor(e){super(),this.config=e}_call(e,...r){return this.post_process(e,...r)}},ur=Ca,Ia=class extends ur{post_process(e,r=null,t=!0){let o=r===null?this.config.single:this.config.pair,n=[],s=[];for(let a of o)"SpecialToken"in a?t&&(n.push(a.SpecialToken.id),s.push(a.SpecialToken.type_id)):"Sequence"in a&&(a.Sequence.id==="A"?(n=we(n,e),s=we(s,new Array(e.length).fill(a.Sequence.type_id))):a.Sequence.id==="B"&&(n=we(n,r),s=we(s,new Array(r.length).fill(a.Sequence.type_id))));return{tokens:n,token_type_ids:s}}},Ta=Ia,qa=class extends ur{post_process(e,r=null){return{tokens:e,tokens_pair:r}}},Na=qa,Wa=class extends ur{constructor(e){super(e),this.sep=e.sep,this.cls=e.cls}post_process(e,r=null,t=!0){t&&(e=we([this.cls[0]],e,[this.sep[0]]));let o=new Array(e.length).fill(0);if(r){let n=[],s=t?[this.sep[0]]:[];e=we(e,n,r,s),o=we(o,new Array(r.length+n.length+s.length).fill(1))}return{tokens:e,token_type_ids:o}}},Ua=Wa,Fa=class extends ur{constructor(e){super(e),this.sep=e.sep,this.cls=e.cls}post_process(e,r,t=!0){t&&(e=we([this.cls[0]],e,[this.sep[0]]));let o=new Array(e.length).fill(0);if(r){let n=t?[this.sep[0]]:[],s=t?[this.sep[0]]:[];e=we(e,n,r,s),o=we(o,new Array(r.length+n.length+s.length).fill(1))}return{tokens:e,token_type_ids:o}}},ja=Fa,Ha=class extends ur{constructor(e){super(e),this.processors=(e.processors??[]).map(r=>qo(r))}post_process(e,r=null,t=!0){let o={tokens:e,tokens_pair:r};for(let n of this.processors)o=n.post_process(o.tokens,o.tokens_pair,t);return o}},Ka=Ha;function Va(e){if(e===null)return null;switch(e.type){case"TemplateProcessing":return new Ta(e);case"ByteLevel":return new Na(e);case"BertProcessing":return new Ua(e);case"RobertaProcessing":return new ja(e);case"Sequence":return new Ka(e);default:throw new Error(`Unknown PostProcessor type: ${e.type}`)}}var qo=Va,Qa=class extends ir{constructor(e){super(),this.config=e,this.added_tokens=[],this.end_of_word_suffix=null,this.trim_offsets="trim_offsets"in e?e.trim_offsets:!1}_call(e){return this.decode(e)}decode(e){return this.decode_chain(e).join("")}},Se=Qa,Xa=class extends Se{constructor(e){super(e),this.byte_decoder=ds,this.text_decoder=new TextDecoder("utf-8",{fatal:!1,ignoreBOM:!0}),this.end_of_word_suffix=null}convert_tokens_to_string(e){let r=e.join(""),t=new Uint8Array([...r].map(o=>this.byte_decoder[o]));return this.text_decoder.decode(t)}decode_chain(e){let r=[],t=[];for(let o of e)this.added_tokens.find(n=>n.content===o)!==void 0?(t.length>0&&(r.push(this.convert_tokens_to_string(t)),t=[]),r.push(o)):t.push(o);return t.length>0&&r.push(this.convert_tokens_to_string(t)),r}},Ya=Xa,Ja=class extends Se{constructor(e){super(e),this.cleanup=e.cleanup}decode_chain(e){return e.map((r,t)=>{if(t!==0){let o=this.config.prefix;o&&r.startsWith(o)?r=r.replace(o,""):r=" "+r}return this.cleanup&&(r=Vr(r)),r})}},Za=Ja,ei=class extends Se{constructor(e){super(e),this.replacement=e.replacement??"\u2581"}decode_chain(e){let r=[];for(let t=0;t<e.length;++t){let o=e[t].replaceAll(this.replacement," ");t==0&&o.startsWith(" ")&&(o=o.substring(1)),r.push(o)}return r}},ri=ei,ti=class extends Se{constructor(e){super(e),this.suffix=e.suffix??""}decode_chain(e){return e.map((r,t)=>r.replaceAll(this.suffix,t===e.length-1?"":" "))}},oi=ti,ni=class extends Se{constructor(e){super(e),this.pad_token=e.pad_token??"",this.word_delimiter_token=e.word_delimiter_token??"",this.cleanup=e.cleanup}convert_tokens_to_string(e){if(e.length===0)return"";let r=[e[0]];for(let n=1;n<e.length;++n)e[n]!==r.at(-1)&&r.push(e[n]);let o=r.filter(n=>n!==this.pad_token).join("");return this.cleanup&&(o=Vr(o).replaceAll(this.word_delimiter_token," ").trim()),o}decode_chain(e){return[this.convert_tokens_to_string(e)]}},si=ni,ai=class extends Se{constructor(e){super(e),this.decoders=(e.decoders??[]).map(r=>No(r))}decode_chain(e){return this.decoders.reduce((r,t)=>t.decode_chain(r),e)}},ii=ai,ui=class extends Se{decode_chain(e){let r=Sr(this.config.pattern),t=this.config.content??"";return r===null?e:e.map(o=>o.replaceAll(r,t))}},li=ui,ci=class extends Se{decode_chain(e){return[e.join("")]}},pi=ci,di=class extends Se{constructor(e){super(e),this.content=e.content??"",this.start=e.start??0,this.stop=e.stop??0}decode_chain(e){return e.map(r=>{let t=0;for(let n=0;n<this.start&&r[n]===this.content;++n){t=n+1;continue}let o=r.length;for(let n=0;n<this.stop;++n){let s=r.length-n-1;if(r[s]===this.content){o=s;continue}else break}return r.slice(t,o)})}},fi=di,mi=class extends Se{constructor(e){super(e),this.text_decoder=new TextDecoder}decode_chain(e){let r=[],t=[];for(let o of e){let n=null;if(o.length===6&&o.startsWith("<0x")&&o.endsWith(">")){let s=parseInt(o.slice(3,5),16);isNaN(s)||(n=s)}if(n!==null)t.push(n);else{if(t.length>0){let s=this.text_decoder.decode(Uint8Array.from(t));r.push(s),t=[]}r.push(o)}}if(t.length>0){let o=this.text_decoder.decode(Uint8Array.from(t));r.push(o),t=[]}return r}},gi=mi;function hi(e){if(e===null)return null;switch(e.type){case"ByteLevel":return new Ya(e);case"WordPiece":return new Za(e);case"Metaspace":return new ri(e);case"BPEDecoder":return new oi(e);case"CTC":return new si(e);case"Sequence":return new ii(e);case"Replace":return new li(e);case"Fuse":return new pi(e);case"Strip":return new fi(e);case"ByteFallback":return new gi(e);default:throw new Error(`Unknown Decoder type: ${e.type}`)}}var No=hi,bi=class{constructor(e,r){let t=Mo(e,"Tokenizer",["model","decoder","post_processor","pre_tokenizer","normalizer"]);if(t)throw new Error(t);let o=Mo(r,"Config");if(o)throw new Error(o);this.tokenizer=e,this.config=r,this.normalizer=Do(this.tokenizer.normalizer),this.pre_tokenizer=Co(this.tokenizer.pre_tokenizer),this.model=Da(this.tokenizer.model,this.config),this.post_processor=qo(this.tokenizer.post_processor),this.decoder=No(this.tokenizer.decoder),this.special_tokens=[],this.all_special_ids=[],this.added_tokens=[];let n=[],s=[];this.added_tokens_map=new Map;for(let a of this.tokenizer.added_tokens){let i=new cs(a);if(this.added_tokens.push(i),this.model.tokens_to_ids.set(i.content,i.id),this.model.vocab[i.id]=i.content,i.special&&(this.special_tokens.push(i.content),this.all_special_ids.push(i.id)),this.added_tokens_map.set(i.content,i),i.normalized&&this.normalizer!==null){let u=this.normalizer(i.content);s.push(u),this.added_tokens_map.set(u,i)}else n.push(i.content)}(this.config.additional_special_tokens??[]).forEach(a=>{this.special_tokens.includes(a)||this.special_tokens.push(a)}),this.decoder&&(this.decoder.added_tokens=this.added_tokens,this.decoder.end_of_word_suffix=this.model.end_of_word_suffix),this.splitter_unnormalized=new xo(n),this.splitter_normalized=new xo(s),this.remove_space=this.config.remove_space,this.clean_up_tokenization_spaces=this.config.clean_up_tokenization_spaces??!0,this.do_lowercase_and_remove_accent=this.config.do_lowercase_and_remove_accent??!1}encode(e,{text_pair:r=null,add_special_tokens:t=!0,return_token_type_ids:o=null}={}){let{tokens:n,token_type_ids:s}=this.tokenize_helper(e,{text_pair:r,add_special_tokens:t}),a=n.map(u=>this.added_tokens_map.get(u)?.id??this.model.tokens_to_ids.get(u)??this.model.unk_token_id),i={ids:a,tokens:n,attention_mask:new Array(a.length).fill(1)};return o&&s&&(i.token_type_ids=s),i}decode(e,r={}){if(!Array.isArray(e)||e.length===0||!bs(e[0]))throw Error("token_ids must be a non-empty array of integers.");let t=e.map(n=>this.model.vocab[Number(n)]??this.model.unk_token);r.skip_special_tokens&&(t=t.filter(n=>!this.special_tokens.includes(n)));let o=this.decoder?this.decoder(t):t.join(" ");return this.decoder&&this.decoder.end_of_word_suffix&&(o=o.replaceAll(this.decoder.end_of_word_suffix," "),r.skip_special_tokens&&(o=o.trim())),(r.clean_up_tokenization_spaces??this.clean_up_tokenization_spaces)&&(o=Vr(o)),o}tokenize(e,{text_pair:r=null,add_special_tokens:t=!1}={}){return this.tokenize_helper(e,{text_pair:r,add_special_tokens:t}).tokens}encode_text(e){if(e===null)return null;let r=this.splitter_unnormalized.split(e);return r.forEach((t,o)=>{let n=this.added_tokens_map.get(t);n&&(n.lstrip&&o>0&&(r[o-1]=r[o-1].trimEnd()),n.rstrip&&o<r.length-1&&(r[o+1]=r[o+1].trimStart()))}),r.flatMap((t,o)=>{if(t.length===0)return[];if(this.added_tokens_map.has(t))return[t];if(this.remove_space===!0&&(t=t.trim().split(/\s+/).join(" ")),this.do_lowercase_and_remove_accent&&(t=ws(t)),this.normalizer!==null&&(t=this.normalizer(t)),t.length===0)return[];let n=this.splitter_normalized.split(t);return n.forEach((s,a)=>{let i=this.added_tokens_map.get(s);i&&(i.lstrip&&a>0&&(n[a-1]=n[a-1].trimEnd()),i.rstrip&&a<n.length-1&&(n[a+1]=n[a+1].trimStart()))}),n.flatMap(s=>{if(s.length===0)return[];if(this.added_tokens_map.has(s))return[s];let a=this.pre_tokenizer!==null?this.pre_tokenizer(s,{section_index:o}):[s];return this.model(a)})})}tokenize_helper(e,{text_pair:r=null,add_special_tokens:t=!0}){let o=this.encode_text(e),n=this.encode_text(r||null);return this.post_processor?this.post_processor(o,n,t):{tokens:we(o??[],n??[])}}token_to_id(e){return this.model.tokens_to_ids.get(e)}id_to_token(e){return this.model.vocab[e]}get_added_tokens_decoder(){let e=new Map;for(let r of this.added_tokens)e.set(r.id,r);return e}get_vocab(e=!0){let r=new Map;for(let t=0;t<this.model.vocab.length;++t){let o=this.model.vocab[t];(e||!this.added_tokens_map.has(o))&&r.set(o,t)}return r}},Wo=bi;var qe={runtime:{shape:{layerTypes:["conv","conv","attention","conv","attention","conv","attention","conv","attention","conv","attention","conv","attention","conv"],hidden:1024,ffn:2560,vocab:65536,heads:16,kvHeads:8,headDim:64,convCache:3,quantBits:4,quantBlock:32},weightsUrl:"/weights.bin",weightBytes:168647680,sections:[{name:"L0.conv.in_proj.quant",offset:0,bytes:1572864},{name:"L0.conv.in_proj.scales",offset:1572864,bytes:393216},{name:"L0.conv.out_proj.quant",offset:1966080,bytes:524288},{name:"L0.conv.out_proj.scales",offset:2490368,bytes:131072},{name:"L0.mlp.gate_up.quant",offset:2621440,bytes:2621440},{name:"L0.mlp.gate_up.scales",offset:5242880,bytes:655360},{name:"L0.mlp.down.quant",offset:5898240,bytes:1310720},{name:"L0.mlp.down.scales",offset:7208960,bytes:327680},{name:"L1.conv.in_proj.quant",offset:7536640,bytes:1572864},{name:"L1.conv.in_proj.scales",offset:9109504,bytes:393216},{name:"L1.conv.out_proj.quant",offset:9502720,bytes:524288},{name:"L1.conv.out_proj.scales",offset:10027008,bytes:131072},{name:"L1.mlp.gate_up.quant",offset:10158080,bytes:2621440},{name:"L1.mlp.gate_up.scales",offset:12779520,bytes:655360},{name:"L1.mlp.down.quant",offset:13434880,bytes:1310720},{name:"L1.mlp.down.scales",offset:14745600,bytes:327680},{name:"L2.attn.qkv.quant",offset:15073280,bytes:1048576},{name:"L2.attn.qkv.scales",offset:16121856,bytes:262144},{name:"L2.attn.o_proj.quant",offset:16384e3,bytes:524288},{name:"L2.attn.o_proj.scales",offset:16908288,bytes:131072},{name:"L2.mlp.gate_up.quant",offset:17039360,bytes:2621440},{name:"L2.mlp.gate_up.scales",offset:19660800,bytes:655360},{name:"L2.mlp.down.quant",offset:20316160,bytes:1310720},{name:"L2.mlp.down.scales",offset:21626880,bytes:327680},{name:"L3.conv.in_proj.quant",offset:21954560,bytes:1572864},{name:"L3.conv.in_proj.scales",offset:23527424,bytes:393216},{name:"L3.conv.out_proj.quant",offset:23920640,bytes:524288},{name:"L3.conv.out_proj.scales",offset:24444928,bytes:131072},{name:"L3.mlp.gate_up.quant",offset:24576e3,bytes:2621440},{name:"L3.mlp.gate_up.scales",offset:27197440,bytes:655360},{name:"L3.mlp.down.quant",offset:27852800,bytes:1310720},{name:"L3.mlp.down.scales",offset:29163520,bytes:327680},{name:"L4.attn.qkv.quant",offset:29491200,bytes:1048576},{name:"L4.attn.qkv.scales",offset:30539776,bytes:262144},{name:"L4.attn.o_proj.quant",offset:30801920,bytes:524288},{name:"L4.attn.o_proj.scales",offset:31326208,bytes:131072},{name:"L4.mlp.gate_up.quant",offset:31457280,bytes:2621440},{name:"L4.mlp.gate_up.scales",offset:34078720,bytes:655360},{name:"L4.mlp.down.quant",offset:34734080,bytes:1310720},{name:"L4.mlp.down.scales",offset:36044800,bytes:327680},{name:"L5.conv.in_proj.quant",offset:36372480,bytes:1572864},{name:"L5.conv.in_proj.scales",offset:37945344,bytes:393216},{name:"L5.conv.out_proj.quant",offset:38338560,bytes:524288},{name:"L5.conv.out_proj.scales",offset:38862848,bytes:131072},{name:"L5.mlp.gate_up.quant",offset:38993920,bytes:2621440},{name:"L5.mlp.gate_up.scales",offset:41615360,bytes:655360},{name:"L5.mlp.down.quant",offset:42270720,bytes:1310720},{name:"L5.mlp.down.scales",offset:43581440,bytes:327680},{name:"L6.attn.qkv.quant",offset:43909120,bytes:1048576},{name:"L6.attn.qkv.scales",offset:44957696,bytes:262144},{name:"L6.attn.o_proj.quant",offset:45219840,bytes:524288},{name:"L6.attn.o_proj.scales",offset:45744128,bytes:131072},{name:"L6.mlp.gate_up.quant",offset:45875200,bytes:2621440},{name:"L6.mlp.gate_up.scales",offset:48496640,bytes:655360},{name:"L6.mlp.down.quant",offset:49152e3,bytes:1310720},{name:"L6.mlp.down.scales",offset:50462720,bytes:327680},{name:"L7.conv.in_proj.quant",offset:50790400,bytes:1572864},{name:"L7.conv.in_proj.scales",offset:52363264,bytes:393216},{name:"L7.conv.out_proj.quant",offset:52756480,bytes:524288},{name:"L7.conv.out_proj.scales",offset:53280768,bytes:131072},{name:"L7.mlp.gate_up.quant",offset:53411840,bytes:2621440},{name:"L7.mlp.gate_up.scales",offset:56033280,bytes:655360},{name:"L7.mlp.down.quant",offset:56688640,bytes:1310720},{name:"L7.mlp.down.scales",offset:57999360,bytes:327680},{name:"L8.attn.qkv.quant",offset:58327040,bytes:1048576},{name:"L8.attn.qkv.scales",offset:59375616,bytes:262144},{name:"L8.attn.o_proj.quant",offset:59637760,bytes:524288},{name:"L8.attn.o_proj.scales",offset:60162048,bytes:131072},{name:"L8.mlp.gate_up.quant",offset:60293120,bytes:2621440},{name:"L8.mlp.gate_up.scales",offset:62914560,bytes:655360},{name:"L8.mlp.down.quant",offset:63569920,bytes:1310720},{name:"L8.mlp.down.scales",offset:64880640,bytes:327680},{name:"L9.conv.in_proj.quant",offset:65208320,bytes:1572864},{name:"L9.conv.in_proj.scales",offset:66781184,bytes:393216},{name:"L9.conv.out_proj.quant",offset:67174400,bytes:524288},{name:"L9.conv.out_proj.scales",offset:67698688,bytes:131072},{name:"L9.mlp.gate_up.quant",offset:67829760,bytes:2621440},{name:"L9.mlp.gate_up.scales",offset:70451200,bytes:655360},{name:"L9.mlp.down.quant",offset:71106560,bytes:1310720},{name:"L9.mlp.down.scales",offset:72417280,bytes:327680},{name:"L10.attn.qkv.quant",offset:72744960,bytes:1048576},{name:"L10.attn.qkv.scales",offset:73793536,bytes:262144},{name:"L10.attn.o_proj.quant",offset:74055680,bytes:524288},{name:"L10.attn.o_proj.scales",offset:74579968,bytes:131072},{name:"L10.mlp.gate_up.quant",offset:74711040,bytes:2621440},{name:"L10.mlp.gate_up.scales",offset:77332480,bytes:655360},{name:"L10.mlp.down.quant",offset:77987840,bytes:1310720},{name:"L10.mlp.down.scales",offset:79298560,bytes:327680},{name:"L11.conv.in_proj.quant",offset:79626240,bytes:1572864},{name:"L11.conv.in_proj.scales",offset:81199104,bytes:393216},{name:"L11.conv.out_proj.quant",offset:81592320,bytes:524288},{name:"L11.conv.out_proj.scales",offset:82116608,bytes:131072},{name:"L11.mlp.gate_up.quant",offset:82247680,bytes:2621440},{name:"L11.mlp.gate_up.scales",offset:84869120,bytes:655360},{name:"L11.mlp.down.quant",offset:85524480,bytes:1310720},{name:"L11.mlp.down.scales",offset:86835200,bytes:327680},{name:"L12.attn.qkv.quant",offset:87162880,bytes:1048576},{name:"L12.attn.qkv.scales",offset:88211456,bytes:262144},{name:"L12.attn.o_proj.quant",offset:88473600,bytes:524288},{name:"L12.attn.o_proj.scales",offset:88997888,bytes:131072},{name:"L12.mlp.gate_up.quant",offset:89128960,bytes:2621440},{name:"L12.mlp.gate_up.scales",offset:91750400,bytes:655360},{name:"L12.mlp.down.quant",offset:92405760,bytes:1310720},{name:"L12.mlp.down.scales",offset:93716480,bytes:327680},{name:"L13.conv.in_proj.quant",offset:94044160,bytes:1572864},{name:"L13.conv.in_proj.scales",offset:95617024,bytes:393216},{name:"L13.conv.out_proj.quant",offset:96010240,bytes:524288},{name:"L13.conv.out_proj.scales",offset:96534528,bytes:131072},{name:"L13.mlp.gate_up.quant",offset:96665600,bytes:2621440},{name:"L13.mlp.gate_up.scales",offset:99287040,bytes:655360},{name:"L13.mlp.down.quant",offset:99942400,bytes:1310720},{name:"L13.mlp.down.scales",offset:101253120,bytes:327680},{name:"head.quant",offset:101580800,bytes:33554432},{name:"head.scales",offset:135135232,bytes:8388608},{name:"head.zero_points",offset:143523840,bytes:1048576},{name:"embed.quant",offset:101580800,bytes:33554432},{name:"embed.scales",offset:135135232,bytes:8388608},{name:"embed.zero_points",offset:143523840,bytes:1048576},{name:"L0.conv.in_proj.norm",offset:144572416,bytes:4096},{name:"L0.conv.core.taps",offset:144576512,bytes:12288},{name:"L0.mlp.gate_up.norm",offset:144588800,bytes:4096},{name:"L1.conv.in_proj.norm",offset:144592896,bytes:4096},{name:"L1.conv.core.taps",offset:144596992,bytes:12288},{name:"L1.mlp.gate_up.norm",offset:144609280,bytes:4096},{name:"L2.attn.qkv.norm",offset:144613376,bytes:4096},{name:"L2.attn.qkv.q_norm",offset:144617472,bytes:256},{name:"L2.attn.qkv.k_norm",offset:144617728,bytes:256},{name:"L2.mlp.gate_up.norm",offset:144617984,bytes:4096},{name:"L3.conv.in_proj.norm",offset:144622080,bytes:4096},{name:"L3.conv.core.taps",offset:144626176,bytes:12288},{name:"L3.mlp.gate_up.norm",offset:144638464,bytes:4096},{name:"L4.attn.qkv.norm",offset:144642560,bytes:4096},{name:"L4.attn.qkv.q_norm",offset:144646656,bytes:256},{name:"L4.attn.qkv.k_norm",offset:144646912,bytes:256},{name:"L4.mlp.gate_up.norm",offset:144647168,bytes:4096},{name:"L5.conv.in_proj.norm",offset:144651264,bytes:4096},{name:"L5.conv.core.taps",offset:144655360,bytes:12288},{name:"L5.mlp.gate_up.norm",offset:144667648,bytes:4096},{name:"L6.attn.qkv.norm",offset:144671744,bytes:4096},{name:"L6.attn.qkv.q_norm",offset:144675840,bytes:256},{name:"L6.attn.qkv.k_norm",offset:144676096,bytes:256},{name:"L6.mlp.gate_up.norm",offset:144676352,bytes:4096},{name:"L7.conv.in_proj.norm",offset:144680448,bytes:4096},{name:"L7.conv.core.taps",offset:144684544,bytes:12288},{name:"L7.mlp.gate_up.norm",offset:144696832,bytes:4096},{name:"L8.attn.qkv.norm",offset:144700928,bytes:4096},{name:"L8.attn.qkv.q_norm",offset:144705024,bytes:256},{name:"L8.attn.qkv.k_norm",offset:144705280,bytes:256},{name:"L8.mlp.gate_up.norm",offset:144705536,bytes:4096},{name:"L9.conv.in_proj.norm",offset:144709632,bytes:4096},{name:"L9.conv.core.taps",offset:144713728,bytes:12288},{name:"L9.mlp.gate_up.norm",offset:144726016,bytes:4096},{name:"L10.attn.qkv.norm",offset:144730112,bytes:4096},{name:"L10.attn.qkv.q_norm",offset:144734208,bytes:256},{name:"L10.attn.qkv.k_norm",offset:144734464,bytes:256},{name:"L10.mlp.gate_up.norm",offset:144734720,bytes:4096},{name:"L11.conv.in_proj.norm",offset:144738816,bytes:4096},{name:"L11.conv.core.taps",offset:144742912,bytes:12288},{name:"L11.mlp.gate_up.norm",offset:144755200,bytes:4096},{name:"L12.attn.qkv.norm",offset:144759296,bytes:4096},{name:"L12.attn.qkv.q_norm",offset:144763392,bytes:256},{name:"L12.attn.qkv.k_norm",offset:144763648,bytes:256},{name:"L12.mlp.gate_up.norm",offset:144763904,bytes:4096},{name:"L13.conv.in_proj.norm",offset:144768e3,bytes:4096},{name:"L13.conv.core.taps",offset:144772096,bytes:12288},{name:"L13.mlp.gate_up.norm",offset:144784384,bytes:4096},{name:"head.norm",offset:144788480,bytes:4096},{name:"rope.cos",offset:144792576,bytes:131072},{name:"rope.sin",offset:144923648,bytes:131072},{name:"L0.mlp.down.kquant",offset:145054720,bytes:1310720},{name:"L0.mlp.down.kscales",offset:146365440,bytes:327680},{name:"L1.mlp.down.kquant",offset:146693120,bytes:1310720},{name:"L1.mlp.down.kscales",offset:148003840,bytes:327680},{name:"L2.mlp.down.kquant",offset:148331520,bytes:1310720},{name:"L2.mlp.down.kscales",offset:149642240,bytes:327680},{name:"L3.mlp.down.kquant",offset:149969920,bytes:1310720},{name:"L3.mlp.down.kscales",offset:151280640,bytes:327680},{name:"L4.mlp.down.kquant",offset:151608320,bytes:1310720},{name:"L4.mlp.down.kscales",offset:152919040,bytes:327680},{name:"L5.mlp.down.kquant",offset:153246720,bytes:1310720},{name:"L5.mlp.down.kscales",offset:154557440,bytes:327680},{name:"L6.mlp.down.kquant",offset:154885120,bytes:1310720},{name:"L6.mlp.down.kscales",offset:156195840,bytes:327680},{name:"L7.mlp.down.kquant",offset:156523520,bytes:1310720},{name:"L7.mlp.down.kscales",offset:157834240,bytes:327680},{name:"L8.mlp.down.kquant",offset:158161920,bytes:1310720},{name:"L8.mlp.down.kscales",offset:159472640,bytes:327680},{name:"L9.mlp.down.kquant",offset:159800320,bytes:1310720},{name:"L9.mlp.down.kscales",offset:161111040,bytes:327680},{name:"L10.mlp.down.kquant",offset:161438720,bytes:1310720},{name:"L10.mlp.down.kscales",offset:162749440,bytes:327680},{name:"L11.mlp.down.kquant",offset:163077120,bytes:1310720},{name:"L11.mlp.down.kscales",offset:164387840,bytes:327680},{name:"L12.mlp.down.kquant",offset:164715520,bytes:1310720},{name:"L12.mlp.down.kscales",offset:166026240,bytes:327680},{name:"L13.mlp.down.kquant",offset:166353920,bytes:1310720},{name:"L13.mlp.down.kscales",offset:167664640,bytes:327680},{name:"rope.cos_ext",offset:167992320,bytes:327680},{name:"rope.sin_ext",offset:16832e4,bytes:327680}],maxQueries:1,maxPositions:2560,samplePartials:256,planPosition:1024,maxTotalPositions:2112,knee:[{shape:"attn_o_proj",geometry:{workgroupSize:128,subgroupSize:32,rowsPerSubgroup:1,subgroupsPerRow:1}},{shape:"mlp_down",geometry:{workgroupSize:128,subgroupSize:32,rowsPerSubgroup:1,subgroupsPerRow:1}},{shape:"attn_qkv",geometry:{workgroupSize:128,subgroupSize:32,rowsPerSubgroup:1,subgroupsPerRow:1}},{shape:"conv_in_proj",geometry:{workgroupSize:128,subgroupSize:32,rowsPerSubgroup:4,subgroupsPerRow:1}},{shape:"mlp_gate_up",geometry:{workgroupSize:128,subgroupSize:32,rowsPerSubgroup:1,subgroupsPerRow:1}},{shape:"head",geometry:{workgroupSize:128,subgroupSize:32,rowsPerSubgroup:4,subgroupsPerRow:1}}],dequantVariant:"carried",flashOptima:[{positions:192,blocks:16},{positions:1024,blocks:16},{positions:2048,blocks:16}],extraFeatures:["timestamp-query","shader-f16","subgroups"],powerPreference:"high-performance",expectedSubgroupSize:32},candidate:{key:"MIX_M40",flash:!0,unfused:!1,padding:0,flashMode:"subgroup",matvec:{attn_o_proj:{geometry:{workgroupSize:512,subgroupSize:32,rowsPerSubgroup:1,subgroupsPerRow:1},variant:"v3_both"},conv_out_proj:{geometry:{workgroupSize:256,subgroupSize:32,rowsPerSubgroup:1,subgroupsPerRow:1},variant:"v3_both"},attn_qkv:{geometry:{workgroupSize:512,subgroupSize:32,rowsPerSubgroup:4,subgroupsPerRow:1},variant:"carried"},conv_in_proj:{geometry:{workgroupSize:512,subgroupSize:32,rowsPerSubgroup:6,subgroupsPerRow:1},variant:"carried"},head:{geometry:{workgroupSize:256,subgroupSize:32,rowsPerSubgroup:8,subgroupsPerRow:1},variant:"v3_both"}},mlpMode:"subgroup",mlpSlices:40,mlpWorkgroup:1024,normSubgroup:!0,convFuse:!0,flashOnline:!0,flashBlocks:32,mlpPair:!0,foldLanes:16,foldSubgroup:!0,mlpHalfDot:"gate",mlpHalfProduct:"down"}};var Le={prompt:[1,1098,809,43861,1726,11099,6502,797,1210,3446,4873,1973,14367,521,14925,779,7921,521,779,766,2259,521,810,779,6289,803,779,12546,896,7410,779,2129,2076,5935,523,2357,779,4976,803,4843,779,3848,14118,988,779,31368,810,779,7980,10357,37370,1626,779,19814,521,810,968,4250,787,61604,779,27651,1320,8557],continuation:[523,941,1726,11099,1236,8366,779,2080,803,12546,896,1320,7410,779,2129,2076,5935,523,509,1098,1726,11099,6502,535,509,1098,809,43861,1726,11099,6502,797,1210,3446,4873,1973,14367,521,14925,779,7921,521,779,766,2259,521,810,779,6289,803,779,12546,896,7410,779,2129,2076,5935,523,509,17978,509,1098,809,43861,1726,11099,6502,797,1210,3446,4873,1973,14367,521,14925,779,7921,521,779,766,2259,521,810,779,6289,803,779,12546,896,7410,779,2129,2076,5935,523,509,17978,509,1098,1726,11099,6502,535,509,1098,809,43861,1726,11099,6502,797,1210,3446,4873,1973,14367,521,14925,779,7921,521,779,766,2259,521,810,779,6289,803,779,12546,896,1320,7410,779,2129,2076,5935,523,509,17978,509,1098,1726,11099,6502,535,509,1098,809,43861,1726,11099,6502,797,1210,3446,4873,1973,4915,521,29379,779,7921,521,779,766,2259,521,810,779,6289,803,779,12546,896,1320,7410,779,2129,2076,5935,523,509,17978,509,1098,1726,11099,6502,535,509,1098,809,43861,1726,11099,6502,797,1210,3446,4873,1973,4915,521,14925,779,7921,521,779,766,2259,521,810,779,6289,803,779,12546,896,1320,7410,779,2129,2076,5935,523,509,17978,509,1098,1726,11099,6502,535,509,1098,809,43861,1726,11099,6502,797,1210,3446,4873,1973,4915,521,29379,779,7921,521,779,766,2259,521,810,779,6289,803,779,12546,896,1320,7410,779,2129,2076,5935,523,7,2,1,708,1098,2773,856,768,14500,803,779,2120,4520,803,779,4277,523,509,1098,4277,29617,779,7083,803,779,809,43861,797,779,36508,5294,523,509,1098,4277,11961,896,809,17307,17662,938,5175,875,20093,810,6155,797,779,5547,523,2216,3176,768,6260,7848,875,12546,811,28062,1149,1959,1892,811,23303,523,509,1098,4277,1236,22587,779,5160,3616,803,809,17307,17662,884,779,36508,5294,523,509,1098,4277,29617,779,3727,4063,803,809,17307,17662,810,1149,3338,4642,523,509,1098,4277,35283,896,779,2780,803,40686,810,14881,809,17307,17662,19968,9646,884,779,6249,810,779,2432,803,809,43861,523,509,1098,4277,1236,29617,779,5888,803,809,17307,17662,875,36508,6155,810,6368,523,509,1098,4277,35761,968,61498,779,7083,803,809,17307,17662,797,779,36508,5294,810,779,1595,875,6806,10961,797,1149,12003,810,1287,30132,523,7,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,558,535,3747,856,779,5410,3560,803,779,997,552,44413,14174,511,5116,797,779,5180,803,779,2704,3008,1784,542,535,941,5410,3560,803,779,997,552,44413,14174,511,5116,797,779,5180,803,779,2704,3008,856,811,3176,2203,810,4768,4418,811,16310,3869,3609,810,7971,523,1470,5260,11552,1559,906,16310,24661,521,19290,521,6879,5180,521,810,8088,6925,523,941,5116,16070,811,1801,7039,2846,779,13207,803,16310,797,779,2704,3008,810,1352,3671,797,40402,779,3374,1090,9552,810,5009,523,14633,521,936,1400,3644,2203,884,16310,4072,521,3316,521,8022,521,810,1312,6879,5906,896,938,8611,811,779,2704,3008,523,941,6341,856,811,3755,768,12336,810,13914,7582,875,7039,9368,797,3481,1209,1337,16310,3869,810,5009,523,7,2,1,2,1,2,1,2,1,2,1,2,1,2,558,535,3747,856,779,5410,1947,803,779,997,552,44413,14174,511,797,779,5180,803,779,2704,3008,1784,542,535,941,5410,1947,803,779,997,552,44413,14174,511,797,779,2704,3008,856,811,3176,768,3255,803,6939,810,5139,3136,7712,521,5777,797,779,5180,803,3709,9552,810,6879,10607,523,1311,14428,906,768,12447,1685,779,2704,3008,810,779,16310,32182,521,51008,779,10607,803,6157,521,16954,521,810,6879,7971,523,1470,5260,779,1517,803,16310,26606,521,31931,2739,521,810,19656,896,938,7968,1593,797,2473,5009,521,1559,906,997,550,6217,768,16310,2512,997,550,6217,2473,2512,810,997,552,44413,2512,1144,938,16944,811,2473,9552,810,3129,523,14633,521,779,3869,14428,906,768,6690,875,17954,6879,5139,810,32588,521,63879,768,5033,803,9071,19504,810,2624,4079,3136,7712,810,8344,998,25848,523,3860,19623,916,16310,3869,810,5009,521,7712,1011,7425,19393,1419,779,4121,1535,803,779,16310,32182,1090,3129,810,779,14674,6879,15300,1685,779,1547,11502,523,7,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,558,535,3747,856,779,5410,1947,803,779,997,552,44413,14174,511,797,779,5180,803,779,2704,3008,1784,542,535,941,5410,1947,803,779,16310,14174,856,811,18655,6939,810,5139,1685,7712,810,8344,998,521,810,811,9679,6879,10607,810,23615,5139,523,1311,14428,906,768,12447,1685,779,1547,13225,521,24793,7712,811,13515,916,8344,998,797,15668,2105,810,811,21904,779,8345,2953,803,16310,5009,810,16954,523,941,16310,14174,856,1593,797,3727,5844,521,1559,906,18077,3869,521,4985,3869,521,810,6544,6939,521,810,10030,768,3069,3671,797,40402,779,6879,9552,810,4396,803,2059,11502,523,14633,521,936,856,1593,797,3727,5844,803,5075,810,3886,521,1559,906,16310,8587,521,21786,521,810,5936,521,811,19461,6879,19393,810,9679,5008,59352,5139,523,7,708,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,558,535,3747,856,779,5410,1947,803,779,997,552,44413,14174,511,797,779,5180,803,779,2704,3008,1784,542,535,941,5410,1947,803,779,16310,14174,856,811,18655,6939,810,5139,1685,7712,810,8344,998,521,810,811,9679,6879,10607,810,23615,5139,523,1311,14428,906,768,12447,1685,779,1547,13225,521,24793,7712,811,13515,916,8344,998,797,15668,2105,810,811,21904,779,8345,2953,803,16310,5009,810,16954,523,941,16310,14174,856,1593,797,3727,5844,521,1559,906,18077,3869,521,4985,3869,521,810,6544,6939,521,810,10030,768,3069,3671,797,40402,779,6879,9552,810,4396,803,2059,11502,523,1311,856,5175,875,7712,811,2111,810,21904,16310,5009,810,16954,811,27927,3020,6939,810,5139,1685,779,1547,4178,523,7,708,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,558,535,3747,856,779,5410,1947,803,779,997,552,44413,14174,511,797,779,5180,803,779,2704,3008,1784,542,535,941,5410,1947,803,779,16310,14174,856,811,18655,6939,810,5139,1685,7712,810,8344,998,521,810,811,9679,6879,10607,810,23615,5139,523,1311,14428,906,768,12447,1685,779,1547,13225,521,24793,7712,811,13515,916,8344,998,797,15668,2105,810,811,21904,779,8345,2953,803,16310,5009,810,16954,523,941,16310,14174,856,1593,797,3727,5844,521,1559,906,18077,3869,521,4985,3869,521,810,6544,6939,521,810,10030,768,3069,3671,797,40402,779,6879,9552,810,4396,803,2059,11502,523,1311,856,5175,875,7712,811,2111,810,21904,16310,5009,810,16954,811,27927,3020,6939,810,5139,1685,779,1547,4178,523,7,708,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,558,535,3747,856,779,5410,1947,803,779,997,552,44413,14174,511,797,779,5180,803,779,2704,3008,1784,542,535,941,5410,1947,803,779,16310,14174,856,811,18655,6939,810,5139,1685,7712,810,8344,998,521,810,811,9679,6879,10607,810,23615,5139,523,1311,14428,906,768,12447,1685,779,1547,13225,521,24793,7712,811,13515,916,8344,998,797,15668,2105,810,811,21904,779,8345,2953,803,16310,5009,810,16954,523,941,16310,14174,856,1593,797,3727,5844,521,1559,906,18077,3869,521,4985,3869,521,810,6544,6939,521,810,10030,768,3069,3671,797,40402,779,6879,9552,810,4396,803,2059,11502,523,1311,856,5175,875,7712,811,2111,810,21904,16310,5009,810,16954,811,27927,3020,6939,810,5139,1685,779,1547,4178,523,7,708,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,1,2,558,535,3747,856,779,5410,1947,803,779,997,552,44413,14174,511,797,779,5180,803]};var Uo={repo:"inductiveML/LFM2.5-230M-MONARCH",revision:"317e0b32ff55684039432eb65ee003a2e4667e97",tokenizerBase:"https://huggingface.co/inductiveML/LFM2.5-230M-MONARCH/resolve/317e0b32ff55684039432eb65ee003a2e4667e97/"};var H=(e,r={})=>postMessage({type:e,...r}),Fo=()=>new Promise(e=>setTimeout(e,0)),Xr=async e=>Array.from(new Uint8Array(await crypto.subtle.digest("SHA-256",e)),r=>r.toString(16).padStart(2,"0")).join(""),ae=null,Br=null,Yr=!1,Ne=!1,lr=null,We=!1,xr="WebGPU device";function cr(){let e=globalThis.__probeHooks;if(e.errors().length||e.lost().length)throw new Error("The GPU reported an error. Unload the model and try again.")}async function jo(){if(!navigator.gpu)return{supported:!1,reason:"WebGPU is unavailable. Open this page in a recent Chrome or Edge browser on a supported computer."};let e=await navigator.gpu.requestAdapter({powerPreference:"high-performance"});if(!e)return{supported:!1,reason:"No WebGPU adapter is available in this browser."};let r=["shader-f16","subgroups"].filter(o=>!e.features.has(o));if(r.length)return{supported:!1,reason:`This decoder needs ${r.join(" and ")}. Your browser does not expose ${r.length===1?"this feature":"these features"}; the model has not been downloaded.`};let t=e.info;return t.subgroupMinSize!==32||t.subgroupMaxSize!==32?{supported:!1,reason:`This kernel is compiled for 32-lane subgroups. This adapter reports ${t.subgroupMinSize??"?"}\u2013${t.subgroupMaxSize??"?"}. No compatible kernel is available here.`}:e.limits.maxComputeInvocationsPerWorkgroup<1024||e.limits.maxComputeWorkgroupSizeX<1024?{supported:!1,reason:"This GPU does not support the 1,024-thread workgroups used by MIX_M40. The model has not been downloaded."}:e.limits.maxStorageBufferBindingSize<qe.runtime.weightBytes?{supported:!1,reason:"This adapter cannot bind the 169 MB model buffer required by this decoder."}:t.isFallbackAdapter?{supported:!1,reason:"A hardware WebGPU adapter is required; a software adapter was detected."}:(We=e.features.has("timestamp-query"),xr=[t.description||t.vendor,t.architecture].filter(Boolean).join(" \xB7 ")||"WebGPU hardware adapter",{supported:!0,device:xr,timestamps:We,subgroup:32})}async function ki(){let e=await fetch("/monarch/model-manifest.json",{signal:lr.signal});if(!e.ok)throw new Error("The model manifest could not be loaded.");let r=await e.json();if(r.weightSha256!=="ca1fea89fd9f3ca7e5d6d5720705c96f457a92edd0cfc811cf8d9391ac784cb0"||r.weightBytes!==qe.runtime.weightBytes)throw new Error("The published model manifest does not match the tested decoder.");let t=new Uint8Array(r.weightBytes),o=0,n=null;try{n=await caches.open("monarch-weights-ca1fea89-v1")}catch{}for(let s of r.chunks){if(Ne)throw new DOMException("Stopped","AbortError");let a=null,i=!1,u=await n?.match(s.url).catch(()=>{});if(u){let c=await u.arrayBuffer();c.byteLength===s.bytes&&await Xr(c)===s.sha256?(a=c,i=!0):await n?.delete(s.url)}if(!a){let c=await fetch(s.url,{signal:lr.signal});if(!c.ok||!c.body)throw new Error(`Model download failed (${c.status}). Try loading again.`);let f=new Uint8Array(s.bytes),k=c.body.getReader(),y=0,P=0;for(;;){let{value:B,done:S}=await k.read();if(S)break;if(y+B.byteLength>f.byteLength)throw new Error("Unexpected model chunk size.");f.set(B,y),y+=B.byteLength,performance.now()-P>80&&(H("progress",{stage:"Downloading weights",loaded:o+y,total:r.weightBytes}),P=performance.now())}if(a=f.buffer,y!==s.bytes||await Xr(a)!==s.sha256)throw new Error("Model integrity check failed. Please retry the download.");try{await n?.put(s.url,new Response(a,{headers:{"Content-Type":"application/octet-stream"}}))}catch{}}t.set(new Uint8Array(a),o),o+=a.byteLength,H("progress",{stage:i?"Reading verified cached weights":"Downloading weights",loaded:o,total:r.weightBytes})}if(o!==r.weightBytes||await Xr(t.buffer)!==r.weightSha256)throw new Error("The complete model checksum did not match.");return t}async function vi(){let e=await jo();if(H("support",e),!e.supported)return;lr=new AbortController,Ne=!1;let r={...qe.runtime,extraFeatures:qe.runtime.extraFeatures.filter(u=>u!=="timestamp-query"||We)};H("progress",{stage:"Preparing your GPU",loaded:0,total:qe.runtime.weightBytes}),await wr(Hr(r),"high-performance",32);let t=await ki();H("progress",{stage:"Loading tokenizer",loaded:t.length,total:t.length});let[o,n]=await Promise.all(["tokenizer.json","tokenizer_config.json"].map(async u=>{let c=await fetch(Uo.tokenizerBase+u,{signal:lr.signal});if(!c.ok)throw new Error("Tokenizer download failed.");return c.json()}));Br=new Wo(o,n),_t([qe.candidate]),H("progress",{stage:"Compiling the 72-dispatch decoder",loaded:t.length,total:t.length});let s=URL.createObjectURL(new Blob([t]));try{ae=(await So({...r,weightsUrl:s})).session}finally{URL.revokeObjectURL(s)}ae.runtime.setArm("MIX_M40"),cr(),H("progress",{stage:"Checking a known 32-token continuation",loaded:t.length,total:t.length}),await kr(ae,{tokenIds:Uint32Array.from(Le.prompt),promptLength:Le.prompt.length,seed:111008,freerun:!0});let a=new Uint32Array(await or(ae.gpu.device,ae.runtime.buffers.sampled,(Le.prompt.length-1)*4,4))[0],i=await Rr(32,Le.prompt.length);if(a!==Le.continuation[0]||i.sampled.some((u,c)=>u!==Le.continuation[c+1]))throw new Error("The decoder failed its known-answer check on this GPU. This device is not supported by the current kernel.");cr(),H("loaded",{device:xr,timestamps:We,verifiedTokens:32,dispatches:ae.runtime.dispatchCount,weightBytes:qe.runtime.weightBytes})}function Rr(e,r){return Bo(ae,{tokens:e,startPosition:r,startStep:r,readbackEvery:32,ringTokens:64,timestamps:We,freerun:!0,seed:111008})}function Jr(e){return e.length?Br.decode(e,{skip_special_tokens:!0}):""}async function Si(e,r){if(!e.trim())throw new Error("Enter a prompt first.");let t=`<|startoftext|><|im_start|>system1012You are a helpful assistant trained by Liquid AI.<|im_end|>1013<|im_start|>user1014${e.trim()}<|im_end|>1015<|im_start|>assistant1016`,o=Br.encode(t,{add_special_tokens:!1}).ids;if(o.length>1024)throw new Error(`Your prompt is ${o.length.toLocaleString()} tokens. Shorten it to fit the demo\u2019s 1,024-token prompt limit.`);let n=Br.token_to_id("<|im_end|>"),s=[],a=performance.now();H("generating",{promptTokens:o.length}),await kr(ae,{tokenIds:Uint32Array.from(o),promptLength:o.length,seed:111008,freerun:!0});let i=new Uint32Array(await or(ae.gpu.device,ae.runtime.buffers.sampled,(o.length-1)*4,4))[0],u=performance.now(),c=u-a,f=0,k=0;i!==n&&s.push(i),H("text",{text:Jr(s),tokens:s.length,tps:null,ttft:c});let y=i===n,P=o.length;for(;!Ne&&!y&&s.length<r;){let B=await Rr(32,P);P+=32,f+=32,k+=Math.max(0,B.gpuBusyMs);for(let v of B.sampled){if(v===n){y=!0;break}if(s.length>=r)break;s.push(v)}cr();let S=performance.now()-u;H("text",{text:Jr(s),tokens:s.length,tps:s.length>1?(s.length-1)*1e3/S:null,ttft:c,gpuTps:We?f*1e3/k:null}),await Fo()}cr(),H("generated",{text:Jr(s),tokens:s.length,tps:s.length>1?(s.length-1)*1e3/(performance.now()-u):null,ttft:c,stopped:Ne,promptTokens:o.length,computedTokens:f+1})}async function Pi(e){if(![192,1024].includes(e))throw new Error("Unsupported benchmark context.");let r=[...Le.prompt,...Le.continuation],t=e-128,o=[],n=()=>kr(ae,{tokenIds:Uint32Array.from(r.slice(0,t)),promptLength:t,seed:111008,freerun:!0});H("benchmarkProgress",{completed:0,total:4,stage:"Warming the decoder"}),await n(),await Rr(64,t);for(let i=0;i<4&&!Ne;i++){await n(),H("benchmarkProgress",{completed:i,total:4,stage:`Measuring run ${i+1} of 4`});let u=await Rr(256,t);if(cr(),o.length&&u.sampled.some((c,f)=>c!==o[0].ids[f]))throw new Error("Benchmark continuations differed between repeats. No valid result is reported.");o.push({wallMs:u.wallMs,tps:256e3/u.wallMs,gpuBusyMs:We?u.gpuBusyMs:null,ids:u.sampled}),H("benchmarkProgress",{completed:i+1,total:4,stage:`Run ${i+1}: ${(256e3/u.wallMs).toFixed(0)} tok/s`,tps:256e3/u.wallMs}),await Fo()}if(Ne){H("cancelled");return}let s=o.map(i=>i.wallMs/256).sort((i,u)=>i-u),a=(s[1]+s[2])/2;H("benchmarkResult",{depth:e,tokens:256,runs:4,ms:a,tps:1e3/a,runTps:o.map(i=>i.tps),spreadPct:100*(s[3]-s[0])/a,device:xr,measuredAt:new Date().toISOString(),timestamps:We,exactRepeats:!0,protocol:"4 \xD7 256 free-running tokens; 64-token warmup; prefill excluded; 32-token readback through a 64-token ring; wall timer includes encoding and queue completion. Context is the center of the 256-token window. This quick browser run is not the original controlled lab confirmation."})}onmessage=async({data:e})=>{if(e.type==="cancel"){Ne=!0,lr?.abort();return}if(e.type==="support"){try{H("support",await jo())}catch(r){H("support",{supported:!1,reason:String(r)})}return}if(!Yr){Yr=!0,Ne=!1;try{e.type==="load"?await vi():e.type==="generate"&&ae?await Si(String(e.prompt??""),Math.min(512,Math.max(32,Number(e.limit)||256))):e.type==="benchmark"&&ae&&await Pi(Number(e.depth))}catch(r){H("error",{message:r instanceof Error?r.message:String(r),fatal:!ae||!!globalThis.__probeHooks.lost().length})}finally{Yr=!1}}};1017