GpuComputeConfig for a given VisionOp and its parameters.
Assembled ops (Threshold, NormalizeInplace, NormalizeRange, RgbaToGray, GrayToRgba) produce a config via ShaderSpec::Assemble + config_from_spec with no .comp file. All other implemented ops reference a .comp path under Portal/Shaders/Vision/. Unimplemented ops return a config with shader_id == INVALID_SHADER.
Stateless; does not depend on or affect m_contexts.
208{
209 switch (op) {
210 case VisionOp::Threshold: {
211 if (const auto* p = std::get_if<ThresholdParams>(¶ms); p && p->channels != ChannelMask::NONE)
212 return { .shader_path =
"threshold_bands.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(ThresholdBandsPC) };
213
214 const auto spec = ShaderSpec::Assemble {}
215 .storage_image("out", BindingDirection::Output)
216 .storage_image("src", BindingDirection::Input)
217 .pc("threshold")
218 .op(KernelOp::CompareGE)
220 .build();
222 }
223 case VisionOp::RgbaToGray: {
224 const auto spec = ShaderSpec::Assemble {}
225 .storage_image("out", BindingDirection::Output)
226 .storage_image("src", BindingDirection::Input)
227 .pc("wr")
228 .pc("wg")
229 .pc("wb")
230 .pc("wa")
231 .op(KernelOp::ChannelDot)
233 .build();
235 }
236 case VisionOp::GrayToRgba: {
237 const auto spec = ShaderSpec::Assemble {}
238 .storage_image("out", BindingDirection::Output)
239 .storage_image("src", BindingDirection::Input)
240 .op(KernelOp::ChannelReplicate)
242 .build();
244 }
245 case VisionOp::GaussianBlur: {
246 const auto spec = ShaderSpec::Assemble {}
247 .tmpl(KernelTemplate::Convolve2D)
248 .storage_image("out", BindingDirection::Output)
249 .storage_image("src", BindingDirection::Input)
255 .build();
257 }
258 case VisionOp::NormalizeRange: {
259 const auto spec = ShaderSpec::Assemble {}
260 .storage_image("out", BindingDirection::Output)
261 .storage_image("src", BindingDirection::Input)
262 .pc("scale")
263 .pc("offset")
264 .op(KernelOp::ScaleOffset)
266 .build();
268 }
269 case VisionOp::NormalizeInplace: {
270 const auto spec = ShaderSpec::Assemble {}
271 .storage_image("out", BindingDirection::Output)
272 .storage_image("src", BindingDirection::Input)
273 .op(KernelOp::Scale)
275 .build();
277 }
278 case VisionOp::Canny: {
279 const auto spec = ShaderSpec::Assemble {}
280 .storage_image("out", BindingDirection::Output)
281 .storage_image("src", BindingDirection::Input)
282 .pc("threshold")
283 .pc("value")
284 .op(KernelOp::CompareGEPreserve)
286 .build();
288 }
289 case VisionOp::RgbaToHsv:
290 return { .shader_path =
"rgba_to_hsv.comp.spv", .workgroup_size =
k_wg2d };
291 case VisionOp::Downsample2x:
292 return { .shader_path =
"downsample_2x.comp.spv", .workgroup_size =
k_wg2d };
293 case VisionOp::FilterSeparable:
294 return { .shader_path =
"filter_separable.comp.spv", .workgroup_size =
k_wg2d };
295 case VisionOp::Sobel:
296 return { .shader_path =
"sobel.comp.spv", .workgroup_size =
k_wg2d };
297 case VisionOp::Scharr:
298 return { .shader_path =
"scharr.comp.spv", .workgroup_size =
k_wg2d };
299 case VisionOp::Erode:
300 return { .shader_path =
"erode.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(MorphPC) };
301 case VisionOp::Dilate:
302 return { .shader_path =
"dilate.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(MorphPC) };
303 case VisionOp::Open:
304 return { .shader_path =
"open.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(MorphPC) };
305 case VisionOp::Close:
306 return { .shader_path =
"close.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(MorphPC) };
307 case VisionOp::MorphGradient:
308 return { .shader_path =
"morph_gradient.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(MorphPC) };
309 case VisionOp::HarrisResponse:
310 return { .shader_path =
"harris_response.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(HarrisPC) };
311 case VisionOp::ExtractPeaks:
312 return { .shader_path =
"extract_peaks.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(ExtractPeaksPC) };
313 case VisionOp::ConnectedComponents:
314 return { .shader_path =
"cc_colorize.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(uint32_t) * 2 };
315 case VisionOp::FindContours:
316 return { .shader_path =
"contour_segments.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(ContourSegmentsPC) };
317 case VisionOp::ThresholdAdaptive:
318 return { .shader_path =
"threshold_adaptive.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(ThresholdAdaptivePC) };
319 case VisionOp::ThresholdOtsu:
320 return { .shader_path = "threshold_otsu.comp.spv", .workgroup_size = { 256, 1, 1 } };
321 case VisionOp::TrackKeypoints:
322 return { .shader_path = "flow_lk.comp.spv", .workgroup_size = { 64, 1, 1 }, .push_constant_size = sizeof(FlowLkPC) };
323 case VisionOp::OpticalFlowDense:
324 return { .shader_path =
"flow_dense.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(FlowDensePC) };
325 case VisionOp::Confine:
326 return { .shader_path =
"vision_confine.comp.spv", .workgroup_size =
k_wg2d, .push_constant_size =
sizeof(ConfinePC) };
327 default:
329 }
330}
@ FLOAT32
32-bit float (standard GPU)
constexpr ShaderID INVALID_SHADER
constexpr std::array< uint32_t, 3 > k_wg2d
Standard 2D workgroup used by all pixel-to-pixel vision shaders.
Portal::Graphics::GpuComputeConfig GpuComputeConfig
GpuComputeConfig config_from_spec(const Portal::Graphics::ShaderSpec &spec)
Derive a GpuComputeConfig from a ShaderSpec.