libraries-openly-fused/fusedkernellibrary/fkl-architecture-overview

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 (ParArch::CPU_MULTITHREAD: the DPP splits its work into chunks on the shared thread pool and never sees the stream; see fkl-implementing-data-parallel-patterns). 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

1. Does correctness require more than one thread to cooperate? Yes → DPP;

no → Op.

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

1. 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 ParArch::GPU_NVIDIA, ParArch::CPU (with a public exec_row, or one row function per phase) and ParArch::CPU_MULTITHREAD (reusing it, under #if defined(FK_CPU_MULTITHREAD_ENABLED) && !defined(NVRTC_COMPILER)) specialisations; keep *Details and the CPU and CPU_MULTITHREAD specs outside any #if defined(__NVCC__) guard. The only exception are the attention DPPs (softmax, flash_attention, flash_attention_mma, flash_decode), which remain GPU-only by the repo owner's decision.

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

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