diff --git a/.gitignore b/.gitignore index 2ee8dcb..284ba4c 100644 --- a/.gitignore +++ b/.gitignore @@ -1,4 +1,3 @@ node_modules .idea .DS_STORE -dist diff --git a/CLAUDE.md b/CLAUDE.md deleted file mode 100644 index 769cf78..0000000 --- a/CLAUDE.md +++ /dev/null @@ -1,284 +0,0 @@ -# CLAUDE.md - -This file provides guidance to Claude Code (claude.ai/code) when working with code in this repository. - -## Development Commands - -```bash -npm run dev # Start Vite dev server -``` - -## Project Structure - -``` -src/ -├── GPUEnv.ts # WebGPU device singleton -├── main.ts # Entry point -├── Trainer.ts # Training loop orchestrator -├── Tester.ts # Test/evaluation harness -├── model/ -│ └── Datasource.ts # Datasource interface -├── MNIST/ -│ ├── MNIST.ts # MNIST model definition -│ └── MNISTDatasource.ts # MNIST data loading -├── layer/ -│ ├── Layer.ts # Base layer interface -│ ├── Linear.ts # Fully connected layer -│ ├── Sequential.ts # Layer container -│ ├── ReLU.ts # ReLU activation layer -│ └── Dropout.ts # Dropout regularization layer -├── math/ -│ └── Utils.ts # Weight initializers (heNormal, heUniform) -├── autograd/ -│ ├── GradientFunction.ts # Interface for backward functions -│ ├── BackwardPass.ts # Reverse topological sort for backprop -│ └── backward/ -│ ├── MatMulBackward.ts -│ ├── BiasAddBackward.ts -│ ├── ReLUBackward.ts -│ ├── DropoutBackward.ts -│ ├── MatAddBackward.ts -│ ├── SoftmaxBackward.ts -│ └── SoftmaxCrossEntropyBackward.ts -├── optimizer/ -│ ├── Optimizer.ts # Optimizer interface with LR scheduling -│ └── SGD.ts # SGD with learning rate scheduling -└── tensor/ - ├── Tensor.ts # GPU-backed tensor with autograd support - ├── TensorManager.ts # Buffer lifecycle management - └── kernel/ - ├── Kernel.ts # Base kernel class - ├── KernelRegistry.ts # Kernel instantiation - ├── MatMulKernel.ts - ├── MatAddKernel.ts - ├── BiasAddKernel.ts - ├── RELUKernel.ts - ├── DropoutKernel.ts - ├── SoftmaxKernel.ts - ├── CrossEntropyKernel.ts - ├── TransposeKernel.ts # Autograd support - ├── SumReduceKernel.ts # Autograd support - ├── ElementwiseMulKernel.ts # Autograd support - ├── ReLUBackwardKernel.ts # Autograd support - ├── SoftmaxBackwardKernel.ts # Autograd support - ├── SoftmaxCEBackwardKernel.ts # Autograd support - ├── ScalarMulKernel.ts # Optimizer support - ├── InplaceAddKernel.ts # Optimizer support - └── SumAllKernel.ts # Optimizer support -``` - -## Architecture - -WebGPU-based tensor computation library for MNIST, built with TypeScript and Vite. - -### Core Components - -- **GPUEnv** (`src/GPUEnv.ts`): Singleton that initializes WebGPU device. Throws on failure. Must call `await GPUEnv.init()` before using `GPUEnv.device`. - -- **Tensor** (`src/tensor/Tensor.ts`): Wraps `GPUBuffer` with shape metadata. Row-major storage. Size computed via `shape.reduce((a, b) => a * b, 1)`. - -- **TensorManager** (`src/tensor/TensorManager.ts`): Manages GPU buffer lifecycle. Handles buffer reuse, 256-byte alignment, readback buffers, and deferred destruction. Use `getTensorBuffer()` to create/reuse tensors by name. - -- **KernelRegistry** (`src/tensor/kernel/KernelRegistry.ts`): Central registry for GPU kernels. Instantiates all kernels with shared device and TensorManager. - -### Kernel Pattern - -All kernels extend `Kernel` base class (`src/tensor/kernel/Kernel.ts`) which creates the shader module and compute pipeline. - -Kernels follow a functional pattern—they return tensors, enabling composition: -```typescript -const out = kernelRegistry.relu.run( - kernelRegistry.matmul.run(t0, t1) -); -``` - -Each kernel: -1. Validates input tensor shapes (must be 2D) -2. Auto-creates output buffer if not provided (via TensorManager) -3. Writes params to uniform buffer -4. Creates bind group, dispatches compute, submits to queue -5. Returns output tensor - -### Available Kernels - -**Forward Kernels:** -- **MatMulKernel**: Tiled 16×16 matrix multiplication with shared memory -- **MatAddKernel**: Element-wise addition of two 2D tensors with matching shapes -- **BiasAddKernel**: Broadcasts 1D bias [N] across rows of 2D input [M,N] -- **RELUKernel**: Element-wise ReLU activation -- **DropoutKernel**: Element-wise multiplication with pre-computed mask -- **SoftmaxKernel**: Per-row softmax with numerical stability (optimized for small N like MNIST's 10 classes) -- **CrossEntropyKernel**: Per-sample cross-entropy loss. Takes predictions [M,N] and one-hot labels [M,N], outputs loss [M,1] - -**Autograd Support Kernels:** -- **TransposeKernel**: Matrix transpose [M,N] → [N,M] -- **SumReduceKernel**: Sum along axis 0, [M,N] → [1,N] -- **ElementwiseMulKernel**: Hadamard product of two tensors -- **ReLUBackwardKernel**: `dX = dY * (X > 0)` -- **SoftmaxBackwardKernel**: Jacobian-vector product for softmax -- **SoftmaxCEBackwardKernel**: Combined softmax + cross-entropy backward: `dLogits = probs - labels` - -**Optimizer Support Kernels:** -- **ScalarMulKernel**: `output = input * scalar` (for learning rate scaling) -- **InplaceAddKernel**: `target += source` (for parameter updates) -- **SumAllKernel**: Reduces all elements to scalar [1,1] (for loss reduction) - -### WebGPU Compute Pattern - -Shaders use 16×16 workgroup size (except Softmax/CrossEntropy which use 256×1). Dispatch: `(ceil(N/16), ceil(M/16), 1)`. - -MatMul uses tiled algorithm with workgroup-local shared memory and barrier synchronization for coalesced memory access. - -### Adding New Kernels - -1. Create class extending `Kernel` in `src/tensor/kernel/` -2. Define WGSL shader as static string -3. Add shape validation, params buffer, bind group creation -4. Register in `KernelRegistry` - -### Layer System - -All layers implement the `Layer` interface (`src/layer/Layer.ts`): -```typescript -interface Layer { - forward(input: Tensor): Tensor; - backward(input: Tensor): void; - parameters(): Tensor[]; -} -``` - -**Available Layers:** - -- **Linear**: Fully connected layer `output = input * W + b` - - Weight layout: `[inputFeatures, outputFeatures]` (row-major) - - Input: `[batch, inputFeatures]`, Output: `[batch, outputFeatures]` - -- **ReLU**: ReLU activation as a layer wrapper - -- **Dropout**: Regularization layer with inverted dropout scaling - - Mask generated on CPU, applied on GPU - -- **Sequential**: Container for stacking layers - ```typescript - const model = new Sequential( - new Linear(tm, kr, {...}), - new ReLU(tm, kr, "relu"), - new Dropout(tm, kr, "dropout", 0.5), - ); - const out = model.forward(input); - ``` - -### Weight Initialization - -`src/math/Utils.ts` provides: -- **heNormal**: He normal initialization `N(0, sqrt(2/fanIn))` -- **heUniform**: He uniform initialization `U(-limit, limit)` where `limit = sqrt(6/fanIn)` - -### Autograd System - -Automatic differentiation via computation graph tracking. - -**Tensor Autograd Fields** (`src/tensor/Tensor.ts`): -```typescript -gradient?: Tensor; // Accumulated gradient -gradFn?: GradientFunction; // Backward function for this node -parents?: Tensor[]; // Input tensors that created this tensor -requiresGradient: boolean; // Whether to track gradients -``` - -**GradientFunction Interface** (`src/autograd/GradientFunction.ts`): -```typescript -interface GradientFunction { - name: string; - savedTensors: Tensor[]; // Tensors saved for backward - backward(gradOutput: Tensor): Tensor[]; // Compute gradients w.r.t. inputs -} -``` - -**Backpropagation** (`src/autograd/BackwardPass.ts`): -1. Build reverse topological order starting from loss -2. Initialize loss gradient to ones -3. Traverse in reverse order, calling `gradFn.backward()` on each node -4. Accumulate gradients in parent tensors via `inplaceAdd` - -**Backward Functions** (`src/autograd/backward/`): -- **MatMulBackward**: `dA = dC @ Bᵀ`, `dB = Aᵀ @ dC` -- **BiasAddBackward**: `dX = dY`, `dBias = sum(dY, axis=0)` -- **ReLUBackward**: `dX = dY * (X > 0)` -- **DropoutBackward**: `dX = dY * mask * scale` -- **SoftmaxCrossEntropyBackward**: `dLogits = probs - labels` - -### Optimizer System - -**Optimizer Interface** (`src/optimizer/Optimizer.ts`): -```typescript -type LRSchedule = - | { type: "constant" } - | { type: "step"; factor: number; everyNSteps: number } - | { type: "exponential"; decayRate: number } - | { type: "cosine"; minLr: number; maxSteps: number }; - -interface Optimizer { - step(batchSizeOverride?: number): void; // Apply gradients to parameters - zeroGrad(): void; // Reset all gradients to zero - getLearningRate(): number; - setLearningRate(lr: number): void; - setSchedule(schedule: LRSchedule): void; -} -``` - -**SGD Optimizer** (`src/optimizer/SGD.ts`): -```typescript -const optimizer = new SGD(model.parameters(), learningRate, tm, kr, batchSize); -optimizer.setSchedule({ type: "cosine", minLr: 0.001, maxSteps: totalSteps }); -optimizer.zeroGrad(); // Before forward pass -// ... forward, backward ... -optimizer.step(currentBatchSize); // Update: param = param - lr * grad / batchSize -``` - -Features: -- Learning rate scheduling (constant, step decay, exponential decay, cosine annealing) -- Batch size normalization of gradients - -### MNIST Model - -**MNIST Class** (`src/MNIST/MNIST.ts`): -```typescript -class MNIST { - readonly model: Sequential; - constructor(tm: TensorManager, kernelRegistry: KernelRegistry, initializer = heUniform); - async readSnapshot(): Promise; // Load pre-trained weights - async restart(): Promise; // Reinitialize weights -} -``` - -The MNIST model is a 2-layer MLP: `784 → 128 (ReLU) → 10` - -### Training/Testing Architecture - -**Trainer** (`src/Trainer.ts`): -- Manages training loop with `requestAnimationFrame` for non-blocking UI -- Supports epoch-based training with configurable batch size -- Handles LR scheduling via cosine annealing -- State machine: `idle → training → finished` (with cancellation support) - -**Tester** (`src/Tester.ts`): -- Evaluates model accuracy on test set -- Non-blocking batch processing via `requestAnimationFrame` -- Computes accuracy by comparing argmax of predictions vs labels - -### Training Loop Pattern - -```typescript -// Using Trainer class -const trainer = new Trainer(tm, kr, mnist, datasource, onFinished, onUpdate); -await trainer.initialize(); -trainer.startTraining(); - -// Or manual loop -optimizer.zeroGrad(); -const logits = model.forward(input, true); -const loss = kr.crossEntropy.run(logits, labelsOneHot); -computeBackwardPass(tm, kr, loss); // Backpropagation -optimizer.step(batchSize); // Update parameters -``` diff --git a/README.md b/README.md deleted file mode 100644 index 1903a8e..0000000 --- a/README.md +++ /dev/null @@ -1,123 +0,0 @@ -# MNIST - -A neural network trained and evaluated entirely in-browser using WebGPU compute shaders. - -**Live Demo**: [hyperandroid.github.io/MNIST](https://hyperandroid.github.io/MNIST/) - -## Motivation - -This project could easily be done with PyTorch or TensorFlow. I built it from scratch anyway. - -Over the past months, I've watched LLMs catch up to tasks I thought were uniquely human. They've become invaluable for making sense of my own thoughts. An unbelievable ability to put into words what I couldn't articulate myself. Thoughts or intuitions, nor matter what, LLM have a magic way of writting them down into words. - -Turns out, I understood nothing about how they worked. Nor how they trained. Nor how they reasoned. Nor why they failed. I was using tools I couldn't see inside, and that felt asymmetrical. - -In an age where knowledge has an amortized cost of zero, understanding has become the only defensible ground. Information is free. Tutorials are everywhere. But if you don't build the intuition yourself, the tools will be more like magic over time. -And eventually your tools will even replace you. Fortunately, I don't define myself by my coding skills. One of my traits is surely to have a hacker mindset, e.g., understand problems from the ground up. - -This project is my attempt to start from fundamentals. An MLP is an old technology, it dates to the 1950s. But by implementing every tensor operation, every backward pass, every gradient update from scratch, I've built something I -actually understand. It's naive. Unoptimized. Limited to two dimensions. Not production-ready. But it trains, it learns, and I know exactly why. - -This is the foundation. Not the destination. - -Big thanks to [TiniTorch](https://mlsysbook.ai/tinytorch/preface.html). This project is the Typescript+WebGPU -implementation of its [Foundation Tier](https://mlsysbook.ai/tinytorch/tiers/foundation.html). - -## Overview - -This project implements a complete deep learning pipeline for MNIST digit classification that runs entirely in the browser. All tensor operations, training, and inference are performed on the GPU using WebGPU compute shaders. - -### Features - -- **In-Browser Training**: Train the model from scratch on 60,000 MNIST images -- **Real-Time Inference**: Draw digits and see predictions in real-time -- **Layer Visualization**: Watch activations flow through each layer as you draw -- **Parameter Visualization**: View learned weight matrices from the pre-trained model -- **97.5% Accuracy**: Achieves high accuracy on the MNIST test set - -## Architecture - -The model is a Multi-Layer Perceptron (MLP): - -``` -Input [784] → Dense [128] → ReLU → Dense [10] → Softmax -``` - -- **Weight Initialization**: He uniform distribution -- **Optimizer**: Stochastic Gradient Descent (no momentum) -- **Learning Rate Schedule**: Cosine annealing - -## Technical Implementation - -### Tensor System - -- GPU-backed tensors with automatic buffer management -- Row-major storage with 256-byte alignment for WebGPU -- Deferred buffer destruction for efficient memory reuse - -### Compute Kernels - -All operations implemented as WGSL compute shaders: - -| Kernel | Description | -|--------|-------------| -| MatMul | Tiled 16x16 matrix multiplication with shared memory | -| BiasAdd | Broadcasts 1D bias across 2D input | -| ReLU | Element-wise activation | -| Softmax | Per-row softmax with numerical stability | -| CrossEntropy | Per-sample cross-entropy loss | -| Transpose | Matrix transpose for backprop | -| SumReduce | Reduction along axis for gradient accumulation | - -### Autograd System - -Automatic differentiation via computation graph: - -1. Forward pass records operations and parent tensors -2. Backward pass traverses graph in reverse topological order -3. Gradients accumulated via in-place addition - -## Project Structure - -``` -src/ -├── tensor/ # GPU tensor and buffer management -│ └── kernel/ # WGSL compute shaders -├── layer/ # Neural network layers (Linear, ReLU, Dropout, Sequential) -├── autograd/ # Backward pass and gradient functions -├── optimizer/ # SGD with learning rate scheduling -├── MNIST/ # Model definition and data loading -│ └── interactive/ # UI components for the demo -└── GPUEnv.ts # WebGPU device initialization -``` - -## Getting Started - -### Requirements - -- A browser with WebGPU support. This model can be trained in my iPhone 15 pro. -- Node.js 18+ - -### Development - -```bash -npm install -npm run dev -``` - -### Build - -```bash -npm run build -npm run preview # Preview production build locally -``` - -### Deploy - -```bash -npm run deploy # Deploy to GitHub Pages -``` - -## Dataset - -MNIST dataset from: https://git-disl.github.io/GTDLBench/datasets/mnist_datasets/ diff --git a/assets/MNIST-BhAB_iSj.js b/assets/MNIST-BhAB_iSj.js new file mode 100644 index 0000000..fdd18a1 --- /dev/null +++ b/assets/MNIST-BhAB_iSj.js @@ -0,0 +1,522 @@ +(function(){const e=document.createElement("link").relList;if(e&&e.supports&&e.supports("modulepreload"))return;for(const a of document.querySelectorAll('link[rel="modulepreload"]'))s(a);new MutationObserver(a=>{for(const t of a)if(t.type==="childList")for(const i of t.addedNodes)i.tagName==="LINK"&&i.rel==="modulepreload"&&s(i)}).observe(document,{childList:!0,subtree:!0});function r(a){const t={};return a.integrity&&(t.integrity=a.integrity),a.referrerPolicy&&(t.referrerPolicy=a.referrerPolicy),a.crossOrigin==="use-credentials"?t.credentials="include":a.crossOrigin==="anonymous"?t.credentials="omit":t.credentials="same-origin",t}function s(a){if(a.ep)return;a.ep=!0;const t=r(a);fetch(a.href,t)}})();class N{static device;static async init(){if(!navigator.gpu)throw new Error("WebGPU not available in this browser/context.");const e=await navigator.gpu.requestAdapter();if(!e)throw new Error("Failed to get GPU adapter.");N.device=await e.requestDevice()}}function d(c,e){return Math.floor((c+e-1)/e)}class p{constructor(e,r,s){this.kr=s,this.module=e.createShaderModule({code:r}),this.pipeline=e.createComputePipeline({layout:"auto",compute:{module:this.module,entryPoint:"main"}})}pipeline;module}class R{constructor(e,r){this.savedTensors=e,this.kr=r}name="MatMulBackward";backward(e){const[r,s]=this.savedTensors;let a=null;if(r.requiresGradient){const i=this.kr.transpose.run(s);a=this.kr.matmul.run(e,i)}let t=null;if(s.requiresGradient){const i=this.kr.transpose.run(r);t=this.kr.matmul.run(i,e)}return[a,t]}}class l extends p{constructor(e,r,s){super(e,l.matmulWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:16,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(4);paramsBuf;run(e,r,s){if(e.shape.length!==2||r.shape.length!==2||s!==void 0&&s.shape.length!==2)throw new Error("MatMul: expected 2D tensors");if(e.shape[1]!==r.shape[0])throw new Error("MatMul: invalid dimensions");const a=e.shape[0],t=e.shape[1],i=r.shape[1];if(s=s??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[a,i]),s.shape[0]!==a||s.shape[1]!==i)throw new Error("MatMul: invalid output dimensions");this.params[0]=a,this.params[1]=i,this.params[2]=t,this.device.queue.writeBuffer(this.paramsBuf,0,this.params);const n=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:s.buffer}},{binding:3,resource:{buffer:this.paramsBuf}}]}),o=d(i,16),u=d(a,16),f=this.device.createCommandEncoder(),h=f.beginComputePass();return h.setPipeline(this.pipeline),h.setBindGroup(0,n),h.dispatchWorkgroups(o,u,1),h.end(),this.device.queue.submit([f.finish()]),(e.requiresGradient||r.requiresGradient)&&(s.requiresGradient=!0,s.parents=[e,r],s.gradFn=new R([e,r],this.kr)),s}static matmulWGSL=` + // C[M,N] = A[M,K] * B[K,N] (row-major flat buffers) + + struct Params { + M : u32, + N : u32, + K : u32, + _pad : u32, + }; + + @group(0) @binding(0) var A : array; + @group(0) @binding(1) var B : array; + @group(0) @binding(2) var C : array; + @group(0) @binding(3) var params : Params; + + const TILE : u32 = 16u; + + var tileA : array; + var tileB : array; + + @compute @workgroup_size(16, 16, 1) + fn main( + @builtin(global_invocation_id) gid : vec3, + @builtin(local_invocation_id) lid : vec3, + ) { + let row : u32 = gid.y; + let col : u32 = gid.x; + + let inBounds : bool = (row < params.M) && (col < params.N); + + let lidx : u32 = lid.y * TILE + lid.x; + + var acc : f32 = 0.0; + let numTiles : u32 = (params.K + TILE - 1u) / TILE; + + for (var t : u32 = 0u; t < numTiles; t = t + 1u) { + let kBase : u32 = t * TILE; + + // Load A tile element (or 0) + let aCol : u32 = kBase + lid.x; + if ((row < params.M) && (aCol < params.K)) { + tileA[lidx] = A[row * params.K + aCol]; + } else { + tileA[lidx] = 0.0; + } + + // Load B tile element (or 0) + let bRow : u32 = kBase + lid.y; + if ((bRow < params.K) && (col < params.N)) { + tileB[lidx] = B[bRow * params.N + col]; + } else { + tileB[lidx] = 0.0; + } + + workgroupBarrier(); + + for (var i : u32 = 0u; i < TILE; i = i + 1u) { + acc = acc + tileA[lid.y * TILE + i] * tileB[i * TILE + lid.x]; + } + + workgroupBarrier(); + } + + // Only write valid output elements + if (inBounds) { + C[row * params.N + col] = acc; + } + } + `}class z{constructor(e,r){this.savedTensors=e,this.kr=r}name="MatAddBackward";backward(e){return[e,e]}}class b extends p{constructor(e,r,s){super(e,b.matAddWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r,s){if(e.shape.length!==2||r.shape.length!==2||s!==void 0&&s.shape.length!==2)throw new Error("MatAdd: expected 2D tensors");if(e.shape[0]!==r.shape[0]||e.shape[1]!==r.shape[1])throw new Error("MatAdd: tensor shapes must match");const a=e.shape[0],t=e.shape[1];this.params[0]=a,this.params[1]=t,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),s=s??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[a,t]);const i=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:s.buffer}},{binding:3,resource:{buffer:this.paramsBuf}}]}),n=d(t,16),o=d(a,16),u=this.device.createCommandEncoder(),f=u.beginComputePass();return f.setPipeline(this.pipeline),f.setBindGroup(0,i),f.dispatchWorkgroups(n,o,1),f.end(),this.device.queue.submit([u.finish()]),(e.requiresGradient||r.requiresGradient)&&(s.requiresGradient=!0,s.parents=[e,r],s.gradFn=new z([e,r],this.kr)),s}static matAddWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var A : array; + @group(0) @binding(1) var B : array; + @group(0) @binding(2) var C : array; + @group(0) @binding(3) var params : Params; + + @compute @workgroup_size(16, 16, 1) + fn main( + @builtin(global_invocation_id) gid : vec3, + ) { + let row : u32 = gid.y; + let col : u32 = gid.x; + + if (row < params.M && col < params.N) { + let idx : u32 = row * params.N + col; + C[idx] = A[idx] + B[idx]; + } + } + `}class I{constructor(e,r){this.savedTensors=e,this.kr=r}name="BiasAddBackward";backward(e){const[r,s]=this.savedTensors,a=r.requiresGradient?e:null,t=s.requiresGradient?this.kr.sumReduce.run(e):null;return[a,t]}}class w extends p{constructor(e,r,s){super(e,w.biasAddWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r,s){if(e.shape.length!==2)throw new Error("BiasAdd: input must be 2D tensor");const a=r.shape.length===1?r.shape[0]:r.shape.length===2&&r.shape[0]===1?r.shape[1]:-1;if(a===-1)throw new Error("BiasAdd: bias must be [N] or [1, N]");if(e.shape[1]!==a)throw new Error(`BiasAdd: input columns (${e.shape[1]}) must match bias size (${a})`);const t=e.shape[0],i=e.shape[1];this.params[0]=t,this.params[1]=i,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),s=s??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[t,i]);const n=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:s.buffer}},{binding:3,resource:{buffer:this.paramsBuf}}]}),o=d(i,16),u=d(t,16),f=this.device.createCommandEncoder(),h=f.beginComputePass();return h.setPipeline(this.pipeline),h.setBindGroup(0,n),h.dispatchWorkgroups(o,u,1),h.end(),this.device.queue.submit([f.finish()]),(e.requiresGradient||r.requiresGradient)&&(s.requiresGradient=!0,s.parents=[e,r],s.gradFn=new I([e,r],this.kr)),s}static biasAddWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var input : array; + @group(0) @binding(1) var bias : array; + @group(0) @binding(2) var output : array; + @group(0) @binding(3) var params : Params; + + @compute @workgroup_size(16, 16, 1) + fn main( + @builtin(global_invocation_id) gid : vec3, + ) { + let row : u32 = gid.y; + let col : u32 = gid.x; + + if (row < params.M && col < params.N) { + let idx : u32 = row * params.N + col; + output[idx] = input[idx] + bias[col]; + } + } + `}class F{constructor(e,r){this.savedTensors=e,this.kr=r}name="ReLUBackward";backward(e){const[r]=this.savedTensors;return[this.kr.reluBackward.run(e,r)]}}class B extends p{constructor(e,r,s){super(e,B.reluWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}static RELU_OUTPUT="relu_out";params=new Uint32Array(2);paramsBuf;run(e,r){if(e.shape.length!==2||r!==void 0&&r.shape.length!==2)throw new Error("RELU: expected 2D tensor");const s=e.shape[0],a=e.shape[1];this.params[0]=s,this.params[1]=a,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),r=r??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[s,a]);const t=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:this.paramsBuf}}]}),i=d(a,16),n=d(s,16),o=this.device.createCommandEncoder(),u=o.beginComputePass();return u.setPipeline(this.pipeline),u.setBindGroup(0,t),u.dispatchWorkgroups(i,n,1),u.end(),this.device.queue.submit([o.finish()]),e.requiresGradient&&(r.requiresGradient=!0,r.parents=[e],r.gradFn=new F([e],this.kr)),r}static reluWGSL=` + // RELU: remove negative values from a tensor. + + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var A : array; + @group(0) @binding(1) var B : array; + @group(0) @binding(2) var params : Params; + + @compute @workgroup_size(16, 16, 1) + fn main( + @builtin(global_invocation_id) gid : vec3, + ) { + let row : u32 = gid.y; + let col : u32 = gid.x; + + let inBounds : bool = (row < params.M) && (col < params.N); + + // Only write valid output elements + if (inBounds) { + let f: f32 = A[row * params.N + col]; + B[row * params.N + col] = max(0f, f); + } + } + `}class Y{constructor(e,r){this.savedTensors=e,this.kr=r}name="SoftmaxBackward";backward(e){const r=this.savedTensors[1];return[this.kr.softmaxBackward.run(e,r)]}}class U extends p{constructor(e,r,s){super(e,U.softmaxWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r){if(e.shape.length!==2)throw new Error("Softmax: expected 2D tensor");const s=e.shape[0],a=e.shape[1];this.params[0]=s,this.params[1]=a,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),r=r??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[s,a]);const t=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:this.paramsBuf}}]}),i=this.device.createCommandEncoder(),n=i.beginComputePass();return n.setPipeline(this.pipeline),n.setBindGroup(0,t),n.dispatchWorkgroups(d(s,256),1,1),n.end(),this.device.queue.submit([i.finish()]),e.requiresGradient&&(r.requiresGradient=!0,r.parents=[e],r.gradFn=new Y([e,r],this.kr)),r}static softmaxWGSL=` + // Softmax: exp(x_i - max) / sum(exp(x_j - max)) + // Applied per row for numerical stability. + + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var A : array; + @group(0) @binding(1) var B : array; + @group(0) @binding(2) var params : Params; + + @compute @workgroup_size(256, 1, 1) + fn main(@builtin(global_invocation_id) gid : vec3) { + let row = gid.x; + if (row >= params.M) { + return; + } + + let N = params.N; + let base = row * N; + + // Find max for numerical stability + var maxVal = A[base]; + for (var i = 1u; i < N; i = i + 1u) { + maxVal = max(maxVal, A[base + i]); + } + + // Compute exp(x - max) and sum + var sum = 0.0; + for (var i = 0u; i < N; i = i + 1u) { + let e = exp(A[base + i] - maxVal); + B[base + i] = e; + sum = sum + e; + } + + // Normalize + let invSum = 1.0 / sum; + for (var i = 0u; i < N; i = i + 1u) { + B[base + i] = B[base + i] * invSum; + } + } + `}class q{constructor(e,r){this.savedTensors=e,this.kr=r}name="SoftmaxCrossEntropyBackward";backward(e){const[r,s]=this.savedTensors;return[this.kr.softmaxCEBackward.run(r,s)]}}class G extends p{constructor(e,r,s){super(e,G.xentropyWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r,s){if(e.shape.length!==2)throw new Error("CrossEntropy: logits must be 2D tensor");if(r.shape.length!==2)throw new Error("CrossEntropy: labels must be 2D tensor");if(e.shape[0]!==r.shape[0]||e.shape[1]!==r.shape[1])throw new Error("CrossEntropy: logits and labels must have same shape");const a=e.shape[0],t=e.shape[1];this.params[0]=a,this.params[1]=t,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),s=s??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[a,1]);const i=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:s.buffer}},{binding:3,resource:{buffer:this.paramsBuf}}]}),n=this.device.createCommandEncoder(),o=n.beginComputePass();return o.setPipeline(this.pipeline),o.setBindGroup(0,i),o.dispatchWorkgroups(d(a,256),1,1),o.end(),this.device.queue.submit([n.finish()]),e.requiresGradient&&(s.requiresGradient=!0,s.parents=[e],s.gradFn=new q([e,r],this.kr)),s}static xentropyWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var logits : array; + @group(0) @binding(1) var labels : array; // one-hot + @group(0) @binding(2) var loss : array; + @group(0) @binding(3) var params : Params; + + @compute @workgroup_size(256, 1, 1) + fn main(@builtin(global_invocation_id) gid : vec3) { + let row = gid.x; + if (row >= params.M) { return; } + + let N = params.N; + let base = row * N; + + // 1) max logit for stability + var m = logits[base]; + for (var i = 1u; i < N; i = i + 1u) { + let z = logits[base + i]; + if (z > m) { m = z; } + } + + // 2) logsumexp + var sumExp = 0.0; + for (var i = 0u; i < N; i = i + 1u) { + sumExp = sumExp + exp(logits[base + i] - m); + } + let logSumExp = log(sumExp) + m; + + // 3) cross entropy: -sum y_i * (z_i - logsumexp) + var ce = 0.0; + for (var i = 0u; i < N; i = i + 1u) { + let y = labels[base + i]; + let z = logits[base + i]; + ce = ce + y * (logSumExp - z); + } + + loss[row] = ce; + } + + `}class W{constructor(e,r){this.savedTensors=e,this.kr=r}name="DropoutBackward";backward(e){const[r]=this.savedTensors;return[this.kr.elemMul.run(e,r)]}}class P extends p{constructor(e,r,s){super(e,P.dropoutWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r,s){if(e.shape.length!==2)throw new Error("Dropout: input must be 2D tensor");if(r.shape.length!==2)throw new Error("Dropout: mask must be 2D tensor");if(e.shape[0]!==r.shape[0]||e.shape[1]!==r.shape[1])throw new Error("Dropout: input and mask shapes must match");const a=e.shape[0],t=e.shape[1];this.params[0]=a,this.params[1]=t,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),s=s??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[a,t]);const i=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:s.buffer}},{binding:3,resource:{buffer:this.paramsBuf}}]}),n=d(t,16),o=d(a,16),u=this.device.createCommandEncoder(),f=u.beginComputePass();return f.setPipeline(this.pipeline),f.setBindGroup(0,i),f.dispatchWorkgroups(n,o,1),f.end(),this.device.queue.submit([u.finish()]),e.requiresGradient&&(s.requiresGradient=!0,s.parents=[e],s.gradFn=new W([r],this.kr)),s}static dropoutWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var input : array; + @group(0) @binding(1) var mask : array; + @group(0) @binding(2) var output : array; + @group(0) @binding(3) var params : Params; + + @compute @workgroup_size(16, 16, 1) + fn main( + @builtin(global_invocation_id) gid : vec3, + ) { + let row : u32 = gid.y; + let col : u32 = gid.x; + + if (row < params.M && col < params.N) { + let idx : u32 = row * params.N + col; + output[idx] = input[idx] * mask[idx]; + } + } + `}class y extends p{constructor(e,r,s){super(e,y.transposeWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r){if(e.shape.length!==2)throw new Error("Transpose: input must be 2D tensor");const s=e.shape[0],a=e.shape[1];if(this.params[0]=s,this.params[1]=a,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),r=r??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[a,s]),r.shape[0]!==a||r.shape[1]!==s)throw new Error("Transpose: output shape must be [N, M]");const t=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:this.paramsBuf}}]}),i=d(a,16),n=d(s,16),o=this.device.createCommandEncoder(),u=o.beginComputePass();return u.setPipeline(this.pipeline),u.setBindGroup(0,t),u.dispatchWorkgroups(i,n,1),u.end(),this.device.queue.submit([o.finish()]),r}static transposeWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var input : array; + @group(0) @binding(1) var output : array; + @group(0) @binding(2) var params : Params; + + @compute @workgroup_size(16, 16, 1) + fn main(@builtin(global_invocation_id) gid : vec3) { + let col : u32 = gid.x; // output column = input row + let row : u32 = gid.y; // output row = input column + + if (row >= params.M || col >= params.N) { + return; + } + + // input[row, col] -> output[col, row] + let inIdx = row * params.N + col; + let outIdx = col * params.M + row; + output[outIdx] = input[inIdx]; + } + `}class S extends p{constructor(e,r,s){super(e,S.sumReduceWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r){if(e.shape.length!==2)throw new Error("SumReduce: input must be 2D tensor");const s=e.shape[0],a=e.shape[1];this.params[0]=s,this.params[1]=a,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),r=r??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[1,a]);const t=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:this.paramsBuf}}]}),i=this.device.createCommandEncoder(),n=i.beginComputePass();return n.setPipeline(this.pipeline),n.setBindGroup(0,t),n.dispatchWorkgroups(d(a,256),1,1),n.end(),this.device.queue.submit([i.finish()]),r}static sumReduceWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var input : array; + @group(0) @binding(1) var output : array; + @group(0) @binding(2) var params : Params; + + @compute @workgroup_size(256, 1, 1) + fn main(@builtin(global_invocation_id) gid : vec3) { + let col = gid.x; + if (col >= params.N) { + return; + } + + var sum = 0.0; + for (var row = 0u; row < params.M; row = row + 1u) { + sum = sum + input[row * params.N + col]; + } + + output[col] = sum; + } + `}class v extends p{constructor(e,r,s){super(e,v.elemMulWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r,s){if(e.shape.length!==2||r.shape.length!==2)throw new Error("ElementwiseMul: inputs must be 2D tensors");if(e.shape[0]!==r.shape[0]||e.shape[1]!==r.shape[1])throw new Error("ElementwiseMul: input shapes must match");const a=e.shape[0],t=e.shape[1];this.params[0]=a,this.params[1]=t,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),s=s??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[a,t]);const i=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:s.buffer}},{binding:3,resource:{buffer:this.paramsBuf}}]}),n=d(t,16),o=d(a,16),u=this.device.createCommandEncoder(),f=u.beginComputePass();return f.setPipeline(this.pipeline),f.setBindGroup(0,i),f.dispatchWorkgroups(n,o,1),f.end(),this.device.queue.submit([u.finish()]),s}static elemMulWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var A : array; + @group(0) @binding(1) var B : array; + @group(0) @binding(2) var C : array; + @group(0) @binding(3) var params : Params; + + @compute @workgroup_size(16, 16, 1) + fn main(@builtin(global_invocation_id) gid : vec3) { + let col = gid.x; + let row = gid.y; + + if (row >= params.M || col >= params.N) { + return; + } + + let idx = row * params.N + col; + C[idx] = A[idx] * B[idx]; + } + `}class x extends p{constructor(e,r,s){super(e,x.reluBackwardWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r,s){if(e.shape.length!==2||r.shape.length!==2)throw new Error("ReLUBackward: inputs must be 2D tensors");if(e.shape[0]!==r.shape[0]||e.shape[1]!==r.shape[1])throw new Error("ReLUBackward: shapes must match");const a=e.shape[0],t=e.shape[1];this.params[0]=a,this.params[1]=t,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),s=s??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[a,t]);const i=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:s.buffer}},{binding:3,resource:{buffer:this.paramsBuf}}]}),n=d(t,16),o=d(a,16),u=this.device.createCommandEncoder(),f=u.beginComputePass();return f.setPipeline(this.pipeline),f.setBindGroup(0,i),f.dispatchWorkgroups(n,o,1),f.end(),this.device.queue.submit([u.finish()]),s}static reluBackwardWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var gradOutput : array; + @group(0) @binding(1) var savedInput : array; + @group(0) @binding(2) var gradInput : array; + @group(0) @binding(3) var params : Params; + + @compute @workgroup_size(16, 16, 1) + fn main(@builtin(global_invocation_id) gid : vec3) { + let col = gid.x; + let row = gid.y; + + if (row >= params.M || col >= params.N) { + return; + } + + let idx = row * params.N + col; + // gradient flows through only where input > 0 + let mask = select(0.0, 1.0, savedInput[idx] > 0.0); + gradInput[idx] = gradOutput[idx] * mask; + } + `}class C extends p{constructor(e,r,s){super(e,C.softmaxCEBackwardWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r,s){if(e.shape.length!==2||r.shape.length!==2)throw new Error("SoftmaxCEBackward: inputs must be 2D tensors");if(e.shape[0]!==r.shape[0]||e.shape[1]!==r.shape[1])throw new Error("SoftmaxCEBackward: shapes must match");const a=e.shape[0],t=e.shape[1];this.params[0]=a,this.params[1]=t,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),s=s??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[a,t]);const i=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:s.buffer}},{binding:3,resource:{buffer:this.paramsBuf}}]}),n=this.device.createCommandEncoder(),o=n.beginComputePass();return o.setPipeline(this.pipeline),o.setBindGroup(0,i),o.dispatchWorkgroups(d(a,256),1,1),o.end(),this.device.queue.submit([n.finish()]),s}static softmaxCEBackwardWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var logits : array; + @group(0) @binding(1) var labels : array; + @group(0) @binding(2) var gradLogits : array; + @group(0) @binding(3) var params : Params; + + @compute @workgroup_size(256, 1, 1) + fn main(@builtin(global_invocation_id) gid : vec3) { + let row = gid.x; + if (row >= params.M) { + return; + } + + let N = params.N; + let base = row * N; + + // 1) Find max for numerical stability + var maxVal = logits[base]; + for (var i = 1u; i < N; i = i + 1u) { + maxVal = max(maxVal, logits[base + i]); + } + + // 2) Compute exp(x - max) and sum + var sum = 0.0; + for (var i = 0u; i < N; i = i + 1u) { + let e = exp(logits[base + i] - maxVal); + gradLogits[base + i] = e; + sum = sum + e; + } + + // 3) Normalize to get softmax, then subtract labels + let invSum = 1.0 / sum; + for (var i = 0u; i < N; i = i + 1u) { + let softmax_i = gradLogits[base + i] * invSum; + gradLogits[base + i] = softmax_i - labels[base + i]; + } + } + `}class k extends p{constructor(e,r,s){super(e,k.softmaxBackwardWGSL,s),this.device=e,this.tm=r,this.kr=s,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r,s){if(e.shape.length!==2||r.shape.length!==2)throw new Error("SoftmaxBackward: inputs must be 2D tensors");if(e.shape[0]!==r.shape[0]||e.shape[1]!==r.shape[1])throw new Error("SoftmaxBackward: shapes must match");const a=e.shape[0],t=e.shape[1];this.params[0]=a,this.params[1]=t,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),s=s??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[a,t]);const i=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:s.buffer}},{binding:3,resource:{buffer:this.paramsBuf}}]}),n=this.device.createCommandEncoder(),o=n.beginComputePass();return o.setPipeline(this.pipeline),o.setBindGroup(0,i),o.dispatchWorkgroups(d(a,256),1,1),o.end(),this.device.queue.submit([n.finish()]),s}static softmaxBackwardWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var gradOutput : array; + @group(0) @binding(1) var softmaxOut : array; + @group(0) @binding(2) var gradInput : array; + @group(0) @binding(3) var params : Params; + + @compute @workgroup_size(256, 1, 1) + fn main(@builtin(global_invocation_id) gid : vec3) { + let row = gid.x; + if (row >= params.M) { + return; + } + + let N = params.N; + let base = row * N; + + // Compute dot = sum(dS * S) for this row + var dot = 0.0; + for (var i = 0u; i < N; i = i + 1u) { + dot = dot + gradOutput[base + i] * softmaxOut[base + i]; + } + + // Compute gradInput = S * (dS - dot) + for (var i = 0u; i < N; i = i + 1u) { + let idx = base + i; + gradInput[idx] = softmaxOut[idx] * (gradOutput[idx] - dot); + } + } + `}class T extends p{constructor(e,r,s){super(e,T.scalarMulWGSL,s),this.device=e,this.tm=r,this.paramsBuf=e.createBuffer({size:16,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Float32Array(4);paramsBuf;run(e,r,s){if(e.shape.length!==2)throw new Error("ScalarMul: input must be 2D tensor");const a=e.shape[0],t=e.shape[1],i=new Uint32Array(this.params.buffer);i[0]=a,i[1]=t,this.params[2]=r,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),s=s??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[a,t]);const n=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:s.buffer}},{binding:2,resource:{buffer:this.paramsBuf}}]}),o=d(t,16),u=d(a,16),f=this.device.createCommandEncoder(),h=f.beginComputePass();return h.setPipeline(this.pipeline),h.setBindGroup(0,n),h.dispatchWorkgroups(o,u,1),h.end(),this.device.queue.submit([f.finish()]),s}static scalarMulWGSL=` + struct Params { + M : u32, + N : u32, + scalar : f32, + _pad : u32, + }; + + @group(0) @binding(0) var input : array; + @group(0) @binding(1) var output : array; + @group(0) @binding(2) var params : Params; + + @compute @workgroup_size(16, 16, 1) + fn main(@builtin(global_invocation_id) gid : vec3) { + let col = gid.x; + let row = gid.y; + + if (row >= params.M || col >= params.N) { + return; + } + + let idx = row * params.N + col; + output[idx] = input[idx] * params.scalar; + } + `}class _ extends p{constructor(e,r,s){super(e,_.inplaceAddWGSL,s),this.device=e,this.tm=r,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r){if(e.shape.length!==2||r.shape.length!==2)throw new Error("InplaceAdd: inputs must be 2D tensors");if(e.shape[0]!==r.shape[0]||e.shape[1]!==r.shape[1])throw new Error("InplaceAdd: tensor shapes must match");const s=e.shape[0],a=e.shape[1];this.params[0]=s,this.params[1]=a,this.device.queue.writeBuffer(this.paramsBuf,0,this.params);const t=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:this.paramsBuf}}]}),i=d(a,16),n=d(s,16),o=this.device.createCommandEncoder(),u=o.beginComputePass();u.setPipeline(this.pipeline),u.setBindGroup(0,t),u.dispatchWorkgroups(i,n,1),u.end(),this.device.queue.submit([o.finish()])}static inplaceAddWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var targetTensor : array; + @group(0) @binding(1) var source : array; + @group(0) @binding(2) var params : Params; + + @compute @workgroup_size(16, 16, 1) + fn main(@builtin(global_invocation_id) gid : vec3) { + let col = gid.x; + let row = gid.y; + + if (row >= params.M || col >= params.N) { + return; + } + + let idx = row * params.N + col; + targetTensor[idx] = targetTensor[idx] + source[idx]; + } + `}class D extends p{constructor(e,r,s){super(e,D.sumAllWGSL,s),this.device=e,this.tm=r,this.paramsBuf=e.createBuffer({size:8,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST})}params=new Uint32Array(2);paramsBuf;run(e,r){if(e.shape.length!==2)throw new Error("SumAll: input must be 2D tensor");const s=e.shape[0],a=e.shape[1];s*a>65536&&console.warn("SumAll: tensor size > 65536, consider using hierarchical reduction"),this.params[0]=s,this.params[1]=a,this.device.queue.writeBuffer(this.paramsBuf,0,this.params),r=r??this.tm.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[1,1]);const i=this.device.createBindGroup({layout:this.pipeline.getBindGroupLayout(0),entries:[{binding:0,resource:{buffer:e.buffer}},{binding:1,resource:{buffer:r.buffer}},{binding:2,resource:{buffer:this.paramsBuf}}]}),n=this.device.createCommandEncoder(),o=n.beginComputePass();return o.setPipeline(this.pipeline),o.setBindGroup(0,i),o.dispatchWorkgroups(1,1,1),o.end(),this.device.queue.submit([n.finish()]),r}static sumAllWGSL=` + struct Params { + M : u32, + N : u32, + }; + + @group(0) @binding(0) var input : array; + @group(0) @binding(1) var output : array; + @group(0) @binding(2) var params : Params; + + var sharedData : array; + + @compute @workgroup_size(256, 1, 1) + fn main( + @builtin(local_invocation_id) lid : vec3, + ) { + let totalSize = params.M * params.N; + let tid = lid.x; + + // Each thread sums a strided portion of the input + var sum = 0.0; + var idx = tid; + while (idx < totalSize) { + sum = sum + input[idx]; + idx = idx + 256u; + } + + sharedData[tid] = sum; + workgroupBarrier(); + + // Parallel reduction in sharedData memory + if (tid < 128u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 128u]; } + workgroupBarrier(); + if (tid < 64u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 64u]; } + workgroupBarrier(); + if (tid < 32u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 32u]; } + workgroupBarrier(); + if (tid < 16u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 16u]; } + workgroupBarrier(); + if (tid < 8u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 8u]; } + workgroupBarrier(); + if (tid < 4u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 4u]; } + workgroupBarrier(); + if (tid < 2u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 2u]; } + workgroupBarrier(); + if (tid < 1u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 1u]; } + workgroupBarrier(); + + // Thread 0 writes the result + if (tid == 0u) { + output[0] = sharedData[0]; + } + } + `}class X{matmul;matadd;biasadd;relu;softmax;crossEntropy;dropout;transpose;sumReduce;elemMul;reluBackward;softmaxBackward;softmaxCEBackward;scalarMul;inplaceAdd;sumAll;constructor(e,r){this.matmul=new l(e,r,this),this.matadd=new b(e,r,this),this.biasadd=new w(e,r,this),this.relu=new B(e,r,this),this.softmax=new U(e,r,this),this.crossEntropy=new G(e,r,this),this.dropout=new P(e,r,this),this.transpose=new y(e,r,this),this.sumReduce=new S(e,r,this),this.elemMul=new v(e,r,this),this.reluBackward=new x(e,r,this),this.softmaxBackward=new k(e,r,this),this.softmaxCEBackward=new C(e,r,this),this.scalarMul=new T(e,r,this),this.inplaceAdd=new _(e,r,this),this.sumAll=new D(e,r,this)}}class L{constructor(e,r,s,a,t=!1){this.name=e,this.buffer=r,this.usage=s,this.shape=a,this.size=a.reduce((i,n)=>i*n,1),this.requiresGradient=t}size;gradient;gradFn=void 0;parents=void 0;requiresGradient;sizeInBytes(){return this.size*4}backward(){if(!this.gradient)throw new Error("Tensor has no gradient")}zeroGrad(e){this.gradient&&e.zeros(this.gradient)}}function E(c,e){return Math.ceil(c/e)*e}class V{constructor(e){this.device=e}tensors=new Map;readback=null;pendingDestroy=[];scopeName="";scopeCounter=0;beginScope(e){this.scopeName=e,this.scopeCounter=0}getScopedTensor(e,r,s){const a=`_${this.scopeName}_${this.scopeCounter++}`;return this.getTensorBuffer(a,e,r,s)}getTensorBuffer(e,r,s,a=void 0){const t=this.tensors.get(e),i=s.reduce((u,f)=>u*f,1)*4;if(t&&t.sizeInBytes()>=i&&t.usage===r){if(a&&this.writeBufferF32(t.buffer,a),t.shape.length===s.length&&t.shape.every((h,O)=>h===s[O]))return t;const f=new L(e,t.buffer,r,s);return this.tensors.set(e,f),f}const n=this.device.createBuffer({size:E(i,256),usage:r});t&&this.pendingDestroy.push(t.buffer);const o=new L(e,n,r,s);return this.tensors.set(e,o),a&&this.writeBufferF32(n,a),o}writeBufferF32(e,r,s=0){this.device.queue.writeBuffer(e,s,r.buffer,r.byteOffset,r.byteLength)}async readBuffer(e,r,s=0){const a=this.ensureReadback(r),t=this.device.createCommandEncoder();t.copyBufferToBuffer(e,s,a,0,r),this.device.queue.submit([t.finish()]),await a.mapAsync(GPUMapMode.READ,0,r);const n=a.getMappedRange(0,r).slice(0);return a.unmap(),new Float32Array(n)}ensureReadback(e){if(e=E(e,256),this.readback!==null&&this.readback.size>=e)return this.readback.buffer;this.readback&&(this.pendingDestroy.push(this.readback.buffer),this.readback=null);const r=this.device.createBuffer({size:e,usage:GPUBufferUsage.COPY_DST|GPUBufferUsage.MAP_READ});return this.readback={buffer:r,size:e},r}ones(e,r="ones"){const s=this.getTensorBuffer(r,GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,e);return this.writeBufferF32(s.buffer,new Float32Array(e.reduce((a,t)=>a*t,1)).fill(1)),s}scopedOnes(e){const r=this.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,e);return this.writeBufferF32(r.buffer,new Float32Array(e.reduce((s,a)=>s*a,1)).fill(1)),r}scopedZeros(e){const r=this.getScopedTensor(GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,e);return this.writeBufferF32(r.buffer,new Float32Array(e.reduce((s,a)=>s*a,1)).fill(0)),r}zeros(e){this.writeBufferF32(e.buffer,new Float32Array(e.size).fill(0))}async flushDestroyQueue(){if(this.pendingDestroy.length!==0){await this.device.queue.onSubmittedWorkDone();for(const e of this.pendingDestroy)try{e.destroy()}catch(r){console.error("Failed to destroy tensor buffer "+r)}this.pendingDestroy.length=0}}async destroyAll(){await this.device.queue.onSubmittedWorkDone();for(const{buffer:e}of this.tensors.values())try{e.destroy()}catch(r){console.error("Failed to destroy tensor buffer "+r)}if(this.tensors.clear(),this.readback){try{this.readback.buffer.destroy()}catch(e){console.error("Failed to destroy readback buffer "+e)}this.readback=null}for(const e of this.pendingDestroy)try{e.destroy()}catch(r){console.error("Failed to destroy scheduled to destroy tensor buffer "+r)}this.pendingDestroy.length=0}}class M{constructor(e,r,s,a,t,i){this.batchSize=e,this.imageData=r,this.oneImageSize=s,this.labelsData=a,this.oneLabelSize=t,this.maxSize=i,this.workingImageBuffer=new Float32Array(s*this.batchSize),this.workingLabelBuffer=new Float32Array(t*this.batchSize);const n=Math.min(this.maxSize,this.imageData.length/s),o=Math.min(this.maxSize,this.labelsData.length/t);if(n!==o)throw new Error("MNIST data source: image and label data have different sizes");this.iteratorSize=o,this.restart()}currentIndex=0;iteratorSize=0;datasetIndices=new Uint32Array(0);workingImageBuffer;workingLabelBuffer;getCurrentIndex(){return this.currentIndex}getSize(){return this.iteratorSize}hasNext(){return this.currentIndex0;s--){const a=Math.floor(Math.random()*(s+1));[r[s],r[a]]=[r[a],r[s]]}return r}}class g{trainData=null;trainLabelsData=null;testData=null;testLabelsData=null;static imageSize=784;testImagesCount=1e4;trainImagesCount=6e4;maxTrainSize=1/0;maxTestSize=1/0;constructor(){}toFloat32(e){const r=new Float32Array(e.length);for(let s=0;se.parameters())}zeroGrad(e){for(const r of this.parameters())r.gradient&&e.writeBufferF32(r.gradient.buffer,new Float32Array(r.size).fill(0))}}class A{constructor(e,r,s){this.tm=e,this.kr=r;const a=[s.inputFeatures,s.outputFeatures];this.name=s.name,this.weights=e.getTensorBuffer(`${s.name}_weights`,GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,a,s.initializer(a,s.inputFeatures)),this.weights.requiresGradient=!0,s.useBias&&(this.bias=e.getTensorBuffer(`${s.name}_bias`,GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[1,s.outputFeatures]),this.bias.requiresGradient=!0)}inputTensor;name;weights;bias;forward(e,r){this.inputTensor=e;const s=this.tm.getTensorBuffer(`${this.name}_mmout`,GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[e.shape[0],this.weights.shape[1]]),a=this.kr.matmul.run(e,this.weights,s);if(!this.bias)return a;const t=this.tm.getTensorBuffer(`${this.name}_sumout`,GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[e.shape[0],this.weights.shape[1]]);return this.kr.biasadd.run(s,this.bias,t)}parameters(){const e=[this.weights];return this.bias!==void 0&&e.push(this.bias),e}}function m(c,e,r){const s=Math.sqrt(6/e),a=c.reduce((t,i)=>t*i,1);r=r??new Float32Array(a);for(let t=0;t{this.capturing=!0}),this.canvas.addEventListener("mouseup",n=>{this.capturing=!1}),this.canvas.addEventListener("mousemove",n=>{this.capturing&&this.setModel(n.offsetX,n.offsetY,n.altKey)}),this.canvas.addEventListener("click",n=>{this.setModel(n.offsetX,n.offsetY,n.altKey)}),this.canvas.addEventListener("touchstart",n=>{n.preventDefault(),this.capturing=!0;const r=n.touches[0],d=this.canvas.getBoundingClientRect();this.setModel(r.clientX-d.left,r.clientY-d.top,!1)}),this.canvas.addEventListener("touchend",n=>{n.preventDefault(),this.capturing=!1}),this.canvas.addEventListener("touchmove",n=>{if(n.preventDefault(),!this.capturing)return;const r=n.touches[0],d=this.canvas.getBoundingClientRect();this.setModel(r.clientX-d.left,r.clientY-d.top,!1)}),this.root.appendChild(document.createTextNode(`${this.rows}x${this.columns}, padded to 28x28`)),this.root.appendChild(document.createElement("br")),this.root.appendChild(document.createElement("br")),this.root.appendChild(this.canvas),this.root.appendChild(document.createElement("br"));const o=document.createElement("button");o.textContent="Clear",o.addEventListener("click",()=>{this.model.fill(0),this.render(this.ctx,this.rows,this.columns,this.model),this.onModelCleared()}),this.root.appendChild(o),a&&(this.canvasOut=document.createElement("canvas"),this.canvasOut.width=28*h,this.canvasOut.height=28*h,this.ctxOut=this.canvasOut.getContext("2d"),this.root.appendChild(document.createElement("br")),this.root.appendChild(this.canvasOut),this.render(this.ctxOut,28,28,this.outmodel))}root;canvas;ctx;canvasOut;ctxOut;model;outmodel;capturing=!1;incColor(t,s,e=1){s>0&&this.increment(t,s-1,.33*e),s0&&this.increment(t-1,s,.33*e),t0&&(os&&(s=o),ni&&(i=n));const a=Math.floor((28-(i-e))/2),c=Math.floor((28-(s-t))/2);for(let o=0;o<=s-t;o++)for(let n=0;n<=i-e;n++)this.outmodel[(o+c)*28+n+a]=this.model[(o+t)*this.columns+n+e]}render(t,s,e,i){t.fillStyle="#111",t.fillRect(0,0,t.canvas.width,t.canvas.height),t.strokeStyle="white";for(let a=0;a{this.drawLayer(e,10,10+i*65,s[i],i>0)})}drawLayer(t,s,e,i,a=!1){let c=-1/0,o=1/0;for(let r=0;rc&&(c=t[r]),t[r]{throw new Error("Failed to load data source: "+l)});const P=new Float32Array(10);let y=!1;new S(async l=>{await p.device.queue.onSubmittedWorkDone(),!y&&(y=!0,await R(l),y=!1)},async()=>{b(P),C.render([],[])});const C=new L;(function(){const l=document.getElementById("row"),t=document.createElement("span");t.style.display="inline-flex",t.style.flexDirection="column",t.style.gap="6px";let s=0;Array.from({length:10},()=>{const e=document.createElement("span");e.classList.add("probability"),e.id=`probability-${s}`;const i=document.createElement("input");i.style.pointerEvents="none",i.type="range",i.min="0",i.max="100",i.value="0",i.id=`probability-slider-${s}`,e.appendChild(i),e.appendChild(document.createTextNode(`${s}`));const a=document.createElement("span");return a.id=`probability-percentage-${s}`,e.appendChild(a),t.appendChild(e),s++,i}),l.appendChild(t)})();async function R(l){const t=m.getTensorBuffer("test_input",GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[1,l.length],l);m.beginScope("test");const s=f.model.forward(t,!1),e=x.softmax.run(s),i=await m.readBuffer(e.buffer,e.sizeInBytes());b(i);const a=[];for(const o of f.model.layers){const n=o.inputTensor,r=await m.readBuffer(n.buffer,n.sizeInBytes());a.push(r)}a.push(i);const c=f.model.layers.map(o=>`${o.name} [${o.inputTensor?.shape}]`);c.push("Predictions. Shape [1,10]. Digit 0..9"),C.render(a,c)}function b(l){let t=-1/0,s=0;for(let e=0;e<10;e++){const i=l[e];i>t&&(t=i,s=e);const a=document.getElementById(`probability-${e}`);a!==null&&(a.style.backgroundColor="default");const c=document.getElementById(`probability-slider-${e}`);c!==null&&(c.value=(100*i).toFixed(2));const o=(i*100).toFixed(2),n=document.getElementById(`probability-percentage-${e}`);n!==null&&(n.innerHTML=`${o}%`)}for(let e=0;e<10;e++){const i=document.getElementById(`probability-percentage-${e}`);i.style.backgroundColor=e===s&&l[s]>0?"green":"rgba(0,0,0,0)"}}function k(l){v(l[0],768,128,l[1],128,1,"First Layer. Dense [768->128]"),v(l[2],128,10,l[3],10,1,"Output Layer. Dense [128->10]")}function v(l,t,s,e,i,a,c){const o=document.getElementById("parameters-container"),n=document.createElement("div");n.className="parameter-images",o.appendChild(n);const r=document.createElement("h4");r.innerHTML=c,n.appendChild(r),w(n,l,t,s,"Weights"),w(n,e,i,a,"Bias")}function w(l,t,s,e,i){const a=document.createElement("canvas"),c=s<768?10:2;a.width=s*c,a.height=e*c;const o=a.getContext("2d"),n=document.createElement("div");n.className="parameter-image",n.appendChild(document.createTextNode(`${i}`)),n.appendChild(a),l.appendChild(n);const r=D(t);for(let d=0;ds&&(s=t[i]),t[i]0&&this.stepCount%this.schedule.everyNSteps===0&&(this.currentLr*=this.schedule.factor);break;case"exponential":this.currentLr=this.baseLr*Math.pow(this.schedule.decayRate,this.stepCount);break;case"cosine":const t=Math.min(this.stepCount/this.schedule.maxSteps,1);this.currentLr=this.schedule.minLr+.5*(this.baseLr-this.schedule.minLr)*(1+Math.cos(Math.PI*t));break}}zeroGrad(){for(const t of this.params)t.gradient&&this.tm.writeBufferF32(t.gradient.buffer,new Float32Array(t.gradient.size).fill(0))}getLearningRate(){return this.currentLr}setLearningRate(t){this.baseLr=t,this.currentLr=t}setSchedule(t){this.schedule=t,this.currentLr=this.baseLr}getStepCount(){return this.stepCount}resetStepCount(){this.stepCount=0,this.currentLr=this.baseLr}}function v(r,t,e){const n=new Set,i=[];function a(s){if(!(n.has(s)||!s.requiresGradient)){n.add(s);for(const o of s.parents??[])a(o);i.push(s)}}a(e);for(const s of i)s.gradient=void 0;e.gradient=r.scopedOnes(e.shape);for(let s=i.length-1;s>=0;s--){const o=i[s];if(!o.gradFn||!o.gradient)continue;const h=o.gradFn.backward(o.gradient);for(let u=0;u<(o.parents??[]).length;u++){const c=o.parents[u],l=h[u];!c.requiresGradient||l===null||(c.gradient===void 0&&(c.gradient=r.getTensorBuffer(`${c.name}_grad`,GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,l.shape,new Float32Array(l.size).fill(0))),t.inplaceAdd.run(c.gradient,l))}}}class ${constructor(t,e,n,i,a,s,o=15,h=32){this.tm=t,this.kernelRegistry=e,this.mnist=n,this.datasource=i,this.onTrainingFinished=a,this.onUpdateData=s,this.epochs=o,this.batchSize=h,this.optimizer=new R(n.model.parameters(),.05,t,e,h)}optimizer;iterator=null;currentEpoch=0;state="idle";saveSnapshot=!1;async initialize(){this.currentEpoch=0;const t=Math.min(this.datasource.trainImagesCount,this.datasource.maxTrainSize),e=Math.ceil(t/this.batchSize);this.optimizer.setSchedule({type:"cosine",minLr:.001,maxSteps:e*this.epochs}),this.iterator=this.datasource.getTrainIterator(this.batchSize),await this.mnist.restart()}async snapshot(){if(this.saveSnapshot)for(const t of this.mnist.model.parameters()){const e=await this.tm.readBuffer(t.buffer,t.sizeInBytes()),n=new Blob([e],{type:"application/octet-stream"}),i=URL.createObjectURL(n),a=document.createElement("a");a.href=i,a.download=`model-${this.currentEpoch}-${t.name}.bin`,a.click(),URL.revokeObjectURL(i)}}finished(){return this.currentEpoch>=this.epochs}async sync(){await m.device.queue.onSubmittedWorkDone()}async trainBatch(t){this.optimizer.zeroGrad();const e=t.size,n=this.tm.getTensorBuffer("input",GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[e,784],t.data),i=this.tm.getTensorBuffer("labels_onehot",GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[e,10],t.labels);this.tm.beginScope("fwd");const a=this.mnist.model.forward(n,!0),s=this.kernelRegistry.crossEntropy.run(a,i);this.tm.beginScope("bwd"),v(this.tm,this.kernelRegistry,s),this.optimizer.step(e),await this.sync(),this.tm.flushDestroyQueue()}async trainStep(){if(!this.iterator)throw new Error("Cannot train step without initializing trainer");if(this.state==="cancelling"){this.state="cancelled",this.iterator=null;return}this.iterator.hasNext()?await this.trainBatch(this.iterator.next()):(this.iterator=this.datasource.getTrainIterator(this.batchSize),this.currentEpoch++,await this.snapshot()),this.finished()?(this.state="finished",this.onUpdateData(this.currentEpoch,this.epochs,this.iterator.getCurrentIndex(),this.iterator.getSize()),this.onTrainingFinished()):(this.onUpdateData(this.currentEpoch,this.epochs,this.iterator.getCurrentIndex(),this.iterator.getSize()),this.state==="training"&&requestAnimationFrame(this.trainStep.bind(this)))}startTraining(){this.state="training",requestAnimationFrame(()=>this.trainStep())}cancelTraining(){this.state==="training"&&(this.state="cancelling")}}class M{constructor(t,e,n,i,a,s,o,h=32){this.tm=t,this.kernelRegistry=e,this.mnist=n,this.datasource=i,this.onTestResult=a,this.onTestFinished=s,this.onWrongGuess=o,this.batchSize=h}state="idle";testCorrect=0;testTotal=0;iterator=null;startTesting(){if(this.state!=="idle")throw new Error("Cannot start testing while already running.");if(!this.iterator)throw new Error("Cannot start testing without initializing. Call initialize() first.");this.state="testing",this.testCorrect=0,this.testTotal=0,requestAnimationFrame(this.testStep.bind(this))}finished(){return this.iterator!==null&&!this.iterator.hasNext()}async testStep(){if(!this.iterator)throw new Error("Cannot test step without initializing tester");if(this.state==="cancelling"){this.state="cancelled",this.iterator=null;return}this.iterator.hasNext()?(await this.testBatch(this.iterator.next()),this.onTestResult(this.testCorrect,this.testTotal,this.iterator.getCurrentIndex(),this.iterator.getSize())):this.onTestFinished(this.testCorrect,this.testTotal),this.finished()?(this.state="finished",this.onTestFinished(this.testCorrect,this.testTotal)):this.state==="testing"&&requestAnimationFrame(this.testStep.bind(this))}async initialize(){this.testTotal=0,this.testCorrect=0,await this.datasource.load("data/mnist").catch(t=>{throw new Error("Failed to load data source: "+t)}),this.iterator=this.datasource.getTestIterator(this.batchSize)}async testBatch(t){const e=this.tm.getTensorBuffer("test_input",GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC|GPUBufferUsage.COPY_DST,[t.size,784],t.data);this.tm.beginScope("test");const n=this.mnist.model.forward(e,!1),i=this.kernelRegistry.softmax.run(n),a=await this.tm.readBuffer(i.buffer,i.sizeInBytes());for(let s=0;so&&(o=l,h=c)}let u=0;for(let c=0;c<10;c++)if(t.labels[s*10+c]>.5){u=c;break}if(this.testTotal++,h===u)this.testCorrect++;else{const c=this.testTotal-this.testCorrect;this.onWrongGuess(t.data.subarray(s*28*28,(s+1)*28*28),h,u,c)}}}cancelTesting(){this.state==="testing"&&(this.state="cancelling")}}await m.init();const g=new G(m.device),b=new x(m.device,g),T=new I(g,b);await T.readSnapshot();const C=new P;await C.load("data/mnist").catch(r=>{throw new Error("Failed to load data source: "+r)});const B=document.getElementById("timer");let f=null,w=0;function E(r){const t=Math.floor(r/6e4),e=Math.floor(r%6e4/1e3),n=r%1e3;return`${t.toString().padStart(2,"0")}:${e.toString().padStart(2,"0")}.${n.toString().padStart(3,"0")}`}function k(){w=performance.now(),f=window.setInterval(()=>{const r=performance.now()-w;B.textContent=E(Math.floor(r))},10)}function F(){if(f!==null){clearInterval(f),f=null;const r=performance.now()-w;B.textContent=E(Math.floor(r))}}async function O(r,t){const e=new $(g,b,T,C,()=>{F(),p.textContent="Done",requestAnimationFrame(()=>L.startTesting())},(n,i,a,s)=>{const o=document.getElementById("out"),h=`Epoch ${n}/${i} (${a}/${s})`;o!==null&&(o.innerHTML=h)},r,t);await e.initialize(),k(),e.startTraining()}const L=new M(g,b,T,C,(r,t,e,n)=>{const i=document.getElementById("outtest"),a=(r/t*100).toFixed(2),s=`Testing: ${e}/${n} - Accuracy: ${a}%`;i!==null&&(i.innerHTML=s)},(r,t)=>{const e=document.getElementById("outtest"),n=(r/t*100).toFixed(2),i=`Test complete: ${r}/${t} (${n}%)`;e!==null&&(e.innerHTML=i)},(r,t,e,n)=>{if(n>=100)return;const i=document.getElementById("errors-container");if(!i)return;const a=document.createElement("div");a.className="error-sample";const s=4,o=document.createElement("canvas");o.width=28*s,o.height=28*s,a.appendChild(o);const h=o.getContext("2d");if(!h)throw new Error("MNISTDatasource: failed to get 2d canvas context");for(let l=0;l<28;l++)for(let d=0;d<28;d++){const U=l*28+d,S=Math.floor(r[U]*255);h.fillStyle=`rgba(${S}, ${S}, ${S}, 1)`,h.fillRect(d*s,l*s,s,s)}const u=document.createElement("span");u.textContent=`Predicted: ${t}`,a.appendChild(u);const c=document.createElement("span");c.textContent=`Actual: ${e}`,a.appendChild(c),i.appendChild(a)});await L.initialize();const p=document.getElementById("start-btn"),y=document.getElementById("epochs-select"),z=document.getElementById("batchsize-select");p.addEventListener("click",async()=>{const r=parseInt(y.value,10),t=parseInt(z.value,10);p.disabled=!0,y.disabled=!0,z.disabled=!0,p.textContent="Training...",await O(r,t)}); diff --git a/guess.html b/guess.html index 43bc803..45b5ee0 100644 --- a/guess.html +++ b/guess.html @@ -183,6 +183,8 @@ } } + +

MNIST Interactive Demo

@@ -214,6 +216,5 @@

Learned Parameters

Weight and bias matrices from the pre-trained model, normalized to 0–1 range for visualization.

- diff --git a/package-lock.json b/package-lock.json deleted file mode 100644 index 396905b..0000000 --- a/package-lock.json +++ /dev/null @@ -1,1888 +0,0 @@ -{ - "name": "mnist", - "version": "1.0.0", - "lockfileVersion": 3, - "requires": true, - "packages": { - "": { - "name": "mnist", - "version": "1.0.0", - "devDependencies": { - "@webgpu/types": "^0.1.69", - "gh-pages": "^6.3.0", - "typescript": "^5.9.3", - "vite": "^7.3.1", - "vite-plugin-static-copy": "^3.1.6" - } - }, - "node_modules/@esbuild/aix-ppc64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/aix-ppc64/-/aix-ppc64-0.27.2.tgz", - "integrity": "sha512-GZMB+a0mOMZs4MpDbj8RJp4cw+w1WV5NYD6xzgvzUJ5Ek2jerwfO2eADyI6ExDSUED+1X8aMbegahsJi+8mgpw==", - "cpu": [ - "ppc64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "aix" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/android-arm": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/android-arm/-/android-arm-0.27.2.tgz", - "integrity": "sha512-DVNI8jlPa7Ujbr1yjU2PfUSRtAUZPG9I1RwW4F4xFB1Imiu2on0ADiI/c3td+KmDtVKNbi+nffGDQMfcIMkwIA==", - "cpu": [ - "arm" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "android" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/android-arm64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/android-arm64/-/android-arm64-0.27.2.tgz", - "integrity": "sha512-pvz8ZZ7ot/RBphf8fv60ljmaoydPU12VuXHImtAs0XhLLw+EXBi2BLe3OYSBslR4rryHvweW5gmkKFwTiFy6KA==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "android" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/android-x64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/android-x64/-/android-x64-0.27.2.tgz", - "integrity": "sha512-z8Ank4Byh4TJJOh4wpz8g2vDy75zFL0TlZlkUkEwYXuPSgX8yzep596n6mT7905kA9uHZsf/o2OJZubl2l3M7A==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "android" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/darwin-arm64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/darwin-arm64/-/darwin-arm64-0.27.2.tgz", - "integrity": "sha512-davCD2Zc80nzDVRwXTcQP/28fiJbcOwvdolL0sOiOsbwBa72kegmVU0Wrh1MYrbuCL98Omp5dVhQFWRKR2ZAlg==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "darwin" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/darwin-x64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/darwin-x64/-/darwin-x64-0.27.2.tgz", - "integrity": "sha512-ZxtijOmlQCBWGwbVmwOF/UCzuGIbUkqB1faQRf5akQmxRJ1ujusWsb3CVfk/9iZKr2L5SMU5wPBi1UWbvL+VQA==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "darwin" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/freebsd-arm64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/freebsd-arm64/-/freebsd-arm64-0.27.2.tgz", - "integrity": "sha512-lS/9CN+rgqQ9czogxlMcBMGd+l8Q3Nj1MFQwBZJyoEKI50XGxwuzznYdwcav6lpOGv5BqaZXqvBSiB/kJ5op+g==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "freebsd" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/freebsd-x64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/freebsd-x64/-/freebsd-x64-0.27.2.tgz", - "integrity": "sha512-tAfqtNYb4YgPnJlEFu4c212HYjQWSO/w/h/lQaBK7RbwGIkBOuNKQI9tqWzx7Wtp7bTPaGC6MJvWI608P3wXYA==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "freebsd" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/linux-arm": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/linux-arm/-/linux-arm-0.27.2.tgz", - "integrity": "sha512-vWfq4GaIMP9AIe4yj1ZUW18RDhx6EPQKjwe7n8BbIecFtCQG4CfHGaHuh7fdfq+y3LIA2vGS/o9ZBGVxIDi9hw==", - "cpu": [ - "arm" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/linux-arm64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/linux-arm64/-/linux-arm64-0.27.2.tgz", - "integrity": "sha512-hYxN8pr66NsCCiRFkHUAsxylNOcAQaxSSkHMMjcpx0si13t1LHFphxJZUiGwojB1a/Hd5OiPIqDdXONia6bhTw==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/linux-ia32": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/linux-ia32/-/linux-ia32-0.27.2.tgz", - "integrity": "sha512-MJt5BRRSScPDwG2hLelYhAAKh9imjHK5+NE/tvnRLbIqUWa+0E9N4WNMjmp/kXXPHZGqPLxggwVhz7QP8CTR8w==", - "cpu": [ - "ia32" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/linux-loong64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/linux-loong64/-/linux-loong64-0.27.2.tgz", - "integrity": "sha512-lugyF1atnAT463aO6KPshVCJK5NgRnU4yb3FUumyVz+cGvZbontBgzeGFO1nF+dPueHD367a2ZXe1NtUkAjOtg==", - "cpu": [ - "loong64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/linux-mips64el": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/linux-mips64el/-/linux-mips64el-0.27.2.tgz", - "integrity": "sha512-nlP2I6ArEBewvJ2gjrrkESEZkB5mIoaTswuqNFRv/WYd+ATtUpe9Y09RnJvgvdag7he0OWgEZWhviS1OTOKixw==", - "cpu": [ - "mips64el" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/linux-ppc64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/linux-ppc64/-/linux-ppc64-0.27.2.tgz", - "integrity": "sha512-C92gnpey7tUQONqg1n6dKVbx3vphKtTHJaNG2Ok9lGwbZil6DrfyecMsp9CrmXGQJmZ7iiVXvvZH6Ml5hL6XdQ==", - "cpu": [ - "ppc64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/linux-riscv64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/linux-riscv64/-/linux-riscv64-0.27.2.tgz", - "integrity": "sha512-B5BOmojNtUyN8AXlK0QJyvjEZkWwy/FKvakkTDCziX95AowLZKR6aCDhG7LeF7uMCXEJqwa8Bejz5LTPYm8AvA==", - "cpu": [ - "riscv64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/linux-s390x": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/linux-s390x/-/linux-s390x-0.27.2.tgz", - "integrity": "sha512-p4bm9+wsPwup5Z8f4EpfN63qNagQ47Ua2znaqGH6bqLlmJ4bx97Y9JdqxgGZ6Y8xVTixUnEkoKSHcpRlDnNr5w==", - "cpu": [ - "s390x" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/linux-x64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/linux-x64/-/linux-x64-0.27.2.tgz", - "integrity": "sha512-uwp2Tip5aPmH+NRUwTcfLb+W32WXjpFejTIOWZFw/v7/KnpCDKG66u4DLcurQpiYTiYwQ9B7KOeMJvLCu/OvbA==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/netbsd-arm64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/netbsd-arm64/-/netbsd-arm64-0.27.2.tgz", - "integrity": "sha512-Kj6DiBlwXrPsCRDeRvGAUb/LNrBASrfqAIok+xB0LxK8CHqxZ037viF13ugfsIpePH93mX7xfJp97cyDuTZ3cw==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "netbsd" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/netbsd-x64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/netbsd-x64/-/netbsd-x64-0.27.2.tgz", - "integrity": "sha512-HwGDZ0VLVBY3Y+Nw0JexZy9o/nUAWq9MlV7cahpaXKW6TOzfVno3y3/M8Ga8u8Yr7GldLOov27xiCnqRZf0tCA==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "netbsd" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/openbsd-arm64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/openbsd-arm64/-/openbsd-arm64-0.27.2.tgz", - "integrity": "sha512-DNIHH2BPQ5551A7oSHD0CKbwIA/Ox7+78/AWkbS5QoRzaqlev2uFayfSxq68EkonB+IKjiuxBFoV8ESJy8bOHA==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "openbsd" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/openbsd-x64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/openbsd-x64/-/openbsd-x64-0.27.2.tgz", - "integrity": "sha512-/it7w9Nb7+0KFIzjalNJVR5bOzA9Vay+yIPLVHfIQYG/j+j9VTH84aNB8ExGKPU4AzfaEvN9/V4HV+F+vo8OEg==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "openbsd" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/openharmony-arm64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/openharmony-arm64/-/openharmony-arm64-0.27.2.tgz", - "integrity": "sha512-LRBbCmiU51IXfeXk59csuX/aSaToeG7w48nMwA6049Y4J4+VbWALAuXcs+qcD04rHDuSCSRKdmY63sruDS5qag==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "openharmony" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/sunos-x64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/sunos-x64/-/sunos-x64-0.27.2.tgz", - "integrity": "sha512-kMtx1yqJHTmqaqHPAzKCAkDaKsffmXkPHThSfRwZGyuqyIeBvf08KSsYXl+abf5HDAPMJIPnbBfXvP2ZC2TfHg==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "sunos" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/win32-arm64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/win32-arm64/-/win32-arm64-0.27.2.tgz", - "integrity": "sha512-Yaf78O/B3Kkh+nKABUF++bvJv5Ijoy9AN1ww904rOXZFLWVc5OLOfL56W+C8F9xn5JQZa3UX6m+IktJnIb1Jjg==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "win32" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/win32-ia32": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/win32-ia32/-/win32-ia32-0.27.2.tgz", - "integrity": "sha512-Iuws0kxo4yusk7sw70Xa2E2imZU5HoixzxfGCdxwBdhiDgt9vX9VUCBhqcwY7/uh//78A1hMkkROMJq9l27oLQ==", - "cpu": [ - "ia32" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "win32" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@esbuild/win32-x64": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/@esbuild/win32-x64/-/win32-x64-0.27.2.tgz", - "integrity": "sha512-sRdU18mcKf7F+YgheI/zGf5alZatMUTKj/jNS6l744f9u3WFu4v7twcUI9vu4mknF4Y9aDlblIie0IM+5xxaqQ==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "win32" - ], - "engines": { - "node": ">=18" - } - }, - "node_modules/@nodelib/fs.scandir": { - "version": "2.1.5", - "resolved": "https://registry.npmjs.org/@nodelib/fs.scandir/-/fs.scandir-2.1.5.tgz", - "integrity": "sha512-vq24Bq3ym5HEQm2NKCr3yXDwjc7vTsEThRDnkp2DK9p1uqLR+DHurm/NOTo0KG7HYHU7eppKZj3MyqYuMBf62g==", - "dev": true, - "license": "MIT", - "dependencies": { - "@nodelib/fs.stat": "2.0.5", - "run-parallel": "^1.1.9" - }, - "engines": { - "node": ">= 8" - } - }, - "node_modules/@nodelib/fs.stat": { - "version": "2.0.5", - "resolved": "https://registry.npmjs.org/@nodelib/fs.stat/-/fs.stat-2.0.5.tgz", - "integrity": "sha512-RkhPPp2zrqDAQA/2jNhnztcPAlv64XdhIp7a7454A5ovI7Bukxgt7MX7udwAu3zg1DcpPU0rz3VV1SeaqvY4+A==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">= 8" - } - }, - "node_modules/@nodelib/fs.walk": { - "version": "1.2.8", - "resolved": "https://registry.npmjs.org/@nodelib/fs.walk/-/fs.walk-1.2.8.tgz", - "integrity": "sha512-oGB+UxlgWcgQkgwo8GcEGwemoTFt3FIO9ababBmaGwXIoBKZ+GTy0pP185beGg7Llih/NSHSV2XAs1lnznocSg==", - "dev": true, - "license": "MIT", - "dependencies": { - "@nodelib/fs.scandir": "2.1.5", - "fastq": "^1.6.0" - }, - "engines": { - "node": ">= 8" - } - }, - "node_modules/@rollup/rollup-android-arm-eabi": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-android-arm-eabi/-/rollup-android-arm-eabi-4.55.1.tgz", - "integrity": "sha512-9R0DM/ykwfGIlNu6+2U09ga0WXeZ9MRC2Ter8jnz8415VbuIykVuc6bhdrbORFZANDmTDvq26mJrEVTl8TdnDg==", - "cpu": [ - "arm" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "android" - ] - }, - "node_modules/@rollup/rollup-android-arm64": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-android-arm64/-/rollup-android-arm64-4.55.1.tgz", - "integrity": "sha512-eFZCb1YUqhTysgW3sj/55du5cG57S7UTNtdMjCW7LwVcj3dTTcowCsC8p7uBdzKsZYa8J7IDE8lhMI+HX1vQvg==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "android" - ] - }, - "node_modules/@rollup/rollup-darwin-arm64": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-darwin-arm64/-/rollup-darwin-arm64-4.55.1.tgz", - "integrity": "sha512-p3grE2PHcQm2e8PSGZdzIhCKbMCw/xi9XvMPErPhwO17vxtvCN5FEA2mSLgmKlCjHGMQTP6phuQTYWUnKewwGg==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "darwin" - ] - }, - "node_modules/@rollup/rollup-darwin-x64": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-darwin-x64/-/rollup-darwin-x64-4.55.1.tgz", - "integrity": "sha512-rDUjG25C9qoTm+e02Esi+aqTKSBYwVTaoS1wxcN47/Luqef57Vgp96xNANwt5npq9GDxsH7kXxNkJVEsWEOEaQ==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "darwin" - ] - }, - "node_modules/@rollup/rollup-freebsd-arm64": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-freebsd-arm64/-/rollup-freebsd-arm64-4.55.1.tgz", - "integrity": "sha512-+JiU7Jbp5cdxekIgdte0jfcu5oqw4GCKr6i3PJTlXTCU5H5Fvtkpbs4XJHRmWNXF+hKmn4v7ogI5OQPaupJgOg==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "freebsd" - ] - }, - "node_modules/@rollup/rollup-freebsd-x64": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-freebsd-x64/-/rollup-freebsd-x64-4.55.1.tgz", - "integrity": "sha512-V5xC1tOVWtLLmr3YUk2f6EJK4qksksOYiz/TCsFHu/R+woubcLWdC9nZQmwjOAbmExBIVKsm1/wKmEy4z4u4Bw==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "freebsd" - ] - }, - "node_modules/@rollup/rollup-linux-arm-gnueabihf": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-arm-gnueabihf/-/rollup-linux-arm-gnueabihf-4.55.1.tgz", - "integrity": "sha512-Rn3n+FUk2J5VWx+ywrG/HGPTD9jXNbicRtTM11e/uorplArnXZYsVifnPPqNNP5BsO3roI4n8332ukpY/zN7rQ==", - "cpu": [ - "arm" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-arm-musleabihf": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-arm-musleabihf/-/rollup-linux-arm-musleabihf-4.55.1.tgz", - "integrity": "sha512-grPNWydeKtc1aEdrJDWk4opD7nFtQbMmV7769hiAaYyUKCT1faPRm2av8CX1YJsZ4TLAZcg9gTR1KvEzoLjXkg==", - "cpu": [ - "arm" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-arm64-gnu": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-arm64-gnu/-/rollup-linux-arm64-gnu-4.55.1.tgz", - "integrity": "sha512-a59mwd1k6x8tXKcUxSyISiquLwB5pX+fJW9TkWU46lCqD/GRDe9uDN31jrMmVP3feI3mhAdvcCClhV8V5MhJFQ==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-arm64-musl": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-arm64-musl/-/rollup-linux-arm64-musl-4.55.1.tgz", - "integrity": "sha512-puS1MEgWX5GsHSoiAsF0TYrpomdvkaXm0CofIMG5uVkP6IBV+ZO9xhC5YEN49nsgYo1DuuMquF9+7EDBVYu4uA==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-loong64-gnu": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-loong64-gnu/-/rollup-linux-loong64-gnu-4.55.1.tgz", - "integrity": "sha512-r3Wv40in+lTsULSb6nnoudVbARdOwb2u5fpeoOAZjFLznp6tDU8kd+GTHmJoqZ9lt6/Sys33KdIHUaQihFcu7g==", - "cpu": [ - "loong64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-loong64-musl": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-loong64-musl/-/rollup-linux-loong64-musl-4.55.1.tgz", - "integrity": "sha512-MR8c0+UxAlB22Fq4R+aQSPBayvYa3+9DrwG/i1TKQXFYEaoW3B5b/rkSRIypcZDdWjWnpcvxbNaAJDcSbJU3Lw==", - "cpu": [ - "loong64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-ppc64-gnu": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-ppc64-gnu/-/rollup-linux-ppc64-gnu-4.55.1.tgz", - "integrity": "sha512-3KhoECe1BRlSYpMTeVrD4sh2Pw2xgt4jzNSZIIPLFEsnQn9gAnZagW9+VqDqAHgm1Xc77LzJOo2LdigS5qZ+gw==", - "cpu": [ - "ppc64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-ppc64-musl": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-ppc64-musl/-/rollup-linux-ppc64-musl-4.55.1.tgz", - "integrity": "sha512-ziR1OuZx0vdYZZ30vueNZTg73alF59DicYrPViG0NEgDVN8/Jl87zkAPu4u6VjZST2llgEUjaiNl9JM6HH1Vdw==", - "cpu": [ - "ppc64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-riscv64-gnu": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-riscv64-gnu/-/rollup-linux-riscv64-gnu-4.55.1.tgz", - "integrity": "sha512-uW0Y12ih2XJRERZ4jAfKamTyIHVMPQnTZcQjme2HMVDAHY4amf5u414OqNYC+x+LzRdRcnIG1YodLrrtA8xsxw==", - "cpu": [ - "riscv64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-riscv64-musl": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-riscv64-musl/-/rollup-linux-riscv64-musl-4.55.1.tgz", - "integrity": "sha512-u9yZ0jUkOED1BFrqu3BwMQoixvGHGZ+JhJNkNKY/hyoEgOwlqKb62qu+7UjbPSHYjiVy8kKJHvXKv5coH4wDeg==", - "cpu": [ - "riscv64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-s390x-gnu": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-s390x-gnu/-/rollup-linux-s390x-gnu-4.55.1.tgz", - "integrity": "sha512-/0PenBCmqM4ZUd0190j7J0UsQ/1nsi735iPRakO8iPciE7BQ495Y6msPzaOmvx0/pn+eJVVlZrNrSh4WSYLxNg==", - "cpu": [ - "s390x" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-x64-gnu": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-x64-gnu/-/rollup-linux-x64-gnu-4.55.1.tgz", - "integrity": "sha512-a8G4wiQxQG2BAvo+gU6XrReRRqj+pLS2NGXKm8io19goR+K8lw269eTrPkSdDTALwMmJp4th2Uh0D8J9bEV1vg==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-linux-x64-musl": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-x64-musl/-/rollup-linux-x64-musl-4.55.1.tgz", - "integrity": "sha512-bD+zjpFrMpP/hqkfEcnjXWHMw5BIghGisOKPj+2NaNDuVT+8Ds4mPf3XcPHuat1tz89WRL+1wbcxKY3WSbiT7w==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "linux" - ] - }, - "node_modules/@rollup/rollup-openbsd-x64": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-openbsd-x64/-/rollup-openbsd-x64-4.55.1.tgz", - "integrity": "sha512-eLXw0dOiqE4QmvikfQ6yjgkg/xDM+MdU9YJuP4ySTibXU0oAvnEWXt7UDJmD4UkYialMfOGFPJnIHSe/kdzPxg==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "openbsd" - ] - }, - "node_modules/@rollup/rollup-openharmony-arm64": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-openharmony-arm64/-/rollup-openharmony-arm64-4.55.1.tgz", - "integrity": "sha512-xzm44KgEP11te3S2HCSyYf5zIzWmx3n8HDCc7EE59+lTcswEWNpvMLfd9uJvVX8LCg9QWG67Xt75AuHn4vgsXw==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "openharmony" - ] - }, - "node_modules/@rollup/rollup-win32-arm64-msvc": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-win32-arm64-msvc/-/rollup-win32-arm64-msvc-4.55.1.tgz", - "integrity": "sha512-yR6Bl3tMC/gBok5cz/Qi0xYnVbIxGx5Fcf/ca0eB6/6JwOY+SRUcJfI0OpeTpPls7f194as62thCt/2BjxYN8g==", - "cpu": [ - "arm64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "win32" - ] - }, - "node_modules/@rollup/rollup-win32-ia32-msvc": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-win32-ia32-msvc/-/rollup-win32-ia32-msvc-4.55.1.tgz", - "integrity": "sha512-3fZBidchE0eY0oFZBnekYCfg+5wAB0mbpCBuofh5mZuzIU/4jIVkbESmd2dOsFNS78b53CYv3OAtwqkZZmU5nA==", - "cpu": [ - "ia32" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "win32" - ] - }, - "node_modules/@rollup/rollup-win32-x64-gnu": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-win32-x64-gnu/-/rollup-win32-x64-gnu-4.55.1.tgz", - "integrity": "sha512-xGGY5pXj69IxKb4yv/POoocPy/qmEGhimy/FoTpTSVju3FYXUQQMFCaZZXJVidsmGxRioZAwpThl/4zX41gRKg==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "win32" - ] - }, - "node_modules/@rollup/rollup-win32-x64-msvc": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/@rollup/rollup-win32-x64-msvc/-/rollup-win32-x64-msvc-4.55.1.tgz", - "integrity": "sha512-SPEpaL6DX4rmcXtnhdrQYgzQ5W2uW3SCJch88lB2zImhJRhIIK44fkUrgIV/Q8yUNfw5oyZ5vkeQsZLhCb06lw==", - "cpu": [ - "x64" - ], - "dev": true, - "license": "MIT", - "optional": true, - "os": [ - "win32" - ] - }, - "node_modules/@types/estree": { - "version": "1.0.8", - "resolved": "https://registry.npmjs.org/@types/estree/-/estree-1.0.8.tgz", - "integrity": "sha512-dWHzHa2WqEXI/O1E9OjrocMTKJl2mSrEolh1Iomrv6U+JuNwaHXsXx9bLu5gG7BUWFIN0skIQJQ/L1rIex4X6w==", - "dev": true, - "license": "MIT" - }, - "node_modules/@webgpu/types": { - "version": "0.1.69", - "resolved": "https://registry.npmjs.org/@webgpu/types/-/types-0.1.69.tgz", - "integrity": "sha512-RPmm6kgRbI8e98zSD3RVACvnuktIja5+yLgDAkTmxLr90BEwdTXRQWNLF3ETTTyH/8mKhznZuN5AveXYFEsMGQ==", - "dev": true, - "license": "BSD-3-Clause" - }, - "node_modules/anymatch": { - "version": "3.1.3", - "resolved": "https://registry.npmjs.org/anymatch/-/anymatch-3.1.3.tgz", - "integrity": "sha512-KMReFUr0B4t+D+OBkjR3KYqvocp2XaSzO55UcB6mgQMd3KbcE+mWTyvVV7D/zsdEbNnV6acZUutkiHQXvTr1Rw==", - "dev": true, - "license": "ISC", - "dependencies": { - "normalize-path": "^3.0.0", - "picomatch": "^2.0.4" - }, - "engines": { - "node": ">= 8" - } - }, - "node_modules/anymatch/node_modules/picomatch": { - "version": "2.3.1", - "resolved": "https://registry.npmjs.org/picomatch/-/picomatch-2.3.1.tgz", - "integrity": "sha512-JU3teHTNjmE2VCGFzuY8EXzCDVwEqB2a8fsIvwaStHhAWJEeVd1o1QD80CU6+ZdEXXSLbSsuLwJjkCBWqRQUVA==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=8.6" - }, - "funding": { - "url": "https://github.com/sponsors/jonschlinkert" - } - }, - "node_modules/array-union": { - "version": "2.1.0", - "resolved": "https://registry.npmjs.org/array-union/-/array-union-2.1.0.tgz", - "integrity": "sha512-HGyxoOTYUyCM6stUe6EJgnd4EoewAI7zMdfqO+kGjnlZmBDz/cR5pf8r/cR4Wq60sL/p0IkcjUEEPwS3GFrIyw==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=8" - } - }, - "node_modules/async": { - "version": "3.2.6", - "resolved": "https://registry.npmjs.org/async/-/async-3.2.6.tgz", - "integrity": "sha512-htCUDlxyyCLMgaM3xXg0C0LW2xqfuQ6p05pCEIsXuyQ+a1koYKTuBMzRNwmybfLgvJDMd0r1LTn4+E0Ti6C2AA==", - "dev": true, - "license": "MIT" - }, - "node_modules/binary-extensions": { - "version": "2.3.0", - "resolved": "https://registry.npmjs.org/binary-extensions/-/binary-extensions-2.3.0.tgz", - "integrity": "sha512-Ceh+7ox5qe7LJuLHoY0feh3pHuUDHAcRUeyL2VYghZwfpkNIy/+8Ocg0a3UuSoYzavmylwuLWQOf3hl0jjMMIw==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=8" - }, - "funding": { - "url": "https://github.com/sponsors/sindresorhus" - } - }, - "node_modules/braces": { - "version": "3.0.3", - "resolved": "https://registry.npmjs.org/braces/-/braces-3.0.3.tgz", - "integrity": "sha512-yQbXgO/OSZVD2IsiLlro+7Hf6Q18EJrKSEsdoMzKePKXct3gvD8oLcOQdIzGupr5Fj+EDe8gO/lxc1BzfMpxvA==", - "dev": true, - "license": "MIT", - "dependencies": { - "fill-range": "^7.1.1" - }, - "engines": { - "node": ">=8" - } - }, - "node_modules/chokidar": { - "version": "3.6.0", - "resolved": "https://registry.npmjs.org/chokidar/-/chokidar-3.6.0.tgz", - "integrity": "sha512-7VT13fmjotKpGipCW9JEQAusEPE+Ei8nl6/g4FBAmIm0GOOLMua9NDDo/DWp0ZAxCr3cPq5ZpBqmPAQgDda2Pw==", - "dev": true, - "license": "MIT", - "dependencies": { - "anymatch": "~3.1.2", - "braces": "~3.0.2", - "glob-parent": "~5.1.2", - "is-binary-path": "~2.1.0", - "is-glob": "~4.0.1", - "normalize-path": "~3.0.0", - "readdirp": "~3.6.0" - }, - "engines": { - "node": ">= 8.10.0" - }, - "funding": { - "url": "https://paulmillr.com/funding/" - }, - "optionalDependencies": { - "fsevents": "~2.3.2" - } - }, - "node_modules/commander": { - "version": "13.1.0", - "resolved": "https://registry.npmjs.org/commander/-/commander-13.1.0.tgz", - "integrity": "sha512-/rFeCpNJQbhSZjGVwO9RFV3xPqbnERS8MmIQzCtD/zl6gpJuV/bMLuN92oG3F7d8oDEHHRrujSXNUr8fpjntKw==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=18" - } - }, - "node_modules/commondir": { - "version": "1.0.1", - "resolved": "https://registry.npmjs.org/commondir/-/commondir-1.0.1.tgz", - "integrity": "sha512-W9pAhw0ja1Edb5GVdIF1mjZw/ASI0AlShXM83UUGe2DVr5TdAPEA1OA8m/g8zWp9x6On7gqufY+FatDbC3MDQg==", - "dev": true, - "license": "MIT" - }, - "node_modules/dir-glob": { - "version": "3.0.1", - "resolved": "https://registry.npmjs.org/dir-glob/-/dir-glob-3.0.1.tgz", - "integrity": "sha512-WkrWp9GR4KXfKGYzOLmTuGVi1UWFfws377n9cc55/tb6DuqyF6pcQ5AbiHEshaDpY9v6oaSr2XCDidGmMwdzIA==", - "dev": true, - "license": "MIT", - "dependencies": { - "path-type": "^4.0.0" - }, - "engines": { - "node": ">=8" - } - }, - "node_modules/email-addresses": { - "version": "5.0.0", - "resolved": "https://registry.npmjs.org/email-addresses/-/email-addresses-5.0.0.tgz", - "integrity": "sha512-4OIPYlA6JXqtVn8zpHpGiI7vE6EQOAg16aGnDMIAlZVinnoZ8208tW1hAbjWydgN/4PLTT9q+O1K6AH/vALJGw==", - "dev": true, - "license": "MIT" - }, - "node_modules/esbuild": { - "version": "0.27.2", - "resolved": "https://registry.npmjs.org/esbuild/-/esbuild-0.27.2.tgz", - "integrity": "sha512-HyNQImnsOC7X9PMNaCIeAm4ISCQXs5a5YasTXVliKv4uuBo1dKrG0A+uQS8M5eXjVMnLg3WgXaKvprHlFJQffw==", - "dev": true, - "hasInstallScript": true, - "license": "MIT", - "bin": { - "esbuild": "bin/esbuild" - }, - "engines": { - "node": ">=18" - }, - "optionalDependencies": { - "@esbuild/aix-ppc64": "0.27.2", - "@esbuild/android-arm": "0.27.2", - "@esbuild/android-arm64": "0.27.2", - "@esbuild/android-x64": "0.27.2", - "@esbuild/darwin-arm64": "0.27.2", - "@esbuild/darwin-x64": "0.27.2", - "@esbuild/freebsd-arm64": "0.27.2", - "@esbuild/freebsd-x64": "0.27.2", - "@esbuild/linux-arm": "0.27.2", - "@esbuild/linux-arm64": "0.27.2", - "@esbuild/linux-ia32": "0.27.2", - "@esbuild/linux-loong64": "0.27.2", - "@esbuild/linux-mips64el": "0.27.2", - "@esbuild/linux-ppc64": "0.27.2", - "@esbuild/linux-riscv64": "0.27.2", - "@esbuild/linux-s390x": "0.27.2", - "@esbuild/linux-x64": "0.27.2", - "@esbuild/netbsd-arm64": "0.27.2", - "@esbuild/netbsd-x64": "0.27.2", - "@esbuild/openbsd-arm64": "0.27.2", - "@esbuild/openbsd-x64": "0.27.2", - "@esbuild/openharmony-arm64": "0.27.2", - "@esbuild/sunos-x64": "0.27.2", - "@esbuild/win32-arm64": "0.27.2", - "@esbuild/win32-ia32": "0.27.2", - "@esbuild/win32-x64": "0.27.2" - } - }, - "node_modules/escape-string-regexp": { - "version": "1.0.5", - "resolved": "https://registry.npmjs.org/escape-string-regexp/-/escape-string-regexp-1.0.5.tgz", - "integrity": "sha512-vbRorB5FUQWvla16U8R/qgaFIya2qGzwDrNmCZuYKrbdSUMG6I1ZCGQRefkRVhuOkIGVne7BQ35DSfo1qvJqFg==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=0.8.0" - } - }, - "node_modules/fast-glob": { - "version": "3.3.3", - "resolved": "https://registry.npmjs.org/fast-glob/-/fast-glob-3.3.3.tgz", - "integrity": "sha512-7MptL8U0cqcFdzIzwOTHoilX9x5BrNqye7Z/LuC7kCMRio1EMSyqRK3BEAUD7sXRq4iT4AzTVuZdhgQ2TCvYLg==", - "dev": true, - "license": "MIT", - "dependencies": { - "@nodelib/fs.stat": "^2.0.2", - "@nodelib/fs.walk": "^1.2.3", - "glob-parent": "^5.1.2", - "merge2": "^1.3.0", - "micromatch": "^4.0.8" - }, - "engines": { - "node": ">=8.6.0" - } - }, - "node_modules/fastq": { - "version": "1.20.1", - "resolved": "https://registry.npmjs.org/fastq/-/fastq-1.20.1.tgz", - "integrity": "sha512-GGToxJ/w1x32s/D2EKND7kTil4n8OVk/9mycTc4VDza13lOvpUZTGX3mFSCtV9ksdGBVzvsyAVLM6mHFThxXxw==", - "dev": true, - "license": "ISC", - "dependencies": { - "reusify": "^1.0.4" - } - }, - "node_modules/fdir": { - "version": "6.5.0", - "resolved": "https://registry.npmjs.org/fdir/-/fdir-6.5.0.tgz", - "integrity": "sha512-tIbYtZbucOs0BRGqPJkshJUYdL+SDH7dVM8gjy+ERp3WAUjLEFJE+02kanyHtwjWOnwrKYBiwAmM0p4kLJAnXg==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=12.0.0" - }, - "peerDependencies": { - "picomatch": "^3 || ^4" - }, - "peerDependenciesMeta": { - "picomatch": { - "optional": true - } - } - }, - "node_modules/filename-reserved-regex": { - "version": "2.0.0", - "resolved": "https://registry.npmjs.org/filename-reserved-regex/-/filename-reserved-regex-2.0.0.tgz", - "integrity": "sha512-lc1bnsSr4L4Bdif8Xb/qrtokGbq5zlsms/CYH8PP+WtCkGNF65DPiQY8vG3SakEdRn8Dlnm+gW/qWKKjS5sZzQ==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=4" - } - }, - "node_modules/filenamify": { - "version": "4.3.0", - "resolved": "https://registry.npmjs.org/filenamify/-/filenamify-4.3.0.tgz", - "integrity": "sha512-hcFKyUG57yWGAzu1CMt/dPzYZuv+jAJUT85bL8mrXvNe6hWj6yEHEc4EdcgiA6Z3oi1/9wXJdZPXF2dZNgwgOg==", - "dev": true, - "license": "MIT", - "dependencies": { - "filename-reserved-regex": "^2.0.0", - "strip-outer": "^1.0.1", - "trim-repeated": "^1.0.0" - }, - "engines": { - "node": ">=8" - }, - "funding": { - "url": "https://github.com/sponsors/sindresorhus" - } - }, - "node_modules/fill-range": { - "version": "7.1.1", - "resolved": "https://registry.npmjs.org/fill-range/-/fill-range-7.1.1.tgz", - "integrity": "sha512-YsGpe3WHLK8ZYi4tWDg2Jy3ebRz2rXowDxnld4bkQB00cc/1Zw9AWnC0i9ztDJitivtQvaI9KaLyKrc+hBW0yg==", - "dev": true, - "license": "MIT", - "dependencies": { - "to-regex-range": "^5.0.1" - }, - "engines": { - "node": ">=8" - } - }, - "node_modules/find-cache-dir": { - "version": "3.3.2", - "resolved": "https://registry.npmjs.org/find-cache-dir/-/find-cache-dir-3.3.2.tgz", - "integrity": "sha512-wXZV5emFEjrridIgED11OoUKLxiYjAcqot/NJdAkOhlJ+vGzwhOAfcG5OX1jP+S0PcjEn8bdMJv+g2jwQ3Onig==", - "dev": true, - "license": "MIT", - "dependencies": { - "commondir": "^1.0.1", - "make-dir": "^3.0.2", - "pkg-dir": "^4.1.0" - }, - "engines": { - "node": ">=8" - }, - "funding": { - "url": "https://github.com/avajs/find-cache-dir?sponsor=1" - } - }, - "node_modules/find-up": { - "version": "4.1.0", - "resolved": "https://registry.npmjs.org/find-up/-/find-up-4.1.0.tgz", - "integrity": "sha512-PpOwAdQ/YlXQ2vj8a3h8IipDuYRi3wceVQQGYWxNINccq40Anw7BlsEXCMbt1Zt+OLA6Fq9suIpIWD0OsnISlw==", - "dev": true, - "license": "MIT", - "dependencies": { - "locate-path": "^5.0.0", - "path-exists": "^4.0.0" - }, - "engines": { - "node": ">=8" - } - }, - "node_modules/fs-extra": { - "version": "11.3.3", - "resolved": "https://registry.npmjs.org/fs-extra/-/fs-extra-11.3.3.tgz", - "integrity": "sha512-VWSRii4t0AFm6ixFFmLLx1t7wS1gh+ckoa84aOeapGum0h+EZd1EhEumSB+ZdDLnEPuucsVB9oB7cxJHap6Afg==", - "dev": true, - "license": "MIT", - "dependencies": { - "graceful-fs": "^4.2.0", - "jsonfile": "^6.0.1", - "universalify": "^2.0.0" - }, - "engines": { - "node": ">=14.14" - } - }, - "node_modules/fsevents": { - "version": "2.3.3", - "resolved": "https://registry.npmjs.org/fsevents/-/fsevents-2.3.3.tgz", - "integrity": "sha512-5xoDfX+fL7faATnagmWPpbFtwh/R77WmMMqqHGS65C3vvB0YHrgF+B1YmZ3441tMj5n63k0212XNoJwzlhffQw==", - "dev": true, - "hasInstallScript": true, - "license": "MIT", - "optional": true, - "os": [ - "darwin" - ], - "engines": { - "node": "^8.16.0 || ^10.6.0 || >=11.0.0" - } - }, - "node_modules/gh-pages": { - "version": "6.3.0", - "resolved": "https://registry.npmjs.org/gh-pages/-/gh-pages-6.3.0.tgz", - "integrity": "sha512-Ot5lU6jK0Eb+sszG8pciXdjMXdBJ5wODvgjR+imihTqsUWF2K6dJ9HST55lgqcs8wWcw6o6wAsUzfcYRhJPXbA==", - "dev": true, - "license": "MIT", - "dependencies": { - "async": "^3.2.4", - "commander": "^13.0.0", - "email-addresses": "^5.0.0", - "filenamify": "^4.3.0", - "find-cache-dir": "^3.3.1", - "fs-extra": "^11.1.1", - "globby": "^11.1.0" - }, - "bin": { - "gh-pages": "bin/gh-pages.js", - "gh-pages-clean": "bin/gh-pages-clean.js" - }, - "engines": { - "node": ">=10" - } - }, - "node_modules/glob-parent": { - "version": "5.1.2", - "resolved": "https://registry.npmjs.org/glob-parent/-/glob-parent-5.1.2.tgz", - "integrity": "sha512-AOIgSQCepiJYwP3ARnGx+5VnTu2HBYdzbGP45eLw1vr3zB3vZLeyed1sC9hnbcOc9/SrMyM5RPQrkGz4aS9Zow==", - "dev": true, - "license": "ISC", - "dependencies": { - "is-glob": "^4.0.1" - }, - "engines": { - "node": ">= 6" - } - }, - "node_modules/globby": { - "version": "11.1.0", - "resolved": "https://registry.npmjs.org/globby/-/globby-11.1.0.tgz", - "integrity": "sha512-jhIXaOzy1sb8IyocaruWSn1TjmnBVs8Ayhcy83rmxNJ8q2uWKCAj3CnJY+KpGSXCueAPc0i05kVvVKtP1t9S3g==", - "dev": true, - "license": "MIT", - "dependencies": { - "array-union": "^2.1.0", - "dir-glob": "^3.0.1", - "fast-glob": "^3.2.9", - "ignore": "^5.2.0", - "merge2": "^1.4.1", - "slash": "^3.0.0" - }, - "engines": { - "node": ">=10" - }, - "funding": { - "url": "https://github.com/sponsors/sindresorhus" - } - }, - "node_modules/graceful-fs": { - "version": "4.2.11", - "resolved": "https://registry.npmjs.org/graceful-fs/-/graceful-fs-4.2.11.tgz", - "integrity": "sha512-RbJ5/jmFcNNCcDV5o9eTnBLJ/HszWV0P73bc+Ff4nS/rJj+YaS6IGyiOL0VoBYX+l1Wrl3k63h/KrH+nhJ0XvQ==", - "dev": true, - "license": "ISC" - }, - "node_modules/ignore": { - "version": "5.3.2", - "resolved": "https://registry.npmjs.org/ignore/-/ignore-5.3.2.tgz", - "integrity": "sha512-hsBTNUqQTDwkWtcdYI2i06Y/nUBEsNEDJKjWdigLvegy8kDuJAS8uRlpkkcQpyEXL0Z/pjDy5HBmMjRCJ2gq+g==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">= 4" - } - }, - "node_modules/is-binary-path": { - "version": "2.1.0", - "resolved": "https://registry.npmjs.org/is-binary-path/-/is-binary-path-2.1.0.tgz", - "integrity": "sha512-ZMERYes6pDydyuGidse7OsHxtbI7WVeUEozgR/g7rd0xUimYNlvZRE/K2MgZTjWy725IfelLeVcEM97mmtRGXw==", - "dev": true, - "license": "MIT", - "dependencies": { - "binary-extensions": "^2.0.0" - }, - "engines": { - "node": ">=8" - } - }, - "node_modules/is-extglob": { - "version": "2.1.1", - "resolved": "https://registry.npmjs.org/is-extglob/-/is-extglob-2.1.1.tgz", - "integrity": "sha512-SbKbANkN603Vi4jEZv49LeVJMn4yGwsbzZworEoyEiutsN3nJYdbO36zfhGJ6QEDpOZIFkDtnq5JRxmvl3jsoQ==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=0.10.0" - } - }, - "node_modules/is-glob": { - "version": "4.0.3", - "resolved": "https://registry.npmjs.org/is-glob/-/is-glob-4.0.3.tgz", - "integrity": "sha512-xelSayHH36ZgE7ZWhli7pW34hNbNl8Ojv5KVmkJD4hBdD3th8Tfk9vYasLM+mXWOZhFkgZfxhLSnrwRr4elSSg==", - "dev": true, - "license": "MIT", - "dependencies": { - "is-extglob": "^2.1.1" - }, - "engines": { - "node": ">=0.10.0" - } - }, - "node_modules/is-number": { - "version": "7.0.0", - "resolved": "https://registry.npmjs.org/is-number/-/is-number-7.0.0.tgz", - "integrity": "sha512-41Cifkg6e8TylSpdtTpeLVMqvSBEVzTttHvERD741+pnZ8ANv0004MRL43QKPDlK9cGvNp6NZWZUBlbGXYxxng==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=0.12.0" - } - }, - "node_modules/jsonfile": { - "version": "6.2.0", - "resolved": "https://registry.npmjs.org/jsonfile/-/jsonfile-6.2.0.tgz", - "integrity": "sha512-FGuPw30AdOIUTRMC2OMRtQV+jkVj2cfPqSeWXv1NEAJ1qZ5zb1X6z1mFhbfOB/iy3ssJCD+3KuZ8r8C3uVFlAg==", - "dev": true, - "license": "MIT", - "dependencies": { - "universalify": "^2.0.0" - }, - "optionalDependencies": { - "graceful-fs": "^4.1.6" - } - }, - "node_modules/locate-path": { - "version": "5.0.0", - "resolved": "https://registry.npmjs.org/locate-path/-/locate-path-5.0.0.tgz", - "integrity": "sha512-t7hw9pI+WvuwNJXwk5zVHpyhIqzg2qTlklJOf0mVxGSbe3Fp2VieZcduNYjaLDoy6p9uGpQEGWG87WpMKlNq8g==", - "dev": true, - "license": "MIT", - "dependencies": { - "p-locate": "^4.1.0" - }, - "engines": { - "node": ">=8" - } - }, - "node_modules/make-dir": { - "version": "3.1.0", - "resolved": "https://registry.npmjs.org/make-dir/-/make-dir-3.1.0.tgz", - "integrity": "sha512-g3FeP20LNwhALb/6Cz6Dd4F2ngze0jz7tbzrD2wAV+o9FeNHe4rL+yK2md0J/fiSf1sa1ADhXqi5+oVwOM/eGw==", - "dev": true, - "license": "MIT", - "dependencies": { - "semver": "^6.0.0" - }, - "engines": { - "node": ">=8" - }, - "funding": { - "url": "https://github.com/sponsors/sindresorhus" - } - }, - "node_modules/merge2": { - "version": "1.4.1", - "resolved": "https://registry.npmjs.org/merge2/-/merge2-1.4.1.tgz", - "integrity": "sha512-8q7VEgMJW4J8tcfVPy8g09NcQwZdbwFEqhe/WZkoIzjn/3TGDwtOCYtXGxA3O8tPzpczCCDgv+P2P5y00ZJOOg==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">= 8" - } - }, - "node_modules/micromatch": { - "version": "4.0.8", - "resolved": "https://registry.npmjs.org/micromatch/-/micromatch-4.0.8.tgz", - "integrity": "sha512-PXwfBhYu0hBCPw8Dn0E+WDYb7af3dSLVWKi3HGv84IdF4TyFoC0ysxFd0Goxw7nSv4T/PzEJQxsYsEiFCKo2BA==", - "dev": true, - "license": "MIT", - "dependencies": { - "braces": "^3.0.3", - "picomatch": "^2.3.1" - }, - "engines": { - "node": ">=8.6" - } - }, - "node_modules/micromatch/node_modules/picomatch": { - "version": "2.3.1", - "resolved": "https://registry.npmjs.org/picomatch/-/picomatch-2.3.1.tgz", - "integrity": "sha512-JU3teHTNjmE2VCGFzuY8EXzCDVwEqB2a8fsIvwaStHhAWJEeVd1o1QD80CU6+ZdEXXSLbSsuLwJjkCBWqRQUVA==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=8.6" - }, - "funding": { - "url": "https://github.com/sponsors/jonschlinkert" - } - }, - "node_modules/nanoid": { - "version": "3.3.11", - "resolved": "https://registry.npmjs.org/nanoid/-/nanoid-3.3.11.tgz", - "integrity": "sha512-N8SpfPUnUp1bK+PMYW8qSWdl9U+wwNWI4QKxOYDy9JAro3WMX7p2OeVRF9v+347pnakNevPmiHhNmZ2HbFA76w==", - "dev": true, - "funding": [ - { - "type": "github", - "url": "https://github.com/sponsors/ai" - } - ], - "license": "MIT", - "bin": { - "nanoid": "bin/nanoid.cjs" - }, - "engines": { - "node": "^10 || ^12 || ^13.7 || ^14 || >=15.0.1" - } - }, - "node_modules/normalize-path": { - "version": "3.0.0", - "resolved": "https://registry.npmjs.org/normalize-path/-/normalize-path-3.0.0.tgz", - "integrity": "sha512-6eZs5Ls3WtCisHWp9S2GUy8dqkpGi4BVSz3GaqiE6ezub0512ESztXUwUB6C6IKbQkY2Pnb/mD4WYojCRwcwLA==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=0.10.0" - } - }, - "node_modules/p-limit": { - "version": "2.3.0", - "resolved": "https://registry.npmjs.org/p-limit/-/p-limit-2.3.0.tgz", - "integrity": "sha512-//88mFWSJx8lxCzwdAABTJL2MyWB12+eIY7MDL2SqLmAkeKU9qxRvWuSyTjm3FUmpBEMuFfckAIqEaVGUDxb6w==", - "dev": true, - "license": "MIT", - "dependencies": { - "p-try": "^2.0.0" - }, - "engines": { - "node": ">=6" - }, - "funding": { - "url": "https://github.com/sponsors/sindresorhus" - } - }, - "node_modules/p-locate": { - "version": "4.1.0", - "resolved": "https://registry.npmjs.org/p-locate/-/p-locate-4.1.0.tgz", - "integrity": "sha512-R79ZZ/0wAxKGu3oYMlz8jy/kbhsNrS7SKZ7PxEHBgJ5+F2mtFW2fK2cOtBh1cHYkQsbzFV7I+EoRKe6Yt0oK7A==", - "dev": true, - "license": "MIT", - "dependencies": { - "p-limit": "^2.2.0" - }, - "engines": { - "node": ">=8" - } - }, - "node_modules/p-map": { - "version": "7.0.4", - "resolved": "https://registry.npmjs.org/p-map/-/p-map-7.0.4.tgz", - "integrity": "sha512-tkAQEw8ysMzmkhgw8k+1U/iPhWNhykKnSk4Rd5zLoPJCuJaGRPo6YposrZgaxHKzDHdDWWZvE/Sk7hsL2X/CpQ==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=18" - }, - "funding": { - "url": "https://github.com/sponsors/sindresorhus" - } - }, - "node_modules/p-try": { - "version": "2.2.0", - "resolved": "https://registry.npmjs.org/p-try/-/p-try-2.2.0.tgz", - "integrity": "sha512-R4nPAVTAU0B9D35/Gk3uJf/7XYbQcyohSKdvAxIRSNghFl4e71hVoGnBNQz9cWaXxO2I10KTC+3jMdvvoKw6dQ==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=6" - } - }, - "node_modules/path-exists": { - "version": "4.0.0", - "resolved": "https://registry.npmjs.org/path-exists/-/path-exists-4.0.0.tgz", - "integrity": "sha512-ak9Qy5Q7jYb2Wwcey5Fpvg2KoAc/ZIhLSLOSBmRmygPsGwkVVt0fZa0qrtMz+m6tJTAHfZQ8FnmB4MG4LWy7/w==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=8" - } - }, - "node_modules/path-type": { - "version": "4.0.0", - "resolved": "https://registry.npmjs.org/path-type/-/path-type-4.0.0.tgz", - "integrity": "sha512-gDKb8aZMDeD/tZWs9P6+q0J9Mwkdl6xMV8TjnGP3qJVJ06bdMgkbBlLU8IdfOsIsFz2BW1rNVT3XuNEl8zPAvw==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=8" - } - }, - "node_modules/picocolors": { - "version": "1.1.1", - "resolved": "https://registry.npmjs.org/picocolors/-/picocolors-1.1.1.tgz", - "integrity": "sha512-xceH2snhtb5M9liqDsmEw56le376mTZkEX/jEb/RxNFyegNul7eNslCXP9FDj/Lcu0X8KEyMceP2ntpaHrDEVA==", - "dev": true, - "license": "ISC" - }, - "node_modules/picomatch": { - "version": "4.0.3", - "resolved": "https://registry.npmjs.org/picomatch/-/picomatch-4.0.3.tgz", - "integrity": "sha512-5gTmgEY/sqK6gFXLIsQNH19lWb4ebPDLA4SdLP7dsWkIXHWlG66oPuVvXSGFPppYZz8ZDZq0dYYrbHfBCVUb1Q==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=12" - }, - "funding": { - "url": "https://github.com/sponsors/jonschlinkert" - } - }, - "node_modules/pkg-dir": { - "version": "4.2.0", - "resolved": "https://registry.npmjs.org/pkg-dir/-/pkg-dir-4.2.0.tgz", - "integrity": "sha512-HRDzbaKjC+AOWVXxAU/x54COGeIv9eb+6CkDSQoNTt4XyWoIJvuPsXizxu/Fr23EiekbtZwmh1IcIG/l/a10GQ==", - "dev": true, - "license": "MIT", - "dependencies": { - "find-up": "^4.0.0" - }, - "engines": { - "node": ">=8" - } - }, - "node_modules/postcss": { - "version": "8.5.6", - "resolved": "https://registry.npmjs.org/postcss/-/postcss-8.5.6.tgz", - "integrity": "sha512-3Ybi1tAuwAP9s0r1UQ2J4n5Y0G05bJkpUIO0/bI9MhwmD70S5aTWbXGBwxHrelT+XM1k6dM0pk+SwNkpTRN7Pg==", - "dev": true, - "funding": [ - { - "type": "opencollective", - "url": "https://opencollective.com/postcss/" - }, - { - "type": "tidelift", - "url": "https://tidelift.com/funding/github/npm/postcss" - }, - { - "type": "github", - "url": "https://github.com/sponsors/ai" - } - ], - "license": "MIT", - "dependencies": { - "nanoid": "^3.3.11", - "picocolors": "^1.1.1", - "source-map-js": "^1.2.1" - }, - "engines": { - "node": "^10 || ^12 || >=14" - } - }, - "node_modules/queue-microtask": { - "version": "1.2.3", - "resolved": "https://registry.npmjs.org/queue-microtask/-/queue-microtask-1.2.3.tgz", - "integrity": "sha512-NuaNSa6flKT5JaSYQzJok04JzTL1CA6aGhv5rfLW3PgqA+M2ChpZQnAC8h8i4ZFkBS8X5RqkDBHA7r4hej3K9A==", - "dev": true, - "funding": [ - { - "type": "github", - "url": "https://github.com/sponsors/feross" - }, - { - "type": "patreon", - "url": "https://www.patreon.com/feross" - }, - { - "type": "consulting", - "url": "https://feross.org/support" - } - ], - "license": "MIT" - }, - "node_modules/readdirp": { - "version": "3.6.0", - "resolved": "https://registry.npmjs.org/readdirp/-/readdirp-3.6.0.tgz", - "integrity": "sha512-hOS089on8RduqdbhvQ5Z37A0ESjsqz6qnRcffsMU3495FuTdqSm+7bhJ29JvIOsBDEEnan5DPu9t3To9VRlMzA==", - "dev": true, - "license": "MIT", - "dependencies": { - "picomatch": "^2.2.1" - }, - "engines": { - "node": ">=8.10.0" - } - }, - "node_modules/readdirp/node_modules/picomatch": { - "version": "2.3.1", - "resolved": "https://registry.npmjs.org/picomatch/-/picomatch-2.3.1.tgz", - "integrity": "sha512-JU3teHTNjmE2VCGFzuY8EXzCDVwEqB2a8fsIvwaStHhAWJEeVd1o1QD80CU6+ZdEXXSLbSsuLwJjkCBWqRQUVA==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=8.6" - }, - "funding": { - "url": "https://github.com/sponsors/jonschlinkert" - } - }, - "node_modules/reusify": { - "version": "1.1.0", - "resolved": "https://registry.npmjs.org/reusify/-/reusify-1.1.0.tgz", - "integrity": "sha512-g6QUff04oZpHs0eG5p83rFLhHeV00ug/Yf9nZM6fLeUrPguBTkTQOdpAWWspMh55TZfVQDPaN3NQJfbVRAxdIw==", - "dev": true, - "license": "MIT", - "engines": { - "iojs": ">=1.0.0", - "node": ">=0.10.0" - } - }, - "node_modules/rollup": { - "version": "4.55.1", - "resolved": "https://registry.npmjs.org/rollup/-/rollup-4.55.1.tgz", - "integrity": "sha512-wDv/Ht1BNHB4upNbK74s9usvl7hObDnvVzknxqY/E/O3X6rW1U1rV1aENEfJ54eFZDTNo7zv1f5N4edCluH7+A==", - "dev": true, - "license": "MIT", - "dependencies": { - "@types/estree": "1.0.8" - }, - "bin": { - "rollup": "dist/bin/rollup" - }, - "engines": { - "node": ">=18.0.0", - "npm": ">=8.0.0" - }, - "optionalDependencies": { - "@rollup/rollup-android-arm-eabi": "4.55.1", - "@rollup/rollup-android-arm64": "4.55.1", - "@rollup/rollup-darwin-arm64": "4.55.1", - "@rollup/rollup-darwin-x64": "4.55.1", - "@rollup/rollup-freebsd-arm64": "4.55.1", - "@rollup/rollup-freebsd-x64": "4.55.1", - "@rollup/rollup-linux-arm-gnueabihf": "4.55.1", - "@rollup/rollup-linux-arm-musleabihf": "4.55.1", - "@rollup/rollup-linux-arm64-gnu": "4.55.1", - "@rollup/rollup-linux-arm64-musl": "4.55.1", - "@rollup/rollup-linux-loong64-gnu": "4.55.1", - "@rollup/rollup-linux-loong64-musl": "4.55.1", - "@rollup/rollup-linux-ppc64-gnu": "4.55.1", - "@rollup/rollup-linux-ppc64-musl": "4.55.1", - "@rollup/rollup-linux-riscv64-gnu": "4.55.1", - "@rollup/rollup-linux-riscv64-musl": "4.55.1", - "@rollup/rollup-linux-s390x-gnu": "4.55.1", - "@rollup/rollup-linux-x64-gnu": "4.55.1", - "@rollup/rollup-linux-x64-musl": "4.55.1", - "@rollup/rollup-openbsd-x64": "4.55.1", - "@rollup/rollup-openharmony-arm64": "4.55.1", - "@rollup/rollup-win32-arm64-msvc": "4.55.1", - "@rollup/rollup-win32-ia32-msvc": "4.55.1", - "@rollup/rollup-win32-x64-gnu": "4.55.1", - "@rollup/rollup-win32-x64-msvc": "4.55.1", - "fsevents": "~2.3.2" - } - }, - "node_modules/run-parallel": { - "version": "1.2.0", - "resolved": "https://registry.npmjs.org/run-parallel/-/run-parallel-1.2.0.tgz", - "integrity": "sha512-5l4VyZR86LZ/lDxZTR6jqL8AFE2S0IFLMP26AbjsLVADxHdhB/c0GUsH+y39UfCi3dzz8OlQuPmnaJOMoDHQBA==", - "dev": true, - "funding": [ - { - "type": "github", - "url": "https://github.com/sponsors/feross" - }, - { - "type": "patreon", - "url": "https://www.patreon.com/feross" - }, - { - "type": "consulting", - "url": "https://feross.org/support" - } - ], - "license": "MIT", - "dependencies": { - "queue-microtask": "^1.2.2" - } - }, - "node_modules/semver": { - "version": "6.3.1", - "resolved": "https://registry.npmjs.org/semver/-/semver-6.3.1.tgz", - "integrity": "sha512-BR7VvDCVHO+q2xBEWskxS6DJE1qRnb7DxzUrogb71CWoSficBxYsiAGd+Kl0mmq/MprG9yArRkyrQxTO6XjMzA==", - "dev": true, - "license": "ISC", - "bin": { - "semver": "bin/semver.js" - } - }, - "node_modules/slash": { - "version": "3.0.0", - "resolved": "https://registry.npmjs.org/slash/-/slash-3.0.0.tgz", - "integrity": "sha512-g9Q1haeby36OSStwb4ntCGGGaKsaVSjQ68fBxoQcutl5fS1vuY18H3wSt3jFyFtrkx+Kz0V1G85A4MyAdDMi2Q==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">=8" - } - }, - "node_modules/source-map-js": { - "version": "1.2.1", - "resolved": "https://registry.npmjs.org/source-map-js/-/source-map-js-1.2.1.tgz", - "integrity": "sha512-UXWMKhLOwVKb728IUtQPXxfYU+usdybtUrK/8uGE8CQMvrhOpwvzDBwj0QhSL7MQc7vIsISBG8VQ8+IDQxpfQA==", - "dev": true, - "license": "BSD-3-Clause", - "engines": { - "node": ">=0.10.0" - } - }, - "node_modules/strip-outer": { - "version": "1.0.1", - "resolved": "https://registry.npmjs.org/strip-outer/-/strip-outer-1.0.1.tgz", - "integrity": "sha512-k55yxKHwaXnpYGsOzg4Vl8+tDrWylxDEpknGjhTiZB8dFRU5rTo9CAzeycivxV3s+zlTKwrs6WxMxR95n26kwg==", - "dev": true, - "license": "MIT", - "dependencies": { - "escape-string-regexp": "^1.0.2" - }, - "engines": { - "node": ">=0.10.0" - } - }, - "node_modules/tinyglobby": { - "version": "0.2.15", - "resolved": "https://registry.npmjs.org/tinyglobby/-/tinyglobby-0.2.15.tgz", - "integrity": "sha512-j2Zq4NyQYG5XMST4cbs02Ak8iJUdxRM0XI5QyxXuZOzKOINmWurp3smXu3y5wDcJrptwpSjgXHzIQxR0omXljQ==", - "dev": true, - "license": "MIT", - "dependencies": { - "fdir": "^6.5.0", - "picomatch": "^4.0.3" - }, - "engines": { - "node": ">=12.0.0" - }, - "funding": { - "url": "https://github.com/sponsors/SuperchupuDev" - } - }, - "node_modules/to-regex-range": { - "version": "5.0.1", - "resolved": "https://registry.npmjs.org/to-regex-range/-/to-regex-range-5.0.1.tgz", - "integrity": "sha512-65P7iz6X5yEr1cwcgvQxbbIw7Uk3gOy5dIdtZ4rDveLqhrdJP+Li/Hx6tyK0NEb+2GCyneCMJiGqrADCSNk8sQ==", - "dev": true, - "license": "MIT", - "dependencies": { - "is-number": "^7.0.0" - }, - "engines": { - "node": ">=8.0" - } - }, - "node_modules/trim-repeated": { - "version": "1.0.0", - "resolved": "https://registry.npmjs.org/trim-repeated/-/trim-repeated-1.0.0.tgz", - "integrity": "sha512-pkonvlKk8/ZuR0D5tLW8ljt5I8kmxp2XKymhepUeOdCEfKpZaktSArkLHZt76OB1ZvO9bssUsDty4SWhLvZpLg==", - "dev": true, - "license": "MIT", - "dependencies": { - "escape-string-regexp": "^1.0.2" - }, - "engines": { - "node": ">=0.10.0" - } - }, - "node_modules/typescript": { - "version": "5.9.3", - "resolved": "https://registry.npmjs.org/typescript/-/typescript-5.9.3.tgz", - "integrity": "sha512-jl1vZzPDinLr9eUt3J/t7V6FgNEw9QjvBPdysz9KfQDD41fQrC2Y4vKQdiaUpFT4bXlb1RHhLpp8wtm6M5TgSw==", - "dev": true, - "license": "Apache-2.0", - "bin": { - "tsc": "bin/tsc", - "tsserver": "bin/tsserver" - }, - "engines": { - "node": ">=14.17" - } - }, - "node_modules/universalify": { - "version": "2.0.1", - "resolved": "https://registry.npmjs.org/universalify/-/universalify-2.0.1.tgz", - "integrity": "sha512-gptHNQghINnc/vTGIk0SOFGFNXw7JVrlRUtConJRlvaw6DuX0wO5Jeko9sWrMBhh+PsYAZ7oXAiOnf/UKogyiw==", - "dev": true, - "license": "MIT", - "engines": { - "node": ">= 10.0.0" - } - }, - "node_modules/vite": { - "version": "7.3.1", - "resolved": "https://registry.npmjs.org/vite/-/vite-7.3.1.tgz", - "integrity": "sha512-w+N7Hifpc3gRjZ63vYBXA56dvvRlNWRczTdmCBBa+CotUzAPf5b7YMdMR/8CQoeYE5LX3W4wj6RYTgonm1b9DA==", - "dev": true, - "license": "MIT", - "dependencies": { - "esbuild": "^0.27.0", - "fdir": "^6.5.0", - "picomatch": "^4.0.3", - "postcss": "^8.5.6", - "rollup": "^4.43.0", - "tinyglobby": "^0.2.15" - }, - "bin": { - "vite": "bin/vite.js" - }, - "engines": { - "node": "^20.19.0 || >=22.12.0" - }, - "funding": { - "url": "https://github.com/vitejs/vite?sponsor=1" - }, - "optionalDependencies": { - "fsevents": "~2.3.3" - }, - "peerDependencies": { - "@types/node": "^20.19.0 || >=22.12.0", - "jiti": ">=1.21.0", - "less": "^4.0.0", - "lightningcss": "^1.21.0", - "sass": "^1.70.0", - "sass-embedded": "^1.70.0", - "stylus": ">=0.54.8", - "sugarss": "^5.0.0", - "terser": "^5.16.0", - "tsx": "^4.8.1", - "yaml": "^2.4.2" - }, - "peerDependenciesMeta": { - "@types/node": { - "optional": true - }, - "jiti": { - "optional": true - }, - "less": { - "optional": true - }, - "lightningcss": { - "optional": true - }, - "sass": { - "optional": true - }, - "sass-embedded": { - "optional": true - }, - "stylus": { - "optional": true - }, - "sugarss": { - "optional": true - }, - "terser": { - "optional": true - }, - "tsx": { - "optional": true - }, - "yaml": { - "optional": true - } - } - }, - "node_modules/vite-plugin-static-copy": { - "version": "3.1.6", - "resolved": "https://registry.npmjs.org/vite-plugin-static-copy/-/vite-plugin-static-copy-3.1.6.tgz", - "integrity": "sha512-dO8Qc71yVCmcrsKrJRgSWmWj9khI7X8fLy8X35/ZFR+Nik8CQ1uUgK7iD2KQc2AQdG51sNegSj8Tb4mDKeNYZA==", - "dev": true, - "license": "MIT", - "dependencies": { - "chokidar": "^3.6.0", - "p-map": "^7.0.4", - "picocolors": "^1.1.1", - "tinyglobby": "^0.2.15" - }, - "engines": { - "node": "^18.0.0 || >=20.0.0" - }, - "peerDependencies": { - "vite": "^5.0.0 || ^6.0.0 || ^7.0.0" - } - } - } -} diff --git a/package.json b/package.json deleted file mode 100644 index 7d86381..0000000 --- a/package.json +++ /dev/null @@ -1,20 +0,0 @@ -{ - "name": "mnist", - "version": "1.0.0", - "description": "", - "main": "index.js", - "scripts": { - "dev": "vite dev", - "build": "vite build", - "preview": "vite preview", - "deploy": "npm run build && gh-pages -d dist" - }, - "devDependencies": { - "@webgpu/types": "^0.1.69", - "gh-pages": "^6.3.0", - "typescript": "^5.9.3", - "vite": "^7.3.1", - "vite-plugin-static-copy": "^3.1.6" - }, - "private": true -} diff --git a/src/GPUEnv.ts b/src/GPUEnv.ts deleted file mode 100644 index 7ad4f8a..0000000 --- a/src/GPUEnv.ts +++ /dev/null @@ -1,21 +0,0 @@ - -export class GPUEnv { - - static device: GPUDevice; - - /** - * initialize the model and get the webgpu device. - */ - static async init() { - if (!navigator.gpu) { - throw new Error("WebGPU not available in this browser/context."); - } - - const adapter = await navigator.gpu.requestAdapter(); - if (!adapter) { - throw new Error("Failed to get GPU adapter."); - } - - GPUEnv.device = await adapter.requestDevice(); - } -} diff --git a/src/MNIST/MNIST.ts b/src/MNIST/MNIST.ts deleted file mode 100644 index bb1b453..0000000 --- a/src/MNIST/MNIST.ts +++ /dev/null @@ -1,75 +0,0 @@ -import {Sequential} from "../layer/Sequential"; -import {Linear} from "../layer/Linear"; -import {TensorManager} from "../tensor/TensorManager"; -import {KernelRegistry} from "../tensor/kernel/KernelRegistry"; -import {heUniform} from "../math/Utils"; -import {ReLU} from "../layer/ReLU"; -import {MNISTDatasource} from "./MNISTDatasource"; - -export class MNIST { - - readonly model: Sequential; - readonly firstLayer: Linear; - readonly secondLayer: Linear; - - constructor( - readonly tm: TensorManager, - kernelRegistry: KernelRegistry, - readonly initializer = heUniform, - ) { - this.firstLayer = new Linear(tm, kernelRegistry, { - name: "first", - inputFeatures: MNISTDatasource.imageSize, - outputFeatures: 128, - useBias: true, - initializer - }); - this.secondLayer = new Linear(tm, kernelRegistry, { - name: "second", - inputFeatures: 128, - outputFeatures: 10, - useBias: true, - initializer - }); - - this.model = new Sequential( - this.firstLayer, - new ReLU(tm, kernelRegistry, "ReLU1"), - this.secondLayer, - ); - } - - async readSnapshot(): Promise { - const path = "data/trained_768_128_10"; - - const firstLayerWeightsR = await fetch(`${path}/model-first_weights.bin`); - const firstWeights = new Float32Array(await firstLayerWeightsR.arrayBuffer()); - const firstLayerBiasR = await fetch(`${path}/model-first_bias.bin`); - const firstBias = new Float32Array(await firstLayerBiasR.arrayBuffer()); - this.tm.writeBufferF32(this.firstLayer.parameters()[0].buffer, firstWeights); - this.tm.writeBufferF32(this.firstLayer.parameters()[1].buffer, firstBias); - - const secondLayerWeightsR = await fetch(`${path}/model-second_weights.bin`); - const secondWeights = new Float32Array(await secondLayerWeightsR.arrayBuffer()); - const secondLayerBiasR = await fetch(`${path}/model-second_bias.bin`); - const secondBias = new Float32Array(await secondLayerBiasR.arrayBuffer()); - this.tm.writeBufferF32(this.secondLayer.parameters()[0].buffer, secondWeights); - this.tm.writeBufferF32(this.secondLayer.parameters()[1].buffer, secondBias); - - return [firstWeights, firstBias, secondWeights, secondBias]; - } - - async restart() { - this.model.zeroGrad(this.tm); - - // Clear bias - this.tm.zeros(this.firstLayer.parameters()[1]); - this.tm.zeros(this.secondLayer.parameters()[1]); - - // Reinitialize weights - const s0 = this.firstLayer.parameters()[0].shape; - this.tm.writeBufferF32(this.firstLayer.parameters()[0].buffer, heUniform(s0, s0[0])); - const s1 = this.secondLayer.parameters()[0].shape; - this.tm.writeBufferF32(this.secondLayer.parameters()[0].buffer, heUniform(s1, s1[0])); - } -} diff --git a/src/MNIST/MNISTDatasource.ts b/src/MNIST/MNISTDatasource.ts deleted file mode 100644 index 0a8d03e..0000000 --- a/src/MNIST/MNISTDatasource.ts +++ /dev/null @@ -1,203 +0,0 @@ -import {Datasource} from "../model/Datasource"; - - -export interface MNISTDataSourceIteratorValue { - data: Float32Array; - labels: Float32Array; - size: number; -} - -export class MNISTDataSourceIterator { - - private currentIndex = 0; - private iteratorSize = 0; - private datasetIndices: Uint32Array = new Uint32Array(0); - - private readonly workingImageBuffer: Float32Array; - private readonly workingLabelBuffer: Float32Array; - - constructor( - readonly batchSize: number, - private imageData: Float32Array, - private oneImageSize: number, - private labelsData: Float32Array, - private oneLabelSize: number, - private maxSize: number, - ) { - this.workingImageBuffer = new Float32Array(oneImageSize * this.batchSize); - this.workingLabelBuffer = new Float32Array(oneLabelSize * this.batchSize); - - const trainImageElemets = Math.min(this.maxSize, this.imageData.length / oneImageSize); - const trainLabelElemets = Math.min(this.maxSize, this.labelsData.length / oneLabelSize); - - if (trainImageElemets !== trainLabelElemets) { - throw new Error("MNIST data source: image and label data have different sizes"); - } - - this.iteratorSize = trainLabelElemets; - - this.restart(); - } - - getCurrentIndex() { - return this.currentIndex; - } - - getSize() { - return this.iteratorSize; - } - - hasNext(): boolean { - return this.currentIndex < this.iteratorSize; - } - - /** - * populate working buffers with the next batch of data. - * @returns the next batch of data. - */ - next(): MNISTDataSourceIteratorValue { - - const bs = Math.min(this.batchSize, this.iteratorSize - this.currentIndex); - - for (let i = 0; i < bs; i++) { - const index = this.datasetIndices[this.currentIndex++]; - - const imageIndex = index * this.oneImageSize; - this.workingImageBuffer.set( - this.imageData.subarray(imageIndex, imageIndex + this.oneImageSize), - i * this.oneImageSize - ); - - const labelIndex = index * this.oneLabelSize; - this.workingLabelBuffer.set( - this.labelsData.subarray(labelIndex, labelIndex + this.oneLabelSize), - i * this.oneLabelSize - ); - } - - return { - data: bs < this.batchSize ? - this.workingImageBuffer.subarray(0, bs * this.oneImageSize) : - this.workingImageBuffer, - labels: bs < this.batchSize ? - this.workingLabelBuffer.subarray(0, bs * this.oneLabelSize) : - this.workingLabelBuffer, - size: bs, - }; - } - - restart(): void { - this.currentIndex = 0; - this.datasetIndices = MNISTDataSourceIterator.getRandomDataset( - this.imageData.length / this.oneImageSize - ); - } - - /** - * generate a random dataset of size N. - * @param size - * @private - */ - static getRandomDataset(size: number): Uint32Array { - const indices = new Uint32Array(size); - for(let i = 0; i < indices.length; i++) { - indices[i] = i; - } - - for (let i = size - 1; i > 0; i--) { - const j = Math.floor(Math.random() * (i + 1)); - [indices[i], indices[j]] = [indices[j], indices[i]]; - } - - return indices; - } - -} - -/** - * A naive MNIST data source. - * It loads in memory all training and test data. - */ -export class MNISTDatasource implements Datasource { - - trainData: Float32Array | null = null; - trainLabelsData: Float32Array | null = null; - - testData: Float32Array | null = null; - testLabelsData: Float32Array | null = null; - - static readonly imageSize = 28 * 28; - testImagesCount = 10000; - - trainImagesCount = 60000; - - maxTrainSize = Infinity; - maxTestSize = Infinity; - - constructor() { - } - - private toFloat32(uint: Uint8Array): Float32Array { - const ret = new Float32Array(uint.length); - for (let i = 0; i < uint.length; i++) { - ret[i] = uint[i] / 255; - } - return ret; - } - - private getIterator( - batchSize: number, - data: Float32Array, - labels: Float32Array, - maxSize: number, - ): MNISTDataSourceIterator { - return new MNISTDataSourceIterator( - batchSize, - data, - MNISTDatasource.imageSize, - labels, - 10, - maxSize, - ); - } - - getTrainIterator(batchSize: number): MNISTDataSourceIterator { - if (!this.trainData || !this.trainLabelsData) { - throw new Error("MNISTDatasource: train data not loaded. Call load() first."); - } - return this.getIterator(batchSize, this.trainData, this.trainLabelsData, this.maxTrainSize); - } - - getTestIterator(batchSize: number): MNISTDataSourceIterator { - if (!this.testData || !this.testLabelsData) { - throw new Error("MNISTDatasource: test data not loaded. Call load() first."); - } - return this.getIterator(batchSize, this.testData, this.testLabelsData, this.maxTestSize); - } - - private onehot(uint: Uint8Array): Float32Array { - const output = new Float32Array(uint.length * 10); - - for (let i = 0; i < uint.length; i++) { - const index = uint[i]; - output[i * 10 + index] = 1.0; - } - - return output; - } - - async load(path: string) { - const trainImagesResponse = await fetch(`${path}/train-images.idx3-ubyte`) - this.trainData = this.toFloat32(new Uint8Array(await trainImagesResponse.arrayBuffer(), 16)); - const trainLabelsResponse = await fetch(`${path}/train-labels.idx1-ubyte`); - this.trainLabelsData = this.onehot(new Uint8Array(await trainLabelsResponse.arrayBuffer(), 8)); - - const testDataResponse = await fetch(`${path}/t10k-images.idx3-ubyte`); - this.testData = this.toFloat32(new Uint8Array(await testDataResponse.arrayBuffer(), 16)); - const testLabelsResponse = await fetch(`${path}/t10k-labels.idx1-ubyte`); - this.testLabelsData = this.onehot(new Uint8Array(await testLabelsResponse.arrayBuffer(), 8)); - - this.testImagesCount = this.testData.length / MNISTDatasource.imageSize; - this.trainImagesCount = this.trainData.length / MNISTDatasource.imageSize; - } -} diff --git a/src/MNIST/Tester.ts b/src/MNIST/Tester.ts deleted file mode 100644 index 4bb19b6..0000000 --- a/src/MNIST/Tester.ts +++ /dev/null @@ -1,154 +0,0 @@ -import {MNISTDatasource, MNISTDataSourceIterator, MNISTDataSourceIteratorValue} from "./MNISTDatasource"; -import {TensorManager} from "../tensor/TensorManager"; -import {KernelRegistry} from "../tensor/kernel/KernelRegistry"; -import {MNIST} from "./MNIST"; -import {GPUEnv} from "../GPUEnv"; - -export type TesterState = "idle" - | "testing" - | "finished" - | "cancelling" - | "cancelled" - ; - -export class Tester { - private state: TesterState = "idle"; - private testCorrect = 0; - private testTotal = 0; - private iterator: MNISTDataSourceIterator | null = null; - - constructor( - readonly tm: TensorManager, - readonly kernelRegistry: KernelRegistry, - readonly mnist: MNIST, - readonly datasource: MNISTDatasource, - readonly onTestResult: (correct: number, total: number, current: number, size: number,) => void, - readonly onTestFinished: (correct: number, total: number,) => void, - readonly onWrongGuess: (imageData: Float32Array, predicted: number, label: number, errorCount: number) => void, - readonly batchSize: number = 32, - ) { - - } - - startTesting() { - if (this.state !== "idle") { - throw new Error("Cannot start testing while already running."); - } - if (!this.iterator) { - throw new Error("Cannot start testing without initializing. Call initialize() first."); - } - - this.state = "testing"; - this.testCorrect = 0; - this.testTotal = 0; - - requestAnimationFrame(this.testStep.bind(this)); - } - - private finished() { - return this.iterator !== null && !this.iterator.hasNext(); - } - - private async testStep() { - - if (!this.iterator) { - throw new Error("Cannot test step without initializing tester"); - } - - if (this.state === "cancelling") { - this.state = "cancelled"; - this.iterator = null; - return; - } - - if (this.iterator.hasNext()) { - await this.testBatch(this.iterator.next()); - this.onTestResult( - this.testCorrect, - this.testTotal, - this.iterator.getCurrentIndex(), - this.iterator.getSize(), - ) - } else { - this.onTestFinished(this.testCorrect, this.testTotal); - } - - if (!this.finished()) { - if (this.state === "testing") { - requestAnimationFrame(this.testStep.bind(this)); - } - } else { - this.state = "finished"; - this.onTestFinished(this.testCorrect, this.testTotal); - } - } - - async initialize() { - - this.testTotal = 0; - this.testCorrect = 0; - - await this.datasource - .load("data/mnist") - .catch((e: Error) => { - throw new Error("Failed to load data source: " + e) - }); - - this.iterator = this.datasource.getTestIterator(this.batchSize); - } - - private async testBatch(data: MNISTDataSourceIteratorValue) { - - const input = this.tm.getTensorBuffer( - "test_input", - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [data.size, 28 * 28], - data.data, - ); - - // Forward pass (isTraining = false to disable dropout) - this.tm.beginScope("test"); - const logits = this.mnist.model.forward(input, false); - const probs = this.kernelRegistry.softmax.run(logits); - - // Read back predictions (readBuffer already syncs GPU work) - const probsData = await this.tm.readBuffer(probs.buffer, probs.sizeInBytes()); - - // Calculate accuracy - for (let i = 0; i < data.size; i++) { - // Find predicted class (argmax of probs) - let maxProb = -Infinity; - let predicted = 0; - for (let j = 0; j < 10; j++) { - const prob = probsData[i * 10 + j]; - if (prob > maxProb) { - maxProb = prob; - predicted = j; - } - } - - // Find actual class (argmax of one-hot labels) - let actual = 0; - for (let j = 0; j < 10; j++) { - if (data.labels[i * 10 + j] > 0.5) { - actual = j; - break; - } - } - - this.testTotal++; - if (predicted === actual) { - this.testCorrect++; - } else { - const errorCount = this.testTotal - this.testCorrect; - this.onWrongGuess(data.data.subarray(i * 28 * 28, (i + 1) * 28 * 28), predicted, actual, errorCount); - } - } - } - - cancelTesting() { - if (this.state === "testing") { - this.state = "cancelling"; - } - } -} diff --git a/src/MNIST/TrainAndTest.ts b/src/MNIST/TrainAndTest.ts deleted file mode 100644 index 3c6d81e..0000000 --- a/src/MNIST/TrainAndTest.ts +++ /dev/null @@ -1,162 +0,0 @@ -import {GPUEnv} from "../GPUEnv"; -import {KernelRegistry} from "../tensor/kernel/KernelRegistry"; -import {TensorManager} from "../tensor/TensorManager"; -import {MNISTDatasource} from "./MNISTDatasource"; -import {MNIST} from "./MNIST"; -import {Trainer} from "./Trainer"; -import {Tester} from "./Tester"; - -await GPUEnv.init() - -const tm = new TensorManager(GPUEnv.device); -const kernelRegistry = new KernelRegistry(GPUEnv.device, tm); - -// createa model, -const mnist = new MNIST(tm, kernelRegistry); -// load pre-trained model. 97.35% accuracy on the test set. -await mnist.readSnapshot(); - -// create train/test datasource -const datasource = new MNISTDatasource(); -await datasource - .load("data/mnist") - .catch((e: Error) => { - throw new Error("Failed to load data source: " + e) - }); - -const timerElement = document.getElementById("timer")!; -let timerInterval: number | null = null; -let startTime: number = 0; - -function formatTime(ms: number): string { - const minutes = Math.floor(ms / 60000); - const seconds = Math.floor((ms % 60000) / 1000); - const millis = ms % 1000; - return `${minutes.toString().padStart(2, "0")}:${seconds.toString().padStart(2, "0")}.${millis.toString().padStart(3, "0")}`; -} - -function startTimer() { - startTime = performance.now(); - timerInterval = window.setInterval(() => { - const elapsed = performance.now() - startTime; - timerElement.textContent = formatTime(Math.floor(elapsed)); - }, 10); -} - -function stopTimer() { - if (timerInterval !== null) { - clearInterval(timerInterval); - timerInterval = null; - const elapsed = performance.now() - startTime; - timerElement.textContent = formatTime(Math.floor(elapsed)); - } -} - -async function train(epochs: number, batchSize: number) { - - const trainer = new Trainer( - tm, - kernelRegistry, - mnist, - datasource, - () => { - stopTimer(); - startBtn.textContent = "Done"; - requestAnimationFrame(() => tester.startTesting()); - }, - (epoch, epochs, current, total) => { - const node = document.getElementById("out"); - const out = `Epoch ${epoch}/${epochs} (${current}/${total})`; - if (node !== null) { - node.innerHTML = out; - } - }, - epochs, - batchSize, - ); - await trainer.initialize(); - startTimer(); - trainer.startTraining(); -} - -const tester = new Tester( - tm, - kernelRegistry, - mnist, - datasource, - (correct, total, current, size) => { - const node = document.getElementById("outtest"); - const accuracy = (correct / total * 100).toFixed(2); - const out = `Testing: ${current}/${size} - Accuracy: ${accuracy}%`; - if (node !== null) { - node.innerHTML = out; - } - }, - (correct: number, total: number) => { - const node = document.getElementById("outtest"); - const accuracy = (correct / total * 100).toFixed(2); - const out = `Test complete: ${correct}/${total} (${accuracy}%)`; - if (node !== null) { - node.innerHTML = out; - } - }, - (imageData: Float32Array, guessed: number, label: number, errorCount: number) => { - - if (errorCount >= 100) { - return; - } - - const container = document.getElementById("errors-container"); - if (!container) return; - - const parent = document.createElement("div"); - parent.className = "error-sample"; - - const px = 4; - const canvas = document.createElement("canvas"); - canvas.width = 28 * px; - canvas.height = 28 * px; - parent.appendChild(canvas); - - const ctx = canvas.getContext("2d"); - if (!ctx) { - throw new Error("MNISTDatasource: failed to get 2d canvas context"); - } - for (let r = 0; r < 28; r++) { - for (let c = 0; c < 28; c++) { - const index = r * 28 + c; - const value = Math.floor(imageData[index] * 255); - ctx.fillStyle = `rgba(${value}, ${value}, ${value}, 1)`; - ctx.fillRect(c * px, r * px, px, px); - } - } - - const label1 = document.createElement("span"); - label1.textContent = `Predicted: ${guessed}`; - parent.appendChild(label1); - - const label2 = document.createElement("span"); - label2.textContent = `Actual: ${label}`; - parent.appendChild(label2); - - container.appendChild(parent); - } -); - -await tester.initialize(); - -const startBtn = document.getElementById("start-btn") as HTMLButtonElement; -const epochsSelect = document.getElementById("epochs-select") as HTMLSelectElement; -const batchSizeSelect = document.getElementById("batchsize-select") as HTMLSelectElement; - -startBtn.addEventListener("click", async () => { - const epochs = parseInt(epochsSelect.value, 10); - const batchSize = parseInt(batchSizeSelect.value, 10); - - startBtn.disabled = true; - epochsSelect.disabled = true; - batchSizeSelect.disabled = true; - startBtn.textContent = "Training..."; - - await train(epochs, batchSize); -}); diff --git a/src/MNIST/Trainer.ts b/src/MNIST/Trainer.ts deleted file mode 100644 index cc42849..0000000 --- a/src/MNIST/Trainer.ts +++ /dev/null @@ -1,189 +0,0 @@ -import {MNISTDatasource, MNISTDataSourceIterator, MNISTDataSourceIteratorValue} from "./MNISTDatasource"; -import {SGD} from "../optimizer/SGD"; -import {Optimizer} from "../optimizer/Optimizer"; -import {TensorManager} from "../tensor/TensorManager"; -import {KernelRegistry} from "../tensor/kernel/KernelRegistry"; -import {GPUEnv} from "../GPUEnv"; -import {computeBackwardPass} from "../autograd/BackwardPass"; -import {MNIST} from "./MNIST"; -import {Tensor} from "../tensor/Tensor"; - -export type TrainerState = "idle" - | "training" - | "finished" - | "cancelling" - | "cancelled" - ; - -export class Trainer { - - private optimizer: Optimizer; - private iterator: MNISTDataSourceIterator | null = null; - - private currentEpoch = 0; - private state: TrainerState = "idle"; - - saveSnapshot = false; - - constructor( - readonly tm: TensorManager, - readonly kernelRegistry: KernelRegistry, - readonly mnist: MNIST, - readonly datasource: MNISTDatasource, - readonly onTrainingFinished: () => void, - readonly onUpdateData: (epoch: number, epochs: number, current: number, total: number) => void, - readonly epochs: number = 15, - readonly batchSize: number = 32, - ) { - this.optimizer = new SGD( - mnist.model.parameters(), - 0.05, - tm, - kernelRegistry, - batchSize - ); - } - - async initialize() { - - this.currentEpoch = 0; - - const trainSize = Math.min( - this.datasource.trainImagesCount, - this.datasource.maxTrainSize - ); - - const stepsPerEpoch = Math.ceil(trainSize / this.batchSize); - - this.optimizer.setSchedule({ - type: "cosine", - minLr: 0.001, - maxSteps: stepsPerEpoch * this.epochs - }); - - this.iterator = this.datasource.getTrainIterator(this.batchSize); - - // restart model. - await this.mnist.restart(); - } - - private async snapshot() { - if (!this.saveSnapshot) { - return; - } - - for (const parameter of this.mnist.model.parameters()) { - const buffer = await this.tm.readBuffer(parameter.buffer, parameter.sizeInBytes()); - const blob = new Blob([buffer], {type: "application/octet-stream"}); - const url = URL.createObjectURL(blob); - const a = document.createElement("a"); - a.href = url; - a.download = `model-${this.currentEpoch}-${parameter.name}.bin`; - a.click(); - URL.revokeObjectURL(url); - } - } - - private finished() { - return this.currentEpoch >= this.epochs; - } - - private async sync() { - await GPUEnv.device.queue.onSubmittedWorkDone(); - } - - private async trainBatch(data: MNISTDataSourceIteratorValue) { - - // 1. Zero gradients - this.optimizer.zeroGrad(); - - // 2. Prepare data - const currentBatchSize = data.size; - - const input = this.tm.getTensorBuffer( - "input", - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [currentBatchSize, 28 * 28], - data.data, - ); - - const labelsOneHot = this.tm.getTensorBuffer( - "labels_onehot", - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [currentBatchSize, 10], - data.labels - ); - - // 3. Forward (begin scope for transient tensors) - this.tm.beginScope("fwd"); - const logits = this.mnist.model.forward(input, true); - const loss = this.kernelRegistry.crossEntropy.run(logits, labelsOneHot); - - // 4. Backward - this.tm.beginScope("bwd"); - computeBackwardPass(this.tm, this.kernelRegistry, loss); - - // 5. Optimize, SGD - this.optimizer.step(currentBatchSize); - - await this.sync(); - - // 6. Flush destroyed buffers to free GPU memory - this.tm.flushDestroyQueue(); - } - - async trainStep() { - - if (!this.iterator) { - throw new Error("Cannot train step without initializing trainer"); - } - - if (this.state === "cancelling") { - this.state = "cancelled"; - this.iterator = null; - return; - } - - if (this.iterator.hasNext()) { - await this.trainBatch(this.iterator.next()); - } else { - this.iterator = this.datasource.getTrainIterator(this.batchSize); - this.currentEpoch++; - await this.snapshot(); - } - - if (!this.finished()) { - this.onUpdateData( - this.currentEpoch, - this.epochs, - this.iterator.getCurrentIndex(), - this.iterator.getSize() - ); - - if (this.state === "training") { - requestAnimationFrame(this.trainStep.bind(this)); - } - } else { - // Note: snapshot() already called at line 150 after final epoch increment - this.state = "finished"; - this.onUpdateData( - this.currentEpoch, - this.epochs, - this.iterator.getCurrentIndex(), - this.iterator.getSize() - ); - this.onTrainingFinished(); - } - } - - startTraining() { - this.state = "training"; - requestAnimationFrame(() => this.trainStep()); - } - - cancelTraining() { - if (this.state === "training") { - this.state = "cancelling"; - } - } -} diff --git a/src/MNIST/interactive/Guess.ts b/src/MNIST/interactive/Guess.ts deleted file mode 100644 index f8065cb..0000000 --- a/src/MNIST/interactive/Guess.ts +++ /dev/null @@ -1,221 +0,0 @@ -import {GPUEnv} from "../../GPUEnv"; -import {KernelRegistry} from "../../tensor/kernel/KernelRegistry"; -import {TensorManager} from "../../tensor/TensorManager"; -import {MNISTDatasource} from "../MNISTDatasource"; -import {MNIST} from "../MNIST"; -import {PaintLayer} from "./PaintLayer"; -import {LayerInputs} from "./LayerInputs"; -import {Tensor} from "../../tensor/Tensor"; - -await GPUEnv.init(); - -const tm = new TensorManager(GPUEnv.device); -const kernelRegistry = new KernelRegistry(GPUEnv.device, tm); - -// create model, -const mnist = new MNIST(tm, kernelRegistry); -// load pre-trained model. 97.35% accuracy on the test set. -const parameters = await mnist.readSnapshot(); -generateParameterImages(parameters); - -// create train/test datasource -const datasource = new MNISTDatasource(); -await datasource - .load("data/mnist") - .catch((e: Error) => { - throw new Error("Failed to load data source: " + e) - }); - -const emptyData = new Float32Array(10); -let working = false; - -const painter = new PaintLayer( - async (data: Float32Array) => { - await GPUEnv.device.queue.onSubmittedWorkDone(); - if (working) { - return; - } - working = true; - await test(data) - working = false; - }, - async () => { - setProbability(emptyData); - layersActivations.render([], []); - }); - -const layersActivations = new LayerInputs(); - -(function() { - - const wrapper0 = document.getElementById("row")!; - - const wrapper = document.createElement("span"); - wrapper.style.display = "inline-flex"; - wrapper.style.flexDirection = "column"; - wrapper.style.gap = "6px"; - - let index = 0; - Array.from({ length: 10 }, () => { - - const line = document.createElement("span"); - line.classList.add("probability"); - line.id = `probability-${index}`; - - const slider = document.createElement("input"); - slider.style.pointerEvents = "none"; - slider.type = "range"; - slider.min = "0"; - slider.max = "100"; - slider.value = "0"; - slider.id = `probability-slider-${index}`; - - line.appendChild(slider); - line.appendChild(document.createTextNode(`${index}`)); - - const percentage = document.createElement("span"); - percentage.id = `probability-percentage-${index}`; - line.appendChild(percentage); - - wrapper.appendChild(line); - - index ++; - - return slider; - }); - - wrapper0.appendChild(wrapper); - -})(); - -async function test(data: Float32Array) { - const input = tm.getTensorBuffer( - "test_input", - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [1, data.length], - data, - ); - - tm.beginScope("test"); - const logits = mnist.model.forward(input, false); - const probs = kernelRegistry.softmax.run(logits); - - const probsData = await tm.readBuffer(probs.buffer, probs.sizeInBytes()); - setProbability(probsData); - - // read back layers inputs. - const outputs: Float32Array[] = []; - for(const layer of mnist.model.layers) { - const it = layer.inputTensor!; - const buffer = await tm.readBuffer(it.buffer, it.sizeInBytes()); - outputs.push(buffer); - } - outputs.push(probsData); - - const names = mnist.model.layers.map(l => `${l.name} [${l.inputTensor?.shape}]`); - names.push("Predictions. Shape [1,10]. Digit 0..9") - layersActivations.render(outputs, names); -} - -function setProbability(probsData: Float32Array) { - let maxProb = -Infinity; - let predicted = 0; - for (let j = 0; j < 10; j++) { - const prob = probsData[j]; - if (prob > maxProb) { - maxProb = prob; - predicted = j; - } - - const node = document.getElementById(`probability-${j}`); - if (node !== null) { - node.style.backgroundColor = "default"; - } - - const nodes = document.getElementById(`probability-slider-${j}`); - if (nodes !== null) { - (nodes as any).value = (100*prob).toFixed(2); - } - - const percentage = (prob * 100).toFixed(2); - const node2 = document.getElementById(`probability-percentage-${j}`); - if (node2 !== null) { - node2.innerHTML = `${percentage}%`; - } - } - - for(let j = 0; j<10; j++) { - const node = document.getElementById(`probability-percentage-${j}`)!; - node.style.backgroundColor = (j===predicted && probsData[predicted]>0) - ? "green" - : "rgba(0,0,0,0)"; - } -} - -function generateParameterImages(data: Float32Array[]) { - generateParameterImage(data[0], 768, 128, data[1], 128, 1, "First Layer. Dense [768->128]"); - generateParameterImage(data[2], 128, 10, data[3], 10, 1, "Output Layer. Dense [128->10]" ); -} - -function generateParameterImage( - weights: Float32Array, rowsw: number, colsw: number, - bias: Float32Array, rowsb: number, colsb: number, - title: string, -) { - const anchor = document.getElementById("parameters-container")!; - const div = document.createElement("div"); - div.className = "parameter-images" - anchor.appendChild(div); - - const h4 = document.createElement("h4"); - h4.innerHTML = title; - div.appendChild(h4); - - generateImage(div, weights, rowsw, colsw, "Weights"); - generateImage(div, bias, rowsb, colsb, "Bias"); -} - -function generateImage(div: HTMLDivElement, parameter: Float32Array, cols: number, rows: number, title: string) { - const canvas = document.createElement("canvas"); - - const scale = cols < 768 ? 10 : 2; - canvas.width = cols * scale; - canvas.height = rows * scale; - const ctx = canvas.getContext("2d")!; - - const container = document.createElement("div"); - container.className = "parameter-image"; - container.appendChild(document.createTextNode(`${title}`)); - container.appendChild(canvas); - - div.appendChild(container); - - - const p = normalize(parameter); - - for (let r = 0; r < rows; r++) { - for (let c = 0; c < cols; c++) { - const index = r * cols + c; - const col = Math.floor(p[index] * 255); - ctx.fillStyle = `rgba(${col}, ${col}, ${col*.8}, 1)`; - ctx.fillRect(c * scale, r * scale, scale, scale); - } - } -} - -function normalize(i: Float32Array): Float32Array { - const data = new Float32Array(i.length); - data.set(i,0); - - let max = -Infinity; - let min = Infinity; - for(let i =0; i max) max = data[i]; - if (data[i] < min) min = data[i]; - } - for(let i =0; i { - this.drawLayer(d, 10, 10 + i * 65, layers[i], i > 0); - }); - } - - private drawLayer(data: Float32Array, x: number, y: number, label: string, outline = false) { - - let max = -Infinity; - let min = Infinity; - for(let i =0; i max) max = data[i]; - if (data[i] < min) min = data[i]; - } - for(let i =0; i void, - private readonly onModelCleared: () => void, - readonly rows: number = 24, - readonly columns: number = 24, - readonly showPadded = false, - ) { - this.root = document.createElement("div"); - const row = document.getElementById("row")!; - row.appendChild(this.root); - - this.canvas = document.createElement("canvas"); - this.canvas.width = this.columns * PX; - this.canvas.height = this.rows * PX; - this.ctx = this.canvas.getContext("2d")!; - - this.model = new Float32Array(rows * columns); - this.outmodel = new Float32Array(28*28); - - this.render(this.ctx, this.rows, this.columns, this.model); - - this.canvas.addEventListener("mousedown", _ => { - this.capturing = true; - }); - - this.canvas.addEventListener("mouseup", _ => { - this.capturing = false; - }); - - this.canvas.addEventListener("mousemove", (e) => { - if (!this.capturing) { - return; - } - - this.setModel(e.offsetX, e.offsetY, e.altKey); - }); - - this.canvas.addEventListener("click", (e) => { - this.setModel(e.offsetX, e.offsetY, e.altKey); - }); - - // Touch support - this.canvas.addEventListener("touchstart", (e) => { - e.preventDefault(); - this.capturing = true; - const touch = e.touches[0]; - const rect = this.canvas.getBoundingClientRect(); - this.setModel(touch.clientX - rect.left, touch.clientY - rect.top, false); - }); - - this.canvas.addEventListener("touchend", (e) => { - e.preventDefault(); - this.capturing = false; - }); - - this.canvas.addEventListener("touchmove", (e) => { - e.preventDefault(); - if (!this.capturing) { - return; - } - const touch = e.touches[0]; - const rect = this.canvas.getBoundingClientRect(); - this.setModel(touch.clientX - rect.left, touch.clientY - rect.top, false); - }); - - this.root.appendChild(document.createTextNode(`${this.rows}x${this.columns}, padded to 28x28`)); - this.root.appendChild(document.createElement("br")); - - this.root.appendChild(document.createElement("br")); - this.root.appendChild(this.canvas); - this.root.appendChild(document.createElement("br")); - - const clearButton = document.createElement("button"); - clearButton.textContent = "Clear"; - clearButton.addEventListener("click", () => { - this.model.fill(0); - this.render(this.ctx, this.rows, this.columns, this.model); - this.onModelCleared(); - }); - this.root.appendChild(clearButton); - - if (showPadded) { - this.canvasOut = document.createElement("canvas"); - this.canvasOut.width = 28 * PX; - this.canvasOut.height = 28 * PX; - this.ctxOut = this.canvasOut.getContext("2d")!; - this.root.appendChild(document.createElement("br")); - this.root.appendChild(this.canvasOut); - this.render(this.ctxOut, 28, 28, this.outmodel); - } - } - - private incColor(x: number, y: number, sign: number = 1) { - - if (y>0) { - this.increment(x,y-1, .33 * sign); - } - if (y0) { - this.increment(x-1,y, .33 * sign); - } - if (x 0) { - if (i < minY) minY = i; - if (i > maxY) maxY = i; - if (j < minX) minX = j; - if (j > maxX) maxX = j; - } - } - } - - // calculate copy offsets - const xoffset = Math.floor((28-(maxX-minX))/2); - const yoffset = Math.floor((28-(maxY-minY))/2); - - // copy from model into outmodel - for (let i = 0; i <= maxY-minY; i++) { - for (let j = 0; j <= maxX-minX; j++) { - this.outmodel[(i+yoffset) * 28 + j + xoffset] = this.model[(i+minY) * this.columns + j + minX]; - } - } - } - - render(ctx: CanvasRenderingContext2D, rows: number, columns: number, model: Float32Array): void { - - ctx.fillStyle = "#111"; - ctx.fillRect(0, 0, ctx.canvas.width, ctx.canvas.height); - ctx.strokeStyle = "white"; - - for (let i = 0; i < rows; i++) { - for (let j = 0; j < columns; j++) { - const v = model[i * columns + j]; - const color = Math.floor(v*255); - ctx.fillStyle = `rgba(${color},${color},${color*.7},1.0)`; - ctx.fillRect(j * PX, i * PX, PX, PX); - ctx.strokeRect(j * PX, i * PX, PX, PX); - } - } - } -} diff --git a/src/autograd/BackwardPass.ts b/src/autograd/BackwardPass.ts deleted file mode 100644 index 2846f97..0000000 --- a/src/autograd/BackwardPass.ts +++ /dev/null @@ -1,67 +0,0 @@ -import {Tensor} from "../tensor/Tensor"; -import {TensorManager} from "../tensor/TensorManager"; -import {KernelRegistry} from "../tensor/kernel/KernelRegistry"; - -export function computeBackwardPass( - tm: TensorManager, - kr: KernelRegistry, - loss: Tensor -): void { - - const visited = new Set(); - const order: Tensor[] = []; - - function sort(t: Tensor): void { - if (visited.has(t) || !t.requiresGradient) { - return; - } - - visited.add(t); - - for (const parent of t.parents ?? []) { - sort(parent); - } - - order.push(t); - } - - sort(loss); - - for (const t of order) { - t.gradient = undefined; - } - - loss.gradient = tm.scopedOnes(loss.shape); - - // order has topological order of parameters. we need inverse topological order to compute gradients. - for (let i = order.length - 1; i >= 0; i--) { - const t = order[i]; - if (!t.gradFn || !t.gradient) { - continue; - } - - const inputGrads = t.gradFn.backward(t.gradient); - - for (let j = 0; j < (t.parents ?? []).length; j++) { - const parent = t.parents![j]; - const inputGrad = inputGrads[j]; - - // Skip if parent doesn't need gradient or gradient wasn't computed - if (!parent.requiresGradient || inputGrad === null) { - continue; - } - - if (parent.gradient === undefined) { - // First gradient - initialize the buffer - parent.gradient = tm.getTensorBuffer( - `${parent.name}_grad`, - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - inputGrad.shape, - new Float32Array(inputGrad.size).fill(0) - ); - } - - kr.inplaceAdd.run(parent.gradient, inputGrad); - } - } -} diff --git a/src/autograd/GradientFunction.ts b/src/autograd/GradientFunction.ts deleted file mode 100644 index 7b86670..0000000 --- a/src/autograd/GradientFunction.ts +++ /dev/null @@ -1,14 +0,0 @@ -import {Tensor} from "../tensor/Tensor"; - -export interface GradientFunction { - - name: string; - savedTensors: Tensor[]; - - /** - * Compute gradients with respect to inputs. - * Returns an array matching the number of parents. - * null entries indicate gradients that weren't computed (parent doesn't require gradient). - */ - backward(gradOutput: Tensor): (Tensor | null)[]; -} diff --git a/src/autograd/backward/BiasAddBackward.ts b/src/autograd/backward/BiasAddBackward.ts deleted file mode 100644 index ec971f1..0000000 --- a/src/autograd/backward/BiasAddBackward.ts +++ /dev/null @@ -1,33 +0,0 @@ -import {GradientFunction} from "../GradientFunction"; -import {Tensor} from "../../tensor/Tensor"; -import {KernelRegistry} from "../../tensor/kernel/KernelRegistry"; - -/** - * Backward for BiasAdd: Y = X + b (broadcast) - * - * Given dL/dY [M, N], computes: - * dL/dX = dL/dY (same shape) - * dL/db = sum(dL/dY, axis=0) (reduce to [1, N] or [N]) - */ -export class BiasAddBackward implements GradientFunction { - readonly name = "BiasAddBackward"; - - constructor( - readonly savedTensors: Tensor[], - readonly kr: KernelRegistry, - ) {} - - backward(gradOutput: Tensor): (Tensor | null)[] { - const [input, bias] = this.savedTensors; - - // dX = gradOutput (identity, same shape) - const gradInput = input.requiresGradient ? gradOutput : null; - - // db = sum(gradOutput, axis=0) -> [1, N] (only compute if bias requires grad) - const gradBias = bias.requiresGradient - ? this.kr.sumReduce.run(gradOutput) - : null; - - return [gradInput, gradBias]; - } -} diff --git a/src/autograd/backward/DropoutBackward.ts b/src/autograd/backward/DropoutBackward.ts deleted file mode 100644 index 20775a1..0000000 --- a/src/autograd/backward/DropoutBackward.ts +++ /dev/null @@ -1,29 +0,0 @@ -import {GradientFunction} from "../GradientFunction"; -import {Tensor} from "../../tensor/Tensor"; -import {KernelRegistry} from "../../tensor/kernel/KernelRegistry"; - -/** - * Backward for Dropout: Y = X * mask - * - * Given dL/dY, computes: - * dL/dX = dL/dY * mask - * - * Uses the same mask from forward pass (inverted dropout scaling included). - */ -export class DropoutBackward implements GradientFunction { - readonly name = "DropoutBackward"; - - constructor( - readonly savedTensors: Tensor[], - readonly kr: KernelRegistry, - ) {} - - backward(gradOutput: Tensor): Tensor[] { - const [mask] = this.savedTensors; - - // dX = gradOutput * mask (same mask used in forward) - const gradInput = this.kr.elemMul.run(gradOutput, mask); - - return [gradInput]; - } -} diff --git a/src/autograd/backward/MatAddBackward.ts b/src/autograd/backward/MatAddBackward.ts deleted file mode 100644 index 990abec..0000000 --- a/src/autograd/backward/MatAddBackward.ts +++ /dev/null @@ -1,26 +0,0 @@ -import {GradientFunction} from "../GradientFunction"; -import {Tensor} from "../../tensor/Tensor"; -import {KernelRegistry} from "../../tensor/kernel/KernelRegistry"; - -/** - * Backward for MatAdd: C = A + B - * - * Given dL/dC, computes: - * dL/dA = dL/dC - * dL/dB = dL/dC - * - * Both gradients are just the upstream gradient (identity). - */ -export class MatAddBackward implements GradientFunction { - readonly name = "MatAddBackward"; - - constructor( - readonly savedTensors: Tensor[], - readonly kr: KernelRegistry, - ) {} - - backward(gradOutput: Tensor): Tensor[] { - // Both inputs get the same gradient - return [gradOutput, gradOutput]; - } -} diff --git a/src/autograd/backward/MatMulBackward.ts b/src/autograd/backward/MatMulBackward.ts deleted file mode 100644 index d78be0a..0000000 --- a/src/autograd/backward/MatMulBackward.ts +++ /dev/null @@ -1,39 +0,0 @@ -import {GradientFunction} from "../GradientFunction"; -import {Tensor} from "../../tensor/Tensor"; -import {KernelRegistry} from "../../tensor/kernel/KernelRegistry"; - -/** - * Backward for MatMul: C = A @ B - * - * Given dL/dC, computes: - * dL/dA = dL/dC @ B^T - * dL/dB = A^T @ dL/dC - */ -export class MatMulBackward implements GradientFunction { - readonly name = "MatMulBackward"; - - constructor( - readonly savedTensors: Tensor[], - readonly kr: KernelRegistry, - ) {} - - backward(gradOutput: Tensor): (Tensor | null)[] { - const [A, B] = this.savedTensors; - - // dA = gradOutput @ B^T (only compute if A requires grad) - let gradA: Tensor | null = null; - if (A.requiresGradient) { - const BT = this.kr.transpose.run(B); - gradA = this.kr.matmul.run(gradOutput, BT); - } - - // dB = A^T @ gradOutput (only compute if B requires grad) - let gradB: Tensor | null = null; - if (B.requiresGradient) { - const AT = this.kr.transpose.run(A); - gradB = this.kr.matmul.run(AT, gradOutput); - } - - return [gradA, gradB]; - } -} diff --git a/src/autograd/backward/ReLUBackward.ts b/src/autograd/backward/ReLUBackward.ts deleted file mode 100644 index 7698838..0000000 --- a/src/autograd/backward/ReLUBackward.ts +++ /dev/null @@ -1,30 +0,0 @@ -import {GradientFunction} from "../GradientFunction"; -import {Tensor} from "../../tensor/Tensor"; -import {KernelRegistry} from "../../tensor/kernel/KernelRegistry"; - -/** - * Backward for ReLU: Y = max(0, X) - * - * Given dL/dY, computes: - * dL/dX = dL/dY * (X > 0) - * - * savedTensors[0] is the original input X from forward pass. - */ -export class ReLUBackward implements GradientFunction { - readonly name = "ReLUBackward"; - - constructor( - readonly savedTensors: Tensor[], - readonly kr: KernelRegistry, - ) {} - - backward(gradOutput: Tensor): Tensor[] { - // savedTensors[0] is the original input X - const [savedInput] = this.savedTensors; - - // dX = gradOutput * (savedInput > 0) - const gradInput = this.kr.reluBackward.run(gradOutput, savedInput); - - return [gradInput]; - } -} diff --git a/src/autograd/backward/SoftmaxBackward.ts b/src/autograd/backward/SoftmaxBackward.ts deleted file mode 100644 index 6893b9a..0000000 --- a/src/autograd/backward/SoftmaxBackward.ts +++ /dev/null @@ -1,30 +0,0 @@ -import {GradientFunction} from "../GradientFunction"; -import {Tensor} from "../../tensor/Tensor"; -import {KernelRegistry} from "../../tensor/kernel/KernelRegistry"; - -/** - * Softmax backward. - * - * dL/dX = S * (dS - sum(dS * S, axis=1, keepdims=True)) - * - * where S = softmax(X) and dS = dL/dSoftmax - */ -export class SoftmaxBackward implements GradientFunction { - readonly name = "SoftmaxBackward"; - - constructor( - readonly savedTensors: Tensor[], - readonly kr: KernelRegistry, - ) {} - - backward(gradOutput: Tensor): Tensor[] { - // savedTensors: [input, softmaxOutput] - const softmaxOut = this.savedTensors[1]; - - // dX = S * (dS - dot) where dot = sum(dS * S, per row) - // This is computed by the softmaxBackward kernel - const gradInput = this.kr.softmaxBackward.run(gradOutput, softmaxOut); - - return [gradInput]; - } -} diff --git a/src/autograd/backward/SoftmaxCrossEntropyBackward.ts b/src/autograd/backward/SoftmaxCrossEntropyBackward.ts deleted file mode 100644 index 6c3ca89..0000000 --- a/src/autograd/backward/SoftmaxCrossEntropyBackward.ts +++ /dev/null @@ -1,35 +0,0 @@ -import {GradientFunction} from "../GradientFunction"; -import {Tensor} from "../../tensor/Tensor"; -import {KernelRegistry} from "../../tensor/kernel/KernelRegistry"; - -/** - * Backward for combined Softmax + CrossEntropy. - * - * Forward: L = CrossEntropy(Softmax(logits), labels) - * - * Given dL/dL = 1 (scalar loss), computes: - * dL/dLogits = Softmax(logits) - labels - * - * This is numerically more stable than computing them separately. - * savedTensors: [logits, labels] - */ -export class SoftmaxCrossEntropyBackward implements GradientFunction { - readonly name = "SoftmaxCrossEntropyBackward"; - - constructor( - readonly savedTensors: Tensor[], - readonly kr: KernelRegistry, - ) {} - - backward(gradOutput: Tensor): Tensor[] { - const [logits, labels] = this.savedTensors; - - // dLogits = softmax(logits) - labels - // The kernel computes softmax internally for numerical stability - // Note: No batch averaging here - the optimizer handles batch normalization - // (SGD divides by batch size when applying gradients) - const gradLogits = this.kr.softmaxCEBackward.run(logits, labels); - - return [gradLogits]; - } -} diff --git a/src/layer/Dropout.ts b/src/layer/Dropout.ts deleted file mode 100644 index aacffbf..0000000 --- a/src/layer/Dropout.ts +++ /dev/null @@ -1,83 +0,0 @@ -import {Layer} from "./Layer"; -import {Tensor} from "../tensor/Tensor"; -import {TensorManager} from "../tensor/TensorManager"; -import {KernelRegistry} from "../tensor/kernel/KernelRegistry"; - -/** - * Dropout layer for regularization. - * - * During training: randomly zeroes elements with probability p, - * scales remaining elements by 1/(1-p) to maintain expected values. - * - * During inference: passes input through unchanged. - * - * Input/Output layout: [batch, features] - */ -export class Dropout implements Layer { - - inputTensor: Tensor | undefined; - - readonly name: string; - readonly p: number; - readonly scale: number; - - private maskData?: Float32Array; - - constructor( - readonly tm: TensorManager, - readonly kr: KernelRegistry, - name: string, - p: number = 0.5, - ) { - if (p < 0 || p >= 1) { - throw new Error("Dropout: p must be in [0, 1)"); - } - - this.name = name; - this.p = p; - this.scale = 1 / (1 - p); - } - - forward(input: Tensor, isTraining: boolean): Tensor { - - this.inputTensor = input; - - if (!isTraining) { - return input; - } - - const M = input.shape[0]; - const N = input.shape[1]; - const size = M * N; - - // Ensure mask buffer exists and is correct size - if (this.maskData === undefined || this.maskData.length < size) { - this.maskData = new Float32Array(size); - } - - // Generate random mask on CPU: 0 with prob p, scale with prob (1-p) - for (let i = 0; i < size; i++) { - this.maskData[i] = Math.random() < this.p ? 0 : this.scale; - } - - // Upload mask to GPU (scoped tensor - only needed for this batch) - const mask = this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, - [M, N], - this.maskData, - ); - - // Output buffer (scoped tensor - only needed for this batch) - const out = this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N], - ); - - return this.kr.dropout.run(input, mask, out); - } - - parameters(): Tensor[] { - return []; // Dropout has no learnable parameters - } - -} diff --git a/src/layer/Layer.ts b/src/layer/Layer.ts deleted file mode 100644 index 6bdf534..0000000 --- a/src/layer/Layer.ts +++ /dev/null @@ -1,14 +0,0 @@ -import {Tensor} from "../tensor/Tensor"; - -/** - * Base interface for all layers. - */ -export interface Layer { - - inputTensor: Tensor | undefined; - name: string; - - forward(input: Tensor, isTraining: boolean): Tensor; - - parameters(): Tensor[]; -} diff --git a/src/layer/Linear.ts b/src/layer/Linear.ts deleted file mode 100644 index 1dfbd70..0000000 --- a/src/layer/Linear.ts +++ /dev/null @@ -1,94 +0,0 @@ -import {Layer} from "./Layer"; -import {Tensor} from "../tensor/Tensor"; -import {TensorManager} from "../tensor/TensorManager"; -import {KernelRegistry} from "../tensor/kernel/KernelRegistry"; - -export type LinearInitializer = { - name: string; - readonly inputFeatures: number; - readonly outputFeatures: number; - readonly useBias: boolean; - initializer: (shape: number[], inputFeatures: number) => Float32Array; -} - -/** - * Linear Layer: output = x*W + b - * - * Weight layout: [inputFeatures, outputFeatures] (row-major) - * Input layout: [batch, inputFeatures] - * Output layout: [batch, outputFeatures] - * - * MatMul computes: input[batch, inputFeatures] * weights[inputFeatures, outputFeatures] - */ -export class Linear implements Layer { - - inputTensor: Tensor | undefined; - - readonly name: string; - - readonly weights: Tensor; - readonly bias?: Tensor; - - constructor( - readonly tm: TensorManager, - readonly kr: KernelRegistry, - init: LinearInitializer, - ) { - const shape = [init.inputFeatures, init.outputFeatures]; - - this.name = init.name; - - this.weights = tm.getTensorBuffer( - `${init.name}_weights`, - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - shape, - init.initializer(shape, init.inputFeatures), - ); - this.weights.requiresGradient = true; - - if (init.useBias) { - this.bias = tm.getTensorBuffer( - `${init.name}_bias`, - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [1, init.outputFeatures], - ); - this.bias.requiresGradient = true; - } - } - - forward(input: Tensor, isTraining: boolean): Tensor { - - this.inputTensor = input; - - const matmulout = this.tm.getTensorBuffer( - `${this.name}_mmout`, - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [input.shape[0], this.weights.shape[1]], - ) - - const mm = this.kr.matmul.run(input, this.weights, matmulout); - if (!this.bias) { - return mm; - } - - const sumout = this.tm.getTensorBuffer( - `${this.name}_sumout`, - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [input.shape[0], this.weights.shape[1]], - ) - - return this.kr.biasadd.run(matmulout, this.bias, sumout); - } - - parameters(): Tensor[] { - const ret = [ - this.weights - ]; - - if (this.bias !== undefined) { - ret.push(this.bias); - } - - return ret; - } -} diff --git a/src/layer/ReLU.ts b/src/layer/ReLU.ts deleted file mode 100644 index 78fe86b..0000000 --- a/src/layer/ReLU.ts +++ /dev/null @@ -1,34 +0,0 @@ -import {Layer} from "./Layer"; -import {Tensor} from "../tensor/Tensor"; -import {KernelRegistry} from "../tensor/kernel/KernelRegistry"; -import {TensorManager} from "../tensor/TensorManager"; - -export class ReLU implements Layer { - - inputTensor: Tensor | undefined; - - constructor( - readonly tm: TensorManager, - readonly kr: KernelRegistry, - readonly name: string = "ReLU", - ) { - - } - - forward(input: Tensor, isTraining: boolean): Tensor { - this.inputTensor = input; - - // Use scoped tensor - kernel overwrites entire buffer, no initialization needed - const output = this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - input.shape, - ); - - return this.kr.relu.run(input, output); - } - - parameters(): Tensor[] { - return []; - } - -} diff --git a/src/layer/Sequential.ts b/src/layer/Sequential.ts deleted file mode 100644 index dff0d25..0000000 --- a/src/layer/Sequential.ts +++ /dev/null @@ -1,44 +0,0 @@ -import {Layer} from "./Layer"; -import {Tensor} from "../tensor/Tensor"; -import {TensorManager} from "../tensor/TensorManager"; - -export class Sequential implements Layer { - - inputTensor: Tensor | undefined; - - readonly layers: Layer[] = []; - - constructor( - ...seq: Layer[] - ) { - for (const l of seq) { - this.layers.push(l); - } - } - - forward(input: Tensor, isTraining: boolean): Tensor { - this.inputTensor = input; - - for (const l of this.layers) { - input = l.forward(input, isTraining); - } - - return input; - } - - parameters(): Tensor[] { - return this.layers.flatMap(l => l.parameters()); - } - - zeroGrad(tm: TensorManager): void { - for (const param of this.parameters()) { - if (param.gradient) { - tm.writeBufferF32( - param.gradient.buffer, - new Float32Array(param.size).fill(0) - ); - } - } - } - -} diff --git a/src/math/Utils.ts b/src/math/Utils.ts deleted file mode 100644 index 2d7fd46..0000000 --- a/src/math/Utils.ts +++ /dev/null @@ -1,50 +0,0 @@ - -export function randn() { - const u1 = 1 - Math.random(); // (0,1] - const u2 = Math.random(); // [0,1) - return Math.sqrt(-2 * Math.log(u1)) * Math.cos(2 * Math.PI * u2); -} - -/** - * - * @param shape - * @param featuresIn are layer input features count - * @param w optional weight tensor - */ -export function heNormal( - shape: number[], - featuresIn: number, - w?: Float32Array -): Float32Array { - const std = Math.sqrt(2 / featuresIn); - const size = shape.reduce((a, b) => a * b, 1); - - w = w ?? new Float32Array(size); - - for (let i = 0; i < size; i++) { - w[i] = randn() * std; - } - return w; -} - -/** - * - * @param shape - * @param fanIn are layer input features count - * @param w optional weight tensor - */ -export function heUniform( - shape: number[], - fanIn: number, - w?: Float32Array, -): Float32Array { - const limit = Math.sqrt(6 / fanIn); - const size = shape.reduce((a, b) => a * b, 1); - - w = w ?? new Float32Array(size); - - for (let i = 0; i < size; i++) { - w[i] = (Math.random() * 2 - 1) * limit; - } - return w; -} diff --git a/src/model/Datasource.ts b/src/model/Datasource.ts deleted file mode 100644 index 7365a3c..0000000 --- a/src/model/Datasource.ts +++ /dev/null @@ -1,4 +0,0 @@ - -export interface Datasource { - load(path: string): Promise; -} diff --git a/src/optimizer/Optimizer.ts b/src/optimizer/Optimizer.ts deleted file mode 100644 index 9a1db45..0000000 --- a/src/optimizer/Optimizer.ts +++ /dev/null @@ -1,13 +0,0 @@ -export type LRSchedule = - | { type: "constant" } - | { type: "step"; factor: number; everyNSteps: number } - | { type: "exponential"; decayRate: number } - | { type: "cosine"; minLr: number; maxSteps: number }; - -export interface Optimizer { - step(batchSizeOverride?: number): void; - zeroGrad(): void; - getLearningRate(): number; - setLearningRate(lr: number): void; - setSchedule(schedule: LRSchedule): void; -} diff --git a/src/optimizer/SGD.ts b/src/optimizer/SGD.ts deleted file mode 100644 index c710d15..0000000 --- a/src/optimizer/SGD.ts +++ /dev/null @@ -1,101 +0,0 @@ -import {LRSchedule, Optimizer} from "./Optimizer"; -import {TensorManager} from "../tensor/TensorManager"; -import {KernelRegistry} from "../tensor/kernel/KernelRegistry"; -import {Tensor} from "../tensor/Tensor"; - -export class SGD implements Optimizer { - private baseLr: number; - private currentLr: number; - private stepCount = 0; - private schedule: LRSchedule = { type: "constant" }; - - constructor( - readonly params: Tensor[], - lr: number, - private tm: TensorManager, - private kr: KernelRegistry, - private batchSize: number, - ) { - this.baseLr = lr; - this.currentLr = lr; - } - - step(batchSizeOverride?: number) { - - this.updateLearningRate(); - - const effectiveBatchSize = batchSizeOverride ?? this.batchSize; - - for (const p of this.params) { - if (!p.gradient) continue; - - // p = p - lr * grad - const update = this.kr.scalarMul.run(p.gradient, -this.currentLr / effectiveBatchSize); - this.kr.inplaceAdd.run(p, update); - } - - this.stepCount++; - } - - private updateLearningRate(): void { - switch (this.schedule.type) { - case "constant": - // No change - break; - - case "step": - // Decay by factor every N steps - if (this.stepCount > 0 && this.stepCount % this.schedule.everyNSteps === 0) { - this.currentLr *= this.schedule.factor; - } - break; - - case "exponential": - // Exponential decay each step: lr = baseLr * (decayRate ^ step) - this.currentLr = this.baseLr * Math.pow(this.schedule.decayRate, this.stepCount); - break; - - case "cosine": - // Cosine annealing: lr oscillates between baseLr and minLr - const progress = Math.min(this.stepCount / this.schedule.maxSteps, 1); - this.currentLr = this.schedule.minLr + - 0.5 * (this.baseLr - this.schedule.minLr) * (1 + Math.cos(Math.PI * progress)); - break; - } - } - - zeroGrad() { - for (const p of this.params) { - if (p.gradient) { - this.tm.writeBufferF32( - p.gradient.buffer, - new Float32Array(p.gradient.size).fill(0) - ); - } - } - } - - getLearningRate(): number { - return this.currentLr; - } - - setLearningRate(lr: number): void { - this.baseLr = lr; - this.currentLr = lr; - } - - setSchedule(schedule: LRSchedule): void { - this.schedule = schedule; - // Reset to base learning rate when schedule changes - this.currentLr = this.baseLr; - } - - getStepCount(): number { - return this.stepCount; - } - - resetStepCount(): void { - this.stepCount = 0; - this.currentLr = this.baseLr; - } -} diff --git a/src/tensor/Tensor.ts b/src/tensor/Tensor.ts deleted file mode 100644 index 45667fb..0000000 --- a/src/tensor/Tensor.ts +++ /dev/null @@ -1,56 +0,0 @@ -import {GradientFunction} from "../autograd/GradientFunction"; -import {TensorManager} from "./TensorManager"; - -/** - * A Tensor is a multidimensional array. - * Right now, it will support 2D arrays. - * - * Tensors are backed by a WebGPU buffer for extra performance. - */ -export class Tensor { - - readonly size: number; - - // Autograd fields - gradient?: Tensor; - gradFn?: GradientFunction = undefined; - parents?: Tensor[] = undefined; - requiresGradient: boolean; - - /** - * Direct build from existing data that is contiguous in memory. Shape if not set, - * will be [1, initalData.length] - * - * @param name - * @param buffer - * @param usage - * @param shape optional shape of the tensor data. - * @param requiresGradient whether this tensor needs gradient computation - */ - constructor( - readonly name: string, - readonly buffer: GPUBuffer, - readonly usage: GPUBufferUsageFlags, - readonly shape: number[], - requiresGradient: boolean = false, - ) { - this.size = shape.reduce((a, b) => a * b, 1); - this.requiresGradient = requiresGradient; - } - - sizeInBytes() { - return this.size * 4; - } - - backward() { - if (!this.gradient) { - throw new Error("Tensor has no gradient"); - } - } - - zeroGrad(tm: TensorManager) { - if (this.gradient) { - tm.zeros(this.gradient); - } - } -} diff --git a/src/tensor/TensorManager.ts b/src/tensor/TensorManager.ts deleted file mode 100644 index 4c4be47..0000000 --- a/src/tensor/TensorManager.ts +++ /dev/null @@ -1,253 +0,0 @@ -import {Tensor} from "./Tensor"; - -export type ReadBackBufferInfo = { - buffer: GPUBuffer, - size: number -} - -function alignTo(n: number, multiple: number) { - return Math.ceil(n / multiple) * multiple; -} - -/** - * Tensor manager is a class responsible of handling tensors. - * It backs tensors with GPU buffers for reusability and performance. - */ -export class TensorManager { - - tensors = new Map(); - readback: ReadBackBufferInfo | null = null; - - private pendingDestroy: GPUBuffer[] = []; - - // Scoped allocation for transient tensors - private scopeName = ""; - private scopeCounter = 0; - - constructor( - readonly device: GPUDevice - ) { - } - - /** - * Begin a named scope for transient tensor allocation. - * Resets the counter so tensors get reusable names like "_fwd_0", "_fwd_1", etc. - */ - beginScope(name: string) { - this.scopeName = name; - this.scopeCounter = 0; - } - - /** - * Get a transient tensor within the current scope. - * Uses sequential naming (_scope_0, _scope_1, ...) that resets each scope, - * enabling buffer reuse across iterations. - */ - getScopedTensor( - usage: GPUBufferUsageFlags, - shape: number[], - initialData?: Float32Array, - ): Tensor { - const name = `_${this.scopeName}_${this.scopeCounter++}`; - return this.getTensorBuffer(name, usage, shape, initialData); - } - - getTensorBuffer( - name: string, - usage: GPUBufferUsageFlags, - shape: number[], - initialData: Float32Array | undefined = undefined, - ): Tensor { - const existing = this.tensors.get(name); - - const sizeBytes = shape.reduce((a, b) => a * b, 1) * 4; - - // this buffer fits in the requested one - reuse buffer but update shape - if (existing && existing.sizeInBytes() >= sizeBytes && existing.usage === usage) { - if (initialData) { - this.writeBufferF32(existing.buffer, initialData); - } - - // Create new tensor with updated shape if shape changed - const shapeMatch = existing.shape.length === shape.length - && existing.shape.every((v, i) => v === shape[i]); - - if (shapeMatch) { - return existing; - } - - const updated = new Tensor(name, existing.buffer, usage, shape); - this.tensors.set(name, updated); - return updated; - } - - const newBuf = this.device.createBuffer({ - size: alignTo(sizeBytes, 256), - usage, - }); - - if (existing) { - this.pendingDestroy.push(existing.buffer); - } - - const tensor = new Tensor(name, newBuf, usage, shape); - this.tensors.set(name, tensor); - - if (initialData) { - this.writeBufferF32(newBuf, initialData); - } - - return tensor; - } - - writeBufferF32(dstBuffer: GPUBuffer, data: Float32Array, dstOffset = 0) { - this.device.queue.writeBuffer( - dstBuffer, dstOffset, - data.buffer, data.byteOffset, data.byteLength - ); - } - - async readBuffer( - srcBuffer: GPUBuffer, - byteLength: number, - srcOffset = 0 - ) { - /* - const rb = this.ensureReadback(byteLength); - - const encoder = this.device.createCommandEncoder(); - encoder.copyBufferToBuffer(srcBuffer, srcOffset, rb, 0, byteLength); - this.device.queue.submit([encoder.finish()]); - - await this.device.queue.onSubmittedWorkDone(); - - await rb.mapAsync(GPUMapMode.READ, 0, byteLength); - const mapped = rb.getMappedRange(0, byteLength); - const copy = mapped.slice(0); - rb.unmap(); - - return new Float32Array(copy); - - */ - const rb = this.ensureReadback(byteLength); - - const encoder = this.device.createCommandEncoder(); - encoder.copyBufferToBuffer(srcBuffer, srcOffset, rb, 0, byteLength); - this.device.queue.submit([encoder.finish()]); - - await rb.mapAsync(GPUMapMode.READ, 0, byteLength); - const mapped = rb.getMappedRange(0, byteLength); - const copy = mapped.slice(0); - rb.unmap(); - - return new Float32Array(copy); - - } - - private ensureReadback(sizeBytes: number) { - sizeBytes = alignTo(sizeBytes, 256); - if (this.readback !== null && this.readback.size >= sizeBytes) { - return this.readback.buffer; - } - - if (this.readback) { - this.pendingDestroy.push(this.readback.buffer); - this.readback = null; - } - - const buffer = this.device.createBuffer({ - size: sizeBytes, - usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, - }); - - this.readback = { - buffer, - size: sizeBytes - }; - - return buffer; - } - - ones(shape: number[], name: string = "ones") { - const buf = this.getTensorBuffer( - name, - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - shape); - this.writeBufferF32(buf.buffer, new Float32Array(shape.reduce((a, b) => a * b, 1)).fill(1)); - return buf; - } - - /** - * Create a scoped tensor filled with ones. - */ - scopedOnes(shape: number[]) { - const buf = this.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - shape); - this.writeBufferF32(buf.buffer, new Float32Array(shape.reduce((a, b) => a * b, 1)).fill(1)); - return buf; - } - - /** - * Create a scoped tensor filled with zeros. - */ - scopedZeros(shape: number[]) { - const buf = this.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - shape); - this.writeBufferF32(buf.buffer, new Float32Array(shape.reduce((a, b) => a * b, 1)).fill(0)); - return buf; - } - - zeros(buf: Tensor) { - this.writeBufferF32(buf.buffer, new Float32Array(buf.size).fill(0)); - } - - async flushDestroyQueue() { - if (this.pendingDestroy.length === 0) { - return; - } - - await this.device.queue.onSubmittedWorkDone(); - for (const b of this.pendingDestroy) { - try { - b.destroy(); - } catch(e) { - console.error("Failed to destroy tensor buffer " + e); - } - } - this.pendingDestroy.length = 0; - } - - async destroyAll() { - // Wait for GPU to finish anything using these buffers. - await this.device.queue.onSubmittedWorkDone(); - - for (const {buffer} of this.tensors.values()) { - try { - buffer.destroy(); - } catch(e) { - console.error("Failed to destroy tensor buffer "+e); - } - } - this.tensors.clear(); - - if (this.readback) { - try { - this.readback.buffer.destroy(); - } catch(e) { - console.error("Failed to destroy readback buffer "+e); - } - this.readback = null; - } - - for (const b of this.pendingDestroy) { - try { - b.destroy(); - } catch(e) { - console.error("Failed to destroy scheduled to destroy tensor buffer "+e); - } - } - this.pendingDestroy.length = 0; - } -} diff --git a/src/tensor/kernel/BiasAddKernel.ts b/src/tensor/kernel/BiasAddKernel.ts deleted file mode 100644 index a03d38b..0000000 --- a/src/tensor/kernel/BiasAddKernel.ts +++ /dev/null @@ -1,123 +0,0 @@ -import {ceilDiv, Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {BiasAddBackward} from "../../autograd/backward/BiasAddBackward"; -import {KernelRegistry} from "./KernelRegistry"; - -/** - * Adds a 1D bias vector to each row of a 2D matrix. - * - * Input: [M, N] - * Bias: [N] (1D) or [1, N] (2D) - * Output: [M, N] where output[i,j] = input[i,j] + bias[j] - */ -export class BiasAddKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, BiasAddKernel.biasAddWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - input: Tensor, - bias: Tensor, - out?: Tensor, - ): Tensor { - - if (input.shape.length !== 2) { - throw new Error("BiasAdd: input must be 2D tensor"); - } - - // Accept bias as [N] or [1, N] - const biasN = bias.shape.length === 1 - ? bias.shape[0] - : (bias.shape.length === 2 && bias.shape[0] === 1 ? bias.shape[1] : -1); - - if (biasN === -1) { - throw new Error("BiasAdd: bias must be [N] or [1, N]"); - } - - if (input.shape[1] !== biasN) { - throw new Error(`BiasAdd: input columns (${input.shape[1]}) must match bias size (${biasN})`); - } - - const M = input.shape[0]; - const N = input.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: input.buffer}}, - {binding: 1, resource: {buffer: bias.buffer}}, - {binding: 2, resource: {buffer: out.buffer}}, - {binding: 3, resource: {buffer: this.paramsBuf}}, - ], - }); - - const wgX = ceilDiv(N, 16); - const wgY = ceilDiv(M, 16); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(wgX, wgY, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - // Autograd: track computation graph - if (input.requiresGradient || bias.requiresGradient) { - out.requiresGradient = true; - out.parents = [input, bias]; - out.gradFn = new BiasAddBackward([input, bias], this.kr!); - } - - return out; - } - - static biasAddWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var input : array; - @group(0) @binding(1) var bias : array; - @group(0) @binding(2) var output : array; - @group(0) @binding(3) var params : Params; - - @compute @workgroup_size(16, 16, 1) - fn main( - @builtin(global_invocation_id) gid : vec3, - ) { - let row : u32 = gid.y; - let col : u32 = gid.x; - - if (row < params.M && col < params.N) { - let idx : u32 = row * params.N + col; - output[idx] = input[idx] + bias[col]; - } - } - `; -} diff --git a/src/tensor/kernel/CrossEntropyKernel.ts b/src/tensor/kernel/CrossEntropyKernel.ts deleted file mode 100644 index 13ed2bc..0000000 --- a/src/tensor/kernel/CrossEntropyKernel.ts +++ /dev/null @@ -1,138 +0,0 @@ -import {Kernel, ceilDiv} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {SoftmaxCrossEntropyBackward} from "../../autograd/backward/SoftmaxCrossEntropyBackward"; -import {KernelRegistry} from "./KernelRegistry"; - -/** - * Cross Entropy Loss kernel (with built-in log-softmax). - * Computes: -sum(labels * log_softmax(logits)) per row. - * - * Input logits: raw logits [M, N] (NOT softmax output). - * Input labels: one-hot encoded [M, N]. - * Output: per-sample loss [M, 1]. - * - * Uses logsumexp for numerical stability. - */ -export class CrossEntropyKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, CrossEntropyKernel.xentropyWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - logits: Tensor, - labels: Tensor, - out?: Tensor, - ): Tensor { - if (logits.shape.length !== 2) { - throw new Error("CrossEntropy: logits must be 2D tensor"); - } - if (labels.shape.length !== 2) { - throw new Error("CrossEntropy: labels must be 2D tensor"); - } - if (logits.shape[0] !== labels.shape[0] || logits.shape[1] !== labels.shape[1]) { - throw new Error("CrossEntropy: logits and labels must have same shape"); - } - - const M = logits.shape[0]; - const N = logits.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, 1], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: logits.buffer}}, - {binding: 1, resource: {buffer: labels.buffer}}, - {binding: 2, resource: {buffer: out.buffer}}, - {binding: 3, resource: {buffer: this.paramsBuf}}, - ], - }); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(ceilDiv(M, 256), 1, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - // Autograd: track computation graph - // Use combined SoftmaxCrossEntropy backward for numerical stability - // Forward computes log-softmax internally; backward computes softmax(logits) - labels - if (logits.requiresGradient) { - out.requiresGradient = true; - out.parents = [logits]; - out.gradFn = new SoftmaxCrossEntropyBackward([logits, labels], this.kr!); - } - - return out; - } - - static xentropyWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var logits : array; - @group(0) @binding(1) var labels : array; // one-hot - @group(0) @binding(2) var loss : array; - @group(0) @binding(3) var params : Params; - - @compute @workgroup_size(256, 1, 1) - fn main(@builtin(global_invocation_id) gid : vec3) { - let row = gid.x; - if (row >= params.M) { return; } - - let N = params.N; - let base = row * N; - - // 1) max logit for stability - var m = logits[base]; - for (var i = 1u; i < N; i = i + 1u) { - let z = logits[base + i]; - if (z > m) { m = z; } - } - - // 2) logsumexp - var sumExp = 0.0; - for (var i = 0u; i < N; i = i + 1u) { - sumExp = sumExp + exp(logits[base + i] - m); - } - let logSumExp = log(sumExp) + m; - - // 3) cross entropy: -sum y_i * (z_i - logsumexp) - var ce = 0.0; - for (var i = 0u; i < N; i = i + 1u) { - let y = labels[base + i]; - let z = logits[base + i]; - ce = ce + y * (logSumExp - z); - } - - loss[row] = ce; - } - - `; -} diff --git a/src/tensor/kernel/DropoutKernel.ts b/src/tensor/kernel/DropoutKernel.ts deleted file mode 100644 index 9a093a7..0000000 --- a/src/tensor/kernel/DropoutKernel.ts +++ /dev/null @@ -1,119 +0,0 @@ -import {ceilDiv, Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {KernelRegistry} from "./KernelRegistry"; -import {DropoutBackward} from "../../autograd/backward/DropoutBackward"; - -/** - * Applies element-wise dropout by multiplying input with a pre-computed mask. - * - * Input: [M, N] - * Mask: [M, N] (values are 0 or scale, where scale = 1/(1-p)) - * Output: [M, N] where output[i,j] = input[i,j] * mask[i,j] - */ -export class DropoutKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, DropoutKernel.dropoutWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - input: Tensor, - mask: Tensor, - out?: Tensor, - ): Tensor { - - if (input.shape.length !== 2) { - throw new Error("Dropout: input must be 2D tensor"); - } - - if (mask.shape.length !== 2) { - throw new Error("Dropout: mask must be 2D tensor"); - } - - if (input.shape[0] !== mask.shape[0] || input.shape[1] !== mask.shape[1]) { - throw new Error("Dropout: input and mask shapes must match"); - } - - const M = input.shape[0]; - const N = input.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: input.buffer}}, - {binding: 1, resource: {buffer: mask.buffer}}, - {binding: 2, resource: {buffer: out.buffer}}, - {binding: 3, resource: {buffer: this.paramsBuf}}, - ], - }); - - const wgX = ceilDiv(N, 16); - const wgY = ceilDiv(M, 16); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(wgX, wgY, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - // Autograd: track computation graph - // Save mask for backward (same mask applied to gradients) - if (input.requiresGradient) { - out.requiresGradient = true; - out.parents = [input]; - out.gradFn = new DropoutBackward([mask], this.kr!); - } - - return out; - } - - static dropoutWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var input : array; - @group(0) @binding(1) var mask : array; - @group(0) @binding(2) var output : array; - @group(0) @binding(3) var params : Params; - - @compute @workgroup_size(16, 16, 1) - fn main( - @builtin(global_invocation_id) gid : vec3, - ) { - let row : u32 = gid.y; - let col : u32 = gid.x; - - if (row < params.M && col < params.N) { - let idx : u32 = row * params.N + col; - output[idx] = input[idx] * mask[idx]; - } - } - `; -} diff --git a/src/tensor/kernel/ElementwiseMulKernel.ts b/src/tensor/kernel/ElementwiseMulKernel.ts deleted file mode 100644 index 54e3d0b..0000000 --- a/src/tensor/kernel/ElementwiseMulKernel.ts +++ /dev/null @@ -1,107 +0,0 @@ -import {ceilDiv, Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {KernelRegistry} from "./KernelRegistry"; - -/** - * Element-wise multiplication (Hadamard product). - * - * Input A: [M, N] - * Input B: [M, N] - * Output: [M, N] where output[i,j] = A[i,j] * B[i,j] - * - * Used for gradient masking (ReLU backward, Dropout backward). - */ -export class ElementwiseMulKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, ElementwiseMulKernel.elemMulWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - a: Tensor, - b: Tensor, - out?: Tensor, - ): Tensor { - if (a.shape.length !== 2 || b.shape.length !== 2) { - throw new Error("ElementwiseMul: inputs must be 2D tensors"); - } - - if (a.shape[0] !== b.shape[0] || a.shape[1] !== b.shape[1]) { - throw new Error("ElementwiseMul: input shapes must match"); - } - - const M = a.shape[0]; - const N = a.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: a.buffer}}, - {binding: 1, resource: {buffer: b.buffer}}, - {binding: 2, resource: {buffer: out.buffer}}, - {binding: 3, resource: {buffer: this.paramsBuf}}, - ], - }); - - const wgX = ceilDiv(N, 16); - const wgY = ceilDiv(M, 16); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(wgX, wgY, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - return out; - } - - static elemMulWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var A : array; - @group(0) @binding(1) var B : array; - @group(0) @binding(2) var C : array; - @group(0) @binding(3) var params : Params; - - @compute @workgroup_size(16, 16, 1) - fn main(@builtin(global_invocation_id) gid : vec3) { - let col = gid.x; - let row = gid.y; - - if (row >= params.M || col >= params.N) { - return; - } - - let idx = row * params.N + col; - C[idx] = A[idx] * B[idx]; - } - `; -} diff --git a/src/tensor/kernel/InplaceAddKernel.ts b/src/tensor/kernel/InplaceAddKernel.ts deleted file mode 100644 index 150d08c..0000000 --- a/src/tensor/kernel/InplaceAddKernel.ts +++ /dev/null @@ -1,96 +0,0 @@ -import {ceilDiv, Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {KernelRegistry} from "./KernelRegistry"; - -/** - * In-place addition: target += source - * - * Target: [M, N] (modified in-place) - * Source: [M, N] - * - * Used in optimizers for: param += update - */ -export class InplaceAddKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - kr: KernelRegistry, - ) { - super(device, InplaceAddKernel.inplaceAddWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - target: Tensor, - source: Tensor, - ): void { - if (target.shape.length !== 2 || source.shape.length !== 2) { - throw new Error("InplaceAdd: inputs must be 2D tensors"); - } - - if (target.shape[0] !== source.shape[0] || target.shape[1] !== source.shape[1]) { - throw new Error("InplaceAdd: tensor shapes must match"); - } - - const M = target.shape[0]; - const N = target.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: target.buffer}}, - {binding: 1, resource: {buffer: source.buffer}}, - {binding: 2, resource: {buffer: this.paramsBuf}}, - ], - }); - - const wgX = ceilDiv(N, 16); - const wgY = ceilDiv(M, 16); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(wgX, wgY, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - } - - static inplaceAddWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var targetTensor : array; - @group(0) @binding(1) var source : array; - @group(0) @binding(2) var params : Params; - - @compute @workgroup_size(16, 16, 1) - fn main(@builtin(global_invocation_id) gid : vec3) { - let col = gid.x; - let row = gid.y; - - if (row >= params.M || col >= params.N) { - return; - } - - let idx = row * params.N + col; - targetTensor[idx] = targetTensor[idx] + source[idx]; - } - `; -} diff --git a/src/tensor/kernel/Kernel.ts b/src/tensor/kernel/Kernel.ts deleted file mode 100644 index 347d486..0000000 --- a/src/tensor/kernel/Kernel.ts +++ /dev/null @@ -1,34 +0,0 @@ -import type {KernelRegistry} from "./KernelRegistry"; - -export function ceilDiv(a: number, b: number) { - return Math.floor((a + b - 1) / b); -} - -/** - * This class has the internal state for the Tensor operations. - * It contains a compute pipeline and a shader module for each operation. - * - * For autograd support, kernels can access the full registry via `kr` - * after it's been set by KernelRegistry.initAutograd(). - */ -export abstract class Kernel { - - readonly pipeline: GPUComputePipeline; - readonly module: GPUShaderModule; - - constructor( - device: GPUDevice, - wgsl: string, - readonly kr: KernelRegistry - ) { - - this.module = device.createShaderModule({code: wgsl}); - this.pipeline = device.createComputePipeline({ - layout: "auto", - compute: { - module: this.module, - entryPoint: "main" - }, - }); - } -} diff --git a/src/tensor/kernel/KernelRegistry.ts b/src/tensor/kernel/KernelRegistry.ts deleted file mode 100644 index 380b092..0000000 --- a/src/tensor/kernel/KernelRegistry.ts +++ /dev/null @@ -1,65 +0,0 @@ -import {MatMulKernel} from "./MatMulKernel"; -import {MatAddKernel} from "./MatAddKernel"; -import {BiasAddKernel} from "./BiasAddKernel"; -import {RELUKernel} from "./RELUKernel"; -import {SoftmaxKernel} from "./SoftmaxKernel"; -import {CrossEntropyKernel} from "./CrossEntropyKernel"; -import {DropoutKernel} from "./DropoutKernel"; -import {TransposeKernel} from "./TransposeKernel"; -import {SumReduceKernel} from "./SumReduceKernel"; -import {ElementwiseMulKernel} from "./ElementwiseMulKernel"; -import {ReLUBackwardKernel} from "./ReLUBackwardKernel"; -import {SoftmaxCEBackwardKernel} from "./SoftmaxCEBackwardKernel"; -import {SoftmaxBackwardKernel} from "./SoftmaxBackwardKernel"; -import {ScalarMulKernel} from "./ScalarMulKernel"; -import {InplaceAddKernel} from "./InplaceAddKernel"; -import {SumAllKernel} from "./SumAllKernel"; -import {TensorManager} from "../TensorManager"; - -export class KernelRegistry { - - readonly matmul: MatMulKernel; - readonly matadd: MatAddKernel; - readonly biasadd: BiasAddKernel; - readonly relu: RELUKernel; - readonly softmax: SoftmaxKernel; - readonly crossEntropy: CrossEntropyKernel; - readonly dropout: DropoutKernel; - - // Autograd support kernels - readonly transpose: TransposeKernel; - readonly sumReduce: SumReduceKernel; - readonly elemMul: ElementwiseMulKernel; - readonly reluBackward: ReLUBackwardKernel; - readonly softmaxBackward: SoftmaxBackwardKernel; - readonly softmaxCEBackward: SoftmaxCEBackwardKernel; - - // Optimizer support kernels - readonly scalarMul: ScalarMulKernel; - readonly inplaceAdd: InplaceAddKernel; - readonly sumAll: SumAllKernel; - - constructor(device: GPUDevice, tm: TensorManager) { - this.matmul = new MatMulKernel(device, tm, this); - this.matadd = new MatAddKernel(device, tm, this); - this.biasadd = new BiasAddKernel(device, tm, this); - this.relu = new RELUKernel(device, tm, this); - this.softmax = new SoftmaxKernel(device, tm, this); - this.crossEntropy = new CrossEntropyKernel(device, tm, this); - this.dropout = new DropoutKernel(device, tm, this); - - // Autograd support kernels - this.transpose = new TransposeKernel(device, tm, this); - this.sumReduce = new SumReduceKernel(device, tm, this); - this.elemMul = new ElementwiseMulKernel(device, tm, this); - this.reluBackward = new ReLUBackwardKernel(device, tm, this); - this.softmaxBackward = new SoftmaxBackwardKernel(device, tm, this); - this.softmaxCEBackward = new SoftmaxCEBackwardKernel(device, tm, this); - - // Optimizer support kernels - this.scalarMul = new ScalarMulKernel(device, tm, this); - this.inplaceAdd = new InplaceAddKernel(device, tm, this); - this.sumAll = new SumAllKernel(device, tm, this); - } - -} diff --git a/src/tensor/kernel/MatAddKernel.ts b/src/tensor/kernel/MatAddKernel.ts deleted file mode 100644 index 565b213..0000000 --- a/src/tensor/kernel/MatAddKernel.ts +++ /dev/null @@ -1,111 +0,0 @@ -import {ceilDiv, Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {MatAddBackward} from "../../autograd/backward/MatAddBackward"; -import {KernelRegistry} from "./KernelRegistry"; - -export class MatAddKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, MatAddKernel.matAddWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - t0: Tensor, - t1: Tensor, - out?: Tensor, - ): Tensor { - - if ( - t0.shape.length !== 2 - || t1.shape.length !== 2 - || (out !== undefined && out.shape.length !== 2) - ) { - throw new Error("MatAdd: expected 2D tensors"); - } - - if (t0.shape[0] !== t1.shape[0] || t0.shape[1] !== t1.shape[1]) { - throw new Error("MatAdd: tensor shapes must match"); - } - - const M = t0.shape[0]; - const N = t0.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: t0.buffer}}, - {binding: 1, resource: {buffer: t1.buffer}}, - {binding: 2, resource: {buffer: out.buffer}}, - {binding: 3, resource: {buffer: this.paramsBuf}}, - ], - }); - - const wgX = ceilDiv(N, 16); - const wgY = ceilDiv(M, 16); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(wgX, wgY, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - // Autograd: track computation graph - if (t0.requiresGradient || t1.requiresGradient) { - out.requiresGradient = true; - out.parents = [t0, t1]; - out.gradFn = new MatAddBackward([t0, t1], this.kr!); - } - - return out; - } - - static matAddWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var A : array; - @group(0) @binding(1) var B : array; - @group(0) @binding(2) var C : array; - @group(0) @binding(3) var params : Params; - - @compute @workgroup_size(16, 16, 1) - fn main( - @builtin(global_invocation_id) gid : vec3, - ) { - let row : u32 = gid.y; - let col : u32 = gid.x; - - if (row < params.M && col < params.N) { - let idx : u32 = row * params.N + col; - C[idx] = A[idx] + B[idx]; - } - } - `; -} diff --git a/src/tensor/kernel/MatMulKernel.ts b/src/tensor/kernel/MatMulKernel.ts deleted file mode 100644 index c3db90e..0000000 --- a/src/tensor/kernel/MatMulKernel.ts +++ /dev/null @@ -1,162 +0,0 @@ -import {ceilDiv, Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {MatMulBackward} from "../../autograd/backward/MatMulBackward"; -import {KernelRegistry} from "./KernelRegistry"; - -export class MatMulKernel extends Kernel { - - private readonly params = new Uint32Array(4); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, MatMulKernel.matmulWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 16, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - t0: Tensor, - t1: Tensor, - out?: Tensor - ): Tensor { - if ( - t0.shape.length !== 2 - || t1.shape.length !== 2 - || (out !== undefined && out.shape.length !== 2) - ) { - throw new Error("MatMul: expected 2D tensors"); - } - - if (t0.shape[1] !== t1.shape[0]) { - throw new Error("MatMul: invalid dimensions"); - } - - const M = t0.shape[0]; - const K = t0.shape[1]; - const N = t1.shape[1]; - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N]); - - if (out.shape[0] !== M || out.shape[1] !== N) { - throw new Error("MatMul: invalid output dimensions"); - } - - this.params[0] = M; - this.params[1] = N; - this.params[2] = K; - - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: t0.buffer}}, - {binding: 1, resource: {buffer: t1.buffer}}, - {binding: 2, resource: {buffer: out.buffer}}, - {binding: 3, resource: {buffer: this.paramsBuf}}, - ], - }); - - // Dispatch - const wgX = ceilDiv(N, 16); - const wgY = ceilDiv(M, 16); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(wgX, wgY, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - // Autograd: track computation graph - if (t0.requiresGradient || t1.requiresGradient) { - out.requiresGradient = true; - out.parents = [t0, t1]; - out.gradFn = new MatMulBackward([t0, t1], this.kr!); - } - - return out; - } - - static matmulWGSL = ` - // C[M,N] = A[M,K] * B[K,N] (row-major flat buffers) - - struct Params { - M : u32, - N : u32, - K : u32, - _pad : u32, - }; - - @group(0) @binding(0) var A : array; - @group(0) @binding(1) var B : array; - @group(0) @binding(2) var C : array; - @group(0) @binding(3) var params : Params; - - const TILE : u32 = 16u; - - var tileA : array; - var tileB : array; - - @compute @workgroup_size(16, 16, 1) - fn main( - @builtin(global_invocation_id) gid : vec3, - @builtin(local_invocation_id) lid : vec3, - ) { - let row : u32 = gid.y; - let col : u32 = gid.x; - - let inBounds : bool = (row < params.M) && (col < params.N); - - let lidx : u32 = lid.y * TILE + lid.x; - - var acc : f32 = 0.0; - let numTiles : u32 = (params.K + TILE - 1u) / TILE; - - for (var t : u32 = 0u; t < numTiles; t = t + 1u) { - let kBase : u32 = t * TILE; - - // Load A tile element (or 0) - let aCol : u32 = kBase + lid.x; - if ((row < params.M) && (aCol < params.K)) { - tileA[lidx] = A[row * params.K + aCol]; - } else { - tileA[lidx] = 0.0; - } - - // Load B tile element (or 0) - let bRow : u32 = kBase + lid.y; - if ((bRow < params.K) && (col < params.N)) { - tileB[lidx] = B[bRow * params.N + col]; - } else { - tileB[lidx] = 0.0; - } - - workgroupBarrier(); - - for (var i : u32 = 0u; i < TILE; i = i + 1u) { - acc = acc + tileA[lid.y * TILE + i] * tileB[i * TILE + lid.x]; - } - - workgroupBarrier(); - } - - // Only write valid output elements - if (inBounds) { - C[row * params.N + col] = acc; - } - } - `; -} diff --git a/src/tensor/kernel/RELUKernel.ts b/src/tensor/kernel/RELUKernel.ts deleted file mode 100644 index b3016f8..0000000 --- a/src/tensor/kernel/RELUKernel.ts +++ /dev/null @@ -1,111 +0,0 @@ -import {ceilDiv, Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {ReLUBackward} from "../../autograd/backward/ReLUBackward"; -import {KernelRegistry} from "./KernelRegistry"; - -export class RELUKernel extends Kernel { - - static readonly RELU_OUTPUT = "relu_out"; - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, RELUKernel.reluWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - t0: Tensor, - out?: Tensor, - ): Tensor { - - if ( - t0.shape.length !== 2 - || (out !== undefined && out.shape.length !== 2) - ) { - throw new Error("RELU: expected 2D tensor"); - } - - const M = t0.shape[0]; - const N = t0.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: t0.buffer}}, - {binding: 1, resource: {buffer: out.buffer}}, - {binding: 2, resource: {buffer: this.paramsBuf}}, - ], - }); - - // Dispatch - const wgX = ceilDiv(N, 16); - const wgY = ceilDiv(M, 16); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(wgX, wgY, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - // Autograd: track computation graph - // Save input for backward (needed to compute mask X > 0) - if (t0.requiresGradient) { - out.requiresGradient = true; - out.parents = [t0]; - out.gradFn = new ReLUBackward([t0], this.kr!); - } - - return out; - } - - static reluWGSL = ` - // RELU: remove negative values from a tensor. - - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var A : array; - @group(0) @binding(1) var B : array; - @group(0) @binding(2) var params : Params; - - @compute @workgroup_size(16, 16, 1) - fn main( - @builtin(global_invocation_id) gid : vec3, - ) { - let row : u32 = gid.y; - let col : u32 = gid.x; - - let inBounds : bool = (row < params.M) && (col < params.N); - - // Only write valid output elements - if (inBounds) { - let f: f32 = A[row * params.N + col]; - B[row * params.N + col] = max(0f, f); - } - } - `; -} diff --git a/src/tensor/kernel/ReLUBackwardKernel.ts b/src/tensor/kernel/ReLUBackwardKernel.ts deleted file mode 100644 index 167e8bf..0000000 --- a/src/tensor/kernel/ReLUBackwardKernel.ts +++ /dev/null @@ -1,110 +0,0 @@ -import {ceilDiv, Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {KernelRegistry} from "./KernelRegistry"; - -/** - * ReLU backward kernel. - * - * Computes: gradInput = gradOutput * (input > 0) - * - * Input gradOutput: [M, N] - * Input savedInput: [M, N] (the original input to ReLU forward) - * Output gradInput: [M, N] - */ -export class ReLUBackwardKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, ReLUBackwardKernel.reluBackwardWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - gradOutput: Tensor, - savedInput: Tensor, - out?: Tensor, - ): Tensor { - if (gradOutput.shape.length !== 2 || savedInput.shape.length !== 2) { - throw new Error("ReLUBackward: inputs must be 2D tensors"); - } - - if (gradOutput.shape[0] !== savedInput.shape[0] || - gradOutput.shape[1] !== savedInput.shape[1]) { - throw new Error("ReLUBackward: shapes must match"); - } - - const M = gradOutput.shape[0]; - const N = gradOutput.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: gradOutput.buffer}}, - {binding: 1, resource: {buffer: savedInput.buffer}}, - {binding: 2, resource: {buffer: out.buffer}}, - {binding: 3, resource: {buffer: this.paramsBuf}}, - ], - }); - - const wgX = ceilDiv(N, 16); - const wgY = ceilDiv(M, 16); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(wgX, wgY, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - return out; - } - - static reluBackwardWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var gradOutput : array; - @group(0) @binding(1) var savedInput : array; - @group(0) @binding(2) var gradInput : array; - @group(0) @binding(3) var params : Params; - - @compute @workgroup_size(16, 16, 1) - fn main(@builtin(global_invocation_id) gid : vec3) { - let col = gid.x; - let row = gid.y; - - if (row >= params.M || col >= params.N) { - return; - } - - let idx = row * params.N + col; - // gradient flows through only where input > 0 - let mask = select(0.0, 1.0, savedInput[idx] > 0.0); - gradInput[idx] = gradOutput[idx] * mask; - } - `; -} diff --git a/src/tensor/kernel/ScalarMulKernel.ts b/src/tensor/kernel/ScalarMulKernel.ts deleted file mode 100644 index 1ff25ec..0000000 --- a/src/tensor/kernel/ScalarMulKernel.ts +++ /dev/null @@ -1,107 +0,0 @@ -import {ceilDiv, Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {KernelRegistry} from "./KernelRegistry"; - -/** - * Multiply a tensor by a scalar. - * - * Input: [M, N] - * Scalar: f32 - * Output: [M, N] where output[i,j] = input[i,j] * scalar - * - * Used in optimizers for: gradient * learning_rate - */ -export class ScalarMulKernel extends Kernel { - - private readonly params = new Float32Array(4); // M, N as u32, scalar as f32, padding - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - kr: KernelRegistry, - ) { - super(device, ScalarMulKernel.scalarMulWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 16, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - input: Tensor, - scalar: number, - out?: Tensor, - ): Tensor { - if (input.shape.length !== 2) { - throw new Error("ScalarMul: input must be 2D tensor"); - } - - const M = input.shape[0]; - const N = input.shape[1]; - - // Pack params: M, N as uint32 view, then scalar as float32 - const uintView = new Uint32Array(this.params.buffer); - uintView[0] = M; - uintView[1] = N; - this.params[2] = scalar; - - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: input.buffer}}, - {binding: 1, resource: {buffer: out.buffer}}, - {binding: 2, resource: {buffer: this.paramsBuf}}, - ], - }); - - const wgX = ceilDiv(N, 16); - const wgY = ceilDiv(M, 16); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(wgX, wgY, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - return out; - } - - static scalarMulWGSL = ` - struct Params { - M : u32, - N : u32, - scalar : f32, - _pad : u32, - }; - - @group(0) @binding(0) var input : array; - @group(0) @binding(1) var output : array; - @group(0) @binding(2) var params : Params; - - @compute @workgroup_size(16, 16, 1) - fn main(@builtin(global_invocation_id) gid : vec3) { - let col = gid.x; - let row = gid.y; - - if (row >= params.M || col >= params.N) { - return; - } - - let idx = row * params.N + col; - output[idx] = input[idx] * params.scalar; - } - `; -} diff --git a/src/tensor/kernel/SoftmaxBackwardKernel.ts b/src/tensor/kernel/SoftmaxBackwardKernel.ts deleted file mode 100644 index 80c663c..0000000 --- a/src/tensor/kernel/SoftmaxBackwardKernel.ts +++ /dev/null @@ -1,118 +0,0 @@ -import {Kernel, ceilDiv} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {KernelRegistry} from "./KernelRegistry"; - -/** - * Softmax backward kernel. - * - * Computes: dX = S * (dS - sum(dS * S, axis=1, keepdims=True)) - * - * where S = softmax output, dS = gradient of loss w.r.t. softmax output - * - * Input gradOutput: [M, N] (dL/dSoftmax) - * Input softmaxOut: [M, N] (softmax output from forward) - * Output gradInput: [M, N] (dL/dX) - */ -export class SoftmaxBackwardKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, SoftmaxBackwardKernel.softmaxBackwardWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - gradOutput: Tensor, - softmaxOut: Tensor, - out?: Tensor, - ): Tensor { - if (gradOutput.shape.length !== 2 || softmaxOut.shape.length !== 2) { - throw new Error("SoftmaxBackward: inputs must be 2D tensors"); - } - - if (gradOutput.shape[0] !== softmaxOut.shape[0] || - gradOutput.shape[1] !== softmaxOut.shape[1]) { - throw new Error("SoftmaxBackward: shapes must match"); - } - - const M = gradOutput.shape[0]; - const N = gradOutput.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: gradOutput.buffer}}, - {binding: 1, resource: {buffer: softmaxOut.buffer}}, - {binding: 2, resource: {buffer: out.buffer}}, - {binding: 3, resource: {buffer: this.paramsBuf}}, - ], - }); - - // One thread per row - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(ceilDiv(M, 256), 1, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - return out; - } - - static softmaxBackwardWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var gradOutput : array; - @group(0) @binding(1) var softmaxOut : array; - @group(0) @binding(2) var gradInput : array; - @group(0) @binding(3) var params : Params; - - @compute @workgroup_size(256, 1, 1) - fn main(@builtin(global_invocation_id) gid : vec3) { - let row = gid.x; - if (row >= params.M) { - return; - } - - let N = params.N; - let base = row * N; - - // Compute dot = sum(dS * S) for this row - var dot = 0.0; - for (var i = 0u; i < N; i = i + 1u) { - dot = dot + gradOutput[base + i] * softmaxOut[base + i]; - } - - // Compute gradInput = S * (dS - dot) - for (var i = 0u; i < N; i = i + 1u) { - let idx = base + i; - gradInput[idx] = softmaxOut[idx] * (gradOutput[idx] - dot); - } - } - `; -} diff --git a/src/tensor/kernel/SoftmaxCEBackwardKernel.ts b/src/tensor/kernel/SoftmaxCEBackwardKernel.ts deleted file mode 100644 index 7cb46a0..0000000 --- a/src/tensor/kernel/SoftmaxCEBackwardKernel.ts +++ /dev/null @@ -1,127 +0,0 @@ -import {ceilDiv, Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {KernelRegistry} from "./KernelRegistry"; - -/** - * Combined Softmax + CrossEntropy backward kernel. - * - * Computes: gradLogits = softmax(logits) - labels - * - * This is the gradient of CrossEntropy(Softmax(logits), labels) w.r.t. logits. - * Numerically stable and efficient as a single operation. - * - * Input logits: [M, N] (raw logits, NOT softmax output) - * Input labels: [M, N] (one-hot encoded) - * Output: [M, N] - */ -export class SoftmaxCEBackwardKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, SoftmaxCEBackwardKernel.softmaxCEBackwardWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - logits: Tensor, - labels: Tensor, - out?: Tensor, - ): Tensor { - if (logits.shape.length !== 2 || labels.shape.length !== 2) { - throw new Error("SoftmaxCEBackward: inputs must be 2D tensors"); - } - - if (logits.shape[0] !== labels.shape[0] || - logits.shape[1] !== labels.shape[1]) { - throw new Error("SoftmaxCEBackward: shapes must match"); - } - - const M = logits.shape[0]; - const N = logits.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: logits.buffer}}, - {binding: 1, resource: {buffer: labels.buffer}}, - {binding: 2, resource: {buffer: out.buffer}}, - {binding: 3, resource: {buffer: this.paramsBuf}}, - ], - }); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(ceilDiv(M, 256), 1, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - return out; - } - - static softmaxCEBackwardWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var logits : array; - @group(0) @binding(1) var labels : array; - @group(0) @binding(2) var gradLogits : array; - @group(0) @binding(3) var params : Params; - - @compute @workgroup_size(256, 1, 1) - fn main(@builtin(global_invocation_id) gid : vec3) { - let row = gid.x; - if (row >= params.M) { - return; - } - - let N = params.N; - let base = row * N; - - // 1) Find max for numerical stability - var maxVal = logits[base]; - for (var i = 1u; i < N; i = i + 1u) { - maxVal = max(maxVal, logits[base + i]); - } - - // 2) Compute exp(x - max) and sum - var sum = 0.0; - for (var i = 0u; i < N; i = i + 1u) { - let e = exp(logits[base + i] - maxVal); - gradLogits[base + i] = e; - sum = sum + e; - } - - // 3) Normalize to get softmax, then subtract labels - let invSum = 1.0 / sum; - for (var i = 0u; i < N; i = i + 1u) { - let softmax_i = gradLogits[base + i] * invSum; - gradLogits[base + i] = softmax_i - labels[base + i]; - } - } - `; -} diff --git a/src/tensor/kernel/SoftmaxKernel.ts b/src/tensor/kernel/SoftmaxKernel.ts deleted file mode 100644 index 3ea06ff..0000000 --- a/src/tensor/kernel/SoftmaxKernel.ts +++ /dev/null @@ -1,126 +0,0 @@ -import {Kernel, ceilDiv} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import type {KernelRegistry} from "./KernelRegistry"; -import {SoftmaxBackward} from "../../autograd/backward/SoftmaxBackward"; - -/** - * Per sample softmax. - * Expected to be used with small Tensors (e.g. MNIST N=10) - * - * This implementation is not optimized for large tensors. - */ -export class SoftmaxKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, SoftmaxKernel.softmaxWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - t0: Tensor, - out?: Tensor, - ): Tensor { - if (t0.shape.length !== 2) { - throw new Error("Softmax: expected 2D tensor"); - } - - const M = t0.shape[0]; - const N = t0.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [M, N], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: t0.buffer}}, - {binding: 1, resource: {buffer: out.buffer}}, - {binding: 2, resource: {buffer: this.paramsBuf}}, - ], - }); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(ceilDiv(M, 256), 1, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - // Autograd: track computation graph - // Save input for backward. Note: Softmax backward is typically - // handled via combined SoftmaxCrossEntropy for numerical stability. - if (t0.requiresGradient) { - out.requiresGradient = true; - out.parents = [t0]; - // Store both input and output for flexible backward computation - out.gradFn = new SoftmaxBackward([t0, out], this.kr!); - } - - return out; - } - - static softmaxWGSL = ` - // Softmax: exp(x_i - max) / sum(exp(x_j - max)) - // Applied per row for numerical stability. - - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var A : array; - @group(0) @binding(1) var B : array; - @group(0) @binding(2) var params : Params; - - @compute @workgroup_size(256, 1, 1) - fn main(@builtin(global_invocation_id) gid : vec3) { - let row = gid.x; - if (row >= params.M) { - return; - } - - let N = params.N; - let base = row * N; - - // Find max for numerical stability - var maxVal = A[base]; - for (var i = 1u; i < N; i = i + 1u) { - maxVal = max(maxVal, A[base + i]); - } - - // Compute exp(x - max) and sum - var sum = 0.0; - for (var i = 0u; i < N; i = i + 1u) { - let e = exp(A[base + i] - maxVal); - B[base + i] = e; - sum = sum + e; - } - - // Normalize - let invSum = 1.0 / sum; - for (var i = 0u; i < N; i = i + 1u) { - B[base + i] = B[base + i] * invSum; - } - } - `; -} diff --git a/src/tensor/kernel/SumAllKernel.ts b/src/tensor/kernel/SumAllKernel.ts deleted file mode 100644 index 1d17054..0000000 --- a/src/tensor/kernel/SumAllKernel.ts +++ /dev/null @@ -1,136 +0,0 @@ -import {Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {KernelRegistry} from "./KernelRegistry"; - -/** - * Sum all elements of a tensor to a scalar. - * - * Input: [M, N] (any 2D tensor) - * Output: [1, 1] containing sum of all elements - * - * Used for loss reduction: mean loss = sum(per_sample_loss) / batch_size - * - * Note: This is a simple single-pass implementation suitable for small tensors. - * For large tensors, a hierarchical reduction would be more efficient. - */ -export class SumAllKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - kr: KernelRegistry, - ) { - super(device, SumAllKernel.sumAllWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - input: Tensor, - out?: Tensor, - ): Tensor { - if (input.shape.length !== 2) { - throw new Error("SumAll: input must be 2D tensor"); - } - - const M = input.shape[0]; - const N = input.shape[1]; - const totalSize = M * N; - - if (totalSize > 65536) { - console.warn("SumAll: tensor size > 65536, consider using hierarchical reduction"); - } - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [1, 1], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: input.buffer}}, - {binding: 1, resource: {buffer: out.buffer}}, - {binding: 2, resource: {buffer: this.paramsBuf}}, - ], - }); - - // Single workgroup does the reduction - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(1, 1, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - return out; - } - - static sumAllWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var input : array; - @group(0) @binding(1) var output : array; - @group(0) @binding(2) var params : Params; - - var sharedData : array; - - @compute @workgroup_size(256, 1, 1) - fn main( - @builtin(local_invocation_id) lid : vec3, - ) { - let totalSize = params.M * params.N; - let tid = lid.x; - - // Each thread sums a strided portion of the input - var sum = 0.0; - var idx = tid; - while (idx < totalSize) { - sum = sum + input[idx]; - idx = idx + 256u; - } - - sharedData[tid] = sum; - workgroupBarrier(); - - // Parallel reduction in sharedData memory - if (tid < 128u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 128u]; } - workgroupBarrier(); - if (tid < 64u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 64u]; } - workgroupBarrier(); - if (tid < 32u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 32u]; } - workgroupBarrier(); - if (tid < 16u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 16u]; } - workgroupBarrier(); - if (tid < 8u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 8u]; } - workgroupBarrier(); - if (tid < 4u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 4u]; } - workgroupBarrier(); - if (tid < 2u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 2u]; } - workgroupBarrier(); - if (tid < 1u) { sharedData[tid] = sharedData[tid] + sharedData[tid + 1u]; } - workgroupBarrier(); - - // Thread 0 writes the result - if (tid == 0u) { - output[0] = sharedData[0]; - } - } - `; -} diff --git a/src/tensor/kernel/SumReduceKernel.ts b/src/tensor/kernel/SumReduceKernel.ts deleted file mode 100644 index 94cedf5..0000000 --- a/src/tensor/kernel/SumReduceKernel.ts +++ /dev/null @@ -1,99 +0,0 @@ -import {Kernel, ceilDiv} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {KernelRegistry} from "./KernelRegistry"; - -/** - * Sum reduction along axis 0 (columns). - * - * Input: [M, N] - * Output: [1, N] where output[0,j] = sum(input[i,j] for all i) - * - * Used for computing bias gradients: db = sum(dOut, axis=0) - */ -export class SumReduceKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, SumReduceKernel.sumReduceWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - input: Tensor, - out?: Tensor, - ): Tensor { - if (input.shape.length !== 2) { - throw new Error("SumReduce: input must be 2D tensor"); - } - - const M = input.shape[0]; - const N = input.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [1, N], - ); - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: input.buffer}}, - {binding: 1, resource: {buffer: out.buffer}}, - {binding: 2, resource: {buffer: this.paramsBuf}}, - ], - }); - - // One thread per column - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(ceilDiv(N, 256), 1, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - return out; - } - - static sumReduceWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var input : array; - @group(0) @binding(1) var output : array; - @group(0) @binding(2) var params : Params; - - @compute @workgroup_size(256, 1, 1) - fn main(@builtin(global_invocation_id) gid : vec3) { - let col = gid.x; - if (col >= params.N) { - return; - } - - var sum = 0.0; - for (var row = 0u; row < params.M; row = row + 1u) { - sum = sum + input[row * params.N + col]; - } - - output[col] = sum; - } - `; -} diff --git a/src/tensor/kernel/TransposeKernel.ts b/src/tensor/kernel/TransposeKernel.ts deleted file mode 100644 index 854a61d..0000000 --- a/src/tensor/kernel/TransposeKernel.ts +++ /dev/null @@ -1,103 +0,0 @@ -import {ceilDiv, Kernel} from "./Kernel"; -import {Tensor} from "../Tensor"; -import {TensorManager} from "../TensorManager"; -import {KernelRegistry} from "./KernelRegistry"; - -/** - * Transpose a 2D matrix. - * - * Input: [M, N] - * Output: [N, M] where output[j,i] = input[i,j] - */ -export class TransposeKernel extends Kernel { - - private readonly params = new Uint32Array(2); - private readonly paramsBuf: GPUBuffer; - - constructor( - readonly device: GPUDevice, - readonly tm: TensorManager, - readonly kr: KernelRegistry, - ) { - super(device, TransposeKernel.transposeWGSL, kr); - - this.paramsBuf = device.createBuffer({ - size: 8, - usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, - }); - } - - run( - input: Tensor, - out?: Tensor, - ): Tensor { - if (input.shape.length !== 2) { - throw new Error("Transpose: input must be 2D tensor"); - } - - const M = input.shape[0]; - const N = input.shape[1]; - - this.params[0] = M; - this.params[1] = N; - this.device.queue.writeBuffer(this.paramsBuf, 0, this.params); - - out = out ?? this.tm.getScopedTensor( - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, - [N, M], - ); - - if (out.shape[0] !== N || out.shape[1] !== M) { - throw new Error("Transpose: output shape must be [N, M]"); - } - - const bindGroup = this.device.createBindGroup({ - layout: this.pipeline.getBindGroupLayout(0), - entries: [ - {binding: 0, resource: {buffer: input.buffer}}, - {binding: 1, resource: {buffer: out.buffer}}, - {binding: 2, resource: {buffer: this.paramsBuf}}, - ], - }); - - const wgX = ceilDiv(N, 16); - const wgY = ceilDiv(M, 16); - - const encoder = this.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(this.pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(wgX, wgY, 1); - pass.end(); - - this.device.queue.submit([encoder.finish()]); - - return out; - } - - static transposeWGSL = ` - struct Params { - M : u32, - N : u32, - }; - - @group(0) @binding(0) var input : array; - @group(0) @binding(1) var output : array; - @group(0) @binding(2) var params : Params; - - @compute @workgroup_size(16, 16, 1) - fn main(@builtin(global_invocation_id) gid : vec3) { - let col : u32 = gid.x; // output column = input row - let row : u32 = gid.y; // output row = input column - - if (row >= params.M || col >= params.N) { - return; - } - - // input[row, col] -> output[col, row] - let inIdx = row * params.N + col; - let outIdx = col * params.M + row; - output[outIdx] = input[inIdx]; - } - `; -} diff --git a/train_and_test.html b/train_and_test.html index f4bca84..9adc056 100644 --- a/train_and_test.html +++ b/train_and_test.html @@ -230,6 +230,8 @@ } } + +

MNIST Train and Test

@@ -289,6 +291,5 @@

Prediction Errors

- diff --git a/tsconfig.json b/tsconfig.json deleted file mode 100644 index 6cfbd93..0000000 --- a/tsconfig.json +++ /dev/null @@ -1,10 +0,0 @@ -{ - "compilerOptions": { - "target": "ES2022", - "module": "ESNext", - "moduleResolution": "bundler", - "strict": true, - "types": ["@webgpu/types"] - }, - "include": ["src"] -} diff --git a/vite.config.ts b/vite.config.ts deleted file mode 100644 index d50e0f2..0000000 --- a/vite.config.ts +++ /dev/null @@ -1,26 +0,0 @@ -import { defineConfig } from 'vite' -import { resolve } from 'path' -import { viteStaticCopy } from 'vite-plugin-static-copy' - -export default defineConfig({ - base: '/MNIST/', - plugins: [ - viteStaticCopy({ - targets: [ - { - src: 'data', - dest: '.', - }, - ], - }), - ], - build: { - rollupOptions: { - input: { - main: resolve(__dirname, 'index.html'), - guess: resolve(__dirname, 'guess.html'), - train_and_test: resolve(__dirname, 'train_and_test.html'), - }, - }, - }, -})