Blaze SDK

A DSL compiler for high-performance ML kernels.

Write plain scalar loops — no intrinsics, no CUDA, no shader code.
The compiler turns them into hardware-optimized kernels across every backend.

127.6 TFLOPS BF16 Matmul
1688 tok/s 9B prefill
98.7 tok/s 9B decode
495 tok/s 27B prefill
2565 GB/s INT8 Matvec · 138% cuBLASLt
1300× Schedule speedup

If these numbers match the workloads you care about, we'd like to hear from you.

How it works

The kernel defines what to compute. The schedule defines how fast.

1

Write

kernel.blz

fn kernel_main(A: array[f32],
    B: array[f32], C: array[f32],
    M: i32, N: i32, K: i32) {
  for (var i = 0i; i < M; i = i + 1i)
    for (var j = 0i; j < N; j = j + 1i)
      for (var k = 0i; k < K; k = k + 1i)
        C[i*N+j] = C[i*N+j]
          + A[i*K+k] * B[k*N+j];
}

Plain scalar loops. No intrinsics.

→
2

Schedule

matmul.lum

schedule MatMul(blaze.loop_nest) {
  tile [64, 64, 64]
  parallel
  unroll j 4
  microkernel f32 [32]
}

Controls tiling, parallelism, vectorization.

→
3

Run

terminal

$ blaze matmul.blz \
    --schedule matmul.lum \
    --emit-mlir -o kernel.o

$ gcc -shared -o kernel.so \
    kernel.o -lblaze_runtime

  1024×1024 f32 matmul
  1045 GFLOPS (99% peak)

Same source. Any backend.

Performance

Measured on Zen 4 (8 cores pinned) and RTX 4090. Qwen numbers use SDK-built kernels and Q4_K_M GGUF weights.

GPU RTX 4090

BF16 Matmul 4096² 127.6 TFLOPS
TF32 Matmul 4096² 43.1 TFLOPS
FlashAttention N=8192 137 TFLOPS
BF16 Matvec 11008×4096 836 GB/s
INT8 Matvec 2565 GB/s
Stories110M 1200 tok/s
Qwen 3.5 9B prefill (256p) 1688 tok/s
Qwen 3.5 9B decode (16t) 98.7 tok/s
Qwen 3.6 27B prefill (2048p) 495 tok/s
Qwen 3.6 27B decode (16t) 18.2 tok/s

CPU Zen 4, 8 cores

F32 Matmul 1024² 1045 GFLOPS
Stories110M 120+ tok/s
Qwen 3.5 9B Q4K CPU 8 tok/s
LLaMA 3.2 3B Q4K+INT8 22.6 tok/s
Same kernel, only the schedule changes
No schedule 0.8 GFLOPS
+ microkernel 1045 GFLOPS (1300×)

Prefill rates are reported with prompt shape: 9B uses 256 prompt tokens and chunk 256; 27B uses 2048 prompt tokens, chunk 2048, and max_seq 4096. Decode rows use 16 generated tokens.

Progressive schedule tuning — 0.8 to 1045 GFLOPS
Schedule impact across matrix sizes

Backends

One kernel source. Five hardware targets.

x86

x86-64

AVX-512 · VNNI · OpenMP

ARM

Apple Silicon

NEON · SME · M-series

GPU

NVIDIA CUDA

Tensor Cores · PTX · CUDA Graph

MTL

Metal

Compute shaders · Apple GPU

ANE

Neural Engine

Experimental · Private API

From loops to hardware

The compiler auto-detects loop nests from plain scalar code — no annotations, no pragmas, no intrinsics. The Lum schedule is the tuning surface.

Quantized LLM inference — Lum schedule
schedule QwenDecode(blaze.loop_nest) {
  // Shape dispatch — each matvec gets optimal tiling
  where bounds = [3584, 18944] {
    tile [1, 3584, 592]
    parallel
    unroll j 4
    microkernel q4k [32]
    weight_format "q4k"
  }
  where bounds = [18944, 3584] {
    tile [1, 18944, 448]
    parallel
    unroll j 4
    microkernel q4k [32]
    weight_format "q4k"
  }
}
GPU Tensor Core schedule
schedule BF16Matmul(blaze.loop_nest) {
  tile [128, 128, 32]
  gpu_threads [16, 16]
  warp_tiles [1, 8]
  tile [8, 8, 1] register
  maxnreg 160
}

// 127.6 TFLOPS at 4096×4096
// m16n8k16 BF16 MMA
// cp.async + ldmatrix.trans