跳到主要内容

内存模型

CUDA 提供了层次化的内存模型，不同类型的内存具有不同的速度、容量和作用域。理解内存模型是优化 CUDA 程序性能的关键。

内存层次结构概览​

CUDA 内存模型包含多种类型的内存空间：

┌─────────────────────────────────────────────────────────────┐
│                      GPU 内存层次结构                         │
├─────────────────────────────────────────────────────────────┤
│  寄存器 (Registers)          - 线程私有，最快                │
│  局部内存 (Local Memory)     - 线程私有，较慢                │
│  共享内存 (Shared Memory)    - Block 内共享，很快            │
│  全局内存 (Global Memory)    - 所有线程可访问，较慢          │
│  常量内存 (Constant Memory)  - 只读，有缓存                  │
│  纹理内存 (Texture Memory)   - 只读，有缓存，支持插值        │
└─────────────────────────────────────────────────────────────┘
内存类型作用域生命周期速度容量
寄存器线程线程最快有限
局部内存线程线程慢大
共享内存BlockBlock很快有限
全局内存Grid应用程序慢大
常量内存Grid应用程序中等64KB
纹理内存Grid应用程序中等大

寄存器（Registers）​

寄存器是 GPU 上最快的内存空间，每个线程拥有独立的寄存器。

自动变量​

核函数中声明的自动变量通常存储在寄存器中：

__global__ void kernel(float *data) {
    int idx = threadIdx.x;
    float temp = data[idx];
    temp = temp * 2.0f;
    data[idx] = temp;
}

这里 idx 和 temp 都存储在寄存器中。

寄存器限制​

每个线程的寄存器数量有限（通常 255 个），过多的寄存器使用会：

  1. 限制 Block 的并发数量
  2. 降低 SM 占用率
  3. 导致寄存器溢出到局部内存

控制寄存器使用​

使用 __launch_bounds__ 限制寄存器使用：

__global__ void __launch_bounds__(256, 2) kernel() {
}

参数含义：

  • 第一个参数：每个 Block 的最大线程数
  • 第二个参数：每个 SM 的最小 Block 数

局部内存（Local Memory）​

局部内存虽然名称暗示「局部」，但实际上位于全局内存中。

使用场景​

以下情况变量会被放置在局部内存：

  1. 寄存器不足时的溢出
  2. 使用动态索引访问的数组
  3. 超过寄存器限制的大型结构体
__global__ void kernel() {
    float largeArray[256];
    int idx = threadIdx.x % 256;
    largeArray[idx] = 1.0f;
}

性能影响​

局部内存访问速度与全局内存相同，应该尽量避免使用。编译器会通过警告提示寄存器溢出。

共享内存（Shared Memory）​

共享内存是位于 GPU 芯片上的高速内存，同一 Block 内的所有线程可以共享访问。

声明方式​

使用 __shared__ 关键字声明：

__global__ void kernel(float *data) {
    __shared__ float sharedData[256];
    
    int idx = threadIdx.x;
    sharedData[idx] = data[idx];
    
    __syncthreads();
    
    float result = sharedData[255 - idx];
    data[idx] = result;
}

动态共享内存​

运行时确定大小的共享内存：

extern __shared__ float dynamicShared[];

__global__ void kernel(float *data, int size) {
    dynamicShared[threadIdx.x] = data[threadIdx.x];
}

kernel<<<grid, block, sharedMemSize>>>(data, size);

共享内存大小​

每个 SM 的共享内存大小因架构而异：

架构每个 SM 共享内存
Fermi48KB
Kepler48KB
Maxwell64KB
Pascal64KB
Volta96KB
Ampere164KB

Bank 冲突​

共享内存被划分为 32 个 Bank（存储体），每个 Bank 宽度为 4 字节。当多个线程访问同一 Bank 的不同地址时，会发生 Bank 冲突。

__shared__ float data[256];

int idx = threadIdx.x;
float val = data[idx];

上面的代码没有 Bank 冲突，因为连续线程访问连续地址。

__shared__ float data[256];

int idx = threadIdx.x * 32;
float val = data[idx];

这段代码会发生严重的 Bank 冲突，因为所有线程访问同一个 Bank。

避免 Bank 冲突​

使用填充（Padding）避免 Bank 冲突：

__shared__ float data[256 + 1];

int idx = threadIdx.x;
float val = data[idx];

全局内存（Global Memory）​

全局内存是 GPU 上容量最大但延迟最高的内存，所有线程都可以访问。

分配与释放​

float *d_data;
cudaMalloc(&d_data, size);
cudaFree(d_data);

合并访问​

当 Warp 中的线程访问连续的内存地址时，GPU 可以将这些访问合并为一次内存事务，大幅提高带宽利用率。

__global__ void goodAccess(float *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[idx];
}

__global__ void badAccess(float *data, int stride) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[idx * stride];
}

内存对齐​

为了获得最佳性能，内存访问应该对齐到 128 字节边界：

cudaMallocPitch(&d_data, &pitch, width * sizeof(float), height);

使用示例​

__global__ void copyKernel(float *dst, float *src, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        dst[idx] = src[idx];
    }
}

常量内存（Constant Memory）​

常量内存是只读内存，有专门的缓存，适合存储只读的常量数据。

声明与使用​

__constant__ float constData[256];

int main() {
    float h_data[256];
    cudaMemcpyToSymbol(constData, h_data, sizeof(h_data));
    
    kernel<<<grid, block>>>();
}

__global__ void kernel() {
    float val = constData[threadIdx.x];
}

特点​

  1. 容量限制：64KB
  2. 有专门的缓存
  3. 适合广播读取（所有线程读取相同地址）
  4. 单次读取可以广播给 Warp 中所有线程

最佳实践​

当所有线程读取相同地址时，常量内存效率最高：

__constant__ float scaleFactor;

__global__ void scaleKernel(float *data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        data[idx] *= scaleFactor;
    }
}

纹理内存（Texture Memory）​

纹理内存是一种特殊的只读内存，具有硬件缓存和插值功能。

特点​

  1. 支持一维、二维、三维纹理
  2. 硬件插值（线性插值）
  3. 边界处理（钳位、循环）
  4. 适合图像处理和空间局部性访问

使用示例​

texture<float, 2, cudaReadModeElementType> texRef;

int main() {
    cudaArray *cuArray;
    cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
    cudaMallocArray(&cuArray, &desc, width, height);
    
    cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);
    
    texRef.addressMode[0] = cudaAddressModeWrap;
    texRef.addressMode[1] = cudaAddressModeWrap;
    texRef.filterMode = cudaFilterModeLinear;
    texRef.normalized = true;
    
    cudaBindTextureToArray(texRef, cuArray);
    
    textureKernel<<<grid, block>>>(width, height);
    
    cudaUnbindTexture(texRef);
    cudaFreeArray(cuArray);
}

__global__ void textureKernel(int width, int height) {
    float x = (float)threadIdx.x / width;
    float y = (float)threadIdx.y / height;
    
    float val = tex2D(texRef, x, y);
}

统一内存（Unified Memory）​

统一内存（CUDA 6.0 引入）提供了一个统一的内存地址空间，CPU 和 GPU 都可以访问。

基本使用​

int main() {
    float *data;
    cudaMallocManaged(&data, N * sizeof(float));
    
    for (int i = 0; i < N; i++) {
        data[i] = 1.0f;
    }
    
    kernel<<<grid, block>>>(data, N);
    cudaDeviceSynchronize();
    
    for (int i = 0; i < N; i++) {
        printf("%f ", data[i]);
    }
    
    cudaFree(data);
}

优点​

  1. 简化编程模型
  2. 自动数据迁移
  3. 减少显式的内存拷贝

注意事项​

  1. 数据迁移有开销
  2. 需要适当的同步
  3. 对于频繁访问的数据，手动管理内存可能更高效

内存访问优化策略​

1. 合并访问​

确保 Warp 中的线程访问连续的内存地址：

__global__ void optimizedKernel(float *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[idx];
}

2. 使用共享内存缓存​

将频繁访问的数据缓存到共享内存：

__global__ void matrixMul(float *A, float *B, float *C, int width) {
    __shared__ float sA[TILE_SIZE][TILE_SIZE];
    __shared__ float sB[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    
    for (int t = 0; t < width / TILE_SIZE; t++) {
        sA[threadIdx.y][threadIdx.x] = A[row * width + t * TILE_SIZE + threadIdx.x];
        sB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * width + col];
        
        __syncthreads();
        
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += sA[threadIdx.y][k] * sB[k][threadIdx.x];
        }
        
        __syncthreads();
    }
    
    C[row * width + col] = sum;
}

3. 避免 Bank 冲突​

使用填充或调整访问模式：

__shared__ float data[TILE_SIZE][TILE_SIZE + 1];

4. 使用只读缓存​

对于只读数据，使用 __ldg 内联函数或纹理内存：

__global__ void kernel(const float *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = __ldg(&data[idx]);
}

小结​

本章介绍了 CUDA 的内存模型：

  1. 寄存器：最快，线程私有
  2. 共享内存：很快，Block 内共享
  3. 全局内存：容量大，延迟高
  4. 常量内存：只读，适合广播
  5. 纹理内存：支持插值和边界处理
  6. 统一内存：简化编程模型

内存访问优化是 CUDA 性能优化的核心。下一章将介绍 核函数详解，深入理解核函数的编写和执行。