bool dev_mask[boxes_num][boxes_num];
// nms kernel ...
bool mask[boxes_num];
for (int i = 0; i < boxes_num; ++i) {
  // 如果当前的基准box被删除，则不用计算当前的结果
  if (mask[i]) continue;
  // 前面的不用算，因为都是置信度高于当前基准box的
  for (int j = i + 1; j < boxes_num; ++j) {
    mask[boxes_num] |= dev_mask[i][j];
  }
}

  int const threadsPerBlock = sizeof(unsigned long long) * 8;
  const int col_blocks = DIVUP(boxes_num, threadsPerBlock);
  std::vector<unsigned long long> remv(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++) {
    // 计算当前box的索引
    int nblock = i / threadsPerBlock;
    int inblock = i % threadsPerBlock;

    // 如果当前box没有被删除，即remv[nblock]的inblock位为0
    if (!(remv[nblock] & (1ULL << inblock))) {
      keep_out[num_to_keep++] = i;
      // p为当前box结果数组的首地址
      unsigned long long *p = &mask_host[0] + i * col_blocks;
      // nblock前面的remv不用更新
      for (int j = nblock; j < col_blocks; j++) {
        remv[j] |= p[j];
      }
    }
  }

  dim3 blocks(DIVUP(boxes_num, threadsPerBlock),
              DIVUP(boxes_num, threadsPerBlock));
  dim3 threads(threadsPerBlock);
  nms_kernel<<<blocks, threads>>>(boxes_num,
                                  nms_overlap_thresh,
                                  boxes_dev,
                                  mask_dev);

__global__ void nms_kernel(const int n_boxes, const float nms_overlap_thresh,
                           const float *dev_boxes, unsigned long long *dev_mask) {
  // row_start表示当前基准box的起始block索引
  // col_start表示当前其他box的起始block索引
  const int row_start = blockIdx.y;
  const int col_start = blockIdx.x;

  // 如果其他box比基准box的置信度大则不用计算
  // if (row_start > col_start) return;

  // row_size是当前block基准box的大小
  // col_size是当前block其他box的大小
  // 因为griddim的大小经过向上取整,所以有些thread是不用计算的
  const int row_size =
        min(n_boxes - row_start * threadsPerBlock, threadsPerBlock);
  const int col_size =
        min(n_boxes - col_start * threadsPerBlock, threadsPerBlock);

  // 每个block内的thread加载对应的box进入共享内存
  // 共享内存访问速度比全局内存快很多，减少了耗时的全局内存访问
  __shared__ float block_boxes[threadsPerBlock * 5];
  if (threadIdx.x < col_size) {
    block_boxes[threadIdx.x * 5 + 0] =
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 0];
    block_boxes[threadIdx.x * 5 + 1] =
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 1];
    block_boxes[threadIdx.x * 5 + 2] =
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 2];
    block_boxes[threadIdx.x * 5 + 3] =
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 3];
    block_boxes[threadIdx.x * 5 + 4] =  
        dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 4];
  }
  __syncthreads();

  if (threadIdx.x < row_size) {
    // 计算当前基准box的索引
    const int cur_box_idx = threadsPerBlock * row_start + threadIdx.x;
    const float *cur_box = dev_boxes + cur_box_idx * 5;
    int i = 0;
    unsigned long long t = 0;
    int start = 0;
    // 如果相等，基准box和其他box是同样范围的，因此基准box不需要和置信度大于自己的其他box计算
    if (row_start == col_start) {
      start = threadIdx.x + 1;
    }
    // 遍历其他box，计算当前基准box和其他box的IOU，如果大于nms阈值，则置对应的位为1，即删除
    for (i = start; i < col_size; i++) {
      if (devIoU(cur_box, block_boxes + i * 5) > nms_overlap_thresh) {
        t |= 1ULL << i;
      }
    }
    const int col_blocks = DIVUP(n_boxes, threadsPerBlock);
    // 结果写入dev_mask
    dev_mask[cur_box_idx * col_blocks + col_start] = t;
  }
}