7/28/2026

pytest, from your first test to your own plugin

Python · Testing

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.

There is a moment in most Python projects where the test suite stops helping. It takes four minutes to run, three tests fail intermittently, and nobody remembers what conftest.py is for. The tool is rarely the problem — pytest has an answer for each of those, and the answers are small.

So this is written as a path rather than a reference. Part 1 gets you productive with nothing but functions and assert. Parts 2 and 3 cover fixtures and parametrization, which is where pytest stops being "unittest with less typing" and starts being a different tool. Part 4 is the boring, high-leverage material: layout, configuration, mocking, CI. Part 5 opens the hood.

Every runnable snippet here was executed against pytest 9.1.1 before publishing, and the terminal output is copied from real runs. Anything version-sensitive is flagged inline.

Part 1 · Basics

Why pytest

Python ships with unittest, a port of Java's JUnit. It works, but it makes you inherit from a base class, remember a different assert method for every comparison, and wrap setup logic in setUp methods that run for every test whether you need them or not.

pytest replaces all of that with plain functions and the plain assert keyword.

unittest
import unittest

class TestMath(unittest.TestCase):
    def setUp(self):
        self.values = [1, 2, 3]

    def test_sum(self):
        self.assertEqual(sum(self.values), 6)

    def test_membership(self):
        self.assertIn(2, self.values)
        self.assertNotIn(9, self.values)
pytest
import pytest

@pytest.fixture
def values():
    return [1, 2, 3]

def test_sum(values):
    assert sum(values) == 6

def test_membership(values):
    assert 2 in values
    assert 9 not in values

The pytest version has no class, no inheritance, and no assertion vocabulary to memorise. It is also strictly more capable: values is a fixture, and unlike setUp it only runs for tests that actually ask for it.

pytest also runs unittest test cases as-is, so adopting it never requires a big-bang migration.

Your first test

install
python -m pip install pytest

Create two files side by side.

slugify.py
import re

def slugify(title: str) -> str:
    """Turn a human title into a URL-safe slug."""
    cleaned = re.sub(r"[^\w\s-]", "", title.lower())
    return re.sub(r"[\s_]+", "-", cleaned).strip("-")
test_slugify.py
from slugify import slugify

def test_lowercases_and_joins_words():
    assert slugify("Hello World") == "hello-world"

def test_strips_punctuation():
    assert slugify("What's new?!") == "whats-new"

def test_collapses_whitespace():
    assert slugify("too    many   spaces") == "too-many-spaces"
$ pytest -q
...                                                            [100%]
3 passed in 0.01s

Three dots, three passing tests. There is no registration step, no test suite object, no if __name__ == "__main__". pytest found the file, found the functions, and ran them.

Run it from the right place

pytest inserts the test file's directory (technically its rootdir for that file) into sys.path, which is why from slugify import slugify works here. For real projects use the layout in Project layout instead of relying on this.

How tests are found

Discovery is convention-driven. pytest walks the directories you give it (or the current one) and collects:

LevelDefault ruleConfig key
Filestest_*.py or *_test.pypython_files
ClassesTest*, and it must not define __init__python_classes
Functionstest*, at module level or inside a collected classpython_functions

Ask pytest what it would run without running anything. This is the single most useful command when a test "isn't running":

dry run
pytest --collect-only -q
Common trap

A test class with an __init__ method is silently skipped — pytest cannot instantiate it, so it warns and moves on. If a whole class mysteriously vanishes, that is almost always why.

The assert statement

pytest rewrites the bytecode of your test modules at import time so that a failing assert reports the value of every subexpression. You get JUnit-grade diagnostics from the plain keyword.

test_report.py
def build_report():
    return {"rows": 41, "status": "ok", "tags": ["a", "b"]}

def test_report_shape():
    report = build_report()
    assert report == {"rows": 42, "status": "ok", "tags": ["a", "c"]}
$ pytest -q
    def test_report_shape():
        report = build_report()
>       assert report == {"rows": 42, "status": "ok", "tags": ["a", "c"]}
E       AssertionError: assert {'rows': 41, ...': ['a', 'b']} == {'rows': 42, ...': ['a', 'c']}
E
E         Omitting 1 identical items, use -vv to show
E         Differing items:
E         {'rows': 41} != {'rows': 42}
E         {'tags': ['a', 'b']} != {'tags': ['a', 'c']}
E         Use -v to get more diff

test_report.py:6: AssertionError
1 failed in 0.02s

pytest knows how to diff dicts, lists, sets, strings and dataclasses. Add -vv when the diff is truncated — it disables the shortening.

Add a message only when it adds information

assert x == y, "x should equal y" is noise; the rewriting already showed you that. A good message supplies context the values cannot: assert resp.ok, f"upstream said {resp.text}".

Exceptions and warnings

Use pytest.raises as a context manager. The test fails if the block does not raise.

test_errors.py
import pytest

def withdraw(balance, amount):
    if amount <= 0:
        raise ValueError(f"amount must be positive, got {amount}")
    if amount > balance:
        raise ValueError("insufficient funds")
    return balance - amount

def test_rejects_negative_amount():
    with pytest.raises(ValueError):
        withdraw(100, -5)

def test_error_message_names_the_amount():
    # match= is a regex applied with re.search against str(exception)
    with pytest.raises(ValueError, match=r"must be positive, got -5"):
        withdraw(100, -5)

def test_inspect_the_exception_object():
    with pytest.raises(ValueError) as excinfo:
        withdraw(100, 500)
    assert "insufficient" in str(excinfo.value)
    assert excinfo.type is ValueError
Assert after the block, not inside it

Code placed after the raising line inside a with pytest.raises(...) block never executes. Anything you want to check about the exception goes after the block using excinfo.value.

Warnings have a mirror-image API:

warnings
import warnings, pytest

def legacy_api():
    warnings.warn("use new_api() instead", DeprecationWarning, stacklevel=2)
    return 42

def test_warns():
    with pytest.warns(DeprecationWarning, match="new_api"):
        assert legacy_api() == 42

Comparing floats

0.1 + 0.2 == 0.3 is False in binary floating point. pytest.approx wraps a value with a tolerance and works on scalars, sequences, dicts and numpy arrays.

approx
from pytest import approx

def test_scalar():
    assert 0.1 + 0.2 == approx(0.3)

def test_explicit_tolerance():
    assert 4.7 == approx(4.7005, abs=1e-3)     # absolute
    assert 1_000_100 == approx(1_000_000, rel=1e-3)  # relative

def test_containers():
    assert [0.1 + 0.2, 1 / 3] == approx([0.3, 0.33333333])
    assert {"loss": 0.1 + 0.2} == approx({"loss": 0.3})

The default is a relative tolerance of 1e-6 with a small absolute floor, which is the right default for most numeric code. Say what you mean when precision actually matters.

Running and selecting tests

CommandWhat it does
pytestEverything under the current directory
pytest tests/unitOne directory
pytest tests/test_api.pyOne file
pytest tests/test_api.py::test_loginOne test — this is a node id
pytest "tests/test_api.py::test_login[admin]"One parametrized case
pytest -k "login and not slow"Substring/boolean match on names
pytest -m integrationMatch a marker
pytest -xStop at the first failure
pytest --maxfail=3Stop after three failures
pytest -q / -v / -vvQuieter / one line per test / no truncated diffs
pytest -rASummary lines for every outcome, including passes
pytest -sDon't capture stdout (see your print calls live)
pytest --lfOnly the tests that failed last run
pytest --ffEverything, but last run's failures first
pytest --swStepwise: stop at a failure, resume there next time

Node ids are copy-pasteable. When a test fails, the id in the output is exactly the argument you need to re-run just that case.

Part 2 · Fixtures

Fixture basics

A fixture is a function that produces something a test needs. Tests request fixtures by naming them as parameters, and pytest supplies them. That is the entire mental model — it is dependency injection with a decorator.

request by name
import pytest

@pytest.fixture
def inventory():
    return {"apple": 3, "pear": 0}

def test_has_apples(inventory):        # ← the parameter name is the fixture name
    assert inventory["apple"] == 3

def test_pears_are_out_of_stock(inventory):
    assert inventory["pear"] == 0

Each test gets a fresh inventory by default, so one test mutating the dict cannot affect another. That isolation is the reason fixtures beat module-level constants.

Fixtures can request other fixtures, forming a graph that pytest resolves for you:

composition
@pytest.fixture
def config():
    return {"currency": "EUR", "vat": 0.21}

@pytest.fixture
def cart(config):                       # fixtures compose
    return {"config": config, "items": []}

def test_cart_uses_config(cart):
    assert cart["config"]["vat"] == 0.21

Setup and teardown

Replace return with yield and everything after the yield becomes teardown. It runs even if the test fails, because pytest wraps it in a finalizer.

yield fixture
import sqlite3
import pytest

@pytest.fixture
def db():
    conn = sqlite3.connect(":memory:")
    conn.execute("CREATE TABLE users (id INTEGER PRIMARY KEY, name TEXT)")
    yield conn                          # ← the test runs here
    conn.close()                        # ← always runs afterwards

def test_insert_and_read_back(db):
    db.execute("INSERT INTO users (name) VALUES (?)", ("ada",))
    assert db.execute("SELECT name FROM users").fetchone() == ("ada",)
Setup failures are errors, not failures

If the code before yield raises, pytest reports the test as an error rather than a failure, and the test body never runs. That distinction in the summary line tells you instantly whether your production code or your scaffolding broke.

Scopes

Expensive setup should not repeat per test. scope controls how long one instance is cached and shared.

ScopeCreated once perTypical use
function defaultTestAnything mutable
classTest classShared object for a group of methods
moduleFileA parsed file, a compiled artifact
packageDirectoryRare; a shared subsystem
sessionWhole runDocker container, server process, big download
session scope
@pytest.fixture(scope="session")
def http_server():
    proc = start_server()               # slow: do it once for the whole run
    yield proc
    proc.terminate()

@pytest.fixture
def client(http_server):                # cheap per-test wrapper over the shared server
    return Client(http_server.url)
The scope rule you will hit eventually

A fixture may only depend on fixtures of equal or wider scope. A session fixture requesting a function fixture is an error, because the narrow one would be destroyed while the wide one still holds it.

And never let a widely scoped fixture hand out something mutable. One test appending to a session-scoped list poisons every later test, producing failures that depend on run order.

conftest.py

Fixtures defined in conftest.py are visible to every test in that directory and below, with no import needed. Nested conftest.py files stack, and the closest definition wins.

tree
tests/
├── conftest.py            # fixtures for everything below
├── unit/
│   └── test_parser.py
└── integration/
    ├── conftest.py        # extra fixtures, only for integration tests
    └── test_api.py
tests/conftest.py
import pytest

@pytest.fixture(scope="session")
def sample_data():
    return {"users": ["ada", "grace"], "version": 3}

@pytest.fixture
def frozen_clock(monkeypatch):
    """Pin time.time() so timestamp assertions are deterministic."""
    import time
    monkeypatch.setattr(time, "time", lambda: 1_700_000_000.0)
    return 1_700_000_000.0
conftest.py is more than a fixture bucket

It is also where hooks, custom CLI options and plugin registration live. pytest imports it automatically — it must never be imported by hand, and it needs no __init__.py.

Where does a fixture come from? Ask:

list available fixtures
pytest --fixtures            # every fixture visible here, with docstrings
pytest --fixtures-per-test tests/test_api.py::test_login

Built-in fixtures

These ship with pytest. Reach for them before writing your own.

tmp_path — a real, empty directory per test

tmp_path
def test_writes_a_csv(tmp_path):
    target = tmp_path / "out.csv"       # tmp_path is a pathlib.Path
    target.write_text("id,name\n1,ada\n")

    assert target.exists()
    assert target.read_text().splitlines()[0] == "id,name"

pytest keeps the last few runs' directories on disk for post-mortem inspection and cleans up older ones. Use tmp_path_factory for a directory shared across a session.

monkeypatch — scoped, auto-undone patching

monkeypatch
def test_reads_token_from_environment(monkeypatch):
    monkeypatch.setenv("API_TOKEN", "secret")
    monkeypatch.delenv("HTTP_PROXY", raising=False)
    assert load_config().token == "secret"

def test_patch_an_attribute(monkeypatch):
    monkeypatch.setattr("myapp.net.fetch", lambda url: {"ok": True})
    assert myapp.summarize("http://x") == "ok"

def test_run_from_another_directory(monkeypatch, tmp_path):
    monkeypatch.chdir(tmp_path)         # undone automatically
    assert Path.cwd() == tmp_path

Every change is reverted when the test ends, including failures. That is the whole point: hand-rolled os.environ[...] = ... leaks into subsequent tests.

capsys and caplog — captured output

capsys / caplog
import logging

def test_prints_a_banner(capsys):
    print("READY")
    captured = capsys.readouterr()      # consumes the buffer
    assert captured.out == "READY\n"
    assert captured.err == ""

def test_logs_a_warning(caplog):
    with caplog.at_level(logging.WARNING):
        logging.getLogger("app").warning("disk almost full")
    assert "disk almost full" in caplog.text
    assert caplog.records[0].levelno == logging.WARNING

Use capfd instead of capsys when the output comes from a subprocess or a C extension.

request — introspection

request
@pytest.fixture
def workspace(request, tmp_path):
    """Name the directory after the test that asked for it."""
    d = tmp_path / request.node.name
    d.mkdir()
    return d

Factory fixtures

When a test needs several of a thing, or needs to choose parameters, return a function instead of a value. Track what you create so teardown stays correct.

factory-as-fixture
import pytest

@pytest.fixture
def make_user(db):
    created = []

    def _make(name, admin=False):
        cur = db.execute(
            "INSERT INTO users (name, admin) VALUES (?, ?)", (name, admin)
        )
        created.append(cur.lastrowid)
        return {"id": cur.lastrowid, "name": name, "admin": admin}

    yield _make

    for user_id in created:             # clean up exactly what this test made
        db.execute("DELETE FROM users WHERE id = ?", (user_id,))

def test_admins_can_see_everyone(make_user):
    admin = make_user("ada", admin=True)
    make_user("grace")
    make_user("alan")
    assert len(visible_users(admin)) == 3

Part 3 · Parametrizing

parametrize

One test function, many cases. Each case is a separate test: it gets its own node id, its own pass/fail, and one failing case does not hide the others.

basic
import pytest

@pytest.mark.parametrize(
    ("title", "expected"),
    [
        ("Hello World", "hello-world"),
        ("What's new?!", "whats-new"),
        ("  padded  ", "padded"),
        ("", ""),
    ],
)
def test_slugify(title, expected):
    assert slugify(title) == expected
$ pytest -v
test_slug.py::test_slugify[Hello World-hello-world] PASSED
test_slug.py::test_slugify[What's new?!-whats-new]   PASSED
test_slug.py::test_slugify[  padded  -padded]        PASSED
test_slug.py::test_slugify[-]                        PASSED

Stacking decorators produces the cartesian product:

3 × 2 = 6 tests
@pytest.mark.parametrize("codec", ["gzip", "zstd", "none"])
@pytest.mark.parametrize("size", [0, 1_000_000])
def test_roundtrip(codec, size):
    payload = b"x" * size
    assert decompress(compress(payload, codec), codec) == payload
Combinatorics grow fast

Stacking is convenient and dangerous. Three stacked decorators of five values each is 125 tests. Prefer an explicit list of meaningful tuples once the product stops being all-interesting.

Readable IDs

Auto-generated ids come from the values, which is unreadable for objects and awkward for long strings. Name your cases — the id is what you will read in CI output and paste on the command line.

ids=
@pytest.mark.parametrize(
    ("payload", "expected"),
    [
        ({"v": 1}, True),
        ({"v": 1, "extra": None}, True),
        ({}, False),
    ],
    ids=["minimal", "with-nulls", "empty"],
)
def test_validate(payload, expected):
    assert validate(payload) is expected

A dict keeps the name and the data in one place, which stops them drifting apart:

dict-driven cases
CASES = {
    "empty-input":     ("",        []),
    "single-token":    ("a",       ["a"]),
    "trailing-comma":  ("a,b,",    ["a", "b"]),
    "quoted-comma":    ('"a,b",c', ["a,b", "c"]),
}

@pytest.mark.parametrize(
    ("raw", "expected"), list(CASES.values()), ids=list(CASES.keys())
)
def test_parse_csv_line(raw, expected):
    assert parse_csv_line(raw) == expected

Per-case marks

pytest.param attaches a mark or an id to a single case, so one known-broken input does not force you to skip the whole function.

pytest.param
@pytest.mark.parametrize(
    ("raw", "expected"),
    [
        ("2024-01-01", date(2024, 1, 1)),
        pytest.param("01/01/2024", date(2024, 1, 1), id="us-format"),
        pytest.param(
            "2024-W01-1", date(2024, 1, 1),
            marks=pytest.mark.xfail(reason="ISO week dates not supported yet"),
        ),
        pytest.param(
            "yesterday", date(2024, 1, 1),
            marks=pytest.mark.skipif(sys.platform == "win32", reason="needs GNU date"),
        ),
    ],
)
def test_parse_date(raw, expected):
    assert parse_date(raw) == expected

Parametrized fixtures

Put params on the fixture and every test using it runs once per value. This is how you sweep a whole suite across backends without touching a single test body.

fixture params
@pytest.fixture(params=["sqlite", "postgres"], ids=["sqlite", "pg"])
def store(request):
    backend = request.param             # ← the current value
    s = open_store(backend)
    yield s
    s.close()

# both of these now run twice, once per backend
def test_put_then_get(store):
    store.put("k", "v")
    assert store.get("k") == "v"

def test_missing_key_returns_none(store):
    assert store.get("nope") is None

Indirect parametrization

indirect=True routes the parameter through a fixture rather than into the test. Use it when the value needs setup before the test can see it.

indirect
@pytest.fixture
def user(request):
    role = request.param                # comes from parametrize, not from the test
    return create_user(role=role)

@pytest.mark.parametrize("user", ["admin", "viewer"], indirect=True)
def test_dashboard_access(user):
    assert user.can_open_dashboard() is (user.role == "admin")

Mix direct and indirect by naming which arguments are indirect:

partial indirect
@pytest.mark.parametrize(
    ("user", "expected"),
    [("admin", True), ("viewer", False)],
    indirect=["user"],                  # only `user` goes through the fixture
)
def test_dashboard_access(user, expected):
    assert user.can_open_dashboard() is expected

Part 4 · Practical

Project layout

The src layout is the one that avoids import surprises: your package is not importable from the repository root, so tests are forced to import the installed package — the same thing your users get.

recommended
myproject/
├── pyproject.toml
├── src/
│   └── myproject/
│       ├── __init__.py
│       └── core.py
└── tests/
    ├── conftest.py
    ├── unit/
    │   └── test_core.py
    └── integration/
        └── test_end_to_end.py
install once, then test
python -m pip install -e ".[dev]"
pytest
Do not put __init__.py in tests/

Without it, test modules are imported as top-level modules — which means two files named test_utils.py in different directories will collide. Either give them distinct names, or add __init__.py files and accept package semantics. The distinct-names route is simpler.

Configuration

One config block removes a lot of repeated typing and makes local runs match CI. pyproject.toml is the modern home; pytest.ini and setup.cfg still work.

pyproject.toml
[tool.pytest.ini_options]
minversion = "8.0"
testpaths = ["tests"]
addopts = [
    "-ra",                 # summary for all non-passing outcomes
    "--strict-markers",    # unknown @pytest.mark.* becomes an error
    "--strict-config",     # typos in this very section become an error
    "--import-mode=importlib",
]
markers = [
    "slow: takes more than a second",
    "integration: needs external services",
]
filterwarnings = [
    "error",                                   # warnings fail the suite ...
    "ignore::DeprecationWarning:third_party.*" # ... except from code you don't own
]
OptionWhy you want it
--strict-markersA typo like @pytest.mark.slwo silently does nothing otherwise
filterwarnings = ["error"]Catches deprecations while they are cheap to fix
testpathsBare pytest stops walking .venv, build, docs
--import-mode=importlibModern import semantics; avoids sys.path surgery

Markers, skip and xfail

Markers are labels. You attach them to tests and select on them later.

marking
import pytest

@pytest.mark.slow
def test_full_reindex():
    ...

@pytest.mark.integration
class TestPaymentGateway:          # applies to every method in the class
    def test_charge(self): ...
    def test_refund(self): ...

pytestmark = pytest.mark.integration   # module-level: applies to the whole file
selecting
pytest -m slow                  # only slow
pytest -m "not slow"            # everything else
pytest -m "integration and not slow"

skip vs xfail

ConstructMeaningWhen
@pytest.mark.skip(reason=…)Never run itTemporarily irrelevant
@pytest.mark.skipif(cond, reason=…)Run only if the condition is falsePlatform / version / optional dependency
pytest.skip(reason)Skip from inside the testDecision needs runtime information
pytest.importorskip("numpy")Skip if the import failsOptional dependency
@pytest.mark.xfail(reason=…)Run it; a failure is expectedKnown bug, with a test that proves it
the useful forms
import sys, pytest

@pytest.mark.skipif(sys.version_info < (3, 12), reason="uses PEP 695 syntax")
def test_new_generics(): ...

def test_needs_a_gpu():
    if not gpu_present():
        pytest.skip("no GPU on this runner")
    ...

pandas = pytest.importorskip("pandas", minversion="2.0")

@pytest.mark.xfail(strict=True, reason="bug #482: rounds half-down")
def test_rounds_half_up():
    assert round_half_up(0.5) == 1
Prefer xfail(strict=True) over skip for known bugs

A strict xfail that starts passing becomes a failure. That is the feature: the day someone fixes the bug, the suite tells you to delete the marker. A skipped test just rots.

Mocking

Two tools, two jobs. monkeypatch replaces an attribute for the duration of a test. unittest.mock additionally records how the replacement was used.

monkeypatch: substitute behaviour
def test_retries_on_timeout(monkeypatch):
    calls = []

    def flaky_get(url):
        calls.append(url)
        if len(calls) < 3:
            raise TimeoutError
        return {"status": "ok"}

    monkeypatch.setattr("myapp.client.get", flaky_get)
    assert fetch_with_retry("http://x")["status"] == "ok"
    assert len(calls) == 3
mock: assert on the interaction
from unittest.mock import patch, MagicMock

def test_sends_exactly_one_email():
    with patch("myapp.mailer.send") as send:
        send.return_value = MagicMock(status=202)
        register_user("ada@example.com")

    send.assert_called_once_with(
        to="ada@example.com", template="welcome"
    )
Patch where it is looked up

If myapp/service.py does from myapp.mailer import send, then patching "myapp.mailer.send" has no effect — service already holds its own reference. Patch "myapp.service.send", the name in the module using it.

Mocking is a smell in proportion to its depth

Mocking a network boundary is good. Mocking five internal collaborators means the test now asserts your implementation's shape rather than its behaviour, and it will break on every refactor while catching no bugs. Consider a fake object or a real in-memory implementation instead.

Output and logging

pytest captures stdout, stderr and logging, then prints them only for failing tests. That is why your print statements seem to vanish on success.

FlagEffect
-sDisable capture entirely — output streams live
--capture=noSame as -s
--log-cli-level=DEBUGStream log records to the terminal as they happen
--show-capture=noHide captured output even on failures

Coverage

pytest-cov
python -m pip install pytest-cov

pytest --cov=myproject --cov-report=term-missing
pytest --cov=myproject --cov-report=html      # then open htmlcov/index.html
pytest --cov=myproject --cov-fail-under=85    # gate in CI
Coverage measures execution, not verification

A test that calls a function and asserts nothing still counts as full coverage of it. Treat the number as a way to find untested code, never as evidence that tested code is correct. Branch coverage (--cov-branch) is a strictly better signal than line coverage.

Speed

find and fix slowness
pytest --durations=10          # the 10 slowest tests, with setup/teardown split out

python -m pip install pytest-xdist
pytest -n auto                 # one worker per CPU
pytest -n 4 --dist loadfile    # keep each file on a single worker

Parallelism exposes hidden coupling: tests that only pass in a particular order will start failing. That is a bug being surfaced, not a bug being introduced.

Debugging a failure

the loop
pytest -x -q                   # 1. stop at the first failure
pytest --lf -x                 # 2. iterate on just that one
pytest --lf -x --pdb           # 3. drop into a debugger at the failure point
pytest --lf -x -vv --tb=long   # 4. full diffs and full traceback
pytest -q                      # 5. confirm the whole suite is green again
FlagTraceback style
--tb=longEvery frame, fully expanded default
--tb=shortOne line per frame — best for CI logs
--tb=lineOne line per failure
--tb=noNames only

breakpoint() works inside tests as long as capture is off (-s), and --pdb opens a debugger at the moment of failure with the whole frame still alive — usually faster than guessing where to put the breakpoint.

CI recipe

GitHub Actions
name: tests
on: [push, pull_request]

jobs:
  test:
    runs-on: ubuntu-latest
    strategy:
      fail-fast: false
      matrix:
        python-version: ["3.10", "3.11", "3.12", "3.13"]
    steps:
      - uses: actions/checkout@v4
      - uses: actions/setup-python@v5
        with:
          python-version: ${{ matrix.python-version }}
          cache: pip
      - run: python -m pip install -e ".[dev]"
      - run: pytest -ra --tb=short --durations=10 --cov --cov-report=xml -n auto

fail-fast: false matters: without it, one failing Python version cancels the others and hides whether the problem is version-specific.

Part 5 · Advanced

Hooks

pytest is a plugin system with a test runner attached. Every phase — collection, setup, running, reporting — publishes hooks, and conftest.py is a plugin that is loaded automatically. Implement a hook by defining a function with the right name.

conftest.py — auto-mark by location
def pytest_collection_modifyitems(config, items):
    """Everything under tests/integration/ gets @pytest.mark.integration."""
    import pytest
    for item in items:
        if "/integration/" in str(item.path):
            item.add_marker(pytest.mark.integration)
conftest.py — expose the outcome to fixtures
import pytest

@pytest.hookimpl(hookwrapper=True)
def pytest_runtest_makereport(item, call):
    outcome = yield
    report = outcome.get_result()
    # stash "rep_setup" / "rep_call" / "rep_teardown" on the test item
    setattr(item, f"rep_{report.when}", report)

@pytest.fixture
def browser(request):
    driver = start_browser()
    yield driver
    if getattr(request.node, "rep_call", None) and request.node.rep_call.failed:
        driver.save_screenshot(f"/tmp/{request.node.name}.png")
    driver.quit()

That pattern — capture an artifact only when the test failed — is the standard answer for browser screenshots, server logs and database dumps.

HookFires
pytest_addoptionOnce at startup, to register CLI flags
pytest_configureAfter config is read; register markers here
pytest_collection_modifyitemsAfter collection — reorder, filter, mark
pytest_generate_testsPer test function, to generate parameters
pytest_runtest_setupBefore each test
pytest_runtest_makereportAfter each phase, with the result
pytest_sessionfinishOnce, at the very end

Custom CLI options

conftest.py
import pytest

def pytest_addoption(parser):
    parser.addoption(
        "--runslow", action="store_true", default=False,
        help="run tests marked slow",
    )
    parser.addoption(
        "--env", action="store", default="staging",
        choices=("staging", "prod"), help="target environment",
    )

def pytest_configure(config):
    config.addinivalue_line("markers", "slow: takes more than a second")

def pytest_collection_modifyitems(config, items):
    if config.getoption("--runslow"):
        return
    skip = pytest.mark.skip(reason="needs --runslow")
    for item in items:
        if "slow" in item.keywords:
            item.add_marker(skip)

@pytest.fixture(scope="session")
def env(request):
    return request.config.getoption("--env")
usage
pytest                      # slow tests are skipped
pytest --runslow            # everything
pytest --env prod

pytest_generate_tests

When the set of cases is only known at runtime — read from a directory, a manifest, a database — generate them in this hook. Each generated case is still a first-class test.

golden files as test cases
# conftest.py
from pathlib import Path

CASES = Path(__file__).parent / "cases"

def pytest_generate_tests(metafunc):
    if "case_file" in metafunc.fixturenames:
        files = sorted(CASES.glob("*.input"))
        metafunc.parametrize(
            "case_file", files, ids=[f.stem for f in files]
        )

# test_render.py
def test_render_matches_golden(case_file):
    expected = case_file.with_suffix(".expected").read_text()
    assert render(case_file.read_text()) == expected

Dropping a new pair of files into cases/ adds a test. No code change, and the id is the file name so failures point straight at the data.

Writing a plugin

A plugin is a module of hooks and fixtures. Once it lives in a package with the pytest11 entry point, installing the package is enough — no imports, no conftest.py.

pytest_timer/plugin.py
import time
import pytest

@pytest.fixture
def timer():
    """Yield an object whose .elapsed is the time spent inside the block."""
    class _Timer:
        start = time.perf_counter()
        @property
        def elapsed(self):
            return time.perf_counter() - self.start
    return _Timer()

def pytest_addoption(parser):
    parser.addoption("--warn-slower-than", type=float, default=None)

@pytest.hookimpl(hookwrapper=True)
def pytest_runtest_call(item):
    started = time.perf_counter()
    yield
    limit = item.config.getoption("--warn-slower-than")
    took = time.perf_counter() - started
    if limit is not None and took > limit:
        item.warn(pytest.PytestWarning(f"{item.name} took {took:.2f}s"))
pyproject.toml
[project.entry-points.pytest11]
timer = "pytest_timer.plugin"

Test your plugin with pytest's own pytester fixture, which runs a nested pytest session in a temporary directory:

pytester
pytest_plugins = ["pytester"]

def test_warns_about_slow_tests(pytester):
    pytester.makepyfile("""
        import time
        def test_slow():
            time.sleep(0.05)
    """)
    result = pytester.runpytest("--warn-slower-than=0.01")
    result.assert_outcomes(passed=1, warnings=1)

Custom assertion helpers

Shared assertion functions clutter tracebacks with frames nobody cares about. Set __tracebackhide__ and the failure points at the caller instead.

__tracebackhide__
import pytest

def assert_valid_invoice(invoice):
    __tracebackhide__ = True            # hide this frame in the report
    if invoice.total < 0:
        pytest.fail(f"negative total: {invoice.total}")
    if invoice.lines and invoice.total != sum(l.amount for l in invoice.lines):
        pytest.fail(
            f"total {invoice.total} != sum of lines "
            f"{sum(l.amount for l in invoice.lines)}"
        )

def test_generated_invoice(order):
    assert_valid_invoice(build_invoice(order))   # ← failure is reported here

To keep pytest's value-introspection inside a helper module, register it for rewriting before it is imported:

conftest.py
import pytest
pytest.register_assert_rewrite("mypackage.testing.helpers")

Why plain assert works

Python's assert throws away everything but the boolean. pytest installs an import hook that rewrites the AST of test modules before they are compiled, replacing each assert with code that stores intermediate values and builds an explanation on failure.

conceptually
# what you write
assert compute(x) == expected

# roughly what gets compiled
@py_left  = compute(x)
@py_right = expected
if not (@py_left == @py_right):
    raise AssertionError(explain("==", @py_left, @py_right))

Three consequences worth knowing:

  • Rewriting applies to test modules, conftest.py, and modules you explicitly register. A plain library module keeps bare AssertionErrors with no detail.
  • Running Python with -O strips assert statements entirely, so never run a suite optimised.
  • Stale __pycache__ from a non-pytest run can shadow rewritten bytecode; pytest --cache-clear or deleting the caches fixes the "my assertions lost their detail" mystery.

Async tests

pytest cannot await a coroutine on its own — an async def test without a plugin is skipped with a warning, which looks like a pass in a hurry. Install one of the two plugins.

pytest-asyncio
# pyproject.toml
# [tool.pytest.ini_options]
# asyncio_mode = "auto"        # no per-test marker needed

import asyncio, pytest

@pytest.fixture
async def client():
    c = await open_client()
    yield c
    await c.aclose()

async def test_fetches_concurrently(client):
    a, b = await asyncio.gather(client.get("/a"), client.get("/b"))
    assert a.status == b.status == 200

anyio is the alternative when you need the same tests to run on both asyncio and trio.

Property-based testing

Instead of listing examples, describe the shape of valid input and assert a property that must hold for all of them. Hypothesis generates cases, and on failure shrinks them to the smallest reproducer.

hypothesis
from hypothesis import given, strategies as st

@given(st.text())
def test_slugify_is_idempotent(s):
    once = slugify(s)
    assert slugify(once) == once

@given(st.lists(st.integers()))
def test_sort_is_a_permutation(xs):
    out = my_sort(xs)
    assert len(out) == len(xs)
    assert sorted(out) == sorted(xs)
    assert all(a <= b for a, b in zip(out, out[1:]))

It composes with pytest normally — @given functions are collected like any other test. Good properties are round-trips (decode(encode(x)) == x), invariants (length, ordering) and equivalence to a slow reference implementation.

Don't mix @given with function-scoped fixtures

Hypothesis calls the test body many times, but a function-scoped fixture is created once for the whole thing. Shared mutable state leaks between generated examples. Build what you need inside the test, or use a factory fixture.

Flaky tests and isolation

A flaky test is worse than no test: it trains people to re-run CI instead of reading it. The usual causes are few, and each has a direct fix.

CauseFix
Shared mutable state across testsNarrow the fixture scope; return copies
Order dependencepytest -p no:randomly to confirm, then fix the coupling
Real clock / real timezoneFreeze time (monkeypatch, freezegun, time-machine)
Unseeded randomnessSeed it, and log the seed
sleep()-based waitingPoll for the condition with a timeout
Dict/set iteration order assumptionsCompare sets, or sort before comparing
Leftover files or env varstmp_path and monkeypatch, never manual mutation
prove isolation
python -m pip install pytest-randomly
pytest                       # shuffles order every run, prints the seed
pytest -p no:randomly        # reproduce the deterministic order
pytest -p randomly --randomly-seed=12345   # reproduce a specific shuffle
Retry plugins are anaesthetic, not treatment

pytest-rerunfailures can keep a pipeline moving, but a test that passes on the second attempt is telling you something real about your system. Retry only at genuinely non-deterministic boundaries (a live network), and never on unit tests.

Reference

Anti-patterns

Logic in the test

A test with branching needs its own tests. Split it into parametrized cases so each path is visible in the report.

avoid
def test_prices():
    for currency in ("EUR", "USD"):
        if currency == "EUR":
            assert fmt(1, currency) == "1,00 €"
        else:
            assert fmt(1, currency) == "$1.00"
prefer
@pytest.mark.parametrize(
    ("currency", "expected"),
    [("EUR", "1,00 €"), ("USD", "$1.00")],
)
def test_prices(currency, expected):
    assert fmt(1, currency) == expected

Asserting the implementation

mock.assert_called_once_with(...) on an internal helper locks in today's call graph. Rename the helper and the test fails while the behaviour is unchanged. Assert on the observable result instead.

One test, many concerns

When a 40-line test fails you learn "something in registration is broken". Several focused tests tell you which part, and the names double as documentation.

Depending on execution order

A test that only passes after another one ran is not a test, it is half of one. Each test must set up everything it needs.

Sleeping

avoid
start_worker()
time.sleep(2)          # slow AND flaky
assert job_done()
prefer
def wait_until(pred, timeout=5, interval=0.05):
    __tracebackhide__ = True
    deadline = time.monotonic() + timeout
    while time.monotonic() < deadline:
        if pred():
            return
        time.sleep(interval)
    pytest.fail(f"condition not met in {timeout}s")

start_worker()
wait_until(job_done)

Cheat sheet

Command line

pytest path::test_nameRun one test
-k "expr" / -m "expr"Select by name / by marker
-x / --maxfail=NStop early
--lf / --ff / --swLast-failed / failures-first / stepwise
-q / -v / -vv / -raVerbosity and outcome summary
-sShow output live
--pdbDebugger at the failure
--tb=shortCompact tracebacks
--collect-onlyList tests without running
--fixturesList available fixtures
--durations=10Slowest tests
-n autoParallel (xdist)
-W errorTurn warnings into failures

API

pytest.fixture(scope=, params=, autouse=, ids=)Define a fixture
pytest.mark.parametrize(argnames, argvalues, ids=, indirect=)Generate cases
pytest.param(*values, id=, marks=)One case with metadata
pytest.raises(Exc, match=)Expect an exception
pytest.warns(W, match=)Expect a warning
pytest.approx(value, rel=, abs=)Float comparison
pytest.skip(reason) / pytest.fail(msg)Skip / fail imperatively
pytest.importorskip(name, minversion=)Optional dependency
pytest.register_assert_rewrite(module)Rewrite asserts in a helper module

Built-in fixtures

tmp_path, tmp_path_factoryTemporary directories
monkeypatchPatch attributes, env vars, cwd — auto-undone
capsys, capfdCaptured stdout/stderr (Python level / fd level)
caplogCaptured log records
requestTest metadata, request.param, config access
pytesterRun a nested pytest — for testing plugins
recwarnRecord all warnings raised

Wrapping up

What actually moves the needle

Reading a tool's full feature list rarely changes how anyone works. If you adopt three things from this post, make them these.

The short list

  • Turn on --strict-markers and filterwarnings = ["error"]. Two lines of config that convert a class of silent failures into loud ones, today, in an existing project.
  • Replace your ad-hoc setup with fixtures, and keep them function-scoped until proven slow. Almost every order-dependent flake traces back to state shared more widely than it needed to be.
  • Reach for parametrize the moment a test grows a loop or an if. Each case gets its own name and its own verdict, which is the difference between "billing is broken" and "billing is broken for zero-quantity line items".

The advanced material matters less often, but it is worth knowing it exists. The day you find yourself copying the same conftest.py into a fourth repository, that is the signal to turn it into a plugin — and by then the entry point is a two-line change.

Further reading

  • docs.pytest.org — the reference; the "How-to guides" section is better than its name suggests
  • Plugin list — around 1,500 plugins, and the one you need probably exists
  • Hypothesis — property-based testing, covered briefly above
  • coverage.py — what pytest-cov wraps, worth reading directly for branch coverage
Every runnable snippet verified against pytest 9.1.1.

7/13/2026

Understanding Python variable type hints (Dict, dict, and mypy)

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.

1. One line, three things at once

That single line is really three pieces glued together.

data  :  Dict[str, int]  =  {}
 │            │              │
name       type hint       actual
            (a note)        value
PieceMeaning
datathe 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

2. = {} is just a dict

A 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']

3. : 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}

4. Is it really not enforced? — break it on purpose

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:

{'alice': 30, 'oops': 'not-int', 123: 99}
Key point: a type hint is not enforced. The Python interpreter simply ignores it at runtime, so breaking it still runs fine.

5. So why write hints at all?

If they're not enforced, why bother? Because they're a "quality tool", not a rule.

  • Humans — reading the code, you instantly see "this is a str→int dict".
  • Editors (IDEs) — autocomplete, and a red squiggle when you misuse it.
  • Static checkers (mypy) — run mypy file.py separately and it catches violations before you run the code.
A hint is "a promise + documentation", not a check. To actually catch violations, run a tool like mypy yourself — the interpreter won't.

6. Dict vs dict — and when did this syntax appear?

The dict itself is old; the type-hint syntax is the newer part.

SyntaxImport needed?Introduced
Dict[str, int] (capital)from typing import DictPython 3.5 (2015)
dict[str, int] (lowercase)none (built-in)Python 3.9 (2020)
dict (plain)nonealways

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 import

Summary

  • name: 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 runs
  • hints are documentation for humans / editors / mypy; the checking is done by mypy, separately

5/06/2026

AMD GPU Programming Primer — Threads, Waves, Tiles & Vector Loads

AMD GPU Programming Primer

Threads · Waves · Memory · Tile Distribution · Vector Loads · MFMA

1. The execution hierarchy: grid → workgroup → wave → thread

A GPU kernel launch is a hierarchy of work units. Bigger units contain smaller ones.

AMD termNVIDIA termWhat it is
GridGridThe whole kernel launch — covers the entire problem.
WorkgroupBlock / threadblockA group of threads on one Compute Unit (CU). Shares LDS (shared memory). Can synchronize via __syncthreads().
Wavefront (wave)Warp64 threads (AMD CDNA) executing the same instruction simultaneously (SIMT).
Thread (work-item)ThreadOne lane in a wave. Has its own thread ID and register state.
GRID (kernel launch — covers the whole problem)
Workgroup 0 (256 threads)
Wave 0 (T0..T63)
Wave 1 (T64..T127)
Wave 2 (T128..T191)
Wave 3 (T192..T255)
Workgroup 1 (256 threads)
Wave 0..3 (64 threads each)
Workgroup N−1
Wave 0..3
Key: 64 threads in a wave always execute the same instruction in lock-step. That is the essence of SIMT (Single Instruction Multiple Threads).

2. Lane vs thread

"Lane" and "thread" are two views of the same physical execution slot.

  • Lane = a hardware ALU slot inside a SIMD unit. There are exactly 64 lanes per SIMD on AMD CDNA.
  • Thread = the software view of one lane. Has its own thread ID and private registers.

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.

3. Hardware: GPU → CU → SIMD → lane

Below the software hierarchy is the physical hardware:

  • GPU contains many CUs (Compute Units). Example: MI300X has 304 CUs.
  • CU contains 4 SIMD units. The 4 SIMDs in a CU operate in parallel.
  • SIMD contains 64 lanes (ALUs) and a register file that can hold up to 8 resident waves.
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
SIMD ≠ instant execution. A SIMD holds up to 8 waves in its register file but executes only one wave per cycle. With multiple waves resident, when one wave waits for memory, the SIMD switches to another. This is latency hiding.

4. Registers, VGPRs & occupancy

Register types

TypeSizeScopeNotes
VGPR (vector GPR)32-bit (4 B)Private per laneUp to 256 per lane per wave. Each lane sees its own VGPR.
SGPR (scalar GPR)32-bit (4 B)Shared by 64 lanesUsed for scalar values like loop counters, addresses.
AGPR (accumulator GPR)32-bit (4 B)Private per laneCDNA-only. Used as MFMA accumulators.

How much register memory does one lane have?

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

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.
Trade-off: using more VGPRs per thread means each thread can hold more data, but fewer waves can be resident, reducing latency hiding.

5. Memory hierarchy: registers → LDS → cache → HBM

GPU memory has multiple levels, similar to CPU cache hierarchy:

LevelSize (per CU / total)Latency (cycles)Managed byCPU analogue
Registers (VGPR/SGPR)~256 KB / CU~1CompilerCPU registers
LDS (shared memory)64 KB / CU~10–30Software (explicit loads/stores)Scratchpad / fast SRAM
L1 cache16 KB / CU~30Hardware (transparent)L1 cache
L2 cache~16 MB total~150HardwareL2 cache
Infinity / L3~256 MB~300HardwareL3 cache
HBM (global memory)192 GB~500–1000HW + softwareDRAM
Key insight: registers are basically free (~1 cycle), HBM is very expensive (~500+ cycles). Performance comes from staging data through LDS and registers, and from hiding HBM latency with high occupancy.

6. Kernel launch <<< grid, block >>>

HIP/CUDA kernel launch syntax:

add_kernel<<< grid, block >>>(A, B, C, N);
ParameterMeaningExample
block (a.k.a. blockSize)Threads per workgroup256 → 4 waves per workgroup
gridNumber of workgroups4 → 4 workgroups total
(implicit)Wave count = blockSize / 64256 / 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

7. Vector loads & the 16-byte rule

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 typeSize (B)vec=1vec=2vec=4vec=8vec=16
fp16 / bf1622 B4 B8 B16 B (max)
fp3244 B8 B16 B (max)
int811 B2 B4 B8 B16 B (max)
Rule: vec × sizeof(dtype) ≤ 16 byte. Larger vec means fewer load instructions to move the same amount of data — faster.

Example GPU instructions

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

8. Walking through a simple kernel

__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.

If occupancy were 4, the SIMD would switch to other waves during the wait, keeping the ALU busy on every cycle. This is why high occupancy matters.

9. MFMA: cooperative matrix multiply

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
Important: MFMA is a wave-level instruction. It cannot be split across waves — one wave executes one MFMA. To compute a larger matrix multiply, multiple waves issue multiple MFMAs (covering different tiles).

CDNA vs RDNA

  • CDNA (data-center: gfx90a, gfx942, gfx950): MFMA available.
  • RDNA (consumer: gfx11, gfx12): no MFMA. Has WMMA (Wave Matrix Multiply-Accumulate) instead with similar idea.

10. Tiles & how a wave fills a tile (X0, Y0, X1, Y1)

What is a tile?

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 │
       └────────────────────────┘

How does a 64-thread wave fill a 64×64 tile?

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:

SymbolMeaning
X1 (= vec)Number of elements one thread loads in one instruction (X direction).
X0Number of threads placed along the X axis.
Y0Number of threads placed along the Y axis.
Y1Number 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

Why the X axis can use vec loads

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

Worked example: 64×64 tile, fp16, vec=4

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
              └─────────────────────────────────────┘

Choosing vec size

vecX0Y0Y1Loads / threadQuality
16416464worst (no vec)
23223232poor
41641616good
88888best (fp16 max)
Larger vec → fewer load instructions → faster, up to the 16-byte hardware limit.

11. Tile distribution patterns: thread / warp / block raked

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.

PatternWho covers one tileWave layout inside tile
thread_raked1 wave (64 threads)n/a (single wave)
warp_rakedMultiple waves cooperate1D stack (waves stripe along one axis)
block_rakedAll waves of the workgroup2D grid (waves arranged in a grid)

warp_raked — 1D wave layout (4 waves stacked along Y)

                X axis (XPerTile)
              ┌─────────────────────────────────┐
              │          Wave 0                 │   each wave covers full X width
              ├─────────────────────────────────┤
              │          Wave 1                 │   1/4 of Y
              ├─────────────────────────────────┤
              │          Wave 2                 │
              ├─────────────────────────────────┤
              │          Wave 3                 │
              └─────────────────────────────────┘

block_raked — 2D wave layout (4 waves in 2×2 grid)

                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 layout1 axis (Y)2 axes (X × Y)
Sub-tile per waveXPerTile × (YPerTile/N)(XPerTile/M) × (YPerTile/M)
X coverage by one wavefullpartial
The choice of pattern affects memory access patterns, MFMA fragment alignment, and register tile shapes. Each is best suited to different scenarios.

12. Cheat sheet

ConceptDefinition
GridThe 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.
LaneOne hardware ALU slot in a SIMD. 64 per SIMD.
SIMDHardware execution unit. 4 per CU. Can hold up to 8 resident waves.
CUCompute Unit. Contains 4 SIMDs and 64 KB LDS. ~304 per MI300X.
VGPRVector register, private per lane. 32-bit. Up to 256 per lane.
SGPRScalar register, shared by all 64 lanes in a wave.
AGPRAccumulator register (CDNA), used for MFMA output.
LDSLocal Data Share = software-managed shared memory. 64 KB per CU.
HBMGlobal memory. Large but slow (~500+ cycle latency).
OccupancyNumber of resident waves on a SIMD (1–8). Higher = better latency hiding.
Latency hidingSIMD switches to another resident wave while one waits on memory.
SIMTSingle Instruction Multiple Threads. All 64 lanes run the same instruction.
MFMAMatrix Fused Multiply-Add. Wave-cooperative matrix multiply (CDNA).
WMMARDNA equivalent of MFMA.
Tile2D chunk of a matrix processed by one workgroup.
XPerTile / YPerTileTile dimensions in elements (algorithm-defined).
vec / X1Elements per thread per load. Constrained by 16-byte limit.
X0, Y0Number of threads placed along X / Y axes.
Y1Y-axis iteration count per thread.
CoalescingAdjacent threads accessing adjacent memory → one wide HBM transaction.
Kernel launchkernel<<< grid, block >>>(...) — grid = workgroups, block = threads/wg.
tile_distribution_patternHow threads/waves are distributed across one tile (thread/warp/block raked).
Key takeaways
  • Threads run in waves of 64; one instruction = one wave-step.
  • Memory is the slow part. High occupancy hides memory latency.
  • Vector loads (up to 16 B) reduce instruction count dramatically.
  • Tile dimensions are chosen by the algorithm; thread layout is derived from them.
  • MFMA is wave-cooperative — cannot be split across waves.

5/03/2026

How to Access Korean-Only Websites from Overseas Using AWS EC2 (Seoul Region VPN)

How to Access Korean-Only Websites from Overseas Using AWS EC2 (Seoul Region VPN)

Some Korean websites (government, financial, public institutions) block access from foreign IP addresses. This guide shows you how to create a quick VPN tunnel through AWS EC2 in Seoul to get a Korean IP address.

What you need: AWS account, AWS CLI installed, Terminal (macOS/Linux)


Step 1: Verify AWS CLI

Make sure AWS CLI is installed and configured:

aws --version
aws sts get-caller-identity --region ap-northeast-2

If you see your Account ID, you're good to go.


Step 2: Create a Key Pair

aws ec2 create-key-pair \
  --key-name kr-proxy-key \
  --region ap-northeast-2 \
  --query 'KeyMaterial' \
  --output text > ~/Desktop/kr-proxy-key.pem

chmod 400 ~/Desktop/kr-proxy-key.pem

Step 3: Create a Security Group

# Create security group
aws ec2 create-security-group \
  --group-name kr-proxy-sg \
  --description "SSH proxy for Korean IP access" \
  --region ap-northeast-2

# Allow SSH inbound (replace sg-xxxxx with your Group ID from above)
aws ec2 authorize-security-group-ingress \
  --group-id sg-xxxxx \
  --protocol tcp \
  --port 22 \
  --cidr 0.0.0.0/0 \
  --region ap-northeast-2

Step 4: Find the Latest AMI

aws ec2 describe-images \
  --owners amazon \
  --filters "Name=name,Values=al2023-ami-2023*-x86_64" "Name=state,Values=available" \
  --query 'Images | sort_by(@, &CreationDate) | [-1].ImageId' \
  --output text \
  --region ap-northeast-2

Note the AMI ID (e.g. ami-09a64de684ce1ac0e).


Step 5: Launch EC2 Instance

aws ec2 run-instances \
  --image-id ami-09a64de684ce1ac0e \
  --instance-type t2.micro \
  --key-name kr-proxy-key \
  --security-group-ids sg-xxxxx \
  --associate-public-ip-address \
  --tag-specifications 'ResourceType=instance,Tags=[{Key=Name,Value=kr-proxy}]' \
  --region ap-northeast-2 \
  --query 'Instances[0].InstanceId' \
  --output text

Note the Instance ID (e.g. i-0df42e446381ecc9a).

t2.micro is free-tier eligible (750 hours/month for 12 months).


Step 6: Get Public IP

Wait about 30 seconds, then:

aws ec2 describe-instances \
  --instance-ids i-xxxxx \
  --region ap-northeast-2 \
  --query 'Reservations[0].Instances[0].[State.Name, PublicIpAddress]' \
  --output text

You should see something like: running 13.125.x.x


Step 7: Open SSH SOCKS5 Tunnel

ssh -D 1080 -N -f \
  -o StrictHostKeyChecking=no \
  -i ~/Desktop/kr-proxy-key.pem \
  ec2-user@13.125.x.x
  • -D 1080 — Creates a SOCKS5 proxy on local port 1080
  • -N — No remote command (tunnel only)
  • -f — Run in background

Step 8: Verify Korean IP

curl --socks5-hostname localhost:1080 https://ifconfig.me

If it returns your EC2's Korean IP (e.g. 13.125.x.x), it's working!


Step 9: Configure Browser Proxy

Option A: macOS System Settings

  1. Open System SettingsNetworkWi-Fi
  2. Click Details...Proxies
  3. Enable SOCKS Proxy
  4. Server: localhost / Port: 1080
  5. Click OK

Option B: Terminal (macOS)

sudo networksetup -setsocksfirewallproxy "Wi-Fi" localhost 1080
sudo networksetup -setsocksfirewallproxystate "Wi-Fi" on

Option C: curl only

curl --socks5-hostname localhost:1080 https://www.target-website.kr

Now open your browser and access the Korean website!


Clean Up (Important!)

When you're done, clean up everything to avoid charges:

1. Close the SSH Tunnel

pkill -f "ssh -D 1080"

2. Turn Off Browser Proxy

macOS GUI: System Settings → Network → Wi-Fi → Details → Proxies → SOCKS Proxy OFF

Terminal:

sudo networksetup -setsocksfirewallproxystate "Wi-Fi" off

# Verify
networksetup -getsocksfirewallproxy "Wi-Fi"
# Should show: Enabled: No

3. Terminate EC2 Instance

aws ec2 terminate-instances \
  --instance-ids i-xxxxx \
  --region ap-northeast-2

4. Delete Security Group

Wait about 30 seconds after termination, then:

aws ec2 delete-security-group \
  --group-id sg-xxxxx \
  --region ap-northeast-2

5. Delete Key Pair

aws ec2 delete-key-pair \
  --key-name kr-proxy-key \
  --region ap-northeast-2

rm ~/Desktop/kr-proxy-key.pem

Cost Summary

Item Cost
t2.micro (free tier) Free (750 hrs/month, 12 months)
t2.micro (after free tier) ~$0.0116/hr (~$8.5/month)
Data transfer Free up to 100GB/month

Tip: If you want to keep the instance for later, Stop it instead of terminating. You only pay for EBS storage (~$0.10/GB/month) while stopped.


Quick Reference: Reconnect Later

If you stopped (not terminated) the instance:

# Start the instance
aws ec2 start-instances --instance-ids i-xxxxx --region ap-northeast-2

# Wait ~30 seconds, get new public IP
aws ec2 describe-instances \
  --instance-ids i-xxxxx \
  --region ap-northeast-2 \
  --query 'Reservations[0].Instances[0].PublicIpAddress' \
  --output text

# Open tunnel
ssh -D 1080 -N -f -i ~/Desktop/kr-proxy-key.pem ec2-user@NEW-IP

# Set proxy
sudo networksetup -setsocksfirewallproxy "Wi-Fi" localhost 1080
sudo networksetup -setsocksfirewallproxystate "Wi-Fi" on

3/08/2026

How to Apply for a D-U-N-S Number

1. What is a D-U-N-S Number?

A D-U-N-S Number (Data Universal Numbering System) is a unique 9-digit business identifier issued by Dun & Bradstreet (D&B). It is used globally to identify and verify businesses for purposes such as:

  • Enrolling in the Apple Developer Program as an organization
  • Applying for government contracts and grants
  • Establishing business credit profiles
  • Vendor and supplier registration
💡 Apple requires a D-U-N-S Number to verify your organization's legal identity before issuing an Apple Developer team account. Individual accounts do not need one.

2. Prerequisites — Documents to Prepare

Before starting the application, gather the following information and documents:

ItemDescription
Business NameFull legal name exactly as registered
Registration NumberTax ID or company registration number
Business AddressFull address including postal code and country
Phone NumberInternational format: +[country code][number]
CEO / Owner InfoFull legal name and title (e.g., Owner, CEO)
WebsiteCompany website URL (if available)
Date FoundedOfficial business establishment date
Employee CountTotal headcount including yourself
Registration CertificateEnglish version — PDF format preferred
⚠️ Non-English documents must be translated. Check your country's official tax authority or business registration portal — many offer an official English-language certificate export option directly.

3. Go to the D&B Application Portal

Navigate to the appropriate portal based on your use case:

Apple Developer Program

👉 https://support.dnb.com/?CUST=APPLEDEV

Use this URL if you are applying specifically to enroll in the Apple Developer Program as an organization.

General D-U-N-S Application

👉 https://www.dnb.com/duns-number/get-a-duns.html

Use this for general business registration, government contracts, or other purposes.

4. Select Your Request Type

On the Apple Developer D&B portal, you will be asked two questions:

Step 4a

Under "What are you?" → select Developer Program

Step 4b

Under "What do you need?" → select Create New DUNS

5. Search for an Existing D-U-N-S

The portal checks whether your business already has a D-U-N-S number before creating a new one.

Step 5a

Enter your business name and full address in the search fields.

Step 5b

Click "Lookup by Name / Address".

Step 5c — Expected Result

For new businesses: you will see "No Match found" — this is completely normal and expected.

Step 5d

Click "click here to submit a request to create a new D-U-N-S" to open the full application form.

✅ If a match is found, your business already has a D-U-N-S. You can claim or update it instead of creating a new one — contact D&B support directly.

6. Fill Out the Application Form

Complete all fields in the Create D-U-N-S Number form. All fields must exactly match your official registration documents.

FieldWhat to Enter
Full Legal Business NameExact name as on registration documents
Business Registration NumberYour tax ID / company ID number
Company PhoneInternational format (e.g., +358501234567)
Street AddressFull street address
CityCity name
State / ProvinceState or province (if applicable)
Postal CodeZIP / postal code
CountryYour country
Business Structuree.g., Sole Proprietorship, LLC, Corporation
CEO / Owner NameFull legal name
CEO / Owner Titlee.g., Owner, CEO, Director
WebsiteCompany website URL
Home-Based BusinessYes / No
Number of EmployeesTotal headcount (including yourself)
Date FoundedOfficial establishment date (MM/DD/YYYY)
⚠️ Exact match required. The business name, address, and registration number must match your official documents precisely. Even minor differences (abbreviations, punctuation) can cause rejection.

Click "Next" to advance to the document attachment page.

7. Attach Supporting Documents

Upload a document to verify your business registration. This is required — submissions without a document attachment are typically rejected.

Accepted Documents

✅ English Business Registration Certificate (strongly preferred)

✅ Official government-issued business license

✅ Any official document showing your business name, registration number, and address

📄 PDF format is recommended. File size is typically not an issue, but keep it under 5MB to be safe. Make sure the document is clearly legible and all key information is visible.

8. Add Notes & Submit

There is usually a free-text Additional Details / Notes field. Always fill this in — it helps the D&B reviewer understand your submission and can speed up approval.

Example — First-time application

Sample note text
This is a new D-U-N-S number request for [Your Business Name], a sole proprietorship registered in [Country] since [Year]. Business is engaged in [industry, e.g. software development]. English Business Registration Certificate is attached.

Example — Resubmission after rejection

Sample note text
Resubmitting after previous rejection (Case #XXXXXXXX) which was closed due to missing English Business Registration Certificate. The English certificate is now attached. Business registration number: [Your Number]. Business started: [Date].

After reviewing all information, click Submit.

💾 Save your case number immediately after submission — you will need it to track your application status or to reference in any resubmission.

9. What Happens Next?

After submission, D&B will review your request. Here is the typical timeline:

ImmediatelyOn-screen confirmation + case number displayed. Check your email for a confirmation from D&B (check spam).
1–2 business daysD&B reviews your submitted information and documents.
3–5 business daysD-U-N-S Number issued and emailed to you.
Up to 30 daysIn some cases D&B may request additional documentation.
📬 D&B will send your D-U-N-S Number by email. Once received, the number is also searchable via the D&B portal. Apple recommends waiting 24–48 hours after receiving your number before using it for developer enrollment.

10. Troubleshooting: Common Rejection Reasons

IssueSolution
Missing English documentObtain an official English translation of your registration certificate and resubmit
Business name mismatchEnsure the name on the form exactly matches your official documents (character for character)
Address not recognizedUse the exact address format from your registration documents
Duplicate D-U-N-S foundContact D&B to claim or update the existing record instead of creating a new one
Case closed, no responseResubmit the form and reference the previous case number in the notes field
Document unreadableRe-export or rescan the document as a clear, high-resolution PDF

11. Using Your D-U-N-S for Apple Developer Enrollment

Once you have your D-U-N-S Number:

Step 2

Select "Enroll as an Organization" (not Individual).

Step 3

Enter your D-U-N-S Number when prompted. Apple will verify it against D&B's database.

Step 4

Allow 24–48 hours after receiving your D-U-N-S before enrolling — the number needs time to propagate in D&B's system.

✅ The Apple verification process typically takes a few days after enrollment. Apple may call your registered business phone number to verify your identity as the organization's legal representative.

Final Checklist

  • Prepared English business registration certificate (PDF)
  • Navigated to the D&B portal (APPLEDEV or general)
  • Selected correct request type (Developer Program → Create New DUNS)
  • Searched for existing D-U-N-S — confirmed no match
  • Filled out all form fields matching official documents exactly
  • Uploaded supporting document (English certificate)
  • Added descriptive notes in the Additional Details field
  • Submitted the form and saved the case number
  • Waiting for D&B email confirmation (3–5 business days)

Last updated: March 2026 · For general informational purposes

Main image by Syawish Rehman on Unsplash