CUDA：将不同线程中的向量堆栈到一维向量

Question

我在 CUDA 中的每个线程都有一个推力向量，我想按顺序堆叠向量（线程 0 中的向量，线程 1 中的向量，......以及线程 n 中的向量）以创建 1d 向量并发回 CPU 。有没有好的方法可以做到这一点？任何帮助表示赞赏。谢谢。

Answer 1

将多个线程中的项目存储到单个向量中的最高效方法是线程交错。假设 4 个线程 (t0-t3) 中的每一个都有 4 个要存储的元素 (e0-e3)。最有效的最终存储模式将是：

t0e0 t1e0 t2e0 t3e0 t0e1 t1e1 t2e1 t3e1 t0e2 t1e2 t2e2 t3e2 t0e3 t1e3 t2e3 t3e3

执行此操作的代码如下所示：

#include <thrust/device_vector.h>
const int nt = 4;
const int ne = 4;
template <typename T>
__global__ void k(T *d){

  T e0 = threadIdx.x+10;
  T e1 = threadIdx.x+20;
  T e2 = threadIdx.x+30;
  T e3 = threadIdx.x+40;
  d[threadIdx.x]      = e0;
  d[threadIdx.x+nt]   = e1;
  d[threadIdx.x+2*nt] = e2;
  d[threadIdx.x+3*nt] = e3;
}

int main(){

  thrust::device_vector<int> d(ne*nt);
  k<<<1,nt>>>(thrust::raw_pointer_cast(d.data()));
}

在您的问题中，您似乎想要这个顺序：

t0e0 t0e1 t0e2 t0e3 t1e0 t1e1 t1e2 t1e3 t2e0 t2e1 t2e2 t2e3 t3e0 t3e1 t3e2 t3e3

该存储模式通常效率较低，但您可以这样实现：

#include <thrust/device_vector.h>
const int nt = 4;
const int ne = 4;
template <typename T>
__global__ void k(T *d){

  T e0 = threadIdx.x+10;
  T e1 = threadIdx.x+20;
  T e2 = threadIdx.x+30;
  T e3 = threadIdx.x+40;
  d[threadIdx.x*nt]    = e0;
  d[threadIdx.x*nt+1]  = e1;
  d[threadIdx.x*nt+2]  = e2;
  d[threadIdx.x*nt+3]  = e3;
}

int main(){

  thrust::device_vector<int> d(ne*nt);
  k<<<1,nt>>>(thrust::raw_pointer_cast(d.data()));
}

这两种情况下的存储效率差异是未合并和合并行为之间的差异，cuda SO 标签上的许多问题都涵盖了这一点，例如 这个。