24 struct PatchExtractPC {
36 struct PointSamplePC {
65 constexpr std::array<uint32_t, 3> k_wg_image { 16, 16, 1 };
66 constexpr std::array<uint32_t, 3> k_wg_reduce { 256, 1, 1 };
68 [[nodiscard]]
GpuBufferBinding binding_at(uint32_t binding, Direction direction, ElementType
type)
70 return { .
set = 0, .binding = binding, .direction = direction, .element_type =
type };
77 [[nodiscard]] std::shared_ptr<TextureExecutionContext> make_image_context()
79 return std::make_shared<TextureExecutionContext>(
81 std::vector<GpuBufferBinding> {
82 binding_at(0, Direction::OUTPUT, ElementType::IMAGE_STORAGE),
83 binding_at(1, Direction::INPUT, ElementType::IMAGE_SAMPLED),
84 binding_at(2, Direction::INPUT, ElementType::IMAGE_SAMPLED),
85 binding_at(3, Direction::INPUT, ElementType::FLOAT32),
86 binding_at(4, Direction::INPUT, ElementType::FLOAT32),
87 binding_at(5, Direction::INPUT, ElementType::FLOAT32),
96 [[nodiscard]] std::shared_ptr<TextureExecutionContext> make_reduce_context()
98 return std::make_shared<TextureExecutionContext>(
100 std::vector<GpuBufferBinding> {
101 binding_at(1, Direction::INPUT, ElementType::IMAGE_SAMPLED),
102 binding_at(3, Direction::INPUT, ElementType::FLOAT32),
103 binding_at(6, Direction::OUTPUT, ElementType::FLOAT32),
112 void use_shader(TextureExecutionContext& ctx,
const char* shader, std::array<uint32_t, 3> workgroup,
size_t push_size)
114 ctx.swap_shader({ .shader_path = shader, .workgroup_size = workgroup, .push_constant_size = push_size });
115 for (
size_t b = 3;
b <= 5; ++
b)
116 ctx.set_binding_data(
b, std::span<const float> {});
117 ctx.stage_image_at(1,
nullptr, ElementType::IMAGE_SAMPLED);
118 ctx.stage_image_at(2,
nullptr, ElementType::IMAGE_SAMPLED);
121 void dispatch_and_wait(TextureExecutionContext& ctx)
124 const auto fence = ctx.dispatch_async({});
125 ctx.clear_output_dimensions();
126 foundry.wait_for_fence(
fence);
127 foundry.release_fence(
fence);
130 [[nodiscard]] std::shared_ptr<Core::VKImage> dispatch_image(TextureExecutionContext& ctx, uint32_t w, uint32_t
h)
132 ctx.prepare_output_image(w,
h);
133 ctx.set_output_dimensions(w,
h);
134 dispatch_and_wait(ctx);
135 return ctx.get_output_image(0);
139 [[nodiscard]] std::vector<float> dispatch_reduce(TextureExecutionContext& ctx, uint32_t invocations,
size_t output_bytes)
141 ctx.set_output_size(6, output_bytes);
142 ctx.set_output_dimensions(invocations, 1);
143 dispatch_and_wait(ctx);
145 const auto result = ctx.collect_result();
146 const auto it = result.aux.find(6);
147 if (it == result.aux.end())
150 std::vector<float>
out(it->second.size() /
sizeof(
float));
151 std::memcpy(
out.data(), it->second.data(),
out.size() *
sizeof(
float));
155 [[nodiscard]] vk::Sampler default_sampler()
160 [[nodiscard]] Kinesis::Vision::FieldSample field_sample(
const float* raw)
163 .mean = { raw[0], raw[1], raw[2], raw[3] },
164 .min_val = { raw[4], raw[5], raw[6], raw[7] },
165 .max_val = { raw[8], raw[9], raw[10], raw[11] },
169 [[nodiscard]] std::array<uint32_t, 4> pixel_rect(
const Kinesis::Vision::BoundingBox& box, uint32_t w, uint32_t
h)
172 static_cast<uint32_t
>(box.x *
static_cast<float>(w)),
173 static_cast<uint32_t
>(box.y *
static_cast<float>(
h)),
174 std::max<uint32_t>(1U,
static_cast<uint32_t
>(box.w *
static_cast<float>(w))),
175 std::max<uint32_t>(1U,
static_cast<uint32_t
>(box.h *
static_cast<float>(
h))),
179 void add_line(std::vector<float>& segments, glm::vec2
a, glm::vec2
b, glm::vec4
color)
181 segments.insert(segments.end(), { a.x, a.y, b.x, b.y, color.r, color.g, color.b, color.a });
184 void add_rect(std::vector<float>& segments,
const Kinesis::Vision::BoundingBox& box, glm::vec4
color)
186 const glm::vec2 tl { box.x, box.y };
187 const glm::vec2 tr { box.x + box.w, box.y };
188 const glm::vec2 br { box.x + box.w, box.y + box.h };
189 const glm::vec2 bl { box.x, box.y + box.h };
190 add_line(segments, tl, tr,
color);
191 add_line(segments, tr, br,
color);
192 add_line(segments, br, bl,
color);
193 add_line(segments, bl, tl,
color);
196 void add_cross(std::vector<float>& segments, glm::vec2 p, glm::vec2 half, glm::vec4
color)
198 add_line(segments, { p.x - half.x, p.y }, { p.x + half.x, p.y },
color);
199 add_line(segments, { p.x, p.y - half.y }, { p.x, p.y + half.y },
color);
236 if (
label != 0 &&
label <= elements.boxes.size())
237 return static_cast<size_t>(
label - 1U);
270 const auto label = elements.boxes[*
index].label_id;
286 std::vector<glm::vec2>
points;
290 points.push_back(k.position);
294 points.push_back(t.position);
307 const auto count = elements.boxes.size();
309 std::vector<float> metrics;
310 metrics.reserve(
count * 6);
311 for (
size_t i = 0; i <
count; ++i) {
312 const auto& box = elements.boxes[i];
313 const uint32_t
label = box.label_id != 0 ? box.label_id :
static_cast<uint32_t
>(i + 1U);
314 const auto contour = std::ranges::find_if(elements.contours,
315 [
label](
const auto& c) { return c.label_id == label; });
317 metrics.push_back(contour != elements.contours.end() ? contour->area : box.w * box.h);
318 metrics.push_back(i < elements.shapes.size() ? elements.shapes[i].compactness : 0.0F);
319 metrics.push_back(i < elements.shapes.size() ? elements.shapes[i].aspect_ratio : 0.0F);
320 metrics.push_back(i < elements.nearest_neighbor_distance_px.size() ? elements.nearest_neighbor_distance_px[i] : 0.0F);
321 metrics.push_back(box.x + box.w * 0.5F);
322 metrics.push_back(box.y + box.h * 0.5F);
326 use_shader(ctx,
"element_select.comp.spv", k_wg_reduce,
sizeof(SelectPC));
327 ctx.set_binding_data(3, std::span<const float>(metrics));
328 ctx.set_push_constants(SelectPC {
329 .count =
static_cast<uint32_t
>(
count),
331 .largest = focus.
largest ? 1U : 0U,
336 const auto raw = dispatch_reduce(ctx, 256, 2 *
sizeof(
float));
340 const auto index =
static_cast<size_t>(std::lround(raw[0]));
341 return index < count ? std::optional<size_t> {
index } : std::nullopt;
345 const std::shared_ptr<Core::VKImage>& source,
354 use_shader(ctx,
"vision_crop.comp.spv", k_wg_image,
sizeof(CropPC));
355 ctx.stage_image(source);
356 ctx.set_push_constants(CropPC {
357 .src_x = region->x, .src_y = region->y, .src_w = region->w, .src_h = region->h,
363 const std::shared_ptr<Core::VKImage>& source,
371 const auto count =
static_cast<uint32_t
>(boxes.size());
373 std::vector<float> rects;
374 rects.reserve(boxes.size() * 4);
375 for (
const auto&
b : boxes)
376 rects.insert(rects.end(), {
b.x,
b.y,
b.w,
b.h });
379 use_shader(ctx,
"vision_region_atlas.comp.spv", k_wg_image,
sizeof(AtlasPC));
380 ctx.stage_image(source);
381 ctx.set_binding_data(3, std::span<const float>(rects));
382 ctx.set_push_constants(AtlasPC { .tile_w =
tile_w, .tile_h =
tile_h, .count =
count });
388 const std::shared_ptr<Core::VKImage>& source,
390 uint32_t w, uint32_t
h)
400 const std::shared_ptr<Core::VKImage>& source,
402 const std::vector<size_t>&
indices,
408 std::vector<float>
points;
409 std::vector<uint32_t> polygons;
410 std::vector<uint32_t> table;
411 table.reserve(
indices.size() * 6);
415 const uint32_t
label = box.label_id != 0 ? box.label_id :
static_cast<uint32_t
>(
index + 1U);
417 const auto first_polygon =
static_cast<uint32_t
>(polygons.size() / 2);
418 uint32_t polygon_count = 0;
419 for (
const auto& contour : elements.
contours) {
420 const bool owned = contour.label_id ==
label || (contour.label_id == 0 && contour.parent_label ==
label);
421 if (!owned || contour.points.size() < 3)
424 polygons.push_back(
static_cast<uint32_t
>(
points.size() / 2));
425 polygons.push_back(
static_cast<uint32_t
>(contour.points.size()));
426 for (
const auto& p : contour.points)
431 table.insert(table.end(), {
432 std::bit_cast<uint32_t>(box.x),
433 std::bit_cast<uint32_t>(box.y),
434 std::bit_cast<uint32_t>(box.w),
435 std::bit_cast<uint32_t>(box.h),
443 if (polygons.empty())
444 polygons.assign(2, 0U);
446 const auto count =
static_cast<uint32_t
>(
indices.size());
449 use_shader(ctx,
"contour_fill.comp.spv", k_wg_image,
sizeof(AtlasPC));
450 ctx.stage_image(source);
451 ctx.set_binding_data(3, std::span<const float>(
points));
452 ctx.set_binding_data(4, std::span<const uint32_t>(polygons));
453 ctx.set_binding_data(5, std::span<const uint32_t>(table));
454 ctx.set_push_constants(AtlasPC { .tile_w =
tile_w, .tile_h =
tile_h, .count =
count });
459 const std::shared_ptr<Core::VKImage>& source,
472 const std::shared_ptr<Core::VKImage>& source,
484 const std::shared_ptr<Core::VKImage>& source,
486 uint32_t source_w, uint32_t source_h)
492 const auto rect = pixel_rect(*region, source_w, source_h);
495 use_shader(ctx,
"region_sample_batch.comp.spv", k_wg_reduce,
sizeof(BatchCountPC));
496 ctx.stage_image(source);
497 ctx.set_binding_data(3, std::span<const uint32_t>(rect));
498 ctx.set_push_constants(BatchCountPC { .count = 1 });
500 const auto raw = dispatch_reduce(ctx, 256, 3 *
sizeof(glm::vec4));
505 const std::shared_ptr<Core::VKImage>& source,
507 uint32_t source_w, uint32_t source_h)
513 const auto count =
static_cast<uint32_t
>(boxes.size());
515 std::vector<uint32_t> rects;
516 rects.reserve(boxes.size() * 4);
517 for (
const auto&
b : boxes) {
518 const auto rect = pixel_rect(
b, source_w, source_h);
519 rects.insert(rects.end(), rect.begin(), rect.end());
523 use_shader(ctx,
"region_sample_batch.comp.spv", k_wg_reduce,
sizeof(BatchCountPC));
524 ctx.stage_image(source);
525 ctx.set_binding_data(3, std::span<const uint32_t>(rects));
526 ctx.set_push_constants(BatchCountPC { .count =
count });
528 const auto raw = dispatch_reduce(ctx, 256U *
count,
static_cast<size_t>(
count) * 3 *
sizeof(glm::vec4));
529 std::vector<Kinesis::Vision::FieldSample>
out;
530 if (raw.size() <
static_cast<size_t>(
count) * 12)
534 for (uint32_t i = 0; i <
count; ++i)
535 out.push_back(field_sample(raw.data() +
static_cast<size_t>(i) * 12));
540 const std::shared_ptr<Core::VKImage>& source,
543 uint32_t source_w, uint32_t source_h)
549 const auto count =
static_cast<uint32_t
>(
points.size());
552 use_shader(ctx,
"point_sample.comp.spv", k_wg_reduce,
sizeof(PointSamplePC));
553 ctx.stage_image(source);
554 ctx.set_binding_data(3, std::span<const float>(
555 reinterpret_cast<const float*
>(
points.data()),
points.size() * 2));
556 ctx.set_push_constants(PointSamplePC {
557 .count =
count, .radius =
radius, .src_w = source_w, .src_h = source_h });
559 const auto raw = dispatch_reduce(ctx,
count,
static_cast<size_t>(
count) *
sizeof(glm::vec4));
560 std::vector<glm::vec4>
out;
561 if (raw.size() <
static_cast<size_t>(
count) * 4)
565 std::memcpy(
out.data(), raw.data(),
static_cast<size_t>(
count) *
sizeof(glm::vec4));
570 const std::shared_ptr<Core::VKImage>& source,
573 uint32_t source_w, uint32_t source_h)
579 const auto count =
static_cast<uint32_t
>(centers.size());
582 use_shader(ctx,
"vision_patch_extract.comp.spv", k_wg_image,
sizeof(PatchExtractPC));
583 ctx.stage_image(source);
584 ctx.set_binding_data(3, std::span<const float>(
585 reinterpret_cast<const float*
>(centers.data()), centers.size() * 2));
586 ctx.set_push_constants(PatchExtractPC {
588 .src_w = source_w, .src_h = source_h });
593 const std::shared_ptr<Core::VKImage>& source,
596 uint32_t source_w, uint32_t source_h)
599 use_shader(ctx,
"histogram_equalize.comp.spv", k_wg_image,
sizeof(TonePC));
600 ctx.stage_image(source);
601 ctx.set_binding_data(3, std::span<const uint32_t>(appearance.
histogram.data(), appearance.
histogram.size()));
602 ctx.set_push_constants(TonePC { .mode =
mode, .gain =
gain, .width = source_w, .height = source_h });
603 return dispatch_image(ctx, source_w, source_h);
607 const std::shared_ptr<Core::VKImage>& source,
609 uint32_t source_w, uint32_t source_h)
618 const std::shared_ptr<Core::VKImage>& source,
621 uint32_t source_w, uint32_t source_h)
627 const float gain = mean > 1e-4F ? target / mean : 1.0F;
632 const std::shared_ptr<Core::VKImage>& source,
634 uint32_t source_w, uint32_t source_h)
640 use_shader(ctx,
"edge_select.comp.spv", k_wg_image,
sizeof(EdgePC));
641 ctx.stage_image(source);
642 ctx.stage_image_at(2, analysis.
detect_edges, ElementType::IMAGE_SAMPLED, default_sampler());
643 ctx.set_push_constants(EdgePC { .threshold =
m_edge_threshold, .out_w = source_w, .out_h = source_h });
644 return dispatch_image(ctx, source_w, source_h);
648 const std::shared_ptr<Core::VKImage>& source,
650 uint32_t source_w, uint32_t source_h)
658 use_shader(ctx,
"flow_mask.comp.spv", k_wg_image,
sizeof(MotionPC));
659 ctx.stage_image(source);
660 ctx.stage_image_at(2, analysis.
estimate_motion, ElementType::IMAGE_SAMPLED, default_sampler());
661 ctx.set_push_constants(MotionPC {
663 .out_w = source_w, .out_h = source_h });
664 return dispatch_image(ctx, source_w, source_h);
668 const std::shared_ptr<Core::VKImage>& source,
670 uint32_t source_w, uint32_t source_h)
672 const glm::vec4 yellow { 1.0F, 1.0F, 0.0F, 1.0F };
673 const glm::vec4 cyan { 0.0F, 1.0F, 1.0F, 1.0F };
674 const glm::vec4 green { 0.0F, 1.0F, 0.0F, 1.0F };
675 const glm::vec4 red { 1.0F, 0.0F, 0.0F, 1.0F };
676 const glm::vec4 white { 1.0F, 1.0F, 1.0F, 1.0F };
677 const glm::vec4 magenta { 1.0F, 0.0F, 1.0F, 1.0F };
679 const glm::vec2 half { 3.0F /
static_cast<float>(source_w), 3.0F /
static_cast<float>(source_h) };
681 std::vector<float> segments;
685 add_rect(segments, box, yellow);
686 for (
const auto& contour : analysis.
find_elements->contours) {
687 const auto n = contour.points.size();
688 for (
size_t i = 0; n > 1 && i < n; ++i)
689 add_line(segments, contour.points[i], contour.points[(i + 1) % n], cyan);
695 add_line(segments, t.previous, t.position, t.tracked ? green : red);
696 add_cross(segments, t.position, half, t.tracked ? green : red);
699 add_cross(segments, analysis.
track_objects->centroid, half, white);
704 add_cross(segments, k.position, half, magenta);
708 const auto count =
static_cast<uint32_t
>(segments.size() / 8);
709 if (segments.empty())
710 segments.assign(8, 0.0F);
713 use_shader(ctx,
"vision_crop.comp.spv", k_wg_image,
sizeof(CropPC));
714 ctx.stage_image(source);
715 ctx.set_binding_data(3, std::span<const float>(segments));
716 ctx.set_push_constants(CropPC {
717 .src_x = 0.0F, .src_y = 0.0F, .src_w = 1.0F, .src_h = 1.0F,
718 .out_w = source_w, .out_h = source_h });
719 auto image = dispatch_image(ctx, source_w, source_h);
725 .shader_path =
"overlay_segments.comp.spv",
726 .workgroup_size = k_wg_image,
727 .push_constant_size =
sizeof(SegmentsPC),
729 ctx.set_push_constants(SegmentsPC { .count =
count, .width = source_w, .height = source_h });
730 ctx.set_output_dimensions(source_w, source_h);
731 dispatch_and_wait(ctx);
733 return ctx.get_output_image(0);
737 const std::shared_ptr<Kakshya::SignalSourceContainer>& source,
743 if (
auto tc = std::dynamic_pointer_cast<Kakshya::TextureContainer>(source))
744 return tc->to_image(
static_cast<uint32_t
>(
index));
746 if (
auto wc = std::dynamic_pointer_cast<Kakshya::WindowContainer>(source))
747 return wc->to_image();
749 auto vc = std::dynamic_pointer_cast<Kakshya::VideoStreamContainer>(source);
753 const uint32_t w = vc->get_width();
754 const uint32_t
h = vc->get_height();
756 const void* raw = source->get_raw_data();
762 w,
h, Portal::Graphics::ImageFormat::RGBA8,
nullptr);
778 const auto supplied = Kakshya::get_metadata_value<std::shared_ptr<Core::VKImage>>(
input.metadata,
"vision_image");
779 const bool has_image = supplied && *supplied;
781 if (!has_image && !
input.data) {
782 output.
metadata[
"error"] = std::string(
"VisionExtractor: missing container");
786 const auto analysis = Kakshya::get_metadata_value<Kinesis::Vision::VisionAnalysis>(
787 input.metadata,
"vision_analysis");
789 output.
metadata[
"error"] = std::string(
"VisionExtractor: missing vision_analysis metadata");
793 const auto index = Kakshya::get_metadata_value<size_t>(
input.metadata,
"container_index").value_or(0);
796 output.
metadata[
"error"] = std::string(
"VisionExtractor: could not resolve image from container");
800 const uint32_t w =
image->get_width();
801 const uint32_t
h =
image->get_height();
842 output.
metadata[
"vision_extraction"] = *selected;
Core::GlobalInputConfig input
std::vector< uint32_t > indices
std::vector< glm::vec2 > * points
std::vector< float > * out
vk::PhysicalDeviceType type
vk::Sampler get_default_linear()
Get a default linear sampler.
static SamplerForge & instance()
void upload_data(const std::shared_ptr< Core::VKImage > &image, const void *data, size_t size)
Upload pixel data to an existing texture.
static TextureLoom & instance()
std::shared_ptr< Core::VKImage > create_2d(uint32_t width, uint32_t height, ImageFormat format=ImageFormat::RGBA8, const void *data=nullptr, uint32_t mip_levels=1)
Create a 2D texture.
GpuExecutionContext specialisation for image compute shaders.
std::shared_ptr< Core::VKImage > select_label(VisionGpuContexts &contexts, const std::shared_ptr< Core::VKImage > &source, uint32_t target_label, uint32_t w, uint32_t h)
Isolate one ConnectedComponents label's silhouette on GPU, through the explicit context set that prod...
Stateful GPU dispatch engine for VisionSequence execution.
OutputMode
Defines how the network's computational results are exposed.
MAYAFLUX_API ShaderFoundry & get_shader_foundry()
Get the global shader compiler instance.
Portal::Graphics::GpuBufferBinding GpuBufferBinding
VisionExtractMode
Which method run_operation dispatches to.
Per-channel mean/min/max over a sampled image region, computed on GPU (Yantra's region_sample....
std::vector< Contour > contours
std::vector< BoundingBox > boxes
FindElements output: contours plus the per-component BoundingBoxes ConnectedComponents already comput...
std::array< uint32_t, 256 > histogram
MeasureAppearance output: whole-frame photometric characterisation, computed on GPU (region_sample....
std::optional< FindElementsAnalysis > find_elements
std::optional< MeasureAppearanceAnalysis > measure_appearance
std::optional< DetectFeaturesAnalysis > detect_features
std::shared_ptr< Core::VKImage > estimate_motion
std::optional< TrackObjectsAnalysis > track_objects
std::shared_ptr< Core::VKImage > detect_edges
std::optional< float > motion_activity
VisionResult::motion_energy from an active EstimateMotion sequence.
Composed analysis output: one populated field per active intent that produced a result this call.
uint32_t set
Descriptor set index.
ElementType
Element type the shader expects in this binding.
Direction
Data flow direction for this binding.
std::unordered_map< std::string, std::any > metadata
Associated metadata.
Input/Output container for computation pipeline data flow with structure preservation.
ElementCriterion criterion
How the extractor chooses among several found elements: rank by criterion, highest first when largest...