blockIdx 与区块执行顺序相关吗?

2024-01-09

之间有什么关系吗blockIdx以及线程块在 GPU 设备上执行的顺序?

我的动机是,我有一个内核,其中多个块将从全局内存中的同一位置读取,如果这些块能够同时运行,那就太好了(因为 L2 缓存命中很好)。在决定如何将这些块组织成网格时,可以肯定地说blockIdx.x=0更有可能同时运行blockIdx.x=1比与blockIdx.x=200?我应该尝试将连续索引分配给从全局内存中同一位置读取的块?

需要明确的是,我并不是在询问块间依赖关系(如这个问题 https://stackoverflow.com/questions/2392250/understanding-cuda-grid-dimensions-block-dimensions-and-threads-organization-s)并且从程序正确性的角度来看,线程块是完全独立的。我已经在使用共享内存来广播块内的数据,并且我无法使块更大。

编辑:再次,我很清楚

线程块需要独立执行:必须能够以任何顺序(并行或串行)执行它们。

并且这些块是完全独立的——它们可以以任何顺序运行并产生相同的输出。我只是问我将块排列到网格中的顺序是否会影响哪些块最终同时运行,因为这确实会通过二级缓存命中率影响性能。


我发现了一篇文章,其中一位 CS 研究人员使用微基准测试对 Fermi 设备上的块调度程序进行逆向工程:

http://cs.rochester.edu/~sree/fermi-tbs/fermi-tbs.html http://cs.rochester.edu/%7Esree/fermi-tbs/fermi-tbs.html

我修改了他的代码以在我的 GPU 设备(GTX 1080,带有 Pascal GP104 GPU)上运行,并随机化运行时间。

Methods

每个块仅包含 1 个线程,并使用足够的共享内存启动,每个 SM 只能驻留 2 个块。内核记录其开始时间(通过获取clock64()),然后运行一段随机时间(该任务是使用进位乘法算法生成随机数)。

GTX 1080 由 4 个图形处理集群 (GPC) 组成,每个集群有 5 个流式多处理器 (SM)。每个 GPC 都有自己的时钟,因此我使用链接中描述的相同方法来确定哪些 SM 属于哪些 GPC,然后减去固定偏移量以将所有时钟值转换为同一时区。

Results

对于一维块网格,我发现块确实是按连续顺序启动的:

我们有 40 个块立即开始(每个 SM 2 个块 * 20 个 SM),后续块在前一个块结束时开始。

对于二维网格,我发现了相同的线性顺序,blockIdx.x是快速维度并且blockIdx.y慢维度:

注意:我在标记这些图时犯了一个严重的拼写错误。 “threadIdx”的所有实例都应替换为“blockIdx”。

And for a 3-d block grid: Block start time for a 3-D block grid

结论

对于一维网格,这些结果与 Pai 博士在链接文章中报告的结果相匹配。然而,对于二维网格,我没有找到任何块执行顺序中空间填充曲线的证据,因此这可能在费米和帕斯卡之间的某个地方发生了变化。

当然,基准测试的常见警告也适用,并且不能保证这不是特定于特定处理器型号的。

Appendix

作为参考,下面的图显示了随机运行时间与固定运行时间的结果:

事实上,我们在随机运行时看到了这种趋势,这让我更有信心这是一个真实的结果,而不仅仅是基准测试任务的一个怪癖。

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

blockIdx 与区块执行顺序相关吗? 的相关文章

  • 使用 CMake 通过 NVCC 传递到 MSVC 的混淆标志

    我有一个 CMake 文件 用于在 Windows 上构建一些 CUDA NVCC MSVC 我正在尝试将 MSVC 警告级别设置为 W4 using add compile options lt
  • cuda简单应用程序适用于32位而不适用于64位

    我的简单 cuda helloworld 应用程序在 Windows 10 上使用 Visual Studio 2015 社区构建 32 位时运行良好 但是 如果我在 64 位中构建它 则不会执行 GPU 特斯拉K40c 工具包 CUDA
  • 加强托管线程和操作系统线程之间的关系(CUDA 用例)

    Problem 我正在尝试创建一个与 net 良好集成的 CUDA 应用程序 设计目标是拥有多个可以从托管代码调用的 CUDA 函数 数据还应该能够在函数调用之间保留在设备上 以便可以将其传递给多个 CUDA 函数 重要的是 每个单独的数据
  • CUDA全局内存事务的成本

    根据 CUDA 5 0 编程指南 如果我同时使用 L1 和 L2 缓存 在 Fermi 或 Kepler 上 则所有全局内存操作都使用 128 字节内存事务完成 但是 如果我仅使用 L2 则使用 32 字节内存事务 第 F 4 2 章 让我
  • C 中带括号和不带括号的循环处理方式不同吗?

    我在调试器中单步执行一些 C CUDA 代码 如下所示 for uint i threadIdx x i lt 8379 i 256 sum d PartialHistograms blockIdx x i HISTOGRAM64 BIN
  • 指定 NVCC 用于编译主机代码的编译器

    运行 nvcc 时 它始终使用 Visual C 编译器 cl exe 我怎样才能让它使用GCC编译器 设置CC环境变量到gcc没有修复它 我在可执行文件帮助输出中也找不到任何选项 在 Windows 上 NVCC 仅支持 Visual C
  • 用于类型比较的 Boost 静态断言

    以下问题给我编译器错误 我不知道如何正确编写它 struct FalseType enum value false struct TrueType enum value true template
  • Cuda Bayer/CFA 去马赛克示例

    我编写了一个 CUDA4 Bayer 去马赛克例程 但它比在 16 核 GTS250 上运行的单线程 CPU 代码慢 块大小是 16 16 图像暗淡是 16 的倍数 但更改此值并不会改善它 我做了什么明显愚蠢的事情吗 calling rou
  • 如何将CUDA时钟周期转换为毫秒?

    我想用一些代码来测量时间within我的内核需要 我已经关注了这个问题 https stackoverflow com questions 11209228 timing different sections in cuda kernel连
  • 同时使用 2 个 GPU 调用 cudaMalloc 时性能较差

    我有一个应用程序 可以在用户系统上的 GPU 之间分配处理负载 基本上 每个 GPU 都有一个 CPU 线程来启动一个GPU处理间隔当由主应用程序线程定期触发时 考虑以下图像 使用 NVIDIA 的 CUDA 分析器工具生成 作为示例GPU
  • 如何在 Visual Studio 2010 中设置 CUDA 编译器标志?

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

    Linux 上的 CUDA 发行版曾经有一个名为version txt例如 CUDA Version 10 2 89 这非常有用 但是 从 CUDA 11 1 开始 该文件不再存在 我如何在 Linux 上通过命令行确定并检查 path t
  • Yocto for Nvidia Jetson 由于 GCC 7 而失败 - 无法计算目标文件的后缀

    我正在尝试将 Yocto 与 meta tegra 一起使用 https github com madisongh meta tegra https github com madisongh meta tegra 为 Nvidia Jets
  • 在 __device/global__ CUDA 内核中动态分配内存

    根据CUDA 编程指南 http developer download nvidia com compute cuda 3 2 prod toolkit docs CUDA C Programming Guide pdf 第 122 页 可
  • 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中有模板化的数学函数吗? [复制]

    这个问题在这里已经有答案了 我一直在寻找 cuda 中的模板化数学函数 但似乎找不到 在普通的 C 中 如果我调用std sqrt它是模板化的 并且将根据参数是浮点数还是双精度数执行不同的版本 我想要这样的 CUDA 设备代码 我的内核将真
  • 如何在cmake中添加cuda源代码的定义

    我使用的是 Visual Studio 2013 Windows 10 CMake 3 5 1 一切都可以使用标准 C 正确编译 例如 CMakeLists txt project Test add definitions D WINDOW
  • cuda中内核的并行执行

    可以说我有三个全局数组 它们已使用 cudaMemcpy 复制到 GPU 中 但 c 中的这些全局数组尚未使用 cudaHostAlloc 分配 以便分配页面锁定的内存 而不是简单的全局分配 int a 100 b 100 c 100 cu
  • __device__ __constant__ 常量

    有什么区别吗 在 CUDA 程序中定义设备常量的最佳方法是什么 在 C 主机 设备程序中 如果我想将常量定义在设备常量内存中 我可以这样做 device constant float a 5 constant float a 5 问题 1
  • CUDA、NPP 滤波器

    CUDA NPP 库支持使用 nppiFilter 8u C1R 命令过滤图像 但不断出现错误 我可以毫无问题地启动并运行 boxFilterNPP 示例代码 eStatusNPP nppiFilterBox 8u C1R oDeviceS

随机推荐

  • Java 中 % 运算符和 IEEEremainder() 方法之间的区别(如果有的话)?

    在 Java 中 使用 运算符获取整数除法 x y 的余数与 Math IEEEremainder x y 方法之间是否存在功能或性能差异 除了 John B 已经指出的类型差异之外 语义上也存在显着差异 Math IEEEremainde
  • 当应用程序在后台运行时,如何对 Core Data iCloud 同步通知采取行动?

    我有一个 iOS 应用程序 它使用 Core Data 和 iCloud 同步 该应用程序运行完美 并且可以跨多个设备同步 作为我的实现的一部分 我的应用程序注册为 NSPersistentStoreCoordinatorStoresDid
  • 如何在 iOS 上读取、修改和写入 PDF (CGPDFDocument)?

    我正在尝试修改 iOS 上现有 PDF 的 PDF 元数据 标题 作者等 虽然很容易找到用于 PDF 解析和从头开始创建 PDF 的示例代码 但似乎没有一种简单的方法可以将现有 PDF 转储到新文件中 并稍微修改它 更具体地说 如何获取阅读
  • 最长最大重复子串

    子串的长度可以是 1 2 3 我试图解决的问题涉及找到出现次数最多的子字符串 所以它基本上分解为寻找具有最大频率的字符 然而 我发现我可以使用后缀树在 O n 中找到最长的重复子串 但是 后缀树返回子字符串 并优先考虑长度 我想找到出现次数
  • 如何在文件顶部添加新的文本行?

    我正在开发一个简单的程序 使Python脚本可执行 并且我正在添加解释器路径的部分 usr bin python 我尝试这样做 但它没有添加新行 而是替换了当前的行and删除下一行的一部分 我做错了什么 我将源代码上传到Ubuntu Pas
  • Swift - 压缩视频文件

    所以 目前我正在使用它来压缩视频 func compressVideo inputURL NSURL outputURL NSURL handler session AVAssetExportSession gt Void let urlA
  • 使用 min/max 时的任何差异都发生在序列而不是元素级别

    我的意思是是否存在以下情况
  • 如何从另一个项目导入.java并在android中导入包

    我们正在大学里开展一个项目 项目 A 该项目必须从多年前制定的另一个项目 项目 B 的课程中扩展出来 这是在 Eclipse 下工作 并为 android 开发应用程序 所以这实际上是我们正在讨论的活动 到目前为止我已经尝试过这个 将项目
  • 嵌套 CSS 样式?

    这样的事情可能吗 imgbox hover ui resizable se some style 或者概念上的等价物 基本上 只有当某个类的元素悬停在上面时 才会出现某个元素within该类应该改变一些风格 你可以这样做 imgbox ho
  • 我们应该在 Rails Factory 中使用 Faker 吗?

    I love Faker https github com stympy faker 我在我的seeds rb一直用真实的数据填充我的开发环境 我也刚刚开始使用工厂女工 https github com thoughtbot factory
  • 在 Android 中使用 color 和 color.darker?

    好的 我的应用程序中有一个整数变量 它是颜色的值 由颜色选择器根据我的喜好设置 现在 我需要使用该颜色和任何颜色的深色版本 现在我知道在标准 Java 中有一个 Color darker 方法 但在 Android 中似乎没有等效的方法 有
  • Linq 中的 SQL LIKE

    在添加这个问题之前 我确实在 stackoverflow 上搜索过类似的问题 但我找不到 互联网上的大多数问题都使用 LIKE 和字符串 例如 LIKE ABC 但我需要与不同表的现有列进行比较 我需要为 select 语句编写一个 lin
  • spring boot org.springframework.beans.factory.BeanCreationException:无法自动装配字段:

    我从 spring boot 开始 遇到一些配置问题 我无法自动装配某些服务 我收到 BeanCreationException 我的应用程序类 SpringBootApplication EnableAutoConfiguration C
  • 活动没有选项菜单

    我试图模拟点击菜单项存在于Toolbar using 机器人电动 使用以下代码 ShadowActivity shadowActivity Shadows shadowOf activity shadowActivity clickMenu
  • 如何查看导致 LINQ to SQL 中的 SubmitChanges 错误的 sql?

    我有一些 LINQ to SQL 有时会抛出 无法在具有唯一索引的对象 dbo Table 中插入重复的键行 IX Indexname 该语句已终止 有什么方法可以打开日志记录或至少调试数据上下文以查看引发错误时正在执行的 sql 吗 Up
  • 如何从文档中删除所有肖像图片

    我正在对文档图像进行 OCR 处理 我想检测所有图片并从文档图像中删除 我想保留文档图像中的表格 一旦我检测到图片 我就会将其删除 然后进行 OCR 我试图找到轮廓试图检测所有更大的区域 不幸的是它也检测到表格 还如何删除在文档图像中保留其
  • 通过 Python 代码使用 SignalR 服务器

    将 Python 与 SignalR 集成有哪些选项 Python 代码是大型第三方产品的一部分 与语言偏好无关 SignalR 服务器提供对现有 NET 产品的订阅 我们希望通过 Python 重用 NET SignalR 服务器 pyt
  • Java中如何随机选择图片?

    我正在制作一个将使用以下代码的程序 JLabel MyImage new JLabel new ImageIcon image1 png 但是 我想让随机图片出现 比如 image2 image3 image4 我该如何做到这一点 不使用L
  • @RequestMapping 在 Spring Boot 内部如何工作?

    RestController RequestMapping employee public class Employee RequestMapping save public void saveEmployee saving employe
  • blockIdx 与区块执行顺序相关吗?

    之间有什么关系吗blockIdx以及线程块在 GPU 设备上执行的顺序 我的动机是 我有一个内核 其中多个块将从全局内存中的同一位置读取 如果这些块能够同时运行 那就太好了 因为 L2 缓存命中很好 在决定如何将这些块组织成网格时 可以肯定