最近应用开发的过程中出现了一个小问题，顺便记录一下原因和方法--数量线程
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占用率的‘上界’，但可能不是‘上确界（小最上界）’，同时，这样做一般来说线程数量太少，没有意义。”
文章结束给大家分享下程序员的一些笑话语录： PC软件体积大，是因为一个PC软件功能往往较多，能够满足你一个方面的需求，而一个iphone软件往往没几行代码，干一件很小的事情，自然需要的软件就多。就像吃西瓜和吃瓜子的来比数目，单位不同啊。