Image Processing Using NPPDx#

This introduction walks through a minimal NPPDx program: compose the ingest and exgest operators, fuse them in a CUDA kernel that converts a packed RGB24 image to a floating-point internal representation, and write the same format back out. The walkthrough is adapted from the 00_introduction/introduction_example.cu sample shipped with NPPDx.

NPPDx follows the same operator programming model as other MathDx device libraries. A compile-time description of the work to perform is built by adding operators together with +. The library validates the description at compile time and generates device code tailored to the specified formats, tile size, block layout, and target architecture.

Terminology#

  • Processing step (image-processing step): ingest, exgest, pointwise, area, or resize work in a fused kernel. Compose each step using MathDx Operators joined with +. Type aliases such as IngestOperation / ExgestOperation name those composed expressions.

  • Fuse: combine processing steps in one kernel, ideally using data on-chip and minimizing global-memory traffic.

  • Fused kernel: a fused sequence of processing steps, run in one __global__ function via execute() calls in order.

  • Local halo: neighborhood support required by one step (for example a 5x5 median needs 2 pixels). Available as a trait on each step.

  • Cumulative halo: minimum additional data required at ingest such that all subsequent processing steps can be run by the fused kernel without boundary condition considerations.

  • Memory halo: optimized halo to improve shared-memory caching and access speed. Cannot be smaller than the ingest Cumulative halo.

  • Tile: rectangular on-chip image region a CUDA block works on, set with TileSize<W, H>. That size is the nominal tile; with halo it becomes the extended tile (nominal plus cumulative halo at ingest). Held in registers or in TileStorage shared memory. As neighborhood steps run, the valid interior shrinks by each step’s local halo. A step may use a tile as:

    • input: often with halo

    • output: with or without halo; may be a different shape for resize

    • temporary: scratch storage (for example resize)

  • Packing format: arrangement of image samples in global memory for ingest and exgest. Selected with InputFormat / OutputFormat and values of packing_format (see Supported memory packing formats).

  • Function: MathDx operator Function<function::Tag> that selects the work a processing step performs (for example function::box_blur). Pair it with the matching parameter operator (see Function tag operator (processing step)).

Register-only API vs Shared-memory API#

The on-chip storage strategy depends upon the fused kernel’s steps and whether they use halos.

  • Register-only API: every step is pointwise and needs no neighborhood support. Ingest loads into registers; compute and exgest stay there. This is the faster processing path when applicable.

    Pointwise fused kernel from ingest format through three pointwise steps to exgest format

    Figure 2 Pointwise fused kernel: ingest format, three pointwise steps, and exgest format. With no neighborhood requirements, this sequence can use either API, but register-only is faster.#

  • Shared-memory API: if any of the steps need a non-zero halo (area filters, resize). Ingest loads the extended tile into shared memory; all steps in the fused kernel – including pointwise ones – run using shared-memory tiles.

    Shared-memory fused kernel with {3, 3, 3, 3} halo and mixture of area and pointwise steps

    Figure 3 Shared-memory fused kernel: ingest format with {3, 3, 3, 3} cumulative halo, 2 area steps, 1 pointwise step, and exgest format.#

execute() signatures for both APIs are listed under Execution Methods (Device API). The walkthrough below starts with the Register-only API (introduction_example.cu); shared-memory setup follows in Shared-memory tiles.

NPPDx Execution model: ingest, compute, exgest#

NPPDx provides processing steps that appear in user kernels as device function calls. Which steps to call and in which order is up to the kernel author. End-to-end image processing is typically structured as:

  1. Ingest: load the required data (plus any required halo) arranged in global memory according to its associated packing format and convert to floating point.

  2. Compute: apply one or more steps in local memory. Pointwise steps (for example color conversion and gamma) can run on either registers or a shared-memory tile. Area steps (blur, median, sharpen, and resize) require a shared-memory tile whenever the step’s halo is non-zero – when neighborhood support extends beyond the nominal tile itself. In a fused kernel, any step that needs a halo forces the whole algorithm to use the Shared-memory API – including pointwise steps in that fused kernel.

  3. Exgest: convert from the intermediate floating-point representation and write output to global memory according to its associated packing format.

In a typical fused kernel, format conversion happens twice: once on the way in, once on the way out, often with processing in the middle. More complex fused kernels may require data from multiple sources and exgest in several packing formats and sizes. This is very efficient since the data is already loaded.

Specifying ingest and exgest#

The first step is to specify the ingest and exgest steps. To that end, NPPDx provides the InputOutput, InputFormat / OutputFormat, TileSize, Block, and SM operators, which can be combined with +:

#include <nppdx.hpp>

template<int SM>
using IngestOperation = decltype(
    nppdx::InputOutput<nppdx::input_output_direction::ingest>() +
    nppdx::InputFormat<nppdx::packing_format::rgb24>() +
    nppdx::TileSize<48, 48>() +
    nppdx::Block() +
    nppdx::SM<SM>());

template<int SM>
using ExgestOperation = decltype(
    nppdx::InputOutput<nppdx::input_output_direction::exgest>() +
    nppdx::OutputFormat<nppdx::packing_format::rgb24>() +
    nppdx::TileSize<48, 48>() +
    nppdx::Block() +
    nppdx::SM<SM>());

IngestOperation specifies the direction (ingest), input format through InputFormat (rgb24), tile size, execution operator, and architecture. ExgestOperation mirrors it with exgest and OutputFormat. IngestOperation and ExgestOperation in the same fused kernel must agree on tile size, execution operator, and architecture.

Each execute() call is backed by a composed expression such as IngestOperation or a Function step in the fused kernel.

A minimal kernel#

The Register-only API keeps the intermediate tile in registers. Shared memory is faster than Global memory, and registers are faster still. This is the fastest and simplest way to connect ingest and exgest inside a user kernel:

template<typename IngestOp, typename ExgestOp>
__global__ void image_processing_kernel(const uint8_t* global_input,
                                        uint8_t* global_output,
                                        int width, int height) {
    using processing_t = nppdx::processing_type_of_t<IngestOp>;
    constexpr size_t elements_per_thread = nppdx::elements_per_thread_of_v<IngestOp>;

    processing_t intermediate_data[elements_per_thread];

    IngestOp().execute(global_input, intermediate_data, width, height);

    // Additional pointwise work can be applied here on intermediate_data.

    ExgestOp().execute(intermediate_data, global_output, width, height);
}

processing_type_of_t is the floating-point type used during processing. elements_per_thread_of_v is the number of intermediate values each thread holds – fixed at compile time so register usage is known before launch.

Querying compile-time traits#

Before launching the kernel, query traits from IngestOperation. Examples below show several of these values:

constexpr auto tile_size             = IngestOperation<800>::suggested_tile_size;
constexpr auto elements_per_thread   = IngestOperation<800>::elements_per_thread;
constexpr auto block_dim             = IngestOperation<800>::block_dim;
constexpr auto format                = nppdx::packing_format_of_v<IngestOperation<800>>;

Op::shared_memory_size is always the dynamic shared-memory requirement for the Shared-memory API for that step (from inputs, outputs, and temp). Composition does not record whether the kernel will use registers or shared memory, so this value is not “zero for Register-only.” Register-only launches pass 0 (or omit the third <<<...>>> argument). Shared-memory examples pass Op::shared_memory_size or shared_memory::compute_total_tile_storage<...>().

Launching the kernel#

IngestOperation exposes launch traits as static members. Prefer calculate_grid_dim over manual grid arithmetic. The introduction example launches as:

const dim3 grid_dim = IngestOperation::calculate_grid_dim(width, height);

image_processing_kernel<IngestOperation, ExgestOperation>
    <<<grid_dim, ingest_block_dim>>>(d_input, d_output, width, height);

Here ingest_block_dim is IngestOperation::block_dim from the trait query above. The Register-only launch omits dynamic shared memory (equivalent to 0). Do not pass IngestOperation::shared_memory_size unless the kernel actually uses the Shared-memory API.

The introduction example uses common::run_example_with_sm to instantiate the correct SM value for the current device.

Fusing three steps: ingest, color_convert, exgest#

The pattern below fuses a three-step image-processing series (ingest, color_convert, exgest) in one kernel, with yuv420p output:

template<typename Ingest, typename Convert, typename Exgest>
__global__ void convert_kernel(const uint8_t* input, uint8_t* output,
                               int width, int height) {
    float tile_data[Ingest::elements_per_thread];

    Ingest().execute(input, tile_data, width, height);
    Convert().execute(tile_data, width, height);
    Exgest().execute(tile_data, output, width, height);
}

template<int SM>
void launch_rgb_to_yuv(const uint8_t* input, uint8_t* output,
                       int width, int height) {
    using Ingest = decltype(nppdx::InputOutput<nppdx::input_output_direction::ingest>() +
                            nppdx::InputFormat<nppdx::packing_format::rgb24>() +
                            nppdx::TileSize<48, 48>() + nppdx::SM<SM>() + nppdx::Block());

    using Convert = decltype(nppdx::Function<nppdx::function::color_convert>() +
                             nppdx::ColorConvert<nppdx::color_space::rgb,
                                                 nppdx::color_space::yuv_bt601,
                                                 nppdx::bit_depth::bpp_8u,
                                                 nppdx::bit_depth::bpp_8u>() +
                             nppdx::TileSize<48, 48>() + nppdx::SM<SM>() + nppdx::Block());

    using Exgest = decltype(nppdx::InputOutput<nppdx::input_output_direction::exgest>() +
                            nppdx::OutputFormat<nppdx::packing_format::yuv420p>() +
                            nppdx::TileSize<48, 48>() + nppdx::SM<SM>() + nppdx::Block());

    dim3 grid = Ingest::calculate_grid_dim(width, height);
    convert_kernel<Ingest, Convert, Exgest>
        <<<grid, Ingest::block_dim>>>(input, output, width, height);
}

Shared-memory tiles#

For neighborhood steps (box blur, resize, etc.), use a shared-memory tile instead of register arrays. IngestOperation (and each Function<...> step in a fused kernel) exposes the TileStorage types and shapes it needs through storage traits (see Shared-memory storage traits).

The introduction_example_shared_memory.cu example uses the same ingest and exgest composition as the Register-only API, but passes a typed shared-memory tile (TileStorage) to execute() instead of a per-thread array. On the host, query the input tile type from the ingest step and use that step’s shared-memory size for the kernel launch:

#include <nppdx/shared_memory.hpp>

using TileStorageType = nppdx::input_storage_of_t<IngestOperation>;
constexpr size_t smem_size = IngestOperation::shared_memory_size;
// equivalent: shared_memory::compute_total_tile_storage<TileStorageType>();

const dim3 grid_dim  = IngestOperation::calculate_grid_dim(width, height);
const dim3 block_dim = IngestOperation::block_dim;

image_processing_kernel_with_channels<IngestOperation, ExgestOperation>
    <<<grid_dim, block_dim, smem_size>>>(d_input, d_output, width, height);

In the kernel, bind the dynamic shared-memory allocation to that storage type and pass the resulting TileStorage object to execute():

template<typename IngestOp, typename ExgestOp>
__global__ void image_processing_kernel_with_channels(const uint8_t* global_input,
                                                      uint8_t* global_output,
                                                      int width, int height) {
    extern __shared__ unsigned char smem[];

    using TileStorageType = nppdx::input_storage_of_t<IngestOp>;
    auto channels = TileStorageType::from_memory(smem);

    IngestOp().execute(global_input, channels, width, height);
    ExgestOp().execute(channels, global_output, width, height);
}

Shared-memory ingest loads the extended tile required by the fused sequence of steps. Boundary padding is applied during ingest so downstream steps do not implement or even be aware of boundary conditions.

Fused kernels (with halos)#

For a fused kernel such as ingest \(\rightarrow\) box blur \(\rightarrow\) exgest, compose the operators for each step (with TileSize, SM, Block, and halo operators as needed), then fuse the steps into one kernel. Query storage types from traits and call execute() in process order – no separate wrapper is required.

template<typename Ingest, typename BoxBlur, typename Exgest,
         typename InputTileStorage, typename BlurTileStorage>
__global__ void box_blur_kernel(const uint8_t* input, uint8_t* output,
                                unsigned int width, unsigned int height) {
    extern __shared__ unsigned char smem[];

    auto tiles = nppdx::shared_memory::slice_into_tile_storage<InputTileStorage, BlurTileStorage>(smem);
    auto input_channels  = std::get<0>(tiles);
    auto output_channels = std::get<1>(tiles);

    Ingest().execute(input, input_channels, width, height);
    BoxBlur().execute(input_channels, output_channels, width, height);
    Exgest().execute(output_channels, output, width, height);
}

Host setup:

using InputTileStorage = nppdx::input_storage_of_t<Ingest>;
using BlurTileStorage  = nppdx::output_storage_of_t<BoxBlur>;

constexpr size_t smem_size =
    nppdx::shared_memory::compute_total_tile_storage<InputTileStorage, BlurTileStorage>();

box_blur_kernel<Ingest, BoxBlur, Exgest, InputTileStorage, BlurTileStorage>
    <<<grid, block, smem_size>>>(d_input, d_output, width, height);

Steps that need three tiles (for example resize) require temporary storage – use temp_storage_of_t<Resize> alongside input_storage_of_t and output_storage_of_t. See 03_area_operation/box_filter.cu and 04_resize/fused_resize.cu.

Note

Kernel Processing order and traits. execute() call order depends on the traits composed on each fused kernel step – halos, tile bounds, shared-memory layout, and validity checks. Clipping to the output format range happens at exgest; intermediate steps use unclipped floating-point tiles. Results for different orderings may be close after final clipping but are not interchangable without testing.

See the numbered examples under example/nppdx/ and the Examples section.

Miscellaneous considerations#

Ingest and exgest consist of interchangable templated formats, see Supported memory packing formats. Just this I/O layer understands how data is packed for RGB, V210, NV12, planar YUV, or another supported layout. Area and pointwise processing steps always see floating-point tiles, and have no knowledge of the underlying format. This also means that new packing formats can easily be added to the library.

Two ingest APIs are provided – see Register-only API vs Shared-memory API above. Briefly: Register-only when every step is pointwise; Shared-memory when any step needs a halo (and then the whole fused kernel uses shared memory).

Exgest clips to the output range by default. It can be omitted in cases where there is only a bit-depth change.

Extensibility and tuning knobs#

NPPDx separates what is configured from what the library derives:

Knob

Role

TileSize<W, H>

User-facing. Primary tuning parameter for occupancy, shared-memory footprint, and arithmetic intensity.

BlockDim<...>

User-facing. CUDA thread-block size for the step (1D: Y and Z must be 1).

SM<arch>

User-facing. Target architecture for code generation.

Cell size

Internal. How ingest, function, and exgest backends map work to threads within a tile; chosen for correctness and performance rather than exposed as a primary API knob.

Local halo

Derived. Neighborhood need for one step; available as a trait on that step (see Terminology).

Cumulative halo

User-facing for Shared-memory fused kernels: minimum additional data at ingest so every processing step can run without boundary-condition handling (see Terminology and Halos and step order).

Internal resource management#

Within a fused kernel, NPPDx tracks the resources the fused kernel needs:

  • Halos and tiles: ingest loads the extended tile; each step’s local halo shrinks the valid interior.

  • Temporary tile: steps that need temporary storage (for example resize) require a temp TileStorage tile.

  • Thread cells: algorithms execute on cells (groups of pixels) so neighboring work can share computation and register storage.

  • Tile bounds for area process: neighborhood steps such as box blur, resize, and related filters use compile-time tile-bound computation so each step knows which shared-memory region is valid without per-step boundary handling at image edges.

Next steps#

For further details on customizing NPPDx, consult: