diff --git a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h index 0e1d7b0..d51d954 100644 --- a/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h +++ b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/MetalSplatRenderer.h @@ -12,11 +12,11 @@ #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" #include "rendering/metal/world/MetalWorld.h" -#include "splatkit/rendering/GpuLayout.h" #include "splatkit/rendering/SplatRenderer.h" namespace splatkit { @@ -102,8 +102,6 @@ class MetalSplatRenderer final : public SplatRenderer { uint32_t takePresentTimes(std::vector* 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..ca20298 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::create() { std::unique_ptr r(new MetalSplatRenderer()); @@ -70,34 +82,30 @@ 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; policy.fallback.lodErrorPixels = 1.0f; policy.fallback.alphaThreshold = kQuadMinAlpha; policy.fallback.subpixelThreshold = kMinPixelRadius; - policy.fallback.sortDepth = compute_->depthBits() == MetalRadixSort::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, // 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/cull/MetalVisibilityPass.h b/packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.h index 189d912..58a6a21 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 #include +#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 { @@ -19,15 +21,14 @@ 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 device, id library, - MetalRadixSort::KeyBits depthBits = MetalRadixSort::KeyBits::Full32); + bool create(id device, id 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 // 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 +61,6 @@ class MetalVisibilityPass { return *static_cast(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 +73,15 @@ class MetalVisibilityPass { std::array, kShDegrees> visibility_{}; std::array, kShDegrees> indexedVisibility_{}; id prepareDraw_ = nil; - MetalRadixSort::KeyBits depthBits_ = MetalRadixSort::KeyBits::Full32; + KeyBits depthBits_ = KeyBits::Full32; uint32_t capacity_ = 0; uint32_t sourceCapacity_ = 0; bool indexedOnly_ = false; id projected_ = nil; - std::array, kSlots> count_{}; - std::array, kSlots> drawArguments_{}; - std::array, kSlots> ranges_{}; - std::array, kSlots> rangeStarts_{}; + std::array, kFramesInFlight> count_{}; + std::array, kFramesInFlight> drawArguments_{}; + std::array, kFramesInFlight> ranges_{}; + std::array, 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..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,14 +8,13 @@ namespace splatkit { -bool MetalVisibilityPass::create(id device, id library, - MetalRadixSort::KeyBits depthBits) { +bool MetalVisibilityPass::create(id device, id library, 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 +26,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 +38,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 +97,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..e08d2ef 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 #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 device, id queue, id 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(frames_ % kSlots); } + uint32_t slot() const { return static_cast(frames_ % kFramesInFlight); } void advance() { ++frames_; } // Writes the camera of `frame` into the slot's uniforms, for a frame drawn at `extent`. @@ -75,43 +74,50 @@ class MetalFrameCompute { id order() const { return sort_.values(); } id 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: - bool createOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, - MetalRadixSort::KeyBits bits) const; + // 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 selectedCount{0}; + std::atomic lodLimitedCount{0}; + std::atomic lodEvaluatedCount{0}; + std::atomic lodReadbacks{0}; // completed selections, so each adapts once + std::atomic drawCount{0}; + std::atomic selectMillis{0}; + std::atomic sortMillis{0}; + }; + + bool createOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, KeyBits bits) const; void adaptLodThreshold(); id device_ = nil; id queue_ = nil; id library_ = nil; MetalFrameLoop& frameLoop_; - std::array, kSlots> uniforms_{}; + std::array, kFramesInFlight> uniforms_{}; MetalVisibilityPass visibility_; MetalRadixSort sort_; bool gpuSort_ = false; std::unique_ptr lod_; - std::array, kSlots> lodReadback_{}; + std::array, 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 // limit so detail thins evenly instead of stopping wherever traversal ran out of room. float lodPixels_ = 1.0f; - std::atomic 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 lastSelectedCount_{0}; - std::atomic lastLodLimitedCount_{0}, lastLodEvaluatedCount_{0}; - std::atomic lastSelectMillis_{0}; - std::atomic lastSortMillis_{0}; - std::atomic lastDrawCount_{0}; + std::shared_ptr report_ = std::make_shared(); }; } // 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 83e2fde..fcf53a4 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(depthBits()), static_cast(depthBits()) / 8); if (!gpuSort_) LOGW("GPU sort unavailable, sorting on the CPU"); @@ -54,9 +54,7 @@ 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_; @@ -77,7 +75,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 +88,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; } @@ -106,7 +104,7 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint if (!createOrdering(visibility, sort, depthBits()) || !reserveOrdering(visibility, sort, static_cast(tree.nodeCount()), lod->budget())) return false; - std::array, kSlots> readback{}; + std::array, kFramesInFlight> readback{}; for (auto& buffer : readback) { buffer = metal::buffer(device_, sizeof(LodCounters)); if (!buffer) return false; @@ -116,11 +114,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 +146,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 +170,26 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint destinationOffset:0 size:sizeof(LodCounters)]; [readback endEncoding]; - std::atomic* selected = &lastSelectedCount_; - std::atomic* milliseconds = &lastSelectMillis_; - std::atomic* limited = &lastLodLimitedCount_; - std::atomic* evaluated = &lastLodEvaluatedCount_; + const std::shared_ptr report = report_; + const std::shared_ptr> failed = frameLoop_.failureLatch(); const bool logLod = frames_ % kLodLogFrames == 0; - std::atomic* readbacks = &lodReadbacks_; const float pixels = lodPixels_; - MetalFrameLoop* loop = &frameLoop_; [selection addCompletedHandler:^(id 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(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 +208,17 @@ bool reserveOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, uint } sort_.encode(sort, visibility_.countBuffer(slot), visibility_.depthBits()); visibility_.encodeIndirect(sort, slot); - std::atomic* sortMillis = &lastSortMillis_; - std::atomic* drawn = &lastDrawCount_; + const std::shared_ptr report = report_; + const std::shared_ptr> failed = frameLoop_.failureLatch(); id countBuffer = visibility_.countBuffer(slot); - MetalFrameLoop* loop = &frameLoop_; [sort addCompletedHandler:^(id 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(countBuffer.contents)); + report->sortMillis.store((done.GPUEndTime - done.GPUStartTime) * 1000.0); + report->drawCount.store(*static_cast(countBuffer.contents)); }]; return sort; } 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 + +#include + +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 + +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 + +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/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* 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> 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 presents_ = std::make_shared(); std::atomic lastGpuMillis_{0}; - std::atomic gpuFailed_{false}; + std::shared_ptr> gpuFailed_ = std::make_shared>(false); std::atomic 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* gpuMillis = &lastGpuMillis_; - std::atomic* failed = &gpuFailed_; + std::atomic* failed = gpuFailed_.get(); auto* worldFrame = &completedWorldFrame_; [cmd addCompletedHandler:^(id done) { if (done.status == MTLCommandBufferStatusError) { 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 device, id 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(); 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 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 #include +#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 device, id library); // Transactional allocation: failure preserves the previous buffers and capacity. bool reserve(uint32_t capacity); diff --git a/packages/splatkit-ios/tests/CMakeLists.txt b/packages/splatkit-ios/tests/CMakeLists.txt index a8ec943..732e5c8 100644 --- a/packages/splatkit-ios/tests/CMakeLists.txt +++ b/packages/splatkit-ios/tests/CMakeLists.txt @@ -1,4 +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/diagnostics/SKPngWriterTest.mm b/packages/splatkit-ios/tests/diagnostics/SKPngWriterTest.mm new file mode 100644 index 0000000..a456b6a --- /dev/null +++ b/packages/splatkit-ios/tests/diagnostics/SKPngWriterTest.mm @@ -0,0 +1,101 @@ +#include + +#import +#import + +#include + +#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 decode(NSString* path, size_t* width, size_t* height) { + CGImageSourceRef source = + CGImageSourceCreateWithURL((__bridge CFURLRef)[NSURL fileURLWithPath:path], nullptr); + if (source == nullptr) return {}; + std::vector 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(*width), static_cast(*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. + // clang-format off + const std::vector 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, + }; + // clang-format on + 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 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 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 shortBuffer(2 * 2 * kBytesPerPixel - 1, 0); + EXPECT_FALSE(writePng(path_, shortBuffer, 2, 2)); + const std::vector 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 bgra(kBytesPerPixel, 0); + NSString* missing = [path_ stringByAppendingPathComponent:@"no-such-directory/out.png"]; + EXPECT_FALSE(writePng(missing, bgra, 1, 1)); +} + +} // namespace +} // namespace splatkit 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 + +#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 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(i % 64) - 31.5f) * 0.02f; + const float y = (static_cast(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 diff --git a/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm b/packages/splatkit-ios/tests/rendering/metal/cull/MetalVisibilityPassTest.mm index d138b71..13e949c 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(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,8 +121,8 @@ explicit VisibilityInput(uint32_t count) { id 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, - nil, 0, nullptr, 0)); + 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(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(visibility.order().contents); const auto* projected = static_cast(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 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( 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 selection, id sort) { } test::Gpu* gpu_ = nullptr; - MetalFrameLoop loop_{MetalFrameCompute::kSlots}; + MetalFrameLoop loop_{kFramesInFlight}; std::unique_ptr 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 selection, id 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 selection, id sort) { const auto* draws = static_cast(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 selection, id 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 +#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 cmd, const MetalWorld* world, uint32_t drawCou } test::Gpu* gpu_ = nullptr; - MetalFrameLoop loop_{MetalFrameCompute::kSlots}; + MetalFrameLoop loop_{kFramesInFlight}; std::unique_ptr compute_; MetalSplatPipeline pipeline_; std::unique_ptr 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 sortOnGpu(MetalRadixSort& v, const std::vector& pairs, - MetalRadixSort::KeyBits bits = MetalRadixSort::KeyBits::Full32) { + KeyBits bits = KeyBits::Full32) { Gpu& gpu = Gpu::get(); auto* keys = static_cast(v.keys().contents); auto* values = static_cast(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); }