build_flow_pyramid()

void MayaFlux::Yantra::VisionGpuExecutor::build_flow_pyramid ( VisionGpuContexts &  contexts, uint32_t  requested_levels, uint32_t  src_channel  )

staticprivate

Build the current frame's pyramid atlas from the working image.

Submitted as one un-awaited dependency sequence, one fused dispatch per level. The level 0 stage carries a hazard on the working image, which orders any later dispatch that overwrites it after this read. The fence is reaped on the next fresh run and in reset().

Parameters

src_channel

Channel of the working image read as intensity: 0 for a gray frame, 2 for the value channel of an HSV frame.

Definition at line 709 of file VisionFlowDispatch.cpp.

{
    auto& flow = contexts.flow;
    auto& state = contexts.flow_state;
    const auto gray = contexts.pass.current;
 
    if (!gray || !gray->is_initialized())
        return;
 
    const uint32_t w = contexts.pass.w;
    const uint32_t h = contexts.pass.h;
    const auto layout = pyramid_layout(w, h, requested_levels, k_flow_min_level_extent);
 
    ensure_flow_resources(contexts);
 
    if (!state.atlas[0] || !(state.layout == layout)) {
        auto& loom = Portal::Graphics::TextureLoom::instance();
        state.atlas[0] = loom.create_storage_image(layout.atlas_w, layout.atlas_h, ImageFormat::RGBA16F);
        state.atlas[1] = loom.create_storage_image(layout.atlas_w, layout.atlas_h, ImageFormat::RGBA16F);
        state.layout = layout;
        state.have_prev = false;
        state.curr = 0;
    }
 
    const auto atlas_a = state.atlas[0];
    const auto atlas_b = state.atlas[1];
    const auto target = state.atlas[state.curr];
 
    std::vector<DependencyStage> stages;
    stages.reserve(layout.levels);
 
    for (uint32_t level = 0; level < layout.levels; ++level) {
        const PyramidLevel dst = layout.level[level];
        const PyramidLevel src = level == 0 ? PyramidLevel { .ox = 0, .oy = 0, .w = w, .h = h } : layout.level[level - 1];
 
        const FlowPyramidPC pc {
            .target_atlas = state.curr,
            .level = level,
            .src_w = src.w,
            .src_h = src.h,
            .src_ox = src.ox,
            .src_oy = src.oy,
            .dst_w = dst.w,
            .dst_h = dst.h,
            .dst_ox = dst.ox,
            .dst_oy = dst.oy,
            .src_channel = src_channel,
        };
 
        stages.push_back({
            .config = { .shader_path = "flow_pyramid.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(FlowPyramidPC) },
            .stage_fn = [gray, atlas_a, atlas_b, pc](GpuDispatchCore& ctx) {
                    ctx.stage_image_at(0, gray, GpuBufferBinding::ElementType::IMAGE_STORAGE);
                    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 = [target, gray, level](GpuDispatchCore&) {
                    std::vector<HazardResource> hazards;
                    hazards.push_back({
                        .binding = { .set = 0, .binding = 3, .direction = GpuBufferBinding::Direction::INPUT_OUTPUT, .element_type = GpuBufferBinding::ElementType::IMAGE_STORAGE },
                        .image = target->get_image(),
                    });
                    if (level == 0) {
                        hazards.push_back({
                            .binding = { .set = 0, .binding = 0, .direction = GpuBufferBinding::Direction::INPUT_OUTPUT, .element_type = GpuBufferBinding::ElementType::IMAGE_STORAGE },
                            .image = gray->get_image(),
                        });
                    }
                    return hazards; },
            .explicit_groups = std::array<uint32_t, 3> { (dst.w + k_wg2d[0] - 1U) / k_wg2d[0], (dst.h + k_wg2d[1] - 1U) / k_wg2d[1], 1U },
        });
    }
 
    auto& foundry = Portal::Graphics::get_shader_foundry();
    if (state.build_fence != Portal::Graphics::INVALID_FENCE) {
        foundry.wait_for_fence(state.build_fence);
        foundry.release_fence(state.build_fence);
        state.build_fence = Portal::Graphics::INVALID_FENCE;
    }
 
    const std::array<uint32_t, 4> cleared {};
    flow.upload_shared_raw(1, 3, reinterpret_cast<const uint8_t*>(cleared.data()), sizeof(cleared));
 
    ExecutionContext pyramid_ctx;
    pyramid_ctx.mode = ExecutionMode::DEPENDENCY;
    DependencyParams pyramid_params;
    pyramid_params.stages = stages;
    pyramid_params.async = true;
    pyramid_ctx.parameters = pyramid_params;
    const auto pyramid_dispatch_result = flow.execute(Datum<> {}, pyramid_ctx);
    state.build_fence = pyramid_dispatch_result.get_metadata<Portal::Graphics::FenceID>("gpu_fence").value_or(Portal::Graphics::INVALID_FENCE);
    state.curr_ready = state.build_fence != Portal::Graphics::INVALID_FENCE;
}

References MayaFlux::Yantra::DependencyParams::async, MayaFlux::Kinesis::Vision::VisionPass< Handle >::current, MayaFlux::Yantra::DEPENDENCY, dst, MayaFlux::Yantra::VisionGpuContexts::flow, MayaFlux::Yantra::VisionGpuContexts::flow_state, MayaFlux::Portal::Graphics::get_shader_foundry(), MayaFlux::Kinesis::Vision::VisionPass< Handle >::h, h, MayaFlux::Portal::Graphics::GpuBufferBinding::IMAGE_STORAGE, MayaFlux::Portal::Graphics::GpuBufferBinding::INPUT_OUTPUT, MayaFlux::Portal::Graphics::TextureLoom::instance(), MayaFlux::Portal::Graphics::INVALID_FENCE, MayaFlux::Yantra::VisionInternal::k_flow_min_level_extent, MayaFlux::Yantra::VisionInternal::k_wg2d, MayaFlux::Yantra::ExecutionContext::mode, MayaFlux::Kinesis::Vision::PyramidLevel::ox, MayaFlux::Yantra::ExecutionContext::parameters, MayaFlux::Yantra::VisionGpuContexts::pass, MayaFlux::Kinesis::Vision::pyramid_layout(), src, MayaFlux::Yantra::DependencyParams::stages, MayaFlux::Yantra::VisionInternal::FlowPyramidPC::target_atlas, and MayaFlux::Kinesis::Vision::VisionPass< Handle >::w.

Here is the call graph for this function: