Extract and rank keypoints from the current response image.
Produces structured host keypoints unless TrackKeypoints follows immediately. In that sequence, tracking extracts detections directly into its GPU buffers.
89{
90 auto& structured_ctx = contexts.structured;
91 auto w = contexts.pass.w;
92 auto h = contexts.pass.h;
94
95 constexpr uint32_t k_max_kp = 4096;
96
97 const auto* tracker = contexts.pass.ahead();
98 if (tracker && tracker->op == VisionOp::TrackKeypoints) {
99 contexts.pass.result.structured = std::monostate {};
100 contexts.pass.result.w = w;
101 contexts.pass.result.h =
h;
102 return;
103 }
104
105 structured_ctx.swap_shader({
106 .shader_path = "extract_peaks.comp.spv",
107 .workgroup_size = { 8, 8, 1 },
108 .push_constant_size = sizeof(ExtractPeaksPC),
109 });
110
111 structured_ctx.set_output_size(1, sizeof(uint32_t));
112 structured_ctx.set_output_size(2, static_cast<size_t>(k_max_kp) * 4 * sizeof(float));
113 structured_ctx.ensure_shared_buffer(0, 3, (static_cast<size_t>(k_max_kp) + 1U) * 4U,
115
116 structured_ctx.stage_image(contexts.pass.current);
117 structured_ctx.set_push_constants(ExtractPeaksPC {
118 .threshold = p.threshold,
119 .nms_radius = p.nms_radius,
120 .width = w,
122 .max_keypoints = k_max_kp,
124 });
125
126 structured_ctx.set_output_dimensions(w,
h);
127 const auto fence = structured_ctx.dispatch_async({});
128 structured_ctx.clear_output_dimensions();
129 foundry.wait_for_fence(
fence);
130 foundry.release_fence(
fence);
131
132 const auto gpu_result = structured_ctx.collect_result();
133
135 if (auto it = gpu_result.aux.find(1); it != gpu_result.aux.end())
136 std::memcpy(&
count, it->second.data(),
sizeof(uint32_t));
138
139 if (p.export_keypoints) {
140 const glm::vec4 header { std::bit_cast<float>(
count), 0.0F, 0.0F, 0.0F };
141 structured_ctx.upload_shared_raw(0, 3, reinterpret_cast<const uint8_t*>(&header), sizeof(glm::vec4));
142 contexts.pass.result.buffers.keypoints = std::make_shared<Portal::Graphics::GpuBufferHandle>(
143 structured_ctx.shared_buffer_handle(0, 3));
144 }
145
146 struct GpuKp {
147 float x, y, response, pad;
148 };
149 std::vector<GpuKp> raw(
count);
151 if (auto it = gpu_result.aux.find(2); it != gpu_result.aux.end())
152 std::memcpy(raw.data(), it->second.data(),
count *
sizeof(GpuKp));
153 }
154
155 std::vector<Kinesis::Vision::Keypoint> kpts;
157 for (const auto& kp : raw) {
158 kpts.push_back({ .position = { kp.x, kp.y },
159 .response = kp.response,
160 .scale = 1.0F,
161 .angle = 0.0F });
162 }
163 std::ranges::sort(kpts, [](
const auto&
a,
const auto&
b) {
return a.response >
b.response; });
164
165 contexts.pass.result.structured = std::move(kpts);
166 contexts.pass.result.w = 0;
167 contexts.pass.result.h = 0;
168}
MAYAFLUX_API ShaderFoundry & get_shader_foundry()
Get the global shader compiler instance.
@ COMPUTE
Storage buffer for compute shaders (device-local)
PixelRect region_rect(const std::optional< Kinesis::Vision::BoundingBox > ®ion, uint32_t w, uint32_t h)
Pixel rectangle of a normalised region over a w x h image, clamped to the image.
enum MayaFlux::Portal::Graphics::GpuBufferBinding::ElementType FLOAT32