如何使用 CUDA thrust 进行分段缩减？

Question

我想将部分缩减结果存储在数组中。

说我有

data[8] = {10,20,30,40,50,60,70,80}

.
如果我将

data

与

chunk_size

的

2

分开，块将是

{10,20}

,

{30,40}

, ... ,

{70,80}

.

如果我以求和为目标，总的减少量将是

360

但我想得到一个

partial_sums = {30,70,110,150}

的数组，它存储每个块的部分和。

到目前为止，我的想法是构建一个迭代器

strided_iterator

，它将访问0、2、...

data[8] = {10,20,30,40,50,60,70,80}

的第一个索引以及类似的东西

thrust::reduce(stride_iterator, stride_iterator + 2,
               partial_sums.begin(),
               thrust::plus<int>());

给出了想要的结果，但不知道如何有效地完成。

对于strided access，

thrust/examples/strided_range.cu

有一个解决方案，但是这似乎不适用于store segmented reductions。

我当然可以用这样的循环来残酷地做，

for (int i = 0; i<4; i++) {
  partial_sums[i] = thrust::reduce(data+2*i, data+2*i+2, 0, thrust::plus<int>());
}

但这种做法正是CUDA thrust力求尽量避免的，对吧？不知何故，我应该能够将所有内容放在一个 Thrust 调用中。