一、CUDA简介
CUDA(Compute Unified Device Architecture)是一种由NVIDIA推出的通用并行计算架构,该架构使GPU(Graphics Processing Unit)能够对复杂的计算问题做性能速度优化。
二、串并行模式
高性能计算的关键是利用多核处理器进行并行计算。
串行模式:将任务分成很多小任务,逐个依次进行。
串并行模式:利用多核处理器同时处理多个子任务(前提是这些子任务不需要相互依赖,比如线程1的计算任务不需要用到线程2的计算结果)。为了加快大任务的计算速度,可以把一些独立的模块分配到不同的处理器上进行同时计算(这就是并行),最后再将这些结果进行整合,完成一次任务计算。
拆分计算模块进行并行加速。强依赖计算采用串行。
CPU需要很强的通用性来处理各种不同的数据类型,比如整型、浮点数等,同时它又必须擅长处理逻辑判断所导致的大量分支跳转和中断处理,而GPU面对的则是类型高度统一的、相互无依赖的大规模数据和不需要被打断的纯净的计算环境。
CPU:擅长流程控制和逻辑处理,不规则数据结构,不可预测存储结构,单线程程序,分支密集型算法。
GPU:擅长数据并行计算,规则数据结构,可预测存储模式。
异构计算:不同体系结构的处理器相互协作完成计算任务。当GPU各个线程完成计算任务后,我们就将GPU计算的结果拷贝到CPU端,完成一次计算任务。
其中蓝色表示串行部分,绿色表示并行部分。
三、GPU参数讲解
线程是程序执行的最基本单元,CUDA的并行计算就是通过成千上万个线程的并行执行来实现的。
CUDA的线程模型从小到大来总结为:
1.Thread:线程,并行的基本单位
2.Thread Block:线程块,互相合作的线程组,线程块有如下几个特点:
(1)允许彼此同步
(2)可以通过共享内存快速交换数据
(3)以1维、2维或者三维组织
3.Grid:一组线程块
(1)以1维、2维组织
(2)共享全局内存
4.Kernel:在GPU上执行的核心程序,这个kernel函数是运行在某个Grid上的。
(1)One kernel 对应 One Grid
(2)每一个block和每一个thread都有自己的ID,我们可以通过相应的索引找到相应的线程块和线程。
a.threadIdx, blockIdx
b.Block ID:1D, 2D or 3D
c.Thread ID:1D, 2D or 3D
型号 | RTX 4090 |
核心代号 | AD102-300 |
架构 | Ada Lovelace |
晶体管数目(亿) | 763 |
核心面积(mm2) | 608.5 |
制程工艺 | TSMC 4N NVIDIA定制工艺 |
GPCs | 11 |
TPCs | 64 |
SMs/CUs | 128 |
Tensor Cores | 512(第四代) |
RT Cores | 128(第三代) |
纹理单元 | 512 |
光栅单元 | 176 |
峰值纹理填充速度 (GT/s) | 1290.2 |
L2缓存(kb) | 73728 |
有效显存速度/显存数据速率(Gbps) | 21 |
显存频率(MHz) | 10501 |
最大显存带宽 (GB/s) | 1008 |
发售时建议零售价 | 12999 |
CUDA Cores/流处理器 | 16384 SP:最基本的处理单元,streaming processor |
Boost频率(MHz) | 2520 |
基础频率(MHz) | 2235 |
显存 | 24GB |
显存位宽(bit) | 384 |
支持 NVIDIA NVLink™ (SLI-Ready) | 否 |
显示器接口 | 3×DisplayPort 1×HDMI |
多显示器 | 4 |
插槽占用 | 3 插槽 |
最高 GPU 温度 (℃) | 90 |
TDP/TGP/TBP (W) | 450 |
推荐电源(W) | 850 |
供电接口 | 16-Pin (含1个4×8-Pin转16针接口适配器) |
PCIe接口 | PCIe 4.0×16 |
四、编程注意事项
在CUDA编程中,核函数(kernel functions)是设计来在GPU上并行执行的函数。当核函数被调用时,它的执行是分配给GPU上的多个CUDA核心(也称为流处理器或SPs)来并行处理的。这包括核函数内的所有操作,包括条件分支语句。
条件分支语句的执行
在CUDA核函数中,条件分支语句(如if
、else
、switch
等)是在GPU上执行的,而不是在CPU上。这意味着每个并行执行的线程(在CUDA中,线程通常组织成线程块(blocks)和网格(grids))都会独立地评估条件分支语句,并根据其自己的局部状态(如线程索引、数据值等)来决定执行哪个分支。
分支发散(Branch Divergence)
虽然条件分支语句是在GPU上执行的,但分支条件的不同结果(即不同线程执行不同的分支)可能导致所谓的“分支发散”(Branch Divergence)。当GPU上的多个线程在同一时间点上需要执行不同的代码路径时,就会发生分支发散。
分支发散会影响GPU的性能,因为GPU硬件通常设计来同时执行大量相同的指令(SIMD,单指令多数据)。当发生分支发散时,GPU必须处理不同的指令流,这可能会降低并行执行的效率。为了最小化分支发散的影响,开发者应当尽量设计算法和数据结构,使得条件分支的预测更加一致,或者通过重新组织代码来减少分支的复杂度。
五、CUDA并行化改造建议
假设算法模型包含了神经网络、控制语句(如if-else
)以及串行迭代计算过程,可以通过如下的分析过程加速你的算法模型:
- 识别可并行化的部分:
- 首先,分析你的算法模型,确定哪些部分是可以并行化的。通常,神经网络的各层计算(如矩阵乘法、激活函数应用等)是高度可并行化的。
- 控制语句(
if-else
)和串行迭代计算过程可能不那么容易并行化,但如果迭代内部的计算可以独立进行,或者可以通过某种方式重新组织以减少依赖关系,则仍然有可能实现部分并行化。
- 设计CUDA核函数:
- 为神经网络的每一层设计一个或多个CUDA核函数。这些核函数将负责在GPU上执行矩阵乘法、激活函数应用等计算任务。
- 如果可能,尝试将控制语句和串行迭代计算过程融入到并行化的框架中。例如,如果迭代内部的计算是独立的,你可以考虑使用不同的线程来处理不同的迭代。
- 优化内存访问:
- 优化CUDA核函数中的内存访问模式,以减少全局内存的访问延迟和带宽限制。这通常涉及到合并内存访问、使用共享内存(如果适用)以及优化数据布局等策略。
- 尽量避免在核函数内部进行不必要的内存分配和释放,因为这会导致额外的开销。
- 处理分支发散:
- 如果你的CUDA核函数中包含
if-else
语句,并且这些语句导致不同的线程执行不同的代码路径,那么可能会出现分支发散。尽量减少分支发散的影响,例如通过重新组织代码、使用位掩码或条件计算等技术来减少分支的复杂度。
- 如果你的CUDA核函数中包含
- 并行和串行部分的协调:
- 如果你的算法模型中既包含并行部分也包含串行部分,你需要仔细设计这些部分的协调机制。这通常涉及到在CPU和GPU之间传输数据、同步线程以及管理计算流程等任务。
- 使用CUDA的流(Streams)和事件(Events)等特性来优化并行和串行部分的执行顺序和效率。
- 性能测试和调优:
- 在GPU上实现你的算法模型后,进行性能测试以评估其加速效果。使用NVIDIA的Nsight Compute、Nsight Systems或CUDA Profiler等工具来分析和优化你的CUDA程序。
- 根据测试结果调整并行策略、优化内存访问模式以及改进算法设计以提高执行效率。
- 性能测试和调优:
- Isaac Gym:NVIDIA的Isaac Gym是一个高效的并行仿真环境,可以在单个GPU上同时运行数千个环境实例。它特别适用于需要大规模并行采样的强化学习任务。(仿真环境的矩阵化改造)
六、一些运行案例说明
python中pytorch用法
PyTorch的Tensor操作和神经网络层是在GPU上通过调用CUDA库来并行执行的,但控制流(如循环、条件判断等)则是由Python解释器在CPU上处理的。这是因为PyTorch(和大多数Python库)的设计是基于高层次的抽象,它们隐藏了底层的CUDA编程细节,以便用户能够以更直观的方式编写代码。
对于一个算法模型,其中存在神经网络、控制语句if else以及串行迭代计算过程。在PyTorch中,使用torch.multiprocessing
来创建多个进程,并将模型和数据移至GPU上时,处理方式会根据不同部分的性质而有所不同:
-
神经网络部分:神经网络的计算(即前向传播、反向传播和参数更新)会发生在指定的GPU上,前提是模型的Tensor和数据都已经被
.to(device)
(其中device
是指向GPU的torch.device
对象)正确地移动到了GPU上。PyTorch会自动管理这些计算,在GPU上执行时利用GPU的并行计算能力。 -
控制语句if else:这些控制语句是由Python的CPU执行的。在PyTorch的上下文中,控制语句用于控制模型的逻辑流程,如条件分支或循环。由于这些操作不涉及到大量的数值计算,因此它们不需要在GPU上执行。相反,Python解释器(在CPU上运行)会处理这些控制语句。
-
串行迭代计算过程:如果这里的“串行迭代计算过程”指的是在Python中编写的循环或迭代操作,那么这些操作也是由CPU执行的。这些迭代过程可能涉及对Tensor的操作,但只要这些Tensor在GPU上,那么涉及Tensor的数值计算(如加法、乘法等)就会被发送到GPU上执行。然而,迭代本身(即循环的迭代逻辑)是在CPU上由Python解释器控制的。
CudaC中
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length)
{
int tid = threadIdx.x;
int a=100;
if(a>1000){
if (tid < length) {
vecres[tid] = vec1[tid] + vec2[tid];
}
}
}
该代码的条件逻辑都是在GPU中执行的。