diff --git a/src/amd/vulkan/meta/radv_meta.h b/src/amd/vulkan/meta/radv_meta.h index d0305ba735d..46531ac7626 100644 --- a/src/amd/vulkan/meta/radv_meta.h +++ b/src/amd/vulkan/meta/radv_meta.h @@ -328,6 +328,16 @@ nir_shader *radv_meta_nir_build_fmask_copy_compute_shader(struct radv_device *de nir_shader *radv_meta_nir_build_fmask_expand_compute_shader(struct radv_device *device, int samples); +enum radv_meta_resolve_type { + RADV_META_DEPTH_RESOLVE, + RADV_META_STENCIL_RESOLVE, +}; +nir_shader *radv_meta_nir_build_resolve_compute_shader(struct radv_device *dev, bool is_integer, bool is_srgb, + int samples); +nir_shader *radv_meta_nir_build_depth_stencil_resolve_compute_shader(struct radv_device *dev, int samples, + enum radv_meta_resolve_type index, + VkResolveModeFlagBits resolve_mode); + uint32_t radv_fill_buffer(struct radv_cmd_buffer *cmd_buffer, const struct radv_image *image, struct radeon_winsys_bo *bo, uint64_t va, uint64_t size, uint32_t value); diff --git a/src/amd/vulkan/meta/radv_meta_resolve_cs.c b/src/amd/vulkan/meta/radv_meta_resolve_cs.c index aa9f86480c6..5c0d997fd88 100644 --- a/src/amd/vulkan/meta/radv_meta_resolve_cs.c +++ b/src/amd/vulkan/meta/radv_meta_resolve_cs.c @@ -7,164 +7,13 @@ #include #include -#include "nir/nir_builder.h" -#include "nir/nir_format_convert.h" - #include "radv_entrypoints.h" #include "radv_formats.h" #include "radv_meta.h" -#include "sid.h" #include "vk_common_entrypoints.h" #include "vk_format.h" #include "vk_shader_module.h" -static nir_def * -radv_meta_build_resolve_srgb_conversion(nir_builder *b, nir_def *input) -{ - unsigned i; - nir_def *comp[4]; - for (i = 0; i < 3; i++) - comp[i] = nir_format_linear_to_srgb(b, nir_channel(b, input, i)); - comp[3] = nir_channels(b, input, 1 << 3); - return nir_vec(b, comp, 4); -} - -static nir_shader * -build_resolve_compute_shader(struct radv_device *dev, bool is_integer, bool is_srgb, int samples) -{ - enum glsl_base_type img_base_type = is_integer ? GLSL_TYPE_UINT : GLSL_TYPE_FLOAT; - const struct glsl_type *sampler_type = glsl_sampler_type(GLSL_SAMPLER_DIM_MS, false, false, img_base_type); - const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, false, img_base_type); - nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_resolve_cs-%d-%s", samples, - is_integer ? "int" : (is_srgb ? "srgb" : "float")); - b.shader->info.workgroup_size[0] = 8; - b.shader->info.workgroup_size[1] = 8; - - nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, sampler_type, "s_tex"); - input_img->data.descriptor_set = 0; - input_img->data.binding = 0; - - nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img"); - output_img->data.descriptor_set = 0; - output_img->data.binding = 1; - - nir_def *global_id = get_global_ids(&b, 2); - - nir_def *src_offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8); - nir_def *dst_offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 8), .range = 16); - - nir_def *src_coord = nir_iadd(&b, global_id, src_offset); - nir_def *dst_coord = nir_iadd(&b, global_id, dst_offset); - - nir_variable *color = nir_local_variable_create(b.impl, glsl_vec4_type(), "color"); - - radv_meta_build_resolve_shader_core(dev, &b, is_integer, samples, input_img, color, src_coord); - - nir_def *outval = nir_load_var(&b, color); - if (is_srgb) - outval = radv_meta_build_resolve_srgb_conversion(&b, outval); - - nir_def *img_coord = nir_vec4(&b, nir_channel(&b, dst_coord, 0), nir_channel(&b, dst_coord, 1), nir_undef(&b, 1, 32), - nir_undef(&b, 1, 32)); - - nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->def, img_coord, nir_undef(&b, 1, 32), outval, - nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D); - return b.shader; -} - -enum { - DEPTH_RESOLVE, - STENCIL_RESOLVE, -}; - -static const char * -get_resolve_mode_str(VkResolveModeFlagBits resolve_mode) -{ - switch (resolve_mode) { - case VK_RESOLVE_MODE_SAMPLE_ZERO_BIT: - return "zero"; - case VK_RESOLVE_MODE_AVERAGE_BIT: - return "average"; - case VK_RESOLVE_MODE_MIN_BIT: - return "min"; - case VK_RESOLVE_MODE_MAX_BIT: - return "max"; - default: - unreachable("invalid resolve mode"); - } -} - -static nir_shader * -build_depth_stencil_resolve_compute_shader(struct radv_device *dev, int samples, int index, - VkResolveModeFlagBits resolve_mode) -{ - enum glsl_base_type img_base_type = index == DEPTH_RESOLVE ? GLSL_TYPE_FLOAT : GLSL_TYPE_UINT; - const struct glsl_type *sampler_type = glsl_sampler_type(GLSL_SAMPLER_DIM_MS, false, true, img_base_type); - const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, true, img_base_type); - - nir_builder b = - radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_resolve_cs_%s-%s-%d", - index == DEPTH_RESOLVE ? "depth" : "stencil", get_resolve_mode_str(resolve_mode), samples); - b.shader->info.workgroup_size[0] = 8; - b.shader->info.workgroup_size[1] = 8; - - nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, sampler_type, "s_tex"); - input_img->data.descriptor_set = 0; - input_img->data.binding = 0; - - nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img"); - output_img->data.descriptor_set = 0; - output_img->data.binding = 1; - - nir_def *global_id = get_global_ids(&b, 3); - - nir_def *offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8); - - nir_def *resolve_coord = nir_iadd(&b, nir_trim_vector(&b, global_id, 2), offset); - - nir_def *img_coord = - nir_vec3(&b, nir_channel(&b, resolve_coord, 0), nir_channel(&b, resolve_coord, 1), nir_channel(&b, global_id, 2)); - - nir_deref_instr *input_img_deref = nir_build_deref_var(&b, input_img); - nir_def *outval = nir_txf_ms_deref(&b, input_img_deref, img_coord, nir_imm_int(&b, 0)); - - if (resolve_mode != VK_RESOLVE_MODE_SAMPLE_ZERO_BIT) { - for (int i = 1; i < samples; i++) { - nir_def *si = nir_txf_ms_deref(&b, input_img_deref, img_coord, nir_imm_int(&b, i)); - - switch (resolve_mode) { - case VK_RESOLVE_MODE_AVERAGE_BIT: - assert(index == DEPTH_RESOLVE); - outval = nir_fadd(&b, outval, si); - break; - case VK_RESOLVE_MODE_MIN_BIT: - if (index == DEPTH_RESOLVE) - outval = nir_fmin(&b, outval, si); - else - outval = nir_umin(&b, outval, si); - break; - case VK_RESOLVE_MODE_MAX_BIT: - if (index == DEPTH_RESOLVE) - outval = nir_fmax(&b, outval, si); - else - outval = nir_umax(&b, outval, si); - break; - default: - unreachable("invalid resolve mode"); - } - } - - if (resolve_mode == VK_RESOLVE_MODE_AVERAGE_BIT) - outval = nir_fdiv_imm(&b, outval, samples); - } - - nir_def *coord = nir_vec4(&b, nir_channel(&b, img_coord, 0), nir_channel(&b, img_coord, 1), - nir_channel(&b, img_coord, 2), nir_undef(&b, 1, 32)); - nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->def, coord, nir_undef(&b, 1, 32), outval, - nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D, .image_array = true); - return b.shader; -} - static VkResult create_layout(struct radv_device *device, VkPipelineLayout *layout_out) { @@ -234,7 +83,7 @@ get_color_resolve_pipeline(struct radv_device *device, struct radv_image_view *s return VK_SUCCESS; } - nir_shader *cs = build_resolve_compute_shader(device, is_integer, is_srgb, samples); + nir_shader *cs = radv_meta_nir_build_resolve_compute_shader(device, is_integer, is_srgb, samples); const VkPipelineShaderStageCreateInfo stage_info = { .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO, @@ -324,7 +173,8 @@ get_depth_stencil_resolve_pipeline(struct radv_device *device, int samples, VkIm VkPipelineLayout *layout_out) { - const int index = aspects == VK_IMAGE_ASPECT_DEPTH_BIT ? DEPTH_RESOLVE : STENCIL_RESOLVE; + const enum radv_meta_resolve_type index = + aspects == VK_IMAGE_ASPECT_DEPTH_BIT ? RADV_META_DEPTH_RESOLVE : RADV_META_STENCIL_RESOLVE; struct radv_resolve_ds_cs_key key; VkResult result; @@ -344,7 +194,7 @@ get_depth_stencil_resolve_pipeline(struct radv_device *device, int samples, VkIm return VK_SUCCESS; } - nir_shader *cs = build_depth_stencil_resolve_compute_shader(device, samples, index, resolve_mode); + nir_shader *cs = radv_meta_nir_build_depth_stencil_resolve_compute_shader(device, samples, index, resolve_mode); const VkPipelineShaderStageCreateInfo stage_info = { .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO, diff --git a/src/amd/vulkan/nir/radv_meta_nir.c b/src/amd/vulkan/nir/radv_meta_nir.c index 52d30682dc6..f36d6b87502 100644 --- a/src/amd/vulkan/nir/radv_meta_nir.c +++ b/src/amd/vulkan/nir/radv_meta_nir.c @@ -11,6 +11,7 @@ #include "ac_surface.h" #include "../meta/radv_meta.h" +#include "nir/nir_format_convert.h" #include "nir_builder.h" nir_shader * @@ -1146,3 +1147,146 @@ radv_meta_nir_build_fmask_expand_compute_shader(struct radv_device *device, int return b.shader; } + +static nir_def * +radv_meta_build_resolve_srgb_conversion(nir_builder *b, nir_def *input) +{ + unsigned i; + nir_def *comp[4]; + for (i = 0; i < 3; i++) + comp[i] = nir_format_linear_to_srgb(b, nir_channel(b, input, i)); + comp[3] = nir_channels(b, input, 1 << 3); + return nir_vec(b, comp, 4); +} + +nir_shader * +radv_meta_nir_build_resolve_compute_shader(struct radv_device *dev, bool is_integer, bool is_srgb, int samples) +{ + enum glsl_base_type img_base_type = is_integer ? GLSL_TYPE_UINT : GLSL_TYPE_FLOAT; + const struct glsl_type *sampler_type = glsl_sampler_type(GLSL_SAMPLER_DIM_MS, false, false, img_base_type); + const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, false, img_base_type); + nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_resolve_cs-%d-%s", samples, + is_integer ? "int" : (is_srgb ? "srgb" : "float")); + b.shader->info.workgroup_size[0] = 8; + b.shader->info.workgroup_size[1] = 8; + + nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, sampler_type, "s_tex"); + input_img->data.descriptor_set = 0; + input_img->data.binding = 0; + + nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img"); + output_img->data.descriptor_set = 0; + output_img->data.binding = 1; + + nir_def *global_id = get_global_ids(&b, 2); + + nir_def *src_offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8); + nir_def *dst_offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 8), .range = 16); + + nir_def *src_coord = nir_iadd(&b, global_id, src_offset); + nir_def *dst_coord = nir_iadd(&b, global_id, dst_offset); + + nir_variable *color = nir_local_variable_create(b.impl, glsl_vec4_type(), "color"); + + radv_meta_build_resolve_shader_core(dev, &b, is_integer, samples, input_img, color, src_coord); + + nir_def *outval = nir_load_var(&b, color); + if (is_srgb) + outval = radv_meta_build_resolve_srgb_conversion(&b, outval); + + nir_def *img_coord = nir_vec4(&b, nir_channel(&b, dst_coord, 0), nir_channel(&b, dst_coord, 1), nir_undef(&b, 1, 32), + nir_undef(&b, 1, 32)); + + nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->def, img_coord, nir_undef(&b, 1, 32), outval, + nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D); + return b.shader; +} + +static const char * +get_resolve_mode_str(VkResolveModeFlagBits resolve_mode) +{ + switch (resolve_mode) { + case VK_RESOLVE_MODE_SAMPLE_ZERO_BIT: + return "zero"; + case VK_RESOLVE_MODE_AVERAGE_BIT: + return "average"; + case VK_RESOLVE_MODE_MIN_BIT: + return "min"; + case VK_RESOLVE_MODE_MAX_BIT: + return "max"; + default: + unreachable("invalid resolve mode"); + } +} + +nir_shader * +radv_meta_nir_build_depth_stencil_resolve_compute_shader(struct radv_device *dev, int samples, + enum radv_meta_resolve_type index, + VkResolveModeFlagBits resolve_mode) +{ + enum glsl_base_type img_base_type = index == RADV_META_DEPTH_RESOLVE ? GLSL_TYPE_FLOAT : GLSL_TYPE_UINT; + const struct glsl_type *sampler_type = glsl_sampler_type(GLSL_SAMPLER_DIM_MS, false, true, img_base_type); + const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, true, img_base_type); + + nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_resolve_cs_%s-%s-%d", + index == RADV_META_DEPTH_RESOLVE ? "depth" : "stencil", + get_resolve_mode_str(resolve_mode), samples); + b.shader->info.workgroup_size[0] = 8; + b.shader->info.workgroup_size[1] = 8; + + nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, sampler_type, "s_tex"); + input_img->data.descriptor_set = 0; + input_img->data.binding = 0; + + nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img"); + output_img->data.descriptor_set = 0; + output_img->data.binding = 1; + + nir_def *global_id = get_global_ids(&b, 3); + + nir_def *offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8); + + nir_def *resolve_coord = nir_iadd(&b, nir_trim_vector(&b, global_id, 2), offset); + + nir_def *img_coord = + nir_vec3(&b, nir_channel(&b, resolve_coord, 0), nir_channel(&b, resolve_coord, 1), nir_channel(&b, global_id, 2)); + + nir_deref_instr *input_img_deref = nir_build_deref_var(&b, input_img); + nir_def *outval = nir_txf_ms_deref(&b, input_img_deref, img_coord, nir_imm_int(&b, 0)); + + if (resolve_mode != VK_RESOLVE_MODE_SAMPLE_ZERO_BIT) { + for (int i = 1; i < samples; i++) { + nir_def *si = nir_txf_ms_deref(&b, input_img_deref, img_coord, nir_imm_int(&b, i)); + + switch (resolve_mode) { + case VK_RESOLVE_MODE_AVERAGE_BIT: + assert(index == RADV_META_DEPTH_RESOLVE); + outval = nir_fadd(&b, outval, si); + break; + case VK_RESOLVE_MODE_MIN_BIT: + if (index == RADV_META_DEPTH_RESOLVE) + outval = nir_fmin(&b, outval, si); + else + outval = nir_umin(&b, outval, si); + break; + case VK_RESOLVE_MODE_MAX_BIT: + if (index == RADV_META_DEPTH_RESOLVE) + outval = nir_fmax(&b, outval, si); + else + outval = nir_umax(&b, outval, si); + break; + default: + unreachable("invalid resolve mode"); + } + } + + if (resolve_mode == VK_RESOLVE_MODE_AVERAGE_BIT) + outval = nir_fdiv_imm(&b, outval, samples); + } + + nir_def *coord = nir_vec4(&b, nir_channel(&b, img_coord, 0), nir_channel(&b, img_coord, 1), + nir_channel(&b, img_coord, 2), nir_undef(&b, 1, 32)); + nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->def, coord, nir_undef(&b, 1, 32), outval, + nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D, .image_array = true); + return b.shader; +}