streaming processor(sp): 最基本的处理单元，streaming processor 最后具体的指令和任务都是在sp上处理的。GPU进行并行计算，也就是很多个sp同时做处理。现在SP的术语已经有点弱化了，而是直接使用thread来代替。一个SP对应一个thread。

Warp：warp是SM调度和执行的基础概念，同时也是一个硬件概念，注意到Warp实际上是一个和硬件相关的概念，通常一个SM中的SP(thread)会分成几个warp(也就是SP在SM中是进行分组的，物理上进行的分组)，每一个WARP中在Tegra中是32个thread.这个WARP中的32个thread(sp)是一起工作的，执行相同的指令，如果没有这么多thread需要工作，那么这个WARP中的一些thread(sp)是不工作的。

每一个线程都有自己的寄存器内存和local memory，一个warp中的线程是同时执行的，也就是当进行并行计算时，线程数尽量为32的倍数，如果线程数不上32的倍数的话；假如是1，则warp会生成一个掩码，当一个指令控制器对一个warp单位的线程发送指令时，32个线程中只有一个线程在真正执行，其他31个 进程会进入静默状态。

每个SM通过使用两个特殊函数(Special Function Unit,SFU)单元进行超越函数和属性插值函数（根据顶点属性来对像素进行插值）计算。SFU用来执行超越函数、插值以及其他特殊运算

Stream：流（Stream）是一系列顺序执行的命令，流之间相对无序或并发的执行他们的命令。

thread-->block-->grid：在利用cuda进行编程时，一个grid分为多个block，而一个block分为多个thread。其中任务划分到是否影响最后的执行效果。划分的依据是任务特性和GPU本身的硬件特性。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

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

从上面的表述中可以总结：

在GPU中最小的硬件单元是SP(这个术语通常使用thread来代替),而硬件上一个SM中的所有SP在物理上是分成了几个WARP(每一个warp包含一些thread),warp中的SP是可以同时工作的，但是执行相同的指令，也就是说取指令单元取一条指令同时发射给WARP中的所有的SP(假设SP都需要工作，否则有些是idle的).可见，在硬件上一个SM->WARPS->threads(sp).

对于软件thread组织来看，因为一个SM中是分WARP的，而一个WARP包含一定数目（比如Tegra 32个)的sp(thread),因此最好按照这个数目来组织thread，否则硬件该warp上有些SP是不工作的。

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

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

GPU中的几种memory及其在系统中的位置：

GPU中由NVIDIA公司生产的GPU引入。不同于CPU中通过SIMD（单指令多数据）来处理；GPU则使用SIMT，SIMT的好处是无需开发者费力把数据凑成合适的矢量长度，并且SIMT允许每个线程有不同的分支。 纯粹使用SIMD不能并行的执行有条件跳转的函数，很显然条件跳转会根据输入数据不同在不同的线程中有不同表现，这个只有利用SIMT才能做到。

下面几张硬件结构简图 便于理解(图片来源于网上）

以上两图可以清晰地表示出sm与sp的关系。

目前市场上的NVIDIA显卡都是基于Tesla架构的，分为G80、G92、GT200三个系列。Tesla体系架构是一块具有可扩展处器数量的处理器阵列。每个GT200 GPU包含240个流处理器（streaming processor,SP），每8个流处理器又组成了一个流多处理器(streaming multiprocessor,SM)，因此共有30个流多处理器。GPU在工作时，工作负载由PCI-E总线从CPU传入GPU显存，按照体系架构的层次自顶向下分发。PCI-E 2.0规范中，每个通道上下行的数据传输速度达到了5.0Gbit/s，这样PCI-E2.0×16插槽能够为上下行数据各提供了5.0*16Gbit/s=10GB/s的带宽，故有效带宽为8GB/s,而PCI-E 3.0规范的上下行数据带宽各为20GB/s。但是由于PCI-E数据封包的影响，实际可用的带宽大约在5-6GB/s（PCI-E 2.0 ×16）。

（multiply-add units，乘加器）。它们能够对符合IEEE标准的单精度浮点数（对应float型）和32-bit整数（对应int型，或者unsigned int型）进行运算。每次运算需要4个时钟周期（SP周期，并非核心周期）。因为使用了四级流水线，因此在每个时钟周期，ALU或MAD都能取出一个warp的32个线程中的8个操作数，在随后的3个时钟周期内进行运算并写回结果。

Thread ：并行运算的基本单位（轻量级的线程）
Block ：由相互合作的一组线程组成。一个block中的thread可以彼此同步，快速交换数据，最多可以同时512个线程。
Grid ：一组Block，有共享全局内存

CUDA在执行的时候是让host里面的一个一个的kernel按照线程网格（Grid）的概念在显卡硬件（GPU）上执行。每一个线程网格又可以包含多个线程块（block），每一个线程块中又可以包含多个线程（thread）。每一个kernel交给每一个Grid来完成。当要执行这些任务的时候，每一个Grid又把任务分成一部分一部分的block，block再分线程来完成。每个Grid中的任务是一定的。二维线程块的索引关系为如下：

memory。每一个时钟周期内，warp（一个block里面一起运行的thread）包含的thread数量是有限的，现在的规定是32个。一个block中含有16个warp。所以一个block中最多含有512个线程.每次Device（就是显卡）只处理一个grid。

此图反应了warp作为调度单位的作用，每次GPU调度一个warp里的32个线程执行同一条指令，其中各

个线程对应的数据资源不同(指令相同但是数据不同)。

此图是一个warp排程的例子。

进行划分时，最好保证每个block里的warp比较合理，那样可以一个sm可以交替执行里面的warp，从而提高

效率，此外，在分配block时，要根据GPU的sm个数，分配出合理的block数，让GPU的sm都利用起来，提

利用率。分配时，也要考虑到同一个线程block的资源问题，不要出现对应的资源不够。

CUDA库包含了很多有用的数学应用，如cuFFT，CUDA RUNTIME其实就是个JIT编译器，动态的将PTX中间代码编译成符合实际平台的硬件代码，并作特定优化。Driver便是直接相应API与GPU打交道的接口了。Nvcc编译器可以将CUDA C编译成纯C或者PTX或者GPU BIN。

在CUDA中程序执行区域分为两部分，CPU和GPU——HOST和DEVICE，任务组织和发送是在CPU里完成的，但并行计算是在GPU里完成，每当CPU遇到需要并行计算的任务，则将要做的运算组织成kernel，然后丢给GPU去执行，当然任务是通过CUDA系统来丢，CUDA在把任务正式提交给GPU前，会对kernel做些处理，让kernel符合GPU体系架构（接下来几个概念是有对应的硬件的），现在先简单的把GPU想想成拥有上百个核的CPU，kernel当成一个要创建为线程的函数，所以CUDA现在要将你的kernel创建出上百个thread，然后将这些thread送到GPU中的各个核上去运行，但为了更好的利用GPU资源，提高并行度，CUDA还要将这些thread加以优化组织，将能利用共有资源的线程组织到一个thread

提前先放个硬件体系图，下面的讲解大家会理解深入很多。

4BYTE用来存放单精度浮点数或者整数，而双精度浮点则需要占用相邻2个entry，注意这些ENTRY是JIT/DRIVER动态分配给thread的。我们在Shader中定义的局部变量一般都是分配在RF中，当RF不够用了就分到Local Memory中，可简单的将Local

协处理器（coprocessor），一种芯片，用于减轻系统微处理器的特定处理任务。

CPU: Center Process Unit的缩写，译为中央处理器。也做叫微处理器。指具有运算器和控制器功能的大规模集成电路。

Unit）图形处理芯片。是显示卡的“心脏”，也就相当于CPU在电脑中的作用，它决定了该显卡的档次和大部分性能，同时也是2D显示卡和3D显示卡的区别依据。2D显示芯片在处理3D图像和特效时主要依赖CPU的处理能力，称为“软加速”。3D显示芯片是将三维图像和特效处理功能集中在显示芯片内，也即所谓的“硬件加速”功能。显示芯片通常是显示卡上最大的芯片。

FPU：（Float Point Unit）浮点运算单元,FPU是专用于浮点运算的处理器，以前的FPU是一种单独芯片，在486之后，英特尔把FPU与集成在CPU之内。

本笔记引用自PyTorch中文文档

• 每个张量tensor都有一个相应的torch.Storage用来保存其数据。会改变tensor的函数操作会用一个下划线后缀来标示。

Tensor类的构造函数：

将函数callable作用于tensor中每一个元素，并将每个元素用callable函数返回值替代。该函数只能在CPU tensor中使用，且不应用在有较高性能要求的代码块

返回与原tensor有相同大小和数据类型的tensor

返回一个内在连续的有相同数据的tensor, 如果原tensor内存连续则返回原tensor。

如果在CPU上没有该tensor，则会返回一个CPU的副本。

返回此对象在GPU内存中的一个副本。若已在CUDA存储中并且在正确的设备上，则不会进行复制并返回原对象。

返回tensor第一个元素的地址。

将该tensor用指定的数值填充。

将tensor投射为半精度浮点类型

如果tensor在内存中是连续的则返回True。

沿着指定的维度重复tensor，不同于expand()，本函数是复制tensor中的数据

将对象投为指定的类型，如果已经是正确的类型，则不会复制并返回原对象。