樟树林

访问global虽然仍然需要最多数百个时钟周期，但是总的来说比cpu的内存还是快多了...如果使用得当的话。CUDA是符合SIMD-PRAM模型的，也就是n台功能相同的处理机（block看成处理机,每个thread看成处理机的一部分更好些），一个容量无限大的共享处理器M。其中M就是global memory，虽然在CUDA中也是有限的，但是比起可怜的shared还是多很多了。

 

global作为共享存储器，最重要的功能就是用来进行block之间的通信。这里能不能叫做通信其实有问题，因为block之间实际上没有通信...CUDA上的通信实际上就是刷新，运行完一个kernel，把所有的需要与其他block交换的数据都写进global，然后再launch一个kernel读入上次生成的数据...如此往复循环。

如果kernel的输入与输出是同一个矩阵，就是一个原地（in-place）操作，下个kernel还是读取这个矩阵。

如果输入矩阵与输出矩阵不同，而且需要反复操作这两个矩阵，就构成了一个乒乓操作，例如第一次kernel使用ping为输入，pong为输出，下次就用pong输入，ping输出。这种方法在迭代等数值计算中经常使用。

 

texure也是全局存储器，速度比global还要更快，但是我不推荐大家在GPGPU中使用texture。这是因为texture是只读的，array中的数据不随对应的矩阵改变而改变，需要反复绑定，而绑定占用的时间不一定比使用texture带来的提高小。大多数情况下，使用global + shared比使用texture要快。当然，在需要大块的只读数据如查找表，索引表时可以使用texture。

 

关于bankconflict

简单的说，矩阵中的数据是按照bank存储的，第i个数据存储在第i%16个bank中。一个block要访问shared memory，只要能够保证以其中相邻的16个线程一组访问thread，每个线程与bank是一一对应就不会产生bank conflict。否则会产生bankconflict，访存时间成倍增加，增加的倍数由一个bank最多被多少个thread同时访问决定。有一种极端情况，就是所有的16个thread同时访问同一bank时反而只需要一个访问周期，此时产生了一次广播。

 

下面有一些小技巧可以避免bank conflict 或者提高global存储器的访问速度

1. 尽量按行操作，需要按列操作时可以先对矩阵进行转置

2. 划分子问题时，使每个block处理的问题宽度恰好为16的整数倍，使得访存可以按照 s_data[tid]=i_data[tid]的形式进行

3. 使用对齐的数据格式，尽量使用nvidia定义的格式如float3,int2等，这些格式本身已经对齐。

4. 当要处理的矩阵宽度不是16的整数倍时，将其补为16的整数倍，或者用malloctopitch而不是malloc。

5. 利用广播，例如s_odata[tid] = tid%16 < 8 ? s_idata[tid] : s_idata[15];会产生8路的块访问冲突

而用：s_odata[tid]=s_idata[15];s_odata[tid]= tid%16 < 8 ? s_idata[tid] : s_data[tid]; 则不会产生块访问冲突

 

在使用CUDA进行GPGPU计算时，global + shared的黄金组合在速度上远远超过了texture，只有在以下两种情况下使用texture：
1.需要图像输出时
2.需要反复随机访问的，但内容不变的大块内容，如索引表，查找表
其他时候尽量使用global + shared

block、thread划分的简单原则：

用不同的block处理完全不相关的数据可以获得最好的性能。此时只需要在block内进行数据交换。例如可以用每个block处理一个独立的矩阵，或者大型矩阵中的一部分（如一行，或者一个子矩阵）。需要处理一列数据时尽量将其转化为按行存储后再处理，以避免bank conflict。

CUDA的thread是一种轻量级线程，比较适合处理小粒度问题，在block一级也仍然只适合处理中小粒度的问题。而使用MPI构造的系统，一般是在比较大的粒度上进行并行操作，以减少数据交换。因此即便是已经经过检验的传统并行算法，在CUDA上移植时也要注意是否有可能进行进一步的划分。

block与thread的划分原则与可以并行处理的最小子问题占用的存储器大小和子问题的数量都有关系。不同尺寸的同一问题，block和thread划分方法往往不同。例如子问题占用存储器资源有以下几种情况：

在同一block的shared memory可以储存多个并行子问题的数据，此时可以尽量用一个block处理相邻的并行子问题，或者若干个子问题合并成一个稍大的子问题。

每个block的shared memory只能储存一个并行子问题的数据，此时就是简单根据输入数据算输出

每个block的shared memory不足以处理一个并行子问题，此时可以使用texture和global作为存储器，或者将子问题进一步分拆为数个串行过程，用若干个kernel完成。
每个block的thread数量以192-256为宜，处理的矩阵或者图像尺寸最好长宽都是16的整数倍，如果可能尽量将其补成16的整数倍。

 

我们这里尽量只介绍使用CUDA进行计算的方法，而少涉及具体的算法。

核内的几种操作:

缩减问题（reduction），推荐使用二叉树的方式来做。
例如：缩减加操作为可以用以下代码描述：
#define ADD     if (tid < 128) { temp[tid] += temp[tid + 128];} __syncthreads();\
    if (tid <  64) { temp[tid] += temp[tid +  64];} __syncthreads();\
    if (tid <  32) { temp[tid] += temp[tid +  32];} __syncthreads();\
    if (tid <  16) { temp[tid] += temp[tid +  16];} __syncthreads();\
    if (tid <   8) { temp[tid] += temp[tid +   8];} __syncthreads();\
    if (tid <   4) { temp[tid] += temp[tid +   4];} __syncthreads();\
    if (tid <   2) { temp[tid] += temp[tid +   2];} __syncthreads();\
    if (tid <   1) { temp[tid] += temp[tid +   1];}\

基于同样的树状结构进行缩减操作，还可以进行排序，累乘，寻找最大最小值，求方差等等

 

邻域问题：
在图像处理中，经常可以遇到邻域问题，如4-邻域，8-邻域等。基本思想就是使用图像中的几个元素按照一定的算子计算输出矩阵中的一个元素
常见的8-邻域算子如:

-1 -1 -1   1 1 1
-1  8 -1   1 1 1
-1 -1 -1 高通滤波器  1 1 1 低通滤波器
等等，可以完成锐化、平滑、去噪，查找边缘等功能
cuda的imagedenoising例子使用的是texture，其实也可以使用shared memory来做。
基本思路是：
1.在原有图像周围扩展像素，向外扩展像素数由邻域决定，如4-邻域和8-邻域扩展1个像素，24-邻域扩展2个像素等
2.每个block处理输出图像中长宽为16x16的一小部分，每个thread处理输出图像中的1个像素。这样每个block的线程数是256，符合nv推荐的192-256的范围。

texture的版本可以直接参考imagedenoising
shared memory的版本的可以参考以下代码（仅作示例，不能运行，处理4-邻域问题，其中i_data为global中的原始图像,func代表使用算子计算的过程）

shared type s_data[18][18];
shared type o_data[16][16];

s_data[tid.x + 1] = i_data[bid.x*16+bid.y*640*18+tid.x-640];
s_data[18*17 + tid.x + 1] = i_data[bid.x*16+bid.y*640*18+tid.x];
s_data[18+tid.x*18] = i_data[bid.x*16+bid.y*640*18+tid.y*640-1];
s_data[35+tid.x*18] = i_data[bid.x*16+bid.y*640*18+tid.y*640+1];
s_data[(tid.x+1)*18+(tid.y)+1]=i_data[bid.x*16+bid.y*640*18+tid.y*640+tid.x]；//中间256个

o_data[tid.x*16+tid.y]=func(
s_data[tid.x*18+tid.y+1],//邻点1
s_data[(tid.x+1)*18+tid.y],//邻点4
s_data[(tid.x+1)*18+tid.y+1],//中心点
s_data[(tid.x+1)*18+tid.y+2],//邻点2
s_data[(tid.x+2)*18+tid.y+2],//邻点3)；

 

投影问题：
这里的投影问题特指通过投影进行维度压缩的问题，如医学成像中常见的radon变换（当然，医学影像处理中更常见的是radon逆变换，但逆变换用傅式变换处理更好，此处不再赘述）。

这一类问题有一些例子：给手照x光片，我们的手是三维的，而最后出来的x光片则是平面的。做核磁共振（CT）时，实际上传感器是一个线阵，一直绕着我们的身体旋转，在不同的角度将我们身体的一个二维的截面投影到一维的传感器上。

以CT为例，实质上是将输入矩阵进行旋转后投影到一行上。设图像与传感器夹角为theta，可以根据tid.x * sin(theta) + tid.y*cos(theta)求出点（tid.x, tid.y）在线阵上的投影。此时结果与输入的数据是一对多的关系，而且随着角度不同，很难确定是由哪些点决定最终输出。
此时可以采用这样的方法：按照输入组织数据，每个block处理一行或者一个子矩阵并各自输出结果，最终将各block结果进行累加。

用这样的方法解决CT造影问题，将输入图像按行输入各个block，各block将旋转后的行存入一个长度为floor(1.414*max(width,height))的行向量
然后将各行输出的结果累加，即可得到最终结果。

 

快进教研室了，工作的同志们也有的开始准备跳槽了。总之，新阶段要开始了。

一上QQ，大家的签名一片愁云惨雾。很多人很迷茫，很多人很郁闷，很多人的母亲也在母亲节附近病倒了。衷心希望大家能够挺过去，家庭成员也都健康。

 

学校这个环境里还是有些资源可以利用的。朋友不多，更多的依靠的是工作伙伴，结果我的情报网已经因为这一届大四和研三的毕业而即将瘫痪，挺讨厌的感觉。

上研究生以后交往圈子小了很多，只是和外面公司有些来往了（仍然见面不多，有几次全是网上联系银行汇款，双方都没见过面）。明明已经好一些了，可是和不大熟的人打招呼的时候还是没法挤出一个象样的微笑。难道比以前更加自闭了么？

自闭也许还不算最可怕的，可怕的是居然对现在的状态无所谓。

既然无所谓，反过来也就没什么可怕的了。

 

但是总是会有些不方便吧。

 

扯远了一点，要进教研室了，也许应该停掉一些打工的工作吧。打工有点鸡肋的感觉，如果我知道现在多赚两千一月以后就少赚三千一月，我当然会义无反顾地不干了，但我不知道，甚至说不定以后我还有可能去干互联网。

 

前途是迷茫的，梦想有可能也是愚蠢的。

 

现在的我两年之内是能够造一枚导弹或者无人机之类的东西的。有些算法和技术实现做出来以后也比较有价值，有些也有变现的潜力，但是得有人给我钱，教研室支持才行。如果干的话，说不定还要把家里的东西也扯过来。总之，这件事很复杂，还要做一些调研工作。假如真的按照自己的来，想必飞机导弹什么的做不出来，工程能力上也能值比一个普通硕士多一些的钱。就是挺麻烦，必须和很多人打交道，而且前途不甚明朗（虽然比在教研室跟着老板干方向要明确一些，干不成最差也是一个普通的硕士）。申请基金支持，要钱也是有可能要到的，但是很烦，而且我通常只是一个人，只有很少的几个人能够提供非常宝贵的帮助（在此鸣谢）。所以如前所述，人缘不好还是不方便。

 

不过，有事没事自己造架飞机导弹这种事情到底也只有我吧？

 

到底怎么办，上不上，还得继续调研，深入调研，回家调研，找老板谈话，plapla

只靠自己还是有风险的，必须偏执，绝对自信，不能被击垮，否则后果就是灾难性的。如果不是pw同志那时候扶了一把，当初也许已经倒下了。

 

有一天我会倒下，但还不是现在，因为生命流过了三分之一，世界还没有一寸属于我。

突击吧...

大四以来一直比较散漫，进了大五更是向没头苍蝇一样乱扑腾。
 
大五应该大悟，可我还在乱扑腾。
 
本科的事情已经全部忘记了，想来当年干的事情不过也是乱扑腾，唯一留下有价值的居然不过是几张证书，一个文凭，外加考了研。不过我本科的主要目标不过就是考研而已，至于樟树同学用于填补空虚的其他成果，连他本人都已经忘记了。可见计划果然很重要...
 
成了研究生，自然也就有了导师。首先“导师”这个词很巧妙，你看，孔子周游列国是他带着一帮学生求学（当然了，孔夫子的桃李满天下，也不排除是他老人家比较关心学生的学习生活情况，挨个的家访了一遍），耶稣也是带着一帮门徒满中东的跑去传教，毛泽东同志更是指出了两万五千里外的延安...总之带路的明显是导师了。而我这种比较容易迷路的，明显暂时还不能当导师...
 
乱扑腾，迷路，没有方向感
不知大家注意没有，鄙人速度很快...
 
最近买了辆车，昨天准备锻炼身体绕二环一圈，结果在第一个转弯处迷失方向，变成自由观光。
最近学了些东西，准备做些东西，结果在选定方向时突发奇想，变成搞飞机。
 
此处的搞飞机是名副其实的搞飞机，但是到了后来就变成了“搞飞机”...
事情是这样的，上学期末某次吃烧烤，大家喝到兴头上，开始讨论造飞机的可行性（喝高了还能做型号论证，不简单）。首先考虑有人驾驶，大家讨论认为造超轻型飞机技术已经比较成熟，拆寝室的床架造飞机结构重量太大，而且大马力发动机不好搞，只是飞起来也没有挑战性（酒精过敏上头了，说话就是有气魄）。
反观无人机方兴未艾，国内军民用市场需求巨大，我校才开了空天学院，我院某教研室做实验就缺架小飞机，研究生院也缺架小飞机...错了，研究生院有创新基金可以申请。发展无人机门槛较低，与我校专业结合较为紧密，应该是有前途的...酒精过敏的我罗列了上述理由，众人曰：善，你上，我们支持你。
事实证明，不可酒后驾车
由此推导，不可酒后论证
 
于是，在地面上扑腾了一个学期以后，终于找到了方向...向天上？！
没人给我方向的时候，我自己找的方向总是很奇怪，也就是本科那些被忘掉的乱扑腾。
 
酒后论证得出的结果应该是一个单片机控制的，由可乐瓶子和竹木架子搭成的能飞的东西，略好于航模，主要用于验证无线电遥控系统和控制率。由于我有近二十年上课画飞机的经验，因此没过多久就出了设计方案，准备动手了。
动手之前，学习一下理论知识。
这个学期有一门叫做软件无线电的课，我觉得很好，受益良多，于是无线电系统就从类似收音机的东西变成了类似火星探测器的东西；
这个学期有一门叫做VERILOG HDL高级数字设计的课程，我觉得很好，受益良多，于是主控机就从单片及变成了基于FPGA的嵌入式系统；
这个学期有一门叫做自适应滤波的课程，我觉得很好，受益良多，于是飞控系统中就从给一个舵机信号，变成了基于卡尔曼滤波的自主飞行控制算法...
这个学期有一门叫做...
人一旦有了什么追求，是很可怕的...和我一起签到重庆的朱凯，已经成为了FPGA和算法强人了，至少比大多数研究生强。
士别三日当刮目相看，如果一个项目几个月可以把人逼成高手，不知道为什么学校不这么干。
如果没有考上研究生，我也可以和他一起去北方试车场开坦克吧？
 
但是，飞机搞到这份上，一个人是搞不出来了...
把方案简化一下，抓住创新点，保留升级潜力，总算一个人有可能弄出来了，结果预算超标。
 
筹钱吧...
于是五一前后又见樟树同学活跃在网站建设和硬件系统设计等战线上（突击！），仿佛又回到了本科乱扑腾的日子，连扑腾的事情都一模一样。真的感觉一点提高都没有，只是变成了熟练工。
五一后本来准备着手搞飞机，结果研究生院老师换岗，要把才建好几个月的招生网重建一遍（今年考研问排名的同学们，你们要是明年考就好了，今年我确实什么都不知道）。所以说某些考核办法真是万恶，逼得老师不得不为了工作成绩搞重复建设...
 
学费很贵，平均下来每个月一千，所以打工是很划不来的。无论如何总算是有了钱搞飞机，但是激情（或者说是酒精的作用，三分钟热度）随着时间也消退了不少....
记得研究生院的某前辈信誓旦旦的要在三年里造一个砖头那么大的手机，现在他要毕业了，可谓成绩斐然前途光明，但手机最终还是不了了之。
数年之后，研究生院的新同志们回忆起我，估计也就是把上面一段话中的“手机”改成“飞机”而已。我准备造的飞机也就只有砖头那么大，比美军在伊拉克用的要好，如果造出来的话。
 
什么叫做乱扑腾，什么叫做搞飞机，我用自己一贯的风格作了诠释：灵感需要一点点疯狂，而我明显过了头
还在乱扑腾，只是我想上天乱扑腾都不行么？
这架飞机如果真做出来会很不错，有很多亮点
比如学校在新校区养的一众鸵鸟乌鸡等等都不会飞，而这架飞机却可以，填补了校内空白。
 
转眼要进教研室了，某些传闻中我们教研室很忙，某些则反之；某些人说我们教研室有些值得搞的东西，有些人说硕士轮不到；导师可能会给一个适合自己的方向吧？也有可能不管你...有种听天由命的感觉。
 
出国的诸位靠自己就找到了正确方向，了不起。看看我，我过去几个月中生命的主线居然是搞飞机。
 
对我来说，搞搞飞机，学学英语这些和口号差不多，已经变成了研究生阶段的乱扑腾了吧...
也许是也许不是，但是如果没有值得拼命的项目，那还不如搞搞飞机娱乐大众。
 
计划？好像有
方向？不管了，突击！

附录：uav计划又更加细化了，也和专业有了更多的联系。多喝了几次酒，uav计划渐渐明朗起来
最近各项计划进行异常顺利，以至于在考试前还能空出十几天。
对考试已经越来越无所谓了，uav计划如果能够有一些技术突破，还是有些价值
技术的价值是不是应该用人命衡量？有时候挺讨厌自己能看到的东西