CUDA与TensorRT(1)之CUDA-C与GPU基础

1. CPU与GPU

1. CPU–延迟导向设计
   内存大：多级缓存结构提高访存速度
   控制复杂：分之预测机制和流水线数据传送
   运算单元强大：整型浮点型复杂运算速度快
2. GPU–吞吐导向设计
   缓存小：提高内存吞吐
   控制简单：没有分支预测和数据前传
   精简运算单元：多 长延时流水线以实现高吞吐量 和 小大量线程来容忍延迟
3. 各自特点
   CPU延迟优先，单条复杂指令比GPU快10倍以上
   GPU吞入优先，单位时间内执行指令数量10倍以上
   显存和内存的联系：都是用来暂存资料的存储空间，只是显存帮GPU，内存帮CPU

2. CUDA和OpenCL

CUDA（Compute Unified Device Architecture)为GPU增加一个易用的编程接口；OpenCL与CUDA作用一致，只是它支持的平台更多，除了GPU外还支持CPU、DSP、FPGA等

a. CUDA编程并行计算整体流程

• Device：GPU
• Host：CPU
• Kernel：GPU上运行的函数

void GPUkernel(float* A, float* B, float* C, int n)
{
	1. // Allocate device memory for A, B, and C
	// copy A and B to device memory 
	2. // Kernel launch code – to have the device
	// to perform the actual vector addition
	3. // copy C from the device memory
	// Free device vectors
}

b. CUDA中的内存模型

硬件上（从内到外）：

• 每个线程处理器（SP）都有自己的寄存器和局部内存，并且只能自己访问。
• 每个多核处理器（SM）内都有自己的共享内存，可以被线程块内所有线程访问。
• 一个GPU的所有SM共有一块全局内存，不用线程块的线程都可使用。

上面的术语对应到软件上为：

• SP–线程，
• SM–线程块，
• 设备端（device）–线程块组合体（grid）

并且一个kernel函数只能由一个grid来执行，且一次只能在一个GPU上执行。
块内的线程通过共享内存、原子操作和屏障同步进行协作，并且不同线程块的线程不能协作；

c.线程块id&线程id

一个Kernel函数运行在一个grid上，当我们编程的时候，要指明哪个线程执行哪些，并且在运行的时候也要去调用固定线程的寄存器和内存，于是就需要地址来对每个线程和线程块定唯一。

dim3 dimGrid(M, N); 
dim3 dimBlock(P, Q, S);
threadId.x = blockIdx.x*blockDim.x+ threadIdx.x; 
threadId.y = blockIdx.y*blockDim.y+ threadIdx.y;

d.线程束（warp）

• SM采用的SIMT(Single-Instruction, Multiple-Thread，单指令多线程）架构，warp（线程束）是最基本的执行单元，一个warp包含32个并行thread，这些thread以不同数据资源执行相同的指令。warp本质上是线程在GPU上运行的最小单元。
• 当一个Kernel被执行时，grid中的线程块被分配到SM上，一个线程块的thread只能在一个SM上调度，SM一般可以调度多个线程块，大量的thread的可能被分到不同的SM上。每个thread拥有它自己的程序计数器和状态寄存器，并且用该线程自己的数据执行指令，这就是所谓的SIMT。
• 由于warp的大小为32，所以block所含的thread大小一般都设置为32的倍数。

3. 并行计算实例：向量相加

主要要做的事：
设备端代码（计算作用）：

• 读写线程的寄存器
• 读写Grid中全局内存
• 读写block中共享内存

主机端代码：

• 申请内存和显存的空间，进行数据的转移和拷贝，最后得到结果到主机端后，进行内存的释放。

a. 使用到的基本函数介绍

#include “cuda.h”

• cudaMalloc()：cudaError_t cudaMalloc (void **devPtr, size_t size)，申请显存函数，其中参数一是申请地址，参数二是申请内存的大小，常用sizeof(float)×N
• cudaFree()：cudaError_t cudaFree ( void* devPtr )，释放显存，参数是指向对象的指针。
• cudaMemcpy()：cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, cudaMemcpyKind kind)
  • 内存数据复制传递，参数一是目的地，参数二是出发地，参数三是大小（同样使用sizeof()×形式）
  • 参数三是拷贝方向：目前支持四种，cudaMemcpyHostToDevice，cudaMemcpyDeviceToHost，cudaMemcpyDeviceToDevice，cudaMemcpyDefault

b.核函数简单介绍及调用

• 在GPU上执行的函数。
• 一般通过标识符__global__修饰。
• 调用通过<<<参数1,参数2>>>，用于说明内核函数中的线程数量，以及线程是如何组织的。
• 以网格（Grid）的形式组织，每个线程格由若干个线程块（block）组成，而每个线程块又由若干个线程（thread）组成。
• 调用时必须声明内核函数的执行参数。
• 在编程时，必须先为kernel函数中用到的数组或变量分配好足够的空间，再调用kernel函数，否则在GPU计算时会发生错误（segmentation fault）。

__ global __ 标志核函数
• 核函数必须返回 void
__ device __ & __ host __ 可以一起用

c. CUDA编译流程

有三种