2 * Copyrigh 2016 Red Hat Inc.
4 * Copyright © 2015 Intel Corporation
6 * Permission is hereby granted, free of charge, to any person obtaining a
7 * copy of this software and associated documentation files (the "Software"),
8 * to deal in the Software without restriction, including without limitation
9 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
10 * and/or sell copies of the Software, and to permit persons to whom the
11 * Software is furnished to do so, subject to the following conditions:
13 * The above copyright notice and this permission notice (including the next
14 * paragraph) shall be included in all copies or substantial portions of the
17 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
18 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
19 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
20 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
21 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
22 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
31 #include "nir/nir_builder.h"
32 #include "util/u_atomic.h"
33 #include "vulkan/vulkan_core.h"
34 #include "radv_acceleration_structure.h"
36 #include "radv_meta.h"
37 #include "radv_private.h"
40 #define TIMESTAMP_NOT_READY UINT64_MAX
42 /* TODO: Add support for mesh/task queries on GFX11 */
43 static const unsigned pipeline_statistics_indices[] = {7, 6, 3, 4, 5, 2, 1, 0, 8, 9, 10};
46 radv_get_pipelinestat_query_size(struct radv_device *device)
48 unsigned num_results = device->physical_device->rad_info.gfx_level >= GFX11 ? 14 : 11;
49 return num_results * 8;
53 radv_store_availability(nir_builder *b, nir_ssa_def *flags, nir_ssa_def *dst_buf,
54 nir_ssa_def *offset, nir_ssa_def *value32)
56 nir_push_if(b, nir_test_mask(b, flags, VK_QUERY_RESULT_WITH_AVAILABILITY_BIT));
58 nir_push_if(b, nir_test_mask(b, flags, VK_QUERY_RESULT_64_BIT));
60 nir_store_ssbo(b, nir_vec2(b, value32, nir_imm_int(b, 0)), dst_buf, offset, .align_mul = 8);
62 nir_push_else(b, NULL);
64 nir_store_ssbo(b, value32, dst_buf, offset);
72 build_occlusion_query_shader(struct radv_device *device)
74 /* the shader this builds is roughly
78 * uint32_t dst_stride;
81 * uint32_t src_stride = 16 * db_count;
83 * location(binding = 0) buffer dst_buf;
84 * location(binding = 1) buffer src_buf;
87 * uint64_t result = 0;
88 * uint64_t src_offset = src_stride * global_id.x;
89 * uint64_t dst_offset = dst_stride * global_id.x;
90 * bool available = true;
91 * for (int i = 0; i < db_count; ++i) {
92 * if (enabled_rb_mask & (1 << i)) {
93 * uint64_t start = src_buf[src_offset + 16 * i];
94 * uint64_t end = src_buf[src_offset + 16 * i + 8];
95 * if ((start & (1ull << 63)) && (end & (1ull << 63)))
96 * result += end - start;
101 * uint32_t elem_size = flags & VK_QUERY_RESULT_64_BIT ? 8 : 4;
102 * if ((flags & VK_QUERY_RESULT_PARTIAL_BIT) || available) {
103 * if (flags & VK_QUERY_RESULT_64_BIT)
104 * dst_buf[dst_offset] = result;
106 * dst_buf[dst_offset] = (uint32_t)result.
108 * if (flags & VK_QUERY_RESULT_WITH_AVAILABILITY_BIT) {
109 * dst_buf[dst_offset + elem_size] = available;
113 nir_builder b = radv_meta_init_shader(device, MESA_SHADER_COMPUTE, "occlusion_query");
114 b.shader->info.workgroup_size[0] = 64;
116 nir_variable *result = nir_local_variable_create(b.impl, glsl_uint64_t_type(), "result");
117 nir_variable *outer_counter =
118 nir_local_variable_create(b.impl, glsl_int_type(), "outer_counter");
119 nir_variable *start = nir_local_variable_create(b.impl, glsl_uint64_t_type(), "start");
120 nir_variable *end = nir_local_variable_create(b.impl, glsl_uint64_t_type(), "end");
121 nir_variable *available = nir_local_variable_create(b.impl, glsl_bool_type(), "available");
122 unsigned enabled_rb_mask = device->physical_device->rad_info.enabled_rb_mask;
123 unsigned db_count = device->physical_device->rad_info.max_render_backends;
125 nir_ssa_def *flags = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .range = 4);
127 nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0);
128 nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1);
130 nir_ssa_def *global_id = get_global_ids(&b, 1);
132 nir_ssa_def *input_stride = nir_imm_int(&b, db_count * 16);
133 nir_ssa_def *input_base = nir_imul(&b, input_stride, global_id);
134 nir_ssa_def *output_stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 4), .range = 8);
135 nir_ssa_def *output_base = nir_imul(&b, output_stride, global_id);
137 nir_store_var(&b, result, nir_imm_int64(&b, 0), 0x1);
138 nir_store_var(&b, outer_counter, nir_imm_int(&b, 0), 0x1);
139 nir_store_var(&b, available, nir_imm_true(&b), 0x1);
143 nir_ssa_def *current_outer_count = nir_load_var(&b, outer_counter);
144 radv_break_on_count(&b, outer_counter, nir_imm_int(&b, db_count));
146 nir_ssa_def *enabled_cond =
147 nir_iand_imm(&b, nir_ishl(&b, nir_imm_int(&b, 1), current_outer_count), enabled_rb_mask);
149 nir_push_if(&b, nir_i2b(&b, enabled_cond));
151 nir_ssa_def *load_offset = nir_imul_imm(&b, current_outer_count, 16);
152 load_offset = nir_iadd(&b, input_base, load_offset);
154 nir_ssa_def *load = nir_load_ssbo(&b, 2, 64, src_buf, load_offset, .align_mul = 16);
156 nir_store_var(&b, start, nir_channel(&b, load, 0), 0x1);
157 nir_store_var(&b, end, nir_channel(&b, load, 1), 0x1);
159 nir_ssa_def *start_done = nir_ilt(&b, nir_load_var(&b, start), nir_imm_int64(&b, 0));
160 nir_ssa_def *end_done = nir_ilt(&b, nir_load_var(&b, end), nir_imm_int64(&b, 0));
162 nir_push_if(&b, nir_iand(&b, start_done, end_done));
164 nir_store_var(&b, result,
165 nir_iadd(&b, nir_load_var(&b, result),
166 nir_isub(&b, nir_load_var(&b, end), nir_load_var(&b, start))),
169 nir_push_else(&b, NULL);
171 nir_store_var(&b, available, nir_imm_false(&b), 0x1);
173 nir_pop_if(&b, NULL);
174 nir_pop_if(&b, NULL);
175 nir_pop_loop(&b, NULL);
177 /* Store the result if complete or if partial results have been requested. */
179 nir_ssa_def *result_is_64bit = nir_test_mask(&b, flags, VK_QUERY_RESULT_64_BIT);
180 nir_ssa_def *result_size =
181 nir_bcsel(&b, result_is_64bit, nir_imm_int(&b, 8), nir_imm_int(&b, 4));
182 nir_push_if(&b, nir_ior(&b, nir_test_mask(&b, flags, VK_QUERY_RESULT_PARTIAL_BIT),
183 nir_load_var(&b, available)));
185 nir_push_if(&b, result_is_64bit);
187 nir_store_ssbo(&b, nir_load_var(&b, result), dst_buf, output_base, .align_mul = 8);
189 nir_push_else(&b, NULL);
191 nir_store_ssbo(&b, nir_u2u32(&b, nir_load_var(&b, result)), dst_buf, output_base,
194 nir_pop_if(&b, NULL);
195 nir_pop_if(&b, NULL);
197 radv_store_availability(&b, flags, dst_buf, nir_iadd(&b, result_size, output_base),
198 nir_b2i32(&b, nir_load_var(&b, available)));
204 build_pipeline_statistics_query_shader(struct radv_device *device)
206 unsigned pipelinestat_block_size = +radv_get_pipelinestat_query_size(device);
208 /* the shader this builds is roughly
212 * uint32_t dst_stride;
213 * uint32_t stats_mask;
214 * uint32_t avail_offset;
217 * uint32_t src_stride = pipelinestat_block_size * 2;
219 * location(binding = 0) buffer dst_buf;
220 * location(binding = 1) buffer src_buf;
223 * uint64_t src_offset = src_stride * global_id.x;
224 * uint64_t dst_base = dst_stride * global_id.x;
225 * uint64_t dst_offset = dst_base;
226 * uint32_t elem_size = flags & VK_QUERY_RESULT_64_BIT ? 8 : 4;
227 * uint32_t elem_count = stats_mask >> 16;
228 * uint32_t available32 = src_buf[avail_offset + 4 * global_id.x];
229 * if (flags & VK_QUERY_RESULT_WITH_AVAILABILITY_BIT) {
230 * dst_buf[dst_offset + elem_count * elem_size] = available32;
232 * if ((bool)available32) {
233 * // repeat 11 times:
234 * if (stats_mask & (1 << 0)) {
235 * uint64_t start = src_buf[src_offset + 8 * indices[0]];
236 * uint64_t end = src_buf[src_offset + 8 * indices[0] +
237 * pipelinestat_block_size]; uint64_t result = end - start; if (flags & VK_QUERY_RESULT_64_BIT)
238 * dst_buf[dst_offset] = result;
240 * dst_buf[dst_offset] = (uint32_t)result.
241 * dst_offset += elem_size;
243 * } else if (flags & VK_QUERY_RESULT_PARTIAL_BIT) {
244 * // Set everything to 0 as we don't know what is valid.
245 * for (int i = 0; i < elem_count; ++i)
246 * dst_buf[dst_base + elem_size * i] = 0;
250 nir_builder b = radv_meta_init_shader(device, MESA_SHADER_COMPUTE, "pipeline_statistics_query");
251 b.shader->info.workgroup_size[0] = 64;
253 nir_variable *output_offset =
254 nir_local_variable_create(b.impl, glsl_int_type(), "output_offset");
255 nir_variable *result =
256 nir_local_variable_create(b.impl, glsl_int64_t_type(), "result");
258 nir_ssa_def *flags = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .range = 4);
259 nir_ssa_def *stats_mask = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 8), .range = 12);
260 nir_ssa_def *avail_offset = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 12), .range = 16);
261 nir_ssa_def *uses_gds = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 16), .range = 20);
263 nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0);
264 nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1);
266 nir_ssa_def *global_id = get_global_ids(&b, 1);
268 nir_variable *input_stride = nir_local_variable_create(b.impl, glsl_int_type(), "input_stride");
269 nir_push_if(&b, nir_ine(&b, uses_gds, nir_imm_int(&b, 0)));
271 nir_store_var(&b, input_stride, nir_imm_int(&b, pipelinestat_block_size * 2 + 8 * 2), 0x1);
273 nir_push_else(&b, NULL);
275 nir_store_var(&b, input_stride, nir_imm_int(&b, pipelinestat_block_size * 2), 0x1);
277 nir_pop_if(&b, NULL);
279 nir_ssa_def *input_base = nir_imul(&b, nir_load_var(&b, input_stride), global_id);
280 nir_ssa_def *output_stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 4), .range = 8);
281 nir_ssa_def *output_base = nir_imul(&b, output_stride, global_id);
283 avail_offset = nir_iadd(&b, avail_offset, nir_imul_imm(&b, global_id, 4));
285 nir_ssa_def *available32 = nir_load_ssbo(&b, 1, 32, src_buf, avail_offset);
287 nir_ssa_def *result_is_64bit = nir_test_mask(&b, flags, VK_QUERY_RESULT_64_BIT);
288 nir_ssa_def *elem_size = nir_bcsel(&b, result_is_64bit, nir_imm_int(&b, 8), nir_imm_int(&b, 4));
289 nir_ssa_def *elem_count = nir_ushr_imm(&b, stats_mask, 16);
291 radv_store_availability(&b, flags, dst_buf,
292 nir_iadd(&b, output_base, nir_imul(&b, elem_count, elem_size)),
295 nir_push_if(&b, nir_i2b(&b, available32));
297 nir_store_var(&b, output_offset, output_base, 0x1);
298 for (int i = 0; i < ARRAY_SIZE(pipeline_statistics_indices); ++i) {
299 nir_push_if(&b, nir_test_mask(&b, stats_mask, BITFIELD64_BIT(i)));
301 nir_ssa_def *start_offset = nir_iadd_imm(&b, input_base, pipeline_statistics_indices[i] * 8);
302 nir_ssa_def *start = nir_load_ssbo(&b, 1, 64, src_buf, start_offset);
304 nir_ssa_def *end_offset =
305 nir_iadd_imm(&b, input_base, pipeline_statistics_indices[i] * 8 + pipelinestat_block_size);
306 nir_ssa_def *end = nir_load_ssbo(&b, 1, 64, src_buf, end_offset);
308 nir_store_var(&b, result, nir_isub(&b, end, start), 0x1);
310 nir_push_if(&b, nir_iand(&b, nir_i2b(&b, uses_gds),
311 nir_ieq(&b, nir_imm_int(&b, 1u << i),
312 nir_imm_int(&b, VK_QUERY_PIPELINE_STATISTIC_GEOMETRY_SHADER_PRIMITIVES_BIT))));
314 /* Compute the GDS result if needed. */
315 nir_ssa_def *gds_start_offset =
316 nir_iadd(&b, input_base, nir_imm_int(&b, pipelinestat_block_size * 2));
317 nir_ssa_def *gds_start = nir_load_ssbo(&b, 1, 64, src_buf, gds_start_offset);
319 nir_ssa_def *gds_end_offset =
320 nir_iadd(&b, input_base, nir_imm_int(&b, pipelinestat_block_size * 2 + 8));
321 nir_ssa_def *gds_end = nir_load_ssbo(&b, 1, 64, src_buf, gds_end_offset);
323 nir_ssa_def *ngg_gds_result = nir_isub(&b, gds_end, gds_start);
325 nir_store_var(&b, result, nir_iadd(&b, nir_load_var(&b, result), ngg_gds_result), 0x1);
327 nir_pop_if(&b, NULL);
330 nir_push_if(&b, result_is_64bit);
332 nir_store_ssbo(&b, nir_load_var(&b, result), dst_buf, nir_load_var(&b, output_offset));
334 nir_push_else(&b, NULL);
336 nir_store_ssbo(&b, nir_u2u32(&b, nir_load_var(&b, result)), dst_buf, nir_load_var(&b, output_offset));
338 nir_pop_if(&b, NULL);
340 nir_store_var(&b, output_offset, nir_iadd(&b, nir_load_var(&b, output_offset), elem_size),
343 nir_pop_if(&b, NULL);
346 nir_push_else(&b, NULL); /* nir_i2b(&b, available32) */
348 nir_push_if(&b, nir_test_mask(&b, flags, VK_QUERY_RESULT_PARTIAL_BIT));
350 /* Stores zeros in all outputs. */
352 nir_variable *counter = nir_local_variable_create(b.impl, glsl_int_type(), "counter");
353 nir_store_var(&b, counter, nir_imm_int(&b, 0), 0x1);
355 nir_loop *loop = nir_push_loop(&b);
357 nir_ssa_def *current_counter = nir_load_var(&b, counter);
358 radv_break_on_count(&b, counter, elem_count);
360 nir_ssa_def *output_elem = nir_iadd(&b, output_base, nir_imul(&b, elem_size, current_counter));
361 nir_push_if(&b, result_is_64bit);
363 nir_store_ssbo(&b, nir_imm_int64(&b, 0), dst_buf, output_elem);
365 nir_push_else(&b, NULL);
367 nir_store_ssbo(&b, nir_imm_int(&b, 0), dst_buf, output_elem);
369 nir_pop_if(&b, NULL);
371 nir_pop_loop(&b, loop);
372 nir_pop_if(&b, NULL); /* VK_QUERY_RESULT_PARTIAL_BIT */
373 nir_pop_if(&b, NULL); /* nir_i2b(&b, available32) */
378 build_tfb_query_shader(struct radv_device *device)
380 /* the shader this builds is roughly
382 * uint32_t src_stride = 32;
384 * location(binding = 0) buffer dst_buf;
385 * location(binding = 1) buffer src_buf;
388 * uint64_t result[2] = {};
389 * bool available = false;
390 * uint64_t src_offset = src_stride * global_id.x;
391 * uint64_t dst_offset = dst_stride * global_id.x;
392 * uint64_t *src_data = src_buf[src_offset];
393 * uint32_t avail = (src_data[0] >> 32) &
394 * (src_data[1] >> 32) &
395 * (src_data[2] >> 32) &
396 * (src_data[3] >> 32);
397 * if (avail & 0x80000000) {
398 * result[0] = src_data[3] - src_data[1];
399 * result[1] = src_data[2] - src_data[0];
402 * uint32_t result_size = flags & VK_QUERY_RESULT_64_BIT ? 16 : 8;
403 * if ((flags & VK_QUERY_RESULT_PARTIAL_BIT) || available) {
404 * if (flags & VK_QUERY_RESULT_64_BIT) {
405 * dst_buf[dst_offset] = result;
407 * dst_buf[dst_offset] = (uint32_t)result;
410 * if (flags & VK_QUERY_RESULT_WITH_AVAILABILITY_BIT) {
411 * dst_buf[dst_offset + result_size] = available;
415 nir_builder b = radv_meta_init_shader(device, MESA_SHADER_COMPUTE, "tfb_query");
416 b.shader->info.workgroup_size[0] = 64;
418 /* Create and initialize local variables. */
419 nir_variable *result =
420 nir_local_variable_create(b.impl, glsl_vector_type(GLSL_TYPE_UINT64, 2), "result");
421 nir_variable *available = nir_local_variable_create(b.impl, glsl_bool_type(), "available");
423 nir_store_var(&b, result, nir_vec2(&b, nir_imm_int64(&b, 0), nir_imm_int64(&b, 0)), 0x3);
424 nir_store_var(&b, available, nir_imm_false(&b), 0x1);
426 nir_ssa_def *flags = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .range = 4);
428 /* Load resources. */
429 nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0);
430 nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1);
432 /* Compute global ID. */
433 nir_ssa_def *global_id = get_global_ids(&b, 1);
435 /* Compute src/dst strides. */
436 nir_ssa_def *input_stride = nir_imm_int(&b, 32);
437 nir_ssa_def *input_base = nir_imul(&b, input_stride, global_id);
438 nir_ssa_def *output_stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 4), .range = 8);
439 nir_ssa_def *output_base = nir_imul(&b, output_stride, global_id);
441 /* Load data from the query pool. */
442 nir_ssa_def *load1 = nir_load_ssbo(&b, 4, 32, src_buf, input_base, .align_mul = 32);
444 nir_load_ssbo(&b, 4, 32, src_buf, nir_iadd_imm(&b, input_base, 16), .align_mul = 16);
446 /* Check if result is available. */
447 nir_ssa_def *avails[2];
448 avails[0] = nir_iand(&b, nir_channel(&b, load1, 1), nir_channel(&b, load1, 3));
449 avails[1] = nir_iand(&b, nir_channel(&b, load2, 1), nir_channel(&b, load2, 3));
450 nir_ssa_def *result_is_available =
451 nir_test_mask(&b, nir_iand(&b, avails[0], avails[1]), 0x80000000);
453 /* Only compute result if available. */
454 nir_push_if(&b, result_is_available);
457 nir_ssa_def *packed64[4];
459 nir_pack_64_2x32(&b, nir_vec2(&b, nir_channel(&b, load1, 0), nir_channel(&b, load1, 1)));
461 nir_pack_64_2x32(&b, nir_vec2(&b, nir_channel(&b, load1, 2), nir_channel(&b, load1, 3)));
463 nir_pack_64_2x32(&b, nir_vec2(&b, nir_channel(&b, load2, 0), nir_channel(&b, load2, 1)));
465 nir_pack_64_2x32(&b, nir_vec2(&b, nir_channel(&b, load2, 2), nir_channel(&b, load2, 3)));
467 /* Compute result. */
468 nir_ssa_def *num_primitive_written = nir_isub(&b, packed64[3], packed64[1]);
469 nir_ssa_def *primitive_storage_needed = nir_isub(&b, packed64[2], packed64[0]);
471 nir_store_var(&b, result, nir_vec2(&b, num_primitive_written, primitive_storage_needed), 0x3);
472 nir_store_var(&b, available, nir_imm_true(&b), 0x1);
474 nir_pop_if(&b, NULL);
476 /* Determine if result is 64 or 32 bit. */
477 nir_ssa_def *result_is_64bit = nir_test_mask(&b, flags, VK_QUERY_RESULT_64_BIT);
478 nir_ssa_def *result_size =
479 nir_bcsel(&b, result_is_64bit, nir_imm_int(&b, 16), nir_imm_int(&b, 8));
481 /* Store the result if complete or partial results have been requested. */
482 nir_push_if(&b, nir_ior(&b, nir_test_mask(&b, flags, VK_QUERY_RESULT_PARTIAL_BIT),
483 nir_load_var(&b, available)));
486 nir_push_if(&b, result_is_64bit);
488 nir_store_ssbo(&b, nir_load_var(&b, result), dst_buf, output_base);
490 nir_push_else(&b, NULL);
492 nir_store_ssbo(&b, nir_u2u32(&b, nir_load_var(&b, result)), dst_buf, output_base);
494 nir_pop_if(&b, NULL);
495 nir_pop_if(&b, NULL);
497 radv_store_availability(&b, flags, dst_buf, nir_iadd(&b, result_size, output_base),
498 nir_b2i32(&b, nir_load_var(&b, available)));
504 build_timestamp_query_shader(struct radv_device *device)
506 /* the shader this builds is roughly
508 * uint32_t src_stride = 8;
510 * location(binding = 0) buffer dst_buf;
511 * location(binding = 1) buffer src_buf;
514 * uint64_t result = 0;
515 * bool available = false;
516 * uint64_t src_offset = src_stride * global_id.x;
517 * uint64_t dst_offset = dst_stride * global_id.x;
518 * uint64_t timestamp = src_buf[src_offset];
519 * if (timestamp != TIMESTAMP_NOT_READY) {
520 * result = timestamp;
523 * uint32_t result_size = flags & VK_QUERY_RESULT_64_BIT ? 8 : 4;
524 * if ((flags & VK_QUERY_RESULT_PARTIAL_BIT) || available) {
525 * if (flags & VK_QUERY_RESULT_64_BIT) {
526 * dst_buf[dst_offset] = result;
528 * dst_buf[dst_offset] = (uint32_t)result;
531 * if (flags & VK_QUERY_RESULT_WITH_AVAILABILITY_BIT) {
532 * dst_buf[dst_offset + result_size] = available;
536 nir_builder b = radv_meta_init_shader(device, MESA_SHADER_COMPUTE, "timestamp_query");
537 b.shader->info.workgroup_size[0] = 64;
539 /* Create and initialize local variables. */
540 nir_variable *result = nir_local_variable_create(b.impl, glsl_uint64_t_type(), "result");
541 nir_variable *available = nir_local_variable_create(b.impl, glsl_bool_type(), "available");
543 nir_store_var(&b, result, nir_imm_int64(&b, 0), 0x1);
544 nir_store_var(&b, available, nir_imm_false(&b), 0x1);
546 nir_ssa_def *flags = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .range = 4);
548 /* Load resources. */
549 nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0);
550 nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1);
552 /* Compute global ID. */
553 nir_ssa_def *global_id = get_global_ids(&b, 1);
555 /* Compute src/dst strides. */
556 nir_ssa_def *input_stride = nir_imm_int(&b, 8);
557 nir_ssa_def *input_base = nir_imul(&b, input_stride, global_id);
558 nir_ssa_def *output_stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 4), .range = 8);
559 nir_ssa_def *output_base = nir_imul(&b, output_stride, global_id);
561 /* Load data from the query pool. */
562 nir_ssa_def *load = nir_load_ssbo(&b, 2, 32, src_buf, input_base, .align_mul = 8);
564 /* Pack the timestamp. */
565 nir_ssa_def *timestamp;
567 nir_pack_64_2x32(&b, nir_vec2(&b, nir_channel(&b, load, 0), nir_channel(&b, load, 1)));
569 /* Check if result is available. */
570 nir_ssa_def *result_is_available = nir_i2b(&b, nir_ine_imm(&b, timestamp, TIMESTAMP_NOT_READY));
572 /* Only store result if available. */
573 nir_push_if(&b, result_is_available);
575 nir_store_var(&b, result, timestamp, 0x1);
576 nir_store_var(&b, available, nir_imm_true(&b), 0x1);
578 nir_pop_if(&b, NULL);
580 /* Determine if result is 64 or 32 bit. */
581 nir_ssa_def *result_is_64bit = nir_test_mask(&b, flags, VK_QUERY_RESULT_64_BIT);
582 nir_ssa_def *result_size =
583 nir_bcsel(&b, result_is_64bit, nir_imm_int(&b, 8), nir_imm_int(&b, 4));
585 /* Store the result if complete or partial results have been requested. */
586 nir_push_if(&b, nir_ior(&b, nir_test_mask(&b, flags, VK_QUERY_RESULT_PARTIAL_BIT),
587 nir_load_var(&b, available)));
590 nir_push_if(&b, result_is_64bit);
592 nir_store_ssbo(&b, nir_load_var(&b, result), dst_buf, output_base);
594 nir_push_else(&b, NULL);
596 nir_store_ssbo(&b, nir_u2u32(&b, nir_load_var(&b, result)), dst_buf, output_base);
598 nir_pop_if(&b, NULL);
600 nir_pop_if(&b, NULL);
602 radv_store_availability(&b, flags, dst_buf, nir_iadd(&b, result_size, output_base),
603 nir_b2i32(&b, nir_load_var(&b, available)));
609 build_pg_query_shader(struct radv_device *device)
611 /* the shader this builds is roughly
613 * uint32_t src_stride = 32;
615 * location(binding = 0) buffer dst_buf;
616 * location(binding = 1) buffer src_buf;
619 * uint64_t result = {};
620 * bool available = false;
621 * uint64_t src_offset = src_stride * global_id.x;
622 * uint64_t dst_offset = dst_stride * global_id.x;
623 * uint64_t *src_data = src_buf[src_offset];
624 * uint32_t avail = (src_data[0] >> 32) &
625 * (src_data[2] >> 32);
626 * if (avail & 0x80000000) {
627 * result = src_data[2] - src_data[0];
629 * uint32_t ngg_gds_result = 0;
630 * ngg_gds_result += src_data[9] - src_data[8];
631 * result += (uint64_t)ngg_gds_result;
635 * uint32_t result_size = flags & VK_QUERY_RESULT_64_BIT ? 16 : 8;
636 * if ((flags & VK_QUERY_RESULT_PARTIAL_BIT) || available) {
637 * if (flags & VK_QUERY_RESULT_64_BIT) {
638 * dst_buf[dst_offset] = result;
640 * dst_buf[dst_offset] = (uint32_t)result;
643 * if (flags & VK_QUERY_RESULT_WITH_AVAILABILITY_BIT) {
644 * dst_buf[dst_offset + result_size] = available;
648 nir_builder b = radv_meta_init_shader(device, MESA_SHADER_COMPUTE, "pg_query");
649 b.shader->info.workgroup_size[0] = 64;
651 /* Create and initialize local variables. */
652 nir_variable *result =
653 nir_local_variable_create(b.impl, glsl_uint64_t_type(), "result");
654 nir_variable *available = nir_local_variable_create(b.impl, glsl_bool_type(), "available");
656 nir_store_var(&b, result, nir_imm_int64(&b, 0), 0x1);
657 nir_store_var(&b, available, nir_imm_false(&b), 0x1);
659 nir_ssa_def *flags = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .range = 16);
661 /* Load resources. */
662 nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0);
663 nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1);
665 /* Compute global ID. */
666 nir_ssa_def *global_id = get_global_ids(&b, 1);
668 /* Compute src/dst strides. */
669 nir_ssa_def *input_stride = nir_imm_int(&b, 32);
670 nir_ssa_def *input_base = nir_imul(&b, input_stride, global_id);
671 nir_ssa_def *output_stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 4), .range = 16);
672 nir_ssa_def *output_base = nir_imul(&b, output_stride, global_id);
674 /* Load data from the query pool. */
675 nir_ssa_def *load1 = nir_load_ssbo(&b, 2, 32, src_buf, input_base, .align_mul = 32);
676 nir_ssa_def *load2 = nir_load_ssbo(
677 &b, 2, 32, src_buf, nir_iadd(&b, input_base, nir_imm_int(&b, 16)), .align_mul = 16);
679 /* Check if result is available. */
680 nir_ssa_def *avails[2];
681 avails[0] = nir_channel(&b, load1, 1);
682 avails[1] = nir_channel(&b, load2, 1);
683 nir_ssa_def *result_is_available =
684 nir_i2b(&b, nir_iand(&b, nir_iand(&b, avails[0], avails[1]), nir_imm_int(&b, 0x80000000)));
686 /* Only compute result if available. */
687 nir_push_if(&b, result_is_available);
690 nir_ssa_def *packed64[2];
692 nir_pack_64_2x32(&b, nir_vec2(&b, nir_channel(&b, load1, 0), nir_channel(&b, load1, 1)));
694 nir_pack_64_2x32(&b, nir_vec2(&b, nir_channel(&b, load2, 0), nir_channel(&b, load2, 1)));
696 /* Compute result. */
697 nir_ssa_def *primitive_storage_needed = nir_isub(&b, packed64[1], packed64[0]);
699 nir_store_var(&b, result, primitive_storage_needed, 0x1);
701 nir_ssa_def *uses_gds = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 16), .range = 20);
702 nir_push_if(&b, nir_i2b(&b, uses_gds));
704 nir_ssa_def *gds_start =
705 nir_load_ssbo(&b, 1, 32, src_buf, nir_iadd(&b, input_base, nir_imm_int(&b, 32)), .align_mul = 4);
706 nir_ssa_def *gds_end =
707 nir_load_ssbo(&b, 1, 32, src_buf, nir_iadd(&b, input_base, nir_imm_int(&b, 36)), .align_mul = 4);
709 nir_ssa_def *ngg_gds_result = nir_isub(&b, gds_end, gds_start);
711 nir_store_var(&b, result, nir_iadd(&b, nir_load_var(&b, result), nir_u2u64(&b, ngg_gds_result)), 0x1);
713 nir_pop_if(&b, NULL);
715 nir_store_var(&b, available, nir_imm_true(&b), 0x1);
717 nir_pop_if(&b, NULL);
719 /* Determine if result is 64 or 32 bit. */
720 nir_ssa_def *result_is_64bit = nir_test_mask(&b, flags, VK_QUERY_RESULT_64_BIT);
721 nir_ssa_def *result_size =
722 nir_bcsel(&b, result_is_64bit, nir_imm_int(&b, 16), nir_imm_int(&b, 8));
724 /* Store the result if complete or partial results have been requested. */
725 nir_push_if(&b, nir_ior(&b, nir_test_mask(&b, flags, VK_QUERY_RESULT_PARTIAL_BIT),
726 nir_load_var(&b, available)));
729 nir_push_if(&b, result_is_64bit);
731 nir_store_ssbo(&b, nir_load_var(&b, result), dst_buf, output_base);
733 nir_push_else(&b, NULL);
735 nir_store_ssbo(&b, nir_u2u32(&b, nir_load_var(&b, result)), dst_buf, output_base);
737 nir_pop_if(&b, NULL);
738 nir_pop_if(&b, NULL);
740 radv_store_availability(&b, flags, dst_buf, nir_iadd(&b, result_size, output_base),
741 nir_b2i32(&b, nir_load_var(&b, available)));
747 radv_device_init_meta_query_state_internal(struct radv_device *device)
750 nir_shader *occlusion_cs = NULL;
751 nir_shader *pipeline_statistics_cs = NULL;
752 nir_shader *tfb_cs = NULL;
753 nir_shader *timestamp_cs = NULL;
754 nir_shader *pg_cs = NULL;
756 mtx_lock(&device->meta_state.mtx);
757 if (device->meta_state.query.pipeline_statistics_query_pipeline) {
758 mtx_unlock(&device->meta_state.mtx);
761 occlusion_cs = build_occlusion_query_shader(device);
762 pipeline_statistics_cs = build_pipeline_statistics_query_shader(device);
763 tfb_cs = build_tfb_query_shader(device);
764 timestamp_cs = build_timestamp_query_shader(device);
765 pg_cs = build_pg_query_shader(device);
767 VkDescriptorSetLayoutCreateInfo occlusion_ds_create_info = {
768 .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,
769 .flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,
771 .pBindings = (VkDescriptorSetLayoutBinding[]){
773 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
774 .descriptorCount = 1,
775 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
776 .pImmutableSamplers = NULL},
778 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
779 .descriptorCount = 1,
780 .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
781 .pImmutableSamplers = NULL},
784 result = radv_CreateDescriptorSetLayout(radv_device_to_handle(device), &occlusion_ds_create_info,
785 &device->meta_state.alloc,
786 &device->meta_state.query.ds_layout);
787 if (result != VK_SUCCESS)
790 VkPipelineLayoutCreateInfo occlusion_pl_create_info = {
791 .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,
793 .pSetLayouts = &device->meta_state.query.ds_layout,
794 .pushConstantRangeCount = 1,
795 .pPushConstantRanges = &(VkPushConstantRange){VK_SHADER_STAGE_COMPUTE_BIT, 0, 20},
799 radv_CreatePipelineLayout(radv_device_to_handle(device), &occlusion_pl_create_info,
800 &device->meta_state.alloc, &device->meta_state.query.p_layout);
801 if (result != VK_SUCCESS)
804 VkPipelineShaderStageCreateInfo occlusion_pipeline_shader_stage = {
805 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
806 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
807 .module = vk_shader_module_handle_from_nir(occlusion_cs),
809 .pSpecializationInfo = NULL,
812 VkComputePipelineCreateInfo occlusion_vk_pipeline_info = {
813 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
814 .stage = occlusion_pipeline_shader_stage,
816 .layout = device->meta_state.query.p_layout,
819 result = radv_CreateComputePipelines(
820 radv_device_to_handle(device), device->meta_state.cache, 1,
821 &occlusion_vk_pipeline_info, NULL, &device->meta_state.query.occlusion_query_pipeline);
822 if (result != VK_SUCCESS)
825 VkPipelineShaderStageCreateInfo pipeline_statistics_pipeline_shader_stage = {
826 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
827 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
828 .module = vk_shader_module_handle_from_nir(pipeline_statistics_cs),
830 .pSpecializationInfo = NULL,
833 VkComputePipelineCreateInfo pipeline_statistics_vk_pipeline_info = {
834 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
835 .stage = pipeline_statistics_pipeline_shader_stage,
837 .layout = device->meta_state.query.p_layout,
840 result = radv_CreateComputePipelines(
841 radv_device_to_handle(device), device->meta_state.cache, 1,
842 &pipeline_statistics_vk_pipeline_info, NULL,
843 &device->meta_state.query.pipeline_statistics_query_pipeline);
844 if (result != VK_SUCCESS)
847 VkPipelineShaderStageCreateInfo tfb_pipeline_shader_stage = {
848 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
849 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
850 .module = vk_shader_module_handle_from_nir(tfb_cs),
852 .pSpecializationInfo = NULL,
855 VkComputePipelineCreateInfo tfb_pipeline_info = {
856 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
857 .stage = tfb_pipeline_shader_stage,
859 .layout = device->meta_state.query.p_layout,
862 result = radv_CreateComputePipelines(
863 radv_device_to_handle(device), device->meta_state.cache, 1,
864 &tfb_pipeline_info, NULL, &device->meta_state.query.tfb_query_pipeline);
865 if (result != VK_SUCCESS)
868 VkPipelineShaderStageCreateInfo timestamp_pipeline_shader_stage = {
869 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
870 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
871 .module = vk_shader_module_handle_from_nir(timestamp_cs),
873 .pSpecializationInfo = NULL,
876 VkComputePipelineCreateInfo timestamp_pipeline_info = {
877 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
878 .stage = timestamp_pipeline_shader_stage,
880 .layout = device->meta_state.query.p_layout,
883 result = radv_CreateComputePipelines(
884 radv_device_to_handle(device), device->meta_state.cache, 1,
885 ×tamp_pipeline_info, NULL, &device->meta_state.query.timestamp_query_pipeline);
886 if (result != VK_SUCCESS)
889 VkPipelineShaderStageCreateInfo pg_pipeline_shader_stage = {
890 .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
891 .stage = VK_SHADER_STAGE_COMPUTE_BIT,
892 .module = vk_shader_module_handle_from_nir(pg_cs),
894 .pSpecializationInfo = NULL,
897 VkComputePipelineCreateInfo pg_pipeline_info = {
898 .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
899 .stage = pg_pipeline_shader_stage,
901 .layout = device->meta_state.query.p_layout,
904 result = radv_CreateComputePipelines(
905 radv_device_to_handle(device), device->meta_state.cache, 1,
906 &pg_pipeline_info, NULL, &device->meta_state.query.pg_query_pipeline);
909 ralloc_free(occlusion_cs);
910 ralloc_free(pipeline_statistics_cs);
913 ralloc_free(timestamp_cs);
914 mtx_unlock(&device->meta_state.mtx);
919 radv_device_init_meta_query_state(struct radv_device *device, bool on_demand)
924 return radv_device_init_meta_query_state_internal(device);
928 radv_device_finish_meta_query_state(struct radv_device *device)
930 if (device->meta_state.query.tfb_query_pipeline)
931 radv_DestroyPipeline(radv_device_to_handle(device),
932 device->meta_state.query.tfb_query_pipeline, &device->meta_state.alloc);
934 if (device->meta_state.query.pipeline_statistics_query_pipeline)
935 radv_DestroyPipeline(radv_device_to_handle(device),
936 device->meta_state.query.pipeline_statistics_query_pipeline,
937 &device->meta_state.alloc);
939 if (device->meta_state.query.occlusion_query_pipeline)
940 radv_DestroyPipeline(radv_device_to_handle(device),
941 device->meta_state.query.occlusion_query_pipeline,
942 &device->meta_state.alloc);
944 if (device->meta_state.query.timestamp_query_pipeline)
945 radv_DestroyPipeline(radv_device_to_handle(device),
946 device->meta_state.query.timestamp_query_pipeline,
947 &device->meta_state.alloc);
949 if (device->meta_state.query.pg_query_pipeline)
950 radv_DestroyPipeline(radv_device_to_handle(device),
951 device->meta_state.query.pg_query_pipeline, &device->meta_state.alloc);
953 if (device->meta_state.query.p_layout)
954 radv_DestroyPipelineLayout(radv_device_to_handle(device), device->meta_state.query.p_layout,
955 &device->meta_state.alloc);
957 if (device->meta_state.query.ds_layout)
958 device->vk.dispatch_table.DestroyDescriptorSetLayout(radv_device_to_handle(device),
959 device->meta_state.query.ds_layout,
960 &device->meta_state.alloc);
964 radv_query_shader(struct radv_cmd_buffer *cmd_buffer, VkPipeline *pipeline,
965 struct radeon_winsys_bo *src_bo, struct radeon_winsys_bo *dst_bo,
966 uint64_t src_offset, uint64_t dst_offset, uint32_t src_stride,
967 uint32_t dst_stride, size_t dst_size, uint32_t count, uint32_t flags,
968 uint32_t pipeline_stats_mask, uint32_t avail_offset, bool uses_gds)
970 struct radv_device *device = cmd_buffer->device;
971 struct radv_meta_saved_state saved_state;
972 struct radv_buffer src_buffer, dst_buffer;
975 VkResult ret = radv_device_init_meta_query_state_internal(device);
976 if (ret != VK_SUCCESS) {
977 vk_command_buffer_set_error(&cmd_buffer->vk, ret);
982 /* VK_EXT_conditional_rendering says that copy commands should not be
983 * affected by conditional rendering.
985 radv_meta_save(&saved_state, cmd_buffer,
986 RADV_META_SAVE_COMPUTE_PIPELINE | RADV_META_SAVE_CONSTANTS |
987 RADV_META_SAVE_DESCRIPTORS | RADV_META_SUSPEND_PREDICATING);
989 uint64_t src_buffer_size = MAX2(src_stride * count, avail_offset + 4 * count - src_offset);
990 uint64_t dst_buffer_size = dst_stride * (count - 1) + dst_size;
992 radv_buffer_init(&src_buffer, device, src_bo, src_buffer_size, src_offset);
993 radv_buffer_init(&dst_buffer, device, dst_bo, dst_buffer_size, dst_offset);
995 radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE,
998 radv_meta_push_descriptor_set(
999 cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, device->meta_state.query.p_layout, 0, /* set */
1000 2, /* descriptorWriteCount */
1001 (VkWriteDescriptorSet[]){
1002 {.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1004 .dstArrayElement = 0,
1005 .descriptorCount = 1,
1006 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
1007 .pBufferInfo = &(VkDescriptorBufferInfo){.buffer = radv_buffer_to_handle(&dst_buffer),
1009 .range = VK_WHOLE_SIZE}},
1010 {.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,
1012 .dstArrayElement = 0,
1013 .descriptorCount = 1,
1014 .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
1015 .pBufferInfo = &(VkDescriptorBufferInfo){.buffer = radv_buffer_to_handle(&src_buffer),
1017 .range = VK_WHOLE_SIZE}}});
1019 /* Encode the number of elements for easy access by the shader. */
1020 pipeline_stats_mask &= 0x7ff;
1021 pipeline_stats_mask |= util_bitcount(pipeline_stats_mask) << 16;
1023 avail_offset -= src_offset;
1027 uint32_t dst_stride;
1028 uint32_t pipeline_stats_mask;
1029 uint32_t avail_offset;
1031 } push_constants = {flags, dst_stride, pipeline_stats_mask, avail_offset, uses_gds};
1033 radv_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer), device->meta_state.query.p_layout,
1034 VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(push_constants), &push_constants);
1036 cmd_buffer->state.flush_bits |= RADV_CMD_FLAG_INV_L2 | RADV_CMD_FLAG_INV_VCACHE;
1038 if (flags & VK_QUERY_RESULT_WAIT_BIT)
1039 cmd_buffer->state.flush_bits |= RADV_CMD_FLUSH_AND_INV_FRAMEBUFFER;
1041 radv_unaligned_dispatch(cmd_buffer, count, 1, 1);
1043 /* Ensure that the query copy dispatch is complete before a potential vkCmdResetPool because
1044 * there is an implicit execution dependency from each such query command to all query commands
1045 * previously submitted to the same queue.
1047 cmd_buffer->active_query_flush_bits |=
1048 RADV_CMD_FLAG_CS_PARTIAL_FLUSH | RADV_CMD_FLAG_INV_L2 | RADV_CMD_FLAG_INV_VCACHE;
1050 radv_buffer_finish(&src_buffer);
1051 radv_buffer_finish(&dst_buffer);
1053 radv_meta_restore(&saved_state, cmd_buffer);
1057 radv_destroy_query_pool(struct radv_device *device, const VkAllocationCallbacks *pAllocator,
1058 struct radv_query_pool *pool)
1060 if (pool->type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR)
1061 radv_pc_deinit_query_pool((struct radv_pc_query_pool *)pool);
1064 device->ws->buffer_destroy(device->ws, pool->bo);
1065 vk_object_base_finish(&pool->base);
1066 vk_free2(&device->vk.alloc, pAllocator, pool);
1069 VKAPI_ATTR VkResult VKAPI_CALL
1070 radv_CreateQueryPool(VkDevice _device, const VkQueryPoolCreateInfo *pCreateInfo,
1071 const VkAllocationCallbacks *pAllocator, VkQueryPool *pQueryPool)
1073 RADV_FROM_HANDLE(radv_device, device, _device);
1075 size_t pool_struct_size = pCreateInfo->queryType == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR
1076 ? sizeof(struct radv_pc_query_pool)
1077 : sizeof(struct radv_query_pool);
1079 struct radv_query_pool *pool = vk_alloc2(&device->vk.alloc, pAllocator, pool_struct_size, 8,
1080 VK_SYSTEM_ALLOCATION_SCOPE_OBJECT);
1083 return vk_error(device, VK_ERROR_OUT_OF_HOST_MEMORY);
1085 vk_object_base_init(&device->vk, &pool->base, VK_OBJECT_TYPE_QUERY_POOL);
1087 pool->type = pCreateInfo->queryType;
1088 pool->pipeline_stats_mask = pCreateInfo->pipelineStatistics;
1090 /* The number of primitives generated by geometry shader invocations is only counted by the
1091 * hardware if GS uses the legacy path. When NGG GS is used, the hardware can't know the number
1092 * of generated primitives and we have to increment it from the shader using a plain GDS atomic.
1094 pool->uses_gds = device->physical_device->use_ngg &&
1095 ((pool->pipeline_stats_mask & VK_QUERY_PIPELINE_STATISTIC_GEOMETRY_SHADER_PRIMITIVES_BIT) ||
1096 pCreateInfo->queryType == VK_QUERY_TYPE_PRIMITIVES_GENERATED_EXT);
1098 switch (pCreateInfo->queryType) {
1099 case VK_QUERY_TYPE_OCCLUSION:
1100 pool->stride = 16 * device->physical_device->rad_info.max_render_backends;
1102 case VK_QUERY_TYPE_PIPELINE_STATISTICS:
1103 pool->stride = radv_get_pipelinestat_query_size(device) * 2;
1104 if (pool->uses_gds) {
1105 /* When the query pool needs GDS (for counting the number of primitives generated by a
1106 * geometry shader with NGG), allocate 2x64-bit values for begin/end.
1108 pool->stride += 8 * 2;
1111 case VK_QUERY_TYPE_TIMESTAMP:
1112 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_COMPACTED_SIZE_KHR:
1113 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_SIZE_KHR:
1114 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_BOTTOM_LEVEL_POINTERS_KHR:
1115 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SIZE_KHR:
1118 case VK_QUERY_TYPE_TRANSFORM_FEEDBACK_STREAM_EXT:
1121 case VK_QUERY_TYPE_PRIMITIVES_GENERATED_EXT:
1123 if (pool->uses_gds && device->physical_device->rad_info.gfx_level < GFX11) {
1124 /* When the hardware can use both the legacy and the NGG paths in the same begin/end pair,
1125 * allocate 2x32-bit values for the GDS counters.
1127 pool->stride += 4 * 2;
1130 case VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR: {
1131 result = radv_pc_init_query_pool(device->physical_device, pCreateInfo,
1132 (struct radv_pc_query_pool *)pool);
1134 if (result != VK_SUCCESS) {
1135 radv_destroy_query_pool(device, pAllocator, pool);
1136 return vk_error(device, result);
1141 unreachable("creating unhandled query type");
1144 pool->availability_offset = pool->stride * pCreateInfo->queryCount;
1145 pool->size = pool->availability_offset;
1146 if (pCreateInfo->queryType == VK_QUERY_TYPE_PIPELINE_STATISTICS)
1147 pool->size += 4 * pCreateInfo->queryCount;
1149 result = device->ws->buffer_create(device->ws, pool->size, 64, RADEON_DOMAIN_GTT,
1150 RADEON_FLAG_NO_INTERPROCESS_SHARING,
1151 RADV_BO_PRIORITY_QUERY_POOL, 0, &pool->bo);
1152 if (result != VK_SUCCESS) {
1153 radv_destroy_query_pool(device, pAllocator, pool);
1154 return vk_error(device, result);
1157 pool->ptr = device->ws->buffer_map(pool->bo);
1159 radv_destroy_query_pool(device, pAllocator, pool);
1160 return vk_error(device, VK_ERROR_OUT_OF_DEVICE_MEMORY);
1163 *pQueryPool = radv_query_pool_to_handle(pool);
1167 VKAPI_ATTR void VKAPI_CALL
1168 radv_DestroyQueryPool(VkDevice _device, VkQueryPool _pool, const VkAllocationCallbacks *pAllocator)
1170 RADV_FROM_HANDLE(radv_device, device, _device);
1171 RADV_FROM_HANDLE(radv_query_pool, pool, _pool);
1176 radv_destroy_query_pool(device, pAllocator, pool);
1179 VKAPI_ATTR VkResult VKAPI_CALL
1180 radv_GetQueryPoolResults(VkDevice _device, VkQueryPool queryPool, uint32_t firstQuery,
1181 uint32_t queryCount, size_t dataSize, void *pData, VkDeviceSize stride,
1182 VkQueryResultFlags flags)
1184 RADV_FROM_HANDLE(radv_device, device, _device);
1185 RADV_FROM_HANDLE(radv_query_pool, pool, queryPool);
1187 VkResult result = VK_SUCCESS;
1189 if (vk_device_is_lost(&device->vk))
1190 return VK_ERROR_DEVICE_LOST;
1192 for (unsigned query_idx = 0; query_idx < queryCount; ++query_idx, data += stride) {
1194 unsigned query = firstQuery + query_idx;
1195 char *src = pool->ptr + query * pool->stride;
1198 switch (pool->type) {
1199 case VK_QUERY_TYPE_TIMESTAMP:
1200 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_COMPACTED_SIZE_KHR:
1201 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_SIZE_KHR:
1202 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_BOTTOM_LEVEL_POINTERS_KHR:
1203 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SIZE_KHR: {
1204 uint64_t const *src64 = (uint64_t const *)src;
1208 value = p_atomic_read(src64);
1209 } while (value == TIMESTAMP_NOT_READY && (flags & VK_QUERY_RESULT_WAIT_BIT));
1211 available = value != TIMESTAMP_NOT_READY;
1213 if (!available && !(flags & VK_QUERY_RESULT_PARTIAL_BIT))
1214 result = VK_NOT_READY;
1216 if (flags & VK_QUERY_RESULT_64_BIT) {
1217 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT))
1218 *(uint64_t *)dest = value;
1221 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT))
1222 *(uint32_t *)dest = (uint32_t)value;
1227 case VK_QUERY_TYPE_OCCLUSION: {
1228 uint64_t const *src64 = (uint64_t const *)src;
1229 uint32_t db_count = device->physical_device->rad_info.max_render_backends;
1230 uint32_t enabled_rb_mask = device->physical_device->rad_info.enabled_rb_mask;
1231 uint64_t sample_count = 0;
1234 for (int i = 0; i < db_count; ++i) {
1235 uint64_t start, end;
1237 if (!(enabled_rb_mask & (1 << i)))
1241 start = p_atomic_read(src64 + 2 * i);
1242 end = p_atomic_read(src64 + 2 * i + 1);
1243 } while ((!(start & (1ull << 63)) || !(end & (1ull << 63))) &&
1244 (flags & VK_QUERY_RESULT_WAIT_BIT));
1246 if (!(start & (1ull << 63)) || !(end & (1ull << 63)))
1249 sample_count += end - start;
1253 if (!available && !(flags & VK_QUERY_RESULT_PARTIAL_BIT))
1254 result = VK_NOT_READY;
1256 if (flags & VK_QUERY_RESULT_64_BIT) {
1257 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT))
1258 *(uint64_t *)dest = sample_count;
1261 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT))
1262 *(uint32_t *)dest = sample_count;
1267 case VK_QUERY_TYPE_PIPELINE_STATISTICS: {
1268 unsigned pipelinestat_block_size = radv_get_pipelinestat_query_size(device);
1269 const uint32_t *avail_ptr =
1270 (const uint32_t *)(pool->ptr + pool->availability_offset + 4 * query);
1271 uint64_t ngg_gds_result = 0;
1274 available = p_atomic_read(avail_ptr);
1275 } while (!available && (flags & VK_QUERY_RESULT_WAIT_BIT));
1277 if (!available && !(flags & VK_QUERY_RESULT_PARTIAL_BIT))
1278 result = VK_NOT_READY;
1280 if (pool->uses_gds) {
1281 /* Compute the result that was copied from GDS. */
1282 const uint64_t *gds_start = (uint64_t *)(src + pipelinestat_block_size * 2);
1283 const uint64_t *gds_stop = (uint64_t *)(src + pipelinestat_block_size * 2 + 8);
1285 ngg_gds_result = gds_stop[0] - gds_start[0];
1288 const uint64_t *start = (uint64_t *)src;
1289 const uint64_t *stop = (uint64_t *)(src + pipelinestat_block_size);
1290 if (flags & VK_QUERY_RESULT_64_BIT) {
1291 uint64_t *dst = (uint64_t *)dest;
1292 dest += util_bitcount(pool->pipeline_stats_mask) * 8;
1293 for (int i = 0; i < ARRAY_SIZE(pipeline_statistics_indices); ++i) {
1294 if (pool->pipeline_stats_mask & (1u << i)) {
1295 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT)) {
1296 *dst = stop[pipeline_statistics_indices[i]] -
1297 start[pipeline_statistics_indices[i]];
1299 if (pool->uses_gds &&
1300 (1u << i) == VK_QUERY_PIPELINE_STATISTIC_GEOMETRY_SHADER_PRIMITIVES_BIT) {
1301 *dst += ngg_gds_result;
1309 uint32_t *dst = (uint32_t *)dest;
1310 dest += util_bitcount(pool->pipeline_stats_mask) * 4;
1311 for (int i = 0; i < ARRAY_SIZE(pipeline_statistics_indices); ++i) {
1312 if (pool->pipeline_stats_mask & (1u << i)) {
1313 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT)) {
1314 *dst = stop[pipeline_statistics_indices[i]] -
1315 start[pipeline_statistics_indices[i]];
1317 if (pool->uses_gds &&
1318 (1u << i) == VK_QUERY_PIPELINE_STATISTIC_GEOMETRY_SHADER_PRIMITIVES_BIT) {
1319 *dst += ngg_gds_result;
1328 case VK_QUERY_TYPE_TRANSFORM_FEEDBACK_STREAM_EXT: {
1329 uint64_t const *src64 = (uint64_t const *)src;
1330 uint64_t num_primitives_written;
1331 uint64_t primitive_storage_needed;
1333 /* SAMPLE_STREAMOUTSTATS stores this structure:
1335 * u64 NumPrimitivesWritten;
1336 * u64 PrimitiveStorageNeeded;
1340 for (int j = 0; j < 4; j++) {
1341 if (!(p_atomic_read(src64 + j) & 0x8000000000000000UL))
1345 if (!available && !(flags & VK_QUERY_RESULT_PARTIAL_BIT))
1346 result = VK_NOT_READY;
1348 num_primitives_written = src64[3] - src64[1];
1349 primitive_storage_needed = src64[2] - src64[0];
1351 if (flags & VK_QUERY_RESULT_64_BIT) {
1352 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT))
1353 *(uint64_t *)dest = num_primitives_written;
1355 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT))
1356 *(uint64_t *)dest = primitive_storage_needed;
1359 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT))
1360 *(uint32_t *)dest = num_primitives_written;
1362 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT))
1363 *(uint32_t *)dest = primitive_storage_needed;
1368 case VK_QUERY_TYPE_PRIMITIVES_GENERATED_EXT: {
1369 uint64_t const *src64 = (uint64_t const *)src;
1370 uint64_t primitive_storage_needed;
1372 /* SAMPLE_STREAMOUTSTATS stores this structure:
1374 * u64 NumPrimitivesWritten;
1375 * u64 PrimitiveStorageNeeded;
1379 if (!(p_atomic_read(src64 + 0) & 0x8000000000000000UL) ||
1380 !(p_atomic_read(src64 + 2) & 0x8000000000000000UL)) {
1384 if (!available && !(flags & VK_QUERY_RESULT_PARTIAL_BIT))
1385 result = VK_NOT_READY;
1387 primitive_storage_needed = src64[2] - src64[0];
1389 if (pool->uses_gds && device->physical_device->rad_info.gfx_level < GFX11) {
1390 uint32_t const *src32 = (uint32_t const *)src;
1392 /* Accumulate the result that was copied from GDS in case NGG shader has been used. */
1393 primitive_storage_needed += src32[9] - src32[8];
1396 if (flags & VK_QUERY_RESULT_64_BIT) {
1397 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT))
1398 *(uint64_t *)dest = primitive_storage_needed;
1401 if (available || (flags & VK_QUERY_RESULT_PARTIAL_BIT))
1402 *(uint32_t *)dest = primitive_storage_needed;
1407 case VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR: {
1408 struct radv_pc_query_pool *pc_pool = (struct radv_pc_query_pool *)pool;
1409 const uint64_t *src64 = (const uint64_t *)src;
1413 for (unsigned i = 0; i < pc_pool->num_passes; ++i)
1414 if (!p_atomic_read(src64 + pool->stride / 8 - i - 1))
1416 } while (!avail && (flags & VK_QUERY_RESULT_WAIT_BIT));
1420 radv_pc_get_results(pc_pool, src64, dest);
1421 dest += pc_pool->num_counters * sizeof(union VkPerformanceCounterResultKHR);
1425 unreachable("trying to get results of unhandled query type");
1428 if (flags & VK_QUERY_RESULT_WITH_AVAILABILITY_BIT) {
1429 if (flags & VK_QUERY_RESULT_64_BIT) {
1430 *(uint64_t *)dest = available;
1432 *(uint32_t *)dest = available;
1441 emit_query_flush(struct radv_cmd_buffer *cmd_buffer, struct radv_query_pool *pool)
1443 if (cmd_buffer->pending_reset_query) {
1444 if (pool->size >= RADV_BUFFER_OPS_CS_THRESHOLD) {
1445 /* Only need to flush caches if the query pool size is
1446 * large enough to be resetted using the compute shader
1447 * path. Small pools don't need any cache flushes
1448 * because we use a CP dma clear.
1450 si_emit_cache_flush(cmd_buffer);
1456 radv_query_result_size(const struct radv_query_pool *pool, VkQueryResultFlags flags)
1458 unsigned values = (flags & VK_QUERY_RESULT_WITH_AVAILABILITY_BIT) ? 1 : 0;
1459 switch (pool->type) {
1460 case VK_QUERY_TYPE_TIMESTAMP:
1461 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_COMPACTED_SIZE_KHR:
1462 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_SIZE_KHR:
1463 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_BOTTOM_LEVEL_POINTERS_KHR:
1464 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SIZE_KHR:
1465 case VK_QUERY_TYPE_OCCLUSION:
1468 case VK_QUERY_TYPE_PIPELINE_STATISTICS:
1469 values += util_bitcount(pool->pipeline_stats_mask);
1471 case VK_QUERY_TYPE_TRANSFORM_FEEDBACK_STREAM_EXT:
1474 case VK_QUERY_TYPE_PRIMITIVES_GENERATED_EXT:
1478 unreachable("trying to get size of unhandled query type");
1480 return values * ((flags & VK_QUERY_RESULT_64_BIT) ? 8 : 4);
1483 VKAPI_ATTR void VKAPI_CALL
1484 radv_CmdCopyQueryPoolResults(VkCommandBuffer commandBuffer, VkQueryPool queryPool,
1485 uint32_t firstQuery, uint32_t queryCount, VkBuffer dstBuffer,
1486 VkDeviceSize dstOffset, VkDeviceSize stride, VkQueryResultFlags flags)
1488 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
1489 RADV_FROM_HANDLE(radv_query_pool, pool, queryPool);
1490 RADV_FROM_HANDLE(radv_buffer, dst_buffer, dstBuffer);
1491 struct radeon_cmdbuf *cs = cmd_buffer->cs;
1492 uint64_t va = radv_buffer_get_va(pool->bo);
1493 uint64_t dest_va = radv_buffer_get_va(dst_buffer->bo);
1494 size_t dst_size = radv_query_result_size(pool, flags);
1495 dest_va += dst_buffer->offset + dstOffset;
1500 radv_cs_add_buffer(cmd_buffer->device->ws, cmd_buffer->cs, pool->bo);
1501 radv_cs_add_buffer(cmd_buffer->device->ws, cmd_buffer->cs, dst_buffer->bo);
1503 /* Workaround engines that forget to properly specify WAIT_BIT because some driver implicitly
1504 * synchronizes before query copy.
1506 if (cmd_buffer->device->instance->flush_before_query_copy)
1507 cmd_buffer->state.flush_bits |= cmd_buffer->active_query_flush_bits;
1509 /* From the Vulkan spec 1.1.108:
1511 * "vkCmdCopyQueryPoolResults is guaranteed to see the effect of
1512 * previous uses of vkCmdResetQueryPool in the same queue, without any
1513 * additional synchronization."
1515 * So, we have to flush the caches if the compute shader path was used.
1517 emit_query_flush(cmd_buffer, pool);
1519 switch (pool->type) {
1520 case VK_QUERY_TYPE_OCCLUSION:
1521 if (flags & VK_QUERY_RESULT_WAIT_BIT) {
1522 unsigned enabled_rb_mask = cmd_buffer->device->physical_device->rad_info.enabled_rb_mask;
1523 uint32_t rb_avail_offset = 16 * util_last_bit(enabled_rb_mask) - 4;
1524 for (unsigned i = 0; i < queryCount; ++i, dest_va += stride) {
1525 unsigned query = firstQuery + i;
1526 uint64_t src_va = va + query * pool->stride + rb_avail_offset;
1528 radeon_check_space(cmd_buffer->device->ws, cs, 7);
1530 /* Waits on the upper word of the last DB entry */
1531 radv_cp_wait_mem(cs, WAIT_REG_MEM_GREATER_OR_EQUAL, src_va, 0x80000000, 0xffffffff);
1534 radv_query_shader(cmd_buffer, &cmd_buffer->device->meta_state.query.occlusion_query_pipeline,
1535 pool->bo, dst_buffer->bo, firstQuery * pool->stride,
1536 dst_buffer->offset + dstOffset, pool->stride, stride, dst_size, queryCount,
1537 flags, 0, 0, false);
1539 case VK_QUERY_TYPE_PIPELINE_STATISTICS:
1540 if (flags & VK_QUERY_RESULT_WAIT_BIT) {
1541 for (unsigned i = 0; i < queryCount; ++i, dest_va += stride) {
1542 unsigned query = firstQuery + i;
1544 radeon_check_space(cmd_buffer->device->ws, cs, 7);
1546 uint64_t avail_va = va + pool->availability_offset + 4 * query;
1548 /* This waits on the ME. All copies below are done on the ME */
1549 radv_cp_wait_mem(cs, WAIT_REG_MEM_EQUAL, avail_va, 1, 0xffffffff);
1553 cmd_buffer, &cmd_buffer->device->meta_state.query.pipeline_statistics_query_pipeline,
1554 pool->bo, dst_buffer->bo, firstQuery * pool->stride, dst_buffer->offset + dstOffset,
1555 pool->stride, stride, dst_size, queryCount, flags, pool->pipeline_stats_mask,
1556 pool->availability_offset + 4 * firstQuery, pool->uses_gds);
1558 case VK_QUERY_TYPE_TIMESTAMP:
1559 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_COMPACTED_SIZE_KHR:
1560 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_SIZE_KHR:
1561 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_BOTTOM_LEVEL_POINTERS_KHR:
1562 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SIZE_KHR:
1563 if (flags & VK_QUERY_RESULT_WAIT_BIT) {
1564 for (unsigned i = 0; i < queryCount; ++i, dest_va += stride) {
1565 unsigned query = firstQuery + i;
1566 uint64_t local_src_va = va + query * pool->stride;
1568 radeon_check_space(cmd_buffer->device->ws, cs, 7);
1570 /* Wait on the high 32 bits of the timestamp in
1571 * case the low part is 0xffffffff.
1573 radv_cp_wait_mem(cs, WAIT_REG_MEM_NOT_EQUAL, local_src_va + 4,
1574 TIMESTAMP_NOT_READY >> 32, 0xffffffff);
1578 radv_query_shader(cmd_buffer, &cmd_buffer->device->meta_state.query.timestamp_query_pipeline,
1579 pool->bo, dst_buffer->bo, firstQuery * pool->stride,
1580 dst_buffer->offset + dstOffset, pool->stride, stride, dst_size, queryCount,
1581 flags, 0, 0, false);
1583 case VK_QUERY_TYPE_TRANSFORM_FEEDBACK_STREAM_EXT:
1584 if (flags & VK_QUERY_RESULT_WAIT_BIT) {
1585 for (unsigned i = 0; i < queryCount; i++) {
1586 unsigned query = firstQuery + i;
1587 uint64_t src_va = va + query * pool->stride;
1589 radeon_check_space(cmd_buffer->device->ws, cs, 7 * 4);
1591 /* Wait on the upper word of all results. */
1592 for (unsigned j = 0; j < 4; j++, src_va += 8) {
1593 radv_cp_wait_mem(cs, WAIT_REG_MEM_GREATER_OR_EQUAL, src_va + 4, 0x80000000,
1599 radv_query_shader(cmd_buffer, &cmd_buffer->device->meta_state.query.tfb_query_pipeline,
1600 pool->bo, dst_buffer->bo, firstQuery * pool->stride,
1601 dst_buffer->offset + dstOffset, pool->stride, stride, dst_size, queryCount,
1602 flags, 0, 0, false);
1604 case VK_QUERY_TYPE_PRIMITIVES_GENERATED_EXT:
1605 if (flags & VK_QUERY_RESULT_WAIT_BIT) {
1606 for (unsigned i = 0; i < queryCount; i++) {
1607 unsigned query = firstQuery + i;
1608 uint64_t src_va = va + query * pool->stride;
1610 radeon_check_space(cmd_buffer->device->ws, cs, 7 * 2);
1612 /* Wait on the upper word of the PrimitiveStorageNeeded result. */
1613 radv_cp_wait_mem(cs, WAIT_REG_MEM_GREATER_OR_EQUAL, src_va + 4, 0x80000000, 0xffffffff);
1614 radv_cp_wait_mem(cs, WAIT_REG_MEM_GREATER_OR_EQUAL, src_va + 20, 0x80000000, 0xffffffff);
1618 radv_query_shader(cmd_buffer, &cmd_buffer->device->meta_state.query.pg_query_pipeline,
1619 pool->bo, dst_buffer->bo, firstQuery * pool->stride,
1620 dst_buffer->offset + dstOffset, pool->stride, stride, dst_size, queryCount,
1621 flags, 0, 0, pool->uses_gds && cmd_buffer->device->physical_device->rad_info.gfx_level < GFX11);
1624 unreachable("trying to get results of unhandled query type");
1629 query_clear_value(VkQueryType type)
1632 case VK_QUERY_TYPE_TIMESTAMP:
1633 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_COMPACTED_SIZE_KHR:
1634 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_SIZE_KHR:
1635 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_BOTTOM_LEVEL_POINTERS_KHR:
1636 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SIZE_KHR:
1637 return (uint32_t)TIMESTAMP_NOT_READY;
1643 VKAPI_ATTR void VKAPI_CALL
1644 radv_CmdResetQueryPool(VkCommandBuffer commandBuffer, VkQueryPool queryPool, uint32_t firstQuery,
1645 uint32_t queryCount)
1647 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
1648 RADV_FROM_HANDLE(radv_query_pool, pool, queryPool);
1649 uint32_t value = query_clear_value(pool->type);
1650 uint32_t flush_bits = 0;
1652 /* Make sure to sync all previous work if the given command buffer has
1653 * pending active queries. Otherwise the GPU might write queries data
1654 * after the reset operation.
1656 cmd_buffer->state.flush_bits |= cmd_buffer->active_query_flush_bits;
1658 flush_bits |= radv_fill_buffer(cmd_buffer, NULL, pool->bo,
1659 radv_buffer_get_va(pool->bo) + firstQuery * pool->stride,
1660 queryCount * pool->stride, value);
1662 if (pool->type == VK_QUERY_TYPE_PIPELINE_STATISTICS) {
1664 radv_fill_buffer(cmd_buffer, NULL, pool->bo,
1665 radv_buffer_get_va(pool->bo) + pool->availability_offset + firstQuery * 4,
1670 /* Only need to flush caches for the compute shader path. */
1671 cmd_buffer->pending_reset_query = true;
1672 cmd_buffer->state.flush_bits |= flush_bits;
1676 VKAPI_ATTR void VKAPI_CALL
1677 radv_ResetQueryPool(VkDevice _device, VkQueryPool queryPool, uint32_t firstQuery,
1678 uint32_t queryCount)
1680 RADV_FROM_HANDLE(radv_query_pool, pool, queryPool);
1682 uint32_t value = query_clear_value(pool->type);
1683 uint32_t *data = (uint32_t *)(pool->ptr + firstQuery * pool->stride);
1684 uint32_t *data_end = (uint32_t *)(pool->ptr + (firstQuery + queryCount) * pool->stride);
1686 for (uint32_t *p = data; p != data_end; ++p)
1689 if (pool->type == VK_QUERY_TYPE_PIPELINE_STATISTICS) {
1690 memset(pool->ptr + pool->availability_offset + firstQuery * 4, 0, queryCount * 4);
1695 event_type_for_stream(unsigned stream)
1700 return V_028A90_SAMPLE_STREAMOUTSTATS;
1702 return V_028A90_SAMPLE_STREAMOUTSTATS1;
1704 return V_028A90_SAMPLE_STREAMOUTSTATS2;
1706 return V_028A90_SAMPLE_STREAMOUTSTATS3;
1711 emit_sample_streamout(struct radv_cmd_buffer *cmd_buffer, uint64_t va, uint32_t index)
1713 struct radeon_cmdbuf *cs = cmd_buffer->cs;
1715 radeon_check_space(cmd_buffer->device->ws, cs, 4);
1717 assert(index < MAX_SO_STREAMS);
1719 radeon_emit(cs, PKT3(PKT3_EVENT_WRITE, 2, 0));
1720 radeon_emit(cs, EVENT_TYPE(event_type_for_stream(index)) | EVENT_INDEX(3));
1721 radeon_emit(cs, va);
1722 radeon_emit(cs, va >> 32);
1726 gfx10_copy_gds_query(struct radv_cmd_buffer *cmd_buffer, uint32_t gds_offset, uint64_t va)
1728 struct radeon_cmdbuf *cs = cmd_buffer->cs;
1730 /* Make sure GDS is idle before copying the value. */
1731 cmd_buffer->state.flush_bits |= RADV_CMD_FLAG_PS_PARTIAL_FLUSH | RADV_CMD_FLAG_INV_L2;
1732 si_emit_cache_flush(cmd_buffer);
1734 radeon_emit(cs, PKT3(PKT3_COPY_DATA, 4, 0));
1735 radeon_emit(cs, COPY_DATA_SRC_SEL(COPY_DATA_GDS) | COPY_DATA_DST_SEL(COPY_DATA_DST_MEM) |
1736 COPY_DATA_WR_CONFIRM);
1737 radeon_emit(cs, gds_offset);
1739 radeon_emit(cs, va);
1740 radeon_emit(cs, va >> 32);
1744 emit_begin_query(struct radv_cmd_buffer *cmd_buffer, struct radv_query_pool *pool, uint64_t va,
1745 VkQueryType query_type, VkQueryControlFlags flags, uint32_t index)
1747 struct radeon_cmdbuf *cs = cmd_buffer->cs;
1748 switch (query_type) {
1749 case VK_QUERY_TYPE_OCCLUSION:
1750 radeon_check_space(cmd_buffer->device->ws, cs, 11);
1752 ++cmd_buffer->state.active_occlusion_queries;
1753 if (cmd_buffer->state.active_occlusion_queries == 1) {
1754 if (flags & VK_QUERY_CONTROL_PRECISE_BIT) {
1755 /* This is the first occlusion query, enable
1756 * the hint if the precision bit is set.
1758 cmd_buffer->state.perfect_occlusion_queries_enabled = true;
1761 radv_set_db_count_control(cmd_buffer, true);
1763 if ((flags & VK_QUERY_CONTROL_PRECISE_BIT) &&
1764 !cmd_buffer->state.perfect_occlusion_queries_enabled) {
1765 /* This is not the first query, but this one
1766 * needs to enable precision, DB_COUNT_CONTROL
1767 * has to be updated accordingly.
1769 cmd_buffer->state.perfect_occlusion_queries_enabled = true;
1771 radv_set_db_count_control(cmd_buffer, true);
1775 if (cmd_buffer->device->physical_device->rad_info.gfx_level >= GFX11) {
1777 BITFIELD64_MASK(cmd_buffer->device->physical_device->rad_info.max_render_backends);
1779 radeon_emit(cs, PKT3(PKT3_EVENT_WRITE, 2, 0));
1780 radeon_emit(cs, EVENT_TYPE(V_028A90_PIXEL_PIPE_STAT_CONTROL) | EVENT_INDEX(1));
1781 radeon_emit(cs, PIXEL_PIPE_STATE_CNTL_COUNTER_ID(0) |
1782 PIXEL_PIPE_STATE_CNTL_STRIDE(2) |
1783 PIXEL_PIPE_STATE_CNTL_INSTANCE_EN_LO(rb_mask));
1784 radeon_emit(cs, PIXEL_PIPE_STATE_CNTL_INSTANCE_EN_HI(rb_mask));
1787 radeon_emit(cs, PKT3(PKT3_EVENT_WRITE, 2, 0));
1789 if (cmd_buffer->device->physical_device->rad_info.gfx_level >= GFX11) {
1790 radeon_emit(cs, EVENT_TYPE(V_028A90_PIXEL_PIPE_STAT_DUMP) | EVENT_INDEX(1));
1792 radeon_emit(cs, EVENT_TYPE(V_028A90_ZPASS_DONE) | EVENT_INDEX(1));
1795 radeon_emit(cs, va);
1796 radeon_emit(cs, va >> 32);
1798 case VK_QUERY_TYPE_PIPELINE_STATISTICS: {
1799 unsigned pipelinestat_block_size = radv_get_pipelinestat_query_size(cmd_buffer->device);
1801 radeon_check_space(cmd_buffer->device->ws, cs, 4);
1803 ++cmd_buffer->state.active_pipeline_queries;
1804 if (cmd_buffer->state.active_pipeline_queries == 1) {
1805 cmd_buffer->state.flush_bits &= ~RADV_CMD_FLAG_STOP_PIPELINE_STATS;
1806 cmd_buffer->state.flush_bits |= RADV_CMD_FLAG_START_PIPELINE_STATS;
1809 radeon_emit(cs, PKT3(PKT3_EVENT_WRITE, 2, 0));
1810 radeon_emit(cs, EVENT_TYPE(V_028A90_SAMPLE_PIPELINESTAT) | EVENT_INDEX(2));
1811 radeon_emit(cs, va);
1812 radeon_emit(cs, va >> 32);
1814 if (pool->uses_gds) {
1815 va += pipelinestat_block_size * 2;
1817 /* pipeline statistics counter for all streams */
1818 gfx10_copy_gds_query(cmd_buffer, RADV_NGG_QUERY_PIPELINE_STAT_OFFSET, va);
1820 /* Record that the command buffer needs GDS. */
1821 cmd_buffer->gds_needed = true;
1823 cmd_buffer->state.active_pipeline_gds_queries++;
1827 case VK_QUERY_TYPE_TRANSFORM_FEEDBACK_STREAM_EXT:
1828 if (cmd_buffer->device->physical_device->use_ngg_streamout) {
1829 /* generated prim counter */
1830 gfx10_copy_gds_query(cmd_buffer, RADV_NGG_QUERY_PRIM_GEN_OFFSET(index), va);
1831 radv_emit_write_data_imm(cs, V_370_ME, va + 4, 0x80000000);
1833 /* written prim counter */
1834 gfx10_copy_gds_query(cmd_buffer, RADV_NGG_QUERY_PRIM_XFB_OFFSET(index), va + 8);
1835 radv_emit_write_data_imm(cs, V_370_ME, va + 12, 0x80000000);
1837 cmd_buffer->state.active_prims_xfb_gds_queries++;
1839 emit_sample_streamout(cmd_buffer, va, index);
1842 case VK_QUERY_TYPE_PRIMITIVES_GENERATED_EXT: {
1843 if (cmd_buffer->device->physical_device->rad_info.gfx_level >= GFX11) {
1844 /* On GFX11+, primitives generated query always use GDS. */
1845 gfx10_copy_gds_query(cmd_buffer, RADV_NGG_QUERY_PRIM_GEN_OFFSET(index), va);
1846 radv_emit_write_data_imm(cs, V_370_ME, va + 4, 0x80000000);
1848 /* Record that the command buffer needs GDS. */
1849 cmd_buffer->gds_needed = true;
1851 cmd_buffer->state.active_prims_gen_gds_queries++;
1853 if (!cmd_buffer->state.active_prims_gen_queries) {
1854 bool old_streamout_enabled = radv_is_streamout_enabled(cmd_buffer);
1856 cmd_buffer->state.active_prims_gen_queries++;
1858 if (old_streamout_enabled != radv_is_streamout_enabled(cmd_buffer)) {
1859 radv_emit_streamout_enable(cmd_buffer);
1862 cmd_buffer->state.active_prims_gen_queries++;
1865 emit_sample_streamout(cmd_buffer, va, index);
1867 if (pool->uses_gds) {
1868 /* generated prim counter */
1869 gfx10_copy_gds_query(cmd_buffer, RADV_NGG_QUERY_PRIM_GEN_OFFSET(index), va + 32);
1871 /* Record that the command buffer needs GDS. */
1872 cmd_buffer->gds_needed = true;
1874 cmd_buffer->state.active_prims_gen_gds_queries++;
1879 case VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR: {
1880 radv_pc_begin_query(cmd_buffer, (struct radv_pc_query_pool *)pool, va);
1884 unreachable("beginning unhandled query type");
1889 emit_end_query(struct radv_cmd_buffer *cmd_buffer, struct radv_query_pool *pool, uint64_t va,
1890 uint64_t avail_va, VkQueryType query_type, uint32_t index)
1892 struct radeon_cmdbuf *cs = cmd_buffer->cs;
1893 switch (query_type) {
1894 case VK_QUERY_TYPE_OCCLUSION:
1895 radeon_check_space(cmd_buffer->device->ws, cs, 14);
1897 cmd_buffer->state.active_occlusion_queries--;
1898 if (cmd_buffer->state.active_occlusion_queries == 0) {
1899 radv_set_db_count_control(cmd_buffer, false);
1901 /* Reset the perfect occlusion queries hint now that no
1902 * queries are active.
1904 cmd_buffer->state.perfect_occlusion_queries_enabled = false;
1907 radeon_emit(cs, PKT3(PKT3_EVENT_WRITE, 2, 0));
1908 if (cmd_buffer->device->physical_device->rad_info.gfx_level >= GFX11) {
1909 radeon_emit(cs, EVENT_TYPE(V_028A90_PIXEL_PIPE_STAT_DUMP) | EVENT_INDEX(1));
1911 radeon_emit(cs, EVENT_TYPE(V_028A90_ZPASS_DONE) | EVENT_INDEX(1));
1913 radeon_emit(cs, va + 8);
1914 radeon_emit(cs, (va + 8) >> 32);
1917 case VK_QUERY_TYPE_PIPELINE_STATISTICS: {
1918 unsigned pipelinestat_block_size = radv_get_pipelinestat_query_size(cmd_buffer->device);
1920 radeon_check_space(cmd_buffer->device->ws, cs, 16);
1922 cmd_buffer->state.active_pipeline_queries--;
1923 if (cmd_buffer->state.active_pipeline_queries == 0) {
1924 cmd_buffer->state.flush_bits &= ~RADV_CMD_FLAG_START_PIPELINE_STATS;
1925 cmd_buffer->state.flush_bits |= RADV_CMD_FLAG_STOP_PIPELINE_STATS;
1927 va += pipelinestat_block_size;
1929 radeon_emit(cs, PKT3(PKT3_EVENT_WRITE, 2, 0));
1930 radeon_emit(cs, EVENT_TYPE(V_028A90_SAMPLE_PIPELINESTAT) | EVENT_INDEX(2));
1931 radeon_emit(cs, va);
1932 radeon_emit(cs, va >> 32);
1934 si_cs_emit_write_event_eop(cs, cmd_buffer->device->physical_device->rad_info.gfx_level,
1935 radv_cmd_buffer_uses_mec(cmd_buffer), V_028A90_BOTTOM_OF_PIPE_TS,
1936 0, EOP_DST_SEL_MEM, EOP_DATA_SEL_VALUE_32BIT, avail_va, 1,
1937 cmd_buffer->gfx9_eop_bug_va);
1939 if (pool->uses_gds) {
1940 va += pipelinestat_block_size + 8;
1942 /* pipeline statistics counter for all streams */
1943 gfx10_copy_gds_query(cmd_buffer, RADV_NGG_QUERY_PIPELINE_STAT_OFFSET, va);
1945 cmd_buffer->state.active_pipeline_gds_queries--;
1949 case VK_QUERY_TYPE_TRANSFORM_FEEDBACK_STREAM_EXT:
1950 if (cmd_buffer->device->physical_device->use_ngg_streamout) {
1951 /* generated prim counter */
1952 gfx10_copy_gds_query(cmd_buffer, RADV_NGG_QUERY_PRIM_GEN_OFFSET(index), va + 16);
1953 radv_emit_write_data_imm(cs, V_370_ME, va + 20, 0x80000000);
1955 /* written prim counter */
1956 gfx10_copy_gds_query(cmd_buffer, RADV_NGG_QUERY_PRIM_XFB_OFFSET(index), va + 24);
1957 radv_emit_write_data_imm(cs, V_370_ME, va + 28, 0x80000000);
1959 cmd_buffer->state.active_prims_xfb_gds_queries--;
1961 emit_sample_streamout(cmd_buffer, va + 16, index);
1964 case VK_QUERY_TYPE_PRIMITIVES_GENERATED_EXT: {
1965 if (cmd_buffer->device->physical_device->rad_info.gfx_level >= GFX11) {
1966 /* On GFX11+, primitives generated query always use GDS. */
1967 gfx10_copy_gds_query(cmd_buffer, RADV_NGG_QUERY_PRIM_GEN_OFFSET(index), va + 16);
1968 radv_emit_write_data_imm(cs, V_370_ME, va + 20, 0x80000000);
1970 cmd_buffer->state.active_prims_gen_gds_queries--;
1972 if (cmd_buffer->state.active_prims_gen_queries == 1) {
1973 bool old_streamout_enabled = radv_is_streamout_enabled(cmd_buffer);
1975 cmd_buffer->state.active_prims_gen_queries--;
1977 if (old_streamout_enabled != radv_is_streamout_enabled(cmd_buffer)) {
1978 radv_emit_streamout_enable(cmd_buffer);
1981 cmd_buffer->state.active_prims_gen_queries--;
1984 emit_sample_streamout(cmd_buffer, va + 16, index);
1986 if (pool->uses_gds) {
1987 /* generated prim counter */
1988 gfx10_copy_gds_query(cmd_buffer, RADV_NGG_QUERY_PRIM_GEN_OFFSET(index), va + 36);
1990 cmd_buffer->state.active_prims_gen_gds_queries--;
1995 case VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR: {
1996 radv_pc_end_query(cmd_buffer, (struct radv_pc_query_pool *)pool, va);
2000 unreachable("ending unhandled query type");
2003 cmd_buffer->active_query_flush_bits |= RADV_CMD_FLAG_PS_PARTIAL_FLUSH |
2004 RADV_CMD_FLAG_CS_PARTIAL_FLUSH | RADV_CMD_FLAG_INV_L2 |
2005 RADV_CMD_FLAG_INV_VCACHE;
2006 if (cmd_buffer->device->physical_device->rad_info.gfx_level >= GFX9) {
2007 cmd_buffer->active_query_flush_bits |=
2008 RADV_CMD_FLAG_FLUSH_AND_INV_CB | RADV_CMD_FLAG_FLUSH_AND_INV_DB;
2012 VKAPI_ATTR void VKAPI_CALL
2013 radv_CmdBeginQueryIndexedEXT(VkCommandBuffer commandBuffer, VkQueryPool queryPool, uint32_t query,
2014 VkQueryControlFlags flags, uint32_t index)
2016 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
2017 RADV_FROM_HANDLE(radv_query_pool, pool, queryPool);
2018 struct radeon_cmdbuf *cs = cmd_buffer->cs;
2019 uint64_t va = radv_buffer_get_va(pool->bo);
2021 radv_cs_add_buffer(cmd_buffer->device->ws, cs, pool->bo);
2023 emit_query_flush(cmd_buffer, pool);
2025 va += pool->stride * query;
2027 emit_begin_query(cmd_buffer, pool, va, pool->type, flags, index);
2030 VKAPI_ATTR void VKAPI_CALL
2031 radv_CmdBeginQuery(VkCommandBuffer commandBuffer, VkQueryPool queryPool, uint32_t query,
2032 VkQueryControlFlags flags)
2034 radv_CmdBeginQueryIndexedEXT(commandBuffer, queryPool, query, flags, 0);
2037 VKAPI_ATTR void VKAPI_CALL
2038 radv_CmdEndQueryIndexedEXT(VkCommandBuffer commandBuffer, VkQueryPool queryPool, uint32_t query,
2041 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
2042 RADV_FROM_HANDLE(radv_query_pool, pool, queryPool);
2043 uint64_t va = radv_buffer_get_va(pool->bo);
2044 uint64_t avail_va = va + pool->availability_offset + 4 * query;
2045 va += pool->stride * query;
2047 /* Do not need to add the pool BO to the list because the query must
2048 * currently be active, which means the BO is already in the list.
2050 emit_end_query(cmd_buffer, pool, va, avail_va, pool->type, index);
2053 * For multiview we have to emit a query for each bit in the mask,
2054 * however the first query we emit will get the totals for all the
2055 * operations, so we don't want to get a real value in the other
2056 * queries. This emits a fake begin/end sequence so the waiting
2057 * code gets a completed query value and doesn't hang, but the
2060 if (cmd_buffer->state.render.view_mask) {
2061 for (unsigned i = 1; i < util_bitcount(cmd_buffer->state.render.view_mask); i++) {
2064 emit_begin_query(cmd_buffer, pool, va, pool->type, 0, 0);
2065 emit_end_query(cmd_buffer, pool, va, avail_va, pool->type, 0);
2070 VKAPI_ATTR void VKAPI_CALL
2071 radv_CmdEndQuery(VkCommandBuffer commandBuffer, VkQueryPool queryPool, uint32_t query)
2073 radv_CmdEndQueryIndexedEXT(commandBuffer, queryPool, query, 0);
2076 VKAPI_ATTR void VKAPI_CALL
2077 radv_CmdWriteTimestamp2(VkCommandBuffer commandBuffer, VkPipelineStageFlags2 stage,
2078 VkQueryPool queryPool, uint32_t query)
2080 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
2081 RADV_FROM_HANDLE(radv_query_pool, pool, queryPool);
2082 bool mec = radv_cmd_buffer_uses_mec(cmd_buffer);
2083 struct radeon_cmdbuf *cs = cmd_buffer->cs;
2084 uint64_t va = radv_buffer_get_va(pool->bo);
2085 uint64_t query_va = va + pool->stride * query;
2087 radv_cs_add_buffer(cmd_buffer->device->ws, cs, pool->bo);
2089 emit_query_flush(cmd_buffer, pool);
2091 int num_queries = 1;
2092 if (cmd_buffer->state.render.view_mask)
2093 num_queries = util_bitcount(cmd_buffer->state.render.view_mask);
2095 ASSERTED unsigned cdw_max = radeon_check_space(cmd_buffer->device->ws, cs, 28 * num_queries);
2097 for (unsigned i = 0; i < num_queries; i++) {
2098 if (stage == VK_PIPELINE_STAGE_2_TOP_OF_PIPE_BIT) {
2099 radeon_emit(cs, PKT3(PKT3_COPY_DATA, 4, 0));
2100 radeon_emit(cs, COPY_DATA_COUNT_SEL | COPY_DATA_WR_CONFIRM |
2101 COPY_DATA_SRC_SEL(COPY_DATA_TIMESTAMP) | COPY_DATA_DST_SEL(V_370_MEM));
2104 radeon_emit(cs, query_va);
2105 radeon_emit(cs, query_va >> 32);
2107 si_cs_emit_write_event_eop(cs, cmd_buffer->device->physical_device->rad_info.gfx_level,
2108 mec, V_028A90_BOTTOM_OF_PIPE_TS, 0, EOP_DST_SEL_MEM,
2109 EOP_DATA_SEL_TIMESTAMP, query_va, 0,
2110 cmd_buffer->gfx9_eop_bug_va);
2112 query_va += pool->stride;
2115 cmd_buffer->active_query_flush_bits |= RADV_CMD_FLAG_PS_PARTIAL_FLUSH |
2116 RADV_CMD_FLAG_CS_PARTIAL_FLUSH | RADV_CMD_FLAG_INV_L2 |
2117 RADV_CMD_FLAG_INV_VCACHE;
2118 if (cmd_buffer->device->physical_device->rad_info.gfx_level >= GFX9) {
2119 cmd_buffer->active_query_flush_bits |=
2120 RADV_CMD_FLAG_FLUSH_AND_INV_CB | RADV_CMD_FLAG_FLUSH_AND_INV_DB;
2123 assert(cmd_buffer->cs->cdw <= cdw_max);
2126 VKAPI_ATTR void VKAPI_CALL
2127 radv_CmdWriteAccelerationStructuresPropertiesKHR(
2128 VkCommandBuffer commandBuffer, uint32_t accelerationStructureCount,
2129 const VkAccelerationStructureKHR *pAccelerationStructures, VkQueryType queryType,
2130 VkQueryPool queryPool, uint32_t firstQuery)
2132 RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);
2133 RADV_FROM_HANDLE(radv_query_pool, pool, queryPool);
2134 struct radeon_cmdbuf *cs = cmd_buffer->cs;
2135 uint64_t pool_va = radv_buffer_get_va(pool->bo);
2136 uint64_t query_va = pool_va + pool->stride * firstQuery;
2138 radv_cs_add_buffer(cmd_buffer->device->ws, cs, pool->bo);
2140 emit_query_flush(cmd_buffer, pool);
2142 ASSERTED unsigned cdw_max =
2143 radeon_check_space(cmd_buffer->device->ws, cs, 6 * accelerationStructureCount);
2145 for (uint32_t i = 0; i < accelerationStructureCount; ++i) {
2146 RADV_FROM_HANDLE(radv_acceleration_structure, accel_struct, pAccelerationStructures[i]);
2147 uint64_t va = accel_struct->va;
2149 switch (queryType) {
2150 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_COMPACTED_SIZE_KHR:
2151 va += offsetof(struct radv_accel_struct_header, compacted_size);
2153 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_SIZE_KHR:
2154 va += offsetof(struct radv_accel_struct_header, serialization_size);
2156 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_BOTTOM_LEVEL_POINTERS_KHR:
2157 va += offsetof(struct radv_accel_struct_header, instance_count);
2159 case VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SIZE_KHR:
2160 va += offsetof(struct radv_accel_struct_header, size);
2163 unreachable("Unhandle accel struct query type.");
2166 radeon_emit(cs, PKT3(PKT3_COPY_DATA, 4, 0));
2167 radeon_emit(cs, COPY_DATA_SRC_SEL(COPY_DATA_SRC_MEM) | COPY_DATA_DST_SEL(COPY_DATA_DST_MEM) |
2168 COPY_DATA_COUNT_SEL | COPY_DATA_WR_CONFIRM);
2169 radeon_emit(cs, va);
2170 radeon_emit(cs, va >> 32);
2171 radeon_emit(cs, query_va);
2172 radeon_emit(cs, query_va >> 32);
2174 query_va += pool->stride;
2177 assert(cmd_buffer->cs->cdw <= cdw_max);