Games101作业7 路径追踪 CUDA

没有杰士邦

已于 2023-06-08 15:17:17 修改

阅读量293

点赞数 1

文章标签： c++ 图形渲染

于 2023-06-08 13:56:35 首次发布

本文链接：https://blog.csdn.net/qq_43636340/article/details/131097811

版权

本文介绍了一篇关于将路径追踪作业从CPU迁移到CUDA并进行优化的过程。作者首先概述了配置和基线性能，然后详细讨论了CUDA化的关键点，包括递归转迭代、线程组织形式的优化，以及使用共享内存和引用传递提高性能。通过这些优化，将原本的渲染时间从30秒降低到约1秒，显著提升了效率。

摘要由CSDN通过智能技术生成

配置

i7-11800h
笔记本版的 RTX3070
CUDA 11.6 vs2019

Baseline

基线为多线程版，我参考的这两个博客：
GAMES101作业7-路径追踪实现过程&代码框架超全解读
 GAMES101作业7及课程总结（重点实现多线程加速，微表面模型材质）
在768 $\times$ 768 $\times$ 16(spp)的设置下，8线程release模式渲染时间为16s

CUDA化

CUDA基础就不介绍了，我也只是了解的程度，迁移这个作业代码到CUDA，注意这几个要点就行：

所有渲染相关函数，包括类成员函数用__device__修饰，其他初始化的函数，如BVH构建过程在主机端跑就行（__host__修饰）
类里面动态分配的内存需要单独用cudaMemcpy拷贝，不能直接拷贝类的对象
Triangle，MeshTriangle这两个类的初始化要在设备端进行，不能直接cudaMemcpy拷贝到设备端，这两个类重写了虚函数，而设备端和主机端编译出来的虚表地址是不一样的。

__global__ void cpy2device(Object* d_object_ptr_)
{
   
   new (d_object_ptr_) MeshTriangle(*(MeshTriangle*)d_object_ptr_);
}
__global__ void create_d_triangles(Triangle* h_triangle_data, Object** d_triangles, int numObjects)
{
   
   int index = blockIdx.x * blockDim.x + threadIdx.x;
   if (index < numObjects) {
   
       d_triangles[index] = new Triangle(h_triangle_data[index]);
   }
}

由于设备端不支持使用STL容器，MeshTriangle这个类用于存储Triangle的vector需要动态初始化，即成员变量保留一个vector指针，然后在构造函数中new，否则在设备端调用MeshTriangle构造函数时会调用不支持的std::vector的构造函数

 std::vector<Triangle> *triangles = nullptr;
 std::vector<Object*> *d_triangles_h_ptr = nullptr;
 Object** d_triangles;

由于我BVH的构建函数recursiveBuild保留了作业本身的版本，也就是在主机端进行的，BVH构建后，BVH叶节点指向的Object对象也是在主机端对象，因此需要一个额外的更新过程来保证BVH加速结构在设备端可用，为实现这一点，还需要将原来的指针型二叉树改为线性存储的二叉树

struct BVHBuildNode {
   
   Bounds3 bounds;
   int leftChildIndex;
   int rightChildIndex;
   Object* object;
   float area;
   __host__ __device__ BVHBuildNode()
       :bounds(Bounds3()),
       leftChildIndex(-1), rightChildIndex(-1),
   object(nullptr), area(0)
   {
   }
};
__host__ void BVHAccel::updateDeviceNodePointer(BVHBuildNode& node, std::vector<Object*>& d_primitives) {
   

   // 如果当前节点是叶节点，则更新 object 指针
   if (node.leftChildIndex == -1 && node.rightChildIndex == -1) {
   
       size_t index = std::distance(primitives.begin(), std::find(primitives.begin(), primitives.end(), node.object));
       // 更新指针为设备内存中的对象
       node.object = d_primitives[index];
   }
   // 递归更新子节点的指针
   if (node.leftChildIndex != -1)
       this->updateDeviceNodePointer(nodes[node.leftChildIndex], d_primitives);
   if (node.rightChildIndex != -1)
       this->updateDeviceNodePointer(nodes[node.rightChildIndex], d_primitives);
};

优化

递归->迭代

如果只是像上面那样，保留源代码框架的算法逻辑放到GPU上跑，效率是很低的，我第一次改完代码渲染时间在30s（同样是768 $\times$ 768 $\times$ 16spp的设置）左右，接近CPU端的两倍了。关于这点，如果在性能探查器看原算法的CPU占用，就会发现我们的光追函数castRay中的光线求交和非直接光照的递归占了很多：
请添加图片描述

CUDA核函数虽然现在支持递归调用，但函数调用需要保存很多临时变量，返回地址，寄存器状态什么的，所以递归对GPU上非常宝贵的寄存器资源很不友好，如果寄存器被迅速耗尽，我们在算法执行中就需要频繁对全局内存进行读写，这通常需要几十到几百个时钟周期，相比只需要几个时钟周期的寄存器读写是非常费时间的