diff --git a/CHANGELOG.md b/CHANGELOG.md index 55c917d..76ec26a 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -16,13 +16,19 @@ after its public API and release process are established. ### Added +- Metal Gaussian compute/image tests now consume a backend-private persistent + attribute store. Camera and metadata changes reuse buffers; matching revision + ranges upload only edited attributes, with GPU version copies protecting + in-flight readers. Transactional updates, source/handle identity, SH layout + changes and a live-byte budget have Apple GPU coverage. Renderer scheduling + and telemetry integration remain open; GPU mode support is unchanged. + - Metal Gaussian gather packs GPU-sorted records into the scalar raster stream and writes indirect draw arguments, including empty-frame resets. The Apple GPU harness runs preparation through indirect raster in one submission and compares color/depth/IDs with CPU-prepared direct draws; ABI, artifact identity - and install checks include the new kernel. Renderer integration and persistent - Gaussian attribute residency remain open; `prefer` still falls back and - `require` rejects. + and install checks include the new kernel. Renderer integration remains open; + `prefer` still falls back and `require` rejects. - Shared Slang Gaussian projection, covariance, SH, culling/compaction and deterministic radix-sort kernels now compile into the Metal Gaussian library. diff --git a/backend/merlin-metal/CMakeLists.txt b/backend/merlin-metal/CMakeLists.txt index 6935793..4f7304b 100644 --- a/backend/merlin-metal/CMakeLists.txt +++ b/backend/merlin-metal/CMakeLists.txt @@ -9,6 +9,7 @@ find_library(MERLIN_COREGRAPHICS_FRAMEWORK CoreGraphics REQUIRED) add_library(merlin-metal STATIC src/backend.mm + src/gaussian_residency.mm src/resource_table.cpp ) merlin_target_defaults(merlin-metal) diff --git a/backend/merlin-metal/src/gaussian_residency.hpp b/backend/merlin-metal/src/gaussian_residency.hpp new file mode 100644 index 0000000..bb868e1 --- /dev/null +++ b/backend/merlin-metal/src/gaussian_residency.hpp @@ -0,0 +1,93 @@ +#pragma once + +#import + +#include + +#include +#include +#include + +namespace merlin::metal { + +// Backend-private immutable attribute versions. The byte budget includes old +// versions and staging retained by unfinished command buffers, not just the +// current scene. Calls are externally serialized on one Metal command queue. +class GaussianResidency { + struct Budget { + std::atomic live{}; + std::uint64_t limit{}; + }; + +public: + struct Buffer { + Buffer() = default; + Buffer(const Buffer&) = delete; + Buffer& operator=(const Buffer&) = delete; + id metal = nil; + std::shared_ptr budget; + ~Buffer(); + }; + using BufferPtr = std::shared_ptr; + + struct Resource { + extraction::GaussianRecord record; + BufferPtr positions; + BufferPtr covariances; + BufferPtr opacities; + BufferPtr radiance; + }; + + struct Update { + // Ordered by the complete generation-bearing resource handle for stable + // frame-wide tie breaking. Metadata belongs to this immutable snapshot. + std::vector resources; + std::uint64_t upload_bytes{}; + std::uint64_t upload_range_count{}; + std::uint64_t device_copy_bytes{}; + std::uint64_t allocation_count{}; + + private: + friend class GaussianResidency; + struct Copy { + BufferPtr source; + BufferPtr destination; + std::uint64_t destination_offset{}; + std::uint64_t bytes{}; + }; + std::vector copies; + std::shared_ptr owner; + std::uint64_t epoch{}; + std::uint64_t source_id{}; + bool encoded{}; + bool committed{}; + }; + + GaussianResidency(id device, std::uint64_t byte_budget); + GaussianResidency(const GaussianResidency&) = delete; + GaussianResidency& operator=(const GaussianResidency&) = delete; + // Preparation is transactional: allocation/validation failure, or dropping + // an unsubmitted update, leaves the resident scene untouched. + std::shared_ptr Prepare(const extraction::FrameSnapshot& snapshot); + // Encode before preparation kernels; automatically retain the update and + // all its buffers until completion, even if the store/frame is destroyed. + void Encode(const std::shared_ptr& update, id command); + // Call only after committing that command to the same serial queue. An + // upload failure invalidates dependent submissions: the owner must Reset + // and report the GPU failure before using residency again. + void Commit(const std::shared_ptr& update); + void Reset(); + [[nodiscard]] std::uint64_t live_bytes() const noexcept; + +private: + BufferPtr Allocate(std::uint64_t bytes, MTLResourceOptions options, Update& update); + void ValidateUpdate(const std::shared_ptr& update) const; + + id device_; + std::shared_ptr budget_; + std::vector resident_; + std::uint64_t source_id_{}; + std::uint64_t epoch_{}; +}; + +} // namespace merlin::metal diff --git a/backend/merlin-metal/src/gaussian_residency.mm b/backend/merlin-metal/src/gaussian_residency.mm new file mode 100644 index 0000000..725e40a --- /dev/null +++ b/backend/merlin-metal/src/gaussian_residency.mm @@ -0,0 +1,192 @@ +#include "gaussian_residency.hpp" + +#include + +#include +#include +#include +#include +#include + +namespace merlin::metal { +namespace { +[[noreturn]] void Fail(render::RendererErrorCode code, const char* message) { + throw render::RendererError(code, "synchronize Metal Gaussian attributes", message); +} + +void Validate(const extraction::GaussianRecord& record) { + if (!record.positions || !record.covariances || !record.opacities || + !record.spherical_harmonics_coefficients || record.spherical_harmonics_degree > 3 || + record.positions->size() > std::numeric_limits::max() || + record.positions->size() != record.covariances->size() || + record.positions->size() != record.opacities->size()) { + Fail(render::RendererErrorCode::InvalidRequest, "Gaussian attribute payload is malformed"); + } + const auto coefficients = (record.spherical_harmonics_degree + 1U) * + (record.spherical_harmonics_degree + 1U); + if (record.spherical_harmonics_coefficients->size() != record.positions->size() * coefficients) { + Fail(render::RendererErrorCode::InvalidRequest, "Gaussian SH payload size is inconsistent"); + } + for (const auto& range : record.particle_ranges) { + if (range.first > record.positions->size() || + range.count > record.positions->size() - range.first) { + Fail(render::RendererErrorCode::InvalidRequest, "Gaussian changed range is out of bounds"); + } + } +} +} // namespace + +GaussianResidency::Buffer::~Buffer() { + if (metal != nil) budget->live.fetch_sub(metal.length, std::memory_order_relaxed); +} + +GaussianResidency::GaussianResidency(id device, std::uint64_t byte_budget) + : device_(device), budget_(std::make_shared()) { + if (!device) Fail(render::RendererErrorCode::InvalidRequest, "Metal device is null"); + budget_->limit = byte_budget; +} + +std::uint64_t GaussianResidency::live_bytes() const noexcept { + return budget_->live.load(std::memory_order_relaxed); +} + +GaussianResidency::BufferPtr GaussianResidency::Allocate( + std::uint64_t bytes, MTLResourceOptions options, Update& update) { + // Even empty attributes have a legal binding; byte-addressed shaders use + // 32-bit offsets. Reject before allocation or narrowing to NSUInteger. + bytes = std::max(bytes, std::uint64_t{16}); + const auto live = live_bytes(); + if (bytes > std::numeric_limits::max() || bytes > device_.maxBufferLength) + Fail(render::RendererErrorCode::Unsupported, "Gaussian attribute exceeds the Metal buffer or shader address limit"); + if (live > budget_->limit || bytes > budget_->limit - live) + Fail(render::RendererErrorCode::ResourceExhausted, "Gaussian residency exceeds the live-byte budget"); + auto result = std::make_shared(); + result->budget = budget_; + result->metal = [device_ newBufferWithLength:bytes options:options]; + if (!result->metal) Fail(render::RendererErrorCode::BackendFailure, "Metal attribute buffer allocation failed"); + budget_->live.fetch_add(result->metal.length, std::memory_order_relaxed); + ++update.allocation_count; + return result; +} + +std::shared_ptr GaussianResidency::Prepare( + const extraction::FrameSnapshot& snapshot) { + auto update = std::make_shared(); + update->owner = budget_; + update->epoch = epoch_; + update->source_id = snapshot.source_id; + std::unordered_map previous; + if (snapshot.source_id == source_id_) { + for (const auto& resource : resident_) previous.emplace(resource.record.gaussian, &resource); + } + update->resources.reserve(snapshot.gaussians.size()); + for (const auto& record : snapshot.gaussians) { + Validate(record); + const auto found = previous.find(record.gaussian); + const auto* old = found == previous.end() ? nullptr : found->second; + const auto coefficients = (record.spherical_harmonics_degree + 1U) * + (record.spherical_harmonics_degree + 1U); + const bool partial = old && snapshot.source_id != 0 && + old->record.revision == record.particle_base_revision && + old->record.positions->size() == record.positions->size() && + old->record.spherical_harmonics_degree == record.spherical_harmonics_degree && + std::any_of(record.particle_ranges.begin(), record.particle_ranges.end(), + [](const auto& range) { return range.count != 0; }); + const auto attribute = [&](const auto& payload, const auto& previous_payload, + std::uint64_t revision, std::uint64_t previous_revision, + BufferPtr previous_buffer, std::uint32_t multiplier) -> BufferPtr { + if (previous_buffer && revision == previous_revision && payload == previous_payload) + return previous_buffer; + using Element = typename std::decay_t::value_type; + const auto bytes = std::uint64_t{payload->size()} * sizeof(Element); + auto destination = Allocate(bytes, MTLResourceStorageModePrivate, *update); + const auto stage = [&](std::uint64_t first, std::uint64_t count) { + const auto size = count * sizeof(Element); + if (!size) return; + auto staging = Allocate(size, MTLResourceStorageModeShared, *update); + std::memcpy(staging->metal.contents, payload->data() + first, size); + update->copies.push_back({staging, destination, first * sizeof(Element), size}); + update->upload_bytes += size; + ++update->upload_range_count; + }; + if (partial && previous_buffer && bytes) { + // Version on the GPU before patching: never overwrite a buffer that an + // earlier command may still read, and never copy unchanged data on CPU. + update->copies.push_back({previous_buffer, destination, 0, bytes}); + update->device_copy_bytes += bytes; + for (const auto& range : record.particle_ranges) + stage(std::uint64_t{range.first} * multiplier, std::uint64_t{range.count} * multiplier); + } else { + stage(0, payload->size()); + } + return destination; + }; + Resource next; + next.record = record; + next.positions = attribute(record.positions, old ? old->record.positions : nullptr, + record.positions_revision, old ? old->record.positions_revision : 0, + old ? old->positions : nullptr, 1); + next.covariances = attribute(record.covariances, old ? old->record.covariances : nullptr, + record.covariance_revision, old ? old->record.covariance_revision : 0, + old ? old->covariances : nullptr, 1); + next.opacities = attribute(record.opacities, old ? old->record.opacities : nullptr, + record.opacity_revision, old ? old->record.opacity_revision : 0, + old ? old->opacities : nullptr, 1); + next.radiance = attribute(record.spherical_harmonics_coefficients, + old ? old->record.spherical_harmonics_coefficients : nullptr, + record.radiance_revision, old ? old->record.radiance_revision : 0, + old ? old->radiance : nullptr, coefficients); + update->resources.push_back(std::move(next)); + } + std::sort(update->resources.begin(), update->resources.end(), + [](const auto& a, const auto& b) { return a.record.gaussian < b.record.gaussian; }); + for (std::size_t i = 1; i < update->resources.size(); ++i) { + if (update->resources[i - 1].record.gaussian == update->resources[i].record.gaussian) + Fail(render::RendererErrorCode::InvalidRequest, "Duplicate Gaussian resource handle"); + } + return update; +} + +void GaussianResidency::ValidateUpdate(const std::shared_ptr& update) const { + if (!update || update->owner != budget_ || update->epoch != epoch_ || update->committed) + Fail(render::RendererErrorCode::InvalidRequest, "Stale or foreign Gaussian residency update"); +} + +void GaussianResidency::Encode(const std::shared_ptr& update, id command) { + ValidateUpdate(update); + if (!command || command.device != device_ || update->encoded || + command.status != MTLCommandBufferStatusNotEnqueued || !command.retainedReferences) + Fail(render::RendererErrorCode::InvalidRequest, "Expected an unsubmitted retaining Metal command buffer"); + if (!update->copies.empty()) { + auto encoder = [command blitCommandEncoder]; + if (!encoder) Fail(render::RendererErrorCode::BackendFailure, "Metal attribute blit encoder allocation failed"); + for (const auto& copy : update->copies) { + [encoder copyFromBuffer:copy.source->metal sourceOffset:0 + toBuffer:copy.destination->metal destinationOffset:copy.destination_offset size:copy.bytes]; + } + [encoder endEncoding]; + } + // Capturing the plan keeps allocation accounting and old versions alive as + // long as the GPU uses them; the plan itself does not retain the command. + __block auto retained = update; + [command addCompletedHandler:^(id) { retained.reset(); }]; + update->encoded = true; +} + +void GaussianResidency::Commit(const std::shared_ptr& update) { + ValidateUpdate(update); + if (!update->encoded) Fail(render::RendererErrorCode::InvalidRequest, "Gaussian update has not been encoded"); + auto resources = update->resources; + resident_.swap(resources); + source_id_ = update->source_id; + update->committed = true; + ++epoch_; +} + +void GaussianResidency::Reset() { + resident_.clear(); + source_id_ = 0; + ++epoch_; +} + +} // namespace merlin::metal diff --git a/docs/design/metal-gaussian-execution.md b/docs/design/metal-gaussian-execution.md index c8bc853..de2d6ad 100644 --- a/docs/design/metal-gaussian-execution.md +++ b/docs/design/metal-gaussian-execution.md @@ -88,8 +88,8 @@ projection/sort policies, transforms, multiple resources, hidden/all-rejected input, workgroup boundaries, asymmetric alpha composition and a prefilled opaque depth attachment. They do not establish general scene or host-presentation parity. This remains a correctness harness, not renderer GPU execution: -persistent attributes, completion-safe frame scheduling and telemetry -integration remain Phase 2 work. Controlled performance captures +renderer integration of persistent attributes, completion-safe frame scheduling +and telemetry remain Phase 2 work. Controlled performance captures and updated native viewport/HgiMetal comparisons remain unfinished; this is not completion of the phase gate below. @@ -130,6 +130,30 @@ update metadata. Replaced buffers and slots remain alive until their last GPU consumer completes. Bound resident and scratch allocation and report failures through the existing diagnostic/fallback contract. +The private Metal attribute store now supplies the compute/image harness with +immutable device-local position, covariance, opacity and SH buffers. It keys +reuse by source identity, the complete resource handle, attribute revisions and +shared payload identity. Camera, transform and visibility changes reuse those +buffers. Matching particle-base revisions and unchanged layouts permit staged +range updates; a GPU copy creates the new attribute version before patching it, +so earlier submissions can continue reading the old version. Revision gaps, +source changes, count changes and SH layout changes upload complete affected +attributes. Unchanged attributes retain their original buffers. + +Preparation is transactional, and a separate commit publishes the resident +scene only after submission. Command completion retains staging and old +versions, including their contribution to an explicit live-byte budget. +Allocation failure leaves the previous scene usable. The caller must invalidate +residency and report a failed upload command before scheduling dependent work. +The harness checks actual bytes, partial SH ranges, removal/generation reuse, +abandoned/invalid updates, budget exhaustion and multiple blocked submissions. +Continuous image comparisons check static, camera, transform, localized edits, +visibility, removal and reintroduction with the existing color/depth/ID tolerance. +This is a backend-private building block: renderer scheduling, scratch reuse, +error recovery and public telemetry integration are still unfinished. It scans +resource metadata, and partial updates currently copy the full changed attribute +on the GPU; it does not yet implement a resource-delta fast path or an arena. + Connect the complete frame path on the GPU: ```text diff --git a/docs/reference/support-matrix.md b/docs/reference/support-matrix.md index fa9b6b6..b8788fc 100644 --- a/docs/reference/support-matrix.md +++ b/docs/reference/support-matrix.md @@ -74,7 +74,7 @@ between separately produced OpenUSD SDKs remain the operator's responsibility. | MaterialXGenSlang material-function prototype | Available: optional `Merlin::MaterialX` emits deterministic graph-only Slang functions and renderer-owned minimum Standard Surface results for constants, image/UV0/world-normal, add/multiply/mix, `base`, `base_color`, `metalness`, `specular_roughness`, and `normal`. Portable library/include fingerprints feed topology-only module keys separated from parameter/resource state. Registered parameter-only and texture/sampler artifacts execute in renderer-owned Vulkan Forward after ABI/reflection checks, reuse pipelines across value and texture-content edits, and report structured fallback/capability telemetry. The same sources retain installed SPIR-V, Metal-target, and reflection evidence. General MaterialX documents, tangent-space normal mapping, production IBL, and Hydra MaterialX ingestion are not claimed | | Material ABI agreement | `merlin.material-abi/v1` is available in Core. A consumer declares the result fields it reads and the geometry inputs it can build, and a module is checked against that rather than in isolation. A compiled artifact's reflected interface is checked back against the module's logical one by name, type, and array size, so SPIR-V and Metal describe one material through their own native bindings and agree with each other by agreeing with the module. The same contract owns the pass-neutrality rule, which `Merlin::MaterialX` applies to its own output; both generated modules and all four of their SPIR-V/Metal artifacts are checked in the test suite, including the dropped, retyped, and undeclared-parameter cases | | Native Metal backend and residency | Available for offscreen Mesh Forward: native device/queue, runtime MSL, buffers/textures/samplers, heap residency, generation-checked argument-buffer tables with conventional fallback, frames-in-flight retirement, basic material/opacity mask, color/depth/primId/instanceId AOVs, CPU readback, capacity diagnostics, and Metal-specific telemetry | -| Metal Gaussian rendering | Available: shared CPU projection/covariance/SH evaluation and sorting, shared Slang raster math compiled into an embedded Metal library, alpha composition with opaque Mesh depth, resource/particle IDs, and immutable prepared-stream reuse. Image tests cover edits, camera/resize changes and in-flight lifetime; the prior embedded-MSL path was checked in the OpenUSD 26.08 native viewport and usdview/HgiMetal GPU-copy display with a 5.8-million-particle stage. The Slang path was also manually checked in usdview/HgiMetal with that stage; the native viewport comparison remains open. GPU preparation/sort/gather and indirect raster have Metal ABI and CPU-reference offscreen image tests without intermediate readback, but are not connected to renderer execution; persistent residency, frame scheduling and tiling remain open. `prefer` falls back and `require` rejects. | +| Metal Gaussian rendering | Available: shared CPU projection/covariance/SH evaluation and sorting, shared Slang raster math compiled into an embedded Metal library, alpha composition with opaque Mesh depth, resource/particle IDs, and immutable prepared-stream reuse. Image tests cover edits, camera/resize changes and in-flight lifetime; the prior embedded-MSL path was checked in the OpenUSD 26.08 native viewport and usdview/HgiMetal GPU-copy display with a 5.8-million-particle stage. The Slang path was also manually checked in usdview/HgiMetal with that stage; the native viewport comparison remains open. GPU preparation/sort/gather and indirect raster have Metal ABI and CPU-reference offscreen image tests without intermediate readback, but are not connected to renderer execution; immutable attribute residency and partial updates are validated in that harness. Renderer integration of residency, frame scheduling and tiling remain open. `prefer` falls back and `require` rejects. | | Native Metal viewport presentation | Available: adapter-owned `CAMetalLayer`, renderer-owned drawable encoding, GPU-only offscreen-to-drawable presentation, resize/frames-in-flight safety, sRGB/Display P3 SDR policy with an explicit future HDR boundary, vsync/drawable-count pacing, Dear ImGui integration, presentation telemetry, and exact offscreen reference parity | | HgiVulkan host presentation bridge | Available: public-driver discovery, Hgi-owned color targets, Tier 0 upload fallback, and selected color GPU copy on validated OpenUSD 26.05/26.08 packages when their imported `hgiVulkan` target is present. Merlin borrows the Hgi Vulkan 1.3 device/graphics queue for its conventional renderer path, exports one color image, copies it with explicit barriers, and releases its lease from Hgi command-buffer completion; depth and id AOVs stay on CPU readback. Runtime comparison evidence covers image parity and no color Map/upload or coarse wait through resize. Merlin-owned Vulkan remains 1.4. Direct sharing reports `public-texture-import-unavailable` because public Hgi cannot import a Merlin-owned `VkImage`, so GPU copy remains selected. | | HgiMetal host presentation bridge | Available on validated OpenUSD 26.05/26.08 packages: Hgi-owned color targets receive same-device Metal GPU copies from leased renderer AOVs with completion-safe resize retirement and Tier 0 fallback. Direct sharing remains rejected because the public host texture-import contract is unavailable. | diff --git a/docs/roadmap/current.md b/docs/roadmap/current.md index a0869f4..7f0f34d 100644 --- a/docs/roadmap/current.md +++ b/docs/roadmap/current.md @@ -23,8 +23,11 @@ ownership, dependency, and fallback contracts. GPU preparation/SH/radix-sort kernels are also shared and checked against the CPU reference on Apple GPU. The offscreen harness now connects GPU gather and indirect raster with color/depth/ID comparisons and no intermediate CPU - readback. Controlled static/motion/edit captures, native viewport/HgiMetal - rechecks, persistent attribute residency and renderer scheduling through + readback. A private attribute store now supplies that harness with immutable + resident buffers, range uploads and completion-safe versions. Continuous + camera/transform/edit images and in-flight byte-budget checks cover the store. + Controlled static/motion/edit captures, native viewport/HgiMetal rechecks, + and renderer integration of residency, frame scheduling and telemetry through raster remain open. - 🚧 Recheck the Hydra CPU-readback floor in usdview after preferring diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 92a412d..8acf950 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -235,6 +235,17 @@ if(TARGET Merlin::Metal) LABELS "metal;gaussian;gpu;compute;cpu-reference;image;aov;indirect" SKIP_RETURN_CODE 77 TIMEOUT 60) + add_executable(merlin-metal-gaussian-residency-tests metal_gaussian_residency_test.mm) + merlin_target_defaults(merlin-metal-gaussian-residency-tests) + target_link_libraries(merlin-metal-gaussian-residency-tests PRIVATE Merlin::Metal) + target_compile_options(merlin-metal-gaussian-residency-tests PRIVATE -fobjc-arc) + set_target_properties(merlin-metal-gaussian-residency-tests PROPERTIES + OBJCXX_STANDARD 20 OBJCXX_STANDARD_REQUIRED ON OBJCXX_EXTENSIONS OFF) + add_test(NAME merlin-metal-gaussian-residency COMMAND merlin-metal-gaussian-residency-tests) + set_tests_properties(merlin-metal-gaussian-residency PROPERTIES + LABELS "metal;gaussian;gpu;residency;lifetime" + SKIP_RETURN_CODE 77 TIMEOUT 60) + add_executable(merlin-metal-resource-table-tests metal_resource_table_test.cpp ) @@ -267,7 +278,7 @@ if(TARGET Merlin::Metal) SKIP_RETURN_CODE 77) # Enable Apple's validation before device creation, so ABI and argument-buffer # mistakes fail the runtime suite instead of depending on a local shell setup. - set_tests_properties(merlin-metal-gaussian merlin-metal-gaussian-compute merlin-metal-backend + set_tests_properties(merlin-metal-gaussian merlin-metal-gaussian-compute merlin-metal-gaussian-residency merlin-metal-backend PROPERTIES ENVIRONMENT "MTL_DEBUG_LAYER=1;MTL_SHADER_VALIDATION=1") endif() diff --git a/tests/metal_gaussian_compute_test.mm b/tests/metal_gaussian_compute_test.mm index 1e85c29..b3cc0b5 100644 --- a/tests/metal_gaussian_compute_test.mm +++ b/tests/metal_gaussian_compute_test.mm @@ -1,5 +1,6 @@ #include "../backend/merlin-metal/src/gaussian_compute_abi.hpp" #include "../backend/merlin-metal/src/gaussian_raster_abi.hpp" +#include "../backend/merlin-metal/src/gaussian_residency.hpp" #include #include @@ -42,6 +43,16 @@ void Near(float actual, float expected, const char* field) { return (count + size - 1) / size; } +merlin::Mat4 Multiply(const merlin::Mat4& a, const merlin::Mat4& b) { + merlin::Mat4 result; + result.values.fill(0); + for (std::size_t column = 0; column < 4; ++column) + for (std::size_t row = 0; row < 4; ++row) + for (std::size_t k = 0; k < 4; ++k) + result.values[column * 4 + row] += a.values[k * 4 + row] * b.values[column * 4 + k]; + return result; +} + id Buffer(id device, std::size_t bytes, const void* data = nullptr) { auto buffer = [device newBufferWithLength:std::max(bytes, std::size_t{16}) options:MTLResourceStorageModeShared]; @@ -169,7 +180,8 @@ void Near(float actual, float expected, const char* field) { // submission. No intermediate readback schedules subsequent GPU stages. std::array, 4> Compare(id device, id queue, const Kernels& kernels, FrameSnapshot snapshot, bool dynamic_count = false, bool compare_image = false, - float opaque_depth = 1.0F) { + float opaque_depth = 1.0F, merlin::metal::GaussianResidency* persistent = nullptr, + std::uint64_t* attribute_upload_bytes = nullptr) { std::vector ordered(snapshot.gaussians.begin(), snapshot.gaussians.end()); std::sort(ordered.begin(), ordered.end(), [](const auto& a, const auto& b) { return a.gaussian < b.gaussian; }); @@ -187,6 +199,11 @@ void Near(float actual, float expected, const char* field) { auto control = Buffer(device, (histogram_offset + padded * 2U + 16U) * sizeof(std::uint32_t)); auto command = [queue commandBuffer]; Require(command != nil, "Metal command allocation failed"); + merlin::metal::GaussianResidency local_residency(device, 64 * 1024 * 1024); + auto& residency = persistent ? *persistent : local_residency; + auto attributes = residency.Prepare(snapshot); + residency.Encode(attributes, command); + if (attribute_upload_bytes) *attribute_upload_bytes = attributes->upload_bytes; std::vector> counters; std::vector> classifications; const auto policy = merlin::extraction::SelectGaussianSortingPolicy(snapshot); @@ -208,7 +225,7 @@ void Near(float actual, float expected, const char* field) { const auto& record = snapshot.gaussians[i]; const auto count = static_cast(record.positions->size()); PrepareConstants constants; - constants.local_to_camera = record.transform; // Fixtures use identity view. + constants.local_to_camera = Multiply(snapshot.view, record.transform); constants.projection = snapshot.projection; constants.viewport_size = {320, 192}; constants.resource_id_low = static_cast(record.gaussian); @@ -224,10 +241,11 @@ void Near(float actual, float expected, const char* field) { if (count && record.visible) { auto encoder = [command computeCommandEncoder]; [encoder setComputePipelineState:kernels.prepare]; - [encoder setBuffer:Upload(device, *record.positions) offset:0 atIndex:0]; - [encoder setBuffer:Upload(device, *record.covariances) offset:0 atIndex:1]; - [encoder setBuffer:Upload(device, *record.opacities) offset:0 atIndex:2]; - [encoder setBuffer:Upload(device, *record.spherical_harmonics_coefficients) offset:0 atIndex:3]; + const auto& resident = attributes->resources[i]; + [encoder setBuffer:resident.positions->metal offset:0 atIndex:0]; + [encoder setBuffer:resident.covariances->metal offset:0 atIndex:1]; + [encoder setBuffer:resident.opacities->metal offset:0 atIndex:2]; + [encoder setBuffer:resident.radiance->metal offset:0 atIndex:3]; [encoder setBuffer:classifications.back() offset:0 atIndex:4]; [encoder setBuffer:prepared offset:base * sizeof(PreparedRecord) atIndex:5]; [encoder setBuffer:counters.back() offset:0 atIndex:6]; @@ -323,6 +341,7 @@ void Near(float actual, float expected, const char* field) { static_cast(cpu_instances.size()), opaque_depth); } [command commit]; + residency.Commit(attributes); [command waitUntilCompleted]; if (command.status != MTLCommandBufferStatusCompleted) throw std::runtime_error(command.error.localizedDescription.UTF8String); @@ -478,6 +497,64 @@ FrameSnapshot RasterFixture(std::uint32_t count) { snapshot.gaussians.assign({record}); return snapshot; } +void CompareResidentFrames(id device, id queue, const Kernels& kernels) { + merlin::metal::GaussianResidency residency(device, 64 * 1024 * 1024); + auto snapshot = RasterFixture(3); + snapshot.source_id = 17; + auto record = snapshot.gaussians[0]; + record.revision = 1; + snapshot.gaussians.assign({record}); + std::uint64_t bytes = 0; + const auto compare = [&] { + Compare(device, queue, kernels, snapshot, false, true, 1.0F, &residency, &bytes); + }; + compare(); + const auto initial_bytes = bytes; + Require(initial_bytes == 3 * (2 * sizeof(merlin::Vec3) + sizeof(merlin::Covariance3) + sizeof(float)), + "Initial resident upload differs from source attributes"); + compare(); + Require(bytes == 0, "Static frame reuploaded source attributes"); + snapshot.view.values[12] = 0.125F; + compare(); + Require(bytes == 0, "Camera motion reuploaded source attributes"); + record.revision = 2; + record.transform.values[13] = -0.125F; + snapshot.gaussians.assign({record}); + compare(); + Require(bytes == 0, "Transform edit reuploaded source attributes"); + + record.revision = record.opacity_revision = record.covariance_revision = 3; + record.particle_base_revision = 2; + record.particle_ranges = {{1, 1}}; + auto opacity = std::make_shared>(*record.opacities); + (*opacity)[1] = 0.25F; + record.opacities = opacity; + auto covariance = std::make_shared>(*record.covariances); + (*covariance)[1].xx *= 2; + record.covariances = covariance; + snapshot.gaussians.assign({record}); + compare(); + const auto partial_bytes = bytes; + Require(partial_bytes == sizeof(float) + sizeof(merlin::Covariance3), + "Localized edit did not upload only changed attribute ranges"); + compare(); + Require(bytes == 0, "Repeated partial-update snapshot uploaded again"); + record.visible = false; + record.revision = 4; + snapshot.gaussians.assign({record}); + compare(); + Require(bytes == 0, "Visibility edit reuploaded source attributes"); + snapshot.gaussians.assign({}); + compare(); + Require(bytes == 0, "Empty frame uploaded source attributes"); + record.visible = true; + record.gaussian += 1ULL << 32; + snapshot.gaussians.assign({record}); + compare(); + Require(bytes == initial_bytes, "New resource generation reused removed attributes"); + std::cout << "Resident image sequence: initial=" << initial_bytes + << " static/camera/transform/visibility=0 partial=" << partial_bytes << " upload bytes\n"; +} } // namespace int main(int argc, char** argv) { @@ -509,6 +586,7 @@ int main(int argc, char** argv) { Pipeline(device, library, @"gaussian_sort_verify"), Pipeline(device, library, @"gaussian_metal_gather"), RasterPipeline(device, library), DepthState(device)}; + CompareResidentFrames(device, queue, kernels); Compare(device, queue, kernels, {}, false, true); const auto image = Compare(device, queue, kernels, RasterFixture(3), false, true); const auto* color = static_cast(image[0].contents); diff --git a/tests/metal_gaussian_residency_test.mm b/tests/metal_gaussian_residency_test.mm new file mode 100644 index 0000000..0f8cb65 --- /dev/null +++ b/tests/metal_gaussian_residency_test.mm @@ -0,0 +1,303 @@ +#include "../backend/merlin-metal/src/gaussian_residency.hpp" +#include + +#include +#include +#include +#include + +namespace { +using merlin::metal::GaussianResidency; +using merlin::extraction::FrameSnapshot; +using merlin::extraction::GaussianRecord; + +void Require(bool value, const char* message) { + if (!value) throw std::runtime_error(message); +} + +template +void Reject(Function function, merlin::render::RendererErrorCode code) { + try { + function(); + } catch (const merlin::render::RendererError& error) { + Require(error.code() == code, "Wrong residency error code"); + return; + } + throw std::runtime_error("Expected residency rejection"); +} + +FrameSnapshot Fixture() { + FrameSnapshot snapshot; + snapshot.source_id = 7; + snapshot.revision = 1; + GaussianRecord record; + record.gaussian = 0x100000002ULL; + record.revision = record.positions_revision = record.covariance_revision = + record.opacity_revision = record.radiance_revision = 1; + record.positions = std::make_shared>(8, merlin::Vec3{1, 2, 3}); + record.covariances = std::make_shared>( + 8, merlin::Covariance3{1, 0, 0, 1, 0, 1}); + record.opacities = std::make_shared>(8, 0.5F); + record.spherical_harmonics_degree = 3; + record.spherical_harmonics_coefficients = std::make_shared>( + 8 * 16, merlin::Vec3{0.1F, 0.2F, 0.3F}); + snapshot.gaussians.assign({record}); + return snapshot; +} + +struct Readback { + id buffer; + const void* expected; + std::size_t bytes; +}; + +std::vector Read(id device, id command, + const GaussianResidency::Resource& resource) { + std::vector reads; + auto encoder = [command blitCommandEncoder]; + const auto read = [&](const auto& payload, const auto& source) { + using Element = typename std::decay_t::value_type; + const auto bytes = payload->size() * sizeof(Element); + if (!bytes) return; + auto buffer = [device newBufferWithLength:bytes options:MTLResourceStorageModeShared]; + Require(buffer != nil, "Readback allocation failed"); + [encoder copyFromBuffer:source->metal sourceOffset:0 toBuffer:buffer destinationOffset:0 size:bytes]; + reads.push_back({buffer, payload->data(), bytes}); + }; + read(resource.record.positions, resource.positions); + read(resource.record.covariances, resource.covariances); + read(resource.record.opacities, resource.opacities); + read(resource.record.spherical_harmonics_coefficients, resource.radiance); + [encoder endEncoding]; + return reads; +} + +void Complete(id command, const std::vector& reads) { + [command waitUntilCompleted]; + Require(command.status == MTLCommandBufferStatusCompleted, "Attribute command failed"); + for (const auto& read : reads) + Require(std::memcmp(read.buffer.contents, read.expected, read.bytes) == 0, + "Resident attribute bytes differ from snapshot"); +} + +void TestResidency(id device, id queue) { + auto initial = Fixture(); + const auto full_bytes = 8 * (sizeof(merlin::Vec3) + sizeof(merlin::Covariance3) + + sizeof(float) + 16 * sizeof(merlin::Vec3)); + GaussianResidency store(device, full_bytes * 4); + auto first = store.Prepare(initial); + Require(first->upload_bytes == full_bytes && first->upload_range_count == 4 && + first->device_copy_bytes == 0 && first->allocation_count == 8, "Wrong initial upload accounting"); + Reject([&] { store.Commit(first); }, merlin::render::RendererErrorCode::InvalidRequest); + auto gate = [device newSharedEvent]; + Require(gate != nil, "Shared event allocation failed"); + struct ReleaseGate { + id event; + ~ReleaseGate() { event.signaledValue = 1; } + } release_gate{gate}; + auto first_command = [queue commandBuffer]; + [first_command encodeWaitForEvent:gate value:1]; + store.Encode(first, first_command); + const auto first_reads = Read(device, first_command, first->resources[0]); + [first_command commit]; + store.Commit(first); + + auto camera = initial; + camera.revision = 2; + camera.view.values[12] = 0.5F; + auto unchanged = store.Prepare(camera); + Require(unchanged->upload_bytes == 0 && unchanged->allocation_count == 0 && + unchanged->resources[0].positions == first->resources[0].positions, + "Camera change reuploaded resident attributes"); + auto metadata = initial.gaussians[0]; + metadata.revision = 2; + metadata.transform.values[12] = 1; + metadata.visible = false; + metadata.sorting_mode = merlin::GaussianSortingMode::CameraDistance; + camera.gaussians.assign({metadata}); + auto metadata_update = store.Prepare(camera); + Require(metadata_update->upload_bytes == 0 && metadata_update->allocation_count == 0 && + !metadata_update->resources[0].record.visible && + metadata_update->resources[0].record.transform.values[12] == 1, + "Metadata edit changed attributes or lost metadata"); + + auto edited = initial; + auto record = initial.gaussians[0]; + record.revision = 2; + record.positions_revision = record.radiance_revision = 2; + record.particle_base_revision = 1; + record.particle_ranges = {{1, 2}, {6, 1}}; + auto positions = std::make_shared>(*record.positions); + auto radiance = std::make_shared>(*record.spherical_harmonics_coefficients); + for (const auto& range : record.particle_ranges) { + for (std::size_t i = range.first; i < range.first + range.count; ++i) { + (*positions)[i] = {4, 5, 6}; + for (std::size_t j = 0; j < 16; ++j) (*radiance)[i * 16 + j] = {0.4F, 0.5F, 0.6F}; + } + } + record.positions = positions; + record.spherical_harmonics_coefficients = radiance; + edited.gaussians.assign({record}); + auto second = store.Prepare(edited); + Require(second->upload_bytes == 3 * 17 * sizeof(merlin::Vec3) && second->upload_range_count == 4 && + second->device_copy_bytes == 8 * 17 * sizeof(merlin::Vec3), "Wrong partial upload accounting"); + Require(second->resources[0].positions != first->resources[0].positions && + second->resources[0].radiance != first->resources[0].radiance && + second->resources[0].covariances == first->resources[0].covariances && + second->resources[0].opacities == first->resources[0].opacities, + "Partial update did not version only changed attributes"); + auto second_command = [queue commandBuffer]; + store.Encode(second, second_command); + const auto second_reads = Read(device, second_command, second->resources[0]); + // Read the old version AFTER the new upload on the GPU as well. + const auto old_reads = Read(device, second_command, first->resources[0]); + [second_command commit]; + store.Commit(second); + Reject([&] { store.Encode(unchanged, [queue commandBuffer]); }, + merlin::render::RendererErrorCode::InvalidRequest); + Reject([&] { store.Commit(second); }, merlin::render::RendererErrorCode::InvalidRequest); + + const auto retained_bytes = store.live_bytes(); + auto replacement = edited; + replacement.source_id = 8; + Reject([&] { store.Prepare(replacement); }, merlin::render::RendererErrorCode::ResourceExhausted); + Require(store.live_bytes() == retained_bytes && store.Prepare(edited)->upload_bytes == 0, + "In-flight budget rejection corrupted residency or leaked allocations"); + + // Release every CPU owner while both submissions are blocked. Command + // completion owns staging, old versions and their live budget accounting. + first.reset(); + second.reset(); + unchanged.reset(); + metadata_update.reset(); + store.Reset(); + Require(store.live_bytes() > full_bytes, "In-flight versions escaped live allocation accounting"); + gate.signaledValue = 1; + Complete(first_command, first_reads); + Complete(second_command, second_reads); + Complete(second_command, old_reads); +} + +void TestReconciliation(id device, id queue) { + auto snapshot = Fixture(); + GaussianResidency store(device, 1024 * 1024); + auto initial = store.Prepare(snapshot); + auto command = [queue commandBuffer]; + store.Encode(initial, command); + [command commit]; + store.Commit(initial); + Complete(command, {}); + const auto full_bytes = initial->upload_bytes; + + auto gap = snapshot; + auto record = snapshot.gaussians[0]; + record.revision = record.opacity_revision = 4; + record.particle_base_revision = 3; + record.particle_ranges = {{1, 1}}; + record.opacities = std::make_shared>(8, 0.75F); + gap.gaussians.assign({record}); + auto update = store.Prepare(gap); + Require(update->upload_bytes == 8 * sizeof(float) && update->device_copy_bytes == 0, + "Revision gap incorrectly applied only dirty ranges"); + command = [queue commandBuffer]; + store.Encode(update, command); + const auto reads = Read(device, command, update->resources[0]); + [command commit]; + store.Commit(update); + Complete(command, reads); + + auto foreign = gap; + foreign.source_id = 8; + Require(store.Prepare(foreign)->upload_bytes == full_bytes, "Foreign source reused resource versions"); + auto reused = gap; + record.gaussian += 1ULL << 32; + reused.gaussians.assign({record}); + Require(store.Prepare(reused)->upload_bytes == full_bytes, "Handle generation reused removed attributes"); + + // Aborted preparation cannot replace committed residency. + auto aborted = store.Prepare(snapshot); + aborted.reset(); + Require(store.Prepare(gap)->upload_bytes == 0, "Aborted plan changed residency"); + GaussianResidency other(device, 1024 * 1024); + Reject([&] { other.Encode(store.Prepare(gap), [queue commandBuffer]); }, + merlin::render::RendererErrorCode::InvalidRequest); + + auto invalid = gap; + record = gap.gaussians[0]; + record.particle_ranges = {{8, 1}}; + invalid.gaussians.assign({record}); + Reject([&] { store.Prepare(invalid); }, merlin::render::RendererErrorCode::InvalidRequest); + record = gap.gaussians[0]; + record.spherical_harmonics_degree = 4; + invalid.gaussians.assign({record}); + Reject([&] { store.Prepare(invalid); }, merlin::render::RendererErrorCode::InvalidRequest); + record = gap.gaussians[0]; + record.opacities.reset(); + invalid.gaussians.assign({record}); + Reject([&] { store.Prepare(invalid); }, merlin::render::RendererErrorCode::InvalidRequest); + invalid.gaussians.assign({gap.gaussians[0], gap.gaussians[0]}); + Reject([&] { store.Prepare(invalid); }, merlin::render::RendererErrorCode::InvalidRequest); + Require(store.Prepare(gap)->upload_bytes == 0, "Invalid plan changed residency"); + + // Count and SH layout changes must upload complete replacement payloads, + // even when a caller supplies a matching base revision and a small range. + auto resized = gap; + record = gap.gaussians[0]; + record.particle_base_revision = record.revision; + record.particle_ranges = {{0, 1}}; + record.positions = std::make_shared>(2); + record.covariances = std::make_shared>(2); + record.opacities = std::make_shared>(2, 0.5F); + record.spherical_harmonics_degree = 1; + record.spherical_harmonics_coefficients = std::make_shared>(8); + resized.gaussians.assign({record}); + auto resize = store.Prepare(resized); + Require(resize->upload_bytes == 2 * (sizeof(merlin::Vec3) + sizeof(merlin::Covariance3) + + sizeof(float) + 4 * sizeof(merlin::Vec3)) && resize->device_copy_bytes == 0, + "Count/SH layout change incorrectly applied partial upload"); + command = [queue commandBuffer]; + store.Encode(resize, command); + const auto resize_reads = Read(device, command, resize->resources[0]); + [command commit]; + store.Commit(resize); + Complete(command, resize_reads); + + auto empty = gap; + empty.gaussians.assign({}); + auto removal = store.Prepare(empty); + command = [queue commandBuffer]; + store.Encode(removal, command); + [command commit]; + store.Commit(removal); + Complete(command, {}); + Require(removal->resources.empty() && removal->upload_bytes == 0, "Removal uploaded attributes"); + Require(store.Prepare(gap)->upload_bytes == full_bytes, "Removed resource reused stale residency"); + + // Failure half way through allocations rolls back both state and accounting. + GaussianResidency bounded(device, 256); + Reject([&] { bounded.Prepare(snapshot); }, merlin::render::RendererErrorCode::ResourceExhausted); + Require(bounded.live_bytes() == 0, "Failed transaction leaked its allocation budget"); +} +} // namespace + +int main() { + @autoreleasepool { + try { + const auto availability = merlin::metal::BackendFactory{}.availability(); + if (!availability.available) { + std::cerr << "skip: " << availability.detail << '\n'; + return 77; + } + auto device = MTLCreateSystemDefaultDevice(); + auto queue = [device newCommandQueue]; + Require(device && queue, "Metal device/queue unavailable"); + TestResidency(device, queue); + TestReconciliation(device, queue); + std::cout << "Metal Gaussian attribute residency, partial updates and lifetime passed\n"; + return 0; + } catch (const std::exception& error) { + std::cerr << error.what() << '\n'; + return 1; + } + } +}