使用 CUDA 模拟管道程序答案

【问题标题】：Simulating pipeline program with CUDA使用 CUDA 模拟管道程序
【发布时间】：2016-04-04 08:16:30
【问题描述】：

假设我有两个数组A 和B 和一个kernel1，它通过将数组分成不同的块并将部分结果写入C 对两个数组进行一些计算（例如向量加法） . kernel1 然后继续这样做，直到处理完数组中的所有元素。

unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
unsigned int gridSize = blockDim.x*gridDim.x;

//iterate through each chunk of gridSize in both A and B
while (i < N) {
     C[i] = A[i] + B[i];
     i += gridSize;
}

说，现在我想在C 上启动一个kernel2 和另一个数据数组D。在计算C 中的第一个块之后，我是否可以开始kernel2 立即？从本质上讲，kernel1 管道它导致kernel2。依赖树看起来像这样

       Result
       /  \
      C    D
    /  \    
   A    B

我曾考虑过使用 CUDA 流，但不确定具体如何。也许将主机纳入计算？

【问题讨论】：

标签： cuda gpgpu gpu

【解决方案1】：

是的，在这种情况下，您可以使用CUDA streams 来管理顺序和依赖关系。

假设您想要重叠复制和计算操作。这通常意味着您将输入数据分解为“块”，并将块复制到设备，然后启动计算操作。每次内核启动都在一个“块”数据上运行。

我们可以在宿主代码中使用循环来管理进程：

// create streams and ping-pong pointer
cudaStream_t stream1, stream2, *st_ptr;
cudaStreamCreate(&stream1); cudaStreamCreate(&stream2);
// assume D is already on device as dev_D
for (int chunkid = 0; chunkid < max; chunkid++){
  //ping-pong streams
  st_ptr = (chunkid % 2)?(&stream1):(&stream2);
  size_t offset = chunkid*chunk_size;
  //copy A and B chunks
  cudaMemcpyAsync(dev_A+offset, A+offset, chksize*sizeof(A_type), cudaMemcpyHostToDevice, *st_ptr);
  cudaMemcpyAsync(dev_B+offset, B+offset, chksize*sizeof(B_type), cudaMemcpyHostToDevice, *st_ptr);
  // then compute C based on A and B
  compute_C_kernel<<<...,*st_ptr>>>(dev_C+offset, dev_A+offset, dev_B+offset, chksize);
  // then compute Result based on C and D
  compute_Result_kernel<<<...,*st_ptr>>>(dev_C+offset, dev_D, chksize);
  // could copy a chunk of Result back to host here with cudaMemcpyAsync on same stream
  }

向同一流发出的所有操作都保证在设备上按顺序（即顺序）执行。发布到单独流的操作可以重叠。因此上述顺序应该：

将 A 块复制到设备
将 B 块复制到设备
启动内核以从 A 和 B 计算 C
启动内核以从 C 和 D 计算结果

上述步骤将针对每个块重复，但连续的块操作将发送到备用流。因此，块 2 的复制操作可以与块 1 的内核操作重叠，等等。

您可以通过查看有关 CUDA 流的演示文稿了解更多信息。 Here 就是一个例子。

较新的设备（Kepler 和 Maxwell）对于见证设备上操作重叠所需的程序发布顺序应该相当灵活。较旧的 (Fermi) 设备可能对发布顺序很敏感。你可以阅读更多关于here

【讨论】：