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 复制到主机并使用 simd for 循环计算偏移量是最简单、最快的方法：

// 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

我认为您正在寻找的算法是 prefix sum。向量上的前缀和生成另一个向量，其中包含输入向量的累积和值。前缀和至少存在两种变体——独占扫描或包含扫描。从概念上讲，它们是相似的。

如果您的 element_sizes 向量已存储在 GPU 全局内存中（根据您的伪代码似乎是这种情况），那么 GPU 上存在运行库函数，您可以在那个时候调用，以生成 memory_offsets 数据（向量），并且 memory_size 值可以从向量中的最后一个值中简单地获得，根据您是否正在做一个包含性的略有变化扫描或独占扫描。

这是一个使用 thrust 的简单示例：

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

CUDA：在和减少期间计算所有部分和的方法

CUDA: method to calculate all partial sums during a sum reduction

cuda

thrust