合理化我的简单 OpenCL 内核中有关全局内存的情况

2024-03-04

const char programSource[] =
        "__kernel void vecAdd(__global int *a, __global int *b, __global int *c)"
        "{"
        "    int gid = get_global_id(0);"
        "for(int i=0; i<10; i++){"
        "    a[gid] = b[gid] + c[gid];}"
        "}";

上面的内核是每个循环执行十次的向量加法。我已经使用编程指南和堆栈溢出来弄清楚全局内存是如何工作的,但我仍然无法通过查看我的代码来弄清楚我是否以良好的方式访问全局内存。我以连续的方式访问它,并且以一致的方式猜测。该卡是否为数组 a、b 和 c 加载 128kb 全局内存块?然后,它是否会在每处理 32 个 gid 索引时为每个数组加载一次 128kb 块? (4*32=128) 看来我没有浪费任何全局内存带宽,对吗?

顺便说一句,计算分析器显示 gld 和 gst 效率为 1.00003,这看起来很奇怪,我认为如果我所有的存储和负载都合并起来,它只会是 1.0。 1.0以上怎么样?


是的,您的内存访问模式几乎是最佳的。每个 halfwarp 访问 16 个连续的 32 位字。此外,访问是 64 字节对齐的,因为缓冲区本身是对齐的,并且每个 halfwarp 的起始索引是 16 的倍数。因此每个 halfwarp 将生成一个 64 字节事务。因此,您不应该通过未合并的访问来浪费内存带宽。

既然您在上一个问题中要求提供示例,那么让我们修改此代码以用于其他(不太理想的访问模式(因为循环实际上没有做任何事情,我将忽略它):

kernel void vecAdd(global int* a, global int* b, global int* c)
{
   int gid = get_global_id(0);
   a[gid+1] = b[gid * 2] + c[gid * 32];
}

首先让我们看看它在计算 1.3 (GT200) 硬件上的工作原理

对于对 a 的写入,这将生成稍微不太理想的模式(遵循由其 id 范围和相应的访问模式标识的 halfwarp):

   gid  | addr. offset | accesses     | reasoning
  0- 15 |     4- 67    | 1x128B       | in aligned 128byte block
 16- 31 |    68-131    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
 32- 47 |   132-195    | 1x128B       | in aligned 128byte block
 48- 63 |   196-256    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access

所以基本上我们浪费了大约一半的带宽(奇数 halfwarp 的访问宽度加倍并没有多大帮助,因为它会生成更多的访问,可以说这并不比浪费更多字节更快)。

对于从 b 读取,线程仅访问数组的偶数元素,因此对于每个 halfwarp,所有访问都位于 128 字节对齐的块中(第一个元素位于 128B 边界,因为对于该元素,gid 是 16 的倍数=>索引是32的倍数,对于4字节元素,这意味着地址偏移量是128B的倍数)。访问模式延伸到整个 128B 块,因此这将为每个 halfwarp 执行 128B 传输,再次减少一半的带宽。

从 c 中读取会产生最坏的情况之一,其中每个线程都在自己的 128B 块中索引,因此每个线程都需要自己的传输,一方面有点序列化场景(尽管不像正常情况那么糟糕,因为硬件应该能够重叠传输)。更糟糕的是,这将为每个线程传输 32B 块,浪费 7/8 的带宽(我们访问 4B/线程,32B/4B=8,因此只利用了 1/8 的带宽)。由于这是朴素矩阵转置的访问模式,因此强烈建议使用本地内存进行这些访问模式(根据经验)。

计算 1.0 (G80)

这里唯一能够创建良好访问的模式是原始模式,示例中的所有模式都将创建完全未合并的访问,浪费 7/8 的带宽(32B 传输/线程,见上文)。对于 G80 硬件,halfwarp 中第 n 个线程不访问第 n 个元素的每次访问都会创建此类未合并的访问

计算 2.0(费米)

在这里,每次对内存的访问都会创建 128B 事务(收集所有数据所需的数量,因此在最坏的情况下为 16x128B),但是这些事务被缓存,使得数据将传输到何处不太明显。目前我们假设缓存足够大,可以容纳所有数据并且不存在冲突,因此每个 128B 缓存行最多会传输一次。让我们进一步假设 halfwarp 是串行执行的,因此我们有确定性的缓存占用。

对 b 的访问仍将始终传输 128B 块(相应内存区域中没有其他线程索引)。对 c 的访问将为每个线程生成 128B 传输(可能是最差的访问模式)。

对于 a 的访问如下(暂时将它们视为读取):

   gid  | offset  | accesses | reasoning
  0- 15 |   4- 67 |  1x128B  | bringing 128B block to cache
 16- 31 |  68-131 |  1x128B  | offsets 68-127 already in cache, bring 128B for 128-131 to cache
 32- 47 | 132-195 |    -     | block already in cache from  last halfwarp
 48- 63 | 196-259 |  1x128B  | offsets 196-255 already in cache, bringing in 256-383

所以对于大型数组来说,理论上a的访问几乎不会浪费带宽。 对于这个例子来说,现实当然没有那么好,因为对 c 的访问会很好地破坏缓存

对于探查器,我假设超过 1.0 的效率只是浮点不准确的结果。

希望有帮助

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

合理化我的简单 OpenCL 内核中有关全局内存的情况 的相关文章

  • 直接在主机上访问设备向量元素的最快方法

    我请您参考以下页面http code google com p thrust wiki QuickStartGuide Vectors http code google com p thrust wiki QuickStartGuide V
  • 支持 Nvidia CUDA 工具包 9.2

    Tensorflow gpu 绑定到 Nvidia CUDA Toolkit 的特定版本的原因是什么 当前版本似乎专门寻找 9 0 并且不适用于任何更高版本 例如 我安装了最新的 Toolkit 9 2 并将其添加到路径中 但 Tensor
  • 如何在没有 nvcc 的情况下在编译时获取 CUDA 工具包版本?

    我在 cpp 文件中对 cuSPARSE 库进行了一些调用 这些调用在旧工具包中不可用 为了支持使用旧工具包的系统 我想使用编译器指令编译不同的代码部分 特别是 我想使用旧工具包的 CSR 格式矩阵和新工具包的 BSR 格式矩阵来求解稀疏三
  • C# 应用程序可以使用多少 RAM? [关闭]

    就目前情况而言 这个问题不太适合我们的问答形式 我们希望答案得到事实 参考资料或专业知识的支持 但这个问题可能会引发辩论 争论 民意调查或扩展讨论 如果您觉得这个问题可以改进并可能重新开放 访问帮助中心 help reopen questi
  • 如何在 opencv 3.0 Beta 中从文件读取 UMat?

    我想用UMat所以我的代码可以使用 OpenCL OpenCV 3 0 0 Beta 在 GPU 和 CPU 上运行 但我找不到将图像文件读入的方法UMat或转换一个Mat to UMat 如何将图像读入UMat 样品用于Mat to UM
  • 在“delete this;”语句期间发生了什么?

    请考虑以下代码 class foo public foo foo void done delete this private int x 以下两个选项中发生了什么 并且有效吗 选项1 void main foo a new foo a gt
  • 找出Linux上一个进程使用了​​多少内存页

    我需要找出进程分配了多少内存页 每个页面是 4096 进程内存使用情况我在查找正确值时遇到一些问题 当我查看 gome system monitor 时 内存映射下有几个值可供选择 Thanks 这样做的目的是将内存使用量除以页数并验证页大
  • CUDA-Kernel 应该根据块大小动态崩溃

    我想做稀疏矩阵 密集向量乘法 假设用于压缩矩阵中条目的唯一存储格式是压缩行存储 CRS 我的内核如下所示 global void krnlSpMVmul1 float data mat int num nonzeroes unsigned
  • 布尔实现的atomicCAS

    我想弄清楚是否存在错误答案 https stackoverflow com a 57444538 11248508 现已删除 关于Cuda like的实现atomicCAS for bool是 答案中的代码 重新格式化 static inl
  • HUGE_VALF 和 INFINITY 常量之间的区别

    在 OpenCL 中 有两个代表无穷大的浮点数学常数 其中之一很简单INFINITY 另一个 HUGE VALF 求值为 无穷大 这两者有什么区别 求值至 无穷大是什么意思 HUGE VALF是一个旧名称 允许不支持无穷大的浮点系统 例如
  • 查找可以为 C# 中的数组分配多少内存

    我正在做一些需要初始化大数组的计算 数组的最大大小决定了我能解决的问题的最大大小 有没有一种方法可以以编程方式确定有多少内存可供使用 例如可能的最大字节数组 Thanks 嗯 依赖单个大数组会带来一系列相关问题 内存碎片 连续块 最大对象大
  • JSON 解析器从大型 JSON 文件中逐条读取

    我有一个巨大的 JSON 文件 1GB 它基本上是以下格式的对象数组 x y p q x1 y1 p1 q1 我想解析这个文件 这样所有的数据都不会加载到内存中 基本上我想获取例如 数组中的前 1000 个对象进行内存处理 然后将接下来的
  • 结构体的内存大小不同?

    为什么第一种情况不是12 测试环境 最新版本的 gcc 和 clang 64 位 Linux struct desc int parts int nr sizeof desc Output 16 struct desc int parts
  • goto 指令对 CUDA 代码中扭曲内发散的影响

    对于CUDA中简单的warp内线程发散 我所知道的是SM选择一个重新收敛点 PC地址 并在两个 多个路径中执行指令 同时禁用未采用该路径的线程的执行效果 例如 在下面的代码中 if threadIdx x lt 16 A do someth
  • Valgrind 输出中的错误摘要?

    我看过一些关于 valgrind 的帖子 但没有一篇帖子帮助我理解 valgrind 输出的解释 我用 valgrind 运行了两个程序 都有内存泄漏 测试 1 的示例输出 20422 LEAK SUMMARY 20422 definite
  • dlib 不使用 CUDA

    我使用 pip 安装了 dlib 我的显卡支持 CUDA 但是在运行 dlib 时 它不使用 GPU 我在 ubuntu 18 04 上工作 Python 3 6 5 default Apr 1 2018 05 46 30 GCC 7 3
  • 如果 free() 知道我的数组的长度,为什么我不能在自己的代码中请求它?

    我知道将动态分配的数组的长度传递给操作它们的函数是一个常见的约定 void initializeAndFree int anArray size t length int main size t arrayLength 0 scanf d
  • 如何获取可用系统内存的大小?

    C NET 中是否可以获取系统可用内存的大小 如果是的话怎么办 Use Microsoft VisualBasic Devices ComputerInfo TotalPhysicalMemory http msdn microsoft c
  • 为什么在 CUDA 中启动 32 倍数的线程?

    我参加了 CUDA 并行编程课程 并且看到了许多 CUDA 线程配置的示例 其中通常将所需的线程数四舍五入到最接近的 32 倍数 我知道线程被分组为 warp 并且如果您启动 1000 个线程 GPU 无论如何都会将其四舍五入到 1024
  • 使用自定义堆的类似 malloc 的函数

    如果我希望使用自定义预分配堆构造类似 malloc 的功能 那么 C 中最好的方法是什么 我的具体问题是 我有一个可映射 类似内存 的设备 已将其放入我的地址空间中 但我需要获得一种更灵活的方式来使用该内存来存储将随着时间的推移分配和释放的

随机推荐

  • C语言中的二维数组如何变成一维数组?

    如果有人可以向我解释以下行为 我将不胜感激 假设我声明一个静态二维数组 float buffer NX NY 现在 如果我想填充这个数组 我注意到可以这样做 initarray buffer NX NY define INITDATAVAL
  • 没有 Redux 的情况下组合Reducer

    我有一个没有 redux 的应用程序 我使用钩子和钩子 useReducer context 处理全局状态 我有 1 个 useReducer 它就像一个 Redux 商店 但要做到这一点我只能发送 1 个减速器 在该减速器中 我拥有所有状
  • 根据数组中的另一个 id 仅对多数组中的第一项进行排序 (PHP)

    我不知道该怎么做 请参阅下面我的数组 我在 while 循环中运行这个数组 需要先找到 attach id 对于每个 topic id 并可以使用 topic id在循环中设置的 正确的输出将是 第一个循环 attach id gt 179
  • 神经网络在一个纪元后趋于平坦

    我正在使用 keras 创建一个卷积神经网络 尝试将图像分类为两个不同的类 并且出于某种原因 在第一个纪元之后 准确性永远不会改变 使用 Keras 后to categorical 我的标签看起来像 0 1 1 0 1 0 0 1 我的模型
  • Amazon SES SMTP Python 用法

    我试图诊断为什么通过 Amazon SES 发送电子邮件无法通过 python 工作 以下示例演示了该问题 其中user and pass设置为适当的凭据 gt gt gt import smtplib gt gt gt s smtplib
  • System.Timers.Timer 与 System.Threading.Timer

    我最近一直在检查一些可能的计时器 并且System Threading Timer https learn microsoft com en us dotnet api system threading timer and System T
  • 即使有标记,pytest-django 也不允许数据库访问

    我很难找出我的设置出了什么问题 我正在尝试测试登录视图 无论我尝试什么 我都会得到 Database access not allowed use the django db mark or the db or transactional
  • 统一处理非托管 API 中的错误代码

    我正在围绕一个相当大的非托管 API 编写一个包装器 几乎每个导入的方法在失败时都会返回一个常见的错误代码 现在 我正在这样做 ErrorCode result Api Method if result ErrorCode SUCCESS
  • Playframework 与 CSRF:“会话中未找到 CSRF 令牌”?

    我正在使用 Playframework 及其内置 CSRF 过滤器和 Security Authenticator 系统制作一个简单的身份验证系统 但我遇到了一个问题 当用户填写登录名 密码并提交输入时 出现以下错误 在会话中找不到 CSR
  • 如何更新 SQL 中游标获取的列

    在进一步讨论之前 是的 我知道与基于集合的操作相比 游标的性能很差 在这种特殊情况下 我在包含 100 条左右记录的临时表上运行游标 并且该临时表始终相当小 因此性能不如灵活性那么重要 我的困难是我无法找到如何更新游标获取的列的示例 以前
  • 在 Git 中运行预提交挂钩。有没有办法验证脚本是否正在运行?

    我想运行 Git 按照博客的建议 我使用了 git init初始化存储库 然后 git在 hooks 目录中存在钩子的位置创建文件夹 然后按照我重命名的脚本的建议pre commit sample as pre commit它不起作用 所以
  • 如何在Qt中保存对话框的状态?

    假设对话框中有复选框 选项等控件 如何在 Qt 中保存对话框的状态 我应该使用 QSettings 还是其他东西 Thanks 我遇到了同样的问题 谷歌搜索并没有太大帮助 所以最后我写了自己的解决方案 我创建了一组函数 用于在创建和销毁时读
  • Keras:为什么损失函数必须为每个批次项返回一个标量,而不仅仅是一个标量?

    我正在 Keras 中编写一个自定义损失函数 但遇到了以下问题 为什么 Keras 损失函数必须为每个批次项返回一个标量 而不是仅返回一个标量 我关心的是整批的累计损失 而不是每件商品的损失 不是吗 我想我已经明白了 fit 有争论samp
  • 更好地理解 C# 泛型

    我查看了一些使用 C 泛型的示例代码 为什么以及何时应该使用它们 所有的例子都很复杂 我需要一个简单 清晰的示例来帮助我开始使用 C 泛型 一个非常简单的例子是通用的List
  • 使用 PHP 5.5 安装 xdebug

    我读了很多答案 但不明白为什么 xdebug 不起作用 php ini xdebug zend extension usr lib php5 20090626 xdebug so php v PHP 5 5 6 1 debphp org p
  • Docker 容器未开始给出“OCI 运行时创建失败”

    我已经安装了Docker版本 https docs docker com release notes 17 12 0 ce 构建 c97c6d6当我尝试启动任何容器时 出现以下错误 docker 来自守护进程的错误响应 OCI运行时创建失败
  • 在像素着色器中计算世界空间坐标

    我有一个像素着色器 我想根据我的世界空间坐标计算每个像素的位置 我该怎么做 我需要什么 我有一个ps input具有 float4 位置的结构 SV POSITION 我认为这很重要 但存储在里面的值似乎有点有趣 我似乎无法弄清楚它们有什么
  • Liferay 连接和压缩 javascript

    我试图弄清楚如何实现所有 js 文件的压缩和缩小 目前我将它们放在一个钩子插件中 html js mycustomjs folder 我知道liferay有自己的机制来压缩javascripts 在barebone jsp或 everyth
  • 在 MacOSX 上,当指定无效的身份验证凭据时,QNetworkAccessManager 会进入无限循环

    在我的跨平台应用程序中 我使用 QNetworkAccessManager 将 HTTP 请求发送到需要身份验证的 HTTP 服务 我最近升级到 QT5 令我完全惊讶的是 在 MacOSX 上 我的应用程序会在某些情况下尽快向我的服务发送大
  • 合理化我的简单 OpenCL 内核中有关全局内存的情况

    const char programSource kernel void vecAdd global int a global int b global int c int gid get global id 0 for int i 0 i