1 /*
2 * Copyright © 2019 Valve Corporation
3 * Copyright © 2018 Red Hat
4 *
5 * SPDX-License-Identifier: MIT
6 */
7
8 #include "radv_formats.h"
9 #include "radv_meta.h"
10 #include "vk_format.h"
11
12 static nir_shader *
build_fmask_expand_compute_shader(struct radv_device * device,int samples)13 build_fmask_expand_compute_shader(struct radv_device *device, int samples)
14 {
15 const struct glsl_type *type = glsl_sampler_type(GLSL_SAMPLER_DIM_MS, false, true, GLSL_TYPE_FLOAT);
16 const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_MS, true, GLSL_TYPE_FLOAT);
17
18 nir_builder b = radv_meta_init_shader(device, MESA_SHADER_COMPUTE, "meta_fmask_expand_cs-%d", samples);
19 b.shader->info.workgroup_size[0] = 8;
20 b.shader->info.workgroup_size[1] = 8;
21
22 nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, type, "s_tex");
23 input_img->data.descriptor_set = 0;
24 input_img->data.binding = 0;
25
26 nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img");
27 output_img->data.descriptor_set = 0;
28 output_img->data.binding = 1;
29 output_img->data.access = ACCESS_NON_READABLE;
30
31 nir_deref_instr *input_img_deref = nir_build_deref_var(&b, input_img);
32 nir_def *output_img_deref = &nir_build_deref_var(&b, output_img)->def;
33
34 nir_def *tex_coord = get_global_ids(&b, 3);
35
36 nir_def *tex_vals[8];
37 for (uint32_t i = 0; i < samples; i++) {
38 tex_vals[i] = nir_txf_ms_deref(&b, input_img_deref, tex_coord, nir_imm_int(&b, i));
39 }
40
41 nir_def *img_coord = nir_vec4(&b, nir_channel(&b, tex_coord, 0), nir_channel(&b, tex_coord, 1),
42 nir_channel(&b, tex_coord, 2), nir_undef(&b, 1, 32));
43
44 for (uint32_t i = 0; i < samples; i++) {
45 nir_image_deref_store(&b, output_img_deref, img_coord, nir_imm_int(&b, i), tex_vals[i], nir_imm_int(&b, 0),
46 .image_dim = GLSL_SAMPLER_DIM_MS, .image_array = true);
47 }
48
49 return b.shader;
50 }
51
52 static VkResult
get_pipeline_layout(struct radv_device * device,VkPipelineLayout * layout_out)53 get_pipeline_layout(struct radv_device *device, VkPipelineLayout *layout_out)
54 {
55 enum radv_meta_object_key_type key = RADV_META_OBJECT_KEY_FMASK_EXPAND;
56
57 const VkDescriptorSetLayoutBinding bindings[] = {
58 {
59 .binding = 0,
60 .descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,
61 .descriptorCount = 1,
62 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
63 },
64 {
65 .binding = 1,
66 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
67 .descriptorCount = 1,
68 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
69 },
70 };
71
72 const VkDescriptorSetLayoutCreateInfo desc_info = {
73 .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
74 .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT,
75 .bindingCount = 2,
76 .pBindings = bindings,
77 };
78
79 return vk_meta_get_pipeline_layout(&device->vk, &device->meta_state.device, &desc_info, NULL, &key, sizeof(key),
80 layout_out);
81 }
82
83 struct radv_fmask_expand_key {
84 enum radv_meta_object_key_type type;
85 uint32_t samples;
86 };
87
88 static VkResult
get_pipeline(struct radv_device * device,uint32_t samples_log2,VkPipeline * pipeline_out,VkPipelineLayout * layout_out)89 get_pipeline(struct radv_device *device, uint32_t samples_log2, VkPipeline *pipeline_out, VkPipelineLayout *layout_out)
90 {
91 const uint32_t samples = 1 << samples_log2;
92 struct radv_fmask_expand_key key;
93 VkResult result;
94
95 result = get_pipeline_layout(device, layout_out);
96 if (result != VK_SUCCESS)
97 return result;
98
99 memset(&key, 0, sizeof(key));
100 key.type = RADV_META_OBJECT_KEY_FMASK_EXPAND;
101 key.samples = samples;
102
103 VkPipeline pipeline_from_cache = vk_meta_lookup_pipeline(&device->meta_state.device, &key, sizeof(key));
104 if (pipeline_from_cache != VK_NULL_HANDLE) {
105 *pipeline_out = pipeline_from_cache;
106 return VK_SUCCESS;
107 }
108
109 nir_shader *cs = build_fmask_expand_compute_shader(device, samples);
110
111 const VkPipelineShaderStageCreateInfo stage_info = {
112 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
113 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
114 .module = vk_shader_module_handle_from_nir(cs),
115 .pName = "main",
116 .pSpecializationInfo = NULL,
117 };
118
119 const VkComputePipelineCreateInfo pipeline_info = {
120 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
121 .stage = stage_info,
122 .flags = 0,
123 .layout = *layout_out,
124 };
125
126 result = vk_meta_create_compute_pipeline(&device->vk, &device->meta_state.device, &pipeline_info, &key, sizeof(key),
127 pipeline_out);
128
129 ralloc_free(cs);
130 return result;
131 }
132
133 void
radv_expand_fmask_image_inplace(struct radv_cmd_buffer * cmd_buffer,struct radv_image * image,const VkImageSubresourceRange * subresourceRange)134 radv_expand_fmask_image_inplace(struct radv_cmd_buffer *cmd_buffer, struct radv_image *image,
135 const VkImageSubresourceRange *subresourceRange)
136 {
137 struct radv_device *device = radv_cmd_buffer_device(cmd_buffer);
138 struct radv_meta_saved_state saved_state;
139 const uint32_t samples = image->vk.samples;
140 const uint32_t samples_log2 = ffs(samples) - 1;
141 unsigned layer_count = vk_image_subresource_layer_count(&image->vk, subresourceRange);
142 struct radv_image_view iview;
143 VkPipelineLayout layout;
144 VkPipeline pipeline;
145 VkResult result;
146
147 result = get_pipeline(device, samples_log2, &pipeline, &layout);
148 if (result != VK_SUCCESS) {
149 vk_command_buffer_set_error(&cmd_buffer->vk, result);
150 return;
151 }
152
153 radv_meta_save(&saved_state, cmd_buffer, RADV_META_SAVE_COMPUTE_PIPELINE | RADV_META_SAVE_DESCRIPTORS);
154
155 radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE, pipeline);
156
157 radv_image_view_init(&iview, device,
158 &(VkImageViewCreateInfo){
159 .sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO,
160 .image = radv_image_to_handle(image),
161 .viewType = radv_meta_get_view_type(image),
162 .format = vk_format_no_srgb(image->vk.format),
163 .subresourceRange =
164 {
165 .aspectMask = subresourceRange->aspectMask,
166 .baseMipLevel = 0,
167 .levelCount = 1,
168 .baseArrayLayer = subresourceRange->baseArrayLayer,
169 .layerCount = layer_count,
170 },
171 },
172 NULL);
173
174 const VkImageSubresourceRange range = vk_image_view_subresource_range(&iview.vk);
175
176 cmd_buffer->state.flush_bits |= radv_dst_access_flush(cmd_buffer, VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT,
177 VK_ACCESS_2_SHADER_READ_BIT, 0, image, &range);
178
179 radv_meta_push_descriptor_set(cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, layout, 0, 2,
180 (VkWriteDescriptorSet[]){{.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
181 .dstBinding = 0,
182 .dstArrayElement = 0,
183 .descriptorCount = 1,
184 .descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,
185 .pImageInfo =
186 (VkDescriptorImageInfo[]){
187 {.sampler = VK_NULL_HANDLE,
188 .imageView = radv_image_view_to_handle(&iview),
189 .imageLayout = VK_IMAGE_LAYOUT_GENERAL},
190 }},
191 {.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
192 .dstBinding = 1,
193 .dstArrayElement = 0,
194 .descriptorCount = 1,
195 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
196 .pImageInfo = (VkDescriptorImageInfo[]){
197 {.sampler = VK_NULL_HANDLE,
198 .imageView = radv_image_view_to_handle(&iview),
199 .imageLayout = VK_IMAGE_LAYOUT_GENERAL},
200 }}});
201
202 radv_unaligned_dispatch(cmd_buffer, image->vk.extent.width, image->vk.extent.height, layer_count);
203
204 radv_image_view_finish(&iview);
205
206 radv_meta_restore(&saved_state, cmd_buffer);
207
208 cmd_buffer->state.flush_bits |=
209 RADV_CMD_FLAG_CS_PARTIAL_FLUSH | radv_src_access_flush(cmd_buffer, VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT,
210 VK_ACCESS_2_SHADER_WRITE_BIT, 0, image, &range);
211
212 /* Re-initialize FMASK in fully expanded mode. */
213 cmd_buffer->state.flush_bits |= radv_init_fmask(cmd_buffer, image, subresourceRange);
214 }
215