help！Thanks.
简单的说，GRID会占用整个GPU资源（也就是一个内核将完全占据GPU，在此期间其他内核不能执行），而BLOCK会被调度到各个SM上执行（一个SM上有8个BLOCK SLOT），而SP就是具体执行的单元，THREAD会在SP上得到执行！
block和sp之间没有具体的对应关系对吧，只是block内的thread在sp上去运行，是这样吗？谢谢您。
虽然TY同学是总版主，但是我必须指出，2#的说法是不正确/不确切的
1：“GRID会占用整个GPU资源（也就是一个内核将完全占据GPU，在此期间其他内核不能执行）”，此说法是不确切的，从fermi开始，GPU上可以同时跑多个小的kernel，以增加GPU的利用率。
2：“而BLOCK会被调度到各个SM上执行（一个SM上有8个BLOCK SLOT）”，此说法是有歧义的。歧义在于，理解1“一个block会被分离到多个SM上运行”，理解2“多个block会被分配到不同的SM上运行”。必须指出，“理解1”是错误的。
确切地说，一个SM上，可以resident多个blocks（具体几个根据架构不同而不同，请参阅programming guide的附录），这些resident 在同一个SM上的blocks共享同该SM的寄存器，shared memory等资源。而SM会从当前的resident blocks里面以warp为单位抓取线程来进行计算。顺便说一下，拿来计算的线程，当前需要计算的指令，可以被SM的scheduler 发射到SP阵列上，也可以发射到SFU，以及还有访存。如果当前的warp由于访存等原因需要等待，将被切换出去，换一个就绪的warp进来继续计算，这样通过线程切换，用计算掩盖了访存等延迟。
然后回答1#的问题
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。
诚如您所说，block和sp并没有直接的关系，具体关系如上。
欢迎莅临CUDA ZONE，祝您新年愉快，编码顺利~
:L俺都说了是简单的说…
谢谢各位的热心解答~Thanks.祝大家新年快乐。
做学问，可以有简单的说，但不能有错误。
ICE担心您有一丝一毫的谬误，从而不辞辛苦，为你指出每一个不恰当的地方。我觉得这是好事。
共勉吧。
学习了，“SM会从当前的resident blocks里面以warp为单位抓取线程来进行计算”这句帮助很大
同样收货很大，+1
请教2个问题：
- 1个block是不是只能resident在1个SM里
- GTX660ti的cuda core是1344，kepler架构，所以应该有7个SM，每个SM有192个SP，这么理解对吗？
- 在GTX660ti上跑一个kernel，如果block number为1，是不是gpu最多负载1/7，这么理解对吗？
一个SMX可以resident多个BLOCK(可以运行一下device query里面有)，但是一个SM里面最大活动线程数是固定的，kepler应该是2048个，如果每个block大小为512，那么可以有4个block被装入SM中。、
不错，GTX660ti是有7个SMX。
您好，目测您这是3个问题（不过无妨）：
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占用率的‘上界’，但可能不是‘上确界（最小上界）’，同时，这样做一般来说线程数量太少，没有意义。”
以上3点回答供您参考，祝您春节愉快~
此外11#的回答“一个SMX可以resident多个BLOCK(可以运行一下device query里面有)，但是一个SM里面最大活动线程数是固定的，kepler应该是2048个，如果每个block大小为512，那么可以有4个block被装入SM中。”是和您提出的第一个问题相关的另外一个问题。您问的问题可以比喻为“一件货物是否可以装入两辆卡车？”，此相关问题为“一辆卡车可以装载几件货物？”。不过此回答您依然可以作为参考。
值得补充的是，线程数量只是一个SM实际能够resident多少个block的一个限制，还有其他限制的。
此外，鉴于您在10#中表示只问两个问题，所以问题3被11#合理地忽略了。这一点也请您今后多加注意。
祝您春节愉快~编码顺利！
谢谢斑竹详尽细致的回答！
您客气了，这是我们职责所在。
欢迎您常来论坛转转！
祝您春节愉快！
请问sm每次抓取一个warp的数量的thread执行，那么如果现在sm中的sp为的数量有几种情况：
1、sp=8，那么sm每次抓取32个但是并行8个执行，共4个周期；
2、sp=32，那么每次抓取32个，并行执行32个，一个周期；
3、sp=48，a、每次抓取32个，执行32个，一个周期，那这样剩余的16个sp是怎么解释？
b、、、、
不知道该怎样理解，辛苦斑竹了
您好，按照我在4#的说法，scheduler按照warp为单位抓取线程来执行该warp当前的指令，我并没有说是只抓取一个warp过来。
SM2.1的一个SM里面除了48sp之外，还有LD/ST单元，还有SFU单元，scheduler会将当前指令为访存的warp发送到LD/ST单元执行，将需要SFU计算的warp送到SFU，会将普通SP计算指令的warp送到SP阵列。
48个SP，不能在一个周期内完成2个warp的SP计算指令（假定该指令的吞吐量为每SP每周期一条，只考虑这两个warp当前的这条SP计算指令），但这实际并不影响。scheduler和其他硬件有自己的机制协调这一过程，比如让32个sp本周期算一个warp，下周期再算一个warp，而剩下16个sp这两个连续的周期完成一个warp；也可能48个sp直接分3组，每个warp扔给一组，2周期算完，等等。理论上总是能协调48sp和32的warp size的，但是NV实际的实现方法是不公开的，同时一般也是无需考虑的。
祝您编码愉快~
斑竹您好：
如果1sm中有192个sp：
1、scheduler每次抓取6个warp，然后各个warp间顺序执行，只不过是在各个warp间有隐藏延时？那就是说增加的sp只是用来隐藏延时？
2、6个warp同时执行？
这方面还是有些困惑，辛苦斑竹帮忙分析一下，谢谢
您好，这个是同时执行的。
所有的SP都在一起干活，而且吞吐量大的指令，如乘法等，每SP每周期可以灌入一个线程的一条该类指令。
此外可能还有发射到SFU和LD/ST UNIT的。
接上楼ICE，所以楼主你不能直接用192 / 32 = 6算。加上其他的，实际可以每次发射给8组执行单元，4个warp(每个2条）。而不能光算SP.