fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-05 05:18:08 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 2

radeonsi: convert copy_image shader to NIR

Browse source 
Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15268>
This commit is contained in:
Mihai Preda 2022-03-07 19:11:16 +02:00 • committed by Marge Bot
parent 79ba1962ac
commit 18722af9d2
4 changed files with 48 additions and 42 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
src/gallium/drivers/radeonsi/si_compute_blit.c
View file

@ -633,7 +633,7 @@ void si_compute_copy_image(struct si_context *sctx, struct pipe_resource *dst, u
si_launch_grid_internal(sctx, &info, sctx->cs_copy_image_1d_array, flags | SI_OP_CS_IMAGE);
} else {
if (!sctx->cs_copy_image)
sctx->cs_copy_image = si_create_copy_image_compute_shader(ctx);
sctx->cs_copy_image = si_create_copy_image_cs(ctx);
/* This is better for access over PCIe. */
if (is_linear) {

2
src/gallium/drivers/radeonsi/si_pipe.h
View file

@ -1517,6 +1517,7 @@ void si_suspend_queries(struct si_context *sctx);
void si_resume_queries(struct si_context *sctx);
/* si_shaderlib_nir.c */
void *si_create_copy_image_cs(struct pipe_context *ctx);
void *si_create_dcc_retile_cs(struct si_context *sctx, struct radeon_surf *surf);
void *gfx9_create_clear_dcc_msaa_cs(struct si_context *sctx, struct si_texture *tex);
@ -1527,7 +1528,6 @@ void *si_create_fixed_func_tcs(struct si_context *sctx);
void *si_create_dma_compute_shader(struct pipe_context *ctx, unsigned num_dwords_per_thread,
bool dst_stream_cache_policy, bool is_copy);
void *si_create_clear_buffer_rmw_cs(struct si_context *sctx);
void *si_create_copy_image_compute_shader(struct pipe_context *ctx);
void *si_create_copy_image_compute_shader_1d_array(struct pipe_context *ctx);
void *si_create_dcc_decompress_cs(struct pipe_context *ctx);
void *si_clear_render_target_shader(struct pipe_context *ctx);

46
src/gallium/drivers/radeonsi/si_shaderlib_nir.c
View file

@ -53,6 +53,52 @@ static void unpack_2x16(nir_builder *b, nir_ssa_def *src, nir_ssa_def **x, nir_s
*y = nir_ushr(b, src, nir_imm_int(b, 16));
}
/* Create a NIR compute shader implementing copy_image.
*
* This is the NIR version of the removed si_create_copy_image_compute_shader() [TGSI].
* It inherits the following note from the TGSI version:
* "Luckily, this works with all texture targets except 1D_ARRAY."
*/
void *si_create_copy_image_cs(struct pipe_context *ctx)
{
struct si_context *sctx = (struct si_context *) ctx;
const nir_shader_compiler_options *options =
sctx->b.screen->get_compiler_options(sctx->b.screen, PIPE_SHADER_IR_NIR, PIPE_SHADER_COMPUTE);
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, "copy_image_cs");
b.shader->info.workgroup_size_variable = true;
b.shader->info.cs.user_data_components_amd = 3;
b.shader->info.num_images = 2;
nir_ssa_def *coord_src = NULL, *coord_dst = NULL;
unpack_2x16(&b, nir_load_user_data_amd(&b), &coord_src, &coord_dst);
nir_ssa_def *ids = get_global_ids(&b, 3);
coord_src = nir_iadd(&b, coord_src, ids);
coord_dst = nir_iadd(&b, coord_dst, ids);
const struct glsl_type *img_type =
glsl_image_type(GLSL_SAMPLER_DIM_2D, /*is_array*/ true, GLSL_TYPE_FLOAT);
nir_variable *img_src = nir_variable_create(b.shader, nir_var_image, img_type, "img_src");
img_src->data.binding = 0;
nir_variable *img_dst = nir_variable_create(b.shader, nir_var_image, img_type, "img_dst");
img_dst->data.binding = 1;
nir_ssa_def *data = nir_image_deref_load(&b, /*num_components*/ 4, /*bit_size*/ 32,
&nir_build_deref_var(&b, img_src)->dest.ssa, coord_src, nir_ssa_undef(&b, 1, 32),
nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D);
nir_image_deref_store(&b,
&nir_build_deref_var(&b, img_dst)->dest.ssa, coord_dst, nir_ssa_undef(&b, 1, 32), data,
nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D);
return create_nir_cs(sctx, &b);
}
void *si_create_dcc_retile_cs(struct si_context *sctx, struct radeon_surf *surf)
{
const nir_shader_compiler_options *options =

40
src/gallium/drivers/radeonsi/si_shaderlib_tgsi.c
View file

@ -429,46 +429,6 @@ void *si_create_query_result_cs(struct si_context *sctx)
return sctx->b.create_compute_state(&sctx->b, &state);
}
/* Create a compute shader implementing copy_image.
* Luckily, this works with all texture targets except 1D_ARRAY.
*/
void *si_create_copy_image_compute_shader(struct pipe_context *ctx)
{
static const char text[] =
"COMP\n"
"PROPERTY CS_USER_DATA_COMPONENTS_AMD 3\n"
"DCL SV[0], THREAD_ID\n"
"DCL SV[1], BLOCK_ID\n"
"DCL SV[2], BLOCK_SIZE\n"
"DCL SV[3], CS_USER_DATA_AMD\n"
"DCL IMAGE[0], 2D_ARRAY, PIPE_FORMAT_R32G32B32A32_FLOAT, WR\n"
"DCL IMAGE[1], 2D_ARRAY, PIPE_FORMAT_R32G32B32A32_FLOAT, WR\n"
"DCL TEMP[0..3], LOCAL\n"
"IMM[0] UINT32 {65535, 16, 0, 0}\n"
"UMAD TEMP[0].xyz, SV[1], SV[2], SV[0]\n" /* threadID.xyz */
"AND TEMP[1].xyz, SV[3], IMM[0].xxxx\n" /* src.xyz */
"UADD TEMP[1].xyz, TEMP[1], TEMP[0]\n" /* src.xyz + threadID.xyz */
"LOAD TEMP[3], IMAGE[0], TEMP[1], 2D_ARRAY, PIPE_FORMAT_R32G32B32A32_FLOAT\n"
"USHR TEMP[2].xyz, SV[3], IMM[0].yyyy\n" /* dst.xyz */
"UADD TEMP[2].xyz, TEMP[2], TEMP[0]\n" /* dst.xyz + threadID.xyz */
"STORE IMAGE[1], TEMP[2], TEMP[3], 2D_ARRAY, PIPE_FORMAT_R32G32B32A32_FLOAT\n"
"END\n";
struct tgsi_token tokens[1024];
struct pipe_compute_state state = {0};
if (!tgsi_text_translate(text, tokens, ARRAY_SIZE(tokens))) {
assert(false);
return NULL;
}
state.ir_type = PIPE_SHADER_IR_TGSI;
state.prog = tokens;
return ctx->create_compute_state(ctx, &state);
}
void *si_create_copy_image_compute_shader_1d_array(struct pipe_context *ctx)
{
static const char text[] =