fix(audio): drop resident safetensors pages after the TTS and ASR uploads (#845)

Co-authored-by: Federico Izzo <[email protected]>

src/core/hip/weight_upload.cpp

@@ -50,6 +50,7 @@ struct WeightUpload::State {
   struct Shard {
     platform::PlatformFile fd{platform::kInvalidFile};
     std::uint64_t size{0};
+    const void* mapping{nullptr};
   };
   struct Slot {
     void* buffer{nullptr};
@@ -212,6 +213,7 @@ std::unique_ptr<WeightUpload> WeightUpload::Create(
       return nullptr;
     }
     shard.size = info.size;
+    shard.mapping = region.data;
   }
   for (auto& slot : s.slots) {
     if (s.host_allocations.Allocate(&slot.buffer, kChunkBytes + kAlignment) !=
@@ -262,6 +264,23 @@ bool WeightUpload::Copy(std::uint32_t shard, std::uint64_t offset,
   return s.Status(error);
 }
 
+bool WeightUpload::CopyMapped(const void* source, std::size_t size,
+                              void* device, std::string* error) {
+  const auto address = reinterpret_cast<std::uintptr_t>(source);
+  for (std::size_t i = 0; i < state_->shards.size(); ++i) {
+    const auto& shard = state_->shards[i];
+    const auto base = reinterpret_cast<std::uintptr_t>(shard.mapping);
+    if (shard.mapping != nullptr && address >= base &&
+        address - base < shard.size) {
+      return Copy(static_cast<std::uint32_t>(i), address - base, size, device,
+                  error);
+    }
+  }
+  if (error)
+    *error = "weight upload source is outside its mappings";
+  return false;
+}
+
 bool WeightUpload::Finish(std::string* error) {
   auto& s = *state_;
   std::unique_lock lock(s.mutex);

src/core/hip/weight_upload.hpp

@@ -24,6 +24,9 @@ public:
 
   bool Copy(std::uint32_t shard, std::uint64_t offset, std::size_t size,
             void* device, std::string* error);
+  /// Resolve a pointer within the supplied mappings without reading its pages.
+  bool CopyMapped(const void* source, std::size_t size, void* device,
+                  std::string* error);
   /// Waits for every queued copy, including on failure.
   bool Finish(std::string* error);
 

src/core/platform/prefault.hpp

@@ -142,6 +142,27 @@ inline void PrefaultMappedRange(const void* data, std::size_t bytes) {
 #endif  // !__HIP_DEVICE_COMPILE__
 }
 
+/// Drop this process's resident pages of a read-only file mapping once its
+/// contents have been copied elsewhere (device upload done and synchronized).
+/// The mapping stays valid: a later read refaults the page from the file, so
+/// this is safe even if a host read remains. Never call it on a range that is
+/// hipHostRegister'ed, where the mapping is the GPU-visible weight memory.
+/// Best effort: failures are ignored, and it is a no-op on Windows.
+inline void ReleaseMappedPages(const void* data, std::size_t bytes) noexcept {
+#if !defined(__HIP_DEVICE_COMPILE__) && !defined(_WIN32)
+  const long page_size = ::sysconf(_SC_PAGESIZE);
+  if (!data || bytes == 0 || page_size <= 0)
+    return;
+  const auto address = reinterpret_cast<std::uintptr_t>(data);
+  const auto skip = address % static_cast<std::size_t>(page_size);
+  (void)::madvise(reinterpret_cast<void*>(address - skip), bytes + skip,
+                  MADV_DONTNEED);
+#else
+  static_cast<void>(data);
+  static_cast<void>(bytes);
+#endif
+}
+
 }  // namespace gufo::platform
 
 #endif  // GUFO_CORE_PLATFORM_PREFETCH_HPP_

src/models/qwen3_asr/hip/audio_encoder_runtime.cpp

@@ -12,6 +12,7 @@
 #include <array>
 #include <cstddef>
 #include <cstdint>
+#include <exception>
 #include <limits>
 #include <memory>
 #include <ranges>
@@ -22,7 +23,7 @@
 #include <utility>
 #include <vector>
 
-#include "src/core/mapped_prefetch.hpp"
+#include "src/core/hip/weight_upload.hpp"
 #include "src/models/qwen3_asr/hip/audio_ops.hpp"
 #include "src/models/qwen3_asr/hip/blas.hpp"
 #include "src/models/qwen3_asr/loader.hpp"
@@ -123,9 +124,9 @@ const Tensor& RequireTensor(const TensorStore& store, std::string_view name,
   return *tensor;
 }
 
-DeviceBuffer<hip_bfloat16> UploadBfloat16Async(
+DeviceBuffer<hip_bfloat16> UploadBfloat16FromFile(
     const TensorStore& store, std::string_view name,
-    std::span<const std::uint64_t> shape, hipStream_t stream) {
+    std::span<const std::uint64_t> shape, gufo::hip::WeightUpload& upload) {
   const Tensor& tensor = RequireTensor(store, name, shape);
   const auto elements = tensor.NumElements();
   if (!elements.has_value()) {
@@ -133,12 +134,26 @@ DeviceBuffer<hip_bfloat16> UploadBfloat16Async(
                              std::string(name));
   }
   DeviceBuffer<hip_bfloat16> result(*elements);
-  RequireHip(hipMemcpyAsync(result.get(), tensor.data, tensor.byte_count,
-                            hipMemcpyHostToDevice, stream),
-             "hipMemcpyAsync Qwen3-ASR weight");
+  std::string error;
+  const bool copied =
+      upload.CopyMapped(tensor.data, tensor.byte_count, result.get(), &error);
+  if (!copied) {
+    // Drain earlier copies before this local destination can be destroyed.
+    (void)upload.Finish(&error);
+    throw std::runtime_error(error);
+  }
   return result;
 }
 
+struct UploadFailureGuard {
+  std::unique_ptr<gufo::hip::WeightUpload>& upload;
+  int exceptions{std::uncaught_exceptions()};
+  ~UploadFailureGuard() {
+    if (std::uncaught_exceptions() > exceptions)
+      upload.reset();
+  }
+};
+
 std::string LayerName(std::size_t layer, std::string_view suffix) {
   return "thinker.audio_tower.layers." + std::to_string(layer) + "." +
          std::string(suffix);
@@ -175,9 +190,14 @@ struct AudioEncoderHipRuntime::Impl {
                    "hipblasSetStream Qwen3-ASR");
     RequireHipblas(hipblasSetAtomicsMode(blas, HIPBLAS_ATOMICS_NOT_ALLOWED),
                    "hipblasSetAtomicsMode Qwen3-ASR");
-    for (const auto& region : model.RegionsFor({"thinker.audio_tower."}))
-      core::PrefaultMappedRange(region.data, region.size);
+    BeginWeightUpload();
     LoadWeights();
+    LoadEncoderWeights();
+    FinishWeightUpload();
+    weight_upload.reset();
+    RequireHip(hipStreamSynchronize(stream),
+               "hipStreamSynchronize Qwen3-ASR audio weights");
+    model.ReleaseHostPages();
   }
 
   ~Impl() {
@@ -192,13 +212,28 @@ struct AudioEncoderHipRuntime::Impl {
   DeviceBuffer<hip_bfloat16> UploadBfloat16(
       const TensorStore& store, std::string_view name,
       std::span<const std::uint64_t> shape) {
-    return UploadBfloat16Async(store, name, shape, stream);
+    return UploadBfloat16FromFile(store, name, shape, *weight_upload);
+  }
+
+  void BeginWeightUpload() {
+    std::string error;
+    weight_upload =
+        gufo::hip::WeightUpload::Create(model.WeightUploadRegions(), &error);
+    if (!weight_upload)
+      throw std::runtime_error(error);
+  }
+
+  void FinishWeightUpload() {
+    std::string error;
+    if (!weight_upload->Finish(&error))
+      throw std::runtime_error(error);
   }
 
   DeviceBuffer<hip_bfloat16> UploadConvolution(const TensorStore& store,
                                                std::string_view name) {
     auto original = UploadBfloat16(
         store, name, std::array<std::uint64_t, 4>{480, 480, 3, 3});
+    FinishWeightUpload();
     DeviceBuffer<hip_bfloat16> packed(original.size());
     LaunchPackConvolutionWeights(original.get(), packed.get(), stream);
     RequireHip(hipGetLastError(), "pack Qwen3-ASR convolution weights");
@@ -238,14 +273,12 @@ struct AudioEncoderHipRuntime::Impl {
     compact.Reset(chunks * kConvTime * kHidden);
   }
 
-  void EnsureEncoderWeights() {
-    if (!layers.empty()) {
-      return;
-    }
+  void LoadEncoderWeights() {
     const TensorStore& store = *model.store;
     layers.reserve(24);
     for (std::size_t layer = 0; layer < 24; ++layer) {
       EncoderLayerWeights weights;
+      UploadFailureGuard drain_before_local_weights{weight_upload};
       weights.q_weight =
           UploadBfloat16(store, LayerName(layer, "self_attn.q_proj.weight"),
                          std::array<std::uint64_t, 2>{kHidden, kHidden});
@@ -414,7 +447,6 @@ struct AudioEncoderHipRuntime::Impl {
                                 hipMemcpyDeviceToHost, stream),
                  "hipMemcpyAsync Qwen3-ASR audio frontend");
     }
-    EnsureEncoderWeights();
     EnsureTransformerCapacity(tokens);
     const std::size_t hidden_elements = tokens * kHidden;
     RequireHip(hipMemcpyAsync(hidden.get(), compact.get(),
@@ -556,6 +588,7 @@ struct AudioEncoderHipRuntime::Impl {
   DeviceBuffer<float> ffn;
   DeviceBuffer<hip_bfloat16> ffn_bfloat16;
   DeviceBuffer<std::uint32_t> finite_flag{1U};
+  std::unique_ptr<gufo::hip::WeightUpload> weight_upload;
 };
 
 AudioEncoderHipRuntime::AudioEncoderHipRuntime(std::unique_ptr<Impl> impl)

src/models/qwen3_asr/hip/text_decoder_runtime.cpp

@@ -24,6 +24,7 @@
 #include <utility>
 #include <vector>
 
+#include "src/core/hip/weight_upload.hpp"
 #include "src/core/mapped_prefetch.hpp"
 #include "src/models/qwen3_asr/hip/blas.hpp"
 #include "src/models/qwen3_asr/hip/text_ops.hpp"
@@ -123,12 +124,15 @@ public:
   DeviceRegion(const DeviceRegion&) = delete;
   DeviceRegion& operator=(const DeviceRegion&) = delete;
 
-  [[nodiscard]] bool Initialize(MappedRegion region, hipStream_t stream) {
+  /// True when the GPU reads the file mapping itself (hipHostRegister
+  /// fallback): its host pages must then stay resident.
+  [[nodiscard]] bool host_mapped() const noexcept { return registered_; }
+
+  [[nodiscard]] bool Initialize(MappedRegion region, const LoadResult& model) {
     host_ = region.data;
     size_ = region.size;
     payload_offset_ = region.payload_offset;
-    core::PrefaultMappedRange(host_, size_);
-    return TryCopy(stream) || TryMap();
+    return TryCopy(model) || TryMap();
   }
 
   [[nodiscard]] const void* Resolve(const std::byte* pointer) const noexcept {
@@ -142,6 +146,7 @@ public:
 
 private:
   [[nodiscard]] bool TryMap() {
+    core::PrefaultMappedRange(host_, size_);
     if (hipHostRegister(const_cast<std::byte*>(host_), size_,
                         hipHostRegisterMapped | hipHostRegisterReadOnly) !=
         hipSuccess) {
@@ -158,7 +163,7 @@ private:
     return false;
   }
 
-  [[nodiscard]] bool TryCopy(hipStream_t stream) {
+  [[nodiscard]] bool TryCopy(const LoadResult& model) {
     constexpr std::size_t kLoadAlignment = 16U;
     const std::size_t shift =
         (kLoadAlignment - (payload_offset_ % kLoadAlignment)) % kLoadAlignment;
@@ -168,8 +173,16 @@ private:
     }
     owns_device_ = true;
     device_ = static_cast<std::byte*>(allocation_) + shift;
-    if (hipMemcpyAsync(device_, host_, size_, hipMemcpyHostToDevice, stream) ==
-        hipSuccess) {
+    // Never expose the file mapping to HIP's pageable-copy registration cache:
+    // its HMM restore work can repopulate pages after MADV_DONTNEED.
+    const bool copied = [&] {
+      std::string error;
+      auto upload =
+          gufo::hip::WeightUpload::Create(model.WeightUploadRegions(), &error);
+      return upload && upload->CopyMapped(host_, size_, device_, &error) &&
+             upload->Finish(&error);
+    }();
+    if (copied) {
       return true;
     }
     (void)hipFree(allocation_);
@@ -302,13 +315,20 @@ struct TextDecoderHipRuntime::Impl {
     for (const MappedRegion region :
          model.RegionsFor({"thinker.model.", "thinker.lm_head."})) {
       auto device_region = std::make_unique<DeviceRegion>();
-      if (!device_region->Initialize(region, stream)) {
+      if (!device_region->Initialize(region, model)) {
         throw std::runtime_error(
             "cannot make Qwen3-ASR safetensors GPU-visible");
       }
       weight_regions.push_back(std::move(device_region));
     }
     LoadWeights();
+    // Complete any derived weight work before dropping the host pages.
+    RequireHip(hipStreamSynchronize(stream),
+               "hipStreamSynchronize Qwen3-ASR text weights");
+    if (std::ranges::none_of(weight_regions, [](const auto& region) {
+          return region->host_mapped();
+        }))
+      model.ReleaseHostPages();
   }
 
   ~Impl() {

src/models/qwen3_asr/loader.cpp

@@ -13,6 +13,7 @@
 
 #include "src/core/json.hpp"
 #include "src/core/platform/mapped_file.hpp"
+#include "src/core/platform/prefault.hpp"
 
 namespace gufo::models::qwen3_asr {
 namespace {
@@ -28,10 +29,11 @@ public:
   MappedFile& operator=(MappedFile&&) = delete;
 
   [[nodiscard]] bool Open(const std::filesystem::path& path) {
-    if (!mapped_.Open(path, platform::MapHint::kSequential, &open_error_))
+    mapped_ = platform::ReadOnlyMappedFile::Open(path, &open_error_);
+    if (!mapped_.valid())
       return false;
-    data_ = mapped_.data;
-    size_ = mapped_.size;
+    data_ = const_cast<std::uint8_t*>(mapped_.data());
+    size_ = mapped_.size();
     return true;
   }
 
@@ -43,6 +45,9 @@ public:
     return static_cast<const std::byte*>(data_);
   }
   [[nodiscard]] std::size_t size() const noexcept { return size_; }
+  platform::PlatformFile file_descriptor() const {
+    return mapped_.PlatformHandle();
+  }
 
 private:
   void Reset() noexcept {
@@ -51,7 +56,7 @@ private:
     size_ = 0;
   }
 
-  platform::RawMappedFile mapped_;
+  platform::ReadOnlyMappedFile mapped_;
   std::string open_error_;
   void* data_{nullptr};
   std::size_t size_{0};
@@ -156,6 +161,7 @@ bool LoadSafetensors(const std::filesystem::path& path, TensorStore* store,
       .data = mapping->data(),
       .size = mapping->size(),
       .payload_offset = static_cast<std::size_t>(payload_offset),
+      .file_descriptor = mapping->file_descriptor(),
   });
   mappings->push_back(std::move(mapping));
   return true;
@@ -163,6 +169,18 @@ bool LoadSafetensors(const std::filesystem::path& path, TensorStore* store,
 
 }  // namespace
 
+void LoadResult::ReleaseHostPages() const noexcept {
+  for (const auto& region : mapped_regions)
+    platform::ReleaseMappedPages(region.data, region.size);
+}
+
+std::vector<core::GgufMappedRegion> LoadResult::WeightUploadRegions() const {
+  std::vector<core::GgufMappedRegion> result;
+  for (const auto& region : mapped_regions)
+    result.push_back({region.data, region.size, region.file_descriptor});
+  return result;
+}
+
 std::vector<MappedRegion> LoadResult::RegionsFor(
     std::initializer_list<std::string_view> prefixes) const {
   std::vector<MappedRegion> result;

src/models/qwen3_asr/loader.hpp

@@ -8,6 +8,7 @@
 #include <string_view>
 #include <vector>
 
+#include "src/core/gguf_reader.hpp"
 #include "src/models/qwen3_asr/config.hpp"
 #include "src/models/qwen3_asr/tensor.hpp"
 
@@ -17,6 +18,7 @@ struct MappedRegion {
   const std::byte* data{nullptr};
   std::size_t size{0};
   std::size_t payload_offset{0};
+  platform::PlatformFile file_descriptor{platform::kInvalidFile};
 };
 
 struct LoadResult {
@@ -31,6 +33,11 @@ struct LoadResult {
   /// the original payload alignment for GPU upload.
   [[nodiscard]] std::vector<MappedRegion> RegionsFor(
       std::initializer_list<std::string_view> prefixes) const;
+
+  /// Drops the resident pages of every mapped shard once the caller's weights
+  /// live in device memory. Pointers in `store` stay valid (pages refault).
+  void ReleaseHostPages() const noexcept;
+  [[nodiscard]] std::vector<core::GgufMappedRegion> WeightUploadRegions() const;
 };
 
 [[nodiscard]] LoadResult LoadModelDirectory(const std::string& model_dir);

src/models/qwen3_tts/hip/speaker_encoder_runtime.cpp

@@ -252,6 +252,8 @@ struct SpeakerEncoderHipRuntime::Impl {
     RequireEncoderHip(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking),
                       "hipStreamCreate Qwen3-TTS speaker encoder");
     LoadWeights();
+    // Weights were converted into owned device buffers by hipMemcpy.
+    model.ReleaseHostPages();
   }
 
   ~Impl() {

src/models/qwen3_tts/hip/speech_decoder_runtime.cpp

@@ -25,6 +25,7 @@
 #include <utility>
 #include <vector>
 
+#include "src/core/hip/weight_upload.hpp"
 #include "src/core/mapped_prefetch.hpp"
 #include "src/models/qwen3_tts/hip/speech_decoder_ops.hpp"
 #include "src/models/qwen3_tts/loader.hpp"
@@ -124,15 +125,19 @@ public:
   DeviceRegion(const DeviceRegion&) = delete;
   DeviceRegion& operator=(const DeviceRegion&) = delete;
 
-  bool Initialize(MappedRegion region, std::string* error) {
+  /// True when the GPU reads the file mapping itself (hipHostRegister
+  /// fallback): its host pages must then stay resident.
+  [[nodiscard]] bool host_mapped() const noexcept { return registered_; }
+
+  bool Initialize(MappedRegion region, const LoadResult& model,
+                  std::string* error) {
     host_ = region.data;
     size_ = region.size;
     payload_offset_ = region.payload_offset;
-    core::PrefaultMappedRange(host_, size_);
     // The supported gfx1151 APU shares physical memory with the CPU, but
     // hipMalloc uses coarse-grained pages that the GPU caches and streams
     // faster than host-registered safetensors pages.
-    if (TryCopy() || TryMap())
+    if (TryCopy(model) || TryMap())
       return true;
     SetError(error,
              "cannot make Qwen3-TTS speech-tokenizer weights GPU-visible");
@@ -150,6 +155,7 @@ public:
 
 private:
   bool TryMap() {
+    core::PrefaultMappedRange(host_, size_);
     if (hipHostRegister(const_cast<std::byte*>(host_), size_,
                         hipHostRegisterMapped | hipHostRegisterReadOnly) !=
         hipSuccess) {
@@ -166,7 +172,7 @@ private:
     return false;
   }
 
-  bool TryCopy() {
+  bool TryCopy(const LoadResult& model) {
     constexpr std::size_t kLoadAlignment = 16;
     const std::size_t shift =
         (kLoadAlignment - (payload_offset_ % kLoadAlignment)) % kLoadAlignment;
@@ -177,7 +183,16 @@ private:
     allocation_ = allocation;
     owns_device_ = true;
     device_ = static_cast<std::byte*>(allocation) + shift;
-    if (hipMemcpy(device_, host_, size_, hipMemcpyHostToDevice) == hipSuccess) {
+    // Never expose the file mapping to HIP's pageable-copy registration cache:
+    // its HMM restore work can repopulate pages after MADV_DONTNEED.
+    const bool copied = [&] {
+      std::string error;
+      auto upload =
+          gufo::hip::WeightUpload::Create(model.WeightUploadRegions(), &error);
+      return upload && upload->CopyMapped(host_, size_, device_, &error) &&
+             upload->Finish(&error);
+    }();
+    if (copied) {
       return true;
     }
     (void)hipFree(allocation_);
@@ -475,13 +490,15 @@ struct SpeechDecoderHipRuntime::Impl {
     }
     RequireHip(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking),
                "hipStreamCreateWithFlags");
-    if (!speech_weights.Initialize(decoder_regions.front(), nullptr)) {
+    if (!speech_weights.Initialize(decoder_regions.front(), model, nullptr)) {
       throw std::runtime_error(
           "cannot initialize Qwen3-TTS speech weight mapping");
     }
     LoadWeights();
     RequireHip(hipStreamSynchronize(stream),
                "synchronize Qwen3-TTS speech weights");
+    if (!speech_weights.host_mapped())
+      model.ReleaseHostPages();
   }
 
   ~Impl() {

src/models/qwen3_tts/hip/speech_encoder_runtime.cpp

@@ -128,6 +128,8 @@ struct SpeechEncoderHipRuntime::Impl {
     RequireEncoderHip(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking),
                       "hipStreamCreate Qwen3-TTS speech encoder");
     LoadWeights();
+    // Weights were converted into owned device buffers by hipMemcpy.
+    model.ReleaseHostPages();
   }
 
   ~Impl() {

src/models/qwen3_tts/hip/talker_runtime.cpp

@@ -28,6 +28,7 @@
 #include <utility>
 #include <vector>
 
+#include "src/core/hip/weight_upload.hpp"
 #include "src/core/mapped_prefetch.hpp"
 #include "src/models/qwen/hip/ops.hpp"
 #include "src/models/qwen3_tts/hip/talker_ops.hpp"
@@ -131,16 +132,20 @@ public:
   DeviceRegion(const DeviceRegion&) = delete;
   DeviceRegion& operator=(const DeviceRegion&) = delete;
 
-  bool Initialize(MappedRegion region, std::string* error) {
+  /// True when the GPU reads the file mapping itself (hipHostRegister
+  /// fallback): its host pages must then stay resident.
+  [[nodiscard]] bool host_mapped() const noexcept { return registered_; }
+
+  bool Initialize(MappedRegion region, const LoadResult& model,
+                  std::string* error) {
     host_ = region.data;
     size_ = region.size;
     payload_offset_ = region.payload_offset;
-    core::PrefaultMappedRange(host_, size_);
     // gfx1151 shares one physical memory pool with the host, but
     // device-resident weights still use coarse-grained pages the GPU caches and
     // streams far faster than host-registered pages, so decode prefers the
     // copy.
-    if (TryCopy() || TryMap())
+    if (TryCopy(model) || TryMap())
       return true;
     SetError(error, "cannot make Qwen3-TTS safetensors GPU-visible");
     return false;
@@ -157,6 +162,7 @@ public:
 
 private:
   bool TryMap() {
+    core::PrefaultMappedRange(host_, size_);
     const hipError_t registered =
         hipHostRegister(const_cast<std::byte*>(host_), size_,
                         hipHostRegisterMapped | hipHostRegisterReadOnly);
@@ -174,7 +180,7 @@ private:
     return false;
   }
 
-  bool TryCopy() {
+  bool TryCopy(const LoadResult& model) {
     // Safetensors headers are an arbitrary length, so the payload rarely starts
     // on a 16-byte boundary. Padding the copy by that much lets every tensor
     // land where the widest packed loads are legal.
@@ -188,7 +194,16 @@ private:
     allocation_ = allocation;
     owns_device_ = true;
     device_ = static_cast<std::byte*>(allocation) + shift;
-    if (hipMemcpy(device_, host_, size_, hipMemcpyHostToDevice) == hipSuccess) {
+    // Never expose the file mapping to HIP's pageable-copy registration cache:
+    // its HMM restore work can repopulate pages after MADV_DONTNEED.
+    const bool copied = [&] {
+      std::string error;
+      auto upload =
+          gufo::hip::WeightUpload::Create(model.WeightUploadRegions(), &error);
+      return upload && upload->CopyMapped(host_, size_, device_, &error) &&
+             upload->Finish(&error);
+    }();
+    if (copied) {
       return true;
     }
     (void)hipFree(allocation_);
@@ -395,7 +410,7 @@ struct TalkerHipRuntime::Impl {
         predictor_token(1) {
     const auto talker_regions = model.RegionsFor("talker.");
     if (talker_regions.size() != 1 ||
-        !main_weights.Initialize(talker_regions.front(), nullptr)) {
+        !main_weights.Initialize(talker_regions.front(), model, nullptr)) {
       throw std::runtime_error(
           "cannot initialize Qwen3-TTS GPU weight mapping");
     }
@@ -410,6 +425,9 @@ struct TalkerHipRuntime::Impl {
     RequireHipblas(hipblasSetAtomicsMode(blas, HIPBLAS_ATOMICS_NOT_ALLOWED),
                    "hipblasSetAtomicsMode");
     LoadWeights();
+    // The bounded uploads above have finished.
+    if (!main_weights.host_mapped())
+      model.ReleaseHostPages();
   }
 
   ~Impl() {

src/models/qwen3_tts/loader.cpp

@@ -13,6 +13,7 @@
 
 #include "src/core/json.hpp"
 #include "src/core/platform/mapped_file.hpp"
+#include "src/core/platform/prefault.hpp"
 
 namespace gufo::models::qwen3_tts {
 namespace {
@@ -28,17 +29,8 @@ public:
   MappedFile(const MappedFile&) = delete;
   MappedFile& operator=(const MappedFile&) = delete;
 
-  MappedFile(MappedFile&& other) noexcept
-      : base_(std::exchange(other.base_, nullptr)), size_(other.size_) {}
-
-  MappedFile& operator=(MappedFile&& other) noexcept {
-    if (this != &other) {
-      Unmap();
-      base_ = std::exchange(other.base_, nullptr);
-      size_ = other.size_;
-    }
-    return *this;
-  }
+  MappedFile(MappedFile&&) = delete;
+  MappedFile& operator=(MappedFile&&) = delete;
 
   [[nodiscard]] const std::string& open_error() const noexcept {
     return open_error_;
@@ -46,14 +38,18 @@ public:
 
   void* base() const { return base_; }
   std::size_t size() const { return size_; }
+  platform::PlatformFile file_descriptor() const {
+    return mapped_.PlatformHandle();
+  }
 
   /// Maps the file read-only. Returns false on failure; the reason is
   /// then available through open_error().
   bool Open(const std::filesystem::path& path) {
-    if (!mapped_.Open(path, platform::MapHint::kNone, &open_error_))
+    mapped_ = platform::ReadOnlyMappedFile::Open(path, &open_error_);
+    if (!mapped_.valid())
       return false;
-    base_ = mapped_.data;
-    size_ = mapped_.size;
+    base_ = const_cast<std::uint8_t*>(mapped_.data());
+    size_ = mapped_.size();
     return true;
   }
 
@@ -64,7 +60,7 @@ private:
     size_ = 0;
   }
 
-  platform::RawMappedFile mapped_;
+  platform::ReadOnlyMappedFile mapped_;
   std::string open_error_;
   void* base_ = nullptr;
   std::size_t size_ = 0;
@@ -177,13 +173,26 @@ bool LoadSafetensorsFile(const std::filesystem::path& path, TensorStore* store,
   mapped_regions->push_back(
       {.data = bytes,
        .size = static_cast<std::size_t>(mapped->size()),
-       .payload_offset = static_cast<std::size_t>(payload_offset)});
+       .payload_offset = static_cast<std::size_t>(payload_offset),
+       .file_descriptor = mapped->file_descriptor()});
   mappings->push_back(std::move(mapped));
   return true;
 }
 
 }  // namespace
 
+void LoadResult::ReleaseHostPages() const noexcept {
+  for (const auto& region : mapped_regions)
+    platform::ReleaseMappedPages(region.data, region.size);
+}
+
+std::vector<core::GgufMappedRegion> LoadResult::WeightUploadRegions() const {
+  std::vector<core::GgufMappedRegion> result;
+  for (const auto& region : mapped_regions)
+    result.push_back({region.data, region.size, region.file_descriptor});
+  return result;
+}
+
 std::vector<MappedRegion> LoadResult::RegionsFor(
     std::string_view prefix) const {
   std::vector<MappedRegion> result;

src/models/qwen3_tts/loader.hpp

@@ -8,6 +8,7 @@
 #include <string_view>
 #include <vector>
 
+#include "src/core/gguf_reader.hpp"
 #include "src/models/qwen3_tts/config.hpp"
 #include "src/models/qwen3_tts/tensor.hpp"
 
@@ -19,6 +20,7 @@ struct MappedRegion {
   // Byte offset of the safetensors payload. Tensor offsets are relative to it,
   // so a device copy can restore vectorized-load alignment by shifting here.
   std::size_t payload_offset = 0;
+  platform::PlatformFile file_descriptor = platform::kInvalidFile;
 };
 
 /// Result of loading the model directory. `tensors` payload pointers stay
@@ -37,6 +39,11 @@ struct LoadResult {
   /// alignment is preserved when these ranges are copied to device memory.
   [[nodiscard]] std::vector<MappedRegion> RegionsFor(
       std::string_view prefix) const;
+
+  /// Drops the resident pages of every mapped file once the caller's weights
+  /// live in device memory. Pointers in `store` stay valid (pages refault).
+  void ReleaseHostPages() const noexcept;
+  [[nodiscard]] std::vector<core::GgufMappedRegion> WeightUploadRegions() const;
 };
 
 /// Loads a quality-qualified Qwen3-TTS 12.5 Hz 1.7B model directory.

tests/core/hip_weight_upload_test.cpp

@@ -1,5 +1,6 @@
 #include <hip/hip_runtime.h>
 
+#include <algorithm>
 #include <array>
 #include <cstdint>
 #include <cstdio>
@@ -55,7 +56,7 @@ int main() {
 
   std::array<gufo::core::GgufMappedRegion, 2> regions;
   for (std::size_t i = 0; i < paths.size(); ++i) {
-    regions[i] = {nullptr, source[i].size(),
+    regions[i] = {source[i].data(), source[i].size(),
                   gufo::platform::OpenReadFile(paths[i])};
     Require(gufo::platform::FileValid(regions[i].file_descriptor),
             "open original shard");
@@ -95,12 +96,18 @@ int main() {
   std::string error;
   auto upload = WeightUpload::Create(regions, &error);
   Require(upload != nullptr, error);
+  // These addresses identify file offsets; their contents must never supply
+  // upload bytes or enter HIP's pageable host-registration cache.
+  for (auto& buffer : source)
+    std::fill(buffer.begin(), buffer.end(), 0);
   for (const auto& s : slices) {
-    Require(
-        upload->Copy(s.shard, s.offset, s.size,
-                     static_cast<char*>(device.data) + s.destination, &error),
-        error);
+    Require(upload->CopyMapped(source[s.shard].data() + s.offset, s.size,
+                               static_cast<char*>(device.data) + s.destination,
+                               &error),
+            error);
   }
+  Require(!upload->CopyMapped(&error, 1, device.data, &error),
+          "reject a pointer outside the source mappings");
   Require(upload->Copy(1, source[1].size(), 0, nullptr, &error),
           "empty range at EOF");
   Require(upload->Finish(&error), error);

tests/core/platform_prefault_test.cpp

@@ -1,10 +1,68 @@
 #include <cassert>
 #include <cstddef>
+#include <cstdio>
+#include <fstream>
+#include <sstream>
+#include <string>
 #include <system_error>
 
 #include "src/core/platform/prefault.hpp"
 
+#if !defined(_WIN32)
+// mincore reports page-cache residency, which MADV_DONTNEED deliberately
+// preserves. smaps measures this process's resident file-mapping pages.
+std::size_t ResidentKiB(const void* memory) {
+  std::ifstream smaps("/proc/self/smaps");
+  assert(smaps);
+  std::string line;
+  bool in_mapping = false;
+  while (std::getline(smaps, line)) {
+    unsigned long begin = 0, end = 0;
+    if (std::sscanf(line.c_str(), "%lx-%lx", &begin, &end) == 2) {
+      const auto address = reinterpret_cast<std::uintptr_t>(memory);
+      in_mapping = begin <= address && address < end;
+    } else if (in_mapping && line.starts_with("Rss:")) {
+      std::istringstream fields(line.substr(4));
+      std::size_t kib = 0;
+      fields >> kib;
+      assert(fields);
+      return kib;
+    }
+  }
+  assert(false && "mapping missing from smaps");
+  return 0;
+}
+
+void CheckReleaseMappedPages() {
+  const auto page_size = ::sysconf(_SC_PAGESIZE);
+  assert(page_size > 0);
+  const auto bytes = static_cast<std::size_t>(page_size) * 4;
+  std::FILE* file = std::tmpfile();
+  assert(file);
+  for (std::size_t i = 0; i < bytes; ++i)
+    assert(std::fputc(static_cast<int>(i % 251), file) != EOF);
+  assert(std::fflush(file) == 0);
+  const auto* memory = static_cast<const unsigned char*>(
+      ::mmap(nullptr, bytes, PROT_READ, MAP_PRIVATE, ::fileno(file), 0));
+  assert(memory != MAP_FAILED);
+  gufo::platform::PrefaultMappedRange(memory, bytes);
+  assert(ResidentKiB(memory) == bytes / 1024);
+  // Include the first partial page and the final byte of the mapping.
+  gufo::platform::ReleaseMappedPages(memory + 1, bytes - 1);
+  assert(ResidentKiB(memory) == 0);
+  for (std::size_t i = 0; i < bytes; ++i)
+    assert(memory[i] == i % 251);
+  assert(ResidentKiB(memory) == bytes / 1024);
+  assert(::munmap(const_cast<unsigned char*>(memory), bytes) == 0);
+  assert(std::fclose(file) == 0);
+}
+#endif
+
 int main() {
+  gufo::platform::ReleaseMappedPages(nullptr, 0);
+#if !defined(_WIN32)
+  CheckReleaseMappedPages();
+#endif
   constexpr std::size_t bytes = 32ULL << 20;
 #if defined(_WIN32)
   auto* memory = static_cast<char*>(