1 /*
2 * Copyright © 2021 Google
3 *
4 * SPDX-License-Identifier: MIT
5 */
6
7 #define AC_SURFACE_INCLUDE_NIR
8 #include "ac_surface.h"
9
10 #include "radv_meta.h"
11 #include "vk_common_entrypoints.h"
12
13 static nir_shader *
build_dcc_retile_compute_shader(struct radv_device * dev,struct radeon_surf * surf)14 build_dcc_retile_compute_shader(struct radv_device *dev, struct radeon_surf *surf)
15 {
16 const struct radv_physical_device *pdev = radv_device_physical(dev);
17 enum glsl_sampler_dim dim = GLSL_SAMPLER_DIM_BUF;
18 const struct glsl_type *buf_type = glsl_image_type(dim, false, GLSL_TYPE_UINT);
19 nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "dcc_retile_compute");
20
21 b.shader->info.workgroup_size[0] = 8;
22 b.shader->info.workgroup_size[1] = 8;
23
24 nir_def *src_dcc_size = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8);
25 nir_def *src_dcc_pitch = nir_channels(&b, src_dcc_size, 1);
26 nir_def *src_dcc_height = nir_channels(&b, src_dcc_size, 2);
27
28 nir_def *dst_dcc_size = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 8), .range = 8);
29 nir_def *dst_dcc_pitch = nir_channels(&b, dst_dcc_size, 1);
30 nir_def *dst_dcc_height = nir_channels(&b, dst_dcc_size, 2);
31 nir_variable *input_dcc = nir_variable_create(b.shader, nir_var_uniform, buf_type, "dcc_in");
32 input_dcc->data.descriptor_set = 0;
33 input_dcc->data.binding = 0;
34 nir_variable *output_dcc = nir_variable_create(b.shader, nir_var_uniform, buf_type, "dcc_out");
35 output_dcc->data.descriptor_set = 0;
36 output_dcc->data.binding = 1;
37
38 nir_def *input_dcc_ref = &nir_build_deref_var(&b, input_dcc)->def;
39 nir_def *output_dcc_ref = &nir_build_deref_var(&b, output_dcc)->def;
40
41 nir_def *coord = get_global_ids(&b, 2);
42 nir_def *zero = nir_imm_int(&b, 0);
43 coord =
44 nir_imul(&b, coord, nir_imm_ivec2(&b, surf->u.gfx9.color.dcc_block_width, surf->u.gfx9.color.dcc_block_height));
45
46 nir_def *src = ac_nir_dcc_addr_from_coord(&b, &pdev->info, surf->bpe, &surf->u.gfx9.color.dcc_equation,
47 src_dcc_pitch, src_dcc_height, zero, nir_channel(&b, coord, 0),
48 nir_channel(&b, coord, 1), zero, zero, zero);
49 nir_def *dst = ac_nir_dcc_addr_from_coord(&b, &pdev->info, surf->bpe, &surf->u.gfx9.color.display_dcc_equation,
50 dst_dcc_pitch, dst_dcc_height, zero, nir_channel(&b, coord, 0),
51 nir_channel(&b, coord, 1), zero, zero, zero);
52
53 nir_def *dcc_val = nir_image_deref_load(&b, 1, 32, input_dcc_ref, nir_vec4(&b, src, src, src, src),
54 nir_undef(&b, 1, 32), nir_imm_int(&b, 0), .image_dim = dim);
55
56 nir_image_deref_store(&b, output_dcc_ref, nir_vec4(&b, dst, dst, dst, dst), nir_undef(&b, 1, 32), dcc_val,
57 nir_imm_int(&b, 0), .image_dim = dim);
58
59 return b.shader;
60 }
61
62 static VkResult
get_pipeline_layout(struct radv_device * device,VkPipelineLayout * layout_out)63 get_pipeline_layout(struct radv_device *device, VkPipelineLayout *layout_out)
64 {
65 enum radv_meta_object_key_type key = RADV_META_OBJECT_KEY_DCC_RETILE;
66
67 const VkDescriptorSetLayoutBinding bindings[] = {
68 {
69 .binding = 0,
70 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
71 .descriptorCount = 1,
72 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
73 },
74 {
75 .binding = 1,
76 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
77 .descriptorCount = 1,
78 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
79 },
80
81 };
82
83 const VkDescriptorSetLayoutCreateInfo desc_info = {
84 .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
85 .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT,
86 .bindingCount = 2,
87 .pBindings = bindings,
88 };
89
90 const VkPushConstantRange pc_range = {
91 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
92 .size = 16,
93 };
94
95 return vk_meta_get_pipeline_layout(&device->vk, &device->meta_state.device, &desc_info, &pc_range, &key, sizeof(key),
96 layout_out);
97 }
98
99 struct radv_dcc_retile_key {
100 enum radv_meta_object_key_type type;
101 uint32_t swizzle;
102 };
103
104 /*
105 * This take a surface, but the only things used are:
106 * - BPE
107 * - DCC equations
108 * - DCC block size
109 *
110 * BPE is always 4 at the moment and the rest is derived from the tilemode.
111 */
112 static VkResult
get_pipeline(struct radv_device * device,struct radv_image * image,VkPipeline * pipeline_out,VkPipelineLayout * layout_out)113 get_pipeline(struct radv_device *device, struct radv_image *image, VkPipeline *pipeline_out,
114 VkPipelineLayout *layout_out)
115 {
116 const unsigned swizzle_mode = image->planes[0].surface.u.gfx9.swizzle_mode;
117 struct radv_dcc_retile_key key;
118 VkResult result;
119
120 result = get_pipeline_layout(device, layout_out);
121 if (result != VK_SUCCESS)
122 return result;
123
124 memset(&key, 0, sizeof(key));
125 key.type = RADV_META_OBJECT_KEY_DCC_RETILE;
126 key.swizzle = swizzle_mode;
127
128 VkPipeline pipeline_from_cache = vk_meta_lookup_pipeline(&device->meta_state.device, &key, sizeof(key));
129 if (pipeline_from_cache != VK_NULL_HANDLE) {
130 *pipeline_out = pipeline_from_cache;
131 return VK_SUCCESS;
132 }
133
134 nir_shader *cs = build_dcc_retile_compute_shader(device, &image->planes[0].surface);
135
136 const VkPipelineShaderStageCreateInfo stage_info = {
137 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
138 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
139 .module = vk_shader_module_handle_from_nir(cs),
140 .pName = "main",
141 .pSpecializationInfo = NULL,
142 };
143
144 const VkComputePipelineCreateInfo pipeline_info = {
145 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
146 .stage = stage_info,
147 .flags = 0,
148 .layout = *layout_out,
149 };
150
151 result = vk_meta_create_compute_pipeline(&device->vk, &device->meta_state.device, &pipeline_info, &key, sizeof(key),
152 pipeline_out);
153
154 ralloc_free(cs);
155 return result;
156 }
157
158 void
radv_retile_dcc(struct radv_cmd_buffer * cmd_buffer,struct radv_image * image)159 radv_retile_dcc(struct radv_cmd_buffer *cmd_buffer, struct radv_image *image)
160 {
161 struct radv_meta_saved_state saved_state;
162 struct radv_device *device = radv_cmd_buffer_device(cmd_buffer);
163 struct radv_buffer buffer;
164 VkPipelineLayout layout;
165 VkPipeline pipeline;
166 VkResult result;
167
168 assert(image->vk.image_type == VK_IMAGE_TYPE_2D);
169 assert(image->vk.array_layers == 1 && image->vk.mip_levels == 1);
170
171 struct radv_cmd_state *state = &cmd_buffer->state;
172
173 result = get_pipeline(device, image, &pipeline, &layout);
174 if (result != VK_SUCCESS) {
175 vk_command_buffer_set_error(&cmd_buffer->vk, result);
176 return;
177 }
178
179 state->flush_bits |= radv_dst_access_flush(cmd_buffer, VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT,
180 VK_ACCESS_2_SHADER_READ_BIT, 0, image, NULL);
181
182 radv_meta_save(&saved_state, cmd_buffer,
183 RADV_META_SAVE_DESCRIPTORS | RADV_META_SAVE_COMPUTE_PIPELINE | RADV_META_SAVE_CONSTANTS);
184
185 radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE, pipeline);
186
187 radv_buffer_init(&buffer, device, image->bindings[0].bo, image->size, image->bindings[0].offset);
188
189 struct radv_buffer_view views[2];
190 VkBufferView view_handles[2];
191 radv_buffer_view_init(views, device,
192 &(VkBufferViewCreateInfo){
193 .sType = VK_STRUCTURE_TYPE_BUFFER_VIEW_CREATE_INFO,
194 .buffer = radv_buffer_to_handle(&buffer),
195 .offset = image->planes[0].surface.meta_offset,
196 .range = image->planes[0].surface.meta_size,
197 .format = VK_FORMAT_R8_UINT,
198 });
199 radv_buffer_view_init(views + 1, device,
200 &(VkBufferViewCreateInfo){
201 .sType = VK_STRUCTURE_TYPE_BUFFER_VIEW_CREATE_INFO,
202 .buffer = radv_buffer_to_handle(&buffer),
203 .offset = image->planes[0].surface.display_dcc_offset,
204 .range = image->planes[0].surface.u.gfx9.color.display_dcc_size,
205 .format = VK_FORMAT_R8_UINT,
206 });
207 for (unsigned i = 0; i < 2; ++i)
208 view_handles[i] = radv_buffer_view_to_handle(&views[i]);
209
210 radv_meta_push_descriptor_set(cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, layout, 0, 2,
211 (VkWriteDescriptorSet[]){
212 {
213 .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
214 .dstBinding = 0,
215 .dstArrayElement = 0,
216 .descriptorCount = 1,
217 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
218 .pTexelBufferView = &view_handles[0],
219 },
220 {
221 .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
222 .dstBinding = 1,
223 .dstArrayElement = 0,
224 .descriptorCount = 1,
225 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
226 .pTexelBufferView = &view_handles[1],
227 },
228 });
229
230 unsigned width = DIV_ROUND_UP(image->vk.extent.width, vk_format_get_blockwidth(image->vk.format));
231 unsigned height = DIV_ROUND_UP(image->vk.extent.height, vk_format_get_blockheight(image->vk.format));
232
233 unsigned dcc_width = DIV_ROUND_UP(width, image->planes[0].surface.u.gfx9.color.dcc_block_width);
234 unsigned dcc_height = DIV_ROUND_UP(height, image->planes[0].surface.u.gfx9.color.dcc_block_height);
235
236 uint32_t constants[] = {
237 image->planes[0].surface.u.gfx9.color.dcc_pitch_max + 1,
238 image->planes[0].surface.u.gfx9.color.dcc_height,
239 image->planes[0].surface.u.gfx9.color.display_dcc_pitch_max + 1,
240 image->planes[0].surface.u.gfx9.color.display_dcc_height,
241 };
242 vk_common_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer), layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, 16,
243 constants);
244
245 radv_unaligned_dispatch(cmd_buffer, dcc_width, dcc_height, 1);
246
247 radv_buffer_view_finish(views);
248 radv_buffer_view_finish(views + 1);
249 radv_buffer_finish(&buffer);
250
251 radv_meta_restore(&saved_state, cmd_buffer);
252
253 state->flush_bits |=
254 RADV_CMD_FLAG_CS_PARTIAL_FLUSH | radv_src_access_flush(cmd_buffer, VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT,
255 VK_ACCESS_2_SHADER_WRITE_BIT, 0, image, NULL);
256 }
257