GPU per-feature kernel scaffolding templates¶

Use the per-backend kernel templates when you write a new GPU feature kernel: they wrap the stream, event and readback lifecycle that every fork-added kernel otherwise re-implements by hand. The templates exist for CUDA, HIP and Metal (ADR-0246; HIP and Metal sections added per ADR-0484). SYCL has no template. The Vulkan backend was removed in ADR-0726.

Which kernels use the template¶

Counts are translation units under core/src/feature/<backend>/ that include kernel_template.h.

The bespoke CUDA and HIP kernels (ADM, VIF, motion) run multi-stream or multi-scale dispatch the template does not model. They stay migration candidates. On Apple Silicon unified memory the device / host_pinned split of CUDA collapses to a single MTLBuffer (MTLResourceStorageModeShared) whose [buffer contents] pointer is cached in VmafMetalKernelBuffer::host_view.

Each kernel migration is a separate PR with its own places=4 cross-backend gate (per ADR-0214). The helpers wrap the steps that historically caused regressions, such as forgetting cuStreamSynchronize before cuStreamDestroy.

CUDA template¶

The CUDA template formalises the async-stream + event lifecycle every fork-added CUDA kernel currently uses. The reference implementation is integer_psnr_cuda.c.

Surface¶

#include "cuda/kernel_template.h"

typedef struct VmafCudaKernelLifecycle {
    CUstream str;       /* private non-blocking stream for readback */
    CUevent  submit;    /* recorded post-launch on picture stream    */
    CUevent  finished;  /* recorded post-readback on str             */
} VmafCudaKernelLifecycle;

typedef struct VmafCudaKernelReadback {
    VmafCudaBuffer *device;       /* device-side accumulator     */
    void           *host_pinned;  /* pinned host readback slot   */
    size_t          bytes;
} VmafCudaKernelReadback;

int  vmaf_cuda_kernel_lifecycle_init(VmafCudaKernelLifecycle *,
                                     VmafCudaState *);

int  vmaf_cuda_kernel_readback_alloc(VmafCudaKernelReadback *,
                                     VmafCudaState *, size_t bytes);

int  vmaf_cuda_kernel_submit_pre_launch(VmafCudaKernelLifecycle *,
                                        VmafCudaState *,
                                        VmafCudaKernelReadback *,
                                        CUstream picture_stream,
                                        CUevent dist_ready_event);

int  vmaf_cuda_kernel_submit_post_record(VmafCudaKernelLifecycle *,
                                         VmafCudaState *);

int  vmaf_cuda_kernel_collect_wait(VmafCudaKernelLifecycle *,
                                   VmafCudaState *);

int  vmaf_cuda_kernel_lifecycle_close(VmafCudaKernelLifecycle *,
                                      VmafCudaState *);

int  vmaf_cuda_kernel_readback_free(VmafCudaKernelReadback *,
                                    VmafCudaState *);

What each helper covers¶

What stays in the kernel TU¶

• The per-metric cuLaunchKernel(...) call (grid dims, kernel parameter pack, function handle).
• The cuModuleLoadData / cuModuleGetFunction chain — kernel binary names and symbol counts vary per metric.
• The host-side reduction and score emission. PSNR's \(10 \log_{10}(\mathrm{peak}^2 / \mathrm{mse})\) is one line; ssimulacra2 has a 6-band pyramid pool. Neither belongs in a shared header.

Migration sketch¶

typedef struct PsnrStateCuda {
    VmafCudaKernelLifecycle  lc;
    VmafCudaKernelReadback   sse;
    /* metric-specific: kernel handles, max constants, dict... */
} PsnrStateCuda;

static int init_fex_cuda(VmafFeatureExtractor *fex, ...)
{
    PsnrStateCuda *s = fex->priv;
    int err = vmaf_cuda_kernel_lifecycle_init(&s->lc, fex->cu_state);
    if (err) return err;
    err = vmaf_cuda_kernel_readback_alloc(&s->sse, fex->cu_state,
                                          sizeof(uint64_t));
    if (err) return err;
    /* metric-specific: module load + function resolve, peak constants */
    return 0;
}

The before/after diff for integer_psnr_cuda.c is roughly −6 LOC of host-side scaffolding per kernel — small, but the win is mostly in the shared error-handling and partial-init unwind paths, not the line count.

HIP template¶

The HIP template mirrors the CUDA template field-for-field (T7-10b / ADR-0212). Unlike the CUDA variant which uses static inline helpers, the HIP helpers are out-of-line (kernel_template.c) because the ROCm HIP driver-loader table was not scaffolded at the time and the bodies require a -ENOSYS guard that inline callers cannot override. The runtime PR (T7-10b) replaced the stub bodies with real hipStreamCreate / hipEventCreate / hipMemcpyAsync calls. The struct shapes and helper signatures are stable.

Surface differences from CUDA¶

What each helper covers¶

Metal template¶

The Metal template mirrors the HIP template with one unified-memory simplification (T8-1b / ADR-0420): on Apple Silicon the device / host_pinned pair collapses to a single MTLBuffer allocated with MTLResourceStorageModeShared. Helpers are out-of-line Objective-C++ in kernel_template.mm (ARC) and bridge uintptr_t slots to id<MTL...> via __bridge_retained / __bridge_transfer.

Surface differences from HIP¶

What each helper covers¶

Lifecycle contract (shared across the templated backends)¶

Every templated backend (CUDA, HIP, Metal) follows the same four-phase sequence:

1. Init — allocate stream/queue, events, device accumulator, host readback slot. Returns 0 or a negative errno. Partial failures roll back in reverse order.
2. Submit — zero the accumulator, wait on the dist-ready event, launch the kernel, record a completion event.
3. Collect — wait for the completion event on the private stream/queue; copy the result to the host readback slot.
4. Close — synchronise the stream/queue, destroy events, free device and host buffers in reverse allocation order.

Backends diverge in how they represent handles (CUstream vs hipStream_t vs MTLCommandQueue) and in whether a device/host buffer split (CUDA, HIP) applies. The four-phase contract itself is invariant.

Migrating an existing kernel¶

Each kernel migration is its own PR, gated by:

1. Netflix golden (CPU only, untouched: the kernel does not run on the CPU lane).
2. /cross-backend-diff at places=4 against the CPU reference, on every Netflix golden YUV pair the kernel is registered against (per ADR-0214).
3. make lint clean on every touched file (hard rule 10 in agent hard rules).

The remaining candidates are the bespoke kernels listed in the coverage table above.

Why per-backend (not cross-backend)¶

The GPU-template scope analysis behind ADR-0246 established that CUDA's async-stream and event model and other APIs' command-buffer and fence models share no concrete shape. A cross-backend abstraction would force a lowest-common-denominator API that captures neither well. The per-backend split keeps each header honest about the platform it targets.

Why helper functions (not macros)¶

The CUDA template uses static inline helpers; HIP and Metal use out-of-line helpers for the reasons in their sections above. Shared trade-offs:

• Debug stepping: cuda-gdb and Nsight step through inline functions; macros expand to a single compound statement that shows up as one line in the source view.
• Type-checking: missing parameters or wrong-type pointers produce compiler errors at the helper site, not at some inscrutable point inside a macro expansion.
• The macros that do pay off live elsewhere: cuda_helper.cuh provides CHECK_CUDA_GOTO / CHECK_CUDA_RETURN, which are macros precisely because their goto label form needs textual substitution. The kernel-template helpers use those macros internally.

See also¶

• ADR-0246 — original CUDA + Vulkan template design decision and alternatives.
• ADR-0484 — HIP and Metal section addition.
• core/src/cuda/AGENTS.md — kernel template invariant row.
• core/src/hip/AGENTS.md — HIP kernel template invariant row.
• core/src/metal/AGENTS.md — Metal kernel template invariant row.
• docs/backends/cuda/overview.md — broader CUDA backend overview.
• docs/backends/vulkan/overview.md — historical Vulkan page (backend removed).
• docs/backends/hip/overview.md — broader HIP backend overview.
• docs/backends/metal/index.md — broader Metal backend overview.

History¶

Vulkan template (backend removed in ADR-0726)¶

The Vulkan template captured the descriptor-pool + pipeline + per-WG int64 partials shape that every Vulkan SSBO-only reduction kernel used. The reference implementation was core/src/feature/vulkan/psnr_vulkan.c (removed in ADR-0726).

Surface¶

#include "vulkan/kernel_template.h"

typedef struct VmafVulkanKernelPipeline {
    VkDescriptorSetLayout dsl;
    VkPipelineLayout      pipeline_layout;
    VkShaderModule        shader;
    VkPipeline            pipeline;
    VkDescriptorPool      desc_pool;
} VmafVulkanKernelPipeline;

typedef struct VmafVulkanKernelSubmit {
    VkCommandBuffer cmd;
    VkFence         fence;
} VmafVulkanKernelSubmit;

typedef struct VmafVulkanKernelPipelineDesc {
    uint32_t                       ssbo_binding_count;
    uint32_t                       push_constant_size;
    const uint32_t                *spv_bytes;
    size_t                         spv_size;
    VkComputePipelineCreateInfo    pipeline_create_info;
    uint32_t                       max_descriptor_sets;
} VmafVulkanKernelPipelineDesc;

int  vmaf_vulkan_kernel_pipeline_create(VmafVulkanContext *,
                                        const VmafVulkanKernelPipelineDesc *,
                                        VmafVulkanKernelPipeline *);

int  vmaf_vulkan_kernel_submit_begin(VmafVulkanContext *,
                                     VmafVulkanKernelSubmit *);

int  vmaf_vulkan_kernel_submit_end_and_wait(VmafVulkanContext *,
                                            VmafVulkanKernelSubmit *);

void vmaf_vulkan_kernel_submit_free(VmafVulkanContext *,
                                    VmafVulkanKernelSubmit *);

void vmaf_vulkan_kernel_pipeline_destroy(VmafVulkanContext *,
                                         VmafVulkanKernelPipeline *);

What each helper covers¶

What stays in the kernel TU¶

• The shader bytecode header (<feature>_spv.h) — generated per-kernel by the subdir('vulkan') glslc chain.
• The push-constant struct layout. PsnrPushConsts and a hypothetical Ssim4VifPushConsts have nothing in common.
• Spec-constant population — the caller fills pipeline_create_info.stage.pSpecializationInfo before calling vmaf_vulkan_kernel_pipeline_create.
• Per-frame buffer alloc, host upload, descriptor-set allocation + binding-write, dispatch grid math, host-side reduction. These shapes diverge enough between kernels that a unified API would be either too narrow (just PSNR's shape) or too generic (callbacks for everything).

Migration sketch¶

typedef struct PsnrVulkanState {
    VmafVulkanContext        *ctx;
    int                       owns_ctx;
    VmafVulkanKernelPipeline  pl;
    /* metric-specific: per-plane buffers, push-const cache, ... */
} PsnrVulkanState;

static int init(VmafFeatureExtractor *fex, ...)
{
    PsnrVulkanState *s = fex->priv;
    /* ... resolve s->ctx ... */
    VmafVulkanKernelPipelineDesc desc = {
        .ssbo_binding_count = 3,
        .push_constant_size = sizeof(PsnrPushConsts),
        .spv_bytes          = psnr_spv,
        .spv_size           = psnr_spv_size,
        .max_descriptor_sets = 12,
        /* caller fills stage.pName + spec_info on pipeline_create_info */
    };
    desc.pipeline_create_info.stage.pName = "main";
    desc.pipeline_create_info.stage.pSpecializationInfo = &spec_info;
    return vmaf_vulkan_kernel_pipeline_create(s->ctx, &desc, &s->pl);
}

static int extract(VmafFeatureExtractor *fex, ...)
{
    PsnrVulkanState *s = fex->priv;
    VmafVulkanKernelSubmit sub;
    int err = vmaf_vulkan_kernel_submit_begin(s->ctx, &sub);
    if (err) return err;

    /* metric-specific: allocate descriptor sets, write bindings,
     * record commands on sub.cmd, etc. */

    err = vmaf_vulkan_kernel_submit_end_and_wait(s->ctx, &sub);
    /* host-side reduce + score emit */
    vmaf_vulkan_kernel_submit_free(s->ctx, &sub);
    return err;
}

static int close_fex(VmafFeatureExtractor *fex)
{
    PsnrVulkanState *s = fex->priv;
    vmaf_vulkan_kernel_pipeline_destroy(s->ctx, &s->pl);
    /* metric-specific frees */
    return 0;
}

The before/after diff for psnr_vulkan.c is roughly −30 LOC — the five vkCreate/vkDestroy pairs collapse into two helper calls each, and the cleanup goto-ladder loses two labels.

Migration coverage record¶

Before the Vulkan removal the page recorded 16 of 20 CUDA kernels, 14 HIP kernels and 8 Metal kernels on the template. The current counts are in the coverage table at the top. ssimulacra2_cuda has since moved onto the template, and integer_motion_cuda runs the diff-first SAD kernel of integer_motion_sad_cuda.c (ADR-1372), not a ping-pong blur ring.