Execution Methods (Device API)#

Each processing step – a composed operator expression such as IngestOperation – exposes two surfaces:

  • Host-side launch traits: block_dim, calculate_grid_dim, elements_per_thread, and related static members (see Traits).

  • Device-side execution methods: execute() overloads called from CUDA device code.

This page catalogs the device API: how to invoke execute() inside a kernel. Terminology (operator, compose, steps, fused kernel) is defined in Terminology. For operator configuration, see Operators.

Calling execute()#

Combine operators with + to form a complete expression for one processing step – for example IngestOperation, or a Function<...> step between ingest and exgest. The expression must include SM, Block(), and either Function<...> with its parameter operator or InputOutput with the matching format operator. Use BlockDim<...> to set the CUDA thread-block size (NPPDx uses 1D thread indexing: Y and Z must be 1). In device code, default-construct that type and call execute():

IngestOp().execute(input, tile_data, width, height);

The example below fuses three processing steps into a fused kernel:

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);
}

Common arguments#

Every execute() overload takes the full image width and height (in pixels), not just the nominal tile size. NPPDx maps CUDA blocks to tiles internally.

Launch configuration is computed on the host from a processing step such as Ingest:

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

Register-only API execute overloads#

Use the Register-only API for the simplest ingest \(\rightarrow\) pointwise work \(\rightarrow\) exgest into a fused kernel. Size intermediate buffers with Op::elements_per_thread (typically float values).

Step

execute signature

Function<...>

execute(processing_type* data, int width, int height)

Ingest

execute(const InputType* input, OutputType* output, int width, int height)

Ingest (planar)

execute(const ImageLayout<InputType, PlaneCount>& input, OutputType* output, int width, int height)

Exgest

execute(const InputType* input, OutputType* output, int width, int height)

Exgest (planar)

execute(InputType* input, const ImageLayout<OutputType, PlaneCount>& output, int width, int height)

With the Register-only API, processing_type is float. Function steps read and write the same register tile in place. Ingest writes the internal tile representation; exgest reads it and packs output.

Shared-memory execute overloads#

Use the shared_memory API for neighborhood steps (blur, median, resize, and similar) in fused kernels that keep the working tile in shared memory. Query storage types with input_storage_of_t, output_storage_of_t, and temp_storage_of_t (see Shared-memory storage traits). When one step’s layout covers the allocation, pass Op::shared_memory_size to the launch; when carving multiple tile buffers from one extern __shared__ allocation, use shared_memory::compute_total_tile_storage. See 00_introduction/introduction_example_shared_memory.cu and 03_area_operation/box_filter.cu.

Step

execute signature

Ingest

execute(const InputType* input, TileStorage<...>& output_channels, int width, int height)

Ingest (planar)

execute(const ImageLayout<...>& input, TileStorage<...>& output_channels, int width, int height)

Exgest

execute(const TileStorage<...>& input_channels, OutputType* output, int width, int height)

Exgest (planar)

execute(const TileStorage<...>& input_channels, const ImageLayout<...>& output, int width, int height)

Function<...> (in-place)

execute(TileStorage<...>& channels, int width, int height)

Function<...> (out-of-place)

execute(const TileStorage<...>& input, TileStorage<...>& output, int width, int height)

Function<...> (with intermediate)

execute(const TileStorage<...>& input, TileStorage<...>& intermediate, TileStorage<...>& output, int width, int height)

The three-tile overload supports steps such as resize that need a separate temporary tile. Pointwise and many neighborhood steps use the in-place or two-tile forms.

Device API rules#

  • Compile-time dispatch: each overload is enabled only for the matching processing step (static_assert rejects ingest signatures on function-only types, and vice versa).

  • Fused kernels: one __global__ function fuses a kernel by chaining execute() calls in the sequence order. A fused kernel may include multiple ingest or exgest steps. Register-only API passes the same register tile between calls; Shared-memory API passes TileStorage references.

  • Synchronization: execute() inserts __syncthreads() between shared-memory steps.

  • Halo and bounds: ingest loads the extended tile; function steps consume their local halo without reimplementing image boundaries. See Halos and step order and Achieving High Performance.

  • Matching processing steps: every + expression in a fused kernel must agree on TileSize, block layout, and SM architecture (see Image Processing Using NPPDx).

See also#