radv: add MSAA support to the comp-to-single fast clear path

Clearing the first sample is enough as long as CMASK is also cleared
to indicate that other samples are also cleared.

I verified that the first sample is always at the beginning of 256B
blocks.

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

src/amd/vulkan/radv_meta_clear.c

@@ -318,8 +318,10 @@ finish_meta_clear_dcc_comp_to_single_state(struct radv_device *device)
 {
    struct radv_meta_state *state = &device->meta_state;
 
-   radv_DestroyPipeline(radv_device_to_handle(device), state->clear_dcc_comp_to_single_pipeline,
-                        &state->alloc);
+   for (uint32_t i = 0; i < 2; i++) {
+      radv_DestroyPipeline(radv_device_to_handle(device),
+                           state->clear_dcc_comp_to_single_pipeline[i], &state->alloc);
+   }
    radv_DestroyPipelineLayout(radv_device_to_handle(device), state->clear_dcc_comp_to_single_p_layout,
                               &state->alloc);
    radv_DestroyDescriptorSetLayout(radv_device_to_handle(device), state->clear_dcc_comp_to_single_ds_layout,
@@ -1147,13 +1149,18 @@ fail:
    return result;
 }
 
+/* Clear DCC using comp-to-single by storing the clear value at the beginning of every 256B block.
+ * For MSAA images, clearing the first sample should be enough as long as CMASK is also cleared.
+ */
 static nir_shader *
-build_clear_dcc_comp_to_single_shader()
+build_clear_dcc_comp_to_single_shader(bool is_msaa)
 {
-   const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, true, GLSL_TYPE_FLOAT);
+   enum glsl_sampler_dim dim = is_msaa ? GLSL_SAMPLER_DIM_MS : GLSL_SAMPLER_DIM_2D;
+   const struct glsl_type *img_type = glsl_image_type(dim, true, GLSL_TYPE_FLOAT);
 
    nir_builder b =
-      nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_clear_dcc_comp_to_single");
+      nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_clear_dcc_comp_to_single-%s",
+                                     is_msaa ? "multisampled" : "singlesampled");
    b.shader->info.workgroup_size[0] = 8;
    b.shader->info.workgroup_size[1] = 8;
    b.shader->info.workgroup_size[2] = 1;
@@ -1190,19 +1197,20 @@ build_clear_dcc_comp_to_single_shader()
                                     nir_channel(&b, clear_values, 1));
 
    /* Store the clear color values. */
+   nir_ssa_def *sample_id = is_msaa ? nir_imm_int(&b, 0) : nir_ssa_undef(&b, 1, 32);
    nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, coord,
-                         nir_imm_int(&b, 0), data, nir_imm_int(&b, 0),
-                         .image_dim = GLSL_SAMPLER_DIM_2D, .image_array = true);
+                         sample_id, data, nir_imm_int(&b, 0),
+                         .image_dim = dim, .image_array = true);
 
    return b.shader;
 }
 
 static VkResult
-create_dcc_comp_to_single_pipeline(struct radv_device *device, VkPipeline *pipeline)
+create_dcc_comp_to_single_pipeline(struct radv_device *device, bool is_msaa, VkPipeline *pipeline)
 {
    struct radv_meta_state *state = &device->meta_state;
    VkResult result;
-   nir_shader *cs = build_clear_dcc_comp_to_single_shader();
+   nir_shader *cs = build_clear_dcc_comp_to_single_shader(is_msaa);
 
    VkPipelineShaderStageCreateInfo shader_stage = {
       .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
@@ -1268,9 +1276,12 @@ init_meta_clear_dcc_comp_to_single_state(struct radv_device *device)
    if (result != VK_SUCCESS)
       goto fail;
 
-   result = create_dcc_comp_to_single_pipeline(device, &state->clear_dcc_comp_to_single_pipeline);
-   if (result != VK_SUCCESS)
-      goto fail;
+   for (uint32_t i = 0; i < 2; i++) {
+      result = create_dcc_comp_to_single_pipeline(device, !!i,
+                                                  &state->clear_dcc_comp_to_single_pipeline[i]);
+      if (result != VK_SUCCESS)
+         goto fail;
+   }
 
 fail:
    return result;
@@ -1513,6 +1524,7 @@ radv_clear_dcc_comp_to_single(struct radv_cmd_buffer *cmd_buffer,
    unsigned bytes_per_pixel = vk_format_get_blocksize(image->vk_format);
    unsigned layer_count = radv_get_layerCount(image, range);
    struct radv_meta_saved_state saved_state;
+   bool is_msaa = image->info.samples > 1;
    struct radv_image_view iview;
    VkFormat format;
 
@@ -1540,7 +1552,7 @@ radv_clear_dcc_comp_to_single(struct radv_cmd_buffer *cmd_buffer,
       &saved_state, cmd_buffer,
       RADV_META_SAVE_DESCRIPTORS | RADV_META_SAVE_COMPUTE_PIPELINE | RADV_META_SAVE_CONSTANTS);
 
-   VkPipeline pipeline = device->meta_state.clear_dcc_comp_to_single_pipeline;
+   VkPipeline pipeline = device->meta_state.clear_dcc_comp_to_single_pipeline[is_msaa];
 
    radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE,
                         pipeline);

src/amd/vulkan/radv_private.h

@@ -475,7 +475,7 @@ struct radv_meta_state {
    VkDescriptorSetLayout copy_vrs_htile_ds_layout;
 
    /* Clear DCC with comp-to-single. */
-   VkPipeline clear_dcc_comp_to_single_pipeline;
+   VkPipeline clear_dcc_comp_to_single_pipeline[2]; /* 0: 1x, 1: 2x/4x/8x */
    VkPipelineLayout clear_dcc_comp_to_single_p_layout;
    VkDescriptorSetLayout clear_dcc_comp_to_single_ds_layout;