CUDA Thrust与原始内核相比如何？

Question

我对GPU编程还不熟悉，不确定如何编写最有效的代码。使用Thrust和编写自定义内核并自己管理内存有什么优缺点？

如果需要详细说明我的目标是什么：我有一个大矩阵，在每个值上需要执行几个向量操作。我知道这个任务需要动态并行性，并且目前有一个自定义内核来遍历矩阵并调用其他内核。我正在考虑是否应该用Thrust调用（例如thrust::for_each）替换内核，以及我是否应该在内核中使用Thrust进行向量操作。

- Reg

Answer 1

在过去的约12个月中，我从主要编写CUDA内核转变为主要使用Thrust，然后又回到了主要编写CUDA内核。一般来说，编写自己的CUDA内核应该提供更好的原始性能，但在简单的测试用例中，差异应该是可以忽略不计的。

Thrust模仿C++ STL，因此它具有与STL相同的优缺点。即，它旨在以一种非常通用的方式操作数据向量。从这个角度来看，Thrust在某些方面比CUDA更好，但不应被视为一种万能解决方案。 Thrust的主要优势在于抽象和可移植性; 您不必考虑块大小，并且很容易编写对设备或主机上的数据都适用的函数对象，而显然CUDA内核只能处理设备内存上的数据。它还具有许多非常有用的算法；不必编写自己的约减或排序算法，因为Thrust提供了这些的非常有效的实现。但是，在底层，您的数据访问模式可能不容易与Thrust设计的匹配，而且Thrust倾向于执行大量的临时内存分配（在性能上下文中通常不好; 您可以通过修改其内存管理模型来缓存这些临时分配，但我不建议这样做，只需编写内核并自己完全控制您的内存使用情况）。

我目前偏好的工作模式是几乎所有事情都使用CUDA，但对于特定算法（例如排序）会使用Thrust的算法，用于原型代码或者希望在主机和设备上同样良好运行的代码。

- Michael

---

仅使用Thrust进行内存管理的想法如何？ - Richard

@Richard，我之前一直在这样做，但是我遇到了一个很严重的问题（我认为是CUDA 10.0的问题？），当我尝试重复使用一个已经显著增长过的thrust向量时，将指针传递给CUDA内核会导致分段错误。我没有费心去调试这个问题，我放弃了向量并开始自己管理所有内存。从长远来看，这更可取，因为我想使用异步设备传输，而这不是你可以使用thrust向量完成的。如果我确定不需要数组增长，那么我仍然使用向量，它们很好用。 - Michael

@Richard：我的想法是，有人需要编写一个像样的内存管理库，无论是仅针对CUDA还是更普遍地，都不会束缚你以特定的方式编写代码。即不使用STL分配器，也不使用需要thrust设备向量等东西。 - einpoklum