From cc8872a72bbd172ea24b3eb6d310455057fe9060 Mon Sep 17 00:00:00 2001
From: Xget7 <juanieltupa@gmail.com>
Date: Tue, 20 Oct 2026 01:20:22 -0300
Subject: [PATCH 1/8] Give the Metal frame-slot count, sort key width and draw
 arguments one home in layout/

The slot count was declared three times (MetalVisibilityPass::kSlots,
MetalFrameCompute::kSlots, MetalSplatRenderer::kFramesInFlight) and the
radix sort owned the key width enum, so cull/ included sort/ and raster/
included cull/ only to read constants.

kFramesInFlight, KeyBits and kDrawBatches/kDrawArgumentBytes now live in
layout/MetalFrameSlots.h, MetalSortKeys.h and MetalDrawArguments.h. Every
per-slot array derives from kFramesInFlight, the cull -> sort and
raster -> cull includes are gone, and the test-local copies of the slot
count in MetalSurfaceTest are dropped. MetalFrameLoopTest keeps its own
count because it acquires exactly two slots by hand.

Behaviour and shader text are unchanged.
---
 .../rendering/metal/MetalSplatRenderer.h      |  3 +--
 .../rendering/metal/MetalSplatRenderer.mm     |  2 +-
 .../metal/cull/MetalVisibilityPass.h          | 24 +++++++++----------
 .../metal/cull/MetalVisibilityPass.mm         | 12 +++++-----
 .../rendering/metal/frame/MetalFrameCompute.h | 13 +++++-----
 .../metal/frame/MetalFrameCompute.mm          | 12 +++++-----
 .../metal/layout/MetalDrawArguments.h         | 14 +++++++++++
 .../rendering/metal/layout/MetalFrameSlots.h  | 11 +++++++++
 .../rendering/metal/layout/MetalSortKeys.h    | 11 +++++++++
 .../metal/raster/MetalSplatPipeline.mm        |  6 ++---
 .../rendering/metal/sort/MetalRadixSort.h     |  4 ++--
 .../metal/cull/MetalVisibilityPassTest.mm     | 18 +++++++-------
 .../metal/frame/MetalFrameComputeTest.mm      | 12 +++++-----
 .../metal/presentation/MetalSurfaceTest.mm    |  5 ++--
 .../metal/raster/MetalSplatPipelineTest.mm    |  2 +-
 .../metal/sort/MetalRadixSortTest.mm          |  6 ++---
 .../tests/support/MetalVisibilityRig.h        |  2 +-
 17 files changed, 94 insertions(+), 63 deletions(-)
 create mode 100644 packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalDrawArguments.h
 create mode 100644 packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalFrameSlots.h
 create mode 100644 packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalSortKeys.h

diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h
index 0e1d7b0..c246a92 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h
@@ -12,6 +12,7 @@
 #include "rendering/metal/device/MetalContext.h"
 #include "rendering/metal/diagnostics/MetalFrameCapture.h"
 #include "rendering/metal/frame/MetalFrameCompute.h"
+#include "rendering/metal/layout/MetalFrameSlots.h"
 #include "rendering/metal/presentation/MetalFrameLoop.h"
 #include "rendering/metal/presentation/MetalSurface.h"
 #include "rendering/metal/raster/MetalSplatPipeline.h"
@@ -102,8 +103,6 @@ class MetalSplatRenderer final : public SplatRenderer {
   uint32_t takePresentTimes(std::vector<int64_t>* times) override;
   const std::string& deviceDescription() const override { return context_->description(); }
 
-  static constexpr uint32_t kFramesInFlight = MetalFrameCompute::kSlots;
-
  private:
   MetalSplatRenderer() = default;
 
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm
index e562a0b..efacea0 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm
@@ -87,7 +87,7 @@
   policy.fallback.lodErrorPixels = 1.0f;
   policy.fallback.alphaThreshold = kQuadMinAlpha;
   policy.fallback.subpixelThreshold = kMinPixelRadius;
-  policy.fallback.sortDepth = compute_->depthBits() == MetalRadixSort::KeyBits::Low16
+  policy.fallback.sortDepth = compute_->depthBits() == KeyBits::Low16
                                   ? SortKeyBits::low16
                                   : SortKeyBits::full32;
   policy.fallback.enableFrustumCulling = true;
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.h
index 189d912..e0b4adf 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.h
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.h
@@ -5,8 +5,10 @@
 #include <array>
 #include <cstdint>
 
+#include "rendering/metal/layout/MetalDrawArguments.h"
+#include "rendering/metal/layout/MetalFrameSlots.h"
 #include "rendering/metal/layout/MetalShaderTypes.h"
-#include "rendering/metal/sort/MetalRadixSort.h"
+#include "rendering/metal/layout/MetalSortKeys.h"
 #include "splatkit/rendering/SplatRenderer.h"
 
 namespace splatkit {
@@ -20,14 +22,14 @@ class MetalVisibilityPass {
   // Both widths sort by linear view depth: Full32 by its float bits, Low16 quantized between
   // finite forward-Z Mat4::perspective planes. Shader key width and radix passes agree.
   bool create(id<MTLDevice> device, id<MTLLibrary> library,
-              MetalRadixSort::KeyBits depthBits = MetalRadixSort::KeyBits::Full32);
+              KeyBits depthBits = KeyBits::Full32);
   // Rebuilds the visibility pipelines for a new key width, keeping every reserved buffer
   // and the radix scratch. The key width and the shader's quantized flag must agree, so
   // both move together. Render thread, while idle. On failure the previous key width is
   // restored.
-  bool reconfigure(MetalRadixSort::KeyBits depthBits);
+  bool reconfigure(KeyBits depthBits);
   // The key width the pipelines were built for.
-  MetalRadixSort::KeyBits depthBits() const { return depthBits_; }
+  KeyBits depthBits() const { return depthBits_; }
   // Space for every source splat, or an explicitly bounded GPU index list when
   // activeCapacity is nonzero. The latter must only be used with indexed encode.
   // Failure preserves the previous allocation. The caller reserves its sort for
@@ -60,10 +62,6 @@ class MetalVisibilityPass {
     return *static_cast<const uint32_t*>(count_[slot].contents);
   }
 
-  static constexpr uint32_t kSlots = 2;
-  static constexpr uint32_t kDrawBatches = 7;
-  static constexpr uint32_t kDrawArgumentBytes = sizeof(MTLDrawPrimitivesIndirectArguments);
-
  private:
   static constexpr uint32_t kThreads = 256;
   static constexpr uint32_t kMaxRanges = 65536;
@@ -76,15 +74,15 @@ class MetalVisibilityPass {
   std::array<id<MTLComputePipelineState>, kShDegrees> visibility_{};
   std::array<id<MTLComputePipelineState>, kShDegrees> indexedVisibility_{};
   id<MTLComputePipelineState> prepareDraw_ = nil;
-  MetalRadixSort::KeyBits depthBits_ = MetalRadixSort::KeyBits::Full32;
+  KeyBits depthBits_ = KeyBits::Full32;
   uint32_t capacity_ = 0;
   uint32_t sourceCapacity_ = 0;
   bool indexedOnly_ = false;
   id<MTLBuffer> projected_ = nil;
-  std::array<id<MTLBuffer>, kSlots> count_{};
-  std::array<id<MTLBuffer>, kSlots> drawArguments_{};
-  std::array<id<MTLBuffer>, kSlots> ranges_{};
-  std::array<id<MTLBuffer>, kSlots> rangeStarts_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> count_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> drawArguments_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> ranges_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> rangeStarts_{};
 };
 
 }  // namespace splatkit
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.mm
index b0264a6..7be179c 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.mm
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.mm
@@ -9,13 +9,13 @@
 namespace splatkit {
 
 bool MetalVisibilityPass::create(id<MTLDevice> device, id<MTLLibrary> library,
-                                 MetalRadixSort::KeyBits depthBits) {
+                                 KeyBits depthBits) {
   device_ = device;
   library_ = library;
   depthBits_ = depthBits;
   if (!createPipelines()) return false;
   bool ok = true;
-  for (uint32_t slot = 0; slot < kSlots; ++slot) {
+  for (uint32_t slot = 0; slot < kFramesInFlight; ++slot) {
     count_[slot] = metal::buffer(device, sizeof(uint32_t));
     drawArguments_[slot] = metal::buffer(device, kDrawBatches * kDrawArgumentBytes);
     ranges_[slot] = metal::buffer(device, size_t{kMaxRanges} * sizeof(SplatRenderer::Range));
@@ -27,9 +27,9 @@
 }
 
 // Only the pipelines depend on the key width; the buffers and the radix scratch are reused.
-bool MetalVisibilityPass::reconfigure(MetalRadixSort::KeyBits depthBits) {
+bool MetalVisibilityPass::reconfigure(KeyBits depthBits) {
   if (device_ == nil || library_ == nil) return false;
-  const MetalRadixSort::KeyBits previousBits = depthBits_;
+  const KeyBits previousBits = depthBits_;
   depthBits_ = depthBits;
   if (!createPipelines()) {
     depthBits_ = previousBits;
@@ -39,7 +39,7 @@
 }
 
 bool MetalVisibilityPass::createPipelines() {
-  const bool quantized = depthBits_ == MetalRadixSort::KeyBits::Low16;
+  const bool quantized = depthBits_ == KeyBits::Low16;
   bool ok = true;
   for (int degree = 0; degree < kShDegrees; ++degree) {
     MTLFunctionConstantValues* constants = [MTLFunctionConstantValues new];
@@ -98,7 +98,7 @@
   const bool indexed = indices != nil;
   if (indexedOnly_ && !indexed) return false;
   if (indexed != (activeCount != nil)) return false;
-  if (slot >= kSlots || rangeCount > kMaxRanges || capacity_ == 0 ||
+  if (slot >= kFramesInFlight || rangeCount > kMaxRanges || capacity_ == 0 ||
       (rangeCount > 0 && ranges == nullptr))
     return false;
   const size_t sortBytes = size_t{capacity_} * sizeof(uint32_t);
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h
index f48a622..ac1d907 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h
@@ -9,6 +9,7 @@
 #include <string>
 
 #include "rendering/metal/cull/MetalVisibilityPass.h"
+#include "rendering/metal/layout/MetalFrameSlots.h"
 #include "rendering/metal/presentation/MetalFrameLoop.h"
 #include "rendering/metal/selection/MetalLodSelection.h"
 #include "rendering/metal/sort/MetalRadixSort.h"
@@ -25,8 +26,6 @@ namespace splatkit {
 // frame loop must outlive it. Encode only; draw() commits what is returned.
 class MetalFrameCompute {
  public:
-  static constexpr uint32_t kSlots = MetalVisibilityPass::kSlots;
-
   MetalFrameCompute(id<MTLDevice> device, id<MTLCommandQueue> queue, id<MTLLibrary> library,
                     MetalFrameLoop& frameLoop);
 
@@ -38,7 +37,7 @@ class MetalFrameCompute {
   // gpuSort() false.
   bool create();
   bool gpuSort() const { return gpuSort_; }
-  MetalRadixSort::KeyBits depthBits() const { return visibility_.depthBits(); }
+  KeyBits depthBits() const { return visibility_.depthBits(); }
 
   // The sort key width and the LOD threshold and splat limit of `policy`. Waits for the frames
   // in flight when something changes. False, with `reason`, when the GPU cannot order or the
@@ -56,7 +55,7 @@ class MetalFrameCompute {
   // The slot the frame being encoded uses: the per-slot uniforms, counts and readbacks of a
   // frame the GPU may still be reading are never touched. advance() moves to the next slot once
   // the frame was submitted.
-  uint32_t slot() const { return static_cast<uint32_t>(frames_ % kSlots); }
+  uint32_t slot() const { return static_cast<uint32_t>(frames_ % kFramesInFlight); }
   void advance() { ++frames_; }
 
   // Writes the camera of `frame` into the slot's uniforms, for a frame drawn at `extent`.
@@ -86,19 +85,19 @@ class MetalFrameCompute {
 
  private:
   bool createOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort,
-                      MetalRadixSort::KeyBits bits) const;
+                      KeyBits bits) const;
   void adaptLodThreshold();
 
   id<MTLDevice> device_ = nil;
   id<MTLCommandQueue> queue_ = nil;
   id<MTLLibrary> library_ = nil;
   MetalFrameLoop& frameLoop_;
-  std::array<id<MTLBuffer>, kSlots> uniforms_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> uniforms_{};
   MetalVisibilityPass visibility_;
   MetalRadixSort sort_;
   bool gpuSort_ = false;
   std::unique_ptr<MetalLodSelection> lod_;
-  std::array<id<MTLBuffer>, kSlots> lodReadback_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> lodReadback_{};
   float lodErrorPixels_ = 1.0f;
   uint32_t lodSplatLimit_ = 0;
   // The threshold selection runs at: the policy's, raised while a cut would exceed the splat
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm
index 83e2fde..fdd2ccd 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm
@@ -37,7 +37,7 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
     : device_(device), queue_(queue), library_(library), frameLoop_(frameLoop) {}
 
 bool MetalFrameCompute::createOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort,
-                                       MetalRadixSort::KeyBits bits) const {
+                                       KeyBits bits) const {
   return visibility.create(device_, library_, bits) && sort.create(device_, library_);
 }
 
@@ -46,7 +46,7 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
     uniform = metal::buffer(device_, sizeof(CameraUniform));
     if (uniform == nil) return false;
   }
-  gpuSort_ = createOrdering(visibility_, sort_, MetalRadixSort::KeyBits::Full32);
+  gpuSort_ = createOrdering(visibility_, sort_, KeyBits::Full32);
   LOGI("Metal depth keys: %u bits, %u radix passes (uint32 scratch)",
        static_cast<uint32_t>(depthBits()), static_cast<uint32_t>(depthBits()) / 8);
   if (!gpuSort_) LOGW("GPU sort unavailable, sorting on the CPU");
@@ -54,9 +54,9 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
 }
 
 bool MetalFrameCompute::configure(const RenderPolicy& policy, std::string* reason) {
-  const MetalRadixSort::KeyBits bits = policy.sortDepth == SortKeyBits::low16
-                                           ? MetalRadixSort::KeyBits::Low16
-                                           : MetalRadixSort::KeyBits::Full32;
+  const KeyBits bits = policy.sortDepth == SortKeyBits::low16
+                                           ? KeyBits::Low16
+                                           : KeyBits::Full32;
   const bool visibilityChanged = bits != depthBits();
   const bool lodChanged =
       policy.lodErrorPixels != lodErrorPixels_ || policy.lodSplatLimit != lodSplatLimit_;
@@ -106,7 +106,7 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
   if (!createOrdering(visibility, sort, depthBits()) ||
       !reserveOrdering(visibility, sort, static_cast<uint32_t>(tree.nodeCount()), lod->budget()))
     return false;
-  std::array<id<MTLBuffer>, kSlots> readback{};
+  std::array<id<MTLBuffer>, kFramesInFlight> readback{};
   for (auto& buffer : readback) {
     buffer = metal::buffer(device_, sizeof(LodCounters));
     if (!buffer) return false;
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalDrawArguments.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalDrawArguments.h
new file mode 100644
index 0000000..618bb0f
--- /dev/null
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalDrawArguments.h
@@ -0,0 +1,14 @@
+#pragma once
+
+#import <Metal/Metal.h>
+
+#include <cstdint>
+
+namespace splatkit {
+
+// The indirect draw arguments per frame slot: the sorted set is split into kDrawBatches draws.
+// PrepareIndirect.metalh writes them with the same batch count.
+inline constexpr uint32_t kDrawBatches = 7;
+inline constexpr uint32_t kDrawArgumentBytes = sizeof(MTLDrawPrimitivesIndirectArguments);
+
+}  // namespace splatkit
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalFrameSlots.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalFrameSlots.h
new file mode 100644
index 0000000..c27d023
--- /dev/null
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalFrameSlots.h
@@ -0,0 +1,11 @@
+#pragma once
+
+#include <cstdint>
+
+namespace splatkit {
+
+// The frames the CPU may encode ahead of the GPU. Every per-frame slot array and the frame
+// loop's semaphore take their size from this.
+inline constexpr uint32_t kFramesInFlight = 2;
+
+}  // namespace splatkit
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalSortKeys.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalSortKeys.h
new file mode 100644
index 0000000..c428619
--- /dev/null
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalSortKeys.h
@@ -0,0 +1,11 @@
+#pragma once
+
+#include <cstdint>
+
+namespace splatkit {
+
+// The depth key width, in bits, that the visibility pass writes and the radix sort orders by.
+// Low16 requires zero upper bits. Keys remain uint32 storage in both widths.
+enum class KeyBits : uint32_t { Low16 = 16, Full32 = 32 };
+
+}  // namespace splatkit
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/raster/MetalSplatPipeline.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/raster/MetalSplatPipeline.mm
index fdd19e2..cdf5cc9 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/raster/MetalSplatPipeline.mm
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/raster/MetalSplatPipeline.mm
@@ -1,6 +1,6 @@
 #include "rendering/metal/raster/MetalSplatPipeline.h"
 
-#include "rendering/metal/cull/MetalVisibilityPass.h"
+#include "rendering/metal/layout/MetalDrawArguments.h"
 #include "splatkit/Log.h"
 
 namespace splatkit {
@@ -120,7 +120,7 @@
       [encoder setVertexBuffer:gpu.projected offset:0 atIndex:1];
       [encoder setVertexBuffer:gpu.order offset:0 atIndex:2];
       id<MTLBuffer> arguments = gpu.arguments;
-      for (uint32_t batch = 0; batch < MetalVisibilityPass::kDrawBatches; ++batch) {
+      for (uint32_t batch = 0; batch < kDrawBatches; ++batch) {
         if (batch > 0) {
           [encoder setRenderPipelineState:maskPipeline_];
           [encoder setDepthStencilState:maskDepth_];
@@ -130,7 +130,7 @@
         [encoder setDepthStencilState:splatDepth_];
         [encoder drawPrimitives:MTLPrimitiveTypeTriangleStrip
                   indirectBuffer:arguments
-            indirectBufferOffset:batch * MetalVisibilityPass::kDrawArgumentBytes];
+            indirectBufferOffset:batch * kDrawArgumentBytes];
       }
     } else {
       const int degree = world->drawDegree(shDegree);
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/sort/MetalRadixSort.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/sort/MetalRadixSort.h
index 43a26bf..4afbbbf 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/sort/MetalRadixSort.h
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/sort/MetalRadixSort.h
@@ -5,6 +5,8 @@
 #include <array>
 #include <cstdint>
 
+#include "rendering/metal/layout/MetalSortKeys.h"
+
 namespace splatkit {
 
 // Stable ascending sort of uint32 key/value pairs. No camera or splat knowledge.
@@ -14,8 +16,6 @@ namespace splatkit {
 // Scratch is shared across frames: all encodes and their consumers use one queue.
 class MetalRadixSort {
  public:
-  // Low16 requires zero upper bits. Keys remain uint32 storage in both modes.
-  enum class KeyBits : uint32_t { Low16 = 16, Full32 = 32 };
   bool create(id<MTLDevice> device, id<MTLLibrary> library);
   // Transactional allocation: failure preserves the previous buffers and capacity.
   bool reserve(uint32_t capacity);
diff --git a/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm b/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm
index d138b71..411d51f 100644
--- a/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm
+++ b/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm
@@ -39,9 +39,9 @@ void SetUp() override {
   VisibilityRig rig;
   ASSERT_TRUE(rig.create());
   ASSERT_TRUE(rig.reserve(16));
-  EXPECT_EQ(rig.pass.depthBits(), MetalRadixSort::KeyBits::Full32);
-  EXPECT_TRUE(rig.pass.reconfigure(MetalRadixSort::KeyBits::Low16));
-  EXPECT_EQ(rig.pass.depthBits(), MetalRadixSort::KeyBits::Low16);
+  EXPECT_EQ(rig.pass.depthBits(), KeyBits::Full32);
+  EXPECT_TRUE(rig.pass.reconfigure(KeyBits::Low16));
+  EXPECT_EQ(rig.pass.depthBits(), KeyBits::Low16);
   EXPECT_EQ(rig.pass.capacity(), 16u);
 }
 
@@ -108,7 +108,7 @@ explicit VisibilityInput(uint32_t count) {
   }
   const auto* draws =
       static_cast<const MTLDrawPrimitivesIndirectArguments*>(visibility.drawArguments(0).contents);
-  for (uint32_t i = 0; i < MetalVisibilityPass::kDrawBatches; ++i) {
+  for (uint32_t i = 0; i < kDrawBatches; ++i) {
     EXPECT_EQ(draws[i].instanceCount, 0u);
   }
 }
@@ -121,7 +121,7 @@ explicit VisibilityInput(uint32_t count) {
   id<MTLCommandBuffer> cmd = [Gpu::get().queue commandBuffer];
   EXPECT_FALSE(visibility.encode(cmd, 0, input.uniforms, input.splats, nil, 0, &range, 1));
   EXPECT_FALSE(visibility.encode(cmd, 0, input.uniforms, input.splats, nil, 0, nullptr, 1));
-  EXPECT_FALSE(visibility.encode(cmd, MetalVisibilityPass::kSlots, input.uniforms, input.splats,
+  EXPECT_FALSE(visibility.encode(cmd, kFramesInFlight, input.uniforms, input.splats,
                                  nil, 0, nullptr, 0));
   [cmd commit];
   [cmd waitUntilCompleted];
@@ -203,7 +203,7 @@ explicit VisibilityInput(uint32_t count) {
   // The batches partition the instances, in order.
   const auto* draw = static_cast<const uint32_t*>(visibility.drawArguments(1).contents);
   uint32_t next = 0;
-  for (uint32_t b = 0; b < MetalVisibilityPass::kDrawBatches; ++b) {
+  for (uint32_t b = 0; b < kDrawBatches; ++b) {
     EXPECT_EQ(draw[b * 4], 4u);        // vertices per instance
     EXPECT_EQ(draw[b * 4 + 3], next);  // base instance
     next += draw[b * 4 + 1];           // instances
@@ -363,7 +363,7 @@ explicit VisibilityInput(uint32_t count) {
   for (bool indexed : {false, true}) {
     SCOPED_TRACE(indexed);
     VisibilityRig visibility;
-    ASSERT_TRUE(visibility.create(MetalRadixSort::KeyBits::Low16));
+    ASSERT_TRUE(visibility.create(KeyBits::Low16));
     ASSERT_TRUE(visibility.reserve(n));
     const SplatRenderer::Range range{0, n};
     auto cmd = [gpu.queue commandBuffer];
@@ -408,7 +408,7 @@ explicit VisibilityInput(uint32_t count) {
   const auto* order = static_cast<const uint32_t*>(visibility.order().contents);
   const auto* projected = static_cast<const ProjectedSplat*>(visibility.projected().contents);
   for (uint32_t n : {1u, 31u, 32u, 33u, 255u, 256u, 257u, 513u, 0u}) {
-    const uint32_t slot = n % MetalVisibilityPass::kSlots;
+    const uint32_t slot = n % kFramesInFlight;
     const SplatRenderer::Range range{0, n};
     id<MTLCommandBuffer> cmd = [gpu.queue commandBuffer];
     ASSERT_TRUE(visibility.encode(cmd, slot, input.uniforms, input.splats, nil, 0,
@@ -430,7 +430,7 @@ explicit VisibilityInput(uint32_t count) {
     const auto* draws = static_cast<const MTLDrawPrimitivesIndirectArguments*>(
         visibility.drawArguments(slot).contents);
     uint32_t next = 0;
-    for (uint32_t b = 0; b < MetalVisibilityPass::kDrawBatches; ++b) {
+    for (uint32_t b = 0; b < kDrawBatches; ++b) {
       EXPECT_EQ(draws[b].vertexCount, 4u);
       EXPECT_EQ(draws[b].vertexStart, 0u);
       EXPECT_EQ(draws[b].baseInstance, next);
diff --git a/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm b/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm
index 0b171ee..180565c 100644
--- a/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm
+++ b/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm
@@ -76,14 +76,14 @@ void run(id<MTLCommandBuffer> selection, id<MTLCommandBuffer> sort) {
   }
 
   test::Gpu* gpu_ = nullptr;
-  MetalFrameLoop loop_{MetalFrameCompute::kSlots};
+  MetalFrameLoop loop_{kFramesInFlight};
   std::unique_ptr<MetalFrameCompute> compute_;
   SplatRenderer::Frame frame_;
 };
 
 TEST_F(MetalFrameComputeTest, StartsWithGpuOrderingAndFullWidthKeys) {
   EXPECT_TRUE(compute_->gpuSort());
-  EXPECT_EQ(compute_->depthBits(), MetalRadixSort::KeyBits::Full32);
+  EXPECT_EQ(compute_->depthBits(), KeyBits::Full32);
   EXPECT_FALSE(compute_->hasLod());
   EXPECT_EQ(compute_->lodPixels(), 1.0f);
   EXPECT_NE(compute_->uniforms(0), nil);
@@ -120,9 +120,9 @@ void run(id<MTLCommandBuffer> selection, id<MTLCommandBuffer> sort) {
   std::string reason;
   EXPECT_TRUE(compute_->configure(policy(1.0f), &reason));
   EXPECT_TRUE(compute_->configure(policy(1.0f, 0, SortKeyBits::low16), &reason)) << reason;
-  EXPECT_EQ(compute_->depthBits(), MetalRadixSort::KeyBits::Low16);
+  EXPECT_EQ(compute_->depthBits(), KeyBits::Low16);
   EXPECT_TRUE(compute_->configure(policy(4.0f, 100), &reason));
-  EXPECT_EQ(compute_->depthBits(), MetalRadixSort::KeyBits::Full32);
+  EXPECT_EQ(compute_->depthBits(), KeyBits::Full32);
   EXPECT_EQ(compute_->lodPixels(), 4.0f);
 }
 
@@ -147,7 +147,7 @@ void run(id<MTLCommandBuffer> selection, id<MTLCommandBuffer> sort) {
   const auto* draws =
       static_cast<const MTLDrawPrimitivesIndirectArguments*>(compute_->drawArguments(0).contents);
   uint32_t next = 0;
-  for (uint32_t b = 0; b < MetalVisibilityPass::kDrawBatches; ++b) {
+  for (uint32_t b = 0; b < kDrawBatches; ++b) {
     EXPECT_EQ(draws[b].baseInstance, next);
     next += draws[b].instanceCount;
   }
@@ -214,7 +214,7 @@ void run(id<MTLCommandBuffer> selection, id<MTLCommandBuffer> sort) {
   // A limit below the refined cut denies the split and raises the threshold; lifting it
   // restores the policy's threshold at once.
   ASSERT_TRUE(compute_->configure(policy(0.1f, 32), &reason));
-  for (uint32_t i = 0; i < 4; ++i) frame(i % MetalFrameCompute::kSlots);
+  for (uint32_t i = 0; i < 4; ++i) frame(i % kFramesInFlight);
   EXPECT_EQ(compute_->lastSelectedCount(), 1u);
   EXPECT_GT(compute_->lastLodLimitedCount(), 0u);
   EXPECT_GT(compute_->lodPixels(), 0.1f);
diff --git a/packages/splatkit-ios/tests/rendering/metal/presentation/MetalSurfaceTest.mm b/packages/splatkit-ios/tests/rendering/metal/presentation/MetalSurfaceTest.mm
index 1cf6731..f362163 100644
--- a/packages/splatkit-ios/tests/rendering/metal/presentation/MetalSurfaceTest.mm
+++ b/packages/splatkit-ios/tests/rendering/metal/presentation/MetalSurfaceTest.mm
@@ -2,14 +2,13 @@
 
 #include <vector>
 
+#include "rendering/metal/layout/MetalFrameSlots.h"
 #include "rendering/metal/presentation/MetalSurface.h"
 #include "support/MetalTestContext.h"
 
 namespace splatkit {
 namespace {
 
-constexpr uint32_t kSlots = 2;
-
 class MetalSurfaceTest : public testing::Test {
  protected:
   void SetUp() override {
@@ -28,7 +27,7 @@ bool scalerSupported() const {
   }
 
   test::Gpu* gpu_ = nullptr;
-  MetalFrameLoop loop_{kSlots};
+  MetalFrameLoop loop_{kFramesInFlight};
   CAMetalLayer* layer_ = nil;
 };
 
diff --git a/packages/splatkit-ios/tests/rendering/metal/raster/MetalSplatPipelineTest.mm b/packages/splatkit-ios/tests/rendering/metal/raster/MetalSplatPipelineTest.mm
index 345372e..6d7d330 100644
--- a/packages/splatkit-ios/tests/rendering/metal/raster/MetalSplatPipelineTest.mm
+++ b/packages/splatkit-ios/tests/rendering/metal/raster/MetalSplatPipelineTest.mm
@@ -72,7 +72,7 @@ void drawCpu(id<MTLCommandBuffer> cmd, const MetalWorld* world, uint32_t drawCou
   }
 
   test::Gpu* gpu_ = nullptr;
-  MetalFrameLoop loop_{MetalFrameCompute::kSlots};
+  MetalFrameLoop loop_{kFramesInFlight};
   std::unique_ptr<MetalFrameCompute> compute_;
   MetalSplatPipeline pipeline_;
   std::unique_ptr<MetalRenderTarget> target_;
diff --git a/packages/splatkit-ios/tests/rendering/metal/sort/MetalRadixSortTest.mm b/packages/splatkit-ios/tests/rendering/metal/sort/MetalRadixSortTest.mm
index 0342433..5802070 100644
--- a/packages/splatkit-ios/tests/rendering/metal/sort/MetalRadixSortTest.mm
+++ b/packages/splatkit-ios/tests/rendering/metal/sort/MetalRadixSortTest.mm
@@ -20,7 +20,7 @@
 
 // Fills the key and value buffers with `pairs`, sorts on the GPU, returns what came out.
 std::vector<Pair> sortOnGpu(MetalRadixSort& v, const std::vector<Pair>& pairs,
-                            MetalRadixSort::KeyBits bits = MetalRadixSort::KeyBits::Full32) {
+                            KeyBits bits = KeyBits::Full32) {
   Gpu& gpu = Gpu::get();
   auto* keys = static_cast<uint32_t*>(v.keys().contents);
   auto* values = static_cast<uint32_t*>(v.values().contents);
@@ -113,11 +113,11 @@ void SetUp() override {
       pairs.back().key = 0;
     }
     const auto expected = sortOnCpu(pairs);
-    EXPECT_EQ(sortOnGpu(sort, pairs, MetalRadixSort::KeyBits::Low16), expected);
+    EXPECT_EQ(sortOnGpu(sort, pairs, KeyBits::Low16), expected);
     EXPECT_EQ(sortOnGpu(sort, pairs), expected);
   }
   const auto ties = randomPairs(70000, 0x7u, 17);
-  EXPECT_EQ(sortOnGpu(sort, ties, MetalRadixSort::KeyBits::Low16), sortOnCpu(ties));
+  EXPECT_EQ(sortOnGpu(sort, ties, KeyBits::Low16), sortOnCpu(ties));
 }
 
 // Not a check, a number: the GPU time of a sort at the scale a phone draws.
diff --git a/packages/splatkit-ios/tests/support/MetalVisibilityRig.h b/packages/splatkit-ios/tests/support/MetalVisibilityRig.h
index f8edcca..d2de066 100644
--- a/packages/splatkit-ios/tests/support/MetalVisibilityRig.h
+++ b/packages/splatkit-ios/tests/support/MetalVisibilityRig.h
@@ -14,7 +14,7 @@ struct VisibilityRig {
   MetalVisibilityPass pass;
   MetalRadixSort sort;
 
-  bool create(MetalRadixSort::KeyBits bits = MetalRadixSort::KeyBits::Full32) {
+  bool create(KeyBits bits = KeyBits::Full32) {
     return pass.create(Gpu::get().device, Gpu::get().library, bits) &&
            sort.create(Gpu::get().device, Gpu::get().library);
   }

From 85f0970a916b311a2d23c6d39f6f0fcfe2197dfa Mon Sep 17 00:00:00 2001
From: Xget7 <juanieltupa@gmail.com>
Date: Tue, 20 Oct 2026 01:20:32 -0300
Subject: [PATCH 2/8] Name the Metal renderer's render-scale, residency, LOD
 error and texture limits

MetalSurface clamps the render scale with kMinRenderScale and
kMaxRenderScale, the same names VulkanSurface uses. MetalSplatRenderer
reports its residency window as kMinResidencyCapacity and
kMaxResidencyCapacity (the Android names), the LOD error range as
kMinLodErrorPixels and kMaxLodErrorPixels, and the texture ceiling as
kAppleFamily7MaxTextureDimension.

The values are unchanged.
---
 .../rendering/metal/MetalSplatRenderer.mm     | 27 ++++++++++++-------
 .../metal/presentation/MetalSurface.h         |  4 +++
 .../metal/presentation/MetalSurface.mm        |  2 +-
 3 files changed, 23 insertions(+), 10 deletions(-)

diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm
index efacea0..08be6c4 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm
@@ -7,6 +7,18 @@
 #include "splatkit/Log.h"
 
 namespace splatkit {
+namespace {
+
+// The residency range reported to the adapters.
+constexpr uint32_t kMinResidencyCapacity = 100'000;
+constexpr uint32_t kMaxResidencyCapacity = 32'000'000;
+// The LOD error threshold range, in pixels, a policy may request.
+constexpr float kMinLodErrorPixels = 0.1f;
+constexpr float kMaxLodErrorPixels = 16.0f;
+// Apple family 7, the minimum this renderer accepts, supports 16384x16384 textures.
+constexpr uint32_t kAppleFamily7MaxTextureDimension = 16384;
+
+}  // namespace
 
 std::unique_ptr<MetalSplatRenderer> MetalSplatRenderer::create() {
   std::unique_ptr<MetalSplatRenderer> r(new MetalSplatRenderer());
@@ -70,17 +82,14 @@
 
 DeviceCapabilities MetalSplatRenderer::deviceCapabilities() const {
   DeviceCapabilities caps;
-  // MetalLodSelection caps a selected cut at 4M nodes; the engine's residency window is
-  // [100k, 32M].
   caps.limits.maxLodCapacitySplats = MetalLodSelection::kMaxBudget;
-  caps.limits.minResidencyCapacitySplats = 100'000;
-  caps.limits.maxResidencyCapacitySplats = 32'000'000;
+  caps.limits.minResidencyCapacitySplats = kMinResidencyCapacity;
+  caps.limits.maxResidencyCapacitySplats = kMaxResidencyCapacity;
   caps.supportsComputeTiles = false;
   caps.supportsHiZOcclusion = false;
   caps.supportsSubgroups = compute_->gpuSort();
-  // Metal exposes no texture-dimension query; Apple family 7, the minimum this renderer
-  // accepts, supports 16384x16384.
-  caps.maxTextureDimension = context_->device() != nil ? 16384u : 0u;
+  // Metal exposes no texture-dimension query.
+  caps.maxTextureDimension = context_->device() != nil ? kAppleFamily7MaxTextureDimension : 0u;
   RenderPolicySupport& policy = caps.policy;
   policy.fallback.raster = RasterStrategy::hardware;
   policy.fallback.tileSize = 16;
@@ -96,8 +105,8 @@
   // alpha, tile size and occlusion stay fixed shader/rasterization choices.
   // The LOD error threshold applies live to a hierarchy world and at its next upload.
   policy.lodErrorPixels = compute_->gpuSort();
-  policy.minLodErrorPixels = 0.1f;
-  policy.maxLodErrorPixels = 16.0f;
+  policy.minLodErrorPixels = kMinLodErrorPixels;
+  policy.maxLodErrorPixels = kMaxLodErrorPixels;
   policy.lodSplatLimit = compute_->gpuSort();
   policy.sortDepth = compute_->gpuSort();
   return caps;
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalSurface.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalSurface.h
index 41675b5..ce9c468 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalSurface.h
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalSurface.h
@@ -20,6 +20,9 @@ namespace splatkit {
 // thread only.
 class MetalSurface {
  public:
+  static constexpr float kMinRenderScale = 0.1f;
+  static constexpr float kMaxRenderScale = 2.0f;
+
   // `compositesInHalfFloat`: the splats are accumulated in a half-float target at every scale,
   // which the GPU order path needs because 8 bits would round away the coverage.
   MetalSurface(id<MTLDevice> device, id<MTLLibrary> library, MetalFrameLoop& frameLoop,
@@ -33,6 +36,7 @@ class MetalSurface {
   bool setLayer(CAMetalLayer* layer);
   // The layer's size in pixels.
   void setDrawableSize(uint32_t width, uint32_t height);
+  // Clamped to [kMinRenderScale, kMaxRenderScale]; below or above 1 draws offscreen.
   void setRenderScale(float scale);
   void setSpatialUpscaling(bool enabled);
   // False when the blit pipeline for the new format could not be built.
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalSurface.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalSurface.mm
index 7679014..a453337 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalSurface.mm
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalSurface.mm
@@ -37,7 +37,7 @@
 }
 
 void MetalSurface::setRenderScale(float scale) {
-  scale = std::clamp(scale, 0.1f, 2.0f);
+  scale = std::clamp(scale, kMinRenderScale, kMaxRenderScale);
   if (scale == renderScale_) return;
   renderScale_ = scale;
   if (upscaler_.active()) frameLoop_.waitIdle();

From 3c0f7a8b14d5e27417a1f6b0d57464f919c1a4d8 Mon Sep 17 00:00:00 2001
From: Xget7 <juanieltupa@gmail.com>
Date: Tue, 20 Oct 2026 01:20:38 -0300
Subject: [PATCH 3/8] Drop the unused GpuLayout.h include from
 MetalSplatRenderer.h

Nothing the header declares comes from GpuLayout.h; the sources that need
its records include it themselves.
---
 .../Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h    | 1 -
 1 file changed, 1 deletion(-)

diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h
index c246a92..d51d954 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h
@@ -17,7 +17,6 @@
 #include "rendering/metal/presentation/MetalSurface.h"
 #include "rendering/metal/raster/MetalSplatPipeline.h"
 #include "rendering/metal/world/MetalWorld.h"
-#include "splatkit/rendering/GpuLayout.h"
 #include "splatkit/rendering/SplatRenderer.h"
 
 namespace splatkit {

From 05fc479e7dddd99e10d1ac55c388c90f036148f1 Mon Sep 17 00:00:00 2001
From: Xget7 <juanieltupa@gmail.com>
Date: Tue, 20 Oct 2026 01:21:03 -0300
Subject: [PATCH 4/8] Test that SKPngWriter writes a decodable PNG and rejects
 bad input

The test writes a 3x2 BGRA image, decodes the file with ImageIO and checks
its size and every channel. A zero size, a buffer one byte short or long,
and an unwritable path all return false, and the first two leave no file.
---
 packages/splatkit-ios/tests/CMakeLists.txt    |  1 +
 .../tests/diagnostics/SKPngWriterTest.mm      | 99 +++++++++++++++++++
 2 files changed, 100 insertions(+)
 create mode 100644 packages/splatkit-ios/tests/diagnostics/SKPngWriterTest.mm

diff --git a/packages/splatkit-ios/tests/CMakeLists.txt b/packages/splatkit-ios/tests/CMakeLists.txt
index a8ec943..5feeb21 100644
--- a/packages/splatkit-ios/tests/CMakeLists.txt
+++ b/packages/splatkit-ios/tests/CMakeLists.txt
@@ -1,4 +1,5 @@
 add_executable(splatkit_ios_tests
+  diagnostics/SKPngWriterTest.mm
   rendering/metal/MetalSplatRendererTest.mm
   rendering/metal/cull/MetalVisibilityPassTest.mm
   rendering/metal/device/MetalContextTest.mm
diff --git a/packages/splatkit-ios/tests/diagnostics/SKPngWriterTest.mm b/packages/splatkit-ios/tests/diagnostics/SKPngWriterTest.mm
new file mode 100644
index 0000000..3f3af75
--- /dev/null
+++ b/packages/splatkit-ios/tests/diagnostics/SKPngWriterTest.mm
@@ -0,0 +1,99 @@
+#include <gtest/gtest.h>
+
+#import <CoreGraphics/CoreGraphics.h>
+#import <ImageIO/ImageIO.h>
+
+#include <vector>
+
+#include "diagnostics/SKPngWriter.h"
+
+namespace splatkit {
+namespace {
+
+constexpr uint32_t kBytesPerPixel = 4;
+
+class SKPngWriterTest : public testing::Test {
+ protected:
+  void SetUp() override {
+    NSString* name = [NSString stringWithFormat:@"splatkit-png-%@.png", [NSUUID UUID].UUIDString];
+    path_ = [NSTemporaryDirectory() stringByAppendingPathComponent:name];
+  }
+  void TearDown() override { [NSFileManager.defaultManager removeItemAtPath:path_ error:nil]; }
+
+  NSString* path_ = nil;
+};
+
+// Decodes `path` to RGBA bytes, empty when the file is not a PNG.
+std::vector<uint8_t> decode(NSString* path, size_t* width, size_t* height) {
+  CGImageSourceRef source =
+      CGImageSourceCreateWithURL((__bridge CFURLRef)[NSURL fileURLWithPath:path], nullptr);
+  if (source == nullptr) return {};
+  std::vector<uint8_t> rgba;
+  CGImageRef image = CGImageSourceCreateImageAtIndex(source, 0, nullptr);
+  if (image != nullptr) {
+    *width = CGImageGetWidth(image);
+    *height = CGImageGetHeight(image);
+    rgba.assign(*width * *height * kBytesPerPixel, 0);
+    CGColorSpaceRef space = CGColorSpaceCreateDeviceRGB();
+    CGContextRef context = CGBitmapContextCreate(rgba.data(), *width, *height, 8,
+                                                 *width * kBytesPerPixel, space,
+                                                 kCGImageAlphaPremultipliedLast);
+    CGContextDrawImage(context, CGRectMake(0, 0, static_cast<CGFloat>(*width),
+                                           static_cast<CGFloat>(*height)),
+                       image);
+    CGContextRelease(context);
+    CGColorSpaceRelease(space);
+    CGImageRelease(image);
+  }
+  CFRelease(source);
+  return rgba;
+}
+
+TEST_F(SKPngWriterTest, WritesADecodablePngOfTheRequestedSizeWithTheSamePixels) {
+  constexpr uint32_t kWidth = 3;
+  constexpr uint32_t kHeight = 2;
+  // BGRA, top row first: blue, green, red / white, black, mid grey.
+  const std::vector<uint8_t> bgra = {
+      255, 0,   0,   255, 0,   255, 0,   255, 0,   0,   255, 255,
+      255, 255, 255, 255, 0,   0,   0,   255, 128, 128, 128, 255,
+  };
+  ASSERT_EQ(bgra.size(), size_t{kWidth} * kHeight * kBytesPerPixel);
+  ASSERT_TRUE(writePng(path_, bgra, kWidth, kHeight));
+
+  size_t width = 0, height = 0;
+  const std::vector<uint8_t> rgba = decode(path_, &width, &height);
+  ASSERT_EQ(width, size_t{kWidth});
+  ASSERT_EQ(height, size_t{kHeight});
+  ASSERT_EQ(rgba.size(), bgra.size());
+  for (size_t pixel = 0; pixel < size_t{kWidth} * kHeight; ++pixel) {
+    const size_t at = pixel * kBytesPerPixel;
+    EXPECT_EQ(rgba[at + 0], bgra[at + 2]) << "red of pixel " << pixel;
+    EXPECT_EQ(rgba[at + 1], bgra[at + 1]) << "green of pixel " << pixel;
+    EXPECT_EQ(rgba[at + 2], bgra[at + 0]) << "blue of pixel " << pixel;
+    EXPECT_EQ(rgba[at + 3], 255) << "alpha of pixel " << pixel;
+  }
+}
+
+TEST_F(SKPngWriterTest, RejectsAZeroSize) {
+  const std::vector<uint8_t> bgra(kBytesPerPixel, 0);
+  EXPECT_FALSE(writePng(path_, bgra, 0, 1));
+  EXPECT_FALSE(writePng(path_, bgra, 1, 0));
+  EXPECT_FALSE([NSFileManager.defaultManager fileExistsAtPath:path_]);
+}
+
+TEST_F(SKPngWriterTest, RejectsABufferThatDoesNotMatchTheSize) {
+  const std::vector<uint8_t> shortBuffer(2 * 2 * kBytesPerPixel - 1, 0);
+  EXPECT_FALSE(writePng(path_, shortBuffer, 2, 2));
+  const std::vector<uint8_t> longBuffer(2 * 2 * kBytesPerPixel + 1, 0);
+  EXPECT_FALSE(writePng(path_, longBuffer, 2, 2));
+  EXPECT_FALSE([NSFileManager.defaultManager fileExistsAtPath:path_]);
+}
+
+TEST_F(SKPngWriterTest, FailsWhenTheFileCannotBeWritten) {
+  const std::vector<uint8_t> bgra(kBytesPerPixel, 0);
+  NSString* missing = [path_ stringByAppendingPathComponent:@"no-such-directory/out.png"];
+  EXPECT_FALSE(writePng(missing, bgra, 1, 1));
+}
+
+}  // namespace
+}  // namespace splatkit

From b40154ab4a3c67bd9b144547f271f127c491a7e2 Mon Sep 17 00:00:00 2001
From: Xget7 <juanieltupa@gmail.com>
Date: Tue, 20 Oct 2026 01:21:03 -0300
Subject: [PATCH 5/8] Test the SKTypeConversions mirrors between the C++ engine
 and the C structs

Policy round trips both sort depths, and toCppPolicy rejects a sort depth of
24 and leaves the output untouched. Policy support, stats, capabilities,
camera state, camera request and lens settings each check every field
against distinct values.
---
 packages/splatkit-ios/tests/CMakeLists.txt    |   1 +
 .../tests/engine/SKTypeConversionsTest.mm     | 261 ++++++++++++++++++
 2 files changed, 262 insertions(+)
 create mode 100644 packages/splatkit-ios/tests/engine/SKTypeConversionsTest.mm

diff --git a/packages/splatkit-ios/tests/CMakeLists.txt b/packages/splatkit-ios/tests/CMakeLists.txt
index 5feeb21..732e5c8 100644
--- a/packages/splatkit-ios/tests/CMakeLists.txt
+++ b/packages/splatkit-ios/tests/CMakeLists.txt
@@ -1,5 +1,6 @@
 add_executable(splatkit_ios_tests
   diagnostics/SKPngWriterTest.mm
+  engine/SKTypeConversionsTest.mm
   rendering/metal/MetalSplatRendererTest.mm
   rendering/metal/cull/MetalVisibilityPassTest.mm
   rendering/metal/device/MetalContextTest.mm
diff --git a/packages/splatkit-ios/tests/engine/SKTypeConversionsTest.mm b/packages/splatkit-ios/tests/engine/SKTypeConversionsTest.mm
new file mode 100644
index 0000000..4882ffb
--- /dev/null
+++ b/packages/splatkit-ios/tests/engine/SKTypeConversionsTest.mm
@@ -0,0 +1,261 @@
+#include <gtest/gtest.h>
+
+#include "engine/SKTypeConversions.h"
+
+namespace splatkit {
+namespace {
+
+RenderPolicy customPolicy() {
+  RenderPolicy policy;
+  policy.raster = RasterStrategy::hybrid;
+  policy.tileSize = 32;
+  policy.lodErrorPixels = 2.5f;
+  policy.lodSplatLimit = 123456;
+  policy.alphaThreshold = 0.125f;
+  policy.subpixelThreshold = 1.5f;
+  policy.enableFrustumCulling = false;
+  policy.enableHiZOcclusion = true;
+  policy.enableEarlyTermination = false;
+  policy.sortDepth = SortKeyBits::low16;
+  return policy;
+}
+
+void expectSamePolicy(const RenderPolicy& actual, const RenderPolicy& expected) {
+  EXPECT_EQ(actual.raster, expected.raster);
+  EXPECT_EQ(actual.tileSize, expected.tileSize);
+  EXPECT_EQ(actual.lodErrorPixels, expected.lodErrorPixels);
+  EXPECT_EQ(actual.lodSplatLimit, expected.lodSplatLimit);
+  EXPECT_EQ(actual.alphaThreshold, expected.alphaThreshold);
+  EXPECT_EQ(actual.subpixelThreshold, expected.subpixelThreshold);
+  EXPECT_EQ(actual.enableFrustumCulling, expected.enableFrustumCulling);
+  EXPECT_EQ(actual.enableHiZOcclusion, expected.enableHiZOcclusion);
+  EXPECT_EQ(actual.enableEarlyTermination, expected.enableEarlyTermination);
+  EXPECT_EQ(actual.sortDepth, expected.sortDepth);
+}
+
+TEST(SKTypeConversionsTest, APolicyRoundTripsThroughTheCStruct) {
+  const RenderPolicy policy = customPolicy();
+  const SKRenderPolicy mirrored = fromCppPolicy(policy);
+  EXPECT_EQ(mirrored.raster, 2u);
+  EXPECT_EQ(mirrored.sortDepth, 16u);
+  EXPECT_EQ(mirrored.enableHiZOcclusion, YES);
+  EXPECT_EQ(mirrored.enableFrustumCulling, NO);
+
+  RenderPolicy back;
+  ASSERT_TRUE(toCppPolicy(mirrored, &back));
+  expectSamePolicy(back, policy);
+}
+
+TEST(SKTypeConversionsTest, The32BitSortDepthRoundTrips) {
+  RenderPolicy policy = customPolicy();
+  policy.sortDepth = SortKeyBits::full32;
+  RenderPolicy back;
+  ASSERT_TRUE(toCppPolicy(fromCppPolicy(policy), &back));
+  EXPECT_EQ(back.sortDepth, SortKeyBits::full32);
+}
+
+TEST(SKTypeConversionsTest, ABadSortDepthIsRejectedAndLeavesTheOutputUntouched) {
+  SKRenderPolicy mirrored = fromCppPolicy(customPolicy());
+  mirrored.sortDepth = 24;
+  const RenderPolicy sentinel;
+  RenderPolicy out = sentinel;
+  EXPECT_FALSE(toCppPolicy(mirrored, &out));
+  expectSamePolicy(out, sentinel);
+}
+
+TEST(SKTypeConversionsTest, PolicySupportKeepsEveryFlagAndRange) {
+  RenderPolicySupport support;
+  support.raster = true;
+  support.tileSize = true;
+  support.lodErrorPixels = true;
+  support.lodSplatLimit = true;
+  support.alphaThreshold = false;
+  support.subpixelThreshold = true;
+  support.enableFrustumCulling = true;
+  support.enableHiZOcclusion = false;
+  support.enableEarlyTermination = true;
+  support.sortDepth = true;
+  support.rasterMask = 5;
+  support.tileSizeMask = 6;
+  support.minLodErrorPixels = 0.25f;
+  support.maxLodErrorPixels = 12.0f;
+  support.minSubpixelThreshold = 0.5f;
+  support.maxSubpixelThreshold = 4.0f;
+
+  const SKRenderPolicySupport out = fromCppSupport(support);
+  EXPECT_EQ(out.raster, YES);
+  EXPECT_EQ(out.tileSize, YES);
+  EXPECT_EQ(out.lodErrorPixels, YES);
+  EXPECT_EQ(out.lodSplatLimit, YES);
+  EXPECT_EQ(out.alphaThreshold, NO);
+  EXPECT_EQ(out.subpixelThreshold, YES);
+  EXPECT_EQ(out.enableFrustumCulling, YES);
+  EXPECT_EQ(out.enableHiZOcclusion, NO);
+  EXPECT_EQ(out.enableEarlyTermination, YES);
+  EXPECT_EQ(out.sortDepth, YES);
+  EXPECT_EQ(out.rasterMask, 5u);
+  EXPECT_EQ(out.tileSizeMask, 6u);
+  EXPECT_EQ(out.minLodErrorPixels, 0.25f);
+  EXPECT_EQ(out.maxLodErrorPixels, 12.0f);
+  EXPECT_EQ(out.minSubpixelThreshold, 0.5f);
+  EXPECT_EQ(out.maxSubpixelThreshold, 4.0f);
+}
+
+TEST(SKTypeConversionsTest, StatsKeepEveryField) {
+  Stats stats;
+  stats.fps = 59.5f;
+  stats.frameMillis = 16.8f;
+  stats.presentTiming = true;
+  stats.frameMillisP95 = 21.0f;
+  stats.lowFps = 41.0f;
+  stats.droppedFrames = 3;
+  stats.gpuMillis = 7.25f;
+  stats.sortMillis = 1.5f;
+  stats.splatCount = 1'000'001;
+  stats.walking = true;
+  stats.motion = true;
+  stats.drawnSplatCount = 777'777;
+  stats.computeTileCount = 11;
+  stats.nonemptyComputeTileCount = 7;
+  stats.hardwareTileCount = 5;
+
+  const SKSplatStats out = fromCppStats(stats);
+  EXPECT_EQ(out.fps, 59.5f);
+  EXPECT_EQ(out.frameMillis, 16.8f);
+  EXPECT_EQ(out.gpuMillis, 7.25f);
+  EXPECT_EQ(out.sortMillis, 1.5f);
+  EXPECT_EQ(out.splatCount, 1'000'001u);
+  EXPECT_EQ(out.walking, YES);
+  EXPECT_EQ(out.motion, YES);
+  EXPECT_EQ(out.drawnSplatCount, 777'777u);
+  EXPECT_EQ(out.computeTileCount, 11u);
+  EXPECT_EQ(out.nonemptyComputeTileCount, 7u);
+  EXPECT_EQ(out.hardwareTileCount, 5u);
+  EXPECT_EQ(out.presentTiming, YES);
+  EXPECT_EQ(out.frameMillisP95, 21.0f);
+  EXPECT_EQ(out.lowFps, 41.0f);
+  EXPECT_EQ(out.droppedFrames, 3u);
+}
+
+TEST(SKTypeConversionsTest, CapabilitiesCarryLimitsFeaturesSupportAndTheFallbackPolicy) {
+  DeviceCapabilities caps;
+  caps.limits.maxLodCapacitySplats = 4'000'000;
+  caps.limits.minResidencyCapacitySplats = 100'000;
+  caps.limits.maxResidencyCapacitySplats = 32'000'000;
+  caps.supportsComputeTiles = true;
+  caps.supportsHiZOcclusion = false;
+  caps.supportsSubgroups = true;
+  caps.maxTextureDimension = 16384;
+  caps.policy.sortDepth = true;
+  caps.policy.maxLodErrorPixels = 16.0f;
+  caps.policy.fallback = customPolicy();
+
+  const SKDeviceCapabilities out = fromCppCapabilities(caps);
+  EXPECT_EQ(out.maxLodCapacitySplats, 4'000'000u);
+  EXPECT_EQ(out.minResidencyCapacitySplats, 100'000u);
+  EXPECT_EQ(out.maxResidencyCapacitySplats, 32'000'000u);
+  EXPECT_EQ(out.supportsComputeTiles, YES);
+  EXPECT_EQ(out.supportsHiZOcclusion, NO);
+  EXPECT_EQ(out.supportsSubgroups, YES);
+  EXPECT_EQ(out.maxTextureDimension, 16384u);
+  EXPECT_EQ(out.policy.sortDepth, YES);
+  EXPECT_EQ(out.policy.maxLodErrorPixels, 16.0f);
+
+  RenderPolicy fallback;
+  ASSERT_TRUE(toCppPolicy(out.fallback, &fallback));
+  expectSamePolicy(fallback, customPolicy());
+}
+
+TEST(SKTypeConversionsTest, ACameraRequestMapsToTheCppRequest) {
+  SKCameraRequest request{};
+  request.mode = SKCameraModeOrbit;
+  request.anchor = {1.0f, -2.0f, 3.5f};
+  request.radius = 4.0f;
+  request.azimuth = 0.5f;
+  request.elevation = -0.25f;
+  request.orbitRadiansPerSecond = -0.75f;
+
+  const CameraRequest out = toCppCameraRequest(request);
+  EXPECT_EQ(out.mode, CameraMode::Orbit);
+  EXPECT_EQ(out.anchor.x, 1.0f);
+  EXPECT_EQ(out.anchor.y, -2.0f);
+  EXPECT_EQ(out.anchor.z, 3.5f);
+  EXPECT_EQ(out.radius, 4.0f);
+  EXPECT_EQ(out.azimuth, 0.5f);
+  EXPECT_EQ(out.elevation, -0.25f);
+  EXPECT_EQ(out.orbitRadiansPerSecond, -0.75f);
+}
+
+TEST(SKTypeConversionsTest, ACameraStateMapsToTheCStruct) {
+  CameraState state;
+  state.mode = CameraMode::Orbit;
+  state.hasAnchor = true;
+  state.anchor = {-1.0f, 2.0f, 0.5f};
+  state.radius = 6.0f;
+  state.azimuth = 1.25f;
+  state.elevation = 0.1f;
+  state.orbitRadiansPerSecond = 0.3f;
+
+  const SKCameraState out = fromCppCameraState(state);
+  EXPECT_EQ(out.mode, SKCameraModeOrbit);
+  EXPECT_EQ(out.hasAnchor, YES);
+  EXPECT_EQ(out.anchor.x, -1.0f);
+  EXPECT_EQ(out.anchor.y, 2.0f);
+  EXPECT_EQ(out.anchor.z, 0.5f);
+  EXPECT_EQ(out.radius, 6.0f);
+  EXPECT_EQ(out.azimuth, 1.25f);
+  EXPECT_EQ(out.elevation, 0.1f);
+  EXPECT_EQ(out.orbitRadiansPerSecond, 0.3f);
+}
+
+TEST(SKTypeConversionsTest, ACameraRequestRoundTripsThroughTheStateMirror) {
+  SKCameraRequest request{};
+  request.mode = SKCameraModeFirstPerson;
+  request.anchor = {0.25f, 0.5f, 0.75f};
+  request.radius = 2.0f;
+  request.azimuth = 3.0f;
+  request.elevation = 0.4f;
+  request.orbitRadiansPerSecond = 1.0f;
+
+  const CameraRequest cpp = toCppCameraRequest(request);
+  CameraState state;
+  state.mode = cpp.mode;
+  state.hasAnchor = true;
+  state.anchor = cpp.anchor;
+  state.radius = cpp.radius;
+  state.azimuth = cpp.azimuth;
+  state.elevation = cpp.elevation;
+  state.orbitRadiansPerSecond = cpp.orbitRadiansPerSecond;
+  const SKCameraState mirrored = fromCppCameraState(state);
+  EXPECT_EQ(mirrored.mode, request.mode);
+  EXPECT_EQ(mirrored.anchor.x, request.anchor.x);
+  EXPECT_EQ(mirrored.anchor.y, request.anchor.y);
+  EXPECT_EQ(mirrored.anchor.z, request.anchor.z);
+  EXPECT_EQ(mirrored.radius, request.radius);
+  EXPECT_EQ(mirrored.azimuth, request.azimuth);
+  EXPECT_EQ(mirrored.elevation, request.elevation);
+  EXPECT_EQ(mirrored.orbitRadiansPerSecond, request.orbitRadiansPerSecond);
+}
+
+TEST(SKTypeConversionsTest, LensSettingsKeepOnlyTheFieldsWhoseFlagIsSet) {
+  SKLensSettings none{};
+  none.fieldOfViewDegrees = 90.0f;
+  none.nearPlane = 1.0f;
+  none.farPlane = 10.0f;
+  const LensSettings ignored = lensSettings(none);
+  EXPECT_FALSE(ignored.fieldOfViewDegrees.has_value());
+  EXPECT_FALSE(ignored.depthRange.has_value());
+
+  SKLensSettings set = none;
+  set.hasFieldOfView = YES;
+  set.hasDepthRange = YES;
+  const LensSettings applied = lensSettings(set);
+  ASSERT_TRUE(applied.fieldOfViewDegrees.has_value());
+  EXPECT_EQ(*applied.fieldOfViewDegrees, 90.0f);
+  ASSERT_TRUE(applied.depthRange.has_value());
+  EXPECT_EQ(applied.depthRange->nearPlane, 1.0f);
+  EXPECT_EQ(applied.depthRange->farPlane, 10.0f);
+}
+
+}  // namespace
+}  // namespace splatkit

From 156ab48cd0b4719e70d758082473752ba6b52f65 Mon Sep 17 00:00:00 2001
From: Xget7 <juanieltupa@gmail.com>
Date: Tue, 20 Oct 2026 01:23:09 -0300
Subject: [PATCH 6/8] Test destroying the Metal renderer with selection and
 sort buffers in flight

Fifty renderers each draw both frame slots on a 4096-leaf hierarchy, so a
selection, a sort and a render buffer are in flight per slot, and are
destroyed straight away. The test is a regression guard for the completion
handlers that capture pointers into MetalFrameCompute; it needs an
AddressSanitizer build to catch a late handler.
---
 .../rendering/metal/MetalSplatRendererTest.mm | 36 +++++++++++++++++++
 1 file changed, 36 insertions(+)

diff --git a/packages/splatkit-ios/tests/rendering/metal/MetalSplatRendererTest.mm b/packages/splatkit-ios/tests/rendering/metal/MetalSplatRendererTest.mm
index 0c97b6a..34ec819 100644
--- a/packages/splatkit-ios/tests/rendering/metal/MetalSplatRendererTest.mm
+++ b/packages/splatkit-ios/tests/rendering/metal/MetalSplatRendererTest.mm
@@ -350,5 +350,41 @@ void SetUp() override {
   }
 }
 
+TEST_F(MetalSplatRendererTest, DestroyingTheRendererWithSelectionAndSortInFlightIsSafe) {
+  constexpr uint32_t kLeaves = 4096;
+  constexpr int kRepeats = 50;
+  splat::LodTree tree;
+  tree.leafCount = kLeaves;
+  tree.nodes.positions = {0, 0, -8};
+  tree.nodes.colors = {1, 1, 1};
+  tree.nodes.alphas = {1};
+  tree.nodes.covariances = {0.04f, 0, 0, 0.04f, 0, 0.04f};
+  tree.layout.push_back({{0, 0, -8}, 0.4f, 1, kLeaves});
+  for (uint32_t i = 0; i < kLeaves; ++i) {
+    const float x = (static_cast<float>(i % 64) - 31.5f) * 0.02f;
+    const float y = (static_cast<float>(i / 64) - 31.5f) * 0.02f;
+    tree.nodes.positions.insert(tree.nodes.positions.end(), {x, y, -8});
+    tree.nodes.colors.insert(tree.nodes.colors.end(), {1, 1, 1});
+    tree.nodes.alphas.push_back(1);
+    tree.nodes.covariances.insert(tree.nodes.covariances.end(), {0.001f, 0, 0, 0.001f, 0, 0.001f});
+    tree.layout.push_back({{x, y, -8}, 0, 0, 0});
+  }
+  SplatRenderer::Frame frame;
+  frame.orderSource = SplatRenderer::OrderSource::gpu;
+  frame.proj = splat::Mat4::perspective(1, 1, 0.1f, 100);
+  for (int repeat = 0; repeat < kRepeats; ++repeat) {
+    auto renderer = MetalSplatRenderer::create();
+    ASSERT_NE(renderer, nullptr);
+    auto layer = [CAMetalLayer layer];
+    layer.drawableSize = CGSizeMake(64, 64);
+    renderer->setLayer(layer);
+    renderer->setDrawableSize(64, 64);
+    ASSERT_TRUE(renderer->uploadLodWorld(tree, 0, kLeaves));
+    // Both frame slots hold a selection, a sort and a render buffer when the renderer goes.
+    for (uint32_t i = 0; i < kFramesInFlight; ++i) ASSERT_TRUE(renderer->draw(frame));
+    renderer.reset();
+  }
+}
+
 }  // namespace
 }  // namespace splatkit

From ece2d5933fe6c1c13d9e6c08c342c3888648bf3d Mon Sep 17 00:00:00 2001
From: Xget7 <juanieltupa@gmail.com>
Date: Tue, 20 Oct 2026 01:23:48 -0300
Subject: [PATCH 7/8] Let the Metal selection and sort handlers outlive the
 frame compute

The selection and sort completion handlers wrote through raw pointers into
MetalFrameCompute and MetalFrameLoop, and ~MetalSplatRenderer waits only on
the frame slots that the render buffer's handler returns. Apple documents
no order between one command buffer's completion handlers and another's, so
a handler can still be running after the renderer has gone. In practice the
driver ran them in commit order here: a 30 ms sleep in the sort handler
under AddressSanitizer did not fail the regression test, which therefore
passes before and after this change.

The handlers now hold a share of a GpuReport that carries the counters, and
of the frame loop's failure latch, so nothing they touch can be freed.
---
 .../rendering/metal/frame/MetalFrameCompute.h | 32 ++++++-----
 .../metal/frame/MetalFrameCompute.mm          | 53 +++++++++----------
 .../metal/presentation/MetalFrameLoop.h       | 12 +++--
 .../metal/presentation/MetalFrameLoop.mm      |  4 +-
 4 files changed, 53 insertions(+), 48 deletions(-)

diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h
index ac1d907..1d5945b 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h
@@ -74,16 +74,29 @@ class MetalFrameCompute {
   id<MTLBuffer> order() const { return sort_.values(); }
   id<MTLBuffer> drawArguments(uint32_t slot) const { return visibility_.drawArguments(slot); }
 
-  double lastSortMillis() const { return lastSortMillis_.load(); }
-  uint32_t lastDrawCount() const { return lastDrawCount_.load(); }
-  uint32_t lastSelectedCount() const { return lastSelectedCount_.load(); }
-  double lastSelectMillis() const { return lastSelectMillis_.load(); }
-  uint32_t lastLodLimitedCount() const { return lod_ ? lastLodLimitedCount_.load() : 0; }
-  uint32_t lastLodEvaluatedCount() const { return lod_ ? lastLodEvaluatedCount_.load() : 0; }
+  double lastSortMillis() const { return report_->sortMillis.load(); }
+  uint32_t lastDrawCount() const { return report_->drawCount.load(); }
+  uint32_t lastSelectedCount() const { return report_->selectedCount.load(); }
+  double lastSelectMillis() const { return report_->selectMillis.load(); }
+  uint32_t lastLodLimitedCount() const { return lod_ ? report_->lodLimitedCount.load() : 0; }
+  uint32_t lastLodEvaluatedCount() const { return lod_ ? report_->lodEvaluatedCount.load() : 0; }
   // The threshold the next selection runs at, at least the policy's.
   float lodPixels() const { return lodPixels_; }
 
  private:
+  // What the selection and sort completion handlers write. The handlers hold a share of it:
+  // Metal documents no order between one command buffer's handlers and another's, so a handler
+  // can still be running after the render buffer's has released the last frame slot.
+  struct GpuReport {
+    std::atomic<uint32_t> selectedCount{0};
+    std::atomic<uint32_t> lodLimitedCount{0};
+    std::atomic<uint32_t> lodEvaluatedCount{0};
+    std::atomic<uint32_t> lodReadbacks{0};  // completed selections, so each adapts once
+    std::atomic<uint32_t> drawCount{0};
+    std::atomic<double> selectMillis{0};
+    std::atomic<double> sortMillis{0};
+  };
+
   bool createOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort,
                       KeyBits bits) const;
   void adaptLodThreshold();
@@ -103,14 +116,9 @@ class MetalFrameCompute {
   // The threshold selection runs at: the policy's, raised while a cut would exceed the splat
   // limit so detail thins evenly instead of stopping wherever traversal ran out of room.
   float lodPixels_ = 1.0f;
-  std::atomic<uint32_t> lodReadbacks_{0};  // completed selections, so each adapts once
   uint32_t lodAdaptedReadback_ = 0;
   uint64_t frames_ = 0;  // submitted so far, which also paces the LOD log
-  std::atomic<uint32_t> lastSelectedCount_{0};
-  std::atomic<uint32_t> lastLodLimitedCount_{0}, lastLodEvaluatedCount_{0};
-  std::atomic<double> lastSelectMillis_{0};
-  std::atomic<double> lastSortMillis_{0};
-  std::atomic<uint32_t> lastDrawCount_{0};
+  std::shared_ptr<GpuReport> report_ = std::make_shared<GpuReport>();
 };
 
 }  // namespace splatkit
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm
index fdd2ccd..5a43d4f 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm
@@ -77,7 +77,7 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
   lodSplatLimit_ = policy.lodSplatLimit;
   lodPixels_ = lodErrorPixels_;
   // Frames are idle here, so every readback so far describes the previous policy.
-  lodAdaptedReadback_ = lodReadbacks_.load();
+  lodAdaptedReadback_ = report_->lodReadbacks.load();
   if (lod_) {
     lod_->setPixelLimit(lodPixels_);
     lod_->setSplatLimit(lodSplatLimit_);
@@ -90,8 +90,8 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
 bool MetalFrameCompute::configureFlatWorld(uint32_t count) {
   if (gpuSort_ && !reserveOrdering(visibility_, sort_, count, 0)) return false;
   lod_.reset();
-  lastSelectedCount_.store(0);
-  lastSelectMillis_.store(0);
+  report_->selectedCount.store(0);
+  report_->selectMillis.store(0);
   return true;
 }
 
@@ -116,11 +116,11 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
   lod_ = std::move(lod);
   lodReadback_ = readback;
   lodPixels_ = lodErrorPixels_;
-  lodAdaptedReadback_ = lodReadbacks_.load();
-  lastLodLimitedCount_.store(0);
-  lastLodEvaluatedCount_.store(0);
-  lastSelectedCount_.store(0);
-  lastSelectMillis_.store(0);
+  lodAdaptedReadback_ = report_->lodReadbacks.load();
+  report_->lodLimitedCount.store(0);
+  report_->lodEvaluatedCount.store(0);
+  report_->selectedCount.store(0);
+  report_->selectMillis.store(0);
   return true;
 }
 
@@ -148,12 +148,12 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
   constexpr float kRaise = 1.1f;
   constexpr float kRelax = 1.03f;
   constexpr float kMaxPixels = 64.0f;
-  const uint32_t readbacks = lodReadbacks_.load();
+  const uint32_t readbacks = report_->lodReadbacks.load();
   if (readbacks == lodAdaptedReadback_) return;
   lodAdaptedReadback_ = readbacks;
-  if (lastLodLimitedCount_.load() > 0) {
+  if (report_->lodLimitedCount.load() > 0) {
     lodPixels_ = std::min(lodPixels_ * kRaise, std::max(kMaxPixels, lodErrorPixels_));
-  } else if (lastSelectedCount_.load() < lod_->limit() / 10 * 9) {
+  } else if (report_->selectedCount.load() < lod_->limit() / 10 * 9) {
     lodPixels_ = std::max(lodPixels_ / kRelax, lodErrorPixels_);
   }
   lod_->setPixelLimit(lodPixels_);
@@ -172,30 +172,26 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
          destinationOffset:0
                       size:sizeof(LodCounters)];
   [readback endEncoding];
-  std::atomic<uint32_t>* selected = &lastSelectedCount_;
-  std::atomic<double>* milliseconds = &lastSelectMillis_;
-  std::atomic<uint32_t>* limited = &lastLodLimitedCount_;
-  std::atomic<uint32_t>* evaluated = &lastLodEvaluatedCount_;
+  const std::shared_ptr<GpuReport> report = report_;
+  const std::shared_ptr<std::atomic<bool>> failed = frameLoop_.failureLatch();
   const bool logLod = frames_ % kLodLogFrames == 0;
-  std::atomic<uint32_t>* readbacks = &lodReadbacks_;
   const float pixels = lodPixels_;
-  MetalFrameLoop* loop = &frameLoop_;
   [selection addCompletedHandler:^(id<MTLCommandBuffer> done) {
     if (done.status == MTLCommandBufferStatusError) {
-      loop->markFailed();
+      failed->store(true);
       LOGE("LOD command failed: %s", done.error.localizedDescription.UTF8String);
       return;
     }
     const auto& counters = *static_cast<const LodCounters*>(selectedCount.contents);
-    selected->store(counters.count);
-    limited->store(counters.limited);
-    evaluated->store(counters.evaluated);
-    readbacks->fetch_add(1);
+    report->selectedCount.store(counters.count);
+    report->lodLimitedCount.store(counters.limited);
+    report->lodEvaluatedCount.store(counters.evaluated);
+    report->lodReadbacks.fetch_add(1);
     if (logLod)
       LOGI("LOD SSE: %u selected, %u evaluated interiors, %u quality-limited refinements, "
            "%.2f px",
            counters.count, counters.evaluated, counters.limited, pixels);
-    milliseconds->store((done.GPUEndTime - done.GPUStartTime) * 1000.0);
+    report->selectMillis.store((done.GPUEndTime - done.GPUStartTime) * 1000.0);
   }];
   return selection;
 }
@@ -214,18 +210,17 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
   }
   sort_.encode(sort, visibility_.countBuffer(slot), visibility_.depthBits());
   visibility_.encodeIndirect(sort, slot);
-  std::atomic<double>* sortMillis = &lastSortMillis_;
-  std::atomic<uint32_t>* drawn = &lastDrawCount_;
+  const std::shared_ptr<GpuReport> report = report_;
+  const std::shared_ptr<std::atomic<bool>> failed = frameLoop_.failureLatch();
   id<MTLBuffer> countBuffer = visibility_.countBuffer(slot);
-  MetalFrameLoop* loop = &frameLoop_;
   [sort addCompletedHandler:^(id<MTLCommandBuffer> done) {
     if (done.status == MTLCommandBufferStatusError) {
-      loop->markFailed();
+      failed->store(true);
       LOGE("visibility command failed: %s", done.error.localizedDescription.UTF8String);
       return;
     }
-    sortMillis->store((done.GPUEndTime - done.GPUStartTime) * 1000.0);
-    drawn->store(*static_cast<const uint32_t*>(countBuffer.contents));
+    report->sortMillis.store((done.GPUEndTime - done.GPUStartTime) * 1000.0);
+    report->drawCount.store(*static_cast<const uint32_t*>(countBuffer.contents));
   }];
   return sort;
 }
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.h
index f374bd8..6309eba 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.h
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.h
@@ -41,15 +41,17 @@ class MetalFrameLoop {
   // Fills `times` with the display times, in nanoseconds, of the frames presented since the
   // last call and returns how many frames were dropped. Zero where the platform reports none.
   uint32_t takePresentTimes(std::vector<int64_t>* times);
-  double lastGpuMillis() const { return gpuFailed_.load() ? 0 : lastGpuMillis_.load(); }
+  double lastGpuMillis() const { return gpuFailed_->load() ? 0 : lastGpuMillis_.load(); }
 
   // A GPU error latches this loop off. Never repeatedly resubmit failed work.
-  bool gpuFailed() const { return gpuFailed_.load(); }
+  bool gpuFailed() const { return gpuFailed_->load(); }
   // For work that fails outside submit(), such as the selection and visibility buffers.
-  void markFailed() { gpuFailed_.store(true); }
+  void markFailed() { gpuFailed_->store(true); }
+  // The latch itself, for completion handlers that can run after the loop is destroyed.
+  std::shared_ptr<std::atomic<bool>> failureLatch() const { return gpuFailed_; }
   // A successful GPU frame of the current world has finished; uploads alone and background
   // frames do not qualify.
-  bool hasCompletedWorldFrame() const { return !gpuFailed_.load() && completedWorldFrame_.load(); }
+  bool hasCompletedWorldFrame() const { return !gpuFailed_->load() && completedWorldFrame_.load(); }
   // The world was replaced.
   void clearCompletedWorldFrame() { completedWorldFrame_.store(false); }
 
@@ -65,7 +67,7 @@ class MetalFrameLoop {
   dispatch_semaphore_t inFlight_;
   std::shared_ptr<PresentLog> presents_ = std::make_shared<PresentLog>();
   std::atomic<double> lastGpuMillis_{0};
-  std::atomic<bool> gpuFailed_{false};
+  std::shared_ptr<std::atomic<bool>> gpuFailed_ = std::make_shared<std::atomic<bool>>(false);
   std::atomic<bool> completedWorldFrame_{false};
 };
 
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.mm
index 4a9067e..f9efd0c 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.mm
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.mm
@@ -22,7 +22,7 @@
 
 bool MetalFrameLoop::acquire() {
   dispatch_semaphore_wait(inFlight_, DISPATCH_TIME_FOREVER);
-  if (gpuFailed_.load()) {
+  if (gpuFailed_->load()) {
     dispatch_semaphore_signal(inFlight_);
     return false;
   }
@@ -66,7 +66,7 @@
   [cmd presentDrawable:drawable];
   dispatch_semaphore_t inFlight = inFlight_;
   std::atomic<double>* gpuMillis = &lastGpuMillis_;
-  std::atomic<bool>* failed = &gpuFailed_;
+  std::atomic<bool>* failed = gpuFailed_.get();
   auto* worldFrame = &completedWorldFrame_;
   [cmd addCompletedHandler:^(id<MTLCommandBuffer> done) {
     if (done.status == MTLCommandBufferStatusError) {

From 4cf7e392946f0abdd2d719134b807ecd84847df9 Mon Sep 17 00:00:00 2001
From: Xget7 <juanieltupa@gmail.com>
Date: Tue, 20 Oct 2026 01:24:09 -0300
Subject: [PATCH 8/8] Apply clang-format to the Metal seam and test changes

---
 .../rendering/metal/MetalSplatRenderer.mm      |  5 ++---
 .../rendering/metal/cull/MetalVisibilityPass.h |  3 +--
 .../metal/cull/MetalVisibilityPass.mm          |  3 +--
 .../rendering/metal/frame/MetalFrameCompute.h  |  3 +--
 .../rendering/metal/frame/MetalFrameCompute.mm |  4 +---
 .../tests/diagnostics/SKPngWriterTest.mm       | 18 ++++++++++--------
 .../metal/cull/MetalVisibilityPassTest.mm      |  4 ++--
 7 files changed, 18 insertions(+), 22 deletions(-)

diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm
index 08be6c4..ca20298 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm
@@ -96,9 +96,8 @@
   policy.fallback.lodErrorPixels = 1.0f;
   policy.fallback.alphaThreshold = kQuadMinAlpha;
   policy.fallback.subpixelThreshold = kMinPixelRadius;
-  policy.fallback.sortDepth = compute_->depthBits() == KeyBits::Low16
-                                  ? SortKeyBits::low16
-                                  : SortKeyBits::full32;
+  policy.fallback.sortDepth =
+      compute_->depthBits() == KeyBits::Low16 ? SortKeyBits::low16 : SortKeyBits::full32;
   policy.fallback.enableFrustumCulling = true;
   policy.fallback.enableEarlyTermination = true;
   // The key width is a public per-instance control. The raster strategy, sub-pixel radius,
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.h
index e0b4adf..58a6a21 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.h
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.h
@@ -21,8 +21,7 @@ class MetalVisibilityPass {
  public:
   // Both widths sort by linear view depth: Full32 by its float bits, Low16 quantized between
   // finite forward-Z Mat4::perspective planes. Shader key width and radix passes agree.
-  bool create(id<MTLDevice> device, id<MTLLibrary> library,
-              KeyBits depthBits = KeyBits::Full32);
+  bool create(id<MTLDevice> device, id<MTLLibrary> library, KeyBits depthBits = KeyBits::Full32);
   // Rebuilds the visibility pipelines for a new key width, keeping every reserved buffer
   // and the radix scratch. The key width and the shader's quantized flag must agree, so
   // both move together. Render thread, while idle. On failure the previous key width is
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.mm
index 7be179c..bc72dff 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.mm
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.mm
@@ -8,8 +8,7 @@
 
 namespace splatkit {
 
-bool MetalVisibilityPass::create(id<MTLDevice> device, id<MTLLibrary> library,
-                                 KeyBits depthBits) {
+bool MetalVisibilityPass::create(id<MTLDevice> device, id<MTLLibrary> library, KeyBits depthBits) {
   device_ = device;
   library_ = library;
   depthBits_ = depthBits;
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h
index 1d5945b..e08d2ef 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h
@@ -97,8 +97,7 @@ class MetalFrameCompute {
     std::atomic<double> sortMillis{0};
   };
 
-  bool createOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort,
-                      KeyBits bits) const;
+  bool createOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, KeyBits bits) const;
   void adaptLodThreshold();
 
   id<MTLDevice> device_ = nil;
diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm
index 5a43d4f..fcf53a4 100644
--- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm
+++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm
@@ -54,9 +54,7 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
 }
 
 bool MetalFrameCompute::configure(const RenderPolicy& policy, std::string* reason) {
-  const KeyBits bits = policy.sortDepth == SortKeyBits::low16
-                                           ? KeyBits::Low16
-                                           : KeyBits::Full32;
+  const KeyBits bits = policy.sortDepth == SortKeyBits::low16 ? KeyBits::Low16 : KeyBits::Full32;
   const bool visibilityChanged = bits != depthBits();
   const bool lodChanged =
       policy.lodErrorPixels != lodErrorPixels_ || policy.lodSplatLimit != lodSplatLimit_;
diff --git a/packages/splatkit-ios/tests/diagnostics/SKPngWriterTest.mm b/packages/splatkit-ios/tests/diagnostics/SKPngWriterTest.mm
index 3f3af75..a456b6a 100644
--- a/packages/splatkit-ios/tests/diagnostics/SKPngWriterTest.mm
+++ b/packages/splatkit-ios/tests/diagnostics/SKPngWriterTest.mm
@@ -35,12 +35,12 @@ void SetUp() override {
     *height = CGImageGetHeight(image);
     rgba.assign(*width * *height * kBytesPerPixel, 0);
     CGColorSpaceRef space = CGColorSpaceCreateDeviceRGB();
-    CGContextRef context = CGBitmapContextCreate(rgba.data(), *width, *height, 8,
-                                                 *width * kBytesPerPixel, space,
-                                                 kCGImageAlphaPremultipliedLast);
-    CGContextDrawImage(context, CGRectMake(0, 0, static_cast<CGFloat>(*width),
-                                           static_cast<CGFloat>(*height)),
-                       image);
+    CGContextRef context =
+        CGBitmapContextCreate(rgba.data(), *width, *height, 8, *width * kBytesPerPixel, space,
+                              kCGImageAlphaPremultipliedLast);
+    CGContextDrawImage(
+        context, CGRectMake(0, 0, static_cast<CGFloat>(*width), static_cast<CGFloat>(*height)),
+        image);
     CGContextRelease(context);
     CGColorSpaceRelease(space);
     CGImageRelease(image);
@@ -53,10 +53,12 @@ void SetUp() override {
   constexpr uint32_t kWidth = 3;
   constexpr uint32_t kHeight = 2;
   // BGRA, top row first: blue, green, red / white, black, mid grey.
+  // clang-format off
   const std::vector<uint8_t> bgra = {
-      255, 0,   0,   255, 0,   255, 0,   255, 0,   0,   255, 255,
-      255, 255, 255, 255, 0,   0,   0,   255, 128, 128, 128, 255,
+      255, 0,   0,   255,   0, 255,   0, 255,   0,   0, 255, 255,
+      255, 255, 255, 255,   0,   0,   0, 255, 128, 128, 128, 255,
   };
+  // clang-format on
   ASSERT_EQ(bgra.size(), size_t{kWidth} * kHeight * kBytesPerPixel);
   ASSERT_TRUE(writePng(path_, bgra, kWidth, kHeight));
 
diff --git a/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm b/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm
index 411d51f..13e949c 100644
--- a/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm
+++ b/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm
@@ -121,8 +121,8 @@ explicit VisibilityInput(uint32_t count) {
   id<MTLCommandBuffer> cmd = [Gpu::get().queue commandBuffer];
   EXPECT_FALSE(visibility.encode(cmd, 0, input.uniforms, input.splats, nil, 0, &range, 1));
   EXPECT_FALSE(visibility.encode(cmd, 0, input.uniforms, input.splats, nil, 0, nullptr, 1));
-  EXPECT_FALSE(visibility.encode(cmd, kFramesInFlight, input.uniforms, input.splats,
-                                 nil, 0, nullptr, 0));
+  EXPECT_FALSE(
+      visibility.encode(cmd, kFramesInFlight, input.uniforms, input.splats, nil, 0, nullptr, 0));
   [cmd commit];
   [cmd waitUntilCompleted];
 }