Page MenuHome
Differential D16990 Diff 60050 source/blender/gpu/metal/mtl_command_buffer.mm

Changeset View

Standalone View

View Options

source/blender/gpu/metal/mtl_command_buffer.mm

Show First 20 Lines • Show All 410 Lines • ▼ Show 20 Linesid<MTLComputeCommandEncoder> MTLCommandBufferManager::ensure_begin_compute_encoder()
if (active_compute_command_encoder_ == nil) { if (active_compute_command_encoder_ == nil) {
active_compute_command_encoder_ = [cmd_buf computeCommandEncoder]; active_compute_command_encoder_ = [cmd_buf computeCommandEncoder];
BLI_assert(active_compute_command_encoder_ != nil); BLI_assert(active_compute_command_encoder_ != nil);
[active_compute_command_encoder_ retain]; [active_compute_command_encoder_ retain];
active_command_encoder_type_ = MTL_COMPUTE_COMMAND_ENCODER; active_command_encoder_type_ = MTL_COMPUTE_COMMAND_ENCODER;
/* Update command buffer encoder heuristics. */ /* Update command buffer encoder heuristics. */
this->register_encoder_counters(); this->register_encoder_counters();
/* Reset RenderPassState to ensure resource bindings are re-applied. */
compute_state_.reset_state();
} }
BLI_assert(active_compute_command_encoder_ != nil); BLI_assert(active_compute_command_encoder_ != nil);
return active_compute_command_encoder_; return active_compute_command_encoder_;
} }
/** \} */ /** \} */
/* -------------------------------------------------------------------- */ /* -------------------------------------------------------------------- */
▲ Show 20 Lines • Show All 64 Lines • ▼ Show 20 Lines
/* Workload Synchronization. */ /* Workload Synchronization. */
bool MTLCommandBufferManager::insert_memory_barrier(eGPUBarrier barrier_bits, bool MTLCommandBufferManager::insert_memory_barrier(eGPUBarrier barrier_bits,
eGPUStageBarrierBits before_stages, eGPUStageBarrierBits before_stages,
eGPUStageBarrierBits after_stages) eGPUStageBarrierBits after_stages)
{ {
/* Only supporting Metal on 10.14 onward anyway - Check required for warnings. */ /* Only supporting Metal on 10.14 onward anyway - Check required for warnings. */
if (@available(macOS 10.14, *)) { if (@available(macOS 10.14, *)) {
/* Apple Silicon does not support memory barriers for RenderCommandEncoder's.
* We do not currently need these due to implicit API guarantees.
* NOTE(Metal): MTLFence/MTLEvent may be required to synchronize work if
* untracked resources are ever used. */
if ([context_.device hasUnifiedMemory] &&
(active_command_encoder_type_ != MTL_COMPUTE_COMMAND_ENCODER)) {
return false;
}
/* Resolve scope. */ /* Resolve scope. */
MTLBarrierScope scope = 0; MTLBarrierScope scope = 0;
if (barrier_bits & GPU_BARRIER_SHADER_IMAGE_ACCESS || if (barrier_bits & GPU_BARRIER_SHADER_IMAGE_ACCESS ||
barrier_bits & GPU_BARRIER_TEXTURE_FETCH) { barrier_bits & GPU_BARRIER_TEXTURE_FETCH) {
scope = scope | MTLBarrierScopeTextures | MTLBarrierScopeRenderTargets; bool is_compute = (active_command_encoder_type_ != MTL_RENDER_COMMAND_ENCODER);
scope |= (is_compute ? 0 : MTLBarrierScopeRenderTargets) | MTLBarrierScopeTextures;
} }
fclemUnsubmitted
Not Done
Inline Actions
barrier_bits & GPU_BARRIER_TEXTURE_FETCH) {
- scope = scope |
- ((active_command_encoder_type_ == MTL_RENDER_COMMAND_ENCODER) ?
- MTLBarrierScopeRenderTargets :
- 0) |
- MTLBarrierScopeTextures;
+ bool is_compute = (active_command_encoder_type_ != MTL_RENDER_COMMAND_ENCODER);
+ scope |= (is_compute ? 0 : MTLBarrierScopeRenderTargets) | MTLBarrierScopeTextures;
}
if (barrier_bits & GPU_BARRIER_SHADER_STORAGE ||

Readability.

fclem: Readability.
if (barrier_bits & GPU_BARRIER_SHADER_STORAGE || if (barrier_bits & GPU_BARRIER_SHADER_STORAGE ||
barrier_bits & GPU_BARRIER_VERTEX_ATTRIB_ARRAY || barrier_bits & GPU_BARRIER_VERTEX_ATTRIB_ARRAY ||
barrier_bits & GPU_BARRIER_ELEMENT_ARRAY || barrier_bits & GPU_BARRIER_UNIFORM) { barrier_bits & GPU_BARRIER_ELEMENT_ARRAY || barrier_bits & GPU_BARRIER_UNIFORM) {
scope = scope | MTLBarrierScopeBuffers; scope = scope | MTLBarrierScopeBuffers;
} }
if (scope != 0) { if (scope != 0) {
/* Issue barrier based on encoder. */ /* Issue barrier based on encoder. */
▲ Show 20 Lines • Show All 86 Lines • ▼ Show 20 Linesvoid MTLRenderPassState::reset_state()
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,
(uint)((fb != nullptr) ? fb->get_width() : 0), (uint)((fb != nullptr) ? fb->get_width() : 0),
(uint)((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_BUFFER_BINDINGS; ubo++) {
this->cached_vertex_buffer_bindings[ubo].is_bytes = false; this->cached_vertex_buffer_bindings[ubo].is_bytes = false;
this->cached_vertex_buffer_bindings[ubo].metal_buffer = nil; this->cached_vertex_buffer_bindings[ubo].metal_buffer = nil;
this->cached_vertex_buffer_bindings[ubo].offset = -1; this->cached_vertex_buffer_bindings[ubo].offset = -1;
this->cached_fragment_buffer_bindings[ubo].is_bytes = false; this->cached_fragment_buffer_bindings[ubo].is_bytes = false;
this->cached_fragment_buffer_bindings[ubo].metal_buffer = nil; this->cached_fragment_buffer_bindings[ubo].metal_buffer = nil;
this->cached_fragment_buffer_bindings[ubo].offset = -1; this->cached_fragment_buffer_bindings[ubo].offset = -1;
} }
/* Reset cached texture and sampler state binding state. */ /* Reset cached texture and sampler state binding state. */
for (int tex = 0; tex < MTL_MAX_TEXTURE_SLOTS; tex++) { for (int tex = 0; tex < MTL_MAX_TEXTURE_SLOTS; tex++) {
this->cached_vertex_texture_bindings[tex].metal_texture = nil; this->cached_vertex_texture_bindings[tex].metal_texture = nil;
this->cached_vertex_sampler_state_bindings[tex].sampler_state = nil; this->cached_vertex_sampler_state_bindings[tex].sampler_state = nil;
this->cached_vertex_sampler_state_bindings[tex].is_arg_buffer_binding = false; this->cached_vertex_sampler_state_bindings[tex].is_arg_buffer_binding = false;
this->cached_fragment_texture_bindings[tex].metal_texture = nil; this->cached_fragment_texture_bindings[tex].metal_texture = nil;
this->cached_fragment_sampler_state_bindings[tex].sampler_state = nil; this->cached_fragment_sampler_state_bindings[tex].sampler_state = nil;
this->cached_fragment_sampler_state_bindings[tex].is_arg_buffer_binding = false; this->cached_fragment_sampler_state_bindings[tex].is_arg_buffer_binding = false;
} }
} }
void MTLComputeState::reset_state()
{
/* Reset Cached pipeline state. */
this->bound_pso = nil;
/* Reset cached resource binding state */
for (int ubo = 0; ubo < MTL_MAX_BUFFER_BINDINGS; ubo++) {
this->cached_compute_buffer_bindings[ubo].is_bytes = false;
this->cached_compute_buffer_bindings[ubo].metal_buffer = nil;
this->cached_compute_buffer_bindings[ubo].offset = -1;
}
/* Reset cached texture and sampler state binding state. */
for (int tex = 0; tex < MTL_MAX_TEXTURE_SLOTS; tex++) {
this->cached_compute_texture_bindings[tex].metal_texture = nil;
this->cached_compute_sampler_state_bindings[tex].sampler_state = nil;
this->cached_compute_sampler_state_bindings[tex].is_arg_buffer_binding = false;
}
}
/* Bind Texture to current RenderCommandEncoder. */ /* Bind Texture to current RenderCommandEncoder. */
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;
} }
} }
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;
} }
} }
void MTLComputeState::bind_compute_texture(id<MTLTexture> tex, uint slot)
{
if (this->cached_compute_texture_bindings[slot].metal_texture != tex) {
id<MTLComputeCommandEncoder> rec = this->cmd.get_active_compute_command_encoder();
BLI_assert(rec != nil);
[rec setTexture:tex atIndex:slot];
[rec useResource:tex
usage:MTLResourceUsageRead | MTLResourceUsageWrite | MTLResourceUsageSample];
this->cached_compute_texture_bindings[slot].metal_texture = tex;
}
}
void MTLRenderPassState::bind_vertex_sampler(MTLSamplerBinding &sampler_binding, void MTLRenderPassState::bind_vertex_sampler(MTLSamplerBinding &sampler_binding,
bool use_argument_buffer_for_samplers, bool use_argument_buffer_for_samplers,
uint slot) uint slot)
{ {
/* Range check. */ /* Range check. */
const MTLShaderInterface *shader_interface = ctx.pipeline_state.active_shader->get_interface(); const MTLShaderInterface *shader_interface = ctx.pipeline_state.active_shader->get_interface();
BLI_assert(slot >= 0); BLI_assert(slot >= 0);
BLI_assert(slot <= shader_interface->get_max_texture_index()); BLI_assert(slot <= shader_interface->get_max_texture_index());
▲ Show 20 Lines • Show All 63 Lines • ▼ Show 20 Lines if (this->cached_fragment_sampler_state_bindings[slot].sampler_state == nil ||
/* Always assign to argument buffer samplers binding array - Efficiently ensures the value in /* Always assign to argument buffer samplers binding array - Efficiently ensures the value in
* the samplers array is always up to date. */ * the samplers array is always up to date. */
ctx.samplers_.mtl_sampler[slot] = sampler_state; ctx.samplers_.mtl_sampler[slot] = sampler_state;
ctx.samplers_.mtl_sampler_flags[slot] = sampler_binding.state; ctx.samplers_.mtl_sampler_flags[slot] = sampler_binding.state;
} }
} }
void MTLComputeState::bind_compute_sampler(MTLSamplerBinding &sampler_binding,
bool use_argument_buffer_for_samplers,
uint slot)
{
/* Range check. */
const MTLShaderInterface *shader_interface = ctx.pipeline_state.active_shader->get_interface();
BLI_assert(slot >= 0);
BLI_assert(slot <= shader_interface->get_max_texture_index());
BLI_assert(slot < MTL_MAX_TEXTURE_SLOTS);
UNUSED_VARS_NDEBUG(shader_interface);
/* If sampler state has not changed for the given slot, we do not need to fetch. */
if (this->cached_compute_sampler_state_bindings[slot].sampler_state == nil ||
!(this->cached_compute_sampler_state_bindings[slot].binding_state ==
sampler_binding.state) ||
use_argument_buffer_for_samplers) {
id<MTLSamplerState> sampler_state = (sampler_binding.state == DEFAULT_SAMPLER_STATE) ?
ctx.get_default_sampler_state() :
ctx.get_sampler_from_state(sampler_binding.state);
if (!use_argument_buffer_for_samplers) {
/* Update binding and cached state. */
id<MTLComputeCommandEncoder> rec = this->cmd.get_active_compute_command_encoder();
BLI_assert(rec != nil);
[rec setSamplerState:sampler_state atIndex:slot];
this->cached_compute_sampler_state_bindings[slot].binding_state = sampler_binding.state;
this->cached_compute_sampler_state_bindings[slot].sampler_state = sampler_state;
}
/* Flag last binding type */
this->cached_compute_sampler_state_bindings[slot].is_arg_buffer_binding =
use_argument_buffer_for_samplers;
/* Always assign to argument buffer samplers binding array - Efficiently ensures the value in
* the samplers array is always up to date. */
ctx.samplers_.mtl_sampler[slot] = sampler_state;
ctx.samplers_.mtl_sampler_flags[slot] = sampler_binding.state;
}
}
void MTLRenderPassState::bind_vertex_buffer(id<MTLBuffer> buffer, uint buffer_offset, uint index) void MTLRenderPassState::bind_vertex_buffer(id<MTLBuffer> buffer, uint buffer_offset, uint index)
{ {
BLI_assert(index >= 0); BLI_assert(index >= 0 && index < MTL_MAX_BUFFER_BINDINGS);
BLI_assert(buffer_offset >= 0); BLI_assert(buffer_offset >= 0);
BLI_assert(buffer != nil); BLI_assert(buffer != nil);
BufferBindingCached &current_vert_ubo_binding = this->cached_vertex_buffer_bindings[index]; BufferBindingCached &current_vert_ubo_binding = this->cached_vertex_buffer_bindings[index];
if (current_vert_ubo_binding.offset != buffer_offset || if (current_vert_ubo_binding.offset != buffer_offset ||
current_vert_ubo_binding.metal_buffer != buffer || current_vert_ubo_binding.is_bytes) { current_vert_ubo_binding.metal_buffer != buffer || current_vert_ubo_binding.is_bytes) {
id<MTLRenderCommandEncoder> rec = this->cmd.get_active_render_command_encoder(); id<MTLRenderCommandEncoder> rec = this->cmd.get_active_render_command_encoder();
Show All 12 Lines if (current_vert_ubo_binding.offset != buffer_offset ||
this->cached_vertex_buffer_bindings[index].is_bytes = false; this->cached_vertex_buffer_bindings[index].is_bytes = false;
this->cached_vertex_buffer_bindings[index].metal_buffer = buffer; this->cached_vertex_buffer_bindings[index].metal_buffer = buffer;
this->cached_vertex_buffer_bindings[index].offset = buffer_offset; this->cached_vertex_buffer_bindings[index].offset = buffer_offset;
} }
} }
void MTLRenderPassState::bind_fragment_buffer(id<MTLBuffer> buffer, uint buffer_offset, uint index) void MTLRenderPassState::bind_fragment_buffer(id<MTLBuffer> buffer, uint buffer_offset, uint index)
{ {
BLI_assert(index >= 0); BLI_assert(index >= 0 && index < MTL_MAX_BUFFER_BINDINGS);
BLI_assert(buffer_offset >= 0); BLI_assert(buffer_offset >= 0);
BLI_assert(buffer != nil); BLI_assert(buffer != nil);
BufferBindingCached &current_frag_ubo_binding = this->cached_fragment_buffer_bindings[index]; BufferBindingCached &current_frag_ubo_binding = this->cached_fragment_buffer_bindings[index];
if (current_frag_ubo_binding.offset != buffer_offset || if (current_frag_ubo_binding.offset != buffer_offset ||
current_frag_ubo_binding.metal_buffer != buffer || current_frag_ubo_binding.is_bytes) { current_frag_ubo_binding.metal_buffer != buffer || current_frag_ubo_binding.is_bytes) {
id<MTLRenderCommandEncoder> rec = this->cmd.get_active_render_command_encoder(); id<MTLRenderCommandEncoder> rec = this->cmd.get_active_render_command_encoder();
Show All 10 Lines if (current_frag_ubo_binding.offset != buffer_offset ||
/* Update Bind-state cache */ /* Update Bind-state cache */
this->cached_fragment_buffer_bindings[index].is_bytes = false; this->cached_fragment_buffer_bindings[index].is_bytes = false;
this->cached_fragment_buffer_bindings[index].metal_buffer = buffer; this->cached_fragment_buffer_bindings[index].metal_buffer = buffer;
this->cached_fragment_buffer_bindings[index].offset = buffer_offset; this->cached_fragment_buffer_bindings[index].offset = buffer_offset;
} }
} }
void MTLComputeState::bind_compute_buffer(id<MTLBuffer> buffer,
uint buffer_offset,
uint index,
bool writeable)
{
BLI_assert(index >= 0 && index < MTL_MAX_BUFFER_BINDINGS);
BLI_assert(buffer_offset >= 0);
BLI_assert(buffer != nil);
BufferBindingCached &current_comp_ubo_binding = this->cached_compute_buffer_bindings[index];
if (current_comp_ubo_binding.offset != buffer_offset ||
current_comp_ubo_binding.metal_buffer != buffer || current_comp_ubo_binding.is_bytes) {
id<MTLComputeCommandEncoder> rec = this->cmd.get_active_compute_command_encoder();
BLI_assert(rec != nil);
if (current_comp_ubo_binding.metal_buffer == buffer) {
/* If buffer is the same, but offset has changed. */
[rec setBufferOffset:buffer_offset atIndex:index];
}
else {
/* Bind Fragment Buffer */
[rec setBuffer:buffer offset:buffer_offset atIndex:index];
}
[rec useResource:buffer
usage:((writeable) ? (MTLResourceUsageRead | MTLResourceUsageWrite) :
MTLResourceUsageRead)];
/* Update Bind-state cache */
this->cached_compute_buffer_bindings[index].is_bytes = false;
this->cached_compute_buffer_bindings[index].metal_buffer = buffer;
this->cached_compute_buffer_bindings[index].offset = buffer_offset;
}
}
void MTLRenderPassState::bind_vertex_bytes(void *bytes, uint length, uint index) void MTLRenderPassState::bind_vertex_bytes(void *bytes, uint length, uint index)
{ {
/* Bytes always updated as source data may have changed. */ /* Bytes always updated as source data may have changed. */
BLI_assert(index >= 0 && index < MTL_MAX_UNIFORM_BUFFER_BINDINGS); BLI_assert(index >= 0 && index < MTL_MAX_BUFFER_BINDINGS);
BLI_assert(length > 0); BLI_assert(length > 0);
BLI_assert(bytes != nullptr); BLI_assert(bytes != nullptr);
if (length < MTL_MAX_SET_BYTES_SIZE) { if (length < MTL_MAX_SET_BYTES_SIZE) {
id<MTLRenderCommandEncoder> rec = this->cmd.get_active_render_command_encoder(); id<MTLRenderCommandEncoder> rec = this->cmd.get_active_render_command_encoder();
[rec setVertexBytes:bytes length:length atIndex:index]; [rec setVertexBytes:bytes length:length atIndex:index];
} }
else { else {
/* We have run over the setBytes limit, bind buffer instead. */ /* We have run over the setBytes limit, bind buffer instead. */
MTLTemporaryBuffer range = MTLTemporaryBuffer range =
ctx.get_scratchbuffer_manager().scratch_buffer_allocate_range_aligned(length, 256); ctx.get_scratchbuffer_manager().scratch_buffer_allocate_range_aligned(length, 256);
memcpy(range.data, bytes, length); memcpy(range.data, bytes, length);
this->bind_vertex_buffer(range.metal_buffer, range.buffer_offset, index); this->bind_vertex_buffer(range.metal_buffer, range.buffer_offset, index);
} }
/* Update Bind-state cache */ /* Update Bind-state cache */
this->cached_vertex_buffer_bindings[index].is_bytes = true; this->cached_vertex_buffer_bindings[index].is_bytes = true;
this->cached_vertex_buffer_bindings[index].metal_buffer = nil; this->cached_vertex_buffer_bindings[index].metal_buffer = nil;
this->cached_vertex_buffer_bindings[index].offset = -1; this->cached_vertex_buffer_bindings[index].offset = -1;
} }
void MTLRenderPassState::bind_fragment_bytes(void *bytes, uint length, uint index) void MTLRenderPassState::bind_fragment_bytes(void *bytes, uint length, uint index)
{ {
/* Bytes always updated as source data may have changed. */ /* Bytes always updated as source data may have changed. */
BLI_assert(index >= 0 && index < MTL_MAX_UNIFORM_BUFFER_BINDINGS); BLI_assert(index >= 0 && index < MTL_MAX_BUFFER_BINDINGS);
BLI_assert(length > 0); BLI_assert(length > 0);
BLI_assert(bytes != nullptr); BLI_assert(bytes != nullptr);
if (length < MTL_MAX_SET_BYTES_SIZE) { if (length < MTL_MAX_SET_BYTES_SIZE) {
id<MTLRenderCommandEncoder> rec = this->cmd.get_active_render_command_encoder(); id<MTLRenderCommandEncoder> rec = this->cmd.get_active_render_command_encoder();
[rec setFragmentBytes:bytes length:length atIndex:index]; [rec setFragmentBytes:bytes length:length atIndex:index];
} }
else { else {
/* We have run over the setBytes limit, bind buffer instead. */ /* We have run over the setBytes limit, bind buffer instead. */
MTLTemporaryBuffer range = MTLTemporaryBuffer range =
ctx.get_scratchbuffer_manager().scratch_buffer_allocate_range_aligned(length, 256); ctx.get_scratchbuffer_manager().scratch_buffer_allocate_range_aligned(length, 256);
memcpy(range.data, bytes, length); memcpy(range.data, bytes, length);
this->bind_fragment_buffer(range.metal_buffer, range.buffer_offset, index); this->bind_fragment_buffer(range.metal_buffer, range.buffer_offset, index);
} }
/* Update Bind-state cache. */ /* Update Bind-state cache. */
this->cached_fragment_buffer_bindings[index].is_bytes = true; this->cached_fragment_buffer_bindings[index].is_bytes = true;
this->cached_fragment_buffer_bindings[index].metal_buffer = nil; this->cached_fragment_buffer_bindings[index].metal_buffer = nil;
this->cached_fragment_buffer_bindings[index].offset = -1; this->cached_fragment_buffer_bindings[index].offset = -1;
} }
void MTLComputeState::bind_compute_bytes(void *bytes, uint length, uint index)
{
/* Bytes always updated as source data may have changed. */
BLI_assert(index >= 0 && index < MTL_MAX_BUFFER_BINDINGS);
BLI_assert(length > 0);
BLI_assert(bytes != nullptr);
if (length < MTL_MAX_SET_BYTES_SIZE) {
id<MTLComputeCommandEncoder> rec = this->cmd.get_active_compute_command_encoder();
[rec setBytes:bytes length:length atIndex:index];
}
else {
/* We have run over the setBytes limit, bind buffer instead. */
MTLTemporaryBuffer range =
ctx.get_scratchbuffer_manager().scratch_buffer_allocate_range_aligned(length, 256);
memcpy(range.data, bytes, length);
this->bind_compute_buffer(range.metal_buffer, range.buffer_offset, index);
}
/* Update Bind-state cache. */
this->cached_compute_buffer_bindings[index].is_bytes = true;
this->cached_compute_buffer_bindings[index].metal_buffer = nil;
this->cached_compute_buffer_bindings[index].offset = -1;
}
void MTLComputeState::bind_pso(id<MTLComputePipelineState> pso)
{
if (this->bound_pso != pso) {
id<MTLComputeCommandEncoder> rec = this->cmd.get_active_compute_command_encoder();
[rec setComputePipelineState:pso];
this->bound_pso = pso;
}
}
/** \} */ /** \} */
} // blender::gpu } // blender::gpu