DEV Community: The Lone Engineer

🎨 Dynamic Texture Loading in NervForge: Async Promises in C++/WASM

The Lone Engineer — Mon, 30 Mar 2026 20:08:39 +0000

New short devlog on a neat quality-of-life improvement: textures in NervForge are now loaded on demand from a remote manifest, with a promise-based async system handling the download and material update.

Note: You can try NervForge directly in your browser: https://nervtech.org/nervforge/

The core idea: instead of bundling all textures upfront, a resource manifest lists available remote files. Selecting a new texture triggers an async download, shows a checkerboard placeholder in the meantime, then swaps in the real texture once ready — and caches it locally for future sessions.

What's covered:

Resource manifest pattern for remote asset discovery
Promise-based async texture creation API (in NervSDK)
Main-thread callback chaining for safe material updates
Local caching strategy to avoid redundant downloads
Bonus: horizon-level black artifact fix on planet rendering

The promise system itself is open source in NervSDK — worth a look if you need lightweight async task handling in a C++/WASM context.

You can check the full DevLog at: https://wiki.nervtech.org/doku.php?id=blog:2026:0330_nvl_dyn_texture_loading

🌳 Building Realistic 3D Bushes: Optimizing Procedural Generation in NervForge

The Lone Engineer — Sat, 14 Mar 2026 09:14:59 +0000

Just published a devlog article on taking NervForge's bush generation from obviously computer-generated to more natural looking plant with a couple of major changes (plus other fixes in the process):

Note: You can also try it yourself directly from here: https://nervtech.org/nervforge/

The biggest challenge? Implementing self-repulsion so branches don't intersect even at root level, without killing performance. Initial naive approach: 23x slower (16ms → 364ms). Final optimized version with deferred insertion levels and center-of-mass approximation: 58ms.

What's covered:

Spawn area distribution (disc/square) for natural root spreading
Architectural refactoring to eliminate widget dependencies
Deferred repulsion levels: real-time for visible structure, batched for details
Spatial grid optimization using weighted center-of-mass per cell
Lock contention solutions for multi-threaded branch generation
Performance profiling and iterative optimization strategies

The video walks through the actual problems encountered, shows the broken versions, and explains the exact solutions implemented. Code snippets, performance comparisons, and visual before/after examples throughout.

Perfect if you're working on procedural generation systems, spatial partitioning, or real-time graphics optimization. (If you think about it, the spatial grid implementation strategy is applicable beyond just tree generation)

You can also check the full DevLog article at this url: https://wiki.nervtech.org/doku.php?id=blog:2026:0314_initial_bush_config

NervForge: Generate Beautiful 3D Trees in Your Browser (Free Tool)

The Lone Engineer — Fri, 06 Mar 2026 22:59:31 +0000

Need procedural trees for your game or 3D project? I've been working on NervForge, a browser-based tool that generates customizable 3D tree models and exports them as glTF files.

What It Does

NervForge gives you control over every aspect of your trees:

Branch levels - Build from simple trunks to complex multi-layered canopies
Natural distribution - Repulsion systems prevent branch collisions
Gnarliness controls - Add realistic twists and curves (because perfectly straight trees look fake)
Customizable foliage - Control leaf size, density, and distribution per branch level
Export to glTF - Drop your trees straight into Blender, Unity, or any 3D application

Why I Built This

I wanted a quick way to generate tree variations without leaving my browser. Adjust a few parameters, tweak the seed value, and you can create an entire forest of unique trees that share the same style.

Current Features

Multi-level branch hierarchy with density controls
Attraction point system for directing growth
Per-level resolution and radius settings
Leaf system with multiple textures
Seed-based generation for variations
Direct glTF export

Try It Yourself

The tool is live and free to use. I walk through all the features in this video tutorial:

Or you can also try it yourself directly from here: https://nervtech.org/nervforge/

What's Next

I'm actively developing NervForge with planned features including:

Proper bark textures (goodbye checkerboard!)
Improved pruning system
More leaf texture options
Additional customization parameters

Reference links:

Your Feedback Matters

This started as a toy project, but I'd love to make it genuinely useful for the gamedev and 3D art communities. If you try it out, let me know what features would make it more valuable for your workflow 😉

Drop any questions or suggestions in the comments below or on the YouTube video!

Side note: Have you built procedural generation tools? What's your approach to creating natural-looking 3D assets? Feel free to share your experience below! 😊

[Tutorial] Procedural Voronoi Texture generation in WGPU

The Lone Engineer — Sat, 11 Oct 2025 07:41:30 +0000

Introduction

Hi everyone! In this new article, I would like to reconsider Voronoi generation in WebGPU but this time from a procedural texture generation perspective as opposed to the original technique I mentioned in a previous youtube recording where I was rendering actual polygons and lines and running a full Delaunay triangulation stage on each frame. And which you can check here actually:

Here is a small overview of the final result we will achieve in this tutorial:

We will see how to display different colors for the different cells,
How to display a circular dot at the center of each cell,
And how to apply time animation on the full voronoi diagram to get this constantly moving effect.
We'll even see how we can dynamically change the size of the diagram
Yet before that of course we will need to cover some simple theory and preliminary computation steps to really understand how this algorithm works.

Before we start

Now before we dive any deeper, I have to mention here that I will roughly be following the same logic as in this ShaderToy Voronoi example by Inigo Quilez
But well, understanding this base code might not be so easy if you are not familiar with the algorithm yet, so we'll provide some explanations in the process.
Anyway, if you don't know Inigo Quilez yet, I would also highly suggest you pay a visit to his website and try to read as much as possible from there, specially the technical articles, because, honestly, this guy is a genius 😄!
Also, just like the last time, I have added a dedicated folder in the NervLandAdventures github repository for this tutorial where you will find the incremental versions of the compute shader we will create here in case you would like to have more careful look at these.

And now, let's get started!

Setting up our Display Surface

The first thing we need to do for this project is to setup a simple surface on which we will render our generated texture.

As before, I will be using a minimal scene construction inside the TerrainView app as our playground, and here is the YAML element I used for this:

  voronoi_surface:
    type: ComputeSurfaceBlueprint
    node: quads0
    transform: subject_pos
    anchor: bottom | center
    shader: texture/voronoi
    texture_width: 1920
    texture_height: 1080
    width: 6.0
    aspect: 16/9

The ComputeSurfaceBlueprint is a new type of element I just introduced, which can be used to specify an arbitrary Compute Shader that will generate the texture to display on the surface quad.

In this case for instance, we will use the shader from the file texture/voronoi.wgsl which we will get to in just a moment.

Since we are generating this texture from nothing, note that we also specify the dimension of the output texture in this config: the texture_width and texture_height

Note: that those texture dimensions are not directly related to the width and height parameters that can be provided to specify the physical size of the quad in the world.

In the ComputeSurfaceBlueprint class we then proceed in a similar way to what was done in the VideoSurfaceBlueprint described in our previous video tutorial (cf. https://www.youtube.com/watch?v=P1jxvLm6SwE): we will allocate a dedicated area on our texture atlas to hold the resulting texture data, and then we will setup some kind of processing to write valid data in that area, and this content will be shown on the final quad with the same mechanism as in the previous recording.

So, first we allocate an area with the appropriate size in our main texture atlas with this code:

    logDEBUG("Creating target BSP area of size {}x{}", _textureWidth,
             _textureHeight);
    auto area = bsp.add_area((U32)_textureWidth, (U32)_textureHeight);
    NVCHK(area.is_valid(), "Invalid BSP area.");

And then we create a ComputeNode with the provided shader, and setup the only binding group we will need for this shader execution:

    logDEBUG("Building compute node for {}", _shader);
    auto cnode = WGPUComputeNode::create(_shader);

    struct Params {
        Vec3u textureSize;
        U32 pad0{0};
        Vec3u origin;
        U32 pad1{0};
    };

    UniformBuffer<Params> params{
        Params{.textureSize = Vec3u(_textureWidth, _textureHeight, 1),
               .origin = Vec3u(area.rect.xmin, area.rect.ymin, area.layer)}};

    auto& tex = bsp.get_texture();

    auto bgrp0 = cnode->setup_bind_grp({
        params.as_ubo(),
        stoTex2dArray(tex),
    });

Of course, the code above is in the NervLand engine so not using pure native WGPU API, but anyway I think you should be able to grasp the main ideas:

The Params struct is used to provide a uniform struct on the shader side: as we need to provide it with the size of the texture we want to generate, and to origin point in the texture atlas where we should start writing the output data.
And next to build the bind group we provide those params as what I call a "uniform buffer object" in the engine, and the underlying texture 2D array from the texture atlas as a writable storage texture 2d array.

Next, we will prepare an actual dispatch for the compute node and the binding group we just created:

    // Need to figure out here how many dispatches to do:
    auto nx = (_textureWidth + 15U) / 16U;
    auto ny = (_textureHeight + 15U) / 16U;

    cnode->add({nx, ny}, {bgrp0});

    scene.get_compute_pass(id_prerender_cpass).add(*cnode);

You see that here we use the texture dimensions to compute the dispatch sizes, and we assume a workgroup size of 16x16x1 in the shader.

And finally we add this compute node to the "prerender compute pass", which is a specific compute pass in my engine executed once at the beginning of each render cycle and before any other rendering steps.

Initial compute shader

Now, let's take a look a the initial version of the voronoi shader:

struct ComputeParams {
    // Size of the texture to generate
    textureSize: vec3u,
    padding0: u32,
    // Origin point of the texture data in the output texture storage
    origin: vec3u,
    padding1: u32,
};

@group(0) @binding(0) var<uniform> params : ComputeParams;
@group(0) @binding(1) var outputTex: texture_storage_2d_array<rgba8unorm,write>;

First we declare the inputs we will use in this shader: as described before, we have a uniform struct "params" providing the texture size and destination origin point.

And as second binding we have the output texture 2D array, which is the full texture atlas, and we should write our computation results into a sub-region of that texture.

Next, I started with this version of the main function:

@compute @workgroup_size(16, 16, 1)
fn main(@builtin(global_invocation_id) id: vec3u) {

    if id.x > params.textureSize.x || id.y > params.textureSize.y {
        // Nothing to write in this case.
        return;
    }

    // Write yellow color:
    let color = vec4f(1.0, 1.0, 0.0, 1.0);

    let layer: u32 = params.origin.z;
    textureStore(outputTex, vec2i(params.origin.xy + id.xy), layer, color);
}

As already mentioned, we use a workgroup size of 16x16x1
We check the texture bounds [very important here as we may have other texture data after those bounds in the atlas and we should not override that]
And for now we simply write a uniform yellow color for all pixels.

If we run the TerrainView app with this scene setup, we get this initial result (which is just as expected 😃):

Then if I simply change the color to red in the shader and save the file

    let color = vec4f(1.0, 0.0, 0.0, 1.0);

I immediately get the quad display updated to red, which means the compute shader is indeed working properly and live reload feature is also functional:

As a final test, we can change the color opacity to 50%:

    let color = vec4f(1.0, 0.0, 0.0, 0.5);

And we will get this result:

So this also seems to work fine, (even if there is some artefact due to the interaction with the ocean post processing shader, but that's something I will fix later on so no worries ;-)).

And this concludes the initial setup of a render surface we needed for our experiment and now we can focus on the Voronoi generation process itself.

A little bit of theory

Yet, before we really get started, we need a little bit of theory on how we will proceed here:

The basic idea when generating Voronoi cells, is that for each pixel/location that we consider in a 2D area, we need to find the closest "reference point" to that pixel.

And then, we can declare that the pixel we are considering belongs to the cell of that specific "reference point" (which is in fact what we can also call the "cell center") and maybe give it a specific color or store its distance to the cell center.

Generally we are given a set of random "reference points" that can move anywhere, and thus it's pretty tricky to follow them all, and for each update cycle and each pixel you would need to reconsider all of them to find the closest.

Yet, we have a way to cheat a bit on this problem as follow:

First we will create a grid on top of all the pixels that we want to consider:

And then, for each of this grid cell we will generate a single reference point at a random location but inside the specific grid cell.

After that, whenever we consider a given pixel, the problem of finding the closest reference point is simplified: we only need to consider the reference points from the cell containing this pixel and all the direct cell neighbors, so that's only 9 reference points to check per pixel!

Now, let's see how we could apply that in our code:
First, we can compute the normalized UV coords on our render quad as follow:

    // Get the normalized UV coords from the pixel coords and texture size:
    let uv = vec2f(id.xy) / vec2f(params.textureSize.xy - 1u);

    // Write the UVs as color:
    let color = vec4f(uv, 0.0, 1.0);

And that will give us this display result:

Which is exactly what we could expect:

The origin of the texture is at the bottom left corner (0,0)
Then when we increment the U coordinate (ie. X/horizontal direction) we get more and more red.
And when we increment the V coordinate (ie. Y/Vertical direction) we get more and more green.

But in fact, if we want to keep the grid cell square, it's easier to work directly with the texture size in pixels and think in terms of "number of pixels" per grid cell. So we can rather use this code:

    // Number of pixels on each dimension to form a grid cell:
    let gridSize: f32 = 64.0;

    // Get our coordinates on the grid:
    let gcoords = vec2f(id.xy) / gridSize;

    // From those coords we can extract the integer part which
    // are the cell coords, and the fractional part which are our pixel
    // local (normalized) coords on the cell:
    let cellCoords = vec2i(floor(gcoords) + 0.5);
    let localCoords = gcoords - floor(gcoords);

    // Write the localCoords as color:
    let color = vec4f(localCoords, 0.0, 1.0);

And this will give us this visual result:

From there I started to think it would be cool to be able to dynamically manipulate the gridSize parameter when building this shader, so why
not add a dedicated control panel for it ?

Adding support for StructuredBuffer creation in config

To introduce this dynamic control support, I first added support to configure and share some WebGPU buffer from my scene config file:

  buffer0:
    type: StructuredBufferBlueprint
    usage: Uniform | CopyDst
    elements:
      - name: TextureSize
        type: Vec3u
      - name: GridSize
        type: F32
        value: 64.0
      - name: Origin
        type: Vec3u

Then I ensured I could get access to/update and use that buffer instead of creating a new one in the ComputeSurface construction process:

     // Instead of creating a buffer here we can retrieve the structured buffer
    // from the engine:
    auto& stbuf =
        WGPUEngine::instance()->structured_buffers().get(_uniformBufferId);
    stbuf.set_element(SID("TextureSize"),
                      Vec3u(_textureWidth, _textureHeight, 1));
    stbuf.set_element(SID("Origin"),
                      Vec3u(area.rect.xmin, area.rect.ymin, area.layer));
    stbuf.update();

    // struct Params {
    //     Vec3u textureSize;
    //     U32 pad0{0};
    //     Vec3u origin;
    //     U32 pad1{0};
    // };

    // UniformBuffer<Params> params{
    //     Params{.textureSize = Vec3u(_textureWidth, _textureHeight, 1),
    //            .origin = Vec3u(area.rect.xmin, area.rect.ymin, area.layer)}};

    auto& tex = bsp.get_texture();

    auto bgrp0 = cnode->setup_bind_grp({
        // params.as_ubo(),
        stbuf.buffer().as_ubo(),
        stoTex2dArray(tex),
    });

And finally I just had to update the ComputeParams struct in the shader file to also provide the gridSize from there:

struct ComputeParams {
    // Size of the texture to generate
    textureSize: vec3u,
    // Number of pixels on each dimension to form a grid cell:
    gridSize: f32,
    // Origin point of the texture data in the output texture storage
    origin: vec3u,
    padding1: u32,
};

And in the main function we can now just use that parameter instead of the hardcoded value:

    // Get our coordinates on the grid:
    let gcoords = vec2f(id.xy) / params.gridSize;

With this first part of the update, if I change the default value for the "gridSize" element to 256.0 for instance in the config file I get this reflected into the shader as expected on the next start:

Building control panel for buffer elements

After adding the support for structured buffers definition in the config, I then further extended the configuration to also support the description of a ControlPanel as follow:

  voronoi_panel:
    type: ControlPanelBlueprint
    icon: icons/settings-sliders.png
    title: Voronoi Settings
    elements:
      - type: SliderRow
        caption: Grid Size:
        text_precision: 2
        use_log_scale: false
        min_value: 1.0
        max_value: 512.0
        target:
          type: StructuredBuffer
          name: buffer0
          element: GridSize

In this element, we specify the ControlPanelBlueprint type, and then we build the panel with a natural flow:

Indicating the icon that should be used for the visibility toggling button of this panel
Then we provide the title to show on the panel
And then a list of elements to display on that panel.
For now we only have 1 element, which is a SliderRow type, allowing us to modify the value of the "GridSize" element of the buffer0 we created before.

With this new component when I run the app I can now dynamically control the GridSize value and the changes are immediately reflected in the shader 🥳:

Adding support for grid display

Okay, now, before moving to the generation the reference points I think we need to improve a bit on the default display we have:

Displaying the local cell UV coords is a bit confusing, and instead I would like to display a simple white background for each cell with black borders around them.

So let's implement this!

First I add a new entry in the buffer0 config to control the border size in pixels:

  buffer0:
    type: StructuredBufferBlueprint
    usage: Uniform | CopyDst
    elements:
      - name: TextureSize
        type: Vec3u
      - name: GridSize
        type: F32
        value: 64.0
      - name: Origin
        type: Vec3u
      - name: BorderSize
        type: F32
        value: 0.0

Next I add another line to control this value from the "Voronoi Settings" panel:

  voronoi_panel:
    type: ControlPanelBlueprint
    icon: icons/settings-sliders.png
    title: Voronoi Settings
    elements:
      - type: SliderRow
        caption: "Grid Size:"
        text_precision: 2
        use_log_scale: false
        min_value: 2.0
        max_value: 512.0
        target:
          type: StructuredBuffer
          name: buffer0
          element: GridSize
      - type: SliderRow
        caption: "Border Size:"
        text_precision: 0
        use_log_scale: false
        min_value: 0.0
        max_value: 8.0
        target:
          type: StructuredBuffer
          name: buffer0
          element: BorderSize

Then I update ComputeParams struct in our shader to reflect the changes we just made to the buffer:

struct ComputeParams {
    // Size of the texture to generate
    textureSize: vec3u,
    // Number of pixels on each dimension to form a grid cell:
    gridSize: f32,
    // Origin point of the texture data in the output texture storage
    origin: vec3u,
    // Border size for the grid display in number of pixels
    borderSize: f32,
};

And finally we use this new parameter to draw black borders on a white background when the provided value is not 0.0:

    // Write the localCoords as color:
    var color = vec4f(localCoords, 0.0, 1.0);

    // Add the border if specified:
    // Compute the "normalized border size" (as borderSize is in pixels):
    let nbsize = params.borderSize / params.gridSize;
    if nbsize > 0.0 {
        let isBorder = localCoords.x < nbsize || localCoords.x > (1.0 - nbsize) || localCoords.y < nbsize || localCoords.y > (1.0 - nbsize);
        let white = vec4f(1.0, 1.0, 1.0, 1.0);
        let black = vec4f(0.0, 0.0, 0.0, 1.0);
        color = select(white, black, isBorder);
    }

With those changes here is the result we get:

Added the reference points

And now, as promised, it's finally time to generate the reference points for each cell.

We will need one pseudo-random but fixed location in each cell, which can be easily achieved using an hash function applied for instance the coordinates of the bottom left corner of each cell as input.

Personally I was already using this hash function in other parts of my code to go from 3D integer to 1D noise:

fn hash(p: vec3i) -> f32 {
    // Better coefficients for 3D -> 1D
    // cf. https://www.reedbeta.com/blog/hash-functions-for-gpu-rendering/
    var n = p.x * 1597 + p.y * 3571 + p.z * 7919;

    var state: u32 = u32(n) * 747796405u + 2891336453u;
    state = ((state >> ((state >> 28u) + 4u)) ^ state) * 277803737u;
    state = (state >> 22u) ^ state;

    return -1.0 + 2.0 * f32(state) / f32(0xffffffffu);
}

And extending that a bit, here is the version Claude came up with for our 2D needs:


fn hash(p: vec2i) -> vec2f {
    // Generate two different hash values from the 2D input
    // Using different prime coefficients for each component
    var n1 = p.x * 1597 + p.y * 3571;
    var n2 = p.x * 2179 + p.y * 5231;

    // First hash component
    var state1: u32 = u32(n1) * 747796405u + 2891336453u;
    state1 = ((state1 >> ((state1 >> 28u) + 4u)) ^ state1) * 277803737u;
    state1 = (state1 >> 22u) ^ state1;

    // Second hash component with different constants
    var state2: u32 = u32(n2) * 1103515245u + 12345u;
    state2 = ((state2 >> ((state2 >> 28u) + 4u)) ^ state2) * 277803737u;
    state2 = (state2 >> 22u) ^ state2;

    // Convert to [-1, 1] range for both components
    var x = -1.0 + 2.0 * f32(state1) / f32(0xffffffffu);
    var y = -1.0 + 2.0 * f32(state2) / f32(0xffffffffu);

    return vec2f(x, y);
}

So we will use that to compute the reference point in the current grid cell:

    let cellCoords = vec2i(floor(gcoords) + 0.5);
    let localCoords = gcoords - floor(gcoords);

    // Hash the cellCoords to get the refPoint coords
    // Normalize the output to [0,1] instead of [-1,1]
    let refPointCoords = hash(cellCoords) * 0.5 + 0.5;

And just to give it a quick check we will add some debug code to display those reference points in our grid:

    // Square Radius in normalized coords space:
    let pointRadiusSqr = params.refPointSize * params.refPointSize;

    // Check the distance to the current point
    // and display the refPoint in each cell in red:
    let dir = localCoords - refPointCoords;
    let d2 = dot(dir, dir);
    let red = vec4f(1.0, 0.0, 0.0, 1.0);
    color = select(color, red, d2 < pointRadiusSqr);

Which will give us this result:

We can notice here that some points are clipped at the border of its cell,
But that's expected for this simple debug display since the reference point will be different when crossing the borders,
Ah, yes, and I also added a parameter to control the radius of the ref point for the debug display here 😉.

Finding the closest ref point

Now we need to find the closest reference point for any given pixel using the 9 closest cells as explained above.

As in the base shader toy example, we will refactor our code to have a dedicated function to find this:

struct RefPointInfos {
    distance: f32,
    gridPos: vec2f,
};

fn find_closest_ref_point(gridCoords: vec2f) -> RefPointInfos {
    // Start with a very large distance:
    var res = RefPointInfos(10.0, vec2f(0.0));

    // From those coords we can extract the integer part which
    // are the cell coords, and the fractional part which are our pixel
    // local (normalized) coords on the cell:
    let cellCoords = vec2i(floor(gridCoords) + 0.5);
    let localCoords = gridCoords - floor(gridCoords);

    for (var j = -1; j <= 1; j++) {
        for (var i = -1; i <= 1; i++) {
            // Get the coord offset to use from our current cell:
            // This "offset" is what we need to add to our current cell
            // coords to get the sibling cells coords. 
            // Note: this was called "g" in the original implementation. 
            // Note2: we use i32 values directly ourselves.
            let cellOffset = vec2i(i, j);

            // hash the sibling cell coords to get its reference point
            // local coords (renormalizing to [0,1] instead of [-1,1])
            // Note: this was called o in the original implementation
            let refPointLocalCoords = hash(cellCoords + cellOffset) * 0.5 + 0.5;

            // Now we need to get the vector from our current pixel location
            // to that reference point. We'll use our current cell origin
            // as "frame origin". So in there the sibling ref point coords
            // are: 
            let refCoords = vec2f(cellOffset) + refPointLocalCoords;

            // So now the "r" vector is:
            let r = refCoords - localCoords;

            // Compute the squared distance:
            // Note: just called "d" in the original implementation:
            let d2 = dot(r, r);

            // And compare to the current best distance:
            // Note: "distance" field actually contains squared distance
            // for now.
            if d2 < res.distance {
                // Replace the best solution:
                res.distance = d2;
                // Get the global grid coords of the ref point:
                res.gridPos = vec2f(cellCoords) + refCoords;
            }
        }
    }

    // Turn squared distance to distance:
    res.distance = sqrt(res.distance);

    return res;
}

This function is a bit long because I added some comments in the code, but the underlying logic is pretty simple in the end.

We prepare an output "res" object with a very large distance to ensure we will find a closer solution,
Then we split the grid coordinates into the cell coordinates part and the local coordinates,
Then we iterate on each sibling cell,
Getting the coordinates of the reference point for each of them,
And computing the squared distance between that reference point and our current grid coordinates.
Updating the best solution if needed,
Finally, we return the distance and location of the closest reference point found.

And then we can use this code to get a specific cell color for each pixel and thus display our voronoi diagram:

    // Find the closest ref point for our grid coords:
    let rp = find_closest_ref_point(gcoords);

    // Generate a color based on the selected refPoint grid position:
    // But avoid floating point changes here on a single cell:
    let coords = floor(rp.gridPos * 100.0);
    var rgb = hash3(coords);

    let d = rp.distance;

    // Darken a bit the color the further we get from the
    // reference point (ie. cell center):
    rgb *= clamp(1.0 - 0.4 * d * d, 0.0, 1.0);

    // Display a black dot instead if we are very close
    // to the reference point:
    rgb *= smoothstep(0.08, 0.09, d);

    let color = vec4f(rgb, 1.0);

We get the closest reference point,
Then we get a fixed rgb color per cell using the fixed local position of the reference point.
Note that in the code above i'm also using an additional helper function to generate a random color from a vec2f input:

fn hash3(p: vec2f) -> vec3f {
    // Simple hash using sin/cos with large multipliers
    // The large numbers help break up patterns
    var r = sin(p.x * 12.9898 + p.y * 78.233) * 43758.5453;
    var g = sin(p.x * 93.9898 + p.y * 67.345) * 27183.1592;
    var b = sin(p.x * 269.5 + p.y * 183.3) * 51829.6737;

    // Take fractional part and ensure [0, 1] range
    return vec3f(abs(fract(r)), abs(fract(g)), abs(fract(b)));
}

This just uses sinus so it is not as good as the PCG hash function we use for the reference points themself,
but since this is only for coloring the cells it will be good enough here.
Also, we could stop with just the computation of the rgb value to get a proper voronoi diagram display
But like in the Inigo Quilez example, we can also then apply a simple darkening radial gradient effect to get some "non flat" effect on each cell.
And we can display the cell centers as black dots using this statement rgb *= smoothstep(0.08, 0.09, d);

Adding time animation

Okay, so this is looking nice already but there is still a critical part missing which is the time animation of the diagram! So let's now add this.

First, I add access to my camera specific UniformBuffer when building the compute node in ComputeSurfaceBlueprint as this will give me a time value I need:

    auto bgrp0 = cnode->setup_bind_grp({
        // params.as_ubo(),
        stbuf.buffer().as_ubo(),
        scene.get_camera().as_ubo(),
        stoTex2dArray(tex),
    });

Then I can reflect this change on this shader part:

@group(0) @binding(0) var<uniform> params : ComputeParams;
@group(0) @binding(1) var<uniform> cam : StdCameraUBO;
@group(0) @binding(2) var outputTex: texture_storage_2d_array<rgba8unorm,write>;

And for reference, here is the content I have for the StdCameraUBO structure:

struct StdCameraUBO {
    projMatrix: mat4x4f,
    viewMatrix: mat4x4f,
    revZProjMatrix: mat4x4f,
    invViewMatrix: mat4x4f,
    invProjMatrix: mat4x4f,
    invRevZProjMatrix: mat4x4f,
    camWorldPos: vec4f,
    winSize: vec2i,
    near_plane: f32,
    far_plane: f32,

    c_log_factor: f32,
    f_log_coef: f32,
    fovy: f32,
    fovx: f32,

    time: f32,
};

=> here you can see that we will provide the time value as a f32 in the last member of the struct.

And finally we use this time parameter to compute a dynamic movedRefCoords vec2f value:

    let refCoords = vec2f(cellOffset) + refPointLocalCoords;

    // Now with time animation as in the original Inigo Quilez version:
    let movedRefCoords = vec2f(cellOffset) + (0.5 + 0.5 * sin(cam.time + 6.2831 * refPointLocalCoords));

    // So now the "r" vector is:
    let r = movedRefCoords - localCoords;

Note that here we replace the refPointLocalCoords with (0.5 + 0.5 * sin(cam.time + 6.2831 * refPointLocalCoords))
This value depends on the current time and the static refPointLocalCoords,
And the output of this will still be in the range [0,1], so our dynamic reference points stays on the grid cell as required.

And here we are! With this simple time parameter addition, we now get our voronoi diagram to animate and change continuously as shown in the first overview video!

Limitations and future work

Now, this animation mechanism is not perfect: if you pay close attention to the movements, you will notice that the cell center are always moving following the same path.

In many cases I don't think this would really be a problem, but still there would be different ways to improve on this point I think:

First we could use a more complex animation function instead of simply sin(time +offset)
Or, more interestingly, one could imagine generating 2 or more reference points per cell to get more random elements to experiment with
Or even better, one could consider 2 grids instead of just one, with a different grid size for the second one maybe and a random continuous translation of that second grid with respect to the first one ??

But anyway, this will be good enough for a first version, so let's just stop here for this time.

As usual, don't forget to check out the shader files on the NervLand Adventures repository if you want to dig deeper into this, and if you have any questions, don't hesitate to leave a comment below!

See you next time 😉! Bye!

FFmpeg video Playback in Native WGPU

The Lone Engineer — Sat, 20 Sep 2025 13:10:49 +0000

Hi there!

In this quick tutorial/overview, we will discuss how to use the updated Dawn API to import external textures into WebGPU, and more precisely how to use this API to import a video stream generated with FFmpeg in a DirectX11 context.

But first: if you prefer to follow this tutorial in a video, there is a compagnon youtube version for it available at:

References

Here are some links relevant to this article:

Microsoft DXGI Format Documentation: https://learn.microsoft.com/en-us/windows/win32/api/dxgiformat/ne-dxgiformat-dxgi_format
NervLand Adventures Repository: https://github.com/roche-emmanuel/nervland_adventures
TerrainView8 reference app: https://nervtech.org/terrainview8/

Introduction

To implement this feature I built a dedicated VideoPlayer component in my NervLand engine, and here is an example of the usage of this component displaying a recording of a gameplay session from "Ghost of Tsushima" inside my TerrainView test application:

Nothing too fancy here: we just start playing the video file on the start of the application, and that's it: no way to pause the playback, rewind, resize the window or anything, but we have to start somewhere, right 😉?

Now, before we go any further into the implementation details, please note that I have also prepared a dedicated folder in the NervLand Adventures repository to share the files relevant for this video import feature:

So if you are interested in this, you could have a look at this repository at this location: https://github.com/roche-emmanuel/nervland_adventures/

Important note: Those source files will not compile out of the box, but this might still be useful as a reference if you want to implement something similar on your side.

Quick Project overview:

Here is a simplified overview of the design:

We start with a regular video file, which in our case is some gameplay recording in mp4 format with x264 encoding.

Then we load that file with the ffmpeg library, configuring the decoder to use hardware acceleration and produce the video image stream directly on the GPU in a DirectX context.

Then we will copy the resulting DirectX Texture into our WebGPU context using the new Shared Texture handling API from Dawn.

And finally we can process the final WGPU Texture as any other texture and display it in our application.

Current limitations:

Now, this is all just a very simple and minimal implementation with a lot of limitations:

In the NervLand code, I have added a VideoPlayer class, which itself contains an abstract VideoDecoder class.

Next I added a concrete implementation called FFMPEGVideoDecoder, which depends on the ffmpeg binaries, and as such, I'm currently only building this module in the native version of my engine. Which means there is currently no support for this video playback feature when I build with Emscripten.

And on top of that, the current implementation I have is currently only provided for Windows, and explicitly requires that we use the DirectX 12 backend on the Dawn side, and the DirectX 11 Video Acceleration format on the ffmpeg side.

So there is still quite a lot to test and investigate on this topic actually 😅.

=> But anyway, now let' get started with some analysis of the code used to render this video clip!

Code Analysis

Setting up the "Video Surface":

In my TerrainView app I'm now starting to progressively implement and use a configuration system to create some additional scene content on the planet.

So, this is not directly part of the Video playback layer, but still, for reference, here it the YAML element I'm using on my side to request the creation of a simple quad at a specific location on the ground, and on which we will render the video stream:

  video_surface:
    type: VideoSurfaceBlueprint
    node: quads0
    reference_point: phantom_rp01
    position: Vec3d(0.0, 2.0, 0.0)
    ypr: Vec3d(0.0, 0.0, 0.0)
    anchor: bottom | center
    video_file: D:\Temp\Videos\Ghost_of_Tsushima_v1.mp4
    width: 3.0
    aspect: 16/9

Next, we will use this VideoSurfaceBlueprint to construct an actual "video surface object" in this method:

void VideoSurfaceBlueprint::construct_blueprint(Scene& scene) {
    // Retrieve the node to which this quad should be attached:
    auto& node = scene.get_node_by_id<QuadsNode>(_node);

    // Get the world position taking the position offset into account:
    auto wpos = calculate_world_position();

    auto& bsp = node.get_texture();
    BSPArea area = create_video_or_fallback_area(bsp);

    // Setup the Quad location:
    setup_quad_location(bsp, area);

    // Add the quad at that position:
    node.add_quad(wpos, _quadLoc, _persistent);
}

In this method, the key step for us is the call to create_video_or_fallback_area(), which, in case the video file is found will lead to the call of this dedicated method:

auto VideoSurfaceBlueprint::create_video_area(WGPUBSPTexture& bsp) -> BSPArea {
    logDEBUG("Loading video file: {}", _videoFile);
    auto player = VideoPlayer::create({.videoFile = _videoFile});
    auto width = player->get_width();
    auto height = player->get_height();
    logDEBUG("Creating target BSP area of size {}x{}", width, height);

    // We should keep the player as a referenced object in the engine:
    auto* eng = WGPUEngine::instance();
    eng->set_shared_object(get_id(), player.get());

    // auto img =
    //     Image::make_checkerboard(16, 32, 16, RGBA8_WHITE,
    //     RGBA8_DARKYELLOW);
    auto img = Image::make_random<RGBA8>(width, height);
    auto area = bsp.add_image(img);
    // area = bsp.add_area(width, height);

    // Assign the target texture and the origin point:
    Vec3u orig(area.rect.xmin, area.rect.ymin, area.layer);
    player->set_target_texture(bsp.get_texture(), orig);

    // Start the video playback:
    player->play();

    return area;
}

As shown above, we will:

First, create a VideoPlayer: auto player = VideoPlayer::create({.videoFile = _videoFile});
Then we need to register the player in the engine to keep a valid reference on it as long as we need it.
Then assign a render location to that player object, which is a fixed location on a given texture that we will later use a source to display texture data on the quad we created for this VideoSurface: player->set_target_texture(bsp.get_texture(), orig);
And we immediately request the player to start the playback with player->play();.

When we create the VideoPlayer object, first we create the dedicated FFMPEGVideoDecoder, and then we request the opening of the video file if it is already specified:

VideoPlayer::VideoPlayer(const VideoPlayerDesc& desc) : _desc(desc) {
    logDEBUG("VideoPlayer initialized.");

    VideoDecoderDesc ddesc{};

#if NV_USE_FFMPEG
    _decoder = nv::create<FFMPEGVideoDecoder>(ddesc);
#endif

    // Check the decoder is valid:
    NVCHK(_decoder != nullptr, "VideoDecoder was not assigned.");

    if (!desc.videoFile.empty()) {
        open_file(desc.videoFile.c_str());
    }
};

The constructor for FFMPEGVideoDecoder is really not doing much: just initializing ffmpeg network support [//which we probably don't really need here anyway//]:

FFMPEGVideoDecoder::FFMPEGVideoDecoder(const VideoDecoderDesc& desc)
    : VideoDecoder(desc) {
    // Initialize FFmpeg (call once globally, but safe to call multiple times)
    avformat_network_init();

    logDEBUG("FFMPEGVideoDecoder initialized");
};

The actual work happens a bit later in the VideoPlayer::open_file() method when we open the input on the decoder itself, and after that when we call the play() function:

auto VideoPlayer::open_file(const char* filename) -> bool {
    if (_decoder == nullptr) {
        logERROR("VideoPlayer: No decoder, cannot open file.");
        return false;
    }

    if (!system_file_exists(filename)) {
        logERROR("VideoPlayer: video file '{}' doesn't exists.", filename);
        return false;
    }

    logDEBUG("Opening video file {}...", filename);
    if (!_decoder->open_input(filename)) {
        logERROR("VideoPlayer: decoder cannot open input.");
        return false;
    }

    _filename = filename;

    if (_desc.playOnOpen) {
        play();
    }

    return true;
};

FFMpeg decoder internals

In the decoder open_input() method we will call the initialize_decoder() method:

auto FFMPEGVideoDecoder::initialize_decoder(const char* filename) -> bool {

    // Open input file
    _formatCtx = avformat_alloc_context();
    if (!_formatCtx) {
        logERROR("Failed to allocate format context");
        return false;
    }

    I32 ret = avformat_open_input(&_formatCtx, filename, nullptr, nullptr);
    if (ret < 0) {
        logDEBUG("Failed to open input file: {}", err2str(ret));
        return false;
    }

    // Retrieve stream information
    ret = avformat_find_stream_info(_formatCtx, nullptr);
    if (ret < 0) {
        logDEBUG("Failed to find stream info: {}", err2str(ret));
        return false;
    }

    // Find video stream
    _videoStreamIdx =
        av_find_best_stream(_formatCtx, AVMEDIA_TYPE_VIDEO, -1, -1, nullptr, 0);
    if (_videoStreamIdx < 0) {
        logDEBUG("No video stream found");
        return false;
    }

    AVStream* video_stream = _formatCtx->streams[_videoStreamIdx];
    const AVCodec* codec =
        avcodec_find_decoder(video_stream->codecpar->codec_id);
    if (codec == nullptr) {
        logDEBUG("Unsupported codec");
        return false;
    }

    // Calculate FPS
    if (video_stream->r_frame_rate.den != 0) {
        _fps = av_q2d(video_stream->r_frame_rate);
    } else if (video_stream->avg_frame_rate.den != 0) {
        _fps = av_q2d(video_stream->avg_frame_rate);
    } else {
        _fps = 25.0; // Default fallback
    }

    // Store frame dimensions
    _frameWidth = video_stream->codecpar->width;
    _frameHeight = video_stream->codecpar->height;

    // Try to setup hardware decoder first
    _isHWAccelerated = false;
    if (_desc.enableHardwareAcceleration && setup_hardware_decoder(codec)) {
        logDEBUG("Hardware acceleration enabled");
        _isHWAccelerated = true;
    } else {
        logDEBUG("Using software decoding");
        // Fallback to software decoder
        _codecCtx = avcodec_alloc_context3(codec);
        if (_codecCtx == nullptr) {
            logDEBUG("Failed to allocate codec context");
            return false;
        }
    }

    // Copy codec parameters
    ret = avcodec_parameters_to_context(_codecCtx, video_stream->codecpar);
    if (ret < 0) {
        logDEBUG("Failed to copy codec parameters: {}", err2str(ret));
        return false;
    }

    // Open codec
    ret = avcodec_open2(_codecCtx, codec, nullptr);
    if (ret < 0) {
        logDEBUG("Failed to open codec: {}", err2str(ret));
        return false;
    }

    return true;
}

This method is pretty large so I'm not going to cover it completely here and you should refer directly to the provided base code if you need to go over this step by step.

Yet, I'd say that the key part for us here is when we do the setup for the hardware acceleration support, which is maybe something unusual or new for some of you (?):

    if (_desc.enableHardwareAcceleration && setup_hardware_decoder(codec)) {
        logDEBUG("Hardware acceleration enabled");
        _isHWAccelerated = true;
    }

This method is where we will need to start setting up the DirectX context for FFmpeg:

auto FFMPEGVideoDecoder::setup_hardware_decoder(const AVCodec* codec) -> bool {
#ifdef DAWN_ENABLE_BACKEND_D3D12
#if NV_FFMPEG_DX_VERSION == 11
    AVHWDeviceType hw_type = AV_HWDEVICE_TYPE_D3D11VA;
    AVPixelFormat px_fmt = AV_PIX_FMT_D3D11;
#else
    AVHWDeviceType hw_type = AV_HWDEVICE_TYPE_D3D12VA;
    AVPixelFormat px_fmt = AV_PIX_FMT_D3D12;
#endif

    AVBufferRef* hwDeviceCtx = av_hwdevice_ctx_alloc(hw_type);
    NVCHK(hwDeviceCtx != nullptr, "Cannot allocate HW Device context for {}.",
          av_hwdevice_get_type_name(hw_type));

    auto* hw_device_ctx = (AVHWDeviceContext*)hwDeviceCtx->data;

// Set the custom device here:
#if NV_FFMPEG_DX_VERSION == 11
    auto* d3d_ctx = (AVD3D11VADeviceContext*)hw_device_ctx->hwctx;
    d3d_ctx->device = DX11Engine::instance().device();
#else
    auto* d3d_ctx = (AVD3D12VADeviceContext*)hw_device_ctx->hwctx;
    d3d_ctx->device = DX12Engine::instance().device();
#endif
    NVCHK(d3d_ctx->device != nullptr,
          "Invalid D3D Device for FFMPEG hw context.");
    d3d_ctx->device->AddRef();
#else
#error "No implementation for FFMPEGVideoDecoder::setup_hardware_decoder yet."
#endif

    // Initialize the context
    int ret = av_hwdevice_ctx_init(hwDeviceCtx);
    NVCHK(ret >= 0, "FFMPEGVideoDecoder: Cannot initialize HW Device context.");

    _hwDeviceCtx = hwDeviceCtx;

    _codecCtx = avcodec_alloc_context3(codec);
    if (_codecCtx != nullptr) {
        _codecCtx->hw_device_ctx = av_buffer_ref(_hwDeviceCtx);
        _hwPixelFormat = px_fmt;
        logDEBUG("Hardware decoder setup successful: {}",
                 av_hwdevice_get_type_name(hw_type));
        return true;
    }

    logDEBUG("No suitable hardware acceleration found.");
    return false;
}

=> I'm sorry this method is also a bit messy because I was experimenting also with a DirectX12 context on the FFmpeg side (but I didn't manage to get this working properly yet when then trying to retrieve the texture in the WGPU context [//which is essentially another DX12 device in fact//])

Anyway:

Here we start with allocating an ffmpeg buffer for the hardware device of the proper type (AV_HWDEVICE_TYPE_D3D11VA for us) with this line: AVBufferRef* hwDeviceCtx = av_hwdevice_ctx_alloc(hw_type);
Next we assign the DX11 device to use on the ffmpeg side into this structure: d3d_ctx->device = DX11Engine::instance().device();
And finally we initialize this hardware device struct: int ret = av_hwdevice_ctx_init(hwDeviceCtx);

And thus, we end with a codec context with hardware acceleration that is now ready to start decoding video frames:

    _codecCtx = avcodec_alloc_context3(codec);
    if (_codecCtx != nullptr) {
        _codecCtx->hw_device_ctx = av_buffer_ref(_hwDeviceCtx);
        _hwPixelFormat = px_fmt;
        logDEBUG("Hardware decoder setup successful: {}",
                 av_hwdevice_get_type_name(hw_type));
        return true;
    }

VideoPlayer play loop

With the ffmpeg decoder ready the processing continues with the call to VideoPlayer::play():

void VideoPlayer::play() {
    NVCHK(_decoder != nullptr, "Invalid decoder.");

    _isPlaying = true;
    _playbackStartTick = SystemTime::tick();
    _lastUpdateTick = -1;
    _playTime = 0.0;
    _currentFrameIndex = 0;
    logDEBUG("Started playing video {}", _filename);

    auto* eng = WGPUEngine::instance();
    _updateCb = eng->add_pre_render_func([this] { update(); });
};

The main idea in the function is simply to register a callback into the engine that will be called on each frame:

void VideoPlayer::update() {
    if (!_isPlaying)
        return;

    auto curTick = SystemTime::tick();
    if (_lastUpdateTick == -1) {
        _lastUpdateTick = curTick;
    }
    auto elapsed = SystemTime::delta_s(_lastUpdateTick, curTick) * _videoSpeed;
    _lastUpdateTick = curTick;
    _playTime += elapsed;

    F64 videoFps = _decoder->get_fps();
    I32 expectedFrameIndex = static_cast<I32>(_playTime * videoFps);

    // Only decode if we need to advance to the next frame:
    if (expectedFrameIndex > _currentFrameIndex) {
        I32 delta = expectedFrameIndex - _currentFrameIndex;
        if (delta > 1) {
            logWARN("Jumping over {} video frames.", delta - 1);
        }

        if (_decoder->decode_next_frame()) {
            _currentFrameIndex = expectedFrameIndex;
            // logDEBUG("Current video frame: {}", _currentFrameIndex);
            // get the texture data:
            _decoder->get_current_frame(_texture, _origin);
        } else {
            // End of video reached
            logDEBUG("No additional frame, stopping playback.");
            stop();
        }
    }
};

In this update method we simply use the current execution time and the video FPS rate to figure out if we should request the next frame to the decoder with _decoder->decode_next_frame().

And if the decoding is successfully, we immediately request the "copy" of that frame into our dedicated WGPU texture (into a subregion of that texture at least).

Getting the current frame from DirectX

The decode_next_frame() function is not that interesting to us: it will simply process a "stream of packets" until a new frame is completed, so let's move directly to get_current_frame() as this step requires some additional work.

For reference, the first part of this method is this:

    if (_isHWAccelerated && hw_frame->format == AV_PIX_FMT_D3D11) {
        auto* d3d_texture = (ID3D11Texture2D*)hw_frame->data[0];
        I64 idx = (I64)(intptr_t)hw_frame->data[1];
        // logDEBUG("DX11 src tex: {}, layer: {}", (const void*)d3d_texture,
        // idx);
        if (_textureInterface == nullptr) {
            init_dx11_texture_interface(d3d_texture);
        }
        convert_nv12_to_rgba(d3d_texture, idx);
    }

And in fact, this is where things start to get a bit tricky as we have a few issues to handle:

1. First, when performing the frame decoding, ffmpeg will produce textures with the format DXGI_FORMAT_NV12, which doesn't have much to do with RGBA unfortunately.
2. Then the generated output texture will not be usable as shader resource directly because it has the D3D11_BIND_DECODER flag which means the texture is optimized for video decoder and this is not compatible with D3D11_BIND_SHADER_RESOURCE
3. And of course, this texture doesn't have any shared flag (by default at least) so it's not possible to share it directly with another DirectX context.
4. Ohh, and in fact this texture is actually a Texture2DArray with the DX11 backend here and the layer to use will change depending on the current frame being generated.

Fixing the BIND_DECODER issue:

To resolve the "BIND_DECODER" issue we need to copy the layer of interest from the input texture into an intermediate texture (still with the NV12 format) before we can perform any processing,

So we create that texture once:

    auto& dx11 = DX11Engine::instance();
    auto* device = dx11.device();

    D3D11_TEXTURE2D_DESC desc;
    srcTex->GetDesc(&desc);

    // Create the intermediate nv12 texture:
    _nv12Texture = dx11.createTexture2D(
        desc.Width, desc.Height, D3D11_BIND_SHADER_RESOURCE, desc.Format);

Then on each copy cycle we will get the source layer of interest copied in this nv12Texture (which by the way also handles the point 4 at the same time):

    auto& dx11 = DX11Engine::instance();
    auto* context = dx11.context();

    // Copy to our intermediate texture since the decoder output cannot be used
    // as shader resource. Copy specific array slice to the single-layer texture
    UINT srcSubresource = D3D11CalcSubresource(0,        // mip level
                                               layerIdx, // array slice
                                               1);       // mip levels
    UINT dstSubresource = 0;                             // Single layer, mip 0

    context->CopySubresourceRegion(_nv12Texture.Get(), dstSubresource, 0, 0,
                                   0, // Dest x, y, z
                                   srcTex, srcSubresource,
                                   nullptr // Copy entire subresource
    );

Converting NV12 to RGBA:

Next we need to convert from the NV12 format to RGBA, so we'll do that on the DirectX 11 side and we'll need another destination texture to store our conversion result which we called rgbaTexture, and while we are at it, we can also build this as a shared texture to be able to import it later in the WGPU context:

    // Get the desc of the source texture:
    D3D11_TEXTURE2D_DESC tdesc;
    srcTex->GetDesc(&tdesc);

    // We expect the format to be NV12:
    NVCHK(tdesc.Format == DXGI_FORMAT_NV12, "Unexpected source tex format: {}",
          tdesc.Format);

    // Now create a corresponding RGBA Teture that is shared:
    logDEBUG("FFMPEGVideoDecoder: Creating target RGBA Texture.");

    auto& dx11 = DX11Engine::instance();
    HANDLE sharedHandle{nullptr};
    _rgbaTexture = dx11.createReadOnlySharedTexture2D(
        &sharedHandle, tdesc.Width, tdesc.Height,
        D3D11_BIND_RENDER_TARGET | D3D11_BIND_SHADER_RESOURCE |
            D3D11_BIND_UNORDERED_ACCESS);

Next we have to execute the compute shader that will convert the nv12 texture into rgba on each cycle:

    // Get output texture dimensions
    D3D11_TEXTURE2D_DESC desc;
    _rgbaTexture->GetDesc(&desc);

    // Set compute shader
    context->CSSetShader(_nv12ToRgbaProgram.computeShader, nullptr, 0);

    // Bind input texture arrays
    ID3D11ShaderResourceView* srvs[] = {_luminanceSRV.Get(), _chromaSRV.Get()};
    context->CSSetShaderResources(0, 2, srvs);

    // Bind output texture
    ID3D11UnorderedAccessView* uavs[] = {_rgbaUAV.Get()};
    context->CSSetUnorderedAccessViews(0, 1, uavs, nullptr);

    // Dispatch compute shader
    UINT groupsX = (desc.Width + 7) / 8;
    UINT groupsY = (desc.Height + 7) / 8;
    context->Dispatch(groupsX, groupsY, 1);

Here we will be rendering on the RGBA Texture with "UnorderedAccess" (ie. the _rgbaUAV resource)

But another interesting thing to notice here is on the shader resources we provide as input: _luminanceSRV and _chromaSRV:

We have a single nv12Texture as input, but given this specific format it will be possible to create those 2 separated SRVs from it.

Here is how we create the luminanceSRV:

    D3D11_SHADER_RESOURCE_VIEW_DESC luminancePlaneDesc{};
    luminancePlaneDesc.Format = DXGI_FORMAT_R8_UNORM;
    luminancePlaneDesc.ViewDimension = D3D11_SRV_DIMENSION_TEXTURE2D;
    luminancePlaneDesc.Texture2D.MipLevels = 1;
    luminancePlaneDesc.Texture2D.MostDetailedMip = 0;

    HRESULT hr = device->CreateShaderResourceView(
        _nv12Texture.Get(), &luminancePlaneDesc, _luminanceSRV.GetAddressOf());
    NVCHK(SUCCEEDED(hr), "Failed to create luminance SRV");

Note here that we use the R8_UNORM format for the view.

And then here is the code used to create the chromaSRV:

    D3D11_SHADER_RESOURCE_VIEW_DESC chromaPlaneDesc{};
    chromaPlaneDesc.Format = DXGI_FORMAT_R8G8_UNORM;
    chromaPlaneDesc.ViewDimension = D3D11_SRV_DIMENSION_TEXTURE2D;
    chromaPlaneDesc.Texture2D.MipLevels = 1;
    chromaPlaneDesc.Texture2D.MostDetailedMip = 0;

    hr = device->CreateShaderResourceView(_nv12Texture.Get(), &chromaPlaneDesc,
                                          _chromaSRV.GetAddressOf());
    NVCHK(SUCCEEDED(hr), "Failed to create chroma SRV");

Note that the only difference here is that we request the R8G8_UNORM format for the view.

And internally this all what is required for the DX11 driver to create the separated luminance/chroma views on the input NV12 texture.

For additional infos on the NV12 format you could check the microsoft learn website which provides some explanations on the logic above:

https://learn.microsoft.com/en-us/windows/win32/api/dxgiformat/ne-dxgiformat-dxgi_format

NV12 conversion shader

Now let's have a look at the NV12 to RGBA conversion shader itself:

First, this shader declares the input/output resources that we will use:

// Input texture arrays - Y plane and UV plane separated
Texture2D<float> lumTexture : register(t0); // Y plane array (single channel)
Texture2D<float2> chromaTexture : register(t1); // UV plane array (two channels)

// Output RGBA texture
RWTexture2D<float4> rgbaTexture : register(u0);

=> So we have the luminance & chroma SRVs, and the writable rgba output texture.

Next we have the core of the logic used to convert NV12 to RGB:

// YUV to RGB conversion matrix (BT.709)
// From Microsoft documentation for 8-bit YUV to RGB888
static const float3x3 YUVtoRGBCoeffMatrix = {1.164383f, 1.164383f,  1.164383f,
                                             0.000000f, -0.391762f, 2.017232f,
                                             1.596027f, -0.812968f, 0.000000f};

float3 ConvertYUVtoRGB(float3 yuv) {
  // Subtract the offset values (16/255 for Y, 128/255 for UV)
  yuv -= float3(0.062745f, 0.501960f, 0.501960f);
  yuv = mul(yuv, YUVtoRGBCoeffMatrix);
  return saturate(yuv);
}

When dealing with NV12 we essentially manipulate YUV triplets, it's just the Y value is provided for each pixel, while the UV values are provided less often/shared between groups of 4 pixels.

But once we have the YUV value for a given pixel, we can easily convert it to RGB with this function.

So then in the main function of our compute shader, we simply collect the proper Y and UV values for our current pixel (given by the id.xy coords):

  // Sample Y component at full resolution from the specified layer
  float Y = lumTexture.Load(uint3(id.xy, 0));

  // Sample UV components at half resolution from the specified layer
  uint2 uvCoord = uint2(id.x / 2, id.y / 2);
  float2 UV = chromaTexture.Load(uint3(uvCoord, 0));

=> Here we only need to be carefull to devide the id.xy coords by 2 to get the UV part.

And finally, we convert this YUV triplet to RGB, and we write the data at the proper texture pixel location:

  // Convert YUV to RGB
  float3 rgb = ConvertYUVtoRGB(float3(Y, UV.x, UV.y));

  // Output RGBA (with full alpha)
  // Note: we flip the image vertically here to match our nervland convention:
  // rgbaTexture[id.xy] = float4(rgb, 1.0f);
  rgbaTexture[uint2(id.x, outputSize.y - id.y - 1)] = float4(rgb, 1.0f);

Copying RGBA texture from DX11 to WGPU:

After the compute shader is executed we have the content we want in RGBA8 format in our DX11 shared rgbaTexture

So, in the second part of the frame update logic we will perform this copy operation with this code:

    SharedTextureMemoryBeginAccessDescriptor beginDesc{};
    beginDesc.initialized = true;
    beginDesc.concurrentRead = false;
    beginDesc.fenceCount = 0;
    beginDesc.fences = nullptr;
    beginDesc.signaledValues = nullptr;

    if (!_sharedTexMem.BeginAccess(_textureInterface, &beginDesc)) {
        logERROR("Cannot begin access to shared texture.");
    }

    // eng->copy_texture(_textureInterface, texture, nullptr, &origin);
    _copyPass->execute();

    SharedTextureMemoryEndAccessState endDesc{};

    if (!_sharedTexMem.EndAccess(_textureInterface, &endDesc)) {
        logERROR("Cannot end access to shared texture.");
    }

So, this is probably starting to feel like unusual code even if you are already used to the Dawn library.

And to be honest, I'm not completely sure this is how it is intended to be used (as it is still pretty hard to find detailed information on this on the net for now), but at least this seems to we working for me 😉.

The key point here is that we want to copy the DX11 rgbaTexture to another container texture in WGPU (which is the location we provided initially if you remember that part):

    // Assign the target texture and the origin point:
    Vec3u orig(area.rect.xmin, area.rect.ymin, area.layer);
    player->set_target_texture(bsp.get_texture(), orig);

    // Start the video playback:
    player->play();

To access the DX11 texture from inside the WGPU context we use that _textureInterface object.
But whenever we try to access that object on the WGPU side, we need to encapsulate that access with a pair of BeginAccess and EndAccess calls on the SharedTextureMemory object used to create that texture interface (I'm not quite sure why as I'm not using keyedMutex when created the DX11 shared texture, but it didn't work for me otherwise).
And last thing to note on this code chunk is that to perform the actual copy operation I use a dedicated compute shader in its compute pass, but I don't think I need to cover this here as this is really a regular and simple WGSL shader copying the source RGBA8 pixels to a specific subregion of a writable storage texture.

Creating Shared texture interface:

But still, now let's have a look at the creation of the texture interface:

    // Open the DX11 texture in Dawn from the shared handle and return it as a
    // WebGPU texture.
    SharedTextureMemoryDXGISharedHandleDescriptor sharedHandleDesc{};
    sharedHandleDesc.handle = sharedHandle;

    SharedTextureMemoryDescriptor desc;
    desc.nextInChain = &sharedHandleDesc;

    auto* eng = WGPUEngine::instance();
    _sharedTexMem = eng->import_shared_texture_memory(&desc);
    // Handle is no longer needed once resources are created.
    ::CloseHandle(sharedHandle);

The first step is to create the SharedTextureMemory object using the DX11 texture sharedHandle
Note that the import_shared_texture_memory method here is a direct call to the method with the same name on the WGPU device.

    TextureDescriptor texDesc{};
    texDesc.usage = TextureUsage::TextureBinding | TextureUsage::CopyDst;
    texDesc.dimension = TextureDimension::e2D;
    texDesc.size = {tdesc.Width, tdesc.Height, 1};
    // Convert DXGI format to WebGPU format as needed
    texDesc.format = convert_dxgi_to_wgpu_format(DXGI_FORMAT_R8G8B8A8_UNORM);

    _textureInterface = _sharedTexMem.CreateTexture(&texDesc);

Once we have the sharedTexMemory object we can create the texture interface from that memory region, which will give a plain old regular wgpu::Texture object representing our DX11 Texture inside WGPU 😎!

And this is it guys 🥳! With all those elements put together we can then extract a stream of video frame from the source video file and get them fully imported into WebGPU with complete hardware acceleration 👍!

Recap and conclusion:

So now, just to recap, here is an overview of the workflow we used to copy those frames:

First we started with a texture 2D array provided by FFMpeg and containing the just decoded frame we asked.
We had to make a first copy of this, selecting only the layer of interest in the process, to produce a texture we could then use in our regular DX11 render or compute pipelines.
Next we use a compute shader to convert from NV12 to RGBA format, writing to a shared DX11 texture as output.
Next we generate an "interface" for that shared DX11 texture into the WGPU context.
Finally we have a simple copy compute shader in WGPU to copy the texture data into another texture which I'm using as a texture atlas in my case.

And here we are for this ffmpeg/webgpu integration tutorial! I hope you learnt something here!

If you'd like to explore this topic further, remember to check out the resources I shared in the NervLand Adventures GitHub repository as I mentioned at the beginning of the video. And if you have any questions, don't hesitate to leave a comment below!

See you next time 🤟👋!

TerrainView7: Full scale planet rendering in WebGPU + Emscripten, now with Precomputed Atmospheric Scattering support

The Lone Engineer — Sun, 13 Apr 2025 12:31:03 +0000

Hi everyone,

Up to this point, I was rather posting short articles on reddit, but for some reason, it seems my account there is just not working anymore (I'm getting a "server error" everytime I try to check my own profile and I can't post/comment anything anymore...), so now, time to give this dev.to network a try 😄!

Anyway, I've been working for a while now on a custom 3D engine I'm building on my own with webGPU, with support for both native compilation (as I'm using the google Dawn library) and WASM compilation with emscripten. I provided a few different tech demo releases already, but today, I just wanted to share the work I did on the latest version, where I'm introducing support for the Atmosphere rendering on my terrain engine.

If your browser supports WebGPU, you can give this demo a try at: https://nervtech.org/terrainview7
And if you want a quick overview on the new features/changes I introduced in this release, you can check the companion video:

Additionally, I also lately recording a simple flythough recording on the terrain, as I think this is starting to look interesting now, and I simply enjoy navigating around ;-):

The atmosphere rendering was implemented from scratches in WebGPU/WGSL, but it is completed based on the implementation provided by Eric Bruneton paper on "Precomputed Atmospheric Scattering" (cf. https://ebruneton.github.io/precomputed_atmospheric_scattering/ ) (plus some optimization in the structure of the code since I can use compute shaders on my side instead of geometry+fragment shaders in the original code).

For the terrain rendering part itself I also use an implementation based on the Proland terrain rendering framework ( https://proland.inria.fr/ ), but with more significant changes (replaced completely the scheduler/tasks concepts, producing gpu tiles in batches with compute shaders, large refactoring of elevation/ortho producers, using texture 2D array atlas for storage, etc...)

If you have any question on this project, please let me know of course ;-)!

Thanks!