diff --git a/hw/xbox/nv2a/pgraph/vk/renderer.h b/hw/xbox/nv2a/pgraph/vk/renderer.h index e3aa1839d4..27b5823d79 100644 --- a/hw/xbox/nv2a/pgraph/vk/renderer.h +++ b/hw/xbox/nv2a/pgraph/vk/renderer.h @@ -320,10 +320,16 @@ typedef struct PGRAPHVkDisplayState { GLuint gl_texture_id; } PGRAPHVkDisplayState; +typedef enum { + COMPUTE_TYPE_DEPTH_STENCIL = 0, + COMPUTE_TYPE_DEPTH_STENCIL_DIRECT = 1, +} ComputeType; + typedef struct ComputePipelineKey { VkFormat host_fmt; bool pack; int workgroup_size; + ComputeType compute_type; } ComputePipelineKey; typedef struct ComputePipeline { @@ -340,6 +346,13 @@ typedef struct PGRAPHVkComputeState { VkPipelineLayout pipeline_layout; Lru pipeline_cache; ComputePipeline *pipeline_cache_entries; + + VkDescriptorPool direct_descriptor_pool; + VkDescriptorSetLayout direct_descriptor_set_layout; + VkDescriptorSet direct_descriptor_sets[1024]; + int direct_descriptor_set_index; + VkPipelineLayout direct_pipeline_layout; + VkSampler direct_depth_sampler; } PGRAPHVkComputeState; typedef struct PGRAPHVkState { @@ -579,6 +592,14 @@ void pgraph_vk_finalize_compute(PGRAPHState *pg); void pgraph_vk_pack_depth_stencil(PGRAPHState *pg, SurfaceBinding *surface, VkCommandBuffer cmd, VkBuffer src, VkBuffer dst, bool downscale); +void pgraph_vk_pack_depth_stencil_direct(PGRAPHState *pg, + SurfaceBinding *surface, + VkCommandBuffer cmd, + VkImageView depth_view, + VkBuffer stencil_buf, + VkDeviceSize stencil_offset, + VkDeviceSize stencil_size, + VkBuffer dst, bool downscale); void pgraph_vk_unpack_depth_stencil(PGRAPHState *pg, SurfaceBinding *surface, VkCommandBuffer cmd, VkBuffer src, VkBuffer dst); diff --git a/hw/xbox/nv2a/pgraph/vk/surface-compute.c b/hw/xbox/nv2a/pgraph/vk/surface-compute.c index 7aced75bd9..26552de47f 100644 --- a/hw/xbox/nv2a/pgraph/vk/surface-compute.c +++ b/hw/xbox/nv2a/pgraph/vk/surface-compute.c @@ -118,6 +118,25 @@ const char *unpack_z24s8_to_d32_sfloat_s8_uint_glsl = " }\n" "}\n"; +static const char *pack_depth_stencil_direct_glsl = + "layout(push_constant) uniform PushConstants { uint width_in, width_out; };\n" + "layout(set = 0, binding = 0) uniform sampler2D depth_tex;\n" + "layout(set = 0, binding = 1) readonly buffer StencilIn { uint stencil_in[]; };\n" + "layout(set = 0, binding = 2) writeonly buffer PackedOut { uint packed_out[]; };\n" + "void main() {\n" + " uint idx_out = gl_GlobalInvocationID.x;\n" + " uint scale = width_in / width_out;\n" + " uint out_x = idx_out % width_out;\n" + " uint out_y = idx_out / width_out;\n" + " uint in_x = out_x * scale;\n" + " uint in_y = out_y * scale;\n" + " uint idx_in = in_y * width_in + in_x;\n" + " float depth = texelFetch(depth_tex, ivec2(in_x, in_y), 0).r;\n" + " uint depth_value = uint(depth * float(0xFFFFFF));\n" + " uint stencil_value = (stencil_in[idx_in / 4] >> ((idx_in % 4) * 8)) & 0xFFu;\n" + " packed_out[idx_out] = depth_value << 8 | stencil_value;\n" + "}\n"; + static gchar *get_compute_shader_glsl(VkFormat host_fmt, bool pack, int workgroup_size) { @@ -169,6 +188,32 @@ static void create_descriptor_pool(PGRAPHState *pg) &r->compute.descriptor_pool)); } +static void create_direct_descriptor_pool(PGRAPHState *pg) +{ + PGRAPHVkState *r = pg->vk_renderer_state; + + VkDescriptorPoolSize pool_sizes[] = { + { + .type = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER, + .descriptorCount = ARRAY_SIZE(r->compute.direct_descriptor_sets), + }, + { + .type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, + .descriptorCount = 2 * ARRAY_SIZE(r->compute.direct_descriptor_sets), + }, + }; + + VkDescriptorPoolCreateInfo pool_info = { + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO, + .poolSizeCount = ARRAY_SIZE(pool_sizes), + .pPoolSizes = pool_sizes, + .maxSets = ARRAY_SIZE(r->compute.direct_descriptor_sets), + .flags = VK_DESCRIPTOR_POOL_CREATE_FREE_DESCRIPTOR_SET_BIT, + }; + VK_CHECK(vkCreateDescriptorPool(r->device, &pool_info, NULL, + &r->compute.direct_descriptor_pool)); +} + static void destroy_descriptor_pool(PGRAPHState *pg) { PGRAPHVkState *r = pg->vk_renderer_state; @@ -177,6 +222,14 @@ static void destroy_descriptor_pool(PGRAPHState *pg) r->compute.descriptor_pool = VK_NULL_HANDLE; } +static void destroy_direct_descriptor_pool(PGRAPHState *pg) +{ + PGRAPHVkState *r = pg->vk_renderer_state; + + vkDestroyDescriptorPool(r->device, r->compute.direct_descriptor_pool, NULL); + r->compute.direct_descriptor_pool = VK_NULL_HANDLE; +} + static void create_descriptor_set_layout(PGRAPHState *pg) { PGRAPHVkState *r = pg->vk_renderer_state; @@ -201,6 +254,40 @@ static void create_descriptor_set_layout(PGRAPHState *pg) &r->compute.descriptor_set_layout)); } +static void create_direct_descriptor_set_layout(PGRAPHState *pg) +{ + PGRAPHVkState *r = pg->vk_renderer_state; + + VkDescriptorSetLayoutBinding bindings[] = { + { + .binding = 0, + .descriptorCount = 1, + .descriptorType = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER, + .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT, + }, + { + .binding = 1, + .descriptorCount = 1, + .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, + .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT, + }, + { + .binding = 2, + .descriptorCount = 1, + .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, + .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT, + }, + }; + + VkDescriptorSetLayoutCreateInfo layout_info = { + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO, + .bindingCount = ARRAY_SIZE(bindings), + .pBindings = bindings, + }; + VK_CHECK(vkCreateDescriptorSetLayout( + r->device, &layout_info, NULL, &r->compute.direct_descriptor_set_layout)); +} + static void destroy_descriptor_set_layout(PGRAPHState *pg) { PGRAPHVkState *r = pg->vk_renderer_state; @@ -210,6 +297,16 @@ static void destroy_descriptor_set_layout(PGRAPHState *pg) r->compute.descriptor_set_layout = VK_NULL_HANDLE; } +static void destroy_direct_descriptor_set_layout(PGRAPHState *pg) +{ + PGRAPHVkState *r = pg->vk_renderer_state; + + vkDestroyDescriptorSetLayout(r->device, + r->compute.direct_descriptor_set_layout, + NULL); + r->compute.direct_descriptor_set_layout = VK_NULL_HANDLE; +} + static void create_descriptor_sets(PGRAPHState *pg) { PGRAPHVkState *r = pg->vk_renderer_state; @@ -228,6 +325,24 @@ static void create_descriptor_sets(PGRAPHState *pg) r->compute.descriptor_sets)); } +static void create_direct_descriptor_sets(PGRAPHState *pg) +{ + PGRAPHVkState *r = pg->vk_renderer_state; + + VkDescriptorSetLayout layouts[ARRAY_SIZE(r->compute.direct_descriptor_sets)]; + for (int i = 0; i < ARRAY_SIZE(layouts); i++) { + layouts[i] = r->compute.direct_descriptor_set_layout; + } + VkDescriptorSetAllocateInfo alloc_info = { + .sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO, + .descriptorPool = r->compute.direct_descriptor_pool, + .descriptorSetCount = ARRAY_SIZE(r->compute.direct_descriptor_sets), + .pSetLayouts = layouts, + }; + VK_CHECK(vkAllocateDescriptorSets(r->device, &alloc_info, + r->compute.direct_descriptor_sets)); +} + static void destroy_descriptor_sets(PGRAPHState *pg) { PGRAPHVkState *r = pg->vk_renderer_state; @@ -240,6 +355,18 @@ static void destroy_descriptor_sets(PGRAPHState *pg) } } +static void destroy_direct_descriptor_sets(PGRAPHState *pg) +{ + PGRAPHVkState *r = pg->vk_renderer_state; + + vkFreeDescriptorSets(r->device, r->compute.direct_descriptor_pool, + ARRAY_SIZE(r->compute.direct_descriptor_sets), + r->compute.direct_descriptor_sets); + for (int i = 0; i < ARRAY_SIZE(r->compute.direct_descriptor_sets); i++) { + r->compute.direct_descriptor_sets[i] = VK_NULL_HANDLE; + } +} + static void create_compute_pipeline_layout(PGRAPHState *pg) { PGRAPHVkState *r = pg->vk_renderer_state; @@ -259,20 +386,74 @@ static void create_compute_pipeline_layout(PGRAPHState *pg) &r->compute.pipeline_layout)); } +static void create_direct_compute_pipeline_layout(PGRAPHState *pg) +{ + PGRAPHVkState *r = pg->vk_renderer_state; + + VkPushConstantRange push_constant_range = { + .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT, + .size = 2 * sizeof(uint32_t), + }; + VkPipelineLayoutCreateInfo pipeline_layout_info = { + .sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, + .setLayoutCount = 1, + .pSetLayouts = &r->compute.direct_descriptor_set_layout, + .pushConstantRangeCount = 1, + .pPushConstantRanges = &push_constant_range, + }; + VK_CHECK(vkCreatePipelineLayout( + r->device, &pipeline_layout_info, NULL, &r->compute.direct_pipeline_layout)); +} + static void destroy_compute_pipeline_layout(PGRAPHVkState *r) { vkDestroyPipelineLayout(r->device, r->compute.pipeline_layout, NULL); r->compute.pipeline_layout = VK_NULL_HANDLE; } -static VkPipeline create_compute_pipeline(PGRAPHVkState *r, const char *glsl) +static void destroy_direct_compute_pipeline_layout(PGRAPHVkState *r) +{ + vkDestroyPipelineLayout(r->device, r->compute.direct_pipeline_layout, NULL); + r->compute.direct_pipeline_layout = VK_NULL_HANDLE; +} + +static void create_direct_depth_sampler(PGRAPHVkState *r) +{ + VkSamplerCreateInfo sampler_info = { + .sType = VK_STRUCTURE_TYPE_SAMPLER_CREATE_INFO, + .magFilter = VK_FILTER_NEAREST, + .minFilter = VK_FILTER_NEAREST, + .mipmapMode = VK_SAMPLER_MIPMAP_MODE_NEAREST, + .addressModeU = VK_SAMPLER_ADDRESS_MODE_CLAMP_TO_EDGE, + .addressModeV = VK_SAMPLER_ADDRESS_MODE_CLAMP_TO_EDGE, + .addressModeW = VK_SAMPLER_ADDRESS_MODE_CLAMP_TO_EDGE, + .anisotropyEnable = VK_FALSE, + .compareEnable = VK_FALSE, + .compareOp = VK_COMPARE_OP_ALWAYS, + .minLod = 0.0f, + .maxLod = 0.0f, + .maxAnisotropy = 1.0f, + .borderColor = VK_BORDER_COLOR_FLOAT_OPAQUE_WHITE, + }; + VK_CHECK(vkCreateSampler(r->device, &sampler_info, NULL, + &r->compute.direct_depth_sampler)); +} + +static void destroy_direct_depth_sampler(PGRAPHVkState *r) +{ + vkDestroySampler(r->device, r->compute.direct_depth_sampler, NULL); + r->compute.direct_depth_sampler = VK_NULL_HANDLE; +} + +static VkPipeline create_compute_pipeline(PGRAPHVkState *r, const char *glsl, + VkPipelineLayout layout) { ShaderModuleInfo *module = pgraph_vk_create_shader_module_from_glsl( r, VK_SHADER_STAGE_COMPUTE_BIT, glsl); VkComputePipelineCreateInfo pipeline_info = { .sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO, - .layout = r->compute.pipeline_layout, + .layout = layout, .stage = (VkPipelineShaderStageCreateInfo){ .sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO, @@ -323,13 +504,17 @@ bool pgraph_vk_compute_needs_finish(PGRAPHVkState *r) { bool need_descriptor_write_reset = (r->compute.descriptor_set_index >= ARRAY_SIZE(r->compute.descriptor_sets)); + bool need_direct_descriptor_write_reset = + (r->compute.direct_descriptor_set_index >= + ARRAY_SIZE(r->compute.direct_descriptor_sets)); - return need_descriptor_write_reset; + return need_descriptor_write_reset || need_direct_descriptor_write_reset; } void pgraph_vk_compute_finish_complete(PGRAPHVkState *r) { r->compute.descriptor_set_index = 0; + r->compute.direct_descriptor_set_index = 0; } static int get_workgroup_size_for_output_units(PGRAPHVkState *r, int output_units) @@ -363,6 +548,7 @@ static ComputePipeline *get_compute_pipeline(PGRAPHVkState *r, VkFormat host_fmt key.host_fmt = host_fmt; key.pack = pack; key.workgroup_size = workgroup_size; + key.compute_type = COMPUTE_TYPE_DEPTH_STENCIL; LruNode *node = lru_lookup(&r->compute.pipeline_cache, fast_hash((void *)&key, sizeof(key)), &key); @@ -525,6 +711,107 @@ void pgraph_vk_unpack_depth_stencil(PGRAPHState *pg, SurfaceBinding *surface, pgraph_vk_end_debug_marker(r, cmd); } +void pgraph_vk_pack_depth_stencil_direct(PGRAPHState *pg, + SurfaceBinding *surface, + VkCommandBuffer cmd, + VkImageView depth_view, + VkBuffer stencil_buf, + VkDeviceSize stencil_offset, + VkDeviceSize stencil_size, + VkBuffer dst, bool downscale) +{ + PGRAPHVkState *r = pg->vk_renderer_state; + + unsigned int input_width = surface->width, input_height = surface->height; + pgraph_apply_scaling_factor(pg, &input_width, &input_height); + + unsigned int output_width = surface->width, output_height = surface->height; + if (!downscale) { + pgraph_apply_scaling_factor(pg, &output_width, &output_height); + } + + size_t output_size = output_width * output_height * 4; + + assert(r->compute.direct_descriptor_set_index < + ARRAY_SIZE(r->compute.direct_descriptor_sets)); + + int ds_idx = r->compute.direct_descriptor_set_index++; + VkDescriptorSet ds = r->compute.direct_descriptor_sets[ds_idx]; + + VkDescriptorImageInfo image_info = { + .imageLayout = VK_IMAGE_LAYOUT_DEPTH_STENCIL_READ_ONLY_OPTIMAL, + .imageView = depth_view, + .sampler = r->compute.direct_depth_sampler, + }; + VkDescriptorBufferInfo stencil_buf_info = { + .buffer = stencil_buf, + .offset = stencil_offset, + .range = stencil_size, + }; + VkDescriptorBufferInfo output_buf_info = { + .buffer = dst, + .offset = 0, + .range = output_size, + }; + VkWriteDescriptorSet writes[] = { + { + .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET, + .dstSet = ds, + .dstBinding = 0, + .descriptorType = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER, + .descriptorCount = 1, + .pImageInfo = &image_info, + }, + { + .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET, + .dstSet = ds, + .dstBinding = 1, + .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, + .descriptorCount = 1, + .pBufferInfo = &stencil_buf_info, + }, + { + .sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET, + .dstSet = ds, + .dstBinding = 2, + .descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, + .descriptorCount = 1, + .pBufferInfo = &output_buf_info, + }, + }; + vkUpdateDescriptorSets(r->device, ARRAY_SIZE(writes), writes, 0, NULL); + + size_t output_units = output_width * output_height; + int workgroup_size = get_workgroup_size_for_output_units(r, output_units); + + ComputePipelineKey key; + memset(&key, 0, sizeof(key)); + key.workgroup_size = workgroup_size; + key.compute_type = COMPUTE_TYPE_DEPTH_STENCIL_DIRECT; + + LruNode *node = lru_lookup(&r->compute.pipeline_cache, + fast_hash((void *)&key, sizeof(key)), &key); + ComputePipeline *pipeline = container_of(node, ComputePipeline, node); + assert(pipeline); + + size_t group_count = output_units / workgroup_size; + + pgraph_vk_begin_debug_marker(r, cmd, RGBA_PINK, __func__); + vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline->pipeline); + vkCmdBindDescriptorSets(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, + r->compute.direct_pipeline_layout, 0, 1, &ds, 0, + NULL); + + uint32_t push_constants[2] = { input_width, output_width }; + assert(sizeof(push_constants) == 8); + vkCmdPushConstants(cmd, r->compute.direct_pipeline_layout, + VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(push_constants), + push_constants); + + vkCmdDispatch(cmd, group_count, 1, 1); + pgraph_vk_end_debug_marker(r, cmd); +} + static void pipeline_cache_entry_init(Lru *lru, LruNode *node, const void *state) { @@ -538,10 +825,25 @@ static void pipeline_cache_entry_init(Lru *lru, LruNode *node, "Warning: Needed compute shader with workgroup size = 1\n"); } - gchar *glsl = get_compute_shader_glsl( - snode->key.host_fmt, snode->key.pack, snode->key.workgroup_size); + gchar *glsl; + VkPipelineLayout layout; + switch (snode->key.compute_type) { + case COMPUTE_TYPE_DEPTH_STENCIL_DIRECT: + glsl = g_strdup_printf( + "#version 450\n" + "layout(local_size_x = %d, local_size_y = 1, local_size_z = 1) in;\n" + "%s", snode->key.workgroup_size, pack_depth_stencil_direct_glsl); + layout = r->compute.direct_pipeline_layout; + break; + case COMPUTE_TYPE_DEPTH_STENCIL: + default: + glsl = get_compute_shader_glsl( + snode->key.host_fmt, snode->key.pack, snode->key.workgroup_size); + layout = r->compute.pipeline_layout; + break; + } assert(glsl); - snode->pipeline = create_compute_pipeline(r, glsl); + snode->pipeline = create_compute_pipeline(r, glsl, layout); g_free(glsl); } @@ -595,6 +897,11 @@ void pgraph_vk_init_compute(PGRAPHState *pg) create_descriptor_set_layout(pg); create_descriptor_sets(pg); create_compute_pipeline_layout(pg); + create_direct_descriptor_pool(pg); + create_direct_descriptor_set_layout(pg); + create_direct_descriptor_sets(pg); + create_direct_compute_pipeline_layout(pg); + create_direct_depth_sampler(r); pipeline_cache_init(r); } @@ -605,8 +912,13 @@ void pgraph_vk_finalize_compute(PGRAPHState *pg) assert(!r->in_command_buffer); pipeline_cache_finalize(r); + destroy_direct_depth_sampler(r); + destroy_direct_compute_pipeline_layout(r); destroy_compute_pipeline_layout(r); + destroy_direct_descriptor_sets(pg); destroy_descriptor_sets(pg); + destroy_direct_descriptor_set_layout(pg); destroy_descriptor_set_layout(pg); + destroy_direct_descriptor_pool(pg); destroy_descriptor_pool(pg); } diff --git a/hw/xbox/nv2a/pgraph/vk/texture.c b/hw/xbox/nv2a/pgraph/vk/texture.c index 3ed3487bf7..63c04bad82 100644 --- a/hw/xbox/nv2a/pgraph/vk/texture.c +++ b/hw/xbox/nv2a/pgraph/vk/texture.c @@ -711,65 +711,206 @@ static void copy_zeta_surface_to_texture(PGRAPHState *pg, SurfaceBinding *surfac scaled_height = surface->height; pgraph_apply_scaling_factor(pg, &scaled_width, &scaled_height); - size_t copied_image_size = - scaled_width * scaled_height * surface->host_fmt.host_bytes_per_pixel; - size_t stencil_buffer_offset = 0; - size_t stencil_buffer_size = 0; + StorageBuffer *dst_storage_buffer = &r->storage_buffers[BUFFER_COMPUTE_DST]; + VkBuffer texture_source_buffer; - int num_regions = 0; - VkBufferImageCopy regions[2]; - regions[num_regions++] = (VkBufferImageCopy){ - .bufferOffset = 0, - .bufferRowLength = 0, // Tightly packed - .bufferImageHeight = 0, // Tightly packed - .imageSubresource.aspectMask = surface->color ? VK_IMAGE_ASPECT_COLOR_BIT : VK_IMAGE_ASPECT_DEPTH_BIT, - .imageSubresource.mipLevel = 0, - .imageSubresource.baseArrayLayer = 0, - .imageSubresource.layerCount = 1, - .imageOffset = (VkOffset3D){0, 0, 0}, - .imageExtent = (VkExtent3D){scaled_width, scaled_height, 1}, - }; - - if (surface->host_fmt.aspect & VK_IMAGE_ASPECT_STENCIL_BIT) { - stencil_buffer_offset = + if (use_compute_to_convert_depth_stencil && + (surface->host_fmt.aspect & VK_IMAGE_ASPECT_STENCIL_BIT)) { + size_t packed_image_size = scaled_width * scaled_height * 4; + size_t stencil_buffer_offset = ROUND_UP(scaled_width * scaled_height * 4, r->device_props.limits.minStorageBufferOffsetAlignment); - stencil_buffer_size = scaled_width * scaled_height; - copied_image_size += stencil_buffer_size; + size_t stencil_buffer_size = scaled_width * scaled_height; + assert(dst_storage_buffer->buffer_size >= + stencil_buffer_offset + stencil_buffer_size); - regions[num_regions++] = (VkBufferImageCopy){ + VkBufferImageCopy stencil_region = { .bufferOffset = stencil_buffer_offset, - .bufferRowLength = 0, // Tightly packed - .bufferImageHeight = 0, // Tightly packed + .bufferRowLength = 0, + .bufferImageHeight = 0, .imageSubresource.aspectMask = VK_IMAGE_ASPECT_STENCIL_BIT, .imageSubresource.mipLevel = 0, .imageSubresource.baseArrayLayer = 0, .imageSubresource.layerCount = 1, - .imageOffset = (VkOffset3D){0, 0, 0}, - .imageExtent = (VkExtent3D){scaled_width, scaled_height, 1}, + .imageOffset = (VkOffset3D){ 0, 0, 0 }, + .imageExtent = (VkExtent3D){ scaled_width, scaled_height, 1 }, }; - } - StorageBuffer *dst_storage_buffer = &r->storage_buffers[BUFFER_COMPUTE_DST]; - assert(dst_storage_buffer->buffer_size >= copied_image_size); - pgraph_vk_transition_image_layout( - pg, cmd, surface->image, surface->host_fmt.vk_format, - VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL, - VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL); + pgraph_vk_transition_image_layout( + pg, cmd, surface->image, surface->host_fmt.vk_format, + VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL, + VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL); - vkCmdCopyImageToBuffer( - cmd, surface->image, VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, - dst_storage_buffer->buffer, - num_regions, regions); + vkCmdCopyImageToBuffer(cmd, surface->image, + VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, + dst_storage_buffer->buffer, 1, &stencil_region); - pgraph_vk_transition_image_layout( - pg, cmd, surface->image, surface->host_fmt.vk_format, - VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, - VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL); + VkImageMemoryBarrier depth_read_barrier = { + .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER, + .oldLayout = VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, + .newLayout = VK_IMAGE_LAYOUT_DEPTH_STENCIL_READ_ONLY_OPTIMAL, + .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .image = surface->image, + .subresourceRange = { + .aspectMask = VK_IMAGE_ASPECT_DEPTH_BIT, + .baseMipLevel = 0, + .levelCount = 1, + .baseArrayLayer = 0, + .layerCount = 1, + }, + .srcAccessMask = VK_ACCESS_TRANSFER_READ_BIT, + .dstAccessMask = VK_ACCESS_SHADER_READ_BIT, + }; + vkCmdPipelineBarrier(cmd, VK_PIPELINE_STAGE_TRANSFER_BIT, + VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 0, NULL, + 0, NULL, 1, &depth_read_barrier); - VkBuffer texture_source_buffer; + VkImageViewCreateInfo depth_view_info = { + .sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO, + .image = surface->image, + .viewType = VK_IMAGE_VIEW_TYPE_2D, + .format = surface->host_fmt.vk_format, + .subresourceRange = { + .aspectMask = VK_IMAGE_ASPECT_DEPTH_BIT, + .baseMipLevel = 0, + .levelCount = 1, + .baseArrayLayer = 0, + .layerCount = 1, + }, + }; + VkImageView depth_view; + VK_CHECK(vkCreateImageView(r->device, &depth_view_info, NULL, + &depth_view)); + + VkBufferMemoryBarrier stencil_barrier = { + .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER, + .srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT, + .dstAccessMask = VK_ACCESS_SHADER_READ_BIT, + .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .buffer = dst_storage_buffer->buffer, + .offset = stencil_buffer_offset, + .size = stencil_buffer_size, + }; + VkBufferMemoryBarrier output_pre_barrier = { + .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER, + .srcAccessMask = VK_ACCESS_TRANSFER_READ_BIT, + .dstAccessMask = VK_ACCESS_SHADER_WRITE_BIT, + .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .buffer = r->storage_buffers[BUFFER_COMPUTE_SRC].buffer, + .size = packed_image_size, + }; + VkBufferMemoryBarrier pre_barriers[] = { + stencil_barrier, + output_pre_barrier, + }; + vkCmdPipelineBarrier(cmd, VK_PIPELINE_STAGE_TRANSFER_BIT, + VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 0, NULL, + ARRAY_SIZE(pre_barriers), pre_barriers, 0, NULL); + + pgraph_vk_pack_depth_stencil_direct( + pg, surface, cmd, depth_view, dst_storage_buffer->buffer, + stencil_buffer_offset, stencil_buffer_size, + r->storage_buffers[BUFFER_COMPUTE_SRC].buffer, false); + + VkBufferMemoryBarrier post_compute_barrier = { + .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER, + .srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT, + .dstAccessMask = VK_ACCESS_TRANSFER_READ_BIT, + .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .buffer = r->storage_buffers[BUFFER_COMPUTE_SRC].buffer, + .size = packed_image_size, + }; + vkCmdPipelineBarrier(cmd, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, + VK_PIPELINE_STAGE_TRANSFER_BIT, 0, 0, NULL, 1, + &post_compute_barrier, 0, NULL); + + VkImageMemoryBarrier depth_restore_barrier = { + .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER, + .oldLayout = VK_IMAGE_LAYOUT_DEPTH_STENCIL_READ_ONLY_OPTIMAL, + .newLayout = VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL, + .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .image = surface->image, + .subresourceRange = { + .aspectMask = VK_IMAGE_ASPECT_DEPTH_BIT | + VK_IMAGE_ASPECT_STENCIL_BIT, + .baseMipLevel = 0, + .levelCount = 1, + .baseArrayLayer = 0, + .layerCount = 1, + }, + .srcAccessMask = VK_ACCESS_SHADER_READ_BIT, + .dstAccessMask = VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_READ_BIT | + VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_WRITE_BIT, + }; + vkCmdPipelineBarrier( + cmd, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, + VK_PIPELINE_STAGE_EARLY_FRAGMENT_TESTS_BIT | + VK_PIPELINE_STAGE_LATE_FRAGMENT_TESTS_BIT, + 0, 0, NULL, 0, NULL, 1, &depth_restore_barrier); + + vkDestroyImageView(r->device, depth_view, NULL); + texture_source_buffer = r->storage_buffers[BUFFER_COMPUTE_SRC].buffer; + } else if (use_compute_to_convert_depth_stencil) { + size_t copied_image_size = + scaled_width * scaled_height * surface->host_fmt.host_bytes_per_pixel; + size_t stencil_buffer_offset = 0; + size_t stencil_buffer_size = 0; + + int num_regions = 0; + VkBufferImageCopy regions[2]; + regions[num_regions++] = (VkBufferImageCopy){ + .bufferOffset = 0, + .bufferRowLength = 0, + .bufferImageHeight = 0, + .imageSubresource.aspectMask = VK_IMAGE_ASPECT_DEPTH_BIT, + .imageSubresource.mipLevel = 0, + .imageSubresource.baseArrayLayer = 0, + .imageSubresource.layerCount = 1, + .imageOffset = (VkOffset3D){ 0, 0, 0 }, + .imageExtent = (VkExtent3D){ scaled_width, scaled_height, 1 }, + }; + + if (surface->host_fmt.aspect & VK_IMAGE_ASPECT_STENCIL_BIT) { + stencil_buffer_offset = + ROUND_UP(scaled_width * scaled_height * 4, + r->device_props.limits.minStorageBufferOffsetAlignment); + stencil_buffer_size = scaled_width * scaled_height; + copied_image_size += stencil_buffer_size; + + regions[num_regions++] = (VkBufferImageCopy){ + .bufferOffset = stencil_buffer_offset, + .bufferRowLength = 0, + .bufferImageHeight = 0, + .imageSubresource.aspectMask = VK_IMAGE_ASPECT_STENCIL_BIT, + .imageSubresource.mipLevel = 0, + .imageSubresource.baseArrayLayer = 0, + .imageSubresource.layerCount = 1, + .imageOffset = (VkOffset3D){ 0, 0, 0 }, + .imageExtent = (VkExtent3D){ scaled_width, scaled_height, 1 }, + }; + } + assert(dst_storage_buffer->buffer_size >= copied_image_size); + + pgraph_vk_transition_image_layout( + pg, cmd, surface->image, surface->host_fmt.vk_format, + VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL, + VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL); + + vkCmdCopyImageToBuffer(cmd, surface->image, + VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, + dst_storage_buffer->buffer, num_regions, + regions); + + pgraph_vk_transition_image_layout( + pg, cmd, surface->image, surface->host_fmt.vk_format, + VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, + VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL); - if (use_compute_to_convert_depth_stencil) { size_t packed_image_size = scaled_width * scaled_height * 4; VkBufferMemoryBarrier pre_pack_src_barrier = { @@ -831,6 +972,38 @@ static void copy_zeta_surface_to_texture(PGRAPHState *pg, SurfaceBinding *surfac texture_source_buffer = r->storage_buffers[BUFFER_COMPUTE_SRC].buffer; } else { + size_t copied_image_size = + scaled_width * scaled_height * surface->host_fmt.host_bytes_per_pixel; + int num_regions = 0; + VkBufferImageCopy regions[2]; + regions[num_regions++] = (VkBufferImageCopy){ + .bufferOffset = 0, + .bufferRowLength = 0, + .bufferImageHeight = 0, + .imageSubresource.aspectMask = VK_IMAGE_ASPECT_DEPTH_BIT, + .imageSubresource.mipLevel = 0, + .imageSubresource.baseArrayLayer = 0, + .imageSubresource.layerCount = 1, + .imageOffset = (VkOffset3D){ 0, 0, 0 }, + .imageExtent = (VkExtent3D){ scaled_width, scaled_height, 1 }, + }; + assert(dst_storage_buffer->buffer_size >= copied_image_size); + + pgraph_vk_transition_image_layout( + pg, cmd, surface->image, surface->host_fmt.vk_format, + VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL, + VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL); + + vkCmdCopyImageToBuffer(cmd, surface->image, + VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, + dst_storage_buffer->buffer, num_regions, + regions); + + pgraph_vk_transition_image_layout( + pg, cmd, surface->image, surface->host_fmt.vk_format, + VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, + VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL); + VkBufferMemoryBarrier barrier = { .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER, .srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT, @@ -852,7 +1025,7 @@ static void copy_zeta_surface_to_texture(PGRAPHState *pg, SurfaceBinding *surfac VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL); texture->current_layout = VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL; - regions[0] = (VkBufferImageCopy){ + VkBufferImageCopy output_region = { .bufferOffset = 0, .bufferRowLength = 0, .bufferImageHeight = 0, @@ -865,7 +1038,7 @@ static void copy_zeta_surface_to_texture(PGRAPHState *pg, SurfaceBinding *surfac }; vkCmdCopyBufferToImage( cmd, texture_source_buffer, texture->image, - VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, 1, regions); + VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, 1, &output_region); VkBufferMemoryBarrier post_copy_barrier = { .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER,