用OpenCL造GPU加速的AviSynth滤镜

这一季度有个叫《月曜日のたわわ》的动画片。原作是比村奇石的漫画,这个作者的漫画有一个特点,和其他用黑白来表现亮度的漫画不同,比村用的是蓝白色调表现。看到这个动画的时候,我在考虑,能不能用一种处理方式,把动画做成漫画那种色调的效果。

思路是很简单的,先拿比村奇石的漫画原作来分析,假设表示同一亮度的像素的 RGB 值是一样的,这样就能得到一个从 YUV 的 Y 通道到 RGB 的映射。结果大概是这样(Y 和 RGB 的范围都是 0 到 255):

R = Y > 85 ? ((Y - 85) / 255 * 340) : 0;
G = Y;
B = Y > 135 ? 255 : Y + 120;

然后要对每个像素都处理。也就是一个很大的循环,然后 1920x1080 个像素全部按照上面的算一次。只是解码然后转成 RGB 进行上面的处理再转回 YV12(不包含拿去编码),速度不到30帧每秒(没有使用 SIMD 指令集)。

因为处理过程每个像素互不影响,如果能以非常多的线程同时运行,那么可以大大加快速度。这种事情是GPU擅长的,所以就思考着用 OpenCL 做这样的运算,看看能有多块的速度。

OpenCL 原始接口实在是太麻烦了,看着就让人提不起干劲。于是往上搜了一圈,找了一个叫做 EasyCL 的库来。它对 OpenCL 原始接口进行封装,然后就能以简洁得多的方式来调用 OpenCL。EasyCL 在编译过程中需要 OpenCL 的头文件和库。到 Intel、Nvidia、AMD 网站看了一圈,感觉 SDK 都好大好麻烦。然而我在自己系统文件夹下看到了 OpenCL.DLL,它的导出函数也和头文件里能对应上。于是就到 GitHub 上直接找 KhronosGroup 发的 OpenCL 的头文件(这里:https://github.com/KhronosGroup/OpenCL-Headers),下下来以后编译能过了。LIB 文件就通过链接的时候报错信息中找不到的那些符号,自己造一个(用这个文章里提到的自制工具)。通过了链接以后,得到了 EasyCL.DLL 和 EasyCL.LIB。

然后利用以下简单的函数就能调用 OpenCL 进行运算了:

EasyCL::isOpenCLAvailable() 可以判断系统能不能调用 OpenCL
EasyCL::createForFirstGpu() 可以创建一个 CL 对象
EasyCL::buildKernelFromString() 可以从源代码字符串创建一个 CL 程序,需要指定入口函数名,Option 是干嘛用的不知道不过可以送一个空字符串进去
CLKernel::in() 可以送从内存拷到显存的参数进去
CLKernel::out() 可以送从显存拷到内存的参数进去,这两个送参数的函数,有些类型是不支持的。不过既然都是指针,可以转成它认的那种指针送进去(数组大小记得也要根据类型的变化而变化)
CLKernel::run_1d() 可以运行程序。两个参数分别是一共有多少任务,然后一次运行多少个任务。前者要是后者的倍数否则它要报错,一次最多能运行多少任务可以通过 EasyCL::getMaxWorkgroupSize() 来获取。

解释一下什么叫“一共多少任务、一次运行多少任务”。比如这个视频是 1920x1080 的,那么一共有 2073600 个像素点,每个像素处理一次的话就一共有 2073600 个任务。通过调用函数得知一次能运行 1024 个任务,那么就可以送 2073600 和 1024 作为参数到 run_1d 函数里去(因为 2073600÷1024 能整除、没有余数)。然后在 OpenCL 代码里用 get_global_id(0) 获取当前是第几个任务,根据任务的编号来决定要做什么任务。(在本文中,编号就是“第几个像素”)

这个库在出现错误的时候,会抛异常给你而不是返回错误码或者什么的。所以要做好异常捕捉,不然异常抛出来程序没捕捉、程序就崩了。

改成用 GPU 处理之后,处理速度从原来的每秒 30 帧不到变成了现在接近 90 帧。效果还是很明显的。

P.S. 从QQ群内聊天得知,微软有个 DirectCompute 能做类似的事情。还有个关键字是 C++ AMP,可以直接在编译时生成拿到 GPU 上跑的代码,需要 DirectX 11 支持。

发表评论

电子邮件地址不会被公开。 必填项已用*标注

您可以使用这些HTML标签和属性: <a href="" title=""> <abbr title=""> <acronym title=""> <b> <blockquote cite=""> <cite> <code> <del datetime=""> <em> <i> <q cite=""> <strike> <strong>