[PATCH vkd3d 06/13] vkd3d: Implement ClearUnorderedAccessViewFloat.
Philip Rebohle
philip.rebohle at tu-dortmund.de
Mon Nov 11 10:03:41 CST 2019
Signed-off-by: Philip Rebohle <philip.rebohle at 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;
--
2.24.0
More information about the wine-devel
mailing list