- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
我的问题是关于 warp 和调度的。我在这里使用 NVIDIA 费米术语。我的观察如下,它们正确吗?
一个。同一 warp 中的线程执行相同的指令。每个 warp 包含 32 个线程。
根据费米白皮书:“费米的双经线调度器选择两个经线,并发出一个从每个 warp 到一组 16 个内核、16 个加载/存储单元或 4 个 SFU 的指令。 "
从这里开始,我认为一个 warp(32 个线程)被调度了两次,因为 32 个内核中的 16 个被组合在一起。每个调度器在一个周期内将一半的 warp 发送到 16 个内核,总而言之,两个调度器在一个周期内将两个 warp-halves 发送到两个 16 核调度组中。换句话说,在这个费米架构中,一个 warp 需要被调度两次,一半一半。如果一个 warp 只包含 SFU 操作,那么这个 warp 需要发出 8 次(32/4),因为一个 SM 中只有 4 个 SFPU。
B.当启动大量线程(比如一维数组,320 个线程)时,连续的线程将自动分组为 10 个线程束,每个线程束有 32 个线程。因此,如果所有线程都在做相同的工作,它们将执行完全相同的指令。那么在这种情况下,所有 warp 总是携带相同的指令。
问题:Q1。哪一部分处理线程分组(成束)?软件还是硬件?如果是硬件,它是 warp 调度程序吗?以及硬件 warp 调度器是如何实现和工作的?
Q2。如果我有 64 个线程,线程 0-15 和 32-47 正在执行相同的指令,而 16-31 和 48-63 执行另一条指令,调度程序是否足够聪明,可以将不连续的线程(具有相同的指令)分组到同一个 warp 中(即,将线程 0-15 和 32-47 分组到同一个 warp,并将线程 16-31 和 48-63 分组到另一个 warp)?
Q3。 warp 大小 (32) 大于调度组大小(16 核)有什么意义?(这是一个硬件问题)因为在这种情况下(Fermi),无论如何都会将 warp 调度两次(在两个周期内)。如果一个 warp 是 16 宽,那么会简单地调度两个 warp(也是两个周期),这似乎和前面的情况一样。我想知道这个组织是否出于性能考虑。
我现在可以想象的是:可以保证同一个 warp 中的线程同步,这有时很有用,或者其他资源(如寄存器和内存)以 warp 大小为基础进行组织。我不确定这是否正确。
最佳答案
纠正一些误解:
A. ...From here, I think a warp(32 threads) is scheduled twice since 16 cores out of 32 are grouped together.
当向一组 16 个核心发出 warp 指令时,整个 warp 都会执行该指令,因为核心被计时两次(Fermi 的“hotclock”),因此每个核心实际上在一个周期内执行两个线程的计算(= 2 个热点)。调度 warp 指令时,整个 warp 都会得到服务。它不需要被安排两次。
B. ...Therefore, if all threads are doing the same work, they will execute exactly the same instruction. Then all warps are always carrying the same instruction in this case.
的确,一个 block 中的所有线程(以及所有线程束)都从相同的指令流执行,但它们不一定执行相同的指令。当然,warp 中的所有线程在任何给定时间都在执行相同的指令。但是 warp 彼此独立执行,因此 block 内的不同 warp 可能在任何给定时间执行来自流的不同指令。 Fermi whitepaper 第 10 页的图表使这一点显而易见。
Q1: Which part handles the threads grouping (into warps)? software or hardware?
这是由硬件完成的,如 hardware implementation 中所述编程指南的部分:“将 block 划分为 warp 的方式始终相同;每个 warp 包含连续的、增加的线程 ID 的线程,第一个 warp 包含线程 0。Thread Hierarchy 描述了线程 ID 与线程索引的关系 block 。”
and how the hardware warp scheduler is implemented and work?
我认为这在任何地方都没有正式记录。 Greg Smith 对此提供了各种解释,您可能希望搜索“user:124092 scheduler”或类似的搜索,以阅读他的一些评论。
Q2. If I have 64 threads, threads 0-15 and 32-47 are executing the same instruction while 16-31 and 48-63 executes another instruction, is the scheduler smart enough to group nonconsecutive threads( with the same instruction) into the same warp (i.e., to group threads 0-15 and 32-47 into the same warp, and to group threads 16-31 and 48-63 into another warp)?
这个问题是基于前面概述的误解。将线程分组到 warp 中不是动态的;它在 threadblock 启动时是固定的,并且遵循上面对 Q1 的回答中描述的方法。此外,线程 0-15 永远不会与 16-31 以外的任何线程一起调度,因为 0-31 包含一个 warp,它在 Fermi 上对于调度目的是不可分割的。
Q3. What's the point to have a warp size(32) larger than the scheduling group size(16 cores)?
同样,我认为这个问题是基于以前的误解。用于为 warp 提供资源的硬件单元在某些功能级别可能存在 16 个单元(或其他数量),但从操作级别来看,warp 被安排为 32 个线程,每条指令为整个扭曲安排,并在一定数量的费米热时钟内一起执行。
关于cuda - 硬件 warp 调度程序如何形成和处理 warp?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/21535471/
这是我关于 Stack Overflow 的第一个问题,这是一个很长的问题。 tl;dr 版本是:我如何使用 thrust::device_vector如果我希望它存储不同类型的对象 DerivedC
我已使用 cudaMalloc 在设备上分配内存并将其传递给内核函数。是否可以在内核完成执行之前从主机访问该内存? 最佳答案 我能想到的在内核仍在执行时启动 memcpy 的唯一方法是在与内核不同的流
是否可以在同一节点上没有支持 CUDA 的设备的情况下编译 CUDA 程序,仅使用 NVIDIA CUDA Toolkit...? 最佳答案 你的问题的答案是肯定的。 nvcc编译器驱动程序与设备的物
我不知道 cuda 不支持引用参数。我的程序中有这两个函数: __global__ void ExtractDisparityKernel ( ExtractDisparity& es)
我正在使用 CUDA 5.0。我注意到编译器将允许我在内核中使用主机声明的 int 常量。但是,它拒绝编译任何使用主机声明的 float 常量的内核。有谁知道这种看似差异的原因? 例如,下面的代码可以
自从 CUDA 9 发布以来,显然可以将不同的线程和 block 分组到同一组中,以便您可以一起管理它们。这对我来说非常有用,因为我需要启动一个包含多个 block 的内核并等待所有 block 都同
我需要在 CUDA 中执行三线性插值。这是问题定义: 给定三个点向量:x[nx]、y[ny]、z[nz] 和一个函数值矩阵func[nx][ny][nz],我想在 x、y 范围之间的一些随机点处找到函
我认为由于 CUDA 可以执行 64 位 128 位加载/存储,因此它可能具有一些用于加/减/等的内在函数。像 float3 这样的向量类型,在像 SSE 这样更少的指令中。 CUDA 有这样的功能吗
我有一个问题,每个线程 block (一维)必须对共享内存内的一个数组进行扫描,并执行几个其他任务。 (该数组最多有 1024 个元素。) 有没有支持这种操作的好库? 我检查了 Thrust 和 Cu
我对线程的形成和执行方式有很多疑惑。 首先,文档将 GPU 线程描述为轻量级线程。假设我希望将两个 100*100 矩阵相乘。如果每个元素都由不同的线程计算,则这将需要 100*100 个线程。但是,
我正在尝试自己解决这个问题,但我不能。 所以我想听听你的建议。 我正在编写这样的内核代码。 VGA 是 GTX 580。 xxxx >> (... threadNum ...) (note. Shar
查看 CUDA Thrust 代码中的内核启动,似乎它们总是使用默认流。我可以让 Thrust 使用我选择的流吗?我在 API 中遗漏了什么吗? 最佳答案 我想在 Thrust 1.8 发布后更新 t
我想知道 CUDA 应用程序的扭曲调度顺序是否是确定性的。 具体来说,我想知道在同一设备上使用相同输入数据多次运行同一内核时,warp 执行的顺序是否会保持不变。如果没有,是否有任何东西可以强制对扭曲
一个 GPU 中可以有多少个 CUDA 网格? 两个网格可以同时存在于 GPU 中吗?还是一台 GPU 设备只有一个网格? Kernel1>(dst1, param1); Kernel1>(dst2,
如果我编译一个计算能力较低的 CUDA 程序,例如 1.3(nvcc 标志 sm_13),并在具有 Compute Capability 2.1 的设备上运行它,它是否会利用 Compute 2.1
固定内存应该可以提高从主机到设备的传输速率(api 引用)。但是我发现我不需要为内核调用 cuMemcpyHtoD 来访问这些值,也不需要为主机调用 cuMemcpyDtoA 来读取值。我不认为这会奏
我希望对 CUDA C 中负载平衡的最佳实践有一些一般性的建议和说明,特别是: 如果经纱中的 1 个线程比其他 31 个线程花费的时间长,它会阻止其他 31 个线程完成吗? 如果是这样,多余的处理能力
CUDA 中是否有像 opencl 一样的内置交叉和点积,所以 cuda 内核可以使用它? 到目前为止,我在规范中找不到任何内容。 最佳答案 您可以在 SDK 的 cutil_math.h 中找到这些
有一些与我要问的问题类似的问题,但我觉得它们都没有触及我真正要寻找的核心。我现在拥有的是一种 CUDA 方法,它需要将两个数组定义到共享内存中。现在,数组的大小由在执行开始后读入程序的变量给出。因此,
经线是 32 根线。 32 个线程是否在多处理器中并行执行? 如果 32 个线程没有并行执行,则扭曲中没有竞争条件。 在经历了一些例子后,我有了这个疑问。 最佳答案 在 CUDA 编程模型中,warp
我是一名优秀的程序员,十分优秀!