radeonsi: "clear_render_target" shader in nir

Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/25221>

diff --git a/src/gallium/drivers/radeonsi/si_compute_blit.c b/src/gallium/drivers/radeonsi/si_compute_blit.c
index 5c5885e..db3978b 100644 (file)
--- 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(ctx);
+         sctx->cs_clear_render_target = si_clear_render_target_shader(sctx);
       shader = sctx->cs_clear_render_target;
 
       info.block[0] = 8;

diff --git a/src/gallium/drivers/radeonsi/si_pipe.h b/src/gallium/drivers/radeonsi/si_pipe.h
index e5b90ca..ba6aec4 100644 (file)
--- a/src/gallium/drivers/radeonsi/si_pipe.h
+++ b/src/gallium/drivers/radeonsi/si_pipe.h
@@ -1641,7 +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 pipe_context *ctx);
+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_12bytes_buffer_shader(struct pipe_context *ctx);
 void *si_create_fmask_expand_cs(struct pipe_context *ctx, unsigned num_samples, bool is_array);

diff --git a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
index b78ed0d..cdbe9e9 100644 (file)
--- a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
+++ b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
@@ -560,3 +560,36 @@ 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)
+{
+   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);
+   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);
+
+   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 94295a0..4775113 100644 (file)
--- a/src/gallium/drivers/radeonsi/si_shaderlib_tgsi.c
+++ b/src/gallium/drivers/radeonsi/si_shaderlib_tgsi.c
@@ -388,40 +388,6 @@ void *si_create_query_result_cs(struct si_context *sctx)
    return sctx->b.create_compute_state(&sctx->b, &state);
 }
 
-void *si_clear_render_target_shader(struct pipe_context *ctx)
-{
-   static const char text[] =
-      "COMP\n"
-      "PROPERTY CS_FIXED_BLOCK_WIDTH 8\n"
-      "PROPERTY CS_FIXED_BLOCK_HEIGHT 8\n"
-      "PROPERTY CS_FIXED_BLOCK_DEPTH 1\n"
-      "DCL SV[0], THREAD_ID\n"
-      "DCL SV[1], BLOCK_ID\n"
-      "DCL IMAGE[0], 2D_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 {8, 1, 0, 0}\n"
-      "MOV TEMP[0].xyz, CONST[0][0].xyzw\n"
-      "UMAD TEMP[1].xyz, SV[1].xyzz, IMM[0].xxyy, SV[0].xyzz\n"
-      "UADD TEMP[2].xyz, TEMP[1].xyzx, TEMP[0].xyzx\n"
-      "MOV TEMP[3].xyzw, CONST[0][1].xyzw\n"
-      "STORE IMAGE[0], TEMP[2].xyzz, 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);
-}
-
 /* TODO: Didn't really test 1D_ARRAY */
 void *si_clear_render_target_shader_1d_array(struct pipe_context *ctx)
 {