https://github.com/ingowald/optix7course마찬가지로 Example 01~03까지에서 다른점만 작성한다.
// main.cpp
#include "SampleRenderer.h"
// our helper library for window handling
#include "glfWindow/GLFWindow.h"
#include <GL/gl.h>
/*! \namespace osc - Optix Siggraph Course */
namespace osc {
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();
}
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;
};
/*! main entry point to this example - initially optix, print hello
world, then exit */
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;
}
} // ::oscExample03과 다르게 마우스를 통한 camera 이동까지 구현되었다.
따라서 SampleWindow의 상속은 GLFCameraWindow으로 바뀌었다.
virtual void render() override
{
if (cameraFrame.modified) {
sample.setCamera(Camera{ cameraFrame.get_from(),
cameraFrame.get_at(),
cameraFrame.get_up() });
cameraFrame.modified = false;
}
sample.render();
}cameraFrame은 GLFWindow.h에서 사용하는 camera 변환이다.
내부적으로 마우스 입력을 받아서 저장하는 것 같다.
이렇게 camera의 위치가 변함에 따라 optix의 ray tracing도 camera 위치가 바뀌어야 한다. 이를 위해 SampleRenderer에서도 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;
}TriangleMesh, Camera 클래스는 SampleRenderer.cpp에서 정의한다.
// SampleRenderer.h
/*! \namespace osc - Optix Siggraph Course */
namespace osc {
struct Camera {
/*! camera position - *from* where we are looking */
vec3f from;
/*! which point we are looking *at* */
vec3f at;
/*! general up-vector */
vec3f up;
};
/*! a simple indexed triangle mesh that our sample renderer will
render */
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 ¢er, const vec3f &size);
std::vector<vec3f> vertex;
std::vector<vec3i> index;
};
/*! a sample OptiX-7 renderer that demonstrates how to set up
context, module, programs, pipeline, SBT, etc, and perform a
valid launch that renders some pixel (using a simple test
pattern, in this case */
class SampleRenderer
{
// ------------------------------------------------------------------
// publicly accessible interface
// ------------------------------------------------------------------
public:
/*! constructor - performs all setup, including initializing
optix, creates module, pipeline, programs, SBT, etc. */
SampleRenderer(const TriangleMesh &model);
/*! render one frame */
void render();
/*! resize frame buffer to given resolution */
void resize(const vec2i &newSize);
/*! download the rendered color buffer */
void downloadPixels(uint32_t h_pixels[]);
/*! set camera to render with */
void setCamera(const Camera &camera);
protected:
// ------------------------------------------------------------------
// internal helper functions
// ------------------------------------------------------------------
/*! helper function that initializes optix and checks for errors */
void initOptix();
/*! creates and configures a optix device context (in this simple
example, only for the primary GPU device) */
void createContext();
/*! creates the module that contains all the programs we are going
to use. in this simple example, we use a single module from a
single .cu file, using a single embedded ptx string */
void createModule();
/*! does all setup for the raygen program(s) we are going to use */
void createRaygenPrograms();
/*! does all setup for the miss program(s) we are going to use */
void createMissPrograms();
/*! does all setup for the hitgroup program(s) we are going to use */
void createHitgroupPrograms();
/*! assembles the full pipeline of all programs */
void createPipeline();
/*! constructs the shader binding table */
void buildSBT();
/*! build an acceleration structure for the given triangle mesh */
OptixTraversableHandle buildAccel(const TriangleMesh &model);
protected:
/*! @{ CUDA device context and stream that optix pipeline will run
on, as well as device properties for this device */
CUcontext cudaContext;
CUstream stream;
cudaDeviceProp deviceProps;
/*! @} */
//! the optix context that our pipeline will run in.
OptixDeviceContext optixContext;
/*! @{ the pipeline we're building */
OptixPipeline pipeline;
OptixPipelineCompileOptions pipelineCompileOptions = {};
OptixPipelineLinkOptions pipelineLinkOptions = {};
/*! @} */
/*! @{ the module that contains out device programs */
OptixModule module;
OptixModuleCompileOptions moduleCompileOptions = {};
/* @} */
/*! vector of all our program(group)s, and the SBT built around
them */
std::vector<OptixProgramGroup> raygenPGs;
CUDABuffer raygenRecordsBuffer;
std::vector<OptixProgramGroup> missPGs;
CUDABuffer missRecordsBuffer;
std::vector<OptixProgramGroup> hitgroupPGs;
CUDABuffer hitgroupRecordsBuffer;
OptixShaderBindingTable sbt = {};
/*! @{ our launch parameters, on the host, and the buffer to store
them on the device */
LaunchParams launchParams;
CUDABuffer launchParamsBuffer;
/*! @} */
CUDABuffer colorBuffer;
/*! the camera we are to render with. */
Camera lastSetCamera;
/*! the model we are going to trace rays against */
const TriangleMesh model;
CUDABuffer vertexBuffer;
CUDABuffer indexBuffer;
//! buffer that keeps the (final, compacted) accel structure
CUDABuffer asBuffer;
};
} // ::osc// SampleRenderer.cpp
//! add aligned cube with front-lower-left corner and size
void TriangleMesh::addCube(const vec3f ¢er, 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]));
}정육면체를 그리기 위한 것이므로 3차원의 점 8개를 vertex 배열에 저장한다. 그리고 정육면체의 정점에 대한 indicies를 통해 정점을 정육면체 형태로 저장하고, index 배열에 저장한다.
// SampleRenderer.cpp
/*! constructor - performs all setup, including initializing
optix, creates module, pipeline, programs, SBT, etc. */
SampleRenderer::SampleRenderer(const TriangleMesh &model)
{
initOptix();
std::cout << "#osc: creating optix context ..." << std::endl;
createContext();
std::cout << "#osc: setting up module ..." << std::endl;
createModule();
std::cout << "#osc: creating raygen programs ..." << std::endl;
createRaygenPrograms();
std::cout << "#osc: creating miss programs ..." << std::endl;
createMissPrograms();
std::cout << "#osc: creating hitgroup programs ..." << std::endl;
createHitgroupPrograms();
launchParams.traversable = buildAccel(model);
std::cout << "#osc: setting up optix pipeline ..." << std::endl;
createPipeline();
std::cout << "#osc: building SBT ..." << std::endl;
buildSBT();
launchParamsBuffer.alloc(sizeof(launchParams));
std::cout << "#osc: context, module, pipeline, etc, all set up ..." << std::endl;
std::cout << GDT_TERMINAL_GREEN;
std::cout << "#osc: Optix 7 Sample fully set up" << std::endl;
std::cout << GDT_TERMINAL_DEFAULT;
}TriangleMesh 객체를 인자로 받아서 Acceleration Structure를 만들고, launchParams의 traversable에 저장한다.
// LaunchPrams.h
#pragma once
#include "gdt/math/vec.h"
#include "optix7.h"
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;
};
} // ::osc여기서 traversable은 rendering object를 BVH와 같은 구조로 만든 것을 얘기하고, acceleration structure은 traversable을 포함한 ray casting에 필요한 더 많은 정보를 포함한다.
// 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;
}Acceleration structure를 만들고, GPU에 있는 traversable에 대한 handler를 반환한다.
먼저 index와 vertex를 위한 device memory에 allocation 후 memcpy를 진행한다.
반환할 optixTraversableHandle을 초기화한다.
OptixBuildInput는 Acceleration Structure를 만들기 위해서 primitive 입력을 저장한다.
OptixBuildInput에서 사용할 수 있는 primitive는 curve, sphere, triangle, custom primitive와 instance화 된 acceleration structure인 instances로 설정할 수 있다.
- OPTIX_BUILD_INPUT_TYPE_INSTANCES
- OPTIX_BUILD_INPUT_TYPE_INSTANCES_POINTERS
- OPTIX_BUILD_INPUT_TYPE_TRIANGLES
- OPTIX_BUILD_INPUT_TYPE_CURVES
- OPTIX_BUILD_INPUT_TYPE_SPHERES
- OPTIX_BUILD_INPUT_TYPE_CUSTOM_PRIMITIVES
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;triangle에 대한 format, stride, amount, device memory address를 입력받는다. vertex와 indicies에 대해서 설정해준다.
triangleInput에 대한 flag를 설정할 수 있다.
설정할 수 있는 flag는 다음과 같이 존재한다
- OPTIX_GEOMETRY_FLAG_NONE
- OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL
- OPTIX_GEOMETRY_FLAG_DISABLE_ANYHIT
OptixBuildInput의 triangleArray에 대한 SBT에 대한 설정을 한다.
numSbtRecords는 ray castring에서 사용될 SBT record 수를 나타낸다.
sbtIndexOffsetBuffer는 SBT record의 offset 정보를 저장하는 버퍼를 가리키는 버퍼를 가리킨다.
sbtIndexOffsetSizeInBytes는 offset 정보가 차지하는 바이트 수를 나타낸다.
sbtIndexOffsetStrideInBytes는 sbt offset 정보간의 간격을 나타낸다.
여기서는 하나의 record를 사용한다. 후에 다른 예제에서 여러개의 record를 사용하는 예제를 보여준다.
OptixAccelBuildOptions는 AS를 build하기 위한 옵션을 설정한다.
OPTIX_BUILD_FLAG_ALLOW_COMPACTION를 통해서 compaction을 허용한다.
build flags는 다음과 같이 존재한다.
- OPTIX_BUILD_FLAG_NONE = 0 ,
- OPTIX_BUILD_FLAG_ALLOW_UPDATE = 1u << 0 ,
- OPTIX_BUILD_FLAG_ALLOW_COMPACTION = 1u << 1 ,
- OPTIX_BUILD_FLAG_PREFER_FAST_TRACE = 1u << 2 ,
- OPTIX_BUILD_FLAG_PREFER_FAST_BUILD = 1u << 3 ,
- OPTIX_BUILD_FLAG_ALLOW_RANDOM_VERTEX_ACCESS = 1u << 4 ,
- OPTIX_BUILD_FLAG_ALLOW_RANDOM_INSTANCE_ACCESS = 1u << 5 ,
- OPTIX_BUILD_FLAG_ALLOW_OPACITY_MICROMAP_UPDATE = 1u << 6 ,
- OPTIX_BUILD_FLAG_ALLOW_DISABLE_OPACITY_MICROMAPS = 1u << 7
motionOptions는 motion blur와 관련된 설정이다. numKeys > 1일 경우 motion은 활성화된다. 만약 비활성화될 경우, timeBegin, timeEnd flag가 무시된다.
operation은 AS build 작업을 지정한다.
업데이트는 같은 AS를 사용하므로 AS build보다 빠르지만, bound를 수정하므로 AS의 품질이 떨어진다.
- OPTIX_BUILD_OPERATION_BUILD
- OPTIX_BUILD_OPERATION_UPDATE
Optix의 AS를 build하기 위해서는 임시 메모리와, AS를 위한 메모리 공간이 필요로 한다. 이를 위해서는 메모리 공간의 크기를 Application에서 입력해줘야 한다. 이를 쉽게 계산해주기 위한 도움을 주는 함수로 optixAccelComputeMemoryUsage 함수가 존재한다.
OptixAccelEmitDesc는 가속 구조를 compact하기 위한 구조체이다. 만약 compact할 필요가 없다면 optixAccelBuild의 인자로 0을 넣으면 된다.
type은 build 중에 어떤 값을 내보낼지에 대한 설정이다.
여기서는 compacted AS의 size를 출력한다. compactedSizeBuffer로 출력되게끔 result를 통해 설정한다.
compacted AS의 size를 출력할 때 uint64로 출력되므로 compactedSizeBuffur의 크기도 sizeof(uint64_t)로 설정했다.
주의할 점은 항상 compacted AS가 기존 AS보다 크게 압축된다는 보장이 없다. 따라서 compactedSizeBuffer의 결과와 기존 AS의 크기를 비교해서 큰 차이가 없다면 compacted AS를 update하는 것이 비효율 적이다.
위에서 말했듯이 optixAccelBuild를 위해서는 임시 메모리와 AS를 위한 메모리 공간이 필요하다. 이를 초기화하고, optixAccelBuild를 통해서 AS를 build 한다.
optixAccelBuild는 asynchronize하므로 CUDA_SYNC_CHECK()를 통해서 cudaDeviceSynchronize()를 호출하고 오류를 체크한다.
그리고 AS를 optixAccelCompact 함수를 통해서 compact하고 handle을 반환하며 종료한다.
마찬가지로 optixAccelCompact 함수는 asynchronize하므로 CUDA_SYNC_CHECK() 매크로를 통해 동기화한다.
// SampleRenderer.cpp
/*! set camera to render with */
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));
}Perspective view에 맞게끔 ray를 발사하기 위해 최대 가로, 세로를 계산해준다. 이전 Example03에서는 orthogonal이며, 2D였으므로 이러한 계산이 필요없었다.
// SampleRenederer.cpp
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);
}resize 함수에서 setCamera 함수가 설정된다.
// devicePrograms.cu
#include <optix_device.h>
#include "LaunchParams.h"
using namespace osc;
namespace osc {
/*! launch parameters in constant memory, filled in by optix upon
optixLaunch (this gets filled in from the buffer we pass to
optixLaunch) */
extern "C" __constant__ LaunchParams optixLaunchParams;
// for this simple example, we have a single ray type
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 ) );
}
//------------------------------------------------------------------------------
// closest hit and anyhit programs for radiance-type rays.
//
// Note eventually we will have to create one pair of those for each
// ray type and each geometry type we want to render; but this
// simple example doesn't use any actual geometries yet, so we only
// create a single, dummy, set of them (we do have to have at least
// one group of them to set up the SBT)
//------------------------------------------------------------------------------
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 */ }
//------------------------------------------------------------------------------
// miss program that gets called for any ray that did not have a
// valid intersection
//
// as with the anyhit/closest hit programs, in this example we only
// need to have _some_ dummy function to set up a valid SBT
// ------------------------------------------------------------------------------
extern "C" __global__ void __miss__radiance()
{
vec3f &prd = *(vec3f*)getPRD<vec3f>();
// set to constant white as background color
prd = vec3f(1.f);
}
//------------------------------------------------------------------------------
// ray gen program - the actual rendering happens in here
//------------------------------------------------------------------------------
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;
}
} // ::oscpayload는 32bit uint만 넘겨줄 수 있다. 이 한계를 넘기 위해서 pointer "주소"를 32bit uint 형태로 변환해서 payload로 넘기는 방식을 사용할 수 있다.
packPointer 함수는 64bit 주소를 32bit uint 2개로 분할하여 저장한다.
unpackPointer 함수는 32bit uint 2개를 가지고 64bit 주소로 다시 변환해준다.
__forceinline__은 컴파일러에게 재귀적으로 inline 하게 지시한다.
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 ) );
}getPRD는 현재 ray의 결과로 나올 payload의 현재 값들을 64bit uint로 갖고 오는 것이다.
이때 optixGetPayload_*로 payload 값을 32bit uint로 갖고온다.
여기서 Payload의 수는 OptixPipelineCompileOptions에서 설정했다.
총 0~31까지 32개의 payload를 불러올 수 있다.
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;
}u0, u1은 반환될 payload의 값이 저장된다.
payload는 32bit uint로만 전달할 수 있기 때문에 vec3f pixelColorPRD의 값으로 u0과 u1의 값을 0으로 초기화해준다.
screen(화면)은 실제로는 박스형태의 grid로 discreate하다. 따라서 박스의 중앙에서 ray를 발사하기 위해서 0.5만큼 더하고 frame의 크기만큼 나누어 0~1로 정규화를 한다.
그리고 camera to screen로의 ray direction을 계산해준다.
이때 -0.5f는 0~1로 정규화 했기 때문이다.
(xy평면에 screen과 camera를 projection하면, screen의 중앙에 camera가 위치하기 때문에)
optixTrace를 통해서 실제로 ray를 발사한다.
tmin, tmax는 ray의 최소, 최대 길이를 얘기한다.
rayTime은 motion blur에서의 시점 t를 얘기한다. 이는 motion blur가 활성화되어 있지 않으면 무시된다.
visibilityMask는 IAS에서 설정된 것으로 같은 비트를 공유하면 collision 되고, 공유하지 않으면 무시된다.
OPTIX_RAY_FLAG_DISABLE_ANYHIT는 any-hit program을 비활성화 한다.
(이는 IAS에서의 flag OPTIX_INSTANCE_FLAG_ENFORCE_ANYHIT를 override한다.)
// devicePrograms.cu
// for this simple example, we have a single ray type
enum { SURFACE_RAY_TYPE=0, RAY_TYPE_COUNT };이 예제에서는 hit program group과 miss program group이 하나 밖에 없기 때문에 0으로 설정된다. 이는 miss program group과 hit program group을 선택할 수 있는 기능이다.
이후 결과로 얻어온 색 정보를 ray가 맡은 위치에 저장한다.
// devicePrograms.cu
extern "C" __global__ void __closesthit__radiance()
{
const int primID = optixGetPrimitiveIndex();
vec3f &prd = *(vec3f*)getPRD<vec3f>();
prd = gdt::randomColor(primID);
}
extern "C" __global__ void __miss__radiance()
{
vec3f &prd = *(vec3f*)getPRD<vec3f>();
// set to constant white as background color
prd = vec3f(1.f);
}closesthit은 ray가 처음 hit 했을 때 callback되는 함수이다. 해당 값을 random color 설정한다.
optixGetPrimitiveIndex는 부딪힌 primitive의 BuildInput에서 넣었던 index를 반환한다.
miss는 ray가 hit가 한 번도 되지 않을 시에 색을 하얀색으로 한다.
