Hardware-Decoded Video Textures in Rust: NVIDIA NVDEC, CUDA, and wgpu-hal

Article

A cube rendered with wgpu, with Big Buck Bunny playing on its faces: the bunny stands in a meadow under a blue sky on the front face, and the same frame shows on the top and left faces. — The finished demo: NVDEC decodes the video, CUDA converts each frame into a shared D3D12 texture, and wgpu samples it on a spinning cube. — Big Buck Bunny © Blender Foundation, CC BY 3.0

I wanted a video playing on a 3D object in a Rust program. That sounds like a solved problem, and in one sense it is: decode a frame with any library, copy its pixels into memory, upload them to the GPU as a texture, repeat 24 times a second. It works. But every frame makes a round trip the GPU didn't need: the decoder (often the GPU's own video hardware) produces pixels in GPU memory, the program copies them down to system memory, and then uploads them straight back up so the renderer can use them.

This post is about building the other version, where that round trip doesn't exist. NVIDIA's hardware decoder writes each frame into GPU memory, a CUDA kernel converts it into an RGBA texture that lives in Direct3D 12, and wgpu, Rust's graphics API, samples that same texture. The decoded pixels never leave the GPU.

The rotating cube above is where it ends up: a Windows program, about 1,100 lines of Rust, C++ and CUDA, playing an MP4 on all six faces. The cube itself is ordinary wgpu, and you could replace it with anything. The work, and this post, is about getting the frames there. I'll go through it in the order I built it, explaining at each step what I needed, why I did it this way, and what the code is doing.

The full project is on GitHub at Lrios4403/nvdec-wgpu-cube. Every code block below is a verbatim excerpt of it, and each block's file name links to those exact lines.

The short version

Media Foundation reads the MP4, NVDEC decodes it. Windows' built-in Source Reader hands over the compressed H.264 or HEVC packets as they are in the file, and NVIDIA's Video Codec SDK decodes them on the GPU's dedicated video engine. No FFmpeg.
CUDA converts NV12 to RGBA in place. The decoder outputs a YUV layout called NV12, which a shader can't sample directly. A small CUDA kernel converts it, writing straight into a D3D12 texture that CUDA has imported as external memory.
wgpu-hal wraps that texture. Two unsafe calls, texture_from_raw and create_texture_from_hal, turn the existing ID3D12Resource into an ordinary wgpu::Texture that you can bind and sample in WGSL like any other.
It's a working example, not a player. Synchronization is a blocking wait each frame, there's no audio, and it's Windows, D3D12 and NVIDIA only.

Part 1: The plan, and why it takes five technologies

Before any code, here's the shape of the thing. The problem splits into stages, and each stage turned out to need a different piece of technology, because no single API does the whole job. That's the main reason this program looks complicated: it isn't that any one part is hard, it's that five parts have to agree with each other.

From MP4 file to cube

Only compressed bytes cross host memory.

Show as text

Media Foundation reads the MP4 and hands over compressed H.264 or HEVC packets with their timestamps. Nothing is decoded yet.
NVIDIA's parser, running on the CPU, splits the packets into pictures and calls back when each one can be decoded and when it's due for display.
NVDEC, the GPU's video decoding engine, decodes each picture into NV12 in GPU memory.
A CUDA kernel converts the NV12 frame to RGBA, writing straight into one of three D3D12 textures.
wgpu wraps those same three textures through wgpu-hal, and the WGSL shader samples the current one onto the cube.
Only compressed bytes pass through host memory. Decoded frames stay in GPU memory from NVDEC to the screen.

Reading the diagram top to bottom, this is what has to happen to show one frame:

Get the compressed video out of the file. An MP4 is a container, not a video. Something has to find the video track inside it and hand over its packets. I use Windows Media Foundation for this.
Turn packets into pictures. That's the decoder's job. I use NVDEC, the video decoding engine on NVIDIA GPUs, through NVIDIA's Video Codec SDK. It leaves each decoded frame in GPU memory.
Convert the frame into something a shader can use. Decoders output YUV, not RGB. A CUDA kernel does the conversion, because the decoded frame is already visible to CUDA.
Put the result where wgpu can see it. This is the step that makes the whole thing work: the kernel writes into a Direct3D 12 texture that both CUDA and wgpu can use, so there's one allocation and no copy.
Draw it. wgpu samples the texture in a WGSL shader. From here on it's ordinary graphics code.

Notice what this is and isn't. It's a no-readback pipeline: the GPU still does real work between decoder and screen (NVDEC writes NV12, the kernel writes RGBA, the renderer samples it), and the demuxing and parsing still run on the CPU. What it avoids is the trip through system memory, and with it the CPU copies and the PCIe transfers that would otherwise happen every frame.

The rest of Part 1 explains each stage in a bit more depth, so the code in Parts 2 to 4 has somewhere to hang.

MP4 is the box, H.264 is what's in it

An MP4 file is a container: it holds one or more tracks (video, audio, subtitles) plus the timing and the settings needed to play them. The video track is compressed with a codec, here H.264 or HEVC. A compressed frame isn't a picture. It's a list of instructions (predictions from other frames, motion vectors, transformed residuals) that a decoder turns back into pixels.

So playing an MP4 takes two separate steps: demuxing (pulling the compressed video samples out of the container, with their timestamps) and decoding (turning those samples into images). NVDEC only does the second. It has never heard of MP4. That's why the first thing I need is a demuxer, and why Media Foundation is in the picture at all: it's already on every Windows machine, it reads MP4, and it can be told to hand over the packets without decoding them.

What NVDEC does, and what it needs from me

NVDEC is the fixed-function video decoder on NVIDIA GPUs. NVIDIA's NVDEC programming guide puts it plainly: “NVDEC runs completely independently of compute/graphics engine.” That's the appeal. Decoding doesn't take CUDA cores or shader time away from the renderer; it happens on separate silicon.

The Video Codec SDK exposes it through an API called NVCUVID. The thing that confused me at first is that it isn't one call that takes packets and gives back pictures. It's two objects. A parser runs on the CPU and reads the compressed bytes to find where each picture starts and ends. A decoder runs on NVDEC and decodes one picture at a time. The parser drives the decoder through callbacks: you feed the parser, and it calls you back when it has something.

The NVDEC calls this program uses
Call	What it does
`cuvidCreateVideoParser`	Creates the parser for one codec, with your three callbacks.
`cuvidParseVideoData`	Feeds the parser one packet of compressed bytes. The callbacks fire from inside this call.
`cuvidDecodePicture`	Called from the decode callback: kicks off decoding of one picture on NVDEC.
`cuvidMapVideoFrame`	Post-processes a decoded picture into an output surface and returns a CUDA pointer to it. Must be unmapped again.

The three callbacks are the key to the whole decode section later, so here's what each is for:

The sequence callback fires when the parser learns the stream's format: codec, size, chroma format, bit depth. That's the first moment you know enough to create the decoder, so that's where it gets created.
The decode callback fires when a complete picture is ready to decode. You call cuvidDecodePicture and NVDEC goes to work.
The display callback fires when a picture is ready to be shown, in display order. That matters because codecs decode B-frames out of order: a frame that depends on a later frame is decoded after it but shown before it. The parser sorts that out and the display callback is where you get the sorted result.

NV12, and why CUDA converts it

Video is stored as luma (brightness) and chroma (color difference), not RGB, because the eye is less sensitive to color detail than to brightness, so codecs store color at lower resolution. NVDEC's output here is NV12: a full-resolution plane of 8-bit Y values, followed by a half-resolution plane of interleaved U and V pairs. Every 2×2 block of pixels shares one U and one V, so a frame takes 1.5 bytes per pixel instead of RGBA's 4.

A fragment shader can't sample NV12 as if it were an RGBA texture, so something has to convert it, and I had two options. I could upload the Y and UV planes as two separate textures and do the math in the fragment shader. Or I could convert once, into RGBA, and keep the shader trivial. I chose the second, for a reason that only becomes clear once you see Step 6: the decoded frame comes back from NVDEC as a CUDA pointer. CUDA already has it. A CUDA kernel, one thread per pixel, is the shortest path from “NV12 I can read” to “RGBA a shader can sample”, and it keeps the WGSL to a single textureSample.

Why wgpu-hal: wgpu has to use memory it didn't allocate

Here's the crux of the whole project. wgpu normally owns every texture it uses: you call device.create_texture(&descriptor), wgpu allocates the memory, and wgpu tracks it from then on. In this program the texture has to already exist before wgpu sees it, because CUDA writes into it, and CUDA can't write into a texture wgpu allocated privately. I need a texture that was created outside wgpu, and then I need wgpu to adopt it.

The high-level wgpu API can't take an existing D3D12 resource. wgpu-hal can. It's the layer underneath wgpu that implements each backend (D3D12, Vulkan, Metal), and wgpu exposes it through a few unsafe functions. Device::as_hal gives me the underlying D3D12 device, so I can create resources on it. wgpu-hal's texture_from_raw wraps a resource I hand it. And Device::create_texture_from_hal registers the wrapped resource with wgpu proper. After that it's an ordinary wgpu::Texture: views, bind groups and textureSample all work as usual.

That's also why Direct3D 12 is involved at all, even though I never write a line of D3D12 rendering. wgpu's Windows backend is D3D12, and CUDA can import D3D12 resources as external memory. A D3D12 texture is the one kind of allocation both sides can use natively. So the plan is: create a D3D12 texture by hand, show it to CUDA, show the same one to wgpu.

Why a C++/CUDA bridge instead of pure Rust

Rust could call NVCUVID, CUDA and Media Foundation directly through FFI bindings, and for a library that's what I'd do. For this project I put all three in one C++ file, native/video_bridge.cu, for two practical reasons. First, I need a CUDA kernel anyway, which means nvcc is already in the build. Second, these are C APIs with lots of structs, callbacks and COM pointers, and they read most naturally in C++. The Rust side then sees only seven small functions.

What the bridge wraps
Interface	Why it's there
NVIDIA Video Codec SDK (NVCUVID)	The parser and the NVDEC decoder
CUDA (driver and runtime APIs)	The context, the imported texture memory and the color kernel
Media Foundation and Direct3D 12	Reading MP4 packets, and creating the shared textures

The trade-off I'm accepting is that all the unsafe interop lives in one file, behind a C ABI that hides Media Foundation and CUDA from Rust entirely. The Rust side can't misuse what it can't see.

Part 2: Set up the project

What you need

Windows 11 x64 and an NVIDIA GPU with NVDEC, with a current driver. The driver provides nvcuda.dll and nvcuvid.dll, the only NVIDIA files the program loads at run time.
Rust stable with the MSVC toolchain, and Visual Studio 2022 Build Tools with the Windows SDK. They provide the C++ compiler nvcc uses underneath, and the D3D12 and Media Foundation headers.
The CUDA Toolkit, which provides nvcc and sets CUDA_PATH.
The NVIDIA Video Codec SDK, downloaded and extracted. It has the NVCUVID headers and nvcuvid.lib. NVIDIA doesn't let you redistribute it, so every reader has to fetch it themselves.
An MP4 with progressive 8-bit 4:2:0 H.264 or HEVC, which is most of them.

What I built and tested with:

Test setup
Item	Version
GPU	NVIDIA GeForce RTX 5080 Laptop GPU, Windows 11
Rust / wgpu	cargo 1.97.1, wgpu 30.0.1
CUDA Toolkit	13.2
Video Codec SDK	13.1.15
Test clip	Big Buck Bunny, 320×180 H.264 (CC BY 3.0, Blender Foundation)

The project

Seven files. If you only read two, read native/video_bridge.cu (everything that talks to Media Foundation, NVDEC and CUDA) and src/video.rs (the Rust side of the bridge and the wgpu-hal import). The rest is setup and a standard wgpu cube that could be any mesh, a fullscreen quad or a material in an engine. To follow along with the files open, clone the repository:

git clone https://github.com/Lrios4403/nvdec-wgpu-cube.git
cd nvdec-wgpu-cube

Cargo.tomlwgpu 30.0.1 with only the DX12 backend
build.rscompiles the CUDA bridge with nvcc and links the SDKs
native/
- video_bridge.cuMedia Foundation, NVDEC, CUDA and the shared textures
- bitstream.hMP4 length prefixes to Annex B
- tests/
  - bitstream_test.cppportable tests for bitstream.h
src/
- main.rsthe wgpu device, bind groups and the cube
- video.rsthe FFI, the wgpu-hal import and frame timing
- cube.wgslvertex and fragment shaders

Step 1: Dependencies

The manifest is short, but two choices in it are deliberate.

Cargo.toml
1[package]
2name = "nvdec-wgpu-cube"
3version = "0.1.0"
4edition = "2021"
5build = "build.rs"
6
7[dependencies]
8wgpu = { version = "=30.0.1", default-features = false, features = ["std", "parking_lot", "dx12", "wgsl"] }
9winit = "0.30"
10pollster = "0.4"
11glam = "0.30"
12bytemuck = { version = "1.22", features = ["derive"] }
13windows = { version = "0.62", features = ["Win32_Graphics_Direct3D12"] }
14
15[build-dependencies]
16cc = "1.2"

wgpu is pinned to an exact version, =30.0.1, not 30. The wgpu-hal calls I make in Step 9 are the kind of thing that changes signature between releases (the number of arguments to texture_from_raw is different in older versions, for example). A pinned version means the code in this post compiles against what it was written for, and a deliberate upgrade means revisiting that one file.

Default features are off, and only dx12 and wgsl are turned back on. This isn't about binary size. The whole design depends on wgpu using D3D12, so I don't want a Vulkan backend even available to be picked by accident. std and parking_lot are just what wgpu needs to function at all.

Of the rest: windows gives me the ID3D12Resource type, which I need to hold the native texture pointer in Rust. cc is the build dependency that will drive nvcc. The others (winit, glam, bytemuck, pollster) are the usual wgpu companions for a window, matrices, vertex bytes and blocking on async setup.

Step 2: Compile CUDA from build.rs

Cargo doesn't know how to compile a .cu file, so a build script has to do it. The cc crate has a CUDA mode that runs nvcc and packages the result as a static library Cargo can link. The script also has to tell the linker about every native library the bridge uses.

build.rs
1use std::{env, path::PathBuf};
2
3fn main() {
  if env::var("CARGO_CFG_TARGET_OS").as_deref() != Ok("windows") {
      panic!("Windows x64 is required (D3D12 + Media Foundation + NVIDIA NVDEC)");
  }
  let cuda = PathBuf::from(env::var_os("CUDA_PATH").expect("Install the NVIDIA CUDA Toolkit (CUDA_PATH)"));
  let sdk = PathBuf::from(env::var_os("VIDEO_CODEC_SDK_DIR").expect("Set VIDEO_CODEC_SDK_DIR to the extracted NVIDIA Video Codec SDK"));
  let nvcuvid_dir = [sdk.join("Lib/win/x64"), sdk.join("Lib/x64")]
      .into_iter()
      .find(|dir| dir.join("nvcuvid.lib").exists())
      .expect("Could not locate SDK nvcuvid.lib (Lib/win/x64 or Lib/x64)");
  let headers = sdk.join("Interface");
  if !headers.join("nvcuvid.h").exists() || !headers.join("cuviddec.h").exists() {
      panic!("Video Codec SDK Interface/nvcuvid.h and cuviddec.h missing");
  }
  println!("cargo:rerun-if-changed=native/video_bridge.cu");
  println!("cargo:rerun-if-changed=native/bitstream.h");
  println!("cargo:rerun-if-env-changed=VIDEO_CODEC_SDK_DIR");
  println!("cargo:rerun-if-env-changed=CUDA_PATH");
  // "shared" links cudart.lib instead of cc's default, cudart_static.lib, which is
  // built against the static C runtime (LIBCMT) and clashes with Rust's /MD one
  // (linker warning LNK4098). With CUDA 13, cudart.lib is itself linked statically
  // and loads the driver's nvcuda.dll at run time, so no CUDA DLL ships with the exe.
  cc::Build::new()
      .cuda(true)
      .cudart("shared")
      .include(&headers)
      .include(cuda.join("include"))
      .file("native/video_bridge.cu")
      .flag("-std=c++17")
      .compile("video_bridge");
  println!("cargo:rustc-link-search=native={}", cuda.join("lib/x64").display());
  println!("cargo:rustc-link-search=native={}", nvcuvid_dir.display());
  // cudart itself is linked by cc (see .cudart above)
  for name in ["nvcuvid", "cuda", "d3d12", "dxgi", "mfplat", "mfreadwrite", "mfuuid", "ole32"] {
      println!("cargo:rustc-link-lib={name}");
  }
  // CUDA 13's cuda.lib is a static driver-API loader (it opens nvcuda.dll at
  // run time) built against the static C runtime. Bind it to the same DLL C
  // runtime as everything else instead of linking a second CRT.
  println!("cargo:rustc-link-arg=/NODEFAULTLIB:LIBCMT");
43}

Walking through it:

Lines 4–15 fail early. If this isn't Windows, or CUDA_PATH or VIDEO_CODEC_SDK_DIR is unset, or the SDK folder is missing its headers or library, the build stops with a message that says what to fix. Without these checks you get a wall of linker errors instead.
Lines 17–20 tell Cargo when to rerun. By default a build script reruns whenever anything in the package changes. Listing the two native files and the two environment variables means editing main.rs doesn't recompile the CUDA.
Lines 25–32 run nvcc on the bridge with the SDK's and CUDA's include folders, and compile it into video_bridge.lib.
Lines 33–37 link the native libraries: nvcuvid for the decoder API, cuda for the CUDA driver API, d3d12 and dxgi, and the three Media Foundation libraries plus ole32 for COM.

GotchaTwo C runtimes in one executable

Two lines took real digging, and they're both about the same problem. Rust links against the DLL version of the Microsoft C runtime. Some of NVIDIA's libraries are built against the static version. Link both and you get warning LNK4098, and a program that may or may not work depending on which runtime's malloc a given pointer came from.

.cudart("shared") on line 27. In CUDA mode, cc quietly links cudart_static.lib by default, the static-runtime build. Asking for "shared" links cudart.lib instead, which matches Rust.
/NODEFAULTLIB:LIBCMT on line 42. In CUDA 13, cuda.lib isn't an import library but a small static loader for the driver, also built against the static runtime. This flag tells the linker to ignore its request for LIBCMT and bind it to the same runtime as everything else.

A side effect worth knowing: both cuda.lib and cudart.lib load the driver's nvcuda.dll at run time, so the finished executable needs the NVIDIA driver but no CUDA Toolkit DLLs. You can copy it to another machine with the same driver and it runs.

Part 3: Decode on the GPU

The Bridge object and its helpers

Everything the native side knows lives in one struct. I'm showing it first so that when a later excerpt reads or writes a field, you've already seen where it came from:

native/video_bridge.cu, lines 78–98
78struct Bridge {
  ComPtr<IMFSourceReader> reader;
  ComPtr<IMFMediaType> compressed_type;
  ComPtr<ID3D12Device> d3d;
  bool com_started = false, mf_started = false;
  bool has_eos = false, flushed = false;
  unsigned width = 0, height = 0;
  unsigned decode_width = 0, decode_height = 0;
  cudaVideoCodec codec_type = cudaVideoCodec_H264;
  std::vector<uint8_t> sequence_header;
  std::vector<uint8_t> bitstream;
  NalFraming framing = NalFraming::Unknown; // Set by the first sample, kept across loops.
  std::deque<CUVIDPARSERDISPINFO> decoded;
  CUcontext cuda_context = nullptr;
  CUvideoparser parser = nullptr;
  CUvideodecoder decoder = nullptr;
  std::array<OutputSlot, kSlots> slots;
  std::string error;
  int current_slot = -1;
  int64_t first_timestamp = -1;
  int64_t last_pts = 0;

Lines 79–81: the Media Foundation reader and the compressed media type (Step 3), and the D3D12 device the textures are created on (Step 7).
Line 82: whether COM and Media Foundation were started, so the destructor only shuts down what was actually started.
Line 83: the end-of-file state machine (Step 6): has_eos is set when the reader hits the end, flushed once the parser has been told.
Lines 84–85: the picture size from the media type, and the coded size from the sequence callback, kept separately so a mid-stream size change can be detected (Step 5).
Lines 86–89: the codec, the parameter sets, the scratch buffer each converted packet is built in (reused every packet, so nothing is allocated per frame), and the stream's NAL framing (Step 3).
Line 90: decoded, the queue of frames ready to show, in display order. The display callback appends to it (Step 5) and the per-frame loop drains it (Step 6).
Lines 91–94: the CUDA context (Step 4), the parser and decoder (Step 5), and the three texture slots (Step 7).
Lines 95–98: the first error message, which slot was written last, and the two timestamp fields that turn the file's absolute timestamps into ones relative to the first frame (Step 6).

Then the helpers. Nearly every line in the steps that follows ends in one of these, so they're worth a minute:

native/video_bridge.cu, lines 100–120
100bool fail(const char* message) {
101    if (error.empty()) error = message;
102    std::fprintf(stderr, "[nvdec] %s\n", message);
103    return false;
104}
105bool check(CUresult result, const char* message) {
106    if (result == CUDA_SUCCESS) return true;
107    const char* name = "unknown";
108    cuGetErrorName(result, &name);
109    error = std::string(message) + ": " + (name ? name : "unknown");
110    std::fprintf(stderr, "[nvdec] %s\n", error.c_str());
111    return false;
112}
113bool check_cuda(cudaError_t result, const char* message) {
114    if (result == cudaSuccess) return true;
115    error = std::string(message) + ": " + cudaGetErrorString(result);
116    std::fprintf(stderr, "[nvdec] %s\n", error.c_str());
117    return false;
118}
119bool push() { return check(cuCtxPushCurrent(cuda_context), "cuCtxPushCurrent"); }
120void pop() { CUcontext old = nullptr; cuCtxPopCurrent(&old); }

Three ways to report an error, because the file talks to three kinds of API. fail takes a plain message, and is what the Media Foundation and D3D12 code uses after checking an HRESULT. check takes a CUDA driver-API result (CUresult, from functions starting with cu, which is what NVCUVID returns too) and looks up its name. check_cuda takes a CUDA runtime-API result (cudaError_t, from functions starting with cuda) and looks up its description. All three print to stderr immediately and return false, which is why call sites throughout the file read as if (!check(...)) return false;. Line 101 is a small but useful rule: only the first error is kept in error, because the first one is the cause and the ones after it are usually consequences. That first message is what Rust gets back through nv_error (Step 8).

push and pop (lines 119–120) make the bridge's CUDA context current on the thread and un-make it. Every method that does CUDA work is wrapped in them. Step 4 explains why, once the context exists; for now, read push() as “about to do CUDA work” and pop() as “done”.

Step 3: Read compressed packets

The goal of this step: get a stream of compressed packets, each with a timestamp, out of an MP4 file, and get them in a form NVIDIA's parser accepts. Media Foundation's Source Reader is the API for reading samples from a media file, and the important thing about it is that it can either decode for you or hand over the stream as it is in the file.

First, the boilerplate: COM and Media Foundation have to be started, and the file opened.

native/video_bridge.cu, lines 139–150
139HRESULT hr = CoInitializeEx(nullptr, COINIT_MULTITHREADED);
140// winit may already have created an STA on the Windows UI thread.
141if (hr == RPC_E_CHANGED_MODE) hr = CoInitializeEx(nullptr, COINIT_APARTMENTTHREADED);
142if (FAILED(hr)) return fail("CoInitializeEx failed");
143com_started = true;
144hr = MFStartup(MF_VERSION);
145if (FAILED(hr)) return fail("MFStartup failed");
146mf_started = true;
147const std::wstring wide = utf8_to_utf16(path);
148if (wide.empty()) return fail("Invalid UTF-8 video path");
149hr = MFCreateSourceReaderFromURL(wide.c_str(), nullptr, reader.GetAddressOf());
150if (FAILED(hr)) return fail("MFCreateSourceReaderFromURL failed (MP4 file required)");

Line 147 calls the one helper that isn't on the Bridge. Rust hands the file path over as a UTF-8 C string (Step 8), and Media Foundation wants a UTF-16 wide string:

native/video_bridge.cu, lines 56–63
56static std::wstring utf8_to_utf16(const char* src) {
57    int chars = MultiByteToWideChar(CP_UTF8, MB_ERR_INVALID_CHARS, src, -1, nullptr, 0);
58    if (chars <= 0) return {};
59    std::wstring out((size_t)chars, 0);
60    if (!MultiByteToWideChar(CP_UTF8, MB_ERR_INVALID_CHARS, src, -1, out.data(), chars)) return {};
61    out.pop_back();
62    return out;
63}

One call to measure, one to convert, and the trailing NUL that MultiByteToWideChar counts is dropped. MB_ERR_INVALID_CHARS makes malformed UTF-8 an error (an empty result, which line 148 turns into a fail) rather than a garbled path.

Now the part that matters: configuring the reader so it does not decode.

native/video_bridge.cu, lines 151–164
151// Disable all other tracks; preserve the *native compressed* video type.
152hr = reader->SetStreamSelection(kAllStreams, FALSE);
153if (FAILED(hr)) return fail("Disable other streams failed");
154hr = reader->SetStreamSelection(kVideoStream, TRUE);
155if (FAILED(hr)) return fail("Select video stream failed");
156hr = reader->GetNativeMediaType(kVideoStream, 0, compressed_type.GetAddressOf());
157if (FAILED(hr)) return fail("Cannot get native compressed video type");
158GUID subtype = {};
159if (FAILED(compressed_type->GetGUID(MF_MT_SUBTYPE, &subtype))) return fail("No video subtype");
160if (subtype == MFVideoFormat_H264) codec_type = cudaVideoCodec_H264;
161else if (subtype == MFVideoFormat_HEVC) codec_type = cudaVideoCodec_HEVC;
162else return fail("Only native compressed H.264/HEVC MP4 is supported");
163hr = reader->SetCurrentMediaType(kVideoStream, nullptr, compressed_type.Get());
164if (FAILED(hr)) return fail("Could not keep native compressed video output");

Three things happen here, and the order matters:

Lines 152–155: select only the video stream. The reader starts with every stream selected. I deselect all of them and reselect the first video stream, so ReadSample never hands me an audio packet.
Lines 156–162: find out what codec it is. The stream's native media type is the type it has in the file, before any decoding. Its subtype GUID tells me whether it's H.264 or HEVC, which I'll need to create the right parser. Anything else is rejected, since NVDEC would need different handling for it.
Lines 163–164: set that native type as the output type. This is the line that makes Media Foundation a demuxer instead of a decoder. If I asked for an RGB or NV12 output type instead, the reader would insert its own software or hardware decoder, and I'd be back to copying pixels around.

Two more pieces of information come from the media type: the frame size, and the codec's parameter sets.

native/video_bridge.cu, lines 165–179
165if (FAILED(MFGetAttributeSize(compressed_type.Get(), MF_MT_FRAME_SIZE, &width, &height)) ||
166    !width || !height || (width & 1) || (height & 1)) {
167    return fail("Missing/unsupported video frame size (even dimensions required)");
168}
169UINT32 header_size = 0;
170if (SUCCEEDED(compressed_type->GetBlobSize(MF_MT_MPEG_SEQUENCE_HEADER, &header_size)) && header_size) {
171    sequence_header.resize(header_size);
172    UINT32 got = 0;
173    if (FAILED(compressed_type->GetBlob(MF_MT_MPEG_SEQUENCE_HEADER,
174            sequence_header.data(), header_size, &got))) return fail("Sequence-header read failed");
175    sequence_header.resize(got);
176    // MF sequence headers are Annex B SPS/PPS or VPS/SPS/PPS, NOT avcC/hvcC.
177    if (!annexb_prefix(sequence_header.data(), sequence_header.size()))
178        return fail("Expected Annex B MF_MT_MPEG_SEQUENCE_HEADER");
179}

The frame size is needed to create the textures in Step 7, before any frame is decoded. The sequence header is something I didn't know I needed until the parser refused to produce frames without it. H.264 and HEVC streams start with parameter sets (the SPS and PPS; HEVC adds a VPS) that describe how the rest of the stream is coded. In an MP4, they're stored in the track's header rather than in the first packet. Media Foundation exposes them as MF_MT_MPEG_SEQUENCE_HEADER, and I feed them to the parser before any packet, in Step 5. Line 177 checks that they're in Annex B form, which brings me to the wrinkle.

NAL units and how they're separated. An H.264 or HEVC stream is a sequence of NAL units (a header byte or two, then a slice of a picture or a parameter set). There are two ways to mark where one ends and the next begins. A raw stream uses Annex B framing: each NAL unit starts with a start code, 00 00 01 or 00 00 00 01. Inside an MP4, NAL units are usually length-prefixed instead: four bytes giving the size, then the data. NVIDIA's own samples demux with FFmpeg and run its mp4toannexb filter before parsing, which converts the second form into the first. Media Foundation's MP4 source happened to deliver Annex B for my test clips, but I didn't want to depend on that, so the bridge accepts both and converts if needed. That logic is in native/bitstream.h:

native/bitstream.h, lines 23–48
23// True when nal_len_bytes-byte lengths tile the whole sample exactly, every
24// NAL unit is non-empty, and every NAL header has its forbidden_zero_bit (the
25// top bit, in both H.264 and HEVC) clear.
26inline bool length_prefixed(const uint8_t* p, size_t n, unsigned nal_len_bytes) {
27    if (!n || !(nal_len_bytes == 1 || nal_len_bytes == 2 || nal_len_bytes == 4)) return false;
28    while (n) {
29        if (n < nal_len_bytes) return false;
30        size_t count = 0;
31        for (unsigned i = 0; i < nal_len_bytes; ++i) count = (count << 8) | p[i];
32        p += nal_len_bytes;
33        n -= nal_len_bytes;
34        if (count == 0 || count > n || (p[0] & 0x80)) return false;
35        p += count;
36        n -= count;
37    }
38    return true;
39}
40
41// Which framing one sample uses. The length check wins: a 4-byte length of 1
42// or 256-511 starts with 00 00 00 01 / 00 00 01 xx, which looks exactly like a
43// start code, while real Annex B data practically never tiles as lengths.
44inline NalFraming detect_framing(const uint8_t* p, size_t n, unsigned nal_len_bytes = 4) {
45    if (length_prefixed(p, n, nal_len_bytes)) return NalFraming::LengthPrefixed;
46    if (annexb_prefix(p, n)) return NalFraming::AnnexB;
47    return NalFraming::Unknown;
48}

GotchaA length prefix that looks like a start code

My first version of this was wrong in a way that's instructive. It checked whether the sample started with a start code and, if so, passed it through unchanged. The trouble is that a 4-byte length of 300 is 00 00 01 2C, which looks exactly like a start code followed by a byte. So any length-prefixed sample whose first NAL unit was 256–511 bytes long (or exactly 1 byte, 00 00 00 01) was sent to the decoder with its length prefixes still in it.

The fix is to make the length-prefixed test the strict one and try it first: the lengths must cover the sample exactly, every NAL unit must be non-empty, and every NAL header's top bit (the forbidden zero bit, which the standards require to be 0) must be clear. Real Annex B data practically never passes that test by accident. And once the first sample has decided the framing, every later sample has to match it.

With that in place, reading one packet looks like this:

native/video_bridge.cu, lines 369–384
369bool read_packet() {
  DWORD flags = 0;
  LONGLONG timestamp = 0;
  ComPtr<IMFSample> sample;
  const HRESULT hr = reader->ReadSample(kVideoStream, 0,
      nullptr, &flags, &timestamp, sample.GetAddressOf());
  if (FAILED(hr)) return fail("IMFSourceReader::ReadSample failed");
  if (flags & (MF_SOURCE_READERF_NATIVEMEDIATYPECHANGED | MF_SOURCE_READERF_CURRENTMEDIATYPECHANGED))
      return fail("Dynamic compressed media-type changes unsupported");
  if (flags & MF_SOURCE_READERF_ENDOFSTREAM) { has_eos = true; return true; }
  if (!sample) return true; // stream tick / no sample
  ComPtr<IMFMediaBuffer> buf;
  if (FAILED(sample->ConvertToContiguousBuffer(buf.GetAddressOf()))) return fail("Cannot obtain compressed sample");
  BYTE* bytes = nullptr;
  DWORD count = 0;
  if (FAILED(buf->Lock(&bytes, nullptr, &count))) return fail("Compressed packet Lock failed");

ReadSample is synchronous: it returns the next sample of the video stream, or a flag. The flags on lines 376–379 cover the cases that aren't a packet: the stream changed format mid-file (rejected; handling it would mean recreating the decoder), the end of the stream (remembered as has_eos; Step 6 uses it), or a gap with no sample. A real sample comes with a timestamp in 100-nanosecond units, which I keep because NVDEC's packet timestamps use the same unit, and it comes as a buffer I have to lock to read.

native/video_bridge.cu, lines 385–408
385// MF's MPEG-4 source typically presents Annex B, but handle MP4 length-prefixed packets too.
386// The first sample decides which; every later sample must match it.
387const NalFraming before = framing;
388const bool converted = annexb_sample(bytes, count, bitstream, framing);
389buf->Unlock();
390if (!converted) {
391    return fail(before == NalFraming::AnnexB ? "Invalid compressed NAL units (stream started as Annex B)"
392        : before == NalFraming::LengthPrefixed ? "Invalid compressed NAL units (stream started with four-byte lengths)"
393        : "Invalid compressed NAL units (expected Annex B or four-byte lengths)");
394}
395if (before == NalFraming::Unknown) {
396    std::printf("[mf] NAL framing: %s\n", framing == NalFraming::AnnexB ? "Annex B" : "four-byte lengths");
397    std::fflush(stdout);
398}
399if (bitstream.size() > UINT32_MAX) return fail("Compressed packet too large");
400CUVIDSOURCEDATAPACKET pkt = {};
401pkt.flags = CUVID_PKT_TIMESTAMP;
402pkt.timestamp = timestamp; // MF sample time and CUVID timestamps both use 100ns ticks.
403pkt.payload = bitstream.data();
404pkt.payload_size = (unsigned long)bitstream.size();
405if (!push()) return false;
406const bool ok = check(cuvidParseVideoData(parser, &pkt), "cuvidParseVideoData(packet)");
407pop();
408return ok;

Lines 387–398 run the framing logic from above and, the first time, print which form the file uses. Lines 400–406 build the NVDEC packet and feed it to the parser: the (possibly converted) bytes, the timestamp, and a flag saying the timestamp is valid. The push() and pop() around cuvidParseVideoData are about CUDA contexts, which Step 4 explains. The point of this function is that by the time it returns, the parser has seen one more packet, and may have fired any of its three callbacks along the way.

bitstream.h has no Windows or CUDA dependencies, which is deliberate: it means the framing logic can be unit-tested on any machine, including the two cases that used to be misread. The test is a plain C++ file:

clang++ -std=c++17 -Wall -Wextra -pedantic native/tests/bitstream_test.cpp -o bitstream_test
./bitstream_test

Step 4: Make CUDA use the GPU wgpu uses

This step exists because of laptops. Mine has an Intel iGPU next to the RTX, and CUDA and wgpu each pick a device on their own. If wgpu renders on one GPU and CUDA writes textures on another, the shared texture in Step 7 means nothing: memory on one GPU isn't visible to the other. So before creating anything, the bridge has to find the CUDA device that is the GPU wgpu chose.

The Rust side does half of this. In main.rs I ask wgpu for a high-performance adapter and refuse anything that isn't NVIDIA:

src/main.rs, lines 84–99
84let mut desc = wgpu::InstanceDescriptor::new_without_display_handle();
85desc.backends = wgpu::Backends::DX12;
86
87let instance = wgpu::Instance::new(desc);
88let surface = instance.create_surface(window.clone())?;
89let adapter = instance.request_adapter(&wgpu::RequestAdapterOptions {
90    power_preference: wgpu::PowerPreference::HighPerformance,
91    compatible_surface: Some(&surface),
92    force_fallback_adapter: false,
93    apply_limit_buckets: false,
94}).await?;
95let adapter_info = adapter.get_info();
96println!("wgpu adapter: {} ({:?})", adapter_info.name, adapter_info.backend);
97if adapter_info.vendor != 0x10DE {
98    return Err("This sample requires wgpu to use an NVIDIA GPU (vendor 0x10DE)".into());
99}

That guarantees wgpu is on an NVIDIA GPU. The bridge then has to find the matching CUDA device. The thing that links them is the LUID, a locally unique ID that Windows assigns to each graphics adapter. D3D12 can report the LUID of the adapter a device was created on, and CUDA can report the LUID of each device it sees. Match them, and you have the same physical GPU on both sides:

native/video_bridge.cu, lines 186–213
186bool init_cuda(ID3D12Device* dev) {
  if (!dev) return fail("Missing wgpu D3D12 device");
  d3d = dev;
  if (!check(cuInit(0), "cuInit")) return false;
  int count = 0;
  if (!check(cuDeviceGetCount(&count), "cuDeviceGetCount")) return false;
  const LUID adapter_luid = dev->GetAdapterLuid();
  CUdevice match = -1;
  for (int i = 0; i < count; ++i) {
      CUdevice candidate = 0;
      char cuda_luid[8] = {};
      unsigned node_mask = 0;
      if (cuDeviceGet(&candidate, i) == CUDA_SUCCESS &&
          cuDeviceGetLuid(cuda_luid, &node_mask, candidate) == CUDA_SUCCESS &&
          std::memcmp(cuda_luid, &adapter_luid, sizeof(cuda_luid)) == 0) {
          match = candidate;
          break;
      }
  }
  if (match < 0) return fail("CUDA device LUID doesn't match the wgpu D3D12 NVIDIA adapter");
  if (!check(cuCtxCreate(
      &cuda_context,
      nullptr,
      0,
      match
  ), "cuCtxCreate")) return false;
  // cuCtxCreate makes this context current; pop so every call explicitly pushes it.
  pop();

Lines 192–205 are the search: ask the D3D12 device for its adapter LUID, then walk CUDA's device list comparing LUIDs byte for byte. Line 206 creates a CUDA context on the match. A context is CUDA's per-device state, and every CUDA call acts on whichever context is “current” on the calling thread.

Lines 212–214 are a small discipline I adopted after getting bitten. cuCtxCreate makes the new context current on the thread, implicitly. I pop it straight back off, and from then on every function that does CUDA work pushes the context first and pops it after, with the push() and pop() helpers from the Bridge section. That way nothing depends on which context happens to be current on the thread, which matters because winit, wgpu and the decoder callbacks all share that thread. You'll see the same pair around every CUDA call from here on: in read_packet, in next, and in the destructor.

Step 5: The parser and the decoder

Now the NVDEC objects. Recall the shape from Part 1: a parser that runs on the CPU and a decoder that runs on the hardware, with the parser driving the decoder through callbacks. The order of creation is forced on me. The parser can be created up front, because all it needs is the codec. The decoder can't, because its settings (picture size, how many surfaces to allocate) depend on the stream, and the parser only learns those from the sequence header. So: create the parser, feed it the sequence header, and create the decoder from inside the sequence callback.

native/video_bridge.cu, lines 325–345
325bool create_parser() {
  if (parser) { if (!check(cuvidDestroyVideoParser(parser), "cuvidDestroyVideoParser")) return false; parser = nullptr; }
  if (decoder) { if (!check(cuvidDestroyDecoder(decoder), "cuvidDestroyDecoder")) return false; decoder = nullptr; }
  decode_width = 0; decode_height = 0;
  CUVIDPARSERPARAMS parameters = {};
  parameters.CodecType = codec_type;
  parameters.ulMaxNumDecodeSurfaces = 20; // Sequence callback can override.
  parameters.ulMaxDisplayDelay = 2; // Frames the parser may hold before the display callback (0 = none).
  parameters.pUserData = this;
  parameters.pfnSequenceCallback = &Bridge::sequence_cb;
  parameters.pfnDecodePicture = &Bridge::decode_cb;
  parameters.pfnDisplayPicture = &Bridge::display_cb;
  if (!check(cuvidCreateVideoParser(&parser, &parameters), "cuvidCreateVideoParser")) return false;
  if (!sequence_header.empty()) {
      CUVIDSOURCEDATAPACKET pkt = {};
      pkt.payload = sequence_header.data();
      pkt.payload_size = (unsigned long)sequence_header.size();
      if (!check(cuvidParseVideoData(parser, &pkt), "cuvidParseVideoData(sequence header)")) return false;
  }
  return true;
345}

Lines 326–328 tear down any existing parser and decoder first, because this function is also what the bridge calls when the clip loops (Step 6). Lines 329–337 create the parser. The three callbacks and pUserData are how C APIs do methods: the parser calls a plain function and passes back whatever pointer I gave it, which is the Bridge itself. ulMaxDisplayDelay is the guide's “max display callback delay”: how many decoded frames the parser may hold before handing one to the display callback (0 means none). The reordering of B-frames is the parser's job either way. Lines 338–343 feed the sequence header from Step 3, which is what triggers the sequence callback for the first time.

The callbacks themselves are static functions that forward to the bridge:

native/video_bridge.cu, lines 264–277
264static int CUDAAPI sequence_cb(void* user, CUVIDEOFORMAT* fmt) {
265    return static_cast<Bridge*>(user)->on_sequence(fmt);
266}
267static int CUDAAPI decode_cb(void* user, CUVIDPICPARAMS* pic) {
268    auto* b = static_cast<Bridge*>(user);
269    return b->decoder && b->check(cuvidDecodePicture(b->decoder, pic), "cuvidDecodePicture") ? 1 : 0;
270}
271static int CUDAAPI display_cb(void* user, CUVIDPARSERDISPINFO* info) {
272    auto* b = static_cast<Bridge*>(user);
273    if (!info) return 1;
274    if (b->decoded.size() >= 64) { b->fail("Decoder display queue overflow"); return 0; }
275    b->decoded.push_back(*info); // Parser delivers display/reordered order.
276    return 1;
277}

The decode callback (lines 267–270) is one line of work: hand the picture to cuvidDecodePicture. The display callback (lines 271–277) doesn't display anything. It just appends the frame's info to the decoded queue from the Bridge section, in display order, and refuses if the queue has somehow grown to 64 entries, which would mean frames aren't being consumed. The per-frame loop in Step 6 drains that queue. Keeping the callback this small is deliberate: it runs from inside cuvidParseVideoData, and I don't want to map frames or run kernels from there.

The sequence callback is the one with real decisions in it. First, what the bridge refuses to handle:

native/video_bridge.cu, lines 278–291
278int on_sequence(CUVIDEOFORMAT* f) {
  if (!f || f->codec != codec_type || f->chroma_format != cudaVideoChromaFormat_420 ||
      f->bit_depth_luma_minus8 != 0 || !f->progressive_sequence) {
      fail("Only progressive 8-bit 4:2:0 video is supported (NV12)"); return 0;
  }
  if (decoder) {
      if (decode_width != f->coded_width || decode_height != f->coded_height) {
          fail("Midstream resolution changes are not supported"); return 0;
      }
      return std::max(
          2u,
          static_cast<unsigned int>(f->min_num_decode_surfaces)
      );
  }

Lines 279–282 limit the program to progressive 8-bit 4:2:0 video. That isn't an NVDEC limit; it's mine. The color kernel in Step 7 assumes NV12, and 10-bit video (P010) or 4:2:2 would need a different output format and a different kernel. Rejecting them with a clear message beats producing garbage. Lines 283–291 handle the case where the sequence callback fires again mid-stream: if the size changed, I give up (handling it would mean recreating the decoder and the textures), and otherwise I just return the surface count.

native/video_bridge.cu, lines 292–323
292CUVIDDECODECAPS caps = {};
293caps.eCodecType = f->codec;
294caps.eChromaFormat = f->chroma_format;
295caps.nBitDepthMinus8 = f->bit_depth_luma_minus8;
296if (!check(cuvidGetDecoderCaps(&caps), "cuvidGetDecoderCaps") || !caps.bIsSupported) {
297    fail("Codec/profile not supported by NVIDIA hardware"); return 0;
298}
299decode_width = f->coded_width;
300decode_height = f->coded_height;
301CUVIDDECODECREATEINFO create = {};
302create.CodecType = f->codec;
303create.ChromaFormat = f->chroma_format;
304create.OutputFormat = cudaVideoSurfaceFormat_NV12;
305create.DeinterlaceMode = cudaVideoDeinterlaceMode_Weave;
306create.ulCreationFlags = cudaVideoCreate_PreferCUVID;
307create.ulNumDecodeSurfaces = std::max(
308    2u,
309    static_cast<unsigned int>(f->min_num_decode_surfaces)
310);
311create.ulNumOutputSurfaces = 2;
312create.ulWidth = f->coded_width;
313create.ulHeight = f->coded_height;
314create.ulMaxWidth = f->coded_width;
315create.ulMaxHeight = f->coded_height;
316create.display_area.left = f->display_area.left;
317create.display_area.top = f->display_area.top;
318create.display_area.right = f->display_area.right;
319create.display_area.bottom = f->display_area.bottom;
320create.ulTargetWidth = width;
321create.ulTargetHeight = height;
322if (!check(cuvidCreateDecoder(&decoder, &create), "cuvidCreateDecoder")) return 0;
323return (int)create.ulNumDecodeSurfaces;

Lines 292–298 ask the driver whether this GPU's NVDEC can decode this codec at this chroma format and bit depth. Older GPUs lack HEVC, for example, and it's better to learn that here than from a cryptic failure later. Lines 301–322 fill in the decoder settings. A few of them are worth understanding:

OutputFormat = NV12 is the format cuvidMapVideoFrame will produce, which is what the kernel expects.
ulNumDecodeSurfaces is how many pictures NVDEC can have in flight. Codecs reference earlier frames, so the decoder has to keep several around. min_num_decode_surfaces is the stream's own minimum, and NVIDIA's guide recommends using it to keep memory use down. Returning that count from the callback (line 323) tells the parser to use it too.
ulNumOutputSurfaces = 2 is a separate pool: the surfaces that mapped frames go into. Two is enough here because I map one frame at a time (Step 6).
The two sizes. ulWidth and ulHeight are the coded size, which can be larger than the picture because codecs store frames in whole blocks (16×16 macroblocks in H.264) and crop the rest. display_area and ulTargetWidth/ulTargetHeight describe the picture I actually want, and NVDEC applies them in its post-processing when a frame is mapped. That's why the textures in Step 7 can be exactly the video's size.

Step 6: The per-frame loop: feed, decode, map

Everything so far was setup. This step is what runs every frame. The bridge exposes it to Rust as nv_next, and the contract is simple: “give me the next frame in display order, converted and sitting in one of the textures, and tell me which texture and what time it's due.” Internally it has two halves. The first feeds the parser until a frame is available:

native/video_bridge.cu, lines 410–430
410// Return 1 on success, -1 on error. out_pts: presentation time in 100ns units.
411int next(int* out_slot, int64_t* out_pts) {
  if (!out_slot || !out_pts) { fail("Null frame result"); return -1; }
  // Read compressed packets only until a display-order frame becomes available.
  // Consume that frame before submitting additional packets: protects NVDEC output surfaces.
  for (unsigned steps = 0; decoded.empty(); ++steps) {
      if (steps > 100000) { fail("No decodable video frames"); return -1; }
      if (!has_eos) {
          if (!read_packet()) return -1;
      } else if (!flushed) {
          CUVIDSOURCEDATAPACKET eos = {};
          eos.flags = CUVID_PKT_ENDOFSTREAM;
          if (!push()) return -1;
          bool ok = check(cuvidParseVideoData(parser, &eos), "cuvidParseVideoData(EOS)");
          pop();
          if (!ok) return -1;
          flushed = true;
      } else {
          if (!rewind_video()) return -1;
      }
  }

The loop on lines 415–430 runs until the display queue from Step 5 has something in it. Each iteration does one of three things, depending on where in the file we are:

Normally, read one more packet (line 418). That may or may not produce a frame: with B-frames, the parser often needs several packets before it can release one in display order, which is exactly why this is a loop and not a single read.
At the end of the file, flush. When Media Foundation reported end-of-stream, the parser may still be holding frames it hasn't released. Sending a packet with the ENDOFSTREAM flag (lines 420–426) tells it no more data is coming, so it releases them all.
Once those are drained, rewind (line 428), which is what makes the clip loop. The helper that does it is below.

Line 416 is a safety valve: a file with no decodable video would otherwise spin here forever.

native/video_bridge.cu, lines 352–368
352bool rewind_video() {
  // Do not destroy the parser until *every* display callback frame was consumed.
  PROPVARIANT seek; PropVariantInit(&seek);
  seek.vt = VT_I8;
  seek.hVal.QuadPart = 0;
  HRESULT hr = reader->SetCurrentPosition(GUID_NULL, seek);
  if (FAILED(hr)) return fail("Media Foundation failed to seek to video beginning");
  has_eos = false;
  flushed = false;
  first_timestamp = -1;
  last_pts = 0;
  decoded.clear();
  if (!push()) return false;
  const bool ok = create_parser();
  pop();
  return ok;
368}

Lines 354–358 seek the Source Reader back to time zero. Media Foundation seeks take a PROPVARIANT holding a 64-bit time in 100 ns units, which is the ceremony on lines 354–356. Lines 359–363 reset the per-playthrough state: the end-of-file flags, the timestamp base (so the second loop also starts at zero), and the display queue, which is empty by now anyway. Line 365 recreates the parser and the decoder through create_parser from Step 5, because a parser can't be rewound; it has to see the sequence header again from the start. The comment on line 353 is why the loop above drains the queue first: the frames in it live in the old decoder's surfaces, and recreating the decoder would destroy them. Notice what is not reset: framing. The file's NAL framing doesn't change between loops, so the first sample's decision stands.

The second half takes the frame at the front of the queue and maps it:

native/video_bridge.cu, lines 431–443
431CUVIDPARSERDISPINFO display = decoded.front();
432decoded.pop_front();
433const int slot_index = (current_slot + 1) % kSlots;
434if (!push()) return -1;
435CUVIDPROCPARAMS process = {};
436process.progressive_frame = display.progressive_frame;
437process.top_field_first = display.top_field_first;
438process.second_field = 0;
439process.unpaired_field = display.repeat_first_field < 0;
440CUdeviceptr device_frame = 0;
441unsigned pitch = 0;
442bool ok = check(cuvidMapVideoFrame(decoder, display.picture_index,
443    &device_frame, &pitch, &process), "cuvidMapVideoFrame");

Line 433 picks which of the three textures this frame will go into: the slot after the current one, wrapping around. Lines 435–443 map the frame. cuvidMapVideoFrame is where NVDEC's post-processing runs: it takes the decoded picture (identified by picture_index), converts it to the NV12 output format, applies the cropping from Step 5, and hands back two things: a CUDA device pointer to the result, and its row pitch. The pitch is the number of bytes from one row to the next, which is usually wider than the image because of alignment, so the kernel can't assume rows are packed.

What I'm holding now is the thing the whole project revolves around: a pointer to a decoded frame, in GPU memory, that CUDA can read. The next excerpt converts it.

native/video_bridge.cu, lines 444–462
444if (ok) {
445    const dim3 threads(16, 16);
446    const dim3 blocks((width + 15) / 16, (height + 15) / 16);
447    nv12_to_rgba<<<blocks, threads>>>(slots[slot_index].surface,
448        reinterpret_cast<const uint8_t*>(device_frame), pitch, width, height);
449    ok = check_cuda(cudaGetLastError(), "CUDA color-convert launch") &&
450         check_cuda(cudaDeviceSynchronize(), "CUDA color-convert synchronize");
451    // Unmap only once the CUDA kernel has finished accessing NVDEC memory.
452    const bool unmapped = check(cuvidUnmapVideoFrame(decoder, device_frame), "cuvidUnmapVideoFrame");
453    ok = unmapped && ok;
454}
455pop();
456if (!ok) return -1;
457current_slot = slot_index;
458if (first_timestamp < 0) first_timestamp = display.timestamp;
459last_pts = std::max<int64_t>(0, display.timestamp - first_timestamp);
460*out_slot = slot_index;
461*out_pts = last_pts;
462return 1;

Lines 445–448 launch the kernel: one thread per pixel, in 16×16 blocks, reading the NV12 frame and writing RGBA into the chosen slot's surface (the kernel and the surface are Step 7). Lines 449–450 wait for it to finish. Then line 452 unmaps the frame.

Lines 457–462 record the result: which slot has the frame, and its timestamp made relative to the first frame, so playback starts at zero no matter what the file's first timestamp is.

Step 7: Three textures that CUDA writes and wgpu reads

The kernel in Step 6 wrote into “the chosen slot's surface”. This step is where those surfaces come from, and it's the step the whole design hinges on. I need a texture that two different APIs can use. CUDA has to be able to write it, as a surface object (CUDA's name for a writable texture). wgpu has to be able to sample it. And there must be one allocation, not two with a copy between them.

One texture, two users

Each of the three slots is a single allocation that CUDA writes and wgpu samples.

Show as text

The bridge creates an RGBA ID3D12Resource with CreateCommittedResource on wgpu's own D3D12 device, in a shared heap with simultaneous access allowed.
For CUDA: CreateSharedHandle gives an NT handle, cudaImportExternalMemory imports the resource, and a CUDA surface object over it is what the kernel writes RGBA pixels into.
For wgpu: the same resource pointer (no handle needed, it's the same device) goes to texture_from_raw and create_texture_from_hal, which give an ordinary wgpu::Texture. The shader samples it through a bind group.
There is one allocation per slot. CUDA and wgpu read and write the same memory, so they must take turns.

The plan from Part 1 was: create the texture in D3D12, because that's the one place both sides can reach. The bridge creates it on wgpu's own D3D12 device, which the Rust side hands over in Step 9. Each texture, with the CUDA objects that will be stacked on it, is one OutputSlot, and the bridge has an array of three of them (line 94 of the Bridge struct):

native/video_bridge.cu, lines 64–76
64struct OutputSlot {
  ComPtr<ID3D12Resource> texture;
  cudaExternalMemory_t memory = nullptr;
  cudaMipmappedArray_t mipmap = nullptr;
  cudaSurfaceObject_t surface = 0;
  void cleanup() {
      if (surface) cudaDestroySurfaceObject(surface);
      if (mipmap) cudaFreeMipmappedArray(mipmap);
      if (memory) cudaDestroyExternalMemory(memory);
      surface = 0; mipmap = nullptr; memory = nullptr;
      texture.Reset();
  }
76};

The four fields are filled in top to bottom by the code below: the D3D12 resource, then CUDA's view of its memory, then an array over that memory, then the surface object the kernel writes through. cleanup destroys them in the reverse order, because each one refers to the one before it, and resets the fields so that calling it twice is harmless. The destructor (Step 9) calls it for each slot.

Here's the creation, which runs once for each of the three slots:

native/video_bridge.cu, lines 216–234
216for (auto& slot : slots) {
  D3D12_HEAP_PROPERTIES heap = {};
  heap.Type = D3D12_HEAP_TYPE_DEFAULT;
  D3D12_RESOURCE_DESC desc = {};
  desc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D;
  desc.Width = width;
  desc.Height = height;
  desc.DepthOrArraySize = 1;
  desc.MipLevels = 1;
  desc.Format = DXGI_FORMAT_R8G8B8A8_UNORM;
  desc.SampleDesc.Count = 1;
  desc.Layout = D3D12_TEXTURE_LAYOUT_UNKNOWN;
  desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_SIMULTANEOUS_ACCESS;
  HRESULT hr = dev->CreateCommittedResource(&heap, D3D12_HEAP_FLAG_SHARED, &desc,
      D3D12_RESOURCE_STATE_COMMON, nullptr, IID_PPV_ARGS(&slot.texture));
  if (FAILED(hr)) { ok = fail("CreateCommittedResource failed"); break; }
  HANDLE shared = nullptr;
  hr = dev->CreateSharedHandle(slot.texture.Get(), nullptr, GENERIC_ALL, nullptr, &shared);
  if (FAILED(hr)) { ok = fail("CreateSharedHandle failed"); break; }

Three settings on the resource make the sharing possible:

D3D12_HEAP_FLAG_SHARED (line 229). Without this, the allocation is private to the D3D12 device and can't be opened by anything else.
ALLOW_SIMULTANEOUS_ACCESS (line 228). D3D12 normally requires explicit “resource state” transitions when a texture changes use (written by one thing, then read by another). This flag opts the texture out of that system, so CUDA and wgpu can both touch it without either knowing the other's state. Step 9 relies on this.
R8G8B8A8_UNORM (line 225). This has to match on all three sides: it's what the kernel writes (uchar4), what CUDA's array format describes (line 245), and what wgpu is told in Step 9 (Rgba8Unorm).

Lines 232–234 then ask D3D12 for an NT handle to the resource. That's the only way CUDA can open a D3D12 resource: through a handle, not a pointer. (wgpu won't need the handle, because it's on the same device and can use the resource pointer directly.) The import into CUDA follows:

native/video_bridge.cu, lines 235–258
235cudaExternalMemoryHandleDesc memory_desc = {};
236memory_desc.type = cudaExternalMemoryHandleTypeD3D12Resource;
237memory_desc.handle.win32.handle = shared;
238memory_desc.size = dev->GetResourceAllocationInfo(0, 1, &desc).SizeInBytes;
239memory_desc.flags = cudaExternalMemoryDedicated;
240const cudaError_t imported = cudaImportExternalMemory(&slot.memory, &memory_desc);
241CloseHandle(shared);
242if (!check_cuda(imported, "cudaImportExternalMemory")) { ok = false; break; }
243cudaExternalMemoryMipmappedArrayDesc mapping = {};
244mapping.offset = 0;
245mapping.formatDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
246mapping.extent = make_cudaExtent(width, height, 1);
247mapping.flags = cudaArraySurfaceLoadStore;
248mapping.numLevels = 1;
249if (!check_cuda(cudaExternalMemoryGetMappedMipmappedArray(&slot.mipmap, slot.memory, &mapping),
250                "cudaExternalMemoryGetMappedMipmappedArray")) { ok = false; break; }
251cudaArray_t array = nullptr;
252if (!check_cuda(cudaGetMipmappedArrayLevel(&array, slot.mipmap, 0),
253                "cudaGetMipmappedArrayLevel")) { ok = false; break; }
254cudaResourceDesc resource = {};
255resource.resType = cudaResourceTypeArray;
256resource.res.array.array = array;
257if (!check_cuda(cudaCreateSurfaceObject(&slot.surface, &resource),
258                "cudaCreateSurfaceObject")) { ok = false; break; }

This is three CUDA objects stacked on top of each other, and each exists for a reason:

External memory (lines 235–242). cudaImportExternalMemory takes the handle and the allocation's size (asked from D3D12, since a texture's allocation is bigger than width × height × 4) and gives CUDA a reference to the same GPU memory. The handle can be closed as soon as the import succeeds; CUDA keeps its own reference.
A mipmapped array over it (lines 243–253). External memory is just bytes to CUDA. To treat it as a 2D image with a known format, I map a CUDA array over it, describing the format (four 8-bit unsigned channels), the size, and that it will be written through a surface. One mip level, since it's a video frame.
A surface object (lines 254–258). This is the handle the kernel actually writes through, with surf2Dwrite. It's what gets passed to the kernel launch in Step 6.

And the kernel that writes into it:

native/video_bridge.cu, lines 32–54
32// Simple SDR BT.709 *limited-range* matrix; HDR/P010 and metadata conversion
33// require a different kernel and are rejected in on_sequence for this example.
34__device__ __forceinline__ unsigned char byte_clamp(float v) {
  return (unsigned char)fminf(255.0f, fmaxf(0.0f, v));
36}
37__global__ void nv12_to_rgba(cudaSurfaceObject_t dst, const uint8_t* nv12,
                           unsigned pitch, unsigned width, unsigned height) {
  const unsigned x = blockIdx.x * blockDim.x + threadIdx.x;
  const unsigned y = blockIdx.y * blockDim.y + threadIdx.y;
  if (x >= width || y >= height) return;
  const uint8_t* y_plane = nv12;
  const uint8_t* uv_plane = nv12 + (size_t)pitch * height;
  const float Y = (float)y_plane[(size_t)y * pitch + x];
  const unsigned uv_index = (y / 2) * pitch + (x / 2) * 2;
  const float U = (float)uv_plane[uv_index] - 128.0f;
  const float V = (float)uv_plane[uv_index + 1] - 128.0f;
  const float c = 1.16438356f * (Y - 16.0f);
  const uchar4 rgba = make_uchar4(
      byte_clamp(c + 1.79274107f * V),
      byte_clamp(c - 0.21324861f * U - 0.53290933f * V),
      byte_clamp(c + 2.11240179f * U), 255);
  surf2Dwrite(rgba, dst, x * sizeof(uchar4), y);
54}

Each thread handles one pixel (lines 39–41). It reads its Y value from the first plane, then finds the U/V pair its 2×2 block shares in the second plane (line 45: halve both coordinates, and the pairs are interleaved, so multiply x by 2). Then it applies the BT.709 limited-range conversion. “Limited range” means the codec stores black as 16 and white as 235, not 0 and 255, which is why the math starts with 1.164 × (Y − 16): that stretches 16–235 back out to 0–255. The constants on lines 50–52 are the standard's coefficients for how much V contributes to red, how much U and V subtract from green, and how much U contributes to blue. The result is clamped to 0–255 and written as one RGBA pixel.

The kernel always assumes BT.709 limited range. That's right for HD video and for my clip, but a real player would read the stream's color metadata and pick the matrix accordingly. I come back to that in the limitations.

Part 4: Hand the frames to wgpu

Step 8: A small C ABI

I want the Rust side to know as little as possible about Media Foundation, NVDEC and CUDA. So the bridge exports seven plain C functions, and the Bridge object is an opaque pointer to Rust: it can hold it and pass it back, but not look inside.

native/video_bridge.cu, lines 466–485
466extern "C" {
467__declspec(dllexport) Bridge* nv_open(const char* path, void* d3d_device) {
468    std::unique_ptr<Bridge> bridge(new (std::nothrow) Bridge());
469    if (!bridge || !bridge->init(path, reinterpret_cast<ID3D12Device*>(d3d_device))) return nullptr;
470    return bridge.release();
471}
472__declspec(dllexport) int nv_width(Bridge* b) { return b ? (int)b->width : 0; }
473__declspec(dllexport) int nv_height(Bridge* b) { return b ? (int)b->height : 0; }
474__declspec(dllexport) void* nv_texture(Bridge* b, int index) {
475    if (!b || index < 0 || index >= kSlots) return nullptr;
476    ID3D12Resource* texture = b->slots[index].texture.Get();
477    if (texture) texture->AddRef(); // Rust from_raw takes this additional reference.
478    return texture;
479}
480__declspec(dllexport) int nv_next(Bridge* b, int* slot, int64_t* pts) {
481    return b ? b->next(slot, pts) : -1;
482}
483__declspec(dllexport) const char* nv_error(Bridge* b) { return b ? b->error.c_str() : "No video bridge"; }
484__declspec(dllexport) void nv_close(Bridge* b) { delete b; }
485}

Most of these are thin. The two that carry the design are nv_next, which runs Step 6 and reports a slot and a timestamp, and nv_texture, which hands out the D3D12 resource for a slot. Line 477 is the subtle part of nv_texture: COM objects are reference-counted, and the bridge calls AddRef before returning the pointer. That gives the Rust side a reference of its own, so the resource stays alive even after the bridge releases its reference in nv_close. Step 9 takes ownership of that extra reference.

On the Rust side, the declarations mirror the exports:

src/video.rs, lines 10–20
10#[repr(C)]
11struct VideoHandle { _opaque: [u8; 0] }
12unsafe extern "C" {
13    fn nv_open(filename: *const c_char, device: *mut c_void) -> *mut VideoHandle;
14    fn nv_width(ptr: *mut VideoHandle) -> i32;
15    fn nv_height(ptr: *mut VideoHandle) -> i32;
16    fn nv_texture(ptr: *mut VideoHandle, slot: i32) -> *mut c_void;
17    fn nv_next(ptr: *mut VideoHandle, slot: *mut i32, pts: *mut i64) -> i32;
18    fn nv_error(ptr: *mut VideoHandle) -> *const c_char;
19    fn nv_close(ptr: *mut VideoHandle);
20}

VideoHandle on lines 10–11 is the Rust idiom for an opaque type: a struct with no fields that can only be used through a pointer. Nothing checks the two declaration lists against each other, so a change to one has to be made in both. That's the cost of the simplicity.

Step 9: Import the textures with wgpu-hal

This is the step the post is named for. Video::open in src/video.rs is where the Rust side and the bridge meet, and it does four things: get the raw D3D12 device out of wgpu, open the bridge on it, decode the first frames, and wrap each slot's resource as a wgpu::Texture.

src/video.rs, lines 54–67
54pub fn open(device: &wgpu::Device, filename: &str) -> Result<Self, String> {
  let hal = unsafe { device.as_hal::<wgpu::hal::api::Dx12>() }
      .ok_or("wgpu must use the D3D12 backend")?;
  let file = CString::new(filename).map_err(|e| e.to_string())?;
  let native = NonNull::new(unsafe { nv_open(file.as_ptr(), hal.raw_device().as_raw()) })
      .ok_or("Media Foundation / Video Codec SDK initialization failed (see stderr)")?;
  // Decode initial + lookahead frame before importing the D3D12 resources into wgpu.
  // On failure, always release the native decoder and CUDA resources.
  let result = (|| {
      let width = unsafe { nv_width(native.as_ptr()) } as u32;
      let height = unsafe { nv_height(native.as_ptr()) } as u32;
      if width == 0 || height == 0 { return Err("Invalid video dimensions".to_owned()); }
      let current = Self::decode(native)?;
      let pending = Self::decode(native)?;

Lines 55–56 are the first unsafe call, and the first place wgpu-hal appears. as_hal reaches under wgpu's abstraction and returns the backend-specific device, here wgpu::hal::dx12::Device. From that I can get the raw ID3D12Device pointer, which is what the bridge needs in Step 7 to create the textures on the same device wgpu renders with. If wgpu weren't using D3D12, as_hal returns None and the program stops here with a clear message. Line 58 opens the bridge, which runs all of Part 3's setup. Lines 66–67 decode two frames before wgpu wraps anything: the frame to show first, and one frame of lookahead (Step 11 explains why there's a lookahead). Self::decode is the Rust wrapper around nv_next, and it returns a DecodedFrame, so here are both:

src/video.rs, lines 22–37
22#[derive(Clone, Copy)]
23struct DecodedFrame { slot: usize, pts_100ns: i64 }
24
25pub struct Video {
26    native: NonNull<VideoHandle>,
27    pub textures: Vec<wgpu::Texture>,
28    current: DecodedFrame,
29    pending: DecodedFrame,
30    start: Instant,
31    next_due: Instant,
32    fallback_interval: Duration,
33}
34
35fn duration_100ns(ticks: i64) -> Duration {
36    Duration::from_nanos((ticks.max(0) as u64).saturating_mul(100))
37}

A DecodedFrame is the answer to one nv_next call: which slot the frame is in and when it's due. Video itself holds the bridge pointer, the three wrapped textures (public, because main.rs builds bind groups from them), the two frames, and the clock: when playback started, when the pending frame is due, and the last known interval between frames, for files whose timestamps don't advance. duration_100ns converts the bridge's 100-nanosecond ticks to a Duration, clamping negative values to zero and saturating rather than overflowing on absurd inputs.

src/video.rs, lines 40–52
40fn decode(handle: NonNull<VideoHandle>) -> Result<DecodedFrame, String> {
  let mut slot = -1;
  let mut pts = -1_i64;
  if unsafe { nv_next(handle.as_ptr(), &mut slot, &mut pts) } != 1 {
      let ptr = unsafe { nv_error(handle.as_ptr()) };
      return Err(if ptr.is_null() { "Native NVDEC decoder error".into() }
          else { unsafe { CStr::from_ptr(ptr) }.to_string_lossy().into_owned() });
  }
  if !(0..SLOTS as i32).contains(&slot) {
      return Err(format!("Native NVDEC returned invalid texture slot {slot}"));
  }
  Ok(DecodedFrame { slot: slot as usize, pts_100ns: pts })
52}

decode is the only place nv_next is called, and it turns the C convention (return 1 or −1, fetch the message separately) into a Rust Result. Lines 44–46 fetch the bridge's first error message (the one fail and check stored), copying it out of C memory before it can change. Lines 48–50 refuse a slot number outside 0..3, which would otherwise index past textures in main.rs. It takes the raw handle rather than &self so that open can call it before Video exists, which is exactly what lines 66–67 do.

The closure starting on line 62 is for cleanup, and I'll come back to it. Inside it is the import:

src/video.rs, lines 69–97
69for slot in 0..SLOTS {
  let resource_ptr = unsafe { nv_texture(native.as_ptr(), slot as i32) };
  if resource_ptr.is_null() { return Err(format!("D3D12 resource slot {slot} missing")); }
  // C++ AddRef's, Rust from_raw takes ownership of the additional COM reference.
  let resource = unsafe { ID3D12Resource::from_raw(resource_ptr) };
  let desc = wgpu::TextureDescriptor {
      label: Some("Video Codec SDK CUDA-shared RGBA texture"),
      size: wgpu::Extent3d { width, height, depth_or_array_layers: 1 },
      mip_level_count: 1,
      sample_count: 1,
      dimension: wgpu::TextureDimension::D2,
      format: wgpu::TextureFormat::Rgba8Unorm,
      usage: wgpu::TextureUsages::TEXTURE_BINDING,
      view_formats: &[],
  };
  let wrapped = unsafe {
      wgpu::hal::dx12::Device::texture_from_raw(
          resource,
          desc.format,
          desc.dimension,
          desc.size,
          desc.mip_level_count,
          desc.sample_count,
      )
  };
  let tex = unsafe { device.create_texture_from_hal::<wgpu::hal::api::Dx12>(
      wrapped, &desc, wgpu::TextureUses::RESOURCE) };
  textures.push(tex);
97}

For each of the three slots:

Lines 70–73: take ownership of the resource. nv_texture returns the raw pointer with an extra reference added (Step 8). ID3D12Resource::from_raw turns that pointer into a Rust value that will call Release when dropped, so the reference count is balanced.
Lines 74–83: describe the texture to wgpu. This descriptor has to match the resource the bridge created. The same format (Rgba8Unorm is R8G8B8A8_UNORM), the same size, one mip level, one sample. The usage is TEXTURE_BINDING only, because wgpu will only ever sample these; it never renders into them or copies from them.
Lines 84–93: wrap the resource in a wgpu-hal texture. texture_from_raw doesn't allocate anything. It just pairs the resource pointer with the metadata wgpu-hal needs to treat it as a texture.
Lines 94–95: register it with wgpu. create_texture_from_hal takes the hal texture and the descriptor and returns an ordinary wgpu::Texture. The third argument, TextureUses::RESOURCE, tells wgpu what state the texture is in right now; more on that below.

Every unsafe in this excerpt is a promise the compiler can't check, and wgpu's documentation for create_texture_from_hal spells out what I'm promising:

- `hal_texture` must be created from this device internal handle
- `hal_texture` must be created respecting `desc`
- `hal_texture` must be initialized
- `initial_state`, if it is not `TextureUses::UNINITIALIZED`, must
  match the actual driver-side layout/state of the wrapped resource at
  the moment of wrap.

It's worth going through these one at a time, because each one was a decision somewhere earlier in the code:

The safety contract, point by point
Requirement	How it holds here
Created from this device	Step 7 creates the resources on the `ID3D12Device` that `as_hal` returned on line 55, and Step 4 matched the CUDA device to that adapter's LUID. Same device, same GPU.
Matches the descriptor	Format, size, mip count and sample count on lines 74–83 are copied from how Step 7 created the resource. If either side changed, this is the line to update.
Initialized	Lines 66–67 decode into the first two slots before the import. The third slot is wrapped before its first write. That's harmless here, because wgpu only ever samples a slot the bridge has reported as containing a frame, but it's the one place this code leans on the contract rather than meeting it outright.
Initial state matches	The resources were created in `COMMON`, and wgpu is told `RESOURCE`. That's consistent for simultaneous-access textures: they're “implicitly promoted from D3D12_RESOURCE_STATE_COMMON to the relevant state on first GPU access” and decay back after each submission (Microsoft's resource state docs). That's why the flag in Step 7 mattered.

The end of open is timing setup and the cleanup I promised to explain:

src/video.rs, lines 98–107
98let start = Instant::now();
99    let fallback_interval = Duration::from_millis(33);
100    let due = if pending.pts_100ns > current.pts_100ns {
101        duration_100ns(pending.pts_100ns)
102    } else { fallback_interval };
103    Ok(Self { native, textures, current, pending, start,
104        next_due: start + due, fallback_interval })
105})();
106if result.is_err() { unsafe { nv_close(native.as_ptr()) }; }
107result

Lines 98–104 start the playback clock and work out when the lookahead frame is due. Line 106 is why the body was wrapped in a closure: if anything after nv_open fails (a slot's resource is missing, a decode errors), the closure returns early with the error, and this line still calls nv_close. Without it, a failed open would leak the decoder, the CUDA context and the textures.

Ownership at the end of this step, because it's easy to lose track of:

The bridge owns the decoder, the CUDA objects and one reference to each D3D12 resource.
Each wgpu::Texture owns another reference to its resource, via from_raw.
Video owns the bridge pointer and the three textures. Its Drop closes the bridge:

src/video.rs, lines 139–145
139impl Drop for Video {
140    fn drop(&mut self) {
141        // Caller waits for GPU idle before dropping Video/State. The COM resource
142        // refs owned by wgpu outlive the bridge if other bind groups retain them.
143        unsafe { nv_close(self.native.as_ptr()) };
144    }
145}

The comment on lines 141–142 is the one rule the caller has to follow: the GPU must be idle before Video is dropped, because nv_close tears down the CUDA side while wgpu might still have a frame in flight that samples the texture. main.rs handles that by waiting for the device in its own Drop, which runs before Video's.

And nv_close is just delete, so what it actually runs is the bridge's destructor:

native/video_bridge.cu, lines 122–136
122~Bridge() {
  // Caller must release every wrapped wgpu texture before nv_close!
  if (cuda_context && push()) {
      if (parser) cuvidDestroyVideoParser(parser);
      if (decoder) cuvidDestroyDecoder(decoder);
      for (auto& slot : slots) slot.cleanup();
      pop();
      cuCtxDestroy(cuda_context);
  }
  reader.Reset();
  compressed_type.Reset();
  d3d.Reset();
  if (mf_started) MFShutdown();
  if (com_started) CoUninitialize();
136}

Teardown in the reverse order of setup. The parser and decoder go first, because they own NVDEC surfaces; then each slot's cleanup from Step 7; then the context itself, which has to outlive everything that was created in it, and which is pushed one last time so the destroy calls have a current context. The COM pointers release themselves when reset. Media Foundation and COM are shut down only if this object started them (the two flags from the Bridge struct). The comment on line 123 is the same rule as above, seen from the C++ side.

Step 10: Sample the video in WGSL

From here on it's ordinary wgpu, which is the point. The three textures came from a video decoder, but nothing in the rendering code knows that. Each slot gets its own bind group, sharing the same uniform buffer and sampler:

src/main.rs, lines 166–176
166let bind_groups: Vec<wgpu::BindGroup> = video.textures.iter().enumerate().map(|(_slot, tex)| {
  let texture_view = tex.create_view(&Default::default());
  device.create_bind_group(&wgpu::BindGroupDescriptor {
      label: Some("NVDEC frame + MVP bindings"), layout: &layout,
      entries: &[
          wgpu::BindGroupEntry { binding: 0, resource: uniform.as_entire_binding() },
          wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&texture_view) },
          wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::Sampler(&sampler) },
      ],
  })
176}).collect();

src/main.rs, lines 276–280
276pass.set_pipeline(&self.pipeline);
277pass.set_bind_group(0, &self.bind_groups[self.video.current_slot()], &[]);
278pass.set_vertex_buffer(0, self.vertices.slice(..));
279pass.set_index_buffer(self.indices.slice(..), wgpu::IndexFormat::Uint16);
280pass.draw_indexed(0..36, 0, 0..1);

And the shader has no idea any of this happened:

src/cube.wgsl
1struct Uniforms { mvp: mat4x4<f32> }
2@group(0) @binding(0) var<uniform> uniforms: Uniforms;
3@group(0) @binding(1) var video: texture_2d<f32>;
4@group(0) @binding(2) var video_sampler: sampler;
5
6struct VertexInput {
7    @location(0) pos: vec3<f32>,
8    @location(1) uv: vec2<f32>,
9};
10struct VertexOutput {
11    @builtin(position) clip: vec4<f32>,
12    @location(0) uv: vec2<f32>,
13};
14
15@vertex
16fn vs_main(v: VertexInput) -> VertexOutput {
17    var result: VertexOutput;
18    result.clip = uniforms.mvp * vec4<f32>(v.pos, 1.0);
19    result.uv = v.uv;
20    return result;
21}
22
23@fragment
24fn fs_main(v: VertexOutput) -> @location(0) vec4<f32> {
25    return textureSample(video, video_sampler, v.uv);
26}

GotchaGamma applied twice: the washed-out video

One thing in this step did go wrong, and it's worth telling because it's easy to repeat. My first version picked an sRGB swapchain format, which is the usual advice for wgpu programs, and the video came out pale and washed out:

The same frame and angle before (left, sRGB swapchain) and after the fix (right, Bgra8Unorm). — Big Buck Bunny © Blender Foundation, CC BY 3.0

The cause took a while to see. The kernel in Step 7 writes display-encoded RGB, the gamma-encoded values a video player would put on screen. It writes them into an Rgba8Unorm texture, and wgpu samples that format as-is, treating the values as linear. Then the sRGB swapchain gamma-encodes everything written to it. So the video was gamma-encoded twice, which pushes midtones toward white. The fix is to render to a non-sRGB swapchain format, so the values reach the screen unchanged, the way a video player would show them:

src/main.rs, lines 110–115
110// The CUDA kernel writes display-encoded (gamma) RGB into Rgba8Unorm textures, so
111// sampled values are already what the screen should show. A non-sRGB target passes
112// them through; an sRGB one would gamma-encode them a second time (washed out).
113let format = caps.formats.iter().copied().find(|f| !f.is_srgb()).unwrap_or(caps.formats[0]);
114println!("Swapchain format: {format:?}");
115let clear = background(format);

That has one knock-on effect. wgpu treats clear colors as linear only on sRGB targets, so the window's background color needs its own conversion to look the same on either kind of format:

src/main.rs, lines 42–50
42/// The window background, rgb(69, 80, 93), as clear values for `format`. Clear colors are
43/// linear on sRGB targets, so the display value is decoded there.
44fn background(format: wgpu::TextureFormat) -> wgpu::Color {
45    let [r, g, b] = [69.0, 80.0, 93.0].map(|v: f64| {
46        let c = v / 255.0;
47        if !format.is_srgb() { c } else if c <= 0.04045 { c / 12.92 } else { ((c + 0.055) / 1.055).powf(2.4) }
48    });
49    wgpu::Color { r, g, b, a: 1.0 }
50}

To check the fix was right rather than just different, I compared the cube's pixels against an independent decode: ffmpeg's frame from the same moment of the clip, converted with the same BT.709 limited-range matrix. The fixed version lands within a few levels of ffmpeg:

Average color of the cube's pixels, and of ffmpeg's frame
	Average RGB	Saturation	Brightness
Before (sRGB swapchain)	237, 231, 227	0.08	0.95
After (Bgra8Unorm)	220, 206, 198	0.17	0.90
ffmpeg's decode	217, 205, 197	0.17	0.88

The cube shows the frame in perspective, so this is an approximate check, but the difference is clear.

Step 11: Frame timing and synchronization

Two clocks are running. The cube redraws about every 16 ms, because that's a reasonable refresh rate for an animation. The video has its own rate, 24 frames per second for my clip, which is about every 42 ms. So on most redraws the video shouldn't change, and on some it should. Video::update runs on every redraw and decides:

src/video.rs, lines 112–136
112pub fn update(&mut self, device: &wgpu::Device) -> Result<bool, String> {
  let now = Instant::now();
  if now < self.next_due { return Ok(false); }
  // Safe first milestone. Serializes CUDA writes against *all* queued wgpu reads.
  // A future revision can replace this with CUDA-imported D3D12 fences per slot.
  device.poll(wgpu::PollType::Wait { submission_index: None, timeout: None })
      .map_err(|e| e.to_string())?;
  let previous_pts = self.current.pts_100ns;
  self.current = self.pending;
  if self.current.pts_100ns <= previous_pts {
      // Video loop: the native reader rewound to the beginning.
      self.start = now;
  }
  let new_pending = Self::decode(self.native)?;
  let delta = new_pending.pts_100ns - self.current.pts_100ns;
  if delta > 0 {
      self.fallback_interval = duration_100ns(delta).min(Duration::from_secs(1));
      self.next_due = self.start + duration_100ns(new_pending.pts_100ns);
  } else {
      // Next packet may be from the next video loop, or contain a duplicate timestamp.
      self.next_due = now + self.fallback_interval;
  }
  self.pending = new_pending;
  Ok(true)
136}

This is where the lookahead frame from Step 9 pays off. At any moment Video holds two frames: the one on screen (current) and the next one (pending), whose timestamp it already knows. Line 114 is the whole scheduling decision: if the pending frame isn't due yet, do nothing. When it is due, lines 120–121 promote it to current, and lines 125–134 decode a new pending frame and work out when that one is due. Knowing the next frame's time in advance is what lets this be a simple comparison instead of a guess.

The timestamps make this possible by passing through every layer unchanged: Media Foundation's sample time became the NVDEC packet timestamp in Step 3, came back out with the decoded frame in Step 6, and the bridge made it relative to the first frame. Lines 121–124 handle the loop: a timestamp that goes backwards means the bridge rewound the clip, so the playback clock restarts. Lines 130–133 handle files with duplicate or missing timestamps by falling back to the last known frame interval.

Then there's the part I've been deferring: synchronization, lines 117–118.

Why this wayWhy every frame waits for the GPU

Here is the problem. The bridge decodes into slot (current + 1) % 3. With one slot on screen and one holding the lookahead, that's the slot that was on screen two frames ago. “Was on screen” is from the CPU's point of view; the GPU runs behind the CPU, and a draw call that samples that slot may not have executed yet. If CUDA writes into it now, that draw samples a half-written frame. Nothing in D3D12, CUDA or wgpu prevents this on its own, because none of them knows what the others are doing with the memory.

The conservative answer, and the one this code uses, is to wait. Before every decode, line 117 waits until the GPU has finished everything wgpu submitted (device.poll(Wait)). And in Step 6, the bridge waits for the kernel to finish before returning (cudaDeviceSynchronize). Between the two waits, CUDA never writes a slot the GPU might be reading, and wgpu never samples a slot CUDA is still writing.

That makes it correct, but not concurrent. Decode, conversion and rendering happen strictly one after another, and at 24 frames per second on a 320×180 clip nobody notices. At 4K it would. The three slots are there for the real fix: give each slot a D3D12 fence, import it into CUDA as an external semaphore, and have each side wait only for the slot it needs. Then decoding could move to its own thread, and three slots would become true triple buffering. I stopped short of that because it's a post of its own, and because the version with the waits is the one where every step is easy to follow.

Part 5: Build and run

Step 12. From a clone of the repository (or your own copy built from the steps above). No developer prompt is needed; cc finds the Visual Studio compiler for nvcc on its own. The only thing to set is where the Video Codec SDK is. In PowerShell:

$env:VIDEO_CODEC_SDK_DIR = "C:\path\to\Video_Codec_SDK_13.1.15"
cargo run --release -- "C:\Videos\sample.mp4"

A working start prints one line per stage, which is useful for seeing how far it got when something fails:

wgpu adapter: NVIDIA GeForce RTX 5080 Laptop GPU (Dx12)
[mf] compressed H.264 320x180, sequence headers=35 bytes
[mf] NAL framing: Annex B
Native Video Codec SDK: H.264/HEVC, Media Foundation compressed packets, 3 GPU textures
Swapchain format: Bgra8Unorm

Then a window opens with the cube, playing the clip on every face. My test clip played for about 90 seconds per run without an error. The process peaked at about 330 MB and used about 17 seconds of CPU time over those 90 seconds.

That clip is Constrained Baseline H.264, which has no B-frames, so it never exercised the reordering I described in Step 5. To check that path, I cut a 4-second High-profile clip from it with ffmpeg (58 of its 96 frames are B-frames) and let it loop. It played through about four loops without errors, which also covers the end-of-stream flush and the rewind in Step 6. I didn't check the frame order frame by frame.

The build also prints a few Compiler family detection failed ... "nvcc" "-E" warnings. That's cc probing nvcc to learn what kind of compiler it is, and it doesn't affect the build.

Common errors

Build errors come from build.rs's checks in Step 2. At run time the bridge prints [nvdec] lines to stderr, and the same message ends the program.

Errors and what they mean
Message	What to do
`Set VIDEO_CODEC_SDK_DIR to the extracted NVIDIA Video Codec SDK`	Set the variable in the same terminal, pointing at the extracted folder.
`Could not locate SDK nvcuvid.lib (Lib/win/x64 or Lib/x64)`	The path points at the wrong level (or at the zip). It should be the folder that contains Interface and Lib.
`Install the NVIDIA CUDA Toolkit (CUDA_PATH)`	Install the CUDA Toolkit, then open a new terminal so CUDA_PATH is set.
`This sample requires wgpu to use an NVIDIA GPU (vendor 0x10DE)`	wgpu picked another adapter, such as an iGPU (Step 4). Check which GPU Windows assigns to the program.
`CUDA device LUID doesn't match the wgpu D3D12 NVIDIA adapter`	CUDA can't see the GPU wgpu chose (Step 4). Check the driver and that the CUDA Toolkit supports the GPU.
`Only native compressed H.264/HEVC MP4 is supported`	The video uses another codec (AV1, VP9 and so on). Re-encode it, or extend the subtype check in Step 3.
`Only progressive 8-bit 4:2:0 video is supported (NV12)`	10-bit, 4:2:2 or interlaced video, rejected in Step 5. See the limitations below.
`Codec/profile not supported by NVIDIA hardware`	cuvidGetDecoderCaps says this GPU's NVDEC can't decode the stream's codec or profile.
`Invalid compressed NAL units (...)`	The packets are neither Annex B nor 4-byte length-prefixed, or the framing changed mid-stream (Step 3).

Part 6: Every file, top to bottom

Parts 2 to 4 followed the data: a packet in, a frame out. This part follows the files instead. If you have the project open next to this post, it's the order to read each file in. The pieces the steps already covered get a pointer back to the step; the pieces they skipped (the helpers, the state, the destructor, the cube geometry, the window loop) are explained here, and each section ends with the whole file. None of this is new design. It's the glue that holds the steps together, and it's where most of the small decisions live.

Cargo.toml and build.rs

Both are short enough that Step 1 and Step 2 already show them in full. The one thing to add is how they relate: Cargo.toml names build.rs as the build script (line 5), and lists cc under [build-dependencies] rather than [dependencies], because it's used by the build script and not by the program. Everything else the program links comes from the rustc-link-lib lines in build.rs, not from Cargo.

native/video_bridge.cu

485 lines, and the file where everything native lives. It's laid out as: includes and constants, the CUDA kernel, two helper types, the Bridge object with its methods in the order they're used, and the C exports at the bottom. Here's the top:

native/video_bridge.cu, lines 1–30
1// Windows/x64 only. No FFmpeg: Media Foundation provides *compressed* packets;
2// NVIDIA Video Codec SDK (nvcuvid) parses/decodes; CUDA writes D3D12 textures.
3#define NOMINMAX
4#define WIN32_LEAN_AND_MEAN
5#include <windows.h>
6#include <mfapi.h>
7#include <mfidl.h>
8#include <mfreadwrite.h>
9#include <wrl/client.h>
10#include <d3d12.h>
11#include <cuda.h>
12#include <cuda_runtime.h>
13#include <nvcuvid.h>
14#include <algorithm>
15#include <array>
16#include <cstdint>
17#include <cstdio>
18#include <cstring>
19#include <deque>
20#include <memory>
21#include <new>
22#include <string>
23#include <vector>
24#include "bitstream.h"
25
26using Microsoft::WRL::ComPtr;
27static constexpr int kSlots = 3;
28// The Source Reader's stream selectors are negative enum values passed as DWORD.
29static constexpr DWORD kVideoStream = static_cast<DWORD>(MF_SOURCE_READER_FIRST_VIDEO_STREAM);
30static constexpr DWORD kAllStreams = static_cast<DWORD>(MF_SOURCE_READER_ALL_STREAMS);

Lines 3–4 are Windows header hygiene: NOMINMAX stops windows.h from defining min and max macros that would break std::max later in the file, and WIN32_LEAN_AND_MEAN leaves out the parts of the Windows API this file doesn't use. The includes then come in groups: Media Foundation (lines 6–8), WRL's ComPtr for reference-counted COM pointers (line 9), D3D12 (line 10), and both CUDA APIs. That last point matters for reading the rest of the file: cuda.h is the driver API (functions starting with cu, returning CUresult), which NVCUVID is built on and which owns the context; cuda_runtime.h is the runtime API (functions starting with cuda, returning cudaError_t), which has the external-memory import and launches the kernel. The bridge uses both, which is why it has two error-checking helpers below.

Lines 27–30 are the constants. kSlots is the number of textures. The two DWORD constants wrap Media Foundation's stream selectors, which are defined as negative enum values but passed to functions that take an unsigned DWORD; casting once here keeps nvcc from warning at every use.

Then, in file order, and where each piece is explained:

native/video_bridge.cu, by line
Lines	What	Where
32–54	nv12_to_rgba, the color kernel	Step 7
56–63	utf8_to_utf16	Step 3
64–76	OutputSlot and its cleanup	Step 7
78–120	The Bridge's fields, error helpers, push and pop	The Bridge object
122–136	The destructor	Step 9
138–184	init_media: open the file, keep it compressed, read the headers	Step 3
186–262	init_cuda: match the GPU, create the context and the three shared textures	Step 4 and Step 7
264–345	The parser callbacks, on_sequence and create_parser	Step 5
346–351	init: the three setup calls in order	Below
352–368	rewind_video	Step 6
369–409	read_packet: one Media Foundation sample into the parser	Step 3
411–463	next: feed until a frame is ready, map it, convert it	Step 6
466–485	The C exports	Step 8

The one method not shown in the steps is the glue that runs the setup:

native/video_bridge.cu, lines 346–351
346bool init(const char* filename, ID3D12Device* device) {
347    if (!init_media(filename) || !init_cuda(device) || !push()) return false;
348    bool ok = create_parser();
349    pop();
350    return ok;
351}

The order is forced: the media type gives the frame size that the textures need, the CUDA context has to exist before the parser is created in it, and the parser is created last. The whole file:

native/video_bridge.cu
1// Windows/x64 only. No FFmpeg: Media Foundation provides *compressed* packets;
2// NVIDIA Video Codec SDK (nvcuvid) parses/decodes; CUDA writes D3D12 textures.
3#define NOMINMAX
4#define WIN32_LEAN_AND_MEAN
5#include <windows.h>
6#include <mfapi.h>
7#include <mfidl.h>
8#include <mfreadwrite.h>
9#include <wrl/client.h>
10#include <d3d12.h>
11#include <cuda.h>
12#include <cuda_runtime.h>
13#include <nvcuvid.h>
14#include <algorithm>
15#include <array>
16#include <cstdint>
17#include <cstdio>
18#include <cstring>
19#include <deque>
20#include <memory>
21#include <new>
22#include <string>
23#include <vector>
24#include "bitstream.h"
25
26using Microsoft::WRL::ComPtr;
27static constexpr int kSlots = 3;
28// The Source Reader's stream selectors are negative enum values passed as DWORD.
29static constexpr DWORD kVideoStream = static_cast<DWORD>(MF_SOURCE_READER_FIRST_VIDEO_STREAM);
30static constexpr DWORD kAllStreams = static_cast<DWORD>(MF_SOURCE_READER_ALL_STREAMS);
31
32// Simple SDR BT.709 *limited-range* matrix; HDR/P010 and metadata conversion
33// require a different kernel and are rejected in on_sequence for this example.
34__device__ __forceinline__ unsigned char byte_clamp(float v) {
  return (unsigned char)fminf(255.0f, fmaxf(0.0f, v));
36}
37__global__ void nv12_to_rgba(cudaSurfaceObject_t dst, const uint8_t* nv12,
                           unsigned pitch, unsigned width, unsigned height) {
  const unsigned x = blockIdx.x * blockDim.x + threadIdx.x;
  const unsigned y = blockIdx.y * blockDim.y + threadIdx.y;
  if (x >= width || y >= height) return;
  const uint8_t* y_plane = nv12;
  const uint8_t* uv_plane = nv12 + (size_t)pitch * height;
  const float Y = (float)y_plane[(size_t)y * pitch + x];
  const unsigned uv_index = (y / 2) * pitch + (x / 2) * 2;
  const float U = (float)uv_plane[uv_index] - 128.0f;
  const float V = (float)uv_plane[uv_index + 1] - 128.0f;
  const float c = 1.16438356f * (Y - 16.0f);
  const uchar4 rgba = make_uchar4(
      byte_clamp(c + 1.79274107f * V),
      byte_clamp(c - 0.21324861f * U - 0.53290933f * V),
      byte_clamp(c + 2.11240179f * U), 255);
  surf2Dwrite(rgba, dst, x * sizeof(uchar4), y);
54}
55
56static std::wstring utf8_to_utf16(const char* src) {
  int chars = MultiByteToWideChar(CP_UTF8, MB_ERR_INVALID_CHARS, src, -1, nullptr, 0);
  if (chars <= 0) return {};
  std::wstring out((size_t)chars, 0);
  if (!MultiByteToWideChar(CP_UTF8, MB_ERR_INVALID_CHARS, src, -1, out.data(), chars)) return {};
  out.pop_back();
  return out;
63}
64struct OutputSlot {
  ComPtr<ID3D12Resource> texture;
  cudaExternalMemory_t memory = nullptr;
  cudaMipmappedArray_t mipmap = nullptr;
  cudaSurfaceObject_t surface = 0;
  void cleanup() {
      if (surface) cudaDestroySurfaceObject(surface);
      if (mipmap) cudaFreeMipmappedArray(mipmap);
      if (memory) cudaDestroyExternalMemory(memory);
      surface = 0; mipmap = nullptr; memory = nullptr;
      texture.Reset();
  }
76};
77
78struct Bridge {
  ComPtr<IMFSourceReader> reader;
  ComPtr<IMFMediaType> compressed_type;
  ComPtr<ID3D12Device> d3d;
  bool com_started = false, mf_started = false;
  bool has_eos = false, flushed = false;
  unsigned width = 0, height = 0;
  unsigned decode_width = 0, decode_height = 0;
  cudaVideoCodec codec_type = cudaVideoCodec_H264;
  std::vector<uint8_t> sequence_header;
  std::vector<uint8_t> bitstream;
  NalFraming framing = NalFraming::Unknown; // Set by the first sample, kept across loops.
  std::deque<CUVIDPARSERDISPINFO> decoded;
  CUcontext cuda_context = nullptr;
  CUvideoparser parser = nullptr;
  CUvideodecoder decoder = nullptr;
  std::array<OutputSlot, kSlots> slots;
  std::string error;
  int current_slot = -1;
  int64_t first_timestamp = -1;
  int64_t last_pts = 0;
99
  bool fail(const char* message) {
      if (error.empty()) error = message;
      std::fprintf(stderr, "[nvdec] %s\n", message);
      return false;
  }
  bool check(CUresult result, const char* message) {
      if (result == CUDA_SUCCESS) return true;
      const char* name = "unknown";
      cuGetErrorName(result, &name);
      error = std::string(message) + ": " + (name ? name : "unknown");
      std::fprintf(stderr, "[nvdec] %s\n", error.c_str());
      return false;
  }
  bool check_cuda(cudaError_t result, const char* message) {
      if (result == cudaSuccess) return true;
      error = std::string(message) + ": " + cudaGetErrorString(result);
      std::fprintf(stderr, "[nvdec] %s\n", error.c_str());
      return false;
  }
  bool push() { return check(cuCtxPushCurrent(cuda_context), "cuCtxPushCurrent"); }
  void pop() { CUcontext old = nullptr; cuCtxPopCurrent(&old); }
121
  ~Bridge() {
      // Caller must release every wrapped wgpu texture before nv_close!
      if (cuda_context && push()) {
          if (parser) cuvidDestroyVideoParser(parser);
          if (decoder) cuvidDestroyDecoder(decoder);
          for (auto& slot : slots) slot.cleanup();
          pop();
          cuCtxDestroy(cuda_context);
      }
      reader.Reset();
      compressed_type.Reset();
      d3d.Reset();
      if (mf_started) MFShutdown();
      if (com_started) CoUninitialize();
  }
137
  bool init_media(const char* path) {
      HRESULT hr = CoInitializeEx(nullptr, COINIT_MULTITHREADED);
      // winit may already have created an STA on the Windows UI thread.
      if (hr == RPC_E_CHANGED_MODE) hr = CoInitializeEx(nullptr, COINIT_APARTMENTTHREADED);
      if (FAILED(hr)) return fail("CoInitializeEx failed");
      com_started = true;
      hr = MFStartup(MF_VERSION);
      if (FAILED(hr)) return fail("MFStartup failed");
      mf_started = true;
      const std::wstring wide = utf8_to_utf16(path);
      if (wide.empty()) return fail("Invalid UTF-8 video path");
      hr = MFCreateSourceReaderFromURL(wide.c_str(), nullptr, reader.GetAddressOf());
      if (FAILED(hr)) return fail("MFCreateSourceReaderFromURL failed (MP4 file required)");
      // Disable all other tracks; preserve the *native compressed* video type.
      hr = reader->SetStreamSelection(kAllStreams, FALSE);
      if (FAILED(hr)) return fail("Disable other streams failed");
      hr = reader->SetStreamSelection(kVideoStream, TRUE);
      if (FAILED(hr)) return fail("Select video stream failed");
      hr = reader->GetNativeMediaType(kVideoStream, 0, compressed_type.GetAddressOf());
      if (FAILED(hr)) return fail("Cannot get native compressed video type");
      GUID subtype = {};
      if (FAILED(compressed_type->GetGUID(MF_MT_SUBTYPE, &subtype))) return fail("No video subtype");
      if (subtype == MFVideoFormat_H264) codec_type = cudaVideoCodec_H264;
      else if (subtype == MFVideoFormat_HEVC) codec_type = cudaVideoCodec_HEVC;
      else return fail("Only native compressed H.264/HEVC MP4 is supported");
      hr = reader->SetCurrentMediaType(kVideoStream, nullptr, compressed_type.Get());
      if (FAILED(hr)) return fail("Could not keep native compressed video output");
      if (FAILED(MFGetAttributeSize(compressed_type.Get(), MF_MT_FRAME_SIZE, &width, &height)) ||
          !width || !height || (width & 1) || (height & 1)) {
          return fail("Missing/unsupported video frame size (even dimensions required)");
      }
      UINT32 header_size = 0;
      if (SUCCEEDED(compressed_type->GetBlobSize(MF_MT_MPEG_SEQUENCE_HEADER, &header_size)) && header_size) {
          sequence_header.resize(header_size);
          UINT32 got = 0;
          if (FAILED(compressed_type->GetBlob(MF_MT_MPEG_SEQUENCE_HEADER,
                  sequence_header.data(), header_size, &got))) return fail("Sequence-header read failed");
          sequence_header.resize(got);
          // MF sequence headers are Annex B SPS/PPS or VPS/SPS/PPS, NOT avcC/hvcC.
          if (!annexb_prefix(sequence_header.data(), sequence_header.size()))
              return fail("Expected Annex B MF_MT_MPEG_SEQUENCE_HEADER");
      }
      std::printf("[mf] compressed %s %ux%u, sequence headers=%zu bytes\n",
          codec_type == cudaVideoCodec_H264 ? "H.264" : "HEVC", width, height, sequence_header.size());
      std::fflush(stdout); // C stdio is fully buffered when redirected; Rust's println! isn't affected
      return true;
  }
185
  bool init_cuda(ID3D12Device* dev) {
      if (!dev) return fail("Missing wgpu D3D12 device");
      d3d = dev;
      if (!check(cuInit(0), "cuInit")) return false;
      int count = 0;
      if (!check(cuDeviceGetCount(&count), "cuDeviceGetCount")) return false;
      const LUID adapter_luid = dev->GetAdapterLuid();
      CUdevice match = -1;
      for (int i = 0; i < count; ++i) {
          CUdevice candidate = 0;
          char cuda_luid[8] = {};
          unsigned node_mask = 0;
          if (cuDeviceGet(&candidate, i) == CUDA_SUCCESS &&
              cuDeviceGetLuid(cuda_luid, &node_mask, candidate) == CUDA_SUCCESS &&
              std::memcmp(cuda_luid, &adapter_luid, sizeof(cuda_luid)) == 0) {
              match = candidate;
              break;
          }
      }
      if (match < 0) return fail("CUDA device LUID doesn't match the wgpu D3D12 NVIDIA adapter");
      if (!check(cuCtxCreate(
          &cuda_context,
          nullptr,
          0,
          match
      ), "cuCtxCreate")) return false;
      // cuCtxCreate makes this context current; pop so every call explicitly pushes it.
      pop();
      if (!push()) return false;
      bool ok = true;
      for (auto& slot : slots) {
          D3D12_HEAP_PROPERTIES heap = {};
          heap.Type = D3D12_HEAP_TYPE_DEFAULT;
          D3D12_RESOURCE_DESC desc = {};
          desc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D;
          desc.Width = width;
          desc.Height = height;
          desc.DepthOrArraySize = 1;
          desc.MipLevels = 1;
          desc.Format = DXGI_FORMAT_R8G8B8A8_UNORM;
          desc.SampleDesc.Count = 1;
          desc.Layout = D3D12_TEXTURE_LAYOUT_UNKNOWN;
          desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_SIMULTANEOUS_ACCESS;
          HRESULT hr = dev->CreateCommittedResource(&heap, D3D12_HEAP_FLAG_SHARED, &desc,
              D3D12_RESOURCE_STATE_COMMON, nullptr, IID_PPV_ARGS(&slot.texture));
          if (FAILED(hr)) { ok = fail("CreateCommittedResource failed"); break; }
          HANDLE shared = nullptr;
          hr = dev->CreateSharedHandle(slot.texture.Get(), nullptr, GENERIC_ALL, nullptr, &shared);
          if (FAILED(hr)) { ok = fail("CreateSharedHandle failed"); break; }
          cudaExternalMemoryHandleDesc memory_desc = {};
          memory_desc.type = cudaExternalMemoryHandleTypeD3D12Resource;
          memory_desc.handle.win32.handle = shared;
          memory_desc.size = dev->GetResourceAllocationInfo(0, 1, &desc).SizeInBytes;
          memory_desc.flags = cudaExternalMemoryDedicated;
          const cudaError_t imported = cudaImportExternalMemory(&slot.memory, &memory_desc);
          CloseHandle(shared);
          if (!check_cuda(imported, "cudaImportExternalMemory")) { ok = false; break; }
          cudaExternalMemoryMipmappedArrayDesc mapping = {};
          mapping.offset = 0;
          mapping.formatDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
          mapping.extent = make_cudaExtent(width, height, 1);
          mapping.flags = cudaArraySurfaceLoadStore;
          mapping.numLevels = 1;
          if (!check_cuda(cudaExternalMemoryGetMappedMipmappedArray(&slot.mipmap, slot.memory, &mapping),
                          "cudaExternalMemoryGetMappedMipmappedArray")) { ok = false; break; }
          cudaArray_t array = nullptr;
          if (!check_cuda(cudaGetMipmappedArrayLevel(&array, slot.mipmap, 0),
                          "cudaGetMipmappedArrayLevel")) { ok = false; break; }
          cudaResourceDesc resource = {};
          resource.resType = cudaResourceTypeArray;
          resource.res.array.array = array;
          if (!check_cuda(cudaCreateSurfaceObject(&slot.surface, &resource),
                          "cudaCreateSurfaceObject")) { ok = false; break; }
      }
      pop();
      return ok;
  }
263
  static int CUDAAPI sequence_cb(void* user, CUVIDEOFORMAT* fmt) {
      return static_cast<Bridge*>(user)->on_sequence(fmt);
  }
  static int CUDAAPI decode_cb(void* user, CUVIDPICPARAMS* pic) {
      auto* b = static_cast<Bridge*>(user);
      return b->decoder && b->check(cuvidDecodePicture(b->decoder, pic), "cuvidDecodePicture") ? 1 : 0;
  }
  static int CUDAAPI display_cb(void* user, CUVIDPARSERDISPINFO* info) {
      auto* b = static_cast<Bridge*>(user);
      if (!info) return 1;
      if (b->decoded.size() >= 64) { b->fail("Decoder display queue overflow"); return 0; }
      b->decoded.push_back(*info); // Parser delivers display/reordered order.
      return 1;
  }
  int on_sequence(CUVIDEOFORMAT* f) {
      if (!f || f->codec != codec_type || f->chroma_format != cudaVideoChromaFormat_420 ||
          f->bit_depth_luma_minus8 != 0 || !f->progressive_sequence) {
          fail("Only progressive 8-bit 4:2:0 video is supported (NV12)"); return 0;
      }
      if (decoder) {
          if (decode_width != f->coded_width || decode_height != f->coded_height) {
              fail("Midstream resolution changes are not supported"); return 0;
          }
          return std::max(
              2u,
              static_cast<unsigned int>(f->min_num_decode_surfaces)
          );
      }
      CUVIDDECODECAPS caps = {};
      caps.eCodecType = f->codec;
      caps.eChromaFormat = f->chroma_format;
      caps.nBitDepthMinus8 = f->bit_depth_luma_minus8;
      if (!check(cuvidGetDecoderCaps(&caps), "cuvidGetDecoderCaps") || !caps.bIsSupported) {
          fail("Codec/profile not supported by NVIDIA hardware"); return 0;
      }
      decode_width = f->coded_width;
      decode_height = f->coded_height;
      CUVIDDECODECREATEINFO create = {};
      create.CodecType = f->codec;
      create.ChromaFormat = f->chroma_format;
      create.OutputFormat = cudaVideoSurfaceFormat_NV12;
      create.DeinterlaceMode = cudaVideoDeinterlaceMode_Weave;
      create.ulCreationFlags = cudaVideoCreate_PreferCUVID;
      create.ulNumDecodeSurfaces = std::max(
          2u,
          static_cast<unsigned int>(f->min_num_decode_surfaces)
      );
      create.ulNumOutputSurfaces = 2;
      create.ulWidth = f->coded_width;
      create.ulHeight = f->coded_height;
      create.ulMaxWidth = f->coded_width;
      create.ulMaxHeight = f->coded_height;
      create.display_area.left = f->display_area.left;
      create.display_area.top = f->display_area.top;
      create.display_area.right = f->display_area.right;
      create.display_area.bottom = f->display_area.bottom;
      create.ulTargetWidth = width;
      create.ulTargetHeight = height;
      if (!check(cuvidCreateDecoder(&decoder, &create), "cuvidCreateDecoder")) return 0;
      return (int)create.ulNumDecodeSurfaces;
  }
  bool create_parser() {
      if (parser) { if (!check(cuvidDestroyVideoParser(parser), "cuvidDestroyVideoParser")) return false; parser = nullptr; }
      if (decoder) { if (!check(cuvidDestroyDecoder(decoder), "cuvidDestroyDecoder")) return false; decoder = nullptr; }
      decode_width = 0; decode_height = 0;
      CUVIDPARSERPARAMS parameters = {};
      parameters.CodecType = codec_type;
      parameters.ulMaxNumDecodeSurfaces = 20; // Sequence callback can override.
      parameters.ulMaxDisplayDelay = 2; // Frames the parser may hold before the display callback (0 = none).
      parameters.pUserData = this;
      parameters.pfnSequenceCallback = &Bridge::sequence_cb;
      parameters.pfnDecodePicture = &Bridge::decode_cb;
      parameters.pfnDisplayPicture = &Bridge::display_cb;
      if (!check(cuvidCreateVideoParser(&parser, &parameters), "cuvidCreateVideoParser")) return false;
      if (!sequence_header.empty()) {
          CUVIDSOURCEDATAPACKET pkt = {};
          pkt.payload = sequence_header.data();
          pkt.payload_size = (unsigned long)sequence_header.size();
          if (!check(cuvidParseVideoData(parser, &pkt), "cuvidParseVideoData(sequence header)")) return false;
      }
      return true;
  }
  bool init(const char* filename, ID3D12Device* device) {
      if (!init_media(filename) || !init_cuda(device) || !push()) return false;
      bool ok = create_parser();
      pop();
      return ok;
  }
  bool rewind_video() {
      // Do not destroy the parser until *every* display callback frame was consumed.
      PROPVARIANT seek; PropVariantInit(&seek);
      seek.vt = VT_I8;
      seek.hVal.QuadPart = 0;
      HRESULT hr = reader->SetCurrentPosition(GUID_NULL, seek);
      if (FAILED(hr)) return fail("Media Foundation failed to seek to video beginning");
      has_eos = false;
      flushed = false;
      first_timestamp = -1;
      last_pts = 0;
      decoded.clear();
      if (!push()) return false;
      const bool ok = create_parser();
      pop();
      return ok;
  }
  bool read_packet() {
      DWORD flags = 0;
      LONGLONG timestamp = 0;
      ComPtr<IMFSample> sample;
      const HRESULT hr = reader->ReadSample(kVideoStream, 0,
          nullptr, &flags, &timestamp, sample.GetAddressOf());
      if (FAILED(hr)) return fail("IMFSourceReader::ReadSample failed");
      if (flags & (MF_SOURCE_READERF_NATIVEMEDIATYPECHANGED | MF_SOURCE_READERF_CURRENTMEDIATYPECHANGED))
          return fail("Dynamic compressed media-type changes unsupported");
      if (flags & MF_SOURCE_READERF_ENDOFSTREAM) { has_eos = true; return true; }
      if (!sample) return true; // stream tick / no sample
      ComPtr<IMFMediaBuffer> buf;
      if (FAILED(sample->ConvertToContiguousBuffer(buf.GetAddressOf()))) return fail("Cannot obtain compressed sample");
      BYTE* bytes = nullptr;
      DWORD count = 0;
      if (FAILED(buf->Lock(&bytes, nullptr, &count))) return fail("Compressed packet Lock failed");
      // MF's MPEG-4 source typically presents Annex B, but handle MP4 length-prefixed packets too.
      // The first sample decides which; every later sample must match it.
      const NalFraming before = framing;
      const bool converted = annexb_sample(bytes, count, bitstream, framing);
      buf->Unlock();
      if (!converted) {
          return fail(before == NalFraming::AnnexB ? "Invalid compressed NAL units (stream started as Annex B)"
              : before == NalFraming::LengthPrefixed ? "Invalid compressed NAL units (stream started with four-byte lengths)"
              : "Invalid compressed NAL units (expected Annex B or four-byte lengths)");
      }
      if (before == NalFraming::Unknown) {
          std::printf("[mf] NAL framing: %s\n", framing == NalFraming::AnnexB ? "Annex B" : "four-byte lengths");
          std::fflush(stdout);
      }
      if (bitstream.size() > UINT32_MAX) return fail("Compressed packet too large");
      CUVIDSOURCEDATAPACKET pkt = {};
      pkt.flags = CUVID_PKT_TIMESTAMP;
      pkt.timestamp = timestamp; // MF sample time and CUVID timestamps both use 100ns ticks.
      pkt.payload = bitstream.data();
      pkt.payload_size = (unsigned long)bitstream.size();
      if (!push()) return false;
      const bool ok = check(cuvidParseVideoData(parser, &pkt), "cuvidParseVideoData(packet)");
      pop();
      return ok;
  }
  // Return 1 on success, -1 on error. out_pts: presentation time in 100ns units.
  int next(int* out_slot, int64_t* out_pts) {
      if (!out_slot || !out_pts) { fail("Null frame result"); return -1; }
      // Read compressed packets only until a display-order frame becomes available.
      // Consume that frame before submitting additional packets: protects NVDEC output surfaces.
      for (unsigned steps = 0; decoded.empty(); ++steps) {
          if (steps > 100000) { fail("No decodable video frames"); return -1; }
          if (!has_eos) {
              if (!read_packet()) return -1;
          } else if (!flushed) {
              CUVIDSOURCEDATAPACKET eos = {};
              eos.flags = CUVID_PKT_ENDOFSTREAM;
              if (!push()) return -1;
              bool ok = check(cuvidParseVideoData(parser, &eos), "cuvidParseVideoData(EOS)");
              pop();
              if (!ok) return -1;
              flushed = true;
          } else {
              if (!rewind_video()) return -1;
          }
      }
      CUVIDPARSERDISPINFO display = decoded.front();
      decoded.pop_front();
      const int slot_index = (current_slot + 1) % kSlots;
      if (!push()) return -1;
      CUVIDPROCPARAMS process = {};
      process.progressive_frame = display.progressive_frame;
      process.top_field_first = display.top_field_first;
      process.second_field = 0;
      process.unpaired_field = display.repeat_first_field < 0;
      CUdeviceptr device_frame = 0;
      unsigned pitch = 0;
      bool ok = check(cuvidMapVideoFrame(decoder, display.picture_index,
          &device_frame, &pitch, &process), "cuvidMapVideoFrame");
      if (ok) {
          const dim3 threads(16, 16);
          const dim3 blocks((width + 15) / 16, (height + 15) / 16);
          nv12_to_rgba<<<blocks, threads>>>(slots[slot_index].surface,
              reinterpret_cast<const uint8_t*>(device_frame), pitch, width, height);
          ok = check_cuda(cudaGetLastError(), "CUDA color-convert launch") &&
               check_cuda(cudaDeviceSynchronize(), "CUDA color-convert synchronize");
          // Unmap only once the CUDA kernel has finished accessing NVDEC memory.
          const bool unmapped = check(cuvidUnmapVideoFrame(decoder, device_frame), "cuvidUnmapVideoFrame");
          ok = unmapped && ok;
      }
      pop();
      if (!ok) return -1;
      current_slot = slot_index;
      if (first_timestamp < 0) first_timestamp = display.timestamp;
      last_pts = std::max<int64_t>(0, display.timestamp - first_timestamp);
      *out_slot = slot_index;
      *out_pts = last_pts;
      return 1;
  }
464};
465
466extern "C" {
467__declspec(dllexport) Bridge* nv_open(const char* path, void* d3d_device) {
  std::unique_ptr<Bridge> bridge(new (std::nothrow) Bridge());
  if (!bridge || !bridge->init(path, reinterpret_cast<ID3D12Device*>(d3d_device))) return nullptr;
  return bridge.release();
471}
472__declspec(dllexport) int nv_width(Bridge* b) { return b ? (int)b->width : 0; }
473__declspec(dllexport) int nv_height(Bridge* b) { return b ? (int)b->height : 0; }
474__declspec(dllexport) void* nv_texture(Bridge* b, int index) {
  if (!b || index < 0 || index >= kSlots) return nullptr;
  ID3D12Resource* texture = b->slots[index].texture.Get();
  if (texture) texture->AddRef(); // Rust from_raw takes this additional reference.
  return texture;
479}
480__declspec(dllexport) int nv_next(Bridge* b, int* slot, int64_t* pts) {
  return b ? b->next(slot, pts) : -1;
482}
483__declspec(dllexport) const char* nv_error(Bridge* b) { return b ? b->error.c_str() : "No video bridge"; }
484__declspec(dllexport) void nv_close(Bridge* b) { delete b; }
485}

native/bitstream.h

A header-only file with no Windows or CUDA in it, on purpose: that's what lets it be tested anywhere. Four functions build on each other. The first two recognize start codes:

native/bitstream.h, lines 1–21
1#pragma once
2#include <cstddef>
3#include <cstdint>
4#include <vector>
5
6// How a stream's samples separate NAL units: Annex B start codes, or the MP4
7// length prefixes (avcC/hvcC style). One track uses one framing throughout.
8enum class NalFraming { Unknown, AnnexB, LengthPrefixed };
9
10// Bytes taken by a start code at p (00 00 01, or more zeros then 01), 0 if none.
11inline size_t start_code_length(const uint8_t* p, size_t n) {
12    size_t zeros = 0;
13    while (zeros < n && p[zeros] == 0) ++zeros;
14    return zeros >= 2 && zeros < n && p[zeros] == 1 ? zeros + 1 : 0;
15}
16
17// A start code followed by at least one byte of NAL unit.
18inline bool annexb_prefix(const uint8_t* p, size_t n) {
19    const size_t prefix = start_code_length(p, n);
20    return prefix && prefix < n;
21}

start_code_length counts leading zero bytes and expects a 01 after at least two of them. Accepting two or more zeros rather than exactly two or three matters: the H.264 byte-stream format allows extra zero_bytes before a start code, and some muxers emit them. annexb_prefix adds the requirement that something follows the start code, so a sample that is only a start code is rejected rather than passed to the decoder as an empty NAL unit.

Lines 23–48 are length_prefixed and detect_framing, which Step 3 showed. The last function is the one the bridge calls:

native/bitstream.h, lines 50–79
50// Normalizes one sample to Annex B in out. With framing Unknown, it's detected
51// from this sample and stored, so the caller keeps it for the rest of the
52// stream; after that, every sample must parse in that framing. Malformed input
53// is rejected rather than guessed at.
54inline bool annexb_sample(const uint8_t* p, size_t n, std::vector<uint8_t>& out,
                        NalFraming& framing, unsigned nal_len_bytes = 4) {
  out.clear();
  if (!n) return false;
  const NalFraming kind = framing == NalFraming::Unknown ? detect_framing(p, n, nal_len_bytes) : framing;
  if (kind == NalFraming::AnnexB) {
      if (!annexb_prefix(p, n)) return false;
      out.assign(p, p + n);
  } else if (kind == NalFraming::LengthPrefixed) {
      if (!length_prefixed(p, n, nal_len_bytes)) return false;
      while (n) {
          size_t count = 0;
          for (unsigned i = 0; i < nal_len_bytes; ++i) count = (count << 8) | p[i];
          p += nal_len_bytes;
          n -= nal_len_bytes;
          out.insert(out.end(), {0, 0, 0, 1});
          out.insert(out.end(), p, p + count);
          p += count;
          n -= count;
      }
  } else {
      return false;
  }
  framing = kind;
  return true;
79}

framing is an in-out parameter. If the caller passes Unknown, line 58 detects the framing from this sample; otherwise the caller's value is used, and the sample has to match it (lines 60 and 63) or the function fails. Annex B input is copied through unchanged (line 61). Length-prefixed input is rebuilt NAL unit by NAL unit, with each length prefix replaced by a 4-byte start code (lines 64–73). Only on success does line 77 write the detected framing back, so a rejected first sample doesn't lock the stream into a bad guess. The whole file:

native/bitstream.h
1#pragma once
2#include <cstddef>
3#include <cstdint>
4#include <vector>
5
6// How a stream's samples separate NAL units: Annex B start codes, or the MP4
7// length prefixes (avcC/hvcC style). One track uses one framing throughout.
8enum class NalFraming { Unknown, AnnexB, LengthPrefixed };
9
10// Bytes taken by a start code at p (00 00 01, or more zeros then 01), 0 if none.
11inline size_t start_code_length(const uint8_t* p, size_t n) {
12    size_t zeros = 0;
13    while (zeros < n && p[zeros] == 0) ++zeros;
14    return zeros >= 2 && zeros < n && p[zeros] == 1 ? zeros + 1 : 0;
15}
16
17// A start code followed by at least one byte of NAL unit.
18inline bool annexb_prefix(const uint8_t* p, size_t n) {
19    const size_t prefix = start_code_length(p, n);
20    return prefix && prefix < n;
21}
22
23// True when nal_len_bytes-byte lengths tile the whole sample exactly, every
24// NAL unit is non-empty, and every NAL header has its forbidden_zero_bit (the
25// top bit, in both H.264 and HEVC) clear.
26inline bool length_prefixed(const uint8_t* p, size_t n, unsigned nal_len_bytes) {
27    if (!n || !(nal_len_bytes == 1 || nal_len_bytes == 2 || nal_len_bytes == 4)) return false;
28    while (n) {
29        if (n < nal_len_bytes) return false;
30        size_t count = 0;
31        for (unsigned i = 0; i < nal_len_bytes; ++i) count = (count << 8) | p[i];
32        p += nal_len_bytes;
33        n -= nal_len_bytes;
34        if (count == 0 || count > n || (p[0] & 0x80)) return false;
35        p += count;
36        n -= count;
37    }
38    return true;
39}
40
41// Which framing one sample uses. The length check wins: a 4-byte length of 1
42// or 256-511 starts with 00 00 00 01 / 00 00 01 xx, which looks exactly like a
43// start code, while real Annex B data practically never tiles as lengths.
44inline NalFraming detect_framing(const uint8_t* p, size_t n, unsigned nal_len_bytes = 4) {
45    if (length_prefixed(p, n, nal_len_bytes)) return NalFraming::LengthPrefixed;
46    if (annexb_prefix(p, n)) return NalFraming::AnnexB;
47    return NalFraming::Unknown;
48}
49
50// Normalizes one sample to Annex B in out. With framing Unknown, it's detected
51// from this sample and stored, so the caller keeps it for the rest of the
52// stream; after that, every sample must parse in that framing. Malformed input
53// is rejected rather than guessed at.
54inline bool annexb_sample(const uint8_t* p, size_t n, std::vector<uint8_t>& out,
55                          NalFraming& framing, unsigned nal_len_bytes = 4) {
56    out.clear();
57    if (!n) return false;
58    const NalFraming kind = framing == NalFraming::Unknown ? detect_framing(p, n, nal_len_bytes) : framing;
59    if (kind == NalFraming::AnnexB) {
60        if (!annexb_prefix(p, n)) return false;
61        out.assign(p, p + n);
62    } else if (kind == NalFraming::LengthPrefixed) {
63        if (!length_prefixed(p, n, nal_len_bytes)) return false;
64        while (n) {
65            size_t count = 0;
66            for (unsigned i = 0; i < nal_len_bytes; ++i) count = (count << 8) | p[i];
67            p += nal_len_bytes;
68            n -= nal_len_bytes;
69            out.insert(out.end(), {0, 0, 0, 1});
70            out.insert(out.end(), p, p + count);
71            p += count;
72            n -= count;
73        }
74    } else {
75        return false;
76    }
77    framing = kind;
78    return true;
79}

native/tests/bitstream_test.cpp

A plain C++ program with asserts. Its helpers build test input without hand-typing hundreds of bytes:

native/tests/bitstream_test.cpp, lines 1–31
1#include "../bitstream.h"
2#include <cassert>
3#include <cstdio>
4
5using Bytes = std::vector<uint8_t>;
6
7// Converts one sample as the first of a new stream.
8static bool convert(const Bytes& in, Bytes& out, unsigned nal_len_bytes = 4) {
9    NalFraming framing = NalFraming::Unknown;
10    return annexb_sample(in.data(), in.size(), out, framing, nal_len_bytes);
11}
12
13// A 4-byte length prefix followed by a NAL unit of `size` bytes (header 0x41).
14static Bytes length_nal(size_t size) {
15    Bytes nal = {uint8_t(size >> 24), uint8_t(size >> 16), uint8_t(size >> 8), uint8_t(size), 0x41};
16    for (size_t i = 1; i < size; ++i) nal.push_back(uint8_t(0x10 + i % 0x60));
17    return nal;
18}
19
20static Bytes annexb_of(const Bytes& lengths_sample) {
21    Bytes out;
22    size_t i = 0;
23    while (i < lengths_sample.size()) {
24        const size_t count = (size_t(lengths_sample[i]) << 24) | (size_t(lengths_sample[i + 1]) << 16) |
25            (size_t(lengths_sample[i + 2]) << 8) | lengths_sample[i + 3];
26        out.insert(out.end(), {0, 0, 0, 1});
27        out.insert(out.end(), lengths_sample.begin() + i + 4, lengths_sample.begin() + i + 4 + count);
28        i += 4 + count;
29    }
30    return out;
31}

convert runs one sample as if it were the first of a stream. length_nal builds a length-prefixed NAL unit of any size; the header byte 0x41 is a valid H.264 NAL header (forbidden bit clear, type 1, a non-IDR slice), which matters because length_prefixed checks that bit. annexb_of is an independent, deliberately simple conversion to compare the real one against.

native/tests/bitstream_test.cpp, lines 33–78
33int main() {
  Bytes out;
35
  // Annex B passes through unchanged, with 4- and 3-byte start codes.
  const Bytes annex = {0,0,0,1,0x67,0x64,0,0,0,1,0x68};
  assert(convert(annex, out) && out == annex);
  const Bytes annex3 = {0,0,1,0x67,0x64,0,0,1,0x68};
  assert(convert(annex3, out) && out == annex3);
41
  // MP4 length prefixes of 4, 2 and 1 bytes become start codes.
  const Bytes expected = {0,0,0,1,0x67,0x64,0,0,0,1,0x68};
  assert(convert({0,0,0,2,0x67,0x64,0,0,0,1,0x68}, out) && out == expected);
  assert(convert({0,2,0x67,0x64,0,1,0x68}, out, 2) && out == expected);
  assert(convert({2,0x67,0x64,1,0x68}, out, 1) && out == expected);
47
  // The old misread: 4-byte lengths of 256-511 start 00 00 01 xx and a length
  // of 1 starts 00 00 00 01, both of which look like start codes.
  Bytes len300 = length_nal(300);
  const Bytes tail = length_nal(5);
  len300.insert(len300.end(), tail.begin(), tail.end());
  assert(len300[0] == 0 && len300[1] == 0 && len300[2] == 1 && annexb_prefix(len300.data(), len300.size()));
  assert(convert(len300, out) && out == annexb_of(len300));
  Bytes len1 = {0,0,0,1,0x09};
  const Bytes slice = length_nal(40);
  len1.insert(len1.end(), slice.begin(), slice.end());
  assert(convert(len1, out) && out == annexb_of(len1));
59
  // The first sample fixes the framing for the stream.
  NalFraming framing = NalFraming::Unknown;
  assert(annexb_sample(annex.data(), annex.size(), out, framing) && framing == NalFraming::AnnexB);
  assert(annexb_sample(len300.data(), len300.size(), out, framing) && out == len300); // kept as Annex B
  framing = NalFraming::Unknown;
  assert(annexb_sample(len1.data(), len1.size(), out, framing) && framing == NalFraming::LengthPrefixed);
  assert(!annexb_sample(annex.data(), annex.size(), out, framing)); // Annex B can't follow lengths
  assert(framing == NalFraming::LengthPrefixed);
68
  // Malformed input is rejected.
  assert(!convert({0,0,0,20,0x68}, out));           // length past the end
  assert(!convert({}, out));                         // empty
  assert(!convert({0,0,0,0}, out));                  // zero length, no start code
  assert(!convert({0,0,1}, out));                    // start code with no NAL unit
  assert(!convert({0,0,0,2,0xE7,0x64,0x22}, out));   // tiles, but forbidden_zero_bit set, and no start code
  assert(!convert({0,0,0,2,0x67,0x64}, out, 3));     // unsupported prefix size
76
  std::puts("bitstream normalization tests passed");
78}

The cases follow the comments: Annex B passes through with 4- and 3-byte start codes; length prefixes of 4, 2 and 1 bytes convert; then the regression cases for the misread from Step 3. Line 53 is the important assertion there: it checks that the 300-byte sample really does start with what looks like a start code, so the test is exercising the ambiguity and not just a normal conversion. Lines 60–67 test the framing lock in both directions, and the last block covers malformed input. The whole file:

native/tests/bitstream_test.cpp
1#include "../bitstream.h"
2#include <cassert>
3#include <cstdio>
4
5using Bytes = std::vector<uint8_t>;
6
7// Converts one sample as the first of a new stream.
8static bool convert(const Bytes& in, Bytes& out, unsigned nal_len_bytes = 4) {
9    NalFraming framing = NalFraming::Unknown;
10    return annexb_sample(in.data(), in.size(), out, framing, nal_len_bytes);
11}
12
13// A 4-byte length prefix followed by a NAL unit of `size` bytes (header 0x41).
14static Bytes length_nal(size_t size) {
15    Bytes nal = {uint8_t(size >> 24), uint8_t(size >> 16), uint8_t(size >> 8), uint8_t(size), 0x41};
16    for (size_t i = 1; i < size; ++i) nal.push_back(uint8_t(0x10 + i % 0x60));
17    return nal;
18}
19
20static Bytes annexb_of(const Bytes& lengths_sample) {
21    Bytes out;
22    size_t i = 0;
23    while (i < lengths_sample.size()) {
24        const size_t count = (size_t(lengths_sample[i]) << 24) | (size_t(lengths_sample[i + 1]) << 16) |
25            (size_t(lengths_sample[i + 2]) << 8) | lengths_sample[i + 3];
26        out.insert(out.end(), {0, 0, 0, 1});
27        out.insert(out.end(), lengths_sample.begin() + i + 4, lengths_sample.begin() + i + 4 + count);
28        i += 4 + count;
29    }
30    return out;
31}
32
33int main() {
34    Bytes out;
35
36    // Annex B passes through unchanged, with 4- and 3-byte start codes.
37    const Bytes annex = {0,0,0,1,0x67,0x64,0,0,0,1,0x68};
38    assert(convert(annex, out) && out == annex);
39    const Bytes annex3 = {0,0,1,0x67,0x64,0,0,1,0x68};
40    assert(convert(annex3, out) && out == annex3);
41
42    // MP4 length prefixes of 4, 2 and 1 bytes become start codes.
43    const Bytes expected = {0,0,0,1,0x67,0x64,0,0,0,1,0x68};
44    assert(convert({0,0,0,2,0x67,0x64,0,0,0,1,0x68}, out) && out == expected);
45    assert(convert({0,2,0x67,0x64,0,1,0x68}, out, 2) && out == expected);
46    assert(convert({2,0x67,0x64,1,0x68}, out, 1) && out == expected);
47
48    // The old misread: 4-byte lengths of 256-511 start 00 00 01 xx and a length
49    // of 1 starts 00 00 00 01, both of which look like start codes.
50    Bytes len300 = length_nal(300);
51    const Bytes tail = length_nal(5);
52    len300.insert(len300.end(), tail.begin(), tail.end());
53    assert(len300[0] == 0 && len300[1] == 0 && len300[2] == 1 && annexb_prefix(len300.data(), len300.size()));
54    assert(convert(len300, out) && out == annexb_of(len300));
55    Bytes len1 = {0,0,0,1,0x09};
56    const Bytes slice = length_nal(40);
57    len1.insert(len1.end(), slice.begin(), slice.end());
58    assert(convert(len1, out) && out == annexb_of(len1));
59
60    // The first sample fixes the framing for the stream.
61    NalFraming framing = NalFraming::Unknown;
62    assert(annexb_sample(annex.data(), annex.size(), out, framing) && framing == NalFraming::AnnexB);
63    assert(annexb_sample(len300.data(), len300.size(), out, framing) && out == len300); // kept as Annex B
64    framing = NalFraming::Unknown;
65    assert(annexb_sample(len1.data(), len1.size(), out, framing) && framing == NalFraming::LengthPrefixed);
66    assert(!annexb_sample(annex.data(), annex.size(), out, framing)); // Annex B can't follow lengths
67    assert(framing == NalFraming::LengthPrefixed);
68
69    // Malformed input is rejected.
70    assert(!convert({0,0,0,20,0x68}, out));           // length past the end
71    assert(!convert({}, out));                         // empty
72    assert(!convert({0,0,0,0}, out));                  // zero length, no start code
73    assert(!convert({0,0,1}, out));                    // start code with no NAL unit
74    assert(!convert({0,0,0,2,0xE7,0x64,0x22}, out));   // tiles, but forbidden_zero_bit set, and no start code
75    assert(!convert({0,0,0,2,0x67,0x64}, out, 3));     // unsupported prefix size
76
77    std::puts("bitstream normalization tests passed");
78}

src/video.rs

145 lines: the Rust side of the bridge. It knows about D3D12 resources and wgpu, and nothing about Media Foundation or CUDA.

src/video.rs, lines 1–8
1use std::{
2    ffi::{c_char, c_void, CStr, CString},
3    ptr::NonNull,
4    time::{Duration, Instant},
5};
6use windows::{core::Interface, Win32::Graphics::Direct3D12::ID3D12Resource};
7
8const SLOTS: usize = 3;

The ffi imports are for the C boundary: CString to pass the path in, CStr to read error messages out, and the raw pointer types. NonNull is a pointer that can't be null, which is how Video holds the bridge once nv_open has succeeded. From the windows crate it needs just ID3D12Resource and the Interface trait that provides from_raw. SLOTS on line 8 has to agree with kSlots in the bridge; nothing enforces that either.

The rest of the file, in order:

src/video.rs, by line
Lines	What	Where
10–20	The extern "C" declarations	Step 8
22–37	DecodedFrame, Video and duration_100ns	Step 9
40–52	decode: nv_next as a Result	Step 9
54–108	open: the bridge, the first two frames and the wgpu-hal import	Step 9
110	current_slot	Step 10
112–136	update: frame timing and the GPU wait	Step 11
139–145	Drop	Step 9

The whole file:

src/video.rs
1use std::{
  ffi::{c_char, c_void, CStr, CString},
  ptr::NonNull,
  time::{Duration, Instant},
5};
6use windows::{core::Interface, Win32::Graphics::Direct3D12::ID3D12Resource};
7
8const SLOTS: usize = 3;
9
10#[repr(C)]
11struct VideoHandle { _opaque: [u8; 0] }
12unsafe extern "C" {
  fn nv_open(filename: *const c_char, device: *mut c_void) -> *mut VideoHandle;
  fn nv_width(ptr: *mut VideoHandle) -> i32;
  fn nv_height(ptr: *mut VideoHandle) -> i32;
  fn nv_texture(ptr: *mut VideoHandle, slot: i32) -> *mut c_void;
  fn nv_next(ptr: *mut VideoHandle, slot: *mut i32, pts: *mut i64) -> i32;
  fn nv_error(ptr: *mut VideoHandle) -> *const c_char;
  fn nv_close(ptr: *mut VideoHandle);
20}
21
22#[derive(Clone, Copy)]
23struct DecodedFrame { slot: usize, pts_100ns: i64 }
24
25pub struct Video {
  native: NonNull<VideoHandle>,
  pub textures: Vec<wgpu::Texture>,
  current: DecodedFrame,
  pending: DecodedFrame,
  start: Instant,
  next_due: Instant,
  fallback_interval: Duration,
33}
34
35fn duration_100ns(ticks: i64) -> Duration {
  Duration::from_nanos((ticks.max(0) as u64).saturating_mul(100))
37}
38
39impl Video {
  fn decode(handle: NonNull<VideoHandle>) -> Result<DecodedFrame, String> {
      let mut slot = -1;
      let mut pts = -1_i64;
      if unsafe { nv_next(handle.as_ptr(), &mut slot, &mut pts) } != 1 {
          let ptr = unsafe { nv_error(handle.as_ptr()) };
          return Err(if ptr.is_null() { "Native NVDEC decoder error".into() }
              else { unsafe { CStr::from_ptr(ptr) }.to_string_lossy().into_owned() });
      }
      if !(0..SLOTS as i32).contains(&slot) {
          return Err(format!("Native NVDEC returned invalid texture slot {slot}"));
      }
      Ok(DecodedFrame { slot: slot as usize, pts_100ns: pts })
  }
53
  pub fn open(device: &wgpu::Device, filename: &str) -> Result<Self, String> {
      let hal = unsafe { device.as_hal::<wgpu::hal::api::Dx12>() }
          .ok_or("wgpu must use the D3D12 backend")?;
      let file = CString::new(filename).map_err(|e| e.to_string())?;
      let native = NonNull::new(unsafe { nv_open(file.as_ptr(), hal.raw_device().as_raw()) })
          .ok_or("Media Foundation / Video Codec SDK initialization failed (see stderr)")?;
      // Decode initial + lookahead frame before importing the D3D12 resources into wgpu.
      // On failure, always release the native decoder and CUDA resources.
      let result = (|| {
          let width = unsafe { nv_width(native.as_ptr()) } as u32;
          let height = unsafe { nv_height(native.as_ptr()) } as u32;
          if width == 0 || height == 0 { return Err("Invalid video dimensions".to_owned()); }
          let current = Self::decode(native)?;
          let pending = Self::decode(native)?;
          let mut textures = Vec::with_capacity(SLOTS);
          for slot in 0..SLOTS {
              let resource_ptr = unsafe { nv_texture(native.as_ptr(), slot as i32) };
              if resource_ptr.is_null() { return Err(format!("D3D12 resource slot {slot} missing")); }
              // C++ AddRef's, Rust from_raw takes ownership of the additional COM reference.
              let resource = unsafe { ID3D12Resource::from_raw(resource_ptr) };
              let desc = wgpu::TextureDescriptor {
                  label: Some("Video Codec SDK CUDA-shared RGBA texture"),
                  size: wgpu::Extent3d { width, height, depth_or_array_layers: 1 },
                  mip_level_count: 1,
                  sample_count: 1,
                  dimension: wgpu::TextureDimension::D2,
                  format: wgpu::TextureFormat::Rgba8Unorm,
                  usage: wgpu::TextureUsages::TEXTURE_BINDING,
                  view_formats: &[],
              };
              let wrapped = unsafe {
                  wgpu::hal::dx12::Device::texture_from_raw(
                      resource,
                      desc.format,
                      desc.dimension,
                      desc.size,
                      desc.mip_level_count,
                      desc.sample_count,
                  )
              };
              let tex = unsafe { device.create_texture_from_hal::<wgpu::hal::api::Dx12>(
                  wrapped, &desc, wgpu::TextureUses::RESOURCE) };
              textures.push(tex);
          }
          let start = Instant::now();
          let fallback_interval = Duration::from_millis(33);
          let due = if pending.pts_100ns > current.pts_100ns {
              duration_100ns(pending.pts_100ns)
          } else { fallback_interval };
          Ok(Self { native, textures, current, pending, start,
              next_due: start + due, fallback_interval })
      })();
      if result.is_err() { unsafe { nv_close(native.as_ptr()) }; }
      result
  }
109
  pub fn current_slot(&self) -> usize { self.current.slot }
111
  pub fn update(&mut self, device: &wgpu::Device) -> Result<bool, String> {
      let now = Instant::now();
      if now < self.next_due { return Ok(false); }
      // Safe first milestone. Serializes CUDA writes against *all* queued wgpu reads.
      // A future revision can replace this with CUDA-imported D3D12 fences per slot.
      device.poll(wgpu::PollType::Wait { submission_index: None, timeout: None })
          .map_err(|e| e.to_string())?;
      let previous_pts = self.current.pts_100ns;
      self.current = self.pending;
      if self.current.pts_100ns <= previous_pts {
          // Video loop: the native reader rewound to the beginning.
          self.start = now;
      }
      let new_pending = Self::decode(self.native)?;
      let delta = new_pending.pts_100ns - self.current.pts_100ns;
      if delta > 0 {
          self.fallback_interval = duration_100ns(delta).min(Duration::from_secs(1));
          self.next_due = self.start + duration_100ns(new_pending.pts_100ns);
      } else {
          // Next packet may be from the next video loop, or contain a duplicate timestamp.
          self.next_due = now + self.fallback_interval;
      }
      self.pending = new_pending;
      Ok(true)
  }
137}
138
139impl Drop for Video {
  fn drop(&mut self) {
      // Caller waits for GPU idle before dropping Video/State. The COM resource
      // refs owned by wgpu outlive the bridge if other bind groups retain them.
      unsafe { nv_close(self.native.as_ptr()) };
  }
145}

src/main.rs

341 lines, and the most ordinary file in the project: a winit window, a wgpu device and a textured cube. The video appears in exactly three places: Video::open in setup, the bind groups, and video.update plus current_slot in draw. Everything else would be the same in any wgpu program.

src/main.rs, lines 1–13
1#![cfg(target_os = "windows")]
2mod video;
3
4use std::{sync::Arc, time::{Duration, Instant}};
5use bytemuck::{Pod, Zeroable};
6use glam::{Mat4, Vec3};
7use wgpu::util::DeviceExt;
8use winit::{
9    application::ApplicationHandler,
10    event::WindowEvent,
11    event_loop::{ActiveEventLoop, ControlFlow, EventLoop},
12    window::{Window, WindowAttributes, WindowId},
13};

Line 1 makes the whole program Windows-only at the compiler level, matching the check in build.rs. Line 2 pulls in video.rs as a module. The rest are the four crates: bytemuck for turning structs into bytes, glam for matrices, wgpu's DeviceExt for create_buffer_init, and winit's 0.30 application API.

src/main.rs, lines 15–40
15#[repr(C)]
16#[derive(Clone, Copy, Pod, Zeroable)]
17struct Vertex { pos: [f32; 3], uv: [f32; 2] }
18#[repr(C)]
19#[derive(Clone, Copy, Pod, Zeroable)]
20struct Uniform { mvp: [[f32; 4]; 4] }
21
22fn cube() -> (Vec<Vertex>, Vec<u16>) {
23    let faces = [
24        [[-1., -1.,  1.], [ 1., -1.,  1.], [ 1.,  1.,  1.], [-1.,  1.,  1.]],
25        [[ 1., -1., -1.], [-1., -1., -1.], [-1.,  1., -1.], [ 1.,  1., -1.]],
26        [[-1.,  1.,  1.], [ 1.,  1.,  1.], [ 1.,  1., -1.], [-1.,  1., -1.]],
27        [[-1., -1., -1.], [ 1., -1., -1.], [ 1., -1.,  1.], [-1., -1.,  1.]],
28        [[ 1., -1.,  1.], [ 1., -1., -1.], [ 1.,  1., -1.], [ 1.,  1.,  1.]],
29        [[-1., -1., -1.], [-1., -1.,  1.], [-1.,  1.,  1.], [-1.,  1., -1.]],
30    ];
31    let uv = [[0., 1.], [1., 1.], [1., 0.], [0., 0.]];
32    let mut vertices = Vec::with_capacity(24);
33    let mut indices = Vec::with_capacity(36);
34    for face in faces {
35        let start = vertices.len() as u16;
36        for i in 0..4 { vertices.push(Vertex { pos: face[i], uv: uv[i] }); }
37        indices.extend_from_slice(&[start, start + 1, start + 2, start, start + 2, start + 3]);
38    }
39    (vertices, indices)
40}

Vertex and Uniform are the two structs that cross into the GPU. #[repr(C)] fixes their memory layout so it matches what the shader expects, and Pod and Zeroable are bytemuck's promise that the struct is plain bytes with no padding or pointers, which is what lets cast_slice and bytes_of hand them to wgpu without copying.

cube builds the geometry. Each of the six faces (lines 24–29) is four corners, listed counter-clockwise as seen from outside. Every face gets the same texture coordinates (line 31): v = 1 at the bottom edge and 0 at the top, because row 0 of the texture is the top of the picture, so the video comes out upright on each face. Lines 34–38 turn each face into four vertices and two triangles (six indices), for 24 vertices and 36 indices in total. The faces share no vertices, because each corner needs a different texture coordinate on each face it belongs to.

Lines 42–50 are background from Step 10. Then the depth buffer:

src/main.rs, lines 52–63
52fn depth_texture(device: &wgpu::Device, size: winit::dpi::PhysicalSize<u32>) -> wgpu::TextureView {
  device.create_texture(&wgpu::TextureDescriptor {
      label: Some("Cube depth"),
      size: wgpu::Extent3d { width: size.width.max(1), height: size.height.max(1), depth_or_array_layers: 1 },
      mip_level_count: 1,
      sample_count: 1,
      dimension: wgpu::TextureDimension::D2,
      format: wgpu::TextureFormat::Depth24Plus,
      usage: wgpu::TextureUsages::RENDER_ATTACHMENT,
      view_formats: &[],
  }).create_view(&Default::default())
63}

The pipeline doesn't cull back faces (line 201 below), so without a depth buffer the far faces of the cube could draw over the near ones depending on order. A Depth24Plus texture the size of the window fixes that, and because it's tied to the window size, resize recreates it. The max(1) guards against a zero-sized texture while the window is minimized.

src/main.rs, lines 65–80
65struct State {
  window: Arc<Window>,
  surface: wgpu::Surface<'static>,
  device: wgpu::Device,
  queue: wgpu::Queue,
  config: wgpu::SurfaceConfiguration,
  pipeline: wgpu::RenderPipeline,
  depth: wgpu::TextureView,
  vertices: wgpu::Buffer,
  indices: wgpu::Buffer,
  uniform: wgpu::Buffer,
  bind_groups: Vec<wgpu::BindGroup>,
  clear: wgpu::Color,
  start: Instant,
  video: video::Video,
80}

State is everything that exists once the window is up. Field order matters at the end: Rust drops fields in declaration order after Drop::drop runs, so video, declared last, is dropped last, after the GPU wait in State's Drop below.

State::new starts on line 83. Lines 84–105 pick the adapter and create the device (Step 4), line 107 opens the video (Step 9; it needs only the device, which is why it can come before the swapchain), and lines 109–115 choose the swapchain format (Step 10). The configuration itself:

src/main.rs, lines 116–129
116let size = window.inner_size();
117let config = wgpu::SurfaceConfiguration {
118    usage: wgpu::TextureUsages::RENDER_ATTACHMENT,
119    format,
120    width: size.width.max(1),
121    height: size.height.max(1),
122    present_mode: wgpu::PresentMode::Fifo,
123    alpha_mode: caps.alpha_modes[0],
124    view_formats: vec![],
125    desired_maximum_frame_latency: 2,
126    color_space: wgpu::SurfaceColorSpace::Auto,
127};
128surface.configure(&device, &config);
129let depth = depth_texture(&device, size);

Fifo is vsync: frames are presented in order at the display's refresh rate, which is the right choice for video. A maximum frame latency of 2 lets the CPU run one frame ahead of the GPU without letting it run away. alpha_modes[0] is whatever the surface supports first, since the window is opaque.

src/main.rs, lines 130–141
130let (v, i) = cube();
131let vertices = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
132    label: Some("Cube VB"), contents: bytemuck::cast_slice(&v), usage: wgpu::BufferUsages::VERTEX,
133});
134let indices = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
135    label: Some("Cube IB"), contents: bytemuck::cast_slice(&i), usage: wgpu::BufferUsages::INDEX,
136});
137let uniform = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
138    label: Some("MVP uniform"),
139    contents: bytemuck::bytes_of(&Uniform { mvp: Mat4::IDENTITY.to_cols_array_2d() }),
140    usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
141});

Three GPU buffers, created with their initial contents. The vertex and index buffers never change. The uniform buffer holds the model-view-projection matrix and is written every frame, so it needs COPY_DST; it starts as the identity matrix and draw overwrites it before the first use.

src/main.rs, lines 142–165
142let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
  label: Some("Video + MVP"),
  entries: &[
      wgpu::BindGroupLayoutEntry {
          binding: 0, visibility: wgpu::ShaderStages::VERTEX,
          ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
          count: None,
      },
      wgpu::BindGroupLayoutEntry {
          binding: 1, visibility: wgpu::ShaderStages::FRAGMENT,
          ty: wgpu::BindingType::Texture { sample_type: wgpu::TextureSampleType::Float { filterable: true }, view_dimension: wgpu::TextureViewDimension::D2, multisampled: false },
          count: None,
      },
      wgpu::BindGroupLayoutEntry {
          binding: 2, visibility: wgpu::ShaderStages::FRAGMENT,
          ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), count: None,
      },
  ],
160});
161let sampler = device.create_sampler(&wgpu::SamplerDescriptor {
  mag_filter: wgpu::FilterMode::Linear,
  min_filter: wgpu::FilterMode::Linear,
  ..Default::default()
165});

The bind group layout is the contract between this file and the shader: binding 0 is a uniform buffer the vertex stage reads, binding 1 is a filterable 2D float texture the fragment stage samples, and binding 2 is the sampler. The numbers match the @binding attributes in cube.wgsl; the filterable and Filtering flags pair with the linear sampler on lines 161–165, which smooths the 320×180 video when a face is drawn larger than that.

Lines 166–176 build one bind group per slot (Step 10). Lines 177–180 compile the shader, which is embedded in the executable with include_str!. Then the pipeline:

src/main.rs, lines 177–212
177let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
  label: Some("Cube WGSL"),
  source: wgpu::ShaderSource::Wgsl(include_str!("cube.wgsl").into()),
180});
181let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
  label: Some("Video cube pipeline layout"), bind_group_layouts: &[Some(&layout)], immediate_size: 0,
183});
184let pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
  label: Some("Video cube pipeline"),
  layout: Some(&pipeline_layout),
  vertex: wgpu::VertexState {
      module: &shader, entry_point: Some("vs_main"),
      buffers: &[Some(wgpu::VertexBufferLayout {
          array_stride: std::mem::size_of::<Vertex>() as u64,
          step_mode: wgpu::VertexStepMode::Vertex,
          attributes: &wgpu::vertex_attr_array![0 => Float32x3, 1 => Float32x2],
      })],
      compilation_options: Default::default(),
  },
  fragment: Some(wgpu::FragmentState {
      module: &shader, entry_point: Some("fs_main"),
      targets: &[Some(wgpu::ColorTargetState { format, blend: None, write_mask: wgpu::ColorWrites::ALL })],
      compilation_options: Default::default(),
  }),
  primitive: wgpu::PrimitiveState { cull_mode: None, ..Default::default() },
  depth_stencil: Some(wgpu::DepthStencilState {
      format: wgpu::TextureFormat::Depth24Plus,
      depth_write_enabled: Some(true),
      depth_compare: Some(wgpu::CompareFunction::Less),
      stencil: Default::default(),
      bias: Default::default(),
  }),
  multisample: Default::default(),
  multiview_mask: None,
  cache: None,
212});

Most of this is wgpu boilerplate, and three lines carry the decisions. Line 192 declares the vertex layout: a Float32x3 at location 0 and a Float32x2 at location 1, which is exactly the Vertex struct and exactly the @locations in the shader. Line 198 renders to the swapchain format chosen in Step 10, with no blending. Line 201 turns face culling off, so the cube draws the same whichever way its faces wind, and lines 202–208 enable the depth test (Less: a fragment wins if it's nearer than what's already there).

src/main.rs, lines 217–223
217fn resize(&mut self, s: winit::dpi::PhysicalSize<u32>) {
218    if s.width == 0 || s.height == 0 { return; }
219    self.config.width = s.width;
220    self.config.height = s.height;
221    self.surface.configure(&self.device, &self.config);
222    self.depth = depth_texture(&self.device, s);
223}

On resize, reconfigure the surface to the new size and rebuild the depth buffer to match. A zero size (minimized) is ignored rather than configured, which wgpu would reject.

src/main.rs, lines 225–234
225fn draw(&mut self) -> Result<(), Box<dyn std::error::Error>> {
  // Presentation timestamps determine *when* to advance the video texture.
  // The cube itself renders at the display refresh cadence.
  self.video.update(&self.device)?;
  let elapsed = self.start.elapsed().as_secs_f32();
  let model = Mat4::from_rotation_y(elapsed * 0.8) * Mat4::from_rotation_x(elapsed * 0.43);
  let view = Mat4::look_at_rh(Vec3::new(3.2, 2.1, 4.6), Vec3::ZERO, Vec3::Y);
  let projection = Mat4::perspective_rh(45f32.to_radians(), self.config.width as f32 / self.config.height as f32, 0.1, 100.0);
  let uniform = Uniform { mvp: (projection * view * model).to_cols_array_2d() };
  self.queue.write_buffer(&self.uniform, 0, bytemuck::bytes_of(&uniform));

draw runs once per redraw. Line 228 is the video's one hook: update decides whether the frame changes (Step 11), and a decode error ends the program through the ?. Lines 229–233 build the matrix: the model spins about two axes at different rates so no face stays hidden, the camera sits at (3.2, 2.1, 4.6) looking at the origin, and the projection is a 45° perspective at the window's aspect ratio. Line 234 uploads it.

src/main.rs, lines 235–256
235let frame = match self.surface.get_current_texture() {
236    wgpu::CurrentSurfaceTexture::Success(frame)
237    | wgpu::CurrentSurfaceTexture::Suboptimal(frame) => frame,
238
239    wgpu::CurrentSurfaceTexture::Timeout
240    | wgpu::CurrentSurfaceTexture::Occluded => {
241        return Ok(());
242    }
243
244    wgpu::CurrentSurfaceTexture::Outdated => {
245        self.surface.configure(&self.device, &self.config);
246        return Ok(());
247    }
248
249    wgpu::CurrentSurfaceTexture::Lost => {
250        panic!("Surface lost; surface recreation required");
251    }
252
253    wgpu::CurrentSurfaceTexture::Validation => {
254        panic!("Surface validation error");
255    }
256};

Getting the next swapchain image can fail in ways that aren't errors. Timeout and Occluded (the window is hidden) just skip this frame. Outdated (the window changed under us) reconfigures and skips. The two cases that panic are the ones a sample program can't recover from; a real application would recreate the surface.

src/main.rs, lines 257–286
257let output = frame.texture.create_view(&Default::default());
  let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("Cube render") });
  {
      let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
          label: Some("Cube main pass"),
          color_attachments: &[Some(wgpu::RenderPassColorAttachment {
              view: &output, resolve_target: None,
              ops: wgpu::Operations { load: wgpu::LoadOp::Clear(self.clear), store: wgpu::StoreOp::Store },
              depth_slice: None,
          })],
          depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
              view: &self.depth,
              depth_ops: Some(wgpu::Operations { load: wgpu::LoadOp::Clear(1.0), store: wgpu::StoreOp::Store }),
              stencil_ops: None,
          }),
          occlusion_query_set: None,
          timestamp_writes: None,
          multiview_mask: None,
      });
      pass.set_pipeline(&self.pipeline);
      pass.set_bind_group(0, &self.bind_groups[self.video.current_slot()], &[]);
      pass.set_vertex_buffer(0, self.vertices.slice(..));
      pass.set_index_buffer(self.indices.slice(..), wgpu::IndexFormat::Uint16);
      pass.draw_indexed(0..36, 0, 0..1);
  }
  self.queue.submit(Some(encoder.finish()));
  self.window.pre_present_notify();
  self.queue.present(frame);
  Ok(())
286}

One render pass: clear the color target to the background from Step 10 and the depth buffer to the far plane, bind the pipeline, bind the current slot's bind group (line 277, the other video hook), set the buffers, and draw the 36 indices. Line 282 submits the commands and line 284 presents the frame; pre_present_notify between them is a winit courtesy that lets the compositor time the frame.

src/main.rs, lines 289–293
289impl Drop for State {
290    fn drop(&mut self) {
291        let _ = self.device.poll(wgpu::PollType::Wait { submission_index: None, timeout: None });
292    }
293}

The GPU-idle wait that Step 9 promised. It runs before any field is dropped, so by the time video goes away and closes the bridge, no draw call is still sampling the textures.

src/main.rs, lines 295–333
295struct App {
  path: String,
  state: Option<State>,
  next_frame: Instant,
299}
300impl ApplicationHandler for App {
  fn resumed(&mut self, event_loop: &ActiveEventLoop) {
      if self.state.is_some() { return; }
      let window = Arc::new(event_loop.create_window(WindowAttributes::default()
          .with_title("wgpu cube — native NVDEC / CUDA / D3D12")
          .with_inner_size(winit::dpi::LogicalSize::new(1024, 768))).unwrap());
      let state = pollster::block_on(State::new(window, &self.path)).expect("Video cube initialization failed");
      self.next_frame = Instant::now();
      self.state = Some(state);
  }
  fn window_event(&mut self, event_loop: &ActiveEventLoop, _id: WindowId, event: WindowEvent) {
      let Some(state) = self.state.as_mut() else { return; };
      match event {
          WindowEvent::CloseRequested => event_loop.exit(),
          WindowEvent::Resized(s) => state.resize(s),
          WindowEvent::RedrawRequested if state.config.width > 0 && state.config.height > 0 => {
              if let Err(e) = state.draw() {
                  eprintln!("Render/decode error: {e}");
                  event_loop.exit();
              }
          }
          _ => {}
      }
  }
  fn about_to_wait(&mut self, event_loop: &ActiveEventLoop) {
      if let Some(state) = &self.state {
          if Instant::now() >= self.next_frame {
              state.window.request_redraw();
              self.next_frame = Instant::now() + Duration::from_millis(16);
          }
          event_loop.set_control_flow(ControlFlow::WaitUntil(self.next_frame));
      }
  }
333}

winit 0.30 drives the program through this handler. resumed (lines 301–309) creates the window and the State; State::new is async because wgpu's adapter and device requests are, and pollster::block_on simply waits for it, since there's nothing else to do yet. The guard on line 302 keeps a second resumed from building everything twice. window_event (lines 310–323) handles the three events that matter: close, resize, and redraw, where a failed draw prints the error and exits. about_to_wait (lines 324–332) is the frame pacing: when 16 ms have passed, request a redraw and schedule the next one, and tell winit to sleep until then rather than spin.

src/main.rs, lines 335–341
335fn main() -> Result<(), Box<dyn std::error::Error>> {
336    let path = std::env::args().nth(1).ok_or("Usage: cargo run --release -- <video.mp4>")?;
337    let event_loop = EventLoop::new()?;
338    let mut app = App { path, state: None, next_frame: Instant::now() };
339    event_loop.run_app(&mut app)?;
340    Ok(())
341}

main takes the video path from the command line and hands the handler to the event loop, which doesn't return until the window closes. The whole file:

src/main.rs
1#![cfg(target_os = "windows")]
2mod video;
3
4use std::{sync::Arc, time::{Duration, Instant}};
5use bytemuck::{Pod, Zeroable};
6use glam::{Mat4, Vec3};
7use wgpu::util::DeviceExt;
8use winit::{
  application::ApplicationHandler,
  event::WindowEvent,
  event_loop::{ActiveEventLoop, ControlFlow, EventLoop},
  window::{Window, WindowAttributes, WindowId},
13};
14
15#[repr(C)]
16#[derive(Clone, Copy, Pod, Zeroable)]
17struct Vertex { pos: [f32; 3], uv: [f32; 2] }
18#[repr(C)]
19#[derive(Clone, Copy, Pod, Zeroable)]
20struct Uniform { mvp: [[f32; 4]; 4] }
21
22fn cube() -> (Vec<Vertex>, Vec<u16>) {
  let faces = [
      [[-1., -1.,  1.], [ 1., -1.,  1.], [ 1.,  1.,  1.], [-1.,  1.,  1.]],
      [[ 1., -1., -1.], [-1., -1., -1.], [-1.,  1., -1.], [ 1.,  1., -1.]],
      [[-1.,  1.,  1.], [ 1.,  1.,  1.], [ 1.,  1., -1.], [-1.,  1., -1.]],
      [[-1., -1., -1.], [ 1., -1., -1.], [ 1., -1.,  1.], [-1., -1.,  1.]],
      [[ 1., -1.,  1.], [ 1., -1., -1.], [ 1.,  1., -1.], [ 1.,  1.,  1.]],
      [[-1., -1., -1.], [-1., -1.,  1.], [-1.,  1.,  1.], [-1.,  1., -1.]],
  ];
  let uv = [[0., 1.], [1., 1.], [1., 0.], [0., 0.]];
  let mut vertices = Vec::with_capacity(24);
  let mut indices = Vec::with_capacity(36);
  for face in faces {
      let start = vertices.len() as u16;
      for i in 0..4 { vertices.push(Vertex { pos: face[i], uv: uv[i] }); }
      indices.extend_from_slice(&[start, start + 1, start + 2, start, start + 2, start + 3]);
  }
  (vertices, indices)
40}
41
42/// The window background, rgb(69, 80, 93), as clear values for `format`. Clear colors are
43/// linear on sRGB targets, so the display value is decoded there.
44fn background(format: wgpu::TextureFormat) -> wgpu::Color {
  let [r, g, b] = [69.0, 80.0, 93.0].map(|v: f64| {
      let c = v / 255.0;
      if !format.is_srgb() { c } else if c <= 0.04045 { c / 12.92 } else { ((c + 0.055) / 1.055).powf(2.4) }
  });
  wgpu::Color { r, g, b, a: 1.0 }
50}
51
52fn depth_texture(device: &wgpu::Device, size: winit::dpi::PhysicalSize<u32>) -> wgpu::TextureView {
  device.create_texture(&wgpu::TextureDescriptor {
      label: Some("Cube depth"),
      size: wgpu::Extent3d { width: size.width.max(1), height: size.height.max(1), depth_or_array_layers: 1 },
      mip_level_count: 1,
      sample_count: 1,
      dimension: wgpu::TextureDimension::D2,
      format: wgpu::TextureFormat::Depth24Plus,
      usage: wgpu::TextureUsages::RENDER_ATTACHMENT,
      view_formats: &[],
  }).create_view(&Default::default())
63}
64
65struct State {
  window: Arc<Window>,
  surface: wgpu::Surface<'static>,
  device: wgpu::Device,
  queue: wgpu::Queue,
  config: wgpu::SurfaceConfiguration,
  pipeline: wgpu::RenderPipeline,
  depth: wgpu::TextureView,
  vertices: wgpu::Buffer,
  indices: wgpu::Buffer,
  uniform: wgpu::Buffer,
  bind_groups: Vec<wgpu::BindGroup>,
  clear: wgpu::Color,
  start: Instant,
  video: video::Video,
80}
81
82impl State {
  async fn new(window: Arc<Window>, path: &str) -> Result<Self, Box<dyn std::error::Error>> {
      let mut desc = wgpu::InstanceDescriptor::new_without_display_handle();
      desc.backends = wgpu::Backends::DX12;
86
      let instance = wgpu::Instance::new(desc);
      let surface = instance.create_surface(window.clone())?;
      let adapter = instance.request_adapter(&wgpu::RequestAdapterOptions {
          power_preference: wgpu::PowerPreference::HighPerformance,
          compatible_surface: Some(&surface),
          force_fallback_adapter: false,
          apply_limit_buckets: false,
      }).await?;
      let adapter_info = adapter.get_info();
      println!("wgpu adapter: {} ({:?})", adapter_info.name, adapter_info.backend);
      if adapter_info.vendor != 0x10DE {
          return Err("This sample requires wgpu to use an NVIDIA GPU (vendor 0x10DE)".into());
      }
      let (device, queue) = adapter.request_device(&wgpu::DeviceDescriptor {
          label: Some("Video cube device"),
          required_features: wgpu::Features::empty(),
          required_limits: wgpu::Limits::default(),
          ..Default::default()
      }).await?;
106
      let video = video::Video::open(&device, path)?;
      println!("Native Video Codec SDK: H.264/HEVC, Media Foundation compressed packets, 3 GPU textures");
      let caps = surface.get_capabilities(&adapter);
      // The CUDA kernel writes display-encoded (gamma) RGB into Rgba8Unorm textures, so
      // sampled values are already what the screen should show. A non-sRGB target passes
      // them through; an sRGB one would gamma-encode them a second time (washed out).
      let format = caps.formats.iter().copied().find(|f| !f.is_srgb()).unwrap_or(caps.formats[0]);
      println!("Swapchain format: {format:?}");
      let clear = background(format);
      let size = window.inner_size();
      let config = wgpu::SurfaceConfiguration {
          usage: wgpu::TextureUsages::RENDER_ATTACHMENT,
          format,
          width: size.width.max(1),
          height: size.height.max(1),
          present_mode: wgpu::PresentMode::Fifo,
          alpha_mode: caps.alpha_modes[0],
          view_formats: vec![],
          desired_maximum_frame_latency: 2,
          color_space: wgpu::SurfaceColorSpace::Auto,
      };
      surface.configure(&device, &config);
      let depth = depth_texture(&device, size);
      let (v, i) = cube();
      let vertices = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
          label: Some("Cube VB"), contents: bytemuck::cast_slice(&v), usage: wgpu::BufferUsages::VERTEX,
      });
      let indices = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
          label: Some("Cube IB"), contents: bytemuck::cast_slice(&i), usage: wgpu::BufferUsages::INDEX,
      });
      let uniform = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
          label: Some("MVP uniform"),
          contents: bytemuck::bytes_of(&Uniform { mvp: Mat4::IDENTITY.to_cols_array_2d() }),
          usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
      });
      let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
          label: Some("Video + MVP"),
          entries: &[
              wgpu::BindGroupLayoutEntry {
                  binding: 0, visibility: wgpu::ShaderStages::VERTEX,
                  ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
                  count: None,
              },
              wgpu::BindGroupLayoutEntry {
                  binding: 1, visibility: wgpu::ShaderStages::FRAGMENT,
                  ty: wgpu::BindingType::Texture { sample_type: wgpu::TextureSampleType::Float { filterable: true }, view_dimension: wgpu::TextureViewDimension::D2, multisampled: false },
                  count: None,
              },
              wgpu::BindGroupLayoutEntry {
                  binding: 2, visibility: wgpu::ShaderStages::FRAGMENT,
                  ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), count: None,
              },
          ],
      });
      let sampler = device.create_sampler(&wgpu::SamplerDescriptor {
          mag_filter: wgpu::FilterMode::Linear,
          min_filter: wgpu::FilterMode::Linear,
          ..Default::default()
      });
      let bind_groups: Vec<wgpu::BindGroup> = video.textures.iter().enumerate().map(|(_slot, tex)| {
          let texture_view = tex.create_view(&Default::default());
          device.create_bind_group(&wgpu::BindGroupDescriptor {
              label: Some("NVDEC frame + MVP bindings"), layout: &layout,
              entries: &[
                  wgpu::BindGroupEntry { binding: 0, resource: uniform.as_entire_binding() },
                  wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&texture_view) },
                  wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::Sampler(&sampler) },
              ],
          })
      }).collect();
      let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
          label: Some("Cube WGSL"),
          source: wgpu::ShaderSource::Wgsl(include_str!("cube.wgsl").into()),
      });
      let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
          label: Some("Video cube pipeline layout"), bind_group_layouts: &[Some(&layout)], immediate_size: 0,
      });
      let pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
          label: Some("Video cube pipeline"),
          layout: Some(&pipeline_layout),
          vertex: wgpu::VertexState {
              module: &shader, entry_point: Some("vs_main"),
              buffers: &[Some(wgpu::VertexBufferLayout {
                  array_stride: std::mem::size_of::<Vertex>() as u64,
                  step_mode: wgpu::VertexStepMode::Vertex,
                  attributes: &wgpu::vertex_attr_array![0 => Float32x3, 1 => Float32x2],
              })],
              compilation_options: Default::default(),
          },
          fragment: Some(wgpu::FragmentState {
              module: &shader, entry_point: Some("fs_main"),
              targets: &[Some(wgpu::ColorTargetState { format, blend: None, write_mask: wgpu::ColorWrites::ALL })],
              compilation_options: Default::default(),
          }),
          primitive: wgpu::PrimitiveState { cull_mode: None, ..Default::default() },
          depth_stencil: Some(wgpu::DepthStencilState {
              format: wgpu::TextureFormat::Depth24Plus,
              depth_write_enabled: Some(true),
              depth_compare: Some(wgpu::CompareFunction::Less),
              stencil: Default::default(),
              bias: Default::default(),
          }),
          multisample: Default::default(),
          multiview_mask: None,
          cache: None,
      });
      Ok(Self { window, surface, device, queue, config, pipeline, depth, vertices, indices,
          uniform, bind_groups, clear, start: Instant::now(), video })
  }
216
  fn resize(&mut self, s: winit::dpi::PhysicalSize<u32>) {
      if s.width == 0 || s.height == 0 { return; }
      self.config.width = s.width;
      self.config.height = s.height;
      self.surface.configure(&self.device, &self.config);
      self.depth = depth_texture(&self.device, s);
  }
224
  fn draw(&mut self) -> Result<(), Box<dyn std::error::Error>> {
      // Presentation timestamps determine *when* to advance the video texture.
      // The cube itself renders at the display refresh cadence.
      self.video.update(&self.device)?;
      let elapsed = self.start.elapsed().as_secs_f32();
      let model = Mat4::from_rotation_y(elapsed * 0.8) * Mat4::from_rotation_x(elapsed * 0.43);
      let view = Mat4::look_at_rh(Vec3::new(3.2, 2.1, 4.6), Vec3::ZERO, Vec3::Y);
      let projection = Mat4::perspective_rh(45f32.to_radians(), self.config.width as f32 / self.config.height as f32, 0.1, 100.0);
      let uniform = Uniform { mvp: (projection * view * model).to_cols_array_2d() };
      self.queue.write_buffer(&self.uniform, 0, bytemuck::bytes_of(&uniform));
      let frame = match self.surface.get_current_texture() {
          wgpu::CurrentSurfaceTexture::Success(frame)
          | wgpu::CurrentSurfaceTexture::Suboptimal(frame) => frame,
238
          wgpu::CurrentSurfaceTexture::Timeout
          | wgpu::CurrentSurfaceTexture::Occluded => {
              return Ok(());
          }
243
          wgpu::CurrentSurfaceTexture::Outdated => {
              self.surface.configure(&self.device, &self.config);
              return Ok(());
          }
248
          wgpu::CurrentSurfaceTexture::Lost => {
              panic!("Surface lost; surface recreation required");
          }
252
          wgpu::CurrentSurfaceTexture::Validation => {
              panic!("Surface validation error");
          }
      };
      let output = frame.texture.create_view(&Default::default());
      let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("Cube render") });
      {
          let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
              label: Some("Cube main pass"),
              color_attachments: &[Some(wgpu::RenderPassColorAttachment {
                  view: &output, resolve_target: None,
                  ops: wgpu::Operations { load: wgpu::LoadOp::Clear(self.clear), store: wgpu::StoreOp::Store },
                  depth_slice: None,
              })],
              depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
                  view: &self.depth,
                  depth_ops: Some(wgpu::Operations { load: wgpu::LoadOp::Clear(1.0), store: wgpu::StoreOp::Store }),
                  stencil_ops: None,
              }),
              occlusion_query_set: None,
              timestamp_writes: None,
              multiview_mask: None,
          });
          pass.set_pipeline(&self.pipeline);
          pass.set_bind_group(0, &self.bind_groups[self.video.current_slot()], &[]);
          pass.set_vertex_buffer(0, self.vertices.slice(..));
          pass.set_index_buffer(self.indices.slice(..), wgpu::IndexFormat::Uint16);
          pass.draw_indexed(0..36, 0, 0..1);
      }
      self.queue.submit(Some(encoder.finish()));
      self.window.pre_present_notify();
      self.queue.present(frame);
      Ok(())
  }
287}
288
289impl Drop for State {
  fn drop(&mut self) {
      let _ = self.device.poll(wgpu::PollType::Wait { submission_index: None, timeout: None });
  }
293}
294
295struct App {
  path: String,
  state: Option<State>,
  next_frame: Instant,
299}
300impl ApplicationHandler for App {
  fn resumed(&mut self, event_loop: &ActiveEventLoop) {
      if self.state.is_some() { return; }
      let window = Arc::new(event_loop.create_window(WindowAttributes::default()
          .with_title("wgpu cube — native NVDEC / CUDA / D3D12")
          .with_inner_size(winit::dpi::LogicalSize::new(1024, 768))).unwrap());
      let state = pollster::block_on(State::new(window, &self.path)).expect("Video cube initialization failed");
      self.next_frame = Instant::now();
      self.state = Some(state);
  }
  fn window_event(&mut self, event_loop: &ActiveEventLoop, _id: WindowId, event: WindowEvent) {
      let Some(state) = self.state.as_mut() else { return; };
      match event {
          WindowEvent::CloseRequested => event_loop.exit(),
          WindowEvent::Resized(s) => state.resize(s),
          WindowEvent::RedrawRequested if state.config.width > 0 && state.config.height > 0 => {
              if let Err(e) = state.draw() {
                  eprintln!("Render/decode error: {e}");
                  event_loop.exit();
              }
          }
          _ => {}
      }
  }
  fn about_to_wait(&mut self, event_loop: &ActiveEventLoop) {
      if let Some(state) = &self.state {
          if Instant::now() >= self.next_frame {
              state.window.request_redraw();
              self.next_frame = Instant::now() + Duration::from_millis(16);
          }
          event_loop.set_control_flow(ControlFlow::WaitUntil(self.next_frame));
      }
  }
333}
334
335fn main() -> Result<(), Box<dyn std::error::Error>> {
  let path = std::env::args().nth(1).ok_or("Usage: cargo run --release -- <video.mp4>")?;
  let event_loop = EventLoop::new()?;
  let mut app = App { path, state: None, next_frame: Instant::now() };
  event_loop.run_app(&mut app)?;
  Ok(())
341}

src/cube.wgsl

26 lines, shown in full in Step 10. Reading it against main.rs: the three @group(0) @binding(n) declarations are the bind group layout, in the same order; VertexInput's two @locations are the vertex buffer layout; vs_main multiplies each vertex by the matrix from the uniform buffer and passes the texture coordinate through; and fs_main samples the video at that coordinate. There's no lighting, no color correction and no hint that the texture is video. Any WGSL that samples a texture would work here unchanged.

Limitations and next steps

This is a demonstration of the interop, not a video player, and I want to be clear about which is which. Here is what the example does and doesn't do:

What the example does and doesn't do
Feature	Status
Hardware H.264 and HEVC decoding	Yes, for progressive 8-bit 4:2:0
Decoded frames stay in GPU memory	Yes
CUDA writing into shared D3D12 textures	Yes
Native textures imported into wgpu	Yes, through wgpu-hal
Video on 3D geometry	Yes
Audio	No
10-bit, HDR, P010	No
Color conversion from the stream's metadata	No: BT.709 limited range is assumed
Asynchronous CUDA and wgpu	No: a blocking wait every frame
Other platforms	No: Windows, D3D12 and NVIDIA only

And to repeat the caveat from Part 1: the demuxing and the bitstream parsing run on the CPU. It's the decoding, the color conversion and the rendering that happen on the GPU.

Color accuracy

The double-gamma bug in Step 10 is fixed, but the kernel still assumes BT.709 limited range for every file. SD video is usually BT.601, HDR is BT.2020 with a different transfer function, and some files are full range. My clip carries no color metadata at all, which is common, so a real player has to guess too, typically from the resolution. Doing this properly means reading the stream's color metadata (Media Foundation exposes it on the media type) and choosing the matrix in the kernel accordingly.

Synchronization

The blocking waits in Step 11 are the biggest cost. The proper version imports a D3D12 fence per slot into CUDA as an external semaphore (the same CUDA interop API has cudaImportExternalSemaphore for this). CUDA signals the slot's fence after the kernel, the render queue waits on it before sampling, and CUDA waits on the render queue's fence before writing the slot again. With a decoder thread and a small frame queue, decode and rendering could overlap.

Audio and timing

There's no audio, so the video clock is the only clock. There's no frame dropping either: if a frame is late, it's shown late. Packet reading is synchronous on the UI thread, so a slow disk read would stall rendering too. A player would read on a worker thread and pace video against the audio clock.

HDR and other formats

10-bit video (P010), 4:2:2, interlaced streams and mid-stream resolution changes are all rejected in Step 5. NVDEC can decode more than this, AV1 included on newer GPUs, but each one needs a different output format, a different kernel, or a decoder that can be recreated on the fly. HDR also needs tone mapping, which is a topic of its own.

Wrapping up

Stepping back from the details, the useful result is small and reusable: an NVDEC-decoded frame as an ordinary, sampleable wgpu::Texture, with no trip through system memory. The cube is just where I happened to show it. The same three textures could feed a fullscreen quad, a material in a game engine, a video wall, or a compute pass. And since the frames are already in CUDA when they're mapped, it's also a starting point for running CUDA processing on decoded video before it's drawn.

If I had to name the three ideas that made it work: the LUID match in Step 4, so CUDA and wgpu are on one GPU; the shared, simultaneous-access texture in Step 7, so both can use one allocation; and wgpu-hal in Step 9, so wgpu can adopt memory it didn't allocate. Everything else is plumbing between those three. The parts that needed fixing were the less obvious ones: how NAL units are framed, two C runtimes in one binary, and gamma applied twice.

Terms used in this post

NVDEC: The hardware video decoder on NVIDIA GPUs, separate from the graphics and compute engines.
NVCUVID: The Video Codec SDK's decoding API: the parser (cuvidCreateVideoParser) and the decoder (cuvidCreateDecoder).
Demuxing: Pulling the compressed samples of one track out of a container such as MP4, with their timestamps.
NAL unit: The basic unit of an H.264 or HEVC stream: a header byte (two in HEVC) followed by a slice of a picture or a parameter set such as the SPS or PPS.
Annex B: The byte-stream format that separates NAL units with start codes (00 00 01). MP4 files usually store them with length prefixes instead.
NV12: An 8-bit 4:2:0 YUV layout: a full-size Y plane, then a half-size plane of interleaved U and V.
PTS: Presentation timestamp: when a frame should be shown. Media Foundation and NVDEC both use 100 ns units.
LUID: The locally unique ID Windows assigns each graphics adapter. D3D12 and CUDA both report it, which is how the bridge matches them.
CUDA context: CUDA's per-device state. Every CUDA call acts on the context that is current on the calling thread.
wgpu-hal: The layer under wgpu that implements each backend. Its unsafe functions let you use native resources.
External memory: CUDA's name for memory allocated by another API, such as a D3D12 resource, and imported through a handle.
Surface object: CUDA's handle for writing to a texture-like array from a kernel.
Simultaneous access: A D3D12 texture flag that lets more than one queue or API use a texture without layout transitions.

References and updates

October 24, 2026: first version.