From 4527dd1ce4784292cd3b8dd3764b9cd843020f9a Mon Sep 17 00:00:00 2001 From: Jason Fielder Date: Fri, 1 Jul 2022 10:30:16 +0200 Subject: 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 --- source/blender/gpu/metal/mtl_backend.mm | 16 +- source/blender/gpu/metal/mtl_command_buffer.mm | 60 +- source/blender/gpu/metal/mtl_common.hh | 5 + source/blender/gpu/metal/mtl_context.hh | 57 +- source/blender/gpu/metal/mtl_context.mm | 50 +- source/blender/gpu/metal/mtl_framebuffer.mm | 2 +- source/blender/gpu/metal/mtl_memory.hh | 476 +++++++++++++ source/blender/gpu/metal/mtl_memory.mm | 880 +++++++++++++++++++++++++ source/blender/gpu/metal/mtl_state.hh | 8 +- source/blender/gpu/metal/mtl_state.mm | 12 +- source/blender/gpu/metal/mtl_texture.hh | 6 +- source/blender/gpu/metal/mtl_texture.mm | 33 +- 12 files changed, 1485 insertions(+), 120 deletions(-) create mode 100644 source/blender/gpu/metal/mtl_memory.hh create mode 100644 source/blender/gpu/metal/mtl_memory.mm (limited to 'source/blender/gpu/metal') 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 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 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 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 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 cmd_buffer_ref = active_command_buffer_; + [cmd_buffer_ref retain]; + [cmd_buffer_ref addCompletedHandler:^(id 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 tex, uint slot) { if (this->cached_vertex_texture_bindings[slot].metal_texture != tex) { - id rec = this->cmd->get_active_render_command_encoder(); + id 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 tex, uint slot) void MTLRenderPassState::bind_fragment_texture(id tex, uint slot) { if (this->cached_fragment_texture_bindings[slot].metal_texture != tex) { - id rec = this->cmd->get_active_render_command_encoder(); + id 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 @@ -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 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 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 sampler_state_cache_[GPU_SAMPLER_MAX] = {0}; + id sampler_state_cache_[GPU_SAMPLER_MAX]; id 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 queue = nil; id 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 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(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 +#include +#include +#include +#include +#include + +#include "mtl_common.hh" + +#include +#include +#include + +@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 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 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 in_use_; + + public: + MTLBuffer(id device, uint64_t size, MTLResourceOptions options, uint alignment = 1); + MTLBuffer(id external_buffer); + ~MTLBuffer(); + + /* Fetch information about backing MTLBuffer. */ + id 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 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 reference_count_; + std::atomic in_free_queue_; + std::recursive_mutex lock_; + + /* Linked list of next MTLSafeFreeList chunk if current chunk is full. */ + std::atomic has_next_pool_; + std::atomic 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 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 per_frame_allocation_count_; + std::atomic allocations_in_pool_; + std::atomic buffers_in_pool_; +#endif + + /* Metal resources. */ + bool ensure_initialised_ = false; + id 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; + using MTLBufferResourceOptions = uint64_t; + + blender::Map buffer_pools_; + blender::Vector 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 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 current_free_list_; + + public: + void init(id 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 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 *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 **pool_search = buffer_pools_.lookup_ptr( + (uint64_t)options); + + if (pool_search != nullptr) { + std::multiset *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 **pool_search = buffer_pools_.lookup_ptr( + (uint64_t)options); + if (pool_search == nullptr) { + std::multiset *pool = + new std::multiset(); + 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 *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 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 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 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(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(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 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(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. */ -- cgit v1.2.3