Optix (01)

Posted on Edited on In Views:
Optix7 Course

Optix7 Course

Main Talk

  • cuda runtime API vs driver API
Runtime API Driver API
cuda* prefix (i.e., cudaStreamCreate) cu* prefi (i.e., cuStreamCreate)
Launch with <<< >>>> Launch kernels with cuLaunch*
CUDA kernels embedded in .o Kernels loaded or compiled at runtime
  • setup
1
2
// initialize cuda
// initialize optix
  • compilation
    • OptiX is built on separate compilation
  • shader
1
2
// C++ shaders, compiled with NVCC
extern "C" __global__ void __raygen__pathtrace() { /**/ }
  • pipeline creation
    • cuda shader -> PTX -> Module(s) -> ProgramGroup(s) -> Pipeline(s)
  • traversable
    • General
      • Geometry Acceleration structure (GAS)
        • Triangles or custom primitives
      • Instance Acceleration structure (IAS)
    • Motion blur
      • Matrix motion transform
        • world coordinates
      • SRT motion transform
        • Scale Rotation Translation
      • Static Transform (rarely used)
        • object-to-world transform

  • Shader Binding Table
    • SBT is just GPU memory
    • Keep them small
  • denoiser
1
2
3
4
5
6
7
8
9
10
11
12
// 1. Create denoiser object
optixDenoiserCreate();
// 2. Select neural network (model)
optixDenoiserSetModel();
// 3. Compute memory resource sizes
optixDenoiserComputeMemoryResources();
// 4. Allocate state and scratch resources
cuMemAlloc();
// 5. Setup denoiser (once)
optixDenoiserSetup();
// 6. Invoke (per image)
optixDenoiserInvoke();

Tutorial

E01-helloOptix

  • setup
1
2
3
4
// init cuda
cudaFree(0);
// init optix
OPTIX_CHECK( optixInit() );
  • 预先将 cuda 代码编译成字符串
    • cmake
      • cuda_compile_ptx 编译成 ptx 代码
      • bin2c 转化为字符串
      • 作为 target 的一部分添加
1
2
// optix7course\example02_pipelineAndRayGen\SampleRenderer.cpp
extern "C" char embedded_ptx_code[];
1
2
3
4
add_executable(ex02_pipelineAndRayGen
 ${embedded_ptx_code}
 # ...
);

E02-pipelineAndRayGen

  • SBT
  • 启动
1
2
3
4
5
6
7
8
inline OptixResult optixLaunch(OptixPipeline                  pipeline,
                               CUstream                       stream,
                               CUdeviceptr                    pipelineParams,
                               size_t                         pipelineParamsSize,
                               const OptixShaderBindingTable* sbt,
                               unsigned int                   width,
                               unsigned int                   height,
                               unsigned int                   depth )
  • 参数传递：pipelineParams
1
2
// cuda device 端
extern "C" __constant__ LaunchParams optixLaunchParams;
  • shader 之间的关系：绿色固定、灰色可编程

E03-inGLFWindow

E04-firstTriangleMesh

  • trace-ray 加速结构构建
    • optixAccelBuild()
    • optixAccelCompact()
    • 作为 launchParams 传递给 cuda
  • PRD (per ray data)
    • 通过 payload 传递，但是传递的数据必须是 32bit-int
      • 0-31：optixGetPayload_0()
    • 在 ray-gen 中获得数据地址，通过 payload 传递给 closest-hit，解包地址写入数据
      • optixTrace 的最后参数就是变长参数，即 payload 的 slot（传递数据 PRD）
1
2
3
4
5
6
7
8
9
10
11
12
// pack and unpack
static __forceinline__ __device__void *unpackPointer(uint32_t i0, uint32_t i1) {
    const uint64_t uptr = static_cast<uint64_t>(i0) << 32 | i1;
    void *ptr = reinterpret_cast<void *>(uptr);
    return ptr;
}

static __forceinline__ __device__ void packPointer(void *ptr, uint32_t &i0, uint32_t &i1) {
    const uint64_t uptr = reinterpret_cast<uint64_t>(ptr);
    i0 = uptr >> 32;
    i1 = uptr & 0x00000000ffffffff;
}

E05-firstSBTData

  • optixGetSbtDataPointer()
    • 从不同的 shader 中能够获取到不同的资源（SBT 生成的时候指定的）

E06-multipleObjects

  • optixAccelComputeMemoryUsage()、optixAccelBuild() 中的每一个 OptixBuildInput 结构体指向一个 object (Mesh)
  • optixGetSbtDataPointer()：不同的 OptixBuildInput 对应不同的返回值
  • optixGetPrimitiveIndex()：返回的是在每一个 OptixBuildInput 内部的索引

E07-firstRealModel

  • 使用 tinyobjloader 加载 obj 模型
  • tinyobjloader 最原始的加载
    • 加载结束包含多个 shape，每一个 shape 中的 mesh 结构包含所有的面片信息
      • 每一个 shape 内部，顶点、法线、纹理坐标、材质贴图索引都重新开始编号（从 0 开始）
    • mesh 包含每一个顶点的索引、每一个面的顶点数、每一个面的材质、（smoothing_group_ids、tags）
    • index 顶点索引包括位置、法向、纹理坐标
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
typedef struct {
    std::string name;
    mesh_t mesh;
    lines_t lines;
    points_t points;
} shape_t;

typedef struct {
    std::vector<index_t> indices;
    // The number of vertices per face. 3 = triangle, 4 = quad, ... Up to 255 vertices per face.
    std::vector<unsigned char> num_face_vertices;
    // per-face material ID
    std::vector<int> material_ids;  
    // per-face smoothing group ID(0 = off. positive value = group id)
    std::vector<unsigned int> smoothing_group_ids;  
    // SubD tag
    std::vector<tag_t> tags;
} mesh_t;

// Index struct to support different indices for vtx/normal/texcoord.
// -1 means not used.
typedef struct {
    int vertex_index;
    int normal_index;
    int texcoord_index;
} index_t;
  • 加载处理
    • 按照材质进行分类，每一个材质构建一个新的索引序列（int），即顶点、法线、纹理共用一个索引
      • 不同于原始的索引，原始的索引是（3 x int），顶点、法线、纹理分开存索引
      • 只有顶点、法线、纹理都相同的顶点（原始索引相同），索引才相同
      • other：也可以按照原始的方式，每一个顶点保存 3 个索引
    • 处理完之后，每一个材质内部的做因都重新编号（都是从 0 开始）
1
2
3
4
5
6
7
8
9
10
11
// 加载处理后的结构体
// vertex.size() = normal.size() = texcoord.size()
struct TriangleMesh {
    std::vector<vec3f> vertex;
    std::vector<vec3f> normal;
    std::vector<vec2f> texcoord;
    std::vector<vec3i> index;

    // material data:
    vec3f diffuse;
};
  • cpp 并行库：tbb
  • 之后和E6一样处理，每一个加载好的 mesh 作为一个 OptixBuildInput

E08-addingTextures

  • 使用 cuda 的 API 进行控制
  • 在模型加载的时候同时加载纹理，TriangleMesh 结构体中加上材质 ID
    • 加载的纹理挂载到 model 上（不是 Mesh 上）
      • 一个 model 上从 0 开始编码
      • 通过文件名进行区分
1
2
3
4
5
struct TriangleMesh {
    // ...
    int diffuseTextureID{-1};
    // ...
};
  • 图片竖直翻转
1
stbi_set_flip_vertically_on_load(true);
  • 一起上传到 device 端，将指针保存到 SBT 结构中
1
2
3
4
5
6
7
8
9
struct TriangleMeshSBTData {
    vec3f color;
    vec3f *vertex;
    vec3f *normal;
    vec2f *texcoord;
    vec3i *index;
    bool hasTexture;
    cudaTextureObject_t texture; // 这里挂载当前 Mesh 对应的纹理贴图
};

E09-shadowRays

  • 增加 shadow ray
  • 需要增加各种 shadow ray 的 program，在如下函数中添加
    • createMissPrograms()
    • createHitgroupPrograms()
  • sbt 也需要两个
  • shadow ray
    • 在 __closesthit__radiance() 里面发射 shadow ray
    • 在 miss shader 中设置 visible，其他的 shader 不需要
    • rayFlags 如下
1
2
3
4
5
6
// For shadow rays: skip any/closest hit shaders and terminate on
// first intersection with anything. The miss shader is used to
// mark if the light was visible.
rayFlags = OPTIX_RAY_FLAG_DISABLE_ANYHIT
    | OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT
    | OPTIX_RAY_FLAG_DISABLE_CLOSESTHIT;

  • 更优实现

    • 技术实现上，我们可以不使用 hit group shader 实现，只使用 miss group shader 就行

    • BLOG

    • 直接在原始的 miss group shader 上实现，通过设置 rayFlags 实现

E10-softShadows

  • 软阴影
  • 面光源
  • 随机数生成器：simple 24-bit linear congruence generator
  • 随机数也放在 PRD 中传递（每一根光线拥有自己的随机数种子）
  • 样例实现
    • spp：raygen 中在像素内发生 \(N\) 条光线
    • shadow ray：在光源上采样 \(M\) 个点，做可见性测试，结果求平均作为最终结果

E11-denoiseColorOnly

1
2
3
4
5
6
7
optixDenoiserCreate();
optixDenoiserComputeMemoryResources();
optixDenoiserSetup();

optixDenoiserInvoke();

optixDenoiserDestroy();

E12-denoiseSeparateChannels

  • 在调用 optixDenoiserInvoke() 的时候加入 OptixDenoiserGuideLayer()
    • albedo, normal
    • 具体数据在 __closesthit__radiance 中记录第一个点的 albedo、normal
  • 同时加上了 tonemapping