diff --git a/src/torchcodec/_core/BetaCudaDeviceInterface.cpp b/src/torchcodec/_core/BetaCudaDeviceInterface.cpp
index 0c7764c43..2d977b9a9 100644
--- a/src/torchcodec/_core/BetaCudaDeviceInterface.cpp
+++ b/src/torchcodec/_core/BetaCudaDeviceInterface.cpp
@@ -266,10 +266,48 @@ std::optional<cudaVideoSurfaceFormat> get_nvdec_surface_format(
 
 // Callback for freeing CUDA memory associated with AVFrame see where it's used
 // for more details.
+// TODO_API_BREAKDOWN P2: Should we align this with the other free callback
+// below?  Why did we use cudaMalloc? Can we just allocate with torch??
 void cuda_buffer_free_callback(void* opaque, [[maybe_unused]] uint8_t* data) {
   cudaFree(opaque);
 }
 
+void standalone_frame_free_callback(
+    [[maybe_unused]] void* opaque,
+    uint8_t* data) {
+  delete reinterpret_cast<StandAloneFrameAttachedData*>(data);
+}
+
+class CudaContextGuard {
+  // There's one CUDA context per process per device. But new threads aren't
+  // bound to a context. The binding often happens automatically when calling
+  // CUDA APIs (like cudaFree), but some APIs like the NVCUVID ones that we use
+  // here aren't automatically binding.
+  // So for a thread to be able to use NVCUVID APIs, it must have a context
+  // bound to it, and we have to enforce that binding manually.
+  // That's what this guard does: it calls cudaFree(nullptr), which is a common
+  // near-free way to force the CUDA runtime to bind the context for the current
+  // thread. And this call must happen within a device guard to make sure we're
+  // binding the context of the device this interface is using.
+  // We must call this guard in every public method of the interface that uses
+  // NVCUVID APIs, because these methods can, in theory, be called from any
+  // thread.
+  // Note that none of this was an issue before when our only entry-point was
+  // the SingleStreamDecoder: all the entry-points were called from the same
+  // thread. Now that we have split the APIs in different blocks (PacketDecoder,
+  // ColorConverter), each of these blocks can be on different threads - and
+  // importantly, they can be created in the main thread (where the context is
+  // bound by our call to initialize_cuda_context_with_pytorch()), but then used
+  // in a different thread that doesn't have the context.
+ public:
+  explicit CudaContextGuard(int device_index) : device_guard_(device_index) {
+    cudaFree(nullptr);
+  }
+
+ private:
+  StableDeviceGuard device_guard_;
+};
+
 } // namespace
 
 BetaCudaDeviceInterface::BetaCudaDeviceInterface(const StableDevice& device)
@@ -278,18 +316,31 @@ BetaCudaDeviceInterface::BetaCudaDeviceInterface(const StableDevice& device)
   STD_TORCH_CHECK(
       device_.type() == kStableCUDA, "Unsupported device: must be CUDA");
 
+  // Note: now that we have the CudaContextGuard, we might not need to do that
+  // anymore. The comment says we need pytorch to create the context - maybe
+  // that's true, but that's a very old comment now.
   initialize_cuda_context_with_pytorch(device_);
 
   nvcuvid_available_ = load_nvcuvid_library();
 }
 
-void BetaCudaDeviceInterface::initialize_video(
-    const AVStream* av_stream,
-    const UniqueDecodingAVFormatContext& av_format_ctx,
+void BetaCudaDeviceInterface::initialize_color_conversion(
     const VideoStreamOptions& video_stream_options,
     const std::vector<std::unique_ptr<Transform>>& transforms,
     const std::optional<FrameDims>& resized_output_dims) {
+  output_dtype_ = video_stream_options.output_dtype;
+  if (cpu_fallback_) {
+    cpu_fallback_->initialize_color_conversion(
+        video_stream_options, transforms, resized_output_dims);
+  }
+}
+
+void BetaCudaDeviceInterface::initialize_video_decoding(
+    const AVStream* av_stream,
+    const UniqueDecodingAVFormatContext& av_format_ctx,
+    const VideoStreamOptions& video_stream_options) {
   STD_TORCH_CHECK(av_stream != nullptr, "AVStream cannot be null");
+  CudaContextGuard context_guard(device_.index());
   rotation_ = rotation_from_degrees(get_rotation_from_stream(av_stream));
   output_dtype_ = video_stream_options.output_dtype;
 
@@ -308,12 +359,8 @@ void BetaCudaDeviceInterface::initialize_video(
     STD_TORCH_CHECK(
         cpu_fallback_ != nullptr, "Failed to create CPU device interface");
     cpu_fallback_->initialize(codec_context_);
-    cpu_fallback_->initialize_video(
-        av_stream,
-        av_format_ctx,
-        video_stream_options,
-        transforms,
-        resized_output_dims);
+    cpu_fallback_->initialize_video_decoding(
+        av_stream, av_format_ctx, video_stream_options);
     return;
   }
 
@@ -393,6 +440,7 @@ void BetaCudaDeviceInterface::send_seqhdr_packet() {
 }
 
 BetaCudaDeviceInterface::~BetaCudaDeviceInterface() {
+  CudaContextGuard context_guard(device_.index());
   if (decoder_) {
     // DALI doesn't seem to do any particular cleanup of the decoder before
     // sending it to the cache, so we probably don't need to do anything either.
@@ -544,6 +592,7 @@ int BetaCudaDeviceInterface::stream_property_change(
 // Moral equivalent of avcodec_send_packet(). Here, we pass the AVPacket down to
 // the NVCUVID parser.
 int BetaCudaDeviceInterface::send_packet(ReferenceAVPacket& packet) {
+  CudaContextGuard context_guard(device_.index());
   if (cpu_fallback_) {
     return cpu_fallback_->send_packet(packet);
   }
@@ -571,6 +620,7 @@ int BetaCudaDeviceInterface::send_packet(ReferenceAVPacket& packet) {
 }
 
 int BetaCudaDeviceInterface::send_eof_packet() {
+  CudaContextGuard context_guard(device_.index());
   if (cpu_fallback_) {
     return cpu_fallback_->send_eof_packet();
   }
@@ -646,6 +696,7 @@ int BetaCudaDeviceInterface::frame_ready_in_display_order(
 
 // Moral equivalent of avcodec_receive_frame().
 int BetaCudaDeviceInterface::receive_frame(UniqueAVFrame& av_frame) {
+  CudaContextGuard context_guard(device_.index());
   if (cpu_fallback_) {
     return cpu_fallback_->receive_frame(av_frame);
   }
@@ -682,12 +733,14 @@ int BetaCudaDeviceInterface::receive_frame(UniqueAVFrame& av_frame) {
   // unmap will cause map to eventually fail. DALI unmaps frames almost
   // immediately  after mapping them: they do the color-conversion in-between,
   // which involves a copy of the data, so that works.
-  // We, OTOH, will do the color-conversion later, outside of ReceiveFrame(). So
-  // we unmap here: just before mapping a new frame. At that point we know that
-  // the previously-mapped frame is no longer needed: it was either
-  // color-converted (with a copy), or that's a frame that was discarded in
-  // SingleStreamDecoder. Either way, the underlying output surface can be
-  // safely re-used.
+  // We, OTOH, will do the color-conversion later, outside of receive_frame().
+  // So we unmap here: just before mapping a new frame. At that point we know
+  // that the previously-mapped frame is no longer needed:
+  // - With SingleStreamDecoder, that frame was either color-converted (with a
+  //   copy), or that's a frame that was discarded in SingleStreamDecoder.
+  //   Either way, the underlying output surface can be safely re-used.
+  // - With the "Blocks" APIs, the PacketDecoder forces a copy in
+  //   make_frame_standalone().
   unmap_previous_frame();
   CUresult result = cuvidMapVideoFrame(
       *decoder_.get(),
@@ -795,15 +848,85 @@ UniqueAVFrame BetaCudaDeviceInterface::convert_cuda_frame_to_av_frame(
       reinterpret_cast<uint8_t*>(frame_ptr + (pitch * even_height));
   av_frame->data[2] = nullptr;
   av_frame->data[3] = nullptr;
-  av_frame->linesize[0] = pitch;
-  av_frame->linesize[1] = pitch;
+  // TODO_API_BREAKDOWN_CUDA P2: Check range before cast?
+  av_frame->linesize[0] = static_cast<int>(pitch);
+  av_frame->linesize[1] = static_cast<int>(pitch);
   av_frame->linesize[2] = 0;
   av_frame->linesize[3] = 0;
 
   return av_frame;
 }
 
+void BetaCudaDeviceInterface::make_frame_standalone(UniqueAVFrame& av_frame) {
+  // Make the frame standalone:
+  // - Crucially, we copy the frame data so that its surface can be unmapped in
+  //   receive_frame() (see comment there).
+  // - We put the frame in a state such that it can be safely used by a
+  //   ColorConverter (i.e. a *different* instance of this
+  //   BetaCudaDeviceInterface): we attach relevant metadata as the
+  //   StandAloneFrameAttachedData struct, which is then used by the
+  //   ColorConverter in convert_cuda_frame_to_av_frame() to perform the
+  //   color-conversion correctly.
+  CudaContextGuard context_guard(device_.index());
+  if (!(av_frame->format == AV_PIX_FMT_P016LE ||
+        av_frame->format == AV_PIX_FMT_NV12)) {
+    // The CPU frames are already standalone, so we don't need to do anything.
+    return;
+  }
+
+  // The amount of bytes an NV12 image takes is:
+  // num_bytes =  len(Y) + len(UV)
+  //           = num_pixels + num_pixels / 2
+  //           = num_pixels * 3 / 2
+  //
+  // where num_pixels = pitch * height, not num_pixels = width * height. The
+  // pitch value also accounts for the data size (uint8 vs uint16) so this is
+  // also correct for P016.
+  int64_t even_height =
+      static_cast<int64_t>(round_up_to_even(av_frame->height));
+  int64_t pitch = static_cast<int64_t>(av_frame->linesize[0]);
+  int64_t num_bytes = pitch * even_height * 3 / 2;
+
+  auto storage =
+      torch::stable::empty({num_bytes}, kStableUInt8, std::nullopt, device_);
+
+  // TODO_API_BREAKDOWN_CUDA P1: I suspect we don't need to wait on the nvdec
+  // stream here, because we can only arrive here from a path where the frame
+  // has already been mapped so its data is available - worth double checking.
+  cudaStream_t current_stream = get_current_cuda_stream(device_.index());
+  cudaError_t err = cudaMemcpyAsync(
+      storage.mutable_data_ptr(),
+      av_frame->data[0],
+      static_cast<size_t>(num_bytes),
+      cudaMemcpyDeviceToDevice,
+      current_stream);
+  STD_TORCH_CHECK(
+      err == cudaSuccess,
+      "Failed to copy NVDEC surface: ",
+      cudaGetErrorString(err));
+
+  // TODO_API_BREAKDOWN_CUDA P2: Should we unmap here? Or let the next
+  // receive_frame() call do it?
+  // unmap_previous_frame();
+
+  auto y_plane = static_cast<uint8_t*>(storage.mutable_data_ptr());
+  av_frame->data[0] = y_plane;
+  av_frame->data[1] = y_plane + (pitch * even_height);
+
+  auto attached_data = new StandAloneFrameAttachedData();
+  attached_data->producer_stream = current_stream;
+  // TODO_API_BREAKDOWN_CUDA P2: We don't *really* need to std::move it I guess?
+  attached_data->storage = std::move(storage);
+  av_frame->opaque_ref = av_buffer_create(
+      reinterpret_cast<uint8_t*>(attached_data),
+      sizeof(StandAloneFrameAttachedData),
+      standalone_frame_free_callback,
+      nullptr,
+      0);
+}
+
 void BetaCudaDeviceInterface::flush() {
+  CudaContextGuard context_guard(device_.index());
   if (cpu_fallback_) {
     cpu_fallback_->flush();
     return;
@@ -811,10 +934,10 @@ void BetaCudaDeviceInterface::flush() {
 
   // The NVCUVID docs mention that after seeking, i.e. when flush() is called,
   // we should send a packet with the CUVID_PKT_DISCONTINUITY flag. The docs
-  // don't say whether this should be an empty packet, or whether it should be a
-  // flag on the next non-empty packet. It doesn't matter: neither work :)
-  // Sending an EOF packet, however, does work. So we do that. And we re-set the
-  // eofSent_ flag to false because that's not a true EOF notification.
+  // don't say whether this should be an empty packet, or whether it should be
+  // a flag on the next non-empty packet. It doesn't matter: neither work :)
+  // Sending an EOF packet, however, does work. So we do that. And we re-set
+  // the eofSent_ flag to false because that's not a true EOF notification.
   send_eof_packet();
   eof_sent_ = false;
 
@@ -827,17 +950,17 @@ void BetaCudaDeviceInterface::flush() {
 UniqueAVFrame BetaCudaDeviceInterface::transfer_cpu_frame_to_gpu(
     const AVFrame& cpu_frame,
     AVPixelFormat target_pix_fmt) {
-  // This is called in the context of the CPU fallback: the frame was decoded on
-  // the CPU, and in this function we convert that frame into NV12 or P016
+  // This is called in the context of the CPU fallback: the frame was decoded
+  // on the CPU, and in this function we convert that frame into NV12 or P016
   // format and send it to the GPU.
   // We do that in 2 steps:
   // - First we convert the input CPU frame into an intermediate NV12/P016 CPU
   //   frame using sws_scale.
   // - Then we allocate GPU memory and copy the CPU frame to the GPU. This
   //   is what we return.
-  // Since NV12/P016 require even dimensions, the returned frame will have even
-  // (rounded up) width and height, even if the original CPU frame had odd
-  // dimensions.
+  // Since NV12/P016 require even dimensions, the returned frame will have
+  // even (rounded up) width and height, even if the original CPU frame had
+  // odd dimensions.
 
   // NV12 = 1 byte per sample, P016 = 2 bytes per sample
   STD_TORCH_CHECK(
@@ -950,9 +1073,9 @@ UniqueAVFrame BetaCudaDeviceInterface::transfer_cpu_frame_to_gpu(
       "Failed to copy frame properties: ",
       get_ffmpeg_error_string_from_error_code(ret));
 
-  // We need to make sure the CUDA memory is freed properly. Since we allocated
-  // it ourselves, FFmpeg doesn't know how to free it. We associate a `free`
-  // callback via opaque_ref that will be called by av_frame_free().
+  // We need to make sure the CUDA memory is freed properly. Since we
+  // allocated it ourselves, FFmpeg doesn't know how to free it. We associate
+  // a `free` callback via opaque_ref that will be called by av_frame_free().
   gpu_frame->opaque_ref = av_buffer_create(
       nullptr, // data - we don't need any
       0, // data size
@@ -970,7 +1093,14 @@ void BetaCudaDeviceInterface::convert_av_frame_to_frame_output(
     const AVFrame& av_frame,
     FrameOutput& frame_output,
     std::optional<torch::stable::Tensor> pre_allocated_output_tensor) {
-  if (cpu_fallback_) {
+  CudaContextGuard context_guard(device_.index());
+  // TODO_API_BREAKDOWN_CUDA P0 is that accurate and safe? Can there be a CPU
+  // NV12 frame in our code? Should we create a helper used in the
+  // make_standalone function too?
+  bool cpu_fallback = av_frame.format != AV_PIX_FMT_NV12 &&
+      av_frame.format != AV_PIX_FMT_P016LE;
+
+  if (cpu_fallback) {
     // When the CPU fallback happens, we'll try to run the color-conversion on
     // GPU by sending those CPU frames to the GPU as NV12 or P016 (See
     // transferCpuFrameToGpu() below). However, it's not always
@@ -982,6 +1112,7 @@ void BetaCudaDeviceInterface::convert_av_frame_to_frame_output(
         av_pix_fmt_desc_get(static_cast<AVPixelFormat>(av_frame.format));
     bool is444 = desc && desc->log2_chroma_w == 0 && desc->log2_chroma_h == 0;
     if (is444) {
+      // TODO_API_BREAKDOWN P1: we need to handle this
       FrameOutput cpu_frame_output;
       cpu_fallback_->convert_av_frame_to_frame_output(
           av_frame, cpu_frame_output);
@@ -1003,37 +1134,48 @@ void BetaCudaDeviceInterface::convert_av_frame_to_frame_output(
   // may round them up to even.
   FrameDims original_dims(av_frame.height, av_frame.width);
 
-  // On the CPU fallback we own the GPU frame we just created; otherwise the
-  // input frame is already what we need, and we only observe it.
   UniqueAVFrame transferred_frame;
-  if (cpu_fallback_) {
+  if (cpu_fallback) {
     AVPixelFormat target_pix_fmt = (output_dtype_ == OutputDtype::FLOAT32)
         ? AV_PIX_FMT_P016LE
         : AV_PIX_FMT_NV12;
     transferred_frame = transfer_cpu_frame_to_gpu(av_frame, target_pix_fmt);
   }
-  const AVFrame& gpu_frame = cpu_fallback_ ? *transferred_frame : av_frame;
+  const AVFrame& gpu_frame = cpu_fallback ? *transferred_frame : av_frame;
 
   STD_TORCH_CHECK(
       gpu_frame.format == AV_PIX_FMT_NV12 ||
           gpu_frame.format == AV_PIX_FMT_P016LE,
       "Expected NV12 or P016LE format frame");
 
-  cudaStream_t nvdec_stream = get_current_cuda_stream(device_.index());
+  // TODO_API_BREAKDOWN P1: Cleanup how we get the attached data? Make it more
+  // robust? Should we couple it to a flag on the interface saying "I'm
+  // color-conversion only, I absolutely expect frames to be standalone"?
+  cudaStream_t producer_stream;
+  if (av_frame.opaque_ref != nullptr &&
+      av_frame.opaque_ref->size == sizeof(StandAloneFrameAttachedData)) {
+    auto attached_data = reinterpret_cast<StandAloneFrameAttachedData*>(
+        av_frame.opaque_ref->data);
+    producer_stream = attached_data->producer_stream;
+  } else {
+    producer_stream = get_current_cuda_stream(device_.index());
+  }
 
+  // TODO_API_BREAKDOWN P1: we don't suppor output_dtype so some of that is not
+  // execrcized.
   auto convert_frame = [&](std::optional<torch::stable::Tensor> pre_alloc)
       -> torch::stable::Tensor {
     bool is_p016 = (gpu_frame.format == AV_PIX_FMT_P016LE);
     int bit_depth = 8;
     if (is_p016) {
-      bit_depth = cpu_fallback_
+      bit_depth = cpu_fallback
           ? codec_context_->bits_per_raw_sample
           : static_cast<int>(video_format_.bit_depth_luma_minus8) + 8;
     }
     return convert_yuv_frame_to_rgb(
         gpu_frame,
         device_,
-        nvdec_stream,
+        producer_stream,
         pre_alloc,
         original_dims,
         is_p016,
diff --git a/src/torchcodec/_core/BetaCudaDeviceInterface.h b/src/torchcodec/_core/BetaCudaDeviceInterface.h
index e2d0e3518..e9af08cdf 100644
--- a/src/torchcodec/_core/BetaCudaDeviceInterface.h
+++ b/src/torchcodec/_core/BetaCudaDeviceInterface.h
@@ -33,6 +33,10 @@
 #include "nvcuvid_include/nvcuvid.h"
 
 namespace facebook::torchcodec {
+struct StandAloneFrameAttachedData {
+  cudaStream_t producer_stream = nullptr;
+  torch::stable::Tensor storage;
+};
 
 class BetaCudaDeviceInterface : public DeviceInterface {
  public:
@@ -41,9 +45,12 @@ class BetaCudaDeviceInterface : public DeviceInterface {
 
   void initialize(const SharedAVCodecContext& codec_context) override;
 
-  void initialize_video(
+  void initialize_video_decoding(
       const AVStream* av_stream,
       const UniqueDecodingAVFormatContext& av_format_ctx,
+      const VideoStreamOptions& video_stream_options) override;
+
+  void initialize_color_conversion(
       const VideoStreamOptions& video_stream_options,
       const std::vector<std::unique_ptr<Transform>>& transforms,
       const std::optional<FrameDims>& resized_output_dims) override;
@@ -91,6 +98,8 @@ class BetaCudaDeviceInterface : public DeviceInterface {
       unsigned int pitch,
       const CUVIDPARSERDISPINFO& disp_info);
 
+  void make_frame_standalone(UniqueAVFrame& av_frame) override;
+
   UniqueAVFrame transfer_cpu_frame_to_gpu(
       const AVFrame& cpu_frame,
       AVPixelFormat target_pix_fmt);
diff --git a/src/torchcodec/_core/ColorConverter.cpp b/src/torchcodec/_core/ColorConverter.cpp
index 2e06dd617..196e32a05 100644
--- a/src/torchcodec/_core/ColorConverter.cpp
+++ b/src/torchcodec/_core/ColorConverter.cpp
@@ -15,10 +15,8 @@
 
 namespace facebook::torchcodec {
 
-ColorConverter::ColorConverter(
-    const StableDevice& device,
-    std::string_view device_variant) {
-  device_interface_ = create_device_interface(device, device_variant);
+ColorConverter::ColorConverter(const StableDevice& device) {
+  device_interface_ = create_device_interface(device);
   STD_TORCH_CHECK(
       device_interface_ != nullptr,
       "Failed to create device interface. This should never happen, please report.");
@@ -27,22 +25,13 @@ ColorConverter::ColorConverter(
   options.output_dtype = OutputDtype::UINT8; // dtype not exposed yet
   options.device = device;
 
-  // No user transforms and no stream: the converter is stream-agnostic and
-  // derives everything it needs from each frame.
-  //
-  // TODO_API_BREAKDOWN Need to refac/rethink all this. It seems unnatural that
-  // the color-converter needs its own device_interface_, but at the same time
-  // the color-conversion *must* be third-party aware, and the only way to
-  // achieve that for now is via the interface.
-  // This will become very relevant when we tackle CUDA, so we can defer until
-  // then. For now this is an OK hack.
+  // TODO_API_BREAKDOWN P1 It seems unnatural that the color-converter needs its
+  // own device_interface_, but at the same time the color-conversion *must* be
+  // third-party aware, and the only way to achieve that for now is via the
+  // interface.
   std::vector<std::unique_ptr<Transform>> no_transforms;
-  device_interface_->initialize_video(
-      /*av_stream=*/nullptr,
-      UniqueDecodingAVFormatContext{},
-      options,
-      no_transforms,
-      /*resized_output_dims=*/std::nullopt);
+  device_interface_->initialize_color_conversion(
+      options, no_transforms, /*resized_output_dims=*/std::nullopt);
 }
 
 torch::stable::Tensor ColorConverter::convert(const AVFrame& av_frame) {
diff --git a/src/torchcodec/_core/ColorConverter.h b/src/torchcodec/_core/ColorConverter.h
index d06f2d4a2..5c66d1f1d 100644
--- a/src/torchcodec/_core/ColorConverter.h
+++ b/src/torchcodec/_core/ColorConverter.h
@@ -18,8 +18,7 @@ namespace facebook::torchcodec {
 class FORCE_PUBLIC_VISIBILITY ColorConverter {
  public:
   explicit ColorConverter(
-      const StableDevice& device = StableDevice(kStableCPU),
-      std::string_view device_variant = "default");
+      const StableDevice& device = StableDevice(kStableCPU));
 
   torch::stable::Tensor convert(const AVFrame& av_frame);
 
diff --git a/src/torchcodec/_core/CpuDeviceInterface.cpp b/src/torchcodec/_core/CpuDeviceInterface.cpp
index b6268adbd..0adb74d09 100644
--- a/src/torchcodec/_core/CpuDeviceInterface.cpp
+++ b/src/torchcodec/_core/CpuDeviceInterface.cpp
@@ -80,20 +80,18 @@ void CpuDeviceInterface::initialize(const SharedAVCodecContext& codec_context) {
   codec_context_ = codec_context;
 }
 
-void CpuDeviceInterface::initialize_video(
+void CpuDeviceInterface::initialize_video_decoding(
     const AVStream* av_stream,
     [[maybe_unused]] const UniqueDecodingAVFormatContext& av_format_ctx,
+    [[maybe_unused]] const VideoStreamOptions& video_stream_options) {
+  STD_TORCH_CHECK(av_stream != nullptr, "avStream is null");
+  time_base_ = av_stream->time_base;
+}
+
+void CpuDeviceInterface::initialize_color_conversion(
     const VideoStreamOptions& video_stream_options,
     const std::vector<std::unique_ptr<Transform>>& transforms,
     const std::optional<FrameDims>& resized_output_dims) {
-  // TODO_API_BREAKDOWN this used to be:
-  // STD_TORCH_CHECK(av_stream != nullptr, "avStream is null");
-  // time_base_ = av_stream->time_base;
-  // but now that avStrean can be null (to create a standalone color converter)
-  // we need this workaround. This is bad, we need to preserve the previous
-  // check somehow.
-  time_base_ = (av_stream != nullptr) ? av_stream->time_base
-                                      : AVRational{1, AV_TIME_BASE};
   av_media_type_ = AVMEDIA_TYPE_VIDEO;
   video_stream_options_ = video_stream_options;
   resized_output_dims_ = resized_output_dims;
diff --git a/src/torchcodec/_core/CpuDeviceInterface.h b/src/torchcodec/_core/CpuDeviceInterface.h
index 2d57eaf7b..556698330 100644
--- a/src/torchcodec/_core/CpuDeviceInterface.h
+++ b/src/torchcodec/_core/CpuDeviceInterface.h
@@ -27,9 +27,12 @@ class CpuDeviceInterface : public DeviceInterface {
 
   virtual void initialize(const SharedAVCodecContext& codec_context) override;
 
-  virtual void initialize_video(
+  virtual void initialize_video_decoding(
       const AVStream* av_stream,
       const UniqueDecodingAVFormatContext& av_format_ctx,
+      const VideoStreamOptions& video_stream_options) override;
+
+  virtual void initialize_color_conversion(
       const VideoStreamOptions& video_stream_options,
       const std::vector<std::unique_ptr<Transform>>& transforms,
       const std::optional<FrameDims>& resized_output_dims) override;
@@ -76,7 +79,11 @@ class CpuDeviceInterface : public DeviceInterface {
       const FrameDims& output_dims) const;
 
   VideoStreamOptions video_stream_options_;
-  AVRational time_base_;
+  // Default used when color conversion runs standalone (no stream to derive it
+  // from, e.g. the ColorConverter block API). initialize_video_decoding()
+  // overrides it from the stream when there is one. Its value doesn't matter on
+  // color-conversion-only mode, but filtergraph still expects it.
+  AVRational time_base_ = {1, AV_TIME_BASE};
   AVPixelFormat output_pixel_format_;
 
   // If the resized output dimensions are present, then we always use those as
diff --git a/src/torchcodec/_core/CudaDeviceInterface.cpp b/src/torchcodec/_core/CudaDeviceInterface.cpp
index 1e8c2c34f..d0e04a545 100644
--- a/src/torchcodec/_core/CudaDeviceInterface.cpp
+++ b/src/torchcodec/_core/CudaDeviceInterface.cpp
@@ -114,17 +114,13 @@ void CudaDeviceInterface::initialize(
   codec_context_ = codec_context;
 }
 
-void CudaDeviceInterface::initialize_video(
+void CudaDeviceInterface::initialize_video_decoding(
     const AVStream* av_stream,
     const UniqueDecodingAVFormatContext& av_format_ctx,
-    const VideoStreamOptions& video_stream_options,
-    [[maybe_unused]] const std::vector<std::unique_ptr<Transform>>& transforms,
-    [[maybe_unused]] const std::optional<FrameDims>& resized_output_dims) {
+    [[maybe_unused]] const VideoStreamOptions& video_stream_options) {
   STD_TORCH_CHECK(av_stream != nullptr, "avStream is null");
   time_base_ = av_stream->time_base;
-  video_stream_options_ = video_stream_options;
 
-  // TODO: Ideally, we should keep all interface implementations independent.
   cpu_interface_ = create_device_interface(kStableCPU);
   STD_TORCH_CHECK(
       cpu_interface_ != nullptr, "Failed to create CPU device interface");
@@ -137,6 +133,13 @@ void CudaDeviceInterface::initialize_video(
       /*resizedOutputDims=*/std::nullopt);
 }
 
+void CudaDeviceInterface::initialize_color_conversion(
+    const VideoStreamOptions& video_stream_options,
+    [[maybe_unused]] const std::vector<std::unique_ptr<Transform>>& transforms,
+    [[maybe_unused]] const std::optional<FrameDims>& resized_output_dims) {
+  video_stream_options_ = video_stream_options;
+}
+
 void CudaDeviceInterface::register_hardware_device_with_codec(
     AVCodecContext* codec_context) {
   STD_TORCH_CHECK(
diff --git a/src/torchcodec/_core/CudaDeviceInterface.h b/src/torchcodec/_core/CudaDeviceInterface.h
index d449ccb97..00cc5e817 100644
--- a/src/torchcodec/_core/CudaDeviceInterface.h
+++ b/src/torchcodec/_core/CudaDeviceInterface.h
@@ -28,14 +28,15 @@ class CudaDeviceInterface : public DeviceInterface {
 
   void initialize(const SharedAVCodecContext& codec_context) override;
 
-  void initialize_video(
+  void initialize_video_decoding(
       const AVStream* av_stream,
       const UniqueDecodingAVFormatContext& av_format_ctx,
+      const VideoStreamOptions& video_stream_options) override;
+
+  void initialize_color_conversion(
       const VideoStreamOptions& video_stream_options,
-      [[maybe_unused]] const std::vector<std::unique_ptr<Transform>>&
-          transforms,
-      [[maybe_unused]] const std::optional<FrameDims>& resized_output_dims)
-      override;
+      const std::vector<std::unique_ptr<Transform>>& transforms,
+      const std::optional<FrameDims>& resized_output_dims) override;
 
   void register_hardware_device_with_codec(
       AVCodecContext* codec_context) override;
diff --git a/src/torchcodec/_core/Demuxer.cpp b/src/torchcodec/_core/Demuxer.cpp
index ba46d6450..9835d0493 100644
--- a/src/torchcodec/_core/Demuxer.cpp
+++ b/src/torchcodec/_core/Demuxer.cpp
@@ -69,6 +69,8 @@ Demuxer::Demuxer(
 }
 
 UniqueAVPacket Demuxer::next_packet() {
+  // TODO_API_BREAKDOWN P2: Not a fan of the ReferenceAVPacket / AutoAVPacket /
+  // UniqueAVPacket dance here. Can we simplify?
   ReferenceAVPacket packet(auto_packet_);
   int status =
       read_next_packet(format_context_.get(), active_stream_index_, packet);
diff --git a/src/torchcodec/_core/DeviceInterface.h b/src/torchcodec/_core/DeviceInterface.h
index ca976fd36..f25efb342 100644
--- a/src/torchcodec/_core/DeviceInterface.h
+++ b/src/torchcodec/_core/DeviceInterface.h
@@ -56,15 +56,33 @@ class DeviceInterface {
   // default sendPacket/receiveFrame/flush implementations.
   virtual void initialize(const SharedAVCodecContext& codec_context) = 0;
 
-  // Initialize the device with parameters specific to video decoding. There is
-  // a default empty implementation.
-  virtual void initialize_video(
+  // Initialize state needed to decode packets into raw AVFrames.
+  virtual void initialize_video_decoding(
       [[maybe_unused]] const AVStream* av_stream,
       [[maybe_unused]] const UniqueDecodingAVFormatContext& av_format_ctx,
+      [[maybe_unused]] const VideoStreamOptions& video_stream_options) {}
+
+  // Initialize state needed to color-convert decoded AVFrames into output
+  // tensors.
+  virtual void initialize_color_conversion(
       [[maybe_unused]] const VideoStreamOptions& video_stream_options,
       [[maybe_unused]] const std::vector<std::unique_ptr<Transform>>&
-          transforms,
-      [[maybe_unused]] const std::optional<FrameDims>& resized_output_dims) {}
+          transforms = {},
+      [[maybe_unused]] const std::optional<FrameDims>& resized_output_dims =
+          std::nullopt) {}
+
+  // Convenience for the combined decode + color-convert path, kept for BC as
+  // it's used by SingleStreamDecoder and out-of-tree interfaces rely on it.
+  void initialize_video(
+      const AVStream* av_stream,
+      const UniqueDecodingAVFormatContext& av_format_ctx,
+      const VideoStreamOptions& video_stream_options,
+      const std::vector<std::unique_ptr<Transform>>& transforms,
+      const std::optional<FrameDims>& resized_output_dims) {
+    initialize_video_decoding(av_stream, av_format_ctx, video_stream_options);
+    initialize_color_conversion(
+        video_stream_options, transforms, resized_output_dims);
+  }
 
   // Initialize the device with parameters specific to audio decoding. There is
   // a default empty implementation.
@@ -138,6 +156,9 @@ class DeviceInterface {
     return avcodec_receive_frame(codec_context_.get(), av_frame.get());
   }
 
+  virtual void make_frame_standalone([[maybe_unused]] UniqueAVFrame& av_frame) {
+  };
+
   // Flush remaining frames from decoder
   virtual void flush() {
     STD_TORCH_CHECK(
@@ -200,7 +221,7 @@ TORCHCODEC_THIRD_PARTY_API bool register_device_interface(
 
 FORCE_PUBLIC_VISIBILITY void validate_device_interface(
     const std::string& device,
-    const std::string& variant);
+    const std::string& variant = "default");
 
 TORCHCODEC_THIRD_PARTY_API std::unique_ptr<DeviceInterface>
 create_device_interface(
diff --git a/src/torchcodec/_core/PacketDecoder.cpp b/src/torchcodec/_core/PacketDecoder.cpp
index 23e726f8e..b9960e92d 100644
--- a/src/torchcodec/_core/PacketDecoder.cpp
+++ b/src/torchcodec/_core/PacketDecoder.cpp
@@ -8,7 +8,7 @@
 
 namespace facebook::torchcodec {
 
-// TODO_API_BREAKDOWN: we should make sure the block APIs can dispatch to
+// TODO_API_BREAKDOWN P1: we should make sure the block APIs can dispatch to
 // third-party extensions - all of them.
 
 SharedAVCodecContext create_and_open_codec_context(
@@ -58,9 +58,8 @@ const AVCodec* find_decoder(
 PacketDecoder::PacketDecoder(
     const Demuxer& demuxer,
     const StableDevice& device,
-    std::string_view device_variant,
     std::optional<int> ffmpeg_thread_count) {
-  device_interface_ = create_device_interface(device, device_variant);
+  device_interface_ = create_device_interface(device);
   STD_TORCH_CHECK(
       device_interface_ != nullptr,
       "Failed to create device interface. This should never happen, please report.");
@@ -71,6 +70,13 @@ PacketDecoder::PacketDecoder(
   codec_context_ = create_and_open_codec_context(
       stream, av_codec, device_interface_.get(), ffmpeg_thread_count);
   device_interface_->initialize(codec_context_);
+
+  VideoStreamOptions options;
+  options.output_dtype = OutputDtype::UINT8; // dtype not exposed yet
+  options.device = device;
+
+  device_interface_->initialize_video_decoding(
+      stream, demuxer.format_context(), options);
 }
 
 int PacketDecoder::send_packet(AVPacket* packet) {
@@ -90,7 +96,11 @@ int PacketDecoder::send_eof() {
 }
 
 int PacketDecoder::receive_frame(UniqueAVFrame& av_frame) {
-  return device_interface_->receive_frame(av_frame);
+  int status = device_interface_->receive_frame(av_frame);
+  if (status == AVSUCCESS) {
+    device_interface_->make_frame_standalone(av_frame);
+  }
+  return status;
 }
 
 } // namespace facebook::torchcodec
diff --git a/src/torchcodec/_core/PacketDecoder.h b/src/torchcodec/_core/PacketDecoder.h
index 929698ac6..a22307874 100644
--- a/src/torchcodec/_core/PacketDecoder.h
+++ b/src/torchcodec/_core/PacketDecoder.h
@@ -33,7 +33,6 @@ class FORCE_PUBLIC_VISIBILITY PacketDecoder {
   explicit PacketDecoder(
       const Demuxer& demuxer,
       const StableDevice& device = StableDevice(kStableCPU),
-      std::string_view device_variant = "default",
       std::optional<int> ffmpeg_thread_count = std::nullopt);
 
   // Feed one packet to the decoder. Borrows `packet` (does not take ownership).
diff --git a/src/torchcodec/_core/color_conversion.cpp b/src/torchcodec/_core/color_conversion.cpp
index 0efa6d699..b05e8273c 100644
--- a/src/torchcodec/_core/color_conversion.cpp
+++ b/src/torchcodec/_core/color_conversion.cpp
@@ -191,7 +191,7 @@ void compute_rgb_to_yuv_matrix(
 torch::stable::Tensor convert_yuv_frame_to_rgb(
     const AVFrame& av_frame,
     const StableDevice& device,
-    cudaStream_t nvdec_stream,
+    cudaStream_t producer_stream,
     std::optional<torch::stable::Tensor> pre_allocated_output_tensor,
     const FrameDims& output_dims,
     bool is_p016,
@@ -221,9 +221,14 @@ torch::stable::Tensor convert_yuv_frame_to_rgb(
         FrameDims(out_height, out_width), device, out_dtype);
   }
 
+  // TODO_API_BREAKDOWN P1: This may not be the semantic that we want: this will
+  // wait for all ongoin work on the producer stream to finish. But maybe the
+  // producer stream produced the frame data a long time ago, and lots of
+  // kernels have been launched on it already. We'd be waiting on those to
+  // finish even though the data we need is already available.
   cudaStream_t stream = get_current_cuda_stream(device.index());
   sync_streams(
-      /*runningStream=*/nvdec_stream, /*waitingStream=*/stream);
+      /*runningStream=*/producer_stream, /*waitingStream=*/stream);
 
   maybe_update_color_matrix(
       cached_color_matrix,
diff --git a/src/torchcodec/_core/color_conversion.h b/src/torchcodec/_core/color_conversion.h
index ad3ae29ef..f2965acf8 100644
--- a/src/torchcodec/_core/color_conversion.h
+++ b/src/torchcodec/_core/color_conversion.h
@@ -86,7 +86,7 @@ void launch_p016_to_rgb16_kernel(
 torch::stable::Tensor convert_yuv_frame_to_rgb(
     const AVFrame& av_frame,
     const StableDevice& device,
-    cudaStream_t nvdec_stream,
+    cudaStream_t producer_stream,
     std::optional<torch::stable::Tensor> pre_allocated_output_tensor,
     const FrameDims& output_dims,
     bool is_p016,
diff --git a/src/torchcodec/_core/custom_ops.cpp b/src/torchcodec/_core/custom_ops.cpp
index 73351c99c..a2d30a113 100644
--- a/src/torchcodec/_core/custom_ops.cpp
+++ b/src/torchcodec/_core/custom_ops.cpp
@@ -77,14 +77,13 @@ STABLE_TORCH_LIBRARY_FRAGMENT(torchcodec_ns, m) {
       "_blocks_create_demuxer(str filename, int? stream_index=None) -> Tensor");
   m.def("_blocks_demuxer_next_packet(Tensor(a!) demuxer) -> (Tensor, bool)");
   m.def(
-      "_blocks_create_packet_decoder(Tensor demuxer, *, int? num_threads=None, str device=\"cpu\", str device_variant=\"default\") -> Tensor");
+      "_blocks_create_packet_decoder(Tensor demuxer, *, int? num_threads=None, str device=\"cpu\") -> Tensor");
   m.def(
       "_blocks_packet_decoder_send_packet(Tensor(a!) decoder, Tensor packet) -> int");
   m.def("_blocks_packet_decoder_send_eof(Tensor(a!) decoder) -> int");
   m.def(
       "_blocks_packet_decoder_receive_frame(Tensor(a!) decoder) -> (Tensor, int, float, float)");
-  m.def(
-      "_blocks_create_color_converter(str device=\"cpu\", str device_variant=\"default\") -> Tensor");
+  m.def("_blocks_create_color_converter(str device=\"cpu\") -> Tensor");
   m.def("_blocks_convert_frame(Tensor(a!) converter, Tensor frame) -> Tensor");
   m.def("_get_key_frame_indices(Tensor(a!) decoder) -> Tensor");
   m.def("get_json_metadata(Tensor(a!) decoder) -> str");
@@ -185,9 +184,7 @@ SingleStreamDecoder* unwrap_tensor_to_get_decoder(
 // (Demuxer / PacketDecoder / ColorConverter / AVPacket / AVFrame). Same trick
 // as wrap_decoder_pointer_to_tensor: the tensor's data pointer IS the raw
 // pointer, and the tensor is the owner. Taking the unique_ptr by value is what
-// makes the ownership transfer explicit at the call site; the unique_ptr's own
-// deleter is what frees the object, so FFmpeg types work here as long as they
-// arrive in their UniqueAVXxx alias.
+// makes the ownership transfer explicit at the call site.
 template <typename T, typename D>
 torch::stable::Tensor wrap_pointer_to_tensor(std::unique_ptr<T, D> ptr) {
   D object_deleter = ptr.get_deleter();
@@ -814,16 +811,15 @@ OpsPacketOutput _blocks_demuxer_next_packet(torch::stable::Tensor& demuxer) {
 torch::stable::Tensor _blocks_create_packet_decoder(
     torch::stable::Tensor& demuxer,
     std::optional<int64_t> num_threads,
-    std::string device,
-    std::string device_variant) {
+    std::string device) {
   Demuxer* demuxer_ptr = unwrap_tensor_to_pointer<Demuxer>(demuxer);
-  validate_device_interface(device, device_variant);
+  validate_device_interface(device);
   std::optional<int> thread_count;
   if (num_threads.has_value()) {
     thread_count = static_cast<int>(num_threads.value());
   }
   auto decoder = std::make_unique<PacketDecoder>(
-      *demuxer_ptr, StableDevice(device), device_variant, thread_count);
+      *demuxer_ptr, StableDevice(device), thread_count);
   return wrap_pointer_to_tensor<PacketDecoder>(std::move(decoder));
 }
 
@@ -831,6 +827,7 @@ int64_t _blocks_packet_decoder_send_packet(
     torch::stable::Tensor& decoder,
     torch::stable::Tensor& packet) {
   PacketDecoder* decoder_ptr = unwrap_tensor_to_pointer<PacketDecoder>(decoder);
+  // TODO_API_BREAKDOWN P1: Do we really need this to be a raw AVPacket*?
   AVPacket* raw_packet = unwrap_tensor_to_pointer<AVPacket>(packet);
   return static_cast<int64_t>(decoder_ptr->send_packet(raw_packet));
 }
@@ -871,12 +868,9 @@ OpsReceiveFrameOutput _blocks_packet_decoder_receive_frame(
       duration_seconds);
 }
 
-torch::stable::Tensor _blocks_create_color_converter(
-    std::string device,
-    std::string device_variant) {
-  validate_device_interface(device, device_variant);
-  auto converter =
-      std::make_unique<ColorConverter>(StableDevice(device), device_variant);
+torch::stable::Tensor _blocks_create_color_converter(std::string device) {
+  validate_device_interface(device);
+  auto converter = std::make_unique<ColorConverter>(StableDevice(device));
   return wrap_pointer_to_tensor<ColorConverter>(std::move(converter));
 }
 
diff --git a/src/torchcodec/decoders/_blocks/_color_converter.py b/src/torchcodec/decoders/_blocks/_color_converter.py
index 69e3bf3c7..786893e95 100644
--- a/src/torchcodec/decoders/_blocks/_color_converter.py
+++ b/src/torchcodec/decoders/_blocks/_color_converter.py
@@ -11,10 +11,10 @@
 
 from ._frame import DecodedFrame
 
-# TODO_API_BREAKDOWN support output_dtype?
-# TODO_API_BREAKDOWN Expose output_shape?
-
-# TODO_API_BREAKDOWN We need to support rotation metadata!!
+# TODO_API_BREAKDOWN FEAT support output_dtype?
+# TODO_API_BREAKDOWN FEAT We need to support rotation metadata!!
+# TODO_API_BREAKDOWN FEAT Implement seeking?
+# TODO_API_BREAKDOWN FEAT Implement range-getting (start/end time) for decoding?
 
 
 class ColorConverter:
@@ -30,8 +30,12 @@ class ColorConverter:
     block is intentionally stream-agnostic.
     """
 
-    def __init__(self):
-        self._handle = _blocks_create_color_converter()
+    # TODO_API_BREAKDOWN P1: device default should be None
+    # TODO_API_BREAKDOWN P1: add checks for coupling between device param of
+    # PacketDecoder and ColorConverter. What if one is CPU and the other is
+    # CUDA? What if they're different CUDA devices? Maybe we should just error.
+    def __init__(self, device="cpu"):
+        self._handle = _blocks_create_color_converter(device=device)
 
     def convert(self, decoded_frame: DecodedFrame) -> Frame:
         data = _blocks_convert_frame(self._handle, decoded_frame._handle)
diff --git a/src/torchcodec/decoders/_blocks/_packet_decoder.py b/src/torchcodec/decoders/_blocks/_packet_decoder.py
index 506096062..7191d0cfc 100644
--- a/src/torchcodec/decoders/_blocks/_packet_decoder.py
+++ b/src/torchcodec/decoders/_blocks/_packet_decoder.py
@@ -17,7 +17,7 @@
 from ._frame import DecodedFrame, Packet
 
 
-# TODO_API_BREAKDOWN revisit every single docstring / comments at some point.
+# TODO_API_BREAKDOWN P2 revisit every single docstring / comments at some point.
 
 
 class PacketDecoder:
@@ -31,8 +31,11 @@ class PacketDecoder:
     on your own threads.
     """
 
-    def __init__(self, demuxer: Demuxer):
-        self._handle = _blocks_create_packet_decoder(demuxer._handle, num_threads=1)
+    # TODO_API_BREAKDOWN P1: device default should be None, here and everywhere else
+    def __init__(self, demuxer: Demuxer, device="cpu"):
+        self._handle = _blocks_create_packet_decoder(
+            demuxer._handle, num_threads=1, device=device
+        )
 
     def _drain(self) -> list[DecodedFrame]:
         frames = []
@@ -53,7 +56,7 @@ def decode(self, packet: Packet) -> list[DecodedFrame]:
             raise RuntimeError(f"Failed to send packet to decoder (status {status})")
         return self._drain()
 
-    # TODO_API_BREAKDOWN maybe this shouldn't be called flush, at least not
+    # TODO_API_BREAKDOWN P2 maybe this shouldn't be called flush, at least not
     # as-is. It's not the same flush as the FFmpeg decoder buffer flush.
     def flush(self) -> list[DecodedFrame]:
         """Signal end-of-stream and return all remaining buffered frames. Call
diff --git a/test/test_decoders.py b/test/test_decoders.py
index 70a2bbf95..9cbc492fd 100644
--- a/test/test_decoders.py
+++ b/test/test_decoders.py
@@ -3275,14 +3275,17 @@ def test_multiple_calls_with_backward_seeks(self):
             assert wav_samples.pts_seconds == audio_samples.pts_seconds
 
 
+def _block_devices():
+    return ("cpu", pytest.param("cuda", marks=pytest.mark.needs_cuda))
+
+
 class TestBlocks:
 
-    def test_block_output_types(self):
+    @pytest.mark.parametrize("device", _block_devices())
+    def test_block_output_types(self, device):
         # Demuxer yields Packets, PacketDecoder yields DecodedFrames, and
         # ColorConverter yields Frames with the expected shape/dtype.
-        demuxer = Demuxer(NASA_VIDEO.path)
-        decoder = PacketDecoder(demuxer)
-        converter = ColorConverter()
+        demuxer, decoder, converter = self._make_blocks(NASA_VIDEO.path, device)
 
         num_packets = 0
         for packet in demuxer:
@@ -3295,6 +3298,7 @@ def test_block_output_types(self):
                 assert frame.data.ndim == 3  # CHW
                 assert frame.data.shape[0] == 3  # channels first
                 assert frame.data.dtype == torch.uint8
+                assert frame.data.device.type == device
                 assert frame.duration_seconds >= 0
 
         assert num_packets > 0
@@ -3351,42 +3355,41 @@ def drain():
 
         return drain()
 
-    def _decoded_frames(self, path):
-        # demux + decode, as a single generator of DecodedFrames (pts order).
+    @staticmethod
+    def _make_blocks(path, device):
         demuxer = Demuxer(path)
-        decoder = PacketDecoder(demuxer)
+        decoder = PacketDecoder(demuxer, device=device)
+        converter = ColorConverter(device=device)
+        return demuxer, decoder, converter
+
+    def _decoded_frames(self, path, device):
+        # demux + decode, as a single generator of DecodedFrames (pts order).
+        demuxer, decoder, _ = self._make_blocks(path, device)
         return self._decode(decoder, self._demux(demuxer))
 
-    def _decode_sequential(self, path):
+    def _decode_sequential(self, path, device):
         # demux -> decode -> color-convert, all on the calling thread.
-        demuxer = Demuxer(path)
-        decoder = PacketDecoder(demuxer)
-        converter = ColorConverter()
+        demuxer, decoder, converter = self._make_blocks(path, device)
         return list(
             self._convert(converter, self._decode(decoder, self._demux(demuxer)))
         )
 
-    def _decode_prefetch_frames(self, path):
+    def _decode_prefetch_frames(self, path, device):
         # [demux + decode] on one thread || [color-convert] on another.
-        demuxer = Demuxer(path)
-        decoder = PacketDecoder(demuxer)
-        converter = ColorConverter()
+        demuxer, decoder, converter = self._make_blocks(path, device)
+
         frames = self.prefetch(self._decode(decoder, self._demux(demuxer)))
         return list(self._convert(converter, frames))
 
-    def _decode_prefetch_packets(self, path):
+    def _decode_prefetch_packets(self, path, device):
         # [demux] on one thread || [decode + color-convert] on another.
-        demuxer = Demuxer(path)
-        decoder = PacketDecoder(demuxer)
-        converter = ColorConverter()
+        demuxer, decoder, converter = self._make_blocks(path, device)
         packets = self.prefetch(self._demux(demuxer))
         return list(self._convert(converter, self._decode(decoder, packets)))
 
-    def _decode_prefetch_packets_and_frames(self, path):
+    def _decode_prefetch_packets_and_frames(self, path, device):
         # [demux] || [decode] || [color-convert], each on its own thread.
-        demuxer = Demuxer(path)
-        decoder = PacketDecoder(demuxer)
-        converter = ColorConverter()
+        demuxer, decoder, converter = self._make_blocks(path, device)
         packets = self.prefetch(self._demux(demuxer))
         frames = self.prefetch(self._decode(decoder, packets))
         return list(self._convert(converter, frames))
@@ -3402,7 +3405,22 @@ def _to_frame_batch(self, frames):
             ),
         )
 
-    @pytest.mark.parametrize("video", (NASA_VIDEO, BT709_FULL_RANGE))
+    # TODO_API_BREAKDOWN P0: We need to test all assets. Generally we need more
+    # tests for all features / edge cases that were eventually fixed.
+    @pytest.mark.parametrize(
+        "video",
+        (
+            NASA_VIDEO,
+            BT709_FULL_RANGE,
+            NASA_VIDEO_HDR,
+            TEST_SRC_2_720P_HDR,
+            TEST_SRC_2_12BIT_HDR,
+            # NVDEC can't decode this one (too small), so on CUDA this covers
+            # the CPU-fallback path: the decoder hands out CPU frames and the
+            # converter has to notice and upload them itself.
+            H265_VIDEO,
+        ),
+    )
     @pytest.mark.parametrize(
         "decode_method",
         (
@@ -3413,10 +3431,12 @@ def _to_frame_batch(self, frames):
         ),
         ids=lambda f: f.__name__.removeprefix("_decode_"),
     )
-    def test_matches_video_decoder(self, video, decode_method):
-        got = self._to_frame_batch(decode_method(self, video.path))
-        ref = VideoDecoder(video.path).get_all_frames()
+    @pytest.mark.parametrize("device", _block_devices())
+    def test_matches_video_decoder(self, video, decode_method, device):
+        got = self._to_frame_batch(decode_method(self, video.path, device))
+        ref = VideoDecoder(video.path, device=device).get_all_frames()
 
+        assert got.data.device.type == device
         assert got.data.shape == ref.data.shape
         torch.testing.assert_close(got.data, ref.data, atol=0, rtol=0)
         torch.testing.assert_close(got.pts_seconds, ref.pts_seconds, atol=0, rtol=0)
@@ -3424,13 +3444,14 @@ def test_matches_video_decoder(self, video, decode_method):
             got.duration_seconds, ref.duration_seconds, atol=0, rtol=0
         )
 
-    def test_color_converter_reused_across_videos(self):
+    @pytest.mark.parametrize("device", _block_devices())
+    def test_color_converter_reused_across_videos(self, device):
         # A single unbound ColorConverter must correctly convert frames from two
         # different videos - here interleaved frame-by-frame, so the converter
         # switches input resolution/format on every call.
-        converter = ColorConverter()
+        converter = ColorConverter(device=device)
         videos = [NASA_VIDEO, BT709_FULL_RANGE]
-        generators = [self._decoded_frames(v.path) for v in videos]
+        generators = [self._decoded_frames(v.path, device) for v in videos]
         outputs = [[] for _ in videos]
 
         done = [False] * len(videos)
@@ -3446,10 +3467,20 @@ def test_color_converter_reused_across_videos(self):
 
         for video, frames in zip(videos, outputs):
             got = self._to_frame_batch(frames)
-            ref = VideoDecoder(video.path).get_all_frames()
+            ref = VideoDecoder(video.path, device=device).get_all_frames()
             assert got.data.shape == ref.data.shape
             torch.testing.assert_close(got.data, ref.data, atol=0, rtol=0)
 
+    @pytest.mark.parametrize("device", _block_devices())
+    def test_set_cuda_backend_is_a_noop(self, device):
+        # The blocks always use the NVDEC CUDA backend. Asking for the "ffmpeg"
+        # one changes nothing, rather than silently producing something else.
+        # TODO_API_BREAKDOWN P2: let's just error?
+        with set_cuda_backend("ffmpeg"):
+            got = self._to_frame_batch(self._decode_sequential(NASA_VIDEO.path, device))
+        ref = self._to_frame_batch(self._decode_sequential(NASA_VIDEO.path, device))
+        torch.testing.assert_close(got.data, ref.data, atol=0, rtol=0)
+
 
 # Small helpers to avoid having to always specify the same skip marks and decode_fn
 def _jpeg_param(*values):