blender
/
blender
129
398
Fork 572
Code Issues 6.2k Pull Requests 746 Projects 19 Wiki Activity

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
This commit is contained in:
Jason Fielder 2022-07-01 10:30:16 +02:00 committed by Clément Foucault
parent 3ffc558341
commit 4527dd1ce4
Notes: blender-bot 2023-02-14 07:30:31 +01:00 
Referenced by issue #96261, Metal Viewport
18 changed files with 1524 additions and 124 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
source/blender/blenlib/BLI_math_base.h
View File

@ -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.
*/

10
source/blender/blenlib/intern/math_base_inline.c
View File

@ -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;

5
source/blender/draw/engines/eevee/eevee_render.c
View File

@ -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);
}
}

5
source/blender/draw/engines/workbench/workbench_render.c
View File

@ -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);

2
source/blender/gpu/CMakeLists.txt
View File

@ -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
)

8
source/blender/gpu/intern/gpu_immediate_util.c
View File

@ -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);

16
source/blender/gpu/metal/mtl_backend.mm
View File

@ -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()

60
source/blender/gpu/metal/mtl_command_buffer.mm
View File

@ -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;

5
source/blender/gpu/metal/mtl_common.hh
View File

@ -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

57
source/blender/gpu/metal/mtl_context.hh
View File

@ -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

50
source/blender/gpu/metal/mtl_context.mm
View File

@ -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();
}

2
source/blender/gpu/metal/mtl_framebuffer.mm
View File

@ -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);

476
source/blender/gpu/metal/mtl_memory.hh Normal file
View File

@ -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

880
source/blender/gpu/metal/mtl_memory.mm Normal file
View File

@ -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

8
source/blender/gpu/metal/mtl_state.hh
View File

@ -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;

12
source/blender/gpu/metal/mtl_state.mm
View File

@ -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();
}

6
source/blender/gpu/metal/mtl_texture.hh
View File

@ -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,

33
source/blender/gpu/metal/mtl_texture.mm
View File

@ -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. */