CUDA块/线程束/线程如何映射到CUDA核心?

177

我已经使用CUDA几周了,但是对于块/线程/线程组的分配仍有一些疑问。

我从教学的角度研究这个架构(大学项目),所以达到最高性能不是我的关注点。

首先,我想了解一下我是否掌握了以下事实:

  1. 程序员编写一个内核,并将其执行组织为线程块网格。

  2. 每个块分配给一个流多处理器(SM)。一旦分配,它就不能迁移到另一个SM。

  3. 每个SM将自己的块拆分成线程束(目前具有最大大小为32的线程)。线程束中的所有线程都在SM的资源上同时执行。

  4. 线程的实际执行由SM中包含的CUDA核心执行。线程和核心之间没有特定的映射。

  5. 如果线程束包含20个线程,但当前只有16个核可用,则线程束将不会运行。

  6. 另一方面,如果块包含48个线程,则它将被拆分为2个线程束,并且只要有足够的内存可用,它们将并行执行。

  7. 如果线程在一个核上启动,然后因为内存访问或长时间的浮点操作而被阻塞,它的执行可能会在另一个核上恢复。

这些正确吗?

现在,我有一块GeForce 560 Ti显卡,根据规格说明,它配备了8个SM,每个SM包含48个CUDA核心(总共384个核心)。

我的目标是确保架构的每个核心都执行相同的指令。假设我的代码不需要更多的寄存器,超过每个SM中可用的寄存器数量,我想象了不同的方法:

  1. 我创建8个具有48个线程的块,以便每个SM都有1个要执行的块。在这种情况下,48个线程将在SM中并行执行(利用所有可用的48个核心)吗?

  2. 如果我启动64个块,每个块有6个线程,是否会产生任何区别?(假设它们将在SMs之间平均映射)

  3. 如果我在计划的工作中“淹没”GPU(例如创建1024个块,每个块有1024个线程),是否可以合理地假设所有核心在某个时刻都将被使用,并且将执行相同的计算(假设线程从未停顿)?

  4. 是否有任何方法使用分析器检查这些情况?

  5. 有没有关于这些内容的参考资料?我读过CUDA编程指南以及“大规模并行处理器编程”和“CUDA应用设计与开发”中专门讲解硬件架构的章节;但我无法得到精确的答案。


我想在注释中添加什么是“CUDA核心”。 “CUDA核心”或“执行单元”是完全流水线整数ALU和FPU,它在一个CUDA线程中每个时钟周期执行一条算术指令。 - Konstantin Burlachenko
2个回答

152

其中两份最佳的参考资料为:

  1. NVIDIA Fermi 计算架构白皮书
  2. GF104 评测

我会尝试回答你的每个问题。

程序员将工作划分为线程,线程划分为线程块,线程块划分为网格。计算工作分配器将线程块分配给流多处理器(SMs)。一旦一个线程块被分配到 SM,就会分配资源(warp 和 shared memory)并将线程分成32个线程一组的 warps。一旦一个warp被分配,它就被称为活动warp。两个warp调度器每个周期选择两个活动warp,并将warp发送到执行单元。有关执行单元和指令调度的更多详细信息,请参见 1 p.7-10 和 2

4'. laneid(warp中的线程索引)和核之间存在映射关系。

5'. 如果一个warp包含少于32个线程,它大多数情况下会像有32个线程一样执行。warp可以有少于32个活动线程,原因有几种:每个块的线程数不可被32整除,程序执行了分歧块,因此未参与当前路径的线程被标记为不活动,或者一个线程在 warp 中退出。

6'. 一个线程块将被分成 WarpsPerBlock = (ThreadsPerBlock + WarpSize - 1) / WarpSize

7'. 执行单元不会因为内存操作而停顿。如果一个资源在指令可以被分配时不可用,那么该指令将在未来再次分配,直到资源可用。Warp会在屏障、内存操作、纹理操作、数据依赖等方面停滞不前... 停滞的Warp不符合被Warp调度器选择的条件。对于Fermi架构,每个周期至少有2个符合条件的Warp非常有用,这样Warp调度器就可以发出一条指令。

有关GTX480和GTX560之间的差异,请参见2

如果您阅读参考资料(几分钟),我认为您会发现您的目标是没有意义的。我将尝试回答您的问题。

1'. 如果您启动kernel<<<8, 48>>>,则会得到8个块,每个块具有32个warp和16个线程的2个warp。不能保证这8个块将分配给不同的SM。如果2个块分配给一个SM,则每个warp调度器可以选择一个warp并执行warp。您只会使用48个核心中的32个。

2'. 48个线程和64个6个线程之间存在很大的差异。假设您的kernel没有分歧,每个线程执行10条指令。

  • 48个线程的8个块=16个warp * 10条指令=160条指令
  • 64个6个线程的块=64个warp * 10条指令=640条指令

为了获得最佳效率,工作的划分应该是32个线程的倍数。硬件不会将来自不同的warp的线程合并。

3'. GTX560可以同时拥有8个SM * 8个块 = 64个块,或者8个SM * 48个线程束 = 512个线程束,如果内核没有用尽寄存器或共享内存。在任何给定的时间,一部分工作只会在SM上活跃。每个SM都有多个执行单元(比CUDA核心更多)。在任何给定的时间使用哪些资源取决于应用程序的线程束调度程序和指令混合。如果不进行TEX操作,则TEX单元将处于空闲状态。如果不执行特殊的浮点运算,则SUFU单元将处于空闲状态。

4'. Parallel Nsight和Visual Profiler显示:

a. 执行IPC

b. 发出IPC

c. 活动线程束每个活动周期

d. 每个活动周期的合格线程束(仅适用于Nsight)

e. 线程束停顿原因(仅适用于Nsight)

f. 每个执行的指令的活动线程数

分析器不显示任何执行单元的利用率百分比。对于GTX560,一个粗略的估计是IssuedIPC / MaxIPC。 对于MaxIPC,假设 GF100(GTX480)为2 GF10x(GTX560)为4,但3是更好的目标。


1
谢谢你的回答。我阅读了参考文献,但是在你的回答中有一些我不理解的地方。在接下来的问题中,我假设我们使用 Fermi 架构,具有 48 个核心(16 个核心 * 3 "核心组"):
  1. 你提到了核心和 laneid 之间的映射。这是什么样的映射?
  2. 根据参考文献,我得知每个 "核心组" 每个时钟周期执行最多半个 warp(16 个线程)。所以理论上如果我们在同一个块中有 48 个线程,它们将被组织成 3 个半 warp 并在 48 个核心上并行执行。我对吗?
- Daedalus
2
CUDA核心是单精度FP单元的数量。以CUDA核心为执行单位的思考方式是不正确的。每个warp有32个线程。这些线程将被分配到一组执行单元(例如16个CUDA核心)。为了在一个时钟周期内发出所有48个核心,两个warp调度器中的一个需要选择一个满足超标量对要求的warp,并且两个指令都需要由CUDA核心执行。此外,另一个warp调度器必须选择一个下一个指令将由CUDA核心执行的warp。 - Greg Smith
4
在你的例子中,每个调度器都会选择一个warp并发出一条指令。在这种情况下,只会使用两组执行单元。为了使用更多的执行单元,其中一个调度器必须进行双重发行。正如参考文献中所示,有多种类型的执行单元(不仅仅是所谓的cuda核心),并且必须满足指令配对规则(未经充分记录)才能进行双重发行。 - Greg Smith
1
@GregSmith,我正在网上搜索有关Fermi架构中每个SM的8个活动块来自何处的信息。这甚至在Fermi白皮书中都没有提到。你有更多相关的参考资料吗? - Greg K.
1
@GregKasapidis 请查看CUDA编程指南(http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities)第G节。 计算能力表13中的“每个多处理器的最大常驻块数”技术规格行。 - Greg Smith
显示剩余6条评论

9

"E. 如果一个warp包含20个线程,但目前只有16个核心可用,该warp将无法运行。"

这是不正确的。你混淆了通常意义上的核心(也用于CPU中)- GPU中“多处理器”的数量,与nVIDIA市场术语中的核心(“我们的显卡有数千个CUDA核心”)。Cuda核心so answer)是一个硬件概念,而线程是一个软件概念。即使只有16个核心可用,您仍然可以运行32个线程。但是,您可能需要2个时钟周期才能使用仅有的16个硬件核心来运行它们。

CUDA核心计数表示每个周期可以执行的单精度浮点或整数线程指令的总数。

warp调度程序负责查找要运行指令的核心(so answer)。

一个warp是32个执行线程的逻辑集合。为了执行单个warp中的单个指令,warp调度程序通常必须调度32个执行单元(或“核心”,尽管“核心”的定义有些宽松)。
一个warp本身只能被调度到SM(多处理器或流式多处理器),并且可以同时运行多达32个线程(取决于SM中的核心数);它不能使用超过一个SM。
在具有计算能力2.x的NVIDIA GPU上,"48个warp"的数字是每个多处理器的活动warp的最大数量(在任何给定周期,可能会选择这些warp来安排工作),这个数字对应于1536 = 48 x 32个线程。
基于this webinar的答案。

@GregSmith:编辑了答案以解决这个问题。你对它的耐心是好的,但是已经过去五年了…… - einpoklum
2
单核心(=多处理器)?我认为问题假设单核心=处理器而不是多处理器。根据您的术语,您的答案是正确的。 - Adarsh
据我所知,要找到允许的最大线程束数,可以使用Cuda Occupency calculator,并从“每个多处理器的最大线程束数”行中读取。 然后,我有一台GTX 770(SDK 3.0),我确定最佳调用方式如下: MaxWrapPerMultiprocessor/WarpAllocationGranuality = BlockPerSm; TotalBlock=BlockPerSm*NbrSm; TotalThreadPerBlock= WarpAllocationGranuality * ThreadPerWarp; 然后:64/4=16;TotalBlock=16*8=128;TotalThreadPerBlock=4*32=128; 然后像这样调用我的cuda函数:myfunc<<<128,128>>>(); - OOM
CUDA核心数表示每个周期可以执行的单精度浮点或整数线程指令的总数。https://dev59.com/BWQn5IYBdhLWcg3wPlKj#16987220。CUDA核心是从硬件角度和软件角度的线程。此外,多处理器是指SM。 - Izana

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接