第4章 GPU 内存管理 - 显存的层次结构（GPU Memory Management）

一、导读

1.1 本章学习目标

通过本章的学习，你将能够：

理解 GPU 内存层次结构的各个层级及其特性

掌握全局内存、共享内存、常量内存、纹理内存的使用方法

理解内存合并访问（Coalesced Access）的原理和重要性

学会使用共享内存优化数据复用和减少全局内存访问

掌握寄存器使用的最佳实践

理解统一内存（Unified Memory）的工作机制和适用场景

能够分析和优化内核函数的内存访问模式

掌握内存带宽的测量和优化方法

1.2 为什么本章如此重要

在 GPU 编程中，内存访问往往是性能瓶颈所在。即使计算本身非常高效，糟糕的内存访问模式也可能使程序慢上数倍甚至数十倍。理解 GPU 的内存层次结构，并学会优化内存访问，是编写高性能 CUDA 程序的关键技能。

与 CPU 不同，GPU 拥有独特的内存架构：

高带宽但高延迟的全局内存：容量大，但访问速度慢
高速但容量有限的共享内存：片上内存，访问速度极快
只读的常量内存和纹理内存：针对特定访问模式优化

每种内存类型都有其适用的场景和最佳实践。掌握这些知识，使你能够根据算法特点选择最合适的内存策略。

此外，内存访问模式对性能的影响远超许多初学者的想象。同样的计算，仅仅因为访问模式不同，性能可能相差数倍。本章将帮助你建立对内存访问模式的直觉，使你能够预测和优化代码的内存行为。

1.3 核心问题引导

在深入学习之前，请思考以下问题：

GPU 有哪些类型的内存？它们各自的特性和用途是什么？

什么是内存合并访问？为什么它对性能如此重要？

共享内存如何帮助提高性能？使用时需要注意什么？

寄存器和共享内存之间有什么区别？

什么是 bank conflict？如何避免？

统一内存简化了编程，但它有性能代价吗？

如何测量程序的内存带宽利用率？

常量内存和纹理内存在什么场景下最有效？

二、核心概念详解

4.1 GPU 内存层次结构概览

GPU 拥有多层次的内存结构，每层在速度、容量和可见性上有所不同：

GPU 内存层次结构（从快到慢）：
┌─────────────────────────────────────────┐
│         Registers (per thread)          │  ← 最快，每线程私有
├─────────────────────────────────────────┤
│      Shared Memory / L1 Cache           │  ← 非常快，块内共享
│         (per Streaming Multiprocessor)   │
├─────────────────────────────────────────┤
│          L2 Cache                       │  ← 较快，全局共享
├─────────────────────────────────────────┤
│     Global Memory (GDDR6 / HBM)         │  ← 较慢，大容量
├─────────────────────────────────────────┤
│     Constant Memory (cached, read-only) │  ← 针对广播访问优化
├─────────────────────────────────────────┤
│     Texture Memory (cached, read-only)  │  ← 针对空间局部性优化
└─────────────────────────────────────────┘

关键特性对比

内存类型	速度	容量	可见性	读/写
寄存器	最快	有限	线程私有	读写
共享内存	非常快	~48-96KB/SM	块内共享	读写
L1 缓存	快	~128KB/SM	SM 内	读写
L2 缓存	中等	~几MB	全局	读写
全局内存	慢	GB 级	全局	读写
常量内存	中等（命中时快）	64KB	全局	只读
纹理内存	中等（命中时快）	部分全局	全局	只读

4.2 全局内存（Global Memory）

全局内存是 GPU 上容量最大、所有线程都可访问的内存。它对应于通过 cudaMalloc 分配的显存。

特性

大容量：现代 GPU 通常有 8-80GB 全局内存
高带宽：GDDR6 可达 500+ GB/s，HBM 可达 1+ TB/s
高延迟：访问需要数百个时钟周期
持久性：内核执行结束后数据仍然保留

分配与访问

cpp

float *d_data;
cudaMalloc((void **)&d_data, size);
// 在内核中直接访问
__global__ void kernel(float *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    data[idx] = ...; // 读写全局内存
}

合并访问（Coalesced Access）

合并访问是 GPU 内存优化的核心概念。当 warp 内的线程访问连续的内存地址时，硬件可以将多个访问合并为少数几个事务，大幅提高带宽利用率。

合并访问示例（warp 内 32 个线程）：
Thread 0: 访问 address[0]
Thread 1: 访问 address[1]
Thread 2: 访问 address[2]
...
Thread 31: 访问 address[31]

→ 硬件合并为 1-2 个 128 字节的事务

非合并访问的代价

如果线程访问不连续的地址，每个访问可能需要单独的事务：

cpp

// 糟糕的访问模式（步幅过大）
__global__ void badAccess(float *data, int stride) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    data[idx * stride] = ...; // 步幅访问导致非合并
}

最佳实践

确保相邻线程访问相邻内存地址
数据结构按行优先组织（C 风格）
避免间接寻址和随机访问

4.3 共享内存（Shared Memory）

共享内存是 GPU 上最快的可读写内存（仅次于寄存器），位于每个 SM 上，由该 SM 上的所有线程块共享。

特性

极低延迟：访问速度与寄存器相当
有限容量：通常 48-96KB 每 SM
块内可见：同一线程块的线程可以共享数据
用户管理：程序员负责数据的加载和同步

声明与使用

cpp

__global__ void kernelWithSharedMem(float *input, float *output, int n) {
    extern __shared__ float sharedData[]; // 动态共享内存
    
    int tid = threadIdx.x;
    int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 从全局内存加载到共享内存
    if (globalIdx < n) {
        sharedData[tid] = input[globalIdx];
    }
    __syncthreads(); // 确保所有线程完成加载
    
    // 使用共享内存进行计算
    if (tid < blockDim.x / 2) {
        sharedData[tid] += sharedData[tid + blockDim.x / 2];
    }
    __syncthreads();
    
    // 写回全局内存
    if (tid == 0) {
        output[blockIdx.x] = sharedData[0];
    }
}

// 启动时指定共享内存大小
kernelWithSharedMem<<<blocks, threads, sharedMemSize>>>(...);

静态 vs 动态共享内存

cpp

// 静态：编译时确定大小
__shared__ float staticShared[256];

// 动态：运行时指定大小
extern __shared__ float dynamicShared[];

Bank Conflict

共享内存被划分为 32 个 bank，每个 bank 可以同时服务一个访问。如果多个线程访问同一个 bank 的不同地址，就会发生 bank conflict，导致访问串行化。

Bank 分布示意：
Address 0 → Bank 0
Address 1 → Bank 1
...
Address 31 → Bank 31
Address 32 → Bank 0
...

// 可能导致 bank conflict 的访问
sharedData[tid * 2]  // 偶数线程都访问 even banks

避免 Bank Conflict 的策略

使用填充（padding）改变映射
重新组织数据布局
使用广播访问（所有线程访问同一地址）

4.4 常量内存（Constant Memory）

常量内存是只读内存，针对广播访问模式优化。当 warp 内所有线程读取同一地址时，性能极佳。

特性

只读：主机端写入，设备端只读
缓存：有专用缓存，命中时速度很快
容量有限：64KB
广播优化：warp 内广播访问只需一个事务

使用方法

cpp

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

// 主机端设置常量数据
float h_constData[256] = {...};
cudaMemcpyToSymbol(constData, h_constData, sizeof(h_constData));

// 设备端读取
__global__ void useConstantMemory(float *output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        output[idx] = constData[idx % 256]; // 广播访问效率高
    }
}

适用场景

查找表（LUT）
滤波器系数
变换矩阵
配置参数

4.5 纹理内存（Texture Memory）

纹理内存原本是用于图形渲染的只读内存，现在也用于通用计算。它针对空间局部性访问模式进行了优化。

特性

只读
内置缓存：针对 2D 空间局部性优化
硬件插值：支持线性插值
边界处理：内置边界模式

使用方法

cpp

texture<float, 2> texRef; // 声明纹理引用

// 绑定数组到纹理
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
cudaArray *cuArray;
cudaMallocArray(&cuArray, &channelDesc, width, height);
cudaMemcpyToArray(cuArray, h_data, size, cudaMemcpyHostToDevice);
texRef.addressMode[0] = cudaAddressModeClamp;
texRef.addressMode[1] = cudaAddressModeClamp;
texRef.filterMode = cudaFilterModeLinear;
texRef.normalized = true;
cudaBindTextureToArray(texRef, cuArray);

// 在内核中使用
__global__ void useTexture(float *output, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < height) {
        float val = tex2D(texRef, (float)x/width, (float)y/height);
        output[y * width + x] = val;
    }
}

适用场景

图像处理
体积渲染
需要插值的查找操作

4.6 寄存器（Registers）

寄存器是每个线程私有的最快存储位置。编译器自动决定哪些变量存储在寄存器中。

特性

最快访问：零延迟
线程私有：其他线程无法访问
数量有限：每个 SM 有固定数量的寄存器
溢出到本地内存：如果寄存器不足，变量会溢出到慢速的本地内存

控制寄存器使用

cpp

// 限制每个线程的最大寄存器数
__launch_bounds__(256, 4) // 256 线程/块，最多 4 块/SM
__global__ void kernel(...) {
    ...
}

// 或使用编译选项
// nvcc --maxrregcount=32 ...

最佳实践

尽量减少局部变量的数量
避免大型数组（会占用大量寄存器或溢出）
让编译器决定寄存器分配（通常是最优的）

4.7 统一内存（Unified Memory）

统一内存（CUDA 6.0 引入）简化了内存管理，允许主机和设备使用相同的指针访问数据。

特性

单一指针：无需分别管理主机和设备指针
按需迁移：数据在首次访问时自动迁移
透明性：简化编程模型

使用方法

cpp

float *data;
cudaMallocManaged(&data, size); // 分配统一内存

// 主机和设备都可以直接使用 data 指针
for (int i = 0; i < n; i++) {
    data[i] = (float)i; // 主机访问
}

kernel<<<blocks, threads>>>(data, n); // 设备访问
cudaDeviceSynchronize(); // 确保设备访问完成

// 主机可以直接读取结果
printf("%f\n", data[0]);

性能考虑

首次访问可能触发页面迁移，产生开销
频繁的主机-设备交替访问会导致页面抖动
对于已知访问模式，手动管理内存可能更高效

预取优化

cpp

// 预取数据到设备
cudaMemPrefetchAsync(data, size, deviceId);

// 预取回主机
cudaMemPrefetchAsync(data, size, cudaCpuDeviceId);

三、程序员的视角

4.1 在实际项目中的应用

矩阵乘法优化

使用共享内存分块技术可以显著提高矩阵乘法的性能：

cpp

__global__ void tiledMatrixMul(float *C, const float *A, const float *B, int N) {
    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
    __shared__ float tileB[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 < (N + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // 加载 tile 到共享内存
        if (row < N && t * TILE_SIZE + threadIdx.x < N)
            tileA[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
        else
            tileA[threadIdx.y][threadIdx.x] = 0.0f;
            
        if (col < N && t * TILE_SIZE + threadIdx.y < N)
            tileB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        else
            tileB[threadIdx.y][threadIdx.x] = 0.0f;
        
        __syncthreads();
        
        // 计算 tile 乘法
        for (int k = 0; k < TILE_SIZE; k++)
            sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
        
        __syncthreads();
    }
    
    if (row < N && col < N)
        C[row * N + col] = sum;
}

图像卷积

使用共享内存存储邻域像素，减少全局内存访问：

cpp

__global__ void convolution(float *output, const float *input, 
                            const float *kernel, int width, int height) {
    __shared__ float tile[TILE_SIZE + KERNEL_RADIUS * 2][TILE_SIZE + KERNEL_RADIUS * 2];
    
    // 加载数据到共享内存（包括边界）
    ...
}

4.2 常见陷阱与误区

误区一：共享内存总是更快

事实：如果使用不当（如 bank conflict），共享内存的性能可能不如良好合并的全局内存访问。

误区二：统一内存没有性能代价

事实：统一内存的页面迁移有开销，频繁的跨处理器访问会导致性能下降。

陷阱：忽视内存对齐

未对齐的内存访问可能导致额外的内存事务。确保数据结构按适当边界对齐。

陷阱：过度使用共享内存

共享内存容量有限，过度使用可能降低 occupancy，反而影响整体性能。

4.3 与其他层级的联系

内存管理与线程模型密切相关：

第3章：线程索引决定了内存访问模式
第5章：共享内存的使用需要同步机制
第8章：内存优化是性能调优的核心

四、代码实战

4.1 演示合并与非合并访问

cpp

#include <stdio.h>
#include <cuda_runtime.h>

// 合并访问
__global__ void coalescedAccess(float *input, float *output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        output[idx] = input[idx] * 2.0f;
    }
}

// 非合并访问（步幅访问）
__global__ void stridedAccess(float *input, float *output, int n, int stride) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx * stride < n) {
        output[idx] = input[idx * stride] * 2.0f;
    }
}

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA error: %s\n", cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while(0)

int main() {
    int n = 1 << 20;
    size_t size = n * sizeof(float);
    
    float *h_input = (float *)malloc(size);
    float *h_output = (float *)malloc(size);
    
    for (int i = 0; i < n; i++) {
        h_input[i] = (float)i;
    }
    
    float *d_input, *d_output;
    CUDA_CHECK(cudaMalloc((void **)&d_input, size));
    CUDA_CHECK(cudaMalloc((void **)&d_output, size));
    CUDA_CHECK(cudaMemcpy(d_input, h_input, size, cudaMemcpyHostToDevice));
    
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));
    
    // 测试合并访问
    CUDA_CHECK(cudaEventRecord(start));
    for (int iter = 0; iter < 100; iter++) {
        coalescedAccess<<<blocksPerGrid, threadsPerBlock>>>(d_input, d_output, n);
    }
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    
    float coalescedTime = 0;
    CUDA_CHECK(cudaEventElapsedTime(&coalescedTime, start, stop));
    printf("Coalesced access time (100 iterations): %.3f ms\n", coalescedTime);
    
    // 测试非合并访问（stride = 32）
    CUDA_CHECK(cudaEventRecord(start));
    for (int iter = 0; iter < 100; iter++) {
        stridedAccess<<<blocksPerGrid, threadsPerBlock>>>(d_input, d_output, n, 32);
    }
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    
    float stridedTime = 0;
    CUDA_CHECK(cudaEventElapsedTime(&stridedTime, start, stop));
    printf("Strided access time (100 iterations): %.3f ms\n", stridedTime);
    
    printf("Slowdown factor: %.2fx\n", stridedTime / coalescedTime);
    
    CUDA_CHECK(cudaFree(d_input));
    CUDA_CHECK(cudaFree(d_output));
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
    free(h_input);
    free(h_output);
    
    return 0;
}

4.2 共享内存归约优化

cpp

#include <stdio.h>
#include <cuda_runtime.h>

// 使用共享内存的并行归约
__global__ void sharedMemReduce(float *input, float *output, int n) {
    extern __shared__ float sharedData[];
    
    int tid = threadIdx.x;
    int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 加载到共享内存
    sharedData[tid] = (globalIdx < n) ? input[globalIdx] : 0.0f;
    __syncthreads();
    
    // 树形归约
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            sharedData[tid] += sharedData[tid + stride];
        }
        __syncthreads();
    }
    
    // 写入结果
    if (tid == 0) {
        output[blockIdx.x] = sharedData[0];
    }
}

// 避免 bank conflict 的归约（使用填充）
__global__ void bankConflictFreeReduce(float *input, float *output, int n) {
    extern __shared__ float sharedData[];
    
    int tid = threadIdx.x;
    int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 使用填充避免 bank conflict
    sharedData[tid] = (globalIdx < n) ? input[globalIdx] : 0.0f;
    __syncthreads();
    
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            sharedData[tid] += sharedData[tid + stride];
        }
        __syncthreads();
    }
    
    if (tid == 0) {
        output[blockIdx.x] = sharedData[0];
    }
}

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA error: %s\n", cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while(0)

int main() {
    int n = 1 << 20;
    size_t size = n * sizeof(float);
    
    float *h_input = (float *)malloc(size);
    float expectedSum = 0.0f;
    
    for (int i = 0; i < n; i++) {
        h_input[i] = 1.0f;
        expectedSum += 1.0f;
    }
    
    float *d_input, *d_output;
    CUDA_CHECK(cudaMalloc((void **)&d_input, size));
    
    int blocksPerGrid = 256;
    int threadsPerBlock = 256;
    size_t outputSize = blocksPerGrid * sizeof(float);
    CUDA_CHECK(cudaMalloc((void **)&d_output, outputSize));
    
    CUDA_CHECK(cudaMemcpy(d_input, h_input, size, cudaMemcpyHostToDevice));
    
    int sharedMemSize = threadsPerBlock * sizeof(float);
    
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));
    
    CUDA_CHECK(cudaEventRecord(start));
    sharedMemReduce<<<blocksPerGrid, threadsPerBlock, sharedMemSize>>>(
        d_input, d_output, n);
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    
    float reduceTime = 0;
    CUDA_CHECK(cudaEventElapsedTime(&reduceTime, start, stop));
    printf("Shared memory reduce time: %.3f ms\n", reduceTime);
    
    float *h_output = (float *)malloc(outputSize);
    CUDA_CHECK(cudaMemcpy(h_output, d_output, outputSize, cudaMemcpyDeviceToHost));
    
    float sum = 0.0f;
    for (int i = 0; i < blocksPerGrid; i++) {
        sum += h_output[i];
    }
    
    printf("Sum: %.0f (expected: %.0f)\n", sum, expectedSum);
    
    CUDA_CHECK(cudaFree(d_input));
    CUDA_CHECK(cudaFree(d_output));
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
    free(h_input);
    free(h_output);
    
    return 0;
}

4.3 常量内存性能测试

cpp

#include <stdio.h>
#include <cuda_runtime.h>

__constant__ float constCoeffs[256];

__global__ void useConstantMemory(float *output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float sum = 0.0f;
        for (int i = 0; i < 256; i++) {
            sum += constCoeffs[i];
        }
        output[idx] = sum;
    }
}

__global__ void useGlobalMemory(float *output, const float *coeffs, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float sum = 0.0f;
        for (int i = 0; i < 256; i++) {
            sum += coeffs[i];
        }
        output[idx] = sum;
    }
}

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA error: %s\n", cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while(0)

int main() {
    int n = 1 << 20;
    size_t size = n * sizeof(float);
    
    float h_coeffs[256];
    for (int i = 0; i < 256; i++) {
        h_coeffs[i] = 1.0f;
    }
    
    CUDA_CHECK(cudaMemcpyToSymbol(constCoeffs, h_coeffs, sizeof(h_coeffs)));
    
    float *d_coeffs;
    CUDA_CHECK(cudaMalloc((void **)&d_coeffs, sizeof(h_coeffs)));
    CUDA_CHECK(cudaMemcpy(d_coeffs, h_coeffs, sizeof(h_coeffs), cudaMemcpyHostToDevice));
    
    float *d_output;
    CUDA_CHECK(cudaMalloc((void **)&d_output, size));
    
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));
    
    // 测试常量内存
    CUDA_CHECK(cudaEventRecord(start));
    for (int iter = 0; iter < 100; iter++) {
        useConstantMemory<<<blocksPerGrid, threadsPerBlock>>>(d_output, n);
    }
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    
    float constTime = 0;
    CUDA_CHECK(cudaEventElapsedTime(&constTime, start, stop));
    printf("Constant memory time (100 iterations): %.3f ms\n", constTime);
    
    // 测试全局内存
    CUDA_CHECK(cudaEventRecord(start));
    for (int iter = 0; iter < 100; iter++) {
        useGlobalMemory<<<blocksPerGrid, threadsPerBlock>>>(d_output, d_coeffs, n);
    }
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    
    float globalTime = 0;
    CUDA_CHECK(cudaEventElapsedTime(&globalTime, start, stop));
    printf("Global memory time (100 iterations): %.3f ms\n", globalTime);
    
    printf("Speedup with constant memory: %.2fx\n", globalTime / constTime);
    
    CUDA_CHECK(cudaFree(d_coeffs));
    CUDA_CHECK(cudaFree(d_output));
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
    
    return 0;
}

4.4 统一内存示例

cpp

#include <stdio.h>
#include <cuda_runtime.h>

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

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA error: %s\n", cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while(0)

int main() {
    int n = 1 << 20;
    size_t size = n * sizeof(float);
    
    // 分配统一内存
    float *data;
    CUDA_CHECK(cudaMallocManaged(&data, size));
    
    // 主机端初始化
    for (int i = 0; i < n; i++) {
        data[i] = (float)i;
    }
    
    // 预取到设备
    int deviceId;
    CUDA_CHECK(cudaGetDevice(&deviceId));
    CUDA_CHECK(cudaMemPrefetchAsync(data, size, deviceId));
    
    // 启动内核
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));
    
    CUDA_CHECK(cudaEventRecord(start));
    unifiedMemKernel<<<blocksPerGrid, threadsPerBlock>>>(data, n);
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    
    float kernelTime = 0;
    CUDA_CHECK(cudaEventElapsedTime(&kernelTime, start, stop));
    printf("Kernel execution time: %.3f ms\n", kernelTime);
    
    // 预取回主机
    CUDA_CHECK(cudaMemPrefetchAsync(data, size, cudaCpuDeviceId));
    CUDA_CHECK(cudaDeviceSynchronize());
    
    // 验证结果
    bool correct = true;
    for (int i = 0; i < 10; i++) {
        float expected = (float)i * 2.0f + 1.0f;
        if (fabs(data[i] - expected) > 1e-5) {
            correct = false;
            break;
        }
    }
    printf("Result verification: %s\n", correct ? "PASSED" : "FAILED");
    
    CUDA_CHECK(cudaFree(data));
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
    
    return 0;
}

五、深入思考

5.1 历史背景

GPU 内存层次结构的演变反映了通用计算需求的增长。早期 GPU 只有简单的全局内存，随着 GPGPU 的发展，共享内存、常量内存等专用内存类型被引入，以支持更广泛的计算模式。

统一内存的出现是对编程复杂性的回应。早期的 CUDA 编程需要手动管理数据传输，容易出错且繁琐。统一内存通过自动化这一过程，降低了入门门槛。

5.2 前沿方向

当前的研究正在探索：

HBM（High Bandwidth Memory）：提供更高带宽的内存技术
NVLink：GPU 之间的高速互连，支持更大的有效内存
内存压缩：硬件级别的内存压缩以提高有效带宽
智能预取：基于机器学习的内存访问预测

5.3 延伸阅读

NVIDIA CUDA C Programming Guide - Memory Model 章节
"Programming Massively Parallel Processors" Chapter 5
NVIDIA Developer Blog 上的内存优化文章

六、实验笔记

实验 1：测量不同内存类型的带宽

目标：量化各种内存类型的实际带宽。

步骤：

编写内核分别测试全局内存、共享内存、常量内存的访问速度

使用 cudaEvent 精确计时

计算有效带宽（字节/秒）

与硬件规格对比

预期结果：共享内存带宽远高于全局内存，常量内存在广播访问时表现优异。

实验 2：分析 bank conflict 的影响

目标：观察 bank conflict 对共享内存性能的影响。

步骤：

编写两个归约内核：一个有 bank conflict，一个无 bank conflict

比较执行时间

使用 nsight compute 分析共享内存访问效率

思考题：在什么情况下 bank conflict 的影响最严重？

七、本章小结

本章详细介绍了 GPU 的内存层次结构。我们学习了各种内存类型的特性、使用方法和优化技巧。

关键要点包括：

全局内存容量大但延迟高，合并访问至关重要
共享内存速度快但容量有限，需注意 bank conflict
常量内存和纹理内存在特定访问模式下效率极高
寄存器是最快的存储，但数量有限
统一内存简化编程，但可能有性能开销

内存优化是 CUDA 性能调优的核心。在下一章中，我们将探讨线程同步机制，这是协调并行线程的基础。

八、习题

列出 GPU 的所有内存类型，并简述各自的特点。

解释什么是内存合并访问，以及为什么它对性能重要。

共享内存和全局内存的主要区别是什么？

什么是 bank conflict？如何避免？

常量内存在什么访问模式下效率最高？

编写一个内核，展示如何使用共享内存实现数据复用。

统一内存的优势和劣势分别是什么？

如果一个内核使用了过多的寄存器，会发生什么？

如何查询设备的内存带宽规格？

设计一个实验来比较合并访问和非合并访问的性能差异。