AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

radeonsi: remove enum si_coherency

Browse Source 
SI_COHERENCY_SHADER is the only one that's used.

Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/31168>
This commit is contained in:
Marek Olšák
2024-08-19 21:43:58 -04:00
parent a009aa7365
commit cfe1f2b841
12 changed files with 78 additions and 124 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/gallium/drivers/radeonsi/gfx11_query.c
+1 -1
View File
@@ -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)
src/gallium/drivers/radeonsi/si_buffer.c
+1 -1
View File
@@ -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);
}
src/gallium/drivers/radeonsi/si_clear.c
+3 -6
View File
@@ -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);
}
}
src/gallium/drivers/radeonsi/si_compute_blit.c
+30 -48
View File
@@ -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,
src/gallium/drivers/radeonsi/si_cp_dma.c
+19 -31
View File
@@ -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;
src/gallium/drivers/radeonsi/si_cp_reg_shadowing.c
+1 -2
View File
@@ -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);
src/gallium/drivers/radeonsi/si_pipe.c
+2 -3
View File
@@ -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.");
}
src/gallium/drivers/radeonsi/si_pipe.h
+10 -18
View File
@@ -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,
src/gallium/drivers/radeonsi/si_query.c
+1 -2
View File
@@ -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);
src/gallium/drivers/radeonsi/si_shader.c
+1 -1
View File
@@ -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 */
src/gallium/drivers/radeonsi/si_test_dma_perf.c
+6 -8
View File
@@ -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);
src/gallium/drivers/radeonsi/si_test_image_copy_region.c
+3 -3
View File
@@ -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,