Ray Tracing in One Weekend in CUDA
RayTracingInOneWeekendCUDA 是 RayTracingInOneWeekend 的CUDA版本实现。
计算能力要大于等于6.0。 Compute capability >= 6.x
因为代码中需要使用
atomicAdd(float*, float)
针对浮点数的原子操作,所以需要GPU的计算能力大于等于6.0。
mkdir build
./run.sh
场景 | 分辨率 | spp | 最大深度 | CPU耗时(s) | GPU耗时(s) |
---|---|---|---|---|---|
场景1 | 800*450 | 100 | 50 | 2019.81 | 26.91 |
电脑配置:
- CPU: Intel(R) Core(TM) i5-10400F CPU @ 2.90GHz
- GPU: NVIDIA Corporation GP104 [GeForce GTX 1080] (rev a1)
- CUDA: 11.4
- nvidia Driver: 470.57.02
在实现时,本仓库CUDA代码尽可能和RayTracingInOneWeekend中的代码相似,其中不得不该的部分主要分为以下几点:
- 改
double
为float
,因为Current GPUs run fastest when they do calculations in single precision. Double precision calculations can be several times slower on some GPUs
; - 将需要在GPU端运行的函数前增加
__device__
函数修饰符号; - 将由CPU端调用,GPU端执行的函数从
类函数
改为全局函数
,并将其他需要在GPU端运行、不适合做为类函数
的函数改为全局函数
; - CUDA不支持智能指针,将C++版本中的智能指针改为普通指针;
- 改为在GPU端使用CUDA库产生随机数;
- 每个光线申请一个GPU线程,每个像素申请一个
vec3
类型的内存/显存空间存储该像素颜色值,单个像素可采样多跟光线,即总光线数=像素数*单像素采样数
。由于一像素需要采样多条光线,每条光线对应的线程都会读写同一个vec3
显存位置的数据,因此需要使用atomicAdd(float*,float)
原子操作;
1. Accelerated Ray Tracing in One Weekend in CUDA
2. 用CUDA加速Ray Tracing in One Weekend- 上
3. 用CUDA加速Ray Tracing in One Weekend- 下