GPU架构基础：GPU并行计算模型、内存层次结构与线程调度机制

在开始使用Nsight进行性能分析之前，必须先建立对GPU底层执行模型的基本认知。CPU与GPU的设计哲学截然不同：CPU追求单线程低延迟，依赖大容量缓存和复杂控制逻辑；GPU则追求大规模并行吞吐量，将更多晶体管用于计算单元。理解这种差异，是后续解读Nsight各项计数器指标的前提。

2.1 GPU并行计算模型：从线程到网格

GPU的并行执行模型是层级化的。最底层的执行单元是线程（Thread），多个线程被组织为线程块（Block），多个线程块再构成网格（Grid）。这种层级结构直接映射到CUDA编程模型，也决定了Nsight中看到的线程调度数据。

• 线程（Thread）：最小的执行单位，每个线程执行相同的核函数（Kernel）代码，但通过内置变量（如 threadIdx）区分处理的数据。
• 线程块（Block）：由数十至上千个线程组成，块内线程可通过共享内存（Shared Memory）协作，并支持同步操作（__syncthreads()）。
• 网格（Grid）：由多个线程块组成，块之间无法直接通信，只能通过全局内存（Global Memory）间接交互。

实际硬件执行时，线程块会被进一步拆分为线程束（Warp）——这是NVIDIA GPU调度的最小硬件单位，通常包含32个线程。一个线程块被分配给某个流多处理器（SM）后，其中的线程会按32个一组被划分为多个Warp，由SM的调度器分发执行。

2.2 内存层次结构：带宽与延迟的权衡

GPU内存层次结构是性能分析的核心。Nsight中大量的性能计数器（如 l1tex__throughput、dram__bytes_read）都对应着不同层级的内存访问。理解各层级的容量、带宽和延迟，才能准确判断瓶颈是“计算密集”还是“访存密集”。

2.3 线程调度机制：Warp的诞生与切换

GPU的线程调度是零开销切换的。当某个Warp因等待内存数据而阻塞时，SM的调度器会立即切换到另一个就绪的Warp执行，这个切换过程不需要保存/恢复寄存器，因为每个Warp的寄存器是物理独立的。

调度机制的核心要点如下：

1. Warp划分：线程块被分配到SM后，立即按连续32个线程划分为Warp。例如，一个256线程的Block会被分为8个Warp。
2. 调度周期：每个时钟周期，SM的调度器从所有就绪Warp中选择一个发出指令。现代GPU（如Ampere架构）每个SM有4个调度器，每个调度器每周期可发射1条指令。
3. 延迟隐藏：GPU通过大量Warp并行来隐藏内存延迟。假设一次全局内存访问需要500周期，如果SM中有16个活跃Warp，当1个Warp等待时，其余15个Warp可以继续执行，从而让计算单元保持忙碌。
4. 分支发散：如果一个Warp内的32个线程执行了不同的分支路径（如 if (tid % 2 == 0)），该Warp会串行执行两个分支，导致部分线程闲置。这是Nsight中“Branch Divergence”指标关注的问题。

Occupancy = 活跃Warp数 / SM最大支持Warp数

2.4 案例：矩阵乘法的内存访问模式分析

以经典的矩阵乘法 C = A × B 为例，观察不同内存访问模式对性能的影响。假设矩阵大小为1024×1024，每个线程计算C中的一个元素。

朴素实现（直接访问全局内存）：

__global__ void matmul_naive(float* A, 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;
}

这种实现中，每个线程需要读取A的一行和B的一列，共2N次全局内存访问。对于1024×1024矩阵，每个线程执行1024次循环，每次循环访问两个全局内存地址。由于相邻线程（同一Warp内）访问的B列地址是连续的（B[k*N + col]，col连续），因此B的访问可以被合并（Coalesced）为高效传输；但A的访问（A[row*N + k]，row相同则地址相同）会导致Warp内所有线程读取同一地址，造成广播，效率尚可。

共享内存优化（分块Tiling）：

__global__ void matmul_tiled(float* A, float* B, float* C, int N) {
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
    float sum = 0.0f;
    for (int tile = 0; tile < N / BLOCK_SIZE; tile++) {
        As[threadIdx.y][threadIdx.x] = A[row * N + tile * BLOCK_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(tile * BLOCK_SIZE + threadIdx.y) * N + col];
        __syncthreads();
        for (int k = 0; k < BLOCK_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

分块后，每个线程块先从全局内存加载一个BLOCK_SIZE×BLOCK_SIZE的子矩阵到共享内存，然后从共享内存读取数据进行计算。这样，全局内存访问次数从2N次降低到2N/BLOCK_SIZE次。例如BLOCK_SIZE=16时，全局内存访问减少16倍。

2.5 误区与复盘：为什么我的GPU利用率很低？

在使用Nsight分析性能时，经常遇到GPU利用率低的情况。以下是几个典型原因及其在Nsight中的表现：