CUDA
待学习：
- Synchronized
- Overlaped
英伟达CUDA介绍
CUDA6中的Unified memory
CUDA简介
CPU & GPU
GPU包括更多的运算核心，其特别适合数据并行的计算密集型任务，如大型矩阵运算，而CPU的运算核心较少，但是其可以实现复杂的逻辑运算，因此其适合控制密集型任务。另外，CPU上的线程是重量级的，上下文切换开销大，但是GPU由于存在很多核心，其线程是轻量级的。因此，基于CPU+GPU的异构计算平台可以优势互补，CPU负责处理逻辑复杂的串行程序，而GPU重点处理数据密集型的并行计算程序，从而发挥最大功效。
CUDA是NVIDIA公司所开发的GPU编程模型，它提供了GPU编程的简易接口，基于CUDA编程可以构建基于GPU计算的应用程序
CUDA编程
CUDA编程模型是一个异构模型，需要CPU和GPU协同工作。在CUDA中，host和device是两个重要的概念，我们用host指代CPU及其内存，而用device指代GPU及其内存。CUDA程序中既包含host程序，又包含device程序，它们分别在CPU和GPU上运行。同时，host与device之间可以进行通信，这样它们之间可以进行数据拷贝。典型的CUDA程序的执行流程如下：
- 分配host内存，并进行数据初始化；
- 分配device内存，并从host将数据拷贝到device上；
- 调用CUDA的核函数在device上完成指定的运算；
- 将device上的运算结果拷贝到host上；
- 释放device和host上分配的内存。
CUDA逻辑层
kernel是CUDA中一个重要的概念，kernel是在device上线程中并行执行的函数，核函数用__global__符号声明，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量，在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号thread ID。
kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间，grid是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。
CUDA物理层
一个kernel实际上会启动很多线程，这些线程是逻辑上并行的，但是在物理层却并不一定。这其实和CPU的多线程有类似之处，多线程如果没有多核支持，在物理层也是无法实现并行的。但是好在GPU存在很多CUDA核心，充分利用CUDA核心可以充分发挥GPU的并行计算能力。
GPU硬件的一个核心组件是SM，前面已经说过，SM是英文名是 Streaming Multiprocessor，翻译过来就是流式多处理器。SM的核心组件包括CUDA核心，共享内存，寄存器等，SM可以并发地执行数百个线程，并发能力就取决于SM所拥有的资源数。当一个kernel被执行时，它的gird中的线程块被分配到SM上，一个线程块只能在一个SM上被调度。SM一般可以调度多个线程块，这要看SM本身的能力。那么有可能一个kernel的各个线程块被分配多个SM，所以grid只是逻辑层，而SM才是执行的物理层。
SM采用的是SIMT (Single-Instruction, Multiple-Thread，单指令多线程)架构，基本的执行单元是线程束（wraps)，线程束包含32个线程，这些线程同时执行相同的指令，但是每个线程都包含自己的指令地址计数器和寄存器状态，也有自己独立的执行路径。所以尽管线程束中的线程同时从同一程序地址执行，但是可能具有不同的行为，比如遇到了分支结构，一些线程可能进入这个分支，但是另外一些有可能不执行，它们只能死等，因为GPU规定线程束中所有线程在同一周期执行相同的指令，线程束分化会导致性能下降。
当线程块被划分到某个SM上时，它将进一步划分为多个线程束，因为这才是SM的基本执行单元，但是一个SM同时并发的线程束数是有限的。这是因为资源限制，SM要为每个线程块分配共享内存，而也要为每个线程束中的线程分配独立的寄存器。所以SM的配置会影响其所支持的线程块和线程束并发数量。总之，就是网格和线程块只是逻辑划分，一个kernel的所有线程其实在物理层是不一定同时并发的。所以kernel的grid和block的配置不同，性能会出现差异，这点是要特别注意的。
还有，由于SM的基本执行单元是包含32个线程的线程束，所以block大小一般要设置为32的倍数。
CPU和GPU间数据传输
Pinned Host Memory
参考
主机(CPU)数据分配默认是分页的。GPU不能直接访问分页主机内存中的数据。
CUDA驱动先分配(allocate)一个临时的page-locked,or pinned host array，再将主机的数据拷贝到pinned array，再转换到device memory。
我们可以直接将主机数据分配到pinned memory中，减少一次传输。在C++中使用cudaMallocHost()函数。
Unified Memory in CUDA 6
参考自英伟达官方文档
CPU和GPU中的数据必须分开放置在各自的存储中，通过程序进行复制传递。
Unified Memory创建了一个pool of managed memory来共享二者之间的数据。只通过同一个指针，GPU和CPU就都可以访问这块managed memory。
实际上，系统会自动在host和device之间迁移管理managed memory中的数据，使得数据看起来像是在一块可共享的内存中。所以我们只需要进行一次内存分配，就可以得到一个能从host和device访问的指向数据的指针。
英伟达CUDA介绍
CUDA简介
CPU & GPU
GPU包括更多的运算核心，其特别适合数据并行的计算密集型任务，如大型矩阵运算，而CPU的运算核心较少，但是其可以实现复杂的逻辑运算，因此其适合控制密集型任务。另外，CPU上的线程是重量级的，上下文切换开销大，但是GPU由于存在很多核心，其线程是轻量级的。因此，基于CPU+GPU的异构计算平台可以优势互补，CPU负责处理逻辑复杂的串行程序，而GPU重点处理数据密集型的并行计算程序，从而发挥最大功效。
CUDA是NVIDIA公司所开发的GPU编程模型，它提供了GPU编程的简易接口，基于CUDA编程可以构建基于GPU计算的应用程序
CUDA编程
CUDA编程模型是一个异构模型，需要CPU和GPU协同工作。在CUDA中，host和device是两个重要的概念，我们用host指代CPU及其内存，而用device指代GPU及其内存。CUDA程序中既包含host程序，又包含device程序，它们分别在CPU和GPU上运行。同时，host与device之间可以进行通信，这样它们之间可以进行数据拷贝。典型的CUDA程序的执行流程如下：
- 分配host内存，并进行数据初始化；
- 分配device内存，并从host将数据拷贝到device上；
- 调用CUDA的核函数在device上完成指定的运算；
- 将device上的运算结果拷贝到host上；
- 释放device和host上分配的内存。
CUDA逻辑层
kernel是CUDA中一个重要的概念，kernel是在device上线程中并行执行的函数，核函数用__global__符号声明，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量，在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号thread ID。
kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间，grid是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。
CUDA物理层
一个kernel实际上会启动很多线程，这些线程是逻辑上并行的，但是在物理层却并不一定。这其实和CPU的多线程有类似之处，多线程如果没有多核支持，在物理层也是无法实现并行的。但是好在GPU存在很多CUDA核心，充分利用CUDA核心可以充分发挥GPU的并行计算能力。
GPU硬件的一个核心组件是SM，前面已经说过，SM是英文名是 Streaming Multiprocessor，翻译过来就是流式多处理器。SM的核心组件包括CUDA核心，共享内存，寄存器等，SM可以并发地执行数百个线程，并发能力就取决于SM所拥有的资源数。当一个kernel被执行时，它的gird中的线程块被分配到SM上，一个线程块只能在一个SM上被调度。SM一般可以调度多个线程块，这要看SM本身的能力。那么有可能一个kernel的各个线程块被分配多个SM，所以grid只是逻辑层，而SM才是执行的物理层。
SM采用的是SIMT (Single-Instruction, Multiple-Thread，单指令多线程)架构，基本的执行单元是线程束（wraps)，线程束包含32个线程，这些线程同时执行相同的指令，但是每个线程都包含自己的指令地址计数器和寄存器状态，也有自己独立的执行路径。所以尽管线程束中的线程同时从同一程序地址执行，但是可能具有不同的行为，比如遇到了分支结构，一些线程可能进入这个分支，但是另外一些有可能不执行，它们只能死等，因为GPU规定线程束中所有线程在同一周期执行相同的指令，线程束分化会导致性能下降。
当线程块被划分到某个SM上时，它将进一步划分为多个线程束，因为这才是SM的基本执行单元，但是一个SM同时并发的线程束数是有限的。这是因为资源限制，SM要为每个线程块分配共享内存，而也要为每个线程束中的线程分配独立的寄存器。所以SM的配置会影响其所支持的线程块和线程束并发数量。总之，就是网格和线程块只是逻辑划分，一个kernel的所有线程其实在物理层是不一定同时并发的。所以kernel的grid和block的配置不同，性能会出现差异，这点是要特别注意的。
还有，由于SM的基本执行单元是包含32个线程的线程束，所以block大小一般要设置为32的倍数。
CPU和GPU间数据传输
Pinned Host Memory
参考
主机(CPU)数据分配默认是分页的。GPU不能直接访问分页主机内存中的数据。
CUDA驱动先分配(allocate)一个临时的page-locked,or pinned host array，再将主机的数据拷贝到pinned array，再转换到device memory。
我们可以直接将主机数据分配到pinned memory中，减少一次传输。在C++中使用cudaMallocHost()函数。
Unified Memory in CUDA 6
参考自英伟达官方文档
CPU和GPU中的数据必须分开放置在各自的存储中，通过程序进行复制传递。
Unified Memory创建了一个pool of managed memory来共享二者之间的数据。只通过同一个指针，GPU和CPU就都可以访问这块managed memory。
实际上，系统会自动在host和device之间迁移管理managed memory中的数据，使得数据看起来像是在一块可共享的内存中。所以我们只需要进行一次内存分配，就可以得到一个能从host和device访问的指向数据的指针。
库函数
cudaDeviceSynchronize()：用在host code中，让CPU等待直到所有GPU操作结束。cudaMemcpy()：只能用于默认的stream0中，自带syn功能，会阻塞调用它的CPU线程直到拷贝完成。cudaMemcpyAsync()：可以接受一个stream参数，当作用在==非默认stream0且操作内存为pinned memory时，可以在拷贝完成前将控制权交还给调用线程。
Stream
只有不同stream中的CUDA操作才能并发交错执行。
当没有设置stream时使用默认的stream，此时host和device完全同步，就相当于在每个cuda操作后面加了一个cudaDeviceSynchronize()
参考资料：英伟达官方PPT，Stream and Concurrency
A sequence of operations that execute in issue-order on GPU
CUDA operations in different streams may run concurrently
CUDA operations from different streams may be interleaved
并发的条件：
- CUDA操作必须在不同的，非零的stream上
- cudaMemcpyAsync必须作用在
pinned memory上(用cudaMallocHost()或者cudaHostAlloc()分配的内存) - 有足够的资源：
- 不同方向的cudaMemcpyAsyncs
- 设备资源(SMEM, registers, blocks等)
当没有设置stream时使用默认的stream，此时host和device完全同步，就相当于在每个cuda操作后面加了一个cudaDeviceSynchronize()（该函数作用？？？）
