程序员的自我修养_Will.zhang

关注多核计算&高性能计算

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

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占用率的‘上界’,但可能不是‘上确界(最小上界)’,同时,这样做一般来说线程数量太少,没有意义。”

阅读更多
个人分类: CUDA
所属专栏: CUDA
想对作者说点什么? 我来说一句

没有更多推荐了,返回首页

加入CSDN,享受更精准的内容推荐,与500万程序员共同成长!
关闭
关闭