Inkling sharpened my reason for building MoEscope. It gives us a frontier-scale example of the routing problem while the project starts with models small enough to inspect.
Thinking Machines Lab describes Inkling as a 975-billion-parameter Mixture-of-Experts model with 41 billion parameters active for each token. In each MoE layer, one token is sent to 6 of 256 routed experts while 2 shared experts run for every token.
That architecture creates two sources of expert work:
The executed workload includes routed experts and always-on shared experts.
The router chooses six experts. The GPU also executes two shared experts for every token. A useful trace format must represent both sources without pretending the router selected the shared experts.
This note maps that model-level statement down to the hardware. I started with a basic question: if CUDA programs NVIDIA GPUs, where do PyTorch, Triton, and Modal fit?
What Inkling changes
The original Phase 1 plan used a tiny top-2 router and OLMoE. That remains the right place to learn and test the trace format on hardware we can use. Inkling has 975 billion parameters, so OLMoE stays our beginner model.
Inkling changes three design requirements:
- The trace schema must support arbitrary top-k values instead of assuming two routed experts per token.
- Shared experts need their own representation because they create work without appearing in the router’s selected IDs.
- Synthetic workloads need honest labels. We can create an “Inkling-shaped” 6-of-256 workload later. Only data captured from the model qualifies as an Inkling trace.
Matching Inkling’s public dimensions does not reproduce its routing distribution.
The playground separates two quantities. tokens × top-k sets the number of
routed-expert calls. The routing distribution sets how unevenly those calls
spread across experts.
Shared experts add tokens × shared experts calls from the model architecture.
The five layers people mix together
PyTorch, Triton, CUDA, NVIDIA GPUs, and Modal can appear in the same installation guide even though each belongs to a different layer.
Express the math.
PyTorch gives us tensors, model layers, automatic differentiation, and familiar operations such as topk. It is where we build the correct reference before touching kernel code.
expert_ids = torch.topk(router_scores, k=2).indicesShape the GPU work.
Triton is a Python-based language and compiler for custom GPU kernels. We describe blocks of tensor work; its compiler lowers them toward the NVIDIA execution stack.
@triton.jit def grouped_gemm(...): # one program handles a tile of expert workControl NVIDIA compute.
CUDA is broader than a language. It is NVIDIA’s programming model, runtime, compiler toolchain, libraries, and vocabulary for launching parallel work and moving memory.
host code → launch kernel → GPU threads → device memoryExecute many tiles in parallel.
The physical GPU contains streaming multiprocessors, compute units, caches, and high-bandwidth memory. Kernel speed depends on how well the requested work fits those resources.
grouped GEMMs → tiles → thread blocks / CTAs → SMsRent the environment.
Modal is not another GPU language. It starts a cloud container with a chosen NVIDIA GPU, installs our dependencies, runs the experiment, and charges for the compute time.
@app.function(gpu="A100-40GB") def benchmark(): ...| Layer | What it is | What MoEscope uses it for |
|---|---|---|
| PyTorch | Tensor and machine-learning framework | Correct model logic and reference output |
| Triton | Python-based GPU kernel language and compiler | Custom grouped-GEMM experiments |
| CUDA | NVIDIA platform, programming model, runtime, toolchain, and libraries | The execution environment underneath our GPU code |
| NVIDIA GPU | Physical accelerator | Runs the compiled parallel work |
| Modal | Cloud compute service | Gives us a reproducible machine with a selected GPU |
Cloudflare Pages publishes the static MoEscope website. Modal supplies the NVIDIA GPU that executes and profiles our later Triton kernels.
PyTorch: describe the tensor computation
PyTorch is where we can express model behavior at a high level:
scores = torch.softmax(router_logits, dim=-1)
weights, expert_ids = torch.topk(scores, k=2, dim=-1)
counts = torch.bincount(expert_ids.flatten(), minlength=num_experts)
This code says what should happen without asking us to manually assign GPU threads, move tiles through shared memory, or write a matrix-multiplication kernel.
PyTorch operations may run on the CPU or dispatch GPU implementations when their tensors live on a CUDA device. For Phase 1, that distinction is useful: we can design and validate the trace logic on a laptop before renting a GPU.
CUDA: the NVIDIA computing platform
CUDA includes more than a programming language. NVIDIA’s programming guide describes CUDA as a parallel computing platform and programming model. It includes:
- a model for organizing parallel threads and memory;
- runtime and driver APIs;
- language support such as CUDA C++;
- compilers and tools;
- optimized libraries such as cuBLAS;
- profilers and debugging tools.
A kernel is a function that runs on the GPU across many parallel threads. CUDA gives programmers direct control over how that work is launched and how memory is used. That control is powerful, but writing fast CUDA C++ kernels requires detailed hardware knowledge.
Triton: write GPU kernels from Python
Triton is a language and compiler for parallel programming. Its Python-based syntax is designed for writing custom deep-learning kernels.
Triton sits above the NVIDIA execution stack. We describe blocks, or tiles, of work in Python, and Triton compiles them for the target GPU.
That makes Triton a useful middle layer:
PyTorch reference
↓ compare correctness
Triton kernel
↓ compile and launch
CUDA / NVIDIA execution stack
↓
GPU hardware
For MoEscope, the interesting operation is grouped matrix multiplication. Every
active expert receives a different number of tokens, so the M dimension of
its matrix multiplication changes:
expert 0: [37 tokens × hidden size] @ expert weights
expert 1: [ 4 tokens × hidden size] @ expert weights
expert 2: [19 tokens × hidden size] @ expert weights
...
The math is straightforward. Scheduling many differently shaped operations efficiently is the systems problem.
Modal rents the machine
Modal is cloud infrastructure. A Modal
function can request an A100-40GB, install a pinned software environment, run
our benchmark, save results, and stop.
Modal supplies the Linux machine and NVIDIA GPU. PyTorch, Triton, and CUDA run on that machine.
The initial compute plan is:
- laptop CPU for trace schema, routing logic, invariants, and tests;
- a pinned Modal A100 for later CUDA and Triton measurements;
- Cloudflare Pages for the public writing and interactive explorer.
What MoEscope will actually observe
MoEscope records one routing event as four separate artifacts:
model state
router scores, selected IDs, routing weights, shared-expert rules
↓
normalized trace
portable facts about what the model decided
↓
compiled workload
tokens per expert, GEMM shapes, tiles, padding, possible waves
↓
benchmark result
correctness, latency distribution, environment, profiler evidence
Benchmark measurements stored inside an input trace prevent a clean replay on another GPU. Recording shared experts as router selections changes what the model decided. Calling synthetic dimensions a real Inkling trace makes the evidence impossible to audit. Separate artifacts avoid all three errors.
Update: Checkpoint 01 is implemented
Since publishing this note, I implemented the first checkpoint.
It uses a synthetic fixture small enough to check by hand:
T0 → E0, E2
T1 → E1, E2
T2 → E2, E0
T3 → E3, E1
Flattening those routed IDs produces eight assignments:
[0, 2, 1, 2, 2, 0, 3, 1]
The implementation derives:
routed histogram [2, 2, 3, 1]
routed calls 4 tokens × top-2 = 8
shared calls 4 tokens × 2 shared experts = 8
total expert calls 16
[2, 2, 3, 1]The fixture contains synthetic data. Its four tokens establish the trace contract before we introduce tensors, model hooks, or a real model checkpoint.
The current implementation has four parts:
- A dependency-free, immutable routing-trace schema.
- The canonical four-token JSON fixture.
- A deterministic summary function that keeps routed and shared work separate.
- Golden and invariant tests for widths, ranges, duplicate IDs, finite weights, versioning, canonical serialization, and the expected histogram.
The schema also rejects unknown fields. A value such as latency_ms cannot be
silently inserted into an input trace; benchmark results belong to a later
artifact.
You can reproduce the checkpoint locally:
uv sync --dev
uv run pytest
uv run python examples/summarize_trace.py
All 13 tests pass. The chronological implementation is preserved in four focused commits, and the accompanying Phase 1 guide explains the routing event, code, and learning path.
Checkpoint 01 stops at serialized traces. Later checkpoints will add:
- tensor-to-trace conversion;
- a real OLMoE trace;
- workload compilation;
- a PyTorch correctness backend;
- a Triton kernel and GPU benchmark.
What Phase 1 requires
Phase 1 ends at a trustworthy routing trace. Triton kernels and CUDA memory details belong to later phases.
We need to understand:
- tensors and their shapes;
- router logits, scores, top-k IDs, and weights;
- expert histograms;
- routed versus shared experts;
- deterministic serialization and validation;
- enough PyTorch to calculate and test those values.
Phase 2 will translate the trace into GPU work, and Phase 3 will replay it. We will profile only after the baseline output matches.
The next checkpoint will compute top-k assignments from the original four score rows and prove that PyTorch reproduces this exact framework-neutral fixture.
References
Primary sources used for this note:
- Thinking Machines Lab, “Inkling: Our Open-Weights Model”, July 15, 2026.
- Thinking Machines Lab, Inkling model card.
- NVIDIA, CUDA Programming Guide.
- Triton, official documentation and Group GEMM tutorial.
- PyTorch, CUDA semantics.
- Modal, GPU acceleration documentation.
- Astro, MDX integration.
- Cloudflare, deploying Astro to Pages.