Path: blob/21.2-virgl/src/amd/vulkan/radv_meta_buffer.c
7129 views
#include "nir/nir_builder.h"1#include "radv_meta.h"23#include "radv_cs.h"4#include "sid.h"56static nir_shader *7build_buffer_fill_shader(struct radv_device *dev)8{9nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_buffer_fill");10b.shader->info.workgroup_size[0] = 64;11b.shader->info.workgroup_size[1] = 1;12b.shader->info.workgroup_size[2] = 1;1314nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);15nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);16nir_ssa_def *block_size =17nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1],18b.shader->info.workgroup_size[2], 0);1920nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);2122nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16));23offset = nir_channel(&b, offset, 0);2425nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0);2627nir_ssa_def *load = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .range = 4);28nir_ssa_def *swizzled_load = nir_swizzle(&b, load, (unsigned[]){0, 0, 0, 0}, 4);2930nir_store_ssbo(&b, swizzled_load, dst_buf, offset, .write_mask = 0xf,31.access = ACCESS_NON_READABLE, .align_mul = 16);3233return b.shader;34}3536static nir_shader *37build_buffer_copy_shader(struct radv_device *dev)38{39nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, NULL, "meta_buffer_copy");40b.shader->info.workgroup_size[0] = 64;41b.shader->info.workgroup_size[1] = 1;42b.shader->info.workgroup_size[2] = 1;4344nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);45nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);46nir_ssa_def *block_size =47nir_imm_ivec4(&b, b.shader->info.workgroup_size[0], b.shader->info.workgroup_size[1],48b.shader->info.workgroup_size[2], 0);4950nir_ssa_def *global_id = nir_iadd(&b, nir_imul(&b, wg_id, block_size), invoc_id);5152nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16));53offset = nir_channel(&b, offset, 0);5455nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0);56nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1);5758nir_ssa_def *load = nir_load_ssbo(&b, 4, 32, src_buf, offset, .align_mul = 16);59nir_store_ssbo(&b, load, dst_buf, offset, .write_mask = 0xf, .access = ACCESS_NON_READABLE,60.align_mul = 16);6162return b.shader;63}6465VkResult66radv_device_init_meta_buffer_state(struct radv_device *device)67{68VkResult result;69nir_shader *fill_cs = build_buffer_fill_shader(device);70nir_shader *copy_cs = build_buffer_copy_shader(device);7172VkDescriptorSetLayoutCreateInfo fill_ds_create_info = {73.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,74.flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,75.bindingCount = 1,76.pBindings = (VkDescriptorSetLayoutBinding[]){77{.binding = 0,78.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,79.descriptorCount = 1,80.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,81.pImmutableSamplers = NULL},82}};8384result = radv_CreateDescriptorSetLayout(radv_device_to_handle(device), &fill_ds_create_info,85&device->meta_state.alloc,86&device->meta_state.buffer.fill_ds_layout);87if (result != VK_SUCCESS)88goto fail;8990VkDescriptorSetLayoutCreateInfo copy_ds_create_info = {91.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO,92.flags = VK_DESCRIPTOR_SET_LAYOUT_CREATE_PUSH_DESCRIPTOR_BIT_KHR,93.bindingCount = 2,94.pBindings = (VkDescriptorSetLayoutBinding[]){95{.binding = 0,96.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,97.descriptorCount = 1,98.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,99.pImmutableSamplers = NULL},100{.binding = 1,101.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,102.descriptorCount = 1,103.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,104.pImmutableSamplers = NULL},105}};106107result = radv_CreateDescriptorSetLayout(radv_device_to_handle(device), ©_ds_create_info,108&device->meta_state.alloc,109&device->meta_state.buffer.copy_ds_layout);110if (result != VK_SUCCESS)111goto fail;112113VkPipelineLayoutCreateInfo fill_pl_create_info = {114.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,115.setLayoutCount = 1,116.pSetLayouts = &device->meta_state.buffer.fill_ds_layout,117.pushConstantRangeCount = 1,118.pPushConstantRanges = &(VkPushConstantRange){VK_SHADER_STAGE_COMPUTE_BIT, 0, 4},119};120121result = radv_CreatePipelineLayout(radv_device_to_handle(device), &fill_pl_create_info,122&device->meta_state.alloc,123&device->meta_state.buffer.fill_p_layout);124if (result != VK_SUCCESS)125goto fail;126127VkPipelineLayoutCreateInfo copy_pl_create_info = {128.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO,129.setLayoutCount = 1,130.pSetLayouts = &device->meta_state.buffer.copy_ds_layout,131.pushConstantRangeCount = 0,132};133134result = radv_CreatePipelineLayout(radv_device_to_handle(device), ©_pl_create_info,135&device->meta_state.alloc,136&device->meta_state.buffer.copy_p_layout);137if (result != VK_SUCCESS)138goto fail;139140VkPipelineShaderStageCreateInfo fill_pipeline_shader_stage = {141.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,142.stage = VK_SHADER_STAGE_COMPUTE_BIT,143.module = vk_shader_module_handle_from_nir(fill_cs),144.pName = "main",145.pSpecializationInfo = NULL,146};147148VkComputePipelineCreateInfo fill_vk_pipeline_info = {149.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,150.stage = fill_pipeline_shader_stage,151.flags = 0,152.layout = device->meta_state.buffer.fill_p_layout,153};154155result = radv_CreateComputePipelines(156radv_device_to_handle(device), radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,157&fill_vk_pipeline_info, NULL, &device->meta_state.buffer.fill_pipeline);158if (result != VK_SUCCESS)159goto fail;160161VkPipelineShaderStageCreateInfo copy_pipeline_shader_stage = {162.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,163.stage = VK_SHADER_STAGE_COMPUTE_BIT,164.module = vk_shader_module_handle_from_nir(copy_cs),165.pName = "main",166.pSpecializationInfo = NULL,167};168169VkComputePipelineCreateInfo copy_vk_pipeline_info = {170.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,171.stage = copy_pipeline_shader_stage,172.flags = 0,173.layout = device->meta_state.buffer.copy_p_layout,174};175176result = radv_CreateComputePipelines(177radv_device_to_handle(device), radv_pipeline_cache_to_handle(&device->meta_state.cache), 1,178©_vk_pipeline_info, NULL, &device->meta_state.buffer.copy_pipeline);179if (result != VK_SUCCESS)180goto fail;181182ralloc_free(fill_cs);183ralloc_free(copy_cs);184return VK_SUCCESS;185fail:186radv_device_finish_meta_buffer_state(device);187ralloc_free(fill_cs);188ralloc_free(copy_cs);189return result;190}191192void193radv_device_finish_meta_buffer_state(struct radv_device *device)194{195struct radv_meta_state *state = &device->meta_state;196197radv_DestroyPipeline(radv_device_to_handle(device), state->buffer.copy_pipeline, &state->alloc);198radv_DestroyPipeline(radv_device_to_handle(device), state->buffer.fill_pipeline, &state->alloc);199radv_DestroyPipelineLayout(radv_device_to_handle(device), state->buffer.copy_p_layout,200&state->alloc);201radv_DestroyPipelineLayout(radv_device_to_handle(device), state->buffer.fill_p_layout,202&state->alloc);203radv_DestroyDescriptorSetLayout(radv_device_to_handle(device), state->buffer.copy_ds_layout,204&state->alloc);205radv_DestroyDescriptorSetLayout(radv_device_to_handle(device), state->buffer.fill_ds_layout,206&state->alloc);207}208209static void210fill_buffer_shader(struct radv_cmd_buffer *cmd_buffer, struct radeon_winsys_bo *bo, uint64_t offset,211uint64_t size, uint32_t value)212{213struct radv_device *device = cmd_buffer->device;214uint64_t block_count = round_up_u64(size, 1024);215struct radv_meta_saved_state saved_state;216217radv_meta_save(218&saved_state, cmd_buffer,219RADV_META_SAVE_COMPUTE_PIPELINE | RADV_META_SAVE_CONSTANTS | RADV_META_SAVE_DESCRIPTORS);220221struct radv_buffer dst_buffer = {.bo = bo, .offset = offset, .size = size};222223radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE,224device->meta_state.buffer.fill_pipeline);225226radv_meta_push_descriptor_set(227cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, device->meta_state.buffer.fill_p_layout,2280, /* set */2291, /* descriptorWriteCount */230(VkWriteDescriptorSet[]){231{.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,232.dstBinding = 0,233.dstArrayElement = 0,234.descriptorCount = 1,235.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,236.pBufferInfo = &(VkDescriptorBufferInfo){.buffer = radv_buffer_to_handle(&dst_buffer),237.offset = 0,238.range = size}}});239240radv_CmdPushConstants(radv_cmd_buffer_to_handle(cmd_buffer),241device->meta_state.buffer.fill_p_layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, 4,242&value);243244radv_CmdDispatch(radv_cmd_buffer_to_handle(cmd_buffer), block_count, 1, 1);245246radv_meta_restore(&saved_state, cmd_buffer);247}248249static void250copy_buffer_shader(struct radv_cmd_buffer *cmd_buffer, struct radeon_winsys_bo *src_bo,251struct radeon_winsys_bo *dst_bo, uint64_t src_offset, uint64_t dst_offset,252uint64_t size)253{254struct radv_device *device = cmd_buffer->device;255uint64_t block_count = round_up_u64(size, 1024);256struct radv_meta_saved_state saved_state;257258radv_meta_save(&saved_state, cmd_buffer,259RADV_META_SAVE_COMPUTE_PIPELINE | RADV_META_SAVE_DESCRIPTORS);260261struct radv_buffer dst_buffer = {.bo = dst_bo, .offset = dst_offset, .size = size};262263struct radv_buffer src_buffer = {.bo = src_bo, .offset = src_offset, .size = size};264265radv_CmdBindPipeline(radv_cmd_buffer_to_handle(cmd_buffer), VK_PIPELINE_BIND_POINT_COMPUTE,266device->meta_state.buffer.copy_pipeline);267268radv_meta_push_descriptor_set(269cmd_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, device->meta_state.buffer.copy_p_layout,2700, /* set */2712, /* descriptorWriteCount */272(VkWriteDescriptorSet[]){273{.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,274.dstBinding = 0,275.dstArrayElement = 0,276.descriptorCount = 1,277.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,278.pBufferInfo = &(VkDescriptorBufferInfo){.buffer = radv_buffer_to_handle(&dst_buffer),279.offset = 0,280.range = size}},281{.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET,282.dstBinding = 1,283.dstArrayElement = 0,284.descriptorCount = 1,285.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,286.pBufferInfo = &(VkDescriptorBufferInfo){.buffer = radv_buffer_to_handle(&src_buffer),287.offset = 0,288.range = size}}});289290radv_CmdDispatch(radv_cmd_buffer_to_handle(cmd_buffer), block_count, 1, 1);291292radv_meta_restore(&saved_state, cmd_buffer);293}294295static bool296radv_prefer_compute_dma(const struct radv_device *device, uint64_t size,297struct radeon_winsys_bo *src_bo, struct radeon_winsys_bo *dst_bo)298{299bool use_compute = size >= RADV_BUFFER_OPS_CS_THRESHOLD;300301if (device->physical_device->rad_info.chip_class >= GFX10 &&302device->physical_device->rad_info.has_dedicated_vram) {303if ((src_bo && !(src_bo->initial_domain & RADEON_DOMAIN_VRAM)) ||304!(dst_bo->initial_domain & RADEON_DOMAIN_VRAM)) {305/* Prefer CP DMA for GTT on dGPUS due to slow PCIe. */306use_compute = false;307}308}309310return use_compute;311}312313uint32_t314radv_fill_buffer(struct radv_cmd_buffer *cmd_buffer, const struct radv_image *image,315struct radeon_winsys_bo *bo, uint64_t offset, uint64_t size, uint32_t value)316{317bool use_compute = radv_prefer_compute_dma(cmd_buffer->device, size, NULL, bo);318uint32_t flush_bits = 0;319320assert(!(offset & 3));321assert(!(size & 3));322323if (use_compute) {324cmd_buffer->state.flush_bits |=325radv_dst_access_flush(cmd_buffer, VK_ACCESS_SHADER_WRITE_BIT, image);326327fill_buffer_shader(cmd_buffer, bo, offset, size, value);328329flush_bits = RADV_CMD_FLAG_CS_PARTIAL_FLUSH | RADV_CMD_FLAG_INV_VCACHE |330radv_src_access_flush(cmd_buffer, VK_ACCESS_SHADER_WRITE_BIT, image);331} else if (size) {332uint64_t va = radv_buffer_get_va(bo);333va += offset;334radv_cs_add_buffer(cmd_buffer->device->ws, cmd_buffer->cs, bo);335si_cp_dma_clear_buffer(cmd_buffer, va, size, value);336}337338return flush_bits;339}340341static void342radv_copy_buffer(struct radv_cmd_buffer *cmd_buffer, struct radeon_winsys_bo *src_bo,343struct radeon_winsys_bo *dst_bo, uint64_t src_offset, uint64_t dst_offset,344uint64_t size)345{346bool use_compute = !(size & 3) && !(src_offset & 3) && !(dst_offset & 3) &&347radv_prefer_compute_dma(cmd_buffer->device, size, src_bo, dst_bo);348349if (use_compute)350copy_buffer_shader(cmd_buffer, src_bo, dst_bo, src_offset, dst_offset, size);351else if (size) {352uint64_t src_va = radv_buffer_get_va(src_bo);353uint64_t dst_va = radv_buffer_get_va(dst_bo);354src_va += src_offset;355dst_va += dst_offset;356357radv_cs_add_buffer(cmd_buffer->device->ws, cmd_buffer->cs, src_bo);358radv_cs_add_buffer(cmd_buffer->device->ws, cmd_buffer->cs, dst_bo);359360si_cp_dma_buffer_copy(cmd_buffer, src_va, dst_va, size);361}362}363364void365radv_CmdFillBuffer(VkCommandBuffer commandBuffer, VkBuffer dstBuffer, VkDeviceSize dstOffset,366VkDeviceSize fillSize, uint32_t data)367{368RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);369RADV_FROM_HANDLE(radv_buffer, dst_buffer, dstBuffer);370371if (fillSize == VK_WHOLE_SIZE)372fillSize = (dst_buffer->size - dstOffset) & ~3ull;373374radv_fill_buffer(cmd_buffer, NULL, dst_buffer->bo, dst_buffer->offset + dstOffset, fillSize,375data);376}377378static void379copy_buffer(struct radv_cmd_buffer *cmd_buffer, struct radv_buffer *src_buffer,380struct radv_buffer *dst_buffer, const VkBufferCopy2KHR *region)381{382bool old_predicating;383384/* VK_EXT_conditional_rendering says that copy commands should not be385* affected by conditional rendering.386*/387old_predicating = cmd_buffer->state.predicating;388cmd_buffer->state.predicating = false;389390radv_copy_buffer(cmd_buffer, src_buffer->bo, dst_buffer->bo,391src_buffer->offset + region->srcOffset, dst_buffer->offset + region->dstOffset,392region->size);393394/* Restore conditional rendering. */395cmd_buffer->state.predicating = old_predicating;396}397398void399radv_CmdCopyBuffer2KHR(VkCommandBuffer commandBuffer, const VkCopyBufferInfo2KHR *pCopyBufferInfo)400{401RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);402RADV_FROM_HANDLE(radv_buffer, src_buffer, pCopyBufferInfo->srcBuffer);403RADV_FROM_HANDLE(radv_buffer, dst_buffer, pCopyBufferInfo->dstBuffer);404405for (unsigned r = 0; r < pCopyBufferInfo->regionCount; r++) {406copy_buffer(cmd_buffer, src_buffer, dst_buffer, &pCopyBufferInfo->pRegions[r]);407}408}409410void411radv_CmdUpdateBuffer(VkCommandBuffer commandBuffer, VkBuffer dstBuffer, VkDeviceSize dstOffset,412VkDeviceSize dataSize, const void *pData)413{414RADV_FROM_HANDLE(radv_cmd_buffer, cmd_buffer, commandBuffer);415RADV_FROM_HANDLE(radv_buffer, dst_buffer, dstBuffer);416bool mec = radv_cmd_buffer_uses_mec(cmd_buffer);417uint64_t words = dataSize / 4;418uint64_t va = radv_buffer_get_va(dst_buffer->bo);419va += dstOffset + dst_buffer->offset;420421assert(!(dataSize & 3));422assert(!(va & 3));423424if (!dataSize)425return;426427if (dataSize < RADV_BUFFER_UPDATE_THRESHOLD) {428si_emit_cache_flush(cmd_buffer);429430radv_cs_add_buffer(cmd_buffer->device->ws, cmd_buffer->cs, dst_buffer->bo);431432radeon_check_space(cmd_buffer->device->ws, cmd_buffer->cs, words + 4);433434radeon_emit(cmd_buffer->cs, PKT3(PKT3_WRITE_DATA, 2 + words, 0));435radeon_emit(cmd_buffer->cs, S_370_DST_SEL(mec ? V_370_MEM : V_370_MEM_GRBM) |436S_370_WR_CONFIRM(1) | S_370_ENGINE_SEL(V_370_ME));437radeon_emit(cmd_buffer->cs, va);438radeon_emit(cmd_buffer->cs, va >> 32);439radeon_emit_array(cmd_buffer->cs, pData, words);440441if (unlikely(cmd_buffer->device->trace_bo))442radv_cmd_buffer_trace_emit(cmd_buffer);443} else {444uint32_t buf_offset;445radv_cmd_buffer_upload_data(cmd_buffer, dataSize, pData, &buf_offset);446radv_copy_buffer(cmd_buffer, cmd_buffer->upload.upload_bo, dst_buffer->bo, buf_offset,447dstOffset + dst_buffer->offset, dataSize);448}449}450451452