我目前正在编写一个使用 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)。