改进 CUDA 中的异步执行

Question

我目前正在编写一个程序，该程序使用 CUDA API 在 GPU 上执行大型模拟。为了提高性能，我尝试同时运行我的内核，然后再次将结果异步复制到主机内存中。代码大致如下：

#define NSTREAMS   8
#define BLOCKDIMX  16
#define BLOCKDIMY  16

void domainUpdate(float* domain_cpu,       // pointer to domain on host
                  float* domain_gpu,       // pointer to domain on device
                  const unsigned int dimX,
                  const unsigned int dimY,
                  const unsigned int dimZ)
{
    dim3 blocks((dimX + BLOCKDIMX - 1) / BLOCKDIMX, (dimY + BLOCKDIMY - 1) / BLOCKDIMY);
    dim3 threads(BLOCKDIMX, BLOCKDIMY);

    for (unsigned int ii = 0; ii < NSTREAMS; ++ii) {

        updateDomain3D<<<blocks,threads, 0, streams[ii]>>>(domain_gpu,
                                                           dimX, 0,  dimX - 1, // dimX, minX, maxX
                                                           dimY, 0,  dimY - 1, // dimY, minY, maxY
                                                           dimZ, dimZ * ii / NSTREAMS,  dimZ * (ii + 1) / NSTREAMS - 1); // dimZ, minZ, maxZ

        unsigned int offset = dimX * dimY * dimZ * ii / NSTREAMS;
        cudaMemcpyAsync(domain_cpu + offset ,
                        domain_gpu+ offset ,
                        sizeof(float) * dimX * dimY * dimZ / NSTREAMS,
                        cudaMemcpyDeviceToHost, streams[ii]);
    }

    cudaDeviceSynchronize();
}

总而言之，它只是一个简单的 for 循环，循环所有流（本例中为 8 个）并划分工作。这实际上是一个更快的交易（高达 30% 的性能提升），尽管可能比我希望的要少。我在 Nvidia 的 Compute Visual Profiler 中分析了一个典型的循环，执行如下：

从图中可以看出，内核确实重叠，但同时运行的内核不会超过两个。我对不同数量的流和不同大小的模拟域尝试了相同的操作，但情况总是如此。

所以我的问题是：有没有办法鼓励/强制 GPU 调度程序同时运行两个以上的东西？或者这是依赖于无法在代码中表示的 GPU 设备的限制？

我的系统规格是：64 位 Windows 7，以及 GeForce GTX 670 显卡（即 Kepler 架构，计算能力 3.0）。

Answer 1

仅当 GPU 有剩余资源可运行第二个内核时，内核才会重叠。 GPU 完全加载后，并行运行更多内核没有任何好处，因此驱动程序不会这样做。