wafer@wafer_ai
🚨 Tenstorrent ships AI accelerator cards for $999 and the entire software stack is open source.
Here's How kernel programming works on their architecture:
their CUDA equivalent is TT-Metalium, a C++ kernel SDK for programming Tenstorrent's Tensix cores.
it's fully open source: the ISA, the kernel APIs, the compiler, the op library, everything.
but: if you know CUDA, TT-Metalium will feel quite different.
in CUDA you write one kernel that handles loading data, computing, and storing results. in TT-Metalium you write three separate kernels: a reader, a compute, and a writer. the reader DMA's tiles from DRAM through a Network-on-Chip into local SRAM. the compute kernel unpacks those tiles, runs matrix or vector math, and packs results. the writer DMA's results back out. three C++ files, three binaries, running concurrently on five tiny RISC-V processors inside each Tensix core.
the three kernels talk to each other through circular buffers - producer-consumer queues in L1 SRAM. the reader pushes a tile into a buffer with
cb_push_back(). the compute kernel blocks on cb_wait_front() until that tile arrives, processes it, pushes the result to an output buffer.
the writer waits on that output buffer and DMA's the result to DRAM. so instead of threads sharing memory and synchronizing with barriers, you have a pipeline: data flows through queues from reader to compute to writer. circular buffers are the entire synchronization mechanism. there are literally no __syncthreads(), no atomics, and no shared memory.
the memory model is also very different from NVIDIA's.
there are zero hardware caches. each Tensix core has 1.5 MB of private SRAM. want data from DRAM? explicitly DMA it. want data from another core's SRAM? specify (x, y, local_address) and issue an async NoC read. every byte of data movement is your responsibility.
on Blackhole that's 210 MB of aggregate SRAM across 140 cores, plus 32 GB GDDR6 at 512 GB/s, and all computation happens on 32x32 tiles that get unpacked into compute registers and packed back out - the hardware handles format conversion between storage formats (like block floating point) and compute formats (FP32) automatically.
as with many other accelerators, CUDA feels like a strong moat here. microbenchmarks show the hardware hitting near-theoretical peak in isolation, but LLM inference only reaches ~50% of peak on Blackhole. 76 of 140 Tensix cores sit idle running forward-compatible Wormhole kernels.
The Register reviewed the $11,999 QuietBox and called it "trapped in a software blackhole." the $1,399 Blackhole p150a performed almost identically to an NVIDIA DGX Spark despite specs suggesting a 2-3x advantage.
deep dive 5/6 by @gpuemi
