用CUDA模拟流水线程序

Question

假设我有两个数组 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 和另一个数据数组 D 上启动一个 kernel2。无论如何，我可以在计算 C 中的第一个块后立即 kernel2 开始吗？本质上，kernel1 管道结果为 kernel2。依赖关系树看起来像这样

       Result
       /  \
      C    D
    /  \    
   A    B

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

Answer 1

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

假设您希望重叠复制和计算操作。这通常意味着您会将输入数据分成 "chunks" 并将数据块复制到设备，然后启动计算操作。每次内核启动都会对 "chunk" 数据进行操作。

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

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

用CUDA模拟流水线程序

Simulating pipeline program with CUDA

cuda

gpu

gpgpu