Mesa (main): radv: use nir_image_deref_{load,store} in the DCC retile compute path

[GitLab Mirror]

Module: Mesa
Branch: main
Commit: ef2e7f765249512658abeda3c36b34d49bb59cca
URL:    http://cgit.freedesktop.org/mesa/mesa/commit/?id=ef2e7f765249512658abeda3c36b34d49bb59cca

Author: Samuel Pitoiset <samuel.pitoiset at gmail.com>
Date:   Thu Aug 26 09:48:00 2021 +0200

radv: use nir_image_deref_{load,store} in the DCC retile compute path

Signed-off-by: Samuel Pitoiset <samuel.pitoiset at gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas at basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12561>

---

 src/amd/vulkan/radv_meta_dcc_retile.c | 33 ++++++++++-----------------------
 1 file changed, 10 insertions(+), 23 deletions(-)

diff --git a/src/amd/vulkan/radv_meta_dcc_retile.c b/src/amd/vulkan/radv_meta_dcc_retile.c
index 0dc41da9bef..943b41ce63f 100644
--- a/src/amd/vulkan/radv_meta_dcc_retile.c
+++ b/src/amd/vulkan/radv_meta_dcc_retile.c
@@ -46,7 +46,8 @@ get_global_ids(nir_builder *b, unsigned num_components)
 static nir_shader *
 build_dcc_retile_compute_shader(struct radv_device *dev, struct radeon_surf *surf)
 {
-   const struct glsl_type *buf_type = glsl_image_type(GLSL_SAMPLER_DIM_BUF, false, GLSL_TYPE_UINT);
+   enum glsl_sampler_dim dim = GLSL_SAMPLER_DIM_BUF;
+   const struct glsl_type *buf_type = glsl_image_type(dim, false, GLSL_TYPE_UINT);
    nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "dcc_retile_compute");
 
    b.shader->info.workgroup_size[0] = 8;
@@ -85,28 +86,14 @@ build_dcc_retile_compute_shader(struct radv_device *dev, struct radeon_surf *sur
       dst_dcc_pitch, dst_dcc_height, zero, nir_channel(&b, coord, 0), nir_channel(&b, coord, 1),
       zero, zero, zero);
 
-   nir_intrinsic_instr *dcc_val =
-      nir_intrinsic_instr_create(b.shader, nir_intrinsic_image_deref_load);
-   dcc_val->num_components = 1;
-   dcc_val->src[0] = nir_src_for_ssa(input_dcc_ref);
-   dcc_val->src[1] = nir_src_for_ssa(nir_vec4(&b, src, src, src, src));
-   dcc_val->src[2] = nir_src_for_ssa(nir_ssa_undef(&b, 1, 32));
-   dcc_val->src[3] = nir_src_for_ssa(nir_imm_int(&b, 0));
-   nir_ssa_dest_init(&dcc_val->instr, &dcc_val->dest, 1, 32, "dcc_val");
-   nir_intrinsic_set_image_dim(dcc_val, GLSL_SAMPLER_DIM_BUF);
-   nir_builder_instr_insert(&b, &dcc_val->instr);
-
-   nir_intrinsic_instr *store =
-      nir_intrinsic_instr_create(b.shader, nir_intrinsic_image_deref_store);
-   store->num_components = 1;
-   store->src[0] = nir_src_for_ssa(output_dcc_ref);
-   store->src[1] = nir_src_for_ssa(nir_vec4(&b, dst, dst, dst, dst));
-   store->src[2] = nir_src_for_ssa(nir_ssa_undef(&b, 1, 32));
-   store->src[3] = nir_src_for_ssa(&dcc_val->dest.ssa);
-   store->src[4] = nir_src_for_ssa(nir_imm_int(&b, 0));
-   nir_intrinsic_set_image_dim(store, GLSL_SAMPLER_DIM_BUF);
-
-   nir_builder_instr_insert(&b, &store->instr);
+   nir_ssa_def *dcc_val = nir_image_deref_load(&b, 1, 32, input_dcc_ref,
+                                               nir_vec4(&b, src, src, src, src),
+                                               nir_ssa_undef(&b, 1, 32), nir_imm_int(&b, 0),
+                                               .image_dim = dim);
+
+   nir_image_deref_store(&b, output_dcc_ref, nir_vec4(&b, dst, dst, dst, dst),
+                         nir_ssa_undef(&b, 1, 32), dcc_val, nir_imm_int(&b, 0), .image_dim = dim);
+
    return b.shader;
 }