From 0ca93e8090b5f88437b85a9278608018294bd0bb Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= Date: Mon, 25 Mar 2024 03:58:45 -0400 Subject: [PATCH] radeonsi: optimize unaligned compute blits If a blit starts on a coordinate that is not at the beginning of a tile (e.g. 8x8), launch extra threads before 0,0,0 to make all following blocks start at the beginning of such tiles. This makes such blits faster. Acked-by: Pierre-Eric Pelloux-Prayer Part-of: --- .../drivers/radeonsi/si_compute_blit.c | 117 ++++++++++++++++++ src/gallium/drivers/radeonsi/si_pipe.h | 4 +- .../drivers/radeonsi/si_shaderlib_nir.c | 22 +++- 3 files changed, 141 insertions(+), 2 deletions(-) diff --git a/src/gallium/drivers/radeonsi/si_compute_blit.c b/src/gallium/drivers/radeonsi/si_compute_blit.c index 8ea56eea969..c2b3c35ec41 100644 --- a/src/gallium/drivers/radeonsi/si_compute_blit.c +++ b/src/gallium/drivers/radeonsi/si_compute_blit.c @@ -1033,10 +1033,16 @@ static bool si_should_blit_clamp_xy(const struct pipe_blit_info *info) return !in_bounds; } +typedef struct { + unsigned x, y, z; +} uvec3; + bool si_compute_blit(struct si_context *sctx, const struct pipe_blit_info *info, bool testing) { struct si_texture *sdst = (struct si_texture *)info->dst.resource; + struct si_texture *ssrc = (struct si_texture *)info->src.resource; bool is_3d_tiling = sdst->surface.thick_tiling; + unsigned dst_samples = MAX2(1, sdst->buffer.b.b.nr_samples); /* Get the channel sizes. */ unsigned max_dst_chan_size = util_format_get_max_channel_size(info->dst.format); unsigned max_src_chan_size = util_format_get_max_channel_size(info->src.format); @@ -1088,8 +1094,117 @@ bool si_compute_blit(struct si_context *sctx, const struct pipe_blit_info *info, unsigned width = info->dst.box.width; unsigned height = info->dst.box.height; unsigned depth = info->dst.box.depth; + + /* Determine the alignment of coordinates of the first thread of each wave. The alignment should be + * to a 256B block or the size of 1 wave, whichever is less, but there are a few exceptions. + */ + uvec3 align; + if (is_3d_tiling) { + /* Thick tiling. */ + /* This is based on GFX11_SW_PATTERN_NIBBLE01, which also matches GFX10. */ + if (sdst->surface.bpe == 1) + align = (uvec3){8, 4, 8}; + else if (sdst->surface.bpe == 2) + align = (uvec3){4, 4, 8}; + else if (sdst->surface.bpe == 4) + align = (uvec3){4, 4, 4}; + else if (sdst->surface.bpe == 8) + align = (uvec3){4, 2, 4}; + else { + /* 16bpp linear source image reads perform better with this. */ + if (ssrc->surface.is_linear) + align = (uvec3){4, 2, 4}; /* align to 512B for linear->tiled */ + else + align = (uvec3){2, 2, 4}; + } + + /* Clamp the alignment to the size of 1 wave. */ + align.x = MIN2(align.x, 4); + align.y = MIN2(align.y, 4); + align.z = MIN2(align.z, 4); + } else if (sdst->surface.is_linear) { + /* 1D blits from linear to linear are faster unaligned. + * 1D image clears don't benefit from any alignment. + */ + if (height == 1 && depth == 1 && ssrc->surface.is_linear) { + align = (uvec3){1, 1, 1}; + } else { + /* Linear blits should use the cache line size instead of 256B alignment. */ + align.x = MIN2(64, sctx->screen->info.tcc_cache_line_size / sdst->surface.bpe); + align.y = 1; + align.z = 1; + } + } else { + /* Thin tiling. */ + if (sctx->gfx_level >= GFX11) { + /* Samples are next to each other on GFX11+. */ + unsigned pix_size = sdst->surface.bpe * dst_samples; + + /* This is based on GFX11_SW_PATTERN_NIBBLE01. */ + if (pix_size == 1) + align = (uvec3){16, 16, 1}; + else if (pix_size == 2) + align = (uvec3){16, 8, 1}; + else if (pix_size == 4) + align = (uvec3){8, 8, 1}; + else if (pix_size == 8) + align = (uvec3){8, 4, 1}; + else if (pix_size == 16) + align = (uvec3){4, 4, 1}; + else if (pix_size == 32) + align = (uvec3){4, 2, 1}; + else if (pix_size == 64) + align = (uvec3){2, 2, 1}; + else + align = (uvec3){2, 1, 1}; /* 16bpp 8xAA */ + } else { + /* This is for 64KB_R_X. (most likely to occur due to DCC) + * It's based on GFX10_SW_64K_R_X_*xaa_RBPLUS_PATINFO (GFX10.3). + * The patterns are GFX10_SW_PATTERN_NIBBLE01[0, 1, 39, 6, 7] for 8bpp-128bpp. + * GFX6-10.1 and other swizzle modes might be similar. + */ + if (sdst->surface.bpe == 1) + align = (uvec3){16, 16, 1}; + else if (sdst->surface.bpe == 2) + align = (uvec3){16, 8, 1}; + else if (sdst->surface.bpe == 4) + align = (uvec3){8, 8, 1}; + else if (sdst->surface.bpe == 8) + align = (uvec3){8, 4, 1}; + else + align = (uvec3){4, 4, 1}; + } + + /* Clamp the alignment to the size of 1 wave. */ + align.x = MIN2(align.x, 8); + align.y = MIN2(align.y, 8); + } + + /* If we don't have much to copy, don't align. The threshold is guessed and isn't covered + * by benchmarking. + */ + if (width <= align.x * 4) + align.x = 1; + if (height <= align.y * 4) + align.y = 1; + if (depth <= align.z * 4) + align.z = 1; + + unsigned start_x, start_y, start_z; unsigned block_x, block_y, block_z; + /* If the blit destination area is unaligned, launch extra threads before 0,0,0 to make it + * aligned. This makes sure that a wave doesn't straddle a DCC block boundary or a cache line + * unnecessarily, so that each cache line is only stored by exactly 1 CU. The shader will skip + * the extra threads. This makes unaligned compute blits faster. + */ + start_x = info->dst.box.x % align.x; + start_y = info->dst.box.y % align.y; + start_z = info->dst.box.z % align.z; + width += start_x; + height += start_y; + depth += start_z; + /* Choose the block (i.e. wave) dimensions based on the copy area size and the image layout * of dst. */ @@ -1138,6 +1253,7 @@ bool si_compute_blit(struct si_context *sctx, const struct pipe_blit_info *info, options.always_true = true; options.wg_dim = wg_dim; + options.has_start_xyz = start_x || start_y || start_z; options.src_is_1d = info->src.resource->target == PIPE_TEXTURE_1D || info->src.resource->target == PIPE_TEXTURE_1D_ARRAY; options.dst_is_1d = info->dst.resource->target == PIPE_TEXTURE_1D || @@ -1196,6 +1312,7 @@ bool si_compute_blit(struct si_context *sctx, const struct pipe_blit_info *info, sctx->cs_user_data[0] = (info->src.box.x & 0xffff) | ((info->dst.box.x & 0xffff) << 16); sctx->cs_user_data[1] = (info->src.box.y & 0xffff) | ((info->dst.box.y & 0xffff) << 16); sctx->cs_user_data[2] = (info->src.box.z & 0xffff) | ((info->dst.box.z & 0xffff) << 16); + sctx->cs_user_data[3] = (start_x & 0xff) | ((start_y & 0xff) << 8) | ((start_z & 0xff) << 16); si_launch_grid_internal_images(sctx, image, 2, &grid, shader, SI_OP_SYNC_BEFORE_AFTER | diff --git a/src/gallium/drivers/radeonsi/si_pipe.h b/src/gallium/drivers/radeonsi/si_pipe.h index f6686987ebb..1628052edc6 100644 --- a/src/gallium/drivers/radeonsi/si_pipe.h +++ b/src/gallium/drivers/radeonsi/si_pipe.h @@ -1639,8 +1639,10 @@ union si_compute_blit_shader_key { struct { /* The key saved in _mesa_hash_table_create_u32_keys() can't be 0. */ bool always_true:1; - /* Declaration modifiers. */ + /* Workgroup settings. */ uint8_t wg_dim:2; /* 1, 2, or 3 */ + bool has_start_xyz:1; + /* Declaration modifiers. */ bool src_is_1d:1; bool dst_is_1d:1; bool src_is_msaa:1; diff --git a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c index c4e321fffd2..1f96eece782 100644 --- a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c +++ b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c @@ -414,7 +414,7 @@ void *si_create_blit_cs(struct si_context *sctx, const union si_compute_blit_sha BITSET_SET(b.shader->info.msaa_images, 1); /* The workgroup size varies depending on the tiling layout and blit dimensions. */ b.shader->info.workgroup_size_variable = true; - b.shader->info.cs.user_data_components_amd = 3; + b.shader->info.cs.user_data_components_amd = options->has_start_xyz ? 4 : 3; const struct glsl_type *img_type[2] = { glsl_image_type(options->src_is_1d ? GLSL_SAMPLER_DIM_1D : @@ -436,6 +436,23 @@ void *si_create_blit_cs(struct si_context *sctx, const union si_compute_blit_sha /* Instructions. */ /* Let's work with 0-based src and dst coordinates (thread IDs) first. */ nir_def *dst_xyz = nir_pad_vector_imm_int(&b, get_global_ids(&b, options->wg_dim), 0, 3); + + /* If the blit area is unaligned, we launched extra threads to make it aligned. + * Skip those threads here. + */ + nir_if *if_positive = NULL; + if (options->has_start_xyz) { + nir_def *start_xyz = nir_channel(&b, nir_load_user_data_amd(&b), 3); + start_xyz = nir_trim_vector(&b, nir_u2u32(&b, nir_unpack_32_4x8(&b, start_xyz)), 3); + + dst_xyz = nir_isub(&b, dst_xyz, start_xyz); + nir_def *is_positive_xyz = nir_ige_imm(&b, dst_xyz, 0); + nir_def *is_positive = nir_iand(&b, nir_channel(&b, is_positive_xyz, 0), + nir_iand(&b, nir_channel(&b, is_positive_xyz, 1), + nir_channel(&b, is_positive_xyz, 2))); + if_positive = nir_push_if(&b, is_positive); + } + nir_def *src_xyz = dst_xyz; /* Flip src coordinates. */ @@ -531,6 +548,9 @@ void *si_create_blit_cs(struct si_context *sctx, const union si_compute_blit_sha nir_image_deref_store(&b, deref_ssa(&b, img_dst), coord_dst, zero, color, zero); } + if (options->has_start_xyz) + nir_pop_if(&b, if_positive); + return create_shader_state(sctx, b.shader); }