CUB (CUDA UnBound)等效推力：：

Question

由于推力库的一些性能问题(更多细节请参见此页 )，我计划重构一个CUDA应用程序，以使用CUB而不是with。特别是，替换thrust::sort_by_key和thrust::inclusive_scan调用)。在我的应用程序中，我需要按键对3个数组进行排序。我就是这样用推力做到的：

thrust::sort_by_key(key_iter, key_iter + numKeys, indices);
thrust::gather_wrapper(indices, indices + numKeys, 
      thrust::make_zip_iterator(thrust::make_tuple(values1Ptr, values2Ptr, values3Ptr)),
      thrust::make_zip_iterator(thrust::make_tuple(valuesOut1Ptr, valuesOut2Ptr, valuesOut3Ptr))
);

哪里

• key iter是一个指向我想要排序的键的thrust::device_ptr
• indices指向设备内存中的序列(从0到numKeys-1)。
• values{1,2,3}Ptr是我想排序的值的device_ptrs
• values{1,2,3}OutPtr是排序值的device_ptrs

使用幼崽SortPairs函数，我可以对单个值缓冲区进行排序，但并不是所有的3个都是一次性的。问题是，我没有看到任何幼崽“聚集”的公用事业。有什么建议吗？

编辑：

我想我可以实现我自己的聚集内核，但是除了这样做之外，还有什么更好的方法吗？

template <typename Index, typename Value> 
__global__ void  gather_kernel(const unsigned int N, const Index * map, 
const Value * src, Value * dst) 
{ 
    unsigned int i = blockDim.x * blockIdx.x + threadIdx.x; 
    if (i < N) 
    { 
        dst[i] = src[map[i]]; 
    } 
} 

这些未合并的负载和商店让我感到兴奋，但如果没有map上已知的结构，这可能是不可避免的。

Answer 1

您想要实现的目标似乎取决于thrust::zip_iterator。你可以

1. 只将thrust::sort_by_key替换为cub::DeviceRadixSort::SortPairs并保留thrust::gather，或
2. 在使用values{1,2,3}之前将cub::DeviceRadixSort::SortPairs压缩到结构数组中

更新

阅读了thrust::gather的实现后，

$CUDA_HOME/include/thrust/system/detail/generic/gather.inl

您可以看到，它只是一个天真的内核，如

__global__ gather(int* index, float* in, float* out, int len) {
  int i=...;
  if (i<len) { out[i] = in[index[i]]; }
}

然后，我认为上面的代码可以被单个内核替换，而无需花费太多的精力。

在这个内核中，您可以首先使用CUB块wize原语cub::BlockRadixSort<...>::SortBlockedToStriped来获取存储在寄存器中的排序索引，然后以thrust::gather的形式执行简单的重排序副本来填充values{1,2,3}Out。

使用SortBlockedToStriped而不是Sort可以在复制values时进行合并写入(但不是为了阅读)。