$\text{GPU}$架构与$\text{CUDA}$编程模型

$\text{1. }\text{GPU}$体系结构

1️⃣计算单元组织架构

结构	功能
$\text{CUDA}$核心	类似$\text{ALU}$(但远没$\text{CPU}$的灵活)，可执行浮点运算/张量运算/光线追踪(高级核心)
$\text{Warp}$	多核心共用一个取指/译码器，按$\text{SIMT}$工作(所有线程指令相同/数据可不同)
$\text{SM}$	包含多组$\text{Warps}$，所有$\text{CUDA}$核心共用一套执行上下文(缓存)$\&$共享内存

2️⃣存储层次架构：

1. 不同$\text{SM}$能够$\text{Access}$相同的$\text{L2 Cache}$
2. 显存与缓存之间的带宽极高，但是相比$\text{GPU}$的运算能力仍然有瓶颈

$\text{2. }$$\text{CUDA}$编程模型

1️⃣$\text{CUDA}$程序简述

1. $\text{CUDA}$程序的两部分

   程序	运行位置	主要职责
   Host程序	$\text{CPU}$	任务管理/数据传输/启动$\text{GPU}$内核
   Device程序	$\text{GPU}$	执行内核/处理数据

2. $\text{Kernel}$即在$\text{GPU}$上运行的函数，如下简单内核定义示例

   //通过__global__关键字声名内核函数
   __global__ void VecAdd(float* A, float* B, float* C)
   {int i = threadIdx.x;C[i] = A[i] + B[i];
   }
   int main()
   {//通过<<<...>>>中参数指定执行kernel的CUDA thread数量VecAdd<<<1, N>>>(A, B, C); 
   }
   

2️⃣线程并行执行架构

1. 线程层次：

   结构	地位	功能
   $\text{Thread}$	并行执行最小单元	执行$\text{Kernel}$的一段代码
   $\text{Warp(32Threads)}$	线程调度的基本单位	所有线程以$\text{SIMD}$方式执行相同指令
   $\text{Block}$	$\text{GPU}$执行线程基本单位	使块内线程内存共享/指令同步
   $\text{Grid}$	并行执行的最大单元	执行整个内核(启动内核时必启动整个$\text{Grid}$)

2. 线程在计算单元的映射：线程层次$\stackrel{层次对应}{\leftrightarrow }\text{GPU}$物理架构
   
   • 注意$\text{SM}$和$\text{Block}$不必$\text{1v1}$对应也可$\text{Nv1}$对应
3. 线程在存储单元的映射

   线程结构	可$\text{Access}$的内存结构	访问速度
   $\text{Thread}$	每线程唯一的$\text{Local Memory}$	极快
   $\text{Block}$	每块唯一的$\text{Shared Memory}$(块中每个线程都可访问)	较快
   所有线程	唯一且共享的$\text{Global Memory}$	较慢

$\text{3. CPU}$与$\text{GPU}$

1️⃣$\text{CPU/}\text{GPU}$结构对比

	$\text{GPU}$	$\text{CPU}$
$\text{ALU}$	功能强但数量少(只占$\text{GPU}$小部)，时钟频率极高	功能弱但数量大，时钟频率低
$\text{Cache}$	容量大并分级，缓存后续访问数据	容量很小，用于提高线程服务
控制	复杂串行逻辑，如流水/分支预测/乱序执行	简单(但大规模)并行逻辑

3️⃣$\text{CPU}\underset{数据/指令传输}{\overset{\text{PCIe}}{\leftrightarrow }}\text{GPU}$交互

设备	逻辑地位	$\text{IO}$模块	任务分配
$\text{GPU}$	外设	$\text{IO Block}$(南桥)	控制逻辑和任务调度
$\text{CPU}$	主机	$\text{Copy Engine}$	执行大量并行计算任务