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.
1229{
1230 switch (op) {
1231 case VisionOp::Threshold: {
1232 const auto spec = ShaderSpec::Assemble {}
1233 .storage_image("out", BindingDirection::Output)
1234 .storage_image("src", BindingDirection::Input)
1235 .pc("threshold")
1236 .op(KernelOp::CompareGE)
1237 .workgroup(k_wg2d[0], k_wg2d[1])
1238 .build();
1240 }
1241 case VisionOp::RgbaToGray: {
1242 const auto spec = ShaderSpec::Assemble {}
1243 .storage_image("out", BindingDirection::Output)
1244 .storage_image("src", BindingDirection::Input)
1245 .pc("wr")
1246 .pc("wg")
1247 .pc("wb")
1248 .pc("wa")
1249 .op(KernelOp::ChannelDot)
1250 .workgroup(k_wg2d[0], k_wg2d[1])
1251 .build();
1253 }
1254 case VisionOp::GrayToRgba: {
1255 const auto spec = ShaderSpec::Assemble {}
1256 .storage_image("out", BindingDirection::Output)
1257 .storage_image("src", BindingDirection::Input)
1258 .op(KernelOp::ChannelReplicate)
1259 .workgroup(k_wg2d[0], k_wg2d[1])
1260 .build();
1262 }
1263 case VisionOp::GaussianBlur: {
1264 const auto spec = ShaderSpec::Assemble {}
1265 .tmpl(KernelTemplate::Convolve2D)
1266 .storage_image("out", BindingDirection::Output)
1267 .storage_image("src", BindingDirection::Input)
1272 .workgroup(k_wg2d[0], k_wg2d[1])
1273 .build();
1275 }
1276 case VisionOp::NormalizeRange: {
1277 const auto spec = ShaderSpec::Assemble {}
1278 .storage_image("out", BindingDirection::Output)
1279 .storage_image("src", BindingDirection::Input)
1280 .pc("scale")
1281 .pc("offset")
1282 .op(KernelOp::ScaleOffset)
1283 .workgroup(k_wg2d[0], k_wg2d[1])
1284 .build();
1286 }
1287 case VisionOp::NormalizeInplace: {
1288 const auto spec = ShaderSpec::Assemble {}
1289 .storage_image("out", BindingDirection::Output)
1290 .storage_image("src", BindingDirection::Input)
1291 .op(KernelOp::Scale)
1292 .workgroup(k_wg2d[0], k_wg2d[1])
1293 .build();
1295 }
1296 case VisionOp::Canny: {
1297 const auto spec = ShaderSpec::Assemble {}
1298 .storage_image("out", BindingDirection::Output)
1299 .storage_image("src", BindingDirection::Input)
1300 .pc("threshold")
1301 .pc("value")
1302 .op(KernelOp::CompareGEPreserve)
1303 .workgroup(k_wg2d[0], k_wg2d[1])
1304 .build();
1306 }
1307 case VisionOp::RgbaToHsv:
1308 return { .shader_path = "rgba_to_hsv.comp.spv", .workgroup_size = k_wg2d };
1309 case VisionOp::Downsample2x:
1310 return { .shader_path = "downsample_2x.comp.spv", .workgroup_size = k_wg2d };
1311 case VisionOp::FilterSeparable:
1312 return { .shader_path = "filter_separable.comp.spv", .workgroup_size = k_wg2d };
1313 case VisionOp::Sobel:
1314 return { .shader_path = "sobel.comp.spv", .workgroup_size = k_wg2d };
1315 case VisionOp::Scharr:
1316 return { .shader_path = "scharr.comp.spv", .workgroup_size = k_wg2d };
1317 case VisionOp::Erode:
1318 return { .shader_path = "erode.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(MorphPC) };
1319 case VisionOp::Dilate:
1320 return { .shader_path = "dilate.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(MorphPC) };
1321 case VisionOp::Open:
1322 return { .shader_path = "open.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(MorphPC) };
1323 case VisionOp::Close:
1324 return { .shader_path = "close.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(MorphPC) };
1325 case VisionOp::MorphGradient:
1326 return { .shader_path = "morph_gradient.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(MorphPC) };
1327 case VisionOp::HarrisResponse:
1328 return { .shader_path = "harris_response.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(HarrisPC) };
1329 case VisionOp::ExtractPeaks:
1330 return { .shader_path = "extract_peaks.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(ExtractPeaksPC) };
1331 case VisionOp::ConnectedComponents:
1332 return { .shader_path = "cc_colorize.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(uint32_t) * 2 };
1333 case VisionOp::FindContours:
1334 return { .shader_path = "contour_segments.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(ContourSegmentsPC) };
1335 case VisionOp::ThresholdAdaptive:
1336 return { .shader_path = "threshold_adaptive.comp.spv", .workgroup_size = k_wg2d, .push_constant_size = sizeof(ThresholdAdaptivePC) };
1337 case VisionOp::ThresholdOtsu:
1338 return { .shader_path = "threshold_otsu.comp.spv", .workgroup_size = { 256, 1, 1 } };
1339 default:
1341 }
1342}
constexpr ShaderID INVALID_SHADER
Portal::Graphics::GpuComputeConfig GpuComputeConfig
GpuComputeConfig config_from_spec(const Portal::Graphics::ShaderSpec &spec)
Derive a GpuComputeConfig from a ShaderSpec.