Simple GPU Path Tracing, Part. 4.1 : Textures

 

In this post, we will be adding the ability to sample textures in our path tracer, and we'll be able to use them for colour, normal mapping, metallic/roughness mapping, and emission mapping.

I'll first show how to load texture data from files, then we will be creating the gpu objects to access those textures from the gpu, and then we will see how to use them in the path tracer.

 


Here's the commit for this post.

Texture Loading

First, we will create a texture struct, and we will store an array of them in our scene, that each material can then reference.We will force uint8 rgba textures for all the textures in the scene.
Here's the struct : 
struct texture
{
    i32 Width = 0;
    i32 Height = 0;
    i32 NumChannels = 0;
    std::vector<u8> Pixels = {};

    void SetFromFile(const std::string &FileName, i32 Width = -1, i32 Height = -1);
    void SetFromPixels(const std::vector<uint8_t> &PixelData, i32 Width = -1, i32 Height = -1);
};

And let's add it to the scene : 
struct scene
{
   
    std::vector<camera> Cameras = {};
    std::vector<instance> Instances = {};
    std::vector<shape> Shapes = {};
    std::vector<material> Materials = {};
    std::vector<texture> Textures = {};
... 

Ok great, now lets see the content of SetFromFile : 
void texture::SetFromFile(const std::string &FileName, int Width = -1, int Height = -1)
{
    int NumChannels=4;
    ImageFromFile(FileName, this->Pixels, Width, Height, NumChannels);
    this->NumChannels = this->Pixels.size() / (Width * Height);
    this->Width = Width;
    this->Height = Height;
}

We call the ImageFromFile function, and here's its body : 

void ImageFromFile(const std::string &FileName, std::vector<uint8_t> &Data, 
                   int &Width, int &Height, int &NumChannels)
{
    int ImgWidth, ImgHeight;
    int channels;

    // Load the image
    stbi_uc* Image = stbi_load(FileName.c_str(), &ImgWidth, &ImgHeight, &channels, 
                               STBI_rgb_alpha);

    if (Image == nullptr) {
        // Handle error (file not found, invalid format, etc.)
        std::cout << "Failed to load image: " << FileName << std::endl;
        return;
    }

    // If there's a requested size, set target size with it. Otherwise, use the image size
    int TargetWidth = (Width != 0) ? Width : ImgWidth;
    int TargetHeight = (Height != 0) ? Height : ImgHeight;

    // If the target size is not the current size, resize the image
    if(TargetWidth != ImgWidth || TargetHeight != ImgHeight)
    {
        // Resize the image using stbir_resize (part of stb_image_resize.h)
        stbi_uc* ResizedImage = new stbi_uc[Width * Height * 4]; // Assuming RGBA format

        int result = stbir_resize_uint8(Image, ImgWidth, ImgHeight, 0, ResizedImage,  
                                        TargetWidth, TargetHeight, 0, 4);
       
        stbi_image_free(Image);

        if (!result) {
            // Handle resize error
            std::cout << "Failed to resize image: " << FileName << std::endl;
            delete[] ResizedImage;
            return;
        }

        // Resize the pixel data, and copy to it
        Data.resize(TargetWidth * TargetHeight * 4);
        memcpy(Data.data(), ResizedImage, TargetWidth * TargetHeight * 4);
        delete[] ResizedImage;
    }
    else
    {
        // Resize the pixel data, and copy to it
        Data.resize(TargetWidth * TargetHeight * 4);
        memcpy(Data.data(), Image, TargetWidth * TargetHeight * 4);
        delete[] Image;
    }
}    

Quite simple usue of stb_image library. Note that we can pass a target size argument to the function, and the function will resize the image to this target size if needed.
This is important because we will be requiring all the textures to be the same size in the scene.
For that, let's add a textureWidth and textureHeight fields in the scene struct :

    u32 TextureWidth = 512;
    u32 TextureHeight = 512;

Cool, let's create some textures now !
In the CreateCornellBox function, I'll add the following : 
    texture &Texture = Scene->Textures.emplace_back();
    Texture.SetFromFile("resources/textures/Debug.jpg", Scene->TextureWidth,  
                        Scene->TextureHeight);
    Scene->TextureNames.push_back("Debug");

    texture &Normal = Scene->Textures.emplace_back();
    Normal.SetFromFile("resources/textures/Normal.jpg", Scene->TextureWidth,  
                        Scene->TextureHeight);
    Scene->TextureNames.push_back("Normal");
   
    texture &Roughness = Scene->Textures.emplace_back();
    Roughness.SetFromFile("resources/textures/Roughness.jpg", Scene->TextureWidth,  
                           Scene->TextureHeight);
    Scene->TextureNames.push_back("Roughness");

Texture GPU objects

Ok now we've got our textures loaded, let's see how to use them on the gpu.
We already have a texture class for openGL, but unfortunately it will not be helpful here. Indeed, in our path tracing kernel, we need a way of accessing all the textures in the scene at any time because any object using textures can be hit. If we start having a lot of textures in the scene, we're not going to be binding each individual oneat every pass.

What we'll do instead is store all the textures of the scene in a single object that we can then sample from the gpu kernel. This object will be different between openGL and cuda, so I'll explain the 2 now : 

OpenGL

it's quite straightforward, we will be using a texture2DArray object that allows to store multiple texture2D, and to sample them in a shader.
For that I'll create a simple class wrapper : 


class textureArray {
public:
    textureArray();
    ~textureArray();
    void CreateTextureArray(int width, int height, int layers);
    void LoadTextureLayer(int layerIndex, const std::vector<uint8_t>& imageData, 
                          int width, int height);
    void Bind(int textureUnit = 0);
    void Unbind() const;

    GLuint TextureID;
};    

Very simple stuff, here's just the function to create the array and to load a layer into it : 
    void textureArray::CreateTextureArray(int Width, int Height, int Layers) {
        glGenTextures(1, &TextureID);
        glBindTexture(GL_TEXTURE_2D_ARRAY, TextureID);

        // Set texture parameters
        glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
        glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
        glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
        glTexParameteri(GL_TEXTURE_2D_ARRAY, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);

        // Allocate storage for the texture array
        glTexImage3D(GL_TEXTURE_2D_ARRAY, 0, GL_RGBA, Width, Height, Layers, 0, 
                     GL_RGBA, GL_UNSIGNED_BYTE, nullptr);
        glBindTexture(GL_TEXTURE_2D_ARRAY, 0);
    }

    void textureArray::LoadTextureLayer(int layerIndex, 
                                        const std::vector<uint8_t>& imageData, 
                                        int Width, int Height) {
        glBindTexture(GL_TEXTURE_2D_ARRAY, TextureID);
        glTexSubImage3D(GL_TEXTURE_2D_ARRAY, 0, 0, 0, layerIndex, Width, Height, 1,  
                        GL_RGBA, GL_UNSIGNED_BYTE, imageData.data());
        glBindTexture(GL_TEXTURE_2D_ARRAY, 0);
    }
 
Ok, now we can store a textureArray in our scene, and fill it with all the scene textures. We will create a function in scene called ReloadTextureArray, that will just do that, and that we will call at the end of our CreateCornellBox() function.
void scene::ReloadTextureArray()
{
#if API==API_GL
    TexArray = std::make_shared<textureArrayGL>();
#elif API==API_CU
   
#endif

    TexArray->CreateTextureArray(TextureWidth, TextureHeight, Textures.size());
    for (size_t i = 0; i < Textures.size(); i++)
    {
        TexArray->LoadTextureLayer(i, Textures[i].Pixels, TextureWidth, TextureHeight);
    }
}
 
 Annnd that's all for openGL !
 

Cuda

Things are a bit more challenging for cuda, as there's no texture array object.
what we could do is use a giant buffer that will store all the texture data, and sample from it in the shader, but we will loose the ability to use texture cache that's very fast at sampling textures in a 2d maner. So what we will do instead is create a giant texture, and store all our small textures inside, like that : 
 
The maximum texture size of most recent gpus is 8192x8192, which means that if we're using 512x512 textures, we can store 16x16=256 textures inside. That will be ok for now.
 
So let's create a textureArrayCu class that will help us doing that : 
 

class textureArrayCu {
public:
    textureArrayCu() = default;
    ~textureArrayCu();
    void CreateTextureArray(int Width, int Height, int Layers);
    void LoadTextureLayer(int layerIndex, const std::vector<uint8_t>& ImageData, int Width, int Height);

    int TotalWidth = 8192;
    int TotalHeight = 8192;
    size_t Pitch;    
    cudaArray* CuArray;
    cudaTextureObject_t  TexObject;
    int Width, Height;
};
 
And there's the content of the 2 important functions : 

    void textureArrayCu::CreateTextureArray(int Width, int Height, int Layers) {
        this->Width = Width;
        this->Height = Height;

        cudaMallocPitch((void**)&CuArray, &Pitch,  TotalWidth*sizeof(uchar4), TotalHeight);
        struct cudaResourceDesc resDesc;
        memset(&resDesc, 0, sizeof(resDesc));
        resDesc.resType = cudaResourceTypePitch2D;
        resDesc.res.pitch2D.devPtr = CuArray;
        resDesc.res.pitch2D.width = TotalWidth;
        resDesc.res.pitch2D.height = TotalHeight;
        resDesc.res.pitch2D.desc = cudaCreateChannelDesc<uchar4>();
        resDesc.res.pitch2D.pitchInBytes = Pitch;
        struct cudaTextureDesc texDesc;
        memset(&texDesc, 0, sizeof(texDesc));
        cudaCreateTextureObject(&TexObject, &resDesc, &texDesc, NULL);        
    }

    void textureArrayCu::LoadTextureLayer(int layerIndex, 
                             const std::vector<uint8_t>& ImageData, 
                             int Width, int Height) {
        static int LayersPerRow = TotalWidth / Width;
        int DestInxX = layerIndex % LayersPerRow;
        int DestInxY = layerIndex / LayersPerRow;
        uint32_t DestX = DestInxX * Width;
        uint32_t DestY = DestInxY * Height;
        uint32_t Dest = (DestY * TotalWidth + DestX) * 4;


        cudaMemcpy2D((uint8_t*)CuArray + Dest, Pitch, ImageData.data(),  
                     512*sizeof(uchar4), 512*sizeof(uchar4), Height, 
                     cudaMemcpyHostToDevice);
    }
 
In CreateTextureArray, we simply create a cudaTextureObject_t with the big texture size.
in LoadTextureLayer, we create a mapping between the passed in layerIndex and the location where it will go in the big texture grid, and copy the data.
 

Using the textures in the path tracer

First, we need to add the texture references into the material struct : 
struct material
{
    glm::vec3 Emission = {};
    float Roughness = 0;
   
    glm::vec3 Colour = {};
    float Metallic = 0;
   
    glm::ivec3 Padding;
    int MaterialType = 0;

    int EmissionTexture = InvalidID;
    int ColourTexture = InvalidID;
    int RoughnessTexture = InvalidID;
    int NormalTexture = InvalidID;
};
 
And we also need to bind those textures to the kernels : 
 
for openGL, that's the line that does that in the Trace() function : 
RTShader->SetTextureArray(Scene->TexArray, 13, "SceneTextures");

For cuda, we do as usual, simply pass the textureArray as an argument to the kernel, add the textureObject into the global scope, and add one line in the INIT() macro to initialize it.

Colour textures

Let's use one of the textures on the back wall : 
BackWallMaterial.ColourTexture = 0;
 
We will now need to evaluate this texture in the kernels.

in the EvalMaterial() function, we will sample the texture at the hit texture coordinate, and modulate the colour with the result. 
First, we need to find the texture coordinates of the hit point : 
FN_DECL vec2 EvalTexCoord(INOUT(sceneIntersection) Isect)
{
    uint Element = Isect.PrimitiveIndex;
    triangleExtraData ExtraData = TriangleExBuffer[Isect.PrimitiveIndex];
    return
        ExtraData.UV1 * Isect.U +
        ExtraData.UV2 * Isect.V +
        ExtraData.UV0 * (1 - Isect.U - Isect.V);    
}
 
Then, we need a function that will sample the texture at a particular coordinate : 

FN_DECL vec4 EvalTexture(int Texture, vec2 UV, bool Linear)
{
    if(Texture == INVALID_ID) return vec4(1, 1, 1, 1);
    vec3 texCoord3D = vec3(UV, Texture);
    vec4 Colour = texture(SceneTextures, texCoord3D);
    if(Linear) Colour = ToLinear(Colour);
    return Colour;
}
Here, note that there's a Linear parameter that tells the function wether to transform from sRGB space to linear space. It's typically the case for colour textures, but not for normal maps for example.
 
Ok so this will work for openGL, but cuda doesn't have a texture() function, and moreover we're not really using a textureArray in cuda, so the 3d texcoord doesn't even apply.
We have to implement the texture() function ourselves in PathTrace.cu, and here's how it looks : 
__device__ vec4 texture(cudaTextureObject_t _SceneTextures, vec3 Coords)
{
    static int NumLayersX = 8192 / 512;
    int LayerInx = Coords.z;
   
    int LocalCoordX = Coords.x * 512;
    int LocalCoordY = Coords.y * 512;

    int XOffset = (LayerInx % NumLayersX) * 512;
    int YOffset = (LayerInx / NumLayersX) * 512;

    int CoordX = XOffset + LocalCoordX;
    int CoordY = YOffset + LocalCoordY;

    uchar4 TexValue = tex2D<uchar4>(_SceneTextures, CoordX, CoordY);
    vec4 TexValueF = vec4((float)TexValue.x / 255.0f, 
                          (float)TexValue.y / 255.0f, 
                         (float)TexValue.z / 255.0f, 
                          (float)TexValue.w / 255.0f);
    return TexValueF;
}

Here, we map from the 3d coordinates to the appropriate position in the big 2d cuda texture object.

In the end, here's the EvalMaterial function with colour textures : 
FN_DECL materialPoint EvalMaterial(INOUT(sceneIntersection) Isect)
{
    material Material = Materials[Isect.MaterialIndex];
    materialPoint Point;

    vec2 TexCoords = EvalTexCoord(Isect);
    vec4 ColourTexture = EvalTexture(Material.ColourTexture, TexCoords, true);

    Point.MaterialType = Material.MaterialType;
    Point.Colour = Material.Colour * vec3(ColourTexture);
    Point.Emission = Material.Emission;
    Point.Roughness = Material.Roughness;
    Point.Roughness = Point.Roughness * Point.Roughness;
    Point.Metallic = Material.Metallic;
    return Point;
}

We can very easily also add support for metallic/roughness textures, and emission textures : 

    vec2 TexCoords = EvalTexCoord(Isect);
    vec4 EmissionTexture = EvalTexture(Material.EmissionTexture, TexCoords, true);    
    vec4 ColourTexture = EvalTexture(Material.ColourTexture, TexCoords, true);
    vec4 RoughnessTexture = EvalTexture(Material.RoughnessTexture, TexCoords, false);
   
    Point.MaterialType = Material.MaterialType;
    Point.Colour = Material.Colour * vec3(ColourTexture);
    Point.Emission = Material.Emission * vec3(EmissionTexture);
   
    Point.Metallic = Material.Metallic * RoughnessTexture.z;
    Point.Roughness = Material.Roughness * RoughnessTexture.y;
    Point.Roughness = Point.Roughness * Point.Roughness;

 And that's it for today !
Here's the result : 
 
 

Commentaires

Enregistrer un commentaire

Articles les plus consultés