Metal: MTLMemoryManager implementation includes functions which manage allocation of MTLBuffer resources.

The memory manager includes both a GPUContext-local manager which allocates per-context resources such as Circular Scratch Buffers for temporary data such as uniform updates and resource staging, and a GPUContext-global memory manager which features a pooled memory allocator for efficient re-use of resources, to reduce CPU-overhead of frequent memory allocations.

These Memory Managers act as a simple interface for use by other Metal backend modules and to coordinate the lifetime of buffers, to ensure that GPU-resident resources are correctly tracked and freed when no longer in use.

Note: This also contains dependent DIFF changes from D15027, though these will be removed once D15027 lands.

Authored by Apple: Michael Parkin-White

Ref T96261

Reviewed By: fclem

Maniphest Tasks: T96261

Differential Revision: https://developer.blender.org/D15277

18 files changed, 1524 insertions, 124 deletions

diff --git a/source/blender/blenlib/BLI_math_base.h b/source/blender/blenlib/BLI_math_base.h
index f072a17f384..c0c4594ddc0 100644
--- a/source/blender/blenlib/BLI_math_base.h
+++ b/source/blender/blenlib/BLI_math_base.h
@@ -221,6 +221,19 @@ MINLINE unsigned int power_of_2_min_u(unsigned int x);
  * with integers, to avoid gradual darkening when rounding down.
  */
 MINLINE int divide_round_i(int a, int b);
+
+/**
+ * Integer division that returns the ceiling, instead of flooring like normal C division.
+ */
+MINLINE uint divide_ceil_u(uint a, uint b);
+MINLINE uint64_t divide_ceil_ul(uint64_t a, uint64_t b);
+
+/**
+ * Returns \a a if it is a multiple of \a b or the next multiple or \a b after \b a .
+ */
+MINLINE uint ceil_to_multiple_u(uint a, uint b);
+MINLINE uint64_t ceil_to_multiple_ul(uint64_t a, uint64_t b);
+
 /**
  * modulo that handles negative numbers, works the same as Python's.
  */
diff --git a/source/blender/blenlib/intern/math_base_inline.c b/source/blender/blenlib/intern/math_base_inline.c
index cb7659a7059..fb71e84c23e 100644
--- a/source/blender/blenlib/intern/math_base_inline.c
+++ b/source/blender/blenlib/intern/math_base_inline.c
@@ -370,6 +370,11 @@ MINLINE uint divide_ceil_u(uint a, uint b)
   return (a + b - 1) / b;
 }
 
+MINLINE uint64_t divide_ceil_ul(uint64_t a, uint64_t b)
+{
+  return (a + b - 1) / b;
+}
+
 /**
  * Returns \a a if it is a multiple of \a b or the next multiple or \a b after \b a .
  */
@@ -378,6 +383,11 @@ MINLINE uint ceil_to_multiple_u(uint a, uint b)
   return divide_ceil_u(a, b) * b;
 }
 
+MINLINE uint64_t ceil_to_multiple_ul(uint64_t a, uint64_t b)
+{
+  return divide_ceil_ul(a, b) * b;
+}
+
 MINLINE int mod_i(int i, int n)
 {
   return (i % n + n) % n;
diff --git a/source/blender/draw/engines/eevee/eevee_render.c b/source/blender/draw/engines/eevee/eevee_render.c
index bef19c589c2..82944f237ea 100644
--- a/source/blender/draw/engines/eevee/eevee_render.c
+++ b/source/blender/draw/engines/eevee/eevee_render.c
@@ -24,6 +24,7 @@
 #include "DEG_depsgraph_query.h"
 
 #include "GPU_capabilities.h"
+#include "GPU_context.h"
 #include "GPU_framebuffer.h"
 #include "GPU_state.h"
 
@@ -646,6 +647,10 @@ void EEVEE_render_draw(EEVEE_Data *vedata, RenderEngine *engine, RenderLayer *rl
     /* XXX Seems to fix TDR issue with NVidia drivers on linux. */
     GPU_finish();
 
+    /* Perform render step between samples to allow
+     * flushing of freed GPUBackend resources. */
+    GPU_render_step();
+
     RE_engine_update_progress(engine, (float)(render_samples++) / (float)tot_sample);
   }
 }
diff --git a/source/blender/draw/engines/workbench/workbench_render.c b/source/blender/draw/engines/workbench/workbench_render.c
index e5dcf6c5624..931f6a2dc92 100644
--- a/source/blender/draw/engines/workbench/workbench_render.c
+++ b/source/blender/draw/engines/workbench/workbench_render.c
@@ -17,6 +17,7 @@
 
 #include "ED_view3d.h"
 
+#include "GPU_context.h"
 #include "GPU_shader.h"
 
 #include "DEG_depsgraph.h"
@@ -188,6 +189,10 @@ void workbench_render(void *ved, RenderEngine *engine, RenderLayer *render_layer
 
   workbench_draw_finish(data);
 
+  /* Perform render step between samples to allow
+   * flushing of freed GPUBackend resources. */
+  GPU_render_step();
+
   /* Write render output. */
   const char *viewname = RE_GetActiveRenderView(engine->re);
   RenderPass *rp = RE_pass_find_by_name(render_layer, RE_PASSNAME_COMBINED, viewname);
diff --git a/source/blender/gpu/CMakeLists.txt b/source/blender/gpu/CMakeLists.txt
index 62d5537772a..9b5ce6e147e 100644
--- a/source/blender/gpu/CMakeLists.txt
+++ b/source/blender/gpu/CMakeLists.txt
@@ -194,6 +194,7 @@ set(METAL_SRC
   metal/mtl_command_buffer.mm
   metal/mtl_debug.mm
   metal/mtl_framebuffer.mm
+  metal/mtl_memory.mm
   metal/mtl_state.mm
   metal/mtl_texture.mm
   metal/mtl_texture_util.mm
@@ -204,6 +205,7 @@ set(METAL_SRC
   metal/mtl_context.hh
   metal/mtl_debug.hh
   metal/mtl_framebuffer.hh
+  metal/mtl_memory.hh
   metal/mtl_state.hh
   metal/mtl_texture.hh
 )
diff --git a/source/blender/gpu/intern/gpu_immediate_util.c b/source/blender/gpu/intern/gpu_immediate_util.c
index a275fd8fc6c..5233ff2dbf6 100644
--- a/source/blender/gpu/intern/gpu_immediate_util.c
+++ b/source/blender/gpu/intern/gpu_immediate_util.c
@@ -142,7 +142,7 @@ static void imm_draw_circle(GPUPrimType prim_type,
                             int nsegments)
 {
   if (prim_type == GPU_PRIM_LINE_LOOP) {
-    /* Note(Metal/AMD): For small primitives, line list more efficient than line strip.. */
+    /* NOTE(Metal/AMD): For small primitives, line list more efficient than line strip.. */
     immBegin(GPU_PRIM_LINES, nsegments * 2);
 
     immVertex2f(shdr_pos, x + (radius_x * cosf(0.0f)), y + (radius_y * sinf(0.0f)));
@@ -333,7 +333,7 @@ static void imm_draw_circle_3D(
     GPUPrimType prim_type, uint pos, float x, float y, float radius, int nsegments)
 {
   if (prim_type == GPU_PRIM_LINE_LOOP) {
-    /* Note(Metal/AMD): For small primitives, line list more efficient than line strip. */
+    /* NOTE(Metal/AMD): For small primitives, line list more efficient than line strip. */
     immBegin(GPU_PRIM_LINES, nsegments * 2);
 
     const float angle = (float)(2 * M_PI) / (float)nsegments;
@@ -386,7 +386,7 @@ void imm_draw_circle_fill_3d(uint pos, float x, float y, float radius, int nsegm
 
 void imm_draw_box_wire_2d(uint pos, float x1, float y1, float x2, float y2)
 {
-  /* Note(Metal/AMD): For small primitives, line list more efficient than line-strip. */
+  /* NOTE(Metal/AMD): For small primitives, line list more efficient than line-strip. */
   immBegin(GPU_PRIM_LINES, 8);
   immVertex2f(pos, x1, y1);
   immVertex2f(pos, x1, y2);
@@ -405,7 +405,7 @@ void imm_draw_box_wire_2d(uint pos, float x1, float y1, float x2, float y2)
 void imm_draw_box_wire_3d(uint pos, float x1, float y1, float x2, float y2)
 {
   /* use this version when GPUVertFormat has a vec3 position */
-  /* Note(Metal/AMD): For small primitives, line list more efficient than line-strip. */
+  /* NOTE(Metal/AMD): For small primitives, line list more efficient than line-strip. */
   immBegin(GPU_PRIM_LINES, 8);
   immVertex3f(pos, x1, y1, 0.0f);
   immVertex3f(pos, x1, y2, 0.0f);
diff --git a/source/blender/gpu/metal/mtl_backend.mm b/source/blender/gpu/metal/mtl_backend.mm
index 81f8f279759..117b8352a0a 100644
--- a/source/blender/gpu/metal/mtl_backend.mm
+++ b/source/blender/gpu/metal/mtl_backend.mm
@@ -127,7 +127,21 @@ void MTLBackend::render_end()
 
 void MTLBackend::render_step()
 {
-  /* Placeholder */
+  /* NOTE(Metal): Primarily called from main thread, but below datastructures
+   * and operations are thread-safe, and GPUContext rendering coordination
+   * is also thread-safe. */
+
+  /* Flush any MTLSafeFreeLists which have previously been released by any MTLContext. */
+  MTLContext::get_global_memory_manager().update_memory_pools();
+
+  /* End existing MTLSafeFreeList and begin new list --
+   * Buffers wont `free` until all associated in-flight command buffers have completed.
+   * Decrement final reference count for ensuring the previous list is certainly
+   * released. */
+  MTLSafeFreeList *cmd_free_buffer_list =
+      MTLContext::get_global_memory_manager().get_current_safe_list();
+  MTLContext::get_global_memory_manager().begin_new_safe_list();
+  cmd_free_buffer_list->decrement_reference();
 }
 
 bool MTLBackend::is_inside_render_boundary()
diff --git a/source/blender/gpu/metal/mtl_command_buffer.mm b/source/blender/gpu/metal/mtl_command_buffer.mm
index 4f6077e8159..f9edd87a73c 100644
--- a/source/blender/gpu/metal/mtl_command_buffer.mm
+++ b/source/blender/gpu/metal/mtl_command_buffer.mm
@@ -19,7 +19,7 @@ namespace blender::gpu {
  * dependencies not being honored for work submitted between
  * different GPUContext's. */
 id<MTLEvent> MTLCommandBufferManager::sync_event = nil;
-unsigned long long MTLCommandBufferManager::event_signal_val = 0;
+uint64_t MTLCommandBufferManager::event_signal_val = 0;
 
 /* Counter for active command buffers. */
 int MTLCommandBufferManager::num_active_cmd_bufs = 0;
@@ -28,10 +28,9 @@ int MTLCommandBufferManager::num_active_cmd_bufs = 0;
 /** \name MTLCommandBuffer initialization and render coordination.
  * \{ */
 
-void MTLCommandBufferManager::prepare(MTLContext *ctx, bool supports_render)
+void MTLCommandBufferManager::prepare(bool supports_render)
 {
-  context_ = ctx;
-  render_pass_state_.prepare(this, ctx);
+  render_pass_state_.reset_state();
 }
 
 void MTLCommandBufferManager::register_encoder_counters()
@@ -54,10 +53,10 @@ id<MTLCommandBuffer> MTLCommandBufferManager::ensure_begin()
       MTLCommandBufferDescriptor *desc = [[MTLCommandBufferDescriptor alloc] init];
       desc.errorOptions = MTLCommandBufferErrorOptionEncoderExecutionStatus;
       desc.retainedReferences = YES;
-      active_command_buffer_ = [context_->queue commandBufferWithDescriptor:desc];
+      active_command_buffer_ = [context_.queue commandBufferWithDescriptor:desc];
     }
     else {
-      active_command_buffer_ = [context_->queue commandBuffer];
+      active_command_buffer_ = [context_.queue commandBuffer];
     }
     [active_command_buffer_ retain];
     MTLCommandBufferManager::num_active_cmd_bufs++;
@@ -67,6 +66,10 @@ id<MTLCommandBuffer> MTLCommandBufferManager::ensure_begin()
       [active_command_buffer_ encodeWaitForEvent:this->sync_event value:this->event_signal_val];
     }
 
+    /* Ensure we begin new Scratch Buffer if we are on a new frame. */
+    MTLScratchBufferManager &mem = context_.memory_manager;
+    mem.ensure_increment_scratch_buffer();
+
     /* Reset Command buffer heuristics. */
     this->reset_counters();
   }
@@ -86,12 +89,15 @@ bool MTLCommandBufferManager::submit(bool wait)
   this->end_active_command_encoder();
   BLI_assert(active_command_encoder_type_ == MTL_NO_COMMAND_ENCODER);
 
+  /* Flush active ScratchBuffer associated with parent MTLContext. */
+  context_.memory_manager.flush_active_scratch_buffer();
+
   /*** Submit Command Buffer. ***/
   /* Strict ordering ensures command buffers are guaranteed to execute after a previous
    * one has completed. Resolves flickering when command buffers are submitted from
    * different MTLContext's. */
   if (MTLCommandBufferManager::sync_event == nil) {
-    MTLCommandBufferManager::sync_event = [context_->device newEvent];
+    MTLCommandBufferManager::sync_event = [context_.device newEvent];
     BLI_assert(MTLCommandBufferManager::sync_event);
     [MTLCommandBufferManager::sync_event retain];
   }
@@ -102,14 +108,27 @@ bool MTLCommandBufferManager::submit(bool wait)
                                       value:MTLCommandBufferManager::event_signal_val];
 
   /* Command buffer lifetime tracking. */
-  /* TODO(Metal): This routine will later be used to track released memory allocations within the
-   * lifetime of a command buffer such that memory is only released once no longer in use. */
-  id<MTLCommandBuffer> cmd_buffer_ref = [active_command_buffer_ retain];
+  /* Increment current MTLSafeFreeList reference counter to flag MTLBuffers freed within
+   * the current command buffer lifetime as used.
+   * This ensures that in-use resources are not prematurely de-referenced and returned to the
+   * available buffer pool while they are in-use by the GPU. */
+  MTLSafeFreeList *cmd_free_buffer_list =
+      MTLContext::get_global_memory_manager().get_current_safe_list();
+  BLI_assert(cmd_free_buffer_list);
+  cmd_free_buffer_list->increment_reference();
+
+  id<MTLCommandBuffer> cmd_buffer_ref = active_command_buffer_;
+  [cmd_buffer_ref retain];
+
   [cmd_buffer_ref addCompletedHandler:^(id<MTLCommandBuffer> cb) {
+    /* Upon command buffer completion, decrement MTLSafeFreeList reference count
+     * to allow buffers no longer in use by this CommandBuffer to be freed. */
+    cmd_free_buffer_list->decrement_reference();
+
     /* Release command buffer after completion callback handled. */
     [cmd_buffer_ref release];
 
-    /* Decrement active cmd buffer count. */
+    /* Decrement count. */
     MTLCommandBufferManager::num_active_cmd_bufs--;
   }];
 
@@ -516,15 +535,6 @@ bool MTLCommandBufferManager::insert_memory_barrier(eGPUBarrier barrier_bits,
 /* -------------------------------------------------------------------- */
 /** \name Render Pass State for active RenderCommandEncoder
  * \{ */
-
-/* Metal Render Pass State. */
-void MTLRenderPassState::prepare(MTLCommandBufferManager *cmd, MTLContext *mtl_ctx)
-{
-  this->cmd = cmd;
-  this->ctx = mtl_ctx;
-  this->reset_state();
-}
-
 /* Reset binding state when a new RenderCommandEncoder is bound, to ensure
  * pipeline resources are re-applied to the new Encoder.
  * NOTE: In Metal, state is only persistent within an MTLCommandEncoder,
@@ -539,12 +549,12 @@ void MTLRenderPassState::reset_state()
   this->last_bound_shader_state.set(nullptr, 0);
 
   /* Other states. */
-  MTLFrameBuffer *fb = this->cmd->get_active_framebuffer();
+  MTLFrameBuffer *fb = this->cmd.get_active_framebuffer();
   this->last_used_stencil_ref_value = 0;
   this->last_scissor_rect = {0,
                              0,
-                             (unsigned long)((fb != nullptr) ? fb->get_width() : 0),
-                             (unsigned long)((fb != nullptr) ? fb->get_height() : 0)};
+                             (uint)((fb != nullptr) ? fb->get_width() : 0),
+                             (uint)((fb != nullptr) ? fb->get_height() : 0)};
 
   /* Reset cached resource binding state */
   for (int ubo = 0; ubo < MTL_MAX_UNIFORM_BUFFER_BINDINGS; ubo++) {
@@ -573,7 +583,7 @@ void MTLRenderPassState::reset_state()
 void MTLRenderPassState::bind_vertex_texture(id<MTLTexture> tex, uint slot)
 {
   if (this->cached_vertex_texture_bindings[slot].metal_texture != tex) {
-    id<MTLRenderCommandEncoder> rec = this->cmd->get_active_render_command_encoder();
+    id<MTLRenderCommandEncoder> rec = this->cmd.get_active_render_command_encoder();
     BLI_assert(rec != nil);
     [rec setVertexTexture:tex atIndex:slot];
     this->cached_vertex_texture_bindings[slot].metal_texture = tex;
@@ -583,7 +593,7 @@ void MTLRenderPassState::bind_vertex_texture(id<MTLTexture> tex, uint slot)
 void MTLRenderPassState::bind_fragment_texture(id<MTLTexture> tex, uint slot)
 {
   if (this->cached_fragment_texture_bindings[slot].metal_texture != tex) {
-    id<MTLRenderCommandEncoder> rec = this->cmd->get_active_render_command_encoder();
+    id<MTLRenderCommandEncoder> rec = this->cmd.get_active_render_command_encoder();
     BLI_assert(rec != nil);
     [rec setFragmentTexture:tex atIndex:slot];
     this->cached_fragment_texture_bindings[slot].metal_texture = tex;
diff --git a/source/blender/gpu/metal/mtl_common.hh b/source/blender/gpu/metal/mtl_common.hh
index 8dda2c43585..28404d91b4b 100644
--- a/source/blender/gpu/metal/mtl_common.hh
+++ b/source/blender/gpu/metal/mtl_common.hh
@@ -4,8 +4,13 @@
 #define __MTL_COMMON
 
 // -- Renderer Options --
+#define MTL_MAX_DRAWABLES 3
 #define MTL_MAX_SET_BYTES_SIZE 4096
 #define MTL_FORCE_WAIT_IDLE 0
 #define MTL_MAX_COMMAND_BUFFERS 64
 
+/* Number of frames for which we retain in-flight resources such as scratch buffers.
+ * Set as number of GPU frames in flight, plus an additioanl value for extra possible CPU frame. */
+#define MTL_NUM_SAFE_FRAMES (MTL_MAX_DRAWABLES + 1)
+
 #endif
diff --git a/source/blender/gpu/metal/mtl_context.hh b/source/blender/gpu/metal/mtl_context.hh
index 1b2af6a584b..4b87b994a3d 100644
--- a/source/blender/gpu/metal/mtl_context.hh
+++ b/source/blender/gpu/metal/mtl_context.hh
@@ -12,7 +12,9 @@
 
 #include "mtl_backend.hh"
 #include "mtl_capabilities.hh"
+#include "mtl_common.hh"
 #include "mtl_framebuffer.hh"
+#include "mtl_memory.hh"
 #include "mtl_texture.hh"
 
 #include <Cocoa/Cocoa.h>
@@ -30,7 +32,6 @@ class MTLContext;
 class MTLCommandBufferManager;
 class MTLShader;
 class MTLUniformBuf;
-class MTLBuffer;
 
 /* Structs containing information on current binding state for textures and samplers. */
 struct MTLTextureBinding {
@@ -56,10 +57,13 @@ struct MTLSamplerBinding {
 struct MTLRenderPassState {
   friend class MTLContext;
 
+  MTLRenderPassState(MTLContext &context, MTLCommandBufferManager &command_buffer_manager)
+      : ctx(context), cmd(command_buffer_manager){};
+
   /* Given a RenderPassState is associated with a live RenderCommandEncoder,
    * this state sits within the MTLCommandBufferManager. */
-  MTLCommandBufferManager *cmd;
-  MTLContext *ctx;
+  MTLContext &ctx;
+  MTLCommandBufferManager &cmd;
 
   /* Caching of resource bindings for active MTLRenderCommandEncoder.
    * In Metal, resource bindings are local to the MTLCommandEncoder,
@@ -110,9 +114,6 @@ struct MTLRenderPassState {
   SamplerStateBindingCached cached_vertex_sampler_state_bindings[MTL_MAX_TEXTURE_SLOTS];
   SamplerStateBindingCached cached_fragment_sampler_state_bindings[MTL_MAX_TEXTURE_SLOTS];
 
-  /* Prepare. */
-  void prepare(MTLCommandBufferManager *cmd, MTLContext *ctx);
-
   /* Reset RenderCommandEncoder binding state. */
   void reset_state();
 
@@ -446,18 +447,6 @@ struct MTLContextGlobalShaderPipelineState {
   float line_width = 1.0f;
 };
 
-/* Metal Buffer */
-struct MTLTemporaryBufferRange {
-  id<MTLBuffer> metal_buffer;
-  void *host_ptr;
-  unsigned long long buffer_offset;
-  unsigned long long size;
-  MTLResourceOptions options;
-
-  void flush();
-  bool requires_flush();
-};
-
 /* Command Buffer Manager - Owned by MTLContext.
  * The MTLCommandBufferManager represents all work associated with
  * a command buffer of a given identity. This manager is a fixed-state
@@ -477,14 +466,14 @@ class MTLCommandBufferManager {
  public:
   /* Event to coordinate sequential execution across all "main" command buffers. */
   static id<MTLEvent> sync_event;
-  static unsigned long long event_signal_val;
+  static uint64_t event_signal_val;
 
   /* Counter for active command buffers. */
   static int num_active_cmd_bufs;
 
  private:
   /* Associated Context and properties. */
-  MTLContext *context_ = nullptr;
+  MTLContext &context_;
   bool supports_render_ = false;
 
   /* CommandBuffer tracking. */
@@ -516,7 +505,9 @@ class MTLCommandBufferManager {
   bool empty_ = true;
 
  public:
-  void prepare(MTLContext *ctx, bool supports_render = true);
+  MTLCommandBufferManager(MTLContext &context)
+      : context_(context), render_pass_state_(context, *this){};
+  void prepare(bool supports_render = true);
 
   /* If wait is true, CPU will stall until GPU work has completed. */
   bool submit(bool wait);
@@ -582,7 +573,7 @@ class MTLContext : public Context {
 
   /* Texture Samplers. */
   /* Cache of generated MTLSamplerState objects based on permutations of `eGPUSamplerState`. */
-  id<MTLSamplerState> sampler_state_cache_[GPU_SAMPLER_MAX] = {0};
+  id<MTLSamplerState> sampler_state_cache_[GPU_SAMPLER_MAX];
   id<MTLSamplerState> default_sampler_state_ = nil;
 
   /* When texture sampler count exceeds the resource bind limit, an
@@ -595,6 +586,7 @@ class MTLContext : public Context {
 
   /* Frame. */
   bool is_inside_frame_ = false;
+  uint current_frame_index_;
 
  public:
   /* Shaders and Pipeline state. */
@@ -604,6 +596,10 @@ class MTLContext : public Context {
   id<MTLCommandQueue> queue = nil;
   id<MTLDevice> device = nil;
 
+  /* Memory Management */
+  MTLScratchBufferManager memory_manager;
+  static MTLBufferPool global_memory_manager;
+
   /* CommandBuffer managers. */
   MTLCommandBufferManager main_command_buffer;
 
@@ -624,7 +620,7 @@ class MTLContext : public Context {
   void memory_statistics_get(int *total_mem, int *free_mem) override;
 
   void debug_group_begin(const char *name, int index) override;
-  void debug_group_end(void) override;
+  void debug_group_end() override;
 
   /*** MTLContext Utility functions. */
   /*
@@ -679,6 +675,21 @@ class MTLContext : public Context {
   {
     return is_inside_frame_;
   }
+
+  uint get_current_frame_index()
+  {
+    return current_frame_index_;
+  }
+
+  MTLScratchBufferManager &get_scratchbuffer_manager()
+  {
+    return this->memory_manager;
+  }
+
+  static MTLBufferPool &get_global_memory_manager()
+  {
+    return MTLContext::global_memory_manager;
+  }
 };
 
 }  // namespace blender::gpu
diff --git a/source/blender/gpu/metal/mtl_context.mm b/source/blender/gpu/metal/mtl_context.mm
index 6ecdb3f48b3..2cf70718b76 100644
--- a/source/blender/gpu/metal/mtl_context.mm
+++ b/source/blender/gpu/metal/mtl_context.mm
@@ -16,44 +16,25 @@ using namespace blender::gpu;
 
 namespace blender::gpu {
 
-/* -------------------------------------------------------------------- */
-/** \name Memory Management
- * \{ */
-
-bool MTLTemporaryBufferRange::requires_flush()
-{
-  /* We do not need to flush shared memory. */
-  return this->options & MTLResourceStorageModeManaged;
-}
-
-void MTLTemporaryBufferRange::flush()
-{
-  if (this->requires_flush()) {
-    BLI_assert(this->metal_buffer);
-    BLI_assert((this->buffer_offset + this->size) <= [this->metal_buffer length]);
-    BLI_assert(this->buffer_offset >= 0);
-    [this->metal_buffer
-        didModifyRange:NSMakeRange(this->buffer_offset, this->size - this->buffer_offset)];
-  }
-}
-
-/** \} */
+/* Global memory mamnager */
+MTLBufferPool MTLContext::global_memory_manager;
 
 /* -------------------------------------------------------------------- */
 /** \name MTLContext
  * \{ */
 
 /* Placeholder functions */
-MTLContext::MTLContext(void *ghost_window)
+MTLContext::MTLContext(void *ghost_window) : memory_manager(*this), main_command_buffer(*this)
 {
   /* Init debug. */
   debug::mtl_debug_init();
 
   /* Initialize command buffer state. */
-  this->main_command_buffer.prepare(this);
+  this->main_command_buffer.prepare();
 
   /* Frame management. */
   is_inside_frame_ = false;
+  current_frame_index_ = 0;
 
   /* Create FrameBuffer handles. */
   MTLFrameBuffer *mtl_front_left = new MTLFrameBuffer(this, "front_left");
@@ -65,9 +46,14 @@ MTLContext::MTLContext(void *ghost_window)
    * initialization). */
   MTLBackend::platform_init(this);
   MTLBackend::capabilities_init(this);
+
   /* Initialize Metal modules. */
+  this->memory_manager.init();
   this->state_manager = new MTLStateManager(this);
 
+  /* Ensure global memory manager is initialied */
+  MTLContext::global_memory_manager.init(this->device);
+
   /* Initialize texture read/update structures. */
   this->get_texture_utils().init();
 
@@ -93,7 +79,7 @@ MTLContext::~MTLContext()
     this->finish();
 
     /* End frame. */
-    if (is_inside_frame_) {
+    if (this->get_inside_frame()) {
       this->end_frame();
     }
   }
@@ -112,7 +98,7 @@ MTLContext::~MTLContext()
 void MTLContext::begin_frame()
 {
   BLI_assert(MTLBackend::get()->is_inside_render_boundary());
-  if (is_inside_frame_) {
+  if (this->get_inside_frame()) {
     return;
   }
 
@@ -122,7 +108,7 @@ void MTLContext::begin_frame()
 
 void MTLContext::end_frame()
 {
-  BLI_assert(is_inside_frame_);
+  BLI_assert(this->get_inside_frame());
 
   /* Ensure pre-present work is committed. */
   this->flush();
@@ -136,20 +122,20 @@ void MTLContext::check_error(const char *info)
   /* TODO(Metal): Implement. */
 }
 
-void MTLContext::activate(void)
+void MTLContext::activate()
 {
   /* TODO(Metal): Implement. */
 }
-void MTLContext::deactivate(void)
+void MTLContext::deactivate()
 {
   /* TODO(Metal): Implement. */
 }
 
-void MTLContext::flush(void)
+void MTLContext::flush()
 {
   /* TODO(Metal): Implement. */
 }
-void MTLContext::finish(void)
+void MTLContext::finish()
 {
   /* TODO(Metal): Implement. */
 }
@@ -180,7 +166,7 @@ id<MTLRenderCommandEncoder> MTLContext::ensure_begin_render_pass()
   BLI_assert(this);
 
   /* Ensure the rendering frame has started. */
-  if (!is_inside_frame_) {
+  if (!this->get_inside_frame()) {
     this->begin_frame();
   }
 
diff --git a/source/blender/gpu/metal/mtl_framebuffer.mm b/source/blender/gpu/metal/mtl_framebuffer.mm
index 22de255bf63..b0a90829c0a 100644
--- a/source/blender/gpu/metal/mtl_framebuffer.mm
+++ b/source/blender/gpu/metal/mtl_framebuffer.mm
@@ -756,7 +756,7 @@ void MTLFrameBuffer::update_attachments(bool update_viewport)
   dirty_attachments_ = false;
 }
 
-void MTLFrameBuffer::apply_state(void)
+void MTLFrameBuffer::apply_state()
 {
   MTLContext *mtl_ctx = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
   BLI_assert(mtl_ctx);
diff --git a/source/blender/gpu/metal/mtl_memory.hh b/source/blender/gpu/metal/mtl_memory.hh
new file mode 100644
index 00000000000..81793b0647c
--- /dev/null
+++ b/source/blender/gpu/metal/mtl_memory.hh
@@ -0,0 +1,476 @@
+
+#pragma once
+
+#include <atomic>
+#include <functional>
+#include <map>
+#include <mutex>
+#include <set>
+#include <unordered_map>
+
+#include "mtl_common.hh"
+
+#include <Cocoa/Cocoa.h>
+#include <Metal/Metal.h>
+#include <QuartzCore/QuartzCore.h>
+
+@class CAMetalLayer;
+@class MTLCommandQueue;
+@class MTLRenderPipelineState;
+
+/* Metal Memory Manager Overview. */
+/*
+ * The Metal Backend Memory manager is designed to provide an interface
+ * for all other MTL_* modules where memory allocation is required.
+ *
+ * Different allocation strategies and datastructures are used depending
+ * on how the data is used by the backend. These aim to optimally handle
+ * system memory and abstract away any complexity from the MTL_* modules
+ * themselves.
+ *
+ * There are two primary allocation modes which can be used:
+ *
+ * ** MTLScratchBufferManager **
+ *
+ *    Each MTLContext owns a ScratchBufferManager which is implemented
+ *    as a pool of circular buffers, designed to handle temporary
+ *    memory allocations which occur on a per-frame basis. The scratch
+ *    buffers allow flushing of host memory to the GPU to be batched.
+ *
+ *    Each frame, the next scratch buffer is reset, then later flushed upon
+ *    command buffer submission.
+ *
+ *    Note: This is allocated per-context due to allocations being tied
+ *    to workload submissions and context-specific submissions.
+ *
+ *    Examples of scratch buffer usage are:
+ *      - Immediate-mode temporary vertex buffers.
+ *      - Shader uniform data updates
+ *      - Staging of data for resource copies, or, data reads/writes.
+ *
+ *  Usage:
+ *
+ *    MTLContext::get_scratchbuffer_manager() - to fetch active manager.
+ *
+ *    MTLTemporaryBuffer scratch_buffer_allocate_range(size)
+ *    MTLTemporaryBuffer scratch_buffer_allocate_range_aligned(size, align)
+ *
+ * ---------------------------------------------------------------------------------
+ *  ** MTLBufferPool **
+ *
+ *    For static and longer-lasting memory allocations, such as those for UBOs,
+ *    Vertex buffers, index buffers, etc; We want an optimal abstraction for
+ *    fetching a MTLBuffer of the desired size and resource options.
+ *
+ *    Memory allocations can be expensive so the MTLBufferPool provides
+ *    functionality to track usage of these buffers and once a buffer
+ *    is no longer in use, it is returned to the buffer pool for use
+ *    by another backend resource.
+ *
+ *    The MTLBufferPool provides functionality for safe tracking of resources,
+ *    as buffers freed on the host side must have their usage by the GPU tracked,
+ *    to ensure they are not prematurely re-used before they have finished being
+ *    used by the GPU.
+ *
+ *    Note: The MTLBufferPool is a global construct which can be fetched from anywhere.
+ *
+ *  Usage:
+ *    MTLContext::get_global_memory_manager();  - static routine to fetch global memory manager.
+ *
+ *    gpu::MTLBuffer *allocate_buffer(size, is_cpu_visibile, bytes=nullptr)
+ *    gpu::MTLBuffer *allocate_buffer_aligned(size, alignment, is_cpu_visibile, bytes=nullptr)
+ */
+
+/* Debug memory statistics: Disabled by Macro rather than guarded for
+ * performance considerations. */
+#define MTL_DEBUG_MEMORY_STATISTICS 0
+
+/* Allows a scratch buffer to temporarily grow beyond its maximum, which allows submission
+ * of one-time-use data packets which are too large. */
+#define MTL_SCRATCH_BUFFER_ALLOW_TEMPORARY_EXPANSION 1
+
+namespace blender::gpu {
+
+/* Forward Declarations. */
+class MTLContext;
+class MTLCommandBufferManager;
+class MTLUniformBuf;
+
+/* -------------------------------------------------------------------- */
+/** \name Memory Management.
+ * \{ */
+
+/* MTLBuffer allocation wrapper. */
+class MTLBuffer {
+
+ private:
+  /* Metal resource. */
+  id<MTLBuffer> metal_buffer_;
+
+  /* Host-visible mapped-memory pointer. Behaviour depends on buffer type:
+   * - Shared buffers: pointer represents base address of MTLBuffer whose data
+   *                   access has shared access by both the CPU and GPU on
+   *                   Unified Memory Architectures (UMA).
+   * - Managed buffer: Host-side mapped buffer region for CPU (Host) access. Managed buffers
+   *                   must be manually flushed to transfer data to GPU-resident buffer.
+   * - Private buffer: Host access is invalid, `data` will be nullptr. */
+  void *data_;
+
+  /* Whether buffer is allocated from an external source. */
+  bool is_external_ = false;
+
+  /* Allocation info. */
+  MTLResourceOptions options_;
+  id<MTLDevice> device_;
+  uint64_t alignment_;
+  uint64_t size_;
+
+  /* Allocated size may be larger than actual size. */
+  uint64_t usage_size_;
+
+  /* Lifetime info - whether the current buffer is actively in use. A buffer
+   * should be in use after it has been allocated. De-allocating the buffer, and
+   * returning it to the free buffer pool will set in_use to false. Using a buffer
+   * while it is not in-use should not be allowed and result in an error. */
+  std::atomic<bool> in_use_;
+
+ public:
+  MTLBuffer(id<MTLDevice> device, uint64_t size, MTLResourceOptions options, uint alignment = 1);
+  MTLBuffer(id<MTLBuffer> external_buffer);
+  ~MTLBuffer();
+
+  /* Fetch information about backing MTLBuffer. */
+  id<MTLBuffer> get_metal_buffer() const;
+  void *get_host_ptr() const;
+  uint64_t get_size_used() const;
+  uint64_t get_size() const;
+
+  /* Flush data to GPU. */
+  void flush();
+  void flush_range(uint64_t offset, uint64_t length);
+  bool requires_flush();
+
+  /* Buffer usage tracking. */
+  void flag_in_use(bool used);
+  bool get_in_use();
+  void set_usage_size(uint64_t size_used);
+
+  /* Debug. */
+  void set_label(NSString *str);
+
+  /* Read properties. */
+  MTLResourceOptions get_resource_options();
+  uint64_t get_alignment();
+
+  /* Resource-local free: For buffers allocated via memory manager,
+   * this will call the context `free_buffer` method to return the buffer to the context memory
+   * pool.
+   *
+   * Otherwise, free will release the associated metal resource.
+   * As a note, calling the destructor will also destroy the buffer and associated metal
+   * resource. */
+  void free();
+
+  /* Safety check to ensure buffers are not used after free. */
+  void debug_ensure_used();
+};
+
+/* View into part of an MTLBuffer. */
+struct MTLBufferRange {
+  id<MTLBuffer> metal_buffer;
+  void *data;
+  uint64_t buffer_offset;
+  uint64_t size;
+  MTLResourceOptions options;
+
+  void flush();
+  bool requires_flush();
+};
+
+/* Circular scratch buffer allocations should be seen as temporary and only used within the
+ * lifetime of the frame. */
+using MTLTemporaryBuffer = MTLBufferRange;
+
+/* Round-Robin Circular-buffer. */
+class MTLCircularBuffer {
+  friend class MTLScratchBufferManager;
+
+ private:
+  MTLContext &own_context_;
+
+  /* Wrapped MTLBuffer allocation handled. */
+  gpu::MTLBuffer *cbuffer_;
+
+  /* Current offset where next allocation will begin. */
+  uint64_t current_offset_;
+
+  /* Whether the Circular Buffer can grow during re-allocation if
+   * the size is exceeded. */
+  bool can_resize_;
+
+  /* Usage information. */
+  uint64_t used_frame_index_;
+  uint64_t last_flush_base_offset_;
+
+ public:
+  MTLCircularBuffer(MTLContext &ctx, uint64_t initial_size, bool allow_grow);
+  ~MTLCircularBuffer();
+  MTLTemporaryBuffer allocate_range(uint64_t alloc_size);
+  MTLTemporaryBuffer allocate_range_aligned(uint64_t alloc_size, uint alignment);
+  void flush();
+
+  /* Reset pointer back to start of circular buffer. */
+  void reset();
+};
+
+/* Wrapper struct used by Memory Manager to sort and compare gpu::MTLBuffer resources inside the
+ * memory pools. */
+struct MTLBufferHandle {
+  gpu::MTLBuffer *buffer;
+  uint64_t buffer_size;
+
+  inline MTLBufferHandle(gpu::MTLBuffer *buf)
+  {
+    this->buffer = buf;
+    this->buffer_size = this->buffer->get_size();
+  }
+
+  inline MTLBufferHandle(uint64_t compare_size)
+  {
+    this->buffer = nullptr;
+    this->buffer_size = compare_size;
+  }
+};
+
+struct CompareMTLBuffer {
+  bool operator()(const MTLBufferHandle &lhs, const MTLBufferHandle &rhs) const
+  {
+    return lhs.buffer_size < rhs.buffer_size;
+  }
+};
+
+/* An MTLSafeFreeList is a temporary list of gpu::MTLBuffers which have
+ * been freed by the high level backend, but are pending GPU work execution before
+ * the gpu::MTLBuffers can be returned to the Memory manager pools.
+ * This list is implemented as a chunked linked-list.
+ *
+ * Only a single MTLSafeFreeList is active at one time and is associated with current command
+ * buffer submissions. If an MTLBuffer is freed during the lifetime of a command buffer, it could
+ * still possibly be in-use and as such, the MTLSafeFreeList will increment its reference count for
+ * each command buffer submitted while the current pool is active.
+ *
+ * -- Reference count is incremented upon MTLCommandBuffer commit.
+ * -- Reference count is decremented in the MTLCommandBuffer completion callback handler.
+ *
+ * A new MTLSafeFreeList will begin each render step (frame). This pooling of buffers, rather than
+ * individual buffer resource tracking reduces performance overhead.
+ *
+ *  * The reference count starts at 1 to ensure that the reference count cannot prematurely reach
+ *  zero until any command buffers have been submitted. This additional decrement happens
+ *  when the next MTLSafeFreeList is created, to allow the existing pool to be released once
+ *  the reference count hits zero after submitted command buffers complete.
+ *
+ * Note: the Metal API independently tracks resources used by command buffers for the purpose of
+ * keeping resources alive while in-use by the driver and CPU, however, this differs from the
+ * MTLSafeFreeList mechanism in the Metal backend, which exists for the purpose of allowing
+ * previously allocated MTLBuffer resources to be re-used. This allows us to save on the expensive
+ * cost of memory allocation.
+ */
+class MTLSafeFreeList {
+  friend class MTLBufferPool;
+
+ private:
+  std::atomic<int> reference_count_;
+  std::atomic<bool> in_free_queue_;
+  std::recursive_mutex lock_;
+
+  /* Linked list of next MTLSafeFreeList chunk if current chunk is full. */
+  std::atomic<int> has_next_pool_;
+  std::atomic<MTLSafeFreeList *> next_;
+
+  /* Lockless list. MAX_NUM_BUFFERS_ within a chunk based on considerations
+   * for performance and memory. */
+  static const int MAX_NUM_BUFFERS_ = 1024;
+  std::atomic<int> current_list_index_;
+  gpu::MTLBuffer *safe_free_pool_[MAX_NUM_BUFFERS_];
+
+ public:
+  MTLSafeFreeList();
+
+  /* Add buffer to Safe Free List, can be called from secondary threads.
+   * Performs a lockless list insert. */
+  void insert_buffer(gpu::MTLBuffer *buffer);
+
+  /* Increments command buffer reference count. */
+  void increment_reference();
+
+  /* Decrement and return of buffers to pool occur on MTLCommandBuffer completion callback thread.
+   */
+  void decrement_reference();
+
+  void flag_in_queue()
+  {
+    in_free_queue_ = true;
+    if (has_next_pool_) {
+      MTLSafeFreeList *next_pool = next_.load();
+      BLI_assert(next_pool != nullptr);
+      next_pool->flag_in_queue();
+    }
+  }
+};
+
+/* MTLBuffer pools. */
+/* Allocating Metal buffers is expensive, so we cache all allocated buffers,
+ * and when requesting a new buffer, find one which fits the required dimensions
+ * from an existing pool of buffers.
+ *
+ * When freeing MTLBuffers, we insert them into the current MTLSafeFreeList, which defers
+ * release of the buffer until the associated command buffers have finished executing.
+ * This prevents a buffer from being re-used while it is still in-use by the GPU.
+ *
+ * * Once command buffers complete, MTLSafeFreeList's associated with the current
+ *   command buffer submission are added to the `completed_safelist_queue_`.
+ *
+ * * At a set point in time, all MTLSafeFreeList's in `completed_safelist_queue_` have their
+ *   MTLBuffers re-inserted into the Memory Manager's pools. */
+class MTLBufferPool {
+
+ private:
+  /* Memory statistics. */
+  long long int total_allocation_bytes_ = 0;
+
+#if MTL_DEBUG_MEMORY_STATISTICS == 1
+  /* Debug statistics. */
+  std::atomic<int> per_frame_allocation_count_;
+  std::atomic<long long int> allocations_in_pool_;
+  std::atomic<long long int> buffers_in_pool_;
+#endif
+
+  /* Metal resources. */
+  bool ensure_initialised_ = false;
+  id<MTLDevice> device_ = nil;
+
+  /* The buffer selection aims to pick a buffer which meets the minimum size requierments.
+   * To do this, we keep an ordered set of all available buffers. If the buffer is larger than the
+   * desired allocation size, we check it aginst `mtl_buffer_size_threshold_factor_`, which defines
+   * what % larger than the original allocation the buffer can be.
+   * - A higher value results in greater re-use of previously allocated buffers of similar sizes.
+   * - A lower value may result in more dynamic allocations, but minimised memory usage for a given
+   *   scenario.
+   * The current value of 1.26 is calibrated for optimal performance and memory utilisation. */
+  static constexpr float mtl_buffer_size_threshold_factor_ = 1.26;
+
+  /* Buffer pools using MTLResourceOptions as key for allocation type.
+   * Aliased as 'uint64_t' for map type compatibility.
+   * - A size-ordered list (MultiSet) of allocated buffers is kept per MTLResourceOptions
+   *   permutation. This allows efficient lookup for buffers of a given requested size.
+   * - MTLBufferHandle wraps a gpu::MTLBuffer pointer to achieve easy size-based sorting
+   *   via CompareMTLBuffer. */
+  using MTLBufferPoolOrderedList = std::multiset<MTLBufferHandle, CompareMTLBuffer>;
+  using MTLBufferResourceOptions = uint64_t;
+
+  blender::Map<MTLBufferResourceOptions, MTLBufferPoolOrderedList *> buffer_pools_;
+  blender::Vector<gpu::MTLBuffer *> allocations_;
+
+  /* Maintain a queue of all MTLSafeFreeList's that have been released
+   * by the GPU and are ready to have their buffers re-inserted into the
+   * MemoryManager pools.
+   * Access to this queue is made thread-safe through safelist_lock_. */
+  std::mutex safelist_lock_;
+  blender::Vector<MTLSafeFreeList *> completed_safelist_queue_;
+
+  /* Current free list, associated with active MTLCommandBuffer submission. */
+  /* MTLBuffer::free() can be called from separate threads, due to usage within animation
+   * system/worker threads. */
+  std::atomic<MTLSafeFreeList *> current_free_list_;
+
+ public:
+  void init(id<MTLDevice> device);
+  ~MTLBufferPool();
+
+  gpu::MTLBuffer *allocate_buffer(uint64_t size, bool cpu_visible, const void *bytes = nullptr);
+  gpu::MTLBuffer *allocate_buffer_aligned(uint64_t size,
+                                          uint alignment,
+                                          bool cpu_visible,
+                                          const void *bytes = nullptr);
+  bool free_buffer(gpu::MTLBuffer *buffer);
+
+  /* Flush MTLSafeFreeList buffers, for completed lists in `completed_safelist_queue_`,
+   * back to memory pools. */
+  void update_memory_pools();
+
+  /* Access and control over active MTLSafeFreeList. */
+  MTLSafeFreeList *get_current_safe_list();
+  void begin_new_safe_list();
+
+  /* Add a completed MTLSafeFreeList to completed_safelist_queue_. */
+  void push_completed_safe_list(MTLSafeFreeList *list);
+
+ private:
+  void ensure_buffer_pool(MTLResourceOptions options);
+  void insert_buffer_into_pool(MTLResourceOptions options, gpu::MTLBuffer *buffer);
+  void free();
+};
+
+/* Scratch buffers are circular-buffers used for temporary data within the current frame.
+ * In order to preserve integrity of contents when having multiple-frames-in-flight,
+ * we cycle through a collection of scratch buffers which are reset upon next use.
+ *
+ * Below are a series of properties, declared to manage scratch buffers. If a scratch buffer
+ * overflows, then the original buffer will be flushed and submitted, with retained references
+ * by usage within the command buffer, and a new buffer will be created.
+ * - The new buffer will grow in size to account for increased demand in temporary memory.
+ */
+class MTLScratchBufferManager {
+
+ private:
+  /* Maximum number of scratch buffers to allocate. This should be the maximum number of
+   * simultaneous frames in flight. */
+  static constexpr uint mtl_max_scratch_buffers_ = MTL_NUM_SAFE_FRAMES;
+
+ public:
+  /* Maximum size of single scratch buffer allocation. When re-sizing, this is the maximum size the
+   * newly allocated buffers will grow to. Larger allocations are possible if
+   * `MTL_SCRATCH_BUFFER_ALLOW_TEMPORARY_EXPANSION` is enabled, but these will instead allocate new
+   * buffers from the memory pools on the fly. */
+  static constexpr uint mtl_scratch_buffer_max_size_ = 128 * 1024 * 1024;
+
+  /* Initial size of circular scratch buffers prior to growth. */
+  static constexpr uint mtl_scratch_buffer_initial_size_ = 16 * 1024 * 1024;
+
+ private:
+  /* Parent MTLContext. */
+  MTLContext &context_;
+  bool initialised_ = false;
+
+  /* Scratch buffer currently in-use. */
+  uint current_scratch_buffer_ = 0;
+
+  /* Scratch buffer pool. */
+  MTLCircularBuffer *scratch_buffers_[mtl_max_scratch_buffers_];
+
+ public:
+  MTLScratchBufferManager(MTLContext &context) : context_(context){};
+  ~MTLScratchBufferManager();
+
+  /* Explicit initialisation and freeing of resources. Init must occur after device creation. */
+  void init();
+  void free();
+
+  /* Allocation functions for creating temporary allocations from active circular buffer. */
+  MTLTemporaryBuffer scratch_buffer_allocate_range(uint64_t alloc_size);
+  MTLTemporaryBuffer scratch_buffer_allocate_range_aligned(uint64_t alloc_size, uint alignment);
+
+  /* Ensure a new scratch buffer is started if we move onto a new frame.
+   * Called when a new command buffer begins. */
+  void ensure_increment_scratch_buffer();
+
+  /* Flush memory for active scratch buffer to GPU.
+   * This call will perform a partial flush of the buffer starting from
+   * the last offset the data was flushed from, to the current offset. */
+  void flush_active_scratch_buffer();
+};
+
+/** \} */
+
+}  // namespace blender::gpu
diff --git a/source/blender/gpu/metal/mtl_memory.mm b/source/blender/gpu/metal/mtl_memory.mm
new file mode 100644
index 00000000000..5c5938997e6
--- /dev/null
+++ b/source/blender/gpu/metal/mtl_memory.mm
@@ -0,0 +1,880 @@
+
+#include "BKE_global.h"
+
+#include "DNA_userdef_types.h"
+
+#include "mtl_context.hh"
+#include "mtl_debug.hh"
+#include "mtl_memory.hh"
+
+using namespace blender;
+using namespace blender::gpu;
+
+namespace blender::gpu {
+
+/* -------------------------------------------------------------------- */
+/** \name Memory Management - MTLBufferPool and MTLSafeFreeList implementations. */
+
+void MTLBufferPool::init(id<MTLDevice> mtl_device)
+{
+  if (!ensure_initialised_) {
+    BLI_assert(mtl_device);
+    ensure_initialised_ = true;
+    device_ = mtl_device;
+
+#if MTL_DEBUG_MEMORY_STATISTICS == 1
+    /* Debug statistics. */
+    per_frame_allocation_count_ = 0;
+    allocations_in_pool_ = 0;
+    buffers_in_pool_ = 0;
+#endif
+
+    /* Free pools -- Create initial safe free pool */
+    BLI_assert(current_free_list_ == nullptr);
+    this->begin_new_safe_list();
+  }
+}
+
+MTLBufferPool::~MTLBufferPool()
+{
+  this->free();
+}
+
+void MTLBufferPool::free()
+{
+
+  for (auto buffer : allocations_) {
+    BLI_assert(buffer);
+    delete buffer;
+  }
+  allocations_.clear();
+
+  for (std::multiset<blender::gpu::MTLBufferHandle, blender::gpu::CompareMTLBuffer> *buffer_pool :
+       buffer_pools_.values()) {
+    delete buffer_pool;
+  }
+  buffer_pools_.clear();
+}
+
+gpu::MTLBuffer *MTLBufferPool::allocate_buffer(uint64_t size, bool cpu_visible, const void *bytes)
+{
+  /* Allocate buffer with default HW-compatible alignemnt of 256 bytes.
+   * See https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf for more. */
+  return this->allocate_buffer_aligned(size, 256, cpu_visible, bytes);
+}
+
+gpu::MTLBuffer *MTLBufferPool::allocate_buffer_aligned(uint64_t size,
+                                                       uint alignment,
+                                                       bool cpu_visible,
+                                                       const void *bytes)
+{
+  /* Check not required. Main GPU module usage considered thread-safe. */
+  // BLI_assert(BLI_thread_is_main());
+
+  /* Calculate aligned size */
+  BLI_assert(alignment > 0);
+  uint64_t aligned_alloc_size = ceil_to_multiple_ul(size, alignment);
+
+  /* Allocate new MTL Buffer */
+  MTLResourceOptions options;
+  if (cpu_visible) {
+    options = ([device_ hasUnifiedMemory]) ? MTLResourceStorageModeShared :
+                                             MTLResourceStorageModeManaged;
+  }
+  else {
+    options = MTLResourceStorageModePrivate;
+  }
+
+  /* Check if we have a suitable buffer */
+  gpu::MTLBuffer *new_buffer = nullptr;
+  std::multiset<MTLBufferHandle, CompareMTLBuffer> **pool_search = buffer_pools_.lookup_ptr(
+      (uint64_t)options);
+
+  if (pool_search != nullptr) {
+    std::multiset<MTLBufferHandle, CompareMTLBuffer> *pool = *pool_search;
+    MTLBufferHandle size_compare(aligned_alloc_size);
+    auto result = pool->lower_bound(size_compare);
+    if (result != pool->end()) {
+      /* Potential buffer found, check if within size threshold requirements. */
+      gpu::MTLBuffer *found_buffer = result->buffer;
+      BLI_assert(found_buffer);
+      BLI_assert(found_buffer->get_metal_buffer());
+
+      uint64_t found_size = found_buffer->get_size();
+
+      if (found_size >= aligned_alloc_size &&
+          found_size <= (aligned_alloc_size * mtl_buffer_size_threshold_factor_)) {
+        MTL_LOG_INFO(
+            "[MemoryAllocator] Suitable Buffer of size %lld found, for requested size: %lld\n",
+            found_size,
+            aligned_alloc_size);
+
+        new_buffer = found_buffer;
+        BLI_assert(!new_buffer->get_in_use());
+
+        /* Remove buffer from free set. */
+        pool->erase(result);
+      }
+      else {
+        MTL_LOG_INFO(
+            "[MemoryAllocator] Buffer of size %lld found, but was incompatible with requested "
+            "size: "
+            "%lld\n",
+            found_size,
+            aligned_alloc_size);
+        new_buffer = nullptr;
+      }
+    }
+  }
+
+  /* Allocate new buffer. */
+  if (new_buffer == nullptr) {
+    new_buffer = new gpu::MTLBuffer(device_, size, options, alignment);
+
+    /* Track allocation in context. */
+    allocations_.append(new_buffer);
+    total_allocation_bytes_ += aligned_alloc_size;
+  }
+  else {
+    /* Re-use suitable buffer. */
+    new_buffer->set_usage_size(aligned_alloc_size);
+
+#if MTL_DEBUG_MEMORY_STATISTICS == 1
+    /* Debug. */
+    allocations_in_pool_ -= new_buffer->get_size();
+    buffers_in_pool_--;
+    BLI_assert(allocations_in_pool_ >= 0);
+#endif
+
+    /* Ensure buffer memory is correctly backed. */
+    BLI_assert(new_buffer->get_metal_buffer());
+  }
+  /* Flag buffer as actively in-use. */
+  new_buffer->flag_in_use(true);
+
+  /* Upload initial data if provided -- Size based on original size param, not aligned size*/
+  if (bytes) {
+    BLI_assert(!(options & MTLResourceStorageModePrivate));
+    BLI_assert(size <= aligned_alloc_size);
+    BLI_assert(size <= [new_buffer->get_metal_buffer() length]);
+    memcpy(new_buffer->get_host_ptr(), bytes, size);
+    new_buffer->flush_range(0, size);
+  }
+
+#if MTL_DEBUG_MEMORY_STATISTICS == 1
+  this->per_frame_allocation_count++;
+#endif
+
+  return new_buffer;
+}
+
+bool MTLBufferPool::free_buffer(gpu::MTLBuffer *buffer)
+{
+  /* Ensure buffer is flagged as in-use. I.e. has not already been returned to memory pools. */
+  bool buffer_in_use = buffer->get_in_use();
+  BLI_assert(buffer_in_use);
+  if (buffer_in_use) {
+
+    /* Fetch active safe pool from atomic ptr. */
+    MTLSafeFreeList *current_pool = this->get_current_safe_list();
+
+    /* Place buffer in safe_free_pool before returning to MemoryManager buffer pools. */
+    BLI_assert(current_pool);
+    current_pool->insert_buffer(buffer);
+    buffer->flag_in_use(false);
+
+    return true;
+  }
+  return false;
+}
+
+void MTLBufferPool::update_memory_pools()
+{
+  /* Ensure thread-safe access to `completed_safelist_queue_`, which contains
+   * the list of MTLSafeFreeList's whose buffers are ready to be
+   * re-inserted into the Memory Manager pools. */
+  safelist_lock_.lock();
+
+#if MTL_DEBUG_MEMORY_STATISTICS == 1
+  int num_buffers_added = 0;
+#endif
+
+  /* Always free oldest MTLSafeFreeList first. */
+  for (int safe_pool_free_index = 0; safe_pool_free_index < completed_safelist_queue_.size();
+       safe_pool_free_index++) {
+    MTLSafeFreeList *current_pool = completed_safelist_queue_[safe_pool_free_index];
+
+    /* Iterate through all MTLSafeFreeList linked-chunks. */
+    while (current_pool != nullptr) {
+      current_pool->lock_.lock();
+      BLI_assert(current_pool);
+      BLI_assert(current_pool->in_free_queue_);
+      int counter = 0;
+      int size = min_ii(current_pool->current_list_index_, MTLSafeFreeList::MAX_NUM_BUFFERS_);
+
+      /* Re-add all buffers within frame index to MemoryManager pools. */
+      while (counter < size) {
+
+        gpu::MTLBuffer *buf = current_pool->safe_free_pool_[counter];
+
+        /* Insert buffer back into open pools. */
+        BLI_assert(buf->get_in_use() == false);
+        this->insert_buffer_into_pool(buf->get_resource_options(), buf);
+        counter++;
+
+#if MTL_DEBUG_MEMORY_STATISTICS == 1
+        num_buffers_added++;
+#endif
+      }
+
+      /* Fetch next MTLSafeFreeList chunk, if any. */
+      MTLSafeFreeList *next_list = nullptr;
+      if (current_pool->has_next_pool_ > 0) {
+        next_list = current_pool->next_.load();
+      }
+
+      /* Delete current MTLSafeFreeList */
+      current_pool->lock_.unlock();
+      delete current_pool;
+      current_pool = nullptr;
+
+      /* Move onto next chunk. */
+      if (next_list != nullptr) {
+        current_pool = next_list;
+      }
+    }
+  }
+
+#if MTL_DEBUG_MEMORY_STATISTICS == 1
+  printf("--- Allocation Stats ---\n");
+  printf("  Num buffers processed in pool (this frame): %u\n", num_buffers_added);
+
+  uint framealloc = (uint)this->per_frame_allocation_count;
+  printf("  Allocations in frame: %u\n", framealloc);
+  printf("  Total Buffers allocated: %u\n", (uint)allocations_.size());
+  printf("  Total Memory allocated: %u MB\n", (uint)total_allocation_bytes_ / (1024 * 1024));
+
+  uint allocs = (uint)(allocations_in_pool_) / 1024 / 2024;
+  printf("  Free memory in pools: %u MB\n", allocs);
+
+  uint buffs = (uint)buffers_in_pool_;
+  printf("  Buffers in pools: %u\n", buffs);
+
+  printf("  Pools %u:\n", (uint)buffer_pools_.size());
+  auto key_iterator = buffer_pools_.keys().begin();
+  auto value_iterator = buffer_pools_.values().begin();
+  while (key_iterator != buffer_pools_.keys().end()) {
+    uint64_t mem_in_pool = 0;
+    uint64_t iters = 0;
+    for (auto it = (*value_iterator)->begin(); it != (*value_iterator)->end(); it++) {
+      mem_in_pool += it->buffer_size;
+      iters++;
+    }
+
+    printf("    Buffers in pool (%u)(%llu): %u (%u MB)\n",
+           (uint)*key_iterator,
+           iters,
+           (uint)((*value_iterator)->size()),
+           (uint)mem_in_pool / 1024 / 1024);
+    ++key_iterator;
+    ++value_iterator;
+  }
+
+  this->per_frame_allocation_count = 0;
+#endif
+
+  /* Clear safe pools list */
+  completed_safelist_queue_.clear();
+  safelist_lock_.unlock();
+}
+
+void MTLBufferPool::push_completed_safe_list(MTLSafeFreeList *safe_list)
+{
+  /* When an MTLSafeFreeList has been released by the GPU, and buffers are ready to
+   * be re-inserted into the MemoryManager pools for future use, add the MTLSafeFreeList
+   * to the `completed_safelist_queue_` for flushing at a controlled point in time. */
+  safe_list->lock_.lock();
+  BLI_assert(safe_list);
+  BLI_assert(safe_list->reference_count_ == 0 &&
+             "Pool must be fully dereferenced by all in-use cmd buffers before returning.\n");
+  BLI_assert(safe_list->in_free_queue_ == false && "Pool must not already be in queue");
+
+  /* Flag MTLSafeFreeList as having been added, and insert into SafeFreePool queue. */
+  safe_list->flag_in_queue();
+  safelist_lock_.lock();
+  completed_safelist_queue_.append(safe_list);
+  safelist_lock_.unlock();
+  safe_list->lock_.unlock();
+}
+
+MTLSafeFreeList *MTLBufferPool::get_current_safe_list()
+{
+  /* Thread-safe access via atomic ptr. */
+  return current_free_list_;
+}
+
+void MTLBufferPool::begin_new_safe_list()
+{
+  safelist_lock_.lock();
+  current_free_list_ = new MTLSafeFreeList();
+  safelist_lock_.unlock();
+}
+
+void MTLBufferPool::ensure_buffer_pool(MTLResourceOptions options)
+{
+  std::multiset<MTLBufferHandle, CompareMTLBuffer> **pool_search = buffer_pools_.lookup_ptr(
+      (uint64_t)options);
+  if (pool_search == nullptr) {
+    std::multiset<MTLBufferHandle, CompareMTLBuffer> *pool =
+        new std::multiset<MTLBufferHandle, CompareMTLBuffer>();
+    buffer_pools_.add_new((uint64_t)options, pool);
+  }
+}
+
+void MTLBufferPool::insert_buffer_into_pool(MTLResourceOptions options, gpu::MTLBuffer *buffer)
+{
+  /* Ensure `safelist_lock_` is locked in calling code before modifying. */
+  BLI_assert(buffer);
+
+  /* Reset usage size to actual size of allocation. */
+  buffer->set_usage_size(buffer->get_size());
+
+  /* Ensure pool exists. */
+  this->ensure_buffer_pool(options);
+
+  /* TODO(Metal): Support purgability - Allow buffer in pool to have its memory taken back by the
+   * OS if needed. As we keep allocations around, they may not actually be in use, but we can
+   * ensure they do not block other apps from using memory. Upon a buffer being needed again, we
+   * can reset this state.
+   *  TODO(Metal): Purgeability state does not update instantly, so this requires a deferral. */
+  BLI_assert(buffer->get_metal_buffer());
+  /* buffer->metal_buffer); [buffer->metal_buffer setPurgeableState:MTLPurgeableStateVolatile]; */
+
+  std::multiset<MTLBufferHandle, CompareMTLBuffer> *pool = buffer_pools_.lookup(options);
+  pool->insert(MTLBufferHandle(buffer));
+
+#if MTL_DEBUG_MEMORY_STATISTICS == 1
+  /* Debug statistics. */
+  allocations_in_pool_ += buffer->size;
+  buffers_in_pool_++;
+#endif
+}
+
+MTLSafeFreeList::MTLSafeFreeList()
+{
+  reference_count_ = 1;
+  in_free_queue_ = false;
+  current_list_index_ = 0;
+  next_ = nullptr;
+  has_next_pool_ = 0;
+}
+
+void MTLSafeFreeList::insert_buffer(gpu::MTLBuffer *buffer)
+{
+  BLI_assert(in_free_queue_ == false);
+
+  /* Lockless list insert. */
+  uint insert_index = current_list_index_++;
+
+  /* If the current MTLSafeFreeList size is exceeded, we ripple down the linked-list chain and
+   * insert the buffer into the next available chunk. */
+  if (insert_index >= MTLSafeFreeList::MAX_NUM_BUFFERS_) {
+
+    /* Check if first caller to generate next pool. */
+    int has_next = has_next_pool_++;
+    if (has_next == 0) {
+      next_ = new MTLSafeFreeList();
+    }
+    MTLSafeFreeList *next_list = next_.load();
+    BLI_assert(next_list);
+    next_list->insert_buffer(buffer);
+
+    /* Clamp index to chunk limit if overflowing. */
+    current_list_index_ = MTLSafeFreeList::MAX_NUM_BUFFERS_;
+    return;
+  }
+
+  safe_free_pool_[insert_index] = buffer;
+}
+
+/* Increments from active GPUContext thread. */
+void MTLSafeFreeList::increment_reference()
+{
+  lock_.lock();
+  BLI_assert(in_free_queue_ == false);
+  reference_count_++;
+  lock_.unlock();
+}
+
+/* Reference decrements and addition to completed list queue can occur from MTLCommandBuffer
+ * completion callback thread. */
+void MTLSafeFreeList::decrement_reference()
+{
+  lock_.lock();
+  BLI_assert(in_free_queue_ == false);
+  int ref_count = reference_count_--;
+
+  if (ref_count == 0) {
+    MTLContext::get_global_memory_manager().push_completed_safe_list(this);
+  }
+  lock_.unlock();
+}
+
+/** \} */
+
+/* -------------------------------------------------------------------- */
+/** \name MTLBuffer wrapper class implementation.
+ * \{ */
+
+/* Construct a gpu::MTLBuffer wrapper around a newly created metal::MTLBuffer. */
+MTLBuffer::MTLBuffer(id<MTLDevice> mtl_device,
+                     uint64_t size,
+                     MTLResourceOptions options,
+                     uint alignment)
+{
+  /* Calculate aligned allocation size. */
+  BLI_assert(alignment > 0);
+  uint64_t aligned_alloc_size = ceil_to_multiple_ul(size, alignment);
+
+  alignment_ = alignment;
+  device_ = mtl_device;
+  is_external_ = false;
+
+  options_ = options;
+  this->flag_in_use(false);
+
+  metal_buffer_ = [device_ newBufferWithLength:aligned_alloc_size options:options];
+  BLI_assert(metal_buffer_);
+  [metal_buffer_ retain];
+
+  size_ = aligned_alloc_size;
+  this->set_usage_size(size_);
+  if (!(options_ & MTLResourceStorageModePrivate)) {
+    data_ = [metal_buffer_ contents];
+  }
+  else {
+    data_ = nullptr;
+  }
+}
+
+MTLBuffer::MTLBuffer(id<MTLBuffer> external_buffer)
+{
+  BLI_assert(external_buffer != nil);
+
+  /* Ensure external_buffer remains referenced while in-use. */
+  metal_buffer_ = external_buffer;
+  [metal_buffer_ retain];
+
+  /* Extract properties. */
+  is_external_ = true;
+  device_ = nil;
+  alignment_ = 1;
+  options_ = [metal_buffer_ resourceOptions];
+  size_ = [metal_buffer_ allocatedSize];
+  this->set_usage_size(size_);
+  data_ = [metal_buffer_ contents];
+  in_use_ = true;
+}
+
+gpu::MTLBuffer::~MTLBuffer()
+{
+  if (metal_buffer_ != nil) {
+    [metal_buffer_ release];
+    metal_buffer_ = nil;
+  }
+}
+
+void gpu::MTLBuffer::free()
+{
+  if (!is_external_) {
+    MTLContext::get_global_memory_manager().free_buffer(this);
+  }
+  else {
+    if (metal_buffer_ != nil) {
+      [metal_buffer_ release];
+      metal_buffer_ = nil;
+    }
+  }
+}
+
+id<MTLBuffer> gpu::MTLBuffer::get_metal_buffer() const
+{
+  return metal_buffer_;
+}
+
+void *gpu::MTLBuffer::get_host_ptr() const
+{
+  BLI_assert(!(options_ & MTLResourceStorageModePrivate));
+  BLI_assert(data_);
+  return data_;
+}
+
+uint64_t gpu::MTLBuffer::get_size() const
+{
+  return size_;
+}
+
+uint64_t gpu::MTLBuffer::get_size_used() const
+{
+  return usage_size_;
+}
+
+bool gpu::MTLBuffer::requires_flush()
+{
+  /* We do not need to flush shared memory, as addressable buffer is shared. */
+  return options_ & MTLResourceStorageModeManaged;
+}
+
+void gpu::MTLBuffer::set_label(NSString *str)
+{
+  metal_buffer_.label = str;
+}
+
+void gpu::MTLBuffer::debug_ensure_used()
+{
+  /* Debug: If buffer is not flagged as in-use, this is a problem. */
+  BLI_assert(in_use_ &&
+             "Buffer should be marked as 'in-use' if being actively used by an instance. Buffer "
+             "has likely already been freed.");
+}
+
+void gpu::MTLBuffer::flush()
+{
+  this->debug_ensure_used();
+  if (this->requires_flush()) {
+    [metal_buffer_ didModifyRange:NSMakeRange(0, size_)];
+  }
+}
+
+void gpu::MTLBuffer::flush_range(uint64_t offset, uint64_t length)
+{
+  this->debug_ensure_used();
+  if (this->requires_flush()) {
+    BLI_assert((offset + length) <= size_);
+    [metal_buffer_ didModifyRange:NSMakeRange(offset, length)];
+  }
+}
+
+void gpu::MTLBuffer::flag_in_use(bool used)
+{
+  in_use_ = used;
+}
+
+bool gpu::MTLBuffer::get_in_use()
+{
+  return in_use_;
+}
+
+void gpu::MTLBuffer::set_usage_size(uint64_t size_used)
+{
+  BLI_assert(size_used > 0 && size_used <= size_);
+  usage_size_ = size_used;
+}
+
+MTLResourceOptions gpu::MTLBuffer::get_resource_options()
+{
+  return options_;
+}
+
+uint64_t gpu::MTLBuffer::get_alignment()
+{
+  return alignment_;
+}
+
+bool MTLBufferRange::requires_flush()
+{
+  /* We do not need to flush shared memory. */
+  return this->options & MTLResourceStorageModeManaged;
+}
+
+void MTLBufferRange::flush()
+{
+  if (this->requires_flush()) {
+    BLI_assert(this->metal_buffer);
+    BLI_assert((this->buffer_offset + this->size) <= [this->metal_buffer length]);
+    BLI_assert(this->buffer_offset >= 0);
+    [this->metal_buffer
+        didModifyRange:NSMakeRange(this->buffer_offset, this->size - this->buffer_offset)];
+  }
+}
+
+/** \} */
+
+/* -------------------------------------------------------------------- */
+/** \name MTLScratchBufferManager and MTLCircularBuffer implementation.
+ * \{ */
+
+MTLScratchBufferManager::~MTLScratchBufferManager()
+{
+  this->free();
+}
+
+void MTLScratchBufferManager::init()
+{
+
+  if (!this->initialised_) {
+    BLI_assert(context_.device);
+
+    /* Initialise Scratch buffers */
+    for (int sb = 0; sb < mtl_max_scratch_buffers_; sb++) {
+      scratch_buffers_[sb] = new MTLCircularBuffer(
+          context_, mtl_scratch_buffer_initial_size_, true);
+      BLI_assert(scratch_buffers_[sb]);
+      BLI_assert(&(scratch_buffers_[sb]->own_context_) == &context_);
+    }
+    current_scratch_buffer_ = 0;
+    initialised_ = true;
+  }
+}
+
+void MTLScratchBufferManager::free()
+{
+  initialised_ = false;
+
+  /* Release Scratch buffers */
+  for (int sb = 0; sb < mtl_max_scratch_buffers_; sb++) {
+    delete scratch_buffers_[sb];
+    scratch_buffers_[sb] = nullptr;
+  }
+  current_scratch_buffer_ = 0;
+}
+
+MTLTemporaryBuffer MTLScratchBufferManager::scratch_buffer_allocate_range(uint64_t alloc_size)
+{
+  return this->scratch_buffer_allocate_range_aligned(alloc_size, 1);
+}
+
+MTLTemporaryBuffer MTLScratchBufferManager::scratch_buffer_allocate_range_aligned(
+    uint64_t alloc_size, uint alignment)
+{
+  /* Ensure scratch buffer allocation alignment adheres to offset alignment requirements. */
+  alignment = max_uu(alignment, 256);
+
+  BLI_assert(current_scratch_buffer_ >= 0 && "Scratch Buffer index not set");
+  MTLCircularBuffer *current_scratch_buff = this->scratch_buffers_[current_scratch_buffer_];
+  BLI_assert(current_scratch_buff != nullptr && "Scratch Buffer does not exist");
+  MTLTemporaryBuffer allocated_range = current_scratch_buff->allocate_range_aligned(alloc_size,
+                                                                                    alignment);
+  BLI_assert(allocated_range.size >= alloc_size && allocated_range.size <= alloc_size + alignment);
+  BLI_assert(allocated_range.metal_buffer != nil);
+  return allocated_range;
+}
+
+void MTLScratchBufferManager::ensure_increment_scratch_buffer()
+{
+  /* Fetch active scratch buffer. */
+  MTLCircularBuffer *active_scratch_buf = scratch_buffers_[current_scratch_buffer_];
+  BLI_assert(&active_scratch_buf->own_context_ == &context_);
+
+  /* Ensure existing scratch buffer is no longer in use. MTL_MAX_SCRATCH_BUFFERS specifies
+   * the number of allocated scratch buffers. This value should be equal to the number of
+   * simultaneous frames in-flight. I.e. the maximal number of scratch buffers which are
+   * simultaneously in-use. */
+  if (active_scratch_buf->used_frame_index_ < context_.get_current_frame_index()) {
+    current_scratch_buffer_ = (current_scratch_buffer_ + 1) % mtl_max_scratch_buffers_;
+    active_scratch_buf = scratch_buffers_[current_scratch_buffer_];
+    active_scratch_buf->reset();
+    BLI_assert(&active_scratch_buf->own_context_ == &context_);
+    MTL_LOG_INFO("Scratch buffer %d reset - (ctx %p)(Frame index: %d)\n",
+                 current_scratch_buffer_,
+                 &context_,
+                 context_.get_current_frame_index());
+  }
+}
+
+void MTLScratchBufferManager::flush_active_scratch_buffer()
+{
+  /* Fetch active scratch buffer and verify context. */
+  MTLCircularBuffer *active_scratch_buf = scratch_buffers_[current_scratch_buffer_];
+  BLI_assert(&active_scratch_buf->own_context_ == &context_);
+  active_scratch_buf->flush();
+}
+
+/* MTLCircularBuffer implementation. */
+MTLCircularBuffer::MTLCircularBuffer(MTLContext &ctx, uint64_t initial_size, bool allow_grow)
+    : own_context_(ctx)
+{
+  BLI_assert(this);
+  MTLResourceOptions options = ([own_context_.device hasUnifiedMemory]) ?
+                                   MTLResourceStorageModeShared :
+                                   MTLResourceStorageModeManaged;
+  cbuffer_ = new gpu::MTLBuffer(own_context_.device, initial_size, options, 256);
+  current_offset_ = 0;
+  can_resize_ = allow_grow;
+  cbuffer_->flag_in_use(true);
+
+  used_frame_index_ = ctx.get_current_frame_index();
+  last_flush_base_offset_ = 0;
+
+  /* Debug label. */
+  if (G.debug & G_DEBUG_GPU) {
+    cbuffer_->set_label(@"Circular Scratch Buffer");
+  }
+}
+
+MTLCircularBuffer::~MTLCircularBuffer()
+{
+  delete cbuffer_;
+}
+
+MTLTemporaryBuffer MTLCircularBuffer::allocate_range(uint64_t alloc_size)
+{
+  return this->allocate_range_aligned(alloc_size, 1);
+}
+
+MTLTemporaryBuffer MTLCircularBuffer::allocate_range_aligned(uint64_t alloc_size, uint alignment)
+{
+  BLI_assert(this);
+
+  /* Ensure alignment of an allocation is aligned to compatible offset boundaries. */
+  BLI_assert(alignment > 0);
+  alignment = max_ulul(alignment, 256);
+
+  /* Align current offset and allocation size to desired alignment */
+  uint64_t aligned_current_offset = ceil_to_multiple_ul(current_offset_, alignment);
+  uint64_t aligned_alloc_size = ceil_to_multiple_ul(alloc_size, alignment);
+  bool can_allocate = (aligned_current_offset + aligned_alloc_size) < cbuffer_->get_size();
+
+  BLI_assert(aligned_current_offset >= current_offset_);
+  BLI_assert(aligned_alloc_size >= alloc_size);
+
+  BLI_assert(aligned_current_offset % alignment == 0);
+  BLI_assert(aligned_alloc_size % alignment == 0);
+
+  /* Recreate Buffer */
+  if (!can_allocate) {
+    uint64_t new_size = cbuffer_->get_size();
+    if (can_resize_) {
+      /* Resize to the maximum of basic resize heuristic OR the size of the current offset +
+       * requested allocation -- we want the buffer to grow to a large enough size such that it
+       * does not need to resize mid-frame. */
+      new_size = max_ulul(
+          min_ulul(MTLScratchBufferManager::mtl_scratch_buffer_max_size_, new_size * 1.2),
+          aligned_current_offset + aligned_alloc_size);
+
+#if MTL_SCRATCH_BUFFER_ALLOW_TEMPORARY_EXPANSION == 1
+      /* IF a requested allocation EXCEEDS the maximum supported size, temporarily allocate up to
+       * this, but shrink down ASAP. */
+      if (new_size > MTLScratchBufferManager::mtl_scratch_buffer_max_size_) {
+
+        /* If new requested allocation is bigger than maximum allowed size, temporarily resize to
+         * maximum allocation size -- Otherwise, clamp the buffer size back down to the defined
+         * maximum */
+        if (aligned_alloc_size > MTLScratchBufferManager::mtl_scratch_buffer_max_size_) {
+          new_size = aligned_alloc_size;
+          MTL_LOG_INFO("Temporarily growing Scratch buffer to %d MB\n",
+                       (int)new_size / 1024 / 1024);
+        }
+        else {
+          new_size = MTLScratchBufferManager::mtl_scratch_buffer_max_size_;
+          MTL_LOG_INFO("Shrinking Scratch buffer back to %d MB\n", (int)new_size / 1024 / 1024);
+        }
+      }
+      BLI_assert(aligned_alloc_size <= new_size);
+#else
+      new_size = min_ulul(MTLScratchBufferManager::mtl_scratch_buffer_max_size_, new_size);
+
+      if (aligned_alloc_size > new_size) {
+        BLI_assert(false);
+
+        /* Cannot allocate */
+        MTLTemporaryBuffer alloc_range;
+        alloc_range.metal_buffer = nil;
+        alloc_range.data = nullptr;
+        alloc_range.buffer_offset = 0;
+        alloc_range.size = 0;
+        alloc_range.options = cbuffer_->options;
+      }
+#endif
+    }
+    else {
+      MTL_LOG_WARNING(
+          "Performance Warning: Reached the end of circular buffer of size: %llu, but cannot "
+          "resize. Starting new buffer\n",
+          cbuffer_->get_size());
+      BLI_assert(aligned_alloc_size <= new_size);
+
+      /* Cannot allocate. */
+      MTLTemporaryBuffer alloc_range;
+      alloc_range.metal_buffer = nil;
+      alloc_range.data = nullptr;
+      alloc_range.buffer_offset = 0;
+      alloc_range.size = 0;
+      alloc_range.options = cbuffer_->get_resource_options();
+    }
+
+    /* Flush current buffer to ensure changes are visible on the GPU. */
+    this->flush();
+
+    /* Discard old buffer and create a new one - Relying on Metal reference counting to track
+     * in-use buffers */
+    MTLResourceOptions prev_options = cbuffer_->get_resource_options();
+    uint prev_alignment = cbuffer_->get_alignment();
+    delete cbuffer_;
+    cbuffer_ = new gpu::MTLBuffer(own_context_.device, new_size, prev_options, prev_alignment);
+    cbuffer_->flag_in_use(true);
+    current_offset_ = 0;
+    last_flush_base_offset_ = 0;
+
+    /* Debug label. */
+    if (G.debug & G_DEBUG_GPU) {
+      cbuffer_->set_label(@"Circular Scratch Buffer");
+    }
+    MTL_LOG_INFO("Resized Metal circular buffer to %llu bytes\n", new_size);
+
+    /* Reset allocation Status. */
+    aligned_current_offset = 0;
+    BLI_assert((aligned_current_offset + aligned_alloc_size) <= cbuffer_->get_size());
+  }
+
+  /* Allocate chunk. */
+  MTLTemporaryBuffer alloc_range;
+  alloc_range.metal_buffer = cbuffer_->get_metal_buffer();
+  alloc_range.data = (void *)((uint8_t *)([alloc_range.metal_buffer contents]) +
+                              aligned_current_offset);
+  alloc_range.buffer_offset = aligned_current_offset;
+  alloc_range.size = aligned_alloc_size;
+  alloc_range.options = cbuffer_->get_resource_options();
+  BLI_assert(alloc_range.data);
+
+  /* Shift offset to match alignment. */
+  current_offset_ = aligned_current_offset + aligned_alloc_size;
+  BLI_assert(current_offset_ <= cbuffer_->get_size());
+  return alloc_range;
+}
+
+void MTLCircularBuffer::flush()
+{
+  BLI_assert(this);
+
+  uint64_t len = current_offset_ - last_flush_base_offset_;
+  if (len > 0) {
+    cbuffer_->flush_range(last_flush_base_offset_, len);
+    last_flush_base_offset_ = current_offset_;
+  }
+}
+
+void MTLCircularBuffer::reset()
+{
+  BLI_assert(this);
+
+  /* If circular buffer has data written to it, offset will be greater than zero. */
+  if (current_offset_ > 0) {
+
+    /* Ensure the circular buffer is no longer being used by an in-flight frame. */
+    BLI_assert((own_context_.get_current_frame_index() >=
+                (used_frame_index_ + MTL_NUM_SAFE_FRAMES - 1)) &&
+               "Trying to reset Circular scratch buffer's while its data is still being used by "
+               "an in-flight frame");
+
+    current_offset_ = 0;
+    last_flush_base_offset_ = 0;
+  }
+
+  /* Update used frame index to current. */
+  used_frame_index_ = own_context_.get_current_frame_index();
+}
+
+/** \} */
+
+}  // blender::gpu
diff --git a/source/blender/gpu/metal/mtl_state.hh b/source/blender/gpu/metal/mtl_state.hh
index 23bf8600ddd..ddb27a444d4 100644
--- a/source/blender/gpu/metal/mtl_state.hh
+++ b/source/blender/gpu/metal/mtl_state.hh
@@ -30,18 +30,18 @@ class MTLStateManager : public StateManager {
  public:
   MTLStateManager(MTLContext *ctx);
 
-  void apply_state(void) override;
-  void force_state(void) override;
+  void apply_state() override;
+  void force_state() override;
 
   void issue_barrier(eGPUBarrier barrier_bits) override;
 
   void texture_bind(Texture *tex, eGPUSamplerState sampler, int unit) override;
   void texture_unbind(Texture *tex) override;
-  void texture_unbind_all(void) override;
+  void texture_unbind_all() override;
 
   void image_bind(Texture *tex, int unit) override;
   void image_unbind(Texture *tex) override;
-  void image_unbind_all(void) override;
+  void image_unbind_all() override;
 
   void texture_unpack_row_length_set(uint len) override;
 
diff --git a/source/blender/gpu/metal/mtl_state.mm b/source/blender/gpu/metal/mtl_state.mm
index cf7fbdba6b9..59c258a0d12 100644
--- a/source/blender/gpu/metal/mtl_state.mm
+++ b/source/blender/gpu/metal/mtl_state.mm
@@ -17,7 +17,7 @@ namespace blender::gpu {
 /** \name MTLStateManager
  * \{ */
 
-void MTLStateManager::mtl_state_init(void)
+void MTLStateManager::mtl_state_init()
 {
   BLI_assert(context_);
   context_->pipeline_state_init();
@@ -36,7 +36,7 @@ MTLStateManager::MTLStateManager(MTLContext *ctx) : StateManager()
   set_mutable_state(mutable_state);
 }
 
-void MTLStateManager::apply_state(void)
+void MTLStateManager::apply_state()
 {
   this->set_state(this->state);
   this->set_mutable_state(this->mutable_state);
@@ -45,7 +45,7 @@ void MTLStateManager::apply_state(void)
   static_cast<MTLFrameBuffer *>(context_->active_fb)->apply_state();
 };
 
-void MTLStateManager::force_state(void)
+void MTLStateManager::force_state()
 {
   /* Little exception for clip distances since they need to keep the old count correct. */
   uint32_t clip_distances = current_.clip_distances;
@@ -548,7 +548,7 @@ void MTLStateManager::issue_barrier(eGPUBarrier barrier_bits)
 
   /* Apple Silicon does not support memory barriers.
    * We do not currently need these due to implicit API guarantees.
-   * Note(Metal): MTLFence/MTLEvent may be required to synchronize work if
+   * NOTE(Metal): MTLFence/MTLEvent may be required to synchronize work if
    * untracked resources are ever used. */
   if ([ctx->device hasUnifiedMemory]) {
     return;
@@ -600,7 +600,7 @@ void MTLStateManager::texture_unbind(Texture *tex_)
   ctx->texture_unbind(mtl_tex);
 }
 
-void MTLStateManager::texture_unbind_all(void)
+void MTLStateManager::texture_unbind_all()
 {
   MTLContext *ctx = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
   BLI_assert(ctx);
@@ -623,7 +623,7 @@ void MTLStateManager::image_unbind(Texture *tex_)
   this->texture_unbind(tex_);
 }
 
-void MTLStateManager::image_unbind_all(void)
+void MTLStateManager::image_unbind_all()
 {
   this->texture_unbind_all();
 }
diff --git a/source/blender/gpu/metal/mtl_texture.hh b/source/blender/gpu/metal/mtl_texture.hh
index 0f908995a93..9387d5af814 100644
--- a/source/blender/gpu/metal/mtl_texture.hh
+++ b/source/blender/gpu/metal/mtl_texture.hh
@@ -237,7 +237,7 @@ class MTLTexture : public Texture {
   void update_sub(
       int mip, int offset[3], int extent[3], eGPUDataFormat type, const void *data) override;
 
-  void generate_mipmap(void) override;
+  void generate_mipmap() override;
   void copy_to(Texture *dst) override;
   void clear(eGPUDataFormat format, const void *data) override;
   void swizzle_set(const char swizzle_mask[4]) override;
@@ -248,7 +248,7 @@ class MTLTexture : public Texture {
   void *read(int mip, eGPUDataFormat type) override;
 
   /* Remove once no longer required -- will just return 0 for now in MTL path*/
-  uint gl_bindcode_get(void) const override;
+  uint gl_bindcode_get() const override;
 
   bool texture_is_baked();
   const char *get_name()
@@ -257,7 +257,7 @@ class MTLTexture : public Texture {
   }
 
  protected:
-  bool init_internal(void) override;
+  bool init_internal() override;
   bool init_internal(GPUVertBuf *vbo) override;
   bool init_internal(const GPUTexture *src,
                      int mip_offset,
diff --git a/source/blender/gpu/metal/mtl_texture.mm b/source/blender/gpu/metal/mtl_texture.mm
index ff2c2fce235..0cb38a3a2b7 100644
--- a/source/blender/gpu/metal/mtl_texture.mm
+++ b/source/blender/gpu/metal/mtl_texture.mm
@@ -478,23 +478,6 @@ void gpu::MTLTexture::update_sub(
     MTLPixelFormat destination_format = gpu_texture_format_to_metal(format_);
     int expected_dst_bytes_per_pixel = get_mtl_format_bytesize(destination_format);
     int destination_num_channels = get_mtl_format_num_components(destination_format);
-    int destination_totalsize = 0;
-    switch (this->dimensions_count()) {
-      case 1:
-        destination_totalsize = expected_dst_bytes_per_pixel * max_ii(expected_update_w, 1);
-        break;
-      case 2:
-        destination_totalsize = expected_dst_bytes_per_pixel * max_ii(expected_update_w, 1) *
-                                max_ii(extent[1], 1);
-        break;
-      case 3:
-        destination_totalsize = expected_dst_bytes_per_pixel * max_ii(expected_update_w, 1) *
-                                max_ii(extent[1], 1) * max_ii(extent[2], 1);
-        break;
-      default:
-        BLI_assert(false);
-        break;
-    }
 
     /* Prepare specialisation struct (For texture update routine). */
     TextureUpdateRoutineSpecialisation compute_specialisation_kernel = {
@@ -568,12 +551,12 @@ void gpu::MTLTexture::update_sub(
 
     /* Prepare staging buffer for data. */
     id<MTLBuffer> staging_buffer = nil;
-    unsigned long long staging_buffer_offset = 0;
+    uint64_t staging_buffer_offset = 0;
 
     /* Fetch allocation from scratch buffer. */
-    MTLTemporaryBufferRange allocation; /* TODO(Metal): Metal Memory manager. */
-    /* = ctx->get_memory_manager().scratch_buffer_allocate_range_aligned(totalsize, 256);*/
-    memcpy(allocation.host_ptr, data, totalsize);
+    MTLTemporaryBuffer allocation =
+        ctx->get_scratchbuffer_manager().scratch_buffer_allocate_range_aligned(totalsize, 256);
+    memcpy(allocation.data, data, totalsize);
     staging_buffer = allocation.metal_buffer;
     staging_buffer_offset = allocation.buffer_offset;
 
@@ -915,7 +898,7 @@ void gpu::MTLTexture::ensure_mipmaps(int miplvl)
   this->mip_range_set(0, mipmaps_);
 }
 
-void gpu::MTLTexture::generate_mipmap(void)
+void gpu::MTLTexture::generate_mipmap()
 {
   /* Fetch Active Context. */
   MTLContext *ctx = reinterpret_cast<MTLContext *>(GPU_context_active_get());
@@ -1230,7 +1213,7 @@ void gpu::MTLTexture::read_internal(int mip,
   destination_buffer = [ctx->device newBufferWithLength:max_ii(total_bytes, 256)
                                                 options:bufferOptions];
   destination_offset = 0;
-  destination_buffer_host_ptr = (void *)((unsigned char *)([destination_buffer contents]) +
+  destination_buffer_host_ptr = (void *)((uint8_t *)([destination_buffer contents]) +
                                          destination_offset);
 
   /* Prepare specialisation struct (For non-trivial texture read routine). */
@@ -1444,12 +1427,12 @@ void gpu::MTLTexture::read_internal(int mip,
 }
 
 /* Remove once no longer required -- will just return 0 for now in MTL path. */
-uint gpu::MTLTexture::gl_bindcode_get(void) const
+uint gpu::MTLTexture::gl_bindcode_get() const
 {
   return 0;
 }
 
-bool gpu::MTLTexture::init_internal(void)
+bool gpu::MTLTexture::init_internal()
 {
   if (format_ == GPU_DEPTH24_STENCIL8) {
     /* Apple Silicon requires GPU_DEPTH32F_STENCIL8 instead of GPU_DEPTH24_STENCIL8. */