[PATCH vkd3d 2/7] vkd3d: Implement d3d12_command_list_ClearUnorderedAccessViewFloat().
Henri Verbeet
hverbeet at codeweavers.com
Mon Nov 25 08:05:37 CST 2019
From: Philip Rebohle <philip.rebohle at tu-dortmund.de>
Signed-off-by: Philip Rebohle <philip.rebohle at tu-dortmund.de>
Signed-off-by: Henri Verbeet <hverbeet at codeweavers.com>
---
This supersedes patches 173323 and 173350.
Makefile.am | 1 +
libs/vkd3d/command.c | 184 ++++++++++++++++++++-
libs/vkd3d/device.c | 6 +
libs/vkd3d/state.c | 222 +++++++++++++++++++++++---
libs/vkd3d/vkd3d_private.h | 38 +++++
libs/vkd3d/vkd3d_shaders.h | 388 +++++++++++++++++++++++++++++++++++++++++++++
tests/d3d12.c | 48 +++---
7 files changed, 836 insertions(+), 51 deletions(-)
create mode 100644 libs/vkd3d/vkd3d_shaders.h
diff --git a/Makefile.am b/Makefile.am
index 2b0e8f3..00a5f58 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -108,6 +108,7 @@ libvkd3d_la_SOURCES = \
libs/vkd3d/vkd3d.map \
libs/vkd3d/vkd3d_main.c \
libs/vkd3d/vkd3d_private.h \
+ libs/vkd3d/vkd3d_shaders.h \
libs/vkd3d/vulkan_procs.h \
libs/vkd3d_version.c
libvkd3d_la_LDFLAGS = $(AM_LDFLAGS) -version-info 2:0:1
diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index 297054b..ade51ba 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -4803,6 +4803,182 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearRenderTargetView(ID3D12Gra
&clear_value, rect_count, rects);
}
+struct vkd3d_uav_clear_pipeline
+{
+ VkDescriptorSetLayout vk_set_layout;
+ VkPipelineLayout vk_pipeline_layout;
+ VkPipeline vk_pipeline;
+ VkExtent3D group_size;
+};
+
+static void vkd3d_uav_clear_state_get_buffer_pipeline(const struct vkd3d_uav_clear_state *state,
+ enum vkd3d_format_type format_type, struct vkd3d_uav_clear_pipeline *info)
+{
+ const struct vkd3d_uav_clear_pipelines *pipelines;
+
+ pipelines = format_type == VKD3D_FORMAT_TYPE_UINT ? &state->pipelines_uint : &state->pipelines_float;
+ info->vk_set_layout = state->vk_set_layout_buffer;
+ info->vk_pipeline_layout = state->vk_pipeline_layout_buffer;
+ info->vk_pipeline = pipelines->buffer;
+ info->group_size = (VkExtent3D){128, 1, 1};
+}
+
+static void vkd3d_uav_clear_state_get_image_pipeline(const struct vkd3d_uav_clear_state *state,
+ VkImageViewType image_view_type, enum vkd3d_format_type format_type, struct vkd3d_uav_clear_pipeline *info)
+{
+ const struct vkd3d_uav_clear_pipelines *pipelines;
+
+ pipelines = format_type == VKD3D_FORMAT_TYPE_UINT ? &state->pipelines_uint : &state->pipelines_float;
+ info->vk_set_layout = state->vk_set_layout_image;
+ info->vk_pipeline_layout = state->vk_pipeline_layout_image;
+
+ switch (image_view_type)
+ {
+ case VK_IMAGE_VIEW_TYPE_1D:
+ info->vk_pipeline = pipelines->image_1d;
+ info->group_size = (VkExtent3D){64, 1, 1};
+ break;
+
+ case VK_IMAGE_VIEW_TYPE_1D_ARRAY:
+ info->vk_pipeline = pipelines->image_1d_array;
+ info->group_size = (VkExtent3D){64, 1, 1};
+ break;
+
+ case VK_IMAGE_VIEW_TYPE_2D:
+ info->vk_pipeline = pipelines->image_2d;
+ info->group_size = (VkExtent3D){8, 8, 1};
+ break;
+
+ case VK_IMAGE_VIEW_TYPE_2D_ARRAY:
+ info->vk_pipeline = pipelines->image_2d_array;
+ info->group_size = (VkExtent3D){8, 8, 1};
+ break;
+
+ case VK_IMAGE_VIEW_TYPE_3D:
+ info->vk_pipeline = pipelines->image_3d;
+ info->group_size = (VkExtent3D){8, 8, 1};
+ break;
+
+ default:
+ ERR("Unhandled view type %#x.\n", image_view_type);
+ info->vk_pipeline = VK_NULL_HANDLE;
+ info->group_size = (VkExtent3D){0, 0, 0};
+ break;
+ }
+}
+
+static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
+ struct d3d12_resource *resource, struct vkd3d_view *view, const VkClearColorValue *clear_colour,
+ unsigned int 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_uav_clear_pipeline pipeline;
+ struct vkd3d_uav_clear_args clear_args;
+ VkDescriptorImageInfo image_info;
+ D3D12_RECT full_rect, curr_rect;
+ VkWriteDescriptorSet write_set;
+
+ 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.colour = *clear_colour;
+
+ 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_buffer(resource))
+ {
+ 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;
+ vkd3d_uav_clear_state_get_buffer_pipeline(&list->device->uav_clear_state,
+ view->format->type, &pipeline);
+ }
+ else
+ {
+ 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;
+ vkd3d_uav_clear_state_get_image_pipeline(&list->device->uav_clear_state,
+ view->info.texture.vk_view_type, view->format->type, &pipeline);
+ }
+
+ 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);
+
+ if (!rect_count)
+ {
+ rects = &full_rect;
+ rect_count = 1;
+ }
+
+ 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)
+ {
+ /* Clamp to the actual resource region and skip empty rectangles. */
+ 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, pipeline.group_size.width),
+ vkd3d_compute_workgroup_count(clear_args.extent.height, pipeline.group_size.height),
+ vkd3d_compute_workgroup_count(layer_count, pipeline.group_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 +5082,17 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewFloat(I
{
struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
struct d3d12_resource *resource_impl;
+ VkClearColorValue colour;
+ struct vkd3d_view *view;
- 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);
resource_impl = unsafe_impl_from_ID3D12Resource(resource);
+ view = d3d12_desc_from_cpu_handle(cpu_handle)->u.view;
+ memcpy(colour.float32, values, sizeof(colour.float32));
- d3d12_command_list_track_resource_usage(list, resource_impl);
+ d3d12_command_list_clear_uav(list, resource_impl, view, &colour, rect_count, rects);
}
static void STDMETHODCALLTYPE d3d12_command_list_DiscardResource(ID3D12GraphicsCommandList1 *iface,
diff --git a/libs/vkd3d/device.c b/libs/vkd3d/device.c
index 0624318..59859ad 100644
--- a/libs/vkd3d/device.c
+++ b/libs/vkd3d/device.c
@@ -2154,6 +2154,7 @@ static ULONG STDMETHODCALLTYPE d3d12_device_Release(ID3D12Device *iface)
vkd3d_private_store_destroy(&device->private_store);
vkd3d_cleanup_format_info(device);
+ vkd3d_uav_clear_state_cleanup(&device->uav_clear_state, device);
vkd3d_destroy_null_resources(&device->null_resources, device);
vkd3d_gpu_va_allocator_cleanup(&device->gpu_va_allocator);
vkd3d_render_pass_cache_cleanup(&device->render_pass_cache, device);
@@ -3447,6 +3448,9 @@ static HRESULT d3d12_device_init(struct d3d12_device *device,
if (FAILED(hr = vkd3d_init_null_resources(&device->null_resources, device)))
goto out_cleanup_format_info;
+ if (FAILED(hr = vkd3d_uav_clear_state_init(&device->uav_clear_state, device)))
+ goto out_destroy_null_resources;
+
vkd3d_render_pass_cache_init(&device->render_pass_cache);
vkd3d_gpu_va_allocator_init(&device->gpu_va_allocator);
@@ -3458,6 +3462,8 @@ static HRESULT d3d12_device_init(struct d3d12_device *device,
return S_OK;
+out_destroy_null_resources:
+ vkd3d_destroy_null_resources(&device->null_resources, device);
out_cleanup_format_info:
vkd3d_cleanup_format_info(device);
out_stop_fence_worker:
diff --git a/libs/vkd3d/state.c b/libs/vkd3d/state.c
index a321fa4..e1f7da9 100644
--- a/libs/vkd3d/state.c
+++ b/libs/vkd3d/state.c
@@ -18,6 +18,7 @@
*/
#include "vkd3d_private.h"
+#include "vkd3d_shaders.h"
/* ID3D12RootSignature */
static inline struct d3d12_root_signature *impl_from_ID3D12RootSignature(ID3D12RootSignature *iface)
@@ -1384,6 +1385,37 @@ static HRESULT create_shader_stage(struct d3d12_device *device,
return S_OK;
}
+static HRESULT vkd3d_create_compute_pipeline(struct d3d12_device *device,
+ const D3D12_SHADER_BYTECODE *code, const struct vkd3d_shader_interface_info *shader_interface,
+ VkPipelineLayout vk_pipeline_layout, VkPipeline *vk_pipeline)
+{
+ const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+ VkComputePipelineCreateInfo pipeline_info;
+ VkResult vr;
+ HRESULT hr;
+
+ pipeline_info.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;
+ pipeline_info.pNext = NULL;
+ pipeline_info.flags = 0;
+ if (FAILED(hr = create_shader_stage(device, &pipeline_info.stage,
+ VK_SHADER_STAGE_COMPUTE_BIT, code, shader_interface, NULL)))
+ return hr;
+ pipeline_info.layout = vk_pipeline_layout;
+ pipeline_info.basePipelineHandle = VK_NULL_HANDLE;
+ pipeline_info.basePipelineIndex = -1;
+
+ vr = VK_CALL(vkCreateComputePipelines(device->vk_device,
+ VK_NULL_HANDLE, 1, &pipeline_info, NULL, vk_pipeline));
+ VK_CALL(vkDestroyShaderModule(device->vk_device, pipeline_info.stage.module, NULL));
+ if (vr < 0)
+ {
+ WARN("Failed to create Vulkan compute pipeline, hr %#x.", hr);
+ return hresult_from_vk_result(vr);
+ }
+
+ return S_OK;
+}
+
static HRESULT d3d12_pipeline_state_init_compute_uav_counters(struct d3d12_pipeline_state *state,
struct d3d12_device *device, const struct d3d12_root_signature *root_signature,
const struct vkd3d_shader_scan_info *shader_info)
@@ -1470,10 +1502,9 @@ static HRESULT d3d12_pipeline_state_init_compute(struct d3d12_pipeline_state *st
const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
struct vkd3d_shader_interface_info shader_interface;
const struct d3d12_root_signature *root_signature;
- VkComputePipelineCreateInfo pipeline_info;
struct vkd3d_shader_scan_info shader_info;
+ VkPipelineLayout vk_pipeline_layout;
struct vkd3d_shader_code dxbc;
- VkResult vr;
HRESULT hr;
int ret;
@@ -1519,36 +1550,18 @@ static HRESULT d3d12_pipeline_state_init_compute(struct d3d12_pipeline_state *st
shader_interface.uav_counters = state->uav_counters;
shader_interface.uav_counter_count = vkd3d_popcount(state->uav_counter_mask);
- pipeline_info.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;
- pipeline_info.pNext = NULL;
- pipeline_info.flags = 0;
- if (FAILED(hr = create_shader_stage(device, &pipeline_info.stage,
- VK_SHADER_STAGE_COMPUTE_BIT, &desc->CS, &shader_interface, NULL)))
- {
- if (state->vk_set_layout)
- VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, state->vk_set_layout, NULL));
- if (state->vk_pipeline_layout)
- VK_CALL(vkDestroyPipelineLayout(device->vk_device, state->vk_pipeline_layout, NULL));
- vkd3d_free(state->uav_counters);
- return hr;
- }
- pipeline_info.layout = state->vk_pipeline_layout
+ vk_pipeline_layout = state->vk_pipeline_layout
? state->vk_pipeline_layout : root_signature->vk_pipeline_layout;
- pipeline_info.basePipelineHandle = VK_NULL_HANDLE;
- pipeline_info.basePipelineIndex = -1;
-
- vr = VK_CALL(vkCreateComputePipelines(device->vk_device, VK_NULL_HANDLE,
- 1, &pipeline_info, NULL, &state->u.compute.vk_pipeline));
- VK_CALL(vkDestroyShaderModule(device->vk_device, pipeline_info.stage.module, NULL));
- if (vr)
+ if (FAILED(hr = vkd3d_create_compute_pipeline(device, &desc->CS, &shader_interface,
+ vk_pipeline_layout, &state->u.compute.vk_pipeline)))
{
- WARN("Failed to create Vulkan compute pipeline, vr %d.\n", vr);
+ WARN("Failed to create Vulkan compute pipeline, hr %#x.\n", hr);
if (state->vk_set_layout)
VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, state->vk_set_layout, NULL));
if (state->vk_pipeline_layout)
VK_CALL(vkDestroyPipelineLayout(device->vk_device, state->vk_pipeline_layout, NULL));
vkd3d_free(state->uav_counters);
- return hresult_from_vk_result(vr);
+ return hr;
}
if (FAILED(hr = vkd3d_private_store_init(&state->private_store)))
@@ -2802,3 +2815,162 @@ VkPipeline d3d12_pipeline_state_get_or_create_pipeline(struct d3d12_pipeline_sta
ERR("Could not get the pipeline compiled by other thread from the cache.\n");
return vk_pipeline;
}
+
+static void vkd3d_uav_clear_pipelines_cleanup(struct vkd3d_uav_clear_pipelines *pipelines,
+ struct d3d12_device *device)
+{
+ const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_3d, NULL));
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_2d_array, NULL));
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_2d, NULL));
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_1d_array, NULL));
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_1d, NULL));
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->buffer, NULL));
+}
+
+void vkd3d_uav_clear_state_cleanup(struct vkd3d_uav_clear_state *state, struct d3d12_device *device)
+{
+ const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+
+ vkd3d_uav_clear_pipelines_cleanup(&state->pipelines_uint, device);
+ vkd3d_uav_clear_pipelines_cleanup(&state->pipelines_float, device);
+
+ VK_CALL(vkDestroyPipelineLayout(device->vk_device, state->vk_pipeline_layout_image, NULL));
+ VK_CALL(vkDestroyPipelineLayout(device->vk_device, state->vk_pipeline_layout_buffer, NULL));
+
+ VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, state->vk_set_layout_image, NULL));
+ VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, state->vk_set_layout_buffer, NULL));
+}
+
+HRESULT vkd3d_uav_clear_state_init(struct vkd3d_uav_clear_state *state, struct d3d12_device *device)
+{
+ struct vkd3d_shader_push_constant_buffer push_constant;
+ struct vkd3d_shader_interface_info shader_interface;
+ struct vkd3d_shader_resource_binding binding;
+ VkDescriptorSetLayoutBinding set_binding;
+ VkPushConstantRange push_constant_range;
+ unsigned int i;
+ HRESULT hr;
+
+ const struct
+ {
+ VkDescriptorSetLayout *set_layout;
+ VkPipelineLayout *pipeline_layout;
+ VkDescriptorType descriptor_type;
+ }
+ set_layouts[] =
+ {
+ {&state->vk_set_layout_buffer, &state->vk_pipeline_layout_buffer, VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER},
+ {&state->vk_set_layout_image, &state->vk_pipeline_layout_image, VK_DESCRIPTOR_TYPE_STORAGE_IMAGE},
+ };
+
+ const struct
+ {
+ VkPipeline *pipeline;
+ VkPipelineLayout *pipeline_layout;
+ D3D12_SHADER_BYTECODE code;
+ }
+ pipelines[] =
+ {
+#define SHADER_CODE(name) {name, sizeof(name)}
+ {&state->pipelines_float.buffer, &state->vk_pipeline_layout_buffer,
+ SHADER_CODE(cs_uav_clear_buffer_float_code)},
+ {&state->pipelines_float.image_1d, &state->vk_pipeline_layout_image,
+ SHADER_CODE(cs_uav_clear_1d_float_code)},
+ {&state->pipelines_float.image_1d_array, &state->vk_pipeline_layout_image,
+ SHADER_CODE(cs_uav_clear_1d_array_float_code)},
+ {&state->pipelines_float.image_2d, &state->vk_pipeline_layout_image,
+ SHADER_CODE(cs_uav_clear_2d_float_code)},
+ {&state->pipelines_float.image_2d_array, &state->vk_pipeline_layout_image,
+ SHADER_CODE(cs_uav_clear_2d_array_float_code)},
+ {&state->pipelines_float.image_3d, &state->vk_pipeline_layout_image,
+ SHADER_CODE(cs_uav_clear_3d_float_code)},
+
+ {&state->pipelines_uint.buffer, &state->vk_pipeline_layout_buffer,
+ SHADER_CODE(cs_uav_clear_buffer_uint_code)},
+ {&state->pipelines_uint.image_1d, &state->vk_pipeline_layout_image,
+ SHADER_CODE(cs_uav_clear_1d_uint_code)},
+ {&state->pipelines_uint.image_1d_array, &state->vk_pipeline_layout_image,
+ SHADER_CODE(cs_uav_clear_1d_array_uint_code)},
+ {&state->pipelines_uint.image_2d, &state->vk_pipeline_layout_image,
+ SHADER_CODE(cs_uav_clear_2d_uint_code)},
+ {&state->pipelines_uint.image_2d_array, &state->vk_pipeline_layout_image,
+ SHADER_CODE(cs_uav_clear_2d_array_uint_code)},
+ {&state->pipelines_uint.image_3d, &state->vk_pipeline_layout_image,
+ SHADER_CODE(cs_uav_clear_3d_uint_code)},
+#undef SHADER_CODE
+ };
+
+ memset(state, 0, sizeof(*state));
+
+ set_binding.binding = 0;
+ set_binding.descriptorCount = 1;
+ set_binding.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
+ set_binding.pImmutableSamplers = NULL;
+
+ binding.type = VKD3D_SHADER_DESCRIPTOR_TYPE_UAV;
+ binding.register_index = 0;
+ binding.shader_visibility = VKD3D_SHADER_VISIBILITY_COMPUTE;
+ binding.binding.set = 0;
+ binding.binding.binding = 0;
+
+ push_constant_range.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
+ push_constant_range.offset = 0;
+ push_constant_range.size = sizeof(struct vkd3d_uav_clear_args);
+
+ push_constant.register_index = 0;
+ push_constant.shader_visibility = VKD3D_SHADER_VISIBILITY_COMPUTE;
+ push_constant.offset = 0;
+ push_constant.size = sizeof(struct vkd3d_uav_clear_args);
+
+ for (i = 0; i < ARRAY_SIZE(set_layouts); ++i)
+ {
+ set_binding.descriptorType = set_layouts[i].descriptor_type;
+
+ if (FAILED(hr = vkd3d_create_descriptor_set_layout(device, 0, 1, &set_binding, set_layouts[i].set_layout)))
+ {
+ ERR("Failed to create descriptor set layout %u, hr %#x.", i, hr);
+ goto fail;
+ }
+
+ if (FAILED(hr = vkd3d_create_pipeline_layout(device, 1, set_layouts[i].set_layout,
+ 1, &push_constant_range, set_layouts[i].pipeline_layout)))
+ {
+ ERR("Failed to create pipeline layout %u, hr %#x.", i, hr);
+ goto fail;
+ }
+ }
+
+ shader_interface.type = VKD3D_SHADER_STRUCTURE_TYPE_SHADER_INTERFACE_INFO;
+ shader_interface.next = NULL;
+ shader_interface.bindings = &binding;
+ shader_interface.binding_count = 1;
+ shader_interface.push_constant_buffers = &push_constant;
+ shader_interface.push_constant_buffer_count = 1;
+ shader_interface.combined_samplers = NULL;
+ shader_interface.combined_sampler_count = 0;
+ shader_interface.uav_counters = NULL;
+ shader_interface.uav_counter_count = 0;
+
+ for (i = 0; i < ARRAY_SIZE(pipelines); ++i)
+ {
+ if (pipelines[i].pipeline_layout == &state->vk_pipeline_layout_buffer)
+ binding.flags = VKD3D_SHADER_BINDING_FLAG_BUFFER;
+ else
+ binding.flags = VKD3D_SHADER_BINDING_FLAG_IMAGE;
+
+ if (FAILED(hr = vkd3d_create_compute_pipeline(device, &pipelines[i].code, &shader_interface,
+ *pipelines[i].pipeline_layout, pipelines[i].pipeline)))
+ {
+ ERR("Failed to create compute pipeline %u, hr %#x.", i, hr);
+ goto fail;
+ }
+ }
+
+ return S_OK;
+
+fail:
+ vkd3d_uav_clear_state_cleanup(state, device);
+ return hr;
+}
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index c6bfe96..933c7c4 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -1059,6 +1059,38 @@ struct vkd3d_format_compatibility_list
VkFormat vk_formats[VKD3D_MAX_COMPATIBLE_FORMAT_COUNT];
};
+struct vkd3d_uav_clear_args
+{
+ VkClearColorValue colour;
+ VkOffset2D offset;
+ VkExtent2D extent;
+};
+
+struct vkd3d_uav_clear_pipelines
+{
+ VkPipeline buffer;
+ VkPipeline image_1d;
+ VkPipeline image_1d_array;
+ VkPipeline image_2d;
+ VkPipeline image_2d_array;
+ VkPipeline image_3d;
+};
+
+struct vkd3d_uav_clear_state
+{
+ VkDescriptorSetLayout vk_set_layout_buffer;
+ VkDescriptorSetLayout vk_set_layout_image;
+
+ VkPipelineLayout vk_pipeline_layout_buffer;
+ VkPipelineLayout vk_pipeline_layout_image;
+
+ struct vkd3d_uav_clear_pipelines pipelines_float;
+ struct vkd3d_uav_clear_pipelines pipelines_uint;
+};
+
+HRESULT vkd3d_uav_clear_state_init(struct vkd3d_uav_clear_state *state, struct d3d12_device *device) DECLSPEC_HIDDEN;
+void vkd3d_uav_clear_state_cleanup(struct vkd3d_uav_clear_state *state, struct d3d12_device *device) DECLSPEC_HIDDEN;
+
/* ID3D12Device */
struct d3d12_device
{
@@ -1104,6 +1136,7 @@ struct d3d12_device
unsigned int format_compatibility_list_count;
const struct vkd3d_format_compatibility_list *format_compatibility_lists;
struct vkd3d_null_resources null_resources;
+ struct vkd3d_uav_clear_state uav_clear_state;
};
HRESULT d3d12_device_create(struct vkd3d_instance *instance,
@@ -1237,6 +1270,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;
diff --git a/libs/vkd3d/vkd3d_shaders.h b/libs/vkd3d/vkd3d_shaders.h
new file mode 100644
index 0000000..b2a90cd
--- /dev/null
+++ b/libs/vkd3d/vkd3d_shaders.h
@@ -0,0 +1,388 @@
+/*
+ * Copyright 2019 Philip Rebohle
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2.1 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library; if not, write to the Free Software
+ * Foundation, Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301, USA
+ */
+
+#ifndef __VKD3D_SHADERS_H
+#define __VKD3D_SHADERS_H
+
+static const uint32_t cs_uav_clear_buffer_float_code[] =
+{
+#if 0
+ RWBuffer<float4> dst;
+
+ struct
+ {
+ float4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(128, 1, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (thread_id.x < u_info.dst_extent.x)
+ dst[u_info.dst_offset.x + thread_id.x] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0xe114ba61, 0xff6a0d0b, 0x7b25c8f4, 0xfcf7cf22, 0x00000001, 0x0000010c, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000b8, 0x00050050, 0x0000002e, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400089c, 0x0011e000, 0x00000000, 0x00005555,
+ 0x0200005f, 0x00020012, 0x02000068, 0x00000001, 0x0400009b, 0x00000080, 0x00000001, 0x00000001,
+ 0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+ 0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+ 0x00000001, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100006, 0x00000000, 0x00208e46, 0x00000000,
+ 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_buffer_uint_code[] =
+{
+#if 0
+ RWBuffer<uint4> dst;
+
+ struct
+ {
+ uint4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(128, 1, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (thread_id.x < u_info.dst_extent.x)
+ dst[u_info.dst_offset.x + thread_id.x] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0x3afd0cfd, 0x5145c166, 0x5b9f76b8, 0xa73775cd, 0x00000001, 0x0000010c, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000b8, 0x00050050, 0x0000002e, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400089c, 0x0011e000, 0x00000000, 0x00004444,
+ 0x0200005f, 0x00020012, 0x02000068, 0x00000001, 0x0400009b, 0x00000080, 0x00000001, 0x00000001,
+ 0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+ 0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+ 0x00000001, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100006, 0x00000000, 0x00208e46, 0x00000000,
+ 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_1d_array_float_code[] =
+{
+#if 0
+ RWTexture1DArray<float4> dst;
+
+ struct
+ {
+ float4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(64, 1, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (thread_id.x < u_info.dst_extent.x)
+ dst[int2(u_info.dst_offset.x + thread_id.x, thread_id.y)] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0x3d73bc2d, 0x2b635f3d, 0x6bf98e92, 0xbe0aa5d9, 0x00000001, 0x0000011c, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000c8, 0x00050050, 0x00000032, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400389c, 0x0011e000, 0x00000000, 0x00005555,
+ 0x0200005f, 0x00020032, 0x02000068, 0x00000001, 0x0400009b, 0x00000040, 0x00000001, 0x00000001,
+ 0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+ 0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+ 0x00000001, 0x04000036, 0x001000e2, 0x00000000, 0x00020556, 0x080000a4, 0x0011e0f2, 0x00000000,
+ 0x00100e46, 0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_1d_array_uint_code[] =
+{
+#if 0
+ RWTexture1DArray<uint4> dst;
+
+ struct
+ {
+ uint4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(64, 1, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (thread_id.x < u_info.dst_extent.x)
+ dst[int2(u_info.dst_offset.x + thread_id.x, thread_id.y)] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0x2f0ca457, 0x72068b34, 0xd9dadc2b, 0xd3178c3e, 0x00000001, 0x0000011c, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000c8, 0x00050050, 0x00000032, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400389c, 0x0011e000, 0x00000000, 0x00004444,
+ 0x0200005f, 0x00020032, 0x02000068, 0x00000001, 0x0400009b, 0x00000040, 0x00000001, 0x00000001,
+ 0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+ 0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+ 0x00000001, 0x04000036, 0x001000e2, 0x00000000, 0x00020556, 0x080000a4, 0x0011e0f2, 0x00000000,
+ 0x00100e46, 0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_1d_float_code[] =
+{
+#if 0
+ RWTexture1D<float4> dst;
+
+ struct
+ {
+ float4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(64, 1, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (thread_id.x < u_info.dst_extent.x)
+ dst[u_info.dst_offset.x + thread_id.x] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0x05266503, 0x4b97006f, 0x01a5cc63, 0xe617d0a1, 0x00000001, 0x0000010c, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000b8, 0x00050050, 0x0000002e, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400109c, 0x0011e000, 0x00000000, 0x00005555,
+ 0x0200005f, 0x00020012, 0x02000068, 0x00000001, 0x0400009b, 0x00000040, 0x00000001, 0x00000001,
+ 0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+ 0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+ 0x00000001, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100006, 0x00000000, 0x00208e46, 0x00000000,
+ 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_1d_uint_code[] =
+{
+#if 0
+ RWTexture1D<uint4> dst;
+
+ struct
+ {
+ uint4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(64, 1, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (thread_id.x < u_info.dst_extent.x)
+ dst[u_info.dst_offset.x + thread_id.x] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0x19d5c8f2, 0x3ca4ac24, 0x9e258499, 0xf0463fd6, 0x00000001, 0x0000010c, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000b8, 0x00050050, 0x0000002e, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400109c, 0x0011e000, 0x00000000, 0x00004444,
+ 0x0200005f, 0x00020012, 0x02000068, 0x00000001, 0x0400009b, 0x00000040, 0x00000001, 0x00000001,
+ 0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+ 0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+ 0x00000001, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100006, 0x00000000, 0x00208e46, 0x00000000,
+ 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_2d_array_float_code[] =
+{
+#if 0
+ RWTexture2DArray<float4> dst;
+
+ struct
+ {
+ float4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(8, 8, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (all(thread_id.xy < u_info.dst_extent.xy))
+ dst[int3(u_info.dst_offset.xy + thread_id.xy, thread_id.z)] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0x924d2d2c, 0xb9166376, 0x99f83871, 0x8ef65025, 0x00000001, 0x00000138, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000e4, 0x00050050, 0x00000039, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400409c, 0x0011e000, 0x00000000, 0x00005555,
+ 0x0200005f, 0x00020072, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+ 0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+ 0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+ 0x00000000, 0x0700001e, 0x00100032, 0x00000000, 0x00020046, 0x00208046, 0x00000000, 0x00000001,
+ 0x04000036, 0x001000c2, 0x00000000, 0x00020aa6, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46,
+ 0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_2d_array_uint_code[] =
+{
+#if 0
+ RWTexture2DArray<uint4> dst;
+
+ struct
+ {
+ uint4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(8, 8, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (all(thread_id.xy < u_info.dst_extent.xy))
+ dst[int3(u_info.dst_offset.xy + thread_id.xy, thread_id.z)] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0xa92219d4, 0xa2c5e47d, 0x0d308500, 0xf32197b4, 0x00000001, 0x00000138, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000e4, 0x00050050, 0x00000039, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400409c, 0x0011e000, 0x00000000, 0x00004444,
+ 0x0200005f, 0x00020072, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+ 0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+ 0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+ 0x00000000, 0x0700001e, 0x00100032, 0x00000000, 0x00020046, 0x00208046, 0x00000000, 0x00000001,
+ 0x04000036, 0x001000c2, 0x00000000, 0x00020aa6, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46,
+ 0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_2d_float_code[] =
+{
+#if 0
+ RWTexture2D<float4> dst;
+
+ struct
+ {
+ float4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(8, 8, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (all(thread_id.xy < u_info.dst_extent.xy))
+ dst[u_info.dst_offset.xy + thread_id.xy] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0x6e735b3f, 0x7348c4fa, 0xb3634e42, 0x50e2d99b, 0x00000001, 0x00000128, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000d4, 0x00050050, 0x00000035, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400189c, 0x0011e000, 0x00000000, 0x00005555,
+ 0x0200005f, 0x00020032, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+ 0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+ 0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+ 0x00000000, 0x0700001e, 0x001000f2, 0x00000000, 0x00020546, 0x00208546, 0x00000000, 0x00000001,
+ 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46, 0x00000000, 0x00208e46, 0x00000000, 0x00000000,
+ 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_2d_uint_code[] =
+{
+#if 0
+ RWTexture2D<uint4> dst;
+
+ struct
+ {
+ uint4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(8, 8, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (all(thread_id.xy < u_info.dst_extent.xy))
+ dst[u_info.dst_offset.xy + thread_id.xy] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0xf01db5dd, 0xc7dc5e55, 0xb017c1a8, 0x55abd52d, 0x00000001, 0x00000128, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000d4, 0x00050050, 0x00000035, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400189c, 0x0011e000, 0x00000000, 0x00004444,
+ 0x0200005f, 0x00020032, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+ 0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+ 0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+ 0x00000000, 0x0700001e, 0x001000f2, 0x00000000, 0x00020546, 0x00208546, 0x00000000, 0x00000001,
+ 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46, 0x00000000, 0x00208e46, 0x00000000, 0x00000000,
+ 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_3d_float_code[] =
+{
+#if 0
+ RWTexture3D<float4> dst;
+
+ struct
+ {
+ float4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(8, 8, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (all(thread_id.xy < u_info.dst_extent.xy))
+ dst[int3(u_info.dst_offset.xy, 0) + thread_id.xyz] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0x5d8f36a0, 0x30fa86a5, 0xfec7f2ef, 0xdfd76cbb, 0x00000001, 0x00000138, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000e4, 0x00050050, 0x00000039, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400289c, 0x0011e000, 0x00000000, 0x00005555,
+ 0x0200005f, 0x00020072, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+ 0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+ 0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+ 0x00000000, 0x0700001e, 0x00100032, 0x00000000, 0x00020046, 0x00208046, 0x00000000, 0x00000001,
+ 0x04000036, 0x001000c2, 0x00000000, 0x00020aa6, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46,
+ 0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_3d_uint_code[] =
+{
+#if 0
+ RWTexture3D<uint4> dst;
+
+ struct
+ {
+ uint4 clear_value;
+ int2 dst_offset;
+ int2 dst_extent;
+ } u_info;
+
+ [numthreads(8, 8, 1)]
+ void main(int3 thread_id : SV_DispatchThreadID)
+ {
+ if (all(thread_id.xy < u_info.dst_extent.xy))
+ dst[int3(u_info.dst_offset.xy, 0) + thread_id.xyz] = u_info.clear_value;
+ }
+#endif
+ 0x43425844, 0x5b9c95b1, 0xc9bde4e3, 0x9aaff806, 0x24a1d264, 0x00000001, 0x00000138, 0x00000003,
+ 0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+ 0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000e4, 0x00050050, 0x00000039, 0x0100086a,
+ 0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400289c, 0x0011e000, 0x00000000, 0x00004444,
+ 0x0200005f, 0x00020072, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+ 0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+ 0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+ 0x00000000, 0x0700001e, 0x00100032, 0x00000000, 0x00020046, 0x00208046, 0x00000000, 0x00000001,
+ 0x04000036, 0x001000c2, 0x00000000, 0x00020aa6, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46,
+ 0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+#endif /* __VKD3D_SHADERS_H */
diff --git a/tests/d3d12.c b/tests/d3d12.c
index 5284138..3b5264f 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -4812,7 +4812,7 @@ static void test_clear_unordered_access_view_buffer(void)
{DXGI_FORMAT_R32_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
{1, 0, 0, 0}, 1, false, true},
{DXGI_FORMAT_R32_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
- {0x3f800000 /* 1.0f */, 0, 0, 0}, 0x3f800000 /* 1.0f */, true, true},
+ {0x3f800000 /* 1.0f */, 0, 0, 0}, 0x3f800000 /* 1.0f */, true},
{DXGI_FORMAT_R16G16_UINT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
{0x1234, 0xabcd, 0, 0}, 0xabcd1234, false, true},
@@ -4822,16 +4822,16 @@ static void test_clear_unordered_access_view_buffer(void)
{DXGI_FORMAT_R16G16_UNORM, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
{0x1234, 0xabcd, 0, 0}, 0xabcd1234, false, true},
{DXGI_FORMAT_R16G16_UNORM, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
- {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0xffff8000, true, true},
+ {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0xffff8000, true},
{DXGI_FORMAT_R16G16_UNORM, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
- {0x40000000 /* 2.0f */, 0 /* 0.0f */, 0, 0}, 0x0000ffff, true, true},
+ {0x40000000 /* 2.0f */, 0 /* 0.0f */, 0, 0}, 0x0000ffff, true},
{DXGI_FORMAT_R16G16_UNORM, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
- {0xbf800000 /* -1.0f */, 0 /* 0.0f */, 0x3f000000 /* 1.0f */, 0x3f000000 /* 1.0f */}, 0, true, true},
+ {0xbf800000 /* -1.0f */, 0 /* 0.0f */, 0x3f000000 /* 1.0f */, 0x3f000000 /* 1.0f */}, 0, true},
{DXGI_FORMAT_R16G16_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
{0x1234, 0xabcd, 0, 0}, 0xabcd1234, false, true},
{DXGI_FORMAT_R16G16_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
- {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0x3c003800, true, true},
+ {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0x3c003800, true},
{DXGI_FORMAT_R8G8B8A8_UINT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
{0x11, 0x22, 0x33, 0x44}, 0x44332211, false, true},
@@ -4845,10 +4845,10 @@ static void test_clear_unordered_access_view_buffer(void)
{DXGI_FORMAT_R11G11B10_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
{0x7ff, 0, 0x3ff, 0}, 0xffc007ff, false, true},
{DXGI_FORMAT_R11G11B10_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
- {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0x40000000 /* 2.0f */, 0}, 0x801e0380, true, true},
+ {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0x40000000 /* 2.0f */, 0}, 0x801e0380, true},
{DXGI_FORMAT_R11G11B10_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
{0x3f000000 /* 1.0f */, 0 /* 0.0f */, 0xbf800000 /* -1.0f */, 0x3f000000 /* 1.0f */},
- 0x00000380, true, true},
+ 0x00000380, true},
};
memset(&desc, 0, sizeof(desc));
@@ -4982,48 +4982,48 @@ static void test_clear_unordered_access_view_image(void)
tests[] =
{
/* Test clearing a specific mip level. */
- {DXGI_FORMAT_R32_FLOAT, 2, 1, 0, 0, 1, 0, {}, {1, 0, 0, 0}, 1, false, false},
- {DXGI_FORMAT_R32_FLOAT, 2, 1, 1, 0, 1, 0, {}, {1, 0, 0, 0}, 1, false, false},
- {DXGI_FORMAT_R32_FLOAT, 2, 1, 0, 0, 1, 0, {}, {0x3f000000, 0, 0, 0}, 0x3f000000, true, true},
- {DXGI_FORMAT_R32_FLOAT, 2, 1, 1, 0, 1, 0, {}, {0x3f000000, 0, 0, 0}, 0x3f000000, true, true},
+ {DXGI_FORMAT_R32_FLOAT, 2, 1, 0, 0, 1, 0, {}, {1, 0, 0, 0}, 1},
+ {DXGI_FORMAT_R32_FLOAT, 2, 1, 1, 0, 1, 0, {}, {1, 0, 0, 0}, 1},
+ {DXGI_FORMAT_R32_FLOAT, 2, 1, 0, 0, 1, 0, {}, {0x3f000000, 0, 0, 0}, 0x3f000000, true},
+ {DXGI_FORMAT_R32_FLOAT, 2, 1, 1, 0, 1, 0, {}, {0x3f000000, 0, 0, 0}, 0x3f000000, true},
/* Test clearing specific array layers. */
- {DXGI_FORMAT_R32_FLOAT, 1, IMAGE_SIZE, 0, 0, IMAGE_SIZE, 0, {}, {1, 0, 0, 0}, 1, false, false},
- {DXGI_FORMAT_R32_FLOAT, 1, IMAGE_SIZE, 0, 3, 2, 0, {}, {1, 0, 0, 0}, 1, false, false},
+ {DXGI_FORMAT_R32_FLOAT, 1, IMAGE_SIZE, 0, 0, IMAGE_SIZE, 0, {}, {1, 0, 0, 0}, 1},
+ {DXGI_FORMAT_R32_FLOAT, 1, IMAGE_SIZE, 0, 3, 2, 0, {}, {1, 0, 0, 0}, 1},
{DXGI_FORMAT_R32_FLOAT, 1, IMAGE_SIZE, 0, 0, IMAGE_SIZE, 0, {},
- {0x3f000000, 0, 0, 0}, 0x3f000000, true, true},
+ {0x3f000000, 0, 0, 0}, 0x3f000000, true},
{DXGI_FORMAT_R32_FLOAT, 1, IMAGE_SIZE, 0, 3, 2, 0, {},
- {0x3f000000, 0, 0, 0}, 0x3f000000, true, true},
+ {0x3f000000, 0, 0, 0}, 0x3f000000, true},
/* Test a single clear rect. */
{DXGI_FORMAT_R32_FLOAT, 1, 1, 0, 0, 1, 1, {{1, 2, IMAGE_SIZE - 4, IMAGE_SIZE - 2}},
{1, 0, 0, 0}, 1, false, true},
{DXGI_FORMAT_R32_FLOAT, 1, 1, 0, 0, 1, 1, {{1, 2, IMAGE_SIZE - 4, IMAGE_SIZE - 2}},
- {0x3f000000, 0, 0, 0}, 0x3f000000, true, true},
+ {0x3f000000, 0, 0, 0}, 0x3f000000, true},
/* Test multiple clear rects. */
{DXGI_FORMAT_R32_FLOAT, 1, 1, 0, 0, 1, 2, {{1, 2, 3, 4}, {5, 6, 7, 8}},
{1, 0, 0, 0}, 1, false, true},
{DXGI_FORMAT_R32_FLOAT, 1, 1, 0, 0, 1, 2, {{1, 2, 3, 4}, {5, 6, 7, 8}},
- {0x3f000000, 0, 0, 0}, 0x3f000000, true, true},
+ {0x3f000000, 0, 0, 0}, 0x3f000000, true},
/* Test uint clears with formats. */
- {DXGI_FORMAT_R16G16_UINT, 1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4}, 0x00020001, false, false},
+ {DXGI_FORMAT_R16G16_UINT, 1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4}, 0x00020001},
{DXGI_FORMAT_R16G16_UINT, 1, 1, 0, 0, 1, 0, {}, {0x12345, 0, 0, 0}, 0x00002345, false, true},
{DXGI_FORMAT_R16G16_UNORM, 1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4}, 0x00020001, false, true},
{DXGI_FORMAT_R16G16_FLOAT, 1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4}, 0x00020001, false, true},
- {DXGI_FORMAT_R8G8B8A8_UINT, 1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4}, 0x04030201, false, false},
+ {DXGI_FORMAT_R8G8B8A8_UINT, 1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4}, 0x04030201},
{DXGI_FORMAT_R8G8B8A8_UINT, 1, 1, 0, 0, 1, 0, {}, {0x123, 0, 0, 0}, 0x00000023, false, true},
{DXGI_FORMAT_R8G8B8A8_UNORM, 1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4}, 0x04030201, false, true},
{DXGI_FORMAT_R11G11B10_FLOAT, 1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4}, 0x00c01001, false, true},
/* Test float clears with formats. */
{DXGI_FORMAT_R16G16_UNORM, 1, 1, 0, 0, 1, 0, {},
- {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0xffff8000, true, true},
+ {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0xffff8000, true},
{DXGI_FORMAT_R16G16_FLOAT, 1, 1, 0, 0, 1, 0, {},
- {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0x3c003800, true, true},
+ {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0x3c003800, true},
{DXGI_FORMAT_R8G8B8A8_UNORM, 1, 1, 0, 0, 1, 0, {},
- {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0x0000ff80, true, true},
+ {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0x0000ff80, true},
{DXGI_FORMAT_R8G8B8A8_UNORM, 1, 1, 0, 0, 1, 0, {},
- {0, 0, 0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */}, 0xff800000, true, true},
+ {0, 0, 0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */}, 0xff800000, true},
{DXGI_FORMAT_R11G11B10_FLOAT, 1, 1, 0, 0, 1, 0, {},
{0x3f000000 /* 1.0f */, 0 /* 0.0f */, 0xbf800000 /* -1.0f */, 0x3f000000 /* 1.0f */},
- 0x00000380, true, true},
+ 0x00000380, true},
};
static const struct
--
2.11.0
More information about the wine-devel
mailing list