TypeScript → CUDA — eight compilation phases transform source into warp-parallel GPU execution at 10 TB/s memory bandwidth. Every HTTP request becomes a SIMT kernel.
Eight phases transform a .ts file into machine code running on GPU hardware. Click any stage to expand detail. Phases 3–5 are proprietary and never published.
.d.ts discardedsymm_dispatchsymm_dispatch_batchAny TypeScript that avoids workers, dynamic eval, native add-ons, or blocking FS APIs. Generics, as casts, and interface declarations produce zero runtime output — all erased by esbuild before SymmParser ever sees the code.
A ~1,500-line CUDA C++ prelude defining SymmCtx, Map, Set, URL, TextEncoder, RegExp (NFA engine), all String.* helpers, Math.* remaps, Promise stubs, and the symm_dispatch_batch kernel entry point.
After cuModuleLoadData the PTX string is immediately zeroed in heap memory. The driver has already JIT-compiled it to a native cubin inside the GPU. Dumping process memory after startup reveals no PTX.
All eight stages above run exactly once, when you invoke sym server.ts — before the server accepts a single connection. Stages ① through ④ are ordinary CPU work (string transforms, parsing, code generation); NVRTC then JIT-compiles the generated CUDA C++ into PTX, and the driver assembles that PTX into native GPU machine code (SASS) resident on the device.
Once server.run() starts blocking, compilation is over. The persistent dispatch kernel is already launched and resident on the SMs, looping forever. Serving a request means: the CPU writes the request into a pinned-memory slot, and GPU threads that are already running compiled machine code pick it up, execute it, and write the response back. There is no interpreter, no bytecode VM, and no per-request recompilation — the "JS" your handler was written in has been permanently reduced to native GPU instructions by the time traffic arrives.
This is also why sym build server.ts exists as a separate command: it lets you do the compile step once, ahead of time, and ship the resulting symmetra.ptx as a standalone artifact — making the compile/runtime split an explicit, first-class boundary rather than an implementation detail.
A real handler, shown at each compilation stage. The same 10 lines of TypeScript become ~80 lines of CUDA C++ with the runtime prelude woven in.
The Product interface is erased by esbuild. Codegen sees a plain JS object literal and inlines field access as direct struct reads.
__device__ dataConstant arrays are placed in __device__ static storage — fast GPU global memory. No heap allocation at runtime.
Tier-1 async (no real I/O) is synchronous on the GPU. async and await keywords are removed; the body executes inline.
Codegen replaces res.json(item) with a direct call to symm_json_obj3 from the prelude — no runtime reflection needed.
Top-to-bottom: user code at the top, bare GPU silicon at the bottom. Each layer has a single responsibility and communicates only with its immediate neighbours. Hover a layer to highlight it.
.ts file, drives esbuild type-strip, invokes SymmParser + Codegen + NvrtcCompiler, starts GpuServerNodePtr tree). Handles classes, generators, for await, computed Symbol.* method names, destructuring, optional chaining.async/await into synchronous or split-kernel CPS, generates symm_dispatch_batch.symm_dispatch and symm_dispatch_batch function handles.preferredBatchSize=256 or fillTimeoutUs=5 ms elapses. All requests share a route → zero warp divergence at the dispatch switch. Pads to warp boundary (×32).symm_dispatch_batch entry point.A single HTTP request travels through 11 phases across CPU, DMA bus, and GPU. The shaded zone (phases 6–9) is where user TypeScript executes. Click any phase for detail.
IO thread wins ::accept() from shared listen socket. No CUDA context touched.
CPU reads headers + body (recvAll). Parses method, path, headers, Content-Length boundary.
CPU scans RouteReg table. Resolves param segments (/users/:id). Returns routeId integer.
Request placed in per-route queue as BatchRequest. Accept thread immediately returns to ::accept().
Scheduler waits for 255 more same-route requests or the 5 ms deadline. Fires immediately on full batch. Pads to warp boundary.
BatchContext::uploadBatch copies request bodies, paths, methods, d_active bitmask into VRAM via 8 async cuMemcpyHtoDAsync calls.
cuLaunchKernel(symm_dispatch_batch, 1,1,1, padded,1,1, …). One block, up to 256 threads. All threads execute same route → zero warp divergence.
Your TypeScript (compiled to CUDA C++) runs. Reads body, writes JSON to resp_buf, sets resp_len + resp_status. __syncwarp() ensures all warp threads finish.
Async copy of d_resp_lens + d_statuses (small, sync first), then per-slot body copy back to RAM.
buildResponse() assembles HTTP/1.1 status line + headers + body string for each slot in the batch.
sendBatchResponses() calls ::send(fd) + CLOSE_SOCKET(fd) per slot. Stream returned to pool.
| Phase | Where | Per batch |
|---|---|---|
| TCP accept + recv | CPU | ~1–10 µs/conn |
| Batch fill window | Scheduler | 0–5 000 µs |
| H2D upload (DMA) | PCIe / NVLink | ~30–80 µs |
| Kernel launch overhead | CUDA Driver | ~2–5 µs |
| User handler (GPU) | GPU SIMT | 10–500 µs |
| D2H download (DMA) | PCIe / NVLink | ~30–80 µs |
| HTTP build + send | CPU | ~1 µs/slot |
When a handler calls @symmetra/net or @symmetra/db it cannot block — the GPU has no scheduler that yields. Instead:
ctx.io_stage = N, then returns.io_stage = N, resuming after the I/O point.Status: architecture designed, VRAM I/O queue format defined in SymmCtx.io_stage; full proxy implementation pending for @symmetra/net and @symmetra/db.
Watch how BatchScheduler fills warp slots. Each row is one warp (32 threads). A full batch = 8 warps = 256 threads. Colour = route. Purple = padding (idle threads). Press Run to animate.
8 warps × 32 threads/warp
Symmetra's batch mode throughput compared to single-threaded Node.js baseline on an A100. All numbers are per GPU card, JSON handler, loopback network.
256 requests share one DMA transfer and one kernel launch overhead. The ~30 µs DMA + ~2 µs launch cost is amortised across all 256 slots. Per-request GPU compute time stays constant; overhead shrinks to <0.15 µs/req.
A request arriving at millisecond 0 may wait up to 5 ms for the fill window to close before GPU execution starts. For latency-critical routes use single-thread mode or reduce fillTimeoutUs.
A100 VRAM delivers ~2 TB/s (HBM2e). CPU RAM delivers ~50 GB/s. In-kernel JSON parsing, KV lookups, and string operations all operate at GPU-memory speed — 40× faster than the same work on CPU hitting L3 cache.
The GPU is deliberately isolated behind a thin abstraction boundary. Swapping backends (NVIDIA → AMD → Intel) requires changing ~4 files; parser, codegen, HTTP stack, and scheduler are untouched.
CUstream → hipStream_t / sycl::queuecuMemAlloc → hipMalloc / sycl::malloc_devicecuLaunchKernel → hipLaunchKernel / SYCL parallel_for__syncwarp() → __syncthreads() / HIP sub-groupHTTP servers are embarrassingly parallel — each request is independent. SIMT hardware runs 32 identical threads in lock-step per warp. When all 32 threads execute the same route handler there is zero branch divergence — full warp throughput.
The GPU is the only commodity hardware sustaining ~2 TB/s HBM bandwidth, letting per-request KV lookups, JSON parsing, and response serialisation happen in tens of microseconds.
Symmetra's SIMT execution model scales linearly with GPU generation. Every CUDA core contributes independently — each thread handles its own HTTP request in parallel, so throughput is directly proportional to CUDA core count and memory bandwidth.
The BatchScheduler collects incoming requests into per-route queues. When a queue reaches 256 requests (or a 5 ms timeout expires), the entire batch is dispatched to the GPU in a single kernel launch.
Inside the kernel each CUDA thread executes the same compiled handler but on its own private data — the SIMT (Single Instruction, Multiple Thread) model. There is no inter-thread synchronisation for independent API calls.
🔑 The 5 ms is the maximum wait to fill a batch — not the execution time. A full 256-request batch executes in approximately 50 µs of GPU time. Under high load the timeout is never reached; batches fill instantly.
* req/s estimates assume 256-thread batches, handler complexity typical of a CRUD endpoint (~50–80 CUDA instructions), sustained 100% queue fill rate. Real-world throughput depends on handler complexity, response payload size, and PCIe/NVLink copy overhead.
Hardening status across all attack surfaces — from the GPU kernel to the TCP edge. Items marked ✅ Done are shipped; items marked 🔲 Missing must be completed before v1.0.0 public release.
cuModuleLoadData✅ Done../, %2e%2e rejection)✅ DoneSecurity infrastructure implemented and covered by our test suites. Symmetra remains beta software — independent security review is planned before any production-readiness claim.