GPU 内存模型

理解 CUDA 内存层次结构对于优化光线追踪性能至关重要。

内存层次结构

内存类型

类型	容量	延迟	作用域	用途
Global	数 GB	400+ cycles	全局	场景数据、输出缓冲
L2 Cache	数 MB	~30 cycles	自动	缓存全局内存访问
Shared	48KB/SM	~5 cycles	块内	BVH 栈、临时数据
Local	256KB/SM	~5 cycles	线程	局部变量
Register	255/线程	1 cycle	线程	计算中间值

内存合并访问

全局内存访问应尽量合并，使同一 warp 的访问落在连续地址：

// 好：合并访问
float value = global_array[threadIdx.x];

// 坏：随机访问
float value = global_array[random_index[threadIdx.x]];

BVH 遍历内存优化

使用 Shared Memory 存储遍历栈：

__global__ void render_kernel(...) {
    // Shared Memory 存储 BVH 栈
    __shared__ int shared_stack[32][64];  // 32 threads, depth 64

    int* my_stack = shared_stack[threadIdx.x];
    // ... 遍历代码
}

场景数据布局

优化数据布局以提高缓存命中率：

// 结构体数组 (AoS) - 不推荐
struct Sphere {
    vec3 center;
    float radius;
    Material material;
};
Sphere spheres[N];

// 数组结构体 (SoA) - 推荐
struct SphereSoA {
    vec3 centers[N];
    float radii[N];
    Material materials[N];
};

内存使用分析

数据	类型	大小估算
球体数组	Global	N × 32 bytes
BVH 节点	Global	(2N-1) × 32 bytes
光线缓冲	Global	W×H × 32 bytes
输出图像	Global	W×H × 12 bytes
BVH 栈	Shared	32 × 64 × 4 bytes

参考资料

• [CUDA Programming Guide] NVIDIA
• [Aila & Karras 2010] "Understanding Ray Traversal Efficiency"