Update to current vkd3d master (a59f198)

15 files changed, 1936 insertions, 448 deletions

diff --git a/Makefile.am b/Makefile.am
index 2b0e8f3f..00a5f589 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -108,6 +108,7 @@ libvkd3d_la_SOURCES = \
 	libs/vkd3d/vkd3d.map \
 	libs/vkd3d/vkd3d_main.c \
 	libs/vkd3d/vkd3d_private.h \
+	libs/vkd3d/vkd3d_shaders.h \
 	libs/vkd3d/vulkan_procs.h \
 	libs/vkd3d_version.c
 libvkd3d_la_LDFLAGS = $(AM_LDFLAGS) -version-info 2:0:1
diff --git a/configure.ac b/configure.ac
index 819e7685..355aaabf 100644
--- a/configure.ac
+++ b/configure.ac
@@ -78,6 +78,8 @@ AC_CHECK_DECL([SpvCapabilityDemoteToHelperInvocationEXT],, [AC_MSG_ERROR([SPIR-V
 # include "vulkan/spirv.h"
 #endif])
 
+AC_CHECK_DECLS([program_invocation_name],,,[#include <errno.h>])
+
 dnl Check for libraries
 m4_ifdef([PKG_PROG_PKG_CONFIG], [PKG_PROG_PKG_CONFIG], [m4_fatal([pkg-config autoconf macros not found.])])
 
diff --git a/include/vkd3d_d3d12.idl b/include/vkd3d_d3d12.idl
index ec8b83db..3bfe47bc 100644
--- a/include/vkd3d_d3d12.idl
+++ b/include/vkd3d_d3d12.idl
@@ -177,6 +177,13 @@ typedef enum D3D12_FORMAT_SUPPORT2
     D3D12_FORMAT_SUPPORT2_MULTIPLANE_OVERLAY = 0x00004000,
 } D3D12_FORMAT_SUPPORT2;
 
+typedef enum D3D12_WRITEBUFFERIMMEDIATE_MODE
+{
+    D3D12_WRITEBUFFERIMMEDIATE_MODE_DEFAULT = 0x0,
+    D3D12_WRITEBUFFERIMMEDIATE_MODE_MARKER_IN = 0x1,
+    D3D12_WRITEBUFFERIMMEDIATE_MODE_MARKER_OUT = 0x2,
+} D3D12_WRITEBUFFERIMMEDIATE_MODE;
+
 interface ID3D12Fence;
 interface ID3D12RootSignature;
 interface ID3D12Heap;
@@ -1657,6 +1664,12 @@ typedef enum D3D12_RESIDENCY_PRIORITY
     D3D12_RESIDENCY_PRIORITY_MAXIMUM = 0xc8000000,
 } D3D12_RESIDENCY_PRIORITY;
 
+typedef struct D3D12_WRITEBUFFERIMMEDIATE_PARAMETER
+{
+    D3D12_GPU_VIRTUAL_ADDRESS Dest;
+    UINT32 Value;
+} D3D12_WRITEBUFFERIMMEDIATE_PARAMETER;
+
 [
     uuid(c4fec28f-7966-4e95-9f94-f431cb56c3b8),
     object,
@@ -2000,6 +2013,21 @@ interface ID3D12GraphicsCommandList1 : ID3D12GraphicsCommandList
             UINT dst_sub_resource_idx, UINT dst_x, UINT dst_y,
             ID3D12Resource *src_resource, UINT src_sub_resource_idx,
             D3D12_RECT *src_rect, DXGI_FORMAT format, D3D12_RESOLVE_MODE mode);
+
+    void SetViewInstanceMask(UINT mask);
+}
+
+[
+    uuid(38c3e585-ff17-412c-9150-4fc6f9d72a28),
+    object,
+    local,
+    pointer_default(unique)
+]
+interface ID3D12GraphicsCommandList2 : ID3D12GraphicsCommandList1
+{
+    void WriteBufferImmediate(UINT count,
+            const D3D12_WRITEBUFFERIMMEDIATE_PARAMETER *parameters,
+            const D3D12_WRITEBUFFERIMMEDIATE_MODE *modes);
 }
 
 typedef enum D3D12_TILE_RANGE_FLAGS
diff --git a/libs/vkd3d-shader/spirv.c b/libs/vkd3d-shader/spirv.c
index a949e4a5..3d88be9e 100644
--- a/libs/vkd3d-shader/spirv.c
+++ b/libs/vkd3d-shader/spirv.c
@@ -2500,6 +2500,9 @@ static bool vkd3d_dxbc_compiler_get_register_name(char *buffer, unsigned int buf
         case VKD3DSPR_JOININSTID:
             snprintf(buffer, buffer_size, "vJoinInstanceId");
             break;
+        case VKD3DSPR_GSINSTID:
+            snprintf(buffer, buffer_size, "vGSInstanceID");
+            break;
         case VKD3DSPR_PATCHCONST:
             snprintf(buffer, buffer_size, "vpc%u", idx);
             break;
@@ -4438,8 +4441,13 @@ static void vkd3d_dxbc_compiler_emit_output(struct vkd3d_dxbc_compiler *compiler
         {
             use_private_variable = true;
             write_mask = VKD3DSP_WRITEMASK_ALL;
+            entry = rb_get(&compiler->symbol_table, &reg_symbol);
         }
     }
+    else if (!use_private_variable && (entry = rb_get(&compiler->symbol_table, &reg_symbol)))
+    {
+        id = RB_ENTRY_VALUE(entry, const struct vkd3d_symbol, entry)->id;
+    }
     else
     {
         if (builtin)
@@ -4487,15 +4495,15 @@ static void vkd3d_dxbc_compiler_emit_output(struct vkd3d_dxbc_compiler *compiler
             vkd3d_spirv_build_op_decorate(builder, id, SpvDecorationPatch, NULL, 0);
 
         vkd3d_dxbc_compiler_decorate_xfb_output(compiler, id, output_component_count, signature_element);
-
-        compiler->output_info[signature_idx].id = id;
-        compiler->output_info[signature_idx].component_type = component_type;
     }
 
+    compiler->output_info[signature_idx].id = id;
+    compiler->output_info[signature_idx].component_type = component_type;
+
     if (use_private_variable)
         storage_class = SpvStorageClassPrivate;
 
-    if ((entry = rb_get(&compiler->symbol_table, &reg_symbol)))
+    if (entry)
         var_id = RB_ENTRY_VALUE(entry, const struct vkd3d_symbol, entry)->id;
     else if (!use_private_variable)
         var_id = id;
diff --git a/libs/vkd3d-shader/vkd3d_shader_private.h b/libs/vkd3d-shader/vkd3d_shader_private.h
index 940cb768..100d5157 100644
--- a/libs/vkd3d-shader/vkd3d_shader_private.h
+++ b/libs/vkd3d-shader/vkd3d_shader_private.h
@@ -838,6 +838,8 @@ static inline enum vkd3d_component_type vkd3d_component_type_from_data_type(
     switch (data_type)
     {
         case VKD3D_DATA_FLOAT:
+        case VKD3D_DATA_UNORM:
+        case VKD3D_DATA_SNORM:
             return VKD3D_TYPE_FLOAT;
         case VKD3D_DATA_UINT:
             return VKD3D_TYPE_UINT;
diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index 0532ec0d..8a7ff668 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -1807,9 +1807,9 @@ HRESULT d3d12_command_allocator_create(struct d3d12_device *device,
 }
 
 /* ID3D12CommandList */
-static inline struct d3d12_command_list *impl_from_ID3D12GraphicsCommandList1(ID3D12GraphicsCommandList1 *iface)
+static inline struct d3d12_command_list *impl_from_ID3D12GraphicsCommandList2(ID3D12GraphicsCommandList2 *iface)
 {
-    return CONTAINING_RECORD(iface, struct d3d12_command_list, ID3D12GraphicsCommandList1_iface);
+    return CONTAINING_RECORD(iface, struct d3d12_command_list, ID3D12GraphicsCommandList2_iface);
 }
 
 static void d3d12_command_list_invalidate_current_framebuffer(struct d3d12_command_list *list)
@@ -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)
@@ -2146,19 +2159,20 @@ static void d3d12_command_list_track_resource_usage(struct d3d12_command_list *l
     }
 }
 
-static HRESULT STDMETHODCALLTYPE d3d12_command_list_QueryInterface(ID3D12GraphicsCommandList1 *iface,
+static HRESULT STDMETHODCALLTYPE d3d12_command_list_QueryInterface(ID3D12GraphicsCommandList2 *iface,
         REFIID iid, void **object)
 {
     TRACE("iface %p, iid %s, object %p.\n", iface, debugstr_guid(iid), object);
 
-    if (IsEqualGUID(iid, &IID_ID3D12GraphicsCommandList1)
+    if (IsEqualGUID(iid, &IID_ID3D12GraphicsCommandList2)
+            || IsEqualGUID(iid, &IID_ID3D12GraphicsCommandList1)
             || IsEqualGUID(iid, &IID_ID3D12GraphicsCommandList)
             || IsEqualGUID(iid, &IID_ID3D12CommandList)
             || IsEqualGUID(iid, &IID_ID3D12DeviceChild)
             || IsEqualGUID(iid, &IID_ID3D12Object)
             || IsEqualGUID(iid, &IID_IUnknown))
     {
-        ID3D12GraphicsCommandList1_AddRef(iface);
+        ID3D12GraphicsCommandList2_AddRef(iface);
         *object = iface;
         return S_OK;
     }
@@ -2169,9 +2183,9 @@ static HRESULT STDMETHODCALLTYPE d3d12_command_list_QueryInterface(ID3D12Graphic
     return E_NOINTERFACE;
 }
 
-static ULONG STDMETHODCALLTYPE d3d12_command_list_AddRef(ID3D12GraphicsCommandList1 *iface)
+static ULONG STDMETHODCALLTYPE d3d12_command_list_AddRef(ID3D12GraphicsCommandList2 *iface)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     ULONG refcount = InterlockedIncrement(&list->refcount);
 
     TRACE("%p increasing refcount to %u.\n", list, refcount);
@@ -2179,9 +2193,9 @@ static ULONG STDMETHODCALLTYPE d3d12_command_list_AddRef(ID3D12GraphicsCommandLi
     return refcount;
 }
 
-static ULONG STDMETHODCALLTYPE d3d12_command_list_Release(ID3D12GraphicsCommandList1 *iface)
+static ULONG STDMETHODCALLTYPE d3d12_command_list_Release(ID3D12GraphicsCommandList2 *iface)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     ULONG refcount = InterlockedDecrement(&list->refcount);
 
     TRACE("%p decreasing refcount to %u.\n", list, refcount);
@@ -2204,66 +2218,66 @@ static ULONG STDMETHODCALLTYPE d3d12_command_list_Release(ID3D12GraphicsCommandL
     return refcount;
 }
 
-static HRESULT STDMETHODCALLTYPE d3d12_command_list_GetPrivateData(ID3D12GraphicsCommandList1 *iface,
+static HRESULT STDMETHODCALLTYPE d3d12_command_list_GetPrivateData(ID3D12GraphicsCommandList2 *iface,
         REFGUID guid, UINT *data_size, void *data)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, guid %s, data_size %p, data %p.\n", iface, debugstr_guid(guid), data_size, data);
 
     return vkd3d_get_private_data(&list->private_store, guid, data_size, data);
 }
 
-static HRESULT STDMETHODCALLTYPE d3d12_command_list_SetPrivateData(ID3D12GraphicsCommandList1 *iface,
+static HRESULT STDMETHODCALLTYPE d3d12_command_list_SetPrivateData(ID3D12GraphicsCommandList2 *iface,
         REFGUID guid, UINT data_size, const void *data)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, guid %s, data_size %u, data %p.\n", iface, debugstr_guid(guid), data_size, data);
 
     return vkd3d_set_private_data(&list->private_store, guid, data_size, data);
 }
 
-static HRESULT STDMETHODCALLTYPE d3d12_command_list_SetPrivateDataInterface(ID3D12GraphicsCommandList1 *iface,
+static HRESULT STDMETHODCALLTYPE d3d12_command_list_SetPrivateDataInterface(ID3D12GraphicsCommandList2 *iface,
         REFGUID guid, const IUnknown *data)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, guid %s, data %p.\n", iface, debugstr_guid(guid), data);
 
     return vkd3d_set_private_data_interface(&list->private_store, guid, data);
 }
 
-static HRESULT STDMETHODCALLTYPE d3d12_command_list_SetName(ID3D12GraphicsCommandList1 *iface, const WCHAR *name)
+static HRESULT STDMETHODCALLTYPE d3d12_command_list_SetName(ID3D12GraphicsCommandList2 *iface, const WCHAR *name)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, name %s.\n", iface, debugstr_w(name, list->device->wchar_size));
 
     return name ? S_OK : E_INVALIDARG;
 }
 
-static HRESULT STDMETHODCALLTYPE d3d12_command_list_GetDevice(ID3D12GraphicsCommandList1 *iface, REFIID iid, void **device)
+static HRESULT STDMETHODCALLTYPE d3d12_command_list_GetDevice(ID3D12GraphicsCommandList2 *iface, REFIID iid, void **device)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, iid %s, device %p.\n", iface, debugstr_guid(iid), device);
 
     return d3d12_device_query_interface(list->device, iid, device);
 }
 
-static D3D12_COMMAND_LIST_TYPE STDMETHODCALLTYPE d3d12_command_list_GetType(ID3D12GraphicsCommandList1 *iface)
+static D3D12_COMMAND_LIST_TYPE STDMETHODCALLTYPE d3d12_command_list_GetType(ID3D12GraphicsCommandList2 *iface)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p.\n", iface);
 
     return list->type;
 }
 
-static HRESULT STDMETHODCALLTYPE d3d12_command_list_Close(ID3D12GraphicsCommandList1 *iface)
+static HRESULT STDMETHODCALLTYPE d3d12_command_list_Close(ID3D12GraphicsCommandList2 *iface)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct vkd3d_vk_device_procs *vk_procs;
     VkResult vr;
 
@@ -2307,7 +2321,7 @@ static HRESULT STDMETHODCALLTYPE d3d12_command_list_Close(ID3D12GraphicsCommandL
 static void d3d12_command_list_reset_state(struct d3d12_command_list *list,
         ID3D12PipelineState *initial_pipeline_state)
 {
-    ID3D12GraphicsCommandList1 *iface = &list->ID3D12GraphicsCommandList1_iface;
+    ID3D12GraphicsCommandList2 *iface = &list->ID3D12GraphicsCommandList2_iface;
 
     memset(list->strides, 0, sizeof(list->strides));
     list->primitive_topology = D3D_PRIMITIVE_TOPOLOGY_POINTLIST;
@@ -2337,14 +2351,14 @@ static void d3d12_command_list_reset_state(struct d3d12_command_list *list,
     memset(list->so_counter_buffers, 0, sizeof(list->so_counter_buffers));
     memset(list->so_counter_buffer_offsets, 0, sizeof(list->so_counter_buffer_offsets));
 
-    ID3D12GraphicsCommandList1_SetPipelineState(iface, initial_pipeline_state);
+    ID3D12GraphicsCommandList2_SetPipelineState(iface, initial_pipeline_state);
 }
 
-static HRESULT STDMETHODCALLTYPE d3d12_command_list_Reset(ID3D12GraphicsCommandList1 *iface,
+static HRESULT STDMETHODCALLTYPE d3d12_command_list_Reset(ID3D12GraphicsCommandList2 *iface,
         ID3D12CommandAllocator *allocator, ID3D12PipelineState *initial_pipeline_state)
 {
     struct d3d12_command_allocator *allocator_impl = unsafe_impl_from_ID3D12CommandAllocator(allocator);
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     HRESULT hr;
 
     TRACE("iface %p, allocator %p, initial_pipeline_state %p.\n",
@@ -2371,7 +2385,7 @@ static HRESULT STDMETHODCALLTYPE d3d12_command_list_Reset(ID3D12GraphicsCommandL
     return hr;
 }
 
-static HRESULT STDMETHODCALLTYPE d3d12_command_list_ClearState(ID3D12GraphicsCommandList1 *iface,
+static HRESULT STDMETHODCALLTYPE d3d12_command_list_ClearState(ID3D12GraphicsCommandList2 *iface,
         ID3D12PipelineState *pipeline_state)
 {
     FIXME("iface %p, pipline_state %p stub!\n", iface, pipeline_state);
@@ -2480,7 +2494,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 +2902,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 +2921,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;
@@ -2949,11 +2988,11 @@ static void d3d12_command_list_check_index_buffer_strip_cut_value(struct d3d12_c
     }
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_DrawInstanced(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_DrawInstanced(ID3D12GraphicsCommandList2 *iface,
         UINT vertex_count_per_instance, UINT instance_count, UINT start_vertex_location,
         UINT start_instance_location)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct vkd3d_vk_device_procs *vk_procs;
 
     TRACE("iface %p, vertex_count_per_instance %u, instance_count %u, "
@@ -2973,11 +3012,11 @@ static void STDMETHODCALLTYPE d3d12_command_list_DrawInstanced(ID3D12GraphicsCom
             instance_count, start_vertex_location, start_instance_location));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_DrawIndexedInstanced(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_DrawIndexedInstanced(ID3D12GraphicsCommandList2 *iface,
         UINT index_count_per_instance, UINT instance_count, UINT start_vertex_location,
         INT base_vertex_location, UINT start_instance_location)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct vkd3d_vk_device_procs *vk_procs;
 
     TRACE("iface %p, index_count_per_instance %u, instance_count %u, start_vertex_location %u, "
@@ -2999,33 +3038,29 @@ static void STDMETHODCALLTYPE d3d12_command_list_DrawIndexedInstanced(ID3D12Grap
             instance_count, start_vertex_location, base_vertex_location, start_instance_location));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_Dispatch(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_Dispatch(ID3D12GraphicsCommandList2 *iface,
         UINT x, UINT y, UINT z)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct vkd3d_vk_device_procs *vk_procs;
 
     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));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_CopyBufferRegion(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_CopyBufferRegion(ID3D12GraphicsCommandList2 *iface,
         ID3D12Resource *dst, UINT64 dst_offset, ID3D12Resource *src, UINT64 src_offset, UINT64 byte_count)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     struct d3d12_resource *dst_resource, *src_resource;
     const struct vkd3d_vk_device_procs *vk_procs;
     VkBufferCopy buffer_copy;
@@ -3305,11 +3340,11 @@ static bool validate_d3d12_box(const D3D12_BOX *box)
             && box->back > box->front;
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_CopyTextureRegion(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_CopyTextureRegion(ID3D12GraphicsCommandList2 *iface,
         const D3D12_TEXTURE_COPY_LOCATION *dst, UINT dst_x, UINT dst_y, UINT dst_z,
         const D3D12_TEXTURE_COPY_LOCATION *src, const D3D12_BOX *src_box)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     struct d3d12_resource *dst_resource, *src_resource;
     const struct vkd3d_format *src_format, *dst_format;
     const struct vkd3d_vk_device_procs *vk_procs;
@@ -3440,10 +3475,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_CopyTextureRegion(ID3D12Graphic
     }
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_CopyResource(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_CopyResource(ID3D12GraphicsCommandList2 *iface,
         ID3D12Resource *dst, ID3D12Resource *src)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     struct d3d12_resource *dst_resource, *src_resource;
     const struct vkd3d_format *src_format, *dst_format;
     const struct vkd3d_vk_device_procs *vk_procs;
@@ -3510,7 +3545,7 @@ static void STDMETHODCALLTYPE d3d12_command_list_CopyResource(ID3D12GraphicsComm
     }
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_CopyTiles(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_CopyTiles(ID3D12GraphicsCommandList2 *iface,
         ID3D12Resource *tiled_resource, const D3D12_TILED_RESOURCE_COORDINATE *tile_region_start_coordinate,
         const D3D12_TILE_REGION_SIZE *tile_region_size, ID3D12Resource *buffer, UINT64 buffer_offset,
         D3D12_TILE_COPY_FLAGS flags)
@@ -3521,11 +3556,11 @@ static void STDMETHODCALLTYPE d3d12_command_list_CopyTiles(ID3D12GraphicsCommand
             buffer, buffer_offset, flags);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_ResolveSubresource(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_ResolveSubresource(ID3D12GraphicsCommandList2 *iface,
         ID3D12Resource *dst, UINT dst_sub_resource_idx,
         ID3D12Resource *src, UINT src_sub_resource_idx, DXGI_FORMAT format)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct vkd3d_format *src_format, *dst_format, *vk_format;
     struct d3d12_resource *dst_resource, *src_resource;
     const struct vkd3d_vk_device_procs *vk_procs;
@@ -3596,10 +3631,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_ResolveSubresource(ID3D12Graphi
             VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, 1, &vk_image_resolve));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_IASetPrimitiveTopology(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_IASetPrimitiveTopology(ID3D12GraphicsCommandList2 *iface,
         D3D12_PRIMITIVE_TOPOLOGY topology)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, topology %#x.\n", iface, topology);
 
@@ -3616,11 +3651,11 @@ static void STDMETHODCALLTYPE d3d12_command_list_IASetPrimitiveTopology(ID3D12Gr
     d3d12_command_list_invalidate_current_pipeline(list);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_RSSetViewports(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_RSSetViewports(ID3D12GraphicsCommandList2 *iface,
         UINT viewport_count, const D3D12_VIEWPORT *viewports)
 {
     VkViewport vk_viewports[D3D12_VIEWPORT_AND_SCISSORRECT_OBJECT_COUNT_PER_PIPELINE];
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct vkd3d_vk_device_procs *vk_procs;
     unsigned int i;
 
@@ -3652,10 +3687,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_RSSetViewports(ID3D12GraphicsCo
     VK_CALL(vkCmdSetViewport(list->vk_command_buffer, 0, viewport_count, vk_viewports));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_RSSetScissorRects(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_RSSetScissorRects(ID3D12GraphicsCommandList2 *iface,
         UINT rect_count, const D3D12_RECT *rects)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     VkRect2D vk_rects[D3D12_VIEWPORT_AND_SCISSORRECT_OBJECT_COUNT_PER_PIPELINE];
     const struct vkd3d_vk_device_procs *vk_procs;
     unsigned int i;
@@ -3680,10 +3715,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_RSSetScissorRects(ID3D12Graphic
     VK_CALL(vkCmdSetScissor(list->vk_command_buffer, 0, rect_count, vk_rects));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_OMSetBlendFactor(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_OMSetBlendFactor(ID3D12GraphicsCommandList2 *iface,
         const FLOAT blend_factor[4])
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct vkd3d_vk_device_procs *vk_procs;
 
     TRACE("iface %p, blend_factor %p.\n", iface, blend_factor);
@@ -3692,10 +3727,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_OMSetBlendFactor(ID3D12Graphics
     VK_CALL(vkCmdSetBlendConstants(list->vk_command_buffer, blend_factor));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_OMSetStencilRef(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_OMSetStencilRef(ID3D12GraphicsCommandList2 *iface,
         UINT stencil_ref)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct vkd3d_vk_device_procs *vk_procs;
 
     TRACE("iface %p, stencil_ref %u.\n", iface, stencil_ref);
@@ -3704,30 +3739,19 @@ static void STDMETHODCALLTYPE d3d12_command_list_OMSetStencilRef(ID3D12GraphicsC
     VK_CALL(vkCmdSetStencilReference(list->vk_command_buffer, VK_STENCIL_FRONT_AND_BACK, stencil_ref));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetPipelineState(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetPipelineState(ID3D12GraphicsCommandList2 *iface,
         ID3D12PipelineState *pipeline_state)
 {
     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;
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     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;
 }
@@ -3770,10 +3794,10 @@ static unsigned int d3d12_find_ds_multiplanar_transition(const D3D12_RESOURCE_BA
     return 0;
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_ResourceBarrier(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_ResourceBarrier(ID3D12GraphicsCommandList2 *iface,
         UINT barrier_count, const D3D12_RESOURCE_BARRIER *barriers)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     bool have_aliasing_barriers = false, have_split_barriers = false;
     const struct vkd3d_vk_device_procs *vk_procs;
     const struct vkd3d_vulkan_info *vk_info;
@@ -4003,13 +4027,13 @@ static void STDMETHODCALLTYPE d3d12_command_list_ResourceBarrier(ID3D12GraphicsC
         WARN("Issuing split barrier(s) on D3D12_RESOURCE_BARRIER_FLAG_END_ONLY.\n");
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_ExecuteBundle(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_ExecuteBundle(ID3D12GraphicsCommandList2 *iface,
         ID3D12GraphicsCommandList *command_list)
 {
     FIXME("iface %p, command_list %p stub!\n", iface, command_list);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetDescriptorHeaps(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetDescriptorHeaps(ID3D12GraphicsCommandList2 *iface,
         UINT heap_count, ID3D12DescriptorHeap *const *heaps)
 {
     TRACE("iface %p, heap_count %u, heaps %p.\n", iface, heap_count, heaps);
@@ -4029,15 +4053,14 @@ 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,
+static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootSignature(ID3D12GraphicsCommandList2 *iface,
         ID3D12RootSignature *root_signature)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_signature %p.\n", iface, root_signature);
 
@@ -4045,10 +4068,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootSignature(ID3D12G
             unsafe_impl_from_ID3D12RootSignature(root_signature));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRootSignature(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRootSignature(ID3D12GraphicsCommandList2 *iface,
         ID3D12RootSignature *root_signature)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_signature %p.\n", iface, root_signature);
 
@@ -4070,10 +4093,10 @@ static void d3d12_command_list_set_descriptor_table(struct d3d12_command_list *l
     bindings->descriptor_table_active_mask |= (uint64_t)1 << index;
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootDescriptorTable(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootDescriptorTable(ID3D12GraphicsCommandList2 *iface,
         UINT root_parameter_index, D3D12_GPU_DESCRIPTOR_HANDLE base_descriptor)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, base_descriptor %#"PRIx64".\n",
             iface, root_parameter_index, base_descriptor.ptr);
@@ -4082,10 +4105,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootDescriptorTable(I
             root_parameter_index, base_descriptor);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRootDescriptorTable(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRootDescriptorTable(ID3D12GraphicsCommandList2 *iface,
         UINT root_parameter_index, D3D12_GPU_DESCRIPTOR_HANDLE base_descriptor)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, base_descriptor %#"PRIx64".\n",
             iface, root_parameter_index, base_descriptor.ptr);
@@ -4107,10 +4130,10 @@ static void d3d12_command_list_set_root_constants(struct d3d12_command_list *lis
             c->stage_flags, c->offset + offset * sizeof(uint32_t), count * sizeof(uint32_t), data));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRoot32BitConstant(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRoot32BitConstant(ID3D12GraphicsCommandList2 *iface,
         UINT root_parameter_index, UINT data, UINT dst_offset)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, data 0x%08x, dst_offset %u.\n",
             iface, root_parameter_index, data, dst_offset);
@@ -4119,10 +4142,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRoot32BitConstant(ID3
             root_parameter_index, dst_offset, 1, &data);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRoot32BitConstant(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRoot32BitConstant(ID3D12GraphicsCommandList2 *iface,
         UINT root_parameter_index, UINT data, UINT dst_offset)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, data 0x%08x, dst_offset %u.\n",
             iface, root_parameter_index, data, dst_offset);
@@ -4131,10 +4154,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRoot32BitConstant(ID
             root_parameter_index, dst_offset, 1, &data);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRoot32BitConstants(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRoot32BitConstants(ID3D12GraphicsCommandList2 *iface,
         UINT root_parameter_index, UINT constant_count, const void *data, UINT dst_offset)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, constant_count %u, data %p, dst_offset %u.\n",
             iface, root_parameter_index, constant_count, data, dst_offset);
@@ -4143,10 +4166,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRoot32BitConstants(ID
             root_parameter_index, dst_offset, constant_count, data);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRoot32BitConstants(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRoot32BitConstants(ID3D12GraphicsCommandList2 *iface,
         UINT root_parameter_index, UINT constant_count, const void *data, UINT dst_offset)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, constant_count %u, data %p, dst_offset %u.\n",
             iface, root_parameter_index, constant_count, data, dst_offset);
@@ -4199,9 +4222,9 @@ static void d3d12_command_list_set_root_cbv(struct d3d12_command_list *list,
 }
 
 static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootConstantBufferView(
-        ID3D12GraphicsCommandList1 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
+        ID3D12GraphicsCommandList2 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, address %#"PRIx64".\n",
             iface, root_parameter_index, address);
@@ -4210,9 +4233,9 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootConstantBufferVie
 }
 
 static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRootConstantBufferView(
-        ID3D12GraphicsCommandList1 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
+        ID3D12GraphicsCommandList2 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, address %#"PRIx64".\n",
             iface, root_parameter_index, address);
@@ -4271,9 +4294,9 @@ static void d3d12_command_list_set_root_descriptor(struct d3d12_command_list *li
 }
 
 static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootShaderResourceView(
-        ID3D12GraphicsCommandList1 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
+        ID3D12GraphicsCommandList2 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, address %#"PRIx64".\n",
             iface, root_parameter_index, address);
@@ -4283,9 +4306,9 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootShaderResourceVie
 }
 
 static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRootShaderResourceView(
-        ID3D12GraphicsCommandList1 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
+        ID3D12GraphicsCommandList2 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, address %#"PRIx64".\n",
             iface, root_parameter_index, address);
@@ -4295,9 +4318,9 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRootShaderResourceVi
 }
 
 static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootUnorderedAccessView(
-        ID3D12GraphicsCommandList1 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
+        ID3D12GraphicsCommandList2 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, address %#"PRIx64".\n",
             iface, root_parameter_index, address);
@@ -4307,9 +4330,9 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetComputeRootUnorderedAccessVi
 }
 
 static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRootUnorderedAccessView(
-        ID3D12GraphicsCommandList1 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
+        ID3D12GraphicsCommandList2 *iface, UINT root_parameter_index, D3D12_GPU_VIRTUAL_ADDRESS address)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
 
     TRACE("iface %p, root_parameter_index %u, address %#"PRIx64".\n",
             iface, root_parameter_index, address);
@@ -4318,10 +4341,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetGraphicsRootUnorderedAccessV
             root_parameter_index, address);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_IASetIndexBuffer(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_IASetIndexBuffer(ID3D12GraphicsCommandList2 *iface,
         const D3D12_INDEX_BUFFER_VIEW *view)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct vkd3d_vk_device_procs *vk_procs;
     struct d3d12_resource *resource;
     enum VkIndexType index_type;
@@ -4356,10 +4379,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_IASetIndexBuffer(ID3D12Graphics
             view->BufferLocation - resource->gpu_address, index_type));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_IASetVertexBuffers(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_IASetVertexBuffers(ID3D12GraphicsCommandList2 *iface,
         UINT start_slot, UINT view_count, const D3D12_VERTEX_BUFFER_VIEW *views)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct vkd3d_null_resources *null_resources;
     struct vkd3d_gpu_va_allocator *gpu_va_allocator;
     VkDeviceSize offsets[ARRAY_SIZE(list->strides)];
@@ -4408,10 +4431,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_IASetVertexBuffers(ID3D12Graphi
         d3d12_command_list_invalidate_current_pipeline(list);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SOSetTargets(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SOSetTargets(ID3D12GraphicsCommandList2 *iface,
         UINT start_slot, UINT view_count, const D3D12_STREAM_OUTPUT_BUFFER_VIEW *views)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     VkDeviceSize offsets[ARRAY_SIZE(list->so_counter_buffers)];
     VkDeviceSize sizes[ARRAY_SIZE(list->so_counter_buffers)];
     VkBuffer buffers[ARRAY_SIZE(list->so_counter_buffers)];
@@ -4473,11 +4496,11 @@ static void STDMETHODCALLTYPE d3d12_command_list_SOSetTargets(ID3D12GraphicsComm
         VK_CALL(vkCmdBindTransformFeedbackBuffersEXT(list->vk_command_buffer, first, count, buffers, offsets, sizes));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_OMSetRenderTargets(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_OMSetRenderTargets(ID3D12GraphicsCommandList2 *iface,
         UINT render_target_descriptor_count, const D3D12_CPU_DESCRIPTOR_HANDLE *render_target_descriptors,
         BOOL single_descriptor_handle, const D3D12_CPU_DESCRIPTOR_HANDLE *depth_stencil_descriptor)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct d3d12_rtv_desc *rtv_desc;
     const struct d3d12_dsv_desc *dsv_desc;
     VkFormat prev_dsv_format;
@@ -4678,12 +4701,12 @@ static void d3d12_command_list_clear(struct d3d12_command_list *list,
     }
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_ClearDepthStencilView(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_ClearDepthStencilView(ID3D12GraphicsCommandList2 *iface,
         D3D12_CPU_DESCRIPTOR_HANDLE dsv, D3D12_CLEAR_FLAGS flags, float depth, UINT8 stencil,
         UINT rect_count, const D3D12_RECT *rects)
 {
     const union VkClearValue clear_value = {.depthStencil = {depth, stencil}};
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct d3d12_dsv_desc *dsv_desc = d3d12_dsv_desc_from_cpu_handle(dsv);
     struct VkAttachmentDescription attachment_desc;
     struct VkAttachmentReference ds_reference;
@@ -4727,10 +4750,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearDepthStencilView(ID3D12Gra
             &clear_value, rect_count, rects);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_ClearRenderTargetView(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_ClearRenderTargetView(ID3D12GraphicsCommandList2 *iface,
         D3D12_CPU_DESCRIPTOR_HANDLE rtv, const FLOAT color[4], UINT rect_count, const D3D12_RECT *rects)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const struct d3d12_rtv_desc *rtv_desc = d3d12_rtv_desc_from_cpu_handle(rtv);
     struct VkAttachmentDescription attachment_desc;
     struct VkAttachmentReference color_reference;
@@ -4781,128 +4804,289 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearRenderTargetView(ID3D12Gra
             &clear_value, rect_count, rects);
 }
 
-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)
+struct vkd3d_uav_clear_pipeline
 {
-    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 d3d12_resource *resource_impl;
-    VkBufferMemoryBarrier buffer_barrier;
-    VkImageMemoryBarrier image_barrier;
-    VkPipelineStageFlags stage_mask;
-    VkImageSubresourceRange range;
-    VkClearColorValue color;
+    VkDescriptorSetLayout vk_set_layout;
+    VkPipelineLayout vk_pipeline_layout;
+    VkPipeline vk_pipeline;
+    VkExtent3D group_size;
+};
 
-    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);
+static void vkd3d_uav_clear_state_get_buffer_pipeline(const struct vkd3d_uav_clear_state *state,
+        enum vkd3d_format_type format_type, struct vkd3d_uav_clear_pipeline *info)
+{
+    const struct vkd3d_uav_clear_pipelines *pipelines;
 
-    vk_procs = &list->device->vk_procs;
-    vk_info = &list->device->vk_info;
+    pipelines = format_type == VKD3D_FORMAT_TYPE_UINT ? &state->pipelines_uint : &state->pipelines_float;
+    info->vk_set_layout = state->vk_set_layout_buffer;
+    info->vk_pipeline_layout = state->vk_pipeline_layout_buffer;
+    info->vk_pipeline = pipelines->buffer;
+    info->group_size = (VkExtent3D){128, 1, 1};
+}
 
-    resource_impl = unsafe_impl_from_ID3D12Resource(resource);
+static void vkd3d_uav_clear_state_get_image_pipeline(const struct vkd3d_uav_clear_state *state,
+        VkImageViewType image_view_type, enum vkd3d_format_type format_type, struct vkd3d_uav_clear_pipeline *info)
+{
+    const struct vkd3d_uav_clear_pipelines *pipelines;
 
-    d3d12_command_list_track_resource_usage(list, resource_impl);
+    pipelines = format_type == VKD3D_FORMAT_TYPE_UINT ? &state->pipelines_uint : &state->pipelines_float;
+    info->vk_set_layout = state->vk_set_layout_image;
+    info->vk_pipeline_layout = state->vk_pipeline_layout_image;
 
-    if (rect_count)
+    switch (image_view_type)
     {
-        FIXME("Clear rects not supported.\n");
-        return;
+        case VK_IMAGE_VIEW_TYPE_1D:
+            info->vk_pipeline = pipelines->image_1d;
+            info->group_size = (VkExtent3D){64, 1, 1};
+            break;
+
+        case VK_IMAGE_VIEW_TYPE_1D_ARRAY:
+            info->vk_pipeline = pipelines->image_1d_array;
+            info->group_size = (VkExtent3D){64, 1, 1};
+            break;
+
+        case VK_IMAGE_VIEW_TYPE_2D:
+            info->vk_pipeline = pipelines->image_2d;
+            info->group_size = (VkExtent3D){8, 8, 1};
+            break;
+
+        case VK_IMAGE_VIEW_TYPE_2D_ARRAY:
+            info->vk_pipeline = pipelines->image_2d_array;
+            info->group_size = (VkExtent3D){8, 8, 1};
+            break;
+
+        case VK_IMAGE_VIEW_TYPE_3D:
+            info->vk_pipeline = pipelines->image_3d;
+            info->group_size = (VkExtent3D){8, 8, 1};
+            break;
+
+        default:
+            ERR("Unhandled view type %#x.\n", image_view_type);
+            info->vk_pipeline = VK_NULL_HANDLE;
+            info->group_size = (VkExtent3D){0, 0, 0};
+            break;
     }
+}
 
+static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
+        struct d3d12_resource *resource, struct vkd3d_view *view, const VkClearColorValue *clear_colour,
+        unsigned int rect_count, const D3D12_RECT *rects)
+{
+    const struct vkd3d_vk_device_procs *vk_procs = &list->device->vk_procs;
+    unsigned int i, miplevel_idx, layer_count;
+    struct vkd3d_uav_clear_pipeline pipeline;
+    struct vkd3d_uav_clear_args clear_args;
+    VkDescriptorImageInfo image_info;
+    D3D12_RECT full_rect, curr_rect;
+    VkWriteDescriptorSet write_set;
+
+    d3d12_command_list_track_resource_usage(list, resource);
     d3d12_command_list_end_current_render_pass(list);
 
-    cpu_descriptor = d3d12_desc_from_cpu_handle(cpu_handle);
+    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_resource_is_buffer(resource_impl))
-    {
-        if (!cpu_descriptor->uav.buffer.size)
-        {
-            FIXME("Not supported for UAV descriptor %p.\n", cpu_descriptor);
-            return;
-        }
+    if (!d3d12_command_allocator_add_view(list->allocator, view))
+        WARN("Failed to add view.\n");
 
-        VK_CALL(vkCmdFillBuffer(list->vk_command_buffer, resource_impl->u.vk_buffer,
-                cpu_descriptor->uav.buffer.offset, cpu_descriptor->uav.buffer.size, values[0]));
+    clear_args.colour = *clear_colour;
 
-        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->uav.buffer.offset;
-        buffer_barrier.size = cpu_descriptor->uav.buffer.size;
+    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;
 
-        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);
+    if (d3d12_resource_is_buffer(resource))
+    {
+        write_set.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER;
+        write_set.pImageInfo = NULL;
+        write_set.pBufferInfo = NULL;
+        write_set.pTexelBufferView = &view->u.vk_buffer_view;
 
-        VK_CALL(vkCmdPipelineBarrier(list->vk_command_buffer,
-                VK_PIPELINE_STAGE_TRANSFER_BIT, stage_mask, 0,
-                0, NULL, 1, &buffer_barrier, 0, NULL));
+        miplevel_idx = 0;
+        layer_count = 1;
+        vkd3d_uav_clear_state_get_buffer_pipeline(&list->device->uav_clear_state,
+                view->format->type, &pipeline);
     }
     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->uav.texture.vk_aspect_mask;
-        range.baseMipLevel = cpu_descriptor->uav.texture.miplevel_idx;
-        range.levelCount = 1;
-        range.baseArrayLayer = cpu_descriptor->uav.texture.layer_idx;
-        range.layerCount = cpu_descriptor->uav.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);
+        image_info.sampler = VK_NULL_HANDLE;
+        image_info.imageView = view->u.vk_image_view;
+        image_info.imageLayout = VK_IMAGE_LAYOUT_GENERAL;
 
-        VK_CALL(vkCmdPipelineBarrier(list->vk_command_buffer,
-                VK_PIPELINE_STAGE_TRANSFER_BIT, stage_mask, 0,
-                0, NULL, 0, NULL, 1, &image_barrier));
+        write_set.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE;
+        write_set.pImageInfo = &image_info;
+        write_set.pBufferInfo = NULL;
+        write_set.pTexelBufferView = NULL;
+
+        miplevel_idx = view->info.texture.miplevel_idx;
+        layer_count = view->info.texture.vk_view_type == VK_IMAGE_VIEW_TYPE_3D
+                ? d3d12_resource_desc_get_depth(&resource->desc, miplevel_idx)
+                : view->info.texture.layer_count;
+        vkd3d_uav_clear_state_get_image_pipeline(&list->device->uav_clear_state,
+                view->info.texture.vk_view_type, view->format->type, &pipeline);
+    }
+
+    if (!(write_set.dstSet = d3d12_command_allocator_allocate_descriptor_set(
+            list->allocator, pipeline.vk_set_layout)))
+    {
+        ERR("Failed to allocate descriptor set.\n");
+        return;
+    }
+
+    VK_CALL(vkUpdateDescriptorSets(list->device->vk_device, 1, &write_set, 0, NULL));
+
+    full_rect.left = 0;
+    full_rect.right = d3d12_resource_desc_get_width(&resource->desc, miplevel_idx);
+    full_rect.top = 0;
+    full_rect.bottom = d3d12_resource_desc_get_height(&resource->desc, miplevel_idx);
+
+    if (!rect_count)
+    {
+        rects = &full_rect;
+        rect_count = 1;
+    }
+
+    VK_CALL(vkCmdBindPipeline(list->vk_command_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline.vk_pipeline));
+
+    VK_CALL(vkCmdBindDescriptorSets(list->vk_command_buffer, VK_PIPELINE_BIND_POINT_COMPUTE,
+            pipeline.vk_pipeline_layout, 0, 1, &write_set.dstSet, 0, NULL));
+
+    for (i = 0; i < rect_count; ++i)
+    {
+        /* Clamp to the actual resource region and skip empty rectangles. */
+        curr_rect.left = max(rects[i].left, full_rect.left);
+        curr_rect.top = max(rects[i].top, full_rect.top);
+        curr_rect.right = min(rects[i].right, full_rect.right);
+        curr_rect.bottom = min(rects[i].bottom, full_rect.bottom);
+
+        if (curr_rect.left >= curr_rect.right || curr_rect.top >= curr_rect.bottom)
+            continue;
+
+        clear_args.offset.x = curr_rect.left;
+        clear_args.offset.y = curr_rect.top;
+        clear_args.extent.width = curr_rect.right - curr_rect.left;
+        clear_args.extent.height = curr_rect.bottom - curr_rect.top;
+
+        VK_CALL(vkCmdPushConstants(list->vk_command_buffer, pipeline.vk_pipeline_layout,
+                VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(clear_args), &clear_args));
+
+        VK_CALL(vkCmdDispatch(list->vk_command_buffer,
+                vkd3d_compute_workgroup_count(clear_args.extent.width, pipeline.group_size.width),
+                vkd3d_compute_workgroup_count(clear_args.extent.height, pipeline.group_size.height),
+                vkd3d_compute_workgroup_count(layer_count, pipeline.group_size.depth)));
+    }
+}
+
+static const struct vkd3d_format *vkd3d_fixup_clear_uav_uint_colour(struct d3d12_device *device,
+        DXGI_FORMAT dxgi_format, VkClearColorValue *colour)
+{
+    switch (dxgi_format)
+    {
+        case DXGI_FORMAT_R11G11B10_FLOAT:
+            colour->uint32[0] = (colour->uint32[0] & 0x7ff)
+                    | ((colour->uint32[1] & 0x7ff) << 11)
+                    | ((colour->uint32[2] & 0x3ff) << 22);
+            return vkd3d_get_format(device, DXGI_FORMAT_R32_UINT, false);
+
+        default:
+            return NULL;
+    }
+}
+
+static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewUint(ID3D12GraphicsCommandList2 *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)
+{
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
+    struct d3d12_device *device = list->device;
+    struct vkd3d_view *view, *uint_view = NULL;
+    struct vkd3d_texture_view_desc view_desc;
+    const struct vkd3d_format *uint_format;
+    struct d3d12_resource *resource_impl;
+    VkClearColorValue colour;
+
+    TRACE("iface %p, gpu_handle %#"PRIx64", cpu_handle %lx, resource %p, values %p, rect_count %u, rects %p.\n",
+            iface, gpu_handle.ptr, cpu_handle.ptr, resource, values, rect_count, rects);
+
+    resource_impl = unsafe_impl_from_ID3D12Resource(resource);
+    view = d3d12_desc_from_cpu_handle(cpu_handle)->u.view;
+    memcpy(colour.uint32, values, sizeof(colour.uint32));
+
+    if (view->format->type != VKD3D_FORMAT_TYPE_UINT)
+    {
+        if (!(uint_format = vkd3d_find_uint_format(device, view->format->dxgi_format))
+                && !(uint_format = vkd3d_fixup_clear_uav_uint_colour(device, view->format->dxgi_format, &colour)))
+        {
+            ERR("Unhandled format %#x.\n", view->format->dxgi_format);
+            return;
+        }
+
+        if (d3d12_resource_is_buffer(resource_impl))
+        {
+            if (!vkd3d_create_buffer_view(device, resource_impl->u.vk_buffer, uint_format,
+                    view->info.buffer.offset, view->info.buffer.size, &uint_view))
+            {
+                ERR("Failed to create buffer view.\n");
+                return;
+            }
+        }
+        else
+        {
+            memset(&view_desc, 0, sizeof(view_desc));
+            view_desc.view_type = view->info.texture.vk_view_type;
+            view_desc.format = uint_format;
+            view_desc.miplevel_idx = view->info.texture.miplevel_idx;
+            view_desc.miplevel_count = 1;
+            view_desc.layer_idx = view->info.texture.layer_idx;
+            view_desc.layer_count = view->info.texture.layer_count;
+
+            if (!vkd3d_create_texture_view(device, resource_impl->u.vk_image, &view_desc, &uint_view))
+            {
+                ERR("Failed to create image view.\n");
+                return;
+            }
+        }
+        view = uint_view;
     }
+
+    d3d12_command_list_clear_uav(list, resource_impl, view, &colour, rect_count, rects);
+
+    if (uint_view)
+        vkd3d_view_decref(uint_view, device);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewFloat(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewFloat(ID3D12GraphicsCommandList2 *iface,
         D3D12_GPU_DESCRIPTOR_HANDLE gpu_handle, D3D12_CPU_DESCRIPTOR_HANDLE cpu_handle, ID3D12Resource *resource,
         const float values[4], UINT rect_count, const D3D12_RECT *rects)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     struct d3d12_resource *resource_impl;
+    VkClearColorValue colour;
+    struct vkd3d_view *view;
 
-    FIXME("iface %p, gpu_handle %#"PRIx64", cpu_handle %lx, resource %p, values %p, rect_count %u, rects %p stub!\n",
+    TRACE("iface %p, gpu_handle %#"PRIx64", cpu_handle %lx, resource %p, values %p, rect_count %u, rects %p.\n",
             iface, gpu_handle.ptr, cpu_handle.ptr, resource, values, rect_count, rects);
 
     resource_impl = unsafe_impl_from_ID3D12Resource(resource);
+    view = d3d12_desc_from_cpu_handle(cpu_handle)->u.view;
+    memcpy(colour.float32, values, sizeof(colour.float32));
 
-    d3d12_command_list_track_resource_usage(list, resource_impl);
+    d3d12_command_list_clear_uav(list, resource_impl, view, &colour, rect_count, rects);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_DiscardResource(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_DiscardResource(ID3D12GraphicsCommandList2 *iface,
         ID3D12Resource *resource, const D3D12_DISCARD_REGION *region)
 {
     FIXME_ONCE("iface %p, resource %p, region %p stub!\n", iface, resource, region);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_BeginQuery(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_BeginQuery(ID3D12GraphicsCommandList2 *iface,
         ID3D12QueryHeap *heap, D3D12_QUERY_TYPE type, UINT index)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     struct d3d12_query_heap *query_heap = unsafe_impl_from_ID3D12QueryHeap(heap);
     const struct vkd3d_vk_device_procs *vk_procs;
     VkQueryControlFlags flags = 0;
@@ -4929,10 +5113,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_BeginQuery(ID3D12GraphicsComman
     VK_CALL(vkCmdBeginQuery(list->vk_command_buffer, query_heap->vk_query_pool, index, flags));
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_EndQuery(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_EndQuery(ID3D12GraphicsCommandList2 *iface,
         ID3D12QueryHeap *heap, D3D12_QUERY_TYPE type, UINT index)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     struct d3d12_query_heap *query_heap = unsafe_impl_from_ID3D12QueryHeap(heap);
     const struct vkd3d_vk_device_procs *vk_procs;
 
@@ -4974,12 +5158,12 @@ static size_t get_query_stride(D3D12_QUERY_TYPE type)
     return sizeof(uint64_t);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_ResolveQueryData(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_ResolveQueryData(ID3D12GraphicsCommandList2 *iface,
         ID3D12QueryHeap *heap, D3D12_QUERY_TYPE type, UINT start_index, UINT query_count,
         ID3D12Resource *dst_buffer, UINT64 aligned_dst_buffer_offset)
 {
     const struct d3d12_query_heap *query_heap = unsafe_impl_from_ID3D12QueryHeap(heap);
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     struct d3d12_resource *buffer = unsafe_impl_from_ID3D12Resource(dst_buffer);
     const struct vkd3d_vk_device_procs *vk_procs;
     unsigned int i, first, count;
@@ -5055,10 +5239,10 @@ static void STDMETHODCALLTYPE d3d12_command_list_ResolveQueryData(ID3D12Graphics
     }
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetPredication(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetPredication(ID3D12GraphicsCommandList2 *iface,
         ID3D12Resource *buffer, UINT64 aligned_buffer_offset, D3D12_PREDICATION_OP operation)
 {
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     struct d3d12_resource *resource = unsafe_impl_from_ID3D12Resource(buffer);
     const struct vkd3d_vulkan_info *vk_info = &list->device->vk_info;
     const struct vkd3d_vk_device_procs *vk_procs;
@@ -5127,19 +5311,19 @@ static void STDMETHODCALLTYPE d3d12_command_list_SetPredication(ID3D12GraphicsCo
     }
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetMarker(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetMarker(ID3D12GraphicsCommandList2 *iface,
         UINT metadata, const void *data, UINT size)
 {
     FIXME("iface %p, metadata %#x, data %p, size %u stub!\n", iface, metadata, data, size);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_BeginEvent(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_BeginEvent(ID3D12GraphicsCommandList2 *iface,
         UINT metadata, const void *data, UINT size)
 {
     FIXME("iface %p, metadata %#x, data %p, size %u stub!\n", iface, metadata, data, size);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_EndEvent(ID3D12GraphicsCommandList1 *iface)
+static void STDMETHODCALLTYPE d3d12_command_list_EndEvent(ID3D12GraphicsCommandList2 *iface)
 {
     FIXME("iface %p stub!\n", iface);
 }
@@ -5148,14 +5332,14 @@ STATIC_ASSERT(sizeof(VkDispatchIndirectCommand) == sizeof(D3D12_DISPATCH_ARGUMEN
 STATIC_ASSERT(sizeof(VkDrawIndexedIndirectCommand) == sizeof(D3D12_DRAW_INDEXED_ARGUMENTS));
 STATIC_ASSERT(sizeof(VkDrawIndirectCommand) == sizeof(D3D12_DRAW_ARGUMENTS));
 
-static void STDMETHODCALLTYPE d3d12_command_list_ExecuteIndirect(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_ExecuteIndirect(ID3D12GraphicsCommandList2 *iface,
         ID3D12CommandSignature *command_signature, UINT max_command_count, ID3D12Resource *arg_buffer,
         UINT64 arg_buffer_offset, ID3D12Resource *count_buffer, UINT64 count_buffer_offset)
 {
     struct d3d12_command_signature *sig_impl = unsafe_impl_from_ID3D12CommandSignature(command_signature);
     struct d3d12_resource *count_impl = unsafe_impl_from_ID3D12Resource(count_buffer);
     struct d3d12_resource *arg_impl = unsafe_impl_from_ID3D12Resource(arg_buffer);
-    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList1(iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
     const D3D12_COMMAND_SIGNATURE_DESC *signature_desc;
     const struct vkd3d_vk_device_procs *vk_procs;
     unsigned int i;
@@ -5232,17 +5416,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;
@@ -5254,7 +5433,7 @@ static void STDMETHODCALLTYPE d3d12_command_list_ExecuteIndirect(ID3D12GraphicsC
     }
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_AtomicCopyBufferUINT(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_AtomicCopyBufferUINT(ID3D12GraphicsCommandList2 *iface,
         ID3D12Resource *dst_buffer, UINT64 dst_offset,
         ID3D12Resource *src_buffer, UINT64 src_offset,
         UINT dependent_resource_count, ID3D12Resource * const *dependent_resources,
@@ -5267,7 +5446,7 @@ static void STDMETHODCALLTYPE d3d12_command_list_AtomicCopyBufferUINT(ID3D12Grap
             dependent_resource_count, dependent_resources, dependent_sub_resource_ranges);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_AtomicCopyBufferUINT64(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_AtomicCopyBufferUINT64(ID3D12GraphicsCommandList2 *iface,
         ID3D12Resource *dst_buffer, UINT64 dst_offset,
         ID3D12Resource *src_buffer, UINT64 src_offset,
         UINT dependent_resource_count, ID3D12Resource * const *dependent_resources,
@@ -5280,20 +5459,20 @@ static void STDMETHODCALLTYPE d3d12_command_list_AtomicCopyBufferUINT64(ID3D12Gr
             dependent_resource_count, dependent_resources, dependent_sub_resource_ranges);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_OMSetDepthBounds(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_OMSetDepthBounds(ID3D12GraphicsCommandList2 *iface,
         FLOAT min, FLOAT max)
 {
     FIXME("iface %p, min %.8e, max %.8e stub!\n", iface, min, max);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_SetSamplePositions(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_SetSamplePositions(ID3D12GraphicsCommandList2 *iface,
         UINT sample_count, UINT pixel_count, D3D12_SAMPLE_POSITION *sample_positions)
 {
     FIXME("iface %p, sample_count %u, pixel_count %u, sample_positions %p stub!\n",
             iface, sample_count, pixel_count, sample_positions);
 }
 
-static void STDMETHODCALLTYPE d3d12_command_list_ResolveSubresourceRegion(ID3D12GraphicsCommandList1 *iface,
+static void STDMETHODCALLTYPE d3d12_command_list_ResolveSubresourceRegion(ID3D12GraphicsCommandList2 *iface,
         ID3D12Resource *dst_resource, UINT dst_sub_resource_idx, UINT dst_x, UINT dst_y,
         ID3D12Resource *src_resource, UINT src_sub_resource_idx,
         D3D12_RECT *src_rect, DXGI_FORMAT format, D3D12_RESOLVE_MODE mode)
@@ -5305,7 +5484,29 @@ static void STDMETHODCALLTYPE d3d12_command_list_ResolveSubresourceRegion(ID3D12
             src_resource, src_sub_resource_idx, src_rect, format, mode);
 }
 
-static const struct ID3D12GraphicsCommandList1Vtbl d3d12_command_list_vtbl =
+static void STDMETHODCALLTYPE d3d12_command_list_SetViewInstanceMask(ID3D12GraphicsCommandList2 *iface, UINT mask)
+{
+    FIXME("iface %p, mask %#x stub!\n", iface, mask);
+}
+
+static void STDMETHODCALLTYPE d3d12_command_list_WriteBufferImmediate(ID3D12GraphicsCommandList2 *iface,
+        UINT count, const D3D12_WRITEBUFFERIMMEDIATE_PARAMETER *parameters,
+        const D3D12_WRITEBUFFERIMMEDIATE_MODE *modes)
+{
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
+    struct d3d12_resource *resource;
+    unsigned int i;
+
+    FIXME("iface %p, count %u, parameters %p, modes %p stub!\n", iface, count, parameters, modes);
+
+    for (i = 0; i < count; ++i)
+    {
+        resource = vkd3d_gpu_va_allocator_dereference(&list->device->gpu_va_allocator, parameters[i].Dest);
+        d3d12_command_list_track_resource_usage(list, resource);
+    }
+}
+
+static const struct ID3D12GraphicsCommandList2Vtbl d3d12_command_list_vtbl =
 {
     /* IUnknown methods */
     d3d12_command_list_QueryInterface,
@@ -5378,6 +5579,9 @@ static const struct ID3D12GraphicsCommandList1Vtbl d3d12_command_list_vtbl =
     d3d12_command_list_OMSetDepthBounds,
     d3d12_command_list_SetSamplePositions,
     d3d12_command_list_ResolveSubresourceRegion,
+    d3d12_command_list_SetViewInstanceMask,
+    /* ID3D12GraphicsCommandList2 methods */
+    d3d12_command_list_WriteBufferImmediate,
 };
 
 static struct d3d12_command_list *unsafe_impl_from_ID3D12CommandList(ID3D12CommandList *iface)
@@ -5385,7 +5589,7 @@ static struct d3d12_command_list *unsafe_impl_from_ID3D12CommandList(ID3D12Comma
     if (!iface)
         return NULL;
     assert(iface->lpVtbl == (struct ID3D12CommandListVtbl *)&d3d12_command_list_vtbl);
-    return CONTAINING_RECORD(iface, struct d3d12_command_list, ID3D12GraphicsCommandList1_iface);
+    return CONTAINING_RECORD(iface, struct d3d12_command_list, ID3D12GraphicsCommandList2_iface);
 }
 
 static HRESULT d3d12_command_list_init(struct d3d12_command_list *list, struct d3d12_device *device,
@@ -5394,7 +5598,7 @@ static HRESULT d3d12_command_list_init(struct d3d12_command_list *list, struct d
 {
     HRESULT hr;
 
-    list->ID3D12GraphicsCommandList1_iface.lpVtbl = &d3d12_command_list_vtbl;
+    list->ID3D12GraphicsCommandList2_iface.lpVtbl = &d3d12_command_list_vtbl;
     list->refcount = 1;
 
     list->type = type;
diff --git a/libs/vkd3d/device.c b/libs/vkd3d/device.c
index 13ebc70a..757d4ac8 100644
--- a/libs/vkd3d/device.c
+++ b/libs/vkd3d/device.c
@@ -2154,6 +2154,7 @@ static ULONG STDMETHODCALLTYPE d3d12_device_Release(ID3D12Device *iface)
         vkd3d_private_store_destroy(&device->private_store);
 
         vkd3d_cleanup_format_info(device);
+        vkd3d_uav_clear_state_cleanup(&device->uav_clear_state, device);
         vkd3d_destroy_null_resources(&device->null_resources, device);
         vkd3d_gpu_va_allocator_cleanup(&device->gpu_va_allocator);
         vkd3d_render_pass_cache_cleanup(&device->render_pass_cache, device);
@@ -2307,8 +2308,8 @@ static HRESULT STDMETHODCALLTYPE d3d12_device_CreateCommandList(ID3D12Device *if
             initial_pipeline_state, &object)))
         return hr;
 
-    return return_interface(&object->ID3D12GraphicsCommandList1_iface,
-            &IID_ID3D12GraphicsCommandList1, riid, command_list);
+    return return_interface(&object->ID3D12GraphicsCommandList2_iface,
+            &IID_ID3D12GraphicsCommandList2, riid, command_list);
 }
 
 /* Direct3D feature levels restrict which formats can be optionally supported. */
@@ -2398,11 +2399,29 @@ done:
     return S_OK;
 }
 
+bool d3d12_device_is_uma(struct d3d12_device *device, bool *coherent)
+{
+    unsigned int i;
+
+    if (coherent)
+        *coherent = true;
+
+    for (i = 0; i < device->memory_properties.memoryTypeCount; ++i)
+    {
+        if (!(device->memory_properties.memoryTypes[i].propertyFlags & VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT))
+            return false;
+        if (coherent && !(device->memory_properties.memoryTypes[i].propertyFlags
+                & VK_MEMORY_PROPERTY_HOST_COHERENT_BIT))
+            *coherent = false;
+    }
+
+    return true;
+}
+
 static HRESULT STDMETHODCALLTYPE d3d12_device_CheckFeatureSupport(ID3D12Device *iface,
         D3D12_FEATURE feature, void *feature_data, UINT feature_data_size)
 {
     struct d3d12_device *device = impl_from_ID3D12Device(iface);
-    unsigned int i;
 
     TRACE("iface %p, feature %#x, feature_data %p, feature_data_size %u.\n",
             iface, feature, feature_data, feature_data_size);
@@ -2443,6 +2462,7 @@ static HRESULT STDMETHODCALLTYPE d3d12_device_CheckFeatureSupport(ID3D12Device *
         case D3D12_FEATURE_ARCHITECTURE:
         {
             D3D12_FEATURE_DATA_ARCHITECTURE *data = feature_data;
+            bool coherent;
 
             if (feature_data_size != sizeof(*data))
             {
@@ -2459,15 +2479,8 @@ static HRESULT STDMETHODCALLTYPE d3d12_device_CheckFeatureSupport(ID3D12Device *
             WARN("Assuming device does not support tile based rendering.\n");
             data->TileBasedRenderer = FALSE;
 
-            data->UMA = TRUE;
-            data->CacheCoherentUMA = TRUE;
-            for (i = 0; i < device->memory_properties.memoryTypeCount; ++i)
-            {
-                if (!(device->memory_properties.memoryTypes[i].propertyFlags & VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT))
-                    data->UMA = FALSE;
-                if (!(device->memory_properties.memoryTypes[i].propertyFlags & VK_MEMORY_PROPERTY_HOST_COHERENT_BIT))
-                    data->CacheCoherentUMA = FALSE;
-            }
+            data->UMA = d3d12_device_is_uma(device, &coherent);
+            data->CacheCoherentUMA = data->UMA ? coherent : FALSE;
 
             TRACE("Tile based renderer %#x, UMA %#x, cache coherent UMA %#x.\n",
                     data->TileBasedRenderer, data->UMA, data->CacheCoherentUMA);
@@ -2886,10 +2899,8 @@ static D3D12_RESOURCE_ALLOCATION_INFO * STDMETHODCALLTYPE d3d12_device_GetResour
         UINT count, const D3D12_RESOURCE_DESC *resource_descs)
 {
     struct d3d12_device *device = impl_from_ID3D12Device(iface);
-    const struct vkd3d_format *format;
     const D3D12_RESOURCE_DESC *desc;
     uint64_t requested_alignment;
-    uint64_t estimated_size;
 
     TRACE("iface %p, info %p, visible_mask 0x%08x, count %u, resource_descs %p.\n",
             iface, info, visible_mask, count, resource_descs);
@@ -2907,15 +2918,12 @@ static D3D12_RESOURCE_ALLOCATION_INFO * STDMETHODCALLTYPE d3d12_device_GetResour
 
     desc = &resource_descs[0];
 
-    if (FAILED(d3d12_resource_validate_desc(desc)))
+    if (FAILED(d3d12_resource_validate_desc(desc, device)))
     {
         WARN("Invalid resource desc.\n");
         goto invalid;
     }
 
-    requested_alignment = desc->Alignment
-            ? desc->Alignment : D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT;
-
     if (desc->Dimension == D3D12_RESOURCE_DIMENSION_BUFFER)
     {
         info->SizeInBytes = desc->Width;
@@ -2929,27 +2937,9 @@ static D3D12_RESOURCE_ALLOCATION_INFO * STDMETHODCALLTYPE d3d12_device_GetResour
             goto invalid;
         }
 
+        requested_alignment = desc->Alignment
+                ? desc->Alignment : D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT;
         info->Alignment = max(info->Alignment, requested_alignment);
-
-        if (info->Alignment < D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT)
-        {
-            if (!(format = vkd3d_format_from_d3d12_resource_desc(device, desc, 0)))
-            {
-                WARN("Invalid format %#x.\n", desc->Format);
-                goto invalid;
-            }
-
-            estimated_size = desc->Width * desc->Height * desc->DepthOrArraySize * format->byte_count;
-            if (estimated_size > D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT)
-                info->Alignment = max(info->Alignment, D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT);
-        }
-    }
-
-    if (desc->Alignment % info->Alignment)
-    {
-        WARN("Invalid resource alignment %#"PRIx64" (required %#"PRIx64").\n",
-                desc->Alignment, info->Alignment);
-        goto invalid;
     }
 
     info->SizeInBytes = align(info->SizeInBytes, info->Alignment);
@@ -2975,11 +2965,43 @@ invalid:
 static D3D12_HEAP_PROPERTIES * STDMETHODCALLTYPE d3d12_device_GetCustomHeapProperties(ID3D12Device *iface,
         D3D12_HEAP_PROPERTIES *heap_properties, UINT node_mask, D3D12_HEAP_TYPE heap_type)
 {
-    FIXME("iface %p, heap_properties %p, node_mask 0x%08x, heap_type %#x stub!\n",
+    struct d3d12_device *device = impl_from_ID3D12Device(iface);
+    bool coherent;
+
+    TRACE("iface %p, heap_properties %p, node_mask 0x%08x, heap_type %#x.\n",
             iface, heap_properties, node_mask, heap_type);
 
     debug_ignored_node_mask(node_mask);
 
+    heap_properties->Type = D3D12_HEAP_TYPE_CUSTOM;
+
+    switch (heap_type)
+    {
+        case D3D12_HEAP_TYPE_DEFAULT:
+            heap_properties->CPUPageProperty = D3D12_CPU_PAGE_PROPERTY_NOT_AVAILABLE;
+            heap_properties->MemoryPoolPreference = d3d12_device_is_uma(device, NULL)
+                    ?  D3D12_MEMORY_POOL_L0 : D3D12_MEMORY_POOL_L1;
+            break;
+
+        case D3D12_HEAP_TYPE_UPLOAD:
+            heap_properties->CPUPageProperty = d3d12_device_is_uma(device, &coherent) && coherent
+                    ? D3D12_CPU_PAGE_PROPERTY_WRITE_BACK : D3D12_CPU_PAGE_PROPERTY_WRITE_COMBINE;
+            heap_properties->MemoryPoolPreference = D3D12_MEMORY_POOL_L0;
+            break;
+
+        case D3D12_HEAP_TYPE_READBACK:
+            heap_properties->CPUPageProperty = D3D12_CPU_PAGE_PROPERTY_WRITE_BACK;
+            heap_properties->MemoryPoolPreference = D3D12_MEMORY_POOL_L0;
+            break;
+
+        default:
+            FIXME("Unhandled heap type %#x.\n", heap_type);
+            break;
+    };
+
+    heap_properties->CreationNodeMask = 1;
+    heap_properties->VisibleNodeMask = 1;
+
     return heap_properties;
 }
 
@@ -3181,7 +3203,7 @@ static void STDMETHODCALLTYPE d3d12_device_GetCopyableFootprints(ID3D12Device *i
         return;
     }
 
-    if (FAILED(d3d12_resource_validate_desc(desc)))
+    if (FAILED(d3d12_resource_validate_desc(desc, device)))
     {
         WARN("Invalid resource desc.\n");
         return;
@@ -3196,13 +3218,6 @@ static void STDMETHODCALLTYPE d3d12_device_GetCopyableFootprints(ID3D12Device *i
         return;
     }
 
-    if (align(desc->Width, format->block_width) != desc->Width
-            || align(desc->Height, format->block_height) != desc->Height)
-    {
-        WARN("Resource size (%"PRIu64"x%u) not aligned to format block size.\n", desc->Width, desc->Height);
-        return;
-    }
-
     offset = 0;
     total = 0;
     for (i = 0; i < sub_resource_count; ++i)
@@ -3403,6 +3418,9 @@ static HRESULT d3d12_device_init(struct d3d12_device *device,
     if (FAILED(hr = vkd3d_init_null_resources(&device->null_resources, device)))
         goto out_cleanup_format_info;
 
+    if (FAILED(hr = vkd3d_uav_clear_state_init(&device->uav_clear_state, device)))
+        goto out_destroy_null_resources;
+
     vkd3d_render_pass_cache_init(&device->render_pass_cache);
     vkd3d_gpu_va_allocator_init(&device->gpu_va_allocator);
 
@@ -3414,6 +3432,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/resource.c b/libs/vkd3d/resource.c
index d6e999eb..f40d9866 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -34,15 +34,6 @@ static inline bool is_cpu_accessible_heap(const D3D12_HEAP_PROPERTIES *propertie
     return true;
 }
 
-static bool is_numa_device(struct d3d12_device *device)
-{
-    unsigned int i;
-    for (i = 0; i < device->memory_properties.memoryTypeCount; ++i)
-        if (!(device->memory_properties.memoryTypes[i].propertyFlags & VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT))
-            return true;
-    return false;
-}
-
 static HRESULT vkd3d_select_memory_type(struct d3d12_device *device, uint32_t memory_type_mask,
         const D3D12_HEAP_PROPERTIES *heap_properties, D3D12_HEAP_FLAGS heap_flags, unsigned int *type_index)
 {
@@ -69,7 +60,7 @@ static HRESULT vkd3d_select_memory_type(struct d3d12_device *device, uint32_t me
         case D3D12_HEAP_TYPE_CUSTOM:
             if (heap_properties->MemoryPoolPreference == D3D12_MEMORY_POOL_UNKNOWN
                     || (heap_properties->MemoryPoolPreference == D3D12_MEMORY_POOL_L1
-                    && (is_cpu_accessible_heap(heap_properties) || !is_numa_device(device))))
+                    && (is_cpu_accessible_heap(heap_properties) || d3d12_device_is_uma(device, NULL))))
             {
                 WARN("Invalid memory pool preference.\n");
                 return E_INVALIDARG;
@@ -873,13 +864,17 @@ 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)
+    {
+        /* Format compatibility rules are more relaxed for UAVs. */
+        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;
@@ -1006,7 +1001,7 @@ HRESULT vkd3d_get_image_allocation_info(struct d3d12_device *device,
     HRESULT hr;
 
     assert(desc->Dimension != D3D12_RESOURCE_DIMENSION_BUFFER);
-    assert(d3d12_resource_validate_desc(desc) == S_OK);
+    assert(d3d12_resource_validate_desc(desc, device) == S_OK);
 
     if (!desc->MipLevels)
     {
@@ -1527,6 +1522,8 @@ static HRESULT STDMETHODCALLTYPE d3d12_resource_GetHeapProperties(ID3D12Resource
         {
             memset(heap_properties, 0, sizeof(*heap_properties));
             heap_properties->Type = D3D12_HEAP_TYPE_DEFAULT;
+            heap_properties->CreationNodeMask = 1;
+            heap_properties->VisibleNodeMask = 1;
         }
         if (flags)
             *flags = D3D12_HEAP_FLAG_NONE;
@@ -1594,8 +1591,68 @@ static void d3d12_validate_resource_flags(D3D12_RESOURCE_FLAGS flags)
         FIXME("Ignoring D3D12_RESOURCE_FLAG_ALLOW_CROSS_ADAPTER.\n");
 }
 
-HRESULT d3d12_resource_validate_desc(const D3D12_RESOURCE_DESC *desc)
+static bool d3d12_resource_validate_texture_format(const D3D12_RESOURCE_DESC *desc,
+        const struct vkd3d_format *format)
 {
+    if (!vkd3d_format_is_compressed(format))
+        return true;
+
+    if (desc->Dimension == D3D12_RESOURCE_DIMENSION_TEXTURE1D && format->block_height > 1)
+    {
+        WARN("1D texture with a format block height > 1.\n");
+        return false;
+    }
+
+    if (align(desc->Width, format->block_width) != desc->Width
+            || align(desc->Height, format->block_height) != desc->Height)
+    {
+        WARN("Invalid size %"PRIu64"x%u for block compressed format %#x.\n",
+                desc->Width, desc->Height, desc->Format);
+        return false;
+    }
+
+    return true;
+}
+
+static bool d3d12_resource_validate_texture_alignment(const D3D12_RESOURCE_DESC *desc,
+        const struct vkd3d_format *format)
+{
+    uint64_t estimated_size;
+
+    if (!desc->Alignment)
+        return true;
+
+    if (desc->Alignment != D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT
+            && desc->Alignment != D3D12_SMALL_RESOURCE_PLACEMENT_ALIGNMENT
+            && (desc->SampleDesc.Count == 1 || desc->Alignment != D3D12_DEFAULT_MSAA_RESOURCE_PLACEMENT_ALIGNMENT))
+    {
+        WARN("Invalid resource alignment %#"PRIx64".\n", desc->Alignment);
+        return false;
+    }
+
+    if (desc->Alignment < D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT)
+    {
+        /* Windows uses the slice size to determine small alignment eligibility. DepthOrArraySize is ignored. */
+        estimated_size = desc->Width * desc->Height * format->byte_count * format->block_byte_count
+                / (format->block_width * format->block_height);
+        if (estimated_size > D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT)
+        {
+            WARN("Invalid resource alignment %#"PRIx64" (required %#x).\n",
+                    desc->Alignment, D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT);
+            return false;
+        }
+    }
+
+    /* The size check for MSAA textures with D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT is probably
+     * not important. The 4MB requirement is no longer universal and Vulkan has no such requirement. */
+
+    return true;
+}
+
+HRESULT d3d12_resource_validate_desc(const D3D12_RESOURCE_DESC *desc, struct d3d12_device *device)
+{
+    const struct vkd3d_format *format;
+
     switch (desc->Dimension)
     {
         case D3D12_RESOURCE_DIMENSION_BUFFER:
@@ -1621,10 +1678,18 @@ HRESULT d3d12_resource_validate_desc(const D3D12_RESOURCE_DESC *desc)
                 WARN("1D texture with a height of %u.\n", desc->Height);
                 return E_INVALIDARG;
             }
-            break;
-
+            /* Fall through. */
         case D3D12_RESOURCE_DIMENSION_TEXTURE2D:
         case D3D12_RESOURCE_DIMENSION_TEXTURE3D:
+            if (!(format = vkd3d_format_from_d3d12_resource_desc(device, desc, 0)))
+            {
+                WARN("Invalid format %#x.\n", desc->Format);
+                return E_INVALIDARG;
+            }
+
+            if (!d3d12_resource_validate_texture_format(desc, format)
+                    || !d3d12_resource_validate_texture_alignment(desc, format))
+                return E_INVALIDARG;
             break;
 
         default:
@@ -1634,8 +1699,6 @@ HRESULT d3d12_resource_validate_desc(const D3D12_RESOURCE_DESC *desc)
 
     d3d12_validate_resource_flags(desc->Flags);
 
-    /* FIXME: Validate alignment for textures. */
-
     return S_OK;
 }
 
@@ -1706,7 +1769,7 @@ static HRESULT d3d12_resource_init(struct d3d12_resource *resource, struct d3d12
     resource->gpu_address = 0;
     resource->flags = 0;
 
-    if (FAILED(hr = d3d12_resource_validate_desc(&resource->desc)))
+    if (FAILED(hr = d3d12_resource_validate_desc(&resource->desc, device)))
         return hr;
 
     switch (desc->Dimension)
@@ -1985,13 +2048,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;
@@ -2002,40 +2066,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)
-    {
-        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)
+    switch (view->type)
     {
-        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,
@@ -2062,7 +2123,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)
@@ -2157,8 +2218,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,
-        VkBuffer vk_buffer, const struct vkd3d_format *format,
+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)
 {
     const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
@@ -2168,13 +2228,16 @@ 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;
     }
 
     object->u.vk_buffer_view = vk_view;
+    object->format = format;
+    object->info.buffer.offset = offset;
+    object->info.buffer.size = size;
     *view = object;
     return true;
 }
@@ -2354,18 +2417,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)
 {
@@ -2413,9 +2464,8 @@ 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,
-        VkImage vk_image, const struct vkd3d_texture_view_desc *desc,
-        struct vkd3d_view **view)
+bool vkd3d_create_texture_view(struct d3d12_device *device, VkImage vk_image,
+        const struct vkd3d_texture_view_desc *desc, struct vkd3d_view **view)
 {
     const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
     const struct vkd3d_format *format = desc->format;
@@ -2444,13 +2494,18 @@ 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;
     }
 
     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;
 }
@@ -2821,16 +2876,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,
@@ -2884,11 +2929,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,
@@ -3033,7 +3073,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/state.c b/libs/vkd3d/state.c
index a321fa47..e1f7da98 100644
--- a/libs/vkd3d/state.c
+++ b/libs/vkd3d/state.c
@@ -18,6 +18,7 @@
  */
 
 #include "vkd3d_private.h"
+#include "vkd3d_shaders.h"
 
 /* ID3D12RootSignature */
 static inline struct d3d12_root_signature *impl_from_ID3D12RootSignature(ID3D12RootSignature *iface)
@@ -1384,6 +1385,37 @@ static HRESULT create_shader_stage(struct d3d12_device *device,
     return S_OK;
 }
 
+static HRESULT vkd3d_create_compute_pipeline(struct d3d12_device *device,
+        const D3D12_SHADER_BYTECODE *code, const struct vkd3d_shader_interface_info *shader_interface,
+        VkPipelineLayout vk_pipeline_layout, VkPipeline *vk_pipeline)
+{
+    const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+    VkComputePipelineCreateInfo pipeline_info;
+    VkResult vr;
+    HRESULT hr;
+
+    pipeline_info.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;
+    pipeline_info.pNext = NULL;
+    pipeline_info.flags = 0;
+    if (FAILED(hr = create_shader_stage(device, &pipeline_info.stage,
+            VK_SHADER_STAGE_COMPUTE_BIT, code, shader_interface, NULL)))
+        return hr;
+    pipeline_info.layout = vk_pipeline_layout;
+    pipeline_info.basePipelineHandle = VK_NULL_HANDLE;
+    pipeline_info.basePipelineIndex = -1;
+
+    vr = VK_CALL(vkCreateComputePipelines(device->vk_device,
+            VK_NULL_HANDLE, 1, &pipeline_info, NULL, vk_pipeline));
+    VK_CALL(vkDestroyShaderModule(device->vk_device, pipeline_info.stage.module, NULL));
+    if (vr < 0)
+    {
+        WARN("Failed to create Vulkan compute pipeline, hr %#x.", hr);
+        return hresult_from_vk_result(vr);
+    }
+
+    return S_OK;
+}
+
 static HRESULT d3d12_pipeline_state_init_compute_uav_counters(struct d3d12_pipeline_state *state,
         struct d3d12_device *device, const struct d3d12_root_signature *root_signature,
         const struct vkd3d_shader_scan_info *shader_info)
@@ -1470,10 +1502,9 @@ static HRESULT d3d12_pipeline_state_init_compute(struct d3d12_pipeline_state *st
     const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
     struct vkd3d_shader_interface_info shader_interface;
     const struct d3d12_root_signature *root_signature;
-    VkComputePipelineCreateInfo pipeline_info;
     struct vkd3d_shader_scan_info shader_info;
+    VkPipelineLayout vk_pipeline_layout;
     struct vkd3d_shader_code dxbc;
-    VkResult vr;
     HRESULT hr;
     int ret;
 
@@ -1519,36 +1550,18 @@ static HRESULT d3d12_pipeline_state_init_compute(struct d3d12_pipeline_state *st
     shader_interface.uav_counters = state->uav_counters;
     shader_interface.uav_counter_count = vkd3d_popcount(state->uav_counter_mask);
 
-    pipeline_info.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;
-    pipeline_info.pNext = NULL;
-    pipeline_info.flags = 0;
-    if (FAILED(hr = create_shader_stage(device, &pipeline_info.stage,
-            VK_SHADER_STAGE_COMPUTE_BIT, &desc->CS, &shader_interface, NULL)))
-    {
-        if (state->vk_set_layout)
-            VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, state->vk_set_layout, NULL));
-        if (state->vk_pipeline_layout)
-            VK_CALL(vkDestroyPipelineLayout(device->vk_device, state->vk_pipeline_layout, NULL));
-        vkd3d_free(state->uav_counters);
-        return hr;
-    }
-    pipeline_info.layout = state->vk_pipeline_layout
+    vk_pipeline_layout = state->vk_pipeline_layout
             ? state->vk_pipeline_layout : root_signature->vk_pipeline_layout;
-    pipeline_info.basePipelineHandle = VK_NULL_HANDLE;
-    pipeline_info.basePipelineIndex = -1;
-
-    vr = VK_CALL(vkCreateComputePipelines(device->vk_device, VK_NULL_HANDLE,
-            1, &pipeline_info, NULL, &state->u.compute.vk_pipeline));
-    VK_CALL(vkDestroyShaderModule(device->vk_device, pipeline_info.stage.module, NULL));
-    if (vr)
+    if (FAILED(hr = vkd3d_create_compute_pipeline(device, &desc->CS, &shader_interface,
+            vk_pipeline_layout, &state->u.compute.vk_pipeline)))
     {
-        WARN("Failed to create Vulkan compute pipeline, vr %d.\n", vr);
+        WARN("Failed to create Vulkan compute pipeline, hr %#x.\n", hr);
         if (state->vk_set_layout)
             VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, state->vk_set_layout, NULL));
         if (state->vk_pipeline_layout)
             VK_CALL(vkDestroyPipelineLayout(device->vk_device, state->vk_pipeline_layout, NULL));
         vkd3d_free(state->uav_counters);
-        return hresult_from_vk_result(vr);
+        return hr;
     }
 
     if (FAILED(hr = vkd3d_private_store_init(&state->private_store)))
@@ -2802,3 +2815,162 @@ VkPipeline d3d12_pipeline_state_get_or_create_pipeline(struct d3d12_pipeline_sta
         ERR("Could not get the pipeline compiled by other thread from the cache.\n");
     return vk_pipeline;
 }
+
+static void vkd3d_uav_clear_pipelines_cleanup(struct vkd3d_uav_clear_pipelines *pipelines,
+        struct d3d12_device *device)
+{
+    const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+
+    VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_3d, NULL));
+    VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_2d_array, NULL));
+    VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_2d, NULL));
+    VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_1d_array, NULL));
+    VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_1d, NULL));
+    VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->buffer, NULL));
+}
+
+void vkd3d_uav_clear_state_cleanup(struct vkd3d_uav_clear_state *state, struct d3d12_device *device)
+{
+    const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+
+    vkd3d_uav_clear_pipelines_cleanup(&state->pipelines_uint, device);
+    vkd3d_uav_clear_pipelines_cleanup(&state->pipelines_float, device);
+
+    VK_CALL(vkDestroyPipelineLayout(device->vk_device, state->vk_pipeline_layout_image, NULL));
+    VK_CALL(vkDestroyPipelineLayout(device->vk_device, state->vk_pipeline_layout_buffer, NULL));
+
+    VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, state->vk_set_layout_image, NULL));
+    VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, state->vk_set_layout_buffer, NULL));
+}
+
+HRESULT vkd3d_uav_clear_state_init(struct vkd3d_uav_clear_state *state, struct d3d12_device *device)
+{
+    struct vkd3d_shader_push_constant_buffer push_constant;
+    struct vkd3d_shader_interface_info shader_interface;
+    struct vkd3d_shader_resource_binding binding;
+    VkDescriptorSetLayoutBinding set_binding;
+    VkPushConstantRange push_constant_range;
+    unsigned int i;
+    HRESULT hr;
+
+    const struct
+    {
+        VkDescriptorSetLayout *set_layout;
+        VkPipelineLayout *pipeline_layout;
+        VkDescriptorType descriptor_type;
+    }
+    set_layouts[] =
+    {
+        {&state->vk_set_layout_buffer, &state->vk_pipeline_layout_buffer, VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER},
+        {&state->vk_set_layout_image,  &state->vk_pipeline_layout_image, VK_DESCRIPTOR_TYPE_STORAGE_IMAGE},
+    };
+
+    const struct
+    {
+        VkPipeline *pipeline;
+        VkPipelineLayout *pipeline_layout;
+        D3D12_SHADER_BYTECODE code;
+    }
+    pipelines[] =
+    {
+#define SHADER_CODE(name) {name, sizeof(name)}
+        {&state->pipelines_float.buffer, &state->vk_pipeline_layout_buffer,
+                SHADER_CODE(cs_uav_clear_buffer_float_code)},
+        {&state->pipelines_float.image_1d, &state->vk_pipeline_layout_image,
+                SHADER_CODE(cs_uav_clear_1d_float_code)},
+        {&state->pipelines_float.image_1d_array, &state->vk_pipeline_layout_image,
+                SHADER_CODE(cs_uav_clear_1d_array_float_code)},
+        {&state->pipelines_float.image_2d, &state->vk_pipeline_layout_image,
+                SHADER_CODE(cs_uav_clear_2d_float_code)},
+        {&state->pipelines_float.image_2d_array, &state->vk_pipeline_layout_image,
+                SHADER_CODE(cs_uav_clear_2d_array_float_code)},
+        {&state->pipelines_float.image_3d, &state->vk_pipeline_layout_image,
+                SHADER_CODE(cs_uav_clear_3d_float_code)},
+
+        {&state->pipelines_uint.buffer, &state->vk_pipeline_layout_buffer,
+                SHADER_CODE(cs_uav_clear_buffer_uint_code)},
+        {&state->pipelines_uint.image_1d, &state->vk_pipeline_layout_image,
+                SHADER_CODE(cs_uav_clear_1d_uint_code)},
+        {&state->pipelines_uint.image_1d_array, &state->vk_pipeline_layout_image,
+                SHADER_CODE(cs_uav_clear_1d_array_uint_code)},
+        {&state->pipelines_uint.image_2d, &state->vk_pipeline_layout_image,
+                SHADER_CODE(cs_uav_clear_2d_uint_code)},
+        {&state->pipelines_uint.image_2d_array, &state->vk_pipeline_layout_image,
+                SHADER_CODE(cs_uav_clear_2d_array_uint_code)},
+        {&state->pipelines_uint.image_3d, &state->vk_pipeline_layout_image,
+                SHADER_CODE(cs_uav_clear_3d_uint_code)},
+#undef SHADER_CODE
+    };
+
+    memset(state, 0, sizeof(*state));
+
+    set_binding.binding = 0;
+    set_binding.descriptorCount = 1;
+    set_binding.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
+    set_binding.pImmutableSamplers = NULL;
+
+    binding.type = VKD3D_SHADER_DESCRIPTOR_TYPE_UAV;
+    binding.register_index = 0;
+    binding.shader_visibility = VKD3D_SHADER_VISIBILITY_COMPUTE;
+    binding.binding.set = 0;
+    binding.binding.binding = 0;
+
+    push_constant_range.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
+    push_constant_range.offset = 0;
+    push_constant_range.size = sizeof(struct vkd3d_uav_clear_args);
+
+    push_constant.register_index = 0;
+    push_constant.shader_visibility = VKD3D_SHADER_VISIBILITY_COMPUTE;
+    push_constant.offset = 0;
+    push_constant.size = sizeof(struct vkd3d_uav_clear_args);
+
+    for (i = 0; i < ARRAY_SIZE(set_layouts); ++i)
+    {
+        set_binding.descriptorType = set_layouts[i].descriptor_type;
+
+        if (FAILED(hr = vkd3d_create_descriptor_set_layout(device, 0, 1, &set_binding, set_layouts[i].set_layout)))
+        {
+            ERR("Failed to create descriptor set layout %u, hr %#x.", i, hr);
+            goto fail;
+        }
+
+        if (FAILED(hr = vkd3d_create_pipeline_layout(device, 1, set_layouts[i].set_layout,
+                1, &push_constant_range, set_layouts[i].pipeline_layout)))
+        {
+            ERR("Failed to create pipeline layout %u, hr %#x.", i, hr);
+            goto fail;
+        }
+    }
+
+    shader_interface.type = VKD3D_SHADER_STRUCTURE_TYPE_SHADER_INTERFACE_INFO;
+    shader_interface.next = NULL;
+    shader_interface.bindings = &binding;
+    shader_interface.binding_count = 1;
+    shader_interface.push_constant_buffers = &push_constant;
+    shader_interface.push_constant_buffer_count = 1;
+    shader_interface.combined_samplers = NULL;
+    shader_interface.combined_sampler_count = 0;
+    shader_interface.uav_counters = NULL;
+    shader_interface.uav_counter_count = 0;
+
+    for (i = 0; i < ARRAY_SIZE(pipelines); ++i)
+    {
+        if (pipelines[i].pipeline_layout == &state->vk_pipeline_layout_buffer)
+            binding.flags = VKD3D_SHADER_BINDING_FLAG_BUFFER;
+        else
+            binding.flags = VKD3D_SHADER_BINDING_FLAG_IMAGE;
+
+        if (FAILED(hr = vkd3d_create_compute_pipeline(device, &pipelines[i].code, &shader_interface,
+                *pipelines[i].pipeline_layout, pipelines[i].pipeline)))
+        {
+            ERR("Failed to create compute pipeline %u, hr %#x.", i, hr);
+            goto fail;
+        }
+    }
+
+    return S_OK;
+
+fail:
+    vkd3d_uav_clear_state_cleanup(state, device);
+    return hr;
+}
diff --git a/libs/vkd3d/utils.c b/libs/vkd3d/utils.c
index 1fc0e9d0..7abfd42e 100644
--- a/libs/vkd3d/utils.c
+++ b/libs/vkd3d/utils.c
@@ -451,6 +451,37 @@ 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)
+            continue;
+
+        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)
@@ -828,7 +859,7 @@ HRESULT vkd3d_load_vk_device_procs(struct vkd3d_vk_device_procs *procs,
     return S_OK;
 }
 
-#ifdef _GNU_SOURCE
+#if HAVE_DECL_PROGRAM_INVOCATION_NAME
 
 bool vkd3d_get_program_name(char program_name[PATH_MAX])
 {
@@ -870,7 +901,7 @@ bool vkd3d_get_program_name(char program_name[PATH_MAX])
     return false;
 }
 
-#endif  /* _GNU_SOURCE */
+#endif  /* HAVE_DECL_PROGRAM_INVOCATION_NAME */
 
 static struct vkd3d_private_data *vkd3d_private_store_get_private_data(
         const struct vkd3d_private_store *store, const GUID *tag)
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index 2d62fdaa..0c031d20 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -205,13 +205,13 @@ HRESULT vkd3d_fence_worker_stop(struct vkd3d_fence_worker *worker,
 struct vkd3d_gpu_va_allocation
 {
     D3D12_GPU_VIRTUAL_ADDRESS base;
-    SIZE_T size;
+    size_t size;
     void *ptr;
 };
 
 struct vkd3d_gpu_va_slab
 {
-    SIZE_T size;
+    size_t size;
     void *ptr;
 };
 
@@ -435,7 +435,7 @@ static inline bool d3d12_resource_is_texture(const struct d3d12_resource *resour
 }
 
 bool d3d12_resource_is_cpu_accessible(const struct d3d12_resource *resource) DECLSPEC_HIDDEN;
-HRESULT d3d12_resource_validate_desc(const D3D12_RESOURCE_DESC *desc) DECLSPEC_HIDDEN;
+HRESULT d3d12_resource_validate_desc(const D3D12_RESOURCE_DESC *desc, struct d3d12_device *device) DECLSPEC_HIDDEN;
 
 HRESULT d3d12_committed_resource_create(struct d3d12_device *device,
         const D3D12_HEAP_PROPERTIES *heap_properties, D3D12_HEAP_FLAGS heap_flags,
@@ -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;
@@ -468,11 +476,44 @@ 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;
 void vkd3d_view_incref(struct vkd3d_view *view) DECLSPEC_HIDDEN;
 
+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;
+
 struct d3d12_desc
 {
     uint32_t magic;
@@ -482,22 +523,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)
@@ -906,7 +931,7 @@ struct vkd3d_pipeline_bindings
 /* ID3D12CommandList */
 struct d3d12_command_list
 {
-    ID3D12GraphicsCommandList1 ID3D12GraphicsCommandList1_iface;
+    ID3D12GraphicsCommandList2 ID3D12GraphicsCommandList2_iface;
     LONG refcount;
 
     D3D12_COMMAND_LIST_TYPE type;
@@ -1051,6 +1076,38 @@ struct vkd3d_format_compatibility_list
     VkFormat vk_formats[VKD3D_MAX_COMPATIBLE_FORMAT_COUNT];
 };
 
+struct vkd3d_uav_clear_args
+{
+    VkClearColorValue colour;
+    VkOffset2D offset;
+    VkExtent2D extent;
+};
+
+struct vkd3d_uav_clear_pipelines
+{
+    VkPipeline buffer;
+    VkPipeline image_1d;
+    VkPipeline image_1d_array;
+    VkPipeline image_2d;
+    VkPipeline image_2d_array;
+    VkPipeline image_3d;
+};
+
+struct vkd3d_uav_clear_state
+{
+    VkDescriptorSetLayout vk_set_layout_buffer;
+    VkDescriptorSetLayout vk_set_layout_image;
+
+    VkPipelineLayout vk_pipeline_layout_buffer;
+    VkPipelineLayout vk_pipeline_layout_image;
+
+    struct vkd3d_uav_clear_pipelines pipelines_float;
+    struct vkd3d_uav_clear_pipelines pipelines_uint;
+};
+
+HRESULT vkd3d_uav_clear_state_init(struct vkd3d_uav_clear_state *state, struct d3d12_device *device) DECLSPEC_HIDDEN;
+void vkd3d_uav_clear_state_cleanup(struct vkd3d_uav_clear_state *state, struct d3d12_device *device) DECLSPEC_HIDDEN;
+
 /* ID3D12Device */
 struct d3d12_device
 {
@@ -1096,12 +1153,14 @@ struct d3d12_device
     unsigned int format_compatibility_list_count;
     const struct vkd3d_format_compatibility_list *format_compatibility_lists;
     struct vkd3d_null_resources null_resources;
+    struct vkd3d_uav_clear_state uav_clear_state;
 };
 
 HRESULT d3d12_device_create(struct vkd3d_instance *instance,
         const struct vkd3d_device_create_info *create_info, struct d3d12_device **device) DECLSPEC_HIDDEN;
 struct vkd3d_queue *d3d12_device_get_vkd3d_queue(struct d3d12_device *device,
         D3D12_COMMAND_LIST_TYPE type) DECLSPEC_HIDDEN;
+bool d3d12_device_is_uma(struct d3d12_device *device, bool *coherent) DECLSPEC_HIDDEN;
 void d3d12_device_mark_as_removed(struct d3d12_device *device, HRESULT reason,
         const char *message, ...) VKD3D_PRINTF_FUNC(3, 4) DECLSPEC_HIDDEN;
 struct d3d12_device *unsafe_impl_from_ID3D12Device(ID3D12Device *iface) DECLSPEC_HIDDEN;
@@ -1183,6 +1242,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;
@@ -1228,6 +1289,11 @@ static inline unsigned int d3d12_resource_desc_get_sub_resource_count(const D3D1
     return d3d12_resource_desc_get_layer_count(desc) * desc->MipLevels;
 }
 
+static inline unsigned int vkd3d_compute_workgroup_count(unsigned int thread_count, unsigned int workgroup_size)
+{
+    return (thread_count + workgroup_size - 1) / workgroup_size;
+}
+
 VkCompareOp vk_compare_op_from_d3d12(D3D12_COMPARISON_FUNC op) DECLSPEC_HIDDEN;
 VkSampleCountFlagBits vk_samples_from_dxgi_sample_desc(const DXGI_SAMPLE_DESC *desc) DECLSPEC_HIDDEN;
 VkSampleCountFlagBits vk_samples_from_sample_count(unsigned int sample_count) DECLSPEC_HIDDEN;
diff --git a/libs/vkd3d/vkd3d_shaders.h b/libs/vkd3d/vkd3d_shaders.h
new file mode 100644
index 00000000..b2a90cdb
--- /dev/null
+++ b/libs/vkd3d/vkd3d_shaders.h
@@ -0,0 +1,388 @@
+/*
+ * Copyright 2019 Philip Rebohle
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2.1 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library; if not, write to the Free Software
+ * Foundation, Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301, USA
+ */
+
+#ifndef __VKD3D_SHADERS_H
+#define __VKD3D_SHADERS_H
+
+static const uint32_t cs_uav_clear_buffer_float_code[] =
+{
+#if 0
+    RWBuffer<float4> dst;
+
+    struct
+    {
+        float4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(128, 1, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (thread_id.x < u_info.dst_extent.x)
+            dst[u_info.dst_offset.x + thread_id.x] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0xe114ba61, 0xff6a0d0b, 0x7b25c8f4, 0xfcf7cf22, 0x00000001, 0x0000010c, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000b8, 0x00050050, 0x0000002e, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400089c, 0x0011e000, 0x00000000, 0x00005555,
+    0x0200005f, 0x00020012, 0x02000068, 0x00000001, 0x0400009b, 0x00000080, 0x00000001, 0x00000001,
+    0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+    0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+    0x00000001, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100006, 0x00000000, 0x00208e46, 0x00000000,
+    0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_buffer_uint_code[] =
+{
+#if 0
+    RWBuffer<uint4> dst;
+
+    struct
+    {
+        uint4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(128, 1, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (thread_id.x < u_info.dst_extent.x)
+            dst[u_info.dst_offset.x + thread_id.x] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0x3afd0cfd, 0x5145c166, 0x5b9f76b8, 0xa73775cd, 0x00000001, 0x0000010c, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000b8, 0x00050050, 0x0000002e, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400089c, 0x0011e000, 0x00000000, 0x00004444,
+    0x0200005f, 0x00020012, 0x02000068, 0x00000001, 0x0400009b, 0x00000080, 0x00000001, 0x00000001,
+    0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+    0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+    0x00000001, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100006, 0x00000000, 0x00208e46, 0x00000000,
+    0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_1d_array_float_code[] =
+{
+#if 0
+    RWTexture1DArray<float4> dst;
+
+    struct
+    {
+        float4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(64, 1, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (thread_id.x < u_info.dst_extent.x)
+            dst[int2(u_info.dst_offset.x + thread_id.x, thread_id.y)] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0x3d73bc2d, 0x2b635f3d, 0x6bf98e92, 0xbe0aa5d9, 0x00000001, 0x0000011c, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000c8, 0x00050050, 0x00000032, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400389c, 0x0011e000, 0x00000000, 0x00005555,
+    0x0200005f, 0x00020032, 0x02000068, 0x00000001, 0x0400009b, 0x00000040, 0x00000001, 0x00000001,
+    0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+    0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+    0x00000001, 0x04000036, 0x001000e2, 0x00000000, 0x00020556, 0x080000a4, 0x0011e0f2, 0x00000000,
+    0x00100e46, 0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_1d_array_uint_code[] =
+{
+#if 0
+    RWTexture1DArray<uint4> dst;
+
+    struct
+    {
+        uint4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(64, 1, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (thread_id.x < u_info.dst_extent.x)
+            dst[int2(u_info.dst_offset.x + thread_id.x, thread_id.y)] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0x2f0ca457, 0x72068b34, 0xd9dadc2b, 0xd3178c3e, 0x00000001, 0x0000011c, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000c8, 0x00050050, 0x00000032, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400389c, 0x0011e000, 0x00000000, 0x00004444,
+    0x0200005f, 0x00020032, 0x02000068, 0x00000001, 0x0400009b, 0x00000040, 0x00000001, 0x00000001,
+    0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+    0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+    0x00000001, 0x04000036, 0x001000e2, 0x00000000, 0x00020556, 0x080000a4, 0x0011e0f2, 0x00000000,
+    0x00100e46, 0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_1d_float_code[] =
+{
+#if 0
+    RWTexture1D<float4> dst;
+
+    struct
+    {
+        float4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(64, 1, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (thread_id.x < u_info.dst_extent.x)
+            dst[u_info.dst_offset.x + thread_id.x] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0x05266503, 0x4b97006f, 0x01a5cc63, 0xe617d0a1, 0x00000001, 0x0000010c, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000b8, 0x00050050, 0x0000002e, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400109c, 0x0011e000, 0x00000000, 0x00005555,
+    0x0200005f, 0x00020012, 0x02000068, 0x00000001, 0x0400009b, 0x00000040, 0x00000001, 0x00000001,
+    0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+    0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+    0x00000001, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100006, 0x00000000, 0x00208e46, 0x00000000,
+    0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_1d_uint_code[] =
+{
+#if 0
+    RWTexture1D<uint4> dst;
+
+    struct
+    {
+        uint4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(64, 1, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (thread_id.x < u_info.dst_extent.x)
+            dst[u_info.dst_offset.x + thread_id.x] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0x19d5c8f2, 0x3ca4ac24, 0x9e258499, 0xf0463fd6, 0x00000001, 0x0000010c, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000b8, 0x00050050, 0x0000002e, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400109c, 0x0011e000, 0x00000000, 0x00004444,
+    0x0200005f, 0x00020012, 0x02000068, 0x00000001, 0x0400009b, 0x00000040, 0x00000001, 0x00000001,
+    0x07000022, 0x00100012, 0x00000000, 0x0002000a, 0x0020802a, 0x00000000, 0x00000001, 0x0304001f,
+    0x0010000a, 0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x0002000a, 0x0020800a, 0x00000000,
+    0x00000001, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100006, 0x00000000, 0x00208e46, 0x00000000,
+    0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_2d_array_float_code[] =
+{
+#if 0
+    RWTexture2DArray<float4> dst;
+
+    struct
+    {
+        float4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(8, 8, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (all(thread_id.xy < u_info.dst_extent.xy))
+            dst[int3(u_info.dst_offset.xy + thread_id.xy, thread_id.z)] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0x924d2d2c, 0xb9166376, 0x99f83871, 0x8ef65025, 0x00000001, 0x00000138, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000e4, 0x00050050, 0x00000039, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400409c, 0x0011e000, 0x00000000, 0x00005555,
+    0x0200005f, 0x00020072, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+    0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+    0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+    0x00000000, 0x0700001e, 0x00100032, 0x00000000, 0x00020046, 0x00208046, 0x00000000, 0x00000001,
+    0x04000036, 0x001000c2, 0x00000000, 0x00020aa6, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46,
+    0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_2d_array_uint_code[] =
+{
+#if 0
+    RWTexture2DArray<uint4> dst;
+
+    struct
+    {
+        uint4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(8, 8, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (all(thread_id.xy < u_info.dst_extent.xy))
+            dst[int3(u_info.dst_offset.xy + thread_id.xy, thread_id.z)] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0xa92219d4, 0xa2c5e47d, 0x0d308500, 0xf32197b4, 0x00000001, 0x00000138, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000e4, 0x00050050, 0x00000039, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400409c, 0x0011e000, 0x00000000, 0x00004444,
+    0x0200005f, 0x00020072, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+    0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+    0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+    0x00000000, 0x0700001e, 0x00100032, 0x00000000, 0x00020046, 0x00208046, 0x00000000, 0x00000001,
+    0x04000036, 0x001000c2, 0x00000000, 0x00020aa6, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46,
+    0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_2d_float_code[] =
+{
+#if 0
+    RWTexture2D<float4> dst;
+
+    struct
+    {
+        float4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(8, 8, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (all(thread_id.xy < u_info.dst_extent.xy))
+            dst[u_info.dst_offset.xy + thread_id.xy] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0x6e735b3f, 0x7348c4fa, 0xb3634e42, 0x50e2d99b, 0x00000001, 0x00000128, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000d4, 0x00050050, 0x00000035, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400189c, 0x0011e000, 0x00000000, 0x00005555,
+    0x0200005f, 0x00020032, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+    0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+    0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+    0x00000000, 0x0700001e, 0x001000f2, 0x00000000, 0x00020546, 0x00208546, 0x00000000, 0x00000001,
+    0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46, 0x00000000, 0x00208e46, 0x00000000, 0x00000000,
+    0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_2d_uint_code[] =
+{
+#if 0
+    RWTexture2D<uint4> dst;
+
+    struct
+    {
+        uint4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(8, 8, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (all(thread_id.xy < u_info.dst_extent.xy))
+            dst[u_info.dst_offset.xy + thread_id.xy] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0xf01db5dd, 0xc7dc5e55, 0xb017c1a8, 0x55abd52d, 0x00000001, 0x00000128, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000d4, 0x00050050, 0x00000035, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400189c, 0x0011e000, 0x00000000, 0x00004444,
+    0x0200005f, 0x00020032, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+    0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+    0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+    0x00000000, 0x0700001e, 0x001000f2, 0x00000000, 0x00020546, 0x00208546, 0x00000000, 0x00000001,
+    0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46, 0x00000000, 0x00208e46, 0x00000000, 0x00000000,
+    0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_3d_float_code[] =
+{
+#if 0
+    RWTexture3D<float4> dst;
+
+    struct
+    {
+        float4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(8, 8, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (all(thread_id.xy < u_info.dst_extent.xy))
+            dst[int3(u_info.dst_offset.xy, 0) + thread_id.xyz] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0x5d8f36a0, 0x30fa86a5, 0xfec7f2ef, 0xdfd76cbb, 0x00000001, 0x00000138, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000e4, 0x00050050, 0x00000039, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400289c, 0x0011e000, 0x00000000, 0x00005555,
+    0x0200005f, 0x00020072, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+    0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+    0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+    0x00000000, 0x0700001e, 0x00100032, 0x00000000, 0x00020046, 0x00208046, 0x00000000, 0x00000001,
+    0x04000036, 0x001000c2, 0x00000000, 0x00020aa6, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46,
+    0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+static const uint32_t cs_uav_clear_3d_uint_code[] =
+{
+#if 0
+    RWTexture3D<uint4> dst;
+
+    struct
+    {
+        uint4 clear_value;
+        int2 dst_offset;
+        int2 dst_extent;
+    } u_info;
+
+    [numthreads(8, 8, 1)]
+    void main(int3 thread_id : SV_DispatchThreadID)
+    {
+        if (all(thread_id.xy < u_info.dst_extent.xy))
+            dst[int3(u_info.dst_offset.xy, 0) + thread_id.xyz] = u_info.clear_value;
+    }
+#endif
+    0x43425844, 0x5b9c95b1, 0xc9bde4e3, 0x9aaff806, 0x24a1d264, 0x00000001, 0x00000138, 0x00000003,
+    0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+    0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000e4, 0x00050050, 0x00000039, 0x0100086a,
+    0x04000059, 0x00208e46, 0x00000000, 0x00000002, 0x0400289c, 0x0011e000, 0x00000000, 0x00004444,
+    0x0200005f, 0x00020072, 0x02000068, 0x00000001, 0x0400009b, 0x00000008, 0x00000008, 0x00000001,
+    0x07000022, 0x00100032, 0x00000000, 0x00020046, 0x00208ae6, 0x00000000, 0x00000001, 0x07000001,
+    0x00100012, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x0304001f, 0x0010000a,
+    0x00000000, 0x0700001e, 0x00100032, 0x00000000, 0x00020046, 0x00208046, 0x00000000, 0x00000001,
+    0x04000036, 0x001000c2, 0x00000000, 0x00020aa6, 0x080000a4, 0x0011e0f2, 0x00000000, 0x00100e46,
+    0x00000000, 0x00208e46, 0x00000000, 0x00000000, 0x01000015, 0x0100003e,
+};
+
+#endif /* __VKD3D_SHADERS_H */
diff --git a/tests/d3d12.c b/tests/d3d12.c
index 858cf991..323ef239 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -1878,6 +1878,49 @@ static void test_create_committed_resource(void)
             &IID_ID3D12Resource, (void **)&resource);
     ok(hr == E_INVALIDARG, "Got unexpected hr %#x.\n", hr);
 
+    heap_properties.Type = D3D12_HEAP_TYPE_DEFAULT;
+    resource_desc.Format = DXGI_FORMAT_BC1_UNORM;
+    hr = ID3D12Device_CreateCommittedResource(device, &heap_properties, D3D12_HEAP_FLAG_NONE,
+            &resource_desc, D3D12_RESOURCE_STATE_COMMON, NULL,
+            &IID_ID3D12Resource, (void **)&resource);
+    ok(hr == S_OK, "Failed to create committed resource, hr %#x.\n", hr);
+    ID3D12Resource_Release(resource);
+
+    resource_desc.Height = 31;
+    hr = ID3D12Device_CreateCommittedResource(device, &heap_properties, D3D12_HEAP_FLAG_NONE,
+            &resource_desc, D3D12_RESOURCE_STATE_COMMON, NULL,
+            &IID_ID3D12Resource, (void **)&resource);
+    ok(hr == E_INVALIDARG, "Got unexpected hr %#x.\n", hr);
+
+    resource_desc.Width = 31;
+    resource_desc.Height = 32;
+    hr = ID3D12Device_CreateCommittedResource(device, &heap_properties, D3D12_HEAP_FLAG_NONE,
+            &resource_desc, D3D12_RESOURCE_STATE_COMMON, NULL,
+            &IID_ID3D12Resource, (void **)&resource);
+    ok(hr == E_INVALIDARG, "Got unexpected hr %#x.\n", hr);
+
+    resource_desc.Width = 30;
+    resource_desc.Height = 30;
+    hr = ID3D12Device_CreateCommittedResource(device, &heap_properties, D3D12_HEAP_FLAG_NONE,
+            &resource_desc, D3D12_RESOURCE_STATE_COMMON, NULL,
+            &IID_ID3D12Resource, (void **)&resource);
+    ok(hr == E_INVALIDARG, "Got unexpected hr %#x.\n", hr);
+
+    resource_desc.Width = 2;
+    resource_desc.Height = 2;
+    hr = ID3D12Device_CreateCommittedResource(device, &heap_properties, D3D12_HEAP_FLAG_NONE,
+            &resource_desc, D3D12_RESOURCE_STATE_COMMON, NULL,
+            &IID_ID3D12Resource, (void **)&resource);
+    ok(hr == E_INVALIDARG, "Got unexpected hr %#x.\n", hr);
+
+    resource_desc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE1D;
+    resource_desc.Width = 32;
+    resource_desc.Height = 1;
+    hr = ID3D12Device_CreateCommittedResource(device, &heap_properties, D3D12_HEAP_FLAG_NONE,
+            &resource_desc, D3D12_RESOURCE_STATE_COMMON, NULL,
+            &IID_ID3D12Resource, (void **)&resource);
+    ok(hr == E_INVALIDARG, "Got unexpected hr %#x.\n", hr);
+
     heap_properties.Type = D3D12_HEAP_TYPE_UPLOAD;
 
     resource_desc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER;
@@ -1991,6 +2034,7 @@ static void test_create_committed_resource(void)
 
 static void test_create_heap(void)
 {
+    D3D12_FEATURE_DATA_ARCHITECTURE architecture;
     D3D12_FEATURE_DATA_D3D12_OPTIONS options;
     D3D12_HEAP_DESC desc, result_desc;
     ID3D12Device *device, *tmp_device;
@@ -2137,6 +2181,53 @@ static void test_create_heap(void)
     refcount = ID3D12Heap_Release(heap);
     ok(!refcount, "ID3D12Heap has %u references left.\n", (unsigned int)refcount);
 
+    memset(&architecture, 0, sizeof(architecture));
+    hr = ID3D12Device_CheckFeatureSupport(device, D3D12_FEATURE_ARCHITECTURE, &architecture, sizeof(architecture));
+    ok(hr == S_OK, "Got unexpected hr %#x.\n", hr);
+    for (i = D3D12_HEAP_TYPE_DEFAULT; i < D3D12_HEAP_TYPE_CUSTOM; ++i)
+    {
+        vkd3d_test_set_context("Test %u\n", i);
+        desc.Properties = ID3D12Device_GetCustomHeapProperties(device, 1, i);
+        ok(desc.Properties.Type == D3D12_HEAP_TYPE_CUSTOM, "Got unexpected heap type %#x.\n", desc.Properties.Type);
+
+        switch (i)
+        {
+            case D3D12_HEAP_TYPE_DEFAULT:
+                ok(desc.Properties.CPUPageProperty == D3D12_CPU_PAGE_PROPERTY_NOT_AVAILABLE,
+                        "Got unexpected CPUPageProperty %#x.\n", desc.Properties.CPUPageProperty);
+                ok(desc.Properties.MemoryPoolPreference == (architecture.UMA
+                        ? D3D12_MEMORY_POOL_L0 : D3D12_MEMORY_POOL_L1),
+                        "Got unexpected MemoryPoolPreference %#x.\n", desc.Properties.MemoryPoolPreference);
+                break;
+
+            case D3D12_HEAP_TYPE_UPLOAD:
+                ok(desc.Properties.CPUPageProperty == (architecture.CacheCoherentUMA
+                        ? D3D12_CPU_PAGE_PROPERTY_WRITE_BACK : D3D12_CPU_PAGE_PROPERTY_WRITE_COMBINE),
+                        "Got unexpected CPUPageProperty %#x.\n", desc.Properties.CPUPageProperty);
+                ok(desc.Properties.MemoryPoolPreference == D3D12_MEMORY_POOL_L0,
+                        "Got unexpected MemoryPoolPreference %#x.\n", desc.Properties.MemoryPoolPreference);
+                break;
+
+            case D3D12_HEAP_TYPE_READBACK:
+                ok(desc.Properties.CPUPageProperty == D3D12_CPU_PAGE_PROPERTY_WRITE_BACK,
+                        "Got unexpected CPUPageProperty %#x.\n", desc.Properties.CPUPageProperty);
+                ok(desc.Properties.MemoryPoolPreference == D3D12_MEMORY_POOL_L0,
+                        "Got unexpected MemoryPoolPreference %#x.\n", desc.Properties.MemoryPoolPreference);
+                break;
+
+            default:
+              ok(0, "Invalid heap type %#x.\n", i);
+              continue;
+        }
+
+        hr = ID3D12Device_CreateHeap(device, &desc, &IID_ID3D12Heap, (void **)&heap);
+        ok(hr == S_OK, "Got unexpected hr %#x.\n", hr);
+        result_desc = ID3D12Heap_GetDesc(heap);
+        check_heap_desc(&result_desc, &desc);
+        ID3D12Heap_Release(heap);
+    }
+    vkd3d_test_set_context(NULL);
+
     is_pool_L1_supported = is_memory_pool_L1_supported(device);
     desc.Properties.Type = D3D12_HEAP_TYPE_CUSTOM;
     desc.Properties.CreationNodeMask = 1;
@@ -4695,7 +4786,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;
@@ -4719,42 +4810,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},
         {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},
         {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},
         {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},
         {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},
         {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},
         {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},
 
         {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},
         {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},
         {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},
         {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},
         {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},
         {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},
         {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},
         {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},
         {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},
+
+        {DXGI_FORMAT_R32_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0, 0, 0, 0}, 0},
+        {DXGI_FORMAT_R32_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {1, 0, 0, 0}, 1},
+        {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},
+
+        {DXGI_FORMAT_R16G16_UINT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x1234, 0xabcd, 0, 0}, 0xabcd1234},
+        {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},
+        {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},
+        {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},
+        {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},
+
+        {DXGI_FORMAT_R16G16_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x1234, 0xabcd, 0, 0}, 0xabcd1234},
+        {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},
+
+        {DXGI_FORMAT_R8G8B8A8_UINT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x11, 0x22, 0x33, 0x44}, 0x44332211},
+        {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},
+        {DXGI_FORMAT_R11G11B10_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x7ff, 0x7ff, 0x3ff, 0}, 0xffffffff},
+        {DXGI_FORMAT_R11G11B10_FLOAT, { 0, BUFFER_SIZE / sizeof(uint32_t), 0, 0, D3D12_BUFFER_UAV_FLAG_NONE},
+                {0x7ff, 0, 0x3ff, 0}, 0xffc007ff},
+        {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},
+        {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},
     };
 
     memset(&desc, 0, sizeof(desc));
@@ -4784,7 +4921,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;
@@ -4810,10 +4947,16 @@ 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,
@@ -4824,7 +4967,8 @@ 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);
         check_readback_data_uint(&rb, &box, clear_value[0], 0);
@@ -4842,6 +4986,283 @@ static void test_clear_unordered_access_view(void)
 #undef BUFFER_SIZE
 }
 
+static void test_clear_unordered_access_view_image(void)
+{
+    unsigned int expected_colour, actual_colour;
+    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
+    static const 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 a specific mip level. */
+        {DXGI_FORMAT_R32_FLOAT,       2, 1, 0, 0, 1, 0, {}, {1,          0, 0, 0}, 1},
+        {DXGI_FORMAT_R32_FLOAT,       2, 1, 1, 0, 1, 0, {}, {1,          0, 0, 0}, 1},
+        {DXGI_FORMAT_R32_FLOAT,       2, 1, 0, 0, 1, 0, {}, {0x3f000000, 0, 0, 0}, 0x3f000000, true},
+        {DXGI_FORMAT_R32_FLOAT,       2, 1, 1, 0, 1, 0, {}, {0x3f000000, 0, 0, 0}, 0x3f000000, true},
+        /* Test clearing specific array layers. */
+        {DXGI_FORMAT_R32_FLOAT,       1, IMAGE_SIZE, 0, 0, IMAGE_SIZE, 0, {}, {1, 0, 0, 0}, 1},
+        {DXGI_FORMAT_R32_FLOAT,       1, IMAGE_SIZE, 0, 3, 2,          0, {}, {1, 0, 0, 0}, 1},
+        {DXGI_FORMAT_R32_FLOAT,       1, IMAGE_SIZE, 0, 0, IMAGE_SIZE, 0, {},
+                {0x3f000000, 0, 0, 0}, 0x3f000000, true},
+        {DXGI_FORMAT_R32_FLOAT,       1, IMAGE_SIZE, 0, 3, 2,          0, {},
+                {0x3f000000, 0, 0, 0}, 0x3f000000, true},
+        /* Test a single clear rect. */
+        {DXGI_FORMAT_R32_FLOAT,       1, 1, 0, 0, 1, 1, {{1, 2, IMAGE_SIZE - 4, IMAGE_SIZE - 2}},
+                {1,          0, 0, 0}, 1},
+        {DXGI_FORMAT_R32_FLOAT,       1, 1, 0, 0, 1, 1, {{1, 2, IMAGE_SIZE - 4, IMAGE_SIZE - 2}},
+                {0x3f000000, 0, 0, 0}, 0x3f000000, true},
+        /* Test multiple clear rects. */
+        {DXGI_FORMAT_R32_FLOAT,       1, 1, 0, 0, 1, 2, {{1, 2, 3, 4}, {5, 6, 7, 8}},
+                {1,          0, 0, 0}, 1},
+        {DXGI_FORMAT_R32_FLOAT,       1, 1, 0, 0, 1, 2, {{1, 2, 3, 4}, {5, 6, 7, 8}},
+                {0x3f000000, 0, 0, 0}, 0x3f000000, true},
+        /* Test uint clears with formats. */
+        {DXGI_FORMAT_R16G16_UINT,     1, 1, 0, 0, 1, 0, {}, {1,       2, 3, 4}, 0x00020001},
+        {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},
+        {DXGI_FORMAT_R16G16_FLOAT,    1, 1, 0, 0, 1, 0, {}, {1,       2, 3, 4}, 0x00020001},
+        {DXGI_FORMAT_R8G8B8A8_UINT,   1, 1, 0, 0, 1, 0, {}, {1,       2, 3, 4}, 0x04030201},
+        {DXGI_FORMAT_R8G8B8A8_UINT,   1, 1, 0, 0, 1, 0, {}, {0x123,   0, 0, 0}, 0x00000023, false, true},
+        {DXGI_FORMAT_R8G8B8A8_UNORM,  1, 1, 0, 0, 1, 0, {}, {1,       2, 3, 4}, 0x04030201},
+        {DXGI_FORMAT_R11G11B10_FLOAT, 1, 1, 0, 0, 1, 0, {}, {1,       2, 3, 4}, 0x00c01001},
+        /* 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},
+        {DXGI_FORMAT_R16G16_FLOAT,    1, 1, 0, 0, 1, 0, {},
+                {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0x3c003800, true},
+        {DXGI_FORMAT_R8G8B8A8_UNORM,  1, 1, 0, 0, 1, 0, {},
+                {0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */, 0, 0}, 0x0000ff80, true},
+        {DXGI_FORMAT_R8G8B8A8_UNORM,  1, 1, 0, 0, 1, 0, {},
+                {0, 0, 0x3f000000 /* 0.5f */, 0x3f800000 /* 1.0f */}, 0xff800000, true},
+        {DXGI_FORMAT_R11G11B10_FLOAT, 1, 1, 0, 0, 1, 0, {},
+                {0x3f000000 /* 1.0f */, 0 /* 0.0f */, 0xbf800000 /* -1.0f */, 0x3f000000 /* 1.0f */},
+                0x00000380, true},
+    };
+
+    static const 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 = IMAGE_SIZE;
+            if (uav_dimensions[d].resource_dim == D3D12_RESOURCE_DIMENSION_TEXTURE1D)
+                resource_desc.Height = 1;
+            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;
+
+            if (FAILED(hr = ID3D12Device_CreateCommittedResource(device, &heap_properties,
+                    D3D12_HEAP_FLAG_NONE, &resource_desc, D3D12_RESOURCE_STATE_UNORDERED_ACCESS,
+                    NULL, &IID_ID3D12Resource, (void **)&texture)))
+            {
+                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_colour = is_inside ? tests[i].expected : clear_value[0];
+                    actual_colour = get_readback_uint(&rb, x, y, z);
+                    success = compare_color(actual_colour, expected_colour, tests[i].is_float ? 1 : 0);
+
+                    todo_if(tests[i].is_todo && expected_colour)
+                    ok(success, "At layer %u, (%u,%u,%u), expected %#x, got %#x.\n",
+                            layer, x, y, z, expected_colour, actual_colour);
+
+                    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;
@@ -29016,7 +29437,6 @@ static void test_resource_allocation_info(void)
     ID3D12Device *device;
     unsigned int i, j;
     ULONG refcount;
-    uint64_t size;
 
     static const unsigned int alignments[] =
     {
@@ -29051,12 +29471,15 @@ static void test_resource_allocation_info(void)
         { 4,  4, 1, 1, DXGI_FORMAT_R8_UINT},
         { 8,  8, 1, 1, DXGI_FORMAT_R8G8B8A8_UNORM},
         {16, 16, 1, 1, DXGI_FORMAT_R8G8B8A8_UNORM},
-        {16, 16, 6, 1, DXGI_FORMAT_R8G8B8A8_UNORM},
+        {16, 16, 1024, 1, DXGI_FORMAT_R8G8B8A8_UNORM},
+        {256, 512, 1, 10, DXGI_FORMAT_BC1_UNORM},
+        {256, 512, 64, 1, DXGI_FORMAT_BC1_UNORM},
 
         {1024, 1024, 1, 1, DXGI_FORMAT_R8G8B8A8_UNORM},
         {1024, 1024, 1, 2, DXGI_FORMAT_R8G8B8A8_UNORM},
         {1024, 1024, 1, 3, DXGI_FORMAT_R8G8B8A8_UNORM},
         {1024, 1024, 1, 0, DXGI_FORMAT_R8G8B8A8_UNORM},
+        {260, 512, 1, 1, DXGI_FORMAT_BC1_UNORM},
     };
 
     if (!(device = create_device()))
@@ -29128,8 +29551,7 @@ static void test_resource_allocation_info(void)
         info = ID3D12Device_GetResourceAllocationInfo(device, 0, 1, &desc);
         ok(info.Alignment >= D3D12_SMALL_RESOURCE_PLACEMENT_ALIGNMENT,
                 "Got unexpected alignment %"PRIu64".\n", info.Alignment);
-        size = desc.Width * desc.Height * desc.DepthOrArraySize * format_size(desc.Format);
-        if (size <= D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT)
+        if (i < 6)
         {
             check_alignment(info.SizeInBytes, info.Alignment);
         }
@@ -32306,6 +32728,89 @@ static void test_bufinfo_instruction(void)
     destroy_test_context(&context);
 }
 
+static void test_write_buffer_immediate(void)
+{
+    D3D12_WRITEBUFFERIMMEDIATE_PARAMETER parameters[2];
+    ID3D12GraphicsCommandList2 *command_list2;
+    D3D12_WRITEBUFFERIMMEDIATE_MODE modes[2];
+    ID3D12GraphicsCommandList *command_list;
+    struct resource_readback rb;
+    struct test_context context;
+    ID3D12CommandQueue *queue;
+    ID3D12Resource *buffer;
+    ID3D12Device *device;
+    unsigned int value;
+    HRESULT hr;
+
+    static const unsigned int data_values[] = {0xdeadbeef, 0xf00baa};
+
+    if (!init_test_context(&context, NULL))
+        return;
+    device = context.device;
+    command_list = context.list;
+    queue = context.queue;
+
+    if (FAILED(hr = ID3D12GraphicsCommandList_QueryInterface(command_list,
+            &IID_ID3D12GraphicsCommandList2, (void **)&command_list2)))
+    {
+        skip("ID3D12GraphicsCommandList2 not implemented.\n");
+        destroy_test_context(&context);
+        return;
+    }
+
+    buffer = create_default_buffer(device, sizeof(data_values),
+            D3D12_RESOURCE_FLAG_NONE, D3D12_RESOURCE_STATE_COPY_DEST);
+    upload_buffer_data(buffer, 0, sizeof(data_values), data_values, queue, command_list);
+    reset_command_list(command_list, context.allocator);
+
+    parameters[0].Dest = ID3D12Resource_GetGPUVirtualAddress(buffer);
+    parameters[0].Value = 0x1020304;
+    parameters[1].Dest = parameters[0].Dest + sizeof(data_values[0]);
+    parameters[1].Value = 0xc0d0e0f;
+    ID3D12GraphicsCommandList2_WriteBufferImmediate(command_list2, ARRAY_SIZE(parameters), parameters, NULL);
+    hr = ID3D12GraphicsCommandList_Close(command_list);
+    ok(hr == S_OK, "Got unexpected hr %#x.\n", hr);
+    exec_command_list(queue, command_list);
+    wait_queue_idle(device, queue);
+    reset_command_list(command_list, context.allocator);
+
+    get_buffer_readback_with_command_list(buffer, DXGI_FORMAT_R32_UINT, &rb, queue, command_list);
+    value = get_readback_uint(&rb, 0, 0, 0);
+    todo ok(value == parameters[0].Value, "Got unexpected value %#x, expected %#x.\n", value, parameters[0].Value);
+    value = get_readback_uint(&rb, 1, 0, 0);
+    todo ok(value == parameters[1].Value, "Got unexpected value %#x, expected %#x.\n", value, parameters[1].Value);
+    release_resource_readback(&rb);
+    reset_command_list(command_list, context.allocator);
+
+    parameters[0].Value = 0x2030405;
+    parameters[1].Value = 0xb0c0d0e;
+    modes[0] = D3D12_WRITEBUFFERIMMEDIATE_MODE_MARKER_IN;
+    modes[1] = D3D12_WRITEBUFFERIMMEDIATE_MODE_MARKER_OUT;
+    ID3D12GraphicsCommandList2_WriteBufferImmediate(command_list2, ARRAY_SIZE(parameters), parameters, modes);
+    hr = ID3D12GraphicsCommandList_Close(command_list);
+    ok(hr == S_OK, "Got unexpected hr %#x.\n", hr);
+    exec_command_list(queue, command_list);
+    wait_queue_idle(device, queue);
+    reset_command_list(command_list, context.allocator);
+
+    get_buffer_readback_with_command_list(buffer, DXGI_FORMAT_R32_UINT, &rb, queue, command_list);
+    value = get_readback_uint(&rb, 0, 0, 0);
+    todo ok(value == parameters[0].Value, "Got unexpected value %#x, expected %#x.\n", value, parameters[0].Value);
+    value = get_readback_uint(&rb, 1, 0, 0);
+    todo ok(value == parameters[1].Value, "Got unexpected value %#x, expected %#x.\n", value, parameters[1].Value);
+    release_resource_readback(&rb);
+    reset_command_list(command_list, context.allocator);
+
+    modes[0] = 0x7fffffff;
+    ID3D12GraphicsCommandList2_WriteBufferImmediate(command_list2, ARRAY_SIZE(parameters), parameters, modes);
+    hr = ID3D12GraphicsCommandList_Close(command_list);
+    todo ok(hr == E_INVALIDARG, "Got unexpected hr %#x.\n", hr);
+
+    ID3D12Resource_Release(buffer);
+    ID3D12GraphicsCommandList2_Release(command_list2);
+    destroy_test_context(&context);
+}
+
 START_TEST(d3d12)
 {
     parse_args(argc, argv);
@@ -32345,7 +32850,8 @@ 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_clear_unordered_access_view_image);
     run_test(test_set_render_targets);
     run_test(test_draw_instanced);
     run_test(test_draw_indexed_instanced);
@@ -32468,4 +32974,5 @@ START_TEST(d3d12)
     run_test(test_early_depth_stencil_tests);
     run_test(test_conditional_rendering);
     run_test(test_bufinfo_instruction);
+    run_test(test_write_buffer_immediate);
 }
diff --git a/tests/d3d12_test_utils.h b/tests/d3d12_test_utils.h
index 024cf87a..5092183a 100644
--- a/tests/d3d12_test_utils.h
+++ b/tests/d3d12_test_utils.h
@@ -288,11 +288,6 @@ static unsigned int format_size(DXGI_FORMAT format)
 {
     switch (format)
     {
-        case DXGI_FORMAT_UNKNOWN:
-        case DXGI_FORMAT_A8_UNORM:
-        case DXGI_FORMAT_R8_UINT:
-        case DXGI_FORMAT_R8_UNORM:
-            return 1;
         case DXGI_FORMAT_R32G32B32A32_FLOAT:
         case DXGI_FORMAT_R32G32B32A32_UINT:
         case DXGI_FORMAT_R8G8_UNORM:
@@ -304,20 +299,26 @@ 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_UNORM:
+        case DXGI_FORMAT_R16G16_UINT:
+        case DXGI_FORMAT_R11G11B10_FLOAT:
         case DXGI_FORMAT_R8G8B8A8_TYPELESS:
         case DXGI_FORMAT_R8G8B8A8_UNORM:
         case DXGI_FORMAT_R8G8B8A8_UNORM_SRGB:
+        case DXGI_FORMAT_R8G8B8A8_UINT:
         case DXGI_FORMAT_B8G8R8A8_UNORM:
             return 4;
         case DXGI_FORMAT_R16_FLOAT:
         case DXGI_FORMAT_R16_UNORM:
         case DXGI_FORMAT_R16_UINT:
             return 2;
-        case DXGI_FORMAT_BC1_UNORM:
-        case DXGI_FORMAT_BC1_UNORM_SRGB:
-        case DXGI_FORMAT_BC4_UNORM:
-        case DXGI_FORMAT_BC4_SNORM:
-            return 8;
+        case DXGI_FORMAT_UNKNOWN:
+        case DXGI_FORMAT_A8_UNORM:
+        case DXGI_FORMAT_R8_UINT:
+        case DXGI_FORMAT_R8_UNORM:
+            return 1;
+
         case DXGI_FORMAT_BC2_UNORM:
         case DXGI_FORMAT_BC2_UNORM_SRGB:
         case DXGI_FORMAT_BC3_UNORM:
@@ -329,6 +330,12 @@ static unsigned int format_size(DXGI_FORMAT format)
         case DXGI_FORMAT_BC7_UNORM:
         case DXGI_FORMAT_BC7_UNORM_SRGB:
             return 16;
+        case DXGI_FORMAT_BC1_UNORM:
+        case DXGI_FORMAT_BC1_UNORM_SRGB:
+        case DXGI_FORMAT_BC4_UNORM:
+        case DXGI_FORMAT_BC4_SNORM:
+            return 8;
+
         default:
             trace("Unhandled format %#x.\n", format);
             return 1;
diff --git a/tests/vkd3d_api.c b/tests/vkd3d_api.c
index a9a7fc10..d4fe79a8 100644
--- a/tests/vkd3d_api.c
+++ b/tests/vkd3d_api.c
@@ -854,7 +854,9 @@ static VkDeviceMemory allocate_vulkan_image_memory(ID3D12Device *device,
 static void test_external_resource_map(void)
 {
     struct vkd3d_image_resource_create_info resource_create_info;
+    D3D12_HEAP_PROPERTIES heap_properties;
     D3D12_GPU_VIRTUAL_ADDRESS gpu_address;
+    D3D12_HEAP_FLAGS heap_flags;
     ID3D12Resource *vk_resource;
     VkDeviceMemory vk_memory;
     ID3D12Device *device;
@@ -898,6 +900,16 @@ static void test_external_resource_map(void)
     gpu_address = ID3D12Resource_GetGPUVirtualAddress(vk_resource);
     ok(!gpu_address, "Got unexpected GPU virtual address %#"PRIx64".\n", gpu_address);
 
+    hr = ID3D12Resource_GetHeapProperties(vk_resource, &heap_properties, &heap_flags);
+    ok(hr == S_OK, "Got unexpected hr %#x.\n", hr);
+    ok(heap_properties.Type == D3D12_HEAP_TYPE_DEFAULT, "Got unexpected heap type %#x.\n", heap_properties.Type);
+    ok(heap_properties.CPUPageProperty == D3D12_CPU_PAGE_PROPERTY_UNKNOWN,
+            "Got unexpected CPU page property %#x.\n", heap_properties.CPUPageProperty);
+    ok(heap_properties.MemoryPoolPreference == D3D12_MEMORY_POOL_UNKNOWN,
+            "Got unexpected memory pool preference %#x.\n", heap_properties.MemoryPoolPreference);
+    ok(!!heap_properties.CreationNodeMask, "Got unexpected node mask %#x.\n", heap_properties.CreationNodeMask);
+    ok(!!heap_properties.VisibleNodeMask, "Got unexpected node mask %#x.\n", heap_properties.VisibleNodeMask);
+
     ID3D12Resource_Release(vk_resource);
     vk_device = vkd3d_get_vk_device(device);
     vkDestroyImage(vk_device, vk_image, NULL);