CUDA全局内存事务的成本

2024-03-31

根据 CUDA 5.0 编程指南,如果我同时使用 L1 和 L2 缓存(在 Fermi 或 Kepler 上),则所有全局内存操作都使用 128 字节内存事务完成。但是,如果我仅使用 L2,则使用 32 字节内存事务(第 F.4.2 章)。

让我们假设所有缓存都是空的。如果我有一个扭曲,每个线程以完美对齐的方式访问单个 4 字节字,这将导致在 L1+L2 情况下产生 1x128B 事务,在仅 L2 情况下产生 4x32B 事务。是对的吗?

我的问题是 - 4 个 32B 事务是否比单个 128B 事务慢? 我对费米之前硬件的直觉表明它会更慢,但也许在较新的硬件上情况不再如此?或者也许我应该只看带宽利用率来判断内存访问的效率?


是的,在缓存模式下,将生成单个 128 字节事务(从 L1 缓存级别可见)。在非缓存模式下,将生成四个 32 字节事务(从 L2 缓存级别可见 - 它仍然是来自来自)在您描述的情况下,对于完全合并的访问,无论缓存或非缓存模式如何,四个 32 字节事务都不会变慢。在任何一种情况下,内存控制器(在给定的 GPU 上)都应该生成相同的事务来满足 warp 的请求。由于内存控制器由多个(最多 6 个)“分区”组成,每个分区都有 64 位宽的路径,因此最终将使用多个内存事务(可能跨越多个分区)来满足请求(4x32 字节或1x128 字节)。跨分区的具体事务数量和组织可能因 GPU 而异(这不是您的问题的一部分,但具有 DDR 泵送内存的 GPU 将为每个内存事务返回每个分区 16 字节,而使用 QDR 泵送内存时,每个内存事务将返回每个分区 32 字节)。这也不是 CUDA 5 特有的。您可能想回顾一下 NVIDIA 的产品之一webinars http://developer.nvidia.com/cuda/gpu-computing-webinars对于此材料,特别是“CUDA 优化:内存带宽有限内核”。即使你不想看video http://developer.download.nvidia.com/CUDA/training/Optimizing_Mem_limited_kernels.mp4,快速回顾一下slides http://developer.download.nvidia.com/CUDA/training/bandwidthlimitedkernels_webinar.pdf会提醒您所谓的“缓存”和“非缓存”访问(这是指 L1)之间的各种差异,并且还会为您提供尝试每种情况所需的编译器开关。

查看幻灯片的另一个原因是它会提醒您在什么情况下可能想要尝试“未缓存”模式。特别是,如果您的 warp 具有分散(未合并)的访问模式,则非缓存模式访问可能会带来改进,因为与 128 字节相比,从内存请求 32 字节数量以满足单个线程的请求时“浪费”更少数量。然而,针对您的最后一个问题,对其进行分析相当困难,因为您的代码可能是有序和无序访问模式的混合。由于非缓存模式是通过编译器开关打开的,因此幻灯片中给出的建议只是“尝试两种方式的代码”,看看哪种运行速度更快。根据我的经验,在非缓存模式下运行很少会带来性能改进。

编辑:抱歉,我的演示文稿链接和标题错误。修复了幻灯片/视频链接和网络研讨会标题。

本文内容由网友自发贡献,版权归原作者所有,本站不承担相应法律责任。如您发现有涉嫌抄袭侵权的内容,请联系:hwhale#tublm.com(使用前将#替换为@)

CUDA全局内存事务的成本 的相关文章

  • 在 Windows 上的 Qt Creator 中编译 Cuda 代码

    几天来我一直在尝试获取在 32 位 Windows 7 系统上运行的 Qt 项目文件 我希望 需要在其中包含 Cuda 代码 这种组合要么非常简单 以至于没有人愿意在网上放一个例子 要么非常困难 似乎没有人成功 不管怎样 我发现的唯一有用的
  • C 中带括号和不带括号的循环处理方式不同吗?

    我在调试器中单步执行一些 C CUDA 代码 如下所示 for uint i threadIdx x i lt 8379 i 256 sum d PartialHistograms blockIdx x i HISTOGRAM64 BIN
  • CUDA错误:在python中使用并行时初始化错误

    我的代码使用 CUDA 但运行速度仍然很慢 因此 我将其更改为使用 python 中的多处理 pool map 并行运行 但我有CUDA ERROR initialization error 这是函数 def step M self ite
  • Ubuntu 11.10/12.04 上的 CUDA“无兼容设备”错误

    一段时间以来 我一直在尝试在我的笔记本电脑上设置 Ubuntu 环境来进行 CUDA 编程 我目前双启动 Windows 8 和 Ubuntu 12 04 并想在 Ubuntu 上安装 CUDA 5 该笔记本电脑配有 GeForce GT
  • cuda中的count3非常慢

    我在 CUDA 中编写了一个小程序 用于计算 C 数组中有多少个 3 并打印它们 include
  • CUDA:如何在设备上填充动态大小的向量并将其内容返回到另一个设备函数?

    我想知道哪种技术可以填充设备上的动态大小数组 int row 在下面的代码中 然后返回其内容 以供另一个设备函数使用 为了将问题置于上下文中 下面的代码尝试使用在 GPU 上运行的高斯 勒让德求积来跨越勒让德多项式基组中的任意函数 incl
  • cudaMallocManaged() 返回“不支持的操作”

    在 CUDA 6 0 中尝试托管内存给了我operation not supported打电话时cudaMallocManaged include cuda runtime h include
  • 在 cuda 的 nvcc 编译器中使用 C++20

    我正在尝试使用std countr zero 函数从
  • CUDA素数生成

    当数据大小增加超过 260k 时 我的 CUDA 程序停止工作 它不打印任何内容 有人能告诉我为什么会发生这种情况吗 这是我的第一个 CUDA 程序 如果我想要更大的素数 如何在 CUDA 上使用大于 long long int 的数据类型
  • cuda cpu功能-gpu内核重叠

    我在尝试开发以练习 CUDA 的 CUDA 应用程序时遇到并发问题 我想通过使用 cudaMemecpyAsync 和 CUDA 内核的异步行为来共享 GPU 和 CPU 之间的工作 但我无法成功重叠 CPU 执行和 GPU 执行 它与主机
  • MPI+CUDA 与纯 MPI 相比有何优势?

    加速应用程序的常用方法是使用 MPI 或更高级别的库 例如在幕后使用 MPI 的 PETSc 并行化应用程序 然而 现在每个人似乎都对使用 CUDA 来并行化他们的应用程序或使用 MPI 和 CUDA 的混合来解决更雄心勃勃 更大的问题感兴
  • 如何优化这个 CUDA 内核

    我已经分析了我的模型 似乎该内核约占我总运行时间的 2 3 我一直在寻找优化它的建议 代码如下 global void calcFlux double concs double fluxes double dt int idx blockI
  • 如何在 Visual Studio 2010 中设置 CUDA 编译器标志?

    经过坚持不懈的得到error identifier atomicAdd is undefined 我找到了编译的解决方案 arch sm 20旗帜 但是如何在 VS 2010 中传递这个编译器标志呢 我已经尝试过如下Project gt P
  • Nvcc 的版本与 CUDA 不同

    我安装了 cuda 7 但是当我点击 nvcc version 时 它打印出 6 5 我想在 GTX 960 卡上安装 Theano 库 但它需要 nvcc 7 0 我尝试重新安装cuda 但它没有更新nvcc 当我运行 apt get i
  • 如何在 CUDA 中执行多个矩阵乘法?

    我有一个方阵数组int M 10 以便M i 定位第一个元素i th 矩阵 我想将所有矩阵相乘M i 通过另一个矩阵N 这样我就收到了方阵数组int P 10 作为输出 我看到有不同的可能性 分配不同元素的计算M i 到不同的线程 例如 我
  • 加速Cuda程序

    要更改哪一部分来加速此代码 代码到底在做什么 global void mat Matrix a Matrix b int tempData new int 2 tempData 0 threadIdx x tempData 1 blockI
  • 有没有一种有效的方法来优化我的序列化代码?

    这个问题缺乏细节 因此 我决定创建另一个问题而不是编辑这个问题 新问题在这里 我可以并行化我的代码吗 还是不值得 https stackoverflow com questions 17937438 can i parallelize my
  • cudaMemcpy() 与 cudaMemcpyFromSymbol()

    我试图找出原因cudaMemcpyFromSymbol 存在 似乎 symbol func 可以做的所有事情 nonSymbol cmd 也可以做 symbol func 似乎可以轻松移动数组或索引的一部分 但这也可以使用 nonSymbo
  • VS 程序在调试模式下崩溃,但在发布模式下不崩溃?

    我正在 VS 2012 中运行以下程序来尝试 Thrust 函数查找 include cuda runtime h include device launch parameters h include
  • 如何运行和理解CUDA Visual Profiler?

    我已经设置了 CUDA 5 0 并且我的 CUDA 项目运行良好 但我不知道如何使用 Visual Profiler 分析我的 CUDA 项目 如何运行它 我还需要安装更多吗 又该如何做呢 我的电脑使用Window 7 64位 CUDA 5

随机推荐