文章目录
本文为PointNet++ CUDA代码阅读系列的第四部分,其他详见:
(一)PointNet++代码梳理
(二)PointNet++中的FPS的CUDA实现
(三)PointNet++中ball query的CUDA实现
(四)PointNet++中的Three_nn的CUDA实现
给定点集known和unknown,Three_nn实现的功能是对于unknown的每个点,找到其在known中最临近的3个点的距离和下标,直接看cu代码,在src/interpolate_gpu.cu中:
__global__ void three_nn_kernel_fast(int b, int n, int m, const float *__restrict__ unknown,
const float *__restrict__ known, float *__restrict__ dist2, int *__restrict__ idx) {
// unknown: (B, N, 3)
// known: (B, M, 3)
// output:
// dist2: (B, N, 3)
// idx: (B, N, 3)
int bs_idx = blockIdx.y; // 找到这个线程处理的batch
int pt_idx = blockIdx.x * blockDim.x + threadIdx.x; // 找到这个线程处理的点
if (bs_idx >= b || pt_idx >= n) return;
unknown += bs_idx * n * 3 + pt_idx * 3; // unknown为指针,定位到这个线程处理的unknown中的点
known += bs_idx * m * 3; // known为指针,定位到这个线程处理的known的batch的起始位置
dist2 += bs_idx * n * 3 + pt_idx * 3; // dist2为指针,定位到这个线程处理的unknown的点所对应输出dist2的起始位置
idx += bs_idx * n * 3 + pt_idx * 3; // idx为指针,定位到这个线程处理的unknown的点所对应输出idx的起始位置
float ux = unknown[0]; // 取出这个线程处理的unknown的点的坐标
float uy = unknown[1];
float uz = unknown[2];
double best1 = 1e40, best2 = 1e40, best3 = 1e40;
int besti1 = 0, besti2 = 0, besti3 = 0;
for (int k = 0; k < m; ++k) { // 遍历known中的每个点,找到最近的3个点的下标和距离
float x = known[k * 3 + 0];
float y = known[k * 3 + 1];
float z = known[k * 3 + 2];
float d = (ux - x) * (ux - x) + (uy - y) * (uy - y) + (uz - z) * (uz - z);
if (d < best1) {
best3 = best2; besti3 = besti2;
best2 = best1; besti2 = besti1;
best1 = d; besti1 = k;
}
else if (d < best2) {
best3 = best2; besti3 = besti2;
best2 = d; besti2 = k;
}
else if (d < best3) {
best3 = d; besti3 = k;
}
}
dist2[0] = best1; dist2[1] = best2; dist2[2] = best3;
idx[0] = besti1; idx[1] = besti2; idx[2] = besti3;
}