4. CUDA进阶：共享内存、同步、流与事件、性能优化初步

4.1 核心原理：从全局内存到共享内存的跨越

在CUDA编程模型中，线程访问数据的内存层级直接决定了内核的执行效率。全局内存（Global Memory）容量大但延迟高达数百周期，而共享内存（Shared Memory）是片上存储，延迟仅为全局内存的十分之一左右，带宽也高得多。共享内存的本质是一个Block内所有线程共享的高速暂存器，它的生命周期与Block一致，Block销毁后数据即失效。

共享内存的典型应用场景包括：数据复用（多个线程读取同一份数据）、块内归约（求和、求最大值）、矩阵分块转置以及卷积/池化操作的滑动窗口。其核心思想是：先将全局内存中的数据协作加载到共享内存，经过__syncthreads()同步后，再从共享内存中高速读取，避免每个线程都直接访问全局内存。

4.2 同步机制：__syncthreads() 与内存可见性

共享内存的使用必须配合同步操作。当多个线程向共享内存写入数据后，其他线程可能尚未完成写入就开始读取，导致数据竞争。__syncthreads() 是块内屏障，它保证：调用点之前的所有共享内存写操作对块内所有线程可见，且所有线程都到达该屏障后才继续执行。

同步的典型模式是“写-同步-读”三段式：

__global__ void kernel(float* input, float* output, int n) {
    __shared__ float tile[256];
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 阶段1：协作加载到共享内存
    if (idx < n) {
        tile[threadIdx.x] = input[idx];
    }
    
    // 阶段2：同步屏障
    __syncthreads();
    
    // 阶段3：从共享内存读取并计算
    float val = tile[threadIdx.x] * 2.0f;
    if (idx < n) {
        output[idx] = val;
    }
}

4.3 流与事件：并发执行与时间测量

CUDA流（Stream）是GPU上按顺序执行的操作序列。默认流（Default Stream）是同步流，而创建多个非默认流可以实现内核执行与数据拷贝的重叠，提升GPU利用率。事件（Event）则用于记录时间戳或标记流中的特定位置，常用于性能分析。

典型的多流并发模式如下：

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 将不同数据块的内核分发到不同流
kernel<<<grid, block, 0, stream1>>>(data1, ...);
kernel<<<grid, block, 0, stream2>>>(data2, ...);

// 事件计时
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, stream1);
// ... 执行操作 ...
cudaEventRecord(stop, stream1);
cudaEventSynchronize(stop);
float ms = 0.0f;
cudaEventElapsedTime(&ms, start, stop);

API	作用	典型使用场景
cudaStreamCreate	创建非默认流	多任务并发
cudaStreamSynchronize	阻塞主机直到流完成	等待特定流结束
cudaEventRecord	在流中插入事件标记	时间测量、流间依赖
cudaEventElapsedTime	计算两个事件间耗时	内核性能分析
cudaStreamWaitEvent	让流等待事件	跨流同步

4.4 性能优化初步：访存合并与占用率

性能优化的第一原则是让全局内存访问尽量合并（Coalesced）。当同一Warp（32个线程）访问连续的内存地址时，硬件将其合并为少数几次事务，大幅减少访存次数。反之，若线程访问跳跃地址，则产生多次事务，性能急剧下降。

第二个关键指标是占用率（Occupancy），即活跃Warp数与硬件最大Warp数的比值。高占用率能隐藏内存延迟，但受限于寄存器数量、共享内存大小和Block维度。使用cudaOccupancyMaxPotentialBlockSize可以自动计算最优Block大小。

4.5 案例：矩阵转置的共享内存优化

以32×32矩阵转置为例，朴素实现中每个线程读取一行写一列，导致写操作完全不合并。优化方案是使用共享内存作为中转：线程协作将矩阵块读入共享内存（读合并），同步后按转置顺序写回（写合并）。

__global__ void transpose_shared(float* in, float* out, int n) {
    __shared__ float tile[32][32];
    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;
    
    // 协作加载（读合并）
    if (x < n && y < n) {
        tile[threadIdx.y][threadIdx.x] = in[y * n + x];
    }
    __syncthreads();
    
    // 转置写回（写合并）
    int x2 = blockIdx.y * 32 + threadIdx.x;
    int y2 = blockIdx.x * 32 + threadIdx.y;
    if (x2 < n && y2 < n) {
        out[y2 * n + x2] = tile[threadIdx.x][threadIdx.y];
    }
}

该优化将转置性能提升约5-10倍，核心收益来自将非合并的写访问转化为合并访问，同时利用共享内存消除了Bank Conflict（通过调整填充或访问模式）。

4.6 误区与复盘

复盘建议：每次优化后使用ncu（NVIDIA Nsight Compute）分析内存吞吐、占用率和Warp Stall原因，用数据驱动下一步优化方向，而非凭经验猜测。