CUDA中基于索引的流压缩和转换

Question

我有一个浮点数组，我想对其执行 stram 压缩操作，如下所示：Parallel Prefix Sum (Scan) with CUDA，然后根据值和地址或原始地址应用转换元素.

例如，我有一个值为 {10,-1, -10, 2} 的数组，我想 return 绝对值大于 5 的所有元素，并应用一个transform 获取值及其在数组中的地址。这里的结果是 {transform(10,0),transform(-10,2)}。

我正在尝试对此使用 thrust，但此代码将运行经常用于大型数组，因此理想情况下它不会使用缓冲区和数组的多次遍历。

是否可以在不分配二级数组并进行多次遍历的情况下做我想做的事情？如果是，这样的代码是否存在于野外？或者至少有人知道我可以编写哪些 thrust 函数或任何其他库来实现我的目标吗？

Answer 1

是的，可以通过单个推力算法调用推力（我假设这就是你所说的 "without ... doing multiple traversals" 的意思）并且没有 "allocating a secondary array".

一种方法是将数据数组加上索引/"address" 数组（通过 thrust::counting_iterator，这避免了分配）传递给 thrust::transform_iterator，它创建了您的 "transform" 操作（结合适当的函子）。

然后您可以将上述转换迭代器传递给适当的 thrust stream compaction algorithm 到 select 您想要的值。

这是一个可能的方法：

$ cat t1044.cu
#include <thrust/device_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/copy.h>
#include <math.h>

#include <iostream>

__host__ __device__ int my_transform(int data, int idx){
  return (data - idx);  //put whatever transform you want here
}

struct my_transform_func : public thrust::unary_function<thrust::tuple<int, int>, int>
{

  __host__ __device__
  int operator()(thrust::tuple<int, int> &t){
    return my_transform(thrust::get<0>(t), thrust::get<1>(t));
    }
};

struct my_test_func
{
  __host__ __device__
  bool operator()(int data){
    return (abs(data) > 5);
    }
};



int main(){

  int data[] = {10,-1,-10,2};
  int dsize = sizeof(data)/sizeof(int);

  thrust::device_vector<int> d_data(data, data+dsize);
  thrust::device_vector<int> d_result(dsize);
  int rsize = thrust::copy_if(thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(d_data.begin(), thrust::counting_iterator<int>(0))), my_transform_func()), thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(d_data.end(), thrust::counting_iterator<int>(dsize))), my_transform_func()), d_data.begin(), d_result.begin(), my_test_func()) - d_result.begin();
  thrust::copy_n(d_result.begin(), rsize, std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  return 0;
}
$ nvcc -o t1044 t1044.cu
$ ./t1044
10,-12,
$

对这种方法的一些可能的批评：

它似乎加载了两次 d_data 元素（一次用于转换操作，一次用于模板）。但是，CUDA 优化编译器可能会识别最终生成的线程代码中的冗余负载，并将其优化掉。
看起来我们正在对每个数据元素执行转换操作，无论我们是否打算将其保存在结果中。再次强调，thrust copy_if 实现实际上可能会推迟数据加载操作，直到做出模板决定之后。如果是这种情况，则转换可能仅在 as-needed 基础上完成。即使总是这样做，这也可能是一个无关紧要的问题，因为许多推力操作往往是 load/store 或内存带宽限制，而不是计算限制。然而，一个有趣的替代方法可能是使用由@m.s 创建的改编。创建应用于 输出迭代器步骤 的转换，这可能会将转换操作限制为仅对实际保存在结果中的数据元素执行，虽然我也没有仔细检查过。
正如下面的评论中所提到的，这种方法确实分配了临时存储（thrust 在幕后这样做，作为 copy_if 操作的一部分）当然我明确地分配了结果的 O(n) 存储。我怀疑推力分配（单个 cudaMalloc）可能也用于 O(n) 存储。虽然完全不需要任何额外存储就可以完成所有要求的事情（并行前缀和、流压缩、数据转换）（所以也许请求是针对 in-place 操作），但我认为以这种方式设计算法可能会对性能产生重大负面影响，如果它完全可行的话（我不清楚并行前缀和可以在绝对没有任何类型的额外存储的情况下实现，更不用说将其与流压缩耦合即并行数据移动）。由于 thrust 释放了它使用的所有此类临时存储，因此不会有太多与频繁使用此方法相关的存储问题。唯一剩下的问题（我猜）是性能。如果性能是一个问题，那么与临时分配相关的时间开销应该通过将上述算法与 thrust custom allocator (also see here) 耦合来消除，这将分配一次所需的最大存储缓冲区，然后 re-use每次使用上述算法时该缓冲区。

CUDA中基于索引的流压缩和转换

Stream compaction and transform based on the index in CUDA

algorithm

optimization

cuda

stream-compaction