65.9K
CodeProject 正在变化。 阅读更多。
Home

Intel® 计算机视觉 SDK 开发人员指南

emptyStarIconemptyStarIconemptyStarIconemptyStarIconemptyStarIcon

0/5 (0投票)

2017年10月24日

CPOL

64分钟阅读

viewsIcon

19414

Intel® 计算机视觉 SDK 是一款用于针对 Intel 片上系统 (SoCs) 开发和优化计算机视觉及图像处理管线的全新软件开发套件。

Intel® 计算机视觉 SDK 是一款用于针对 Intel 片上系统 (SoCs) 开发和优化计算机视觉及图像处理管线的全新软件开发套件。

该产品包含六个主要组件

  1. Intel 优化实现的 Khronos* OpenVX* 1.1 API,并通过附加 API(请参阅 Intel 对 OpenVX* API 的扩展:OpenCL™ 自定义内核Intel 对 OpenVX* API 的扩展:高级分块 章)和内核进行了扩展,包括对卷积神经网络 (CNN) 的支持(请参阅 Intel 对 OpenVX* 原始操作的支持扩展 章)。
  2. 预编译且完全验证的社区 OpenCV* 3.3 二进制文件,增加了人脸识别功能,更快速、更准确,详情请参阅 Intel OpenCV 中的高级人脸识别功能 章。
  3. Vision Algorithm Designer (VAD) 是一款跨平台 IDE 工具,用于 OpenVX* 图的可视化开发和分析。如需了解更多详情,请参阅 VAD 开发者指南
  4. 深度学习模型优化器工具,可帮助您部署卷积神经网络。如需了解更多详情,请参阅 深度学习模型优化器开发者指南
  5. 深度学习推理引擎,具有统一的(OpenVX-无关)API,可将推理集成到应用程序逻辑中。如需了解更多详情,请参阅 深度学习推理引擎开发者指南
  6. 示例应用程序(可作为单独的软件包获取)。如需了解更多详情,请参阅 Intel(R) 计算机视觉 SDK 示例应用程序 主题。

该产品与 Intel® VTune™ 兼容(请参阅 Intel(R) VTune(TM) 工具支持 部分)。

使用 SDK 的一般开发流程是将低级原语组合起来构建高度优化的计算机视觉管线,这些管线可以独立使用或与其他代码集成。

支持的平台

Intel® 计算机视觉 SDK 在以下目标和开发环境中获得官方支持和验证

开发/目标 操作系统 平台 操作系统
64 位/32 位
开发和目标 Ubuntu* 16.04, CentOS* 7.2, Microsoft Windows* 10 第六代 Intel® Core™ 处理器 64 位
目标 Yocto* MR3 下一代 Intel® Atom™ 处理器(前身为 Apollo Lake) 64 位

Intel 的 OpenVX* 实现:主要功能

性能

  • Intel 的 OpenVX* 实现提供多线程(使用 Intel® Threading Building Blocks)和矢量化(使用 Intel® Integrated Performance Primitives)的 CPU 内核。此外,GPU 支持由优化的 OpenCL 实现提供。
  • 为了对本地值执行大多数读/写操作,该实现支持输入、中间和输出数据的自动分块。

可扩展性

异构性

  • 支持任务并行和数据并行,以最大化利用 CPU、GPU 和 IPU 等计算资源。
  • 通用的系统级设备亲和性,以及用于通过目标概念协调单个节点的精细 API。有关详细信息,请参阅 使用 Intel® 计算机视觉 SDK 进行异构计算 章。

注意: 请注意,IPU 支持处于实验阶段。有关详细信息,请参阅 关于 IPU 支持的一般说明 部分。

Intel 的 OpenVX 实现使用模块化方法,具有通用运行时和独立的设备特定插件。

Intel® CV SDK OpenVX* Architecture

此版本有什么新内容?

  • CPU 的 CNN 内核实现现在由 Intel® Math Kernel Library (Intel® MKL) 提供支持(针对 32 位浮点路径)。有关更多详细信息,请参阅 系统级设备亲和性 部分。
  • OpenCV 功能,例如新的性能分析机制和新的“dnn”模块,将在 OpenCV 发行日志中进行描述。
  • 工具方面也有重大改进,包括 Vision Algorithm Designer (VAD) 的多项改进。有关更多详细信息,请参阅 VAD 开发者指南
  • 模型优化器工具进一步简化了卷积神经网络的部署。模型优化器引入了对 TensorFlow* 框架的 Alpha 支持,以实现 TensorFlow 模型和数据到 OpenVX 图和格式的转换。还可以注册您自己的实现以处理尚未被 OpenVX 支持的层。 Intel 对 OpenVX* 原始操作的支持扩展 部分包含有关 CNN 支持的更多参考。

入门

有关 Intel® 计算机视觉 SDK 的安装和环境设置的详细信息,请参阅 快速入门指南快速入门指南 还提供了有关使用 OpenVX* 构建应用程序的更多详细信息,以及一些 Intel OpenCV 集成技巧。

OpenVX* 基础知识

OpenVX* 是 Khronos* 的一项新标准,提供了一组优化的原语,用于低级图像处理和计算机视觉原语。OpenVX 规范提供了一种编写可在多个供应商和平台以及多种硬件类型之间移植的代码的方法。

OpenVX 资源和执行抽象使硬件供应商能够专注于特定平台进行实现优化。同时,标准化允许软件开发人员在无需立即了解针对特定硬件进行优化的所有细节的情况下,就开始基于高性能供应商实现进行算法创新。

基于图的编程接口旨在创建使用图概念的视觉算法。该标准定义了图形式和执行语义,以满足计算机视觉开发人员的需求。图接口的主要优点之一是底层实现能够积极优化完整的视觉管线,而不是特定的函数。

跨 CPU、GPU 和固定功能硬件等各种引擎的性能可移植性是 OpenVX 规范的设计目标之一。

有关更多详细信息,请参阅以下主题

OpenVX* 编程框架

OpenVX* 定义了一个 C 应用程序编程接口 (API),用于构建、验证和协调图执行,以及访问内存对象。图抽象使 OpenVX 实现能够为底层加速架构优化执行。

图定义了一个显式数据流,使用节点对象数据对象(如图像和数组)来描述图的数据流。创建图后,需要进行图验证。之后,同一个图可以多次执行并更新输入。

OpenVX 还定义了vxu 实用库,使您可以将每个 OpenVX 预定义函数作为直接可调用的 C 函数使用,而无需先创建图。使用 vxu 库构建的应用程序无法获得图所带来的优化好处,但 vxu 库可以作为使用 OpenVX 的最简单方式,以及移植现有视觉应用程序的第一步。

随着计算机视觉领域仍在快速发展,OpenVX 提供了一个可扩展机制,使开发人员可以定义函数并将其添加到图中。

OpenVX* 对象

本章简要介绍了理解基本 OpenVX* 应用程序开发所需的一些重要 OpenVX* 对象。有关更多详细信息,请参阅 OpenVX* 1.1 规范。

OpenVX 上下文是 OpenVX* 对象的容器。您必须创建一个上下文才能创建对象。程序可能拥有多个上下文,原因各不相同。但是,跨上下文共享对象非常有限。

OpenVX 内核是计算机视觉函数的抽象表示,例如“Sobel 梯度”或“Lucas Kanade 特征跟踪”。内核通过限定名称(例如 "org.khronos.openvxsobel_3x3")或枚举来引用,其功能由规范和符合性测试定义。

OpenVX 节点是内核的实例。节点封装了内核(功能)以及该功能在图数据流中的位置。后者通过分配节点参数来定义,例如,一个节点的输出成为另一个节点的输入。每个创建的节点都必须属于一个图。

OpenVX 图是节点及其连接的集合。图必须是定向的(仅单向)且无环的(不回环)。OpenVX* 支持两种图执行模式:同步/阻塞模式(使用 vxProcessGraph)和异步模式(通过 vxScheduleGraphvxWaitGraph)。在这两种模式下,单个图内节点的执行顺序由数据流定义,这是保证的。但是,在异步模式下,图执行没有定义顺序(每个都通过 vxScheduleGraph 调用指定)。要控制执行顺序,您必须相应地插入 vxWaitGraph

OpenVX* 中的图依赖数据对象来连接节点。OpenVX 图像是一个不透明的数据对象,包含像素。有多种类型的图像。例如,为了启用某些优化,中间图像可以声明为虚拟图像。此类图像无法在图外部进行读写访问,因为没有关联的可访问内存。基本上,虚拟图像是定义 OpenVX 节点依赖性的存根(就像 C/C++ 中的临时变量一样),并且这些可以被 OpenVX 图编译器消除。

创建简单的 OpenVX* 图应用程序

本节介绍了一个图设计和执行的简单示例。图基于的应用程序使用 Intel® 计算机视觉 SDK 中的标准 OpenVX* 内核,从模糊的输入图像计算梯度幅度和梯度相位。下图以 DAG 的形式描述了该算法。

Example OpenVX* Graph

在基于 OpenVX* 的应用程序中,开发人员需要在应用程序开始时创建上下文和图。

//The code to create context and graph 
vx_context context = vxCreateContext();
vx_graph graph = vxCreateGraph(context);

图像、虚拟图像、数组和其他数据对象是连接图节点以及将数据传入和传出图的手段。

//The code to create input and output image data objects
vx_image images[] = { 
  vxCreateImage(context, 640, 480,VX_DF_IMAGE_UYVY), 
  vxCreateImage(context, 640, 480,VX_DF_IMAGE_U8),
  vxCreateImage(context, 640, 480,VX_DF_IMAGE_U8),
};
//The code to create virtual images that connect nodes
vx_image virts[] = { 
  vxCreateVirtualImage(graph, 0, 0,VX_DF_IMAGE_VIRT), 
  vxCreateVirtualImage(graph, 0, 0,VX_DF_IMAGE_VIRT),
  vxCreateVirtualImage(graph, 0, 0,VX_DF_IMAGE_VIRT),
  vxCreateVirtualImage(graph, 0, 0,VX_DF_IMAGE_VIRT)
};

请注意,对于中间(虚拟)图像,您无需指定类型或大小,因为图编译器能够自动推断这些数据。正如我们之前讨论过的,这些实际上是没有任何关联的可访问内存的存根。

还有其他特殊类型的图像,例如,存储所有像素中单个统一值的统一图像。您还可以创建一个引用已分配内存的图像对象(请参阅 vxCreateImageFromHandle),与其他 API 的互操作性 部分。

图的数据流是通过在图中添加节点并将数据对象作为输入和输出提供给这些节点来创建的。在上图所示的图中,示例 UYVY 图像是非虚拟输入图像。它作为参数传递给提取亮度信息的节点 vxChannelExtractNode,然后连接到模糊节点 vxGaussian3x3Node。模糊节点连接到计算梯度的 vxSobel3x3Node 节点。进而,梯度节点连接到同时生成幅度和相位的 vxMagnitudeNodevxPhaseNode 节点。

请注意,两个输出节点都生成非虚拟图像,否则这些图像将无法被主机代码访问。

//Create nodes and connect them in the graph
vxChannelExtractNode(graph, images[0],VX_CHANNEL_Y, virts[0]);
vxGaussian3x3Node(graph, virts[0], virts[1];
vxSobel3x3Node (graph, virts[1], virts[2], virts[3]);
vxMagnitudeNode(graph, virts[2], virts[3], images[1]);
vxPhaseNode (graph, virts[2], virts[3], images[2]);

OpenVX* 图在执行之前必须经过验证过程。

//Verify the graph
status = vxVerifyGraph (graph);
//Finally let’s execute the graph and wait for completion:
//Execute the graph synchronously       
vxProcessGraph(graph);

最后一步是释放与图、上下文、图像等相关的资源。

vxReleaseGraph(&graph);
...
vxReleaseContext(&context);

有关包含错误处理的更完整示例,请参阅 Intel(R) 计算机视觉 SDK 示例应用程序 部分所述的示例。

请注意,要编译代码,您需要包含 OpenVX 的头文件。

#include <VX/vx.h>

您还应该将代码链接到 OpenVX 库。有关必要的环境设置的详细信息,请参阅快速入门指南

将用户代码包装为 OpenVX* 用户内核

OpenVX* 允许您将自己的算法添加到处理管线中。

要成为真正的 OpenVX 内核,您的用户代码必须实现某些 API,例如验证函数(在 vxVerifyGraph 期间调用,有关代码示例,请参阅 Intel(R) 计算机视觉 SDK 示例应用程序 部分)。其他一些回调函数,如内核初始化/反初始化,是可选的。所有这些回调函数以及内核函数本身都提供给 vxAddKernel

VX_API vx_kernel vxAddUserKernel(vx_context context,
                             const vx_char name[VX_MAX_KERNEL_NAME],
                             vx_enum enumeration,
                             vx_kernel_f func_ptr,
                             vx_uint32 numParams,
                             vx_kernel_validate_f validate,
                             vx_kernel_initialize_f init,
                             vx_kernel_deinitialize_f deinit);

然后,内核接受的每个参数都通过 vxAddParameterToKernel 指定,例如

vx_kernel kernel = vxAddKernel(context,…);
vxAddParameterToKernel(kernel, 0, VX_INPUT,  VX_TYPE_IMAGE…);
vxAddParameterToKernel(kernel, 1, VX_OUTPUT, VX_TYPE_IMAGE…);

可以通过 vxSetKernelAttribute 来调整内核的某些运行时行为。

最后,通过调用 vxFinalizeKernel 完成内核发布。

有关用户内核实现细节的完整示例,请参阅视频稳定 SDK 示例(<SDK_SAMPLES_ROOT>/samples),以及 Intel(R) 计算机视觉 SDK 示例应用程序 部分。

从用户内核库添加用户内核

为了实现用户内核在应用程序之间的某种重用,通常将内核分组到动态加载的库中。以下是相关 OpenVX API 的一个子集。

名称 描述
vxLoadKernels 将指定的内核模块加载到 OpenVX* 上下文中。加载文件的格式可能是动态链接库,但这取决于具体实现。
vxPublishKernels 由 vxLoadKernels 加载的模块应实现的入口点。同样,在卸载内核库时会调用 vxUnpublishKernels。
vxGetKernelByEnum 使用枚举获取内核的引用。
vxGetKernelByName 通过名称获取内核的引用。
vxQueryKernel 获取有关参数数量、枚举值等的信息。
vxCreateGenericNode 为给定的内核创建节点对象的引用。
vxSetParameterByIndex 设置参数以完成节点。

函数 vxLoadKernels 将模块加载到给定的上下文中,并调用 vxPublishKernels 来发布已添加内核及其参数的信息列表。

在内核发布后,您可以使用 vxGetKernelByEnumvxGetKernelByName 来查询内核,然后为内核对象创建一个通用节点。最后,必须使用 vxSetParameterByIndex 来设置参数,同时也可以查询每个内核的参数列表。

//Loads the "libmy_ovx" library with user-kernel for division 
vx_status s = vxLoadKernels(ctx, "libmy_ovx") ;
//Get the kernel, by name
vx_kernel kernel = vxGetKernelByName(ctx, "my_div");
//query the kernel parameters
vxQueryKernel(kernel,VX_KERNEL_PARAMETERS, &np, sizeof(np));
printf("Number of Kernel parameters is: %d\n",np);
for(size_t j = 0; j < np; ++j)
{
    vx_parameter p = vxGetKernelParameterByIndex(kernel, j);
    vx_type_e t;
    vxQueryParameter(p,VX_PARAMETER_TYPE, &t, sizeof(t));
    assert(VX_TYPE_DF_IMAGE==t); //the kernel accepts just images 
}
vx_node div_node = vxCreateGenericNode(graph, kernel);
//Set the arguments to the node, connecting the node to the graph
vx_image img_in0 = vxCreateImage(ctx, 640, 480,VX_DF_IMAGE_U8);
vx_image img_in1 = vxCreateImage(ctx, 640, 480,VX_DF_IMAGE_U8);
vx_image img_out = vxCreateImage(ctx, 640, 480,VX_DF_IMAGE_U8);
vxSetParameterByIndex(div_node, 0, (vx_reference)img_in0);
vxSetParameterByIndex(div_node, 1, (vx_reference)img_in1);
vxSetParameterByIndex(div_node, 2, (vx_reference)img_out);

或者,当用户库头文件可用时,您可以使用常规函数声明创建节点。

//in my_div.h:
vx_node vxDivNode(vx_graph g,vx_image i0,vx_image i1,vx_image out);
//in main.cpp:
#include "my_div.h"
...
vxDivNode(graph, in0, in1, out);

关于用户内核的一般建议

将遗留代码与 OpenVX* 结合的最直接方法是执行 OpenVX 图,等待完成,在主机上访问其数据,然后执行用户代码处理数据。然后可以继续 OpenVX 执行,依此类推。有关 OpenCV 互操作性示例,请参阅以下章节。

有时将自己的代码包含为图节点是有意义的。以下是一些用户内核技巧。

  • 与其他图中的内核一样,同一个用户定义的内核可能有多个节点实例。这些实例可能并行执行。因此,您的用户内核代码应是线程安全的。
  • 用户内核可以充当隐式内存屏障。在大多数情况下,OpenVX 实现应完成所有先前计算并将更新的对象内存,而这些对象是用户内核所依赖的。这并不完全有利于性能。因此,如果您在热点中看到用户内核( Intel(R) VTune(TM) 工具支持 章),请考虑以下几点。
    • 如果内核足够重,请利用用户内核内部的并行性,例如通过并行化代码。强烈建议使用 Intel® Threading Building Blocks 作为并行性的首选 API,因为它将保证与 OpenVX 运行时其余部分的组合性。具体来说,在用户内核中使用 TBB 有助于避免线程过度订阅。

    • 考虑高级分块扩展(Intel 对 OpenVX* API 的扩展:高级分块 章)。这样,您就可以指定用户内核的分块版本,并让运行时以最佳方式调度事物,包括并行性。

    • 合并内核以处理本地数据(内核融合)通常比任何自动方法都能提供更好的性能。

最后,如果您的遗留代码是 OpenCL™,那么将 OpenCL API 的调用包装成用户内核相对简单。与其他类型的 OpenVX 用户代码一样,请注意当多个用户内核实例可能并行存在时可能出现的问题。例如,您的代码需要为每个用户节点保留单独的 cl_kernel 实例。这可以避免因多个线程对同一个 cl_kernel 设置不同参数而产生的冲突。此外,使用 OpenCL 的用户节点的开销可能会更加明显,因为涉及与 GPU 的通信。最终,这种方法缺乏数据共享,因为 OpenCL 上下文和队列在 OpenVX 运行时和您的代码之间不共享。

总之,与其使用用户内核包装 OpenCL 调用,不如考虑使用自定义 OpenCL 内核(Intel 对 OpenVX* API 的扩展:OpenCL™ 自定义内核 部分)来利用现有的 OpenCL 内核。

追求性能

Intel OpenVX* 性能承诺

OpenVX* 方法将传统的单次函数加速扩展为包含图的概念。图暴露了传统方法可能无法提供或不明显地暴露的优化可能性。例如,在 Intel OpenVX 实现中,共享数据的不同内核不必使用全局(慢)内存。相反,自动分块将数据放入缓存。类似地,虽然并行性并未直接在图中表达,但独立的 数据流从其结构中提取出来。

OpenVX 还创建了一个逻辑模型,其中 Intel SoC 的 IP 块完全共享系统资源(如内存);代码可以无缝地调度到最适合执行它的块上。

除了全局图级优化之外,OpenVX 视觉函数的性能还通过使用针对特定平台高度优化的实现来解决。对于 CPU,这通过 Intel® Integrated Performance Primitives (Intel® IPP) 来实现,该库具有针对不同架构的代码分支。对于 GPU,则使用成熟的 OpenCL* 即时编译到特定架构的堆栈。

为了达到良好的性能,必须充分理解 OpenVX 计算模型中的隐含权衡。本节描述了 OpenVX 在性能方面的通用注意事项。

使用 OpenVX* 图模式估算性能

OpenVX 支持一种称为即时模式的单函数执行模型。

注意: 请注意,使用视觉函数的即时模式(前缀为 vxu*,例如 vxuColorConvert)实际上仍然意味着在后台使用图(每个图仅包含一个函数)。因此,图验证和其他成本(如内存分配)将包含在计时中,并且不会分摊到多个节点/迭代。

即时模式仍然可以作为中间步骤,例如,在将应用程序从 OpenCV 移植到 OpenVX 时(请参阅 与 OpenCV 互操作性的示例 部分)。

注意图验证的开销

图验证步骤是一项重量级操作,应避免在“紧循环”执行时进行。请注意,更改图输入的元数据(例如,大小或类型)可能会使图无效。有关更新数据的技巧,请参阅 OpenVX* 图像的映射/取消映射 部分。

将 OpenVX 性能与本机代码进行比较

将 OpenVX 性能与本机代码(例如,C/C++ 或 OpenCL)进行比较时,请确保两个版本尽可能相似。

  • 包装完全相同的操作集。
  • 估算执行时间时,请勿包含图验证。图验证旨在分摊到图执行的多次迭代中。
  • 单独跟踪数据传输成本(读取/写入图像、数组等)。此外,尽可能使用数据映射,因为它更接近常规代码中数据传递的方式(通过指针)。
  • 要求相同的精度。

启用每个节点的性能分析

到目前为止,我们讨论了图的整体性能。为了获得每个节点的性能数据,OpenVX* 1.1 规范明确要求在应用程序中启用性能分析。有一个专门的指令来执行此操作。

vx_status res = vxDirective(context, VX_DIRECTIVE_ENABLE_PERFORMANCE);

注意: 每个节点的性能分析是按上下文启用的。由于它可能引入某些开销,请在生产代码中以及/或在测量整体图性能时将其禁用。

启用分析后,您可以获取节点的性能信息。

vx_perf_t perf;
vxQueryNode(node, VX_NODE_ATTRIBUTE_PERFORMANCE, &perf, sizeof(perf));
printf("Average exec time for the %h node: %0.3lfms\n", node, (vx_float32)perf.avg/1000000.0f);

这正是 Intel® 计算机视觉 SDK 算法设计器 (VAD) 等工具为每个节点收集/公开的数据,例如,请参阅自动对比度示例(在 Intel(R) 计算机视觉 SDK 示例应用程序 部分),以及VAD 开发者指南。类似地,一些 SDK 示例(例如车道检测)也在简化的时间线上显示了这些数据。

注意:请注意,要获取 GPU 上运行的节点的性能数据,您需要设置以下环境变量。

$export VX_CL_PERFORMANCE_REPORTING=1

通用经验法则

Intel 对 OpenVX* API 的扩展:原始操作

OpenVX* 基本视觉内核

Intel® 计算机视觉 SDK 附带了标准的 OpenVX* 计算机视觉函数集。有关这些函数的 API 详细信息,请参阅Khronos* Group – OpenVX™ 规范 1.1视觉函数 部分。

CPU 和 GPU OpenVX 目标都支持核心视觉函数,而扩展则高度依赖于设备。

Intel 的 OpenVX* 内核扩展

Intel® 计算机视觉 SDK 将实验性和已批准的(Intel 的“供应商”和/或 Khronos*)扩展与 OpenVX 分开。您可以在专用的 Intel® 计算机视觉 SDK 内核扩展文档(可在 SDK 安装目录的 /documentation/ 文件夹中找到)中找到函数原型和参数详细信息等信息。

注意:与真正的供应商扩展不同,对被定义为实验性扩展的内核的支持可能会在不通知的情况下更改。

有关实验性扩展声明,请参阅 <VX/vx_intel_volatile.h> 头文件。

支持卷积神经网络 (CNN)

卷积神经网络 (CNN) 在计算机视觉领域引起了极大的兴趣。近年来,使用 CNN 进行图像识别取得了巨大进步。CV SDK 支持官方 Khronos OpenVX* 1.1 神经网络扩展(请参阅 Khronos* 神经网络扩展 主题)。这使得 CNN 的识别任务可以在嵌入式环境中加速。有关 API 和新数据对象参考,请参阅 SDK /documentation/ 目录中的文档。

注意:强烈建议您在 GPU 目标上尝试带有 CNN 节点的图(请参阅 将单个节点调度到不同目标 章)。

部署 CNN 可能是一项艰巨的任务,因为最终目标环境可能与训练环境不同。训练通常在高配数据中心进行,使用 Caffe*、Torch* 等流行训练框架。相比之下,评分(或推理)可以在嵌入式设备上进行。这些环境的限制要大得多,因此通常使用专用的评分 API 可以获得更好的性能。这些限制使得部署和训练环境在数据类型支持和用于评分的 API 方面都看起来非常不同。

Intel 的模型优化器工具旨在解决这些特定问题。它简化了带有 OpenVX 的卷积神经网络的部署。例如,通过模型优化器,您可以自动将 Caffe 布局和数据转换为 OpenVX 图和格式。有关使用模型优化器生成图的代码示例,请参阅 Intel(R) 计算机视觉 SDK 示例应用程序 章。

不同目标(设备)的内核扩展支持

请注意,虽然 CPU 和 GPU 都完全支持 OpenVX* 1.1 规范中定义的,但它们支持的扩展集不同。除非您使用异构支持(使用 Intel® 计算机视觉 SDK 进行异构计算 章),否则支持差异对您来说是透明的:只需将所需的节点添加到图中,运行时就会将其分配给支持的设备。

注意: 通常,作为后备设备,CPU 应正常支持所有扩展。但是,此版本有一个显著例外:双边滤波,仅在 GPU 上可用。

注意: IPU 扩展没有 CPU 等效项(因此,如果 IPU 设备不可用,则没有自动后备)。

请仔细考虑您的实际目标平台设备和支持的扩展的可用性。例如,支持的开发系统缺少 IPU 设备(仅在基于下一代 Intel® Atom™ 处理器(原 Apollo Lake)的系统中可用)。

Khronos* 神经网络扩展

Khronos OpenVX* 神经网络扩展指定了一种在 OpenVX 图中执行基于 CNN 的推理的架构。该扩展定义了一个多维张量对象数据结构,可用于连接表示为 OpenVX 节点的神经网络层,以创建灵活的 CNN 拓扑。OpenVX 神经网络层类型包括卷积、池化、全连接、归一化、softmax 和激活。该扩展允许在同一个 OpenVX 图中将神经网络推理与传统的视觉处理操作混合。

有关扩展 API 的详细信息,请参阅 Khronos* OpenVX™ 神经网络扩展 文档。

与其他 API 的互操作性

OpenVX* 与 Intel® Media Server Studio 或 OpenGL* 等媒体/图形 API 之间共享数据的通用方法是基于共享系统内存。也就是说,在您的代码中,您应首先将数据从 API 映射或复制到 CPU 地址空间。然后,生成的指针可以包装成 OpenVX 图像。或者,可以将数据(指针引用)复制到 OpenVX 对象,以解锁原始数据(并将所有权返回给另一个 API)。

类似地,可以在主机上访问 OpenVX 数据,并将其位复制到或包装到另一个 API 的数据对象中。但请注意,大多数 API 不支持具有像素之间跨距的内存布局,而这在 OpenVX 中是可能的,因此您的数据共享例程应尊重跨距。

无论如何,要更新 OpenVX 图像或从中读取数据,您必须使用映射/取消映射方法。

有关更多详细信息,请参阅以下主题

使用 OpenVX* 图像包装系统内存

考虑通用互操作流程。

  1. 将资源映射到 CPU。
  2. 创建包装内存的 OpenVX 图像。
    // populate the OpenVX structure specifying the image layout
    vx_imagepatch_addressing_t frameFormat;
    //dimension in x, in pixels
    frameFormat.dim_x =...;
    //dimension in y, in pixels
    frameFormat.dim_y = ..;
    //stride in x, in bytes
    frameFormat.stride_x = ...;
    //stride in y, in bytes
    frameFormat.stride_y = ...;
       void* frameData = ...;
       //creating input image, by wrapping the host side 
     pointer
            vx_image im = vxCreateImageFromHandle(
                context,
                VX_DF_IMAGE_RGB,
                &frameFormat,
                &frameData,
                VX_MEMORY_TYPE_HOST);
  3. 在 OpenVX 中使用缓冲区。
  4. 在 OpenVX 图使用完映射的资源位后,应销毁 OpenVX 图像,然后才能安全地取消映射资源,将数据所有权返回给另一个 API。

注意: 在 OpenVX* 使用期间,原始资源(及其映射的数据)不能被应用程序使用。

有关共享图像位与 OpenCV 的示例,请参阅 Auto-Contrast SDK 示例,在 Intel(R) 计算机视觉 SDK 示例应用程序 部分。

OpenVX* 图像的映射/取消映射

您的代码不需要在每次新数据(例如,来自相机)到达时重新创建 OpenVX* 图像。将新图像设置为图输入可能会使图无效,从而触发隐式验证。

此外,大多数您可能希望与之进行互操作的 API 并不保证每次都会返回相同的资源映射指针。在许多情况下,创建常规 OpenVX 图像然后每次更新图像数据时进行数据复制会更简单。图像与其他 OpenVX 资源(例如,数组或矩阵)一样,可以通过映射/取消映射进行操作。

//create an image in a regular way
vx_image image = vxCreateImage(context, 640, 480,VX_DF_IMAGE_RGB);
vx_imagepatch_addressing_t addr;
vx_map_id map_id;
void* ptr = 0;
//request a pointer in the system memory system
vx_status status = vxMapImagePatch(image, &rect, 0,&map_id, &addr, &ptr, usage, VX_MEMORY_TYPE_HOST, 0);
//manipulate image data (referenced by the ptr), e.g. update the values
...
//commit back the changes and pass the ownership back to the OpenVX
status = vxUnmapImagePatch(image, map_id);

与 OpenCV 互操作性的示例

与使用专用地址空间和/或特殊数据布局(例如,压缩的 OpenGL* 纹理)的 API 不同,常规 OpenCV 数据对象(如 cv::Mat)位于常规系统内存中。也就是说,内存可以与 OpenVX* 共享,并且只需传输数据所有权。

//Capture from default webcam 
cv::VideoCapture inp;
inp.open(0);
cv::Mat matBGR, mat RGB;
//read first frame so the ‘matRGB’ allocates data, etc
inp.read(matBGR);
//convert the format, as OpenCV’s default is BGR
cv::cvtColor(matBGR, matRGB, cv::COLOR_BGR2RGB);
// wrap into the OpenVX
vx_imagepatch_addressing_t addr;
addr.dim_x = matRGB.cols;
addr.dim_y = matRGB.rows;
addr.stride_x = matRGB.elemSize();
addr.stride_y = matRGB.step;
void *ptr[] = { matRGB.data };
vx_image image=vxCreateImageFromHandle(ctx,…,&addr,ptr, VX_MEMORY_TYPE_HOST);
//Capture frames
while(1)
{
    // Per spec, mapping of the image created from handle should return 
    // same address (and same memory layout)
    void*p = NULL;
    vx_imagepatch_addressing_t addr;
    vx_rectangle_t rect = { 0, 0, matRGB.cols, matRGB.rows };
    //Request pointer to the system mem
    vx_map_id map_id;
    vxMapImagePatch(image, &rect, 0, &map_id, &addr, &p, VX_WRITE_ONLY );
    //Read the next image
    inp.read(matBGR);
    cv::cvtColor(matBGR, matRGB, cv::COLOR_BGR2RGB);
    // Notify the OpenVX we are done
    vxUnmapImagePatch(image, map_id);
    //Do processing with OpenVX
    ...
}

请注意,原始 cv::Mat 不能同时被应用程序和 OpenVX 使用。

其他重要注意事项

  • 确保 OpenCV cv::Mat 对象在相应 vx_image 对象完成工作之前不会被销毁(或调整大小等)。
  • 请注意,默认情况下,OpenCV 对彩色图像使用 BGR 色彩空间,而 OpenVX 只支持 RGB 通道顺序,因此您可能需要进行格式转换,请参阅上面的示例。

有关 OpenCV 捕获、OpenVX* 处理和 OpenCV 渲染的详细信息,请参阅 Intel(R) 计算机视觉 SDK 示例应用程序 部分的 SDK 示例。

OpenCV 处理 OpenVX* 图像的示例

通常,标准的 OpenVX* 内核和扩展不足以实现特定的计算机视觉管线。类似地,vx_image 可以使用 OpenCV 访问,使用相同的 vxMapImagePatch/vxUnmapImagePatch 函数(请参阅 与 OpenCV 互操作性的示例 部分)。

void callOpenCVGaussianBlur(vx_image input, vx_image output)
{
    vx_uint32 width = 0, height = 0;
    vxQueryImage(input, VX_IMAGE_WIDTH, &width, sizeof(width));
    vxQueryImage(input, VX_IMAGE_HEIGHT, &height, sizeof(height));
    vx_rectangle_t rect = { 0u, 0u, width, height };
    vx_imagepatch_addressing_t input_addr = {};
    void *input_ptr = nullptr;
    vx_map_id map_id_in;
    vxMapImagePatch(input, &rect, 0, &map_id_in, &input_addr, &input_ptr,  VX_READ_ONLY);
    vx_imagepatch_addressing_t output_addr = {};
    void *output_ptr = nullptr;
    vx_map_id map_id_out;
    vxMapImagePatch(output, &rect, 0, & map_id_out , &output_addr, &output_ptr, VX_WRITE_ONLY);
    cv::Mat mat_input(input_addr.dim_y * VX_SCALE_UNITY / input_addr.scale_y,
                      input_addr.dim_x * VX_SCALE_UNITY / input_addr.scale_x,
                      CV_8UC1,
                      input_ptr,
                      input_addr.stride_y);
    cv::Mat mat_output(output_addr.dim_y * VX_SCALE_UNITY / output_addr.scale_y,
                       output_addr.dim_x * VX_SCALE_UNITY / output_addr.scale_x,
                       CV_8UC1,
                       output_ptr,
                       output_addr.stride_y);
    cv::GaussianBlur(mat_input, mat_output, cv::Size(5, 5), 1.5);
    vxUnmapImagePatch(input, map_id_in);
}

此外,您可以使上述函数完全可插入 OpenVX 图。具体来说,OpenVX* 支持用户定义的函数,这些函数作为节点从图内部执行。如多个示例(请参阅 Intel(R) 计算机视觉 SDK 示例应用程序 章)中所述,要定义用户内核,您只需实现几个回调函数以及执行例程本身。

与 OpenCV 透明 API 的互操作性

OpenCV 3.0 引入了新的透明 API,它为 OpenCV 函数启用了基于 OpenCL* 的代码路径。T-API 基于 cv::Umat 对象,该对象抽象了数据表示。请注意,所有对 cv::Umat 对象的访问都应通过 cv::UMat 方法进行,而不是直接通过数据指针。

要混合 cv::UMatvx_image 对象,您可以使用上述的 vx_image->cv::Mat 映射方法,然后使用 OpenCV copyTo 方法将数据从映射的 cv::Matcv::UMat 对象传播。

void callOpenCV_TAPI(vx_image input, vx_image output)
{
    vx_uint32 width = 0, height = 0;
    vxQueryImage(input, VX_IMAGE_WIDTH, &width, sizeof(width));
    vxQueryImage(input, VX_IMAGE_HEIGHT, &height, sizeof(height));
    vx_rectangle_t rect = { 0u, 0u, width, height };
    vx_imagepatch_addressing_t input_addr = {};
    void *input_ptr = nullptr;
    vx_map_id map_id_in;
    vxMapImagePatch(input, &rect, 0, &map_id_in, &input_addr, &input_ptr, VX_READ_ONLY);
   vx_imagepatch_addressing_t output_addr = {};
   void *output_ptr = nullptr;
   vx_map_id map_id_out;
   vxMapImagePatch(output, &rect, 0, & map_id_out , &output_addr, &output_ptr, VX_WRITE_ONLY);
    cv::Mat mat_input(input_addr.dim_y * VX_SCALE_UNITY / input_addr.scale_y,
                      input_addr.dim_x * VX_SCALE_UNITY / input_addr.scale_x,
                      CV_8UC1,
                      input_ptr,
                      input_addr.stride_y);
    cv::Mat mat_output(
                output_addr.dim_y * VX_SCALE_UNITY / output_addr.scale_y, 
                output_addr.dim_x * VX_SCALE_UNITY / output_addr.scale_x,
                       CV_8UC1,
                       output_ptr,
                       output_addr.stride_y);
    cv::UMat umat_input;
    mat_input.copyTo(umat_input);
    cv::UMat umat_tmp, umat_output;
    cv::GaussianBlur(umat_input, umat_tmp, cv::Size(5, 5), 1.5);
    cv::Canny(umat_tmp, umat_output, 50, 100);
    umat_output.copyTo(mat_output);
    vxUnmapImagePatch(input, map_id_in);
}

OpenCV 和 OpenVX* 参数解释上的差异

从 OpenCV 移植到 OpenVX* 时,存在 API 之间参数匹配直接的情况。但是,参数含义不同的情况也存在。例如,OpenCV 和 OpenVX 都提供 WarpAffineWarpPerspective 函数,但存在以下细微差别:

  • OpenVX 的 WarpAffineWarpPerspective 实现对应 OpenCV 调用的 cv::WARP_INVERSE_MAP 变体。
  • OpenCV 和 OpenVX 使用不同的公式来访问变换矩阵。为了将变换矩阵转换为 OpenVX,必须对 OpenCV 矩阵进行转置。

同样,OpenCV 的 Filter2D 和 OpenVX 的 Convolution 都是滤波例程,可以互换使用,并附带以下说明:

  • OpenVX Convolution 公式与 OpenCV 的 Filter2D 不同。为了在 OpenVX 中获得相同的结果,原始 OpenCV“内核”参数必须同时围绕两个轴翻转。
  • OpenVX Convolution 使用 16 位有符号整数数据类型作为“内核”参数,并附加了一个尺度参数,调用 OpenCV 时需要相应地进行转换。

Intel 对 OpenVX* API 的扩展:OpenCL™ 自定义内核

自定义 OpenCL™ 内核是 Intel OpenVX* 扩展的简称,它允许将常规 OpenCL 内核中的代码用作 OpenVX 用户内核。有关代码示例,请参阅 Intel(R) 计算机视觉 SDK 示例应用程序 部分。

注意: 此扩展是一项实验性功能。它可能会在不通知的情况下发生更改。OpenCL 自定义内核仅支持 GPU 目标。

当前扩展的 API 仍然限制了 OpenCL 内核可以访问的 OpenVX 参数。

OpenVX* 对象 OpenCL™ 内核参数(OpenCL™ C 代码)

vx_image

global uchar* ptr,

uint width, // 输出图像宽度

uint height, // 输出图像高度

uint pixelStride, // 像素跨距(字节)

uint rowPitch, // 行跨距(字节)

// 对每个平面重复(对于多平面图像)

vx_array

global uchar* ptr,

uint32 items, // 项目数量

uint32 itemSize, // 项目大小(字节)

uint32 stride // 项目跨距(字节)

vx_matrix

global uchar* ptr, // 指向矩阵项目的指针

uint32 rows, // 行数

uint32 cols // 列数

vx_tensor

global dtype* ptr, // 数据指针(short/float/half)

long offset, // ptr 偏移量,实际内容
// 开始(以字节为单位)

int num_dims, // 张量的维度数

global int* dims, // 维度大小

global int* strides // 维度跨距

vx_scalar (VX_TYPE_INT16)

short

vx_scalar (VX_TYPE_INT32)

int

vx_scalar (VX_TYPE_INT64)

long

vx_scalar (VX_TYPE_UINT16)

ushort

vx_scalar (VX_TYPE_UINT32)

uint

vx_scalar (VX_TYPE_UINT64)

ulong

vx_scalar (VX_TYPE_FLOAT32)

float

vx_scalar (VX_TYPE_FLOAT64)

double

vx_scalar (VX_ENUM)

int

vx_threshold

int value_or_upper, // BINARY 类型的值
// RANGE 类型的 upper

int lower, // RANGE 类型定义

int true_value, // VX_THRESHOLD_TRUE_VALUE 属性

int false_value // VX_THRESHOLD_FALSE_VALUE 属性

当应用程序将 OpenVX 数据对象(如 vx_imagevx_array)设置为节点的参数时,OpenVX 框架会将实际数据指针传递给 OpenCL 内核调用。所有 work-item 内置函数(如 get_global_id())都受支持。因此,程序员负责确保正确的基于 ID 的指针计算以及从 OpenCL 内核对 OpenVX 数据对象的合法访问。

有关扩展 API 的实际应用,请参阅 Intel(R) 计算机视觉 SDK 示例应用程序 部分。您可以在其中找到专门的教程,详细解释 OpenCL 自定义内核。

对于支持图像分块的自定义 OpenCL 内核(通过 vxIntelAddDeviceTilingKernel 添加到上下文),vx_image 的参数翻译不同。

OpenVX* 对象 OpenCL™ 分块内核参数(C 代码)

vx_image

global uchar* // ptr,

int tile_x, // 分块的 x 坐标

int tile_y, // 分块的 y 坐标

int tile_width, // 分块宽度

int tile_height, // 分块高度

int image_width, // 完整图像宽度

int image_height, // 完整图像高度

int pixelStride, // 像素跨距(字节)

int rowStride, // 行跨距(字节)

// 对每个平面重复(对于多平面图像)

注意: ptr 指向分块的起始位置,而不是完整图像的起始位置。不应访问位于给定分块外部的像素。

此外,应将内核属性 VX_KERNEL_INPUT_NEIGHBORHOOD_INTEL 设置为相应的 vx_neighborhood_size_intel_t 对象。这是由运行时用于确定给定输出分块所需的输入分块的大小和位置。

例如,以下代码片段演示了 5x5 滤波器所需的邻域。

vx_neighborhood_size_intel_t n;
n.top = -2;  //requires 2 additional input pixel above
n.left = -2; //requires 2 additional input pixels to the left
n.right = 2; //requires 2 additional input pixels to the right
n.bottom = 2; //requires 2 additional input pixels below
vxSetKernelAttribute(kernel, 
                     VX_KERNEL_INPUT_NEIGHBORHOOD_INTEL, 
                     &n,
                     sizeof(vx_neighborhood_size_intel_t));

注意: 对于导致输入分块的一部分位于输入图像有效区域之外的邻域输入分块,该分块将被裁剪为仅包含与输入图像有效区域相交的部分。

Intel 对 OpenVX* API 的扩展:高级分块

OpenVX* 高级分块扩展是 Intel 的 API 扩展。它旨在支持需要复杂分块依赖项的内核,超越了 Khronos OpenVX* 用户内核分块扩展的一对一映射。

有关详细的代码示例,请参阅 Intel(R) 计算机视觉 SDK 示例应用程序 部分中的Census-Transform SDK 示例说明。

高级分块扩展支持以下内容:

  • 需要输入/输出分块之间非平凡映射的内核,例如旋转或缩放。
  • 需要串行生成输出分块的内核(如下面的错误扩散)。
  • 输出分块依赖于多个输入图像的不同分块的内核。

使用以下接口将高级分块内核添加到 OpenVX 上下文。

vx_kernel vxAddAdvancedTilingKernelIntel(vx_context context,
        vx_char name[VX_MAX_KERNEL_NAME], vx_enum enumeration,
        vx_advanced_tiling_kernel_intel_f kernel_func_ptr,
        vx_mapping_intel_f mapping_func_ptr,
        vx_uint32 num_params,
        vx_kernel_validate_f validate,
        vx_kernel_initialize_f initialize,
        vx_kernel_deinitialize_f deinitialize,
        vx_kernel_preprocess_intel_f preprocess,
        vx_kernel_postprocess_intel_f postprocess,
        vx_kernel_set_tile_dimensions_intel_f settiledimensions
        vx_kernel_tile_dimensions_initialize_intel_f tiledimensionsinit);

有关更多信息,请参阅以下主题:

支持高级分块的内核的强制回调

注册高级分块内核需要为 vxAddAdvancedTilingKernelIntel 定义和指定以下回调函数:

  • 输入和输出验证函数。在 vxVerifyGraph 期间调用,用于验证每个输入和输出参数(由索引定义)。
    typedef vx_status (*vx_kernel_validate_f) (vx_node node, const vx_reference parameters[], 
    vx_uint32 num, vx_meta_format metas[]);

    此函数是任何用户内核的常规回调,如 OpenVX* 1.1 规范所定义。

  • 高级分块内核函数。在图执行期间,每次调用此函数都会处理输出图像中包含的每个分块。
    typedef vx_status (*vx_advanced_tiling_kernel_intel_f)(vx_node node, void  * parameters[], 
    vx_uint32 num, void *tile_memory, vx_size tile_memory_size);
  • 分块映射函数,用于在 vxVerifyGraph 中的图验证期间描述输出分块与输入分块之间的依赖关系。此函数用于设置输入分块(srcRectOut)的属性,以生成给定的输出分块(dstRectIn),该参数作为 const 参数传递给分块映射函数。

    如果内核有多个 vx_image 输入参数,此分块映射函数将分别对每个输入调用。param_num 参数指定运行时请求分块映射的输入参数索引。

    typedef vx_status (*vx_mapping_intel_f) (vx_node node, vx_reference parameters[], 
    const vx_tile_t_attributes_intel_t* dstRectIn, vx_tile_t_attributes_intel_t* srcRectOut,vx_uint32 param_num);

有关代码示例,请参阅 整合:示例代码 主题,其中包含分块内核的代码示例。

支持高级分块的内核的可选回调

高级分块内核可以选择提供以下回调函数,如果设置了这些函数,OpenVX* 运行时将在 vxVerifyGraphvxProcessGraph 中调用它们:

  • 初始化
    typedef vx_status (*vx_kernel_initialize_f)(vx_node node, const vx_reference 
     *parameters, vx_uint32 num); 

    如果设置,将在 vxVerifyGraph 期间调用一次,所有输入和输出参数都已验证后。

  • 反初始化

    typedef vx_status (*vx_kernel_deinitialize_f)(vx_node node, const vx_reference 
     *parameters, vx_uint32 num);

    如果设置,将在节点销毁期间调用此函数。

    上述两个函数是任何用户内核的常规回调,如 OpenVX* 1.0.1/1.1 规范所定义。有关使用详情,请参阅 Khronos* 的标准规范,有关代码示例,请参阅 整合:示例代码 部分。

  • 预处理

    typedef vx_status (*vx_kernel_preprocess_intel_f)(vx_node node, const 
     vx_reference *parameters, vx_uint32 num, void * tile_memory[], vx_uint32 
     num_tile_memory_elements,vx_size tile_memory_size);

    如果设置,此函数将在每次用户调用 vxProcessGraph 开始时调用一次,在图中的任何节点开始处理之前。它旨在用于执行每次 vxProcessGraph 所需的任何初始化,这可能对某些内核是必需的。

  • 后处理

    typedef vx_status (*vx_kernel_postprocess_intel_f)(vx_node node, const 
     vx_reference *parameters, vx_uint32 num, void * tile_memory[], vx_uint32 
     num_tile_memory_elements,vx_size tile_memory_size);

    如果设置,此函数将在每次用户调用 vxProcessGraph 结束时调用一次,在图中的所有节点都完成处理之后。它旨在用于执行每次 vxProcessGraph 所需的任何反初始化或数据聚合,这可能对某些内核是必需的。

    注意: 由于此函数在图中的所有节点都完成处理后调用,因此不建议在此函数中设置输出参数的值,如果该参数在同一图中的另一个节点中被用作输入参数。

  • 设置分块尺寸

    typedef vx_status (*vx_kernel_set_tile_dimensions_intel_f)(vx_node node, 
     const vx_reference *parameters, vx_uint32 param_num, const vx_tile_block_size_intel_t 
     *current_tile_dimensions, vx_tile_block_size_intel_t *updated_tile_dimensions);

    如果设置,此函数将在 vxVerifyGraph 中调用,让内核能够用自定义分块大小覆盖运行时选择的当前输出分块大小。例如,一个内核可能被设计为仅与宽度等于输出图像大小的分块一起工作。该函数需要使用内核参数和当前分块尺寸的组合来设置更新后的分块尺寸,当前分块尺寸作为输入参数传递。

  • 初始化(作为分块尺寸的函数)

    typedef vx_status (*vx_kernel_tile_dimensions_initialize_intel_f)(vx_node 
     node, const vx_reference *parameters, vx_uint32 param_num, const vx_tile_block_size_intel_t 
     *current_tile_dimensions);

    如果设置,此函数将在 vxVerifyGraph 中调用,让内核能够根据当前设置的分块尺寸执行一些初始化。例如,一个内核可能希望设置一个节点属性,如 VX_KERNEL_TILE_MEMORY_SIZE_INTEL,作为当前设置的分块尺寸的函数。

vxVerifyGraph 期间的高级分块顺序图

以下是 vxVerifyGraph 中所需和可选的高级分块内核回调的调用序列图:

vxProcessGraph 期间的高级分块顺序图

以下是 vxProcessGraph 中所需和可选的高级分块内核函数的调用序列图:

自定义高级分块内核的行为

除了在内核注册时传递给 vxAddAdvancedTilingKernelIntel 的函数回调之外,还可以使用 vxSetKernelAttribute 设置许多内核属性。

注意:所有本节描述的属性都必须在调用 vxFinalizeKernel 之前设置。

例如,由 vx_node_attribute_tiling_intel_e 枚举定义的属性系列。

  • VX_KERNEL_SERIAL_TYPE_INTEL,用于指示该内核的各个输出分块必须“按顺序”生成。

    执行顺序通过提供 vx_serial_type_intel_e 类型的参数来定义。

    vx_serial_type_intel_e v = <specific order, below>;
            vxSetKernelAttribute(k,VX_KERNEL_SERIAL_TYPE_INTEL,&v,sizeof(v));

    可能的执行顺序是:

  • VX_SERIAL_LEFT_TO_RIGHT_TOP_TO_BOTTOM_INTEL

    为了生成给定的输出分块,必须先完成其“左边”(X 方向)的分块。如果给定输出分块的 x 坐标为 0,则必须先完成上一行分块的最右侧分块。坐标为 (0,0) 的分块是唯一没有依赖项的输出分块。

  • VX_SERIAL_LEFTTOP_TO_RIGHTBOTTOM_INTEL

    为了生成给定的输出分块,必须先完成其“左边”(X 方向)的分块,“上方”(Y 方向)的分块以及“左上方”的分块。如果给定输出分块的 x 坐标为 0,则该分块仅依赖于“上方”的分块。如果给定输出分块的 y 坐标为 0,则该分块仅依赖于“左边”的分块。坐标为 (0,0) 的分块是唯一没有依赖项的输出分块。

高级分块内核的就地处理

两个专用内核属性(VX_KERNEL_INPLACE_KERNEL_INTELVX_KERNEL_INPLACE_PARAMS_FUNC_INTEL)共同控制中间缓冲区的用法。

通常,OpenVX* 运行时会为节点的输入和输出维护单独的中间缓冲区。对于某些内核,可以使用相同的缓冲区来节省内存并提高性能。这被称为“就地”(in-place)缓冲区。

  • VX_KERNEL_INPLACE_PARAMS_FUNC_INTEL 设置为 vx_inplace_params_intel_f 类型的函数指针。
    typedef vx_int32 (*vx_inplace_params_intel_f) (vx_reference parameters[], 
             vx_uint32 output_image_index);

    此函数的目的是通知运行时,哪个输入 vx_image 可以与给定的输出 vx_image 共享。传递了给定输出 vx_image 参数的索引。如果给定的输出 vx_image 无法与任何输入 vx_image 共享,则该函数应返回 -1

  • VX_KERNEL_INPLACE_KERNEL_INTEL 设置为 vx_advanced_tiling_kernel_intel_f 类型的函数指针。这与(在内核注册时传递给 vxAddAdvancedTilingKernelIntel 的)高级内核函数类型相同。在许多情况下,无论是否使用“就地”缓冲区,都可以使用相同的内核函数。

即使设置了这两个属性,OpenVX* 运行时也可能在某些情况下忽略“就地”缓冲区的请求。例如,如果输入 vx_image 不是虚拟图像,或者输入缓冲区与其他节点共享。

使用专用线程暂存内存的高级分块内核

内核函数 vx_advanced_tiling_kernel_intel_f 可能被多个运行时线程同时调用,以并行处理许多分块。内核可能需要一些暂存内存,专供每个工作线程使用,以辅助处理,或用作线程特定存储。

如果专用线程暂存内存的所需大小在调用 vxFinalizeKernel 之前已知,则可以将以下内核属性设置为以字节为单位的专用线程暂存内存大小。

VX_KERNEL_TILE_MEMORY_SIZE_INTEL

如果专用线程暂存内存的所需大小不是给定内核所有实例的常量,而是取决于在 vxVerifyGraph 之前才能知道的参数或属性,则可以将以下节点属性设置为以字节为单位的专用线程暂存内存大小。

VX_NODE_TILE_MEMORY_SIZE_INTEL

如果设置,运行时将为每个运行时线程分配一个缓存对齐的缓冲区,该线程可能在 vxProcessGraph 期间调用 vx_advanced_tiling_kernel_intel_f。对于每次调用高级分块内核函数,运行时将传递 tile_memory(线程专用暂存内存缓冲区的起始指针)以及 tile_memory_size(分配的缓冲区大小)。

typedef vx_status (*vx_advanced_tiling_kernel_intel_f)(vx_node node, 
 void * parameters[], vx_uint32 num, void * tile_memory, vx_size tile_memory_size);

如果上述“分块内存大小”属性均未设置,则 vx_advanced_tiling_kernel_intel_f 将以 tile_memory 指针等于 NULL,tile_memory_size 等于 0 的形式调用。

除了将专用线程暂存内存传递给 vx_advanced_tiling_kernel_intel_f 之外,缓冲区还会传递给可选的回调函数 vx_kernel_preprocess_intel_fvx_kernel_postprocess_intel_f

typedef vx_status (*vx_kernel_preprocess_intel_f)(vx_node node, const 
 vx_reference *parameters, vx_uint32 num, void * tile_memory[], vx_uint32 
 num_tile_memory_elements,vx_size tile_memory_size);
typedef vx_status (*vx_kernel_postprocess_intel_f)(vx_node node, const 
 vx_reference *parameters, vx_uint32 num, void * tile_memory[], vx_uint32 
 num_tile_memory_elements,vx_size tile_memory_size);

由于这些函数仅为每个 vxProcessGraph 调用一次,因此每个线程专用缓冲区的指针将作为参数传递。tile_memory 是指针数组。此数组的每个元素对应于特定线程的专用暂存内存缓冲区。num_tile_memory_elements 指定 tile_memory 数组中的元素数量。tile_memory_size 指定缓冲区的分配大小。

访问“预处理”和“后处理”中的线程特定暂存缓冲区可能在许多方面很有用。一个具体的例子是自定义直方图内核。此直方图内核可能会选择将直方图数据累积到线程特定的直方图中,因为这比在 vxProcessGraph 期间同步对单个直方图的访问更有效。在 vx_kernel_preprocess_intel_f 中,内核需要确保线程特定的直方图已初始化为零。在 vx_kernel_postprocess_intel_f 中,内核将通过累积每个线程特定直方图条目的结果来生成最终直方图。

整合:示例代码

本节总结了实现高级分块用户节点的事件函数的代码片段。示例节点实现了图像转置操作。

首先,任何用户节点都需要根据规范的输入参数验证函数(此处省略了错误处理)。

static vx_status TranspValidator (vx_node node, const vx_reference parameters[], vx_uint32 num, vx_meta_format metas[])
{
    if (num!=2) //the node accepts single input (image) and single ouput only
        return VX_ERROR_INVALID_PARAMETERS;
    vx_status status = VX_FAILURE;
    vx_df_image format;
    vxQueryImage((vx_image)parameters[0],VX_IMAGE_FORMAT,&format,   sizeof(format)); //input image format
    if (format != VX_DF_IMAGE_U8)//the node supports only U8 images
    {
        return VX_INVALID_FORMAT;
    }
 
    vx_uint32 in_width = 0, in_height = 0;
    const size_t sz = sizeof(vx_uint32);
    //Query the width & height of the input image 
    vxQueryImage((vx_image)parameters[0],VX_IMAGE_WIDTH, &in_width, sz);
    vxQueryImage((vx_image)parameters[0],VX_IMAGE_HEIGHT,&in_height,sz);
    //Swap the width & height as required by transpose, accordingly
    vx_uint32 out_width = in_height;
    vx_uint32 out_height = in_width;
    vx_df_image format = VX_DF_IMAGE_U8;
    status = vxSetMetaFormatAttribute(meta[1], VX_IMAGE_WIDTH, &out_width, sz);
    status |= vxSetMetaFormatAttribute(meta[1], VX_IMAGE_HEIGHT, &out_height, sz);
    status |= vxSetMetaFormatAttribute(meta[1], VX_IMAGE_FORMAT, &format, sizeof(format));
    return status;
}

此函数验证(单个)输入图像的格式是否正确。它还验证输出参数,并设置定义输出图像的“元格式属性”。对于转置,输出宽度等于输入高度,输出高度等于输入宽度。验证器还指定输出图像的格式与输入相同(VX_DF_IMAGE_U8)。

分块映射函数是高级分块的第一个特定回调。

static vx_status TranspTileMapping(vx_node node, vx_reference params[],
        const vx_tile_t_attributes_intel_t* dstRectIn,
        vx_tile_t_attributes_intel_t* srcRectOut,
        vx_uint32 param_num)
{
    //Given the output tile attributes (dstRectIn),
    //set the input tile attributes which are required (srcRectOut)
    srcRectOut->x = dstRectIn->y;
    srcRectOut->y = dstRectIn->x;
    srcRectOut->tile_block.width  = dstRectIn->tile_block.height;
    srcRectOut->tile_block.height = dstRectIn->tile_block.width;
    return VX_SUCCESS;
}

此函数设置生成给定输出分块所需的输入分块。对于转置,映射是简单的(交换)。

最后一个必需的函数是内核函数,它实际执行转置计算。

static vx_status TranspAdvancedTilingKernel(vx_node node,
                   void * parameters[],
                   vx_uint32 num,
                   void * tile_memory,
                   vx_size tile_memory_size)
{
    vx_tile_t *pInTileIn  = (vx_tile_intel_t *)parameters[0];<!--?rh-implicit_p ????-->
    vx_tile_t *pInTileOut = (vx_tile_intel_t *)parameters[1];
    const vx_uint8 *pSrc = pInTileIn->base[0];
    vx_uint8 *pDst = pInTileOut->base[0];
        
    const int srcStride = pInTileIn->addr[0].stride_y;
    const int dstStride = pInTileOut->addr[0].stride_y;
    const int output_tile_width = pInTileOut->addr[0].dim_x;
    const int output_tile_height = pInTileOut->addr[0].dim_y;
    vx_uint8 *pDstRow = pDst;
    //for each output line
    for(int y = 0; y < output_tile_height; y++)
    {
        vx_uint8 *pSrcColumn = pSrc + y*sizeof(vx_uint8);
        //for each output pixel for this line
        for(int int x = 0; x < output_tile_width; x++)
        {
            pDstRow[x] = *pSrcColumn;
            pSrcColumn += srcStride;
        }
        pDstRow += dstStride;
    }
    return VX_SUCCESS;
}

运行时将传递要处理的输出分块的指针,以及通过分块映射函数请求的输入分块的指针。请注意源缓冲区和目标缓冲区正确使用了 stride_y。分块内核不应假设缓冲区在系统内存中的布局方式,并应遵守跨距。

//create a kernel via the vxAddAdvancedTilingKernelIntel interface
vx_kernel kernel = vxAddAdvancedTilingKernelIntel
        (context, 
        (vx_char *)"AdvancedTilingTranspose",//name
        TRANSPOSE_ID,                       //enumeration
        TranspAdvancedTilingKernel,         //kernel_func_ptr
        TranspTileMapping,                  //mapping_func_ptr
        2,                                  //num_params
        TranspValidator,               //input/output validate
        NULL,                               //initialize
        NULL,                               //deinitialize
        NULL,                               //preprocess
        NULL,                               //postprocess
        NULL,                               //settiledimensions
    NULL);                              //tiledimensionsinit
//add the parameters for this kernel
vxAddParameterToKernel(kernel, 0, VX_INPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED);
vxAddParameterToKernel(kernel, 1, VX_OUTPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED);
//done with publishing the kernel
vxFinalizeKernel(kernel);

高级分块错误扩散示例代码

错误扩散是一种常用的技术,用于将多级图像转换为二值图像。在此示例中,我们将实现一种错误扩散版本,该版本接受单通道、每像素 8 位(范围 0255)的图像作为输入,并生成每像素 8 位(值仅为 0255)的单通道图像。

基本思想是将每个输入像素值与固定阈值(本例中为 127)进行比较。如果像素值高于阈值,则相应的输出像素设置为 255,否则设置为 0。计算“误差”值的方法如下:

Error = InputPixelValue – OutputPixelValue

例如,假设给定输入像素值为 96。此值低于阈值 127,因此输出像素值设置为 0。在这种情况下,误差为 (96 – 0) = 96

然后,为每个像素计算出的误差会扩散到相邻的输入像素。对于我们实现的 Floyd-Steinberg 错误扩散,用于将误差扩散到邻近像素的“系数”为:

  X 7/16
3/16 5/16 1/16

使用上面的例子,(7/16)*96 将扩散到右侧的邻近输入像素,(1/16*96) 将扩散到下方的邻近输入像素,依此类推。

以下图像显示了左侧的输入灰度图像和右侧的错误扩散结果。

错误扩散技术是一个很好的例子,可以说明为什么需要高级分块 API 支持的某些可选回调函数,因为:

  • 要处理给定的输入像素,必须完成相邻输入像素的处理。这意味着我们需要设置输出分块串行化属性。
  • 我们需要一个专用的“错误缓冲区”,我们将在“初始化”函数中分配它。因此,我们将在“反初始化”函数中释放此缓冲区。
  • 我们需要在每个 vxProcessGraph 之前初始化错误缓冲区为 0;因此,我们必须实现“预处理”函数。

为简便起见,我们省略了输入和输出验证器的代码,因为它们与我们在上一个示例中定义的类似。此外,在错误扩散的情况下,输入分块与输出分块之间存在一对一的对应关系;因此,我们省略了分块映射函数,该函数是微不足道的。

以下代码片段显示了“初始化”回调函数的实现。

vx_status ErrorDiffusionInitialize(vx_node node,
const vx_reference *parameters, vx_uint32 num)
{
  // we are going to allocate a floating
  // point error buffer, such that there is an error entry
  // for each pixel.
  vx_image input = (vx_image)parameters[0];
  vx_int32 width, height;
  vxQueryImage(input, VX_IMAGE_WIDTH, &width, sizeof(width));
  vxQueryImage(input, VX_IMAGE_HEIGHT, &height, sizeof(height));
  //we pad image with 2 pixels, to prevent memory
  // access violations on the right and left edges of the image.
  width += 2;
  //we add 1 to the height, to prevent memory
  // access violations on the bottom edge of the image
  height += 1;
  vx_float32 *pErrorBuffer = 
        (vx_float32*)malloc(width*height*sizeof(vx_float32));
  if(!pErrorBuffer )
  {
    return VX_ERROR_NO_MEMORY;
  }
  //free previously set local ptr for this node
  vx_float32 *p = 0;
  vxQueryNode(node, VX_NODE_LOCAL_DATA_PTR, &p, sizeof(p));
  if( p) free(p);
  //set the 'local data ptr' for this node to the new errors buffer.
  return vxSetNodeAttribute(node, VX_NODE_LOCAL_DATA_PTR,                                      
        &pErrorBuffer, sizeof(pErrorBuffer));
}

在错误扩散的情况下,“初始化”函数(在每次用户调用 vxVerifyGraph 时调用一次)会分配我们的错误缓冲区。该缓冲区在运行时用于将错误传播到相邻像素。我们将一个节点属性,“本地数据指针”,设置为已分配的缓冲区。请注意,检查以前分配的缓冲区是否存在以防止内存泄漏是一个好习惯。此数据指针可以在内核函数中使用 vxQueryNode 检索。

“反初始化”仅在节点销毁时调用,因此连续两次用户调用 vxVerifyGraph 具有相同的图意味着两次连续调用我们的“初始化”函数,此时可能已设置了先前分配的缓冲区。

以下代码片段显示了“反初始化”回调函数的实现。

vx_status ErrorDiffusionDeinitialize(vx_node node,
const vx_reference* parameters, vx_uint32 num)
{
  vx_float32 *pErrorBuffer = 0;
  vxQueryNode(node, VX_NODE_LOCAL_DATA_PTR, &pErrorBuffer, sizeof
  pErrorBuffer));
  if( pErrorBuffer )
  {
    free(pErrorBuffer);
    pErrorBuffer = 0;
    // set the local data ptr to 0
    vxSetNodeAttribute(node, VX_NODE_LOCAL_DATA_PTR, 
                       &pErrorBuffer, sizeof(pErrorBuffer));
  }
  return VX_SUCCESS;
}

“反初始化”函数在节点销毁时调用一次。我们必须释放“初始化”函数中分配的错误缓冲区。请注意,必须将节点的“本地数据指针”设置回 0,以防止运行时尝试也释放该指针。

以下代码片段显示了“设置分块尺寸”回调函数。

vx_status ErrorDiffusionSetTileDimensions(vx_node node,
                        const vx_reference *parameters,
                        vx_uint32 num,
                        const vx_tile_block_size_intel_t *current_tile_dimensions,
                        vx_tile_block_size_intel_t *updated_tile_dimensions)
{
  vx_image input = (vx_image)parameters[0];
  vx_int32 width;
  vxQueryImage(input, VX_IMAGE_WIDTH, &width, sizeof(width));
  //Set the desired tile width to the entire input image width
  updated_tile_dimensions->width = width;
  //Keep the height as the current tile height
  updated_tile_dimensions->height = current_tile_dimensions->height;
  return VX_SUCCESS;
}

以下代码片段显示了“预处理”回调函数的实现。

vx_status ErrorDiffusionPreProcess(vx_node node, 
const vx_reference *parameters, vx_uint32 num, void * tile_memory[], vx_uint32 num_tile_memory_elements,vx_size tile_memory_size)
{
  vx_image input = (vx_image)parameters[0];
  vx_int32 width, height;
  vxQueryImage(input, VX_IMAGE_WIDTH, &width, sizeof(width));
  vxQueryImage(input, VX_IMAGE_HEIGHT, &height, sizeof(height));
  vx_float32 *pErrorBuffer = 0;
  size_t sz = sizeof(pErrorBuffer);
        
  vxQueryNode(node, VX_NODE_LOCAL_DATA_PTR, &pErrorBuffer, sz);
  if( !pErrorBuffer )
    return VX_ERROR_NOT_ALLOCATED;
  // patch our width & height (following the logic of the Initialize function) 
  width  += 2;
  height += 1;
  //initialize our error buffer to all 0's
  memset(pErrorBuffer, 0, width*height*sz);
  return VX_SUCCESS;
}

“预处理”回调函数在每次用户调用 vxProcessGraph 开始时调用,在任何节点开始处理之前。我们使用此函数重新初始化我们的错误缓冲区。

为简洁起见,我们省略了错误扩散内核本身的计算代码。

static vx_status ErrorDiffusionAdvancedTilingKernel
(                         vx_node node,void * parameters[],
                         vx_uint32 num,void * tile_memory,
                         vx_size tile_memory_size)
{
  vx_tile_t *pInTileIn = (vx_tile_t *)parameters[0];
  vx_tile_t *pInTileOut = (vx_tile_t *)parameters[1];
  ...
  return VX_SUCCESS;
}

最后,要创建高级分块内核,应调用 vxAddAdvancedTilingKernelIntel 接口,如下面的代码片段所示。

//create a kernel via the vxAddAdvancedTilingKernelIntel interface
  vx_kernel kernel = vxAddAdvancedTilingKernelIntel(context,
                     (vx_char *)"ErrorDiffusion",           //name
                     ERRORDIFFUSION_ID,                   //enumeration
                     ErrorDiffusionAdvancedTilingKernel, //kernel_func_ptr
                     ErrorDiffusionTileMapping,       //mapping_func_ptr
                     2,                               //num_params
                     ErrorDiffusionInputValidator,    //input validate
                     ErrorDiffusionOutputValidator,   //output validate
                     ErrorDiffusionInitialize,        //initialize
                     ErrorDiffusionDeinitialize,      //deinitialize
                     ErrorDiffusionPreProcess,        //preprocess
                     NULL,                            //postprocess
                     ErrorDiffusionSetTileDimensions, //settiledimensions
                     NULL                             //tiledimensionsinit);      
  //specifiy the parameters for this kernel
  vxAddParameterToKernel(kernel, 0, VX_INPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED);
  vxAddParameterToKernel(kernel, 1, VX_OUTPUT, VX_TYPE_IMAGE, VX_PARAMETER_STATE_REQUIRED);
  //set the serial attribute, to produce output tiles serially
  vx_intel_serial_type_intel_e serial_type = VX_SERIAL_LEFTTOP_TO_RIGHTBOTTOM_INTEL;
  vxSetKernelAttribute(kernel, VX_KERNEL_SERIAL_TYPE_INTEL, 
                       &serial_type, sizeof(serial_type));
  //done with publishing
  vxFinalizeKernel(kernel);

使用 Intel® 计算机视觉 SDK 进行异构计算

Intel® 计算机视觉 SDK (Intel® CV SDK) 支持 OpenVX* 1.1 API,用于将图的各个节点调度到特定的加速器或目标。

目前,只有 CPU 和 GPU 被完全支持作为目标。对 IPU 目标的支持有限(因为该设备支持的内核数量非常有限)。

Intel® CV SDK 提供了一个简单的实用程序,可以列出支持的目标及其支持的内核。有关更多详细信息,请参阅 README 文件,位于 <SDK_ROOT>/samples/kernel_enumerator

有关更多信息,请参阅以下主题:

关于 IPU 支持的一般说明

IPU 目标以有限的方式支持。

  • 它仅支持 OpenVX 核心视觉函数的一部分,并且存在进一步的限制(请参阅 Intel 对 OpenVX* 原始操作的支持扩展 部分)。具体来说,仅支持以下标准内核的 3x3 版本:Sobel、Median、Box、Dilate、Erode、Gaussian 和 Canny L1/L2。此外,还支持 Canny 5x5 (L1/L2)。此外,请参阅Intel® 计算机视觉 SDK 内核扩展,了解有关 IPU 支持的扩展的更多信息。
  • 当前实现仅是功能性的,而非高性能的。
  • 唯一的经验法则是,IPU 相对适合图像滤波器,而不适合通用的复杂内核。
  • 在异构场景中使用 IPU 可能会导致数据通信开销和不必要的同步。
  • 同时使用 IPU 进行计算与常规图像捕获(例如,通过 libCamHal)可能会导致未定义行为。

系统级设备亲和性

可以通过 shell 环境变量将执行定向到特定设备,例如,以下行仅启用 GPU。

$ export VX_INTEL_ALLOWED_TARGETS="gpu"

注意: Intel® MKL 驱动的 CPU CNN 内核实现构成了一个单独的目标“mkldnn”。它仅加速常规(32 位)浮点数,而不加速半精度(16 位)浮点数。

几个重要细节

  • 默认情况下(未指定 VX_INTEL_ALLOWED_TARGETS 时),允许所有目标。
  • OpenVX 运行时将尝试尊重目标定义的顺序(例如,如果第一个目标不支持某个扩展,则尝试第二个目标,依此类推)。
  • 建议尝试性能场景(特别是对于包含 CNN 节点的图形,请参阅 OpenVX* 原始对象 Intel 扩展 部分),并使 GPU 目标在亲和性列表中获得最高优先级。

    $ export VX_INTEL_ALLOWED_TARGETS="ipu,cpu,mkldnn,gpu"

  • 如果您的代码使用的扩展不受允许的目标支持,则图执行将失败。

注意:为避免因目标不支持某些视觉功能而导致失败,请始终将“cpu”添加到允许的目标列表中。另一种选择是使用更精细(按节点)的控制,这在下一节中进行了描述。

将单个节点调度到不同目标

Intel® CV SDK 支持 OpenVX* 1.1 的“target”API,用于将 OpenVX 图的单个节点分配给特定目标,从而覆盖运行时默认设置。这使得 Intel 平台能够实现异构使用,以更好地利用硬件。

您可能希望将节点调度到特定目标以提高性能或降低功耗。此外,如果图中存在足够的并行性,图的多个分支(例如单个 RGB 通道)可以调度到不同的目标,然后同时执行。

注意:有关 Targets API 的详细信息和代码示例,请参阅位于 <SDK_ROOT_DIR>/samples/hetero_basic OpenVX Heterogeneous Basic Sample

注意:如前所述,只有 CPU 和 GPU 被支持为完整的目标。IPU 支持是有限的。

注意:Intel® MKL 支持的 CPU 的 CNN 内核实现构成了一个单独的目标“mkldnn”。包含 CNN 节点的图形的目标推荐顺序如下:
$ export VX_INTEL_ALLOWED_TARGETS="cpu,mkldnn,gpu"

以下代码片段说明了目标 API 的应用:

// Instantiate a node in a graph and set its target 
node = vxAccumulateImageNode(graph, input, output);
status = vxSetNodeTarget(node, VX_TARGET_GPU_INTEL, NULL);
if (status == VX_ERROR_NOT_SUPPORTED)
{
        // the node is not supported by the target
   // so the vxAssignNodeAffinity had no effect
         // and the implementation falls back to the default target
}

注意:所有对图节点调用 vxSetNodeTarget 的操作都应在对该图调用 vxVerifyGraph 之前完成。

以下是设置节点亲和性的重要注意事项:

  • 如果未为节点指定目标,则默认目标是 CPU。
  • 如果设置目标失败(例如,如果节点不受给定目标的支持),则实现将回退到默认目标(CPU)。

异构执行的通用技巧

  • 正如在 追求性能 部分所讨论的,从单个内核的执行时间推断性能结论可能会产生误导。在大多数情况下,您分配给加速器的子图越大,通信成本就越能得到更好的摊销。
  • 通常,GPU 在大型图像上的性能更好。因此,如果工作量太小(执行时间 < 1ms),则应在 CPU 设备上运行图或融合内核。
  • 请注意,使用 GPU 目标会引入一次性开销(约几秒钟)来编译 OpenCL™ 内核。编译发生在 OpenVX* 上下文创建时,不影响执行。
  • 一个典型的起始策略是首先测试仅 CPU 和仅 GPU 的场景(第 9.2 节)。请注意某些节点不受特定目标支持的情况(请参阅“内核扩展”文档中的内核支持矩阵)。在这种情况下,唯一的方法是单独调度节点,并通过将子图(或图中的独立分支)调度到不同目标来寻找最佳分割。
  • 对于图中混合使用 CPU 和 GPU 目标的场景,建议尝试启用 GPU 磁贴(默认为 OFF)的选项。这可能会释放两个设备之间额外的(数据)并行性。

    $ export VX_CL_TILED_MODE=1

    对于仅 GPU 的场景,应将该选项肯定地还原为 OFF。

  • 建议进行性能分析(下一章)以确定应首先卸载到附加目标的“热点”节点。同时,为了最小化调度和其他运行时开销,通常更有效率地将一定大小的内核序列卸载到附加目标,而不是单独的内核。
  • 请注意,GPU 可能忙于其他任务(如渲染),同样,CPU 也可能负责处理常规的操作系统例程。
  • 设备性能可能会受到动态频率调整的影响。例如,同时运行长内核可能会最终导致一个或两个设备停止使用 Intel® Turbo Boost Technology。这可能会导致整体性能下降,甚至与单设备场景相比。
  • 同样,即使在仅 GPU 的场景下,高中断率和与主机的频繁同步也可能提高 CPU 的频率并降低 GPU 的频率。

Intel® VTune™ 工具支持

Intel® VTune™ 工具完全支持在运行 Ubuntu* 的开发系统上收集和查看性能数据。从常规的 Intel® VTune™ Amplifier XE for Linux* 开始。

Intel VTune 还提供 OpenVX API 跟踪。这对 CPU(及其传统的性能计数器,如函数中花费的时间)和 GPU(通过 Intel VTune 中的 OpenCL™ 支持)都是可能的,如下所示。SDK 附带了一些解释性能工作流程的示例,包括 VTune 中的分析。请参阅 Intel(R) Computer Vision SDK Sample Applications 部分。

在使用 Intel VTune 分析应用程序时,请确保为该工具设置所有环境变量(例如,指向 CV SDK 的变量)。

另一种选择是直接从用于运行 OpenVX 应用程序的设置下运行 Intel VTune,以便该工具完全继承环境变量。

Intel® Vtune™ 示例分析(针对 CPU 目标)

为了在 VTune 中获取 CPU 的完全限定内核名称,您应该设置以下环境变量:

$export VX_INTEL_ENABLE_VTUNE_INSTRUMENTATION=1

同样,在工具中,当选择分析类型时,请确保选中“Analyze user tasks, events, and counters”选项,以下是“Basic Hotspots”分析的示例:

Setting "Analyze user tasks, events, and counters" option for the VTune(TM)

考虑下面的简单 OpenVX* 图(有关详细信息,请参阅 Auto-Contrast SDK 示例)。

                             >|channel extract|->Y->|equalize_hist|-
                            /                                        \
RGB(in)->|convert|->NV12->|channel extract|->----U------------- ->|combine|->|convert|->RGB(out)
                            \                                        /
                             >|channel extract|->V-------------------

根据 VTune™ 的数据,在 CPU(Intel® Core™ 开发机)上运行 100 次迭代时的热点如下:

CPU time breakdown of the Auto-Contrast SDK Sample, as seen in "Bottom-up" view in Intel® VTune™. Notice that primary grouping is by the "Task domain"

将完全限定的 OpenVX* 内核名称,如 org.khronos.openvx.color_convert(也归类为 NV12->RGB 转换),与示例图中的特定(第一个)节点关联起来相对容易。在这个简单的图中,颜色转换节点几乎完全占据了执行时间。

请注意,“TBB node functor call”只是底层节点执行时间的聚合,因为运行时使用 Intel® TBB 并行地按块运行它们。这在另一个 VTune 视图中清晰可见:

CPU tiles, mapped to the (TBB) threads in the VTune Platform pane (heavily zoomed)

User Tasks”选项(如果未设置)将完全禁用图块信息的显示。用户任务分析是通过专用开关(参见上图)启用的。有关与图块执行相关的分析示例,请参阅 Video Stabilization Census Transform SDK 示例。

除非您有其他内核可以替换顶级热点(或者如果无法通过更改内核参数来减轻计算需求),否则一种可能的步骤是将顶级热点卸载到 GPU 目标(有关更多详细信息,请参阅 将单个节点调度到不同目标 部分)。

Intel® Vtune™ 示例分析(针对 GPU 目标)

Intel® VTune™ 为 OpenCL™ API 跟踪和相关的硬件计数器提供了丰富的支持。由于 GPU 的 OpenVX* 实现依赖于成熟的 OpenCL 堆栈,因此在工具中关联 API 相对容易。

具体来说,考虑上一章的图示例。当仅在 GPU 上执行时(请参阅 系统级设备亲和性 部分)并且在 VTune 下运行时,这是 VTune 生成的摘要:

GPU Hotspots in Intel® VTune™

在这里,您可以看到,与 CPU(上一节)类似,颜色转换是最耗时的,其次是直方图。下一个详细级别(包括设备上的实际执行时间、驱动程序开销、复制成本等)可以通过 OpenCL 分析和队列信息来收集。有关更多信息,请参阅 VTune™ 中 OpenCL 支持的通用培训。

OpenCL Queue for the GPU Executing the OpenVX Graph (“Platform” view in Intel® VTune™)

分析远程目标

VTune™ Amplifier for Systems 2017(Intel® System Studio 2017 的一部分)支持 Yocto* 目标的远程性能分析。在 https://software.intel.com/en-us/articles/vtune-embedded/#profiling_on_target 上可以找到一些关于 Yocto 启用的有用说明。请注意,对于基本指标,您可以跳过说明页面中的 SEP 构建步骤。

Intel® 计算机视觉 SDK 示例应用程序

Intel® CV SDK 附带示例,可在 Intel CV SDK 示例页面 https://software.intel.com/en-us/computer-vision-sdk-support/code-samples 在线获取。

  • OpenCL™ Custom Kernel in OpenVX* Sample 演示了一个实验性的 Intel 扩展,该扩展允许将常规 OpenCL 内核的代码用作 OpenVX* 用户内核。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/ocl_custom_kernel 目录下的 OpenCL™ Custom Kernel in OpenVX* Sample – Developer Guide README

  • Heterogeneous Basics with OpenVX* sample 使用 Targets API 将 OpenVX 图的部分调度到不同的计算单元(例如,CPU、GPU 或 IPU)。
  • Auto-Contrast Sample 通过直方图均衡实现简单的对比度归一化。该示例的重点是 OpenVX 开发的基本步骤。它还展示了与 OpenCV 的基本互操作性,OpenCV 用于加载/保存和显示图像。该示例还解释了使用 Intel® VTune™ 工具进行基本分析。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/auto_contrast 目录下的 Auto-Contrast OpenVX* Sample – Developer Guide README

  • CNN AlexNet Sample 通过在原始 Caffe* 数据上运行 Intel CV SDK 模型优化器工具来构建 AlexNet 的 OpenVX 图(使用 Khronos* CNN OpenVX* 扩展)。然后,该图用于对给定图像的区域执行推理,并报告前 5 个结果。

    注意:此示例仅针对 Linux*。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/cnn_alexnet 目录下的 CNN AlexNet OpenVX* Sample – Developer Guide README

  • CNN Custom Kernel Sample 通过在原始 Caffe 数据上运行 Intel CV SDK 模型优化器工具来构建 FCN8s* 模型的 OpenVX 图。生成的图用于对输入图像进行语义分割。FCN8s 具有 OpenVX 运行时默认不支持的某些层。该示例解释了如何在 OpenCL 中实现缺失的层并注册实现,以便模型优化器在解析 Caffe 模型时生成正确的调用。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/cnn_custom_kernel 目录下的 CNN Custom Layer for OpenVX* Sample – Developer Guide README

    此外,<SDK_SAMPLES_ROOT>/samples/cnn_custom_kernels_lib 包含更多内核示例(Yolo*SSD* 特定)。

    注意:此示例仅针对 Linux*。

  • Video-Stabilization Sample 展示了更高级的 OpenVX 功能,如延迟和用户节点。该示例实现了基于 Harris 角点的特征点检测和光流点跟踪的流水线。解释了用户节点的基础知识。OpenCV 用于从文件读取和显示输入视频、调试可视化,并最终显示结果。该示例还展示了在异构场景中的优势,其中一些节点被调度到 GPU,并讨论了性能,包括基本的 Intel® VTune™ 分析。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/video_stabilization 目录下的 Video Stabilization OpenVX* Sample – Developer Guide README

  • Census-Transform Sample 实现了一个著名的 CENTRIST 视觉描述符计算算法,该算法用于场景分类和对象识别任务。此示例介绍了 Intel OpenVX 高级图块扩展用于用户节点的代码示例。与之前的示例一样,它通过数据共享展示了与 OpenCV 的基本互操作性。OpenCV 用于从视频文件中读取数据和可视化结果。此外,该示例还实现了 CENTRIST 作为自定义 OpenCL 内核(请参阅 OpenCV 和 OpenVX* 参数的解释差异 主题),以展示其在 GPU 上运行的优势。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/census_transform 目录下的 Census Transform OpenVX* Sample – Developer Guide README

  • GStreamer*-OpenVX* Interoperability Sample 展示了如何开发使用 OpenVX 进行典型计算机视觉任务的 GStreamer 插件。OpenVX 代码封装在 GStreamer 插件模板处理函数中,成为更大的 GStreamer 媒体处理流水线的一部分。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/gstovx_plugin 目录下的 GStreamer* - OpenVX* Interoperability Sample – Developer Guide README

    注意:此示例仅针对 Linux*。

  • Face Detection Sample 演示了如何使用 OpenVX* 来实现基于分类器级联检测的人脸检测算法。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/face_detection 目录下的 Face Detection OpenVX* Workload – Developer Guide README

  • Motion Detection Sample 使用 OpenVX* 开发运动检测应用程序。具体来说,它实现了一个基于背景减法 MOG2、膨胀、腐蚀和连通分量标记的简化运动检测算法。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/motion_detection 目录下的 Motion Detection OpenVX* Workload – Developer Guide README

  • Lane Detection Sample 在 OpenVX 中实现了一个简单的车道检测流水线。该示例使用核心 OpenVX 视觉函数,并结合 Intel 的扩展(如 HoughLineP)来计算来自前视摄像机视频的车道边界。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/lane_detection 目录下的 Lane Detection OpenVX* Sample – Developer Guide README

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/hetero_basic 目录下的 Heterogeneous Basics with OpenVX* Sample – Developer Guide README

  • Kernel Enumerator Sample 是一个命令行实用程序,用于查询 OpenVX* 运行时中的内核扩展列表、支持的目标以及内核列表。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/kernel_enumerator 目录下的 README 文件。

  • Color Copy Pipeline Sample 是一个打印和成像流水线示例,除了 CPU 之外,还利用了 IPU 和 GPU,以展示产品的异构优势,并进行性能讨论,包括基本的 Intel® VTune 分析。

    有关更多详细信息,请参阅 <SDK_SAMPLES_ROOT>/samples/color_copy_pipeline 目录下的 Copy Pipeline OpenVX Sample – Developer Guide README

    注意:此示例仅针对 Linux*。

OpenVX* 性能提示

OpenVX* 的经验法则是通用“计算”优化适用,因为从优化角度来看,OpenVX 只是另一个用于计算的 API。因此,诸如提供足够的并行松弛度、避免冗余复制或过度的同步等通用技巧仍然适用。这些在异构场景中尤为突出,因为涉及与加速器的通信延迟。

有一些特定的操作可以优化 OpenVX 应用程序的性能。这些操作如下:

  • 尽可能使用虚拟图像,因为这可以解锁许多图编译器优化。
  • 尽可能优先使用标准节点和/或扩展而不是用户内核节点(它们充当内存和执行屏障,阻碍性能)。这为流水线管理器提供了更大的灵活性来优化图执行。
  • 如果您仍需要实现用户节点,请基于高级图块扩展(请参阅 Intel's Extensions to the OpenVX* API: Advanced Tiling 章节)。
  • 如果应用程序包含独立的图,请使用 vxScheduleGraph API 调用并行运行这些图。
  • 为调度器提供足够的并行松弛度,不要将工作(例如图像)分成太多小块。考虑内核融合。
  • 对于图像,使用适合应用程序精度需求的最小数据类型(例如,32->16->8 位)。
  • 考虑异构执行(请参阅 Heterogeneous Computing with Intel® Computer Vision SDK 章节)。
  • 您可以创建一个引用已分配内存的 OpenVX 图像对象(vxCreateImageFromHandle)。要实现与 GPU 的零拷贝,必须对外部分配的内存进行对齐。有关更多详细信息,请参阅 https://software.intel.com/en-us/node/540453
  • 注意(通常是禁止性的)vxVerifyGraph 延迟成本。例如,以一种不会在参数更新时触发验证的方式构建图。请注意,与输入图像的 Map/Unmap(请参阅 Map/Unmap for OpenVX* Images 部分)不同,设置具有不同元数据(大小、类型等)的新图像几乎肯定会触发验证,从而可能增加显著的开销。

Intel 的 OpenCV 中的高级人脸功能

Photography Vision Library (PVL) 是 Intel 开发的一套最先进的计算机视觉和成像算法。PVL 中的每个组件都在 Intel 架构上进行了性能和功耗优化。PVL 是 OpenCV 的一部分,随 Intel® CV SDK 一起提供。

PVL 中有几种计算机视觉算法,现在都通过 OpenCV API 公开:

  • 人脸检测:PVL 提供快速轻量级的人脸检测,准确度处于行业领先水平。可以自动检测和跟踪人脸,支持广泛的旋转角度和长距离。在 FDDB 基准测试中,它也显示了顶级的人脸检测率。
  • 眨眼检测:PVL 实时检测并索引人的眨眼。眨眼检测可用于拍照时触发拍摄(眨眼)或丢弃主体眨眼的快照。它还可以应用于驾驶安全管理系统,以提醒疲劳驾驶员。
  • 微笑检测:通过分析 PVL 中人脸的变化和运动,可以检测出人的微笑强度。微笑检测可应用于数码相机和手机,作为微笑自动快门机制,当主体微笑时自动触发拍摄功能。
  • 人脸识别:PVL 以行业领先的准确度提供实时高效的人脸识别。在 FRGC 基准测试中,它显示了顶级的人脸识别率。您可以使用 PVL 人脸识别算法按人对照片进行排序或浏览。它还可以用于生物识别认证。

FaceDetector

FaceDetector 类检测人脸特征,包括人脸矩形、眼睛位置、嘴巴位置、眨眼状态和微笑状态。可以有选择地检测人脸特征。您可以只检测人脸矩形,或者检测人脸矩形、眼睛位置和眨眼状态。人脸特征之间存在依赖关系。要检测眼睛位置或嘴巴位置,需要人脸矩形。要检测微笑状态或眨眼状态,还需要人脸矩形和眼睛位置。即使未指定,也会检测所需的人脸特征。

许多参数会影响人脸检测结果。根据您设置的参数,检测到的人脸数量可能会有所不同。当您尝试检测越来越多的人脸时,检测需要更长的时间。检测可以有两种模式:普通模式和跟踪模式。在普通模式下,您可以一次性获取所有检测到的人脸。在跟踪模式下,您需要多次运行检测才能检测到所有的人脸。这样,跟踪模式下的人脸检测相对非常快。如果您想检测静态图像中的人脸,应使用普通模式。但是,如果您想检测视频流中的人脸,通常建议使用跟踪模式。

  1. 首先,您必须创建一个 FaceDetector 类的实例。FaceDetector 没有公共构造函数。而是通过调用静态方法来创建实例:

    cv::Ptr<cv::pvl::FaceDetector> fd = cv::pvl::FaceDetector::create();

    FaceDetector 具有以下参数:

    名称 描述 默认值 最小值/最大值
    最大可检测人脸数 可检测人脸的最大数量 32 1/32
    RIP 平面内旋转 135 0/180
    ROP 平面外旋转 90 0/90
    眨眼阈值 用于评估眨眼的阈值 50 1/99
    微笑阈值 用于评估微笑的阈值 48 1/99
    最小人脸尺寸 最小人脸像素尺寸 64 32/无
    跟踪模式 跟踪模式

    每个参数都有一个 get/set 方法,某些参数的 min/max 值被声明为枚举类型。您可以参考头文件。

    以下示例演示了如何从图像中检测人脸地标。

  2. 要更改 FaceDetector 的设置,请为每个参数调用 set 方法:

    fd->setMaxDetectableFaces(15);<!--?rh-implicit_p ????-->
    fd->setRIPAngleRange(45);
    fd->setROPAngleRange(90);
    fd->setBlinkThreshold(80);
    fd->setSmileThreshold(70);
    fd->setMinFaceSize(50);
    fd->setTrackingModeEnabled(false); // normal mode
  3. 要检测人脸矩形,您需要一个 Face 类的向量。Face 类包含人脸地标和人脸分类的信息。人脸地标包含人脸矩形、眼睛和嘴巴的信息。此外,人脸分类还包含眨眼和微笑信息。所有信息都可以通过 FaceDetector 的检测方法获取。也可以由您自己手动设置信息。如果您有自己的人脸地标和分类算法,可以将其与 FaceDetector 和 FaceRecognizer 一起使用。

    std::vector<cv::pvl::Face> faces;

    FaceDetector 的所有方法都应该使用灰度图像,该图像只有一个通道且深度为 8 位。

    cv::Mat grayImg = cv::imread("image.jpg", cv::IMREAD_GRAYSCALE);
    fd->detectFaceRect(grayImg, faces);
  4. 如前所述,人脸特征之间存在依赖关系。如果您想获取眨眼状态或微笑状态,则应检测眼睛位置。

    for (uint i = 0; i < faces.size(); ++i)
    {
            fd->detectEye(grayImg, faces[i]);
            fd->detectMouth(grayImg, faces[i]);
            fd->detectBlink(grayImg, faces[i]);
            fd->detectSmile(grayImg, faces[i]);
    }

整合:人脸检测示例代码

以下演示了如何在跟踪模式下使用网络摄像头检测人脸地标:

VideoCapture cam(0);
Ptr<FaceDetector> fd = FaceDetector::create();
Mat frame, grayedFrame;
vector<Face> faces;
fd->setTrackingModeEnabled(true); // enable tracking mode
while (true)
{
  cam >> frame;
  cvtColor(frame, grayedFrame, COLOR_BGR2GRAY);
  fd->detectFaceRect(grayedFrame, faces);
  for (uint i = 0; i < faces.size(); ++i)
  {
    fd->detectEye(grayedFrame, faces[i]);
    fd->detectMouth(grayedFrame, faces[i]);
    fd->detectBlink(grayedFrame, faces[i]);
    fd->detectSmile(grayedFrame, faces[i]);
  }
  for (uint i = 0; i < faces.size(); ++i)
  {
    const Face& face = faces[i];
    Rect faceRect = face.get<Rect>(Face::FACE_RECT);
    Point leftEyePos = face.get<Point>(Face::LEFT_EYE_POS);
    Point rightEyePos = face.get<Point>(Face::RIGHT_EYE_POS);
    Point mouthPos = face.get<Point>(Face::MOUTH_POS);
    bool closingLeft = face.get<bool>(Face::CLOSING_LEFT_EYE);
    bool closingRight = face.get<bool>(Face::CLOSING_RIGHT_EYE);
    bool smiling = face.get<bool>(Face::SMILING);
  }
  if (cv::waitKey(5) > 0)
    break;
}

在最后一个 for 循环块中,您可以看到如何使用模板方法获取 Face 类的。您也可以使用 set 模板方法设置值。以下是您可以设置和获取的 Face 类的值:

枚举 类型 描述
FACE_RECT cv::Rect 人脸矩形。没有默认值。(类型:cv::Rect)
RIP_ANGLE int 人脸的平面内旋转角度。默认值为 0。
ROP_ANGLE int 人脸的平面外旋转(偏航)角度。
FACE_RECT_CONFIDENCE int 表示人脸与典型人脸有多接近的值。默认值为 100。
TRAKING_ID int 人脸的跟踪 ID。默认值为 -1。
LEFT_EYE_POS cv::Point 左眼位置。没有默认值。
RIGHT_EYE_POS cv::Point 右眼位置。没有默认值。
EYE_POS_CONFIDENCE int 表示眼睛位置与典型眼睛位置有多接近的值。默认值为 100。
MOUTH_POS cv::Point 嘴巴位置。没有默认值。
MOUTH_POS_CONFIDENCE int 表示嘴巴位置与典型嘴巴位置有多接近的值。默认值为 100。
CLOSING_LEFT_EYE bool 指示人是否闭上左眼。没有默认值。
CLOSING_RIGHT_EYE bool 指示人是否闭上右眼。没有默认值。
LEFT_BLINK_SCORE int 左眼得分。该值范围从 0 到 100。0 表示张开的眼睛,100 表示完全闭合的眼睛。没有默认值。
RIGHT_BLINK_SCORE int 右眼得分。该值范围从 0 到 100。0 表示张开的眼睛,100 表示完全闭合的眼睛。没有默认值。
SMILING bool 指示人是否在微笑。没有默认值。
SMILE_SCORE int 微笑得分在 0 到 100 之间。0 表示不微笑,100 表示完全微笑。没有默认值。

Face 类提供了其他方法来设置成员变量,当您只想调用 FaceDetector 的特定检测方法时,这些变量非常方便。

以下示例演示了如何在有来自其他面部检测算法检测到的人脸矩形时检测眼睛位置:

cv::Rect faceRect;
// the face rectangle acquired by your own algorithm
cv::Ptr<cv::pvl::FaceDetector> fd = cv::pvl::FaceDetector::create();
cv::pvl::Face face;
face.setFaceRectInfo(faceRect);
fd->detectEye(grayImg, face);

Face 类的 setFaceRectInfo() 方法有许多参数,如 RIP 角度、ROP 角度、置信度和跟踪 ID。如果您不知道这些值,将使用默认值。

FaceRecognizer

FaceRecognizer 类识别输入图像或帧中的给定人脸。要识别面孔,应首先构建内部面孔数据库。您可以将面孔注册到此类实例的内部数据库中。您也可以从内部数据库中注销已注册的面孔。已注册的面孔可以存储在文件中,您可以在创建实例时加载存储的文件。除非您明确将已注册的面孔存储到文件中,否则当实例的析构函数调用时,内部数据库将被删除。

人脸识别有两种不同的模式:普通模式和跟踪模式。在普通模式下,它始终尝试处理整个输入图像以识别面孔。在跟踪模式下,如果给定的人脸与前几帧的人脸相同并且已识别,则不会再次识别。如果性能是重要因素,建议使用跟踪模式。

  1. 您应该通过其静态工厂方法创建 FaceRecognizer 类的实例。FaceRecognizer 不提供任何公共构造函数。
    cv::Ptr<cv::pvl::FaceRecognizer> fr = cv::pvl::FaceRecognizer::create();

    如果您已经有一个保存的 DB xml 文件,您可以加载它,它包含了创建过程。

    cv::Ptr<cv::pvl::FaceRecognizer> fr = cv::Algorithm::load<cv::pvl::FaceRecognizer>("saved_db.xml");
  2. 您可以将人脸注册到此类实例的内部数据库中。要注册人脸,需要进行人脸检测以获取 Face 值。

    std::Mat imgGray;<!--?rh-implicit_p ????-->
    
    std::vector<cv::pvl::Face> faces;
    int personID = fr->createNewPersonID();
    ...
    //fill imgGray
    //need face detect process to fill faces
    ...
    fr->registerFace(imgGray, faces, personID);

    如果您想将人脸数据保存到外部 xml 文件,您应该在 registerFace() 中将第四个参数设置为 true。

    fr->registerFace(imgGray, faces, personID, true);
  3. 人脸识别的目的是获取已注册数据库中匹配的 personID。如果识别出人脸,则 personID 为正值。如果*不*识别任何人脸,则 personID-10000,表示未知人员。

    std::Mat imgGray;
    std::vector<cv::pvl::Face> faces;
    std::vector<int> personIDs;
    ...
    //fill imgGray
    //need face detect process to fill faces
    ...
    fr->recognize(imgGray, faces, personIDs);
  4. 您可以顺序识别摄像头或视频文件等连续帧中的人脸。在这种情况下,识别速度可能更重要。您可以为这种顺序输入启用跟踪模式。人脸检测器和人脸识别器都应启用跟踪模式。

    fd->setTrackingModeEnabled(true);
    fr->setTrackingModeEnabled(true);
    //need face detect process to fill faces
    ...
    fr->recognize(imgGray, faces, personIDs);
  5. 5. 当您想将人脸数据保存到外部文件时,可以使用 save() API。此 API 仅保存已注册且在 registerFace()saveToFile 参数设置为 true 的人脸。

    fr->save("save_db.xml");
  6. 设置跟踪模式下每帧最多可识别的人脸数。此值不得超过 MAX_SUPPORTED_FACES_FOR_RECOGNITION_IN_TRACKING (8)。默认值为 MAX_SUPPORTED_FACES_FOR_RECOGNITION_IN_TRACKING (8)

    fr->setMaxFacesInTracking(5);

整合:人脸识别示例代码

以下示例代码显示了使用摄像头输入的が基本用例:

VideoCapture cam(0);
Mat frame, grayedFrame;
Ptr<FaceDetector> fd = FaceDetector::create();
Ptr<FaceRecognizer> fr = FaceRecognizer::create();
vector<Face> faces;
Mat image = imread("my_face.jpg", cv::IMREAD_GRAYSCALE);
vector<int> personIDs;
// Register a face by one image with my face
fd->detectFaceRect(image, faces);
fr->registerFace(image, faces[0], fr->createNewPersonID(), true);
//turn on tracking
fd->setTrackingModeEnabled(true);
fr->setTrackingModeEnabled(true);
while (true)
{
    cam >> frame;
    cvtColor(frame, grayedFrame, COLOR_BGR2GRAY);
    // Detect faces first and then try to recognize the faces
    fd->detectFaceRect(grayedFrame, faces);
    if (!faces.empty())
    {
        fr->recognize(grayedFrame, faces, personIDs);
        for (int i = 0; i < personIDs.size(); ++i)
        {
            // recognized
        }
        if (waitKey(5) > 0)
            break;
    }
    fr->save("frdb.xml");
}

相关信息

有关 Intel® 计算机视觉 SDK 的要求、已知问题和限制,请参阅 Intel® 计算机视觉 SDK 版本说明文档。

有关更多信息,请参阅以下文档:

获取帮助和支持

您可以在在线服务中心提交问题:http://www.intel.com/supporttickets。使用“请求支持”,并指定产品 ID:computer vision。

使用此页面上的链接访问有关 Intel® 计算机视觉 SDK 的其他信息:

© . All rights reserved.