sorting - CUDA:在排序功能中使用比较器时，Thrust 如何管理内存？

Question

我有一个 char array我想作为参数传递给比较器的 10 个字符，该比较器将被 Thrust 的排序函数使用。

为了为这个数组分配内存，我使用 cudaMalloc .然而cudaMalloc在全局内存中分配内存，所以每当一个线程想要从这个数组中读取数据时，它就必须访问全局内存。

但是这个数组很小，我相信如果将它存储到某个共享内存甚至每个线程的寄存器中会更有效。但是，是否可以通过 Thrust 实现这一点，如果是，如何实现？

这是比较器:

struct comp{
   int *data_to_sort;
   char *helpingArray;

   comp(int *data_ptr) this->data_to_sort = data_ptr;

   __host__ __device__
      bool operator()(const int&a, const int&b){

            //use helpingArray to do some comparisons and
           // return true/false accordingly

      }
 }

然后我为 helpingArray 分配内存在全局内存中并将其作为参数传递给 Comparator struct 到排序函数。

请注意 data_to_sort数组存储在全局内存中，因为它包含需要排序的数据，我们无法避免这种情况的发生。

这很好用，排序方法比 cpu 排序方法快，但是我相信如果我避免存储 helpingArray在全局内存中，排序方法将变得更快。

Answer 1

我同意放置 helpingArray进入全局内存几乎没有意义，并且至少在某种程度上降低了性能。执行内核的推力后端是“封闭的”，并且不公开共享内存或寄存器等内核级功能，因此不能直接使用这些功能。

话虽如此，您可能可以做两件事来改善这一点。第一个是像这样声明你的仿函数:

struct comp{
   char helpingArray[10];

   __host__ __device__
      bool operator()(const int&a, const int&b){ ... }
 }

您可以填写 helpingArray在将仿函数传递给您正在使用的推力算法之前的主机代码中(请注意，仿函数是按值传递的，因此这是完全合法的)。在这种情况下， helpingArray可能会在线程本地内存中结束。这样做可能会或可能不会提高性能。当然，它极大地简化了支持事物所需的主机代码。

另一种选择是声明 helpingArray在 __constant__ 内存中，只需在仿函数内部引用它。如果每个线程的访问模式是统一的，那么由于常量缓存，这样做可能会提高性能。