1 /*
2 * Copyright © 2016 Red Hat.
3 * Copyright © 2016 Bas Nieuwenhuizen
4 *
5 * Permission is hereby granted, free of charge, to any person obtaining a
6 * copy of this software and associated documentation files (the "Software"),
7 * to deal in the Software without restriction, including without limitation
8 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
9 * and/or sell copies of the Software, and to permit persons to whom the
10 * Software is furnished to do so, subject to the following conditions:
11 *
12 * The above copyright notice and this permission notice (including the next
13 * paragraph) shall be included in all copies or substantial portions of the
14 * Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
19 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
21 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
22 * IN THE SOFTWARE.
23 */
24 #include "nir/nir_builder.h"
25 #include "radv_meta.h"
26
27 /*
28 * GFX queue: Compute shader implementation of image->buffer copy
29 * Compute queue: implementation also of buffer->image, image->image, and image clear.
30 */
31
32 static nir_shader *
build_nir_itob_compute_shader(struct radv_device * dev,bool is_3d)33 build_nir_itob_compute_shader(struct radv_device *dev, bool is_3d)
34 {
35 enum glsl_sampler_dim dim = is_3d ? GLSL_SAMPLER_DIM_3D : GLSL_SAMPLER_DIM_2D;
36 const struct glsl_type *sampler_type = glsl_sampler_type(dim, false, false, GLSL_TYPE_FLOAT);
37 const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_BUF, false, GLSL_TYPE_FLOAT);
38 nir_builder b =
39 radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, is_3d ? "meta_itob_cs_3d" : "meta_itob_cs");
40 b.shader->info.workgroup_size[0] = 8;
41 b.shader->info.workgroup_size[1] = 8;
42 nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, sampler_type, "s_tex");
43 input_img->data.descriptor_set = 0;
44 input_img->data.binding = 0;
45
46 nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img");
47 output_img->data.descriptor_set = 0;
48 output_img->data.binding = 1;
49
50 nir_ssa_def *global_id = get_global_ids(&b, is_3d ? 3 : 2);
51
52 nir_ssa_def *offset =
53 nir_load_push_constant(&b, is_3d ? 3 : 2, 32, nir_imm_int(&b, 0), .range = is_3d ? 12 : 8);
54 nir_ssa_def *stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 12), .range = 16);
55
56 nir_ssa_def *img_coord = nir_iadd(&b, global_id, offset);
57 nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa;
58
59 nir_tex_instr *tex = nir_tex_instr_create(b.shader, 3);
60 tex->sampler_dim = dim;
61 tex->op = nir_texop_txf;
62 tex->src[0].src_type = nir_tex_src_coord;
63 tex->src[0].src = nir_src_for_ssa(nir_trim_vector(&b, img_coord, 2 + is_3d));
64 tex->src[1].src_type = nir_tex_src_lod;
65 tex->src[1].src = nir_src_for_ssa(nir_imm_int(&b, 0));
66 tex->src[2].src_type = nir_tex_src_texture_deref;
67 tex->src[2].src = nir_src_for_ssa(input_img_deref);
68 tex->dest_type = nir_type_float32;
69 tex->is_array = false;
70 tex->coord_components = is_3d ? 3 : 2;
71
72 nir_ssa_dest_init(&tex->instr, &tex->dest, 4, 32, "tex");
73 nir_builder_instr_insert(&b, &tex->instr);
74
75 nir_ssa_def *pos_x = nir_channel(&b, global_id, 0);
76 nir_ssa_def *pos_y = nir_channel(&b, global_id, 1);
77
78 nir_ssa_def *tmp = nir_imul(&b, pos_y, stride);
79 tmp = nir_iadd(&b, tmp, pos_x);
80
81 nir_ssa_def *coord = nir_vec4(&b, tmp, tmp, tmp, tmp);
82
83 nir_ssa_def *outval = &tex->dest.ssa;
84 nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, coord,
85 nir_ssa_undef(&b, 1, 32), outval, nir_imm_int(&b, 0),
86 .image_dim = GLSL_SAMPLER_DIM_BUF);
87
88 return b.shader;
89 }
90
91 /* Image to buffer - don't write use image accessors */
92 static VkResult
radv_device_init_meta_itob_state(struct radv_device * device)93 radv_device_init_meta_itob_state(struct radv_device *device)
94 {
95 VkResult result;
96 nir_shader *cs = build_nir_itob_compute_shader(device, false);
97 nir_shader *cs_3d = build_nir_itob_compute_shader(device, true);
98
99 /*
100 * two descriptors one for the image being sampled
101 * one for the buffer being written.
102 */
103 VkDescriptorSetLayoutCreateInfo ds_create_info = {
104 .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
105 .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,
106 .bindingCount = 2,
107 .pBindings = (VkDescriptorSetLayoutBinding[]){
108 {.binding = 0,
109 .descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,
110 .descriptorCount = 1,
111 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
112 .pImmutableSamplers = NULL},
113 {.binding = 1,
114 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
115 .descriptorCount = 1,
116 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
117 .pImmutableSamplers = NULL},
118 }};
119
120 result = radv_CreateDescriptorSetLayout(radv_device_to_handle(device), &ds_create_info,
121 &device->meta_state.alloc,
122 &device->meta_state.itob.img_ds_layout);
123 if (result != VK_SUCCESS)
124 goto fail;
125
126 VkPipelineLayoutCreateInfo pl_create_info = {
127 .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,
128 .setLayoutCount = 1,
129 .pSetLayouts = &device->meta_state.itob.img_ds_layout,
130 .pushConstantRangeCount = 1,
131 .pPushConstantRanges = &(VkPushConstantRange){VK_SHADER_STAGE_COMPUTE_BIT, 0, 16},
132 };
133
134 result =
135 radv_CreatePipelineLayout(radv_device_to_handle(device), &pl_create_info,
136 &device->meta_state.alloc, &device->meta_state.itob.img_p_layout);
137 if (result != VK_SUCCESS)
138 goto fail;
139
140 /* compute shader */
141
142 VkPipelineShaderStageCreateInfo pipeline_shader_stage = {
143 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
144 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
145 .module = vk_shader_module_handle_from_nir(cs),
146 .pName = "main",
147 .pSpecializationInfo = NULL,
148 };
149
150 VkComputePipelineCreateInfo vk_pipeline_info = {
151 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
152 .stage = pipeline_shader_stage,
153 .flags = 0,
154 .layout = device->meta_state.itob.img_p_layout,
155 };
156
157 result = radv_CreateComputePipelines(radv_device_to_handle(device),
158 radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,
159 &vk_pipeline_info, NULL, &device->meta_state.itob.pipeline);
160 if (result != VK_SUCCESS)
161 goto fail;
162
163 VkPipelineShaderStageCreateInfo pipeline_shader_stage_3d = {
164 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
165 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
166 .module = vk_shader_module_handle_from_nir(cs_3d),
167 .pName = "main",
168 .pSpecializationInfo = NULL,
169 };
170
171 VkComputePipelineCreateInfo vk_pipeline_info_3d = {
172 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
173 .stage = pipeline_shader_stage_3d,
174 .flags = 0,
175 .layout = device->meta_state.itob.img_p_layout,
176 };
177
178 result = radv_CreateComputePipelines(
179 radv_device_to_handle(device), radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,
180 &vk_pipeline_info_3d, NULL, &device->meta_state.itob.pipeline_3d);
181 if (result != VK_SUCCESS)
182 goto fail;
183
184 ralloc_free(cs_3d);
185 ralloc_free(cs);
186
187 return VK_SUCCESS;
188 fail:
189 ralloc_free(cs);
190 ralloc_free(cs_3d);
191 return result;
192 }
193
194 static void
radv_device_finish_meta_itob_state(struct radv_device * device)195 radv_device_finish_meta_itob_state(struct radv_device *device)
196 {
197 struct radv_meta_state *state = &device->meta_state;
198
199 radv_DestroyPipelineLayout(radv_device_to_handle(device), state->itob.img_p_layout,
200 &state->alloc);
201 device->vk.dispatch_table.DestroyDescriptorSetLayout(radv_device_to_handle(device),
202 state->itob.img_ds_layout, &state->alloc);
203 radv_DestroyPipeline(radv_device_to_handle(device), state->itob.pipeline, &state->alloc);
204 radv_DestroyPipeline(radv_device_to_handle(device), state->itob.pipeline_3d, &state->alloc);
205 }
206
207 static nir_shader *
build_nir_btoi_compute_shader(struct radv_device * dev,bool is_3d)208 build_nir_btoi_compute_shader(struct radv_device *dev, bool is_3d)
209 {
210 enum glsl_sampler_dim dim = is_3d ? GLSL_SAMPLER_DIM_3D : GLSL_SAMPLER_DIM_2D;
211 const struct glsl_type *buf_type =
212 glsl_sampler_type(GLSL_SAMPLER_DIM_BUF, false, false, GLSL_TYPE_FLOAT);
213 const struct glsl_type *img_type = glsl_image_type(dim, false, GLSL_TYPE_FLOAT);
214 nir_builder b =
215 radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, is_3d ? "meta_btoi_cs_3d" : "meta_btoi_cs");
216 b.shader->info.workgroup_size[0] = 8;
217 b.shader->info.workgroup_size[1] = 8;
218 nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, buf_type, "s_tex");
219 input_img->data.descriptor_set = 0;
220 input_img->data.binding = 0;
221
222 nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img");
223 output_img->data.descriptor_set = 0;
224 output_img->data.binding = 1;
225
226 nir_ssa_def *global_id = get_global_ids(&b, is_3d ? 3 : 2);
227
228 nir_ssa_def *offset =
229 nir_load_push_constant(&b, is_3d ? 3 : 2, 32, nir_imm_int(&b, 0), .range = is_3d ? 12 : 8);
230 nir_ssa_def *stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 12), .range = 16);
231
232 nir_ssa_def *pos_x = nir_channel(&b, global_id, 0);
233 nir_ssa_def *pos_y = nir_channel(&b, global_id, 1);
234
235 nir_ssa_def *buf_coord = nir_imul(&b, pos_y, stride);
236 buf_coord = nir_iadd(&b, buf_coord, pos_x);
237
238 nir_ssa_def *coord = nir_iadd(&b, global_id, offset);
239 nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa;
240
241 nir_tex_instr *tex = nir_tex_instr_create(b.shader, 3);
242 tex->sampler_dim = GLSL_SAMPLER_DIM_BUF;
243 tex->op = nir_texop_txf;
244 tex->src[0].src_type = nir_tex_src_coord;
245 tex->src[0].src = nir_src_for_ssa(buf_coord);
246 tex->src[1].src_type = nir_tex_src_lod;
247 tex->src[1].src = nir_src_for_ssa(nir_imm_int(&b, 0));
248 tex->src[2].src_type = nir_tex_src_texture_deref;
249 tex->src[2].src = nir_src_for_ssa(input_img_deref);
250 tex->dest_type = nir_type_float32;
251 tex->is_array = false;
252 tex->coord_components = 1;
253
254 nir_ssa_dest_init(&tex->instr, &tex->dest, 4, 32, "tex");
255 nir_builder_instr_insert(&b, &tex->instr);
256
257 nir_ssa_def *outval = &tex->dest.ssa;
258
259 nir_ssa_def *img_coord = nir_vec4(&b, nir_channel(&b, coord, 0),
260 nir_channel(&b, coord, 1),
261 is_3d ? nir_channel(&b, coord, 2) : nir_ssa_undef(&b, 1, 32),
262 nir_ssa_undef(&b, 1, 32));
263
264 nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, img_coord,
265 nir_ssa_undef(&b, 1, 32), outval, nir_imm_int(&b, 0), .image_dim = dim);
266
267 return b.shader;
268 }
269
270 /* Buffer to image - don't write use image accessors */
271 static VkResult
radv_device_init_meta_btoi_state(struct radv_device * device)272 radv_device_init_meta_btoi_state(struct radv_device *device)
273 {
274 VkResult result;
275 nir_shader *cs = build_nir_btoi_compute_shader(device, false);
276 nir_shader *cs_3d = build_nir_btoi_compute_shader(device, true);
277 /*
278 * two descriptors one for the image being sampled
279 * one for the buffer being written.
280 */
281 VkDescriptorSetLayoutCreateInfo ds_create_info = {
282 .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
283 .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,
284 .bindingCount = 2,
285 .pBindings = (VkDescriptorSetLayoutBinding[]){
286 {.binding = 0,
287 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
288 .descriptorCount = 1,
289 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
290 .pImmutableSamplers = NULL},
291 {.binding = 1,
292 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
293 .descriptorCount = 1,
294 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
295 .pImmutableSamplers = NULL},
296 }};
297
298 result = radv_CreateDescriptorSetLayout(radv_device_to_handle(device), &ds_create_info,
299 &device->meta_state.alloc,
300 &device->meta_state.btoi.img_ds_layout);
301 if (result != VK_SUCCESS)
302 goto fail;
303
304 VkPipelineLayoutCreateInfo pl_create_info = {
305 .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,
306 .setLayoutCount = 1,
307 .pSetLayouts = &device->meta_state.btoi.img_ds_layout,
308 .pushConstantRangeCount = 1,
309 .pPushConstantRanges = &(VkPushConstantRange){VK_SHADER_STAGE_COMPUTE_BIT, 0, 16},
310 };
311
312 result =
313 radv_CreatePipelineLayout(radv_device_to_handle(device), &pl_create_info,
314 &device->meta_state.alloc, &device->meta_state.btoi.img_p_layout);
315 if (result != VK_SUCCESS)
316 goto fail;
317
318 /* compute shader */
319
320 VkPipelineShaderStageCreateInfo pipeline_shader_stage = {
321 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
322 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
323 .module = vk_shader_module_handle_from_nir(cs),
324 .pName = "main",
325 .pSpecializationInfo = NULL,
326 };
327
328 VkComputePipelineCreateInfo vk_pipeline_info = {
329 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
330 .stage = pipeline_shader_stage,
331 .flags = 0,
332 .layout = device->meta_state.btoi.img_p_layout,
333 };
334
335 result = radv_CreateComputePipelines(radv_device_to_handle(device),
336 radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,
337 &vk_pipeline_info, NULL, &device->meta_state.btoi.pipeline);
338 if (result != VK_SUCCESS)
339 goto fail;
340
341 VkPipelineShaderStageCreateInfo pipeline_shader_stage_3d = {
342 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
343 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
344 .module = vk_shader_module_handle_from_nir(cs_3d),
345 .pName = "main",
346 .pSpecializationInfo = NULL,
347 };
348
349 VkComputePipelineCreateInfo vk_pipeline_info_3d = {
350 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
351 .stage = pipeline_shader_stage_3d,
352 .flags = 0,
353 .layout = device->meta_state.btoi.img_p_layout,
354 };
355
356 result = radv_CreateComputePipelines(
357 radv_device_to_handle(device), radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,
358 &vk_pipeline_info_3d, NULL, &device->meta_state.btoi.pipeline_3d);
359
360 ralloc_free(cs_3d);
361 ralloc_free(cs);
362
363 return VK_SUCCESS;
364 fail:
365 ralloc_free(cs_3d);
366 ralloc_free(cs);
367 return result;
368 }
369
370 static void
radv_device_finish_meta_btoi_state(struct radv_device * device)371 radv_device_finish_meta_btoi_state(struct radv_device *device)
372 {
373 struct radv_meta_state *state = &device->meta_state;
374
375 radv_DestroyPipelineLayout(radv_device_to_handle(device), state->btoi.img_p_layout,
376 &state->alloc);
377 device->vk.dispatch_table.DestroyDescriptorSetLayout(radv_device_to_handle(device),
378 state->btoi.img_ds_layout, &state->alloc);
379 radv_DestroyPipeline(radv_device_to_handle(device), state->btoi.pipeline, &state->alloc);
380 radv_DestroyPipeline(radv_device_to_handle(device), state->btoi.pipeline_3d, &state->alloc);
381 }
382
383 /* Buffer to image - special path for R32G32B32 */
384 static nir_shader *
build_nir_btoi_r32g32b32_compute_shader(struct radv_device * dev)385 build_nir_btoi_r32g32b32_compute_shader(struct radv_device *dev)
386 {
387 const struct glsl_type *buf_type =
388 glsl_sampler_type(GLSL_SAMPLER_DIM_BUF, false, false, GLSL_TYPE_FLOAT);
389 const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_BUF, false, GLSL_TYPE_FLOAT);
390 nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_btoi_r32g32b32_cs");
391 b.shader->info.workgroup_size[0] = 8;
392 b.shader->info.workgroup_size[1] = 8;
393 nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, buf_type, "s_tex");
394 input_img->data.descriptor_set = 0;
395 input_img->data.binding = 0;
396
397 nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img");
398 output_img->data.descriptor_set = 0;
399 output_img->data.binding = 1;
400
401 nir_ssa_def *global_id = get_global_ids(&b, 2);
402
403 nir_ssa_def *offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8);
404 nir_ssa_def *pitch = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 8), .range = 12);
405 nir_ssa_def *stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 12), .range = 16);
406
407 nir_ssa_def *pos_x = nir_channel(&b, global_id, 0);
408 nir_ssa_def *pos_y = nir_channel(&b, global_id, 1);
409
410 nir_ssa_def *buf_coord = nir_imul(&b, pos_y, stride);
411 buf_coord = nir_iadd(&b, buf_coord, pos_x);
412
413 nir_ssa_def *img_coord = nir_iadd(&b, global_id, offset);
414
415 nir_ssa_def *global_pos = nir_iadd(&b, nir_imul(&b, nir_channel(&b, img_coord, 1), pitch),
416 nir_imul_imm(&b, nir_channel(&b, img_coord, 0), 3));
417
418 nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa;
419
420 nir_tex_instr *tex = nir_tex_instr_create(b.shader, 3);
421 tex->sampler_dim = GLSL_SAMPLER_DIM_BUF;
422 tex->op = nir_texop_txf;
423 tex->src[0].src_type = nir_tex_src_coord;
424 tex->src[0].src = nir_src_for_ssa(buf_coord);
425 tex->src[1].src_type = nir_tex_src_lod;
426 tex->src[1].src = nir_src_for_ssa(nir_imm_int(&b, 0));
427 tex->src[2].src_type = nir_tex_src_texture_deref;
428 tex->src[2].src = nir_src_for_ssa(input_img_deref);
429 tex->dest_type = nir_type_float32;
430 tex->is_array = false;
431 tex->coord_components = 1;
432 nir_ssa_dest_init(&tex->instr, &tex->dest, 4, 32, "tex");
433 nir_builder_instr_insert(&b, &tex->instr);
434
435 nir_ssa_def *outval = &tex->dest.ssa;
436
437 for (int chan = 0; chan < 3; chan++) {
438 nir_ssa_def *local_pos = nir_iadd_imm(&b, global_pos, chan);
439
440 nir_ssa_def *coord = nir_vec4(&b, local_pos, local_pos, local_pos, local_pos);
441
442 nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, coord,
443 nir_ssa_undef(&b, 1, 32), nir_channel(&b, outval, chan),
444 nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_BUF);
445 }
446
447 return b.shader;
448 }
449
450 static VkResult
radv_device_init_meta_btoi_r32g32b32_state(struct radv_device * device)451 radv_device_init_meta_btoi_r32g32b32_state(struct radv_device *device)
452 {
453 VkResult result;
454 nir_shader *cs = build_nir_btoi_r32g32b32_compute_shader(device);
455
456 VkDescriptorSetLayoutCreateInfo ds_create_info = {
457 .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
458 .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,
459 .bindingCount = 2,
460 .pBindings = (VkDescriptorSetLayoutBinding[]){
461 {.binding = 0,
462 .descriptorType = VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER,
463 .descriptorCount = 1,
464 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
465 .pImmutableSamplers = NULL},
466 {.binding = 1,
467 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
468 .descriptorCount = 1,
469 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
470 .pImmutableSamplers = NULL},
471 }};
472
473 result = radv_CreateDescriptorSetLayout(radv_device_to_handle(device), &ds_create_info,
474 &device->meta_state.alloc,
475 &device->meta_state.btoi_r32g32b32.img_ds_layout);
476 if (result != VK_SUCCESS)
477 goto fail;
478
479 VkPipelineLayoutCreateInfo pl_create_info = {
480 .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,
481 .setLayoutCount = 1,
482 .pSetLayouts = &device->meta_state.btoi_r32g32b32.img_ds_layout,
483 .pushConstantRangeCount = 1,
484 .pPushConstantRanges = &(VkPushConstantRange){VK_SHADER_STAGE_COMPUTE_BIT, 0, 16},
485 };
486
487 result = radv_CreatePipelineLayout(radv_device_to_handle(device), &pl_create_info,
488 &device->meta_state.alloc,
489 &device->meta_state.btoi_r32g32b32.img_p_layout);
490 if (result != VK_SUCCESS)
491 goto fail;
492
493 /* compute shader */
494
495 VkPipelineShaderStageCreateInfo pipeline_shader_stage = {
496 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
497 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
498 .module = vk_shader_module_handle_from_nir(cs),
499 .pName = "main",
500 .pSpecializationInfo = NULL,
501 };
502
503 VkComputePipelineCreateInfo vk_pipeline_info = {
504 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
505 .stage = pipeline_shader_stage,
506 .flags = 0,
507 .layout = device->meta_state.btoi_r32g32b32.img_p_layout,
508 };
509
510 result = radv_CreateComputePipelines(
511 radv_device_to_handle(device), radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,
512 &vk_pipeline_info, NULL, &device->meta_state.btoi_r32g32b32.pipeline);
513
514 fail:
515 ralloc_free(cs);
516 return result;
517 }
518
519 static void
radv_device_finish_meta_btoi_r32g32b32_state(struct radv_device * device)520 radv_device_finish_meta_btoi_r32g32b32_state(struct radv_device *device)
521 {
522 struct radv_meta_state *state = &device->meta_state;
523
524 radv_DestroyPipelineLayout(radv_device_to_handle(device), state->btoi_r32g32b32.img_p_layout,
525 &state->alloc);
526 device->vk.dispatch_table.DestroyDescriptorSetLayout(
527 radv_device_to_handle(device), state->btoi_r32g32b32.img_ds_layout, &state->alloc);
528 radv_DestroyPipeline(radv_device_to_handle(device), state->btoi_r32g32b32.pipeline,
529 &state->alloc);
530 }
531
532 static nir_shader *
build_nir_itoi_compute_shader(struct radv_device * dev,bool is_3d,int samples)533 build_nir_itoi_compute_shader(struct radv_device *dev, bool is_3d, int samples)
534 {
535 bool is_multisampled = samples > 1;
536 enum glsl_sampler_dim dim = is_3d ? GLSL_SAMPLER_DIM_3D
537 : is_multisampled ? GLSL_SAMPLER_DIM_MS
538 : GLSL_SAMPLER_DIM_2D;
539 const struct glsl_type *buf_type = glsl_sampler_type(dim, false, false, GLSL_TYPE_FLOAT);
540 const struct glsl_type *img_type = glsl_image_type(dim, false, GLSL_TYPE_FLOAT);
541 nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE,
542 is_3d ? "meta_itoi_cs_3d-%d" : "meta_itoi_cs-%d", samples);
543 b.shader->info.workgroup_size[0] = 8;
544 b.shader->info.workgroup_size[1] = 8;
545 nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, buf_type, "s_tex");
546 input_img->data.descriptor_set = 0;
547 input_img->data.binding = 0;
548
549 nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img");
550 output_img->data.descriptor_set = 0;
551 output_img->data.binding = 1;
552
553 nir_ssa_def *global_id = get_global_ids(&b, is_3d ? 3 : 2);
554
555 nir_ssa_def *src_offset =
556 nir_load_push_constant(&b, is_3d ? 3 : 2, 32, nir_imm_int(&b, 0), .range = is_3d ? 12 : 8);
557 nir_ssa_def *dst_offset =
558 nir_load_push_constant(&b, is_3d ? 3 : 2, 32, nir_imm_int(&b, 12), .range = is_3d ? 24 : 20);
559
560 nir_ssa_def *src_coord = nir_iadd(&b, global_id, src_offset);
561 nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa;
562
563 nir_ssa_def *dst_coord = nir_iadd(&b, global_id, dst_offset);
564
565 nir_tex_instr *tex_instr[8];
566 for (uint32_t i = 0; i < samples; i++) {
567 tex_instr[i] = nir_tex_instr_create(b.shader, is_multisampled ? 4 : 3);
568
569 nir_tex_instr *tex = tex_instr[i];
570 tex->sampler_dim = dim;
571 tex->op = is_multisampled ? nir_texop_txf_ms : nir_texop_txf;
572 tex->src[0].src_type = nir_tex_src_coord;
573 tex->src[0].src = nir_src_for_ssa(nir_trim_vector(&b, src_coord, 2 + is_3d));
574 tex->src[1].src_type = nir_tex_src_lod;
575 tex->src[1].src = nir_src_for_ssa(nir_imm_int(&b, 0));
576 tex->src[2].src_type = nir_tex_src_texture_deref;
577 tex->src[2].src = nir_src_for_ssa(input_img_deref);
578 if (is_multisampled) {
579 tex->src[3].src_type = nir_tex_src_ms_index;
580 tex->src[3].src = nir_src_for_ssa(nir_imm_int(&b, i));
581 }
582 tex->dest_type = nir_type_float32;
583 tex->is_array = false;
584 tex->coord_components = is_3d ? 3 : 2;
585
586 nir_ssa_dest_init(&tex->instr, &tex->dest, 4, 32, "tex");
587 nir_builder_instr_insert(&b, &tex->instr);
588 }
589
590 nir_ssa_def *img_coord = nir_vec4(&b, nir_channel(&b, dst_coord, 0),
591 nir_channel(&b, dst_coord, 1),
592 is_3d ? nir_channel(&b, dst_coord, 2) : nir_ssa_undef(&b, 1, 32),
593 nir_ssa_undef(&b, 1, 32));
594
595 for (uint32_t i = 0; i < samples; i++) {
596 nir_ssa_def *outval = &tex_instr[i]->dest.ssa;
597 nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, img_coord,
598 nir_imm_int(&b, i), outval, nir_imm_int(&b, 0), .image_dim = dim);
599 }
600
601 return b.shader;
602 }
603
604 static VkResult
create_itoi_pipeline(struct radv_device * device,int samples,VkPipeline * pipeline)605 create_itoi_pipeline(struct radv_device *device, int samples, VkPipeline *pipeline)
606 {
607 struct radv_meta_state *state = &device->meta_state;
608 nir_shader *cs = build_nir_itoi_compute_shader(device, false, samples);
609 VkResult result;
610
611 VkPipelineShaderStageCreateInfo pipeline_shader_stage = {
612 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
613 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
614 .module = vk_shader_module_handle_from_nir(cs),
615 .pName = "main",
616 .pSpecializationInfo = NULL,
617 };
618
619 VkComputePipelineCreateInfo vk_pipeline_info = {
620 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
621 .stage = pipeline_shader_stage,
622 .flags = 0,
623 .layout = state->itoi.img_p_layout,
624 };
625
626 result = radv_CreateComputePipelines(radv_device_to_handle(device),
627 radv_pipeline_cache_to_handle(&state->cache), 1,
628 &vk_pipeline_info, NULL, pipeline);
629 ralloc_free(cs);
630 return result;
631 }
632
633 /* image to image - don't write use image accessors */
634 static VkResult
radv_device_init_meta_itoi_state(struct radv_device * device)635 radv_device_init_meta_itoi_state(struct radv_device *device)
636 {
637 VkResult result;
638
639 /*
640 * two descriptors one for the image being sampled
641 * one for the buffer being written.
642 */
643 VkDescriptorSetLayoutCreateInfo ds_create_info = {
644 .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
645 .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,
646 .bindingCount = 2,
647 .pBindings = (VkDescriptorSetLayoutBinding[]){
648 {.binding = 0,
649 .descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,
650 .descriptorCount = 1,
651 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
652 .pImmutableSamplers = NULL},
653 {.binding = 1,
654 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
655 .descriptorCount = 1,
656 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
657 .pImmutableSamplers = NULL},
658 }};
659
660 result = radv_CreateDescriptorSetLayout(radv_device_to_handle(device), &ds_create_info,
661 &device->meta_state.alloc,
662 &device->meta_state.itoi.img_ds_layout);
663 if (result != VK_SUCCESS)
664 goto fail;
665
666 VkPipelineLayoutCreateInfo pl_create_info = {
667 .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,
668 .setLayoutCount = 1,
669 .pSetLayouts = &device->meta_state.itoi.img_ds_layout,
670 .pushConstantRangeCount = 1,
671 .pPushConstantRanges = &(VkPushConstantRange){VK_SHADER_STAGE_COMPUTE_BIT, 0, 24},
672 };
673
674 result =
675 radv_CreatePipelineLayout(radv_device_to_handle(device), &pl_create_info,
676 &device->meta_state.alloc, &device->meta_state.itoi.img_p_layout);
677 if (result != VK_SUCCESS)
678 goto fail;
679
680 for (uint32_t i = 0; i < MAX_SAMPLES_LOG2; i++) {
681 uint32_t samples = 1 << i;
682 result = create_itoi_pipeline(device, samples, &device->meta_state.itoi.pipeline[i]);
683 if (result != VK_SUCCESS)
684 goto fail;
685 }
686
687 nir_shader *cs_3d = build_nir_itoi_compute_shader(device, true, 1);
688
689 VkPipelineShaderStageCreateInfo pipeline_shader_stage_3d = {
690 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
691 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
692 .module = vk_shader_module_handle_from_nir(cs_3d),
693 .pName = "main",
694 .pSpecializationInfo = NULL,
695 };
696
697 VkComputePipelineCreateInfo vk_pipeline_info_3d = {
698 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
699 .stage = pipeline_shader_stage_3d,
700 .flags = 0,
701 .layout = device->meta_state.itoi.img_p_layout,
702 };
703
704 result = radv_CreateComputePipelines(
705 radv_device_to_handle(device), radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,
706 &vk_pipeline_info_3d, NULL, &device->meta_state.itoi.pipeline_3d);
707 ralloc_free(cs_3d);
708
709 return VK_SUCCESS;
710 fail:
711 return result;
712 }
713
714 static void
radv_device_finish_meta_itoi_state(struct radv_device * device)715 radv_device_finish_meta_itoi_state(struct radv_device *device)
716 {
717 struct radv_meta_state *state = &device->meta_state;
718
719 radv_DestroyPipelineLayout(radv_device_to_handle(device), state->itoi.img_p_layout,
720 &state->alloc);
721 device->vk.dispatch_table.DestroyDescriptorSetLayout(radv_device_to_handle(device),
722 state->itoi.img_ds_layout, &state->alloc);
723
724 for (uint32_t i = 0; i < MAX_SAMPLES_LOG2; ++i) {
725 radv_DestroyPipeline(radv_device_to_handle(device), state->itoi.pipeline[i], &state->alloc);
726 }
727
728 radv_DestroyPipeline(radv_device_to_handle(device), state->itoi.pipeline_3d, &state->alloc);
729 }
730
731 static nir_shader *
build_nir_itoi_r32g32b32_compute_shader(struct radv_device * dev)732 build_nir_itoi_r32g32b32_compute_shader(struct radv_device *dev)
733 {
734 const struct glsl_type *type =
735 glsl_sampler_type(GLSL_SAMPLER_DIM_BUF, false, false, GLSL_TYPE_FLOAT);
736 const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_BUF, false, GLSL_TYPE_FLOAT);
737 nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_itoi_r32g32b32_cs");
738 b.shader->info.workgroup_size[0] = 8;
739 b.shader->info.workgroup_size[1] = 8;
740 nir_variable *input_img = nir_variable_create(b.shader, nir_var_uniform, type, "input_img");
741 input_img->data.descriptor_set = 0;
742 input_img->data.binding = 0;
743
744 nir_variable *output_img =
745 nir_variable_create(b.shader, nir_var_image, img_type, "output_img");
746 output_img->data.descriptor_set = 0;
747 output_img->data.binding = 1;
748
749 nir_ssa_def *global_id = get_global_ids(&b, 2);
750
751 nir_ssa_def *src_offset = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 0), .range = 12);
752 nir_ssa_def *dst_offset = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 12), .range = 24);
753
754 nir_ssa_def *src_stride = nir_channel(&b, src_offset, 2);
755 nir_ssa_def *dst_stride = nir_channel(&b, dst_offset, 2);
756
757 nir_ssa_def *src_img_coord = nir_iadd(&b, global_id, src_offset);
758 nir_ssa_def *dst_img_coord = nir_iadd(&b, global_id, dst_offset);
759
760 nir_ssa_def *src_global_pos =
761 nir_iadd(&b, nir_imul(&b, nir_channel(&b, src_img_coord, 1), src_stride),
762 nir_imul_imm(&b, nir_channel(&b, src_img_coord, 0), 3));
763
764 nir_ssa_def *dst_global_pos =
765 nir_iadd(&b, nir_imul(&b, nir_channel(&b, dst_img_coord, 1), dst_stride),
766 nir_imul_imm(&b, nir_channel(&b, dst_img_coord, 0), 3));
767
768 for (int chan = 0; chan < 3; chan++) {
769 /* src */
770 nir_ssa_def *src_local_pos = nir_iadd_imm(&b, src_global_pos, chan);
771 nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa;
772
773 nir_tex_instr *tex = nir_tex_instr_create(b.shader, 3);
774 tex->sampler_dim = GLSL_SAMPLER_DIM_BUF;
775 tex->op = nir_texop_txf;
776 tex->src[0].src_type = nir_tex_src_coord;
777 tex->src[0].src = nir_src_for_ssa(src_local_pos);
778 tex->src[1].src_type = nir_tex_src_lod;
779 tex->src[1].src = nir_src_for_ssa(nir_imm_int(&b, 0));
780 tex->src[2].src_type = nir_tex_src_texture_deref;
781 tex->src[2].src = nir_src_for_ssa(input_img_deref);
782 tex->dest_type = nir_type_float32;
783 tex->is_array = false;
784 tex->coord_components = 1;
785 nir_ssa_dest_init(&tex->instr, &tex->dest, 4, 32, "tex");
786 nir_builder_instr_insert(&b, &tex->instr);
787
788 nir_ssa_def *outval = &tex->dest.ssa;
789
790 /* dst */
791 nir_ssa_def *dst_local_pos = nir_iadd_imm(&b, dst_global_pos, chan);
792
793 nir_ssa_def *dst_coord =
794 nir_vec4(&b, dst_local_pos, dst_local_pos, dst_local_pos, dst_local_pos);
795
796 nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, dst_coord,
797 nir_ssa_undef(&b, 1, 32), nir_channel(&b, outval, 0),
798 nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_BUF);
799 }
800
801 return b.shader;
802 }
803
804 /* Image to image - special path for R32G32B32 */
805 static VkResult
radv_device_init_meta_itoi_r32g32b32_state(struct radv_device * device)806 radv_device_init_meta_itoi_r32g32b32_state(struct radv_device *device)
807 {
808 VkResult result;
809 nir_shader *cs = build_nir_itoi_r32g32b32_compute_shader(device);
810
811 VkDescriptorSetLayoutCreateInfo ds_create_info = {
812 .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
813 .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,
814 .bindingCount = 2,
815 .pBindings = (VkDescriptorSetLayoutBinding[]){
816 {.binding = 0,
817 .descriptorType = VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER,
818 .descriptorCount = 1,
819 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
820 .pImmutableSamplers = NULL},
821 {.binding = 1,
822 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
823 .descriptorCount = 1,
824 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
825 .pImmutableSamplers = NULL},
826 }};
827
828 result = radv_CreateDescriptorSetLayout(radv_device_to_handle(device), &ds_create_info,
829 &device->meta_state.alloc,
830 &device->meta_state.itoi_r32g32b32.img_ds_layout);
831 if (result != VK_SUCCESS)
832 goto fail;
833
834 VkPipelineLayoutCreateInfo pl_create_info = {
835 .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,
836 .setLayoutCount = 1,
837 .pSetLayouts = &device->meta_state.itoi_r32g32b32.img_ds_layout,
838 .pushConstantRangeCount = 1,
839 .pPushConstantRanges = &(VkPushConstantRange){VK_SHADER_STAGE_COMPUTE_BIT, 0, 24},
840 };
841
842 result = radv_CreatePipelineLayout(radv_device_to_handle(device), &pl_create_info,
843 &device->meta_state.alloc,
844 &device->meta_state.itoi_r32g32b32.img_p_layout);
845 if (result != VK_SUCCESS)
846 goto fail;
847
848 /* compute shader */
849
850 VkPipelineShaderStageCreateInfo pipeline_shader_stage = {
851 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
852 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
853 .module = vk_shader_module_handle_from_nir(cs),
854 .pName = "main",
855 .pSpecializationInfo = NULL,
856 };
857
858 VkComputePipelineCreateInfo vk_pipeline_info = {
859 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
860 .stage = pipeline_shader_stage,
861 .flags = 0,
862 .layout = device->meta_state.itoi_r32g32b32.img_p_layout,
863 };
864
865 result = radv_CreateComputePipelines(
866 radv_device_to_handle(device), radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,
867 &vk_pipeline_info, NULL, &device->meta_state.itoi_r32g32b32.pipeline);
868
869 fail:
870 ralloc_free(cs);
871 return result;
872 }
873
874 static void
radv_device_finish_meta_itoi_r32g32b32_state(struct radv_device * device)875 radv_device_finish_meta_itoi_r32g32b32_state(struct radv_device *device)
876 {
877 struct radv_meta_state *state = &device->meta_state;
878
879 radv_DestroyPipelineLayout(radv_device_to_handle(device), state->itoi_r32g32b32.img_p_layout,
880 &state->alloc);
881 device->vk.dispatch_table.DestroyDescriptorSetLayout(
882 radv_device_to_handle(device), state->itoi_r32g32b32.img_ds_layout, &state->alloc);
883 radv_DestroyPipeline(radv_device_to_handle(device), state->itoi_r32g32b32.pipeline,
884 &state->alloc);
885 }
886
887 static nir_shader *
build_nir_cleari_compute_shader(struct radv_device * dev,bool is_3d,int samples)888 build_nir_cleari_compute_shader(struct radv_device *dev, bool is_3d, int samples)
889 {
890 bool is_multisampled = samples > 1;
891 enum glsl_sampler_dim dim = is_3d ? GLSL_SAMPLER_DIM_3D
892 : is_multisampled ? GLSL_SAMPLER_DIM_MS
893 : GLSL_SAMPLER_DIM_2D;
894 const struct glsl_type *img_type = glsl_image_type(dim, false, GLSL_TYPE_FLOAT);
895 nir_builder b = radv_meta_init_shader(
896 dev, MESA_SHADER_COMPUTE, is_3d ? "meta_cleari_cs_3d-%d" : "meta_cleari_cs-%d", samples);
897 b.shader->info.workgroup_size[0] = 8;
898 b.shader->info.workgroup_size[1] = 8;
899
900 nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img");
901 output_img->data.descriptor_set = 0;
902 output_img->data.binding = 0;
903
904 nir_ssa_def *global_id = get_global_ids(&b, 2);
905
906 nir_ssa_def *clear_val = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 16);
907 nir_ssa_def *layer = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 16), .range = 20);
908
909 nir_ssa_def *comps[4];
910 comps[0] = nir_channel(&b, global_id, 0);
911 comps[1] = nir_channel(&b, global_id, 1);
912 comps[2] = layer;
913 comps[3] = nir_ssa_undef(&b, 1, 32);
914 global_id = nir_vec(&b, comps, 4);
915
916 for (uint32_t i = 0; i < samples; i++) {
917 nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, global_id,
918 nir_imm_int(&b, i), clear_val, nir_imm_int(&b, 0), .image_dim = dim);
919 }
920
921 return b.shader;
922 }
923
924 static VkResult
create_cleari_pipeline(struct radv_device * device,int samples,VkPipeline * pipeline)925 create_cleari_pipeline(struct radv_device *device, int samples, VkPipeline *pipeline)
926 {
927 nir_shader *cs = build_nir_cleari_compute_shader(device, false, samples);
928 VkResult result;
929
930 VkPipelineShaderStageCreateInfo pipeline_shader_stage = {
931 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
932 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
933 .module = vk_shader_module_handle_from_nir(cs),
934 .pName = "main",
935 .pSpecializationInfo = NULL,
936 };
937
938 VkComputePipelineCreateInfo vk_pipeline_info = {
939 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
940 .stage = pipeline_shader_stage,
941 .flags = 0,
942 .layout = device->meta_state.cleari.img_p_layout,
943 };
944
945 result = radv_CreateComputePipelines(radv_device_to_handle(device),
946 radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,
947 &vk_pipeline_info, NULL, pipeline);
948 ralloc_free(cs);
949 return result;
950 }
951
952 static VkResult
radv_device_init_meta_cleari_state(struct radv_device * device)953 radv_device_init_meta_cleari_state(struct radv_device *device)
954 {
955 VkResult result;
956
957 /*
958 * two descriptors one for the image being sampled
959 * one for the buffer being written.
960 */
961 VkDescriptorSetLayoutCreateInfo ds_create_info = {
962 .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
963 .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,
964 .bindingCount = 1,
965 .pBindings = (VkDescriptorSetLayoutBinding[]){
966 {.binding = 0,
967 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
968 .descriptorCount = 1,
969 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
970 .pImmutableSamplers = NULL},
971 }};
972
973 result = radv_CreateDescriptorSetLayout(radv_device_to_handle(device), &ds_create_info,
974 &device->meta_state.alloc,
975 &device->meta_state.cleari.img_ds_layout);
976 if (result != VK_SUCCESS)
977 goto fail;
978
979 VkPipelineLayoutCreateInfo pl_create_info = {
980 .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,
981 .setLayoutCount = 1,
982 .pSetLayouts = &device->meta_state.cleari.img_ds_layout,
983 .pushConstantRangeCount = 1,
984 .pPushConstantRanges = &(VkPushConstantRange){VK_SHADER_STAGE_COMPUTE_BIT, 0, 20},
985 };
986
987 result =
988 radv_CreatePipelineLayout(radv_device_to_handle(device), &pl_create_info,
989 &device->meta_state.alloc, &device->meta_state.cleari.img_p_layout);
990 if (result != VK_SUCCESS)
991 goto fail;
992
993 for (uint32_t i = 0; i < MAX_SAMPLES_LOG2; i++) {
994 uint32_t samples = 1 << i;
995 result = create_cleari_pipeline(device, samples, &device->meta_state.cleari.pipeline[i]);
996 if (result != VK_SUCCESS)
997 goto fail;
998 }
999
1000 nir_shader *cs_3d = build_nir_cleari_compute_shader(device, true, 1);
1001
1002 /* compute shader */
1003 VkPipelineShaderStageCreateInfo pipeline_shader_stage_3d = {
1004 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
1005 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
1006 .module = vk_shader_module_handle_from_nir(cs_3d),
1007 .pName = "main",
1008 .pSpecializationInfo = NULL,
1009 };
1010
1011 VkComputePipelineCreateInfo vk_pipeline_info_3d = {
1012 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
1013 .stage = pipeline_shader_stage_3d,
1014 .flags = 0,
1015 .layout = device->meta_state.cleari.img_p_layout,
1016 };
1017
1018 result = radv_CreateComputePipelines(
1019 radv_device_to_handle(device), radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,
1020 &vk_pipeline_info_3d, NULL, &device->meta_state.cleari.pipeline_3d);
1021 ralloc_free(cs_3d);
1022
1023 return VK_SUCCESS;
1024 fail:
1025 return result;
1026 }
1027
1028 static void
radv_device_finish_meta_cleari_state(struct radv_device * device)1029 radv_device_finish_meta_cleari_state(struct radv_device *device)
1030 {
1031 struct radv_meta_state *state = &device->meta_state;
1032
1033 radv_DestroyPipelineLayout(radv_device_to_handle(device), state->cleari.img_p_layout,
1034 &state->alloc);
1035 device->vk.dispatch_table.DestroyDescriptorSetLayout(radv_device_to_handle(device),
1036 state->cleari.img_ds_layout, &state->alloc);
1037
1038 for (uint32_t i = 0; i < MAX_SAMPLES_LOG2; ++i) {
1039 radv_DestroyPipeline(radv_device_to_handle(device), state->cleari.pipeline[i], &state->alloc);
1040 }
1041
1042 radv_DestroyPipeline(radv_device_to_handle(device), state->cleari.pipeline_3d, &state->alloc);
1043 }
1044
1045 /* Special path for clearing R32G32B32 images using a compute shader. */
1046 static nir_shader *
build_nir_cleari_r32g32b32_compute_shader(struct radv_device * dev)1047 build_nir_cleari_r32g32b32_compute_shader(struct radv_device *dev)
1048 {
1049 const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_BUF, false, GLSL_TYPE_FLOAT);
1050 nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_cleari_r32g32b32_cs");
1051 b.shader->info.workgroup_size[0] = 8;
1052 b.shader->info.workgroup_size[1] = 8;
1053
1054 nir_variable *output_img = nir_variable_create(b.shader, nir_var_image, img_type, "out_img");
1055 output_img->data.descriptor_set = 0;
1056 output_img->data.binding = 0;
1057
1058 nir_ssa_def *global_id = get_global_ids(&b, 2);
1059
1060 nir_ssa_def *clear_val = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 0), .range = 12);
1061 nir_ssa_def *stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 12), .range = 16);
1062
1063 nir_ssa_def *global_x = nir_channel(&b, global_id, 0);
1064 nir_ssa_def *global_y = nir_channel(&b, global_id, 1);
1065
1066 nir_ssa_def *global_pos =
1067 nir_iadd(&b, nir_imul(&b, global_y, stride), nir_imul_imm(&b, global_x, 3));
1068
1069 for (unsigned chan = 0; chan < 3; chan++) {
1070 nir_ssa_def *local_pos = nir_iadd_imm(&b, global_pos, chan);
1071
1072 nir_ssa_def *coord = nir_vec4(&b, local_pos, local_pos, local_pos, local_pos);
1073
1074 nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img)->dest.ssa, coord,
1075 nir_ssa_undef(&b, 1, 32), nir_channel(&b, clear_val, chan),
1076 nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_BUF);
1077 }
1078
1079 return b.shader;
1080 }
1081
1082 static VkResult
radv_device_init_meta_cleari_r32g32b32_state(struct radv_device * device)1083 radv_device_init_meta_cleari_r32g32b32_state(struct radv_device *device)
1084 {
1085 VkResult result;
1086 nir_shader *cs = build_nir_cleari_r32g32b32_compute_shader(device);
1087
1088 VkDescriptorSetLayoutCreateInfo ds_create_info = {
1089 .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
1090 .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,
1091 .bindingCount = 1,
1092 .pBindings = (VkDescriptorSetLayoutBinding[]){
1093 {.binding = 0,
1094 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
1095 .descriptorCount = 1,
1096 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
1097 .pImmutableSamplers = NULL},
1098 }};
1099
1100 result = radv_CreateDescriptorSetLayout(radv_device_to_handle(device), &ds_create_info,
1101 &device->meta_state.alloc,
1102 &device->meta_state.cleari_r32g32b32.img_ds_layout);
1103 if (result != VK_SUCCESS)
1104 goto fail;
1105
1106 VkPipelineLayoutCreateInfo pl_create_info = {
1107 .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,
1108 .setLayoutCount = 1,
1109 .pSetLayouts = &device->meta_state.cleari_r32g32b32.img_ds_layout,
1110 .pushConstantRangeCount = 1,
1111 .pPushConstantRanges = &(VkPushConstantRange){VK_SHADER_STAGE_COMPUTE_BIT, 0, 16},
1112 };
1113
1114 result = radv_CreatePipelineLayout(radv_device_to_handle(device), &pl_create_info,
1115 &device->meta_state.alloc,
1116 &device->meta_state.cleari_r32g32b32.img_p_layout);
1117 if (result != VK_SUCCESS)
1118 goto fail;
1119
1120 /* compute shader */
1121 VkPipelineShaderStageCreateInfo pipeline_shader_stage = {
1122 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
1123 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
1124 .module = vk_shader_module_handle_from_nir(cs),
1125 .pName = "main",
1126 .pSpecializationInfo = NULL,
1127 };
1128
1129 VkComputePipelineCreateInfo vk_pipeline_info = {
1130 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
1131 .stage = pipeline_shader_stage,
1132 .flags = 0,
1133 .layout = device->meta_state.cleari_r32g32b32.img_p_layout,
1134 };
1135
1136 result = radv_CreateComputePipelines(
1137 radv_device_to_handle(device), radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,
1138 &vk_pipeline_info, NULL, &device->meta_state.cleari_r32g32b32.pipeline);
1139
1140 fail:
1141 ralloc_free(cs);
1142 return result;
1143 }
1144
1145 static void
radv_device_finish_meta_cleari_r32g32b32_state(struct radv_device * device)1146 radv_device_finish_meta_cleari_r32g32b32_state(struct radv_device *device)
1147 {
1148 struct radv_meta_state *state = &device->meta_state;
1149
1150 radv_DestroyPipelineLayout(radv_device_to_handle(device), state->cleari_r32g32b32.img_p_layout,
1151 &state->alloc);
1152 device->vk.dispatch_table.DestroyDescriptorSetLayout(
1153 radv_device_to_handle(device), state->cleari_r32g32b32.img_ds_layout, &state->alloc);
1154 radv_DestroyPipeline(radv_device_to_handle(device), state->cleari_r32g32b32.pipeline,
1155 &state->alloc);
1156 }
1157
1158 void
radv_device_finish_meta_bufimage_state(struct radv_device * device)1159 radv_device_finish_meta_bufimage_state(struct radv_device *device)
1160 {
1161 radv_device_finish_meta_itob_state(device);
1162 radv_device_finish_meta_btoi_state(device);
1163 radv_device_finish_meta_btoi_r32g32b32_state(device);
1164 radv_device_finish_meta_itoi_state(device);
1165 radv_device_finish_meta_itoi_r32g32b32_state(device);
1166 radv_device_finish_meta_cleari_state(device);
1167 radv_device_finish_meta_cleari_r32g32b32_state(device);
1168 }
1169
1170 VkResult
radv_device_init_meta_bufimage_state(struct radv_device * device)1171 radv_device_init_meta_bufimage_state(struct radv_device *device)
1172 {
1173 VkResult result;
1174
1175 result = radv_device_init_meta_itob_state(device);
1176 if (result != VK_SUCCESS)
1177 return result;
1178
1179 result = radv_device_init_meta_btoi_state(device);
1180 if (result != VK_SUCCESS)
1181 return result;
1182
1183 result = radv_device_init_meta_btoi_r32g32b32_state(device);
1184 if (result != VK_SUCCESS)
1185 return result;
1186
1187 result = radv_device_init_meta_itoi_state(device);
1188 if (result != VK_SUCCESS)
1189 return result;
1190
1191 result = radv_device_init_meta_itoi_r32g32b32_state(device);
1192 if (result != VK_SUCCESS)
1193 return result;
1194
1195 result = radv_device_init_meta_cleari_state(device);
1196 if (result != VK_SUCCESS)
1197 return result;
1198
1199 result = radv_device_init_meta_cleari_r32g32b32_state(device);
1200 if (result != VK_SUCCESS)
1201 return result;
1202
1203 return VK_SUCCESS;
1204 }
1205
1206 static void
create_iview(struct radv_cmd_buffer * cmd_buffer,struct radv_meta_blit2d_surf * surf,struct radv_image_view * iview,VkFormat format,VkImageAspectFlagBits aspects)1207 create_iview(struct radv_cmd_buffer *cmd_buffer, struct radv_meta_blit2d_surf *surf,
1208 struct radv_image_view *iview, VkFormat format, VkImageAspectFlagBits aspects)
1209 {
1210 if (format == VK_FORMAT_UNDEFINED)
1211 format = surf->format;
1212
1213 radv_image_view_init(iview, cmd_buffer->device,
1214 &(VkImageViewCreateInfo){
1215 .sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO,
1216 .image = radv_image_to_handle(surf->image),
1217 .viewType = radv_meta_get_view_type(surf->image),
1218 .format = format,
1219 .subresourceRange = {.aspectMask = aspects,
1220 .baseMipLevel = surf->level,
1221 .levelCount = 1,
1222 .baseArrayLayer = surf->layer,
1223 .layerCount = 1},
1224 },
1225 0, &(struct radv_image_view_extra_create_info){
1226 .disable_compression = surf->disable_compression,
1227 });
1228 }
1229
1230 static void
create_bview(struct radv_cmd_buffer * cmd_buffer,struct radv_buffer * buffer,unsigned offset,VkFormat format,struct radv_buffer_view * bview)1231 create_bview(struct radv_cmd_buffer *cmd_buffer, struct radv_buffer *buffer, unsigned offset,
1232 VkFormat format, struct radv_buffer_view *bview)
1233 {
1234 radv_buffer_view_init(bview, cmd_buffer->device,
1235 &(VkBufferViewCreateInfo){
1236 .sType = VK_STRUCTURE_TYPE_BUFFER_VIEW_CREATE_INFO,
1237 .flags = 0,
1238 .buffer = radv_buffer_to_handle(buffer),
1239 .format = format,
1240 .offset = offset,
1241 .range = VK_WHOLE_SIZE,
1242 });
1243 }
1244
1245 static void
create_buffer_from_image(struct radv_cmd_buffer * cmd_buffer,struct radv_meta_blit2d_surf * surf,VkBufferUsageFlagBits usage,VkBuffer * buffer)1246 create_buffer_from_image(struct radv_cmd_buffer *cmd_buffer, struct radv_meta_blit2d_surf *surf,
1247 VkBufferUsageFlagBits usage, VkBuffer *buffer)
1248 {
1249 struct radv_device *device = cmd_buffer->device;
1250 struct radv_device_memory mem;
1251
1252 radv_device_memory_init(&mem, device, surf->image->bindings[0].bo);
1253
1254 radv_CreateBuffer(radv_device_to_handle(device),
1255 &(VkBufferCreateInfo){
1256 .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO,
1257 .flags = 0,
1258 .size = surf->image->size,
1259 .usage = usage,
1260 .sharingMode = VK_SHARING_MODE_EXCLUSIVE,
1261 },
1262 NULL, buffer);
1263
1264 radv_BindBufferMemory2(radv_device_to_handle(device), 1,
1265 (VkBindBufferMemoryInfo[]){{
1266 .sType = VK_STRUCTURE_TYPE_BIND_BUFFER_MEMORY_INFO,
1267 .buffer = *buffer,
1268 .memory = radv_device_memory_to_handle(&mem),
1269 .memoryOffset = surf->image->bindings[0].offset,
1270 }});
1271
1272 radv_device_memory_finish(&mem);
1273 }
1274
1275 static void
create_bview_for_r32g32b32(struct radv_cmd_buffer * cmd_buffer,struct radv_buffer * buffer,unsigned offset,VkFormat src_format,struct radv_buffer_view * bview)1276 create_bview_for_r32g32b32(struct radv_cmd_buffer *cmd_buffer, struct radv_buffer *buffer,
1277 unsigned offset, VkFormat src_format, struct radv_buffer_view *bview)
1278 {
1279 VkFormat format;
1280
1281 switch (src_format) {
1282 case VK_FORMAT_R32G32B32_UINT:
1283 format = VK_FORMAT_R32_UINT;
1284 break;
1285 case VK_FORMAT_R32G32B32_SINT:
1286 format = VK_FORMAT_R32_SINT;
1287 break;
1288 case VK_FORMAT_R32G32B32_SFLOAT:
1289 format = VK_FORMAT_R32_SFLOAT;
1290 break;
1291 default:
1292 unreachable("invalid R32G32B32 format");
1293 }
1294
1295 radv_buffer_view_init(bview, cmd_buffer->device,
1296 &(VkBufferViewCreateInfo){
1297 .sType = VK_STRUCTURE_TYPE_BUFFER_VIEW_CREATE_INFO,
1298 .flags = 0,
1299 .buffer = radv_buffer_to_handle(buffer),
1300 .format = format,
1301 .offset = offset,
1302 .range = VK_WHOLE_SIZE,
1303 });
1304 }
1305
1306 static unsigned
get_image_stride_for_r32g32b32(struct radv_cmd_buffer * cmd_buffer,struct radv_meta_blit2d_surf * surf)1307 get_image_stride_for_r32g32b32(struct radv_cmd_buffer *cmd_buffer,
1308 struct radv_meta_blit2d_surf *surf)
1309 {
1310 unsigned stride;
1311
1312 if (cmd_buffer->device->physical_device->rad_info.gfx_level >= GFX9) {
1313 stride = surf->image->planes[0].surface.u.gfx9.surf_pitch;
1314 } else {
1315 stride = surf->image->planes[0].surface.u.legacy.level[0].nblk_x * 3;
1316 }
1317
1318 return stride;
1319 }
1320
1321 static void
itob_bind_descriptors(struct radv_cmd_buffer * cmd_buffer,struct radv_image_view * src,struct radv_buffer_view * dst)1322 itob_bind_descriptors(struct radv_cmd_buffer *cmd_buffer, struct radv_image_view *src,
1323 struct radv_buffer_view *dst)
1324 {
1325 struct radv_device *device = cmd_buffer->device;
1326
1327 radv_meta_push_descriptor_set(
1328 cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, device->meta_state.itob.img_p_layout, 0, /* set */
1329 2, /* descriptorWriteCount */
1330 (VkWriteDescriptorSet[]){
1331 {.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1332 .dstBinding = 0,
1333 .dstArrayElement = 0,
1334 .descriptorCount = 1,
1335 .descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,
1336 .pImageInfo =
1337 (VkDescriptorImageInfo[]){
1338 {
1339 .sampler = VK_NULL_HANDLE,
1340 .imageView = radv_image_view_to_handle(src),
1341 .imageLayout = VK_IMAGE_LAYOUT_GENERAL,
1342 },
1343 }},
1344 {
1345 .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1346 .dstBinding = 1,
1347 .dstArrayElement = 0,
1348 .descriptorCount = 1,
1349 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
1350 .pTexelBufferView = (VkBufferView[]){radv_buffer_view_to_handle(dst)},
1351 }});
1352 }
1353
1354 void
radv_meta_image_to_buffer(struct radv_cmd_buffer * cmd_buffer,struct radv_meta_blit2d_surf * src,struct radv_meta_blit2d_buffer * dst,unsigned num_rects,struct radv_meta_blit2d_rect * rects)1355 radv_meta_image_to_buffer(struct radv_cmd_buffer *cmd_buffer, struct radv_meta_blit2d_surf *src,
1356 struct radv_meta_blit2d_buffer *dst, unsigned num_rects,
1357 struct radv_meta_blit2d_rect *rects)
1358 {
1359 VkPipeline pipeline = cmd_buffer->device->meta_state.itob.pipeline;
1360 struct radv_device *device = cmd_buffer->device;
1361 struct radv_image_view src_view;
1362 struct radv_buffer_view dst_view;
1363
1364 create_iview(cmd_buffer, src, &src_view, VK_FORMAT_UNDEFINED, src->aspect_mask);
1365 create_bview(cmd_buffer, dst->buffer, dst->offset, dst->format, &dst_view);
1366 itob_bind_descriptors(cmd_buffer, &src_view, &dst_view);
1367
1368 if (src->image->vk.image_type == VK_IMAGE_TYPE_3D)
1369 pipeline = cmd_buffer->device->meta_state.itob.pipeline_3d;
1370
1371 radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE,
1372 pipeline);
1373
1374 for (unsigned r = 0; r < num_rects; ++r) {
1375 unsigned push_constants[4] = {rects[r].src_x, rects[r].src_y, src->layer, dst->pitch};
1376 radv_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer),
1377 device->meta_state.itob.img_p_layout, VK_SHADER_STAGE_COMPUTE_BIT, 0,
1378 16, push_constants);
1379
1380 radv_unaligned_dispatch(cmd_buffer, rects[r].width, rects[r].height, 1);
1381 }
1382
1383 radv_image_view_finish(&src_view);
1384 radv_buffer_view_finish(&dst_view);
1385 }
1386
1387 static void
btoi_r32g32b32_bind_descriptors(struct radv_cmd_buffer * cmd_buffer,struct radv_buffer_view * src,struct radv_buffer_view * dst)1388 btoi_r32g32b32_bind_descriptors(struct radv_cmd_buffer *cmd_buffer, struct radv_buffer_view *src,
1389 struct radv_buffer_view *dst)
1390 {
1391 struct radv_device *device = cmd_buffer->device;
1392
1393 radv_meta_push_descriptor_set(
1394 cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, device->meta_state.btoi_r32g32b32.img_p_layout,
1395 0, /* set */
1396 2, /* descriptorWriteCount */
1397 (VkWriteDescriptorSet[]){
1398 {
1399 .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1400 .dstBinding = 0,
1401 .dstArrayElement = 0,
1402 .descriptorCount = 1,
1403 .descriptorType = VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER,
1404 .pTexelBufferView = (VkBufferView[]){radv_buffer_view_to_handle(src)},
1405 },
1406 {
1407 .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1408 .dstBinding = 1,
1409 .dstArrayElement = 0,
1410 .descriptorCount = 1,
1411 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
1412 .pTexelBufferView = (VkBufferView[]){radv_buffer_view_to_handle(dst)},
1413 }});
1414 }
1415
1416 static void
radv_meta_buffer_to_image_cs_r32g32b32(struct radv_cmd_buffer * cmd_buffer,struct radv_meta_blit2d_buffer * src,struct radv_meta_blit2d_surf * dst,unsigned num_rects,struct radv_meta_blit2d_rect * rects)1417 radv_meta_buffer_to_image_cs_r32g32b32(struct radv_cmd_buffer *cmd_buffer,
1418 struct radv_meta_blit2d_buffer *src,
1419 struct radv_meta_blit2d_surf *dst, unsigned num_rects,
1420 struct radv_meta_blit2d_rect *rects)
1421 {
1422 VkPipeline pipeline = cmd_buffer->device->meta_state.btoi_r32g32b32.pipeline;
1423 struct radv_device *device = cmd_buffer->device;
1424 struct radv_buffer_view src_view, dst_view;
1425 unsigned dst_offset = 0;
1426 unsigned stride;
1427 VkBuffer buffer;
1428
1429 /* This special btoi path for R32G32B32 formats will write the linear
1430 * image as a buffer with the same underlying memory. The compute
1431 * shader will copy all components separately using a R32 format.
1432 */
1433 create_buffer_from_image(cmd_buffer, dst, VK_BUFFER_USAGE_STORAGE_TEXEL_BUFFER_BIT, &buffer);
1434
1435 create_bview(cmd_buffer, src->buffer, src->offset, src->format, &src_view);
1436 create_bview_for_r32g32b32(cmd_buffer, radv_buffer_from_handle(buffer), dst_offset, dst->format,
1437 &dst_view);
1438 btoi_r32g32b32_bind_descriptors(cmd_buffer, &src_view, &dst_view);
1439
1440 radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE,
1441 pipeline);
1442
1443 stride = get_image_stride_for_r32g32b32(cmd_buffer, dst);
1444
1445 for (unsigned r = 0; r < num_rects; ++r) {
1446 unsigned push_constants[4] = {
1447 rects[r].dst_x,
1448 rects[r].dst_y,
1449 stride,
1450 src->pitch,
1451 };
1452
1453 radv_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer),
1454 device->meta_state.btoi_r32g32b32.img_p_layout,
1455 VK_SHADER_STAGE_COMPUTE_BIT, 0, 16, push_constants);
1456
1457 radv_unaligned_dispatch(cmd_buffer, rects[r].width, rects[r].height, 1);
1458 }
1459
1460 radv_buffer_view_finish(&src_view);
1461 radv_buffer_view_finish(&dst_view);
1462 radv_DestroyBuffer(radv_device_to_handle(device), buffer, NULL);
1463 }
1464
1465 static void
btoi_bind_descriptors(struct radv_cmd_buffer * cmd_buffer,struct radv_buffer_view * src,struct radv_image_view * dst)1466 btoi_bind_descriptors(struct radv_cmd_buffer *cmd_buffer, struct radv_buffer_view *src,
1467 struct radv_image_view *dst)
1468 {
1469 struct radv_device *device = cmd_buffer->device;
1470
1471 radv_meta_push_descriptor_set(
1472 cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, device->meta_state.btoi.img_p_layout, 0, /* set */
1473 2, /* descriptorWriteCount */
1474 (VkWriteDescriptorSet[]){
1475 {
1476 .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1477 .dstBinding = 0,
1478 .dstArrayElement = 0,
1479 .descriptorCount = 1,
1480 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
1481 .pTexelBufferView = (VkBufferView[]){radv_buffer_view_to_handle(src)},
1482 },
1483 {.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1484 .dstBinding = 1,
1485 .dstArrayElement = 0,
1486 .descriptorCount = 1,
1487 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
1488 .pImageInfo = (VkDescriptorImageInfo[]){
1489 {
1490 .sampler = VK_NULL_HANDLE,
1491 .imageView = radv_image_view_to_handle(dst),
1492 .imageLayout = VK_IMAGE_LAYOUT_GENERAL,
1493 },
1494 }}});
1495 }
1496
1497 void
radv_meta_buffer_to_image_cs(struct radv_cmd_buffer * cmd_buffer,struct radv_meta_blit2d_buffer * src,struct radv_meta_blit2d_surf * dst,unsigned num_rects,struct radv_meta_blit2d_rect * rects)1498 radv_meta_buffer_to_image_cs(struct radv_cmd_buffer *cmd_buffer,
1499 struct radv_meta_blit2d_buffer *src, struct radv_meta_blit2d_surf *dst,
1500 unsigned num_rects, struct radv_meta_blit2d_rect *rects)
1501 {
1502 VkPipeline pipeline = cmd_buffer->device->meta_state.btoi.pipeline;
1503 struct radv_device *device = cmd_buffer->device;
1504 struct radv_buffer_view src_view;
1505 struct radv_image_view dst_view;
1506
1507 if (dst->image->vk.format == VK_FORMAT_R32G32B32_UINT ||
1508 dst->image->vk.format == VK_FORMAT_R32G32B32_SINT ||
1509 dst->image->vk.format == VK_FORMAT_R32G32B32_SFLOAT) {
1510 radv_meta_buffer_to_image_cs_r32g32b32(cmd_buffer, src, dst, num_rects, rects);
1511 return;
1512 }
1513
1514 create_bview(cmd_buffer, src->buffer, src->offset, src->format, &src_view);
1515 create_iview(cmd_buffer, dst, &dst_view, VK_FORMAT_UNDEFINED, dst->aspect_mask);
1516 btoi_bind_descriptors(cmd_buffer, &src_view, &dst_view);
1517
1518 if (dst->image->vk.image_type == VK_IMAGE_TYPE_3D)
1519 pipeline = cmd_buffer->device->meta_state.btoi.pipeline_3d;
1520 radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE,
1521 pipeline);
1522
1523 for (unsigned r = 0; r < num_rects; ++r) {
1524 unsigned push_constants[4] = {
1525 rects[r].dst_x,
1526 rects[r].dst_y,
1527 dst->layer,
1528 src->pitch,
1529 };
1530 radv_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer),
1531 device->meta_state.btoi.img_p_layout, VK_SHADER_STAGE_COMPUTE_BIT, 0,
1532 16, push_constants);
1533
1534 radv_unaligned_dispatch(cmd_buffer, rects[r].width, rects[r].height, 1);
1535 }
1536
1537 radv_image_view_finish(&dst_view);
1538 radv_buffer_view_finish(&src_view);
1539 }
1540
1541 static void
itoi_r32g32b32_bind_descriptors(struct radv_cmd_buffer * cmd_buffer,struct radv_buffer_view * src,struct radv_buffer_view * dst)1542 itoi_r32g32b32_bind_descriptors(struct radv_cmd_buffer *cmd_buffer, struct radv_buffer_view *src,
1543 struct radv_buffer_view *dst)
1544 {
1545 struct radv_device *device = cmd_buffer->device;
1546
1547 radv_meta_push_descriptor_set(
1548 cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, device->meta_state.itoi_r32g32b32.img_p_layout,
1549 0, /* set */
1550 2, /* descriptorWriteCount */
1551 (VkWriteDescriptorSet[]){
1552 {
1553 .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1554 .dstBinding = 0,
1555 .dstArrayElement = 0,
1556 .descriptorCount = 1,
1557 .descriptorType = VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER,
1558 .pTexelBufferView = (VkBufferView[]){radv_buffer_view_to_handle(src)},
1559 },
1560 {
1561 .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1562 .dstBinding = 1,
1563 .dstArrayElement = 0,
1564 .descriptorCount = 1,
1565 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
1566 .pTexelBufferView = (VkBufferView[]){radv_buffer_view_to_handle(dst)},
1567 }});
1568 }
1569
1570 static void
radv_meta_image_to_image_cs_r32g32b32(struct radv_cmd_buffer * cmd_buffer,struct radv_meta_blit2d_surf * src,struct radv_meta_blit2d_surf * dst,unsigned num_rects,struct radv_meta_blit2d_rect * rects)1571 radv_meta_image_to_image_cs_r32g32b32(struct radv_cmd_buffer *cmd_buffer,
1572 struct radv_meta_blit2d_surf *src,
1573 struct radv_meta_blit2d_surf *dst, unsigned num_rects,
1574 struct radv_meta_blit2d_rect *rects)
1575 {
1576 VkPipeline pipeline = cmd_buffer->device->meta_state.itoi_r32g32b32.pipeline;
1577 struct radv_device *device = cmd_buffer->device;
1578 struct radv_buffer_view src_view, dst_view;
1579 unsigned src_offset = 0, dst_offset = 0;
1580 unsigned src_stride, dst_stride;
1581 VkBuffer src_buffer, dst_buffer;
1582
1583 /* 96-bit formats are only compatible to themselves. */
1584 assert(dst->format == VK_FORMAT_R32G32B32_UINT || dst->format == VK_FORMAT_R32G32B32_SINT ||
1585 dst->format == VK_FORMAT_R32G32B32_SFLOAT);
1586
1587 /* This special itoi path for R32G32B32 formats will write the linear
1588 * image as a buffer with the same underlying memory. The compute
1589 * shader will copy all components separately using a R32 format.
1590 */
1591 create_buffer_from_image(cmd_buffer, src, VK_BUFFER_USAGE_UNIFORM_TEXEL_BUFFER_BIT, &src_buffer);
1592 create_buffer_from_image(cmd_buffer, dst, VK_BUFFER_USAGE_STORAGE_TEXEL_BUFFER_BIT, &dst_buffer);
1593
1594 create_bview_for_r32g32b32(cmd_buffer, radv_buffer_from_handle(src_buffer), src_offset,
1595 src->format, &src_view);
1596 create_bview_for_r32g32b32(cmd_buffer, radv_buffer_from_handle(dst_buffer), dst_offset,
1597 dst->format, &dst_view);
1598 itoi_r32g32b32_bind_descriptors(cmd_buffer, &src_view, &dst_view);
1599
1600 radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE,
1601 pipeline);
1602
1603 src_stride = get_image_stride_for_r32g32b32(cmd_buffer, src);
1604 dst_stride = get_image_stride_for_r32g32b32(cmd_buffer, dst);
1605
1606 for (unsigned r = 0; r < num_rects; ++r) {
1607 unsigned push_constants[6] = {
1608 rects[r].src_x, rects[r].src_y, src_stride, rects[r].dst_x, rects[r].dst_y, dst_stride,
1609 };
1610 radv_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer),
1611 device->meta_state.itoi_r32g32b32.img_p_layout,
1612 VK_SHADER_STAGE_COMPUTE_BIT, 0, 24, push_constants);
1613
1614 radv_unaligned_dispatch(cmd_buffer, rects[r].width, rects[r].height, 1);
1615 }
1616
1617 radv_buffer_view_finish(&src_view);
1618 radv_buffer_view_finish(&dst_view);
1619 radv_DestroyBuffer(radv_device_to_handle(device), src_buffer, NULL);
1620 radv_DestroyBuffer(radv_device_to_handle(device), dst_buffer, NULL);
1621 }
1622
1623 static void
itoi_bind_descriptors(struct radv_cmd_buffer * cmd_buffer,struct radv_image_view * src,struct radv_image_view * dst)1624 itoi_bind_descriptors(struct radv_cmd_buffer *cmd_buffer, struct radv_image_view *src,
1625 struct radv_image_view *dst)
1626 {
1627 struct radv_device *device = cmd_buffer->device;
1628
1629 radv_meta_push_descriptor_set(
1630 cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, device->meta_state.itoi.img_p_layout, 0, /* set */
1631 2, /* descriptorWriteCount */
1632 (VkWriteDescriptorSet[]){{.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1633 .dstBinding = 0,
1634 .dstArrayElement = 0,
1635 .descriptorCount = 1,
1636 .descriptorType = VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE,
1637 .pImageInfo =
1638 (VkDescriptorImageInfo[]){
1639 {
1640 .sampler = VK_NULL_HANDLE,
1641 .imageView = radv_image_view_to_handle(src),
1642 .imageLayout = VK_IMAGE_LAYOUT_GENERAL,
1643 },
1644 }},
1645 {.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1646 .dstBinding = 1,
1647 .dstArrayElement = 0,
1648 .descriptorCount = 1,
1649 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
1650 .pImageInfo = (VkDescriptorImageInfo[]){
1651 {
1652 .sampler = VK_NULL_HANDLE,
1653 .imageView = radv_image_view_to_handle(dst),
1654 .imageLayout = VK_IMAGE_LAYOUT_GENERAL,
1655 },
1656 }}});
1657 }
1658
1659 void
radv_meta_image_to_image_cs(struct radv_cmd_buffer * cmd_buffer,struct radv_meta_blit2d_surf * src,struct radv_meta_blit2d_surf * dst,unsigned num_rects,struct radv_meta_blit2d_rect * rects)1660 radv_meta_image_to_image_cs(struct radv_cmd_buffer *cmd_buffer, struct radv_meta_blit2d_surf *src,
1661 struct radv_meta_blit2d_surf *dst, unsigned num_rects,
1662 struct radv_meta_blit2d_rect *rects)
1663 {
1664 struct radv_device *device = cmd_buffer->device;
1665 struct radv_image_view src_view, dst_view;
1666 uint32_t samples = src->image->info.samples;
1667 uint32_t samples_log2 = ffs(samples) - 1;
1668
1669 if (src->format == VK_FORMAT_R32G32B32_UINT || src->format == VK_FORMAT_R32G32B32_SINT ||
1670 src->format == VK_FORMAT_R32G32B32_SFLOAT) {
1671 radv_meta_image_to_image_cs_r32g32b32(cmd_buffer, src, dst, num_rects, rects);
1672 return;
1673 }
1674
1675 u_foreach_bit(i, dst->aspect_mask) {
1676 unsigned aspect_mask = 1u << i;
1677 VkFormat depth_format = 0;
1678 if (aspect_mask == VK_IMAGE_ASPECT_STENCIL_BIT)
1679 depth_format = vk_format_stencil_only(dst->image->vk.format);
1680 else if (aspect_mask == VK_IMAGE_ASPECT_DEPTH_BIT)
1681 depth_format = vk_format_depth_only(dst->image->vk.format);
1682
1683 create_iview(cmd_buffer, src, &src_view, depth_format, aspect_mask);
1684 create_iview(cmd_buffer, dst, &dst_view, depth_format, aspect_mask);
1685
1686 itoi_bind_descriptors(cmd_buffer, &src_view, &dst_view);
1687
1688 VkPipeline pipeline = cmd_buffer->device->meta_state.itoi.pipeline[samples_log2];
1689 if (src->image->vk.image_type == VK_IMAGE_TYPE_3D ||
1690 dst->image->vk.image_type == VK_IMAGE_TYPE_3D)
1691 pipeline = cmd_buffer->device->meta_state.itoi.pipeline_3d;
1692 radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE,
1693 pipeline);
1694
1695 for (unsigned r = 0; r < num_rects; ++r) {
1696 unsigned push_constants[6] = {
1697 rects[r].src_x, rects[r].src_y, src->layer, rects[r].dst_x, rects[r].dst_y, dst->layer,
1698 };
1699 radv_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer),
1700 device->meta_state.itoi.img_p_layout, VK_SHADER_STAGE_COMPUTE_BIT, 0,
1701 24, push_constants);
1702
1703 radv_unaligned_dispatch(cmd_buffer, rects[r].width, rects[r].height, 1);
1704 }
1705
1706 radv_image_view_finish(&src_view);
1707 radv_image_view_finish(&dst_view);
1708 }
1709 }
1710
1711 static void
cleari_r32g32b32_bind_descriptors(struct radv_cmd_buffer * cmd_buffer,struct radv_buffer_view * view)1712 cleari_r32g32b32_bind_descriptors(struct radv_cmd_buffer *cmd_buffer, struct radv_buffer_view *view)
1713 {
1714 struct radv_device *device = cmd_buffer->device;
1715
1716 radv_meta_push_descriptor_set(
1717 cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, device->meta_state.cleari_r32g32b32.img_p_layout,
1718 0, /* set */
1719 1, /* descriptorWriteCount */
1720 (VkWriteDescriptorSet[]){{
1721 .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1722 .dstBinding = 0,
1723 .dstArrayElement = 0,
1724 .descriptorCount = 1,
1725 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER,
1726 .pTexelBufferView = (VkBufferView[]){radv_buffer_view_to_handle(view)},
1727 }});
1728 }
1729
1730 static void
radv_meta_clear_image_cs_r32g32b32(struct radv_cmd_buffer * cmd_buffer,struct radv_meta_blit2d_surf * dst,const VkClearColorValue * clear_color)1731 radv_meta_clear_image_cs_r32g32b32(struct radv_cmd_buffer *cmd_buffer,
1732 struct radv_meta_blit2d_surf *dst,
1733 const VkClearColorValue *clear_color)
1734 {
1735 VkPipeline pipeline = cmd_buffer->device->meta_state.cleari_r32g32b32.pipeline;
1736 struct radv_device *device = cmd_buffer->device;
1737 struct radv_buffer_view dst_view;
1738 unsigned stride;
1739 VkBuffer buffer;
1740
1741 /* This special clear path for R32G32B32 formats will write the linear
1742 * image as a buffer with the same underlying memory. The compute
1743 * shader will clear all components separately using a R32 format.
1744 */
1745 create_buffer_from_image(cmd_buffer, dst, VK_BUFFER_USAGE_STORAGE_TEXEL_BUFFER_BIT, &buffer);
1746
1747 create_bview_for_r32g32b32(cmd_buffer, radv_buffer_from_handle(buffer), 0, dst->format,
1748 &dst_view);
1749 cleari_r32g32b32_bind_descriptors(cmd_buffer, &dst_view);
1750
1751 radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE,
1752 pipeline);
1753
1754 stride = get_image_stride_for_r32g32b32(cmd_buffer, dst);
1755
1756 unsigned push_constants[4] = {
1757 clear_color->uint32[0],
1758 clear_color->uint32[1],
1759 clear_color->uint32[2],
1760 stride,
1761 };
1762
1763 radv_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer),
1764 device->meta_state.cleari_r32g32b32.img_p_layout,
1765 VK_SHADER_STAGE_COMPUTE_BIT, 0, 16, push_constants);
1766
1767 radv_unaligned_dispatch(cmd_buffer, dst->image->info.width, dst->image->info.height, 1);
1768
1769 radv_buffer_view_finish(&dst_view);
1770 radv_DestroyBuffer(radv_device_to_handle(device), buffer, NULL);
1771 }
1772
1773 static void
cleari_bind_descriptors(struct radv_cmd_buffer * cmd_buffer,struct radv_image_view * dst_iview)1774 cleari_bind_descriptors(struct radv_cmd_buffer *cmd_buffer, struct radv_image_view *dst_iview)
1775 {
1776 struct radv_device *device = cmd_buffer->device;
1777
1778 radv_meta_push_descriptor_set(cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE,
1779 device->meta_state.cleari.img_p_layout, 0, /* set */
1780 1, /* descriptorWriteCount */
1781 (VkWriteDescriptorSet[]){
1782 {.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1783 .dstBinding = 0,
1784 .dstArrayElement = 0,
1785 .descriptorCount = 1,
1786 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
1787 .pImageInfo =
1788 (VkDescriptorImageInfo[]){
1789 {
1790 .sampler = VK_NULL_HANDLE,
1791 .imageView = radv_image_view_to_handle(dst_iview),
1792 .imageLayout = VK_IMAGE_LAYOUT_GENERAL,
1793 },
1794 }},
1795 });
1796 }
1797
1798 void
radv_meta_clear_image_cs(struct radv_cmd_buffer * cmd_buffer,struct radv_meta_blit2d_surf * dst,const VkClearColorValue * clear_color)1799 radv_meta_clear_image_cs(struct radv_cmd_buffer *cmd_buffer, struct radv_meta_blit2d_surf *dst,
1800 const VkClearColorValue *clear_color)
1801 {
1802 struct radv_device *device = cmd_buffer->device;
1803 struct radv_image_view dst_iview;
1804 uint32_t samples = dst->image->info.samples;
1805 uint32_t samples_log2 = ffs(samples) - 1;
1806
1807 if (dst->format == VK_FORMAT_R32G32B32_UINT || dst->format == VK_FORMAT_R32G32B32_SINT ||
1808 dst->format == VK_FORMAT_R32G32B32_SFLOAT) {
1809 radv_meta_clear_image_cs_r32g32b32(cmd_buffer, dst, clear_color);
1810 return;
1811 }
1812
1813 create_iview(cmd_buffer, dst, &dst_iview, VK_FORMAT_UNDEFINED, dst->aspect_mask);
1814 cleari_bind_descriptors(cmd_buffer, &dst_iview);
1815
1816 VkPipeline pipeline = cmd_buffer->device->meta_state.cleari.pipeline[samples_log2];
1817 if (dst->image->vk.image_type == VK_IMAGE_TYPE_3D)
1818 pipeline = cmd_buffer->device->meta_state.cleari.pipeline_3d;
1819
1820 radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE,
1821 pipeline);
1822
1823 unsigned push_constants[5] = {
1824 clear_color->uint32[0],
1825 clear_color->uint32[1],
1826 clear_color->uint32[2],
1827 clear_color->uint32[3],
1828 dst->layer,
1829 };
1830
1831 radv_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer),
1832 device->meta_state.cleari.img_p_layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, 20,
1833 push_constants);
1834
1835 radv_unaligned_dispatch(cmd_buffer, dst->image->info.width, dst->image->info.height, 1);
1836
1837 radv_image_view_finish(&dst_iview);
1838 }
1839