op_connected_components()

void MayaFlux::Yantra::VisionGpuExecutor::op_connected_components ( VisionGpuContexts &  contexts, const Kinesis::Vision::ConnectedComponentsParams &  p  )

private

Label components in the current image on the GPU.

Produces component counts and bounds for structured output. When FindContours follows immediately, the labels stay on the GPU for that step and component bounds are not read back to the host.

Definition at line 10 of file VisionComponentsDispatch.cpp.

{
    auto& cc_pipeline = contexts.component_contours;
    auto w = contexts.pass.w;
    auto h = contexts.pass.h;
    auto& foundry = Portal::Graphics::get_shader_foundry();
 
    const auto seed_input = contexts.pass.current;
 
    const uint32_t block_width = (w + 1U) / 2U;
    const uint32_t block_height = (h + 1U) / 2U;
    cc_pipeline.ensure_shared_buffer(0, 2, static_cast<size_t>(block_width) * block_height, GpuBufferBinding::ElementType::UINT32);
    cc_pipeline.ensure_shared_buffer(0, 3, 1, GpuBufferBinding::ElementType::UINT32);
    cc_pipeline.ensure_shared_buffer(0, 4, static_cast<size_t>(block_width) * block_height, GpuBufferBinding::ElementType::UINT32);
    cc_pipeline.ensure_shared_buffer(0, 5, 1, GpuBufferBinding::ElementType::UINT32);
    cc_pipeline.ensure_shared_buffer(0, 6, static_cast<size_t>(w) * h, GpuBufferBinding::ElementType::UINT32,
        Portal::Graphics::BufferUsageHint::COMPUTE);
    cc_pipeline.ensure_shared_buffer(0, 7, static_cast<size_t>(k_max_components) * 2, GpuBufferBinding::ElementType::UINT32);
    cc_pipeline.ensure_shared_buffer(0, 8, static_cast<size_t>(k_max_components) * 2, GpuBufferBinding::ElementType::UINT32);
    cc_pipeline.ensure_shared_buffer(0, 9, k_max_components, GpuBufferBinding::ElementType::UINT32);
    cc_pipeline.ensure_shared_buffer(0, 10, 6, GpuBufferBinding::ElementType::UINT32,
        Portal::Graphics::BufferUsageHint::INDIRECT);
 
    const CCBlockInitPC init_pc { .width = w, .height = h, .block_width = block_width, .block_height = block_height };
    const CCMergePC merge_pc { .width = w, .height = h, .block_width = block_width, .block_height = block_height };
    const auto* next = contexts.pass.ahead();
    const bool contours_follow = next && next->op == VisionOp::FindContours;
    const uint32_t export_labels = (p.export_labels || contours_follow) ? 1U : 0U;
 
    const CCFinalLabelPC final_pc {
        .width = w,
        .height = h,
        .block_width = block_width,
        .block_height = block_height,
        .max_components = k_max_components,
        .export_labels = export_labels
    };
 
    const std::array<uint32_t, 3> block_groups {
        (block_width + k_wg2d[0] - 1U) / k_wg2d[0],
        (block_height + k_wg2d[1] - 1U) / k_wg2d[1],
        1U
    };
 
    const uint32_t reset_elements = std::max(block_width * block_height, k_max_components);
    const auto buffer_hazards = [](GpuDispatchCore& ctx, std::initializer_list<size_t> bindings) {
        std::vector<HazardResource> hazards;
        hazards.reserve(bindings.size());
        for (const size_t binding : bindings) {
            hazards.push_back(ctx.shared_buffer_hazard(
                { .set = 0, .binding = static_cast<uint32_t>(binding), .direction = GpuBufferBinding::Direction::INPUT_OUTPUT, .element_type = GpuBufferBinding::ElementType::UINT32 }));
        }
        return hazards;
    };
 
    std::shared_ptr<Core::VKImage> cc_image;
    std::vector<DependencyStage> cc_stages;
    cc_stages.reserve(5);
 
    cc_stages.push_back({
        .config = { .shader_path = "cc_reset.comp.spv", .workgroup_size = { 256, 1, 1 }, .push_constant_size = sizeof(CCResetPC) },
        .stage_fn = [=](GpuDispatchCore& ctx) { ctx.set_push_constants(CCResetPC { .lut_size = block_width * block_height, .max_components = k_max_components }); },
        .hazard_fn = [=](GpuDispatchCore& ctx) { return buffer_hazards(ctx, { 4, 5, 7, 8, 9 }); },
        .explicit_groups = std::array<uint32_t, 3> { (reset_elements + 255U) / 256U, 1U, 1U },
    });
 
    cc_stages.push_back({
        .config = { .shader_path = "cc_block_init.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(CCBlockInitPC) },
        .stage_fn = [seed_input, init_pc](GpuDispatchCore& ctx) {
                ctx.stage_image_at(1, seed_input, GpuBufferBinding::ElementType::IMAGE_STORAGE);
                ctx.set_push_constants(init_pc); },
        .hazard_fn = [=](GpuDispatchCore& ctx) { return buffer_hazards(ctx, { 2 }); },
        .explicit_groups = block_groups,
    });
 
    cc_stages.push_back({
        .config = { .shader_path = "cc_merge.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(CCMergePC) },
        .stage_fn = [seed_input, merge_pc](GpuDispatchCore& ctx) {
                ctx.stage_image_at(1, seed_input, GpuBufferBinding::ElementType::IMAGE_STORAGE);
                ctx.set_push_constants(merge_pc); },
        .hazard_fn = [=](GpuDispatchCore& ctx) { return buffer_hazards(ctx, { 2 }); },
        .explicit_groups = block_groups,
    });
 
    cc_stages.push_back({
        .config = { .shader_path = "cc_compress.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(CCCompressPC) },
        .stage_fn = [=](GpuDispatchCore& ctx) { ctx.set_push_constants(CCCompressPC { .block_width = block_width, .block_height = block_height }); },
        .hazard_fn = [=](GpuDispatchCore& ctx) { return buffer_hazards(ctx, { 2, 4 }); },
        .explicit_groups = block_groups,
    });
 
    cc_stages.push_back({
        .config = { .shader_path = "cc_final_label.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(CCFinalLabelPC) },
        .stage_fn = [seed_input, final_pc, w, h, &cc_pipeline, &cc_image](GpuDispatchCore& ctx) {
            ctx.stage_image_at(1, seed_input, GpuBufferBinding::ElementType::IMAGE_STORAGE);
            ctx.set_push_constants(final_pc);
            cc_pipeline.prepare_output_image(w, h);
            cc_image = cc_pipeline.get_output_image(0);
        },
        .explicit_groups = std::array<uint32_t, 3> { (w + k_wg2d[0] - 1U) / k_wg2d[0], (h + k_wg2d[1] - 1U) / k_wg2d[1], 1U },
    });
 
    ExecutionContext cc_ctx;
    cc_ctx.mode = ExecutionMode::DEPENDENCY;
    DependencyParams cc_params;
    cc_params.stages = cc_stages;
    cc_params.async = true;
    cc_ctx.parameters = cc_params;
    const auto cc_dispatch_result = cc_pipeline.execute(Datum<> {}, cc_ctx);
    const auto fence = cc_dispatch_result.get_metadata<FenceID>("gpu_fence").value_or(INVALID_FENCE);
    foundry.wait_for_fence(fence);
    foundry.release_fence(fence);
    cc_pipeline.clear_output_dimensions();
 
    contexts.pass.result.images.component_colors = p.with_colors ? cc_image : nullptr;
 
    if (p.export_label_buffer) {
        contexts.pass.result.buffers.labels = std::make_shared<Portal::Graphics::GpuBufferHandle>(
            cc_pipeline.shared_buffer_handle(0, 6));
    }
 
    if (contours_follow && !p.export_boxes)
        return;
 
    uint32_t compact_count = 0;
    cc_pipeline.download_shared(0, 5, &compact_count, sizeof(uint32_t));
    compact_count = std::min(compact_count, k_max_components);
 
    Kinesis::Vision::ComponentResult cc_result;
    cc_result.count = compact_count;
    cc_result.boxes.reserve(compact_count);
 
    if (compact_count > 0) {
        std::vector<glm::uvec2> bmin(compact_count);
        std::vector<glm::uvec2> bmax(compact_count);
        std::vector<uint32_t> bcount(compact_count);
        cc_pipeline.download_shared(0, 7, bmin.data(), bmin.size() * sizeof(glm::uvec2));
        cc_pipeline.download_shared(0, 8, bmax.data(), bmax.size() * sizeof(glm::uvec2));
        cc_pipeline.download_shared(0, 9, bcount.data(), bcount.size() * sizeof(uint32_t));
 
        const float inv_w = 1.0F / static_cast<float>(w);
        const float inv_h = 1.0F / static_cast<float>(h);
 
        for (uint32_t i = 0; i < compact_count; ++i) {
            if (bcount[i] == 0)
                continue;
            const float x = static_cast<float>(bmin[i].x) * inv_w;
            const float y = static_cast<float>(bmin[i].y) * inv_h;
            const float bw = static_cast<float>(bmax[i].x - bmin[i].x + 1) * inv_w;
            const float bh = static_cast<float>(bmax[i].y - bmin[i].y + 1) * inv_h;
            cc_result.boxes.push_back({ .x = x, .y = y, .w = bw, .h = bh, .confidence = 1.0F, .label_id = i + 1 });
        }
    }
 
    contexts.pass.result.component_boxes = cc_result.boxes;
 
    if (contours_follow)
        return;
 
    contexts.pass.result.structured = std::move(cc_result);
    contexts.pass.result.w = 0;
    contexts.pass.result.h = 0;
}

References MayaFlux::Kinesis::Vision::VisionPass< Handle >::ahead(), MayaFlux::Yantra::DependencyParams::async, MayaFlux::Yantra::VisionInternal::CCCompressPC::block_width, MayaFlux::Kinesis::Vision::ComponentResult::boxes, MayaFlux::Kinesis::Vision::VisionResult::buffers, MayaFlux::Kinesis::Vision::VisionResult::component_boxes, MayaFlux::Kinesis::Vision::ImageOutputs::component_colors, MayaFlux::Yantra::VisionGpuContexts::component_contours, MayaFlux::Portal::Graphics::COMPUTE, MayaFlux::Kinesis::Vision::ComponentResult::count, MayaFlux::Kinesis::Vision::VisionPass< Handle >::current, MayaFlux::Yantra::DEPENDENCY, MayaFlux::Kinesis::Vision::ConnectedComponentsParams::export_boxes, MayaFlux::Kinesis::Vision::ConnectedComponentsParams::export_label_buffer, MayaFlux::Kinesis::Vision::ConnectedComponentsParams::export_labels, fence, MayaFlux::Portal::Graphics::get_shader_foundry(), MayaFlux::Kinesis::Vision::VisionResult::h, MayaFlux::Kinesis::Vision::VisionPass< Handle >::h, h, MayaFlux::Portal::Graphics::GpuBufferBinding::IMAGE_STORAGE, MayaFlux::Kinesis::Vision::VisionResult::images, MayaFlux::Portal::Graphics::INDIRECT, MayaFlux::Portal::Graphics::INVALID_FENCE, MayaFlux::Yantra::VisionInternal::k_max_components, MayaFlux::Yantra::VisionInternal::k_wg2d, MayaFlux::Kinesis::Vision::BufferOutputs::labels, MayaFlux::Yantra::VisionInternal::CCResetPC::lut_size, MayaFlux::Yantra::ExecutionContext::mode, MayaFlux::Kinesis::Vision::VisionStep::op, MayaFlux::Yantra::ExecutionContext::parameters, MayaFlux::Yantra::VisionGpuContexts::pass, MayaFlux::Kinesis::Vision::VisionPass< Handle >::result, MayaFlux::Yantra::GpuDispatchCore::set_push_constants(), MayaFlux::Yantra::GpuDispatchCore::shared_buffer_hazard(), MayaFlux::Yantra::GpuDispatchCore::stage_image_at(), MayaFlux::Yantra::DependencyParams::stages, MayaFlux::Kinesis::Vision::VisionResult::structured, MayaFlux::Portal::Graphics::GpuBufferBinding::UINT32, MayaFlux::Kinesis::Vision::VisionResult::w, MayaFlux::Kinesis::Vision::VisionPass< Handle >::w, MayaFlux::Yantra::VisionInternal::CCBlockInitPC::width, MayaFlux::Yantra::VisionInternal::CCMergePC::width, MayaFlux::Yantra::VisionInternal::CCFinalLabelPC::width, and MayaFlux::Kinesis::Vision::ConnectedComponentsParams::with_colors.

Here is the call graph for this function: