How does aNeural Networkthink inSilicon?

$ trace_inference --depth=full
# Follow a single token from a Python string to a tensor in HBM,
# through a Tensor Core, back to a vocabulary probability.

Start Journey View End-to-End

STAGES

↓

SOFTWARE → SILICON

LEVELS

INFERENCE PIPELINE

LIVE

framework

runtime

isa

hardware

driver

SCROLL TO TRACE

STAGE_01

The Input

Raw text becomes discrete tokens, then dense floating-point vectors. A vocabulary lookup is the first lookup table in the network.

Tokenization

01.A

[FRAMEWORK]

INPUT STRING12 chars

Hello, world

TOKENS (id, str)0/3

// awaiting input...

# A tokenizer maps strings to integer token IDs. GPT-2 uses Byte-Pair Encoding (BPE) — iteratively merge the most frequent adjacent byte pair.

# Python

from transformers import AutoTokenizer

tk = AutoTokenizer.from_pretrained("gpt2")

ids = tk("Hello, world")["input_ids"]

→ [15496, 11, 995] // 3 token IDs

Vocab

01.B

[FRAMEWORK]

VOCAB LOOKUP (BPE merge table)[FRAMEWORK]

000"Hello"

001","

002" world"

003"!"

004"?"

005"The"

006" a"

007" is"

008" of"

009" and"

010" to"

011" in"

012" that"

013" it"

014" for"

015" with"

016" as"

017" on"

018" be"

019" by"

Embedding Tensor

01.C

[FRAMEWORK]

BATCH_SIZEB = 2

1234

SEQ_LENN = 6

D_MODELD = 16

(real GPT-2: 768 → 12 layers)

DTYPE

32-bit float · IEEE 754

EMBEDDING TENSOR

FP32768 B

shape = [2, 6, 16]

b=0

b=1

PARAMS

192

FOOTPRINT

768 B

DTYPE

memory_footprint = B × N × D × 4 bytes

= 2 × 6 × 16 × 4 = 768 B

STAGE_02

The Layer

A single Linear layer is one matrix multiplication and one bias add, followed by a non-linearity. The matmul is the entire reason for the rest of this page.

Forward Pass

02.A

[FRAMEWORK]

LINEAR LAYER

FRAMEWORK

nn.Linear (PyTorch)

Dense (Keras) / Dense (JAX)

ACTIVATION

σ(Y)

ReLU · GELU · SiLU

# forward

Y = X @ W.T + b # [B, N, D_out]

A = gelu(Y) # activation

Activation

02.B

[FRAMEWORK]

Non-linear element-wise function

ReLU

max(0, x)

GELU

x·Φ(x)

During inference, the activation kernel is a separate GPU launch — small but adds latency.

FLOPs Calculator

02.C

[FRAMEWORK]

D_IN (input dim)8

D_OUT (output dim)8

BATCH2

SEQ_LEN6

COMPUTE COST

1.54K FLOPs

# formula

FLOPS = 2 × B × N × D_in × D_out

# expand

= 2 × 2 × 6 × 8 × 8

= 1,536 multiply-adds

WEIGHTS

BIASES

PARAMS TOTAL

STAGE_03

Matrix Multiplication

C[i][j] = Σₖ A[i][k] × B[k][j]. Naively O(n³) — but real hardware runs tiled, vectorized, and on dedicated silicon.

Dot-Product Stepping

03.A

[FRAMEWORK]

A [4×4]

B [4×4]

C [4×4]

128

120

197

154

108

138

126

132

183

181

133

# row of A (highlighted cyan)

A[0, :] = [10277]

# col of B (highlighted amber)

B[:, 0] = [4919]

# elementwise multiply + sum (dot product)

10×4+2×9+7×1+7×9=128

SPEED1.0s

C[0][0] = 128

Naive vs Tiled

03.B

[RUNTIME]

Naive reads each element O(n) times. Tiling loads a 2×2 block into cache and reuses each element 2×.

# naive (cache-unfriendly)

for i in range(N):
  for j in range(N):
    for k in range(N):
      C[i][j] += A[i][k] * B[k][j]

# tiled (cache-blocked) — used by cuBLAS

for ii in range(0, N, T):
  for jj in range(0, N, T):
    # load A[ii:ii+T, :] → SMEM
    for kk in range(0, N, T):
      # load B[:, kk:kk+T] → SMEM
      # 2×2 micro-kernel → registers

TILE

2×2

BLOCKS

2×2

REUSE

×2

Memory Pattern

03.C

[ISA]

Row-major (C) vs Column-major (Fortran, cuBLAS). Stride matters.

Row-Major (PyTorch)

Column-Major (cuBLAS)

# access pattern for C[i][j] += A[i][k] * B[k][j]

A[i][k]: stride 1, hot in L1

B[k][j]: stride N, ✗ cache miss

→ transpose B first (cuBLAS does this)

⚡ This single op is ~90% of total inference compute.

STAGE_04

The Math Library

PyTorch never writes a matrix-multiply kernel. It dispatches down a hardware-aware stack to BLAS, then to intrinsics, then to silicon.

GPU Path

04.A

[RUNTIME]

Tracing: torch.matmul(x, w).cuda() · click any layer to expand

CPU Path

04.B

[RUNTIME]

Tracing: torch.matmul(x, w).cpu() · click any layer to expand

Kernel Launch Overhead

04.C

[DRIVER]

HOST → DEVICE

~5 μs

cudaLaunchKernel copies args, validates, signals GPU

CPU ↔ GPU (PCIe 4.0 x16)

~32 GB/s

H→D / D→H transfer bound; often the dominant cost

NVLink (multi-GPU)

~900 GB/s

Hopper NVLink: 18 links × 50 GB/s

ℹ TensorRT fuses kernels: one matmul + one bias-add + one GeLU = single kernel launch. CUDA Graphs capture whole forward passes to remove launch overhead entirely.

STAGE_05

The GPU Execution Model

The GPU is a hierarchy of parallelism. A single matmul launches thousands of threads, mapped across Streaming Multiprocessors.

GPU Die

05.A

[HARDWARE]

GPC → TPC → SM · hover any block

GPC

Graphics Processing Cluster

TPC

Texture Processing Cluster

Streaming Multiprocessor

Inside an SM

05.B

[HARDWARE]

The SM is the atomic unit. 100+ warps can be co-resident, scheduled in groups of 4. Hover any block.

SIMT

Single Instruction, Multiple Threads

WARP

32 threads, lockstep

Thread Hierarchy

05.C

[RUNTIME]

Grid

All blocks for one kernel

dim3 grid(8, 8)

Block

64+

Runs on one SM, has shared memory

dim3 block(32, 8)

Warp

32 threads

SIMT — same instruction, lockstep

Always 32 threads

Thread

One program counter, own registers

tid = threadIdx.x

Occupancy Calculator

05.D

[RUNTIME]

Occupancy = active warps / max warps per SM

THREADS / BLOCK256

REGS / THREAD32

SHARED MEM / BLOCK (KB)48

Threads

blocks/SM

Registers

blocks/SM

Shared Mem

blocks/SM

Bottleneck

Shared Mem

limits to 4

THEORETICAL OCCUPANCY50%

1,024 active threads / 2048 max

STAGE_06

Tensor Cores

Specialized matrix-multiply-accumulate silicon. A single Tensor Core does 128 FP16 FMA per cycle — and Hopper's FP8 doubles that.

MMA in One Cycle

06.A

[HARDWARE]

A 4×4 (FP16)

B 4×4 (FP16)

D 4×4 (FP32)

⚡ One Tensor Core MMA: D = A × B + C — completes in 1 clock cycle

Ampere/Hopper: 16×8×16 fragment size per warp instruction. 4× such units per SM.

PTX / SASS

06.B

[ISA]

The Tensor Core is programmed via warp-level matrix intrinsics.

# CUDA C++ (CUDA 12+)

#include <cuda/barrier>

#include <cuda/mma>

using nvcuda::wmma;

// 16×16×16 fragment

fragment matrix_a<16, 16, 16, half, row_major> a_frag;

fragment matrix_b<16, 16, 16, half, col_major> b_frag;

fragment accumulator<16, 16, 16, float> c_frag;

load_matrix_sync(a_frag, A, 16);

load_matrix_sync(b_frag, B, 16);

mma_sync(c_frag, a_frag, b_frag, c_frag);

# Generated SASS (Hopper)

HMMA.16832.F32 a0, a1, b0, d0, d1

ℹ torch.backends.cuda.matmul.allow_tf32 = True routes FP32 matmul through TF32 Tensor Cores — silent ~8× speedup.

Throughput by Precision

06.C

[HARDWARE]

Numbers are H100 SXM5, peak TFLOPS (dense)

FP32 CUDA Core

60 TFLOPS

TF32 Tensor Core

500 TFLOPS

FP16 Tensor Core

1000 TFLOPS

BF16 Tensor Core

1000 TFLOPS

INT8 Tensor Core

2000 TFLOPS

FP8 Tensor Core

4000 TFLOPS

NVIDIA

Tensor Core

Volta → Hopper

AMD

Matrix Core

CDNA · RDNA 3+

Intel

AMX Tile

Sapphire Rapids+

STAGE_07

The CPU Path

On CPU, parallelism comes from wide SIMD registers and many cores, not thousands of threads. AVX-512 + AMX tile are the equivalent of a Tensor Core.

AVX-512 SIMD Lanes

07.A

[HARDWARE]

AVX-512: 16 × FP32 multiply-add in 1 cycle

× FMA

↓

⚡ Single instruction:

_mm512_fmadd_ps(a, b, c) // a*b+c, all 16 lanes

→ 16 multiplies + 16 adds in 2 cycles (2 ports)

Cache Hierarchy

07.B

[HARDWARE]

Registers

~1 KB · per thread

0 cycles

32–64 KB · per core, ~5 ns

4 cycles

1–2 MB · per core, ~4 ns

12 cycles

32–128 MB · shared, ~14 ns

40 cycles

DRAM (DDR5)

32–256 GB · ~70 ns

200 cycles

⚠ L1 miss → L2 miss → L3 miss → DRAM miss. Each step is roughly 10× the previous. NUMA: cross-socket access is 2–3× slower.

CPU Die

07.C

[HARDWARE]

Mesh interconnect ties cores to a shared L3 slice. Hover any block.

x86

AVX-512 / AMX

ARM

NEON / SVE / SME

Apple

AMX / M-series GPU

Pipeline

07.D

[ISA]

Out-of-order execution: 5+ GHz, ~14 stage pipeline, branch predictor hits 95%+.

# Intrinsics (C++)

__m512 a = _mm512_load_ps(A);

__m512 b = _mm512_load_ps(B);

__m512 c = _mm512_load_ps(C);

c = _mm512_fmadd_ps(a, b, c);

_mm512_store_ps(C, c);

# Generated asm (gcc -O3 -mavx512f)

vmovaps (%rdi),%zmm0

vmovaps (%rsi),%zmm1

vfmadd231ps (%rdx),%zmm1,%zmm0,%zmm0

vmovaps %zmm0,(%rdx)

FMA

Fused Multiply-Add

OoO

Out-of-Order Execution

ROB

Reorder Buffer ~512

IPC

~4 instructions/cycle

STAGE_08

Memory & the Roofline

The fastest matmul is the one that didn't have to read memory. Modern inference is bounded by data movement, not FLOPs.

Roofline Model

08.A

[ISA]

ARITHMETIC INTENSITY (FLOPs / byte)200 F/B

DTYPE

16-bit float · 2 B/elem · 1000 TFLOPS peak

# classification

AI = 200 → MEMORY-BOUND

At ridge point (299 F/B ) the chip saturates both.

Memory Waterfall

08.B

[HARDWARE]

Bandwidth increases 4–5× per tier. Latency decreases 10×. Hover any row.

HBM3 (GPU)

80 GB

3.35 TB/s

L2 (GPU)

50 MB

5.5 TB/s

L1 / SMEM

228 KB

~19 TB/s

Registers

~256 KB/SM

TB+/s

ALU / Tensor

32–64 lanes

compute

ℹ Quantization (FP32→INT8) shrinks bytes by 4× → AI goes up 4× → moves the point right on the roofline.

Interconnect

08.C

[DRIVER]

PCIe 5.0 ×16

128 GB/s

CPU ↔ GPU

NVLink 4

900 GB/s

Multi-GPU

Infinity Fabric

~36 GB/s

AMD CPU ↔ GPU

CXL 2.0

~64 GB/s

CPU ↔ pooled memory

STAGE_09

End-to-End Timeline

One forward pass of a single transformer block, from Python call to Python string. Each bar is a real kernel, each color a kind of work.

Gantt: One Token Forward Pass (GPT-2 ~125M)

09.A

[DRIVER]

0 μs50 μs100 μs150 μs200 μs

HOST

LAUNCH

TRANSFER

COMPUTE

SYNC

↑ CLICK ANY BAR
TO INSPECT KERNEL
+ HARDWARE

STAGE_10

The Output

The matmul produced a vector of 50,257 raw numbers. Softmax turns them into a probability distribution. Sampling picks the next token.

Raw Logits

10.A

[FRAMEWORK]

logits = X @ W_lm + b_lm // [batch, seq, 50,257]

0the

1.33

1 a

-0.99

2 of

2.66

3 and

2.75

4 to

-2.19

5 in

-5.80

6 is

-2.78

7 that

6.34

8 it

4.93

9 for

0.84

10 with

-2.49

11 as

-0.68

12 on

2.43

13 be

0.95

14 by

-0.81

15 are

-4.84

16 was

0.28

17 at

5.39

18 have

4.17

19 has

-1.17

20Hello

-3.32

21 world

2.67

22 !

2.08

23 ?

-0.45

24 ...

-4.16

25 yes

-2.69

26 no

3.16

27 maybe

5.96

28 I

5.02

29 you

-3.43

30 silicon

-3.62

31 neural

0.40

32 network

1.93

33 think

-1.20

34 compute

-3.45

35 tensor

2.82

36 core

4.96

37 data

4.82

38 model

-0.95

39 inference

-5.12

seed = 42

Softmax → Sample

10.B

[FRAMEWORK]

P(i) = exp(logit_i / T) / Σ exp(logit_j / T)

TEMPERATURET = 1.00

sharp· T → 0 greedy· T → ∞ uniformflat

#1 that

29.42%

#2 maybe

20.25%

#3 at

11.42%

#4 I

7.89%

#5 core

7.44%

#6 it

7.22%

#7 data

6.45%

#8 have

3.37%

TOP_K8

TOP_P0.90

Decoded Output

10.C

[FRAMEWORK]

SELECTED TOKEN

id = #7

" that"

p = 29.42%

→

PROMPT (cumulative)

"The neural network " that

→ Next forward pass: this token is appended to input, KV cache is reused, one new token is generated. Repeat.

VOCAB

50,257

ENTROPY

2.14 nats

TOP_P COVERAGE

90%

SAMPLE SPACE

7 tokens

TRACE COMPLETE

The round trip took microseconds.

A single token touched a tokenizer, an embedding table, twelve transformer blocks, a softmax, and a vocabulary lookup. The hardware traversed Python, CUDA driver, PTX, SASS, Tensor Cores, SMEM, registers, HBM, and back. Every stage exists for a reason — and removing any of them makes the network 10–1000× slower.

STAGES=10LAYERS=PYTHON → SILICONCOMPONENTS=80+DONE