4

我有char array10 个字符,我想将其作为参数传递给比较器,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结构的参数传递给排序函数。

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

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

4

1 回答 1

1

我同意放入helpingArray全局内存没有什么意义,并且至少在某种程度上降低了性能。执行内核的推力后端是“封闭的”,不会暴露内核级功能,如共享内存或寄存器,因此不能直接使用。

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

struct comp{
   char helpingArray[10];

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

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

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

于 2013-06-23T15:57:03.307 回答