From 7a65b3f006ad1c77736dd9618c9e8cb2ed12d910 Mon Sep 17 00:00:00 2001 From: Iago Toral Quiroga Date: Fri, 28 Oct 2022 12:07:07 +0200 Subject: [PATCH] v3dv: reimplement occlusion queries MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Our implementation was mostly CPU-based, with things such as query resets and result copying handled in the CPU, as well as some aspects of query availability tracking. This new implementation handles all GPU-side query functions by dispatching compute shaders to push the work to the GPU. This involves query availability, reset and result copying. For now, only occlusion queries are managed this way. Performance queries can also be implemented in a similar fashion in the future with some additional work, however, for timestamp queries our only option to improve this would be to execute the actual timestamp in the kernel, since we can't take a timestamp from a shader. Reviewed-by: Alejandro Piñeiro Part-of: --- src/broadcom/vulkan/v3dv_cmd_buffer.c | 158 +-- src/broadcom/vulkan/v3dv_device.c | 8 + src/broadcom/vulkan/v3dv_private.h | 116 ++- src/broadcom/vulkan/v3dv_query.c | 1291 ++++++++++++++++++++++-- src/broadcom/vulkan/v3dv_queue.c | 26 +- src/broadcom/vulkan/v3dvx_cmd_buffer.c | 6 +- 6 files changed, 1379 insertions(+), 226 deletions(-) diff --git a/src/broadcom/vulkan/v3dv_cmd_buffer.c b/src/broadcom/vulkan/v3dv_cmd_buffer.c index ffa66689665..7d5aaf19bae 100644 --- a/src/broadcom/vulkan/v3dv_cmd_buffer.c +++ b/src/broadcom/vulkan/v3dv_cmd_buffer.c @@ -565,22 +565,41 @@ v3dv_cmd_buffer_create_cpu_job(struct v3dv_device *device, } static void -cmd_buffer_add_cpu_jobs_for_pending_state(struct v3dv_cmd_buffer *cmd_buffer) +cmd_buffer_emit_end_query_cpu(struct v3dv_cmd_buffer *cmd_buffer, + struct v3dv_query_pool *pool, + uint32_t query, uint32_t count) +{ + assert(pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR); + + struct v3dv_job *job = + v3dv_cmd_buffer_create_cpu_job(cmd_buffer->device, + V3DV_JOB_TYPE_CPU_END_QUERY, + cmd_buffer, -1); + v3dv_return_if_oom(cmd_buffer, NULL); + + job->cpu.query_end.pool = pool; + job->cpu.query_end.query = query; + job->cpu.query_end.count = count; + list_addtail(&job->list_link, &cmd_buffer->jobs); +} + +static void +cmd_buffer_add_jobs_for_pending_state(struct v3dv_cmd_buffer *cmd_buffer) { struct v3dv_cmd_buffer_state *state = &cmd_buffer->state; if (state->query.end.used_count > 0) { - const uint32_t query_count = state->query.end.used_count; - for (uint32_t i = 0; i < query_count; i++) { + const uint32_t count = state->query.end.used_count; + for (uint32_t i = 0; i < count; i++) { assert(i < state->query.end.used_count); - struct v3dv_job *job = - v3dv_cmd_buffer_create_cpu_job(cmd_buffer->device, - V3DV_JOB_TYPE_CPU_END_QUERY, - cmd_buffer, -1); - v3dv_return_if_oom(cmd_buffer, NULL); - - job->cpu.query_end = state->query.end.states[i]; - list_addtail(&job->list_link, &cmd_buffer->jobs); + struct v3dv_end_query_info *info = &state->query.end.states[i]; + if (info->pool->query_type == VK_QUERY_TYPE_OCCLUSION) { + v3dv_cmd_buffer_emit_set_query_availability(cmd_buffer, info->pool, + info->query, info->count, 1); + } else { + cmd_buffer_emit_end_query_cpu(cmd_buffer, info->pool, + info->query, info->count); + } } state->query.end.used_count = 0; } @@ -650,14 +669,14 @@ v3dv_cmd_buffer_finish_job(struct v3dv_cmd_buffer *cmd_buffer) cmd_buffer->state.job = NULL; /* If we have recorded any state with this last GPU job that requires to - * emit CPU jobs after the job is completed, add them now. The only - * exception is secondary command buffers inside a render pass, because in + * emit jobs after the job is completed, add them now. The only exception + * is secondary command buffers inside a render pass, because in * that case we want to defer this until we finish recording the primary * job into which we execute the secondary. */ if (cmd_buffer->vk.level == VK_COMMAND_BUFFER_LEVEL_PRIMARY || !cmd_buffer->state.pass) { - cmd_buffer_add_cpu_jobs_for_pending_state(cmd_buffer); + cmd_buffer_add_jobs_for_pending_state(cmd_buffer); } } @@ -765,7 +784,7 @@ v3dv_job_init(struct v3dv_job *job, cmd_buffer->state.dirty = ~0; cmd_buffer->state.dirty_descriptor_stages = ~0; - /* Honor inheritance of occlussion queries in secondaries if requested */ + /* Honor inheritance of occlusion queries in secondaries if requested */ if (cmd_buffer->vk.level == VK_COMMAND_BUFFER_LEVEL_SECONDARY && cmd_buffer->state.inheritance.occlusion_query_enable) { cmd_buffer->state.dirty &= ~V3DV_CMD_DIRTY_OCCLUSION_QUERY; @@ -3495,34 +3514,6 @@ v3dv_CmdSetColorWriteEnableEXT(VkCommandBuffer commandBuffer, state->dirty |= V3DV_CMD_DIRTY_COLOR_WRITE_ENABLE; } -void -v3dv_cmd_buffer_reset_queries(struct v3dv_cmd_buffer *cmd_buffer, - struct v3dv_query_pool *pool, - uint32_t first, - uint32_t count) -{ - /* Resets can only happen outside a render pass instance so we should not - * be in the middle of job recording. - */ - assert(cmd_buffer->state.pass == NULL); - assert(cmd_buffer->state.job == NULL); - - assert(first < pool->query_count); - assert(first + count <= pool->query_count); - - struct v3dv_job *job = - v3dv_cmd_buffer_create_cpu_job(cmd_buffer->device, - V3DV_JOB_TYPE_CPU_RESET_QUERIES, - cmd_buffer, -1); - v3dv_return_if_oom(cmd_buffer, NULL); - - job->cpu.query_reset.pool = pool; - job->cpu.query_reset.first = first; - job->cpu.query_reset.count = count; - - list_addtail(&job->list_link, &cmd_buffer->jobs); -} - void v3dv_cmd_buffer_ensure_array_state(struct v3dv_cmd_buffer *cmd_buffer, uint32_t slot_size, @@ -3562,8 +3553,9 @@ v3dv_cmd_buffer_begin_query(struct v3dv_cmd_buffer *cmd_buffer, /* FIXME: we only support one active occlusion query for now */ assert(cmd_buffer->state.query.active_query.bo == NULL); - cmd_buffer->state.query.active_query.bo = pool->queries[query].bo; - cmd_buffer->state.query.active_query.offset = pool->queries[query].offset; + cmd_buffer->state.query.active_query.bo = pool->occlusion.bo; + cmd_buffer->state.query.active_query.offset = + pool->queries[query].occlusion.offset; cmd_buffer->state.dirty |= V3DV_CMD_DIRTY_OCCLUSION_QUERY; break; case VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR: { @@ -3591,22 +3583,25 @@ v3dv_cmd_buffer_schedule_end_query(struct v3dv_cmd_buffer *cmd_buffer, uint32_t query) { assert(query < pool->query_count); + assert(pool->query_type == VK_QUERY_TYPE_OCCLUSION || + pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR); + /* For occlusion queries in the middle of a render pass we don't want to + * split the current job at the EndQuery just to emit query availability, + * instead we queue this state in the command buffer and we emit it when + * we finish the current job. + */ if (cmd_buffer->state.pass && - pool->query_type != VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR) { - /* Queue the EndQuery in the command buffer state, we will create a CPU - * job to flag all of these queries as possibly available right after the - * render pass job in which they have been recorded. - */ + pool->query_type == VK_QUERY_TYPE_OCCLUSION) { struct v3dv_cmd_buffer_state *state = &cmd_buffer->state; v3dv_cmd_buffer_ensure_array_state(cmd_buffer, - sizeof(struct v3dv_end_query_cpu_job_info), + sizeof(struct v3dv_end_query_info), state->query.end.used_count, &state->query.end.alloc_count, (void **) &state->query.end.states); v3dv_return_if_oom(cmd_buffer, NULL); - struct v3dv_end_query_cpu_job_info *info = + struct v3dv_end_query_info *info = &state->query.end.states[state->query.end.used_count++]; info->pool = pool; @@ -3633,20 +3628,15 @@ v3dv_cmd_buffer_schedule_end_query(struct v3dv_cmd_buffer *cmd_buffer, info->count = util_bitcount(subpass->view_mask); } } else { - /* Otherwise, schedule the CPU job immediately */ - struct v3dv_job *job = - v3dv_cmd_buffer_create_cpu_job(cmd_buffer->device, - V3DV_JOB_TYPE_CPU_END_QUERY, - cmd_buffer, -1); - v3dv_return_if_oom(cmd_buffer, NULL); - - job->cpu.query_end.pool = pool; - job->cpu.query_end.query = query; - - /* Multiview queries cannot cross subpass boundaries */ - job->cpu.query_end.count = 1; - - list_addtail(&job->list_link, &cmd_buffer->jobs); + /* Otherwise, schedule the end query job immediately. + * + * Multiview queries cannot cross subpass boundaries, so query count is + * always 1. + */ + if (pool->query_type == VK_QUERY_TYPE_OCCLUSION) + v3dv_cmd_buffer_emit_set_query_availability(cmd_buffer, pool, query, 1, 1); + else + cmd_buffer_emit_end_query_cpu(cmd_buffer, pool, query, 1); } } @@ -3699,42 +3689,6 @@ void v3dv_cmd_buffer_end_query(struct v3dv_cmd_buffer *cmd_buffer, } } -void -v3dv_cmd_buffer_copy_query_results(struct v3dv_cmd_buffer *cmd_buffer, - struct v3dv_query_pool *pool, - uint32_t first, - uint32_t count, - struct v3dv_buffer *dst, - uint32_t offset, - uint32_t stride, - VkQueryResultFlags flags) -{ - /* Copies can only happen outside a render pass instance so we should not - * be in the middle of job recording. - */ - assert(cmd_buffer->state.pass == NULL); - assert(cmd_buffer->state.job == NULL); - - assert(first < pool->query_count); - assert(first + count <= pool->query_count); - - struct v3dv_job *job = - v3dv_cmd_buffer_create_cpu_job(cmd_buffer->device, - V3DV_JOB_TYPE_CPU_COPY_QUERY_RESULTS, - cmd_buffer, -1); - v3dv_return_if_oom(cmd_buffer, NULL); - - job->cpu.query_copy_results.pool = pool; - job->cpu.query_copy_results.first = first; - job->cpu.query_copy_results.count = count; - job->cpu.query_copy_results.dst = dst; - job->cpu.query_copy_results.offset = offset; - job->cpu.query_copy_results.stride = stride; - job->cpu.query_copy_results.flags = flags; - - list_addtail(&job->list_link, &cmd_buffer->jobs); -} - void v3dv_cmd_buffer_add_tfu_job(struct v3dv_cmd_buffer *cmd_buffer, struct drm_v3d_submit_tfu *tfu) diff --git a/src/broadcom/vulkan/v3dv_device.c b/src/broadcom/vulkan/v3dv_device.c index f1dae7b6c6a..ce4a2948e97 100644 --- a/src/broadcom/vulkan/v3dv_device.c +++ b/src/broadcom/vulkan/v3dv_device.c @@ -2113,6 +2113,10 @@ v3dv_CreateDevice(VkPhysicalDevice physicalDevice, goto fail; } + result = v3dv_query_allocate_resources(device); + if (result != VK_SUCCESS) + goto fail; + *pDevice = v3dv_device_to_handle(device); return VK_SUCCESS; @@ -2123,6 +2127,8 @@ fail: queue_finish(&device->queue); destroy_device_meta(device); v3dv_pipeline_cache_finish(&device->default_pipeline_cache); + v3dv_event_free_resources(device); + v3dv_query_free_resources(device); vk_device_finish(&device->vk); vk_free(&device->vk.alloc, device); @@ -2141,6 +2147,8 @@ v3dv_DestroyDevice(VkDevice _device, v3dv_event_free_resources(device); mtx_destroy(&device->events.lock); + v3dv_query_free_resources(device); + destroy_device_meta(device); v3dv_pipeline_cache_finish(&device->default_pipeline_cache); diff --git a/src/broadcom/vulkan/v3dv_private.h b/src/broadcom/vulkan/v3dv_private.h index 19318e4bb74..d05f0638a7e 100644 --- a/src/broadcom/vulkan/v3dv_private.h +++ b/src/broadcom/vulkan/v3dv_private.h @@ -556,6 +556,31 @@ struct v3dv_device { VkPipeline wait_event_pipeline; } events; + /* Query handling resources. + * + * Our implementation of occlusion queries uses a BO per pool to keep track + * of the per-query availability state and dispatches compute shaders to + * handle GPU query functions that read and write that state. This struct + * holds Vulkan resources that can be shared across all query pools to + * implement this. This framework may be extended in the future to handle + * more query types. + */ + struct { + VkDescriptorSetLayout buf_descriptor_set_layout; + + /* Set query availability */ + VkPipelineLayout avail_pipeline_layout; + VkPipeline avail_pipeline; + + /* Reset query availability and clear occlusion counters */ + VkPipelineLayout reset_occlusion_pipeline_layout; + VkPipeline reset_occlusion_pipeline; + + /* Copy query results */ + VkPipelineLayout copy_pipeline_layout; + VkPipeline copy_pipeline; + } queries; + struct v3dv_pipeline_cache default_pipeline_cache; /* GL_SHADER_STATE_RECORD needs to specify default attribute values. The @@ -1026,7 +1051,7 @@ struct v3dv_reset_query_cpu_job_info { uint32_t count; }; -struct v3dv_end_query_cpu_job_info { +struct v3dv_end_query_info { struct v3dv_query_pool *pool; uint32_t query; @@ -1218,7 +1243,7 @@ struct v3dv_job { /* Job specs for CPU jobs */ union { struct v3dv_reset_query_cpu_job_info query_reset; - struct v3dv_end_query_cpu_job_info query_end; + struct v3dv_end_query_info query_end; struct v3dv_copy_query_results_cpu_job_info query_copy_results; struct v3dv_copy_buffer_to_image_cpu_job_info copy_buffer_to_image; struct v3dv_csd_indirect_cpu_job_info csd_indirect; @@ -1449,7 +1474,7 @@ struct v3dv_cmd_buffer_state { struct { uint32_t used_count; uint32_t alloc_count; - struct v3dv_end_query_cpu_job_info *states; + struct v3dv_end_query_info *states; } end; struct { @@ -1498,13 +1523,19 @@ struct v3dv_descriptor { }; struct v3dv_query { + /* Used by queries where we implement result copying in the CPU so we can + * tell if the relevant jobs have been submitted for execution. Currently + * these are all but occlusion queries. + */ bool maybe_available; + union { - /* Used by GPU queries (occlusion) */ + /* Used by occlusion queries */ struct { - struct v3dv_bo *bo; + /* Offset of this query in the occlusion query counter BO */ uint32_t offset; - }; + } occlusion; + /* Used by CPU queries (timestamp) */ uint64_t value; @@ -1516,7 +1547,36 @@ struct v3dv_query { struct v3dv_query_pool { struct vk_object_base base; - struct v3dv_bo *bo; /* Only used with GPU queries (occlusion) */ + /* Availability state for each query in the pool. Only used with occlusion + * queries for now, but could be used by other query types in the future. + */ + struct v3dv_bo *avail_bo; + + /* Per-pool Vulkan resources required to implement GPU-side query + * functions (only occlusion queries for now). + */ + struct { + /* Buffer to access query availability state */ + VkBuffer avail_buf; + VkDeviceMemory avail_mem; + + /* Buffer to access occlusion query results */ + VkBuffer res_buf; + VkDeviceMemory res_mem; + + VkDescriptorPool descriptor_pool; + + /* Two descriptor sets: one for accessing the availability buffer and + * another for the buffer with the occlusion query results. + */ + VkDescriptorSet descriptor_sets[2]; + } meta; + + /* Only used with occlusion queries */ + struct { + /* BO with the occlusion counters */ + struct v3dv_bo *bo; + } occlusion; /* Only used with performance queries */ struct { @@ -1537,18 +1597,29 @@ struct v3dv_query_pool { struct v3dv_query *queries; }; -VkResult v3dv_get_query_pool_results(struct v3dv_device *device, - struct v3dv_query_pool *pool, - uint32_t first, - uint32_t count, - void *data, - VkDeviceSize stride, - VkQueryResultFlags flags); +VkResult +v3dv_query_allocate_resources(struct v3dv_device *decice); -void v3dv_reset_query_pools(struct v3dv_device *device, - struct v3dv_query_pool *query_pool, - uint32_t first, - uint32_t last); +void +v3dv_query_free_resources(struct v3dv_device *decice); + +VkResult v3dv_get_query_pool_results_cpu(struct v3dv_device *device, + struct v3dv_query_pool *pool, + uint32_t first, + uint32_t count, + void *data, + VkDeviceSize stride, + VkQueryResultFlags flags); + +void v3dv_reset_query_pool_cpu(struct v3dv_device *device, + struct v3dv_query_pool *query_pool, + uint32_t first, + uint32_t last); + +void v3dv_cmd_buffer_emit_set_query_availability(struct v3dv_cmd_buffer *cmd_buffer, + struct v3dv_query_pool *pool, + uint32_t query, uint32_t count, + uint8_t availability); typedef void (*v3dv_cmd_buffer_private_obj_destroy_cb)(VkDevice device, uint64_t pobj, @@ -1597,6 +1668,10 @@ struct v3dv_cmd_buffer { /* The current descriptor pool for texel buffer copy sources */ VkDescriptorPool dspool; } texel_buffer_copy; + struct { + /* The current descriptor pool for the copy query results output buffer */ + VkDescriptorPool dspool; + } query; } meta; /* List of jobs in the command buffer. For primary command buffers it @@ -1625,11 +1700,6 @@ void v3dv_cmd_buffer_meta_state_pop(struct v3dv_cmd_buffer *cmd_buffer, uint32_t dirty_dynamic_state, bool needs_subpass_resume); -void v3dv_cmd_buffer_reset_queries(struct v3dv_cmd_buffer *cmd_buffer, - struct v3dv_query_pool *pool, - uint32_t first, - uint32_t count); - void v3dv_cmd_buffer_begin_query(struct v3dv_cmd_buffer *cmd_buffer, struct v3dv_query_pool *pool, uint32_t query, diff --git a/src/broadcom/vulkan/v3dv_query.c b/src/broadcom/vulkan/v3dv_query.c index 1f8aa523491..1a6a2e87d39 100644 --- a/src/broadcom/vulkan/v3dv_query.c +++ b/src/broadcom/vulkan/v3dv_query.c @@ -24,6 +24,7 @@ #include "v3dv_private.h" #include "util/timespec.h" +#include "compiler/nir/nir_builder.h" static const char *v3dv_counters[][3] = { {"FEP", "FEP-valid-primitives-no-rendered-pixels", "[FEP] Valid primitives that result in no rendered pixels, for all rendered tiles"}, @@ -167,6 +168,198 @@ kperfmon_destroy(struct v3dv_device *device, } } +/** + * Creates a VkBuffer (and VkDeviceMemory) to access a BO. + */ +static VkResult +create_vk_storage_buffer(struct v3dv_device *device, + struct v3dv_bo *bo, + VkBuffer *vk_buf, + VkDeviceMemory *vk_mem) +{ + VkDevice vk_device = v3dv_device_to_handle(device); + + VkBufferCreateInfo buf_info = { + .sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, + .size = bo->size, + .usage = VK_BUFFER_USAGE_STORAGE_BUFFER_BIT, + }; + VkResult result = v3dv_CreateBuffer(vk_device, &buf_info, NULL, vk_buf); + if (result != VK_SUCCESS) + return result; + + struct v3dv_device_memory *mem = + vk_object_zalloc(&device->vk, NULL, sizeof(*mem), + VK_OBJECT_TYPE_DEVICE_MEMORY); + if (!mem) + return VK_ERROR_OUT_OF_HOST_MEMORY; + + mem->bo = bo; + mem->type = &device->pdevice->memory.memoryTypes[0]; + + *vk_mem = v3dv_device_memory_to_handle(mem); + VkBindBufferMemoryInfo bind_info = { + .sType = VK_STRUCTURE_TYPE_BIND_BUFFER_MEMORY_INFO, + .buffer = *vk_buf, + .memory = *vk_mem, + .memoryOffset = 0, + }; + v3dv_BindBufferMemory2(vk_device, 1, &bind_info); + + return VK_SUCCESS; +} + +static void +destroy_vk_storage_buffer(struct v3dv_device *device, + VkBuffer *vk_buf, + VkDeviceMemory *vk_mem) +{ + if (*vk_mem) { + vk_object_free(&device->vk, NULL, v3dv_device_memory_from_handle(*vk_mem)); + *vk_mem = VK_NULL_HANDLE; + } + + v3dv_DestroyBuffer(v3dv_device_to_handle(device), *vk_buf, NULL); + *vk_buf = VK_NULL_HANDLE; +} + +/** + * Allocates descriptor sets to access query pool BOs (availability and + * occlusion query results) from Vulkan pipelines. + */ +static VkResult +create_pool_descriptors(struct v3dv_device *device, + struct v3dv_query_pool *pool) +{ + assert(pool->query_type == VK_QUERY_TYPE_OCCLUSION); + VkDevice vk_device = v3dv_device_to_handle(device); + + VkDescriptorPoolSize pool_size = { + .type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, + .descriptorCount = 2, + }; + VkDescriptorPoolCreateInfo pool_info = { + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO, + .flags = VK_DESCRIPTOR_POOL_CREATE_FREE_DESCRIPTOR_SET_BIT, + .maxSets = 2, + .poolSizeCount = 1, + .pPoolSizes = &pool_size, + }; + VkResult result = + v3dv_CreateDescriptorPool(vk_device, &pool_info, NULL, + &pool->meta.descriptor_pool); + + if (result != VK_SUCCESS) + return result; + + VkDescriptorSetLayout set_layouts[2] = { + device->queries.buf_descriptor_set_layout, + device->queries.buf_descriptor_set_layout + }; + VkDescriptorSetAllocateInfo alloc_info = { + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO, + .descriptorPool = pool->meta.descriptor_pool, + .descriptorSetCount = 2, + .pSetLayouts = set_layouts, + }; + result = v3dv_AllocateDescriptorSets(vk_device, &alloc_info, + pool->meta.descriptor_sets); + if (result != VK_SUCCESS) + return result; + + VkDescriptorBufferInfo desc_buf_info[2] = { + { + .buffer = pool->meta.avail_buf, + .offset = 0, + .range = VK_WHOLE_SIZE, + }, + { + .buffer = pool->meta.res_buf, + .offset = 0, + .range = VK_WHOLE_SIZE, + }, + }; + + VkWriteDescriptorSet writes[2] = { + { + .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET, + .dstSet = pool->meta.descriptor_sets[0], + .dstBinding = 0, + .dstArrayElement = 0, + .descriptorCount = 1, + .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, + .pBufferInfo = &desc_buf_info[0], + }, + { + .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET, + .dstSet = pool->meta.descriptor_sets[1], + .dstBinding = 0, + .dstArrayElement = 0, + .descriptorCount = 1, + .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, + .pBufferInfo = &desc_buf_info[1], + }, + }; + v3dv_UpdateDescriptorSets(vk_device, 2, writes, 0, NULL); + + return VK_SUCCESS; +} + +static void +destroy_pool_descriptors(struct v3dv_device *device, + struct v3dv_query_pool *pool) +{ + assert(pool->query_type == VK_QUERY_TYPE_OCCLUSION); + + v3dv_FreeDescriptorSets(v3dv_device_to_handle(device), + pool->meta.descriptor_pool, + 2, pool->meta.descriptor_sets); + pool->meta.descriptor_sets[0] = VK_NULL_HANDLE; + pool->meta.descriptor_sets[1] = VK_NULL_HANDLE; + + v3dv_DestroyDescriptorPool(v3dv_device_to_handle(device), + pool->meta.descriptor_pool, NULL); + pool->meta.descriptor_pool = VK_NULL_HANDLE; +} + +static VkResult +pool_create_meta_resources(struct v3dv_device *device, + struct v3dv_query_pool *pool) +{ + VkResult result; + + if (pool->query_type != VK_QUERY_TYPE_OCCLUSION) + return VK_SUCCESS; + + result = create_vk_storage_buffer(device, pool->occlusion.bo, + &pool->meta.res_buf, + &pool->meta.res_mem); + if (result != VK_SUCCESS) + return result; + + result = create_vk_storage_buffer(device, pool->avail_bo, + &pool->meta.avail_buf, + &pool->meta.avail_mem); + + result = create_pool_descriptors(device, pool); + if (result != VK_SUCCESS) + return result; + + return VK_SUCCESS; +} + +static void +pool_destroy_meta_resources(struct v3dv_device *device, + struct v3dv_query_pool *pool) +{ + if (pool->query_type != VK_QUERY_TYPE_OCCLUSION) + return; + + destroy_pool_descriptors(device, pool); + destroy_vk_storage_buffer(device, &pool->meta.avail_buf, &pool->meta.avail_mem); + destroy_vk_storage_buffer(device, &pool->meta.res_buf, &pool->meta.res_mem); +} + VKAPI_ATTR VkResult VKAPI_CALL v3dv_CreateQueryPool(VkDevice _device, const VkQueryPoolCreateInfo *pCreateInfo, @@ -208,15 +401,28 @@ v3dv_CreateQueryPool(VkDevice _device, */ const uint32_t query_groups = DIV_ROUND_UP(pool->query_count, 16); const uint32_t bo_size = query_groups * 1024; - pool->bo = v3dv_bo_alloc(device, bo_size, "query", true); - if (!pool->bo) { + pool->occlusion.bo = v3dv_bo_alloc(device, bo_size, "query:r", true); + if (!pool->occlusion.bo) { result = vk_error(device, VK_ERROR_OUT_OF_DEVICE_MEMORY); goto fail; } - if (!v3dv_bo_map(device, pool->bo, bo_size)) { + if (!v3dv_bo_map(device, pool->occlusion.bo, bo_size)) { result = vk_error(device, VK_ERROR_OUT_OF_DEVICE_MEMORY); goto fail; } + + /* For now we only use the availability BO with occlusion queries, but + * in the future we may want to use this with more query types. + */ + pool->avail_bo = v3dv_bo_alloc(device, pool->query_count, "query:a", true); + if (!pool->avail_bo) { + result = vk_error(device, VK_ERROR_OUT_OF_DEVICE_MEMORY); + goto fail; + } + if (!v3dv_bo_map(device, pool->avail_bo, pool->avail_bo->size)) { + result = vk_error(device, VK_ERROR_OUT_OF_HOST_MEMORY); + goto fail; + } break; } case VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR: { @@ -243,14 +449,14 @@ v3dv_CreateQueryPool(VkDevice _device, unreachable("Unsupported query type"); } + /* Initialize queries in the pool */ for (; query_idx < pool->query_count; query_idx++) { pool->queries[query_idx].maybe_available = false; switch (pool->query_type) { case VK_QUERY_TYPE_OCCLUSION: { const uint32_t query_group = query_idx / 16; const uint32_t query_offset = query_group * 1024 + (query_idx % 16) * 4; - pool->queries[query_idx].bo = pool->bo; - pool->queries[query_idx].offset = query_offset; + pool->queries[query_idx].occlusion.offset = query_offset; break; } case VK_QUERY_TYPE_TIMESTAMP: @@ -272,6 +478,11 @@ v3dv_CreateQueryPool(VkDevice _device, } } + /* Create meta resources */ + result = pool_create_meta_resources(device, pool); + if (result != VK_SUCCESS) + goto fail; + *pQueryPool = v3dv_query_pool_to_handle(pool); return VK_SUCCESS; @@ -282,10 +493,13 @@ fail: vk_sync_destroy(&device->vk, pool->queries[j].perf.last_job_sync); } - if (pool->bo) - v3dv_bo_free(device, pool->bo); + if (pool->avail_bo) + v3dv_bo_free(device, pool->avail_bo); + if (pool->occlusion.bo) + v3dv_bo_free(device, pool->occlusion.bo); if (pool->queries) vk_free2(&device->vk.alloc, pAllocator, pool->queries); + pool_destroy_meta_resources(device, pool); vk_object_free(&device->vk, pAllocator, pool); return result; @@ -302,8 +516,11 @@ v3dv_DestroyQueryPool(VkDevice _device, if (!pool) return; - if (pool->bo) - v3dv_bo_free(device, pool->bo); + if (pool->avail_bo) + v3dv_bo_free(device, pool->avail_bo); + + if (pool->occlusion.bo) + v3dv_bo_free(device, pool->occlusion.bo); if (pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR) { for (uint32_t i = 0; i < pool->query_count; i++) { @@ -315,6 +532,8 @@ v3dv_DestroyQueryPool(VkDevice _device, if (pool->queries) vk_free2(&device->vk.alloc, pAllocator, pool->queries); + pool_destroy_meta_resources(device, pool); + vk_object_free(&device->vk, pAllocator, pool); } @@ -332,16 +551,34 @@ write_to_buffer(void *dst, uint32_t idx, bool do_64bit, uint64_t value) static VkResult query_wait_available(struct v3dv_device *device, + struct v3dv_query_pool *pool, struct v3dv_query *q, - VkQueryType query_type) + uint32_t query_idx) { + /* For occlusion queries we prefer to poll the availability BO in a loop + * to waiting on the occlusion query results BO, because the latter would + * make us wait for any job running occlusion queries, even if those queries + * do not involve the one we want to wait on. + */ + if (pool->query_type == VK_QUERY_TYPE_OCCLUSION) { + uint8_t *q_addr = ((uint8_t *) pool->avail_bo->map) + query_idx; + while (*q_addr == 0) + usleep(250); + return VK_SUCCESS; + } + + /* For other queries we need to wait for the queue to signal that + * the query has been submitted for execution before anything else. + */ + assert(pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR || + pool->query_type == VK_QUERY_TYPE_TIMESTAMP); + + VkResult result = VK_SUCCESS; if (!q->maybe_available) { struct timespec timeout; timespec_get(&timeout, TIME_UTC); timespec_add_msec(&timeout, &timeout, 2000); - VkResult result = VK_SUCCESS; - mtx_lock(&device->query_mutex); while (!q->maybe_available) { if (vk_device_is_lost(&device->vk)) { @@ -362,16 +599,77 @@ query_wait_available(struct v3dv_device *device, if (result != VK_SUCCESS) return result; + + /* For performance queries, we also need to wait for the relevant syncobj + * to be signaled to ensure completion of the GPU work. + */ + if (pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR && + vk_sync_wait(&device->vk, q->perf.last_job_sync, + 0, VK_SYNC_WAIT_COMPLETE, UINT64_MAX) != VK_SUCCESS) { + return vk_device_set_lost(&device->vk, "Query job wait failed"); + } } - if (query_type == VK_QUERY_TYPE_OCCLUSION && - !v3dv_bo_wait(device, q->bo, 0xffffffffffffffffull)) - return vk_device_set_lost(&device->vk, "Query BO wait failed: %m"); + return result; +} - if (query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR && +static VkResult +query_check_available(struct v3dv_device *device, + struct v3dv_query_pool *pool, + struct v3dv_query *q, + uint32_t query_idx) +{ + /* For occlusion and performance queries we check the availability BO */ + if (pool->query_type == VK_QUERY_TYPE_OCCLUSION) { + const uint8_t *q_addr = ((uint8_t *) pool->avail_bo->map) + query_idx; + return (*q_addr != 0) ? VK_SUCCESS : VK_NOT_READY; + } + + /* For other queries we need to check if the queue has submitted the query + * for execution at all. + */ + assert(pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR || + pool->query_type == VK_QUERY_TYPE_TIMESTAMP); + if (!q->maybe_available) + return VK_NOT_READY; + + /* For performance queries, we also need to check if the relevant GPU job + * has completed. + */ + if (pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR && vk_sync_wait(&device->vk, q->perf.last_job_sync, - 0, VK_SYNC_WAIT_COMPLETE, UINT64_MAX) != VK_SUCCESS) - return vk_device_set_lost(&device->vk, "Query job wait failed"); + 0, VK_SYNC_WAIT_COMPLETE, 0) != VK_SUCCESS) { + return VK_NOT_READY; + } + + return VK_SUCCESS; +} + +static VkResult +query_is_available(struct v3dv_device *device, + struct v3dv_query_pool *pool, + uint32_t query, + bool do_wait, + bool *available) +{ + struct v3dv_query *q = &pool->queries[query]; + + assert(pool->query_type != VK_QUERY_TYPE_OCCLUSION || + (pool->occlusion.bo && pool->occlusion.bo->map)); + + if (do_wait) { + VkResult result = query_wait_available(device, pool, q, query); + if (result != VK_SUCCESS) { + *available = false; + return result; + } + + *available = true; + } else { + VkResult result = query_check_available(device, pool, q, query); + assert(result == VK_SUCCESS || result == VK_NOT_READY); + *available = (result == VK_SUCCESS); + } return VK_SUCCESS; } @@ -390,9 +688,10 @@ write_occlusion_query_result(struct v3dv_device *device, return VK_ERROR_DEVICE_LOST; struct v3dv_query *q = &pool->queries[query]; - assert(q->bo && q->bo->map); + assert(pool->occlusion.bo && pool->occlusion.bo->map); - const uint8_t *query_addr = ((uint8_t *) q->bo->map) + q->offset; + const uint8_t *query_addr = + ((uint8_t *) pool->occlusion.bo->map) + q->occlusion.offset; write_to_buffer(data, slot, do_64bit, (uint64_t) *((uint32_t *)query_addr)); return VK_SUCCESS; } @@ -449,26 +748,6 @@ write_performance_query_result(struct v3dv_device *device, return VK_SUCCESS; } -static VkResult -query_check_available(struct v3dv_device *device, - struct v3dv_query *q, - VkQueryType query_type) -{ - if (!q->maybe_available) - return VK_NOT_READY; - - if (query_type == VK_QUERY_TYPE_OCCLUSION && - !v3dv_bo_wait(device, q->bo, 0)) - return VK_NOT_READY; - - if (query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR && - vk_sync_wait(&device->vk, q->perf.last_job_sync, - 0, VK_SYNC_WAIT_COMPLETE, 0) != VK_SUCCESS) - return VK_NOT_READY; - - return VK_SUCCESS; -} - static VkResult write_query_result(struct v3dv_device *device, struct v3dv_query_pool *pool, @@ -492,35 +771,6 @@ write_query_result(struct v3dv_device *device, } } -static VkResult -query_is_available(struct v3dv_device *device, - struct v3dv_query_pool *pool, - uint32_t query, - bool do_wait, - bool *available) -{ - struct v3dv_query *q = &pool->queries[query]; - - assert(pool->query_type != VK_QUERY_TYPE_OCCLUSION || - (q->bo && q->bo->map)); - - if (do_wait) { - VkResult result = query_wait_available(device, q, pool->query_type); - if (result != VK_SUCCESS) { - *available = false; - return result; - } - - *available = true; - } else { - VkResult result = query_check_available(device, q, pool->query_type); - assert(result == VK_SUCCESS || result == VK_NOT_READY); - *available = (result == VK_SUCCESS); - } - - return VK_SUCCESS; -} - static uint32_t get_query_result_count(struct v3dv_query_pool *pool) { @@ -536,13 +786,13 @@ get_query_result_count(struct v3dv_query_pool *pool) } VkResult -v3dv_get_query_pool_results(struct v3dv_device *device, - struct v3dv_query_pool *pool, - uint32_t first, - uint32_t count, - void *data, - VkDeviceSize stride, - VkQueryResultFlags flags) +v3dv_get_query_pool_results_cpu(struct v3dv_device *device, + struct v3dv_query_pool *pool, + uint32_t first, + uint32_t count, + void *data, + VkDeviceSize stride, + VkQueryResultFlags flags) { assert(first < pool->query_count); assert(first + count <= pool->query_count); @@ -605,8 +855,168 @@ v3dv_GetQueryPoolResults(VkDevice _device, V3DV_FROM_HANDLE(v3dv_device, device, _device); V3DV_FROM_HANDLE(v3dv_query_pool, pool, queryPool); - return v3dv_get_query_pool_results(device, pool, firstQuery, queryCount, - pData, stride, flags); + return v3dv_get_query_pool_results_cpu(device, pool, firstQuery, queryCount, + pData, stride, flags); +} + +/* Emits a series of vkCmdDispatchBase calls to execute all the workgroups + * required to handle a number of queries considering per-dispatch limits. + */ +static void +cmd_buffer_emit_dispatch_queries(struct v3dv_cmd_buffer *cmd_buffer, + uint32_t query_count) +{ + VkCommandBuffer vk_cmd_buffer = v3dv_cmd_buffer_to_handle(cmd_buffer); + + uint32_t dispatched = 0; + const uint32_t max_batch_size = 65535; + while (dispatched < query_count) { + uint32_t batch_size = MIN2(query_count - dispatched, max_batch_size); + v3dv_CmdDispatchBase(vk_cmd_buffer, dispatched, 0, 0, batch_size, 1, 1); + dispatched += batch_size; + } +} + +void +v3dv_cmd_buffer_emit_set_query_availability(struct v3dv_cmd_buffer *cmd_buffer, + struct v3dv_query_pool *pool, + uint32_t query, uint32_t count, + uint8_t availability) +{ + assert(pool->query_type == VK_QUERY_TYPE_OCCLUSION || + pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR); + + struct v3dv_device *device = cmd_buffer->device; + VkCommandBuffer vk_cmd_buffer = v3dv_cmd_buffer_to_handle(cmd_buffer); + + /* We are about to emit a compute job to set query availability and we need + * to ensure this executes after the graphics work using the queries has + * completed. + */ + VkMemoryBarrier2 barrier = { + .sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER_2, + .srcStageMask = VK_PIPELINE_STAGE_2_COLOR_ATTACHMENT_OUTPUT_BIT, + .dstStageMask = VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT, + }; + VkDependencyInfo barrier_info = { + .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO, + .memoryBarrierCount = 1, + .pMemoryBarriers = &barrier, + }; + v3dv_cmd_buffer_emit_pipeline_barrier(cmd_buffer, &barrier_info); + + /* Dispatch queries */ + v3dv_cmd_buffer_meta_state_push(cmd_buffer, true); + + v3dv_CmdBindPipeline(vk_cmd_buffer, + VK_PIPELINE_BIND_POINT_COMPUTE, + device->queries.avail_pipeline); + + v3dv_CmdBindDescriptorSets(vk_cmd_buffer, + VK_PIPELINE_BIND_POINT_COMPUTE, + device->queries.avail_pipeline_layout, + 0, 1, &pool->meta.descriptor_sets[0], + 0, NULL); + + struct { + uint32_t query; + uint8_t availability; + } push_data = { query, availability }; + v3dv_CmdPushConstants(vk_cmd_buffer, + device->queries.avail_pipeline_layout, + VK_SHADER_STAGE_COMPUTE_BIT, + 0, 5, &push_data); + cmd_buffer_emit_dispatch_queries(cmd_buffer, count); + + v3dv_cmd_buffer_meta_state_pop(cmd_buffer, 0, false); +} + +static void +cmd_buffer_emit_reset_occlusion_query_pool(struct v3dv_cmd_buffer *cmd_buffer, + struct v3dv_query_pool *pool, + uint32_t query, uint32_t count) +{ + struct v3dv_device *device = cmd_buffer->device; + VkCommandBuffer vk_cmd_buffer = v3dv_cmd_buffer_to_handle(cmd_buffer); + + /* Ensure the GPU is done with the queries in the graphics queue before + * we reset in the compute queue. + */ + VkMemoryBarrier2 barrier = { + .sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER_2, + .srcStageMask = VK_PIPELINE_STAGE_2_COLOR_ATTACHMENT_OUTPUT_BIT, + .dstStageMask = VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT, + }; + VkDependencyInfo barrier_info = { + .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO, + .memoryBarrierCount = 1, + .pMemoryBarriers = &barrier, + }; + v3dv_cmd_buffer_emit_pipeline_barrier(cmd_buffer, &barrier_info); + + /* Emit compute reset */ + v3dv_cmd_buffer_meta_state_push(cmd_buffer, true); + + v3dv_CmdBindPipeline(vk_cmd_buffer, + VK_PIPELINE_BIND_POINT_COMPUTE, + device->queries.reset_occlusion_pipeline); + + v3dv_CmdBindDescriptorSets(vk_cmd_buffer, + VK_PIPELINE_BIND_POINT_COMPUTE, + device->queries.reset_occlusion_pipeline_layout, + 0, 2, pool->meta.descriptor_sets, + 0, NULL); + + v3dv_CmdPushConstants(vk_cmd_buffer, + device->queries.reset_occlusion_pipeline_layout, + VK_SHADER_STAGE_COMPUTE_BIT, + 0, 4, &query); + + cmd_buffer_emit_dispatch_queries(cmd_buffer, count); + + v3dv_cmd_buffer_meta_state_pop(cmd_buffer, 0, false); + + /* Ensure future work in the graphics queue using the queries doesn't start + * before the reset completed. + */ + barrier = (VkMemoryBarrier2) { + .sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER_2, + .srcStageMask = VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT, + .dstStageMask = VK_PIPELINE_STAGE_2_EARLY_FRAGMENT_TESTS_BIT, + }; + barrier_info = (VkDependencyInfo) { + .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO, + .memoryBarrierCount = 1, + .pMemoryBarriers = &barrier, + }; + v3dv_cmd_buffer_emit_pipeline_barrier(cmd_buffer, &barrier_info); +} + +static void +cmd_buffer_emit_reset_query_pool(struct v3dv_cmd_buffer *cmd_buffer, + struct v3dv_query_pool *pool, + uint32_t first, uint32_t count) +{ + assert(pool->query_type == VK_QUERY_TYPE_OCCLUSION); + cmd_buffer_emit_reset_occlusion_query_pool(cmd_buffer, pool, first, count); +} + +static void +cmd_buffer_emit_reset_query_pool_cpu(struct v3dv_cmd_buffer *cmd_buffer, + struct v3dv_query_pool *pool, + uint32_t first, uint32_t count) +{ + assert(pool->query_type != VK_QUERY_TYPE_OCCLUSION); + + struct v3dv_job *job = + v3dv_cmd_buffer_create_cpu_job(cmd_buffer->device, + V3DV_JOB_TYPE_CPU_RESET_QUERIES, + cmd_buffer, -1); + v3dv_return_if_oom(cmd_buffer, NULL); + job->cpu.query_reset.pool = pool; + job->cpu.query_reset.first = first; + job->cpu.query_reset.count = count; + list_addtail(&job->list_link, &cmd_buffer->jobs); } VKAPI_ATTR void VKAPI_CALL @@ -618,7 +1028,228 @@ v3dv_CmdResetQueryPool(VkCommandBuffer commandBuffer, V3DV_FROM_HANDLE(v3dv_cmd_buffer, cmd_buffer, commandBuffer); V3DV_FROM_HANDLE(v3dv_query_pool, pool, queryPool); - v3dv_cmd_buffer_reset_queries(cmd_buffer, pool, firstQuery, queryCount); + /* Resets can only happen outside a render pass instance so we should not + * be in the middle of job recording. + */ + assert(cmd_buffer->state.pass == NULL); + assert(cmd_buffer->state.job == NULL); + + assert(firstQuery < pool->query_count); + assert(firstQuery + queryCount <= pool->query_count); + + /* We can reset occlusion queries in the GPU, but for other query types + * we emit a CPU job that will call v3dv_reset_query_pool_cpu when executed + * in the queue. + */ + if (pool->query_type == VK_QUERY_TYPE_OCCLUSION) { + cmd_buffer_emit_reset_query_pool(cmd_buffer, pool, firstQuery, queryCount); + } else { + cmd_buffer_emit_reset_query_pool_cpu(cmd_buffer, pool, + firstQuery, queryCount); + } +} + +/** + * Creates a descriptor pool so we can create a descriptors for the destination + * buffers of vkCmdCopyQueryResults for queries where this is implemented in + * the GPU. + */ +static VkResult +create_storage_buffer_descriptor_pool(struct v3dv_cmd_buffer *cmd_buffer) +{ + /* If this is not the first pool we create one for this command buffer + * size it based on the size of the currently exhausted pool. + */ + uint32_t descriptor_count = 32; + if (cmd_buffer->meta.query.dspool != VK_NULL_HANDLE) { + struct v3dv_descriptor_pool *exhausted_pool = + v3dv_descriptor_pool_from_handle(cmd_buffer->meta.query.dspool); + descriptor_count = MIN2(exhausted_pool->max_entry_count * 2, 1024); + } + + /* Create the descriptor pool */ + cmd_buffer->meta.query.dspool = VK_NULL_HANDLE; + VkDescriptorPoolSize pool_size = { + .type = VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER, + .descriptorCount = descriptor_count, + }; + VkDescriptorPoolCreateInfo info = { + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO, + .maxSets = descriptor_count, + .poolSizeCount = 1, + .pPoolSizes = &pool_size, + .flags = 0, + }; + VkResult result = + v3dv_CreateDescriptorPool(v3dv_device_to_handle(cmd_buffer->device), + &info, + &cmd_buffer->device->vk.alloc, + &cmd_buffer->meta.query.dspool); + + if (result == VK_SUCCESS) { + assert(cmd_buffer->meta.query.dspool != VK_NULL_HANDLE); + const VkDescriptorPool vk_pool = cmd_buffer->meta.query.dspool; + + v3dv_cmd_buffer_add_private_obj( + cmd_buffer, (uintptr_t) vk_pool, + (v3dv_cmd_buffer_private_obj_destroy_cb)v3dv_DestroyDescriptorPool); + + struct v3dv_descriptor_pool *pool = + v3dv_descriptor_pool_from_handle(vk_pool); + pool->is_driver_internal = true; + } + + return result; +} + +static VkResult +allocate_storage_buffer_descriptor_set(struct v3dv_cmd_buffer *cmd_buffer, + VkDescriptorSet *set) +{ + /* Make sure we have a descriptor pool */ + VkResult result; + if (cmd_buffer->meta.query.dspool == VK_NULL_HANDLE) { + result = create_storage_buffer_descriptor_pool(cmd_buffer); + if (result != VK_SUCCESS) + return result; + } + assert(cmd_buffer->meta.query.dspool != VK_NULL_HANDLE); + + /* Allocate descriptor set */ + struct v3dv_device *device = cmd_buffer->device; + VkDevice vk_device = v3dv_device_to_handle(device); + VkDescriptorSetAllocateInfo info = { + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO, + .descriptorPool = cmd_buffer->meta.query.dspool, + .descriptorSetCount = 1, + .pSetLayouts = &device->queries.buf_descriptor_set_layout, + }; + result = v3dv_AllocateDescriptorSets(vk_device, &info, set); + + /* If we ran out of pool space, grow the pool and try again */ + if (result == VK_ERROR_OUT_OF_POOL_MEMORY) { + result = create_storage_buffer_descriptor_pool(cmd_buffer); + if (result == VK_SUCCESS) { + info.descriptorPool = cmd_buffer->meta.query.dspool; + result = v3dv_AllocateDescriptorSets(vk_device, &info, set); + } + } + + return result; +} + +static void +cmd_buffer_emit_copy_query_pool_results(struct v3dv_cmd_buffer *cmd_buffer, + struct v3dv_query_pool *pool, + uint32_t first, uint32_t count, + struct v3dv_buffer *buf, + uint32_t offset, uint32_t stride, + VkQueryResultFlags flags) +{ + struct v3dv_device *device = cmd_buffer->device; + VkDevice vk_device = v3dv_device_to_handle(device); + VkCommandBuffer vk_cmd_buffer = v3dv_cmd_buffer_to_handle(cmd_buffer); + + /* FIXME: do we need this barrier? Since vkCmdEndQuery should've been called + * and that already waits maybe we don't (since this is serialized + * in the compute queue with EndQuery anyway). + */ + if (flags & VK_QUERY_RESULT_WAIT_BIT) { + VkMemoryBarrier2 barrier = { + .sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER_2, + .srcStageMask = VK_PIPELINE_STAGE_2_COLOR_ATTACHMENT_OUTPUT_BIT, + .dstStageMask = VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT, + }; + VkDependencyInfo barrier_info = { + .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO, + .memoryBarrierCount = 1, + .pMemoryBarriers = &barrier, + }; + v3dv_cmd_buffer_emit_pipeline_barrier(cmd_buffer, &barrier_info); + } + + /* Allocate and setup descriptor set for output buffer */ + VkDescriptorSet out_buf_descriptor_set; + VkResult result = + allocate_storage_buffer_descriptor_set(cmd_buffer, + &out_buf_descriptor_set); + if (result != VK_SUCCESS) { + fprintf(stderr, "vkCmdCopyQueryPoolResults failed: " + "could not allocate descriptor.\n"); + return; + } + + VkDescriptorBufferInfo desc_buf_info = { + .buffer = v3dv_buffer_to_handle(buf), + .offset = 0, + .range = VK_WHOLE_SIZE, + }; + VkWriteDescriptorSet write = { + .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET, + .dstSet = out_buf_descriptor_set, + .dstBinding = 0, + .dstArrayElement = 0, + .descriptorCount = 1, + .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, + .pBufferInfo = &desc_buf_info, + }; + v3dv_UpdateDescriptorSets(vk_device, 1, &write, 0, NULL); + + /* Dispatch copy */ + v3dv_cmd_buffer_meta_state_push(cmd_buffer, true); + + v3dv_CmdBindPipeline(vk_cmd_buffer, + VK_PIPELINE_BIND_POINT_COMPUTE, + device->queries.copy_pipeline); + + VkDescriptorSet sets[3] = { + pool->meta.descriptor_sets[0], + pool->meta.descriptor_sets[1], + out_buf_descriptor_set, + }; + v3dv_CmdBindDescriptorSets(vk_cmd_buffer, + VK_PIPELINE_BIND_POINT_COMPUTE, + device->queries.copy_pipeline_layout, + 0, 3, sets, 0, NULL); + + struct { + uint32_t first, offset, stride, flags; + } push_data = { first, offset, stride, flags }; + v3dv_CmdPushConstants(vk_cmd_buffer, + device->queries.copy_pipeline_layout, + VK_SHADER_STAGE_COMPUTE_BIT, + 0, sizeof(push_data), &push_data); + + cmd_buffer_emit_dispatch_queries(cmd_buffer, count); + + v3dv_cmd_buffer_meta_state_pop(cmd_buffer, 0, false); +} + +static void +cmd_buffer_emit_copy_query_pool_results_cpu(struct v3dv_cmd_buffer *cmd_buffer, + struct v3dv_query_pool *pool, + uint32_t first, + uint32_t count, + struct v3dv_buffer *dst, + uint32_t offset, + uint32_t stride, + VkQueryResultFlags flags) +{ + struct v3dv_job *job = + v3dv_cmd_buffer_create_cpu_job(cmd_buffer->device, + V3DV_JOB_TYPE_CPU_COPY_QUERY_RESULTS, + cmd_buffer, -1); + v3dv_return_if_oom(cmd_buffer, NULL); + + job->cpu.query_copy_results.pool = pool; + job->cpu.query_copy_results.first = first; + job->cpu.query_copy_results.count = count; + job->cpu.query_copy_results.dst = dst; + job->cpu.query_copy_results.offset = offset; + job->cpu.query_copy_results.stride = stride; + job->cpu.query_copy_results.flags = flags; + + list_addtail(&job->list_link, &cmd_buffer->jobs); } VKAPI_ATTR void VKAPI_CALL @@ -635,9 +1266,30 @@ v3dv_CmdCopyQueryPoolResults(VkCommandBuffer commandBuffer, V3DV_FROM_HANDLE(v3dv_query_pool, pool, queryPool); V3DV_FROM_HANDLE(v3dv_buffer, dst, dstBuffer); - v3dv_cmd_buffer_copy_query_results(cmd_buffer, pool, - firstQuery, queryCount, - dst, dstOffset, stride, flags); + /* Copies can only happen outside a render pass instance so we should not + * be in the middle of job recording. + */ + assert(cmd_buffer->state.pass == NULL); + assert(cmd_buffer->state.job == NULL); + + assert(firstQuery < pool->query_count); + assert(firstQuery + queryCount <= pool->query_count); + + /* For occlusion queries we implement the copy in the GPU but for other + * queries we emit a CPU job that will call v3dv_get_query_pool_results_cpu + * when executed in the queue. + */ + if (pool->query_type == VK_QUERY_TYPE_OCCLUSION) { + cmd_buffer_emit_copy_query_pool_results(cmd_buffer, pool, + firstQuery, queryCount, + dst, (uint32_t) dstOffset, + (uint32_t) stride, flags); + } else { + cmd_buffer_emit_copy_query_pool_results_cpu(cmd_buffer, pool, + firstQuery, queryCount, + dst, (uint32_t)dstOffset, + (uint32_t) stride, flags); + } } VKAPI_ATTR void VKAPI_CALL @@ -664,10 +1316,10 @@ v3dv_CmdEndQuery(VkCommandBuffer commandBuffer, } void -v3dv_reset_query_pools(struct v3dv_device *device, - struct v3dv_query_pool *pool, - uint32_t first, - uint32_t count) +v3dv_reset_query_pool_cpu(struct v3dv_device *device, + struct v3dv_query_pool *pool, + uint32_t first, + uint32_t count) { mtx_lock(&device->query_mutex); @@ -677,7 +1329,13 @@ v3dv_reset_query_pools(struct v3dv_device *device, q->maybe_available = false; switch (pool->query_type) { case VK_QUERY_TYPE_OCCLUSION: { - const uint8_t *q_addr = ((uint8_t *) q->bo->map) + q->offset; + /* Reset availability */ + uint8_t *base_addr = ((uint8_t *) pool->avail_bo->map) + first; + memset(base_addr, 0, count); + + /* Reset occlusion counter */ + const uint8_t *q_addr = + ((uint8_t *) pool->occlusion.bo->map) + q->occlusion.offset; uint32_t *counter = (uint32_t *) q_addr; *counter = 0; break; @@ -708,7 +1366,7 @@ v3dv_ResetQueryPool(VkDevice _device, V3DV_FROM_HANDLE(v3dv_device, device, _device); V3DV_FROM_HANDLE(v3dv_query_pool, pool, queryPool); - v3dv_reset_query_pools(device, pool, firstQuery, queryCount); + v3dv_reset_query_pool_cpu(device, pool, firstQuery, queryCount); } VKAPI_ATTR VkResult VKAPI_CALL @@ -776,3 +1434,460 @@ VKAPI_ATTR void VKAPI_CALL v3dv_ReleaseProfilingLockKHR(VkDevice device) { } + +static inline void +nir_set_query_availability(nir_builder *b, + nir_ssa_def *buf, + nir_ssa_def *query_idx, + nir_ssa_def *avail) +{ + nir_ssa_def *offset = query_idx; /* we use 1B per query */ + nir_store_ssbo(b, avail, buf, offset, .write_mask = 0x1, .align_mul = 1); +} + +static inline nir_ssa_def * +nir_get_query_availability(nir_builder *b, + nir_ssa_def *buf, + nir_ssa_def *query_idx) +{ + nir_ssa_def *offset = query_idx; /* we use 1B per query */ + nir_ssa_def *avail = + nir_load_ssbo(b, 1, 8, buf, offset, .align_mul = 1); + return nir_i2i32(b, avail); +} + +static nir_shader * +get_set_query_availability_cs() +{ + const nir_shader_compiler_options *options = v3dv_pipeline_get_nir_options(); + nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, + "set query availability cs"); + + /* We rely on supergroup packing to maximize SIMD lane occupancy */ + b.shader->info.workgroup_size[0] = 1; + b.shader->info.workgroup_size[1] = 1; + b.shader->info.workgroup_size[2] = 1; + + nir_ssa_def *buf = + nir_vulkan_resource_index(&b, 2, 32, nir_imm_int(&b, 0), + .desc_set = 0, + .binding = 0, + .desc_type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER); + + /* This assumes a local size of 1 and a horizontal-only dispatch. If we + * ever change any of these parameters we need to update how we compute the + * query index here. + */ + nir_ssa_def *wg_id = nir_channel(&b, nir_load_workgroup_id(&b, 32), 0); + + nir_ssa_def *base_query_idx = + nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 0, .range = 4); + + nir_ssa_def *avail = + nir_load_push_constant(&b, 1, 8, nir_imm_int(&b, 0), .base = 4, .range = 1); + + nir_ssa_def *query_idx = nir_iadd(&b, base_query_idx, wg_id); + nir_set_query_availability(&b, buf, query_idx, avail); + + return b.shader; +} + +static inline nir_ssa_def * +nir_get_occlusion_counter_offset(nir_builder *b, nir_ssa_def *query_idx) +{ + nir_ssa_def *query_group = nir_udiv(b, query_idx, nir_imm_int(b, 16)); + nir_ssa_def *query_group_offset = nir_umod(b, query_idx, nir_imm_int(b, 16)); + nir_ssa_def *offset = + nir_iadd(b, nir_imul(b, query_group, nir_imm_int(b, 1024)), + nir_imul(b, query_group_offset, nir_imm_int(b, 4))); + return offset; +} + +static inline void +nir_reset_occlusion_counter(nir_builder *b, + nir_ssa_def *buf, + nir_ssa_def *query_idx) +{ + nir_ssa_def *offset = nir_get_occlusion_counter_offset(b, query_idx); + nir_ssa_def *zero = nir_imm_int(b, 0); + nir_store_ssbo(b, zero, buf, offset, .write_mask = 0x1, .align_mul = 4); +} + +static inline nir_ssa_def * +nir_read_occlusion_counter(nir_builder *b, + nir_ssa_def *buf, + nir_ssa_def *query_idx) +{ + nir_ssa_def *offset = nir_get_occlusion_counter_offset(b, query_idx); + return nir_load_ssbo(b, 1, 32, buf, offset, .access = 0, .align_mul = 4); +} + +static nir_shader * +get_reset_occlusion_query_cs() +{ + const nir_shader_compiler_options *options = v3dv_pipeline_get_nir_options(); + nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, + "reset occlusion query cs"); + + /* We rely on supergroup packing to maximize SIMD lane occupancy */ + b.shader->info.workgroup_size[0] = 1; + b.shader->info.workgroup_size[1] = 1; + b.shader->info.workgroup_size[2] = 1; + + nir_ssa_def *buf_avail = + nir_vulkan_resource_index(&b, 2, 32, nir_imm_int(&b, 0), + .desc_set = 0, + .binding = 0, + .desc_type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER); + + nir_ssa_def *buf_res = + nir_vulkan_resource_index(&b, 2, 32, nir_imm_int(&b, 0), + .desc_set = 1, + .binding = 0, + .desc_type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER); + + /* This assumes a local size of 1 and a horizontal-only dispatch. If we + * ever change any of these parameters we need to update how we compute the + * query index here. + */ + nir_ssa_def *wg_id = nir_channel(&b, nir_load_workgroup_id(&b, 32), 0); + + nir_ssa_def *base_query_idx = + nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 0, .range = 4); + + nir_ssa_def *query_idx = nir_iadd(&b, base_query_idx, wg_id); + + nir_set_query_availability(&b, buf_avail, query_idx, nir_imm_intN_t(&b, 0, 8)); + nir_reset_occlusion_counter(&b, buf_res, query_idx); + + return b.shader; +} + +static void +write_query_buffer(nir_builder *b, + nir_ssa_def *buf, + nir_ssa_def **offset, + nir_ssa_def *value, + nir_ssa_def *flag_64bit) +{ + nir_ssa_def *do_64bit_cond = nir_ine_imm(b, flag_64bit, 0); + nir_if *if_stmt = nir_push_if(b, do_64bit_cond); + /* Create a 64-bit value using a vec2 with the .Y component set to 0 + * so we can write a 64-bit value in a single store. + */ + nir_ssa_def *value64 = nir_vec2(b, value, nir_imm_int(b, 0)); + nir_store_ssbo(b, value64, buf, *offset, .write_mask = 0x3, .align_mul = 8); + nir_push_else(b, NULL); + nir_store_ssbo(b, value, buf, *offset, .write_mask = 0x1, .align_mul = 4); + nir_pop_if(b, if_stmt); + + *offset = nir_bcsel(b, do_64bit_cond, + nir_iadd(b, *offset, nir_imm_int(b, 8)), + nir_iadd(b, *offset, nir_imm_int(b, 4))); +} + +static nir_shader * +get_copy_query_results_cs() +{ + const nir_shader_compiler_options *options = v3dv_pipeline_get_nir_options(); + nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options, + "copy query results cs"); + + /* We rely on supergroup packing to maximize SIMD lane occupancy */ + b.shader->info.workgroup_size[0] = 1; + b.shader->info.workgroup_size[1] = 1; + b.shader->info.workgroup_size[2] = 1; + + nir_ssa_def *buf_avail = + nir_vulkan_resource_index(&b, 2, 32, nir_imm_int(&b, 0), + .desc_set = 0, + .binding = 0, + .desc_type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER); + + nir_ssa_def *buf_res = + nir_vulkan_resource_index(&b, 2, 32, nir_imm_int(&b, 0), + .desc_set = 1, + .binding = 0, + .desc_type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER); + + nir_ssa_def *buf_out = + nir_vulkan_resource_index(&b, 2, 32, nir_imm_int(&b, 0), + .desc_set = 2, + .binding = 0, + .desc_type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER); + + /* Read push constants */ + nir_ssa_def *base_query_idx = + nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 0, .range = 4); + + nir_ssa_def *base_offset_out = + nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 4, .range = 4); + + nir_ssa_def *stride = + nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 8, .range = 4); + + nir_ssa_def *flags = + nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 12, .range = 4); + + /* Check copy flags */ + nir_ssa_def *flag_64bit = + nir_iand(&b, flags, nir_imm_int(&b, (uint32_t)VK_QUERY_RESULT_64_BIT)); + + nir_ssa_def *flag_partial = + nir_iand(&b, flags, nir_imm_int(&b, (uint32_t)VK_QUERY_RESULT_PARTIAL_BIT)); + + nir_ssa_def *flag_avail = + nir_iand(&b, flags, nir_imm_int(&b, (uint32_t)VK_QUERY_RESULT_WITH_AVAILABILITY_BIT)); + + /* This assumes a local size of 1 and a horizontal-only dispatch. If we + * ever change any of these parameters we need to update how we compute the + * query index here. + */ + nir_ssa_def *wg_id = nir_channel(&b, nir_load_workgroup_id(&b, 32), 0); + nir_ssa_def *query_idx = nir_iadd(&b, base_query_idx, wg_id); + + /* Read query availability */ + nir_ssa_def *avail = nir_get_query_availability(&b, buf_avail, query_idx); + + /* Read query result */ + nir_ssa_def *query_res = nir_read_occlusion_counter(&b, buf_res, query_idx); + + /* Write output buffer */ + nir_ssa_def *offset = + nir_iadd(&b, base_offset_out, nir_imul(&b, wg_id, stride)); + nir_ssa_def *write_result = nir_ior(&b, flag_partial, avail); + + /* Query result */ + nir_if *if_stmt = nir_push_if(&b, nir_ine_imm(&b, write_result, 0)); + write_query_buffer(&b, buf_out, &offset, query_res, flag_64bit); + nir_pop_if(&b, if_stmt); + + /* Availability */ + if_stmt = nir_push_if(&b, nir_ine_imm(&b, flag_avail, 0)); + write_query_buffer(&b, buf_out, &offset, avail, flag_64bit); + nir_pop_if(&b, if_stmt); + + return b.shader; +} + +static bool +create_query_pipelines(struct v3dv_device *device) +{ + VkResult result; + VkPipeline pipeline; + + /* Set layout: single storage buffer */ + if (!device->queries.buf_descriptor_set_layout) { + VkDescriptorSetLayoutBinding descriptor_set_layout_binding = { + .binding = 0, + .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, + .descriptorCount = 1, + .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT, + }; + VkDescriptorSetLayoutCreateInfo descriptor_set_layout_info = { + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO, + .bindingCount = 1, + .pBindings = &descriptor_set_layout_binding, + }; + result = + v3dv_CreateDescriptorSetLayout(v3dv_device_to_handle(device), + &descriptor_set_layout_info, + &device->vk.alloc, + &device->queries.buf_descriptor_set_layout); + if (result != VK_SUCCESS) + return false; + } + + /* Set availability pipeline. + * + * Pipeline layout: + * - 1 storage buffer for the BO with the query availability. + * - 2 push constants: + * 0B: base query index (4 bytes). + * 4B: availability (1 byte). + */ + if (!device->queries.avail_pipeline_layout) { + VkPipelineLayoutCreateInfo pipeline_layout_info = { + .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, + .setLayoutCount = 1, + .pSetLayouts = &device->queries.buf_descriptor_set_layout, + .pushConstantRangeCount = 1, + .pPushConstantRanges = + &(VkPushConstantRange) { VK_SHADER_STAGE_COMPUTE_BIT, 0, 5 }, + }; + + result = + v3dv_CreatePipelineLayout(v3dv_device_to_handle(device), + &pipeline_layout_info, + &device->vk.alloc, + &device->queries.avail_pipeline_layout); + + if (result != VK_SUCCESS) + return false; + } + + if (!device->queries.avail_pipeline) { + nir_shader *set_query_availability_cs_nir = get_set_query_availability_cs(); + result = v3dv_create_compute_pipeline_from_nir(device, + set_query_availability_cs_nir, + device->queries.avail_pipeline_layout, + &pipeline); + ralloc_free(set_query_availability_cs_nir); + if (result != VK_SUCCESS) + return false; + + device->queries.avail_pipeline = pipeline; + } + + /* Reset occlusion query pipeline. + * + * Pipeline layout: + * - 1 storage buffer for the BO with the query availability. + * - 1 storage buffer for the BO with the occlusion query results. + * - Push constants: + * 0B: base query index (4B) + */ + if (!device->queries.reset_occlusion_pipeline_layout) { + VkDescriptorSetLayout set_layouts[2] = { + device->queries.buf_descriptor_set_layout, + device->queries.buf_descriptor_set_layout, + }; + VkPipelineLayoutCreateInfo pipeline_layout_info = { + .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, + .setLayoutCount = 2, + .pSetLayouts = set_layouts, + .pushConstantRangeCount = 1, + .pPushConstantRanges = + &(VkPushConstantRange) { VK_SHADER_STAGE_COMPUTE_BIT, 0, 4 }, + }; + + result = + v3dv_CreatePipelineLayout(v3dv_device_to_handle(device), + &pipeline_layout_info, + &device->vk.alloc, + &device->queries.reset_occlusion_pipeline_layout); + + if (result != VK_SUCCESS) + return false; + } + + if (!device->queries.reset_occlusion_pipeline) { + nir_shader *reset_occlusion_query_cs_nir = get_reset_occlusion_query_cs(); + result = v3dv_create_compute_pipeline_from_nir( + device, + reset_occlusion_query_cs_nir, + device->queries.reset_occlusion_pipeline_layout, + &pipeline); + ralloc_free(reset_occlusion_query_cs_nir); + if (result != VK_SUCCESS) + return false; + + device->queries.reset_occlusion_pipeline = pipeline; + } + + /* Copy query results pipeline. + * + * Pipeline layout: + * - 1 storage buffer for the BO with the query availability. + * - 1 storage buffer for the BO with the occlusion query results. + * - 1 storage buffer for the output. + * - Push constants: + * 0B: base query index (4B) + * 4B: offset into output buffer (4B) + * 8B: stride (4B) + * 12B: copy flags (4B) + */ + if (!device->queries.copy_pipeline_layout) { + VkDescriptorSetLayout set_layouts[3] = { + device->queries.buf_descriptor_set_layout, + device->queries.buf_descriptor_set_layout, + device->queries.buf_descriptor_set_layout + }; + VkPipelineLayoutCreateInfo pipeline_layout_info = { + .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, + .setLayoutCount = 3, + .pSetLayouts = set_layouts, + .pushConstantRangeCount = 1, + .pPushConstantRanges = + &(VkPushConstantRange) { VK_SHADER_STAGE_COMPUTE_BIT, 0, 16 }, + }; + + result = + v3dv_CreatePipelineLayout(v3dv_device_to_handle(device), + &pipeline_layout_info, + &device->vk.alloc, + &device->queries.copy_pipeline_layout); + + if (result != VK_SUCCESS) + return false; + } + + if (!device->queries.copy_pipeline) { + nir_shader *copy_query_results_cs_nir = get_copy_query_results_cs(); + result = v3dv_create_compute_pipeline_from_nir(device, + copy_query_results_cs_nir, + device->queries.copy_pipeline_layout, + &pipeline); + ralloc_free(copy_query_results_cs_nir); + if (result != VK_SUCCESS) + return false; + + device->queries.copy_pipeline = pipeline; + } + + return true; +} + +static void +destroy_query_pipelines(struct v3dv_device *device) +{ + VkDevice _device = v3dv_device_to_handle(device); + + /* Availability pipeline */ + v3dv_DestroyPipeline(_device, device->queries.avail_pipeline, + &device->vk.alloc); + device->queries.avail_pipeline = VK_NULL_HANDLE; + v3dv_DestroyPipelineLayout(_device, device->queries.avail_pipeline_layout, + &device->vk.alloc); + device->queries.avail_pipeline_layout = VK_NULL_HANDLE; + + /* Reset occlusion pipeline */ + v3dv_DestroyPipeline(_device, device->queries.reset_occlusion_pipeline, + &device->vk.alloc); + device->queries.reset_occlusion_pipeline = VK_NULL_HANDLE; + v3dv_DestroyPipelineLayout(_device, + device->queries.reset_occlusion_pipeline_layout, + &device->vk.alloc); + device->queries.reset_occlusion_pipeline_layout = VK_NULL_HANDLE; + + /* Copy pipeline */ + v3dv_DestroyPipeline(_device, device->queries.copy_pipeline, + &device->vk.alloc); + device->queries.copy_pipeline = VK_NULL_HANDLE; + v3dv_DestroyPipelineLayout(_device, device->queries.copy_pipeline_layout, + &device->vk.alloc); + device->queries.copy_pipeline_layout = VK_NULL_HANDLE; + + v3dv_DestroyDescriptorSetLayout(_device, + device->queries.buf_descriptor_set_layout, + &device->vk.alloc); + device->queries.buf_descriptor_set_layout = VK_NULL_HANDLE; +} + +/** + * Allocates device resources for implementing certain types of queries. + */ +VkResult +v3dv_query_allocate_resources(struct v3dv_device *device) +{ + if (!create_query_pipelines(device)) + return vk_error(device, VK_ERROR_OUT_OF_HOST_MEMORY); + + return VK_SUCCESS; +} + +void +v3dv_query_free_resources(struct v3dv_device *device) +{ + destroy_query_pipelines(device); +} diff --git a/src/broadcom/vulkan/v3dv_queue.c b/src/broadcom/vulkan/v3dv_queue.c index 1b2bee4b5d4..579e4e58e8b 100644 --- a/src/broadcom/vulkan/v3dv_queue.c +++ b/src/broadcom/vulkan/v3dv_queue.c @@ -135,7 +135,7 @@ handle_reset_query_cpu_job(struct v3dv_queue *queue, struct v3dv_job *job, * we handle those in the CPU. */ if (info->pool->query_type == VK_QUERY_TYPE_OCCLUSION) - v3dv_bo_wait(job->device, info->pool->bo, PIPE_TIMEOUT_INFINITE); + v3dv_bo_wait(job->device, info->pool->occlusion.bo, PIPE_TIMEOUT_INFINITE); if (info->pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR) { struct vk_sync_wait waits[info->count]; @@ -160,7 +160,7 @@ handle_reset_query_cpu_job(struct v3dv_queue *queue, struct v3dv_job *job, return result; } - v3dv_reset_query_pools(job->device, info->pool, info->first, info->count); + v3dv_reset_query_pool_cpu(job->device, info->pool, info->first, info->count); return VK_SUCCESS; } @@ -218,12 +218,15 @@ handle_end_query_cpu_job(struct v3dv_job *job, uint32_t counter_pass_idx) mtx_lock(&job->device->query_mutex); - struct v3dv_end_query_cpu_job_info *info = &job->cpu.query_end; + struct v3dv_end_query_info *info = &job->cpu.query_end; struct v3dv_queue *queue = &job->device->queue; int err = 0; int fd = -1; + assert(info->pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR || + info->pool->query_type == VK_QUERY_TYPE_TIMESTAMP); + if (info->pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR) { result = export_perfmon_last_job_sync(queue, job, &fd); @@ -268,6 +271,9 @@ handle_copy_query_results_cpu_job(struct v3dv_job *job) struct v3dv_copy_query_results_cpu_job_info *info = &job->cpu.query_copy_results; + assert(info->pool->query_type == VK_QUERY_TYPE_PERFORMANCE_QUERY_KHR || + info->pool->query_type == VK_QUERY_TYPE_TIMESTAMP); + assert(info->dst && info->dst->mem && info->dst->mem->bo); struct v3dv_bo *bo = info->dst->mem->bo; @@ -278,13 +284,13 @@ handle_copy_query_results_cpu_job(struct v3dv_job *job) uint8_t *offset = ((uint8_t *) bo->map) + info->offset + info->dst->mem_offset; - v3dv_get_query_pool_results(job->device, - info->pool, - info->first, - info->count, - offset, - info->stride, - info->flags); + v3dv_get_query_pool_results_cpu(job->device, + info->pool, + info->first, + info->count, + offset, + info->stride, + info->flags); return VK_SUCCESS; } diff --git a/src/broadcom/vulkan/v3dvx_cmd_buffer.c b/src/broadcom/vulkan/v3dvx_cmd_buffer.c index 02af003e7e3..8ad5ef3bce6 100644 --- a/src/broadcom/vulkan/v3dvx_cmd_buffer.c +++ b/src/broadcom/vulkan/v3dvx_cmd_buffer.c @@ -1598,17 +1598,17 @@ cmd_buffer_copy_secondary_end_query_state(struct v3dv_cmd_buffer *primary, const uint32_t total_state_count = p_state->query.end.used_count + s_state->query.end.used_count; v3dv_cmd_buffer_ensure_array_state(primary, - sizeof(struct v3dv_end_query_cpu_job_info), + sizeof(struct v3dv_end_query_info), total_state_count, &p_state->query.end.alloc_count, (void **) &p_state->query.end.states); v3dv_return_if_oom(primary, NULL); for (uint32_t i = 0; i < s_state->query.end.used_count; i++) { - const struct v3dv_end_query_cpu_job_info *s_qstate = + const struct v3dv_end_query_info *s_qstate = &secondary->state.query.end.states[i]; - struct v3dv_end_query_cpu_job_info *p_qstate = + struct v3dv_end_query_info *p_qstate = &p_state->query.end.states[p_state->query.end.used_count++]; p_qstate->pool = s_qstate->pool;