📅 创建时间：2026-06-03 🏷️ 标签：#CUDA #GPU #Warp #Block #SharedMemory #BankConflict #Coalescing #PTX #SASS #TensorCore 📚 前置知识：[[13-codegen-architecture]]（代码生成架构） 📚 相关知识：[[08-operator-fusion]]（算子融合） [[17-kernel-primer]]（Kernel 开发入门）

CUDA 后端：合并访存与 Tensor Core / CUDA Backends, Memory Coalescing, and Tensor Cores

┌──────────────────────────────────────────────────────────────────────────────┐ │ 情境描述 │ ├──────────────────────────────────────────────────────────────────────────────┤ │ 你的 GEMM kernel 比 cuBLAS 慢 3 倍。用 Nsight Compute 分析后发现： │ │ - Shared Memory 利用率只有 30%（你配了 48KB，只用了 15KB） │ │ - 大量 bank conflict 导致 serialization │ │ - Global memory access 没有 coalescing，线程访问散列地址 │ │ │ │ 你的 kernel 代码看起来没问题，但硬件层面的行为一团糟。 │ └──────────────────────────────────────────────────────────────────────────────┘

第1节：GPU 执行模型——SIMT 架构详解

1.1 GPU vs CPU：设计哲学的差异

CPU 和 GPU 有根本不同的设计哲学：

┌─────────────────────────────────────────────────────────────────────────────┐
│                      CPU vs GPU 架构对比                                    │
├─────────────────────────────────────────────────────────────────────────────┤
│                                                                             │
│  CPU (延迟优化):                                                            │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │  ┌─────┐ ┌─────┐ ┌─────┐ ┌─────┐   Core 0-3: 复杂，深度流水线         │   │
│  │  │ ALU │ │ FPU │ │ LSU │ │Branch│  每个核都很"聪明"                     │   │
│  │  └──┬──┘ └──┬──┘ └──┬──┘ └──┬──┘                                      │   │
│  │     └────────┬────────┘                                               │   │
│  │        ┌──────┴──────┐                                                  │   │
│  │        │ Large Cache │  ← 大 Cache 减少内存访问                        │   │
│  │        └─────────────┘                                                  │   │
│  │                    ↓                                                    │   │
│  │        ┌─────────────────┐                                              │   │
│  │        │   DDR4/DDR5     │                                              │   │
│  │        │   Memory        │                                              │   │
│  │        └─────────────────┘                                              │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│                                                                             │
│  GPU (吞吐量优化):                                                          │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │  ┌─────────────────────────────────────────────────────────────┐    │   │
│  │  │  SM 0      SM 1      SM 2      SM 3    ...  SM 79            │    │   │
│  │  │  ┌───┐    ┌───┐    ┌───┐    ┌───┐         ┌───┐             │    │   │
│  │  │  │SP │    │SP │    │SP │    │SP │  × 4   │SP │             │    │   │
│  │  │  │×8 │    │×8 │    │×8 │    │×8 │        │×8 │             │    │   │
│  │  │  └───┘    └───┘    └───┘    └───┘        └───┘             │    │   │
│  │  │  大规模并行：80 SM × 64 threads/SM = 5120 线程并行           │    │   │
│  │  └─────────────────────────────────────────────────────────────┘    │   │
│  │                              ↓                                        │   │
│  │  ┌─────────────────────────────────────────────────────────────┐    │   │
│  │  │                    HBM2 / HBM3 Memory                       │    │   │
│  │  │              带宽：~2 TB/s（相比 CPU 的 ~100 GB/s）         │    │   │
│  │  └─────────────────────────────────────────────────────────────┘    │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│                                                                             │
└─────────────────────────────────────────────────────────────────────────────┘

1.2 CUDA 编程模型层级

python

class CUDAHierarchy:
    """
    CUDA 层级结构
    
    Grid → Block → Warp → Thread
    
    CUDA 设备内存层级:
    - Global Memory: 所有线程可访问，容量大，延迟高
    - Shared Memory: Block 内线程共享，容量小，延迟低
    - Register: 线程私有，速度最快
    - Local Memory: 线程私有，溢出到 Global Memory
    """
    
    # Grid 级别
    # 整个 kernel 启动的全局范围
    # gridDim: Grid 维度 (最多 3D)
    grid_dim = (num_blocks_x, num_blocks_y, num_blocks_z)
    
    # Block 级别
    # 一组线程，共享 Shared Memory
    # blockIdx: Block 在 Grid 中的索引
    # blockDim: Block 的维度 (最多 3D)
    block_dim = (threads_x, threads_y, threads_z)
    max_threads_per_block = 1024
    
    # Warp 级别
    # 32 个线程为一组，同步执行 (SIMT)
    # warpSize = 32
    # 同一 Warp 内的线程执行同一条指令
    
    # Thread 级别
    # 最基本的执行单元
    # threadIdx: 线程在 Block 中的索引
    # 每个线程有独立的寄存器和 Local Memory

1.3 SM (Streaming Multiprocessor) 内部结构

┌─────────────────────────────────────────────────────────────────────────────┐
│                    NVIDIA SM (Streaming Multiprocessor)                      │
│                           以 Volta/Ampere 为例                                │
├─────────────────────────────────────────────────────────────────────────────┤
│                                                                             │
│                              Warp Scheduler × 4                              │
│                              ↓         ↓         ↓         ↓                │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │                     Dispatch Unit × 4                                │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│                              ↓                                               │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │  ┌─────────┐  ┌─────────┐  ┌─────────┐  ┌─────────┐                │   │
│  │  │ FP32    │  │ FP32    │  │ INT32   │  │ FP64    │  ← 执行单元    │   │
│  │  │ CUDA    │  │ CUDA    │  │ ALU     │  │ FMA     │                │   │
│  │  │ Core ×32│  │ Core ×32│  │ ×32     │  │ ×16     │                │   │
│  │  └─────────┘  └─────────┘  └─────────┘  └─────────┘                │   │
│  │                                                                      │   │
│  │  ┌─────────┐  ┌─────────┐  ┌─────────┐  ┌─────────────────────┐    │   │
│  │  │ LD/ST   │  │ LD/ST   │  │Tensor   │  │ Special Function    │    │   │
│  │  │ Unit ×16│  │ Unit ×16│  │Core×4   │  │ Unit ×2            │    │   │
│  │  │         │  │         │  │(WMMA)   │  │ (Transcendental)    │    │   │
│  │  └─────────┘  └─────────┘  └─────────┘  └─────────────────────┘    │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│                              ↓                                               │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │                        Register File: 65536 × 32-bit               │   │
│  │                           (2 MB per SM)                              │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│                              ↓                                               │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │                   L1 Data Cache / Shared Memory                    │   │
│  │                    128 KB (可配置比例)                               │   │
│  │        Shared Memory: 48 KB / 64 KB / 96 KB / 100 KB              │   │
│  │        L1 Cache: 32 KB / 16 KB / ...                               │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│                              ↓                                               │
│                           L2 Cache (所有 SM 共享)                             │
│                     A100: 40 MB, H100: 50 MB                               │
│                                                                             │
└─────────────────────────────────────────────────────────────────────────────┘

第2节：GPU 内存层级——性能的关键

2.1 内存层级与延迟

内存类型	容量	延迟	带宽	作用域
Register	~256 KB/SM	1 cycle	~16 TB/s	线程私有
Local Memory	~1 MB/SM	~400 cycles	~2 TB/s	线程私有
Shared Memory	48-100 KB/SM	~1 cycle	~2 TB/s	Block 内
L1 Cache	32-128 KB/SM	~30 cycles	~2 TB/s	SM 内
L2 Cache	20-50 MB (GPU)	~200 cycles	~2 TB/s	GPU 全局
Global Memory (HBM)	16-80 GB	~400 cycles	~2 TB/s	全局

2.2 Shared Memory 配置

cuda

// Shared Memory 配置
// 每个 SM 有 128 KB 的 L1/Shared Memory 空间
// 可以配置为不同的比例

// 方式1: kernel 级别配置
__global__ void kernel_with_shared(float* data) {
    // 动态分配 Shared Memory
    extern __shared__ float shared_data[];
    
    // 静态分配 Shared Memory
    __shared__ float static_shared[1024];
    
    // 使用 shared memory
    // ...
}

// 启动时配置: cudaKernelLaunch(kernel, blocks, threads, shared_size)
// shared_size = 48 * 1024; // 48 KB


// 方式2: 使用 cudaFuncSetAttribute
cudaFuncSetAttribute(kernel_with_shared, 
                     cudaFuncAttributePreferredSharedMemoryCarveout,
                     cudaSharedmemConfigDynamicSegmentSize);  // 更多 Shared Memory


// 方式3: 编译时配置
// --dparm=sharedMemPerBlock=49152  (48 KB)


// 推荐配置场景
// GEMM/Conv: 48 KB Shared Memory (更少的 L1 Cache)
// 普通计算: 32 KB Shared Memory / 32 KB L1 Cache

2.3 Global Memory 访问模式

cuda

// Coalesced vs Non-coalesced 访问

// ========== 正确：Coalesced 访问 ==========
// 每个线程访问连续地址，合并为少量 memory transaction
__global__ void coalesced_write(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (tid < n) {
        // 线程 tid 写入 data[tid]
        // 连续线程访问连续地址 ✓
        data[tid] = (float)tid;
    }
}

// 访问模式示意:
// Thread 0: data[0], data[16], data[32]  (stride = blockDim)
// Thread 1: data[1], data[17], data[33]
// ...
// 每个 Warp (32 threads) 访问连续 128 bytes → 合并为 1-2 次 transaction


// ========== 错误：Non-coalesced 访问 ==========
__global__ void non_coalesced_write(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int offset = threadIdx.x * blockDim.x;  // 错误！
    
    if (offset < n) {
        // 线程 tid 写入 data[offset]
        // 连续线程访问不相邻的地址 ✗
        // Thread 0: data[0], Thread 1: data[16], Thread 2: data[32]...
        // 地址不连续，无法合并
        data[offset] = (float)tid;
    }
}


// ========== 另一个常见错误：跨行访问 ==========
__global__ void wrong_matrix_access(float* A, float* B, int N) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (row < N && col < N) {
        // A 按行存储
        // A[row, col] 在内存中连续
        float a = A[row * N + col];  // ✓ Coalesced
        
        // B 按行存储，但我们按列访问
        float b = B[col * N + row];  // ✗ Non-coalesced!
        // 当 row != col 时，访问地址跳跃很大
    }
}

第3节：Bank Conflict——Shared Memory 的陷阱

3.1 Shared Memory Bank 架构

┌─────────────────────────────────────────────────────────────────────────────┐
│                    Shared Memory Bank Conflict                               │
├─────────────────────────────────────────────────────────────────────────────┤
│                                                                             │
│  Shared Memory 被分成多个 Bank（类似银行）                                    │
│  每个 Bank 每周期可以处理一次访问                                            │
│                                                                             │
│  Bank 数量 (每 SM):                                                         │
│  - Kepler/ Maxwell/ Pascal: 32 banks (每 bank 32-bit)                      │
│  - Volta/ Turing/ Ampere/Ada/Hopper: 32 banks (每 bank 32-bit)              │
│                                                                             │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │  Bank 0: addr 0, 32, 64, 96, ...                                   │   │
│  │  Bank 1: addr 1, 33, 65, 97, ...                                   │   │
│  │  Bank 2: addr 2, 34, 66, 98, ...                                   │   │
│  │  ...                                                               │   │
│  │  Bank 31: addr 31, 63, 95, 127, ...                               │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│                                                                             │
│  冲突情况:                                                                  │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │  No Conflict: 不同 bank → 可以并行访问 ✓                            │   │
│  │  Bank Conflict: 同一 warp 访问同一 bank → 串行化 ✗                  │   │
│  │  Broadcast: 同一 bank + 同一地址 → 广播 ✓                           │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│                                                                             │
└─────────────────────────────────────────────────────────────────────────────┘

3.2 Bank Conflict 示例

cuda

// ========== Bank Conflict 示例 ==========

// 问题代码：所有线程访问同一 bank 的不同地址
__global__ void bank_conflict_example(float* data) {
    __shared__ float shared[32];  // 32 个 float，刚好覆盖所有 bank
    
    int tid = threadIdx.x;
    shared[tid] = data[tid];  // 每个线程访问 shared[tid]
    // tid=0 → bank 0, tid=1 → bank 1, ...
    // 线程 0 和线程 32 都访问 bank 0 → 2-way conflict
    // ...
    
    __syncthreads();
    
    // 读取：再次访问同一 bank
    data[tid] = shared[tid];
}


// ========== 解决：Padding ==========
__global__ void no_bank_conflict(float* data) {
    // 在每行末尾添加 padding
    // 32 threads + 1 padding float = 33 floats per row
    __shared__ float shared[32 * 33];  // Padding
    
    int tid = threadIdx.x;
    int row = blockIdx.x;
    int col = tid;
    
    // 每个线程访问 shared[row * 33 + col]
    // 现在每个线程访问不同的 bank（因为列索引不同）
    shared[row * 33 + col] = data[row * 32 + col];
    
    __syncthreads();
    
    data[row * 32 + col] = shared[row * 33 + col];
}


// ========== 矩阵乘法中的 Bank Conflict ==========
__global__ void gemm_with_conflict(float* C, const float* A, const float* B, int N) {
    __shared__ float As[32][32];  // 32x32 tile
    __shared__ float Bs[32][32];
    
    int row = blockIdx.y * 32 + threadIdx.y;
    int col = blockIdx.x * 32 + threadIdx.x;
    
    float sum = 0.0f;
    
    for (int k = 0; k < N; k += 32) {
        // 加载 A: thread (tx, ty) 访问 As[ty][tx]
        // 32 个线程按行访问
        As[threadIdx.y][threadIdx.x] = A[row * N + (k + threadIdx.x)];
        
        // 加载 B: thread (tx, ty) 访问 Bs[ty][tx]
        // 32 个线程按行访问
        // 但 B 在 global memory 是按列存储的
        Bs[threadIdx.y][threadIdx.x] = B[(k + threadIdx.y) * N + col];
        
        __syncthreads();
        
        // 计算
        for (int i = 0; i < 32; i++) {
            sum += As[threadIdx.y][i] * Bs[i][threadIdx.x];
        }
        
        __syncthreads();
    }
    
    C[row * N + col] = sum;
}


// ========== 解决：交换线程索引 ==========
__global__ void gemm_no_conflict(float* C, const float* A, const float* B, int N) {
    __shared__ float As[32][33];  // 添加 padding
    __shared__ float Bs[32][33];
    
    int row = blockIdx.y * 32 + threadIdx.y;
    int col = blockIdx.x * 32 + threadIdx.x;
    
    float sum = 0.0f;
    
    for (int k = 0; k < N; k += 32) {
        // 加载 A: thread (tx, ty) 访问 As[tx][ty]
        // 这样每个线程访问不同列的同一行
        As[threadIdx.x][threadIdx.y] = A[row * N + (k + threadIdx.x)];
        
        // 加载 B: thread (tx, ty) 访问 Bs[tx][ty]
        // 同样交换索引
        Bs[threadIdx.x][threadIdx.y] = B[(k + threadIdx.y) * N + col];
        
        __syncthreads();
        
        for (int i = 0; i < 32; i++) {
            sum += As[i][threadIdx.y] * Bs[threadIdx.x][i];
        }
        
        __syncthreads();
    }
    
    C[row * N + col] = sum;
}

第4节：Tensor Core——矩阵计算加速器

4.1 Tensor Core 架构

python

class TensorCoreEvolution:
    """
    Tensor Core 演进
    """
    
    # ========== Volta (V100) ==========
    volta = {
        'shape': '16×16×16',
        'precision': 'FP16',
        'mma_instructions': 'HMMA (Half-precision Matrix Multiply-Accumulate)',
        'throughput': '125 TFLOPS (FP16)',
        'note': '首次引入，每个 SM 2 个 Tensor Core'
    }
    
    # ========== Turing (T4, RTX 20xx) ==========
    turing = {
        'shape': '16×16×16',
        'precision': 'FP16, INT8, INT4, INT1',
        'new_precision': 'FP16 + FP32 accumulate',  # 累加到 FP32
        'throughput': '65-130 TFLOPS (FP16)'
    }
    
    # ========== Ampere (A100, RTX 30xx) ==========
    ampere = {
        'shape': '16×16×16',
        'precision': 'FP16, BF16, TF32, FP64, INT8, INT4, INT1',
        'new_precision': 'TF32 (19-bit), BF16',
        'warp_specialization': True,  # 新功能
        'throughput': '312 TFLOPS (FP16), 156 TFLOPS (BF16)'
    }
    
    # ========== Hopper (H100) ==========
    hopper = {
        'shape': '16×16×16',
        'new_feature': 'WGMMA (Warp Group MMA)',
        'precision': 'FP8 (新), FP16, BF16, FP64, INT8',
        'async_copy': True,  # 异步拷贝
        'throughput': '495 TFLOPS (FP16), 989 TFLOPS (FP8)'
    }

4.2 WMMA API 使用

cuda

// 使用 WMMA (Warp-level Matrix Multiply-Accumulate) API
#include <mma.h>

// 定义矩阵布局
using namespace nvcuda::wmma;

// GEMM 使用 Tensor Core
__global__ void wmma_gemm(half* C, const half* A, const half* B, int M, int N, int K) {
    
    // Warp 级别的坐标
    int warpM = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
    int warpN = blockIdx.y;
    
    // 每个 warp 处理 16×16 的输出块（Tensor Core 的 tile 大小）
    // 输出矩阵 C 的 tile 形状
    layout_t layoutC = wmma::mem_row_major;
    
    // 定义 fragment（Tensor Core 操作的数据格式）
    fragment<matrix_a, 16, 16, 16, half, row_major> a_frag;
    fragment<matrix_b, 16, 16, 16, half, col_major> b_frag;
    fragment<accumulator, 16, 16, 16, half> c_frag;
    
    // 初始化累加器为 0
    fill_fragment(c_frag, 0.0f);
    
    // K 维度循环
    for (int k = 0; k < K; k += 16) {
        // 加载 A tile
        // A tile: 16 行 × 16 列，从 A[warpM*16, k] 开始
        int a_row = warpM * 16;
        int a_col = k;
        wmma::load_matrix_sync(a_frag, &A[a_row * K + a_col], K);
        
        // 加载 B tile
        // B tile: 16 行 × 16 列，从 B[k, warpN*16] 开始
        int b_row = k;
        int b_col = warpN * 16;
        wmma::load_matrix_sync(b_frag, &B[b_row * N + b_col], N);
        
        // 执行矩阵乘加
        // C += A × B
        wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
    }
    
    // 保存结果
    int c_row = warpM * 16;
    int c_col = warpN * 16;
    wmma::store_matrix_sync(&C[c_row * N + c_col], c_frag, N, layoutC);
}


// 使用 BF16 的 Tensor Core
__global__ void wmma_bf16_gemm(float* C, const __nv_bfloat16* A, 
                                const __nv_bfloat16* B, int M, int N, int K) {
    
    int warpM = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
    int warpN = blockIdx.y;
    
    // BF16 A, FP32 累加（高精度）
    fragment<matrix_a, 16, 16, 16, __nv_bfloat16, row_major> a_frag;
    fragment<matrix_b, 16, 16, 16, __nv_bfloat16, col_major> b_frag;
    fragment<accumulator, 16, 16, 16, float> c_frag;  // FP32 累加
    
    fill_fragment(c_frag, 0.0f);
    
    for (int k = 0; k < K; k += 16) {
        int a_row = warpM * 16;
        int a_col = k;
        wmma::load_matrix_sync(a_frag, &A[a_row * K + a_col], K);
        
        int b_row = k;
        int b_col = warpN * 16;
        wmma::load_matrix_sync(b_frag, &B[b_row * N + b_col], N);
        
        // MMA 同步
        wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
    }
    
    // 保存时转换为 FP32
    int c_row = warpM * 16;
    int c_col = warpN * 16;
    wmma::store_matrix_sync(&C[c_row * N + c_col], c_frag, N, wmma::mem_row_major);
}

第5节：CUDA 编译链路——从源码到 SASS

5.1 编译流程

┌─────────────────────────────────────────────────────────────────────────────┐
│                        CUDA 编译链路                                          │
├─────────────────────────────────────────────────────────────────────────────┤
│                                                                             │
│  源文件 (.cu / .cpp)                                                         │
│       ↓                                                                      │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │                     NVCC 前端 (C++ Frontend)                         │   │
│  │  - CUDA C++ 语法解析                                                  │   │
│  │  - Device/Host 代码分离                                               │   │
│  │  - PTX Generation                                                    │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│       ↓                                                                      │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │                        PTX (Parallel Thread eXecution)               │   │
│  │  - Virtual ISA (虚拟机指令集)                                         │   │
│  │  - 类似于 RISC (单地址，3地址指令)                                     │   │
│  │  - 示例: ld.global.f32, fma.rn.f32, exit                            │   │
│  │  - 可跨 GPU 架构兼容                                                  │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│       ↓                                                                      │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │                        SASS (Streaming ASSembler)                    │   │
│  │  - 硬件 ISA (真实指令集)                                              │   │
│  │  - 示例: LDG.E.64, FFMA, EXIT                                        │   │
│  │  - 特定于 GPU 架构 (Sm_70, Sm_80, Sm_86, Sm_89, Sm_90)               │   │
│  │  - 二进制格式                                                        │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│       ↓                                                                      │
│  ┌─────────────────────────────────────────────────────────────────────┐   │
│  │                    GPU Binary (cubin)                                │   │
│  │  - 可加载到 GPU 执行                                                   │   │
│  │  - 包含 kernel 代码、常量、符号表                                       │   │
│  └─────────────────────────────────────────────────────────────────────┘   │
│                                                                             │
└─────────────────────────────────────────────────────────────────────────────┘

5.2 PTX 示例

cuda

// CUDA 源码
__global__ void vector_add(float* c, const float* a, const float* b, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

asm

// 对应的 PTX (可以通过 nvcc -ptx 查看)
//
// .visible .entry _Z12vector_addPfPKfS0_i(
//     .param .u64 _Z12vector_addPfPKfS0_i_param_0,
//     .param .u64 _Z12vector_addPfPKfS0_i_param_1,
//     .param .u64 _Z12vector_addPfPKfS0_i_param_2,
//     .param .u32 _Z12vector_addPfPKfS0_i_param_3
// )
{
    .reg .pred %p<4>;
    .reg .f32 %f<7>;
    .reg .b32 %r<5>;
    .reg .b64 %rd<10>;
    
    ld.param.u64 %rd1, [_Z12vector_addPfPKfS0_i_param_0];  // c 指针
    ld.param.u64 %rd2, [_Z12vector_addPfPKfS0_i_param_1];  // a 指针
    ld.param.u64 %rd3, [_Z12vector_addPfPKfS0_i_param_2];  // b 指针
    ld.param.u32 %r2, [_Z12vector_addPfPKfS0_i_param_3];    // n
    
    mov.u32 %r3, %ctaid.x;       // blockIdx.x
    mov.u32 %r4, %ntid.x;        // blockDim.x
    mov.u32 %r5, %tid.x;         // threadIdx.x
    mad.lo.s32 %r1, %r3, %r4, %r5;  // idx = blockIdx.x * blockDim.x + threadIdx.x
    
    setp.ge.s32 %p1, %r1, %r2;   // if (idx >= n) goto $exit
    @%p1 bra $exit;
    
    // c[idx] = a[idx] + b[idx]
    cvta.to.global.u64 %rd4, %rd2;    // 转换为全局地址
    shl.b32 %r6, %r1, 2;              // idx * 4 (float 大小)
    add.s64 %rd5, %rd4, %r6;         // &a[idx]
    ld.global.f32 %f1, [%rd5];       // 加载 a[idx]
    
    cvta.to.global.u64 %rd6, %rd3;
    add.s64 %rd7, %rd6, %r6;
    ld.global.f32 %f2, [%rd7];      // 加载 b[idx]
    
    add.f32 %f3, %f1, %f2;           // f3 = f1 + f2
    
    cvta.to.global.u64 %rd8, %rd1;
    add.s64 %rd9, %rd8, %r6;
    st.global.f32 [%rd9], %f3;      // c[idx] = f3
    
$exit:
    ret;
}

5.3 nvcc 编译选项

bash

# 基础编译选项
nvcc -o kernel kernel.cu                    # 输出可执行文件
nvcc -c -o kernel.o kernel.cu               # 编译为目标文件
nvcc -ptx -o kernel.ptx kernel.cu            # 生成 PTX

# 优化选项
nvcc -O3 kernel.cu                          # 最高优化级别
nvcc -use_fast_math kernel.cu               # 使用快速数学函数（低精度）
nvcc --maxrregcount=64 kernel.cu            # 最大寄存器数

# 架构选项
nvcc -arch=sm_80 kernel.cu                  # 针对 Ampere (A100)
nvcc -arch=sm_90a kernel.cu                 # 针对 Hopper (H100)
nvcc -code=sm_80,sm_86 kernel.cu           # 多架构支持

# 常见选项组合
# 生产部署
nvcc -O3 --use_fast_math -lineinfo \
     -arch=sm_80 -code=sm_80 \
     -maxrregcount=128 \
     -o kernel.cu

# 调试
nvcc -G -g -lineinfo \
     -arch=sm_80 \
     -o kernel_debug.cu

第6节：性能分析与优化检查清单

6.1 Nsight Compute 分析

bash

# 使用 Nsight Compute 分析 kernel
ncu --set full ./matrix_mul

# 查看特定指标
ncu --metrics sm__throughput.avg.pct_of_peak_sustained, \\
              sm__warps_per_active_cycle_pct, \\
              dram__bytes.sum, \\
              sm__sass_average_data_bytes_per_sector_fetched_pipe_lsu_op_ld_global \\
              ./matrix_mul

# 导出报告
ncu --export report.ncu-rep ./matrix_mul

6.2 内存效率分析表

指标	理想值	差值	优化方向
Global Memory 效率	>80%	<50%	Coalesced 访问
Shared Memory 效率	>90%	<60%	减少 bank conflict
L1/Tex Cache 命中率	>80%	<50%	数据复用
Warp 活跃度	100%	<50%	减少 divergence
寄存器压力	<64/thread	>128	减少寄存器使用

6.3 优化检查清单

cuda

// GPU Kernel 优化检查清单

// 1. 内存访问
// □ Global Memory 访问是否 coalesced?
// □ 避免跨行矩阵访问
// □ 使用异步内存操作 (cudaMemcpyAsync)
// □ 考虑 Unified Memory vs 分离内存

// 2. Shared Memory
// □ Shared Memory 配置是否合适?
// □ 是否有 bank conflict?
// □ 数据是否在 shared memory 中复用?

// 3. 指令效率
// □ 是否使用 Tensor Core?
// □ 是否可以使用 half/BF16?
// □ 是否使用 FMA 指令?
// □ 是否有寄存器溢出?

// 4. 并行度
// □ Block 数量是否足够 (每个 SM 至少 2-4 blocks)?
// □ Warp 是否满编?
// □ 是否有 branch divergence?

// 5. 异步操作
// □ 计算和内存传输是否可以重叠?
// □ 是否使用了 Double Buffering?
// □ CUDA Streams 是否正确使用?

__global__ void optimized_kernel(float* C, const float* A, const float* B, int N) {
    // 优化后的 GEMM Kernel 模板
    // 1. 使用 shared memory tile
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
    
    // 2. 线程索引计算
    int bx = blockIdx.x;
    int by = blockIdx.y;
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    
    int row = by * BLOCK_SIZE + ty;
    int col = bx * BLOCK_SIZE + tx;
    
    // 3. 累加器
    float sum = 0.0f;
    
    // 4. 主循环
    for (int m = 0; m < (N + BLOCK_SIZE - 1) / BLOCK_SIZE; m++) {
        // Coalesced 加载 A
        int a_col = m * BLOCK_SIZE + tx;
        if (a_col < N && row < N) {
            As[ty][tx] = A[row * N + a_col];
        } else {
            As[ty][tx] = 0.0f;
        }
        
        // Coalesced 加载 B
        int b_row = m * BLOCK_SIZE + ty;
        if (b_row < N && col < N) {
            Bs[ty][tx] = B[b_row * N + col];
        } else {
            Bs[ty][tx] = 0.0f;
        }
        
        __syncthreads();
        
        // 计算
        for (int k = 0; k < BLOCK_SIZE; k++) {
            sum += As[ty][k] * Bs[k][tx];
        }
        
        __syncthreads();
    }
    
    // 写入结果
    if (row < N && col < N) {
        C[row * N + col] = sum;
    }
}

升华

┌─────────────────────────────────────────────────────────────────────────────────┐ │ GPU 优化的核心原则 │ ├─────────────────────────────────────────────────────────────────────────────────┤ │ │ │ 1. 并行是根本：GPU 设计的核心是大规模并行，不充分利用就是浪费 │ │ │ │ 2. 内存带宽是瓶颈：Global Memory 带宽有限，所有优化都要围绕这个约束 │ │ │ │ 3. Coalescing 是关键：连续线程访问连续地址，硬件合并内存访问 │ │ │ │ 4. Shared Memory 是加速器：正确使用可以减少 Global Memory 访问 │ │ │ │ 5. Tensor Core 是未来：矩阵运算必须用 Tensor Core 才能达到硬件峰值 │ │ │ └─────────────────────────────────────────────────────────────────────────────────┘

"AI 可查 vs 必须理解"清单

必须理解（不理解就等于不会）：

🔴 Warp 是 GPU 的基本执行单位——32 线程同步执行，不知道就无法理解 divergence
🔴 Coalesced 访问的原理——连续线程访问连续地址，否则带宽利用率极低
🔴 Bank Conflict 的本质——同一 warp 访问同 bank 导致串行化
🔴 Shared Memory 的用途——Block 内线程共享数据，减少 Global Memory 访问

AI 可查（知道去哪查就行）：

✅ 具体 GPU 架构的参数（如 A100 的 L2 Cache = 40 MB）——查 NVIDIA 官方文档
✅ Tensor Core WMMA API 的具体语法——每个版本略有变化
✅ Nsight Compute 的具体指标含义——官方文档非常详细
✅ PTX/SASS 指令的具体编码——需要时查官方 ISA 文档

学习状态：🟡 开始学习

CUDA 后端：合并访存与 Tensor Core / CUDA Backends, Memory Coalescing, and Tensor Cores ​

第1节：GPU 执行模型——SIMT 架构详解 ​

1.1 GPU vs CPU：设计哲学的差异 ​

1.2 CUDA 编程模型层级 ​

1.3 SM (Streaming Multiprocessor) 内部结构 ​

第2节：GPU 内存层级——性能的关键 ​

2.1 内存层级与延迟 ​

2.2 Shared Memory 配置 ​

2.3 Global Memory 访问模式 ​

第3节：Bank Conflict——Shared Memory 的陷阱 ​

3.1 Shared Memory Bank 架构 ​

3.2 Bank Conflict 示例 ​

第4节：Tensor Core——矩阵计算加速器 ​

4.1 Tensor Core 架构 ​

4.2 WMMA API 使用 ​

第5节：CUDA 编译链路——从源码到 SASS ​

5.1 编译流程 ​

5.2 PTX 示例 ​

5.3 nvcc 编译选项 ​

第6节：性能分析与优化检查清单 ​

6.1 Nsight Compute 分析 ​

6.2 内存效率分析表 ​

6.3 优化检查清单 ​

升华 ​

"AI 可查 vs 必须理解"清单 ​

CUDA 后端：合并访存与 Tensor Core / CUDA Backends, Memory Coalescing, and Tensor Cores

第1节：GPU 执行模型——SIMT 架构详解

1.1 GPU vs CPU：设计哲学的差异

1.2 CUDA 编程模型层级

1.3 SM (Streaming Multiprocessor) 内部结构

第2节：GPU 内存层级——性能的关键

2.1 内存层级与延迟

2.2 Shared Memory 配置

2.3 Global Memory 访问模式

第3节：Bank Conflict——Shared Memory 的陷阱

3.1 Shared Memory Bank 架构

3.2 Bank Conflict 示例

第4节：Tensor Core——矩阵计算加速器

4.1 Tensor Core 架构

4.2 WMMA API 使用

第5节：CUDA 编译链路——从源码到 SASS

5.1 编译流程

5.2 PTX 示例

5.3 nvcc 编译选项

第6节：性能分析与优化检查清单

6.1 Nsight Compute 分析

6.2 内存效率分析表

6.3 优化检查清单

升华

"AI 可查 vs 必须理解"清单