◆ op_track_keypoints()

bool MayaFlux::Yantra::VisionGpuExecutor::op_track_keypoints	(	VisionGpuContexts &	contexts,
		const Kinesis::Vision::VisionStep &	step
	)
private
Track the previous frame's keypoints into the current frame.
Adjacency mirrors ConnectedComponents into FindContours: ExtractPeaks must immediately precede, and its work is folded into this step so the detections are written straight into the flow context's buffers. The whole step is one dependency sequence: peaks, one Lucas-Kanade dispatch per pyramid level from coarse to fine, and the selection phases that build the next point list. Tracks persist: survivors carry over with their id and age, and new detections fill only the freed capacity, strongest first, one per grid cell. The first frame after a reset runs peaks and selection only.
Returns: True when the step was deferred and the run must suspend.
Definition at line 311 of file VisionFlowDispatch.cpp.
{
    auto& state = contexts.flow_state;
    auto& flow = contexts.flow;
    auto& foundry = Portal::Graphics::get_shader_foundry();
 
    contexts.pass.result.structured = std::vector<Kinesis::Vision::TrackResult> {};
    contexts.pass.result.w = 0;
    contexts.pass.result.h = 0;
 
    const auto* peaks = contexts.pass.behind();
    if (!peaks || peaks->op != VisionOp::ExtractPeaks) {
        MF_ERROR(Journal::Component::Yantra, Journal::Context::ComputeMatrix,
            "run_gpu: TrackKeypoints requires ExtractPeaks as the immediately preceding step");
        return false;
    }
 
    if (!state.curr_ready || !state.atlas[0]) {
        MF_ERROR(Journal::Component::Yantra, Journal::Context::ComputeMatrix,
            "run_gpu: TrackKeypoints requires an earlier RgbaToGray step to supply the gray frame");
        return false;
    }
 
    const auto& p = std::get<TrackKeypointsParams>(step.params);
    const auto& pk = std::get<ExtractPeaksParams>(peaks->params);
    const auto& layout = state.layout;
    const uint32_t w = contexts.pass.w;
    const uint32_t h = contexts.pass.h;
    const uint32_t max_points = std::clamp(p.max_points, 1U, k_flow_max_points);
 
    if (w != layout.level[0].w || h != layout.level[0].h) {
        MF_ERROR(Journal::Component::Yantra, Journal::Context::ComputeMatrix,
            "run_gpu: TrackKeypoints image is {}x{} but the gray frame was captured at {}x{}",
            w, h, layout.level[0].w, layout.level[0].h);
        return false;
    }
 
    const bool export_tracks = p.export_tracks;
    const uint32_t duplicate_cutoff = (state.last_track_count > 0 && state.last_export)
        ? k_flow_duplicate_energy + 1U
        : 0U;
 
    const auto response = contexts.pass.current;
    const auto atlas_a = state.atlas[0];
    const auto atlas_b = state.atlas[1];
 
    const ExtractPeaksPC peaks_pc {
        .threshold = pk.threshold,
        .nms_radius = pk.nms_radius,
        .width = w,
        .height = h,
        .max_keypoints = k_flow_max_points,
        .rect = region_rect(pk.region, w, h),
    };
 
    std::vector<DependencyStage> stages;
    stages.reserve(layout.levels * (p.forward_backward_threshold > 0.0F ? 2U : 1U) + 8U);
 
    stages.push_back({
        .config = { .shader_path = "extract_peaks.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(ExtractPeaksPC) },
        .stage_fn = [response, peaks_pc](GpuDispatchCore& ctx) {
                ctx.stage_image_at(0, response, GpuBufferBinding::ElementType::IMAGE_STORAGE);
                ctx.set_push_constants(peaks_pc); },
        .hazard_fn = [](GpuDispatchCore& ctx) { return flow_hazards(ctx, { k_flow_det_count, k_flow_det_points }); },
        .explicit_groups = std::array<uint32_t, 3> { (w + k_wg2d[0] - 1U) / k_wg2d[0], (h + k_wg2d[1] - 1U) / k_wg2d[1], 1U },
    });
 
    const auto base_w = static_cast<float>(layout.level[0].w);
    const auto base_h = static_cast<float>(layout.level[0].h);
    float cell = std::max(p.min_distance, 1.0F);
    auto grid_w = static_cast<uint32_t>(std::ceil(base_w / cell));
    auto grid_h = static_cast<uint32_t>(std::ceil(base_h / cell));
    while (static_cast<uint64_t>(grid_w) * grid_h > k_flow_grid_capacity) {
        cell *= 1.1F;
        grid_w = static_cast<uint32_t>(std::ceil(base_w / cell));
        grid_h = static_cast<uint32_t>(std::ceil(base_h / cell));
    }
    const uint32_t grid_cells = grid_w * grid_h;
    const uint32_t clear_groups = std::max(1U, (grid_cells + k_flow_select_local - 1U) / k_flow_select_local);
    const uint32_t candidate_groups = (k_flow_max_points + k_flow_select_local - 1U) / k_flow_select_local;
    const uint32_t have_prev = state.have_prev ? 1U : 0U;
 
    const auto add_select = [&](FlowSelectPhase phase, uint32_t groups, std::vector<FlowBufferSpec> touched, std::optional<FlowArgs> args = std::nullopt) {
        const FlowSelectPC pc {
            .phase = static_cast<uint32_t>(phase),
            .max_points = max_points,
            .have_prev = have_prev,
            .grid_w = grid_w,
            .grid_h = grid_h,
            .grid_cells = grid_cells,
            .cell = cell,
            .base_w = base_w,
            .base_h = base_h,
            .publish_slot = state.export_slot,
            .duplicate_cutoff = duplicate_cutoff,
        };
        stages.push_back({
            .config = { .shader_path = "flow_select.comp.spv", .workgroup_size = { k_flow_select_local, 1, 1 }, .push_constant_size = sizeof(FlowSelectPC) },
            .stage_fn = [pc](GpuDispatchCore& ctx) { ctx.set_push_constants(pc); },
            .hazard_fn = [touched = std::move(touched)](GpuDispatchCore& ctx) { return flow_hazards(ctx, touched); },
            .explicit_groups = std::array<uint32_t, 3> { groups, 1U, 1U },
            .indirect_groups = args ? std::optional<IndirectGroupsSource> { flow_args_source(*args) } : std::nullopt,
        });
    };
 
    add_select(FlowSelectPhase::ARGS, 1U, {});
 
    const auto add_lk = [&](bool backward) {
        for (uint32_t level = layout.levels; level-- > 0;) {
            const PyramidLevel lv = layout.level[level];
            const bool finest = level == 0;
 
            const FlowLkPC pc {
                .curr_atlas = state.curr,
                .level = level,
                .coarsest = level + 1U == layout.levels ? 1U : 0U,
                .pad0 = 0,
                .lvl_ox = lv.ox,
                .lvl_oy = lv.oy,
                .lvl_w = lv.w,
                .lvl_h = lv.h,
                .base_w = layout.level[0].w,
                .base_h = layout.level[0].h,
                .window_radius = std::min(p.window_radius, k_flow_max_radius),
                .max_iterations = p.max_iterations,
                .eigen_threshold = p.eigen_threshold,
                .error_threshold = p.error_threshold,
                .max_points = max_points,
                .pad1 = 0,
                .backward = backward ? 1U : 0U,
                .forward_backward_threshold = p.forward_backward_threshold,
            };
 
            stages.push_back({
                .config = { .shader_path = "flow_lk.comp.spv", .workgroup_size = { 64, 1, 1 }, .push_constant_size = sizeof(FlowLkPC) },
                .stage_fn = [atlas_a, atlas_b, pc](GpuDispatchCore& ctx) {
                        ctx.stage_image_at(3, atlas_a, GpuBufferBinding::ElementType::IMAGE_STORAGE);
                        ctx.stage_image_at(4, atlas_b, GpuBufferBinding::ElementType::IMAGE_STORAGE);
                        ctx.set_push_constants(pc); },
                .hazard_fn = [finest, backward](GpuDispatchCore& ctx) {
                        if (finest || backward)
                            return flow_hazards(ctx, { k_flow_tracks, k_flow_meta, k_flow_prev_count, k_flow_prev_points });
                        return flow_hazards(ctx, { k_flow_tracks, k_flow_meta }); },
                .explicit_groups = std::array<uint32_t, 3> { k_flow_max_points, 1U, 1U },
                .indirect_groups = flow_args_source(FlowArgs::LK),
            });
        }
    };
    if (state.have_prev) {
        add_lk(false);
        if (p.forward_backward_threshold > 0.0F)
            add_lk(true);
    }
 
    add_select(FlowSelectPhase::CLEAR, clear_groups, { k_flow_grid, k_flow_histogram, k_flow_counters });
    add_select(FlowSelectPhase::SURVIVORS, candidate_groups, { k_flow_next_points, k_flow_counters, k_flow_grid }, FlowArgs::RETAINED);
    add_select(FlowSelectPhase::HISTOGRAM, candidate_groups, { k_flow_histogram }, FlowArgs::DETECTED);
    add_select(FlowSelectPhase::THRESHOLD, 1U, { k_flow_counters });
    add_select(FlowSelectPhase::ADD_STRONG, candidate_groups, { k_flow_next_points, k_flow_counters, k_flow_grid }, FlowArgs::DETECTED);
    add_select(FlowSelectPhase::ADD_MARGINAL, candidate_groups, { k_flow_next_points, k_flow_counters, k_flow_grid }, FlowArgs::DETECTED);
    add_select(FlowSelectPhase::COMMIT, 1U, {});
    if (export_tracks)
        add_select(FlowSelectPhase::PUBLISH, candidate_groups, { k_flow_export[0], k_flow_export[1] }, FlowArgs::PUBLISH);
 
    ExecutionContext track_ctx;
    track_ctx.mode = ExecutionMode::DEPENDENCY;
    DependencyParams track_params;
    track_params.stages = stages;
    track_params.async = true;
    track_ctx.parameters = track_params;
    const auto track_dispatch_result = flow.execute(Datum<> {}, track_ctx);
    const auto fence = track_dispatch_result.get_metadata<Portal::Graphics::FenceID>("gpu_fence").value_or(Portal::Graphics::INVALID_FENCE);
    if (fence == Portal::Graphics::INVALID_FENCE) {
        MF_ERROR(Journal::Component::Yantra, Journal::Context::ComputeMatrix,
            "run_gpu: TrackKeypoints failed to submit its dispatch sequence");
        return false;
    }
    state.export_pending = export_tracks;
    state.host_pending = p.host_tracks;
 
    if (step.deferred) {
        contexts.suspended.fence = fence;
        contexts.suspended.finalize = [](VisionGpuContexts& c) { finish_track(c); };
        return true;
    }
 
    foundry.wait_for_fence(fence);
    foundry.release_fence(fence);
    finish_track(contexts);
    return false;
}
Here is the call graph for this function: