From fd45a320e89bed985ea8c19d3060e9c73570efaf Mon Sep 17 00:00:00 2001 From: Xget7 Date: Tue, 20 Oct 2026 01:54:25 -0300 Subject: [PATCH 1/7] Pin the Metal completion handler lifetime with a frame compute test that cannot pass without the fix --- .../rendering/metal/MetalSplatRendererTest.mm | 4 ++- .../metal/frame/MetalFrameComputeTest.mm | 34 +++++++++++++++++++ 2 files changed, 37 insertions(+), 1 deletion(-) diff --git a/packages/splatkit-ios/tests/rendering/metal/MetalSplatRendererTest.mm b/packages/splatkit-ios/tests/rendering/metal/MetalSplatRendererTest.mm index 34ec819..b1c9fcc 100644 --- a/packages/splatkit-ios/tests/rendering/metal/MetalSplatRendererTest.mm +++ b/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; diff --git a/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm b/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm index 180565c..5f7ed5b 100644 --- a/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm +++ b/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm @@ -238,5 +238,39 @@ void run(id selection, id 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(kFramesInFlight); + auto compute = + std::make_unique(gpu_->device, gpu_->queue, gpu_->library, *loop); + ASSERT_TRUE(compute->create()); + ASSERT_TRUE(compute->configureLodWorld(tree, 64)); + compute->updateCamera(frame_, 0, kExtent, false); + id selection = compute->encodeSelection(0); + id sort = compute->encodeVisibilityAndSort(frame_, 0, *world); + ASSERT_NE(selection, nil); + ASSERT_NE(sort, nil); + + const std::shared_ptr> failed = loop->failureLatch(); + compute.reset(); + loop.reset(); + // Only the two handlers and this test still hold the latch. + ASSERT_EQ(failed.use_count(), 3); + + [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 From 8f5e9374e449545d515a818ac31403a37f5d1cc0 Mon Sep 17 00:00:00 2001 From: Xget7 Date: Tue, 20 Oct 2026 01:54:25 -0300 Subject: [PATCH 2/7] Drop markFailed and size the frame loop tests from kFramesInFlight --- .../metal/presentation/MetalFrameLoop.h | 5 ++- .../metal/presentation/MetalFrameLoopTest.mm | 34 +++++++++---------- 2 files changed, 18 insertions(+), 21 deletions(-) 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 6309eba..b4f91cd 100644 --- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.h +++ b/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> failureLatch() const { return gpuFailed_; } // A successful GPU frame of the current world has finished; uploads alone and background // frames do not qualify. diff --git a/packages/splatkit-ios/tests/rendering/metal/presentation/MetalFrameLoopTest.mm b/packages/splatkit-ios/tests/rendering/metal/presentation/MetalFrameLoopTest.mm index c75ecd2..c5c336d 100644 --- a/packages/splatkit-ios/tests/rendering/metal/presentation/MetalFrameLoopTest.mm +++ b/packages/splatkit-ios/tests/rendering/metal/presentation/MetalFrameLoopTest.mm @@ -3,13 +3,12 @@ #include #include +#include "rendering/metal/layout/MetalFrameSlots.h" #include "rendering/metal/presentation/MetalFrameLoop.h" namespace splatkit { namespace { -constexpr uint32_t kSlots = 2; - class MetalFrameLoopTest : public testing::Test { protected: void SetUp() override { @@ -49,26 +48,25 @@ static bool idleWithin(MetalFrameLoop& loop, std::chrono::milliseconds limit) { }; TEST_F(MetalFrameLoopTest, AnAcquiredSlotHoldsWaitIdleUntilItIsReleased) { - MetalFrameLoop loop(kSlots); - ASSERT_TRUE(loop.acquire()); - ASSERT_TRUE(loop.acquire()); + MetalFrameLoop loop(kFramesInFlight); + for (uint32_t i = 0; i < kFramesInFlight; ++i) ASSERT_TRUE(loop.acquire()); auto idle = std::async(std::launch::async, [&] { loop.waitIdle(); }); EXPECT_EQ(idle.wait_for(std::chrono::milliseconds(100)), std::future_status::timeout); - loop.release(); - EXPECT_EQ(idle.wait_for(std::chrono::milliseconds(100)), std::future_status::timeout); + for (uint32_t i = 0; i + 1 < kFramesInFlight; ++i) { + loop.release(); + EXPECT_EQ(idle.wait_for(std::chrono::milliseconds(100)), std::future_status::timeout); + } loop.release(); EXPECT_EQ(idle.wait_for(std::chrono::seconds(5)), std::future_status::ready); // waitIdle gives every slot back. - ASSERT_TRUE(loop.acquire()); - ASSERT_TRUE(loop.acquire()); - loop.release(); - loop.release(); + for (uint32_t i = 0; i < kFramesInFlight; ++i) ASSERT_TRUE(loop.acquire()); + for (uint32_t i = 0; i < kFramesInFlight; ++i) loop.release(); } TEST_F(MetalFrameLoopTest, AFailedGpuRefusesFramesAndHoldsNoSlot) { - MetalFrameLoop loop(kSlots); + MetalFrameLoop loop(kFramesInFlight); EXPECT_FALSE(loop.gpuFailed()); - loop.markFailed(); + loop.failureLatch()->store(true); EXPECT_TRUE(loop.gpuFailed()); EXPECT_FALSE(loop.acquire()); EXPECT_TRUE(idleWithin(loop, std::chrono::seconds(5))); @@ -76,7 +74,7 @@ static bool idleWithin(MetalFrameLoop& loop, std::chrono::milliseconds limit) { } TEST_F(MetalFrameLoopTest, ASubmittedFrameRecordsItsTimeAndReleasesItsSlotAfterTheCallback) { - MetalFrameLoop loop(kSlots); + MetalFrameLoop loop(kFramesInFlight); dispatch_semaphore_t done = dispatch_semaphore_create(0); MTLCommandBufferStatus status = MTLCommandBufferStatusNotEnqueued; ASSERT_TRUE(loop.acquire()); @@ -91,7 +89,7 @@ static bool idleWithin(MetalFrameLoop& loop, std::chrono::milliseconds limit) { } TEST_F(MetalFrameLoopTest, ABackgroundFrameDoesNotCountAsAWorldFrame) { - MetalFrameLoop loop(kSlots); + MetalFrameLoop loop(kFramesInFlight); dispatch_semaphore_t done = dispatch_semaphore_create(0); MTLCommandBufferStatus status = MTLCommandBufferStatusNotEnqueued; ASSERT_TRUE(loop.acquire()); @@ -102,7 +100,7 @@ static bool idleWithin(MetalFrameLoop& loop, std::chrono::milliseconds limit) { } TEST_F(MetalFrameLoopTest, AFailureAfterAWorldFrameWithdrawsIt) { - MetalFrameLoop loop(kSlots); + MetalFrameLoop loop(kFramesInFlight); dispatch_semaphore_t done = dispatch_semaphore_create(0); MTLCommandBufferStatus status = MTLCommandBufferStatusNotEnqueued; ASSERT_TRUE(loop.acquire()); @@ -110,13 +108,13 @@ static bool idleWithin(MetalFrameLoop& loop, std::chrono::milliseconds limit) { ASSERT_EQ(dispatch_semaphore_wait(done, dispatch_time(DISPATCH_TIME_NOW, 5 * NSEC_PER_SEC)), 0); EXPECT_TRUE(idleWithin(loop, std::chrono::seconds(5))); ASSERT_TRUE(loop.hasCompletedWorldFrame()); - loop.markFailed(); + loop.failureLatch()->store(true); EXPECT_FALSE(loop.hasCompletedWorldFrame()); EXPECT_EQ(loop.lastGpuMillis(), 0.0); } TEST_F(MetalFrameLoopTest, TakingPresentTimesClearsTheCallersVectorFirst) { - MetalFrameLoop loop(kSlots); + MetalFrameLoop loop(kFramesInFlight); std::vector times = {1, 2, 3}; EXPECT_EQ(loop.takePresentTimes(×), 0u); EXPECT_TRUE(times.empty()); From b60e65397bfd672b24dafea3f909a58e4394b3e3 Mon Sep 17 00:00:00 2001 From: Xget7 Date: Tue, 20 Oct 2026 01:54:25 -0300 Subject: [PATCH 3/7] Name the fallback policy and default LOD literals and use the SortKeyBits values --- .../Sources/SplatKitCore/engine/SKTypeConversions.mm | 5 +++-- .../SplatKitCore/rendering/metal/MetalSplatRenderer.mm | 6 ++++-- .../SplatKitCore/rendering/metal/frame/MetalFrameCompute.h | 6 ++++-- 3 files changed, 11 insertions(+), 6 deletions(-) diff --git a/packages/splatkit-ios/Sources/SplatKitCore/engine/SKTypeConversions.mm b/packages/splatkit-ios/Sources/SplatKitCore/engine/SKTypeConversions.mm index 49a6248..0f0bebc 100644 --- a/packages/splatkit-ios/Sources/SplatKitCore/engine/SKTypeConversions.mm +++ b/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(p.sortDepth); + if (sortDepth != SortKeyBits::low16 && sortDepth != SortKeyBits::full32) return false; out->raster = static_cast(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; } diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm index ca20298..0ab2320 100644 --- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.mm +++ b/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 = 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 e08d2ef..27fdab4 100644 --- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h +++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h @@ -97,6 +97,8 @@ class MetalFrameCompute { std::atomic sortMillis{0}; }; + static constexpr float kDefaultLodPixels = 1.0f; + bool createOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, KeyBits bits) const; void adaptLodThreshold(); @@ -110,11 +112,11 @@ class MetalFrameCompute { bool gpuSort_ = false; std::unique_ptr lod_; std::array, 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 report_ = std::make_shared(); From e3bf62a8d7e685ba9c7310a1d456a613cf831d70 Mon Sep 17 00:00:00 2001 From: Xget7 Date: Tue, 20 Oct 2026 01:54:25 -0300 Subject: [PATCH 4/7] Trim the surface and slot comments and state the render handler's last action --- .../SplatKitCore/rendering/metal/layout/MetalFrameSlots.h | 3 +-- .../rendering/metal/presentation/MetalFrameLoop.mm | 2 ++ .../SplatKitCore/rendering/metal/presentation/MetalSurface.h | 2 +- 3 files changed, 4 insertions(+), 3 deletions(-) diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalFrameSlots.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalFrameSlots.h index c27d023..ac4d90b 100644 --- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/layout/MetalFrameSlots.h +++ b/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 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 f9efd0c..cd5ee5e 100644 --- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalFrameLoop.mm +++ b/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* gpuMillis = &lastGpuMillis_; std::atomic* failed = gpuFailed_.get(); 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 ce9c468..efd14b7 100644 --- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/presentation/MetalSurface.h +++ b/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. From fcc8fa1956e68f08277d1d6498c636a0bcf787ff Mon Sep 17 00:00:00 2001 From: Xget7 Date: Tue, 20 Oct 2026 01:54:25 -0300 Subject: [PATCH 5/7] Remove the camera request test that copied fields by hand --- .../tests/engine/SKTypeConversionsTest.mm | 29 ------------------- 1 file changed, 29 deletions(-) diff --git a/packages/splatkit-ios/tests/engine/SKTypeConversionsTest.mm b/packages/splatkit-ios/tests/engine/SKTypeConversionsTest.mm index 4882ffb..24dac45 100644 --- a/packages/splatkit-ios/tests/engine/SKTypeConversionsTest.mm +++ b/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; From 9420d9c9cb8546610c080d44033ed5145ec86dbd Mon Sep 17 00:00:00 2001 From: Xget7 Date: Tue, 20 Oct 2026 01:54:25 -0300 Subject: [PATCH 6/7] Test that the prepare shader writes exactly kDrawBatches argument records --- .../metal/cull/MetalVisibilityPassTest.mm | 44 +++++++++++++++++++ 1 file changed, 44 insertions(+) diff --git a/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm b/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm index 13e949c..1cc8365 100644 --- a/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm +++ b/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 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 initial((kDrawBatches + kGuardRecords) * kWords, kSentinel); + id draws = [gpu.device newBufferWithBytes:initial.data() + length:initial.size() * sizeof(uint32_t) + options:MTLResourceStorageModeShared]; + const uint32_t splats = 64; + id count = [gpu.device newBufferWithBytes:&splats + length:sizeof(splats) + options:MTLResourceStorageModeShared]; + id cmd = [gpu.queue commandBuffer]; + id 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(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(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 splats; id uniforms; From 39f0632d4ac226c3de098cf79dacb9a5bd9d4753 Mon Sep 17 00:00:00 2001 From: Xget7 Date: Tue, 20 Oct 2026 02:18:54 -0300 Subject: [PATCH 7/7] Let the GpuReport own the failure latch so the handler test pins both captures --- .../rendering/metal/frame/MetalFrameCompute.h | 13 +++++++++---- .../rendering/metal/frame/MetalFrameCompute.mm | 12 +++++++----- .../rendering/metal/frame/MetalFrameComputeTest.mm | 5 +++-- 3 files changed, 19 insertions(+), 11 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 27fdab4..baa8664 100644 --- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h +++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.h @@ -7,6 +7,7 @@ #include #include #include +#include #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> failureLatch) + : failed(std::move(failureLatch)) {} + const std::shared_ptr> failed; std::atomic selectedCount{0}; std::atomic lodLimitedCount{0}; std::atomic lodEvaluatedCount{0}; @@ -119,7 +124,7 @@ class MetalFrameCompute { float lodPixels_ = kDefaultLodPixels; uint32_t lodAdaptedReadback_ = 0; uint64_t frames_ = 0; // submitted so far, which also paces the LOD log - std::shared_ptr report_ = std::make_shared(); + std::shared_ptr report_; }; } // 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 fcf53a4..c49a2dd 100644 --- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm +++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/frame/MetalFrameCompute.mm @@ -34,7 +34,11 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint MetalFrameCompute::MetalFrameCompute(id device, id queue, id library, MetalFrameLoop& frameLoop) - : device_(device), queue_(queue), library_(library), frameLoop_(frameLoop) {} + : device_(device), + queue_(queue), + library_(library), + frameLoop_(frameLoop), + report_(std::make_shared(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 report = report_; - const std::shared_ptr> failed = frameLoop_.failureLatch(); const bool logLod = frames_ % kLodLogFrames == 0; const float pixels = lodPixels_; [selection addCompletedHandler:^(id 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 report = report_; - const std::shared_ptr> failed = frameLoop_.failureLatch(); id countBuffer = visibility_.countBuffer(slot); [sort addCompletedHandler:^(id done) { if (done.status == MTLCommandBufferStatusError) { - failed->store(true); + report->failed->store(true); LOGE("visibility command failed: %s", done.error.localizedDescription.UTF8String); return; } diff --git a/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm b/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm index 5f7ed5b..f854b5c 100644 --- a/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm +++ b/packages/splatkit-ios/tests/rendering/metal/frame/MetalFrameComputeTest.mm @@ -260,8 +260,9 @@ void run(id selection, id sort) { const std::shared_ptr> failed = loop->failureLatch(); compute.reset(); loop.reset(); - // Only the two handlers and this test still hold the latch. - ASSERT_EQ(failed.use_count(), 3); + // 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];