在 ROS 上下文中使用基于 CUDA 的函数的正确方法

Question

我正在开发一个基于 ROS 的管道，其主要功能是订阅图像主题并连续执行特征检测、匹配等功能。为了使这个管道更快，我尝试使用基于 CUDA 的检测和匹配模块作为我的包的一部分。在这个问题的上下文中，我假设一个简单的管道，我订阅一个图像主题，并且在每次图像可用时调用的订阅者回调中，调用不同类的两个成员函数：一个用于检测，另一个匹配，每个都包含自己的 CUDA 内核。有点类似于循环执行这两个函数。

第一个函数获取图像数据并计算特征关键点和描述符，然后将它们返回给主机。然后我将这些描述符复制回 GPU 内存，我需要将它们与属于参考图像的另一组描述符进行匹配。

例如，独立匹配函数如下所示：

// Detection module returns a struct featureData, containing keypoints and descriptors 
// in featureData.kps and featureData.desc

uint64_t* d_desc;
cudaMalloc(&d_desc, 64 * featureData.kps.size());
cudaMemcpy(d_desc, &featureData., 64 * (featureData.kps.size()), cudaMemcpyHostToDevice);

cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);

// Create texture object for descriptors

struct cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeLinear;
resDesc.res.linear.devPtr = d_desc;
resDesc.res.linear.desc.f = cudaChannelFormatKindUnsigned;
resDesc.res.linear.desc.x = 32;
resDesc.res.linear.desc.y = 32;
resDesc.res.linear.sizeInBytes = 64 * featureData.kps.size();

struct cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = cudaAddressModeBorder;
texDesc.addressMode[1] = cudaAddressModeBorder;
texDesc.filterMode = cudaFilterModePoint;
texDesc.readMode = cudaReadModeElementType;
texDesc.normalizedCoords = 0;
cudaTextureObject_t tex_q = 0;
cudaCreateTextureObject(&tex_q, &resDesc, &texDesc, nullptr);

// Allocate space for match results
int* d_matches;
cudaMalloc(&d_matches, 4 * featureData.kps.size());

// Launch the matching kernel
CUDAmatch(d_descRef, static_cast<int>(refData.kps.size()), tex_q, static_cast<int>(featureData.kps.size()), d_matches, threshold);

// d_descRef is memory pointed to by a uint64_t* for the reference descriptors.

在这种情况下，我有几个问题，因为这是我第一次涉足基于 GPU 的开发。

1. 当需要执行匹配时，描述符等被复制到设备内存中，然后将结果复制回来。我是否应该在每次执行匹配后释放这些设备内存指针并在下一个回调中重新分配（cudaMalloc()）？描述符的长度将根据检测到的特征数量而变化。或者有没有更有效的方法来只分配一次内存并重用它？
2. 检测和匹配函数还使用cudaResourceDesc 和cudaTextureDesc 等对象，这些对象将在每次执行结束时超出范围，因此应该被销毁。我应该以任何其他特定方式处理它们吗？
3. 我假设在执行这两个函数后我需要cudaDeviceSynchronize()。我说的对吗？
4. 我能否安全地将“参考”描述符保留在 GPU 内存中，仅在需要时更新它们？

Answer 1

我是否应该在每次执行匹配后释放这些设备内存指针并在下一个回调中重新分配 (cudaMalloc())？

可能不会。这似乎是不必要且耗时的。

或者有没有更有效的方法来只分配一次内存并重用它？

大概吧。例如，您可以确定可能需要的最大大小，为其分配，然后将指向它的指针传递到您的事件处理循环中，并重用分配。

检测和匹配函数还使用了诸如 cudaResourceDesc 和 cudaTextureDesc 之类的对象，这些对象在每次执行结束时都会超出范围，因此应该被销毁。我应该以任何其他特定方式处理它们吗？

同样，您可以在更高的范围内创建它们，并将对它们的引用传递到您的事件处理系统中。但是，我认为这里的主要消费者将是填充纹理的数据副本以及纹理的绑定。无论如何，这些都必须重复（大概）。但是，如果您为纹理对象的后备存储进行了底层分配，那么您可能不需要为此重新分配，请参阅之前的 cmets。

我假设在执行这两个函数后我需要 cudaDeviceSynchronize()。我说的对吗？

对我来说，需要它并不明显。您还没有展示一个完整的示例，但是如果在您的函数结束时，存在从设备到主机的数据副本，那可能就足够了。 cudaMemcpy 是一个阻塞函数。如果您在 TX1/TX2 等物理/逻辑统一的内存情况下运行，那么是的，您可能需要一个同步点来确保数据在主机代码中使用之前是有效的。

我能否安全地将“参考”描述符留在 GPU 内存中，仅在需要时更新它们？

我不知道为什么不。由cudaMalloc 创建的分配在应用程序终止或使用cudaFree 显式释放之前不会“超出范围”。如果您将数据复制到这样的分配中，它应该在您的应用程序期间保持不变，除非您以某种方式覆盖它（或释放底层分配）。