Module: Mesa Branch: main Commit: 3f44a8321f38890fefc1c553ad80810b61611e18 URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=3f44a8321f38890fefc1c553ad80810b61611e18
Author: Ganesh Belgur Ramachandra <[email protected]> Date: Thu Sep 14 07:33:35 2023 -0500 radeonsi: "clear_render_target_1d_array" shader in nir Reviewed-by: Marek Olšák <[email protected]> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/25221> --- src/gallium/drivers/radeonsi/si_compute_blit.c | 4 +-- src/gallium/drivers/radeonsi/si_pipe.h | 3 +- src/gallium/drivers/radeonsi/si_shaderlib_nir.c | 32 +++++++++++++++++----- src/gallium/drivers/radeonsi/si_shaderlib_tgsi.c | 35 ------------------------ 4 files changed, 28 insertions(+), 46 deletions(-) diff --git a/src/gallium/drivers/radeonsi/si_compute_blit.c b/src/gallium/drivers/radeonsi/si_compute_blit.c index db3978bfb4f..a25c0a1d19c 100644 --- a/src/gallium/drivers/radeonsi/si_compute_blit.c +++ b/src/gallium/drivers/radeonsi/si_compute_blit.c @@ -1000,7 +1000,7 @@ void si_compute_clear_render_target(struct pipe_context *ctx, struct pipe_surfac if (dstsurf->texture->target != PIPE_TEXTURE_1D_ARRAY) { if (!sctx->cs_clear_render_target) - sctx->cs_clear_render_target = si_clear_render_target_shader(sctx); + sctx->cs_clear_render_target = si_clear_render_target_shader(sctx, PIPE_TEXTURE_2D_ARRAY); shader = sctx->cs_clear_render_target; info.block[0] = 8; @@ -1013,7 +1013,7 @@ void si_compute_clear_render_target(struct pipe_context *ctx, struct pipe_surfac info.grid[2] = num_layers; } else { if (!sctx->cs_clear_render_target_1d_array) - sctx->cs_clear_render_target_1d_array = si_clear_render_target_shader_1d_array(ctx); + sctx->cs_clear_render_target_1d_array = si_clear_render_target_shader(sctx, PIPE_TEXTURE_1D_ARRAY); shader = sctx->cs_clear_render_target_1d_array; info.block[0] = 64; diff --git a/src/gallium/drivers/radeonsi/si_pipe.h b/src/gallium/drivers/radeonsi/si_pipe.h index ba6aec40d95..4c94aa67982 100644 --- a/src/gallium/drivers/radeonsi/si_pipe.h +++ b/src/gallium/drivers/radeonsi/si_pipe.h @@ -1641,8 +1641,7 @@ void *si_get_blitter_vs(struct si_context *sctx, enum blitter_attrib_type type, 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_clear_render_target_shader(struct si_context *sctx); -void *si_clear_render_target_shader_1d_array(struct pipe_context *ctx); +void *si_clear_render_target_shader(struct si_context *sctx, enum pipe_texture_target type); void *si_clear_12bytes_buffer_shader(struct pipe_context *ctx); void *si_create_fmask_expand_cs(struct pipe_context *ctx, unsigned num_samples, bool is_array); void *si_create_query_result_cs(struct si_context *sctx); diff --git a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c index cdbe9e9a5b5..555d98b61c0 100644 --- a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c +++ b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c @@ -561,35 +561,53 @@ void *si_create_blit_cs(struct si_context *sctx, const union si_compute_blit_sha return create_shader_state(sctx, b.shader); } -void *si_clear_render_target_shader(struct si_context *sctx) +void *si_clear_render_target_shader(struct si_context *sctx, enum pipe_texture_target type) { + nir_def *address; + enum glsl_sampler_dim sampler_type; + 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, "clear_render_target"); - b.shader->info.workgroup_size[0] = 8; - b.shader->info.workgroup_size[1] = 8; - b.shader->info.workgroup_size[2] = 1; b.shader->info.num_ubos = 1; b.shader->info.num_images = 1; b.shader->num_uniforms = 2; - const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, true, GLSL_TYPE_FLOAT); + switch (type) { + case PIPE_TEXTURE_1D_ARRAY: + b.shader->info.workgroup_size[0] = 64; + b.shader->info.workgroup_size[1] = 1; + b.shader->info.workgroup_size[2] = 1; + sampler_type = GLSL_SAMPLER_DIM_1D; + address = get_global_ids(&b, 2); + break; + case PIPE_TEXTURE_2D_ARRAY: + b.shader->info.workgroup_size[0] = 8; + b.shader->info.workgroup_size[1] = 8; + b.shader->info.workgroup_size[2] = 1; + sampler_type = GLSL_SAMPLER_DIM_2D; + address = get_global_ids(&b, 3); + break; + default: + unreachable("unsupported texture target type"); + } + + const struct glsl_type *img_type = glsl_image_type(sampler_type, true, GLSL_TYPE_FLOAT); nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "image"); output_img->data.image.format = PIPE_FORMAT_R32G32B32A32_FLOAT; nir_def *zero = nir_imm_int(&b, 0); nir_def *ubo = nir_load_ubo(&b, 4, 32, zero, zero, .range_base = 0, .range = 16); - nir_def *address = get_global_ids(&b, 3); address = nir_iadd(&b, address, ubo); nir_def *coord = nir_pad_vector(&b, address, 4); nir_def *data = nir_load_ubo(&b, 4, 32, zero, nir_imm_int(&b, 16), .range_base = 16, .range = 16); nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->def, coord, zero, data, zero, - .image_dim = GLSL_SAMPLER_DIM_2D, .image_array = true); + .image_dim = sampler_type, .image_array = true); return create_shader_state(sctx, b.shader); } diff --git a/src/gallium/drivers/radeonsi/si_shaderlib_tgsi.c b/src/gallium/drivers/radeonsi/si_shaderlib_tgsi.c index 4775113d465..1468bdc508c 100644 --- a/src/gallium/drivers/radeonsi/si_shaderlib_tgsi.c +++ b/src/gallium/drivers/radeonsi/si_shaderlib_tgsi.c @@ -388,41 +388,6 @@ void *si_create_query_result_cs(struct si_context *sctx) return sctx->b.create_compute_state(&sctx->b, &state); } -/* TODO: Didn't really test 1D_ARRAY */ -void *si_clear_render_target_shader_1d_array(struct pipe_context *ctx) -{ - static const char text[] = - "COMP\n" - "PROPERTY CS_FIXED_BLOCK_WIDTH 64\n" - "PROPERTY CS_FIXED_BLOCK_HEIGHT 1\n" - "PROPERTY CS_FIXED_BLOCK_DEPTH 1\n" - "DCL SV[0], THREAD_ID\n" - "DCL SV[1], BLOCK_ID\n" - "DCL IMAGE[0], 1D_ARRAY, PIPE_FORMAT_R32G32B32A32_FLOAT, WR\n" - "DCL CONST[0][0..1]\n" // 0:xyzw 1:xyzw - "DCL TEMP[0..3], LOCAL\n" - "IMM[0] UINT32 {64, 1, 0, 0}\n" - "MOV TEMP[0].xy, CONST[0][0].xzzw\n" - "UMAD TEMP[1].xy, SV[1].xyzz, IMM[0].xyyy, SV[0].xyzz\n" - "UADD TEMP[2].xy, TEMP[1].xyzx, TEMP[0].xyzx\n" - "MOV TEMP[3].xyzw, CONST[0][1].xyzw\n" - "STORE IMAGE[0], TEMP[2].xyzz, TEMP[3], 1D_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_clear_12bytes_buffer_shader(struct pipe_context *ctx) { static const char text[] = "COMP\n"
