打算写一系列faster-rcnn的阅读笔记,侧重于程序实现的细节问题。包括安装,版本选择,编译,数据读取,事无巨细的一一细说。没有规划,没有顺序,没有时间节点,信马由缰,有空就写一点。
我的理解,faster-crnn是目标检测的一个界碑,虽然核心思想源于rcnn和fast-rcnn,但相对以前的算法,性能和效果都实现了质的飞跃,而后面的算法与faster-rcnn有很大的相似性,没有质的飞跃,只是针对特定问题做了不同的优化。好了,这只是我目前的理解,真正的评价完全超出了我的能力范围,就此打住。
要学习一个算法,我想,应该分两个方面:第一是高阶的算法和思想层面,第二是具体实现层面。算法和思想有点像地图,实现层面有点像去实地查看。
对于算法和思想层面最好的是阅读论文,但是对我目前的能力来讲,阅读论文我只能保证每个字都认识,我也希望有一天只要阅读论文就知道了一切。读不懂论文,那就读论文的解读,这方面已很多非常好的文章,我从他们那里学到很多,非常感谢他们,比如:
白裳:一文读懂Faster RCNNzhuanlan.zhihu.com![47b4f5fc3df978b9e775f9b6a7a98f46.png](https://i-blog.csdnimg.cn/blog_migrate/d547d7792b5618ea6bf8930f4235cc0f.jpeg)
![37257c4a002088e2e93e70ce62a98e2a.png](https://i-blog.csdnimg.cn/blog_migrate/85b018198f3bb8e239c7a587930489e8.png)
当然,相对于算法和思想,实现细节并不重要,但对于学习,还是仔仔细细,踏踏实实为好。仔细学习一个经典算法,其他类似的算法只要看看论文就知道大概怎么实现了,所以对学习来讲,慢就是快。
本文的主要内容是nms的具体实现,通过阅读GPU代码,学习CUDA编程。
简单讲,NMS就是去重,重叠度大的矩形框中,保留分数值比较大的一个。IOU(交并比,即重叠度)和NMS(非极大抑制)是什么(参见https://zhuanlan.zhihu.com/p/54709759),自行搜索,不再一一列举。
![12ab9c5c266a9ba84c61d460dd9df4ad.png](https://i-blog.csdnimg.cn/blog_migrate/e56bbe2f41475b54c8d1fd2719f9fb6c.jpeg)
一、先看cpu版的代码,后面再看gpu版,cpu版比较容易懂,调用也很容易,
nms_cpu.py(来源https://github.com/jwyang/faster-rcnn.pytorch):
from
二、CUDA版
cpu版验证和理解了算法,下面来看看GPU实现加速,速度大概可以提升50X。在2080it上大概是115ms比3ms。
nms各种实现的benchmark请看:
fmscole/benchmarkgithub.com![95dfeb12ee2b067d0de25f6d19840242.png](https://i-blog.csdnimg.cn/blog_migrate/d17829219cd6dcbffb1a7e964860d7a6.png)
主控函数分两部分,第一部分计算mask,还需要的标0,不需要的(重复度大的)标1,第二部分是根据mask,选出留下来的候选框。
先简要说一下CUDA编程模型:
GPU之所以能够加速,是因为并行计算,即每个线程负责计算一个数据,充分利用GPU计算核心超多(几千个)的优势。
(1)每个计算核心相互独立,运行同一段代码,这段代码称为核函数;
(2)每个核心有自己的身份id,线程的身份id是两个三维数组:(blockIdx.x,blockIdx.y,blockIdx.z)-(threadIdx.x,threadIdx.y,threadIdx.z)。
身份id被另两个三维数组grid(gridDim.x,gridDim.y,gridDim.z)和block(blockDim.x,blockDim.y,blockDim.z)确定范围。
总共有gridDim.x×gridDim.y×gridDim.z个block,
每个block有blockDim.x×blockDim.y×blockDim.z个thread。
有了线程的身份id,经过恰当的安排,让身份id(核函数可以获取):(blockIdx.x,blockIdx.y,blockIdx.z)-(threadIdx.x,threadIdx.y,threadIdx.z)对应到一个数据,就可以实现一个线程计算一个数据,至于如何对应,开发人员得好好安排,可以 说这是CUDA开发的一个核心问题。
gridDim.x、blockIdx.x这些是核函数可以获取的,gridDim.x等于多少,调用核函数的时候就要定一下来。看代码:
dim3
这里的threadsPerBlock=8*8=64,
当boxes_num=12030时,DIVUP(12030, 64)=12030/64+12030%64>0=188
在调用核函数的时候,通过<<<blocks, threads>>>(#这是cu语法,不是标准C语言)把线程数量安排传递进去,核函数里就有
gridDim.x=188,gridDim.y=188,gridDim.z=1;
blockDim.x=64,blockDim.y=1,blockDim.z=1;
0<=blockIdx.x<188,
0<=blockIdx.y<188,
blockIdx.z=0,
0<=threadIdx.x<64,
threadIdx.y=threadIdx.z=0,
这样就启动了2,262,016个(两百多万个线程)来计算,两百多万看起来吓人,对GPU来书毫无负担!每个线程计算不超过64个值,后面再讲。
(3)这里的grid(a,b,c),block(x,y,z)值是多少,由程序设计人员根据问题来定,在调用核函数时就要确定下来,但有一个基本限制block(x,y,z)中的x×y×z<=1024(这个值随GPU版本确定,起码nvidia 1080,2080都是这样);
(4)block中的线程每32个thread为一束,绝对同步:比如if-else语句,这32个线程中有的满足if条件,有的满足else。满足else的那部分线程不能直接进入,而是要等满足if的那部分线程运行完毕才进入else部分,而满足if的那部分线程现在也不能结束,而是要等else部分线程运行完毕,大家才能同时结束。for语句也是一样。因此GPU计算尽可能不要有分支语句。
不是说不能用if和for,该用还得用,用的时候要知道付出的代价。否则实现了减速都不知道为了啥。
不同的线程束之间不同步,如果同步需要请__syncthreads();
如果设置block(1),即一个block只安排一个线程呢?事实上GPU还是要启动32个线程,另外31个陪跑。
因此block(x,y,z)中的x×y×z应该为32的倍数。不过32×32=1024了。
(5)要并行计算,前提是数据之间没有相互依赖,有前后依赖的部分只能放在同一个核函数里计算;
先看控制部分代码,这部分做的事情就是:
1,在GPU上分配内存,把数据传到GPU
2,调用核函数,计算mask;
3,把数据传回来,
4,根据mask把获取保留下来的候选框。
nms_kernel.cu(来源https://github.com/jwyang/faster-rcnn.pytorch):
void
计算mask的核函数是核心代码,什么是mask呢?从内存里讲是一段连续内存,但应该把它想象成一个矩阵:
比如候选框个数为boxes_num=12030时,由于要计算任意两个候选框之间的IOU是否大于阈值,因此,我们需要建一个12030*12030的矩阵,1表示两个候选框的IOU大于0.7,0表示不大于。这样任意两个候选框之间的关系都可以用这个12030行*12030列(行表示本框,列表示其他框)的矩阵保存。
但是,这里有好几处值得改进的地方:
1)对角线上的值不需要计算,因为意味着自己与自己计算IOU,没意义;
2)上三角与下三角是对称的,只需要用到上三角即可;
3)更重要的是,为了保存0或1的值,真需要建这么大的矩阵吗?事实上每连续的64个0、1刚好构成一个无符号的64位整数(unsigned long long),我们只用一个整数表示即可,这样内存减少64倍!这中间设计到位运算,这不是问题。原本需要一个12030*12030的一个数组来表示这个矩阵,现在只要12030*188的一个unsigned long long数组就可以映射这个12030行,188列的矩阵。
比如:这长度为1230*188数组中的第一组188个整数(相当于矩阵的第一行)记录的是第一个框(按分数排序之后)与其他所有框之间的重叠关系,第一个整数的第一位表示与自己,由于自己与自己不参与计算,所以这个值一定为0.
__global__
再总结一下,总共规划了(188,188,1)个block,每个block有(64,1,1)个线程,第(c,r,0)-(t,0,0)号线程负责计算的数据是:r×64+t号候选框与c×64~c×64+63号这64个候选框之间的重叠关系,并存进(r×64+t)×188+c这个整数里。当c,r遍历完188,t遍历完64,这188×188×64个线程就计算完了任意两个候选框之间的关系。
由于cuda里每个线程计算一个数据,相当于cpu里的循环。CPU里的循环是一个一个的算,而cuda里是同时在算。
本质上,这里gird(188,188,1)和block(64,1,1)来代替了三重循环:
for(int i=0;i<188;i++)
for(int j=0;j<188;j++)
for(int k=0;k<64;k++)
。。。。。。
只要把这循环中的i,j,k替换为blockIdx.x,blockIdx.y,threadIdx.x即可。
CUDA编程就是循环替代!——这是我目前的理解。
回过头来看把候选框复制进共享内存这部分,
block_boxes[threadIdx.x * 5 + 0] =
dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 0];
可以看出这里面只用到了blockIdx.x和threadIdx.x,注意到了没,没用到blockIdx.y,而blockIdx.y的范围是0~187,什么意思,意味着这段代码被重复执行了188次!毕竟,我们启动了188×64×188个线程,而数据只有12030个数据,因此有188个线程执行的是相同的数据!不知道我理解的对不对,望大佬帮我指正。
计算好了mask数组之后,计算保留下来的目标框:
//...................
最后,让我吃惊的是,你竟然看到了这里,陪我唠叨了这么久,那么不妨再看下一篇,等一等,先点个赞再走,哈哈:
无用:二、faster-rcnn源码阅读:ROIPool的CUDA编程zhuanlan.zhihu.com