0

我有一个浮点数组，我想对其执行 stram 压缩操作，就像这里介绍的那样：Parallel Prefix Sum (Scan) with CUDA，然后根据值和地址或原始元素应用转换。

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

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

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

4

1 回答 1

4

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

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

然后，您会将上述变换迭代器传递给适当的推力流压缩算法以选择您想要的值。

这是一种可能的方法：

$ 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实现实际上可能会将数据加载操作推迟到做出模板决定之后。如果是这种情况，则可能仅根据需要进行转换。即使总是这样做，这也可能是一个无关紧要的问题，因为许多推力操作往往受加载/存储或内存带宽限制，而不是计算限制。然而，一个有趣的替代方法可能是使用 @ms这里创建的适配，它创建一个应用于输出迭代器步骤的转换，这可能会将转换操作限制为仅对实际保存在结果中的数据元素执行，尽管我也没有仔细检查过。

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

于 2016-01-14T15:19:14.803 回答