cuda中内核的并行执行

Question

假设我有三个全局数组，它们已使用 cudaMemcpy 复制到 gpu 中，但 c 中的这些 gloabl 数组尚未使用 cudaHostAlloc 分配，以便分配页面锁定的内存，而不是它们是简单的全局分配。

int a[100],b [100],c[100];
cudaMemcpy(d_a,a,100*sizeof(int),cudaMemcpyHostToDevice);
cudaMemcpy(d_b,b,100*sizeof(int),cudaMemcpyHostToDevice);
cudaMemcpy(d_c,c,100*sizeof(int),cudaMemcpyHostToDevice);

现在我有 10 个内核，它们在单独的流中启动，以便同时运行，其中一些使用在 gpu 中复制的全局数组。 现在这些内核运行了 1000 次迭代。 他们不必在迭代期间将任何内容复制回主机。

但问题是它们不是并行执行，而是采用串行方式。

  cudaStream_t stream[3];

   for(int i=0;i<3;i++)cudaStreamCreate (&stream[i]);

   for(int i=0;i<100;i++){
       kernel1<<<blocks,threads,0,stream[0]>>>(d_a,d_b);
       kernel2<<<blocks,threads,0,strea[1]>>(d_b,d_c);
      kernal3<<<blocks,threads,0,stream[2]>>>(d_c,d_a);
    cudaDeviceSynchronize();
  }

我不明白为什么？

Answer 1

以这种方式发布的内核：

   for(int i=0;i<100;i++){
    kernel1<<<blocks,threads>>>(d_a,d_b);
    kernel2<<<blocks,threads>>>(d_b,d_c);
    kernal3<<<blocks,threads>>>(d_c,d_a);
    cudaDeviceSynchronize();
  }

将始终连续运行。为了让内核同时运行，它们必须发给separate CUDA streams。还有其他要求。阅读documentation。

您需要创建一些 CUDA 流，然后像这样启动您的内核：

   cudaStream_t stream1, stream2, stream3;
   cudaStreamCreate(&stream1); cudaStreamCreate(&stream2); cudaStreamCreate(&stream3);

   for(int i=0;i<100;i++){
    kernel1<<<blocks,threads,0,stream1>>>(d_a,d_b);
    kernel2<<<blocks,threads,0,stream2>>>(d_b,d_c);
    kernal3<<<blocks,threads,0,stream3>>>(d_c,d_a);
    cudaDeviceSynchronize();
  }

实际见证并发内核执行通常还需要资源利用率有限的内核。如果一个给定的内核将“填满”机器，由于大量的块，或每个块的线程，或共享内存使用，或其他一些资源使用，那么你实际上不会见证并发；机器里没有空间了。

您可能还想查看一些 CUDA 示例代码，例如 simpleStreams 和 concurrentKernels。