4. CUDA内存管理：全局内存、共享内存、常量内存、纹理内存

内存管理，说白了就是CUDA编程的命门。我见过太多新手，核函数写得飞起，一跑就崩，最后发现是内存访问出了问题。今天咱们就把这块硬骨头啃下来。

4.1 全局内存：最常用，也最慢

全局内存是GPU上最大的存储空间，所有线程都能访问。但大是有代价的——访问延迟高达几百个时钟周期。

我刚开始做项目时，写过一个矩阵转置的核函数。线程按行读取，按列写入，结果性能惨不忍睹。后来改成共享内存做中转，速度直接翻了5倍。你想想看，这就是内存访问模式的重要性。

// 糟糕的访问模式：按列访问
__global__ void bad_access(float *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[idx * 1024];  // 跨步访问，无法合并
}

// 好的访问模式：按行访问
__global__ void good_access(float *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[idx];  // 连续访问，完美合并
}

4.2 共享内存：片上高速缓存

共享内存是GPU上的"L1缓存"，由程序员手动管理。延迟只有几个时钟周期，比全局内存快两个数量级。

我个人习惯把共享内存当作"线程间的数据交换站"。同一个block内的线程可以通过它共享数据，避免重复访问全局内存。

// 共享内存的典型用法：矩阵分块
__global__ void matmul_shared(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;
}

4.3 常量内存：只读且广播

常量内存适合存储那些所有线程都用到的只读数据。比如物理常数、查找表之类的。它有专门的缓存，访问延迟低，而且能广播给所有线程。

我曾经在做一个流体模拟项目时，把物理参数放在常量内存里。一开始用的全局内存，每次读取都要几十个周期。换成常量内存后，基本零延迟。嗯，这里要注意：常量内存只有64KB，别想着存大数据。

// 声明常量内存
__constant__ float params[256];

// 主机端拷贝
cudaMemcpyToSymbol(params, host_params, sizeof(float) * 256);

// 设备端使用
__global__ void compute(float *data) {
    int idx = threadIdx.x;
    data[idx] = data[idx] * params[0] + params[1];  // 自动广播
}

4.4 纹理内存：为空间局部性而生

纹理内存听起来玄乎，其实就是为了处理"二维或三维数据"而优化的。它内置了硬件插值、边界处理等功能。做图像处理、科学计算时特别好用。

// 纹理内存声明
texture<float, 2, cudaReadModeElementType> tex;

// 绑定纹理
cudaArray *array;
cudaMallocArray(&array, &desc, width, height);
cudaMemcpyToArray(array, 0, 0, host_data, size, cudaMemcpyHostToDevice);
cudaBindTextureToArray(tex, array);

// 核函数中使用
__global__ void texture_kernel(float *output) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    output[y * width + x] = tex2D(tex, x + 0.5f, y + 0.5f);  // 硬件插值
}

4.5 内存分配与释放

CUDA的内存管理遵循"谁分配谁释放"的原则。我见过不少新手忘记释放内存，导致显存泄漏。说实话，调试这种问题特别头疼。

函数	用途	注意事项
cudaMalloc	分配全局内存	返回设备指针
cudaFree	释放全局内存	必须配对使用
cudaMallocHost	分配页锁定内存	提高DMA传输速度
cudaFreeHost	释放页锁定内存	别用free()释放

4.6 内存拷贝：主机与设备的数据桥梁

内存拷贝是CUDA编程中最容易出错的环节。为什么？因为主机内存和设备内存是物理隔离的，不能直接访问。

// 标准的内存拷贝
float *d_data;
cudaMalloc(&d_data, size);
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);  // 主机到设备
// ... 计算 ...
cudaMemcpy(h_result, d_data, size, cudaMemcpyDeviceToHost);  // 设备到主机
cudaFree(d_data);

我建议使用异步拷贝来隐藏传输延迟。特别是当你有多个数据流时，可以让计算和传输重叠进行。

知识体系总览

这张图把CUDA的内存层次结构梳理清楚了。从上到下，容量递减但速度递增。实际开发中，你要根据数据的使用模式，选择最合适的内存类型。