如何使用thrust和CUDA流异步地将内存从主机复制到设备

Question

我希望使用Thrust将内存从主机复制到设备上，就像下面这样：

thrust::host_vector<float> h_vec(1 << 28);
thrust::device_vector<float> d_vec(1 << 28);
thrust::copy(h_vec.begin(), h_vec.end(), d_vec.begin());

使用CUDA流与使用流将内存从设备复制到设备时类似。

cudaStream_t s;
cudaStreamCreate(&s);

thrust::device_vector<float> d_vec1(1 << 28), d_vec2(1 << 28);
thrust::copy(thrust::cuda::par.on(s), d_vec1.begin(), d_vec1.end(), d_vec2.begin());

cudaStreamSynchronize(s);
cudaStreamDestroy(s);

问题在于我无法将CUDA的执行策略设置为指定从主机到设备时要使用的流，因为这种情况下，thrust会假定两个向量都存储在设备上。有没有解决这个问题的方法？我正在使用来自GitHub的最新thrust版本（version.h文件中标明为1.8）。

- Tom

---

我读到的公告让我感觉流被实现用于底层内核调用，并非在整个thrust中都使用。如果您确实使用了流从主机向设备的传输，那么您可能需要在主机上使用pinned allocator。因此，我认为您所要求的可以通过使用thrust向量和cudaMemcpyAsync来完成。 - Robert Crovella

是的，正如Robert建议的那样，您应该直接使用cudaMemcpyAsync。 - Jared Hoberock

截至今天（2016年5月），我在这里的文档中找到了第一个条目：http://thrust.github.io/doc/group__copying.html#ga3e43fb8472db501412452fa27b931ee2 ，真的让人感到不安。它说我们可以将thrust :: copy（thrust :: cuda :: par.on（cudaStream），HostPtr，HostPtr + size，DevicePtr）编写为有效的语法，而不会异步地将您的副本发布到传递给参数的特定流...... - Tobbey

Answer 1

正如在评论中所指出的，我不认为使用 thrust::copy 直接实现这一目标是可能的。然而，在 thrust 应用程序中，我们可以使用 cudaMemcpyAsync 实现异步复制和复制与计算重叠的目标。

以下是一个具体的示例：

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/system/cuda/experimental/pinned_allocator.h>
#include <thrust/system/cuda/execution_policy.h>
#include <thrust/fill.h>
#include <thrust/sequence.h>
#include <thrust/for_each.h>
#include <iostream>

// DSIZE determines duration of H2D and D2H transfers
#define DSIZE (1048576*8)
// SSIZE,LSIZE determine duration of kernel launched by thrust
#define SSIZE (1024*512)
#define LSIZE 1
// KSIZE determines size of thrust kernels (number of threads per block)
#define KSIZE 64
#define TV1 1
#define TV2 2

typedef int mytype;
typedef thrust::host_vector<mytype, thrust::cuda::experimental::pinned_allocator<mytype> > pinnedVector;

struct sum_functor
{
  mytype *dptr;
  sum_functor(mytype* _dptr) : dptr(_dptr) {};
  __host__ __device__ void operator()(mytype &data) const
    {
      mytype result = data;
      for (int j = 0; j < LSIZE; j++)
        for (int i = 0; i < SSIZE; i++)
          result += dptr[i];
      data = result;
    }
};

int main(){

  pinnedVector hi1(DSIZE);
  pinnedVector hi2(DSIZE);
  pinnedVector ho1(DSIZE);
  pinnedVector ho2(DSIZE);
  thrust::device_vector<mytype> di1(DSIZE);
  thrust::device_vector<mytype> di2(DSIZE);
  thrust::device_vector<mytype> do1(DSIZE);
  thrust::device_vector<mytype> do2(DSIZE);
  thrust::device_vector<mytype> dc1(KSIZE);
  thrust::device_vector<mytype> dc2(KSIZE);

  thrust::fill(hi1.begin(), hi1.end(),  TV1);
  thrust::fill(hi2.begin(), hi2.end(),  TV2);
  thrust::sequence(do1.begin(), do1.end());
  thrust::sequence(do2.begin(), do2.end());

  cudaStream_t s1, s2;
  cudaStreamCreate(&s1); cudaStreamCreate(&s2);

  cudaMemcpyAsync(thrust::raw_pointer_cast(di1.data()), thrust::raw_pointer_cast(hi1.data()), di1.size()*sizeof(mytype), cudaMemcpyHostToDevice, s1);
  cudaMemcpyAsync(thrust::raw_pointer_cast(di2.data()), thrust::raw_pointer_cast(hi2.data()), di2.size()*sizeof(mytype), cudaMemcpyHostToDevice, s2);

  thrust::for_each(thrust::cuda::par.on(s1), do1.begin(), do1.begin()+KSIZE, sum_functor(thrust::raw_pointer_cast(di1.data())));
  thrust::for_each(thrust::cuda::par.on(s2), do2.begin(), do2.begin()+KSIZE, sum_functor(thrust::raw_pointer_cast(di2.data())));

  cudaMemcpyAsync(thrust::raw_pointer_cast(ho1.data()), thrust::raw_pointer_cast(do1.data()), do1.size()*sizeof(mytype), cudaMemcpyDeviceToHost, s1);
  cudaMemcpyAsync(thrust::raw_pointer_cast(ho2.data()), thrust::raw_pointer_cast(do2.data()), do2.size()*sizeof(mytype), cudaMemcpyDeviceToHost, s2);

  cudaDeviceSynchronize();
  for (int i=0; i < KSIZE; i++){
    if (ho1[i] != ((LSIZE*SSIZE*TV1) + i)) { std::cout << "mismatch on stream 1 at " << i << " was: " << ho1[i] << " should be: " << ((DSIZE*TV1)+i) << std::endl; return 1;}
    if (ho2[i] != ((LSIZE*SSIZE*TV2) + i)) { std::cout << "mismatch on stream 2 at " << i << " was: " << ho2[i] << " should be: " << ((DSIZE*TV2)+i) << std::endl; return 1;}
    }
  std::cout << "Success!" << std::endl;
  return 0;
}

我的测试用例使用的是RHEL5.5，Quadro5000和cuda 6.5RC。这个示例旨在让thrust创建非常小的内核（仅有一个线程块，只要KSIZE很小，比如32或64），以使从thrust::for_each创建的内核能够并发运行。

当我分析这段代码时，我看到：

这表明我们正在实现thrust内核之间以及复制操作和thrust内核之间的适当重叠，并且在内核完成时异步进行数据复制。请注意，cudaDeviceSynchronize()操作“填充”时间轴，表明所有异步操作（数据复制、thrust函数）都是异步发出的，控制在任何操作开始之前就返回到主机线程。所有这些都是完全实现主机、GPU和数据复制操作之间并发的预期且正确的行为。

- Robert Crovella

---

在thrust::host_vector上使用cudaMemcpyAsync时，是否需要使用pinned_allocator才能使其正常工作？如果我使用标准的thrust::host_vector会发生什么？ - m.s.

标准的 host_vector 使用未固定页面（即未锁页）的分配器。 这意味着当您尝试执行 cudaMemcpyAsync 时，操作不会是异步的。如果您正在尝试将该操作与其他内容重叠，它也不会重叠。请参见这里和这里。 - Robert Crovella

请注意，随CUDA 7一起发布的Thrust版本存在问题，在某些情况下会阻止正确地将Thrust内核发出到流中。解决方法是：1. 将CUDA 7上的Thrust更新为当前的开发版本（其中包括修复此问题的修复程序），或2. 回退到CUDA 6.5（或升级到可用时的某个未来的CUDA工具包版本）。 - Robert Crovella

这个问题在7.5中已经修复了吗？还是仍然存在？ - Bar

是的，CUDA 7.5附带的推力版本已经修复了这个问题。 - Robert Crovella

Answer 2

这里有一个使用 thrust::cuda::experimental::pinned_allocator<T> 的示例：

一个已工作的示例：

// Compile with:
// nvcc --std=c++11 mem_async.cu -o mem_async

#include <cuda.h>
#include <cuda_runtime.h>
#include <cufft.h>

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/fill.h>
#include <thrust/system/cuda/experimental/pinned_allocator.h>

#define LEN 1024

int main(int argc, char *argv[]) {
    thrust::host_vector<float, thrust::cuda::experimental::pinned_allocator<float>> h_vec(LEN);
    thrust::device_vector<float> d_vec(LEN);

    thrust::fill(d_vec.begin(), d_vec.end(), -1.0);

    cudaMemcpyAsync(thrust::raw_pointer_cast(h_vec.data()),
                    thrust::raw_pointer_cast(d_vec.data()),
                    d_vec.size()*sizeof(float),
                    cudaMemcpyDeviceToHost);

    // Comment out this line to see what happens.
    cudaDeviceSynchronize();

    std::cout << h_vec[0] << std::endl;
}

注释掉同步步骤，由于异步内存传输，应该会在控制台上打印0。

- Andreas Yankopolus