Doppler · Doe · Cerebras

From WebGPU to CSL.

Doppler supplies the browser reference. Doe generates kernels in CSL, the Cerebras programming language. The research question: can the generated code reproduce the same model computation?

Selected-logit checks and compilation are recorded. Full-model hardware execution and performance are not established by this page.

01 · EVIDENCE

What the records show

Bounded checks from the May 6, 2026 archive. Selected-logit agreement is not full-model parity.

SELECTED LOGIT

Gemma selected token

One output-head score from frozen WebGPU state: token 3730 ( blue). Recorded difference: 0.0087417.

SELECTED LOGIT

Qwen selected token

Final norm plus one output-head row from frozen WebGPU state: token 760. Recorded difference: 0.0133286.

COMPILES

Compilation records

The archive includes generated-target compilation records for both models. Compilation does not establish successful inference.

PACKAGED

Hardware commands

Both model paths have packaged run commands. A prepared command is not a returned hardware result.

CEILING

Local simulator boundary

The recorded Gemma lm_head tile reaches compute completion, then stops at device-to-host copyback. It does not produce a complete model transcript.

The archive also retains reference outputs, generated kernels, run scripts, and a verifier.

02 · CHECK LOOP

The comparison path

Reference, generated code, execution, and comparison are separate evidence steps. The complete hardware loop remains a target.

REFERENCE

Browser model

JavaScript and WGSL model code that runs locally.

DOCS Program Bundle

REFERENCE

Doppler bundle

Portable graph, kernels, shapes, parameters, and expected output.

HOSTED Doppler demo

Doe TSIR

Tiled Spatial IR describes shapes and operations. Live CSL execution wiring remains incomplete.

SOURCE TSIR tree

TARGET

Generated CSL

PE kernels, routing, memory movement, and run order emitted by Doe.

DOCS CSL architecture

CHECK

Diff against browser

Hardware or simulator output compared to the browser reference with a recorded tolerance.

SCHEMA Proof file

Inspect the Gemma splice in Forge: a recorded reference and browser-executed tail, not a WSE run.

03 · TECHNICAL MAP

WGSL concepts mapped to CSL

How GPU concepts map to a spatial target.

WGSL and JS	Generated CSL and run script	Why it lowers
`@workgroup_size` and workgroup tiles	PE grid and tile dimensions	Parallel topology is explicit instead of inferred.
Storage buffers and typed array views	Named CSL symbols plus scripted input/output movement	Dtype, shape, and visibility stay auditable.
JavaScript dispatch call	Scripted launch on Cerebras hardware	The run boundary is recorded in the proof file.
`workgroup_id` and `local_invocation_id`	PE coordinates and tile indices	Indexing intent survives lowering.
Barrier or subgroup collective	Fabric collective or explicit host step	Synchronization becomes a declared target decision.
Kernel source + JS wrapper	Generated CSL plus proof file	The proof file binds source identity, generated code, and output comparison.

DOCS Doppler bundle SOURCE Doe TSIR DOCS CSL architecture SOURCE emit_csl_matmul.zig

04 · SIDE BY SIDE

Visible kernel lowering

RoPE, tiled matmul, and output packing, with source links.

RoPE SOURCE rope.wgsl

let cos_val = freqs_cos[freq_idx];
let sin_val = freqs_sin[freq_idx];
let x0 = input[base_idx + first_idx];
let x1 = input[base_idx + second_idx];

let y0 = x0 * cos_val - x1 * sin_val;
let y1 = x0 * sin_val + x1 * cos_val;

input[base_idx + first_idx] = y0;
input[base_idx + second_idx] = y1;

RoPE SOURCE rope.csl

const q0: f32 = buf_out[q_base + a + 0];
const q1: f32 = buf_out[q_base + a + 1];
const qc: f32 = rope_cos_at(kv_len_per_head, d);
const qs_: f32 = rope_sin_at(kv_len_per_head, d);

q_rot[a + 0] = (qc * q0) - (qs_ * q1);
q_rot[a + 1] = (qs_ * q0) + (qc * q1);

WGSL uses 16×16 workgroups; CSL uses a P×P processing-element grid with explicit broadcasts. Doppler selects variants through a JSON rule map.

Matmul (tiled f16) SOURCE matmul_f16_tiled.wgsl

const TILE_M: u32 = 64u;
const TILE_K: u32 = 16u;
var<workgroup> tileA: array<f16, 1024>;

@compute @workgroup_size(16, 16, 1)
fn main(...) {
    for (var t: u32 = 0u; t < num_tiles; t = t + 1u) {
        // load tile, barrier, accumulate 4x4 register tile
    }
}

Matmul (SUMMA on P×P PE grid) SOURCE emit_csl_matmul.zig

param Mt: i16; param Kt: i16; param Nt: i16; param P: u16;

const mpi_x = @import_module("<collectives_2d/pe>",
                              .{ .dim_params = c2d_params.x });

var A_tile = @zeros([Mt * Kt]f32);
var C_tile = @zeros([Mt * Nt]f32);

mpi_x.broadcast(step, &A_tile, Mt * Kt, x_done_id);

The output-packing pass groups f16 values into 128-word chunks for export.

RMSNorm post-emission pack SOURCE emit_csl_rmsnorm_pack.zig

for (@range(i16, 128)) |word| {
    const base = (@as(u32, word) + chunk_offset) * 2;
    const lo: u32 = @as(u32, @bitcast(u16, output[base]));
    const hi: u32 = @as(u32, @bitcast(u16, output[base + 1]));
    output_chunk_0000[@as(u32, word)] = lo | (hi << 16);
}

05 · HARDWARE CHECK

What would close the loop

A real-session transcript and a returned WSE receipt, tied to the same model, weights, inputs, and generated code. Performance needs a separate measurement.

REFERENCE

Browser reference

Freeze the inputs, model, browser outputs, and tolerances.

TARGET

Generated CSL

Complete live lowering and real-session execution, not just target compilation.

VERIFY

Evidence bundle

The linked May archive is historical. Its own metadata identifies the source revision; it is not a current release declaration.

RETURN

Returned trace

Record tokens, logits, and KV state, or name the exact boundary that failed.

ARCHIVE May 6 archive DOCS Hardware runbook DOCS Bundle README DOCS Archive pointer SOURCE Gemma run command SOURCE Qwen run command SOURCE Verifier

CSL compiler research. Inference and validation, not training. · github.com/clocksmith/doppler · github.com/doe-gpu/doe · x@d4da.com