diff --git a/.github/workflows/check.yml b/.github/workflows/check.yml index 60478de..fd69394 100644 --- a/.github/workflows/check.yml +++ b/.github/workflows/check.yml @@ -142,8 +142,11 @@ jobs: odin check wgpu/glfw-triangle -target:windows_amd64 $FLAGS odin check wgpu/glfw-triangle -target:js_wasm32 $FLAGS - odin check wgpu/sdl3-triangle -target:windows_amd64 $FLAGS - odin check wgpu/sdl3-triangle -target:js_wasm32 $FLAGS + odin check wgpu/sdl3/triangle -target:windows_amd64 $FLAGS + odin check wgpu/sdl3/triangle -target:js_wasm32 $FLAGS + + odin check wgpu/sdl3/game-of-life -target:windows_amd64 $FLAGS + odin check wgpu/sdl3/game-of-life -target:js_wasm32 $FLAGS odin check win32/game_of_life -target:windows_amd64 $FLAGS odin check win32/open_window -target:windows_amd64 $FLAGS diff --git a/wgpu/sdl3/game-of-life/README.md b/wgpu/sdl3/game-of-life/README.md new file mode 100644 index 0000000..454bc3c --- /dev/null +++ b/wgpu/sdl3/game-of-life/README.md @@ -0,0 +1,829 @@ +# Conway's Game of Life - Architecture Documentation + +## Overview + +This is a GPU-accelerated implementation of Conway's Game of Life using WebGPU, written in Odin. The application runs on both desktop (SDL3) and web (WebAssembly) platforms with a unified codebase. + +> **Note:** This is based on the SDL3 trinagle example and a port of the tutorial you can find here: https://codelabs.developers.google.com/your-first-webgpu-app#0. The original tutorial used JavaScript. + +**Key Features:** +- 32×32 cell grid (1,024 cells) +- 5 Hz simulation update rate (200ms intervals) +- Compute shader implementation of Game of Life rules +- Ping-pong buffer architecture for efficient GPU updates +- Simple, direct architecture following Odin idioms + +**Learning Goals:** +- Using Odin to implement something previously done in JavaScript +- Working with Odin WGPU bindings: https://pkg.odin-lang.org/vendor/wgpu/ +- Platform abstraction using Odin's build tags + +--- +![](img/conways-gol.png) + +--- + +## Project Structure + +The codebase follows a simple 3-file structure: + +```mermaid +graph LR + A[main.odin] --> B[os_desktop.odin
#+build !js] + A --> C[os_web.odin
#+build js] + + D[shaders/] --> E[compute.wgsl] + D --> F[render.wgsl] + + style A fill:#4CAF50,color:#fff + style B fill:#FF9800,color:#fff + style C fill:#2196F3,color:#fff + style D fill:#9C27B0,color:#fff +``` + +### File Overview + +| File | LOC | Purpose | +|------|-----|---------| +| `main.odin` | 466 | Core GPU logic, state management, rendering | +| `os_desktop.odin` | 104 | SDL3 platform layer (synchronous) | +| `os_web.odin` | 122 | WASM platform layer (asynchronous) | +| **Total** | **692** | **Complete application** | + +### Build Tags + +Odin's build tag system selects the appropriate platform file at compile time: + +- **Desktop build:** `odin build .` → Uses `os_desktop.odin` (excludes `os_web.odin`) +- **Web build:** `odin build . -target:js_wasm32` → Uses `os_web.odin` (excludes `os_desktop.odin`) + +--- + +## Architecture Overview + +```mermaid +flowchart TD + Start([Program Start]) --> Init[main.odin::main] + Init --> OSInit[os_init
Platform-specific] + OSInit --> GPUInit[init_gpu
Create instance & surface] + GPUInit --> ReqAdapter[os_request_adapter_and_device
Platform callbacks] + + ReqAdapter --> |Desktop: Sync| DesktopAdapter[on_adapter_sync] + ReqAdapter --> |Web: Async| WebAdapter[on_adapter callback] + + DesktopAdapter --> DesktopDevice[on_device_sync] + WebAdapter --> WebDevice[on_device callback] + + DesktopDevice --> Complete[complete_gpu_init] + WebDevice --> Complete + + Complete --> CreatePipelines[Create Pipelines
Render & Compute] + CreatePipelines --> CreateBuffers[Create Buffers
Vertex, Uniform, Storage] + CreateBuffers --> Run[os_run
Start event loop] + + Run --> |Desktop| DesktopLoop[SDL Event Loop
Calculate dt] + Run --> |Web| WebLoop[Browser step
Receives dt] + + DesktopLoop --> Frame[frame] + WebLoop --> Frame + + Frame --> Update[update_simulation
Accumulate dt] + Update --> Compute{Time for
update?} + Compute --> |Yes| ComputePass[run_compute_pass
Execute shader] + Compute --> |No| Skip[Skip compute] + ComputePass --> Render[Render Pass
Draw cells] + Skip --> Render + Render --> Present[Present Frame] + Present --> |Loop| DesktopLoop + Present --> |Loop| WebLoop + + style Init fill:#4CAF50,color:#fff + style Complete fill:#4CAF50,color:#fff + style DesktopLoop fill:#FF9800,color:#fff + style WebLoop fill:#2196F3,color:#fff + style Frame fill:#9C27B0,color:#fff +``` + +--- + +## main.odin - Core Application Logic + +The main file contains all GPU-related code and the application state. + +### Key Components + +#### 1. Configuration Constants +```odin +WIDTH :: 512 // Window width +HEIGHT :: 512 // Window height +GRID_SIZE :: 32 // 32×32 grid +WORKGROUP_SIZE :: 8 // Compute shader workgroup size +UPDATE_INTERVAL_MILLISECONDS :: 200.0 // 5 Hz update rate +``` + +#### 2. Application State +```odin +App_State :: struct { + ctx: runtime.Context + + // WebGPU core + instance, surface, adapter, device: wgpu.* + queue: wgpu.Queue + config: wgpu.SurfaceConfiguration + + // Pipelines & layouts + pipeline_layout, bind_group_layout: wgpu.* + render_module, compute_module: wgpu.ShaderModule + render_pipeline: wgpu.RenderPipeline + compute_pipeline: wgpu.ComputePipeline + + // Buffers + vertex_buffer, uniform_buffer: wgpu.Buffer + cell_state_storage: [2]wgpu.Buffer // Ping-pong + bind_groups: [2]wgpu.BindGroup + + // Simulation + step_index: u64 + did_compute, do_update: bool + last_tick: time.Tick + accumulator: time.Duration +} +``` + +#### 3. Initialization Flow + +```mermaid +sequenceDiagram + participant M as main() + participant OS as os_* + participant GPU as GPU Init + + M->>OS: os_init() + M->>GPU: init_gpu() + GPU->>GPU: Create instance + GPU->>OS: os_get_surface() + OS-->>GPU: Surface + GPU->>OS: os_request_adapter_and_device() + + Note over OS: Platform-specific async/sync + + OS->>GPU: complete_gpu_init(device) + GPU->>GPU: create_bind_group_layout() + GPU->>GPU: create_render_pipeline() + GPU->>GPU: create_compute_pipeline() + GPU->>GPU: create_buffers_and_bind_groups() + GPU->>OS: os_run() +``` + +#### 4. Pipeline Creation + +**Render Pipeline:** +- Vertex shader positions cell instances in grid +- Fragment shader colors cells (green/black) +- Reads from storage buffer to determine cell state + +**Compute Pipeline:** +- Implements Conway's Game of Life rules +- Reads from one storage buffer (current state) +- Writes to another storage buffer (next state) +- Runs at 5 Hz (every 200ms) + +```mermaid +graph LR + A[Vertex Buffer
Cell Quad] --> B[Render Pipeline] + C[Uniform Buffer
Grid Size] --> B + D[Storage Buffer A
Cell States] --> B + D --> E[Compute Pipeline] + E --> F[Storage Buffer B
New States] + F --> B + + style B fill:#4CAF50,color:#fff + style E fill:#9C27B0,color:#fff +``` + +#### 5. Buffer Architecture (Ping-Pong) + +```mermaid +graph TD + subgraph "Frame N" + A[Storage Buffer A
Current State] --> C[Compute Shader] + C --> B[Storage Buffer B
Next State] + B --> R1[Render Pipeline
Read from B] + end + + subgraph "Frame N+1" + B2[Storage Buffer B
Current State] --> C2[Compute Shader] + C2 --> A2[Storage Buffer A
Next State] + A2 --> R2[Render Pipeline
Read from A] + end + + R1 -.-> B2 + R2 -.-> A + + style C fill:#9C27B0,color:#fff + style C2 fill:#9C27B0,color:#fff +``` + +**Key insight:** +- Compute writes to buffer `(step + 1) % 2` +- Render reads from buffer `(step + 1) % 2` (the latest computed state) + +#### 6. Simulation Timing + +```mermaid +flowchart LR + A[Frame Called] --> B[update_simulation dt] + B --> C{accumulator >= 200ms?} + C --> |Yes| D[Set do_update = true
Reset accumulator] + C --> |No| E[Set do_update = false] + D --> F[run_compute_pass] + E --> F + F --> G{do_update?} + G --> |Yes| H[Execute Compute Shader
Increment step_index] + G --> |No| I[Skip compute] + H --> J[Render Pass] + I --> J +``` + +**Desktop timing:** `dt` calculated using `SDL.GetPerformanceCounter()` +**Web timing:** `dt` provided by browser (seconds since last frame) + +#### 7. Frame Rendering + +```odin +frame :: proc "c" (dt: f32) { + update_simulation(dt) + + // Acquire surface texture + // Create command encoder + + run_compute_pass(encoder) // Conditional based on timing + + // Render pass + // - Clear to dark blue + // - Bind render pipeline + // - Bind group: (step + 1) % 2 (read latest) + // - Draw instances: GRID_SIZE * GRID_SIZE + + // Submit and present + + if did_compute { + step_index += 1 + } +} +``` + +--- + +## os_desktop.odin - SDL3 Platform Layer + +Provides synchronous initialization and blocking event loop for desktop platforms. + +### Key Components + +```mermaid +flowchart TD + A[os_init] --> B[SDL.Init] + B --> C[SDL.CreateWindow] + + D[os_get_surface] --> E[sdl3glue.GetSurface] + + F[os_get_framebuffer_size] --> G[SDL.GetWindowSizeInPixels] + + H[os_request_adapter_and_device] --> I[wgpu.InstanceRequestAdapter
callback: on_adapter_sync] + I --> J[on_adapter_sync
fires immediately] + J --> K[wgpu.AdapterRequestDevice
callback: on_device_sync] + K --> L[on_device_sync
fires immediately] + L --> M[complete_gpu_init] + + N[os_run] --> O[SDL Event Loop] + O --> P[Calculate dt
SDL.GetPerformanceCounter] + P --> Q[SDL.PollEvent] + Q --> R{Event Type} + R --> |QUIT| S[Exit] + R --> |KEY_DOWN ESCAPE| S + R --> |WINDOW_RESIZED| T[resize] + R --> |Continue| U[frame dt] + T --> U + U --> P + S --> V[cleanup
SDL.DestroyWindow
SDL.Quit] + + style A fill:#FF9800,color:#fff + style N fill:#FF9800,color:#fff + style J fill:#4CAF50,color:#fff + style L fill:#4CAF50,color:#fff +``` + +### Synchronous Callbacks + +On desktop, WebGPU callbacks fire **immediately** (synchronously): + +```odin +wgpu.InstanceRequestAdapter(...) +// Callback fires before this line executes +// adapter is already available +``` + +### Event Loop + +```odin +os_run :: proc() { + last := SDL.GetPerformanceCounter() + + for running { + now = SDL.GetPerformanceCounter() + dt = f32((now - last) * 1000) / f32(SDL.GetPerformanceFrequency()) + last = now + + for SDL.PollEvent(&event) { + #partial switch event.type { + case .QUIT: running = false + case .KEY_DOWN: + if event.key.scancode == .ESCAPE { + running = false + } + case .WINDOW_RESIZED: resize() + } + } + + frame(dt) // dt in milliseconds + } +} +``` + +--- + +## os_web.odin - WASM Platform Layer + +Provides asynchronous initialization and browser-driven event loop for web platforms. + +### Key Components + +```mermaid +flowchart TD + A[os_init] --> B[js.add_window_event_listener
Resize] + + C[os_get_surface] --> D[wgpu.InstanceCreateSurface
Canvas selector: #wgpu-canvas] + + E[os_get_framebuffer_size] --> F[js.get_bounding_client_rect
body] + F --> G[Apply device_pixel_ratio] + + H[os_request_adapter_and_device] --> I[wgpu.InstanceRequestAdapter
inline callback: on_adapter] + I -.Async.-> J[on_adapter
fires when ready] + J --> K[wgpu.AdapterRequestDevice
inline callback: on_device] + K -.Async.-> L[on_device
fires when ready] + L --> M[complete_gpu_init] + + N[os_run] --> O[Set device_ready = true] + + P[step dt
@export] --> Q{device_ready?} + Q --> |No| R[return true] + Q --> |Yes| S[frame dt] + S --> T[return true] + + U[size_callback] --> V{device_ready?} + V --> |Yes| W[resize] + V --> |No| X[return] + + style A fill:#2196F3,color:#fff + style N fill:#2196F3,color:#fff + style P fill:#2196F3,color:#fff + style J fill:#4CAF50,color:#fff + style L fill:#4CAF50,color:#fff +``` + +### Asynchronous Callbacks + +On web, WebGPU callbacks fire **asynchronously** (when browser completes operation): + +```odin +wgpu.InstanceRequestAdapter(...) +// Function returns immediately +// Callback fires later (100-500ms typical) +``` + +**Critical:** Callbacks must be defined **inline** in the same scope for WASM: + +```odin +os_request_adapter_and_device :: proc() { + wgpu.InstanceRequestAdapter( + state.instance, + &{compatibleSurface = state.surface}, + {callback = on_adapter}, + ) + + // Define callback inline + on_adapter :: proc "c" (...) { + context = state.ctx + // ... adapter handling + wgpu.AdapterRequestDevice(..., {callback = on_device}) + + // Nested inline callback + on_device :: proc "c" (...) { + context = state.ctx + // ... device handling + complete_gpu_init(device) + } + } +} +``` + +### Browser Integration + +```mermaid +sequenceDiagram + participant B as Browser + participant W as WASM Module + participant G as GPU + + B->>W: Load & Initialize + W->>G: Request Adapter + Note over W,G: Async - returns immediately + + G-->>W: on_adapter callback + W->>G: Request Device + Note over W,G: Async - returns immediately + + G-->>W: on_device callback + W->>W: complete_gpu_init + W->>W: os_run (set ready flag) + + loop 60 FPS + B->>W: step(dt) + W->>W: frame(dt) + W-->>B: return true + end +``` + +### Exported Functions + +```odin +@(export) +step :: proc(dt: f32) -> bool { + context = state.ctx + + if !device_ready { + return true // Still initializing + } + + frame(dt) // dt in seconds + return true +} + +@(fini) +cleanup_on_exit :: proc "contextless" () { + cleanup() + js.remove_window_event_listener(.Resize, nil, size_callback) +} +``` + +--- + +## Shader Architecture + +### compute.wgsl - Game of Life Rules + +```mermaid +flowchart LR + A[Invocation ID
x, y] --> B[Read 8 Neighbors
from cellStateIn] + B --> C[Count Active
Neighbors] + C --> D{Current Cell
Active?} + D --> |Yes| E{neighbors == 2
or 3?} + D --> |No| F{neighbors == 3?} + E --> |Yes| G[Write 1
cellStateOut] + E --> |No| H[Write 0
cellStateOut] + F --> |Yes| G + F --> |No| H + + style C fill:#9C27B0,color:#fff + style G fill:#4CAF50,color:#fff + style H fill:#f44336,color:#fff +``` + +**Workgroup Configuration:** +- Size: 8×8 (64 threads per workgroup) +- Dispatch: 4×4 workgroups for 32×32 grid +- Total: 1,024 parallel executions + +**Bindings:** +- `@binding(0)`: Uniform buffer (grid size) +- `@binding(1)`: Read-only storage (current cell states) +- `@binding(2)`: Storage (output cell states) + +### render.wgsl - Cell Visualization + +```mermaid +flowchart TD + A[Instance ID] --> B[Calculate Grid Position
x = id % GRID_SIZE
y = id / GRID_SIZE] + B --> C[Read Cell State
cellStateIn instance] + C --> D[Scale Vertex Position
by 1/GRID_SIZE] + D --> E[Translate to Grid Cell
offset by x, y] + E --> F[Output Position] + + G[Fragment Shader] --> H{Cell Active?} + H --> |Yes| I[Green: 0, 0.6, 0] + H --> |No| J[Black: 0, 0, 0] + + style F fill:#4CAF50,color:#fff + style I fill:#4CAF50,color:#fff + style J fill:#424242,color:#fff +``` + +**Vertex Shader:** +- Input: Cell quad vertices (-0.8 to 0.8) +- Instance rendering: Draw GRID_SIZE² instances +- Each instance represents one cell + +**Fragment Shader:** +- Simple color selection based on cell state +- Active cells: Green `(0, 0.6, 0)` +- Inactive cells: Black `(0, 0, 0)` + +--- + +## Platform Differences + +### Desktop (SDL3) + +| Aspect | Implementation | +|--------|----------------| +| **Windowing** | SDL3 native window | +| **Event Loop** | Blocking `SDL.PollEvent()` | +| **Frame Timing** | Manual via `SDL.GetPerformanceCounter()` | +| **WebGPU Init** | Synchronous callbacks | +| **Delta Time** | Milliseconds (calculated) | +| **Exit Handling** | Window close or Escape key | + +### Web (WASM) + +| Aspect | Implementation | +|--------|----------------| +| **Windowing** | HTML5 Canvas (`#wgpu-canvas`) | +| **Event Loop** | Browser `requestAnimationFrame` calls `step()` | +| **Frame Timing** | Provided by browser | +| **WebGPU Init** | Asynchronous callbacks (inline) | +| **Delta Time** | Seconds (from browser) | +| **Exit Handling** | Browser tab close | + +### Callback Timing Diagram + +```mermaid +sequenceDiagram + participant C as Caller + participant W as WebGPU + participant CB as Callback + + Note over C,CB: DESKTOP (Synchronous) + C->>W: RequestAdapter + W->>CB: Callback fires immediately + CB->>C: Returns to caller + Note over C: Adapter ready here + + Note over C,CB: WEB (Asynchronous) + C->>W: RequestAdapter + W-->>C: Returns immediately + Note over C: Adapter NOT ready yet + Note over W: Browser processes... + W->>CB: Callback fires later + Note over CB: Adapter ready NOW +``` + +--- + +## Build System + +### Desktop Build + +```bash +odin build . -out:game-of-life -vet -strict-style -vet-tabs -disallow-do -warnings-as-errors +./sdl3-Game-of-life +``` + +**What happens:** +1. Compiler includes `main.odin` and `os_desktop.odin` +2. `os_web.odin` excluded via `#+build js` tag +3. Links SDL3 and WebGPU native libraries +4. Creates native executable + +### Web Build + +```bash +odin build . -target:js_wasm32 -out:web/game_of_life.wasm +``` + +**What happens:** +1. Compiler includes `main.odin` and `os_web.odin` +2. `os_desktop.odin` excluded via `#+build !js` tag +3. Generates WebAssembly module +4. Exports `step()` function for browser + +### Build Tags Explanation + +```odin +// os_desktop.odin +#+build !js // Include when NOT building for JavaScript +package main + +// os_web.odin +#+build js // Include ONLY when building for JavaScript +package main +``` + +Build tags are **file-level** in Odin - you cannot use them inline within functions. + +--- + +## Data Flow + +### Initialization Flow + +```mermaid +graph TD + A[main] --> B[os_init] + B --> C[init_gpu] + C --> D[Create instance] + D --> E[os_get_surface] + E --> F[os_request_adapter_and_device] + + F --> G{Platform?} + G --> |Desktop| H[Sync callbacks] + G --> |Web| I[Async callbacks] + + H --> J[complete_gpu_init] + I --> J + + J --> K[Configure surface] + K --> L[create_bind_group_layout] + L --> M[create_render_pipeline] + M --> N[create_compute_pipeline] + N --> O[create_buffers_and_bind_groups] + O --> P[Initialize with random cells] + P --> Q[os_run] + + style A fill:#4CAF50,color:#fff + style J fill:#4CAF50,color:#fff + style Q fill:#FF9800,color:#fff +``` + +### Frame Flow + +```mermaid +graph TD + A[frame dt] --> B[update_simulation] + B --> C[accumulator += dt] + C --> D{accumulator >= 200ms?} + D --> |Yes| E[do_update = true
accumulator = 0] + D --> |No| F[do_update = false] + + E --> G[Get surface texture] + F --> G + G --> H[Create command encoder] + + H --> I[run_compute_pass] + I --> J{do_update?} + J --> |Yes| K[Begin compute pass] + J --> |No| L[Skip] + + K --> M[Set compute pipeline] + M --> N[Bind group: step % 2] + N --> O[Dispatch workgroups] + O --> P[did_compute = true] + + P --> Q[Begin render pass] + L --> Q + Q --> R[Clear to dark blue] + R --> S[Set render pipeline] + S --> T[Bind group: step+1 % 2] + T --> U[Set vertex buffer] + U --> V[Draw GRID_SIZE²] + V --> W[End render pass] + + W --> X[Submit & Present] + X --> Y{did_compute?} + Y --> |Yes| Z[step_index++] + Y --> |No| AA[Keep step_index] + + style E fill:#4CAF50,color:#fff + style K fill:#9C27B0,color:#fff + style P fill:#9C27B0,color:#fff + style Z fill:#4CAF50,color:#fff +``` + +--- + +## Performance Characteristics + +### GPU Workload + +| Operation | Frequency | GPU Load | +|-----------|-----------|----------| +| **Compute Shader** | 5 Hz | 1,024 threads (32×32 grid) | +| **Render Pass** | 60 FPS | 1,024 instances, 6 vertices each | +| **Buffer Updates** | 0 Hz | No CPU→GPU transfers after init | + +### Memory Usage + +| Resource | Size | Count | Total | +|----------|------|-------|-------| +| **Vertex Buffer** | 48 bytes | 1 | 48 B | +| **Uniform Buffer** | 8 bytes | 1 | 8 B | +| **Storage Buffers** | 4 KB | 2 | 8 KB | +| **Bind Groups** | - | 2 | - | +| **Pipelines** | - | 2 | - | + +**Total GPU memory:** ~8 KB (excluding shader bytecode and pipeline state) + +### CPU Load + +| Platform | Per Frame | Notes | +|----------|-----------|-------| +| **Desktop** | Minimal | Event polling, dt calculation | +| **Web** | Minimal | Browser calls `step()` | + +**Key insight:** After initialization, all simulation logic runs on GPU. CPU only submits command buffers. + +--- + +## Troubleshooting + +### Desktop Issues + +**Problem:** Window opens but cells don't update +**Solution:** Check that delta time is being calculated (not passing `0`) + +**Problem:** Window doesn't open +**Solution:** Ensure SDL3 is installed and linked correctly + +**Problem:** WebGPU errors +**Solution:** Check that your GPU supports WebGPU (Metal on macOS, D3D12 on Windows, Vulkan on Linux) + +### Web Issues + +**Problem:** Black screen, no errors +**Solution:** Check browser console - likely async callbacks not firing + +**Problem:** "WebGPU not supported" +**Solution:** Use Chrome/Edge with WebGPU enabled, serve over `http://127.0.0.1` (not `file://` or IPV6 [::]: ) + +**Problem:** Callbacks never fire +**Solution:** Ensure callbacks are defined **inline** in the same scope as registration + +**Problem:** Canvas size wrong +**Solution:** Check `device_pixel_ratio` and canvas CSS + +### Common Issues + +**Problem:** Compilation errors with build tags +**Solution:** Build tags must be at top of file, before `package` declaration + +**Problem:** Linking errors +**Solution:** Ensure `vendor:wgpu`, `vendor:sdl3` are available in your Odin installation + +--- + +## Future Enhancements + +Potential improvements while maintaining simplicity: + +1. **Interactive Controls** + - Mouse click to toggle cells + - Space bar to pause/resume + - R key to randomize grid (matchin the Win32 example) + +2. **Adjustable Parameters** + - Grid size selection (16×16, 32×32, 64×64) + - Update rate slider + - Color themes + +3. **Patterns** + - Load predefined patterns (glider, blinker, etc.) + - Save/load grid states + +4. **Performance** + - Larger grids (128×128, 256×256) + - Multiple compute passes per frame + - Benchmarking mode + +**Constraint:** Keep 3-file structure, avoid over-abstraction + +--- + +## References + +- **Original Tutorial:** https://codelabs.developers.google.com/your-first-webgpu-app +- **Odin Language:** https://odin-lang.org +- **Odin WGPU Bindings:** https://pkg.odin-lang.org/vendor/wgpu/ +- **SDL3:** https://wiki.libsdl.org/SDL3/ +- **WebGPU Spec:** https://gpuweb.github.io/gpuweb/ +- **Conway's Game of Life:** https://en.wikipedia.org/wiki/Conway%27s_Game_of_Life + +--- + +## License + +This code is part of the Odin examples repository. + +--- + +*Last updated: October 13, 2025* + diff --git a/wgpu/sdl3/game-of-life/build_web.bat b/wgpu/sdl3/game-of-life/build_web.bat new file mode 100644 index 0000000..50d77df --- /dev/null +++ b/wgpu/sdl3/game-of-life/build_web.bat @@ -0,0 +1,13 @@ +REM NOTE: changing this requires changing the same values in the `web/index.html`. +set INITIAL_MEMORY_PAGES=2000 +set MAX_MEMORY_PAGES=65536 + +set PAGE_SIZE=65536 +set /a INITIAL_MEMORY_BYTES=%INITIAL_MEMORY_PAGES% * %PAGE_SIZE% +set /a MAX_MEMORY_BYTES=%MAX_MEMORY_PAGES% * %PAGE_SIZE% + +call odin.exe build . -target:js_wasm32 -out:web/game_of_life.wasm -o:size -extra-linker-flags:"--export-table --import-memory --initial-memory=%INITIAL_MEMORY_BYTES% --max-memory=%MAX_MEMORY_BYTES%" + +for /f "delims=" %%i in ('odin.exe root') do set "ODIN_ROOT=%%i" +copy "%ODIN_ROOT%\vendor\wgpu\wgpu.js" "web\wgpu.js" +copy "%ODIN_ROOT%\core\sys\wasm\js\odin.js" "web\odin.js" diff --git a/wgpu/sdl3/game-of-life/build_web.sh b/wgpu/sdl3/game-of-life/build_web.sh new file mode 100755 index 0000000..bd5f419 --- /dev/null +++ b/wgpu/sdl3/game-of-life/build_web.sh @@ -0,0 +1,20 @@ +#!/usr/bin/env bash + +set -euxo pipefail + +# NOTE: changing this requires changing the same values in the `web/index.html`. +INITIAL_MEMORY_PAGES=2000 +MAX_MEMORY_PAGES=65536 + +INITIAL_MEMORY_BYTES=$(expr $INITIAL_MEMORY_PAGES \* $MAX_MEMORY_PAGES) +MAX_MEMORY_BYTES=$(expr $MAX_MEMORY_PAGES \* $MAX_MEMORY_PAGES) + +ODIN_ROOT=$(odin root) +ODIN_JS="$ODIN_ROOT/core/sys/wasm/js/odin.js" +WGPU_JS="$ODIN_ROOT/vendor/wgpu/wgpu.js" + +odin build . -target:js_wasm32 -out:web/game_of_life.wasm -o:size \ + -extra-linker-flags:"--export-table --import-memory --initial-memory=$INITIAL_MEMORY_BYTES --max-memory=$MAX_MEMORY_BYTES" + +cp $ODIN_JS web/odin.js +cp $WGPU_JS web/wgpu.js diff --git a/wgpu/sdl3/game-of-life/img/conways-gol.png b/wgpu/sdl3/game-of-life/img/conways-gol.png new file mode 100644 index 0000000..235ad15 Binary files /dev/null and b/wgpu/sdl3/game-of-life/img/conways-gol.png differ diff --git a/wgpu/sdl3/game-of-life/main.odin b/wgpu/sdl3/game-of-life/main.odin new file mode 100644 index 0000000..23d95c1 --- /dev/null +++ b/wgpu/sdl3/game-of-life/main.odin @@ -0,0 +1,469 @@ +// Conway's Game of Life - WebGPU/ODIN Implementation +// Ported to SDL3/WASM Odin from : https://codelabs.developers.google.com/your-first-webgpu-app#0 +// by Jason Coleman, 2025 +// ============================================== +// GPU-accelerated cellular automaton running on both desktop (SDL3) and web (WASM). +// +// Architecture: +// - Single file contains all GPU initialisation, pipelines, buffers, simulation, and rendering +// - Platform-specific code separated in os_desktop.odin and os_web.odin (selected by build tags) +// - Compute shader implements Game of Life rules with ping-pong buffers +// - Fixed 5 Hz simulation update rate (200ms intervals) +// +// Build: +// Desktop: odin build . -out:game-of-life -vet -strict-style -vet-tabs -disallow-do -warnings-as-errors +// Web: odin build . -target:js_wasm32 -out:web/game_of_life.wasm +// +package main + +import "base:runtime" +import "core:fmt" +import "core:math/rand" +import "core:time" +import "vendor:wgpu" + +//============================================================================= +// Configuration Constants +//============================================================================= + +WIDTH :: 512 +HEIGHT :: 512 +GRID_SIZE :: 32 +WORKGROUP_SIZE :: 8 +UPDATE_INTERVAL_MILLISECONDS :: 200.0 // 5 Hz update rate + +//============================================================================= +// Application State +//============================================================================= + +App_State :: struct { + // Context + ctx: runtime.Context, + + // WebGPU core + instance: wgpu.Instance, + surface: wgpu.Surface, + adapter: wgpu.Adapter, + device: wgpu.Device, + queue: wgpu.Queue, + config: wgpu.SurfaceConfiguration, + + // Pipelines & layouts + pipeline_layout: wgpu.PipelineLayout, + bind_group_layout: wgpu.BindGroupLayout, + render_module: wgpu.ShaderModule, + compute_module: wgpu.ShaderModule, + render_pipeline: wgpu.RenderPipeline, + compute_pipeline: wgpu.ComputePipeline, + + // Buffers + vertex_buffer: wgpu.Buffer, + vertex_count: u32, + vertex_buffer_size: u64, + uniform_buffer: wgpu.Buffer, + cell_state_storage: [2]wgpu.Buffer, + bind_groups: [2]wgpu.BindGroup, + + // Simulation state + step_index: u64, + did_compute: bool, + do_update: bool, + + // Timing + last_tick: time.Tick, + accumulator: time.Duration, +} + +state: App_State + +//============================================================================= +// Entry Point +//============================================================================= + +main :: proc() { + state.ctx = context + os_init() + init_gpu() +} + +//============================================================================= +// GPU Initialisation +//============================================================================= + +init_gpu :: proc() { + state.instance = wgpu.CreateInstance(nil) + if state.instance == nil { + panic("WebGPU is not supported") + } + + state.surface = os_get_surface(state.instance) + + // Platform-specific initialisation (os_desktop.odin or os_web.odin) + os_request_adapter_and_device() +} + +// Called by platform code after device is acquired +complete_gpu_init :: proc(device: wgpu.Device) { + state.device = device + state.queue = wgpu.DeviceGetQueue(state.device) + + // Configure surface + width, height := os_get_framebuffer_size() + state.config = wgpu.SurfaceConfiguration { + device = state.device, + usage = {.RenderAttachment}, + format = .BGRA8Unorm, + width = width, + height = height, + presentMode = .Fifo, + alphaMode = .Opaque, + } + wgpu.SurfaceConfigure(state.surface, &state.config) + + // Create all GPU resources + create_bind_group_layout() + create_render_pipeline() + create_compute_pipeline() + create_buffers_and_bind_groups() + + // Initialize timing + state.last_tick = time.tick_now() + state.accumulator = 0 + + // Start platform loop + os_run() +} + +on_device_error :: proc "c" ( + device: ^wgpu.Device, + errorType: wgpu.ErrorType, + message: string, + userdata1, userdata2: rawptr, +) { + context = state.ctx + panic(message) +} + +//============================================================================= +// Pipeline Creation +//============================================================================= + +create_bind_group_layout :: proc() { + // Binding 0: uniform vec2f (grid size) + b0 := wgpu.BindGroupLayoutEntry { + binding = 0, + visibility = {.Vertex, .Fragment, .Compute}, + buffer = {type = .Uniform, minBindingSize = size_of(f32) * 2}, + } + + // Binding 1: read-only storage (cell state input) + b1 := wgpu.BindGroupLayoutEntry { + binding = 1, + visibility = {.Vertex, .Compute}, + buffer = {type = .ReadOnlyStorage}, + } + + // Binding 2: storage (cell state output - compute only) + b2 := wgpu.BindGroupLayoutEntry { + binding = 2, + visibility = {.Compute}, + buffer = {type = .Storage}, + } + + entries := [3]wgpu.BindGroupLayoutEntry{b0, b1, b2} + + state.bind_group_layout = wgpu.DeviceCreateBindGroupLayout( + state.device, + &{label = "Cell State Bind Group Layout", entryCount = 3, entries = &entries[0]}, + ) + + layouts := [1]wgpu.BindGroupLayout{state.bind_group_layout} + state.pipeline_layout = wgpu.DeviceCreatePipelineLayout( + state.device, + &{bindGroupLayoutCount = 1, bindGroupLayouts = &layouts[0]}, + ) +} + +create_render_pipeline :: proc() { + // Vertex buffer for cell quad + verts := []f32{-0.8, -0.8, 0.8, -0.8, 0.8, 0.8, -0.8, -0.8, 0.8, 0.8, -0.8, 0.8} + + state.vertex_count = u32(len(verts) / 2) + state.vertex_buffer_size = u64(len(verts) * size_of(f32)) + + state.vertex_buffer = wgpu.DeviceCreateBuffer( + state.device, + &{label = "Cell Quad Vertices", usage = {.Vertex, .CopyDst}, size = state.vertex_buffer_size}, + ) + + wgpu.QueueWriteBuffer(state.queue, state.vertex_buffer, 0, raw_data(verts), len(verts) * size_of(f32)) + + // Load shader + state.render_module = wgpu.DeviceCreateShaderModule( + state.device, + &{ + label = "Render Shader", + nextInChain = &wgpu.ShaderSourceWGSL{sType = .ShaderSourceWGSL, code = #load("shaders/render.wgsl")}, + }, + ) + + // Create pipeline + state.render_pipeline = wgpu.DeviceCreateRenderPipeline( + state.device, + &{ + label = "Cell Render Pipeline", + layout = state.pipeline_layout, + vertex = { + module = state.render_module, + entryPoint = "vertexMain", + bufferCount = 1, + buffers = &wgpu.VertexBufferLayout { + arrayStride = size_of(f32) * 2, + stepMode = .Vertex, + attributes = &wgpu.VertexAttribute{shaderLocation = 0, format = .Float32x2}, + attributeCount = 1, + }, + }, + fragment = &{ + module = state.render_module, + entryPoint = "fragmentMain", + targetCount = 1, + targets = &wgpu.ColorTargetState{format = .BGRA8Unorm, writeMask = wgpu.ColorWriteMaskFlags_All}, + }, + primitive = {topology = .TriangleList}, + multisample = {count = 1, mask = 0xFFFFFFFF}, + }, + ) +} + +create_compute_pipeline :: proc() { + state.compute_module = wgpu.DeviceCreateShaderModule( + state.device, + &{ + label = "Compute Shader", + nextInChain = &wgpu.ShaderSourceWGSL{sType = .ShaderSourceWGSL, code = #load("shaders/compute.wgsl")}, + }, + ) + + state.compute_pipeline = wgpu.DeviceCreateComputePipeline( + state.device, + &{ + label = "Game of Life Compute Pipeline", + layout = state.pipeline_layout, + compute = {module = state.compute_module, entryPoint = "computeMain"}, + }, + ) +} + +//============================================================================= +// Buffer Creation +//============================================================================= + +create_buffers_and_bind_groups :: proc() { + // Uniform buffer (grid size) + grid_data := [2]f32{GRID_SIZE, GRID_SIZE} + state.uniform_buffer = wgpu.DeviceCreateBuffer( + state.device, + &{label = "Grid Uniform", usage = {.Uniform, .CopyDst}, size = size_of(grid_data)}, + ) + wgpu.QueueWriteBuffer(state.queue, state.uniform_buffer, 0, &grid_data[0], size_of(grid_data)) + + // Storage buffers (ping-pong for cell states) + cell_count := GRID_SIZE * GRID_SIZE + cell_bytes := u64(size_of(u32) * cell_count) + + for i in 0 ..< 2 { + state.cell_state_storage[i] = wgpu.DeviceCreateBuffer( + state.device, + &{ + label = i == 0 ? "Cell State A" : "Cell State B", + usage = {.Storage, .CopyDst}, + size = cell_bytes, + }, + ) + } + + // Initialise both buffers with random data + { + cells: [GRID_SIZE * GRID_SIZE]u32 + context = runtime.default_context() + rand.reset(u64(time.now()._nsec)) + + for i in 0 ..< GRID_SIZE * GRID_SIZE { + cells[i] = cast(u32)rand.int31_max(2) + } + + wgpu.QueueWriteBuffer(state.queue, state.cell_state_storage[0], 0, &cells, uint(cell_bytes)) + wgpu.QueueWriteBuffer(state.queue, state.cell_state_storage[1], 0, &cells, uint(cell_bytes)) + } + + // Create bind groups (ping-pong) + for i in 0 ..< 2 { + read_buffer := state.cell_state_storage[i] + write_buffer := state.cell_state_storage[(i + 1) % 2] + + entries := [3]wgpu.BindGroupEntry{ + {binding = 0, buffer = state.uniform_buffer, size = size_of(f32) * 2}, + {binding = 1, buffer = read_buffer, size = cell_bytes}, + {binding = 2, buffer = write_buffer, size = cell_bytes}, + } + + state.bind_groups[i] = wgpu.DeviceCreateBindGroup( + state.device, + &{ + label = i == 0 ? "Bind Group 0" : "Bind Group 1", + layout = state.bind_group_layout, + entryCount = 3, + entries = &entries[0], + }, + ) + } +} + +//============================================================================= +// Simulation Logic +//============================================================================= + +update_simulation :: proc(dt: f32) { + // Convert dt to Duration (web uses seconds, desktop uses milliseconds) + dt_duration: time.Duration + when ODIN_OS == .JS { + dt_duration = time.Duration(f64(dt) * f64(time.Second)) + } else { + dt_duration = time.Duration(f64(dt) * f64(time.Millisecond)) + } + + state.accumulator += dt_duration + accumulator_ms := time.duration_milliseconds(state.accumulator) + state.do_update = accumulator_ms >= UPDATE_INTERVAL_MILLISECONDS + + if state.do_update { + state.accumulator = 0 + } +} + +run_compute_pass :: proc(encoder: wgpu.CommandEncoder) { + state.did_compute = false + + if !state.do_update { + return + } + + cpass := wgpu.CommandEncoderBeginComputePass(encoder) + defer wgpu.ComputePassEncoderRelease(cpass) + + wgpu.ComputePassEncoderSetPipeline(cpass, state.compute_pipeline) + wgpu.ComputePassEncoderSetBindGroup(cpass, 0, state.bind_groups[state.step_index % 2]) + + workgroups := u32((GRID_SIZE + WORKGROUP_SIZE - 1) / WORKGROUP_SIZE) + wgpu.ComputePassEncoderDispatchWorkgroups(cpass, workgroups, workgroups, 1) + wgpu.ComputePassEncoderEnd(cpass) + + state.did_compute = true +} + +//============================================================================= +// Rendering +//============================================================================= + +resize :: proc "c" () { + context = state.ctx + state.config.width, state.config.height = os_get_framebuffer_size() + wgpu.SurfaceConfigure(state.surface, &state.config) +} + +frame :: proc "c" (dt: f32) { + context = state.ctx + + update_simulation(dt) + + // Get surface texture + surface_texture := wgpu.SurfaceGetCurrentTexture(state.surface) + + switch surface_texture.status { + case .SuccessOptimal, .SuccessSuboptimal: + // OK + case .Timeout, .Outdated, .Lost: + if surface_texture.texture != nil { + wgpu.TextureRelease(surface_texture.texture) + } + resize() + return + case .OutOfMemory, .DeviceLost, .Error: + fmt.panicf("[Game of Life] surface error: %v", surface_texture.status) + } + defer wgpu.TextureRelease(surface_texture.texture) + + frame_view := wgpu.TextureCreateView(surface_texture.texture, nil) + defer wgpu.TextureViewRelease(frame_view) + + encoder := wgpu.DeviceCreateCommandEncoder(state.device) + defer wgpu.CommandEncoderRelease(encoder) + + // Compute pass + run_compute_pass(encoder) + + // Render pass + { + render_pass := wgpu.CommandEncoderBeginRenderPass( + encoder, + &{ + label = "Main Render Pass", + colorAttachmentCount = 1, + colorAttachments = &wgpu.RenderPassColorAttachment { + view = frame_view, + loadOp = .Clear, + storeOp = .Store, + depthSlice = wgpu.DEPTH_SLICE_UNDEFINED, + clearValue = {0.1, 0.2, 0.3, 1.0}, + }, + }, + ) + defer wgpu.RenderPassEncoderRelease(render_pass) + + wgpu.RenderPassEncoderSetPipeline(render_pass, state.render_pipeline) + + // Use opposite bind group to read latest computed state + bind_group_index := (state.step_index + 1) % 2 + wgpu.RenderPassEncoderSetBindGroup(render_pass, 0, state.bind_groups[bind_group_index]) + + wgpu.RenderPassEncoderSetVertexBuffer(render_pass, 0, state.vertex_buffer, 0, state.vertex_buffer_size) + wgpu.RenderPassEncoderDraw(render_pass, state.vertex_count, GRID_SIZE * GRID_SIZE, 0, 0) + wgpu.RenderPassEncoderEnd(render_pass) + } + + command_buffer := wgpu.CommandEncoderFinish(encoder, nil) + defer wgpu.CommandBufferRelease(command_buffer) + + wgpu.QueueSubmit(state.queue, {command_buffer}) + wgpu.SurfacePresent(state.surface) + + // Advance simulation + if state.did_compute { + state.step_index += 1 + } +} + +//============================================================================= +// Cleanup +//============================================================================= + +cleanup :: proc() { + wgpu.RenderPipelineRelease(state.render_pipeline) + wgpu.ComputePipelineRelease(state.compute_pipeline) + wgpu.PipelineLayoutRelease(state.pipeline_layout) + wgpu.BindGroupLayoutRelease(state.bind_group_layout) + wgpu.ShaderModuleRelease(state.render_module) + wgpu.ShaderModuleRelease(state.compute_module) + wgpu.BufferRelease(state.vertex_buffer) + wgpu.BufferRelease(state.uniform_buffer) + wgpu.BufferRelease(state.cell_state_storage[0]) + wgpu.BufferRelease(state.cell_state_storage[1]) + wgpu.BindGroupRelease(state.bind_groups[0]) + wgpu.BindGroupRelease(state.bind_groups[1]) + wgpu.QueueRelease(state.queue) + wgpu.DeviceRelease(state.device) + wgpu.AdapterRelease(state.adapter) + wgpu.SurfaceRelease(state.surface) + wgpu.InstanceRelease(state.instance) +} diff --git a/wgpu/sdl3/game-of-life/os_desktop.odin b/wgpu/sdl3/game-of-life/os_desktop.odin new file mode 100644 index 0000000..c7950cf --- /dev/null +++ b/wgpu/sdl3/game-of-life/os_desktop.odin @@ -0,0 +1,117 @@ +#+build !js +// Desktop Platform Layer - SDL3 +// ============================== +// Provides platform-specific initialisation and event loop for desktop builds. +// +package main + +import "core:fmt" +import SDL "vendor:sdl3" +import "vendor:wgpu" +import glue "vendor:wgpu/sdl3glue" + +window: ^SDL.Window + +os_init :: proc() { + if !SDL.Init({.VIDEO}) { + fmt.panicf("Failed to initialise SDL: %s", SDL.GetError()) + } + + window = SDL.CreateWindow("Conway's Game of Life", WIDTH, HEIGHT, {.RESIZABLE}) + if window == nil { + fmt.panicf("Failed to create window: %s", SDL.GetError()) + } +} + +os_get_surface :: proc(instance: wgpu.Instance) -> wgpu.Surface { + return glue.GetSurface(instance, window) +} + +os_get_framebuffer_size :: proc() -> (width: u32, height: u32) { + w, h: i32 + if SDL.GetWindowSizeInPixels(window, &w, &h) { + return u32(w), u32(h) + } + return WIDTH, HEIGHT +} + +os_request_adapter_and_device :: proc() { + // Request adapter (synchronous on desktop) + wgpu.InstanceRequestAdapter( + state.instance, + &{compatibleSurface = state.surface, powerPreference = .HighPerformance}, + {callback = on_adapter_sync}, + ) +} + +on_adapter_sync :: proc "c" ( + status: wgpu.RequestAdapterStatus, + adapter: wgpu.Adapter, + message: string, + userdata1, userdata2: rawptr, +) { + context = state.ctx + + if status != .Success || adapter == nil { + fmt.panicf("Failed to get WebGPU adapter: [%v] %s", status, message) + } + + state.adapter = adapter + + // Request device + desc := wgpu.DeviceDescriptor{} + desc.uncapturedErrorCallbackInfo = {callback = on_device_error} + + wgpu.AdapterRequestDevice(state.adapter, &desc, {callback = on_device_sync}) +} + +on_device_sync :: proc "c" ( + status: wgpu.RequestDeviceStatus, + device: wgpu.Device, + message: string, + userdata1, userdata2: rawptr, +) { + context = state.ctx + + if status != .Success || device == nil { + fmt.panicf("Failed to get WebGPU device: [%v] %s", status, message) + } + + complete_gpu_init(device) +} + +os_run :: proc() { + // Main event loop with delta time tracking + event: SDL.Event + running := true + + last := SDL.GetPerformanceCounter() + now: u64 + dt: f32 + + for running { + // Calculate delta time in milliseconds + now = SDL.GetPerformanceCounter() + dt = f32((now - last) * 1000) / f32(SDL.GetPerformanceFrequency()) + last = now + + for SDL.PollEvent(&event) { + #partial switch event.type { + case .QUIT: + running = false + case .WINDOW_RESIZED, .WINDOW_PIXEL_SIZE_CHANGED: + resize() + case .KEY_DOWN: + if event.key.scancode == .ESCAPE { + running = false + } + } + } + + frame(dt) + } + + cleanup() + SDL.DestroyWindow(window) + SDL.Quit() +} diff --git a/wgpu/sdl3/game-of-life/os_web.odin b/wgpu/sdl3/game-of-life/os_web.odin new file mode 100644 index 0000000..6e15396 --- /dev/null +++ b/wgpu/sdl3/game-of-life/os_web.odin @@ -0,0 +1,122 @@ +#+build js +// Web Platform Layer - WASM +// ========================== +// Provides platform-specific initialisation and browser event loop for web builds. +// +package main + +import "base:runtime" +import "core:fmt" +import "core:sys/wasm/js" +import "vendor:wgpu" + +device_ready: bool = false + +os_init :: proc() { + ok := js.add_window_event_listener(.Resize, nil, size_callback) + assert(ok) +} + +os_get_surface :: proc(instance: wgpu.Instance) -> wgpu.Surface { + return wgpu.InstanceCreateSurface( + instance, + &wgpu.SurfaceDescriptor { + nextInChain = &wgpu.SurfaceSourceCanvasHTMLSelector { + sType = .SurfaceSourceCanvasHTMLSelector, + selector = "#wgpu-canvas", + }, + }, + ) +} + +os_get_framebuffer_size :: proc() -> (width: u32, height: u32) { + rect := js.get_bounding_client_rect("body") + dpi := js.device_pixel_ratio() + return u32(f64(rect.width) * dpi), u32(f64(rect.height) * dpi) +} + +os_request_adapter_and_device :: proc() { + fmt.println("[1/3] Requesting adapter (async - waiting for browser)...") + + // Define callbacks inline - crucial for async to work in WASM + wgpu.InstanceRequestAdapter( + state.instance, + &{compatibleSurface = state.surface}, + {callback = on_adapter}, + ) + + // Inline callback for adapter acquisition + on_adapter :: proc "c" ( + status: wgpu.RequestAdapterStatus, + adapter: wgpu.Adapter, + message: string, + userdata1, userdata2: rawptr, + ) { + context = state.ctx + + if status != .Success || adapter == nil { + fmt.panicf("Failed to get WebGPU adapter: [%v] %s", status, message) + } + + fmt.println("[OK] [1/3] Adapter acquired") + state.adapter = adapter + + desc := wgpu.DeviceDescriptor{} + desc.uncapturedErrorCallbackInfo = {callback = on_device_error} + + fmt.println("[2/3] Requesting device (async)...") + wgpu.AdapterRequestDevice(adapter, &desc, {callback = on_device}) + } + + // Inline callback for device acquisition + on_device :: proc "c" ( + status: wgpu.RequestDeviceStatus, + device: wgpu.Device, + message: string, + userdata1, userdata2: rawptr, + ) { + context = state.ctx + + if status != .Success || device == nil { + fmt.panicf("Failed to get WebGPU device: [%v] %s", status, message) + } + + fmt.println("[OK] [2/3] Device acquired") + complete_gpu_init(device) + } + + fmt.println("Waiting for browser callbacks...") +} + +os_run :: proc() { + device_ready = true + fmt.println("[OK] [3/3] WebGPU initialised - Game of Life ready!") +} + +// step is called by the browser runtime at 60 FPS +@(export) +step :: proc(dt: f32) -> bool { + context = state.ctx + + // Wait for async GPU initialization + if !device_ready { + return true + } + + frame(dt) + return true +} + +@(fini) +cleanup_on_exit :: proc "contextless" () { + context = runtime.default_context() + cleanup() + js.remove_window_event_listener(.Resize, nil, size_callback) +} + +size_callback :: proc(e: js.Event) { + if !device_ready { + return + } + resize() +} diff --git a/wgpu/sdl3/game-of-life/shaders/compute.wgsl b/wgpu/sdl3/game-of-life/shaders/compute.wgsl new file mode 100644 index 0000000..0942f5a --- /dev/null +++ b/wgpu/sdl3/game-of-life/shaders/compute.wgsl @@ -0,0 +1,32 @@ +@group(0) @binding(0) var grid : vec2f; +@group(0) @binding(1) var cellStateIn : array; +@group(0) @binding(2) var cellStateOut : array; + +fn cellIndex(cell : vec2u) -> u32 { + return (cell.y % u32(grid.y)) * u32(grid.x) + (cell.x % u32(grid.x)); +} + +fn cellActive(x : u32, y : u32) -> u32 { + return cellStateIn[cellIndex(vec2(x, y))]; +} + +@compute @workgroup_size(8, 8) +fn computeMain(@builtin(global_invocation_id) cell : vec3u) { + let activeNeighbors = + cellActive(cell.x + 1u, cell.y + 1u) + + cellActive(cell.x + 1u, cell.y + 0u) + + cellActive(cell.x + 1u, cell.y - 1u) + + cellActive(cell.x + 0u, cell.y - 1u) + + cellActive(cell.x - 1u, cell.y - 1u) + + cellActive(cell.x - 1u, cell.y + 0u) + + cellActive(cell.x - 1u, cell.y + 1u) + + cellActive(cell.x + 0u, cell.y + 1u); + + let i = cellIndex(cell.xy); + + switch activeNeighbors { + case 2u: { cellStateOut[i] = cellStateIn[i]; } + case 3u: { cellStateOut[i] = 1u; } + default: { cellStateOut[i] = 0u; } + } +} diff --git a/wgpu/sdl3/game-of-life/shaders/render.wgsl b/wgpu/sdl3/game-of-life/shaders/render.wgsl new file mode 100644 index 0000000..9051b2b --- /dev/null +++ b/wgpu/sdl3/game-of-life/shaders/render.wgsl @@ -0,0 +1,33 @@ +struct VertexInput { + @location(0) pos : vec2f, + @builtin(instance_index) instance : u32, +}; + +struct VertexOutput { + @builtin(position) pos : vec4f, + @location(0) cell : vec2f, +}; + +@group(0) @binding(0) var grid : vec2f; +@group(0) @binding(1) var cellState : array; + +@vertex +fn vertexMain(input : VertexInput) -> VertexOutput { + let i = f32(input.instance); + let cell = vec2f(i % grid.x, floor(i / grid.x)); + let state = f32(cellState[input.instance]); + + let cellOffset = cell / grid * 2.0; + let gridPos = (input.pos * state + 1.0) / grid - 1.0 + cellOffset; + + var out : VertexOutput; + out.pos = vec4f(gridPos, 0.0, 1.0); + out.cell = cell; + return out; +} + +@fragment +fn fragmentMain(input : VertexOutput) -> @location(0) vec4f { + let c = input.cell / grid; + return vec4f(c, 1.0 - c.x, 1.0); +} diff --git a/wgpu/sdl3/game-of-life/web/index.html b/wgpu/sdl3/game-of-life/web/index.html new file mode 100644 index 0000000..99560bb --- /dev/null +++ b/wgpu/sdl3/game-of-life/web/index.html @@ -0,0 +1,35 @@ + + + + + + WGPU WASM Game of Life + + + + + + + + + diff --git a/wgpu/sdl3-triangle/build_web.bat b/wgpu/sdl3/triangle/build_web.bat similarity index 100% rename from wgpu/sdl3-triangle/build_web.bat rename to wgpu/sdl3/triangle/build_web.bat diff --git a/wgpu/sdl3-triangle/build_web.sh b/wgpu/sdl3/triangle/build_web.sh similarity index 100% rename from wgpu/sdl3-triangle/build_web.sh rename to wgpu/sdl3/triangle/build_web.sh diff --git a/wgpu/sdl3-triangle/main.odin b/wgpu/sdl3/triangle/main.odin similarity index 100% rename from wgpu/sdl3-triangle/main.odin rename to wgpu/sdl3/triangle/main.odin diff --git a/wgpu/sdl3-triangle/os_js.odin b/wgpu/sdl3/triangle/os_js.odin similarity index 100% rename from wgpu/sdl3-triangle/os_js.odin rename to wgpu/sdl3/triangle/os_js.odin diff --git a/wgpu/sdl3-triangle/os_sdl3.odin b/wgpu/sdl3/triangle/os_sdl3.odin similarity index 100% rename from wgpu/sdl3-triangle/os_sdl3.odin rename to wgpu/sdl3/triangle/os_sdl3.odin diff --git a/wgpu/sdl3-triangle/web/index.html b/wgpu/sdl3/triangle/web/index.html similarity index 100% rename from wgpu/sdl3-triangle/web/index.html rename to wgpu/sdl3/triangle/web/index.html