diff --git a/RHEL-53423.patch b/RHEL-53423.patch new file mode 100644 index 0000000..4a9c256 --- /dev/null +++ b/RHEL-53423.patch @@ -0,0 +1,2565 @@ +From 9797443fc2d655b8333ab80aec9d17a06706fe53 Mon Sep 17 00:00:00 2001 +From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= +Date: Wed, 20 Nov 2024 10:34:27 -0500 +Subject: [PATCH 01/27] ac/surface: adjust HiZ enablement + +Acked-by: Pierre-Eric Pelloux-Prayer +Part-of: +--- + src/amd/common/ac_surface.c | 2 +- + 1 file changed, 1 insertion(+), 1 deletion(-) + +diff --git a/src/amd/common/ac_surface.c b/src/amd/common/ac_surface.c +index 3ecb5d57be4..69f722bc4cc 100644 +--- a/src/amd/common/ac_surface.c ++++ b/src/amd/common/ac_surface.c +@@ -3022,7 +3022,7 @@ static bool gfx12_compute_hiz_his_info(struct ac_addrlib *addrlib, const struct + { + assert(surf_in->flags.depth != surf_in->flags.stencil); + +- if (surf->flags & RADEON_SURF_NO_HTILE || (info->gfx_level == GFX12 && info->chip_rev == 0)) ++ if (surf->flags & RADEON_SURF_NO_HTILE || (info->gfx_level == GFX12 && info->chip_rev <= 1)) + return true; + + ADDR3_COMPUTE_SURFACE_INFO_OUTPUT out = {0}; +-- +2.48.1 + + +From 51f92a3545b42d4250da95e9623576b65e8260b9 Mon Sep 17 00:00:00 2001 +From: Pierre-Eric Pelloux-Prayer +Date: Thu, 21 Nov 2024 11:46:02 +0100 +Subject: [PATCH 02/27] radeonsi/gfx12: disable display dcc for front buffer + rendering +MIME-Version: 1.0 +Content-Type: text/plain; charset=UTF-8 +Content-Transfer-Encoding: 8bit + +Same logic as other chips, except we need to reallocate the texture +as we can't disable dcc. + +Reviewed-by: Marek Olšák +Part-of: +--- + src/gallium/drivers/radeonsi/si_texture.c | 33 ++++++++++++++++------- + 1 file changed, 24 insertions(+), 9 deletions(-) + +diff --git a/src/gallium/drivers/radeonsi/si_texture.c b/src/gallium/drivers/radeonsi/si_texture.c +index ad234b2f018..ae65b725600 100644 +--- a/src/gallium/drivers/radeonsi/si_texture.c ++++ b/src/gallium/drivers/radeonsi/si_texture.c +@@ -797,16 +797,31 @@ static bool si_texture_get_handle(struct pipe_screen *screen, struct pipe_contex + assert(tex->surface.tile_swizzle == 0); + } + +- /* Since shader image stores don't support DCC on GFX8, +- * disable it for external clients that want write +- * access. ++ const bool debug_disable_dcc = sscreen->debug_flags & DBG(NO_EXPORTED_DCC); ++ /* Since shader image stores don't support DCC on GFX9 and older, ++ * disable it for external clients that want write access. + */ +- if (sscreen->debug_flags & DBG(NO_EXPORTED_DCC) || +- (usage & PIPE_HANDLE_USAGE_SHADER_WRITE && !tex->is_depth && tex->surface.meta_offset) || +- /* Displayable DCC requires an explicit flush. */ +- (!(usage & PIPE_HANDLE_USAGE_EXPLICIT_FLUSH) && +- si_displayable_dcc_needs_explicit_flush(tex))) { +- if (si_texture_disable_dcc(sctx, tex)) { ++ const bool shader_write = sscreen->info.gfx_level <= GFX9 && ++ usage & PIPE_HANDLE_USAGE_SHADER_WRITE && ++ !tex->is_depth && ++ tex->surface.meta_offset; ++ /* Another reason to disable display dcc is front buffer rendering. ++ * This can happens with Xorg. If the ddx driver uses GBM_BO_USE_FRONT_RENDERING, ++ * there's nothing to do because the texture is not using DCC. ++ * If the flag isn't set, we have to infer it to get correct rendering. ++ */ ++ const bool front_buffer_rendering = !(usage & PIPE_HANDLE_USAGE_EXPLICIT_FLUSH) && ++ tex->buffer.b.b.bind & PIPE_BIND_SCANOUT; ++ ++ /* If display dcc requires a retiling step, drop dcc. */ ++ const bool explicit_flush = !(usage & PIPE_HANDLE_USAGE_EXPLICIT_FLUSH) && ++ si_displayable_dcc_needs_explicit_flush(tex); ++ ++ if (debug_disable_dcc || shader_write || front_buffer_rendering || explicit_flush) { ++ if (sscreen->info.gfx_level >= GFX12) { ++ si_reallocate_texture_inplace(sctx, tex, PIPE_BIND_CONST_BW, false); ++ update_metadata = true; ++ } else if (si_texture_disable_dcc(sctx, tex)) { + update_metadata = true; + /* si_texture_disable_dcc flushes the context */ + flush = false; +-- +2.48.1 + + +From 89d3f1a550ade706459e6c2f0649608c084387ff Mon Sep 17 00:00:00 2001 +From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= +Date: Mon, 16 Dec 2024 03:48:37 -0500 +Subject: [PATCH 03/27] radeonsi/gfx12: set DB_RENDER_OVERRIDE based on stencil + state + +Acked-by: Pierre-Eric Pelloux-Prayer +Part-of: +--- + src/gallium/drivers/radeonsi/si_gfx_cs.c | 1 + + src/gallium/drivers/radeonsi/si_state.c | 8 ++++++++ + src/gallium/drivers/radeonsi/si_state.h | 2 ++ + 3 files changed, 11 insertions(+) + +diff --git a/src/gallium/drivers/radeonsi/si_gfx_cs.c b/src/gallium/drivers/radeonsi/si_gfx_cs.c +index 9b11eaf678e..31814bc94d2 100644 +--- a/src/gallium/drivers/radeonsi/si_gfx_cs.c ++++ b/src/gallium/drivers/radeonsi/si_gfx_cs.c +@@ -318,6 +318,7 @@ void si_set_tracked_regs_to_clear_state(struct si_context *ctx) + ctx->tracked_regs.reg_value[SI_TRACKED_SPI_PS_INPUT_ENA] = 0; + ctx->tracked_regs.reg_value[SI_TRACKED_SPI_PS_INPUT_ADDR] = 0; + ++ ctx->tracked_regs.reg_value[SI_TRACKED_DB_RENDER_OVERRIDE] = 0; + ctx->tracked_regs.reg_value[SI_TRACKED_DB_EQAA] = 0; + ctx->tracked_regs.reg_value[SI_TRACKED_DB_RENDER_OVERRIDE2] = 0; + ctx->tracked_regs.reg_value[SI_TRACKED_DB_SHADER_CONTROL] = 0; +diff --git a/src/gallium/drivers/radeonsi/si_state.c b/src/gallium/drivers/radeonsi/si_state.c +index 6bb1037a79b..018be3dbe4f 100644 +--- a/src/gallium/drivers/radeonsi/si_state.c ++++ b/src/gallium/drivers/radeonsi/si_state.c +@@ -1573,6 +1573,12 @@ static void *si_create_dsa_state(struct pipe_context *ctx, + S_028090_TESTMASK_BF(state->stencil[1].valuemask); + dsa->db_stencil_write_mask = S_028094_WRITEMASK(state->stencil[0].writemask) | + S_028094_WRITEMASK_BF(state->stencil[1].writemask); ++ ++ bool force_s_valid = state->stencil[0].zpass_op != state->stencil[0].zfail_op || ++ (state->stencil[1].enabled && ++ state->stencil[1].zpass_op != state->stencil[1].zfail_op); ++ dsa->db_render_override = S_02800C_FORCE_STENCIL_READ(1) | ++ S_02800C_FORCE_STENCIL_VALID(force_s_valid); + } + + bool zfunc_is_ordered = +@@ -1608,6 +1614,8 @@ static void si_pm4_emit_dsa(struct si_context *sctx, unsigned index) + if (sctx->gfx_level >= GFX12) { + radeon_begin(&sctx->gfx_cs); + gfx12_begin_context_regs(); ++ gfx12_opt_set_context_reg(R_02800C_DB_RENDER_OVERRIDE, SI_TRACKED_DB_RENDER_OVERRIDE, ++ state->db_render_override); + gfx12_opt_set_context_reg(R_028070_DB_DEPTH_CONTROL, SI_TRACKED_DB_DEPTH_CONTROL, + state->db_depth_control); + if (state->stencil_enabled) { +diff --git a/src/gallium/drivers/radeonsi/si_state.h b/src/gallium/drivers/radeonsi/si_state.h +index a984c7d6918..f964fa21606 100644 +--- a/src/gallium/drivers/radeonsi/si_state.h ++++ b/src/gallium/drivers/radeonsi/si_state.h +@@ -130,6 +130,7 @@ struct si_state_dsa { + unsigned spi_shader_user_data_ps_alpha_ref; + unsigned db_stencil_read_mask; + unsigned db_stencil_write_mask; ++ unsigned db_render_override; /* only gfx12 */ + + /* 0 = without stencil buffer, 1 = when both Z and S buffers are present */ + struct si_dsa_order_invariance order_invariance[2]; +@@ -314,6 +315,7 @@ enum si_tracked_reg + SI_TRACKED_SPI_PS_INPUT_ENA, + SI_TRACKED_SPI_PS_INPUT_ADDR, + ++ SI_TRACKED_DB_RENDER_OVERRIDE, + SI_TRACKED_DB_EQAA, + SI_TRACKED_DB_RENDER_OVERRIDE2, + SI_TRACKED_DB_SHADER_CONTROL, +-- +2.48.1 + + +From 9a0377096ab33b12c10913ecc2b5a1358da29923 Mon Sep 17 00:00:00 2001 +From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= +Date: Tue, 6 Aug 2024 16:23:01 -0400 +Subject: [PATCH 04/27] radeonsi: add a new PM4 helper radeon_event_write + +Reviewed-by: Pierre-Eric Pelloux-Prayer +Part-of: +--- + src/gallium/drivers/radeonsi/si_build_pm4.h | 11 +++ + src/gallium/drivers/radeonsi/si_compute.c | 7 +- + src/gallium/drivers/radeonsi/si_gfx_cs.c | 89 +++++++------------ + src/gallium/drivers/radeonsi/si_perfcounter.c | 21 ++--- + src/gallium/drivers/radeonsi/si_query.c | 3 +- + src/gallium/drivers/radeonsi/si_state.c | 24 +++-- + .../drivers/radeonsi/si_state_draw.cpp | 21 ++--- + .../drivers/radeonsi/si_state_shaders.cpp | 13 +-- + .../drivers/radeonsi/si_state_streamout.c | 3 +- + 9 files changed, 74 insertions(+), 118 deletions(-) + +diff --git a/src/gallium/drivers/radeonsi/si_build_pm4.h b/src/gallium/drivers/radeonsi/si_build_pm4.h +index 15259aa329b..630b5885ff3 100644 +--- a/src/gallium/drivers/radeonsi/si_build_pm4.h ++++ b/src/gallium/drivers/radeonsi/si_build_pm4.h +@@ -503,6 +503,17 @@ + } \ + } while (0) + ++/* Other packet helpers. */ ++#define radeon_event_write(event_type) do { \ ++ unsigned __event_type = (event_type); \ ++ radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); \ ++ radeon_emit(EVENT_TYPE(__event_type) | \ ++ EVENT_INDEX(__event_type == V_028A90_VS_PARTIAL_FLUSH || \ ++ __event_type == V_028A90_PS_PARTIAL_FLUSH || \ ++ __event_type == V_028A90_CS_PARTIAL_FLUSH ? 4 : \ ++ __event_type == V_028A90_PIXEL_PIPE_STAT_CONTROL ? 1 : 0)); \ ++} while (0) ++ + /* This should be evaluated at compile time if all parameters are constants. */ + static ALWAYS_INLINE unsigned + si_get_user_data_base(enum amd_gfx_level gfx_level, enum si_has_tess has_tess, +diff --git a/src/gallium/drivers/radeonsi/si_compute.c b/src/gallium/drivers/radeonsi/si_compute.c +index 727face471e..01aa75d222a 100644 +--- a/src/gallium/drivers/radeonsi/si_compute.c ++++ b/src/gallium/drivers/radeonsi/si_compute.c +@@ -1114,10 +1114,9 @@ static void si_emit_dispatch_packets(struct si_context *sctx, const struct pipe_ + radeon_emit(dispatch_initiator); + } + +- if (unlikely(sctx->sqtt_enabled && sctx->gfx_level >= GFX9)) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_THREAD_TRACE_MARKER) | EVENT_INDEX(0)); +- } ++ if (unlikely(sctx->sqtt_enabled && sctx->gfx_level >= GFX9)) ++ radeon_event_write(V_028A90_THREAD_TRACE_MARKER); ++ + radeon_end(); + } + +diff --git a/src/gallium/drivers/radeonsi/si_gfx_cs.c b/src/gallium/drivers/radeonsi/si_gfx_cs.c +index 31814bc94d2..d8dccd9b008 100644 +--- a/src/gallium/drivers/radeonsi/si_gfx_cs.c ++++ b/src/gallium/drivers/radeonsi/si_gfx_cs.c +@@ -160,8 +160,7 @@ void si_flush_gfx_cs(struct si_context *ctx, unsigned flags, struct pipe_fence_h + */ + if ((ctx->gfx_level == GFX11 || ctx->gfx_level == GFX11_5) && ctx->has_tessellation) { + radeon_begin(cs); +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_SQ_NON_EVENT) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_SQ_NON_EVENT); + radeon_end(); + } + +@@ -797,10 +796,8 @@ void gfx10_emit_cache_flush(struct si_context *ctx, struct radeon_cmdbuf *cs) + + radeon_begin(cs); + +- if (flags & SI_CONTEXT_VGT_FLUSH) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_VGT_FLUSH) | EVENT_INDEX(0)); +- } ++ if (flags & SI_CONTEXT_VGT_FLUSH) ++ radeon_event_write(V_028A90_VGT_FLUSH); + + if (flags & SI_CONTEXT_INV_ICACHE) + gcr_cntl |= S_586_GLI_INV(V_586_GLI_ALL); +@@ -837,19 +834,15 @@ void gfx10_emit_cache_flush(struct si_context *ctx, struct radeon_cmdbuf *cs) + } + + if (flags & (SI_CONTEXT_FLUSH_AND_INV_CB | SI_CONTEXT_FLUSH_AND_INV_DB)) { +- if (ctx->gfx_level < GFX12 && flags & SI_CONTEXT_FLUSH_AND_INV_CB) { +- /* Flush CMASK/FMASK/DCC. Will wait for idle later. */ +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_FLUSH_AND_INV_CB_META) | EVENT_INDEX(0)); +- } ++ /* Flush CMASK/FMASK/DCC. Will wait for idle later. */ ++ if (ctx->gfx_level < GFX12 && flags & SI_CONTEXT_FLUSH_AND_INV_CB) ++ radeon_event_write(V_028A90_FLUSH_AND_INV_CB_META); + + /* Gfx11 can't flush DB_META and should use a TS event instead. */ ++ /* Flush HTILE. Will wait for idle later. */ + if (ctx->gfx_level < GFX12 && ctx->gfx_level != GFX11 && +- flags & SI_CONTEXT_FLUSH_AND_INV_DB) { +- /* Flush HTILE. Will wait for idle later. */ +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_FLUSH_AND_INV_DB_META) | EVENT_INDEX(0)); +- } ++ flags & SI_CONTEXT_FLUSH_AND_INV_DB) ++ radeon_event_write(V_028A90_FLUSH_AND_INV_DB_META); + + /* First flush CB/DB, then L1/L2. */ + gcr_cntl |= S_586_SEQ(V_586_SEQ_FORWARD); +@@ -870,21 +863,18 @@ void gfx10_emit_cache_flush(struct si_context *ctx, struct radeon_cmdbuf *cs) + } else { + /* Wait for graphics shaders to go idle if requested. */ + if (flags & SI_CONTEXT_PS_PARTIAL_FLUSH) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_PS_PARTIAL_FLUSH) | EVENT_INDEX(4)); ++ radeon_event_write(V_028A90_PS_PARTIAL_FLUSH); + /* Only count explicit shader flushes, not implicit ones. */ + ctx->num_vs_flushes++; + ctx->num_ps_flushes++; + } else if (flags & SI_CONTEXT_VS_PARTIAL_FLUSH) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_VS_PARTIAL_FLUSH) | EVENT_INDEX(4)); ++ radeon_event_write(V_028A90_VS_PARTIAL_FLUSH); + ctx->num_vs_flushes++; + } + } + + if (flags & SI_CONTEXT_CS_PARTIAL_FLUSH && ctx->compute_is_busy) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_CS_PARTIAL_FLUSH | EVENT_INDEX(4))); ++ radeon_event_write(V_028A90_CS_PARTIAL_FLUSH); + ctx->num_cs_flushes++; + ctx->compute_is_busy = false; + } +@@ -1026,12 +1016,10 @@ void gfx10_emit_cache_flush(struct si_context *ctx, struct radeon_cmdbuf *cs) + } + + if (flags & SI_CONTEXT_START_PIPELINE_STATS && ctx->pipeline_stats_enabled != 1) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_PIPELINESTAT_START) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_PIPELINESTAT_START); + ctx->pipeline_stats_enabled = 1; + } else if (flags & SI_CONTEXT_STOP_PIPELINE_STATS && ctx->pipeline_stats_enabled != 0) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_PIPELINESTAT_STOP) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_PIPELINESTAT_STOP); + ctx->pipeline_stats_enabled = 0; + } + radeon_end(); +@@ -1092,16 +1080,13 @@ void gfx6_emit_cache_flush(struct si_context *sctx, struct radeon_cmdbuf *cs) + + radeon_begin(cs); + +- if (flags & SI_CONTEXT_FLUSH_AND_INV_CB) { +- /* Flush CMASK/FMASK/DCC. SURFACE_SYNC will wait for idle. */ +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_FLUSH_AND_INV_CB_META) | EVENT_INDEX(0)); +- } +- if (flags & (SI_CONTEXT_FLUSH_AND_INV_DB | SI_CONTEXT_FLUSH_AND_INV_DB_META)) { +- /* Flush HTILE. SURFACE_SYNC will wait for idle. */ +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_FLUSH_AND_INV_DB_META) | EVENT_INDEX(0)); +- } ++ /* Flush CMASK/FMASK/DCC. SURFACE_SYNC will wait for idle. */ ++ if (flags & SI_CONTEXT_FLUSH_AND_INV_CB) ++ radeon_event_write(V_028A90_FLUSH_AND_INV_CB_META); ++ ++ /* Flush HTILE. SURFACE_SYNC will wait for idle. */ ++ if (flags & (SI_CONTEXT_FLUSH_AND_INV_DB | SI_CONTEXT_FLUSH_AND_INV_DB_META)) ++ radeon_event_write(V_028A90_FLUSH_AND_INV_DB_META); + + /* Wait for shader engines to go idle. + * VS and PS waits are unnecessary if SURFACE_SYNC is going to wait +@@ -1109,36 +1094,28 @@ void gfx6_emit_cache_flush(struct si_context *sctx, struct radeon_cmdbuf *cs) + */ + if (!flush_cb_db) { + if (flags & SI_CONTEXT_PS_PARTIAL_FLUSH) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_PS_PARTIAL_FLUSH) | EVENT_INDEX(4)); +- /* Only count explicit shader flushes, not implicit ones +- * done by SURFACE_SYNC. +- */ ++ radeon_event_write(V_028A90_PS_PARTIAL_FLUSH); ++ /* Only count explicit shader flushes, not implicit ones done by SURFACE_SYNC. */ + sctx->num_vs_flushes++; + sctx->num_ps_flushes++; + } else if (flags & SI_CONTEXT_VS_PARTIAL_FLUSH) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_VS_PARTIAL_FLUSH) | EVENT_INDEX(4)); ++ radeon_event_write(V_028A90_VS_PARTIAL_FLUSH); + sctx->num_vs_flushes++; + } + } + + if (flags & SI_CONTEXT_CS_PARTIAL_FLUSH && sctx->compute_is_busy) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_CS_PARTIAL_FLUSH) | EVENT_INDEX(4)); ++ radeon_event_write(V_028A90_CS_PARTIAL_FLUSH); + sctx->num_cs_flushes++; + sctx->compute_is_busy = false; + } + + /* VGT state synchronization. */ +- if (flags & SI_CONTEXT_VGT_FLUSH) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_VGT_FLUSH) | EVENT_INDEX(0)); +- } +- if (flags & SI_CONTEXT_VGT_STREAMOUT_SYNC) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_VGT_STREAMOUT_SYNC) | EVENT_INDEX(0)); +- } ++ if (flags & SI_CONTEXT_VGT_FLUSH) ++ radeon_event_write(V_028A90_VGT_FLUSH); ++ ++ if (flags & SI_CONTEXT_VGT_STREAMOUT_SYNC) ++ radeon_event_write(V_028A90_VGT_STREAMOUT_SYNC); + + radeon_end(); + +@@ -1267,14 +1244,12 @@ void gfx6_emit_cache_flush(struct si_context *sctx, struct radeon_cmdbuf *cs) + + if (flags & SI_CONTEXT_START_PIPELINE_STATS && sctx->pipeline_stats_enabled != 1) { + radeon_begin(cs); +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_PIPELINESTAT_START) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_PIPELINESTAT_START); + radeon_end(); + sctx->pipeline_stats_enabled = 1; + } else if (flags & SI_CONTEXT_STOP_PIPELINE_STATS && sctx->pipeline_stats_enabled != 0) { + radeon_begin(cs); +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_PIPELINESTAT_STOP) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_PIPELINESTAT_STOP); + radeon_end(); + sctx->pipeline_stats_enabled = 0; + } +diff --git a/src/gallium/drivers/radeonsi/si_perfcounter.c b/src/gallium/drivers/radeonsi/si_perfcounter.c +index 9cbd08648af..785c98a9dc9 100644 +--- a/src/gallium/drivers/radeonsi/si_perfcounter.c ++++ b/src/gallium/drivers/radeonsi/si_perfcounter.c +@@ -114,8 +114,7 @@ static void si_pc_emit_start(struct si_context *sctx, struct si_resource *buffer + radeon_begin(cs); + radeon_set_uconfig_reg(R_036020_CP_PERFMON_CNTL, + S_036020_PERFMON_STATE(V_036020_CP_PERFMON_STATE_DISABLE_AND_RESET)); +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_PERFCOUNTER_START) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_PERFCOUNTER_STOP); + radeon_set_uconfig_reg(R_036020_CP_PERFMON_CNTL, + S_036020_PERFMON_STATE(V_036020_CP_PERFMON_STATE_START_COUNTING)); + radeon_end(); +@@ -132,13 +131,10 @@ static void si_pc_emit_stop(struct si_context *sctx, struct si_resource *buffer, + si_cp_wait_mem(sctx, cs, va, 0, 0xffffffff, WAIT_REG_MEM_EQUAL); + + radeon_begin(cs); +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_PERFCOUNTER_SAMPLE) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_PERFCOUNTER_SAMPLE); + +- if (!sctx->screen->info.never_send_perfcounter_stop) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_PERFCOUNTER_STOP) | EVENT_INDEX(0)); +- } ++ if (!sctx->screen->info.never_send_perfcounter_stop) ++ radeon_event_write(V_028A90_PERFCOUNTER_STOP); + + radeon_set_uconfig_reg( + R_036020_CP_PERFMON_CNTL, +@@ -158,8 +154,7 @@ void si_pc_emit_spm_start(struct radeon_cmdbuf *cs) + S_036020_PERFMON_STATE(V_036020_CP_PERFMON_STATE_DISABLE_AND_RESET) | + S_036020_SPM_PERFMON_STATE(V_036020_STRM_PERFMON_STATE_START_COUNTING)); + /* Start windowed performance counters. */ +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_PERFCOUNTER_START) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_PERFCOUNTER_START); + radeon_set_sh_reg(R_00B82C_COMPUTE_PERFCOUNT_ENABLE, S_00B82C_PERFCOUNT_ENABLE(1)); + + radeon_end(); +@@ -171,10 +166,8 @@ void si_pc_emit_spm_stop(struct radeon_cmdbuf *cs, bool never_stop_sq_perf_count + radeon_begin(cs); + + /* Stop windowed performance counters. */ +- if (!never_send_perfcounter_stop) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_PERFCOUNTER_STOP) | EVENT_INDEX(0)); +- } ++ if (!never_send_perfcounter_stop) ++ radeon_event_write(V_028A90_PERFCOUNTER_STOP); + + radeon_set_sh_reg(R_00B82C_COMPUTE_PERFCOUNT_ENABLE, S_00B82C_PERFCOUNT_ENABLE(0)); + +diff --git a/src/gallium/drivers/radeonsi/si_query.c b/src/gallium/drivers/radeonsi/si_query.c +index 18c19b58063..eca647505ae 100644 +--- a/src/gallium/drivers/radeonsi/si_query.c ++++ b/src/gallium/drivers/radeonsi/si_query.c +@@ -980,8 +980,7 @@ static void si_query_hw_do_emit_stop(struct si_context *sctx, struct si_query_hw + + radeon_begin(cs); + if (sctx->screen->use_ngg && query->flags & SI_QUERY_EMULATE_GS_COUNTERS) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_VS_PARTIAL_FLUSH) | EVENT_INDEX(4)); ++ radeon_event_write(V_028A90_VS_PARTIAL_FLUSH); + + if (--sctx->num_pipeline_stat_emulated_queries == 0) { + si_set_internal_shader_buffer(sctx, SI_GS_QUERY_BUF, NULL); +diff --git a/src/gallium/drivers/radeonsi/si_state.c b/src/gallium/drivers/radeonsi/si_state.c +index 018be3dbe4f..8af8bae58ae 100644 +--- a/src/gallium/drivers/radeonsi/si_state.c ++++ b/src/gallium/drivers/radeonsi/si_state.c +@@ -62,8 +62,7 @@ static void si_emit_cb_render_state(struct si_context *sctx, unsigned index) + sctx->last_cb_target_mask = cb_target_mask; + + radeon_begin(cs); +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_BREAK_BATCH) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_BREAK_BATCH); + radeon_end(); + } + +@@ -3191,10 +3190,9 @@ static void gfx6_emit_framebuffer_state(struct si_context *sctx, unsigned index) + S_028208_BR_X(state->width) | S_028208_BR_Y(state->height)); + + if (sctx->screen->dpbb_allowed && +- sctx->screen->pbb_context_states_per_bin > 1) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_BREAK_BATCH) | EVENT_INDEX(0)); +- } ++ sctx->screen->pbb_context_states_per_bin > 1) ++ radeon_event_write(V_028A90_BREAK_BATCH); ++ + radeon_end(); + + si_update_display_dcc_dirty(sctx); +@@ -3341,10 +3339,9 @@ static void gfx11_dgpu_emit_framebuffer_state(struct si_context *sctx, unsigned + gfx11_end_packed_context_regs(); + + if (sctx->screen->dpbb_allowed && +- sctx->screen->pbb_context_states_per_bin > 1) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_BREAK_BATCH) | EVENT_INDEX(0)); +- } ++ sctx->screen->pbb_context_states_per_bin > 1) ++ radeon_event_write(V_028A90_BREAK_BATCH); ++ + radeon_end(); + + si_update_display_dcc_dirty(sctx); +@@ -3478,10 +3475,9 @@ static void gfx12_emit_framebuffer_state(struct si_context *sctx, unsigned index + gfx12_end_context_regs(); + + if (sctx->screen->dpbb_allowed && +- sctx->screen->pbb_context_states_per_bin > 1) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_BREAK_BATCH) | EVENT_INDEX(0)); +- } ++ sctx->screen->pbb_context_states_per_bin > 1) ++ radeon_event_write(V_028A90_BREAK_BATCH); ++ + radeon_end(); + + sctx->framebuffer.dirty_cbufs = 0; +diff --git a/src/gallium/drivers/radeonsi/si_state_draw.cpp b/src/gallium/drivers/radeonsi/si_state_draw.cpp +index bd2c936306b..81565b7694e 100644 +--- a/src/gallium/drivers/radeonsi/si_state_draw.cpp ++++ b/src/gallium/drivers/radeonsi/si_state_draw.cpp +@@ -1212,16 +1212,6 @@ void si_emit_buffered_compute_sh_regs(struct si_context *sctx) + + #endif + +-#define EMIT_SQTT_END_DRAW \ +- do { \ +- if (GFX_VERSION >= GFX9 && unlikely(sctx->sqtt_enabled)) { \ +- radeon_begin(&sctx->gfx_cs); \ +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); \ +- radeon_emit(EVENT_TYPE(V_028A90_THREAD_TRACE_MARKER) | EVENT_INDEX(0)); \ +- radeon_end(); \ +- } \ +- } while (0) +- + template ALWAYS_INLINE + static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw_info *info, +@@ -1638,10 +1628,8 @@ static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw + radeon_emit(0); + radeon_emit(V_0287F0_DI_SRC_SEL_AUTO_INDEX | use_opaque); + +- for (unsigned i = 0; i < 3; i++) { +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_SQ_NON_EVENT) | EVENT_INDEX(0)); +- } ++ for (unsigned i = 0; i < 3; i++) ++ radeon_event_write(V_028A90_SQ_NON_EVENT); + } else if (increment_draw_id) { + for (unsigned i = 0; i < num_draws; i++) { + if (i > 0) { +@@ -1675,9 +1663,10 @@ static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw + } + } + } +- radeon_end(); + +- EMIT_SQTT_END_DRAW; ++ if (GFX_VERSION >= GFX9 && unlikely(sctx->sqtt_enabled)) ++ radeon_event_write(V_028A90_THREAD_TRACE_MARKER); ++ radeon_end(); + } + + /* Return false if not bound. */ +diff --git a/src/gallium/drivers/radeonsi/si_state_shaders.cpp b/src/gallium/drivers/radeonsi/si_state_shaders.cpp +index 3eb56675313..a3a1b613337 100644 +--- a/src/gallium/drivers/radeonsi/si_state_shaders.cpp ++++ b/src/gallium/drivers/radeonsi/si_state_shaders.cpp +@@ -4081,12 +4081,10 @@ static void si_emit_vgt_flush(struct radeon_cmdbuf *cs) + radeon_begin(cs); + + /* This is required before VGT_FLUSH. */ +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_VS_PARTIAL_FLUSH) | EVENT_INDEX(4)); ++ radeon_event_write(V_028A90_VS_PARTIAL_FLUSH); + + /* VGT_FLUSH is required even if VGT is idle. It resets VGT pointers. */ +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_VGT_FLUSH) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_VGT_FLUSH); + radeon_end(); + } + +@@ -4973,11 +4971,8 @@ static void si_emit_spi_ge_ring_state(struct si_context *sctx, unsigned index) + + radeon_begin(&sctx->gfx_cs); + /* Required before writing tessellation config registers. */ +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_VS_PARTIAL_FLUSH) | EVENT_INDEX(4)); +- +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_VGT_FLUSH) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_VS_PARTIAL_FLUSH); ++ radeon_event_write(V_028A90_VGT_FLUSH); + + if (sctx->gfx_level >= GFX7) { + radeon_set_uconfig_reg_seq(R_030938_VGT_TF_RING_SIZE, 3); +diff --git a/src/gallium/drivers/radeonsi/si_state_streamout.c b/src/gallium/drivers/radeonsi/si_state_streamout.c +index 8047965ff84..9a92b4a4f3d 100644 +--- a/src/gallium/drivers/radeonsi/si_state_streamout.c ++++ b/src/gallium/drivers/radeonsi/si_state_streamout.c +@@ -259,8 +259,7 @@ static void si_flush_vgt_streamout(struct si_context *sctx) + radeon_set_config_reg(reg_strmout_cntl, 0); + } + +- radeon_emit(PKT3(PKT3_EVENT_WRITE, 0, 0)); +- radeon_emit(EVENT_TYPE(V_028A90_SO_VGTSTREAMOUT_FLUSH) | EVENT_INDEX(0)); ++ radeon_event_write(V_028A90_SO_VGTSTREAMOUT_FLUSH); + + radeon_emit(PKT3(PKT3_WAIT_REG_MEM, 5, 0)); + radeon_emit(WAIT_REG_MEM_EQUAL); /* wait until the register is equal to the reference value */ +-- +2.48.1 + + +From eefc15a11186a5b806ed5b2fcfbd2c9e1ad3d8e9 Mon Sep 17 00:00:00 2001 +From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= +Date: Mon, 16 Dec 2024 03:59:13 -0500 +Subject: [PATCH 05/27] radeonsi/gfx12: adjust HiZ/HiS logic + +Acked-by: Pierre-Eric Pelloux-Prayer +Part-of: +--- + src/amd/common/ac_gpu_info.c | 1 + + src/amd/common/ac_surface.c | 8 ++- + src/gallium/drivers/radeonsi/si_build_pm4.h | 14 +++++ + .../drivers/radeonsi/si_debug_options.h | 1 + + src/gallium/drivers/radeonsi/si_pipe.h | 1 + + src/gallium/drivers/radeonsi/si_state.c | 51 ++++++++++++---- + .../drivers/radeonsi/si_state_draw.cpp | 61 +++++++++++++++---- + 7 files changed, 110 insertions(+), 27 deletions(-) + +diff --git a/src/amd/common/ac_gpu_info.c b/src/amd/common/ac_gpu_info.c +index 0d475cdab63..dd43af8dbb6 100644 +--- a/src/amd/common/ac_gpu_info.c ++++ b/src/amd/common/ac_gpu_info.c +@@ -593,6 +593,7 @@ static void handle_env_var_force_family(struct radeon_info *info) + info->gfx_level = ac_get_gfx_level(i); + info->family_id = ac_get_family_id(i); + info->family_overridden = true; ++ info->chip_rev = 1; + return; + } + } +diff --git a/src/amd/common/ac_surface.c b/src/amd/common/ac_surface.c +index 69f722bc4cc..66d773921c4 100644 +--- a/src/amd/common/ac_surface.c ++++ b/src/amd/common/ac_surface.c +@@ -3022,7 +3022,7 @@ static bool gfx12_compute_hiz_his_info(struct ac_addrlib *addrlib, const struct + { + assert(surf_in->flags.depth != surf_in->flags.stencil); + +- if (surf->flags & RADEON_SURF_NO_HTILE || (info->gfx_level == GFX12 && info->chip_rev <= 1)) ++ if (surf->flags & RADEON_SURF_NO_HTILE || (info->gfx_level == GFX12 && info->chip_rev == 0)) + return true; + + ADDR3_COMPUTE_SURFACE_INFO_OUTPUT out = {0}; +@@ -3079,7 +3079,11 @@ static bool gfx12_compute_miptree(struct ac_addrlib *addrlib, const struct radeo + surf->surf_alignment_log2 = MAX2(surf->surf_alignment_log2, util_logbase2(out.baseAlign)); + surf->surf_size = surf->u.gfx9.zs.stencil_offset + out.surfSize; + +- return gfx12_compute_hiz_his_info(addrlib, info, surf, &surf->u.gfx9.zs.his, in); ++ if (info->chip_rev >= 2 && ++ !gfx12_compute_hiz_his_info(addrlib, info, surf, &surf->u.gfx9.zs.his, in)) ++ return false; ++ ++ return true; + } + + surf->u.gfx9.surf_slice_size = out.sliceSize; +diff --git a/src/gallium/drivers/radeonsi/si_build_pm4.h b/src/gallium/drivers/radeonsi/si_build_pm4.h +index 630b5885ff3..c792d0fd2f8 100644 +--- a/src/gallium/drivers/radeonsi/si_build_pm4.h ++++ b/src/gallium/drivers/radeonsi/si_build_pm4.h +@@ -514,6 +514,20 @@ + __event_type == V_028A90_PIXEL_PIPE_STAT_CONTROL ? 1 : 0)); \ + } while (0) + ++#define radeon_emit_alt_hiz_logic() do { \ ++ static_assert(GFX_VERSION == GFX12 || !ALT_HIZ_LOGIC, ""); \ ++ if (GFX_VERSION == GFX12 && ALT_HIZ_LOGIC) { \ ++ radeon_emit(PKT3(PKT3_RELEASE_MEM, 6, 0)); \ ++ radeon_emit(S_490_EVENT_TYPE(V_028A90_BOTTOM_OF_PIPE_TS) | S_490_EVENT_INDEX(5)); \ ++ radeon_emit(0); /* DST_SEL, INT_SEL = no write confirm, DATA_SEL = no data */ \ ++ radeon_emit(0); /* ADDRESS_LO */ \ ++ radeon_emit(0); /* ADDRESS_HI */ \ ++ radeon_emit(0); /* DATA_LO */ \ ++ radeon_emit(0); /* DATA_HI */ \ ++ radeon_emit(0); /* INT_CTXID */ \ ++ } \ ++} while (0) ++ + /* This should be evaluated at compile time if all parameters are constants. */ + static ALWAYS_INLINE unsigned + si_get_user_data_base(enum amd_gfx_level gfx_level, enum si_has_tess has_tess, +diff --git a/src/gallium/drivers/radeonsi/si_debug_options.h b/src/gallium/drivers/radeonsi/si_debug_options.h +index ba2c2336ee0..a5057c4700e 100644 +--- a/src/gallium/drivers/radeonsi/si_debug_options.h ++++ b/src/gallium/drivers/radeonsi/si_debug_options.h +@@ -23,6 +23,7 @@ OPT_BOOL(zerovram, false, "Zero all VRAM allocations") + OPT_BOOL(clear_lds, false, "Clear LDS at the end of shaders. Might decrease performance.") + OPT_BOOL(cache_rb_gl2, false, "Enable GL2 caching for CB and DB.") + OPT_BOOL(optimize_io, true, "Run nir_opt_varyings in the GLSL linker.") ++OPT_BOOL(alt_hiz_logic, false, "Enable alternative HiZ logic") + + #undef OPT_BOOL + #undef OPT_INT +diff --git a/src/gallium/drivers/radeonsi/si_pipe.h b/src/gallium/drivers/radeonsi/si_pipe.h +index a7f5b544051..0002a3056c9 100644 +--- a/src/gallium/drivers/radeonsi/si_pipe.h ++++ b/src/gallium/drivers/radeonsi/si_pipe.h +@@ -398,6 +398,7 @@ struct si_texture { + bool can_sample_z : 1; + bool can_sample_s : 1; + bool need_flush_after_depth_decompression: 1; ++ bool force_disable_hiz_his : 1; + + /* We need to track DCC dirtiness, because st/dri usually calls + * flush_resource twice per frame (not a bug) and we don't wanna +diff --git a/src/gallium/drivers/radeonsi/si_state.c b/src/gallium/drivers/radeonsi/si_state.c +index 8af8bae58ae..a647defc26c 100644 +--- a/src/gallium/drivers/radeonsi/si_state.c ++++ b/src/gallium/drivers/radeonsi/si_state.c +@@ -1713,6 +1713,20 @@ static void si_bind_dsa_state(struct pipe_context *ctx, void *state) + si_mark_atom_dirty(sctx, &sctx->atoms.s.stencil_ref); + } + ++ struct pipe_surface *zssurf = sctx->framebuffer.state.zsbuf; ++ struct si_texture *zstex = (struct si_texture*)(zssurf ? zssurf->texture : NULL); ++ ++ if (sctx->gfx_level == GFX12 && !sctx->screen->options.alt_hiz_logic && ++ sctx->framebuffer.has_stencil && dsa->stencil_enabled && !zstex->force_disable_hiz_his) { ++ zstex->force_disable_hiz_his = true; ++ si_mark_atom_dirty(sctx, &sctx->atoms.s.framebuffer); ++ ++ if (sctx->framebuffer.has_hiz_his) { ++ sctx->framebuffer.has_hiz_his = false; ++ si_mark_atom_dirty(sctx, &sctx->atoms.s.msaa_config); ++ } ++ } ++ + if (old_dsa->alpha_func != dsa->alpha_func) { + si_ps_key_update_dsa(sctx); + si_update_ps_inputs_read_or_disabled(sctx); +@@ -2814,9 +2828,14 @@ static void si_set_framebuffer_state(struct pipe_context *ctx, + if (util_format_has_stencil(util_format_description(zstex->buffer.b.b.format))) + sctx->framebuffer.has_stencil = true; + ++ if (sctx->gfx_level == GFX12 && !sctx->screen->options.alt_hiz_logic && ++ sctx->framebuffer.has_stencil && sctx->queued.named.dsa->stencil_enabled) ++ zstex->force_disable_hiz_his = true; ++ + if (sctx->gfx_level >= GFX12) { +- sctx->framebuffer.has_hiz_his = zstex->surface.u.gfx9.zs.hiz.offset || +- zstex->surface.u.gfx9.zs.his.offset; ++ sctx->framebuffer.has_hiz_his = (zstex->surface.u.gfx9.zs.hiz.offset || ++ zstex->surface.u.gfx9.zs.his.offset) && ++ !zstex->force_disable_hiz_his; + } + } + +@@ -3443,18 +3462,24 @@ static void gfx12_emit_framebuffer_state(struct si_context *sctx, unsigned index + gfx12_set_context_reg(R_028034_DB_STENCIL_READ_BASE_HI, zb->ds.db_stencil_base >> 32); + gfx12_set_context_reg(R_028038_DB_STENCIL_WRITE_BASE, zb->ds.db_stencil_base); + gfx12_set_context_reg(R_02803C_DB_STENCIL_WRITE_BASE_HI, zb->ds.db_stencil_base >> 32); +- gfx12_set_context_reg(R_028B94_PA_SC_HIZ_INFO, zb->ds.u.gfx12.hiz_info); +- gfx12_set_context_reg(R_028B98_PA_SC_HIS_INFO, zb->ds.u.gfx12.his_info); + +- if (zb->ds.u.gfx12.hiz_info) { +- gfx12_set_context_reg(R_028B9C_PA_SC_HIZ_BASE, zb->ds.u.gfx12.hiz_base); +- gfx12_set_context_reg(R_028BA0_PA_SC_HIZ_BASE_EXT, zb->ds.u.gfx12.hiz_base >> 32); +- gfx12_set_context_reg(R_028BA4_PA_SC_HIZ_SIZE_XY, zb->ds.u.gfx12.hiz_size_xy); +- } +- if (zb->ds.u.gfx12.his_info) { +- gfx12_set_context_reg(R_028BA8_PA_SC_HIS_BASE, zb->ds.u.gfx12.his_base); +- gfx12_set_context_reg(R_028BAC_PA_SC_HIS_BASE_EXT, zb->ds.u.gfx12.his_base >> 32); +- gfx12_set_context_reg(R_028BB0_PA_SC_HIS_SIZE_XY, zb->ds.u.gfx12.his_size_xy); ++ if (tex->force_disable_hiz_his) { ++ gfx12_set_context_reg(R_028B94_PA_SC_HIZ_INFO, S_028B94_SURFACE_ENABLE(0)); ++ gfx12_set_context_reg(R_028B98_PA_SC_HIS_INFO, S_028B98_SURFACE_ENABLE(0)); ++ } else { ++ gfx12_set_context_reg(R_028B94_PA_SC_HIZ_INFO, zb->ds.u.gfx12.hiz_info); ++ gfx12_set_context_reg(R_028B98_PA_SC_HIS_INFO, zb->ds.u.gfx12.his_info); ++ ++ if (zb->ds.u.gfx12.hiz_info) { ++ gfx12_set_context_reg(R_028B9C_PA_SC_HIZ_BASE, zb->ds.u.gfx12.hiz_base); ++ gfx12_set_context_reg(R_028BA0_PA_SC_HIZ_BASE_EXT, zb->ds.u.gfx12.hiz_base >> 32); ++ gfx12_set_context_reg(R_028BA4_PA_SC_HIZ_SIZE_XY, zb->ds.u.gfx12.hiz_size_xy); ++ } ++ if (zb->ds.u.gfx12.his_info) { ++ gfx12_set_context_reg(R_028BA8_PA_SC_HIS_BASE, zb->ds.u.gfx12.his_base); ++ gfx12_set_context_reg(R_028BAC_PA_SC_HIS_BASE_EXT, zb->ds.u.gfx12.his_base >> 32); ++ gfx12_set_context_reg(R_028BB0_PA_SC_HIS_SIZE_XY, zb->ds.u.gfx12.his_size_xy); ++ } + } + } else if (sctx->framebuffer.dirty_zsbuf) { + gfx12_set_context_reg(R_028018_DB_Z_INFO, +diff --git a/src/gallium/drivers/radeonsi/si_state_draw.cpp b/src/gallium/drivers/radeonsi/si_state_draw.cpp +index 81565b7694e..bc4cf308936 100644 +--- a/src/gallium/drivers/radeonsi/si_state_draw.cpp ++++ b/src/gallium/drivers/radeonsi/si_state_draw.cpp +@@ -823,6 +823,11 @@ enum si_has_sh_pairs_packed { + HAS_SH_PAIRS_PACKED_ON, + }; + ++enum si_alt_hiz_logic { ++ ALT_HIZ_LOGIC_OFF, ++ ALT_HIZ_LOGIC_ON, ++}; ++ + template ALWAYS_INLINE + static bool num_instanced_prims_less_than(const struct pipe_draw_indirect_info *indirect, + enum mesa_prim prim, +@@ -1213,7 +1218,8 @@ void si_emit_buffered_compute_sh_regs(struct si_context *sctx) + #endif + + template ALWAYS_INLINE ++ si_is_draw_vertex_state IS_DRAW_VERTEX_STATE, si_has_sh_pairs_packed HAS_SH_PAIRS_PACKED, ++ si_alt_hiz_logic ALT_HIZ_LOGIC> ALWAYS_INLINE + static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw_info *info, + unsigned drawid_base, + const struct pipe_draw_indirect_info *indirect, +@@ -1405,6 +1411,8 @@ static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw + radeon_emit((sh_base_reg + SI_SGPR_BASE_VERTEX * 4 - SI_SH_REG_OFFSET) >> 2); + radeon_emit((sh_base_reg + SI_SGPR_START_INSTANCE * 4 - SI_SH_REG_OFFSET) >> 2); + radeon_emit(di_src_sel); ++ ++ radeon_emit_alt_hiz_logic(); + } else { + uint64_t count_va = 0; + +@@ -1430,6 +1438,8 @@ static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw + radeon_emit(count_va >> 32); + radeon_emit(indirect->stride); + radeon_emit(di_src_sel); ++ ++ radeon_emit_alt_hiz_logic(); + } + } else { + if (sctx->last_instance_count == SI_INSTANCE_COUNT_UNKNOWN || +@@ -1549,6 +1559,8 @@ static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw + radeon_emit(va >> 32); + radeon_emit(draws[i].count); + radeon_emit(V_0287F0_DI_SRC_SEL_DMA); /* NOT_EOP disabled */ ++ ++ radeon_emit_alt_hiz_logic(); + } + if (num_draws > 1) { + BITSET_CLEAR(sctx->tracked_regs.reg_saved_mask, tracked_base_vertex_reg); /* BaseVertex */ +@@ -1568,6 +1580,8 @@ static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw + radeon_emit(va >> 32); + radeon_emit(draws[i].count); + radeon_emit(V_0287F0_DI_SRC_SEL_DMA); /* NOT_EOP disabled */ ++ ++ radeon_emit_alt_hiz_logic(); + } + if (num_draws > 1) { + BITSET_CLEAR(sctx->tracked_regs.reg_saved_mask, tracked_base_vertex_reg + 1); /* DrawID */ +@@ -1588,6 +1602,8 @@ static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw + radeon_emit(va >> 32); + radeon_emit(draws[i].count); + radeon_emit(V_0287F0_DI_SRC_SEL_DMA); /* NOT_EOP disabled */ ++ ++ radeon_emit_alt_hiz_logic(); + } + if (num_draws > 1) { + BITSET_CLEAR(sctx->tracked_regs.reg_saved_mask, tracked_base_vertex_reg); /* BaseVertex */ +@@ -1615,6 +1631,8 @@ static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw + radeon_emit(V_0287F0_DI_SRC_SEL_DMA | + S_0287F0_NOT_EOP(GFX_VERSION >= GFX10 && GFX_VERSION < GFX12 && + i < num_draws - 1)); ++ ++ radeon_emit_alt_hiz_logic(); + } + } + } +@@ -1628,6 +1646,7 @@ static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw + radeon_emit(0); + radeon_emit(V_0287F0_DI_SRC_SEL_AUTO_INDEX | use_opaque); + ++ radeon_emit_alt_hiz_logic(); + for (unsigned i = 0; i < 3; i++) + radeon_event_write(V_028A90_SQ_NON_EVENT); + } else if (increment_draw_id) { +@@ -1643,6 +1662,8 @@ static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw + radeon_emit(PKT3(PKT3_DRAW_INDEX_AUTO, 1, render_cond_bit)); + radeon_emit(draws[i].count); + radeon_emit(V_0287F0_DI_SRC_SEL_AUTO_INDEX | use_opaque); ++ ++ radeon_emit_alt_hiz_logic(); + } + if (num_draws > 1 && (IS_DRAW_VERTEX_STATE || !sctx->num_vs_blit_sgprs)) { + BITSET_CLEAR(sctx->tracked_regs.reg_saved_mask, tracked_base_vertex_reg); /* BaseVertex */ +@@ -1656,6 +1677,8 @@ static void si_emit_draw_packets(struct si_context *sctx, const struct pipe_draw + radeon_emit(PKT3(PKT3_DRAW_INDEX_AUTO, 1, render_cond_bit)); + radeon_emit(draws[i].count); + radeon_emit(V_0287F0_DI_SRC_SEL_AUTO_INDEX | use_opaque); ++ ++ radeon_emit_alt_hiz_logic(); + } + if (num_draws > 1 && (IS_DRAW_VERTEX_STATE || !sctx->num_vs_blit_sgprs)) { + BITSET_CLEAR(sctx->tracked_regs.reg_saved_mask, tracked_base_vertex_reg); /* BaseVertex */ +@@ -2012,7 +2035,7 @@ static void si_emit_all_states(struct si_context *sctx, uint64_t skip_atom_mask) + + template ALWAYS_INLINE ++ util_popcnt POPCNT, si_alt_hiz_logic ALT_HIZ_LOGIC> ALWAYS_INLINE + static void si_draw(struct pipe_context *ctx, + const struct pipe_draw_info *info, + unsigned drawid_offset, +@@ -2312,7 +2335,8 @@ static void si_draw(struct pipe_context *ctx, + return; + } + +- si_emit_draw_packets ++ si_emit_draw_packets + (sctx, info, drawid_offset, indirect, draws, num_draws, indexbuf, + index_size, index_offset, instance_count); + /* <-- CUs start to get busy here if we waited. */ +@@ -2362,7 +2386,7 @@ static void si_draw(struct pipe_context *ctx, + } + + template ++ si_has_sh_pairs_packed HAS_SH_PAIRS_PACKED, si_alt_hiz_logic ALT_HIZ_LOGIC> + static void si_draw_vbo(struct pipe_context *ctx, + const struct pipe_draw_info *info, + unsigned drawid_offset, +@@ -2370,12 +2394,14 @@ static void si_draw_vbo(struct pipe_context *ctx, + const struct pipe_draw_start_count_bias *draws, + unsigned num_draws) + { +- si_draw ++ si_draw + (ctx, info, drawid_offset, indirect, draws, num_draws, NULL, 0); + } + + template ++ si_has_sh_pairs_packed HAS_SH_PAIRS_PACKED, util_popcnt POPCNT, ++ si_alt_hiz_logic ALT_HIZ_LOGIC> + static void si_draw_vertex_state(struct pipe_context *ctx, + struct pipe_vertex_state *vstate, + uint32_t partial_velem_mask, +@@ -2391,7 +2417,8 @@ static void si_draw_vertex_state(struct pipe_context *ctx, + dinfo.instance_count = 1; + dinfo.index.resource = state->b.input.indexbuf; + +- si_draw ++ si_draw + (ctx, &dinfo, 0, NULL, draws, num_draws, vstate, partial_velem_mask); + + if (info.take_vertex_state_ownership) +@@ -2453,18 +2480,28 @@ static void si_init_draw_vbo(struct si_context *sctx) + if (!NGG && GFX_VERSION >= GFX11) + return; + +- if (GFX_VERSION >= GFX11 && GFX_VERSION < GFX12 && sctx->screen->info.has_set_sh_pairs_packed) { ++ if (GFX_VERSION == GFX12 && sctx->screen->options.alt_hiz_logic) { ++ sctx->draw_vbo[HAS_TESS][HAS_GS][NGG] = ++ si_draw_vbo; ++ ++ sctx->draw_vertex_state[HAS_TESS][HAS_GS][NGG] = ++ si_draw_vertex_state; ++ } else if (GFX_VERSION >= GFX11 && GFX_VERSION < GFX12 && ++ sctx->screen->info.has_set_sh_pairs_packed) { + sctx->draw_vbo[HAS_TESS][HAS_GS][NGG] = +- si_draw_vbo; ++ si_draw_vbo; + + sctx->draw_vertex_state[HAS_TESS][HAS_GS][NGG] = +- si_draw_vertex_state; ++ si_draw_vertex_state; + } else { + sctx->draw_vbo[HAS_TESS][HAS_GS][NGG] = +- si_draw_vbo; ++ si_draw_vbo; + + sctx->draw_vertex_state[HAS_TESS][HAS_GS][NGG] = +- si_draw_vertex_state; ++ si_draw_vertex_state; + } + } + +-- +2.48.1 + + +From a4f5130de6d737aa1264d3a981f95b81a9f13d05 Mon Sep 17 00:00:00 2001 +From: Pierre-Eric Pelloux-Prayer +Date: Thu, 21 Nov 2024 13:32:55 +0100 +Subject: [PATCH 06/27] radeonsi: disable DCC for PIPE_BIND_USE_FRONT_RENDERING +MIME-Version: 1.0 +Content-Type: text/plain; charset=UTF-8 +Content-Transfer-Encoding: 8bit + +Front rendering and (display) DCC are causing artifacts on screen. + +si_texture_get_handle deals with this problem by disabling dcc, but +we can make it simpler by not allocating DCC at all when this flag +is set. + +Reviewed-by: Marek Olšák +Part-of: +--- + src/gallium/drivers/radeonsi/si_texture.c | 19 +++++++++++++++++++ + 1 file changed, 19 insertions(+) + +diff --git a/src/gallium/drivers/radeonsi/si_texture.c b/src/gallium/drivers/radeonsi/si_texture.c +index ae65b725600..9cc11a81669 100644 +--- a/src/gallium/drivers/radeonsi/si_texture.c ++++ b/src/gallium/drivers/radeonsi/si_texture.c +@@ -240,6 +240,8 @@ static int si_init_surface(struct si_screen *sscreen, struct radeon_surf *surfac + + if (modifier == DRM_FORMAT_MOD_INVALID && + (ptex->bind & PIPE_BIND_CONST_BW || ++ ptex->bind & PIPE_BIND_PROTECTED || ++ ptex->bind & PIPE_BIND_USE_FRONT_RENDERING || + sscreen->debug_flags & DBG(NO_DCC) || + (ptex->bind & PIPE_BIND_SCANOUT && sscreen->debug_flags & DBG(NO_DISPLAY_DCC)))) + flags |= RADEON_SURF_DISABLE_DCC; +@@ -289,6 +291,9 @@ static int si_init_surface(struct si_screen *sscreen, struct radeon_surf *surfac + if (ptex->bind & PIPE_BIND_CONST_BW) + flags |= RADEON_SURF_DISABLE_DCC; + ++ if (ptex->bind & PIPE_BIND_USE_FRONT_RENDERING) ++ flags |= RADEON_SURF_DISABLE_DCC; ++ + switch (sscreen->info.gfx_level) { + case GFX8: + /* Stoney: 128bpp MSAA textures randomly fail piglit tests with DCC. */ +@@ -1638,6 +1643,20 @@ si_modifier_supports_resource(struct pipe_screen *screen, + struct si_screen *sscreen = (struct si_screen *)screen; + uint32_t max_width, max_height; + ++ if (((templ->bind & PIPE_BIND_LINEAR) || sscreen->debug_flags & DBG(NO_TILING)) && ++ modifier != DRM_FORMAT_MOD_LINEAR) ++ return false; ++ ++ if ((templ->bind & PIPE_BIND_USE_FRONT_RENDERING) && ac_modifier_has_dcc(modifier)) ++ return false; ++ ++ /* Protected content doesn't support DCC on GFX12. */ ++ if (sscreen->info.gfx_level >= GFX12 && templ->bind & PIPE_BIND_PROTECTED && ++ IS_AMD_FMT_MOD(modifier) && ++ AMD_FMT_MOD_GET(TILE_VERSION, modifier) >= AMD_FMT_MOD_TILE_VER_GFX12 && ++ AMD_FMT_MOD_GET(DCC, modifier)) ++ return false; ++ + ac_modifier_max_extent(&sscreen->info, modifier, &max_width, &max_height); + return templ->width0 <= max_width && templ->height0 <= max_height; + } +-- +2.48.1 + + +From 35435069c1063180814aff055365dce939d95e60 Mon Sep 17 00:00:00 2001 +From: Pierre-Eric Pelloux-Prayer +Date: Wed, 29 Jan 2025 18:10:12 +0100 +Subject: [PATCH 07/27] radeonsi: update si_need_gfx_cs_space upper bound +MIME-Version: 1.0 +Content-Type: text/plain; charset=UTF-8 +Content-Transfer-Encoding: 8bit + +radeon_emit_alt_hiz_logic can add 8 extra dw per draw. + +Fixes: cdecbee9225 ("radeonsi/gfx12: adjust HiZ/HiS logic") +Reviewed-by: Marek Olšák +Part-of: +--- + src/gallium/drivers/radeonsi/si_compute.c | 2 +- + src/gallium/drivers/radeonsi/si_cp_dma.c | 2 +- + src/gallium/drivers/radeonsi/si_perfcounter.c | 2 +- + src/gallium/drivers/radeonsi/si_pipe.h | 12 ++++++++++-- + src/gallium/drivers/radeonsi/si_query.c | 6 +++--- + src/gallium/drivers/radeonsi/si_state_draw.cpp | 2 +- + 6 files changed, 17 insertions(+), 9 deletions(-) + +diff --git a/src/gallium/drivers/radeonsi/si_compute.c b/src/gallium/drivers/radeonsi/si_compute.c +index 01aa75d222a..d2c2bc16500 100644 +--- a/src/gallium/drivers/radeonsi/si_compute.c ++++ b/src/gallium/drivers/radeonsi/si_compute.c +@@ -1218,7 +1218,7 @@ static void si_launch_grid(struct pipe_context *ctx, const struct pipe_grid_info + } + } + +- si_need_gfx_cs_space(sctx, 0); ++ si_need_gfx_cs_space(sctx, 0, 0); + + /* If we're using a secure context, determine if cs must be secure or not */ + if (unlikely(radeon_uses_secure_bos(sctx->ws))) { +diff --git a/src/gallium/drivers/radeonsi/si_cp_dma.c b/src/gallium/drivers/radeonsi/si_cp_dma.c +index 89ee80e14f5..f0ba0796377 100644 +--- a/src/gallium/drivers/radeonsi/si_cp_dma.c ++++ b/src/gallium/drivers/radeonsi/si_cp_dma.c +@@ -146,7 +146,7 @@ static void si_cp_dma_prepare(struct si_context *sctx, struct pipe_resource *dst + bool *is_first, unsigned *packet_flags) + { + if (!(user_flags & SI_OP_CPDMA_SKIP_CHECK_CS_SPACE)) +- si_need_gfx_cs_space(sctx, 0); ++ si_need_gfx_cs_space(sctx, 0, 0); + + /* This must be done after need_cs_space. */ + if (dst) +diff --git a/src/gallium/drivers/radeonsi/si_perfcounter.c b/src/gallium/drivers/radeonsi/si_perfcounter.c +index 785c98a9dc9..806a842a08d 100644 +--- a/src/gallium/drivers/radeonsi/si_perfcounter.c ++++ b/src/gallium/drivers/radeonsi/si_perfcounter.c +@@ -277,7 +277,7 @@ static void si_pc_query_resume(struct si_context *sctx, struct si_query *squery) + + if (!si_query_buffer_alloc(sctx, &query->buffer, NULL, query->result_size)) + return; +- si_need_gfx_cs_space(sctx, 0); ++ si_need_gfx_cs_space(sctx, 0, 0); + + if (query->shaders) + si_pc_emit_shaders(&sctx->gfx_cs, query->shaders); +diff --git a/src/gallium/drivers/radeonsi/si_pipe.h b/src/gallium/drivers/radeonsi/si_pipe.h +index 0002a3056c9..f132e182b4f 100644 +--- a/src/gallium/drivers/radeonsi/si_pipe.h ++++ b/src/gallium/drivers/radeonsi/si_pipe.h +@@ -2026,11 +2026,19 @@ static inline bool util_rast_prim_is_lines_or_triangles(unsigned prim) + return ((1 << prim) & (UTIL_ALL_PRIM_LINE_MODES | UTIL_ALL_PRIM_TRIANGLE_MODES)) != 0; + } + +-static inline void si_need_gfx_cs_space(struct si_context *ctx, unsigned num_draws) ++static inline void si_need_gfx_cs_space(struct si_context *ctx, unsigned num_draws, ++ unsigned extra_dw_per_draw) + { + struct radeon_cmdbuf *cs = &ctx->gfx_cs; ++ /* Don't count the needed CS space exactly and just use an upper bound. ++ * ++ * Also reserve space for stopping queries at the end of IB, because ++ * the number of active queries is unlimited in theory. ++ */ ++ unsigned reserve_dw = 2048 + ctx->num_cs_dw_queries_suspend + ++ num_draws * (10 + extra_dw_per_draw); + +- if (!ctx->ws->cs_check_space(cs, si_get_minimum_num_gfx_cs_dwords(ctx, num_draws))) ++ if (!ctx->ws->cs_check_space(cs, reserve_dw)) + si_flush_gfx_cs(ctx, RADEON_FLUSH_ASYNC_START_NEXT_GFX_IB_NOW, NULL); + } + +diff --git a/src/gallium/drivers/radeonsi/si_query.c b/src/gallium/drivers/radeonsi/si_query.c +index eca647505ae..67dd732f615 100644 +--- a/src/gallium/drivers/radeonsi/si_query.c ++++ b/src/gallium/drivers/radeonsi/si_query.c +@@ -918,7 +918,7 @@ static void si_query_hw_emit_start(struct si_context *sctx, struct si_query_hw * + si_update_prims_generated_query_state(sctx, query->b.type, 1); + si_update_hw_pipeline_stats(sctx, query->b.type, 1); + +- si_need_gfx_cs_space(sctx, 0); ++ si_need_gfx_cs_space(sctx, 0, 0); + + va = query->buffer.buf->gpu_address + query->buffer.results_end; + si_query_hw_do_emit_start(sctx, query, query->buffer.buf, va); +@@ -1014,7 +1014,7 @@ static void si_query_hw_emit_stop(struct si_context *sctx, struct si_query_hw *q + + /* The queries which need begin already called this in begin_query. */ + if (query->flags & SI_QUERY_HW_FLAG_NO_START) { +- si_need_gfx_cs_space(sctx, 0); ++ si_need_gfx_cs_space(sctx, 0, 0); + if (!si_query_buffer_alloc(sctx, &query->buffer, si_query_hw_prepare_buffer, + query->result_size)) + return; +@@ -1726,7 +1726,7 @@ void si_resume_queries(struct si_context *sctx) + struct si_query *query; + + /* Check CS space here. Resuming must not be interrupted by flushes. */ +- si_need_gfx_cs_space(sctx, 0); ++ si_need_gfx_cs_space(sctx, 0, 0); + + LIST_FOR_EACH_ENTRY (query, &sctx->active_queries, active_list) + query->ops->resume(sctx, query); +diff --git a/src/gallium/drivers/radeonsi/si_state_draw.cpp b/src/gallium/drivers/radeonsi/si_state_draw.cpp +index bc4cf308936..8b57ddb3b70 100644 +--- a/src/gallium/drivers/radeonsi/si_state_draw.cpp ++++ b/src/gallium/drivers/radeonsi/si_state_draw.cpp +@@ -2060,7 +2060,7 @@ static void si_draw(struct pipe_context *ctx, + else if (GFX_VERSION < GFX12) + gfx11_decompress_textures(sctx, u_bit_consecutive(0, SI_NUM_GRAPHICS_SHADERS)); + +- si_need_gfx_cs_space(sctx, num_draws); ++ si_need_gfx_cs_space(sctx, num_draws, ALT_HIZ_LOGIC ? 8 : 0); + + if (u_trace_perfetto_active(&sctx->ds.trace_context)) + trace_si_begin_draw(&sctx->trace); +-- +2.48.1 + + +From 1e95d1b0fec92e3cba014e32ac731705eb6e8425 Mon Sep 17 00:00:00 2001 +From: Samuel Pitoiset +Date: Wed, 15 Jan 2025 12:00:43 +0100 +Subject: [PATCH 08/27] radv: disable GFX12+ support + +It's mostly broken, but 25.0+ should be good enough. + +Signed-off-by: Samuel Pitoiset +Part-of: +--- + src/amd/vulkan/winsys/amdgpu/radv_amdgpu_winsys.c | 13 +++++++++++++ + 1 file changed, 13 insertions(+) + +diff --git a/src/amd/vulkan/winsys/amdgpu/radv_amdgpu_winsys.c b/src/amd/vulkan/winsys/amdgpu/radv_amdgpu_winsys.c +index 6ed5718e695..b6ce7ba6ac5 100644 +--- a/src/amd/vulkan/winsys/amdgpu/radv_amdgpu_winsys.c ++++ b/src/amd/vulkan/winsys/amdgpu/radv_amdgpu_winsys.c +@@ -22,12 +22,25 @@ + #include "vk_drm_syncobj.h" + #include "xf86drm.h" + ++static bool ++radv_is_gpu_supported(const struct radeon_info *info) ++{ ++ /* GFX12 isn't supported. */ ++ if (info->gfx_level >= GFX12) ++ return false; ++ ++ return true; ++} ++ + static bool + do_winsys_init(struct radv_amdgpu_winsys *ws, int fd) + { + if (!ac_query_gpu_info(fd, ws->dev, &ws->info, true)) + return false; + ++ if (!radv_is_gpu_supported(&ws->info)) ++ return false; ++ + /* + * Override the max submits on video queues. + * If you submit multiple session contexts in the same IB sequence the +-- +2.48.1 + + +From 317d71daef2588ff76c62aa54a5f1920e717a465 Mon Sep 17 00:00:00 2001 +From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= +Date: Wed, 18 Dec 2024 19:18:30 -0500 +Subject: [PATCH 09/27] radeonsi: fix a gfx10.3 regression due to a gfx12 + change + +This fixes: + Assertion `!"BITSET_TEST_RANGE: bit range crosses word boundary"' failed. + +Fixes: e3cef02c245 - radeonsi/gfx12: set DB_RENDER_OVERRIDE based on stencil state + +Reviewed-by: Qiang Yu +Part-of: +--- + src/gallium/drivers/radeonsi/si_gfx_cs.c | 7 +++++-- + src/gallium/drivers/radeonsi/si_state.h | 2 +- + 2 files changed, 6 insertions(+), 3 deletions(-) + +diff --git a/src/gallium/drivers/radeonsi/si_gfx_cs.c b/src/gallium/drivers/radeonsi/si_gfx_cs.c +index d8dccd9b008..9782e1afaa1 100644 +--- a/src/gallium/drivers/radeonsi/si_gfx_cs.c ++++ b/src/gallium/drivers/radeonsi/si_gfx_cs.c +@@ -317,7 +317,6 @@ void si_set_tracked_regs_to_clear_state(struct si_context *ctx) + ctx->tracked_regs.reg_value[SI_TRACKED_SPI_PS_INPUT_ENA] = 0; + ctx->tracked_regs.reg_value[SI_TRACKED_SPI_PS_INPUT_ADDR] = 0; + +- ctx->tracked_regs.reg_value[SI_TRACKED_DB_RENDER_OVERRIDE] = 0; + ctx->tracked_regs.reg_value[SI_TRACKED_DB_EQAA] = 0; + ctx->tracked_regs.reg_value[SI_TRACKED_DB_RENDER_OVERRIDE2] = 0; + ctx->tracked_regs.reg_value[SI_TRACKED_DB_SHADER_CONTROL] = 0; +@@ -368,7 +367,11 @@ void si_set_tracked_regs_to_clear_state(struct si_context *ctx) + ctx->tracked_regs.reg_value[SI_TRACKED_VGT_GS_VERT_ITEMSIZE_2] = 0; + ctx->tracked_regs.reg_value[SI_TRACKED_VGT_GS_VERT_ITEMSIZE_3] = 0; + +- ctx->tracked_regs.reg_value[SI_TRACKED_SPI_VS_OUT_CONFIG] = 0; ++ if (ctx->gfx_level >= GFX12) ++ ctx->tracked_regs.reg_value[SI_TRACKED_DB_RENDER_OVERRIDE] = 0; ++ else ++ ctx->tracked_regs.reg_value[SI_TRACKED_SPI_VS_OUT_CONFIG] = 0; ++ + ctx->tracked_regs.reg_value[SI_TRACKED_VGT_PRIMITIVEID_EN] = 0; + ctx->tracked_regs.reg_value[SI_TRACKED_CB_DCC_CONTROL] = 0; + +diff --git a/src/gallium/drivers/radeonsi/si_state.h b/src/gallium/drivers/radeonsi/si_state.h +index f964fa21606..75159c475e6 100644 +--- a/src/gallium/drivers/radeonsi/si_state.h ++++ b/src/gallium/drivers/radeonsi/si_state.h +@@ -315,7 +315,6 @@ enum si_tracked_reg + SI_TRACKED_SPI_PS_INPUT_ENA, + SI_TRACKED_SPI_PS_INPUT_ADDR, + +- SI_TRACKED_DB_RENDER_OVERRIDE, + SI_TRACKED_DB_EQAA, + SI_TRACKED_DB_RENDER_OVERRIDE2, + SI_TRACKED_DB_SHADER_CONTROL, +@@ -371,6 +370,7 @@ enum si_tracked_reg + SI_TRACKED_VGT_GS_VERT_ITEMSIZE_3, /* GFX6-10 (GFX11+ can reuse this slot) */ + + SI_TRACKED_SPI_VS_OUT_CONFIG, /* GFX6-11 */ ++ SI_TRACKED_DB_RENDER_OVERRIDE = SI_TRACKED_SPI_VS_OUT_CONFIG, /* GFX12+ (slot reused) */ + SI_TRACKED_VGT_PRIMITIVEID_EN, /* GFX6-11 */ + SI_TRACKED_CB_DCC_CONTROL, /* GFX8-11 */ + SI_TRACKED_DB_STENCIL_READ_MASK, /* GFX12+ */ +-- +2.48.1 + + +From 5e4f6e7f2c644219f998672b90c6de9146ca6952 Mon Sep 17 00:00:00 2001 +From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= +Date: Wed, 8 Jan 2025 15:00:50 -0500 +Subject: [PATCH 10/27] radeonsi/gfx12: enable alt_hiz_logic + +Reviewed-by: Pierre-Eric Pelloux-Prayer +Part-of: +--- + src/gallium/drivers/radeonsi/si_debug_options.h | 2 +- + 1 file changed, 1 insertion(+), 1 deletion(-) + +diff --git a/src/gallium/drivers/radeonsi/si_debug_options.h b/src/gallium/drivers/radeonsi/si_debug_options.h +index a5057c4700e..74b9b59a455 100644 +--- a/src/gallium/drivers/radeonsi/si_debug_options.h ++++ b/src/gallium/drivers/radeonsi/si_debug_options.h +@@ -23,7 +23,7 @@ OPT_BOOL(zerovram, false, "Zero all VRAM allocations") + OPT_BOOL(clear_lds, false, "Clear LDS at the end of shaders. Might decrease performance.") + OPT_BOOL(cache_rb_gl2, false, "Enable GL2 caching for CB and DB.") + OPT_BOOL(optimize_io, true, "Run nir_opt_varyings in the GLSL linker.") +-OPT_BOOL(alt_hiz_logic, false, "Enable alternative HiZ logic") ++OPT_BOOL(alt_hiz_logic, true, "Enable alternative HiZ logic") + + #undef OPT_BOOL + #undef OPT_INT +-- +2.48.1 + + +From 074851b27e9b59bf9b450907e3231824495e589a Mon Sep 17 00:00:00 2001 +From: Pierre-Eric Pelloux-Prayer +Date: Mon, 10 Feb 2025 12:27:48 +0100 +Subject: [PATCH 11/27] radeonsi: disable dcc when external shader stores are + used +MIME-Version: 1.0 +Content-Type: text/plain; charset=UTF-8 +Content-Transfer-Encoding: 8bit + +See comment. + +Fixes: 666a6eb871d ("radeonsi/gfx12: disable display dcc for front buffer rendering") +Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/12552 +Reviewed-by: Marek Olšák +Part-of: +--- + src/gallium/drivers/radeonsi/si_texture.c | 8 +++++--- + 1 file changed, 5 insertions(+), 3 deletions(-) + +diff --git a/src/gallium/drivers/radeonsi/si_texture.c b/src/gallium/drivers/radeonsi/si_texture.c +index 9cc11a81669..54e2bfe7e53 100644 +--- a/src/gallium/drivers/radeonsi/si_texture.c ++++ b/src/gallium/drivers/radeonsi/si_texture.c +@@ -803,10 +803,12 @@ static bool si_texture_get_handle(struct pipe_screen *screen, struct pipe_contex + } + + const bool debug_disable_dcc = sscreen->debug_flags & DBG(NO_EXPORTED_DCC); +- /* Since shader image stores don't support DCC on GFX9 and older, +- * disable it for external clients that want write access. ++ /* Disable DCC for external clients that might use shader image stores. ++ * They don't support DCC on GFX9 and older. GFX10/10.3 is also problematic ++ * if the view formats between clients are incompatible or if DCC clear is ++ * used. + */ +- const bool shader_write = sscreen->info.gfx_level <= GFX9 && ++ const bool shader_write = sscreen->info.gfx_level < GFX11 && + usage & PIPE_HANDLE_USAGE_SHADER_WRITE && + !tex->is_depth && + tex->surface.meta_offset; +-- +2.48.1 + + +From ec81fa2026aa8760158dcf4520bc307f627db40f Mon Sep 17 00:00:00 2001 +From: Rhys Perry +Date: Thu, 25 Jul 2024 21:39:25 +0100 +Subject: [PATCH 12/27] aco: split CounterMap off from VGPRCounterMap + +Signed-off-by: Rhys Perry +Reviewed-by: Georg Lehmann +Part-of: +--- + src/amd/compiler/aco_insert_NOPs.cpp | 38 ++++++++++++++++------------ + 1 file changed, 22 insertions(+), 16 deletions(-) + +diff --git a/src/amd/compiler/aco_insert_NOPs.cpp b/src/amd/compiler/aco_insert_NOPs.cpp +index a6290a7f457..3f4b5d6c85f 100644 +--- a/src/amd/compiler/aco_insert_NOPs.cpp ++++ b/src/amd/compiler/aco_insert_NOPs.cpp +@@ -164,14 +164,14 @@ struct NOP_ctx_gfx10 { + } + }; + +-template struct VGPRCounterMap { ++template struct CounterMap { + public: + int base = 0; +- BITSET_DECLARE(resident, 256); +- int val[256]; ++ BITSET_DECLARE(resident, Size); ++ int val[Size]; + + /* Initializes all counters to Max. */ +- VGPRCounterMap() { BITSET_ZERO(resident); } ++ CounterMap() { BITSET_ZERO(resident); } + + /* Increase all counters, clamping at Max. */ + void inc() { base++; } +@@ -185,11 +185,12 @@ public: + + void set(PhysReg reg, unsigned bytes) + { +- if (reg.reg() < 256) ++ if (reg.reg() < Start) + return; + +- for (unsigned i = 0; i < DIV_ROUND_UP(bytes, 4); i++) +- set(reg.reg() - 256 + i); ++ unsigned size = MIN2(DIV_ROUND_UP(bytes, 4), Start + Size - reg.reg()); ++ for (unsigned i = 0; i < size; i++) ++ set(reg.reg() - Start + i); + } + + /* Reset all counters to Max. */ +@@ -201,11 +202,12 @@ public: + + void reset(PhysReg reg, unsigned bytes) + { +- if (reg.reg() < 256) ++ if (reg.reg() < Start) + return; + +- for (unsigned i = 0; i < DIV_ROUND_UP(bytes, 4); i++) +- BITSET_CLEAR(resident, reg.reg() - 256 + i); ++ unsigned size = MIN2(DIV_ROUND_UP(bytes, 4), Start + Size - reg.reg()); ++ for (unsigned i = 0; i < size; i++) ++ BITSET_CLEAR(resident, reg.reg() - Start + i); + } + + uint8_t get(unsigned idx) +@@ -215,14 +217,14 @@ public: + + uint8_t get(PhysReg reg, unsigned offset = 0) + { +- assert(reg.reg() >= 256); +- return get(reg.reg() - 256 + offset); ++ assert(reg.reg() >= Start); ++ return get(reg.reg() - Start + offset); + } + +- void join_min(const VGPRCounterMap& other) ++ void join_min(const CounterMap& other) + { + unsigned i; +- BITSET_FOREACH_SET (i, other.resident, 256) { ++ BITSET_FOREACH_SET (i, other.resident, Size) { + if (BITSET_TEST(resident, i)) + val[i] = MIN2(val[i] + base, other.val[i] + other.base) - base; + else +@@ -231,13 +233,13 @@ public: + BITSET_OR(resident, resident, other.resident); + } + +- bool operator==(const VGPRCounterMap& other) const ++ bool operator==(const CounterMap& other) const + { + if (!BITSET_EQUAL(resident, other.resident)) + return false; + + unsigned i; +- BITSET_FOREACH_SET (i, other.resident, 256) { ++ BITSET_FOREACH_SET (i, other.resident, Size) { + if (!BITSET_TEST(resident, i)) + return false; + if (val[i] + base != other.val[i] + other.base) +@@ -245,8 +247,12 @@ public: + } + return true; + } ++ ++ unsigned size() const { return Size; } + }; + ++template using VGPRCounterMap = CounterMap<256, 256, Max>; ++ + struct NOP_ctx_gfx11 { + /* VcmpxPermlaneHazard */ + bool has_Vcmpx = false; +-- +2.48.1 + + +From 8cfd8daf2e0ed1783f3aeac6a5274bc70e1d103a Mon Sep 17 00:00:00 2001 +From: Rhys Perry +Date: Thu, 25 Jul 2024 21:36:29 +0100 +Subject: [PATCH 13/27] aco: minor CounterMap::operator== fix + +I don't think this matters for how we use CounterMap::operator==. + +The BITSET_TEST() was unnecessary because of the BITSET_EQUAL above. + +Signed-off-by: Rhys Perry +Reviewed-by: Georg Lehmann +Part-of: +--- + src/amd/compiler/aco_insert_NOPs.cpp | 4 +--- + 1 file changed, 1 insertion(+), 3 deletions(-) + +diff --git a/src/amd/compiler/aco_insert_NOPs.cpp b/src/amd/compiler/aco_insert_NOPs.cpp +index 3f4b5d6c85f..d6df9190864 100644 +--- a/src/amd/compiler/aco_insert_NOPs.cpp ++++ b/src/amd/compiler/aco_insert_NOPs.cpp +@@ -240,9 +240,7 @@ public: + + unsigned i; + BITSET_FOREACH_SET (i, other.resident, Size) { +- if (!BITSET_TEST(resident, i)) +- return false; +- if (val[i] + base != other.val[i] + other.base) ++ if (MIN2(val[i] + base, Max) != MIN2(other.val[i] + other.base, Max)) + return false; + } + return true; +-- +2.48.1 + + +From 62f14df1ca6435f30e13f44e7246fa763706cd9e Mon Sep 17 00:00:00 2001 +From: Rhys Perry +Date: Thu, 25 Jul 2024 21:40:02 +0100 +Subject: [PATCH 14/27] aco: workaround VALUReadSGPRHazard + +fossil-db (gfx1200): +Totals from 65112 (82.01% of 79395) affected shaders: +Instrs: 41732906 -> 42987198 (+3.01%); split: -0.00%, +3.01% +CodeSize: 222451964 -> 226942644 (+2.02%); split: -0.01%, +2.03% +Latency: 290411063 -> 290944688 (+0.18%); split: -0.00%, +0.18% +InvThroughput: 45854913 -> 45910275 (+0.12%); split: -0.00%, +0.12% + +Signed-off-by: Rhys Perry +Reviewed-by: Georg Lehmann +Part-of: +--- + src/amd/compiler/aco_insert_NOPs.cpp | 91 +++++++++++++++++++++++++--- + 1 file changed, 84 insertions(+), 7 deletions(-) + +diff --git a/src/amd/compiler/aco_insert_NOPs.cpp b/src/amd/compiler/aco_insert_NOPs.cpp +index d6df9190864..1bf53050e78 100644 +--- a/src/amd/compiler/aco_insert_NOPs.cpp ++++ b/src/amd/compiler/aco_insert_NOPs.cpp +@@ -273,6 +273,10 @@ struct NOP_ctx_gfx11 { + /* WMMAHazards */ + std::bitset<256> vgpr_written_by_wmma; + ++ /* VALUReadSGPRHazard */ ++ std::bitset sgpr_read_by_valu; /* SGPR pairs, excluding null, exec, m0 and scc */ ++ CounterMap<0, m0.reg(), 11> sgpr_read_by_valu_then_wr_by_salu; ++ + void join(const NOP_ctx_gfx11& other) + { + has_Vcmpx |= other.has_Vcmpx; +@@ -287,6 +291,8 @@ struct NOP_ctx_gfx11 { + sgpr_read_by_valu_as_lanemask_then_wr_by_salu |= + other.sgpr_read_by_valu_as_lanemask_then_wr_by_salu; + vgpr_written_by_wmma |= other.vgpr_written_by_wmma; ++ sgpr_read_by_valu |= other.sgpr_read_by_valu; ++ sgpr_read_by_valu_then_wr_by_salu.join_min(other.sgpr_read_by_valu_then_wr_by_salu); + } + + bool operator==(const NOP_ctx_gfx11& other) +@@ -302,7 +308,9 @@ struct NOP_ctx_gfx11 { + sgpr_read_by_valu_as_lanemask == other.sgpr_read_by_valu_as_lanemask && + sgpr_read_by_valu_as_lanemask_then_wr_by_salu == + other.sgpr_read_by_valu_as_lanemask_then_wr_by_salu && +- vgpr_written_by_wmma == other.vgpr_written_by_wmma; ++ vgpr_written_by_wmma == other.vgpr_written_by_wmma && ++ sgpr_read_by_valu == other.sgpr_read_by_valu && ++ sgpr_read_by_valu_then_wr_by_salu == other.sgpr_read_by_valu_then_wr_by_salu; + } + }; + +@@ -1527,6 +1535,48 @@ handle_instruction_gfx11(State& state, NOP_ctx_gfx11& ctx, aco_ptr& + } + } + } ++ } else { ++ /* VALUReadSGPRHazard ++ * VALU reads SGPR and later written by SALU cannot safely be read by VALU/SALU. ++ */ ++ if (instr->isVALU() || instr->isSALU()) { ++ unsigned expiry_count = instr->isSALU() ? 10 : 11; ++ for (Operand& op : instr->operands) { ++ if (sa_sdst == 0) ++ break; ++ ++ for (unsigned i = 0; i < op.size(); i++) { ++ unsigned reg = op.physReg() + i; ++ if (reg < ctx.sgpr_read_by_valu_then_wr_by_salu.size() && ++ ctx.sgpr_read_by_valu_then_wr_by_salu.get(reg) < expiry_count) { ++ bld.sopp(aco_opcode::s_waitcnt_depctr, 0xfffe); ++ sa_sdst = 0; ++ break; ++ } ++ } ++ } ++ } ++ ++ if (sa_sdst == 0) ++ ctx.sgpr_read_by_valu_then_wr_by_salu.reset(); ++ else if (instr->isSALU() && !instr->isSOPP()) ++ ctx.sgpr_read_by_valu_then_wr_by_salu.inc(); ++ ++ if (instr->isVALU()) { ++ for (const Operand& op : instr->operands) { ++ for (unsigned i = 0; i < DIV_ROUND_UP(op.size(), 2); i++) { ++ unsigned reg = (op.physReg() / 2) + i; ++ if (reg < ctx.sgpr_read_by_valu.size()) ++ ctx.sgpr_read_by_valu.set(reg); ++ } ++ } ++ } else if (instr->isSALU() && !instr->definitions.empty()) { ++ for (unsigned i = 0; i < instr->definitions[0].size(); i++) { ++ unsigned def_reg = instr->definitions[0].physReg() + i; ++ if ((def_reg / 2) < ctx.sgpr_read_by_valu.size() && ctx.sgpr_read_by_valu[def_reg / 2]) ++ ctx.sgpr_read_by_valu_then_wr_by_salu.set(def_reg); ++ } ++ } + } + + /* LdsDirectVMEMHazard +@@ -1683,6 +1733,15 @@ resolve_all_gfx11(State& state, NOP_ctx_gfx11& ctx, + } + } + ++ /* VALUReadSGPRHazard */ ++ if (state.program->gfx_level >= GFX12) { ++ for (unsigned i = 0; i < ctx.sgpr_read_by_valu_then_wr_by_salu.size(); i++) { ++ if (ctx.sgpr_read_by_valu_then_wr_by_salu.get(i) < 11) ++ waitcnt_depctr &= 0xfffe; ++ } ++ ctx.sgpr_read_by_valu_then_wr_by_salu.reset(); ++ } ++ + /* LdsDirectVMEMHazard */ + if (ctx.vgpr_used_by_vmem_load.any() || ctx.vgpr_used_by_vmem_store.any() || + ctx.vgpr_used_by_ds.any() || ctx.vgpr_used_by_vmem_sample.any() || +@@ -1758,7 +1817,7 @@ handle_block(Program* program, Ctx& ctx, Block& block) + + template Handle, ResolveAll Resolve> + void +-mitigate_hazards(Program* program) ++mitigate_hazards(Program* program, Ctx initial_ctx = Ctx()) + { + std::vector all_ctx(program->blocks.size()); + std::stack> loop_header_indices; +@@ -1767,6 +1826,9 @@ mitigate_hazards(Program* program) + Block& block = program->blocks[i]; + Ctx& ctx = all_ctx[i]; + ++ if (i == 0 || (block.kind & block_kind_resume)) ++ ctx = initial_ctx; ++ + if (block.kind & block_kind_loop_header) { + loop_header_indices.push(i); + } else if (block.kind & block_kind_loop_exit) { +@@ -1864,14 +1926,29 @@ required_export_priority(Program* program) + void + insert_NOPs(Program* program) + { +- if (program->gfx_level >= GFX11) +- mitigate_hazards(program); +- else if (program->gfx_level >= GFX10_3) ++ if (program->gfx_level >= GFX11) { ++ NOP_ctx_gfx11 initial_ctx; ++ ++ bool has_previous_part = ++ program->is_epilog || program->info.vs.has_prolog || program->info.ps.has_prolog || ++ (program->info.merged_shader_compiled_separately && program->stage.sw != SWStage::VS && ++ program->stage.sw != SWStage::TES) || program->stage == raytracing_cs; ++ if (program->gfx_level >= GFX12 && has_previous_part) { ++ /* resolve_all_gfx11 can't resolve VALUReadSGPRHazard entirely. We have to assume that any ++ * SGPR might have been read by VALU if there was a previous shader part. ++ */ ++ initial_ctx.sgpr_read_by_valu.flip(); ++ } ++ ++ mitigate_hazards(program, ++ initial_ctx); ++ } else if (program->gfx_level >= GFX10_3) { + ; /* no hazards/bugs to mitigate */ +- else if (program->gfx_level >= GFX10) ++ } else if (program->gfx_level >= GFX10) { + mitigate_hazards(program); +- else ++ } else { + mitigate_hazards(program); ++ } + + if (program->gfx_level == GFX11_5 && (program->stage.hw == AC_HW_NEXT_GEN_GEOMETRY_SHADER || + program->stage.hw == AC_HW_PIXEL_SHADER)) +-- +2.48.1 + + +From d6d58210c9658d759622022c7d196849c8f4e28d Mon Sep 17 00:00:00 2001 +From: Rhys Perry +Date: Tue, 26 Nov 2024 12:00:35 +0000 +Subject: [PATCH 15/27] aco/gfx12: insert wait between VMEM WaW + +https://github.com/llvm/llvm-project/pull/105549 + +fossil-db (gfx1200): +Totals from 1783 (2.25% of 79395) affected shaders: +Instrs: 7398391 -> 7404566 (+0.08%); split: -0.00%, +0.08% +CodeSize: 38862456 -> 38886364 (+0.06%); split: -0.00%, +0.06% +Latency: 83191513 -> 84211504 (+1.23%); split: -0.00%, +1.23% +InvThroughput: 15185936 -> 15345744 (+1.05%); split: -0.01%, +1.06% + +Signed-off-by: Rhys Perry +Part-of: +--- + src/amd/compiler/aco_insert_waitcnt.cpp | 8 ++++++-- + src/amd/compiler/tests/test_insert_waitcnt.cpp | 5 +++++ + 2 files changed, 11 insertions(+), 2 deletions(-) + +diff --git a/src/amd/compiler/aco_insert_waitcnt.cpp b/src/amd/compiler/aco_insert_waitcnt.cpp +index ae079dcd755..b19f49e8571 100644 +--- a/src/amd/compiler/aco_insert_waitcnt.cpp ++++ b/src/amd/compiler/aco_insert_waitcnt.cpp +@@ -385,9 +385,13 @@ check_instr(wait_ctx& ctx, wait_imm& wait, alu_delay_info& delay, Instruction* i + + wait_imm reg_imm = it->second.imm; + +- /* Vector Memory reads and writes return in the order they were issued */ ++ /* Vector Memory reads and writes decrease the counter in the order they were issued. ++ * Before GFX12, they also write VGPRs in order if they're of the same type. ++ * TODO: We can do this for GFX12 and different types for GFX11 if we know that the two ++ * VMEM loads do not write the same lanes. Since GFX11, we track VMEM operations on the ++ * linear CFG, so this is difficult */ + uint8_t vmem_type = get_vmem_type(ctx.gfx_level, instr); +- if (vmem_type) { ++ if (vmem_type && ctx.gfx_level < GFX12) { + wait_event event = get_vmem_event(ctx, instr, vmem_type); + wait_type type = (wait_type)(ffs(ctx.info->get_counters_for_event(event)) - 1); + if ((it->second.events & ctx.info->events[type]) == event && +diff --git a/src/amd/compiler/tests/test_insert_waitcnt.cpp b/src/amd/compiler/tests/test_insert_waitcnt.cpp +index edc34bfc7f8..c69772c5d05 100644 +--- a/src/amd/compiler/tests/test_insert_waitcnt.cpp ++++ b/src/amd/compiler/tests/test_insert_waitcnt.cpp +@@ -192,6 +192,7 @@ BEGIN_TEST(insert_waitcnt.waw.vmem_types) + + //>> p_unit_test 0 + //! v1: %0:v[4] = buffer_load_dword %0:s[0-3], %0:v[0], 0 ++ //~gfx12! s_wait_loadcnt imm:0 + //! v1: %0:v[4] = buffer_load_dword %0:s[0-3], %0:v[0], 0 + bld.pseudo(aco_opcode::p_unit_test, Operand::c32(0)); + bld.mubuf(aco_opcode::buffer_load_dword, def_v4, desc_s4, op_v0, Operand::zero(), 0, false); +@@ -220,6 +221,7 @@ BEGIN_TEST(insert_waitcnt.waw.vmem_types) + + //>> p_unit_test 3 + //! v1: %0:v[4] = image_sample %0:s[8-15], %0:s[0-3], v1: undef, %0:v[0] 1d ++ //~gfx12! s_wait_samplecnt imm:0 + //! v1: %0:v[4] = image_sample %0:s[8-15], %0:s[0-3], v1: undef, %0:v[0] 1d + bld.reset(program->create_and_insert_block()); + bld.pseudo(aco_opcode::p_unit_test, Operand::c32(3)); +@@ -249,6 +251,7 @@ BEGIN_TEST(insert_waitcnt.waw.vmem_types) + + //>> p_unit_test 6 + //! v1: %0:v[4] = image_bvh64_intersect_ray %0:s[0-3], s4: undef, v1: undef, %0:v[16-26] 1d ++ //~gfx12! s_wait_bvhcnt imm:0 + //! v1: %0:v[4] = image_bvh64_intersect_ray %0:s[0-3], s4: undef, v1: undef, %0:v[16-26] 1d + bld.reset(program->create_and_insert_block()); + bld.pseudo(aco_opcode::p_unit_test, Operand::c32(6)); +@@ -300,6 +303,7 @@ BEGIN_TEST(insert_waitcnt.waw.vmem_types) + //>> BB11 + //! /* logical preds: BB9, BB10, / linear preds: BB9, BB10, / kind: uniform, */ + //! p_unit_test 9 ++ //~gfx12! s_wait_loadcnt imm:0 + //! v1: %0:v[4] = buffer_load_dword %0:s[0-3], %0:v[0], 0 + bld.pseudo(aco_opcode::p_unit_test, Operand::c32(9)); + bld.mubuf(aco_opcode::buffer_load_dword, def_v4, desc_s4, op_v0, Operand::zero(), 0, false); +@@ -326,6 +330,7 @@ BEGIN_TEST(insert_waitcnt.waw.vmem_types) + //! /* logical preds: BB12, BB13, / linear preds: BB12, BB13, / kind: uniform, */ + //! p_unit_test 10 + //~gfx11! s_waitcnt vmcnt(0) ++ //~gfx12! s_wait_loadcnt imm:0 + //~gfx12! s_wait_samplecnt imm:0 + //! v1: %0:v[4] = buffer_load_dword %0:s[0-3], %0:v[0], 0 + bld.pseudo(aco_opcode::p_unit_test, Operand::c32(10)); +-- +2.48.1 + + +From 41edfd3cd9470ab4e54efb0b3af9eab7242f5e5f Mon Sep 17 00:00:00 2001 +From: Rhys Perry +Date: Wed, 27 Nov 2024 14:51:32 +0000 +Subject: [PATCH 16/27] aco: force linear for event_vmem_sample and + event_vmem_bvh + +I don't know if this issue affects GFX12, but workaround it anyway to be +safe. + +fossil-db (gfx1200): +Totals from 3463 (4.36% of 79395) affected shaders: +Instrs: 9794280 -> 9833253 (+0.40%); split: -0.00%, +0.40% +CodeSize: 52306040 -> 52457988 (+0.29%); split: -0.01%, +0.30% +Latency: 90549385 -> 93617517 (+3.39%); split: -0.00%, +3.39% +InvThroughput: 13189030 -> 13602942 (+3.14%); split: -0.00%, +3.14% + +Signed-off-by: Rhys Perry +Reviewed-by: Georg Lehmann +Part-of: +--- + src/amd/compiler/aco_insert_waitcnt.cpp | 3 ++- + 1 file changed, 2 insertions(+), 1 deletion(-) + +diff --git a/src/amd/compiler/aco_insert_waitcnt.cpp b/src/amd/compiler/aco_insert_waitcnt.cpp +index b19f49e8571..66ce8603ea4 100644 +--- a/src/amd/compiler/aco_insert_waitcnt.cpp ++++ b/src/amd/compiler/aco_insert_waitcnt.cpp +@@ -735,7 +735,8 @@ insert_wait_entry(wait_ctx& ctx, Definition def, wait_event event, uint8_t vmem_ + * Also, follow linear control flow for ALU because it's unlikely that the hardware does per-lane + * dependency checks. + */ +- uint32_t ds_vmem_events = event_lds | event_gds | event_vmem | event_flat; ++ uint32_t ds_vmem_events = ++ event_lds | event_gds | event_vmem | event_vmem_sample | event_vmem_bvh | event_flat; + uint32_t alu_events = event_trans | event_valu | event_salu; + bool force_linear = ctx.gfx_level >= GFX11 && (event & (ds_vmem_events | alu_events)); + +-- +2.48.1 + + +From c56875aa56504eab06cc07245b78f9acfb3a1bf6 Mon Sep 17 00:00:00 2001 +From: Georg Lehmann +Date: Thu, 5 Dec 2024 11:33:08 +0100 +Subject: [PATCH 17/27] aco/gfx12+: do not use v_pack_b32_f16 to pack untyped + data +MIME-Version: 1.0 +Content-Type: text/plain; charset=UTF-8 +Content-Transfer-Encoding: 8bit + +GFX12 removed IEEE_MODE, and made its signalling NaN quieting the default. + +Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/12251 +Reviewed-by: Daniel Schürmann +Part-of: +--- + src/amd/compiler/aco_lower_to_hw_instr.cpp | 8 +++++++- + 1 file changed, 7 insertions(+), 1 deletion(-) + +diff --git a/src/amd/compiler/aco_lower_to_hw_instr.cpp b/src/amd/compiler/aco_lower_to_hw_instr.cpp +index ed955a32a1e..bceaa0d25e9 100644 +--- a/src/amd/compiler/aco_lower_to_hw_instr.cpp ++++ b/src/amd/compiler/aco_lower_to_hw_instr.cpp +@@ -1562,9 +1562,15 @@ do_pack_2x16(lower_context* ctx, Builder& bld, Definition def, Operand lo, Opera + return; + } + ++ /* v_pack_b32_f16 can be used for bit exact copies if: ++ * - fp16 input denorms are enabled, otherwise they get flushed to zero ++ * - signalling input NaNs are kept, which is the case with IEEE_MODE=0 ++ * GFX12+ always quiets signalling NaNs, IEEE_MODE was removed ++ */ + bool can_use_pack = (ctx->block->fp_mode.denorm16_64 & fp_denorm_keep_in) && + (ctx->program->gfx_level >= GFX10 || +- (ctx->program->gfx_level >= GFX9 && !lo.isLiteral() && !hi.isLiteral())); ++ (ctx->program->gfx_level >= GFX9 && !lo.isLiteral() && !hi.isLiteral())) && ++ ctx->program->gfx_level < GFX12; + + if (can_use_pack) { + Instruction* instr = bld.vop3(aco_opcode::v_pack_b32_f16, def, lo, hi); +-- +2.48.1 + + +From 04d7f9ed9e5b0cd98c32a9d8b35c92698e427dfa Mon Sep 17 00:00:00 2001 +From: Rhys Perry +Date: Wed, 4 Dec 2024 13:36:04 +0000 +Subject: [PATCH 18/27] aco: don't CSE p_shader_cycles_hi_lo_hi + +Signed-off-by: Rhys Perry +Reviewed-by: Georg Lehmann +Fixes: fae2a85d57a4 ("aco/gfx12: implement subgroup shader clock") +Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/12243 +Part-of: +--- + src/amd/compiler/aco_opt_value_numbering.cpp | 1 + + 1 file changed, 1 insertion(+) + +diff --git a/src/amd/compiler/aco_opt_value_numbering.cpp b/src/amd/compiler/aco_opt_value_numbering.cpp +index e040221be61..e0c8eeffbbb 100644 +--- a/src/amd/compiler/aco_opt_value_numbering.cpp ++++ b/src/amd/compiler/aco_opt_value_numbering.cpp +@@ -306,6 +306,7 @@ can_eliminate(aco_ptr& instr) + if (instr->definitions.empty() || instr->opcode == aco_opcode::p_phi || + instr->opcode == aco_opcode::p_linear_phi || + instr->opcode == aco_opcode::p_pops_gfx9_add_exiting_wave_id || ++ instr->opcode == aco_opcode::p_shader_cycles_hi_lo_hi || + instr->definitions[0].isNoCSE()) + return false; + +-- +2.48.1 + + +From 8de2e82ca04d973c70eec66c650809a89000e90d Mon Sep 17 00:00:00 2001 +From: Georg Lehmann +Date: Mon, 9 Dec 2024 13:40:49 +0100 +Subject: [PATCH 19/27] aco/gfx12: don't assume memory operations complete in + order + +Reviewed-by: Rhys Perry +Part-of: +--- + src/amd/compiler/aco_insert_waitcnt.cpp | 4 ++-- + 1 file changed, 2 insertions(+), 2 deletions(-) + +diff --git a/src/amd/compiler/aco_insert_waitcnt.cpp b/src/amd/compiler/aco_insert_waitcnt.cpp +index 66ce8603ea4..22822d92f24 100644 +--- a/src/amd/compiler/aco_insert_waitcnt.cpp ++++ b/src/amd/compiler/aco_insert_waitcnt.cpp +@@ -449,9 +449,9 @@ perform_barrier(wait_ctx& ctx, wait_imm& imm, memory_sync_info sync, unsigned se + if (bar_scope_lds <= subgroup_scope) + events &= ~event_lds; + +- /* in non-WGP, the L1 (L0 on GFX10+) cache keeps all memory operations ++ /* Until GFX12, in non-WGP, the L1 (L0 on GFX10+) cache keeps all memory operations + * in-order for the same workgroup */ +- if (!ctx.program->wgp_mode && sync.scope <= scope_workgroup) ++ if (ctx.gfx_level < GFX12 && !ctx.program->wgp_mode && sync.scope <= scope_workgroup) + events &= ~(event_vmem | event_vmem_store | event_smem); + + if (events) +-- +2.48.1 + + +From 054163a6be4c130b5a95fb5fa2d39d64ea8ee259 Mon Sep 17 00:00:00 2001 +From: Qiang Yu +Date: Mon, 9 Dec 2024 10:01:21 +0800 +Subject: [PATCH 20/27] aco: enable gfx12 support for radeonsi + +Part-of: +--- + src/amd/compiler/aco_interface.cpp | 1 + + 1 file changed, 1 insertion(+) + +diff --git a/src/amd/compiler/aco_interface.cpp b/src/amd/compiler/aco_interface.cpp +index ab518aa14ef..8508cc0445e 100644 +--- a/src/amd/compiler/aco_interface.cpp ++++ b/src/amd/compiler/aco_interface.cpp +@@ -422,6 +422,7 @@ aco_is_gpu_supported(const struct radeon_info* info) + case GFX10_3: + case GFX11: + case GFX11_5: ++ case GFX12: + return true; + default: + return false; +-- +2.48.1 + + +From a3006214823ffacb34383fad792b3184a171c135 Mon Sep 17 00:00:00 2001 +From: Samuel Pitoiset +Date: Fri, 3 Jan 2025 08:29:34 -0800 +Subject: [PATCH 21/27] aco: fix VS prologs on GFX12 + +MTBUF/MUBUF instructions must use zero for SOFFSET, use const_offset +instead. + +Signed-off-by: Samuel Pitoiset +Part-of: +--- + .../compiler/aco_instruction_selection.cpp | 30 +++++++++++++++---- + 1 file changed, 25 insertions(+), 5 deletions(-) + +diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp +index 93a88c74d30..5e4806f608f 100644 +--- a/src/amd/compiler/aco_instruction_selection.cpp ++++ b/src/amd/compiler/aco_instruction_selection.cpp +@@ -12903,8 +12903,16 @@ load_unaligned_vs_attrib(Builder& bld, PhysReg dst, Operand desc, Operand index, + } else { + for (unsigned i = 0; i < size; i++) { + Definition def(i ? scratch.advance(i * 4 - 4) : dst, v1); +- bld.mubuf(aco_opcode::buffer_load_ubyte, def, desc, index, Operand::c32(offset + i), 0, +- false, true); ++ unsigned soffset = 0, const_offset = 0; ++ ++ if (bld.program->gfx_level >= GFX12) { ++ const_offset = offset + i; ++ } else { ++ soffset = offset + i; ++ } ++ ++ bld.mubuf(aco_opcode::buffer_load_ubyte, def, desc, index, Operand::c32(soffset), ++ const_offset, false, true); + } + } + +@@ -13082,6 +13090,17 @@ select_vs_prolog(Program* program, const struct aco_vs_prolog_info* pinfo, ac_sh + for (unsigned j = 0; j < (vtx_info->chan_byte_size ? vtx_info->num_channels : 1); j++) { + bool post_shuffle = pinfo->post_shuffle & (1u << loc); + unsigned offset = vtx_info->chan_byte_size * (post_shuffle && j < 3 ? 2 - j : j); ++ unsigned soffset = 0, const_offset = 0; ++ ++ /* We need to use soffset on GFX6-7 to avoid being considered ++ * out-of-bounds when offset>=stride. GFX12 doesn't support a ++ * non-zero constant soffset. ++ */ ++ if (program->gfx_level >= GFX12) { ++ const_offset = offset; ++ } else { ++ soffset = offset; ++ } + + if ((pinfo->unaligned_mask & (1u << loc)) && vtx_info->chan_byte_size <= 4) + load_unaligned_vs_attrib(bld, dest.advance(j * 4u), Operand(cur_desc, s4), +@@ -13089,11 +13108,12 @@ select_vs_prolog(Program* program, const struct aco_vs_prolog_info* pinfo, ac_sh + else if (vtx_info->chan_byte_size == 8) + bld.mtbuf(aco_opcode::tbuffer_load_format_xy, + Definition(dest.advance(j * 8u), v2), Operand(cur_desc, s4), +- fetch_index, Operand::c32(offset), dfmt, nfmt, 0, false, true); ++ fetch_index, Operand::c32(soffset), dfmt, nfmt, const_offset, false, ++ true); + else + bld.mtbuf(aco_opcode::tbuffer_load_format_x, Definition(dest.advance(j * 4u), v1), +- Operand(cur_desc, s4), fetch_index, Operand::c32(offset), dfmt, nfmt, +- 0, false, true); ++ Operand(cur_desc, s4), fetch_index, Operand::c32(soffset), dfmt, nfmt, ++ const_offset, false, true); + } + + unsigned slots = vtx_info->chan_byte_size == 8 && vtx_info->num_channels > 2 ? 2 : 1; +-- +2.48.1 + + +From 9aeed54b7f1050648612f59844c4cfbedfe1b182 Mon Sep 17 00:00:00 2001 +From: Samuel Pitoiset +Date: Wed, 8 Jan 2025 02:03:09 -0800 +Subject: [PATCH 22/27] aco: always use ds_bpermute for shuffle/rotate on GFX12 + +ds_bpermute supports both 32 and 64 lanes now. + +Signed-off-by: Samuel Pitoiset +Part-of: +--- + src/amd/compiler/aco_instruction_selection.cpp | 13 +++++++------ + 1 file changed, 7 insertions(+), 6 deletions(-) + +diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp +index 5e4806f608f..8ab9c18ee2a 100644 +--- a/src/amd/compiler/aco_instruction_selection.cpp ++++ b/src/amd/compiler/aco_instruction_selection.cpp +@@ -189,9 +189,9 @@ emit_bpermute(isel_context* ctx, Builder& bld, Temp index, Temp data) + + return bld.pseudo(aco_opcode::p_bpermute_readlane, bld.def(v1), bld.def(bld.lm), + bld.def(bld.lm, vcc), index_op, input_data); +- } else if (ctx->options->gfx_level >= GFX10 && ctx->program->wave_size == 64) { +- +- /* GFX10 wave64 mode: emulate full-wave bpermute */ ++ } else if (ctx->options->gfx_level >= GFX10 && ctx->options->gfx_level <= GFX11_5 && ++ ctx->program->wave_size == 64) { ++ /* GFX10-11.5 wave64 mode: emulate full-wave bpermute */ + Temp index_is_lo = + bld.vopc(aco_opcode::v_cmp_ge_u32, bld.def(bld.lm), Operand::c32(31u), index); + Builder::Result index_is_lo_split = +@@ -221,7 +221,7 @@ emit_bpermute(isel_context* ctx, Builder& bld, Temp index, Temp data) + same_half); + } + } else { +- /* GFX8-9 or GFX10 wave32: bpermute works normally */ ++ /* wave32 or GFX8-9, GFX12+: bpermute works normally */ + Temp index_x4 = bld.vop2(aco_opcode::v_lshlrev_b32, bld.def(v1), Operand::c32(2u), index); + return bld.ds(aco_opcode::ds_bpermute_b32, bld.def(v1), index_x4, data); + } +@@ -8747,8 +8747,9 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr) + Temp tid = emit_mbcnt(ctx, bld.tmp(v1)); + Temp src_lane = bld.vadd32(bld.def(v1), tid, delta); + +- if (ctx->program->gfx_level >= GFX10 && cluster_size == 32) { +- /* ds_bpermute is restricted to 32 lanes on GFX10+. */ ++ if (ctx->program->gfx_level >= GFX10 && ctx->program->gfx_level <= GFX11_5 && ++ cluster_size == 32) { ++ /* ds_bpermute is restricted to 32 lanes on GFX10-GFX11.5. */ + Temp index_x4 = + bld.vop2(aco_opcode::v_lshlrev_b32, bld.def(v1), Operand::c32(2u), src_lane); + tmp = bld.ds(aco_opcode::ds_bpermute_b32, bld.def(v1), index_x4, src); +-- +2.48.1 + + +From 4929c358e0d44cb25cb486cc634ee96536a674db Mon Sep 17 00:00:00 2001 +From: Georg Lehmann +Date: Sun, 19 Jan 2025 10:17:12 +0100 +Subject: [PATCH 23/27] aco: update is_dual_issue_capable for gfx11.5+ +MIME-Version: 1.0 +Content-Type: text/plain; charset=UTF-8 +Content-Transfer-Encoding: 8bit + +Reviewed-by: Daniel Schürmann +Reviewed-by: Daniel Schürmann +Part-of: +--- + src/amd/compiler/aco_statistics.cpp | 7 ++++++- + 1 file changed, 6 insertions(+), 1 deletion(-) + +diff --git a/src/amd/compiler/aco_statistics.cpp b/src/amd/compiler/aco_statistics.cpp +index 552bfd2cee4..00934b89ceb 100644 +--- a/src/amd/compiler/aco_statistics.cpp ++++ b/src/amd/compiler/aco_statistics.cpp +@@ -142,7 +142,12 @@ is_dual_issue_capable(const Program& program, const Instruction& instr) + } + return false; + } +- default: return false; ++ default: ++ if (instr.isVINTERP_INREG()) ++ return program.gfx_level >= GFX11_5; ++ if (instr.isVOPC() && instr_info.classes[(int)instr.opcode] == instr_class::valu32) ++ return program.gfx_level == GFX11_5; ++ return false; + } + } + +-- +2.48.1 + + +From 9f0671af164968432750d7b1be5da6d838765179 Mon Sep 17 00:00:00 2001 +From: Samuel Pitoiset +Date: Thu, 12 Sep 2024 15:46:29 +0200 +Subject: [PATCH 24/27] aco,radv,radeonsi: move has_epilog to the fragment + shader info + +Signed-off-by: Samuel Pitoiset +Part-of: +--- + src/amd/compiler/aco_assembler.cpp | 3 ++- + src/amd/compiler/aco_instruction_selection.cpp | 8 ++++---- + src/amd/compiler/aco_interface.cpp | 2 +- + src/amd/compiler/aco_shader_info.h | 2 +- + src/amd/vulkan/radv_aco_shader_info.h | 2 +- + src/amd/vulkan/radv_cmd_buffer.c | 6 +++--- + src/amd/vulkan/radv_pipeline.c | 2 +- + src/amd/vulkan/radv_pipeline_graphics.c | 2 +- + src/amd/vulkan/radv_shader_args.c | 2 +- + src/amd/vulkan/radv_shader_info.c | 6 +++--- + src/amd/vulkan/radv_shader_info.h | 3 ++- + src/gallium/drivers/radeonsi/si_shader_aco.c | 2 +- + 12 files changed, 21 insertions(+), 19 deletions(-) + +diff --git a/src/amd/compiler/aco_assembler.cpp b/src/amd/compiler/aco_assembler.cpp +index 714648b252c..cf57a7e95c6 100644 +--- a/src/amd/compiler/aco_assembler.cpp ++++ b/src/amd/compiler/aco_assembler.cpp +@@ -1727,7 +1727,8 @@ emit_program(Program* program, std::vector& code, std::vectorinfo.merged_shader_compiled_separately; + + /* Prolog has no exports. */ +- if (!program->is_prolog && !program->info.has_epilog && !is_separately_compiled_ngg_vs_or_es && ++ if (!program->is_prolog && !program->info.ps.has_epilog && ++ !is_separately_compiled_ngg_vs_or_es && + (program->stage.hw == AC_HW_VERTEX_SHADER || program->stage.hw == AC_HW_PIXEL_SHADER || + program->stage.hw == AC_HW_NEXT_GEN_GEOMETRY_SHADER)) + fix_exports(ctx, code, program); +diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp +index 8ab9c18ee2a..987ba7d22ae 100644 +--- a/src/amd/compiler/aco_instruction_selection.cpp ++++ b/src/amd/compiler/aco_instruction_selection.cpp +@@ -177,7 +177,7 @@ emit_bpermute(isel_context* ctx, Builder& bld, Temp index, Temp data) + const bool avoid_shared_vgprs = + ctx->options->gfx_level >= GFX10 && ctx->options->gfx_level < GFX11 && + ctx->program->wave_size == 64 && +- (ctx->program->info.has_epilog || ctx->program->info.merged_shader_compiled_separately || ++ (ctx->program->info.ps.has_epilog || ctx->program->info.merged_shader_compiled_separately || + ctx->program->info.vs.has_prolog || ctx->stage == raytracing_cs); + + if (ctx->options->gfx_level <= GFX7 || avoid_shared_vgprs) { +@@ -5611,7 +5611,7 @@ store_output_to_temps(isel_context* ctx, nir_intrinsic_instr* instr) + idx++; + } + +- if (ctx->stage == fragment_fs && ctx->program->info.has_epilog && base >= FRAG_RESULT_DATA0) { ++ if (ctx->stage == fragment_fs && ctx->program->info.ps.has_epilog && base >= FRAG_RESULT_DATA0) { + unsigned index = base - FRAG_RESULT_DATA0; + + if (nir_intrinsic_src_type(instr) == nir_type_float16) { +@@ -11927,7 +11927,7 @@ select_shader(isel_context& ctx, nir_shader* nir, const bool need_startpgm, cons + nir_function_impl* func = nir_shader_get_entrypoint(nir); + visit_cf_list(&ctx, &func->body); + +- if (ctx.program->info.has_epilog) { ++ if (ctx.program->info.ps.has_epilog) { + if (ctx.stage == fragment_fs) { + if (ctx.options->is_opengl) + create_fs_end_for_epilog(&ctx); +@@ -11965,7 +11965,7 @@ select_shader(isel_context& ctx, nir_shader* nir, const bool need_startpgm, cons + append_logical_end(ctx.block); + ctx.block->kind |= block_kind_uniform; + +- if ((!program->info.has_epilog && !is_first_stage_of_merged_shader) || ++ if ((!program->info.ps.has_epilog && !is_first_stage_of_merged_shader) || + (nir->info.stage == MESA_SHADER_TESS_CTRL && program->gfx_level >= GFX9)) { + Builder(program, ctx.block).sopp(aco_opcode::s_endpgm); + } +diff --git a/src/amd/compiler/aco_interface.cpp b/src/amd/compiler/aco_interface.cpp +index 8508cc0445e..64396cfbc75 100644 +--- a/src/amd/compiler/aco_interface.cpp ++++ b/src/amd/compiler/aco_interface.cpp +@@ -277,7 +277,7 @@ aco_compile_shader(const struct aco_compiler_options* options, const struct aco_ + /* OpenGL combine multi shader parts into one continous code block, + * so only last part need the s_endpgm instruction. + */ +- bool append_endpgm = !(options->is_opengl && info->has_epilog); ++ bool append_endpgm = !(options->is_opengl && info->ps.has_epilog); + unsigned exec_size = emit_program(program.get(), code, &symbols, append_endpgm); + + if (program->collect_statistics) +diff --git a/src/amd/compiler/aco_shader_info.h b/src/amd/compiler/aco_shader_info.h +index 1cc2bc11db4..7f6fe7a5e9e 100644 +--- a/src/amd/compiler/aco_shader_info.h ++++ b/src/amd/compiler/aco_shader_info.h +@@ -102,7 +102,6 @@ struct aco_shader_info { + bool has_ngg_early_prim_export; + bool image_2d_view_of_3d; + unsigned workgroup_size; +- bool has_epilog; /* Only for TCS or PS. */ + bool merged_shader_compiled_separately; /* GFX9+ */ + struct ac_arg next_stage_pc; + struct ac_arg epilog_pc; /* Vulkan only */ +@@ -127,6 +126,7 @@ struct aco_shader_info { + uint32_t num_interp; + unsigned spi_ps_input_ena; + unsigned spi_ps_input_addr; ++ bool has_epilog; + + /* OpenGL only */ + struct ac_arg alpha_reference; +diff --git a/src/amd/vulkan/radv_aco_shader_info.h b/src/amd/vulkan/radv_aco_shader_info.h +index b9665b2b20f..912d6b6ca4e 100644 +--- a/src/amd/vulkan/radv_aco_shader_info.h ++++ b/src/amd/vulkan/radv_aco_shader_info.h +@@ -31,7 +31,7 @@ radv_aco_convert_shader_info(struct aco_shader_info *aco_info, const struct radv + ASSIGN_FIELD(has_ngg_culling); + ASSIGN_FIELD(has_ngg_early_prim_export); + ASSIGN_FIELD(workgroup_size); +- ASSIGN_FIELD(has_epilog); ++ ASSIGN_FIELD(ps.has_epilog); + ASSIGN_FIELD(merged_shader_compiled_separately); + ASSIGN_FIELD(vs.tcs_in_out_eq); + ASSIGN_FIELD(vs.tcs_temp_only_input_mask); +diff --git a/src/amd/vulkan/radv_cmd_buffer.c b/src/amd/vulkan/radv_cmd_buffer.c +index d0cf1f1123b..671b35cb972 100644 +--- a/src/amd/vulkan/radv_cmd_buffer.c ++++ b/src/amd/vulkan/radv_cmd_buffer.c +@@ -8096,7 +8096,7 @@ radv_bind_fragment_shader(struct radv_cmd_buffer *cmd_buffer, const struct radv_ + cmd_buffer->state.dirty |= RADV_CMD_DIRTY_DB_SHADER_CONTROL; + + /* Re-emit the PS epilog when a new fragment shader is bound. */ +- if (ps->info.has_epilog) ++ if (ps->info.ps.has_epilog) + cmd_buffer->state.emitted_ps_epilog = NULL; + } + +@@ -10703,7 +10703,7 @@ radv_emit_all_graphics_states(struct radv_cmd_buffer *cmd_buffer, const struct r + struct radv_shader_part *ps_epilog = NULL; + + if (cmd_buffer->state.shaders[MESA_SHADER_FRAGMENT] && +- cmd_buffer->state.shaders[MESA_SHADER_FRAGMENT]->info.has_epilog) { ++ cmd_buffer->state.shaders[MESA_SHADER_FRAGMENT]->info.ps.has_epilog) { + if ((cmd_buffer->state.emitted_graphics_pipeline != cmd_buffer->state.graphics_pipeline || + ((cmd_buffer->state.dirty & (RADV_CMD_DIRTY_GRAPHICS_SHADERS | RADV_CMD_DIRTY_FRAMEBUFFER)) || + (cmd_buffer->state.dirty_dynamic & +@@ -10921,7 +10921,7 @@ radv_bind_graphics_shaders(struct radv_cmd_buffer *cmd_buffer) + } + + const struct radv_shader *ps = cmd_buffer->state.shaders[MESA_SHADER_FRAGMENT]; +- if (ps && !ps->info.has_epilog) { ++ if (ps && !ps->info.ps.has_epilog) { + uint32_t col_format = 0, cb_shader_mask = 0; + if (radv_needs_null_export_workaround(device, ps, 0)) + col_format = V_028714_SPI_SHADER_32_R; +diff --git a/src/amd/vulkan/radv_pipeline.c b/src/amd/vulkan/radv_pipeline.c +index e07737cd742..738fc736b2c 100644 +--- a/src/amd/vulkan/radv_pipeline.c ++++ b/src/amd/vulkan/radv_pipeline.c +@@ -551,7 +551,7 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_graphics_stat + .use_aco = !radv_use_llvm_for_stage(pdev, stage->stage), + .uses_discard = true, + .alpha_func = COMPARE_FUNC_ALWAYS, +- .no_color_export = stage->info.has_epilog, ++ .no_color_export = stage->info.ps.has_epilog, + .no_depth_export = stage->info.ps.exports_mrtz_via_epilog, + + .bc_optimize_for_persp = G_0286CC_PERSP_CENTER_ENA(stage->info.ps.spi_ps_input_ena) && +diff --git a/src/amd/vulkan/radv_pipeline_graphics.c b/src/amd/vulkan/radv_pipeline_graphics.c +index ed380341ba5..daab8ce2ed6 100644 +--- a/src/amd/vulkan/radv_pipeline_graphics.c ++++ b/src/amd/vulkan/radv_pipeline_graphics.c +@@ -3150,7 +3150,7 @@ radv_graphics_pipeline_init(struct radv_graphics_pipeline *pipeline, struct radv + radv_pipeline_init_dynamic_state(device, pipeline, &gfx_state.vk, pCreateInfo); + + const struct radv_shader *ps = pipeline->base.shaders[MESA_SHADER_FRAGMENT]; +- if (ps && !ps->info.has_epilog) { ++ if (ps && !ps->info.ps.has_epilog) { + pipeline->spi_shader_col_format = ps->info.ps.spi_shader_col_format; + pipeline->cb_shader_mask = ps->info.ps.cb_shader_mask; + } +diff --git a/src/amd/vulkan/radv_shader_args.c b/src/amd/vulkan/radv_shader_args.c +index 10f1ee3ad07..479bc49c803 100644 +--- a/src/amd/vulkan/radv_shader_args.c ++++ b/src/amd/vulkan/radv_shader_args.c +@@ -846,7 +846,7 @@ declare_shader_args(const struct radv_device *device, const struct radv_graphics + case MESA_SHADER_FRAGMENT: + declare_global_input_sgprs(gfx_level, info, user_sgpr_info, args); + +- if (info->has_epilog) { ++ if (info->ps.has_epilog) { + add_ud_arg(args, 1, AC_ARG_INT, &args->epilog_pc, AC_UD_EPILOG_PC); + } + +diff --git a/src/amd/vulkan/radv_shader_info.c b/src/amd/vulkan/radv_shader_info.c +index c8c0b77201d..a1050a9ec07 100644 +--- a/src/amd/vulkan/radv_shader_info.c ++++ b/src/amd/vulkan/radv_shader_info.c +@@ -972,9 +972,9 @@ gather_shader_info_fs(const struct radv_device *device, const nir_shader *nir, + info->ps.spi_ps_input_addr &= C_02865C_COVERAGE_TO_SHADER_SELECT; + } + +- info->has_epilog = gfx_state->ps.has_epilog && info->ps.colors_written; ++ info->ps.has_epilog = gfx_state->ps.has_epilog && info->ps.colors_written; + +- if (!info->has_epilog) { ++ if (!info->ps.has_epilog) { + info->ps.mrt0_is_dual_src = gfx_state->ps.epilog.mrt0_is_dual_src; + info->ps.spi_shader_col_format = gfx_state->ps.epilog.spi_shader_col_format; + +@@ -988,7 +988,7 @@ gather_shader_info_fs(const struct radv_device *device, const nir_shader *nir, + (info->ps.color0_written & 0x8) && (info->ps.writes_z || info->ps.writes_stencil || info->ps.writes_sample_mask); + + info->ps.exports_mrtz_via_epilog = +- info->has_epilog && gfx_state->ps.exports_mrtz_via_epilog && export_alpha_and_mrtz; ++ info->ps.has_epilog && gfx_state->ps.exports_mrtz_via_epilog && export_alpha_and_mrtz; + + if (!info->ps.exports_mrtz_via_epilog) { + info->ps.writes_mrt0_alpha = gfx_state->ms.alpha_to_coverage_via_mrtz && export_alpha_and_mrtz; +diff --git a/src/amd/vulkan/radv_shader_info.h b/src/amd/vulkan/radv_shader_info.h +index e072c9ffd5b..0ffa972b078 100644 +--- a/src/amd/vulkan/radv_shader_info.h ++++ b/src/amd/vulkan/radv_shader_info.h +@@ -106,7 +106,6 @@ struct radv_shader_info { + uint32_t user_data_0; + bool inputs_linked; + bool outputs_linked; +- bool has_epilog; /* Only for TCS or PS */ + bool merged_shader_compiled_separately; /* GFX9+ */ + bool force_indirect_desc_sets; + +@@ -213,6 +212,8 @@ struct radv_shader_info { + bool load_provoking_vtx; + bool load_rasterization_prim; + bool force_sample_iter_shading_rate; ++ bool uses_fbfetch_output; ++ bool has_epilog; + } ps; + struct { + bool uses_grid_size; +diff --git a/src/gallium/drivers/radeonsi/si_shader_aco.c b/src/gallium/drivers/radeonsi/si_shader_aco.c +index 67ba140ed67..232104954fb 100644 +--- a/src/gallium/drivers/radeonsi/si_shader_aco.c ++++ b/src/gallium/drivers/radeonsi/si_shader_aco.c +@@ -100,7 +100,7 @@ si_fill_aco_shader_info(struct si_shader *shader, struct aco_shader_info *info, + info->ps.spi_ps_input_ena = shader->config.spi_ps_input_ena; + info->ps.spi_ps_input_addr = shader->config.spi_ps_input_addr; + info->ps.alpha_reference = args->alpha_reference; +- info->has_epilog = !shader->is_monolithic; ++ info->ps.has_epilog = !shader->is_monolithic; + break; + default: + break; +-- +2.48.1 + + +From 82939c7825d3a45d5e992e10491e43183e511a9d Mon Sep 17 00:00:00 2001 +From: Rhys Perry +Date: Thu, 24 Oct 2024 11:01:46 +0100 +Subject: [PATCH 25/27] aco,radv,radeonsi: add aco_shader_info::ps::has_prolog + +Signed-off-by: Rhys Perry +Reviewed-by: Georg Lehmann +Part-of: +--- + src/amd/compiler/aco_shader_info.h | 1 + + src/amd/vulkan/radv_aco_shader_info.h | 1 + + src/gallium/drivers/radeonsi/si_shader_aco.c | 1 + + 3 files changed, 3 insertions(+) + +diff --git a/src/amd/compiler/aco_shader_info.h b/src/amd/compiler/aco_shader_info.h +index 7f6fe7a5e9e..ad615e19992 100644 +--- a/src/amd/compiler/aco_shader_info.h ++++ b/src/amd/compiler/aco_shader_info.h +@@ -126,6 +126,7 @@ struct aco_shader_info { + uint32_t num_interp; + unsigned spi_ps_input_ena; + unsigned spi_ps_input_addr; ++ bool has_prolog; + bool has_epilog; + + /* OpenGL only */ +diff --git a/src/amd/vulkan/radv_aco_shader_info.h b/src/amd/vulkan/radv_aco_shader_info.h +index 912d6b6ca4e..9a0b4d562e0 100644 +--- a/src/amd/vulkan/radv_aco_shader_info.h ++++ b/src/amd/vulkan/radv_aco_shader_info.h +@@ -41,6 +41,7 @@ radv_aco_convert_shader_info(struct aco_shader_info *aco_info, const struct radv + ASSIGN_FIELD(cs.uses_full_subgroups); + aco_info->ps.spi_ps_input_ena = radv->ps.spi_ps_input_ena; + aco_info->ps.spi_ps_input_addr = radv->ps.spi_ps_input_addr; ++ aco_info->ps.has_prolog = false; + aco_info->gfx9_gs_ring_lds_size = radv->gs_ring_info.lds_size; + aco_info->is_trap_handler_shader = radv->type == RADV_SHADER_TYPE_TRAP_HANDLER; + aco_info->image_2d_view_of_3d = radv_key->image_2d_view_of_3d; +diff --git a/src/gallium/drivers/radeonsi/si_shader_aco.c b/src/gallium/drivers/radeonsi/si_shader_aco.c +index 232104954fb..dd0b3759a2b 100644 +--- a/src/gallium/drivers/radeonsi/si_shader_aco.c ++++ b/src/gallium/drivers/radeonsi/si_shader_aco.c +@@ -100,6 +100,7 @@ si_fill_aco_shader_info(struct si_shader *shader, struct aco_shader_info *info, + info->ps.spi_ps_input_ena = shader->config.spi_ps_input_ena; + info->ps.spi_ps_input_addr = shader->config.spi_ps_input_addr; + info->ps.alpha_reference = args->alpha_reference; ++ info->ps.has_prolog = !shader->is_monolithic; + info->ps.has_epilog = !shader->is_monolithic; + break; + default: +-- +2.48.1 + + +From af418b1dcedd56326872945c18dff6dae79ebfdf Mon Sep 17 00:00:00 2001 +From: Samuel Pitoiset +Date: Thu, 12 Sep 2024 15:44:43 +0200 +Subject: [PATCH 26/27] radv,radeonsi: remove remaining occurrences of TCS + epilog + +TCS epilog has been removed few months ago. + +Signed-off-by: Samuel Pitoiset +Part-of: +--- + src/amd/vulkan/radv_shader_args.c | 8 -------- + src/gallium/drivers/radeonsi/si_shader_aco.c | 1 - + 2 files changed, 9 deletions(-) + +diff --git a/src/amd/vulkan/radv_shader_args.c b/src/amd/vulkan/radv_shader_args.c +index 479bc49c803..100797c7389 100644 +--- a/src/amd/vulkan/radv_shader_args.c ++++ b/src/amd/vulkan/radv_shader_args.c +@@ -683,10 +683,6 @@ declare_shader_args(const struct radv_device *device, const struct radv_graphics + add_ud_arg(args, 1, AC_ARG_INT, &args->tcs_offchip_layout, AC_UD_TCS_OFFCHIP_LAYOUT); + } + +- if (info->has_epilog) { +- add_ud_arg(args, 1, AC_ARG_INT, &args->epilog_pc, AC_UD_EPILOG_PC); +- } +- + ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.tcs_patch_id); + ac_add_arg(&args->ac, AC_ARG_VGPR, 1, AC_ARG_INT, &args->ac.tcs_rel_ids); + +@@ -703,10 +699,6 @@ declare_shader_args(const struct radv_device *device, const struct radv_graphics + add_ud_arg(args, 1, AC_ARG_INT, &args->tcs_offchip_layout, AC_UD_TCS_OFFCHIP_LAYOUT); + } + +- if (info->has_epilog) { +- add_ud_arg(args, 1, AC_ARG_INT, &args->epilog_pc, AC_UD_EPILOG_PC); +- } +- + ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->ac.tess_offchip_offset); + ac_add_arg(&args->ac, AC_ARG_SGPR, 1, AC_ARG_INT, &args->ac.tcs_factor_offset); + if (args->explicit_scratch_args) { +diff --git a/src/gallium/drivers/radeonsi/si_shader_aco.c b/src/gallium/drivers/radeonsi/si_shader_aco.c +index dd0b3759a2b..dacfe15bb5a 100644 +--- a/src/gallium/drivers/radeonsi/si_shader_aco.c ++++ b/src/gallium/drivers/radeonsi/si_shader_aco.c +@@ -88,7 +88,6 @@ si_fill_aco_shader_info(struct si_shader *shader, struct aco_shader_info *info, + case MESA_SHADER_TESS_CTRL: + info->vs.tcs_in_out_eq = key->ge.opt.same_patch_vertices; + info->vs.tcs_temp_only_input_mask = sel->info.tcs_vgpr_only_inputs; +- info->has_epilog = !shader->is_monolithic; + info->tcs.pass_tessfactors_by_reg = sel->info.tessfactors_are_def_in_all_invocs; + info->tcs.patch_stride = si_get_tcs_out_patch_stride(&sel->info); + info->tcs.tcs_offchip_layout = args->tcs_offchip_layout; +-- +2.48.1 + + +From fdd1b46c2683b34bb7171038999e0dc4620a9296 Mon Sep 17 00:00:00 2001 +From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= +Date: Wed, 8 Jan 2025 15:03:32 -0500 +Subject: [PATCH 27/27] radeonsi/gfx12: use ACO if LLVM is 19 or older + +LLVM 19 is missing a SALU hazard fix. + +(cherry-picked + adapted from 239840556f99aaa1b71c450a3b389bd9be1a24f3) + +Reviewed-by: Pierre-Eric Pelloux-Prayer +Part-of: +--- + src/gallium/drivers/radeonsi/si_pipe.c | 9 ++++++++- + 1 file changed, 8 insertions(+), 1 deletion(-) + +diff --git a/src/gallium/drivers/radeonsi/si_pipe.c b/src/gallium/drivers/radeonsi/si_pipe.c +index 4e063c1ce3e..f3f23ab0dd0 100644 +--- a/src/gallium/drivers/radeonsi/si_pipe.c ++++ b/src/gallium/drivers/radeonsi/si_pipe.c +@@ -1198,7 +1198,14 @@ static struct pipe_screen *radeonsi_screen_create_impl(struct radeon_winsys *ws, + sscreen->info.register_shadowing_required = true; + + #if AMD_LLVM_AVAILABLE +- sscreen->use_aco = (sscreen->debug_flags & DBG(USE_ACO)); ++ /* For GFX11.5, LLVM < 19 is missing a workaround that can cause GPU hangs. ACO is the only ++ * alternative that has the workaround and is always available. Same for GFX12. ++ */ ++ if ((sscreen->info.gfx_level == GFX12 && LLVM_VERSION_MAJOR < 20) || ++ (sscreen->info.gfx_level == GFX11_5 && LLVM_VERSION_MAJOR < 19)) ++ sscreen->use_aco = true; ++ else ++ sscreen->use_aco = sscreen->debug_flags & DBG(USE_ACO); + #else + sscreen->use_aco = true; + #endif +-- +2.48.1 + diff --git a/mesa.spec b/mesa.spec index f5c8832..cd80113 100644 --- a/mesa.spec +++ b/mesa.spec @@ -68,7 +68,7 @@ Name: mesa Summary: Mesa graphics libraries %global ver 24.2.8 Version: %{lua:ver = string.gsub(rpm.expand("%{ver}"), "-", "~"); print(ver)} -Release: 1%{?dist} +Release: 2%{?dist} License: MIT AND BSD-3-Clause AND SGI-B-2.0 URL: http://www.mesa3d.org @@ -110,6 +110,11 @@ BuildRequires: wayland-devel Patch10: gnome-shell-glthread-disable.patch +# AMD Navi4x support: +# Backport fixes for radeonsi and disable GFX12 on radv +# https://issues.redhat.com/browse/RHEL-53419 +Patch11: RHEL-53423.patch + # Build our own version but keep the dependency for the RPM macros BuildRequires: meson BuildRequires: gcc @@ -849,6 +854,11 @@ popd %endif %changelog +* Thu Feb 13 2025 José Expósito - 24.2.8-2 +- AMD Navi4x support + Backport fixes for radeonsi and disable GFX12 on radv + Resolves: https://issues.redhat.com/browse/RHEL-53419 + * Thu Nov 28 2024 José Expósito - 24.2.8-1 - Update to 24.2.8 Resolves: https://issues.redhat.com/browse/RHEL-53868