| name | fkl-architecture-overview |
| description | Use this skill to classify any FKL work as an Operation (Op), Data Parallel Pattern (DPP), or constexpr_lib. Trigger this before writing a new algorithm, when adapting an existing kernel to FKL standards, or when addressing PR review comments like 'should be a DPP' or 'not FKL'. |
| version | 0.1.0 |
| metadata | {"hermes":{"tags":["GPU","CUDA","FusedKernelLibrary","Architecture"]}} |
FKL Architecture Overview
A fast classification guide for the Fused Kernel Library: given an algorithm,
decide whether it is an Operation or a Data Parallel Pattern (DPP), or if
it should be split into a DPP and several Operations and route its data correctly.
The algorithm, or parts of it, could also be a function in constexpr_lib.
It does NOT replace fkl-implementing-operations or
fkl-implementing-data-parallel-patterns (those are
authoritative and go deeper) — it is the 60-second "what am I building?" filter
to apply before reaching for them, plus the most common review-fix mappings.
When to Use
- Before writing a new algorithm, to decide Op vs DPP vs cxp::.
- When a review says "this should be a DPP, not an Operation or standard Kernel",
"this is not FKL / I can't fuse this", or "use standard Ops".
- When adapting an existing kernel to FKL standards, and having to refactor how
a kernel reads, computes, and writes, using IOps.
The single classifying rule
An Operation is strictly ALWAYS single-thread code. If more than one thread
must participate in the code, that code belongs in a DPP.
- Op — code that defines how a single thread has to modify the data. One thread's work; stateless struct,
static
exec(). Single-thread.
- DPP — code that defined thread behavior. Anything needing >1 thread to cooperate: shared memory,
__syncthreads, __shfl, warp-collective mma.sync, cp.async tile staging,
CPU threads in case of CPU multithreaded implementations of the DPP.
A DPP routes data through IOps and never modifies the data directly, it always
does it using IOps, except when using tensor cores and specialized hardware that
determines both thread behavior and modifications over data.
- IOp (Instantiable Operation) — an Op plus its params, produced by
::build(), invoked as IOp::Operation::exec(...). Exceptionally, and only
to avoid code duplication, can use Operation::exec(...) without having an IOp and
only inside the implementation of another Op.
Operation types
The full IOp-form table (with exact exec() signatures) lives in
fkl-implementing-operations. The types:
ReadType, WriteType, UnaryType, BinaryType, ReadBackType, IncompleteReadBackType, TernaryType, IncompleteTernaryType, MidWriteType, OpenType, ClosedType.
Procedure
- Does correctness require more than one thread to cooperate? Yes → DPP;
no → Op.
- A DPP takes its read as a Read IOp (prologue), its write as a Write
IOp (epilogue), and intermediate compute as IOps. Ops will always be
passed as IOp instances to the DPPs. Never as template parameters.
- DPP shape:
exec(details, readIOp, computeIOp, writeIOp) — IOps are separate
exec parameters, not packed into a struct; a *Details struct carries only
external scalars (dims, identity, mask). Provide both ParArch::GPU_NVIDIA
and ParArch::CPU specialisations; keep *Details and the CPU spec outside
any #if defined(__NVCC__) guard.
- For an output written through a fused epilogue, fuse the epilogue chain with
the destination write (
epilogue.then(D)) and invoke the whole fused IOp:
Output::Operation::exec(thread, value, output). The first link of the
.then() chain must be a compute op (Cast<T,T> for identity), not a Read.
Pitfalls
mma.sync is warp-collective → it must be used inside a DPP, never an Op.__nv_bfloat16 has no VectorTraits, so PerThreadRead/PerThreadWrite
will not instantiate for it — supply a small Read/Write Op over a RawPtr in
tests that use bf16.
Review-fix mapping
| Review comment | Fix |
|---|
| "not FKL" / "breaks the FKL philosophy" | make it a DPP composing IOps + standard Ops |
| "should be a DPP, not an Operation/kernel" | wrap the kernel as a DPP taking Read/Write IOps |
| "this is cudaGraphs, no FKL Ops/DPPs" | rebuild on the FKL execution model; benchmark the DPP path |
| "faking the IOps, not using operator|" | compose with the real | / .then() fusion |
| "epilogue done by a single thread" | parallelise the epilogue with shared memory / a warp-cooperative DPP |
| "Max/Min/Sum should be standard Ops as args" | pass Add/Max/Min as IOps in exec function parameters |
| "duplicates code" | reuse an existing Op or a prior primitive |
Verification
For any proposed change you should be able to state, in one line: (a) Op vs DPP
and why (the single-thread rule), and (b) which IOps carry its read,
compute, and write. Cross-check the type names against the live table in
fkl-implementing-operations/SKILL.md.
Background: "The Fused Kernel Library: A C++ API to Develop Highly-Efficient GPU
Libraries" (arXiv:2508.07071).
