文件结构
文件结构如下:
- hdr_loader.h : 用作外部文件的导入。
- main.cpp : 其中包括启动初始化cuda、glew、glfw、新建窗口、添加着色器、创建对于图形界面的各种操作(键盘、鼠标、滑轮等事件)、创建环境的材质、处理相机的移动等。
- volume_kernel.h: 其中定义了
Kernel_params
数据结构。 - volume_kernel.cu : 需要用到cuda来加速的整个工程的核心部分: 光追过程的具体细节。
Tips:注意此处对于.cu文件,需要对整个工程进行一些配置:
- 对工程中的链接器,在附加依赖项中,添加
cudart.lib
; - 对该工程设置生成依赖项,在生成自定义中选择对应的CUDA版本;
- 修改该.cu文件的属性,选择项类型为 CUDA C/C++;
代码理解
volume_kernel.h
struct Kernel_params {
// Display
uint2 resolution;
float exposure_scale;
unsigned int *display_buffer;
// Progressive rendering state
unsigned int iteration;
float3 *accum_buffer;
// Limit on path length
unsigned int max_interactions; //最大追踪次数
// Camera
float3 cam_pos; //相机位置
float3 cam_dir; //相机方向
float3 cam_right; //?
float3 cam_up; //?
float cam_focal; //?
//Environment
unsigned int environment_type;
cudaTextureObject_t env_tex;
// Volume definition
unsigned int volume_type;
float max_extinction; // 衰减系数 + 空碰撞系数 majorant
float albedo; // sigma / kappa
};
extern "C" __global__ void volume_rt_kernel(
const Kernel_params kernel_params);
该部分定义主要数据结构,该结构中包括有:display、渲染状态、光线追踪时的路径深度、相机、环境、体定义(volume definition)等相关的参数。
volume_kernel.cu
函数列表
//3d vector math
__device__ inline float3 operator+(const float3& a, const float3& b)
__device__ inline float3 operator-(const float3& a, const float3& b)
__device__ inline float3 operator*(const float3& a, const float s)
__device__ inline float3 operator/(const float3& a, const float s)
__device__ inline void operator+=(float3& a, const float3& b)
__device__ inline void operator*=(float3& a, const float& s)
__device__ inline float3 normalize(const float3 &d)
//core function
// 判断是否与体边框相交
__device__ inline bool intersect_volume_box(
float &tmin, const float3 &raypos, const float3 &raydir)
// 判断是否还在体中
__device__ inline bool in_volume(
const float3 &pos)
// 获得衰减系数(?)
__device__ inline float get_extinction(
const Kernel_params &kernel_params,
const float3 &p)
// delta tracking 过程;对距离采样
__device__ inline bool sample_interaction(
Rand_state &rand_state,
float3 &ray_pos,
const float3 &ray_dir,
const Kernel_params &kernel_param)
// 包含了光线在体中的一些过程:是否超过追踪最大深度;用轮盘赌决定是否结束光线;如果是散射情况,采样下一个散射方向
__device__ inline float3 trace_volume(
Rand_state &rand_state,
float3 &ray_pos,
float3 &ray_dir,
const Kernel_params &kernel_params)
Delta tracking 过程
__device__ inline bool sample_interaction(
Rand_state &rand_state,
float3 &ray_pos,
const float3 &ray_dir,
const Kernel_params &kernel_param)
{
float t = 0.0f;
float3 pos;
do {
t -= logf(1.0f - rand(&rand_state)) / kernel_params.max_extinction;
pos = ray_pos + ray_dir * t;
if (!in_volume(pos))
return false;
} while (get_extinction(kernel_params, pos) < rand(&rand_state) * kernel_params.max_extinction);
ray_pos = pos;
return true;
}
对相位函数(phase function)采样
- 注意kernel函数是需要进行封装才能添加extern “C”来进行声明并调用的,cpp文件中是不能直接调用 __global__函数的,因为编译器是无法解析符号<<<……>>>以及blockIdx、threadIdx等,因此 global 函数只能放在.cu文件函数里面定义和被调用
- 用
extern “C”
修饰的原因,是因为.cu是扩展C,但由于cuda-c只是部分属于c++,所以需要更改工程的配置属性,将.cu文件定义为类C语言编译(NVCC.exe编译)。 - 在.cpp中是不能include
.cu
文件的,就好比cpp中不能include.cpp
一样。
REF
CUDA学习笔记四