From cfe1f2b841a254083127cb8c875fcccde73b7890 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= Date: Mon, 19 Aug 2024 21:43:58 -0400 Subject: [PATCH] radeonsi: remove enum si_coherency SI_COHERENCY_SHADER is the only one that's used. Reviewed-by: Pierre-Eric Pelloux-Prayer Part-of: --- src/gallium/drivers/radeonsi/gfx11_query.c | 2 +- src/gallium/drivers/radeonsi/si_buffer.c | 2 +- src/gallium/drivers/radeonsi/si_clear.c | 9 +-- .../drivers/radeonsi/si_compute_blit.c | 78 +++++++------------ src/gallium/drivers/radeonsi/si_cp_dma.c | 50 +++++------- .../drivers/radeonsi/si_cp_reg_shadowing.c | 3 +- src/gallium/drivers/radeonsi/si_pipe.c | 5 +- src/gallium/drivers/radeonsi/si_pipe.h | 28 +++---- src/gallium/drivers/radeonsi/si_query.c | 3 +- src/gallium/drivers/radeonsi/si_shader.c | 2 +- .../drivers/radeonsi/si_test_dma_perf.c | 14 ++-- .../radeonsi/si_test_image_copy_region.c | 6 +- 12 files changed, 78 insertions(+), 124 deletions(-) diff --git a/src/gallium/drivers/radeonsi/gfx11_query.c b/src/gallium/drivers/radeonsi/gfx11_query.c index b6e5e34100a..ef33b4b2e40 100644 --- a/src/gallium/drivers/radeonsi/gfx11_query.c +++ b/src/gallium/drivers/radeonsi/gfx11_query.c @@ -393,7 +393,7 @@ static void gfx11_sh_query_get_result_resource(struct si_context *sctx, struct s /* ssbo[2] is either tmp_buffer or resource */ assert(ssbo[2].buffer); si_launch_grid_internal_ssbos(sctx, &grid, sctx->sh_query_result_shader, - SI_OP_SYNC_PS_BEFORE | SI_OP_SYNC_AFTER, SI_COHERENCY_SHADER, + SI_OP_SYNC_PS_BEFORE | SI_OP_SYNC_AFTER, 3, ssbo, (1 << 2) | (ssbo[1].buffer ? 1 << 1 : 0)); if (qbuf == query->last) diff --git a/src/gallium/drivers/radeonsi/si_buffer.c b/src/gallium/drivers/radeonsi/si_buffer.c index 9f017236be1..38f449dd553 100644 --- a/src/gallium/drivers/radeonsi/si_buffer.c +++ b/src/gallium/drivers/radeonsi/si_buffer.c @@ -195,7 +195,7 @@ bool si_alloc_resource(struct si_screen *sscreen, struct si_resource *res) uint32_t value = 0; si_clear_buffer(ctx, &res->b.b, 0, res->bo_size, &value, 4, SI_OP_SYNC_AFTER, - SI_COHERENCY_SHADER, SI_AUTO_SELECT_CLEAR_METHOD); + SI_AUTO_SELECT_CLEAR_METHOD); si_put_aux_context_flush(&sscreen->aux_context.general); } diff --git a/src/gallium/drivers/radeonsi/si_clear.c b/src/gallium/drivers/radeonsi/si_clear.c index 31e24998b87..7728917fd44 100644 --- a/src/gallium/drivers/radeonsi/si_clear.c +++ b/src/gallium/drivers/radeonsi/si_clear.c @@ -88,8 +88,7 @@ void si_execute_clears(struct si_context *sctx, struct si_clear_info *info, } if (info[i].is_dcc_msaa) { - gfx9_clear_dcc_msaa(sctx, info[i].resource, info[i].clear_value, - flags, SI_COHERENCY_SHADER); + gfx9_clear_dcc_msaa(sctx, info[i].resource, info[i].clear_value, flags); continue; } @@ -97,13 +96,11 @@ void si_execute_clears(struct si_context *sctx, struct si_clear_info *info, if (info[i].writemask != 0xffffffff) { si_compute_clear_buffer_rmw(sctx, info[i].resource, info[i].offset, info[i].size, - info[i].clear_value, info[i].writemask, - flags, SI_COHERENCY_SHADER); + info[i].clear_value, info[i].writemask, flags); } else { /* Compute shaders are much faster on both dGPUs and APUs. Don't use CP DMA. */ si_clear_buffer(sctx, info[i].resource, info[i].offset, info[i].size, - &info[i].clear_value, 4, flags, SI_COHERENCY_SHADER, - SI_COMPUTE_CLEAR_METHOD); + &info[i].clear_value, 4, flags, SI_COMPUTE_CLEAR_METHOD); } } diff --git a/src/gallium/drivers/radeonsi/si_compute_blit.c b/src/gallium/drivers/radeonsi/si_compute_blit.c index 35b00a30dab..7724c9dd328 100644 --- a/src/gallium/drivers/radeonsi/si_compute_blit.c +++ b/src/gallium/drivers/radeonsi/si_compute_blit.c @@ -12,15 +12,10 @@ #include "util/u_pack_color.h" #include "ac_nir_meta.h" -unsigned si_get_flush_flags(struct si_context *sctx, enum si_coherency coher, - enum si_cache_policy cache_policy) +unsigned si_get_flush_flags(struct si_context *sctx, enum si_cache_policy cache_policy) { - switch (coher) { - default: - case SI_COHERENCY_SHADER: - return SI_CONTEXT_INV_SCACHE | SI_CONTEXT_INV_VCACHE | - (cache_policy == L2_BYPASS ? SI_CONTEXT_INV_L2 : 0); - } + return SI_CONTEXT_INV_SCACHE | SI_CONTEXT_INV_VCACHE | + (cache_policy == L2_BYPASS ? SI_CONTEXT_INV_L2 : 0); } static bool si_is_buffer_idle(struct si_context *sctx, struct si_resource *buf, @@ -139,12 +134,12 @@ static void si_launch_grid_internal(struct si_context *sctx, const struct pipe_g } void si_launch_grid_internal_ssbos(struct si_context *sctx, struct pipe_grid_info *info, - void *shader, unsigned flags, enum si_coherency coher, - unsigned num_buffers, const struct pipe_shader_buffer *buffers, + void *shader, unsigned flags, unsigned num_buffers, + const struct pipe_shader_buffer *buffers, unsigned writeable_bitmask) { if (!(flags & SI_OP_SKIP_CACHE_INV_BEFORE)) { - sctx->flags |= si_get_flush_flags(sctx, coher, L2_LRU); + sctx->flags |= si_get_flush_flags(sctx, L2_LRU); si_mark_atom_dirty(sctx, &sctx->atoms.s.cache_flush); } @@ -166,19 +161,16 @@ void si_launch_grid_internal_ssbos(struct si_context *sctx, struct pipe_grid_inf true /* don't update bind_history to prevent unnecessary syncs later */); si_launch_grid_internal(sctx, info, shader, flags); - /* Do additional cache flushing if needed: - * - CP, CB, DB don't use L2 on GFX6-8. If the coherency is not "shader", flush L2 now. - * - CP doesn't use L2 on GFX12. + /* We must set TC_L2_dirty because: + * - GFX6,12: CP DMA doesn't use L2. + * - GFX6-7,12: Index buffer reads don't use L2. + * - GFX6-8,12: CP doesn't use L2. + * - GFX6-8: CB/DB don't use L2. * - * Set TC_L2_dirty if not flushing now. + * TC_L2_dirty is checked explicitly when buffers are used in those cases to enforce coherency. */ - if (flags & SI_OP_SYNC_AFTER && coher != SI_COHERENCY_SHADER && sctx->gfx_level <= GFX8) { - sctx->flags |= SI_CONTEXT_WB_L2; - si_mark_atom_dirty(sctx, &sctx->atoms.s.cache_flush); - } else { - while (writeable_bitmask) - si_resource(buffers[u_bit_scan(&writeable_bitmask)].buffer)->TC_L2_dirty = true; - } + while (writeable_bitmask) + si_resource(buffers[u_bit_scan(&writeable_bitmask)].buffer)->TC_L2_dirty = true; /* Restore states. */ sctx->b.set_shader_buffers(&sctx->b, PIPE_SHADER_COMPUTE, 0, num_buffers, saved_sb, @@ -209,9 +201,8 @@ set_work_size(struct pipe_grid_info *info, unsigned block_x, unsigned block_y, u * The clear value has 32 bits. */ void si_compute_clear_buffer_rmw(struct si_context *sctx, struct pipe_resource *dst, - unsigned dst_offset, unsigned size, - uint32_t clear_value, uint32_t writebitmask, - unsigned flags, enum si_coherency coher) + unsigned dst_offset, unsigned size, uint32_t clear_value, + uint32_t writebitmask, unsigned flags) { assert(dst_offset % 4 == 0); assert(size % 4 == 0); @@ -236,8 +227,7 @@ void si_compute_clear_buffer_rmw(struct si_context *sctx, struct pipe_resource * if (!sctx->cs_clear_buffer_rmw) sctx->cs_clear_buffer_rmw = si_create_clear_buffer_rmw_cs(sctx); - si_launch_grid_internal_ssbos(sctx, &info, sctx->cs_clear_buffer_rmw, flags, coher, - 1, &sb, 0x1); + si_launch_grid_internal_ssbos(sctx, &info, sctx->cs_clear_buffer_rmw, flags, 1, &sb, 0x1); } /** @@ -256,8 +246,7 @@ bool si_compute_clear_copy_buffer(struct si_context *sctx, struct pipe_resource unsigned dst_offset, struct pipe_resource *src, unsigned src_offset, unsigned size, const uint32_t *clear_value, unsigned clear_value_size, - unsigned flags, enum si_coherency coher, - unsigned dwords_per_thread, bool fail_if_slow) + unsigned flags, unsigned dwords_per_thread, bool fail_if_slow) { assert(dst->target != PIPE_BUFFER || dst_offset + size <= dst->width0); assert(!src || src_offset + size <= src->width0); @@ -312,15 +301,14 @@ bool si_compute_clear_copy_buffer(struct si_context *sctx, struct pipe_resource struct pipe_grid_info grid = {}; set_work_size(&grid, dispatch.workgroup_size, 1, 1, dispatch.num_threads, 1, 1); - si_launch_grid_internal_ssbos(sctx, &grid, shader, flags, coher, dispatch.num_ssbos, sb, + si_launch_grid_internal_ssbos(sctx, &grid, shader, flags, dispatch.num_ssbos, sb, is_copy ? 0x2 : 0x1); return true; } void si_clear_buffer(struct si_context *sctx, struct pipe_resource *dst, uint64_t offset, uint64_t size, uint32_t *clear_value, - uint32_t clear_value_size, unsigned flags, - enum si_coherency coher, enum si_clear_method method) + uint32_t clear_value_size, unsigned flags, enum si_clear_method method) { if (!size) return; @@ -341,7 +329,7 @@ void si_clear_buffer(struct si_context *sctx, struct pipe_resource *dst, if (method != SI_CP_DMA_CLEAR_METHOD && si_compute_clear_copy_buffer(sctx, dst, offset, NULL, 0, size, clear_value, - clear_value_size, flags, coher, 0, + clear_value_size, flags, 0, method == SI_AUTO_SELECT_CLEAR_METHOD)) return; @@ -349,8 +337,7 @@ void si_clear_buffer(struct si_context *sctx, struct pipe_resource *dst, if (aligned_size) { assert(clear_value_size == 4); assert(!(flags & SI_OP_CS_RENDER_COND_ENABLE)); - si_cp_dma_clear_buffer(sctx, &sctx->gfx_cs, dst, offset, aligned_size, *clear_value, - flags, coher); + si_cp_dma_clear_buffer(sctx, &sctx->gfx_cs, dst, offset, aligned_size, *clear_value, flags); } offset += aligned_size; @@ -378,8 +365,7 @@ static void si_pipe_clear_buffer(struct pipe_context *ctx, struct pipe_resource int clear_value_size) { si_clear_buffer((struct si_context *)ctx, dst, offset, size, (uint32_t *)clear_value, - clear_value_size, SI_OP_SYNC_BEFORE_AFTER, SI_COHERENCY_SHADER, - SI_AUTO_SELECT_CLEAR_METHOD); + clear_value_size, SI_OP_SYNC_BEFORE_AFTER, SI_AUTO_SELECT_CLEAR_METHOD); } void si_copy_buffer(struct si_context *sctx, struct pipe_resource *dst, struct pipe_resource *src, @@ -388,15 +374,13 @@ void si_copy_buffer(struct si_context *sctx, struct pipe_resource *dst, struct p if (!size) return; - enum si_coherency coher = SI_COHERENCY_SHADER; - si_improve_sync_flags(sctx, dst, src, &flags); if (si_compute_clear_copy_buffer(sctx, dst, dst_offset, src, src_offset, size, NULL, 0, flags, - coher, 0, true)) + 0, true)) return; - si_cp_dma_copy_buffer(sctx, dst, src, dst_offset, src_offset, size, flags, coher); + si_cp_dma_copy_buffer(sctx, dst, src, dst_offset, src_offset, size, flags); } void si_compute_shorten_ubyte_buffer(struct si_context *sctx, struct pipe_resource *dst, struct pipe_resource *src, @@ -422,8 +406,7 @@ void si_compute_shorten_ubyte_buffer(struct si_context *sctx, struct pipe_resour sb[1].buffer_offset = src_offset; sb[1].buffer_size = count; - si_launch_grid_internal_ssbos(sctx, &info, sctx->cs_ubyte_to_ushort, flags, SI_COHERENCY_SHADER, - 2, sb, 0x1); + si_launch_grid_internal_ssbos(sctx, &info, sctx->cs_ubyte_to_ushort, flags, 2, sb, 0x1); } static void si_launch_grid_internal_images(struct si_context *sctx, @@ -542,14 +525,13 @@ void si_retile_dcc(struct si_context *sctx, struct si_texture *tex) struct pipe_grid_info info = {}; set_work_size(&info, 8, 8, 1, width, height, 1); - si_launch_grid_internal_ssbos(sctx, &info, *shader, SI_OP_SYNC_BEFORE, - SI_COHERENCY_SHADER, 1, &sb, 0x1); + si_launch_grid_internal_ssbos(sctx, &info, *shader, SI_OP_SYNC_BEFORE, 1, &sb, 0x1); /* Don't flush caches. L2 will be flushed by the kernel fence. */ } void gfx9_clear_dcc_msaa(struct si_context *sctx, struct pipe_resource *res, uint32_t clear_value, - unsigned flags, enum si_coherency coher) + unsigned flags) { struct si_texture *tex = (struct si_texture*)res; @@ -588,7 +570,7 @@ void gfx9_clear_dcc_msaa(struct si_context *sctx, struct pipe_resource *res, uin struct pipe_grid_info info = {}; set_work_size(&info, 8, 8, 1, width, height, depth); - si_launch_grid_internal_ssbos(sctx, &info, *shader, flags, coher, 1, &sb, 0x1); + si_launch_grid_internal_ssbos(sctx, &info, *shader, flags, 1, &sb, 0x1); } /* Expand FMASK to make it identity, so that image stores can ignore it. */ @@ -656,7 +638,7 @@ void si_compute_expand_fmask(struct pipe_context *ctx, struct pipe_resource *tex si_clear_buffer(sctx, tex, stex->surface.fmask_offset, stex->surface.fmask_size, (uint32_t *)&fmask_expand_values[log_fragments][log_samples - 1], log_fragments >= 2 && log_samples == 4 ? 8 : 4, SI_OP_SYNC_AFTER, - SI_COHERENCY_SHADER, SI_AUTO_SELECT_CLEAR_METHOD); + SI_AUTO_SELECT_CLEAR_METHOD); } void si_compute_clear_image_dcc_single(struct si_context *sctx, struct si_texture *tex, diff --git a/src/gallium/drivers/radeonsi/si_cp_dma.c b/src/gallium/drivers/radeonsi/si_cp_dma.c index 2dc2f249bdd..140eeb7e496 100644 --- a/src/gallium/drivers/radeonsi/si_cp_dma.c +++ b/src/gallium/drivers/radeonsi/si_cp_dma.c @@ -21,12 +21,9 @@ #define CP_DMA_PFP_SYNC_ME (1 << 4) #define CP_DMA_SRC_IS_GDS (1 << 5) -static enum si_cache_policy get_cache_policy(struct si_context *sctx, enum si_coherency coher) +static enum si_cache_policy get_cache_policy(struct si_context *sctx) { - if (sctx->gfx_level >= GFX7 && coher == SI_COHERENCY_SHADER) - return L2_LRU; - - return L2_BYPASS; + return sctx->gfx_level >= GFX7 ? L2_LRU : L2_BYPASS; } /* The max number of bytes that can be copied per packet. */ @@ -148,7 +145,7 @@ void si_cp_dma_wait_for_idle(struct si_context *sctx, struct radeon_cmdbuf *cs) static void si_cp_dma_prepare(struct si_context *sctx, struct pipe_resource *dst, struct pipe_resource *src, unsigned byte_count, - uint64_t remaining_size, unsigned user_flags, enum si_coherency coher, + uint64_t remaining_size, unsigned user_flags, bool *is_first, unsigned *packet_flags) { if (!(user_flags & SI_OP_CPDMA_SKIP_CHECK_CS_SPACE)) @@ -176,22 +173,18 @@ static void si_cp_dma_prepare(struct si_context *sctx, struct pipe_resource *dst /* Do the synchronization after the last dma, so that all data * is written to memory. */ - if (user_flags & SI_OP_SYNC_AFTER && byte_count == remaining_size) { - *packet_flags |= CP_DMA_SYNC; - - if (coher == SI_COHERENCY_SHADER) - *packet_flags |= CP_DMA_PFP_SYNC_ME; - } + if (user_flags & SI_OP_SYNC_AFTER && byte_count == remaining_size) + *packet_flags |= CP_DMA_SYNC | CP_DMA_PFP_SYNC_ME; } void si_cp_dma_clear_buffer(struct si_context *sctx, struct radeon_cmdbuf *cs, struct pipe_resource *dst, uint64_t offset, uint64_t size, - unsigned value, unsigned user_flags, enum si_coherency coher) + unsigned value, unsigned user_flags) { struct si_resource *sdst = si_resource(dst); uint64_t va = (sdst ? sdst->gpu_address : 0) + offset; bool is_first = true; - enum si_cache_policy cache_policy = get_cache_policy(sctx, coher); + enum si_cache_policy cache_policy = get_cache_policy(sctx); assert(!sctx->screen->info.cp_sdma_ge_use_system_memory_scope); assert(size && size % 4 == 0); @@ -212,7 +205,7 @@ void si_cp_dma_clear_buffer(struct si_context *sctx, struct radeon_cmdbuf *cs, util_range_add(dst, &sdst->valid_buffer_range, offset, offset + size); if (!(user_flags & SI_OP_SKIP_CACHE_INV_BEFORE)) - sctx->flags |= si_get_flush_flags(sctx, coher, cache_policy); + sctx->flags |= si_get_flush_flags(sctx, cache_policy); } if (sctx->flags) @@ -233,8 +226,7 @@ void si_cp_dma_clear_buffer(struct si_context *sctx, struct radeon_cmdbuf *cs, if (!byte_count) continue; - si_cp_dma_prepare(sctx, dst, NULL, byte_count, size, user_flags, coher, &is_first, - &dma_flags); + si_cp_dma_prepare(sctx, dst, NULL, byte_count, size, user_flags, &is_first, &dma_flags); /* Emit the clear packet. */ si_emit_cp_dma(sctx, cs, va, value, byte_count, dma_flags, cache_policy); @@ -246,9 +238,7 @@ void si_cp_dma_clear_buffer(struct si_context *sctx, struct radeon_cmdbuf *cs, if (sdst && cache_policy != L2_BYPASS) sdst->TC_L2_dirty = true; - /* If it's not a framebuffer fast clear... */ - if (coher == SI_COHERENCY_SHADER) - sctx->num_cp_dma_calls++; + sctx->num_cp_dma_calls++; } /** @@ -258,8 +248,7 @@ void si_cp_dma_clear_buffer(struct si_context *sctx, struct radeon_cmdbuf *cs, * \param size Remaining size to the CP DMA alignment. */ static void si_cp_dma_realign_engine(struct si_context *sctx, unsigned size, unsigned user_flags, - enum si_coherency coher, enum si_cache_policy cache_policy, - bool *is_first) + enum si_cache_policy cache_policy, bool *is_first) { uint64_t va; unsigned dma_flags = 0; @@ -283,7 +272,7 @@ static void si_cp_dma_realign_engine(struct si_context *sctx, unsigned size, uns } si_cp_dma_prepare(sctx, &sctx->scratch_buffer->b.b, &sctx->scratch_buffer->b.b, size, size, - user_flags, coher, is_first, &dma_flags); + user_flags, is_first, &dma_flags); va = sctx->scratch_buffer->gpu_address; si_emit_cp_dma(sctx, &sctx->gfx_cs, va, va + SI_CPDMA_ALIGNMENT, size, dma_flags, cache_policy); @@ -297,14 +286,14 @@ static void si_cp_dma_realign_engine(struct si_context *sctx, unsigned size, uns */ void si_cp_dma_copy_buffer(struct si_context *sctx, struct pipe_resource *dst, struct pipe_resource *src, uint64_t dst_offset, uint64_t src_offset, - unsigned size, unsigned user_flags, enum si_coherency coher) + unsigned size, unsigned user_flags) { uint64_t main_dst_offset, main_src_offset; unsigned skipped_size = 0; unsigned realign_size = 0; unsigned gds_flags = (dst ? 0 : CP_DMA_DST_IS_GDS) | (src ? 0 : CP_DMA_SRC_IS_GDS); bool is_first = true; - enum si_cache_policy cache_policy = get_cache_policy(sctx, coher); + enum si_cache_policy cache_policy = get_cache_policy(sctx); assert(!sctx->screen->info.cp_sdma_ge_use_system_memory_scope); assert(size); @@ -366,7 +355,7 @@ void si_cp_dma_copy_buffer(struct si_context *sctx, struct pipe_resource *dst, sctx->flags |= SI_CONTEXT_PS_PARTIAL_FLUSH; if ((dst || src) && !(user_flags & SI_OP_SKIP_CACHE_INV_BEFORE)) - sctx->flags |= si_get_flush_flags(sctx, coher, cache_policy); + sctx->flags |= si_get_flush_flags(sctx, cache_policy); if (sctx->flags) si_mark_atom_dirty(sctx, &sctx->atoms.s.cache_flush); @@ -401,7 +390,7 @@ void si_cp_dma_copy_buffer(struct si_context *sctx, struct pipe_resource *dst, continue; si_cp_dma_prepare(sctx, dst, src, byte_count, size + skipped_size + realign_size, user_flags, - coher, &is_first, &dma_flags); + &is_first, &dma_flags); si_emit_cp_dma(sctx, &sctx->gfx_cs, main_dst_offset, main_src_offset, byte_count, dma_flags, cache_policy); @@ -416,16 +405,15 @@ void si_cp_dma_copy_buffer(struct si_context *sctx, struct pipe_resource *dst, unsigned dma_flags = gds_flags; si_cp_dma_prepare(sctx, dst, src, skipped_size, skipped_size + realign_size, user_flags, - coher, &is_first, &dma_flags); + &is_first, &dma_flags); si_emit_cp_dma(sctx, &sctx->gfx_cs, dst_offset, src_offset, skipped_size, dma_flags, cache_policy); } /* Finally, realign the engine if the size wasn't aligned. */ - if (realign_size) { - si_cp_dma_realign_engine(sctx, realign_size, user_flags, coher, cache_policy, &is_first); - } + if (realign_size) + si_cp_dma_realign_engine(sctx, realign_size, user_flags, cache_policy, &is_first); if (dst && cache_policy != L2_BYPASS) si_resource(dst)->TC_L2_dirty = true; diff --git a/src/gallium/drivers/radeonsi/si_cp_reg_shadowing.c b/src/gallium/drivers/radeonsi/si_cp_reg_shadowing.c index 6130f4090a5..ddac2a06cdf 100644 --- a/src/gallium/drivers/radeonsi/si_cp_reg_shadowing.c +++ b/src/gallium/drivers/radeonsi/si_cp_reg_shadowing.c @@ -58,8 +58,7 @@ void si_init_cp_reg_shadowing(struct si_context *sctx) if (sctx->shadowing.registers) { /* We need to clear the shadowed reg buffer. */ si_cp_dma_clear_buffer(sctx, &sctx->gfx_cs, &sctx->shadowing.registers->b.b, - 0, sctx->shadowing.registers->bo_size, 0, SI_OP_SYNC_AFTER, - SI_COHERENCY_SHADER); + 0, sctx->shadowing.registers->bo_size, 0, SI_OP_SYNC_AFTER); /* Create the shadowing preamble. (allocate enough dwords because the preamble is large) */ struct si_pm4_state *shadowing_preamble = si_pm4_create_sized(sctx->screen, 256, false); diff --git a/src/gallium/drivers/radeonsi/si_pipe.c b/src/gallium/drivers/radeonsi/si_pipe.c index 4f6dbbf1b28..3b9fdb2a98f 100644 --- a/src/gallium/drivers/radeonsi/si_pipe.c +++ b/src/gallium/drivers/radeonsi/si_pipe.c @@ -830,8 +830,7 @@ static struct pipe_context *si_create_context(struct pipe_screen *screen, unsign */ uint32_t clear_value = 0; si_clear_buffer(sctx, sctx->null_const_buf.buffer, 0, sctx->null_const_buf.buffer->width0, - &clear_value, 4, SI_OP_SYNC_AFTER, SI_COHERENCY_SHADER, - SI_CP_DMA_CLEAR_METHOD); + &clear_value, 4, SI_OP_SYNC_AFTER, SI_CP_DMA_CLEAR_METHOD); } if (!(flags & SI_CONTEXT_FLAG_AUX)) { @@ -1088,7 +1087,7 @@ static void si_test_vmfault(struct si_screen *sscreen, uint64_t test_flags) si_resource(buf)->gpu_address = 0; /* cause a VM fault */ if (test_flags & DBG(TEST_VMFAULT_CP)) { - si_cp_dma_copy_buffer(sctx, buf, buf, 0, 4, 4, SI_OP_SYNC_BEFORE_AFTER, SI_COHERENCY_SHADER); + si_cp_dma_copy_buffer(sctx, buf, buf, 0, 4, 4, SI_OP_SYNC_BEFORE_AFTER); ctx->flush(ctx, NULL, 0); puts("VM fault test: CP - done."); } diff --git a/src/gallium/drivers/radeonsi/si_pipe.h b/src/gallium/drivers/radeonsi/si_pipe.h index 5e726d34763..0348de7d865 100644 --- a/src/gallium/drivers/radeonsi/si_pipe.h +++ b/src/gallium/drivers/radeonsi/si_pipe.h @@ -270,11 +270,6 @@ enum si_cache_policy L2_LRU, /* same as SLC=0 */ }; -enum si_coherency -{ - SI_COHERENCY_SHADER, -}; - #define SI_BIND_CONSTANT_BUFFER_SHIFT 0 #define SI_BIND_SHADER_BUFFER_SHIFT 6 #define SI_BIND_IMAGE_BUFFER_SHIFT 12 @@ -1486,19 +1481,17 @@ void si_destroy_compute(struct si_compute *program); #define SI_OP_FAIL_IF_SLOW (1 << 9) #define SI_OP_IS_NESTED (1 << 10) -unsigned si_get_flush_flags(struct si_context *sctx, enum si_coherency coher, - enum si_cache_policy cache_policy); +unsigned si_get_flush_flags(struct si_context *sctx, enum si_cache_policy cache_policy); bool si_should_blit_clamp_to_edge(const struct pipe_blit_info *info, unsigned coord_mask); void si_launch_grid_internal_ssbos(struct si_context *sctx, struct pipe_grid_info *info, - void *shader, unsigned flags, enum si_coherency coher, - unsigned num_buffers, const struct pipe_shader_buffer *buffers, + void *shader, unsigned flags, unsigned num_buffers, + const struct pipe_shader_buffer *buffers, unsigned writeable_bitmask); bool si_compute_clear_copy_buffer(struct si_context *sctx, struct pipe_resource *dst, unsigned dst_offset, struct pipe_resource *src, unsigned src_offset, unsigned size, const uint32_t *clear_value, unsigned clear_value_size, - unsigned flags, enum si_coherency coher, - unsigned dwords_per_thread, bool fail_if_slow); + unsigned flags, unsigned dwords_per_thread, bool fail_if_slow); enum si_clear_method { SI_CP_DMA_CLEAR_METHOD, SI_COMPUTE_CLEAR_METHOD, @@ -1507,11 +1500,10 @@ enum si_clear_method { void si_clear_buffer(struct si_context *sctx, struct pipe_resource *dst, uint64_t offset, uint64_t size, uint32_t *clear_value, uint32_t clear_value_size, unsigned flags, - enum si_coherency coher, enum si_clear_method method); + enum si_clear_method method); void si_compute_clear_buffer_rmw(struct si_context *sctx, struct pipe_resource *dst, - unsigned dst_offset, unsigned size, - uint32_t clear_value, uint32_t writebitmask, - unsigned flags, enum si_coherency coher); + unsigned dst_offset, unsigned size, uint32_t clear_value, + uint32_t writebitmask, unsigned flags); void si_copy_buffer(struct si_context *sctx, struct pipe_resource *dst, struct pipe_resource *src, uint64_t dst_offset, uint64_t src_offset, unsigned size, unsigned flags); void si_compute_shorten_ubyte_buffer(struct si_context *sctx, struct pipe_resource *dst, struct pipe_resource *src, @@ -1521,7 +1513,7 @@ void si_compute_clear_image_dcc_single(struct si_context *sctx, struct si_textur const union pipe_color_union *color, unsigned flags); void si_retile_dcc(struct si_context *sctx, struct si_texture *tex); void gfx9_clear_dcc_msaa(struct si_context *sctx, struct pipe_resource *res, uint32_t clear_value, - unsigned flags, enum si_coherency coher); + unsigned flags); void si_compute_expand_fmask(struct pipe_context *ctx, struct pipe_resource *tex); bool si_compute_clear_image(struct si_context *sctx, struct pipe_resource *tex, enum pipe_format format, unsigned level, const struct pipe_box *box, @@ -1540,10 +1532,10 @@ void si_init_compute_blit_functions(struct si_context *sctx); void si_cp_dma_wait_for_idle(struct si_context *sctx, struct radeon_cmdbuf *cs); void si_cp_dma_clear_buffer(struct si_context *sctx, struct radeon_cmdbuf *cs, struct pipe_resource *dst, uint64_t offset, uint64_t size, - unsigned value, unsigned user_flags, enum si_coherency coher); + unsigned value, unsigned user_flags); void si_cp_dma_copy_buffer(struct si_context *sctx, struct pipe_resource *dst, struct pipe_resource *src, uint64_t dst_offset, uint64_t src_offset, - unsigned size, unsigned user_flags, enum si_coherency coher); + unsigned size, unsigned user_flags); void si_cp_write_data(struct si_context *sctx, struct si_resource *buf, unsigned offset, unsigned size, unsigned dst_sel, unsigned engine, const void *data); void si_cp_copy_data(struct si_context *sctx, struct radeon_cmdbuf *cs, unsigned dst_sel, diff --git a/src/gallium/drivers/radeonsi/si_query.c b/src/gallium/drivers/radeonsi/si_query.c index 7b5f072c2d6..316a7bfa83e 100644 --- a/src/gallium/drivers/radeonsi/si_query.c +++ b/src/gallium/drivers/radeonsi/si_query.c @@ -1645,8 +1645,7 @@ static void si_query_hw_get_result_resource(struct si_context *sctx, struct si_q si_cp_wait_mem(sctx, &sctx->gfx_cs, va, 0x80000000, 0x80000000, WAIT_REG_MEM_EQUAL); } si_launch_grid_internal_ssbos(sctx, &grid, sctx->query_result_shader, - SI_OP_SYNC_AFTER, SI_COHERENCY_SHADER, - 3, ssbo, 0x4); + SI_OP_SYNC_AFTER, 3, ssbo, 0x4); } si_restore_qbo_state(sctx, &saved_state); diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index 14ed0515895..6bf55088ae2 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -982,7 +982,7 @@ static void post_upload_binary(struct si_screen *sscreen, struct si_shader *shad * them available. */ si_cp_dma_copy_buffer(upload_ctx, &shader->bo->b.b, staging, 0, staging_offset, - binary_size, SI_OP_SYNC_AFTER, SI_COHERENCY_SHADER); + binary_size, SI_OP_SYNC_AFTER); upload_ctx->flags |= SI_CONTEXT_INV_ICACHE | SI_CONTEXT_INV_L2; #if 0 /* debug: validate whether the copy was successful */ diff --git a/src/gallium/drivers/radeonsi/si_test_dma_perf.c b/src/gallium/drivers/radeonsi/si_test_dma_perf.c index d824b74b9b9..3549a32e0ed 100644 --- a/src/gallium/drivers/radeonsi/si_test_dma_perf.c +++ b/src/gallium/drivers/radeonsi/si_test_dma_perf.c @@ -228,7 +228,7 @@ void si_test_dma_perf(struct si_screen *sscreen) continue; } si_cp_dma_copy_buffer(sctx, dst, src, dst_offset, src_offset, size, - SI_OP_SYNC_BEFORE_AFTER, SI_COHERENCY_SHADER); + SI_OP_SYNC_BEFORE_AFTER); } else { /* CP DMA clears must be aligned to 4 bytes. */ if (dst_offset % 4 || size % 4 || @@ -239,16 +239,15 @@ void si_test_dma_perf(struct si_screen *sscreen) } assert(clear_value_size == 4); si_cp_dma_clear_buffer(sctx, &sctx->gfx_cs, dst, dst_offset, size, - clear_value[0], SI_OP_SYNC_BEFORE_AFTER, - SI_COHERENCY_SHADER); + clear_value[0], SI_OP_SYNC_BEFORE_AFTER); } } else { /* Compute */ success &= si_compute_clear_copy_buffer(sctx, dst, dst_offset, src, src_offset, size, clear_value, clear_value_size, - SI_OP_SYNC_BEFORE_AFTER, SI_COHERENCY_SHADER, - dwords_per_thread, false); + SI_OP_SYNC_BEFORE_AFTER, dwords_per_thread, + false); } sctx->flags |= SI_CONTEXT_INV_L2; @@ -478,8 +477,7 @@ void si_test_clear_buffer(struct si_screen *sscreen) bool done = si_compute_clear_copy_buffer(sctx, dst, dst_offset, NULL, 0, op_size, (uint32_t*)clear_value, clear_value_size, - SI_OP_SYNC_BEFORE_AFTER, SI_COHERENCY_SHADER, - dwords_per_thread, false); + SI_OP_SYNC_BEFORE_AFTER, dwords_per_thread, false); if (done) { pipe_buffer_read(ctx, dst, 0, buf_size, read_dst_buffer); @@ -584,7 +582,7 @@ void si_test_copy_buffer(struct si_screen *sscreen) bool done = si_compute_clear_copy_buffer(sctx, dst, dst_offset, src, src_offset, op_size, NULL, 0, SI_OP_SYNC_BEFORE_AFTER, - SI_COHERENCY_SHADER, dwords_per_thread, false); + dwords_per_thread, false); if (done) { pipe_buffer_read(ctx, dst, 0, buf_size, read_dst_buffer); diff --git a/src/gallium/drivers/radeonsi/si_test_image_copy_region.c b/src/gallium/drivers/radeonsi/si_test_image_copy_region.c index 3e8003aa46c..198a33ebdaf 100644 --- a/src/gallium/drivers/radeonsi/si_test_image_copy_region.c +++ b/src/gallium/drivers/radeonsi/si_test_image_copy_region.c @@ -538,7 +538,7 @@ void si_test_image_copy_region(struct si_screen *sscreen) /* clear dst pixels */ uint32_t zero = 0; si_clear_buffer(sctx, dst, 0, sdst->surface.surf_size, &zero, 4, SI_OP_SYNC_BEFORE_AFTER, - SI_COHERENCY_SHADER, SI_AUTO_SELECT_CLEAR_METHOD); + SI_AUTO_SELECT_CLEAR_METHOD); for (j = 0; j < num_partial_copies; j++) { int width, height, depth; @@ -716,9 +716,9 @@ void si_test_blit(struct si_screen *sscreen, unsigned test_flags) /* clear dst pixels */ uint32_t zero = 0; si_clear_buffer(sctx, gfx_dst, 0, ((struct si_texture *)gfx_dst)->surface.surf_size, &zero, - 4, SI_OP_SYNC_BEFORE_AFTER, SI_COHERENCY_SHADER, SI_AUTO_SELECT_CLEAR_METHOD); + 4, SI_OP_SYNC_BEFORE_AFTER, SI_AUTO_SELECT_CLEAR_METHOD); si_clear_buffer(sctx, comp_dst, 0, ((struct si_texture *)comp_dst)->surface.surf_size, &zero, - 4, SI_OP_SYNC_BEFORE_AFTER, SI_COHERENCY_SHADER, SI_AUTO_SELECT_CLEAR_METHOD); + 4, SI_OP_SYNC_BEFORE_AFTER, SI_AUTO_SELECT_CLEAR_METHOD); /* TODO: These two fix quite a lot of BCn cases. */ /*si_clear_buffer(sctx, gfx_src, 0, ((struct si_texture *)gfx_src)->surface.surf_size, &zero,