Simple GPU Path Tracing, Part. 2.1 : Acceleration structure

 
So now that we have a scene representation, let's create an acceleration structure out of it that will be optimized for fast ray tracing on the gpu.
There are many possibilities for creating an acceleration structure. The one I chose and that is commonly used for path tracing is Bounding Volume Hierarchy. It's a tree-based spatial data structure that divides 3d space into a tree, based on the bounding volume of each shape.



The BVH

The implementation of the BVH comes from the excellent series of blog posts by Jacco Bikker.
I can't recommend them enough to really understand how BVHs work.
There are lots of other resources on BVH, I'll just add the amazing pbrt chapter on that.

I won't go into details in the implementation of the BVH, as it's quite literally the same as Jacco Bikker's, and you can read all about it in his blog posts.
I'll document a bit the BVH code that you can find here.

Bottom line is that we can now create a BVH and all its gpu buffers with the function 
std::shared_ptr<sceneBVH> CreateBVH(std::shared_ptr<scene> Scene)

so that's what we do after we create the scene : 
Scene = CreateCornellBox();
BVH = CreateBVH(Scene);


Ray Tracing

So now that we have a BVH and all its associated gpu buffers, we can start ray tracing in a gpu kernel !
First, we will have to bind all those gpu buffers to the kernel.
For OpenGL, that will happen with binding functions. For Cuda, they will just get passed as arguments to the kernel.
 

OpenGL

We use the SetSSBO() function of shaderGL to bind the buffers to the desired binding points in the shader : 
void shaderGL::SetSSBO(std::shared_ptr<bufferGL> Buffer, int BindingPoint)
{
    glBindBuffer(GL_SHADER_STORAGE_BUFFER, Buffer->BufferID);
    glBindBufferBase(GL_SHADER_STORAGE_BUFFER, BindingPoint, Buffer->BufferID);
}

Once that is done, we can bind all the BVH and camera gpu buffers to the kernel before running it in the App.cu udpate loop :
PathTracingShader->Use();
PathTracingShader->SetTexture(0, RenderTexture->TextureID, GL_READ_WRITE);
PathTracingShader->SetSSBO(BVH->TrianglesBuffer, 1);
PathTracingShader->SetSSBO(BVH->TrianglesExBuffer, 2);
PathTracingShader->SetSSBO(BVH->BVHBuffer, 3);
PathTracingShader->SetSSBO(BVH->IndicesBuffer, 4);
PathTracingShader->SetSSBO(BVH->IndexDataBuffer, 5);
PathTracingShader->SetSSBO(BVH->TLASInstancesBuffer, 6);
PathTracingShader->SetSSBO(BVH->TLASNodeBuffer, 7);        
PathTracingShader->SetSSBO(Scene->CamerasBuffer, 8);
PathTracingShader->Dispatch(Window->Width / 16 + 1, Window->Height / 16 +1, 1);

Cool! Now we can access all these datas in the gl shader, we just need to declare them at the start.
I'll add them into the Inputs.glsl file that already contained the input texture.

layout(std430, binding = 8) buffer CameraBuffer {
  camera Cameras[];
};

// BVH
layout(std430, binding = 1)  buffer triangleBuffer
{
    triangle TriangleBuffer[];
};

layout(std430, binding = 2)  buffer triangleExBuffer
{
    triangleExtraData TriangleExBuffer[];
};

layout(std430, binding = 3)  buffer bvhBuffer
{
    bvhNode BVHBuffer[];
};

layout(std430, binding = 4)  buffer indicesBuffer
{
    uint IndicesBuffer[];
};

layout(std430, binding = 5)  buffer indexDataBuffer
{
    indexData IndexDataBuffer[];
};

layout(std430, binding = 6)  buffer tlasInstancesBuffer
{
    bvhInstance TLASInstancesBuffer[];
};

layout(std430, binding = 7)  buffer tlasNodes
{
    tlasNode TLASNodes[];
};

Ray tracing the BVH

Great stuff, Now we can finally start tracing rays! 
So first of all, we need to generate a ray based on the pixel we're on.
We basically generate a ray as if the camera was positioned at the origin, and looking forward. We then transform the origin and direction vectors by the camera transform matrix, and that's it !
That's some quite basic vector maths that I'm not going to explain here.
 
ray GetRay(vec2 ImageUV)
{
    camera Camera = Cameras[0];

    // Point on the film
    vec3 Q = vec3(
        (0.5f - ImageUV.x),
        (ImageUV.y - 0.5f),
        1
    );
    vec3 RayDirection = -normalize(Q);
    vec3 PointOnLens = vec3 (0,0,0);

    //Transform the ray direction and origin
    ray Ray = MakeRay(
        TransformPoint(Camera.Frame, PointOnLens),
        TransformDirection(Camera.Frame, RayDirection),
        vec3(0)
    );
    return Ray;
}
 
For now, we assume the camera film is 1 unit away from the camera. We calculate a point on that film based on UV coordinates, and that will define our direction.
We will make a more advanced ray generation model based on physical camera later.
 
Now that we have a ray, we are ready to test intersections on the BVH.
We will use a sceneIntersection structure that will be passed around by the BVH routines, and that will store the hit informations : 
struct sceneIntersection
{
    float Distance;
    uint InstanceIndex;
    uint PrimitiveIndex;
    float U;
    float V;
};
 
It stores the distance of the hit point from the camera, and some indices of the shape and of the triangle that was hit.
The U and V fields represent the hit position as barycentric coordinates of the triangle that was hit.
Barycentric coordinates of a triangle are a way of defining a point on a triangle.
It's usually composed of 3 coordinates, each one being a weight for one of the 3 vertices of the triangle.

Here, you see that we can find the position of any point inside the triangle by multiplying each vertex of the triangle with numbers a, b, and c. Those 3 abc weights are the barycentric coordinates.

This will allow us to interpolate the attributes of each of the 3 vertices to get an exact value at the position where the ray hit the triangle.

Now that we have a ray, we're going to test intersections of the ray against the top level acceleration structure of the bvh. this will return all the instances that we're possibly going to hit with our ray.
Then, we can intersect the bottom level acceleration structures of those instances, and find the closest one from the camera. We then fill in the sceneIntersection structure with the hit information that we can then use for shading. That sounds quite simple, but the intersection routines can be a bit complicated!

TLAS intersection
The following function is the main function to call for intersecting a ray with the scene.
void IntersectTLAS(ray Ray, INOUT(sceneIntersection) Isect) 

it's using a stack-based tree traversal algorithm. It will start intersecting the ray with the bounding box of the root of the tlas.
If it hits it, it will try intersecting the ray with its 2 child's bounding boxes. 
If it hits one of the 2 bounding boxes, it will add it to the stack, and visit it later.
If it hits the 2 bounding boxes, it will add only the closest one to the stack and visit it later.
If it doens't hit any of the 2 bounding boxes, we decrement the stack index and we can visit the next node.

Eventually, when we find a node that's a leaf node (Meaning it has no children), we will intersect the instance that's contained in the node. Our bvh model can only store 1 instance per node.
 
Here's the code of IntersectTLAS.

Instance Intersection
 
To intersect an instance, we transform the ray to the instance local space, and use that ray to intersect the blas (=bottom level acceleration structure) of the instance.
The blas intersection function works very similarly to the tlas, except that when it's reaching a leaf node, it will test the intersection with the triangles that are contained in the node.

When the ray hits a triangle, the RayTriangleIntersection function will fill in the sceneIntersection structure with all the needed information.

That's pretty much all we need to know for bvh intersection, again I invite you to go read the excellent blog post series by Jacco Bikker here.

Basic Ray Tracing

So let's cast a ray in the scene, and set the pixel colour with the colour of the shape that we hit : 
        sceneIntersection Isect;
        Isect.Distance = 1e30f;

        IntersectTLAS(Ray, Isect);
        if(Isect.Distance < 1e30f)
        {
            triangleExtraData ExtraData = TriangleExBuffer[Isect.PrimitiveIndex];    
            vec4 Colour =
                ExtraData.Colour1 * Isect.U +
                ExtraData.Colour2 * Isect.V +
                ExtraData.Colour0 * (1 - Isect.U - Isect.V);
            imageStore(RenderImage, ivec2(GLOBAL_ID()), Colour);
        }

We first create a sceneIntersection object, and initialize its distance with a high value.
Then, we check intersections with the scene by calling IntersectTLAS.
if the distance returned in the sceneIntersection object is less than the high value we set, it means that it's been updated during the intersection test, meaning that the ray hit something!

To find the colour of the shape we hit, we have to use the PrimitiveIndex of the sceneIntersection object. This will store the triangle index in the big triangles buffer that contains all the triangles in the scene.
We can fetch triangle data with this PrimitiveIndex value, here we're fetching the colour of each vertex of the triangle that was hit. We can then interpolate between the 3 colours of the 3 vertices using the barycentric coordinates to get the exact colour at the hit position.

And that's it ! We can now view the scene in the most basic way, just displaying the colour of each shape in the scene : 

 
 
Amazing ! It's quite an ugly looking Cornell box, but hey it's something. I made up some tones of gray for the walls and the boxes, just so that we can distinguish the objects more clearly.

Cuda

Now we also need to make some changes in order to get the cuda back end to work as well.
First, we need to pass the buffers as arguments to the cuda kernel call : 
TraceKernel<<<gridSize, blockSize>>>
((glm::vec4*)RenderBuffer->Data, Window->Width, Window->Height,
(triangle*)BVH->TrianglesBuffer->Data, (triangleExtraData*) BVH->TrianglesExBuffer->Data,
(bvhNode*) BVH->BVHBuffer->Data, (u32*) BVH->IndicesBuffer->Data, 
(indexData*) BVH->IndexDataBuffer->Data, (bvhInstance*)BVH->TLASInstancesBuffer->Data,
(tlasNode*) BVH->TLASNodeBuffer->Data, (camera*)Scene->CamerasBuffer->Data);

 
Then, we need to change a bit our macros in PathTrace.cu to accept those arguments : 
 
__device__ u32 Width;
__device__ u32 Height;
__device__ triangle *TriangleBuffer;
__device__ triangleExtraData *TriangleExBuffer;
__device__ bvhNode *BVHBuffer;
__device__ u32 *IndicesBuffer;
__device__ indexData *IndexDataBuffer;
__device__ bvhInstance *TLASInstancesBuffer;
__device__ tlasNode *TLASNodes;
__device__ camera *Cameras;

#define MAIN() \
__global__ void TraceKernel(glm::vec4 *RenderImage, int _Width, int _Height, \
                            triangle *_AllTriangles, triangleExtraData *_AllTrianglesEx, bvhNode *_AllBVHNodes, u32 *_AllTriangleIndices, indexData *_IndexData, bvhInstance *_Instances, tlasNode *_TLASNodes,\
                            camera *_Cameras)

#define INIT() \
    Width = _Width; \
    Height = _Height; \
    TriangleBuffer = _AllTriangles; \
    TriangleExBuffer = _AllTrianglesEx; \
    BVHBuffer = _AllBVHNodes; \
    IndicesBuffer = _AllTriangleIndices; \
    IndexDataBuffer = _IndexData; \
    TLASInstancesBuffer = _Instances; \
    TLASNodes = _TLASNodes; \
    Cameras = _Cameras; \ 
 
 
So we add all those buffers in the global scope, we add the arguments in the MAIN() function, and then we initialize the global scope variables in the INIT() macro.

There are 2 more macros that we need to use : 
- One for declaring functions, because in cuda kernels, all the functions that are used from a kernel need to be declared with a __device__ prefix. We create a FN_DECL macro that will just do that, and that we will use at the start of each function definition : 
#define FN_DECL __device__
 
Example of function definition with this macro : 
FN_DECL vec3 TransformDirection(mat4 A, vec3 B)
{
    vec4 Res = A * vec4(B, 0);
    return normalize(vec3(Res));
}

- Another one for using inout arguments inside functions. In glsl, we can use the inout keyword in the function declaration, but in cuda, we will pass the variables by reference to achieve the same result.

So here is this macro for cuda :
#define INOUT(Type) Type &
and for glsl :
#define INOUT(Type) inout Type

Annnd that's it ! we don't have to make any other change, the PathTracingCode will work for both backends.

Next time, we're going to write a very simple path tracer, which is very exciting ! we will finally get some pretty pixels after all this time.

Links

 

Commentaires

Enregistrer un commentaire

Articles les plus consultés