cuda第一次計算耗時 nms的cuda實現解析

nms是不太好在cuda上實現的，因為nms的計算過程是有依賴關係的，比如a,b,c三個置信度由大到小的檢測框，如果iou(a,b)大於閾值，那麼bc的iou則不必再計算了，由於依賴關係會限制cuda的發揮，因此cuda實現中，所有的box之間的iou都是需要計算的，不過有的結果是不需要的。

**中nms的輸出是長度為

bool dev_mask[boxes_num][boxes_num];
// nms kernel ...
bool mask[boxes_num];
for (int i = 0; i < boxes_num; ++i) 
}

作者用unsigned long long則是將64位的unsigned long long的每乙個bit當做乙個box的結果，bit位1則刪除，為0則保留，dev_mask變為

再來看看作者的實現

int const threadsperblock = sizeof(unsigned long long) * 8;
const int col_blocks = divup(boxes_num, threadsperblock);
std::vectorremv(col_blocks);
memset(&remv[0], 0, sizeof(unsigned long long) * col_blocks);  
int num_to_keep = 0;
for (int i = 0; i < boxes_num; i++) }}

再來看看nms kernel的實現，首先kernel的grid是2d的，block是1d的

dim3 blocks(divup(boxes_num, threadsperblock),
divup(boxes_num, threadsperblock));
dim3 threads(threadsperblock);
nms_kernel<<>>(boxes_num,
nms_overlap_thresh,
boxes_dev,
mask_dev);

其中每乙個thread計算dev_mask中的乙個值，可以通過乙個示意圖更好地理解並行的邏輯

其中第i行的block計算以第

block(0,0) 計算[0,threadsperblock](基準) x [0,threadsperblock](其他)的dev_mask

block(0,1) 計算[0,threadsperblock] x [threadsperblock,threadsperblock*2]的dev_mask

block(0,2) 計算[0,threadsperblock] x [threadsperblock*2,threadsperblock*3]的dev_mask

block(i,j) 計算

[i*threadsperblock, (i+1)*threadsperblock]（基準）x [j*threadsperblock,(j+1)*threadsperblock](其他)的devmask結果，如果把dev_mask看成二維陣列更好理解，dev_mask的第一維是基準box的索引，第二維是其他box的索引

更新:

從更簡單的形式開始分析會更好理解，比如只有乙個block，block的大小是2d的，為nxn，n為檢測框數量，假設n=6，並行邏輯如下圖

每個thread計算一對box的iou（只用計算上三角部分）作者的實現則是將6x6的分成3x3的小block，每個block的thread為2，即

這樣做的好處是減少cuda資源占用，原來需要36個thread，現在只要18個thread（每個thread要計算2對box的iou）並且每個小的block的其他box可以存入共享記憶體，減少全域性記憶體的訪問

__global__ void nms_kernel(const int n_boxes, const float nms_overlap_thresh,
const float *dev_boxes, unsigned long long *dev_mask) 
__syncthreads();
if (threadidx.x < row_size) 
// 遍歷其他box，計算當前基準box和其他box的iou，如果大於nms閾值，則置對應的位為1，即刪除
for (i = start; i < col_size; i++) 
}const int col_blocks = divup(n_boxes, threadsperblock);
// 結果寫入dev_mask
dev_mask[cur_box_idx * col_blocks + col_start] = t;
}}

cuda第一次計算耗時 nms的cuda實現解析

第一次飛行

第一次面試

第一次哭泣

cuda第一次計算耗時 nms的cuda實現解析

第一次飛行

第一次面試

第一次哭泣

相關推薦