|
|
文礼：专栏文章目录
上一篇我们用Nvidia的Optix绘制了一副只有背景的光追图。它实际上是在Optix的RayGen阶段，也就是发射“光线”的阶段直接返回一个固定的颜色作为光追的结果。
而实际的光追，是要通过求光线与场景的交点，然后计算此交点处的颜色来实现的。这里面包含三个大的方面：
其中Optix通过Nvidia硬件所提供的高效加速部分，是上面的第二步：计算光线与场景的交点。但是为了高效地实现这个求交，Optix对求交所用到的数据格式进行了严格的定义。并且将这些数据尽可能地组织在GPU板载内存（片上存储+显存）空间，以及GPU自身的寄存器当中，利用GPU的高带宽，来加速这个过程。
其实在我们之前的CUDA版本的路径追踪器当中，我们也是将整个BVH结构通过CUDA命令拷贝到了显存当中，并且将相关的求交计算通过CUDA编译成为GPU端的代码，在GPU端并列执行。然而，无论是我们的BVH数据组织结构，还是相应的求交代码路径，其实都是平移的CPU端的设计，并没有将其针对GPU端的访存和执行特点进行优化。
Optix则是将这一部分完全以在GPU端执行为前提，进行了高度的数据结构和算法优化。代价是，我们无法直接使用我们在CPU端所写的数据结构，而是要转而使用Optix内置的结构。
比如，我们原本在CPU端构建场景的代码如下：
// World
auto random_scene() {
My::HitableList<float_precision> world;
auto material_ground = std::make_shared<lambertian>(color({0.5, 0.5, 0.5}));
world.add(std::make_shared<My::Sphere<float_precision, std::shared_ptr<material>>>(
(float_precision)1000.0, point3({0, -1000, -1.0}), material_ground));
for (int a = -11; a < 11; a++) {
for (int b = -11; b < 11; b++) {
auto choose_mat = My::random_f<float_precision>();
point3 center(
{a + (float_precision)0.9 * My::random_f<float_precision>(),
(float_precision)0.2,
b + (float_precision)0.9 * My::random_f<float_precision>()});
if (Length(center - point3({4, 0.2, 0})) > 0.9) {
std::shared_ptr<material> sphere_material;
if (choose_mat < 0.8) {
// diffuse
color albedo = My::random_v<float_precision, 3>() *
My::random_v<float_precision, 3>();
sphere_material = std::make_shared<lambertian>(albedo);
world.add(
std::make_shared<My::Sphere<float_precision, std::shared_ptr<material>>>(
(float_precision)0.2, center, sphere_material));
} else if (choose_mat < 0.95) {
// metal
auto albedo = My::random_v<float_precision, 3>(0.5, 1);
auto fuzz = My::random_f<float_precision>(0, 0.5);
sphere_material = std::make_shared<metal>(albedo, fuzz);
world.add(
std::make_shared<My::Sphere<float_precision, std::shared_ptr<material>>>(
(float_precision)0.2, center, sphere_material));
} else {
// glass
sphere_material = std::make_shared<dielectric>((float_precision)1.5);
world.add(
std::make_shared<My::Sphere<float_precision, std::shared_ptr<material>>>(
(float_precision)0.2, center, sphere_material));
}
}
}
}
auto material_1 = std::make_shared<dielectric>((float_precision)1.5);
world.add(std::make_shared<My::Sphere<float_precision, std::shared_ptr<material>>>(
(float_precision)1.0, point3({0, 1, 0}), material_1));
auto material_2 = std::make_shared<lambertian>(color({0.4, 0.2, 0.1}));
world.add(std::make_shared<My::Sphere<float_precision, std::shared_ptr<material>>>(
(float_precision)1.0, point3({-4, 1, 0}), material_2));
auto material_3 = std::make_shared<metal>(color({0.7, 0.6, 0.5}), (float_precision)0.0);
world.add(std::make_shared<My::Sphere<float_precision, std::shared_ptr<material>>>(
(float_precision)1.0, point3({4, 1, 0}), material_3));
return world;
}
这段代码可以原封不动地用于我们的CPU版本和CUDA版本的路径追踪器，但是对于Optix版本的路径追踪器，我们就必须如下重写，以使用Optix内置的数据结构：
// accel handling
constexpr int scene_obj_num = 22 * 22 + 1 + 3;
OptixTraversableHandle gas_handle;
CUdeviceptr d_gas_output_buffer;
{
std::array<float3, scene_obj_num> sphereVertex;
std::array<float, scene_obj_num> sphereRadius;
std::array<uint16_t, scene_obj_num> g_mat_indices;
int index = 0;
// ground
g_mat_indices[index] = index;
sphereVertex[index] = make_float3(0, -1000, -1);
sphereRadius[index++] = 1000.f;
// random spheres
for (int a = -11; a < 11; a++) {
for (int b = -11; b < 11; b++) {
g_mat_indices[index] = index;
sphereVertex[index] = make_float3(a + 0.9 * RND, 0.2f, b + 0.9 * RND);
sphereRadius[index++] = 0.2f;
}
}
// three big balls
g_mat_indices[index] = index;
sphereVertex[index] = make_float3(0, 1, 0);
sphereRadius[index++] = 1.0f;
g_mat_indices[index] = index;
sphereVertex[index] = make_float3(-4, 1, 0);
sphereRadius[index++] = 1.0f;
g_mat_indices[index] = index;
sphereVertex[index] = make_float3(4, 1, 0);
sphereRadius[index++] = 1.0f;
assert(index == scene_obj_num);
size_t sphereVertexSize = sizeof(sphereVertex[0]) * sphereVertex.size();
size_t sphereRadiusSize = sizeof(sphereRadius[0]) * sphereRadius.size();
CUdeviceptr d_vertex_buffer;
checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_vertex_buffer), sphereVertexSize));
checkCudaErrors(cudaMemcpy(reinterpret_cast<void *>(d_vertex_buffer), sphereVertex.data(), sphereVertexSize, cudaMemcpyHostToDevice));
CUdeviceptr d_radius_buffer;
checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_radius_buffer), sphereRadiusSize));
checkCudaErrors(cudaMemcpy(reinterpret_cast<void *>(d_radius_buffer), sphereRadius.data(), sphereRadiusSize, cudaMemcpyHostToDevice));
CUdeviceptr d_mat_indices;
const size_t mat_indices_size_in_bytes = g_mat_indices.size() * sizeof(uint16_t);
checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_mat_indices), mat_indices_size_in_bytes));
checkCudaErrors(cudaMemcpy(
reinterpret_cast<void *>(d_mat_indices),
g_mat_indices.data(),
mat_indices_size_in_bytes,
cudaMemcpyHostToDevice
));
std::array<uint32_t, scene_obj_num> sphere_input_flags;
sphere_input_flags.fill(OPTIX_GEOMETRY_FLAG_DISABLE_ANYHIT);
OptixBuildInput sphere_input = {};
sphere_input.type = OPTIX_BUILD_INPUT_TYPE_SPHERES;
sphere_input.sphereArray.vertexBuffers = &d_vertex_buffer;
sphere_input.sphereArray.numVertices = sphereVertex.size();
sphere_input.sphereArray.radiusBuffers = &d_radius_buffer;
sphere_input.sphereArray.flags = sphere_input_flags.data();
sphere_input.sphereArray.numSbtRecords = scene_obj_num;
sphere_input.sphereArray.sbtIndexOffsetBuffer = d_mat_indices;
sphere_input.sphereArray.sbtIndexOffsetSizeInBytes = sizeof(uint16_t);
sphere_input.sphereArray.sbtIndexOffsetStrideInBytes= sizeof(uint16_t);
OptixAccelBuildOptions accel_options = {};
accel_options.buildFlags = OPTIX_BUILD_FLAG_ALLOW_COMPACTION | OPTIX_BUILD_FLAG_ALLOW_RANDOM_VERTEX_ACCESS;
accel_options.operation = OPTIX_BUILD_OPERATION_BUILD;
OptixAccelBufferSizes gas_buffer_sizes;
checkOptiXErrors(optixAccelComputeMemoryUsage(context, &accel_options, &sphere_input, 1, &gas_buffer_sizes));
CUdeviceptr d_temp_buffer_gas;
checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_temp_buffer_gas), gas_buffer_sizes.tempSizeInBytes));
// non-compacted output
CUdeviceptr d_buffer_temp_output_gas_and_compacted_size;
size_t compactedSizeOffset = My::roundUp<size_t>(gas_buffer_sizes.outputSizeInBytes, 8ull);
checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_buffer_temp_output_gas_and_compacted_size), compactedSizeOffset + 8));
OptixAccelEmitDesc emitProperty = {};
emitProperty.type = OPTIX_PROPERTY_TYPE_COMPACTED_SIZE;
emitProperty.result = (CUdeviceptr)((char*)d_buffer_temp_output_gas_and_compacted_size + compactedSizeOffset);
checkOptiXErrors(optixAccelBuild(context,
0, // CUDA stream
&accel_options,
&sphere_input,
1, // num build inputs
d_temp_buffer_gas, gas_buffer_sizes.tempSizeInBytes,
d_buffer_temp_output_gas_and_compacted_size, gas_buffer_sizes.outputSizeInBytes,
&gas_handle,
&emitProperty,
1));
d_gas_output_buffer = d_buffer_temp_output_gas_and_compacted_size;
checkCudaErrors(cudaFree((void *)d_temp_buffer_gas));
checkCudaErrors(cudaFree((void *)d_mat_indices));
checkCudaErrors(cudaFree((void *)d_vertex_buffer));
checkCudaErrors(cudaFree((void *)d_radius_buffer));
size_t compacted_gas_size;
checkCudaErrors(cudaMemcpy(&compacted_gas_size, (void *)emitProperty.result, sizeof(size_t), cudaMemcpyDeviceToHost));
if(compacted_gas_size < gas_buffer_sizes.outputSizeInBytes) {
checkCudaErrors(cudaMalloc(reinterpret_cast<void **>(&d_gas_output_buffer), compacted_gas_size));
// use handle as input and output
checkOptiXErrors(optixAccelCompact(context, 0, gas_handle, d_gas_output_buffer, compacted_gas_size, &gas_handle));
checkCudaErrors(cudaFree((void *)d_buffer_temp_output_gas_and_compacted_size));
} else {
d_gas_output_buffer = d_buffer_temp_output_gas_and_compacted_size;
}
}
首先可以看到，我们使用了Optix内置的，sphereArray的数据结构，来存储我们的488个球，并构建加速结构GAS。这个GAS数据结构的细节，Nvidia对其进行了隐藏。但是从构建过程可以看出，GAS当中BVH树的本体，实际上应该只是存储了包围盒。然后每个节点（对应一个包围盒）带一个payload，包含下面三个字段：
而圆心坐标、半径、材质都是分别存储在另外开设的缓冲区当中（同样位于显存）。各个索引指向了对应数据在这些缓冲区当中的位置。
这其实就是一种典型的SoA的数据组织方式。我们将多个相同类型的对象的同一个属性的值组织在地址连续的空间当中。因为GPU的并行执行方式决定了，在一个线程组（wavefront）当中的所有线程在某个特定时刻总是执行着相同的命令。所以如果1号线程在访问1号圆的圆心，那么2号线程也一定是在访问2号圆的圆心。所以对于GPU来说，SoA的数据组织形式往往比AoS的形式在访问效率方面要高效得多。因为缓存效率会高很多。
好的，现在我们完成了场景的构建。Optix的内置框架代码会处理对于整个BVH结构的遍历以及求交。在找到交点时，Optix会按照之前文章当中引用的那张图，调取相应的shader进行交点色的计算。如果没有找到交点，Optix也会调用对应的shader进行背景的计算。
不过Optix虽然提供了这个光线和场景求交的框架，但是在这个过程当中所需要传递的一些参数，比如用以返回交点处颜色的变量，当光线在交点处发生反射和折射时，反射光线和折射光线的参数等等，都是需要我们自己定义的。
或者比照之前CPU和CUDA版本的路径追踪器来说的话，我们需要按照Optix的方式，重新定义Ray和Hit，这两个数据结构。
因为无论是Ray还是Hit其实都是属于计算迭代过程当中的中间变量，我们并没有永久存储它们的需要。Optix当中是直接采用了寄存器的方式来保存这些参数。而我们需要在调用真正实现光追的optixTrace方法之前，完成对于这些参数的展开和重组。
static __forceinline__ __device__ void trace(
OptixTraversableHandle handle,
vec3 ray_origin,
vec3 ray_direction,
float tmin,
float tmax,
MyTracePayload& prd ) {
unsigned int p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10;
p0 = __float_as_uint(prd.attenuation[0]);
p1 = __float_as_uint(prd.attenuation[1]);
p2 = __float_as_uint(prd.attenuation[2]);
p3 = __float_as_uint(prd.scatter_ray_origin[0]);
p4 = __float_as_uint(prd.scatter_ray_origin[1]);
p5 = __float_as_uint(prd.scatter_ray_origin[2]);
p6 = __float_as_uint(prd.scatter_ray_direction[0]);
p7 = __float_as_uint(prd.scatter_ray_direction[1]);
p8 = __float_as_uint(prd.scatter_ray_direction[2]);
p9 = prd.max_depth;
p10 = (prd.done) ? 1 : 0;
optixTrace(
handle,
_f(ray_origin),
_f(ray_direction),
tmin,
tmax,
0.0f, // rayTime
OptixVisibilityMask( 1 ),
OPTIX_RAY_FLAG_NONE,
0, // SBT offset
0, // SBT stride
0, // missSBTIndex
p0, p1, p2,
p3, p4, p5,
p6, p7, p8,
p9, p10 );
prd.attenuation[0] = __uint_as_float(p0);
prd.attenuation[1] = __uint_as_float(p1);
prd.attenuation[2] = __uint_as_float(p2);
prd.scatter_ray_origin[0] = __uint_as_float(p3);
prd.scatter_ray_origin[1] = __uint_as_float(p4);
prd.scatter_ray_origin[2] = __uint_as_float(p5);
prd.scatter_ray_direction[0] = __uint_as_float(p6);
prd.scatter_ray_direction[1] = __uint_as_float(p7);
prd.scatter_ray_direction[2] = __uint_as_float(p8);
prd.max_depth = p9;
prd.done = (p10 == 1) ? true : false;
}
在上面这个例子当中，我们一共是使用了11个32bit的寄存器，来“热”跟踪我们光追所需的参数。由于这些寄存器都被OptiX声明为无符号整数类型，所以我们也会需要一些简单的“小魔法”，将其它数据类型转为无符号整数类型。注意这里的类型转换仅仅是C++语义层面上的。我们实际上对于数据是没有进行任何的操作。
另外一个要点就是，optixTrace当中再次调用optixTrace，这个是有很大限制的。之前我们写CUDA的路径追踪程序也是已经遇到过类似的问题。因为GPU当中并没有如同CPU那样的调用堆栈，所以递归算法总的来说在GPU上是不适用的。
好在路径追踪的过程当中因为蒙特卡洛方法的导入，对于每一个像素我们都会发射比较多的&#34;光线&#34;。那么，对于每一条光线，我们只需要让其在每个交点，按照交点处的表面光学特性，在满足其分布的情况（BRDF和BSDF）下，随机地选取一个方向跟下去就可以了。那这样的话，利用光路的可逆性，对于每一条光线，我们只需要求得最近的交点，然后按此处的光学特性对光线进行吸收（降低其能量）并按照BRDF/BSDF计算出反射光线/折射光线的方向，然后沿着这条新的光线继续追踪，如此循环反复，并把结果不断用乘法混合就好了。
这样我们就完成了第二步，光线与场景的求交以及BVH的迭代。
最后就是决定交点的颜色。这是通过3种类型的shader来完成的。
首先是上一篇也用过的，发射光线的shader。这个shader用来根据屏幕像素的位置，生成&#34;光线&#34;并启动路径追踪。同时，负责合成光线路径上每次反射/折射的结果：
extern &#34;C&#34;
__global__ void __raygen__rg() {
uint3 launch_index = optixGetLaunchIndex();
// RayGenData* rtData = (RayGenData*)optixGetSbtDataPointer();
unsigned int i = launch_index.x;
unsigned int j = launch_index.y;
unsigned int pixel_index = j * params.image->Width + i;
curandStateMRG32k3a* local_rand_state = &params.rand_state[pixel_index];
int num_of_samples = params.num_of_samples;
vec3 col = {0.f, 0.f, 0.f};
for (int s = 0; s < num_of_samples; s++) {
float u = float(i + curand_uniform(local_rand_state)) / params.image->Width;
float v = float(j + curand_uniform(local_rand_state)) / params.image->Height;
ray r = params.cam->get_ray(u, v, local_rand_state);
vec3 attenuation = {1.f, 1.f, 1.f};
MyTracePayload payload;
payload.attenuation = attenuation;
payload.scatter_ray_origin = r.getOrigin();
payload.scatter_ray_direction = r.getDirection();
payload.max_depth = params.max_depth;
payload.done = false;
do {
trace( params.handle,
payload.scatter_ray_origin,
payload.scatter_ray_direction,
0.001f, // tmin
FLT_MAX, // tmax
payload);
attenuation = attenuation * payload.attenuation;
} while (!payload.done);
col += attenuation;
}
col = col / (float)num_of_samples;
vec3* pOutputBuffer = reinterpret_cast<vec3*>(params.image->data);
pOutputBuffer[pixel_index] = My::Linear2SRGB(col);
}
而对于交点处的计算，称为closehit shader，则是一个颇为标准的BRDF/BSDF的蒙特卡洛采样计算过程：
extern &#34;C&#34;
__global__ void __closesthit__ch() {
auto payload = getPayload();
if(payload.max_depth <= 0) {
setPayload({0.f, 0.f, 0.f}, false);
return;
}
uint3 launch_index = optixGetLaunchIndex();
float t_hit = optixGetRayTmax();
unsigned int i = launch_index.x;
unsigned int j = launch_index.y;
unsigned int pixel_index = j * params.image->Width + i;
curandStateMRG32k3a* local_rand_state = &params.rand_state[pixel_index];
const vec3 ray_orig = _V(optixGetWorldRayOrigin());
const vec3 ray_dir = _V(optixGetWorldRayDirection());
const unsigned int prim_idx = optixGetPrimitiveIndex();
const OptixTraversableHandle gas = optixGetGASTraversableHandle();
const unsigned int sbtGASIndex = optixGetSbtGASIndex();
float4 q;
// sphere center (q.x, q.y, q.z), sphere radius q.w
optixGetSphereData(gas, prim_idx, 0, 0.f, &q); // don&#39;t know why but set sbtGASIndex cause all-zero q when sbtGASIndex = 2
vec3 world_raypos = ray_orig + t_hit * ray_dir;
vec3 obj_raypos = _V(optixTransformPointFromWorldToObjectSpace(_f(world_raypos)));
vec3 obj_normal = (obj_raypos - _V({q.x, q.y, q.z})) / q.w;
vec3 world_normal = _V(optixTransformNormalFromObjectToWorldSpace(_f(obj_normal)));
My::Normalize(world_normal);
ray scattered;
ray r_in(ray_orig, ray_dir);
hit_record rec;
rec.set(t_hit, world_raypos, world_normal, optixIsFrontFaceHit(), nullptr);
color attenuation;
HitGroupData* hg_data = reinterpret_cast<HitGroupData*>(optixGetSbtDataPointer() + sbtGASIndex * sizeof(HitGroupSbtRecord));
bool b_scattered;
switch(hg_data->material_type) {
case Material::MAT_DIFFUSE:
{
b_scattered = lambertian::scatter_static(r_in, rec, attenuation, scattered, local_rand_state, hg_data->base_color);
}
break;
case Material::MAT_METAL:
{
b_scattered = metal::scatter_static(r_in, rec, attenuation, scattered, local_rand_state, hg_data->base_color, hg_data->fuzz);
}
break;
case Material::MAT_DIELECTRIC:
{
b_scattered = dielectric::scatter_static(r_in, rec, attenuation, scattered, local_rand_state, hg_data->ir);
}
break;
}
setPayload(attenuation, scattered, b_scattered, payload.max_depth - 1);
}
最后，miss shader，负责光线未能与场景有交的时候的颜色计算（也就是背景色的计算）
extern &#34;C&#34;
__global__ void __miss__ms() {
MissData* missData = reinterpret_cast<MissData*>(optixGetSbtDataPointer());
vec3 unit_direction = _V(optixGetWorldRayDirection());
float t = 0.5f * (unit_direction[1] + 1.0f);
vec3 c = (1.0f - t) * color({1, 1, 1}) + t * (missData->bg_color);
setPayload(c, false);
}
代码
Test/OptixTest.cu · Chen, Wenli/GameEngineFromScratch - 码云 - 开源中国 (gitee.com)
Asset/Shaders/CUDA/OptixTest.shader.cu · Chen, Wenli/GameEngineFromScratch - 码云 - 开源中国 (gitee.com) |
|
发表于 2023-3-9 02:57:57