[PATCH vkd3d 00/13] vkd3d: Implement compute-shader based UAV clears.

[Philip Rebohle]

This series addresses the following issues with the current ClearUnorderedAccessView* implementation:
- Supports arbitrary buffer formats, not just R32_{UINT,TYPELESS}.
- Handles images with non-UINT formats in ClearUnorderedAccessViewUint.
- Avoids transfer stage barriers by always using a compute shader.
- Implements ClearUnorderedAccessViewFloat, which was previously not supported.
- Supports clear rects.

It should be possible to reuse a lot of the preliminary work for future meta operations
(such as depth<->color clears mentioned on https://wiki.winehq.org/Vkd3d_known_issues).

As for shaders, I did try to hook up glslangValidator to the build system so that they
could be compiled at build time, but without success, so I put the binaries into their
own header file.

Unfortunately, the number of permutations is fairly high (12 in this case), and it seems
impossible to generate a single SPIR-V binary with multiple entry points (GLSL requires
the entry point to be 'main' so there can only be one, and none of the easily accessible
HLSL compilers seem to support it either).

An alternative would be to generate the SPIR-V shaders on the fly, but the SPIR-V
builder from vkd3d-shader is private so this would require a lot of additional work.

Note that this may conflict with the following pending patches (haven't checked in detail):
- 171883 "vkd3d: Allocate one large buffer for a heap and offset into it."
- 172639 "vkd3d: Store a copy of Vulkan view object handles in descriptors."

[PATCH vkd3d 01/13] vkd3d: Store view properties in vkd3d_view.

[Philip Rebohle]

The additional data is needed to implement UAV clears.

Moving this out of d3d12_desc also helps make copying and
traversing descriptor arrays more CPU cache-friendly.

Signed-off-by: Philip Rebohle 
---
 libs/vkd3d/command.c       | 16 ++++++++--------
 libs/vkd3d/resource.c      | 23 ++++++++---------------
 libs/vkd3d/vkd3d_private.h | 32 ++++++++++++++++----------------
 3 files changed, 32 insertions(+), 39 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index 0532ec0..f88f05d 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -4818,14 +4818,14 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewUint(ID

if (d3d12_resource_is_buffer(resource_impl))
     {
-        if (!cpu_descriptor->uav.buffer.size)
+        if (cpu_descriptor->u.view->format->vk_format != VK_FORMAT_R32_UINT)
         {
             FIXME("Not supported for UAV descriptor %p.\n", cpu_descriptor);
             return;
         }

VK_CALL(vkCmdFillBuffer(list->vk_command_buffer, resource_impl->u.vk_buffer,
-                cpu_descriptor->uav.buffer.offset, cpu_descriptor->uav.buffer.size, values[0]));
+                cpu_descriptor->u.view->info.buffer.offset, cpu_descriptor->u.view->info.buffer.size, values[0]));

buffer_barrier.sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER;
         buffer_barrier.pNext = NULL;
@@ -4833,8 +4833,8 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewUint(ID
         buffer_barrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
         buffer_barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
         buffer_barrier.buffer = resource_impl->u.vk_buffer;
-        buffer_barrier.offset = cpu_descriptor->uav.buffer.offset;
-        buffer_barrier.size = cpu_descriptor->uav.buffer.size;
+        buffer_barrier.offset = cpu_descriptor->u.view->info.buffer.offset;
+        buffer_barrier.size = cpu_descriptor->u.view->info.buffer.size;

vk_barrier_parameters_from_d3d12_resource_state(D3D12_RESOURCE_STATE_UNORDERED_ACCESS, 0,
                 resource_impl, list->vk_queue_flags, vk_info, &buffer_barrier.dstAccessMask, &stage_mask, NULL);
@@ -4850,11 +4850,11 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewUint(ID
         color.uint32[2] = values[2];
         color.uint32[3] = values[3];

-        range.aspectMask = cpu_descriptor->uav.texture.vk_aspect_mask;
-        range.baseMipLevel = cpu_descriptor->uav.texture.miplevel_idx;
+        range.aspectMask = cpu_descriptor->u.view->format->vk_aspect_mask;
+        range.baseMipLevel = cpu_descriptor->u.view->info.texture.miplevel_idx;
         range.levelCount = 1;
-        range.baseArrayLayer = cpu_descriptor->uav.texture.layer_idx;
-        range.layerCount = cpu_descriptor->uav.texture.layer_count;
+        range.baseArrayLayer = cpu_descriptor->u.view->info.texture.layer_idx;
+        range.layerCount = cpu_descriptor->u.view->info.texture.layer_count;

VK_CALL(vkCmdClearColorImage(list->vk_command_buffer,
                 resource_impl->u.vk_image, VK_IMAGE_LAYOUT_GENERAL, &color, 1, &range));
diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index e93d50b..b4ad846 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -2166,6 +2166,9 @@ static bool vkd3d_create_buffer_view(struct d3d12_device *device,
     }

object->u.vk_buffer_view = vk_view;
+    object->format = format;
+    object->info.buffer.offset = offset;
+    object->info.buffer.size = size;
     *view = object;
     return true;
 }
@@ -2442,6 +2445,11 @@ static bool vkd3d_create_texture_view(struct d3d12_device *device,
     }

object->u.vk_image_view = vk_view;
+    object->format = format;
+    object->info.texture.vk_view_type = desc->view_type;
+    object->info.texture.miplevel_idx = desc->miplevel_idx;
+    object->info.texture.layer_idx = desc->layer_idx;
+    object->info.texture.layer_count = desc->layer_count;
     *view = object;
     return true;
 }
@@ -2812,16 +2820,6 @@ static void vkd3d_create_buffer_uav(struct d3d12_desc *descriptor, struct d3d12_
             d3d12_desc_destroy(descriptor, device);
         }
     }
-
-    /* FIXME: Clears are implemented only for R32_UINT buffer UAVs. */
-    if ((desc->Format == DXGI_FORMAT_R32_TYPELESS && (desc->u.Buffer.Flags & VKD3D_VIEW_RAW_BUFFER))
-            || desc->Format == DXGI_FORMAT_R32_UINT)
-    {
-        const struct vkd3d_format *format = vkd3d_get_format(device, DXGI_FORMAT_R32_UINT, false);
-
-        descriptor->uav.buffer.offset = desc->u.Buffer.FirstElement * format->byte_count;
-        descriptor->uav.buffer.size = desc->u.Buffer.NumElements * format->byte_count;
-    }
 }

static void vkd3d_create_texture_uav(struct d3d12_desc *descriptor,
@@ -2875,11 +2873,6 @@ static void vkd3d_create_texture_uav(struct d3d12_desc *descriptor,
     descriptor->magic = VKD3D_DESCRIPTOR_MAGIC_UAV;
     descriptor->vk_descriptor_type = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE;
     descriptor->u.view = view;
-
-    descriptor->uav.texture.vk_aspect_mask = vkd3d_desc.format->vk_aspect_mask;
-    descriptor->uav.texture.miplevel_idx = vkd3d_desc.miplevel_idx;
-    descriptor->uav.texture.layer_idx = vkd3d_desc.layer_idx;
-    descriptor->uav.texture.layer_count = vkd3d_desc.layer_count;
 }

void d3d12_desc_create_uav(struct d3d12_desc *descriptor, struct d3d12_device *device,
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index daa521d..2d3f6f8 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -468,6 +468,22 @@ struct vkd3d_view
         VkSampler vk_sampler;
     } u;
     VkBufferView vk_counter_view;
+    const struct vkd3d_format *format;
+    union
+    {
+        struct
+        {
+            VkDeviceSize offset;
+            VkDeviceSize size;
+        } buffer;
+        struct
+        {
+            VkImageViewType vk_view_type;
+            unsigned int miplevel_idx;
+            unsigned int layer_idx;
+            unsigned int layer_count;
+        } texture;
+    } info;
 };

void vkd3d_view_decref(struct vkd3d_view *view, struct d3d12_device *device) DECLSPEC_HIDDEN;
@@ -482,22 +498,6 @@ struct d3d12_desc
         VkDescriptorBufferInfo vk_cbv_info;
         struct vkd3d_view *view;
     } u;
-
-    union
-    {
-        struct
-        {
-            VkDeviceSize offset;
-            VkDeviceSize size;
-        } buffer;
-        struct
-        {
-            VkImageAspectFlags vk_aspect_mask;
-            unsigned int miplevel_idx;
-            unsigned int layer_idx;
-            unsigned int layer_count;
-        } texture;
-    } uav;
 };

static inline struct d3d12_desc *d3d12_desc_from_cpu_handle(D3D12_CPU_DESCRIPTOR_HANDLE cpu_handle)

[PATCH vkd3d 01/13] vkd3d: Store view properties in vkd3d_view.

[Conor McCarthy]

On Tue, Nov 12, 2019 at 2:40 AM Philip Rebohle <
[e-mail]> wrote:
> Moving this out of d3d12_desc also helps make copying and
> traversing descriptor arrays more CPU cache-friendly.

Making descriptors smaller would help, but the most likely cause of a cache
miss is accessing two objects located far apart in memory. Reading data
from a descriptor and then other data via a view pointer creates this
situation. I made a new version of the cache coherence patch which
eliminates use of vkd3d_view for CBV, SRV and UAV descriptors and only
stores a refcount on the heap, but haven't sent it because it conflicts
with a pending patch. It only gains 0.5% fps, so if new implementations
need an expanded view struct instead then that takes precedence. CPU
caching is worth taking into account though.

Conor

[PATCH vkd3d 01/13] vkd3d: Store view properties in vkd3d_view.

[Philip Rebohle]

Storing the Vulkan view handle in the descriptor should still work with 
this if it's useful for descriptor updates (in fact it can even be done 
without increasing the descriptor size since the VkDescriptorBufferInfo 
struct in the union pushes the whole thing to 32 bytes anyway), but the 
previous implementation stored a lot of data in the descriptor that 
wasn't really needed at all most of the time. ClearUAV and friends are 
rare and heavy enough that, dereferencing the view pointer really isn't 
a performance concern at all.

Am 12.11.19 um 05:19 schrieb Conor McCarthy:
> On Tue, Nov 12, 2019 at 2:40 AM Philip Rebohle 
> <[e-mail] > 
> wrote:
> Moving this out of d3d12_desc also helps make copying and
> traversing descriptor arrays more CPU cache-friendly.
>
>
> Making descriptors smaller would help, but the most likely cause of a 
> cache miss is accessing two objects located far apart in memory. Reading 
> data from a descriptor and then other data via a view pointer creates 
> this situation. I made a new version of the cache coherence patch which 
> eliminates use of vkd3d_view for CBV, SRV and UAV descriptors and only 
> stores a refcount on the heap, but haven't sent it because it conflicts 
> with a pending patch. It only gains 0.5% fps, so if new implementations 
> need an expanded view struct instead then that takes precedence. CPU 
> caching is worth taking into account though.
>
> Conor

[PATCH vkd3d 01/13] vkd3d: Store view properties in vkd3d_view.

[Henri Verbeet]

Signed-off-by: Henri Verbeet

[PATCH vkd3d 02/13] vkd3d: Store view object type in vkd3d_view.

[Philip Rebohle]

Currently, vkd3d_view_destroy_descriptor assumes image views
by default, but we need to be able to attach buffer views to
command allocators for UAV clears.

Signed-off-by: Philip Rebohle 
---
 libs/vkd3d/resource.c      | 44 ++++++++++++++++++--------------------
 libs/vkd3d/vkd3d_private.h |  8 +++++++
 2 files changed, 29 insertions(+), 23 deletions(-)

diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index b4ad846..8615464 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -1976,13 +1976,14 @@ ULONG vkd3d_resource_decref(ID3D12Resource *resource)
 }

/* CBVs, SRVs, UAVs */
-static struct vkd3d_view *vkd3d_view_create(void)
+static struct vkd3d_view *vkd3d_view_create(enum vkd3d_view_type type)
 {
     struct vkd3d_view *view;

if ((view = vkd3d_malloc(sizeof(*view))))
     {
         view->refcount = 1;
+        view->type = type;
         view->vk_counter_view = VK_NULL_HANDLE;
     }
     return view;
@@ -1993,40 +1994,37 @@ void vkd3d_view_incref(struct vkd3d_view *view)
     InterlockedIncrement(&view->refcount);
 }

-static void vkd3d_view_destroy_descriptor(struct vkd3d_view *view,
-        const struct d3d12_desc *descriptor, struct d3d12_device *device)
+static void vkd3d_view_destroy(struct vkd3d_view *view, struct d3d12_device *device)
 {
     const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;

TRACE("Destroying view %p.\n", view);

-    if (!descriptor)
+    switch (view->type)
     {
-        VK_CALL(vkDestroyImageView(device->vk_device, view->u.vk_image_view, NULL));
-    }
-    else if (descriptor->magic == VKD3D_DESCRIPTOR_MAGIC_SRV || descriptor->magic == VKD3D_DESCRIPTOR_MAGIC_UAV)
-    {
-        if (descriptor->vk_descriptor_type == VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER
-                || descriptor->vk_descriptor_type == VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER)
+        case VKD3D_VIEW_TYPE_BUFFER:
             VK_CALL(vkDestroyBufferView(device->vk_device, view->u.vk_buffer_view, NULL));
-        else
+            break;
+        case VKD3D_VIEW_TYPE_IMAGE:
             VK_CALL(vkDestroyImageView(device->vk_device, view->u.vk_image_view, NULL));
-
-        if (view->vk_counter_view)
-            VK_CALL(vkDestroyBufferView(device->vk_device, view->vk_counter_view, NULL));
-    }
-    else if (descriptor->magic == VKD3D_DESCRIPTOR_MAGIC_SAMPLER)
-    {
-        VK_CALL(vkDestroySampler(device->vk_device, view->u.vk_sampler, NULL));
+            break;
+        case VKD3D_VIEW_TYPE_SAMPLER:
+            VK_CALL(vkDestroySampler(device->vk_device, view->u.vk_sampler, NULL));
+            break;
+        default:
+            WARN("Unhandled view type %d.\n", view->type);
     }

+    if (view->vk_counter_view)
+        VK_CALL(vkDestroyBufferView(device->vk_device, view->vk_counter_view, NULL));
+
     vkd3d_free(view);
 }

void vkd3d_view_decref(struct vkd3d_view *view, struct d3d12_device *device)
 {
     if (!InterlockedDecrement(&view->refcount))
-        vkd3d_view_destroy_descriptor(view, NULL, device);
+        vkd3d_view_destroy(view, device);
 }

void d3d12_desc_write_atomic(struct d3d12_desc *dst, const struct d3d12_desc *src,
@@ -2053,7 +2051,7 @@ void d3d12_desc_write_atomic(struct d3d12_desc *dst, const struct d3d12_desc *sr

/* Destroy the view after unlocking to reduce wait time. */
     if (destroy_desc.u.view)
-        vkd3d_view_destroy_descriptor(destroy_desc.u.view, &destroy_desc, device);
+        vkd3d_view_destroy(destroy_desc.u.view, device);
 }

static void d3d12_desc_destroy(struct d3d12_desc *descriptor, struct d3d12_device *device)
@@ -2159,7 +2157,7 @@ static bool vkd3d_create_buffer_view(struct d3d12_device *device,
     if (!vkd3d_create_vk_buffer_view(device, vk_buffer, format, offset, size, &vk_view))
         return false;

-    if (!(object = vkd3d_view_create()))
+    if (!(object = vkd3d_view_create(VKD3D_VIEW_TYPE_BUFFER)))
     {
         VK_CALL(vkDestroyBufferView(device->vk_device, vk_view, NULL));
         return false;
@@ -2438,7 +2436,7 @@ static bool vkd3d_create_texture_view(struct d3d12_device *device,
         return false;
     }

-    if (!(object = vkd3d_view_create()))
+    if (!(object = vkd3d_view_create(VKD3D_VIEW_TYPE_IMAGE)))
     {
         VK_CALL(vkDestroyImageView(device->vk_device, vk_view, NULL));
         return false;
@@ -3017,7 +3015,7 @@ void d3d12_desc_create_sampler(struct d3d12_desc *sampler,
         FIXME("Ignoring border color {%.8e, %.8e, %.8e, %.8e}.\n",
                 desc->BorderColor[0], desc->BorderColor[1], desc->BorderColor[2], desc->BorderColor[3]);

-    if (!(view = vkd3d_view_create()))
+    if (!(view = vkd3d_view_create(VKD3D_VIEW_TYPE_SAMPLER)))
         return;

if (d3d12_create_sampler(device, desc->Filter, desc->AddressU,
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index 2d3f6f8..84b5ff2 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -458,9 +458,17 @@ HRESULT vkd3d_create_buffer(struct d3d12_device *device,
 HRESULT vkd3d_get_image_allocation_info(struct d3d12_device *device,
         const D3D12_RESOURCE_DESC *desc, D3D12_RESOURCE_ALLOCATION_INFO *allocation_info) DECLSPEC_HIDDEN;

+enum vkd3d_view_type
+{
+    VKD3D_VIEW_TYPE_BUFFER,
+    VKD3D_VIEW_TYPE_IMAGE,
+    VKD3D_VIEW_TYPE_SAMPLER,
+};
+
 struct vkd3d_view
 {
     LONG refcount;
+    enum vkd3d_view_type type;
     union
     {
         VkBufferView vk_buffer_view;

[PATCH vkd3d 02/13] vkd3d: Store view object type in vkd3d_view.

[Henri Verbeet]

Signed-off-by: Henri Verbeet

[PATCH vkd3d 03/13] vkd3d: Defer compute pipeline update to Dispatch time.

[Philip Rebohle]

Needed to support compute-based clear and copy operations.

Signed-off-by: Philip Rebohle 
---
 libs/vkd3d/command.c | 73 +++++++++++++++++++++++---------------------
 1 file changed, 39 insertions(+), 34 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index f88f05d..5db2007 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -2480,7 +2480,26 @@ static bool d3d12_command_list_update_current_framebuffer(struct d3d12_command_l
     return true;
 }

-static bool d3d12_command_list_update_current_pipeline(struct d3d12_command_list *list)
+static bool d3d12_command_list_update_compute_pipeline(struct d3d12_command_list *list)
+{
+    const struct vkd3d_vk_device_procs *vk_procs = &list->device->vk_procs;
+
+    if (list->current_pipeline != VK_NULL_HANDLE)
+        return true;
+
+    if (!d3d12_pipeline_state_is_compute(list->state))
+    {
+        WARN("Pipeline state %p is not a compute pipeline.\n", list->state);
+        return false;
+    }
+
+    VK_CALL(vkCmdBindPipeline(list->vk_command_buffer, list->state->vk_bind_point, list->state->u.compute.vk_pipeline));
+    list->current_pipeline = list->state->u.compute.vk_pipeline;
+
+    return true;
+}
+
+static bool d3d12_command_list_update_graphics_pipeline(struct d3d12_command_list *list)
 {
     const struct vkd3d_vk_device_procs *vk_procs = &list->device->vk_procs;
     VkRenderPass vk_render_pass;
@@ -2869,6 +2888,18 @@ static void d3d12_command_list_update_descriptors(struct d3d12_command_list *lis
     d3d12_command_list_update_uav_counter_descriptors(list, bind_point);
 }

+static bool d3d12_command_list_update_compute_state(struct d3d12_command_list *list)
+{
+    d3d12_command_list_end_current_render_pass(list);
+
+    if (!d3d12_command_list_update_compute_pipeline(list))
+        return false;
+
+    d3d12_command_list_update_descriptors(list, VK_PIPELINE_BIND_POINT_COMPUTE);
+
+    return true;
+}
+
 static bool d3d12_command_list_begin_render_pass(struct d3d12_command_list *list)
 {
     const struct vkd3d_vk_device_procs *vk_procs = &list->device->vk_procs;
@@ -2876,13 +2907,7 @@ static bool d3d12_command_list_begin_render_pass(struct d3d12_command_list *list
     struct VkRenderPassBeginInfo begin_desc;
     VkRenderPass vk_render_pass;

-    if (!list->state)
-    {
-        WARN("Pipeline state is NULL.\n");
-        return false;
-    }
-
-    if (!d3d12_command_list_update_current_pipeline(list))
+    if (!d3d12_command_list_update_graphics_pipeline(list))
         return false;
     if (!d3d12_command_list_update_current_framebuffer(list))
         return false;
@@ -3007,18 +3032,14 @@ static void STDMETHODCALLTYPE d3d12_command_list_Dispatch(ID3D12GraphicsCommandL

TRACE("iface %p, x %u, y %u, z %u.\n", iface, x, y, z);

-    if (list->state->vk_bind_point != VK_PIPELINE_BIND_POINT_COMPUTE)
+    if (!d3d12_command_list_update_compute_state(list))
     {
-        WARN("Pipeline state %p has bind point %#x.\n", list->state, list->state->vk_bind_point);
+        WARN("Failed to update compute state, ignoring dispatch.\n");
         return;
     }

vk_procs = &list->device->vk_procs;

-    d3d12_command_list_end_current_render_pass(list);
-
-    d3d12_command_list_update_descriptors(list, VK_PIPELINE_BIND_POINT_COMPUTE);
-
     VK_CALL(vkCmdDispatch(list->vk_command_buffer, x, y, z));
 }

@@ -3709,25 +3730,14 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetPipelineState(ID3D12Graphics
 {
     struct d3d12_pipeline_state *state = unsafe_impl_from_ID3D12PipelineState(pipeline_state);
     struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
-    const struct vkd3d_vk_device_procs *vk_procs;

TRACE("iface %p, pipeline_state %p.\n", iface, pipeline_state);

if (list->state == state)
         return;

-    vk_procs = &list->device->vk_procs;
-
     d3d12_command_list_invalidate_bindings(list, state);
-
-    if (d3d12_pipeline_state_is_compute(state))
-    {
-        VK_CALL(vkCmdBindPipeline(list->vk_command_buffer, state->vk_bind_point, state->u.compute.vk_pipeline));
-    }
-    else
-    {
-        d3d12_command_list_invalidate_current_pipeline(list);
-    }
+    d3d12_command_list_invalidate_current_pipeline(list);

list->state = state;
 }
@@ -5232,17 +5242,12 @@ static void STDMETHODCALLTYPE d3d12_command_list_ExecuteIndirect(ID3D12GraphicsC
                     break;
                 }

-                if (list->state->vk_bind_point != VK_PIPELINE_BIND_POINT_COMPUTE)
+                if (!d3d12_command_list_update_compute_state(list))
                 {
-                    WARN("Pipeline state %p has bind point %#x, ignoring dispatch.\n",
-                            list->state, list->state->vk_bind_point);
-                    break;
+                    WARN("Failed to update compute state, ignoring dispatch.\n");
+                    return;
                 }

-                d3d12_command_list_end_current_render_pass(list);
-
-                d3d12_command_list_update_descriptors(list, VK_PIPELINE_BIND_POINT_COMPUTE);
-
                 VK_CALL(vkCmdDispatchIndirect(list->vk_command_buffer,
                         arg_impl->u.vk_buffer, arg_buffer_offset));
                 break;

[PATCH vkd3d 03/13] vkd3d: Defer compute pipeline update to Dispatch time.

[Henri Verbeet]

Signed-off-by: Henri Verbeet

[PATCH vkd3d 04/13] vkd3d: Factor out root parameter invalidation.

[Philip Rebohle]

Signed-off-by: Philip Rebohle 
---
 libs/vkd3d/command.c | 18 +++++++++++++++---
 1 file changed, 15 insertions(+), 3 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index 5db2007..297054b 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -1872,6 +1872,19 @@ static void d3d12_command_list_invalidate_bindings(struct d3d12_command_list *li
     }
 }

+static void d3d12_command_list_invalidate_root_parameters(struct d3d12_command_list *list,
+        VkPipelineBindPoint bind_point)
+{
+    struct vkd3d_pipeline_bindings *bindings = &list->pipeline_bindings[bind_point];
+
+    if (!bindings->root_signature)
+        return;
+
+    bindings->descriptor_set = VK_NULL_HANDLE;
+    bindings->descriptor_table_dirty_mask = bindings->descriptor_table_active_mask & bindings->root_signature->descriptor_table_mask;
+    bindings->push_descriptor_dirty_mask = bindings->push_descriptor_active_mask & bindings->root_signature->push_descriptor_mask;
+}
+
 static bool vk_barrier_parameters_from_d3d12_resource_state(unsigned int state, unsigned int stencil_state,
         const struct d3d12_resource *resource, VkQueueFlags vk_queue_flags, const struct vkd3d_vulkan_info *vk_info,
         VkAccessFlags *access_mask, VkPipelineStageFlags *stage_flags, VkImageLayout *image_layout)
@@ -4039,9 +4052,8 @@ static void d3d12_command_list_set_root_signature(struct d3d12_command_list *lis
         return;

bindings->root_signature = root_signature;
-    bindings->descriptor_set = VK_NULL_HANDLE;
-    bindings->descriptor_table_dirty_mask = bindings->descriptor_table_active_mask & root_signature->descriptor_table_mask;
-    bindings->push_descriptor_dirty_mask = bindings->push_descriptor_active_mask & root_signature->push_descriptor_mask;
+
+    d3d12_command_list_invalidate_root_parameters(list, bind_point);
 }

static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootSignature(ID3D12GraphicsCommandList1 *iface,

[PATCH vkd3d 04/13] vkd3d: Factor out root parameter invalidation.

[Henri Verbeet]

Signed-off-by: Henri Verbeet

[PATCH vkd3d 05/13] vkd3d: Introduce meta ops for shader-based UAV clears.

[Philip Rebohle]

Signed-off-by: Philip Rebohle 
---
 Makefile.am                    |   2 +
 libs/vkd3d/device.c            |   6 +
 libs/vkd3d/meta.c              | 325 +++++++++++++
 libs/vkd3d/vkd3d_private.h     |  62 +++
 libs/vkd3d/vkd3d_spv_shaders.h | 830 +++++++++++++++++++++++++++++++++
 5 files changed, 1225 insertions(+)
 create mode 100644 libs/vkd3d/meta.c
 create mode 100644 libs/vkd3d/vkd3d_spv_shaders.h

diff --git a/Makefile.am b/Makefile.am
index 2b0e8f3..75d2a97 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -102,12 +102,14 @@ libvkd3d_la_SOURCES = \
    include/vkd3d_unknown.idl \
    libs/vkd3d/command.c \
    libs/vkd3d/device.c \
+	libs/vkd3d/meta.c \
    libs/vkd3d/resource.c \
    libs/vkd3d/state.c \
    libs/vkd3d/utils.c \
    libs/vkd3d/vkd3d.map \
    libs/vkd3d/vkd3d_main.c \
    libs/vkd3d/vkd3d_private.h \
+	libs/vkd3d/vkd3d_spv_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/device.c b/libs/vkd3d/device.c
index 0624318..114d671 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_meta_ops_destroy(&device->meta_ops, 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_meta_ops_init(&device->meta_ops, 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/meta.c b/libs/vkd3d/meta.c
new file mode 100644
index 0000000..9a18cf6
--- /dev/null
+++ b/libs/vkd3d/meta.c
@@ -0,0 +1,325 @@
+#include "vkd3d_private.h"
+#include "vkd3d_spv_shaders.h"
+
+#define SPIRV_CODE(name) name, sizeof(name)
+
+static VkResult vkd3d_create_shader_module(struct d3d12_device *device,
+        size_t code_size, const uint32_t *code, VkShaderModule *module)
+{
+    const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+    VkShaderModuleCreateInfo shader_module_info;
+
+    shader_module_info.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO;
+    shader_module_info.pNext = NULL;
+    shader_module_info.flags = 0;
+    shader_module_info.codeSize = code_size;
+    shader_module_info.pCode = code;
+
+    return VK_CALL(vkCreateShaderModule(device->vk_device, &shader_module_info, NULL, module));
+}
+
+static VkResult vkd3d_create_descriptor_set_layout(struct d3d12_device *device,
+        uint32_t binding_count, const VkDescriptorSetLayoutBinding *bindings, VkDescriptorSetLayout *set_layout)
+{
+    const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+    VkDescriptorSetLayoutCreateInfo set_layout_info;
+
+    set_layout_info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO;
+    set_layout_info.pNext = NULL;
+    set_layout_info.flags = 0;
+    set_layout_info.bindingCount = binding_count;
+    set_layout_info.pBindings = bindings;
+
+    return VK_CALL(vkCreateDescriptorSetLayout(device->vk_device, &set_layout_info, NULL, set_layout));
+}
+
+static VkResult vkd3d_create_pipeline_layout(struct d3d12_device *device,
+        uint32_t set_layout_count, const VkDescriptorSetLayout *set_layouts,
+        uint32_t push_constant_range_count, const VkPushConstantRange *push_constant_ranges,
+        VkPipelineLayout *pipeline_layout)
+{
+    const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+    VkPipelineLayoutCreateInfo pipeline_layout_info;
+
+    pipeline_layout_info.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO;
+    pipeline_layout_info.pNext = NULL;
+    pipeline_layout_info.flags = 0;
+    pipeline_layout_info.setLayoutCount = set_layout_count;
+    pipeline_layout_info.pSetLayouts = set_layouts;
+    pipeline_layout_info.pushConstantRangeCount = push_constant_range_count;
+    pipeline_layout_info.pPushConstantRanges = push_constant_ranges;
+
+    return VK_CALL(vkCreatePipelineLayout(device->vk_device, &pipeline_layout_info, NULL, pipeline_layout));
+}
+
+static VkResult vkd3d_create_compute_pipeline(struct d3d12_device *device,
+        size_t code_size, const uint32_t *code, VkPipelineLayout layout,
+        const VkSpecializationInfo *specialization_info, VkPipeline *pipeline)
+{
+    const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+    VkComputePipelineCreateInfo pipeline_info;
+    VkResult vr;
+
+    pipeline_info.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;
+    pipeline_info.pNext = NULL;
+    pipeline_info.flags = 0;
+    pipeline_info.layout = layout;
+    pipeline_info.basePipelineHandle = VK_NULL_HANDLE;
+    pipeline_info.basePipelineIndex = -1;
+
+    pipeline_info.stage.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO;
+    pipeline_info.stage.pNext = NULL;
+    pipeline_info.stage.flags = 0;
+    pipeline_info.stage.stage = VK_SHADER_STAGE_COMPUTE_BIT;
+    pipeline_info.stage.pName = "main";
+    pipeline_info.stage.pSpecializationInfo = specialization_info;
+
+    if ((vr = vkd3d_create_shader_module(device, code_size, code, &pipeline_info.stage.module)) < 0)
+    {
+        ERR("Failed to create shader module, vr %d.", vr);
+        return vr;
+    }
+
+    vr = VK_CALL(vkCreateComputePipelines(device->vk_device, VK_NULL_HANDLE, 1, &pipeline_info, NULL, pipeline));
+    VK_CALL(vkDestroyShaderModule(device->vk_device, pipeline_info.stage.module, NULL));
+
+    return vr;
+}
+
+HRESULT vkd3d_clear_uav_ops_init(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+        struct d3d12_device *device)
+{
+    VkDescriptorSetLayoutBinding set_binding;
+    VkPushConstantRange push_constant_range;
+    unsigned int i;
+    VkResult vr;
+
+    struct {
+      VkDescriptorSetLayout *set_layout;
+      VkPipelineLayout *pipeline_layout;
+      VkDescriptorType descriptor_type;
+    }
+    set_layouts[] =
+    {
+      { &meta_clear_uav_ops->vk_set_layout_buffer, &meta_clear_uav_ops->vk_pipeline_layout_buffer, VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER },
+      { &meta_clear_uav_ops->vk_set_layout_image,  &meta_clear_uav_ops->vk_pipeline_layout_image,  VK_DESCRIPTOR_TYPE_STORAGE_IMAGE },
+    };
+
+    struct {
+      VkPipeline *pipeline;
+      VkPipelineLayout *pipeline_layout;
+      const uint32_t *code;
+      size_t code_size;
+    }
+    pipelines[] =
+    {
+      { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_buffer,
+        &meta_clear_uav_ops->vk_pipeline_layout_buffer,
+        SPIRV_CODE(cs_clear_uav_buffer_float_spv) },
+      { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_image_1d,
+        &meta_clear_uav_ops->vk_pipeline_layout_image,
+        SPIRV_CODE(cs_clear_uav_image_1d_float_spv) },
+      { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_image_1d_array,
+        &meta_clear_uav_ops->vk_pipeline_layout_image,
+        SPIRV_CODE(cs_clear_uav_image_1d_array_float_spv) },
+      { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_image_2d,
+        &meta_clear_uav_ops->vk_pipeline_layout_image,
+        SPIRV_CODE(cs_clear_uav_image_2d_float_spv) },
+      { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_image_2d_array,
+        &meta_clear_uav_ops->vk_pipeline_layout_image,
+        SPIRV_CODE(cs_clear_uav_image_2d_array_float_spv) },
+      { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_image_3d,
+        &meta_clear_uav_ops->vk_pipeline_layout_image,
+        SPIRV_CODE(cs_clear_uav_image_3d_float_spv) },
+      { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_buffer,
+        &meta_clear_uav_ops->vk_pipeline_layout_buffer,
+        SPIRV_CODE(cs_clear_uav_buffer_uint_spv) },
+      { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_image_1d,
+        &meta_clear_uav_ops->vk_pipeline_layout_image,
+        SPIRV_CODE(cs_clear_uav_image_1d_uint_spv) },
+      { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_image_1d_array,
+        &meta_clear_uav_ops->vk_pipeline_layout_image,
+        SPIRV_CODE(cs_clear_uav_image_1d_array_uint_spv) },
+      { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_image_2d,
+        &meta_clear_uav_ops->vk_pipeline_layout_image,
+        SPIRV_CODE(cs_clear_uav_image_2d_uint_spv) },
+      { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_image_2d_array,
+        &meta_clear_uav_ops->vk_pipeline_layout_image,
+        SPIRV_CODE(cs_clear_uav_image_2d_array_uint_spv) },
+      { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_image_3d,
+        &meta_clear_uav_ops->vk_pipeline_layout_image,
+        SPIRV_CODE(cs_clear_uav_image_3d_uint_spv) },
+    };
+
+    memset(meta_clear_uav_ops, 0, sizeof(*meta_clear_uav_ops));
+
+    set_binding.binding = 0;
+    set_binding.descriptorCount = 1;
+    set_binding.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
+    set_binding.pImmutableSamplers = NULL;
+
+    push_constant_range.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
+    push_constant_range.offset = 0;
+    push_constant_range.size = sizeof(struct vkd3d_clear_uav_args);
+
+    for (i = 0; i < ARRAY_SIZE(set_layouts); i++)
+    {
+        set_binding.descriptorType = set_layouts[i].descriptor_type;
+
+        vr = vkd3d_create_descriptor_set_layout(device, 1, &set_binding, set_layouts[i].set_layout);
+
+        if (vr < 0)
+        {
+            ERR("Failed to create descriptor set layout %u, vr %d.", i, vr);
+            goto fail;
+        }
+
+        vr = vkd3d_create_pipeline_layout(device, 1, set_layouts[i].set_layout,
+                1, &push_constant_range, set_layouts[i].pipeline_layout);
+
+        if (vr < 0)
+        {
+            ERR("Failed to create pipeline layout %u, vr %d.", i, vr);
+            goto fail;
+        }
+    }
+
+    for (i = 0; i < ARRAY_SIZE(pipelines); i++)
+    {
+        if ((vr = vkd3d_create_compute_pipeline(device, pipelines[i].code_size, pipelines[i].code,
+                *pipelines[i].pipeline_layout, NULL, pipelines[i].pipeline)) < 0)
+        {
+            ERR("Failed to create compute pipeline %u, vr %d.", i, vr);
+            goto fail;
+        }
+    }
+
+    return S_OK;
+fail:
+    vkd3d_clear_uav_ops_destroy(meta_clear_uav_ops, device);
+    return hresult_from_vk_result(vr);
+}
+
+void vkd3d_clear_uav_ops_destroy(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+        struct d3d12_device *device) {
+    const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+    unsigned int i;
+
+    struct vkd3d_clear_uav_pipelines* pipeline_sets[] =
+    {
+        &meta_clear_uav_ops->clear_float,
+        &meta_clear_uav_ops->clear_uint,
+    };
+
+    VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, meta_clear_uav_ops->vk_set_layout_buffer, NULL));
+    VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, meta_clear_uav_ops->vk_set_layout_image, NULL));
+
+    VK_CALL(vkDestroyPipelineLayout(device->vk_device, meta_clear_uav_ops->vk_pipeline_layout_buffer, NULL));
+    VK_CALL(vkDestroyPipelineLayout(device->vk_device, meta_clear_uav_ops->vk_pipeline_layout_image, NULL));
+
+    for (i = 0; i < ARRAY_SIZE(pipeline_sets); i++)
+    {
+        VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_buffer, NULL));
+        VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_image_1d, NULL));
+        VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_image_2d, NULL));
+        VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_image_3d, NULL));
+        VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_image_1d_array, NULL));
+        VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_image_2d_array, NULL));
+    }
+}
+
+struct vkd3d_clear_uav_pipeline vkd3d_clear_uav_ops_get_clear_buffer_pipeline(const struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+        bool as_uint)
+{
+    struct vkd3d_clear_uav_pipeline info;
+
+    const struct vkd3d_clear_uav_pipelines *pipelines = as_uint
+            ? &meta_clear_uav_ops->clear_uint
+            : &meta_clear_uav_ops->clear_float;
+
+    info.vk_set_layout = meta_clear_uav_ops->vk_set_layout_buffer;
+    info.vk_pipeline_layout = meta_clear_uav_ops->vk_pipeline_layout_buffer;
+    info.vk_pipeline = pipelines->vk_pipeline_clear_buffer;
+    return info;
+}
+
+struct vkd3d_clear_uav_pipeline vkd3d_clear_uav_ops_get_clear_image_pipeline(const struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+        VkImageViewType image_view_type, bool as_uint)
+{
+    struct vkd3d_clear_uav_pipeline info;
+
+    const struct vkd3d_clear_uav_pipelines *pipelines = as_uint
+            ? &meta_clear_uav_ops->clear_uint
+            : &meta_clear_uav_ops->clear_float;
+
+    info.vk_set_layout = meta_clear_uav_ops->vk_set_layout_image;
+    info.vk_pipeline_layout = meta_clear_uav_ops->vk_pipeline_layout_image;
+
+    switch (image_view_type)
+    {
+        case VK_IMAGE_VIEW_TYPE_1D:
+            info.vk_pipeline = pipelines->vk_pipeline_clear_image_1d;
+            break;
+        case VK_IMAGE_VIEW_TYPE_2D:
+            info.vk_pipeline = pipelines->vk_pipeline_clear_image_2d;
+            break;
+        case VK_IMAGE_VIEW_TYPE_3D:
+            info.vk_pipeline = pipelines->vk_pipeline_clear_image_3d;
+            break;
+        case VK_IMAGE_VIEW_TYPE_1D_ARRAY:
+            info.vk_pipeline = pipelines->vk_pipeline_clear_image_1d_array;
+            break;
+        case VK_IMAGE_VIEW_TYPE_2D_ARRAY:
+            info.vk_pipeline = pipelines->vk_pipeline_clear_image_2d_array;
+            break;
+        default:
+            ERR("Unhandled view type %d.\n", image_view_type);
+            info.vk_pipeline = VK_NULL_HANDLE;
+    }
+
+    return info;
+}
+
+VkExtent3D vkd3d_get_clear_image_uav_workgroup_size(VkImageViewType view_type)
+{
+    switch (view_type)
+    {
+        case VK_IMAGE_VIEW_TYPE_1D:
+        case VK_IMAGE_VIEW_TYPE_1D_ARRAY:
+        {
+            VkExtent3D result = { 64, 1, 1 };
+            return result;
+        }
+        case VK_IMAGE_VIEW_TYPE_2D:
+        case VK_IMAGE_VIEW_TYPE_2D_ARRAY:
+        case VK_IMAGE_VIEW_TYPE_3D:
+        {
+            VkExtent3D result = { 8, 8, 1 };
+            return result;
+        }
+        default:
+        {
+            VkExtent3D result = { 0, 0, 0 };
+            ERR("Unhandled view type %d.\n", view_type);
+            return result;
+        }
+    }
+}
+
+HRESULT vkd3d_meta_ops_init(struct vkd3d_meta_ops *meta_ops, struct d3d12_device *device)
+{
+    HRESULT hr;
+
+    memset(meta_ops, 0, sizeof(*meta_ops));
+
+    if (FAILED(hr = vkd3d_clear_uav_ops_init(&meta_ops->clear_uav, device)))
+        return hr;
+
+    return S_OK;
+}
+
+HRESULT vkd3d_meta_ops_destroy(struct vkd3d_meta_ops *meta_ops, struct d3d12_device *device)
+{
+    vkd3d_clear_uav_ops_destroy(&meta_ops->clear_uav, device);
+    return S_OK;
+}
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index 84b5ff2..d0224d3 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -1059,6 +1059,67 @@ struct vkd3d_format_compatibility_list
     VkFormat vk_formats[VKD3D_MAX_COMPATIBLE_FORMAT_COUNT];
 };

+/* meta operations */
+struct vkd3d_clear_uav_args
+{
+    VkClearColorValue clear_color;
+    VkOffset2D offset;
+    VkExtent2D extent;
+};
+
+struct vkd3d_clear_uav_pipelines
+{
+    VkPipeline vk_pipeline_clear_buffer;
+    VkPipeline vk_pipeline_clear_image_1d;
+    VkPipeline vk_pipeline_clear_image_2d;
+    VkPipeline vk_pipeline_clear_image_3d;
+    VkPipeline vk_pipeline_clear_image_1d_array;
+    VkPipeline vk_pipeline_clear_image_2d_array;
+};
+
+struct vkd3d_clear_uav_ops
+{
+    VkDescriptorSetLayout vk_set_layout_buffer;
+    VkDescriptorSetLayout vk_set_layout_image;
+
+    VkPipelineLayout vk_pipeline_layout_buffer;
+    VkPipelineLayout vk_pipeline_layout_image;
+
+    struct vkd3d_clear_uav_pipelines clear_float;
+    struct vkd3d_clear_uav_pipelines clear_uint;
+};
+
+struct vkd3d_clear_uav_pipeline
+{
+    VkDescriptorSetLayout vk_set_layout;
+    VkPipelineLayout vk_pipeline_layout;
+    VkPipeline vk_pipeline;
+};
+
+HRESULT vkd3d_clear_uav_ops_init(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+        struct d3d12_device *device) DECLSPEC_HIDDEN;
+void vkd3d_clear_uav_ops_destroy(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+        struct d3d12_device *device) DECLSPEC_HIDDEN;
+struct vkd3d_clear_uav_pipeline vkd3d_clear_uav_ops_get_clear_buffer_pipeline(const struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+        bool as_uint) DECLSPEC_HIDDEN;
+struct vkd3d_clear_uav_pipeline vkd3d_clear_uav_ops_get_clear_image_pipeline(const struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+        VkImageViewType image_view_type, bool as_uint) DECLSPEC_HIDDEN;
+VkExtent3D vkd3d_get_clear_image_uav_workgroup_size(VkImageViewType view_type) DECLSPEC_HIDDEN;
+
+inline VkExtent3D vkd3d_get_clear_buffer_uav_workgroup_size()
+{
+    VkExtent3D result = { 128, 1, 1 };
+    return result;
+}
+
+struct vkd3d_meta_ops
+{
+    struct vkd3d_clear_uav_ops clear_uav;
+};
+
+HRESULT vkd3d_meta_ops_init(struct vkd3d_meta_ops *meta_ops, struct d3d12_device *device) DECLSPEC_HIDDEN;
+HRESULT vkd3d_meta_ops_destroy(struct vkd3d_meta_ops *meta_ops, struct d3d12_device *device) DECLSPEC_HIDDEN;
+
 /* ID3D12Device */
 struct d3d12_device
 {
@@ -1104,6 +1165,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_meta_ops meta_ops;
 };

HRESULT d3d12_device_create(struct vkd3d_instance *instance,
diff --git a/libs/vkd3d/vkd3d_spv_shaders.h b/libs/vkd3d/vkd3d_spv_shaders.h
new file mode 100644
index 0000000..bf86f6b
--- /dev/null
+++ b/libs/vkd3d/vkd3d_spv_shaders.h
@@ -0,0 +1,830 @@
+#ifndef __VKD3D_SPV_SHADERS_H
+#define __VKD3D_SPV_SHADERS_H
+
+const uint32_t cs_clear_uav_buffer_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 128) in;
+
+layout(binding = 0)
+writeonly uniform imageBuffer dst;
+
+layout(push_constant)
+uniform u_info_t {
+  vec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  int thread_id = int(gl_GlobalInvocationID.x);
+  
+  if (thread_id < u_info.dst_extent.x)
+    imageStore(dst, u_info.dst_offset.x + thread_id, u_info.clear_value);
+}
+#endif
+    0x07230203,0x00010000,0x00080007,0x00000031,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x0000002f,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+    0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+    0x00000000,0x0000000c,0x00060010,0x00000004,0x00000011,0x00000080,0x00000001,0x00000001,
+    0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+    0x00000008,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000c,0x475f6c67,0x61626f6c,
+    0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000016,0x6e695f75,0x745f6f66,
+    0x00000000,0x00060006,0x00000016,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+    0x00000016,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000016,0x00000002,
+    0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000018,0x6e695f75,0x00006f66,0x00030005,
+    0x00000023,0x00747364,0x00040047,0x0000000c,0x0000000b,0x0000001c,0x00050048,0x00000016,
+    0x00000000,0x00000023,0x00000000,0x00050048,0x00000016,0x00000001,0x00000023,0x00000010,
+    0x00050048,0x00000016,0x00000002,0x00000023,0x00000018,0x00030047,0x00000016,0x00000002,
+    0x00040047,0x00000023,0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,
+    0x00030047,0x00000023,0x00000019,0x00040047,0x00000030,0x0000000b,0x00000019,0x00020013,
+    0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+    0x00040020,0x00000007,0x00000007,0x00000006,0x00040015,0x00000009,0x00000020,0x00000000,
+    0x00040017,0x0000000a,0x00000009,0x00000003,0x00040020,0x0000000b,0x00000001,0x0000000a,
+    0x0004003b,0x0000000b,0x0000000c,0x00000001,0x0004002b,0x00000009,0x0000000d,0x00000000,
+    0x00040020,0x0000000e,0x00000001,0x00000009,0x00030016,0x00000013,0x00000020,0x00040017,
+    0x00000014,0x00000013,0x00000004,0x00040017,0x00000015,0x00000006,0x00000002,0x0005001e,
+    0x00000016,0x00000014,0x00000015,0x00000015,0x00040020,0x00000017,0x00000009,0x00000016,
+    0x0004003b,0x00000017,0x00000018,0x00000009,0x0004002b,0x00000006,0x00000019,0x00000002,
+    0x00040020,0x0000001a,0x00000009,0x00000006,0x00020014,0x0000001d,0x00090019,0x00000021,
+    0x00000013,0x00000005,0x00000000,0x00000000,0x00000000,0x00000002,0x00000000,0x00040020,
+    0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,0x00000023,0x00000000,0x0004002b,
+    0x00000006,0x00000025,0x00000001,0x0004002b,0x00000006,0x0000002a,0x00000000,0x00040020,
+    0x0000002b,0x00000009,0x00000014,0x0004002b,0x00000009,0x0000002e,0x00000080,0x0004002b,
+    0x00000009,0x0000002f,0x00000001,0x0006002c,0x0000000a,0x00000030,0x0000002e,0x0000002f,
+    0x0000002f,0x00050036,0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,
+    0x0004003b,0x00000007,0x00000008,0x00000007,0x00050041,0x0000000e,0x0000000f,0x0000000c,
+    0x0000000d,0x0004003d,0x00000009,0x00000010,0x0000000f,0x0004007c,0x00000006,0x00000011,
+    0x00000010,0x0003003e,0x00000008,0x00000011,0x0004003d,0x00000006,0x00000012,0x00000008,
+    0x00060041,0x0000001a,0x0000001b,0x00000018,0x00000019,0x0000000d,0x0004003d,0x00000006,
+    0x0000001c,0x0000001b,0x000500b1,0x0000001d,0x0000001e,0x00000012,0x0000001c,0x000300f7,
+    0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,0x0000001f,
+    0x0004003d,0x00000021,0x00000024,0x00000023,0x00060041,0x0000001a,0x00000026,0x00000018,
+    0x00000025,0x0000000d,0x0004003d,0x00000006,0x00000027,0x00000026,0x0004003d,0x00000006,
+    0x00000028,0x00000008,0x00050080,0x00000006,0x00000029,0x00000027,0x00000028,0x00050041,
+    0x0000002b,0x0000002c,0x00000018,0x0000002a,0x0004003d,0x00000014,0x0000002d,0x0000002c,
+    0x00040063,0x00000024,0x00000029,0x0000002d,0x000200f9,0x00000020,0x000200f8,0x00000020,
+    0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_buffer_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 128) in;
+
+layout(binding = 0)
+writeonly uniform uimageBuffer dst;
+
+layout(push_constant)
+uniform u_info_t {
+  uvec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  int thread_id = int(gl_GlobalInvocationID.x);
+  
+  if (thread_id < u_info.dst_extent.x)
+    imageStore(dst, u_info.dst_offset.x + thread_id, u_info.clear_value);
+}
+#endif
+    0x07230203,0x00010000,0x00080007,0x00000030,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x0000002f,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+    0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+    0x00000000,0x0000000c,0x00060010,0x00000004,0x00000011,0x00000080,0x00000001,0x00000001,
+    0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+    0x00000008,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000c,0x475f6c67,0x61626f6c,
+    0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000015,0x6e695f75,0x745f6f66,
+    0x00000000,0x00060006,0x00000015,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+    0x00000015,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000015,0x00000002,
+    0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000017,0x6e695f75,0x00006f66,0x00030005,
+    0x00000022,0x00747364,0x00040047,0x0000000c,0x0000000b,0x0000001c,0x00050048,0x00000015,
+    0x00000000,0x00000023,0x00000000,0x00050048,0x00000015,0x00000001,0x00000023,0x00000010,
+    0x00050048,0x00000015,0x00000002,0x00000023,0x00000018,0x00030047,0x00000015,0x00000002,
+    0x00040047,0x00000022,0x00000022,0x00000000,0x00040047,0x00000022,0x00000021,0x00000000,
+    0x00030047,0x00000022,0x00000019,0x00040047,0x0000002f,0x0000000b,0x00000019,0x00020013,
+    0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+    0x00040020,0x00000007,0x00000007,0x00000006,0x00040015,0x00000009,0x00000020,0x00000000,
+    0x00040017,0x0000000a,0x00000009,0x00000003,0x00040020,0x0000000b,0x00000001,0x0000000a,
+    0x0004003b,0x0000000b,0x0000000c,0x00000001,0x0004002b,0x00000009,0x0000000d,0x00000000,
+    0x00040020,0x0000000e,0x00000001,0x00000009,0x00040017,0x00000013,0x00000009,0x00000004,
+    0x00040017,0x00000014,0x00000006,0x00000002,0x0005001e,0x00000015,0x00000013,0x00000014,
+    0x00000014,0x00040020,0x00000016,0x00000009,0x00000015,0x0004003b,0x00000016,0x00000017,
+    0x00000009,0x0004002b,0x00000006,0x00000018,0x00000002,0x00040020,0x00000019,0x00000009,
+    0x00000006,0x00020014,0x0000001c,0x00090019,0x00000020,0x00000009,0x00000005,0x00000000,
+    0x00000000,0x00000000,0x00000002,0x00000000,0x00040020,0x00000021,0x00000000,0x00000020,
+    0x0004003b,0x00000021,0x00000022,0x00000000,0x0004002b,0x00000006,0x00000024,0x00000001,
+    0x0004002b,0x00000006,0x00000029,0x00000000,0x00040020,0x0000002a,0x00000009,0x00000013,
+    0x0004002b,0x00000009,0x0000002d,0x00000080,0x0004002b,0x00000009,0x0000002e,0x00000001,
+    0x0006002c,0x0000000a,0x0000002f,0x0000002d,0x0000002e,0x0000002e,0x00050036,0x00000002,
+    0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000007,0x00000008,
+    0x00000007,0x00050041,0x0000000e,0x0000000f,0x0000000c,0x0000000d,0x0004003d,0x00000009,
+    0x00000010,0x0000000f,0x0004007c,0x00000006,0x00000011,0x00000010,0x0003003e,0x00000008,
+    0x00000011,0x0004003d,0x00000006,0x00000012,0x00000008,0x00060041,0x00000019,0x0000001a,
+    0x00000017,0x00000018,0x0000000d,0x0004003d,0x00000006,0x0000001b,0x0000001a,0x000500b1,
+    0x0000001c,0x0000001d,0x00000012,0x0000001b,0x000300f7,0x0000001f,0x00000000,0x000400fa,
+    0x0000001d,0x0000001e,0x0000001f,0x000200f8,0x0000001e,0x0004003d,0x00000020,0x00000023,
+    0x00000022,0x00060041,0x00000019,0x00000025,0x00000017,0x00000024,0x0000000d,0x0004003d,
+    0x00000006,0x00000026,0x00000025,0x0004003d,0x00000006,0x00000027,0x00000008,0x00050080,
+    0x00000006,0x00000028,0x00000026,0x00000027,0x00050041,0x0000002a,0x0000002b,0x00000017,
+    0x00000029,0x0004003d,0x00000013,0x0000002c,0x0000002b,0x00040063,0x00000023,0x00000028,
+    0x0000002c,0x000200f9,0x0000001f,0x000200f8,0x0000001f,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_1d_array_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 64) in;
+
+layout(binding = 0)
+writeonly uniform image1DArray dst;
+
+layout(push_constant)
+uniform u_info_t {
+  vec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+  
+  if (thread_id.x < u_info.dst_extent.x)
+    imageStore(dst, ivec2(u_info.dst_offset.x + thread_id.x, thread_id.y), u_info.clear_value);
+}
+#endif
+    0x07230203,0x00010000,0x00080007,0x00000036,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x0000002c,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+    0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+    0x00000000,0x0000000d,0x00060010,0x00000004,0x00000011,0x00000040,0x00000001,0x00000001,
+    0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+    0x00000009,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,
+    0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000017,0x6e695f75,0x745f6f66,
+    0x00000000,0x00060006,0x00000017,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+    0x00000017,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000017,0x00000002,
+    0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000019,0x6e695f75,0x00006f66,0x00030005,
+    0x00000024,0x00747364,0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000017,
+    0x00000000,0x00000023,0x00000000,0x00050048,0x00000017,0x00000001,0x00000023,0x00000010,
+    0x00050048,0x00000017,0x00000002,0x00000023,0x00000018,0x00030047,0x00000017,0x00000002,
+    0x00040047,0x00000024,0x00000022,0x00000000,0x00040047,0x00000024,0x00000021,0x00000000,
+    0x00030047,0x00000024,0x00000019,0x00040047,0x00000035,0x0000000b,0x00000019,0x00020013,
+    0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+    0x00040017,0x00000007,0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,
+    0x00040015,0x0000000a,0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,
+    0x00040020,0x0000000c,0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,
+    0x0004002b,0x0000000a,0x00000010,0x00000000,0x00040020,0x00000011,0x00000007,0x00000006,
+    0x00030016,0x00000014,0x00000020,0x00040017,0x00000015,0x00000014,0x00000004,0x00040017,
+    0x00000016,0x00000006,0x00000002,0x0005001e,0x00000017,0x00000015,0x00000016,0x00000016,
+    0x00040020,0x00000018,0x00000009,0x00000017,0x0004003b,0x00000018,0x00000019,0x00000009,
+    0x0004002b,0x00000006,0x0000001a,0x00000002,0x00040020,0x0000001b,0x00000009,0x00000006,
+    0x00020014,0x0000001e,0x00090019,0x00000022,0x00000014,0x00000000,0x00000000,0x00000001,
+    0x00000000,0x00000002,0x00000000,0x00040020,0x00000023,0x00000000,0x00000022,0x0004003b,
+    0x00000023,0x00000024,0x00000000,0x0004002b,0x00000006,0x00000026,0x00000001,0x0004002b,
+    0x0000000a,0x0000002c,0x00000001,0x0004002b,0x00000006,0x00000030,0x00000000,0x00040020,
+    0x00000031,0x00000009,0x00000015,0x0004002b,0x0000000a,0x00000034,0x00000040,0x0006002c,
+    0x0000000b,0x00000035,0x00000034,0x0000002c,0x0000002c,0x00050036,0x00000002,0x00000004,
+    0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,
+    0x0004003d,0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,
+    0x0003003e,0x00000009,0x0000000f,0x00050041,0x00000011,0x00000012,0x00000009,0x00000010,
+    0x0004003d,0x00000006,0x00000013,0x00000012,0x00060041,0x0000001b,0x0000001c,0x00000019,
+    0x0000001a,0x00000010,0x0004003d,0x00000006,0x0000001d,0x0000001c,0x000500b1,0x0000001e,
+    0x0000001f,0x00000013,0x0000001d,0x000300f7,0x00000021,0x00000000,0x000400fa,0x0000001f,
+    0x00000020,0x00000021,0x000200f8,0x00000020,0x0004003d,0x00000022,0x00000025,0x00000024,
+    0x00060041,0x0000001b,0x00000027,0x00000019,0x00000026,0x00000010,0x0004003d,0x00000006,
+    0x00000028,0x00000027,0x00050041,0x00000011,0x00000029,0x00000009,0x00000010,0x0004003d,
+    0x00000006,0x0000002a,0x00000029,0x00050080,0x00000006,0x0000002b,0x00000028,0x0000002a,
+    0x00050041,0x00000011,0x0000002d,0x00000009,0x0000002c,0x0004003d,0x00000006,0x0000002e,
+    0x0000002d,0x00050050,0x00000016,0x0000002f,0x0000002b,0x0000002e,0x00050041,0x00000031,
+    0x00000032,0x00000019,0x00000030,0x0004003d,0x00000015,0x00000033,0x00000032,0x00040063,
+    0x00000025,0x0000002f,0x00000033,0x000200f9,0x00000021,0x000200f8,0x00000021,0x000100fd,
+    0x00010038
+};
+
+const uint32_t cs_clear_uav_image_1d_array_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 64) in;
+
+layout(binding = 0)
+writeonly uniform uimage1DArray dst;
+
+layout(push_constant)
+uniform u_info_t {
+  uvec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+  
+  if (thread_id.x < u_info.dst_extent.x)
+    imageStore(dst, ivec2(u_info.dst_offset.x + thread_id.x, thread_id.y), u_info.clear_value);
+}
+#endif
+    0x07230203,0x00010000,0x00080007,0x00000035,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x0000002c,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+    0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+    0x00000000,0x0000000d,0x00060010,0x00000004,0x00000011,0x00000040,0x00000001,0x00000001,
+    0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+    0x00000009,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,
+    0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000016,0x6e695f75,0x745f6f66,
+    0x00000000,0x00060006,0x00000016,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+    0x00000016,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000016,0x00000002,
+    0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000018,0x6e695f75,0x00006f66,0x00030005,
+    0x00000023,0x00747364,0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000016,
+    0x00000000,0x00000023,0x00000000,0x00050048,0x00000016,0x00000001,0x00000023,0x00000010,
+    0x00050048,0x00000016,0x00000002,0x00000023,0x00000018,0x00030047,0x00000016,0x00000002,
+    0x00040047,0x00000023,0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,
+    0x00030047,0x00000023,0x00000019,0x00040047,0x00000034,0x0000000b,0x00000019,0x00020013,
+    0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+    0x00040017,0x00000007,0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,
+    0x00040015,0x0000000a,0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,
+    0x00040020,0x0000000c,0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,
+    0x0004002b,0x0000000a,0x00000010,0x00000000,0x00040020,0x00000011,0x00000007,0x00000006,
+    0x00040017,0x00000014,0x0000000a,0x00000004,0x00040017,0x00000015,0x00000006,0x00000002,
+    0x0005001e,0x00000016,0x00000014,0x00000015,0x00000015,0x00040020,0x00000017,0x00000009,
+    0x00000016,0x0004003b,0x00000017,0x00000018,0x00000009,0x0004002b,0x00000006,0x00000019,
+    0x00000002,0x00040020,0x0000001a,0x00000009,0x00000006,0x00020014,0x0000001d,0x00090019,
+    0x00000021,0x0000000a,0x00000000,0x00000000,0x00000001,0x00000000,0x00000002,0x00000000,
+    0x00040020,0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,0x00000023,0x00000000,
+    0x0004002b,0x00000006,0x00000025,0x00000001,0x0004002b,0x0000000a,0x0000002b,0x00000001,
+    0x0004002b,0x00000006,0x0000002f,0x00000000,0x00040020,0x00000030,0x00000009,0x00000014,
+    0x0004002b,0x0000000a,0x00000033,0x00000040,0x0006002c,0x0000000b,0x00000034,0x00000033,
+    0x0000002b,0x0000002b,0x00050036,0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,
+    0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,0x0004003d,0x0000000b,0x0000000e,
+    0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,0x0003003e,0x00000009,0x0000000f,
+    0x00050041,0x00000011,0x00000012,0x00000009,0x00000010,0x0004003d,0x00000006,0x00000013,
+    0x00000012,0x00060041,0x0000001a,0x0000001b,0x00000018,0x00000019,0x00000010,0x0004003d,
+    0x00000006,0x0000001c,0x0000001b,0x000500b1,0x0000001d,0x0000001e,0x00000013,0x0000001c,
+    0x000300f7,0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,
+    0x0000001f,0x0004003d,0x00000021,0x00000024,0x00000023,0x00060041,0x0000001a,0x00000026,
+    0x00000018,0x00000025,0x00000010,0x0004003d,0x00000006,0x00000027,0x00000026,0x00050041,
+    0x00000011,0x00000028,0x00000009,0x00000010,0x0004003d,0x00000006,0x00000029,0x00000028,
+    0x00050080,0x00000006,0x0000002a,0x00000027,0x00000029,0x00050041,0x00000011,0x0000002c,
+    0x00000009,0x0000002b,0x0004003d,0x00000006,0x0000002d,0x0000002c,0x00050050,0x00000015,
+    0x0000002e,0x0000002a,0x0000002d,0x00050041,0x00000030,0x00000031,0x00000018,0x0000002f,
+    0x0004003d,0x00000014,0x00000032,0x00000031,0x00040063,0x00000024,0x0000002e,0x00000032,
+    0x000200f9,0x00000020,0x000200f8,0x00000020,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_1d_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 64) in;
+
+layout(binding = 0)
+writeonly uniform image1D dst;
+
+layout(push_constant)
+uniform u_info_t {
+  vec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+  
+  if (thread_id.x < u_info.dst_extent.x)
+    imageStore(dst, u_info.dst_offset.x + thread_id.x, u_info.clear_value);
+}
+#endif
+    0x07230203,0x00010000,0x00080007,0x00000033,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x0000002c,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+    0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+    0x00000000,0x0000000d,0x00060010,0x00000004,0x00000011,0x00000040,0x00000001,0x00000001,
+    0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+    0x00000009,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,
+    0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000017,0x6e695f75,0x745f6f66,
+    0x00000000,0x00060006,0x00000017,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+    0x00000017,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000017,0x00000002,
+    0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000019,0x6e695f75,0x00006f66,0x00030005,
+    0x00000024,0x00747364,0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000017,
+    0x00000000,0x00000023,0x00000000,0x00050048,0x00000017,0x00000001,0x00000023,0x00000010,
+    0x00050048,0x00000017,0x00000002,0x00000023,0x00000018,0x00030047,0x00000017,0x00000002,
+    0x00040047,0x00000024,0x00000022,0x00000000,0x00040047,0x00000024,0x00000021,0x00000000,
+    0x00030047,0x00000024,0x00000019,0x00040047,0x00000032,0x0000000b,0x00000019,0x00020013,
+    0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+    0x00040017,0x00000007,0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,
+    0x00040015,0x0000000a,0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,
+    0x00040020,0x0000000c,0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,
+    0x0004002b,0x0000000a,0x00000010,0x00000000,0x00040020,0x00000011,0x00000007,0x00000006,
+    0x00030016,0x00000014,0x00000020,0x00040017,0x00000015,0x00000014,0x00000004,0x00040017,
+    0x00000016,0x00000006,0x00000002,0x0005001e,0x00000017,0x00000015,0x00000016,0x00000016,
+    0x00040020,0x00000018,0x00000009,0x00000017,0x0004003b,0x00000018,0x00000019,0x00000009,
+    0x0004002b,0x00000006,0x0000001a,0x00000002,0x00040020,0x0000001b,0x00000009,0x00000006,
+    0x00020014,0x0000001e,0x00090019,0x00000022,0x00000014,0x00000000,0x00000000,0x00000000,
+    0x00000000,0x00000002,0x00000000,0x00040020,0x00000023,0x00000000,0x00000022,0x0004003b,
+    0x00000023,0x00000024,0x00000000,0x0004002b,0x00000006,0x00000026,0x00000001,0x0004002b,
+    0x00000006,0x0000002c,0x00000000,0x00040020,0x0000002d,0x00000009,0x00000015,0x0004002b,
+    0x0000000a,0x00000030,0x00000040,0x0004002b,0x0000000a,0x00000031,0x00000001,0x0006002c,
+    0x0000000b,0x00000032,0x00000030,0x00000031,0x00000031,0x00050036,0x00000002,0x00000004,
+    0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,
+    0x0004003d,0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,
+    0x0003003e,0x00000009,0x0000000f,0x00050041,0x00000011,0x00000012,0x00000009,0x00000010,
+    0x0004003d,0x00000006,0x00000013,0x00000012,0x00060041,0x0000001b,0x0000001c,0x00000019,
+    0x0000001a,0x00000010,0x0004003d,0x00000006,0x0000001d,0x0000001c,0x000500b1,0x0000001e,
+    0x0000001f,0x00000013,0x0000001d,0x000300f7,0x00000021,0x00000000,0x000400fa,0x0000001f,
+    0x00000020,0x00000021,0x000200f8,0x00000020,0x0004003d,0x00000022,0x00000025,0x00000024,
+    0x00060041,0x0000001b,0x00000027,0x00000019,0x00000026,0x00000010,0x0004003d,0x00000006,
+    0x00000028,0x00000027,0x00050041,0x00000011,0x00000029,0x00000009,0x00000010,0x0004003d,
+    0x00000006,0x0000002a,0x00000029,0x00050080,0x00000006,0x0000002b,0x00000028,0x0000002a,
+    0x00050041,0x0000002d,0x0000002e,0x00000019,0x0000002c,0x0004003d,0x00000015,0x0000002f,
+    0x0000002e,0x00040063,0x00000025,0x0000002b,0x0000002f,0x000200f9,0x00000021,0x000200f8,
+    0x00000021,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_1d_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 64) in;
+
+layout(binding = 0)
+writeonly uniform uimage1D dst;
+
+layout(push_constant)
+uniform u_info_t {
+  uvec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+  
+  if (thread_id.x < u_info.dst_extent.x)
+    imageStore(dst, u_info.dst_offset.x + thread_id.x, u_info.clear_value);
+}
+#endif
+    0x07230203,0x00010000,0x00080007,0x00000032,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x0000002c,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+    0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+    0x00000000,0x0000000d,0x00060010,0x00000004,0x00000011,0x00000040,0x00000001,0x00000001,
+    0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+    0x00000009,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,
+    0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000016,0x6e695f75,0x745f6f66,
+    0x00000000,0x00060006,0x00000016,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+    0x00000016,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000016,0x00000002,
+    0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000018,0x6e695f75,0x00006f66,0x00030005,
+    0x00000023,0x00747364,0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000016,
+    0x00000000,0x00000023,0x00000000,0x00050048,0x00000016,0x00000001,0x00000023,0x00000010,
+    0x00050048,0x00000016,0x00000002,0x00000023,0x00000018,0x00030047,0x00000016,0x00000002,
+    0x00040047,0x00000023,0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,
+    0x00030047,0x00000023,0x00000019,0x00040047,0x00000031,0x0000000b,0x00000019,0x00020013,
+    0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+    0x00040017,0x00000007,0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,
+    0x00040015,0x0000000a,0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,
+    0x00040020,0x0000000c,0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,
+    0x0004002b,0x0000000a,0x00000010,0x00000000,0x00040020,0x00000011,0x00000007,0x00000006,
+    0x00040017,0x00000014,0x0000000a,0x00000004,0x00040017,0x00000015,0x00000006,0x00000002,
+    0x0005001e,0x00000016,0x00000014,0x00000015,0x00000015,0x00040020,0x00000017,0x00000009,
+    0x00000016,0x0004003b,0x00000017,0x00000018,0x00000009,0x0004002b,0x00000006,0x00000019,
+    0x00000002,0x00040020,0x0000001a,0x00000009,0x00000006,0x00020014,0x0000001d,0x00090019,
+    0x00000021,0x0000000a,0x00000000,0x00000000,0x00000000,0x00000000,0x00000002,0x00000000,
+    0x00040020,0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,0x00000023,0x00000000,
+    0x0004002b,0x00000006,0x00000025,0x00000001,0x0004002b,0x00000006,0x0000002b,0x00000000,
+    0x00040020,0x0000002c,0x00000009,0x00000014,0x0004002b,0x0000000a,0x0000002f,0x00000040,
+    0x0004002b,0x0000000a,0x00000030,0x00000001,0x0006002c,0x0000000b,0x00000031,0x0000002f,
+    0x00000030,0x00000030,0x00050036,0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,
+    0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,0x0004003d,0x0000000b,0x0000000e,
+    0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,0x0003003e,0x00000009,0x0000000f,
+    0x00050041,0x00000011,0x00000012,0x00000009,0x00000010,0x0004003d,0x00000006,0x00000013,
+    0x00000012,0x00060041,0x0000001a,0x0000001b,0x00000018,0x00000019,0x00000010,0x0004003d,
+    0x00000006,0x0000001c,0x0000001b,0x000500b1,0x0000001d,0x0000001e,0x00000013,0x0000001c,
+    0x000300f7,0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,
+    0x0000001f,0x0004003d,0x00000021,0x00000024,0x00000023,0x00060041,0x0000001a,0x00000026,
+    0x00000018,0x00000025,0x00000010,0x0004003d,0x00000006,0x00000027,0x00000026,0x00050041,
+    0x00000011,0x00000028,0x00000009,0x00000010,0x0004003d,0x00000006,0x00000029,0x00000028,
+    0x00050080,0x00000006,0x0000002a,0x00000027,0x00000029,0x00050041,0x0000002c,0x0000002d,
+    0x00000018,0x0000002b,0x0004003d,0x00000014,0x0000002e,0x0000002d,0x00040063,0x00000024,
+    0x0000002a,0x0000002e,0x000200f9,0x00000020,0x000200f8,0x00000020,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_2d_array_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8) in;
+
+layout(binding = 0)
+writeonly uniform image2DArray dst;
+
+layout(push_constant)
+uniform u_info_t {
+  vec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+  
+  if (all(lessThan(thread_id.xy, u_info.dst_extent.xy)))
+    imageStore(dst, ivec3(u_info.dst_offset.xy + thread_id.xy, thread_id.z), u_info.clear_value);
+}
+#endif
+    0x07230203,0x00010000,0x00080007,0x0000003a,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+    0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+    0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+    0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+    0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+    0x496e6f69,0x00000044,0x00050005,0x00000015,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+    0x00000015,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000015,0x00000001,
+    0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000015,0x00000002,0x5f747364,0x65747865,
+    0x0000746e,0x00040005,0x00000017,0x6e695f75,0x00006f66,0x00030005,0x00000024,0x00747364,
+    0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000015,0x00000000,0x00000023,
+    0x00000000,0x00050048,0x00000015,0x00000001,0x00000023,0x00000010,0x00050048,0x00000015,
+    0x00000002,0x00000023,0x00000018,0x00030047,0x00000015,0x00000002,0x00040047,0x00000024,
+    0x00000022,0x00000000,0x00040047,0x00000024,0x00000021,0x00000000,0x00030047,0x00000024,
+    0x00000019,0x00040047,0x00000039,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+    0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+    0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+    0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+    0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+    0x00000006,0x00000002,0x00030016,0x00000013,0x00000020,0x00040017,0x00000014,0x00000013,
+    0x00000004,0x0005001e,0x00000015,0x00000014,0x00000010,0x00000010,0x00040020,0x00000016,
+    0x00000009,0x00000015,0x0004003b,0x00000016,0x00000017,0x00000009,0x0004002b,0x00000006,
+    0x00000018,0x00000002,0x00040020,0x00000019,0x00000009,0x00000010,0x00020014,0x0000001c,
+    0x00040017,0x0000001d,0x0000001c,0x00000002,0x00090019,0x00000022,0x00000013,0x00000001,
+    0x00000000,0x00000001,0x00000000,0x00000002,0x00000000,0x00040020,0x00000023,0x00000000,
+    0x00000022,0x0004003b,0x00000023,0x00000024,0x00000000,0x0004002b,0x00000006,0x00000026,
+    0x00000001,0x0004002b,0x0000000a,0x0000002c,0x00000002,0x00040020,0x0000002d,0x00000007,
+    0x00000006,0x0004002b,0x00000006,0x00000033,0x00000000,0x00040020,0x00000034,0x00000009,
+    0x00000014,0x0004002b,0x0000000a,0x00000037,0x00000008,0x0004002b,0x0000000a,0x00000038,
+    0x00000001,0x0006002c,0x0000000b,0x00000039,0x00000037,0x00000037,0x00000038,0x00050036,
+    0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,
+    0x00000009,0x00000007,0x0004003d,0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,
+    0x0000000f,0x0000000e,0x0003003e,0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,
+    0x00000009,0x0007004f,0x00000010,0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,
+    0x00050041,0x00000019,0x0000001a,0x00000017,0x00000018,0x0004003d,0x00000010,0x0000001b,
+    0x0000001a,0x000500b1,0x0000001d,0x0000001e,0x00000012,0x0000001b,0x0004009b,0x0000001c,
+    0x0000001f,0x0000001e,0x000300f7,0x00000021,0x00000000,0x000400fa,0x0000001f,0x00000020,
+    0x00000021,0x000200f8,0x00000020,0x0004003d,0x00000022,0x00000025,0x00000024,0x00050041,
+    0x00000019,0x00000027,0x00000017,0x00000026,0x0004003d,0x00000010,0x00000028,0x00000027,
+    0x0004003d,0x00000007,0x00000029,0x00000009,0x0007004f,0x00000010,0x0000002a,0x00000029,
+    0x00000029,0x00000000,0x00000001,0x00050080,0x00000010,0x0000002b,0x00000028,0x0000002a,
+    0x00050041,0x0000002d,0x0000002e,0x00000009,0x0000002c,0x0004003d,0x00000006,0x0000002f,
+    0x0000002e,0x00050051,0x00000006,0x00000030,0x0000002b,0x00000000,0x00050051,0x00000006,
+    0x00000031,0x0000002b,0x00000001,0x00060050,0x00000007,0x00000032,0x00000030,0x00000031,
+    0x0000002f,0x00050041,0x00000034,0x00000035,0x00000017,0x00000033,0x0004003d,0x00000014,
+    0x00000036,0x00000035,0x00040063,0x00000025,0x00000032,0x00000036,0x000200f9,0x00000021,
+    0x000200f8,0x00000021,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_2d_array_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8) in;
+
+layout(binding = 0)
+writeonly uniform uimage2DArray dst;
+
+layout(push_constant)
+uniform u_info_t {
+  uvec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+  
+  if (all(lessThan(thread_id.xy, u_info.dst_extent.xy)))
+    imageStore(dst, ivec3(u_info.dst_offset.xy + thread_id.xy, thread_id.z), u_info.clear_value);
+}
+#endif
+    0x07230203,0x00010000,0x00080007,0x00000039,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+    0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+    0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+    0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+    0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+    0x496e6f69,0x00000044,0x00050005,0x00000014,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+    0x00000014,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000014,0x00000001,
+    0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000014,0x00000002,0x5f747364,0x65747865,
+    0x0000746e,0x00040005,0x00000016,0x6e695f75,0x00006f66,0x00030005,0x00000023,0x00747364,
+    0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000014,0x00000000,0x00000023,
+    0x00000000,0x00050048,0x00000014,0x00000001,0x00000023,0x00000010,0x00050048,0x00000014,
+    0x00000002,0x00000023,0x00000018,0x00030047,0x00000014,0x00000002,0x00040047,0x00000023,
+    0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,0x00030047,0x00000023,
+    0x00000019,0x00040047,0x00000038,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+    0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+    0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+    0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+    0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+    0x00000006,0x00000002,0x00040017,0x00000013,0x0000000a,0x00000004,0x0005001e,0x00000014,
+    0x00000013,0x00000010,0x00000010,0x00040020,0x00000015,0x00000009,0x00000014,0x0004003b,
+    0x00000015,0x00000016,0x00000009,0x0004002b,0x00000006,0x00000017,0x00000002,0x00040020,
+    0x00000018,0x00000009,0x00000010,0x00020014,0x0000001b,0x00040017,0x0000001c,0x0000001b,
+    0x00000002,0x00090019,0x00000021,0x0000000a,0x00000001,0x00000000,0x00000001,0x00000000,
+    0x00000002,0x00000000,0x00040020,0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,
+    0x00000023,0x00000000,0x0004002b,0x00000006,0x00000025,0x00000001,0x0004002b,0x0000000a,
+    0x0000002b,0x00000002,0x00040020,0x0000002c,0x00000007,0x00000006,0x0004002b,0x00000006,
+    0x00000032,0x00000000,0x00040020,0x00000033,0x00000009,0x00000013,0x0004002b,0x0000000a,
+    0x00000036,0x00000008,0x0004002b,0x0000000a,0x00000037,0x00000001,0x0006002c,0x0000000b,
+    0x00000038,0x00000036,0x00000036,0x00000037,0x00050036,0x00000002,0x00000004,0x00000000,
+    0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,0x0004003d,
+    0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,0x0003003e,
+    0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,0x00000009,0x0007004f,0x00000010,
+    0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,0x00050041,0x00000018,0x00000019,
+    0x00000016,0x00000017,0x0004003d,0x00000010,0x0000001a,0x00000019,0x000500b1,0x0000001c,
+    0x0000001d,0x00000012,0x0000001a,0x0004009b,0x0000001b,0x0000001e,0x0000001d,0x000300f7,
+    0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,0x0000001f,
+    0x0004003d,0x00000021,0x00000024,0x00000023,0x00050041,0x00000018,0x00000026,0x00000016,
+    0x00000025,0x0004003d,0x00000010,0x00000027,0x00000026,0x0004003d,0x00000007,0x00000028,
+    0x00000009,0x0007004f,0x00000010,0x00000029,0x00000028,0x00000028,0x00000000,0x00000001,
+    0x00050080,0x00000010,0x0000002a,0x00000027,0x00000029,0x00050041,0x0000002c,0x0000002d,
+    0x00000009,0x0000002b,0x0004003d,0x00000006,0x0000002e,0x0000002d,0x00050051,0x00000006,
+    0x0000002f,0x0000002a,0x00000000,0x00050051,0x00000006,0x00000030,0x0000002a,0x00000001,
+    0x00060050,0x00000007,0x00000031,0x0000002f,0x00000030,0x0000002e,0x00050041,0x00000033,
+    0x00000034,0x00000016,0x00000032,0x0004003d,0x00000013,0x00000035,0x00000034,0x00040063,
+    0x00000024,0x00000031,0x00000035,0x000200f9,0x00000020,0x000200f8,0x00000020,0x000100fd,
+    0x00010038
+};
+
+const uint32_t cs_clear_uav_image_2d_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8) in;
+
+layout(binding = 0)
+writeonly uniform image2D dst;
+
+layout(push_constant)
+uniform u_info_t {
+  vec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+  
+  if (all(lessThan(thread_id.xy, u_info.dst_extent.xy)))
+    imageStore(dst, u_info.dst_offset.xy + thread_id.xy, u_info.clear_value);
+}
+#endif
+    0x07230203,0x00010000,0x00080007,0x00000033,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+    0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+    0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+    0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+    0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+    0x496e6f69,0x00000044,0x00050005,0x00000015,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+    0x00000015,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000015,0x00000001,
+    0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000015,0x00000002,0x5f747364,0x65747865,
+    0x0000746e,0x00040005,0x00000017,0x6e695f75,0x00006f66,0x00030005,0x00000024,0x00747364,
+    0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000015,0x00000000,0x00000023,
+    0x00000000,0x00050048,0x00000015,0x00000001,0x00000023,0x00000010,0x00050048,0x00000015,
+    0x00000002,0x00000023,0x00000018,0x00030047,0x00000015,0x00000002,0x00040047,0x00000024,
+    0x00000022,0x00000000,0x00040047,0x00000024,0x00000021,0x00000000,0x00030047,0x00000024,
+    0x00000019,0x00040047,0x00000032,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+    0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+    0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+    0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+    0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+    0x00000006,0x00000002,0x00030016,0x00000013,0x00000020,0x00040017,0x00000014,0x00000013,
+    0x00000004,0x0005001e,0x00000015,0x00000014,0x00000010,0x00000010,0x00040020,0x00000016,
+    0x00000009,0x00000015,0x0004003b,0x00000016,0x00000017,0x00000009,0x0004002b,0x00000006,
+    0x00000018,0x00000002,0x00040020,0x00000019,0x00000009,0x00000010,0x00020014,0x0000001c,
+    0x00040017,0x0000001d,0x0000001c,0x00000002,0x00090019,0x00000022,0x00000013,0x00000001,
+    0x00000000,0x00000000,0x00000000,0x00000002,0x00000000,0x00040020,0x00000023,0x00000000,
+    0x00000022,0x0004003b,0x00000023,0x00000024,0x00000000,0x0004002b,0x00000006,0x00000026,
+    0x00000001,0x0004002b,0x00000006,0x0000002c,0x00000000,0x00040020,0x0000002d,0x00000009,
+    0x00000014,0x0004002b,0x0000000a,0x00000030,0x00000008,0x0004002b,0x0000000a,0x00000031,
+    0x00000001,0x0006002c,0x0000000b,0x00000032,0x00000030,0x00000030,0x00000031,0x00050036,
+    0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,
+    0x00000009,0x00000007,0x0004003d,0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,
+    0x0000000f,0x0000000e,0x0003003e,0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,
+    0x00000009,0x0007004f,0x00000010,0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,
+    0x00050041,0x00000019,0x0000001a,0x00000017,0x00000018,0x0004003d,0x00000010,0x0000001b,
+    0x0000001a,0x000500b1,0x0000001d,0x0000001e,0x00000012,0x0000001b,0x0004009b,0x0000001c,
+    0x0000001f,0x0000001e,0x000300f7,0x00000021,0x00000000,0x000400fa,0x0000001f,0x00000020,
+    0x00000021,0x000200f8,0x00000020,0x0004003d,0x00000022,0x00000025,0x00000024,0x00050041,
+    0x00000019,0x00000027,0x00000017,0x00000026,0x0004003d,0x00000010,0x00000028,0x00000027,
+    0x0004003d,0x00000007,0x00000029,0x00000009,0x0007004f,0x00000010,0x0000002a,0x00000029,
+    0x00000029,0x00000000,0x00000001,0x00050080,0x00000010,0x0000002b,0x00000028,0x0000002a,
+    0x00050041,0x0000002d,0x0000002e,0x00000017,0x0000002c,0x0004003d,0x00000014,0x0000002f,
+    0x0000002e,0x00040063,0x00000025,0x0000002b,0x0000002f,0x000200f9,0x00000021,0x000200f8,
+    0x00000021,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_2d_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8) in;
+
+layout(binding = 0)
+writeonly uniform uimage2D dst;
+
+layout(push_constant)
+uniform u_info_t {
+  uvec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+  
+  if (all(lessThan(thread_id.xy, u_info.dst_extent.xy)))
+    imageStore(dst, u_info.dst_offset.xy + thread_id.xy, u_info.clear_value);
+}
+#endif
+    0x07230203,0x00010000,0x00080007,0x00000032,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+    0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+    0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+    0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+    0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+    0x496e6f69,0x00000044,0x00050005,0x00000014,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+    0x00000014,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000014,0x00000001,
+    0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000014,0x00000002,0x5f747364,0x65747865,
+    0x0000746e,0x00040005,0x00000016,0x6e695f75,0x00006f66,0x00030005,0x00000023,0x00747364,
+    0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000014,0x00000000,0x00000023,
+    0x00000000,0x00050048,0x00000014,0x00000001,0x00000023,0x00000010,0x00050048,0x00000014,
+    0x00000002,0x00000023,0x00000018,0x00030047,0x00000014,0x00000002,0x00040047,0x00000023,
+    0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,0x00030047,0x00000023,
+    0x00000019,0x00040047,0x00000031,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+    0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+    0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+    0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+    0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+    0x00000006,0x00000002,0x00040017,0x00000013,0x0000000a,0x00000004,0x0005001e,0x00000014,
+    0x00000013,0x00000010,0x00000010,0x00040020,0x00000015,0x00000009,0x00000014,0x0004003b,
+    0x00000015,0x00000016,0x00000009,0x0004002b,0x00000006,0x00000017,0x00000002,0x00040020,
+    0x00000018,0x00000009,0x00000010,0x00020014,0x0000001b,0x00040017,0x0000001c,0x0000001b,
+    0x00000002,0x00090019,0x00000021,0x0000000a,0x00000001,0x00000000,0x00000000,0x00000000,
+    0x00000002,0x00000000,0x00040020,0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,
+    0x00000023,0x00000000,0x0004002b,0x00000006,0x00000025,0x00000001,0x0004002b,0x00000006,
+    0x0000002b,0x00000000,0x00040020,0x0000002c,0x00000009,0x00000013,0x0004002b,0x0000000a,
+    0x0000002f,0x00000008,0x0004002b,0x0000000a,0x00000030,0x00000001,0x0006002c,0x0000000b,
+    0x00000031,0x0000002f,0x0000002f,0x00000030,0x00050036,0x00000002,0x00000004,0x00000000,
+    0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,0x0004003d,
+    0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,0x0003003e,
+    0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,0x00000009,0x0007004f,0x00000010,
+    0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,0x00050041,0x00000018,0x00000019,
+    0x00000016,0x00000017,0x0004003d,0x00000010,0x0000001a,0x00000019,0x000500b1,0x0000001c,
+    0x0000001d,0x00000012,0x0000001a,0x0004009b,0x0000001b,0x0000001e,0x0000001d,0x000300f7,
+    0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,0x0000001f,
+    0x0004003d,0x00000021,0x00000024,0x00000023,0x00050041,0x00000018,0x00000026,0x00000016,
+    0x00000025,0x0004003d,0x00000010,0x00000027,0x00000026,0x0004003d,0x00000007,0x00000028,
+    0x00000009,0x0007004f,0x00000010,0x00000029,0x00000028,0x00000028,0x00000000,0x00000001,
+    0x00050080,0x00000010,0x0000002a,0x00000027,0x00000029,0x00050041,0x0000002c,0x0000002d,
+    0x00000016,0x0000002b,0x0004003d,0x00000013,0x0000002e,0x0000002d,0x00040063,0x00000024,
+    0x0000002a,0x0000002e,0x000200f9,0x00000020,0x000200f8,0x00000020,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_3d_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in;
+
+layout(binding = 0)
+writeonly uniform image3D dst;
+
+layout(push_constant)
+uniform u_info_t {
+  vec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+  
+  if (all(lessThan(thread_id.xy, u_info.dst_extent)))
+    imageStore(dst, ivec3(u_info.dst_offset.xy, 0) + thread_id.xyz, u_info.clear_value);
+}
+
+#endif
+    0x07230203,0x00010000,0x00080007,0x00000035,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+    0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+    0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+    0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+    0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+    0x496e6f69,0x00000044,0x00050005,0x00000015,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+    0x00000015,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000015,0x00000001,
+    0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000015,0x00000002,0x5f747364,0x65747865,
+    0x0000746e,0x00040005,0x00000017,0x6e695f75,0x00006f66,0x00030005,0x00000024,0x00747364,
+    0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000015,0x00000000,0x00000023,
+    0x00000000,0x00050048,0x00000015,0x00000001,0x00000023,0x00000010,0x00050048,0x00000015,
+    0x00000002,0x00000023,0x00000018,0x00030047,0x00000015,0x00000002,0x00040047,0x00000024,
+    0x00000022,0x00000000,0x00040047,0x00000024,0x00000021,0x00000000,0x00030047,0x00000024,
+    0x00000019,0x00040047,0x00000034,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+    0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+    0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+    0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+    0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+    0x00000006,0x00000002,0x00030016,0x00000013,0x00000020,0x00040017,0x00000014,0x00000013,
+    0x00000004,0x0005001e,0x00000015,0x00000014,0x00000010,0x00000010,0x00040020,0x00000016,
+    0x00000009,0x00000015,0x0004003b,0x00000016,0x00000017,0x00000009,0x0004002b,0x00000006,
+    0x00000018,0x00000002,0x00040020,0x00000019,0x00000009,0x00000010,0x00020014,0x0000001c,
+    0x00040017,0x0000001d,0x0000001c,0x00000002,0x00090019,0x00000022,0x00000013,0x00000002,
+    0x00000000,0x00000000,0x00000000,0x00000002,0x00000000,0x00040020,0x00000023,0x00000000,
+    0x00000022,0x0004003b,0x00000023,0x00000024,0x00000000,0x0004002b,0x00000006,0x00000026,
+    0x00000001,0x0004002b,0x00000006,0x00000029,0x00000000,0x00040020,0x0000002f,0x00000009,
+    0x00000014,0x0004002b,0x0000000a,0x00000032,0x00000008,0x0004002b,0x0000000a,0x00000033,
+    0x00000001,0x0006002c,0x0000000b,0x00000034,0x00000032,0x00000032,0x00000033,0x00050036,
+    0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,
+    0x00000009,0x00000007,0x0004003d,0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,
+    0x0000000f,0x0000000e,0x0003003e,0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,
+    0x00000009,0x0007004f,0x00000010,0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,
+    0x00050041,0x00000019,0x0000001a,0x00000017,0x00000018,0x0004003d,0x00000010,0x0000001b,
+    0x0000001a,0x000500b1,0x0000001d,0x0000001e,0x00000012,0x0000001b,0x0004009b,0x0000001c,
+    0x0000001f,0x0000001e,0x000300f7,0x00000021,0x00000000,0x000400fa,0x0000001f,0x00000020,
+    0x00000021,0x000200f8,0x00000020,0x0004003d,0x00000022,0x00000025,0x00000024,0x00050041,
+    0x00000019,0x00000027,0x00000017,0x00000026,0x0004003d,0x00000010,0x00000028,0x00000027,
+    0x00050051,0x00000006,0x0000002a,0x00000028,0x00000000,0x00050051,0x00000006,0x0000002b,
+    0x00000028,0x00000001,0x00060050,0x00000007,0x0000002c,0x0000002a,0x0000002b,0x00000029,
+    0x0004003d,0x00000007,0x0000002d,0x00000009,0x00050080,0x00000007,0x0000002e,0x0000002c,
+    0x0000002d,0x00050041,0x0000002f,0x00000030,0x00000017,0x00000029,0x0004003d,0x00000014,
+    0x00000031,0x00000030,0x00040063,0x00000025,0x0000002e,0x00000031,0x000200f9,0x00000021,
+    0x000200f8,0x00000021,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_3d_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in;
+
+layout(binding = 0)
+writeonly uniform uimage3D dst;
+
+layout(push_constant)
+uniform u_info_t {
+  uvec4 clear_value;
+  ivec2 dst_offset;
+  ivec2 dst_extent;
+} u_info;
+
+void main() {
+  ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+  
+  if (all(lessThan(thread_id.xy, u_info.dst_extent)))
+    imageStore(dst, ivec3(u_info.dst_offset.xy, 0) + thread_id.xyz, u_info.clear_value);
+}
+
+#endif
+    0x07230203,0x00010000,0x00080007,0x00000034,0x00000000,0x00020011,0x00000001,0x00020011,
+    0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+    0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+    0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+    0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+    0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+    0x496e6f69,0x00000044,0x00050005,0x00000014,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+    0x00000014,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000014,0x00000001,
+    0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000014,0x00000002,0x5f747364,0x65747865,
+    0x0000746e,0x00040005,0x00000016,0x6e695f75,0x00006f66,0x00030005,0x00000023,0x00747364,
+    0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000014,0x00000000,0x00000023,
+    0x00000000,0x00050048,0x00000014,0x00000001,0x00000023,0x00000010,0x00050048,0x00000014,
+    0x00000002,0x00000023,0x00000018,0x00030047,0x00000014,0x00000002,0x00040047,0x00000023,
+    0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,0x00030047,0x00000023,
+    0x00000019,0x00040047,0x00000033,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+    0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+    0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+    0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+    0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+    0x00000006,0x00000002,0x00040017,0x00000013,0x0000000a,0x00000004,0x0005001e,0x00000014,
+    0x00000013,0x00000010,0x00000010,0x00040020,0x00000015,0x00000009,0x00000014,0x0004003b,
+    0x00000015,0x00000016,0x00000009,0x0004002b,0x00000006,0x00000017,0x00000002,0x00040020,
+    0x00000018,0x00000009,0x00000010,0x00020014,0x0000001b,0x00040017,0x0000001c,0x0000001b,
+    0x00000002,0x00090019,0x00000021,0x0000000a,0x00000002,0x00000000,0x00000000,0x00000000,
+    0x00000002,0x00000000,0x00040020,0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,
+    0x00000023,0x00000000,0x0004002b,0x00000006,0x00000025,0x00000001,0x0004002b,0x00000006,
+    0x00000028,0x00000000,0x00040020,0x0000002e,0x00000009,0x00000013,0x0004002b,0x0000000a,
+    0x00000031,0x00000008,0x0004002b,0x0000000a,0x00000032,0x00000001,0x0006002c,0x0000000b,
+    0x00000033,0x00000031,0x00000031,0x00000032,0x00050036,0x00000002,0x00000004,0x00000000,
+    0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,0x0004003d,
+    0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,0x0003003e,
+    0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,0x00000009,0x0007004f,0x00000010,
+    0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,0x00050041,0x00000018,0x00000019,
+    0x00000016,0x00000017,0x0004003d,0x00000010,0x0000001a,0x00000019,0x000500b1,0x0000001c,
+    0x0000001d,0x00000012,0x0000001a,0x0004009b,0x0000001b,0x0000001e,0x0000001d,0x000300f7,
+    0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,0x0000001f,
+    0x0004003d,0x00000021,0x00000024,0x00000023,0x00050041,0x00000018,0x00000026,0x00000016,
+    0x00000025,0x0004003d,0x00000010,0x00000027,0x00000026,0x00050051,0x00000006,0x00000029,
+    0x00000027,0x00000000,0x00050051,0x00000006,0x0000002a,0x00000027,0x00000001,0x00060050,
+    0x00000007,0x0000002b,0x00000029,0x0000002a,0x00000028,0x0004003d,0x00000007,0x0000002c,
+    0x00000009,0x00050080,0x00000007,0x0000002d,0x0000002b,0x0000002c,0x00050041,0x0000002e,
+    0x0000002f,0x00000016,0x00000028,0x0004003d,0x00000013,0x00000030,0x0000002f,0x00040063,
+    0x00000024,0x0000002d,0x00000030,0x000200f9,0x00000020,0x000200f8,0x00000020,0x000100fd,
+    0x00010038
+};
+
+#endif  /* __VKD3D_SPV_SHADERS_H */

[PATCH vkd3d 06/13] vkd3d: Implement ClearUnorderedAccessViewFloat.

[Philip Rebohle]

Signed-off-by: Philip Rebohle 
---
 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;

[PATCH vkd3d 07/13] vkd3d: Add helper to find bit-compatible UINT format for a DXGI format.

[Philip Rebohle]

Needed to imlpement ClearUnorderedAccessViewUint for non-UINT formats.

Signed-off-by: Philip Rebohle 
---
 libs/vkd3d/utils.c         | 33 +++++++++++++++++++++++++++++++++
 libs/vkd3d/vkd3d_private.h |  2 ++
 2 files changed, 35 insertions(+)

diff --git a/libs/vkd3d/utils.c b/libs/vkd3d/utils.c
index 1fc0e9d..fb608d2 100644
--- a/libs/vkd3d/utils.c
+++ b/libs/vkd3d/utils.c
@@ -451,6 +451,39 @@ const struct vkd3d_format *vkd3d_get_format(const struct d3d12_device *device,
     return NULL;
 }

+const struct vkd3d_format *vkd3d_find_uint_format(const struct d3d12_device *device,
+        DXGI_FORMAT dxgi_format)
+{
+    DXGI_FORMAT typeless_format = DXGI_FORMAT_UNKNOWN;
+    const struct vkd3d_format *vkd3d_format;
+    unsigned int i;
+
+    for (i = 0; i < ARRAY_SIZE(vkd3d_format_compatibility_info); i++)
+    {
+        if (vkd3d_format_compatibility_info[i].format == dxgi_format)
+        {
+            typeless_format = vkd3d_format_compatibility_info[i].typeless_format;
+            break;
+        }
+    }
+
+    if (!typeless_format)
+        return NULL;
+
+    for (i = 0; i < ARRAY_SIZE(vkd3d_format_compatibility_info); i++)
+    {
+        if (vkd3d_format_compatibility_info[i].typeless_format == typeless_format)
+        {
+            vkd3d_format = vkd3d_get_format(device, vkd3d_format_compatibility_info[i].format, false);
+
+            if (vkd3d_format->type == VKD3D_FORMAT_TYPE_UINT)
+                return vkd3d_format;
+        }
+    }
+
+    return NULL;
+}
+
 void vkd3d_format_copy_data(const struct vkd3d_format *format, const uint8_t *src,
         unsigned int src_row_pitch, unsigned int src_slice_pitch, uint8_t *dst, unsigned int dst_row_pitch,
         unsigned int dst_slice_pitch, unsigned int w, unsigned int h, unsigned int d)
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index d8b4f2f..6f3ac09 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -1254,6 +1254,8 @@ void vkd3d_format_copy_data(const struct vkd3d_format *format, const uint8_t *sr

const struct vkd3d_format *vkd3d_get_format(const struct d3d12_device *device,
         DXGI_FORMAT dxgi_format, bool depth_stencil) DECLSPEC_HIDDEN;
+const struct vkd3d_format *vkd3d_find_uint_format(const struct d3d12_device *device,
+        DXGI_FORMAT dxgi_format) DECLSPEC_HIDDEN;

HRESULT vkd3d_init_format_info(struct d3d12_device *device) DECLSPEC_HIDDEN;
 void vkd3d_cleanup_format_info(struct d3d12_device *device) DECLSPEC_HIDDEN;

[PATCH vkd3d 08/13] vkd3d: Set MUTABLE_FORMAT_BIT for UAV images with non-UINT format.

[Philip Rebohle]

Needed to support ClearUnorderedAccessViewUint for all formats.

Signed-off-by: Philip Rebohle 
---
 libs/vkd3d/resource.c | 11 +++++++----
 1 file changed, 7 insertions(+), 4 deletions(-)

diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index 8615464..a1c7201 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -864,13 +864,16 @@ static HRESULT vkd3d_create_image(struct d3d12_device *device,
     image_info.sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO;
     image_info.pNext = NULL;
     image_info.flags = 0;
-    if (!(desc->Flags & D3D12_RESOURCE_FLAG_ALLOW_DEPTH_STENCIL) && format->type == VKD3D_FORMAT_TYPE_TYPELESS)
+    if (desc->Flags & D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS)
+    {
+        if (format->type != VKD3D_FORMAT_TYPE_UINT)
+            image_info.flags |= VK_IMAGE_CREATE_MUTABLE_FORMAT_BIT;
+    }
+    else if (!(desc->Flags & D3D12_RESOURCE_FLAG_ALLOW_DEPTH_STENCIL) && format->type == VKD3D_FORMAT_TYPE_TYPELESS)
     {
         image_info.flags |= VK_IMAGE_CREATE_MUTABLE_FORMAT_BIT;

-        /* Format compatibility rules are more relaxed for UAVs. */
-        if (!(desc->Flags & D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS)
-                && (compat_list = vkd3d_get_format_compatibility_list(device, desc->Format)))
+        if ((compat_list = vkd3d_get_format_compatibility_list(device, desc->Format)))
         {
             format_list.sType = VK_STRUCTURE_TYPE_IMAGE_FORMAT_LIST_CREATE_INFO_KHR;
             format_list.pNext = NULL;

[PATCH vkd3d 09/13] vkd3d: Make create_buffer_view and create_texture_view public.

[Philip Rebohle]

Needed for ClearUnorderedAccessViewUint.

Signed-off-by: Philip Rebohle 
---
 libs/vkd3d/resource.c      | 16 ++--------------
 libs/vkd3d/vkd3d_private.h | 19 +++++++++++++++++++
 2 files changed, 21 insertions(+), 14 deletions(-)

diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index a1c7201..361c1fe 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -2149,7 +2149,7 @@ static bool vkd3d_create_vk_buffer_view(struct d3d12_device *device,
     return vr == VK_SUCCESS;
 }

-static bool vkd3d_create_buffer_view(struct d3d12_device *device,
+bool vkd3d_create_buffer_view(struct d3d12_device *device,
         VkBuffer vk_buffer, const struct vkd3d_format *format,
         VkDeviceSize offset, VkDeviceSize size, struct vkd3d_view **view)
 {
@@ -2349,18 +2349,6 @@ static void vk_component_mapping_compose(VkComponentMapping *dst, const VkCompon
     dst->a = swizzle_vk_component(&a, a.a, b->a);
 }

-struct vkd3d_texture_view_desc
-{
-    VkImageViewType view_type;
-    const struct vkd3d_format *format;
-    unsigned int miplevel_idx;
-    unsigned int miplevel_count;
-    unsigned int layer_idx;
-    unsigned int layer_count;
-    VkComponentMapping components;
-    bool allowed_swizzle;
-};
-
 static bool init_default_texture_view_desc(struct vkd3d_texture_view_desc *desc,
         struct d3d12_resource *resource, DXGI_FORMAT view_format)
 {
@@ -2408,7 +2396,7 @@ static bool init_default_texture_view_desc(struct vkd3d_texture_view_desc *desc,
     return true;
 }

-static bool vkd3d_create_texture_view(struct d3d12_device *device,
+bool vkd3d_create_texture_view(struct d3d12_device *device,
         VkImage vk_image, const struct vkd3d_texture_view_desc *desc,
         struct vkd3d_view **view)
 {
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index 6f3ac09..3ab443b 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -494,6 +494,25 @@ struct vkd3d_view
     } info;
 };

+struct vkd3d_texture_view_desc
+{
+    VkImageViewType view_type;
+    const struct vkd3d_format *format;
+    unsigned int miplevel_idx;
+    unsigned int miplevel_count;
+    unsigned int layer_idx;
+    unsigned int layer_count;
+    VkComponentMapping components;
+    bool allowed_swizzle;
+};
+
+bool vkd3d_create_buffer_view(struct d3d12_device *device,
+        VkBuffer vk_buffer, const struct vkd3d_format *format,
+        VkDeviceSize offset, VkDeviceSize size, struct vkd3d_view **view) DECLSPEC_HIDDEN;
+bool vkd3d_create_texture_view(struct d3d12_device *device,
+        VkImage vk_image, const struct vkd3d_texture_view_desc *desc,
+        struct vkd3d_view **view) DECLSPEC_HIDDEN;
+
 void vkd3d_view_decref(struct vkd3d_view *view, struct d3d12_device *device) DECLSPEC_HIDDEN;
 void vkd3d_view_incref(struct vkd3d_view *view) DECLSPEC_HIDDEN;

[PATCH vkd3d 10/13] vkd3d: Re-implement ClearUnorderedAccessViewUint.

[Philip Rebohle]

Addresses the following limitations of the previous implementation:
- Only R32_{UINT,TYPELESS} were supported for buffers.
- Clearing an image UAV did not behave correctly for images with non-UINT formats.
- Due to the use of transfer operations, extra memory barriers were needed.

If necessary, this will create a temporary view with a bit-compatible
UINT format for the resource in order to perform a bit-exact clear.

Signed-off-by: Philip Rebohle 
---
 libs/vkd3d/command.c | 113 ++++++++++++++++---------------------------
 1 file changed, 41 insertions(+), 72 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index 3648ea4..e5fe0ad 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -4930,96 +4930,65 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewUint(ID
         const UINT values[4], UINT rect_count, const D3D12_RECT *rects)
 {
     struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
-    const struct vkd3d_vk_device_procs *vk_procs;
-    const struct vkd3d_vulkan_info *vk_info;
-    const struct d3d12_desc *cpu_descriptor;
+    struct vkd3d_view *base_view, *uint_view;
+    struct vkd3d_texture_view_desc view_desc;
+    const struct vkd3d_format *uint_format;
     struct d3d12_resource *resource_impl;
-    VkBufferMemoryBarrier buffer_barrier;
-    VkImageMemoryBarrier image_barrier;
-    VkPipelineStageFlags stage_mask;
-    VkImageSubresourceRange range;
     VkClearColorValue color;

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);

-    vk_procs = &list->device->vk_procs;
-    vk_info = &list->device->vk_info;
+    memcpy(color.uint32, values, sizeof(color.uint32));

resource_impl = unsafe_impl_from_ID3D12Resource(resource);

-    d3d12_command_list_track_resource_usage(list, resource_impl);
+    base_view = d3d12_desc_from_cpu_handle(cpu_handle)->u.view;
+    uint_view = NULL;

-    if (rect_count)
+    if (base_view->format->type != VKD3D_FORMAT_TYPE_UINT)
     {
-        FIXME("Clear rects not supported.\n");
-        return;
-    }
-
-    d3d12_command_list_end_current_render_pass(list);
+        uint_format = vkd3d_find_uint_format(list->device, base_view->format->dxgi_format);

-    cpu_descriptor = d3d12_desc_from_cpu_handle(cpu_handle);
-
-    if (d3d12_resource_is_buffer(resource_impl))
-    {
-        if (cpu_descriptor->u.view->format->vk_format != VK_FORMAT_R32_UINT)
+        if (!uint_format)
         {
-            FIXME("Not supported for UAV descriptor %p.\n", cpu_descriptor);
+            ERR("Unhandled format %d.\n", base_view->format->dxgi_format);
             return;
         }

-        VK_CALL(vkCmdFillBuffer(list->vk_command_buffer, resource_impl->u.vk_buffer,
-                cpu_descriptor->u.view->info.buffer.offset, cpu_descriptor->u.view->info.buffer.size, values[0]));
-
-        buffer_barrier.sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER;
-        buffer_barrier.pNext = NULL;
-        buffer_barrier.srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT;
-        buffer_barrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
-        buffer_barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
-        buffer_barrier.buffer = resource_impl->u.vk_buffer;
-        buffer_barrier.offset = cpu_descriptor->u.view->info.buffer.offset;
-        buffer_barrier.size = cpu_descriptor->u.view->info.buffer.size;
-
-        vk_barrier_parameters_from_d3d12_resource_state(D3D12_RESOURCE_STATE_UNORDERED_ACCESS, 0,
-                resource_impl, list->vk_queue_flags, vk_info, &buffer_barrier.dstAccessMask, &stage_mask, NULL);
-
-        VK_CALL(vkCmdPipelineBarrier(list->vk_command_buffer,
-                VK_PIPELINE_STAGE_TRANSFER_BIT, stage_mask, 0,
-                0, NULL, 1, &buffer_barrier, 0, NULL));
+        if (d3d12_resource_is_texture(resource_impl))
+        {
+            memset(&view_desc, 0, sizeof(view_desc));
+            view_desc.view_type = base_view->info.texture.vk_view_type;
+            view_desc.format = uint_format;
+            view_desc.miplevel_idx = base_view->info.texture.miplevel_idx;
+            view_desc.miplevel_count = 1;
+            view_desc.layer_idx = base_view->info.texture.layer_idx;
+            view_desc.layer_count = base_view->info.texture.layer_count;
+            view_desc.allowed_swizzle = false;
+
+            if (!vkd3d_create_texture_view(list->device, resource_impl->u.vk_image, &view_desc, &uint_view))
+            {
+                ERR("Failed to create image view.\n");
+                return;
+            }
+        }
+        else
+        {
+            if (!vkd3d_create_buffer_view(list->device, resource_impl->u.vk_buffer, uint_format,
+                    base_view->info.buffer.offset, base_view->info.buffer.size, &uint_view))
+            {
+                ERR("Failed to create buffer view.\n");
+                return;
+            }
+        }
     }
-    else
-    {
-        color.uint32[0] = values[0];
-        color.uint32[1] = values[1];
-        color.uint32[2] = values[2];
-        color.uint32[3] = values[3];
-
-        range.aspectMask = cpu_descriptor->u.view->format->vk_aspect_mask;
-        range.baseMipLevel = cpu_descriptor->u.view->info.texture.miplevel_idx;
-        range.levelCount = 1;
-        range.baseArrayLayer = cpu_descriptor->u.view->info.texture.layer_idx;
-        range.layerCount = cpu_descriptor->u.view->info.texture.layer_count;
-
-        VK_CALL(vkCmdClearColorImage(list->vk_command_buffer,
-                resource_impl->u.vk_image, VK_IMAGE_LAYOUT_GENERAL, &color, 1, &range));
-
-        image_barrier.sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER;
-        image_barrier.pNext = NULL;
-        image_barrier.srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT;
-        image_barrier.oldLayout = VK_IMAGE_LAYOUT_GENERAL;
-        image_barrier.newLayout = VK_IMAGE_LAYOUT_GENERAL;
-        image_barrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
-        image_barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
-        image_barrier.image = resource_impl->u.vk_image;
-        image_barrier.subresourceRange = range;
-
-        vk_barrier_parameters_from_d3d12_resource_state(D3D12_RESOURCE_STATE_UNORDERED_ACCESS, 0,
-                resource_impl, list->vk_queue_flags, vk_info, &image_barrier.dstAccessMask, &stage_mask, NULL);

-        VK_CALL(vkCmdPipelineBarrier(list->vk_command_buffer,
-                VK_PIPELINE_STAGE_TRANSFER_BIT, stage_mask, 0,
-                0, NULL, 0, NULL, 1, &image_barrier));
-    }
+    d3d12_command_list_clear_unordered_access_view(list, resource_impl,
+      uint_view ? uint_view : base_view, &color, rect_count, rects);
+
+    if (uint_view)
+        vkd3d_view_decref(uint_view, list->device);
 }

static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewFloat(ID3D12GraphicsCommandList1 *iface,

[PATCH vkd3d 11/13] vkd3d: Handle R11G11B10_FLOAT in ClearUnorderedAccessViewUint.

[Philip Rebohle]

There is no bit-compatible UINT format, so we'll use R32_UINT.

Signed-off-by: Philip Rebohle 
---
 libs/vkd3d/command.c | 19 ++++++++++++++++++-
 1 file changed, 18 insertions(+), 1 deletion(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index e5fe0ad..0745cfc 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -4925,6 +4925,22 @@ static void d3d12_command_list_clear_unordered_access_view(struct d3d12_command_
     }
 }

+static const struct vkd3d_format *vkd3d_fixup_clear_uav_uint_color(struct d3d12_device *device,
+        DXGI_FORMAT dxgi_format, const UINT values[4], VkClearColorValue *color)
+{
+    switch (dxgi_format)
+    {
+        case DXGI_FORMAT_R11G11B10_FLOAT:
+            color->uint32[0] = (values[0] & 0x7FF)
+                    | ((values[1] & 0x7FF) << 11)
+                    | ((values[2] & 0x3FF) << 22);
+            return vkd3d_get_format(device, DXGI_FORMAT_R32_UINT, false);
+
+        default:
+            return NULL;
+    }
+}
+
 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)
@@ -4950,7 +4966,8 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewUint(ID
     {
         uint_format = vkd3d_find_uint_format(list->device, base_view->format->dxgi_format);

-        if (!uint_format)
+        if (!uint_format && !(uint_format = vkd3d_fixup_clear_uav_uint_color(
+                list->device, base_view->format->dxgi_format, values, &color)))
         {
             ERR("Unhandled format %d.\n", base_view->format->dxgi_format);
             return;

[PATCH vkd3d 12/13] tests: Test more view formats for ClearUnorderedAccessView{Uint, Float}.

[Philip Rebohle]

Signed-off-by: Philip Rebohle 
---
 tests/d3d12.c | 107 ++++++++++++++++++++++++++++++++++++++------------
 1 file changed, 82 insertions(+), 25 deletions(-)

diff --git a/tests/d3d12.c b/tests/d3d12.c
index 0f843b4..6cf2427 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -4743,7 +4743,7 @@ static void test_clear_render_target_view(void)
     destroy_test_context(&context);
 }

-static void test_clear_unordered_access_view(void)
+static void test_clear_unordered_access_view_buffer(void)
 {
     D3D12_UNORDERED_ACCESS_VIEW_DESC uav_desc;
     ID3D12DescriptorHeap *cpu_heap, *gpu_heap;
@@ -4767,42 +4767,88 @@ static void test_clear_unordered_access_view(void)
         DXGI_FORMAT format;
         D3D12_BUFFER_UAV buffer_uav;
         unsigned int values[4];
+        unsigned int expected;
+        bool is_float;
+        bool is_todo;
     }
     tests[] =
     {
         {DXGI_FORMAT_R32_UINT, { 0, BUFFER_SIZE / sizeof(uint32_t),      0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
-                {0, 0, 0, 0}},
+                {0, 0, 0, 0}, 0, false, false},
         {DXGI_FORMAT_R32_UINT, {64, BUFFER_SIZE / sizeof(uint32_t) - 64, 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
-                {0, 0, 0, 0}},
+                {0, 0, 0, 0}, 0, false, false},
         {DXGI_FORMAT_R32_UINT, { 0, BUFFER_SIZE / sizeof(uint32_t),      0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
-                {1, 0, 0, 0}},
+                {1, 0, 0, 0}, 1, false, false},
         {DXGI_FORMAT_R32_UINT, {64, BUFFER_SIZE / sizeof(uint32_t) - 64, 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
-                {2, 0, 0, 0}},
+                {2, 0, 0, 0}, 2, false, false},
         {DXGI_FORMAT_R32_UINT, {64, BUFFER_SIZE / sizeof(uint32_t) - 64, 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
-                {3, 0, 0, 0}},
+                {3, 0, 0, 0}, 3, false, false},
         {DXGI_FORMAT_R32_UINT, {64, BUFFER_SIZE / sizeof(uint32_t) - 64, 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
-                {4, 2, 3, 4}},
+                {4, 2, 3, 4}, 4, false, false},
         {DXGI_FORMAT_R32_UINT, { 0, BUFFER_SIZE / sizeof(uint32_t) - 10, 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
-                {5, 0, 0, 0}},
+                {5, 0, 0, 0}, 5, false, false},

{DXGI_FORMAT_R32_TYPELESS, { 0, BUFFER_SIZE / sizeof(uint32_t),      0, 0, D3D12_BUFFER_UAV_FLAG_RAW},
-                {0, 0, 0, 0}},
+                {0, 0, 0, 0}, 0, false, false},
         {DXGI_FORMAT_R32_TYPELESS, {64, BUFFER_SIZE / sizeof(uint32_t) - 64, 0, 0, D3D12_BUFFER_UAV_FLAG_RAW},
-                {0, 0, 0, 0}},
+                {0, 0, 0, 0}, 0, false, false},
         {DXGI_FORMAT_R32_TYPELESS, { 0, BUFFER_SIZE / sizeof(uint32_t),      0, 0, D3D12_BUFFER_UAV_FLAG_RAW},
-                {6, 0, 0, 0}},
+                {6, 0, 0, 0}, 6, false, false},
         {DXGI_FORMAT_R32_TYPELESS, {64, BUFFER_SIZE / sizeof(uint32_t) - 64, 0, 0, D3D12_BUFFER_UAV_FLAG_RAW},
-                {7, 0, 0, 0}},
+                {7, 0, 0, 0}, 7, false, false},
         {DXGI_FORMAT_R32_TYPELESS, {64, BUFFER_SIZE / sizeof(uint32_t) - 64, 0, 0, D3D12_BUFFER_UAV_FLAG_RAW},
-                {8, 0, 0, 0}},
+                {8, 0, 0, 0}, 8, false, false},
         {DXGI_FORMAT_R32_TYPELESS, {64, BUFFER_SIZE / sizeof(uint32_t) - 64, 0, 0, D3D12_BUFFER_UAV_FLAG_RAW},
-                {9, 1, 1, 1}},
+                {9, 1, 1, 1}, 9, false, false},
         {DXGI_FORMAT_R32_TYPELESS, {64, BUFFER_SIZE / sizeof(uint32_t) - 64, 0, 0, D3D12_BUFFER_UAV_FLAG_RAW},
-                {~0u, 0, 0, 0}},
+                {~0u, 0, 0, 0}, ~0u, false, false},
         {DXGI_FORMAT_R32_TYPELESS, { 0, BUFFER_SIZE / sizeof(uint32_t) - 10, 0, 0, D3D12_BUFFER_UAV_FLAG_RAW},
-                {10, 0, 0, 0}},
+                {10, 0, 0, 0}, 10, false, false},
         {DXGI_FORMAT_R32_TYPELESS, { 0, BUFFER_SIZE / sizeof(uint32_t) - 9,  0, 0, D3D12_BUFFER_UAV_FLAG_RAW},
-                {11, 0, 0, 0}},
+                {11, 0, 0, 0}, 11, false, false},
+
+        {DXGI_FORMAT_R32_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0, 0, 0, 0}, 0, false, false},
+        {DXGI_FORMAT_R32_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {1, 0, 0, 0}, 1, false, false},
+        {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, false},
+
+        {DXGI_FORMAT_R16G16_UINT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x1234, 0xABCD, 0, 0}, 0xABCD1234, false, false},
+        {DXGI_FORMAT_R16G16_UINT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x10000, 0, 0, 0}, 0, false, true},
+
+        {DXGI_FORMAT_R16G16_UNORM, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x1234, 0xABCD, 0, 0}, 0xABCD1234, false, false},
+        {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, false},
+        {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, false},
+        {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, false},
+
+        {DXGI_FORMAT_R16G16_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x1234, 0xABCD, 0, 0}, 0xABCD1234, false, false},
+        {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, false},
+
+        {DXGI_FORMAT_R8G8B8A8_UINT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x11, 0x22, 0x33, 0x44}, 0x44332211, false, false},
+        {DXGI_FORMAT_R8G8B8A8_UINT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x100, 0, 0, 0}, 0, false, true},
+
+        {DXGI_FORMAT_R11G11B10_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0, 0, 0, 0}, 0, false, false},
+        {DXGI_FORMAT_R11G11B10_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x7FF, 0x7FF, 0x3FF, 0}, 0xFFFFFFFF, false, false},
+        {DXGI_FORMAT_R11G11B10_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x7FF, 0, 0x3FF, 0}, 0xFFC007FF, false, false},
+        {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, false},
+        {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, false},
+
     };

memset(&desc, 0, sizeof(desc));
@@ -4832,7 +4878,7 @@ static void test_clear_unordered_access_view(void)
                 D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_UNORDERED_ACCESS);

for (j = 0; j < ARRAY_SIZE(clear_value); ++j)
-            clear_value[j] = tests[i].values[j] ? 0 : ~0u;
+            clear_value[j] = tests[i].expected ? 0 : ~0u;

memset(&uav_desc, 0, sizeof(uav_desc));
         uav_desc.Format = DXGI_FORMAT_R32_UINT;
@@ -4858,10 +4904,20 @@ static void test_clear_unordered_access_view(void)

uav_barrier(command_list, buffer);

-        ID3D12GraphicsCommandList_ClearUnorderedAccessViewUint(command_list,
-                get_gpu_descriptor_handle(&context, gpu_heap, 0),
-                get_cpu_descriptor_handle(&context, cpu_heap, 0),
-                buffer, tests[i].values, 0, NULL);
+        if (tests[i].is_float)
+        {
+            ID3D12GraphicsCommandList_ClearUnorderedAccessViewFloat(command_list,
+                    get_gpu_descriptor_handle(&context, gpu_heap, 0),
+                    get_cpu_descriptor_handle(&context, cpu_heap, 0),
+                    buffer, (const float *)tests[i].values, 0, NULL);
+        }
+        else
+        {
+            ID3D12GraphicsCommandList_ClearUnorderedAccessViewUint(command_list,
+                    get_gpu_descriptor_handle(&context, gpu_heap, 0),
+                    get_cpu_descriptor_handle(&context, cpu_heap, 0),
+                    buffer, tests[i].values, 0, NULL);
+        }

set_box(&box, 0, 0, 0, 1, 1, 1);
         transition_resource_state(command_list, buffer,
@@ -4872,9 +4928,10 @@ static void test_clear_unordered_access_view(void)
         check_readback_data_uint(&rb, &box, clear_value[0], 0);
         box.left = uav_desc.Buffer.FirstElement;
         box.right = uav_desc.Buffer.FirstElement + uav_desc.Buffer.NumElements;
-        check_readback_data_uint(&rb, &box, tests[i].values[0], 0);
+        todo_if(tests[i].is_todo)
+        check_readback_data_uint(&rb, &box, tests[i].expected, tests[i].is_float ? 1 : 0);
         box.left = uav_desc.Buffer.FirstElement + uav_desc.Buffer.NumElements;
-        box.right = BUFFER_SIZE / format_size(uav_desc.Format);
+        box.right = BUFFER_SIZE / sizeof(uint32_t);
         check_readback_data_uint(&rb, &box, clear_value[0], 0);
         release_resource_readback(&rb);

@@ -32393,7 +32450,7 @@ START_TEST(d3d12)
     run_test(test_fence_values);
     run_test(test_clear_depth_stencil_view);
     run_test(test_clear_render_target_view);
-    run_test(test_clear_unordered_access_view);
+    run_test(test_clear_unordered_access_view_buffer);
     run_test(test_set_render_targets);
     run_test(test_draw_instanced);
     run_test(test_draw_indexed_instanced);

[PATCH vkd3d 13/13] tests: Add test for ClearUnorderedAccessView{Uint, Float} for images.

[Philip Rebohle]

Signed-off-by: Philip Rebohle 
---
 tests/d3d12.c            | 262 +++++++++++++++++++++++++++++++++++++++
 tests/d3d12_test_utils.h |   5 +
 2 files changed, 267 insertions(+)

diff --git a/tests/d3d12.c b/tests/d3d12.c
index 6cf2427..d1e8b91 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -4947,6 +4947,267 @@ static void test_clear_unordered_access_view_buffer(void)
 #undef BUFFER_SIZE
 }

+static void test_clear_unordered_access_view_image(void)
+{
+    unsigned int expected_color, actual_color;
+    D3D12_UNORDERED_ACCESS_VIEW_DESC uav_desc;
+    ID3D12DescriptorHeap *cpu_heap, *gpu_heap;
+    ID3D12GraphicsCommandList *command_list;
+    unsigned int i, j, d, p, x, y, z, layer;
+    D3D12_HEAP_PROPERTIES heap_properties;
+    unsigned int image_size, image_depth;
+    D3D12_RESOURCE_DESC resource_desc;
+    struct test_context_desc desc;
+    struct test_context context;
+    struct resource_readback rb;
+    ID3D12CommandQueue *queue;
+    bool is_inside, success;
+    ID3D12Resource *texture;
+    ID3D12Device *device;
+    UINT clear_value[4];
+    HRESULT hr;
+
+#define IMAGE_SIZE 16
+    struct {
+        DXGI_FORMAT format;
+        unsigned int image_mips;
+        unsigned int image_layers;
+        unsigned int mip_level;
+        unsigned int first_layer;
+        unsigned int layer_count;
+        unsigned int rect_count;
+        RECT clear_rects[2];
+        unsigned int values[4];
+        unsigned int expected;
+        bool is_float;
+        bool is_todo;
+    }
+    tests[] =
+    {
+        /* test clearing 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, false },
+        { DXGI_FORMAT_R32_FLOAT, 2, 1, 1, 0, 1, 0, {}, {0x3f000000, 0, 0, 0}, 0x3f000000, true, false },
+        /* 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, {}, {0x3f000000, 0, 0, 0}, 0x3f000000, true, false },
+        { DXGI_FORMAT_R32_FLOAT, 1, IMAGE_SIZE, 0, 3, 2,          0, {}, {0x3f000000, 0, 0, 0}, 0x3f000000, true, false },
+        /* test 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, false },
+        { DXGI_FORMAT_R32_FLOAT, 1, 1, 0, 0, 1, 1, {{1, 2, IMAGE_SIZE - 4, IMAGE_SIZE - 2}}, {0x3f000000, 0, 0, 0}, 0x3f000000, true,  false },
+        /* 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, false },
+        { DXGI_FORMAT_R32_FLOAT, 1, 1, 0, 0, 1, 2, {{1, 2, 3, 4}, {5, 6, 7, 8}}, {0x3f000000, 0, 0, 0}, 0x3f000000, true,  false },
+        /* 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, {}, {0x12345, 0, 0, 0}, 0x00002345, false, true  },
+        { DXGI_FORMAT_R16G16_UNORM,     1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4},       0x00020001, false, false },
+        { DXGI_FORMAT_R16G16_FLOAT,     1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4},       0x00020001, false, false },
+        { 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, {}, {0x123, 0, 0, 0},   0x00000023, false, true  },
+        { DXGI_FORMAT_R8G8B8A8_UNORM,   1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4},       0x04030201, false, false },
+        { DXGI_FORMAT_R11G11B10_FLOAT,  1, 1, 0, 0, 1, 0, {}, {1, 2, 3, 4},       0x00c01001, false, false },
+        /* 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, false },
+        { DXGI_FORMAT_R16G16_FLOAT,     1, 1, 0, 0, 1, 0, {}, {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0x3c003800, true, false },
+        { DXGI_FORMAT_R8G8B8A8_UNORM,   1, 1, 0, 0, 1, 0, {}, {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0x0000ff80, true, false },
+        { DXGI_FORMAT_R8G8B8A8_UNORM,   1, 1, 0, 0, 1, 0, {}, {0, 0, 0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */}, 0xff800000, true, false },
+        { DXGI_FORMAT_R11G11B10_FLOAT,  1, 1, 0, 0, 1, 0, {}, {0x3f000000 /* 1.0f */, 0 /* 0.0f */, 0xbf800000 /* -1.0f */, 0x3f000000 /* 1.0f */}, 0x00000380, true, false},
+    };
+
+    struct {
+        D3D12_RESOURCE_DIMENSION resource_dim;
+        D3D12_UAV_DIMENSION view_dim;
+        bool is_layered;
+    }
+    uav_dimensions[] =
+    {
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, D3D12_UAV_DIMENSION_TEXTURE2D,      false },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, D3D12_UAV_DIMENSION_TEXTURE2DARRAY, true  },
+        /* expected behaviour with partial layer coverage is unclear */
+        { D3D12_RESOURCE_DIMENSION_TEXTURE3D, D3D12_UAV_DIMENSION_TEXTURE3D,      false },
+    };
+
+    memset(&desc, 0, sizeof(desc));
+    desc.no_render_target = true;
+    if (!init_test_context(&context, &desc))
+        return;
+    device = context.device;
+    command_list = context.list;
+    queue = context.queue;
+
+    cpu_heap = create_cpu_descriptor_heap(device, D3D12_DESCRIPTOR_HEAP_TYPE_CBV_SRV_UAV, 2);
+    gpu_heap = create_gpu_descriptor_heap(device, D3D12_DESCRIPTOR_HEAP_TYPE_CBV_SRV_UAV, 2);
+
+    memset(&heap_properties, 0, sizeof(heap_properties));
+    heap_properties.Type = D3D12_HEAP_TYPE_DEFAULT;
+
+    for (d = 0; d < ARRAY_SIZE(uav_dimensions); ++d)
+    {
+        for (i = 0; i < ARRAY_SIZE(tests); ++i)
+        {
+            vkd3d_test_set_context("Dim %u, Test %u", d, i);
+
+            if (tests[i].image_layers > 1 && !uav_dimensions[d].is_layered)
+                continue;
+
+            resource_desc.Dimension = uav_dimensions[d].resource_dim;
+            resource_desc.Alignment = D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT;
+            resource_desc.Width = IMAGE_SIZE;
+            resource_desc.Height = uav_dimensions[d].resource_dim == D3D12_RESOURCE_DIMENSION_TEXTURE1D ? 1 : IMAGE_SIZE;
+            resource_desc.DepthOrArraySize = tests[i].image_layers;
+            resource_desc.MipLevels = tests[i].image_mips;
+            resource_desc.Format = tests[i].format;
+            resource_desc.SampleDesc.Count = 1;
+            resource_desc.SampleDesc.Quality = 0;
+            resource_desc.Layout = D3D12_TEXTURE_LAYOUT_UNKNOWN;
+            resource_desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS;
+
+            hr = ID3D12Device_CreateCommittedResource(device,
+                    &heap_properties, D3D12_HEAP_FLAG_NONE, &resource_desc,
+                    D3D12_RESOURCE_STATE_UNORDERED_ACCESS, NULL,
+                    &IID_ID3D12Resource, (void **)&texture);
+            
+            if (FAILED(hr))
+            {
+                skip("Failed to create texture, hr %#x.\n", hr);
+                continue;
+            }
+
+            uav_desc.Format = tests[i].format;
+            uav_desc.ViewDimension = uav_dimensions[d].view_dim;
+
+            for (j = 0; j < 2; j++)
+            {
+                unsigned int first_layer = j ? 0 : tests[i].first_layer;
+                unsigned int layer_count = j ? tests[i].image_layers : tests[i].layer_count;
+
+                switch (uav_desc.ViewDimension)
+                {
+                    case D3D12_UAV_DIMENSION_TEXTURE1D:
+                        uav_desc.Texture1D.MipSlice = tests[i].mip_level;
+                        break;
+
+                    case D3D12_UAV_DIMENSION_TEXTURE1DARRAY:
+                        uav_desc.Texture1DArray.MipSlice = tests[i].mip_level;
+                        uav_desc.Texture1DArray.FirstArraySlice = first_layer;
+                        uav_desc.Texture1DArray.ArraySize = layer_count;
+                        break;
+
+                    case D3D12_UAV_DIMENSION_TEXTURE2D:
+                        uav_desc.Texture2D.MipSlice = tests[i].mip_level;
+                        uav_desc.Texture2D.PlaneSlice = 0;
+                        break;
+
+                    case D3D12_UAV_DIMENSION_TEXTURE2DARRAY:
+                        uav_desc.Texture2DArray.MipSlice = tests[i].mip_level;
+                        uav_desc.Texture2DArray.FirstArraySlice = first_layer;
+                        uav_desc.Texture2DArray.ArraySize = layer_count;
+                        uav_desc.Texture2DArray.PlaneSlice = 0;
+                        break;
+
+                    case D3D12_UAV_DIMENSION_TEXTURE3D:
+                        uav_desc.Texture3D.MipSlice = tests[i].mip_level;
+                        uav_desc.Texture3D.FirstWSlice = first_layer;
+                        uav_desc.Texture3D.WSize = layer_count;
+                        break;
+
+                    default:
+                        continue;
+                }
+
+                ID3D12Device_CreateUnorderedAccessView(device, texture, NULL, &uav_desc, get_cpu_descriptor_handle(&context, cpu_heap, j));
+                ID3D12Device_CreateUnorderedAccessView(device, texture, NULL, &uav_desc, get_cpu_descriptor_handle(&context, gpu_heap, j));
+            }
+
+            for (j = 0; j < 4; j++)
+                clear_value[j] = tests[i].expected ? 0u : ~0u;
+
+            ID3D12GraphicsCommandList_ClearUnorderedAccessViewUint(command_list,
+                    get_gpu_descriptor_handle(&context, gpu_heap, 1),
+                    get_cpu_descriptor_handle(&context, cpu_heap, 1),
+                    texture, clear_value, 0, NULL);
+
+            uav_barrier(command_list, texture);
+
+            if (tests[i].is_float)
+            {
+                ID3D12GraphicsCommandList_ClearUnorderedAccessViewFloat(command_list,
+                        get_gpu_descriptor_handle(&context, gpu_heap, 0),
+                        get_cpu_descriptor_handle(&context, cpu_heap, 0),
+                        texture, (const float *)tests[i].values, tests[i].rect_count, tests[i].clear_rects);
+            }
+            else
+            {
+                ID3D12GraphicsCommandList_ClearUnorderedAccessViewUint(command_list,
+                        get_gpu_descriptor_handle(&context, gpu_heap, 0),
+                        get_cpu_descriptor_handle(&context, cpu_heap, 0),
+                        texture, tests[i].values, tests[i].rect_count, tests[i].clear_rects);
+            }
+
+            transition_resource_state(command_list, texture, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+            image_depth = uav_dimensions[d].resource_dim == D3D12_RESOURCE_DIMENSION_TEXTURE3D
+                    ? max(tests[i].image_layers >> tests[i].mip_level, 1u) : 1;
+            image_size = max(IMAGE_SIZE >> tests[i].mip_level, 1u);
+
+            for (layer = 0; layer < tests[i].image_layers / image_depth; ++layer)
+            {
+                get_texture_readback_with_command_list(texture,
+                        tests[i].mip_level + (layer * tests[i].image_mips),
+                        &rb, queue, command_list);
+
+                for (p = 0; p < image_depth * image_size * image_size; ++p)
+                {
+                    x = p % image_size;
+                    y = (p / image_size) % image_size;
+                    z = p / (image_size * image_size);
+
+                    is_inside = tests[i].rect_count == 0;
+
+                    for (j = 0; j < tests[i].rect_count; ++j)
+                    {
+                        if (y >= tests[i].clear_rects[j].top && y < tests[i].clear_rects[j].bottom
+                                && x >= tests[i].clear_rects[j].left && x < tests[i].clear_rects[j].right)
+                        {
+                            is_inside = true;
+                            break;
+                        }
+                    }
+
+                    if (uav_dimensions[d].resource_dim == D3D12_RESOURCE_DIMENSION_TEXTURE3D)
+                        is_inside = is_inside && z >= tests[i].first_layer && z < tests[i].first_layer + tests[i].layer_count;
+                    else
+                        is_inside = is_inside && layer >= tests[i].first_layer && layer < tests[i].first_layer + tests[i].layer_count;
+
+                    expected_color = is_inside ? tests[i].expected : clear_value[0];
+                    actual_color = get_readback_uint(&rb, x, y, z);
+                    success = compare_color(actual_color, expected_color, tests[i].is_float ? 1 : 0);
+
+                    todo_if(tests[i].is_todo)
+                    ok(success, "At layer %u, (%u,%u,%u), expected %#x, got %#x.\n",
+                            layer, x, y, z, expected_color, actual_color);
+
+                    if (!success)
+                        break;
+                }
+
+                release_resource_readback(&rb);
+                reset_command_list(command_list, context.allocator);
+            }
+
+            ID3D12Resource_Release(texture);
+        }
+    }
+
+    ID3D12DescriptorHeap_Release(cpu_heap);
+    ID3D12DescriptorHeap_Release(gpu_heap);
+    destroy_test_context(&context);
+#undef IMAGE_SIZE
+}
+
 static void test_set_render_targets(void)
 {
     ID3D12DescriptorHeap *dsv_heap, *rtv_heap;
@@ -32451,6 +32712,7 @@ START_TEST(d3d12)
     run_test(test_clear_depth_stencil_view);
     run_test(test_clear_render_target_view);
     run_test(test_clear_unordered_access_view_buffer);
+    run_test(test_clear_unordered_access_view_image);
     run_test(test_set_render_targets);
     run_test(test_draw_instanced);
     run_test(test_draw_indexed_instanced);
diff --git a/tests/d3d12_test_utils.h b/tests/d3d12_test_utils.h
index 024cf87..824d78d 100644
--- a/tests/d3d12_test_utils.h
+++ b/tests/d3d12_test_utils.h
@@ -304,10 +304,15 @@ static unsigned int format_size(DXGI_FORMAT format)
         case DXGI_FORMAT_R32_FLOAT:
         case DXGI_FORMAT_R32_UINT:
         case DXGI_FORMAT_R32_SINT:
+        case DXGI_FORMAT_R16G16_FLOAT:
+        case DXGI_FORMAT_R16G16_UINT:
+        case DXGI_FORMAT_R16G16_UNORM:
         case DXGI_FORMAT_R8G8B8A8_TYPELESS:
+        case DXGI_FORMAT_R8G8B8A8_UINT:
         case DXGI_FORMAT_R8G8B8A8_UNORM:
         case DXGI_FORMAT_R8G8B8A8_UNORM_SRGB:
         case DXGI_FORMAT_B8G8R8A8_UNORM:
+        case DXGI_FORMAT_R11G11B10_FLOAT:
             return 4;
         case DXGI_FORMAT_R16_FLOAT:
         case DXGI_FORMAT_R16_UNORM:

[Sveinar Søpler]

Sveinar

On 11.11.2019 17:03, Philip Rebohle wrote:
> Note that this may conflict with the following pending patches (haven't checked in detail):
> - 171883 "vkd3d: Allocate one large buffer for a heap and offset into it."
> - 172639 "vkd3d: Store a copy of Vulkan view object handles in descriptors."

It does conflict with pending - 172677 "vkd3d: Support RS 1.0 VOLATILE 
descriptors."