改善CUDA中的异步执行

我目前正在编写一个程序，使用CUDA API在GPU上执行大型模拟。为了加速性能，我试图同时运行我的内核，然后再次将结果异步复制到主机内存中。代码看起来大致是这样的：改善CUDA中的异步执行

#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(); 
}

总而言之，这只是一个简单的循环，遍历所有的流（8在这种情况下）和分工。这实际上是一个更快的交易（高达30％的性能增益），尽管可能比我希望的要少。我分析了Nvidia的计算视觉探查一个典型的周期，并执行如下：

CUDA API trace in the Compute Visual Profiler

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

所以我的问题是：有没有办法鼓励/强制GPU调度程序在同一时间运行两件事？或者，这是否取决于代码中无法表示的GPU设备的限制？我的系统规格为：64位Windows 7和GeForce GTX 670图形卡（即开普勒架构，计算能力3.0）。

来源

2013-04-25 Yellow

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

来源

2013-04-25 11:30:37 tera

即使有一个很小的内核，就像一些块一样，同一时间内也不会有超过两个内核运行。所以GPU的物理尺寸不可能是整个故事，可以吗？ – Yellow 2013-04-25 12:18:16

是的，它可以。什么是“小内核”？多少块？每块有多少个线程？他们使用共享内存吗？寄存器？除非你分析了内核的资源利用率，否则你不知道可以运行多少个内核。 Windows（GPU处于WDDM模式时）也可以通过批量GPU活动来影响并发性。 GPU不限于同时运行两件事。 – 2013-04-25 13:14:51

这是一个很好的观点，我没有完全想到所有的共享内存和注册要求，我不明白这会影响性能。我尝试的一个'小'内核是例如8x8块和16x16线程。其中，理论上适合GPU的负载更多，我会说。它每个线程使用33个寄存器，每块大约2 kB共享内存。这很多吗？ – Yellow 2013-04-26 13:19:05

改善CUDA中的异步执行

回答

相关问题