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:
@@ -221,6 +221,19 @@ MINLINE unsigned int power_of_2_min_u(unsigned int x);
|
|||||||
* with integers, to avoid gradual darkening when rounding down.
|
* with integers, to avoid gradual darkening when rounding down.
|
||||||
*/
|
*/
|
||||||
MINLINE int divide_round_i(int a, int b);
|
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.
|
* modulo that handles negative numbers, works the same as Python's.
|
||||||
*/
|
*/
|
||||||
|
@@ -370,6 +370,11 @@ MINLINE uint divide_ceil_u(uint a, uint b)
|
|||||||
return (a + b - 1) / 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 .
|
* 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;
|
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)
|
MINLINE int mod_i(int i, int n)
|
||||||
{
|
{
|
||||||
return (i % n + n) % n;
|
return (i % n + n) % n;
|
||||||
|
@@ -24,6 +24,7 @@
|
|||||||
#include "DEG_depsgraph_query.h"
|
#include "DEG_depsgraph_query.h"
|
||||||
|
|
||||||
#include "GPU_capabilities.h"
|
#include "GPU_capabilities.h"
|
||||||
|
#include "GPU_context.h"
|
||||||
#include "GPU_framebuffer.h"
|
#include "GPU_framebuffer.h"
|
||||||
#include "GPU_state.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. */
|
/* XXX Seems to fix TDR issue with NVidia drivers on linux. */
|
||||||
GPU_finish();
|
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);
|
RE_engine_update_progress(engine, (float)(render_samples++) / (float)tot_sample);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
@@ -17,6 +17,7 @@
|
|||||||
|
|
||||||
#include "ED_view3d.h"
|
#include "ED_view3d.h"
|
||||||
|
|
||||||
|
#include "GPU_context.h"
|
||||||
#include "GPU_shader.h"
|
#include "GPU_shader.h"
|
||||||
|
|
||||||
#include "DEG_depsgraph.h"
|
#include "DEG_depsgraph.h"
|
||||||
@@ -188,6 +189,10 @@ void workbench_render(void *ved, RenderEngine *engine, RenderLayer *render_layer
|
|||||||
|
|
||||||
workbench_draw_finish(data);
|
workbench_draw_finish(data);
|
||||||
|
|
||||||
|
/* Perform render step between samples to allow
|
||||||
|
* flushing of freed GPUBackend resources. */
|
||||||
|
GPU_render_step();
|
||||||
|
|
||||||
/* Write render output. */
|
/* Write render output. */
|
||||||
const char *viewname = RE_GetActiveRenderView(engine->re);
|
const char *viewname = RE_GetActiveRenderView(engine->re);
|
||||||
RenderPass *rp = RE_pass_find_by_name(render_layer, RE_PASSNAME_COMBINED, viewname);
|
RenderPass *rp = RE_pass_find_by_name(render_layer, RE_PASSNAME_COMBINED, viewname);
|
||||||
|
@@ -194,6 +194,7 @@ set(METAL_SRC
|
|||||||
metal/mtl_command_buffer.mm
|
metal/mtl_command_buffer.mm
|
||||||
metal/mtl_debug.mm
|
metal/mtl_debug.mm
|
||||||
metal/mtl_framebuffer.mm
|
metal/mtl_framebuffer.mm
|
||||||
|
metal/mtl_memory.mm
|
||||||
metal/mtl_state.mm
|
metal/mtl_state.mm
|
||||||
metal/mtl_texture.mm
|
metal/mtl_texture.mm
|
||||||
metal/mtl_texture_util.mm
|
metal/mtl_texture_util.mm
|
||||||
@@ -204,6 +205,7 @@ set(METAL_SRC
|
|||||||
metal/mtl_context.hh
|
metal/mtl_context.hh
|
||||||
metal/mtl_debug.hh
|
metal/mtl_debug.hh
|
||||||
metal/mtl_framebuffer.hh
|
metal/mtl_framebuffer.hh
|
||||||
|
metal/mtl_memory.hh
|
||||||
metal/mtl_state.hh
|
metal/mtl_state.hh
|
||||||
metal/mtl_texture.hh
|
metal/mtl_texture.hh
|
||||||
)
|
)
|
||||||
|
@@ -142,7 +142,7 @@ static void imm_draw_circle(GPUPrimType prim_type,
|
|||||||
int nsegments)
|
int nsegments)
|
||||||
{
|
{
|
||||||
if (prim_type == GPU_PRIM_LINE_LOOP) {
|
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);
|
immBegin(GPU_PRIM_LINES, nsegments * 2);
|
||||||
|
|
||||||
immVertex2f(shdr_pos, x + (radius_x * cosf(0.0f)), y + (radius_y * sinf(0.0f)));
|
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)
|
GPUPrimType prim_type, uint pos, float x, float y, float radius, int nsegments)
|
||||||
{
|
{
|
||||||
if (prim_type == GPU_PRIM_LINE_LOOP) {
|
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);
|
immBegin(GPU_PRIM_LINES, nsegments * 2);
|
||||||
|
|
||||||
const float angle = (float)(2 * M_PI) / (float)nsegments;
|
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)
|
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);
|
immBegin(GPU_PRIM_LINES, 8);
|
||||||
immVertex2f(pos, x1, y1);
|
immVertex2f(pos, x1, y1);
|
||||||
immVertex2f(pos, x1, y2);
|
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)
|
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 */
|
/* 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);
|
immBegin(GPU_PRIM_LINES, 8);
|
||||||
immVertex3f(pos, x1, y1, 0.0f);
|
immVertex3f(pos, x1, y1, 0.0f);
|
||||||
immVertex3f(pos, x1, y2, 0.0f);
|
immVertex3f(pos, x1, y2, 0.0f);
|
||||||
|
@@ -127,7 +127,21 @@ void MTLBackend::render_end()
|
|||||||
|
|
||||||
void MTLBackend::render_step()
|
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()
|
bool MTLBackend::is_inside_render_boundary()
|
||||||
|
@@ -19,7 +19,7 @@ namespace blender::gpu {
|
|||||||
* dependencies not being honored for work submitted between
|
* dependencies not being honored for work submitted between
|
||||||
* different GPUContext's. */
|
* different GPUContext's. */
|
||||||
id<MTLEvent> MTLCommandBufferManager::sync_event = nil;
|
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. */
|
/* Counter for active command buffers. */
|
||||||
int MTLCommandBufferManager::num_active_cmd_bufs = 0;
|
int MTLCommandBufferManager::num_active_cmd_bufs = 0;
|
||||||
@@ -28,10 +28,9 @@ int MTLCommandBufferManager::num_active_cmd_bufs = 0;
|
|||||||
/** \name MTLCommandBuffer initialization and render coordination.
|
/** \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_.reset_state();
|
||||||
render_pass_state_.prepare(this, ctx);
|
|
||||||
}
|
}
|
||||||
|
|
||||||
void MTLCommandBufferManager::register_encoder_counters()
|
void MTLCommandBufferManager::register_encoder_counters()
|
||||||
@@ -54,10 +53,10 @@ id<MTLCommandBuffer> MTLCommandBufferManager::ensure_begin()
|
|||||||
MTLCommandBufferDescriptor *desc = [[MTLCommandBufferDescriptor alloc] init];
|
MTLCommandBufferDescriptor *desc = [[MTLCommandBufferDescriptor alloc] init];
|
||||||
desc.errorOptions = MTLCommandBufferErrorOptionEncoderExecutionStatus;
|
desc.errorOptions = MTLCommandBufferErrorOptionEncoderExecutionStatus;
|
||||||
desc.retainedReferences = YES;
|
desc.retainedReferences = YES;
|
||||||
active_command_buffer_ = [context_->queue commandBufferWithDescriptor:desc];
|
active_command_buffer_ = [context_.queue commandBufferWithDescriptor:desc];
|
||||||
}
|
}
|
||||||
else {
|
else {
|
||||||
active_command_buffer_ = [context_->queue commandBuffer];
|
active_command_buffer_ = [context_.queue commandBuffer];
|
||||||
}
|
}
|
||||||
[active_command_buffer_ retain];
|
[active_command_buffer_ retain];
|
||||||
MTLCommandBufferManager::num_active_cmd_bufs++;
|
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];
|
[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. */
|
/* Reset Command buffer heuristics. */
|
||||||
this->reset_counters();
|
this->reset_counters();
|
||||||
}
|
}
|
||||||
@@ -86,12 +89,15 @@ bool MTLCommandBufferManager::submit(bool wait)
|
|||||||
this->end_active_command_encoder();
|
this->end_active_command_encoder();
|
||||||
BLI_assert(active_command_encoder_type_ == MTL_NO_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. ***/
|
/*** Submit Command Buffer. ***/
|
||||||
/* Strict ordering ensures command buffers are guaranteed to execute after a previous
|
/* Strict ordering ensures command buffers are guaranteed to execute after a previous
|
||||||
* one has completed. Resolves flickering when command buffers are submitted from
|
* one has completed. Resolves flickering when command buffers are submitted from
|
||||||
* different MTLContext's. */
|
* different MTLContext's. */
|
||||||
if (MTLCommandBufferManager::sync_event == nil) {
|
if (MTLCommandBufferManager::sync_event == nil) {
|
||||||
MTLCommandBufferManager::sync_event = [context_->device newEvent];
|
MTLCommandBufferManager::sync_event = [context_.device newEvent];
|
||||||
BLI_assert(MTLCommandBufferManager::sync_event);
|
BLI_assert(MTLCommandBufferManager::sync_event);
|
||||||
[MTLCommandBufferManager::sync_event retain];
|
[MTLCommandBufferManager::sync_event retain];
|
||||||
}
|
}
|
||||||
@@ -102,14 +108,27 @@ bool MTLCommandBufferManager::submit(bool wait)
|
|||||||
value:MTLCommandBufferManager::event_signal_val];
|
value:MTLCommandBufferManager::event_signal_val];
|
||||||
|
|
||||||
/* Command buffer lifetime tracking. */
|
/* Command buffer lifetime tracking. */
|
||||||
/* TODO(Metal): This routine will later be used to track released memory allocations within the
|
/* Increment current MTLSafeFreeList reference counter to flag MTLBuffers freed within
|
||||||
* lifetime of a command buffer such that memory is only released once no longer in use. */
|
* the current command buffer lifetime as used.
|
||||||
id<MTLCommandBuffer> cmd_buffer_ref = [active_command_buffer_ retain];
|
* 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) {
|
[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. */
|
/* Release command buffer after completion callback handled. */
|
||||||
[cmd_buffer_ref release];
|
[cmd_buffer_ref release];
|
||||||
|
|
||||||
/* Decrement active cmd buffer count. */
|
/* Decrement count. */
|
||||||
MTLCommandBufferManager::num_active_cmd_bufs--;
|
MTLCommandBufferManager::num_active_cmd_bufs--;
|
||||||
}];
|
}];
|
||||||
|
|
||||||
@@ -516,15 +535,6 @@ bool MTLCommandBufferManager::insert_memory_barrier(eGPUBarrier barrier_bits,
|
|||||||
/* -------------------------------------------------------------------- */
|
/* -------------------------------------------------------------------- */
|
||||||
/** \name Render Pass State for active RenderCommandEncoder
|
/** \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
|
/* Reset binding state when a new RenderCommandEncoder is bound, to ensure
|
||||||
* pipeline resources are re-applied to the new Encoder.
|
* pipeline resources are re-applied to the new Encoder.
|
||||||
* NOTE: In Metal, state is only persistent within an MTLCommandEncoder,
|
* 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);
|
this->last_bound_shader_state.set(nullptr, 0);
|
||||||
|
|
||||||
/* Other states. */
|
/* 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_used_stencil_ref_value = 0;
|
||||||
this->last_scissor_rect = {0,
|
this->last_scissor_rect = {0,
|
||||||
0,
|
0,
|
||||||
(unsigned long)((fb != nullptr) ? fb->get_width() : 0),
|
(uint)((fb != nullptr) ? fb->get_width() : 0),
|
||||||
(unsigned long)((fb != nullptr) ? fb->get_height() : 0)};
|
(uint)((fb != nullptr) ? fb->get_height() : 0)};
|
||||||
|
|
||||||
/* Reset cached resource binding state */
|
/* Reset cached resource binding state */
|
||||||
for (int ubo = 0; ubo < MTL_MAX_UNIFORM_BUFFER_BINDINGS; ubo++) {
|
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)
|
void MTLRenderPassState::bind_vertex_texture(id<MTLTexture> tex, uint slot)
|
||||||
{
|
{
|
||||||
if (this->cached_vertex_texture_bindings[slot].metal_texture != tex) {
|
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);
|
BLI_assert(rec != nil);
|
||||||
[rec setVertexTexture:tex atIndex:slot];
|
[rec setVertexTexture:tex atIndex:slot];
|
||||||
this->cached_vertex_texture_bindings[slot].metal_texture = tex;
|
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)
|
void MTLRenderPassState::bind_fragment_texture(id<MTLTexture> tex, uint slot)
|
||||||
{
|
{
|
||||||
if (this->cached_fragment_texture_bindings[slot].metal_texture != tex) {
|
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);
|
BLI_assert(rec != nil);
|
||||||
[rec setFragmentTexture:tex atIndex:slot];
|
[rec setFragmentTexture:tex atIndex:slot];
|
||||||
this->cached_fragment_texture_bindings[slot].metal_texture = tex;
|
this->cached_fragment_texture_bindings[slot].metal_texture = tex;
|
||||||
|
@@ -4,8 +4,13 @@
|
|||||||
#define __MTL_COMMON
|
#define __MTL_COMMON
|
||||||
|
|
||||||
// -- Renderer Options --
|
// -- Renderer Options --
|
||||||
|
#define MTL_MAX_DRAWABLES 3
|
||||||
#define MTL_MAX_SET_BYTES_SIZE 4096
|
#define MTL_MAX_SET_BYTES_SIZE 4096
|
||||||
#define MTL_FORCE_WAIT_IDLE 0
|
#define MTL_FORCE_WAIT_IDLE 0
|
||||||
#define MTL_MAX_COMMAND_BUFFERS 64
|
#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
|
#endif
|
||||||
|
@@ -12,7 +12,9 @@
|
|||||||
|
|
||||||
#include "mtl_backend.hh"
|
#include "mtl_backend.hh"
|
||||||
#include "mtl_capabilities.hh"
|
#include "mtl_capabilities.hh"
|
||||||
|
#include "mtl_common.hh"
|
||||||
#include "mtl_framebuffer.hh"
|
#include "mtl_framebuffer.hh"
|
||||||
|
#include "mtl_memory.hh"
|
||||||
#include "mtl_texture.hh"
|
#include "mtl_texture.hh"
|
||||||
|
|
||||||
#include <Cocoa/Cocoa.h>
|
#include <Cocoa/Cocoa.h>
|
||||||
@@ -30,7 +32,6 @@ class MTLContext;
|
|||||||
class MTLCommandBufferManager;
|
class MTLCommandBufferManager;
|
||||||
class MTLShader;
|
class MTLShader;
|
||||||
class MTLUniformBuf;
|
class MTLUniformBuf;
|
||||||
class MTLBuffer;
|
|
||||||
|
|
||||||
/* Structs containing information on current binding state for textures and samplers. */
|
/* Structs containing information on current binding state for textures and samplers. */
|
||||||
struct MTLTextureBinding {
|
struct MTLTextureBinding {
|
||||||
@@ -56,10 +57,13 @@ struct MTLSamplerBinding {
|
|||||||
struct MTLRenderPassState {
|
struct MTLRenderPassState {
|
||||||
friend class MTLContext;
|
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,
|
/* Given a RenderPassState is associated with a live RenderCommandEncoder,
|
||||||
* this state sits within the MTLCommandBufferManager. */
|
* this state sits within the MTLCommandBufferManager. */
|
||||||
MTLCommandBufferManager *cmd;
|
MTLContext &ctx;
|
||||||
MTLContext *ctx;
|
MTLCommandBufferManager &cmd;
|
||||||
|
|
||||||
/* Caching of resource bindings for active MTLRenderCommandEncoder.
|
/* Caching of resource bindings for active MTLRenderCommandEncoder.
|
||||||
* In Metal, resource bindings are local to the MTLCommandEncoder,
|
* 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_vertex_sampler_state_bindings[MTL_MAX_TEXTURE_SLOTS];
|
||||||
SamplerStateBindingCached cached_fragment_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. */
|
/* Reset RenderCommandEncoder binding state. */
|
||||||
void reset_state();
|
void reset_state();
|
||||||
|
|
||||||
@@ -446,18 +447,6 @@ struct MTLContextGlobalShaderPipelineState {
|
|||||||
float line_width = 1.0f;
|
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.
|
/* Command Buffer Manager - Owned by MTLContext.
|
||||||
* The MTLCommandBufferManager represents all work associated with
|
* The MTLCommandBufferManager represents all work associated with
|
||||||
* a command buffer of a given identity. This manager is a fixed-state
|
* a command buffer of a given identity. This manager is a fixed-state
|
||||||
@@ -477,14 +466,14 @@ class MTLCommandBufferManager {
|
|||||||
public:
|
public:
|
||||||
/* Event to coordinate sequential execution across all "main" command buffers. */
|
/* Event to coordinate sequential execution across all "main" command buffers. */
|
||||||
static id<MTLEvent> sync_event;
|
static id<MTLEvent> sync_event;
|
||||||
static unsigned long long event_signal_val;
|
static uint64_t event_signal_val;
|
||||||
|
|
||||||
/* Counter for active command buffers. */
|
/* Counter for active command buffers. */
|
||||||
static int num_active_cmd_bufs;
|
static int num_active_cmd_bufs;
|
||||||
|
|
||||||
private:
|
private:
|
||||||
/* Associated Context and properties. */
|
/* Associated Context and properties. */
|
||||||
MTLContext *context_ = nullptr;
|
MTLContext &context_;
|
||||||
bool supports_render_ = false;
|
bool supports_render_ = false;
|
||||||
|
|
||||||
/* CommandBuffer tracking. */
|
/* CommandBuffer tracking. */
|
||||||
@@ -516,7 +505,9 @@ class MTLCommandBufferManager {
|
|||||||
bool empty_ = true;
|
bool empty_ = true;
|
||||||
|
|
||||||
public:
|
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. */
|
/* If wait is true, CPU will stall until GPU work has completed. */
|
||||||
bool submit(bool wait);
|
bool submit(bool wait);
|
||||||
@@ -582,7 +573,7 @@ class MTLContext : public Context {
|
|||||||
|
|
||||||
/* Texture Samplers. */
|
/* Texture Samplers. */
|
||||||
/* Cache of generated MTLSamplerState objects based on permutations of `eGPUSamplerState`. */
|
/* 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;
|
id<MTLSamplerState> default_sampler_state_ = nil;
|
||||||
|
|
||||||
/* When texture sampler count exceeds the resource bind limit, an
|
/* When texture sampler count exceeds the resource bind limit, an
|
||||||
@@ -595,6 +586,7 @@ class MTLContext : public Context {
|
|||||||
|
|
||||||
/* Frame. */
|
/* Frame. */
|
||||||
bool is_inside_frame_ = false;
|
bool is_inside_frame_ = false;
|
||||||
|
uint current_frame_index_;
|
||||||
|
|
||||||
public:
|
public:
|
||||||
/* Shaders and Pipeline state. */
|
/* Shaders and Pipeline state. */
|
||||||
@@ -604,6 +596,10 @@ class MTLContext : public Context {
|
|||||||
id<MTLCommandQueue> queue = nil;
|
id<MTLCommandQueue> queue = nil;
|
||||||
id<MTLDevice> device = nil;
|
id<MTLDevice> device = nil;
|
||||||
|
|
||||||
|
/* Memory Management */
|
||||||
|
MTLScratchBufferManager memory_manager;
|
||||||
|
static MTLBufferPool global_memory_manager;
|
||||||
|
|
||||||
/* CommandBuffer managers. */
|
/* CommandBuffer managers. */
|
||||||
MTLCommandBufferManager main_command_buffer;
|
MTLCommandBufferManager main_command_buffer;
|
||||||
|
|
||||||
@@ -624,7 +620,7 @@ class MTLContext : public Context {
|
|||||||
void memory_statistics_get(int *total_mem, int *free_mem) override;
|
void memory_statistics_get(int *total_mem, int *free_mem) override;
|
||||||
|
|
||||||
void debug_group_begin(const char *name, int index) 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. */
|
/*** MTLContext Utility functions. */
|
||||||
/*
|
/*
|
||||||
@@ -679,6 +675,21 @@ class MTLContext : public Context {
|
|||||||
{
|
{
|
||||||
return is_inside_frame_;
|
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
|
} // namespace blender::gpu
|
||||||
|
@@ -16,44 +16,25 @@ using namespace blender::gpu;
|
|||||||
|
|
||||||
namespace blender::gpu {
|
namespace blender::gpu {
|
||||||
|
|
||||||
/* -------------------------------------------------------------------- */
|
/* Global memory mamnager */
|
||||||
/** \name Memory Management
|
MTLBufferPool MTLContext::global_memory_manager;
|
||||||
* \{ */
|
|
||||||
|
|
||||||
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)];
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
/** \} */
|
|
||||||
|
|
||||||
/* -------------------------------------------------------------------- */
|
/* -------------------------------------------------------------------- */
|
||||||
/** \name MTLContext
|
/** \name MTLContext
|
||||||
* \{ */
|
* \{ */
|
||||||
|
|
||||||
/* Placeholder functions */
|
/* Placeholder functions */
|
||||||
MTLContext::MTLContext(void *ghost_window)
|
MTLContext::MTLContext(void *ghost_window) : memory_manager(*this), main_command_buffer(*this)
|
||||||
{
|
{
|
||||||
/* Init debug. */
|
/* Init debug. */
|
||||||
debug::mtl_debug_init();
|
debug::mtl_debug_init();
|
||||||
|
|
||||||
/* Initialize command buffer state. */
|
/* Initialize command buffer state. */
|
||||||
this->main_command_buffer.prepare(this);
|
this->main_command_buffer.prepare();
|
||||||
|
|
||||||
/* Frame management. */
|
/* Frame management. */
|
||||||
is_inside_frame_ = false;
|
is_inside_frame_ = false;
|
||||||
|
current_frame_index_ = 0;
|
||||||
|
|
||||||
/* Create FrameBuffer handles. */
|
/* Create FrameBuffer handles. */
|
||||||
MTLFrameBuffer *mtl_front_left = new MTLFrameBuffer(this, "front_left");
|
MTLFrameBuffer *mtl_front_left = new MTLFrameBuffer(this, "front_left");
|
||||||
@@ -65,9 +46,14 @@ MTLContext::MTLContext(void *ghost_window)
|
|||||||
* initialization). */
|
* initialization). */
|
||||||
MTLBackend::platform_init(this);
|
MTLBackend::platform_init(this);
|
||||||
MTLBackend::capabilities_init(this);
|
MTLBackend::capabilities_init(this);
|
||||||
|
|
||||||
/* Initialize Metal modules. */
|
/* Initialize Metal modules. */
|
||||||
|
this->memory_manager.init();
|
||||||
this->state_manager = new MTLStateManager(this);
|
this->state_manager = new MTLStateManager(this);
|
||||||
|
|
||||||
|
/* Ensure global memory manager is initialied */
|
||||||
|
MTLContext::global_memory_manager.init(this->device);
|
||||||
|
|
||||||
/* Initialize texture read/update structures. */
|
/* Initialize texture read/update structures. */
|
||||||
this->get_texture_utils().init();
|
this->get_texture_utils().init();
|
||||||
|
|
||||||
@@ -93,7 +79,7 @@ MTLContext::~MTLContext()
|
|||||||
this->finish();
|
this->finish();
|
||||||
|
|
||||||
/* End frame. */
|
/* End frame. */
|
||||||
if (is_inside_frame_) {
|
if (this->get_inside_frame()) {
|
||||||
this->end_frame();
|
this->end_frame();
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
@@ -112,7 +98,7 @@ MTLContext::~MTLContext()
|
|||||||
void MTLContext::begin_frame()
|
void MTLContext::begin_frame()
|
||||||
{
|
{
|
||||||
BLI_assert(MTLBackend::get()->is_inside_render_boundary());
|
BLI_assert(MTLBackend::get()->is_inside_render_boundary());
|
||||||
if (is_inside_frame_) {
|
if (this->get_inside_frame()) {
|
||||||
return;
|
return;
|
||||||
}
|
}
|
||||||
|
|
||||||
@@ -122,7 +108,7 @@ void MTLContext::begin_frame()
|
|||||||
|
|
||||||
void MTLContext::end_frame()
|
void MTLContext::end_frame()
|
||||||
{
|
{
|
||||||
BLI_assert(is_inside_frame_);
|
BLI_assert(this->get_inside_frame());
|
||||||
|
|
||||||
/* Ensure pre-present work is committed. */
|
/* Ensure pre-present work is committed. */
|
||||||
this->flush();
|
this->flush();
|
||||||
@@ -136,20 +122,20 @@ void MTLContext::check_error(const char *info)
|
|||||||
/* TODO(Metal): Implement. */
|
/* TODO(Metal): Implement. */
|
||||||
}
|
}
|
||||||
|
|
||||||
void MTLContext::activate(void)
|
void MTLContext::activate()
|
||||||
{
|
{
|
||||||
/* TODO(Metal): Implement. */
|
/* TODO(Metal): Implement. */
|
||||||
}
|
}
|
||||||
void MTLContext::deactivate(void)
|
void MTLContext::deactivate()
|
||||||
{
|
{
|
||||||
/* TODO(Metal): Implement. */
|
/* TODO(Metal): Implement. */
|
||||||
}
|
}
|
||||||
|
|
||||||
void MTLContext::flush(void)
|
void MTLContext::flush()
|
||||||
{
|
{
|
||||||
/* TODO(Metal): Implement. */
|
/* TODO(Metal): Implement. */
|
||||||
}
|
}
|
||||||
void MTLContext::finish(void)
|
void MTLContext::finish()
|
||||||
{
|
{
|
||||||
/* TODO(Metal): Implement. */
|
/* TODO(Metal): Implement. */
|
||||||
}
|
}
|
||||||
@@ -180,7 +166,7 @@ id<MTLRenderCommandEncoder> MTLContext::ensure_begin_render_pass()
|
|||||||
BLI_assert(this);
|
BLI_assert(this);
|
||||||
|
|
||||||
/* Ensure the rendering frame has started. */
|
/* Ensure the rendering frame has started. */
|
||||||
if (!is_inside_frame_) {
|
if (!this->get_inside_frame()) {
|
||||||
this->begin_frame();
|
this->begin_frame();
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@@ -756,7 +756,7 @@ void MTLFrameBuffer::update_attachments(bool update_viewport)
|
|||||||
dirty_attachments_ = false;
|
dirty_attachments_ = false;
|
||||||
}
|
}
|
||||||
|
|
||||||
void MTLFrameBuffer::apply_state(void)
|
void MTLFrameBuffer::apply_state()
|
||||||
{
|
{
|
||||||
MTLContext *mtl_ctx = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
|
MTLContext *mtl_ctx = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
|
||||||
BLI_assert(mtl_ctx);
|
BLI_assert(mtl_ctx);
|
||||||
|
476
source/blender/gpu/metal/mtl_memory.hh
Normal file
476
source/blender/gpu/metal/mtl_memory.hh
Normal 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
880
source/blender/gpu/metal/mtl_memory.mm
Normal 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
|
@@ -30,18 +30,18 @@ class MTLStateManager : public StateManager {
|
|||||||
public:
|
public:
|
||||||
MTLStateManager(MTLContext *ctx);
|
MTLStateManager(MTLContext *ctx);
|
||||||
|
|
||||||
void apply_state(void) override;
|
void apply_state() override;
|
||||||
void force_state(void) override;
|
void force_state() override;
|
||||||
|
|
||||||
void issue_barrier(eGPUBarrier barrier_bits) override;
|
void issue_barrier(eGPUBarrier barrier_bits) override;
|
||||||
|
|
||||||
void texture_bind(Texture *tex, eGPUSamplerState sampler, int unit) override;
|
void texture_bind(Texture *tex, eGPUSamplerState sampler, int unit) override;
|
||||||
void texture_unbind(Texture *tex) 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_bind(Texture *tex, int unit) override;
|
||||||
void image_unbind(Texture *tex) 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;
|
void texture_unpack_row_length_set(uint len) override;
|
||||||
|
|
||||||
|
@@ -17,7 +17,7 @@ namespace blender::gpu {
|
|||||||
/** \name MTLStateManager
|
/** \name MTLStateManager
|
||||||
* \{ */
|
* \{ */
|
||||||
|
|
||||||
void MTLStateManager::mtl_state_init(void)
|
void MTLStateManager::mtl_state_init()
|
||||||
{
|
{
|
||||||
BLI_assert(context_);
|
BLI_assert(context_);
|
||||||
context_->pipeline_state_init();
|
context_->pipeline_state_init();
|
||||||
@@ -36,7 +36,7 @@ MTLStateManager::MTLStateManager(MTLContext *ctx) : StateManager()
|
|||||||
set_mutable_state(mutable_state);
|
set_mutable_state(mutable_state);
|
||||||
}
|
}
|
||||||
|
|
||||||
void MTLStateManager::apply_state(void)
|
void MTLStateManager::apply_state()
|
||||||
{
|
{
|
||||||
this->set_state(this->state);
|
this->set_state(this->state);
|
||||||
this->set_mutable_state(this->mutable_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();
|
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. */
|
/* Little exception for clip distances since they need to keep the old count correct. */
|
||||||
uint32_t clip_distances = current_.clip_distances;
|
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.
|
/* Apple Silicon does not support memory barriers.
|
||||||
* We do not currently need these due to implicit API guarantees.
|
* 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. */
|
* untracked resources are ever used. */
|
||||||
if ([ctx->device hasUnifiedMemory]) {
|
if ([ctx->device hasUnifiedMemory]) {
|
||||||
return;
|
return;
|
||||||
@@ -600,7 +600,7 @@ void MTLStateManager::texture_unbind(Texture *tex_)
|
|||||||
ctx->texture_unbind(mtl_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()));
|
MTLContext *ctx = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
|
||||||
BLI_assert(ctx);
|
BLI_assert(ctx);
|
||||||
@@ -623,7 +623,7 @@ void MTLStateManager::image_unbind(Texture *tex_)
|
|||||||
this->texture_unbind(tex_);
|
this->texture_unbind(tex_);
|
||||||
}
|
}
|
||||||
|
|
||||||
void MTLStateManager::image_unbind_all(void)
|
void MTLStateManager::image_unbind_all()
|
||||||
{
|
{
|
||||||
this->texture_unbind_all();
|
this->texture_unbind_all();
|
||||||
}
|
}
|
||||||
|
@@ -237,7 +237,7 @@ class MTLTexture : public Texture {
|
|||||||
void update_sub(
|
void update_sub(
|
||||||
int mip, int offset[3], int extent[3], eGPUDataFormat type, const void *data) override;
|
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 copy_to(Texture *dst) override;
|
||||||
void clear(eGPUDataFormat format, const void *data) override;
|
void clear(eGPUDataFormat format, const void *data) override;
|
||||||
void swizzle_set(const char swizzle_mask[4]) 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;
|
void *read(int mip, eGPUDataFormat type) override;
|
||||||
|
|
||||||
/* Remove once no longer required -- will just return 0 for now in MTL path*/
|
/* 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();
|
bool texture_is_baked();
|
||||||
const char *get_name()
|
const char *get_name()
|
||||||
@@ -257,7 +257,7 @@ class MTLTexture : public Texture {
|
|||||||
}
|
}
|
||||||
|
|
||||||
protected:
|
protected:
|
||||||
bool init_internal(void) override;
|
bool init_internal() override;
|
||||||
bool init_internal(GPUVertBuf *vbo) override;
|
bool init_internal(GPUVertBuf *vbo) override;
|
||||||
bool init_internal(const GPUTexture *src,
|
bool init_internal(const GPUTexture *src,
|
||||||
int mip_offset,
|
int mip_offset,
|
||||||
|
@@ -478,23 +478,6 @@ void gpu::MTLTexture::update_sub(
|
|||||||
MTLPixelFormat destination_format = gpu_texture_format_to_metal(format_);
|
MTLPixelFormat destination_format = gpu_texture_format_to_metal(format_);
|
||||||
int expected_dst_bytes_per_pixel = get_mtl_format_bytesize(destination_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_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). */
|
/* Prepare specialisation struct (For texture update routine). */
|
||||||
TextureUpdateRoutineSpecialisation compute_specialisation_kernel = {
|
TextureUpdateRoutineSpecialisation compute_specialisation_kernel = {
|
||||||
@@ -568,12 +551,12 @@ void gpu::MTLTexture::update_sub(
|
|||||||
|
|
||||||
/* Prepare staging buffer for data. */
|
/* Prepare staging buffer for data. */
|
||||||
id<MTLBuffer> staging_buffer = nil;
|
id<MTLBuffer> staging_buffer = nil;
|
||||||
unsigned long long staging_buffer_offset = 0;
|
uint64_t staging_buffer_offset = 0;
|
||||||
|
|
||||||
/* Fetch allocation from scratch buffer. */
|
/* Fetch allocation from scratch buffer. */
|
||||||
MTLTemporaryBufferRange allocation; /* TODO(Metal): Metal Memory manager. */
|
MTLTemporaryBuffer allocation =
|
||||||
/* = ctx->get_memory_manager().scratch_buffer_allocate_range_aligned(totalsize, 256);*/
|
ctx->get_scratchbuffer_manager().scratch_buffer_allocate_range_aligned(totalsize, 256);
|
||||||
memcpy(allocation.host_ptr, data, totalsize);
|
memcpy(allocation.data, data, totalsize);
|
||||||
staging_buffer = allocation.metal_buffer;
|
staging_buffer = allocation.metal_buffer;
|
||||||
staging_buffer_offset = allocation.buffer_offset;
|
staging_buffer_offset = allocation.buffer_offset;
|
||||||
|
|
||||||
@@ -915,7 +898,7 @@ void gpu::MTLTexture::ensure_mipmaps(int miplvl)
|
|||||||
this->mip_range_set(0, mipmaps_);
|
this->mip_range_set(0, mipmaps_);
|
||||||
}
|
}
|
||||||
|
|
||||||
void gpu::MTLTexture::generate_mipmap(void)
|
void gpu::MTLTexture::generate_mipmap()
|
||||||
{
|
{
|
||||||
/* Fetch Active Context. */
|
/* Fetch Active Context. */
|
||||||
MTLContext *ctx = reinterpret_cast<MTLContext *>(GPU_context_active_get());
|
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)
|
destination_buffer = [ctx->device newBufferWithLength:max_ii(total_bytes, 256)
|
||||||
options:bufferOptions];
|
options:bufferOptions];
|
||||||
destination_offset = 0;
|
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);
|
destination_offset);
|
||||||
|
|
||||||
/* Prepare specialisation struct (For non-trivial texture read routine). */
|
/* 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. */
|
/* 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;
|
return 0;
|
||||||
}
|
}
|
||||||
|
|
||||||
bool gpu::MTLTexture::init_internal(void)
|
bool gpu::MTLTexture::init_internal()
|
||||||
{
|
{
|
||||||
if (format_ == GPU_DEPTH24_STENCIL8) {
|
if (format_ == GPU_DEPTH24_STENCIL8) {
|
||||||
/* Apple Silicon requires GPU_DEPTH32F_STENCIL8 instead of GPU_DEPTH24_STENCIL8. */
|
/* Apple Silicon requires GPU_DEPTH32F_STENCIL8 instead of GPU_DEPTH24_STENCIL8. */
|
||||||
|
Reference in New Issue
Block a user