1、GPU架构概述

GPU是一种众核架构，非常适合解决大规模的并行计算。
GPU是CPU的协处理器，必须通过PCIe总线与基于CPU的主机（Host）相连来进行操作，形成异构架构，如下图所示。其中CPU为主机端（Host），负责逻辑控制、数据分发，GPU为设备端（Device），负责并行数据的密集型计算。其中，ALU为算数运算单元。
GPU架构是围绕一个流式多处理器（SM）的可扩展阵列搭建的。下图是英伟达公司的Fermi架构SM的示意图，SM的关键组件包括GPU核心、共享内存/一级缓存、寄存器文件、加载/存储单元、特殊功能单元和线程束调度器。
GPU的每个SM都支持数百个线程并发执行，每个GPU有多个SM，这表示一个GPU可以并发执行数千个线程。
当启动一个内核网格（Grid）时，它的线程块（block）被分布在了可用的SM上执行。
多个block可能会被分配到同一个SM上。

2、CUDA 并行计算

CUDA 是英伟达公司推出的通用并行计算平台和编程模型，它利用英伟达的GPU能够实现并行计算。
CUDA可以通过CUDA加速库、编译器指令、应用编程接口以及标准程序语言的扩展（包括C、C++、Fortan、Python）来使用。

CUDA线程模型

线程（Thread）是GPU的最小执行单元，能够完成一个逻辑操作，每个线程都有自己的指令地址计数器和寄存器状态，利用自身的数据执行当前的指令。
而线程束（Warp）是GPU的基本执行单元，包括32个线程，GPU每次调用线程都是以线程束为单位的，在一个线程束中，所有的线程按照单指令多线程（SIMT）方式执行，即所有线程执行相同的指令。
多个线程束位于一个最高维度为3的线程块（Block）中，同一个线程块中的所有线程，都可以使用共享内存来进行通信、同步。
线程块又通过一个最高维度为3的网格（Grid）来管理。

CUDA 内存模型

CUDA内存模型中，有两种类型的存储器：不可编程存储器和可编程存储器，
- 前者并不对开发人员开放存取接口，包括一级缓存和二级缓存；
- 后者可以显式地控制数据在内存空间中的存取，包括寄存器、共享内存、本地内存、常量内存、纹理内存以及全局内存。
寄存器：
- 速度最快，分配于每个线程中，数量有限，如果一个核函数使用了超过了限定数量的寄存器，将会溢出到本地内存，降低算法性能。
本地内存：
- 本地内存用来存放寄存器溢出的内存，本地内存访问符合高效内存访问要求。
共享内存：
- 比本地内存和全局内存有更高的带宽和更低的延迟。它由线程块分配，生命周期伴随着线程块，线程块中的每个线程都可以共享其存储空间。
- 一个块内的线程可以通过共享内存进行通信合作。常用的方式是将全局内存读进共享内存中，而读取的方式是每个线程负责读取某一个位置的数据，读完之后块内的所有线程都能够使用整个共享内存中的数据。
- 读取全局内存到共享内存时要注意同步，在CUDA C中使用线程同步函数__syncthreads()来实现同步。
- 核函数中存储在共享内存的变量通过修饰符__shared__修饰。
常量内存：
- 用修饰符__constant__修饰
- 必须在全局空间内和所有核函数之外声明，对同一编译单元的所有线程核函数可见
- 只读
纹理内存：
- 只读
- 适合访问二维数据
全局内存：
- 是GPU中最大、延迟最高、最常使用的内存，贯穿程序的整个生命周期。
- 对全局内存访问时，必须注意内存访问的两个特性：对齐内存访问和合并内存访问。
- 当一个线程束中全部的32个线程访问一个连续的内存块时，满足合并访问，效率非常高。