由于工作需要我需要分析Chrome的图片解码流程。众所周知，Chrome是建立在开源的Chromium项目上的。最近的一个多月时间里，我基本把所有心思都放在Chromium上了，而且不得不说，学习并分析开源项目的代码对一个程序员的提高确实蛮大的。这篇博文我会记录一下学习过程中我遇到的一些问题，并分享学习中我所参考的几篇优秀的Chromium代码分析文章。

我们先来看下面的代码 （来源于”12个有趣的C语言问答”）:

#include <stdio.h>  
  
int main(int argc, char *argv[])  
{  
    int flag = 0;  
    char passwd[10];  
  
    memset(passwd,0,sizeof(passwd));  
  
    strcpy(passwd, argv[1]);  
  
    if(0 == strcmp("LinuxGeek", passwd))  
    {  
        flag = 1;  
    }  
  
    if(flag)  
    {  
        printf("\n Password cracked \n");  
    }  
    else  
    {  
        printf("\n Incorrect passwd \n");  
  
    }  
    return 0;  
}  

这段关于密码验证的代码简单明了，除却argv[1]为空（用户没有给出参数）可能导致的程序崩溃问题，还有一个致命缺陷。

读者可以尝试编译以上代码并运行。当然，在参数为LinuxGeek时，控制台会正常输出显示密码正确；但是，问题出在参数的位数大于10的情况。

xp对此疑惑不解（直接说不懂C语言就好了嘛！），于是浏览此贴评论，才恍然大悟：

以下来自oschina用户vingzhang

“flag和passwd变量是存在栈上的，栈的扩展是从高地址到低地址。压栈的时候，先压int型的flag占四个字节，然后压passwd，占十个字节。举个例子，如果0x60fe0102表示passwd开始存放的地址，那么flag开始存放的地址是0x60fe010c，高10个字节。在intel x86的平台上，字节序是小端，也就是0x60fe010c标记的那个字节的存放是flag的最低位的一个字节数据，比如，flag的值为0xffffff01，那么01就存放在0x60fe010c标记的那个字节上。

攻击的目的，是让flag的值为非零，就可以绕过密码的验证，让打印为“ Passwd cracked ”。所以，只要输入的字节超过10之后的那四个字节任意一个非零就可以了。”

简单来说，用户使用strcpy函数时，如果输入的数据量大于了缓冲buffer数据的额定大小，就会覆写掉末端之后的数据。在上面的情况来说，覆盖掉了flag本来的值，而这个值不为0，于是如果输入长度大于10的时候，程序总会提示”密码已破解”。

通过命令行输出，我们可以更清楚的看出passwd和flag在栈上的地址：

printf("flag addr: 0x%x passwd[0] addr: 0x%x passwd[10] passwd: 0x%x \n", &flag, passwd, passwd + 10);  
//flag addr: 0x28ff0c passwd[0] addr: 0x28ff02 passwd[10] passwd: 0x28ff0c`

延伸阅读 StackSmashing

栈攻击，是攻击者有意的利用溢出栈缓存的手段达到访问隐藏的计算机内存的一种恶意攻击方法。这种手段非常恶劣。栈攻击通常基于C或者C++以下的特征来达到目的：

基于栈的语言将数据和返回值地址放在同样的栈上;

语言本身允许程序越过数组边缘读取储存在栈上的数据;

语言允许储存在栈上的数据被执行;

被攻击的程序有攻击者所需要的东西：例如，root权限;

攻击者将数据加载到把数据放到栈数组的程序。数据量过大，超过了程序开始设置的额定长度大小，但是程序本身没有检测输入数据的大小，于是多余的数据会覆盖掉数组结尾之后的数据，直到覆盖掉函数返回地址的值。这样攻击者就覆写了真实的返回地址，以自己的攻击程序代替。当这个函数返回时，函数并没有返回给它的发起者，而是直接返回到了储存在栈上的攻击代码位置，这样攻击者就能让这段程序执行这个程序所能允许的任何功能了;

栈攻击的威胁可以通过写一个有一个后台运行的权限受限的daemon程序来动态检查数组长度的语言来避免。像是C语言这种不检查数组越界问题的语言，你可以通过system call来显示的分配和释放页面文件作为缓存内存区来使用，这样对于这个缓存区所有的越界操作都会被硬件捕捉到。

参加OpenCV的OpenCL模块（以下称OCL）移植工作已经有2个月了。这里我说移植而不是开发，是因为大部分OCL模块的函数是从已经很成熟的GPU模块直接移植过来的。因此，目前阶段OCL模块所支持的函数接口是GPU模块的一个子集，但由于运行平台差别问题，在某些函数上有些细微不同。

OpenCV的版本控制系统已经转移到了git上面（见OpenCV on GitHub），而最新的trunk的master分支也正式加入了OCL模块。今天逛OpenCV的开发者社区时，我发现有人提问在OpenCV库中如何进行使用OCL模块的函数；回答问题的同时，考虑到网上还没有针对OpenCV的OCL模块的资料，我决定写一篇文章简单介绍下OCL模块以方便开发者使用。

今天在OpenCL的开发过程中遇到了对浮点数的原子运算(atomic operations)的问题。OpenCL spec中只提供了对于32位或64位整数的原子运算；对于浮点数，我们就得另辟蹊径了。

因为OpenCL在语法上跟CUDA非常类似，我们可以参考一下CUDA C Programming Guide上面关于浮点数原子加法的例子，如下：

__device__ double atomicAdd(double* address, double val)  
{  
  unsigned long long int* address_as_ull = (unsigned long long int*)address;  
  unsigned long long int old = *address_as_ull, assumed;  
  do {  
    assumed = old;  
    old = atomicCAS(address_as_ull, assumed,   
                    __double_as_longlong(val + __longlong_as_double(assumed)));  
  } while (assumed != old);  
  return __longlong_as_double(old);  
}

atomicCAS是cuda中一个把compare和swap组合起来的函数。对应的OpenCL函数是atom_cmpxchg。

有经验的读者会注意到cmpxchg也存在于Intel的汇编指令集，而这条指令常常用来实现琐无关的线程等待机制。 具体可参考：锁无关的(Lock-Free)数据结构——在避免死锁的同时确保线程继续。

转化成OpenCL中的内联函数，float版本：

inline void AtomicAdd(volatile __global float *source, const float operand) {  
    union {  
        unsigned int intVal;  
        float floatVal;  
    } newVal;  
    union {  
        unsigned int intVal;  
        float floatVal;  
    } prevVal;  
    do {  
        prevVal.floatVal = *source;  
        newVal.floatVal = prevVal.floatVal + operand;  
    } while (atomic_cmpxchg((volatile __global unsigned int *)source,   
                             prevVal.intVal, newVal.intVal)   
                             != prevVal.intVal);  
}  

对于乘法和除法，可以把其中关键运算的那一行替换

newVal.floatVal = prevVal.floatVal + operand;

替换为

AtomicMul(): newVal.floatVal = prevVal.floatVal * operand; //乘法  
AtomicMad(source,operand1,operand2): newVal.floatVal = mad(operand1,operand2,prevVal.floatVal); //乘后相加  
AtomicDiv(): newVal.floatVal = prevVal.floatVal / operand;  //除法  

不过，浮点数的原子运算效率非常低，所以实际应用中应尽量避免。 来源

1. OpenCL 1.1: Atomic operations on floating point values

2. OpenCLで浮動小数のatomic addをしたい