Follow up the Metal seams review: pin the handler lifetime and tidy the minors

packages/splatkit-ios/Sources/SplatKitCore/engine/SKTypeConversions.mm

@@ -18,7 +18,8 @@ SKRenderPolicy fromCppPolicy(const RenderPolicy& p) {
 }
 
 bool toCppPolicy(SKRenderPolicy p, RenderPolicy* out) {
-  if (p.sortDepth != 16 && p.sortDepth != 32) return false;
+  const auto sortDepth = static_cast<SortKeyBits>(p.sortDepth);
+  if (sortDepth != SortKeyBits::low16 && sortDepth != SortKeyBits::full32) return false;
   out->raster = static_cast<RasterStrategy>(p.raster);
   out->tileSize = p.tileSize;
   out->lodErrorPixels = p.lodErrorPixels;
@@ -27,7 +28,7 @@ bool toCppPolicy(SKRenderPolicy p, RenderPolicy* out) {
   out->enableFrustumCulling = p.enableFrustumCulling == YES;
   out->enableHiZOcclusion = p.enableHiZOcclusion == YES;
   out->enableEarlyTermination = p.enableEarlyTermination == YES;
-  out->sortDepth = p.sortDepth == 16 ? SortKeyBits::low16 : SortKeyBits::full32;
+  out->sortDepth = sortDepth;
   out->lodSplatLimit = p.lodSplatLimit;
   return true;
 }

packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm

@@ -15,6 +15,8 @@
 // The LOD error threshold range, in pixels, a policy may request.
 constexpr float kMinLodErrorPixels = 0.1f;
 constexpr float kMaxLodErrorPixels = 16.0f;
+constexpr uint32_t kFallbackTileSize = 16;
+constexpr float kFallbackLodErrorPixels = 1.0f;
 // Apple family 7, the minimum this renderer accepts, supports 16384x16384 textures.
 constexpr uint32_t kAppleFamily7MaxTextureDimension = 16384;
 
@@ -92,8 +94,8 @@
   caps.maxTextureDimension = context_->device() != nil ? kAppleFamily7MaxTextureDimension : 0u;
   RenderPolicySupport& policy = caps.policy;
   policy.fallback.raster = RasterStrategy::hardware;
-  policy.fallback.tileSize = 16;
-  policy.fallback.lodErrorPixels = 1.0f;
+  policy.fallback.tileSize = kFallbackTileSize;
+  policy.fallback.lodErrorPixels = kFallbackLodErrorPixels;
   policy.fallback.alphaThreshold = kQuadMinAlpha;
   policy.fallback.subpixelThreshold = kMinPixelRadius;
   policy.fallback.sortDepth =

packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h

@@ -7,6 +7,7 @@
 #include <cstdint>
 #include <memory>
 #include <string>
+#include <utility>
 
 #include "rendering/metal/cull/MetalVisibilityPass.h"
 #include "rendering/metal/layout/MetalFrameSlots.h"
@@ -84,10 +85,14 @@ class MetalFrameCompute {
   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.
+  // What the selection and sort completion handlers write, with the frame loop's failure latch
+  // they set. 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 {
+    explicit GpuReport(std::shared_ptr<std::atomic<bool>> failureLatch)
+        : failed(std::move(failureLatch)) {}
+    const std::shared_ptr<std::atomic<bool>> failed;
     std::atomic<uint32_t> selectedCount{0};
     std::atomic<uint32_t> lodLimitedCount{0};
     std::atomic<uint32_t> lodEvaluatedCount{0};
@@ -97,6 +102,8 @@ class MetalFrameCompute {
     std::atomic<double> sortMillis{0};
   };
 
+  static constexpr float kDefaultLodPixels = 1.0f;
+
   bool createOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, KeyBits bits) const;
   void adaptLodThreshold();
 
@@ -110,14 +117,14 @@ class MetalFrameCompute {
   bool gpuSort_ = false;
   std::unique_ptr<MetalLodSelection> lod_;
   std::array<id<MTLBuffer>, kFramesInFlight> lodReadback_{};
-  float lodErrorPixels_ = 1.0f;
+  float lodErrorPixels_ = kDefaultLodPixels;
   uint32_t lodSplatLimit_ = 0;
   // 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;
+  float lodPixels_ = kDefaultLodPixels;
   uint32_t lodAdaptedReadback_ = 0;
   uint64_t frames_ = 0;  // submitted so far, which also paces the LOD log
-  std::shared_ptr<GpuReport> report_ = std::make_shared<GpuReport>();
+  std::shared_ptr<GpuReport> report_;
 };
 
 }  // namespace splatkit

packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm

@@ -34,7 +34,11 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
 
 MetalFrameCompute::MetalFrameCompute(id<MTLDevice> device, id<MTLCommandQueue> queue,
                                      id<MTLLibrary> library, MetalFrameLoop& frameLoop)
-    : device_(device), queue_(queue), library_(library), frameLoop_(frameLoop) {}
+    : device_(device),
+      queue_(queue),
+      library_(library),
+      frameLoop_(frameLoop),
+      report_(std::make_shared<GpuReport>(frameLoop.failureLatch())) {}
 
 bool MetalFrameCompute::createOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort,
                                        KeyBits bits) const {
@@ -171,12 +175,11 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
                       size:sizeof(LodCounters)];
   [readback endEncoding];
   const std::shared_ptr<GpuReport> report = report_;
-  const std::shared_ptr<std::atomic<bool>> failed = frameLoop_.failureLatch();
   const bool logLod = frames_ % kLodLogFrames == 0;
   const float pixels = lodPixels_;
   [selection addCompletedHandler:^(id<MTLCommandBuffer> done) {
     if (done.status == MTLCommandBufferStatusError) {
-      failed->store(true);
+      report->failed->store(true);
       LOGE("LOD command failed: %s", done.error.localizedDescription.UTF8String);
       return;
     }
@@ -209,11 +212,10 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint
   sort_.encode(sort, visibility_.countBuffer(slot), visibility_.depthBits());
   visibility_.encodeIndirect(sort, slot);
   const std::shared_ptr<GpuReport> report = report_;
-  const std::shared_ptr<std::atomic<bool>> failed = frameLoop_.failureLatch();
   id<MTLBuffer> countBuffer = visibility_.countBuffer(slot);
   [sort addCompletedHandler:^(id<MTLCommandBuffer> done) {
     if (done.status == MTLCommandBufferStatusError) {
-      failed->store(true);
+      report->failed->store(true);
       LOGE("visibility command failed: %s", done.error.localizedDescription.UTF8String);
       return;
     }

packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalFrameSlots.h

@@ -4,8 +4,7 @@
 
 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.
+// The frames the CPU may encode ahead of the GPU.
 inline constexpr uint32_t kFramesInFlight = 2;
 
 }  // namespace splatkit

packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.h

@@ -45,9 +45,8 @@ class MetalFrameLoop {
 
   // A GPU error latches this loop off. Never repeatedly resubmit failed work.
   bool gpuFailed() const { return gpuFailed_->load(); }
-  // For work that fails outside submit(), such as the selection and visibility buffers.
-  void markFailed() { gpuFailed_->store(true); }
-  // The latch itself, for completion handlers that can run after the loop is destroyed.
+  // The latch itself. Work that fails outside submit(), such as the selection and visibility
+  // buffers, sets it from 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.

packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.mm

@@ -64,6 +64,8 @@
   }];
 #endif
   [cmd presentDrawable:drawable];
+  // The handler touches the loop's members through these raw pointers, so signalling the
+  // semaphore must stay its last action: ~MetalFrameLoop returns only after waitIdle().
   dispatch_semaphore_t inFlight = inFlight_;
   std::atomic<double>* gpuMillis = &lastGpuMillis_;
   std::atomic<bool>* failed = gpuFailed_.get();

packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalSurface.h

@@ -36,7 +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.
+  // Clamped to [kMinRenderScale, kMaxRenderScale].
   void setRenderScale(float scale);
   void setSpatialUpscaling(bool enabled);
   // False when the blit pipeline for the new format could not be built.

packages/splatkit-ios/tests/engine/SKTypeConversionsTest.mm

@@ -208,35 +208,6 @@ void expectSamePolicy(const RenderPolicy& actual, const RenderPolicy& expected)
   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;

packages/splatkit-ios/tests/rendering/metal/MetalSplatRendererTest.mm

@@ -350,7 +350,9 @@ void SetUp() override {
   }
 }
 
-TEST_F(MetalSplatRendererTest, DestroyingTheRendererWithSelectionAndSortInFlightIsSafe) {
+// A smoke test: it fails only on a crash or a hang. MetalFrameComputeTest pins the handler
+// lifetime this relies on.
+TEST_F(MetalSplatRendererTest, DestroyingTheRendererWithFramesInFlightDoesNotCrash) {
   constexpr uint32_t kLeaves = 4096;
   constexpr int kRepeats = 50;
   splat::LodTree tree;

packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm

@@ -45,6 +45,50 @@ void SetUp() override {
   EXPECT_EQ(rig.pass.capacity(), 16u);
 }
 
+TEST(MetalPrepareDrawArgumentsTest, WritesExactlyTheDrawBatches) {
+  auto& gpu = Gpu::get();
+  id<MTLComputePipelineState> prepare =
+      metal::pipeline(gpu.device, gpu.library, "prepareDrawArguments");
+  ASSERT_NE(prepare, nil);
+  constexpr uint32_t kSentinel = 0xdeadbeefu;
+  constexpr uint32_t kGuardRecords = 3;
+  constexpr uint32_t kWords = kDrawArgumentBytes / sizeof(uint32_t);
+  std::vector<uint32_t> initial((kDrawBatches + kGuardRecords) * kWords, kSentinel);
+  id<MTLBuffer> draws = [gpu.device newBufferWithBytes:initial.data()
+                                                length:initial.size() * sizeof(uint32_t)
+                                               options:MTLResourceStorageModeShared];
+  const uint32_t splats = 64;
+  id<MTLBuffer> count = [gpu.device newBufferWithBytes:&splats
+                                                length:sizeof(splats)
+                                               options:MTLResourceStorageModeShared];
+  id<MTLCommandBuffer> cmd = [gpu.queue commandBuffer];
+  id<MTLComputeCommandEncoder> encoder = [cmd computeCommandEncoder];
+  [encoder setComputePipelineState:prepare];
+  [encoder setBuffer:count offset:0 atIndex:0];
+  [encoder setBuffer:draws offset:0 atIndex:1];
+  [encoder dispatchThreadgroups:MTLSizeMake(1, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
+  [encoder endEncoding];
+  [cmd commit];
+  [cmd waitUntilCompleted];
+  ASSERT_EQ(cmd.status, MTLCommandBufferStatusCompleted);
+
+  const auto* records = static_cast<const MTLDrawPrimitivesIndirectArguments*>(draws.contents);
+  // 64 splats split at 1/64, 2/64, 4/64 ... of the set.
+  const uint32_t expectedInstances[kDrawBatches] = {1, 1, 2, 4, 8, 16, 32};
+  uint32_t start = 0;
+  for (uint32_t b = 0; b < kDrawBatches; ++b) {
+    EXPECT_EQ(records[b].vertexCount, 4u);
+    EXPECT_EQ(records[b].instanceCount, expectedInstances[b]);
+    EXPECT_EQ(records[b].vertexStart, 0u);
+    EXPECT_EQ(records[b].baseInstance, start);
+    start += expectedInstances[b];
+  }
+  const auto* words = static_cast<const uint32_t*>(draws.contents);
+  for (uint32_t i = kDrawBatches * kWords; i < initial.size(); ++i) {
+    EXPECT_EQ(words[i], kSentinel) << "word " << i << " is past the last draw batch";
+  }
+}
+
 struct VisibilityInput {
   id<MTLBuffer> splats;
   id<MTLBuffer> uniforms;

packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm

@@ -238,5 +238,40 @@ void run(id<MTLCommandBuffer> selection, id<MTLCommandBuffer> sort) {
   EXPECT_EQ(compute_->lastLodLimitedCount(), 0u);
 }
 
+// The completion handlers run when the GPU finishes, which can be after the renderer that
+// encoded them is gone. They are encoded and left uncommitted here, so the compute and the loop
+// are certainly destroyed before any handler runs.
+TEST_F(MetalFrameComputeTest,
+       SelectionAndSortHandlersKeepTheStateTheyWriteAliveAfterTheComputeIsGone) {
+  auto tree = sixtyFourLeaves();
+  auto world = MetalWorld::upload(gpu_->device, tree.nodes, 0);
+  ASSERT_NE(world, nullptr);
+  auto loop = std::make_unique<MetalFrameLoop>(kFramesInFlight);
+  auto compute =
+      std::make_unique<MetalFrameCompute>(gpu_->device, gpu_->queue, gpu_->library, *loop);
+  ASSERT_TRUE(compute->create());
+  ASSERT_TRUE(compute->configureLodWorld(tree, 64));
+  compute->updateCamera(frame_, 0, kExtent, false);
+  id<MTLCommandBuffer> selection = compute->encodeSelection(0);
+  id<MTLCommandBuffer> sort = compute->encodeVisibilityAndSort(frame_, 0, *world);
+  ASSERT_NE(selection, nil);
+  ASSERT_NE(sort, nil);
+
+  const std::shared_ptr<std::atomic<bool>> failed = loop->failureLatch();
+  compute.reset();
+  loop.reset();
+  // The report owns the latch and only the handlers still hold the report, so the latch has a
+  // second owner exactly while a handler owns the report.
+  ASSERT_GE(failed.use_count(), 2);
+
+  [selection commit];
+  [sort commit];
+  [selection waitUntilCompleted];
+  [sort waitUntilCompleted];
+  EXPECT_EQ(selection.status, MTLCommandBufferStatusCompleted);
+  EXPECT_EQ(sort.status, MTLCommandBufferStatusCompleted);
+  EXPECT_FALSE(failed->load());
+}
+
 }  // namespace
 }  // namespace splatkit