c – 改进CUDA中的异步执行

我目前正在编写一个程序,使用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,streams[ii]>>>(domain_gpu,dimX,dimX - 1,// dimX,minX,maxX
                                                           dimY,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).

解决方法

仅当GPU有剩余资源运行第二个内核时,内核才会重叠.一旦GPU完全加载,并行运行更多内核就没有收获,因此驱动程序不会这样做.