用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;

然后要对每个像素都处理。也就是一个很大的循环,然后 1920×1080 个像素全部按照上面的算一次。只是解码然后转成 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() 来获取。

解释一下什么叫“一共多少任务、一次运行多少任务”。比如这个视频是 1920×1080 的,那么一共有 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 支持。

VC++里4、8、12和16字节的成员指针

背景信息简单提一下……大概就是有两次被这事情坑了,我一直以为它只是一个普通指针的大小,32位下是4字节。

第一次,是在前一家公司的时候,遇到一个应用场景,一个头文件需要同时在C和C++里用,里面存有一个成员函数的指针,和void*一起放在union里,是专门给C++用的;C的话这个地方就用条件编译语句换成单个void*。结果是我离开项目组之后,后面用到这段代码的人发现这个成员指针有的时候比void*大。但是那个时候已经不在组里了,也不知道具体那边应用场景是啥样,最后简单把void*替换成char[32]了事,更多的也没在意。

第二次,是在现在这家公司的时候,有个同事在某个头文件里定义的类里面包含了另一个类的成员函数指针,结果这个头文件在两个不同的cpp文件里包含,这个类的大小是不一样的……访问成员变量的时候偏移也不同,数据就坏了。但是cpp里头文件包含的顺序调换一下就能好……同事说是那个成员函数指针的大小会不同,不过那会儿工作有点忙,结果是没放在心上,后来就忘了这事儿了。

最后一次是网友Advance在代码开了最高警告级别的时候,有一个成员指针被报没有对齐到16字节边界,然后就怀疑这个指针的大小。我突然想起之前同事的代码在不同文件访问变量所用的偏移不同的事情,就打算简单调查一下。

继续阅读

VC++中多重继承的指针转换

C++是支持多重继承的。不是像其他一些面向对象语言那样是“一个类实现多个接口”,而是真的多重继承。C++又是能从源代码生成机器码的语言,这种情况下,就要使得 在某个类的指针上调用对象的方法或访问成员变量(不管这个指针指向的就是这个类或者是这个类的子类)都能用同样的机器码实现,就会出现同一个对象,在用不同类型的指针表示它的时候,指针变量保存的地址会不同。它是把一个这样的对象分成几个部分,每个部分都是一个完整的父类型代表的对象,这样当父类型指向对应部分的时候,就能够完全像对待父类型的对象那样对待它。比如有类A和B,然后C同时继承了A和B,那么在C对象的内存空间里,就会有一块可以用A*指,还有另一块可以用B*指。 继续阅读

绿色守护治不了腾讯地图但阻止运行能

安卓手机上,我之前一直在用绿色守护来对付平时使用中不必要的后台程序。在屏幕关闭后几分钟,在绿色守护列表中的程序就会被自动关闭。这样它们就不能一直推送奇奇怪怪的消息给我,以及没有特别深刻体会的据说能够加速手机运行(因为后面跑的服务少了)。

但是在一次批量的程序升级之后,手机出现了奇怪的状况:时不时就振动一下,但是打开发现又没有短信或者通知或者电话。就是振一下、振一下。闹鬼似的。

继续阅读

服务器迁移

其实想要换一台网络状况好一些的服务器已经很久了,这事情想着想着拖着拖着,随着一次又一次给原来的服务器续费而计划终止。这几天突然又想回某不存在的140字符SNS网站上上玩了,正好这会儿服务器续费通知来了,又逢春节放假,会有时间折腾,就订购了一台在美国洛杉矶的KVM VPS。(之前问了客服,客服说在中国连过来比较快的是他们的洛杉矶机房服务器)

因为换了KVM,操作系统比起原来的OpenVZ能选的多了不少。我安装了NetBSD在上面。想玩玩/尝试这个系统挺长时间了。折腾了一晚上,服务器架起来原来博客的文件全部复制过来、数据库导过来,域名解析更新之后马上就能用了。体感上现在这个服务器要快不少。

配置服务器其实不应该要一个晚上的,结果NetBSD系统自带一个HTTPD(不是Apache),给搞混了,编辑了半天Apache的配置文件结果就是没反应(因为启动的是那个HTTPD)。最后发现问题所在的时候倒是很快,编辑即生效,也能看到日志。

博客都一年多没更新了,服务器日志里面还有RSS访问的记录,虽然不知道是谁不过确实给人继续写下去的动力:多多少少还是有人看对吧。虽然是就那么几个……

等原来的服务都弄过来了,下一步计划是搞新的域名,争取能够让谷歌搜索到。现在的域名可能哪里配置有问题,谷歌是访问不了的,抓取不到数据。

最近感觉工作中提升有限,晚上想自己学点什么。之前在玩的手游梅露可物语似乎越来越接近弃坑了,如果真弃坑了,至少每天晚场公会战的半小时是空出来了。不过不幸的是怎么又把灰姑娘捡回来玩了。

在上海工作、住、玩

在大学本科毕业之前,我到过一次上海。那次是第一次,主要还是因为之前根本就没有出过省(我是福建省的),然后毕业设计毕业答辩那边又没有问题了,就是等着拿证。因为之后是读研,所以也不急着工作那边的事情,于是就有了很多的空闲时间,然后就想趁着这个机会出去走走。

那次去了以后呢,感觉这地方好像不错(没见过世面吧,毕竟之前没出过省),然后又因为好多网友都在这里,作为一个平时和日常生活中的人交流不太多的人来说,也算是一个比较有吸引力的方面。所以现在就在上海找了工作,暂时先在这里。

继续阅读

x86逻辑地址转物理地址实验

之前一篇文章关于往NetBSD添加系统调用,是为了做这个实验准备的。

想要研究x86架构下的内存管理,最重要的参考文档应该是英特尔官方关于内存管理方面的解释了。在 Intel? 64 and IA-32 Architectures Software Developer’s Manual Volume 3A: System Programming Guide, Part 1 文档的第三个章节,详细解释了整个翻译过程,其中涉及到的寄存器和各种标志位的作用。长是比较长,而且又是英语的,但是说得很清楚。

而实验则是验证自己理解是不是正确的一个很好的方式。理论看了半天,也知道怎么算,套公式好像也能算得出来,但是却不知道有没有算对。有真实的环境,一测试一验证,就知道对不对了。如果自己理解有误,可以及早发现和纠正。

继续阅读

给NetBSD添加系统调用

NetBSD本来官方文档里( http://www.netbsd.org/docs/internals/en/chap-processes.html ,3.2.5节 )就有添加系统调用的相关指导的,但是实际根据那个指导操作的时候,会发现指导文档过期得有点严重,已经无法按照上面说的来达到目的了。然后我在其他地方找到了关于OpenBSD的教学( http://www.onlamp.com/pub/a/bsd/2003/10/09/adding_system_calls.html ),同样可以用在NetBSD上。选用NetBSD是因为它安装包小,内核编译又快,一来减少下载时间,二来减少编译时间。不过如果x86-32和x86-64都想玩过去的话,总下载量也还是会有七百来兆的:两片ISO(每个300多兆)和一个内核代码包(40兆)。

添加系统调用本来的目的是为了做“x86平台内存分页机制”实验的时候来执行一些特权操作的。这里用的系统是NetBSD 6.1.4 i386版本。

继续阅读

在*nix挂载Windows下的文件夹

最初是为了学习Unix程序设计的时候,能够更加方便地使用虚拟机:因为我将NetBSD系统安装在虚拟机中。这样的话虽然能够通过FTP等方式传文件,但总之不是那么方便。如果能直接在Windows下共享这个文件夹,然后虚拟机里面用mount挂载的话,就方便很多,文件啥的全就都在Windows下了,平时日常用的是Windows,所以对我来说是方便。

继续阅读

WDK编译boost(1.34.1)::regex

Windows Driver Kit里面的STL库啊啥的,基本是停留在VC6时代。不过也有一个好处,可以链接到msvcrt.dll、msvcp60.dll,因为这两个文件大家都有,所以可以减小最后生成的可执行文件的体积。然后比起VC6.0,WDK可以直接从微软那边下载来,而VC6直到最后也没像2005、2008、2010、2012、2013那样有个Express的免费版啥的。于是如果想自己弄了简单的小东西如果想要体积小、要链接到msvcrt.dll,感觉WDK是一个挺好的选择……

继续阅读