在了解了CUDA编程的核心思想后，我们便可以开始阅读nms_kernel.cu文件了，先直接上源码（部分简单的已经注释），如下：

1. // ------------------------------------------------------------------
2. // Faster R-CNN
3. // Copyright (c) 2015 Microsoft
4. // Licensed under The MIT License [see fast-rcnn/LICENSE for details]
5. // Written by Shaoqing Ren
6. // ------------------------------------------------------------------
7. #include "gpu_nms.hpp"
8. #include <vector>
9. #include <iostream>
10. //cudaError_t是cuda中的一个类，用于记录cuda错误（所有的cuda函数，几乎都会返回一个cudaError_t）
11. #define CUDA_CHECK(condition) \
12. /* Code block avoids redefinition of cudaError_t error */ \
13. do { \
14. cudaError_t error = condition; \
15. if (error != cudaSuccess) { \
16. std::cout << cudaGetErrorString(error) << std::endl; \
17. } \
18. } while (0)
19. //DIVUP即实现除法的向上取整
20. #define DIVUP(m,n) ((m) / (n) + ((m) % (n) > 0))
21. //unsigned long long类型是目前C语言中精度最高的数据类型，为64位精度
22. //threadsPerBlock即自定义的每个Block所含有的线程数目（每个Block的线程数不宜太多，也不宜太少）
23. int const threadsPerBlock = sizeof(unsigned long long) * 8; //其实threadsPerBlock = 64
24. __device__ inline float devIoU(float const * const a, float const * const b) {
25. float left = max(a[0], b[0]), right = min(a[2], b[2]);
26. float top = max(a[1], b[1]), bottom = min(a[3], b[3]);
27. float width = max(right - left + 1, 0.f), height = max(bottom - top + 1, 0.f);
28. float interS = width * height;
29. float Sa = (a[2] - a[0] + 1) * (a[3] - a[1] + 1);
30. float Sb = (b[2] - b[0] + 1) * (b[3] - b[1] + 1);
31. return interS / (Sa + Sb - interS);
32. }
33. //nms kernel
34. /*
35. 参数n_boxes:边界框数目
36. 参数nms_overlap_thresh:交并比阈值
37. */
38. __global__ void nms_kernel(const int n_boxes, const float nms_overlap_thresh,
39. const float *dev_boxes, unsigned long long *dev_mask) {
40. const int row_start = blockIdx.y;
41. const int col_start = blockIdx.x;
42. // if (row_start > col_start) return;
43. const int row_size =
44. min(n_boxes - row_start * threadsPerBlock, threadsPerBlock);
45. const int col_size =
46. min(n_boxes - col_start * threadsPerBlock, threadsPerBlock);
47. __shared__ float block_boxes[threadsPerBlock * 5]; //共享内存
48. if (threadIdx.x < col_size) {
49. block_boxes[threadIdx.x * 5 + 0] =
50. dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 0];
51. block_boxes[threadIdx.x * 5 + 1] =
52. dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 1];
53. block_boxes[threadIdx.x * 5 + 2] =
54. dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 2];
55. block_boxes[threadIdx.x * 5 + 3] =
56. dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 3];
57. block_boxes[threadIdx.x * 5 + 4] =
58. dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 4];
59. }
60. __syncthreads(); //同步线程
61. if (threadIdx.x < row_size) {
62. const int cur_box_idx = threadsPerBlock * row_start + threadIdx.x;
63. const float *cur_box = dev_boxes + cur_box_idx * 5;
64. int i = 0;
65. unsigned long long t = 0;
66. int start = 0;
67. if (row_start == col_start) {
68. start = threadIdx.x + 1;
69. }
70. for (i = start; i < col_size; i++) {
71. if (devIoU(cur_box, block_boxes + i * 5) > nms_overlap_thresh) {
72. t |= 1ULL << i; //1ULL = unsigned long long型的数字1（最高位为第64位）
73. }
74. }
75. const int col_blocks = DIVUP(n_boxes, threadsPerBlock);
76. dev_mask[cur_box_idx * col_blocks + col_start] = t;
77. }
78. }
79. //设置哪个GPU用于nms
80. void _set_device(int device_id) {
81. int current_device;
82. CUDA_CHECK(cudaGetDevice(&current_device)); //获取当前GPU序号
83. if (current_device == device_id) {
84. return;
85. }
86. // The call to cudaSetDevice must come before any calls to Get, which
87. // may perform initialization using the GPU.
88. CUDA_CHECK(cudaSetDevice(device_id)); //设置device_id号GPU生效
89. }
90. void _nms(int* keep_out, int* num_out, const float* boxes_host, int boxes_num,
91. int boxes_dim, float nms_overlap_thresh, int device_id) {
92. _set_device(device_id);
93. float* boxes_dev = NULL;
94. unsigned long long* mask_dev = NULL;
95. const int col_blocks = DIVUP(boxes_num, threadsPerBlock);
96. CUDA_CHECK(cudaMalloc(&boxes_dev,
97. boxes_num * boxes_dim * sizeof(float)));
98. CUDA_CHECK(cudaMemcpy(boxes_dev,
99. boxes_host,
100. boxes_num * boxes_dim * sizeof(float),
101. cudaMemcpyHostToDevice));
102. CUDA_CHECK(cudaMalloc(&mask_dev,
103. boxes_num * col_blocks * sizeof(unsigned long long)));
104. dim3 blocks(DIVUP(boxes_num, threadsPerBlock),
105. DIVUP(boxes_num, threadsPerBlock));
106. dim3 threads(threadsPerBlock);
107. nms_kernel<<<blocks, threads>>>(boxes_num,
108. nms_overlap_thresh,
109. boxes_dev,
110. mask_dev);
111. std::vector<unsigned long long> mask_host(boxes_num * col_blocks);
112. CUDA_CHECK(cudaMemcpy(&mask_host[0],
113. mask_dev,
114. sizeof(unsigned long long) * boxes_num * col_blocks,
115. cudaMemcpyDeviceToHost));
116. std::vector<unsigned long long> remv(col_blocks);
117. memset(&remv[0], 0, sizeof(unsigned long long) * col_blocks);
118. int num_to_keep = 0;
119. for (int i = 0; i < boxes_num; i++) {
120. int nblock = i / threadsPerBlock;
121. int inblock = i % threadsPerBlock;
122. if (!(remv[nblock] & (1ULL << inblock))) {
123. keep_out[num_to_keep++] = i;
124. unsigned long long *p = &mask_host[0] + i * col_blocks;
125. for (int j = nblock; j < col_blocks; j++) {
126. remv[j] |= p[j];
127. }
128. }
129. }
130. *num_out = num_to_keep;
131. CUDA_CHECK(cudaFree(boxes_dev));
132. CUDA_CHECK(cudaFree(mask_dev));
133. }

1.devIoU()函数

1. //devIoU计算两个边界框之间的交并比
2. //__device__是CUDA中的限定词，具体含义如下图
3. //float const * const a表示a是常量指针常量，即a是一个指针常量（不可修改的指针），指向一个常量
4. __device__ inline float devIoU(float const * const a, float const * const b) {
5. float left = max(a[0], b[0]), right = min(a[2], b[2]);
6. float top = max(a[1], b[1]), bottom = min(a[3], b[3]);
7. float width = max(right - left + 1, 0.f), height = max(bottom - top + 1, 0.f);
8. float interS = width * height; //交集
9. float Sa = (a[2] - a[0] + 1) * (a[3] - a[1] + 1); //边界框a的面积
10. float Sb = (b[2] - b[0] + 1) * (b[3] - b[1] + 1); //边界框b的面积
11. return interS / (Sa + Sb - interS);
12. }

       <p>可由计算能力为3的设备调用（callable from the device for devices of compute capability 3）</p></td><td style="width:173px;">返回必须为void型（即不能范围任何其他类型）</td></tr><tr><td style="width:137px;">__device__</td><td style="width:240px;">在设备上执行（GPU）</td><td style="width:249px;">只能被设备（device）调用</td><td style="width:173px;">&nbsp;</td></tr><tr><td style="width:137px;">__host__</td><td style="width:240px;">在主机上执行（CPU）</td><td style="width:249px;">只能被主机调用</td><td style="width:173px;">可以省略</td></tr></tbody></table></div><p>注：</p>

__device__ ：声明一个函数是设备上执行的，仅可以从设备调用；
__global__： 在设备上执行，可以从主机调用；
__host__ ： 声明的函数是在主机上执行的，仅可从主机调用；
__device__和__global__函数不支持递归；
__device__和__global__函数不能声明静态变量在它们内部。

2.nms_kernel()函数

1. //nms kernel（CUDA编程中的核函数）
2. /*
3. 参数n_boxes:边界框数目
4. 参数nms_overlap_thresh:交并比阈值
5. 参数dev_boxes:存储边界框信息，每五位组成一个边界框信息，[left.x,left.y,right.x,right.y,class]
6. 参数dev_mask:存储边界框间的交并比是否超过上述阈值的信息，以ULL类型进行表示，与哪个框交并比超过阈值，相应位置1，否则置0（输出参数）
7. */
8. __global__ void nms_kernel(const int n_boxes, const float nms_overlap_thresh,
9. const float *dev_boxes, unsigned long long *dev_mask) {
10. const int row_start = blockIdx.y; //当前调用的block的y坐标（实际是一个索引）
11. const int col_start = blockIdx.x; //当前调用的block的x坐标
12. // if (row_start > col_start) return;
13. //min()的目的是防止从dev_boxes中读取数据越界（原因是n_boxes不一定被threadsPerBlock整除）
14. //实际上只有最后一个block中所需要的线程数目可能小于threadsPerBlock，其余均等于threadsPerBlock
15. const int row_size =
16. min(n_boxes - row_start * threadsPerBlock, threadsPerBlock);
17. const int col_size =
18. min(n_boxes - col_start * threadsPerBlock, threadsPerBlock);
19. //__shared__限定词，即每个block中的所有线程共享内存
20. __shared__ float block_boxes[threadsPerBlock * 5]; //数字5即边界框的5个信息
21. if (threadIdx.x < col_size) {
22. block_boxes[threadIdx.x * 5 + 0] =
23. dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 0]; //left.x
24. block_boxes[threadIdx.x * 5 + 1] =
25. dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 1]; //left.y
26. block_boxes[threadIdx.x * 5 + 2] =
27. dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 2]; //right.x
28. block_boxes[threadIdx.x * 5 + 3] =
29. dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 3]; //right.y
30. block_boxes[threadIdx.x * 5 + 4] =
31. dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 4]; //class
32. }
33. __syncthreads(); //同步线程(使得当前block中的所有线程均读取到相应边界框信息后再执行后面的代码)
34. //以下代码实现某一边界框与其余所有边界框（删去了部分重复）进行交并比的阈值判断
35. if (threadIdx.x < row_size) {
36. const int cur_box_idx = threadsPerBlock * row_start + threadIdx.x; //当前选中的边界框索引
37. const float *cur_box = dev_boxes + cur_box_idx * 5; //当前选中的边界框信息首地址索引
38. int i = 0;
39. unsigned long long t = 0; //用于记录与当前边界框交并比情况，大于阈值相应位置1
40. int start = 0;
41. if (row_start == col_start) { //如果当前边界框所处的block与要比较的边界框所处的block相同，则start不从0开始，减少重复计算
42. start = threadIdx.x + 1;
43. }
44. for (i = start; i < col_size; i++) {
45. if (devIoU(cur_box, block_boxes + i * 5) > nms_overlap_thresh) {
46. t |= 1ULL << i; //1ULL = unsigned long long型的数字1（最高位为第64位）；每一位就代表一个边界框索引，如果大于阈值，则该位置1
47. }
48. }
49. const int col_blocks = DIVUP(n_boxes, threadsPerBlock);
50. dev_mask[cur_box_idx * col_blocks + col_start] = t; //存入当前边界框与当前选定的block中的64个边界框的交并比比较情况，用于后续的nms
51. }
52. }

此函数在理解上可能会有一定困难，以下我以图像的方式稍生动一点来说明该函数在干什么。

A.函数输入的dev_boxes中存储的内容如下图（每一个边界框都有5个信息按顺序存储着）：

B.函数输出的dev_mask中存储的内容如下图（threadsPerBlock即每个block所含有的线程数）：

图中是数字，拿0x11为例说明如下：

0x11 = 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0001 0001（共64位）

其中第5位为1，表示当前矩形框与当前选中的block中的第5号矩形框的交并比大于设定的阈值。

由于0x10处于box1所在的第1位，更具体一步表示就是box1与第一个block中的64个边界框中的第5个（即box5）的交并比大于设定的阈值。

注：上述所有索引都从1开始（算法本身是从0开始）;图中的符号表示向上取整。

C. 具体在干什么

作者的一大巧妙之处是将二维的block中的两维都表示为dev_boxes，即blockIdx.x和blockIdx.y名义上是block的索引，但实际上表示的是将dev_boxes分块后的块索引，如下图：

代码在干的事就是取当前blockIdx.y块中的第threadIdx.x个边界框与当前blockIdx.x块中的所有边界框进行交并比上的判断，由此为后续nms做准备。

但是为了降低部分重复计算，如(box1, box2)和(box2, box1)这成对的重复计算，采用如下代码：

1. if (row_start == col_start) { //如果当前边界框所处的block与要比较的边界框所处的block相同，则start不从0开始，减少重复计算
2. start = threadIdx.x + 1;
3. }

但细心的你们一定会发现，其实上述代码只避免了相同块中的重复计算，对于不同块之间仍旧存在重复计算，例如(box1, box65)和(box65, box1)，其中box1属于第blockIdx.y = 0块，box65属于blockIdx.y = 1 块。（当然重复计算并不会影响后续的nms，但会消耗时间）

3._nms()函数

1. //此函数实际上的__host__类型，真正实现nms
2. /*
3. 参数keep_out:int型指针，用于存储所有保留下来的边界框索引
4. 参数num_out:保留下的边界框数目
5. 参数:boxes_host:输入参数，存储着边界框信息，来自于主机
6. 参数boxes_num:输入的边界框数目
7. 参数boxes_dim:边界框维度（一般为5，即左上角、右下角和类别）
8. 参数nms_overlap_thresh:交并比阈值，用于nms
9. 参数device_id:GPU设备号
10. */
11. void _nms(int* keep_out, int* num_out, const float* boxes_host, int boxes_num,
12. int boxes_dim, float nms_overlap_thresh, int device_id) {
13. _set_device(device_id); //设置相应设备
14. float* boxes_dev = NULL;
15. unsigned long long* mask_dev = NULL;
16. const int col_blocks = DIVUP(boxes_num, threadsPerBlock); //向上取整，即当前输入分块后的块数目
17. CUDA_CHECK(cudaMalloc(&boxes_dev,
18. boxes_num * boxes_dim * sizeof(float))); //开辟显存
19. CUDA_CHECK(cudaMemcpy(boxes_dev,
20. boxes_host,
21. boxes_num * boxes_dim * sizeof(float),
22. cudaMemcpyHostToDevice)); //将host输入的数据送入到boxes_dev中
23. CUDA_CHECK(cudaMalloc(&mask_dev,
24. boxes_num * col_blocks * sizeof(unsigned long long)));
25. dim3 blocks(DIVUP(boxes_num, threadsPerBlock),
26. DIVUP(boxes_num, threadsPerBlock)); //所设置的block为二维block，两维的大小相同
27. dim3 threads(threadsPerBlock); //每一个block中的线程为一维，均为threadsPerBlock条线程
28. nms_kernel<<<blocks, threads>>>(boxes_num,
29. nms_overlap_thresh,
30. boxes_dev,
31. mask_dev); //调用上述定义的核函数获取交并比情况
32. std::vector<unsigned long long> mask_host(boxes_num * col_blocks);
33. CUDA_CHECK(cudaMemcpy(&mask_host[0],
34. mask_dev,
35. sizeof(unsigned long long) * boxes_num * col_blocks,
36. cudaMemcpyDeviceToHost)); //从device中处理好的数据送回mask_host，进行后续CPU计算
37. std::vector<unsigned long long> remv(col_blocks); //存储要移除的边界框索引
38. memset(&remv[0], 0, sizeof(unsigned long long) * col_blocks); //初始化为0
39. //以下正式开始进行nms，思想和CPU版本有所不同，但本质是一样的
40. //由于输入此函数的boxes_host是按置信度从高到低排过序，所以第一个边界框肯定会存入keep_out中
41. int num_to_keep = 0;
42. for (int i = 0; i < boxes_num; i++) {
43. int nblock = i / threadsPerBlock; //当前边界框输入哪一个block
44. int inblock = i % threadsPerBlock; //当前边界框输入对应block中的第几个
45. //当i = 0时，remv[0] = 0（初始值），但由于第一个边界框肯定要存入keep_out中，所以没问题
46. if (!(remv[nblock] & (1ULL << inblock))) { //判断当前边界框与前面保留下来的边界框之间的交并比是否大于阈值
47. keep_out[num_to_keep++] = i; //如果不大于阈值，则当前边界框应该保留
48. unsigned long long *p = &mask_host[0] + i * col_blocks;
49. for (int j = nblock; j < col_blocks; j++) {
50. remv[j] |= p[j]; //预存入后续所有边界框是否要被移除的信息（相应位为1则移除）
51. }
52. }
53. }
54. *num_out = num_to_keep;
55. CUDA_CHECK(cudaFree(boxes_dev));
56. CUDA_CHECK(cudaFree(mask_dev));
57. }

此函数的nms部分可能较难理解（越是没几行的代码越是难以理解），我就举个例子引导一下大家的思维：

假如当前的i = 0，即取到box1，根据nms的原理可知，box1肯定会保留下来（因为它的置信度最高），即!(remv[nblock] & (1ULL << inblock)) = true一定得成立（故remv的所有元素要初始化为0，原因便在于此），由此会进入到if中执行里面的代码。

这时关键就来了，作者通过按位或操作来快速形成要移除的边界框索引，即如下代码：

1. unsigned long long *p = &mask_host[0] + i * col_blocks;
2. for (int j = nblock; j < col_blocks; j++) {
3. remv[j] |= p[j]; //预存入后续所有边界框是否要被移除的信息（相应位为1则移除）
4. }

所谓的要移除的边界框索引是指：如果remv[n]中的某一位的值为1，则第n个block中对应的该位所对应的边界框需要被移除，因为该边界框与保留下来的某一边界框的交并比已经超过了所设定的阈值。

好了，回到当前的box1，因为所有的边界框都被分配到了相应的块（block）中，所以remv数组的大小为col_blocks，而通过循环按位或后，remv中存储的是box1与其余边界框的交并比比较情况，也即要移除的边界框索引。

当 i = 1时，如果remv[0]的第2位（从1开始）为1，则不进入if，即直接移除不保留；如果为0，则进入if，保留box2的索引，以及更新remv。更新过程就是将box1的dev_mask中的内容（也即当前的remv）与box2的dev_mask中的内容进行按位或，意思就是如果box3与更新后的remv中的对应为吻合，则我们不需要管是和box1还是box2的交并比超过了阈值，直接将其移除即可。

后面的过程依此类推。

NMS算法的GPU实现（使用CUDA加速计算）相关推荐

1. CUDA加速计算矩阵乘法进阶玩法(共享内存)

   CUDA加速计算矩阵乘法&进阶玩法~共享内存 一.基础版矩阵乘法 二.为什么可以利用共享内存加速矩阵乘法 1.CUDA内存读写速度比较 2.申请共享内存 三.改进版矩阵乘法(利用共享内存) 一 ...

2. CUDA加速计算的基础C/C++

   本文是Nvidia 90美金的课程笔记 无论是从出色的性能,还是从易用性来看,CUDA计算平台都是加速计算的制胜法宝.CUDA 提供了一种可扩展 C.C++.Python 和 Fortran 等语言的 ...

3. 利用gpu加速神经网络算法,外接gpu 训练神经网络

   神经网络做图像分类一定要用到gpu吗? GPU最大的价值一直是"accelerating"(加速),GPU不是取代CPU,而是利用GPU的并行计算架构,来将并行计算的负载放到GPU ...

4. FFmpeg在Intel GPU上的硬件加速与优化

   英特尔提供了一套基于VA-API/Media SDK的硬件加速方案,通过在FFmpeg中集成Intel GPU的媒体硬件加速能力,为用户提供更多的收益.本文来自英特尔资深软件开发工程师赵军在LiveV ...

5. gpu处理信号_GPU显卡不仅用来打游戏那么简单，它还可以用于通用加速计算

   如今,显卡不仅在工作站.个人PC中变得非常重要,而且在数据中心也处于举足轻重的地位.CPU负责通用计算.GPU负责加速计算已经成为绝大数数据中心一种常态.用于加速计算的GPU专用处理器,它将计算密集型 ...

6. MATLAB上的GPU加速计算

   概述 怎样在MATLAB上做GPU计算呢?主要分为三个步骤:数据的初始化.对GPU数据进行操作.把GPU上的数据回传给CPU 一.数据的初始化 首先要进行数据的初始化.有两种方法可以进行初始化:一是先 ...

7. 用好CUDA加速 6款视频软件评测与指南

   从2008年下半年开始和GTX280的发布,NVIDIA的GPU从传统的单一3D渲染角色快速像通用并行处理器+3D渲染角色转变.近一年来,基于NVIDIA CUDA架构GPU的应用情况已经非常清晰.基 ...

8. 使用c++onnxruntime部署yolov5模型并使用CUDA加速（超详细）

   文章目录 前言 1.Yolo简介 2.onnxruntime简介 3.Yolov5模型训练及转换 4.利用cmake向C++部署该onnx模型 总结 前言 接到一个项目,需要用c++和单片机通信,还要 ...

9. 《GPU高性能编程CUDA实战》中代码整理

   CUDA架构专门为GPU计算设计了一种全新的模块,目的是减轻早期GPU计算中存在的一些限制,而正是这些限制使得之前的GPU在通用计算中没有得到广泛的应用. 使用CUDA C来编写代码的前提条件包括:( ...

最新文章

1. 东芝发布15nm SG5固态硬盘 容量高达1TB
2. 【安卓开发】Layout Inflation不能这么用
3. nyist 541最强DE 战斗力
4. firefox英文网页乱码解决方法
5. Spark GraphX
6. centos7安装oracle12c 二
7. lnmp 1.4 mysql_lnmp1.4配置https教程
8. leetcode python3 简单题232. Implement Queue using Stacks
9. 设计模式—工厂模式（思维导图）
10. windows 下安装 mysql
11. Microsoft Visual Studio 2012 旗舰版 镜像 ISO 官方下载地址 旗舰版 序列号 SN VS2012_ULT_chs.iso
12. 简单的病毒制作及破解方法
13. php模拟登陆校园教务网络管理系统
14. ALLEGRO 任意形状铜皮倒圆角
15. Kafka bootstrap.servers
16. FME 2011预览：新特性 IFMEWorkspaceRunner
17. CCLE（Cancer Cell Line Encyclopedia）数据库使用
18. 对一个整形数组进行顺序排列
19. Gtk-ERROR **: GTK+ 2.x symbols detected. Using GTK+ 2.x and GTK+ 3 in the same process is not suppor
20. 西瓜视频直播助手下载与安装过程 0523

热门文章

1. 新书《区块链实战》中信出版：破除迷思 推动技术落地 | ArcBlock 新闻
2. Samba服务概述、配置、及其搭建
3. Mysql 语句操作索引SQL语句
4. 考研计算机学校张雪峰,“网红”张雪峰考研宣讲会圆满落幕
5. Fresh Diagnose 8.65-电脑硬件信息查看和硬件性能基准测试
6. python数据图表可视化GUI,python做可视化数据图表
7. 存储过程中使用休眠(sleep)函数
8. 处理Unity中使用HTC VIVE PRo的SRWork插件若干问题
9. mybatis最佳实践
10. 电脑彻底卸载双系统中的Ubuntu ThinkPad X1 Carbon