op_threshold_otsu()

GpuVisionPass::Completed MayaFlux::Yantra::VisionGpuExecutor::op_threshold_otsu ( VisionGpuContexts &  contexts, const Kinesis::Vision::OtsuParams &  p  )

private

Apply an Otsu threshold to the current image.

Computes the histogram and threshold on the GPU, then replaces the working image with the binary result and publishes it as VisionResult::images.threshold_otsu.

Returns

The resulting image and the image it was derived from.

Definition at line 37 of file VisionImageDispatch.cpp.

{
    auto& pixel_ctx = contexts.pixel;
    auto w = contexts.pass.w;
    auto h = contexts.pass.h;
    auto& foundry = Portal::Graphics::get_shader_foundry();
 
    const auto otsu_input = contexts.pass.current;
 
    pixel_ctx.ensure_shared_buffer(0, 3, 768, GpuBufferBinding::ElementType::UINT32,
        Portal::Graphics::BufferUsageHint::COMPUTE);
    pixel_ctx.ensure_shared_buffer(0, 4, 3, GpuBufferBinding::ElementType::UINT32,
        Portal::Graphics::BufferUsageHint::COMPUTE);
 
    const std::array<uint32_t, 768> histogram_reset {};
    pixel_ctx.upload_shared_raw(0, 3, reinterpret_cast<const uint8_t*>(histogram_reset.data()), sizeof(histogram_reset));
 
    const auto channels = static_cast<uint32_t>(p.channels);
    const OtsuHistPC hist_pc { .width = w, .height = h, .channels = channels };
    const OtsuApplyPC apply_pc { .channels = channels };
    const uint32_t select_groups = channels == 0U ? 1U : 3U;
    const auto otsu_hazard = [](uint32_t binding) {
        return [binding](GpuDispatchCore& ctx) {
            return std::vector<HazardResource> {
                ctx.shared_buffer_hazard({ .set = 0, .binding = binding, .direction = GpuBufferBinding::Direction::INPUT_OUTPUT, .element_type = GpuBufferBinding::ElementType::UINT32 }),
            };
        };
    };
    const std::array<uint32_t, 3> image_groups { (w + k_wg2d[0] - 1U) / k_wg2d[0], (h + k_wg2d[1] - 1U) / k_wg2d[1], 1U };
 
    std::shared_ptr<Core::VKImage> thresholded;
    std::vector<DependencyStage> otsu_stages;
    otsu_stages.reserve(3);
    otsu_stages.push_back({
        .config = { .shader_path = "otsu_histogram.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(OtsuHistPC) },
        .stage_fn = [&pixel_ctx, otsu_input, hist_pc](GpuDispatchCore& ctx) {
                pixel_ctx.stage_image(otsu_input);
                ctx.set_push_constants(hist_pc); },
        .hazard_fn = otsu_hazard(3),
        .explicit_groups = image_groups,
    });
    otsu_stages.push_back({
        .config = { .shader_path = "otsu_select.comp.spv", .workgroup_size = { 256, 1, 1 } },
        .stage_fn = [&pixel_ctx, otsu_input](GpuDispatchCore&) { pixel_ctx.stage_image(otsu_input); },
        .hazard_fn = otsu_hazard(4),
        .explicit_groups = std::array<uint32_t, 3> { select_groups, 1U, 1U },
    });
    otsu_stages.push_back({
        .config = { .shader_path = "otsu_apply.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(OtsuApplyPC) },
        .stage_fn = [&pixel_ctx, &thresholded, otsu_input, w, h, apply_pc](GpuDispatchCore& ctx) {
                pixel_ctx.stage_image(otsu_input);
                ctx.set_push_constants(apply_pc);
                pixel_ctx.prepare_output_image(w, h);
                thresholded = pixel_ctx.get_output_image(0); },
        .explicit_groups = image_groups,
    });
 
    {
        ExecutionContext otsu_ctx;
        otsu_ctx.mode = ExecutionMode::DEPENDENCY;
        DependencyParams otsu_params;
        otsu_params.stages = otsu_stages;
        otsu_params.async = true;
        otsu_ctx.parameters = otsu_params;
        const auto result = pixel_ctx.execute(Datum<> {}, otsu_ctx);
        const auto f = result.get_metadata<FenceID>("gpu_fence").value_or(INVALID_FENCE);
        foundry.wait_for_fence(f);
        foundry.release_fence(f);
    }
 
    contexts.pass.result.images.threshold_otsu = thresholded;
    contexts.pass.current = thresholded;
    contexts.pass.result.structured = std::monostate {};
 
    return { .output = thresholded, .input = otsu_input };
}

References MayaFlux::Yantra::DependencyParams::async, MayaFlux::Kinesis::Vision::OtsuParams::channels, channels, MayaFlux::Yantra::VisionInternal::OtsuApplyPC::channels, MayaFlux::Portal::Graphics::COMPUTE, MayaFlux::Kinesis::Vision::VisionPass< Handle >::current, MayaFlux::Yantra::DEPENDENCY, MayaFlux::Portal::Graphics::get_shader_foundry(), MayaFlux::Kinesis::Vision::VisionPass< Handle >::h, h, MayaFlux::Kinesis::Vision::VisionResult::images, MayaFlux::Portal::Graphics::GpuBufferBinding::INPUT_OUTPUT, MayaFlux::Portal::Graphics::INVALID_FENCE, MayaFlux::Yantra::VisionInternal::k_wg2d, MayaFlux::Yantra::ExecutionContext::mode, MayaFlux::Yantra::ExecutionContext::parameters, MayaFlux::Yantra::VisionGpuContexts::pass, MayaFlux::Yantra::VisionGpuContexts::pixel, MayaFlux::Kinesis::Vision::VisionPass< Handle >::result, MayaFlux::Yantra::DependencyParams::stages, MayaFlux::Kinesis::Vision::VisionResult::structured, MayaFlux::Kinesis::Vision::ImageOutputs::threshold_otsu, MayaFlux::Portal::Graphics::GpuBufferBinding::UINT32, MayaFlux::Kinesis::Vision::VisionPass< Handle >::w, and MayaFlux::Yantra::VisionInternal::OtsuHistPC::width.

Here is the call graph for this function: