cuda

前言

CUDA 编程需要理解线程与硬件的物理映射，并优化显存（Global Memory）访问以提升性能。

---

一、 核心映射

理解 Grid/Block/Thread 在物理资源上的占用情况。

软件层级 (Software)	硬件真身 (Hardware)	说明
Grid	Device (GPU)	对应显存 (Global Memory)。所有线程可见，但访问延迟较高。
Block	SM (流多处理器)	Block 被分配给 SM 后会驻留直到结束。Block 间无法通信，Block 内可通过 Shared Memory 通信。
Thread	CUDA Core (SP)	最小执行单位。拥有私有的 Registers (极速访问)。

Warp (线程束)

Warp 是 GPU 执行的基本单位。

• 执行机制：GPU 采用 SIMT (Single Instruction, Multiple Threads) 架构，32 个线程 (Warp) 同时执行相同指令。
• **分支发散 (Warp Divergence)**：如果 if-else 分支导致 Warp 内线程执行路径不同，硬件将串行化执行各分支，降低并行效率。

---

二、 从 1D 到 Memory Coalescing

1. 1D 索引计算

公式 i = blockIdx.x * blockDim.x + threadIdx.x 用于将并行线程映射到线性内存地址。

• Thread ID：作为指针算术 (Pointer Arithmetic) 的偏移量。
• Thread 0 -> BaseAddr + 0
• Thread 1 -> BaseAddr + 4 bytes

2. 内存合并 (Memory Coalescing)

当一个 Warp 发起内存请求时，硬件会尝试合并访问：

• **合并访问 (Coalesced)**：线程访问连续地址 (k, k+1, k+2...)。硬件可将 32 个请求合并为一个 128-byte 的事务。
• **非连续访问 (Strided/Random)**：地址跳跃或乱序。硬件需发射多个独立事务，导致带宽浪费。

结论：应确保相邻 threadIdx 访问相邻内存地址。

---

三、 Grid-Stride Loop

Grid-Stride Loop 模式用于处理数据量超过线程总数的情况，并提高代码复用性。

__global__ void vectorAdd(const float *A, const float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x; 

    while (i < N) {
        C[i] = A[i] + B[i];
        i += stride; 
    }
}

• 解耦：Kernel 执行不受 Grid 大小限制。
• 效率：减少线程创建与销毁的开销。

---

四、 2D 矩阵与 Flattening

1. 坐标映射 (Mapping)

利用 CUDA 的 2D 索引计算坐标：

• **Col (列 / x)**：blockIdx.x * blockDim.x + threadIdx.x
• **Row (行 / y)**：blockIdx.y * blockDim.y + threadIdx.y

2. 地址压扁 (Flattening)

C 语言矩阵通常采用行主序 (Row-Major) 存储。

3. Naive GEMM 实现

以下是基础的矩阵乘法实现。

__global__ void matrixMulNaive(const float *A, const float *B, float *C, int N) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

    if (row < N && col < N) {
        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;
    }
}

4. 性能分析

• **内存受限 (Memory Bound)**：计算一个元素需要多次访问 Global Memory。
• 延迟：Global Memory 访问延迟较高，导致计算单元等待数据。

---

五、 Roofline 模型

Roofline 模型用于分析应用程序在特定硬件上的性能瓶颈。

1. 核心指标

• 算力峰值 ()：硬件每秒能完成的最大浮点运算次数 (FLOPS)。
• 带宽峰值 ()：硬件每秒能完成的最大内存交换量 (Bytes/s)。
• 计算强度 ()：也称为算术强度 (Arithmetic Intensity)，指每字节内存交换所完成的浮点运算次数。
  

2. 性能模型

应用程序的可达性能 受限于算力和带宽：

3. 瓶颈分析

• 带宽受限 (Memory Bound)：当 时，。此时性能受限于内存带宽，优化方向为减少内存访问或提高内存访问效率（如 Coalescing）。
• 计算受限 (Compute Bound)：当 时，。此时性能受限于硬件算力，优化方向为提高计算并行度或使用更高效的指令。