CUDA并行架构

一、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 SM：多个SP加上其他的一些资源组成一个streaming multiprocessor，也叫GPU大核
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语句，并且这些语句导致不同的线程执行不同的代码路径，那么可能会出现分支发散。尽量减少分支发散的影响，例如通过重新组织代码、使用位掩码或条件计算等技术来减少分支的复杂度。
并行和串行部分的协调：
- 如果你的算法模型中既包含并行部分也包含串行部分，你需要仔细设计这些部分的协调机制。这通常涉及到在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中执行的。