🚀 Fused Kernel Library (FKL)

Redefining Data Parallel portability, performance, and programmability using standard C++20.

---

💡 What is FKL?

Built for speed. The number one purpose of the Fused Kernel Library is to deliver uncompromising, programmer-friendly fast code.

The Fused Kernel Library (FKL) is a C++20 header-only framework designed to perform automatic GPU kernel fusion. Instead of relying on custom external compilers or domain-specific languages, FKL leverages modern C++ capabilities to fuse sequential and parallel operations into a single, highly optimized kernel at compile time.

Currently supporting CPU (single-thread and multithreaded), CUDA, and ROCm backends (with an architecture designed to easily adopt other GPU languages), FKL transforms memory-bound operations into compute-bound powerhouses by keeping data in registers and eliminating redundant global VRAM reads and writes.

✨ Key Features

• 🔥 Generic Vertical Fusion: Combine operations sequentially without rewriting thread handling. The compiler treats consecutive operations as if written inline, unlocking massive performance optimizations.
• ⏪ Backwards Vertical Fusion: Read and compute only what you need. Apply complex pipelines (e.g., YUV decode $\rightarrow$ crop $\rightarrow$ resize $\rightarrow$ normalize) in a single pass. Data stays in GPU registers until the final write.
• ⚡ Horizontal Fusion: Process multiple data planes in parallel within the same GPU kernel using blockIdx.z, maximizing memory bandwidth for small data payloads.
• 🔀 Divergent Horizontal Fusion: A novel approach that executes completely different kernels in parallel across the same grid. This allows different Streaming Multiprocessor (SM) components to be saturated simultaneously.
• 🤝 Closed-Source Friendly: Integrate FKL into proprietary codebases easily. Wrap your custom CUDA or HIP kernels in FKL's InstantiableOperation interface to fuse them without exposing your internal source code.
• 🧼 Pure Imperative Ergonomics: Despite relying on compile-time metaprogramming under the hood, consuming FKL requires zero template gymnastics. Pipelines are declared as clean, flat, sequential arguments (Operation::build(params)).
• 🔬 Non-Invasive Inline Probing: Debug intermediate register data effortlessly. Drop a MidWrite at any point in the pipeline to capture in-flight values to memory for CPU inspection without breaking fusion.
• 🖥️ Trivial CPU Debugging: Because operations are hardware-agnostic, you can seamlessly switch to the CPU backend to step through your pipeline logic using standard C++ debuggers (like GDB or Visual Studio) before deploying to the GPU. The multithreaded CPU backend (ParArch::CPU_MULTITHREAD) then runs the same pipelines on all the CPU cores, with the same results as the single-thread one.

📊 Benchmarks

One executeOperations() call. One kernel. This is what happens to the clock:

50× faster than OpenCV (CUDA)	34× faster than NVIDIA NPP	5.9× faster than a hand-written NPPDx kernel	164× in production, 50 crops across four 4K frames
3.1 µs 4K NV12 → crop → resize → normalize, end to end	9 KB read from DRAM instead of a 12.4 MB frame	11× faster than OpenCV on 32 CPU threads, same source code	10 lines of pipeline code, versus 86 for the NPPDx kernel

All of it measured, on real pipelines, with validated outputs. Across everything we have benchmarked so far, FKL is 2× to more than 1,000× faster than the equivalent library calls, and the bigger the GPU or the longer the pipeline, the larger the gain.

How these numbers were measured

• 4K NV12 pipeline (first group of bars, and the tiles): a 3840×2160 NV12 frame → RGB → N crops of 20×10 px → bilinear resize with border to 64×64 → multiply, subtract, divide → float3 tensor, for N = 1 to 100. RTX 4090, CUDA 13.3, OpenCV 4.14 with CUDA modules, NPP 13.1, NPPDx 0.1.2. End-to-end latency including each library's host-side work, median of 1,000 interleaved iterations; speedups are per-iteration means at the best N for FKL (the worst N still gives 26× over OpenCV, 26× over NPP and 4.4× over NPPDx). Counting GPU time alone, the FKL kernel is up to 102×, 73× and 15.5× faster, respectively; when intermediates can't stay in L2 between kernels, as on a busy GPU, up to 122× and 89× over OpenCV and NPP. The CPU figure compares FKL's multithreaded CPU backend with OpenCV's pthreads backend on a 16-core Ryzen 9 5950X. From the Automatic Backwards Vertical Fusion paper (2026, in preparation); code in bvf-paper-bench, with the comparison with NPPDx in its fkl-vs-nppdx folder.
• Production figures are from the PUMPS+AI 2025 poster: the AutomaticTV pipelines before and after moving to FKL (NV12 → warp/crop → resize → normalize → tensor), on the GPUs named in the chart.
• Vertical & horizontal fusion figures are from The Fused Kernel Library paper, measured on an RTX 4090 through the cvGPUSpeedup (OpenCV-like) and FastNPP (NPP-like) wrappers: chains of fused arithmetic on a 4096×2160 image versus one OpenCV-CUDA kernel per operation; a batched crop → resize → color convert → multiply → subtract → divide → split pipeline versus NPP, for batches of 10 to 150; and a synthetic stress test with up to 10,000 fused operations on a batch of 50, which also runs 2,527× faster than OpenCV-CUDA captured in a CUDA Graph.

---

🧬 Vector Types and HIP/CUDA Interoperability

FKL owns its vector types and operators in the fk namespace on every backend.

⚠️ Crucial Namespace Guidelines:

• Use fk::float3, fk::uchar4, etc. outside that namespace, even with using namespace fk;. HIP and CUDA declare different types with the same names in the global namespace.
• Do NOT put using namespace fk; at global scope in CUDA translation units or headers. nvcc appends host-registration code that uses unqualified CUDA types (like uint3), making names ambiguous. Keep using-directives inside functions or qualify FKL names explicitly.
• Native HIP/CUDA vector types are NOT FKL vector types. Convert values explicitly by component at APIs requiring a native vector (e.g., ::float3 native{value.x, value.y, value.z}). Do not reinterpret vector pointers, as layouts differ (e.g., fk::double3 has size/alignment 32/16, while AMD Clang's ::double3 is 24/8).

Internal FKL expressions do not need backend-specific casts. FKL implements almost all standard C++ operators for its vector types, strictly following scalar C++ type promotion conventions per component. Furthermore, it seamlessly supports mixing scalar and vector types in binary operators where it logically makes sense (e.g., multiplying a vector by a scalar).

💻 Show Me The Code

Traditional GPU programming requires launching multiple sequential kernels to process an image. With FKL, you define the operations and fuse them lazily.

Here is an example that extracts 5 crops from an image, resizes them, applies arithmetic changes, and writes them contiguously to a Tensor—all in one kernel launch.

#include <fused_kernel/core/execution_model/memory_operations.h>
#include <fused_kernel/algorithms/basic_ops/arithmetic.h>
#include <fused_kernel/algorithms/image_processing/crop.h>
#include <fused_kernel/algorithms/image_processing/color_conversion.h>
#include <fused_kernel/algorithms/image_processing/resize.h>
#include <fused_kernel/fused_kernel.h>

void preprocess() {
    using namespace fk;

    // 1. Create the fkl CUDA stream
    Stream stream;
    
    // 2. Get the input image
    const Ptr2D<fk::uchar3> inputImage = getGPUSourceImage(stream);
    
    // 3. Define the crops on the source image
    constexpr std::array<Rect, BATCH> crops{
        Rect(300, 125, 60, 40),
        Rect(400, 125, 60, 40),
        Rect(530, 270, 130, 140),
        Rect(560, 115, 100, 35),
        Rect(572, 196, 40, 15)
    };
    
    // 4. We want a Tensor of contiguous memory for all crops as output
    Tensor<fk::uchar3> output(outputSize.width, outputSize.height, BATCH);

    // 5. CREATING AND EXECUTING YOUR FUSED CUDA/HIP KERNEL
    // No execution happens until `executeOperations` is called.
    // Types determine the compile-time kernel code; parameters are passed at runtime.
    executeOperations<TransformDPP<>>(
        stream,
        PerThreadRead<ND::_2D, fk::uchar3>::build(inputImage.ptr()),
        Crop<>::build(crops),
        Resize<InterpolationType::INTER_LINEAR, AspectRatio::PRESERVE_AR>::build(outputSize, backgroundColor),
        Mul<fk::float3>::build(make_<fk::float3>(2.f, 2.f, 2.f)),
        Sub<fk::float3>::build(make_set<fk::float3>(128.f)),
        SaturateCast<fk::float3, fk::uchar3>::build(),
        TensorWrite<fk::uchar3>::build(output.ptr())
    );

    stream.sync();
}

🔍 Under the Hood: How it Works

1. PerThreadRead & Crop: Defines a 3D threadblock configuration based on the BATCH size. FKL automatically calculates that only the useful threads for each crop size will perform reads.
2. Resize: Re-evaluates the thread grid to 60x60x5. Each thread asks the Crop operation for the specific pixels it needs to interpolate the output.
3. Continuation Operations (Mul, Sub, ColorConversion): Applied element-wise entirely in GPU registers.
4. TensorWrite: Finally writes the optimized, contiguous data to global memory.

Result: A highly efficient, variadic template kernel that compiles down to a single optimized footprint.

On all CPU cores: the same pipeline runs multithreaded with executeOperations<TransformDPP<ParArch::CPU_MULTITHREAD>> on a Stream_<ParArch::CPU_MULTITHREAD>, which behaves like a CUDA stream: calls are asynchronous, and stream.sync() waits for the work and rethrows its errors. Its data must be in host memory (MemType::Host) and stay valid until sync(). This backend uses the thread-pool git submodule: clone with git clone --recursive (or run git submodule update --init), otherwise CMake silently disables it (option ENABLE_CPU_MULTITHREAD). The attention DPPs are the exception: they remain GPU-only.

🧘 Developer Experience: Common Myths vs. Reality

Myth	Reality in FKL
"It's a template metaprogramming library, so the user API must be complex."	False. All complexity is hidden behind static .build() factories. Chaining operations is as simple as: executeOperations(stream, Op1::build(...), Op2::build(...));. You write standard imperative-looking C++.
"Fused kernels are a black box and impossible to debug."	False. You can tap intermediate register values at any point by dropping a MidWrite or custom probe operation directly into the operation list without breaking the pipeline.
"GPU kernels are impossible to step through line-by-line."	False. Since operations work across all backends, you can simply compile with the CPU backend to step through your pipeline using standard debuggers (GDB, MSVC) to verify logic before deploying to the GPU.
"Template errors will be unreadable."	True and False. While C++ template errors are notoriously verbose, modern AI tools can instantly translate them and explain exactly where the issue lies. Furthermore, because FKL constructs arguments via modular .build() calls, type mismatches are often caught early and localized to your specific operation rather than buried deep inside internal instantiation trees.

---

🚀 Try It Live (Zero Setup)

You can explore how FKL generates highly optimized assembly without installing anything:

• 🛠️ Compiler Explorer (Godbolt): Try FKL v0.1.13-LTS
• 📓 Google Colab: Run the FKL Playground

---

📦 Using FKL in Your CMake Project

FKL is header-only. Clone it with git clone --recursive (the thread-pool submodule enables the multithreaded CPU backend), then either add it to your project:

set(BUILD_TEST OFF)   # FKL's tests are not needed
set(BUILD_UTEST OFF)
add_subdirectory(FusedKernelLibrary)
target_link_libraries(my_app PRIVATE FKL::FKL)

or install it and find it:

cmake -B build -S FusedKernelLibrary -DBUILD_TEST=OFF -DBUILD_UTEST=OFF
cmake --install build --prefix /path/to/install

find_package(FKL 0.2 REQUIRED)   # configure with -DCMAKE_PREFIX_PATH=/path/to/install
target_link_libraries(my_app PRIVATE FKL::FKL)

FKL::FKL provides the include directories and requires C++20: your project has to enable CXX (e.g. project(my_app LANGUAGES CXX CUDA)), and CMake then compiles your C++, CUDA and HIP sources as C++20, unless you set CUDA_STANDARD/HIP_STANDARD, which must then be 20 or higher. When FKL was configured with the multithreaded CPU backend (ENABLE_CPU_MULTITHREAD, ON when the submodule is checked out), it also adds FK_CPU_MULTITHREAD_ENABLED and the threads library. Before 1.0, find_package(FKL 0.2) only accepts 0.2.x versions, because a minor version can break the API.

---

📚 Research & Publications

FKL is built on rigorously tested academic methodology. If you use FKL in your research, please refer to our publications:

• 📄 Preprint Journal Paper (arXiv): Methodology for GPU Kernel Fusion (Pending IEEE approval)
• 📄 Automatic Backwards Vertical Fusion (2026, in preparation): the fusion model behind the ReadBack operations (crop, resize, warp, border, deinterlace). Benchmarks: bvf-paper-bench, including the comparison with NPPDx (fkl-vs-nppdx).
• 🖼️ NVIDIA GTC 2025: .
• 🏆 PUMPS + AI 2025: Award-winning continuation poster at the Barcelona Supercomputing Center ![PUMPS 2025 Poster].
• PUMPS + AI 2026: Explanation of how we avoid performance damaging local memory accesses in a systematic way. ![PUMPS 2026 Poster].

---

🛠️ Testing & Support

FKL is an Apache 2.0 OpenSource project maintained by contributors in their spare time (and supported by Grup Mediapro S.L.U. during work hours).

While we provide no formal guarantees or free support, we actively test against:

CUDA Builds

• Ubuntu 24.04: g++ 13 + CUDA 13.4
• Windows 11: Visual Studio 2026 (14.44 toolset) + CUDA 13.0
• Windows: Visual Studio 2026 (14.51 toolset) + CUDA 13.4

ROCm Builds (x86_64 only)

• Ubuntu 26.04: ROCm 10.1.0

CPU Backend (single-thread and multithreaded)

• Linux: g++ 13 or clang++-23
• Windows: Visual Studio 2026 (14.51 toolset)

Note on Versioning: The main branch is where active API development happens (minimum C++20, subject to breaking changes). For production stability, please use the LTS-C++17 branch (Note: ROCm is not supported on the LTS branch).