The study blog from marearts.com
Computer Vision & Machine Learning Research Laboratory
def sensor_id(sensor: Sensor | int | None) -> int has three bars in one signature and
looks like bit arithmetic. It is not arithmetic and it is not new syntax to memorise: each bar is the
word or, and the whole annotation is one sentence — "a Sensor, or an
int, or nothing".
MareArts ANPR is a license-plate recognition app for iOS and Android, with a web viewer and a desktop viewer on the same account. Plates are read on the phone. Cloud MMC (make, model, colour, type, side, nation), team sync, rules, maps, and stats sit on marearts.com.
Two acronyms one letter apart, describing promises made at two different moments. An API is what you agree to when you write the call. An ABI is what your compiled bytes agree to when they meet somebody else's compiled bytes — and unlike the API, breaking it usually produces no error at all.
C forgets the name of every field the moment it finishes compiling. The X-macro is how a C
programmer gets those names back — by writing the list once and replaying it as an enum, a string
table, a descriptor table, and a set of registration calls that run before main.
def append_to(item, target=[]) looks like it starts from an empty list every time. It does
not. One list is created when the def statement runs, and every call for the rest of the
program's life shares it.
You define a function inside a loop, collect the functions, call them later — and every single one returns the last value. The usual fix is a strange-looking line where the same name appears twice. Here is what is actually going on, and why the fix works.
Most pytest tutorials stop right after assert 1 + 1 == 2, which is roughly where the
interesting part begins. This one keeps going — through fixtures, parametrization and the
configuration that makes a suite pleasant to live with, and out the other side into hooks and
plugins.
When you read Python code, sooner or later you hit a line like this:
data: Dict[str, int] = {}The : Dict[str, int] part looks strange the first time. In this post we'll
figure out exactly what that one line means — and whether a "type hint" is actually
enforced — using tiny runnable examples.
That single line is really three pieces glued together.
data : Dict[str, int] = {}
│ │ │
name type hint actual
(a note) value| Piece | Meaning |
|---|---|
data | the variable name |
: Dict[str, int] | a type hint — "keys are str, values are int". Just a note for humans/editors |
= {} | the actual value = an empty dict |
At runtime it is exactly the same as:
data = {} # drop the hint and this is all that's left= {} is just a dictA dict is a "key → value" store. It has nothing to do with type hints — it's a basic built-in type that has always existed.
data = {} # empty dict
data["gemm"] = "A" # key (str) -> value
data["conv2d"] = "B"
print(data) # {'gemm': 'A', 'conv2d': 'B'}
print(data["gemm"]) # A
print(list(data)) # ['gemm', 'conv2d']: Dict[str, int] is a "type hint"Dict[str, int] is a note saying "a dict whose keys are strings and values
are integers". Python does not enforce that note when it runs.
from typing import Dict
ages: Dict[str, int] = {} # at runtime this is just ages = {}
ages["alice"] = 30
ages["bob"] = 25
print(ages) # {'alice': 30, 'bob': 25}Violating the hint does not raise an error. Let's prove it.
from typing import Dict
ages: Dict[str, int] = {} # promise: "str keys, int values"
ages["alice"] = 30 # keeps the promise
ages["oops"] = "not-int" # value is a string -> breaks it
ages[123] = 99 # key is an int -> breaks it
print(ages)Output:
If they're not enforced, why bother? Because they're a "quality tool", not a rule.
mypy file.py separately and it catches
violations before you run the code.mypy yourself — the interpreter won't.
Dict vs dict — and when did this syntax appear?The dict itself is old; the type-hint syntax is the newer part.
| Syntax | Import needed? | Introduced |
|---|---|---|
Dict[str, int] (capital) | from typing import Dict | Python 3.5 (2015) |
dict[str, int] (lowercase) | none (built-in) | Python 3.9 (2020) |
dict (plain) | none | always |
Annotating a variable with x: int = 0 has been possible since Python 3.6 (2016).
Both forms mean the same thing, so new code usually prefers the lowercase dict[...].
scores: dict[str, float] = {} # 3.9+ : lowercase, no importname: Type = value == name = value + (a type note)= {} is just an empty dict (old, ordinary syntax): Dict[str, int] is a hint that is not enforced — breaking it still runsmypy; the checking is done by mypy, separatelyA GPU kernel launch is a hierarchy of work units. Bigger units contain smaller ones.
| AMD term | NVIDIA term | What it is |
|---|---|---|
| Grid | Grid | The whole kernel launch — covers the entire problem. |
| Workgroup | Block / threadblock | A group of threads on one Compute Unit (CU). Shares LDS (shared memory). Can synchronize via __syncthreads(). |
| Wavefront (wave) | Warp | 64 threads (AMD CDNA) executing the same instruction simultaneously (SIMT). |
| Thread (work-item) | Thread | One lane in a wave. Has its own thread ID and register state. |
"Lane" and "thread" are two views of the same physical execution slot.
One lane runs one thread at a time. They are 1:1 within an executing wave.
1 wave (= 64 threads) running on 1 SIMD:
Lane 0 ↔ Thread 0 (running my_function with tid=0)
Lane 1 ↔ Thread 1 (running my_function with tid=1)
Lane 2 ↔ Thread 2
...
Lane 63 ↔ Thread 63
All 64 lanes execute the same instruction at the same cycle.
Below the software hierarchy is the physical hardware:
GPU ├─ CU 0 │ ├─ SIMD 0 (64 lanes, ≤ 8 resident waves) │ ├─ SIMD 1 (64 lanes, ≤ 8 resident waves) │ ├─ SIMD 2 (64 lanes, ≤ 8 resident waves) │ ├─ SIMD 3 (64 lanes, ≤ 8 resident waves) │ └─ LDS (shared memory, 64 KB) ├─ CU 1 ├─ ... └─ CU 303 ← total 304 CUs on MI300X
| Type | Size | Scope | Notes |
|---|---|---|---|
| VGPR (vector GPR) | 32-bit (4 B) | Private per lane | Up to 256 per lane per wave. Each lane sees its own VGPR. |
| SGPR (scalar GPR) | 32-bit (4 B) | Shared by 64 lanes | Used for scalar values like loop counters, addresses. |
| AGPR (accumulator GPR) | 32-bit (4 B) | Private per lane | CDNA-only. Used as MFMA accumulators. |
1 lane × 256 VGPRs × 4 bytes = 1 KB per lane 1 wave (64 lanes) × 1 KB = 64 KB total register file used by one wave
Occupancy = number of waves resident on a SIMD (1 to 8). Higher occupancy enables better latency hiding.
If a wave uses 256 VGPRs/lane → only 1 wave fits in SIMD → occupancy 1 If a wave uses 128 VGPRs/lane → 2 waves fit → occupancy 2 If a wave uses 32 VGPRs/lane → 8 waves fit → occupancy 8 (max) More resident waves = SIMD can switch when one wave stalls on memory.
GPU memory has multiple levels, similar to CPU cache hierarchy:
| Level | Size (per CU / total) | Latency (cycles) | Managed by | CPU analogue |
|---|---|---|---|---|
| Registers (VGPR/SGPR) | ~256 KB / CU | ~1 | Compiler | CPU registers |
| LDS (shared memory) | 64 KB / CU | ~10–30 | Software (explicit loads/stores) | Scratchpad / fast SRAM |
| L1 cache | 16 KB / CU | ~30 | Hardware (transparent) | L1 cache |
| L2 cache | ~16 MB total | ~150 | Hardware | L2 cache |
| Infinity / L3 | ~256 MB | ~300 | Hardware | L3 cache |
| HBM (global memory) | 192 GB | ~500–1000 | HW + software | DRAM |
<<< grid, block >>>HIP/CUDA kernel launch syntax:
add_kernel<<< grid, block >>>(A, B, C, N);
| Parameter | Meaning | Example |
|---|---|---|
block (a.k.a. blockSize) | Threads per workgroup | 256 → 4 waves per workgroup |
grid | Number of workgroups | 4 → 4 workgroups total |
| (implicit) | Wave count = blockSize / 64 | 256 / 64 = 4 waves per workgroup |
You don't pick the wave count directly — it is derived from blockSize. The hardware always groups threads into waves of 64 on CDNA.
add_kernel<<< grid=4, block=256 >>>(...) Total threads = 4 × 256 = 1024 Total waves = 1024 / 64 = 16 Total workgroups = 4 Each workgroup → one CU Each wave → one SIMD inside that CU
A single load instruction can pull up to 16 bytes into a thread's registers. This is the hardware limit on AMD CDNA.
The number of elements per load (called vec size) depends on the data type:
| Data type | Size (B) | vec=1 | vec=2 | vec=4 | vec=8 | vec=16 |
|---|---|---|---|---|---|---|
| fp16 / bf16 | 2 | 2 B | 4 B | 8 B | 16 B (max) | — |
| fp32 | 4 | 4 B | 8 B | 16 B (max) | — | — |
| int8 | 1 | 1 B | 2 B | 4 B | 8 B | 16 B (max) |
vec × sizeof(dtype) ≤ 16 byte. Larger vec means fewer load instructions to move the same amount of data — faster.
fp16, vec=1 (2 B): global_load_ushort v0, v[1:2] ; load 2 bytes (1 fp16) fp16, vec=4 (8 B): global_load_dwordx2 v[0:1], v[2:3] ; load 8 bytes (4 fp16) fp16, vec=8 (16 B): global_load_dwordx4 v[0:3], v[4:5] ; load 16 bytes (8 fp16) - MAX fp32, vec=4 (16 B): global_load_dwordx4 v[0:3], v[4:5] ; load 16 bytes (4 fp32) - MAX
__global__ void add_kernel(float* A, float* B, float* C, int N) {
int tid = blockIdx.x * blockDim.x + threadIdx.x; // global thread ID
if (tid < N) {
float a = A[tid]; // load (HBM → register)
float b = B[tid]; // load (HBM → register)
float c = a + b; // ALU (register-to-register, ~1 cycle)
C[tid] = c; // store (register → HBM)
}
}
add_kernel<<< 4, 256 >>>(A, B, C, 1024);
What happens per cycle (assuming occupancy 1, the worst case):
cycle: 1 2..99 100 101..199 200 201
---- ------ ---- -------- ---- -----
inst: load wait load wait add store
A (idle) B (idle)
↓
ALU only busy 2 cycles out of 200.
Instructions in this kernel: 4. Actual cycles: ~200. The reason: each HBM load takes ~100 cycles to complete, even though issuing it takes 1 cycle. With occupancy 1, the SIMD has nothing else to do but wait.
MFMA (Matrix Fused Multiply-Add) instructions are wave-cooperative: all 64 lanes work together to compute a small matrix multiply (e.g., 16×16).
v_mfma_f32_16x16x16_f16 acc, a_frag, b_frag, c_frag
64 lanes cooperatively compute D = A × B + C
where A is 16×16 fp16, B is 16×16 fp16, D is 16×16 fp32
Each lane holds a small piece of A, B, and accumulates a small piece of D.
The hardware exchanges data between lanes during execution.
Latency: ~8–32 cycles (NOT 1 cycle), but throughput is enormous:
16×16×16 = 4,096 multiply-adds per instruction per wave
A tile is a 2D chunk of a matrix that one workgroup (or one wave) processes. GPU kernels divide a big problem into many small tiles.
Big matrix (e.g., 1024 × 1024)
divided into tiles of 64 × 64:
X →
┌────────────────────────┐
Y │ t0 t1 t2 ... t15 │
↓ │ t16 t17 ... │ 16 × 16 = 256 tiles
│ ... │ each handled by one workgroup
│ t240 ... t255 │
└────────────────────────┘
One 64×64 tile = 4,096 elements. One wave = 64 threads. Each thread is responsible for 4096/64 = 64 elements.
Those 64 elements per thread are split between two axes:
| Symbol | Meaning |
|---|---|
| X1 (= vec) | Number of elements one thread loads in one instruction (X direction). |
| X0 | Number of threads placed along the X axis. |
| Y0 | Number of threads placed along the Y axis. |
| Y1 | Number of times each thread iterates along the Y axis. |
Constraints:
X0 × Y0 = 64 ← total threads (wave size) X0 × X1 = XPerTile ← X axis fully covered Y0 × Y1 = YPerTile ← Y axis fully covered (with iteration) X1 × sizeof ≤ 16 byte ← hardware load limit
Memory is 1D, but we view it as 2D (row-major):
memory: [a][b][c][d] [e][f][g][h] [i][j][k][l] [m][n][o][p]
─────────── ─────────── ─────────── ───────────
row 0 row 1 row 2 row 3
X direction: addresses +1 (contiguous) → one instruction can load 4/8/16 bytes
Y direction: addresses +width (strided) → needs separate instructions per row
X1 = 4 (vec) X0 = XPerTile / X1 = 64 / 4 = 16 Y0 = wave / X0 = 64 / 16 = 4 Y1 = YPerTile / Y0 = 64 / 4 = 16 Per thread: X1 × Y1 = 4 × 16 = 64 elements Per wave: 64 threads × 64 elements = 4,096 elements ✓ Load count: 16 (per thread) → 1024 total vec loads → 4,096 elements
X axis (16 threads × 4 vec = 64 cols)
┌─────────────────────────────────────┐
row 0..3 │ T0 T1 T2 T3 T4 T5 ... T14 T15 │ ← Y0=0, Y1=0..3
row 4..7 │ T0 T1 T2 T3 ... │ ← Y0=0, Y1=4..7
... Y1 iterates 16 times
row 16.. │ T16 T17 ... │ ← Y0=1
row 32.. │ T32 T33 ... │ ← Y0=2
row 48..63 │ T48 T49 ... T63 │ ← Y0=3
└─────────────────────────────────────┘
| vec | X0 | Y0 | Y1 | Loads / thread | Quality |
|---|---|---|---|---|---|
| 1 | 64 | 1 | 64 | 64 | worst (no vec) |
| 2 | 32 | 2 | 32 | 32 | poor |
| 4 | 16 | 4 | 16 | 16 | good |
| 8 | 8 | 8 | 8 | 8 | best (fp16 max) |
The same tile can be distributed across threads in several ways. The choice depends on the algorithm, the data layout, and the matrix instruction (MFMA) shape.
| Pattern | Who covers one tile | Wave layout inside tile |
|---|---|---|
thread_raked | 1 wave (64 threads) | n/a (single wave) |
warp_raked | Multiple waves cooperate | 1D stack (waves stripe along one axis) |
block_raked | All waves of the workgroup | 2D grid (waves arranged in a grid) |
X axis (XPerTile)
┌─────────────────────────────────┐
│ Wave 0 │ each wave covers full X width
├─────────────────────────────────┤
│ Wave 1 │ 1/4 of Y
├─────────────────────────────────┤
│ Wave 2 │
├─────────────────────────────────┤
│ Wave 3 │
└─────────────────────────────────┘
X axis
┌───────────────┬─────────────┐
│ Wave 0 │ Wave 1 │
│ (X 0..63) │ (X 64..127) │
├───────────────┼─────────────┤
│ Wave 2 │ Wave 3 │
│ (X 0..63) │ (X 64..127) │
└───────────────┴─────────────┘
| warp_raked (1D) | block_raked (2D) | |
|---|---|---|
| Wave layout | 1 axis (Y) | 2 axes (X × Y) |
| Sub-tile per wave | XPerTile × (YPerTile/N) | (XPerTile/M) × (YPerTile/M) |
| X coverage by one wave | full | partial |
| Concept | Definition |
|---|---|
| Grid | The whole kernel launch. |
| Workgroup (block) | Group of threads on one CU. Shares LDS. |
| Wave (warp) | 64 threads (CDNA) executing in lockstep (SIMT). |
| Thread (work-item) | One software unit; runs on one lane. |
| Lane | One hardware ALU slot in a SIMD. 64 per SIMD. |
| SIMD | Hardware execution unit. 4 per CU. Can hold up to 8 resident waves. |
| CU | Compute Unit. Contains 4 SIMDs and 64 KB LDS. ~304 per MI300X. |
| VGPR | Vector register, private per lane. 32-bit. Up to 256 per lane. |
| SGPR | Scalar register, shared by all 64 lanes in a wave. |
| AGPR | Accumulator register (CDNA), used for MFMA output. |
| LDS | Local Data Share = software-managed shared memory. 64 KB per CU. |
| HBM | Global memory. Large but slow (~500+ cycle latency). |
| Occupancy | Number of resident waves on a SIMD (1–8). Higher = better latency hiding. |
| Latency hiding | SIMD switches to another resident wave while one waits on memory. |
| SIMT | Single Instruction Multiple Threads. All 64 lanes run the same instruction. |
| MFMA | Matrix Fused Multiply-Add. Wave-cooperative matrix multiply (CDNA). |
| WMMA | RDNA equivalent of MFMA. |
| Tile | 2D chunk of a matrix processed by one workgroup. |
| XPerTile / YPerTile | Tile dimensions in elements (algorithm-defined). |
| vec / X1 | Elements per thread per load. Constrained by 16-byte limit. |
| X0, Y0 | Number of threads placed along X / Y axes. |
| Y1 | Y-axis iteration count per thread. |
| Coalescing | Adjacent threads accessing adjacent memory → one wide HBM transaction. |
| Kernel launch | kernel<<< grid, block >>>(...) — grid = workgroups, block = threads/wg. |
| tile_distribution_pattern | How threads/waves are distributed across one tile (thread/warp/block raked). |