SM,SP和GRID,BLOCK,THREAD之间的对应关系是什么?

6 篇文章 7 订阅
SM,SP是硬件结构
GRID,BLOCK,THREAD是软件概念


从硬件角度讲,一个GPU由多个SM组成(当然还有其他部分),一个SM包含有多个SP(以及还有寄存器资源,shared memory资源,L1cache,scheduler,SPU,LD/ST单元等等),1.x硬件,一个SM包含8个SP,2.0是32个,2.1是48个,3.0和3.5是192个。以及SP目前也称为CUDA CORE,而SM目前也称为MP,在KEPLER架构(SM3.0和3.5)下也称为SMX。


从软件角度讲,CUDA因为是SIMT的形式,GRID,block,thread是thread的组织形式。最小的逻辑单位是一个thread,最小的硬件执行单位是thread warp(简称warp),若干个thread(典型值是128~512个)组成一个block,block被加载到SM上运行,多个block组成整体的GRID。

这里为什么要有一个中间的层次block呢?这是因为CUDA通过这个概念,提供了细粒度的通信手段,因为block是加载在SM上运行的,所以可以利用SM提供的shared memory和__syncthreads()功能实现线程同步和通信,这带来了很多好处。而block之间,除了结束kernel之外是无法同步的,一般也不保证运行先后顺序,这是因为CUDA程序要保证在不同规模(不同SM数量)的GPU上都可以运行,必须具备规模的可扩展性,因此block之间不能有依赖。

这就是CUDA的两级并行结构。

总而言之,一个kernel对应一个GRID,该GRID又包含若干个block,block内包含若干个thread。GRID跑在GPU上的时候,可能是独占一个GPU的,也可能是多个kernel并发占用一个GPU的(需要fermi及更新的GPU架构支持)。

block是resident在SM上的,一个SM可能有一个或多个resident blocks,需要具体根据资源占用分析。

thread以warp为单位被SM的scheduler 发射到SP或者其他单元,如SFU,LD/ST unit执行相关操作,需要等待的warp会被切出(依然是resident 状态),以空出执行单元给其他warps。


那么有问题 

1. 1个block是不是只能resident在1个SM里
2. GTX660ti的cuda core是1344,kepler架构,所以应该有7个SM,每个SM有192个SP,这么理解对吗?
3. 在GTX660ti上跑一个kernel,如果block number为1,是不是gpu最多负载1/7,这么理解对吗?


ice大神这样回复


1:是的,您可以这样辅助考虑,如果一个block要使用shared memory,此时注意到shared memory是SM上的资源,不同的SM上shared memory是不通信的,也不能互相借用。所以,可以反证,一个block只能resident在一个SM上。

2:kepler架构下的SM(又称SMX)是拥有192个SP(又称CUDA CORE)没错,因而使用总的SP数量除以192即得 SM数量。您的理解是正确的。

3:这个问题说起来稍微有点复杂,因为这个和该block使用资源的情况有关,一般情况下,是无法达到1/7的,也就是说只上一个block的话很可能一个SM都跑不满,(比如这个block里面线程数量非常少,或者线程数量中等但是依然无法掩盖其他的延迟等)同时一个block最大只能有1024个线程,这对于GPU计算还是少了些。
简单地借用一个数学的概念来说明可能更为明了:“1/7是您GPU占用率的‘上界’,但可能不是‘上确界(最小上界)’,同时,这样做一般来说线程数量太少,没有意义。”

  • 2
    点赞
  • 10
    收藏
    觉得还不错? 一键收藏
  • 0
    评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值