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全局内存事务的成本 的相关文章

  • 无法在 CUDA 中执行设备内核

    我正在尝试在全局内核中调用设备内核 我的全局内核是矩阵乘法 我的设备内核正在查找乘积矩阵每列中的最大值和索引 以下是代码 device void MaxFunction float Pd float max int x threadIdx
  • CUDA 常量内存是否应该被均匀地访问?

    我的 CUDA 应用程序的恒定内存小于 8KB 既然它都会被缓存 我是否需要担心每个线程访问相同的地址以进行优化 如果是 如何确保所有线程同时访问同一地址 既然它都会被缓存 我是否需要担心每个线程访问相同的地址以进行优化 是的 这缓存本身每
  • Cuda 6.5 找不到 - libGLU。 (在 ubuntu 14.04 64 位上)

    我已经在我的ubuntu上安装了cuda 6 5 我的显卡是 GTX titan 当我想要制作 cuda 样本之一时 模拟 粒子 我收到这条消息 gt gt gt WARNING libGLU so not found refer to C
  • 尝试构建我的 CUDA 程序时出现错误 MSB4062

    当我尝试构建我的第一个 GPU 程序时 出现以下错误 有什么建议可能会出什么问题吗 错误 1 错误 MSB4062 Nvda Build CudaTasks SanitizePaths 任务 无法从程序集 C Program 加载 文件 M
  • 如何在cmake中添加cuda源代码的定义

    我使用的是 Visual Studio 2013 Windows 10 CMake 3 5 1 一切都可以使用标准 C 正确编译 例如 CMakeLists txt project Test add definitions D WINDOW
  • 在 cudaFree() 之前需要 cudaDeviceSynchronize() 吗?

    CUDA 版本 10 1 帕斯卡 GPU 所有命令都发送到默认流 void ptr cudaMalloc ptr launch kernel lt lt lt gt gt gt ptr cudaDeviceSynchronize Is th
  • cuda中内核的并行执行

    可以说我有三个全局数组 它们已使用 cudaMemcpy 复制到 GPU 中 但 c 中的这些全局数组尚未使用 cudaHostAlloc 分配 以便分配页面锁定的内存 而不是简单的全局分配 int a 100 b 100 c 100 cu
  • 从 CUDA 设备写入输出文件

    我是 CUDA 编程的新手 正在将 C 代码重写为并行 CUDA 新代码 有没有一种方法可以直接从设备写入输出数据文件 而无需将数组从设备复制到主机 我假设如果cuPrintf存在 一定有地方可以写一个cuFprintf 抱歉 如果答案已经
  • CUDA - 将 CPU 变量传输到 GPU __constant__ 变量

    与 CUDA 的任何事情一样 最基本的事情有时也是最难的 所以 我只想将变量从 CPU 复制到 GPUconstant变量 我很难过 这就是我所拥有的 constant int contadorlinhasx d int main int
  • VS 程序在调试模式下崩溃,但在发布模式下不崩溃?

    我正在 VS 2012 中运行以下程序来尝试 Thrust 函数查找 include cuda runtime h include device launch parameters h include
  • CUDA 中指令重放的其他原因

    这是我从 nvprof CUDA 5 5 获得的输出 Invocations Metric Name Metric Description Min Max Avg Device Tesla K40c 0 Kernel MyKernel do
  • 如何运行和理解CUDA Visual Profiler?

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

    我想使用我的两张显卡通过 CUDA Thrust 进行计算 我有两张显卡 在单卡上运行对于两张卡都适用 即使我在 std vector 中存储两个 device vector 也是如此 如果我同时使用两张卡 循环中的第一个周期将起作用并且不
  • CUDA 的嵌套循环

    我想将我的 C 代码移植到 CUDA 主要计算部分包含3个for嵌套循环 for int i 0 i lt Nx i for int j 0 j
  • 无法在 CUDA 中找到 1 到 100 数字的简单和?

    我正在研究使用 CUDA 的图像处理算法 在我的算法中 我想使用 CUDA 内核找到图像所有像素的总和 所以我在cuda中制作了内核方法 来测量16位灰度图像的所有像素的总和 但我得到了错误的答案 所以我在cuda中编写了一个简单的程序来查
  • 为什么 cudaGLSetGLDevice 失败,即使它是在 main 函数的第一行中调用的

    我想使用 OpenGL 和 CUDA 之间的互操作性 我知道 正如一些教程所说 第一步是选择设备 但是 当我在主函数的第一行中调用 cudaGLSetGLDevice 0 时 程序退出并显示信息 cudaSafeCall 运行时 API 错
  • 在 Cuda 中简单添加两个 int,结果始终相同

    我开始了学习Cuda的旅程 我正在玩一些 hello world 类型的 cuda 代码 但它不起作用 我不知道为什么 代码非常简单 取两个整数并将它们添加到 GPU 上并返回结果 但无论我将数字更改为什么 我都会得到相同的结果 如果数学那
  • 对 CUDA 操作进行计时

    我需要计算 CUDA 内核执行的时间 最佳实践指南说我们可以使用事件或标准计时函数 例如clock 在Windows中 我的问题是使用这两个函数给出了完全不同的结果 事实上 与实践中的实际速度相比 事件给出的结果似乎是巨大的 我实际上需要这
  • CUDA-Kernel 应该根据块大小动态崩溃

    我想做稀疏矩阵 密集向量乘法 假设用于压缩矩阵中条目的唯一存储格式是压缩行存储 CRS 我的内核如下所示 global void krnlSpMVmul1 float data mat int num nonzeroes unsigned
  • 将数据从 GPU 复制到 CPU - CUDA

    我在将数据从 GPU 复制到 CPU 时遇到问题 一开始我在 GPU 空间中创建变量 device float gpu array 在此 GPU 函数中 我想将数据从 od fS gi 值 0 43 复制到 gpu array global

随机推荐