GG_Bruse的博客

CUDA是并行计算的平台和类C编程模型，能很容易的实现并行算法。只需配备NVIDIA GPU，就可以在许多设备上运行并行程序。

nvidia-smi 提供了GPU技术参数、使用温度和电源管理等重要信息。

假设要启动一些线程块且每个线程块中均包含 256 个线程（32 的倍数），并需运行 1000 个并行任务（此处使用极小的数量以便于说明），则任何数量的线程块均无法在网格中精确生成 1000 个总线程，因为没有任何整数值在乘以 32 后可以恰好等于1000。以下是编写执行配置的惯用方法示例，适用于 N 和线程块中的线程数已知，但无法保证网格中的线程数和 N 之间完全匹配的情况。GPU 函数称为核函数，核函数通过执行配置启动，执行配置定义了网格中的块数以及每个块中的线程数，网格中的每个块均包含相同数量的线程。

SM（Streaming Multiprocessors）是GPU架构中非常重要的部分，GPU硬件的并行性就是由SM决定的。GPU中每个SM都设计成支持数以百计的线程并行执行，并且每个GPU都包含了很多的SM，所以GPU支持成百上千的线程并行执行。当一个kernel启动后，thread会被分配到这些SM中执行。大量的thread可能会被分配到不同的SM，但是同一个block中的thread必然在同一个SM中并行执行CUDA采用 Single Instruction Multiple Thread（

SM 中 warp调度器 每个 cycle 会挑选 active warp 送去执行，一个被选中的 warp 被称为 selected warp，没被选中，但是已经做好准备被执行的称为 Eligible warp，没准备好要执行的称为 Stalled warp。在GPU编程中，预测变量是一个用于指示分支执行路径的变量，其值为1或0。同一个 warp 中的 thread 必须执行相同的指令，若这些线程在遇到控制流语句时，进入了不同的分支，那么同一时刻除了正在执行的分支外，其余分支都会被阻塞，十分影响性能。

从高级语言层面是无法看出性能提升的原因的，需从 low-level instruction 层面去分析，第二段代码循环次数减少了一半，而循环体两句语句的读写操作的执行在CPU上是可以同时执行互相独立的，所以相对第一段，第二段性能要好。Fermi 和 Kepler 每个 block 的最大 thread 数目都是1024，文中的 kernel 的迭代次数都是基于 blockDim 的，所以完全展开循环是可行的。由于使用的SIMT模式，warp内的thread 是有一个隐式的同步过程的。stride > 0;

Local，Global，Contant，Texture为片外DRAM，其中Global，Constant，Texture内存在Host端代码声明，所有线程可见SM拥有私有的Registers和Shared Memory（其实还有SM私有的L1 cache以及共有的L2 cache），Constant和Texture内存有专有的Caches(片上)

GPU对片外DRAM的访问延迟大，带宽低，导致其成为很多应用的性能瓶颈。因此对DRAM访问的进一步优化可以有效改善程序性能。优化之前，首先看一下GPU对内存的访问模式如上图所示，DRAM内存的读写在物理上是从 片外DRAM -> 片上Cache -> 寄存器。其中，片外DRAM到片上Cache是主要性能瓶颈。DRAM 到 Cache 之间的一次传输(transaction)设计为32、64或者128字节，并且内存地址按照32、64或者128字节的间隔对齐。

GPU有两种类型的内存global memory 就是一块很大的板载内存，具有相对较高的延迟。shared memory 是较小的片上内存，具有相对较低的延迟，并且 shared memory 可以提供比 global memory 高得多的带宽。可将其当作一个可编程的cache共享内存(SMEM)是 GPU 的一个关键部件。物理上，每个 SM 都有一个小的低延迟内存池，这个内存池被当前正在该 SM 上执行的线程块中的所有线程所共享。

常量内存对内核代码而言是只读的，但它对主机而言即是可读的又是可写的。常量内存位于设备的DRAM上（和全局内存一样）。有一个专用的片上缓存，从每个SM的常量缓存中读取的延迟，比直接从常量内存中读取要低得多。每个SM常量内存缓存大小的限制为64KBconstant Memory 的获取方式不同于其它的 GPU 内存，对于 constant Memory 来说，最佳获取方式是 warp 中的 32 个thread获取 constant Memory 中的同一个地址。

kernel level，即一个 kernel 或者一个 task 由许多 thread 并行的执行在GPU上。Stream的概念是相对于后者来说的，Grid level是指多个 kernel 在一个 device 上同时执行（一个 Grid 中的 Block 可以在多个 SM 中执行）

CUDA 为许多常用编程语言提供扩展，如 C、C++、Python 和 Fortran 等语言。CUDA 加速程序的文件扩展名是.cu下面包含两个函数，第一个函数将在 CPU 上运行，第二个将在 GPU 上运行__global__ 关键字表明以下函数将在 GPU 上运行并可全局调用将在 CPU 上执行的代码称为主机代码，而将在 GPU 上运行的代码称为设备代码注意返回类型为 void，使用 __global__ 关键字定义的函数要求返回 void 类型。