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;