Share Gaussian preparation and radix sort with Metal compute

CHANGELOG.md

@@ -10,6 +10,13 @@ after its public API and release process are established.
 
 ### Added
 
+- Shared Slang Gaussian projection, covariance, SH, culling/compaction and
+  deterministic radix-sort kernels now compile into the Metal Gaussian library.
+  Native Metal buffer bindings and a scalar 64-byte prepared-record ABI have
+  reflection, install-package and Apple GPU comparisons against the CPU
+  reference. These kernels are not yet connected to Metal renderer scheduling;
+  interactive rendering still uses CPU preparation and sorting.
+
 - Metal Gaussian reference rasterization now uses shared Slang ellipse/alpha
   math with Vulkan. Slang 2026.8.x and Xcode compile an embedded metallib during
   the build; installed consumers do not compile Gaussian shaders at runtime.
@@ -283,6 +290,11 @@ after its public API and release process are established.
 
 ### Fixed
 
+- Shared GPU Gaussian preparation now applies the inverse local-to-camera
+  matrix when evaluating directional SH. The previous inverse-transpose
+  calculation changed radiance under rotation and nonuniform scale; a Metal
+  compute regression compares transformed degree-three SH with the CPU path.
+
 - Zooming into a Gaussian scene no longer floods the view with a single
   color in usdview or the development viewport. CPU and GPU preparation
   cull a kernel whose center lies in front of the camera's near plane, as

backend/merlin-metal/shaders/gaussian-prepare-metal.slang

@@ -0,0 +1,27 @@
+#include "../../../core/merlin-render-backend/shaders/gaussian-prepare-abi.slang"
+
+// Metal uses one explicit buffer namespace for constants and storage.
+[[vk::binding(7, 0)]]
+ConstantBuffer<GaussianPrepareConstants> gaussian_prepare_constants : register(b7);
+
+// Byte-addressed inputs match the existing tightly packed arena payloads:
+// float3 positions, six-float covariance, float opacity, and float3 SH terms.
+[[vk::binding(0, 0)]] ByteAddressBuffer gaussian_positions : register(t0);
+[[vk::binding(1, 0)]] ByteAddressBuffer gaussian_covariances : register(t1);
+[[vk::binding(2, 0)]] ByteAddressBuffer gaussian_opacities : register(t2);
+[[vk::binding(3, 0)]] ByteAddressBuffer gaussian_radiance : register(t3);
+[[vk::binding(4, 0)]] RWStructuredBuffer<uint> gaussian_candidate_results : register(u4);
+[[vk::binding(5, 0)]] RWByteAddressBuffer gaussian_prepared_records : register(u5);
+[[vk::binding(6, 0)]] RWStructuredBuffer<GaussianPrepareDispatchCounters> gaussian_prepare_counters : register(u6);
+
+// Raw records keep the Vulkan 64-byte ABI despite Metal float3 alignment.
+void StorePreparedRecord(uint index, GaussianPreparedRecord record)
+{
+    uint offset = index * 64u;
+    gaussian_prepared_records.Store4(offset, asuint(float4(record.center_pixels, record.radius_pixels, record.depth)));
+    gaussian_prepared_records.Store4(offset + 16u, asuint(float4(record.inverse_conic, record.opacity)));
+    gaussian_prepared_records.Store4(offset + 32u, asuint(float4(record.radiance, record.sort_key)));
+    gaussian_prepared_records.Store4(offset + 48u, uint4(record.resource_id_low, record.resource_id_high, record.particle_id, record.padding));
+}
+
+#include "../../../core/merlin-render-backend/shaders/gaussian-prepare-common.slang"

backend/merlin-metal/shaders/gaussian-sort-metal.slang

@@ -0,0 +1,37 @@
+#include "../../../core/merlin-render-backend/shaders/gaussian-sort-abi.slang"
+
+[[vk::binding(4, 0)]]
+ConstantBuffer<GaussianSortConstants> gaussian_sort_constants : register(b4);
+
+[[vk::binding(0, 0)]] RWStructuredBuffer<GaussianSortElement> gaussian_sort_source : register(u0);
+[[vk::binding(1, 0)]] RWStructuredBuffer<GaussianSortElement> gaussian_sort_destination : register(u1);
+// Verification words, per-resource visible counts, then the digit-major
+// histogram and its scan levels.
+[[vk::binding(2, 0)]] RWStructuredBuffer<uint> gaussian_sort_scan : register(u2);
+[[vk::binding(3, 0)]] ByteAddressBuffer gaussian_sort_prepared_records : register(t3);
+
+// Match the raw 64-byte stream written by Metal preparation.
+[ForceInline]
+GaussianPreparedRecord LoadSortPreparedRecord(uint index)
+{
+    uint offset = index * 64u;
+    float4 geometry = asfloat(gaussian_sort_prepared_records.Load4(offset));
+    float4 conic = asfloat(gaussian_sort_prepared_records.Load4(offset + 16u));
+    float4 radiance = asfloat(gaussian_sort_prepared_records.Load4(offset + 32u));
+    uint4 identity = gaussian_sort_prepared_records.Load4(offset + 48u);
+    GaussianPreparedRecord record;
+    record.center_pixels = geometry.xy;
+    record.radius_pixels = geometry.z;
+    record.depth = geometry.w;
+    record.inverse_conic = conic.xyz;
+    record.opacity = conic.w;
+    record.radiance = radiance.xyz;
+    record.sort_key = radiance.w;
+    record.resource_id_low = identity.x;
+    record.resource_id_high = identity.y;
+    record.particle_id = identity.z;
+    record.padding = identity.w;
+    return record;
+}
+
+#include "../../../core/merlin-render-backend/shaders/gaussian-sort-common.slang"

backend/merlin-metal/src/gaussian_compute_abi.hpp

@@ -0,0 +1,91 @@
+#pragma once
+
+#include <cstddef>
+#include <cstdint>
+
+#include <merlin/core/types.hpp>
+
+// Private host records for the shared Gaussian compute kernels. Metal accesses
+// prepared records through byte-addressed buffers to preserve the scalar ABI.
+namespace merlin::metal::gaussian_compute {
+
+struct alignas(16) PrepareConstants {
+  Mat4 local_to_camera;
+  Mat4 projection;
+  Vec2 viewport_size;
+  float sigma_extent{3.0F};
+  float minimum_variance_pixels{0.25F};
+  std::uint32_t resource_id_low{};
+  std::uint32_t resource_id_high{};
+  std::uint32_t particle_count{};
+  std::uint32_t coefficients_per_particle{};
+  std::uint32_t spherical_harmonics_degree{};
+  std::uint32_t projection_mode{};
+  std::uint32_t sorting_mode{};
+  std::uint32_t padding{};
+};
+
+struct alignas(16) PreparedRecord {
+  Vec2 center_pixels;
+  float radius_pixels{};
+  float depth{};
+  Vec3 inverse_conic;
+  float opacity{};
+  Vec3 radiance;
+  float sort_key{};
+  std::uint32_t resource_id_low{};
+  std::uint32_t resource_id_high{};
+  std::uint32_t particle_id{};
+  std::uint32_t padding{};
+};
+
+struct PrepareCounters {
+  std::uint32_t candidate_count{};
+  std::uint32_t visible_count{};
+  std::uint32_t opacity_culled_count{};
+  std::uint32_t frustum_culled_count{};
+  std::uint32_t invalid_culled_count{};
+  std::uint32_t padding[3]{};
+};
+
+struct SortElement {
+  std::uint32_t key_low{};
+  std::uint32_t key_high{};
+  std::uint32_t value{};
+};
+
+struct SortConstants {
+  std::uint32_t element_count{};
+  std::uint32_t block_count{};
+  std::uint32_t digit_shift{};
+  std::uint32_t digit_word{};
+  std::uint32_t scan_offset{};
+  std::uint32_t scan_count{};
+  std::uint32_t scan_sums_offset{};
+  std::uint32_t candidate_base{};
+  std::uint32_t prepared_base{};
+  std::uint32_t visible_count_offset{};
+  std::uint32_t count_word{};
+  std::uint32_t flags{};
+};
+
+static_assert(sizeof(PrepareConstants) == 176);
+static_assert(offsetof(PrepareConstants, projection) == 64);
+static_assert(offsetof(PrepareConstants, viewport_size) == 128);
+static_assert(offsetof(PrepareConstants, particle_count) == 152);
+static_assert(offsetof(PrepareConstants, sorting_mode) == 168);
+static_assert(sizeof(PreparedRecord) == 64);
+static_assert(offsetof(PreparedRecord, inverse_conic) == 16);
+static_assert(offsetof(PreparedRecord, opacity) == 28);
+static_assert(offsetof(PreparedRecord, radiance) == 32);
+static_assert(offsetof(PreparedRecord, sort_key) == 44);
+static_assert(offsetof(PreparedRecord, resource_id_low) == 48);
+static_assert(offsetof(PreparedRecord, particle_id) == 56);
+static_assert(sizeof(PrepareCounters) == 32);
+static_assert(sizeof(SortElement) == 12);
+static_assert(sizeof(SortConstants) == 48);
+static_assert(offsetof(SortConstants, scan_offset) == 16);
+static_assert(offsetof(SortConstants, visible_count_offset) == 36);
+static_assert(offsetof(SortConstants, flags) == 44);
+
+} // namespace merlin::metal::gaussian_compute