cuda 怎么读_一、faster-rcnn源码阅读:nms的CUDA编程

打算写一系列faster-rcnn的阅读笔记,侧重于程序实现的细节问题。包括安装,版本选择,编译,数据读取,事无巨细的一一细说。没有规划,没有顺序,没有时间节点,信马由缰,有空就写一点。

我的理解,faster-crnn是目标检测的一个界碑,虽然核心思想源于rcnn和fast-rcnn,但相对以前的算法,性能和效果都实现了质的飞跃,而后面的算法与faster-rcnn有很大的相似性,没有质的飞跃,只是针对特定问题做了不同的优化。好了,这只是我目前的理解,真正的评价完全超出了我的能力范围,就此打住。

要学习一个算法,我想,应该分两个方面:第一是高阶的算法和思想层面,第二是具体实现层面。算法和思想有点像地图,实现层面有点像去实地查看。

对于算法和思想层面最好的是阅读论文,但是对我目前的能力来讲,阅读论文我只能保证每个字都认识,我也希望有一天只要阅读论文就知道了一切。读不懂论文,那就读论文的解读,这方面已很多非常好的文章,我从他们那里学到很多,非常感谢他们,比如:

白裳:一文读懂Faster RCNN​zhuanlan.zhihu.com
47b4f5fc3df978b9e775f9b6a7a98f46.png
Jacqueline:【目标检测】Faster R-CNN代码解析​zhuanlan.zhihu.com
https://zhuanlan.zhihu.com/p/32404424​zhuanlan.zhihu.com
37257c4a002088e2e93e70ce62a98e2a.png

当然,相对于算法和思想,实现细节并不重要,但对于学习,还是仔仔细细,踏踏实实为好。仔细学习一个经典算法,其他类似的算法只要看看论文就知道大概怎么实现了,所以对学习来讲,慢就是快。

本文的主要内容是nms的具体实现,通过阅读GPU代码,学习CUDA编程。

简单讲,NMS就是去重,重叠度大的矩形框中,保留分数值比较大的一个。IOU(交并比,即重叠度)和NMS(非极大抑制)是什么(参见https://zhuanlan.zhihu.com/p/54709759),自行搜索,不再一一列举。

12ab9c5c266a9ba84c61d460dd9df4ad.png

一、先看cpu版的代码,后面再看gpu版,cpu版比较容易懂,调用也很容易,

nms_cpu.py(来源https://github.com/jwyang/faster-rcnn.pytorch):

from 

二、CUDA版

cpu版验证和理解了算法,下面来看看GPU实现加速,速度大概可以提升50X。在2080it上大概是115ms比3ms。

nms各种实现的benchmark请看:

fmscole/benchmark​github.com
95dfeb12ee2b067d0de25f6d19840242.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
  • 0
    点赞
  • 0
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值