griefith/test2
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

Metal: Remove redundant synchronization operations

Browse Source 
Remove both compute barriers and useResource calls
as explicit resources bound via setTexture and
setComputeBuffer are implicitly tracked by the Metal
API anyway, so these calls increase complexity, without
altering correctness

Authored by Apple: Michael Parkin-White

Pull Request: https://projects.blender.org/blender/blender/pulls/121598
This commit is contained in:
Jason Fielder
2024-05-17 13:38:55 +02:00
committed by Jeroen Bakker
parent d4bf23771d
commit 47ada34324
5 changed files with 20 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -583,12 +583,16 @@ bool MTLCommandBufferManager::insert_memory_barrier(eGPUBarrier barrier_bits,
* synchronization using raster order groups, or, prefer compute to avoid subsequent passes
* re-loading pass attachments which are not needed. */
const bool is_tile_based_arch = (GPU_platform_architecture() == GPU_ARCHITECTURE_TBDR);
if (is_tile_based_arch && (active_command_encoder_type_ != MTL_COMPUTE_COMMAND_ENCODER)) {
if (is_tile_based_arch) {
if (active_command_encoder_type_ == MTL_RENDER_COMMAND_ENCODER) {
/* Break render pass to ensure final pass results are visible to subsequent calls. */
end_active_command_encoder();
return true;
}
return false;
else {
/* Skip all barriers for compute and blit passes as Metal will resolve these dependencies. */
return false;
}
}
/* Resolve scope. */
@@ -771,8 +775,6 @@ void MTLComputeState::bind_compute_texture(id<MTLTexture> tex, uint slot)
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;
}
@@ -964,10 +966,7 @@ void MTLRenderPassState::bind_fragment_buffer(id<MTLBuffer> buffer,
}
}
void MTLComputeState::bind_compute_buffer(id<MTLBuffer> buffer,
uint64_t buffer_offset,
uint index,
bool writeable)
void MTLComputeState::bind_compute_buffer(id<MTLBuffer> buffer, uint64_t buffer_offset, uint index)
{
BLI_assert(index >= 0 && index < MTL_MAX_BUFFER_BINDINGS);
BLI_assert(buffer_offset >= 0);
@@ -989,9 +988,6 @@ void MTLComputeState::bind_compute_buffer(id<MTLBuffer> buffer,
/* Bind Compute 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;

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

@@ -180,10 +180,7 @@ class MTLComputeState {
bool use_argument_buffer_for_samplers,
uint slot);
/* Buffer binding (ComputeCommandEncoder). */
void bind_compute_buffer(id<MTLBuffer> buffer,
uint64_t buffer_offset,
uint index,
bool writeable = false);
void bind_compute_buffer(id<MTLBuffer> buffer, uint64_t buffer_offset, uint index);
void bind_compute_bytes(const void *bytes, uint64_t length, uint index);
};

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

@@ -1547,7 +1547,7 @@ bool MTLContext::ensure_buffer_bindings(
/* Bind Compute SSBO. */
if (bool(ssbo.stage_mask & ShaderStage::COMPUTE)) {
BLI_assert(buffer_bind_index >= 0 && buffer_bind_index < MTL_MAX_BUFFER_BINDINGS);
cs.bind_compute_buffer(ssbo_buffer, 0, buffer_bind_index, true);
cs.bind_compute_buffer(ssbo_buffer, 0, buffer_bind_index);
}
}
else {

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

@@ -300,7 +300,7 @@ void MTLStorageBuf::clear(uint32_t clear_value)
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
cs.bind_pso(pso);
cs.bind_compute_bytes(&clear_value, sizeof(uint32_t), 0);
cs.bind_compute_buffer(metal_buffer_->get_metal_buffer(), 0, 1, true);
cs.bind_compute_buffer(metal_buffer_->get_metal_buffer(), 0, 1);
[compute_encoder dispatchThreads:MTLSizeMake(size_in_bytes_ / sizeof(uint32_t), 1, 1)
threadsPerThreadgroup:MTLSizeMake(128, 1, 1)];
}

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

@@ -804,7 +804,7 @@ void gpu::MTLTexture::update_sub(
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
cs.bind_pso(pso);
cs.bind_compute_bytes(&params, sizeof(params), 0);
cs.bind_compute_buffer(staging_buffer, 0, 1, true);
cs.bind_compute_buffer(staging_buffer, 0, 1);
cs.bind_compute_texture(texture_handle, 0);
[compute_encoder
dispatchThreads:MTLSizeMake(extent[0], 1, 1) /* Width, Height, Layer */
@@ -824,7 +824,7 @@ void gpu::MTLTexture::update_sub(
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
cs.bind_pso(pso);
cs.bind_compute_bytes(&params, sizeof(params), 0);
cs.bind_compute_buffer(staging_buffer, 0, 1, true);
cs.bind_compute_buffer(staging_buffer, 0, 1);
cs.bind_compute_texture(texture_handle, 0);
[compute_encoder
dispatchThreads:MTLSizeMake(extent[0], extent[1], 1) /* Width, layers, nil */
@@ -891,7 +891,7 @@ void gpu::MTLTexture::update_sub(
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
cs.bind_pso(pso);
cs.bind_compute_bytes(&params, sizeof(params), 0);
cs.bind_compute_buffer(staging_buffer, 0, 1, true);
cs.bind_compute_buffer(staging_buffer, 0, 1);
cs.bind_compute_texture(texture_handle, 0);
[compute_encoder
dispatchThreads:MTLSizeMake(
@@ -912,7 +912,7 @@ void gpu::MTLTexture::update_sub(
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
cs.bind_pso(pso);
cs.bind_compute_bytes(&params, sizeof(params), 0);
cs.bind_compute_buffer(staging_buffer, 0, 1, true);
cs.bind_compute_buffer(staging_buffer, 0, 1);
cs.bind_compute_texture(texture_handle, 0);
[compute_encoder dispatchThreads:MTLSizeMake(extent[0],
extent[1],
@@ -955,7 +955,7 @@ void gpu::MTLTexture::update_sub(
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
cs.bind_pso(pso);
cs.bind_compute_bytes(&params, sizeof(params), 0);
cs.bind_compute_buffer(staging_buffer, 0, 1, true);
cs.bind_compute_buffer(staging_buffer, 0, 1);
cs.bind_compute_texture(texture_handle, 0);
[compute_encoder
dispatchThreads:MTLSizeMake(
@@ -1767,7 +1767,7 @@ void gpu::MTLTexture::read_internal(int mip,
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
cs.bind_pso(pso);
cs.bind_compute_bytes(&params, sizeof(params), 0);
cs.bind_compute_buffer(destination_buffer, 0, 1, true);
cs.bind_compute_buffer(destination_buffer, 0, 1);
cs.bind_compute_texture(read_texture, 0);
[compute_encoder dispatchThreads:MTLSizeMake(width, 1, 1) /* Width, Height, Layer */
threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];
@@ -1817,7 +1817,7 @@ void gpu::MTLTexture::read_internal(int mip,
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
cs.bind_pso(pso);
cs.bind_compute_bytes(&params, sizeof(params), 0);
cs.bind_compute_buffer(destination_buffer, 0, 1, true);
cs.bind_compute_buffer(destination_buffer, 0, 1);
cs.bind_compute_texture(read_texture, 0);
[compute_encoder dispatchThreads:MTLSizeMake(width, height, 1) /* Width, Height, Layer */
threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];
@@ -1860,7 +1860,7 @@ void gpu::MTLTexture::read_internal(int mip,
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
cs.bind_pso(pso);
cs.bind_compute_bytes(&params, sizeof(params), 0);
cs.bind_compute_buffer(destination_buffer, 0, 1, true);
cs.bind_compute_buffer(destination_buffer, 0, 1);
cs.bind_compute_texture(read_texture, 0);
[compute_encoder dispatchThreads:MTLSizeMake(width, height, 1) /* Width, Height, Layer */
threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];
@@ -1910,7 +1910,7 @@ void gpu::MTLTexture::read_internal(int mip,
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
cs.bind_pso(pso);
cs.bind_compute_bytes(&params, sizeof(params), 0);
cs.bind_compute_buffer(destination_buffer, 0, 1, true);
cs.bind_compute_buffer(destination_buffer, 0, 1);
cs.bind_compute_texture(read_texture, 0);
[compute_encoder
dispatchThreads:MTLSizeMake(width, height, depth) /* Width, Height, Layer */
@@ -1954,7 +1954,7 @@ void gpu::MTLTexture::read_internal(int mip,
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
cs.bind_pso(pso);
cs.bind_compute_bytes(&params, sizeof(params), 0);
cs.bind_compute_buffer(destination_buffer, 0, 1, true);
cs.bind_compute_buffer(destination_buffer, 0, 1);
cs.bind_compute_texture(read_texture, 0);
[compute_encoder
dispatchThreads:MTLSizeMake(width, height, depth) /* Width, Height, Layer */