二值化程序终极优化

论坛>计算机科学>软件综合

文号

t18011

阅读

10927

二值化程序终极优化

小俊

2009/07/29

软件综合

IP:香港

优化前：

    int x, y;
    unsigned long offset;
    unsigned char r, g, b, mi, ma;

    for(y = 0; y < HEIGHT; y ++)
        for(x = 0; x < WIDTH; x ++)
        {
            offset = y * WIDTH + x;

            b = in[offset * 3];
            g = in[offset * 3 + 1];
            r = in[offset * 3 + 2];
            mi = min(r, min(g, b));
            ma = max(r, max(g, b));

            if(((unsigned short)ma + (unsigned short)mi) > THRESHOLD * 2)
                out[offset] = 255;
            else
                out[offset] = 0;
        }

优化后：

    int x, y;
    unsigned long offset = 0, p;
    unsigned char r, g, b, mi, ma;

    for(y = 0; y < HEIGHT; y ++)
        for(x = 0; x < WIDTH; x ++)
        {
            p = ((unsigned long *)in)[offset];

            b = ((unsigned char *)(&p))[0];
            g = ((unsigned char *)(&p))[1];
            r = ((unsigned char *)(&p))[2];

            mi = min(r, min(g, b));
            ma = max(r, max(g, b));

            if(((unsigned short)ma + (unsigned short)mi) > THRESHOLD * 2)
                out[offset] = 255;
            else
                out[offset] = 0;
            offset ++;
        }

优化方式：

1、把原来每个点占3个字节（R, G, B）的存储方式改为4个字节存储方式，这样的话访存时可以对齐。
2、用一个long型变量一次取一个点，再分给R, G, B，避免多次访存操作。
3、offset那里稍改了一下，其实对性能没什么影响。

+1000

科创币

破93

2009/07/29

我屈服 = =

来自：

计算机科学/软件综合

小俊

作者

17年6个月前IP:未同步

133129

1楼

优化前：

1280×800：5.8ms
3264×2448：42.6ms

1280×800（CUDA@GT200）：0.27ms
3264×2448（CUDA@GT200）：1.73ms

优化后：

1280×800：5.5ms
3264×2448：41.2ms

1280×800（CUDA@GT200）：0.16ms
3264×2448（CUDA@GT200）：0.83ms

引用

小俊

作者

17年6个月前IP:未同步

133132

2楼

CPU由于二级缓存比较大，访存不对齐以及多次字节型访存操作其实多数情况都会转化到访问二级缓存，优化后节省的访存周期不多，而且从3字节存储改为4字节存储后，访存的空间增大了，所以优化有限。

GPU没有二级缓存，访存不对齐会严重影响其性能，所以尽管优化后需要访问的空间比原来更多，但反而获得翻倍的性能。

引用

darkstorm

17年6个月前IP:未同步

133160

3楼

[s:252] 不错不错....
CUDA强大
40字节40字节40字节40字节

引用

chiataimakra 十步芳草

17年6个月前IP:未同步

133374

4楼

考虑相当细致,赞扬,虽然我不会用CUDA但是我觉得这及其符合我的习惯(奔三上面用万恶的VB做几千次循环的时候那个慢啊,93你拿着E系列处理器是可以理解但是无法体会的),优化变量之间互相读取的关系与减少“倒脚”次数。

干嘛不来当斑竹呢。。。

引用

小俊

作者

17年6个月前IP:未同步

133386

5楼

引用第4楼chiataimakra于2009-07-29 20:48发表的 :
考虑相当细致,赞扬,虽然我不会用CUDA但是我觉得这及其符合我的习惯(奔三上面用万恶的VB做几千次循环的时候那个慢啊,93你拿着E系列处理器是可以理解但是无法体会的),优化变量之间互相读取的关系与减少“倒脚”次数。

干嘛不来当斑竹呢。。。

1L贴的是用普通CPP实现的方法，CUDA的程序我没贴上来不过也差不多这个样子。

还有一个更进一步优化的想法，明天用普通的CPP和CUDA各实现一下试试，看看效率能提高多少。

虽然这个RGB2L二值化算法非常简单，但是往往是越简单的算法，执行的数据量越多，优化后性能提升越明显。例如，有统计表明，MPEG解码过程中，花费时间最多的步骤，不是huffman解码、预测帧恢复这些较复杂的部分，反而是内存拷贝、色度空间转换（YUV转RGB）、IDCT这些非常简单的步骤，因为这些步骤处理的数据量是整个解码过程中最大的。

引用

93°

17年6个月前IP:未同步

133389

6楼

引用第4楼chiataimakra于2009-07-29 20:48发表的 :
考虑相当细致,赞扬,虽然我不会用CUDA但是我觉得这及其符合我的习惯(奔三上面用万恶的VB做几千次循环的时候那个慢啊,93你拿着E系列处理器是可以理解但是无法体会的),优化变量之间互相读取的关系与减少“倒脚”次数。

干嘛不来当斑竹呢。。。

我一秒做几十次的几十万次for循环，怎么会体会不到呢。。

引用

小俊

作者

17年6个月前IP:未同步

133396

7楼

引用第6楼破93于2009-07-29 21:06发表的 :

我一秒做几十次的几十万次for循环，怎么会体会不到呢。。

我用CUDA的话没有循环，不过相当于开了几百万个线程……

引用

93°

17年6个月前IP:未同步

133400

8楼

还好 = = 我的手指追踪好多东西要算，用指针还是比较快的。。

引用

novakon

17年6个月前IP:未同步

133405

9楼

本人言论纯属胡扯

引用第8楼破93于09-07-29 21:14发表的 :
还好 = = 我的手指追踪好多东西要算，用指针还是比较快的。。

针使织者,必得线法操之。法不正,则针织不速。

引用

rc5

17年6个月前IP:未同步

133649

10楼

b = ((unsigned char *)(&p))[0];
g = ((unsigned char *)(&p))[1];
r = ((unsigned char *)(&p))[2];
*********************************
要终极，这三个是不是应该写成:

unsigned char *tmp = (unsigned char *)(&p);
b = *tmp++;
g = *tmp++;
r = *tmp;

引用

小俊

作者

17年6个月前IP:未同步

133662

11楼

引用第10楼rc5于2009-07-30 08:57发表的 :
b = ((unsigned char *)(&p))[0];
g = ((unsigned char *)(&p))[1];
r = ((unsigned char *)(&p))[2];
*********************************
要终极，这三个是不是应该写成:
.......

两者区别不大，类型转换仅对编译器有意义的，在汇编上都是直接从内存取数的指令。

引用

rc5

17年6个月前IP:未同步

133697

12楼

谈到汇编就应该写成:

b = (unsigned char) p;
g = (unsigned char) (p>>8);
r = (unsigned char) (p>>16);

完全是寄存器操作。

引用

小俊

作者

17年6个月前IP:未同步

133699

13楼

引用第12楼rc5于2009-07-30 15:55发表的 :
谈到汇编就应该写成:

b = (unsigned char) p;
g = (unsigned char) (p>>8);
r = (unsigned char) (p>>16);
.......

都是一样的。你看看反汇编：

p、r、g、b都是在寄存器里的，用指针跟用移位是一个效果。

引用

rc5

17年6个月前IP:未同步

133702

14楼

编译器强了，人就懒了，哈哈。

要是俺的话， in和out两个缓冲区肯定是用两个unsigned long指针，r,g,b是三个unsigned char指针，在for循环外就指到p上面各字节，这样循环里除了给两个缓冲区指针++外和给p赋值，就不会出现r,g,b的任何赋值

引用

小俊

作者

17年6个月前IP:未同步

133728

15楼

引用第14楼rc5于2009-07-30 16:14发表的 :
编译器强了，人就懒了，哈哈。

要是俺的话， in和out两个缓冲区肯定是用两个unsigned long指针，r,g,b是三个unsigned char指针，在for循环外就指到p上面各字节，这样循环里除了给两个缓冲区指针++外和给p赋值，就不会出现r,g,b的任何赋值

r,g,b本来就已经在寄存器中咯。就算用指针，运算时也是要从内存读到寄存器中。

引用

小俊

作者

17年6个月前IP:未同步

133730

16楼

顺便贴上CUDA版的优化代码：

#define THREAD_N    256
#define THRESHOLD    127

__global__ static void binarizeKernel(unsigned char *in, unsigned char *out)
{
    const unsigned long offset = (blockIdx.x * THREAD_N + threadIdx.x);
    unsigned long p = ((unsigned long *)in)[offset];

    unsigned char b = ((unsigned char *)(&p))[0];
    unsigned char g = ((unsigned char *)(&p))[1];
    unsigned char r = ((unsigned char *)(&p))[2];

    unsigned char mi = __min(r, __min(g, b));
    unsigned char ma = __max(r, __max(g, b));

    out[offset] = (((unsigned short)ma + (unsigned short)mi) > THRESHOLD * 2) ? 255 : 0;
}

优化后性能提高超过1倍。

引用

rc5

17年6个月前IP:未同步

134244

17楼

再琢磨琢磨能不能用64位来玩。

还有一终极办法-查表：
显然lz是输入2^24个值，输出0和FF两种值，
所以建立一个2^24位的表，0，1代表0:ff, 查一下就搞定了。
这个表也不过2^24/8= 2MB多一点。

估计速度还能翻一翻

引用

小俊

作者

17年6个月前IP:未同步

134289

18楼

查表不是一个好方法，尤其是在CUDA上。
查表会占用内存带宽，而且这是离散的访存，会造成很长的latency，这时CPU或GPU都不知道能够计算多少次min、max咯。
如果算法比较复杂的话就适合用查表。

引用

rc5

17年6个月前IP:未同步

134300

19楼

还能优化一点，你的那两个min和max要做4次if和4次赋值
如果你自己用if在三值中同时找最大和最小，应该是2-3个if和3次赋值就可以了。

当然代码就长了，还要看编译器有没有本事把4个min,max优化成同样效果

引用

小俊

作者

17年6个月前IP:未同步

134436

20楼

用CUDA的话，GPU本身有最值运算指令。
CPU的话我要去看看反汇编。

引用