2. CUDA基础回顾：GPU架构、CUDA编程模型、内存层次结构、线程组织与调度

在深入CUDA图优化与TensorRT动态形状之前，必须夯实CUDA编程的核心基础。GPU并非“更快的CPU”，而是一种大规模并行吞吐量处理器。理解其硬件架构、编程模型、内存层次与线程调度机制，是后续分析图捕获、内核融合与动态形状优化的前提。

2.1 GPU硬件架构：从SM到CUDA Core

现代NVIDIA GPU（如Ampere、Hopper、Ada Lovelace架构）的核心计算单元是流式多处理器（Streaming Multiprocessor, SM）。每个SM包含若干CUDA Core（FP32/INT32单元）、Tensor Core（用于矩阵乘加）、特殊函数单元（SFU）、共享内存（Shared Memory）、寄存器文件（Register File）以及调度器（Warp Scheduler）。

以Ampere架构的GA102为例，一个SM包含128个CUDA Core、4个Tensor Core、4个Warp Scheduler，以及最大228KB可配置的L1/共享内存。GPU由多个GPC（图形处理簇）和TPC（纹理处理簇）组成，每个TPC包含两个SM。硬件层级关系如下：

GPU (例如 A100)
└── GPC (图形处理簇) ×8
    └── TPC (纹理处理簇) ×2
        └── SM (流式多处理器) ×2
            ├── CUDA Core ×128 (FP32)
            ├── Tensor Core ×4
            ├── Warp Scheduler ×4
            ├── Register File (64K × 32-bit)
            └── Shared Memory / L1 Cache (最大228KB)

2.2 CUDA编程模型：网格、线程块与线程

CUDA采用层次化编程模型，将并行计算抽象为三层：网格（Grid）→ 线程块（Block）→ 线程（Thread）。内核函数（Kernel）以网格形式启动，每个网格包含若干线程块，每个线程块包含若干线程。这种层次结构直接映射到硬件：线程块被调度到SM上执行，线程块内的线程通过共享内存协作。

线程和线程块可以使用dim3类型定义多维索引（最多三维），便于处理图像、体素等数据。典型的内核启动语法如下：

// 定义线程块和网格维度
dim3 blockDim(16, 16);      // 线程块：16×16 = 256线程
dim3 gridDim((width + 15) / 16, (height + 15) / 16);

// 启动内核
myKernel<<<gridDim, blockDim>>>(d_input, d_output, width, height);

// 内核内部获取索引
__global__ void myKernel(const float* in, float* out, int w, int h) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < w && y < h) {
        int idx = y * w + x;
        out[idx] = in[idx] * 2.0f;
    }
}

2.3 内存层次结构：从全局内存到寄存器

CUDA内存模型是性能优化的核心。不同内存的容量、延迟、带宽差异巨大，合理使用内存层次是CUDA优化的关键。下表总结了主要内存类型及其特征：

2.4 线程组织与调度机制

GPU的线程调度遵循SIMT（单指令多线程）模型。硬件以Warp（32线程）为单位进行调度。当一个Warp中的线程执行条件分支时，会产生分支发散（Branch Divergence），导致串行执行不同路径，降低效率。

调度流程如下：

1. 网格启动：主机端通过<<<gridDim, blockDim>>>启动内核，GPU工作调度器将线程块分配到空闲SM。
2. Warp划分：每个线程块被划分为多个Warp（线程ID连续32个线程为一组）。例如，256线程的块被分为8个Warp。
3. Warp调度：SM内的Warp Scheduler每个周期选择一个就绪的Warp发射指令。当Warp等待内存访问时，调度器切换到其他Warp，隐藏延迟。
4. 执行：Warp中的32个线程在同一周期执行相同指令，但操作不同的数据。

// 线程索引计算示例
__global__ void exampleKernel(float* data, int pitch) {
    // 线性索引
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 2D索引
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int idx_2d = y * pitch + x;
    
    // Warp内线程ID
    int lane = threadIdx.x % 32;
    // Warp ID（块内）
    int warp_id = threadIdx.x / 32;
}

2.5 案例：矩阵乘法的CUDA实现与优化思路

矩阵乘法（GEMM）是深度学习中最核心的运算。以C = A × B为例，朴素实现和优化实现差异巨大。

朴素实现（直接映射）：

__global__ void matmul_naive(const float* A, const float* B, float* C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    float sum = 0.0f;
    for (int k = 0; k < N; ++k) {
        sum += A[row * N + k] * B[k * N + col];
    }
    C[row * N + col] = sum;
}

该实现每个线程计算一个输出元素，但每次循环都访问全局内存，内存带宽成为瓶颈。优化思路包括：

• 共享内存分块（Tiling）：将A和B的块加载到共享内存，减少全局内存访问次数。
• 寄存器缓存：每个线程计算多个输出元素，复用寄存器中的数据。
• 避免Bank Conflict：调整共享内存访问模式，使同一Warp的线程访问不同Bank。
• 向量化加载：使用float4一次加载4个浮点数，提高内存吞吐。

2.6 常见误区与复盘

if (threadIdx.x % 2 == 0) {
    // 偶数线程执行
} else {
    // 奇数线程执行
}