Signed-off-by: Philip Rebohle philip.rebohle@tu-dortmund.de --- libs/vkd3d/command.c | 132 ++++++++++++++++++++++++++++++++++++- libs/vkd3d/vkd3d_private.h | 5 ++ 2 files changed, 135 insertions(+), 2 deletions(-)
diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c index 297054b..3648ea4 100644 --- a/libs/vkd3d/command.c +++ b/libs/vkd3d/command.c @@ -4803,6 +4803,128 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearRenderTargetView(ID3D12Gra &clear_value, rect_count, rects); }
+static void d3d12_command_list_clear_unordered_access_view(struct d3d12_command_list *list, + struct d3d12_resource *resource, struct vkd3d_view *view, const VkClearColorValue *clear_color, + UINT rect_count, const D3D12_RECT *rects) +{ + const struct vkd3d_vk_device_procs *vk_procs = &list->device->vk_procs; + unsigned int i, miplevel_idx, layer_count; + struct vkd3d_clear_uav_pipeline pipeline; + struct vkd3d_clear_uav_args clear_args; + VkDescriptorImageInfo image_info; + D3D12_RECT full_rect, curr_rect; + VkWriteDescriptorSet write_set; + VkExtent3D workgroup_size; + + d3d12_command_list_track_resource_usage(list, resource); + d3d12_command_list_end_current_render_pass(list); + + d3d12_command_list_invalidate_current_pipeline(list); + d3d12_command_list_invalidate_bindings(list, list->state); + d3d12_command_list_invalidate_root_parameters(list, VK_PIPELINE_BIND_POINT_COMPUTE); + + if (!d3d12_command_allocator_add_view(list->allocator, view)) + { + WARN("Failed to add view.\n"); + } + + clear_args.clear_color = *clear_color; + + write_set.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET; + write_set.pNext = NULL; + write_set.dstBinding = 0; + write_set.dstArrayElement = 0; + write_set.descriptorCount = 1; + + if (d3d12_resource_is_texture(resource)) + { + image_info.sampler = VK_NULL_HANDLE; + image_info.imageView = view->u.vk_image_view; + image_info.imageLayout = VK_IMAGE_LAYOUT_GENERAL; + + write_set.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE; + write_set.pImageInfo = &image_info; + write_set.pBufferInfo = NULL; + write_set.pTexelBufferView = NULL; + + miplevel_idx = view->info.texture.miplevel_idx; + layer_count = view->info.texture.vk_view_type == VK_IMAGE_VIEW_TYPE_3D + ? d3d12_resource_desc_get_depth(&resource->desc, miplevel_idx) + : view->info.texture.layer_count; + pipeline = vkd3d_clear_uav_ops_get_clear_image_pipeline( + &list->device->meta_ops.clear_uav, view->info.texture.vk_view_type, + view->format->type == VKD3D_FORMAT_TYPE_UINT); + workgroup_size = vkd3d_get_clear_image_uav_workgroup_size(view->info.texture.vk_view_type); + } + else + { + write_set.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER; + write_set.pImageInfo = NULL; + write_set.pBufferInfo = NULL; + write_set.pTexelBufferView = &view->u.vk_buffer_view; + + miplevel_idx = 0; + layer_count = 1; + pipeline = vkd3d_clear_uav_ops_get_clear_buffer_pipeline( + &list->device->meta_ops.clear_uav, + view->format->type == VKD3D_FORMAT_TYPE_UINT); + workgroup_size = vkd3d_get_clear_buffer_uav_workgroup_size(); + } + + if (!(write_set.dstSet = d3d12_command_allocator_allocate_descriptor_set( + list->allocator, pipeline.vk_set_layout))) + { + ERR("Failed to allocate descriptor set.\n"); + return; + } + + VK_CALL(vkUpdateDescriptorSets(list->device->vk_device, 1, &write_set, 0, NULL)); + + full_rect.left = 0; + full_rect.right = d3d12_resource_desc_get_width(&resource->desc, miplevel_idx); + full_rect.top = 0; + full_rect.bottom = d3d12_resource_desc_get_height(&resource->desc, miplevel_idx); + + /* clear full resource if no rects are specified */ + curr_rect = full_rect; + + VK_CALL(vkCmdBindPipeline(list->vk_command_buffer, + VK_PIPELINE_BIND_POINT_COMPUTE, pipeline.vk_pipeline)); + + VK_CALL(vkCmdBindDescriptorSets(list->vk_command_buffer, + VK_PIPELINE_BIND_POINT_COMPUTE, pipeline.vk_pipeline_layout, + 0, 1, &write_set.dstSet, 0, NULL)); + + for (i = 0; i < rect_count || !i; i++) + { + if (rect_count) + { + /* clamp to actual resource region and skip empty rects */ + curr_rect.left = max(rects[i].left, full_rect.left); + curr_rect.top = max(rects[i].top, full_rect.top); + curr_rect.right = min(rects[i].right, full_rect.right); + curr_rect.bottom = min(rects[i].bottom, full_rect.bottom); + + if (curr_rect.left >= curr_rect.right || curr_rect.top >= curr_rect.bottom) + continue; + } + + clear_args.offset.x = curr_rect.left; + clear_args.offset.y = curr_rect.top; + clear_args.extent.width = curr_rect.right - curr_rect.left; + clear_args.extent.height = curr_rect.bottom - curr_rect.top; + + VK_CALL(vkCmdPushConstants(list->vk_command_buffer, + pipeline.vk_pipeline_layout, VK_SHADER_STAGE_COMPUTE_BIT, + 0, sizeof(clear_args), &clear_args)); + + VK_CALL(vkCmdDispatch(list->vk_command_buffer, + vkd3d_compute_workgroup_count(clear_args.extent.width, workgroup_size.width), + vkd3d_compute_workgroup_count(clear_args.extent.height, workgroup_size.height), + vkd3d_compute_workgroup_count(layer_count, workgroup_size.depth))); + } +} + static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewUint(ID3D12GraphicsCommandList1 *iface, D3D12_GPU_DESCRIPTOR_HANDLE gpu_handle, D3D12_CPU_DESCRIPTOR_HANDLE cpu_handle, ID3D12Resource *resource, const UINT values[4], UINT rect_count, const D3D12_RECT *rects) @@ -4906,13 +5028,19 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewFloat(I { struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface); struct d3d12_resource *resource_impl; + struct vkd3d_view *view; + VkClearColorValue color;
- FIXME("iface %p, gpu_handle %#"PRIx64", cpu_handle %lx, resource %p, values %p, rect_count %u, rects %p stub!\n", + TRACE("iface %p, gpu_handle %#"PRIx64", cpu_handle %lx, resource %p, values %p, rect_count %u, rects %p.\n", iface, gpu_handle.ptr, cpu_handle.ptr, resource, values, rect_count, rects);
+ memcpy(color.float32, values, sizeof(color.float32)); + resource_impl = unsafe_impl_from_ID3D12Resource(resource); + view = d3d12_desc_from_cpu_handle(cpu_handle)->u.view;
- d3d12_command_list_track_resource_usage(list, resource_impl); + d3d12_command_list_clear_unordered_access_view(list, resource_impl, + view, &color, rect_count, rects); }
static void STDMETHODCALLTYPE d3d12_command_list_DiscardResource(ID3D12GraphicsCommandList1 *iface, diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h index d0224d3..d8b4f2f 100644 --- a/libs/vkd3d/vkd3d_private.h +++ b/libs/vkd3d/vkd3d_private.h @@ -1299,6 +1299,11 @@ static inline unsigned int d3d12_resource_desc_get_sub_resource_count(const D3D1 return d3d12_resource_desc_get_layer_count(desc) * desc->MipLevels; }
+static inline unsigned int vkd3d_compute_workgroup_count(unsigned int thread_count, unsigned int workgroup_size) +{ + return (thread_count + workgroup_size - 1) / workgroup_size; +} + VkCompareOp vk_compare_op_from_d3d12(D3D12_COMPARISON_FUNC op) DECLSPEC_HIDDEN; VkSampleCountFlagBits vk_samples_from_dxgi_sample_desc(const DXGI_SAMPLE_DESC *desc) DECLSPEC_HIDDEN; VkSampleCountFlagBits vk_samples_from_sample_count(unsigned int sample_count) DECLSPEC_HIDDEN;