1. 为什么要在VS2017里折腾OpenCL加OpenCV这套组合
先把场景说清楚。你手上有一个图像处理或者计算机视觉的项目,跑在CPU上的OpenCV已经满足不了速度要求了,你想把一部分计算密集的环节丢给GPU去做。这时候有两条路:一条是直接用CUDA,另一条是走OpenCL。CUDA的问题在于它绑死了某个显卡厂商,换一台机器可能就跑不起来;OpenCL的好处是跨平台、跨厂商,CPU、GPU、甚至一些专用加速器都能当设备用。所以如果你的项目需要交付到不同硬件环境,OpenCL往往是更稳妥的选择。
那为什么还要把OpenCV拉进来?因为纯OpenCL写图像处理,从读图、格式转换、内存对齐到结果保存,每一步都要自己手写kernel,开发效率很低。OpenCV本身有大量的图像处理算法,而且它内部有一个叫做Transparent API(简称T-API)的机制,也就是UMat这个类。你只要把Mat换成UMat,OpenCV在底层会自动尝试调用OpenCL来加速,不需要你写一行kernel代码。这就形成了一个很舒服的分工:常规图像操作交给OpenCV的UMat自动加速,特殊算法自己写OpenCL kernel,两者共享同一个OpenCL上下文,数据不用来回拷贝。
听起来很美,但真正动手搭环境的时候,坑比想象中多。VS2017这个版本本身就有点尴尬——它不像VS2015那么老,也不像VS2019/2022那么新,很多库的预编译包对它的支持处于一个“能用但文档少”的状态。再加上OpenCL的SDK、OpenCV的编译选项、链接顺序、运行时dll路径这几件事只要有一件没弄对,就是一堆LNK2019或者运行时崩溃。我前后在不同机器上搭过好几遍这套环境,下面把完整的过程和踩过的坑都摊开讲。
这篇文章适合两类人:一类是刚接触GPU加速、想在Windows上快速把环境跑通的开发者;另一类是用过OpenCV但没碰过OpenCL、想搞清楚T-API背后到底怎么回事的人。我会尽量把每一步的“为什么”讲清楚,而不是只丢一堆配置截图。
2. 搭建前的整体思路与组件依赖关系梳理
2.1 三个核心组件各自扮演什么角色
在动手之前,先把这套环境里的三个主角的关系理清楚,不然后面配置选项一多容易晕。
OpenCL运行时(ICD Loader):这是操作系统层面的东西,负责在运行时找到机器上所有支持OpenCL的设备(显卡驱动、CPU运行时等)。它本身不是编译器,也不是SDK,就是一个“调度员”。Windows上它通常由显卡驱动安装,注册表里会记录各个厂商的ICD(Installable Client Driver)路径。
OpenCL SDK:这是给开发者用的,包含头文件(CL/cl.h等)、导入库(OpenCL.lib)以及一些工具。注意,SDK和运行时是两回事——你可以装了SDK但机器上没有可用的OpenCL设备,那样程序编译能过但运行时找不到设备。
OpenCV:图像处理库。关键在于它编译时是否开启了WITH_OPENCL选项。官方发布的预编译包(也就是那个exe自解压出来的)通常是带OpenCL支持的,但具体版本差异很大,有些版本默认关闭。这一点后面会专门讲怎么验证。
三者关系可以这样理解:OpenCV在运行时通过ICD Loader去枚举设备,如果找到可用设备且自己编译时开了OpenCL,就会启用T-API加速;你自己写的OpenCL代码则直接调用SDK提供的API,同样通过ICD Loader访问设备。所以理想情况下,OpenCV和你自己的kernel应该共享同一个上下文,避免数据在CPU和GPU之间反复搬运。
2.2 为什么推荐“自己编译OpenCV”而不是直接用预编译包
很多人图省事,直接下载OpenCV的Windows预编译包,解压后配置一下包含目录和库目录就用了。对于纯CPU的Mat操作,这样确实没问题。但一旦涉及OpenCL,预编译包有几个隐患:
第一,你无法确认它编译时WITH_OPENCL到底是ON还是OFF。虽然可以用cv::ocl::haveOpenCL()在运行时检测,但如果检测出来是false,你就得换包或者自己编译,反而更折腾。
第二,预编译包用的编译器版本、运行时库(MT还是MD)可能和你的项目不一致,链接时容易出现RuntimeLibrary不匹配的警告甚至错误。
第三,如果你后续想用OpenCV的ocl模块里一些高级功能,或者想调整OpenCL相关的编译宏,预编译包给不了你灵活性。
所以我的建议是:既然都走到OpenCL这一步了,索性用CMake自己编译一遍OpenCV。虽然编译一次要花不少时间(取决于机器性能,半小时到两小时都有可能),但编译出来的库和你的VS2017、你的运行时库完全匹配,后面省心得多。而且编译过程中你能清楚看到WITH_OPENCL这一项的状态。
2.3 版本选择的几个现实考量
版本这块没有绝对的最优解,但有几个经验性的建议。
OpenCV方面,3.4.x系列和4.x系列对OpenCL的支持都比较成熟。3.4.x的好处是稳定、资料多,很多老项目还在用;4.x的好处是新特性多,比如4.5之后对OpenCL的一些优化。如果你没有特别的版本依赖,选一个4.x的中期版本(比如4.5.x)比较稳妥,太新的版本有时候和VS2017的兼容性反而会有小问题。
OpenCL SDK方面,如果你用的是某家显卡厂商的SDK,注意它的版本要和驱动匹配。SDK版本太新而驱动太老,或者反过来,都可能出问题。一个比较保险的做法是:先确认显卡驱动是较新的版本,然后去厂商官网下载对应世代的SDK。如果你不想绑定特定厂商,也可以用通用的OpenCL头文件和库,但那样就少了厂商提供的调试工具。
VS2017本身,建议装最新的更新版本(15.9.x),早期版本在C++17支持和一些链接器行为上有差异,可能影响OpenCV的编译。
3. OpenCV源码编译:CMake配置里那些容易点错的选项
3.1 源码获取与目录规划
先从源码说起。去OpenCV的官方仓库或者发布页下载源码包,同时别忘了下载opencv_contrib(如果你需要那些额外的模块,比如aruco、xfeatures2d等)。即使你暂时不用contrib,一起编译进去也没坏处,免得以后想用还得重编。
目录规划上,我习惯这样安排:
D:\dev\ opencv-4.5.5\ <- 源码 opencv_contrib-4.5.5\ <- contrib源码 opencv-build\ <- CMake构建目录 opencv-install\ <- 安装目录(编译产物放这里)把源码、构建、安装三个目录分开,好处是构建失败时可以整个删掉opencv-build重来,不影响源码;安装目录独立出来,项目里引用路径也干净。
3.2 CMake关键选项逐条说明
打开CMake GUI,设置好source和build目录,点Configure,选择Visual Studio 15 2017 Win64。这里注意,一定要选Win64,除非你的项目确实是32位的。现在还在用32位环境的情况很少了,而且64位下内存寻址空间大,处理大图像不容易出问题。
第一次Configure会跑一段时间,然后出来一堆红色选项。下面这几个是必须关注的:
| 选项 | 推荐值 | 原因 |
|---|---|---|
WITH_OPENCL | ON | 这是核心,不开这个T-API就是摆设 |
WITH_OPENCL_D3D11_NV | 按需 | 如果你用NVIDIA显卡且需要D3D11互操作才开,一般关掉 |
WITH_OPENCLAMDFFT | OFF | AMD的FFT库,除非你明确要用 |
WITH_OPENCLAMDBLAS | OFF | 同上 |
BUILD_opencv_world | 按需 | 把所有模块打包成一个库,方便但体积大 |
OPENCV_EXTRA_MODULES_PATH | 指向contrib的modules目录 | 要用contrib就设 |
BUILD_SHARED_LIBS | ON | 动态库,方便更新;静态库也行但体积大 |
CMAKE_INSTALL_PREFIX | 你的install目录 | 编译完install的目标路径 |
WITH_OPENCL这一项如果Configure之后显示为ON,说明CMake在你的系统里找到了OpenCL的头文件和库。如果它自动变成了OFF,或者你手动打开后Configure报错说找不到OpenCL,那就说明SDK没装好或者路径没配好,得先解决这个。
3.3 编译过程中的资源消耗与时间预估
配置好之后点Generate,然后用VS2017打开生成的OpenCV.sln。编译的时候有几个点要注意。
第一,选择正确的配置。Debug和Release都要编的话,时间翻倍。如果你只做Release开发,可以只编Release。但建议至少把Release编出来,Debug版在调试OpenCL相关代码时有用。
第二,并行编译。VS里可以设置/MP选项,或者用cmake --build . --config Release -- /m这样的命令来多核编译。不过OpenCV的某些模块编译时内存占用很大,如果你机器内存不够(比如16G以下),并行度开太高反而会因为内存不足导致编译失败。我一般设成CPU核心数的一半左右比较稳。
第三,opencv_world如果开了,链接阶段会慢一些,因为要把所有模块合并。不开的话模块多,但每个模块编译快。
编译完成后,在VS里对INSTALL工程点“仅生成”,就会把编译产物拷贝到CMAKE_INSTALL_PREFIX指定的目录。最终你会得到include、x64/vc15/lib、x64/vc15/bin这样的结构。
3.4 验证OpenCL是否真的被编进去了
编译完别急着写代码,先验证一下。写一个最小的程序:
#include <opencv2/core.hpp> #include <iostream> int main() { std::cout << "OpenCL available: " << cv::ocl::haveOpenCL() << std::endl; if (cv::ocl::haveOpenCL()) { cv::ocl::Context ctx; if (ctx.create(cv::ocl::Device::TYPE_ALL)) { std::cout << "Device count: " << ctx.ndevices() << std::endl; for (int i = 0; i < ctx.ndevices(); i++) { cv::ocl::Device dev = ctx.device(i); std::cout << "Device " << i << ": " << dev.name() << std::endl; } } } return 0; }如果haveOpenCL()返回1,并且能列出设备,说明OpenCV的OpenCL支持是正常的。如果返回0,要么是编译时没开,要么是运行时找不到ICD。这时候可以用一个叫clinfo的小工具(OpenCL SDK里通常带)来确认系统层面能不能看到设备。如果clinfo能看到设备但OpenCV看不到,那问题多半出在OpenCV的编译配置上。
4. VS2017项目配置:从包含目录到链接顺序的完整设置
4.1 项目属性里的目录配置
新建一个空的C++控制台项目,然后右键项目属性。需要配置的地方集中在“VC++目录”和“链接器”两块。
包含目录里加上:
D:\dev\opencv-install\include D:\dev\opencv-install\include\opencv2库目录里加上:
D:\dev\opencv-install\x64\vc15\lib注意vc15这个目录名对应的是VS2017(MSVC工具集版本1910系列)。如果你用的是VS2019,目录名会是vc16。这个对应关系别搞错,用错了链接时会报找不到库。
4.2 附加依赖项与Debug/Release的区分
链接器的“输入”->“附加依赖项”里,Release配置下加:
opencv_world455.libDebug配置下加:
opencv_world455d.lib注意那个d后缀,Debug版库带d,Release版不带。如果你在Debug配置下链接了Release库,或者反过来,轻则警告重则崩溃。这个错误很常见,尤其是从别人那里拷贝项目配置的时候。
如果你没开opencv_world,那依赖项就要一个个列,比如opencv_core455.lib、opencv_imgproc455.lib等等。模块多的时候容易漏,所以除非有特殊需求,我一般建议开opencv_world省事。
4.3 运行时dll的放置策略
编译能过不代表运行能跑。opencv_world455.dll这个文件必须在运行时能被找到。有三种做法:
第一种,把dll所在目录(opencv-install\x64\vc15\bin)加到系统PATH环境变量里。好处是一劳永逸,坏处是如果你有多个OpenCV版本会冲突。
第二种,把dll拷贝到你的exe输出目录。简单直接,但每次清理重建都要重新拷。
第三种,在VS的项目属性里配置“调试”->“环境”,加上PATH=%PATH%;D:\dev\opencv-install\x64\vc15\bin。这样只在调试时生效,不影响系统环境。我比较推荐这种,尤其是同时维护多个项目的时候。
4.4 一个容易忽略的点:OpenCL库的链接
如果你自己的代码里直接调用了OpenCL的API(比如clCreateContext、clBuildProgram),那还需要链接OpenCL.lib。这个库通常在SDK的lib\x64目录下。把它加到库目录和附加依赖项里。
但这里有个细节:OpenCV内部已经动态加载了OpenCL运行时(通过ICD Loader),它不一定需要你显式链接OpenCL.lib。如果你只是用OpenCV的UMat,不直接调OpenCL API,那不链接OpenCL.lib也能跑。但如果你想自己写kernel,那就必须链接。
还有一个更隐蔽的问题:如果你链接了OpenCL.lib,但运行时找不到对应的dll(比如OpenCL.dll不在PATH里),程序启动时就会直接报错退出,错误信息可能是“找不到OpenCL.dll”。而OpenCL.dll通常由显卡驱动安装到系统目录,一般不会有问题,但如果你在虚拟机或者没有独显的机器上跑,就可能缺失。
5. 跑通第一个OpenCL+OpenCV程序:从UMat到自定义kernel
5.1 用UMat验证T-API加速是否生效
先写一个最简单的UMat例子,确认OpenCV的自动加速在工作:
#include <opencv2/core.hpp> #include <opencv2/imgproc.hpp> #include <opencv2/imgcodecs.hpp> #include <iostream> #include <chrono> int main() { cv::Mat img = cv::imread("test.jpg", cv::IMREAD_GRAYSCALE); if (img.empty()) { std::cerr << "Failed to load image" << std::endl; return -1; } // 用UMat做一次高斯模糊 cv::UMat uimg, uresult; img.copyTo(uimg); auto t1 = std::chrono::high_resolution_clock::now(); cv::GaussianBlur(uimg, uresult, cv::Size(31, 31), 0); auto t2 = std::chrono::high_resolution_clock::now(); std::cout << "UMat GaussianBlur time: " << std::chrono::duration_cast<std::chrono::milliseconds>(t2 - t1).count() << " ms" << std::endl; // 对比Mat版本 cv::Mat result; t1 = std::chrono::high_resolution_clock::now(); cv::GaussianBlur(img, result, cv::Size(31, 31), 0); t2 = std::chrono::high_resolution_clock::now(); std::cout << "Mat GaussianBlur time: " << std::chrono::duration_cast<std::chrono::milliseconds>(t2 - t1).count() << " ms" << std::endl; return 0; }跑一下,如果UMat版本明显快于Mat版本(在大图上差距会更明显),说明T-API在起作用。如果两者时间差不多,可能是图像太小、核太小,或者OpenCL根本没启用。可以用cv::ocl::setUseOpenCL(true)强制开启,再用cv::ocl::useOpenCL()确认状态。
5.2 自己写一个OpenCL kernel并与OpenCV共享上下文
这一步是很多人卡住的地方。OpenCV内部有自己的OpenCL上下文,你自己创建的上下文和它的不是同一个。如果两边各用各的上下文,数据在两者之间传递就要经过CPU内存中转,失去了加速的意义。
正确做法是复用OpenCV的上下文。OpenCV提供了cv::ocl::Context::getDefault()来获取默认上下文,以及cv::ocl::Device::getDefault()获取默认设备。你可以用这些来创建自己的cl_command_queue和cl_program。
#include <opencv2/core/ocl.hpp> #include <CL/cl.h> #include <iostream> int main() { if (!cv::ocl::haveOpenCL()) { std::cerr << "OpenCL not available" << std::endl; return -1; } cv::ocl::Context& ctx = cv::ocl::Context::getDefault(); cl_context clCtx = (cl_context)ctx.ptr(); cv::ocl::Device dev = cv::ocl::Device::getDefault(); cl_device_id clDev = (cl_device_id)dev.ptr(); // 创建命令队列 cl_int err; cl_command_queue queue = clCreateCommandQueue(clCtx, clDev, 0, &err); if (err != CL_SUCCESS) { std::cerr << "Failed to create queue: " << err << std::endl; return -1; } // 一个简单的kernel:把每个像素值加10 const char* src = R"( __kernel void addTen(__global const uchar* in, __global uchar* out, const int total) { int gid = get_global_id(0); if (gid < total) { out[gid] = in[gid] + 10; } } )"; cl_program prog = clCreateProgramWithSource(clCtx, 1, &src, nullptr, &err); err = clBuildProgram(prog, 1, &clDev, nullptr, nullptr, nullptr); if (err != CL_SUCCESS) { size_t logSize; clGetProgramBuildInfo(prog, clDev, CL_PROGRAM_BUILD_LOG, 0, nullptr, &logSize); std::string log(logSize, '\0'); clGetProgramBuildInfo(prog, clDev, CL_PROGRAM_BUILD_LOG, logSize, &log[0], nullptr); std::cerr << "Build failed:\n" << log << std::endl; return -1; } cl_kernel kernel = clCreateKernel(prog, "addTen", &err); // 准备数据 const int N = 1024 * 1024; std::vector<uchar> hostIn(N, 100); std::vector<uchar> hostOut(N, 0); cl_mem bufIn = clCreateBuffer(clCtx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, N, hostIn.data(), &err); cl_mem bufOut = clCreateBuffer(clCtx, CL_MEM_WRITE_ONLY, N, nullptr, &err); clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufIn); clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufOut); clSetKernelArg(kernel, 2, sizeof(int), &N); size_t globalSize = N; err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &globalSize, nullptr, 0, nullptr, nullptr); clFinish(queue); clEnqueueReadBuffer(queue, bufOut, CL_TRUE, 0, N, hostOut.data(), 0, nullptr, nullptr); std::cout << "First pixel: " << (int)hostOut[0] << std::endl; // 清理 clReleaseMemObject(bufIn); clReleaseMemObject(bufOut); clReleaseKernel(kernel); clReleaseProgram(prog); clReleaseCommandQueue(queue); return 0; }这段代码的关键在于ctx.ptr()和dev.ptr()这两个转换。OpenCV的ocl::Context和ocl::Device内部封装了cl_context和cl_device_id,通过ptr()可以拿到原始句柄。这样创建的buffer和OpenCV的UMat在同一个上下文里,理论上可以通过cl_mem和UMat之间的转换来共享数据,避免拷贝。
5.3 数据在UMat和cl_mem之间的传递
OpenCV的UMat内部有一个cl_mem句柄,可以通过UMat::handle()获取。但要注意,这个句柄的类型是ACCESS_RW还是ACCESS_READ取决于UMat的状态。如果你要自己写kernel处理UMat的数据,可以这样做:
cv::UMat uimg = ...; // 假设已经填充了数据 cl_mem clBuf = (cl_mem)uimg.handle(cv::ACCESS_RW);拿到cl_mem之后,就可以直接把它作为kernel参数传入,不需要额外的拷贝。但这里有个同步问题:OpenCV可能在后台异步执行了一些操作,你在用这个buffer之前需要确保之前的操作已经完成。OpenCV提供了cv::ocl::finish()来等待所有队列完成。
反过来,如果你有一个cl_mem,想把它包装成UMat,可以用cv::UMat的构造函数或者cv::ocl::convertFromBuffer。不过这个转换在不同OpenCV版本里API有差异,用之前最好查一下对应版本的文档。
6. 那些让我熬夜的报错:常见问题与排查路径
6.1 LNK2019无法解析的外部符号
这是最常见的链接错误。可能的原因有几个:
一是库目录或附加依赖项配错了。检查vc15目录是否存在,库文件名是否和你的OpenCV版本一致。有时候编译出来的是opencv_world455.lib,但你写成了opencv_world45.lib,少个5就找不到。
二是Debug/Release配置不匹配。Debug配置下链接了不带d的库,或者反过来。这个错误有时候不会直接报LNK2019,而是报一些奇怪的运行时错误,所以配置的时候要仔细。
三是运行时库选项不一致。OpenCV编译时用的可能是/MD(多线程DLL),而你的项目用的是/MT(多线程静态)。这会导致符号冲突。解决办法是在项目属性里把“C/C++”->“代码生成”->“运行时库”改成和OpenCV一致。自己编译OpenCV的话,可以在CMake里设置BUILD_WITH_STATIC_CRT为OFF来强制用/MD。
6.2 运行时提示找不到OpenCL.dll
这个错误通常发生在程序启动阶段。原因可能是:
机器上没有安装任何OpenCL运行时。这种情况在虚拟机或者只有集成显卡且驱动很老的机器上会出现。解决办法是安装对应硬件的驱动,或者安装一个通用的OpenCL CPU运行时。
OpenCL.dll不在PATH里。虽然它通常在系统目录,但如果你用的是某个SDK自带的OpenCL.dll,而那个目录没加到PATH,就会找不到。检查一下SDK的bin目录是否在PATH中。
32位/64位不匹配。你的程序是64位的,但PATH里只有32位的OpenCL.dll,或者反过来。用Dependency Walker或者dumpbin /dependents看一下exe依赖的dll架构。
6.3 clBuildProgram返回CL_BUILD_PROGRAM_FAILURE
这个错误说明kernel代码编译失败了。一定要把build log打出来看,不要只看错误码。常见的kernel编译错误包括:
语法错误,比如少了个分号、括号不匹配。OpenCL C的语法和C99很接近,但有一些限制,比如不支持函数指针、递归等。
用了设备不支持的特性。比如某些设备不支持double类型,或者不支持某些扩展。可以在kernel开头加#pragma OPENCL EXTENSION来启用扩展,但要先确认设备支持。
全局工作大小不是本地工作大小的整数倍。虽然很多设备能自动处理,但有些设备会报错。保险的做法是把global size向上取整到local size的倍数,然后在kernel里用边界检查。
6.4 UMat操作结果和Mat不一致
这个问题比较隐蔽。有时候用UMat算出来的结果和Mat有细微差异,尤其是在浮点运算上。原因可能是:
GPU的浮点精度和CPU不同。GPU通常遵循IEEE 754,但某些操作(比如sin、cos)的实现可能有细微差异。如果对精度要求极高,可能需要在kernel里用更高精度的算法,或者干脆用CPU。
OpenCL的快速数学选项。有些设备默认开启-cl-fast-relaxed-math,这会降低精度换取速度。可以在build options里关掉这个选项。
数据竞争。如果你自己写的kernel有多个work-item写同一块内存,结果会不确定。检查kernel里是否有写冲突。
6.5 性能没有提升甚至变慢
这是最让人沮丧的情况。可能的原因:
数据传输开销大于计算开销。如果你的图像很小,或者kernel计算量很小,那么数据在CPU和GPU之间搬运的时间可能比计算本身还长。解决办法是尽量让数据留在GPU上,多个kernel连续处理,最后再读回结果。
kernel没有优化好。比如全局内存访问没有合并、用了太多分支、work-group大小设置不合理等。可以用厂商提供的profiler工具来分析。
OpenCV的T-API没有真正启用。有些操作在特定条件下会回退到CPU。可以用cv::ocl::useOpenCL()和cv::ocl::Device::getDefault()来确认状态,也可以设置环境变量OPENCV_OPENCL_RUNTIME来查看详细日志。
7. 几个让开发效率翻倍的小习惯
7.1 用CMake管理自己的项目而不是手动配VS
虽然这篇文章讲的是VS2017里的配置,但我强烈建议你自己的项目也用CMake来管理。写一个CMakeLists.txt,用find_package(OpenCV REQUIRED)来定位OpenCV,用find_package(OpenCL REQUIRED)来定位OpenCL。这样换机器或者换OpenCV版本时,只需要改一下OpenCV_DIR和OpenCL_ROOT这两个变量,不用在VS属性页里一页页翻。
CMake还能自动处理Debug/Release的库名差异,你只需要写target_link_libraries(myapp ${OpenCV_LIBS} OpenCL::OpenCL),剩下的交给CMake。
7.2 把OpenCL kernel放在单独的文件里
不要图省事把kernel写成C++代码里的字符串。虽然那样编译时不用管文件路径,但kernel代码没有语法高亮、没有缩进、改起来很痛苦。更好的做法是把kernel写在.cl文件里,程序运行时读取文件内容再clCreateProgramWithSource。这样你可以用任何文本编辑器编辑kernel,甚至可以用一些支持OpenCL语法的编辑器。
如果担心运行时找不到.cl文件,可以在CMake里用configure_file把kernel文件拷贝到输出目录,或者直接把kernel内容编译进程序(用一些工具把.cl转成C数组)。
7.3 用事件来做性能分析而不是靠感觉
OpenCL提供了事件机制来精确测量kernel执行时间。在clEnqueueNDRangeKernel之后,可以创建一个事件,然后用clWaitForEvents等待完成,再用clGetEventProfilingInfo获取开始和结束时间戳。这样得到的时间比用std::chrono在CPU端测量准确得多,因为后者包含了命令入队和等待的时间。
记得在创建命令队列时加上CL_QUEUE_PROFILING_ENABLE标志,否则profiling信息不可用。
7.4 保留一个纯CPU的回退路径
不管GPU加速多诱人,永远保留一个纯CPU的实现作为回退。原因很简单:不是所有部署环境都有可用的OpenCL设备。如果你的程序在没有OpenCL的机器上直接崩溃,用户体验会很差。用cv::ocl::haveOpenCL()检测,如果没有就自动走CPU路径。OpenCV的UMat本身就有这个机制——如果没有OpenCL,UMat会自动在CPU上模拟,只是没有加速效果。但你自己写的kernel就需要手动处理了。
7.5 注意驱动更新带来的行为变化
显卡驱动更新有时会改变OpenCL的运行行为,比如之前能跑的kernel突然编译失败,或者性能突然下降。这在用厂商SDK时尤其常见。建议在项目里记录一下开发和测试时用的驱动版本,部署时尽量保持一致。如果必须升级驱动,升级后要重新跑一遍测试用例。
8. 从这套环境延伸出去还能做什么
环境搭好之后,能做的事情就多了。一个很自然的方向是把OpenCV的T-API和自定义kernel结合起来做流水线:比如用UMat做预处理(去噪、缩放),然后用自己的kernel做特征提取或滤波,最后再用UMat做后处理。整个流程数据都在GPU上,只在开始和结束时有两次传输。
另一个方向是探索OpenCV的dnn模块。OpenCV 4.x的DNN模块支持OpenCL后端,可以用GPU加速神经网络推理。虽然性能比不上专门的推理框架,但对于轻量级模型和快速原型来说很方便。
还有就是多设备协同。如果你的机器上既有集成显卡又有独立显卡,OpenCL可以把它们都枚举出来。你可以把不同的任务分配到不同的设备上并行执行。不过这涉及到多队列和多上下文的协调,复杂度会高不少,建议先把单设备跑顺了再考虑。
最后提醒一句:OpenCL的生态这些年变化挺大,有些厂商在逐渐把重心转向其他计算框架。但对于需要跨平台、跨厂商的场合,OpenCL仍然是目前最通用的选择。把这套环境搭好、把T-API和自定义kernel的配合摸清楚,这个技能在相当长一段时间内都用得上。