引言

上一节我们搭建了optix基础架构,但是使用的是一个无模型、无射线,更无光照着色的场景。本节我们来逐步引入模型和射线求交的内容。但是注意:本节实现的光追只弹射一次射线,将该射线第一次求交得到的颜色直接作为像素值,并非路径追踪的算法。

Part1. 使用GLFW窗口

本节我们来复现一下example03。

之前我们渲染出来图片以后,使用stb_image_write将图片直接输出到本地。为了更方便地预览渲染结果,甚至日后实现编辑场景功能,我们必须将渲染结果搬到窗口上显示。

我们需要在main.cpp中引用glfWindow/GLFWindow.h 和GL/gl.h这两个头文件,首先将glfwindow文件夹复制到工程,接下来我们右键项目-属性-VC++目录-包含目录,添加一条xxx\3rdParty\glfw\include;在库目录添加glfwindow.lib和glfw3.lib的所在路径;然后在链接-输入里添加glfwindow.lib和glfw3.lib。(如果找不到glfwindow.lib和glfw3.lib,就去编译一下官方的源码或者在网上下一个。这部分就是opengl配置环境那一套了不想多言,出错就百度吧)

接下来我们来定义一下glfw窗口:

struct SampleWindow : public GLFWindow
    {
        SampleWindow(const std::string& title)
            : GLFWindow(title)
        {}

        virtual void render() override
        {
            sample.render();
        }

        virtual void draw() override
        {
            sample.downloadPixels(pixels.data());
            if (fbTexture == 0)
                glGenTextures(1, &fbTexture);

            glBindTexture(GL_TEXTURE_2D, fbTexture);
            GLenum texFormat = GL_RGBA;
            GLenum texelType = GL_UNSIGNED_BYTE;
            glTexImage2D(GL_TEXTURE_2D, 0, texFormat, fbSize.x, fbSize.y, 0, GL_RGBA,
                texelType, pixels.data());

            glDisable(GL_LIGHTING);
            glColor3f(1, 1, 1);

            glMatrixMode(GL_MODELVIEW);
            glLoadIdentity();

            glEnable(GL_TEXTURE_2D);
            glBindTexture(GL_TEXTURE_2D, fbTexture);
            glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
            glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);

            glDisable(GL_DEPTH_TEST);

            glViewport(0, 0, fbSize.x, fbSize.y);

            glMatrixMode(GL_PROJECTION);
            glLoadIdentity();
            glOrtho(0.f, (float)fbSize.x, 0.f, (float)fbSize.y, -1.f, 1.f);

            glBegin(GL_QUADS);
            {
                glTexCoord2f(0.f, 0.f);
                glVertex3f(0.f, 0.f, 0.f);

                glTexCoord2f(0.f, 1.f);
                glVertex3f(0.f, (float)fbSize.y, 0.f);

                glTexCoord2f(1.f, 1.f);
                glVertex3f((float)fbSize.x, (float)fbSize.y, 0.f);

                glTexCoord2f(1.f, 0.f);
                glVertex3f((float)fbSize.x, 0.f, 0.f);
            }
            glEnd();
        }

        virtual void resize(const vec2i& newSize)
        {
            fbSize = newSize;
            sample.resize(newSize);
            pixels.resize(newSize.x * newSize.y);
        }

        vec2i                 fbSize;
        GLuint                fbTexture{ 0 };
        SampleRenderer        sample;
        std::vector<uint32_t> pixels;
    };

这一部分就是opengl的基础了,即使你对它的语法和api不熟,光看函数名和变量名也能猜出来这是干嘛的。首先窗体是一个继承了GlfWindow的类,他的成员包括我们的渲染器SampleRenderer、从设备下载下来的像素pixels、尺寸fbSize以及要绑定在窗口上的2D纹理fbTexture。

一般我们使用一个标题字符串来构造窗口,然后我们可以分别override它的render()、draw()、resize()方法。

glfw的render()方法就是一个渲染回调,因为OpenGL窗口定期要渲染一次画面,本质就是定期调用一次render()方法。在override时,我们就让该窗口定期调用一次SampleRenderer的渲染函数。

glfw的draw()方法就是一个绘制回调,在窗口上绘制像素。这里区分一下渲染和绘制:渲染是后台进行的,绘制是在窗口前台进行的。render未完成的时候也可以中途把render的当前结果拿来draw,这样就能看到绘制过程了。在override时,我们首先通过download函数从管线拿到目前渲染的结果,然后在glfw窗口中定义视口、绑定2D纹理、定位2D纹理,这样就能在窗口中观看到2D纹理上显示的场景渲染结果了。

glfw的resize()顾名思义,当窗口改变大小时,我们要让渲染器的缓冲区尺寸也跟着改变大小。

extern "C" int main(int ac, char** av)
    {
        try {
            SampleWindow* window = new SampleWindow("Optix 7 Course Example");
            window->run();

        }
        catch (std::runtime_error& e) {
            std::cout << GDT_TERMINAL_RED << "FATAL ERROR: " << e.what()
                << GDT_TERMINAL_DEFAULT << std::endl;
            exit(1);
        }
        return 0;
    }

稍微改动一下main函数,这回我们直接实例化一个SampleWindow,让他运行即可,效果如下:

Part.2 三维空间三角形渲染

截至目前,渲染都是2D的。现在我们希望复现example04,将渲染升至三维空间,首先就要考虑要引入哪些新元素:①相机;②三角形;③方体(毕竟单个三角形依然是平面的,多个三角形组成方体就显得立体了);④场景结构(将场景中所有三角面打包构成可加速的求交结构);⑤raygen、miss、hit shader的真正定义(之前几乎全是空的)。

1、缓冲区内容补充

我们定义的LaunchParams缓冲区本质是为gpu端的shader代码服务的,根据前面各种shader的定义可以知道,Raygen shader需要拿到相机位置才能计算第一根发射的光线;hit shader需要拿到场景中的几何结构才能完成求交。因此,我们需要向LaunchParams中补充相机和场景几何结构:

namespace osc {
  using namespace gdt;
  
  struct LaunchParams
  {
    struct {
      uint32_t *colorBuffer;
      vec2i     size;
    } frame;
    
    struct {
      vec3f position;
      vec3f direction;
      vec3f horizontal;
      vec3f vertical;
    } camera;

    OptixTraversableHandle traversable;
  };

}

camera中定义了position(相机位置)、direction(相机面向的方向)、horizontal(描述相机视口水平的方向和长度)、vertical(描述相机视口竖直的方向和长度)。之所以需要horizontal和vertical这两个参量是因为它们能告诉Raygen shader 视口一共要发射多少像素、从哪里发射。

而traversable就是我们的场景几何结构了,类型是optix定义的OptixTraversableHandle类型,这是一个经过特殊处理(为了求交加速)的包含了场景所有三角形或其他形状的数据结构。后面会讲这个东西怎么用。

2、相机

在窗口中,我们鼠标操控的是glfwindows的内置相机,这个内置相机是包括position、direction、up三个方位向量的经典相机,并且还有一些其他的冗余信息。如果我们直接把glfwindows的相机传送给LaunchParams,会显得非常臃肿。所以我们再在SampleRenderer.h定义一个额外的相机Camera,参量也是position、direction、up,让他去把glfwindows内置相机的参数下载下来,再调用咱们自写的setCamera()函数,把参数换算给LaunchParams的camera。

SampleRenderer.h:

struct Camera {
    vec3f from;
    vec3f at;
    vec3f up;
  };

class SampleRenderer{
...
void setCamera(const Camera& camera);
...
Camera lastSetCamera;
}

setCamera就是根据当前的glfwindows内置相机,或我们新定义的相机,计算出position、direction、horizontal、vertical的值并更新到LaunchParams的camera中。这里额外定义了一个lastSetCamera变量,这个是当窗口触发resize的时候,需要临时更新LaunchParams的camera,这时省的再去glfw里拿新相机了,直接使用上一次setCamera时存的值即可。

SampleRenderer.cpp:

  void SampleRenderer::render()
  {
      // sanity check: make sure we launch only after first resize is
      // already done:
      if (launchParams.frame.size.x == 0) return;

      launchParamsBuffer.upload(&launchParams, 1);

      OPTIX_CHECK(optixLaunch(/*! pipeline we're launching launch: */
          pipeline, stream,
          /*! parameters and SBT */
          launchParamsBuffer.d_pointer(),
          launchParamsBuffer.sizeInBytes,
          &sbt,
          /*! dimensions of the launch: */
          launchParams.frame.size.x,
          launchParams.frame.size.y,
          1
      ));
      // sync - make sure the frame is rendered before we download and
      // display (obviously, for a high-performance application you
      // want to use streams and double-buffering, but for this simple
      // example, this will have to do)
      CUDA_SYNC_CHECK();
  }
void SampleRenderer::setCamera(const Camera &camera)
  {
    lastSetCamera = camera;
    launchParams.camera.position  = camera.from;
    launchParams.camera.direction = normalize(camera.at-camera.from);
    const float cosFovy = 0.66f;
    const float aspect = launchParams.frame.size.x / float(launchParams.frame.size.y);
    launchParams.camera.horizontal
      = cosFovy * aspect * normalize(cross(launchParams.camera.direction,
                                           camera.up));
    launchParams.camera.vertical
      = cosFovy * normalize(cross(launchParams.camera.horizontal,
                                  launchParams.camera.direction));
  }
  void SampleRenderer::resize(const vec2i &newSize)
  {
    // if window minimized
    if (newSize.x == 0 | newSize.y == 0) return;
    
    // resize our cuda frame buffer
    colorBuffer.resize(newSize.x*newSize.y*sizeof(uint32_t));

    // update the launch parameters that we'll pass to the optix
    // launch:
    launchParams.frame.size  = newSize;
    launchParams.frame.colorBuffer = (uint32_t*)colorBuffer.d_pointer();

    // and re-set the camera, since aspect may have changed
    setCamera(lastSetCamera);
  }

render()函数里主要纠正一些变量名;新增的setCamera()代码还是比较通俗的。cosFovy是视口的纵向张角的cosine值,cosFovy越大-视野越广。注意每次调用setCamera后预存一下lastSetCamera,这样当窗口发生resize的时候,resize函数直接用lastSetCamera来调用setCamera就很方便了。

main.cpp:

struct SampleWindow : public GLFWindow
    {
        SampleWindow(const std::string& title)
            : GLFWindow(title)
        {}

        virtual void render() override
        {
            if (cameraFrame.modified) {
                sample.setCamera(Camera{ cameraFrame.get_from(),
                                         cameraFrame.get_at(),
                                         cameraFrame.get_up() });
                cameraFrame.modified = false;
            }
            sample.render();
        }

3、三角模型

无论是C++光栅化渲染器还是C++光线追踪渲染器,我都已经定义过多次三角模型了,所以这里我不想太多费口舌了。。。

首先我们在SampleRenderer.h中定义一个TriangleMesh类,用来保存一个或多个模型的所有顶点、三角面(后续还会有法线)信息。

SampleRenderer.h:

#include "gdt/math/AffineSpace.h"
struct TriangleMesh {
    /*! add a unit cube (subject to given xfm matrix) to the current
        triangleMesh */
    void addUnitCube(const affine3f &xfm);
    
    //! add aligned cube aith front-lower-left corner and size
    void addCube(const vec3f &center, const vec3f &size);
    
    std::vector<vec3f> vertex;
    std::vector<vec3i> index;
  };

如上代码所示,vertex就保存模型的所有顶点,index就保存模型三角面的索引(比如0,1,2,1,2,3的索引代表012构成第一个三角形,123构成第二个三角形)。为了方便构造立方体实例,定义了个addCube和addUnitCube函数。addCube就是给定立方体的位置和尺寸,得到立方体的三个轴的方向和长度,然后调用addUnitCube;addUnitCube就是根据轴来计算每个顶点的位置,然后构成索引面,将顶点和面放到该TriangleMesh的vertex和index组里。

SampleRenderer.cpp:

 void TriangleMesh::addCube(const vec3f& center, const vec3f& size)
  {
      affine3f xfm;
      xfm.p = center - 0.5f * size;
      xfm.l.vx = vec3f(size.x, 0.f, 0.f);
      xfm.l.vy = vec3f(0.f, size.y, 0.f);
      xfm.l.vz = vec3f(0.f, 0.f, size.z);
      addUnitCube(xfm);
  }

  /*! add a unit cube (subject to given xfm matrix) to the current
      triangleMesh */
  void TriangleMesh::addUnitCube(const affine3f& xfm)
  {
      int firstVertexID = (int)vertex.size();
      vertex.push_back(xfmPoint(xfm, vec3f(0.f, 0.f, 0.f)));
      vertex.push_back(xfmPoint(xfm, vec3f(1.f, 0.f, 0.f)));
      vertex.push_back(xfmPoint(xfm, vec3f(0.f, 1.f, 0.f)));
      vertex.push_back(xfmPoint(xfm, vec3f(1.f, 1.f, 0.f)));
      vertex.push_back(xfmPoint(xfm, vec3f(0.f, 0.f, 1.f)));
      vertex.push_back(xfmPoint(xfm, vec3f(1.f, 0.f, 1.f)));
      vertex.push_back(xfmPoint(xfm, vec3f(0.f, 1.f, 1.f)));
      vertex.push_back(xfmPoint(xfm, vec3f(1.f, 1.f, 1.f)));


      int indices[] = { 0,1,3, 2,3,0,
                       5,7,6, 5,6,4,
                       0,4,5, 0,5,1,
                       2,3,7, 2,7,6,
                       1,5,7, 1,7,3,
                       4,0,2, 4,2,6
      };
      for (int i = 0; i < 12; i++)
          index.push_back(firstVertexID + vec3i(indices[3 * i + 0],
              indices[3 * i + 1],
              indices[3 * i + 2]));
  }

xfm是用来定义立方体的轴心和三维轴的。得到轴以后,调用的addUnitCube()就可以根据轴心和三维轴算出立方体每个顶点的位置(这种设计模式有点晦涩,大家可以根据自己喜好重写一下这些逻辑)。

要注意的是,一个TriangleMesh未必只有一个立方体,因此在添加新立方体的时候,先要获取前面已经有多少vertex了,此时你新立方体真正的顶点索引=前面立方体vertex的总数+当前立方体内该顶点的索引。

4、场景加速结构构造

首先我们在SampleRenderer.h中继续补充一个函数和几个参量:

SampleRenderer.h:

  class SampleRenderer
  {
    ...
    OptixTraversableHandle buildAccel(const TriangleMesh &model);
    ...
    const TriangleMesh model;
    CUDABuffer vertexBuffer;
    CUDABuffer indexBuffer;
    //! buffer that keeps the (final, compacted) accel structure
    CUDABuffer asBuffer;
  }

首先实例化一个模型实例model,然后申请三个gpu端的buffer,用来分别临时存储顶点数据、索引数据和加速结构数据。后面解释这几个buffer怎么用。

SampleRenderer.cpp:

OptixTraversableHandle SampleRenderer::buildAccel(const TriangleMesh &model)
  {
    // upload the model to the device: the builder
    vertexBuffer.alloc_and_upload(model.vertex);
    indexBuffer.alloc_and_upload(model.index);
    
    OptixTraversableHandle asHandle { 0 };
    
    // ==================================================================
    // triangle inputs
    // ==================================================================
    OptixBuildInput triangleInput = {};
    triangleInput.type
      = OPTIX_BUILD_INPUT_TYPE_TRIANGLES;

    // create local variables, because we need a *pointer* to the
    // device pointers
    CUdeviceptr d_vertices = vertexBuffer.d_pointer();
    CUdeviceptr d_indices  = indexBuffer.d_pointer();
      
    triangleInput.triangleArray.vertexFormat        = OPTIX_VERTEX_FORMAT_FLOAT3;
    triangleInput.triangleArray.vertexStrideInBytes = sizeof(vec3f);
    triangleInput.triangleArray.numVertices         = (int)model.vertex.size();
    triangleInput.triangleArray.vertexBuffers       = &d_vertices;
    
    triangleInput.triangleArray.indexFormat         = OPTIX_INDICES_FORMAT_UNSIGNED_INT3;
    triangleInput.triangleArray.indexStrideInBytes  = sizeof(vec3i);
    triangleInput.triangleArray.numIndexTriplets    = (int)model.index.size();
    triangleInput.triangleArray.indexBuffer         = d_indices;
    
    uint32_t triangleInputFlags[1] = { 0 };
    
    // in this example we have one SBT entry, and no per-primitive
    // materials:
    triangleInput.triangleArray.flags               = triangleInputFlags;
    triangleInput.triangleArray.numSbtRecords               = 1;
    triangleInput.triangleArray.sbtIndexOffsetBuffer        = 0; 
    triangleInput.triangleArray.sbtIndexOffsetSizeInBytes   = 0; 
    triangleInput.triangleArray.sbtIndexOffsetStrideInBytes = 0; 
      
    // ==================================================================
    // BLAS setup
    // ==================================================================
    
    OptixAccelBuildOptions accelOptions = {};
    accelOptions.buildFlags             = OPTIX_BUILD_FLAG_NONE
      | OPTIX_BUILD_FLAG_ALLOW_COMPACTION
      ;
    accelOptions.motionOptions.numKeys  = 1;
    accelOptions.operation              = OPTIX_BUILD_OPERATION_BUILD;
    
    OptixAccelBufferSizes blasBufferSizes;
    OPTIX_CHECK(optixAccelComputeMemoryUsage
                (optixContext,
                 &accelOptions,
                 &triangleInput,
                 1,  // num_build_inputs
                 &blasBufferSizes
                 ));
    
    // ==================================================================
    // prepare compaction
    // ==================================================================
    
    CUDABuffer compactedSizeBuffer;
    compactedSizeBuffer.alloc(sizeof(uint64_t));
    
    OptixAccelEmitDesc emitDesc;
    emitDesc.type   = OPTIX_PROPERTY_TYPE_COMPACTED_SIZE;
    emitDesc.result = compactedSizeBuffer.d_pointer();
    
    // ==================================================================
    // execute build (main stage)
    // ==================================================================
    
    CUDABuffer tempBuffer;
    tempBuffer.alloc(blasBufferSizes.tempSizeInBytes);
    
    CUDABuffer outputBuffer;
    outputBuffer.alloc(blasBufferSizes.outputSizeInBytes);
      
    OPTIX_CHECK(optixAccelBuild(optixContext,
                                /* stream */0,
                                &accelOptions,
                                &triangleInput,
                                1,  
                                tempBuffer.d_pointer(),
                                tempBuffer.sizeInBytes,
                                
                                outputBuffer.d_pointer(),
                                outputBuffer.sizeInBytes,
                                
                                &asHandle,
                                
                                &emitDesc,1
                                ));
    CUDA_SYNC_CHECK();
    
    // ==================================================================
    // perform compaction
    // ==================================================================
    uint64_t compactedSize;
    compactedSizeBuffer.download(&compactedSize,1);
    
    asBuffer.alloc(compactedSize);
    OPTIX_CHECK(optixAccelCompact(optixContext,
                                  /*stream:*/0,
                                  asHandle,
                                  asBuffer.d_pointer(),
                                  asBuffer.sizeInBytes,
                                  &asHandle));
    CUDA_SYNC_CHECK();
    
    // ==================================================================
    // aaaaaand .... clean up
    // ==================================================================
    outputBuffer.free(); // << the UNcompacted, temporary output buffer
    tempBuffer.free();
    compactedSizeBuffer.free();

    return asHandle;
  }

上面这段代码就是本文的精华部分了。但说句实话,基本又全是套路式地调用api...所以我也是一知半解。尝试解读一下这段代码,首先我们要将模型的顶点和索引上传到vertex和index两个buffer中,然后定义一个triangleInput,将这两大缓冲区送入注册绑定;同时每个物体要通过sbt对shader进行绑定,不过这个example没有特殊材质shader,所有的三角面最后都绑定一个相同的hit group shader。

上述信息都绑定好后,便开始配置加速结构、创建加速结构、申请加速结构的缓冲空间,然后执行加速结构(optixAccelBuild),再进行一步压缩,最终得到了OptixTraversableHandle类型的完整加速结构,存在asHandle中。

在渲染器初始化的时候,我们需要在着色器实例已完成、管线搭建之前构造好这个加速结构,并将它配置进LaunchParams中。

SampleRenderer.cpp:

SampleRenderer::SampleRenderer()
  {
    ...
    createHitgroupPrograms();

    launchParams.traversable = buildAccel(model);

    std::cout << "#osc: setting up optix pipeline ..." << std::endl;
    createPipeline();
    ...
  }

5、着色器

虽然说本节的精华部分在加速结构的构造上,但实际上最有含金量的内容仍然是我们的shader部分。在前面几个example中shader基本都是空白,没有任何光追相关的行为,这一节我们就要写入正儿八经的光追shader。

首先我们补充几个基础函数:

devicePrograms.cu:

enum { SURFACE_RAY_TYPE=0, RAY_TYPE_COUNT };
  
  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;
  }

  template<typename T>
  static __forceinline__ __device__ T *getPRD()
  { 
    const uint32_t u0 = optixGetPayload_0();
    const uint32_t u1 = optixGetPayload_1();
    return reinterpret_cast<T*>( unpackPointer( u0, u1 ) );
  }

SURFACE_RAY_TYPE就是选择射线类型,射线类型不仅包括传统的计算光色的射线,还有计算阴影的射线。这里我们只使用传统默认的光色射线。

unpackPainter和packPainter就是对某个指针中保存的地址进行高32位和低32位的拆分和合并。之所以需要拆分,是因为我们的计算机是64位的所以指针也是64位的,然而gpu的寄存器是32位的,因此只能将指针拆分成两部分存进gpu的寄存器。

这里简单解释下payload,payload就类似一个负载寄存器,负责在不同shader之间传递信息。后面我们会在Raygen Shader中申请一个颜色指针来存储最终的光追颜色,当开始tracing后,会将这个颜色指针拆分(pack)写入0号和1号寄存器,当Hit Shader和Miss Shader想往里面写东西时,就可以通过上面代码中的getPRD函数获得0号和1号寄存器中的值,将其unpack便得到了那个颜色指针,然后就可以往这个颜色指针里写内容了。

接下来我们分别编写一下Hit、Miss、RayGen shader。首先是Hit Shader:

extern "C" __global__ void __closesthit__radiance()
  {
    const int   primID = optixGetPrimitiveIndex();
    vec3f &prd = *(vec3f*)getPRD<vec3f>();
    prd = gdt::randomColor(primID);
  }
  
  extern "C" __global__ void __anyhit__radiance()
  { /*! for this simple example, this will remain empty */ }

因为场景没有半透明的物体,因此AnyHit暂时不需要写东西。ClosetHit必然要写了(因为它相当于BSDF函数),这一节我们不需要任何复杂光照模型,直接根据三角面索引给一个随机颜色就好了:首先通过optixGetPrimitiveIndex获得当前击中的三角面索引并以此生成一个随机颜色,然后获得把颜色放进0号、1号寄存器,这样就完成了着色过程。

接下来是Miss Shader:

extern "C" __global__ void __miss__radiance()
  {
    vec3f &prd = *(vec3f*)getPRD<vec3f>();
    // set to constant white as background color
    prd = vec3f(1.f);
  }

这个更简单,直接返回纯色(1,1,1)。

最后是最复杂的Raygen Shader:

extern "C" __global__ void __raygen__renderFrame()
  {
    // compute a test pattern based on pixel ID
    const int ix = optixGetLaunchIndex().x;
    const int iy = optixGetLaunchIndex().y;

    const auto &camera = optixLaunchParams.camera;

    // our per-ray data for this example. what we initialize it to
    // won't matter, since this value will be overwritten by either
    // the miss or hit program, anyway
    vec3f pixelColorPRD = vec3f(0.f);

    // the values we store the PRD pointer in:
    uint32_t u0, u1;
    packPointer( &pixelColorPRD, u0, u1 );

    // normalized screen plane position, in [0,1]^2
    const vec2f screen(vec2f(ix+.5f,iy+.5f)
                       / vec2f(optixLaunchParams.frame.size));
    
    // generate ray direction
    vec3f rayDir = normalize(camera.direction
                             + (screen.x - 0.5f) * camera.horizontal
                             + (screen.y - 0.5f) * camera.vertical);

    optixTrace(optixLaunchParams.traversable,
               camera.position,
               rayDir,
               0.f,    // tmin
               1e20f,  // tmax
               0.0f,   // rayTime
               OptixVisibilityMask( 255 ),
               OPTIX_RAY_FLAG_DISABLE_ANYHIT,//OPTIX_RAY_FLAG_NONE,
               SURFACE_RAY_TYPE,             // SBT offset
               RAY_TYPE_COUNT,               // SBT stride
               SURFACE_RAY_TYPE,             // missSBTIndex 
               u0, u1 );

    const int r = int(255.99f*pixelColorPRD.x);
    const int g = int(255.99f*pixelColorPRD.y);
    const int b = int(255.99f*pixelColorPRD.z);

    // convert to 32-bit rgba value (we explicitly set alpha to 0xff
    // to make stb_image_write happy ...
    const uint32_t rgba = 0xff000000
      | (r<<0) | (g<<8) | (b<<16);

    // and write to frame buffer ...
    const uint32_t fbIndex = ix+iy*optixLaunchParams.frame.size.x;
    optixLaunchParams.frame.colorBuffer[fbIndex] = rgba;
  }

因为我们现在已经有了相机、有了视口,所以当我们枚举每一个像素的时候,我们可以算出这个像素从哪里、向什么方向发射第一根射线。首先通过vec2f(ix+.5f,iy+.5f) / vec2f(optixLaunchParams.frame.size)算出当前像素的x、y坐标位于视口的多少百分比的位置,然后根据这个百分比算出这个像素发射的光线向什么方向偏。

这里定义了一个pixelColorPRD来存储该像素的最终颜色,通过packPointer将其分为u0、u1两部分,在调用optixTrace这个开始渲染的终极函数时,u0和u1就会自动被写入gpu的0号和1号寄存器中,方便hit和miss shader往里面写入颜色结果。

关于optixTrace函数其实没啥想说的,optix最精华的封装就在这里了,把相机、场景、射线属性、寄存信息(u0和u1)传值进去就ok了,天花乱坠的求交便在这里发生。等光追结束了,把我们存的颜色指针pixelColorPRD中的rgb值算出来,写入LaunchParams的frame中即可。

当然还是老样子,写完.cu立刻编译cuda,编译完得到ptx机器码后用python将其parse成MyShader.c文件放到目录下,这样SampleRenderer.cpp中的embedded_ptx_code[]就能获取到机器码。(当然因为我们main.cpp的代码还没改,所以此时编译以后出现runtime error是正常的,不要紧张。当务之急就是拿到ptx机器码)

6、渲染

回到main.cpp函数,定义一下我们的场景。

main.cpp:

struct SampleWindow : public GLFCameraWindow
    {
        SampleWindow(const std::string& title,
            const TriangleMesh& model,
            const Camera& camera,
            const float worldScale)
            : GLFCameraWindow(title, camera.from, camera.at, camera.up, worldScale),
            sample(model)
        {
        }

        virtual void render() override
        {
            if (cameraFrame.modified) {
                sample.setCamera(Camera{ cameraFrame.get_from(),
                                         cameraFrame.get_at(),
                                         cameraFrame.get_up() });
                cameraFrame.modified = false;
            }
            sample.render();
        }

首先glfw继承的窗口要改成GLFCameraWindow,因为这个类带有相机,并且可以直接用鼠标操纵相机,非常方便。构造函数需要稍作修改,需要提前把场景中的模型、相机、世界尺度这些信息传入,给GLFwindow自带的camera传入一个初始值。

接下来改写一下render函数,当鼠标交互场景导致GLFwindow的相机发生变化时,就会触发SampleRenderer的setCamera函数,同步更新到LaunchParams的camera中。

extern "C" int main(int ac, char** av)
    {
        try {
            TriangleMesh model;
            // 100x100 thin ground plane
            model.addCube(vec3f(0.f, -1.5f, 0.f), vec3f(10.f, .1f, 10.f));
            // a unit cube centered on top of that
            model.addCube(vec3f(0.f, 0.f, 0.f), vec3f(2.f, 2.f, 2.f));

            Camera camera = { /*from*/vec3f(-10.f,2.f,-12.f),
                /* at */vec3f(0.f,0.f,0.f),
                /* up */vec3f(0.f,1.f,0.f) };

            // something approximating the scale of the world, so the
            // camera knows how much to move for any given user interaction:
            const float worldScale = 10.f;

            SampleWindow* window = new SampleWindow("Optix 7 Course Example",
                model, camera, worldScale);
            window->run();

        }
        catch (std::runtime_error& e) {
            std::cout << GDT_TERMINAL_RED << "FATAL ERROR: " << e.what()
                << GDT_TERMINAL_DEFAULT << std::endl;
            exit(1);
        }
        return 0;
    }

在main.cpp中,我们来定义一下场景的最初相机和场景模型。自定义的部分相当简单我就不多赘述了,大家可以根据喜好搭搭积木。