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

Question

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

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

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

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

Answer 1

是的，可以通过单个推力算法调用推力(我假设这就是您所说的“无需...进行多次遍历”的意思)并且无需“分配辅助数组”。

一种方法是将数据数组加上索引/“地址”数组(通过 thrust::counting_iterator，避免分配)传递给 thrust::transform_iterator 创建你的“转换”操作(与适当的仿函数结合)。

然后您可以将上述转换迭代器传递给适当的 thrust stream compaction algorithm选择所需的值。

这是一种可能的方法:

$ 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,
$

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

1. 它似乎加载了两次 d_data 元素(一次用于转换操作，一次用于模板)。但是，CUDA 优化编译器可能会识别最终生成的线程代码中的冗余负载，并将其优化掉。

2. 看起来我们正在对每个数据元素执行转换操作，无论我们是否打算将其保存在结果中。再一次，推力 copy_if 实现实际上可能会推迟数据加载操作，直到做出模板决定之后。如果真是这样，那么转换可能只在需要的基础上完成。即使总是这样做，这也可能是一个无关紧要的问题，因为许多推力操作往往受加载/存储或内存带宽限制，而不是计算限制。然而，一个有趣的替代方法可能是使用@m.s 创建的改编。 here它创建了一个应用于输出迭代器步骤的转换，这可能会将转换操作限制为仅对实际保存在结果中的数据元素执行，尽管我也没有仔细检查过.

3. 正如下面评论中提到的，这种方法确实分配了临时存储空间(推力在幕后这样做，作为 copy_if 操作的一部分)，当然我明确地分配了 O (n) 结果的存储。我怀疑推力分配(单个 cudaMalloc)可能也用于 O(n) 存储。虽然完全不需要任何类型的额外存储就可以完成所有要求的事情(并行前缀和、流压缩、数据转换)(所以也许请求是针对就地操作)，但我认为制作一个如果这种算法完全可行的话，它可能会对性能产生重大的负面影响(我不清楚并行前缀和可以在绝对没有任何类型的额外存储的情况下实现，更不用说将其与流压缩耦合，即数据平行移动)。由于 thrust 释放了它使用的所有此类临时存储，因此不会有太多与频繁使用此方法相关的存储问题。唯一剩下的问题(我猜)是性能。如果性能是一个问题，那么与临时分配相关的时间开销应该通过将上述算法与 thrust custom allocator 耦合来大部分消除。 (另见 here )，这将分配一次所需的最大存储缓冲区，然后在每次使用上述算法时重新使用该缓冲区。