CUDA内存优化技术详解

在CUDA编程中，内存管理是决定程序性能的关键因素之一。本文将详细介绍CUDA内存层次结构及其优化策略。

CUDA内存层次结构

CUDA提供了多个级别的内存，每一级都有不同的访问特性和性能特征：

1. 寄存器 (Registers)

每个线程独享
访问速度最快
容量有限（通常32KB/SM）
编译器自动分配

2. 共享内存 (Shared Memory)

每个线程块内共享
访问速度仅次于寄存器
容量通常48KB或96KB/SM（可配置）
需要显式声明和使用

__global__ void example_kernel()
{
    __shared__ float sdata[256];  // 声明共享内存
    // ...
}

3. 全局内存 (Global Memory)

所有线程可访问
容量最大（GB级别）
访问速度相对较慢
但带宽可以通过合并访问优化

4. 常量内存 (Constant Memory)

只读，缓存优化
容量限制为64KB
适合存储不变参数

5. 纹理内存 (Texture Memory)

只读，缓存优化
支持硬件插值
适合图像处理算法

内存访问模式优化

合并访问 (Coalesced Access)

// 优化的访问模式
__global__ void optimized_copy(float* input, float* output)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    output[idx] = input[idx];  // 连续地址访问
}

// 非优化的访问模式
__global__ void unoptimized_copy(float* input, float* output, int stride)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    output[idx] = input[idx * stride];  // 跳跃访问
}

银行冲突 (Bank Conflicts)

共享内存被分为32个bank，当同一warp的线程访问相同bank时发生冲突：

__global__ void bank_conflict_example()
{
    __shared__ float sdata[32][33];  // 添加填充避免冲突
    
    int tid = threadIdx.x;
    // 避免这种情况：多个线程访问同一列
    // sdata[0][tid] -> 可能造成银行冲突
    sdata[tid][0] -> 无冲突访问
}

内存优化实践技巧

1. 使用共享内存减少全局内存访问

__global__ void matrix_multiply_optimized(float* A, float* B, float* C, int N)
{
    __shared__ float As[16][16];
    __shared__ float Bs[16][16];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    // 计算输出子矩阵的索引
    int row = by * 16 + ty;
    int col = bx * 16 + tx;
    
    float sum = 0.0f;
    
    // 分块计算，利用共享内存
    for (int m = 0; m < (N + 15)/16; ++m) {
        // 加载数据到共享内存
        if (row < N && m * 16 + tx < N)
            As[ty][tx] = A[row * N + m * 16 + tx];
        else
            As[ty][tx] = 0.0f;
            
        if (col < N && m * 16 + ty < N)
            Bs[ty][tx] = B[(m * 16 + ty) * N + col];
        else
            Bs[ty][tx] = 0.0f;
        
        __syncthreads();
        
        // 计算部分乘积
        for (int k = 0; k < 16; ++k)
            sum += As[ty][k] * Bs[k][tx];
        
        __syncthreads();
    }
    
    if (row < N && col < N)
        C[row * N + col] = sum;
}

2. 使用统一内存 (Unified Memory)

void unified_memory_example(int N)
{
    float *d_data;
    cudaMallocManaged(&d_data, N * sizeof(float));
    
    // 主机和设备端都可以访问
    for (int i = 0; i < N; i++) {
        d_data[i] = i;  // 在主机端初始化
    }
    
    kernel<<<blocks, threads>>>(d_data, N);
    cudaDeviceSynchronize();
    
    cudaFree(d_data);
}

3. 流和异步内存拷贝

void async_memory_transfer()
{
    const int N = 1024 * 1024;
    size_t size = N * sizeof(float);
    
    float *h_data = (float*)malloc(size);
    float *d_data;
    cudaMalloc(&d_data, size);
    
    // 创建CUDA流
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    
    // 异步内存拷贝
    cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);
    
    // 在设备上执行计算
    kernel<<<blocks, threads, 0, stream>>>(d_data, N);
    
    // 同步流
    cudaStreamSynchronize(stream);
    
    cudaStreamDestroy(stream);
    cudaFree(d_data);
    free(h_data);
}