Tidy the Metal seams, name the renderer's literals and test the iOS mirrors

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<int64_t>* times) override;
   const std::string& deviceDescription() const override { return context_->description(); }
 
-  static constexpr uint32_t kFramesInFlight = MetalFrameCompute::kSlots;
-
  private:
   MetalSplatRenderer() = default;
 

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

@@ -7,6 +7,18 @@
 #include "splatkit/Log.h"
 
 namespace splatkit {
+namespace {
+
+// The residency range reported to the adapters.
+constexpr uint32_t kMinResidencyCapacity = 100'000;
+constexpr uint32_t kMaxResidencyCapacity = 32'000'000;
+// The LOD error threshold range, in pixels, a policy may request.
+constexpr float kMinLodErrorPixels = 0.1f;
+constexpr float kMaxLodErrorPixels = 16.0f;
+// Apple family 7, the minimum this renderer accepts, supports 16384x16384 textures.
+constexpr uint32_t kAppleFamily7MaxTextureDimension = 16384;
+
+}  // namespace
 
 std::unique_ptr<MetalSplatRenderer> MetalSplatRenderer::create() {
   std::unique_ptr<MetalSplatRenderer> r(new MetalSplatRenderer());
@@ -70,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;

packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.h

@@ -5,8 +5,10 @@
 #include <array>
 #include <cstdint>
 
+#include "rendering/metal/layout/MetalDrawArguments.h"
+#include "rendering/metal/layout/MetalFrameSlots.h"
 #include "rendering/metal/layout/MetalShaderTypes.h"
-#include "rendering/metal/sort/MetalRadixSort.h"
+#include "rendering/metal/layout/MetalSortKeys.h"
 #include "splatkit/rendering/SplatRenderer.h"
 
 namespace splatkit {
@@ -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<MTLDevice> device, id<MTLLibrary> library,
-              MetalRadixSort::KeyBits depthBits = MetalRadixSort::KeyBits::Full32);
+  bool create(id<MTLDevice> device, id<MTLLibrary> library, KeyBits depthBits = KeyBits::Full32);
   // Rebuilds the visibility pipelines for a new key width, keeping every reserved buffer
   // and the radix scratch. The key width and the shader's quantized flag must agree, so
   // both move together. Render thread, while idle. On failure the previous key width is
   // 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<const uint32_t*>(count_[slot].contents);
   }
 
-  static constexpr uint32_t kSlots = 2;
-  static constexpr uint32_t kDrawBatches = 7;
-  static constexpr uint32_t kDrawArgumentBytes = sizeof(MTLDrawPrimitivesIndirectArguments);
-
  private:
   static constexpr uint32_t kThreads = 256;
   static constexpr uint32_t kMaxRanges = 65536;
@@ -76,15 +73,15 @@ class MetalVisibilityPass {
   std::array<id<MTLComputePipelineState>, kShDegrees> visibility_{};
   std::array<id<MTLComputePipelineState>, kShDegrees> indexedVisibility_{};
   id<MTLComputePipelineState> prepareDraw_ = nil;
-  MetalRadixSort::KeyBits depthBits_ = MetalRadixSort::KeyBits::Full32;
+  KeyBits depthBits_ = KeyBits::Full32;
   uint32_t capacity_ = 0;
   uint32_t sourceCapacity_ = 0;
   bool indexedOnly_ = false;
   id<MTLBuffer> projected_ = nil;
-  std::array<id<MTLBuffer>, kSlots> count_{};
-  std::array<id<MTLBuffer>, kSlots> drawArguments_{};
-  std::array<id<MTLBuffer>, kSlots> ranges_{};
-  std::array<id<MTLBuffer>, kSlots> rangeStarts_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> count_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> drawArguments_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> ranges_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> rangeStarts_{};
 };
 
 }  // namespace splatkit

packages/splatkit-ios/Sources/SplatKitCore/rendering/metal/cull/MetalVisibilityPass.mm

@@ -8,14 +8,13 @@
 
 namespace splatkit {
 
-bool MetalVisibilityPass::create(id<MTLDevice> device, id<MTLLibrary> library,
-                                 MetalRadixSort::KeyBits depthBits) {
+bool MetalVisibilityPass::create(id<MTLDevice> device, id<MTLLibrary> 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);

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

@@ -9,6 +9,7 @@
 #include <string>
 
 #include "rendering/metal/cull/MetalVisibilityPass.h"
+#include "rendering/metal/layout/MetalFrameSlots.h"
 #include "rendering/metal/presentation/MetalFrameLoop.h"
 #include "rendering/metal/selection/MetalLodSelection.h"
 #include "rendering/metal/sort/MetalRadixSort.h"
@@ -25,8 +26,6 @@ namespace splatkit {
 // frame loop must outlive it. Encode only; draw() commits what is returned.
 class MetalFrameCompute {
  public:
-  static constexpr uint32_t kSlots = MetalVisibilityPass::kSlots;
-
   MetalFrameCompute(id<MTLDevice> device, id<MTLCommandQueue> queue, id<MTLLibrary> library,
                     MetalFrameLoop& frameLoop);
 
@@ -38,7 +37,7 @@ class MetalFrameCompute {
   // gpuSort() false.
   bool create();
   bool gpuSort() const { return gpuSort_; }
-  MetalRadixSort::KeyBits depthBits() const { return visibility_.depthBits(); }
+  KeyBits depthBits() const { return visibility_.depthBits(); }
 
   // The sort key width and the LOD threshold and splat limit of `policy`. Waits for the frames
   // in flight when something changes. False, with `reason`, when the GPU cannot order or the
@@ -56,7 +55,7 @@ class MetalFrameCompute {
   // The slot the frame being encoded uses: the per-slot uniforms, counts and readbacks of a
   // frame the GPU may still be reading are never touched. advance() moves to the next slot once
   // the frame was submitted.
-  uint32_t slot() const { return static_cast<uint32_t>(frames_ % kSlots); }
+  uint32_t slot() const { return static_cast<uint32_t>(frames_ % kFramesInFlight); }
   void advance() { ++frames_; }
 
   // Writes the camera of `frame` into the slot's uniforms, for a frame drawn at `extent`.
@@ -75,43 +74,50 @@ class MetalFrameCompute {
   id<MTLBuffer> order() const { return sort_.values(); }
   id<MTLBuffer> drawArguments(uint32_t slot) const { return visibility_.drawArguments(slot); }
 
-  double lastSortMillis() const { return lastSortMillis_.load(); }
-  uint32_t lastDrawCount() const { return lastDrawCount_.load(); }
-  uint32_t lastSelectedCount() const { return lastSelectedCount_.load(); }
-  double lastSelectMillis() const { return lastSelectMillis_.load(); }
-  uint32_t lastLodLimitedCount() const { return lod_ ? lastLodLimitedCount_.load() : 0; }
-  uint32_t lastLodEvaluatedCount() const { return lod_ ? lastLodEvaluatedCount_.load() : 0; }
+  double lastSortMillis() const { return report_->sortMillis.load(); }
+  uint32_t lastDrawCount() const { return report_->drawCount.load(); }
+  uint32_t lastSelectedCount() const { return report_->selectedCount.load(); }
+  double lastSelectMillis() const { return report_->selectMillis.load(); }
+  uint32_t lastLodLimitedCount() const { return lod_ ? report_->lodLimitedCount.load() : 0; }
+  uint32_t lastLodEvaluatedCount() const { return lod_ ? report_->lodEvaluatedCount.load() : 0; }
   // The threshold the next selection runs at, at least the policy's.
   float lodPixels() const { return lodPixels_; }
 
  private:
-  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<uint32_t> selectedCount{0};
+    std::atomic<uint32_t> lodLimitedCount{0};
+    std::atomic<uint32_t> lodEvaluatedCount{0};
+    std::atomic<uint32_t> lodReadbacks{0};  // completed selections, so each adapts once
+    std::atomic<uint32_t> drawCount{0};
+    std::atomic<double> selectMillis{0};
+    std::atomic<double> sortMillis{0};
+  };
+
+  bool createOrdering(MetalVisibilityPass& visibility, MetalRadixSort& sort, KeyBits bits) const;
   void adaptLodThreshold();
 
   id<MTLDevice> device_ = nil;
   id<MTLCommandQueue> queue_ = nil;
   id<MTLLibrary> library_ = nil;
   MetalFrameLoop& frameLoop_;
-  std::array<id<MTLBuffer>, kSlots> uniforms_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> uniforms_{};
   MetalVisibilityPass visibility_;
   MetalRadixSort sort_;
   bool gpuSort_ = false;
   std::unique_ptr<MetalLodSelection> lod_;
-  std::array<id<MTLBuffer>, kSlots> lodReadback_{};
+  std::array<id<MTLBuffer>, kFramesInFlight> lodReadback_{};
   float lodErrorPixels_ = 1.0f;
   uint32_t lodSplatLimit_ = 0;
   // The threshold selection runs at: the policy's, raised while a cut would exceed the splat
   // limit so detail thins evenly instead of stopping wherever traversal ran out of room.
   float lodPixels_ = 1.0f;
-  std::atomic<uint32_t> lodReadbacks_{0};  // completed selections, so each adapts once
   uint32_t lodAdaptedReadback_ = 0;
   uint64_t frames_ = 0;  // submitted so far, which also paces the LOD log
-  std::atomic<uint32_t> lastSelectedCount_{0};
-  std::atomic<uint32_t> lastLodLimitedCount_{0}, lastLodEvaluatedCount_{0};
-  std::atomic<double> lastSelectMillis_{0};
-  std::atomic<double> lastSortMillis_{0};
-  std::atomic<uint32_t> lastDrawCount_{0};
+  std::shared_ptr<GpuReport> report_ = std::make_shared<GpuReport>();
 };
 
 }  // namespace splatkit