OptiX入门（二）

发布于 396 天前

本人初学者，如有错误和更好的表述，请指出

这次我们看optixTriangle程序

同样分为optixTriangle.cpp、optixTriangle.h、optixTriangle.cu文件

optixTriangle.h文件

其中包含的是参数结构Params和sbt的record信息。

这里注意到有一个hitgroup，这个可以认为是intersection、any hit、closest hit的集合，主要加速编译和光追执行速度。

struct Params
{
    uchar4*                image; //一维数组图像
    unsigned int           image_width; //图像宽
    unsigned int           image_height; //图像高
    float3                 cam_eye; //摄像机位置
    float3                 cam_u, cam_v, cam_w; //以摄像机为原点的笛卡尔坐标轴
    OptixTraversableHandle handle; //用于遍历时加速的句柄，就当做一个参数就好
};

optixTriangle.cpp文件

与optixHello.cpp相近的内容就不写了。

创建context

创建加速结构Acceleration structures

// accel handling
OptixTraversableHandle gas_handle;
CUdeviceptr            d_gas_output_buffer;
{
    // Use default options for simplicity.  In a real use case we would want to
    // enable compaction, etc
    OptixAccelBuildOptions accel_options = {};
    accel_options.buildFlags = OPTIX_BUILD_FLAG_NONE;
    accel_options.operation  = OPTIX_BUILD_OPERATION_BUILD;

    // Triangle build input: simple list of three vertices
    const std::array<float3, 3> vertices =
    { {
          { -0.5f, -0.5f, 0.0f },
          {  0.5f, -0.5f, 0.0f },
          {  0.0f,  0.5f, 0.0f }
    } }; //三角形三个顶点坐标

    const size_t vertices_size = sizeof( float3 )*vertices.size();
    CUdeviceptr d_vertices=0; //分配GPU内存空间vertexbuffer并复制数据
    CUDA_CHECK( cudaMalloc( reinterpret_cast<void**>( &d_vertices ), vertices_size ) );
    CUDA_CHECK( cudaMemcpy(
                reinterpret_cast<void*>( d_vertices ),
                vertices.data(),
                vertices_size,
                cudaMemcpyHostToDevice
                ) );

    // Our build input is a simple list of non-indexed triangle vertices
    const uint32_t triangle_input_flags[1] = { OPTIX_GEOMETRY_FLAG_NONE };
    OptixBuildInput triangle_input = {};
    triangle_input.type                        = OPTIX_BUILD_INPUT_TYPE_TRIANGLES;
    triangle_input.triangleArray.vertexFormat  = OPTIX_VERTEX_FORMAT_FLOAT3;
    triangle_input.triangleArray.numVertices   = static_cast<uint32_t>( vertices.size() );
    triangle_input.triangleArray.vertexBuffers = &d_vertices;
    triangle_input.triangleArray.flags         = triangle_input_flags;
    triangle_input.triangleArray.numSbtRecords = 1;

    OptixAccelBufferSizes gas_buffer_sizes; //加速结构所用到的buffer
    OPTIX_CHECK( optixAccelComputeMemoryUsage( //计算这个buffer所需要的大小
                context,
                &accel_options,
                &triangle_input,
                1, // Number of build inputs
                &gas_buffer_sizes
                ) );
    CUdeviceptr d_temp_buffer_gas; //构建加速结构临时所需要的buffer空间
    CUDA_CHECK( cudaMalloc(
                reinterpret_cast<void**>( &d_temp_buffer_gas ),
                gas_buffer_sizes.tempSizeInBytes
                ) );
    CUDA_CHECK( cudaMalloc(
                reinterpret_cast<void**>( &d_gas_output_buffer ),
                gas_buffer_sizes.outputSizeInBytes
                ) );

    OPTIX_CHECK( optixAccelBuild( //创建加速结构
                context,
                0,                  // CUDA stream
                &accel_options,
                &triangle_input,
                1,                  // num build inputs
                d_temp_buffer_gas,
                gas_buffer_sizes.tempSizeInBytes,
                d_gas_output_buffer,
                gas_buffer_sizes.outputSizeInBytes,
                &gas_handle,
                nullptr,            // emitted property list
                0                   // num emitted properties
                ) );

    // We can now free the scratch space buffer used during build and the vertex
    // inputs, since they are not needed by our trivial shading method
    CUDA_CHECK( cudaFree( reinterpret_cast<void*>( d_temp_buffer_gas ) ) );
    CUDA_CHECK( cudaFree( reinterpret_cast<void*>( d_vertices        ) ) );
}

创建module

注意到这里的module两个参数不一样。

pipeline_compile_options.numPayloadValues      = 3;
pipeline_compile_options.numAttributeValues    = 3;

payload是optixTrace和程序之间传递的信息。

attribute是intersection与any hit/closest-hit之间传递的信息。

这里的3意味着是3个32 bit的数据。

程序使用了payload，因此numPayloadValues不可以变，而numAttributeValues我试着改成2也可以（注意numAttributeValues最小为2）。

创建program groups

创建raygen、miss、hitgroup三个ProgramGroup。

hitgroup创建时有

hitgroup_prog_group_desc.hitgroup.moduleCH            = module;

这里moduleCH的CH就是closest-hit。

创建pipeline

创建shader binding table

这里往miss record中设置了背景色。

launch

设置参数后调用了camera的UVWFrame函数，点进去可以看到

void Camera::UVWFrame(float3& U, float3& V, float3& W) const
{
    W = m_lookat - m_eye; // Do not normalize W -- it implies focal length
    float wlen = length(W);
    U = normalize(cross(W, m_up));
    V = normalize(cross(U, W));

    float vlen = wlen * tanf(0.5f * m_fovY * M_PIf / 180.0f);
    V *= vlen;
    float ulen = vlen * m_aspectRatio;
    U *= ulen;
}

W是摄像机朝向，U是摄像机的正上方，UVW形成了一个三维笛卡尔直角坐标系。

显示结果

清理资源

optixTriangle.cu文件

看raygen函数。

extern "C" __global__ void __raygen__rg()
{
    // Lookup our location within the launch grid
    const uint3 idx = optixGetLaunchIndex(); //获取当前的pixel坐标
    const uint3 dim = optixGetLaunchDimensions(); //获取整个场景的坐标
    // Map our launch idx to a screen location and create a ray from the camera
    // location through the screen
    float3 ray_origin, ray_direction;
    computeRay( idx, dim, ray_origin, ray_direction ); //计算光线起点和朝向

    // Trace the ray against our scene hierarchy
    unsigned int p0, p1, p2;
    optixTrace(
            params.handle,
            ray_origin,
            ray_direction,
            0.0f,                // Min intersection distance
            1e16f,               // Max intersection distance
            0.0f,                // rayTime -- used for motion blur
            OptixVisibilityMask( 255 ), // Specify always visible
            OPTIX_RAY_FLAG_NONE,
            0,                   // SBT offset   -- See SBT discussion
            1,                   // SBT stride   -- See SBT discussion
            0,                   // missSBTIndex -- See SBT discussion
            p0, p1, p2 );
    float3 result;
    result.x = __uint_as_float( p0 );
    result.y = __uint_as_float( p1 );
    result.z = __uint_as_float( p2 );

    // Record results in our output raster
    params.image[idx.y * params.image_width + idx.x] = make_color( result ); 
}

这个optixGetLaunchIndex和optixGetLaunchDimensions不知道是什么意思，我们可以打印出来看看结果。

printf("%d %d %d %d %d %d\n",idx.x,idx.y,idx.z,dim.x,dim.y,dim.z);

截取部分结果分析。

691 41 0 1024 768 1
688 42 0 1024 768 1
692 32 0 1024 768 1
693 32 0 1024 768 1

1024和768分别是整个窗口的宽高，那么可以理解这个idx就是当前的pixel，而且是并行执行的，因为坐标在跳变，而dim就是整个窗口的长宽高。

那我们就可以分析下computeRay函数了。

static __forceinline__ __device__ void computeRay( uint3 idx, uint3 dim, float3& origin, float3& direction ) //device指的是在GPU端执行
{
    const float3 U = params.cam_u;
    const float3 V = params.cam_v;
    const float3 W = params.cam_w;
    const float2 d = 2.0f * make_float2(
            static_cast<float>( idx.x ) / static_cast<float>( dim.x ),
            static_cast<float>( idx.y ) / static_cast<float>( dim.y )
            ) - 1.0f; //2*x-1即将[0,1]区间的比值变为[-1,1]的位置，即该像素在整个窗口的位置

    origin    = params.cam_eye;
    direction = normalize( d.x * U + d.y * V + W ); //从(0,0,2)摄像机处发射向XOY平面
}

进行optixTrace后，在空间中会触发miss函数或closesthit函数。

如果没有与物体相交，触发miss函数，直接填充背景色。

extern "C" __global__ void __miss__ms()
{
    MissData* miss_data  = reinterpret_cast<MissData*>( optixGetSbtDataPointer() );
    setPayload(  miss_data->bg_color );
}

如果与物体相交，触发closesthit函数。

extern "C" __global__ void __closesthit__ch()
{
    // When built-in triangle intersection is used, a number of fundamental
    // attributes are provided by the OptiX API, indlucing barycentric coordinates.
    const float2 barycentrics = optixGetTriangleBarycentrics();

    setPayload( make_float3( barycentrics, 1.0f ) );
}

这里的optixGetTriangleBarycentrics是求该位置在当前三角形的重心坐标，是一个float2类型，说明是u*AB+v*AC的向量类型，setpayload函数就是写一个信息以便后续调用。

static __forceinline__ __device__ void setPayload( float3 p )
{
    optixSetPayload_0( __float_as_uint( p.x ) );
    optixSetPayload_1( __float_as_uint( p.y ) );
    optixSetPayload_2( __float_as_uint( p.z ) );
}

那么在raygen中获取到payload便可保存结果。

float3 result;
result.x = __uint_as_float( p0 );
result.y = __uint_as_float( p1 );
result.z = __uint_as_float( p2 );

// Record results in our output raster
params.image[idx.y * params.image_width + idx.x] = make_color( result );

码字不易，点个赞吧

总结

生成三角形后，光线往每一个pixel中发出射线，如果打到三角形上，则计算重心坐标生成颜色，否则填充背景色。

参考资料

NVIDIA OptiX 8.0 - Programming Guide

图形学