CUDA:在求和过程中计算所有部分和的方法

Question

我在CUDA中一次又一次地遇到这个问题。我已经为一组元素做了一些GPU计算。这导致了一些具有线性意义的值(例如，就内存而言(：

element_sizes = [ 10, 100, 23, 45 ]

现在，对于GPU计算的下一阶段，我需要以下值：

memory_size = sum(element_sizes)
memory_offsets = [ 0, 10, 110, 133 ]

我可以使用NVIDIA提供的简化代码在GPU上计算80 gbps的memory_size。但是，我不能使用这段代码，因为它使用的是一种分支技术，不构成内存偏移数组。我尝试了很多方法，但我发现，简单地通过elements_sizes复制到主机并使用simdfor循环计算偏移量是最简单、最快的方法：

// in pseudo code
host_element_sizes = copy_to_host(element_sizes);
host_offsets = (... *) malloc(...);
int total_size = 0;
for(int i = 0; i < ...; ...){
host_offsets[i] = total_size;
total_size += host_element_sizes[i];
}
device_offsets = (... *) device_malloc(...);
device_offsets = copy_to_device(host_offsets,...);

---

然而，我现在已经做了很多次了，它开始成为一个瓶颈。这似乎是一个典型的问题，但我没有找到解决办法。

CUDA程序员解决这个问题的预期方法是什么？

Answer 1

我认为您正在寻找的算法是前缀和。一个向量上的前缀和产生另一个向量，该向量包含输入向量的累积和值。前缀和至少存在于两种变体中——排他性扫描或包容性扫描。从概念上讲，它们是相似的。

如果你的element_sizes矢量已经存储在GPU全局内存中(这似乎是基于你的伪代码的情况(，那么你可以在GPU上调用一些库函数，以产生memory_offsets数据(矢量(，而memory_size值可以从矢量中的最后一个值中轻松获得，根据您是进行包容性扫描还是排他性扫描，略有不同。

下面是一个使用推力的琐碎示例：

$ cat t319.cu
#include <thrust/scan.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/copy.h>
#include <iostream>

int main(){
const int element_sizes[] = { 10, 100, 23, 45 };
const int ds = sizeof(element_sizes)/sizeof(element_sizes[0]);
thrust::device_vector<int> dv_es(element_sizes, element_sizes+ds);
thrust::device_vector<int> dv_mo(ds);
thrust::exclusive_scan(dv_es.begin(), dv_es.end(), dv_mo.begin());
std::cout << "element_sizes:" << std::endl;
thrust::copy_n(dv_es.begin(), ds, std::ostream_iterator<int>(std::cout, ","));
std::cout << std::endl << "memory_offsets:" << std::endl;
thrust::copy_n(dv_mo.begin(), ds, std::ostream_iterator<int>(std::cout, ","));
std::cout << std::endl << "memory_size:" << std::endl << dv_es[ds-1] + dv_mo[ds-1] << std::endl;
}
$ nvcc -o t319 t319.cu
$ ./t319
element_sizes:
10,100,23,45,
memory_offsets:
0,10,110,133,
memory_size:
178
$