线程结构 & 索引
概念
- 存储带宽：在一定时间内从DRAM读出或写入的数据量
- 延迟：响应一个获取内存的请求所花费的时间，这个时间通常是上百个处理器周期
- SM, SP, Grid, Block, thread, warp：
SM,SP是硬件结构； Grid, Block, thread是软件概念
从硬件角度讲：SM(流多处理器)，一个SM可以看作是一个多线程的CPU核，一个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个。这意味着每个SM在任何时刻都能同时运行这些数目的硬件线程。
从软件角度讲：Grid,Block,Thread是线程(thread)的组织形式，最小的逻辑单位是thread,若干个（典型值是128~512个）线程组成一个block，block被加载到SM上运行（每个SM能调度若干个block），多个block组成整体的grid。最小的硬件执行单位是warp（线程束），当前一个线程束包含32个线程。
总而言之，一个kernel对应一个grid，该grid又包含若干个block，block内包含若干个thread。grid跑在GPU上的时候，可能是独占一个GPU的，也可能是多个kernel并发占用一个GPU的（需要fermi及更新的GPU架构支持）。
- 一些问题：
1. 1个block是不是只能resident在1个SM里？
是的，您可以这样辅助考虑，如果一个block要使用shared memory，此时注意到shared memory是SM上的资源，不同的SM上shared memory是不通信的，也不能互相借用。所以，可以反证，一个block只能resident在一个SM上。
2. GTX660ti的cuda core是1344，kepler架构，所以应该有7个SM，每个SM有192个SP，这么理解对吗？
对的
3. 在GTX660ti上跑一个kernel，如果block number为1，是不是gpu最多负载1/7，这么理解对吗？
这个问题说起来稍微有点复杂，因为这个和该block使用资源的情况有关，一般情况下，是无法达到1/7的，也就是说只上一个block的话很可能一个SM都跑不满，（比如这个block里面线程数量非常少，或者线程数量中等但是依然无法掩盖其他的延迟等）同时一个block最大只能有1024个线程，这对于GPU计算还是少了些。
简单地借用一个数学的概念来说明可能更为明了：“1/7是您GPU占用率的‘上界’，但可能不是‘上确界（最小上界）’，同时，这样做一般来说线程数量太少，没有意义。”
寄存器的用法
CPU与GPU映射寄存器的方式不同。CPU通过使用寄存器重命名和栈来执行多线程；GPU利用多线程隐藏了内存获取与指令执行带来的延迟，GPU不使用寄存器重命名机制，而是致力于为每一个线程都分配真实的寄存器。因此当需要上下文切换时，所需要的操作就是将指向当前寄存器组的选择器（或指针）更新，以指向下一个执行的线程束的寄存器组，因此几乎是零开销。
每个SM能调度若干个线程块。在SM层，线程块即若干个独立线程束的逻辑组。编译时会计算出每个内核线程所需要的寄存器数目。所有的线程块都具有相同大小，并拥有已知数目的线程，每个线程块需要的寄存器数目也就是已知和固定的。因此，GPU就能为在硬件上调度的线程块分配固定数目的寄存器。
然而在线程层，这些细节对程序员是完全透明的。如果一个内核函数中的每个线程需要的寄存器过多，在每个SM中GPU能够调度的线程块的数量就会受到限制，因此总的可以执行的线程数量就会受到限制。开启的线程数量过少就会赵成硬件无法被充分利用，性能急剧下降，但开启过多又意味着资源可能短缺，调度到SM上的线程块数量会减少。
根据硬件的不同，每个SM可供所有线程使用的寄存器空间大小也不同，分别有8KB,16KB,32KB以及64KB.牢记，每个线程中的每个变量会占用一个寄存器
总之，在编写GPU程序是，为了充分发挥硬件性能和程序速度，应考虑好寄存器数，线程数，SM调度的线程块数等的关系。同时最大化地利用寄存器。
索引
内建变量
-
threadIdx(.x/.y/.z代表几维索引)：线程所在block中各个维度上的线程号
-
blockIdx(.x/.y/.z代表几维索引)：块所在grid中各个维度上的块号
-
blockDim(.x/.y/.z代表各维度上block的大小)：block的大小即block中线程的数量，blockDim.x代表块中x轴上的线程数量，blockDim.y代表块中y轴上的线程数量，blockDim.z代表块中z轴上的线程数量
-
gridDim(.x/.y/.z代表个维度上grid的大小)：grid的大小即grid中block的数量，gridDim.x代表grid中x轴上块的数量，gridDim.y代表grid中y轴上块的数量，gridDim.z代表grid中z轴上块的数量
定义grid、block大小：
dim3 numBlock(m,n)
dim3 threadPerBlock(i,j)
则 blockDim.x=i;
blockDim.y=j;
gridDim.x=m;
gridDim.y=n
索引
1、 grid划分成1维，block划分为1维
int threadId = blockIdx.x *blockDim.x + threadIdx.x;
2、 grid划分成1维，block划分为2维
int threadId = blockIdx.x * blockDim.x * blockDim.y+ threadIdx.y * blockDim.x + threadIdx.x;
3、 grid划分成1维，block划分为3维
int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z
+ threadIdx.z * blockDim.y * blockDim.x
+ threadIdx.y * blockDim.x + threadIdx.x;
4、 grid划分成2维，block划分为1维
int blockId = blockIdx.y * gridDim.x + blockIdx.x;
int threadId = blockId * blockDim.x + threadIdx.x;
5、 grid划分成2维，block划分为2维
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y)
+ (threadIdx.y * blockDim.x) + threadIdx.x;
6、 grid划分成2维，block划分为3维
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x) + threadIdx.x;
7、 grid划分成3维，block划分为1维
int blockId = blockIdx.x + blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * blockDim.x + threadIdx.x;
8、 grid划分成3维，block划分为2维
int blockId = blockIdx.x + blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * (blockDim.x * blockDim.y)
+ (threadIdx.y * blockDim.x) + threadIdx.x;
9、 grid划分成3维，block划分为3维
int blockId = blockIdx.x + blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x) + threadIdx.x;
