菜单
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'.
用 Codex 或 Claude 帮你安装 复制这段 Prompt，粘贴到 Codex、Claude 或其他助手里，让它检查 Skill 页面并帮你完成安装。
基于 SOC 职业分类
Build the FusedKernelLibrary tests and run them. Covers CMake configuration for Linux
Write user code with the Fused Kernel Library (FKL) — compose fused GPU/CPU pipelines with executeOperations, manage Ptr2D/Tensor data, streams, and runtime parameters. Use when writing an application or library that calls FKL, when porting OpenCV-style image pipelines to fused kernels, or when you need the canonical pipeline patterns (DNN preprocessing, multi-ROI crop, color conversion).
Wrap FKL from another language (Python, Rust, Julia...) using an AST/graph approach and JIT compilation. Covers why NVRTC fails, the Pointer Array ABI (void** for aligned parameters and I/O), generating C++ source strings for IOps and static DPP execution, zero-copy interop, and cross-platform shared library generation (.so/.dll). Use when building or extending a language binding for FKL.
Implement a new Operation for the Fused Kernel Library — Unary, Binary, Ternary, Read, Write, ReadBack or IncompleteReadBack structs, with Parent aliases, DECLARE_*_PARENT macros, build() overloads and unit tests. Use when adding an algorithm to FKL (new arithmetic, color conversion, geometric transform, memory pattern) or when a FusedOperation/alias fails to compile.
FKL data structures — Ptr2D, Tensor, TensorT, RawPtr, PtrDims, MemType, constructors and memory layouts (packed, planar CHW, transposed T3D). Use when allocating or wrapping GPU memory for FKL pipelines, when interfacing external pointers (torch/cupy buffers), or when a Tensor/TensorT constructor or pitch issue appears.
Choose and combine FKL's four fusion techniques — Vertical Fusion (VF), Backwards Vertical Fusion (BVF), Horizontal Fusion (HF) and Divergent Horizontal Fusion (DHF) — as described in the paper (arXiv:2508.07071). Use when deciding how to structure a pipeline for maximum fusion, when batching ROIs or images, or when different data planes need different processing in one kernel.
| 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"]}} |
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.
An Operation is strictly ALWAYS single-thread code. If more than one thread must participate in the code, that code belongs in a DPP.
exec(). Single-thread.__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.::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.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.
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.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.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 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 |
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).