【问题标题】:Data Pre-fetching to Overlap Memory Access and Computation CUDA数据预取以重叠内存访问和计算 CUDA
【发布时间】:2014-05-15 17:21:14
【问题描述】:

为了优化我的 cuda 内核的性能,我尝试将全局内存中的数据预取(或预读)到我的内核中,以尝试重叠内存带宽和计算。

我的实现基于以下逻辑:

// Original code
for (i = 0; i < N; i++) {
    sum += array[i];
}

// Code with pre-fetch
temp = array[0];
for (i = 0; i < N-1; i++) {
    temp2 = array[i+1];
    sum += temp;
    temp = temp2;
}
sum += temp;

这个想法是尝试同时执行 temp2 = array[i+1] 和 sum+=temp,这样我就可以将当前元素的执行与下一个元素的数据获取重叠。

与原来的相比,这样的实现在加速方面产生了非常小的提升(大约 3%)。我想知道是否有明确的方法告诉 cuda 重叠执行这两行代码?如何确保这些行同时运行,从而在读取全局内存和执行某些计算之间实现良好的重叠

感谢您抽出宝贵时间阅读我的问题。

【问题讨论】:

    标签: cuda


    【解决方案1】:

    您的示例只能显示性能的最小提升(如果有的话),因为这两个操作(从全局内存加载和执行求和)具有非常不同的速度:访问全局内存要慢得多。

    典型的方法是将数据从全局内存加载到共享内存中,而不是对位于那里的数据执行大量计算(注意:共享内存相当于共享 L1 缓存,在性能:与 L1 缓存的主要区别在于它必须被显式寻址)。

    典型的例子(见this link in the Cuda C Programming Guide)是矩阵乘法:你在共享内存中加载两个矩阵的N×N方块(即N*N个元素),并执行部分乘积(即N*N*N 次操作):每个元素 N 次操作。

    【讨论】:

    • 非常感谢您的回复@Sigismondo!但是,我的应用程序每次全局写入只需要 1 次全局读取,因此使用共享内存只会使其变慢(因为这将涉及 1 次全局内存读取 + 1 次共享内存读取每次写入操作)。有没有其他方法可以提高此类应用程序的性能?
    • 实际上并不多...如果您的数据来自 CPU,主要瓶颈将是通过 PCI Express 总线(cudaMemcpy() 的)传输。如果没有,您可以尝试通过正确合并访问来尽力而为,但由于内存瓶颈,您将远离 GPU 的理论性能。
    猜你喜欢
    • 1970-01-01
    • 2016-09-30
    • 1970-01-01
    • 2015-03-24
    • 2017-03-01
    • 2020-10-23
    • 2018-09-22
    • 2018-08-28
    • 1970-01-01
    相关资源
    最近更新 更多