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

这次我们看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 handlingOptixTraversableHandle 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 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( &d_vertices ), vertices_size ) );    CUDA_CHECK( cudaMemcpy(                reinterpret_cast( 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( 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( &d_temp_buffer_gas ),                gas_buffer_sizes.tempSizeInBytes                ) );    CUDA_CHECK( cudaMalloc(                reinterpret_cast( &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( d_temp_buffer_gas ) ) );    CUDA_CHECK( cudaFree( reinterpret_cast( 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 1688 42 0 1024 768 1692 32 0 1024 768 1693 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( idx.x ) / static_cast( dim.x ),            static_cast( idx.y ) / static_cast( 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( 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 rasterparams.image[idx.y * params.image_width + idx.x] = make_color( result );

执行结果：

码字不易，点个赞吧

总结

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

参考资料

NVIDIA OptiX 8.0 – Programming Guide